Skip to content

Commit 9ab5362

Browse files
chsiggchsigg
committed
[mlir][gpu] NFC: switch occurrences of gpu.launch_func to custom format.
Reviewed By: herhut Differential Revision: https://reviews.llvm.org/D89929
1 parent 7ae0033 commit 9ab5362

File tree

14 files changed

+62
-34
lines changed

14 files changed

+62
-34
lines changed

‎mlir/test/Conversion/GPUCommon/lower-launch-func-to-gpu-runtime-calls.mlir

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,13 +20,14 @@ module attributes {gpu.container_module} {
2020
func @foo(%buffer: memref<?xf32>) {
2121
%c8 = constant 8 : index
2222
%c32 = constant 32 : i32
23-
"gpu.launch_func"(%c8, %c8, %c8, %c8, %c8, %c8, %c32, %buffer) {
24-
kernel = @kernel_module::@kernel
25-
} : (index, index, index, index, index, index, i32, memref<?xf32>) -> ()
23+
gpu.launch_func @kernel_module::@kernel
24+
blocks in (%c8, %c8, %c8)
25+
threads in (%c8, %c8, %c8)
26+
args(%c32 : i32, %buffer : memref<?xf32>)
2627
return
2728
}
2829

29-
// CHECK: [[C8:%.*]] = llvm.mlir.constant(8 : index) : !llvm.i64
30+
// CHECK: [[C8:%.*]] = llvm.mlir.constant(8 : index) : !llvm.i64
3031
// CHECK: [[ADDRESSOF:%.*]] = llvm.mlir.addressof @[[GLOBAL]]
3132
// CHECK: [[C0:%.*]] = llvm.mlir.constant(0 : index)
3233
// CHECK: [[BINARY:%.*]] = llvm.getelementptr [[ADDRESSOF]]{{\[}}[[C0]], [[C0]]]

‎mlir/test/Conversion/GPUToSPIRV/builtins.mlir

Lines changed: 16 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,8 @@
33
module attributes {gpu.container_module} {
44
func @builtin() {
55
%c0 = constant 1 : index
6-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_x} : (index, index, index, index, index, index) -> ()
6+
gpu.launch_func @kernels::@builtin_workgroup_id_x
7+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
78
return
89
}
910

@@ -26,7 +27,8 @@ module attributes {gpu.container_module} {
2627
module attributes {gpu.container_module} {
2728
func @builtin() {
2829
%c0 = constant 1 : index
29-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_y} : (index, index, index, index, index, index) -> ()
30+
gpu.launch_func @kernels::@builtin_workgroup_id_y
31+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
3032
return
3133
}
3234

@@ -49,7 +51,8 @@ module attributes {gpu.container_module} {
4951
module attributes {gpu.container_module} {
5052
func @builtin() {
5153
%c0 = constant 1 : index
52-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_z} : (index, index, index, index, index, index) -> ()
54+
gpu.launch_func @kernels::@builtin_workgroup_id_z
55+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
5356
return
5457
}
5558

@@ -72,7 +75,8 @@ module attributes {gpu.container_module} {
7275
module attributes {gpu.container_module} {
7376
func @builtin() {
7477
%c0 = constant 1 : index
75-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_x} : (index, index, index, index, index, index) -> ()
78+
gpu.launch_func @kernels::@builtin_workgroup_size_x
79+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
7680
return
7781
}
7882

@@ -96,7 +100,8 @@ module attributes {gpu.container_module} {
96100
module attributes {gpu.container_module} {
97101
func @builtin() {
98102
%c0 = constant 1 : index
99-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_y} : (index, index, index, index, index, index) -> ()
103+
gpu.launch_func @kernels::@builtin_workgroup_size_y
104+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
100105
return
101106
}
102107

@@ -117,7 +122,8 @@ module attributes {gpu.container_module} {
117122
module attributes {gpu.container_module} {
118123
func @builtin() {
119124
%c0 = constant 1 : index
120-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_z} : (index, index, index, index, index, index) -> ()
125+
gpu.launch_func @kernels::@builtin_workgroup_size_z
126+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
121127
return
122128
}
123129

@@ -138,7 +144,8 @@ module attributes {gpu.container_module} {
138144
module attributes {gpu.container_module} {
139145
func @builtin() {
140146
%c0 = constant 1 : index
141-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_local_id_x} : (index, index, index, index, index, index) -> ()
147+
gpu.launch_func @kernels::@builtin_local_id_x
148+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
142149
return
143150
}
144151

@@ -161,7 +168,8 @@ module attributes {gpu.container_module} {
161168
module attributes {gpu.container_module} {
162169
func @builtin() {
163170
%c0 = constant 1 : index
164-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_num_workgroups_x} : (index, index, index, index, index, index) -> ()
171+
gpu.launch_func @kernels::@builtin_num_workgroups_x
172+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
165173
return
166174
}
167175

‎mlir/test/Conversion/GPUToSPIRV/if.mlir

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,9 @@ module attributes {
77
} {
88
func @main(%arg0 : memref<10xf32>, %arg1 : i1) {
99
%c0 = constant 1 : index
10-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = @kernels::@kernel_simple_selection} : (index, index, index, index, index, index, memref<10xf32>, i1) -> ()
10+
gpu.launch_func @kernels::@kernel_simple_selection
11+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
12+
args(%arg0 : memref<10xf32>, %arg1 : i1)
1113
return
1214
}
1315

‎mlir/test/Conversion/GPUToSPIRV/load-store.mlir

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,10 @@ module attributes {
1515
%1 = subi %c4, %c0_0 : index
1616
%c1_1 = constant 1 : index
1717
%c1_2 = constant 1 : index
18-
"gpu.launch_func"(%0, %c1_2, %c1_2, %1, %c1_2, %c1_2, %arg0, %arg1, %arg2, %c0, %c0_0, %c1, %c1_1) {kernel = @kernels::@load_store_kernel} : (index, index, index, index, index, index, memref<12x4xf32>, memref<12x4xf32>, memref<12x4xf32>, index, index, index, index) -> ()
18+
gpu.launch_func @kernels::@load_store_kernel
19+
blocks in (%0, %c1_2, %c1_2) threads in (%1, %c1_2, %c1_2)
20+
args(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>, %arg2 : memref<12x4xf32>,
21+
%c0 : index, %c0_0 : index, %c1 : index, %c1_1 : index)
1922
return
2023
}
2124

‎mlir/test/Conversion/GPUToSPIRV/loop.mlir

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,9 @@ module attributes {
77
} {
88
func @loop(%arg0 : memref<10xf32>, %arg1 : memref<10xf32>) {
99
%c0 = constant 1 : index
10-
"gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = @kernels::@loop_kernel} : (index, index, index, index, index, index, memref<10xf32>, memref<10xf32>) -> ()
10+
gpu.launch_func @kernels::@loop_kernel
11+
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
12+
args(%arg0 : memref<10xf32>, %arg1 : memref<10xf32>)
1113
return
1214
}
1315

‎mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,9 @@ module attributes {
2222
%0 = "op"() : () -> (f32)
2323
%1 = "op"() : () -> (memref<12xf32, 11>)
2424
%cst = constant 1 : index
25-
"gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure }
26-
: (index, index, index, index, index, index, f32, memref<12xf32, 11>) -> ()
25+
gpu.launch_func @kernels::@basic_module_structure
26+
blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
27+
args(%0 : f32, %1 : memref<12xf32, 11>)
2728
return
2829
}
2930
}

‎mlir/test/Conversion/GPUToSPIRV/simple.mlir

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,9 @@ module attributes {gpu.container_module} {
1818
%0 = "op"() : () -> (f32)
1919
%1 = "op"() : () -> (memref<12xf32>)
2020
%cst = constant 1 : index
21-
"gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure }
22-
: (index, index, index, index, index, index, f32, memref<12xf32>) -> ()
21+
gpu.launch_func @kernels::@basic_module_structure
22+
blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
23+
args(%0 : f32, %1 : memref<12xf32>)
2324
return
2425
}
2526
}
@@ -63,8 +64,9 @@ module attributes {gpu.container_module} {
6364
%0 = "op"() : () -> (f32)
6465
%1 = "op"() : () -> (memref<12xf32>)
6566
%cst = constant 1 : index
66-
"gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@missing_entry_point_abi }
67-
: (index, index, index, index, index, index, f32, memref<12xf32>) -> ()
67+
gpu.launch_func @kernels::@missing_entry_point_abi
68+
blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
69+
args(%0 : f32, %1 : memref<12xf32>)
6870
return
6971
}
7072
}

‎mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,10 @@ module attributes {gpu.container_module} {
2626
func @foo() {
2727
%0 = alloc() : memref<12xf32>
2828
%c1 = constant 1 : index
29-
"gpu.launch_func"(%c1, %c1, %c1, %c1, %c1, %c1, %0) {kernel = @kernels::@kernel} : (index, index, index, index, index, index, memref<12xf32>) -> ()
29+
gpu.launch_func @kernels::@kernel
30+
blocks in(%c1, %c1, %c1)
31+
threads in(%c1, %c1, %c1)
32+
args(%0 : memref<12xf32>)
3033
return
3134
}
3235
}

‎mlir/test/mlir-vulkan-runner/addf.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -37,8 +37,9 @@ module attributes {
3737

3838
%cst1 = constant 1 : index
3939
%cst8 = constant 8 : index
40-
"gpu.launch_func"(%cst8, %cst1, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add }
41-
: (index, index, index, index, index, index, memref<8xf32>, memref<8xf32>, memref<8xf32>) -> ()
40+
gpu.launch_func @kernels::@kernel_add
41+
blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
42+
args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
4243
%arg6 = memref_cast %arg5 : memref<?xf32> to memref<*xf32>
4344
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
4445
return

‎mlir/test/mlir-vulkan-runner/addi.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -36,8 +36,9 @@ module attributes {
3636

3737
%cst1 = constant 1 : index
3838
%cst8 = constant 8 : index
39-
"gpu.launch_func"(%cst8, %cst8, %cst8, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_addi }
40-
: (index, index, index, index, index, index, memref<8xi32>, memref<8x8xi32>, memref<8x8x8xi32>) -> ()
39+
gpu.launch_func @kernels::@kernel_addi
40+
blocks in (%cst8, %cst8, %cst8) threads in (%cst1, %cst1, %cst1)
41+
args(%arg0 : memref<8xi32>, %arg1 : memref<8x8xi32>, %arg2 : memref<8x8x8xi32>)
4142
%arg6 = memref_cast %arg5 : memref<?x?x?xi32> to memref<*xi32>
4243
call @print_memref_i32(%arg6) : (memref<*xi32>) -> ()
4344
return

‎mlir/test/mlir-vulkan-runner/addi8.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -37,8 +37,9 @@ module attributes {
3737

3838
%cst1 = constant 1 : index
3939
%cst8 = constant 8 : index
40-
"gpu.launch_func"(%cst8, %cst8, %cst8, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_addi }
41-
: (index, index, index, index, index, index, memref<8xi8>, memref<8x8xi8>, memref<8x8x8xi32>) -> ()
40+
gpu.launch_func @kernels::@kernel_addi
41+
blocks in (%cst8, %cst8, %cst8) threads in (%cst1, %cst1, %cst1)
42+
args(%arg0 : memref<8xi8>, %arg1 : memref<8x8xi8>, %arg2 : memref<8x8x8xi32>)
4243
%arg6 = memref_cast %arg5 : memref<?x?x?xi32> to memref<*xi32>
4344
call @print_memref_i32(%arg6) : (memref<*xi32>) -> ()
4445
return

‎mlir/test/mlir-vulkan-runner/mulf.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -38,8 +38,9 @@ module attributes {
3838

3939
%cst1 = constant 1 : index
4040
%cst4 = constant 4 : index
41-
"gpu.launch_func"(%cst4, %cst4, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_mul }
42-
: (index, index, index, index, index, index, memref<4x4xf32>, memref<4x4xf32>, memref<4x4xf32>) -> ()
41+
gpu.launch_func @kernels::@kernel_mul
42+
blocks in (%cst4, %cst4, %cst1) threads in(%cst1, %cst1, %cst1)
43+
args(%arg0 : memref<4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<4x4xf32>)
4344
%arg6 = memref_cast %arg5 : memref<?x?xf32> to memref<*xf32>
4445
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
4546
return

‎mlir/test/mlir-vulkan-runner/subf.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -40,8 +40,9 @@ module attributes {
4040
%cst1 = constant 1 : index
4141
%cst4 = constant 4 : index
4242
%cst8 = constant 8 : index
43-
"gpu.launch_func"(%cst8, %cst4, %cst4, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_sub }
44-
: (index, index, index, index, index, index, memref<8x4x4xf32>, memref<4x4xf32>, memref<8x4x4xf32>) -> ()
43+
gpu.launch_func @kernels::@kernel_sub
44+
blocks in (%cst8, %cst4, %cst4) threads in (%cst1, %cst1, %cst1)
45+
args(%arg0 : memref<8x4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<8x4x4xf32>)
4546
%arg6 = memref_cast %arg5 : memref<?x?x?xf32> to memref<*xf32>
4647
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
4748
return

‎mlir/test/mlir-vulkan-runner/time.mlir

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -44,8 +44,9 @@ module attributes {
4444

4545
%cst1 = constant 1 : index
4646
%cst128 = constant 128 : index
47-
"gpu.launch_func"(%cst128, %cst1, %cst1, %cst128, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add }
48-
: (index, index, index, index, index, index, memref<16384xf32>, memref<16384xf32>, memref<16384xf32>) -> ()
47+
gpu.launch_func @kernels::@kernel_add
48+
blocks in (%cst128, %cst1, %cst1) threads in (%cst128, %cst1, %cst1)
49+
args(%arg0 : memref<16384xf32>, %arg1 : memref<16384xf32>, %arg2 : memref<16384xf32>)
4950
%arg6 = memref_cast %arg5 : memref<?xf32> to memref<*xf32>
5051
return
5152
}

0 commit comments

Comments
 (0)