[flang-commits] [flang] [flang][cuda] Allocate adjustable automatic arrays in unified/managed memory (PR #212965)
via flang-commits
flang-commits at lists.llvm.org
Thu Jul 30 13:49:19 PDT 2026
https://github.com/khaki3 updated https://github.com/llvm/llvm-project/pull/212965
>From 4d3e177c7c1c1149d843283c9f92305e3a921004 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Thu, 30 Jul 2026 01:29:13 -0700
Subject: [PATCH 1/3] [flang][cuda] Allocate adjustable autos in
unified/managed memory
Under -gpu=mem:unified|managed, tag dynamic-extent automatic arrays so
lowering uses cuf.alloc/cuf.free; fixed-size automatic arrays stay on
the stack.
---
flang/lib/Semantics/resolve-names.cpp | 31 ++++++++++
.../CUDA/cuda-gpu-unified-automatic-array.f90 | 58 +++++++++++++++++++
2 files changed, 89 insertions(+)
create mode 100644 flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index 83d92f253e626..8e510250f5ab8 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -10648,6 +10648,37 @@ void ResolveNamesVisitor::FinishSpecificationPart(
context().languageFeatures().IsEnabled(
common::LanguageFeature::CudaPinned))
object->set_cudaDataAttr(common::CUDADataAttr::Pinned);
+ } else if (!object->cudaDataAttr() && !IsDummy(symbol) &&
+ !IsAllocatable(symbol) && !IsPointer(symbol) && !IsSaved(symbol) &&
+ !IsCUDADeviceContext(&symbol.owner()) &&
+ object->shape().IsExplicitShape()) {
+ // Under -gpu=mem:unified|managed, allocate adjustable / VLA automatic
+ // arrays in CUDA unified/managed memory (fixed-size automatic arrays
+ // stay on the stack). Tag those locals so lowering uses
+ // cuf.alloc/cuf.free. Unlike the allocatable managed tagging above,
+ // this does not require -fcuda: OpenACC + -gpu=mem:unified relies on
+ // it, and cuf.alloc does not go through the CUDA Fortran
+ // managed-descriptor pipeline that motivated the -fcuda gate.
+ auto boundIsNonConstant{[](const Bound &b) {
+ return !b.isExplicit() || !b.GetExplicit() ||
+ !evaluate::IsConstantExpr(*b.GetExplicit());
+ }};
+ bool hasDynamicExtent{false};
+ for (const ShapeSpec &ss : object->shape()) {
+ if (boundIsNonConstant(ss.lbound()) ||
+ boundIsNonConstant(ss.ubound())) {
+ hasDynamicExtent = true;
+ break;
+ }
+ }
+ if (hasDynamicExtent) {
+ if (context().languageFeatures().IsEnabled(
+ common::LanguageFeature::CudaUnified))
+ object->set_cudaDataAttr(common::CUDADataAttr::Unified);
+ else if (context().languageFeatures().IsEnabled(
+ common::LanguageFeature::CudaManaged))
+ object->set_cudaDataAttr(common::CUDADataAttr::Managed);
+ }
}
}
}
diff --git a/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90 b/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
new file mode 100644
index 0000000000000..07638b0ca8dfa
--- /dev/null
+++ b/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
@@ -0,0 +1,58 @@
+! RUN: bbc -emit-hlfir -gpu=unified %s -o - | FileCheck %s
+! RUN: bbc -emit-hlfir -gpu=managed %s -o - | FileCheck %s --check-prefix=MANAGED
+
+! Under -gpu=mem:unified|managed, allocate adjustable / VLA automatic arrays
+! in CUDA unified/managed memory. Fixed-size automatic arrays stay on the
+! stack so host pointers remain shared under unified memory (e.g. OpenACC).
+
+module m_adj
+ integer :: nx = 32
+end module
+
+! CHECK-LABEL: func.func @_QPvla(
+! CHECK: %[[ALLOC:.*]] = cuf.alloc !fir.array<?xf32>, %{{.*}} : index {bindc_name = "a", data_attr = #cuf.cuda<unified>, uniq_name = "_QFvlaEa"}
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOC]](%{{.*}}) {data_attr = #cuf.cuda<unified>, uniq_name = "_QFvlaEa"}
+! CHECK: cuf.free %[[DECL]]#1 : !fir.ref<!fir.array<?xf32>> {data_attr = #cuf.cuda<unified>}
+! MANAGED-LABEL: func.func @_QPvla(
+! MANAGED: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<managed>
+! MANAGED: cuf.free %{{.*}} : !fir.ref<!fir.array<?xf32>> {data_attr = #cuf.cuda<managed>}
+subroutine vla(n)
+ integer :: n
+ real :: a(n)
+ a(1) = 1.0
+end subroutine
+
+! CHECK-LABEL: func.func @_QPadjustable(
+! CHECK: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<unified>
+! CHECK: cuf.free %{{.*}} {data_attr = #cuf.cuda<unified>}
+! MANAGED-LABEL: func.func @_QPadjustable(
+! MANAGED: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<managed>
+subroutine adjustable
+ use m_adj
+ real :: a(0:(nx+1)/2)
+ a(0) = 0.0
+end subroutine
+
+! Fixed-size automatic arrays must remain ordinary stack allocations.
+! CHECK-LABEL: func.func @_QPfixed(
+! CHECK-NOT: cuf.alloc
+! CHECK: fir.alloca !fir.array<128xf32>
+! CHECK-NOT: cuf.free
+! MANAGED-LABEL: func.func @_QPfixed(
+! MANAGED-NOT: cuf.alloc
+! MANAGED: fir.alloca !fir.array<128xf32>
+subroutine fixed
+ real :: a(128)
+ a(1) = 1.0
+end subroutine
+
+! Dummy adjustable arrays are caller-allocated; do not retag them.
+! CHECK-LABEL: func.func @_QPdummy_adj(
+! CHECK-NOT: cuf.alloc
+! MANAGED-LABEL: func.func @_QPdummy_adj(
+! MANAGED-NOT: cuf.alloc
+subroutine dummy_adj(a, n)
+ integer :: n
+ real :: a(n)
+ a(1) = 1.0
+end subroutine
>From 136e4b632561e9c8401b1afd99f58d58e519e2fa Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Thu, 30 Jul 2026 01:38:51 -0700
Subject: [PATCH 2/3] [flang][cuda] Check unified/managed features before shape
walk
---
flang/lib/Semantics/resolve-names.cpp | 36 +++++++++++++--------------
1 file changed, 18 insertions(+), 18 deletions(-)
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index 8e510250f5ab8..a3514ccf72e5c 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -10659,26 +10659,26 @@ void ResolveNamesVisitor::FinishSpecificationPart(
// this does not require -fcuda: OpenACC + -gpu=mem:unified relies on
// it, and cuf.alloc does not go through the CUDA Fortran
// managed-descriptor pipeline that motivated the -fcuda gate.
- auto boundIsNonConstant{[](const Bound &b) {
- return !b.isExplicit() || !b.GetExplicit() ||
- !evaluate::IsConstantExpr(*b.GetExplicit());
- }};
- bool hasDynamicExtent{false};
- for (const ShapeSpec &ss : object->shape()) {
- if (boundIsNonConstant(ss.lbound()) ||
- boundIsNonConstant(ss.ubound())) {
- hasDynamicExtent = true;
- break;
+ std::optional<common::CUDADataAttr> attr;
+ if (context().languageFeatures().IsEnabled(
+ common::LanguageFeature::CudaUnified))
+ attr = common::CUDADataAttr::Unified;
+ else if (context().languageFeatures().IsEnabled(
+ common::LanguageFeature::CudaManaged))
+ attr = common::CUDADataAttr::Managed;
+ if (attr) {
+ auto boundIsNonConstant{[](const Bound &b) {
+ return !b.isExplicit() || !b.GetExplicit() ||
+ !evaluate::IsConstantExpr(*b.GetExplicit());
+ }};
+ for (const ShapeSpec &ss : object->shape()) {
+ if (boundIsNonConstant(ss.lbound()) ||
+ boundIsNonConstant(ss.ubound())) {
+ object->set_cudaDataAttr(*attr);
+ break;
+ }
}
}
- if (hasDynamicExtent) {
- if (context().languageFeatures().IsEnabled(
- common::LanguageFeature::CudaUnified))
- object->set_cudaDataAttr(common::CUDADataAttr::Unified);
- else if (context().languageFeatures().IsEnabled(
- common::LanguageFeature::CudaManaged))
- object->set_cudaDataAttr(common::CUDADataAttr::Managed);
- }
}
}
}
>From db6bcd0eedfe4aa45b3809bdf8d19117d442f472 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Thu, 30 Jul 2026 13:48:45 -0700
Subject: [PATCH 3/3] [flang][cuda] Route dynamic autos through
malloc_unified/free_unified
Drop implicit cudaDataAttr tagging for adjustable/VLA automatic arrays.
Under -gpu=mem:unified|managed, promote those locals to fir.allocmem and
lower host heap allocations via the indirect unified/managed allocators.
---
.../Optimizer/Dialect/Support/FIRContext.h | 14 +++
.../flang/Optimizer/Transforms/MemoryUtils.h | 7 ++
flang/lib/Lower/Bridge.cpp | 7 ++
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 84 +++++++++++++-----
.../Optimizer/Dialect/Support/FIRContext.cpp | 41 +++++++++
.../Transforms/AllocationPlacement.cpp | 5 ++
.../Optimizer/Transforms/MemoryAllocation.cpp | 3 +
.../lib/Optimizer/Transforms/MemoryUtils.cpp | 65 ++++++++++++++
.../lib/Optimizer/Transforms/StackArrays.cpp | 7 ++
flang/lib/Semantics/resolve-names.cpp | 31 -------
.../test/Fir/CUDA/cuda-heap-alloc-managed.fir | 44 ++++++++++
.../test/Fir/CUDA/cuda-heap-alloc-unified.fir | 86 +++++++++++++++++++
.../CUDA/cuda-gpu-unified-automatic-array.f90 | 42 ++++-----
13 files changed, 360 insertions(+), 76 deletions(-)
create mode 100644 flang/test/Fir/CUDA/cuda-heap-alloc-managed.fir
create mode 100644 flang/test/Fir/CUDA/cuda-heap-alloc-unified.fir
diff --git a/flang/include/flang/Optimizer/Dialect/Support/FIRContext.h b/flang/include/flang/Optimizer/Dialect/Support/FIRContext.h
index 79337584d6d67..27eee3a101e29 100644
--- a/flang/include/flang/Optimizer/Dialect/Support/FIRContext.h
+++ b/flang/include/flang/Optimizer/Dialect/Support/FIRContext.h
@@ -121,6 +121,20 @@ void setIsPIE(mlir::ModuleOp mod, bool value);
/// Get whether the module is compiled as a position-independent executable.
bool getIsPIE(mlir::ModuleOp mod);
+/// Host heap allocator selected under -gpu=mem:unified|managed, recorded on the
+/// module by lowering and consumed by the allocation placement passes.
+enum class CudaHeapAllocMode { None, Unified, Managed };
+
+void setCudaHeapAllocMode(mlir::ModuleOp mod, CudaHeapAllocMode mode);
+CudaHeapAllocMode getCudaHeapAllocMode(mlir::ModuleOp mod);
+
+/// Same attribute on a fir.allocmem/fir.freemem pair: this allocation uses the
+/// indirect runtime entry points (`malloc_unified`/`free_unified`, ...) instead
+/// of libc. Only pairs created together may be marked, since the allocator and
+/// the deallocator must match.
+void setCudaHeapAllocMode(mlir::Operation *op, CudaHeapAllocMode mode);
+CudaHeapAllocMode getCudaHeapAllocMode(mlir::Operation *op);
+
/// Helper for determining the target from the host, etc. Tools may use this
/// function to provide a consistent interpretation of the `--target=<string>`
/// command-line option.
diff --git a/flang/include/flang/Optimizer/Transforms/MemoryUtils.h b/flang/include/flang/Optimizer/Transforms/MemoryUtils.h
index 92a519cd0c838..62438c7c44dd3 100644
--- a/flang/include/flang/Optimizer/Transforms/MemoryUtils.h
+++ b/flang/include/flang/Optimizer/Transforms/MemoryUtils.h
@@ -57,6 +57,13 @@ bool replaceAllocas(mlir::RewriterBase &rewriter, mlir::Operation *parentOp,
MustRewriteCallBack, AllocaRewriterCallBack,
DeallocCallBack);
+/// Under -gpu=mem:unified|managed, move the dynamically sized fir.alloca of
+/// \p func to fir.allocmem/fir.freemem pairs marked for the unified/managed
+/// allocator. Does nothing for device code, which keeps its stack allocations.
+/// Returns true if the function was modified.
+bool promoteDynamicAllocasToCudaHeap(mlir::RewriterBase &rewriter,
+ mlir::Operation *func);
+
} // namespace fir
#endif // FORTRAN_OPTIMIZER_TRANSFORMS_MEMORYUTILS_H
diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index ed8b256f47fd4..22da804d879b5 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -6991,6 +6991,13 @@ Fortran::lower::LoweringBridge::LoweringBridge(
fir::setIsPIE(*module, cgOpts.IsPIE);
if (cgOpts.RecordCommandLine)
fir::setCommandline(*module, *cgOpts.RecordCommandLine);
+ // Under -gpu=mem:unified|managed, host heap allocations use the matching
+ // indirect runtime allocators (malloc_unified / malloc_managed).
+ if (languageFeatures.IsEnabled(Fortran::common::LanguageFeature::CudaUnified))
+ fir::setCudaHeapAllocMode(*module, fir::CudaHeapAllocMode::Unified);
+ else if (languageFeatures.IsEnabled(
+ Fortran::common::LanguageFeature::CudaManaged))
+ fir::setCudaHeapAllocMode(*module, fir::CudaHeapAllocMode::Managed);
}
Fortran::lower::LoweringBridge::~LoweringBridge() {
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 0e88524126f11..a9c6cef6624f3 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -22,6 +22,7 @@
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/Dialect/Support/FIRContext.h"
#include "flang/Optimizer/Support/DataLayout.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Support/TypeCode.h"
@@ -1300,8 +1301,7 @@ template <typename ModuleOp>
static mlir::SymbolRefAttr
getMallocInModule(ModuleOp mod, fir::AllocMemOp op,
mlir::ConversionPatternRewriter &rewriter,
- mlir::Type indexType) {
- static constexpr char mallocName[] = "malloc";
+ mlir::Type indexType, llvm::StringRef mallocName) {
if (auto mallocFunc =
mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(mallocName))
return mlir::SymbolRefAttr::get(mallocFunc);
@@ -1318,22 +1318,61 @@ getMallocInModule(ModuleOp mod, fir::AllocMemOp op,
return mlir::SymbolRefAttr::get(mallocDecl);
}
+/// Allocator entry points for an allocation marked by the allocation placement
+/// passes under -gpu=mem:unified|managed. Only marked fir.allocmem/fir.freemem
+/// pairs are routed: memory that the Fortran runtime allocated must keep being
+/// released by libc free, and vice versa.
+static llvm::StringRef getHeapAllocName(mlir::Operation *op,
+ llvm::StringRef plain,
+ llvm::StringRef unified,
+ llvm::StringRef managed) {
+ // Device modules keep libc names; the indirect entry points are host-side.
+ if (op->getParentOfType<mlir::gpu::GPUModuleOp>())
+ return plain;
+ switch (fir::getCudaHeapAllocMode(op)) {
+ case fir::CudaHeapAllocMode::Unified:
+ return unified;
+ case fir::CudaHeapAllocMode::Managed:
+ return managed;
+ case fir::CudaHeapAllocMode::None:
+ return plain;
+ }
+ llvm_unreachable("unexpected CudaHeapAllocMode");
+}
+
+static llvm::StringRef getHostHeapMallocName(mlir::Operation *op) {
+ return getHeapAllocName(op, "malloc", "malloc_unified", "malloc_managed");
+}
+
+static llvm::StringRef getHostHeapFreeName(mlir::Operation *op) {
+ return getHeapAllocName(op, "free", "free_unified", "free_managed");
+}
+
+static llvm::StringRef getHostHeapAlignedAllocName(mlir::Operation *op) {
+ return getHeapAllocName(op, "aligned_alloc", "aligned_alloc_unified",
+ "aligned_alloc_managed");
+}
+
+static llvm::StringRef getHostHeapPosixMemalignName(mlir::Operation *op) {
+ return getHeapAllocName(op, "posix_memalign", "posix_memalign_unified",
+ "posix_memalign_managed");
+}
+
/// Return the LLVMFuncOp corresponding to the standard malloc call.
static mlir::SymbolRefAttr getMalloc(fir::AllocMemOp op,
mlir::ConversionPatternRewriter &rewriter,
mlir::Type indexType) {
+ llvm::StringRef name = getHostHeapMallocName(op);
if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
- return getMallocInModule(mod, op, rewriter, indexType);
+ return getMallocInModule(mod, op, rewriter, indexType, name);
auto mod = op->getParentOfType<mlir::ModuleOp>();
- return getMallocInModule(mod, op, rewriter, indexType);
+ return getMallocInModule(mod, op, rewriter, indexType, name);
}
template <typename ModuleOp>
-static mlir::SymbolRefAttr
-getAlignedAllocInModule(ModuleOp mod, fir::AllocMemOp op,
- mlir::ConversionPatternRewriter &rewriter,
- mlir::Type indexType) {
- static constexpr char alignedAllocName[] = "aligned_alloc";
+static mlir::SymbolRefAttr getAlignedAllocInModule(
+ ModuleOp mod, fir::AllocMemOp op, mlir::ConversionPatternRewriter &rewriter,
+ mlir::Type indexType, llvm::StringRef alignedAllocName) {
if (auto func =
mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(alignedAllocName))
return mlir::SymbolRefAttr::get(func);
@@ -1353,18 +1392,17 @@ getAlignedAllocInModule(ModuleOp mod, fir::AllocMemOp op,
static mlir::SymbolRefAttr
getAlignedAlloc(fir::AllocMemOp op, mlir::ConversionPatternRewriter &rewriter,
mlir::Type indexType) {
+ llvm::StringRef name = getHostHeapAlignedAllocName(op);
if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
- return getAlignedAllocInModule(mod, op, rewriter, indexType);
+ return getAlignedAllocInModule(mod, op, rewriter, indexType, name);
auto mod = op->getParentOfType<mlir::ModuleOp>();
- return getAlignedAllocInModule(mod, op, rewriter, indexType);
+ return getAlignedAllocInModule(mod, op, rewriter, indexType, name);
}
template <typename ModuleOp>
-static mlir::SymbolRefAttr
-getPosixMemalignInModule(ModuleOp mod, fir::AllocMemOp op,
- mlir::ConversionPatternRewriter &rewriter,
- mlir::Type indexType) {
- static constexpr char posixMemalignName[] = "posix_memalign";
+static mlir::SymbolRefAttr getPosixMemalignInModule(
+ ModuleOp mod, fir::AllocMemOp op, mlir::ConversionPatternRewriter &rewriter,
+ mlir::Type indexType, llvm::StringRef posixMemalignName) {
if (auto func =
mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(posixMemalignName))
return mlir::SymbolRefAttr::get(func);
@@ -1386,10 +1424,11 @@ getPosixMemalignInModule(ModuleOp mod, fir::AllocMemOp op,
static mlir::SymbolRefAttr
getPosixMemalign(fir::AllocMemOp op, mlir::ConversionPatternRewriter &rewriter,
mlir::Type indexType) {
+ llvm::StringRef name = getHostHeapPosixMemalignName(op);
if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
- return getPosixMemalignInModule(mod, op, rewriter, indexType);
+ return getPosixMemalignInModule(mod, op, rewriter, indexType, name);
auto mod = op->getParentOfType<mlir::ModuleOp>();
- return getPosixMemalignInModule(mod, op, rewriter, indexType);
+ return getPosixMemalignInModule(mod, op, rewriter, indexType, name);
}
/// Return value of the stride in bytes between adjacent elements
@@ -1526,8 +1565,8 @@ struct AllocMemOpConversion : public fir::FIROpConversion<fir::AllocMemOp> {
template <typename ModuleOp>
static mlir::SymbolRefAttr
getFreeInModule(ModuleOp mod, fir::FreeMemOp op,
- mlir::ConversionPatternRewriter &rewriter) {
- static constexpr char freeName[] = "free";
+ mlir::ConversionPatternRewriter &rewriter,
+ llvm::StringRef freeName) {
// Check if free already defined in the module.
if (auto freeFunc =
mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(freeName))
@@ -1548,10 +1587,11 @@ getFreeInModule(ModuleOp mod, fir::FreeMemOp op,
static mlir::SymbolRefAttr getFree(fir::FreeMemOp op,
mlir::ConversionPatternRewriter &rewriter) {
+ llvm::StringRef name = getHostHeapFreeName(op);
if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
- return getFreeInModule(mod, op, rewriter);
+ return getFreeInModule(mod, op, rewriter, name);
auto mod = op->getParentOfType<mlir::ModuleOp>();
- return getFreeInModule(mod, op, rewriter);
+ return getFreeInModule(mod, op, rewriter, name);
}
static unsigned getDimension(mlir::LLVM::LLVMArrayType ty) {
diff --git a/flang/lib/Optimizer/Dialect/Support/FIRContext.cpp b/flang/lib/Optimizer/Dialect/Support/FIRContext.cpp
index 16757f934c8d7..66ae5c4c653ac 100644
--- a/flang/lib/Optimizer/Dialect/Support/FIRContext.cpp
+++ b/flang/lib/Optimizer/Dialect/Support/FIRContext.cpp
@@ -248,6 +248,47 @@ void fir::setIsPIE(mlir::ModuleOp mod, bool value) {
bool fir::getIsPIE(mlir::ModuleOp mod) { return mod->hasAttr(isPIEName); }
+static constexpr const char *cudaHeapAllocModeName = "fir.cuda_heap_alloc";
+
+static void setCudaHeapAllocModeOn(mlir::Operation *op,
+ fir::CudaHeapAllocMode mode) {
+ if (mode == fir::CudaHeapAllocMode::None) {
+ if (op->hasAttr(cudaHeapAllocModeName))
+ op->removeAttr(cudaHeapAllocModeName);
+ return;
+ }
+ llvm::StringRef value =
+ mode == fir::CudaHeapAllocMode::Unified ? "unified" : "managed";
+ op->setAttr(cudaHeapAllocModeName,
+ mlir::StringAttr::get(op->getContext(), value));
+}
+
+static fir::CudaHeapAllocMode getCudaHeapAllocModeOf(mlir::Operation *op) {
+ if (auto attr = op->getAttrOfType<mlir::StringAttr>(cudaHeapAllocModeName)) {
+ if (attr.getValue() == "unified")
+ return fir::CudaHeapAllocMode::Unified;
+ if (attr.getValue() == "managed")
+ return fir::CudaHeapAllocMode::Managed;
+ }
+ return fir::CudaHeapAllocMode::None;
+}
+
+void fir::setCudaHeapAllocMode(mlir::ModuleOp mod, CudaHeapAllocMode mode) {
+ setCudaHeapAllocModeOn(mod.getOperation(), mode);
+}
+
+fir::CudaHeapAllocMode fir::getCudaHeapAllocMode(mlir::ModuleOp mod) {
+ return getCudaHeapAllocModeOf(mod.getOperation());
+}
+
+void fir::setCudaHeapAllocMode(mlir::Operation *op, CudaHeapAllocMode mode) {
+ setCudaHeapAllocModeOn(op, mode);
+}
+
+fir::CudaHeapAllocMode fir::getCudaHeapAllocMode(mlir::Operation *op) {
+ return getCudaHeapAllocModeOf(op);
+}
+
std::string fir::determineTargetTriple(llvm::StringRef triple) {
// Treat "" or "default" as stand-ins for the default machine.
if (triple.empty() || triple == "default")
diff --git a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
index d070376dd5637..2e14bf2960473 100644
--- a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
+++ b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
@@ -198,6 +198,11 @@ void AllocationPlacementPass::runOnOperation() {
if (func.empty())
return;
+ // Done first: the pairs it creates are marked fir.must_be_heap, so the
+ // placement decisions below leave them alone.
+ mlir::IRRewriter cudaHeapRewriter(&getContext());
+ fir::promoteDynamicAllocasToCudaHeap(cudaHeapRewriter, func.getOperation());
+
fir::AllocationPlacementThresholds baseThresholds;
baseThresholds.stackArrays = stackArrays;
baseThresholds.smallArrayThresholdBytes = smallArrayThresholdBytes;
diff --git a/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp b/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp
index fd1d566ca2825..6c4342b141050 100644
--- a/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp
+++ b/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp
@@ -9,6 +9,7 @@
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/Dialect/Support/FIRContext.h"
#include "flang/Optimizer/Transforms/MemoryUtils.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -131,6 +132,8 @@ class MemoryAllocationOpt
// If func is a declaration, skip it.
if (func.empty())
return;
+ mlir::IRRewriter cudaHeapRewriter(context);
+ fir::promoteDynamicAllocasToCudaHeap(cudaHeapRewriter, func.getOperation());
auto tryReplacing = [&](fir::AllocaOp alloca) {
bool res = !keepStackAllocation(alloca, options);
if (res) {
diff --git a/flang/lib/Optimizer/Transforms/MemoryUtils.cpp b/flang/lib/Optimizer/Transforms/MemoryUtils.cpp
index d1b27d9872ea1..0e5155dd0a361 100644
--- a/flang/lib/Optimizer/Transforms/MemoryUtils.cpp
+++ b/flang/lib/Optimizer/Transforms/MemoryUtils.cpp
@@ -8,7 +8,12 @@
#include "flang/Optimizer/Transforms/MemoryUtils.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
+#include "flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h"
+#include "flang/Optimizer/Dialect/FIRAttr.h"
+#include "flang/Optimizer/Dialect/Support/FIRContext.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/Dominance.h"
#include "llvm/ADT/STLExtras.h"
@@ -309,3 +314,63 @@ bool fir::replaceAllocas(mlir::RewriterBase &rewriter,
rewriter.restoreInsertionPoint(insertPoint);
return replacedAllRequestedAlloca;
}
+
+/// Device code keeps its stack allocations: the unified/managed entry points
+/// are host-only, and a kernel-side heap allocation would be a large
+/// regression over a device stack array.
+static bool isDeviceCode(mlir::Operation *func, mlir::ModuleOp mod) {
+ if (func->getParentOfType<mlir::gpu::GPUModuleOp>())
+ return true;
+ if (auto procAttr =
+ func->getAttrOfType<cuf::ProcAttributeAttr>(cuf::getProcAttrName()))
+ return procAttr.getValue() != cuf::ProcAttribute::Host;
+ if (mlir::acc::isAccRoutine(func))
+ return true;
+ if (auto offloadMod =
+ llvm::dyn_cast<mlir::omp::OffloadModuleInterface>(mod.getOperation()))
+ return offloadMod.getIsTargetDevice();
+ return false;
+}
+
+bool fir::promoteDynamicAllocasToCudaHeap(mlir::RewriterBase &rewriter,
+ mlir::Operation *func) {
+ auto mod = func->getParentOfType<mlir::ModuleOp>();
+ if (!mod)
+ return false;
+ fir::CudaHeapAllocMode mode = fir::getCudaHeapAllocMode(mod);
+ if (mode == fir::CudaHeapAllocMode::None || isDeviceCode(func, mod))
+ return false;
+
+ bool changed = false;
+ // Named locals only: automatic arrays and automatic character. Compiler
+ // temporaries do not need unified memory and would turn a stack save/restore
+ // into a malloc/free pair, possibly per loop iteration.
+ auto mustReplace = [](fir::AllocaOp alloca) {
+ std::optional<llvm::StringRef> uniqName = alloca.getUniqName();
+ return alloca.isDynamic() && uniqName && !uniqName->empty();
+ };
+ auto genAllocmem = [&](mlir::OpBuilder &builder, fir::AllocaOp alloca,
+ bool) -> mlir::Value {
+ auto name = [](std::optional<llvm::StringRef> opt) -> llvm::StringRef {
+ return opt ? *opt : llvm::StringRef{};
+ };
+ auto heap = fir::AllocMemOp::create(
+ builder, alloca.getLoc(), alloca.getInType(),
+ name(alloca.getUniqName()), name(alloca.getBindcName()),
+ alloca.getTypeparams(), alloca.getShape());
+ fir::setCudaHeapAllocMode(heap.getOperation(), mode);
+ // Keep the placement passes from sinking it back to the stack: the
+ // allocator is chosen here and the matching free is emitted below.
+ heap->setAttr(fir::MustBeHeapAttr::getAttrName(),
+ fir::MustBeHeapAttr::get(builder.getContext(), true));
+ changed = true;
+ return heap;
+ };
+ auto genFreemem = [&](mlir::Location loc, mlir::OpBuilder &builder,
+ mlir::Value allocmem) {
+ auto free = fir::FreeMemOp::create(builder, loc, allocmem);
+ fir::setCudaHeapAllocMode(free.getOperation(), mode);
+ };
+ fir::replaceAllocas(rewriter, func, mustReplace, genAllocmem, genFreemem);
+ return changed;
+}
diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp
index 77861e67a07b1..dbd4f5f08b701 100644
--- a/flang/lib/Optimizer/Transforms/StackArrays.cpp
+++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp
@@ -15,6 +15,7 @@
#include "flang/Optimizer/Dialect/FIRType.h"
#include "flang/Optimizer/Dialect/Support/FIRContext.h"
#include "flang/Optimizer/Support/DataLayout.h"
+#include "flang/Optimizer/Transforms/MemoryUtils.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h"
#include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h"
@@ -746,6 +747,12 @@ llvm::StringRef StackArraysPass::getDescription() const {
void StackArraysPass::runOnOperation() {
mlir::func::FuncOp func = getOperation();
+ // -fstack-arrays does not apply to the automatic arrays that
+ // -gpu=mem:unified|managed must place in unified/managed memory. Done before
+ // the analysis below, which skips the fir.must_be_heap pairs it creates.
+ mlir::IRRewriter cudaHeapRewriter(&getContext());
+ fir::promoteDynamicAllocasToCudaHeap(cudaHeapRewriter, func.getOperation());
+
auto &analysis = getAnalysis<fir::StackArraysAnalysisWrapper>();
const fir::StackArraysAnalysisWrapper::AllocMemMap *candidateOps =
analysis.getCandidateOps(func);
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index a3514ccf72e5c..83d92f253e626 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -10648,37 +10648,6 @@ void ResolveNamesVisitor::FinishSpecificationPart(
context().languageFeatures().IsEnabled(
common::LanguageFeature::CudaPinned))
object->set_cudaDataAttr(common::CUDADataAttr::Pinned);
- } else if (!object->cudaDataAttr() && !IsDummy(symbol) &&
- !IsAllocatable(symbol) && !IsPointer(symbol) && !IsSaved(symbol) &&
- !IsCUDADeviceContext(&symbol.owner()) &&
- object->shape().IsExplicitShape()) {
- // Under -gpu=mem:unified|managed, allocate adjustable / VLA automatic
- // arrays in CUDA unified/managed memory (fixed-size automatic arrays
- // stay on the stack). Tag those locals so lowering uses
- // cuf.alloc/cuf.free. Unlike the allocatable managed tagging above,
- // this does not require -fcuda: OpenACC + -gpu=mem:unified relies on
- // it, and cuf.alloc does not go through the CUDA Fortran
- // managed-descriptor pipeline that motivated the -fcuda gate.
- std::optional<common::CUDADataAttr> attr;
- if (context().languageFeatures().IsEnabled(
- common::LanguageFeature::CudaUnified))
- attr = common::CUDADataAttr::Unified;
- else if (context().languageFeatures().IsEnabled(
- common::LanguageFeature::CudaManaged))
- attr = common::CUDADataAttr::Managed;
- if (attr) {
- auto boundIsNonConstant{[](const Bound &b) {
- return !b.isExplicit() || !b.GetExplicit() ||
- !evaluate::IsConstantExpr(*b.GetExplicit());
- }};
- for (const ShapeSpec &ss : object->shape()) {
- if (boundIsNonConstant(ss.lbound()) ||
- boundIsNonConstant(ss.ubound())) {
- object->set_cudaDataAttr(*attr);
- break;
- }
- }
- }
}
}
}
diff --git a/flang/test/Fir/CUDA/cuda-heap-alloc-managed.fir b/flang/test/Fir/CUDA/cuda-heap-alloc-managed.fir
new file mode 100644
index 0000000000000..c6fc559ce5088
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-heap-alloc-managed.fir
@@ -0,0 +1,44 @@
+// RUN: fir-opt --memory-allocation-opt %s | FileCheck %s --check-prefix=HEAP
+// RUN: fir-opt --fir-to-llvm-ir %s | FileCheck %s --check-prefix=LLVM
+
+// Same routing as cuda-heap-alloc-unified.fir, with the managed entry points.
+
+// Declarations are emitted at the top of the module, before any function.
+// LLVM-DAG: llvm.func @malloc_managed(i64) -> !llvm.ptr
+// LLVM-DAG: llvm.func @free_managed(!llvm.ptr)
+// LLVM-DAG: llvm.func @malloc(i64) -> !llvm.ptr
+// LLVM-DAG: llvm.func @free(!llvm.ptr)
+
+module attributes {fir.cuda_heap_alloc = "managed"} {
+
+// HEAP-LABEL: func.func @vla(
+// HEAP: %[[MEM:.*]] = fir.allocmem !fir.array<?xf32>, %{{.*}} {bindc_name = "a", fir.cuda_heap_alloc = "managed", fir.must_be_heap = true, uniq_name = "_QFvlaEa"}
+// HEAP: fir.freemem %[[MEM]] {fir.cuda_heap_alloc = "managed"} : !fir.heap<!fir.array<?xf32>>
+func.func @vla(%arg0: !fir.ref<i32>) {
+ %0 = fir.load %arg0 : !fir.ref<i32>
+ %1 = fir.convert %0 : (i32) -> index
+ %2 = fir.alloca !fir.array<?xf32>, %1 {bindc_name = "a", uniq_name = "_QFvlaEa"}
+ return
+}
+
+// LLVM-LABEL: llvm.func @marked_heap(
+// LLVM: llvm.call @malloc_managed(
+// LLVM: llvm.call @free_managed(
+func.func @marked_heap(%n: index) {
+ %0 = fir.allocmem !fir.array<?xf32>, %n {fir.cuda_heap_alloc = "managed"}
+ fir.freemem %0 {fir.cuda_heap_alloc = "managed"} : !fir.heap<!fir.array<?xf32>>
+ return
+}
+
+// LLVM-LABEL: llvm.func @unmarked_heap(
+// LLVM-NOT: llvm.call @malloc_managed(
+// LLVM: llvm.call @malloc(
+// LLVM-NOT: llvm.call @free_managed(
+// LLVM: llvm.call @free(
+func.func @unmarked_heap(%n: index) {
+ %0 = fir.allocmem !fir.array<?xf32>, %n
+ fir.freemem %0 : !fir.heap<!fir.array<?xf32>>
+ return
+}
+
+}
diff --git a/flang/test/Fir/CUDA/cuda-heap-alloc-unified.fir b/flang/test/Fir/CUDA/cuda-heap-alloc-unified.fir
new file mode 100644
index 0000000000000..7a043cf9737a4
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-heap-alloc-unified.fir
@@ -0,0 +1,86 @@
+// RUN: fir-opt --memory-allocation-opt %s | FileCheck %s --check-prefix=HEAP
+// RUN: fir-opt --stack-arrays %s | FileCheck %s --check-prefix=HEAP
+// RUN: fir-opt --fir-to-llvm-ir %s | FileCheck %s --check-prefix=LLVM
+
+// Under fir.cuda_heap_alloc = "unified", named automatic arrays move to the
+// heap and are marked. Only marked allocations use malloc_unified: memory the
+// Fortran runtime allocated must keep being released by libc free.
+
+// Declarations are emitted at the top of the module, before any function.
+// LLVM-DAG: llvm.func @malloc_unified(i64) -> !llvm.ptr
+// LLVM-DAG: llvm.func @free_unified(!llvm.ptr)
+// LLVM-DAG: llvm.func @malloc(i64) -> !llvm.ptr
+// LLVM-DAG: llvm.func @free(!llvm.ptr)
+
+module attributes {fir.cuda_heap_alloc = "unified"} {
+
+// HEAP-LABEL: func.func @vla(
+// HEAP: %[[MEM:.*]] = fir.allocmem !fir.array<?xf32>, %{{.*}} {bindc_name = "a", fir.cuda_heap_alloc = "unified", fir.must_be_heap = true, uniq_name = "_QFvlaEa"}
+// HEAP: fir.freemem %[[MEM]] {fir.cuda_heap_alloc = "unified"} : !fir.heap<!fir.array<?xf32>>
+func.func @vla(%arg0: !fir.ref<i32>) {
+ %0 = fir.load %arg0 : !fir.ref<i32>
+ %1 = fir.convert %0 : (i32) -> index
+ %2 = fir.alloca !fir.array<?xf32>, %1 {bindc_name = "a", uniq_name = "_QFvlaEa"}
+ return
+}
+
+// Automatic character is an automatic too.
+// HEAP-LABEL: func.func @autochar(
+// HEAP: fir.allocmem !fir.char<1,?>(%{{.*}} : index) {{{.*}}fir.cuda_heap_alloc = "unified"
+func.func @autochar(%arg0: index) {
+ %0 = fir.alloca !fir.char<1,?>(%arg0 : index) {bindc_name = "s", uniq_name = "_QFautocharEs"}
+ return
+}
+
+// Fixed-size automatics stay on the stack.
+// HEAP-LABEL: func.func @fixed(
+// HEAP: fir.alloca !fir.array<128xf32>
+// HEAP-NOT: fir.allocmem
+func.func @fixed() {
+ %0 = fir.alloca !fir.array<128xf32> {bindc_name = "a", uniq_name = "_QFfixedEa"}
+ return
+}
+
+// Compiler temporaries are not automatics: turning them into malloc/free would
+// cost an allocation per loop iteration.
+// HEAP-LABEL: func.func @anon_temp(
+// HEAP: fir.alloca !fir.array<?xf32>
+// HEAP-NOT: fir.allocmem
+func.func @anon_temp(%arg0: index) {
+ %0 = fir.alloca !fir.array<?xf32>, %arg0
+ return
+}
+
+// Device code keeps its stack allocation: the entry points are host-only.
+// HEAP-LABEL: func.func @device_vla(
+// HEAP: fir.alloca !fir.array<?xf32>
+// HEAP-NOT: fir.allocmem
+func.func @device_vla(%arg0: index) attributes {cuf.proc_attr = #cuf.cuda_proc<global>} {
+ %0 = fir.alloca !fir.array<?xf32>, %arg0 {bindc_name = "a", uniq_name = "_QFdevice_vlaEa"}
+ return
+}
+
+// HEAP-LABEL: func.func @marked_heap(
+// LLVM-LABEL: llvm.func @marked_heap(
+// LLVM: llvm.call @malloc_unified(
+// LLVM: llvm.call @free_unified(
+func.func @marked_heap(%n: index) {
+ %0 = fir.allocmem !fir.array<?xf32>, %n {fir.cuda_heap_alloc = "unified"}
+ fir.freemem %0 {fir.cuda_heap_alloc = "unified"} : !fir.heap<!fir.array<?xf32>>
+ return
+}
+
+// An unmarked pair belongs to libc: this is the shape of a buffer the Fortran
+// runtime allocated and lowered code releases.
+// LLVM-LABEL: llvm.func @unmarked_heap(
+// LLVM-NOT: llvm.call @malloc_unified(
+// LLVM: llvm.call @malloc(
+// LLVM-NOT: llvm.call @free_unified(
+// LLVM: llvm.call @free(
+func.func @unmarked_heap(%n: index) {
+ %0 = fir.allocmem !fir.array<?xf32>, %n
+ fir.freemem %0 : !fir.heap<!fir.array<?xf32>>
+ return
+}
+
+}
diff --git a/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90 b/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
index 07638b0ca8dfa..c3c15c98ceed0 100644
--- a/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
+++ b/flang/test/Lower/CUDA/cuda-gpu-unified-automatic-array.f90
@@ -1,21 +1,23 @@
-! RUN: bbc -emit-hlfir -gpu=unified %s -o - | FileCheck %s
-! RUN: bbc -emit-hlfir -gpu=managed %s -o - | FileCheck %s --check-prefix=MANAGED
+! RUN: bbc -emit-hlfir -gpu=unified %s -o - | FileCheck %s --check-prefixes=CHECK,UNIFIED
+! RUN: bbc -emit-hlfir -gpu=managed %s -o - | FileCheck %s --check-prefixes=CHECK,MANAGED
+! RUN: bbc -emit-hlfir %s -o - | FileCheck %s --check-prefixes=CHECK,NOFLAG
-! Under -gpu=mem:unified|managed, allocate adjustable / VLA automatic arrays
-! in CUDA unified/managed memory. Fixed-size automatic arrays stay on the
-! stack so host pointers remain shared under unified memory (e.g. OpenACC).
+! Under -gpu=mem:unified|managed, dynamic automatic arrays are later moved to
+! the heap and allocated with malloc_unified / malloc_managed. Lowering only
+! records the mode on the module; symbols stay unmarked (no cudaDataAttr).
+
+! UNIFIED: module attributes {{{.*}}fir.cuda_heap_alloc = "unified"
+! MANAGED: module attributes {{{.*}}fir.cuda_heap_alloc = "managed"
+! NOFLAG-NOT: fir.cuda_heap_alloc
module m_adj
integer :: nx = 32
end module
! CHECK-LABEL: func.func @_QPvla(
-! CHECK: %[[ALLOC:.*]] = cuf.alloc !fir.array<?xf32>, %{{.*}} : index {bindc_name = "a", data_attr = #cuf.cuda<unified>, uniq_name = "_QFvlaEa"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOC]](%{{.*}}) {data_attr = #cuf.cuda<unified>, uniq_name = "_QFvlaEa"}
-! CHECK: cuf.free %[[DECL]]#1 : !fir.ref<!fir.array<?xf32>> {data_attr = #cuf.cuda<unified>}
-! MANAGED-LABEL: func.func @_QPvla(
-! MANAGED: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<managed>
-! MANAGED: cuf.free %{{.*}} : !fir.ref<!fir.array<?xf32>> {data_attr = #cuf.cuda<managed>}
+! CHECK-NOT: cuf.alloc
+! CHECK-NOT: data_attr = #cuf.cuda
+! CHECK: fir.alloca !fir.array<?xf32>
subroutine vla(n)
integer :: n
real :: a(n)
@@ -23,34 +25,28 @@ subroutine vla(n)
end subroutine
! CHECK-LABEL: func.func @_QPadjustable(
-! CHECK: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<unified>
-! CHECK: cuf.free %{{.*}} {data_attr = #cuf.cuda<unified>}
-! MANAGED-LABEL: func.func @_QPadjustable(
-! MANAGED: cuf.alloc !fir.array<?xf32>, %{{.*}} : index {{{.*}}data_attr = #cuf.cuda<managed>
+! CHECK-NOT: cuf.alloc
+! CHECK-NOT: data_attr = #cuf.cuda
+! CHECK: fir.alloca !fir.array<?xf32>
subroutine adjustable
use m_adj
real :: a(0:(nx+1)/2)
a(0) = 0.0
end subroutine
-! Fixed-size automatic arrays must remain ordinary stack allocations.
+! Fixed-size automatic arrays remain ordinary stack allocations.
! CHECK-LABEL: func.func @_QPfixed(
! CHECK-NOT: cuf.alloc
! CHECK: fir.alloca !fir.array<128xf32>
-! CHECK-NOT: cuf.free
-! MANAGED-LABEL: func.func @_QPfixed(
-! MANAGED-NOT: cuf.alloc
-! MANAGED: fir.alloca !fir.array<128xf32>
subroutine fixed
real :: a(128)
a(1) = 1.0
end subroutine
-! Dummy adjustable arrays are caller-allocated; do not retag them.
+! Dummy adjustable arrays are caller-allocated.
! CHECK-LABEL: func.func @_QPdummy_adj(
! CHECK-NOT: cuf.alloc
-! MANAGED-LABEL: func.func @_QPdummy_adj(
-! MANAGED-NOT: cuf.alloc
+! CHECK-NOT: fir.alloca !fir.array<?xf32>
subroutine dummy_adj(a, n)
integer :: n
real :: a(n)
More information about the flang-commits
mailing list