Remove old implementation of XeTile-Blocking which

generates 4D representations
intel · Dec 16, 2024 · d9bc716 · d9bc716
1 parent eb22052
commit d9bc716
Show file tree

Hide file tree

Showing 26 changed files with 1,607 additions and 3,330 deletions.
diff --git a/include/imex/Dialect/XeTile/Transforms/Passes.td b/include/imex/Dialect/XeTile/Transforms/Passes.td
@@ -91,10 +91,7 @@ def XeTileBlocking : Pass<"xetile-blocking", "::mlir::gpu::GPUModuleOp">{
   let options = [
      Option<"device", "device", "std::string",
             /*default=*/"\"pvc\"",
-            "gpu platform architecture where these ops are running">,
-     Option<"EnableTransform", "enable-2d-transform", "bool",
-            /*default=*/"false",
-            "Using 2D transform or 4D Conversion.">
+            "gpu platform architecture where these ops are running">
  ];
 }
 

diff --git a/lib/Dialect/XeTile/Transforms/Blocking.cpp b/lib/Dialect/XeTile/Transforms/Blocking.cpp
diff --git a/lib/Dialect/XeTile/Transforms/BlockingAnalysis.cpp b/lib/Dialect/XeTile/Transforms/BlockingAnalysis.cpp
@@ -4,8 +4,6 @@
 
 #include "imex/Dialect/XeTile/Transforms/BlockingAnalysis.h"
 
