Integrate LLVM at llvm/llvm-project@c329cc59d91b

third_party/llvm/workspace.bzl

@@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive")
 
 def repo(name):
     """Imports LLVM."""
-    LLVM_COMMIT = "7ccd92e5e6e5c622b2b571d396fff9016241a8f1"
-    LLVM_SHA256 = "0365ff939e7ba4876437f45318ac4140e7b278fd8b2c5dc1e78c4dd04104c831"
+    LLVM_COMMIT = "c329cc59d91b3f9774383e5e933d812539f17d1e"
+    LLVM_SHA256 = "8c4e9bc4e6f99d922ecd2eb9946fe92975704bb5187d609eec33e8087fbd5d50"
 
     tf_http_archive(
         name = name,

third_party/shardy/workspace.bzl

@@ -3,8 +3,8 @@
 load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
 
 def repo():
-    SHARDY_COMMIT = "7991023acf075463d524ee94491f634514c5aa32"
-    SHARDY_SHA256 = "c007b67fd21e484848d9ec3b334abea31a895b71d998ba11a3a0b6e60826dd2a"
+    SHARDY_COMMIT = "71aa6159555e8da42911a0d60cf9225d72d42980"
+    SHARDY_SHA256 = "055f595eff9eb09cf39bcac3696d1b14f287c1dd6262c221ecf4f1c7f1b5128c"
 
     tf_http_archive(
         name = "shardy",

third_party/triton/common/llvm_cl893905698.patch

@@ -0,0 +1,146 @@
+--- a/include/triton/Dialect/TritonNvidiaGPU/IR/Dialect.h
++++ b/include/triton/Dialect/TritonNvidiaGPU/IR/Dialect.h
+@@ -70,6 +70,7 @@
+
+ struct TensorMemory : public SideEffects::Resource::Base<TensorMemory> {
+   StringRef getName() const final { return "<TensorMemory>"; }
++  Resource* getParent() const override { return nullptr; }
+ };
+
+ struct TMemAllocation {
+
+--- a/lib/Dialect/Triton/IR/Ops.cpp
++++ b/lib/Dialect/Triton/IR/Ops.cpp
+@@ -214,11 +214,10 @@
+ }
+
+ //-- DotOp --
+-LogicalResult
+-DotOp::inferReturnTypes(MLIRContext *context, std::optional<Location> location,
+-                        ValueRange operands, DictionaryAttr attributes,
+-                        OpaqueProperties properties, RegionRange regions,
+-                        SmallVectorImpl<Type> &inferredReturnTypes) {
++LogicalResult DotOp::inferReturnTypes(
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   // type is the same as the accumulator
+   auto accTy = cast<RankedTensorType>(operands[2].getType());
+   inferredReturnTypes.push_back(accTy);
+@@ -481,11 +480,10 @@
+   return success();
+ }
+
+-LogicalResult
+-ReduceOp::inferReturnTypes(MLIRContext *context, std::optional<Location> loc,
+-                           ValueRange operands, DictionaryAttr attributes,
+-                           OpaqueProperties properties, RegionRange regions,
+-                           SmallVectorImpl<Type> &inferredReturnTypes) {
++LogicalResult ReduceOp::inferReturnTypes(
++    MLIRContext* context, std::optional<Location> loc, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   Properties *prop = properties.as<Properties *>();
+   int axis = prop->axis.getInt();
+   for (auto arg : operands) {
+@@ -636,11 +634,10 @@
+   ScanOp::build(builder, state, inferredReturnTypes, operands, axis, reverse);
+ }
+
+-LogicalResult
+-ScanOp::inferReturnTypes(MLIRContext *context, std::optional<Location> location,
+-                         ValueRange operands, DictionaryAttr attributes,
+-                         OpaqueProperties properties, RegionRange regions,
+-                         SmallVectorImpl<Type> &inferredReturnTypes) {
++LogicalResult ScanOp::inferReturnTypes(
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   for (auto arg : operands)
+     inferredReturnTypes.push_back(arg.getType());
+   return success();
+@@ -737,9 +734,9 @@
+ }
+
+ LogicalResult UnsplatOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> location, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   auto dstTy = cast<RankedTensorType>(operands[0].getType()).getElementType();
+   inferredReturnTypes.push_back(dstTy);
+   return success();
+@@ -747,9 +744,9 @@
+
+ //-- ExpandDimsOp --
+ LogicalResult ExpandDimsOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> loc, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> loc, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   // infer shape
+   auto arg = operands[0];
+   auto argTy = cast<RankedTensorType>(arg.getType());
+@@ -1344,9 +1341,9 @@
+ }
+
+ LogicalResult GatherOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> location, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   GatherOpAdaptor adaptor(operands, attributes, properties, regions);
+   auto indicesType = cast<RankedTensorType>(adaptor.getIndices().getType());
+   auto srcType = cast<RankedTensorType>(adaptor.getSrc().getType());
+
+--- a/lib/Dialect/TritonNvidiaGPU/IR/Ops.cpp
++++ b/lib/Dialect/TritonNvidiaGPU/IR/Ops.cpp
+@@ -49,9 +49,9 @@
+
+ // -- WarpGroupDotOp --
+ LogicalResult WarpGroupDotOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> location, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   // type is the same as the accumulator
+   auto accTy = cast<RankedTensorType>(operands[2].getType());
+   inferredReturnTypes.push_back(accTy);
+@@ -156,9 +156,9 @@
+
+ // -- WarpGroupDotWaitOp --
+ LogicalResult WarpGroupDotWaitOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> location, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   for (Value operand : operands)
+     inferredReturnTypes.push_back(operand.getType());
+   return success();
+
+--- a/third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp
++++ b/third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp
+@@ -851,9 +851,9 @@
+ // prefetch. Values are the byte offsets added to the base pointer for each
+ // prefetch instruction.
+ LogicalResult TDMPrefetchOp::inferReturnTypes(
+-    MLIRContext *context, std::optional<Location> location, ValueRange operands,
+-    DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions,
+-    SmallVectorImpl<Type> &inferredReturnTypes) {
++    MLIRContext* context, std::optional<Location> location, ValueRange operands,
++    DictionaryAttr attributes, mlir::PropertyRef properties,
++    RegionRange regions, SmallVectorImpl<Type>& inferredReturnTypes) {
+   TDMPrefetchOp::Adaptor ad(operands, attributes, properties, regions);
+
+   // If returnOffsets is not set the op will not return any results
+

third_party/triton/common/series.bzl

@@ -29,5 +29,6 @@ common_patch_list = [
     "//third_party/triton:common/wgmma_pipeline_fix.patch",
     "//third_party/triton:common/nvdisasm_bin_path.patch",
     "//third_party/triton:common/llvm_cl887809531.patch",
+    "//third_party/triton:common/llvm_cl893905698.patch",
     # Add new patches just above this line
 ]

xla/backends/gpu/codegen/emitters/ir/xla_gpu_ops.cc

@@ -139,7 +139,7 @@ SmallVector<Type, 2> inferReductionInitTypes(TypeRange input_types) {
 
 LogicalResult ReduceOp::inferReturnTypes(
     MLIRContext* context, std::optional<Location> location, ValueRange operands,
-    mlir::DictionaryAttr attributes, mlir::OpaqueProperties properties,
+    mlir::DictionaryAttr attributes, mlir::PropertyRef properties,
     mlir::RegionRange regions,
     mlir::SmallVectorImpl<Type>& inferredReturnTypes) {
   ReduceOp::Adaptor adaptor(operands, attributes, properties, regions);