https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/229256
>From c606e2254e71db763358471aedb64535dc590d06 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 --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 7 + clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 381 +++++++++++++++--- clang/test/CIR/CodeGenOpenMP/for-loop-forms.c | 228 +++++++++++ clang/test/CIR/CodeGenOpenMP/parallel-for.c | 54 +++ clang/test/CIR/CodeGenOpenMP/pragma-omp-for.c | 207 ++++++++++ .../target-parallel-for-llvm-device.c | 84 ++++ .../target-parallel-for-llvm-host.c | 83 ++++ .../CIR/CodeGenOpenMP/target-parallel-for.c | 127 ++++++ .../target-parallel-llvm-device.c | 85 ++++ .../CodeGenOpenMP/target-parallel-llvm-host.c | 94 +++++ 10 files changed, 1283 insertions(+), 67 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-llvm-device.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-for.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index 36d583cfe9fbed..f4dbbc58fc7eed 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -673,6 +673,13 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return createBuiltinIntCast(src.getLoc(), src, newTy); } + /// Casts \p src to the builtin MLIR integer type with the same width as + /// its own CIR integer type. + mlir::Value createBuiltinIntCast(mlir::Location loc, mlir::Value src) { + auto cirIntTy = mlir::cast<cir::IntType>(src.getType()); + return createBuiltinIntCast(loc, src, getIntegerType(cirIntTy.getWidth())); + } + mlir::Value createIntToPtr(mlir::Value src, mlir::Type newTy) { return createCast(cir::CastKind::int_to_ptr, src, newTy); } diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp index 9d937f5ed67f90..71e943085bcde3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -105,6 +105,27 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, return res; } +/// Checks, emits clauses for, and emits an omp.parallel for the `parallel` +/// directive at \p parallelItem, deferring to \p emitBody for the region's +/// contents. +template <typename DirectiveTy> +static mlir::LogicalResult emitParallel( + CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue, + omp::ConstructQueue::const_iterator parallelItem, mlir::Location begin, + mlir::Location end, llvm::function_ref<mlir::LogicalResult()> emitBody) { + if (mlir::failed(checkSynthesizedClauses(cgf, s, parallelItem))) + return mlir::failure(); + + mlir::omp::ParallelOperands parallelOps; + if (mlir::failed(emitParallelClauses(cgf, cgf.getCIRGenModule(), + cgf.getBuilder(), begin, + parallelItem->clauses, parallelOps))) + return mlir::failure(); + + return emitParallelOp(cgf, s, queue, parallelItem, begin, end, parallelOps, + emitBody); +} + mlir::LogicalResult CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { mlir::Location begin = getLoc(s.getBeginLoc()); @@ -114,22 +135,216 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s); omp::ConstructQueue::const_iterator item = queue.begin(); - if (mlir::failed(checkSynthesizedClauses(*this, s, item))) + return emitParallel( + *this, s, queue, item, begin, end, [&]() -> mlir::LogicalResult { + // emitStmt() rejects CapturedStmt directly; the + // parent construct must unwrap it, so emit the + // inner statement instead. + const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel); + return emitStmt(cs->getCapturedStmt(), + /*useCurrentScope=*/true); + }); +} + +/// Emits a single Sema-generated pre-init statement for an OpenMP loop +/// directive: a VarDecl's DeclStmt is lowered directly (it has no +/// initializer CIRGenFunction::emitStmt can handle on its own), everything +/// else goes through the general statement emitter. +static mlir::LogicalResult emitPreinit(CIRGenFunction &cgf, + const Stmt *preInit) { + if (const auto *declStmt = dyn_cast<DeclStmt>(preInit)) { + for (const Decl *d : declStmt->decls()) + cgf.emitVarDecl(cast<VarDecl>(*d)); + return mlir::success(); + } + return cgf.emitStmt(preInit, /*useCurrentScope=*/true); +} + +/// 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(); + + if (const auto *compound = dyn_cast<CompoundStmt>(preInits)) { + for (const Stmt *stmt : compound->body()) + if (emitPreinit(cgf, stmt).failed()) + return mlir::failure(); + return mlir::success(); + } + + return emitPreinit(cgf, preInits); +} + +/// 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 the `for` directive; 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) { + // getGlobalValue() is an O(1) cache lookup, unlike + // mlir::SymbolTable::lookupSymbolIn()'s O(n) module scan. + unsigned uniquingCounter = 0; + llvm::SmallString<64> name = mlir::SymbolTable::generateSymbolName<64>( + (baseName + ".privatizer").str(), + [&](llvm::StringRef candidate) { + return cgm.getGlobalValue(candidate) != nullptr; + }, + uniquingCounter); + + mlir::ModuleOp mod = cgm.getModule(); + mlir::OpBuilder b(mod.getContext()); + b.setInsertionPointToEnd(mod.getBody()); + auto privatizer = mlir::omp::PrivateClauseOp::create( + b, loc, mlir::TypeRange{}, b.getStringAttr(name), + mlir::TypeAttr::get(elemTy)); + cgm.insertGlobalSymbol(privatizer); + return privatizer; +} + +/// 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` directive 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::ParallelOperands clauseOps; - if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin, - item->clauses, clauseOps))) + mlir::omp::WsloopOperands clauseOps; + if (emitWsloopClauses(cgf, cgm, builder, loc, item->clauses, clauseOps) + .failed()) return mlir::failure(); - return emitParallelOp( - *this, s, queue, item, begin, end, clauseOps, - [&]() -> mlir::LogicalResult { - // emitStmt() rejects CapturedStmt directly; the parent construct - // must unwrap it, so emit the inner statement instead. - const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel); - return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true); - }); + 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 = 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); + + 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 @@ -181,8 +396,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 +432,22 @@ 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 directive"); + + return emitParallel( + *this, s, queue, parallelItem, begin, end, [&]() -> mlir::LogicalResult { + omp::ConstructQueue::const_iterator forItem = std::next(parallelItem); + assert(forItem != queue.end() && forItem->id == llvm::omp::OMPD_for && + "expected a 'for' directive nested in 'parallel'"); + return emitOMPWorksharingLoop(*this, s, forItem); + }); } mlir::LogicalResult CIRGenFunction::emitOMPParallelForSimdDirective( const OMPParallelForSimdDirective &s) { @@ -387,6 +616,28 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, return res; } +/// Checks, emits clauses for, and emits an omp.target for the `target` +/// directive at \p targetItem, deferring to \p emitBody for the region's +/// contents. +template <typename DirectiveTy> +static mlir::LogicalResult emitTarget( + CIRGenFunction &cgf, const DirectiveTy &s, const omp::ConstructQueue &queue, + omp::ConstructQueue::const_iterator targetItem, mlir::Location begin, + mlir::Location end, llvm::function_ref<mlir::LogicalResult()> emitBody) { + if (mlir::failed(checkSynthesizedClauses(cgf, s, targetItem))) + return mlir::failure(); + + mlir::omp::TargetExtOperands targetOps; + llvm::SmallVector<const VarDecl *> mapSyms; + if (mlir::failed(emitTargetClauses(cgf, cgf.getCIRGenModule(), + cgf.getBuilder(), begin, + targetItem->clauses, targetOps, mapSyms))) + return mlir::failure(); + + return emitTargetOp(cgf, s, queue, targetItem, begin, end, targetOps, mapSyms, + emitBody); +} + mlir::LogicalResult CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) { mlir::Location begin = getLoc(s.getBeginLoc()); @@ -396,22 +647,12 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) { omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s); omp::ConstructQueue::const_iterator item = queue.begin(); - if (mlir::failed(checkSynthesizedClauses(*this, s, item))) - return mlir::failure(); - - mlir::omp::TargetExtOperands clauseOps; - llvm::SmallVector<const VarDecl *> mapSyms; - if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin, - item->clauses, clauseOps, mapSyms))) - return mlir::failure(); - - return emitTargetOp(*this, s, queue, item, begin, end, clauseOps, mapSyms, - [&]() -> mlir::LogicalResult { - const CapturedStmt *cs = - s.getCapturedStmt(llvm::omp::OMPD_target); - return emitStmt(cs->getCapturedStmt(), - /*useCurrentScope=*/true); - }); + return emitTarget( + *this, s, queue, item, begin, end, [&]() -> mlir::LogicalResult { + const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target); + return emitStmt(cs->getCapturedStmt(), + /*useCurrentScope=*/true); + }); } mlir::LogicalResult CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) { @@ -456,49 +697,55 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective( 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(); + "expected 'target' to be the outermost directive"); - return emitTargetOp( - *this, s, queue, targetItem, begin, end, targetOps, mapSyms, - [&]() -> mlir::LogicalResult { + return emitTarget( + *this, s, queue, targetItem, begin, end, [&]() -> 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 { - const CapturedStmt *cs = - s.getCapturedStmt(llvm::omp::OMPD_parallel); - return emitStmt(cs->getCapturedStmt(), - /*useCurrentScope=*/true); - }); + "expected a 'parallel' directive nested in 'target'"); + + return emitParallel(*this, s, queue, parallelItem, begin, end, + [&]() -> mlir::LogicalResult { + const CapturedStmt *cs = + s.getCapturedStmt(llvm::omp::OMPD_parallel); + return emitStmt(cs->getCapturedStmt(), + /*useCurrentScope=*/true); + }); }); } 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 directive"); + + return emitTarget( + *this, s, queue, targetItem, begin, end, [&]() -> mlir::LogicalResult { + omp::ConstructQueue::const_iterator parallelItem = + std::next(targetItem); + assert(parallelItem != queue.end() && + parallelItem->id == llvm::omp::OMPD_parallel && + "expected a 'parallel' directive nested in 'target'"); + + return emitParallel( + *this, s, queue, parallelItem, begin, end, + [&]() -> mlir::LogicalResult { + omp::ConstructQueue::const_iterator forItem = + std::next(parallelItem); + assert(forItem != queue.end() && + forItem->id == llvm::omp::OMPD_for && + "expected a 'for' directive 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..9875d18cf198f6 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/for-loop-forms.c @@ -0,0 +1,228 @@ +// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s + +void during(int); +int getLB(void); + +// 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. + +// Decreasing loop: `i > 0; i--` (induction variable on the left, plain +// unary 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 compound assignment: `i -= 2`. +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 `i = i - 2` canonical assignment form. +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 `i = 2 + i` 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 on the left of the comparison +// (`20 > i`). +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 +} + +// Loop with a side-effecting init expression (a function call, `getLB()`). +void side_effecting_init() { + // CHECK: 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]] + // 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 +} + +// Loop with a pointer induction variable (`int *p`). +void pointer_induction_var() { + // CHECK: cir.func{{.*}}@{{.*}}pointer_induction_var + int a[10]; +#pragma omp for + for (int *p = a; p < a + 10; ++p) { + during(*p); + } + // 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 +} + +// Pre-declared induction variable (plain assignment, no DeclStmt) with a +// `!=` condition. +void predeclared_var_and_not_equal_cond() { + // CHECK: cir.func{{.*}}@{{.*}}predeclared_var_and_not_equal_cond + int i; +#pragma omp for + for (i = 0; i != 10; i++) { + during(i); + } + // 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 +} + +// Same pre-declared counter and `!=` condition, inside a combined +// `parallel for`. +void predeclared_var_inside_parallel_for() { + // CHECK: 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..76f31ac0786996 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/parallel-for.c @@ -0,0 +1,54 @@ +// 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 { + + // 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. + // 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 + + // 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 the 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-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c new file mode 100644 index 00000000000000..428174a1d3bf2c --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-device.c @@ -0,0 +1,84 @@ +// Two-step host-BC -> device pipeline that mirrors the offloading driver. +// +// Step 1: Host compilation to bitcode (provides offload entry info to device pass). +// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc +// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc +// +// Step 2: Device compilation using host BC. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \ +// RUN: -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \ +// RUN: -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void during(int); + +void target_parallel_for() { +#pragma omp target +#pragma omp parallel +#pragma omp for + for (int i = 0; i < 10; i++) { + during(i); + } +} + +void combined_target_parallel_for() { +#pragma omp target parallel for + for (int i = 0; i < 10; i++) { + during(i); + } +} + +// 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_target_deinit() + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par +// LLVM: call void @__kmpc_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par, +// LLVM: call void @__kmpc_barrier( + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par(i32 +// LLVM: call void @during(i32 + +// 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_target_deinit() + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par +// LLVM: call void @__kmpc_for_static_loop_4u(ptr {{.*}}, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par, +// LLVM: call void @__kmpc_barrier( + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par(i32 +// LLVM: call void @during(i32 + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_for_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 0, i32 0) +// OGCG: call void @__kmpc_target_deinit() + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}_omp_outlined( +// OGCG: call void @__kmpc_for_static_init_4( +// OGCG: call void @during(i32 +// OGCG: call void @__kmpc_for_static_fini( +// OGCG: call void @__kmpc_barrier( + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_combined_target_parallel_for_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 0, i32 0) +// OGCG: call void @__kmpc_target_deinit() + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}_omp_outlined( +// OGCG: call void @__kmpc_for_static_init_4( +// OGCG: call void @during(i32 +// OGCG: call void @__kmpc_for_static_fini( diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c new file mode 100644 index 00000000000000..41428595851922 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for-llvm-host.c @@ -0,0 +1,83 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void during(int); + +void target_parallel_for() { +#pragma omp target +#pragma omp parallel +#pragma omp for + for (int i = 0; i < 10; i++) { + during(i); + } +} + +void combined_target_parallel_for() { +#pragma omp target parallel for + for (int i = 0; i < 10; i++) { + during(i); + } +} + +// The explicit target/parallel/for nesting and the combined `target parallel +// for` spelling lower identically: an omp.target forks an omp.parallel via +// __kmpc_fork_call, whose body statically schedules the omp.wsloop via +// __kmpc_for_static_init_4/_4u. +// +// LLVM-LABEL: define {{.*}} void @target_parallel_for( +// 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: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}( + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par) + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}..omp_par( +// LLVM: call void @__kmpc_for_static_init_4u( +// LLVM: call void @__kmpc_for_static_fini( +// LLVM: call void @__kmpc_barrier( +// LLVM: call void @during(i32 + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par) + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}..omp_par( +// LLVM: call void @__kmpc_for_static_init_4u( +// LLVM: call void @__kmpc_for_static_fini( +// LLVM: call void @__kmpc_barrier( +// LLVM: call void @during(i32 + +// OGCG-LABEL: define {{.*}} void @target_parallel_for( +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}( + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}( +// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}.omp_outlined) + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_for_l{{.*}}.omp_outlined( +// OGCG: call void @__kmpc_for_static_init_4( +// OGCG: call void @during(i32 +// OGCG: call void @__kmpc_for_static_fini( + +// OGCG-LABEL: define {{.*}} void @combined_target_parallel_for( +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}( + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}( +// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 0, ptr @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}.omp_outlined) + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_combined_target_parallel_for_l{{.*}}.omp_outlined( +// OGCG: call void @__kmpc_for_static_init_4( +// OGCG: call void @during(i32 +// OGCG: call void @__kmpc_for_static_fini( 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..d9ba201a312b57 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-for.c @@ -0,0 +1,127 @@ +// 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); + +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: %[[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 +// 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 + +// 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 + +// 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); + } +} + +// 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} diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c new file mode 100644 index 00000000000000..4ce2badf7b182a --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-device.c @@ -0,0 +1,85 @@ +// Two-step host-BC -> device pipeline that mirrors the offloading driver. +// +// Step 1: Host compilation to bitcode (provides offload entry info to device pass). +// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc +// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc +// +// Step 2: Device compilation using host BC. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \ +// RUN: -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \ +// RUN: -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void use(int); + +void target_parallel(int x) { +#pragma omp target parallel map(tofrom : x) + { + use(x); + } +} + +void target_parallel_proc_bind(int x) { +#pragma omp target parallel proc_bind(spread) map(tofrom : x) + { + use(x); + } +} + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// 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_l{{.*}}..omp_par, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par.wrapper, ptr {{.*}}, i64 1, i32 0) +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par( +// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8 +// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: ret void + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// 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_proc_bind_l{{.*}}..omp_par, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par.wrapper, ptr {{.*}}, i64 1, i32 0) +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par( +// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8 +// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 1, i32 0) +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}_omp_outlined( +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: call void @__kmpc_parallel_60(ptr {{.*}}, i32 {{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}_omp_outlined, ptr null, ptr {{.*}}, i64 1, i32 0) +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}_omp_outlined( +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: ret void diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c new file mode 100644 index 00000000000000..fee411bc509e9f --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-parallel-llvm-host.c @@ -0,0 +1,94 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void use(int); + +void target_parallel(int x) { +#pragma omp target parallel map(tofrom : x) + { + use(x); + } +} + +void target_parallel_proc_bind(int x) { +#pragma omp target parallel proc_bind(spread) map(tofrom : x) + { + use(x); + } +} + +// Host wrappers +// +// LLVM-LABEL: define {{.*}} void @target_parallel( +// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]] +// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4 +// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}(ptr %[[X_ADDR]], ptr null) + +// LLVM-LABEL: define {{.*}} void @target_parallel_proc_bind( +// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]] +// LLVM: %[[X_ADDR:.*]] = alloca i32, align 4 +// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}(ptr %[[X_ADDR]], ptr null) + +// Outlined target functions: `target parallel` forks an omp.parallel region +// via __kmpc_fork_call; the proc_bind clause lowers to a __kmpc_push_proc_bind +// call right before the fork. +// +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par, ptr {{.*}}) + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}..omp_par( +// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8 +// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: call void @__kmpc_push_proc_bind(ptr {{.*}}, i32 {{.*}}, i32 4) +// LLVM: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par, ptr {{.*}}) + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}..omp_par( +// LLVM: %[[P:.*]] = load ptr, ptr %{{.*}}, align 8 +// LLVM: %[[V:.*]] = load i32, ptr %[[P]], align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: ret void + +// OGCG interleaves host wrapper and outlined function per target region. +// +// OGCG-LABEL: define {{.*}} void @target_parallel(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}( +// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_l{{.*}}.omp_outlined, ptr {{.*}}) + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_l{{.*}}.omp_outlined( +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: ret void + +// OGCG-LABEL: define {{.*}} void @target_parallel_proc_bind(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}( +// OGCG: call void @__kmpc_push_proc_bind(ptr {{.*}}, i32 {{.*}}, i32 4) +// OGCG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr {{.*}}, i32 1, ptr @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}.omp_outlined, ptr {{.*}}) + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_parallel_proc_bind_l{{.*}}.omp_outlined( +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: ret void _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
