[flang-commits] [flang] a3ac598 - [flang][cuda] Defer data-transfer conversion in OpenACC routines (#225906)
via flang-commits
flang-commits at lists.llvm.org
Wed Sep 23 19:02:17 PDT 2026
Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-09-24T02:02:07Z
New Revision: a3ac598a8782df337f477672979d4dd53dc601b0
URL: https://github.com/llvm/llvm-project/commit/a3ac598a8782df337f477672979d4dd53dc601b0
DIFF: https://github.com/llvm/llvm-project/commit/a3ac598a8782df337f477672979d4dd53dc601b0.diff
LOG: [flang][cuda] Defer data-transfer conversion in OpenACC routines (#225906)
Add an option so cuf-convert can leave cuf.data_transfer in the
host copy of an OpenACC routine until host/device specialization.
After that, transfers in a specialized device body become
assignments. Place allocas inside IsolatedFromAbove offload
regions so later FIR lowering does not store into a temporary
defined outside acc.compute_region.
Added:
flang/test/Fir/CUDA/cuda-data-transfer-defer-acc-routine.mlir
Modified:
flang/include/flang/Optimizer/Transforms/Passes.td
flang/lib/Optimizer/Builder/FIRBuilder.cpp
flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
Removed:
################################################################################
diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index 6dbaa49e0ad7e..3034fbf9f386b 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -588,8 +588,15 @@ def CUFAllocationConversion : Pass<"cuf-allocation-convert", "mlir::ModuleOp"> {
def CUFOpConversion : Pass<"cuf-convert", "mlir::ModuleOp"> {
let summary = "Convert some CUF operations to runtime calls";
- let dependentDialects = ["fir::FIROpsDialect", "mlir::gpu::GPUDialect",
- "mlir::DLTIDialect"];
+ let dependentDialects = ["fir::FIROpsDialect", "hlfir::hlfirDialect",
+ "mlir::gpu::GPUDialect", "mlir::DLTIDialect"];
+ let options = [Option<
+ "deferAccRoutineDataTransfers", "defer-acc-routine-data-transfers",
+ "bool", /*default=*/"false",
+ "Do not convert cuf.data_transfer operations in the host copy of an "
+ "OpenACC routine. The routine is later specialized for host and device, "
+ "after which each copy can be handled according to its execution "
+ "context.">];
}
def CUFOpConversionLate : Pass<"cuf-convert-late", "mlir::ModuleOp"> {
diff --git a/flang/lib/Optimizer/Builder/FIRBuilder.cpp b/flang/lib/Optimizer/Builder/FIRBuilder.cpp
index 5793926504112..b808dbfa267cf 100644
--- a/flang/lib/Optimizer/Builder/FIRBuilder.cpp
+++ b/flang/lib/Optimizer/Builder/FIRBuilder.cpp
@@ -290,6 +290,12 @@ mlir::Block *fir::getAllocaBlock(mlir::Region ®ion) {
mlir::dyn_cast<mlir::acc::ComputeRegionOpInterface>(parent))
return accComputeRegionIface.getAllocaBlock();
+ // Offload regions are isolated from above, so allocas cannot be hoisted
+ // past them.
+ if (auto accOffloadRegionIface =
+ mlir::dyn_cast<mlir::acc::OffloadRegionOpInterface>(parent))
+ return &accOffloadRegionIface.getOffloadRegion().front();
+
if (auto ompOutlineableIface =
mlir::dyn_cast<mlir::omp::OutlineableOpenMPOpInterface>(parent))
return ompOutlineableIface.getAllocaBlock();
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
index 7a3576d37a030..d537e9442c8e1 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
@@ -14,6 +14,7 @@
#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
+#include "flang/Optimizer/HLFIR/HLFIRDialect.h"
#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "flang/Optimizer/Support/DataLayout.h"
#include "flang/Optimizer/Transforms/Passes.h"
@@ -130,6 +131,21 @@ static mlir::Value getShapeFromDecl(mlir::Value src) {
return mlir::Value{};
}
+// hlfir.assign rejects a raw !fir.ref<!fir.array<?xT>> because a dynamic-size
+// array is not an HLFIR variable unless it is boxed. Use the transfer shape
+// (or a declare's shape) to build a descriptor.
+static mlir::Value asHLFIREntity(mlir::PatternRewriter &rewriter,
+ mlir::Location loc, mlir::Value val,
+ mlir::Value shape) {
+ if (hlfir::isFortranEntity(val))
+ return val;
+ mlir::Type unwrapped = fir::unwrapRefType(val.getType());
+ if (!shape)
+ shape = getShapeFromDecl(val);
+ auto boxTy = fir::BoxType::get(unwrapped);
+ return fir::EmboxOp::create(rewriter, loc, boxTy, val, shape);
+}
+
static mlir::Value emboxSrc(mlir::PatternRewriter &rewriter,
cuf::DataTransferOp op,
const mlir::SymbolTable &symtab,
@@ -215,6 +231,28 @@ struct CUFDataTransferOpConversion
mlir::Type dstTy = fir::unwrapRefType(op.getDst().getType());
mlir::Location loc = op.getLoc();
+ // A transfer left in the accelerator copy of an OpenACC routine describes
+ // an assignment that the device executes itself. A host runtime copy makes
+ // no sense there, so turn it back into a plain assignment.
+ if (inDeviceContext(op) && mlir::acc::isSpecializedAccRoutine(
+ op->getParentOfType<mlir::func::FuncOp>())) {
+ mlir::Value src = op.getSrc();
+ mlir::Value dst = op.getDst();
+ if (fir::isa_trivial(srcTy) && fir::isa_ref_type(dst.getType()) &&
+ fir::isa_trivial(dstTy)) {
+ if (fir::isa_ref_type(src.getType()))
+ src = fir::LoadOp::create(rewriter, loc, src);
+ src = createConvertOp(rewriter, loc, dstTy, src);
+ fir::StoreOp::create(rewriter, loc, src, dst);
+ } else {
+ mlir::Value shape = op.getShape();
+ src = asHLFIREntity(rewriter, loc, src, shape);
+ dst = asHLFIREntity(rewriter, loc, dst, shape);
+ hlfir::AssignOp::create(rewriter, loc, src, dst);
+ }
+ rewriter.eraseOp(op);
+ return mlir::success();
+ }
unsigned mode = 0;
if (op.getTransferKind() == cuf::DataTransferKind::HostDevice) {
mode = kHostToDevice;
@@ -614,9 +652,16 @@ class CUFOpConversion : public fir::impl::CUFOpConversionBase<CUFOpConversion> {
fir::LLVMTypeConverter typeConverter(module, /*applyTBAA=*/false,
/*forceUnifiedTBAATree=*/false, *dl);
target.addLegalDialect<fir::FIROpsDialect, mlir::arith::ArithDialect,
- mlir::gpu::GPUDialect>();
+ mlir::gpu::GPUDialect, hlfir::hlfirDialect>();
target.addLegalOp<cuf::StreamCastOp>();
target.addLegalOp<cuf::DeviceAddressOp>();
+ target.addDynamicallyLegalOp<cuf::DataTransferOp>(
+ [&](cuf::DataTransferOp op) {
+ if (!deferAccRoutineDataTransfers)
+ return false;
+ auto funcOp = op->getParentOfType<mlir::func::FuncOp>();
+ return funcOp && mlir::acc::isAccRoutine(funcOp);
+ });
target.addLegalOp<cuf::DeviceIsActiveOp>();
cuf::populateCUFToFIRConversionPatterns(typeConverter, *dl, symtab,
patterns);
diff --git a/flang/test/Fir/CUDA/cuda-data-transfer-defer-acc-routine.mlir b/flang/test/Fir/CUDA/cuda-data-transfer-defer-acc-routine.mlir
new file mode 100644
index 0000000000000..84ac0e2f7a3c7
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-data-transfer-defer-acc-routine.mlir
@@ -0,0 +1,101 @@
+// RUN: fir-opt --cuf-convert="defer-acc-routine-data-transfers=true" %s | FileCheck %s --check-prefix=DEFER
+// RUN: fir-opt --cuf-convert %s | FileCheck %s --check-prefix=CONVERT
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
+ func.func @host(%dst: !fir.ref<i32>) {
+ %c1_i32 = arith.constant 1 : i32
+ cuf.data_transfer %c1_i32 to %dst {transfer_kind = #cuf.cuda_transfer<host_device>} : i32, !fir.ref<i32>
+ return
+ }
+
+ func.func @acc_routine(%dst: !fir.ref<i32>) attributes {acc.routine_info = #acc.routine_info<[@routine]>} {
+ %c1_i32 = arith.constant 1 : i32
+ cuf.data_transfer %c1_i32 to %dst {transfer_kind = #cuf.cuda_transfer<host_device>} : i32, !fir.ref<i32>
+ return
+ }
+
+ func.func @device_specialized() attributes {acc.specialized_routine = #acc.specialized_routine<@routine, <seq>, "device_specialized">} {
+ acc.compute_region {
+ %dst = fir.alloca i32
+ %c1_i32 = arith.constant 1 : i32
+ cuf.data_transfer %c1_i32 to %dst {transfer_kind = #cuf.cuda_transfer<host_device>} : i32, !fir.ref<i32>
+ acc.yield
+ } <{origin = "acc.routine"}>
+ return
+ }
+
+ func.func @device_specialized_logical() attributes {acc.specialized_routine = #acc.specialized_routine<@routine, <seq>, "device_specialized_logical">} {
+ acc.compute_region {
+ %dst = fir.alloca !fir.logical<4>
+ %true = arith.constant true
+ cuf.data_transfer %true to %dst {transfer_kind = #cuf.cuda_transfer<host_device>} : i1, !fir.ref<!fir.logical<4>>
+ acc.yield
+ } <{origin = "acc.routine"}>
+ return
+ }
+
+ // dst = src for explicit-shape arrays (host src, device dst) in an OpenACC
+ // routine. The host copy keeps a transfer; the specialized device body
+ // becomes an assignment.
+ func.func @acc_routine_copy_array() attributes {acc.routine_info = #acc.routine_info<[@routine]>} {
+ %n = arith.constant 10 : index
+ %shape = fir.shape %n : (index) -> !fir.shape<1>
+ %src = fir.alloca !fir.array<?xf32>, %n
+ %dst = fir.alloca !fir.array<?xf32>, %n
+ cuf.data_transfer %src to %dst, %shape : !fir.shape<1> {transfer_kind = #cuf.cuda_transfer<host_device>} : !fir.ref<!fir.array<?xf32>>, !fir.ref<!fir.array<?xf32>>
+ return
+ }
+
+ func.func @device_specialized_copy_array() attributes {acc.specialized_routine = #acc.specialized_routine<@routine, <seq>, "device_specialized_copy_array">} {
+ acc.compute_region {
+ %n = arith.constant 10 : index
+ %shape = fir.shape %n : (index) -> !fir.shape<1>
+ %src = fir.alloca !fir.array<?xf32>, %n
+ %dst = fir.alloca !fir.array<?xf32>, %n
+ cuf.data_transfer %src to %dst, %shape : !fir.shape<1> {transfer_kind = #cuf.cuda_transfer<host_device>} : !fir.ref<!fir.array<?xf32>>, !fir.ref<!fir.array<?xf32>>
+ acc.yield
+ } <{origin = "acc.routine"}>
+ return
+ }
+}
+
+// DEFER-LABEL: func.func @host(
+// DEFER-NOT: cuf.data_transfer
+// DEFER: fir.call @_FortranACUFDataTransferPtrPtr
+
+// DEFER-LABEL: func.func @acc_routine(
+// DEFER: cuf.data_transfer
+// DEFER-NOT: fir.call @_FortranACUFDataTransferPtrPtr
+
+// DEFER-LABEL: func.func @acc_routine_copy_array(
+// DEFER: cuf.data_transfer
+// DEFER-NOT: fir.call @_FortranACUFDataTransferPtrPtr
+
+// CONVERT-LABEL: func.func @host(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT: fir.call @_FortranACUFDataTransferPtrPtr
+
+// CONVERT-LABEL: func.func @acc_routine(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT: fir.call @_FortranACUFDataTransferPtrPtr
+
+// CONVERT-LABEL: func.func @device_specialized(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT-NOT: fir.call @_FortranACUFDataTransferPtrPtr
+// CONVERT: fir.store %{{.*}} to %{{.*}} : !fir.ref<i32>
+
+// CONVERT-LABEL: func.func @device_specialized_logical(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT: %[[CVT:.*]] = fir.convert %{{.*}} : (i1) -> !fir.logical<4>
+// CONVERT: fir.store %[[CVT]] to %{{.*}} : !fir.ref<!fir.logical<4>>
+
+// CONVERT-LABEL: func.func @acc_routine_copy_array(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT: fir.call @_FortranACUFDataTransferPtrPtr
+
+// CONVERT-LABEL: func.func @device_specialized_copy_array(
+// CONVERT-NOT: cuf.data_transfer
+// CONVERT-NOT: fir.call @_FortranACUFDataTransferPtrPtr
+// CONVERT: %[[SRC:.*]] = fir.embox %{{.*}}(%{{.*}}) : (!fir.ref<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xf32>>
+// CONVERT: %[[DST:.*]] = fir.embox %{{.*}}(%{{.*}}) : (!fir.ref<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xf32>>
+// CONVERT: hlfir.assign %[[SRC]] to %[[DST]] : !fir.box<!fir.array<?xf32>>, !fir.box<!fir.array<?xf32>>
More information about the flang-commits
mailing list