[clang] [llvm] [Clang][OpenMP] Don't mark outlined region functions norecurse (PR #218862)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 26 01:37:21 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
Author: Larry Meadows (lfmeadow)
<details>
<summary>Changes</summary>
Clang marks every outlined OpenMP region function `norecurse`, in both
`emitOutlinedFunctionPrologue` and `emitOutlinedFunctionPrologueAggregate`, and
does the same for the GPU parallel wrapper. For a `parallel` region that is not
true: the body is reached through the function pointer the runtime is handed,
and if the body opens another parallel region the runtime's parallel entry is
re-entered, thus recursive.
This PR drops the three `setDoesNotRecurse()` calls.
This supersedes #<!-- -->218637, which dropped the attribute in the Attributor when
specialization exposed the cycle; thanks @<!-- -->nikic for pointing at the annotation
instead.
## Testing
Tested on `main` at `7380050ae1a8`, gfx90a.
- `clang/test`: 48936 passed, 27 expectedly failed, 0 failed. `clang/test/OpenMP`
is 1594 passed, 13 unsupported.
- `offload/test`, all three configurations (amdgpu, host, unit): 1126 passed,
357 unsupported, 4 expectedly failed, 0 failed.
- `nested_parallel_reduction.c` prints `aa = 6300, expected 6300` with the fix
and faults with `OFFLOAD ERROR: memory access fault` without it.
---
Full diff: https://github.com/llvm/llvm-project/pull/218862.diff
9 Files Affected:
- (modified) clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp (-1)
- (modified) clang/lib/CodeGen/CGStmtOpenMP.cpp (-2)
- (modified) clang/test/OpenMP/amdgcn-attributes.cpp (+2-2)
- (modified) clang/test/OpenMP/parallel_if_codegen_PR51349.cpp (+4-4)
- (modified) clang/test/utils/update_cc_test_checks/Inputs/generated-funcs-regex.c.expected (+1-1)
- (modified) clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.generated.all.expected (+1-2)
- (modified) clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.no-generated.all.expected (+1-2)
- (modified) clang/test/utils/update_cc_test_checks/Inputs/replace-value-regex-across-runs.c.expected (+1-1)
- (added) offload/test/offloading/nested_parallel_reduction.c (+39)
``````````diff
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 90ae21283d031..e785ee635af70 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -2034,7 +2034,6 @@ llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
- Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index a761ac6bbf816..a5e73f6c917ca 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -632,7 +632,6 @@ static llvm::Function *emitOutlinedFunctionPrologue(
if (CD->isNothrow())
F->setDoesNotThrow();
- F->setDoesNotRecurse();
// Always inline the outlined function if optimizations are enabled.
if (CGM.getCodeGenOpts().OptimizationLevel != 0) {
@@ -745,7 +744,6 @@ static llvm::Function *emitOutlinedFunctionPrologueAggregate(
CGM.SetInternalFunctionAttributes(CD, F, FuncInfo);
if (CD->isNothrow())
F->setDoesNotThrow();
- F->setDoesNotRecurse();
CGF.StartFunction(CD, Ctx.VoidTy, F, FuncInfo, Args, Loc, Loc);
Address ContextAddr = CGF.GetAddrOfLocalVar(CD->getContextParam());
diff --git a/clang/test/OpenMP/amdgcn-attributes.cpp b/clang/test/OpenMP/amdgcn-attributes.cpp
index 789c8d6d72218..a1921f5cd3f61 100644
--- a/clang/test/OpenMP/amdgcn-attributes.cpp
+++ b/clang/test/OpenMP/amdgcn-attributes.cpp
@@ -31,8 +31,8 @@ int callable(int x) {
return x + 1;
}
-// DEFAULT: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
-// NOIEEE: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "amdgpu-ieee"="false" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// DEFAULT: attributes #0 = { convergent mustprogress noinline nounwind optnone "amdgpu-flat-work-group-size"="1,42" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// NOIEEE: attributes #0 = { convergent mustprogress noinline nounwind optnone "amdgpu-flat-work-group-size"="1,42" "amdgpu-ieee"="false" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// DEFAULT: attributes #2 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
// NOIEEE: attributes #2 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
diff --git a/clang/test/OpenMP/parallel_if_codegen_PR51349.cpp b/clang/test/OpenMP/parallel_if_codegen_PR51349.cpp
index 996498c25f024..2dbfca5366350 100644
--- a/clang/test/OpenMP/parallel_if_codegen_PR51349.cpp
+++ b/clang/test/OpenMP/parallel_if_codegen_PR51349.cpp
@@ -30,7 +30,7 @@ void foo() {
// CHECK-NEXT: ret void
//
//
-// CHECK: Function Attrs: noinline norecurse nounwind
+// CHECK: Function Attrs: noinline nounwind
// CHECK-LABEL: define internal void @_Z3foov.omp_outlined(
// CHECK-SAME: ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -41,7 +41,7 @@ void foo() {
// CHECK-NEXT: ret void
//
//
-// CHECK: Function Attrs: alwaysinline norecurse nounwind
+// CHECK: Function Attrs: alwaysinline nounwind
// CHECK-LABEL: define internal void @_Z3foov.omp_outlined.1(
// CHECK-SAME: ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -68,7 +68,7 @@ void foo() {
// CHECK-NOINLINE-NEXT: ret void
//
//
-// CHECK-NOINLINE: Function Attrs: noinline norecurse nounwind
+// CHECK-NOINLINE: Function Attrs: noinline nounwind
// CHECK-NOINLINE-LABEL: define internal void @_Z3foov.omp_outlined(
// CHECK-NOINLINE-SAME: ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK-NOINLINE-NEXT: [[ENTRY:.*:]]
@@ -79,7 +79,7 @@ void foo() {
// CHECK-NOINLINE-NEXT: ret void
//
//
-// CHECK-NOINLINE: Function Attrs: alwaysinline norecurse nounwind
+// CHECK-NOINLINE: Function Attrs: alwaysinline nounwind
// CHECK-NOINLINE-LABEL: define internal void @_Z3foov.omp_outlined.1(
// CHECK-NOINLINE-SAME: ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3:[0-9]+]] {
// CHECK-NOINLINE-NEXT: [[ENTRY:.*:]]
diff --git a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs-regex.c.expected b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs-regex.c.expected
index 9f1639da282d7..04b4f15b1d331 100644
--- a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs-regex.c.expected
+++ b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs-regex.c.expected
@@ -21,7 +21,7 @@ void foo(int a)
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
// CHECK-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8
-// CHECK-NEXT: call void @{{__omp_offloading_[a-z0-9]+_[a-z0-9]+_foo_l[0-9]+}}(i64 [[TMP1]], ptr null) #[[ATTR3:[0-9]+]]
+// CHECK-NEXT: call void @{{__omp_offloading_[a-z0-9]+_[a-z0-9]+_foo_l[0-9]+}}(i64 [[TMP1]], ptr null) #[[ATTR2:[0-9]+]]
// CHECK-NEXT: call void @{{__test_offloading_[a-z0-9]+_[a-z0-9]+_bar_l[0-9]+}}()
// CHECK-NEXT: ret void
//
diff --git a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.generated.all.expected b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.generated.all.expected
index 29a7bafe183cb..e9acc9ccbe71f 100644
--- a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.generated.all.expected
+++ b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.generated.all.expected
@@ -239,8 +239,7 @@ void foo(void) {
//
//.
// OMP: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
-// OMP: attributes #[[ATTR1:[0-9]+]] = { noinline norecurse nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
-// OMP: attributes #[[ATTR2:[0-9]+]] = { nounwind }
+// OMP: attributes #[[ATTR1:[0-9]+]] = { nounwind }
//.
// NOOMP: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
//.
diff --git a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.no-generated.all.expected b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.no-generated.all.expected
index 7b4afbb5df927..92c7af7eb594f 100644
--- a/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.no-generated.all.expected
+++ b/clang/test/utils/update_cc_test_checks/Inputs/generated-funcs.c.no-generated.all.expected
@@ -110,8 +110,7 @@ void foo(void) {
}
//.
// OMP: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
-// OMP: attributes #[[ATTR1:[0-9]+]] = { noinline norecurse nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
-// OMP: attributes #[[ATTR2:[0-9]+]] = { nounwind }
+// OMP: attributes #[[ATTR1:[0-9]+]] = { nounwind }
//.
// NOOMP: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
//.
diff --git a/clang/test/utils/update_cc_test_checks/Inputs/replace-value-regex-across-runs.c.expected b/clang/test/utils/update_cc_test_checks/Inputs/replace-value-regex-across-runs.c.expected
index 608eac6df699c..697433fc1c398 100644
--- a/clang/test/utils/update_cc_test_checks/Inputs/replace-value-regex-across-runs.c.expected
+++ b/clang/test/utils/update_cc_test_checks/Inputs/replace-value-regex-across-runs.c.expected
@@ -6,7 +6,7 @@
// CHECK-LABEL: @foo(
// CHECK-NEXT: entry:
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_foo_l7(ptr null) #[[ATTR2:[0-9]+]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_foo_l7(ptr null) #[[ATTR1:[0-9]+]]
// CHECK-NEXT: ret void
//
void foo(void) {
diff --git a/offload/test/offloading/nested_parallel_reduction.c b/offload/test/offloading/nested_parallel_reduction.c
new file mode 100644
index 0000000000000..862aeba057326
--- /dev/null
+++ b/offload/test/offloading/nested_parallel_reduction.c
@@ -0,0 +1,39 @@
+// Both parallel regions have to be nested for this to reproduce, and only at
+// -O2 and above; either one alone is correct at every optimization level.
+
+// RUN: %libomptarget-compileopt-generic
+// RUN: %libomptarget-run-generic | %fcheck-generic
+
+// REQUIRES: gpu
+
+#include <stdio.h>
+
+#define N 5
+
+int main(void) {
+ long aa = 0;
+ int ng = 6, cmom = 4, nxyz = 5;
+
+#pragma omp target teams distribute num_teams(nxyz) thread_limit(4) \
+ map(tofrom : aa)
+ for (int gid = 0; gid < nxyz; gid++) {
+#pragma omp parallel for collapse(2)
+ for (unsigned g = 0; g < ng; g++)
+ for (unsigned l = 0; l < cmom - 1; l++) {
+ int a = 0;
+ for (int ii = 0; ii < N + 2; ii++) {
+#pragma omp parallel for reduction(+ : a)
+ for (int i = 0; i < N; i++)
+ a += i;
+ }
+#pragma omp atomic
+ aa += a;
+ }
+ }
+
+ long expected = (long)ng * (cmom - 1) * nxyz * (N * (N - 1) / 2) * (N + 2);
+ printf("aa = %ld, expected %ld\n", aa, expected);
+ return aa != expected;
+}
+
+// CHECK: aa = 6300, expected 6300
``````````
</details>
https://github.com/llvm/llvm-project/pull/218862
More information about the llvm-commits
mailing list