[flang-commits] [flang] 4242d15 - [flang][cuda] Update syntax of fir.cuda_kernel_launch to match fir.call (#85814)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2024-03-19T13:15:58-07:00
New Revision: 4242d15e68580d1d1f7c4d782a59e6dffb698ca8

URL: https://github.com/llvm/llvm-project/commit/4242d15e68580d1d1f7c4d782a59e6dffb698ca8
DIFF: https://github.com/llvm/llvm-project/commit/4242d15e68580d1d1f7c4d782a59e6dffb698ca8.diff

LOG: [flang][cuda] Update syntax of fir.cuda_kernel_launch to match fir.call (#85814)

`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.

Added: 
    

Modified: 
    flang/include/flang/Optimizer/Dialect/FIROps.td
    flang/test/Lower/CUDA/cuda-kernel-calls.cuf

Removed: 
    


################################################################################
diff  --git a/flang/include/flang/Optimizer/Dialect/FIROps.td b/flang/include/flang/Optimizer/Dialect/FIROps.td
index 6e520d111701f0..8a7e36e42457f5 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