[clang] 88156f8 - [clang][OpenMP] Don't use fused dist schedule for teams loop emitted as distribute (#228129)

[via cfe-commits]

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,