Skip to content
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

Force inline par_for_inner #967

Merged
merged 3 commits into from
Oct 24, 2023
Merged

Force inline par_for_inner #967

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
081b14c
INLINE to FORCEINLINE on all par_for_inner functions
jdolence Oct 24, 2023
5145bd8
update changelog
jdolence Oct 24, 2023
790cebb
update copyright
jdolence Oct 24, 2023
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
1 change: 1 addition & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@
- [[PR 890]](https://github.com/parthenon-hpc-lab/parthenon/pull/890) Fix bugs in sparse communication and prolongation

### Infrastructure (changes irrelevant to downstream codes)
- [[PR 967]](https://github.com/parthenon-hpc-lab/parthenon/pull/967) Change INLINE to FORCEINLINE on par_for_inner overloads
- [[PR 938]](https://github.com/parthenon-hpc-lab/parthenon/pull/938) Restructure buffer packing/unpacking kernel hierarchical parallelism
- [[PR 944]](https://github.com/parthenon-hpc-lab/parthenon/pull/944) Move sparse pack identifier creation to descriptor
- [[PR 904]](https://github.com/parthenon-hpc-lab/parthenon/pull/904) Move to prolongation/restriction in one for AMR and communicate non-cell centered fields
Expand Down
54 changes: 27 additions & 27 deletions src/kokkos_abstraction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -707,7 +707,7 @@ inline void par_for_outer(OuterLoopPatternTeams, const std::string &name,

// Inner parallel loop using TeamThreadRange
template <typename Function>
KOKKOS_INLINE_FUNCTION void
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ll, const int lu,
const int ml, const int mu, const int nl, const int nu, const int kl,
const int ku, const int jl, const int ju, const int il, const int iu,
Expand Down Expand Up @@ -741,7 +741,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ll, const i
});
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ml, const int mu,
const int nl, const int nu, const int kl, const int ku, const int jl,
const int ju, const int il, const int iu, const Function &function) {
Expand Down Expand Up @@ -770,7 +770,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ml, const i
});
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int nl, const int nu,
const int kl, const int ku, const int jl, const int ju, const int il,
const int iu, const Function &function) {
Expand All @@ -795,10 +795,10 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int nl, const i
});
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
const int kl, const int ku, const int jl,
const int ju, const int il, const int iu,
const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int kl, const int ku,
const int jl, const int ju, const int il, const int iu,
const Function &function) {
const int Nk = ku - kl + 1;
const int Nj = ju - jl + 1;
const int Ni = iu - il + 1;
Expand All @@ -815,9 +815,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m
});
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
const int jl, const int ju, const int il,
const int iu, const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int jl, const int ju,
const int il, const int iu, const Function &function) {
const int Nj = ju - jl + 1;
const int Ni = iu - il + 1;
const int NjNi = Nj * Ni;
Expand All @@ -828,22 +828,22 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m
});
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
const int il, const int iu,
const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR,
team_mbr_t team_member, const int il,
const int iu, const Function &function) {
Kokkos::parallel_for(Kokkos::TeamThreadRange(team_member, il, iu + 1), function);
}
// Inner parallel loop using TeamVectorRange
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTVR, team_mbr_t team_member,
const int il, const int iu,
const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternTVR,
team_mbr_t team_member, const int il,
const int iu, const Function &function) {
Kokkos::parallel_for(Kokkos::TeamVectorRange(team_member, il, iu + 1), function);
}

// Inner parallel loop using FOR SIMD
template <typename Function>
KOKKOS_INLINE_FUNCTION void
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int nl, const int nu,
const int kl, const int ku, const int jl, const int ju, const int il,
const int iu, const Function &function) {
Expand All @@ -859,10 +859,10 @@ par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int nl, con
}
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
const int kl, const int ku, const int jl,
const int ju, const int il, const int iu,
const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int kl, const int ku,
const int jl, const int ju, const int il, const int iu,
const Function &function) {
for (int k = kl; k <= ku; ++k) {
for (int j = jl; j <= ju; ++j) {
#pragma omp simd
Expand All @@ -873,9 +873,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te
}
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
const int jl, const int ju, const int il,
const int iu, const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void
par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int jl, const int ju,
const int il, const int iu, const Function &function) {
for (int j = jl; j <= ju; ++j) {
#pragma omp simd
for (int i = il; i <= iu; i++) {
Expand All @@ -884,9 +884,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te
}
}
template <typename Function>
KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
const int il, const int iu,
const Function &function) {
KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor,
team_mbr_t team_member, const int il,
const int iu, const Function &function) {
#pragma omp simd
for (int i = il; i <= iu; i++) {
function(i);
Expand Down
Loading