https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/229256
>From 4f655336227464ddf2069c073f2030f30b40acb7 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 | 280 +++++++++++++++++- clang/test/CIR/CodeGenOpenMP/for-loop-forms.c | 161 ++++++++++ .../CIR/CodeGenOpenMP/for-loop-review-fixes.c | 123 ++++++++ clang/test/CIR/CodeGenOpenMP/parallel-for.c | 62 ++++ clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c | 207 +++++++++++++ .../CIR/CodeGenOpenMP/target-parallel-for.c | 146 +++++++++ 6 files changed, 971 insertions(+), 8 deletions(-) create mode 100644 clang/test/CIR/CodeGenOpenMP/for-loop-forms.c create mode 100644 clang/test/CIR/CodeGenOpenMP/for-loop-review-fixes.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 9d937f5ed67f90..23d4880e61f17b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -132,6 +132,203 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { }); } +/// 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 that iterates a normalized `[0, tripCount)` range +/// and recomputes the real induction variable each iteration using the given +/// update expression. +static mlir::LogicalResult +emitOMPLoopNest(CIRGenFunction &cgf, const ForStmt &forStmt, mlir::Value lb, + mlir::Value ub, mlir::Value step, const VarDecl *ivDecl, + const Expr *update) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Location loc = cgf.getLoc(forStmt.getSourceRange()); + + auto loopNestOp = mlir::omp::LoopNestOp::create( + builder, loc, /*collapse_num_loops=*/1, lb, ub, step, + /*loop_inclusive=*/false, /*tile_sizes=*/nullptr); + mlir::Block *block = new mlir::Block(); + loopNestOp.getRegion().push_back(block); + block->addArgument(lb.getType(), loc); + + mlir::OpBuilder::InsertionGuard guard(builder); + builder.setInsertionPointToStart(block); + + // Store the normalized counter block argument into the normalized loop + // variable's storage location, then let Sema's update expression recompute + // the real induction variable from it. + mlir::Value iv = block->getArgument(0); + Address ivAddr = cgf.getAddrOfLocalVar(ivDecl); + mlir::Value civVal = + builder.createBuiltinIntCast(loc, iv, ivAddr.getElementType()); + builder.createStore(loc, civVal, ivAddr); + cgf.emitIgnoredExpr(update); + + mlir::LogicalResult bodyRes = mlir::success(); + if (forStmt.getBody()) + if (cgf.emitStmt(forStmt.getBody(), /*useCurrentScope=*/true).failed()) + bodyRes = mlir::failure(); + + mlir::omp::YieldOp::create(builder, cgf.getLoc(forStmt.getEndLoc())); + return bodyRes; +} + +/// Evaluates the clauses allowed on 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); +} + +/// Creates a module-level `omp.private` op (a "privatizer") for a scalar +/// of CIR type elemTy. +static mlir::omp::PrivateClauseOp +createCounterPrivatizer(CIRGenModule &cgm, mlir::Location loc, + llvm::StringRef baseName, mlir::Type elemTy) { + mlir::ModuleOp mod = cgm.getModule(); + std::string name = (baseName + ".privatizer").str(); + for (unsigned suffix = 0; mlir::SymbolTable::lookupSymbolIn(mod, name); + ++suffix) + name = (baseName + ".privatizer." + llvm::Twine(suffix)).str(); + + mlir::OpBuilder b(mod.getContext()); + b.setInsertionPointToEnd(mod.getBody()); + return mlir::omp::PrivateClauseOp::create(b, loc, mlir::TypeRange{}, + b.getStringAttr(name), + mlir::TypeAttr::get(elemTy)); +} + +/// Registers a loop counter as predetermined-private on the omp.wsloop op. +static Address addPrivateCounter(CIRGenFunction &cgf, CIRGenModule &cgm, + mlir::Location loc, + const VarDecl *privateVd, + mlir::omp::WsloopOperands &clauseOps) { + cgf.emitVarDecl(*privateVd); + Address moldAddr = cgf.getAddrOfLocalVar(privateVd); + + mlir::omp::PrivateClauseOp privatizer = createCounterPrivatizer( + cgm, loc, privateVd->getName(), moldAddr.getElementType()); + + clauseOps.privateVars.push_back(moldAddr.getPointer()); + clauseOps.privateSyms.push_back( + mlir::FlatSymbolRefAttr::get(privatizer.getSymNameAttr())); + return moldAddr; +} + +/// Lowers an OMPLoopDirective's `for` leaf to an omp.wsloop + omp.loop_nest. +static mlir::LogicalResult +emitOMPWorksharingLoop(CIRGenFunction &cgf, const OMPLoopDirective &s, + omp::ConstructQueue::const_iterator item) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + CIRGenModule &cgm = cgf.getCIRGenModule(); + mlir::Location loc = cgf.getLoc(s.getBeginLoc()); + + if (mlir::failed(checkSynthesizedClauses(cgf, s, item))) + return mlir::failure(); + + mlir::omp::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()); + + if (emitPreinits(cgf, s.getPreInits()).failed()) + return mlir::failure(); + + // Allocate storage for Sema's normalized 0-based loop counter. + const auto *ivDecl = + cast<VarDecl>(cast<DeclRefExpr>(s.getIterationVariable())->getDecl()); + cgf.emitVarDecl(*ivDecl); + llvm::SmallVector<CIRGenFunction::DeclMapRevertingRAII, 4> counterShadows; + llvm::SmallVector<std::pair<const VarDecl *, Address>, 4> counterMolds; + counterShadows.reserve(s.counters().size()); + counterMolds.reserve(s.counters().size()); + for (auto [counter, privateCounter] : + llvm::zip_equal(s.counters(), s.private_counters())) { + const auto *vd = cast<VarDecl>(cast<DeclRefExpr>(counter)->getDecl()); + const auto *privateVd = + cast<VarDecl>(cast<DeclRefExpr>(privateCounter)->getDecl()); + counterShadows.emplace_back(cgf, vd); + counterMolds.emplace_back( + vd, addPrivateCounter(cgf, cgm, loc, privateVd, clauseOps)); + } + + mlir::Value tripCountCir = cgf.emitScalarExpr(s.getNumIterations()); + auto cirIntType = mlir::cast<cir::IntType>(tripCountCir.getType()); + mlir::Value zero = + 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); + + auto wsloopOp = mlir::omp::WsloopOp::create(builder, loc, clauseOps); + mlir::Block *innerBlock = new mlir::Block(); + // Ensures the block arguments match BlockArgOpenMPOpInterface's expected + // order. + for (auto &counterMold : counterMolds) + innerBlock->addArgument(counterMold.second.getPointer().getType(), loc); + wsloopOp.getRegion().push_back(innerBlock); + + mlir::OpBuilder::InsertionGuard guard(builder); + builder.setInsertionPointToStart(innerBlock); + + // Bind each counter to its private block argument. + for (unsigned idx = 0; idx < counterMolds.size(); ++idx) { + const VarDecl *vd = counterMolds[idx].first; + Address moldAddr = counterMolds[idx].second; + Address privAddr(innerBlock->getArgument(idx), moldAddr.getElementType(), + moldAddr.getAlignment()); + cgf.replaceAddrOfLocalVar(vd, privAddr); + cgf.symbolTable.insert(vd, privAddr.getPointer()); + } + + return emitOMPLoopNest(cgf, *forStmt, zero, tripCount, one, ivDecl, + s.updates()[0]); +} + mlir::LogicalResult CIRGenFunction::emitOMPTaskwaitDirective(const OMPTaskwaitDirective &s) { getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPTaskwaitDirective"); @@ -181,8 +378,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 +414,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 +716,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 00000000000000..9166dca58d9207 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c @@ -0,0 +1,161 @@ +// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s + +void during(int); + +// All the loop forms below normalize to an omp.loop_nest iterating a +// `[0, tripCount)` range, where tripCount is computed from the original +// bounds/step via Sema's helper expression (OMPLoopDirective::getNumIterations()), +// and the user's real induction variable "i" is recomputed each iteration +// from the normalized counter via Sema's update expression +// (OMPLoopDirective::updates()). This is what makes these forms "just work" +// without the CIR codegen having to pattern-match each one individually. + +// 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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) { + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // The update expression recomputes `i = 20 - 1 * iv`. + // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i + // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i + // CHECK: %[[C1:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C1]] : !s32i + // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i + // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during +} + +// Decreasing loop using 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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) { + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // The update expression recomputes `i = 20 - 2 * iv`. + // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i + // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i + // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i + // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i + // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i + // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during +} + +// Decreasing loop using the `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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) { + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // The update expression recomputes `i = 20 - 2 * iv`. + // CHECK: %[[C20:.*]] = cir.const #cir.int<20> : !s32i + // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i + // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i + // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i + // CHECK: %[[SUB:.*]] = cir.sub nsw %[[C20]], %[[MUL]] : !s32i + // CHECK: cir.store align(4) %[[SUB]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during +} + +// Increasing loop using the `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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) { + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // The update expression recomputes `i = 0 + 2 * iv`. + // CHECK: %[[C0:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i + // CHECK: %[[C2:.*]] = cir.const #cir.int<2> : !s32i + // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C2]] : !s32i + // CHECK: %[[ADD:.*]] = cir.add nsw %[[C0]], %[[MUL]] : !s32i + // CHECK: cir.store align(4) %[[ADD]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during +} + +// Increasing loop with the upper bound 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: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[STEP_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[STEP:.*]] = cir.builtin_int_cast %[[STEP_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[STEP]]) { + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // The update expression recomputes `i = 0 + 1 * iv`. + // CHECK: %[[C0:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[IV_RELOAD:.*]] = cir.load align(4) %[[IV_ALLOCA]] : !cir.ptr<!s32i>, !s32i + // CHECK: %[[C1:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[MUL:.*]] = cir.mul nsw %[[IV_RELOAD]], %[[C1]] : !s32i + // CHECK: %[[ADD:.*]] = cir.add nsw %[[C0]], %[[MUL]] : !s32i + // CHECK: cir.store align(4) %[[ADD]], %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during +} diff --git a/clang/test/CIR/CodeGenOpenMP/for-loop-review-fixes.c b/clang/test/CIR/CodeGenOpenMP/for-loop-review-fixes.c new file mode 100644 index 00000000000000..2661c4d7db6392 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/for-loop-review-fixes.c @@ -0,0 +1,123 @@ +// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s + +// Regression tests for https://github.com/llvm/llvm-project/pull/229256 +// review feedback: the worksharing-loop codegen used to pattern-match the +// raw for-statement AST to extract the loop's lower/upper bound and step, +// which (a) evaluated the loop's initializer expression twice, (b) crashed +// on pointer induction variables, and (c) silently produced no diagnostic +// for valid-but-unhandled forms (a pre-declared induction variable used +// with plain assignment, or a `!=` condition). The codegen now iterates a +// normalized `[0, tripCount)` range and recomputes the real induction +// variable via Sema's already-validated update expression, which is +// agnostic to all of these forms. + +void during(int); +int getLB(void); + +// The loop's initializer must be evaluated exactly once, not once by the +// (now-removed) bound-extraction code and again when re-emitting the +// for-statement's init. +void side_effecting_init() { + // CHECK-LABEL: cir.func{{.*}}@{{.*}}side_effecting_init +#pragma omp for + for (int i = getLB(); i < 10; i++) { + during(i); + } + // CHECK: %[[CAPTURE:.*]] = cir.alloca ".capture_expr." align(4) init : !cir.ptr<!s32i> + // CHECK: %[[CALL:.*]] = cir.call @{{.*}}getLB{{.*}}() : {{.*}} -> !s32i + // CHECK-NEXT: cir.store{{.*}} %[[CALL]], %[[CAPTURE]] + // The call result is only ever reloaded afterwards, never called again. + // CHECK-NOT: cir.call @{{.*}}getLB + // CHECK: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) { + // CHECK: omp.loop_nest + // CHECK-NOT: cir.call @{{.*}}getLB + // CHECK: cir.call @{{.*}}during +} + +// Pointer induction variables must not crash: the loop_nest's own bounds +// are always a plain integer trip count, regardless of the induction +// variable's type, and the real pointer is recomputed each iteration via +// pointer arithmetic (cir.ptr_stride), not integer arithmetic. +void pointer_induction_var() { + // CHECK-LABEL: cir.func{{.*}}@{{.*}}pointer_induction_var + int a[10]; +#pragma omp for + for (int *p = a; p < a + 10; ++p) { + during(*p); + } + // "p" is registered as predetermined-private on the wsloop's `private` + // clause: %[[P_ALLOCA]] (never read; the recipe has no init/copy regions) + // is the "mold" operand, and the loop body uses the matching block + // argument instead, giving each thread its own storage. + // CHECK: %[[P_ALLOCA:.*]] = cir.alloca "p" align(8) : !cir.ptr<!cir.ptr<!s32i>> + // CHECK: omp.wsloop private(@{{.*}} %[[P_ALLOCA]] -> %[[P_PRIV:.*]] : !cir.ptr<!cir.ptr<!s32i>>) { + // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i64 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) { + // CHECK: %[[P_BASE:.*]] = cir.load{{.*}} : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i> + // CHECK: %[[P_NEXT:.*]] = cir.ptr_stride %[[P_BASE]], %{{.*}} : (!cir.ptr<!s32i>, !s64i) -> !cir.ptr<!s32i> + // CHECK: cir.store{{.*}} %[[P_NEXT]], %[[P_PRIV]] + // CHECK: cir.call @{{.*}}during +} + +// A pre-declared induction variable (no DeclStmt in the for-statement's +// init, just a plain assignment) combined with a `!=` condition: both are +// valid OpenMP canonical loop forms that must codegen successfully, not +// silently fail with no diagnostic. +void predeclared_var_and_not_equal_cond() { + // CHECK-LABEL: cir.func{{.*}}@{{.*}}predeclared_var_and_not_equal_cond + int i; +#pragma omp for + for (i = 0; i != 10; i++) { + during(i); + } + // "i"'s original alloca (from `int i;`) is unused by the loop itself: + // OpenMP requires loop counters to be predetermined private even when, + // as here, they're pre-declared and could otherwise alias shared storage + // from an enclosing scope (e.g. inside an enclosing omp.parallel region). + // The wsloop's `private` clause is what guarantees that; its "mold" + // operand is a fresh, never-read placeholder (the recipe has no + // init/copy regions) rather than "i"'s original alloca -- using the + // original alloca directly would be invalid whenever this loop sits + // inside an omp.target region, since omp.target (unlike omp.parallel) is + // IsolatedFromAbove and cannot reference values from outside its region. + // The loop body uses the private clause's block argument throughout. + // CHECK: cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_MOLD:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: omp.wsloop private(@{{.*}} %[[I_MOLD]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) { + // CHECK: cir.store{{.*}} %{{.*}}, %[[IV_ALLOCA]] + // CHECK: cir.store{{.*}} %{{.*}}, %[[I_PRIV]] + // CHECK: cir.call @{{.*}}during +} + +// A pre-declared induction variable used inside a combined `parallel for`: +// the counter's storage lives in an enclosing scope that is itself captured +// (by reference, as shared data) into the omp.parallel region, so it must +// not be used directly as per-thread storage for the loop -- the +// omp.wsloop's `private` clause is what guarantees that: the outer "i" +// alloca is never referenced at all here (a fresh, local, never-read mold +// placeholder is used instead -- see predeclared_var_and_not_equal_cond +// above for why), and the real per-thread storage is allocated later, +// inside the per-thread code the shared OpenMPIRBuilder-based lowering +// generates for the (lowered) omp.parallel region, by which point there is +// one execution of the wsloop's private-clause allocation per thread. +// Getting this wrong doesn't show up as a verifier or compile error -- it +// compiles fine but races at runtime, with multiple threads corrupting +// each other's "i". +void predeclared_var_inside_parallel_for() { + // CHECK-LABEL: cir.func{{.*}}@{{.*}}predeclared_var_inside_parallel_for + int i; +#pragma omp parallel for + for (i = 0; i != 10; i++) { + during(i); + } + // CHECK: cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: omp.parallel { + // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_MOLD:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + // CHECK: omp.wsloop private(@{{.*}} %[[I_MOLD]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) { + // CHECK: cir.store{{.*}} %{{.*}}, %[[IV_ALLOCA]] + // CHECK: cir.store{{.*}} %{{.*}}, %[[I_PRIV]] + // CHECK: cir.call @{{.*}}during +} diff --git a/clang/test/CIR/CodeGenOpenMP/parallel-for.c b/clang/test/CIR/CodeGenOpenMP/parallel-for.c new file mode 100644 index 00000000000000..6a87f9f5fd306e --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/parallel-for.c @@ -0,0 +1,62 @@ +// 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 normalized 0-based counter and the real induction variable each get + // their own alloca before the wsloop; "i"'s alloca has no `init` because + // its value is now produced by the update expression below, not by + // directly emitting the for-statement's init. + // CHECK: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i> + // CHECK: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i> + + // The loop bounds are normalized to a `[0, tripCount)` range; tripCount is + // computed from the original bounds/step via Sema's helper expression. + // CHECK: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i + // CHECK: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[LBM1:.*]] = cir.sub nsw %[[C0_CIR]], %[[C1_CIR]] : !s32i + // CHECK: %[[C1B:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[LBM1P1:.*]] = cir.add nsw %[[LBM1]], %[[C1B]] : !s32i + // CHECK: %[[SPAN:.*]] = cir.sub nsw %[[C10_CIR]], %[[LBM1P1]] : !s32i + // CHECK: %[[C1C:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[TRIPCOUNT_CIR:.*]] = cir.div %[[SPAN]], %[[C1C]] : !s32i + // CHECK: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i + // CHECK: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 + // CHECK: %[[ONE_CIR:.*]] = cir.const #cir.int<1> : !s32i + // CHECK: %[[ONE:.*]] = cir.builtin_int_cast %[[ONE_CIR]] : !s32i -> i32 + // CHECK: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + + // "i" is registered as predetermined-private on the wsloop's `private` + // clause: %[[I_ALLOCA]] (never read; the recipe below has no init/copy + // regions) is the "mold" operand, and the loop body uses the matching + // block argument instead, giving each thread its own storage. + // CHECK: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) { + + // The normalized counter block argument is stored into the normalized + // counter's alloca, then Sema's update expression (`i = 0 + 1 * iv`) + // recomputes the real induction variable from it. + // CHECK: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i + // CHECK: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> + // CHECK: cir.call @{{.*}}during + + // CHECK: omp.yield + // CHECK: } + // CHECK: } + // CHECK: omp.terminator + // The parallel is a non-innermost leaf of the combined construct. + // CHECK: } {omp.combined} +} diff --git a/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c b/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c new file mode 100644 index 00000000000000..160f3238358539 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c @@ -0,0 +1,207 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +void before(int); +void during(int); +void after(int); + +void emit_simple_for() { + int j = 5; + before(j); +#pragma omp parallel + { +#pragma omp for + for (int i = 0; i < 10; i++) { + during(j); + } + } + after(j); +} + +// LLVM-LABEL: define{{.*}} void @emit_simple_for() +// LLVM: %[[STRUCT_ARG:.*]] = alloca { ptr }, align 8 +// LLVM: %[[J_SLOT:.*]] = alloca i32, align 4 +// LLVM: store i32 5, ptr %[[J_SLOT]], align 4 +// LLVM: call void @before(i32 +// LLVM: %[[GEP0:.*]] = getelementptr { ptr }, ptr %[[STRUCT_ARG]], i32 0, i32 0 +// LLVM-NEXT: store ptr %[[J_SLOT]], ptr %[[GEP0]], align 8 +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call( +// LLVM-SAME: ptr @{{.*}}, i32 1, ptr @emit_simple_for..omp_par, ptr %[[STRUCT_ARG]]) +// LLVM: %[[J_VAL:.*]] = load i32, ptr %[[J_SLOT]], align 4 +// LLVM-NEXT: call void @after(i32 noundef %[[J_VAL]]) +// LLVM-NEXT: ret void + +// OGCG-LABEL: define{{.*}} void @emit_simple_for() +// OGCG: %[[J:.*]] = alloca i32, align 4 +// OGCG: store i32 5, ptr %[[J]], align 4 +// OGCG: %[[J_VAL:.*]] = load i32, ptr %[[J]], align 4 +// OGCG-NEXT: call void @before(i32 noundef %[[J_VAL]]) +// OGCG-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call( +// OGCG-SAME: ptr @{{.*}}, i32 1, ptr @emit_simple_for.omp_outlined, ptr %[[J]]) +// OGCG: %[[J_VAL2:.*]] = load i32, ptr %[[J]], align 4 +// OGCG-NEXT: call void @after(i32 noundef %[[J_VAL2]]) +// OGCG-NEXT: ret void + +// LLVM-LABEL: define internal void @emit_simple_for..omp_par( +// LLVM-SAME: ptr noalias %tid.addr, ptr noalias %zero.addr, ptr %[[STRUCT:.*]]) +// LLVM: %[[GEP_J:.*]] = getelementptr { ptr }, ptr %[[STRUCT]], i32 0, i32 0 +// LLVM-NEXT: %[[J_PTR:.*]] = load ptr, ptr %[[GEP_J]], align 8 +// LLVM: %p.lastiter = alloca i32, align 4 +// LLVM-NEXT: %p.lowerbound = alloca i32, align 4 +// LLVM-NEXT: %p.upperbound = alloca i32, align 4 +// LLVM-NEXT: %p.stride = alloca i32, align 4 +// LLVM: store i32 0, ptr %p.lowerbound, align 4 +// LLVM-NEXT: store i32 9, ptr %p.upperbound, align 4 +// LLVM-NEXT: store i32 1, ptr %p.stride, align 4 +// LLVM: %[[TID:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr @{{.*}}) +// LLVM-NEXT: call void @__kmpc_for_static_init_4u( +// LLVM-SAME: ptr @{{.*}}, i32 %[[TID]], i32 34, +// LLVM-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.stride, +// LLVM-SAME: i32 1, i32 0) +// LLVM: %omp_loop.iv = phi i32 [ 0, %omp_loop.preheader ], [ %omp_loop.next, %omp_loop.inc ] +// LLVM: %omp_loop.cmp = icmp ult i32 %omp_loop.iv, %{{.*}} +// LLVM-NEXT: br i1 %omp_loop.cmp, label %omp_loop.body, label %omp_loop.exit +// LLVM: omp_loop.exit: +// LLVM-NEXT: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID]]) +// LLVM: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}}) +// LLVM: omp_loop.body: +// LLVM: %[[J_VAL:.*]] = load i32, ptr %[[J_PTR]], align 4 +// LLVM-NEXT: call void @during(i32 noundef %[[J_VAL]]) +// LLVM: %omp_loop.next = add nuw i32 %omp_loop.iv, 1 +// LLVM: ret void + +// OGCG-LABEL: define internal void @emit_simple_for.omp_outlined( +// OGCG-SAME: ptr noalias noundef %.global_tid., ptr noalias noundef %.bound_tid., +// OGCG-SAME: ptr noundef nonnull align 4 dereferenceable(4) %j) +// OGCG: %.global_tid..addr = alloca ptr, align 8 +// OGCG-NEXT: %.bound_tid..addr = alloca ptr, align 8 +// OGCG-NEXT: %j.addr = alloca ptr, align 8 +// OGCG: %.omp.lb = alloca i32, align 4 +// OGCG-NEXT: %.omp.ub = alloca i32, align 4 +// OGCG-NEXT: %.omp.stride = alloca i32, align 4 +// OGCG-NEXT: %.omp.is_last = alloca i32, align 4 +// OGCG: store ptr %.global_tid., ptr %.global_tid..addr, align 8 +// OGCG: store ptr %j, ptr %j.addr, align 8 +// OGCG: %[[J_LOADED:.*]] = load ptr, ptr %j.addr, align 8 +// OGCG: store i32 0, ptr %.omp.lb, align 4 +// OGCG-NEXT: store i32 9, ptr %.omp.ub, align 4 +// OGCG-NEXT: store i32 1, ptr %.omp.stride, align 4 +// OGCG-NEXT: store i32 0, ptr %.omp.is_last, align 4 +// OGCG: %[[GTID_VAL_PTR:.*]] = load ptr, ptr %.global_tid..addr, align 8 +// OGCG-NEXT: %[[TID:.*]] = load i32, ptr %[[GTID_VAL_PTR]], align 4 +// OGCG-NEXT: call void @__kmpc_for_static_init_4( +// OGCG-SAME: ptr @{{.*}}, i32 %[[TID]], i32 34, +// OGCG-SAME: ptr %.omp.is_last, ptr %.omp.lb, ptr %.omp.ub, ptr %.omp.stride, +// OGCG-SAME: i32 1, i32 1) +// OGCG: %[[UB_VAL:.*]] = load i32, ptr %.omp.ub, align 4 +// OGCG-NEXT: %[[CMP:.*]] = icmp sgt i32 %[[UB_VAL]], 9 +// OGCG-NEXT: br i1 %[[CMP]], label %cond.true, label %cond.false +// OGCG: cond.end: +// OGCG: %[[COND:.*]] = phi i32 [ 9, %cond.true ], [ %{{.*}}, %cond.false ] +// OGCG-NEXT: store i32 %[[COND]], ptr %.omp.ub, align 4 +// OGCG: omp.inner.for.cond: +// OGCG: %[[CMP1:.*]] = icmp sle i32 %{{.*}}, %{{.*}} +// OGCG-NEXT: br i1 %[[CMP1]], label %omp.inner.for.body, label %omp.inner.for.end +// OGCG: omp.inner.for.body: +// OGCG: %[[J_VAL:.*]] = load i32, ptr %[[J_LOADED]], align 4 +// OGCG-NEXT: call void @during(i32 noundef %[[J_VAL]]) +// OGCG: omp.loop.exit: +// OGCG: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID]]) +// OGCG-NEXT: call void @__kmpc_barrier(ptr @{{.*}}, i32 %[[TID]]) +// OGCG-NEXT: ret void + +// OGCG: declare void @__kmpc_for_static_init_4(ptr, i32, i32, ptr, ptr, ptr, ptr, i32, i32) +// OGCG: declare void @__kmpc_for_static_fini(ptr, i32) +// OGCG: declare void @__kmpc_barrier(ptr, i32) +// OGCG: declare {{.*}}void @__kmpc_fork_call(ptr, i32, ptr, ...) + +void emit_for_with_vars() { + int j = 5; + before(j); +#pragma omp parallel + { + int lb = 1; + long ub = 10; + short step = 1; +#pragma omp for + for (int i = 0; i < ub; i = i + step) { + during(j); + } + } + after(j); +} + +// LLVM-LABEL: define{{.*}} void @emit_for_with_vars() +// LLVM: %[[STRUCT_ARG:.*]] = alloca { ptr }, align 8 +// LLVM: call void @before(i32 +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call( +// LLVM-SAME: ptr @{{.*}}, i32 1, ptr @emit_for_with_vars..omp_par, ptr %[[STRUCT_ARG]]) +// LLVM: call void @after(i32 +// LLVM: ret void + +// OGCG-LABEL: define{{.*}} void @emit_for_with_vars() +// OGCG: %[[J:.*]] = alloca i32, align 4 +// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call( +// OGCG-SAME: ptr @{{.*}}, i32 1, ptr @emit_for_with_vars.omp_outlined, ptr %[[J]]) +// OGCG: ret void + +// LLVM-LABEL: define internal void @emit_for_with_vars..omp_par( +// LLVM-SAME: ptr noalias %tid.addr, ptr noalias %zero.addr, ptr %[[STRUCT:.*]]) +// LLVM: store i32 1, ptr %{{.*}}, align 4 +// LLVM: store i64 10, ptr %{{.*}}, align 8 +// LLVM: store i16 1, ptr %{{.*}}, align 2 +// LLVM: %{{.*}} = sext i16 %{{.*}} to i32 +// LLVM: %{{.*}} = sext i32 %{{.*}} to i64 +// LLVM: %{{.*}} = sext i16 %{{.*}} to i64 +// LLVM: %{{.*}} = sdiv i64 +// LLVM: %omp_loop.tripcount = select i1 %{{.*}}, i32 0, i32 %{{.*}} +// LLVM: store i32 0, ptr %p.lowerbound, align 4 +// LLVM-NEXT: %[[UB:.*]] = sub i32 %omp_loop.tripcount, 1 +// LLVM-NEXT: store i32 %[[UB]], ptr %p.upperbound, align 4 +// LLVM-NEXT: store i32 1, ptr %p.stride, align 4 +// LLVM: %[[TID2:omp_global_thread_num.*]] = call i32 @__kmpc_global_thread_num(ptr @{{.*}}) +// LLVM-NEXT: call void @__kmpc_for_static_init_4u( +// LLVM-SAME: ptr @{{.*}}, i32 %[[TID2]], i32 34, +// LLVM-SAME: ptr %p.lastiter, ptr %p.lowerbound, ptr %p.upperbound, ptr %p.stride, +// LLVM-SAME: i32 1, i32 0) +// LLVM: %omp_loop.iv = phi i32 +// LLVM: icmp ult i32 %omp_loop.iv, %{{.*}} +// LLVM: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %[[TID2]]) +// LLVM: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}}) +// LLVM: call void @during(i32 +// LLVM: ret void + +// OGCG-LABEL: define internal void @emit_for_with_vars.omp_outlined( +// OGCG-SAME: ptr noalias noundef %.global_tid., ptr noalias noundef %.bound_tid., +// OGCG-SAME: ptr noundef nonnull align 4 dereferenceable(4) %j) +// OGCG: %lb = alloca i32, align 4 +// OGCG-NEXT: %ub = alloca i64, align 8 +// OGCG-NEXT: %step = alloca i16, align 2 +// OGCG: store i32 1, ptr %lb, align 4 +// OGCG-NEXT: store i64 10, ptr %ub, align 8 +// OGCG-NEXT: store i16 1, ptr %step, align 2 +// OGCG: %{{.*}} = sext i16 %{{.*}} to i32 +// OGCG: %{{.*}} = sext i32 %{{.*}} to i64 +// OGCG: %{{.*}} = sext i16 %{{.*}} to i64 +// OGCG: %{{.*}} = sdiv i64 +// OGCG: %[[PRECOND:.*]] = icmp slt i64 0, %{{.*}} +// OGCG-NEXT: br i1 %[[PRECOND]], label %omp.precond.then, label %omp.precond.end +// OGCG: omp.precond.then: +// OGCG: store i32 0, ptr %.omp.lb, align 4 +// OGCG: store i32 1, ptr %.omp.stride, align 4 +// OGCG: call void @__kmpc_for_static_init_4( +// OGCG-SAME: ptr @{{.*}}, i32 %{{.*}}, i32 34, +// OGCG-SAME: i32 1, i32 1) +// OGCG: call void @during(i32 +// OGCG: call void @__kmpc_for_static_fini(ptr @{{.*}}, i32 %{{.*}}) +// OGCG: omp.precond.end: +// OGCG: call void @__kmpc_barrier(ptr @{{.*}}, i32 %{{.*}}) +// OGCG-NEXT: ret void + +// LLVM: declare void @__kmpc_for_static_init_4u(ptr, i32, i32, ptr, ptr, ptr, ptr, i32, i32) +// LLVM: declare void @__kmpc_for_static_fini(ptr, i32) +// LLVM: declare i32 @__kmpc_global_thread_num(ptr) +// LLVM: declare void @__kmpc_barrier(ptr, i32) +// LLVM: declare {{.*}}void @__kmpc_fork_call(ptr, i32, ptr, ...) diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c new file mode 100644 index 00000000000000..6a7cdac9e8faf9 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c @@ -0,0 +1,146 @@ +// 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 { + +// The normalized 0-based counter and the real induction variable each get +// their own alloca before the wsloop; "i"'s alloca has no `init` because +// its value is now produced by the update expression below, not by +// directly emitting the for-statement's init. +// 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: %[[I_ALLOCA]] (never read; the recipe below has no init/copy +// regions) is the "mold" operand, and the loop body uses the matching +// block argument instead, giving each thread its own storage. +// CIR-HOST: omp.wsloop private(@{{.*}} %[[I_ALLOCA]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { +// CIR-HOST-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) { +// CIR-HOST: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i +// CIR-HOST: cir.store align(4) %[[IV_CIR]], %[[IV_ALLOCA]] : !s32i, !cir.ptr<!s32i> +// CIR-HOST: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> +// CIR-HOST: cir.call @{{.*}}during +// CIR-HOST: omp.yield +// CIR-HOST: } +// CIR-HOST: } +// CIR-HOST: omp.terminator +// CIR-HOST: omp.terminator +// CIR-HOST: } + +// CIR-DEVICE: cir.func{{.*}}@target_parallel_for +// CIR-DEVICE: omp.target kernel_type(generic) { +// CIR-DEVICE: omp.parallel { + +// The two allocas and their address-space casts can be emitted in either +// relative order, so match them unordered (CHECK-DAG) rather than assuming +// a specific interleaving. +// CIR-DEVICE-DAG: %[[IV_ALLOCA:.*]] = cir.alloca ".omp.iv" align(4) : !cir.ptr<!s32i, target_address_space(5)> +// CIR-DEVICE-DAG: %[[IV_CAST:.*]] = cir.cast address_space %[[IV_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> +// CIR-DEVICE-DAG: %[[I_ALLOCA:.*]] = cir.alloca "i" align(4) : !cir.ptr<!s32i, target_address_space(5)> +// CIR-DEVICE-DAG: %[[I_CAST:.*]] = cir.cast address_space %[[I_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> +// CIR-DEVICE: %[[C10_CIR:.*]] = cir.const #cir.int<10> : !s32i +// CIR-DEVICE: %[[C0_CIR:.*]] = cir.const #cir.int<0> : !s32i +// CIR-DEVICE: %[[C1_CIR:.*]] = cir.const #cir.int<1> : !s32i +// CIR-DEVICE: %[[TRIPCOUNT_CIR:.*]] = cir.div %{{.*}}, %{{.*}} : !s32i +// CIR-DEVICE: %[[ZERO_CIR:.*]] = cir.const #cir.int<0> : !s32i +// CIR-DEVICE: %[[ZERO:.*]] = cir.builtin_int_cast %[[ZERO_CIR]] : !s32i -> i32 +// CIR-DEVICE: %[[ONE_CIR:.*]] = cir.const #cir.int<1> : !s32i +// CIR-DEVICE: %[[ONE:.*]] = cir.builtin_int_cast %[[ONE_CIR]] : !s32i -> i32 +// CIR-DEVICE: %[[TRIPCOUNT:.*]] = cir.builtin_int_cast %[[TRIPCOUNT_CIR]] : !s32i -> i32 + +// "i" is registered as predetermined-private on the wsloop's `private` +// clause: %[[I_CAST]] (never read; the recipe below has no init/copy +// regions) is the "mold" operand, and the loop body uses the matching +// block argument instead, giving each thread its own storage. +// CIR-DEVICE: omp.wsloop private(@{{.*}} %[[I_CAST]] -> %[[I_PRIV:.*]] : !cir.ptr<!s32i>) { +// CIR-DEVICE-NEXT: omp.loop_nest (%[[IV:.*]]) : i32 = (%[[ZERO]]) to (%[[TRIPCOUNT]]) step (%[[ONE]]) { +// CIR-DEVICE: %[[IV_CIR:.*]] = cir.builtin_int_cast %[[IV]] : i32 -> !s32i +// CIR-DEVICE: cir.store align(4) %[[IV_CIR]], %[[IV_CAST]] : !s32i, !cir.ptr<!s32i> +// CIR-DEVICE: cir.store align(4) %{{.*}}, %[[I_PRIV]] : !s32i, !cir.ptr<!s32i> +// CIR-DEVICE: cir.call @{{.*}}during +// CIR-DEVICE: omp.yield +// CIR-DEVICE: } +// CIR-DEVICE: } +// CIR-DEVICE: omp.terminator +// CIR-DEVICE: omp.terminator +// CIR-DEVICE: } + +// The combined `target parallel for` directive decomposes into `target`, +// `parallel` and `for` leaves and lowers to the same nesting as the explicit +// target/parallel/for above: an omp.wsloop + omp.loop_nest inside omp.parallel +// inside omp.target. +void combined_target_parallel_for() { +#pragma omp target parallel for + for (int i = 0; i < 10; i++) { + during(i); + } +} + +// 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) : !cir.ptr<!s32i> +// CIR-HOST: omp.wsloop private(@{{.*}} %[[CI_ALLOCA]] -> %{{.*}} : !cir.ptr<!s32i>) { +// CIR-HOST-NEXT: omp.loop_nest (%[[CIV:.*]]) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) { +// CIR-HOST: cir.call @{{.*}}during +// CIR-HOST: omp.yield +// CIR-HOST: } +// CIR-HOST: } +// CIR-HOST: omp.terminator +// CIR-HOST: } {omp.combined} +// CIR-HOST: omp.terminator +// CIR-HOST: } {omp.combined} + +// CIR-DEVICE: cir.func{{.*}}@combined_target_parallel_for +// CIR-DEVICE: omp.target kernel_type(generic) { +// CIR-DEVICE: omp.parallel { +// CIR-DEVICE: omp.wsloop private(@{{.*}} %{{.*}} -> %{{.*}} : !cir.ptr<!s32i>) { +// CIR-DEVICE-NEXT: omp.loop_nest (%{{.*}}) : i32 = (%{{.*}}) to (%{{.*}}) step (%{{.*}}) { +// CIR-DEVICE: cir.call @{{.*}}during +// CIR-DEVICE: omp.yield +// CIR-DEVICE: } +// CIR-DEVICE: } +// CIR-DEVICE: omp.terminator +// CIR-DEVICE: } {omp.combined} +// CIR-DEVICE: omp.terminator +// CIR-DEVICE: } {omp.combined} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
