[mlir][gpu] NFC: switch occurrences of gpu.launch_func to custom format.

[csigg]

Rebase.

[herhut]

Thanks for cleaning this up!

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

llvm.func @kernel(%arg0: !llvm.i32, %arg1: !llvm.ptr<float>,
         %arg5: !llvm.i64) attributes {gpu.kernel} {
       llvm.return
     }
   }
 
   func @foo(%buffer: memref<?xf32>) {
     %c8 = constant 8 : index
     %c32 = constant 32 : i32
-    "gpu.launch_func"(%c8, %c8, %c8, %c8, %c8, %c8, %c32, %buffer) {
-      kernel = @kernel_module::@kernel
-    } : (index, index, index, index, index, index, i32, memref<?xf32>) -> ()
+    gpu.launch_func @kernel_module::@kernel
+        blocks in (%c8, %c8, %c8)
+        threads in (%c8, %c8, %c8)
+        args(%c32 : i32, %buffer : memref<?xf32>)
     return
   }
 
-  // CHECK: [[C8:%.*]] = llvm.mlir.constant(8 : index) : !llvm.i64   
+  // CHECK: [[C8:%.*]] = llvm.mlir.constant(8 : index) : !llvm.i64
   // CHECK: [[ADDRESSOF:%.*]] = llvm.mlir.addressof @[[GLOBAL]]
   // CHECK: [[C0:%.*]] = llvm.mlir.constant(0 : index)
   // CHECK: [[BINARY:%.*]] = llvm.getelementptr [[ADDRESSOF]]{{\[}}[[C0]], [[C0]]]
   // CHECK-SAME: -> !llvm.ptr<i8>
 
   // CHECK: [[MODULE:%.*]] = llvm.call @mgpuModuleLoad([[BINARY]])
   // CHECK: [[FUNC:%.*]] = llvm.call @mgpuModuleGetFunction([[MODULE]], {{.*}})
 

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_x} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_id_x
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_y} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_id_y
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_z} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_id_z
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_x} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_size_x
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // Note that this ignores the workgroup size specification in gpu.launch.
       // We may want to define gpu.workgroup_size and convert it to the entry
       // point ABI we want here.
       // CHECK: spv.constant 32 : i32
       %0 = "gpu.block_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_y} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_size_y
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.constant 4 : i32
       %0 = "gpu.block_dim"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_z} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_workgroup_size_z
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.constant 1 : i32
       %0 = "gpu.block_dim"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_local_id_x} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_local_id_x
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_local_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.thread_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
-    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_num_workgroups_x} : (index, index, index, index, index, index) -> ()
+    gpu.launch_func @kernels::@builtin_num_workgroups_x
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {

mlir/test/Conversion/GPUToSPIRV/if.mlir

 // RUN: mlir-opt -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
   func @main(%arg0 : memref<10xf32>, %arg1 : i1) {
     %c0 = constant 1 : index
-    "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) -> ()
+    gpu.launch_func @kernels::@kernel_simple_selection
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
+        args(%arg0 : memref<10xf32>, %arg1 : i1)
     return
   }
 
   gpu.module @kernels {
     // CHECK-LABEL: @kernel_simple_selection
     gpu.func @kernel_simple_selection(%arg2 : memref<10xf32>, %arg3 : i1) kernel
     attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       %value = constant 0.0 : f32

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

func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
     %c12 = constant 12 : index
     %0 = subi %c12, %c0 : index
     %c1 = constant 1 : index
     %c0_0 = constant 0 : index
     %c4 = constant 4 : index
     %1 = subi %c4, %c0_0 : index
     %c1_1 = constant 1 : index
     %c1_2 = constant 1 : index
-    "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) -> ()
+    gpu.launch_func @kernels::@load_store_kernel
+        blocks in (%0, %c1_2, %c1_2) threads in (%1, %c1_2, %c1_2)
+        args(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>, %arg2 : memref<12x4xf32>, 
+             %c0 : index, %c0_0 : index, %c1 : index, %c1_1 : index)
     return
   }
 
   // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     // CHECK-DAG: spv.globalVariable @[[NUMWORKGROUPSVAR:.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable @[[$WORKGROUPIDVAR:.*]] built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>

mlir/test/Conversion/GPUToSPIRV/loop.mlir

 // RUN: mlir-opt -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
   func @loop(%arg0 : memref<10xf32>, %arg1 : memref<10xf32>) {
     %c0 = constant 1 : index
-    "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>) -> ()
+    gpu.launch_func @kernels::@loop_kernel
+        blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
+        args(%arg0 : memref<10xf32>, %arg1 : memref<10xf32>)
     return
   }
 
   gpu.module @kernels {
     gpu.func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) kernel
     attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: %[[LB:.*]] = spv.constant 4 : i32
       %lb = constant 4 : index

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

gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel
       gpu.return
     }
   }
 
   func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32, 11>)
     %cst = constant 1 : index
