[llvm-branch-commits] [clang] [CIR][OpenMP] Add support for 'distribute parallel for' with 'target … (PR #234258)
Jan Leyonberg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Oct 21 09:02:49 PDT 2026
https://github.com/jsjodin created https://github.com/llvm/llvm-project/pull/234258
…teams'
Adds support for handling the full '[target] teams distribute parallel for'. Introduces setting the omp.composite attibute for the 'distribute parallel for' composite construct.
Assisted-by: Cursor / Claude Sonnet 5 High
>From dce02e4d63224b4ce008145c4ad9798a2903d8e9 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 7 Oct 2026 14:43:58 -0400
Subject: [PATCH] [CIR][OpenMP] Add support for 'distribute parallel for' with
'target teams'
Adds support for handling the full '[target] teams distribute parallel for'.
Introduces setting the omp.composite attibute for the 'distribute parallel for'
composite construct.
Assisted-by: Cursor / Claude Sonnet 5 High
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 334 +++++++++--
.../CodeGenOpenMP/distribute-parallel-for.c | 213 +++++++
.../CIR/CodeGenOpenMP/not-yet-implemented.c | 13 +
.../target-teams-distribute-parallel-for.c | 556 ++++++++++++++++++
4 files changed, 1078 insertions(+), 38 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenMP/distribute-parallel-for.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-teams-distribute-parallel-for.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 61558a4b0b5200..406b5674b6a630 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -99,16 +99,31 @@ findNestedOMPLoopDirective(const OMPExecutableDirective &s) {
return findNestedOMPLoopDirective(*nested);
}
-/// Returns true if \p s is a target SPMD construct: a combined
-/// `target parallel for`, or an explicitly nested `target` whose body is a
-/// parallel directive.
+/// Returns true if \p s is a target SPMD construct.
static bool isTargetSPMD(const OMPExecutableDirective &s) {
switch (s.getDirectiveKind()) {
case llvm::omp::OMPD_target_parallel_for:
+ case llvm::omp::OMPD_target_teams_distribute_parallel_for:
return true;
case llvm::omp::OMPD_target: {
const OMPExecutableDirective *nested = getSingleNestedOMPDirective(s);
- return nested && isOpenMPParallelDirective(nested->getDirectiveKind());
+ if (!nested)
+ return false;
+ if (isOpenMPParallelDirective(nested->getDirectiveKind()))
+ return true;
+ if (nested->getDirectiveKind() != llvm::omp::OMPD_teams)
+ return false;
+ const OMPExecutableDirective *distribute =
+ getSingleNestedOMPDirective(*nested);
+ return distribute && distribute->getDirectiveKind() ==
+ llvm::omp::OMPD_distribute_parallel_for;
+ }
+ case llvm::omp::OMPD_target_teams: {
+ // `target teams` (combined) with a separately, explicitly nested
+ // `distribute parallel for` (composite, split off).
+ const OMPExecutableDirective *distribute = getSingleNestedOMPDirective(s);
+ return distribute && distribute->getDirectiveKind() ==
+ llvm::omp::OMPD_distribute_parallel_for;
}
default:
return false;
@@ -161,27 +176,43 @@ emitParallelClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
llvm::omp::Directive::OMPD_parallel);
}
+/// Emits an omp.parallel op. When \p isComposite is set, this omp.parallel is
+/// a leaf of a composite construct, e.g. the `parallel` in `distribute
+/// parallel for`, and gets the `composite` attribute set. \p counterMolds is
+/// only used by composite callers. Used to create privatized induction
+/// variable(s) that can be used inside the loop nest.
template <typename DirectiveTy>
-static mlir::LogicalResult
-emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
- const omp::ConstructQueue &queue,
- omp::ConstructQueue::const_iterator item, mlir::Location begin,
- mlir::Location end, const mlir::omp::ParallelOperands &clauseOps,
- llvm::function_ref<mlir::LogicalResult()> emitBody) {
+static mlir::LogicalResult emitParallelOp(
+ CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator item, mlir::Location begin,
+ mlir::Location end, const mlir::omp::ParallelOperands &clauseOps,
+ llvm::function_ref<mlir::LogicalResult()> emitBody,
+ bool isComposite = false,
+ llvm::ArrayRef<std::pair<const VarDecl *, Address>> counterMolds = {}) {
CIRGenBuilderTy &builder = cgf.getBuilder();
CIRGenModule &cgm = cgf.getCIRGenModule();
auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
- if (!omp::isLastItemInQueue(item, queue) ||
- hasCombinableNestedLeaf(static_cast<const OMPExecutableDirective &>(s)))
- parallelOp.setCombined(true);
mlir::Block &block = parallelOp.getRegion().emplaceBlock();
+ for (auto &counterMold : counterMolds)
+ block.addArgument(counterMold.second.getPointer().getType(), begin);
+
mlir::OpBuilder::InsertionGuard guard(builder);
builder.setInsertionPointToEnd(&block);
CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+ // Bind each composite loop counter to its private block argument.
+ for (auto [idx, counterMold] : llvm::enumerate(counterMolds)) {
+ const VarDecl *vd = counterMold.first;
+ Address moldAddr = counterMold.second;
+ Address privAddr(block.getArgument(idx), moldAddr.getElementType(),
+ moldAddr.getAlignment());
+ cgf.replaceAddrOfLocalVar(vd, privAddr);
+ cgf.symbolTable.insert(vd, privAddr.getPointer());
+ }
+
// hasCancel() reports a nested cancel/cancellation point directive in the
// body, not a clause on `parallel`, so it can't be an emitParallelClauses()
// NYI check.
@@ -195,6 +226,19 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
mlir::LogicalResult res = emitBody();
mlir::omp::TerminatorOp::create(builder, end);
+
+ // omp.composite/omp.combined describe the nested leaves emitted by the body
+ // above, so they are only attached once that succeeded: a body that bailed
+ // out leaves an empty region, which omp.composite's verifier rejects with
+ // an error that would mask the original diagnostic.
+ if (res.succeeded()) {
+ if (isComposite)
+ parallelOp.setComposite(true);
+ else if (!omp::isLastItemInQueue(item, queue) ||
+ hasCombinableNestedLeaf(
+ static_cast<const OMPExecutableDirective &>(s)))
+ parallelOp.setCombined(true);
+ }
return res;
}
@@ -452,6 +496,37 @@ emitHostEvalLoopBounds(CIRGenFunction &cgf, const OMPLoopDirective &s) {
return hev;
}
+/// Computes an OMPLoopDirective's normalized `[0, tripCount)` bounds, as
+/// builtin-integer values ready for an omp.loop_nest.
+///
+/// If \p consumingHostEval is true, the enclosing omp.target already
+/// host_eval'd this loop's trip count, so that value is reused here instead
+/// of re-emitting the trip-count expression. Otherwise, the caller must have
+/// already run preinits (see emitPreinits) before calling this function.
+static void emitOMPLoopNestBounds(CIRGenFunction &cgf,
+ const OMPLoopDirective &s,
+ bool consumingHostEval, mlir::Value &zero,
+ mlir::Value &tripCount, mlir::Value &one) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+
+ if (consumingHostEval) {
+ // Consume the host_eval bounds forwarded by the enclosing omp.target.
+ cgf.ompHostEvalBounds->applied = true;
+ zero = cgf.ompHostEvalBounds->zero;
+ tripCount = cgf.ompHostEvalBounds->tripCount;
+ one = cgf.ompHostEvalBounds->one;
+ } else {
+ mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
+ auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
+ zero = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 0));
+ one = builder.createBuiltinIntCast(loc,
+ builder.getConstInt(loc, cirIntType, 1));
+ tripCount = builder.createBuiltinIntCast(loc, tripCountCir);
+ }
+}
+
/// Lowers an OMPLoopDirective to an omp.loop_nest wrapped by an \p OpTy
/// (omp.wsloop or omp.distribute).
template <typename OpTy, typename OperandsTy>
@@ -492,21 +567,7 @@ static mlir::LogicalResult emitOMPLoop(CIRGenFunction &cgf,
}
mlir::Value zero, tripCount, one;
- if (consumingHostEval) {
- // Consume the host_eval bounds forwarded by the enclosing omp.target.
- cgf.ompHostEvalBounds->applied = true;
- zero = cgf.ompHostEvalBounds->zero;
- tripCount = cgf.ompHostEvalBounds->tripCount;
- one = cgf.ompHostEvalBounds->one;
- } else {
- mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
- auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
- zero = builder.createBuiltinIntCast(
- loc, builder.getConstInt(loc, cirIntType, 0));
- one = builder.createBuiltinIntCast(loc,
- builder.getConstInt(loc, cirIntType, 1));
- tripCount = builder.createBuiltinIntCast(loc, tripCountCir);
- }
+ emitOMPLoopNestBounds(cgf, s, consumingHostEval, zero, tripCount, one);
auto loopOp = OpTy::create(builder, loc, clauseOps);
mlir::Block *innerBlock = new mlir::Block();
@@ -585,6 +646,160 @@ emitOMPDistributeLoop(CIRGenFunction &cgf, const OMPLoopDirective &s,
return emitOMPLoop<mlir::omp::DistributeOp>(cgf, s, clauseOps);
}
+/// Lowers the `distribute parallel for` composite construct's loop leaves
+/// into the stacked omp.distribute + omp.wsloop sharing a single
+/// omp.loop_nest, with both wrappers marked composite. This runs inside the
+/// composite omp.parallel created by the caller (emitOMPDistributeParallelFor
+/// below), which privatizes the shared loop counter(s) on itself (matching
+/// Flang -- see distribute-parallel-do.f90): the two loop wrappers here stay
+/// bare, consuming the same private block argument(s) already bound by the
+/// enclosing omp.parallel.
+static mlir::LogicalResult emitOMPDistributeParallelForLoop(
+ CIRGenFunction &cgf, const OMPLoopDirective &s,
+ const mlir::omp::DistributeOperands &distributeClauseOps,
+ const mlir::omp::WsloopOperands &wsloopClauseOps) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+
+ const CapturedStmt *capturedStmt = s.getInnermostCapturedStmt();
+ const auto *forStmt = cast<ForStmt>(capturedStmt->getCapturedStmt());
+
+ // Allocate storage for Sema's synthetic 0-based loop counter here, inside
+ // the composite omp.parallel's region (whose insertion point the caller
+ // has already switched to): like the real per-thread counter(s) bound by
+ // the enclosing omp.parallel, this storage must live inside the parallel
+ // region to get a per-thread copy once the region is outlined, rather
+ // than racing across threads if it were hoisted outside like the
+ // counter(s)' mold allocas are.
+ const auto *ivDecl =
+ cast<VarDecl>(cast<DeclRefExpr>(s.getIterationVariable())->getDecl());
+ cgf.emitVarDecl(*ivDecl);
+
+ // The counter(s) were already privatized on the enclosing omp.parallel by
+ // the caller, so unlike emitOMPLoop this never runs preinits
+ // itself: that already happened (or didn't, if host_eval is being
+ // consumed) before this composite's omp.parallel was created.
+ bool consumingHostEval =
+ cgf.ompHostEvalBounds && !cgf.ompHostEvalBounds->applied;
+ mlir::Value zero, tripCount, one;
+ emitOMPLoopNestBounds(cgf, s, consumingHostEval, zero, tripCount, one);
+
+ // Stack the composite loop wrappers around a single shared omp.loop_nest:
+ // omp.distribute {composite} > omp.wsloop {composite} > omp.loop_nest.
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ auto distributeOp =
+ mlir::omp::DistributeOp::create(builder, loc, distributeClauseOps);
+ distributeOp.getRegion().push_back(new mlir::Block());
+ builder.setInsertionPointToStart(&distributeOp.getRegion().front());
+
+ auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, wsloopClauseOps);
+ wsloopOp.getRegion().push_back(new mlir::Block());
+ builder.setInsertionPointToStart(&wsloopOp.getRegion().front());
+
+ // Composite is only attached once the loop body has been emitted: a body
+ // that bailed out leaves an empty region, which omp.composite's verifier
+ // rejects with an error that would mask the original diagnostic.
+ mlir::LogicalResult res = emitOMPLoopNest(cgf, *forStmt, zero, tripCount, one,
+ ivDecl, s.updates()[0]);
+ if (res.succeeded()) {
+ distributeOp.setComposite(/*val=*/true);
+ wsloopOp.setComposite(/*val=*/true);
+ }
+ return res;
+}
+
+/// Emits the `distribute parallel for` composite construct's `distribute`,
+/// `parallel` and `for` leaves: \p distributeItem (the first of the three,
+/// located in \p queue) and the two leaves following it. Lowers to a
+/// composite omp.parallel wrapping the composite omp.distribute + omp.wsloop
+/// + omp.loop_nest stack (see emitOMPDistributeParallelForLoop). Shared by
+/// emitOMPDistributeParallelForDirective and its `teams`/`target teams`
+/// combinations below, which call this once they've peeled off their own
+/// (non-composite) outer wrapper(s).
+///
+/// Templated on \p DirectiveTy (rather than just taking an OMPLoopDirective,
+/// like emitOMPDistributeParallelForLoop below does) so the concrete
+/// directive type reaches emitParallelOp, which needs it: hasCancel() and
+/// getTaskReductionRefExpr() are declared per-concrete-class, not on the
+/// common OMPLoopDirective base.
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitOMPDistributeParallelFor(CIRGenFunction &cgf, const DirectiveTy &s,
+ const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator distributeItem,
+ mlir::Location begin, mlir::Location end) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+
+ assert(distributeItem->id == llvm::omp::OMPD_distribute &&
+ "expected 'distribute' to be the first directive of the composite "
+ "construct");
+ omp::ConstructQueue::const_iterator parallelItem = std::next(distributeItem);
+ assert(parallelItem != queue.end() &&
+ parallelItem->id == llvm::omp::OMPD_parallel &&
+ "expected a 'parallel' directive nested in 'distribute'");
+ omp::ConstructQueue::const_iterator forItem = std::next(parallelItem);
+ assert(forItem != queue.end() && forItem->id == llvm::omp::OMPD_for &&
+ "expected a 'for' directive nested in 'parallel'");
+
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, distributeItem)))
+ return mlir::failure();
+ mlir::omp::DistributeOperands distributeOps;
+ if (emitDistributeClauses(cgf, cgm, builder, loc, distributeItem->clauses,
+ distributeOps)
+ .failed())
+ return mlir::failure();
+
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, parallelItem)))
+ return mlir::failure();
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(cgf, cgm, builder, loc,
+ parallelItem->clauses, parallelOps)))
+ return mlir::failure();
+
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, forItem)))
+ return mlir::failure();
+ mlir::omp::WsloopOperands wsloopOps;
+ if (emitWsloopClauses(cgf, cgm, builder, loc, forItem->clauses, wsloopOps)
+ .failed())
+ return mlir::failure();
+
+ // If the enclosing omp.target already host_eval'd this loop's trip count,
+ // the preinits below are skipped, matching emitOMPLoop.
+ bool consumingHostEval =
+ cgf.ompHostEvalBounds && !cgf.ompHostEvalBounds->applied;
+ if (!consumingHostEval && emitPreinits(cgf, s.getPreInits()).failed())
+ return mlir::failure();
+
+ // The composite loop's shared counter(s) are privatized on the enclosing
+ // omp.parallel (matching Flang: see distribute-parallel-do.f90), not on
+ // the omp.distribute/omp.wsloop stack below, which stays bare. (The
+ // synthetic 0-based `.omp.iv` counter is allocated separately, inside the
+ // omp.parallel region -- see emitOMPDistributeParallelForLoop.)
+ llvm::SmallVector<CIRGenFunction::DeclMapRevertingRAII, 4> counterShadows;
+ llvm::SmallVector<std::pair<const VarDecl *, Address>, 4> counterMolds;
+ counterShadows.reserve(s.counters().size());
+ counterMolds.reserve(s.counters().size());
+ for (auto [counter, privateCounter] :
+ llvm::zip_equal(s.counters(), s.private_counters())) {
+ const auto *vd = cast<VarDecl>(cast<DeclRefExpr>(counter)->getDecl());
+ const auto *privateVd =
+ cast<VarDecl>(cast<DeclRefExpr>(privateCounter)->getDecl());
+ counterShadows.emplace_back(cgf, vd);
+ counterMolds.emplace_back(
+ vd, addPrivateCounter(cgf, cgm, loc, privateVd, parallelOps));
+ }
+
+ return emitParallelOp(
+ cgf, s, queue, parallelItem, begin, end, parallelOps,
+ [&]() -> mlir::LogicalResult {
+ return emitOMPDistributeParallelForLoop(cgf, s, distributeOps,
+ wsloopOps);
+ },
+ /*isComposite=*/true, counterMolds);
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPTaskwaitDirective(const OMPTaskwaitDirective &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPTaskwaitDirective");
@@ -1113,9 +1328,15 @@ CIRGenFunction::emitOMPDistributeDirective(const OMPDistributeDirective &s) {
}
mlir::LogicalResult CIRGenFunction::emitOMPDistributeParallelForDirective(
const OMPDistributeParallelForDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPDistributeParallelForDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ // `distribute parallel for` is a composite construct: a composite
+ // omp.parallel wrapping the composite omp.distribute + omp.wsloop stack.
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ return emitOMPDistributeParallelFor(*this, s, queue, queue.begin(), begin,
+ end);
}
mlir::LogicalResult CIRGenFunction::emitOMPDistributeParallelForSimdDirective(
const OMPDistributeParallelForSimdDirective &s) {
@@ -1196,9 +1417,24 @@ CIRGenFunction::emitOMPTeamsDistributeParallelForSimdDirective(
}
mlir::LogicalResult CIRGenFunction::emitOMPTeamsDistributeParallelForDirective(
const OMPTeamsDistributeParallelForDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPTeamsDistributeParallelForDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ // `teams distribute parallel for` nests the composite `distribute parallel
+ // for` stack inside an omp.teams.
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator teamsItem = queue.begin();
+ assert(teamsItem->id == llvm::omp::OMPD_teams &&
+ "expected 'teams' to be the outermost directive");
+
+ return emitTeams(*this, s, queue, teamsItem, begin, end,
+ [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator distributeItem =
+ std::next(teamsItem);
+ return emitOMPDistributeParallelFor(
+ *this, s, queue, distributeItem, begin, end);
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPTeamsGenericLoopDirective(
const OMPTeamsGenericLoopDirective &s) {
@@ -1266,10 +1502,32 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetTeamsDistributeDirective(
mlir::LogicalResult
CIRGenFunction::emitOMPTargetTeamsDistributeParallelForDirective(
const OMPTargetTeamsDistributeParallelForDirective &s) {
- getCIRGenModule().errorNYI(
- s.getSourceRange(),
- "OpenMP OMPTargetTeamsDistributeParallelForDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ // `target teams distribute parallel for` nests the composite `distribute
+ // parallel for` stack inside omp.teams inside an (SPMD) omp.target.
+ 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 {
+ omp::ConstructQueue::const_iterator distributeItem =
+ std::next(teamsItem);
+ return emitOMPDistributeParallelFor(
+ *this, s, queue, distributeItem, begin, end);
+ });
+ });
}
mlir::LogicalResult
CIRGenFunction::emitOMPTargetTeamsDistributeParallelForSimdDirective(
diff --git a/clang/test/CIR/CodeGenOpenMP/distribute-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/distribute-parallel-for.c
new file mode 100644
index 00000000000000..e09a54519e5f94
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/distribute-parallel-for.c
@@ -0,0 +1,213 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s --check-prefix=CIR
+
+// Lowering to LLVM IR: mostly identical between CIR and OGCG, modulo how
+// each path schedules the distribute and for chunking (see the comments
+// below); they also 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 during(int);
+
+// `distribute parallel for` is a composite construct. It lowers to a
+// composite omp.parallel wrapping the composite omp.distribute + omp.wsloop
+// stack around a single shared omp.loop_nest. Each loop wrapper (plus the
+// enclosing omp.parallel) carries the omp.composite attribute required by
+// the verifier. Unlike the standalone `distribute`/`for` case, the shared loop
+// counter is privatized on the omp.parallel itself (matching Flang's
+// lowering of the same construct), not on omp.distribute/omp.wsloop, which
+// stay bare. Here the composite is closely nested in a standalone `teams`
+// region; the `teams` directive is not the innermost directive of its own
+// queue (`distribute parallel for`'s decomposition still follows it), so the
+// omp.teams carries omp.combined just like the plain `teams`+`distribute`
+// nesting.
+void distribute_parallel_for() {
+ // CIR: cir.func{{.*}}@distribute_parallel_for
+#pragma omp teams
+ {
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+ }
+
+ // CIR: omp.teams {
+
+ // The real counter's "mold" alloca must dominate omp.parallel (it becomes
+ // one of its operands), so it is emitted directly inside omp.teams, before
+ // omp.parallel.
+ // CIR: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CIR: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+
+ // The synthetic normalized counter is written on every iteration, so,
+ // unlike the mold above, it must live inside omp.parallel's region to
+ // become per-thread once the region is outlined.
+ // CIR: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CIR: %[[C0:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+ // CIR: %[[C1:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+ // CIR: %[[C10:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+
+ // CIR: omp.distribute {
+ // CIR-NEXT: omp.wsloop {
+ // CIR-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[C0]]) to (%[[C10]]) step (%[[C1]]) {
+ // CIR: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CIR: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // CIR: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CIR: cir.load align(4) %[[I_PRIV]]
+ // CIR: cir.call @{{.*}}during
+ // CIR: omp.yield
+ // CIR: }
+ // CIR: } {omp.composite}
+ // CIR: } {omp.composite}
+ // CIR: omp.terminator
+ // CIR: } {omp.composite}
+ // CIR: omp.terminator
+ // CIR: } {omp.combined}
+}
+
+// The combined `teams distribute parallel for` directive nests the composite
+// `distribute parallel for` stack inside an omp.teams. The `teams` directive
+// is not the innermost directive, so it carries the omp.combined attribute,
+// while the loop wrappers and their enclosing omp.parallel are marked
+// omp.composite.
+void teams_distribute_parallel_for() {
+ // CIR: cir.func{{.*}}@teams_distribute_parallel_for
+#pragma omp teams distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+
+ // CIR: omp.teams {
+ // CIR: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CIR: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CIR: omp.distribute {
+ // CIR-NEXT: omp.wsloop {
+ // CIR-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+ // CIR: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CIR: cir.call @{{.*}}during
+ // CIR: omp.yield
+ // CIR: }
+ // CIR: } {omp.composite}
+ // CIR: } {omp.composite}
+ // CIR: omp.terminator
+ // CIR: } {omp.composite}
+ // CIR: omp.terminator
+ // CIR: } {omp.combined}
+}
+
+// Host wrapper: __kmpc_fork_teams is identical on both paths (only the
+// pointer to the teams-outlined function differs, absorbed by the
+// wildcard below).
+//
+// LLVM-LABEL: define {{.*}} void @distribute_parallel_for(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @distribute_parallel_for{{.*}})
+
+// The CIR path's teams-outlined function just forks the parallel region
+// directly via __kmpc_fork_call, since CIR's composite lowering privatizes
+// the shared counter on the omp.parallel itself rather than chunking inside
+// the teams-outlined function the way OGCG does (compare the OGCG blocks
+// further down, which do their own distribute chunking here instead).
+// LLVM-CIR-LABEL: define internal void @distribute_parallel_for..omp_par{{.*}}(
+// LLVM-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @distribute_parallel_for..omp_par{{.*}})
+
+// CIR schedules both the distribute and the for chunking in a single
+// __kmpc_dist_for_static_init_4u call (schedule kind 34, "distribute +
+// static for"), ending with a barrier (the composite's implicit `nowait`
+// doesn't apply across the whole construct the way plain `distribute`'s
+// does). CIR's exit block happens to print before its body block even
+// though body always runs first, so we check the actual branch targets
+// below instead of relying on the printed order (same quirk as
+// teams-distribute.c).
+// LLVM-CIR-LABEL: define internal void @distribute_parallel_for..omp_par{{.*}}(
+// LLVM-CIR: %[[TID:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-CIR-NEXT: call void @__kmpc_dist_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.distupperbound, 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(
+// LLVM-CIR: omp_loop.body:
+// LLVM-CIR: call void @during(i32
+
+// The OGCG path instead schedules the distribute chunking (schedule kind
+// 92) in the teams-outlined function, forking the parallel region once per
+// *distribute chunk* via __kmpc_fork_call; the forked function then
+// schedules its own static-for chunking (schedule kind 34) over just that
+// chunk, with no barrier at the end.
+// LLVM-OGCG-LABEL: define internal void @distribute_parallel_for.omp_outlined(
+// LLVM-OGCG: store i32 0, ptr %.omp.comb.lb, align 4
+// LLVM-OGCG-NEXT: store i32 9, ptr %.omp.comb.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: %[[TID3:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-OGCG-SAME: ptr {{.*}}, i32 %[[TID3]], i32 92,
+// LLVM-OGCG-SAME: ptr %.omp.is_last, ptr %.omp.comb.lb, ptr %.omp.comb.ub, ptr %.omp.stride,
+// LLVM-OGCG-SAME: i32 1, i32 1)
+// LLVM-OGCG: omp.inner.for.body:
+// LLVM-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @distribute_parallel_for.omp_outlined.omp_outlined,
+// LLVM-OGCG: omp.loop.exit:
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID3]])
+// LLVM-OGCG-NEXT: ret void
+
+// LLVM-OGCG-LABEL: define internal void @distribute_parallel_for.omp_outlined.omp_outlined(
+// LLVM-OGCG: %[[TID4:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-OGCG-SAME: ptr {{.*}}, i32 %[[TID4]], i32 34,
+// LLVM-OGCG: omp.inner.for.body:
+// LLVM-OGCG: call void @during(i32
+// LLVM-OGCG: omp.loop.exit:
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID4]])
+// LLVM-OGCG-NEXT: ret void
+
+// `teams_distribute_parallel_for` lowers to the exact same shape as
+// `distribute_parallel_for` above; we check the full CFG again here instead
+// of just the call sequence, so a future change that alters this shape for
+// only one of the two spellings still gets caught.
+// LLVM-LABEL: define {{.*}} void @teams_distribute_parallel_for(
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @teams_distribute_parallel_for{{.*}})
+
+// LLVM-CIR-LABEL: define internal void @teams_distribute_parallel_for..omp_par{{.*}}(
+// LLVM-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @teams_distribute_parallel_for..omp_par{{.*}})
+
+// LLVM-CIR-LABEL: define internal void @teams_distribute_parallel_for..omp_par{{.*}}(
+// LLVM-CIR: %[[TID2:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-CIR-NEXT: call void @__kmpc_dist_for_static_init_4u(
+// LLVM-CIR-SAME: ptr {{.*}}, i32 %[[TID2]], i32 34,
+// LLVM-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.distupperbound, 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 %[[TID2]])
+// LLVM-CIR: call void @__kmpc_barrier(
+// LLVM-CIR: omp_loop.body:
+// LLVM-CIR: call void @during(i32
+
+// `teams_distribute_parallel_for` lowers to the exact same shape as
+// `distribute_parallel_for` above on the OGCG path too.
+// LLVM-OGCG-LABEL: define internal void @teams_distribute_parallel_for.omp_outlined(
+// LLVM-OGCG: %[[TID5:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-OGCG-SAME: ptr {{.*}}, i32 %[[TID5]], i32 92,
+// LLVM-OGCG: omp.inner.for.body:
+// LLVM-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @teams_distribute_parallel_for.omp_outlined.omp_outlined,
+// LLVM-OGCG: omp.loop.exit:
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID5]])
+// LLVM-OGCG-NEXT: ret void
+
+// LLVM-OGCG-LABEL: define internal void @teams_distribute_parallel_for.omp_outlined.omp_outlined(
+// LLVM-OGCG: %[[TID6:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-OGCG-SAME: ptr {{.*}}, i32 %[[TID6]], i32 34,
+// LLVM-OGCG: omp.inner.for.body:
+// LLVM-OGCG: call void @during(i32
+// LLVM-OGCG: omp.loop.exit:
+// LLVM-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID6]])
+// LLVM-OGCG-NEXT: ret void
diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
index fd58fae5ea4ec0..2f65390cda368d 100644
--- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
+++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
@@ -51,4 +51,17 @@ void do_things() {
#pragma omp teams distribute collapse(1)
for (int j = 0; j < 10; j++) {
}
+
+ // The composite `distribute parallel for` construct routes clauses to its
+ // `distribute` and `for` directives, both of which report every clause as
+ // not-yet-implemented.
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP DISTRIBUTE 'dist_schedule' clause}}
+#pragma omp teams distribute parallel for dist_schedule(static)
+ for (int j = 0; j < 10; j++) {
+ }
+
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP FOR 'schedule' clause}}
+#pragma omp teams distribute parallel for schedule(static)
+ for (int j = 0; j < 10; j++) {
+ }
}
diff --git a/clang/test/CIR/CodeGenOpenMP/target-teams-distribute-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/target-teams-distribute-parallel-for.c
new file mode 100644
index 00000000000000..0b3e7e556ee8a7
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-teams-distribute-parallel-for.c
@@ -0,0 +1,556 @@
+// 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: the Tripcount host_eval and
+// __tgt_target_kernel call are identical on both paths, so LLVM-HOST
+// checks both; the outlined helper functions differ, 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 during(int);
+
+void explicit_target_teams_distribute_parallel_for() {
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// CIR-HOST: cir.func{{.*}}@explicit_target_teams_distribute_parallel_for
+// CIR-HOST: %[[LB:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: %[[UB:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: %[[STEP:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: omp.target kernel_type(spmd) host_eval(%[[LB]] -> %[[ALB:.*]], %[[UB]] -> %[[AUB:.*]], %[[STEP]] -> %[[ASTEP:.*]] : i32, i32, i32) {
+// CIR-HOST: omp.teams {
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+// CIR-HOST: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-HOST: omp.distribute {
+// CIR-HOST-NEXT: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ALB]]) to (%[[AUB]]) step (%[[ASTEP]]) {
+// CIR-HOST: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+
+// CIR-DEVICE: cir.func{{.*}}@explicit_target_teams_distribute_parallel_for
+// CIR-DEVICE: omp.target kernel_type(spmd) {
+// CIR-DEVICE: omp.teams {
+// CIR-DEVICE: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i{{.*}}>
+// CIR-DEVICE: omp.parallel private(@{{.*}} %{{.*}} -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-DEVICE: omp.distribute {
+// CIR-DEVICE-NEXT: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+// CIR-DEVICE: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+
+void target_teams_distribute_parallel_for() {
+#pragma omp target teams distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// CIR-HOST: cir.func{{.*}}@target_teams_distribute_parallel_for
+// CIR-HOST: %[[LB:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: %[[UB:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: %[[STEP:.*]] = cir.builtin_int_cast %{{.*}} : !s32i -> i32
+// CIR-HOST: omp.target kernel_type(spmd) host_eval(%[[LB]] -> %[[ALB:.*]], %[[UB]] -> %[[AUB:.*]], %[[STEP]] -> %[[ASTEP:.*]] : i32, i32, i32) {
+// CIR-HOST: omp.teams {
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+// CIR-HOST: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-HOST: omp.distribute {
+// CIR-HOST-NEXT: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ALB]]) to (%[[AUB]]) step (%[[ASTEP]]) {
+// CIR-HOST: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+
+// CIR-DEVICE: cir.func{{.*}}@target_teams_distribute_parallel_for
+// CIR-DEVICE: omp.target kernel_type(spmd) {
+// CIR-DEVICE: omp.teams {
+// CIR-DEVICE: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i{{.*}}>
+// CIR-DEVICE: omp.parallel private(@{{.*}} %{{.*}} -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-DEVICE: omp.distribute {
+// CIR-DEVICE-NEXT: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+// CIR-DEVICE: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+
+void target_teams_split_distribute_parallel_for() {
+#pragma omp target teams
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// CIR-HOST: cir.func{{.*}}@target_teams_split_distribute_parallel_for
+// CIR-HOST: omp.target kernel_type(spmd) host_eval(%{{.*}} -> %{{.*}}, %{{.*}} -> %{{.*}}, %{{.*}} -> %{{.*}} : i32, i32, i32) {
+// CIR-HOST: omp.teams {
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+// CIR-HOST: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-HOST: omp.distribute {
+// CIR-HOST-NEXT: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest
+// CIR-HOST: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+
+// CIR-DEVICE: cir.func{{.*}}@target_teams_split_distribute_parallel_for
+// CIR-DEVICE: omp.target kernel_type(spmd) {
+// CIR-DEVICE: omp.teams {
+// CIR-DEVICE: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i{{.*}}>
+// CIR-DEVICE: omp.parallel private(@{{.*}} %{{.*}} -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-DEVICE: omp.distribute {
+// CIR-DEVICE-NEXT: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest
+// CIR-DEVICE: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+
+void target_split_teams_distribute_parallel_for() {
+#pragma omp target
+#pragma omp teams distribute parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// CIR-HOST: cir.func{{.*}}@target_split_teams_distribute_parallel_for
+// CIR-HOST: omp.target kernel_type(spmd) host_eval(%{{.*}} -> %{{.*}}, %{{.*}} -> %{{.*}}, %{{.*}} -> %{{.*}} : i32, i32, i32) {
+// CIR-HOST: omp.teams {
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+// CIR-HOST: omp.parallel private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-HOST: omp.distribute {
+// CIR-HOST-NEXT: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest
+// CIR-HOST: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.composite}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+
+// CIR-DEVICE: cir.func{{.*}}@target_split_teams_distribute_parallel_for
+// CIR-DEVICE: omp.target kernel_type(spmd) {
+// CIR-DEVICE: omp.teams {
+// CIR-DEVICE: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i{{.*}}>
+// CIR-DEVICE: omp.parallel private(@{{.*}} %{{.*}} -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-DEVICE: omp.distribute {
+// CIR-DEVICE-NEXT: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest
+// CIR-DEVICE: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.composite}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+
+// LLVM-HOST-LABEL: define {{.*}} void @explicit_target_teams_distribute_parallel_for(
+// LLVM-HOST: %[[TRIPCOUNT_GEP:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST: store i64 10, ptr %[[TRIPCOUNT_GEP]], align 8
+// LLVM-HOST: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-CIR-LABEL: define {{.*}} void @target_teams_distribute_parallel_for(
+// LLVM-HOST-CIR: %[[TRIPCOUNT_GEP2:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-CIR: store i64 10, ptr %[[TRIPCOUNT_GEP2]], align 8
+// LLVM-HOST-CIR: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-CIR-LABEL: define {{.*}} void @target_teams_split_distribute_parallel_for(
+// LLVM-HOST-CIR: %[[TRIPCOUNT_GEP3:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-CIR: store i64 10, ptr %[[TRIPCOUNT_GEP3]], align 8
+// LLVM-HOST-CIR: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-CIR-LABEL: define {{.*}} void @target_split_teams_distribute_parallel_for(
+// LLVM-HOST-CIR: %[[TRIPCOUNT_GEP4:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-CIR: store i64 10, ptr %[[TRIPCOUNT_GEP4]], align 8
+// LLVM-HOST-CIR: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: %[[TID:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_dist_for_static_init_4u(
+// LLVM-HOST-CIR-SAME: ptr {{.*}}, i32 %[[TID]], i32 34,
+// LLVM-HOST-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.distupperbound, ptr %p.stride,
+// LLVM-HOST-CIR-SAME: i32 1, i32 0)
+// LLVM-HOST-CIR: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM-HOST-CIR: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-HOST-CIR-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM-HOST-CIR: omp_loop.exit:
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID]])
+// LLVM-HOST-CIR: call void @__kmpc_barrier(
+// LLVM-HOST-CIR: omp_loop.body:
+// LLVM-HOST-CIR: call void @during(i32
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined(
+// LLVM-HOST-OGCG: store i32 0, ptr %.omp.comb.lb, align 4
+// LLVM-HOST-OGCG-NEXT: store i32 9, ptr %.omp.comb.ub, align 4
+// LLVM-HOST-OGCG-NEXT: store i32 1, ptr %.omp.stride, align 4
+// LLVM-HOST-OGCG-NEXT: store i32 0, ptr %.omp.is_last, align 4
+// LLVM-HOST-OGCG: %[[TID3:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID3]], i32 92,
+// LLVM-HOST-OGCG-SAME: ptr %.omp.is_last, ptr %.omp.comb.lb, ptr %.omp.comb.ub, ptr %.omp.stride,
+// LLVM-HOST-OGCG-SAME: i32 1, i32 1)
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined,
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID3]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID4:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID4]], i32 34,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void @during(i32
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID4]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define {{.*}} void @target_teams_distribute_parallel_for(
+// LLVM-HOST-OGCG: %[[TRIPCOUNT_GEP2:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-OGCG: store i64 10, ptr %[[TRIPCOUNT_GEP2]], align 8
+// LLVM-HOST-OGCG: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: %[[TID2:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_dist_for_static_init_4u(
+// LLVM-HOST-CIR-SAME: ptr {{.*}}, i32 %[[TID2]], i32 34,
+// LLVM-HOST-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.distupperbound, ptr %p.stride,
+// LLVM-HOST-CIR-SAME: i32 1, i32 0)
+// LLVM-HOST-CIR: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM-HOST-CIR: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-HOST-CIR-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM-HOST-CIR: omp_loop.exit:
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID2]])
+// LLVM-HOST-CIR: call void @__kmpc_barrier(
+// LLVM-HOST-CIR: omp_loop.body:
+// LLVM-HOST-CIR: call void @during(i32
+
+// `target_teams_distribute_parallel_for` lowers to the exact same shape as
+// `explicit_target_teams_distribute_parallel_for` above on the OGCG path
+// too.
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID5:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID5]], i32 92,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined,
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID5]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID6:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID6]], i32 34,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void @during(i32
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID6]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define {{.*}} void @target_teams_split_distribute_parallel_for(
+// LLVM-HOST-OGCG: %[[TRIPCOUNT_GEP3:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-OGCG: store i64 10, ptr %[[TRIPCOUNT_GEP3]], align 8
+// LLVM-HOST-OGCG: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: %[[TID7:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_dist_for_static_init_4u(
+// LLVM-HOST-CIR-SAME: ptr {{.*}}, i32 %[[TID7]], i32 34,
+// LLVM-HOST-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.distupperbound, ptr %p.stride,
+// LLVM-HOST-CIR-SAME: i32 1, i32 0)
+// LLVM-HOST-CIR: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM-HOST-CIR: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-HOST-CIR-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM-HOST-CIR: omp_loop.exit:
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID7]])
+// LLVM-HOST-CIR: call void @__kmpc_barrier(
+// LLVM-HOST-CIR: omp_loop.body:
+// LLVM-HOST-CIR: call void @during(i32
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID8:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID8]], i32 92,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined,
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID8]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID9:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID9]], i32 34,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void @during(i32
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID9]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define {{.*}} void @target_split_teams_distribute_parallel_for(
+// LLVM-HOST-OGCG: %[[TRIPCOUNT_GEP4:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM-HOST-OGCG: store i64 10, ptr %[[TRIPCOUNT_GEP4]], align 8
+// LLVM-HOST-OGCG: call i32 @__tgt_target_kernel(
+
+// LLVM-HOST-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l
+// LLVM-HOST: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}})
+
+// LLVM-HOST-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}(
+// LLVM-HOST-CIR: %[[TID10:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr {{.*}})
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_dist_for_static_init_4u(
+// LLVM-HOST-CIR-SAME: ptr {{.*}}, i32 %[[TID10]], i32 34,
+// LLVM-HOST-CIR-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.distupperbound, ptr %p.stride,
+// LLVM-HOST-CIR-SAME: i32 1, i32 0)
+// LLVM-HOST-CIR: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM-HOST-CIR: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-HOST-CIR-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM-HOST-CIR: omp_loop.exit:
+// LLVM-HOST-CIR-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID10]])
+// LLVM-HOST-CIR: call void @__kmpc_barrier(
+// LLVM-HOST-CIR: omp_loop.body:
+// LLVM-HOST-CIR: call void @during(i32
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID11:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID11]], i32 92,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 2, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined,
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID11]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// LLVM-HOST-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}.omp_outlined.omp_outlined(
+// LLVM-HOST-OGCG: %[[TID12:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// LLVM-HOST-OGCG-SAME: ptr {{.*}}, i32 %[[TID12]], i32 34,
+// LLVM-HOST-OGCG: omp.inner.for.body:
+// LLVM-HOST-OGCG: call void @during(i32
+// LLVM-HOST-OGCG: omp.loop.exit:
+// LLVM-HOST-OGCG-NEXT: call void @__kmpc_for_static_fini(ptr {{.*}}, i32 %[[TID12]])
+// LLVM-HOST-OGCG-NEXT: ret void
+
+// Device kernel entry: __kmpc_target_init/__kmpc_target_deinit and the
+// user_code.entry/worker.exit split are identical on both paths.
+//
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null,
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_distribute_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}..omp_par
+// LLVM-DEVICE-CIR: call void @__kmpc_barrier(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 91,
+// LLVM-DEVICE-OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined, ptr null,
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_fini(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_explicit_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 33,
+// LLVM-DEVICE-OGCG: call void @during(i32
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_fini(
+
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null,
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_distribute_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}..omp_par
+// LLVM-DEVICE-CIR: call void @__kmpc_barrier(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 91,
+// LLVM-DEVICE-OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined, ptr null,
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_fini(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 33,
+// LLVM-DEVICE-OGCG: call void @during(i32
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_fini(
+
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null,
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_distribute_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}..omp_par
+// LLVM-DEVICE-CIR: call void @__kmpc_barrier(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 91,
+// LLVM-DEVICE-OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined, ptr null,
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_fini(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_teams_split_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 33,
+// LLVM-DEVICE-OGCG: call void @during(i32
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_fini(
+
+// LLVM-DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l
+// LLVM-DEVICE: call i32 @__kmpc_target_init(
+// LLVM-DEVICE: user_code.entry:
+// LLVM-DEVICE-CIR: call void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-OGCG: call void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE: call void @__kmpc_target_deinit()
+// LLVM-DEVICE: ret void
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null,
+
+// LLVM-DEVICE-CIR-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}..omp_par{{.*}}(
+// LLVM-DEVICE-CIR: call void @__kmpc_distribute_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}..omp_par
+// LLVM-DEVICE-CIR: call void @__kmpc_barrier(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 91,
+// LLVM-DEVICE-OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined, ptr null,
+// LLVM-DEVICE-OGCG: call void @__kmpc_distribute_static_fini(
+
+// LLVM-DEVICE-OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_split_teams_distribute_parallel_for_l{{.*}}_omp_outlined_omp_outlined(
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_init_4(ptr {{.*}}, i32 {{.*}}, i32 33,
+// LLVM-DEVICE-OGCG: call void @during(i32
+// LLVM-DEVICE-OGCG: call void @__kmpc_for_static_fini(
More information about the llvm-branch-commits
mailing list