https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/205325
>From 13b93b2ec6e0d365105cf69e7d4a6134780280c3 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Thu, 27 Aug 2026 06:33:46 -0500 Subject: [PATCH 1/2] [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 1d39c1a8d4b77..122d065dd218c 100644 --- a/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp +++ b/flang/lib/Lower/OpenMP/DataSharingProcessor.cpp @@ -326,6 +326,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 d65924f2a8b3b..8296513ec559b 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 7001189b4c2ce..253793e1a8ec5 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -6592,11 +6592,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: @@ -6693,6 +6714,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(); } @@ -6702,9 +6735,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 e23cda41ded81..683422782d2c5 100644 --- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp +++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp @@ -2081,13 +2081,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) { @@ -2096,7 +2110,8 @@ static LogicalResult copyFirstPrivateVars( }); if (!needsFirstprivate) - return success(); + return emitPrivatizationBarrier(op, builder, moduleTranslation, + insertBarrier); llvm::BasicBlock *copyBlock = splitBB(builder, /*CreateBranch=*/true, "omp.private.copy"); @@ -2135,24 +2150,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 @@ -2161,9 +2170,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> @@ -2417,7 +2426,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()))) @@ -3334,7 +3343,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()))) @@ -3804,7 +3813,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()))) @@ -4103,10 +4112,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(); @@ -4699,7 +4708,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()))) @@ -4812,7 +4821,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(); @@ -4930,7 +4940,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()))) @@ -5169,7 +5179,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()); @@ -8641,10 +8651,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(); @@ -9515,7 +9525,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 04458af9654c3..4a60d3ceabcee 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 ca9de0321b9ee293901784427b6f50e6a05b1796 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> Date: Thu, 27 Aug 2026 07:19:57 -0500 Subject: [PATCH 2/2] [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 a761ac6bbf816..31cc5a4c86115 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. @@ -2443,6 +2475,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, @@ -2646,12 +2684,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); @@ -3708,6 +3760,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); @@ -6285,7 +6396,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())) @@ -6310,9 +6426,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); @@ -6418,11 +6536,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 7bdc79d86ea0a..1038c4e4cb6b4 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 253793e1a8ec5..d5354abe5b74a 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -1056,7 +1056,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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
