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