-    "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure }
-        : (index, index, index, index, index, index, f32, memref<12xf32, 11>) -> ()
+    gpu.launch_func @kernels::@basic_module_structure
+        blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
+        args(%0 : f32, %1 : memref<12xf32, 11>)
     return
   }
 }

mlir/test/Conversion/GPUToSPIRV/simple.mlir

gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>) kernel
       gpu.return
     }
   }
 
   func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32>)
     %cst = constant 1 : index
-    "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure }
-        : (index, index, index, index, index, index, f32, memref<12xf32>) -> ()
+    gpu.launch_func @kernels::@basic_module_structure
+        blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
+        args(%0 : f32, %1 : memref<12xf32>)
     return
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
gpu.func @missing_entry_point_abi(%arg0 : f32, %arg1 : memref<12xf32>) kernel {
       gpu.return
     }
   }
 
   func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32>)
     %cst = constant 1 : index
-    "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@missing_entry_point_abi }
-        : (index, index, index, index, index, index, f32, memref<12xf32>) -> ()
+    gpu.launch_func @kernels::@missing_entry_point_abi
+        blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
+        args(%0 : f32, %1 : memref<12xf32>)
     return
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {

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

module attributes {gpu.container_module} {
   gpu.module @kernels {
     gpu.func @kernel(%arg0: memref<12xf32>) kernel {
       gpu.return
     }
   }
   func @foo() {
     %0 = alloc() : memref<12xf32>
     %c1 = constant 1 : index
-    "gpu.launch_func"(%c1, %c1, %c1, %c1, %c1, %c1, %0) {kernel = @kernels::@kernel} : (index, index, index, index, index, index, memref<12xf32>) -> ()
+    gpu.launch_func @kernels::@kernel
+        blocks in(%c1, %c1, %c1)
+        threads in(%c1, %c1, %c1)
+        args(%0 : memref<12xf32>)
     return
   }
 }

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

func @main() {
     %arg4 = memref_cast %arg1 : memref<8xf32> to memref<?xf32>
     %arg5 = memref_cast %arg2 : memref<8xf32> to memref<?xf32>
     call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
     call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
     call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()
 
     %cst1 = constant 1 : index
     %cst8 = constant 8 : index
-    "gpu.launch_func"(%cst8, %cst1, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add }
-        : (index, index, index, index, index, index, memref<8xf32>, memref<8xf32>, memref<8xf32>) -> ()
+    gpu.launch_func @kernels::@kernel_add
+        blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
+        args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
     %arg6 = memref_cast %arg5 : memref<?xf32> to memref<*xf32>
     call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
     return
   }
   func @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
   func @print_memref_f32(%ptr : memref<*xf32>)
 }
 

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

func @main() {
     %arg4 = memref_cast %arg1 : memref<8x8xi32> to memref<?x?xi32>
     %arg5 = memref_cast %arg2 : memref<8x8x8xi32> to memref<?x?x?xi32>
     call @fillResource1DInt(%arg3, %value1) : (memref<?xi32>, i32) -> ()
     call @fillResource2DInt(%arg4, %value2) : (memref<?x?xi32>, i32) -> ()
     call @fillResource3DInt(%arg5, %value0) : (memref<?x?x?xi32>, i32) -> ()
 
     %cst1 = constant 1 : index
     %cst8 = constant 8 : index
-    "gpu.launch_func"(%cst8, %cst8, %cst8, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_addi }
-        : (index, index, index, index, index, index, memref<8xi32>, memref<8x8xi32>, memref<8x8x8xi32>) -> ()
+    gpu.launch_func @kernels::@kernel_addi
+        blocks in (%cst8, %cst8, %cst8) threads in (%cst1, %cst1, %cst1)
+        args(%arg0 : memref<8xi32>, %arg1 : memref<8x8xi32>, %arg2 : memref<8x8x8xi32>)
     %arg6 = memref_cast %arg5 : memref<?x?x?xi32> to memref<*xi32>
     call @print_memref_i32(%arg6) : (memref<*xi32>) -> ()
     return
   }
   func @fillResource1DInt(%0 : memref<?xi32>, %1 : i32)
   func @fillResource2DInt(%0 : memref<?x?xi32>, %1 : i32)
   func @fillResource3DInt(%0 : memref<?x?x?xi32>, %1 : i32)
   func @print_memref_i32(%ptr : memref<*xi32>)
 }
 

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

