[flang-commits] [flang] [flang] Add a pass to get OpenACC device ptr for CUDA kernel (PR #212299)
via flang-commits
flang-commits at lists.llvm.org
Wed Jul 29 12:54:45 PDT 2026
https://github.com/yebinchon updated https://github.com/llvm/llvm-project/pull/212299
>From 768cd898457bc1b7f9829517b38845de5ea7d775 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 09:50:52 -0700
Subject: [PATCH 01/10] [flang] Add a pass to get device pointer set up by
OpenACC for CUDA kernel
---
.../include/flang/Optimizer/OpenACC/Passes.td | 28 +++
.../Transforms/ACCDevicePtrToCUFKernel.cpp | 206 ++++++++++++++++++
.../OpenACC/Transforms/CMakeLists.txt | 2 +
3 files changed, 236 insertions(+)
create mode 100644 flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
diff --git a/flang/include/flang/Optimizer/OpenACC/Passes.td b/flang/include/flang/Optimizer/OpenACC/Passes.td
index 0c726b7e5cb86..ea220c818eadd 100644
--- a/flang/include/flang/Optimizer/OpenACC/Passes.td
+++ b/flang/include/flang/Optimizer/OpenACC/Passes.td
@@ -95,4 +95,32 @@ def ACCOptimizeFirstprivateMap
let dependentDialects = ["mlir::acc::OpenACCDialect", "fir::FIROpsDialect"];
}
+def ACCDevicePtrToCUFKernel
+ : Pass<"acc-device-ptr-to-cuf-kernel", "mlir::ModuleOp"> {
+ let summary = "Pass device addresses to CUDA Fortran kernels launched inside "
+ "OpenACC data regions";
+ let description = [{
+ A CUDA Fortran kernel launched inside an OpenACC data region must receive
+ the device address of any host variable that OpenACC has made present, not
+ the host address. Otherwise the kernel dereferences a host pointer, which is
+ only valid on shared-memory/unified-addressing systems and is illegal on a
+ device with a separate address space.
+
+ For each cuf.kernel_launch whose reference arguments resolve to a variable
+ mapped by an enclosing acc.data region, this pass wraps the launch in an
+ acc.host_data construct with acc.use_device operands for the mapped
+ variables, and rebuilds the launch argument addressing on top of the
+ use_device result so the kernel receives the device pointer.
+
+ Both directly-addressed variables (e.g. static arrays, whose data address is
+ the mapped varPtr) and descriptor-based variables (e.g. allocatables and
+ pointers, whose data address is box_addr(load(<descriptor>))) are handled.
+
+ This must run before the pass that lowers cuf.kernel_launch to
+ gpu.launch_func.
+ }];
+ let dependentDialects = ["mlir::acc::OpenACCDialect", "fir::FIROpsDialect",
+ "cuf::CUFDialect"];
+}
+
#endif // FORTRAN_OPTIMIZER_OPENACC_PASSES
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
new file mode 100644
index 0000000000000..863a2490412a9
--- /dev/null
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -0,0 +1,206 @@
+//===- ACCDevicePtrToCUFKernel.cpp --------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// A CUDA Fortran kernel launched inside an OpenACC data
+// region must receive the device address of any host variable that OpenACC has
+// made present, not the host address. This pass wraps each
+// cuf.kernel_launch in an acc.host_data construct with acc.use_device operands
+// for the mapped host variables, and rebuilds the launch's argument addressing
+// on top of the use_device result. The host_data/use_device lowering then
+// materializes the present-table device pointer, and any array-section
+// addressing is recomputed on the device pointer.
+//
+//===----------------------------------------------------------------------===//
+
+#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
+#include "flang/Optimizer/Dialect/FIRDialect.h"
+#include "flang/Optimizer/Dialect/FIROps.h"
+#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/HLFIR/HLFIROps.h"
+#include "flang/Optimizer/OpenACC/Passes.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/IRMapping.h"
+#include "mlir/Pass/Pass.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SetVector.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace fir {
+namespace acc {
+#define GEN_PASS_DEF_ACCDEVICEPTRTOCUFKERNEL
+#include "flang/Optimizer/OpenACC/Passes.h.inc"
+} // namespace acc
+} // namespace fir
+
+using namespace mlir;
+
+namespace {
+
+/// Walk down an addressing chain to the underlying variable that OpenACC maps
+/// as a data-clause pointer
+static Value getMappedVar(Value value) {
+ while (true) {
+ if (auto convert = value.getDefiningOp<fir::ConvertOp>()) {
+ value = convert.getValue();
+ continue;
+ }
+ if (auto coor = value.getDefiningOp<fir::ArrayCoorOp>()) {
+ value = coor.getMemref();
+ continue;
+ }
+ if (auto coor = value.getDefiningOp<fir::CoordinateOp>()) {
+ value = coor.getRef();
+ continue;
+ }
+ if (auto designate = value.getDefiningOp<hlfir::DesignateOp>()) {
+ value = designate.getMemref();
+ continue;
+ }
+ // Descriptor-based (allocatable/pointer) variables: the data address is
+ // extracted from the descriptor via box_addr(load(<descriptor ref>)). Peel
+ // both so the walk reaches the descriptor variable, which is what OpenACC
+ // maps as the data-clause varPtr for such variables.
+ if (auto boxAddr = value.getDefiningOp<fir::BoxAddrOp>()) {
+ value = boxAddr.getVal();
+ continue;
+ }
+ if (auto load = value.getDefiningOp<fir::LoadOp>()) {
+ // Only a load that produces a descriptor is part of the addressing
+ // chain; scalar loads are ordinary values, not addressing steps.
+ if (mlir::isa<fir::BaseBoxType>(load.getType())) {
+ value = load.getMemref();
+ continue;
+ }
+ }
+ if (isa_and_nonnull<fir::DeclareOp, hlfir::DeclareOp>(
+ value.getDefiningOp()))
+ return value;
+ return {};
+ }
+}
+
+/// Checks if mappedVar is present due to an enclosing acc.data region.
+static bool isMappedInEnclosingAccData(Value mappedVar,
+ cuf::KernelLaunchOp launch) {
+ if(!mappedVar)
+ return false;
+ for (auto dataOp = launch->getParentOfType<acc::DataOp>(); dataOp;
+ dataOp = dataOp->getParentOfType<acc::DataOp>()) {
+ for (Value dataOperand : dataOp.getDataClauseOperands()) {
+ if (Value hostVar = acc::getVar(dataOperand.getDefiningOp()))
+ if (getMappedVar(hostVar) == mappedVar)
+ return true;
+ }
+ }
+ return false;
+}
+
+/// Reconstructs the addressing chain that produced `value` from `mappedVar`,
+/// substituting `deviceVar` for `mappedVar`. Only addressing ops are cloned;
+/// everything else (constants, shapes, ...) is reused as a live-in. New ops are
+/// created at `builder`'s current insertion point.
+static Value rebuildOnDevice(OpBuilder &builder, Value value, Value mappedVar,
+ Value deviceVar) {
+ if (value == mappedVar)
+ return deviceVar;
+
+ Operation *def = value.getDefiningOp();
+ if (!def || !isa<fir::ConvertOp, fir::ArrayCoorOp, fir::CoordinateOp,
+ hlfir::DesignateOp, fir::BoxAddrOp, fir::LoadOp>(def))
+ return value;
+
+ // Mirror getMappedVar: only a descriptor load is part of the addressing
+ // chain and must be rebuilt on the device descriptor; any other load is a
+ // live-in and is reused as-is.
+ if (auto load = dyn_cast<fir::LoadOp>(def))
+ if (!mlir::isa<fir::BaseBoxType>(load.getType()))
+ return value;
+
+ IRMapping map;
+ for (Value operand : def->getOperands())
+ map.map(operand, rebuildOnDevice(builder, operand, mappedVar, deviceVar));
+ return builder.clone(*def, map)->getResult(0);
+}
+
+class ACCDevicePtrToCUFKernel
+ : public fir::acc::impl::ACCDevicePtrToCUFKernelBase<
+ ACCDevicePtrToCUFKernel> {
+public:
+ using fir::acc::impl::ACCDevicePtrToCUFKernelBase<
+ ACCDevicePtrToCUFKernel>::ACCDevicePtrToCUFKernelBase;
+
+ void runOnOperation() override {
+ llvm::SmallVector<cuf::KernelLaunchOp> launches;
+ getOperation().walk(
+ [&](cuf::KernelLaunchOp launch) { launches.push_back(launch); });
+
+ for (cuf::KernelLaunchOp launch : launches)
+ rewriteLaunch(launch);
+ }
+
+private:
+ void rewriteLaunch(cuf::KernelLaunchOp launch) {
+ // Collect kernel arguments that are references to a host variable made
+ // present by an enclosing acc.data region.
+ struct MappedArg {
+ OpOperand *operand;
+ Value mappedVar;
+ };
+ llvm::SmallVector<MappedArg> mappedArgs;
+ llvm::SetVector<Value> mappedVars;
+
+ for (OpOperand &operand : launch.getArgsMutable()) {
+ Value arg = operand.get();
+ if (!fir::isa_ref_type(arg.getType()))
+ continue;
+ Value mappedVar = getMappedVar(arg);
+ if (!isMappedInEnclosingAccData(mappedVar, launch))
+ continue;
+ mappedArgs.push_back({&operand, mappedVar});
+ mappedVars.insert(mappedVar);
+ }
+
+ if (mappedArgs.empty())
+ return;
+
+ OpBuilder builder(launch);
+ Location loc = launch.getLoc();
+
+ // One acc.use_device per distinct mapped variable, emitted before the
+ // launch so it dominates the host_data region.
+ llvm::DenseMap<Value, Value> deviceVars;
+ llvm::SmallVector<Value> dataOperands;
+ for (Value mappedVar : mappedVars) {
+ Value deviceVar = acc::UseDeviceOp::create(builder, loc, mappedVar,
+ /*structured=*/true,
+ /*implicit=*/false)
+ .getAccVar();
+ deviceVars[mappedVar] = deviceVar;
+ dataOperands.push_back(deviceVar);
+ }
+
+ // Wrap the launch in an acc.host_data region.
+ auto hostData =
+ acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{}, dataOperands);
+ Block *body = builder.createBlock(&hostData.getRegion());
+ builder.setInsertionPointToStart(body);
+ Operation *terminator = acc::TerminatorOp::create(builder, loc);
+ launch->moveBefore(terminator);
+
+ // Recompute each mapped argument's address on the device pointer.
+ builder.setInsertionPoint(launch);
+ for (MappedArg &mappedArg : mappedArgs) {
+ Value arg = mappedArg.operand->get();
+ Value deviceVar = deviceVars[mappedArg.mappedVar];
+ mappedArg.operand->assign(
+ rebuildOnDevice(builder, arg, mappedArg.mappedVar, deviceVar));
+ }
+ }
+};
+
+} // namespace
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/CMakeLists.txt b/flang/lib/Optimizer/OpenACC/Transforms/CMakeLists.txt
index 5bf4e629861cf..7dd2468dd27c6 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/CMakeLists.txt
+++ b/flang/lib/Optimizer/OpenACC/Transforms/CMakeLists.txt
@@ -1,5 +1,6 @@
add_flang_library(FIROpenACCTransforms
ACCDeclareActionConversion.cpp
+ ACCDevicePtrToCUFKernel.cpp
ACCInitializeFIRAnalyses.cpp
ACCOptimizeFirstprivateMap.cpp
ACCRecipeBufferization.cpp
@@ -9,6 +10,7 @@ add_flang_library(FIROpenACCTransforms
FIROpenACCPassesIncGen
LINK_LIBS
+ CUFDialect
FIRAnalysis
FIRBuilder
FIRDialect
>From f6997f1e9ef803d3edeb98bf1ec1a86bbc3ee339 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 11:04:56 -0700
Subject: [PATCH 02/10] add test and format
---
.../Transforms/ACCDevicePtrToCUFKernel.cpp | 2 +-
.../Fir/OpenACC/device-ptr-to-cuf-kernel.mlir | 92 +++++++++++++++++++
2 files changed, 93 insertions(+), 1 deletion(-)
create mode 100644 flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 863a2490412a9..d7f832514c75e 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -87,7 +87,7 @@ static Value getMappedVar(Value value) {
/// Checks if mappedVar is present due to an enclosing acc.data region.
static bool isMappedInEnclosingAccData(Value mappedVar,
cuf::KernelLaunchOp launch) {
- if(!mappedVar)
+ if (!mappedVar)
return false;
for (auto dataOp = launch->getParentOfType<acc::DataOp>(); dataOp;
dataOp = dataOp->getParentOfType<acc::DataOp>()) {
diff --git a/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
new file mode 100644
index 0000000000000..712d3862f168c
--- /dev/null
+++ b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
@@ -0,0 +1,92 @@
+// RUN: fir-opt %s --acc-device-ptr-to-cuf-kernel -split-input-file | FileCheck %s
+
+// A CUF kernel launched inside an acc.data region that maps a directly-addressed
+// (static) array. The kernel argument is the mapped variable itself, so it is
+// replaced by the acc.use_device result.
+func.func @static_array() {
+ %c1 = arith.constant 1 : i32
+ %0 = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QFEa"}
+ %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %2 = acc.copyin varPtr(%1 : !fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>> {name = "a"}
+ acc.data dataOperands(%2 : !fir.ref<!fir.array<100xi32>>) {
+ cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
+ acc.terminator
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @static_array
+// CHECK: %[[DECL:.*]] = fir.declare
+// CHECK: acc.data
+// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[DECL]] : !fir.ref<!fir.array<100xi32>>)
+// CHECK: acc.host_data dataOperands(%[[DEV]]
+// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[DEV]]) : (!fir.ref<!fir.array<100xi32>>)
+// CHECK: acc.terminator
+
+// -----
+
+// A CUF kernel launched inside an acc.data region that maps a descriptor-based
+// (allocatable) variable. OpenACC maps the descriptor, so the data address is
+// recomputed as box_addr(load(<device descriptor>)) on the acc.use_device
+// result.
+func.func @descriptor_array() {
+ %c1 = arith.constant 1 : i32
+ %0 = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "h", uniq_name = "_QFEh"}
+ %1 = fir.declare %0 {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFEh"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+ %2 = acc.create varPtr(%1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "h"}
+ acc.data dataOperands(%2 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) {
+ %3 = fir.load %1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+ %4 = fir.box_addr %3 : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+ %5 = fir.convert %4 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%5) : (!fir.ref<!fir.array<?xi32>>)
+ acc.terminator
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @descriptor_array
+// CHECK: %[[DECL:.*]] = fir.declare
+// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[DECL]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
+// CHECK: acc.host_data dataOperands(%[[DEV]]
+// CHECK: %[[LOAD:.*]] = fir.load %[[DEV]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+// CHECK: %[[ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+// CHECK: %[[CONV:.*]] = fir.convert %[[ADDR]] : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[CONV]]) : (!fir.ref<!fir.array<?xi32>>)
+
+// -----
+
+// No enclosing acc.data region: the launch must be left untouched.
+func.func @no_enclosing_data() {
+ %c1 = arith.constant 1 : i32
+ %0 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEa"}
+ %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
+ return
+}
+
+// CHECK-LABEL: func.func @no_enclosing_data
+// CHECK-NOT: acc.use_device
+// CHECK-NOT: acc.host_data
+// CHECK: cuf.kernel_launch @kernel
+
+// -----
+
+// The enclosing acc.data maps a different variable than the one launched: no
+// rewrite.
+func.func @unmapped_arg() {
+ %c1 = arith.constant 1 : i32
+ %0 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEa"}
+ %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %2 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEb"}
+ %3 = fir.declare %2 {uniq_name = "_QFEb"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %4 = acc.copyin varPtr(%3 : !fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>> {name = "b"}
+ acc.data dataOperands(%4 : !fir.ref<!fir.array<100xi32>>) {
+ cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
+ acc.terminator
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @unmapped_arg
+// CHECK-NOT: acc.use_device
+// CHECK-NOT: acc.host_data
>From 33db9af34453d7da620f7a66ca5e89c5a98f4075 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 11:06:12 -0700
Subject: [PATCH 03/10] fix test
---
.../test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir | 14 ++++++++++----
1 file changed, 10 insertions(+), 4 deletions(-)
diff --git a/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
index 712d3862f168c..b5fa8124f1a40 100644
--- a/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
+++ b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
@@ -5,8 +5,10 @@
// replaced by the acc.use_device result.
func.func @static_array() {
%c1 = arith.constant 1 : i32
+ %c100 = arith.constant 100 : index
%0 = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QFEa"}
- %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %sh = fir.shape %c100 : (index) -> !fir.shape<1>
+ %1 = fir.declare %0(%sh) {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xi32>>
%2 = acc.copyin varPtr(%1 : !fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>> {name = "a"}
acc.data dataOperands(%2 : !fir.ref<!fir.array<100xi32>>) {
cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
@@ -58,8 +60,10 @@ func.func @descriptor_array() {
// No enclosing acc.data region: the launch must be left untouched.
func.func @no_enclosing_data() {
%c1 = arith.constant 1 : i32
+ %c100 = arith.constant 100 : index
%0 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEa"}
- %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %sh = fir.shape %c100 : (index) -> !fir.shape<1>
+ %1 = fir.declare %0(%sh) {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xi32>>
cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
return
}
@@ -75,10 +79,12 @@ func.func @no_enclosing_data() {
// rewrite.
func.func @unmapped_arg() {
%c1 = arith.constant 1 : i32
+ %c100 = arith.constant 100 : index
+ %sh = fir.shape %c100 : (index) -> !fir.shape<1>
%0 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEa"}
- %1 = fir.declare %0 {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %1 = fir.declare %0(%sh) {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xi32>>
%2 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEb"}
- %3 = fir.declare %2 {uniq_name = "_QFEb"} : (!fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>>
+ %3 = fir.declare %2(%sh) {uniq_name = "_QFEb"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xi32>>
%4 = acc.copyin varPtr(%3 : !fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>> {name = "b"}
acc.data dataOperands(%4 : !fir.ref<!fir.array<100xi32>>) {
cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%1) : (!fir.ref<!fir.array<100xi32>>)
>From 83c16b46de9c6ed4245313ea57fd4288d45e8e80 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 12:16:03 -0700
Subject: [PATCH 04/10] remove unneeded dependent dialects
---
flang/include/flang/Optimizer/OpenACC/Passes.td | 3 +--
.../Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp | 2 +-
2 files changed, 2 insertions(+), 3 deletions(-)
diff --git a/flang/include/flang/Optimizer/OpenACC/Passes.td b/flang/include/flang/Optimizer/OpenACC/Passes.td
index cafe269dcae0c..a2c5e4677fcd2 100644
--- a/flang/include/flang/Optimizer/OpenACC/Passes.td
+++ b/flang/include/flang/Optimizer/OpenACC/Passes.td
@@ -133,8 +133,7 @@ def ACCDevicePtrToCUFKernel
This must run before the pass that lowers cuf.kernel_launch to
gpu.launch_func.
}];
- let dependentDialects = ["mlir::acc::OpenACCDialect", "fir::FIROpsDialect",
- "cuf::CUFDialect"];
+ let dependentDialects = ["mlir::acc::OpenACCDialect"];
}
#endif // FORTRAN_OPTIMIZER_OPENACC_PASSES
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index d7f832514c75e..67440fbe71f0d 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -1,4 +1,4 @@
-//===- ACCDevicePtrToCUFKernel.cpp --------------------------------------===//
+//===- ACCDevicePtrToCUFKernel.cpp ---------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
>From e9df71e543c157f9e3f14c27c15138f8b2cdad7f Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 15:21:29 -0700
Subject: [PATCH 05/10] remove hardcoded ops for peeling; use
getDominatingDataClauses instead of parent searching
---
.../Transforms/ACCDevicePtrToCUFKernel.cpp | 78 +++++++++----------
1 file changed, 35 insertions(+), 43 deletions(-)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 67440fbe71f0d..019921065ad6f 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -24,9 +24,12 @@
#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "flang/Optimizer/OpenACC/Passes.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/OpenACC/OpenACCUtils.h"
+#include "mlir/IR/Dominance.h"
#include "mlir/IR/IRMapping.h"
#include "mlir/Pass/Pass.h"
#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/DenseSet.h"
#include "llvm/ADT/SetVector.h"
#include "llvm/ADT/SmallVector.h"
@@ -45,30 +48,18 @@ namespace {
/// as a data-clause pointer
static Value getMappedVar(Value value) {
while (true) {
- if (auto convert = value.getDefiningOp<fir::ConvertOp>()) {
- value = convert.getValue();
- continue;
- }
- if (auto coor = value.getDefiningOp<fir::ArrayCoorOp>()) {
- value = coor.getMemref();
- continue;
- }
- if (auto coor = value.getDefiningOp<fir::CoordinateOp>()) {
- value = coor.getRef();
- continue;
- }
- if (auto designate = value.getDefiningOp<hlfir::DesignateOp>()) {
- value = designate.getMemref();
+ Operation *def = value.getDefiningOp();
+ // Stop at the variable OpenACC maps as the data-clause pointer.
+ if (isa_and_nonnull<fir::DeclareOp, hlfir::DeclareOp>(def))
+ return value;
+ if (auto view = dyn_cast_or_null<fir::FortranObjectViewOpInterface>(def)) {
+ value = view.getViewSource(cast<OpResult>(value));
continue;
}
// Descriptor-based (allocatable/pointer) variables: the data address is
// extracted from the descriptor via box_addr(load(<descriptor ref>)). Peel
// both so the walk reaches the descriptor variable, which is what OpenACC
// maps as the data-clause varPtr for such variables.
- if (auto boxAddr = value.getDefiningOp<fir::BoxAddrOp>()) {
- value = boxAddr.getVal();
- continue;
- }
if (auto load = value.getDefiningOp<fir::LoadOp>()) {
// Only a load that produces a descriptor is part of the addressing
// chain; scalar loads are ordinary values, not addressing steps.
@@ -77,27 +68,23 @@ static Value getMappedVar(Value value) {
continue;
}
}
- if (isa_and_nonnull<fir::DeclareOp, hlfir::DeclareOp>(
- value.getDefiningOp()))
- return value;
return {};
}
}
-/// Checks if mappedVar is present due to an enclosing acc.data region.
-static bool isMappedInEnclosingAccData(Value mappedVar,
- cuf::KernelLaunchOp launch) {
- if (!mappedVar)
- return false;
- for (auto dataOp = launch->getParentOfType<acc::DataOp>(); dataOp;
- dataOp = dataOp->getParentOfType<acc::DataOp>()) {
- for (Value dataOperand : dataOp.getDataClauseOperands()) {
- if (Value hostVar = acc::getVar(dataOperand.getDefiningOp()))
- if (getMappedVar(hostVar) == mappedVar)
- return true;
- }
- }
- return false;
+/// Collects the host variables made present by an
+/// OpenACC data directive that dominates `launch`. The result is
+/// computed once per launch, then queried per argument.
+static llvm::DenseSet<Value>
+collectPresentAccVars(cuf::KernelLaunchOp launch, DominanceInfo &domInfo,
+ PostDominanceInfo &postDomInfo) {
+ llvm::DenseSet<Value> presentVars;
+ for (Value dataClause :
+ acc::getDominatingDataClauses(launch, domInfo, postDomInfo))
+ if (Value hostVar = acc::getVar(dataClause.getDefiningOp()))
+ if (Value mappedVar = getMappedVar(hostVar))
+ presentVars.insert(mappedVar);
+ return presentVars;
}
/// Reconstructs the addressing chain that produced `value` from `mappedVar`,
@@ -110,16 +97,15 @@ static Value rebuildOnDevice(OpBuilder &builder, Value value, Value mappedVar,
return deviceVar;
Operation *def = value.getDefiningOp();
- if (!def || !isa<fir::ConvertOp, fir::ArrayCoorOp, fir::CoordinateOp,
- hlfir::DesignateOp, fir::BoxAddrOp, fir::LoadOp>(def))
- return value;
-
+ bool isViewStep = isa_and_nonnull<fir::FortranObjectViewOpInterface>(def);
+ bool isDescriptorLoad = false;
// Mirror getMappedVar: only a descriptor load is part of the addressing
// chain and must be rebuilt on the device descriptor; any other load is a
// live-in and is reused as-is.
- if (auto load = dyn_cast<fir::LoadOp>(def))
- if (!mlir::isa<fir::BaseBoxType>(load.getType()))
- return value;
+ if (auto load = dyn_cast_or_null<fir::LoadOp>(def))
+ isDescriptorLoad = isa<fir::BaseBoxType>(load.getType());
+ if (!isViewStep && !isDescriptorLoad)
+ return value;
IRMapping map;
for (Value operand : def->getOperands())
@@ -154,12 +140,18 @@ class ACCDevicePtrToCUFKernel
llvm::SmallVector<MappedArg> mappedArgs;
llvm::SetVector<Value> mappedVars;
+ DominanceInfo domInfo;
+ PostDominanceInfo postDomInfo;
+ llvm::DenseSet<Value> presentAccVars =
+ collectPresentAccVars(launch, domInfo, postDomInfo);
+
for (OpOperand &operand : launch.getArgsMutable()) {
Value arg = operand.get();
if (!fir::isa_ref_type(arg.getType()))
continue;
Value mappedVar = getMappedVar(arg);
- if (!isMappedInEnclosingAccData(mappedVar, launch))
+ // Check if mappedVar is present due to an enclosing OpenACC data region.
+ if (!mappedVar || !presentAccVars.contains(mappedVar))
continue;
mappedArgs.push_back({&operand, mappedVar});
mappedVars.insert(mappedVar);
>From cfcb21f6b1f9435a3d141826a6b1e0afcfc10289 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Mon, 27 Jul 2026 16:48:45 -0700
Subject: [PATCH 06/10] drop isa_ref_type check
---
.../Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 019921065ad6f..a596dfb6d1c7d 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -147,8 +147,6 @@ class ACCDevicePtrToCUFKernel
for (OpOperand &operand : launch.getArgsMutable()) {
Value arg = operand.get();
- if (!fir::isa_ref_type(arg.getType()))
- continue;
Value mappedVar = getMappedVar(arg);
// Check if mappedVar is present due to an enclosing OpenACC data region.
if (!mappedVar || !presentAccVars.contains(mappedVar))
>From 91b472ebad7a55f9f4ab52776091c791d2bf228f Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Tue, 28 Jul 2026 10:26:34 -0700
Subject: [PATCH 07/10] add conditional
---
.../lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp | 1 +
1 file changed, 1 insertion(+)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index a596dfb6d1c7d..7acf798e4f99f 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -177,6 +177,7 @@ class ACCDevicePtrToCUFKernel
// Wrap the launch in an acc.host_data region.
auto hostData =
acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{}, dataOperands);
+ hostData.setIfPresent(true);
Block *body = builder.createBlock(&hostData.getRegion());
builder.setInsertionPointToStart(body);
Operation *terminator = acc::TerminatorOp::create(builder, loc);
>From 6d52d495e6b58aa0a5134f46a98d9f5f241a2ea2 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Tue, 28 Jul 2026 15:07:49 -0700
Subject: [PATCH 08/10] pass kernel argument to acc.use_deviceinstead of
reconstructing addressing chain
---
.../Transforms/ACCDevicePtrToCUFKernel.cpp | 65 +++++--------------
.../Fir/OpenACC/device-ptr-to-cuf-kernel.mlir | 55 +++++++++++++---
2 files changed, 61 insertions(+), 59 deletions(-)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 7acf798e4f99f..4a95e3b9fcbf4 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -87,32 +87,6 @@ collectPresentAccVars(cuf::KernelLaunchOp launch, DominanceInfo &domInfo,
return presentVars;
}
-/// Reconstructs the addressing chain that produced `value` from `mappedVar`,
-/// substituting `deviceVar` for `mappedVar`. Only addressing ops are cloned;
-/// everything else (constants, shapes, ...) is reused as a live-in. New ops are
-/// created at `builder`'s current insertion point.
-static Value rebuildOnDevice(OpBuilder &builder, Value value, Value mappedVar,
- Value deviceVar) {
- if (value == mappedVar)
- return deviceVar;
-
- Operation *def = value.getDefiningOp();
- bool isViewStep = isa_and_nonnull<fir::FortranObjectViewOpInterface>(def);
- bool isDescriptorLoad = false;
- // Mirror getMappedVar: only a descriptor load is part of the addressing
- // chain and must be rebuilt on the device descriptor; any other load is a
- // live-in and is reused as-is.
- if (auto load = dyn_cast_or_null<fir::LoadOp>(def))
- isDescriptorLoad = isa<fir::BaseBoxType>(load.getType());
- if (!isViewStep && !isDescriptorLoad)
- return value;
-
- IRMapping map;
- for (Value operand : def->getOperands())
- map.map(operand, rebuildOnDevice(builder, operand, mappedVar, deviceVar));
- return builder.clone(*def, map)->getResult(0);
-}
-
class ACCDevicePtrToCUFKernel
: public fir::acc::impl::ACCDevicePtrToCUFKernelBase<
ACCDevicePtrToCUFKernel> {
@@ -133,12 +107,8 @@ class ACCDevicePtrToCUFKernel
void rewriteLaunch(cuf::KernelLaunchOp launch) {
// Collect kernel arguments that are references to a host variable made
// present by an enclosing acc.data region.
- struct MappedArg {
- OpOperand *operand;
- Value mappedVar;
- };
- llvm::SmallVector<MappedArg> mappedArgs;
- llvm::SetVector<Value> mappedVars;
+ llvm::SmallVector<OpOperand *> deviceOperands;
+ llvm::SetVector<Value> deviceArgs;
DominanceInfo domInfo;
PostDominanceInfo postDomInfo;
@@ -151,45 +121,40 @@ class ACCDevicePtrToCUFKernel
// Check if mappedVar is present due to an enclosing OpenACC data region.
if (!mappedVar || !presentAccVars.contains(mappedVar))
continue;
- mappedArgs.push_back({&operand, mappedVar});
- mappedVars.insert(mappedVar);
+ deviceOperands.push_back(&operand);
+ deviceArgs.insert(arg);
}
- if (mappedArgs.empty())
+ if (deviceOperands.empty())
return;
OpBuilder builder(launch);
Location loc = launch.getLoc();
- // One acc.use_device per distinct mapped variable, emitted before the
+ // One acc.use_device per distinct kernel argument, emitted before the
// launch so it dominates the host_data region.
- llvm::DenseMap<Value, Value> deviceVars;
- llvm::SmallVector<Value> dataOperands;
- for (Value mappedVar : mappedVars) {
- Value deviceVar = acc::UseDeviceOp::create(builder, loc, mappedVar,
+ llvm::DenseMap<Value, Value> deviceArgsMap;
+ llvm::SmallVector<Value> hostDataOperands;
+ for (Value arg : deviceArgs) {
+ Value deviceVar = acc::UseDeviceOp::create(builder, loc, arg,
/*structured=*/true,
/*implicit=*/false)
.getAccVar();
- deviceVars[mappedVar] = deviceVar;
- dataOperands.push_back(deviceVar);
+ deviceArgsMap[arg] = deviceVar;
+ hostDataOperands.push_back(deviceVar);
}
// Wrap the launch in an acc.host_data region.
auto hostData =
- acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{}, dataOperands);
+ acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{}, hostDataOperands);
hostData.setIfPresent(true);
Block *body = builder.createBlock(&hostData.getRegion());
builder.setInsertionPointToStart(body);
Operation *terminator = acc::TerminatorOp::create(builder, loc);
launch->moveBefore(terminator);
- // Recompute each mapped argument's address on the device pointer.
- builder.setInsertionPoint(launch);
- for (MappedArg &mappedArg : mappedArgs) {
- Value arg = mappedArg.operand->get();
- Value deviceVar = deviceVars[mappedArg.mappedVar];
- mappedArg.operand->assign(
- rebuildOnDevice(builder, arg, mappedArg.mappedVar, deviceVar));
+ for (OpOperand *operand : deviceOperands) {
+ operand->assign(deviceArgsMap[operand->get()]);
}
}
};
diff --git a/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
index b5fa8124f1a40..786073a01547d 100644
--- a/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
+++ b/flang/test/Fir/OpenACC/device-ptr-to-cuf-kernel.mlir
@@ -2,7 +2,7 @@
// A CUF kernel launched inside an acc.data region that maps a directly-addressed
// (static) array. The kernel argument is the mapped variable itself, so it is
-// replaced by the acc.use_device result.
+// wrapped in acc.use_device and substituted directly.
func.func @static_array() {
%c1 = arith.constant 1 : i32
%c100 = arith.constant 100 : index
@@ -23,14 +23,49 @@ func.func @static_array() {
// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[DECL]] : !fir.ref<!fir.array<100xi32>>)
// CHECK: acc.host_data dataOperands(%[[DEV]]
// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[DEV]]) : (!fir.ref<!fir.array<100xi32>>)
-// CHECK: acc.terminator
+// CHECK: attributes {ifPresent}
+
+// -----
+
+// A CUF kernel launched with an array section. The launch argument is already an
+// interior pointer (&a(3)) computed on the host base. The pass wraps that
+// interior pointer directly in acc.use_device and substitutes it; the section
+// addressing is left untouched (not rebuilt on a device base).
+func.func @array_section() {
+ %c1 = arith.constant 1 : i32
+ %c2_i32 = arith.constant 2 : i32
+ %c100 = arith.constant 100 : index
+ %c3 = arith.constant 3 : index
+ %0 = fir.alloca !fir.array<100xi32> {uniq_name = "_QFEa"}
+ %sh = fir.shape %c100 : (index) -> !fir.shape<1>
+ %1 = fir.declare %0(%sh) {uniq_name = "_QFEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xi32>>
+ %2 = acc.copyin varPtr(%1 : !fir.ref<!fir.array<100xi32>>) -> !fir.ref<!fir.array<100xi32>> {name = "a"}
+ acc.data dataOperands(%2 : !fir.ref<!fir.array<100xi32>>) {
+ %3 = fir.array_coor %1(%sh) %c3 : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, index) -> !fir.ref<i32>
+ %4 = fir.convert %3 : (!fir.ref<i32>) -> !fir.ref<!fir.array<?xi32>>
+ cuf.kernel_launch @kernel<<<%c1, %c1, %c1, %c1, %c1, %c1>>>(%4, %c2_i32) : (!fir.ref<!fir.array<?xi32>>, i32)
+ acc.terminator
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @array_section
+// CHECK: %[[DECL:.*]] = fir.declare
+// CHECK: acc.data
+// CHECK: %[[COOR:.*]] = fir.array_coor %[[DECL]]
+// CHECK: %[[CONV:.*]] = fir.convert %[[COOR]] : (!fir.ref<i32>) -> !fir.ref<!fir.array<?xi32>>
+// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[CONV]] : !fir.ref<!fir.array<?xi32>>)
+// CHECK: acc.host_data dataOperands(%[[DEV]]
+// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[DEV]], {{.*}}) : (!fir.ref<!fir.array<?xi32>>, i32)
+// CHECK: attributes {ifPresent}
// -----
// A CUF kernel launched inside an acc.data region that maps a descriptor-based
-// (allocatable) variable. OpenACC maps the descriptor, so the data address is
-// recomputed as box_addr(load(<device descriptor>)) on the acc.use_device
-// result.
+// (allocatable) variable. OpenACC maps the descriptor, but the launch already
+// extracts the data address on the host via box_addr(load(<descriptor>)). The
+// pass wraps that data pointer in acc.use_device and substitutes it directly;
+// the descriptor addressing is left on the host descriptor and not rebuilt.
func.func @descriptor_array() {
%c1 = arith.constant 1 : i32
%0 = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "h", uniq_name = "_QFEh"}
@@ -48,12 +83,14 @@ func.func @descriptor_array() {
// CHECK-LABEL: func.func @descriptor_array
// CHECK: %[[DECL:.*]] = fir.declare
-// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[DECL]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
-// CHECK: acc.host_data dataOperands(%[[DEV]]
-// CHECK: %[[LOAD:.*]] = fir.load %[[DEV]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+// CHECK: acc.data
+// CHECK: %[[LOAD:.*]] = fir.load %[[DECL]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
// CHECK: %[[ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
// CHECK: %[[CONV:.*]] = fir.convert %[[ADDR]] : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
-// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[CONV]]) : (!fir.ref<!fir.array<?xi32>>)
+// CHECK: %[[DEV:.*]] = acc.use_device varPtr(%[[CONV]] : !fir.ref<!fir.array<?xi32>>)
+// CHECK: acc.host_data dataOperands(%[[DEV]]
+// CHECK: cuf.kernel_launch @kernel<<<{{.*}}>>>(%[[DEV]]) : (!fir.ref<!fir.array<?xi32>>)
+// CHECK: attributes {ifPresent}
// -----
>From b185118f033c26fe018610c4ab5d406bb35507c1 Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Tue, 28 Jul 2026 18:26:35 -0700
Subject: [PATCH 09/10] format
---
.../Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 4a95e3b9fcbf4..83bdd4da2e090 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -145,8 +145,8 @@ class ACCDevicePtrToCUFKernel
}
// Wrap the launch in an acc.host_data region.
- auto hostData =
- acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{}, hostDataOperands);
+ auto hostData = acc::HostDataOp::create(builder, loc, /*ifCond=*/Value{},
+ hostDataOperands);
hostData.setIfPresent(true);
Block *body = builder.createBlock(&hostData.getRegion());
builder.setInsertionPointToStart(body);
>From be64fb6c3a889974b1445d783257fa1ed37bc2be Mon Sep 17 00:00:00 2001
From: Yebin Chon <ychon at nvidia.com>
Date: Wed, 29 Jul 2026 12:54:27 -0700
Subject: [PATCH 10/10] add support for char args
---
.../Transforms/ACCDevicePtrToCUFKernel.cpp | 32 +++++++++++++++++--
1 file changed, 30 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
index 83bdd4da2e090..fc95b6730c5b7 100644
--- a/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
+++ b/flang/lib/Optimizer/OpenACC/Transforms/ACCDevicePtrToCUFKernel.cpp
@@ -56,6 +56,12 @@ static Value getMappedVar(Value value) {
value = view.getViewSource(cast<OpResult>(value));
continue;
}
+ // TODO: drop this case once (if) fir.emboxchar implements the
+ // FortranObjectViewOpInterface.
+ if (auto embox = dyn_cast_or_null<fir::EmboxCharOp>(def)) {
+ value = embox.getMemref();
+ continue;
+ }
// Descriptor-based (allocatable/pointer) variables: the data address is
// extracted from the descriptor via box_addr(load(<descriptor ref>)). Peel
// both so the walk reaches the descriptor variable, which is what OpenACC
@@ -110,6 +116,13 @@ class ACCDevicePtrToCUFKernel
llvm::SmallVector<OpOperand *> deviceOperands;
llvm::SetVector<Value> deviceArgs;
+ // Special case: a boxchar argument is not pointer-like, so acc.use_device
+ // cannot take it. Translate its base address and rebuild the boxchar on the
+ // device address. Keyed by the launch operand so the general path below is
+ // untouched. TODO: remove once fir.emboxchar is a
+ // FortranObjectViewOpInterface.
+ llvm::DenseMap<OpOperand *, fir::EmboxCharOp> boxCharArgs;
+
DominanceInfo domInfo;
PostDominanceInfo postDomInfo;
llvm::DenseSet<Value> presentAccVars =
@@ -122,7 +135,12 @@ class ACCDevicePtrToCUFKernel
if (!mappedVar || !presentAccVars.contains(mappedVar))
continue;
deviceOperands.push_back(&operand);
- deviceArgs.insert(arg);
+ if (auto embox = arg.getDefiningOp<fir::EmboxCharOp>()) {
+ boxCharArgs[&operand] = embox;
+ deviceArgs.insert(embox.getMemref()); // translate the base pointer
+ } else {
+ deviceArgs.insert(arg); // general case
+ }
}
if (deviceOperands.empty())
@@ -154,7 +172,17 @@ class ACCDevicePtrToCUFKernel
launch->moveBefore(terminator);
for (OpOperand *operand : deviceOperands) {
- operand->assign(deviceArgsMap[operand->get()]);
+ if (fir::EmboxCharOp embox = boxCharArgs.lookup(operand)) {
+ // Char case: rebuild the boxchar on the device base address. Must be
+ // emitted before the launch so the rebuilt value dominates its use.
+ builder.setInsertionPoint(launch);
+ operand->assign(fir::EmboxCharOp::create(
+ builder, embox.getLoc(), embox.getType(),
+ deviceArgsMap[embox.getMemref()], embox.getLen()));
+ } else {
+ // General case
+ operand->assign(deviceArgsMap[operand->get()]);
+ }
}
}
};
More information about the flang-commits
mailing list