https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/205325
>From e5c4294ffa6e31df3d8950191f16cd14c3705dee Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Thu, 27 Aug 2026 06:33:46 -0500 Subject: [PATCH 1/3] [OpenMP] Honor barrier and last iteration flags in loop lowering Emitting loop-free kernels for no-loop target regions in Clang requires the shared OpenMP lowering to honor flags that reach it today and are then discarded. applyWorkshareLoop takes a NeedsBarrier flag, but the device path drops it, so a worksharing loop without nowait emits no barrier at its exit. The device path also never sets the canonical loop's last iteration variable, which the linear clause finalization reads. Forward the flag to applyWorkshareLoopTarget and compute the last iteration in the loop body, mirroring how the host runtime reports it. Auditing the surrounding lowering for the same class of problem turned up one more. The barrier that follows privatization is emitted as part of the firstprivate copy region, so a construct with lastprivate and no firstprivate never gets one. Emit it independently of the copy region. Flang skips it for taskloop, where the write-back already happens after the reads. --- .../lib/Lower/OpenMP/DataSharingProcessor.cpp | 7 ++ .../OpenMP/privatization-barrier.f90 | 38 ++++++++ flang/test/Lower/OpenMP/taskloop.f90 | 15 +++ .../llvm/Frontend/OpenMP/OMPIRBuilder.h | 16 ++-- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 42 +++++++- .../OpenMP/OpenMPToLLVMIRTranslation.cpp | 96 ++++++++++--------- mlir/test/Target/LLVMIR/omptarget-wsloop.mlir | 40 ++++++++ 7 files changed, 201 insertions(+), 53 deletions(-) create mode 100644 flang/test/Integration/OpenMP/privatization-barrier.f90 diff --git a/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp b/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp index 8d0d191058cfb..3e63c068e1cdb 100644 --- a/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp +++ b/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp @@ -372,6 +372,13 @@ bool DataSharingProcessor::needBarrier() { // initialization of firstprivate variables and post-update of lastprivate // variables. // Emit implicit barrier for linear clause in the OpenMPIRBuilder. + // Skip for taskloop: the write-back only happens after the reads are done. + + const auto *ompEval = eval.getIf<parser::OpenMPConstruct>(); + if (ompEval && llvm::omp::allTaskloopSet.test( + parser::omp::GetOmpDirectiveName(*ompEval).v)) + return false; + for (const semantics::Symbol *sym : allPrivatizedSymbols) { if (sym->test(semantics::Symbol::Flag::OmpLastPrivate) && (sym->test(semantics::Symbol::Flag::OmpFirstPrivate) || diff --git a/flang/test/Integration/OpenMP/privatization-barrier.f90 b/flang/test/Integration/OpenMP/privatization-barrier.f90 new file mode 100644 index 0000000000000..06ab26df2521e --- /dev/null +++ b/flang/test/Integration/OpenMP/privatization-barrier.f90 @@ -0,0 +1,38 @@ +!===----------------------------------------------------------------------===! +! This directory can be used to add Integration tests involving multiple +! stages of the compiler (for eg. from Fortran to LLVM IR). It should not +! contain executable tests. We should only add tests here sparingly and only +! if there is no other way to test. Repeat this message in each test that is +! added to this directory and sub-directories. +!===----------------------------------------------------------------------===! + +! RUN: %flang_fc1 -fopenmp -emit-llvm %s -o - | FileCheck %s +! RUN: %if amdgpu-registered-target %{ %flang_fc1 -triple amdgcn-amd-amdhsa -emit-llvm -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s %} + +subroutine lastprivate_allocatable_barrier_host + integer, allocatable :: a + integer :: i + !$omp parallel do lastprivate(a) + do i = 1, 10 + a = i + end do + !$omp end parallel do +end subroutine + +subroutine lastprivate_allocatable_barrier_device + integer, allocatable :: a + integer :: i + allocate(a) + !$omp target parallel do lastprivate(a) + do i = 1, 10 + a = i + end do + !$omp end target parallel do +end subroutine + +! CHECK-LABEL: define internal void @{{.*}}lastprivate_allocatable_barrier_{{(host|device)}} +! CHECK: call void @__kmpc_barrier +! CHECK-NEXT: br label %omp.wsloop.region +! CHECK: call void @__kmpc_barrier +! CHECK-NEXT: br label %omp_loop.after +! CHECK-LABEL: define{{.*}}void @{{.*}}lastprivate_allocatable_barrier_device diff --git a/flang/test/Lower/OpenMP/taskloop.f90 b/flang/test/Lower/OpenMP/taskloop.f90 index 94e4e2947fa71..c9ce984db7e8e 100644 --- a/flang/test/Lower/OpenMP/taskloop.f90 +++ b/flang/test/Lower/OpenMP/taskloop.f90 @@ -2,6 +2,9 @@ ! RUN: bbc -emit-hlfir %openmp_flags -fopenmp-version=50 -o - %s 2>&1 | FileCheck %s ! RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=50 -o - %s 2>&1 | FileCheck %s +! CHECK-LABEL: omp.private {type = firstprivate} +! CHECK-SAME: @[[FIRST_LAST_PRIVATE_X:.*]] : i32 + ! CHECK-LABEL: omp.private ! CHECK-SAME: {type = private} @[[LAST_PRIVATE_I:.*]] : i32 @@ -260,3 +263,15 @@ subroutine omp_taskloop_lastprivate() ! CHECK: omp.terminator !$omp end taskloop end subroutine omp_taskloop_lastprivate + +! CHECK-LABEL: func @_QPomp_taskloop_first_and_lastprivate() +subroutine omp_taskloop_first_and_lastprivate() + integer x + x = 0 + ! CHECK: omp.taskloop.context private(@[[FIRST_LAST_PRIVATE_X]] {{.*}}) { + !$omp taskloop firstprivate(x) lastprivate(x) + do i = 1, 100 + x = x + 1 + end do + !$omp end taskloop +end subroutine omp_taskloop_first_and_lastprivate diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 5f3da79cf53b1..de2cda0dd81ef 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -1178,15 +1178,19 @@ class OpenMPIRBuilder { /// \param CLI A descriptor of the canonical loop to workshare. /// \param AllocaIP An insertion point for Alloca instructions usable in the /// preheader of the loop. + /// \param NeedsBarrier Indicates whether a barrier must be inserted after + /// the loop. /// \param LoopType Information about type of loop worksharing. /// It corresponds to type of loop workshare OpenMP pragma. /// \param NoLoop If true, no-loop code is generated. + /// \param NeedsLastIter If true, the last iteration variable is emitted. /// /// \returns Point where to insert code after the workshare construct. - InsertPointTy applyWorkshareLoopTarget(DebugLoc DL, CanonicalLoopInfo *CLI, - InsertPointTy AllocaIP, - omp::WorksharingLoopType LoopType, - bool NoLoop); + InsertPointOrErrorTy + applyWorkshareLoopTarget(DebugLoc DL, CanonicalLoopInfo *CLI, + InsertPointTy AllocaIP, bool NeedsBarrier, + omp::WorksharingLoopType LoopType, bool NoLoop, + bool NeedsLastIter); /// Modifies the canonical loop to be a statically-scheduled workshare loop. /// @@ -1344,8 +1348,8 @@ class OpenMPIRBuilder { /// \param NoLoop If true, no-loop code is generated. /// \param HasDistSchedule Defines if the clause being lowered is /// dist_schedule as this is handled slightly differently - /// /// \param DistScheduleChunkSize The chunk size for dist_schedule loop + /// \param NeedsLastIter If true, the last iteration variable is emitted. /// /// \returns Point where to insert code after the workshare construct. LLVM_ABI InsertPointOrErrorTy applyWorkshareLoop( @@ -1358,7 +1362,7 @@ class OpenMPIRBuilder { omp::WorksharingLoopType LoopType = omp::WorksharingLoopType::ForStaticLoop, bool NoLoop = false, bool HasDistSchedule = false, - Value *DistScheduleChunkSize = nullptr); + Value *DistScheduleChunkSize = nullptr, bool NeedsLastIter = false); /// Tile a loop nest. /// diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 3f80c858d033e..615a63fa62181 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -6595,11 +6595,32 @@ static void workshareLoopTargetCallback( CLI->invalidate(); } -OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::applyWorkshareLoopTarget( +OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::applyWorkshareLoopTarget( DebugLoc DL, CanonicalLoopInfo *CLI, InsertPointTy AllocaIP, - WorksharingLoopType LoopType, bool NoLoop) { + bool NeedsBarrier, WorksharingLoopType LoopType, bool NoLoop, + bool NeedsLastIter) { uint32_t SrcLocStrSize; Constant *SrcLocStr = getOrCreateSrcLocStr(DL, SrcLocStrSize); + + // Mirrors host runtime reporting of last iteration by in-body computation. + if (NeedsLastIter) { + Type *I32Type = Type::getInt32Ty(M.getContext()); + Builder.restoreIP(AllocaIP); + AllocaInst *PLastIter = + Builder.CreateAlloca(I32Type, nullptr, "p.lastiter"); + Builder.CreateStore(ConstantInt::get(I32Type, 0), PLastIter); + CLI->setLastIter(PLastIter); + + Builder.SetInsertPoint(CLI->getBody(), + CLI->getBody()->getFirstInsertionPt()); + Value *TripCount = CLI->getTripCount(); + Value *LastIter = + Builder.CreateSub(TripCount, ConstantInt::get(TripCount->getType(), 1)); + Value *IsLast = + Builder.CreateICmpEQ(CLI->getIndVar(), LastIter, "omp.is_last_iter"); + Builder.CreateStore(Builder.CreateZExt(IsLast, I32Type), PLastIter); + } + IdentFlag Flag = IdentFlag(0); switch (LoopType) { case WorksharingLoopType::ForStaticLoop: @@ -6696,6 +6717,18 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::applyWorkshareLoopTarget( LoopType, NoLoop); }; addOutlineInfo(std::move(OI)); + + if (NeedsBarrier) { + Builder.SetInsertPoint(CLI->getExit(), + CLI->getExit()->getTerminator()->getIterator()); + InsertPointOrErrorTy BarrierIP = + createBarrier(LocationDescription(Builder.saveIP(), DL), + omp::Directive::OMPD_for, /* ForceSimpleCall */ false, + /* CheckCancelFlag*/ false); + if (!BarrierIP) + return BarrierIP.takeError(); + } + return CLI->getAfterIP(); } @@ -6705,9 +6738,10 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::applyWorkshareLoop( bool HasSimdModifier, bool HasMonotonicModifier, bool HasNonmonotonicModifier, bool HasOrderedClause, WorksharingLoopType LoopType, bool NoLoop, bool HasDistSchedule, - Value *DistScheduleChunkSize) { + Value *DistScheduleChunkSize, bool NeedsLastIter) { if (Config.isTargetDevice()) - return applyWorkshareLoopTarget(DL, CLI, AllocaIP, LoopType, NoLoop); + return applyWorkshareLoopTarget(DL, CLI, AllocaIP, NeedsBarrier, LoopType, + NoLoop, NeedsLastIter); OMPScheduleType EffectiveScheduleType = computeOpenMPScheduleType( SchedKind, ChunkSize, HasSimdModifier, HasMonotonicModifier, HasNonmonotonicModifier, HasOrderedClause, DistScheduleChunkSize); diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp index 1cbad3312235c..ab3968001e794 100644 --- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp +++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp @@ -2083,13 +2083,27 @@ static bool opIsInSingleThread(mlir::Operation *op) { return false; } -static LogicalResult copyFirstPrivateVars( - mlir::Operation *op, llvm::IRBuilderBase &builder, - LLVM::ModuleTranslation &moduleTranslation, - SmallVectorImpl<llvm::Value *> &moldVars, - ArrayRef<llvm::Value *> llvmPrivateVars, - SmallVectorImpl<omp::PrivateClauseOp> &privateDecls, bool insertBarrier, - llvm::DenseMap<Value, Value> *mappedPrivateVars = nullptr) { +static LogicalResult +emitPrivatizationBarrier(mlir::Operation *op, llvm::IRBuilderBase &builder, + LLVM::ModuleTranslation &moduleTranslation, + bool insertBarrier) { + if (!insertBarrier || opIsInSingleThread(op)) + return success(); + + llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder(); + llvm::OpenMPIRBuilder::InsertPointOrErrorTy res = + ompBuilder->createBarrier(builder, llvm::omp::OMPD_barrier); + return handleError(res, *op); +} + +static LogicalResult +completePrivateVars(mlir::Operation *op, llvm::IRBuilderBase &builder, + LLVM::ModuleTranslation &moduleTranslation, + SmallVectorImpl<llvm::Value *> &moldVars, + ArrayRef<llvm::Value *> llvmPrivateVars, + SmallVectorImpl<omp::PrivateClauseOp> &privateDecls, + bool insertBarrier, + llvm::DenseMap<Value, Value> *mappedPrivateVars = nullptr) { // Apply copy region for firstprivate. bool needsFirstprivate = llvm::any_of(privateDecls, [](omp::PrivateClauseOp &privOp) { @@ -2098,7 +2112,8 @@ static LogicalResult copyFirstPrivateVars( }); if (!needsFirstprivate) - return success(); + return emitPrivatizationBarrier(op, builder, moduleTranslation, + insertBarrier); llvm::BasicBlock *copyBlock = splitBB(builder, /*CreateBranch=*/true, "omp.private.copy"); @@ -2137,24 +2152,18 @@ static LogicalResult copyFirstPrivateVars( moduleTranslation.forgetMapping(copyRegion); } - if (insertBarrier && !opIsInSingleThread(op)) { - llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder(); - llvm::OpenMPIRBuilder::InsertPointOrErrorTy res = - ompBuilder->createBarrier(builder, llvm::omp::OMPD_barrier); - if (failed(handleError(res, *op))) - return failure(); - } - - return success(); + return emitPrivatizationBarrier(op, builder, moduleTranslation, + insertBarrier); } -static LogicalResult copyFirstPrivateVars( - mlir::Operation *op, llvm::IRBuilderBase &builder, - LLVM::ModuleTranslation &moduleTranslation, - SmallVectorImpl<mlir::Value> &mlirPrivateVars, - ArrayRef<llvm::Value *> llvmPrivateVars, - SmallVectorImpl<omp::PrivateClauseOp> &privateDecls, bool insertBarrier, - llvm::DenseMap<Value, Value> *mappedPrivateVars = nullptr) { +static LogicalResult +completePrivateVars(mlir::Operation *op, llvm::IRBuilderBase &builder, + LLVM::ModuleTranslation &moduleTranslation, + SmallVectorImpl<mlir::Value> &mlirPrivateVars, + ArrayRef<llvm::Value *> llvmPrivateVars, + SmallVectorImpl<omp::PrivateClauseOp> &privateDecls, + bool insertBarrier, + llvm::DenseMap<Value, Value> *mappedPrivateVars = nullptr) { llvm::SmallVector<llvm::Value *> moldVars(mlirPrivateVars.size()); llvm::transform(mlirPrivateVars, moldVars.begin(), [&](mlir::Value mlirVar) { // map copyRegion rhs arg @@ -2163,9 +2172,9 @@ static LogicalResult copyFirstPrivateVars( assert(moldVar); return moldVar; }); - return copyFirstPrivateVars(op, builder, moduleTranslation, moldVars, - llvmPrivateVars, privateDecls, insertBarrier, - mappedPrivateVars); + return completePrivateVars(op, builder, moduleTranslation, moldVars, + llvmPrivateVars, privateDecls, insertBarrier, + mappedPrivateVars); } template <typename T> @@ -2419,7 +2428,7 @@ convertOmpScope(omp::ScopeOp &scopeOp, llvm::IRBuilderBase &builder, .failed()) return llvm::make_error<PreviouslyReportedError>(); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( scopeOp, builder, moduleTranslation, privateVarsInfo.mlirVars, privateVarsInfo.llvmVars, privateVarsInfo.privatizers, scopeOp.getPrivateNeedsBarrier()))) @@ -3336,7 +3345,7 @@ convertOmpTaskOp(omp::TaskOp taskOp, llvm::IRBuilderBase &builder, // firstprivate copy region setInsertPointForPossiblyEmptyBlock(builder, copyBlock); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( taskOp, builder, moduleTranslation, privateVarsInfo.mlirVars, taskStructMgr.getLLVMPrivateVarGEPs(), privateVarsInfo.privatizers, taskOp.getPrivateNeedsBarrier()))) @@ -3806,7 +3815,7 @@ convertOmpTaskloopContextOp(omp::TaskloopContextOp contextOp, // firstprivate copy region setInsertPointForPossiblyEmptyBlock(builder, copyBlock); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( contextOp, builder, moduleTranslation, privateVarsInfo.mlirVars, taskStructMgr.getLLVMPrivateVarGEPs(), privateVarsInfo.privatizers, contextOp.getPrivateNeedsBarrier()))) @@ -4105,10 +4114,10 @@ convertOmpTaskloopContextOp(omp::TaskloopContextOp contextOp, // through a stack allocated structure. } - if (failed(copyFirstPrivateVars(contextOp.getOperation(), builder, - moduleTranslation, srcGEPs, destGEPs, - privateVarsInfo.privatizers, - contextOp.getPrivateNeedsBarrier()))) + if (failed(completePrivateVars(contextOp.getOperation(), builder, + moduleTranslation, srcGEPs, destGEPs, + privateVarsInfo.privatizers, + contextOp.getPrivateNeedsBarrier()))) return llvm::make_error<PreviouslyReportedError>(); return builder.saveIP(); @@ -4701,7 +4710,7 @@ convertOmpWsloop(Operation &opInst, llvm::IRBuilderBase &builder, .failed()) return failure(); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( wsloopOp, builder, moduleTranslation, privateVarsInfo.mlirVars, privateVarsInfo.llvmVars, privateVarsInfo.privatizers, wsloopOp.getPrivateNeedsBarrier()))) @@ -4814,7 +4823,8 @@ convertOmpWsloop(Operation &opInst, llvm::IRBuilderBase &builder, convertToScheduleKind(schedule), chunk, isSimd, scheduleMod == omp::ScheduleModifier::monotonic, scheduleMod == omp::ScheduleModifier::nonmonotonic, isOrdered, - workshareLoopType, noLoopMode, hasDistSchedule, distScheduleChunk); + workshareLoopType, noLoopMode, hasDistSchedule, distScheduleChunk, + !wsloopOp.getLinearVars().empty()); if (failed(handleError(wsloopIP, opInst))) return failure(); @@ -4932,7 +4942,7 @@ convertOmpParallel(omp::ParallelOp opInst, llvm::IRBuilderBase &builder, .failed()) return llvm::make_error<PreviouslyReportedError>(); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( opInst, builder, moduleTranslation, privateVarsInfo.mlirVars, privateVarsInfo.llvmVars, privateVarsInfo.privatizers, opInst.getPrivateNeedsBarrier()))) @@ -5171,7 +5181,7 @@ convertOmpSimd(Operation &opInst, llvm::IRBuilderBase &builder, .failed()) return failure(); - // No call to copyFirstPrivateVars because FIRSTPRIVATE is not allowed for + // No call to completePrivateVars because FIRSTPRIVATE is not allowed for // SIMD. assert(afterAllocas.get()->getSinglePredecessor()); @@ -8645,10 +8655,10 @@ convertOmpDistribute(Operation &opInst, llvm::IRBuilderBase &builder, .failed()) return llvm::make_error<PreviouslyReportedError>(); - if (failed(copyFirstPrivateVars( - distributeOp, builder, moduleTranslation, privVarsInfo.mlirVars, - privVarsInfo.llvmVars, privVarsInfo.privatizers, - distributeOp.getPrivateNeedsBarrier()))) + if (failed(completePrivateVars(distributeOp, builder, moduleTranslation, + privVarsInfo.mlirVars, privVarsInfo.llvmVars, + privVarsInfo.privatizers, + distributeOp.getPrivateNeedsBarrier()))) return llvm::make_error<PreviouslyReportedError>(); llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder(); @@ -9529,7 +9539,7 @@ convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder, .failed()) return llvm::make_error<PreviouslyReportedError>(); - if (failed(copyFirstPrivateVars( + if (failed(completePrivateVars( targetOp, builder, moduleTranslation, privateVarsInfo.mlirVars, privateVarsInfo.llvmVars, privateVarsInfo.privatizers, targetOp.getPrivateNeedsBarrier(), &mappedPrivateVars))) diff --git a/mlir/test/Target/LLVMIR/omptarget-wsloop.mlir b/mlir/test/Target/LLVMIR/omptarget-wsloop.mlir index acbb3ec916113..a14efc9ae84f8 100644 --- a/mlir/test/Target/LLVMIR/omptarget-wsloop.mlir +++ b/mlir/test/Target/LLVMIR/omptarget-wsloop.mlir @@ -29,6 +29,32 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo } llvm.return } + + llvm.func @target_wsloop_nowait(%arg0: !llvm.ptr) attributes {omp.declare_target = #omp.declaretarget<device_type = any, capture_clause = to>} { + %loop_ub = llvm.mlir.constant(9 : i32) : i32 + %loop_lb = llvm.mlir.constant(0 : i32) : i32 + %loop_step = llvm.mlir.constant(1 : i32) : i32 + omp.wsloop nowait { + omp.loop_nest (%loop_cnt) : i32 = (%loop_lb) to (%loop_ub) inclusive step (%loop_step) { + %gep = llvm.getelementptr %arg0[0, %loop_cnt] : (!llvm.ptr, i32) -> !llvm.ptr, !llvm.array<10 x i32> + llvm.store %loop_cnt, %gep : i32, !llvm.ptr + omp.yield + } + } + llvm.return + } + + llvm.func @target_wsloop_linear(%arg0: !llvm.ptr) attributes {omp.declare_target = #omp.declaretarget<device_type = any, capture_clause = to>} { + %loop_ub = llvm.mlir.constant(9 : i32) : i32 + %loop_lb = llvm.mlir.constant(0 : i32) : i32 + %loop_step = llvm.mlir.constant(1 : i32) : i32 + omp.wsloop linear(%arg0 : !llvm.ptr = %loop_step : i32) linear_var_types([i32]) { + omp.loop_nest (%loop_cnt) : i32 = (%loop_lb) to (%loop_ub) inclusive step (%loop_step) { + omp.yield + } + } + llvm.return + } } // CHECK: define hidden void @[[FUNC0:.*]](ptr %[[ARG0:.*]]) @@ -38,6 +64,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo // CHECK: store ptr %[[ARG0]], ptr addrspace(5) %[[GEP]], align 8 // CHECK: %[[NUM_THREADS:.*]] = call i32 @omp_get_num_threads() // CHECK: call void @__kmpc_for_static_loop_4u(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), ptr @[[LOOP_BODY_FN:.*]], ptr %[[STRUCTARG_ASCAST]], i32 10, i32 %[[NUM_THREADS]], i32 0, i8 0) +// CHECK: call void @__kmpc_barrier{{.*}}[[IMPL_WSBAR:@[0-9]+]] to ptr // CHECK: define internal void @[[LOOP_BODY_FN]](i32 %[[LOOP_CNT:.*]], ptr %[[LOOP_BODY_ARG:.*]]) // CHECK: %[[GEP2:.*]] = getelementptr { ptr }, ptr %[[LOOP_BODY_ARG]], i32 0, i32 0 @@ -49,3 +76,16 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo // CHECK: call void @__kmpc_for_static_loop_4u(ptr addrspacecast (ptr addrspace(1) @[[GLOB2:[0-9]+]] to ptr), ptr @[[LOOP_EMPTY_BODY_FN:.*]], ptr null, i32 10, i32 %[[NUM_THREADS:.*]], i32 0, i8 0) // CHECK: define internal void @[[LOOP_EMPTY_BODY_FN]](i32 %[[LOOP_CNT:.*]]) + +// CHECK: define hidden void @target_wsloop_nowait(ptr %{{.*}}) +// CHECK: call void @__kmpc_for_static_loop_4u +// CHECK-NOT: call void @__kmpc_barrier{{.*}}[[IMPL_WSBAR]] to ptr + +// CHECK: define hidden void @target_wsloop_linear(ptr %{{.*}}) +// CHECK: store i32 0, ptr{{.*}} %p.lastiter +// CHECK: call void @__kmpc_for_static_loop_4u +// CHECK: call void @__kmpc_barrier{{.*}}[[IMPL_WSBAR]] to ptr + +// CHECK: define internal void @target_wsloop_linear +// CHECK: %omp.is_last_iter = icmp eq +// CHECK: store{{.*}}_p.lastiter >From 0078cb54ff35b2c91eb857bb25c5b26eac86eee0 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Thu, 27 Aug 2026 07:19:57 -0500 Subject: [PATCH 2/3] [OpenMP] Add no-loop SPMD kernel emission to Clang 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 SPMD_NO_LOOP tag is emitted with or without the IRBuilder and makes the runtime size the launch to one thread per iteration. The loop-free body itself needs the IRBuilder and returns early from the regular worksharing path, so cases like lastprivate get their own handling to stay equivalent to it. 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 | 22 ++- clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 4 + clang/lib/CodeGen/CGStmtOpenMP.cpp | 146 +++++++++++++++++-- clang/lib/CodeGen/CodeGenFunction.h | 7 +- clang/test/OpenMP/irbuilder_target_no_loop.c | 146 +++++++++++++++++++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +- offload/test/offloading/target-no-loop.c | 139 ++++++++++++++++++ 8 files changed, 456 insertions(+), 16 deletions(-) create mode 100644 clang/test/OpenMP/irbuilder_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 90ae21283d031..fd8ee7097b46a 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,17 @@ 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>(); +} + 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 a05a38ca18e74..27a3ed0543a43 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -66,6 +66,38 @@ static bool canEmitGPUFusedDistSchedule(const CodeGenModule &CGM, !S.getSingleClause<OMPOrderedClause>(); } +static bool isLoopCounter(const OMPLoopDirective &S, const Expr *E) { + const VarDecl *Canonical = + cast<VarDecl>(cast<DeclRefExpr>(E->IgnoreParenImpCasts())->getDecl()) + ->getCanonicalDecl(); + for (const Expr *C : S.counters()) + if (cast<VarDecl>(cast<DeclRefExpr>(C)->getDecl())->getCanonicalDecl() == + Canonical) + return true; + return false; +} + +static bool +hasLoopLastprivateVariable(const OMPLoopDirective &S, + llvm::function_ref<bool(const Expr *)> Pred) { + for (const auto *C : S.getClausesOfKind<OMPLastprivateClause>()) { + for (const Expr *E : C->varlist()) { + if (Pred(E)) + return true; + } + } + return false; +} + +static bool canEmitGPUNoLoopKernel(CodeGenModule &CGM, + const OMPLoopDirective &S) { + const auto *D = dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&S); + return CGM.getLangOpts().OpenMPIRBuilder && 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. @@ -2447,6 +2479,12 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) { emitCapturedStmtCall(*this, DistanceClosure, {CountAddr.getPointer()}); llvm::Value *DistVal = Builder.CreateLoad(CountAddr, ".count"); + // Privatize pending lastprivate variables after trip count is computed. + if (auto Privatize = OMPCanonicalLoopPendingPrivatization) { + OMPCanonicalLoopPendingPrivatization = nullptr; + Privatize(); + } + // Emit the loop structure. llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder(); auto BodyGen = [&, this](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP, @@ -2650,12 +2688,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); @@ -3712,6 +3764,65 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF, HasCancel = D->hasCancel(); } CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel); + + CodeGenModule &CGM = CGF.CGM; + if (canEmitGPUNoLoopKernel(CGM, S)) { + // Privatize pending variables, deferring lastprivate privatization until + // after the trip count is computed from the originals. + CodeGenFunction::OMPPrivateScope PrivateScope(CGF); + CGF.EmitOMPPrivateClause(S, PrivateScope); + CGF.EmitOMPPrivateLoopCounters(S, PrivateScope, /*OnlyUnresolved=*/true); + (void)PrivateScope.Privatize(); + + CodeGenFunction::OMPPrivateScope LastprivateScope(CGF); + (void)CGF.EmitOMPLastprivateClauseInit(S, LastprivateScope); + auto PendingPrivatizations = [&] { + (void)LastprivateScope.Privatize(); + if (isOpenMPTargetExecutionDirective(EKind)) + CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S); + }; + CGF.OMPCanonicalLoopPendingPrivatization = PendingPrivatizations; + + const Stmt *Inner = S.getRawStmt(); + llvm::CanonicalLoopInfo *CLI = + CGF.EmitOMPCollapsedCanonicalLoopNest(Inner, 1); + assert(!CGF.OMPCanonicalLoopPendingPrivatization && + "canonical loop did not apply pending privatization"); + + llvm::OpenMPIRBuilder::InsertPointTy AllocaIP( + CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator()); + llvm::OpenMPIRBuilder &OMPBuilder = + CGM.getOpenMPRuntime().getOMPBuilder(); + bool NeedsLastprivateFinalCopy = hasLoopLastprivateVariable( + S, [&S](const Expr *E) { return !isLoopCounter(S, E); }); + + cantFail(OMPBuilder.applyWorkshareLoop( + CGF.Builder.getCurrentDebugLocation(), CLI, AllocaIP, + /*NeedsBarrier=*/!S.getSingleClause<OMPNowaitClause>() || + NeedsLastprivateFinalCopy, + 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, + /*NeedsLastIter=*/NeedsLastprivateFinalCopy)); + + // Emit final copy for lastprivate variables, excluding counters deferred + // to the distribute level. + if (NeedsLastprivateFinalCopy) { + llvm::Value *LastIter = CLI->getLastIter(); + assert(LastIter && "workshare loop did not publish last iteration"); + CGF.EmitOMPLastprivateClauseFinal( + S, /*NoFinals=*/true, + CGF.Builder.CreateIsNotNull(CGF.Builder.CreateAlignedLoad( + CGF.Int32Ty, LastIter, CharUnits::fromQuantity(4), + ".omp.is_last"))); + } + return; + } + CGF.EmitOMPWorksharingLoop(S, S.getPrevEnsureUpperBound(), emitDistributeParallelForInnerBounds, emitDistributeParallelForDispatchBounds); @@ -6289,7 +6400,12 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, !isOpenMPParallelDirective(S.getDirectiveKind()) && !isOpenMPTeamsDirective(S.getDirectiveKind())) EmitOMPReductionClauseInit(S, LoopScope); - HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope); + + // Defer lastprivate init to the parallel region for no-loop + const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S); + if (!NoLoopKernel) + HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope); + EmitOMPPrivateLoopCounters(S, LoopScope); (void)LoopScope.Privatize(); if (isOpenMPTargetExecutionDirective(S.getDirectiveKind())) @@ -6314,9 +6430,11 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, 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); @@ -6422,11 +6540,17 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, CodeGenLoop); } } - if (isOpenMPSimdDirective(S.getDirectiveKind())) { - EmitOMPSimdFinal(S, [IL, &S](CodeGenFunction &CGF) { - return CGF.Builder.CreateIsNotNull( - CGF.EmitLoadOfScalar(IL, S.getBeginLoc())); - }); + if (isOpenMPSimdDirective(S.getDirectiveKind()) || + (NoLoopKernel && hasLoopLastprivateVariable(S, [&S](const Expr *E) { + return isLoopCounter(S, E); + }))) { + 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.h b/clang/lib/CodeGen/CodeGenFunction.h index 3c8188c4cefdd..006b5b2747fca 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -316,6 +316,10 @@ class CodeGenFunction : public CodeGenTypeCache { /// nest would extend. SmallVector<llvm::CanonicalLoopInfo *, 4> OMPLoopNestStack; + /// Privatization to be applied for canonical loops once the trip count has + /// been computed. + llvm::function_ref<void()> OMPCanonicalLoopPendingPrivatization = nullptr; + /// Stack to track the controlled convergence tokens. SmallVector<llvm::ConvergenceControlInst *, 4> ConvergenceTokenStack; @@ -4131,7 +4135,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/irbuilder_target_no_loop.c b/clang/test/OpenMP/irbuilder_target_no_loop.c new file mode 100644 index 0000000000000..7f9c17f3a585a --- /dev/null +++ b/clang/test/OpenMP/irbuilder_target_no_loop.c @@ -0,0 +1,146 @@ +// 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: -fopenmp-enable-irbuilder -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: -fopenmp-assume-teams-oversubscription \ +// RUN: -fopenmp-assume-threads-oversubscription \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ +// RUN: --check-prefixes=LAUNCH --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-enable-irbuilder -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: -fopenmp-enable-irbuilder -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: -fopenmp-enable-irbuilder -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_lastprivate_counter(int *array) { + int i; +#pragma omp target teams distribute parallel for lastprivate(i) + for (i = 0; i < 1024; ++i) + array[i] = i + 1; +} + +void no_loop_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 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 no_loop_lastprivate_scalar_nowait(int *array) { + int last = 0; +#pragma omp target teams distribute parallel for lastprivate(last) nowait + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; + last = i; + } +} + +// 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: @__kmpc_barrier +// 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: @__kmpc_barrier +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_counter{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: store i32 1024, ptr %i +// NOLOOP-NEXT: @__kmpc_free_shared(ptr %i{{.*}}) +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}lastprivate_counter{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_barrier +// NOLOOP: omp_loop.after: +// NOLOOP-NEXT: ret void + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_scalar{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}}) +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}lastprivate_scalar{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_barrier +// NOLOOP: store {{.*}}, ptr %last. +// NOLOOP-NEXT: %.omp.lastprivate.done + +// 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 + +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_scalar_nowait{{.*}}) +// NOLOOP: omp.loop.exit: +// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}}) +// NOLOOP-NEXT: ret void +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}lastprivate_scalar_nowait{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_barrier +// NOLOOP: store {{.*}}, ptr %last. +// NOLOOP-NEXT: %.omp.lastprivate.done + +// LAUNCH-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 6 + +// SPMD-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 2 diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 615a63fa62181..5df890b58f38c 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -1055,7 +1055,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..fc1a86239f075 --- /dev/null +++ b/offload/test/offloading/target-no-loop.c @@ -0,0 +1,139 @@ +// clang-format off +// C counterpart of fortran/target-no-loop.f90. + +// RUN: %libomptarget-compile-generic -O3 -fopenmp-assume-threads-oversubscription -fopenmp-assume-teams-oversubscription -fopenmp-enable-irbuilder +// 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; + + // No-loop kernel with a lastprivate variable (final value from thread executing last iteration) + for (int i = 0; i < 1024; ++i) + array[i] = 1; + + int last = -1; +#pragma omp target teams distribute parallel for lastprivate(last) + for (int i = 0; i < 1024; ++i) { + array[i] = i + 1; + last = i; + } + errors += check_errors(array); +#ifdef ASSERT_LASTPRIVATE + if (last != 1023) + ++errors; +#endif + + // No-loop kernel with a lastprivate loop counter (final value from evaluation at distribution level) + for (int i = 0; i < 1024; ++i) + array[i] = 1; + int iv = -1; +#pragma omp target teams distribute parallel for lastprivate(iv) + for (iv = 0; iv < 1024; ++iv) + array[iv] = iv + 1; + errors += check_errors(array); +#ifdef ASSERT_LASTPRIVATE + if (iv != 1024) + ++errors; +#endif + + // No-loop simd kernel with a lastprivate loop counter + for (int i = 0; i < 1024; ++i) + array[i] = 1; + int simd_iv = -1; +#pragma omp target teams distribute parallel for simd lastprivate(simd_iv) + for (simd_iv = 0; simd_iv < 1024; ++simd_iv) + array[simd_iv] = simd_iv + 1; + errors += check_errors(array); +#ifdef ASSERT_LASTPRIVATE + if (simd_iv != 1024) + ++errors; +#endif + + 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: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}} +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}} +// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}} +// CHECK: number of errors: 0 >From d43cd065b76caf5bd1e5d22a34227059e7c44bd5 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Tue, 8 Sep 2026 07:22:50 -0500 Subject: [PATCH 3/3] [OpenMP] Emit no-loop kernels without the IRBuilder flag Clang could only emit the loop-free body of a no-loop kernel under -fopenmp-enable-irbuilder, since the OMPCanonicalLoop node it consumed is built by the parser only under that flag. Rebuild the same pieces in the no-loop branch instead, the trip count, the canonical loop and the body, and share the iteration space setup with the worksharing and distribute paths. Finalization follows: it now triggers on the builder holding outlining work for the function rather than on the flag, which is what gets the parallel region outlined in a default build. The tests lose the flag and irbuilder_target_no_loop.c is renamed accordingly. --- clang/lib/CodeGen/CGStmtOpenMP.cpp | 119 ++++++++++-------- clang/lib/CodeGen/CodeGenFunction.cpp | 18 +-- clang/lib/CodeGen/CodeGenFunction.h | 4 - ...lder_target_no_loop.c => target_no_loop.c} | 16 +-- .../llvm/Frontend/OpenMP/OMPIRBuilder.h | 8 ++ offload/test/offloading/target-no-loop.c | 2 +- 6 files changed, 86 insertions(+), 81 deletions(-) rename clang/test/OpenMP/{irbuilder_target_no_loop.c => target_no_loop.c} (86%) diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 27a3ed0543a43..a3ac4b1645c7a 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -92,7 +92,7 @@ hasLoopLastprivateVariable(const OMPLoopDirective &S, static bool canEmitGPUNoLoopKernel(CodeGenModule &CGM, const OMPLoopDirective &S) { const auto *D = dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&S); - return CGM.getLangOpts().OpenMPIRBuilder && S.getLoopsNumber() == 1 && + return S.getLoopsNumber() == 1 && CGM.getOpenMPRuntime().canPromoteToNoLoop(S) && !S.getSingleClause<OMPScheduleClause>() && !S.getSingleClause<OMPDistScheduleClause>() && !(D && D->hasCancel()); @@ -2405,6 +2405,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"); @@ -2479,18 +2497,8 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) { emitCapturedStmtCall(*this, DistanceClosure, {CountAddr.getPointer()}); llvm::Value *DistVal = Builder.CreateLoad(CountAddr, ".count"); - // Privatize pending lastprivate variables after trip count is computed. - if (auto Privatize = OMPCanonicalLoopPendingPrivatization) { - OMPCanonicalLoopPendingPrivatization = nullptr; - Privatize(); - } - // 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(); @@ -2501,14 +2509,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. @@ -3745,6 +3750,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, @@ -3767,27 +3788,37 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF, CodeGenModule &CGM = CGF.CGM; if (canEmitGPUNoLoopKernel(CGM, S)) { - // Privatize pending variables, deferring lastprivate privatization until - // after the trip count is computed from the originals. + // 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, /*OnlyUnresolved=*/true); + CGF.EmitOMPPrivateLoopCounters(S, PrivateScope); (void)PrivateScope.Privatize(); CodeGenFunction::OMPPrivateScope LastprivateScope(CGF); (void)CGF.EmitOMPLastprivateClauseInit(S, LastprivateScope); - auto PendingPrivatizations = [&] { - (void)LastprivateScope.Privatize(); - if (isOpenMPTargetExecutionDirective(EKind)) - CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S); - }; - CGF.OMPCanonicalLoopPendingPrivatization = PendingPrivatizations; + (void)LastprivateScope.Privatize(); - const Stmt *Inner = S.getRawStmt(); + 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 = - CGF.EmitOMPCollapsedCanonicalLoopNest(Inner, 1); - assert(!CGF.OMPCanonicalLoopPendingPrivatization && - "canonical loop did not apply pending privatization"); + 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()); @@ -3901,19 +3932,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(); @@ -4008,6 +4027,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. @@ -6329,19 +6349,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(); @@ -6427,6 +6435,7 @@ 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(); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 7e7f9a072f765..5eaab8e42c5d1 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -103,16 +103,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 006b5b2747fca..26917a6dc0a8c 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -316,10 +316,6 @@ class CodeGenFunction : public CodeGenTypeCache { /// nest would extend. SmallVector<llvm::CanonicalLoopInfo *, 4> OMPLoopNestStack; - /// Privatization to be applied for canonical loops once the trip count has - /// been computed. - llvm::function_ref<void()> OMPCanonicalLoopPendingPrivatization = nullptr; - /// Stack to track the controlled convergence tokens. SmallVector<llvm::ConvergenceControlInst *, 4> ConvergenceTokenStack; diff --git a/clang/test/OpenMP/irbuilder_target_no_loop.c b/clang/test/OpenMP/target_no_loop.c similarity index 86% rename from clang/test/OpenMP/irbuilder_target_no_loop.c rename to clang/test/OpenMP/target_no_loop.c index 7f9c17f3a585a..d754331f3f87b 100644 --- a/clang/test/OpenMP/irbuilder_target_no_loop.c +++ b/clang/test/OpenMP/target_no_loop.c @@ -8,35 +8,27 @@ // RUN: -fopenmp-host-ir-file-path %t-host.bc \ // RUN: -fopenmp-assume-teams-oversubscription \ // RUN: -fopenmp-assume-threads-oversubscription \ -// RUN: -fopenmp-enable-irbuilder -emit-llvm %s -o - | FileCheck %s \ +// 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: -fopenmp-assume-teams-oversubscription \ -// RUN: -fopenmp-assume-threads-oversubscription \ // RUN: -emit-llvm %s -o - | FileCheck %s \ -// RUN: --check-prefixes=LAUNCH --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-enable-irbuilder -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: -fopenmp-enable-irbuilder -emit-llvm %s -o - | FileCheck %s \ +// 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: -fopenmp-enable-irbuilder -emit-llvm %s -o - | FileCheck %s \ +// RUN: -emit-llvm %s -o - | FileCheck %s \ // RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u // expected-no-diagnostics @@ -141,6 +133,4 @@ void no_loop_lastprivate_scalar_nowait(int *array) { // NOLOOP: store {{.*}}, ptr %last. // NOLOOP-NEXT: %.omp.lastprivate.done -// LAUNCH-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 6 - // SPMD-COUNT-6: _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 de2cda0dd81ef..034ed40ce93b6 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/offload/test/offloading/target-no-loop.c b/offload/test/offloading/target-no-loop.c index fc1a86239f075..ae6c31c966b19 100644 --- a/offload/test/offloading/target-no-loop.c +++ b/offload/test/offloading/target-no-loop.c @@ -1,7 +1,7 @@ // clang-format off // C counterpart of fortran/target-no-loop.f90. -// RUN: %libomptarget-compile-generic -O3 -fopenmp-assume-threads-oversubscription -fopenmp-assume-teams-oversubscription -fopenmp-enable-irbuilder +// 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
