[flang-commits] [flang] [llvm] [flang-rt][CUF] Mirror the execution environment to the device at CUFInit (PR #224491)

[Eugene Epshteyn via flang-commits]

https://github.com/eugeneepshteyn created https://github.com/llvm/llvm-project/pull/224491

The device copy of `executionEnvironment` is statically initialized and never configured: `ExecutionEnvironment::Configure()` — which parses `FLANG_RT_COPYOUT_MODIFIED_ONLY`, `FORT_CHECK_POINTER_DEALLOCATION`, `FORT_NO_EMPTY_ALLOCATION`, `FLANG_RT_DEBUG`, … — runs only on the host. Offload-compiled runtime code reads those fields on the device (`assign.cpp`: `copyOutModifiedOnly`; `pointer.cpp`: `checkPointerDeallocation`; `descriptor.cpp`: `noEmptyAllocation`; `work-queue.cpp`: `internalDebugging`), so on the device every one of these environment variables is silently ignored and the defaults are frozen in.

This change mirrors the host-side `executionEnvironment` — already configured, since lowering generates the `CUFInit` call after `ProgramStart` — into the device image's copy at `CUFInit` time, making the documented environment variables take effect in device code the same way they do on the host.

Mechanics:

* `cudaGetSymbolAddress` resolves the device copy through the CUDA runtime's symbol registration. When the program links no device-side flang-rt (host-only runtime build, or no device code), the symbol is not registered; that case is detected (`cudaErrorInvalidSymbol`) and treated as "nothing to sync", with the sticky error cleared.
* The whole struct is copied. The pointer-valued members (`argv`, `envp`) are meaningless on the device but are never dereferenced there; copying the whole object keeps the mirror maintenance-free as fields are added.
* The copy targets the current device, consistent with the rest of `CUFInit` (`cudaDeviceSetLimit`).

Also documents the behavior in `flang/docs/RuntimeEnvironment.md`.

Follow-up to the fix that compiles `ShallowCopyModifiedSuffix` for the device: with this PR, `FLANG_RT_COPYOUT_MODIFIED_ONLY=0` restores the unconditional copy-out on the device too, instead of being ignored.

Testing: CUDA offload build (`FLANG_RT_EXPERIMENTAL_OFFLOAD_SUPPORT=CUDA`, nvcc, sm_90) and the default host build compile cleanly. Runtime verification of the mirror requires a device-linked Fortran main; offload runtimes that load the device image through the CUDA driver API instead of the CUDA runtime (so `CUFInit`'s registration-based lookup cannot see the symbol) need to perform the equivalent mirror themselves at module-load time via `cuModuleGetGlobal`.


>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] [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