[flang-commits] [flang] [llvm] [flang-rt][CUF] Mirror the execution environment to the device at CUFInit (PR #224491)
Eugene Epshteyn via flang-commits
flang-commits at lists.llvm.org
Fri Sep 18 06:20:28 PDT 2026
https://github.com/eugeneepshteyn updated https://github.com/llvm/llvm-project/pull/224491
>From 033c8567afff7a8851faa0b11428ceead9e58abc Mon Sep 17 00:00:00 2001
From: Eugene Epshteyn <eepshteyn at nvidia.com>
Date: Thu, 17 Sep 2026 18:50:16 -0700
Subject: [PATCH 1/2] [flang-rt][CUF] Mirror the execution environment to the
device at CUFInit
The device copy of executionEnvironment is statically initialized and
never configured: ExecutionEnvironment::Configure() runs only on the
host, so every environment variable it parses is silently ignored by
runtime code executing on the device. Offload-compiled runtime code
does read these fields on the device: copyOutModifiedOnly (assign.cpp),
checkPointerDeallocation (pointer.cpp), noEmptyAllocation
(descriptor.cpp), and internalDebugging (work-queue.cpp) all stay
frozen at their defaults regardless of FLANG_RT_COPYOUT_MODIFIED_ONLY,
FORT_CHECK_POINTER_DEALLOCATION, FORT_NO_EMPTY_ALLOCATION, or
FLANG_RT_DEBUG.
Mirror the host execution environment - already configured, since
lowering generates the CUFInit call after ProgramStart - into the
device image's copy at CUFInit time, so these variables have the same
effect in device code as on the host. The device copy is resolved
through the CUDA runtime's symbol registration (cudaGetSymbolAddress);
when the program links no device-side flang-rt the symbol is not
registered, which is detected and treated as nothing to sync. The
whole struct is copied: the pointer-valued members are meaningless on
the device but never dereferenced there, and a whole-object copy stays
correct as fields are added.
An offload runtime that loads the device image through the CUDA driver
API instead of the CUDA runtime cannot be reached by this mirror; it
must perform the equivalent copy itself at module-load time (e.g. via
cuModuleGetGlobal). Documented in flang/docs/RuntimeEnvironment.md.
---
flang-rt/lib/cuda/init.cpp | 19 +++++++++++++++++++
flang/docs/RuntimeEnvironment.md | 10 ++++++++++
2 files changed, 29 insertions(+)
diff --git a/flang-rt/lib/cuda/init.cpp b/flang-rt/lib/cuda/init.cpp
index d79bffc32424d..c99874705f4d6 100644
--- a/flang-rt/lib/cuda/init.cpp
+++ b/flang-rt/lib/cuda/init.cpp
@@ -21,5 +21,24 @@ void RTDEF(CUFInit)() {
CUDA_REPORT_IF_ERROR(cudaDeviceSetLimit(cudaLimitStackSize,
Fortran::runtime::executionEnvironment.cudaStackLimit));
}
+ // Mirror the host execution environment (already configured: lowering
+ // generates the CUFInit call after ProgramStart) into the device image's
+ // copy, so that the environment variables documented in
+ // flang/docs/RuntimeEnvironment.md have the same effect in runtime code
+ // compiled for the device as they have on the host. Without this, the
+ // device copy stays at its default-initialized values. When the program
+ // links no device-side flang-rt, the symbol is not registered with the
+ // CUDA runtime; treat that as nothing to sync.
+ void *devicePtr{nullptr};
+ if (cudaGetSymbolAddress(
+ &devicePtr, &Fortran::runtime::executionEnvironment) == cudaSuccess) {
+ CUDA_REPORT_IF_ERROR(
+ cudaMemcpy(devicePtr, &Fortran::runtime::executionEnvironment,
+ sizeof(Fortran::runtime::executionEnvironment),
+ cudaMemcpyHostToDevice));
+ } else {
+ // Clear the sticky cudaErrorInvalidSymbol.
+ (void)cudaGetLastError();
+ }
}
}
diff --git a/flang/docs/RuntimeEnvironment.md b/flang/docs/RuntimeEnvironment.md
index 16e75d3ab1e9c..0e584a0e31fe8 100644
--- a/flang/docs/RuntimeEnvironment.md
+++ b/flang/docs/RuntimeEnvironment.md
@@ -20,6 +20,16 @@ library.
The following environment variables can affect the behavior of
Fortran programs during execution.
+Some of them also control runtime code that executes on an offload device
+(for example, `FLANG_RT_COPYOUT_MODIFIED_ONLY` affects copy-out performed
+in device code). The runtime reads the variables on the host at program
+startup, and CUDA Fortran program initialization mirrors the resulting
+settings into the device image, so such a variable has the same effect in
+device code as on the host. An offload runtime that loads the device image
+through the CUDA driver API instead must mirror the settings itself when
+it loads the image; otherwise the device copy of the execution environment
+stays at its default values.
+
## `DEFAULT_UTF8=1`
Set `DEFAULT_UTF8` to cause formatted external input to assume UTF-8
>From 16f018be663ce294f88e512083317c3667adefcb Mon Sep 17 00:00:00 2001
From: Eugene Epshteyn <eepshteyn at nvidia.com>
Date: Fri, 18 Sep 2026 06:20:11 -0700
Subject: [PATCH 2/2] [flang-rt][CUF] Use a setter kernel for the device
environment mirror
The device copy of executionEnvironment has no host-side shadow
registration: RT_VAR_ATTRS applies __device__ only under __CUDA_ARCH__,
so the host compilation pass sees a plain host global and nvcc emits no
registration for the device variable. cudaGetSymbolAddress therefore
always fails with cudaErrorInvalidSymbol, and the mirror in the previous
commit silently did nothing in every configuration - found by runtime
verification on an sm_100 GPU.
Mirror through a channel that does not need registration: a __global__
setter kernel that receives the new value as a kernel parameter and
assigns it to the device global (resolved by the RDC device link against
the offload-built runtime). The kernel lives in a new
lib/cuda/environment.cpp, compiled as CUDA when
FLANG_RT_EXPERIMENTAL_OFFLOAD_SUPPORT=CUDA and as a plain C++ no-op
otherwise, so CUFInit can call it unconditionally.
The kernel is launched with a sanitized local copy of the host
environment: argc, argv, and envp are host memory with no meaning on
the device, so they are zeroed rather than leaked into the device
image. Everything else - including the fields device runtime code
demonstrably reads (listDirectedOutputLineLengthLimit,
defaultOutputRoundingMode, noStopMessage, defaultUTF8,
checkPointerDeallocation, truncateStream, noEmptyAllocation,
copyOutModifiedOnly, internalDebugging) - is mirrored as configured.
Verified on an sm_100 GPU with a standalone harness that replays the
CUF startup sequence (ProgramStart, CUFInit) against the offload-built
runtime and snapshots the device copy from a kernel: without the mirror
the device copy stays at defaults while the host parses
FLANG_RT_COPYOUT_MODIFIED_ONLY=0; with it, FORT_FMT_RECL=132,
NO_STOP_MESSAGE=1, DEFAULT_UTF8=1, FORT_CHECK_POINTER_DEALLOCATION=0,
FORT_NO_EMPTY_ALLOCATION=1, FLANG_RT_DEBUG=1, and
FLANG_RT_COPYOUT_MODIFIED_ONLY=0 all reach the device, and the device
copy reads argc=0, argv=envp=null.
---
flang-rt/lib/cuda/CMakeLists.txt | 14 +++++++
flang-rt/lib/cuda/environment.cpp | 53 +++++++++++++++++++++++++
flang-rt/lib/cuda/init.cpp | 17 ++------
flang/docs/RuntimeEnvironment.md | 9 +++--
flang/include/flang/Runtime/CUDA/init.h | 5 +++
5 files changed, 81 insertions(+), 17 deletions(-)
create mode 100644 flang-rt/lib/cuda/environment.cpp
diff --git a/flang-rt/lib/cuda/CMakeLists.txt b/flang-rt/lib/cuda/CMakeLists.txt
index 118fd20f39e5f..cc1d74aa5df1d 100644
--- a/flang-rt/lib/cuda/CMakeLists.txt
+++ b/flang-rt/lib/cuda/CMakeLists.txt
@@ -10,6 +10,7 @@ add_flangrt_library(flang_rt.cuda STATIC SHARED
allocatable.cpp
allocator.cpp
descriptor.cpp
+ environment.cpp
init.cpp
kernel.cpp
memmove-function.cpp
@@ -28,6 +29,19 @@ add_flangrt_library(flang_rt.cuda STATIC SHARED
INSTALL_WITH_TOOLCHAIN
)
+if (FLANG_RT_EXPERIMENTAL_OFFLOAD_SUPPORT STREQUAL "CUDA")
+ # environment.cpp launches a device kernel that mirrors the execution
+ # environment into the device image's copy; compile it as CUDA so the
+ # kernel exists. In builds without device offload it compiles as plain
+ # C++ into a no-op (see the guard in the file).
+ enable_language(CUDA)
+ set_source_files_properties(environment.cpp PROPERTIES LANGUAGE CUDA)
+ if (TARGET flang_rt.cuda.static)
+ set_target_properties(flang_rt.cuda.static
+ PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+ endif ()
+endif ()
+
# For the static library, link-in the static dependencies as well.
if (TARGET flang_rt.cuda.static)
target_link_libraries(flang_rt.cuda.static PUBLIC
diff --git a/flang-rt/lib/cuda/environment.cpp b/flang-rt/lib/cuda/environment.cpp
new file mode 100644
index 0000000000000..e6fda58caa02e
--- /dev/null
+++ b/flang-rt/lib/cuda/environment.cpp
@@ -0,0 +1,53 @@
+//===-- lib/cuda/environment.cpp --------------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "flang-rt/runtime/environment.h"
+#include "flang-rt/runtime/terminator.h"
+#include "flang/Runtime/CUDA/common.h"
+#include "flang/Runtime/CUDA/init.h"
+
+#include "cuda_runtime.h"
+
+// When Flang-RT is built with device offload support, this file is compiled
+// as CUDA (see lib/cuda/CMakeLists.txt) and mirrors the host-side execution
+// environment into the device image's copy. The device copy of
+// executionEnvironment is a plain __device__ variable with no host-side
+// shadow registration (RT_VAR_ATTRS applies __device__ only under
+// __CUDA_ARCH__), so cudaMemcpyToSymbol/cudaGetSymbolAddress cannot reach
+// it; a setter kernel taking the new value by parameter is the channel that
+// works. Compiled as plain C++ (no device offload), the sync is a no-op.
+
+#if defined(__CUDACC__) || defined(__CUDA__)
+
+namespace {
+__global__ void SetDeviceExecutionEnvironment(
+ Fortran::runtime::ExecutionEnvironment env) {
+ Fortran::runtime::executionEnvironment = env;
+}
+} // namespace
+
+extern "C" void RTDEF(CUFSyncExecutionEnvironment)() {
+ // In the host pass of this translation unit RT_VAR_ATTRS is empty, so this
+ // names the host copy, already configured by ProgramStart.
+ Fortran::runtime::ExecutionEnvironment env{
+ Fortran::runtime::executionEnvironment};
+ // The command line and the host environment table are host memory and have
+ // no meaning on the device; never leak host pointers into the device image.
+ env.argc = 0;
+ env.argv = nullptr;
+ env.envp = nullptr;
+ SetDeviceExecutionEnvironment<<<1, 1>>>(env);
+ CUDA_REPORT_IF_ERROR(cudaGetLastError());
+ CUDA_REPORT_IF_ERROR(cudaDeviceSynchronize());
+}
+
+#else // plain C++ build: no device-side flang-rt to sync
+
+extern "C" void RTDEF(CUFSyncExecutionEnvironment)() {}
+
+#endif
diff --git a/flang-rt/lib/cuda/init.cpp b/flang-rt/lib/cuda/init.cpp
index c99874705f4d6..1c87cd0c533e6 100644
--- a/flang-rt/lib/cuda/init.cpp
+++ b/flang-rt/lib/cuda/init.cpp
@@ -26,19 +26,8 @@ void RTDEF(CUFInit)() {
// copy, so that the environment variables documented in
// flang/docs/RuntimeEnvironment.md have the same effect in runtime code
// compiled for the device as they have on the host. Without this, the
- // device copy stays at its default-initialized values. When the program
- // links no device-side flang-rt, the symbol is not registered with the
- // CUDA runtime; treat that as nothing to sync.
- void *devicePtr{nullptr};
- if (cudaGetSymbolAddress(
- &devicePtr, &Fortran::runtime::executionEnvironment) == cudaSuccess) {
- CUDA_REPORT_IF_ERROR(
- cudaMemcpy(devicePtr, &Fortran::runtime::executionEnvironment,
- sizeof(Fortran::runtime::executionEnvironment),
- cudaMemcpyHostToDevice));
- } else {
- // Clear the sticky cudaErrorInvalidSymbol.
- (void)cudaGetLastError();
- }
+ // device copy stays at its default-initialized values. A no-op unless
+ // Flang-RT was built with device offload support (see environment.cpp).
+ RTNAME(CUFSyncExecutionEnvironment)();
}
}
diff --git a/flang/docs/RuntimeEnvironment.md b/flang/docs/RuntimeEnvironment.md
index 0e584a0e31fe8..0daee75057309 100644
--- a/flang/docs/RuntimeEnvironment.md
+++ b/flang/docs/RuntimeEnvironment.md
@@ -23,9 +23,12 @@ Fortran programs during execution.
Some of them also control runtime code that executes on an offload device
(for example, `FLANG_RT_COPYOUT_MODIFIED_ONLY` affects copy-out performed
in device code). The runtime reads the variables on the host at program
-startup, and CUDA Fortran program initialization mirrors the resulting
-settings into the device image, so such a variable has the same effect in
-device code as on the host. An offload runtime that loads the device image
+startup, and — when Flang-RT is built with device offload support — CUDA
+Fortran program initialization mirrors the resulting settings into the
+device image, so such a variable has the same effect in device code as on
+the host. The command line and the host environment table (`argc`, `argv`,
+`envp`) are deliberately not mirrored: they are host memory and have no
+meaning on the device. An offload runtime that loads the device image
through the CUDA driver API instead must mirror the settings itself when
it loads the image; otherwise the device copy of the execution environment
stays at its default values.
diff --git a/flang/include/flang/Runtime/CUDA/init.h b/flang/include/flang/Runtime/CUDA/init.h
index d1c8fcbf7d587..4168d526000fc 100644
--- a/flang/include/flang/Runtime/CUDA/init.h
+++ b/flang/include/flang/Runtime/CUDA/init.h
@@ -14,6 +14,11 @@
extern "C" {
void RTDECL(CUFInit)();
+
+// Mirrors the host executionEnvironment (sanitized: no host argc/argv/envp)
+// into the device image's copy. Called by CUFInit; a no-op unless Flang-RT
+// was built with device offload support.
+void RTDECL(CUFSyncExecutionEnvironment)();
}
#endif // FORTRAN_RUNTIME_CUDA_INIT_H_
More information about the flang-commits
mailing list