[Mlir-commits] [mlir] [mlir][GPU] Fix double spaces in tests after ODS printer fix. NFC. (PR #185325)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun Mar 8 15:28:18 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-gpu
Author: Jakub Kuderski (kuhar)
<details>
<summary>Changes</summary>
Follow-up to #<!-- -->184253. The ODS attr/type printer fix removed the leading space from generated print() methods. Update tests that checked for the old double-space output of GPU ops using GPU_DimensionAttr and GPU_MmaElementwiseOpAttr.
---
Patch is 43.66 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/185325.diff
25 Files Affected:
- (modified) mlir/test/Conversion/VectorToGPU/vector-to-mma-ops.mlir (+6-6)
- (modified) mlir/test/Dialect/Affine/ops.mlir (+4-4)
- (modified) mlir/test/Dialect/GPU/subgroup-mma-vector-unroll.mlir (+4-4)
- (modified) mlir/test/Dialect/GPU/subgroupId-rewrite.mlir (+5-5)
- (modified) mlir/test/Dialect/GPU/transform-gpu.mlir (+38-38)
- (modified) mlir/test/Dialect/SparseTensor/GPU/gpu_matmul.mlir (+4-4)
- (modified) mlir/test/Dialect/SparseTensor/GPU/gpu_matvec.mlir (+4-4)
- (modified) mlir/test/Dialect/Vector/vector-warp-distribute.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/cga_cluster.mlir (+9-9)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir (+1-1)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir (+1-1)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x128_stride_noswizzle.mlir (+7-7)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x64_swizzle128b.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir (+1-1)
- (modified) mlir/test/Integration/GPU/LevelZero/gpu-addf32-to-spirv.mlir (+3-3)
- (modified) mlir/test/Integration/GPU/LevelZero/gpu-addi64-to-spirv.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/LevelZero/gpu-memcpy-addf32-to-spirv.mlir (+3-3)
- (modified) mlir/test/Integration/GPU/LevelZero/gpu-reluf32-to-spirv.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/SYCL/gpu-addf32-to-spirv.mlir (+3-3)
- (modified) mlir/test/Integration/GPU/SYCL/gpu-addi64-to-spirv.mlir (+2-2)
- (modified) mlir/test/Integration/GPU/SYCL/gpu-memcpy-addf32-to-spirv.mlir (+3-3)
- (modified) mlir/test/Integration/GPU/SYCL/gpu-reluf32-to-spirv.mlir (+2-2)
- (modified) mlir/test/python/dialects/gpu/dialect.py (+4-4)
``````````diff
diff --git a/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops.mlir b/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops.mlir
index 32065035b6f21..5b316f3e9e219 100644
--- a/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops.mlir
+++ b/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops.mlir
@@ -463,7 +463,7 @@ func.func @matmul_mixed_signedness_int8(%arg0: memref<16x32xi8>, %arg1: memref<1
// CHECK-LABEL: func @cast_f16_to_f32_write
// CHECK: %[[COMPUTE:.+]] = gpu.subgroup_mma_compute
-// CHECK: %[[EXT:.+]] = gpu.subgroup_mma_elementwise extf %[[COMPUTE]] : (!gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf32, "COp">
+// CHECK: %[[EXT:.+]] = gpu.subgroup_mma_elementwise extf %[[COMPUTE]] : (!gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf32, "COp">
// CHECK: gpu.subgroup_mma_store_matrix %[[EXT]]
func.func @cast_f16_to_f32_write(%arg0: memref<16x16xf16>, %arg1: memref<16x16xf16>, %arg2: memref<16x16xf16>, %arg3: memref<16x16xf32>) {
%c0 = arith.constant 0 : index
@@ -485,7 +485,7 @@ func.func @cast_f16_to_f32_write(%arg0: memref<16x16xf16>, %arg1: memref<16x16xf
// CHECK-LABEL: func @cast_f32_to_f16_write
// CHECK: %[[COMPUTE:.+]] = gpu.subgroup_mma_compute
-// CHECK: %[[EXT:.+]] = gpu.subgroup_mma_elementwise truncf %[[COMPUTE]] : (!gpu.mma_matrix<16x16xf32, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
+// CHECK: %[[EXT:.+]] = gpu.subgroup_mma_elementwise truncf %[[COMPUTE]] : (!gpu.mma_matrix<16x16xf32, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
// CHECK: gpu.subgroup_mma_store_matrix %[[EXT]]
func.func @cast_f32_to_f16_write(%arg0: memref<16x16xf32>, %arg1: memref<16x16xf32>, %arg2: memref<16x16xf32>, %arg3: memref<16x16xf16>) {
%c0 = arith.constant 0 : index
@@ -536,10 +536,10 @@ func.func @fold_transpose_into_transfer_read(%alloc: memref<64x128xf16>, %vector
// CHECK-LABEL: func @cast_f16_to_f32_read
// CHECK: %[[A:.+]] = gpu.subgroup_mma_load_matrix {{.+}} {leadDimension = 16 : index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "AOp">
// CHECK: %[[C:.+]] = gpu.subgroup_mma_load_matrix {{.+}} {leadDimension = 16 : index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "COp">
-// CHECK: %[[AE:.+]] = gpu.subgroup_mma_elementwise extf %[[A]] : (!gpu.mma_matrix<16x16xf16, "AOp">) -> !gpu.mma_matrix<16x16xf32, "AOp">
-// CHECK: %[[CE:.+]] = gpu.subgroup_mma_elementwise extf %[[C]] : (!gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf32, "COp">
+// CHECK: %[[AE:.+]] = gpu.subgroup_mma_elementwise extf %[[A]] : (!gpu.mma_matrix<16x16xf16, "AOp">) -> !gpu.mma_matrix<16x16xf32, "AOp">
+// CHECK: %[[CE:.+]] = gpu.subgroup_mma_elementwise extf %[[C]] : (!gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf32, "COp">
// CHECK: %[[B:.+]] = gpu.subgroup_mma_load_matrix {{.+}} {leadDimension = 16 : index, transpose} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "BOp">
-// CHECK: %[[BE:.+]] = gpu.subgroup_mma_elementwise extf %[[B]] : (!gpu.mma_matrix<16x16xf16, "BOp">) -> !gpu.mma_matrix<16x16xf32, "BOp">
+// CHECK: %[[BE:.+]] = gpu.subgroup_mma_elementwise extf %[[B]] : (!gpu.mma_matrix<16x16xf16, "BOp">) -> !gpu.mma_matrix<16x16xf32, "BOp">
// CHECK: gpu.subgroup_mma_compute %[[AE]], %[[BE]], %[[CE]]
func.func @cast_f16_to_f32_read(%arg0: memref<16x16xf16>, %arg1: memref<16x16xf16>, %arg2: memref<16x16xf16>, %arg3: memref<16x16xf32>) {
%c0 = arith.constant 0 : index
@@ -582,7 +582,7 @@ func.func @test_unsupported(%arg0: vector<4x4xi32>, %arg1: vector<4x4xi32>, %arg
// CHECK-LABEL: func @addf
// CHECK: %[[A:.+]] = gpu.subgroup_mma_load_matrix {{.+}} {leadDimension = 16 : index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "COp">
// CHECK: %[[B:.+]] = gpu.subgroup_mma_load_matrix {{.+}} {leadDimension = 16 : index, transpose} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "COp">
-// CHECK: %[[C:.+]] = gpu.subgroup_mma_elementwise addf %[[A]], %[[B]] : (!gpu.mma_matrix<16x16xf16, "COp">, !gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
+// CHECK: %[[C:.+]] = gpu.subgroup_mma_elementwise addf %[[A]], %[[B]] : (!gpu.mma_matrix<16x16xf16, "COp">, !gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
// CHECK: gpu.subgroup_mma_store_matrix %[[C]]
func.func @addf(%arg0: memref<16x16xf16>, %arg1: memref<16x16xf16>, %arg2: memref<16x16xf16>) {
%c0 = arith.constant 0 : index
diff --git a/mlir/test/Dialect/Affine/ops.mlir b/mlir/test/Dialect/Affine/ops.mlir
index 8a3f41d1d9b05..0992d392bcd12 100644
--- a/mlir/test/Dialect/Affine/ops.mlir
+++ b/mlir/test/Dialect/Affine/ops.mlir
@@ -328,7 +328,7 @@ module {
%c1 = arith.constant 1 : index
gpu.launch blocks(%arg0, %arg1, %arg2) in (%arg6 = %c1, %arg7 = %c1, %arg8 = %c1)
threads(%arg3, %arg4, %arg5) in (%arg9 = %c1, %arg10 = %c1, %arg11 = %c1) {
- %thread_id_x = gpu.thread_id x
+ %thread_id_x = gpu.thread_id x
%c128 = arith.constant 128 : index
affine.for %arg12 = %thread_id_x to %c128 step 8 {
}
@@ -338,7 +338,7 @@ module {
}
}
-// CHECK: %[[THREAD_ID:.*]] = gpu.thread_id x
+// CHECK: %[[THREAD_ID:.*]] = gpu.thread_id x
// CHECK: %[[VAL:.*]] = arith.constant 128 : index
// CHECK: affine.for %{{.*}} = %[[THREAD_ID]] to %[[VAL]] step 8 {
@@ -357,7 +357,7 @@ module {
%dim = memref.dim %arg0, %c3 : memref<?x?xf32>
%c0 = arith.constant 0 : index
affine.for %arg3 = %c0 to %dim step 32 {
- %thread_id_x = gpu.thread_id x
+ %thread_id_x = gpu.thread_id x
%0 = affine.apply #map()[%thread_id_x]
%c128 = arith.constant 128 : index
affine.for %arg4 = %0 to %c128 step 8 {
@@ -374,7 +374,7 @@ module {
// CHECK: %[[VAL_2:.*]] = memref.dim %[[VAL_0]], %[[VAL_1]] : memref<?x?xf32>
// CHECK: %[[VAL_3:.*]] = arith.constant 0 : index
// CHECK: affine.for %[[VAL_4:.*]] = %[[VAL_3]] to %[[VAL_2]] step 32 {
-// CHECK: %[[VAL_5:.*]] = gpu.thread_id x
+// CHECK: %[[VAL_5:.*]] = gpu.thread_id x
// CHECK: %[[VAL_6:.*]] = affine.apply #[[$ATTR_0]](){{\[}}%[[VAL_5]]]
// CHECK: %[[VAL_7:.*]] = arith.constant 128 : index
// CHECK: affine.for %{{.*}} = %[[VAL_6]] to %[[VAL_7]] step 8 {
diff --git a/mlir/test/Dialect/GPU/subgroup-mma-vector-unroll.mlir b/mlir/test/Dialect/GPU/subgroup-mma-vector-unroll.mlir
index 03aba89c11afc..8b0a62e9c6387 100644
--- a/mlir/test/Dialect/GPU/subgroup-mma-vector-unroll.mlir
+++ b/mlir/test/Dialect/GPU/subgroup-mma-vector-unroll.mlir
@@ -7,8 +7,8 @@ func.func @matmul(%lhs: memref<32x32xf32>, %rhs: memref<32x32xf32>, %out: memref
%c16 = arith.constant 16 : index
%c32 = arith.constant 32 : index
%cst_0 = arith.constant 0.000000e+00 : f32
- %3 = gpu.thread_id x
- %4 = gpu.thread_id y
+ %3 = gpu.thread_id x
+ %4 = gpu.thread_id y
%5 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%4]
%6 = affine.apply affine_map<()[s0] -> ((s0 floordiv 32) * 16)>()[%3]
// CHECK: scf.for {{.*}} -> (vector<16x16xf32>) {
@@ -58,8 +58,8 @@ func.func @gathered_matmul(%lhs: memref<32x32xf32>, %rhs: memref<32x32xf32>, %ou
%cst_1 = arith.constant dense<[0, 1, 2, 3]> : vector<4xindex>
%cst_2 = arith.constant dense<1> : vector<4x4xindex>
%alloc = memref.alloc() {alignment = 64 : i64} : memref<32x32xf32>
- %3 = gpu.thread_id x
- %4 = gpu.thread_id y
+ %3 = gpu.thread_id x
+ %4 = gpu.thread_id y
%5 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%4]
%6 = affine.apply affine_map<()[s0] -> ((s0 floordiv 32) * 16)>()[%3]
// CHECK: scf.for {{.*}} -> (vector<16x16xf32>) {
diff --git a/mlir/test/Dialect/GPU/subgroupId-rewrite.mlir b/mlir/test/Dialect/GPU/subgroupId-rewrite.mlir
index 386793ad88649..0d4f4d590bb4e 100644
--- a/mlir/test/Dialect/GPU/subgroupId-rewrite.mlir
+++ b/mlir/test/Dialect/GPU/subgroupId-rewrite.mlir
@@ -5,11 +5,11 @@
func.func @subgroupId(%sz : index, %mem: memref<index, 1>) {
gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %sz, %grid_y = %sz, %grid_z = %sz)
threads(%tx, %ty, %tz) in (%block_x = %sz, %block_y = %sz, %block_z = %sz) {
- // CHECK: %[[DIMX:.*]] = gpu.block_dim x
- // CHECK-NEXT: %[[DIMY:.*]] = gpu.block_dim y
- // CHECK-NEXT: %[[TIDX:.*]] = gpu.thread_id x
- // CHECK-NEXT: %[[TIDY:.*]] = gpu.thread_id y
- // CHECK-NEXT: %[[TIDZ:.*]] = gpu.thread_id z
+ // CHECK: %[[DIMX:.*]] = gpu.block_dim x
+ // CHECK-NEXT: %[[DIMY:.*]] = gpu.block_dim y
+ // CHECK-NEXT: %[[TIDX:.*]] = gpu.thread_id x
+ // CHECK-NEXT: %[[TIDY:.*]] = gpu.thread_id y
+ // CHECK-NEXT: %[[TIDZ:.*]] = gpu.thread_id z
// CHECK-NEXT: %[[T0:.*]] = arith.muli %[[DIMY]], %[[TIDZ]] : index
// CHECK-NEXT: %[[T1:.*]] = arith.addi %[[T0]], %[[TIDY]] : index
// CHECK-NEXT: %[[T2:.*]] = arith.muli %[[DIMX]], %[[T1]] : index
diff --git a/mlir/test/Dialect/GPU/transform-gpu.mlir b/mlir/test/Dialect/GPU/transform-gpu.mlir
index 465e8fdd66422..7e4a02109227a 100644
--- a/mlir/test/Dialect/GPU/transform-gpu.mlir
+++ b/mlir/test/Dialect/GPU/transform-gpu.mlir
@@ -12,8 +12,8 @@ func.func @blocks_3d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream :
%c7 = arith.constant 7 : index
%one = arith.constant 1 : index
// CHECK: gpu.launch
-// CHECK: %[[BLKX:.*]] = gpu.block_id x
-// CHECK: %[[BLKY:.*]] = gpu.block_id y
+// CHECK: %[[BLKX:.*]] = gpu.block_id x
+// CHECK: %[[BLKY:.*]] = gpu.block_id y
// CHECK: memref.load %[[ARGX]][%[[BLKX]], %[[BLKY]]]
// CHECK: memref.load %[[ARGY]][%[[BLKX]], %[[BLKY]]]
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
@@ -59,8 +59,8 @@ func.func @warpgroup_3d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream
// CHECK-DAG: %[[C512:.*]] = arith.constant 512 : index
// CHECK: gpu.launch
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDY:.*]] = gpu.thread_id y
// CHECK-DAG: %[[WG:.*]] = affine.apply #[[$MAP]]()[%[[TIDX]]]
// CHECK-DAG: %[[CMPX:.*]] = arith.cmpi ult, %[[TIDX]], %[[C384]] : index
// CHECK-DAG: %[[CMPY:.*]] = arith.cmpi ult, %[[TIDY]], %[[C1]] : index
@@ -112,8 +112,8 @@ func.func @warp_3d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream : !g
// CHECK-DAG: %[[c64:.*]] = arith.constant 64 : index
// CHECK: gpu.launch
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDY:.*]] = gpu.thread_id y
// CHECK-DAG: %[[W:.*]] = affine.apply #[[$MAP]]()[%[[TIDX]]]
// CHECK-DAG: %[[CMPX:.*]] = arith.cmpi ult, %[[TIDX]], %[[C32]] : index
// CHECK-DAG: %[[CMPY:.*]] = arith.cmpi ult, %[[TIDY]], %[[C3]] : index
@@ -162,8 +162,8 @@ func.func @threads_3d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream :
// CHECK: %[[C9:.*]] = arith.constant 9 : index
// CHECK: %[[C7:.*]] = arith.constant 7 : index
// CHECK: gpu.launch async [%{{.*}}] blocks(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C1]], %{{.*}} = %[[C1]], %{{.*}} = %[[C1]]) threads(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C12]], %{{.*}} = %[[C9]], %{{.*}} = %[[C1]])
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDY:.*]] = gpu.thread_id y
// CHECK: arith.cmpi ult, %[[TIDX]], %[[C9]] : index
// CHECK: arith.cmpi ult, %[[TIDY]], %[[C7]] : index
// CHECK: memref.load %[[ARGX]][%[[TIDY]], %[[TIDX]]]
@@ -215,10 +215,10 @@ func.func @saxpy4d(%x: !type4d, %y: !type4d, %alpha : f32) -> !type4d {
// CHECK: %[[C4:.*]] = arith.constant 4 : index
// CHECK: %[[C1:.*]] = arith.constant 1 : index
// CHECK: gpu.launch blocks(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C32]], %{{.*}} = %[[C64]], %{{.*}} = %[[C1]]) threads(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C32]], %{{.*}} = %[[C4]], %{{.*}} = %[[C1]])
-// CHECK: %[[BLKX:.*]] = gpu.block_id x
-// CHECK: %[[BLKY:.*]] = gpu.block_id y
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK: %[[BLKX:.*]] = gpu.block_id x
+// CHECK: %[[BLKY:.*]] = gpu.block_id y
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDY:.*]] = gpu.thread_id y
// CHECK: memref.load %[[ARGX]][%[[BLKX]], %[[BLKY]], %[[TIDY]], %[[TIDX]]]
// CHECK: memref.load %[[ARGY]][%[[BLKX]], %[[BLKY]], %[[TIDY]], %[[TIDX]]]
scf.forall (%i, %j) in (%c32, %c64) {
@@ -288,7 +288,7 @@ func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token)
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
{
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
// CHECK: memref.load %[[ARGX]][%[[TIDX]], %[[TIDX]]]
// CHECK: memref.load %[[ARGY]][%[[TIDX]], %[[TIDX]]]
scf.forall (%i) in (%c32) {
@@ -322,7 +322,7 @@ func.func @saxpy3d_fold_id_z(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %s
%c9 = arith.constant 9 : index
%c7 = arith.constant 7 : index
// CHECK: %[[C0:.+]] = arith.constant 0 : index
-// CHECK-NOT: gpu.thread_id z
+// CHECK-NOT: gpu.thread_id z
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
{
@@ -373,9 +373,9 @@ func.func @warpgroup_linear(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %st
// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : index
// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index
-// CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
-// CHECK-DAG: %[[TIDZ:.*]] = gpu.thread_id z
+// CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK-DAG: %[[TIDZ:.*]] = gpu.thread_id z
// CHECK-DAG: %[[WIDLIN:.*]] = affine.apply #[[$MAPWGLIN]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
// CHECK-DAG: %[[WIDX:.*]] = affine.apply #[[$MAPWGX]]()[%[[TIDX]], %[[TIDY]]]
// CHECK-DAG: %[[WIDY:.*]] = affine.apply #[[$MAPWGY]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
@@ -429,9 +429,9 @@ func.func @warp_linear(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream
// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index
// CHECK-DAG: %[[C192:.*]] = arith.constant 192 : index
-// CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
-// CHECK-DAG: %[[TIDZ:.*]] = gpu.thread_id z
+// CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK-DAG: %[[TIDZ:.*]] = gpu.thread_id z
// CHECK-DAG: %[[WIDLIN:.*]] = affine.apply #[[$MAPWLIN]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
// CHECK-DAG: %[[WIDX:.*]] = affine.apply #[[$MAPWX]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
// CHECK-DAG: %[[WIDY:.*]] = affine.apply #[[$MAPWY]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
@@ -495,8 +495,8 @@ func.func @map_multi_level_linear(%x: !type, %y: !type, %t: !type1d, %alpha : f3
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
{
- // CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
- // CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
+ // CHECK-DAG: %[[TIDX:.*]] = gpu.thread_id x
+ // CHECK-DAG: %[[TIDY:.*]] = gpu.thread_id y
scf.forall (%i, %j) in (%c7, %c9) {
%4 = memref.load %x[%i, %j] : !type
%5 = memref.load %y[%i, %j] : !type
@@ -563,9 +563,9 @@ func.func @block_linear_existing_launch(
// CHECK-DAG: %[[C12:.*]] = arith.constant 12 : index
// CHECK-DAG: %[[C63:.*]] = arith.constant 63 : index
// CHECK: gpu.launch async [{{.*}}] blocks({{.*}}) in (%{{.*}} = %[[C12]], %{{.*}} = %[[C9]], %{{.*}} = %[[C1]]) threads
-// CHECK-DAG: %[[BIDX:.*]] = gpu.block_id x
-// CHECK-DAG: %[[BIDY:.*]] = gpu.block_id y
-// CHECK-DAG: %[[BIDZ:.*]] = gpu.block_id z
+// CHECK-DAG: %[[BIDX:.*]] = gpu.block_id x
+// CHECK-DAG: %[[BIDY:.*]] = gpu.block_id y
+// CHECK-DAG: %[[BIDZ:.*]] = gpu.block_id z
// CHECK-DAG: %[[BIDLIN:.*]] = affine.apply #[[$MAPBLIN]]()[%[[BIDX]], %[[BIDY]], %[[BIDZ]]]
// CHECK-DAG: %[[BLX:.*]] = affine.apply #[[$MAPBX]]()[%[[BIDX]], %[[BIDY]], %[[BIDZ]]]
// CHECK-DAG: %[[BLY:.*]] = affine.apply #[[$MAPBY]]()[%[[BIDX]], %[[BIDY]], %[[BIDZ]]]
@@ -617,9 +617,9 @@ func.func @block_linear_generate_launch(
// CHECK-DAG: %[[C7:.*]] = arith.constant 7 : index
// CHECK-DAG: %[[C9:.*]] = arith.constant 9 : index
// CHECK: gpu.launch blocks({{.*}}) in (%{{.*}} = %[[C7]], %{{.*}} = %[[C9]], %{{.*}} = %[[C1]]) threads
-// CHECK-DAG: %[[BIDX:.*]] = gpu.block_id x
-// CHECK-DAG: %[[BIDY:.*]] = gpu.block_id y
-// CHECK-DAG: %[[BIDZ:.*]] = gpu.block_id z
+// CHECK-DAG: %[[BIDX:.*]] = gpu.block_id x
+// CHECK-DAG: %[[BIDY:.*]] = gpu.block_id y
+// CHECK-DAG: %[[BIDZ:.*]] = gpu.block_id z
// CHECK-DAG: %[[BLX:.*]] = affine.apply #[[$MAPBX]]()[%[[BIDX]]]
// CHECK-DAG: %[[BLY:.*]] = affine.apply #[[$MAPBY]]()[%[[BIDX]], %[[BIDY]], %[[BIDZ]]]
// CHECK: memref.load %[[ARGX]][%[[BLX]], %[[BLY]]]
@@ -659,14 +659,14 @@ func.func @simple_fill(%arg0: memref<128xf32>) -> memref<128xf32> {
// CHECK: %[[C8:.*]] = arith.constant 8 : index
// CHECK: gpu.launch
scf.forall (%arg1) in (1) {
-// CHECK: %[[BIDX:.*]] = gpu.block_id x
+// CHECK: %[[BIDX:.*]] = gpu.block_id x
// CHECK: %[[BLX:.*]] = affine.apply #[[$MAPB]]()[%[[BIDX]]]
%0 = affine.apply #map(%arg1)
%subview = memref.subview %arg0[%0] [128] [1] : memref<128xf32> to memref<128xf32, strided<[1], offset: ?>>
scf.forall (%arg2) in (4) {
-// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[TIDY:.*]] = gpu.thread_id y
-// CHECK: %[[TIDZ:.*]] = gpu.thread_id z
+// CHECK: %[[TIDX:.*]] = gpu.thread_id x
+// CHECK: %[[TIDY:.*]] = gpu.thread_id y
+// CHECK: %[[TIDZ:.*]] = gpu.thread_id z
// CHECK: %[[THX:.*]] = affine.apply #[[$MAPW]]()[%[[TIDX]], %[[TIDY]], %[[TIDZ]]]
// CHECK-NOT: scf.if
// CHECK: memref.subview %{{.*}}[%[[THX]]]
@@ -709,7 +709,7 @@ func.func @simple_fill(%arg0: memref<128x256xf32>) -> memref<128x256xf32> {
// CHECK: %[[C6:.*]] = arith.constant 6 : index
// CHECK: gpu.launch
scf.forall (%arg1) in (1) {
- // CHECK: %[[BIDX:.*]] = gpu.block_id x
+ // CHECK: %[[BIDX:.*]] = gpu.block_id x
// CHECK: %[[BLX:.*]] = affine.apply #[[$MAPB]]()[%[[BIDX]]]
%0 = affine.apply #map(%arg1)
%subview = memref.subview %arg0[%0, 0] [128, 256] [1, 1]
@@ -719,8 +719,8 @@ func.func @simple_fill(%arg0: memref<128x256xf32>) -> memref<128x256xf32> {
// involving threadIdx.x/y by the map_nested_forall_to_threads
// transformation. This results in a if (linear_thread_id < 6) conditional.
scf.forall (%arg2, %arg3) in (2, 3) {
- // CHECK: %[[TIDX:.*]] = gpu.thread_id x
- // CHECK: %[[TIDY:.*]] = gpu.thread_id y
+ // CHECK: %[[TIDX:.*]] = gpu.thread_id x
+ // CHECK: %[[TIDY:.*]] = gpu.thread_id y
// CHECK: %[[LID:.*]] = affine.apply #[[$MAPLANE]]()[%[[TIDX]], %[[TIDY]]]
// CHECK: %[[COND:.*]] = arith.cmpi ult, %[[LID]], %[[C6]]
// CHECK: scf.if %[[COND]]
@@ -777,7 +777,7 @@ func.func @simple_fill(%arg0: memref<128xf32>) -> memref<128xf32> {
// CHECK: gpu.launch
scf.forall (%arg1) in (1) {
-// CHECK: %[[BIDX:.*]] = gpu.block_id x
+// CHECK: %[[BIDX:.*]] = gpu.block_id x
// CHECK: %[[BLX:.*]] = affine.apply #[[$MAPB]]()[%[[BIDX]]]
%0 = affine.apply #map(%arg1)
%subview = memref.subview %arg0[%0] [128] [1] : memref<128xf32> to memre...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/185325
More information about the Mlir-commits
mailing list