[Backend] Bump to llvm/llvm-project@49d5bb0ad0cb (#8766)

This updates LLVM to pick up fixes

* llvm/llvm-project#165692

To unblock ASAN breakages.

---------

Co-authored-by: ravil-mobile <[email protected]>
Co-authored-by: enjustli <[email protected]>
Co-authored-by: Thomas Raoux <[email protected]>

Author: 4 people

cmake/llvm-hash.txt

@@ -1 +1 @@
-f6ded0be897e2878612dd903f7e8bb85448269e5
+49d5bb0ad0cb31410184c462801c5049ad671517

lib/Dialect/TritonGPU/IR/Ops.cpp

@@ -929,8 +929,9 @@ void WarpSpecializeOp::getSuccessorRegions(
     return;
   }
   // And the default region branches transparently back to the parent.
-  assert(src.getRegionOrNull() == &getDefaultRegion());
-  successors.push_back(RegionSuccessor(getResults()));
+  assert(src.getTerminatorPredecessorOrNull()->getParentRegion() ==
+         &getDefaultRegion());
+  successors.push_back(RegionSuccessor(getOperation(), getResults()));
 }
 
 LogicalResult WarpSpecializeOp::verify() {

lib/Dialect/TritonGPU/Transforms/WarpSpecialization/PartitionLoops.cpp

@@ -338,7 +338,7 @@ void cloneOpsInBlock(Block *block, SmallVector<WarpGroupBuilder> &builders,
               builder.mapping.lookupOrDefault(yieldOp.getOperand(i)));
         }
 
