[clang] [llvm] [clang][OpenMP] Add no-loop SPMD kernel promotion (PR #224041)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 16 08:03:27 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
Author: Nicole Aschenbrenner (nicebert)
<details>
<summary>Changes</summary>
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.
Split out of #<!-- -->205325.
---
Patch is 25.58 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/224041.diff
10 Files Affected:
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.h (+6)
- (modified) clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp (+21-3)
- (modified) clang/lib/CodeGen/CGOpenMPRuntimeGPU.h (+4)
- (modified) clang/lib/CodeGen/CGStmtOpenMP.cpp (+121-40)
- (modified) clang/lib/CodeGen/CodeGenFunction.cpp (+10-8)
- (modified) clang/lib/CodeGen/CodeGenFunction.h (+2-1)
- (added) clang/test/OpenMP/target_no_loop.c (+83)
- (modified) llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h (+8)
- (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+1-1)
- (added) offload/test/offloading/target-no-loop.c (+91)
``````````diff
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..4d19fa91172a8 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);
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..0fe3de1780cce
--- /dev/null
+++ b/clang/test/OpenMP/target_no_loop.c
@@ -0,0 +1,83 @@
+// 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(...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/224041
More information about the cfe-commits
mailing list