https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224041
>From c7db8c2a51903fe86ece3a0d2b310d2e1c046a86 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Tue, 15 Sep 2026 10:34:44 -0500 Subject: [PATCH 1/4] [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 ae55295cadc14..9a93f5c5c2e82 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 e7322bdd0deca..12e4f718d64ed 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -763,9 +763,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; @@ -1158,6 +1163,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 50efbfd1577a5..d40d15767552e 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h @@ -146,6 +146,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 0913f76018a58..1f57c8292e5b2 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -75,6 +75,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. @@ -2574,6 +2583,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"); @@ -2649,11 +2676,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(); @@ -2664,14 +2687,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. @@ -2861,12 +2881,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); @@ -3915,6 +3949,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, @@ -3934,6 +3984,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->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); @@ -4012,19 +4111,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(); @@ -4119,6 +4206,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. @@ -6469,19 +6557,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(); @@ -6540,6 +6616,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(); @@ -6562,12 +6640,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); @@ -6674,10 +6755,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 62dade6cf6f49..78e37b7b2b853 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -88,16 +88,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 490267aaefd86..caefc275771ff 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -4185,7 +4185,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 5acbb69248309..cfc84226b34a7 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -2734,6 +2734,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 b6a11bc3c805e..4f98f59d71e31 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..8407cffc13020 --- /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: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 +// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 3x 16 {{.*}} +// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 {{.*}} +// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 1x 16 +// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} Generic-SPMD mode +// CHECK: info: #Args: 2 Teams x Thrds: 16x 16 {{.*}} +// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}} +// CHECK: number of errors: 0 >From b404e8490fcc7069b98d53382ad7529f7e0657e7 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Fri, 25 Sep 2026 04:27:56 -0500 Subject: [PATCH 2/4] Remove unused OnlyUnresolved parameter --- clang/lib/CodeGen/CGStmtOpenMP.cpp | 18 ++---------------- clang/lib/CodeGen/CodeGenFunction.h | 3 +-- 2 files changed, 3 insertions(+), 18 deletions(-) diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 1f57c8292e5b2..251241b38f918 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -2881,26 +2881,12 @@ static void emitAlignedClause(CodeGenFunction &CGF, } void CodeGenFunction::EmitOMPPrivateLoopCounters( - const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope, - bool OnlyUnresolved) { + const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope) { if (!HaveInsertPoint()) return; auto I = S.private_counters().begin(); for (const Expr *E : S.counters()) { - 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 *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl()); const auto *PrivateVD = cast<VarDecl>(cast<DeclRefExpr>(*I)->getDecl()); // Emit var without initialization. AutoVarEmission VarEmission = EmitAutoVarAlloca(*PrivateVD); diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index caefc275771ff..490267aaefd86 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -4185,8 +4185,7 @@ 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, - bool OnlyUnresolved = false); + OMPPrivateScope &LoopScope); /// Helper for the OpenMP loop directives. void EmitOMPLoopBody(const OMPLoopDirective &D, JumpDest LoopExit); >From 4fa0c8c4ced823567061c4a0c8f899414c5c2c32 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Tue, 6 Oct 2026 10:57:41 -0500 Subject: [PATCH 3/4] Share one no-loop eligibility check between exec mode and body The kernel is tagged SPMD_NO_LOOP by a single check, and the body is emitted loop-free exactly when the kernel carries that tag. Add tests for the clauses that block promotion. --- clang/lib/CodeGen/CGOpenMPRuntime.h | 9 ++-- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 53 +++++++++++++-------- clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 14 ++++-- clang/lib/CodeGen/CGStmtOpenMP.cpp | 13 +----- clang/test/OpenMP/target_no_loop.c | 59 +++++++++++++++++++++++- 5 files changed, 109 insertions(+), 39 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h index 9a93f5c5c2e82..29cc5c0bc2711 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/clang/lib/CodeGen/CGOpenMPRuntime.h @@ -669,11 +669,10 @@ 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; - } + /// Check whether the target kernel being emitted is tagged SPMD_NO_LOOP, to + /// complete the promotion to a "no-loop" SPMD kernel, mirroring Flang's MLIR + /// promotion path. + virtual bool canPromoteToNoLoop() const { return false; } /// Get call to __kmpc_alloc_shared virtual std::pair<llvm::Value *, llvm::Value *> diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index 12e4f718d64ed..af342eaec2da4 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -706,6 +706,32 @@ static bool supportsSPMDExecutionMode(ASTContext &Ctx, "Unknown programming model for OpenMP directive on NVPTX target."); } +static bool isNoLoopEligible(ASTContext &Ctx, const OMPExecutableDirective &D) { + const LangOptions &LangOpts = Ctx.getLangOpts(); + if (!LangOpts.OpenMPTeamSubscription || !LangOpts.OpenMPThreadSubscription) + return false; + + OpenMPDirectiveKind DKind = D.getDirectiveKind(); + if (DKind != OMPD_target_teams_distribute_parallel_for && + DKind != OMPD_target_teams_distribute_parallel_for_simd) + return false; + + const auto &LD = cast<OMPLoopDirective>(D); + // Do not filter out 'ordered' since Sema rejects it on these directives. + if (LD.getLoopsNumber() != 1 || LD.hasClausesOfKind<OMPNumTeamsClause>() || + LD.hasClausesOfKind<OMPReductionClause>() || + LD.hasClausesOfKind<OMPLastprivateClause>() || + LD.hasClausesOfKind<OMPLinearClause>() || + LD.getSingleClause<OMPScheduleClause>() || + LD.getSingleClause<OMPDistScheduleClause>()) + return false; + + if (const auto *TTD = + dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&LD)) + return !TTD->hasCancel(); + return true; +} + void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName, llvm::Function *&OutlinedFn, @@ -745,6 +771,7 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D, emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry, CodeGen); IsInTTDRegion = false; + KernelAttrs = {}; } void CGOpenMPRuntimeGPU::emitBareKernelEnvironment( @@ -762,19 +789,19 @@ void CGOpenMPRuntimeGPU::emitBareKernelEnvironment( void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D, CodeGenFunction &CGF, EntryFunctionState &EST, bool IsSPMD) { - llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs; - if (IsSPMD && canPromoteToNoLoop(D)) - Attrs.ExecFlags = + KernelAttrs = {}; + if (IsSPMD && isNoLoopEligible(CGM.getContext(), D)) + KernelAttrs.ExecFlags = llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP; else - Attrs.ExecFlags = + KernelAttrs.ExecFlags = IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC; - computeMinAndMaxThreadsAndTeams(D, CGF, Attrs); + computeMinAndMaxThreadsAndTeams(D, CGF, KernelAttrs); CGBuilderTy &Bld = CGF.Builder; - Bld.restoreIP(OMPBuilder.createTargetInit(Bld, Attrs)); + Bld.restoreIP(OMPBuilder.createTargetInit(Bld, KernelAttrs)); if (!IsSPMD) emitGenericVarsProlog(CGF, EST.Loc); } @@ -863,6 +890,7 @@ void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D, emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry, CodeGen); IsInTTDRegion = false; + KernelAttrs = {}; } void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction( @@ -1163,19 +1191,6 @@ 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 d40d15767552e..669fc34ceecb1 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h @@ -146,9 +146,13 @@ 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; + /// Check whether the target kernel being emitted is tagged SPMD_NO_LOOP, to + /// complete the promotion to a "no-loop" SPMD kernel, mirroring Flang's MLIR + /// promotion path. + bool canPromoteToNoLoop() const override { + return KernelAttrs.ExecFlags == + llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP; + } /// Get call to __kmpc_alloc_shared std::pair<llvm::Value *, llvm::Value *> @@ -382,6 +386,10 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime { /// - otherwise. bool IsInTTDRegion = false; + /// The default attributes of the target region being emitted, filled in by + /// the kernel prologue and read back while the region's body is emitted. + llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs KernelAttrs; + /// Map between an outlined function and its wrapper. llvm::DenseMap<llvm::Function *, llvm::Function *> WrapperFunctionsMap; diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 251241b38f918..9a5ab88a91f6b 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -75,15 +75,6 @@ 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. @@ -3972,7 +3963,7 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF, CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel); CodeGenModule &CGM = CGF.CGM; - if (canEmitGPUNoLoopKernel(CGM, S)) { + if (CGM.getOpenMPRuntime().canPromoteToNoLoop()) { // Prepare the loop variables and their privatization. emitLoopIterationspaceVars(CGF, S); OMPLoopScope PreInitScope(CGF, S); @@ -6603,7 +6594,7 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, !isOpenMPTeamsDirective(S.getDirectiveKind())) EmitOMPReductionClauseInit(S, LoopScope); - const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S); + const bool NoLoopKernel = CGM.getOpenMPRuntime().canPromoteToNoLoop(); HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope); EmitOMPPrivateLoopCounters(S, LoopScope); (void)LoopScope.Privatize(); diff --git a/clang/test/OpenMP/target_no_loop.c b/clang/test/OpenMP/target_no_loop.c index 1caffb04a8c76..ecaa5e1c02494 100644 --- a/clang/test/OpenMP/target_no_loop.c +++ b/clang/test/OpenMP/target_no_loop.c @@ -51,9 +51,66 @@ void no_loop_nowait(int *array) { array[i] = i + 1; } +void spmd_collapse2(int *array) { +#pragma omp target teams distribute parallel for collapse(2) + for (int i = 0; i < 32; ++i) + for (int j = 0; j < 32; ++j) + array[i * 32 + j] = i + j; +} + +void spmd_schedule_static(int *array) { +#pragma omp target teams distribute parallel for schedule(static) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void spmd_schedule_guided(int *array) { +#pragma omp target teams distribute parallel for schedule(guided) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void spmd_dist_schedule(int *array) { +#pragma omp target teams distribute parallel for dist_schedule(static) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void spmd_cancel(int *array) { +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; +#pragma omp cancel for + } +} + +void spmd_lastprivate_scalar(int *array) { + int last = 0; +#pragma omp target teams distribute parallel for lastprivate(last) + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; + last = i; + } +} + +void spmd_linear(int *array) { + int i; +#pragma omp target teams distribute parallel for simd linear(i) + for (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: spmd_collapse2_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_schedule_static_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_schedule_guided_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_dist_schedule_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_cancel_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_lastprivate_scalar_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// NOLOOP: spmd_linear_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 + // NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_l{{[0-9]+}}{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: ret void @@ -78,4 +135,4 @@ void no_loop_nowait(int *array) { // NOLOOP: omp_loop.after: // NOLOOP-NEXT: ret void -// SPMD-COUNT-3: _kernel_environment {{.*}} i8 0, i8 1, i8 2 +// SPMD-COUNT-10: _kernel_environment {{.*}} i8 0, i8 1, i8 2 >From 47f0274faab008b9b34bddd0e366249eeb51fafb Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Wed, 7 Oct 2026 06:20:58 -0500 Subject: [PATCH 4/4] Restructure no-loop test and cover missing cases Group the kernels by whether they are promoted and add cases to cover gaps pointed out in review. --- clang/test/OpenMP/target_no_loop.c | 138 --------------------- clang/test/OpenMP/target_no_loop.cpp | 173 +++++++++++++++++++++++++++ 2 files changed, 173 insertions(+), 138 deletions(-) delete mode 100644 clang/test/OpenMP/target_no_loop.c create mode 100644 clang/test/OpenMP/target_no_loop.cpp diff --git a/clang/test/OpenMP/target_no_loop.c b/clang/test/OpenMP/target_no_loop.c deleted file mode 100644 index ecaa5e1c02494..0000000000000 --- a/clang/test/OpenMP/target_no_loop.c +++ /dev/null @@ -1,138 +0,0 @@ -// 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; -} - -void spmd_collapse2(int *array) { -#pragma omp target teams distribute parallel for collapse(2) - for (int i = 0; i < 32; ++i) - for (int j = 0; j < 32; ++j) - array[i * 32 + j] = i + j; -} - -void spmd_schedule_static(int *array) { -#pragma omp target teams distribute parallel for schedule(static) - for (int i = 0; i < 1024; ++i) - array[i] = i + 1; -} - -void spmd_schedule_guided(int *array) { -#pragma omp target teams distribute parallel for schedule(guided) - for (int i = 0; i < 1024; ++i) - array[i] = i + 1; -} - -void spmd_dist_schedule(int *array) { -#pragma omp target teams distribute parallel for dist_schedule(static) - for (int i = 0; i < 1024; ++i) - array[i] = i + 1; -} - -void spmd_cancel(int *array) { -#pragma omp target teams distribute parallel for - for (int i = 0; i < 1024; ++i) { - array[i] = i + 1; -#pragma omp cancel for - } -} - -void spmd_lastprivate_scalar(int *array) { - int last = 0; -#pragma omp target teams distribute parallel for lastprivate(last) - for (int i = 0; i < 1024; ++i) { - array[i] = i + 1; - last = i; - } -} - -void spmd_linear(int *array) { - int i; -#pragma omp target teams distribute parallel for simd linear(i) - for (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: spmd_collapse2_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_schedule_static_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_schedule_guided_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_dist_schedule_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_cancel_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_lastprivate_scalar_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// NOLOOP: spmd_linear_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 - -// 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-10: _kernel_environment {{.*}} i8 0, i8 1, i8 2 diff --git a/clang/test/OpenMP/target_no_loop.cpp b/clang/test/OpenMP/target_no_loop.cpp new file mode 100644 index 0000000000000..1859b26ae3fb8 --- /dev/null +++ b/clang/test/OpenMP/target_no_loop.cpp @@ -0,0 +1,173 @@ +// 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 promotable_no_loop(int *array, int n) { +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for simd + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for nowait + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for + for (int i = 0; i < n; ++i) + array[i] = i + 1; + + { + int tmp; +#pragma omp target teams distribute parallel for private(tmp) + for (int i = 0; i < 1024; ++i) { + tmp = i + 1; + array[i] = tmp; + } + } + + { + auto set = [&array](int i) { array[i] = i + 1; }; +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) + set(i); + } +} + +void non_promotable(int *array) { +#pragma omp target teams distribute parallel for collapse(2) + for (int i = 0; i < 32; ++i) + for (int j = 0; j < 32; ++j) + array[i * 32 + j] = i + j; + +#pragma omp target teams distribute parallel for schedule(static) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for schedule(guided) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for dist_schedule(static) + for (int i = 0; i < 1024; ++i) + array[i] = i + 1; + +#pragma omp target teams distribute parallel for + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; +#pragma omp cancel for + } + + { + int last = 0; +#pragma omp target teams distribute parallel for lastprivate(last) + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; + last = i; + } + } + + { + int i; +#pragma omp target teams distribute parallel for simd linear(i) + for (i = 0; i < 1024; ++i) + array[i] = i + 1; + } +} + +// NOLOOP-COUNT-6: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 +// NOLOOP-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 + +// SPMD-COUNT-6: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// SPMD-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 + +// no clause +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l37_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l37_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// simd +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l41_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: store i32 1024, ptr %i +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l41_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// nowait +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l45_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l45_{{.*}}, 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 + +// runtime trip count +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l49_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: br label %omp.precond.end +// NOLOOP: [[LAST:%.*]] = load i32, ptr %.capture_expr.1.ascast +// NOLOOP-NEXT: [[TC:%.*]] = add nsw i32 [[LAST]], 1 +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l49_{{.*}}, i32 [[TC]], i32 %{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// private +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l55_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: store ptr %tmp1.ascast, ptr addrspace(5) %gep_tmp1.ascast +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l55_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void +// NOLOOP: store i32 {{.*}}, ptr %loadgep_tmp1.ascast + +// lambda +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l64_{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: ret void +// NOLOOP: [[SET:%.*]] = load ptr, ptr %set.addr.ascast +// NOLOOP-NEXT: [[FIELD:%.*]] = getelementptr inbounds nuw %class.anon, ptr [[SET]], i32 0, i32 0 +// NOLOOP-NEXT: store ptr %array.addr.ascast, ptr [[FIELD]] +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l64_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
