[flang-commits] [flang] [flang][cuda] Use -1 for grid values when * is used (PR #115534)

[Valentin Clement バレンタイン クレメン via flang-commits]

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/115534

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. 

>From c7a0cf6097bd8a99d9df07b993d6f820ba9f09c7 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Fri, 8 Nov 2024 11:01:04 -0800
Subject: [PATCH] [flang][cuda] Use -1 for grid values when * is used

---
 flang/lib/Lower/ConvertCall.cpp             | 11 +++++++++--
 flang/test/Lower/CUDA/cuda-kernel-calls.cuf |  5 ++++-
 2 files changed, 13 insertions(+), 3 deletions(-)

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