[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
Fri Oct 16 08:06:26 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/231982
>From 32225e83eba4fcbaf6da23c15f75713c45635dd6 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 ++-
.../CodeGenOpenMP/target-teams-llvm-device.c | 86 ++++++++++++++
.../CodeGenOpenMP/target-teams-llvm-host.c | 91 +++++++++++++++
clang/test/CIR/CodeGenOpenMP/target-teams.c | 69 ++++++++++++
clang/test/CIR/CodeGenOpenMP/teams-llvm.c | 65 +++++++++++
clang/test/CIR/CodeGenOpenMP/teams.c | 39 +++++++
7 files changed, 463 insertions(+), 8 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-teams-llvm-device.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-teams-llvm-host.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-teams.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/teams-llvm.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 9dffa8f2846ee0..6eddbd238ac544 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 c4a97e2b76d05b..ef28330b2a478f 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-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-teams-llvm-device.c
new file mode 100644
index 00000000000000..4c8ce1fd019d27
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-teams-llvm-device.c
@@ -0,0 +1,86 @@
+// Two-step host-BC -> device pipeline that mirrors the offloading driver.
+//
+// Step 1: Host compilation to bitcode (provides offload entry info to device pass).
+// 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
+//
+// Step 2: Device compilation using 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-prefix=LLVM
+// 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-prefix=OGCG
+
+void use(int);
+
+void explicit_target_teams(int x) {
+#pragma omp target map(tofrom : x)
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+void target_teams(int x) {
+#pragma omp target teams map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_explicit_target_teams_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM: call void @__kmpc_target_deinit()
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_teams_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// LLVM: call void @__kmpc_target_deinit()
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_explicit_target_teams_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}_omp_outlined(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// OGCG: call void @__kmpc_target_deinit()
+// OGCG: ret void
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}_omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_teams_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}_omp_outlined(ptr {{.*}}, ptr {{.*}}, ptr {{.*}})
+// OGCG: call void @__kmpc_target_deinit()
+// OGCG: ret void
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}_omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
diff --git a/clang/test/CIR/CodeGenOpenMP/target-teams-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-teams-llvm-host.c
new file mode 100644
index 00000000000000..768b82aab9e1c4
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-teams-llvm-host.c
@@ -0,0 +1,91 @@
+// 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-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void use(int);
+
+void explicit_target_teams(int x) {
+#pragma omp target map(tofrom : x)
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+void target_teams(int x) {
+#pragma omp target teams map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+// Host wrappers
+//
+// LLVM-LABEL: define {{.*}} void @explicit_target_teams(
+// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// LLVM-LABEL: define {{.*}} void @target_teams(
+// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// Outlined target functions: both the explicit `target`+`teams` nesting and
+// the combined `target teams` spelling lower identically, forking an
+// omp.teams region via __kmpc_fork_teams.
+//
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par, ptr {{.*}})
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}..omp_par(
+// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par, ptr {{.*}})
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}..omp_par(
+// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// OGCG interleaves host wrapper and outlined function per target region.
+//
+// OGCG-LABEL: define {{.*}} void @explicit_target_teams(i32
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_explicit_target_teams_l
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}.omp_outlined, ptr {{.*}})
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_l{{.*}}.omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
+
+// OGCG-LABEL: define {{.*}} void @target_teams(i32
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_target_teams_l
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_teams_l{{.*}}.omp_outlined, ptr {{.*}})
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_l{{.*}}.omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
diff --git a/clang/test/CIR/CodeGenOpenMP/target-teams.c b/clang/test/CIR/CodeGenOpenMP/target-teams.c
new file mode 100644
index 00000000000000..3907947d12e153
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-teams.c
@@ -0,0 +1,69 @@
+// 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
+
+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);
+ }
+}
diff --git a/clang/test/CIR/CodeGenOpenMP/teams-llvm.c b/clang/test/CIR/CodeGenOpenMP/teams-llvm.c
new file mode 100644
index 00000000000000..a7221765f8bd6c
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/teams-llvm.c
@@ -0,0 +1,65 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void use(int);
+
+void host_teams(int x) {
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+void teams_nested_parallel_for(void) {
+#pragma omp teams
+ {
+#pragma omp parallel for
+ for (int i = 0; i < 10; i++) {
+ use(i);
+ }
+ }
+}
+
+// LLVM-LABEL: define {{.*}} void @host_teams(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @host_teams..omp_par, ptr {{.*}})
+
+// LLVM-LABEL: define internal void @host_teams..omp_par(
+// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// The nested `parallel for` forks its own threads via __kmpc_fork_call from
+// inside the teams-outlined function, so the wsloop's static-init/fini and
+// 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..omp_par.1)
+
+// LLVM-LABEL: define internal void @teams_nested_parallel_for..omp_par.1(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @teams_nested_parallel_for..omp_par)
+
+// LLVM-LABEL: define internal void @teams_nested_parallel_for..omp_par(
+// LLVM: call void @__kmpc_for_static_init_4u(
+// LLVM: call void @__kmpc_for_static_fini(
+// LLVM: call void @__kmpc_barrier(
+// LLVM: call void @use(i32
+
+// OGCG-LABEL: define {{.*}} void @host_teams(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 1, ptr @host_teams.omp_outlined, ptr {{.*}})
+
+// OGCG-LABEL: define internal void @host_teams.omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
+
+// OGCG-LABEL: define {{.*}} void @teams_nested_parallel_for(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @teams_nested_parallel_for.omp_outlined)
+
+// OGCG-LABEL: define internal void @teams_nested_parallel_for.omp_outlined(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @teams_nested_parallel_for.omp_outlined.omp_outlined)
+
+// OGCG-LABEL: define internal void @teams_nested_parallel_for.omp_outlined.omp_outlined(
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG: call void @use(i32
+// OGCG: call void @__kmpc_for_static_fini(
diff --git a/clang/test/CIR/CodeGenOpenMP/teams.c b/clang/test/CIR/CodeGenOpenMP/teams.c
new file mode 100644
index 00000000000000..d8b4d1436d0a79
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/teams.c
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
+
+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) {
+ // CHECK: cir.func{{.*}}@host_teams
+ // CHECK: omp.teams {
+ // CHECK: cir.call @use
+ // CHECK: omp.terminator
+ // CHECK: }
+#pragma omp teams
+ {
+ use(x);
+ }
+}
+
+void teams_nested_parallel_for(void) {
+ // CHECK: cir.func{{.*}}@teams_nested_parallel_for
+ // CHECK: omp.teams {
+ // CHECK: omp.parallel {
+ // CHECK: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest
+ // CHECK: cir.call @use
+ // CHECK: omp.yield
+ // CHECK: }
+ // CHECK: omp.terminator
+ // CHECK: } {omp.combined}
+ // CHECK: omp.terminator
+ // CHECK: } {omp.combined}
+#pragma omp teams
+ {
+#pragma omp parallel for
+ for (int i = 0; i < 10; i++) {
+ use(i);
+ }
+ }
+}
More information about the llvm-branch-commits
mailing list