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

[via flang-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir-openacc

Author: nvptm

<details>
<summary>Changes</summary>

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

---
Full diff: https://github.com/llvm/llvm-project/pull/210757.diff


3 Files Affected:

- (added) flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir (+23) 
- (modified) mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td (+4) 
- (modified) mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp (+19-2) 


``````````diff
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 f815245882630..463cd5a5b8043 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 only when the existing GPU symbol is structurally equivalent to
@@ -104,19 +119,21 @@ class ACCDeclareGPUModuleInsertion
         // op type or different definition).
         if (existing->getName() != globalOp.getName() ||
             !OperationEquivalence::isEquivalentTo(
-                existing, &globalOp,
+                existing, deviceGlobal,
                 OperationEquivalence::ignoreValueEquivalence,
                 /*markEquivalent=*/nullptr,
                 OperationEquivalence::IgnoreLocations)) {
+          deviceGlobal->destroy();
           accSupport.emitNYI(globalOp.getLoc(),
                              llvm::Twine("duplicate global symbol '") +
                                  name.getValue() + "' in gpu module");
           return failure();
         }
+        deviceGlobal->destroy();
         continue;
       }
 
-      gpuSymTable.insert(globalOp.clone());
+      gpuSymTable.insert(deviceGlobal);
     }
     return success();
   }

``````````

</details>


https://github.com/llvm/llvm-project/pull/210757