From 426c8195a31abd7c5d4531f6955fad147fb075a4 Mon Sep 17 00:00:00 2001
From: Felix Reissmann <felix.rm.153@gmail.com>
Date: Wed, 30 Oct 2024 13:26:21 +0100
Subject: [PATCH] Migrate to llvm version 19.1.3.

---
 cinnamon/include/cinm-mlir/Dialect/UPMEM/IR/UPMEMOps.td       | 4 ++--
 cinnamon/lib/Conversion/CnmToGPU/CnmToGPU.cpp                 | 1 -
 cinnamon/lib/Conversion/CnmToUPMEM/CnmToUPMEM.cpp             | 2 +-
 cinnamon/lib/Conversion/UPMEMToLLVM/UPMEMToLLVM.cpp           | 4 ++--
 cinnamon/lib/Dialect/Cnm/Transforms/SPIRVAttachAttributes.cpp | 2 +-
 cinnamon/lib/Dialect/UPMEM/IR/InferIntRangeInterfaceImpls.cpp | 1 -
 6 files changed, 6 insertions(+), 8 deletions(-)

diff --git a/cinnamon/include/cinm-mlir/Dialect/UPMEM/IR/UPMEMOps.td b/cinnamon/include/cinm-mlir/Dialect/UPMEM/IR/UPMEMOps.td
index bb3124d..87f6d7b 100644
--- 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
index cb797e5..cdced48 100644
--- 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
index a40a897..0c498a5 100644
--- 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
index 6f8f15e..2c8b5d6 100644
--- 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
index 6230dd7..db2a12c 100644
--- 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
index 09cc04e..3e4fb9b 100644
--- 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);