[flang][cuda][NFC] Add test for <<<*, block>>> lowering

[clementval]

Chevron syntax has been update to allow * to be used for the grid value. Make sure we set the three grid values to -1, 1, 1 in lowering.

[llvmbot]

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

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

Changes

Chevron syntax has been update to allow * to be used for the grid value. Make sure we set the three grid values to -1 in lowering.

---

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

2 Files Affected:

• (modified) flang/lib/Lower/ConvertCall.cpp (+9-2)
• (modified) flang/test/Lower/CUDA/cuda-kernel-calls.cuf (+4-1)

diff --git a/flang/lib/Lower/ConvertCall.cpp b/flang/lib/Lower/ConvertCall.cpp
index 9f5b58590fb79e..eaf5a25e4390ef 100644
--- a/flang/lib/Lower/ConvertCall.cpp
+++ b/flang/lib/Lower/ConvertCall.cpp
@@ -541,8 +541,15 @@ std::pair<fir::ExtendedValue, bool> Fortran::lower::genCallOpAndResult(
           loc, i32Ty,
           fir::getBase(converter.genExprValue(
               caller.getCallDescription().chevrons()[0], stmtCtx)));
-      grid_y = one;
-      grid_z = one;
+      auto gridXValue = fir::getIntIfConstant(grid_x);
+      if (gridXValue && *gridXValue < 0) {
+        // Call using * for grid size.
+        grid_y = grid_x;
+        grid_z = grid_x;
+      } else {
+        grid_y = one;
+        grid_z = one;
+      }
     } else {
       auto dim3Addr = converter.genExprAddr(
           caller.getCallDescription().chevrons()[0], stmtCtx);
diff --git a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
index 82d1a61f8e157c..08ec2f433f5838 100644
--- a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
+++ b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf
@@ -47,7 +47,10 @@ contains
 ! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel0<<<%c10{{.*}}, %c1{{.*}}, %c1{{.*}}, %c20{{.*}}, %c1{{.*}}, %c1{{.*}}, %c2{{.*}}, %c0{{.*}}>>>()
 
     call dev_kernel1<<<1, 32>>>(a)
-! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel1<<<%c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}>>>(%1#1) : (!fir.ref<f32>)
+! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel1<<<%c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}>>>(%{{.*}}) : (!fir.ref<f32>)
+
+    call dev_kernel1<<<*, 32>>>(a)
+! CHECK: cuf.kernel_launch @_QMtest_callPdev_kernel1<<<%c-1{{.*}}, %c-1{{.*}}, %c-1{{.*}}, %c32{{.*}}, %c1{{.*}}, %c1{{.*}}>>>(%{{.*}})
   end
 
 end