[clang] [CIR][OpenMP] Add support for the OpenMP 'for' directive (PR #229256)
Jan Leyonberg via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 15 08:03:33 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/229256
>From c606e2254e71db763358471aedb64535dc590d06 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
---
.../CIR/Dialect/Builder/CIRBaseBuilder.h | 7 +
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 381 +++++++++++++++---
clang/test/CIR/CodeGenOpenMP/for-loop-forms.c | 228 +++++++++++
clang/test/CIR/CodeGenOpenMP/parallel-for.c | 54 +++
clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c | 207 ++++++++++
.../target-parallel-for-llvm-device.c | 84 ++++
.../target-parallel-for-llvm-host.c | 83 ++++
.../CIR/CodeGenOpenMP/target-parallel-for.c | 127 ++++++
.../target-parallel-llvm-device.c | 85 ++++
.../CodeGenOpenMP/target-parallel-llvm-host.c | 94 +++++
10 files changed, 1283 insertions(+), 67 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-llvm-device.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c
diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
index 36d583cfe9fbed..f4dbbc58fc7eed 100644
--- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
+++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
@@ -673,6 +673,13 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
return createBuiltinIntCast(src.getLoc(), src, newTy);
}
+ /// Casts \p src to the builtin MLIR integer type with the same width as
+ /// its own CIR integer type.
+ mlir::Value createBuiltinIntCast(mlir::Location loc, mlir::Value src) {
+ auto cirIntTy = mlir::cast<cir::IntType>(src.getType());
+ return createBuiltinIntCast(loc, src, getIntegerType(cirIntTy.getWidth()));
+ }
+
mlir::Value createIntToPtr(mlir::Value src, mlir::Type newTy) {
return createCast(cir::CastKind::int_to_ptr, src, newTy);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 9d937f5ed67f90..71e943085bcde3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -105,6 +105,27 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
return res;
}
+/// Checks, emits clauses for, and emits an omp.parallel for the `parallel`
+/// directive at \p parallelItem, deferring to \p emitBody for the region's
+/// contents.
+template <typename DirectiveTy>
+static mlir::LogicalResult emitParallel(
+ CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator parallelItem, mlir::Location begin,
+ mlir::Location end, llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, parallelItem)))
+ return mlir::failure();
+
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(cgf, cgf.getCIRGenModule(),
+ cgf.getBuilder(), begin,
+ parallelItem->clauses, parallelOps)))
+ return mlir::failure();
+
+ return emitParallelOp(cgf, s, queue, parallelItem, begin, end, parallelOps,
+ emitBody);
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
@@ -114,22 +135,216 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
omp::ConstructQueue::const_iterator item = queue.begin();
- if (mlir::failed(checkSynthesizedClauses(*this, s, item)))
+ return emitParallel(
+ *this, s, queue, item, begin, end, [&]() -> mlir::LogicalResult {
+ // emitStmt() rejects CapturedStmt directly; the
+ // parent construct must unwrap it, so emit the
+ // inner statement instead.
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
+}
+
+/// Emits a single Sema-generated pre-init statement for an OpenMP loop
+/// directive: a VarDecl's DeclStmt is lowered directly (it has no
+/// initializer CIRGenFunction::emitStmt can handle on its own), everything
+/// else goes through the general statement emitter.
+static mlir::LogicalResult emitPreinit(CIRGenFunction &cgf,
+ const Stmt *preInit) {
+ if (const auto *declStmt = dyn_cast<DeclStmt>(preInit)) {
+ for (const Decl *d : declStmt->decls())
+ cgf.emitVarDecl(cast<VarDecl>(*d));
+ return mlir::success();
+ }
+ return cgf.emitStmt(preInit, /*useCurrentScope=*/true);
+}
+
+/// 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();
+
+ if (const auto *compound = dyn_cast<CompoundStmt>(preInits)) {
+ for (const Stmt *stmt : compound->body())
+ if (emitPreinit(cgf, stmt).failed())
+ return mlir::failure();
+ return mlir::success();
+ }
+
+ return emitPreinit(cgf, preInits);
+}
+
+/// Emits an omp.loop_nest that iterates a normalized `[0, tripCount)` range
+/// and recomputes the real induction variable each iteration using the given
+/// update expression.
+static mlir::LogicalResult
+emitOMPLoopNest(CIRGenFunction &cgf, const ForStmt &forStmt, mlir::Value lb,
+ mlir::Value ub, mlir::Value step, const VarDecl *ivDecl,
+ const Expr *update) {
+ 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=*/false, /*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 normalized counter block argument into the normalized loop
+ // variable's storage location, then let Sema's update expression recompute
+ // the real induction variable from it.
+ mlir::Value iv = block->getArgument(0);
+ Address ivAddr = cgf.getAddrOfLocalVar(ivDecl);
+ mlir::Value civVal =
+ builder.createBuiltinIntCast(loc, iv, ivAddr.getElementType());
+ builder.createStore(loc, civVal, ivAddr);
+ cgf.emitIgnoredExpr(update);
+
+ 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 the `for` directive; 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);
+}
+
+/// Creates a module-level `omp.private` op (a "privatizer") for a scalar
+/// of CIR type elemTy.
+static mlir::omp::PrivateClauseOp
+createCounterPrivatizer(CIRGenModule &cgm, mlir::Location loc,
+ llvm::StringRef baseName, mlir::Type elemTy) {
+ // getGlobalValue() is an O(1) cache lookup, unlike
+ // mlir::SymbolTable::lookupSymbolIn()'s O(n) module scan.
+ unsigned uniquingCounter = 0;
+ llvm::SmallString<64> name = mlir::SymbolTable::generateSymbolName<64>(
+ (baseName + ".privatizer").str(),
+ [&](llvm::StringRef candidate) {
+ return cgm.getGlobalValue(candidate) != nullptr;
+ },
+ uniquingCounter);
+
+ mlir::ModuleOp mod = cgm.getModule();
+ mlir::OpBuilder b(mod.getContext());
+ b.setInsertionPointToEnd(mod.getBody());
+ auto privatizer = mlir::omp::PrivateClauseOp::create(
+ b, loc, mlir::TypeRange{}, b.getStringAttr(name),
+ mlir::TypeAttr::get(elemTy));
+ cgm.insertGlobalSymbol(privatizer);
+ return privatizer;
+}
+
+/// Registers a loop counter as predetermined-private on the omp.wsloop op.
+static Address addPrivateCounter(CIRGenFunction &cgf, CIRGenModule &cgm,
+ mlir::Location loc, const VarDecl *privateVd,
+ mlir::omp::WsloopOperands &clauseOps) {
+ cgf.emitVarDecl(*privateVd);
+ Address moldAddr = cgf.getAddrOfLocalVar(privateVd);
+
+ mlir::omp::PrivateClauseOp privatizer = createCounterPrivatizer(
+ cgm, loc, privateVd->getName(), moldAddr.getElementType());
+
+ clauseOps.privateVars.push_back(moldAddr.getPointer());
+ clauseOps.privateSyms.push_back(
+ mlir::FlatSymbolRefAttr::get(privatizer.getSymNameAttr()));
+ return moldAddr;
+}
+
+/// Lowers an OMPLoopDirective's `for` directive to an omp.wsloop +
+/// omp.loop_nest.
+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::ParallelOperands clauseOps;
- if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin,
- item->clauses, clauseOps)))
+ mlir::omp::WsloopOperands clauseOps;
+ if (emitWsloopClauses(cgf, cgm, builder, loc, item->clauses, clauseOps)
+ .failed())
return mlir::failure();
- return emitParallelOp(
- *this, s, queue, item, begin, end, clauseOps,
- [&]() -> mlir::LogicalResult {
- // emitStmt() rejects CapturedStmt directly; the parent construct
- // must unwrap it, so emit the inner statement instead.
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
- return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
- });
+ const CapturedStmt *capturedStmt = s.getInnermostCapturedStmt();
+ const auto *forStmt = cast<ForStmt>(capturedStmt->getCapturedStmt());
+
+ if (emitPreinits(cgf, s.getPreInits()).failed())
+ return mlir::failure();
+
+ // Allocate storage for Sema's normalized 0-based loop counter.
+ const auto *ivDecl =
+ cast<VarDecl>(cast<DeclRefExpr>(s.getIterationVariable())->getDecl());
+ cgf.emitVarDecl(*ivDecl);
+ 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, clauseOps));
+ }
+
+ mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
+ auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
+ mlir::Value zero = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 0));
+ mlir::Value one = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 1));
+ mlir::Value tripCount = builder.createBuiltinIntCast(loc, tripCountCir);
+
+ auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, clauseOps);
+ mlir::Block *innerBlock = new mlir::Block();
+ // Ensures the block arguments match BlockArgOpenMPOpInterface's expected
+ // order.
+ for (auto &counterMold : counterMolds)
+ innerBlock->addArgument(counterMold.second.getPointer().getType(), loc);
+ wsloopOp.getRegion().push_back(innerBlock);
+
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(innerBlock);
+
+ // Bind each counter to its private block argument.
+ for (unsigned idx = 0; idx < counterMolds.size(); ++idx) {
+ const VarDecl *vd = counterMolds[idx].first;
+ Address moldAddr = counterMolds[idx].second;
+ Address privAddr(innerBlock->getArgument(idx), moldAddr.getElementType(),
+ moldAddr.getAlignment());
+ cgf.replaceAddrOfLocalVar(vd, privAddr);
+ cgf.symbolTable.insert(vd, privAddr.getPointer());
+ }
+
+ return emitOMPLoopNest(cgf, *forStmt, zero, tripCount, one, ivDecl,
+ s.updates()[0]);
}
mlir::LogicalResult
@@ -181,8 +396,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 +432,22 @@ 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 directive");
+
+ return emitParallel(
+ *this, s, queue, parallelItem, begin, end, [&]() -> mlir::LogicalResult {
+ 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'");
+ return emitOMPWorksharingLoop(*this, s, forItem);
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPParallelForSimdDirective(
const OMPParallelForSimdDirective &s) {
@@ -387,6 +616,28 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s,
return res;
}
+/// Checks, emits clauses for, and emits an omp.target for the `target`
+/// directive at \p targetItem, deferring to \p emitBody for the region's
+/// contents.
+template <typename DirectiveTy>
+static mlir::LogicalResult emitTarget(
+ CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator targetItem, mlir::Location begin,
+ mlir::Location end, llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ if (mlir::failed(checkSynthesizedClauses(cgf, s, targetItem)))
+ return mlir::failure();
+
+ mlir::omp::TargetExtOperands targetOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+ if (mlir::failed(emitTargetClauses(cgf, cgf.getCIRGenModule(),
+ cgf.getBuilder(), begin,
+ targetItem->clauses, targetOps, mapSyms)))
+ return mlir::failure();
+
+ return emitTargetOp(cgf, s, queue, targetItem, begin, end, targetOps, mapSyms,
+ emitBody);
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
@@ -396,22 +647,12 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
omp::ConstructQueue::const_iterator item = queue.begin();
- if (mlir::failed(checkSynthesizedClauses(*this, s, item)))
- return mlir::failure();
-
- mlir::omp::TargetExtOperands clauseOps;
- llvm::SmallVector<const VarDecl *> mapSyms;
- if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin,
- item->clauses, clauseOps, mapSyms)))
- return mlir::failure();
-
- return emitTargetOp(*this, s, queue, item, begin, end, clauseOps, mapSyms,
- [&]() -> mlir::LogicalResult {
- const CapturedStmt *cs =
- s.getCapturedStmt(llvm::omp::OMPD_target);
- return emitStmt(cs->getCapturedStmt(),
- /*useCurrentScope=*/true);
- });
+ return emitTarget(
+ *this, s, queue, item, begin, end, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult
CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) {
@@ -456,49 +697,55 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective(
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();
+ "expected 'target' to be the outermost directive");
- return emitTargetOp(
- *this, s, queue, targetItem, begin, end, targetOps, mapSyms,
- [&]() -> mlir::LogicalResult {
+ return emitTarget(
+ *this, s, queue, targetItem, begin, end, [&]() -> 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 {
- const CapturedStmt *cs =
- s.getCapturedStmt(llvm::omp::OMPD_parallel);
- return emitStmt(cs->getCapturedStmt(),
- /*useCurrentScope=*/true);
- });
+ "expected a 'parallel' directive nested in 'target'");
+
+ return emitParallel(*this, s, queue, parallelItem, begin, end,
+ [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs =
+ s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
});
}
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 directive");
+
+ return emitTarget(
+ *this, s, queue, targetItem, begin, end, [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator parallelItem =
+ std::next(targetItem);
+ assert(parallelItem != queue.end() &&
+ parallelItem->id == llvm::omp::OMPD_parallel &&
+ "expected a 'parallel' directive nested in 'target'");
+
+ return emitParallel(
+ *this, s, queue, parallelItem, begin, end,
+ [&]() -> mlir::LogicalResult {
+ 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'");
+ 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 00000000000000..9875d18cf198f6
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c
@@ -0,0 +1,228 @@
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
+
+void during(int);
+int getLB(void);
+
+// All the loop forms below normalize to an omp.loop_nest iterating a
+// `[0, tripCount)` range, where tripCount is computed from the original
+// bounds/step.
+
+// Decreasing loop: `i > 0; i--` (induction variable on the left, plain
+// unary decrement).
+void dec_unary() {
+ // CHECK: cir.func{{.*}}@{{.*}}dec_unary
+#pragma omp for
+ for (int i = 20; i > 0; i--) {
+ during(i);
+ }
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) {
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // The update expression recomputes `i = 20 - 1 * iv`.
+ // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i
+ // CHECK: %[[C1:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C1]] : !s32i
+ // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i
+ // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Decreasing loop using compound assignment: `i -= 2`.
+void dec_compound() {
+ // CHECK: cir.func{{.*}}@{{.*}}dec_compound
+#pragma omp for
+ for (int i = 20; i > 0; i -= 2) {
+ during(i);
+ }
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) {
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // The update expression recomputes `i = 20 - 2 * iv`.
+ // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i
+ // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i
+ // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i
+ // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Decreasing loop using the `i = i - 2` canonical assignment form.
+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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) {
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // The update expression recomputes `i = 20 - 2 * iv`.
+ // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i
+ // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i
+ // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i
+ // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i
+ // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Increasing loop using the `i = 2 + i` 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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) {
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // The update expression recomputes `i = 0 + 2 * iv`.
+ // CHECK: %[[C0:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i
+ // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i
+ // CHECK: %[[ADD:.*]] = cir.add nsw %[[C0]], %[[MUL]] : !s32i
+ // CHECK: cir.store align(4) %[[ADD]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Increasing loop with the upper bound on the left of the comparison
+// (`20 > i`).
+void ub_on_left() {
+ // CHECK: cir.func{{.*}}@{{.*}}ub_on_left
+#pragma omp for
+ for (int i = 0; 20 > i; i++) {
+ during(i);
+ }
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) {
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // The update expression recomputes `i = 0 + 1 * iv`.
+ // CHECK: %[[C0:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i
+ // CHECK: %[[C1:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C1]] : !s32i
+ // CHECK: %[[ADD:.*]] = cir.add nsw %[[C0]], %[[MUL]] : !s32i
+ // CHECK: cir.store align(4) %[[ADD]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Loop with a side-effecting init expression (a function call, `getLB()`).
+void side_effecting_init() {
+ // CHECK: cir.func{{.*}}@{{.*}}side_effecting_init
+#pragma omp for
+ for (int i = getLB(); i < 10; i++) {
+ during(i);
+ }
+ // CHECK: %[[CAPTURE:.*]] = cir.alloca ".capture_expr." align(4) init : !cir.ptr<!s32i>
+ // CHECK: %[[CALL:.*]] = cir.call @{{.*}}getLB{{.*}}() : {{.*}} -> !s32i
+ // CHECK-NEXT: cir.store{{.*}} %[[CALL]], %[[CAPTURE]]
+ // CHECK-NOT: cir.call @{{.*}}getLB
+ // CHECK: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) {
+ // CHECK: omp.loop_nest
+ // CHECK-NOT: cir.call @{{.*}}getLB
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Loop with a pointer induction variable (`int *p`).
+void pointer_induction_var() {
+ // CHECK: cir.func{{.*}}@{{.*}}pointer_induction_var
+ int a[10];
+#pragma omp for
+ for (int *p = a; p < a + 10; ++p) {
+ during(*p);
+ }
+ // CHECK: %[[P_ALLOCA:.*]] = cir.alloca "p" align(8) : !cir.ptr<!cir.ptr<!s32i>>
+ // CHECK: omp.wsloop private(@{{.*}} %[[P_ALLOCA]] -> %[[P_PRIV:.*]] : !cir.ptr<!cir.ptr<!s32i>>) {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i64 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+ // CHECK: %[[P_BASE:.*]] = cir.load{{.*}} : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
+ // CHECK: %[[P_NEXT:.*]] = cir.ptr_stride %[[P_BASE]], %{{.*}} : (!cir.ptr<!s32i>, !s64i) -> !cir.ptr<!s32i>
+ // CHECK: cir.store{{.*}} %[[P_NEXT]], %[[P_PRIV]]
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Pre-declared induction variable (plain assignment, no DeclStmt) with a
+// `!=` condition.
+void predeclared_var_and_not_equal_cond() {
+ // CHECK: cir.func{{.*}}@{{.*}}predeclared_var_and_not_equal_cond
+ int i;
+#pragma omp for
+ for (i = 0; i != 10; i++) {
+ during(i);
+ }
+ // CHECK: cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_MOLD:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_MOLD]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+ // CHECK: cir.store{{.*}} %{{.*}}, %[[IV_ALLOCA]]
+ // CHECK: cir.store{{.*}} %{{.*}}, %[[I_PRIV]]
+ // CHECK: cir.call @{{.*}}during
+}
+
+// Same pre-declared counter and `!=` condition, inside a combined
+// `parallel for`.
+void predeclared_var_inside_parallel_for() {
+ // CHECK: cir.func{{.*}}@{{.*}}predeclared_var_inside_parallel_for
+ int i;
+#pragma omp parallel for
+ for (i = 0; i != 10; i++) {
+ during(i);
+ }
+ // CHECK: cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: omp.parallel {
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_MOLD:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_MOLD]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
+ // CHECK: cir.store{{.*}} %{{.*}}, %[[IV_ALLOCA]]
+ // CHECK: cir.store{{.*}} %{{.*}}, %[[I_PRIV]]
+ // CHECK: cir.call @{{.*}}during
+}
diff --git a/clang/test/CIR/CodeGenOpenMP/parallel-for.c b/clang/test/CIR/CodeGenOpenMP/parallel-for.c
new file mode 100644
index 00000000000000..76f31ac0786996
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/parallel-for.c
@@ -0,0 +1,54 @@
+// 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 {
+
+ // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+ // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+
+ // The loop bounds are normalized to a `[0, tripCount)` range; tripCount is
+ // computed from the original bounds/step.
+ // CHECK: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+ // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[LBM1:.*]] = cir.sub nsw %[[C0_CIR]], %[[C1_CIR]] : !s32i
+ // CHECK: %[[C1B:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[LBM1P1:.*]] = cir.add nsw %[[LBM1]], %[[C1B]] : !s32i
+ // CHECK: %[[SPAN:.*]] = cir.sub nsw %[[C10_CIR]], %[[LBM1P1]] : !s32i
+ // CHECK: %[[C1C:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %[[SPAN]], %[[C1C]] : !s32i
+ // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+ // CHECK: %[[ONE_CIR:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK: %[[ONE:.*]] = cir.builtin_int_cast %[[ONE_CIR]] : !s32i -> i32
+ // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+
+ // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) {
+
+ // The normalized counter block argument is stored into the normalized
+ // counter's alloca, then the update expression (`i = 0 + 1 * iv`)
+ // recomputes the real induction variable from it.
+ // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+ // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+ // CHECK: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !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 00000000000000..160f3238358539
--- /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-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c
new file mode 100644
index 00000000000000..428174a1d3bf2c
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c
@@ -0,0 +1,84 @@
+// Two-step host-BC -> device pipeline that mirrors the offloading driver.
+//
+// Step 1: Host compilation to bitcode (provides offload entry info to device pass).
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc
+//
+// Step 2: Device compilation using host BC.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void during(int);
+
+void target_parallel_for() {
+#pragma omp target
+#pragma omp parallel
+#pragma omp for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+void combined_target_parallel_for() {
+#pragma omp target parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_for_l
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par{{.*}}, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par{{.*}}.wrapper, ptr {{.*}}, i64 0, i32 0)
+// LLVM: call void @__kmpc_target_deinit()
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par
+// LLVM: call void @__kmpc_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par,
+// LLVM: call void @__kmpc_barrier(
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par(i32
+// LLVM: call void @during(i32
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_combined_target_parallel_for_l
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par{{.*}}, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par{{.*}}.wrapper, ptr {{.*}}, i64 0, i32 0)
+// LLVM: call void @__kmpc_target_deinit()
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par
+// LLVM: call void @__kmpc_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par,
+// LLVM: call void @__kmpc_barrier(
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par(i32
+// LLVM: call void @during(i32
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_for_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 0, i32 0)
+// OGCG: call void @__kmpc_target_deinit()
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}_omp_outlined(
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG: call void @during(i32
+// OGCG: call void @__kmpc_for_static_fini(
+// OGCG: call void @__kmpc_barrier(
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_combined_target_parallel_for_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 0, i32 0)
+// OGCG: call void @__kmpc_target_deinit()
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}_omp_outlined(
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG: call void @during(i32
+// OGCG: call void @__kmpc_for_static_fini(
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
new file mode 100644
index 00000000000000..41428595851922
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
@@ -0,0 +1,83 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void during(int);
+
+void target_parallel_for() {
+#pragma omp target
+#pragma omp parallel
+#pragma omp for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+void combined_target_parallel_for() {
+#pragma omp target parallel for
+ for (int i = 0; i < 10; i++) {
+ during(i);
+ }
+}
+
+// The explicit target/parallel/for nesting and the combined `target parallel
+// for` spelling lower identically: an omp.target forks an omp.parallel via
+// __kmpc_fork_call, whose body statically schedules the omp.wsloop via
+// __kmpc_for_static_init_4/_4u.
+//
+// LLVM-LABEL: define {{.*}} void @target_parallel_for(
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}(
+
+// LLVM-LABEL: define {{.*}} void @combined_target_parallel_for(
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}(
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par)
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par(
+// LLVM: call void @__kmpc_for_static_init_4u(
+// LLVM: call void @__kmpc_for_static_fini(
+// LLVM: call void @__kmpc_barrier(
+// LLVM: call void @during(i32
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par)
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par(
+// LLVM: call void @__kmpc_for_static_init_4u(
+// LLVM: call void @__kmpc_for_static_fini(
+// LLVM: call void @__kmpc_barrier(
+// LLVM: call void @during(i32
+
+// OGCG-LABEL: define {{.*}} void @target_parallel_for(
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}(
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}.omp_outlined)
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}.omp_outlined(
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG: call void @during(i32
+// OGCG: call void @__kmpc_for_static_fini(
+
+// OGCG-LABEL: define {{.*}} void @combined_target_parallel_for(
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}(
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}.omp_outlined)
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}.omp_outlined(
+// OGCG: call void @__kmpc_for_static_init_4(
+// OGCG: call void @during(i32
+// OGCG: call void @__kmpc_for_static_fini(
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 00000000000000..d9ba201a312b57
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
@@ -0,0 +1,127 @@
+// 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);
+
+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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i>
+// CIR-HOST: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i>
+
+// The loop bounds are normalized to a `[0, tripCount)` range; tripCount is
+// computed from the original bounds/step.
+// CIR-HOST: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+// CIR-HOST: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-HOST: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-HOST: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+// CIR-HOST: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-HOST: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+// CIR-HOST: %[[ONE_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-HOST: %[[ONE:.*]] = cir.builtin_int_cast %[[ONE_CIR]] : !s32i -> i32
+// CIR-HOST: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+
+// CIR-HOST: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-HOST-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) {
+// CIR-HOST: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+// CIR-HOST: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i>
+// 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: }
+// 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 {
+
+// The two allocas and their address-space casts can be emitted in either
+// relative order, so match them unordered (CHECK-DAG) rather than assuming
+// a specific interleaving.
+// CIR-DEVICE-DAG: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i, target_address_space(5)>
+// CIR-DEVICE-DAG: %[[IV_CAST:.*]] = cir.cast address_space %[[IV_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
+// CIR-DEVICE-DAG: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i, target_address_space(5)>
+// CIR-DEVICE-DAG: %[[I_CAST:.*]] = cir.cast address_space %[[I_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
+// CIR-DEVICE: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i
+// CIR-DEVICE: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-DEVICE: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-DEVICE: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i
+// CIR-DEVICE: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-DEVICE: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32
+// CIR-DEVICE: %[[ONE_CIR:.*]] = cir.const #cir.int<1> : !s32i
+// CIR-DEVICE: %[[ONE:.*]] = cir.builtin_int_cast %[[ONE_CIR]] : !s32i -> i32
+// CIR-DEVICE: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32
+
+// CIR-DEVICE: omp.wsloop private(@{{.*}} %[[I_CAST]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) {
+// CIR-DEVICE-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) {
+// CIR-DEVICE: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i
+// CIR-DEVICE: cir.store align(4) %[[IV_CIR]], %[[IV_CAST]] : !s32i, !cir.ptr<!s32i>
+// 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: }
+// 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);
+ }
+}
+
+// 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) : !cir.ptr<!s32i>
+// CIR-HOST: omp.wsloop private(@{{.*}} %[[CI_ALLOCA]] -> %{{.*}} : !cir.ptr<!s32i>) {
+// 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 private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) {
+// 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}
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c
new file mode 100644
index 00000000000000..4ce2badf7b182a
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c
@@ -0,0 +1,85 @@
+// Two-step host-BC -> device pipeline that mirrors the offloading driver.
+//
+// Step 1: Host compilation to bitcode (provides offload entry info to device pass).
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc
+//
+// Step 2: Device compilation using host BC.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \
+// RUN: -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void use(int);
+
+void target_parallel(int x) {
+#pragma omp target parallel map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+void target_parallel_proc_bind(int x) {
+#pragma omp target parallel proc_bind(spread) map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par.wrapper, ptr {{.*}}, i64 1, i32 0)
+// LLVM: call void @__kmpc_target_deinit()
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par.wrapper, ptr {{.*}}, i64 1, i32 0)
+// LLVM: call void @__kmpc_target_deinit()
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 1, i32 0)
+// OGCG: call void @__kmpc_target_deinit()
+// OGCG: ret void
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}_omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
+
+// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l
+// OGCG: call i32 @__kmpc_target_init(
+// OGCG: user_code.entry:
+// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 1, i32 0)
+// OGCG: call void @__kmpc_target_deinit()
+// OGCG: ret void
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}_omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c
new file mode 100644
index 00000000000000..fee411bc509e9f
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c
@@ -0,0 +1,94 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=OGCG
+
+void use(int);
+
+void target_parallel(int x) {
+#pragma omp target parallel map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+void target_parallel_proc_bind(int x) {
+#pragma omp target parallel proc_bind(spread) map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+// Host wrappers
+//
+// LLVM-LABEL: define {{.*}} void @target_parallel(
+// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// LLVM-LABEL: define {{.*}} void @target_parallel_proc_bind(
+// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]]
+// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4
+// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4
+// LLVM: call i32 @__tgt_target_kernel(
+// LLVM: omp_offload.failed:
+// LLVM: call void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}(ptr %[[X_ADDR]], ptr null)
+
+// Outlined target functions: `target parallel` forks an omp.parallel region
+// via __kmpc_fork_call; the proc_bind clause lowers to a __kmpc_push_proc_bind
+// call right before the fork.
+//
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par, ptr {{.*}})
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l
+// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr
+// LLVM: call void @__kmpc_push_proc_bind(ptr {{.*}}, i32 {{.*}}, i32 4)
+// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par, ptr {{.*}})
+
+// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par(
+// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8
+// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4
+// LLVM: call void @use(i32 {{.*}} %[[V]])
+// LLVM: ret void
+
+// OGCG interleaves host wrapper and outlined function per target region.
+//
+// OGCG-LABEL: define {{.*}} void @target_parallel(i32
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_l
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}(
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}.omp_outlined, ptr {{.*}})
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}.omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
+
+// OGCG-LABEL: define {{.*}} void @target_parallel_proc_bind(i32
+// OGCG: call i32 @__tgt_target_kernel(
+// OGCG: omp_offload.failed:
+// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}(
+// OGCG: call void @__kmpc_push_proc_bind(ptr {{.*}}, i32 {{.*}}, i32 4)
+// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}.omp_outlined, ptr {{.*}})
+
+// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}.omp_outlined(
+// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4
+// OGCG: call void @use(i32 {{.*}} %[[V]])
+// OGCG: ret void
More information about the cfe-commits
mailing list