fix insert allocs

Signed-off-by: dchigarev <[email protected]>

Author: dchigarev

lib/Transforms/CMakeLists.txt

@@ -24,7 +24,11 @@ add_mlir_library(IMEXTransforms
   MLIRSupport
   MLIRTransformUtils
   MLIRVectorTransforms
+  IMEXXeTileDialect
+  IMEXRegionDialect
 
   DEPENDS
   IMEXTransformsPassIncGen
+  IMEXXeTilePassIncGen
+  MLIRRegionOpsIncGen
 )

lib/Transforms/InsertGPUAllocs.cpp

@@ -44,7 +44,7 @@ class InsertGPUAllocsPass final
     : public imex::impl::InsertGPUAllocsBase<InsertGPUAllocsPass> {
 
 public:
-  explicit InsertGPUAllocsPass() : m_clientAPI("vulkan") {}
+  explicit InsertGPUAllocsPass() : m_clientAPI("opencl") {}
   explicit InsertGPUAllocsPass(const mlir::StringRef &clientAPI)
       : m_clientAPI(clientAPI) {}
 
@@ -411,6 +411,15 @@ class InsertGPUAllocsPass final
             use.set(newAlloc.getResult());
           }
         }
+
+        // remove 'memref.dealloc' (it's later replaced with gpu.dealloc)
+        auto memory = alloc->getResult(0);
+        for (auto u : memory.getUsers()) {
+          if (auto dealloc = mlir::dyn_cast<mlir::memref::DeallocOp>(u)) {
+            dealloc.erase();
+          }
+        }
+
         alloc.replaceAllUsesWith(allocResult);
         builder.create<mlir::gpu::DeallocOp>(loc, std::nullopt, allocResult);
         alloc.erase();

lib/Transforms/SetSPIRVAbiAttribute.cpp

@@ -32,7 +32,7 @@ namespace {
 class SetSPIRVAbiAttributePass
     : public imex::impl::SetSPIRVAbiAttributeBase<SetSPIRVAbiAttributePass> {
 public:
-  explicit SetSPIRVAbiAttributePass() { m_clientAPI = "vulkan"; }
+  explicit SetSPIRVAbiAttributePass() { m_clientAPI = "opencl"; }
   explicit SetSPIRVAbiAttributePass(const mlir::StringRef &clientAPI)
       : m_clientAPI(clientAPI) {}