[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 16 08:03:36 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/229259
>From 2052c5695a7a7d8337b663dbe174c5ef493df145 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 and other clauses can be implemented correctly.
Assisted-by: Cursor / Claude Sonnet 5 High
---
clang/lib/CIR/CodeGen/CIRGenFunction.h | 26 +++
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 164 ++++++++++++++++--
.../target-parallel-for-llvm-device.c | 4 +-
.../target-parallel-for-llvm-host.c | 11 ++
.../CIR/CodeGenOpenMP/target-parallel-for.c | 42 ++---
5 files changed, 211 insertions(+), 36 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 8be794a660224..10a834d15baac 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -2811,6 +2811,32 @@ 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;
+
+ /// An RAII object to set (and then restore) ompHostEvalBounds for the
+ /// extent of a nested loop directive's emission.
+ struct OMPHostEvalBoundsRAII {
+ CIRGenFunction &cgf;
+ std::optional<OMPHostEvalBounds> oldBounds;
+
+ OMPHostEvalBoundsRAII(CIRGenFunction &cgf,
+ std::optional<OMPHostEvalBounds> newBounds)
+ : cgf(cgf), oldBounds(std::move(cgf.ompHostEvalBounds)) {
+ cgf.ompHostEvalBounds = std::move(newBounds);
+ }
+ ~OMPHostEvalBoundsRAII() { cgf.ompHostEvalBounds = std::move(oldBounds); }
+ };
+
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 71e943085bcde..9dffa8f2846ee 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();
@@ -273,6 +336,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 = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 0));
+ hev.tripCount = builder.createBuiltinIntCast(loc, tripCountCir);
+ hev.one = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 1));
+ return hev;
+}
+
/// Lowers an OMPLoopDirective's `for` directive to an omp.wsloop +
/// omp.loop_nest.
static mlir::LogicalResult
@@ -293,7 +378,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.
@@ -314,13 +404,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 = 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);
+ mlir::Value zero, tripCount, one;
+ if (consumingHostEval) {
+ // Consume the host_eval bounds forwarded by the enclosing omp.target.
+ cgf.ompHostEvalBounds->applied = true;
+ zero = cgf.ompHostEvalBounds->zero;
+ tripCount = cgf.ompHostEvalBounds->tripCount;
+ one = cgf.ompHostEvalBounds->one;
+ } else {
+ mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations());
+ auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType());
+ zero = builder.createBuiltinIntCast(
+ loc, builder.getConstInt(loc, cirIntType, 0));
+ one = builder.createBuiltinIntCast(loc,
+ builder.getConstInt(loc, cirIntType, 1));
+ tripCount = builder.createBuiltinIntCast(loc, tripCountCir);
+ }
auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, clauseOps);
mlir::Block *innerBlock = new mlir::Block();
@@ -582,15 +681,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);
@@ -599,13 +720,28 @@ 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> newHostEvalBounds;
+ if (hostEval) {
+ llvm::MutableArrayRef<mlir::BlockArgument> hostEvalBlockArgs =
+ argIface.getHostEvalBlockArgs();
+ newHostEvalBounds = CIRGenFunction::OMPHostEvalBounds{
+ hostEvalBlockArgs[0], hostEvalBlockArgs[1], hostEvalBlockArgs[2],
+ /*applied=*/false};
}
+ CIRGenFunction::OMPHostEvalBoundsRAII hostEvalGuard(
+ cgf, std::move(newHostEvalBounds));
mlir::LogicalResult res = emitBody();
mlir::omp::TerminatorOp::create(builder, end);
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c
index 428174a1d3bf2..694534331019c 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c
@@ -37,7 +37,7 @@ void combined_target_parallel_for() {
// 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_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null, ptr {{.*}}, i64 0, i32 0)
// LLVM: call void @__kmpc_target_deinit()
// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par
@@ -50,7 +50,7 @@ void combined_target_parallel_for() {
// 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_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par{{.*}}, ptr null, 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
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
index 4142859585192..d3bac47d5e0d2 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c
@@ -28,12 +28,19 @@ void combined_target_parallel_for() {
// __kmpc_fork_call, whose body statically schedules the omp.wsloop via
// __kmpc_for_static_init_4/_4u.
//
+// This is a target SPMD construct, so the host's kernel-launch arguments
+// carry the host_eval'd trip count (10) in the Tripcount field; OGCG below
+// leaves that field 0, since it doesn't forward a host-computed trip count.
// LLVM-LABEL: define {{.*}} void @target_parallel_for(
+// LLVM: %[[TRIPCOUNT_GEP:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM: store i64 10, ptr %[[TRIPCOUNT_GEP]], align 8
// 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: %[[TRIPCOUNT_GEP:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// LLVM: store i64 10, ptr %[[TRIPCOUNT_GEP]], align 8
// LLVM: call i32 @__tgt_target_kernel(
// LLVM: omp_offload.failed:
// LLVM: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}(
@@ -57,6 +64,8 @@ void combined_target_parallel_for() {
// LLVM: call void @during(i32
// OGCG-LABEL: define {{.*}} void @target_parallel_for(
+// OGCG: %[[TRIPCOUNT_GEP:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// OGCG: store i64 0, ptr %[[TRIPCOUNT_GEP]], align 8
// OGCG: call i32 @__tgt_target_kernel(
// OGCG: omp_offload.failed:
// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}(
@@ -70,6 +79,8 @@ void combined_target_parallel_for() {
// OGCG: call void @__kmpc_for_static_fini(
// OGCG-LABEL: define {{.*}} void @combined_target_parallel_for(
+// OGCG: %[[TRIPCOUNT_GEP:.*]] = getelementptr inbounds nuw %struct.__tgt_kernel_arguments, ptr %kernel_args, i32 0, i32 8
+// OGCG: store i64 0, ptr %[[TRIPCOUNT_GEP]], align 8
// OGCG: call i32 @__tgt_target_kernel(
// OGCG: omp_offload.failed:
// OGCG: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}(
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
index d9ba201a312b5..812ae3c79e864 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c
@@ -11,6 +11,8 @@
void during(int);
+// 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
@@ -21,26 +23,24 @@ void target_parallel_for() {
}
// 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
+// The trip count is computed from the original bounds/step, 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: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_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 {
+
+// 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>
// 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>
@@ -53,7 +53,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
@@ -63,9 +63,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
@@ -89,7 +89,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++) {
@@ -98,11 +100,11 @@ void combined_target_parallel_for() {
}
// 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: }
@@ -113,7 +115,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