Skip to content

[clang][OpenMP] Emit lastprivate final copies in no-loop kernels #224045

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
nicebert wants to merge 1 commit into from
+156 −8
Open

[clang][OpenMP] Emit lastprivate final copies in no-loop kernels #224045

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 1 addition & 3 deletions clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1158,9 +1158,7 @@ bool CGOpenMPRuntimeGPU::canPromoteToNoLoop(
DKind == OMPD_target_teams_distribute_parallel_for_simd) &&
LangOpts.OpenMPTeamSubscription && LangOpts.OpenMPThreadSubscription &&
!D.hasClausesOfKind<OMPNumTeamsClause>() &&
!D.hasClausesOfKind<OMPReductionClause>() &&
!D.hasClausesOfKind<OMPLastprivateClause>() &&
!D.hasClausesOfKind<OMPLinearClause>();
!D.hasClausesOfKind<OMPReductionClause>();
}

std::pair<llvm::Value *, llvm::Value *>
Expand Down
57 changes: 53 additions & 4 deletions clang/lib/CodeGen/CGStmtOpenMP.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,29 @@ 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);
Expand Down Expand Up @@ -3772,6 +3795,10 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
CGF.EmitOMPPrivateLoopCounters(S, PrivateScope);
(void)PrivateScope.Privatize();

CodeGenFunction::OMPPrivateScope LastprivateScope(CGF);
(void)CGF.EmitOMPLastprivateClauseInit(S, LastprivateScope);
(void)LastprivateScope.Privatize();

if (isOpenMPTargetExecutionDirective(EKind))
CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);

Expand All @@ -3795,17 +3822,33 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
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>(),
/*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));
/*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;
}

Expand Down Expand Up @@ -6364,8 +6407,11 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
!isOpenMPTeamsDirective(S.getDirectiveKind()))
EmitOMPReductionClauseInit(S, LoopScope);

// Defer lastprivate init to the parallel region for no-loop
const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S);
HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
if (!NoLoopKernel)
HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);

EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
Expand Down Expand Up @@ -6501,7 +6547,10 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
CodeGenLoop);
}
}
if (isOpenMPSimdDirective(S.getDirectiveKind())) {
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)
Expand Down
55 changes: 54 additions & 1 deletion clang/test/OpenMP/target_no_loop.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,12 +45,37 @@ void no_loop_simd(int *array) {
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

Expand All @@ -69,6 +94,25 @@ void no_loop_nowait(int *array) {
// 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
Expand All @@ -78,4 +122,13 @@ void no_loop_nowait(int *array) {
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void

// SPMD-COUNT-3: _kernel_environment {{.*}} i8 0, i8 1, i8 2
// 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

// SPMD-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 2
48 changes: 48 additions & 0 deletions offload/test/offloading/target-no-loop.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,48 @@ int main(void) {
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;
}
Expand All @@ -88,4 +130,10 @@ int main(void) {
// 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
Loading