Skip to content

Commit

Permalink
Integrate LLVM at llvm/llvm-project@ac2a2816e3fe
Browse files Browse the repository at this point in the history
Updates LLVM usage to match
[ac2a2816e3fe](llvm/llvm-project@ac2a2816e3fe)

PiperOrigin-RevId: 680981749
  • Loading branch information
Google-ML-Automation committed Oct 1, 2024
1 parent ed705dc commit 877928f
Show file tree
Hide file tree
Showing 14 changed files with 3,585 additions and 44 deletions.
407 changes: 407 additions & 0 deletions third_party/llvm/generated.patch

Large diffs are not rendered by default.

4 changes: 2 additions & 2 deletions third_party/llvm/workspace.bzl
Original file line number Diff line number Diff line change
Expand Up @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive")

def repo(name):
"""Imports LLVM."""
LLVM_COMMIT = "29b92d07746fac26cd64c914bc9c5c3833974f6d"
LLVM_SHA256 = "3e8e93e3749454af4b64f7f34b792a4748b62fc533bca1703d33b2b04e34eb70"
LLVM_COMMIT = "ac2a2816e3fe934998e5f950c9426fca0796929d"
LLVM_SHA256 = "bfcec3c9a582ab71c1e4a212e5a7cfe2c9a8164331a0b728da5c8a99d81c46de"

