https://github.com/lfmeadow created https://github.com/llvm/llvm-project/pull/218862
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. >From c2691e868fb82b9b955e23b46469fbdf42d14da9 Mon Sep 17 00:00:00 2001 From: Larry Meadows <[email protected]> Date: Wed, 26 Aug 2026 03:23:29 -0500 Subject: [PATCH] [Clang][OpenMP] Don't mark outlined region functions norecurse An outlined parallel region body is reached through the runtime's indirect call, and if it opens another parallel region the runtime entry is re-entered while the first call is still live, so both the entry and the region body sit inside a dynamic call graph cycle. norecurse rules that out for every function occurring in such a cycle, so the annotation is wrong for these functions, and rpo-function-attrs propagates it from them onto the device runtime's parallel entry once internalization makes it internal. Co-authored-by: Cursor <[email protected]> --- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 1 - clang/lib/CodeGen/CGStmtOpenMP.cpp | 2 - clang/test/OpenMP/amdgcn-attributes.cpp | 4 +- .../OpenMP/parallel_if_codegen_PR51349.cpp | 8 ++-- .../Inputs/generated-funcs-regex.c.expected | 2 +- .../generated-funcs.c.generated.all.expected | 3 +- ...enerated-funcs.c.no-generated.all.expected | 3 +- ...replace-value-regex-across-runs.c.expected | 2 +- .../offloading/nested_parallel_reduction.c | 39 +++++++++++++++++++ 9 files changed, 49 insertions(+), 15 deletions(-) create mode 100644 offload/test/offloading/nested_parallel_reduction.c 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
