https://github.com/ro-i updated https://github.com/llvm/llvm-project/pull/228230
>From cca4917cf92bd5d3b6693e53975f3530c511d06b Mon Sep 17 00:00:00 2001 From: Robert Imschweiler <[email protected]> Date: Thu, 1 Oct 2026 15:01:47 -0500 Subject: [PATCH] [OpenMP][clang] Properly handle nested "teams loop" Treat a nested "teams loop" directive like a "teams distribute" since it's lowered as such. Claude assisted with this patch. --- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 5 +++ .../teams_generic_loop_nested_exec_mode.cpp | 37 +++++++++++++++++++ .../offloading/teams_loop_nested_in_target.c | 36 ++++++++++++++++++ offload/test/offloading/thread_limit.c | 4 +- 4 files changed, 81 insertions(+), 1 deletion(-) create mode 100644 clang/test/OpenMP/teams_generic_loop_nested_exec_mode.cpp create mode 100644 offload/test/offloading/teams_loop_nested_in_target.c diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index e7322bdd0deca9..5fd63b2fcb288d 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -527,6 +527,11 @@ static bool hasNestedSPMDDirective(ASTContext &Ctx, OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind(); switch (D.getDirectiveKind()) { case OMPD_target: + // A nested 'teams loop' is emitted as 'teams distribute' (see + // CodeGenFunction::EmitOMPTeamsGenericLoopDirective), which does not + // support SPMD mode. + if (DKind == OMPD_teams_loop) + return false; if (isOpenMPParallelDirective(DKind)) return true; if (DKind == OMPD_teams) { diff --git a/clang/test/OpenMP/teams_generic_loop_nested_exec_mode.cpp b/clang/test/OpenMP/teams_generic_loop_nested_exec_mode.cpp new file mode 100644 index 00000000000000..4790b201467fec --- /dev/null +++ b/clang/test/OpenMP/teams_generic_loop_nested_exec_mode.cpp @@ -0,0 +1,37 @@ +// REQUIRES: amdgpu-registered-target + +// A 'teams loop' nested in a 'target' region is emitted as 'teams distribute', +// so the kernel must use generic mode, unlike a nested 'teams distribute +// parallel for'. + +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgpu-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgpu-amd-amdhsa -fopenmp-targets=amdgpu-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s + +// expected-no-diagnostics + +// The exec mode is the third field of the configuration environment: generic +// is 1, SPMD is 2. +// CHECK: @{{.*}}teams_loop{{.*}}_kernel_environment = {{.*}} { %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 1, +// CHECK: @{{.*}}teams_loop_no_reduction{{.*}}_kernel_environment = {{.*}} { %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 1, +// CHECK: @{{.*}}teams_distribute_parallel_for{{.*}}_kernel_environment = {{.*}} { %struct.ConfigurationEnvironmentTy { i8 0, i8 1, i8 2, + +void teams_loop(double *y, double &c, int N) { +#pragma omp target map(to: y[0:N]) map(tofrom: c) +#pragma omp teams loop reduction(+:c) + for (int i = 0; i < N; i++) + c += y[i]; +} + +void teams_loop_no_reduction(int *z, int N) { +#pragma omp target map(tofrom: z[0:N]) +#pragma omp teams loop + for (int i = 0; i < N; i++) + z[i] += 1; +} + +void teams_distribute_parallel_for(int *z, int N) { +#pragma omp target map(tofrom: z[0:N]) +#pragma omp teams distribute parallel for + for (int i = 0; i < N; i++) + z[i] += 1; +} diff --git a/offload/test/offloading/teams_loop_nested_in_target.c b/offload/test/offloading/teams_loop_nested_in_target.c new file mode 100644 index 00000000000000..3bc08cc6f0f506 --- /dev/null +++ b/offload/test/offloading/teams_loop_nested_in_target.c @@ -0,0 +1,36 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// RUN: %libomptarget-compileopt-run-and-check-generic + +// A 'teams loop' nested in a 'target' region is emitted as 'teams distribute' +// and must not be executed in SPMD mode. + +#include <stdio.h> + +#define N (1 << 20) + +int Y[N]; + +int main() { + for (int I = 0; I < N; ++I) + Y[I] = 1; + + long Sum = 0; +#pragma omp target map(to : Y) map(tofrom : Sum) +#pragma omp teams loop reduction(+ : Sum) + for (int I = 0; I < N; ++I) + Sum += Y[I]; + // CHECK: reduction: 1048576 + printf("reduction: %ld\n", Sum); + +#pragma omp target map(tofrom : Y) +#pragma omp teams loop + for (int I = 0; I < N; ++I) + Y[I] += 1; + Sum = 0; + for (int I = 0; I < N; ++I) + Sum += Y[I]; + // CHECK: increments: 2097152 + printf("increments: %ld\n", Sum); + + return 0; +} diff --git a/offload/test/offloading/thread_limit.c b/offload/test/offloading/thread_limit.c index 34c952175f8b64..4fc1c103d993df 100644 --- a/offload/test/offloading/thread_limit.c +++ b/offload/test/offloading/thread_limit.c @@ -12,7 +12,9 @@ int main() { int n = 1 << 20; int th = 12; int te = n / th; -// DEFAULT: 12 (MaxFlatWorkGroupSize: +// The nested 'teams loop' is executed in generic mode, which adds a warp for +// the main thread (64 or 32 threads). +// DEFAULT: {{76|44}} (MaxFlatWorkGroupSize: #pragma omp target #pragma omp teams loop num_teams(te), thread_limit(th) for (int i = 0; i < n; i++) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