tf_http_archive(
name = name,
Expand Down
2,419 changes: 2,391 additions & 28 deletions third_party/shardy/temporary.patch

Large diffs are not rendered by default.

285 changes: 285 additions & 0 deletions third_party/stablehlo/temporary.patch
Original file line number Diff line number Diff line change
@@ -1,3 +1,196 @@
diff --ruN a/stablehlo/examples/c++/ExampleAdd.cpp b/stablehlo/examples/c++/ExampleAdd.cpp
--- stablehlo/examples/c++/ExampleAdd.cpp
+++ stablehlo/examples/c++/ExampleAdd.cpp
@@ -18,7 +18,7 @@
#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/LogicalResult.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Block.h"
#include "mlir/IR/Builders.h"
@@ -43,7 +43,7 @@
mlir::ModuleOp::create(mlir::UnknownLoc::get(&context));
module->getContext()->loadDialect<mlir::func::FuncDialect>();
module->getContext()->loadDialect<mlir::stablehlo::StablehloDialect>();
- module->getContext()->loadDialect<mlir::quant::QuantizationDialect>();
+ module->getContext()->loadDialect<mlir::quant::QuantDialect>();
module->setName("test_module");

/** create function **/
diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp b/stablehlo/stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp
--- stablehlo/stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp
+++ stablehlo/stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp
@@ -17,7 +17,7 @@
#include <utility>

#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/Dialect/Tosa/IR/TosaOps.h"
#include "mlir/Dialect/Tosa/Utils/ConversionUtils.h"
#include "mlir/Dialect/Tosa/Utils/QuantUtils.h"
diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp b/stablehlo/stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp
--- stablehlo/stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp
+++ stablehlo/stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp
@@ -18,7 +18,7 @@
#include <utility>

#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/Dialect/Tosa/IR/TosaOps.h"
#include "mlir/Dialect/Tosa/Utils/ConversionUtils.h"
#include "mlir/Dialect/Tosa/Utils/QuantUtils.h"
diff --ruN a/stablehlo/stablehlo/dialect/Base.cpp b/stablehlo/stablehlo/dialect/Base.cpp
--- stablehlo/stablehlo/dialect/Base.cpp
+++ stablehlo/stablehlo/dialect/Base.cpp
@@ -31,7 +31,7 @@
#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/Dialect/Shape/IR/Shape.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
diff --ruN a/stablehlo/stablehlo/dialect/ChloOps.h b/stablehlo/stablehlo/dialect/ChloOps.h
--- stablehlo/stablehlo/dialect/ChloOps.h
+++ stablehlo/stablehlo/dialect/ChloOps.h
@@ -20,7 +20,7 @@
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/StringRef.h"
#include "mlir/Bytecode/BytecodeOpInterface.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinTypes.h"
diff --ruN a/stablehlo/stablehlo/dialect/Register.cpp b/stablehlo/stablehlo/dialect/Register.cpp
--- stablehlo/stablehlo/dialect/Register.cpp
+++ stablehlo/stablehlo/dialect/Register.cpp
@@ -17,7 +17,7 @@
#include "stablehlo/dialect/Register.h"

#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/Dialect/SparseTensor/IR/SparseTensor.h"
#include "mlir/IR/DialectRegistry.h"
#include "stablehlo/dialect/ChloOps.h"
@@ -30,7 +30,7 @@
void registerAllDialects(mlir::DialectRegistry &registry) {
// clang-format off
registry.insert<mlir::func::FuncDialect,
- mlir::quant::QuantizationDialect,
+ mlir::quant::QuantDialect,
mlir::sparse_tensor::SparseTensorDialect>();
registry.insert<mlir::chlo::ChloDialect,
mlir::stablehlo::StablehloDialect,
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp
--- stablehlo/stablehlo/dialect/StablehloOps.cpp
+++ stablehlo/stablehlo/dialect/StablehloOps.cpp
@@ -52,7 +52,7 @@
#include "llvm/Support/Regex.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Complex/IR/Complex.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/Dialect/Shape/IR/Shape.h"
#include "mlir/Dialect/SparseTensor/IR/SparseTensor.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h"
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.h b/stablehlo/stablehlo/dialect/StablehloOps.h
--- stablehlo/stablehlo/dialect/StablehloOps.h
+++ stablehlo/stablehlo/dialect/StablehloOps.h
@@ -21,7 +21,7 @@
#include <optional>

#include "llvm/ADT/StringRef.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/Dialect/Shape/IR/Shape.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Builders.h"
diff --ruN a/stablehlo/stablehlo/dialect/TypeInference.cpp b/stablehlo/stablehlo/dialect/TypeInference.cpp
--- stablehlo/stablehlo/dialect/TypeInference.cpp
+++ stablehlo/stablehlo/dialect/TypeInference.cpp
@@ -52,7 +52,7 @@
#include "llvm/Support/Regex.h"
#include "llvm/Support/raw_ostream.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
diff --ruN a/stablehlo/stablehlo/dialect/VhloTypes.cpp b/stablehlo/stablehlo/dialect/VhloTypes.cpp
--- stablehlo/stablehlo/dialect/VhloTypes.cpp
+++ stablehlo/stablehlo/dialect/VhloTypes.cpp
@@ -20,7 +20,7 @@
#include "llvm/ADT/SmallVectorExtras.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/TypeSwitch.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/Dialect/Shape/IR/Shape.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinTypes.h"
diff --ruN a/stablehlo/stablehlo/reference/Api.cpp b/stablehlo/stablehlo/reference/Api.cpp
--- stablehlo/stablehlo/reference/Api.cpp
+++ stablehlo/stablehlo/reference/Api.cpp
@@ -31,7 +31,7 @@
#include "llvm/Support/Path.h"
#include "llvm/Support/SourceMgr.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/BuiltinTypeInterfaces.h"
diff --ruN a/stablehlo/stablehlo/tests/CheckOps.h b/stablehlo/stablehlo/tests/CheckOps.h
--- stablehlo/stablehlo/tests/CheckOps.h
+++ stablehlo/stablehlo/tests/CheckOps.h
@@ -17,7 +17,7 @@
#define STABLEHLO_DIALECT_CHECKOPS_H_

#include "mlir/Bytecode/BytecodeOpInterface.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Dialect.h"
diff --ruN a/stablehlo/stablehlo/tests/ops_stablehlo_quantized.mlir b/stablehlo/stablehlo/tests/ops_stablehlo_quantized.mlir
--- stablehlo/stablehlo/tests/ops_stablehlo_quantized.mlir
+++ stablehlo/stablehlo/tests/ops_stablehlo_quantized.mlir
@@ -1338,24 +1338,24 @@

// -----

+// expected-error@+1 {{scale out of expressed type range}}
func.func @quantized_element_type_c6(%arg0: tensor<1x2x!quant.uniform<i4:f16, 10.550400e+04>>) {
- // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}} 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values, but got 'tensor<1x2x!quant.uniform<i4:f16, 1.055040e+05>>'}}
%0 = stablehlo.add %arg0, %arg0 : tensor<1x2x!quant.uniform<i4:f16, 10.550400e+04>>
func.return
}

// -----

+// expected-error@+1 {{scale out of expressed type range}}
func.func @quantized_element_type_c6(%arg0: tensor<1x2x!quant.uniform<i4:f16, 4.960464e-08>>) {
- // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}} 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values, but got 'tensor<1x2x!quant.uniform<i4:f16, 4.9604639999999998E-8>>'}}
%0 = stablehlo.add %arg0, %arg0 : tensor<1x2x!quant.uniform<i4:f16, 4.960464e-08>>
func.return
}

// -----

+// expected-error@+1 {{illegal quantized dimension: -1}}
func.func @quantized_element_type_c11(%arg0: tensor<1x5x2x!quant.uniform<i8<-128:127>:f32:-1, {0.1:-30, 0.1:-30}>>) {
- // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}} 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values, but got 'tensor<1x5x2x!quant.uniform<i8:f32:-1, {1.000000e-01:-30,1.000000e-01:-30}>>'}}
%0 = stablehlo.add %arg0, %arg0 : tensor<1x5x2x!quant.uniform<i8<-128:127>:f32:-1, {0.1:-30, 0.1:-30}>>
func.return
}
diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_create_compatibility_expander.mlir b/stablehlo/stablehlo/tests/transforms/stablehlo_create_compatibility_expander.mlir
--- stablehlo/stablehlo/tests/transforms/stablehlo_create_compatibility_expander.mlir
+++ stablehlo/stablehlo/tests/transforms/stablehlo_create_compatibility_expander.mlir
Expand Down Expand Up @@ -427,6 +620,51 @@ diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_create_compatibility
+ }) : (tensor<?x2x4x7x9xi32>, tensor<4x?x5x2xi32>, tensor<4x?x5x8xi32>) -> tensor<?x2x4x7x9xi32>
+ func.return %0 : tensor<?x2x4x7x9xi32>
+}
diff --ruN a/stablehlo/stablehlo/tools/StablehloTranslateMain.cpp b/stablehlo/stablehlo/tools/StablehloTranslateMain.cpp
--- stablehlo/stablehlo/tools/StablehloTranslateMain.cpp
+++ stablehlo/stablehlo/tools/StablehloTranslateMain.cpp
@@ -24,7 +24,7 @@
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/LogicalResult.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/DialectRegistry.h"
@@ -237,7 +237,7 @@
},
[](DialectRegistry &registry) {
registry.insert<func::FuncDialect>();
- registry.insert<quant::QuantizationDialect>();
+ registry.insert<quant::QuantDialect>();
registry.insert<stablehlo::check::CheckDialect>();
registry.insert<stablehlo::interpreter::InterpreterDialect>();
registry.insert<stablehlo::StablehloDialect>();
diff --ruN a/stablehlo/stablehlo/transforms/Passes.h b/stablehlo/stablehlo/transforms/Passes.h
--- stablehlo/stablehlo/transforms/Passes.h
+++ stablehlo/stablehlo/transforms/Passes.h
@@ -19,7 +19,7 @@
#include <memory>

#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
#include "mlir/Dialect/Shape/IR/Shape.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/Pass/Pass.h"
diff --ruN a/stablehlo/stablehlo/transforms/Passes.td b/stablehlo/stablehlo/transforms/Passes.td
--- stablehlo/stablehlo/transforms/Passes.td
+++ stablehlo/stablehlo/transforms/Passes.td
@@ -68,7 +68,7 @@
let summary = "Legalize VHLO to StableHLO.";
let dependentDialects = [
"mlir::func::FuncDialect",
- "mlir::quant::QuantizationDialect",
+ "mlir::quant::QuantDialect",
"mlir::shape::ShapeDialect",
"mlir::stablehlo::StablehloDialect",
];
diff --ruN a/stablehlo/stablehlo/transforms/StablehloCreateCompatibilityExpander.cpp b/stablehlo/stablehlo/transforms/StablehloCreateCompatibilityExpander.cpp
--- stablehlo/stablehlo/transforms/StablehloCreateCompatibilityExpander.cpp
+++ stablehlo/stablehlo/transforms/StablehloCreateCompatibilityExpander.cpp
Expand Down Expand Up @@ -614,4 +852,51 @@ diff --ruN a/stablehlo/stablehlo/transforms/StablehloCreateCompatibilityExpander

newScatterOp.getUpdateComputation().takeBody(op.getUpdateComputation());
rewriter.replaceOp(op, newScatterOp.getResults());
diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
--- stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
+++ stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
@@ -15,7 +15,7 @@

#include "llvm/ADT/SmallVector.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/Operation.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Transforms/DialectConversion.h" // Include for TypeConverter
diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeQuantToMath.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeQuantToMath.cpp
--- stablehlo/stablehlo/transforms/StablehloLegalizeQuantToMath.cpp
+++ stablehlo/stablehlo/transforms/StablehloLegalizeQuantToMath.cpp
@@ -24,8 +24,8 @@
#include "llvm/ADT/SmallVector.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/Func/Transforms/FuncConversions.h"
-#include "mlir/Dialect/Quant/QuantOps.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/Quant.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinTypeInterfaces.h"
@@ -1331,7 +1331,7 @@
populateReturnOpTypeConversionPattern(patterns, converter);

ConversionTarget target(*op->getContext());
- target.addIllegalDialect<quant::QuantizationDialect>();
+ target.addIllegalDialect<quant::QuantDialect>();
auto isLegal = [&converter](Operation *op) {
return converter.isLegal(op);
};
diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeQuantizedOpToQDQ.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeQuantizedOpToQDQ.cpp
--- stablehlo/stablehlo/transforms/StablehloLegalizeQuantizedOpToQDQ.cpp
+++ stablehlo/stablehlo/transforms/StablehloLegalizeQuantizedOpToQDQ.cpp
@@ -17,7 +17,7 @@

#include "llvm/ADT/STLExtras.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/Quant/QuantTypes.h"
+#include "mlir/Dialect/Quant/IR/QuantTypes.h"
#include "mlir/IR/BuiltinTypeInterfaces.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/IR/TypeRange.h"

74 changes: 74 additions & 0 deletions third_party/triton/llvm_integration/cl680875920.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,74 @@

--- a/lib/Conversion/TritonGPUToLLVM/ControlFlowOpToLLVM.cpp 2024-03-19 09:23:43.000000000 -0700
+++ b/lib/Conversion/TritonGPUToLLVM/ControlFlowOpToLLVM.cpp 2024-10-01 02:58:18.000000000 -0700
@@ -104,9 +104,26 @@
this->getTypeConverter()->packFunctionResults(resultTypes)))
return nullptr;
}
+ // Add LLVMOp Bundle Attrs
+ // https://github.com/llvm/llvm-project/blob/main/flang/lib/Optimizer/CodeGen/CodeGen.cpp#L113-L131
+ llvm::SmallVector<mlir::NamedAttribute> newAttrs;
+ newAttrs.reserve(callOp->getAttrs().size() + 2);
+
+ for (mlir::NamedAttribute attr : callOp->getAttrs()) {
+ if (attr.getName() != "operandSegmentSizes")
+ newAttrs.push_back(attr);
+ }
+
+ newAttrs.push_back(rewriter.getNamedAttr(
+ "operandSegmentSizes",
+ rewriter.getDenseI32ArrayAttr(
+ {static_cast<int>(promotedOperands.size()), 0})));
+ newAttrs.push_back(rewriter.getNamedAttr(
+ "op_bundle_sizes", rewriter.getDenseI32ArrayAttr({})));
+
auto newCallOp = rewriter.create<LLVM::CallOp>(
callOp.getLoc(), packedResult ? TypeRange(packedResult) : TypeRange(),
- promotedOperands, callOp->getAttrs());
+ promotedOperands, newAttrs);
return newCallOp;
}


