[clang] [CIR][OpenMP] Add support for the OpenMP 'for' directive (PR #229256)
Jan Leyonberg via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 5 15:50:49 PDT 2026
https://github.com/jsjodin created https://github.com/llvm/llvm-project/pull/229256
This patch adds support for wsloop in ClangIR: the `for` directive and its combined forms `parallel for` and `target parallel for`. This is lowered to an omp.wsloop + omp.loop_nest, nested utilizing the existing queue-based decomposition.
Assisted-by: Cursor / Claude Sonnet 5 High
>From 7cbbd123cf6719effd936da934cbb361c792df60 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Fri, 2 Oct 2026 10:55:22 -0400
Subject: [PATCH] [CIR][OpenMP] Add support for the OpenMP 'for' directive
This patch adds support for wsloop in ClangIR: the `for` directive and its
combined forms `parallel for` and `target parallel for`. This is lowered to an
omp.wsloop + omp.loop_nest, nested utilizing the existing queue-based
decomposition.
Assisted-by: Cursor / Claude Sonnet 5 High
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 381 +++++++++++++++++-
clang/test/CIR/CodeGenOpenMP/for-loop-forms.c | 93 +++++
clang/test/CIR/CodeGenOpenMP/parallel-for.c | 41 ++
clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c | 207 ++++++++++
.../CIR/CodeGenOpenMP/target-parallel-for.c | 117 ++++++
5 files changed, 831 insertions(+), 8 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenMP/for-loop-forms.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/parallel-for.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 9d937f5ed67f9..beb431129834e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -132,6 +132,304 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
});
}
+/// Casts a CIR value to the given CIR integer type, loading through a
+/// pointer first if needed.
+static mlir::Value ensureCIRIntType(CIRGenBuilderTy &builder,
+ mlir::Location loc, mlir::Value cirValue,
+ cir::IntType targetCIRType) {
+ if (mlir::isa<cir::PointerType>(cirValue.getType()))
+ cirValue = cir::LoadOp::create(builder, loc, cirValue).getResult();
+
+ if (cirValue.getType() == targetCIRType)
+ return cirValue;
+
+ return builder.createCast(loc, cir::CastKind::integral, cirValue,
+ targetCIRType);
+}
+
+/// Converts a CIR integer value to the equivalent builtin MLIR integer type.
+static mlir::Value cirIntToBuiltinInt(CIRGenBuilderTy &builder,
+ mlir::Location loc,
+ mlir::Value cirValue) {
+ auto cirIntType = mlir::cast<cir::IntType>(cirValue.getType());
+ mlir::Type builtinIntType = builder.getIntegerType(cirIntType.getWidth());
+ return builder.createBuiltinIntCast(loc, cirValue, builtinIntType);
+}
+
+/// Emits the Sema-generated pre-init statements for an OpenMP loop directive.
+static mlir::LogicalResult emitPreinits(CIRGenFunction &cgf,
+ const Stmt *preInits) {
+ if (!preInits)
+ return mlir::success();
+
+ llvm::SmallVector<const Stmt *> stmts;
+ if (const auto *compound = dyn_cast<CompoundStmt>(preInits))
+ llvm::append_range(stmts, compound->body());
+ else
+ stmts.push_back(preInits);
+
+ for (const Stmt *stmt : stmts) {
+ if (const auto *declStmt = dyn_cast<DeclStmt>(stmt)) {
+ for (const Decl *d : declStmt->decls())
+ cgf.emitVarDecl(cast<VarDecl>(*d));
+ } else {
+ if (cgf.emitStmt(stmt, /*useCurrentScope=*/true).failed())
+ return mlir::failure();
+ }
+ }
+ return mlir::success();
+}
+
+/// Emits an omp.loop_nest for the worksharing loop `forStmt`
+static mlir::LogicalResult emitOMPLoopNest(CIRGenFunction &cgf,
+ const ForStmt &forStmt,
+ mlir::Value lb, mlir::Value ub,
+ mlir::Value step, bool inclusive,
+ const VarDecl *inductionVar) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(forStmt.getSourceRange());
+
+ auto loopNestOp = mlir::omp::LoopNestOp::create(
+ builder, loc, /*collapse_num_loops=*/1, lb, ub, step,
+ /*loop_inclusive=*/inclusive, /*tile_sizes=*/nullptr);
+ mlir::Block *block = new mlir::Block();
+ loopNestOp.getRegion().push_back(block);
+ block->addArgument(lb.getType(), loc);
+
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(block);
+
+ // Store the induction variable block argument into the loop variable alloca,
+ // converting back from the builtin integer to the CIR integer type.
+ mlir::Value iv = block->getArgument(0);
+ Address inductionAddr = cgf.getAddrOfLocalVar(inductionVar);
+ mlir::Value civVal =
+ builder.createBuiltinIntCast(loc, iv, inductionAddr.getElementType());
+ builder.createStore(loc, civVal, inductionAddr);
+
+ mlir::LogicalResult bodyRes = mlir::success();
+ if (forStmt.getBody())
+ if (cgf.emitStmt(forStmt.getBody(), /*useCurrentScope=*/true).failed())
+ bodyRes = mlir::failure();
+
+ mlir::omp::YieldOp::create(builder, cgf.getLoc(forStmt.getEndLoc()));
+ return bodyRes;
+}
+
+/// Evaluates the clauses allowed on an omp.wsloop leaf; none are supported
+/// yet, so every eligible clause is reported as NYI.
+static mlir::LogicalResult
+emitWsloopClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
+ CIRGenBuilderTy &builder, mlir::Location loc,
+ llvm::ArrayRef<const OMPClause *> clauses,
+ mlir::omp::WsloopOperands &clauseOps) {
+ OpenMPClauseEmitter ce(cgf, cgm, builder, loc, clauses);
+ return ce.emitNYI</*supported=*/>(
+ /*nyi=*/OpenMPNYIClauseList<
+ OMPAllocateClause, OMPCollapseClause, OMPFirstprivateClause,
+ OMPLastprivateClause, OMPLinearClause, OMPNowaitClause,
+ OMPOrderClause, OMPOrderedClause, OMPPrivateClause,
+ OMPReductionClause, OMPScheduleClause>{},
+ llvm::omp::Directive::OMPD_for);
+}
+
+/// Emits the loop's lower bound from the induction variable's initializer
+/// (`int i = <expr>`). Returns failure if the variable has no initializer.
+static mlir::FailureOr<mlir::Value>
+emitLoopLowerBound(CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ mlir::Location loc, const VarDecl *varDecl,
+ cir::IntType cirIntType) {
+ if (!varDecl->hasInit())
+ return mlir::failure();
+ mlir::Value v = cgf.emitScalarExpr(varDecl->getInit());
+ return ensureCIRIntType(builder, loc, v, cirIntType);
+}
+
+/// Returns true if the given expression (after stripping parens and implicit
+/// casts) is a reference to `varDecl`.
+static bool refersToVar(const Expr *e, const VarDecl *varDecl) {
+ const auto *ref = dyn_cast<DeclRefExpr>(e->IgnoreParenImpCasts());
+ return ref && ref->getDecl() == varDecl;
+}
+
+/// The loop's upper bound, and whether it is inclusive (`<=`/`>=`) or
+/// exclusive (`<`/`>`).
+struct LoopUpperBound {
+ mlir::Value value;
+ bool inclusive;
+};
+
+/// Emits the loop's upper bound from the controlling comparison
+/// (`var < ub`, `ub < var`, and the `<=`/`>`/`>=` equivalents, with the
+/// induction variable on either side). Returns failure if the condition
+/// isn't one of these forms.
+static mlir::FailureOr<LoopUpperBound>
+emitLoopUpperBound(CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ mlir::Location loc, const ForStmt &forStmt,
+ const VarDecl *varDecl, cir::IntType cirIntType) {
+ const auto *condBinOp = dyn_cast_or_null<BinaryOperator>(forStmt.getCond());
+ if (!condBinOp)
+ return mlir::failure();
+ BinaryOperatorKind op = condBinOp->getOpcode();
+ if (op != BO_LT && op != BO_LE && op != BO_GT && op != BO_GE)
+ return mlir::failure();
+ bool inclusive = (op == BO_LE || op == BO_GE);
+
+ const Expr *boundExpr;
+ if (refersToVar(condBinOp->getLHS(), varDecl))
+ boundExpr = condBinOp->getRHS();
+ else if (refersToVar(condBinOp->getRHS(), varDecl))
+ boundExpr = condBinOp->getLHS();
+ else
+ return mlir::failure();
+
+ mlir::Value v = cgf.emitScalarExpr(boundExpr);
+ return LoopUpperBound{ensureCIRIntType(builder, loc, v, cirIntType),
+ inclusive};
+}
+
+/// Emits the loop's step from the induction variable's increment expression
+/// (`i++`, `--i`, `i += <expr>`, `i -= <expr>`, `i = i + <expr>`,
+/// `i = <expr> + i`, or `i = i - <expr>`). These are the only increment
+/// forms OpenMP's canonical loop form allows, so one of them always
+/// matches.
+static mlir::Value emitLoopStep(CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ mlir::Location loc, const ForStmt &forStmt,
+ const VarDecl *varDecl,
+ cir::IntType cirIntType) {
+ if (const auto *unary = dyn_cast_or_null<UnaryOperator>(forStmt.getInc())) {
+ if (unary->isIncrementDecrementOp() &&
+ refersToVar(unary->getSubExpr(), varDecl))
+ return builder.getConstInt(loc, cirIntType,
+ unary->isIncrementOp() ? 1 : -1);
+ } else if (const auto *binOp =
+ dyn_cast_or_null<BinaryOperator>(forStmt.getInc())) {
+ BinaryOperatorKind op = binOp->getOpcode();
+ const Expr *stepExpr = nullptr;
+ bool negate = false;
+ if ((op == BO_AddAssign || op == BO_SubAssign) &&
+ refersToVar(binOp->getLHS(), varDecl)) {
+ stepExpr = binOp->getRHS();
+ negate = (op == BO_SubAssign);
+ } else if (op == BO_Assign && refersToVar(binOp->getLHS(), varDecl)) {
+ if (const auto *sub =
+ dyn_cast<BinaryOperator>(binOp->getRHS()->IgnoreParenImpCasts());
+ sub && sub->isAdditiveOp()) {
+ bool isAdd = sub->getOpcode() == BO_Add;
+ if (refersToVar(sub->getLHS(), varDecl)) {
+ stepExpr = sub->getRHS();
+ negate = !isAdd;
+ } else if (isAdd && refersToVar(sub->getRHS(), varDecl)) {
+ stepExpr = sub->getLHS();
+ }
+ }
+ }
+ if (stepExpr) {
+ mlir::Value v = cgf.emitScalarExpr(stepExpr);
+ mlir::Value step = ensureCIRIntType(builder, loc, v, cirIntType);
+ if (negate)
+ step = ensureCIRIntType(builder, loc, builder.createNeg(loc, step),
+ cirIntType);
+ return step;
+ }
+ }
+ llvm_unreachable("ForStmt increment must be a canonical OpenMP form, "
+ "already validated by Sema");
+}
+
+/// The loop's lower/upper bounds and step, as CIR integers (no induction
+/// variable alloca involved), plus whether the upper bound is inclusive.
+struct OMPLoopBounds {
+ mlir::Value lowerBound;
+ LoopUpperBound upperBound;
+ mlir::Value step;
+};
+
+/// Emits pre-inits and computes the loop's bounds/step as CIR integers (no
+/// induction variable alloca). Delegates to emitLoopLowerBound/
+/// emitLoopUpperBound/emitLoopStep, which are independent of one another.
+static mlir::FailureOr<OMPLoopBounds>
+computeOMPLoopBounds(CIRGenFunction &cgf, const OMPLoopDirective &s,
+ const ForStmt &forStmt, const VarDecl *inductionVar) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+
+ if (emitPreinits(cgf, s.getPreInits()).failed())
+ return mlir::failure();
+
+ QualType loopVarQType = inductionVar->getType();
+ auto cirIntType = mlir::cast<cir::IntType>(cgf.convertType(loopVarQType));
+
+ mlir::FailureOr<mlir::Value> lowerBound =
+ emitLoopLowerBound(cgf, builder, loc, inductionVar, cirIntType);
+ if (mlir::failed(lowerBound))
+ return mlir::failure();
+
+ mlir::FailureOr<LoopUpperBound> upperBound =
+ emitLoopUpperBound(cgf, builder, loc, forStmt, inductionVar, cirIntType);
+ if (mlir::failed(upperBound))
+ return mlir::failure();
+
+ mlir::Value step =
+ emitLoopStep(cgf, builder, loc, forStmt, inductionVar, cirIntType);
+
+ return OMPLoopBounds{*lowerBound, *upperBound, step};
+}
+
+/// Lowers an OMPLoopDirective's `for` leaf to an omp.wsloop + omp.loop_nest.
+/// `for` is always innermost, so unlike emitParallelOp/emitTargetOp this
+/// never needs to mark the op as combined.
+static mlir::LogicalResult
+emitOMPWorksharingLoop(CIRGenFunction &cgf, const OMPLoopDirective &s,
+ omp::ConstructQueue::const_iterator item) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, item)))
+ return mlir::failure();
+
+ mlir::omp::WsloopOperands clauseOps;
+ if (emitWsloopClauses(cgf, cgm, builder, loc, item->clauses, clauseOps)
+ .failed())
+ return mlir::failure();
+
+ const CapturedStmt *capturedStmt = s.getInnermostCapturedStmt();
+ const auto *forStmt = cast<ForStmt>(capturedStmt->getCapturedStmt());
+
+ const auto *declStmt = dyn_cast_or_null<DeclStmt>(forStmt->getInit());
+ const auto *varDecl =
+ declStmt ? dyn_cast<VarDecl>(declStmt->getSingleDecl()) : nullptr;
+ if (!varDecl)
+ return mlir::failure();
+
+ mlir::FailureOr<OMPLoopBounds> bounds =
+ computeOMPLoopBounds(cgf, s, *forStmt, varDecl);
+ if (mlir::failed(bounds))
+ return mlir::failure();
+
+ if (forStmt->getInit())
+ if (cgf.emitStmt(forStmt->getInit(), /*useCurrentScope=*/true).failed())
+ return mlir::failure();
+
+ // omp.loop_nest requires IntLikeType operands, not CIR integer types.
+ mlir::Value builtinLB = cirIntToBuiltinInt(builder, loc, bounds->lowerBound);
+ mlir::Value builtinUB =
+ cirIntToBuiltinInt(builder, loc, bounds->upperBound.value);
+ mlir::Value builtinStep = cirIntToBuiltinInt(builder, loc, bounds->step);
+
+ auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, clauseOps);
+ mlir::Block *innerBlock = new mlir::Block();
+ wsloopOp.getRegion().push_back(innerBlock);
+
+ // The for-init was already emitted above, so the induction variable alloca
+ // lives outside the loop region.
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(innerBlock);
+ return emitOMPLoopNest(cgf, *forStmt, builtinLB, builtinUB, builtinStep,
+ bounds->upperBound.inclusive, varDecl);
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPTaskwaitDirective(const OMPTaskwaitDirective &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPTaskwaitDirective");
@@ -181,8 +479,9 @@ CIRGenFunction::emitOMPFuseDirective(const OMPFuseDirective &s) {
}
mlir::LogicalResult
CIRGenFunction::emitOMPForDirective(const OMPForDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPForDirective");
- return mlir::failure();
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ return emitOMPWorksharingLoop(*this, s, queue.begin());
}
mlir::LogicalResult
CIRGenFunction::emitOMPForSimdDirective(const OMPForSimdDirective &s) {
@@ -216,9 +515,31 @@ CIRGenFunction::emitOMPCriticalDirective(const OMPCriticalDirective &s) {
}
mlir::LogicalResult
CIRGenFunction::emitOMPParallelForDirective(const OMPParallelForDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPParallelForDirective");
- 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 parallelItem = queue.begin();
+ assert(parallelItem->id == llvm::omp::OMPD_parallel &&
+ "expected 'parallel' to be the outermost leaf");
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, parallelItem)))
+ return mlir::failure();
+
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin,
+ parallelItem->clauses, parallelOps)))
+ return mlir::failure();
+
+ return emitParallelOp(
+ *this, s, queue, parallelItem, begin, end, parallelOps,
+ [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator forItem = std::next(parallelItem);
+ assert(forItem != queue.end() && forItem->id == llvm::omp::OMPD_for &&
+ "expected a 'for' leaf nested in 'parallel'");
+ return emitOMPWorksharingLoop(*this, s, forItem);
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPParallelForSimdDirective(
const OMPParallelForSimdDirective &s) {
@@ -496,9 +817,53 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective(
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelForDirective(
const OMPTargetParallelForDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPTargetParallelForDirective");
- 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 leaf");
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, targetItem)))
+ return mlir::failure();
+
+ mlir::omp::TargetExtOperands targetOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+ if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin,
+ targetItem->clauses, targetOps, mapSyms)))
+ return mlir::failure();
+
+ return emitTargetOp(
+ *this, s, queue, targetItem, begin, end, targetOps, mapSyms,
+ [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator parallelItem =
+ std::next(targetItem);
+ assert(parallelItem != queue.end() &&
+ parallelItem->id == llvm::omp::OMPD_parallel &&
+ "expected a 'parallel' leaf nested in 'target'");
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, parallelItem)))
+ return mlir::failure();
+
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder,
+ begin, parallelItem->clauses,
+ parallelOps)))
+ return mlir::failure();
+
+ return emitParallelOp(
+ *this, s, queue, parallelItem, begin, end, parallelOps,
+ [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator forItem =
+ std::next(parallelItem);
+ assert(forItem != queue.end() &&
+ forItem->id == llvm::omp::OMPD_for &&
+ "expected a 'for' leaf nested in 'parallel'");
+ return emitOMPWorksharingLoop(*this, s, forItem);
+ });
+ });
}
mlir::LogicalResult
CIRGenFunction::emitOMPTaskLoopDirective(const OMPTaskLoopDirective &s) {
diff --git a/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c b/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c
new file mode 100644
index 0000000000000..1c9f1686a6943
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c
@@ -0,0 +1,93 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
+
+void during(int);
+
+// Decreasing loop using the common `i > 0; i--` idiom: the induction
+// variable is on the left of the comparison and the increment is a plain
+// decrement.
+void dec_unary() {
+ // CHECK: cir.func{{.*}}@{{.*}}dec_unary
+#pragma omp for
+ for (int i = 20; i > 0; i--) {
+ during(i);
+ }
+ // CHECK: %[[C20_CIR:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[CM1_CIR:.*]] = cir.const #cir.int<-1> : !s32i
+ // CHECK: %[[C20:.*]] = cir.builtin_int_cast %[[C20_CIR]] : !s32i -> i32
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[CM1:.*]] = cir.builtin_int_cast %[[CM1_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%[[C20]]) to (%[[C0]]) step (%[[CM1]]) {
+}
+
+// Decreasing loop using a compound assignment: the step must be negated,
+// not just taken verbatim from the RHS of `-=`.
+void dec_compound() {
+ // CHECK: cir.func{{.*}}@{{.*}}dec_compound
+#pragma omp for
+ for (int i = 20; i > 0; i -= 2) {
+ during(i);
+ }
+ // CHECK: %[[C20_CIR:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[CM2_CIR:.*]] = cir.const #cir.int<-2> : !s32i
+ // CHECK: %[[C20:.*]] = cir.builtin_int_cast %[[C20_CIR]] : !s32i -> i32
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[CM2:.*]] = cir.builtin_int_cast %[[CM2_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%[[C20]]) to (%[[C0]]) step (%[[CM2]]) {
+}
+
+// Decreasing loop using the `var = var - incr` canonical assignment form:
+// the step must be negated here too.
+void dec_assign_sub() {
+ // CHECK: cir.func{{.*}}@{{.*}}dec_assign_sub
+#pragma omp for
+ for (int i = 20; i > 0; i = i - 2) {
+ during(i);
+ }
+ // CHECK: %[[C20_CIR:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[CM2_CIR:.*]] = cir.const #cir.int<-2> : !s32i
+ // CHECK: %[[C20:.*]] = cir.builtin_int_cast %[[C20_CIR]] : !s32i -> i32
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[CM2:.*]] = cir.builtin_int_cast %[[CM2_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%[[C20]]) to (%[[C0]]) step (%[[CM2]]) {
+}
+
+// Increasing loop using the `var = incr + var` canonical assignment form.
+void inc_assign_add_rev() {
+ // CHECK: cir.func{{.*}}@{{.*}}inc_assign_add_rev
+#pragma omp for
+ for (int i = 0; i < 20; i = 2 + i) {
+ during(i);
+ }
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[C20_CIR:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[C2_CIR:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[C20:.*]] = cir.builtin_int_cast %[[C20_CIR]] : !s32i -> i32
+ // CHECK: %[[C2:.*]] = cir.builtin_int_cast %[[C2_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%[[C0]]) to (%[[C20]]) step (%[[C2]]) {
+}
+
+// Increasing loop with the upper bound written on the left of the
+// comparison (`ub > var`), which is also valid canonical loop form.
+void ub_on_left() {
+ // CHECK: cir.func{{.*}}@{{.*}}ub_on_left
+#pragma omp for
+ for (int i = 0; 20 > i; i++) {
+ during(i);
+ }
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[C20_CIR:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[C20:.*]] = cir.builtin_int_cast %[[C20_CIR]] : !s32i -> i32
+ // CHECK: %[[C1:.*]] = cir.builtin_int_cast %[[C1_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%[[C0]]) to (%[[C20]]) step (%[[C1]]) {
+}
diff --git a/clang/test/CIR/CodeGenOpenMP/parallel-for.c b/clang/test/CIR/CodeGenOpenMP/parallel-for.c
new file mode 100644
index 0000000000000..208f9c1861f67
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/parallel-for.c
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
+
+void during(int);
+
+// The combined `parallel for` directive decomposes into a `parallel` leaf and a
+// `for` leaf. It lowers to an omp.wsloop + omp.loop_nest nested directly inside
+// an omp.parallel, mirroring the separate `parallel` / `for` nesting.
+void parallel_for() {
+ // CHECK: cir.func{{.*}}@{{.*}}parallel_for
+#pragma omp parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+
+ // CHECK: omp.parallel {
+
+ // The induction variable alloca is emitted before the wsloop.
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) init : !cir.ptr<!s32i>
+
+ // CIR constants for the loop bounds cast to builtin integers.
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+ // CHECK: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+ // CHECK: %[[C10:.*]] = cir.builtin_int_cast %[[C10_CIR]] : !s32i -> i32
+ // CHECK: %[[C1:.*]] = cir.builtin_int_cast %[[C1_CIR]] : !s32i -> i32
+
+ // CHECK: omp.wsloop {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[C0]]) to (%[[C10]]) step (%[[C1]]) {
+
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[I_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+
+ // CHECK: omp.yield
+ // CHECK: }
+ // CHECK: }
+ // CHECK: omp.terminator
+ // The parallel is a non-innermost leaf of the combined construct.
+ // CHECK: } {omp.combined}
+}
diff --git a/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c b/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c
new file mode 100644
index 0000000000000..160f323835853
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c
@@ -0,0 +1,207 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -emit-llvm %s -o %t.ll
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
+
+void before(int);
+void during(int);
+void after(int);
+
+void emit_simple_for() {
+ int j = 5;
+ before(j);
+#pragma omp parallel
+ {
+#pragma omp for
+ for (int i = 0; i < 10; i++) {
+ during(j);
+ }
+ }
+ after(j);
+}
+
+// LLVM-LABEL: define{{.*}} void @emit_simple_for()
+// LLVM: %[[STRUCT_ARG:.*]] = alloca { ptr }, align 8
+// LLVM: %[[J_SLOT:.*]] = alloca i32, align 4
+// LLVM: store i32 5, ptr %[[J_SLOT]], align 4
+// LLVM: call void @before(i32
+// LLVM: %[[GEP0:.*]] = getelementptr { ptr }, ptr %[[STRUCT_ARG]], i32 0, i32 0
+// LLVM-NEXT: store ptr %[[J_SLOT]], ptr %[[GEP0]], align 8
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(
+// LLVM-SAME: ptr @{{.*}}, i32 1, ptr @emit_simple_for..omp_par, ptr %[[STRUCT_ARG]])
+// LLVM: %[[J_VAL:.*]] = load i32, ptr %[[J_SLOT]], align 4
+// LLVM-NEXT: call void @after(i32 noundef %[[J_VAL]])
+// LLVM-NEXT: ret void
+
+// OGCG-LABEL: define{{.*}} void @emit_simple_for()
+// OGCG: %[[J:.*]] = alloca i32, align 4
+// OGCG: store i32 5, ptr %[[J]], align 4
+// OGCG: %[[J_VAL:.*]] = load i32, ptr %[[J]], align 4
+// OGCG-NEXT: call void @before(i32 noundef %[[J_VAL]])
+// OGCG-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(
+// OGCG-SAME: ptr @{{.*}}, i32 1, ptr @emit_simple_for.omp_outlined, ptr %[[J]])
+// OGCG: %[[J_VAL2:.*]] = load i32, ptr %[[J]], align 4
+// OGCG-NEXT: call void @after(i32 noundef %[[J_VAL2]])
+// OGCG-NEXT: ret void
+
+// LLVM-LABEL: define internal void @emit_simple_for..omp_par(
+// LLVM-SAME: ptr noalias %tid.addr, ptr noalias %zero.addr, ptr %[[STRUCT:.*]])
+// LLVM: %[[GEP_J:.*]] = getelementptr { ptr }, ptr %[[STRUCT]], i32 0, i32 0
+// LLVM-NEXT: %[[J_PTR:.*]] = load ptr, ptr %[[GEP_J]], align 8
+// LLVM: %p.lastiter = alloca i32, align 4
+// LLVM-NEXT: %p.lowerbound = alloca i32, align 4
+// LLVM-NEXT: %p.upperbound = alloca i32, align 4
+// LLVM-NEXT: %p.stride = alloca i32, align 4
+// LLVM: store i32 0, ptr %p.lowerbound, align 4
+// LLVM-NEXT: store i32 9, ptr %p.upperbound, align 4
+// LLVM-NEXT: store i32 1, ptr %p.stride, align 4
+// LLVM: %[[TID:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr @{{.*}})
+// LLVM-NEXT: call void @__kmpc_for_static_init_4u(
+// LLVM-SAME: ptr @{{.*}}, i32 %[[TID]], i32 34,
+// LLVM-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.stride,
+// LLVM-SAME: i32 1, i32 0)
+// LLVM: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ]
+// LLVM: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit
+// LLVM: omp_loop.exit:
+// LLVM-NEXT: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID]])
+// LLVM: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}})
+// LLVM: omp_loop.body:
+// LLVM: %[[J_VAL:.*]] = load i32, ptr %[[J_PTR]], align 4
+// LLVM-NEXT: call void @during(i32 noundef %[[J_VAL]])
+// LLVM: %omp_loop.next = add nuw i32 %omp_loop.iv, 1
+// LLVM: ret void
+
+// OGCG-LABEL: define internal void @emit_simple_for.omp_outlined(
+// OGCG-SAME: ptr noalias noundef %.global_tid., ptr noalias noundef %.bound_tid.,
+// OGCG-SAME: ptr noundef nonnull align 4 dereferenceable(4) %j)
+// OGCG: %.global_tid..addr = alloca ptr, align 8
+// OGCG-NEXT: %.bound_tid..addr = alloca ptr, align 8
+// OGCG-NEXT: %j.addr = alloca ptr, align 8
+// OGCG: %.omp.lb = alloca i32, align 4
+// OGCG-NEXT: %.omp.ub = alloca i32, align 4
+// OGCG-NEXT: %.omp.stride = alloca i32, align 4
+// OGCG-NEXT: %.omp.is_last = alloca i32, align 4
+// OGCG: store ptr %.global_tid., ptr %.global_tid..addr, align 8
+// OGCG: store ptr %j, ptr %j.addr, align 8
+// OGCG: %[[J_LOADED:.*]] = load ptr, ptr %j.addr, align 8
+// OGCG: store i32 0, ptr %.omp.lb, align 4
+// OGCG-NEXT: store i32 9, ptr %.omp.ub, align 4
+// OGCG-NEXT: store i32 1, ptr %.omp.stride, align 4
+// OGCG-NEXT: store i32 0, ptr %.omp.is_last, align 4
+// OGCG: %[[GTID_VAL_PTR:.*]] = load ptr, ptr %.global_tid..addr, align 8
+// OGCG-NEXT: %[[TID:.*]] = load i32, ptr %[[GTID_VAL_PTR]], align 4
+// OGCG-NEXT: call void @__kmpc_for_static_init_4(
+// OGCG-SAME: ptr @{{.*}}, i32 %[[TID]], i32 34,
+// OGCG-SAME: ptr %.omp.is_last, ptr %.omp.lb, ptr %.omp.ub, ptr %.omp.stride,
+// OGCG-SAME: i32 1, i32 1)
+// OGCG: %[[UB_VAL:.*]] = load i32, ptr %.omp.ub, align 4
+// OGCG-NEXT: %[[CMP:.*]] = icmp sgt i32 %[[UB_VAL]], 9
+// OGCG-NEXT: br i1 %[[CMP]], label %cond.true, label %cond.false
+// OGCG: cond.end:
+// OGCG: %[[COND:.*]] = phi i32 [ 9, %cond.true ], [ %{{.*}}, %cond.false ]
+// OGCG-NEXT: store i32 %[[COND]], ptr %.omp.ub, align 4
+// OGCG: omp.inner.for.cond:
+// OGCG: %[[CMP1:.*]] = icmp sle i32 %{{.*}}, %{{.*}}
+// OGCG-NEXT: br i1 %[[CMP1]], label %omp.inner.for.body, label %omp.inner.for.end
+// OGCG: omp.inner.for.body:
+// OGCG: %[[J_VAL:.*]] = load i32, ptr %[[J_LOADED]], align 4
+// OGCG-NEXT: call void @during(i32 noundef %[[J_VAL]])
+// OGCG: omp.loop.exit:
+// OGCG: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID]])
+// OGCG-NEXT: call void @__kmpc_barrier(ptr @{{.*}}, i32 %[[TID]])
+// OGCG-NEXT: ret void
+
+// OGCG: declare void @__kmpc_for_static_init_4(ptr, i32, i32, ptr, ptr, ptr, ptr, i32, i32)
+// OGCG: declare void @__kmpc_for_static_fini(ptr, i32)
+// OGCG: declare void @__kmpc_barrier(ptr, i32)
+// OGCG: declare {{.*}}void @__kmpc_fork_call(ptr, i32, ptr, ...)
+
+void emit_for_with_vars() {
+ int j = 5;
+ before(j);
+#pragma omp parallel
+ {
+ int lb = 1;
+ long ub = 10;
+ short step = 1;
+#pragma omp for
+ for (int i = 0; i < ub; i = i + step) {
+ during(j);
+ }
+ }
+ after(j);
+}
+
+// LLVM-LABEL: define{{.*}} void @emit_for_with_vars()
+// LLVM: %[[STRUCT_ARG:.*]] = alloca { ptr }, align 8
+// LLVM: call void @before(i32
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(
+// LLVM-SAME: ptr @{{.*}}, i32 1, ptr @emit_for_with_vars..omp_par, ptr %[[STRUCT_ARG]])
+// LLVM: call void @after(i32
+// LLVM: ret void
+
+// OGCG-LABEL: define{{.*}} void @emit_for_with_vars()
+// OGCG: %[[J:.*]] = alloca i32, align 4
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(
+// OGCG-SAME: ptr @{{.*}}, i32 1, ptr @emit_for_with_vars.omp_outlined, ptr %[[J]])
+// OGCG: ret void
+
+// LLVM-LABEL: define internal void @emit_for_with_vars..omp_par(
+// LLVM-SAME: ptr noalias %tid.addr, ptr noalias %zero.addr, ptr %[[STRUCT:.*]])
+// LLVM: store i32 1, ptr %{{.*}}, align 4
+// LLVM: store i64 10, ptr %{{.*}}, align 8
+// LLVM: store i16 1, ptr %{{.*}}, align 2
+// LLVM: %{{.*}} = sext i16 %{{.*}} to i32
+// LLVM: %{{.*}} = sext i32 %{{.*}} to i64
+// LLVM: %{{.*}} = sext i16 %{{.*}} to i64
+// LLVM: %{{.*}} = sdiv i64
+// LLVM: %omp_loop.tripcount = select i1 %{{.*}}, i32 0, i32 %{{.*}}
+// LLVM: store i32 0, ptr %p.lowerbound, align 4
+// LLVM-NEXT: %[[UB:.*]] = sub i32 %omp_loop.tripcount, 1
+// LLVM-NEXT: store i32 %[[UB]], ptr %p.upperbound, align 4
+// LLVM-NEXT: store i32 1, ptr %p.stride, align 4
+// LLVM: %[[TID2:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr @{{.*}})
+// LLVM-NEXT: call void @__kmpc_for_static_init_4u(
+// LLVM-SAME: ptr @{{.*}}, i32 %[[TID2]], i32 34,
+// LLVM-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.stride,
+// LLVM-SAME: i32 1, i32 0)
+// LLVM: %omp_loop.iv = phi i32
+// LLVM: icmp ult i32 %omp_loop.iv, %{{.*}}
+// LLVM: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID2]])
+// LLVM: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}})
+// LLVM: call void @during(i32
+// LLVM: ret void
+
+// OGCG-LABEL: define internal void @emit_for_with_vars.omp_outlined(
+// OGCG-SAME: ptr noalias noundef %.global_tid., ptr noalias noundef %.bound_tid.,
+// OGCG-SAME: ptr noundef nonnull align 4 dereferenceable(4) %j)
+// OGCG: %lb = alloca i32, align 4
+// OGCG-NEXT: %ub = alloca i64, align 8
+// OGCG-NEXT: %step = alloca i16, align 2
+// OGCG: store i32 1, ptr %lb, align 4
+// OGCG-NEXT: store i64 10, ptr %ub, align 8
+// OGCG-NEXT: store i16 1, ptr %step, align 2
+// OGCG: %{{.*}} = sext i16 %{{.*}} to i32
+// OGCG: %{{.*}} = sext i32 %{{.*}} to i64
+// OGCG: %{{.*}} = sext i16 %{{.*}} to i64
+// OGCG: %{{.*}} = sdiv i64
+// OGCG: %[[PRECOND:.*]] = icmp slt i64 0, %{{.*}}
+// OGCG-NEXT: br i1 %[[PRECOND]], label %omp.precond.then, label %omp.precond.end
+// OGCG: omp.precond.then:
+// OGCG: store i32 0, ptr %.omp.lb, align 4
+// OGCG: store i32 1, ptr %.omp.stride, align 4
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG-SAME: ptr @{{.*}}, i32 %{{.*}}, i32 34,
+// OGCG-SAME: i32 1, i32 1)
+// OGCG: call void @during(i32
+// OGCG: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %{{.*}})
+// OGCG: omp.precond.end:
+// OGCG: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}})
+// OGCG-NEXT: ret void
+
+// LLVM: declare void @__kmpc_for_static_init_4u(ptr, i32, i32, ptr, ptr, ptr, ptr, i32, i32)
+// LLVM: declare void @__kmpc_for_static_fini(ptr, i32)
+// LLVM: declare i32 @__kmpc_global_thread_num(ptr)
+// LLVM: declare void @__kmpc_barrier(ptr, i32)
+// LLVM: declare {{.*}}void @__kmpc_fork_call(ptr, i32, ptr, ...)
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
new file mode 100644
index 0000000000000..e7a0a8898c6a5
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
@@ -0,0 +1,117 @@
+// 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): allocas live in the private address space.
+// 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 during(int);
+
+// The legal nesting of target, parallel and for lowers to an omp.wsloop +
+// omp.loop_nest inside omp.parallel inside omp.target. The worksharing loop
+// bounds and induction variable are cast between CIR and builtin integers
+// with cir.builtin_int_cast on both the host and the GPU device.
+void target_parallel_for() {
+#pragma omp target
+#pragma omp parallel
+#pragma omp for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// CIR-HOST: cir.func{{.*}}@target_parallel_for
+// CIR-HOST: omp.target kernel_type(generic) {
+// CIR-HOST: omp.parallel {
+
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) init : !cir.ptr<!s32i>
+// CIR-HOST: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-HOST: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+// CIR-HOST: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-HOST: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+// CIR-HOST: %[[C10:.*]] = cir.builtin_int_cast %[[C10_CIR]] : !s32i -> i32
+// CIR-HOST: %[[C1:.*]] = cir.builtin_int_cast %[[C1_CIR]] : !s32i -> i32
+
+// CIR-HOST: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[C0]]) to (%[[C10]]) step (%[[C1]]) {
+// CIR-HOST: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+// CIR-HOST: cir.store align(4) %[[IV_CIR]], %[[I_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: }
+// CIR-HOST: omp.terminator
+// CIR-HOST: omp.terminator
+// CIR-HOST: }
+
+// CIR-DEVICE: cir.func{{.*}}@target_parallel_for
+// CIR-DEVICE: omp.target kernel_type(generic) {
+// CIR-DEVICE: omp.parallel {
+
+// CIR-DEVICE: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) init : !cir.ptr<!s32i, target_address_space(5)>
+// CIR-DEVICE: %[[I_CAST:.*]] = cir.cast address_space %[[I_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
+// CIR-DEVICE: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-DEVICE: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+// CIR-DEVICE: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-DEVICE: %[[C0:.*]] = cir.builtin_int_cast %[[C0_CIR]] : !s32i -> i32
+// CIR-DEVICE: %[[C10:.*]] = cir.builtin_int_cast %[[C10_CIR]] : !s32i -> i32
+// CIR-DEVICE: %[[C1:.*]] = cir.builtin_int_cast %[[C1_CIR]] : !s32i -> i32
+
+// CIR-DEVICE: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[C0]]) to (%[[C10]]) step (%[[C1]]) {
+// CIR-DEVICE: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+// CIR-DEVICE: cir.store align(4) %[[IV_CIR]], %[[I_CAST]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: }
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: }
+
+// The combined `target parallel for` directive decomposes into `target`,
+// `parallel` and `for` leaves and lowers to the same nesting as the explicit
+// target/parallel/for above: an omp.wsloop + omp.loop_nest inside omp.parallel
+// inside omp.target.
+void combined_target_parallel_for() {
+#pragma omp target parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// The `target` and `parallel` are non-innermost leaves of the combined
+// construct, so both carry the omp.combined attribute (unlike the explicitly
+// nested directives above).
+// CIR-HOST: cir.func{{.*}}@combined_target_parallel_for
+// CIR-HOST: omp.target kernel_type(generic) {
+// CIR-HOST: omp.parallel {
+// CIR-HOST: %[[CI_ALLOCA:.*]] = cir.alloca "i" align(4) init : !cir.ptr<!s32i>
+// CIR-HOST: omp.wsloop {
+// CIR-HOST-NEXT: omp.loop_nest (%[[CIV:.*]]) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+// CIR-HOST: cir.call @{{.*}}during
+// CIR-HOST: omp.yield
+// CIR-HOST: }
+// CIR-HOST: }
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+// CIR-HOST: omp.terminator
+// CIR-HOST: } {omp.combined}
+
+// CIR-DEVICE: cir.func{{.*}}@combined_target_parallel_for
+// CIR-DEVICE: omp.target kernel_type(generic) {
+// CIR-DEVICE: omp.parallel {
+// CIR-DEVICE: omp.wsloop {
+// CIR-DEVICE-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+// CIR-DEVICE: cir.call @{{.*}}during
+// CIR-DEVICE: omp.yield
+// CIR-DEVICE: }
+// CIR-DEVICE: }
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
+// CIR-DEVICE: omp.terminator
+// CIR-DEVICE: } {omp.combined}
More information about the cfe-commits
mailing list