[flang-commits] [flang] [mlir] [flang][acc] Emit unified declare globals as extern. (PR #210757)
via flang-commits
flang-commits at lists.llvm.org
Mon Jul 20 09:47:40 PDT 2026
https://github.com/nvptm created https://github.com/llvm/llvm-project/pull/210757
Emit external device symbols for OpenACC declare create and copyin globals in -gpu=mem:unified; preserve device definitions for device_resident globals.
>From 3b8fb7e41b597e46bdf353e3e80d289a691829cc Mon Sep 17 00:00:00 2001
From: nvpm <pmathew at nvidia.com>
Date: Fri, 17 Jul 2026 17:05:49 -0700
Subject: [PATCH 1/3] [acc] Emit unified declare globals as extern
---
.../acc-declare-gpu-module-insertion.fir | 23 +++++++++++++++++++
.../mlir/Dialect/OpenACC/Transforms/Passes.td | 4 ++++
.../ACCDeclareGPUModuleInsertion.cpp | 17 +++++++++++++-
3 files changed, 43 insertions(+), 1 deletion(-)
create mode 100644 flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
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..07e9f0ea02da2
--- /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})" | 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 {{.*}} : !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..0b02c731e9b6f 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 ®ion : globalOp.getRegions()) {
+ region.dropAllReferences();
+ region.getBlocks().clear();
+ }
+}
+
class ACCDeclareGPUModuleInsertion
: public acc::impl::ACCDeclareGPUModuleInsertionBase<
ACCDeclareGPUModuleInsertion> {
@@ -116,7 +125,13 @@ class ACCDeclareGPUModuleInsertion
continue;
}
- gpuSymTable.insert(globalOp.clone());
+ 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);
+ gpuSymTable.insert(deviceGlobal);
}
return success();
}
>From 5567e2ba9e971d5874ce745788e89a88875194e4 Mon Sep 17 00:00:00 2001
From: nvpm <pmathew at nvidia.com>
Date: Mon, 20 Jul 2026 07:50:06 -0700
Subject: [PATCH 2/3] Make unified declare global insertion idempotent
---
.../OpenACC/acc-declare-gpu-module-insertion.fir | 4 ++--
.../Transforms/ACCDeclareGPUModuleInsertion.cpp | 16 +++++++++-------
2 files changed, 11 insertions(+), 9 deletions(-)
diff --git a/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
index 07e9f0ea02da2..b27094467f9cd 100644
--- a/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
+++ b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
@@ -1,4 +1,4 @@
-// RUN: fir-opt %s --pass-pipeline="builtin.module(acc-declare-gpu-module-insertion{cuda-unified=true})" | FileCheck %s
+// 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.
@@ -15,7 +15,7 @@ module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", gpu.conta
}
// CHECK-LABEL: gpu.module @acc_gpu_module {
-// CHECK: fir.global @shared {{.*}} : !fir.array<7xf32>
+// 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
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
index 0b02c731e9b6f..e6ba250927a7c 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
@@ -106,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
@@ -113,24 +119,20 @@ 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;
}
- 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);
gpuSymTable.insert(deviceGlobal);
}
return success();
>From 99b3e270c780c5f0c2c0c22cfaa92e4283280f2b Mon Sep 17 00:00:00 2001
From: nvpm <pmathew at nvidia.com>
Date: Mon, 20 Jul 2026 07:55:10 -0700
Subject: [PATCH 3/3] Format
---
.../Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
index e6ba250927a7c..463cd5a5b8043 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
@@ -110,7 +110,7 @@ class ACCDeclareGPUModuleInsertion
auto declareAttr =
globalOp.getAttrOfType<acc::DeclareAttr>(acc::getDeclareAttrName());
if (cudaUnified && declareAttr.getDataClause().getValue() !=
- acc::DataClause::acc_declare_device_resident)
+ acc::DataClause::acc_declare_device_resident)
makeDeviceGlobalDeclaration(*deviceGlobal);
if (Operation *existing = gpuSymTable.lookup(name.getValue())) {
More information about the flang-commits
mailing list