[flang-commits] [flang] b310b1a - [flang][cuda] Register device/constant globals as device-resident under -gpu=mem:unified (#208336)

[via flang-commits]

Author: Matsu
Date: 2026-07-08T15:48:30-07:00
New Revision: b310b1ab76f24a0e7a88c8ded76f54b816e5d542

URL: https://github.com/llvm/llvm-project/commit/b310b1ab76f24a0e7a88c8ded76f54b816e5d542
DIFF: https://github.com/llvm/llvm-project/commit/b310b1ab76f24a0e7a88c8ded76f54b816e5d542.diff

LOG: [flang][cuda] Register device/constant globals as device-resident under -gpu=mem:unified (#208336)

```fortran
module m
  integer, device   :: aaa;  bind(c, name='aaa_from_c') :: aaa
  integer, constant :: zzz;  bind(c, name='zzz_from_c') :: zzz
end module
```
```c
extern int aaa_from_c;
#pragma acc declare create(aaa_from_c)   // same for zzz_from_c
```

In this code, under `-gpu=mem:unified` a module-scope device/constant
global is registered only as a CUDA variable. Another translation unit
that declares the same symbol on the host can make it be treated as host
memory, so the device symbol is unresolved at run time.

Fix: emit `cuf.register_variable_static` for device/constant globals
under `mem:unified` so they are additionally registered as
device-resident.

Added: 
    

Modified: 
    flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
    flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
    flang/test/Fir/CUDA/cuda-constructor-2.f90

Removed: 
    


################################################################################
diff  --git a/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td b/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
index d5bf5503d3ced..e08f9c5e23cde 100644
--- a/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
+++ b/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
@@ -334,6 +334,26 @@ def cuf_RegisterKernelOp : cuf_Op<"register_kernel", []> {
   }];
 }
 
+def cuf_RegisterVariableStaticOp : cuf_Op<"register_variable_static", []> {
+  let summary = "Register a device or constant global as device-resident";
+
+  let description = [{
+    Emitted under -gpu=mem:unified for a module-scope device or constant global.
+    Registers the global with the runtime as device-resident so that a matching
+    host-side declaration of the same symbol is not treated as host memory.
+  }];
+
+  let arguments = (ins
+    SymbolRefAttr:$global,
+    StrAttr:$varName,
+    I64Attr:$size
+  );
+
+  let assemblyFormat = [{
+    $global `(` $varName `,` $size `)` attr-dict
+  }];
+}
+
 def cuf_DeviceAddressOp : cuf_Op<"device_address", []> {
   let summary = "Get the device address from a host symbol";
 

diff  --git a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
index 5e8136af7991e..97f1c0c8762e8 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
@@ -124,6 +124,26 @@ static mlir::Value computeGlobalSize(fir::FirOpBuilder &builder,
   return builder.createIntegerConstant(loc, idxTy, *size);
 }
 
+/// Storage size in bytes of \p globalOp as a raw integer (see computeGlobalSize
+/// for the box vs declared-type handling).
+static uint64_t getGlobalSizeInBytes(mlir::Location loc,
+                                     const mlir::DataLayout &dl,
+                                     const fir::KindMapping &kindMap,
+                                     fir::LLVMTypeConverter &typeConverter,
+                                     fir::GlobalOp globalOp) {
+  std::optional<uint64_t> size;
+  if (auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(globalOp.getType())) {
+    mlir::Type structTy = typeConverter.convertBoxTypeAsStruct(boxTy);
+    size = dl.getTypeSizeInBits(structTy) / 8;
+  }
+  if (!size) {
+    size = fir::getTypeSizeAndAlignmentOrCrash(loc, globalOp.getType(), dl,
+                                               kindMap)
+               .first;
+  }
+  return *size;
+}
+
 /// Emit a call to a CUF registration runtime function with the canonical
 /// (module, addr, name, size) signature, where addr is the address of \p
 /// addrGlobal taken via fir.address_of and name/size describe \p nameGlobal.
@@ -288,6 +308,18 @@ struct CUFAddConstructor
                                       typeConverter, registeredMod, func,
                                       /*addrGlobal=*/globalOp,
                                       /*nameGlobal=*/globalOp);
+              // Under -gpu=mem:unified, also register the global as
+              // device-resident so a matching host symbol from another
+              // translation unit is not treated as host memory.
+              if (cudaUnified) {
+                uint64_t szBytes = getGlobalSizeInBytes(
+                    loc, *dl, kindMap, typeConverter, globalOp);
+                cuf::RegisterVariableStaticOp::create(
+                    builder, loc,
+                    mlir::SymbolRefAttr::get(ctx, globalOp.getSymName()),
+                    builder.getStringAttr(globalOp.getSymName()),
+                    builder.getI64IntegerAttr(szBytes));
+              }
             }
           } break;
           default:

diff  --git a/flang/test/Fir/CUDA/cuda-constructor-2.f90 b/flang/test/Fir/CUDA/cuda-constructor-2.f90
index bb2a98c294558..26ebcf73d6dbe 100644
--- a/flang/test/Fir/CUDA/cuda-constructor-2.f90
+++ b/flang/test/Fir/CUDA/cuda-constructor-2.f90
@@ -43,6 +43,10 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<!llvm.ptr, dense<
 // CHECK-DAG: %[[BOX:.*]] = fir.address_of(@_QMmtestsEndev) : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
 // CHECK-DAG: %[[BOXREF:.*]] = fir.convert %[[BOX]] : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.ref<i8>
 // CHECK-DAG: fir.call @_FortranACUFRegisterVariable(%[[MODULE:.*]], %[[BOXREF]], %{{.*}}, %{{.*}})
+// Under -gpu=mem:unified, device/constant globals are also registered as
+// device-resident so a matching host symbol is not treated as host memory.
+// UNIFIED-DAG: cuf.register_variable_static @_QMmtestsEn("_QMmtestsEn", 20)
+// UNIFIED-DAG: cuf.register_variable_static @_QMmtestsEndev("_QMmtestsEndev",
 // CHECK-NOT: fir.call @_FortranACUFInitModule
 
 // -----