Author: Robert Imschweiler Date: 2026-10-01T20:52:10+02:00 New Revision: 88156f8635bee29338aa5d2589318d9857b1236a
URL: https://github.com/llvm/llvm-project/commit/88156f8635bee29338aa5d2589318d9857b1236a DIFF: https://github.com/llvm/llvm-project/commit/88156f8635bee29338aa5d2589318d9857b1236a.diff LOG: [clang][OpenMP] Don't use fused dist schedule for teams loop emitted as distribute (#228129) Fix teams loop reductions lowered as 'distribute' lose their loop. Claude assisted with this patch. Added: clang/test/OpenMP/teams_generic_loop_reduction_distribute_codegen.cpp Modified: clang/lib/CodeGen/CGStmtOpenMP.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index b5e4b01109cb3..3ca9d7ee30d17 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -56,6 +56,15 @@ getEffectiveDirectiveKind(const OMPExecutableDirective &S); static bool canEmitGPUFusedDistSchedule(const CodeGenModule &CGM, const OMPLoopDirective &S, OpenMPDirectiveKind DKind) { + // 'teams loop' is always emitted as 'distribute', and 'target teams loop' + // only becomes 'distribute parallel for' if canBeParallelFor() holds. + // Without the inner worksharing loop, there is nothing that would schedule + // the iteration space if the outer distribute loop is omitted. + if (DKind == OMPD_teams_loop) + return false; + if (const auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(&S); + TTLD && !TTLD->canBeParallelFor()) + return false; // Reduction-only for now. Non-reduction cases might follow in the future, but // need more analysis for maximum profit. return CGM.getLangOpts().OpenMPIsTargetDevice && CGM.getTriple().isGPU() && diff --git a/clang/test/OpenMP/teams_generic_loop_reduction_distribute_codegen.cpp b/clang/test/OpenMP/teams_generic_loop_reduction_distribute_codegen.cpp new file mode 100644 index 0000000000000..9be8527fcdf20 --- /dev/null +++ b/clang/test/OpenMP/teams_generic_loop_reduction_distribute_codegen.cpp @@ -0,0 +1,63 @@ +// REQUIRES: amdgpu-registered-target + +// Check that a 'teams loop' with a reduction that is emitted as 'distribute' +// (rather than 'distribute parallel for') keeps its distribute loop on the +// GPU. The fused distribute schedule omits the outer distribute loop and relies +// on an inner worksharing loop, which only exists for '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 + +#pragma omp declare target +double getY(const double *y, int i) { return y[i]; } +#pragma omp end declare target + +// The call to getY prevents 'target teams loop' from being emitted as +// 'distribute parallel for'. +double target_teams_loop_call(const double *y, int N) { + double check = 0.0; +#pragma omp target teams loop reduction(+:check) map(to: y[0:N]) + for (int i = 0; i < N; i++) + check += getY(y, i); + return check; +} + +// 'teams loop' is always emitted as 'distribute'. +double teams_loop(const double *y, int N) { + double check = 0.0; +#pragma omp target map(to: y[0:N]) map(tofrom: check) +#pragma omp teams loop reduction(+:check) + for (int i = 0; i < N; i++) + check += y[i]; + return check; +} + +// Emitted as 'distribute parallel for', so the fused schedule is used. +double target_teams_loop_parallel_for(const double *y, int N) { + double check = 0.0; +#pragma omp target teams loop reduction(+:check) map(to: y[0:N]) + for (int i = 0; i < N; i++) + check += y[i]; + return check; +} + +// CHECK-LABEL: define internal void @{{.*}}target_teams_loop_call{{.*}}_l{{[0-9]+}}_omp_outlined( +// CHECK: call void @__kmpc_distribute_static_init_4( +// CHECK: omp.inner.for.body: +// CHECK: call noundef double @_Z4getYPKdi( +// CHECK: call void @__kmpc_distribute_static_fini( +// CHECK: call i32 @__kmpc_gpu_xteam_reduce_nowait( + +// CHECK-LABEL: define internal void @{{.*}}teams_loop{{.*}}_l{{[0-9]+}}_omp_outlined( +// CHECK: call void @__kmpc_distribute_static_init_4( +// CHECK: omp.inner.for.body: +// CHECK: call void @__kmpc_distribute_static_fini( +// CHECK: call i32 @__kmpc_gpu_xteam_reduce_nowait( + +// CHECK-LABEL: define internal void @{{.*}}target_teams_loop_parallel_for{{.*}}_l{{[0-9]+}}_omp_outlined( +// CHECK-NOT: call void @__kmpc_distribute_static_init +// CHECK: call void @__kmpc_parallel_60( +// CHECK-LABEL: define internal void @{{.*}}target_teams_loop_parallel_for{{.*}}_l{{[0-9]+}}_omp_outlined_omp_outlined( +// CHECK: call void @__kmpc_for_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 93, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
