[BACKEND] Update LLVM version to llvm/llvm-project@7a33569 (triton-lang#7378)

Author: gflegar

cmake/llvm-hash.txt

@@ -1 +1 @@
-570885128351868c1308bb22e8ca351d318bc4a1
+7a33569510535f0b917a2e50f644bf57490aee24

include/triton/Conversion/TritonGPUToLLVM/PatternTritonGPUOpToLLVM.h

@@ -95,7 +95,8 @@ void populateSPMDOpToLLVMPattern(LLVMTypeConverter &typeConverter,
 void populateFuncOpConversionPattern(LLVMTypeConverter &typeConverter,
                                      RewritePatternSet &patterns,
                                      const TargetInfoBase &targetInfo,
-                                     PatternBenefit benefit);
+                                     PatternBenefit benefit,
+                                     SymbolTableCollection *symbolTables);
 
 void populatePrintOpToLLVMPattern(LLVMTypeConverter &typeConverter,
                                   RewritePatternSet &patterns,

lib/Conversion/TritonGPUToLLVM/FuncOpToLLVM.cpp

@@ -27,8 +27,10 @@ using namespace mlir::triton;
 /// information.
 struct FuncOpConversion : public ConvertOpToLLVMPattern<triton::FuncOp> {
   FuncOpConversion(LLVMTypeConverter &converter,
-                   const TargetInfoBase &targetInfo, PatternBenefit benefit)
-      : ConvertOpToLLVMPattern(converter, benefit), targetInfo(targetInfo) {}
+                   const TargetInfoBase &targetInfo, PatternBenefit benefit,
+                   SymbolTableCollection *symbolTables)
+      : ConvertOpToLLVMPattern(converter, benefit), targetInfo(targetInfo),
+        symbolTables(symbolTables) {}
 
   /// Only retain those attributes that are not constructed by
   /// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out argument
@@ -146,7 +148,7 @@ struct FuncOpConversion : public ConvertOpToLLVMPattern<triton::FuncOp> {
 
     FailureOr<LLVM::LLVMFuncOp> maybeNewFuncOp =
         mlir::convertFuncOpToLLVMFuncOp(amendedFuncOp, rewriter,
-                                        *getTypeConverter());
+                                        *getTypeConverter(), symbolTables);
     if (failed(maybeNewFuncOp)) {
       return failure();
     }
@@ -196,12 +198,16 @@ struct FuncOpConversion : public ConvertOpToLLVMPattern<triton::FuncOp> {
 
 private:
   const TargetInfoBase &targetInfo;
+  // Store a pointer to the single, pass-wide symbol table
+  SymbolTableCollection *symbolTables;
 };
 
 } // namespace
 
 void mlir::triton::populateFuncOpConversionPattern(
     LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
-    const TargetInfoBase &targetInfo, PatternBenefit benefit) {
-  patterns.add<FuncOpConversion>(typeConverter, targetInfo, benefit);
+    const TargetInfoBase &targetInfo, PatternBenefit benefit,
+    SymbolTableCollection *symbolTables) {
+  patterns.add<FuncOpConversion>(typeConverter, targetInfo, benefit,
+                                 symbolTables);
 }

third_party/amd/lib/TritonAMDGPUToLLVM/TritonGPUToLLVM.cpp

@@ -112,7 +112,8 @@ struct ConvertTritonAMDGPUToLLVM
       TritonLLVMFunctionConversionTarget funcTarget(*context);
       RewritePatternSet funcPatterns(context);
       mlir::triton::populateFuncOpConversionPattern(
-          typeConverter, funcPatterns, targetInfo, patternBenefitDefault);
+          typeConverter, funcPatterns, targetInfo, patternBenefitDefault,
+          /*symTable=*/nullptr);
       mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter,
                                                             funcPatterns);
       if (failed(