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..e761a1f4bd --- /dev/null +++ b/include/RAJA/util/SubView.hpp @@ -0,0 +1,268 @@ +/*! + ****************************************************************************** + * + * \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/for_each.hpp" +#include "RAJA/util/macros.hpp" +#include "RAJA/util/types.hpp" +#include "camp/tuple.hpp" +#include "camp/array.hpp" + +namespace RAJA +{ + +template +struct RangeSlice { + 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; + } + + 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 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; + } +}; + +template +struct FixedSlice { + IndexType idx_; + + static constexpr bool reduces_dimension = true; + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType map_index(const IndexType&) const { + 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 1; + } + +}; + +template +struct NoSlice { + 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& layout) const { + return layout.template get_dim_size(); + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return 1; + } +}; + +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 start_ + stride_ * idx; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType size(const LayoutType&) const { + return (end_ - start_) / stride_ + 1; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr IndexType stride() const { + return stride_; + } +}; + +template +RAJA_INLINE RAJA_HOST_DEVICE constexpr auto make_slice_to_parent_index_map() { + IndexType sub_idx = 0; + IndexType i = 0; + camp::array map = {}; + ((map[i++] = (Slices::reduces_dimension ? -1 : sub_idx++)), ...); + return map; +} + +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 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 SubRegion, IndexType> { + + using IndexLinear = IndexType; + + const LayoutType& parent_; + camp::tuple slices_; + + static inline constexpr size_t num_slices_ = sizeof...(Slices); + + static inline constexpr + IndexType n_dims = ((Slices::reduces_dimension == false ? 1 : 0) + ...); + + static inline constexpr + camp::array slice_to_parent_map_ = + make_slice_to_parent_index_map(); + + static inline constexpr + camp::array parent_to_slice_map_ = + make_parent_to_slice_index_map(); + + 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() const { + return parent_; + } + + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slices() const { + return slices_; + } + + template + RAJA_INLINE RAJA_HOST_DEVICE constexpr auto& get_slice() const { + return camp::get(slices_); + } + + 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(parent_); + prod_dims *= (dim_size == 0) ? 1 : dim_size; + }); + + return prod_dims; + } + + 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(parent_); + }); + + return prod_dims; + } + + 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_); + } + + template + 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) == n_dims, "Wrong number of indices"); + + camp::array arr{idxs...}; + camp::array parent_indices; + + for_each_tuple_index( slices_, + [&](auto slice, auto 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. + constexpr IndexType dummy_value = -1; + parent_indices[index] = slice.map_index(dummy_value); + } + }); + + return camp::apply(parent_, parent_indices); + } +}; + +template +SubRegion(LayoutType, Slices...) -> SubRegion>; + +} // namespace RAJA + +#endif \ No newline at end of file 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/include/RAJA/util/for_each.hpp b/include/RAJA/util/for_each.hpp index 708d3115e8..dc05c83715 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,12 +39,26 @@ namespace RAJA namespace detail { +// compile time expansion applying func to each of the indices +RAJA_SUPPRESS_HD_WARN +template +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 + 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 -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) { @@ -54,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)...}; @@ -67,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 @@ -79,8 +96,32 @@ 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 +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)), + std::integral_constant{}), 0)...}; + RAJA_UNUSED_VAR(seq_unused_array); + + return func; +} + } // namespace detail +RAJA_SUPPRESS_HD_WARN +template +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)); +} + /*! \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 @@ -88,8 +129,8 @@ 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> +constexpr RAJA_HOST_DEVICE RAJA_INLINE + camp::concepts::enable_if_t> for_each(Container&& c, UnaryFunc func) { using std::begin; @@ -104,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)); } @@ -116,13 +157,27 @@ 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), 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 +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), + camp::make_idx_seq_t>::value> {}); +} + } // namespace RAJA #endif 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) \ 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-subview.cpp b/test/unit/view-layout/test-subview.cpp new file mode 100644 index 0000000000..08dbe258fa --- /dev/null +++ b/test/unit/view-layout/test-subview.cpp @@ -0,0 +1,376 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// 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/SubView.hpp" +#include "RAJA/util/types.hpp" +#include "RAJA_test-base.hpp" +#include "RAJA_unit-test-forone.hpp" + +using namespace RAJA; + +/* helper to create a RAJA View with a sliced SubLayout */ +template +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_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...)); +} + +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}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[1:3] + auto sv = TypeParam{}(view, RangeSlice<>{1,3}); + + EXPECT_EQ(sv(0), 2); + EXPECT_EQ(sv(1), 3); + EXPECT_EQ(sv(2), 4); + + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 3); + +} + +TYPED_TEST(SubViewTest, RangeStartSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[2:] + auto sv = TypeParam{}(view, RangeStartSlice<>{2}); + + EXPECT_EQ(sv(0), 3); + EXPECT_EQ(sv(1), 4); + EXPECT_EQ(sv(2), 5); + + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 3); +} + +TYPED_TEST(SubViewTest, StridedSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[0:3:2] + auto sv = make_view_with_sublayout(view, StridedSlice<>{0,3,2}); + + EXPECT_EQ(sv(0), 1); + EXPECT_EQ(sv(1), 3); + + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 2); +} + +TYPED_TEST(SubViewTest, FixedSubView1D) +{ + + Index_type a[] = {1,2,3,4,5}; + + View> view(&a[0], Layout<1>(5)); + + // sv = View[1] + auto sv = TypeParam{}(view, FixedSlice<>{1}); + + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 1); + EXPECT_EQ(sr.size_noproj(), 1); +} + +TYPED_TEST(SubViewTest, 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 = 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); + + 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); + +} + +TYPED_TEST(SubViewTest, 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 = TypeParam{}(view, RangeSlice<>{1,2}, FixedSlice<>{1}); + + EXPECT_EQ(sv(0), 5); + EXPECT_EQ(sv(1), 8); + + 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); + +} + +TYPED_TEST(SubViewTest, 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 = TypeParam{}(view, FixedSlice<>{1}, NoSlice{}); + + EXPECT_EQ(sv(0), 4); + EXPECT_EQ(sv(1), 5); + EXPECT_EQ(sv(2), 6); + + 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); +} + +TYPED_TEST(SubViewTest, 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 = TypeParam{}(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); + + 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); + +} + +TYPED_TEST(SubViewTest, 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 = TypeParam{}(view, RangeSlice<>{1,2}, StridedSlice<>{1,5,2}); + + 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); + + 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); + +} + +TYPED_TEST(SubViewTest, 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 = TypeParam{}(view, RangeSlice<>{1,2}, NoSlice{}); + + + auto& sr = TypeParam::get_subregion(sv); + EXPECT_EQ(sr.size(), 6); + + // sv2 = sv[0:1,1:2] + auto sv2 =TypeParam{}(sv, RangeSlice<>{0,1}, RangeSlice<>{1,2}); + + EXPECT_EQ(sv2(0,0), 5); + EXPECT_EQ(sv2(0,1), 6); + + EXPECT_EQ(sv2(1,0), 8); + EXPECT_EQ(sv2(1,1), 9); + + EXPECT_EQ(TypeParam::get_subregion(sv2).size(), 4); + + // sv3 = sv2[:,1] + auto sv3 =TypeParam{}(sv2, NoSlice{}, FixedSlice<>{1}); + + EXPECT_EQ(sv3(0), 6); + EXPECT_EQ(sv3(1), 9); + + EXPECT_EQ(TypeParam::get_subregion(sv3).size(), 2); + +} + +TEST(SubViewMultiViewTest, 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_multiview_with_sublayout(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); + + // sv2 = MultiView[1,1:2] + auto sv2 = make_subview_with_layout(view, FixedSlice<>{1}, RangeSlice<>{1,2}); + + // the parent layout is a MultiView + EXPECT_EQ(sv2.get_parent().get_layout().size(), 4); + EXPECT_EQ(sv2.size(), 2); + + 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() { +// #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 +// } + +// TYPED_TEST(SubViewTest, RangeFirstDimSubView2DGPU) +// { +// test_subviewGPU(); +// } \ No newline at end of file