-extern bool Enable2DBlockingTransform;
-
 namespace llvm {
 using imex::Block;
 // Implementation of llvm::DenseMapInfo for Block, required for
@@ -741,7 +739,7 @@ void BlockingAnalysisImpl::visitCreateMaskOp(
 
   auto lattice = results[0]->getValue();
 
-  if (Enable2DBlockingTransform && !op->use_empty() && !lattice.isInitialized())
+  if (!op->use_empty() && !lattice.isInitialized())
     return;
 
   BlockingRequests &def = getLatticeElement(op->getResult(0))->getValue();
@@ -752,11 +750,9 @@ void BlockingAnalysisImpl::visitCreateMaskOp(
 
   // TODO: need to enable the following code after 2D lowering in
   // GPUToSPIRV is enabled.
-  // if (Enable2DBlockingTransform) {
-  //   for (auto &req : lattice.getRequests()) {
-  //     block[0] = std::max(block[0], req[0]);
-  //     block[1] = std::min(block[1], req[1]);
-  //   }
+  // for (auto &req : lattice.getRequests()) {
+  //   block[0] = std::max(block[0], req[0]);
+  //   block[1] = std::min(block[1], req[1]);
   // }
   def.updateDefBlock(block);
 }

diff --git a/test/Conversion/XeTileToXeGPU/gemm_preop.mlir b/test/Conversion/XeTileToXeGPU/gemm_preop.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" --cse \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking --cse \
 // RUN: --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 #map = affine_map<() -> (0)>
 #map1 = affine_map<() -> (64)>

diff --git a/test/Conversion/XeTileToXeGPU/sg_gemm_1k_1k_1k_f16_f32.mlir b/test/Conversion/XeTileToXeGPU/sg_gemm_1k_1k_1k_f16_f32.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" --canonicalize \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking --canonicalize \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -o -| FileCheck %s
 // CHECK-LABEL: gpu.module @test_kernel {
 gpu.module @test_kernel {

diff --git a/test/Conversion/XeTileToXeGPU/sg_gemm_1k_1k_1k_f16_f32_slm.mlir b/test/Conversion/XeTileToXeGPU/sg_gemm_1k_1k_1k_f16_f32_slm.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" --canonicalize \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking --canonicalize \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -o -| FileCheck %s
 
 

diff --git a/test/Conversion/XeTileToXeGPU/sg_init_tile.mlir b/test/Conversion/XeTileToXeGPU/sg_init_tile.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 
 gpu.module @test_kernel {

diff --git a/test/Conversion/XeTileToXeGPU/sg_load_tile.mlir b/test/Conversion/XeTileToXeGPU/sg_load_tile.mlir
@@ -1,6 +1,5 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
-
 gpu.module @test_kernel {
   //CHECK: gpu.func @sg_load_tile(%[[arg0:.*]]: memref<1024x1024xf16>, %[[arg1:.*]]: memref<1024x1024xf16>, %[[arg2:.*]]: memref<1024x1024xf32>) {
   gpu.func @sg_load_tile(%a: memref<1024x1024xf16>, %b: memref<1024x1024xf16>, %c: memref<1024x1024xf32>) {

diff --git a/test/Conversion/XeTileToXeGPU/sg_mixed_scf.mlir b/test/Conversion/XeTileToXeGPU/sg_mixed_scf.mlir
@@ -1,8 +1,6 @@
 
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse --canonicalize %s -verify-diagnostics -o -| FileCheck %s
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
-// RUN: --cse --convert-xetile-to-xegpu="enable-2d-transform=true" --cse --canonicalize %s -verify-diagnostics -o -| FileCheck %s
 
 //CHECK-LABEL: gpu.module @postop_reduce_m
 gpu.module @postop_reduce_m attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Addresses, Float16Buffer, Int64, Int16, Int8, Bfloat16ConversionINTEL, Kernel, Linkage, Vector16, GenericPointer, Groups, Float16, Float64, AtomicFloat32AddEXT, ExpectAssumeKHR, VectorAnyINTEL], [SPV_INTEL_bfloat16_conversion, SPV_EXT_shader_atomic_float_add, SPV_KHR_expect_assume, SPV_INTEL_vector_compute]>, api=OpenCL, #spirv.resource_limits<>>} {

diff --git a/test/Conversion/XeTileToXeGPU/sg_scattered_ops.mlir b/test/Conversion/XeTileToXeGPU/sg_scattered_ops.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse --canonicalize --cse %s -verify-diagnostics -o -| FileCheck %s
 
 gpu.module @test {

diff --git a/test/Conversion/XeTileToXeGPU/sg_scf_for.mlir b/test/Conversion/XeTileToXeGPU/sg_scf_for.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 gpu.module @test_kernel {
   //CHECK: gpu.func @sglevel_tiled_gemm(%[[arg0:.*]]: memref<1024x1024xf16>, %[[arg1:.*]]: memref<1024x1024xf16>)

diff --git a/test/Conversion/XeTileToXeGPU/sg_softmax.mlir b/test/Conversion/XeTileToXeGPU/sg_softmax.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 gpu.module @test_kernel {
     //CHECK-LABEL: @sglevel_softmax_dim_0

diff --git a/test/Conversion/XeTileToXeGPU/sg_store_tile.mlir b/test/Conversion/XeTileToXeGPU/sg_store_tile.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 
 gpu.module @test_kernel {

diff --git a/test/Conversion/XeTileToXeGPU/sg_tile_mma.mlir b/test/Conversion/XeTileToXeGPU/sg_tile_mma.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-init-duplicate --xetile-blocking \
 // RUN: --convert-xetile-to-xegpu %s -verify-diagnostics -o -| FileCheck %s
 gpu.module @test_kernel {
   //CHECK: s_tiled_gemm(%[[arg0:.*]]: memref<1024x1024xf16>, %[[arg1:.*]]: memref<1024x1024xf16>)

diff --git a/test/Conversion/XeTileToXeGPU/test_blocking.mlir b/test/Conversion/XeTileToXeGPU/test_blocking.mlir
diff --git a/test/Conversion/XeTileToXeGPU/test_order.mlir b/test/Conversion/XeTileToXeGPU/test_order.mlir
@@ -1,4 +1,4 @@
-// RUN: imex-opt --split-input-file --xetile-canonicalization --xetile-blocking="enable-2d-transform=true" \
+// RUN: imex-opt --split-input-file --xetile-canonicalization --xetile-blocking \
 // RUN: --cse --convert-xetile-to-xegpu --cse %s -verify-diagnostics -o -| FileCheck %s
 
 // CHECK-LABEL: @test_func