From d59ecf0eb5d89abd974cadd87fc7ef5997271dde Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Tue, 11 Mar 2025 13:15:35 +0100 Subject: [PATCH 1/6] attempt at fixing Kokkos SYCL loop configuration (a proper benchmark would be needed). --- src/loop.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/loop.hpp b/src/loop.hpp index 8ad345de..73afe53e 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -51,6 +51,8 @@ typedef Kokkos::TeamPolicy<>::member_type member_type; constexpr LoopPattern defaultLoop = LoopPattern::RANGE; #elif defined(KOKKOS_ENABLE_HIP) constexpr LoopPattern defaultLoop = LoopPattern::RANGE; + #elif defined(KOKKOS_ENABLE_SYCL) + constexpr LoopPattern defaultLoop = LoopPattern::RANGE; #elif defined(KOKKOS_ENABLE_SERIAL) constexpr LoopPattern defaultLoop = LoopPattern::SIMDFOR; #else From ca9c8b0302c424f0f0b6439c4568979d9c21ef56 Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Tue, 11 Mar 2025 13:22:26 +0100 Subject: [PATCH 2/6] tabs --- src/loop.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/loop.hpp b/src/loop.hpp index 73afe53e..262c6ac1 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -51,7 +51,7 @@ typedef Kokkos::TeamPolicy<>::member_type member_type; constexpr LoopPattern defaultLoop = LoopPattern::RANGE; #elif defined(KOKKOS_ENABLE_HIP) constexpr LoopPattern defaultLoop = LoopPattern::RANGE; - #elif defined(KOKKOS_ENABLE_SYCL) + #elif defined(KOKKOS_ENABLE_SYCL) constexpr LoopPattern defaultLoop = LoopPattern::RANGE; #elif defined(KOKKOS_ENABLE_SERIAL) constexpr LoopPattern defaultLoop = LoopPattern::SIMDFOR; From dd702ab3e6d8a54b13e632773a00f78ab4b7247b Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Wed, 12 Mar 2025 14:27:16 -0400 Subject: [PATCH 3/6] add a static_assert so that the error observed on syscl doesn't go unnoticed when a new architecture is tried --- src/loop.hpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/src/loop.hpp b/src/loop.hpp index 262c6ac1..b6a5a9a1 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -310,6 +310,12 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { + // Check that Idefix Arrays can be assigned from SIMD loop + static_assert( + Kokkos::SpaceAccessibility::accessible, + "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." + ); for (auto n = NB; n < NE; n++) for (auto k = KB; k < KE; k++) for (auto j = JB; j < JE; j++) From 9d7c589b43f935ac10768cf49072cea597fb2e4d Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Wed, 12 Mar 2025 14:32:08 -0400 Subject: [PATCH 4/6] do this for all loops --- src/loop.hpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/src/loop.hpp b/src/loop.hpp index b6a5a9a1..06238011 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -128,6 +128,12 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { + // Check that Idefix Arrays can be assigned from SIMD loop + static_assert( + Kokkos::SpaceAccessibility::accessible, + "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." + ); for (auto j = JB; j < JE; j++) #pragma omp simd for (auto i = IB; i < IE; i++) @@ -212,6 +218,12 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { + // Check that Idefix Arrays can be assigned from SIMD loop + static_assert( + Kokkos::SpaceAccessibility::accessible, + "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." + ); for (auto k = KB; k < KE; k++) for (auto j = JB; j < JE; j++) #pragma omp simd From 217a812a9a2bf82b676004ecee5b21eaf043796b Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Wed, 12 Mar 2025 14:46:13 -0400 Subject: [PATCH 5/6] catch misbehaved loops earlier (fix cuda!), and be more explicit. --- src/loop.hpp | 29 +++++++++++++---------------- 1 file changed, 13 insertions(+), 16 deletions(-) diff --git a/src/loop.hpp b/src/loop.hpp index 06238011..20933ec9 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -34,6 +34,12 @@ typedef Kokkos::TeamPolicy<>::member_type member_type; // Check if the user requested a specific loop unrolling strategy #if defined(LOOP_PATTERN_SIMD) + // Check that Idefix Arrays can be assigned from SIMD loop + static_assert( + Kokkos::SpaceAccessibility::accessible, + "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." + ); constexpr LoopPattern defaultLoop = LoopPattern::SIMDFOR; #elif defined(LOOP_PATTERN_1DRANGE) constexpr LoopPattern defaultLoop = LoopPattern::RANGE; @@ -54,6 +60,13 @@ typedef Kokkos::TeamPolicy<>::member_type member_type; #elif defined(KOKKOS_ENABLE_SYCL) constexpr LoopPattern defaultLoop = LoopPattern::RANGE; #elif defined(KOKKOS_ENABLE_SERIAL) + // Check that Idefix Arrays can be assigned from SIMD loop + static_assert( + Kokkos::SpaceAccessibility::accessible, + "Idefix Arrays cannot be accessed from Host, but Device is unknown/untested. " + "Ask the developers to add support." + ); constexpr LoopPattern defaultLoop = LoopPattern::SIMDFOR; #else #warning "Unknown target architeture: default to MDrange" @@ -129,11 +142,6 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { // Check that Idefix Arrays can be assigned from SIMD loop - static_assert( - Kokkos::SpaceAccessibility::accessible, - "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." - ); for (auto j = JB; j < JE; j++) #pragma omp simd for (auto i = IB; i < IE; i++) @@ -219,11 +227,6 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { // Check that Idefix Arrays can be assigned from SIMD loop - static_assert( - Kokkos::SpaceAccessibility::accessible, - "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." - ); for (auto k = KB; k < KE; k++) for (auto j = JB; j < JE; j++) #pragma omp simd @@ -322,12 +325,6 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { - // Check that Idefix Arrays can be assigned from SIMD loop - static_assert( - Kokkos::SpaceAccessibility::accessible, - "Idefix Arrays cannot be accessed from SIMD loop. You should try another loop pattern." - ); for (auto n = NB; n < NE; n++) for (auto k = KB; k < KE; k++) for (auto j = JB; j < JE; j++) From e58f72452254ef5650fd3529452af56eae1337f1 Mon Sep 17 00:00:00 2001 From: Geoffroy Lesur Date: Wed, 12 Mar 2025 16:32:48 -0400 Subject: [PATCH 6/6] cleaning up --- src/loop.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/loop.hpp b/src/loop.hpp index 20933ec9..42192fa8 100644 --- a/src/loop.hpp +++ b/src/loop.hpp @@ -141,7 +141,6 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { - // Check that Idefix Arrays can be assigned from SIMD loop for (auto j = JB; j < JE; j++) #pragma omp simd for (auto i = IB; i < IE; i++) @@ -226,7 +225,6 @@ inline void idefix_for(const std::string & NAME, // SIMD FOR loops } else if constexpr(defaultLoop == LoopPattern::SIMDFOR) { - // Check that Idefix Arrays can be assigned from SIMD loop for (auto k = KB; k < KE; k++) for (auto j = JB; j < JE; j++) #pragma omp simd