From 0e97668f227781259239bc3bdb8c9f000e786160 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Fri, 12 Sep 2025 18:26:48 -0700 Subject: [PATCH 01/13] first commit. WIP. --- include/RAJA/RAJA.hpp | 2 +- include/RAJA/util/SubView.hpp | 75 ++++++++++++++++++++++ test/unit/view-layout/test-indexlayout.cpp | 14 ++++ 3 files changed, 90 insertions(+), 1 deletion(-) create mode 100644 include/RAJA/util/SubView.hpp diff --git a/include/RAJA/RAJA.hpp b/include/RAJA/RAJA.hpp index abc9f30738..8cb3abb40b 100644 --- a/include/RAJA/RAJA.hpp +++ b/include/RAJA/RAJA.hpp @@ -118,7 +118,7 @@ #include "RAJA/util/StaticLayout.hpp" #include "RAJA/util/IndexLayout.hpp" #include "RAJA/util/View.hpp" - +#include "RAJA/util/SubView.hpp" // // View for sequences of objects diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp new file mode 100644 index 0000000000..2436d5678c --- /dev/null +++ b/include/RAJA/util/SubView.hpp @@ -0,0 +1,75 @@ +#include "camp/number.hpp" +#include "camp/tuple.hpp" +#include "camp/array.hpp" +#include +#include +#include +#include +#include +#include + +// Slice descriptors +struct Range { long start, end; }; +struct Fixed { long idx; }; + +// Helper to count non-fixed dimensions at compile time +template +constexpr size_t count_ranges() { + return (0 + ... + (std::is_same_v ? 1 : 0)); +} + +// Helper to get the i-th element from a tuple +template +decltype(auto) get_elem(Tuple&& tup) { + return camp::get(std::forward(tup)); +} + +template +auto array_to_tuple_impl(const camp::array& arr, std::index_sequence) { + return camp::make_tuple(arr[Is]...); +} + +template +auto array_to_tuple(const camp::array& arr) { + return array_to_tuple_impl(arr, std::make_index_sequence{}); +} + +// The generic subview class +template +class SubView { + ViewType view_; + camp::tuple slices_; + + // Helper to map subview indices to parent indices + template + auto map_indices(IndexType* idxs, camp::idx_seq) const { + camp::array parent_indices{}; + size_t range_idx = 0; + ( + ( + parent_indices[Is] = [&] { + const auto& s = get_elem(slices_); + if constexpr (std::is_same_v, Range>) { + return s.start + idxs[range_idx++]; + } else { + return s.idx; + } + }() + ), ... + ); + return parent_indices; + } + +public: + SubView(ViewType view, Slices... slices) + : view_(view), slices_(slices...) {} + + template + auto operator()(Idxs... idxs) const { + constexpr size_t nidx = count_ranges(); + static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); + camp::array arr{static_cast(idxs)...}; + auto parent_indices = map_indices(arr.data(), camp::make_idx_seq_t()); + return camp::apply(view_, array_to_tuple(parent_indices)); + } +}; diff --git a/test/unit/view-layout/test-indexlayout.cpp b/test/unit/view-layout/test-indexlayout.cpp index c15304d3ed..c286f52cdc 100644 --- a/test/unit/view-layout/test-indexlayout.cpp +++ b/test/unit/view-layout/test-indexlayout.cpp @@ -343,3 +343,17 @@ TEST(IndexLayout, MultiView1DLayout) } +TEST(IndexLayout, SubView) +{ + +long a[] = {1,2,3,4,5}; + +RAJA::View> view(&a[0], RAJA::Layout<1>(5)); + +auto sv = SubView(view, Range{1,3}); + +long idx = 2; +std::cout << sv(idx) << std::endl; + +} + From 7ebc87a6b853003dc86b378122fc0ffba9c14ed6 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Thu, 18 Sep 2025 13:06:24 -0700 Subject: [PATCH 02/13] Added tests --- test/unit/view-layout/CMakeLists.txt | 4 ++++ test/unit/view-layout/test-indexlayout.cpp | 14 -------------- 2 files changed, 4 insertions(+), 14 deletions(-) diff --git a/test/unit/view-layout/CMakeLists.txt b/test/unit/view-layout/CMakeLists.txt index 90141fbf82..635145ce8b 100644 --- a/test/unit/view-layout/CMakeLists.txt +++ b/test/unit/view-layout/CMakeLists.txt @@ -28,3 +28,7 @@ raja_add_test( raja_add_test( NAME test-indexlayout SOURCES test-indexlayout.cpp) + +raja_add_test( + NAME test-subview + SOURCES test-subview.cpp) diff --git a/test/unit/view-layout/test-indexlayout.cpp b/test/unit/view-layout/test-indexlayout.cpp index c286f52cdc..c15304d3ed 100644 --- a/test/unit/view-layout/test-indexlayout.cpp +++ b/test/unit/view-layout/test-indexlayout.cpp @@ -343,17 +343,3 @@ TEST(IndexLayout, MultiView1DLayout) } -TEST(IndexLayout, SubView) -{ - -long a[] = {1,2,3,4,5}; - -RAJA::View> view(&a[0], RAJA::Layout<1>(5)); - -auto sv = SubView(view, Range{1,3}); - -long idx = 2; -std::cout << sv(idx) << std::endl; - -} - From 67d9653caa85d59b7817b293d27d80562920e5ac Mon Sep 17 00:00:00 2001 From: gberg617 Date: Thu, 18 Sep 2025 16:50:12 -0700 Subject: [PATCH 03/13] small fixes --- include/RAJA/util/SubView.hpp | 87 ++++++++++++++++++++++------------- 1 file changed, 54 insertions(+), 33 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 2436d5678c..99e74af697 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -1,75 +1,96 @@ +/*! + ****************************************************************************** + * + * \file + * + * \brief RAJA header file defining the SubView class + * + ****************************************************************************** + */ + +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC +// and RAJA project contributors. See the RAJA/LICENSE file for details. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#ifndef RAJA_SUBVIEW_HPP +#define RAJA_SUBVIEW_HPP + +#include "RAJA/util/types.hpp" #include "camp/number.hpp" #include "camp/tuple.hpp" #include "camp/array.hpp" -#include -#include -#include -#include -#include -#include + +namespace RAJA +{ // Slice descriptors -struct Range { long start, end; }; -struct Fixed { long idx; }; +struct RangeSlice { long start, end; }; +struct FixedSlice { long idx; }; +struct NoSlice { }; // Helper to count non-fixed dimensions at compile time template -constexpr size_t count_ranges() { - return (0 + ... + (std::is_same_v ? 1 : 0)); -} - -// Helper to get the i-th element from a tuple -template -decltype(auto) get_elem(Tuple&& tup) { - return camp::get(std::forward(tup)); +RAJA_INLINE RAJA_HOST_DEVICE constexpr size_t count_nonfixed_dims() { + return (0 + ... + ((std::is_same_v || + std::is_same_v) ? 1 : 0)); } -template -auto array_to_tuple_impl(const camp::array& arr, std::index_sequence) { +template +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple_impl(const camp::array& arr, camp::idx_seq) { return camp::make_tuple(arr[Is]...); } template -auto array_to_tuple(const camp::array& arr) { - return array_to_tuple_impl(arr, std::make_index_sequence{}); +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple(const camp::array& arr) { + return array_to_tuple_impl(arr, camp::make_idx_seq_t{}); } -// The generic subview class template class SubView { ViewType view_; camp::tuple slices_; - // Helper to map subview indices to parent indices template - auto map_indices(IndexType* idxs, camp::idx_seq) const { - camp::array parent_indices{}; - size_t range_idx = 0; + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_indices(IndexType* idxs, camp::idx_seq) const { + camp::array parent_indices{}; + RAJA::Index_type idx = 0; + + // For each slice, map subview index to parent index ( ( parent_indices[Is] = [&] { - const auto& s = get_elem(slices_); - if constexpr (std::is_same_v, Range>) { - return s.start + idxs[range_idx++]; - } else { + const auto& s = camp::get(slices_); + if constexpr (std::is_same_v, RangeSlice>) { + return s.start + idxs[idx++]; + } else if constexpr (std::is_same_v, FixedSlice>) { return s.idx; + } else { + return idxs[idx++]; } }() ), ... ); + return parent_indices; } public: - SubView(ViewType view, Slices... slices) + RAJA_INLINE RAJA_HOST_DEVICE SubView(ViewType view, Slices... slices) : view_(view), slices_(slices...) {} template - auto operator()(Idxs... idxs) const { - constexpr size_t nidx = count_ranges(); + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto operator()(Idxs... idxs) const { + constexpr size_t nidx = count_nonfixed_dims(); static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); - camp::array arr{static_cast(idxs)...}; + camp::array arr{idxs...}; auto parent_indices = map_indices(arr.data(), camp::make_idx_seq_t()); return camp::apply(view_, array_to_tuple(parent_indices)); } }; + +} // namespace RAJA + +#endif \ No newline at end of file From b942833cce8daddd44ba1fc61db488a559058d27 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Fri, 19 Sep 2025 13:59:13 -0700 Subject: [PATCH 04/13] updates to test --- test/unit/view-layout/test-subview.cpp | 119 +++++++++++++++++++++++++ 1 file changed, 119 insertions(+) create mode 100644 test/unit/view-layout/test-subview.cpp diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp new file mode 100644 index 0000000000..15eb22c8ad --- /dev/null +++ b/test/unit/view-layout/test-subview.cpp @@ -0,0 +1,119 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC +// and RAJA project contributors. See the RAJA/LICENSE file for details. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#include +#include "RAJA/policy/PolicyBase.hpp" +#include "RAJA/util/types.hpp" +#include "RAJA_test-base.hpp" +#include "RAJA_unit-test-forone.hpp" + +using namespace RAJA; + +TEST(SubView, RangeSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[1:3] + auto sv = SubView(view, RangeSlice{1,3}); + + EXPECT_EQ(sv(0), 2); + EXPECT_EQ(sv(1), 3); + EXPECT_EQ(sv(2), 4); + +} + +TEST(SubView, RangeSubView2D) +{ + + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1:2,1:2] + auto sv = SubView(view, RangeSlice{1,2}, RangeSlice{1,2}); + + EXPECT_EQ(sv(0,0), 5); + EXPECT_EQ(sv(0,1), 6); + EXPECT_EQ(sv(1,0), 8); + EXPECT_EQ(sv(1,1), 9); + +} + +TEST(SubView, RangeFixedSubView2D) +{ + + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1:2,1] + auto sv = SubView(view, RangeSlice{1,2}, FixedSlice{1}); + + EXPECT_EQ(sv(0), 5); + EXPECT_EQ(sv(1), 8); + +} + +TEST(SubView, FixedFirstDimSubView2D) +{ + + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1,:] + auto sv = SubView(view, FixedSlice{1}, NoSlice{}); + + EXPECT_EQ(sv(0), 4); + EXPECT_EQ(sv(1), 5); + EXPECT_EQ(sv(2), 6); + +} + +TEST(SubView, RangeFirstDimSubView2D) +{ + + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1:2,:] + auto sv = SubView(view, RangeSlice{1,2}, NoSlice{}); + + EXPECT_EQ(sv(0,0), 4); + EXPECT_EQ(sv(0,1), 5); + EXPECT_EQ(sv(0,2), 6); + + EXPECT_EQ(sv(1,0), 7); + EXPECT_EQ(sv(1,1), 8); + EXPECT_EQ(sv(1,2), 9); + +} + +void test_subviewGPU() { +#if defined(RAJA_ENABLE_HIP) + forone([=] __host__ __device__ () { + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1:2,:] + auto sv = SubView(view, RangeSlice{1,2}, NoSlice{}); + + //printf("sv(0,0): %ld\n", sv(0,0)); + + }); +#endif +} + +TEST(SubView, RangeFirstDimSubView2DGPU) +{ + test_subviewGPU(); +} \ No newline at end of file From 5ba920d6543ab88b21577a5b29ef83f673e69e65 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Mon, 29 Sep 2025 18:03:43 -0700 Subject: [PATCH 05/13] various changes/fixes to address comments --- include/RAJA/util/SubView.hpp | 82 +++++++++++++++++--------- test/unit/view-layout/test-subview.cpp | 41 ++++++------- 2 files changed, 74 insertions(+), 49 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 99e74af697..97a5fc88dc 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -27,15 +27,42 @@ namespace RAJA { // Slice descriptors -struct RangeSlice { long start, end; }; -struct FixedSlice { long idx; }; -struct NoSlice { }; + +template +struct RangeSlice { + IndexType start, end; + + static constexpr bool reduces_dimension = false; + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { + return start + idx; + } +}; + +template +struct FixedSlice { + IndexType idx; + + static constexpr bool reduces_dimension = true; + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType&) const { + return idx; + } +}; + +struct NoSlice { + static constexpr bool reduces_dimension = false; + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { + return idx; + } +}; // Helper to count non-fixed dimensions at compile time template RAJA_INLINE RAJA_HOST_DEVICE constexpr size_t count_nonfixed_dims() { - return (0 + ... + ((std::is_same_v || - std::is_same_v) ? 1 : 0)); + return (!Slices::reduces_dimension + ...); } template @@ -48,45 +75,42 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple(const camp::array{}); } -template +template class SubView { ViewType view_; camp::tuple slices_; + std::array map_; - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_indices(IndexType* idxs, camp::idx_seq) const { - camp::array parent_indices{}; - RAJA::Index_type idx = 0; + RAJA_INLINE RAJA_HOST_DEVICE constexpr void make_subview_index_map() { + size_t sub_idx = 0; + size_t i = 0; + ((map_[i++] = (Slices::reduces_dimension ? -1 : sub_idx++)), ...); + } + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_subview_idx_to_parent(IndexType* idxs) const { + return camp::get(slices_).map_index(idxs[map_[I]]); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_indices(IndexType* idxs, camp::idx_seq) const { // For each slice, map subview index to parent index - ( - ( - parent_indices[Is] = [&] { - const auto& s = camp::get(slices_); - if constexpr (std::is_same_v, RangeSlice>) { - return s.start + idxs[idx++]; - } else if constexpr (std::is_same_v, FixedSlice>) { - return s.idx; - } else { - return idxs[idx++]; - } - }() - ), ... - ); - - return parent_indices; + return camp::array{(map_subview_idx_to_parent(idxs))...}; } public: - RAJA_INLINE RAJA_HOST_DEVICE SubView(ViewType view, Slices... slices) - : view_(view), slices_(slices...) {} + + RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) + : view_(view), slices_(slices...) { make_subview_index_map(); } template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto operator()(Idxs... idxs) const { + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { constexpr size_t nidx = count_nonfixed_dims(); static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); + camp::array arr{idxs...}; auto parent_indices = map_indices(arr.data(), camp::make_idx_seq_t()); + return camp::apply(view_, array_to_tuple(parent_indices)); } }; diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index 15eb22c8ad..48897e1a0c 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -7,6 +7,7 @@ #include #include "RAJA/policy/PolicyBase.hpp" +#include "RAJA/util/SubView.hpp" #include "RAJA/util/types.hpp" #include "RAJA_test-base.hpp" #include "RAJA_unit-test-forone.hpp" @@ -21,7 +22,7 @@ TEST(SubView, RangeSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[1:3] - auto sv = SubView(view, RangeSlice{1,3}); + auto sv = SubView(view, RangeSlice<>{1,3}); EXPECT_EQ(sv(0), 2); EXPECT_EQ(sv(1), 3); @@ -37,7 +38,7 @@ TEST(SubView, RangeSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1:2] - auto sv = SubView(view, RangeSlice{1,2}, RangeSlice{1,2}); + auto sv = SubView(view, RangeSlice<>{1,2}, RangeSlice<>{1,2}); EXPECT_EQ(sv(0,0), 5); EXPECT_EQ(sv(0,1), 6); @@ -54,7 +55,7 @@ TEST(SubView, RangeFixedSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1] - auto sv = SubView(view, RangeSlice{1,2}, FixedSlice{1}); + auto sv = SubView(view, RangeSlice<>{1,2}, FixedSlice<>{1}); EXPECT_EQ(sv(0), 5); EXPECT_EQ(sv(1), 8); @@ -69,7 +70,7 @@ TEST(SubView, FixedFirstDimSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1,:] - auto sv = SubView(view, FixedSlice{1}, NoSlice{}); + auto sv = SubView(view, FixedSlice<>{1}, NoSlice{}); EXPECT_EQ(sv(0), 4); EXPECT_EQ(sv(1), 5); @@ -85,7 +86,7 @@ TEST(SubView, RangeFirstDimSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,:] - auto sv = SubView(view, RangeSlice{1,2}, NoSlice{}); + auto sv = SubView(view, RangeSlice<>{1,2}, NoSlice{}); EXPECT_EQ(sv(0,0), 4); EXPECT_EQ(sv(0,1), 5); @@ -97,23 +98,23 @@ TEST(SubView, RangeFirstDimSubView2D) } -void test_subviewGPU() { -#if defined(RAJA_ENABLE_HIP) - forone([=] __host__ __device__ () { - Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; +// void test_subviewGPU() { +// #if defined(RAJA_ENABLE_HIP) +// forone([=] __host__ __device__ () { +// Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; - View> view(&a[0][0], Layout<2>(3,3)); +// View> view(&a[0][0], Layout<2>(3,3)); - // sv = View[1:2,:] - auto sv = SubView(view, RangeSlice{1,2}, NoSlice{}); +// // sv = View[1:2,:] +// auto sv = SubView(view, RangeSlice<>{1,2}, NoSlice{}); - //printf("sv(0,0): %ld\n", sv(0,0)); +// //printf("sv(0,0): %ld\n", sv(0,0)); - }); -#endif -} +// }); +// #endif +// } -TEST(SubView, RangeFirstDimSubView2DGPU) -{ - test_subviewGPU(); -} \ No newline at end of file +// TEST(SubView, RangeFirstDimSubView2DGPU) +// { +// test_subviewGPU(); +// } \ No newline at end of file From ffc1d4cf3bca67fa0691de1535112d4f669fffe2 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Fri, 10 Oct 2025 13:52:00 -0700 Subject: [PATCH 06/13] implemented for_each_index --- include/RAJA/util/SubView.hpp | 30 ++++++++++++++------------ include/RAJA/util/for_each.hpp | 27 ++++++++++++++++++++++- test/unit/view-layout/test-subview.cpp | 2 ++ 3 files changed, 44 insertions(+), 15 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 97a5fc88dc..6a8fb24dea 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -18,6 +18,7 @@ #ifndef RAJA_SUBVIEW_HPP #define RAJA_SUBVIEW_HPP +#include "RAJA/util/for_each.hpp" #include "RAJA/util/types.hpp" #include "camp/number.hpp" #include "camp/tuple.hpp" @@ -28,6 +29,11 @@ namespace RAJA // Slice descriptors +// template +// RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { +// return start + idx; +// } + template struct RangeSlice { IndexType start, end; @@ -78,7 +84,7 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple(const camp::array class SubView { ViewType view_; - camp::tuple slices_; + camp::tuple slices_; std::array map_; RAJA_INLINE RAJA_HOST_DEVICE constexpr void make_subview_index_map() { @@ -87,17 +93,6 @@ class SubView { ((map_[i++] = (Slices::reduces_dimension ? -1 : sub_idx++)), ...); } - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_subview_idx_to_parent(IndexType* idxs) const { - return camp::get(slices_).map_index(idxs[map_[I]]); - } - - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto map_indices(IndexType* idxs, camp::idx_seq) const { - // For each slice, map subview index to parent index - return camp::array{(map_subview_idx_to_parent(idxs))...}; - } - public: RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) @@ -108,8 +103,15 @@ class SubView { constexpr size_t nidx = count_nonfixed_dims(); static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); - camp::array arr{idxs...}; - auto parent_indices = map_indices(arr.data(), camp::make_idx_seq_t()); + camp::array arr{idxs...}; + camp::array parent_indices; + + for_each_index( + [&](auto I) { + constexpr auto Ival = I.value; + auto slice = camp::get(slices_); + parent_indices[Ival] = slice.map_index(arr[map_[Ival]]); + }); return camp::apply(view_, array_to_tuple(parent_indices)); } diff --git a/include/RAJA/util/for_each.hpp b/include/RAJA/util/for_each.hpp index 708d3115e8..3ab2020342 100644 --- a/include/RAJA/util/for_each.hpp +++ b/include/RAJA/util/for_each.hpp @@ -24,6 +24,9 @@ #include #include "camp/list.hpp" +#include "camp/concepts.hpp" +#include "camp/number.hpp" +#include "camp/tuple.hpp" #include "RAJA/pattern/detail/algorithm.hpp" @@ -36,6 +39,20 @@ namespace RAJA namespace detail { +// compile time expansion applying func to each of the indices +RAJA_SUPPRESS_HD_WARN +template +RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_index(camp::idx_seq, + UnaryFunc func) +{ + // braced init lists are evaluated in order + // create integral_constant type to allow UnaryFunc arguments to be used in compile-time context + int seq_unused_array[] = {0, (func(std::integral_constant{}), 0)...}; + RAJA_UNUSED_VAR(seq_unused_array); + + return func; +} + // runtime loop applying func to each element in the range in order RAJA_SUPPRESS_HD_WARN template @@ -81,6 +98,14 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, } // namespace detail +RAJA_SUPPRESS_HD_WARN +template +RAJA_HOST_DEVICE RAJA_INLINE +UnaryFunc for_each_index(UnaryFunc func) +{ + return detail::for_each_index(camp::make_idx_seq_t(), std::move(func)); +} + /*! \brief Apply func to all the elements in the given range in order using a sequential for loop in O(N) operations and O(1) extra memory @@ -89,7 +114,7 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, RAJA_SUPPRESS_HD_WARN template RAJA_HOST_DEVICE RAJA_INLINE - concepts::enable_if_t> + camp::concepts::enable_if_t> for_each(Container&& c, UnaryFunc func) { using std::begin; diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index 48897e1a0c..91f42ed424 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -98,6 +98,8 @@ TEST(SubView, RangeFirstDimSubView2D) } +// To Do tests: subview of subview, "sliding winow" modify slice entries, GPU tests. + // void test_subviewGPU() { // #if defined(RAJA_ENABLE_HIP) // forone([=] __host__ __device__ () { From e899696b7c0cda5028be54bbc1debac522b6e73a Mon Sep 17 00:00:00 2001 From: gberg617 Date: Fri, 10 Oct 2025 18:05:03 -0700 Subject: [PATCH 07/13] various more simplifications --- include/RAJA/util/SubView.hpp | 42 +++++++++++++++++++--------------- include/RAJA/util/for_each.hpp | 29 +++++++++++++++++++++++ 2 files changed, 52 insertions(+), 19 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 6a8fb24dea..a3844f7f10 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -56,21 +56,15 @@ struct FixedSlice { } }; +template struct NoSlice { static constexpr bool reduces_dimension = false; - template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { return idx; } }; -// Helper to count non-fixed dimensions at compile time -template -RAJA_INLINE RAJA_HOST_DEVICE constexpr size_t count_nonfixed_dims() { - return (!Slices::reduces_dimension + ...); -} - template RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple_impl(const camp::array& arr, camp::idx_seq) { return camp::make_tuple(arr[Is]...); @@ -81,42 +75,52 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple(const camp::array{}); } -template -class SubView { +template +class SubView; + +template +class SubView, IndexType> { ViewType view_; camp::tuple slices_; std::array map_; - RAJA_INLINE RAJA_HOST_DEVICE constexpr void make_subview_index_map() { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { size_t sub_idx = 0; - size_t i = 0; - ((map_[i++] = (Slices::reduces_dimension ? -1 : sub_idx++)), ...); + std::array map; + + for_each_tuple_index( slices_, + [&](auto slice, auto index) { + map[index] = decltype(slice)::reduces_dimension ? -1 : sub_idx++; + }); + + return map; } public: RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) - : view_(view), slices_(slices...) { make_subview_index_map(); } + : view_(view), slices_(slices...), map_(make_subview_index_map()) { } template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { - constexpr size_t nidx = count_nonfixed_dims(); + constexpr size_t nidx = ((Slices::reduces_dimension == false ? 1 : 0) + ...); static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); camp::array arr{idxs...}; camp::array parent_indices; - for_each_index( - [&](auto I) { - constexpr auto Ival = I.value; - auto slice = camp::get(slices_); - parent_indices[Ival] = slice.map_index(arr[map_[Ival]]); + for_each_tuple_index( slices_, + [&](auto slice, auto index) { + parent_indices[index] = slice.map_index(arr[map_[index]]); }); return camp::apply(view_, array_to_tuple(parent_indices)); } }; +template +SubView(ViewType, Slices...) -> SubView>; + } // namespace RAJA #endif \ No newline at end of file diff --git a/include/RAJA/util/for_each.hpp b/include/RAJA/util/for_each.hpp index 3ab2020342..1e8813aa1e 100644 --- a/include/RAJA/util/for_each.hpp +++ b/include/RAJA/util/for_each.hpp @@ -96,6 +96,21 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, return func; } +// compile time expansion applying func to a each type in the tuple in order +RAJA_SUPPRESS_HD_WARN +template +RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, + BinaryFunc func, + camp::idx_seq) +{ + using camp::get; + // braced init lists are evaluated in order + int seq_unused_array[] = {0, (func(get(std::forward(t)), Is), 0)...}; + RAJA_UNUSED_VAR(seq_unused_array); + + return func; +} + } // namespace detail RAJA_SUPPRESS_HD_WARN @@ -148,6 +163,20 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, UnaryFunc func) camp::make_idx_seq_t>::value> {}); } +/*! + \brief Apply func to each object in the given tuple or tuple like type as well as + the index of the tuple element in order using a compile-time expansion in O(N) + operations and O(1) extra memory +*/ +RAJA_SUPPRESS_HD_WARN +template +RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, BinaryFunc func) +{ + return detail::for_each_tuple_index( + std::forward(t), std::move(func), + camp::make_idx_seq_t>::value> {}); +} + } // namespace RAJA #endif From d8a9081a113a2aea32c5a81e731d7e9b92182f64 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Fri, 10 Oct 2025 19:00:01 -0700 Subject: [PATCH 08/13] added set/get slice methods and test --- include/RAJA/util/SubView.hpp | 20 +++++++++++++++++++- test/unit/view-layout/test-subview.cpp | 10 ++++++++++ 2 files changed, 29 insertions(+), 1 deletion(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index a3844f7f10..05878498da 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -81,7 +81,7 @@ class SubView; template class SubView, IndexType> { ViewType view_; - camp::tuple slices_; + camp::tuple slices_; std::array map_; RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { @@ -101,6 +101,24 @@ class SubView, IndexType> { RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) : view_(view), slices_(slices...), map_(make_subview_index_map()) { } + RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slices(Slices... slices) { + slices_ = camp::tuple(slices...); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slice(Slice slice) { + camp::get(slices_) = slice; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slices() { + return slices_; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slice() { + return camp::get(slices_); + } + template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { constexpr size_t nidx = ((Slices::reduces_dimension == false ? 1 : 0) + ...); diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index 91f42ed424..34851efe2f 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -96,6 +96,16 @@ TEST(SubView, RangeFirstDimSubView2D) EXPECT_EQ(sv(1,1), 8); EXPECT_EQ(sv(1,2), 9); + sv.set_slice<0>(RangeSlice<>{0,1}); + + EXPECT_EQ(sv(0,0), 1); + EXPECT_EQ(sv(0,1), 2); + EXPECT_EQ(sv(0,2), 3); + + EXPECT_EQ(sv(1,0), 4); + EXPECT_EQ(sv(1,1), 5); + EXPECT_EQ(sv(1,2), 6); + } // To Do tests: subview of subview, "sliding winow" modify slice entries, GPU tests. From 6aa5662fd03ec565309ce17e796b61ff95298df2 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Tue, 14 Oct 2025 12:25:17 -0700 Subject: [PATCH 09/13] made simplifications to map and parent_indices --- include/RAJA/util/SubView.hpp | 66 +++++++++++++++------------------- include/RAJA/util/for_each.hpp | 41 ++++++++++----------- include/RAJA/util/macros.hpp | 2 +- 3 files changed, 50 insertions(+), 59 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 05878498da..545c2b5e80 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -19,8 +19,8 @@ #define RAJA_SUBVIEW_HPP #include "RAJA/util/for_each.hpp" +#include "RAJA/util/macros.hpp" #include "RAJA/util/types.hpp" -#include "camp/number.hpp" #include "camp/tuple.hpp" #include "camp/array.hpp" @@ -29,18 +29,13 @@ namespace RAJA // Slice descriptors -// template -// RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { -// return start + idx; -// } - template struct RangeSlice { IndexType start, end; static constexpr bool reduces_dimension = false; - RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { return start + idx; } }; @@ -51,7 +46,7 @@ struct FixedSlice { static constexpr bool reduces_dimension = true; - RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType&) const { + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType&) const { return idx; } }; @@ -60,19 +55,18 @@ template struct NoSlice { static constexpr bool reduces_dimension = false; - RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(IndexType& idx) const { + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { return idx; } }; -template -RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple_impl(const camp::array& arr, camp::idx_seq) { - return camp::make_tuple(arr[Is]...); -} - -template -RAJA_INLINE RAJA_HOST_DEVICE constexpr auto array_to_tuple(const camp::array& arr) { - return array_to_tuple_impl(arr, camp::make_idx_seq_t{}); +template +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { + IndexType sub_idx = 0; + IndexType i = 0; + camp::array map = {}; + ((map[i++] = (Slices::reduces_dimension ? -1 : sub_idx++)), ...); + return map; } template @@ -82,24 +76,14 @@ template class SubView, IndexType> { ViewType view_; camp::tuple slices_; - std::array map_; - - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { - size_t sub_idx = 0; - std::array map; - - for_each_tuple_index( slices_, - [&](auto slice, auto index) { - map[index] = decltype(slice)::reduces_dimension ? -1 : sub_idx++; - }); - - return map; - } + static inline constexpr size_t num_slices_ = sizeof...(Slices); + static inline constexpr IndexType num_sub_indices_ = ((Slices::reduces_dimension == false ? 1 : 0) + ...); + static inline constexpr camp::array map_ = make_subview_index_map(); public: RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) - : view_(view), slices_(slices...), map_(make_subview_index_map()) { } + : view_(view), slices_(slices...) { } RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slices(Slices... slices) { slices_ = camp::tuple(slices...); @@ -121,18 +105,24 @@ class SubView, IndexType> { template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { - constexpr size_t nidx = ((Slices::reduces_dimension == false ? 1 : 0) + ...); - static_assert(sizeof...(idxs) == nidx, "Wrong number of indices for subview"); + static_assert(sizeof...(idxs) == num_sub_indices_, "Wrong number of indices for subview"); - camp::array arr{idxs...}; - camp::array parent_indices; + camp::array arr{idxs...}; + camp::array parent_indices; for_each_tuple_index( slices_, - [&](auto slice, auto index) { - parent_indices[index] = slice.map_index(arr[map_[index]]); + [&](auto slice, auto index) { + if (map_[index] >= 0) { + parent_indices[index] = slice.map_index(arr[map_[index]]); + } else { + // map_index will not need index values for dimension-reducing slices + // so we pass a "dummy" value. + constexpr IndexType dummy_value = -1; + parent_indices[index] = slice.map_index(dummy_value); + } }); - return camp::apply(view_, array_to_tuple(parent_indices)); + return camp::apply(view_, parent_indices); } }; diff --git a/include/RAJA/util/for_each.hpp b/include/RAJA/util/for_each.hpp index 1e8813aa1e..dc05c83715 100644 --- a/include/RAJA/util/for_each.hpp +++ b/include/RAJA/util/for_each.hpp @@ -42,8 +42,8 @@ namespace detail // compile time expansion applying func to each of the indices RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_index(camp::idx_seq, - UnaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_index(camp::idx_seq, + UnaryFunc func) { // braced init lists are evaluated in order // create integral_constant type to allow UnaryFunc arguments to be used in compile-time context @@ -56,9 +56,9 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_index(camp::idx_seq, // runtime loop applying func to each element in the range in order RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each(Iter begin, - Iter end, - UnaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each(Iter begin, + Iter end, + UnaryFunc func) { for (; begin != end; ++begin) { @@ -71,8 +71,8 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each(Iter begin, // compile time expansion applying func to a each type in the list in order RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const&, - UnaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const&, + UnaryFunc func) { // braced init lists are evaluated in order int seq_unused_array[] = {0, (func(Ts {}), 0)...}; @@ -84,9 +84,9 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const&, // compile time expansion applying func to a each type in the tuple in order RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, - UnaryFunc func, - camp::idx_seq) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, + UnaryFunc func, + camp::idx_seq) { using camp::get; // braced init lists are evaluated in order @@ -99,13 +99,14 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, // compile time expansion applying func to a each type in the tuple in order RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, - BinaryFunc func, - camp::idx_seq) +constexpr RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, + BinaryFunc func, + camp::idx_seq) { using camp::get; // braced init lists are evaluated in order - int seq_unused_array[] = {0, (func(get(std::forward(t)), Is), 0)...}; + int seq_unused_array[] = {0, (func(get(std::forward(t)), + std::integral_constant{}), 0)...}; RAJA_UNUSED_VAR(seq_unused_array); return func; @@ -115,7 +116,7 @@ RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_index(UnaryFunc func) { return detail::for_each_index(camp::make_idx_seq_t(), std::move(func)); @@ -128,7 +129,7 @@ UnaryFunc for_each_index(UnaryFunc func) */ RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE +constexpr RAJA_HOST_DEVICE RAJA_INLINE camp::concepts::enable_if_t> for_each(Container&& c, UnaryFunc func) { @@ -144,8 +145,8 @@ RAJA_HOST_DEVICE RAJA_INLINE */ RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const& c, - UnaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const& c, + UnaryFunc func) { return detail::for_each_type(c, std::move(func)); } @@ -156,7 +157,7 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_type(camp::list const& c, */ RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, UnaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, UnaryFunc func) { return detail::for_each_tuple( std::forward(t), std::move(func), @@ -170,7 +171,7 @@ RAJA_HOST_DEVICE RAJA_INLINE UnaryFunc for_each_tuple(Tuple&& t, UnaryFunc func) */ RAJA_SUPPRESS_HD_WARN template -RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, BinaryFunc func) +constexpr RAJA_HOST_DEVICE RAJA_INLINE BinaryFunc for_each_tuple_index(Tuple&& t, BinaryFunc func) { return detail::for_each_tuple_index( std::forward(t), std::move(func), diff --git a/include/RAJA/util/macros.hpp b/include/RAJA/util/macros.hpp index a41179c1a3..e04e915a59 100644 --- a/include/RAJA/util/macros.hpp +++ b/include/RAJA/util/macros.hpp @@ -115,7 +115,7 @@ ******************************************************************************* */ template -RAJA_HOST_DEVICE RAJA_INLINE void RAJA_UNUSED_VAR(T&&...) noexcept +constexpr RAJA_HOST_DEVICE RAJA_INLINE void RAJA_UNUSED_VAR(T&&...) noexcept {} #define RAJA_DIVIDE_CEILING_INT(dividend, divisor) \ From 17d24e698decf2068deb1aff1938a8af939a6f3a Mon Sep 17 00:00:00 2001 From: gberg617 Date: Tue, 14 Oct 2025 14:18:02 -0700 Subject: [PATCH 10/13] started adding various size/dim methods --- include/RAJA/util/SubView.hpp | 103 ++++++++++++++++++++++++++++++++-- 1 file changed, 98 insertions(+), 5 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index 545c2b5e80..b9a38a3caa 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -21,34 +21,55 @@ #include "RAJA/util/for_each.hpp" #include "RAJA/util/macros.hpp" #include "RAJA/util/types.hpp" +#include "camp/helpers.hpp" +#include "camp/number/number.hpp" #include "camp/tuple.hpp" #include "camp/array.hpp" +#include "camp/type_traits/is_same.hpp" namespace RAJA { -// Slice descriptors template struct RangeSlice { - IndexType start, end; + IndexType start_, end_; static constexpr bool reduces_dimension = false; RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { - return start + idx; + return start_ + idx; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType&) const { + return (end_ - start_ + 1); + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return 0; } }; template struct FixedSlice { - IndexType idx; + IndexType idx_; static constexpr bool reduces_dimension = true; RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType&) const { - return idx; + return idx_; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType&) const { + return 1; } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return 0; + } + }; template @@ -58,6 +79,35 @@ struct NoSlice { RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { return idx; } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType& layout) const { + return layout.template get_dim_size(); + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return 0; + } +}; + +template +struct StridedSlice { + IndexType start_, end_, stride_; + + static constexpr bool reduces_dimension = false; + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { + return idx; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType&) const { + return (end_ - start_ + 1) / stride_; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return stride_; + } }; template @@ -103,6 +153,49 @@ class SubView, IndexType> { return camp::get(slices_); } + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size() { + + IndexType prod_dims = 1; + for_each_tuple_index( slices_, + [&](auto slice, auto index) { + const IndexType dim_size = slice.template size(view_.get_layout()); + prod_dims *= (dim_size == 0) ? 1 : dim_size; + }); + + return prod_dims; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size_noproj() { + + IndexType prod_dims = 1; + for_each_tuple_index( slices_, + [&](auto slice, auto index) { + prod_dims *= slice.template size(view_.get_layout()); + }); + + return prod_dims; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_size() { + return camp::get(slices_).template size(view_.get_layout()); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_stride() { + return camp::get(slices_).template stride(); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_layout() { + return view_.get_layout(); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_data() { + return view_.get_data(); + } + template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { static_assert(sizeof...(idxs) == num_sub_indices_, "Wrong number of indices for subview"); From 45206acea8b47b8d0074cc8f5a0a278d0d4604d1 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Tue, 14 Oct 2025 14:55:52 -0700 Subject: [PATCH 11/13] started implementing SubLayout --- include/RAJA/util/SubView.hpp | 73 ++++++++++++++++---------- test/unit/view-layout/test-subview.cpp | 10 +++- 2 files changed, 53 insertions(+), 30 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index b9a38a3caa..bab75b1e1e 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -21,16 +21,12 @@ #include "RAJA/util/for_each.hpp" #include "RAJA/util/macros.hpp" #include "RAJA/util/types.hpp" -#include "camp/helpers.hpp" -#include "camp/number/number.hpp" #include "camp/tuple.hpp" #include "camp/array.hpp" -#include "camp/type_traits/is_same.hpp" namespace RAJA { - template struct RangeSlice { IndexType start_, end_; @@ -97,7 +93,7 @@ struct StridedSlice { static constexpr bool reduces_dimension = false; RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { - return idx; + return start_ + stride_ * idx; } template @@ -119,12 +115,12 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { return map; } -template -class SubView; +template +class SubLayout; -template -class SubView, IndexType> { - ViewType view_; +template +class SubLayout, IndexType> { + const LayoutType& layout_; camp::tuple slices_; static inline constexpr size_t num_slices_ = sizeof...(Slices); static inline constexpr IndexType num_sub_indices_ = ((Slices::reduces_dimension == false ? 1 : 0) + ...); @@ -132,8 +128,8 @@ class SubView, IndexType> { public: - RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) - : view_(view), slices_(slices...) { } + RAJA_INLINE RAJA_HOST_DEVICE constexpr SubLayout(const LayoutType& layout, Slices... slices) + : layout_(layout), slices_(slices...) { } RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slices(Slices... slices) { slices_ = camp::tuple(slices...); @@ -158,7 +154,7 @@ class SubView, IndexType> { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - const IndexType dim_size = slice.template size(view_.get_layout()); + const IndexType dim_size = slice.template size(layout_); prod_dims *= (dim_size == 0) ? 1 : dim_size; }); @@ -170,7 +166,7 @@ class SubView, IndexType> { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - prod_dims *= slice.template size(view_.get_layout()); + prod_dims *= slice.template size(layout_); }); return prod_dims; @@ -178,7 +174,7 @@ class SubView, IndexType> { template RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_size() { - return camp::get(slices_).template size(view_.get_layout()); + return camp::get(slices_).template size(layout_); } template @@ -186,18 +182,8 @@ class SubView, IndexType> { return camp::get(slices_).template stride(); } - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_layout() { - return view_.get_layout(); - } - - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_data() { - return view_.get_data(); - } - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto operator()(Idxs... idxs) const { static_assert(sizeof...(idxs) == num_sub_indices_, "Wrong number of indices for subview"); camp::array arr{idxs...}; @@ -215,12 +201,43 @@ class SubView, IndexType> { } }); - return camp::apply(view_, parent_indices); + return parent_indices; + } +}; + +template +class SubView; + +template +class SubView, IndexType> { + ViewType view_; + SubLayout, IndexType> sublayout_; + +public: + + RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) + : view_(view), sublayout_(view.get_layout(), slices...) { } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_layout() { + return view_.get_layout(); + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_sublayout() { + return sublayout_; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_data() { + return view_.get_data(); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { + return camp::apply(view_, sublayout_(idxs...)); } }; template -SubView(ViewType, Slices...) -> SubView>; +SubView(ViewType, Slices...) -> SubView>; } // namespace RAJA diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index 34851efe2f..b5e563368a 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -76,6 +76,8 @@ TEST(SubView, FixedFirstDimSubView2D) EXPECT_EQ(sv(1), 5); EXPECT_EQ(sv(2), 6); + EXPECT_EQ(sv.get_sublayout().size(), 3); + } TEST(SubView, RangeFirstDimSubView2D) @@ -96,7 +98,9 @@ TEST(SubView, RangeFirstDimSubView2D) EXPECT_EQ(sv(1,1), 8); EXPECT_EQ(sv(1,2), 9); - sv.set_slice<0>(RangeSlice<>{0,1}); + EXPECT_EQ(sv.get_sublayout().size(), 6); + + sv.get_sublayout().set_slice<0>(RangeSlice<>{0,1}); EXPECT_EQ(sv(0,0), 1); EXPECT_EQ(sv(0,1), 2); @@ -106,9 +110,11 @@ TEST(SubView, RangeFirstDimSubView2D) EXPECT_EQ(sv(1,1), 5); EXPECT_EQ(sv(1,2), 6); + EXPECT_EQ(sv.get_sublayout().size(), 6); + } -// To Do tests: subview of subview, "sliding winow" modify slice entries, GPU tests. +// To Do tests: subview of subview // void test_subviewGPU() { // #if defined(RAJA_ENABLE_HIP) From 2f45d02f2f437a7e72a2e71cc19eabb9ecc9b3a8 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Thu, 16 Oct 2025 17:39:42 -0700 Subject: [PATCH 12/13] first attempt at changing to SubLayout --- include/RAJA/util/SubView.hpp | 154 +++++++++------- test/unit/view-layout/test-subview.cpp | 241 ++++++++++++++++++++++--- 2 files changed, 305 insertions(+), 90 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index bab75b1e1e..f846d34138 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -43,7 +43,27 @@ struct RangeSlice { } RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { - return 0; + return 1; + } +}; + +template +struct RangeStartSlice { + IndexType start_; + + static constexpr bool reduces_dimension = false; + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType& idx) const { + return start_ + idx; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType& layout) const { + return (layout.size() - start_); + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return 1; } }; @@ -63,7 +83,7 @@ struct FixedSlice { } RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { - return 0; + return 1; } }; @@ -82,7 +102,7 @@ struct NoSlice { } RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { - return 0; + return 1; } }; @@ -98,7 +118,7 @@ struct StridedSlice { template RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType&) const { - return (end_ - start_ + 1) / stride_; + return (end_ - start_) / stride_ + 1; } RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { @@ -107,7 +127,7 @@ struct StridedSlice { }; template -RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_slice_to_parent_index_map() { IndexType sub_idx = 0; IndexType i = 0; camp::array map = {}; @@ -115,84 +135,111 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_subview_index_map() { return map; } -template -class SubLayout; +template +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_parent_to_slice_index_map() { + IndexType sub_idx = 0; + IndexType i = 0; + camp::array map = {}; + + auto process_slice = [&](auto slice_type) constexpr { + if constexpr (!decltype(slice_type)::reduces_dimension) { + map[sub_idx++] = i++; + } else { + i++; + } + }; + + (process_slice(Slices{}), ...); + + return map; +} + +template +struct SubLayout; + +template +struct SubLayout, IndexType> { -template -class SubLayout, IndexType> { - const LayoutType& layout_; + using IndexLinear = IndexType; + + const LayoutType& parent_layout_; camp::tuple slices_; + static inline constexpr size_t num_slices_ = sizeof...(Slices); - static inline constexpr IndexType num_sub_indices_ = ((Slices::reduces_dimension == false ? 1 : 0) + ...); - static inline constexpr camp::array map_ = make_subview_index_map(); -public: + static inline constexpr + IndexType n_dims = ((Slices::reduces_dimension == false ? 1 : 0) + ...); - RAJA_INLINE RAJA_HOST_DEVICE constexpr SubLayout(const LayoutType& layout, Slices... slices) - : layout_(layout), slices_(slices...) { } + static inline constexpr + camp::array slice_to_parent_map_ = + make_slice_to_parent_index_map(); - RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slices(Slices... slices) { - slices_ = camp::tuple(slices...); - } + static inline constexpr + camp::array parent_to_slice_map_ = + make_parent_to_slice_index_map(); - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr void set_slice(Slice slice) { - camp::get(slices_) = slice; + RAJA_INLINE RAJA_HOST_DEVICE constexpr SubLayout(const LayoutType& parent_layout, Slices... slices) + : parent_layout_(parent_layout), slices_(slices...) { } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_parent_layout() const { + return parent_layout_; } - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slices() { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slices() const { return slices_; } template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slice() { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slice() const { return camp::get(slices_); } - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size() { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size() const { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - const IndexType dim_size = slice.template size(layout_); + const IndexType dim_size = slice.template size(parent_layout_); prod_dims *= (dim_size == 0) ? 1 : dim_size; }); return prod_dims; } - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size_noproj() { + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto size_noproj() const { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - prod_dims *= slice.template size(layout_); + prod_dims *= slice.template size(parent_layout_); }); return prod_dims; } template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_size() { - return camp::get(slices_).template size(layout_); + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_size() const { + constexpr auto SliceDim = parent_to_slice_map_[DIM]; + return camp::get(slices_).template size(parent_layout_); } template - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_stride() { - return camp::get(slices_).template stride(); + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_stride() const { + constexpr auto SliceDim = parent_to_slice_map_[DIM]; + return camp::get(slices_).stride(); } template RAJA_INLINE RAJA_HOST_DEVICE constexpr auto operator()(Idxs... idxs) const { - static_assert(sizeof...(idxs) == num_sub_indices_, "Wrong number of indices for subview"); + static_assert(sizeof...(idxs) == n_dims, "Wrong number of indices for subview"); - camp::array arr{idxs...}; + camp::array arr{idxs...}; camp::array parent_indices; for_each_tuple_index( slices_, [&](auto slice, auto index) { - if (map_[index] >= 0) { - parent_indices[index] = slice.map_index(arr[map_[index]]); + if (slice_to_parent_map_[index] >= 0) { + parent_indices[index] = slice.map_index(arr[slice_to_parent_map_[index]]); } else { // map_index will not need index values for dimension-reducing slices // so we pass a "dummy" value. @@ -201,43 +248,12 @@ class SubLayout, IndexType> { } }); - return parent_indices; - } -}; - -template -class SubView; - -template -class SubView, IndexType> { - ViewType view_; - SubLayout, IndexType> sublayout_; - -public: - - RAJA_INLINE RAJA_HOST_DEVICE constexpr SubView(ViewType view, Slices... slices) - : view_(view), sublayout_(view.get_layout(), slices...) { } - - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_layout() { - return view_.get_layout(); - } - - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_sublayout() { - return sublayout_; - } - - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_data() { - return view_.get_data(); - } - - template - RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType operator()(Idxs... idxs) const { - return camp::apply(view_, sublayout_(idxs...)); + return camp::apply(parent_layout_, parent_indices); } }; -template -SubView(ViewType, Slices...) -> SubView>; +template +SubLayout(LayoutType, Slices...) -> SubLayout>; } // namespace RAJA diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index b5e563368a..8e5eb35a62 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -14,6 +14,18 @@ using namespace RAJA; +template +auto make_subview(ViewType view, Slices... slices) { + using SubLayoutType = SubLayout>; + return View(view.get_data(), SubLayoutType(view.get_layout(), slices...)); +} + +template +auto make_submultiview(ViewType view, Slices... slices) { + using SubLayoutType = SubLayout>; + return MultiView(view.get_data(), SubLayoutType(view.get_layout(), slices...)); +} + TEST(SubView, RangeSubView1D) { @@ -22,73 +34,159 @@ TEST(SubView, RangeSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[1:3] - auto sv = SubView(view, RangeSlice<>{1,3}); + auto sv = make_subview(view, RangeSlice<>{1,3}); + auto sublayout = sv.get_layout(); EXPECT_EQ(sv(0), 2); EXPECT_EQ(sv(1), 3); EXPECT_EQ(sv(2), 4); + EXPECT_EQ(sublayout(0), 1); + EXPECT_EQ(sublayout(1), 2); + EXPECT_EQ(sublayout(2), 3); + + EXPECT_EQ(sublayout.size(), 3); +} + +TEST(SubView, RangeStartSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[2:] + auto sv = make_subview(view, RangeStartSlice<>{2}); + auto sublayout = sv.get_layout(); + + EXPECT_EQ(sv(0), 3); + EXPECT_EQ(sv(1), 4); + EXPECT_EQ(sv(2), 5); + + EXPECT_EQ(sublayout.size(), 3); +} + +TEST(SubView, StridedSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[0:3:2] + auto sv = make_subview(view, StridedSlice<>{0,3,2}); + auto sublayout = sv.get_layout(); + + EXPECT_EQ(sv(0), 1); + EXPECT_EQ(sv(1), 3); + + EXPECT_EQ(sublayout(0), 0); + EXPECT_EQ(sublayout(1), 2); + + EXPECT_EQ(sublayout.size(), 2); +} + +TEST(SubView, FixedSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[1] + auto sv = make_subview(view, FixedSlice<>{1}); + auto sublayout = sv.get_layout(); + + // 0D views don't work + // EXPECT_EQ(sv(), 2); + EXPECT_EQ(sublayout(), 1); + + EXPECT_EQ(sublayout.size(), 1); + EXPECT_EQ(sublayout.size_noproj(), 1); } TEST(SubView, RangeSubView2D) { - Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + Index_type a[3][3] = {{1,2,3}, + {4,5,6}, + {7,8,9}}; View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1:2] - auto sv = SubView(view, RangeSlice<>{1,2}, RangeSlice<>{1,2}); + auto sv = make_subview(view, RangeSlice<>{1,2}, RangeSlice<>{1,2}); + auto sublayout = sv.get_layout(); EXPECT_EQ(sv(0,0), 5); EXPECT_EQ(sv(0,1), 6); EXPECT_EQ(sv(1,0), 8); EXPECT_EQ(sv(1,1), 9); + EXPECT_EQ(sublayout.size(), 4); + EXPECT_EQ(sublayout.get_dim_size<0>(), 2); + EXPECT_EQ(sublayout.get_dim_size<1>(), 2); + EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + EXPECT_EQ(sublayout.get_dim_stride<1>(), 1); + } TEST(SubView, RangeFixedSubView2D) { - Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + Index_type a[3][3] = {{1,2,3}, + {4,5,6}, + {7,8,9}}; View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1] - auto sv = SubView(view, RangeSlice<>{1,2}, FixedSlice<>{1}); + auto sv = make_subview(view, RangeSlice<>{1,2}, FixedSlice<>{1}); + auto sublayout = sv.get_layout(); EXPECT_EQ(sv(0), 5); EXPECT_EQ(sv(1), 8); + EXPECT_EQ(sublayout.size(), 2); + EXPECT_EQ(sublayout.get_dim_size<0>(), 2); + EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + } TEST(SubView, FixedFirstDimSubView2D) { - Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + Index_type a[3][3] = {{1,2,3}, + {4,5,6}, + {7,8,9}}; View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1,:] - auto sv = SubView(view, FixedSlice<>{1}, NoSlice{}); + auto sv = make_subview(view, FixedSlice<>{1}, NoSlice{}); + auto sublayout = sv.get_layout(); EXPECT_EQ(sv(0), 4); EXPECT_EQ(sv(1), 5); EXPECT_EQ(sv(2), 6); - EXPECT_EQ(sv.get_sublayout().size(), 3); - + EXPECT_EQ(sublayout.size(), 3); + EXPECT_EQ(sublayout.get_dim_size<0>(), 3); + EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); } TEST(SubView, RangeFirstDimSubView2D) { - Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + Index_type a[3][3] = {{1,2,3}, + {4,5,6}, + {7,8,9}}; View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,:] - auto sv = SubView(view, RangeSlice<>{1,2}, NoSlice{}); + auto sv = make_subview(view, RangeSlice<>{1,2}, NoSlice{}); + auto sublayout = sv.get_layout(); EXPECT_EQ(sv(0,0), 4); EXPECT_EQ(sv(0,1), 5); @@ -98,23 +196,124 @@ TEST(SubView, RangeFirstDimSubView2D) EXPECT_EQ(sv(1,1), 8); EXPECT_EQ(sv(1,2), 9); - EXPECT_EQ(sv.get_sublayout().size(), 6); + EXPECT_EQ(sublayout.size(), 6); + EXPECT_EQ(sublayout.get_dim_size<0>(), 2); + EXPECT_EQ(sublayout.get_dim_size<1>(), 3); + EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + EXPECT_EQ(sublayout.get_dim_stride<1>(), 1); - sv.get_sublayout().set_slice<0>(RangeSlice<>{0,1}); +} - EXPECT_EQ(sv(0,0), 1); - EXPECT_EQ(sv(0,1), 2); - EXPECT_EQ(sv(0,2), 3); +TEST(SubView, RangeFirstDimStridedSecondDimSubView2D) +{ + + Index_type a[3][6] = {{1,2,3,4,5,6}, + {7,8,9,10,11,12}, + {13,14,15,16,17,18}}; + + View> view(&a[0][0], Layout<2>(3,6)); + + // sv = View[1:2,1:5:2] + auto sv = make_subview(view, RangeSlice<>{1,2}, StridedSlice<>{1,5,2}); + auto sublayout = sv.get_layout(); + + EXPECT_EQ(sv(0,0), 8); + EXPECT_EQ(sv(0,1), 10); + EXPECT_EQ(sv(0,2), 12); + + EXPECT_EQ(sv(1,0), 14); + EXPECT_EQ(sv(1,1), 16); + EXPECT_EQ(sv(1,2), 18); + + EXPECT_EQ(sublayout.size(), 6); + EXPECT_EQ(sublayout.get_dim_size<0>(), 2); + EXPECT_EQ(sublayout.get_dim_size<1>(), 3); + EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + EXPECT_EQ(sublayout.get_dim_stride<1>(), 2); + +} + +TEST(SubView, SubViewOfSubView2D) +{ + + Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; + + View> view(&a[0][0], Layout<2>(3,3)); + + // sv = View[1:2,:] + auto sv = make_subview(view, RangeSlice<>{1,2}, NoSlice{}); + + + EXPECT_EQ(sv.get_layout().size(), 6); + + // sv2 = sv[0:1,1:2] + auto sv2 = make_subview(sv, RangeSlice<>{0,1}, RangeSlice<>{1,2}); - EXPECT_EQ(sv(1,0), 4); - EXPECT_EQ(sv(1,1), 5); - EXPECT_EQ(sv(1,2), 6); + EXPECT_EQ(sv2(0,0), 5); + EXPECT_EQ(sv2(0,1), 6); - EXPECT_EQ(sv.get_sublayout().size(), 6); + EXPECT_EQ(sv2(1,0), 8); + EXPECT_EQ(sv2(1,1), 9); + + EXPECT_EQ(sv2.get_layout().size(), 4); + + // sv3 = sv2[:,1] + auto sv3 = make_subview(sv2, NoSlice{}, FixedSlice<>{1}); + + EXPECT_EQ(sv3(0), 6); + EXPECT_EQ(sv3(1), 9); + + EXPECT_EQ(sv3.get_layout().size(), 2); } -// To Do tests: subview of subview +TEST(SubView, SubViewOfMultiView2D) +{ + + Index_type data_squared[4]; + Index_type data_cubed[4]; + + for (int i = 0; i < 4; i ++ ) { + data_squared[i] = i*i; + } + + for (int i = 0; i < 4; i ++ ) { + data_cubed[i] = i*i*i; + } + + Index_type* data_array[2]; + data_array[0] = data_squared; + data_array[1] = data_cubed; + + Index_type index_list[4] = {3,1,2,0}; + + auto index_tuple = make_index_tuple(IndexList<>{&index_list[0]}); + auto index_layout = make_index_layout(index_tuple, 4); + + auto view = MultiView > >(data_array, index_layout); + + // sv = MultiView[:,1:2] + auto sv = make_submultiview(view, RangeSlice<>{1,2}); + + EXPECT_EQ(sv.get_layout().size(), 2); + EXPECT_EQ(sv.get_layout().get_dim_size<0>(), 2); + + EXPECT_EQ(sv(0,0), 1); + EXPECT_EQ(sv(0,1), 4); + + EXPECT_EQ(sv(1,0), 1); + EXPECT_EQ(sv(1,1), 8); + + // sv = MultiView[:,2] + auto sv2 = make_submultiview(view, FixedSlice<>{2}); + + // this size does not look right + EXPECT_EQ(sv2.get_layout().size(), 1); + + EXPECT_EQ(sv2(0), 4); + EXPECT_EQ(sv2(1), 8); + +} // void test_subviewGPU() { // #if defined(RAJA_ENABLE_HIP) From 1e928d3bfa70ae98b2757cada6bca18b1fad0938 Mon Sep 17 00:00:00 2001 From: gberg617 Date: Mon, 17 Nov 2025 17:25:44 -0800 Subject: [PATCH 13/13] made some changes to testing subview/sublayouts --- include/RAJA/util/SubView.hpp | 34 +++-- include/RAJA/util/TypedViewBase.hpp | 17 ++- include/RAJA/util/View.hpp | 27 ++++ test/unit/view-layout/test-subview.cpp | 199 +++++++++++++++---------- 4 files changed, 183 insertions(+), 94 deletions(-) diff --git a/include/RAJA/util/SubView.hpp b/include/RAJA/util/SubView.hpp index f846d34138..e761a1f4bd 100644 --- a/include/RAJA/util/SubView.hpp +++ b/include/RAJA/util/SubView.hpp @@ -155,14 +155,22 @@ RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_parent_to_slice_index_map() { } template -struct SubLayout; +struct SubRegion; + +/* SubLayout is a semantic alias for a SubRegion whose parent is a layout */ +template +using SubLayout = SubRegion; + +/* SubView is a semantic alias for a SubRegion whose parent is a view */ +template +using SubView = SubRegion; template -struct SubLayout, IndexType> { +struct SubRegion, IndexType> { using IndexLinear = IndexType; - const LayoutType& parent_layout_; + const LayoutType& parent_; camp::tuple slices_; static inline constexpr size_t num_slices_ = sizeof...(Slices); @@ -178,11 +186,11 @@ struct SubLayout, IndexType> { camp::array parent_to_slice_map_ = make_parent_to_slice_index_map(); - RAJA_INLINE RAJA_HOST_DEVICE constexpr SubLayout(const LayoutType& parent_layout, Slices... slices) - : parent_layout_(parent_layout), slices_(slices...) { } + RAJA_INLINE RAJA_HOST_DEVICE constexpr SubRegion(const LayoutType& parent, Slices... slices) + : parent_(parent), slices_(slices...) { } - RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_parent_layout() const { - return parent_layout_; + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_parent() const { + return parent_; } RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slices() const { @@ -199,7 +207,7 @@ struct SubLayout, IndexType> { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - const IndexType dim_size = slice.template size(parent_layout_); + const IndexType dim_size = slice.template size(parent_); prod_dims *= (dim_size == 0) ? 1 : dim_size; }); @@ -211,7 +219,7 @@ struct SubLayout, IndexType> { IndexType prod_dims = 1; for_each_tuple_index( slices_, [&](auto slice, auto index) { - prod_dims *= slice.template size(parent_layout_); + prod_dims *= slice.template size(parent_); }); return prod_dims; @@ -220,7 +228,7 @@ struct SubLayout, IndexType> { template RAJA_INLINE RAJA_HOST_DEVICE constexpr auto get_dim_size() const { constexpr auto SliceDim = parent_to_slice_map_[DIM]; - return camp::get(slices_).template size(parent_layout_); + return camp::get(slices_).template size(parent_); } template @@ -231,7 +239,7 @@ struct SubLayout, IndexType> { template RAJA_INLINE RAJA_HOST_DEVICE constexpr auto operator()(Idxs... idxs) const { - static_assert(sizeof...(idxs) == n_dims, "Wrong number of indices for subview"); + static_assert(sizeof...(idxs) == n_dims, "Wrong number of indices"); camp::array arr{idxs...}; camp::array parent_indices; @@ -248,12 +256,12 @@ struct SubLayout, IndexType> { } }); - return camp::apply(parent_layout_, parent_indices); + return camp::apply(parent_, parent_indices); } }; template -SubLayout(LayoutType, Slices...) -> SubLayout>; +SubRegion(LayoutType, Slices...) -> SubRegion>; } // namespace RAJA diff --git a/include/RAJA/util/TypedViewBase.hpp b/include/RAJA/util/TypedViewBase.hpp index 359783845f..770048a925 100644 --- a/include/RAJA/util/TypedViewBase.hpp +++ b/include/RAJA/util/TypedViewBase.hpp @@ -669,16 +669,31 @@ class ViewBase constexpr layout_type const& get_layout() const { return m_layout; } RAJA_HOST_DEVICE - RAJA_INLINE constexpr linear_index_type size() const { return m_layout.size(); } + RAJA_HOST_DEVICE + RAJA_INLINE + constexpr linear_index_type size_noproj() const { return m_layout.size_noproj(); } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr linear_index_type get_dim_stride() const + { + return m_layout.template get_dim_stride(); + } + template RAJA_HOST_DEVICE RAJA_INLINE constexpr linear_index_type get_dim_size() const { return m_layout.template get_dim_size(); } + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr linear_index_type get_dim_begin() const + { + return m_layout.template get_dim_begin(); + } + template RAJA_HOST_DEVICE RAJA_INLINE constexpr view_return_type_t; using nc_pointer_type = NonConstPointerType; using NonConstView = @@ -228,6 +229,32 @@ struct MultiView return layout; } + RAJA_HOST_DEVICE + RAJA_INLINE + constexpr linear_index_type size() const { return layout.size(); } + + RAJA_HOST_DEVICE + RAJA_INLINE + constexpr linear_index_type size_noproj() const { return layout.size_noproj(); } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr linear_index_type get_dim_stride() const + { + return layout.template get_dim_stride(); + } + + template + RAJA_HOST_DEVICE RAJA_INLINE constexpr linear_index_type get_dim_size() const + { + return layout.template get_dim_size(); + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr linear_index_type get_dim_begin() const + { + return layout.template get_dim_begin(); + } + RAJA_HOST_DEVICE RAJA_INLINE constexpr pointer_type get_data() const { return pointer_type(data); diff --git a/test/unit/view-layout/test-subview.cpp b/test/unit/view-layout/test-subview.cpp index 8e5eb35a62..08dbe258fa 100644 --- a/test/unit/view-layout/test-subview.cpp +++ b/test/unit/view-layout/test-subview.cpp @@ -14,19 +14,56 @@ using namespace RAJA; +/* helper to create a RAJA View with a sliced SubLayout */ template -auto make_subview(ViewType view, Slices... slices) { +auto make_view_with_sublayout(ViewType& view, Slices... slices) { using SubLayoutType = SubLayout>; return View(view.get_data(), SubLayoutType(view.get_layout(), slices...)); } +/* helper to create a sliced SubView without modifying underlying layout */ template -auto make_submultiview(ViewType view, Slices... slices) { +auto make_subview_with_layout(ViewType& view, Slices... slices) { + using SubViewType = SubView>; + return SubViewType(view, slices...); +} + +template +auto make_multiview_with_sublayout(ViewType& view, Slices... slices) { using SubLayoutType = SubLayout>; return MultiView(view.get_data(), SubLayoutType(view.get_layout(), slices...)); } -TEST(SubView, RangeSubView1D) +struct UseViewWithSubLayout { + + template + auto operator()(ViewType& view, Slices... slices) const { + return make_view_with_sublayout(view, slices...); + } + + template + static auto& get_subregion(ViewType& sv) { return sv.get_layout(); } + +}; + +struct UseSubViewWithLayout { + + template + auto operator()(ViewType& view, Slices... slices) const { + return make_subview_with_layout(view, slices...); + } + + template + static auto& get_subregion(ViewType& sv) { return sv; } +}; + +template +class SubViewTest : public ::testing::Test {}; + +using FactoryTypes = ::testing::Types; +TYPED_TEST_SUITE(SubViewTest, FactoryTypes); + +TYPED_TEST(SubViewTest, RangeSubView1D) { Index_type a[] = {1,2,3,4,5}; @@ -34,21 +71,18 @@ TEST(SubView, RangeSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[1:3] - auto sv = make_subview(view, RangeSlice<>{1,3}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeSlice<>{1,3}); EXPECT_EQ(sv(0), 2); EXPECT_EQ(sv(1), 3); EXPECT_EQ(sv(2), 4); - EXPECT_EQ(sublayout(0), 1); - EXPECT_EQ(sublayout(1), 2); - EXPECT_EQ(sublayout(2), 3); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 3); - EXPECT_EQ(sublayout.size(), 3); } -TEST(SubView, RangeStartSubView1D) +TYPED_TEST(SubViewTest, RangeStartSubView1D) { Index_type a[] = {1,2,3,4,5}; @@ -56,17 +90,17 @@ TEST(SubView, RangeStartSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[2:] - auto sv = make_subview(view, RangeStartSlice<>{2}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeStartSlice<>{2}); EXPECT_EQ(sv(0), 3); EXPECT_EQ(sv(1), 4); EXPECT_EQ(sv(2), 5); - EXPECT_EQ(sublayout.size(), 3); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 3); } -TEST(SubView, StridedSubView1D) +TYPED_TEST(SubViewTest, StridedSubView1D) { Index_type a[] = {1,2,3,4,5}; @@ -74,19 +108,16 @@ TEST(SubView, StridedSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[0:3:2] - auto sv = make_subview(view, StridedSlice<>{0,3,2}); - auto sublayout = sv.get_layout(); + auto sv = make_view_with_sublayout(view, StridedSlice<>{0,3,2}); EXPECT_EQ(sv(0), 1); EXPECT_EQ(sv(1), 3); - EXPECT_EQ(sublayout(0), 0); - EXPECT_EQ(sublayout(1), 2); - - EXPECT_EQ(sublayout.size(), 2); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 2); } -TEST(SubView, FixedSubView1D) +TYPED_TEST(SubViewTest, FixedSubView1D) { Index_type a[] = {1,2,3,4,5}; @@ -94,18 +125,14 @@ TEST(SubView, FixedSubView1D) View> view(&a[0], Layout<1>(5)); // sv = View[1] - auto sv = make_subview(view, FixedSlice<>{1}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, FixedSlice<>{1}); - // 0D views don't work - // EXPECT_EQ(sv(), 2); - EXPECT_EQ(sublayout(), 1); - - EXPECT_EQ(sublayout.size(), 1); - EXPECT_EQ(sublayout.size_noproj(), 1); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 1); + EXPECT_EQ(sr.size_noproj(), 1); } -TEST(SubView, RangeSubView2D) +TYPED_TEST(SubViewTest, RangeSubView2D) { Index_type a[3][3] = {{1,2,3}, @@ -115,23 +142,23 @@ TEST(SubView, RangeSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1:2] - auto sv = make_subview(view, RangeSlice<>{1,2}, RangeSlice<>{1,2}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeSlice<>{1,2}, RangeSlice<>{1,2}); EXPECT_EQ(sv(0,0), 5); EXPECT_EQ(sv(0,1), 6); EXPECT_EQ(sv(1,0), 8); EXPECT_EQ(sv(1,1), 9); - EXPECT_EQ(sublayout.size(), 4); - EXPECT_EQ(sublayout.get_dim_size<0>(), 2); - EXPECT_EQ(sublayout.get_dim_size<1>(), 2); - EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); - EXPECT_EQ(sublayout.get_dim_stride<1>(), 1); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 4); + EXPECT_EQ(sr.template get_dim_size<0>(), 2); + EXPECT_EQ(sr.template get_dim_size<1>(), 2); + EXPECT_EQ(sr.template get_dim_stride<0>(), 1); + EXPECT_EQ(sr.template get_dim_stride<1>(), 1); } -TEST(SubView, RangeFixedSubView2D) +TYPED_TEST(SubViewTest, RangeFixedSubView2D) { Index_type a[3][3] = {{1,2,3}, @@ -141,19 +168,19 @@ TEST(SubView, RangeFixedSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,1] - auto sv = make_subview(view, RangeSlice<>{1,2}, FixedSlice<>{1}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeSlice<>{1,2}, FixedSlice<>{1}); EXPECT_EQ(sv(0), 5); EXPECT_EQ(sv(1), 8); - EXPECT_EQ(sublayout.size(), 2); - EXPECT_EQ(sublayout.get_dim_size<0>(), 2); - EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 2); + EXPECT_EQ(sr.template get_dim_size<0>(), 2); + EXPECT_EQ(sr.template get_dim_stride<0>(), 1); } -TEST(SubView, FixedFirstDimSubView2D) +TYPED_TEST(SubViewTest, FixedFirstDimSubView2D) { Index_type a[3][3] = {{1,2,3}, @@ -163,19 +190,19 @@ TEST(SubView, FixedFirstDimSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1,:] - auto sv = make_subview(view, FixedSlice<>{1}, NoSlice{}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, FixedSlice<>{1}, NoSlice{}); EXPECT_EQ(sv(0), 4); EXPECT_EQ(sv(1), 5); EXPECT_EQ(sv(2), 6); - EXPECT_EQ(sublayout.size(), 3); - EXPECT_EQ(sublayout.get_dim_size<0>(), 3); - EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 3); + EXPECT_EQ(sr.template get_dim_size<0>(), 3); + EXPECT_EQ(sr.template get_dim_stride<0>(), 1); } -TEST(SubView, RangeFirstDimSubView2D) +TYPED_TEST(SubViewTest, RangeFirstDimSubView2D) { Index_type a[3][3] = {{1,2,3}, @@ -185,8 +212,7 @@ TEST(SubView, RangeFirstDimSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,:] - auto sv = make_subview(view, RangeSlice<>{1,2}, NoSlice{}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeSlice<>{1,2}, NoSlice{}); EXPECT_EQ(sv(0,0), 4); EXPECT_EQ(sv(0,1), 5); @@ -196,15 +222,16 @@ TEST(SubView, RangeFirstDimSubView2D) EXPECT_EQ(sv(1,1), 8); EXPECT_EQ(sv(1,2), 9); - EXPECT_EQ(sublayout.size(), 6); - EXPECT_EQ(sublayout.get_dim_size<0>(), 2); - EXPECT_EQ(sublayout.get_dim_size<1>(), 3); - EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); - EXPECT_EQ(sublayout.get_dim_stride<1>(), 1); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 6); + EXPECT_EQ(sr.template get_dim_size<0>(), 2); + EXPECT_EQ(sr.template get_dim_size<1>(), 3); + EXPECT_EQ(sr.template get_dim_stride<0>(), 1); + EXPECT_EQ(sr.template get_dim_stride<1>(), 1); } -TEST(SubView, RangeFirstDimStridedSecondDimSubView2D) +TYPED_TEST(SubViewTest, RangeFirstDimStridedSecondDimSubView2D) { Index_type a[3][6] = {{1,2,3,4,5,6}, @@ -214,8 +241,7 @@ TEST(SubView, RangeFirstDimStridedSecondDimSubView2D) View> view(&a[0][0], Layout<2>(3,6)); // sv = View[1:2,1:5:2] - auto sv = make_subview(view, RangeSlice<>{1,2}, StridedSlice<>{1,5,2}); - auto sublayout = sv.get_layout(); + auto sv = TypeParam{}(view, RangeSlice<>{1,2}, StridedSlice<>{1,5,2}); EXPECT_EQ(sv(0,0), 8); EXPECT_EQ(sv(0,1), 10); @@ -225,15 +251,16 @@ TEST(SubView, RangeFirstDimStridedSecondDimSubView2D) EXPECT_EQ(sv(1,1), 16); EXPECT_EQ(sv(1,2), 18); - EXPECT_EQ(sublayout.size(), 6); - EXPECT_EQ(sublayout.get_dim_size<0>(), 2); - EXPECT_EQ(sublayout.get_dim_size<1>(), 3); - EXPECT_EQ(sublayout.get_dim_stride<0>(), 1); - EXPECT_EQ(sublayout.get_dim_stride<1>(), 2); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 6); + EXPECT_EQ(sr.template get_dim_size<0>(), 2); + EXPECT_EQ(sr.template get_dim_size<1>(), 3); + EXPECT_EQ(sr.template get_dim_stride<0>(), 1); + EXPECT_EQ(sr.template get_dim_stride<1>(), 2); } -TEST(SubView, SubViewOfSubView2D) +TYPED_TEST(SubViewTest, SubViewOfSubView2D) { Index_type a[3][3] = {{1,2,3},{4,5,6},{7,8,9}}; @@ -241,13 +268,14 @@ TEST(SubView, SubViewOfSubView2D) View> view(&a[0][0], Layout<2>(3,3)); // sv = View[1:2,:] - auto sv = make_subview(view, RangeSlice<>{1,2}, NoSlice{}); + auto sv = TypeParam{}(view, RangeSlice<>{1,2}, NoSlice{}); - EXPECT_EQ(sv.get_layout().size(), 6); + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 6); // sv2 = sv[0:1,1:2] - auto sv2 = make_subview(sv, RangeSlice<>{0,1}, RangeSlice<>{1,2}); + auto sv2 =TypeParam{}(sv, RangeSlice<>{0,1}, RangeSlice<>{1,2}); EXPECT_EQ(sv2(0,0), 5); EXPECT_EQ(sv2(0,1), 6); @@ -255,19 +283,19 @@ TEST(SubView, SubViewOfSubView2D) EXPECT_EQ(sv2(1,0), 8); EXPECT_EQ(sv2(1,1), 9); - EXPECT_EQ(sv2.get_layout().size(), 4); + EXPECT_EQ(TypeParam::get_subregion(sv2).size(), 4); // sv3 = sv2[:,1] - auto sv3 = make_subview(sv2, NoSlice{}, FixedSlice<>{1}); + auto sv3 =TypeParam{}(sv2, NoSlice{}, FixedSlice<>{1}); EXPECT_EQ(sv3(0), 6); EXPECT_EQ(sv3(1), 9); - EXPECT_EQ(sv3.get_layout().size(), 2); + EXPECT_EQ(TypeParam::get_subregion(sv3).size(), 2); } -TEST(SubView, SubViewOfMultiView2D) +TEST(SubViewMultiViewTest, SubViewOfMultiView2D) { Index_type data_squared[4]; @@ -293,7 +321,7 @@ TEST(SubView, SubViewOfMultiView2D) auto view = MultiView > >(data_array, index_layout); // sv = MultiView[:,1:2] - auto sv = make_submultiview(view, RangeSlice<>{1,2}); + auto sv = make_multiview_with_sublayout(view, RangeSlice<>{1,2}); EXPECT_EQ(sv.get_layout().size(), 2); EXPECT_EQ(sv.get_layout().get_dim_size<0>(), 2); @@ -304,15 +332,26 @@ TEST(SubView, SubViewOfMultiView2D) EXPECT_EQ(sv(1,0), 1); EXPECT_EQ(sv(1,1), 8); - // sv = MultiView[:,2] - auto sv2 = make_submultiview(view, FixedSlice<>{2}); + // sv2 = MultiView[1,1:2] + auto sv2 = make_subview_with_layout(view, FixedSlice<>{1}, RangeSlice<>{1,2}); - // this size does not look right - EXPECT_EQ(sv2.get_layout().size(), 1); + // the parent layout is a MultiView + EXPECT_EQ(sv2.get_parent().get_layout().size(), 4); + EXPECT_EQ(sv2.size(), 2); - EXPECT_EQ(sv2(0), 4); + EXPECT_EQ(sv2(0), 1); EXPECT_EQ(sv2(1), 8); + // sv3 = MultiView[:,2] + auto sv3 = make_multiview_with_sublayout(view, FixedSlice<>{2}); + + // this size corresponds to the sliced sublayout (0D) + // which is sliced from the original MultiView's 1D layout + EXPECT_EQ(sv3.get_layout().size(), 1); + + EXPECT_EQ(sv3(0), 4); + EXPECT_EQ(sv3(1), 8); + } // void test_subviewGPU() { @@ -331,7 +370,7 @@ TEST(SubView, SubViewOfMultiView2D) // #endif // } -// TEST(SubView, RangeFirstDimSubView2DGPU) +// TYPED_TEST(SubViewTest, RangeFirstDimSubView2DGPU) // { // test_subviewGPU(); // } \ No newline at end of file