[BACKEND] Apply patches for LLVM version up to llvm/llvm-project@2bc22ea (#8707)

Author: khasanovaa

lib/Dialect/TritonGPU/IR/Ops.cpp

@@ -933,8 +933,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() {

python/src/llvm.cc

@@ -53,7 +53,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;

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

@@ -702,7 +702,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 ";
@@ -787,10 +787,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/nvidia/lib/TritonNVIDIAGPUToLLVM/ElementwiseOpToLLVM.cpp

@@ -650,7 +650,7 @@ struct ExpOpConversionApprox
     Value prod = b.fmul(f32_ty, operands[0][0], b.f32_val(log2e));
 
     Type resultTy = operands[0][0].getType();
-    StringRef name = "llvm.nvvm.ex2.approx.f";
+    StringRef name = "llvm.nvvm.ex2.approx.f32";
     auto callOp =
         LLVM::createLLVMIntrinsicCallOp(rewriter, loc, name, resultTy, {prod});
     return {callOp.getResult(0)};