[flang-commits] [flang] [llvm] [flang][cuda] Add entry points for cudastreamsynchronize (PR #181932)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Tue Feb 17 15:03:07 PST 2026
https://github.com/clementval created https://github.com/llvm/llvm-project/pull/181932
None
>From 2e42a466738451f24cd63a963e79f1b799b149e4 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Tue, 17 Feb 2026 15:02:11 -0800
Subject: [PATCH] [flang][cuda] Add entry points for cudastreamsynchronize
---
flang-rt/lib/cuda/stream.cpp | 9 +++++
.../Optimizer/Builder/CUDAIntrinsicCall.h | 4 +++
flang/include/flang/Runtime/CUDA/stream.h | 2 ++
.../Optimizer/Builder/CUDAIntrinsicCall.cpp | 36 +++++++++++++++++++
flang/module/cuda_runtime_api.f90 | 10 ++++++
flang/test/Lower/CUDA/cuda-default-stream.cuf | 12 +++++++
6 files changed, 73 insertions(+)
diff --git a/flang-rt/lib/cuda/stream.cpp b/flang-rt/lib/cuda/stream.cpp
index af9e63f6a85b1..217aa9501e510 100644
--- a/flang-rt/lib/cuda/stream.cpp
+++ b/flang-rt/lib/cuda/stream.cpp
@@ -29,6 +29,15 @@ int RTDECL(CUFSetDefaultStream)(cudaStream_t stream) {
}
cudaStream_t RTDECL(CUFGetDefaultStream)() { return defaultStream; }
+
+int RTDECL(CUFStreamSynchronize)(cudaStream_t stream) {
+ return cudaStreamSynchronize(stream);
+}
+
+int RTDECL(CUFStreamSynchronizeNull)() {
+ cudaStream_t defaultStream = 0;
+ return cudaStreamSynchronize(defaultStream);
+}
}
} // namespace Fortran::runtime::cuda
diff --git a/flang/include/flang/Optimizer/Builder/CUDAIntrinsicCall.h b/flang/include/flang/Optimizer/Builder/CUDAIntrinsicCall.h
index 3e23a4dfa0203..5e8fbcd1b93ac 100644
--- a/flang/include/flang/Optimizer/Builder/CUDAIntrinsicCall.h
+++ b/flang/include/flang/Optimizer/Builder/CUDAIntrinsicCall.h
@@ -61,6 +61,10 @@ struct CUDAIntrinsicLibrary : IntrinsicLibrary {
llvm::ArrayRef<fir::ExtendedValue>);
mlir::Value genCUDAGetDefaultStreamNull(mlir::Type,
llvm::ArrayRef<mlir::Value>);
+ fir::ExtendedValue
+ genCUDAStreamSynchronize(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
+ mlir::Value genCUDAStreamSynchronizeNull(mlir::Type,
+ llvm::ArrayRef<mlir::Value>);
void genFenceProxyAsync(llvm::ArrayRef<fir::ExtendedValue>);
template <const char *fctName, int extent>
fir::ExtendedValue genLDXXFunc(mlir::Type,
diff --git a/flang/include/flang/Runtime/CUDA/stream.h b/flang/include/flang/Runtime/CUDA/stream.h
index ce97d0d067217..686b828d4e145 100644
--- a/flang/include/flang/Runtime/CUDA/stream.h
+++ b/flang/include/flang/Runtime/CUDA/stream.h
@@ -21,6 +21,8 @@ extern "C" {
int RTDECL(CUFSetDefaultStream)(cudaStream_t);
cudaStream_t RTDECL(CUFGetDefaultStream)();
+int RTDECL(CUFStreamSynchronize)(cudaStream_t);
+int RTDECL(CUFStreamSynchronizeNull)();
}
} // namespace Fortran::runtime::cuda
diff --git a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
index bbc353634cd42..e5b07aa3861a2 100644
--- a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
@@ -403,6 +403,16 @@ static constexpr IntrinsicHandler cudaHandlers[]{
&CI::genCUDASetDefaultStream),
{{{"stream", asValue}}},
/*isElemental=*/false},
+ {"cudastreamsynchronize",
+ static_cast<CUDAIntrinsicLibrary::ExtendedGenerator>(
+ &CI::genCUDAStreamSynchronize),
+ {{{"stream", asValue}}},
+ /*isElemental=*/false},
+ {"cudastreamsynchronizenull",
+ static_cast<CUDAIntrinsicLibrary::ElementalGenerator>(
+ &CI::genCUDAStreamSynchronizeNull),
+ {},
+ /*isElemental=*/false},
{"fence_proxy_async",
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(
&CI::genFenceProxyAsync),
@@ -1161,6 +1171,32 @@ fir::ExtendedValue CUDAIntrinsicLibrary::genCUDASetDefaultStreamArray(
return call.getResult(0);
}
+// CUDASTREAMSYNCHRONIZE
+fir::ExtendedValue CUDAIntrinsicLibrary::genCUDAStreamSynchronize(
+ mlir::Type resTy, llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 1);
+ mlir::Value stream = fir::getBase(args[0]);
+ mlir::Type i64Ty = builder.getI64Type();
+ auto ctx = builder.getContext();
+ mlir::FunctionType ftype = mlir::FunctionType::get(ctx, {i64Ty}, {resTy});
+ auto funcOp =
+ builder.createFunction(loc, RTNAME_STRING(CUFStreamSynchronize), ftype);
+ auto call = fir::CallOp::create(builder, loc, funcOp, {stream});
+ return call.getResult(0);
+}
+
+// CUDASTREAMSYNCHRONIZENULL
+mlir::Value CUDAIntrinsicLibrary::genCUDAStreamSynchronizeNull(
+ mlir::Type resTy, llvm::ArrayRef<mlir::Value> args) {
+ assert(args.size() == 0);
+ auto ctx = builder.getContext();
+ mlir::FunctionType ftype = mlir::FunctionType::get(ctx, {}, {resTy});
+ auto funcOp = builder.createFunction(
+ loc, RTNAME_STRING(CUFStreamSynchronizeNull), ftype);
+ auto call = fir::CallOp::create(builder, loc, funcOp, {});
+ return call.getResult(0);
+}
+
// CUDAGETDEFAULTSTREAMARG
fir::ExtendedValue CUDAIntrinsicLibrary::genCUDAGetDefaultStreamArg(
mlir::Type resultType, llvm::ArrayRef<fir::ExtendedValue> args) {
diff --git a/flang/module/cuda_runtime_api.f90 b/flang/module/cuda_runtime_api.f90
index 7c6968cabc373..437bb53d8b27d 100644
--- a/flang/module/cuda_runtime_api.f90
+++ b/flang/module/cuda_runtime_api.f90
@@ -36,4 +36,14 @@ integer function cudasetstreamarray(devptr, stream)
end function
end interface
+interface cudaStreamSynchronize
+ integer function cudastreamsynchronize(stream)
+ import cuda_stream_kind
+ !DIR$ IGNORE_TKR (K) stream
+ integer(kind=cuda_stream_kind), value :: stream
+ end function
+ integer function cudastreamsynchronizenull()
+ end function
+end interface
+
end module cuda_runtime_api
diff --git a/flang/test/Lower/CUDA/cuda-default-stream.cuf b/flang/test/Lower/CUDA/cuda-default-stream.cuf
index beacb409f44f6..5bc78af22e84b 100644
--- a/flang/test/Lower/CUDA/cuda-default-stream.cuf
+++ b/flang/test/Lower/CUDA/cuda-default-stream.cuf
@@ -38,3 +38,15 @@ end subroutine
! CHECK: %{{.*}} = fir.call @_FortranACUFSetDefaultStream(%{{.*}}) fastmath<contract> : (i64) -> i32
! CHECK: %{{.*}} = fir.call @_FortranACUFGetDefaultStream() fastmath<contract> : () -> i64
! CHECK: %{{.*}} = fir.call @_FortranACUFGetDefaultStream() fastmath<contract> : () -> i64
+
+subroutine stream_synchronize
+ use cuda_runtime_api
+ integer(kind=cuda_stream_kind) :: strm
+ integer :: istat
+ istat = cudastreamsynchronize(strm)
+ istat = cudastreamsynchronize()
+end subroutine
+
+! CHECK-LABEL: func.func @_QPstream_synchronize()
+! CHECK: %{{.*}} = fir.call @_FortranACUFStreamSynchronize(%{{.*}}) fastmath<contract> : (i64) -> i32
+! CHECK: %{{.*}} = fir.call @_FortranACUFStreamSynchronizeNull() fastmath<contract> : () -> i32
More information about the flang-commits
mailing list