[flang-commits] [flang] [flang][cuda] do not honor -fstack-arrays in device procedures (PR #223040)
Zhen Wang via flang-commits
flang-commits at lists.llvm.org
Fri Sep 11 17:55:27 PDT 2026
https://github.com/wangzpgi updated https://github.com/llvm/llvm-project/pull/223040
>From 11559969f9a7218860c5850cf22bf12966b99cd0 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Fri, 11 Sep 2026 12:49:34 -0700
Subject: [PATCH 1/2] [flang][cuda] do not honor -fstack-arrays in device
procedures
---
.../Optimizer/Support/AllocationPolicy.h | 16 ++++--
flang/lib/Lower/CallInterface.cpp | 21 ++++++--
.../Optimizer/Support/AllocationPolicy.cpp | 37 ++++++++++----
.../lib/Optimizer/Transforms/StackArrays.cpp | 9 ++++
.../Lower/CUDA/cuda-allocation-policy.cuf | 50 +++++++++++++++++++
.../allocation-policy-precedence.fir | 47 +++++++++++++++++
.../stack-arrays-policy-opt-out.fir | 35 +++++++++++++
7 files changed, 198 insertions(+), 17 deletions(-)
create mode 100644 flang/test/Lower/CUDA/cuda-allocation-policy.cuf
create mode 100644 flang/test/Transforms/stack-arrays-policy-opt-out.fir
diff --git a/flang/include/flang/Optimizer/Support/AllocationPolicy.h b/flang/include/flang/Optimizer/Support/AllocationPolicy.h
index cb4e57628744f..f43dc9d5b55c2 100644
--- a/flang/include/flang/Optimizer/Support/AllocationPolicy.h
+++ b/flang/include/flang/Optimizer/Support/AllocationPolicy.h
@@ -133,12 +133,22 @@ AllocationPolicy getCommandLineAllocationPolicy(bool stackArrays);
/// any policy already recorded there.
void setAllocationPolicy(mlir::ModuleOp mod, const AllocationPolicy &policy);
+/// Record \p policy on \p op, which is meant to be a function that needs a
+/// policy of its own, narrower than the module one.
+void setAllocationPolicy(mlir::Operation *op, const AllocationPolicy &policy);
+
/// Get the policy recorded on \p mod, or the defaults if none was recorded.
AllocationPolicy getAllocationPolicy(mlir::ModuleOp mod);
-/// Get the policy in effect for \p op, which is the one recorded on its
-/// enclosing ModuleOp. Returns the defaults if \p op is not inside a module or
-/// if no policy was recorded.
+/// Get the policy recorded directly on \p op, without looking at its parents.
+/// Use this to tell "this operation opted out" from "nothing was recorded
+/// anywhere", which the defaults cannot express.
+std::optional<AllocationPolicy> getLocalAllocationPolicy(mlir::Operation *op);
+
+/// Get the policy in effect for \p op: the one recorded on the innermost
+/// enclosing operation that carries one, usually the enclosing function if it
+/// has its own policy, otherwise the ModuleOp. Returns the defaults if no
+/// policy was recorded.
AllocationPolicy getAllocationPolicy(mlir::Operation *op);
} // namespace fir
diff --git a/flang/lib/Lower/CallInterface.cpp b/flang/lib/Lower/CallInterface.cpp
index 64960b5e09a98..0709986f692a0 100644
--- a/flang/lib/Lower/CallInterface.cpp
+++ b/flang/lib/Lower/CallInterface.cpp
@@ -20,6 +20,7 @@
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/Dialect/FIROpsSupport.h"
+#include "flang/Optimizer/Support/AllocationPolicy.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Support/Utils.h"
#include "flang/Semantics/symbol.h"
@@ -695,10 +696,22 @@ setCUDAAttributes(mlir::func::FuncOp func,
std::optional<Fortran::evaluate::characteristics::Procedure>
characteristic) {
if (characteristic && characteristic->cudaSubprogramAttrs) {
- func.getOperation()->setAttr(
- cuf::getProcAttrName(),
- cuf::getProcAttribute(func.getContext(),
- *characteristic->cudaSubprogramAttrs));
+ auto procAttr = cuf::getProcAttribute(func.getContext(),
+ *characteristic->cudaSubprogramAttrs);
+ func.getOperation()->setAttr(cuf::getProcAttrName(), procAttr);
+ // -fstack-arrays cannot be honored in device code: the device stack is
+ // orders of magnitude smaller, and an automatic array that fits the host
+ // stack easily overflows it. host_device is the host copy of the routine.
+ cuf::ProcAttribute proc = procAttr.getValue();
+ if (proc != cuf::ProcAttribute::Host &&
+ proc != cuf::ProcAttribute::HostDevice) {
+ fir::AllocationPolicy policy =
+ fir::getAllocationPolicy(func.getOperation());
+ if (policy.stackArrays) {
+ policy.stackArrays = false;
+ fir::setAllocationPolicy(func.getOperation(), policy);
+ }
+ }
}
if (sym) {
diff --git a/flang/lib/Optimizer/Support/AllocationPolicy.cpp b/flang/lib/Optimizer/Support/AllocationPolicy.cpp
index ad8eca26f8473..fa0036db9758b 100644
--- a/flang/lib/Optimizer/Support/AllocationPolicy.cpp
+++ b/flang/lib/Optimizer/Support/AllocationPolicy.cpp
@@ -91,12 +91,17 @@ fir::AllocationPolicy fir::getCommandLineAllocationPolicy(bool stackArrays) {
return policy;
}
+void fir::setAllocationPolicy(mlir::Operation *op,
+ const fir::AllocationPolicy &policy) {
+ op->setAttr(allocationPolicyName, fir::AllocationPolicyAttr::get(
+ op->getContext(), policy.stackArrays,
+ policy.smallArrayThresholdBytes,
+ policy.totalStackLimitBytes));
+}
+
void fir::setAllocationPolicy(mlir::ModuleOp mod,
const fir::AllocationPolicy &policy) {
- mod->setAttr(allocationPolicyName, fir::AllocationPolicyAttr::get(
- mod.getContext(), policy.stackArrays,
- policy.smallArrayThresholdBytes,
- policy.totalStackLimitBytes));
+ setAllocationPolicy(mod.getOperation(), policy);
}
fir::AllocationPolicy fir::getAllocationPolicy(mlir::ModuleOp mod) {
@@ -109,11 +114,23 @@ fir::AllocationPolicy fir::getAllocationPolicy(mlir::ModuleOp mod) {
attr.getTotalStackLimit()};
}
+std::optional<fir::AllocationPolicy>
+fir::getLocalAllocationPolicy(mlir::Operation *op) {
+ auto attr =
+ op->getAttrOfType<fir::AllocationPolicyAttr>(allocationPolicyName);
+ if (!attr)
+ return std::nullopt;
+ return fir::AllocationPolicy{attr.getStackArrays(),
+ attr.getSmallArrayThreshold(),
+ attr.getTotalStackLimit()};
+}
+
fir::AllocationPolicy fir::getAllocationPolicy(mlir::Operation *op) {
- auto mod = mlir::dyn_cast<mlir::ModuleOp>(op);
- if (!mod)
- mod = op->getParentOfType<mlir::ModuleOp>();
- if (!mod)
- return fir::AllocationPolicy{};
- return getAllocationPolicy(mod);
+ // The innermost policy wins, so that a function can narrow the module one
+ // (e.g. device code, where the stack is a scarce resource).
+ for (mlir::Operation *cur = op; cur; cur = cur->getParentOp())
+ if (std::optional<fir::AllocationPolicy> policy =
+ fir::getLocalAllocationPolicy(cur))
+ return *policy;
+ return fir::AllocationPolicy{};
}
diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp
index 8c2b39c40762c..7d4e6d49642d3 100644
--- a/flang/lib/Optimizer/Transforms/StackArrays.cpp
+++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp
@@ -14,6 +14,7 @@
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
#include "flang/Optimizer/Dialect/Support/FIRContext.h"
+#include "flang/Optimizer/Support/AllocationPolicy.h"
#include "flang/Optimizer/Support/DataLayout.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h"
@@ -774,6 +775,14 @@ llvm::StringRef StackArraysPass::getDescription() const {
void StackArraysPass::runOnOperation() {
mlir::func::FuncOp func = getOperation();
+ // 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.
+ if (std::optional<fir::AllocationPolicy> policy =
+ fir::getLocalAllocationPolicy(func))
+ if (!policy->stackArrays)
+ return;
+
auto &analysis = getAnalysis<fir::StackArraysAnalysisWrapper>();
const fir::StackArraysAnalysisWrapper::AllocMemMap *candidateOps =
analysis.getCandidateOps(func);
diff --git a/flang/test/Lower/CUDA/cuda-allocation-policy.cuf b/flang/test/Lower/CUDA/cuda-allocation-policy.cuf
new file mode 100644
index 0000000000000..8749d6b78e7da
--- /dev/null
+++ b/flang/test/Lower/CUDA/cuda-allocation-policy.cuf
@@ -0,0 +1,50 @@
+! Test that -fstack-arrays is not applied to device code. The device stack is
+! far smaller than the host one, so lowering records a policy of its own on
+! device and global procedures to keep their automatic arrays on the heap.
+
+! RUN: %flang_fc1 -emit-fir %s -o - | FileCheck %s --check-prefix=DEFAULT
+! RUN: %flang_fc1 -emit-fir -fstack-arrays %s -o - | FileCheck %s
+
+module m
+contains
+ attributes(global) subroutine kernel(a, n)
+ integer, device :: a(*)
+ integer, value :: n
+ integer :: auto(n)
+ do i = 1, n
+ auto(i) = i
+ end do
+ a(1) = sum(auto(1:n))
+ end subroutine
+
+ attributes(device) subroutine devsub(a, n)
+ integer, device :: a(*)
+ integer, value :: n
+ integer :: auto(n)
+ do i = 1, n
+ auto(i) = i
+ end do
+ a(1) = sum(auto(1:n))
+ end subroutine
+
+ attributes(host) subroutine hostsub(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
+end module
+
+! Without -fstack-arrays the module policy already keeps automatic arrays on the
+! heap, so no procedure needs a policy of its own.
+! DEFAULT-NOT: func.func{{.*}}fir.allocation_policy
+
+! CHECK: func.func @_QMmPkernel({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false,
+! CHECK: func.func @_QMmPdevsub({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false,
+
+! The host procedure keeps the module policy, so -fstack-arrays still applies
+! to it.
+! CHECK: func.func @_QMmPhostsub({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<host>}
diff --git a/flang/test/Transforms/allocation-policy-precedence.fir b/flang/test/Transforms/allocation-policy-precedence.fir
index 7f3a9803b8758..10e99cd6ac7de 100644
--- a/flang/test/Transforms/allocation-policy-precedence.fir
+++ b/flang/test/Transforms/allocation-policy-precedence.fir
@@ -60,3 +60,50 @@ func.func @without_policy_attribute() {
return
}
}
+
+// -----
+
+// A policy on the function narrows the module one: the function policy wins for
+// the allocations inside it, the module policy still applies elsewhere.
+//
+// <2000xi32> is 8000 bytes, above the option threshold too, so the function
+// sends it to the heap in both runs.
+
+// CHECK-LABEL: func.func @function_policy
+// CHECK: fir.allocmem !fir.array<2000xi32>
+// CHECK-LABEL: func.func @module_policy
+// CHECK: fir.alloca !fir.array<2000xi32>
+// OPTION-LABEL: func.func @function_policy
+// OPTION: fir.allocmem !fir.array<2000xi32>
+// OPTION-LABEL: func.func @module_policy
+// OPTION: fir.alloca !fir.array<2000xi32>
+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"} {
+func.func @function_policy() attributes {fir.allocation_policy =
+ #fir.allocation_policy<stack_arrays = false,
+ small_array_threshold = 64,
+ total_stack_limit = 4194304>} {
+ %0 = fir.allocmem !fir.array<2000xi32>
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %r = fir.convert %0 : (!fir.heap<!fir.array<2000xi32>>) -> !fir.ref<!fir.array<2000xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<2000xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<2000xi32>>
+ return
+}
+func.func @module_policy() {
+ %0 = fir.allocmem !fir.array<2000xi32>
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %r = fir.convert %0 : (!fir.heap<!fir.array<2000xi32>>) -> !fir.ref<!fir.array<2000xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<2000xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<2000xi32>>
+ return
+}
+}
diff --git a/flang/test/Transforms/stack-arrays-policy-opt-out.fir b/flang/test/Transforms/stack-arrays-policy-opt-out.fir
new file mode 100644
index 0000000000000..52f75a7da398a
--- /dev/null
+++ b/flang/test/Transforms/stack-arrays-policy-opt-out.fir
@@ -0,0 +1,35 @@
+// Test that a function opting out of -fstack-arrays in its own policy keeps its
+// heap allocations, while a function without a policy of its own is unaffected.
+
+// RUN: fir-opt --stack-arrays %s | FileCheck %s
+
+// CHECK-LABEL: func.func @opted_out
+// CHECK: fir.allocmem !fir.array<42xi32>
+// CHECK-NOT: fir.alloca
+func.func @opted_out() attributes {fir.allocation_policy =
+ #fir.allocation_policy<stack_arrays = false,
+ small_array_threshold = 1024,
+ total_stack_limit = 4194304>} {
+ %0 = fir.allocmem !fir.array<42xi32>
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %r = fir.convert %0 : (!fir.heap<!fir.array<42xi32>>) -> !fir.ref<!fir.array<42xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<42xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<42xi32>>
+ return
+}
+
+// CHECK-LABEL: func.func @no_policy
+// CHECK: fir.alloca !fir.array<42xi32>
+// CHECK-NOT: fir.allocmem
+func.func @no_policy() {
+ %0 = fir.allocmem !fir.array<42xi32>
+ %c0 = arith.constant 0 : index
+ %v = arith.constant 0 : i32
+ %r = fir.convert %0 : (!fir.heap<!fir.array<42xi32>>) -> !fir.ref<!fir.array<42xi32>>
+ %e = fir.coordinate_of %r, %c0 : (!fir.ref<!fir.array<42xi32>>, index) -> !fir.ref<i32>
+ fir.store %v to %e : !fir.ref<i32>
+ fir.freemem %0 : !fir.heap<!fir.array<42xi32>>
+ return
+}
>From a7b973554b538f029bf0838338775c2bd3a27500 Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Fri, 11 Sep 2026 17:36:54 -0700
Subject: [PATCH 2/2] record the device allocation policy unconditionally
---
flang/lib/Lower/CallInterface.cpp | 13 ++++++-------
flang/test/Lower/CUDA/cuda-allocatable.cuf | 2 +-
.../Lower/CUDA/cuda-allocation-policy.cuf | 19 +++++++++++--------
flang/test/Lower/CUDA/cuda-atomicadd.cuf | 8 ++++----
flang/test/Lower/CUDA/cuda-cdevloc.cuf | 2 +-
flang/test/Lower/CUDA/cuda-cluster.cuf | 6 +++---
flang/test/Lower/CUDA/cuda-device-proc.cuf | 4 ++--
flang/test/Lower/CUDA/cuda-mod.cuf | 2 +-
flang/test/Lower/CUDA/cuda-module-use.cuf | 4 ++--
flang/test/Lower/CUDA/cuda-proc-attribute.cuf | 16 ++++++++--------
flang/test/Lower/CUDA/cuda-shared.cuf | 2 +-
.../test/Lower/CUDA/cuda-synchronization.cuf | 2 +-
12 files changed, 41 insertions(+), 39 deletions(-)
diff --git a/flang/lib/Lower/CallInterface.cpp b/flang/lib/Lower/CallInterface.cpp
index 0709986f692a0..9885345bed3a5 100644
--- a/flang/lib/Lower/CallInterface.cpp
+++ b/flang/lib/Lower/CallInterface.cpp
@@ -699,18 +699,17 @@ setCUDAAttributes(mlir::func::FuncOp func,
auto procAttr = cuf::getProcAttribute(func.getContext(),
*characteristic->cudaSubprogramAttrs);
func.getOperation()->setAttr(cuf::getProcAttrName(), procAttr);
- // -fstack-arrays cannot be honored in device code: the device stack is
- // orders of magnitude smaller, and an automatic array that fits the host
- // stack easily overflows it. host_device is the host copy of the routine.
+ // -fstack-arrays cannot be honored in device code: the device stack is far
+ // smaller, and an automatic array that fits the host stack overflows it.
+ // Recorded unconditionally so the opt-out is explicit in the IR, as on the
+ // module. host_device is the host copy of the routine.
cuf::ProcAttribute proc = procAttr.getValue();
if (proc != cuf::ProcAttribute::Host &&
proc != cuf::ProcAttribute::HostDevice) {
fir::AllocationPolicy policy =
fir::getAllocationPolicy(func.getOperation());
- if (policy.stackArrays) {
- policy.stackArrays = false;
- fir::setAllocationPolicy(func.getOperation(), policy);
- }
+ policy.stackArrays = false;
+ fir::setAllocationPolicy(func.getOperation(), policy);
}
}
diff --git a/flang/test/Lower/CUDA/cuda-allocatable.cuf b/flang/test/Lower/CUDA/cuda-allocatable.cuf
index 3b641c72d004e..7e03a56e5d94b 100644
--- a/flang/test/Lower/CUDA/cuda-allocatable.cuf
+++ b/flang/test/Lower/CUDA/cuda-allocatable.cuf
@@ -183,7 +183,7 @@ attributes(global) subroutine sub8()
deallocate(a)
end subroutine
-! CHECK-LABEL: func.func @_QPsub8() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPsub8() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %[[DESC:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xf32>>> {bindc_name = "a", uniq_name = "_QFsub8Ea"}
! CHECK: %[[A:.*]]:2 = hlfir.declare %[[DESC]] {data_attr = #cuf.cuda<device>, fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFsub8Ea"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>)
! CHECK: %[[HEAP:.*]] = fir.allocmem !fir.array<?xf32>, %{{.*}} {alignment = 64 : i64, fir.must_be_heap = true, uniq_name = "_QFsub8Ea.alloc"}
diff --git a/flang/test/Lower/CUDA/cuda-allocation-policy.cuf b/flang/test/Lower/CUDA/cuda-allocation-policy.cuf
index 8749d6b78e7da..a5983c16e62ae 100644
--- a/flang/test/Lower/CUDA/cuda-allocation-policy.cuf
+++ b/flang/test/Lower/CUDA/cuda-allocation-policy.cuf
@@ -1,9 +1,11 @@
! Test that -fstack-arrays is not applied to device code. The device stack is
! far smaller than the host one, so lowering records a policy of its own on
-! device and global procedures to keep their automatic arrays on the heap.
+! device and global procedures to keep their automatic arrays on the heap. 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 %s -o - | FileCheck %s --check-prefix=DEFAULT
-! RUN: %flang_fc1 -emit-fir -fstack-arrays %s -o - | FileCheck %s
+! RUN: %flang_fc1 -emit-fir %s -o - | FileCheck %s --check-prefixes=CHECK,DEFAULT
+! RUN: %flang_fc1 -emit-fir -fstack-arrays %s -o - | FileCheck %s --check-prefixes=CHECK,STACK
module m
contains
@@ -38,12 +40,13 @@ contains
end subroutine
end module
-! Without -fstack-arrays the module policy already keeps automatic arrays on the
-! heap, so no procedure needs a policy of its own.
-! DEFAULT-NOT: func.func{{.*}}fir.allocation_policy
+! 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
-! CHECK: func.func @_QMmPkernel({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false,
-! CHECK: func.func @_QMmPdevsub({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false,
+! Device procedures opt out of it in both cases.
+! CHECK: func.func @_QMmPkernel({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false
+! CHECK: func.func @_QMmPdevsub({{.*}}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.allocation_policy = #fir.allocation_policy<stack_arrays = false
! The host procedure keeps the module policy, so -fstack-arrays still applies
! to it.
diff --git a/flang/test/Lower/CUDA/cuda-atomicadd.cuf b/flang/test/Lower/CUDA/cuda-atomicadd.cuf
index 573e01242c78f..dff029d33930d 100644
--- a/flang/test/Lower/CUDA/cuda-atomicadd.cuf
+++ b/flang/test/Lower/CUDA/cuda-atomicadd.cuf
@@ -7,7 +7,7 @@ attributes(global) subroutine test_atomicaddvector_r2()
tmp1 = atomicAddVector(a, tmp2)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r2() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r2() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>
attributes(global) subroutine test_atomicaddvector_r4()
@@ -15,7 +15,7 @@ attributes(global) subroutine test_atomicaddvector_r4()
tmp1 = atomicAddVector(a, tmp2)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
attributes(global) subroutine test_atomicadd_r2x4()
@@ -23,7 +23,7 @@ attributes(global) subroutine test_atomicadd_r2x4()
tmp1 = atomicaddreal4x2(a, tmp2)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_atomicadd_r2x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_atomicadd_r2x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
attributes(global) subroutine test_atomicadd_r4x4()
@@ -31,5 +31,5 @@ attributes(global) subroutine test_atomicadd_r4x4()
tmp1 = atomicaddreal4x4(a, tmp2)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_atomicadd_r4x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_atomicadd_r4x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: atom.add.v4.f32
diff --git a/flang/test/Lower/CUDA/cuda-cdevloc.cuf b/flang/test/Lower/CUDA/cuda-cdevloc.cuf
index d663e6eda478b..aaf08a2a53091 100644
--- a/flang/test/Lower/CUDA/cuda-cdevloc.cuf
+++ b/flang/test/Lower/CUDA/cuda-cdevloc.cuf
@@ -7,7 +7,7 @@ attributes(global) subroutine testcdevloc(a)
end
! CHECK-LABEL: func.func @_QPtestcdevloc(
-! CHECK-SAME: %[[A_ARG:.*]]: !fir.ref<!fir.array<10xi32>> {cuf.data_attr = #cuf.cuda<device>, fir.bindc_name = "a"}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-SAME: %[[A_ARG:.*]]: !fir.ref<!fir.array<10xi32>> {cuf.data_attr = #cuf.cuda<device>, fir.bindc_name = "a"}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %[[A:.*]]:2 = hlfir.declare %[[A_ARG]](%{{.*}}) dummy_scope %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtestcdevlocEa"} : (!fir.ref<!fir.array<10xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<10xi32>>, !fir.ref<!fir.array<10xi32>>)
! CHECK: %[[A1:.*]] = hlfir.designate %[[A]]#0 (%c1{{.*}}) : (!fir.ref<!fir.array<10xi32>>, index) -> !fir.ref<i32>
! CHECK: %[[BOX:.*]] = fir.embox %[[A1]] : (!fir.ref<i32>) -> !fir.box<i32>
diff --git a/flang/test/Lower/CUDA/cuda-cluster.cuf b/flang/test/Lower/CUDA/cuda-cluster.cuf
index 78cca15b11dab..71e8a91506912 100644
--- a/flang/test/Lower/CUDA/cuda-cluster.cuf
+++ b/flang/test/Lower/CUDA/cuda-cluster.cuf
@@ -7,7 +7,7 @@ attributes(global) subroutine test_this_cluster()
cluster = this_cluster()
end subroutine
-! CHECK-LABEL: func.func @_QPtest_this_cluster() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_this_cluster() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %{{.*}} = fir.alloca !fir.type<_QMcooperative_groupsTcluster_group
! CHECK: %[[RES:.*]] = fir.alloca !fir.type<_QMcooperative_groupsTcluster_group{_QMcooperative_groupsTcluster_group.handle:!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>,size:i32,rank:i32}>
! CHECK: %[[RANK:.*]] = nvvm.read.ptx.sreg.cluster.ctarank : i32
@@ -22,7 +22,7 @@ attributes(global) subroutine test_cluster_dim_blocks()
clusterDim = cluster_dim_blocks()
end subroutine
-! CHECK-LABEL: func.func @_QPtest_cluster_dim_blocks() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_cluster_dim_blocks() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %[[X:.*]] = nvvm.read.ptx.sreg.cluster.nctaid.x : i32
! CHECK: %[[COORD_X:.*]] = fir.coordinate_of %{{.*}}, x : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>
! CHECK: fir.store %[[X]] to %[[COORD_X]] : !fir.ref<i32>
@@ -40,7 +40,7 @@ attributes(global) subroutine test_cluster_block_index()
blockIndex = cluster_block_index()
end subroutine
-! CHECK-LABEL: func.func @_QPtest_cluster_block_index() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPtest_cluster_block_index() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %[[X:.*]] = nvvm.read.ptx.sreg.cluster.ctaid.x : i32
! CHECK: %[[X1:.*]] = arith.addi %[[X]], %c1{{.*}} : i32
! CHECK: %[[COORD_X:.*]] = fir.coordinate_of %{{.*}}, x : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 97fb1ceefd8d2..c70c7e661f3c3 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -100,7 +100,7 @@ attributes(global) subroutine devsub()
ad = __ll2double_rz(al)
end
-! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: nvvm.barrier
! CHECK: nvvm.bar.warp.sync %c1{{.*}} : i32
! CHECK: %{{.*}} = nvvm.barrier.reduction #nvvm.reduction<and> %c1{{.*}} -> i32
@@ -779,4 +779,4 @@ end subroutine
! CHECK-LABEL: func.func @_QPcall_c_kernel()
! CHECK: cuf.kernel_launch @c_kernel<<<{{.*}}>>>()
-! CHECK: func.func private @c_kernel() attributes {cuf.proc_attr = #cuf.cuda_proc<global>, fir.bindc_name = "c_kernel", fir.proc_attrs = #fir.proc_attrs<bind_c>}
+! CHECK: func.func private @c_kernel() attributes {cuf.proc_attr = #cuf.cuda_proc<global>,{{.*}}fir.bindc_name = "c_kernel", fir.proc_attrs = #fir.proc_attrs<bind_c>}
diff --git a/flang/test/Lower/CUDA/cuda-mod.cuf b/flang/test/Lower/CUDA/cuda-mod.cuf
index cb1a6d8ad28ec..b6f41925102aa 100644
--- a/flang/test/Lower/CUDA/cuda-mod.cuf
+++ b/flang/test/Lower/CUDA/cuda-mod.cuf
@@ -10,6 +10,6 @@ contains
end
end module
-! CHECK: func.func @_QMcuf_modPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<device>}
+! CHECK: func.func @_QMcuf_modPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<device>{{.*}}}
! CHECK: fir.global @_QMcuf_modEmd {data_attr = #cuf.cuda<device>} : f32
diff --git a/flang/test/Lower/CUDA/cuda-module-use.cuf b/flang/test/Lower/CUDA/cuda-module-use.cuf
index 130fefab24d90..e1f5ae71fe5f3 100644
--- a/flang/test/Lower/CUDA/cuda-module-use.cuf
+++ b/flang/test/Lower/CUDA/cuda-module-use.cuf
@@ -17,9 +17,9 @@ attributes(device) subroutine sub2()
call devsub()
end
-! CHECK-LABEL: func.func @_QPsub2() attributes {cuf.proc_attr = #cuf.cuda_proc<device>}
+! CHECK-LABEL: func.func @_QPsub2() attributes {cuf.proc_attr = #cuf.cuda_proc<device>{{.*}}}
! CHECK: fir.call @_QMcuf_modPdevsub()
! CHECK-LABEL: fir.global @_QMcuf_modEmd {data_attr = #cuf.cuda<device>} : f32
-! CHECK-LABEL: func.func private @_QMcuf_modPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<device>}
+! CHECK-LABEL: func.func private @_QMcuf_modPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<device>{{.*}}}
diff --git a/flang/test/Lower/CUDA/cuda-proc-attribute.cuf b/flang/test/Lower/CUDA/cuda-proc-attribute.cuf
index dc3c925f888b6..b528b251b39b2 100644
--- a/flang/test/Lower/CUDA/cuda-proc-attribute.cuf
+++ b/flang/test/Lower/CUDA/cuda-proc-attribute.cuf
@@ -7,7 +7,7 @@ attributes(host) subroutine sub_host(); end
! CHECK: func.func @_QPsub_host() attributes {cuf.proc_attr = #cuf.cuda_proc<host>}
attributes(device) subroutine sub_device(); end
-! CHECK: func.func @_QPsub_device() attributes {cuf.proc_attr = #cuf.cuda_proc<device>}
+! CHECK: func.func @_QPsub_device() attributes {cuf.proc_attr = #cuf.cuda_proc<device>{{.*}}}
attributes(host) attributes(device) subroutine sub_host_device; end
! CHECK: func.func @_QPsub_host_device() attributes {cuf.proc_attr = #cuf.cuda_proc<host_device>}
@@ -16,16 +16,16 @@ attributes(device) attributes(host) subroutine sub_device_host; end
! CHECK: func.func @_QPsub_device_host() attributes {cuf.proc_attr = #cuf.cuda_proc<host_device>}
attributes(global) subroutine sub_global(); end
-! CHECK: func.func @_QPsub_global() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: func.func @_QPsub_global() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
attributes(grid_global) subroutine sub_grid_global(); end
-! CHECK: func.func @_QPsub_grid_global() attributes {cuf.proc_attr = #cuf.cuda_proc<grid_global>}
+! CHECK: func.func @_QPsub_grid_global() attributes {cuf.proc_attr = #cuf.cuda_proc<grid_global>{{.*}}}
attributes(host) integer function fct_host(); end
! CHECK: func.func @_QPfct_host() -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<host>}
attributes(device) integer function fct_device(); end
-! CHECK: func.func @_QPfct_device() -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>}
+! CHECK: func.func @_QPfct_device() -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>{{.*}}}
attributes(host) attributes(device) integer function fct_host_device; end
! CHECK: func.func @_QPfct_host_device() -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<host_device>}
@@ -34,13 +34,13 @@ attributes(device) attributes(host) integer function fct_device_host; end
! CHECK: func.func @_QPfct_device_host() -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<host_device>}
attributes(global) launch_bounds(1) subroutine sub_lbounds0(); end
-! CHECK: func.func @_QPsub_lbounds0() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: func.func @_QPsub_lbounds0() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
attributes(global) launch_bounds(1, 2) subroutine sub_lbounds1(); end
-! CHECK: func.func @_QPsub_lbounds1() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64, minBPM = 2 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: func.func @_QPsub_lbounds1() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64, minBPM = 2 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
attributes(global) launch_bounds(1, 2, 3) subroutine sub_lbounds2(); end
-! CHECK: func.func @_QPsub_lbounds2() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64, minBPM = 2 : i64, upperBoundClusterSize = 3 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: func.func @_QPsub_lbounds2() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 1 : i64, minBPM = 2 : i64, upperBoundClusterSize = 3 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
attributes(global) cluster_dims(1, 2, 3) subroutine sub_clusterdims1(); end
-! CHECK: func.func @_QPsub_clusterdims1() attributes {cuf.cluster_dims = #cuf.cluster_dims<x = 1 : i64, y = 2 : i64, z = 3 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: func.func @_QPsub_clusterdims1() attributes {cuf.cluster_dims = #cuf.cluster_dims<x = 1 : i64, y = 2 : i64, z = 3 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
diff --git a/flang/test/Lower/CUDA/cuda-shared.cuf b/flang/test/Lower/CUDA/cuda-shared.cuf
index f41011df06ae7..4fe6e29a77256 100644
--- a/flang/test/Lower/CUDA/cuda-shared.cuf
+++ b/flang/test/Lower/CUDA/cuda-shared.cuf
@@ -7,6 +7,6 @@ attributes(global) subroutine sharedmem()
s(t) = t
end subroutine
-! CHECK-LABEL: func.func @_QPsharedmem() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPsharedmem() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: %{{.*}} = cuf.shared_memory !fir.array<32xf32> {bindc_name = "s", uniq_name = "_QFsharedmemEs"} -> !fir.ref<!fir.array<32xf32>>
! CHECK-NOT: cuf.free
diff --git a/flang/test/Lower/CUDA/cuda-synchronization.cuf b/flang/test/Lower/CUDA/cuda-synchronization.cuf
index 7b6d4183be09b..ce7af8343abc0 100644
--- a/flang/test/Lower/CUDA/cuda-synchronization.cuf
+++ b/flang/test/Lower/CUDA/cuda-synchronization.cuf
@@ -8,7 +8,7 @@ attributes(global) subroutine sync()
call threadfence_system()
end subroutine
-! CHECK-LABEL: func.func @_QPsync() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK-LABEL: func.func @_QPsync() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: nvvm.memory.barrier gpu
! CHECK: nvvm.memory.barrier cta
! CHECK: nvvm.memory.barrier sys
More information about the flang-commits
mailing list