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 <[email protected]> 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 (%{{.*}}) { _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