-        builder.create<scf::YieldOp>(op->getLoc(), newYieldOperands);
+        scf::YieldOp::create(builder, op->getLoc(), newYieldOperands);
       }
     } else {
       assert(hasPartition(op));
@@ -449,7 +449,7 @@ LogicalResult triton::gpu::partitionLoop(scf::ForOp loop) {
   for (auto [b, region, partition] : llvm::zip(
            builders, wgOp.getPartitionRegions(), partitions.getPartitions())) {
     if (!llvm::is_contained(getPartitionIds(loop), b.partitionId)) {
-      b.create<nvws::WarpGroupYieldOp>(wgOp.getLoc(), SmallVector<Value>{});
+      nvws::WarpGroupYieldOp::create(b, wgOp.getLoc(), SmallVector<Value>{});
       continue;
     }
     auto newForOp = *region.front().getOps<scf::ForOp>().begin();

python/src/llvm.cc

@@ -57,7 +57,6 @@ createTargetMachine(llvm::Module *module, std::string proc,
   bool disableLLVMOpt = mlir::triton::tools::getBoolEnv("DISABLE_LLVM_OPT");
   if (enable_fp_fusion)
     opt.AllowFPOpFusion = llvm::FPOpFusion::Fast;
-  opt.UnsafeFPMath = false;
   opt.NoInfsFPMath = false;
   opt.NoNaNsFPMath = true;
   opt.TrapUnreachable = true;

test/Conversion/tritonnvidiagpu_to_llvm.mlir

@@ -215,9 +215,9 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.targ
 // CHECK-LABEL: async_copy_mbarrier_arrive
 module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
   tt.func public @async_copy_mbarrier_arrive(%arg0: !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>)  attributes { noinline = false } {
-    // CHECK: nvvm.cp.async.mbarrier.arrive.shared %{{.*}} : !llvm.ptr<3>
+    // CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} : !llvm.ptr<3>
     ttng.async_copy_mbarrier_arrive %arg0 : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>
-    // CHECK: nvvm.cp.async.mbarrier.arrive.shared %{{.*}} {noinc = true} : !llvm.ptr<3>
+    // CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} {noinc = true} : !llvm.ptr<3>
     ttng.async_copy_mbarrier_arrive %arg0 { noIncrement } : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>
     tt.return
   }

third_party/amd/include/Analysis/RangeAnalysis.h

@@ -84,7 +84,7 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
   /// the loop operands and all users and all users of the results of the loop.
   void visitRegionSuccessors(
       ProgramPoint *point, RegionBranchOpInterface branch,
-      RegionBranchPoint successor,
+      RegionSuccessor successor,
       ArrayRef<dataflow::AbstractSparseLattice *> abstractLattices) override;
 
   /// Collect all operands that participate in assumptions (see description of

third_party/amd/lib/Analysis/RangeAnalysis.cpp

@@ -630,7 +630,7 @@ void TritonIntegerRangeAnalysis::initializeFuncOp(tt::FuncOp op) {
 
 void TritonIntegerRangeAnalysis::visitRegionSuccessors(
     ProgramPoint *point, RegionBranchOpInterface branch,
-    RegionBranchPoint successor,
+    RegionSuccessor successor,
     ArrayRef<dataflow::AbstractSparseLattice *> abstractLattices) {
   LLVM_DEBUG({
     DBGS() << "Visit Region Succesors of ";
@@ -715,10 +715,11 @@ void TritonIntegerRangeAnalysis::visitRegionSuccessors(
         if (!inputs.empty()) {
           firstIndex = cast<OpResult>(inputs.front()).getResultNumber();
         }
-        visitNonControlFlowArguments(branch,
-                                     RegionSuccessor(branch->getResults().slice(
-                                         firstIndex, inputs.size())),
-                                     lattices, firstIndex);
+        visitNonControlFlowArguments(
+            branch,
+            RegionSuccessor(
+                branch, branch->getResults().slice(firstIndex, inputs.size())),
+            lattices, firstIndex);
       } else {
         if (!inputs.empty()) {
           firstIndex = cast<BlockArgument>(inputs.front()).getArgNumber();

third_party/amd/lib/TritonAMDGPUToLLVM/MemoryOpToLLVM.cpp

@@ -481,8 +481,8 @@ class LocalBarrierOpConversion
     //   amdgpu::MemoryCounterWaitOp will lower s_waitcnt
     // - s_barrier syncronizes the execution for the CTA
     auto dsAttr = rewriter.getI32IntegerAttr(0);
-    rewriter.create<amdgpu::MemoryCounterWaitOp>(
-        op->getLoc(), /* load= */ nullptr, /* store= */ nullptr,
+    amdgpu::MemoryCounterWaitOp::create(
+        rewriter, op->getLoc(), /* load= */ nullptr, /* store= */ nullptr,
         /* ds= */ dsAttr);
     rewriter.replaceOpWithNewOp<ROCDL::SBarrierOp>(op);
 

third_party/nvidia/lib/Dialect/NVWS/Transforms/LowerAref.cpp

@@ -305,10 +305,9 @@ void createTMALoad(triton::nvws::DescriptorLoadOp op, PatternRewriter &rewriter,
       }
     }
   }
-  auto newLoadOp =
-      rewriter.create<triton::nvidia_gpu::AsyncTMACopyGlobalToLocalOp>(
-          op.getLoc(), op.getDesc(), indices, barrierAlloc, op.getResult(),
-          pred);
+  auto newLoadOp = triton::nvidia_gpu::AsyncTMACopyGlobalToLocalOp::create(
+      rewriter, op.getLoc(), op.getDesc(), indices, barrierAlloc,
+      op.getResult(), pred);
   assignStageCluster(newLoadOp, getPartitionWsTagIds(op), getStageCluster(op),
                      rewriter);
 };

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/ConvertWarpSpecializeToLLVM.cpp

@@ -91,8 +91,8 @@ static void createBarrier(TritonLLVMIRRewriter &b, unsigned barIdx,
   if (numThreads == 32)
     LLVM::NVIDIA::createSyncWarp(b.getLoc(), b);
   else
-    NVVM::BarrierOp::create(b, b.getLoc(), b.i32_val(barIdx),
-                            b.i32_val(numThreads));
+    NVVM::BarrierOp::create(b, b.getLoc(), TypeRange{}, b.i32_val(barIdx),
+                            b.i32_val(numThreads), {}, Value{});
 }
 
 static void createAllBarrier(TritonLLVMIRRewriter &b, unsigned barIdx) {