[flang-commits] [flang] b0f5017 - [flang][cuda][NFC] Move CUFDeviceIsActive function to the right place (#223453)
via flang-commits
flang-commits at lists.llvm.org
Tue Sep 15 11:09:53 PDT 2026
Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-09-15T11:09:47-07:00
New Revision: b0f50174f52d02eb68a1194ed93ec89507ef2fe1
URL: https://github.com/llvm/llvm-project/commit/b0f50174f52d02eb68a1194ed93ec89507ef2fe1
DIFF: https://github.com/llvm/llvm-project/commit/b0f50174f52d02eb68a1194ed93ec89507ef2fe1.diff
LOG: [flang][cuda][NFC] Move CUFDeviceIsActive function to the right place (#223453)
The definition and declaration were done in different files (descriptor
and allocator). Move all to allocatable as this check is used for the
automatic deallocation.
Added:
Modified:
flang-rt/lib/cuda/allocatable.cpp
flang-rt/lib/cuda/allocator.cpp
flang-rt/unittests/Runtime/CUDA/Allocatable.cpp
flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp
flang/include/flang/Optimizer/Builder/Runtime/CUDA/Descriptor.h
flang/include/flang/Optimizer/Builder/Runtime/CUDA/Support.h
flang/include/flang/Runtime/CUDA/allocatable.h
flang/include/flang/Runtime/CUDA/descriptor.h
flang/lib/Optimizer/Builder/Runtime/CUDA/Descriptor.cpp
flang/lib/Optimizer/Builder/Runtime/CUDA/Support.cpp
Removed:
################################################################################
diff --git a/flang-rt/lib/cuda/allocatable.cpp b/flang-rt/lib/cuda/allocatable.cpp
index 0a7828f8016d5..3e16dda7c3e45 100644
--- a/flang-rt/lib/cuda/allocatable.cpp
+++ b/flang-rt/lib/cuda/allocatable.cpp
@@ -16,10 +16,64 @@
#include "flang/Runtime/CUDA/memmove-function.h"
#include "flang/Runtime/allocatable.h"
+#include "cuda.h"
#include "cuda_runtime.h"
namespace Fortran::runtime::cuda {
+static bool deviceContextTornDown() {
+ // Keep cudaGetLastError transparent: consume probe-only sticky errors when
+ // the slot started clean, never discarding a pre-existing user error.
+ cudaError_t priorErr{cudaPeekAtLastError()};
+ // Prefer cleanup when state cannot be proven torn down (avoids leaks).
+ bool tornDown{false};
+ int device{0};
+ if (cudaGetDevice(&device) == cudaSuccess) {
+ // Driver API reports primary-context state without lazily creating one;
+ // resolve via cudart to avoid a libcuda link (current device only).
+ using GetStateFn = CUresult(CUDAAPI *)(CUdevice, unsigned *, int *);
+ static GetStateFn getState{[]() -> GetStateFn {
+ void *fn{nullptr};
+ // Prefer ByVersion(driver): unversioned lookup uses the runtime version
+ // and fails when the runtime is newer than the driver.
+ int driverVersion{0};
+ if (cudaDriverGetVersion(&driverVersion) == cudaSuccess &&
+ cudaGetDriverEntryPointByVersion("cuDevicePrimaryCtxGetState", &fn,
+ static_cast<unsigned>(driverVersion), cudaEnableDefault,
+ nullptr) == cudaSuccess &&
+ fn) {
+ return reinterpret_cast<GetStateFn>(fn);
+ }
+ if (cudaGetDriverEntryPoint("cuDevicePrimaryCtxGetState", &fn,
+ cudaEnableDefault, nullptr) == cudaSuccess &&
+ fn) {
+ return reinterpret_cast<GetStateFn>(fn);
+ }
+ return nullptr;
+ }()};
+ if (getState) {
+ unsigned flags{0};
+ int active{0};
+ if (getState(device, &flags, &active) == CUDA_SUCCESS) {
+ tornDown = active == 0;
+ // A sticky error (e.g. an illegal kernel memory access) leaves the
+ // primary context active but unusable: later calls all fail, so
+ // scope-exit frees would abort an otherwise successful program. A
+ // null free is a no-op that surfaces this without creating a context.
+ if (!tornDown && cudaFree(nullptr) != cudaSuccess) {
+ tornDown = true;
+ }
+ }
+ }
+ } else {
+ tornDown = true;
+ }
+ if (priorErr == cudaSuccess && cudaPeekAtLastError() != cudaSuccess) {
+ (void)cudaGetLastError();
+ }
+ return tornDown;
+}
+
extern "C" {
RT_EXT_API_GROUP_BEGIN
@@ -109,6 +163,8 @@ int RTDEF(CUFAllocatableDeallocate)(Descriptor &desc, bool hasStat,
RT_EXT_API_GROUP_END
+bool RTDEF(CUFDeviceIsActive)() { return !deviceContextTornDown(); }
+
} // extern "C"
} // namespace Fortran::runtime::cuda
diff --git a/flang-rt/lib/cuda/allocator.cpp b/flang-rt/lib/cuda/allocator.cpp
index 410a175e09056..816659be35883 100644
--- a/flang-rt/lib/cuda/allocator.cpp
+++ b/flang-rt/lib/cuda/allocator.cpp
@@ -24,59 +24,6 @@
namespace Fortran::runtime::cuda {
-static bool deviceContextTornDown() {
- // Keep cudaGetLastError transparent: consume probe-only sticky errors when
- // the slot started clean, never discarding a pre-existing user error.
- cudaError_t priorErr{cudaPeekAtLastError()};
- // Prefer cleanup when state cannot be proven torn down (avoids leaks).
- bool tornDown{false};
- int device{0};
- if (cudaGetDevice(&device) == cudaSuccess) {
- // Driver API reports primary-context state without lazily creating one;
- // resolve via cudart to avoid a libcuda link (current device only).
- using GetStateFn = CUresult(CUDAAPI *)(CUdevice, unsigned *, int *);
- static GetStateFn getState{[]() -> GetStateFn {
- void *fn{nullptr};
- // Prefer ByVersion(driver): unversioned lookup uses the runtime version
- // and fails when the runtime is newer than the driver.
- int driverVersion{0};
- if (cudaDriverGetVersion(&driverVersion) == cudaSuccess &&
- cudaGetDriverEntryPointByVersion("cuDevicePrimaryCtxGetState", &fn,
- static_cast<unsigned>(driverVersion), cudaEnableDefault,
- nullptr) == cudaSuccess &&
- fn) {
- return reinterpret_cast<GetStateFn>(fn);
- }
- if (cudaGetDriverEntryPoint("cuDevicePrimaryCtxGetState", &fn,
- cudaEnableDefault, nullptr) == cudaSuccess &&
- fn) {
- return reinterpret_cast<GetStateFn>(fn);
- }
- return nullptr;
- }()};
- if (getState) {
- unsigned flags{0};
- int active{0};
- if (getState(device, &flags, &active) == CUDA_SUCCESS) {
- tornDown = active == 0;
- // A sticky error (e.g. an illegal kernel memory access) leaves the
- // primary context active but unusable: later calls all fail, so
- // scope-exit frees would abort an otherwise successful program. A
- // null free is a no-op that surfaces this without creating a context.
- if (!tornDown && cudaFree(nullptr) != cudaSuccess) {
- tornDown = true;
- }
- }
- }
- } else {
- tornDown = true;
- }
- if (priorErr == cudaSuccess && cudaPeekAtLastError() != cudaSuccess) {
- (void)cudaGetLastError();
- }
- return tornDown;
-}
-
struct DeviceAllocation {
void *ptr;
std::size_t size;
@@ -207,8 +154,6 @@ void RTDEF(CUFRegisterAllocator)() {
kUnifiedAllocatorPos, {&CUFAllocUnified, CUFFreeUnified});
}
-bool RTDEF(CUFDeviceIsActive)() { return !deviceContextTornDown(); }
-
cudaStream_t RTDECL(CUFGetAssociatedStream)(void *p) {
int pos = findAsyncDeviceAllocation(p);
if (pos >= 0) {
diff --git a/flang-rt/unittests/Runtime/CUDA/Allocatable.cpp b/flang-rt/unittests/Runtime/CUDA/Allocatable.cpp
index 9ca1bac3ec8f1..66e2bef3c1625 100644
--- a/flang-rt/unittests/Runtime/CUDA/Allocatable.cpp
+++ b/flang-rt/unittests/Runtime/CUDA/Allocatable.cpp
@@ -13,6 +13,7 @@
#include "flang-rt/runtime/descriptor.h"
#include "flang-rt/runtime/stat.h"
#include "flang-rt/runtime/terminator.h"
+#include "flang/Runtime/CUDA/allocatable.h"
#include "flang/Runtime/CUDA/allocator.h"
#include "flang/Runtime/CUDA/common.h"
#include "flang/Runtime/CUDA/descriptor.h"
@@ -244,3 +245,11 @@ TEST(AllocatableAsyncTest, DestroyStreamTest) {
cudaDeviceSynchronize();
EXPECT_EQ(cudaSuccess, cudaGetLastError());
}
+
+TEST(AllocatableCUFTest, DeviceIsActiveKeepsLastErrorClean) {
+ // CUFDeviceIsActive() probes primary-context state (including a version-
+ // skew fallback). It must not leave a sticky cudaGetLastError behind.
+ (void)cudaGetLastError(); // start from a clean error state
+ (void)RTNAME(CUFDeviceIsActive)();
+ EXPECT_EQ(cudaGetLastError(), cudaSuccess);
+}
diff --git a/flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp b/flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp
index cf0de6f25b4d6..f1f931e87a86e 100644
--- a/flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp
+++ b/flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp
@@ -72,11 +72,3 @@ TEST(AllocatableCUFTest, DescriptorAllocationTest) {
EXPECT_TRUE(desc != nullptr);
RTNAME(CUFFreeDescriptor)(desc);
}
-
-TEST(AllocatableCUFTest, DeviceIsActiveKeepsLastErrorClean) {
- // CUFDeviceIsActive() probes primary-context state (including a version-
- // skew fallback). It must not leave a sticky cudaGetLastError behind.
- (void)cudaGetLastError(); // start from a clean error state
- (void)RTNAME(CUFDeviceIsActive)();
- EXPECT_EQ(cudaGetLastError(), cudaSuccess);
-}
diff --git a/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Descriptor.h b/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Descriptor.h
index 84061e658da41..bdeb7574012c6 100644
--- a/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Descriptor.h
+++ b/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Descriptor.h
@@ -31,10 +31,6 @@ void genSyncGlobalDescriptor(fir::FirOpBuilder &builder, mlir::Location loc,
void genDescriptorCheckSection(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value desc);
-/// Generate a call returning (as i1) whether the device's primary context is
-/// alive, to guard scope-exit frees against a user cudaDeviceReset().
-mlir::Value genDeviceIsActive(fir::FirOpBuilder &builder, mlir::Location loc);
-
} // namespace fir::runtime::cuda
#endif // FORTRAN_OPTIMIZER_BUILDER_RUNTIME_CUDA_DESCRIPTOR_H_
diff --git a/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Support.h b/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Support.h
index e74773460c5e1..df1c90eb53515 100644
--- a/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Support.h
+++ b/flang/include/flang/Optimizer/Builder/Runtime/CUDA/Support.h
@@ -11,6 +11,7 @@
namespace mlir {
class Location;
+class Value;
} // namespace mlir
namespace fir {
@@ -22,6 +23,10 @@ namespace fir::runtime::cuda {
/// Generate runtime call to synchronize the CUDA device.
void genCUDADeviceSynchronize(fir::FirOpBuilder &builder, mlir::Location loc);
+/// Generate a call returning (as i1) whether the device's primary context is
+/// alive, to guard scope-exit frees against a user cudaDeviceReset().
+mlir::Value genDeviceIsActive(fir::FirOpBuilder &builder, mlir::Location loc);
+
} // namespace fir::runtime::cuda
#endif // FORTRAN_OPTIMIZER_BUILDER_RUNTIME_CUDA_SUPPORT_H_
diff --git a/flang/include/flang/Runtime/CUDA/allocatable.h b/flang/include/flang/Runtime/CUDA/allocatable.h
index d5a649594ae92..4cbe9d7405f5c 100644
--- a/flang/include/flang/Runtime/CUDA/allocatable.h
+++ b/flang/include/flang/Runtime/CUDA/allocatable.h
@@ -51,6 +51,10 @@ int RTDECL(CUFAllocatableDeallocate)(Descriptor &, bool hasStat = false,
const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
int sourceLine = 0);
+/// True unless the device's primary context has been torn down (e.g. by a user
+/// cudaDeviceReset()); used to skip scope-exit device frees safely.
+bool RTDECL(CUFDeviceIsActive)();
+
} // extern "C"
} // namespace Fortran::runtime::cuda
diff --git a/flang/include/flang/Runtime/CUDA/descriptor.h b/flang/include/flang/Runtime/CUDA/descriptor.h
index be4c760b649c8..06e4a4649db1b 100644
--- a/flang/include/flang/Runtime/CUDA/descriptor.h
+++ b/flang/include/flang/Runtime/CUDA/descriptor.h
@@ -41,10 +41,6 @@ void RTDECL(CUFSyncGlobalDescriptor)(
void RTDECL(CUFDescriptorCheckSection)(
const Descriptor *, const char *sourceFile = nullptr, int sourceLine = 0);
-/// True unless the device's primary context has been torn down (e.g. by a user
-/// cudaDeviceReset()); used to skip scope-exit device frees safely.
-bool RTDECL(CUFDeviceIsActive)();
-
} // extern "C"
} // namespace Fortran::runtime::cuda
diff --git a/flang/lib/Optimizer/Builder/Runtime/CUDA/Descriptor.cpp b/flang/lib/Optimizer/Builder/Runtime/CUDA/Descriptor.cpp
index 4e61a756aed7d..8f03c037d22de 100644
--- a/flang/lib/Optimizer/Builder/Runtime/CUDA/Descriptor.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/CUDA/Descriptor.cpp
@@ -1,5 +1,4 @@
-
-//===-- Allocatable.cpp -- Allocatable statements lowering ----------------===//
+//===-- Descriptor.cpp -- descriptor related runtime calls ----------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -47,11 +46,3 @@ void fir::runtime::cuda::genDescriptorCheckSection(fir::FirOpBuilder &builder,
builder, loc, fTy, desc, sourceFile, sourceLine)};
fir::CallOp::create(builder, loc, func, args);
}
-
-mlir::Value fir::runtime::cuda::genDeviceIsActive(fir::FirOpBuilder &builder,
- mlir::Location loc) {
- mlir::func::FuncOp func =
- fir::runtime::getRuntimeFunc<mkRTKey(CUFDeviceIsActive)>(loc, builder);
- auto call = fir::CallOp::create(builder, loc, func, mlir::ValueRange{});
- return builder.createConvert(loc, builder.getI1Type(), call.getResult(0));
-}
diff --git a/flang/lib/Optimizer/Builder/Runtime/CUDA/Support.cpp b/flang/lib/Optimizer/Builder/Runtime/CUDA/Support.cpp
index 7ad55afc0ba80..eff0068ebb634 100644
--- a/flang/lib/Optimizer/Builder/Runtime/CUDA/Support.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/CUDA/Support.cpp
@@ -9,12 +9,13 @@
#include "flang/Optimizer/Builder/Runtime/CUDA/Support.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/Runtime/RTBuilder.h"
-
-using namespace fir::runtime::cuda;
+#include "flang/Runtime/CUDA/allocatable.h"
static constexpr llvm::StringRef kCudaDeviceSynchronizeName =
"_QPcudadevicesynchronize";
+using namespace Fortran::runtime::cuda;
+
void fir::runtime::cuda::genCUDADeviceSynchronize(fir::FirOpBuilder &builder,
mlir::Location loc) {
mlir::func::FuncOp func =
@@ -33,3 +34,11 @@ void fir::runtime::cuda::genCUDADeviceSynchronize(fir::FirOpBuilder &builder,
call.setProcedureAttrsAttr(fir::FortranProcedureFlagsEnumAttr::get(
builder.getContext(), fir::FortranProcedureFlagsEnum::intrinsic));
}
+
+mlir::Value fir::runtime::cuda::genDeviceIsActive(fir::FirOpBuilder &builder,
+ mlir::Location loc) {
+ mlir::func::FuncOp func =
+ fir::runtime::getRuntimeFunc<mkRTKey(CUFDeviceIsActive)>(loc, builder);
+ auto call = fir::CallOp::create(builder, loc, func, mlir::ValueRange{});
+ return builder.createConvert(loc, builder.getI1Type(), call.getResult(0));
+}
More information about the flang-commits
mailing list