Migrate to llvm version 19.1.3.

tud-ccc · Oct 30, 2024 · 426c819 · 426c819
1 parent 105bfaa
commit 426c819
Show file tree

Hide file tree

Showing 6 changed files with 6 additions and 8 deletions.
diff --git a/cinnamon/include/cinm-mlir/Dialect/UPMEM/IR/UPMEMOps.td b/cinnamon/include/cinm-mlir/Dialect/UPMEM/IR/UPMEMOps.td
@@ -33,7 +33,7 @@ include "mlir/IR/RegionKindInterface.td"
 
 class UPMEM_IndexOp<string mnemonic, list<Trait> traits = []> :
     UPMEM_Op<mnemonic, !listconcat(traits, [
-        Pure, DeclareOpInterfaceMethods<InferIntRangeInterface>])>, 
+        Pure, DeclareOpInterfaceMethods<InferIntRangeInterface, ["inferResultRanges"]>])>, 
         Results<(outs Index)> {
   let assemblyFormat = "attr-dict";
 }
@@ -128,7 +128,7 @@ def UPMEM_PrivateWRAMAllocOp : UPMEM_Op<"pwram_alloc", [
 def UPMEM_LaunchOp : UPMEM_Op<"launch", [
       AutomaticAllocationScope, AttrSizedOperandSegments, UPMEM_AsyncOpInterface,
       IsolatedFromAbove,
-      DeclareOpInterfaceMethods<InferIntRangeInterface>]> {
+      DeclareOpInterfaceMethods<InferIntRangeInterface, ["inferResultRanges"]>]> {
   let summary = "UPMEM kernel launch operation";
 
   let arguments = (ins 

diff --git a/cinnamon/lib/Conversion/CnmToGPU/CnmToGPU.cpp b/cinnamon/lib/Conversion/CnmToGPU/CnmToGPU.cpp
@@ -282,7 +282,6 @@ struct ConvertCnmToGPUPass
 
     RewritePatternSet patterns(&getContext());
     populateCnmToGPUConversionPatterns(patterns, &getContext());
-    populateReconcileUnrealizedCastsPatterns(patterns);
 
     ConversionTarget target(getContext());
     target.addIllegalDialect<cnm::CnmDialect>();

diff --git a/cinnamon/lib/Conversion/CnmToUPMEM/CnmToUPMEM.cpp b/cinnamon/lib/Conversion/CnmToUPMEM/CnmToUPMEM.cpp
@@ -24,6 +24,7 @@
 #include <mlir/IR/ValueRange.h>
 #include <mlir/Support/LLVM.h>
 #include <mlir/Transforms/DialectConversion.h>
+#include <mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h>
 
 #define GEN_PASS_DEF_CONVERTCNMTOUPMEMPASS
 #include "cinm-mlir/Conversion/CnmPasses.h.inc"
@@ -355,7 +356,6 @@ struct ConvertCnmToUPMEMPass
 
     RewritePatternSet patterns(&getContext());
     populateCnmToUPMEMConversionPatterns(converter, patterns);
-    populateReconcileUnrealizedCastsPatterns(patterns);
     populateFinalBufferizationPatterns(patterns);
 
     ConversionTarget target(getContext());

diff --git a/cinnamon/lib/Conversion/UPMEMToLLVM/UPMEMToLLVM.cpp b/cinnamon/lib/Conversion/UPMEMToLLVM/UPMEMToLLVM.cpp
@@ -43,6 +43,7 @@
 #include <mlir/IR/Visitors.h>
 #include <mlir/Support/LogicalResult.h>
 #include <mlir/Transforms/DialectConversion.h>
+#include <mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h>
 #include <optional>
 
 namespace mlir {
@@ -315,7 +316,7 @@ outlineAffineMap(ImplicitLocOpBuilder &rewriter,
   affineMapFun->setAttr("upmem.generated_from", AffineMapAttr::get(*linearMap));
 
   rewriter = ImplicitLocOpBuilder::atBlockBegin(rewriter.getLoc(),
-                                                affineMapFun.addEntryBlock());
+                                                affineMapFun.addEntryBlock(rewriter));
   Value arg = affineMapFun.getArgument(0);
   // affine expects to deal with index type only
   arg = createOrFoldUnrealizedConversionCast(rewriter.getLoc(), rewriter,
@@ -570,7 +571,6 @@ struct ConvertUPMEMToLLVMPass
     RewritePatternSet patterns(&getContext());
     populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
     populateUPMEMToLLVMConversionPatterns(converter, patterns);
-    populateReconcileUnrealizedCastsPatterns(patterns);
 
     ConversionTarget target(getContext());
     target.addIllegalDialect<upmem::UPMEMDialect>();

diff --git a/cinnamon/lib/Dialect/Cnm/Transforms/SPIRVAttachAttributes.cpp b/cinnamon/lib/Dialect/Cnm/Transforms/SPIRVAttachAttributes.cpp
@@ -51,7 +51,7 @@ void CnmSPIRVAttachKernelEntryPointAttributePass::runOnOperation() {
         const DenseI32ArrayAttr workgroup_size = DenseI32ArrayAttr::get(&getContext(), {1, 1, 1});
         const std::optional<int> subgroup_size;
 
-        gpuFunc->setAttr("spirv.entry_point_abi", spirv::EntryPointABIAttr::get(&getContext(), workgroup_size, subgroup_size));
+        gpuFunc->setAttr("spirv.entry_point_abi", spirv::EntryPointABIAttr::get(&getContext(), workgroup_size, subgroup_size, {}));
     });
 }
 

diff --git a/cinnamon/lib/Dialect/UPMEM/IR/InferIntRangeInterfaceImpls.cpp b/cinnamon/lib/Dialect/UPMEM/IR/InferIntRangeInterfaceImpls.cpp
@@ -65,7 +65,6 @@ static std::optional<uint64_t> getKnownLaunchDim(Op op, LaunchDims type) {
   return std::nullopt;
 }
 
-
 void RankDimOp::inferResultRanges(ArrayRef<ConstantIntRanges>,
                                   SetIntRangeFn setResultRange) {
   uint64_t max = getKnownLaunchDim(*this, LaunchDims::Rank).value_or(kMaxRankDim);