[BACKEND] Update LLVM version to llvm/llvm-project@7752e0a (#6735)

Removed `addArgumentMaterialization` since that method was removed from
MLIR in
llvm/llvm-project@23e3cbb

Author: gflegar

cmake/llvm-hash.txt

@@ -1 +1 @@
-092b6e73e651469527662443b592f98f442ece72
+7752e0a10b25da2f2eadbed10606bd5454dbca05

lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp

@@ -295,7 +295,7 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
             offset);
       }
       auto vecAddr = b.gep(sharedPtrTy, elemTy, smemBase, offset);
-      vecAddr.setInbounds(true);
+      vecAddr.setNoWrapFlags(mlir::LLVM::GEPNoWrapFlags::inbounds);
       return vecAddr;
     };
 

lib/Conversion/TritonGPUToLLVM/Utility.cpp

@@ -399,7 +399,7 @@ Value getSmemVecAddr(const LinearLayout &regLayout,
   }
   auto ptrTy = smemBase.getType();
   auto vecAddr = b.gep(ptrTy, elemLlvmTy, smemBase, smemOffset);
-  vecAddr.setInbounds(true);
+  vecAddr.setNoWrapFlags(mlir::LLVM::GEPNoWrapFlags::inbounds);
   return vecAddr;
 }
 

lib/Conversion/TritonToTritonGPU/TritonGPUConversion.cpp

@@ -55,15 +55,6 @@ TritonGPUTypeConverter::TritonGPUTypeConverter(MLIRContext *context,
   //
   // Materializations
   //
-  // This will be called when (newArgType != origArgType)
-  // This will create newArg, and map(origArg, newArg)
-  addArgumentMaterialization([](OpBuilder &builder, RankedTensorType tensorType,
-                                ValueRange inputs, Location loc) -> Value {
-    llvm_unreachable("Argument rematerialization should not happen in Triton "
-                     "-> TritonGPU conversion");
-    return {};
-  });
-
   // If the origValue still has live user(s), use this to
   // convert origValue to newValue
   addSourceMaterialization([=](OpBuilder &builder, RankedTensorType tensorType,

third_party/amd/lib/TritonAMDGPUToLLVM/BuiltinFuncToLLVM.cpp

@@ -204,7 +204,7 @@ struct ConvertBuiltinFuncToLLVM
     ModuleOp mod = getOperation();
 
     GreedyRewriteConfig config;
-    config.enableRegionSimplification = GreedySimplifyRegionLevel::Aggressive;
+    config.setRegionSimplificationLevel(GreedySimplifyRegionLevel::Aggressive);
 
     RewritePatternSet patterns(context);
     patterns.add<CallOpConversion>(context, this->ftz);

third_party/amd/lib/TritonAMDGPUToLLVM/TritonGPUToLLVM.cpp

@@ -11,6 +11,7 @@
 #include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
 #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h"
 #include "mlir/Conversion/UBToLLVM/UBToLLVM.h"
+#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
@@ -24,6 +25,7 @@
 #include "triton/Dialect/Triton/IR/Dialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"
+#include "llvm/TargetParser/TargetParser.h"
 
 #include "third_party/proton/dialect/include/TritonProtonToLLVM/PatternTritonProtonOpToLLVM.h"
 
@@ -86,6 +88,15 @@ struct ConvertTritonAMDGPUToLLVM
       mod.emitError("unsupported target: '") << this->arch.getValue() << "'";
       return signalPassFailure();
     }
+    llvm::StringRef chipset =
+        llvm::AMDGPU::getArchNameAMDGCN(targetInfo.getGPUKind());
+    llvm::FailureOr<mlir::amdgpu::Chipset> maybeChipset =
+        mlir::amdgpu::Chipset::parse(chipset);
+    if (failed(maybeChipset)) {
+      mlir::emitError(mlir::UnknownLoc::get(&getContext()),
+                      "Invalid chipset name: " + chipset);
+      return signalPassFailure();
+    }
 
     mlir::LowerToLLVMOptions option(context);
     option.overrideIndexBitwidth(32);
@@ -210,8 +221,8 @@ struct ConvertTritonAMDGPUToLLVM
     mlir::populateMathToLLVMConversionPatterns(typeConverter, patterns);
 
     // Native lowering patterns
-    mlir::populateGpuToROCDLConversionPatterns(typeConverter, patterns,
-                                               mlir::gpu::amd::HIP);
+    mlir::populateGpuToROCDLConversionPatterns(
+        typeConverter, patterns, mlir::gpu::amd::HIP, *maybeChipset);
 
     mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter,
                                                           patterns);

third_party/amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp

@@ -13,7 +13,7 @@
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Pass/Pass.h"
-#include "mlir/Transforms/OneToNTypeConversion.h"
+#include "mlir/Transforms/DialectConversion.h"
 #include "triton/Analysis/Utility.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"
 #include "triton/Dialect/Triton/IR/Types.h"

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/MemoryOpToLLVM.cpp

@@ -191,7 +191,7 @@ LogicalResult lowerDistributedToSharedStmatrix(
             .second;
     Value offset = b.xor_(regBase, b.i32_val(regIdx));
     auto vecAddr = b.gep(smemPtrTy, llvmElemTy, smemBase, offset);
-    vecAddr.setInbounds(true);
+    vecAddr.setNoWrapFlags(mlir::LLVM::GEPNoWrapFlags::inbounds);
     SmallVector<Value> inValsVec;
     for (int j = 0; j < srcVec; j++)
       inValsVec.push_back(srcVals[i + j]);

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TargetInfo.cpp

@@ -230,7 +230,7 @@ void TargetInfo::storeDShared(RewriterBase &rewriter, Location loc, Value ptr,
     SmallVector<Value> vals = unpackLLVector(loc, val, rewriter);
     for (int i = 0; i < vec / maxVec; i++) {
       auto newPtr = b.gep(ptr.getType(), elemTy, ptr, b.i32_val(i * maxVec),
-                          /*inbounds=*/true);
+                          mlir::LLVM::GEPNoWrapFlags::inbounds);
       storeDShared(
           rewriter, loc, newPtr, ctaId,
           packLLVector(loc, ArrayRef(vals).slice(i * maxVec, maxVec), rewriter),
@@ -343,7 +343,7 @@ Value TargetInfo::loadDShared(RewriterBase &rewriter, Location loc, Value ptr,
     SmallVector<Value> vals;
     for (int i = 0; i < vec / maxVec; i++) {
       auto newPtr = b.gep(ptr.getType(), elemTy, ptr, b.i32_val(i * maxVec),
-                          /*inbounds=*/true);
+                          mlir::LLVM::GEPNoWrapFlags::inbounds);
       auto newVal = loadDShared(rewriter, loc, newPtr, ctaId,
                                 vec_ty(elemTy, maxVec), pred);
       for (Value v : unpackLLVector(loc, newVal, rewriter)) {