func @main() {
     %arg4 = memref_cast %arg1 : memref<8x8xi8> to memref<?x?xi8>
     %arg5 = memref_cast %arg2 : memref<8x8x8xi32> to memref<?x?x?xi32>
     call @fillResource1DInt8(%arg3, %value1) : (memref<?xi8>, i8) -> ()
     call @fillResource2DInt8(%arg4, %value2) : (memref<?x?xi8>, i8) -> ()
     call @fillResource3DInt(%arg5, %value0) : (memref<?x?x?xi32>, i32) -> ()
 
     %cst1 = constant 1 : index
     %cst8 = constant 8 : index
-    "gpu.launch_func"(%cst8, %cst8, %cst8, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_addi }
-        : (index, index, index, index, index, index, memref<8xi8>, memref<8x8xi8>, memref<8x8x8xi32>) -> ()
+    gpu.launch_func @kernels::@kernel_addi
+        blocks in (%cst8, %cst8, %cst8) threads in (%cst1, %cst1, %cst1)
+        args(%arg0 : memref<8xi8>, %arg1 : memref<8x8xi8>, %arg2 : memref<8x8x8xi32>)
     %arg6 = memref_cast %arg5 : memref<?x?x?xi32> to memref<*xi32>
     call @print_memref_i32(%arg6) : (memref<*xi32>) -> ()
     return
   }
   func @fillResource1DInt8(%0 : memref<?xi8>, %1 : i8)
   func @fillResource2DInt8(%0 : memref<?x?xi8>, %1 : i8)
   func @fillResource3DInt(%0 : memref<?x?x?xi32>, %1 : i32)
   func @print_memref_i32(%ptr : memref<*xi32>)
 }

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

func @main() {
     %arg4 = memref_cast %arg1 : memref<4x4xf32> to memref<?x?xf32>
     %arg5 = memref_cast %arg2 : memref<4x4xf32> to memref<?x?xf32>
     call @fillResource2DFloat(%arg3, %value1) : (memref<?x?xf32>, f32) -> ()
     call @fillResource2DFloat(%arg4, %value2) : (memref<?x?xf32>, f32) -> ()
     call @fillResource2DFloat(%arg5, %value0) : (memref<?x?xf32>, f32) -> ()
 
     %cst1 = constant 1 : index
     %cst4 = constant 4 : index
-    "gpu.launch_func"(%cst4, %cst4, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_mul }
-        : (index, index, index, index, index, index, memref<4x4xf32>, memref<4x4xf32>, memref<4x4xf32>) -> ()
+    gpu.launch_func @kernels::@kernel_mul
+        blocks in (%cst4, %cst4, %cst1) threads in(%cst1, %cst1, %cst1)
+        args(%arg0 : memref<4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<4x4xf32>)
     %arg6 = memref_cast %arg5 : memref<?x?xf32> to memref<*xf32>
     call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
     return
   }
   func @fillResource2DFloat(%0 : memref<?x?xf32>, %1 : f32)
   func @print_memref_f32(%ptr : memref<*xf32>)
 }
 

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

func @main() {
     %arg5 = memref_cast %arg2 : memref<8x4x4xf32> to memref<?x?x?xf32>
     call @fillResource3DFloat(%arg3, %value1) : (memref<?x?x?xf32>, f32) -> ()
     call @fillResource2DFloat(%arg4, %value2) : (memref<?x?xf32>, f32) -> ()
     call @fillResource3DFloat(%arg5, %value0) : (memref<?x?x?xf32>, f32) -> ()
 
     %cst1 = constant 1 : index
     %cst4 = constant 4 : index
     %cst8 = constant 8 : index
-    "gpu.launch_func"(%cst8, %cst4, %cst4, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_sub }
-        : (index, index, index, index, index, index, memref<8x4x4xf32>, memref<4x4xf32>, memref<8x4x4xf32>) -> ()
+    gpu.launch_func @kernels::@kernel_sub
+        blocks in (%cst8, %cst4, %cst4) threads in (%cst1, %cst1, %cst1)
+        args(%arg0 : memref<8x4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<8x4x4xf32>)
     %arg6 = memref_cast %arg5 : memref<?x?x?xf32> to memref<*xf32>
     call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()
     return
   }
   func @fillResource2DFloat(%0 : memref<?x?xf32>, %1 : f32)
   func @fillResource3DFloat(%0 : memref<?x?x?xf32>, %1 : f32)
   func @print_memref_f32(%ptr : memref<*xf32>)
 }

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

func @main() {
     %arg4 = memref_cast %arg1 : memref<16384xf32> to memref<?xf32>
     %arg5 = memref_cast %arg2 : memref<16384xf32> to memref<?xf32>
     call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
     call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
     call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()
 
     %cst1 = constant 1 : index
     %cst128 = constant 128 : index
-    "gpu.launch_func"(%cst128, %cst1, %cst1, %cst128, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add }
-        : (index, index, index, index, index, index, memref<16384xf32>, memref<16384xf32>, memref<16384xf32>) -> ()
+    gpu.launch_func @kernels::@kernel_add
+        blocks in (%cst128, %cst1, %cst1) threads in (%cst128, %cst1, %cst1)
+        args(%arg0 : memref<16384xf32>, %arg1 : memref<16384xf32>, %arg2 : memref<16384xf32>)
     %arg6 = memref_cast %arg5 : memref<?xf32> to memref<*xf32>
     return
   }
   func @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
   func @print_memref_f32(%ptr : memref<*xf32>)
 }