[flang][cuda] Update syntax of fir.cuda_kernel_launch to match fir.call

[clementval]

fir.cuda_kernel_launch represents a call to a cuda kernel with the chervon syntax. Its assembly format is meant to match fir.call. This patch updates the format to match the syntax closer for args and their types.

[llvmbot]

@llvm/pr-subscribers-flang-fir-hlfir

Author: Valentin Clement (バレンタイン クレメン) (clementval)

Changes

fir.cuda_kernel_launch represents a call to a cuda kernel with the chervon syntax. Its assembly format is meant to match fir.call. This patch updates the format to match the syntax closer for args and their types.

---

Full diff: https://github.com/llvm/llvm-project/pull/85814.diff

2 Files Affected:

• (modified) flang/include/flang/Optimizer/Dialect/FIROps.td (+1-1)
• (modified) flang/test/Lower/CUDA/cuda-kernel-calls.cuf (+1-1)

diff --git a/flang/include/flang/Optimizer/Dialect/FIROps.td b/flang/include/flang/Optimizer/Dialect/FIROps.td
index f4792637f481c0..db36302df0745b 100644
--- a/flang/include/flang/Optimizer/Dialect/FIROps.td
+++ b/flang/include/flang/Optimizer/Dialect/FIROps.td
@@ -2466,7 +2466,7 @@ def fir_CUDAKernelLaunch : fir_Op<"cuda_kernel_launch", [CallOpInterface,
   let assemblyFormat = [{
     $callee `<` `<` `<` $grid_x `,` $grid_y `,` $grid_z `,`$block_x `,`
         $block_y `,` $block_z ( `,` $bytes^ ( `,` $stream^ )? )? `>` `>` `>`
-        `` `(` ( $args^ `:` type($args) )? `)` attr-dict
+        `` `(` $args `)` `:` `(` type($args) `)` attr-dict
   }];
 
   let extraClassDeclaration = [{
diff --git a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
index 55b5246e59a006..f4327b3261751f 100644
--- a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
+++ b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
@@ -46,7 +46,7 @@ contains
 ! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel0<<<%c10{{.*}}, %c1{{.*}}, %c1{{.*}}, %c20{{.*}}, %c1{{.*}}, %c1{{.*}}, %c2{{.*}}, %c0{{.*}}>>>()
 
     call dev_kernel1<<<1, 32>>>(a)
-! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel1<<<%c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}>>>(%1#1 : !fir.ref<f32>)
+! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel1<<<%c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}>>>(%1#1) : (!fir.ref<f32>)
   end
 
 end

[jeanPerier]

LGTM

[vzakhari]

LGTM

[schweitzpgi]

+1