Author: Robert Imschweiler Date: 2026-08-28T14:16:08+02:00 New Revision: d3480d8ca8d2616f8e484086c57bac1791d64ec0
URL: https://github.com/llvm/llvm-project/commit/d3480d8ca8d2616f8e484086c57bac1791d64ec0 DIFF: https://github.com/llvm/llvm-project/commit/d3480d8ca8d2616f8e484086c57bac1791d64ec0.diff LOG: [clang][OpenMP] Fix thread count upper bound for target directives (#218018) Fix the UpperBound-computing ternary, which was inverted: std::min is only reachable when UpperBound is 0, where it is a no-op, so any clause with a bound to combine simply overwrites it. Two consequences: - A clause can raise the bound above a smaller one. 'target teams distribute parallel for thread_limit(8) num_threads(64)' gets a bound of 64. - A non-constant clause pins the bound at 0 and discards every constant clause after it, losing the bound entirely. Additionally, let a construct split over separate 'target', 'teams' and worksharing directives find its num_threads clause, so it agrees with the combined spelling. Note that the codegen changes in `clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp` and `clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp` are incidental. `CGOpenMPRuntime::emitNumThreadsForTargetDirective` special-cases the value 1 for passing to the host, which is why the codegen changes for 1 but not for 5. Claude assisted with this patch. Added: Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp clang/test/OpenMP/thread_limit_gpu.c Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 9f8822a37bf71..1a1479b4b9b7e 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6649,6 +6649,20 @@ llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective( return llvm::ConstantInt::getSigned(CGF.Int32Ty, MinNT); } +/// Merge the thread count upper bound \p Val into \p UpperBound. +/// +/// \p UpperBound is -1 while no thread limiting clause has been seen, 0 once +/// one has been seen whose value is not known at compile time, and otherwise +/// the smallest constant bound found so far. +/// +/// Thread limiting clauses compose by taking the minimum, so a constant bound +/// stays valid whatever the clauses that are not compile time constants +/// evaluate to. That makes it correct to replace the 0 marker with \p Val, and +/// necessary to keep a clause from raising a smaller bound found earlier. +static void mergeThreadCountUpperBound(int32_t &UpperBound, int32_t Val) { + UpperBound = UpperBound > 0 ? std::min(UpperBound, Val) : Val; +} + /// Check for a num threads constant value (stored in \p DefaultVal), or /// expression (stored in \p E). If the value is conditional (via an if-clause), /// store the condition in \p CondVal. If \p E, and \p CondVal respectively, are @@ -6712,11 +6726,8 @@ static void getNumThreads(CodeGenFunction &CGF, const CapturedStmt *CS, const Expr *NTExpr = NumThreadsClause->getNumThreads().front(); if (NTExpr->isIntegerConstantExpr(CGF.getContext())) if (auto Constant = NTExpr->getIntegerConstantExpr(CGF.getContext())) - UpperBound = - UpperBound - ? Constant->getZExtValue() - : std::min(UpperBound, - static_cast<int32_t>(Constant->getZExtValue())); + mergeThreadCountUpperBound( + UpperBound, static_cast<int32_t>(Constant->getZExtValue())); // If we haven't found a upper bound, remember we saw a thread limiting // clause. if (UpperBound == -1) @@ -6760,9 +6771,8 @@ const Expr *CGOpenMPRuntime::getNumThreadsExprForTargetDirective( auto CheckForConstExpr = [&](const Expr *E, const Expr **EPtr) { if (E->isIntegerConstantExpr(CGF.getContext())) { if (auto Constant = E->getIntegerConstantExpr(CGF.getContext())) - UpperBound = UpperBound ? Constant->getZExtValue() - : std::min(UpperBound, - int32_t(Constant->getZExtValue())); + mergeThreadCountUpperBound( + UpperBound, static_cast<int32_t>(Constant->getZExtValue())); } // If we haven't found a upper bound, remember we saw a thread limiting // clause. @@ -6817,6 +6827,17 @@ const Expr *CGOpenMPRuntime::getNumThreadsExprForTargetDirective( if (isOpenMPTeamsDirective(Dir->getDirectiveKind()) && !isOpenMPDistributeDirective(Dir->getDirectiveKind())) { CS = Dir->getInnermostCapturedStmt(); + // Now that the 'teams' level has been peeled off, the remainder is + // shaped like a 'target teams' region, so pick up the num_threads of + // the directive nested in it the same way the OMPD_target_teams case + // below does. Without this the upper bound of a construct written as + // 'target' / 'teams' / 'distribute parallel for' would stay at the + // default, while every combined spelling of the same construct honors + // the clause. Only the bound is taken here: passing null for the + // expression and the condition keeps this from emitting anything, so + // the value the host passes to the kernel launch is left as it was. + getNumThreads(CGF, CS, /*E=*/nullptr, UpperBound, UpperBoundOnly, + /*CondVal=*/nullptr); const Stmt *Child = CGOpenMPRuntime::getSingleCompoundChild( CGF.getContext(), CS->getCapturedStmt()); Dir = dyn_cast_or_null<OMPExecutableDirective>(Child); diff --git a/clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp b/clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp index a07ce8902ed52..0ef2a2784a393 100644 --- a/clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp +++ b/clang/test/OpenMP/distribute_parallel_for_num_threads_codegen.cpp @@ -3116,10 +3116,10 @@ int main() { // CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK1-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK1-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK1-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK1: omp_offload.failed: @@ -4441,10 +4441,10 @@ int main() { // CHECK5-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK5-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK5-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK5-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK5-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK5-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK5-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK5-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK5-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK5-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK5-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK5: omp_offload.failed: @@ -5766,10 +5766,10 @@ int main() { // CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK9-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK9-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK9-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK9-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK9-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK9-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK9: omp_offload.failed: @@ -7091,10 +7091,10 @@ int main() { // CHECK13-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK13-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK13-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK13-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK13-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK13-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK13-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK13-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK13-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK13-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK13-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK13: omp_offload.failed: diff --git a/clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp b/clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp index b502884ae51f8..201ce6e4da6ca 100644 --- a/clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp +++ b/clang/test/OpenMP/distribute_parallel_for_simd_num_threads_codegen.cpp @@ -748,10 +748,10 @@ int main() { // CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK1-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK1-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK1-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK1: omp_offload.failed: @@ -2496,10 +2496,10 @@ int main() { // CHECK5-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK5-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK5-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK5-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK5-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK5-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK5-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK5-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK5-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK5-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK5-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK5: omp_offload.failed: @@ -3905,10 +3905,10 @@ int main() { // CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK9-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK9-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK9-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK9-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK9-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK9-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK9: omp_offload.failed: @@ -5653,10 +5653,10 @@ int main() { // CHECK13-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 // CHECK13-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4 // CHECK13-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 -// CHECK13-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 +// CHECK13-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4 // CHECK13-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 // CHECK13-NEXT: store i32 0, ptr [[TMP17]], align 4 -// CHECK13-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) +// CHECK13-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainI1SLi1EEiv_l52.region_id, ptr [[KERNEL_ARGS]]) // CHECK13-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 // CHECK13-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] // CHECK13: omp_offload.failed: diff --git a/clang/test/OpenMP/thread_limit_gpu.c b/clang/test/OpenMP/thread_limit_gpu.c index 1716e16465eca..fcf87ae282cf9 100644 --- a/clang/test/OpenMP/thread_limit_gpu.c +++ b/clang/test/OpenMP/thread_limit_gpu.c @@ -26,6 +26,58 @@ void foo(int N) { #pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 84, 86)))) num_threads(20) for (int i = 0; i < N; ++i) ; + // A construct split over separate 'target', 'teams' and worksharing + // directives describes the same kernel as the combined spelling below and + // must end up with the same thread bounds. +#pragma omp target +#pragma omp teams +#pragma omp distribute parallel for num_threads(19) + for (int i = 0; i < N; ++i) + ; +#pragma omp target teams distribute parallel for num_threads(19) + for (int i = 0; i < N; ++i) + ; + // thread_limit bounds the size of the contention group, so a num_threads + // clause asking for more than that cannot raise the bound. Both spellings + // again have to agree. +#pragma omp target +#pragma omp teams thread_limit(8) +#pragma omp distribute parallel for num_threads(64) + for (int i = 0; i < N; ++i) + ; +#pragma omp target teams distribute parallel for thread_limit(8) num_threads(64) + for (int i = 0; i < N; ++i) + ; + // The other way round: a num_threads below the thread_limit is the bound. A + // 'target' wrapping a combined 'teams distribute parallel for' is scanned for + // num_threads before the thread_limit clause is known, so the thread_limit + // must not overwrite the smaller value it already found. +#pragma omp target +#pragma omp teams distribute parallel for num_threads(5) thread_limit(9) + for (int i = 0; i < N; ++i) + ; +#pragma omp target teams distribute parallel for num_threads(5) thread_limit(9) + for (int i = 0; i < N; ++i) + ; + // A constant thread_limit still bounds the kernel when num_threads is not a + // constant, and must not be lost because the non-constant clause was seen + // first. +#pragma omp target +#pragma omp teams distribute parallel for num_threads(N) thread_limit(7) + for (int i = 0; i < N; ++i) + ; +#pragma omp target teams distribute parallel for num_threads(N) thread_limit(7) + for (int i = 0; i < N; ++i) + ; + // And the other way round: a constant num_threads still bounds the kernel + // when the thread_limit is not a constant. +#pragma omp target +#pragma omp teams distribute parallel for num_threads(6) thread_limit(N) + for (int i = 0; i < N; ++i) + ; +#pragma omp target teams distribute parallel for num_threads(6) thread_limit(N) + for (int i = 0; i < N; ++i) + ; } #endif @@ -35,18 +87,38 @@ void foo(int N) { // CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l20({{.*}}) #[[ATTR3:.+]] { // CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l23({{.*}}) #[[ATTR4:.+]] { // CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l26({{.*}}) #[[ATTR5:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l32({{.*}}) #[[SPLIT:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l37({{.*}}) #[[SPLIT]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l43({{.*}}) #[[CLAMP:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l48({{.*}}) #[[CLAMP]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l55({{.*}}) #[[SMALLER:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l59({{.*}}) #[[SMALLER]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l65({{.*}}) #[[DYN_NT:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l69({{.*}}) #[[DYN_NT]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l74({{.*}}) #[[DYN_TL:.+]] { +// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l78({{.*}}) #[[DYN_TL]] { // CHECK-AMDGPU: attributes #[[ATTR1]] = { {{.*}} "amdgpu-flat-work-group-size"="1,256" {{.*}} } // CHECK-AMDGPU: attributes #[[ATTR2]] = { {{.*}} "amdgpu-flat-work-group-size"="1,4" {{.*}} } // CHECK-AMDGPU: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" {{.*}} } // CHECK-AMDGPU: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" {{.*}} } // CHECK-AMDGPU: attributes #[[ATTR5]] = { {{.*}} "amdgpu-flat-work-group-size"="1,20" "amdgpu-max-num-workgroups"="86,1,1" {{.*}} } +// CHECK-AMDGPU: attributes #[[SPLIT]] = { {{.*}} "amdgpu-flat-work-group-size"="1,19" {{.*}} } +// CHECK-AMDGPU: attributes #[[CLAMP]] = { {{.*}} "amdgpu-flat-work-group-size"="1,8" {{.*}} } +// CHECK-AMDGPU: attributes #[[SMALLER]] = { {{.*}} "amdgpu-flat-work-group-size"="1,5" {{.*}} } +// CHECK-AMDGPU: attributes #[[DYN_NT]] = { {{.*}} "amdgpu-flat-work-group-size"="1,7" {{.*}} } +// CHECK-AMDGPU: attributes #[[DYN_TL]] = { {{.*}} "amdgpu-flat-work-group-size"="1,6" {{.*}} } // CHECK-SPIRV: attributes #[[ATTR1]] = { {{.*}} "omp_target_thread_limit"="256" {{.*}} } // CHECK-SPIRV: attributes #[[ATTR2]] = { {{.*}} "omp_target_thread_limit"="4" {{.*}} } // CHECK-SPIRV: attributes #[[ATTR3]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="42" {{.*}} } // CHECK-SPIRV: attributes #[[ATTR4]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="22" {{.*}} } // CHECK-SPIRV: attributes #[[ATTR5]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="20" {{.*}} } +// CHECK-SPIRV: attributes #[[SPLIT]] = { {{.*}} "omp_target_thread_limit"="19" {{.*}} } +// CHECK-SPIRV: attributes #[[CLAMP]] = { {{.*}} "omp_target_thread_limit"="8" {{.*}} } +// CHECK-SPIRV: attributes #[[SMALLER]] = { {{.*}} "omp_target_thread_limit"="5" {{.*}} } +// CHECK-SPIRV: attributes #[[DYN_NT]] = { {{.*}} "omp_target_thread_limit"="7" {{.*}} } +// CHECK-SPIRV: attributes #[[DYN_TL]] = { {{.*}} "omp_target_thread_limit"="6" {{.*}} } // CHECK-AMDGPU-FLAG: attributes #[[ATTR1]] = { // CHECK-AMDGPU-FLAG-NOT: amdgpu-flat-work-group-size @@ -56,6 +128,11 @@ void foo(int N) { // CHECK-AMDGPU-FLAG: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" {{.*}} } // CHECK-AMDGPU-FLAG: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" {{.*}} } // CHECK-AMDGPU-FLAG: attributes #[[ATTR5]] = { {{.*}} "amdgpu-flat-work-group-size"="1,20" "amdgpu-max-num-workgroups"="86,1,1" {{.*}} } +// CHECK-AMDGPU-FLAG: attributes #[[SPLIT]] = { {{.*}} "amdgpu-flat-work-group-size"="1,19" {{.*}} } +// CHECK-AMDGPU-FLAG: attributes #[[CLAMP]] = { {{.*}} "amdgpu-flat-work-group-size"="1,8" {{.*}} } +// CHECK-AMDGPU-FLAG: attributes #[[SMALLER]] = { {{.*}} "amdgpu-flat-work-group-size"="1,5" {{.*}} } +// CHECK-AMDGPU-FLAG: attributes #[[DYN_NT]] = { {{.*}} "amdgpu-flat-work-group-size"="1,7" {{.*}} } +// CHECK-AMDGPU-FLAG: attributes #[[DYN_TL]] = { {{.*}} "amdgpu-flat-work-group-size"="1,6" {{.*}} } // CHECK-SPIRV-FLAG: attributes #[[ATTR1]] = { // CHECK-SPIRV-FLAG-NOT: omp_target_thread_limit @@ -64,3 +141,8 @@ void foo(int N) { // CHECK-SPIRV-FLAG: attributes #[[ATTR3]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="42" {{.*}} } // CHECK-SPIRV-FLAG: attributes #[[ATTR4]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="22" {{.*}} } // CHECK-SPIRV-FLAG: attributes #[[ATTR5]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="20" {{.*}} } +// CHECK-SPIRV-FLAG: attributes #[[SPLIT]] = { {{.*}} "omp_target_thread_limit"="19" {{.*}} } +// CHECK-SPIRV-FLAG: attributes #[[CLAMP]] = { {{.*}} "omp_target_thread_limit"="8" {{.*}} } +// CHECK-SPIRV-FLAG: attributes #[[SMALLER]] = { {{.*}} "omp_target_thread_limit"="5" {{.*}} } +// CHECK-SPIRV-FLAG: attributes #[[DYN_NT]] = { {{.*}} "omp_target_thread_limit"="7" {{.*}} } +// CHECK-SPIRV-FLAG: attributes #[[DYN_TL]] = { {{.*}} "omp_target_thread_limit"="6" {{.*}} } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