--- a/third_party/amd/lib/TritonAMDGPUToLLVM/BuiltinFuncToLLVM.cpp 2024-09-25 10:13:59.000000000 -0700
+++ b/third_party/amd/lib/TritonAMDGPUToLLVM/BuiltinFuncToLLVM.cpp 2024-09-30 23:51:44.000000000 -0700
@@ -190,7 +190,8 @@
auto name = StringAttr::get(callOp.getContext(), "llvm.amdgcn.rcp.f32");
LLVM::FastmathFlagsAttr defaultFlags{};
auto rcpOp = rewriter.create<LLVM::CallIntrinsicOp>(
- loc, returnType, name, operands[1], defaultFlags);
+ loc, returnType, name, operands[1], defaultFlags,
+ ::llvm::ArrayRef<::mlir::ValueRange>{} /*op_bundle_operands*/);

replacementOp = rewriter.create<LLVM::FMulOp>(
loc, returnType, operands[0], rcpOp->getResult(0), defaultFlags);


--- a/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/WMMA.cpp 2024-08-20 03:28:55.000000000 -0700
+++ b/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/WMMA.cpp 2024-09-30 23:51:44.000000000 -0700
@@ -219,7 +219,8 @@
}
auto wmmaIntrinsic = rewriter.create<mlir::LLVM::CallIntrinsicOp>(
loc, TypeRange{valC.getType()}, StringAttr::get(loc.getContext(), name),
- operands, defaultFlags);
+ operands, defaultFlags,
+ ::llvm::ArrayRef<::mlir::ValueRange>{} /*op_bundle_operands*/);

return wmmaIntrinsic.getResult(0);
}


--- a/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.cpp 2024-09-16 13:44:40.000000000 -0700
+++ b/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.cpp 2024-09-30 23:51:44.000000000 -0700
@@ -72,7 +72,10 @@
auto stringAttr = rewriter.getStringAttr("llvm.amdgcn.ballot");
SmallVector<Value> operands = {cmp};
Value asmResult =
- rewriter.create<LLVM::CallIntrinsicOp>(loc, type, stringAttr, operands)
+ rewriter
+ .create<LLVM::CallIntrinsicOp>(
+ loc, type, stringAttr, operands, ::mlir::LLVM::FastmathFlags{},
+ ::llvm::ArrayRef<::mlir::ValueRange>{} /*op_bundle_operands*/)
->getResult(0);
return asmResult;
}
1 change: 1 addition & 0 deletions third_party/triton/llvm_integration/series.bzl
Original file line number Diff line number Diff line change
Expand Up @@ -8,5 +8,6 @@ LLVM nor MLIR integrator, please do not add any patches to this list.
"""

llvm_patch_list = [
"//third_party/triton:llvm_integration/cl680875920.patch",
# Add new patches just above this line
]
Loading

0 comments on commit 877928f

Please sign in to comment.