[flang-commits] [flang] [flang] Do not honor -fstack-arrays inside offload regions (PR #227537)
Zhen Wang via flang-commits
flang-commits at lists.llvm.org
Wed Sep 30 09:55:16 PDT 2026
https://github.com/wangzpgi updated https://github.com/llvm/llvm-project/pull/227537
>From c8fbc72535fb14f02c77dc596edc2126c6e011ab Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Tue, 29 Sep 2026 14:40:39 -0700
Subject: [PATCH 1/6] [flang] Do not honor -fstack-arrays inside offload
regions
---
.../Transforms/AllocationPlacement.cpp | 12 +-
.../lib/Optimizer/Transforms/StackArrays.cpp | 14 ++-
.../allocation-placement-offload-region.fir | 106 ++++++++++++++++++
.../Transforms/stack-arrays-alloca-scope.fir | 27 +++--
.../stack-arrays-offload-region.fir | 38 +++++++
5 files changed, 181 insertions(+), 16 deletions(-)
create mode 100644 flang/test/Transforms/allocation-placement-offload-region.fir
create mode 100644 flang/test/Transforms/stack-arrays-offload-region.fir
diff --git a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
index 1e6c7f900159d..de6028b92dae5 100644
--- a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
+++ b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
@@ -17,6 +17,7 @@
//===----------------------------------------------------------------------===//
#include "StackArrays.h"
+#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Dialect/FIRAttr.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
@@ -207,12 +208,19 @@ void AllocationPlacementPass::runOnOperation() {
: (allocmem.hasLenParams() || allocmem.hasShapeOperands());
info.byteSize = getConstantByteSize(op, dl, kindMap);
+ // -fstack-arrays cannot be honored in an offload region either: like a
+ // device procedure, it runs on the device stack, which is far smaller than
+ // the host one. The size based part of the policy still applies.
+ fir::AllocationPolicy policy = basePolicy;
+ if (policy.stackArrays && cuf::isExecutingOnDevice(op))
+ policy.stackArrays = false;
+
// A hook, if provided, fully overrides the default policy; it may delegate
// back to decideAllocationPlacement after adjusting the policy.
fir::AllocationPlacement placement =
placementHook
- ? placementHook(info, basePolicy, stackBytesUsed)
- : fir::decideAllocationPlacement(info, basePolicy, stackBytesUsed);
+ ? placementHook(info, policy, stackBytesUsed)
+ : fir::decideAllocationPlacement(info, policy, stackBytesUsed);
// Account for the decision in the running stack budget.
if (endsUpOnStack(placement, info.isCurrentlyOnStack) && info.byteSize)
diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp
index 7d4e6d49642d3..575c853b0bc21 100644
--- a/flang/lib/Optimizer/Transforms/StackArrays.cpp
+++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "StackArrays.h"
+#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/LowLevelIntrinsics.h"
#include "flang/Optimizer/Dialect/FIRAttr.h"
@@ -791,14 +792,17 @@ void StackArraysPass::runOnOperation() {
return;
}
- if (candidateOps->empty())
- return;
- runCount += candidateOps->size();
-
+ // An offload region runs on the device stack, which is far smaller than the
+ // host one, so its allocations stay on the heap like in a device procedure.
llvm::SmallVector<mlir::Operation *> opsToConvert;
opsToConvert.reserve(candidateOps->size());
for (auto [op, _] : *candidateOps)
- opsToConvert.push_back(op);
+ if (!cuf::isExecutingOnDevice(op))
+ opsToConvert.push_back(op);
+
+ if (opsToConvert.empty())
+ return;
+ runCount += opsToConvert.size();
mlir::MLIRContext &context = getContext();
mlir::RewritePatternSet patterns(&context);
diff --git a/flang/test/Transforms/allocation-placement-offload-region.fir b/flang/test/Transforms/allocation-placement-offload-region.fir
new file mode 100644
index 0000000000000..c0fd11c15b12e
--- /dev/null
+++ b/flang/test/Transforms/allocation-placement-offload-region.fir
@@ -0,0 +1,106 @@
+// Test that -fstack-arrays is not honored inside an offload region: the region
+// runs on the device stack, which is far smaller than the host one, so its
+// runtime-sized temporaries stay on the heap while the same allocation in host
+// code goes on the stack. The size based part of the policy still applies, so
+// a small constant-size temporary goes on the device stack.
+
+// RUN: fir-opt --allocation-placement %s | FileCheck %s
+
+module attributes {fir.allocation_policy =
+ #fir.allocation_policy<stack_arrays = true,
+ small_array_threshold = 1024,
+ total_stack_limit = 4194304>,
+ fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "",
+ llvm.data_layout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"} {
+
+// CHECK-LABEL: func.func @dynamic_temp_in_host_code
+// CHECK: fir.alloca !fir.array<?xi32>, %{{.*}}
+// CHECK-NOT: fir.allocmem
+func.func @dynamic_temp_in_host_code(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ return
+}
+
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_kernels
+// CHECK: acc.kernels {
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_kernels(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ acc.kernels {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ acc.terminator
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_parallel
+// CHECK: acc.parallel {
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_parallel(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ acc.parallel {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ acc.yield
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_serial
+// CHECK: acc.serial {
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_serial(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ acc.serial {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ acc.yield
+ }
+ return
+}
+
+// <10xi32> is 40 bytes, under the 1024 byte threshold: the size based policy
+// still puts it on the (device) stack.
+// CHECK-LABEL: func.func @small_temp_in_acc_parallel
+// CHECK: acc.parallel {
+// CHECK-NEXT: fir.alloca !fir.array<10xi32>
+// CHECK-NOT: fir.allocmem
+func.func @small_temp_in_acc_parallel() {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ acc.parallel {
+ %0 = fir.allocmem !fir.array<10xi32>
+ %r = fir.convert %0 : (!fir.heap<!fir.array<10xi32>>) -> !fir.ref<!fir.array<10xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<10xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<10xi32>>
+ acc.yield
+ }
+ return
+}
+}
diff --git a/flang/test/Transforms/stack-arrays-alloca-scope.fir b/flang/test/Transforms/stack-arrays-alloca-scope.fir
index 342f09ccfaf6d..9debacd40dd03 100644
--- a/flang/test/Transforms/stack-arrays-alloca-scope.fir
+++ b/flang/test/Transforms/stack-arrays-alloca-scope.fir
@@ -1,9 +1,15 @@
-// RUN: fir-opt --stack-arrays %s | FileCheck %s
+// RUN: fir-opt --allocation-placement %s | FileCheck %s
// Test that stack allocations are created in the block where the enclosing
// construct expects them (fir::getAllocaBlock) and are not hoisted out of it:
// each concurrent execution of a construct modelling parallelism needs its own
// storage.
+//
+// -fstack-arrays is not honored inside offload regions, so the default size
+// based policy is used: the 168 byte <42xi32> temporaries go on the stack,
+// runtime-sized ones stay on the heap.
+
+module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"} {
// The allocation has no operand but must still stay inside the compute
// construct instead of being hoisted to the function entry block.
@@ -63,8 +69,8 @@ func.func @acc_kernels_no_operand() {
// CHECK: acc.kernels {
// CHECK-NEXT: fir.alloca !fir.array<42xi32>
-// The extent operand is available before the compute construct: the allocation
-// can only be hoisted up to the beginning of the construct.
+// A runtime-sized allocation inside a compute construct stays on the heap, even
+// when its extent is available before the construct.
func.func @acc_parallel_operand_outside(%n: index) {
%c0 = arith.constant 0 : index
%c0_i32 = arith.constant 0 : i32
@@ -80,13 +86,13 @@ func.func @acc_parallel_operand_outside(%n: index) {
return
}
// CHECK-LABEL: func.func @acc_parallel_operand_outside(
-// CHECK: %[[SIZE:.*]] = arith.addi
// CHECK-NOT: fir.alloca
// CHECK: acc.parallel {
-// CHECK-NEXT: fir.alloca !fir.array<?xi32>, %[[SIZE]]
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
-// The extent operand is defined inside the compute construct: the allocation is
-// placed right after it, as it would be in a plain function.
+// Same when the extent is defined inside the compute construct.
func.func @acc_parallel_operand_inside(%n: index) {
%c0 = arith.constant 0 : index
%c0_i32 = arith.constant 0 : i32
@@ -104,8 +110,10 @@ func.func @acc_parallel_operand_inside(%n: index) {
// CHECK-LABEL: func.func @acc_parallel_operand_inside(
// CHECK-NOT: fir.alloca
// CHECK: acc.parallel {
-// CHECK-NEXT: %[[SIZE:.*]] = arith.addi
-// CHECK-NEXT: fir.alloca !fir.array<?xi32>, %[[SIZE]]
+// CHECK-NEXT: arith.addi
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
// Hoisting out of a sequential loop is still done, but only up to the beginning
// of the compute construct.
@@ -212,3 +220,4 @@ func.func @acc_data_no_operand() {
// CHECK-LABEL: func.func @acc_data_no_operand()
// CHECK: fir.alloca !fir.array<42xi32>
// CHECK: acc.data
+}
diff --git a/flang/test/Transforms/stack-arrays-offload-region.fir b/flang/test/Transforms/stack-arrays-offload-region.fir
new file mode 100644
index 0000000000000..d14b82b7ab921
--- /dev/null
+++ b/flang/test/Transforms/stack-arrays-offload-region.fir
@@ -0,0 +1,38 @@
+// Test that stack-arrays leaves the heap allocations of an offload region
+// alone: the region runs on the device stack, which is far smaller than the
+// host one. The same allocation in host code is moved to the stack.
+
+// RUN: fir-opt --stack-arrays %s | FileCheck %s
+
+// CHECK-LABEL: func.func @dynamic_temp_in_host_code
+// CHECK: fir.alloca !fir.array<?xi32>, %{{.*}}
+// CHECK-NOT: fir.allocmem
+func.func @dynamic_temp_in_host_code(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ return
+}
+
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_parallel
+// CHECK: acc.parallel {
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_parallel(%n: index) {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ acc.parallel {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ acc.yield
+ }
+ return
+}
>From 9ab1e40bec0c8f6178c396cc486708952af67b2f Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Tue, 29 Sep 2026 18:39:04 -0700
Subject: [PATCH 2/6] add acc routine
---
.../Transforms/AllocationPlacement.cpp | 13 +++++++---
.../lib/Optimizer/Transforms/StackArrays.cpp | 6 ++++-
.../allocation-placement-offload-region.fir | 26 +++++++++++++++----
.../stack-arrays-offload-region.fir | 22 +++++++++++++---
4 files changed, 54 insertions(+), 13 deletions(-)
diff --git a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
index de6028b92dae5..9c61a074e9db0 100644
--- a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
+++ b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
@@ -32,6 +32,7 @@
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/Diagnostics.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
@@ -172,6 +173,10 @@ void AllocationPlacementPass::runOnOperation() {
return;
}
+ // An acc routine is compiled for the device as well, so it is device code
+ // for the purpose of -fstack-arrays even before it is specialized.
+ bool isAccRoutine = mlir::acc::isAccRoutine(func);
+
// Walk allocations in deterministic program order, maintaining the running
// per-function stack budget while collecting the conversions to perform.
std::size_t stackBytesUsed = 0;
@@ -208,11 +213,11 @@ void AllocationPlacementPass::runOnOperation() {
: (allocmem.hasLenParams() || allocmem.hasShapeOperands());
info.byteSize = getConstantByteSize(op, dl, kindMap);
- // -fstack-arrays cannot be honored in an offload region either: like a
- // device procedure, it runs on the device stack, which is far smaller than
- // the host one. The size based part of the policy still applies.
+ // -fstack-arrays cannot be honored in an offload region or an acc routine
+ // either: like a device procedure, they run on the device stack, which is
+ // far smaller than the host one. The size based policy still applies.
fir::AllocationPolicy policy = basePolicy;
- if (policy.stackArrays && cuf::isExecutingOnDevice(op))
+ if (policy.stackArrays && (isAccRoutine || cuf::isExecutingOnDevice(op)))
policy.stackArrays = false;
// A hook, if provided, fully overrides the default policy; it may delegate
diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp
index 575c853b0bc21..06cf3fff7bdcf 100644
--- a/flang/lib/Optimizer/Transforms/StackArrays.cpp
+++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp
@@ -25,6 +25,7 @@
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/Diagnostics.h"
@@ -778,11 +779,14 @@ void StackArraysPass::runOnOperation() {
// This pass only runs under -fstack-arrays, so honor a function that opted
// out in its own policy (device code, where the stack is tiny). Functions
- // without a policy of their own are left to the module setting.
+ // without a policy of their own are left to the module setting. An acc
+ // routine is compiled for the device as well and is treated the same way.
if (std::optional<fir::AllocationPolicy> policy =
fir::getLocalAllocationPolicy(func))
if (!policy->stackArrays)
return;
+ if (mlir::acc::isAccRoutine(func))
+ return;
auto &analysis = getAnalysis<fir::StackArraysAnalysisWrapper>();
const fir::StackArraysAnalysisWrapper::AllocMemMap *candidateOps =
diff --git a/flang/test/Transforms/allocation-placement-offload-region.fir b/flang/test/Transforms/allocation-placement-offload-region.fir
index c0fd11c15b12e..f53c90a242ed9 100644
--- a/flang/test/Transforms/allocation-placement-offload-region.fir
+++ b/flang/test/Transforms/allocation-placement-offload-region.fir
@@ -1,8 +1,8 @@
-// Test that -fstack-arrays is not honored inside an offload region: the region
-// runs on the device stack, which is far smaller than the host one, so its
-// runtime-sized temporaries stay on the heap while the same allocation in host
-// code goes on the stack. The size based part of the policy still applies, so
-// a small constant-size temporary goes on the device stack.
+// Test that -fstack-arrays is not honored inside an offload region or an acc
+// routine: they run on the device stack, which is far smaller than the host
+// one, so their runtime-sized temporaries stay on the heap while the same
+// allocation in host code goes on the stack. The size based part of the policy
+// still applies, so a small constant-size temporary goes on the device stack.
// RUN: fir-opt --allocation-placement %s | FileCheck %s
@@ -103,4 +103,20 @@ func.func @small_temp_in_acc_parallel() {
}
return
}
+
+// An acc routine is compiled for the device as well: its allocations stay on
+// the heap even before it is specialized.
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_routine
+// CHECK: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_routine(%n: index) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>} {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ return
+}
}
diff --git a/flang/test/Transforms/stack-arrays-offload-region.fir b/flang/test/Transforms/stack-arrays-offload-region.fir
index d14b82b7ab921..45ca5d76bf93a 100644
--- a/flang/test/Transforms/stack-arrays-offload-region.fir
+++ b/flang/test/Transforms/stack-arrays-offload-region.fir
@@ -1,6 +1,6 @@
-// Test that stack-arrays leaves the heap allocations of an offload region
-// alone: the region runs on the device stack, which is far smaller than the
-// host one. The same allocation in host code is moved to the stack.
+// Test that stack-arrays leaves the heap allocations of an offload region and
+// of an acc routine alone: they run on the device stack, which is far smaller
+// than the host one. The same allocation in host code is moved to the stack.
// RUN: fir-opt --stack-arrays %s | FileCheck %s
@@ -36,3 +36,19 @@ func.func @dynamic_temp_in_acc_parallel(%n: index) {
}
return
}
+
+// An acc routine is compiled for the device as well: its allocations stay on
+// the heap even before it is specialized.
+// CHECK-LABEL: func.func @dynamic_temp_in_acc_routine
+// CHECK: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_acc_routine(%n: index) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>} {
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ return
+}
>From 1b7108ff46324794de234bb2cb36d3dadbab5900 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Tue, 29 Sep 2026 19:47:28 -0700
Subject: [PATCH 3/6] Fold the acc routine check into the base allocation
policy
---
.../Optimizer/Transforms/AllocationPlacement.cpp | 16 ++++++++--------
1 file changed, 8 insertions(+), 8 deletions(-)
diff --git a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
index 9c61a074e9db0..10f7f615e8abb 100644
--- a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
+++ b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
@@ -154,6 +154,10 @@ void AllocationPlacementPass::runOnOperation() {
smallArrayThresholdBytes);
fir::overrideIfExplicitlySet(basePolicy.totalStackLimitBytes,
totalStackLimitBytes);
+ // An acc routine is compiled for the device as well, so -fstack-arrays
+ // cannot be honored in it, as in a device procedure.
+ if (mlir::acc::isAccRoutine(func))
+ basePolicy.stackArrays = false;
auto module = func->getParentOfType<mlir::ModuleOp>();
std::optional<mlir::DataLayout> dl =
@@ -173,10 +177,6 @@ void AllocationPlacementPass::runOnOperation() {
return;
}
- // An acc routine is compiled for the device as well, so it is device code
- // for the purpose of -fstack-arrays even before it is specialized.
- bool isAccRoutine = mlir::acc::isAccRoutine(func);
-
// Walk allocations in deterministic program order, maintaining the running
// per-function stack budget while collecting the conversions to perform.
std::size_t stackBytesUsed = 0;
@@ -213,11 +213,11 @@ void AllocationPlacementPass::runOnOperation() {
: (allocmem.hasLenParams() || allocmem.hasShapeOperands());
info.byteSize = getConstantByteSize(op, dl, kindMap);
- // -fstack-arrays cannot be honored in an offload region or an acc routine
- // either: like a device procedure, they run on the device stack, which is
- // far smaller than the host one. The size based policy still applies.
+ // -fstack-arrays cannot be honored in an offload region either: like a
+ // device procedure, it runs on the device stack, which is far smaller than
+ // the host one. The size based part of the policy still applies.
fir::AllocationPolicy policy = basePolicy;
- if (policy.stackArrays && (isAccRoutine || cuf::isExecutingOnDevice(op)))
+ if (policy.stackArrays && cuf::isExecutingOnDevice(op))
policy.stackArrays = false;
// A hook, if provided, fully overrides the default policy; it may delegate
>From e7448f38f9bdc0dfc2bb14036edd089bbd9b0093 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Tue, 29 Sep 2026 20:09:28 -0700
Subject: [PATCH 4/6] Add cuf.kernel cases to the offload-region stack-arrays
tests
---
.../allocation-placement-offload-region.fir | 21 +++++++++++++++++++
.../stack-arrays-offload-region.fir | 21 +++++++++++++++++++
2 files changed, 42 insertions(+)
diff --git a/flang/test/Transforms/allocation-placement-offload-region.fir b/flang/test/Transforms/allocation-placement-offload-region.fir
index f53c90a242ed9..a711857cf0792 100644
--- a/flang/test/Transforms/allocation-placement-offload-region.fir
+++ b/flang/test/Transforms/allocation-placement-offload-region.fir
@@ -84,6 +84,27 @@ func.func @dynamic_temp_in_acc_serial(%n: index) {
return
}
+// CHECK-LABEL: func.func @dynamic_temp_in_cuf_kernel
+// CHECK: cuf.kernel
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_cuf_kernel(%n: index) {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c1_i32 = arith.constant 1 : i32
+ %v = arith.constant 0 : i32
+ cuf.kernel<<<%c1_i32, %c1_i32>>> (%i : index) = (%c1 : index) to (%n : index) step (%c1 : index) {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ "fir.end"() : () -> ()
+ }
+ return
+}
+
// <10xi32> is 40 bytes, under the 1024 byte threshold: the size based policy
// still puts it on the (device) stack.
// CHECK-LABEL: func.func @small_temp_in_acc_parallel
diff --git a/flang/test/Transforms/stack-arrays-offload-region.fir b/flang/test/Transforms/stack-arrays-offload-region.fir
index 45ca5d76bf93a..3c7a4b2071cf1 100644
--- a/flang/test/Transforms/stack-arrays-offload-region.fir
+++ b/flang/test/Transforms/stack-arrays-offload-region.fir
@@ -37,6 +37,27 @@ func.func @dynamic_temp_in_acc_parallel(%n: index) {
return
}
+// CHECK-LABEL: func.func @dynamic_temp_in_cuf_kernel
+// CHECK: cuf.kernel
+// CHECK-NEXT: fir.allocmem !fir.array<?xi32>, %{{.*}}
+// CHECK: fir.freemem
+// CHECK-NOT: fir.alloca
+func.func @dynamic_temp_in_cuf_kernel(%n: index) {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c1_i32 = arith.constant 1 : i32
+ %v = arith.constant 0 : i32
+ cuf.kernel<<<%c1_i32, %c1_i32>>> (%i : index) = (%c1 : index) to (%n : index) step (%c1 : index) {
+ %0 = fir.allocmem !fir.array<?xi32>, %n
+ %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
+ "fir.end"() : () -> ()
+ }
+ return
+}
+
// An acc routine is compiled for the device as well: its allocations stay on
// the heap even before it is specialized.
// CHECK-LABEL: func.func @dynamic_temp_in_acc_routine
>From ef3d1b59cce3f36461642598adebf87568c1d627 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Wed, 30 Sep 2026 09:16:03 -0700
Subject: [PATCH 5/6] Record the device allocation policy on acc routines in
lowering
---
flang/lib/Lower/OpenACC.cpp | 4 ++
.../Transforms/AllocationPlacement.cpp | 5 ---
.../lib/Optimizer/Transforms/StackArrays.cpp | 6 +--
.../OpenACC/acc-routine-allocation-policy.f90 | 39 +++++++++++++++++++
.../allocation-placement-offload-region.fir | 25 +++---------
.../stack-arrays-offload-region.fir | 22 ++---------
6 files changed, 52 insertions(+), 49 deletions(-)
create mode 100644 flang/test/Lower/OpenACC/acc-routine-allocation-policy.f90
diff --git a/flang/lib/Lower/OpenACC.cpp b/flang/lib/Lower/OpenACC.cpp
index 0193265cf940b..4668e18479897 100644
--- a/flang/lib/Lower/OpenACC.cpp
+++ b/flang/lib/Lower/OpenACC.cpp
@@ -25,6 +25,7 @@
#include "flang/Lower/Support/Utils.h"
#include "flang/Lower/SymbolMap.h"
#include "flang/Optimizer/Builder/BoxValue.h"
+#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/HLFIRTools.h"
#include "flang/Optimizer/Builder/IntrinsicCall.h"
@@ -4795,6 +4796,9 @@ static void attachRoutineInfo(mlir::func::FuncOp func,
func.getOperation()->setAttr(
mlir::acc::getRoutineInfoAttrName(),
mlir::acc::RoutineInfoAttr::get(func.getContext(), routines));
+ // The routine is compiled for the device as well, where -fstack-arrays
+ // cannot be honored: record the same policy as for a device procedure.
+ cuf::setDeviceAllocationPolicy(func.getOperation());
}
static mlir::ArrayAttr
diff --git a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
index 10f7f615e8abb..de6028b92dae5 100644
--- a/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
+++ b/flang/lib/Optimizer/Transforms/AllocationPlacement.cpp
@@ -32,7 +32,6 @@
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/Diagnostics.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
@@ -154,10 +153,6 @@ void AllocationPlacementPass::runOnOperation() {
smallArrayThresholdBytes);
fir::overrideIfExplicitlySet(basePolicy.totalStackLimitBytes,
totalStackLimitBytes);
- // An acc routine is compiled for the device as well, so -fstack-arrays
- // cannot be honored in it, as in a device procedure.
- if (mlir::acc::isAccRoutine(func))
- basePolicy.stackArrays = false;
auto module = func->getParentOfType<mlir::ModuleOp>();
std::optional<mlir::DataLayout> dl =
diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp
index 06cf3fff7bdcf..575c853b0bc21 100644
--- a/flang/lib/Optimizer/Transforms/StackArrays.cpp
+++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp
@@ -25,7 +25,6 @@
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/Diagnostics.h"
@@ -779,14 +778,11 @@ void StackArraysPass::runOnOperation() {
// This pass only runs under -fstack-arrays, so honor a function that opted
// out in its own policy (device code, where the stack is tiny). Functions
- // without a policy of their own are left to the module setting. An acc
- // routine is compiled for the device as well and is treated the same way.
+ // without a policy of their own are left to the module setting.
if (std::optional<fir::AllocationPolicy> policy =
fir::getLocalAllocationPolicy(func))
if (!policy->stackArrays)
return;
- if (mlir::acc::isAccRoutine(func))
- return;
auto &analysis = getAnalysis<fir::StackArraysAnalysisWrapper>();
const fir::StackArraysAnalysisWrapper::AllocMemMap *candidateOps =
diff --git a/flang/test/Lower/OpenACC/acc-routine-allocation-policy.f90 b/flang/test/Lower/OpenACC/acc-routine-allocation-policy.f90
new file mode 100644
index 0000000000000..aa2510ddda2c1
--- /dev/null
+++ b/flang/test/Lower/OpenACC/acc-routine-allocation-policy.f90
@@ -0,0 +1,39 @@
+! Test that -fstack-arrays is not applied to acc routines. They are compiled for
+! the device as well, where the stack is far smaller than on the host, so
+! lowering records the same policy as for a device procedure. The policy is
+! recorded whether or not -fstack-arrays was requested, so that the opt-out is
+! explicit in the IR.
+
+! RUN: %flang_fc1 -emit-fir -fopenacc %s -o - | FileCheck %s --check-prefixes=CHECK,DEFAULT
+! RUN: %flang_fc1 -emit-fir -fopenacc -fstack-arrays %s -o - | FileCheck %s --check-prefixes=CHECK,STACK
+
+subroutine routine_seq(a, n)
+ !$acc routine seq
+ integer :: a(*)
+ integer :: n
+ integer :: auto(n)
+ do i = 1, n
+ auto(i) = i
+ end do
+ a(1) = sum(auto(1:n))
+end subroutine
+
+subroutine host_sub(a, n)
+ integer :: a(*)
+ integer :: n
+ integer :: auto(n)
+ do i = 1, n
+ auto(i) = i
+ end do
+ a(1) = sum(auto(1:n))
+end subroutine
+
+! The module policy follows -fstack-arrays.
+! DEFAULT: module attributes {{.*}}fir.allocation_policy = #fir.allocation_policy<stack_arrays = false
+! STACK: module attributes {{.*}}fir.allocation_policy = #fir.allocation_policy<stack_arrays = true
+
+! The acc routine opts out of it in both cases.
+! CHECK: func.func @_QProutine_seq({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false
+
+! The host procedure keeps the module policy.
+! CHECK: func.func @_QPhost_sub({{.*}}) {
diff --git a/flang/test/Transforms/allocation-placement-offload-region.fir b/flang/test/Transforms/allocation-placement-offload-region.fir
index a711857cf0792..50294b4d0fd49 100644
--- a/flang/test/Transforms/allocation-placement-offload-region.fir
+++ b/flang/test/Transforms/allocation-placement-offload-region.fir
@@ -1,8 +1,8 @@
-// Test that -fstack-arrays is not honored inside an offload region or an acc
-// routine: they run on the device stack, which is far smaller than the host
-// one, so their runtime-sized temporaries stay on the heap while the same
-// allocation in host code goes on the stack. The size based part of the policy
-// still applies, so a small constant-size temporary goes on the device stack.
+// Test that -fstack-arrays is not honored inside an offload region: the region
+// runs on the device stack, which is far smaller than the host one, so its
+// runtime-sized temporaries stay on the heap while the same allocation in host
+// code goes on the stack. The size based part of the policy still applies, so
+// a small constant-size temporary goes on the device stack.
// RUN: fir-opt --allocation-placement %s | FileCheck %s
@@ -125,19 +125,4 @@ func.func @small_temp_in_acc_parallel() {
return
}
-// An acc routine is compiled for the device as well: its allocations stay on
-// the heap even before it is specialized.
-// CHECK-LABEL: func.func @dynamic_temp_in_acc_routine
-// CHECK: fir.allocmem !fir.array<?xi32>, %{{.*}}
-// CHECK-NOT: fir.alloca
-func.func @dynamic_temp_in_acc_routine(%n: index) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>} {
- %c0 = arith.constant 0 : index
- %v = arith.constant 0 : i32
- %0 = fir.allocmem !fir.array<?xi32>, %n
- %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
- %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
- fir.store %v to %e : !fir.ref<i32>
- fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
- return
-}
}
diff --git a/flang/test/Transforms/stack-arrays-offload-region.fir b/flang/test/Transforms/stack-arrays-offload-region.fir
index 3c7a4b2071cf1..1a92522e0f3a0 100644
--- a/flang/test/Transforms/stack-arrays-offload-region.fir
+++ b/flang/test/Transforms/stack-arrays-offload-region.fir
@@ -1,6 +1,6 @@
-// Test that stack-arrays leaves the heap allocations of an offload region and
-// of an acc routine alone: they run on the device stack, which is far smaller
-// than the host one. The same allocation in host code is moved to the stack.
+// Test that stack-arrays leaves the heap allocations of an offload region
+// alone: the region runs on the device stack, which is far smaller than the
+// host one. The same allocation in host code is moved to the stack.
// RUN: fir-opt --stack-arrays %s | FileCheck %s
@@ -57,19 +57,3 @@ func.func @dynamic_temp_in_cuf_kernel(%n: index) {
}
return
}
-
-// An acc routine is compiled for the device as well: its allocations stay on
-// the heap even before it is specialized.
-// CHECK-LABEL: func.func @dynamic_temp_in_acc_routine
-// CHECK: fir.allocmem !fir.array<?xi32>, %{{.*}}
-// CHECK-NOT: fir.alloca
-func.func @dynamic_temp_in_acc_routine(%n: index) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>} {
- %c0 = arith.constant 0 : index
- %v = arith.constant 0 : i32
- %0 = fir.allocmem !fir.array<?xi32>, %n
- %r = fir.convert %0 : (!fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
- %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
- fir.store %v to %e : !fir.ref<i32>
- fir.freemem %0 : !fir.heap<!fir.array<?xi32>>
- return
-}
>From b24ce511f40ba2a75b74fc29d9d3f289d66282e5 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Wed, 30 Sep 2026 09:28:38 -0700
Subject: [PATCH 6/6] Loosen acc routine attribute checks for the allocation
policy
---
.../acc-routine-bind-clone-signature.f90 | 2 +-
.../acc-routine-bind-devtype-filter.f90 | 10 ++++-----
.../acc-routine-bind-devtype-undeclared.f90 | 2 +-
.../acc-routine-bind-string-undeclared.f90 | 2 +-
.../OpenACC/acc-routine-bind-undeclared.f90 | 2 +-
.../Lower/OpenACC/acc-routine-multi-name.f90 | 10 ++++-----
.../test/Lower/OpenACC/acc-routine-named.f90 | 4 ++--
.../Lower/OpenACC/acc-routine-use-module.f90 | 2 +-
flang/test/Lower/OpenACC/acc-routine.f90 | 22 +++++++++----------
flang/test/Lower/OpenACC/acc-routine02.f90 | 2 +-
flang/test/Lower/OpenACC/acc-routine03.f90 | 4 ++--
flang/test/Lower/OpenACC/acc-routine04.f90 | 4 ++--
12 files changed, 33 insertions(+), 33 deletions(-)
diff --git a/flang/test/Lower/OpenACC/acc-routine-bind-clone-signature.f90 b/flang/test/Lower/OpenACC/acc-routine-bind-clone-signature.f90
index 181e1883a68d1..63b9cb3cf17d5 100644
--- a/flang/test/Lower/OpenACC/acc-routine-bind-clone-signature.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-bind-clone-signature.f90
@@ -24,4 +24,4 @@ subroutine aclear(y)
! The decorated routine's signature (assumed-shape array descriptor):
! CHECK: func.func private @_QPaclear(!fir.box<!fir.array<?xf32>>
! The bind target is declared with the same type, proving the clone:
-! CHECK: func.func private @_QPaclear_seq(!fir.box<!fir.array<?xf32>>) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_SEQ_ROUTINE]]]>}
+! CHECK: func.func private @_QPaclear_seq(!fir.box<!fir.array<?xf32>>) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_SEQ_ROUTINE]]]>{{.*}}}
diff --git a/flang/test/Lower/OpenACC/acc-routine-bind-devtype-filter.f90 b/flang/test/Lower/OpenACC/acc-routine-bind-devtype-filter.f90
index 725dec744b01b..c17668364b4ff 100644
--- a/flang/test/Lower/OpenACC/acc-routine-bind-devtype-filter.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-bind-devtype-filter.f90
@@ -17,8 +17,8 @@ subroutine s_bind_devtype_filter(n, x)
! CHECK-DAG: acc.routine @{{.*}} func(@_QPfoo) bind(@_QPfoo_n [#acc.device_type<nvidia>], @_QPfoo_m [#acc.device_type<multicore>]) worker ([#acc.device_type<multicore>]) vector ([#acc.device_type<nvidia>])
! CHECK-DAG: acc.routine @[[FOO_N_ROUTINE:.*]] func(@_QPfoo_n) vector ([#acc.device_type<nvidia>]){{$}}
! CHECK-DAG: acc.routine @[[FOO_M_ROUTINE:.*]] func(@_QPfoo_m) worker ([#acc.device_type<multicore>]){{$}}
-! CHECK-DAG: func.func private @_QPfoo_n({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_N_ROUTINE]]]>}
-! CHECK-DAG: func.func private @_QPfoo_m({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_M_ROUTINE]]]>}
+! CHECK-DAG: func.func private @_QPfoo_n({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_N_ROUTINE]]]>{{.*}}}
+! CHECK-DAG: func.func private @_QPfoo_m({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_M_ROUTINE]]]>{{.*}}}
subroutine s_bind_devtype_merged_target(n, x)
integer :: n, i
@@ -33,7 +33,7 @@ subroutine s_bind_devtype_merged_target(n, x)
! CHECK-DAG: acc.routine @{{.*}} func(@_QPfoo_merge) bind(@_QPfoo_dev [#acc.device_type<nvidia>], @_QPfoo_dev [#acc.device_type<multicore>]) worker ([#acc.device_type<multicore>]) vector ([#acc.device_type<nvidia>])
! CHECK-DAG: acc.routine @[[FOO_DEV_ROUTINE:.*]] func(@_QPfoo_dev) worker ([#acc.device_type<multicore>]) vector ([#acc.device_type<nvidia>]){{$}}
-! CHECK-DAG: func.func private @_QPfoo_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_DEV_ROUTINE]]]>}
+! CHECK-DAG: func.func private @_QPfoo_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[FOO_DEV_ROUTINE]]]>{{.*}}}
subroutine s_bind_before_modality(n, x)
integer :: n, i
@@ -49,5 +49,5 @@ subroutine s_bind_before_modality(n, x)
! CHECK-DAG: acc.routine @{{.*}} func(@_QPbar) bind(@_QPbar_n [#acc.device_type<nvidia>], @_QPbar_m [#acc.device_type<multicore>]) vector ([#acc.device_type<nvidia>]) seq ([#acc.device_type<multicore>])
! CHECK-DAG: acc.routine @[[BAR_N_ROUTINE:.*]] func(@_QPbar_n) vector ([#acc.device_type<nvidia>]){{$}}
! CHECK-DAG: acc.routine @[[BAR_M_ROUTINE:.*]] func(@_QPbar_m) seq ([#acc.device_type<multicore>]){{$}}
-! CHECK-DAG: func.func private @_QPbar_n({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[BAR_N_ROUTINE]]]>}
-! CHECK-DAG: func.func private @_QPbar_m({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[BAR_M_ROUTINE]]]>}
+! CHECK-DAG: func.func private @_QPbar_n({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[BAR_N_ROUTINE]]]>{{.*}}}
+! CHECK-DAG: func.func private @_QPbar_m({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[BAR_M_ROUTINE]]]>{{.*}}}
diff --git a/flang/test/Lower/OpenACC/acc-routine-bind-devtype-undeclared.f90 b/flang/test/Lower/OpenACC/acc-routine-bind-devtype-undeclared.f90
index 74e13ff3a5bff..a2add365202ca 100644
--- a/flang/test/Lower/OpenACC/acc-routine-bind-devtype-undeclared.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-bind-devtype-undeclared.f90
@@ -15,4 +15,4 @@ subroutine s_bind_devtype(n, x)
! CHECK: acc.routine @[[ACLEAR_DEV_ROUTINE:.*]] func(@_QPaclear_dev) seq
! CHECK: acc.routine @{{.*}} func(@_QPaclear){{.*}}@_QPaclear_dev
-! CHECK: func.func private @_QPaclear_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_DEV_ROUTINE]]]>}
+! CHECK: func.func private @_QPaclear_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_DEV_ROUTINE]]]>{{.*}}}
diff --git a/flang/test/Lower/OpenACC/acc-routine-bind-string-undeclared.f90 b/flang/test/Lower/OpenACC/acc-routine-bind-string-undeclared.f90
index 6647c6986b8b1..eed98594bb3f4 100644
--- a/flang/test/Lower/OpenACC/acc-routine-bind-string-undeclared.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-bind-string-undeclared.f90
@@ -5,7 +5,7 @@
! CHECK: acc.routine @[[ACLEAR_DEV_ROUTINE:.*]] func(@aclear_dev) seq
! CHECK: acc.routine @{{.*}} func(@_QPaclear) bind("aclear_dev") seq
-! CHECK: func.func private @aclear_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_DEV_ROUTINE]]]>}
+! CHECK: func.func private @aclear_dev({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_DEV_ROUTINE]]]>{{.*}}}
! CHECK-SAME: loc("{{.*}}acc-routine-bind-string-undeclared.f90":{{[0-9]+}}:{{[0-9]+}})
! CHECK-NOT: func.func private @aclear_dev{{.*}}loc(unknown)
diff --git a/flang/test/Lower/OpenACC/acc-routine-bind-undeclared.f90 b/flang/test/Lower/OpenACC/acc-routine-bind-undeclared.f90
index c80e062dc8782..d559aa520e366 100644
--- a/flang/test/Lower/OpenACC/acc-routine-bind-undeclared.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-bind-undeclared.f90
@@ -4,7 +4,7 @@
! CHECK: acc.routine @[[ACLEAR_SEQ_ROUTINE:.*]] func(@_QPaclear_seq) seq
! CHECK: acc.routine @{{.*}} func(@_QPaclear) bind(@_QPaclear_seq) seq
-! CHECK: func.func private @_QPaclear_seq({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_SEQ_ROUTINE]]]>}
+! CHECK: func.func private @_QPaclear_seq({{.*}}) attributes {acc.routine_info = #acc.routine_info<[@[[ACLEAR_SEQ_ROUTINE]]]>{{.*}}}
! CHECK-SAME: loc("{{.*}}acc-routine-bind-undeclared.f90":{{[0-9]+}}:{{[0-9]+}})
! CHECK-NOT: func.func private @_QPaclear_seq{{.*}}loc(unknown)
diff --git a/flang/test/Lower/OpenACC/acc-routine-multi-name.f90 b/flang/test/Lower/OpenACC/acc-routine-multi-name.f90
index 6ba46f2da7fbb..d293eaa5b48b7 100644
--- a/flang/test/Lower/OpenACC/acc-routine-multi-name.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-multi-name.f90
@@ -23,30 +23,30 @@ subroutine seq1()
end subroutine
! CHECK-LABEL: func.func @_QMacc_multi_routinesPseq1()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_seq1]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_seq1]]]>{{.*}}}
subroutine seq2()
end subroutine
! CHECK-LABEL: func.func @_QMacc_multi_routinesPseq2()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_seq2]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_seq2]]]>{{.*}}}
subroutine gang1()
end subroutine
! CHECK-LABEL: func.func @_QMacc_multi_routinesPgang1()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang1]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang1]]]>{{.*}}}
subroutine gang2()
end subroutine
! CHECK-LABEL: func.func @_QMacc_multi_routinesPgang2()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang2]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang2]]]>{{.*}}}
subroutine gang3()
end subroutine
! CHECK-LABEL: func.func @_QMacc_multi_routinesPgang3()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang3]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r_gang3]]]>{{.*}}}
end module
diff --git a/flang/test/Lower/OpenACC/acc-routine-named.f90 b/flang/test/Lower/OpenACC/acc-routine-named.f90
index 24d47e58b6e1b..3a8a2a21a195a 100644
--- a/flang/test/Lower/OpenACC/acc-routine-named.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-named.f90
@@ -15,13 +15,13 @@ subroutine acc1()
end subroutine
! CHECK-LABEL: func.func @_QMacc_routinesPacc1()
-! CHECK-SAME:attributes {acc.routine_info = #acc.routine_info<[@[[r1]]]>}
+! CHECK-SAME:attributes {acc.routine_info = #acc.routine_info<[@[[r1]]]>{{.*}}}
subroutine acc2()
!$acc routine(acc2)
end subroutine
! CHECK-LABEL: func.func @_QMacc_routinesPacc2()
-! CHECK-SAME:attributes {acc.routine_info = #acc.routine_info<[@[[r0]]]>}
+! CHECK-SAME:attributes {acc.routine_info = #acc.routine_info<[@[[r0]]]>{{.*}}}
end module
diff --git a/flang/test/Lower/OpenACC/acc-routine-use-module.f90 b/flang/test/Lower/OpenACC/acc-routine-use-module.f90
index 059324230a746..79316d405f77b 100644
--- a/flang/test/Lower/OpenACC/acc-routine-use-module.f90
+++ b/flang/test/Lower/OpenACC/acc-routine-use-module.f90
@@ -19,5 +19,5 @@ subroutine caller(aa)
!$acc end serial
end subroutine
!CHECK: }
- !CHECK: func.func private @_QMmod1Pcallee(!fir.ref<i32>) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>}
+ !CHECK: func.func private @_QMmod1Pcallee(!fir.ref<i32>) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>{{.*}}}
end module
\ No newline at end of file
diff --git a/flang/test/Lower/OpenACC/acc-routine.f90 b/flang/test/Lower/OpenACC/acc-routine.f90
index c281ca5dfc287..37c2cc17240f5 100644
--- a/flang/test/Lower/OpenACC/acc-routine.f90
+++ b/flang/test/Lower/OpenACC/acc-routine.f90
@@ -24,56 +24,56 @@ subroutine acc_routine1()
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine1()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r00]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r00]]]>{{.*}}}
subroutine acc_routine2()
!$acc routine seq
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine2()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r01]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r01]]]>{{.*}}}
subroutine acc_routine3()
!$acc routine gang
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine3()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r02]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r02]]]>{{.*}}}
subroutine acc_routine4()
!$acc routine vector
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine4()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r03]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r03]]]>{{.*}}}
subroutine acc_routine5()
!$acc routine worker
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine5()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r04]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r04]]]>{{.*}}}
subroutine acc_routine6()
!$acc routine nohost
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine6()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r05]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r05]]]>{{.*}}}
subroutine acc_routine7()
!$acc routine gang(dim:1)
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine7()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r06]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r06]]]>{{.*}}}
subroutine acc_routine8()
!$acc routine bind("routine8_")
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine8()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r07]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r07]]]>{{.*}}}
subroutine acc_routine9a()
end subroutine
@@ -83,14 +83,14 @@ subroutine acc_routine9()
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine9()
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r08]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r08]]]>{{.*}}}
function acc_routine10()
!$acc routine(acc_routine10) seq
end function
! CHECK-LABEL: func.func @_QPacc_routine10() -> f32
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r09]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r09]]]>{{.*}}}
subroutine acc_routine11(a)
real :: a
@@ -98,7 +98,7 @@ subroutine acc_routine11(a)
end subroutine
! CHECK-LABEL: func.func @_QPacc_routine11(%arg0: !fir.ref<f32> {fir.bindc_name = "a"})
-! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r10]]]>}
+! CHECK-SAME: attributes {acc.routine_info = #acc.routine_info<[@[[r10]]]>{{.*}}}
subroutine acc_routine12()
diff --git a/flang/test/Lower/OpenACC/acc-routine02.f90 b/flang/test/Lower/OpenACC/acc-routine02.f90
index dd07cba4b20e3..1c350c1ee33c8 100644
--- a/flang/test/Lower/OpenACC/acc-routine02.f90
+++ b/flang/test/Lower/OpenACC/acc-routine02.f90
@@ -17,4 +17,4 @@ program test
! CHECK-LABEL: acc.routine @acc_routine_0 func(@_QPsub1)
-! CHECK: func.func @_QPsub1(%ar{{.*}}: !fir.ref<!fir.array<?xf32>> {fir.bindc_name = "a"}, %arg1: !fir.ref<i32> {fir.bindc_name = "n"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>}
+! CHECK: func.func @_QPsub1(%ar{{.*}}: !fir.ref<!fir.array<?xf32>> {fir.bindc_name = "a"}, %arg1: !fir.ref<i32> {fir.bindc_name = "n"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>{{.*}}}
diff --git a/flang/test/Lower/OpenACC/acc-routine03.f90 b/flang/test/Lower/OpenACC/acc-routine03.f90
index 3fc307746849f..0c2754ab3fcbf 100644
--- a/flang/test/Lower/OpenACC/acc-routine03.f90
+++ b/flang/test/Lower/OpenACC/acc-routine03.f90
@@ -31,5 +31,5 @@ subroutine sub2(a)
! CHECK: acc.routine @acc_routine_1 func(@_QPsub2) worker nohost
! CHECK: acc.routine @acc_routine_0 func(@_QPsub1) bind(@_QPsub2) worker
-! CHECK: func.func @_QPsub1(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "a"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>}
-! CHECK: func.func @_QPsub2(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "a"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_1]>}
+! CHECK: func.func @_QPsub1(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "a"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>{{.*}}}
+! CHECK: func.func @_QPsub2(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "a"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_1]>{{.*}}}
diff --git a/flang/test/Lower/OpenACC/acc-routine04.f90 b/flang/test/Lower/OpenACC/acc-routine04.f90
index 470440728d2f5..fb8009ebe160c 100644
--- a/flang/test/Lower/OpenACC/acc-routine04.f90
+++ b/flang/test/Lower/OpenACC/acc-routine04.f90
@@ -29,6 +29,6 @@ subroutine sub2()
! CHECK: acc.routine @acc_routine_1 func(@_QFPsub2) seq
! CHECK: acc.routine @acc_routine_0 func(@_QMdummy_modPsub1) seq
-! CHECK: func.func @_QMdummy_modPsub1(%arg0: !fir.ref<i32> {fir.bindc_name = "i"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>}
+! CHECK: func.func @_QMdummy_modPsub1(%arg0: !fir.ref<i32> {fir.bindc_name = "i"}) attributes {acc.routine_info = #acc.routine_info<[@acc_routine_0]>{{.*}}}
! CHECK: func.func @_QQmain() attributes {fir.bindc_name = "TEST_ACC_ROUTINE"}
-! CHECK: func.func private @_QFPsub2() attributes {acc.routine_info = #acc.routine_info<[@acc_routine_1]>, fir.host_symbol = @_QQmain, llvm.linkage = #llvm.linkage<internal>}
+! CHECK: func.func private @_QFPsub2() attributes {acc.routine_info = #acc.routine_info<[@acc_routine_1]>, fir.allocation_policy = {{.*}}, fir.host_symbol = @_QQmain, llvm.linkage = #llvm.linkage<internal>}
More information about the flang-commits
mailing list