[llvm-branch-commits] [clang] [CIR][OpenMP] Add support for host_eval so that SPMD kernels can be used (PR #229259)
Jan Leyonberg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Oct 9 12:46:06 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/229259
>From 3a1598ce31191c1c645ef1237fc06ba4780d9c9a Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 7 Oct 2026 13:13:03 -0400
Subject: [PATCH] [CIR][OpenMP] Add support for host_eval so that SPMD kernels
can be used
This patch adds support for host_eval so that SPMD kernels and in the future
num_threads etc. can be implemented correctly.
Assisted-by: Cursor / Claude Sonnet 5 High
---
clang/lib/CIR/CodeGen/CIRGenFunction.h | 12 ++
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 165 ++++++++++++++++--
.../CIR/CodeGenOpenMP/target-parallel-for.c | 51 +++---
3 files changed, 190 insertions(+), 38 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index c6427d60a37926..0c040a2da566d8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -2753,6 +2753,18 @@ class CIRGenFunction : public CIRGenTypeCache {
// OpenMP Emission
//===--------------------------------------------------------------------===//
public:
+ /// The enclosing omp.target's host-evaluated loop bounds -- a normalized
+ /// [0, tripCount) range matching emitOMPLoopNest's convention -- forwarded
+ /// as host_eval block arguments and consumed once by the nested
+ /// omp.wsloop.
+ struct OMPHostEvalBounds {
+ mlir::Value zero;
+ mlir::Value tripCount;
+ mlir::Value one;
+ bool applied = false;
+ };
+ std::optional<OMPHostEvalBounds> ompHostEvalBounds;
+
mlir::LogicalResult emitOMPScopeDirective(const OMPScopeDirective &s);
mlir::LogicalResult emitOMPErrorDirective(const OMPErrorDirective &s);
mlir::LogicalResult emitOMPParallelDirective(const OMPParallelDirective &s);
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index f94467456950b8..8afe97370f1f9b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -51,6 +51,68 @@ checkSynthesizedClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
return res;
}
+/// Returns \p s's nested OpenMP directive -- its body after unwrapping any
+/// single-statement compounds -- or null if the body isn't itself a single
+/// OpenMP directive.
+static const OMPExecutableDirective *
+getSingleNestedOMPDirective(const OMPExecutableDirective &s) {
+ const Stmt *body =
+ s.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
+ /*IgnoreCaptured=*/true);
+ return dyn_cast<OMPExecutableDirective>(body);
+}
+
+/// Returns true if \p dir is a leaf construct kind that participates in
+/// CIR's "combined" op marking (see hasCombinableNestedLeaf).
+static bool isCombinableLeaf(llvm::omp::Directive dir) {
+ switch (dir) {
+ case llvm::omp::OMPD_target:
+ case llvm::omp::OMPD_parallel:
+ case llvm::omp::OMPD_for:
+ return true;
+ default:
+ return false;
+ }
+}
+
+/// Returns true if \p s's body is itself an OpenMP directive that combines
+/// with \p s for codegen purposes, e.g. explicitly-nested `target` +
+/// `parallel` written as separate pragmas rather than `target parallel`.
+static bool hasCombinableNestedLeaf(const OMPExecutableDirective &s) {
+ const OMPExecutableDirective *nested = getSingleNestedOMPDirective(s);
+ if (!nested)
+ return false;
+ return isCombinableLeaf(
+ llvm::omp::getLeafConstructsOrSelf(nested->getDirectiveKind()).front());
+}
+
+/// Finds the loop directive nested (possibly transitively) inside \p s.
+static const OMPLoopDirective *
+findNestedOMPLoopDirective(const OMPExecutableDirective &s) {
+ const OMPExecutableDirective *nested = getSingleNestedOMPDirective(s);
+ if (!nested)
+ return nullptr;
+ if (const auto *loopDir = dyn_cast<OMPLoopDirective>(nested))
+ return loopDir;
+ return findNestedOMPLoopDirective(*nested);
+}
+
+/// Returns true if \p s is a target SPMD construct: a combined
+/// `target parallel for`, or an explicitly nested `target` whose body is a
+/// parallel directive.
+static bool isTargetSPMD(const OMPExecutableDirective &s) {
+ switch (s.getDirectiveKind()) {
+ case llvm::omp::OMPD_target_parallel_for:
+ return true;
+ case llvm::omp::OMPD_target: {
+ const OMPExecutableDirective *nested = getSingleNestedOMPDirective(s);
+ return nested && isOpenMPParallelDirective(nested->getDirectiveKind());
+ }
+ default:
+ return false;
+ }
+}
+
static mlir::LogicalResult
emitParallelClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
CIRGenBuilderTy &builder, mlir::Location loc,
@@ -80,7 +142,8 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
CIRGenModule &cgm = cgf.getCIRGenModule();
auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
- if (!omp::isLastItemInQueue(item, queue))
+ if (!omp::isLastItemInQueue(item, queue) ||
+ hasCombinableNestedLeaf(static_cast<const OMPExecutableDirective &>(s)))
parallelOp.setCombined(true);
mlir::Block &block = parallelOp.getRegion().emplaceBlock();
@@ -255,6 +318,28 @@ static Address addPrivateCounter(CIRGenFunction &cgf, CIRGenModule &cgm,
return moldAddr;
}
+/// Computes a target SPMD loop's normalized `[0, tripCount)` bounds at the
+/// host insertion point, as builtin-integer values for the enclosing
+/// omp.target's host_eval operands. All three are forwarded even though
+/// only the trip count is runtime-dependent, since the dialect requires an
+/// SPMD loop_nest's bounds to be host-evaluated in full.
+static std::optional<CIRGenFunction::OMPHostEvalBounds>
+emitHostEvalLoopBounds(CIRGenFunction &cgf, const OMPLoopDirective &s) {
+ if (emitPreinits(cgf, s.getPreInits()).failed())
+ return std::nullopt;
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(s.getBeginLoc());
+ mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
+ auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
+ CIRGenFunction::OMPHostEvalBounds hev;
+ hev.zero =
+ cirIntToBuiltinInt(builder, loc, builder.getConstInt(loc, cirIntType, 0));
+ hev.tripCount = cirIntToBuiltinInt(builder, loc, tripCountCir);
+ hev.one =
+ cirIntToBuiltinInt(builder, loc, builder.getConstInt(loc, cirIntType, 1));
+ return hev;
+}
+
/// Lowers an OMPLoopDirective's `for` leaf to an omp.wsloop + omp.loop_nest.
static mlir::LogicalResult
emitOMPWorksharingLoop(CIRGenFunction &cgf, const OMPLoopDirective &s,
@@ -274,7 +359,12 @@ emitOMPWorksharingLoop(CIRGenFunction &cgf, const OMPLoopDirective &s,
const CapturedStmt *capturedStmt = s.getInnermostCapturedStmt();
const auto *forStmt = cast<ForStmt>(capturedStmt->getCapturedStmt());
- if (emitPreinits(cgf, s.getPreInits()).failed())
+ // True when the enclosing omp.target already computed this loop's trip
+ // count (see emitHostEvalLoopBounds) and forwarded it via
+ // cgf.ompHostEvalBounds.
+ bool consumingHostEval =
+ cgf.ompHostEvalBounds && !cgf.ompHostEvalBounds->applied;
+ if (!consumingHostEval && emitPreinits(cgf, s.getPreInits()).failed())
return mlir::failure();
// Allocate storage for Sema's normalized 0-based loop counter.
@@ -295,13 +385,22 @@ emitOMPWorksharingLoop(CIRGenFunction &cgf, const OMPLoopDirective &s,
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 =
- cirIntToBuiltinInt(builder, loc, builder.getConstInt(loc, cirIntType, 0));
- mlir::Value one =
- cirIntToBuiltinInt(builder, loc, builder.getConstInt(loc, cirIntType, 1));
- mlir::Value tripCount = cirIntToBuiltinInt(builder, loc, tripCountCir);
+ mlir::Value zero, tripCount, one;
+ if (consumingHostEval) {
+ // Consume the host_eval bounds forwarded by the enclosing omp.target.
+ cgf.ompHostEvalBounds->applied = true;
+ zero = cgf.ompHostEvalBounds->zero;
+ tripCount = cgf.ompHostEvalBounds->tripCount;
+ one = cgf.ompHostEvalBounds->one;
+ } else {
+ mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
+ auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
+ zero = cirIntToBuiltinInt(builder, loc,
+ builder.getConstInt(loc, cirIntType, 0));
+ one = cirIntToBuiltinInt(builder, loc,
+ builder.getConstInt(loc, cirIntType, 1));
+ tripCount = cirIntToBuiltinInt(builder, loc, tripCountCir);
+ }
auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, clauseOps);
mlir::Block *innerBlock = new mlir::Block();
@@ -572,15 +671,37 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s,
if (mlir::failed(emitOMPTargetImplicitCaptures(cgf, s, mapSyms)))
return mlir::failure();
- // Use generic for now.
+ const auto &execDir = static_cast<const OMPExecutableDirective &>(s);
+ bool isSPMD = isTargetSPMD(execDir);
+
+ std::optional<CIRGenFunction::OMPHostEvalBounds> hostEval;
+ if (isSPMD && !cgf.getCIRGenModule().getLangOpts().OpenMPIsTargetDevice) {
+ const auto *loopDir = dyn_cast<OMPLoopDirective>(&execDir);
+ if (!loopDir)
+ loopDir = findNestedOMPLoopDirective(execDir);
+ if (!loopDir || !(hostEval = emitHostEvalLoopBounds(cgf, *loopDir))) {
+ cgf.getCIRGenModule().errorNYI(
+ s.getSourceRange(), "OpenMP target host-evaluated loop bounds");
+ return mlir::failure();
+ }
+ clauseOps.hostEvalVars.push_back(hostEval->zero);
+ clauseOps.hostEvalVars.push_back(hostEval->tripCount);
+ clauseOps.hostEvalVars.push_back(hostEval->one);
+ }
+
clauseOps.kernelType = mlir::omp::TargetExecModeAttr::get(
- &cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic);
+ &cgf.getMLIRContext(), isSPMD ? mlir::omp::TargetExecMode::spmd
+ : mlir::omp::TargetExecMode::generic);
auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
- if (!omp::isLastItemInQueue(item, queue))
+ if (!omp::isLastItemInQueue(item, queue) || hasCombinableNestedLeaf(execDir))
targetOp.setCombined(true);
+ // Block arguments must be added in the order BlockArgOpenMPOpInterface
+ // expects: host_eval arguments precede map arguments.
mlir::Block &block = targetOp.getRegion().emplaceBlock();
+ for (mlir::Value hostEvalVar : clauseOps.hostEvalVars)
+ block.addArgument(hostEvalVar.getType(), begin);
for (mlir::Value mapVar : clauseOps.mapVars)
block.addArgument(mapVar.getType(), begin);
@@ -589,17 +710,33 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s,
CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+ auto argIface = mlir::cast<mlir::omp::BlockArgOpenMPOpInterface>(*targetOp);
+ llvm::MutableArrayRef<mlir::BlockArgument> mapBlockArgs =
+ argIface.getMapBlockArgs();
llvm::SmallVector<std::pair<const VarDecl *, Address>> savedAddrs;
for (auto [idx, vd] : llvm::enumerate(mapSyms)) {
Address origAddr = cgf.getAddrOfLocalVar(vd);
savedAddrs.push_back({vd, origAddr});
- mlir::Value blockArg = block.getArgument(idx);
- cgf.replaceAddrOfLocalVar(vd, Address(blockArg, origAddr.getAlignment()));
+ cgf.replaceAddrOfLocalVar(
+ vd, Address(mapBlockArgs[idx], origAddr.getAlignment()));
+ }
+
+ // Forward the host_eval block arguments to the nested loop's emission.
+ std::optional<CIRGenFunction::OMPHostEvalBounds> savedHostEvalBounds =
+ std::move(cgf.ompHostEvalBounds);
+ cgf.ompHostEvalBounds.reset();
+ if (hostEval) {
+ llvm::MutableArrayRef<mlir::BlockArgument> hostEvalBlockArgs =
+ argIface.getHostEvalBlockArgs();
+ cgf.ompHostEvalBounds = CIRGenFunction::OMPHostEvalBounds{
+ hostEvalBlockArgs[0], hostEvalBlockArgs[1], hostEvalBlockArgs[2],
+ /*applied=*/false};
}
mlir::LogicalResult res = emitBody();
mlir::omp::TerminatorOp::create(builder, end);
+ cgf.ompHostEvalBounds = std::move(savedHostEvalBounds);
for (auto &[vd, addr] : savedAddrs)
cgf.replaceAddrOfLocalVar(vd, addr);
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
index 2295beb2b3e6d8..5223f2c6a70f63 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
@@ -15,6 +15,9 @@ void during(int);
// 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.
+//
+// This is a target SPMD construct, so the host also computes the loop's
+// trip count and forwards it via host_eval block arguments.
void target_parallel_for() {
#pragma omp target
#pragma omp parallel
@@ -25,7 +28,17 @@ void target_parallel_for() {
}
// CIR-HOST: cir.func{{.*}}@target_parallel_for
-// CIR-HOST: omp.target kernel_type(generic) {
+
+// The trip count is computed from the original bounds/step via Sema's
+// helper expression, then 0, the trip count, and 1 are each cast to a
+// builtin integer and forwarded as the omp.target's host_eval operands.
+// 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: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_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: omp.target kernel_type(spmd) host_eval(%[[ZERO]] -> %[[AZERO:.*]], %[[TRIPCOUNT]] -> %[[ATRIPCOUNT:.*]], %[[ONE]] -> %[[AONE:.*]] : i32, i32, i32) {
// CIR-HOST: omp.parallel {
// The normalized 0-based counter and the real induction variable each get
@@ -35,22 +48,10 @@ void target_parallel_for() {
// 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 via Sema's helper expression.
-// 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
-
// "i" is registered as predetermined-private on the wsloop's `private`
// clause (see parallel-for.c for why %[[I_ALLOCA]] is never read).
// 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-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[AZERO]]) to (%[[ATRIPCOUNT]]) step (%[[AONE]]) {
// 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>
@@ -63,7 +64,7 @@ void target_parallel_for() {
// CIR-HOST: }
// CIR-DEVICE: cir.func{{.*}}@target_parallel_for
-// CIR-DEVICE: omp.target kernel_type(generic) {
+// CIR-DEVICE: omp.target kernel_type(spmd) {
// CIR-DEVICE: omp.parallel {
// The two allocas and their address-space casts can be emitted in either
@@ -73,9 +74,9 @@ void target_parallel_for() {
// 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
+
+// The device recomputes the same normalized `[0, tripCount)` bounds itself
+// (no host_eval here -- see the file comment above).
// 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
@@ -101,7 +102,9 @@ void target_parallel_for() {
// 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.
+// inside omp.target. This is a target SPMD construct, so the omp.target is
+// marked kernel_type(spmd) with host_eval'd loop bounds, just like the
+// explicitly nested directives above.
void combined_target_parallel_for() {
#pragma omp target parallel for
for (int i = 0; i < 10; i++) {
@@ -110,14 +113,14 @@ void combined_target_parallel_for() {
}
// 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).
+// construct, so both carry the omp.combined attribute, just like the
+// explicitly nested directives above.
// CIR-HOST: cir.func{{.*}}@combined_target_parallel_for
-// CIR-HOST: omp.target kernel_type(generic) {
+// CIR-HOST: omp.target kernel_type(spmd) host_eval(%[[ZERO:.*]] -> %[[AZERO:.*]], %[[TRIPCOUNT:.*]] -> %[[ATRIPCOUNT:.*]], %[[ONE:.*]] -> %[[AONE:.*]] : i32, i32, i32) {
// 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-NEXT: omp.loop_nest (%[[CIV:.*]]) : i32 = (%[[AZERO]]) to (%[[ATRIPCOUNT]]) step (%[[AONE]]) {
// CIR-HOST: cir.call @{{.*}}during
// CIR-HOST: omp.yield
// CIR-HOST: }
@@ -128,7 +131,7 @@ void combined_target_parallel_for() {
// CIR-HOST: } {omp.combined}
// CIR-DEVICE: cir.func{{.*}}@combined_target_parallel_for
-// CIR-DEVICE: omp.target kernel_type(generic) {
+// CIR-DEVICE: omp.target kernel_type(spmd) {
// CIR-DEVICE: omp.parallel {
// CIR-DEVICE: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) {
// CIR-DEVICE-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) {
More information about the llvm-branch-commits
mailing list