[flang-commits] [flang] 5df4aa1 - [flang][acc] Emit unified declare globals as extern. (#210757)

[via flang-commits]

Author: nvptm
Date: 2026-07-21T12:59:01-07:00
New Revision: 5df4aa19aa1e6308509356db19b5e46884d47361

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

LOG: [flang][acc] Emit unified declare globals as extern. (#210757)

Emit external device symbols for OpenACC declare create and copyin
globals in -gpu=mem:unified; preserve device definitions for
device_resident globals.

Added: 
    flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir

Modified: 
    mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
    mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp

Removed: 
    


################################################################################
diff  --git a/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
new file mode 100644
index 0000000000000..b27094467f9cd
--- /dev/null
+++ b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
@@ -0,0 +1,23 @@
+// RUN: fir-opt %s --pass-pipeline="builtin.module(acc-declare-gpu-module-insertion{cuda-unified=true},acc-declare-gpu-module-insertion{cuda-unified=true})" | FileCheck %s
+
+// Globals registered as host variables must be declarations in device code.
+// device_resident globals remain device definitions.
+module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", gpu.container_module} {
+  fir.global internal @shared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<7xf32> {
+    %0 = fir.zero_bits !fir.array<7xf32>
+    fir.has_value %0 : !fir.array<7xf32>
+  }
+  fir.global @initialized_shared(dense<1.0> : vector<1xf32>) {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.array<1xf32>
+  fir.global @resident {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>} : !fir.array<7xf32> {
+    %0 = fir.zero_bits !fir.array<7xf32>
+    fir.has_value %0 : !fir.array<7xf32>
+  }
+}
+
+// CHECK-LABEL: gpu.module @acc_gpu_module {
+// CHECK: fir.global @shared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<7xf32>
+// CHECK-NOT: fir.has_value
+// CHECK: fir.global @initialized_shared {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.array<1xf32>
+// CHECK-NOT: fir.has_value
+// CHECK: fir.global @resident {{.*}} : !fir.array<7xf32> {
+// CHECK: fir.has_value

diff  --git a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
index 04723c5db2d12..fbc832a37f478 100644
--- a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
@@ -157,6 +157,10 @@ def ACCDeclareGPUModuleInsertion : Pass<"acc-declare-gpu-module-insertion", "mli
     that device code can reference them.
   }];
   let dependentDialects = ["mlir::acc::OpenACCDialect", "mlir::gpu::GPUDialect"];
+  let options = [
+    Option<"cudaUnified", "cuda-unified", "bool", "false",
+           "Emit declaration-only device globals for CUDA unified memory.">
+  ];
 }
 
 def ACCLegalizeSerial : Pass<"acc-legalize-serial", "mlir::func::FuncOp"> {

diff  --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
index 705b4f9df7770..51e2222208d0d 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
@@ -77,6 +77,15 @@ static bool hasAccDeclareGlobals(ModuleOp mod) {
   return false;
 }
 
+static void makeDeviceGlobalDeclaration(Operation &globalOp) {
+  globalOp.removeAttr("initVal");
+  globalOp.removeAttr("linkName");
+  for (Region &region : globalOp.getRegions()) {
+    region.dropAllReferences();
+    region.getBlocks().clear();
+  }
+}
+
 class ACCDeclareGPUModuleInsertion
     : public acc::impl::ACCDeclareGPUModuleInsertionBase<
           ACCDeclareGPUModuleInsertion> {
@@ -97,6 +106,12 @@ class ACCDeclareGPUModuleInsertion
         continue;
 
       StringAttr name = symOp.getNameAttr();
+      Operation *deviceGlobal = globalOp.clone();
+      auto declareAttr =
+          globalOp.getAttrOfType<acc::DeclareAttr>(acc::getDeclareAttrName());
+      if (cudaUnified && declareAttr.getDataClause().getValue() !=
+                             acc::DataClause::acc_declare_device_resident)
+        makeDeviceGlobalDeclaration(*deviceGlobal);
 
       if (Operation *existing = gpuSymTable.lookup(name.getValue())) {
         // Reuse when structurally equivalent ignoring locations and discardable
@@ -104,11 +119,12 @@ class ACCDeclareGPUModuleInsertion
         // true definition mismatch is a conflict.
         if (existing->getName() != globalOp.getName() ||
             !OperationEquivalence::isEquivalentTo(
-                existing, &globalOp,
+                existing, deviceGlobal,
                 OperationEquivalence::ignoreValueEquivalence,
                 /*markEquivalent=*/nullptr,
                 OperationEquivalence::IgnoreLocations |
                     OperationEquivalence::IgnoreDiscardableAttrs)) {
+          deviceGlobal->destroy();
           accSupport.emitNYI(globalOp.getLoc(),
                              llvm::Twine("duplicate global symbol '") +
                                  name.getValue() + "' in gpu module");
@@ -120,10 +136,11 @@ class ACCDeclareGPUModuleInsertion
           if (Attribute declareAttr =
                   globalOp.getAttr(acc::getDeclareAttrName()))
             existing->setAttr(acc::getDeclareAttrName(), declareAttr);
+        deviceGlobal->destroy();
         continue;
       }
 
-      gpuSymTable.insert(globalOp.clone());
+      gpuSymTable.insert(deviceGlobal);
     }
     return success();
   }