[flang-commits] [flang] 7e38793 - [flang][cuda] Make sure stream is a i64 reference (#157957)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2025-09-10T22:08:46-07:00
New Revision: 7e38793795006df01f77e0ad44fed3a9198e2d2a

URL: https://github.com/llvm/llvm-project/commit/7e38793795006df01f77e0ad44fed3a9198e2d2a
DIFF: https://github.com/llvm/llvm-project/commit/7e38793795006df01f77e0ad44fed3a9198e2d2a.diff

LOG: [flang][cuda] Make sure stream is a i64 reference (#157957)

When the stream is a scalar constant, it is lowered as i32. Stream needs
to be i64 to pass the verifier. Detect and update the stream reference
when it is i32.

Added: 
    flang/test/Lower/CUDA/cuda-stream.cuf

Modified: 
    flang/lib/Lower/ConvertCall.cpp

Removed: 
    


################################################################################
diff  --git a/flang/lib/Lower/ConvertCall.cpp b/flang/lib/Lower/ConvertCall.cpp
index 3951401ebed37..e82d4ea0904f1 100644
--- a/flang/lib/Lower/ConvertCall.cpp
+++ b/flang/lib/Lower/ConvertCall.cpp
@@ -639,9 +639,18 @@ Fortran::lower::genCallOpAndResult(
               caller.getCallDescription().chevrons()[2], stmtCtx)));
 
     mlir::Value stream; // stream is optional.
-    if (caller.getCallDescription().chevrons().size() > 3)
+    if (caller.getCallDescription().chevrons().size() > 3) {
       stream = fir::getBase(converter.genExprAddr(
           caller.getCallDescription().chevrons()[3], stmtCtx));
+      if (!fir::unwrapRefType(stream.getType()).isInteger(64)) {
+        auto i64Ty = mlir::IntegerType::get(builder.getContext(), 64);
+        mlir::Value newStream = builder.createTemporary(loc, i64Ty);
+        mlir::Value load = fir::LoadOp::create(builder, loc, stream);
+        mlir::Value conv = fir::ConvertOp::create(builder, loc, i64Ty, load);
+        fir::StoreOp::create(builder, loc, conv, newStream);
+        stream = newStream;
+      }
+    }
 
     cuf::KernelLaunchOp::create(builder, loc, funcType.getResults(),
                                 funcSymbolAttr, grid_x, grid_y, grid_z, block_x,

diff  --git a/flang/test/Lower/CUDA/cuda-stream.cuf b/flang/test/Lower/CUDA/cuda-stream.cuf
new file mode 100644
index 0000000000000..a58ab4ed4235a
--- /dev/null
+++ b/flang/test/Lower/CUDA/cuda-stream.cuf
@@ -0,0 +1,15 @@
+! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s
+
+attributes(global) subroutine sharedmem()
+  real, shared :: s(*)
+  integer :: t
+  t = threadIdx%x
+  s(t) = t
+end subroutine
+
+program test
+  call sharedmem<<<1, 1, 1024, 0>>>()
+end
+
+! CHECK-LABEL: func.func @_QQmain()
+! CHECK: cuf.kernel_launch @_QPsharedmem<<<%c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c1{{.*}}, %c1024{{.*}}, %{{.*}} : !fir.ref<i64>>>>()