From 5fbf00d929e2e436fd20dc6f88f3ab8d68d5fb2d Mon Sep 17 00:00:00 2001 From: Valentin Clement Date: Wed, 10 Sep 2025 14:04:07 -0700 Subject: [PATCH 1/2] [flang][cuda] Make sure stream is a i64 reference --- flang/lib/Lower/ConvertCall.cpp | 11 ++++++++++- flang/test/Lower/CUDA/cuda-stream.cuf | 15 +++++++++++++++ 2 files changed, 25 insertions(+), 1 deletion(-) create mode 100644 flang/test/Lower/CUDA/cuda-stream.cuf diff --git a/flang/lib/Lower/ConvertCall.cpp b/flang/lib/Lower/ConvertCall.cpp index 3951401ebed37..fbbc8b1fa83cb 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 = fir::AllocaOp::create(builder, 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>>>() From 88c283203fd499a30969688b648190522be937ce Mon Sep 17 00:00:00 2001 From: Valentin Clement Date: Wed, 10 Sep 2025 14:24:55 -0700 Subject: [PATCH 2/2] Use builder.createTemporary --- flang/lib/Lower/ConvertCall.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/flang/lib/Lower/ConvertCall.cpp b/flang/lib/Lower/ConvertCall.cpp index fbbc8b1fa83cb..e82d4ea0904f1 100644 --- a/flang/lib/Lower/ConvertCall.cpp +++ b/flang/lib/Lower/ConvertCall.cpp @@ -644,7 +644,7 @@ Fortran::lower::genCallOpAndResult( caller.getCallDescription().chevrons()[3], stmtCtx)); if (!fir::unwrapRefType(stream.getType()).isInteger(64)) { auto i64Ty = mlir::IntegerType::get(builder.getContext(), 64); - mlir::Value newStream = fir::AllocaOp::create(builder, loc, i64Ty); + 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);