[Backend] Bump to llvm/llvm-project@8957e64a20fc (triton-lang#7138)

antiagainst · paultrojahnamd · commit 94d755908719 · 2025-07-10T15:20:49.000+02:00
This picks up a bug fix for AMDGPU scaled mfma: * llvm/llvm-project#142493 Along the way we need to fix API changes: * `getBackwardSlice` now returns `LogicalResult`; adding `(void)` to ignore for now. * `LLVM::InlineAsmOp` now has a TailCallKind parameter. * A few conversion ops in AMD backend changed return types and parameters.
diff --git a/cmake/llvm-hash.txt b/cmake/llvm-hash.txt
@@ -1 +1 @@
-e12cbd8339b89563059c2bb2a312579b652560d0
+8957e64a20fc7f4277565c6cfe3e555c119783ce
diff --git a/lib/Analysis/Utility.cpp b/lib/Analysis/Utility.cpp
@@ -973,7 +973,7 @@ SetVector<Operation *> multiRootGetSlice(Operation *op,
     BackwardSliceOptions opt;
     opt.omitBlockArguments = true;
     opt.filter = backwardFilter;
-    getBackwardSlice(currentOp, &backwardSlice, opt);
+    (void)getBackwardSlice(currentOp, &backwardSlice, opt);
     slice.insert(backwardSlice.begin(), backwardSlice.end());
 
     // Compute and insert the forwardSlice starting from currentOp.
diff --git a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp
@@ -298,7 +298,7 @@ struct ElementwiseInlineAsmOpConversion
                 /*asm_string=*/op.getAsmString(),
                 /*constraints=*/op.getConstraints(),
                 /*has_side_effects=*/!op.getPure(),
-                /*is_align_stack=*/false,
+                /*is_align_stack=*/false, LLVM::TailCallKind::None,
                 /*asm_dialect=*/
                 LLVM::AsmDialectAttr::get(rewriter.getContext(),
                                           LLVM::AsmDialect::AD_ATT),
diff --git a/lib/Dialect/TritonGPU/Transforms/AccelerateMatmul.cpp b/lib/Dialect/TritonGPU/Transforms/AccelerateMatmul.cpp
@@ -235,7 +235,7 @@ static int computeOrigBitWidth(Value x) {
   mlir::BackwardSliceOptions opt;
   opt.omitBlockArguments = true;
   opt.filter = bwdFilter;
-  getBackwardSlice(x, &slice, opt);
+  (void)getBackwardSlice(x, &slice, opt);
 
   // TODO: This heuristic may be a bit too coarse and may need improving
   // If the chain contains a fp4 to fp16/bf16 conversion, then the original
diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeliner/ScheduleLoops.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeliner/ScheduleLoops.cpp
@@ -209,7 +209,7 @@ CoarseSchedule::Cluster schedulePrologueAndEpilogue(scf::ForOp forOp,
       BackwardSliceOptions opt;
       opt.omitBlockArguments = true;
       opt.omitUsesFromAbove = false;
-      getBackwardSlice((Operation *)op, &backwardSlice, opt);
+      (void)getBackwardSlice((Operation *)op, &backwardSlice, opt);
 
       for (auto op : backwardSlice) {
         if (auto ifOp = dyn_cast<scf::IfOp>(op)) {
diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeliner/TC05MMAPipeline.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeliner/TC05MMAPipeline.cpp
@@ -632,7 +632,7 @@ FailureOr<scf::ForOp> preProcessLoopForTC05MMAPipelining(scf::ForOp forOp,
     SetVector<Operation *> backwardSlice;
     BackwardSliceOptions opt;
     opt.omitBlockArguments = true;
-    getBackwardSlice(mmaOp, &backwardSlice, opt);
+    (void)getBackwardSlice(mmaOp, &backwardSlice, opt);
     if (llvm::any_of(backwardSlice, [&](Operation *op) {
           return op->hasAttr(kPipelineStageAttrName);
         })) {
diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeliner/WGMMAPipeline.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeliner/WGMMAPipeline.cpp
@@ -207,7 +207,7 @@ static void threadValuesThroughWait(ttng::WarpGroupDotWaitOp wait,
       return op->getBlock() == wait->getBlock();
     };
     SetVector<Operation *> slice;
-    getBackwardSlice(v, &slice, options);
+    (void)getBackwardSlice(v, &slice, options);
   }
 
   for (ttng::WarpGroupDotOp dot : asyncDots) {
diff --git a/lib/Dialect/TritonGPU/Transforms/WSCodePartition.cpp b/lib/Dialect/TritonGPU/Transforms/WSCodePartition.cpp
@@ -925,7 +925,7 @@ void reorderProducerOps(SmallVector<Channel *> &channels) {
       BackwardSliceOptions opt;
       opt.omitBlockArguments = true;
       SetVector<Operation *> backwardSlice;
-      getBackwardSlice(channel->getSrcOp(), &backwardSlice, opt);
+      (void)getBackwardSlice(channel->getSrcOp(), &backwardSlice, opt);
       for (auto &op : backwardSlice) {
         if (op->getBlock() == block)
           op->moveBefore(channel->getSrcOp());
diff --git a/lib/Dialect/TritonGPU/Transforms/WSTaskPartition.cpp b/lib/Dialect/TritonGPU/Transforms/WSTaskPartition.cpp
@@ -97,8 +97,8 @@ void doPartition(triton::FuncOp &funcOp, unsigned numConsumerGroups) {
     if (!dotOp)
       continue;
     SetVector<Operation *> backwardSlice;
-    getBackwardSlice(dotOp.getA(), &backwardSlice, opt);
-    getBackwardSlice(dotOp.getB(), &backwardSlice, opt);
+    (void)getBackwardSlice(dotOp.getA(), &backwardSlice, opt);
+    (void)getBackwardSlice(dotOp.getB(), &backwardSlice, opt);
     for (auto depOp : backwardSlice) {
       if (isa<ExperimentalDescriptorLoadOp>(depOp)) {
         producerOps.insert(depOp);
diff --git a/lib/Dialect/TritonNvidiaGPU/Transforms/KeepAccInTMem.cpp b/lib/Dialect/TritonNvidiaGPU/Transforms/KeepAccInTMem.cpp
@@ -50,7 +50,7 @@ class TMEMToGlobal : public OpRewritePattern<triton::StoreOp> {
     mlir::BackwardSliceOptions opt;
     opt.omitBlockArguments = true;
     opt.filter = bwdFilter;
-    getBackwardSlice(data, &slice, opt);
+    (void)getBackwardSlice(data, &slice, opt);
     Attribute encoding;
     for (auto op : slice) {
       if (auto tmemLoad = dyn_cast<ttng::TMEMLoadOp>(op)) {
diff --git a/python/test/unit/language/test_core.py b/python/test/unit/language/test_core.py
@@ -6686,7 +6686,7 @@ def test_tl_range_num_stages(device):
             if capability[0] >= 8:
                 ptx = pgm.asm['ptx']
                 # check that the loop got pipelined with the right number of stages.
-                assert 'cp.async.wait_group 6' in ptx
+                assert 'cp.async.wait_group \t6' in ptx
 
 
 def test_tl_range_fuse():
diff --git a/test/Triton/reproducer.mlir b/test/Triton/reproducer.mlir
@@ -17,4 +17,4 @@ module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-
 #-}
 
 // CHECK: Pass Manager with
-// CHECK-NEXT: convert-triton-gpu-to-llvm
+// CHECK: convert-triton-gpu-to-llvm
diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/GCNAsmFormat.cpp b/third_party/amd/lib/TritonAMDGPUToLLVM/GCNAsmFormat.cpp
@@ -82,6 +82,7 @@ mlir::Value GCNBuilder::launch(RewriterBase &rewriter, Location loc, Type resTy,
       getConstraints(),             // constraints
       hasSideEffect,                // has_side_effects
       isAlignStack,                 // is_align_stack
+      LLVM::TailCallKind::None,
       LLVM::AsmDialectAttr::get(ctx,
                                 LLVM::AsmDialect::AD_ATT), // asm_dialect
       ArrayAttr::get(ctx, attrs)                           // operand_attrs
diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp b/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp
@@ -91,7 +91,7 @@ bool isChainDotTail(tt::DotOpInterface dotOp) {
   Operation *opA = dotOp.getA().getDefiningOp();
   if (!opA)
     return false;
-  getBackwardSlice(opA, &bwdSlices, bwdOpt);
+  (void)getBackwardSlice(opA, &bwdSlices, bwdOpt);
   if (llvm::find_if(bwdSlices, [](Operation *op) {
         return isa<tt::DotOpInterface>(op);
       }) != bwdSlices.end())
diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/BlockPingpong.cpp b/third_party/amd/lib/TritonAMDGPUTransforms/BlockPingpong.cpp
@@ -114,7 +114,7 @@ void Pingponger::moveOpAndPredecessorsUpSameBlock(Operation *op) {
       return op->getBlock() == checkedOp->getBlock() &&
              checkedOp->isBeforeInBlock(op);
     };
-    getBackwardSlice(op, &backwardSlice, opt);
+    (void)getBackwardSlice(op, &backwardSlice, opt);
     for (auto predOp : backwardSlice)
       appendOp(predOp);
     appendOp(op);
diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/ReorderInstructions.cpp b/third_party/amd/lib/TritonAMDGPUTransforms/ReorderInstructions.cpp
@@ -214,7 +214,7 @@ static void scheduleGlobalLoadLocalStore(Operation *parentOp) {
       // Only move ops residing in the same block.
       return defBlock == block;
     };
-    mlir::getBackwardSlice(op, &backwardSet, options);
+    (void)mlir::getBackwardSlice(op, &backwardSet, options);
     backwardSet.insert(op);
 
     // Don't move a local_store if its source is a load from
diff --git a/third_party/nvidia/hopper/lib/Transforms/WarpSpecialization/WSCodePartition.cpp b/third_party/nvidia/hopper/lib/Transforms/WarpSpecialization/WSCodePartition.cpp
diff --git a/third_party/nvidia/hopper/lib/Transforms/WarpSpecialization/WSDataPartition.cpp b/third_party/nvidia/hopper/lib/Transforms/WarpSpecialization/WSDataPartition.cpp
diff --git a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/PTXAsmFormat.cpp b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/PTXAsmFormat.cpp

Original file line number	Diff line number	Diff line change
`@@ -1 +1 @@`
`1`		`-e12cbd8339b89563059c2bb2a312579b652560d0`
	`1`	`+8957e64a20fc7f4277565c6cfe3e555c119783ce`
Original file line number	Diff line number	Diff line change
`@@ -207,7 +207,7 @@ static void threadValuesThroughWait(ttng::WarpGroupDotWaitOp wait,`
`207`	`207`	`return op->getBlock() == wait->getBlock();`
`208`	`208`	`};`
`209`	`209`	`SetVector<Operation *> slice;`
`210`		`- getBackwardSlice(v, &slice, options);`
	`210`	`+ (void)getBackwardSlice(v, &slice, options);`
`211`	`211`	`}`
`212`	`212`
`213`	`213`	`for (ttng::WarpGroupDotOp dot : asyncDots) {`