https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224041
>From 34338f853176b085795e59098a7300e1d5aea1cc Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Tue, 15 Sep 2026 10:34:44 -0500 Subject: [PATCH] [clang][OpenMP] Add no-loop SPMD kernel promotion A target teams distribute parallel for that is guaranteed a thread for every iteration does not need the loop around its body. Flang already drops it and runs the region as a no-loop kernel. Enable the same optimization for Clang through mirroring Flang's MLIR promotion using OpenMPIRBuilder. The kernel is tagged SPMD_NO_LOOP, so the runtime sizes the grid to the iteration space, and the body is emitted without a loop around it. The canonical loop it consumes is reconstructed in the no-loop branch rather than taken from an OMPCanonicalLoop node, so the promotion does not require -fopenmp-enable-irbuilder. Restrict offload entry creation to module level finalize, preventing asserts on missing offload entries from nested CodeGenFunction finalizing before module completion. --- clang/lib/CodeGen/CGOpenMPRuntime.h | 6 + clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 24 ++- clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 4 + clang/lib/CodeGen/CGStmtOpenMP.cpp | 172 +++++++++++++----- clang/lib/CodeGen/CodeGenFunction.cpp | 18 +- clang/lib/CodeGen/CodeGenFunction.h | 3 +- clang/test/OpenMP/target_no_loop.c | 81 +++++++++ .../llvm/Frontend/OpenMP/OMPIRBuilder.h | 8 + llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +- offload/test/offloading/target-no-loop.c | 91 +++++++++ 10 files changed, 352 insertions(+), 57 deletions(-) create mode 100644 clang/test/OpenMP/target_no_loop.c create mode 100644 offload/test/offloading/target-no-loop.c diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h index 73bbcc05f515a..fb78f90e33878 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/clang/lib/CodeGen/CGOpenMPRuntime.h @@ -669,6 +669,12 @@ class CGOpenMPRuntime { return false; }; + /// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel, + /// mirroring Flang's MLIR promotion path. + virtual bool canPromoteToNoLoop(const OMPExecutableDirective &D) const { + return false; + } + /// Get call to __kmpc_alloc_shared virtual std::pair<llvm::Value *, llvm::Value *> getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) { diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index e785ee635af70..264d78b0f5245 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -751,9 +751,14 @@ void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D, CodeGenFunction &CGF, EntryFunctionState &EST, bool IsSPMD) { llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs; - Attrs.ExecFlags = - IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD - : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC; + if (IsSPMD && canPromoteToNoLoop(D)) + Attrs.ExecFlags = + llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP; + else + Attrs.ExecFlags = + IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD + : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC; + computeMinAndMaxThreadsAndTeams(D, CGF, Attrs); CGBuilderTy &Bld = CGF.Builder; @@ -1145,6 +1150,19 @@ bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF, return llvm::is_contained(I->getSecond().DelayedVariableLengthDecls, VD); } +bool CGOpenMPRuntimeGPU::canPromoteToNoLoop( + const OMPExecutableDirective &D) const { + OpenMPDirectiveKind DKind = D.getDirectiveKind(); + const LangOptions &LangOpts = CGM.getLangOpts(); + return (DKind == OMPD_target_teams_distribute_parallel_for || + DKind == OMPD_target_teams_distribute_parallel_for_simd) && + LangOpts.OpenMPTeamSubscription && LangOpts.OpenMPThreadSubscription && + !D.hasClausesOfKind<OMPNumTeamsClause>() && + !D.hasClausesOfKind<OMPReductionClause>() && + !D.hasClausesOfKind<OMPLastprivateClause>() && + !D.hasClausesOfKind<OMPLinearClause>(); +} + std::pair<llvm::Value *, llvm::Value *> CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) { diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h index 3a7ee5456a9d2..d1f682780b9bd 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h @@ -141,6 +141,10 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime { bool isDelayedVariableLengthDecl(CodeGenFunction &CGF, const VarDecl *VD) const override; + /// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel, + /// mirroring Flang's MLIR promotion path. + bool canPromoteToNoLoop(const OMPExecutableDirective &D) const override; + /// Get call to __kmpc_alloc_shared std::pair<llvm::Value *, llvm::Value *> getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) override; diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 7c3b30c6cedc0..2836911518398 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -66,6 +66,15 @@ static bool canEmitGPUFusedDistSchedule(const CodeGenModule &CGM, !S.getSingleClause<OMPOrderedClause>(); } +static bool canEmitGPUNoLoopKernel(CodeGenModule &CGM, + const OMPLoopDirective &S) { + const auto *D = dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&S); + return S.getLoopsNumber() == 1 && + CGM.getOpenMPRuntime().canPromoteToNoLoop(S) && + !S.getSingleClause<OMPScheduleClause>() && + !S.getSingleClause<OMPDistScheduleClause>() && !(D && D->hasCancel()); +} + namespace { /// Lexical scope for OpenMP executable constructs, that handles correct codegen /// for captured expressions. @@ -2371,6 +2380,24 @@ emitCapturedStmtCall(CodeGenFunction &ParentCGF, EmittedClosureTy Cap, return ParentCGF.Builder.CreateCall(Cap.first, EffectiveArgs); } +static llvm::CanonicalLoopInfo * +createCanonicalLoop(CodeGenFunction &CGF, llvm::Value *TripCount, + llvm::function_ref<void(llvm::Value *)> BodyGen) { + auto BodyGenCB = [&](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP, + llvm::Value *IndVar) { + CGF.Builder.restoreIP(CodeGenIP); + BodyGen(IndVar); + return llvm::Error::success(); + }; + + llvm::OpenMPIRBuilder &OMPBuilder = + CGF.CGM.getOpenMPRuntime().getOMPBuilder(); + llvm::CanonicalLoopInfo *CL = cantFail( + OMPBuilder.createCanonicalLoop(CGF.Builder, BodyGenCB, TripCount)); + CGF.Builder.restoreIP(CL->getAfterIP()); + return CL; +} + llvm::CanonicalLoopInfo * CodeGenFunction::EmitOMPCollapsedCanonicalLoopNest(const Stmt *S, int Depth) { assert(Depth == 1 && "Nested loops with OpenMPIRBuilder not yet implemented"); @@ -2446,11 +2473,7 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) { llvm::Value *DistVal = Builder.CreateLoad(CountAddr, ".count"); // Emit the loop structure. - llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder(); - auto BodyGen = [&, this](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP, - llvm::Value *IndVar) { - Builder.restoreIP(CodeGenIP); - + auto BodyGen = [&, this](llvm::Value *IndVar) { // Emit the loop body: Convert the logical iteration number to the loop // variable and emit the body. const DeclRefExpr *LoopVarRef = S->getLoopVarRef(); @@ -2461,14 +2484,11 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) { RunCleanupsScope BodyScope(*this); EmitStmt(BodyStmt); - return llvm::Error::success(); }; - llvm::CanonicalLoopInfo *CL = - cantFail(OMPBuilder.createCanonicalLoop(Builder, BodyGen, DistVal)); + llvm::CanonicalLoopInfo *CL = createCanonicalLoop(*this, DistVal, BodyGen); // Finish up the loop. - Builder.restoreIP(CL->getAfterIP()); ForScope.ForceCleanup(); // Remember the CanonicalLoopInfo for parent AST nodes consuming it. @@ -2648,12 +2668,26 @@ static void emitAlignedClause(CodeGenFunction &CGF, } void CodeGenFunction::EmitOMPPrivateLoopCounters( - const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope) { + const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope, + bool OnlyUnresolved) { if (!HaveInsertPoint()) return; auto I = S.private_counters().begin(); for (const Expr *E : S.counters()) { - const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl()); + const auto *DRE = cast<DeclRefExpr>(E); + const auto *VD = cast<VarDecl>(DRE->getDecl()); + // Skip counters that already resolve, mirroring EmitDeclRefLValue's + // handling for these cases. + if (OnlyUnresolved) { + const VarDecl *Canonical = VD->getCanonicalDecl(); + if (!DRE->refersToEnclosingVariableOrCapture() || !CapturedStmtInfo || + LocalDeclMap.count(Canonical) || + CapturedStmtInfo->lookup(Canonical)) { + ++I; + continue; + } + } + const auto *PrivateVD = cast<VarDecl>(cast<DeclRefExpr>(*I)->getDecl()); // Emit var without initialization. AutoVarEmission VarEmission = EmitAutoVarAlloca(*PrivateVD); @@ -3691,6 +3725,22 @@ static void emitDistributeParallelForDistributeInnerBoundParams( CapturedVars.push_back(UBCast); } +static void emitLoopIterationspaceVars(CodeGenFunction &CGF, + const OMPLoopDirective &S) { + // Emit the loop iteration variable. + const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable()); + CGF.EmitVarDecl(*cast<VarDecl>(IVExpr->getDecl())); + + // Emit the iterations count variable. + // If it is not a variable, Sema decided to calculate iterations count on each + // iteration (e.g., it is foldable into a constant). + if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) { + CGF.EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl())); + // Emit calculation of the iterations count. + CGF.EmitIgnoredExpr(S.getCalcLastIteration()); + } +} + static void emitInnerParallelForWhenCombined(CodeGenFunction &CGF, const OMPLoopDirective &S, @@ -3710,6 +3760,55 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF, HasCancel = D->hasCancel(); } CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel); + + CodeGenModule &CGM = CGF.CGM; + if (canEmitGPUNoLoopKernel(CGM, S)) { + // Prepare the loop variables and their privatization. + emitLoopIterationspaceVars(CGF, S); + OMPLoopScope PreInitScope(CGF, S); + + CodeGenFunction::OMPPrivateScope PrivateScope(CGF); + CGF.EmitOMPPrivateClause(S, PrivateScope); + CGF.EmitOMPPrivateLoopCounters(S, PrivateScope); + (void)PrivateScope.Privatize(); + + if (isOpenMPTargetExecutionDirective(EKind)) + CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S); + + // Rebuild what the OMPCanonicalLoop node supplies under the IRBuilder + // flag: iteration count, loop, index-to-variable mapping, body. + const Expr *NumIterations = S.getNumIterations(); + llvm::Value *TripCount = CGF.EmitScalarConversion( + CGF.EmitScalarExpr(NumIterations), NumIterations->getType(), + S.getIterationVariable()->getType(), S.getBeginLoc()); + llvm::CanonicalLoopInfo *CLI = + createCanonicalLoop(CGF, TripCount, [&CGF, &S](llvm::Value *IndVar) { + llvm::BasicBlock *BodyExit = llvm::splitBBWithSuffix( + CGF.Builder, /*CreateBranch=*/false, ".cont"); + CGF.EmitStoreOfScalar(IndVar, + CGF.EmitLValue(S.getIterationVariable())); + emitOMPLoopBodyWithStopPoint(CGF, S, CodeGenFunction::JumpDest()); + CGF.Builder.CreateBr(BodyExit); + }); + + llvm::OpenMPIRBuilder::InsertPointTy AllocaIP( + CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator()); + llvm::OpenMPIRBuilder &OMPBuilder = + CGM.getOpenMPRuntime().getOMPBuilder(); + + cantFail(OMPBuilder.applyWorkshareLoop( + CGF.Builder.getCurrentDebugLocation(), CLI, AllocaIP, + /*NeedsBarrier=*/!S.getSingleClause<OMPNowaitClause>(), + llvm::omp::OMP_SCHEDULE_Default, + /*ChunkSize=*/nullptr, /*HasSimdModifier=*/false, + /*HasMonotonicModifier=*/false, /*HasNonmonotonicModifier=*/false, + /*HasOrderedClause=*/false, + llvm::omp::WorksharingLoopType::DistributeForStaticLoop, + /*NoLoop=*/true, /*HasDistSchedule=*/false, + /*DistScheduleChunkSize=*/nullptr)); + return; + } + CGF.EmitOMPWorksharingLoop(S, S.getPrevEnsureUpperBound(), emitDistributeParallelForInnerBounds, emitDistributeParallelForDispatchBounds); @@ -3788,19 +3887,7 @@ bool CodeGenFunction::EmitOMPWorksharingLoop( const OMPLoopDirective &S, Expr *EUB, const CodeGenLoopBoundsTy &CodeGenLoopBounds, const CodeGenDispatchBoundsTy &CGDispatchBounds) { - // Emit the loop iteration variable. - const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable()); - const auto *IVDecl = cast<VarDecl>(IVExpr->getDecl()); - EmitVarDecl(*IVDecl); - - // Emit the iterations count variable. - // If it is not a variable, Sema decided to calculate iterations count on each - // iteration (e.g., it is foldable into a constant). - if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) { - EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl())); - // Emit calculation of the iterations count. - EmitIgnoredExpr(S.getCalcLastIteration()); - } + emitLoopIterationspaceVars(*this, S); CGOpenMPRuntime &RT = CGM.getOpenMPRuntime(); @@ -3895,6 +3982,7 @@ bool CodeGenFunction::EmitOMPWorksharingLoop( HasChunkSizeOne = (EvaluatedChunk.getLimitedValue() == 1); } } + const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable()); const unsigned IVSize = getContext().getTypeSize(IVExpr->getType()); const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation(); // OpenMP 4.5, 2.7.1 Loop Construct, Description. @@ -6216,19 +6304,7 @@ void CodeGenFunction::EmitOMPScanDirective(const OMPScanDirective &S) { void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, const CodeGenLoopTy &CodeGenLoop, Expr *IncExpr) { - // Emit the loop iteration variable. - const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable()); - const auto *IVDecl = cast<VarDecl>(IVExpr->getDecl()); - EmitVarDecl(*IVDecl); - - // Emit the iterations count variable. - // If it is not a variable, Sema decided to calculate iterations count on each - // iteration (e.g., it is foldable into a constant). - if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) { - EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl())); - // Emit calculation of the iterations count. - EmitIgnoredExpr(S.getCalcLastIteration()); - } + emitLoopIterationspaceVars(*this, S); CGOpenMPRuntime &RT = CGM.getOpenMPRuntime(); @@ -6287,6 +6363,8 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, !isOpenMPParallelDirective(S.getDirectiveKind()) && !isOpenMPTeamsDirective(S.getDirectiveKind())) EmitOMPReductionClauseInit(S, LoopScope); + + const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S); HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope); EmitOMPPrivateLoopCounters(S, LoopScope); (void)LoopScope.Privatize(); @@ -6309,12 +6387,15 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, CGM.getOpenMPRuntime().getDefaultDistScheduleAndChunk( *this, S, ScheduleKind, Chunk); } + const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable()); const unsigned IVSize = getContext().getTypeSize(IVExpr->getType()); const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation(); - // GPU fused schedule: omit the outer distribute loop and let the inner - // worksharing loop schedule the flattened team/thread iteration space. - if (canEmitGPUFusedDistSchedule(CGM, S, S.getDirectiveKind())) { + // omit the outer distribute loop and let the inner worksharing loop + // schedule the flattened team/thread iteration space, necessary for + // GPU fused schedule and no-loop optimization + if (canEmitGPUFusedDistSchedule(CGM, S, S.getDirectiveKind()) || + NoLoopKernel) { JumpDest LoopExit = getJumpDestInCurrentScope(createBasicBlock("omp.loop.exit")); CodeGenLoop(*this, S, LoopExit); @@ -6421,10 +6502,13 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, } } if (isOpenMPSimdDirective(S.getDirectiveKind())) { - EmitOMPSimdFinal(S, [IL, &S](CodeGenFunction &CGF) { - return CGF.Builder.CreateIsNotNull( - CGF.EmitLoadOfScalar(IL, S.getBeginLoc())); - }); + EmitOMPSimdFinal( + S, [IL, &S, NoLoopKernel](CodeGenFunction &CGF) -> llvm::Value * { + if (NoLoopKernel) + return nullptr; + return CGF.Builder.CreateIsNotNull( + CGF.EmitLoadOfScalar(IL, S.getBeginLoc())); + }); } if (isOpenMPSimdDirective(S.getDirectiveKind()) && !isOpenMPParallelDirective(S.getDirectiveKind()) && diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index fe39235fcd4f6..fa4404c66ca70 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -104,16 +104,18 @@ CodeGenFunction::~CodeGenFunction() { assert(DeferredDeactivationCleanupStack.empty() && "missed to deactivate a cleanup"); - if (getLangOpts().OpenMP && CurFn) + if (getLangOpts().OpenMP && CurFn) { CGM.getOpenMPRuntime().functionFinished(*this); - // If we have an OpenMPIRBuilder we want to finalize functions (incl. - // outlining etc) at some point. Doing it once the function codegen is done - // seems to be a reasonable spot. We do it here, as opposed to the deletion - // time of the CodeGenModule, because we have to ensure the IR has not yet - // been "emitted" to the outside, thus, modifications are still sensible. - if (CGM.getLangOpts().OpenMPIRBuilder && CurFn) - CGM.getOpenMPRuntime().getOMPBuilder().finalize(CurFn); + // Finalizing (incl. outlining etc) once the function codegen is done, as + // opposed to the deletion time of the CodeGenModule, ensures the IR has + // not yet been "emitted" to the outside, thus, modifications are still + // sensible. + llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder(); + if (CGM.getLangOpts().OpenMPIRBuilder || + OMPBuilder.hasPendingOutlines(CurFn)) + OMPBuilder.finalize(CurFn); + } } // Map the LangOption for exception behavior into diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index 3c8188c4cefdd..26917a6dc0a8c 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -4131,7 +4131,8 @@ class CodeGenFunction : public CodeGenTypeCache { JumpDest getOMPCancelDestination(OpenMPDirectiveKind Kind); /// Emit initial code for loop counters of loop-based directives. void EmitOMPPrivateLoopCounters(const OMPLoopDirective &S, - OMPPrivateScope &LoopScope); + OMPPrivateScope &LoopScope, + bool OnlyUnresolved = false); /// Helper for the OpenMP loop directives. void EmitOMPLoopBody(const OMPLoopDirective &D, JumpDest LoopExit); diff --git a/clang/test/OpenMP/target_no_loop.c b/clang/test/OpenMP/target_no_loop.c new file mode 100644 index 0000000000000..1caffb04a8c76 --- /dev/null +++ b/clang/test/OpenMP/target_no_loop.c @@ -0,0 +1,81 @@ +// REQUIRES: amdgpu-registered-target + +// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc + +// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-host.bc \ +// RUN: -fopenmp-assume-teams-oversubscription \ +// RUN: -fopenmp-assume-threads-oversubscription \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ +// RUN: --check-prefixes=NOLOOP + +// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-host.bc \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ +// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u + +// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-host.bc \ +// RUN: -fopenmp-assume-teams-oversubscription \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ +// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u + +// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-host.bc \ +// RUN: -fopenmp-assume-threads-oversubscription \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ +// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u + +// expected-no-diagnostics + +void no_loop(int *array) { +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void no_loop_simd(int *array) { +#pragma omp target teams distribute parallel for simd + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void no_loop_nowait(int *array) { +#pragma omp target teams distribute parallel for nowait + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +// NOLOOP: no_loop_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 +// NOLOOP: no_loop_simd_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_l{{[0-9]+}}{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_l{{[0-9]+}}{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_simd{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: store i32 1024, ptr %i +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_simd{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_nowait{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_nowait{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.exit: +// NOLOOP-NEXT: br label %omp_loop.after +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// SPMD-COUNT-3: _kernel_environment {{.*}} i8 0, i8 1, i8 2 diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index a13832d69b6a0..ccf01e504269f 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -2731,6 +2731,14 @@ class OpenMPIRBuilder { OutlineInfos.emplace_back(std::move(OI)); } + /// Return true if \p Fn has a region registered for outlining that has not + /// been processed yet. + bool hasPendingOutlines(const Function *Fn) const { + return any_of(OutlineInfos, [Fn](const std::unique_ptr<OutlineInfo> &OI) { + return OI->getFunction() == Fn; + }); + } + /// An ordered map of auto-generated variables to their unique names. /// It stores variables with the following names: 1) ".gomp_critical_user_" + /// <critical_section_name> + ".var" for "omp critical" directives; 2) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8b196c461f195..4fee93b9d90a4 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -1063,7 +1063,7 @@ void OpenMPIRBuilder::finalize(Function *Fn) { "OMPIRBuilder finalization \n"; }; - if (!OffloadInfoManager.empty()) + if (!Fn && !OffloadInfoManager.empty()) createOffloadEntriesAndInfoMetadata(ErrorReportFn); // Rewrite uses of globals to their replacement declare target globals if diff --git a/offload/test/offloading/target-no-loop.c b/offload/test/offloading/target-no-loop.c new file mode 100644 index 0000000000000..59558c67152e2 --- /dev/null +++ b/offload/test/offloading/target-no-loop.c @@ -0,0 +1,91 @@ +// clang-format off +// C counterpart of fortran/target-no-loop.f90. + +// RUN: %libomptarget-compile-generic -O3 -fopenmp-assume-threads-oversubscription -fopenmp-assume-teams-oversubscription +// RUN: env LIBOMPTARGET_INFO=16 OMP_NUM_TEAMS=16 OMP_TEAMS_THREAD_LIMIT=16 %libomptarget-run-generic 2>&1 | %fcheck-generic +// REQUIRES: gpu +// XFAIL: intelgpu + +#include <stdio.h> + +static int check_errors(int *array) { + int errors = 0; + for (int i = 0; i < 1024; ++i) + if (array[i] != i + 1) + ++errors; + return errors; +} + +int main(void) { + int array[1024]; + int errors = 0; + int red; + + for (int i = 0; i < 1024; ++i) + array[i] = 1; + + // No-loop kernel +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + errors += check_errors(array); + + // SPMD kernel (num_teams clause blocks promotion to no-loop) + for (int i = 0; i < 1024; ++i) + array[i] = 1; +#pragma omp target teams distribute parallel for num_teams(3) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + errors += check_errors(array); + + // No-loop kernel + for (int i = 0; i < 1024; ++i) + array[i] = 1; +#pragma omp target teams distribute parallel for num_threads(64) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + errors += check_errors(array); + + // SPMD kernel + for (int i = 0; i < 1024; ++i) + array[i] = 1; +#pragma omp target parallel for + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + errors += check_errors(array); + + // Generic kernel + for (int i = 0; i < 1024; ++i) + array[i] = 1; +#pragma omp target teams distribute + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + errors += check_errors(array); + + // SPMD kernel (reduction clause blocks promotion to no-loop) + for (int i = 0; i < 1024; ++i) + array[i] = 1; + red = 0; +#pragma omp target teams distribute parallel for reduction(+ : red) + for (int i = 0; i < 1024; ++i) + red += array[i]; + if (red != 1024) + ++errors; + + printf("number of errors: %d\n", errors); + return 0; +} + +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 3x 16 {{.*}} +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 {{.*}} +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 1x 16 +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} Generic-SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 16x 16 {{.*}} +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}} +// CHECK: number of errors: 0 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
