From 081b14c6379a4adbb7beacae0cd549f78f36681b Mon Sep 17 00:00:00 2001 From: jdolence Date: Tue, 24 Oct 2023 12:10:21 -0600 Subject: [PATCH 1/3] INLINE to FORCEINLINE on all par_for_inner functions --- src/kokkos_abstraction.hpp | 54 +++++++++++++++++++------------------- 1 file changed, 27 insertions(+), 27 deletions(-) diff --git a/src/kokkos_abstraction.hpp b/src/kokkos_abstraction.hpp index cabae1273b6f..f265fd22cc32 100644 --- a/src/kokkos_abstraction.hpp +++ b/src/kokkos_abstraction.hpp @@ -707,7 +707,7 @@ inline void par_for_outer(OuterLoopPatternTeams, const std::string &name, // Inner parallel loop using TeamThreadRange template -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, @@ -741,7 +741,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ll, const i }); } template -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) { @@ -770,7 +770,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ml, const i }); } template -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) { @@ -795,10 +795,10 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int nl, const i }); } template -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; @@ -815,9 +815,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m }); } template -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; @@ -828,22 +828,22 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m }); } template -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 -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 -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) { @@ -859,10 +859,10 @@ par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int nl, con } } template -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 @@ -873,9 +873,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te } } template -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++) { @@ -884,9 +884,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te } } template -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); From 5145bd8887fc01f9d49562a24e2fe14d5e08d847 Mon Sep 17 00:00:00 2001 From: jdolence Date: Tue, 24 Oct 2023 12:19:41 -0600 Subject: [PATCH 2/3] update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 49d929c46ab5..978bccbf19c0 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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 From 790cebb6488d4fe7b24bef31c7a4de4bbb1a5a3e Mon Sep 17 00:00:00 2001 From: jdolence Date: Tue, 24 Oct 2023 12:43:20 -0600 Subject: [PATCH 3/3] update copyright --- src/kokkos_abstraction.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/kokkos_abstraction.hpp b/src/kokkos_abstraction.hpp index f265fd22cc32..bfaf6a247a66 100644 --- a/src/kokkos_abstraction.hpp +++ b/src/kokkos_abstraction.hpp @@ -1,9 +1,9 @@ //======================================================================================== // Parthenon performance portable AMR framework -// Copyright(C) 2020-2022 The Parthenon collaboration +// Copyright(C) 2020-2023 The Parthenon collaboration // Licensed under the 3-clause BSD License, see LICENSE file for details //======================================================================================== -// (C) (or copyright) 2020-2022. Triad National Security, LLC. All rights reserved. +// (C) (or copyright) 2020-2023. Triad National Security, LLC. All rights reserved. // // This program was produced under U.S. Government contract 89233218CNA000001 // for Los Alamos National Laboratory (LANL), which is operated by Triad