[llvm-branch-commits] [clang] [CIR][OpenMP] Add support for omp teams (PR #231982)
Jan Leyonberg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Oct 19 13:51:16 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/231982
>From 4db562ac8039a0e3fbfc4ca5e5edb90371e734f3 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Fri, 2 Oct 2026 13:07:09 -0400
Subject: [PATCH] [CIR][OpenMP] Add support for omp teams
This patch adds support for the OpenMP 'teams' directive, both standalone
and as the combined `target teams` construct. isCombinableLeaf now also
recognizes teams, so an explicitly-nested `target` + `teams` pair gets
the same omp.combined marking as the combined case.
Assisted-by: Cursor / Claude Sonnet 5 High
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 105 ++++++++++-
.../CIR/CodeGenOpenMP/not-yet-implemented.c | 16 +-
clang/test/CIR/CodeGenOpenMP/target-teams.c | 176 ++++++++++++++++++
clang/test/CIR/CodeGenOpenMP/teams.c | 108 +++++++++++
4 files changed, 397 insertions(+), 8 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-teams.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/teams.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 9dffa8f2846ee..6eddbd238ac54 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -68,6 +68,7 @@ static bool isCombinableLeaf(llvm::omp::Directive dir) {
switch (dir) {
case llvm::omp::OMPD_target:
case llvm::omp::OMPD_parallel:
+ case llvm::omp::OMPD_teams:
case llvm::omp::OMPD_for:
return true;
default:
@@ -189,6 +190,68 @@ static mlir::LogicalResult emitParallel(
emitBody);
}
+/// Evaluates the clauses allowed on an omp.teams leaf; none are supported
+/// yet, so every eligible clause is reported as NYI.
+static mlir::LogicalResult
+emitTeamsClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
+ CIRGenBuilderTy &builder, mlir::Location loc,
+ llvm::ArrayRef<const OMPClause *> clauses,
+ mlir::omp::TeamsOperands &clauseOps) {
+ OpenMPClauseEmitter ce(cgf, cgm, builder, loc, clauses);
+ return ce.emitNYI</*supported=*/>(
+ /*nyi=*/OpenMPNYIClauseList<
+ OMPAllocateClause, OMPDefaultClause, OMPDynGroupprivateClause,
+ OMPFirstprivateClause, OMPIfClause, OMPNumTeamsClause,
+ OMPPrivateClause, OMPReductionClause, OMPSharedClause,
+ OMPThreadLimitClause, OMPXAttributeClause>{},
+ llvm::omp::Directive::OMPD_teams);
+}
+
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitTeamsOp(CIRGenFunction &cgf, const DirectiveTy &s,
+ const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator item, mlir::Location begin,
+ mlir::Location end, const mlir::omp::TeamsOperands &clauseOps,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+
+ auto teamsOp = mlir::omp::TeamsOp::create(builder, begin, clauseOps);
+ if (!omp::isLastItemInQueue(item, queue) ||
+ hasCombinableNestedLeaf(static_cast<const OMPExecutableDirective &>(s)))
+ teamsOp.setCombined(true);
+
+ mlir::Block &block = teamsOp.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(&block);
+
+ CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+
+ mlir::LogicalResult res = emitBody();
+ mlir::omp::TerminatorOp::create(builder, end);
+ return res;
+}
+
+/// Checks, emits clauses for, and emits an omp.teams for the `teams`
+/// directive at \p teamsItem, deferring to \p emitBody for the region's
+/// contents.
+template <typename DirectiveTy>
+static mlir::LogicalResult emitTeams(
+ CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator teamsItem, mlir::Location begin,
+ mlir::Location end, llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, teamsItem)))
+ return mlir::failure();
+
+ mlir::omp::TeamsOperands teamsOps;
+ if (mlir::failed(emitTeamsClauses(cgf, cgf.getCIRGenModule(),
+ cgf.getBuilder(), begin, teamsItem->clauses,
+ teamsOps)))
+ return mlir::failure();
+
+ return emitTeamsOp(cgf, s, queue, teamsItem, begin, end, teamsOps, emitBody);
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
@@ -792,8 +855,19 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
}
mlir::LogicalResult
CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPTeamsDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator item = queue.begin();
+
+ return emitTeams(
+ *this, s, queue, item, begin, end, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_teams);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPCancellationPointDirective(
const OMPCancellationPointDirective &s) {
@@ -1041,9 +1115,30 @@ mlir::LogicalResult CIRGenFunction::emitOMPTeamsGenericLoopDirective(
}
mlir::LogicalResult
CIRGenFunction::emitOMPTargetTeamsDirective(const OMPTargetTeamsDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPTargetTeamsDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator targetItem = queue.begin();
+ assert(targetItem->id == llvm::omp::OMPD_target &&
+ "expected 'target' to be the outermost directive");
+
+ return emitTarget(
+ *this, s, queue, targetItem, begin, end, [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator teamsItem = std::next(targetItem);
+ assert(teamsItem != queue.end() &&
+ teamsItem->id == llvm::omp::OMPD_teams &&
+ "expected a 'teams' directive nested in 'target'");
+
+ return emitTeams(*this, s, queue, teamsItem, begin, end,
+ [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs =
+ s.getCapturedStmt(llvm::omp::OMPD_teams);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetTeamsDistributeDirective(
const OMPTargetTeamsDistributeDirective &s) {
diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
index c4a97e2b76d05..ef28330b2a478 100644
--- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
+++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
@@ -10,8 +10,8 @@ void do_things() {
{}
int i;
- // A leaf that reports a not-yet-implemented clause emits no op at all, rather
- // than one that silently ignores the clause.
+ // A directive that reports a not-yet-implemented clause emits no op at all,
+ // rather than one that silently ignores the clause.
int a, b;
// expected-error at +2{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'shared' clause}}
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'firstprivate' clause}}
@@ -19,7 +19,7 @@ void do_things() {
{}
// A clause routed through construct decomposition but not yet emittable must
- // still be diagnosed by the leaf emitter's NYI handling.
+ // still be diagnosed by the directive emitter's NYI handling.
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP TARGET 'private' clause}}
#pragma omp target private(i)
{}
@@ -29,4 +29,14 @@ void do_things() {
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP synthesized 'firstprivate' clause from construct decomposition}}
#pragma omp target
{ i = 1; }
+
+ // The teams construct supports no clauses yet, so every `teams` clause is
+ // reported as not-yet-implemented.
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP TEAMS 'num_teams' clause}}
+#pragma omp teams num_teams(4)
+ {}
+
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP TEAMS 'thread_limit' clause}}
+#pragma omp teams thread_limit(8)
+ {}
}
diff --git a/clang/test/CIR/CodeGenOpenMP/target-teams.c b/clang/test/CIR/CodeGenOpenMP/target-teams.c
new file mode 100644
index 0000000000000..fdeb4f8ef8286
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-teams.c
@@ -0,0 +1,176 @@
+// REQUIRES: amdgpu-registered-target
+
+// Host compilation (x86 host, AMDGPU offload target).
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -emit-cir -fclangir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR-HOST
+
+// Device compilation (AMDGPU).
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -emit-cir -fclangir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR-DEVICE
+
+// Host-side lowering to LLVM IR: identical on both paths, so LLVM-HOST
+// checks both. The only difference is where the second function's helper
+// functions end up relative to the first function's, checked below under
+// LLVM-HOST-CIR/LLVM-HOST-OGCG.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM-HOST,LLVM-HOST-CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM-HOST,LLVM-HOST-OGCG
+
+// Device-side lowering to LLVM IR: compiled in two steps to mirror the
+// offloading driver.
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM-DEVICE,LLVM-DEVICE-CIR
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM-DEVICE,LLVM-DEVICE-OGCG
+
+void use(int);
+
+// Explicit `target` + `teams` nesting lowers the same as the combined
+// `target teams` spelling below, so both get omp.combined.
+void explicit_target_teams(int x) {
+ // CIR-HOST: cir.func{{.*}}@explicit_target_teams
+ // CIR-HOST: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(tofrom) {{.*}} name("x")
+ // CIR-HOST: omp.target kernel_type(generic) map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) {
+ // CIR-HOST: omp.teams {
+ // CIR-HOST: %[[LOAD:.*]] = cir.load align(4) %[[ARG]]
+ // CIR-HOST: cir.call @use(%[[LOAD]])
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: } {omp.combined}
+
+ // CIR-DEVICE: cir.func{{.*}}@explicit_target_teams
+ // CIR-DEVICE: omp.target kernel_type(generic) {{.*}} {
+ // CIR-DEVICE: omp.teams {
+ // CIR-DEVICE: cir.call @use
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: }
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: } {omp.combined}
+#pragma omp target map(tofrom : x)
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+// Combined `target teams` lowers the same as the explicit nesting above,
+// with omp.combined on the omp.target.
+void target_teams(int x) {
+ // CIR-HOST: cir.func{{.*}}@target_teams
+ // CIR-HOST: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(tofrom) {{.*}} name("x")
+ // CIR-HOST: omp.target kernel_type(generic) map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) {
+ // CIR-HOST: omp.teams {
+ // CIR-HOST: %[[LOAD:.*]] = cir.load align(4) %[[ARG]]
+ // CIR-HOST: cir.call @use(%[[LOAD]])
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: } {omp.combined}
+
+ // CIR-DEVICE: cir.func{{.*}}@target_teams
+ // CIR-DEVICE: omp.target kernel_type(generic) {{.*}} {
+ // CIR-DEVICE: omp.teams {
+ // CIR-DEVICE: cir.call @use
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: }
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: } {omp.combined}
+#pragma omp target teams map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+/// LLVM-HOST-LABEL: define {{.*}} void @explicit_target_teams(
+// LLVM-HOST-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM-HOST: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM-HOST: store i32 {{.*}}, ptr %[[X_ADDR]], align 4
+// LLVM-HOST: call i32 @__tgt_target_kernel(
+// LLVM-HOST: omp_offload.failed:
+// LLVM-HOST: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// CIR groups all its helper functions together up front, since the
+// MLIR-to-LLVM translation emits them during the lowering pass itself;
+// OGCG's direct AST walk emits each one in source order instead,
+// interleaved with other constructs. The checks are the same either way,
+// just placed where each path actually emits them (compare the
+// LLVM-HOST-OGCG copy of this block further down).
+// LLVM-HOST-CIR-LABEL: define {{.*}} void @target_teams(
+// LLVM-HOST-CIR-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM-HOST-CIR: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM-HOST-CIR: store i32 {{.*}}, ptr %[[X_ADDR]], align 4
+// LLVM-HOST-CIR: call i32 @__tgt_target_kernel(
+// LLVM-HOST-CIR: omp_offload.failed:
+// LLVM-HOST-CIR: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}(
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}, ptr {{.*}})
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}(
+// LLVM-HOST: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM-HOST: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM-HOST: call void @use(i32 {{.*}} %[[V]])
+// LLVM-HOST: ret void
+
+// OGCG's copy of target_teams's host wrapper: identical content to the
+// LLVM-HOST-CIR block above, just emitted after explicit_target_teams's
+// outlined functions instead of before them.
+// LLVM-HOST-OGCG-LABEL: define {{.*}} void @target_teams(
+// LLVM-HOST-OGCG-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM-HOST-OGCG: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM-HOST-OGCG: store i32 {{.*}}, ptr %[[X_ADDR]], align 4
+// LLVM-HOST-OGCG: call i32 @__tgt_target_kernel(
+// LLVM-HOST-OGCG: omp_offload.failed:
+// LLVM-HOST-OGCG: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_teams_l{{.*}}, ptr {{.*}})
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(
+// LLVM-HOST: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM-HOST: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM-HOST: call void @use(i32 {{.*}} %[[V]])
+// LLVM-HOST: ret void
+
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_explicit_target_teams_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}_omp_outlined(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par(
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM-DEVICE: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM-DEVICE: call void @use(i32 {{.*}} %[[V]])
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_teams_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}_omp_outlined(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par(
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM-DEVICE: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM-DEVICE: call void @use(i32 {{.*}} %[[V]])
+// LLVM-DEVICE: ret void
diff --git a/clang/test/CIR/CodeGenOpenMP/teams.c b/clang/test/CIR/CodeGenOpenMP/teams.c
new file mode 100644
index 0000000000000..088610c44c50f
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/teams.c
@@ -0,0 +1,108 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s --check-prefix=CIR
+
+// Lowering to LLVM IR: mostly identical between CIR and OGCG; they just name
+// the outlined functions differently.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-OGCG
+
+void use(int);
+
+// A standalone `teams` directive lowers to an omp.teams op whose region holds
+// the associated structured block.
+void host_teams(int x) {
+ // CIR: cir.func{{.*}}@host_teams
+ // CIR: omp.teams {
+ // CIR: cir.call @use
+ // CIR: omp.terminator
+ // CIR: }
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+// LLVM-LABEL: define {{.*}} void @host_teams(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @host_teams{{.*}}, ptr {{.*}})
+
+// LLVM-CIR-LABEL: define internal void @host_teams..omp_par(
+// LLVM-OGCG-LABEL: define internal void @host_teams.omp_outlined(
+// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+void teams_nested_parallel_for(void) {
+ // CIR: cir.func{{.*}}@teams_nested_parallel_for
+ // CIR: omp.teams {
+ // CIR: omp.parallel {
+ // CIR: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) {
+ // CIR-NEXT: omp.loop_nest
+ // CIR: cir.call @use
+ // CIR: omp.yield
+ // CIR: }
+ // CIR: omp.terminator
+ // CIR: } {omp.combined}
+ // CIR: omp.terminator
+ // CIR: } {omp.combined}
+#pragma omp teams
+ {
+#pragma omp parallel for
+ for (int i = 0; i < 10; i++) {
+ use(i);
+ }
+ }
+}
+
+// The nested `parallel for` forks its own threads via __kmpc_fork_call from
+// inside the teams-outlined function, so the wsloop's init/fini/barrier
+// live in a second, further-outlined function.
+// LLVM-LABEL: define {{.*}} void @teams_nested_parallel_for(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @teams_nested_parallel_for{{.*}})
+
+// LLVM-CIR-LABEL: define internal void @teams_nested_parallel_for..omp_par.1(
+// LLVM-OGCG-LABEL: define internal void @teams_nested_parallel_for.omp_outlined(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @teams_nested_parallel_for{{.*}})
+
+// OGCG skips the trailing barrier after the worksharing loop (the
+// implicit join at the end of the parallel region already synchronizes);
+// CIR doesn't yet apply that optimization, so it emits one either way.
+// LLVM-CIR-LABEL: define internal void @teams_nested_parallel_for..omp_par(
+// LLVM-CIR: %[[TID:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-CIR-NEXT: call void @__kmpc_for_static_init_4u(
+// LLVM-CIR-SAME: ptr {{.*}}, i32 %[[TID]], i32 34,
+// LLVM-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.stride,
+// LLVM-CIR-SAME: i32 1, i32 0)
+// LLVM-CIR: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM-CIR: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-CIR-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM-CIR: omp_loop.exit:
+// LLVM-CIR-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID]])
+// LLVM-CIR: call void @__kmpc_barrier(ptr {{.*}}, i32 %{{.*}})
+// LLVM-CIR: omp_loop.body:
+// LLVM-CIR: call void @use(i32
+
+// LLVM-OGCG-LABEL: define internal void @teams_nested_parallel_for.omp_outlined.omp_outlined(
+// LLVM-OGCG: store i32 0, ptr %.omp.lb, align 4
+// LLVM-OGCG-NEXT: store i32 9, ptr %.omp.ub, align 4
+// LLVM-OGCG-NEXT: store i32 1, ptr %.omp.stride, align 4
+// LLVM-OGCG-NEXT: store i32 0, ptr %.omp.is_last, align 4
+// LLVM-OGCG: %[[TID2:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-OGCG-SAME: ptr {{.*}}, i32 %[[TID2]], i32 34,
+// LLVM-OGCG-SAME: ptr %.omp.is_last, ptr %.omp.lb, ptr %.omp.ub, ptr %.omp.stride,
+// LLVM-OGCG-SAME: i32 1, i32 1)
+// LLVM-OGCG: %[[UB:.*]] = load i32, ptr %.omp.ub, align 4
+// LLVM-OGCG-NEXT: %[[CMP:.*]] = icmp sgt i32 %[[UB]], 9
+// LLVM-OGCG-NEXT: br i1 %[[CMP]], label %cond.true, label %cond.false
+// LLVM-OGCG: cond.end:
+// LLVM-OGCG: %[[COND:.*]] = phi i32 [ 9, %cond.true ], [ %{{.*}}, %cond.false ]
+// LLVM-OGCG-NEXT: store i32 %[[COND]], ptr %.omp.ub, align 4
+// LLVM-OGCG: omp.inner.for.cond:
+// LLVM-OGCG: %[[CMP1:.*]] = icmp sle i32 %{{.*}}, %{{.*}}
+// LLVM-OGCG-NEXT: br i1 %[[CMP1]], label %omp.inner.for.body, label %omp.inner.for.end
+// LLVM-OGCG: omp.inner.for.body:
+// LLVM-OGCG: call void @use(i32
+// LLVM-OGCG: omp.loop.exit:
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID2]])
+// LLVM-OGCG-NEXT: ret void
More information about the llvm-branch-commits
mailing list