diff --git a/include/RAJA/index/IndexValue.hpp b/include/RAJA/index/IndexValue.hpp index 6f1fc7336f..6edfea9fc8 100644 --- a/include/RAJA/index/IndexValue.hpp +++ b/include/RAJA/index/IndexValue.hpp @@ -24,6 +24,7 @@ #include +#include "RAJA/util/concepts.hpp" #include "RAJA/util/macros.hpp" #include "RAJA/util/types.hpp" @@ -58,6 +59,12 @@ struct IndexValue : public IndexValueBase RAJA_INLINE IndexValue& operator=(IndexValue const&) = default; RAJA_INLINE IndexValue& operator=(IndexValue&&) = default; + RAJA_INLINE IndexValue& operator=(const value_type& v) + { + value = v; + return *this; + } + /*! * \brief Explicit constructor. * \param v Initial value @@ -302,19 +309,100 @@ convertIndex_helper(typename FROM::IndexValueType const val) namespace type_traits { -template -struct is_instance_of_index_value : std::is_base_of, T> +template +struct is_instance_of_index_value : std::is_base_of, T> {}; -template +template constexpr bool is_instance_of_index_value_v = - is_instance_of_index_value::value; + is_instance_of_index_value::value; } // namespace type_traits namespace concepts { template -concept IndexValued = type_traits::is_instance_of_index_value_v; +concept IndexValued = + type_traits::is_instance_of_index_value_v; + +template +concept Index = concepts::Integral || concepts::IndexValued; +} // namespace concepts + +template +RAJA_HOST_DEVICE RAJA_INLINE TYPE operator+(typename TYPE::value_type lhs, + TYPE rhs) +{ + return TYPE(lhs + *rhs); +} + +template +RAJA_HOST_DEVICE RAJA_INLINE TYPE operator-(typename TYPE::value_type lhs, + TYPE rhs) +{ + return TYPE(lhs - *rhs); +} + +template +RAJA_HOST_DEVICE RAJA_INLINE TYPE operator*(typename TYPE::value_type lhs, + TYPE rhs) +{ + return TYPE(lhs * *rhs); +} + +template +RAJA_HOST_DEVICE RAJA_INLINE TYPE operator/(typename TYPE::value_type lhs, + TYPE rhs) +{ + return TYPE(lhs / *rhs); +} + +template +RAJA_HOST_DEVICE RAJA_INLINE TYPE operator%(typename TYPE::value_type lhs, + TYPE rhs) +{ + return TYPE(lhs % *rhs); +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator<(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs < *rhs; +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator<=(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs <= *rhs; +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator>(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs > *rhs; +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator>=(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs >= *rhs; +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator==(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs == *rhs; +} + +template +RAJA_HOST_DEVICE RAJA_INLINE bool operator!=(typename TYPE::value_type lhs, + TYPE rhs) +{ + return lhs != *rhs; } /*! @@ -402,12 +490,15 @@ using make_signed_t = #define RAJA_INDEX_VALUE(TYPE, NAME) \ class TYPE : public ::RAJA::IndexValue \ { \ - using parent = ::RAJA::IndexValue; \ - \ public: \ + using parent = ::RAJA::IndexValue; \ + using parent::operator=; \ + using parent::operator*; \ + using parent::operator++; \ + using parent::operator--; \ using IndexValueType = TYPE; \ RAJA_HOST_DEVICE RAJA_INLINE TYPE() : parent::IndexValue() {} \ - RAJA_HOST_DEVICE RAJA_INLINE explicit TYPE(::RAJA::Index_type v) \ + RAJA_HOST_DEVICE RAJA_INLINE TYPE(::RAJA::Index_type v) \ : parent::IndexValue(v) \ {} \ static inline std::string getName() { return NAME; } \ @@ -423,10 +514,15 @@ using make_signed_t = class TYPE : public ::RAJA::IndexValue \ { \ public: \ + using parent = RAJA::IndexValue; \ + using parent::operator=; \ + using parent::operator*; \ + using parent::operator++; \ + using parent::operator--; \ RAJA_HOST_DEVICE RAJA_INLINE TYPE() \ : RAJA::IndexValue::IndexValue() \ {} \ - RAJA_HOST_DEVICE RAJA_INLINE explicit TYPE(IDXT v) \ + RAJA_HOST_DEVICE RAJA_INLINE TYPE(IDXT v) \ : RAJA::IndexValue::IndexValue(v) \ {} \ static inline std::string getName() { return NAME; } \ diff --git a/include/RAJA/index/RangeSegment.hpp b/include/RAJA/index/RangeSegment.hpp index bb91cfcc00..73946f02cf 100644 --- a/include/RAJA/index/RangeSegment.hpp +++ b/include/RAJA/index/RangeSegment.hpp @@ -94,7 +94,7 @@ namespace RAJA * ****************************************************************************** */ -template>> struct TypedRangeSegment { @@ -105,8 +105,6 @@ struct TypedRangeSegment // static_assert(std::is_signed::value, "TypedRangeSegment DiffT requires signed type."); - static_assert(!std::is_floating_point::value, - "TypedRangeSegment Type must be non floating point."); //@{ //! @name Types used in implementation based on template parameters. @@ -139,6 +137,21 @@ struct TypedRangeSegment m_end(begin > end ? m_begin : iterator(end)) {} + //! This constructor exists whenever an IndexValued index is passed to a + //! RangeSegment So that the class may be constructed directly with values of + //! the underlying storage type + template + requires((concepts::IndexValued || concepts::IndexValued) && + std::is_convertible_v, StripStorageT> && + std::is_convertible_v, StripStorageT>) + RAJA_HOST_DEVICE constexpr TypedRangeSegment(BeginT begin, EndT end) + : m_begin(iterator(StripStorageT(stripIndexType(begin)))), + m_end(StripStorageT(stripIndexType(begin)) > + StripStorageT(stripIndexType(end)) + ? m_begin + : iterator(StripStorageT(stripIndexType(end)))) + {} + //! Disable compiler generated constructor RAJA_HOST_DEVICE TypedRangeSegment() = delete; diff --git a/include/RAJA/pattern/params/params_base.hpp b/include/RAJA/pattern/params/params_base.hpp index 20daf5c820..7cd87047eb 100644 --- a/include/RAJA/pattern/params/params_base.hpp +++ b/include/RAJA/pattern/params/params_base.hpp @@ -10,22 +10,90 @@ #ifndef RAJA_PARAMS_BASE #define RAJA_PARAMS_BASE +#include + +#include "RAJA/index/IndexValue.hpp" + namespace RAJA { namespace expt { +namespace detail +{ + +template +using valloc_index_storage_t = + std::conditional_t::value, + RAJA::strip_index_type_t, + IndexType>; + +template +RAJA_HOST_DEVICE constexpr valloc_index_storage_t +strip_valloc_index(IndexType const& index) +{ + if constexpr (std::is_base_of::value) + { + return RAJA::stripIndexType(index); + } + else + { + return index; + } +} + +template +RAJA_HOST_DEVICE constexpr IndexType +restore_valloc_index(valloc_index_storage_t const& index) +{ + if constexpr (std::is_base_of::value) + { + return IndexType(index); + } + else + { + return index; + } +} + +template +RAJA_HOST_DEVICE constexpr valloc_index_storage_t +default_valloc_index() +{ + using stored_index_type = valloc_index_storage_t; + + if constexpr (std::is_constructible::value) + { + return stored_index_type(-1); + } + else + { + return stored_index_type {}; + } +} + +} // namespace detail template struct ValLoc { - using index_type = IndexType; - using value_type = T; + using index_type = IndexType; + using stored_index_type = detail::valloc_index_storage_t; + using value_type = T; ValLoc() = default; RAJA_HOST_DEVICE constexpr explicit ValLoc(value_type v) : val(v) {} - RAJA_HOST_DEVICE constexpr ValLoc(value_type v, index_type l) : val(v), loc(l) + RAJA_HOST_DEVICE constexpr ValLoc(value_type v, index_type l) + : val(v), + loc(detail::strip_valloc_index(l)) + {} + + template::value>* = nullptr> + RAJA_HOST_DEVICE constexpr ValLoc(value_type v, stored_index_type l) + : val(v), + loc(l) {} ValLoc(ValLoc const&) = default; @@ -45,20 +113,31 @@ struct ValLoc RAJA_HOST_DEVICE constexpr const value_type& getVal() const { return val; } - RAJA_HOST_DEVICE constexpr const index_type& getLoc() const { return loc; } + RAJA_HOST_DEVICE constexpr index_type getLoc() const + { + return detail::restore_valloc_index(loc); + } + + RAJA_HOST_DEVICE constexpr const stored_index_type& getStoredLoc() const + { + return loc; + } RAJA_HOST_DEVICE void set(T inval, IndexType inindex) { val = inval; - loc = inindex; + loc = detail::strip_valloc_index(inindex); } RAJA_HOST_DEVICE void setVal(T inval) { val = inval; } - RAJA_HOST_DEVICE void setLoc(IndexType inindex) { loc = inindex; } + RAJA_HOST_DEVICE void setLoc(IndexType inindex) + { + loc = detail::strip_valloc_index(inindex); + } value_type val; - index_type loc = -1; + stored_index_type loc = detail::default_valloc_index(); }; template class Op> diff --git a/include/RAJA/pattern/params/reducer.hpp b/include/RAJA/pattern/params/reducer.hpp index 4279a6a70c..bc7e03022e 100644 --- a/include/RAJA/pattern/params/reducer.hpp +++ b/include/RAJA/pattern/params/reducer.hpp @@ -146,8 +146,9 @@ struct Reducer, ValLoc, ValLoc>, ValOp, Op>> : public ForallParamBase { using target_value_type = T; - using target_index_type = I; using value_type = ValLoc; + using target_index_type = typename value_type::index_type; + using stored_index_type = typename value_type::stored_index_type; using op = Op; using VOp = ValOp, Op>; @@ -168,7 +169,7 @@ struct Reducer, ValLoc, ValLoc>, target_index_type* index_in) : m_valop(VOp {}), target_value(data_in), - target_index(index_in) + external_target_index(index_in) {} Reducer(Reducer const&) = default; @@ -183,18 +184,26 @@ struct Reducer, ValLoc, ValLoc>, // Points to either dual value and index defined by the user, or value and // index within a ValLoc defined by the user - target_value_type* target_value = nullptr; - target_index_type* target_index = nullptr; + target_value_type* target_value = nullptr; + stored_index_type* target_index = nullptr; + target_index_type* external_target_index = nullptr; // combineTarget() performs the final op on the target data and location in // param_resolve() RAJA_HOST_DEVICE void combineTarget(value_type in) { // Create a different temp ValLoc solely for combining - value_type temp(*target_value, *target_index); + value_type temp(*target_value, getTargetIndex()); temp = op {}(temp, in); *target_value = temp.val; - *target_index = temp.loc; + if (target_index != nullptr) + { + *target_index = temp.getStoredLoc(); + } + if (external_target_index != nullptr) + { + *external_target_index = temp.getLoc(); + } } RAJA_HOST_DEVICE @@ -222,6 +231,17 @@ struct Reducer, ValLoc, ValLoc>, using ARG_LIST_T = typename ARG_TUP_T::TList; static constexpr size_t num_lambda_args = camp::tuple_size::value; + +private: + RAJA_HOST_DEVICE target_index_type getTargetIndex() const + { + if (target_index != nullptr) + { + return RAJA::expt::detail::restore_valloc_index( + *target_index); + } + return *external_target_index; + } }; } // namespace detail diff --git a/include/RAJA/util/Span.hpp b/include/RAJA/util/Span.hpp index 3119431100..aef426a15d 100644 --- a/include/RAJA/util/Span.hpp +++ b/include/RAJA/util/Span.hpp @@ -20,11 +20,7 @@ #ifndef RAJA_SPAN_HPP #define RAJA_SPAN_HPP -#include -#include - #include "RAJA/index/IndexValue.hpp" -#include "RAJA/util/concepts.hpp" #include "RAJA/util/macros.hpp" #include "RAJA/util/types.hpp" #include "camp/concepts.hpp" @@ -61,19 +57,13 @@ namespace RAJA * compile time extents * */ - -namespace concepts -{ -template -concept SpanIndex = concepts::Integral || concepts::IndexValued; -}; - -template +template struct Span { using element_type = typename std::iterator_traits::value_type; using value_type = camp::decay; using size_type = IndexType; + using offset_type = RAJA::strip_index_type_t; using difference_type = std::ptrdiff_t; using reference = element_type&; using const_reference = const element_type&; @@ -87,7 +77,7 @@ struct Span constexpr RAJA_HOST_DEVICE Span(iterator begin, size_type size) : m_begin {begin}, - m_end {begin + size} + m_end {begin + RAJA::stripIndexType(size)} {} constexpr RAJA_HOST_DEVICE RAJA_INLINE iterator begin() { return m_begin; } @@ -160,7 +150,7 @@ struct Span constexpr RAJA_HOST_DEVICE RAJA_INLINE reference operator[](size_type i) const { - return data()[i]; + return data()[RAJA::stripIndexType(i)]; } constexpr RAJA_HOST_DEVICE RAJA_INLINE iterator data() const @@ -197,8 +187,11 @@ struct Span constexpr RAJA_HOST_DEVICE RAJA_INLINE Span slice(size_type begin, size_type length) const { - auto start = m_begin + begin; - auto end = start + length > m_end ? m_end : start + length; + offset_type stripped_begin = RAJA::stripIndexType(begin); + offset_type stripped_length = RAJA::stripIndexType(length); + auto start = m_begin + stripped_begin; + auto end = start + stripped_length > m_end ? m_end + : start + stripped_length; return Span(start, end); } diff --git a/test/functional/forall/CombiningAdapter/test-forall-CombiningAdapter.cpp.in b/test/functional/forall/CombiningAdapter/test-forall-CombiningAdapter.cpp.in index 70fcefd657..4b9986a91e 100644 --- a/test/functional/forall/CombiningAdapter/test-forall-CombiningAdapter.cpp.in +++ b/test/functional/forall/CombiningAdapter/test-forall-CombiningAdapter.cpp.in @@ -30,7 +30,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@ForallCombiningAdapter@DIMENSION@Types = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/forall/indexset-view/test-forall-indexset-view.cpp.in b/test/functional/forall/indexset-view/test-forall-indexset-view.cpp.in index 15f464b88e..1adf50d694 100644 --- a/test/functional/forall/indexset-view/test-forall-indexset-view.cpp.in +++ b/test/functional/forall/indexset-view/test-forall-indexset-view.cpp.in @@ -32,7 +32,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@ForallIndexSetViewTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/forall/indexset/test-forall-indexset.cpp.in b/test/functional/forall/indexset/test-forall-indexset.cpp.in index 24748ee380..7589001204 100644 --- a/test/functional/forall/indexset/test-forall-indexset.cpp.in +++ b/test/functional/forall/indexset/test-forall-indexset.cpp.in @@ -32,7 +32,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@ForallIndexSetTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/forall/region/test-forall-region.cpp.in b/test/functional/forall/region/test-forall-region.cpp.in index 0a255b9851..b5d37562bb 100644 --- a/test/functional/forall/region/test-forall-region.cpp.in +++ b/test/functional/forall/region/test-forall-region.cpp.in @@ -47,7 +47,7 @@ using OpenMPForallRegionExecPols = // Cartesian product of types used in parameterized tests // using @REGION_BACKEND@ForallRegionTypes = - Test< camp::cartesian_product>::Types; @@ -60,4 +60,3 @@ INSTANTIATE_TYPED_TEST_SUITE_P(@REGION_BACKEND@, @REGION_BACKEND@ForallRegionTypes); - diff --git a/test/functional/forall/resource-indexset/test-forall-resource-indexset.cpp.in b/test/functional/forall/resource-indexset/test-forall-resource-indexset.cpp.in index 0803f9e663..7445ebff39 100644 --- a/test/functional/forall/resource-indexset/test-forall-resource-indexset.cpp.in +++ b/test/functional/forall/resource-indexset/test-forall-resource-indexset.cpp.in @@ -32,7 +32,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@ForallResourceIndexSetTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/forall/segment-view/test-forall-segment-view.cpp.in b/test/functional/forall/segment-view/test-forall-segment-view.cpp.in index 7447087c71..49f34113b5 100644 --- a/test/functional/forall/segment-view/test-forall-segment-view.cpp.in +++ b/test/functional/forall/segment-view/test-forall-segment-view.cpp.in @@ -30,7 +30,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@ForallSegmentViewTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/hyperplane/test-kernel-hyperplane-2D.cpp.in b/test/functional/kernel/hyperplane/test-kernel-hyperplane-2D.cpp.in index 3bafd5f7fd..9fcafe29fb 100644 --- a/test/functional/kernel/hyperplane/test-kernel-hyperplane-2D.cpp.in +++ b/test/functional/kernel/hyperplane/test-kernel-hyperplane-2D.cpp.in @@ -179,7 +179,7 @@ using HipKernelHyperplaneExecPols = // Cartesian product of types used in parameterized tests // using @BACKEND@KernelHyperplane = - Test< camp::cartesian_product #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template std::enable_if_t CallKernel(DATA_TYPE& trip_count, DATA_TYPE& oob_count, - RAJA::View>& WorkView, - const int idim, - const int jdim, - const int groups) + RAJA::View>>& WorkView, + const INDEX_TYPE idim, + const INDEX_TYPE jdim, + const INDEX_TYPE groups) { RAJA::TypedRangeSegment Grange( 0, groups ); RAJA::TypedRangeSegment Irange( 0, idim ); @@ -34,7 +56,7 @@ CallKernel(DATA_TYPE& trip_count, [=] RAJA_HOST_DEVICE (INDEX_TYPE g, INDEX_TYPE ii, INDEX_TYPE jj, RAJA::expt::ValOp& _trip_count, RAJA::expt::ValOp& _oob_count ) { - if ((int)g < 0 || (int)g >= groups || (int)ii < 0 || (int)ii >= idim || (int)jj < 0 || (int)jj >= jdim) { + if (get_val(g) >= get_val(groups) || get_val(ii) >= get_val(idim) || get_val(jj) >= get_val(jdim)) { _oob_count += 1; } @@ -59,10 +81,10 @@ template CallKernel(DATA_TYPE& _trip_count, DATA_TYPE& _oob_count, - RAJA::View>& WorkView, - const int idim, - const int jdim, - const int groups) + RAJA::View>>& WorkView, + const INDEX_TYPE idim, + const INDEX_TYPE jdim, + const INDEX_TYPE groups) { RAJA::TypedRangeSegment Grange( 0, groups ); RAJA::TypedRangeSegment Irange( 0, idim ); @@ -73,7 +95,7 @@ CallKernel(DATA_TYPE& _trip_count, RAJA::kernel ( RAJA::make_tuple( Grange, Irange, Jrange ), [=] RAJA_HOST_DEVICE ( INDEX_TYPE g, INDEX_TYPE ii, INDEX_TYPE jj ) { - if ((int)g < 0 || (int)g >= groups || (int)ii < 0 || (int)ii >= idim || (int)jj < 0 || (int)jj >= jdim) { + if (get_val(g) >= get_val(groups) || get_val(ii) >= get_val(idim) || get_val(jj) >= get_val(jdim)) { oob_count += 1; } @@ -97,7 +119,7 @@ CallKernel(DATA_TYPE& _trip_count, template -void KernelHyperplane2DTestImpl(const int groups, const int idim, const int jdim) +void KernelHyperplane2DTestImpl(const INDEX_TYPE groups, const INDEX_TYPE idim, const INDEX_TYPE jdim) { // This test traverses "groups" 2D arrays, and modifies values in a 1D hyperplane manner. @@ -116,14 +138,16 @@ void KernelHyperplane2DTestImpl(const int groups, const int idim, const int jdim &test_array ); - RAJA::View> HostView( test_array, groups, idim, jdim ); - RAJA::View> WorkView( work_array, groups, idim, jdim ); - RAJA::View> CheckView( check_array, groups, idim, jdim ); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType HostView( test_array, groups, idim, jdim ); + ViewType WorkView( work_array, groups, idim, jdim ); + ViewType CheckView( check_array, groups, idim, jdim ); // initialize array - std::iota( test_array, test_array + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); DATA_TYPE trip_count(0); DATA_TYPE oob_count(0); @@ -131,15 +155,15 @@ void KernelHyperplane2DTestImpl(const int groups, const int idim, const int jdim // perform array arithmetic with a 1D hyperplane, in either the I or J direction CallKernel(trip_count, oob_count, WorkView, idim, jdim, groups); - ASSERT_EQ((INDEX_TYPE)trip_count, (INDEX_TYPE)groups * idim * jdim); + ASSERT_EQ((INDEX_TYPE)trip_count, groups * idim * jdim); ASSERT_EQ((INDEX_TYPE)oob_count, (INDEX_TYPE)0); - work_res.memcpy( check_array, work_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array, work_array, sizeof(DATA_TYPE) * get_val(array_length) ); // perform array arithmetic on the CPU - for (int g = 0; g < groups; ++g) { - for (int i = 0; i < idim; ++i) { - for (int j = 0; j < jdim; ++j) { + for (INDEX_TYPE g(0); g < groups; ++g) { + for (INDEX_TYPE i(0); i < idim; ++i) { + for (INDEX_TYPE j(0); j < jdim; ++j) { DATA_TYPE left = 1; if (i > 0) { left = HostView(g, i - 1, j); @@ -155,9 +179,9 @@ void KernelHyperplane2DTestImpl(const int groups, const int idim, const int jdim } } - for (int g = 0; g < groups; ++g) { - for (int i = 0; i < idim; ++i) { - for (int j = 0; j < jdim; ++j) { + for (INDEX_TYPE g(0); g < groups; ++g) { + for (INDEX_TYPE i(0); i < idim; ++i) { + for (INDEX_TYPE j(0); j < jdim; ++j) { ASSERT_FLOAT_EQ(CheckView(g, i, j), HostView(g, i, j)); } } diff --git a/test/functional/kernel/hyperplane/tests/test-kernel-hyperplane-3D.hpp b/test/functional/kernel/hyperplane/tests/test-kernel-hyperplane-3D.hpp index 01ed31ebef..e6597e7a5a 100644 --- a/test/functional/kernel/hyperplane/tests/test-kernel-hyperplane-3D.hpp +++ b/test/functional/kernel/hyperplane/tests/test-kernel-hyperplane-3D.hpp @@ -13,22 +13,44 @@ #include #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template std::enable_if_t CallKernel(DATA_TYPE& trip_count, DATA_TYPE& oob_count, - RAJA::View>& WorkView, - const int idim, - const int jdim, - const int kdim, - const int groups) + RAJA::View>>& WorkView, + const INDEX_TYPE idim, + const INDEX_TYPE jdim, + const INDEX_TYPE kdim, + const INDEX_TYPE groups) { // perform array arithmetic with a 2D J-K hyperplane - RAJA::TypedRangeSegment Grange( 0, groups ); - RAJA::TypedRangeStrideSegment Irange( 0, idim, 1 ); - RAJA::TypedRangeStrideSegment Jrange( jdim-1, -1, -1 ); - RAJA::TypedRangeStrideSegment Krange( 0, kdim, 1 ); + RAJA::TypedRangeSegment Grange( 0, get_val(groups) ); + RAJA::TypedRangeStrideSegment Irange( 0, get_val(idim), 1 ); + RAJA::TypedRangeStrideSegment Jrange( get_val(jdim) - 1, -1, -1 ); + RAJA::TypedRangeStrideSegment Krange( 0, get_val(kdim), 1 ); RAJA::kernel_param ( RAJA::make_tuple( Grange, Irange, Jrange, Krange ), @@ -39,7 +61,8 @@ CallKernel(DATA_TYPE& trip_count, [=] RAJA_HOST_DEVICE ( INDEX_TYPE g, INDEX_TYPE ii, INDEX_TYPE jj, INDEX_TYPE kk, RAJA::expt::ValOp& _trip_count, RAJA::expt::ValOp& _oob_count ) { - if (g < 0 || g >= groups || ii < 0 || ii >= idim || jj < 0 || jj >= jdim || kk < 0 || kk >= kdim) { + if (get_val(g) >= get_val(groups) || get_val(ii) >= get_val(idim) || + get_val(jj) >= get_val(jdim) || get_val(kk) >= get_val(kdim)) { _oob_count += 1; } @@ -68,23 +91,24 @@ template CallKernel(DATA_TYPE& _trip_count, DATA_TYPE& _oob_count, - RAJA::View>& WorkView, - const int idim, - const int jdim, - const int kdim, - const int groups) + RAJA::View>>& WorkView, + const INDEX_TYPE idim, + const INDEX_TYPE jdim, + const INDEX_TYPE kdim, + const INDEX_TYPE groups) { RAJA::ReduceSum trip_count (_trip_count); RAJA::ReduceSum oob_count (_oob_count); // perform array arithmetic with a 2D J-K hyperplane - RAJA::TypedRangeSegment Grange( 0, groups ); - RAJA::TypedRangeStrideSegment Irange( 0, idim, 1 ); - RAJA::TypedRangeStrideSegment Jrange( jdim-1, -1, -1 ); - RAJA::TypedRangeStrideSegment Krange( 0, kdim, 1 ); + RAJA::TypedRangeSegment Grange( 0, get_val(groups) ); + RAJA::TypedRangeStrideSegment Irange( 0, get_val(idim), 1 ); + RAJA::TypedRangeStrideSegment Jrange( get_val(jdim - 1), -1, -1 ); + RAJA::TypedRangeStrideSegment Krange( 0, get_val(kdim), 1 ); RAJA::kernel ( RAJA::make_tuple( Grange, Irange, Jrange, Krange ), [=] RAJA_HOST_DEVICE ( INDEX_TYPE g, INDEX_TYPE ii, INDEX_TYPE jj, INDEX_TYPE kk ) { - if (g < 0 || g >= groups || ii < 0 || ii >= idim || jj < 0 || jj >= jdim || kk < 0 || kk >= kdim) { + if (get_val(g) >= get_val(groups) || get_val(ii) >= get_val(idim) || + get_val(jj) >= get_val(jdim) || get_val(kk) >= get_val(kdim)) { oob_count += 1; } @@ -113,18 +137,19 @@ CallKernel(DATA_TYPE& _trip_count, template std::enable_if_t>::value> -KernelHyperplane3DTestImpl(const int RAJA_UNUSED_ARG(groups), const int RAJA_UNUSED_ARG(idim), const int RAJA_UNUSED_ARG(jdim), const int RAJA_UNUSED_ARG(kdim)) +KernelHyperplane3DTestImpl(const INDEX_TYPE RAJA_UNUSED_ARG(groups), const INDEX_TYPE RAJA_UNUSED_ARG(idim), const INDEX_TYPE RAJA_UNUSED_ARG(jdim), const INDEX_TYPE RAJA_UNUSED_ARG(kdim)) { // do nothing for unsigned index types } template std::enable_if_t>::value> -KernelHyperplane3DTestImpl(const int groups, const int idimin, const int jdimin, const int kdimin) +KernelHyperplane3DTestImpl(const INDEX_TYPE groups, const INDEX_TYPE idimin, const INDEX_TYPE jdimin, const INDEX_TYPE kdimin) { // This test traverses "groups" number of 3D arrays, and modifies values in a 2D hyperplane manner. - int idim, jdim, kdim; + using raw_index_type = VAL_T; + raw_index_type idim, jdim, kdim; if ( std::is_same::value ) { // Restrict to a small data size for better float precision. @@ -134,18 +159,22 @@ KernelHyperplane3DTestImpl(const int groups, const int idimin, const int jdimin, } else { - idim = idimin; - jdim = jdimin; - kdim = kdimin; + idim = get_val(idimin); + jdim = get_val(jdimin); + kdim = get_val(kdimin); } + INDEX_TYPE idim_t(idim); + INDEX_TYPE jdim_t(jdim); + INDEX_TYPE kdim_t(kdim); + camp::resources::Resource work_res{WORKING_RES::get_default()}; DATA_TYPE * work_array; DATA_TYPE * check_array; DATA_TYPE * test_array; - INDEX_TYPE array_length = groups * idim * jdim * kdim; + INDEX_TYPE array_length = groups * idim_t * jdim_t * kdim_t; allocateForallTestData ( array_length, work_res, @@ -154,30 +183,32 @@ KernelHyperplane3DTestImpl(const int groups, const int idimin, const int jdimin, &test_array ); - RAJA::View> HostView( test_array, groups, idim, jdim, kdim ); - RAJA::View> WorkView( work_array, groups, idim, jdim, kdim ); - RAJA::View> CheckView( check_array, groups, idim, jdim, kdim ); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType HostView( test_array, groups, idim_t, jdim_t, kdim_t ); + ViewType WorkView( work_array, groups, idim_t, jdim_t, kdim_t ); + ViewType CheckView( check_array, groups, idim_t, jdim_t, kdim_t ); // initialize array - std::iota( test_array, test_array + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); DATA_TYPE trip_count(0); DATA_TYPE oob_count(0); - CallKernel(trip_count, oob_count, WorkView, idim, jdim, kdim, groups); + CallKernel(trip_count, oob_count, WorkView, idim_t, jdim_t, kdim_t, groups); - ASSERT_EQ((INDEX_TYPE)trip_count, (INDEX_TYPE)groups * idim * jdim * kdim); + ASSERT_EQ((INDEX_TYPE)trip_count, groups * idim_t * jdim_t * kdim_t); ASSERT_EQ((INDEX_TYPE)oob_count, (INDEX_TYPE)0); - work_res.memcpy( check_array, work_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array, work_array, sizeof(DATA_TYPE) * get_val(array_length) ); // perform array arithmetic on the CPU - for (int g = 0; g < groups; ++g) { - for (int i = 0; i < idim; ++i) { - for (int j = jdim - 1; j >= 0; --j) { - for (int k = 0; k < kdim; ++k) { + for (INDEX_TYPE g(0); g < groups; ++g) { + for (INDEX_TYPE i(0); i < idim_t; ++i) { + for (INDEX_TYPE j(jdim_t - 1); j >= INDEX_TYPE(0); --j) { + for (INDEX_TYPE k(0); k < kdim_t; ++k) { DATA_TYPE left = 1; if (i > 0) { left = HostView(g, i - 1, j, k); @@ -199,10 +230,10 @@ KernelHyperplane3DTestImpl(const int groups, const int idimin, const int jdimin, } } - for (int g = 0; g < groups; ++g) { - for (int i = 0; i < idim; ++i) { - for (int j = 0; j < jdim; ++j) { - for (int k = 0; k < kdim; ++k) { + for (INDEX_TYPE g(0); g < groups; ++g) { + for (INDEX_TYPE i(0); i < idim_t; ++i) { + for (INDEX_TYPE j(0); j < jdim_t; ++j) { + for (INDEX_TYPE k(0); k < kdim_t; ++k) { ASSERT_FLOAT_EQ(CheckView(g, i, j, k), HostView(g, i, j, k)); } } diff --git a/test/functional/kernel/nested-loop-segment-types/test-kernel-nested-loop-segments.cpp.in b/test/functional/kernel/nested-loop-segment-types/test-kernel-nested-loop-segments.cpp.in index 827c4ca7d2..73d504cc58 100644 --- a/test/functional/kernel/nested-loop-segment-types/test-kernel-nested-loop-segments.cpp.in +++ b/test/functional/kernel/nested-loop-segment-types/test-kernel-nested-loop-segments.cpp.in @@ -27,7 +27,7 @@ // Sequential execution policy types using SequentialKernelExecPols = camp::list< - + RAJA::KernelPolicy< RAJA::statement::For<0, RAJA::seq_exec, RAJA::statement::For<1, RAJA::seq_exec, @@ -42,7 +42,7 @@ using SequentialKernelExecPols = camp::list< #if defined(RAJA_ENABLE_OPENMP) using OpenMPKernelExecPols = camp::list< - + RAJA::KernelPolicy< RAJA::statement::For<0, RAJA::omp_parallel_for_exec, RAJA::statement::For<1, RAJA::seq_exec, @@ -174,7 +174,7 @@ using SyclKernelExecPols = camp::list< // Cartesian product of types used in parameterized tests // using @BACKEND@KernelNestedLoopsSegmentTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/nested-loop-segment-types/tests/test-kernel-nested-loops-segment-types.hpp b/test/functional/kernel/nested-loop-segment-types/tests/test-kernel-nested-loops-segment-types.hpp index bb9e9065a8..a0259e91cf 100644 --- a/test/functional/kernel/nested-loop-segment-types/tests/test-kernel-nested-loops-segment-types.hpp +++ b/test/functional/kernel/nested-loop-segment-types/tests/test-kernel-nested-loops-segment-types.hpp @@ -18,6 +18,27 @@ #include #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} template void KernelNestedLoopsSegmentTypesTestImpl( @@ -61,24 +82,25 @@ void KernelNestedLoopsSegmentTypesTestImpl( &check_array, &test_array); - RAJA::View< DATA_TYPE, RAJA::Layout<3> > work_view(work_array, - dim1, dim2, dim3); - RAJA::View< DATA_TYPE, RAJA::Layout<3> > test_view(test_array, - dim1, dim2, dim3); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + + ViewType work_view(work_array, dim1, dim2, dim3); + ViewType test_view(test_array, dim1, dim2, dim3); memset( static_cast(test_array), 0, - sizeof(DATA_TYPE) * RAJA::stripIndexType(data_len) ); + sizeof(DATA_TYPE) * get_val(data_len) ); working_res.memcpy(work_array, test_array, - sizeof(DATA_TYPE) * RAJA::stripIndexType(data_len)); + sizeof(DATA_TYPE) * get_val(data_len)); if ( !zero_legth_segment ) { for (IDX_TYPE i1 = 0; i1 < idx1_len; ++i1) { for (IDX_TYPE i2 = 0; i2 < idx2_len; ++i2) { for (IDX_TYPE i3 = 0; i3 < idx3_len; ++i3) { - auto ii1 = RAJA::stripIndexType(i1); - auto ii2 = RAJA::stripIndexType(i2); - auto ii3 = RAJA::stripIndexType(i3); + auto ii1 = get_val(i1); + auto ii2 = get_val(i2); + auto ii3 = get_val(i3); test_view( s1_idx[ii1], s2_idx[ii2], s3_idx[ii3] ) = static_cast( RAJA::stripIndexType( s1_idx[ii1] + s2_idx[ii2] + s3_idx[ii3]) ); @@ -115,7 +137,7 @@ void KernelNestedLoopsSegmentTypesTestImpl( } working_res.memcpy(check_array, work_array, - sizeof(DATA_TYPE) * RAJA::stripIndexType(data_len)); + sizeof(DATA_TYPE) * get_val(data_len)); for (IDX_TYPE i = 0; i < data_len; ++i) { auto ii = RAJA::stripIndexType(i); @@ -140,6 +162,7 @@ TYPED_TEST_P(KernelNestedLoopsSegmentTypesTest, NestedLoopsSegmentTypesKernel) using IDX_TYPE = typename camp::at>::type; using WORKING_RES = typename camp::at>::type; using EXEC_POLICY = typename camp::at>::type; + using raw_idx_type = VAL_T; camp::resources::Resource working_res{WORKING_RES::get_default()}; @@ -153,7 +176,7 @@ TYPED_TEST_P(KernelNestedLoopsSegmentTypesTest, NestedLoopsSegmentTypesKernel) RAJA::TypedRangeSegment s1( 0, 69 ); RAJA::getIndices(s1_idx, s1); - RAJA::TypedRangeStrideSegment s2( 3, 188, 2 ); + RAJA::TypedRangeStrideSegment s2( raw_idx_type(3), raw_idx_type(188), raw_idx_type(2) ); RAJA::getIndices(s2_idx, s2); IDX_TYPE last = IDX_TYPE(427); @@ -223,7 +246,7 @@ TYPED_TEST_P(KernelNestedLoopsSegmentTypesTest, NestedLoopsSegmentTypesKernel) perm); // Zero-length range stride segment - RAJA::TypedRangeStrideSegment s5( 3, 3, 2 ); + RAJA::TypedRangeStrideSegment s5( raw_idx_type(3), raw_idx_type(3), raw_idx_type(2) ); std::vector s5_idx; RAJA::getIndices(s5_idx, s5); diff --git a/test/functional/kernel/nested-loop-view-types/test-kernel-nested-loop-view.cpp.in b/test/functional/kernel/nested-loop-view-types/test-kernel-nested-loop-view.cpp.in index c7d9b318b8..94acc89ca6 100644 --- a/test/functional/kernel/nested-loop-view-types/test-kernel-nested-loop-view.cpp.in +++ b/test/functional/kernel/nested-loop-view-types/test-kernel-nested-loop-view.cpp.in @@ -309,7 +309,7 @@ camp::list< // Cartesian product of types used in parameterized tests // using @BACKEND@KernelNestesLoop@TESTTYPE@@DIM@Types = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView2D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView2D.hpp index 331a1cff1a..0ebe38a689 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView2D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView2D.hpp @@ -10,6 +10,20 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_OFFSETVIEW2D_HPP__ #define __TEST_KERNEL_NESTEDLOOP_OFFSETVIEW2D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; +} template void KernelOffsetView2DTestImpl(std::array dim, @@ -45,14 +59,20 @@ void KernelOffsetView2DTestImpl(std::array dim, } - RAJA::OffsetLayout<2> layout = - RAJA::make_offset_layout<2>( {{offset_lo.at(0), offset_lo.at(1)}}, - {{offset_lo.at(0) + dim.at(0), - offset_lo.at(1) + dim.at(1)}} ); - RAJA::View< IDX_TYPE, RAJA::OffsetLayout<2> > view(working_array, layout); + using raw_idx_type = VAL_T; + using LayoutType = + RAJA::TypedOffsetLayout>; + + LayoutType layout({{raw_idx_type(offset_lo.at(0)), + raw_idx_type(offset_lo.at(1))}}, + {{raw_idx_type(offset_lo.at(0) + dim.at(0)), + raw_idx_type(offset_lo.at(1) + dim.at(1))}}); + RAJA::View view(working_array, layout); - RAJA::TypedRangeSegment iseg( offset_lo.at(0), offset_hi.at(0)); - RAJA::TypedRangeSegment jseg( offset_lo.at(1), offset_hi.at(1)); + RAJA::TypedRangeSegment iseg( IDX_TYPE(offset_lo.at(0)), + IDX_TYPE(offset_hi.at(0))); + RAJA::TypedRangeSegment jseg( IDX_TYPE(offset_lo.at(1)), + IDX_TYPE(offset_hi.at(1))); RAJA::kernel( RAJA::make_tuple( iseg, jseg ), diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView3D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView3D.hpp index 4161cf81a5..601eed6b99 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView3D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-OffsetView3D.hpp @@ -10,6 +10,20 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_OFFSETVIEW3D_HPP__ #define __TEST_KERNEL_NESTEDLOOP_OFFSETVIEW3D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; +} template void KernelOffsetView3DTestImpl(std::array dim, @@ -50,19 +64,26 @@ void KernelOffsetView3DTestImpl(std::array dim, } - RAJA::OffsetLayout<3> layout = - RAJA::make_offset_layout<3>( {{offset_lo.at(0), - offset_lo.at(1), - offset_lo.at(2)}}, - {{offset_lo.at(0) + dim.at(0), - offset_lo.at(1) + dim.at(1), - offset_lo.at(2) + dim.at(2)}} ); + using raw_idx_type = VAL_T; + using LayoutType = + RAJA::TypedOffsetLayout>; + + LayoutType layout({{raw_idx_type(offset_lo.at(0)), + raw_idx_type(offset_lo.at(1)), + raw_idx_type(offset_lo.at(2))}}, + {{raw_idx_type(offset_lo.at(0) + dim.at(0)), + raw_idx_type(offset_lo.at(1) + dim.at(1)), + raw_idx_type(offset_lo.at(2) + dim.at(2))}}); - RAJA::View< IDX_TYPE, RAJA::OffsetLayout<3> > view(working_array, layout); + RAJA::View view(working_array, layout); - RAJA::TypedRangeSegment iseg( offset_lo.at(0), offset_hi.at(0)); - RAJA::TypedRangeSegment jseg( offset_lo.at(1), offset_hi.at(1)); - RAJA::TypedRangeSegment kseg( offset_lo.at(2), offset_hi.at(2)); + RAJA::TypedRangeSegment iseg( IDX_TYPE(offset_lo.at(0)), + IDX_TYPE(offset_hi.at(0))); + RAJA::TypedRangeSegment jseg( IDX_TYPE(offset_lo.at(1)), + IDX_TYPE(offset_hi.at(1))); + RAJA::TypedRangeSegment kseg( IDX_TYPE(offset_lo.at(2)), + IDX_TYPE(offset_hi.at(2))); RAJA::kernel( RAJA::make_tuple( iseg, jseg, kseg ), diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView2D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView2D.hpp index 8b5de9aaf1..966af4aee6 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView2D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView2D.hpp @@ -10,10 +10,33 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_PERMUTEDOFFSETVIEW2D_HPP__ #define __TEST_KERNEL_NESTEDLOOP_PERMUTEDOFFSETVIEW2D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template void KernelPermutedOffsetView2DTestImpl(std::array dim, std::array perm) { + using raw_idx_type = VAL_T; camp::resources::Resource working_res{WORKING_RES::get_default()}; IDX_TYPE* A_work_array; IDX_TYPE* A_check_array; @@ -78,8 +101,8 @@ void KernelPermutedOffsetView2DTestImpl(std::array dim, for (RAJA::idx_t i = 0; i < Nint_outer; ++i) { for (RAJA::idx_t j = 0; j < Nint_inner; ++j) { - int A_idx = j + Nint_inner * i; - int B_idx = (j + 1) + Ntot_inner * (i + 1); + RAJA::idx_t A_idx = j + Nint_inner * i; + RAJA::idx_t B_idx = (j + 1) + Ntot_inner * (i + 1); A_test_array[A_idx] = B_test_array[B_idx] + // C B_test_array[B_idx - Ntot_inner] + // S @@ -98,20 +121,21 @@ void KernelPermutedOffsetView2DTestImpl(std::array dim, RAJA::Layout<2> A_layout = RAJA::make_permuted_layout( {{Nint_len.at(0), Nint_len.at(1)}}, perm ); - RAJA::View< IDX_TYPE, RAJA::OffsetLayout<2> > B_view(B_work_array, B_layout); - RAJA::View< IDX_TYPE, RAJA::Layout<2> > A_view(A_work_array, A_layout); - - RAJA::TypedRangeSegment iseg( 0, Nint_len.at(0) ); - RAJA::TypedRangeSegment jseg( 0, Nint_len.at(1) ); + RAJA::TypedRangeSegment iseg( 0, IDX_TYPE(Nint_len.at(0)) ); + RAJA::TypedRangeSegment jseg( 0, IDX_TYPE(Nint_len.at(1)) ); RAJA::kernel( RAJA::make_tuple( iseg, jseg ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE j) { - - A_view(i, j) = B_view(i, j) + - B_view(i - 1, j) + B_view(i + 1, j) + - B_view(i, j - 1) + B_view(i, j + 1); - + auto ii = raw_idx_type(get_val(i)); + auto jj = raw_idx_type(get_val(j)); + auto a_idx = RAJA::stripIndexType(A_layout(ii, jj)); + A_work_array[a_idx] = + B_work_array[RAJA::stripIndexType(B_layout(ii, jj))] + + B_work_array[RAJA::stripIndexType(B_layout(ii - raw_idx_type(1), jj))] + + B_work_array[RAJA::stripIndexType(B_layout(ii + raw_idx_type(1), jj))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj - raw_idx_type(1)))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj + raw_idx_type(1)))]; } ); diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView3D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView3D.hpp index e8db6957e3..f086e5900a 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView3D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedOffsetView3D.hpp @@ -10,10 +10,33 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_PERMUTEDOFFSETVIEW3D_HPP__ #define __TEST_KERNEL_NESTEDLOOP_PERMUTEDOFFSETVIEW3D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template void KernelPermutedOffsetView3DTestImpl(std::array dim, std::array perm) { + using raw_idx_type = VAL_T; camp::resources::Resource working_res{WORKING_RES::get_default()}; IDX_TYPE* A_work_array; IDX_TYPE* A_check_array; @@ -86,8 +109,8 @@ void KernelPermutedOffsetView3DTestImpl(std::array dim, for (RAJA::idx_t j = 0; j < Nint_middle; ++j) { for (RAJA::idx_t k = 0; k < Nint_inner; ++k) { - int A_idx = k + j * Nint_inner + i * Nint_inner * Nint_middle; - int B_idx = + RAJA::idx_t A_idx = k + j * Nint_inner + i * Nint_inner * Nint_middle; + RAJA::idx_t B_idx = (k + 1) + (j + 1) * Ntot_inner + (i + 1) * Ntot_inner * Ntot_middle; A_test_array[A_idx] = @@ -115,20 +138,25 @@ void KernelPermutedOffsetView3DTestImpl(std::array dim, Nint_len.at(1), Nint_len.at(2)}}, perm ); - RAJA::View< IDX_TYPE, RAJA::OffsetLayout<3> > B_view(B_work_array, B_layout); - RAJA::View< IDX_TYPE, RAJA::Layout<3> > A_view(A_work_array, A_layout); - - RAJA::TypedRangeSegment iseg( 0, Nint_len.at(0) ); - RAJA::TypedRangeSegment jseg( 0, Nint_len.at(1) ); - RAJA::TypedRangeSegment kseg( 0, Nint_len.at(2) ); + RAJA::TypedRangeSegment iseg( 0, IDX_TYPE(Nint_len.at(0)) ); + RAJA::TypedRangeSegment jseg( 0, IDX_TYPE(Nint_len.at(1)) ); + RAJA::TypedRangeSegment kseg( 0, IDX_TYPE(Nint_len.at(2)) ); RAJA::kernel( RAJA::make_tuple( iseg, jseg, kseg ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE j, IDX_TYPE k) { - A_view(i, j, k) = B_view(i, j, k) + - B_view(i - 1, j, k) + B_view(i + 1, j, k) + - B_view(i, j - 1, k) + B_view(i, j + 1, k) + - B_view(i, j, k - 1) + B_view(i, j, k + 1); + auto ii = raw_idx_type(get_val(i)); + auto jj = raw_idx_type(get_val(j)); + auto kk = raw_idx_type(get_val(k)); + auto a_idx = RAJA::stripIndexType(A_layout(ii, jj, kk)); + A_work_array[a_idx] = + B_work_array[RAJA::stripIndexType(B_layout(ii, jj, kk))] + + B_work_array[RAJA::stripIndexType(B_layout(ii - raw_idx_type(1), jj, kk))] + + B_work_array[RAJA::stripIndexType(B_layout(ii + raw_idx_type(1), jj, kk))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj - raw_idx_type(1), kk))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj + raw_idx_type(1), kk))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj, kk - raw_idx_type(1)))] + + B_work_array[RAJA::stripIndexType(B_layout(ii, jj, kk + raw_idx_type(1)))]; } ); diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView2D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView2D.hpp index c97663c2db..82bc205cb0 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView2D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView2D.hpp @@ -10,6 +10,27 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_PERMUTEDVIEW2D_HPP__ #define __TEST_KERNEL_NESTEDLOOP_PERMUTEDVIEW2D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} template void KernelPermutedView2DTestImpl(std::array dim, @@ -35,20 +56,20 @@ void KernelPermutedView2DTestImpl(std::array dim, working_res.memcpy(working_array, test_array, sizeof(IDX_TYPE) * N); - int mod_val = dim.at( perm.at(1) ); + using raw_idx_type = VAL_T; + raw_idx_type mod_val = get_val(dim.at( perm.at(1) )); for (RAJA::idx_t ii = 0; ii < N; ++ii) { test_array[ii] = static_cast(ii % mod_val); } RAJA::Layout<2> layout = RAJA::make_permuted_layout(dim_strip, perm); - RAJA::View< IDX_TYPE, RAJA::Layout<2, int> > view(working_array, layout); RAJA::kernel( - RAJA::make_tuple( RAJA::TypedRangeSegment(0, dim_strip.at(0)), - RAJA::TypedRangeSegment(0, dim_strip.at(1)) ), + RAJA::make_tuple( RAJA::TypedRangeSegment(0, IDX_TYPE(dim_strip.at(0))), + RAJA::TypedRangeSegment(0, IDX_TYPE(dim_strip.at(1))) ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE j) { - int val = RAJA::stripIndexType(layout(i, j)) % mod_val; - view(i, j) = static_cast(val); + auto linear = RAJA::stripIndexType(layout(get_val(i), get_val(j))); + working_array[linear] = static_cast(linear % mod_val); } ); diff --git a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView3D.hpp b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView3D.hpp index 0a6a99a7cc..043cfb39b3 100644 --- a/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView3D.hpp +++ b/test/functional/kernel/nested-loop-view-types/tests/test-kernel-nested-loop-PermutedView3D.hpp @@ -10,6 +10,27 @@ #ifndef __TEST_KERNEL_NESTEDLOOP_PERMUTEDVIEW3D_HPP_ #define __TEST_KERNEL_NESTEDLOOP_PERMUTEDVIEW3D_HPP_ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} template void KernelPermutedView3DTestImpl(std::array dim, @@ -36,21 +57,21 @@ void KernelPermutedView3DTestImpl(std::array dim, working_res.memcpy(working_array, test_array, sizeof(IDX_TYPE) * N); - int mod_val = dim.at( perm.at(1) ) * dim.at( perm.at(2) ); + using raw_idx_type = VAL_T; + raw_idx_type mod_val = get_val(dim.at(perm.at(1))) * get_val(dim.at(perm.at(2))); for (RAJA::idx_t ii = 0; ii < N; ++ii) { test_array[ii] = static_cast(ii % mod_val); } RAJA::Layout<3> layout = RAJA::make_permuted_layout(dim_strip, perm); - RAJA::View< IDX_TYPE, RAJA::Layout<3, int> > view(working_array, layout); RAJA::kernel( - RAJA::make_tuple( RAJA::TypedRangeSegment(0, dim_strip.at(0)), - RAJA::TypedRangeSegment(0, dim_strip.at(1)), - RAJA::TypedRangeSegment(0, dim_strip.at(2)) ), + RAJA::make_tuple( RAJA::TypedRangeSegment(0, IDX_TYPE(dim_strip.at(0))), + RAJA::TypedRangeSegment(0, IDX_TYPE(dim_strip.at(1))), + RAJA::TypedRangeSegment(0, IDX_TYPE(dim_strip.at(2))) ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE j, IDX_TYPE k) { - int val = RAJA::stripIndexType(layout(i, j, k)) % mod_val; - view(i, j, k) = static_cast(val); + auto linear = RAJA::stripIndexType(layout(get_val(i), get_val(j), get_val(k))); + working_array[linear] = static_cast(linear % mod_val); } ); diff --git a/test/functional/kernel/reduce-loc/test-kernel-reduceloc.cpp.in b/test/functional/kernel/reduce-loc/test-kernel-reduceloc.cpp.in index cc3b2ee7e1..6917615f57 100644 --- a/test/functional/kernel/reduce-loc/test-kernel-reduceloc.cpp.in +++ b/test/functional/kernel/reduce-loc/test-kernel-reduceloc.cpp.in @@ -245,7 +245,7 @@ using HipKernelLocForallPols = // Cartesian product of types used in parameterized tests // using @LOC_BACKEND@KernelLocTypes = - Test< camp::cartesian_product +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMax2DTestImpl(const int xdim, const int ydim) +void KernelLocMax2DTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,60 +52,48 @@ void KernelLocMax2DTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType WorkView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx; + CheckView(zz, xx) = get_val(zz * xdim + xx); } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::ReduceMaxLoc maxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); + RAJA::ReduceMaxLoc> maxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - maxloc_reducer.maxloc(workarr2D[r][c], Index2D(c, r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + maxloc_reducer.maxloc(WorkView(r, c), Index2D(c, r)); }); // CPU answer - RAJA::ReduceMaxLoc checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); + RAJA::ReduceMaxLoc> checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkmaxloc_reducer.maxloc(checkarr2D[r][c], Index2D(c, r)); + checkmaxloc_reducer.maxloc(CheckView(r, c), Index2D(c, r)); } }); - Index2D raja_loc = maxloc_reducer.getLoc(); + Index2D raja_loc = maxloc_reducer.getLoc(); DATA_TYPE raja_max = (DATA_TYPE)maxloc_reducer.get(); - Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); + Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); DATA_TYPE checkraja_max = (DATA_TYPE)checkmaxloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_max, (DATA_TYPE)raja_max); @@ -97,12 +105,6 @@ void KernelLocMax2DTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DView.hpp b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DView.hpp index 6395a3b712..657d592507 100644 --- a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DView.hpp +++ b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DView.hpp @@ -10,17 +10,37 @@ #ifndef __TEST_KERNEL_REDUCELOC_MAX2DVIEW_HPP__ #define __TEST_KERNEL_REDUCELOC_MAX2DVIEW_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMax2DViewTestImpl(const int xdim, const int ydim) +void KernelLocMax2DViewTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,62 +52,48 @@ void KernelLocMax2DViewTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType ArrView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx; + CheckView(zz, xx) = get_val(zz * xdim + xx); } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::View> ArrView(work_array, xdim, ydim); - - RAJA::ReduceMaxLoc maxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); + RAJA::ReduceMaxLoc> maxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - maxloc_reducer.maxloc(ArrView(r, c), Index2D(c, r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + maxloc_reducer.maxloc(ArrView(r, c), Index2D(c, r)); }); // CPU answer - RAJA::ReduceMaxLoc checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); + RAJA::ReduceMaxLoc> checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkmaxloc_reducer.maxloc(checkarr2D[r][c], Index2D(c, r)); + checkmaxloc_reducer.maxloc(CheckView(r, c), Index2D(c, r)); } }); - Index2D raja_loc = maxloc_reducer.getLoc(); + Index2D raja_loc = maxloc_reducer.getLoc(); DATA_TYPE raja_max = (DATA_TYPE)maxloc_reducer.get(); - Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); + Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); DATA_TYPE checkraja_max = (DATA_TYPE)checkmaxloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_max, (DATA_TYPE)raja_max); @@ -99,12 +105,6 @@ void KernelLocMax2DViewTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DViewTuple.hpp b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DViewTuple.hpp index dbe5c475e8..9a26682a8e 100644 --- a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DViewTuple.hpp +++ b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Max2DViewTuple.hpp @@ -10,17 +10,37 @@ #ifndef __TEST_KERNEL_REDUCELOC_MAX2DVIEWTUPLE_HPP__ #define __TEST_KERNEL_REDUCELOC_MAX2DVIEWTUPLE_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMax2DViewTupleTestImpl(const int xdim, const int ydim) +void KernelLocMax2DViewTupleTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,36 +52,24 @@ void KernelLocMax2DViewTupleTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType ArrView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx; + CheckView(zz, xx) = get_val(zz * xdim + xx); } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); #if defined(RAJA_ENABLE_TARGET_OPENMP) //#pragma omp target data map(to:work_array[0:array_length]) @@ -70,30 +78,28 @@ void KernelLocMax2DViewTupleTestImpl(const int xdim, const int ydim) RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::View> ArrView(work_array, xdim, ydim); - RAJA::tuple LocTup(0, 0); RAJA::ReduceMaxLoc> maxloc_reducer((DATA_TYPE)0, LocTup); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - maxloc_reducer.maxloc(ArrView(r, c), RAJA::make_tuple((DATA_TYPE)c, (DATA_TYPE)r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + maxloc_reducer.maxloc(ArrView(r, c), RAJA::make_tuple((DATA_TYPE)get_val(c), (DATA_TYPE)get_val(r))); }); // CPU answer - RAJA::ReduceMaxLoc checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); + RAJA::ReduceMaxLoc> checkmaxloc_reducer((DATA_TYPE)0, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkmaxloc_reducer.maxloc(checkarr2D[r][c], Index2D(c, r)); + checkmaxloc_reducer.maxloc(CheckView(r, c), Index2D(c, r)); } }); RAJA::tuple raja_loc = maxloc_reducer.getLoc(); DATA_TYPE raja_max = (DATA_TYPE)maxloc_reducer.get(); - Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); + Index2D checkraja_loc = checkmaxloc_reducer.getLoc(); DATA_TYPE checkraja_max = (DATA_TYPE)checkmaxloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_max, (DATA_TYPE)raja_max); @@ -105,12 +111,6 @@ void KernelLocMax2DViewTupleTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2D.hpp b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2D.hpp index 30dc977023..8667b7cf34 100644 --- a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2D.hpp +++ b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2D.hpp @@ -10,17 +10,37 @@ #ifndef __TEST_KERNEL_REDUCELOC_MIN2D_HPP__ #define __TEST_KERNEL_REDUCELOC_MIN2D_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMin2DTestImpl(const int xdim, const int ydim) +void KernelLocMin2DTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,60 +52,48 @@ void KernelLocMin2DTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType WorkView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx + 1; + CheckView(zz, xx) = get_val(zz * xdim + xx) + 1; } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::ReduceMinLoc minloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); + RAJA::ReduceMinLoc> minloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - minloc_reducer.minloc(workarr2D[r][c], Index2D(c, r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + minloc_reducer.minloc(WorkView(r, c), Index2D(c, r)); }); // CPU answer - RAJA::ReduceMinLoc checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); + RAJA::ReduceMinLoc> checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkminloc_reducer.minloc(checkarr2D[r][c], Index2D(c, r)); + checkminloc_reducer.minloc(CheckView(r, c), Index2D(c, r)); } }); - Index2D raja_loc = minloc_reducer.getLoc(); + Index2D raja_loc = minloc_reducer.getLoc(); DATA_TYPE raja_min = (DATA_TYPE)minloc_reducer.get(); - Index2D checkraja_loc = checkminloc_reducer.getLoc(); + Index2D checkraja_loc = checkminloc_reducer.getLoc(); DATA_TYPE checkraja_min = (DATA_TYPE)checkminloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_min, (DATA_TYPE)raja_min); @@ -97,12 +105,6 @@ void KernelLocMin2DTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DView.hpp b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DView.hpp index 87386976ae..c8e6d246c7 100644 --- a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DView.hpp +++ b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DView.hpp @@ -10,17 +10,37 @@ #ifndef __TEST_KERNEL_REDUCELOC_MIN2DVIEW_HPP__ #define __TEST_KERNEL_REDUCELOC_MIN2DVIEW_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMin2DViewTestImpl(const int xdim, const int ydim) +void KernelLocMin2DViewTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,62 +52,48 @@ void KernelLocMin2DViewTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType ArrView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx + 1; + CheckView(zz, xx) = get_val(zz * xdim + xx) + 1; } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::View> ArrView(work_array, xdim, ydim); - - RAJA::ReduceMinLoc minloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); + RAJA::ReduceMinLoc> minloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - minloc_reducer.minloc(ArrView(r, c), Index2D(c, r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + minloc_reducer.minloc(ArrView(r, c), Index2D(c, r)); }); // CPU answer - RAJA::ReduceMinLoc checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); + RAJA::ReduceMinLoc> checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkminloc_reducer.minloc(checkarr2D[r][c], Index2D(c, r)); + checkminloc_reducer.minloc(CheckView(r, c), Index2D(c, r)); } }); - Index2D raja_loc = minloc_reducer.getLoc(); + Index2D raja_loc = minloc_reducer.getLoc(); DATA_TYPE raja_min = (DATA_TYPE)minloc_reducer.get(); - Index2D checkraja_loc = checkminloc_reducer.getLoc(); + Index2D checkraja_loc = checkminloc_reducer.getLoc(); DATA_TYPE checkraja_min = (DATA_TYPE)checkminloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_min, (DATA_TYPE)raja_min); @@ -99,12 +105,6 @@ void KernelLocMin2DViewTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DViewTuple.hpp b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DViewTuple.hpp index f80b4fea1b..90dafc9d73 100644 --- a/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DViewTuple.hpp +++ b/test/functional/kernel/reduce-loc/tests/test-kernel-reduceloc-Min2DViewTuple.hpp @@ -10,17 +10,37 @@ #ifndef __TEST_KERNEL_REDUCELOC_MIN2DVIEWTUPLE_HPP__ #define __TEST_KERNEL_REDUCELOC_MIN2DVIEWTUPLE_HPP__ +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } + +} + template -void KernelLocMin2DViewTupleTestImpl(const int xdim, const int ydim) +void KernelLocMin2DViewTupleTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,64 +52,50 @@ void KernelLocMin2DViewTupleTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - // set rows to point to check and work _arrays RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType ArrView(work_array, xdim, ydim); + ViewType CheckView(check_array, xdim, ydim); // initializing values RAJA::forall(seg, [=] (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx(0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = zz*xdim + xx + 1; + CheckView(zz, xx) = get_val(zz * xdim + xx) + 1; } - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - RAJA::View> ArrView(work_array, xdim, ydim); - RAJA::tuple LocTup(0, 0); RAJA::ReduceMinLoc> minloc_reducer((DATA_TYPE)1024, LocTup); RAJA::kernel(RAJA::make_tuple(colrange, rowrange), - [=] RAJA_HOST_DEVICE (int c, int r) { - minloc_reducer.minloc(ArrView(r, c), RAJA::make_tuple((DATA_TYPE)c, (DATA_TYPE)r)); + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, INDEX_TYPE r) { + minloc_reducer.minloc(ArrView(r, c), RAJA::make_tuple((DATA_TYPE)get_val(c), (DATA_TYPE)get_val(r))); }); // CPU answer - RAJA::ReduceMinLoc checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); + RAJA::ReduceMinLoc> checkminloc_reducer((DATA_TYPE)1024, Index2D(0, 0)); RAJA::forall(colrange, [=] (INDEX_TYPE c) { - for( int r = 0; r < ydim; ++r) + for (INDEX_TYPE r(0); r < ydim; ++r) { - checkminloc_reducer.minloc(checkarr2D[r][c], Index2D(c, r)); + checkminloc_reducer.minloc(CheckView(r, c), Index2D(c, r)); } }); RAJA::tuple raja_loc = minloc_reducer.getLoc(); DATA_TYPE raja_min = (DATA_TYPE)minloc_reducer.get(); - Index2D checkraja_loc = checkminloc_reducer.getLoc(); + Index2D checkraja_loc = checkminloc_reducer.getLoc(); DATA_TYPE checkraja_min = (DATA_TYPE)checkminloc_reducer.get(); ASSERT_DOUBLE_EQ((DATA_TYPE)checkraja_min, (DATA_TYPE)raja_min); @@ -101,12 +107,6 @@ void KernelLocMin2DViewTupleTestImpl(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/reduce-params-multi-lambda/test-kernel-reduce-params-multi-lambda.cpp.in b/test/functional/kernel/reduce-params-multi-lambda/test-kernel-reduce-params-multi-lambda.cpp.in index 8e5f4cd83f..5a060a2e5b 100644 --- a/test/functional/kernel/reduce-params-multi-lambda/test-kernel-reduce-params-multi-lambda.cpp.in +++ b/test/functional/kernel/reduce-params-multi-lambda/test-kernel-reduce-params-multi-lambda.cpp.in @@ -253,7 +253,7 @@ using HipKernelLocForallPols = // Cartesian product of types used in parameterized tests // using @BACKEND@KernelLocTypes = - Test< camp::cartesian_product +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelParamReduceMultiLambda(const int xdim, const int ydim) +void KernelParamReduceMultiLambda(const INDEX_TYPE xdim_t, const INDEX_TYPE ydim_t) { + using raw_index_type = VAL_T; + camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; DATA_TYPE * work_array; DATA_TYPE * check_array; DATA_TYPE * test_array; // square 2D array, xdim x ydim - INDEX_TYPE array_length = xdim * ydim; + INDEX_TYPE array_length = xdim_t * ydim_t; allocateForallTestData ( array_length, work_res, @@ -32,42 +53,30 @@ void KernelParamReduceMultiLambda(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - - // set rows to point to check and work _arrays - RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE (INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] RAJA_HOST_DEVICE (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); + using LayoutType = + RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType WorkView(work_array, ydim_t, xdim_t); + ViewType CheckView(check_array, ydim_t, xdim_t); // initializing values - RAJA::forall(seg, [=] RAJA_HOST_DEVICE (INDEX_TYPE zz) + for (INDEX_TYPE row(0); row < ydim_t; ++row) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( raw_index_type col = 0; col < get_val(xdim_t); ++col ) { - checkarr2D[zz][xx] = (zz*xdim + xx ) % 100 + 1; + CheckView(row, INDEX_TYPE(col)) = + (get_val(row) * get_val(xdim_t) + col) % 100 + 1; } // Make a unique min - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim_t - INDEX_TYPE(1), xdim_t - INDEX_TYPE(1)) = 0; // Make a unique max - checkarr2D[ydim/2][xdim/2] = 101; - }); + CheckView(ydim_t / INDEX_TYPE(2), xdim_t / INDEX_TYPE(2)) = 101; + } - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); - RAJA::TypedRangeSegment colrange(0, xdim); - RAJA::TypedRangeSegment rowrange(0, ydim); + RAJA::TypedRangeSegment colrange(0, xdim_t); + RAJA::TypedRangeSegment rowrange(0, ydim_t); using VALOP_DATA_TYPE_SUM = RAJA::expt::ValOp; using VALOP_DATA_TYPE_MIN = RAJA::expt::ValOp; @@ -104,28 +113,28 @@ void KernelParamReduceMultiLambda(const int xdim, const int ydim) // not init or resolving reducers in between lambda calls. RAJA::expt::Reduce(&mutual_sum) ), - [=] RAJA_HOST_DEVICE (int c, - int r, + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, + INDEX_TYPE r, VALOP_DATA_TYPE_SUM &_sum, VALOP_DATA_TYPE_MIN &_min, VALOP_DATA_TYPE_MAX &_max, VALOP_DATA_TYPE_SUM &_mutual_sum) { - _sum += workarr2D[r][c]; - _mutual_sum += workarr2D[r][c]; - _min.min(workarr2D[r][c]); - _max.max(workarr2D[r][c]); + _sum += WorkView(r, c); + _mutual_sum += WorkView(r, c); + _min.min(WorkView(r, c)); + _max.max(WorkView(r, c)); }, - [=] RAJA_HOST_DEVICE (int c, - int r, + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, + INDEX_TYPE r, VALOP_DATA_TYPE_SUM &_sum, VALOP_DATA_TYPE_MIN &_min, VALOP_DATA_TYPE_MAX &_max, double, VALOP_DATA_TYPE_SUM &_mutual_sum) { - _sum += workarr2D[r][c]; - _mutual_sum += workarr2D[r][c]; - _min.min(workarr2D[r][c]); - _max.max(workarr2D[r][c]); + _sum += WorkView(r, c); + _mutual_sum += WorkView(r, c); + _min.min(WorkView(r, c)); + _max.max(WorkView(r, c)); } ); @@ -141,11 +150,12 @@ void KernelParamReduceMultiLambda(const int xdim, const int ydim) VALOP_DATA_TYPE_MIN &_min, VALOP_DATA_TYPE_MAX &_max ) { - for (int c = 0; c < xdim; ++c) + for (raw_index_type c = 0; c < get_val(xdim_t); ++c) { - _sum += checkarr2D[r][c]; - _min.min(checkarr2D[r][c]); - _max.max(checkarr2D[r][c]); + auto c_idx = INDEX_TYPE(c); + _sum += CheckView(r, c_idx); + _min.min(CheckView(r, c_idx)); + _max.max(CheckView(r, c_idx)); } }); @@ -169,12 +179,6 @@ void KernelParamReduceMultiLambda(const int xdim, const int ydim) check_array, test_array ); - - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } @@ -192,9 +196,9 @@ TYPED_TEST_P(KernelParamReduceMultiLambdaTest, ParamReduceKernel) using FORALL_POLICY = typename camp::at>::type; using EXEC_POLICY = typename camp::at>::type; - KernelParamReduceMultiLambda(10, 10); - KernelParamReduceMultiLambda(151, 151); - KernelParamReduceMultiLambda(362, 362); + KernelParamReduceMultiLambda(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelParamReduceMultiLambda(INDEX_TYPE(151), INDEX_TYPE(151)); + KernelParamReduceMultiLambda(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelParamReduceMultiLambdaTest, diff --git a/test/functional/kernel/reduce-params/test-kernel-basic-param.cpp.in b/test/functional/kernel/reduce-params/test-kernel-basic-param.cpp.in index f7e11e75de..b0140e4692 100644 --- a/test/functional/kernel/reduce-params/test-kernel-basic-param.cpp.in +++ b/test/functional/kernel/reduce-params/test-kernel-basic-param.cpp.in @@ -245,7 +245,7 @@ using HipKernelLocForallPols = // Cartesian product of types used in parameterized tests // using @LOC_BACKEND@KernelLocTypes = - Test< camp::cartesian_product +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +auto get_val(T index_val) { return index_val; } + +} + template -void KernelParamReduceTestImpl(const int xdim, const int ydim) +void KernelParamReduceTestImpl(const INDEX_TYPE xdim, const INDEX_TYPE ydim) { camp::resources::Resource work_res{WORKING_RES::get_default()}; - DATA_TYPE ** workarr2D; - DATA_TYPE ** checkarr2D; - DATA_TYPE ** testarr2D; - DATA_TYPE * work_array; - DATA_TYPE * check_array; - DATA_TYPE * test_array; + DATA_TYPE* work_array; + DATA_TYPE* check_array; + DATA_TYPE* test_array; // square 2D array, xdim x ydim INDEX_TYPE array_length = xdim * ydim; @@ -32,64 +52,53 @@ void KernelParamReduceTestImpl(const int xdim, const int ydim) &test_array ); - allocateForallTestData ( ydim, - work_res, - &workarr2D, - &checkarr2D, - &testarr2D - ); - + // using index_setup_type = VAL_T; // set rows to point to check and work _arrays - RAJA::TypedRangeSegment seg(0,ydim); - RAJA::forall(seg, [=] RAJA_HOST_DEVICE(INDEX_TYPE zz) - { - workarr2D[zz] = work_array + zz * ydim; - }); - - RAJA::forall(seg, [=] RAJA_HOST_DEVICE (INDEX_TYPE zz) - { - checkarr2D[zz] = check_array + zz * ydim; - }); - + RAJA::TypedRangeSegment seg(0, ydim); + using LayoutType = RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType TestView (test_array, xdim, ydim ); + ViewType WorkView (work_array, xdim, ydim ); + ViewType CheckView (check_array, xdim, ydim ); // initializing values RAJA::forall(seg, [=] RAJA_HOST_DEVICE (INDEX_TYPE zz) { - for ( int xx = 0; xx < xdim; ++xx ) + for ( INDEX_TYPE xx (0); xx < xdim; ++xx ) { - checkarr2D[zz][xx] = (zz*xdim + xx ) % 100 + 1; + CheckView(zz, xx) = get_val(zz * xdim + xx ) % 100 + 1; } // Make a unique min - checkarr2D[ydim-1][xdim-1] = 0; + CheckView(ydim - 1, xdim - 1) = 0; // Make a unique max - checkarr2D[ydim/2][xdim/2] = 101; + CheckView(ydim / 2, xdim / 2) = 101; }); - work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * array_length); + work_res.memcpy(work_array, check_array, sizeof(DATA_TYPE) * get_val(array_length)); RAJA::TypedRangeSegment colrange(0, xdim); RAJA::TypedRangeSegment rowrange(0, ydim); - using VALLOC_DATA_TYPE = RAJA::expt::ValLoc; + using VALLOC_DATA_TYPE = RAJA::expt::ValLoc>; using VALOP_DATA_TYPE_SUM = RAJA::expt::ValOp; using VALOP_DATA_TYPE_MIN = RAJA::expt::ValOp; using VALOP_DATA_TYPE_MAX = RAJA::expt::ValOp; - using VALOPLOC_DATA_TYPE_MIN = RAJA::expt::ValLocOp; - using VALOPLOC_DATA_TYPE_MAX = RAJA::expt::ValLocOp; + using VALOPLOC_DATA_TYPE_MIN = RAJA::expt::ValLocOp, RAJA::operators::minimum>; + using VALOPLOC_DATA_TYPE_MAX = RAJA::expt::ValLocOp, RAJA::operators::maximum>; - VALLOC_DATA_TYPE seq_minloc(std::numeric_limits::max(), Index2D(-1,-1)); - VALLOC_DATA_TYPE seq_maxloc(std::numeric_limits::min(), Index2D(-1,-1)); - Index2D seq_minloc2(-1, -1); - Index2D seq_maxloc2(-1, -1); + VALLOC_DATA_TYPE seq_minloc(std::numeric_limits::max(), Index2D(-1,-1)); + VALLOC_DATA_TYPE seq_maxloc(std::numeric_limits::min(), Index2D(-1,-1)); + Index2D seq_minloc2(-1, -1); + Index2D seq_maxloc2(-1, -1); DATA_TYPE seq_sum = 0; DATA_TYPE seq_min = std::numeric_limits::max(); DATA_TYPE seq_max = std::numeric_limits::min(); DATA_TYPE seq_min2 = std::numeric_limits::max(); DATA_TYPE seq_max2 = std::numeric_limits::min(); - VALLOC_DATA_TYPE minloc(std::numeric_limits::max(), Index2D(-1,-1)); - VALLOC_DATA_TYPE maxloc(std::numeric_limits::min(), Index2D(-1,-1)); - Index2D minloc2(-1, -1); - Index2D maxloc2(-1, -1); + VALLOC_DATA_TYPE minloc(std::numeric_limits::max(), Index2D(-1,-1)); + VALLOC_DATA_TYPE maxloc(std::numeric_limits::min(), Index2D(-1,-1)); + Index2D minloc2(-1, -1); + Index2D maxloc2(-1, -1); DATA_TYPE sum = 0; DATA_TYPE min2 = std::numeric_limits::max(); DATA_TYPE max2 = std::numeric_limits::min(); @@ -109,8 +118,8 @@ void KernelParamReduceTestImpl(const int xdim, const int ydim) RAJA::expt::ReduceLoc(&min2, &minloc2), RAJA::expt::ReduceLoc(&max2, &maxloc2) ), - [=] RAJA_HOST_DEVICE (int c, - int r, + [=] RAJA_HOST_DEVICE (INDEX_TYPE c, + INDEX_TYPE r, VALOP_DATA_TYPE_SUM &_sum, VALOP_DATA_TYPE_MIN &_min, VALOP_DATA_TYPE_MAX &_max, @@ -118,15 +127,15 @@ void KernelParamReduceTestImpl(const int xdim, const int ydim) VALOPLOC_DATA_TYPE_MAX &_maxloc, VALOPLOC_DATA_TYPE_MIN &_minloc2, VALOPLOC_DATA_TYPE_MAX &_maxloc2) { - _sum += workarr2D[r][c]; - _min.min(workarr2D[r][c]); - _max.max(workarr2D[r][c]); + _sum += WorkView(r,c); + _min.min(WorkView(r,c)); + _max.max(WorkView(r,c)); // loc - _minloc.minloc(workarr2D[r][c], Index2D(c, r)); - _maxloc.maxloc(workarr2D[r][c], Index2D(c, r)); - _minloc2.minloc(workarr2D[r][c], Index2D(c, r)); - _maxloc2.maxloc(workarr2D[r][c], Index2D(c, r)); + _minloc.minloc(WorkView(r,c), Index2D(c, r)); + _maxloc.maxloc(WorkView(r,c), Index2D(c, r)); + _minloc2.minloc(WorkView(r,c), Index2D(c, r)); + _maxloc2.maxloc(WorkView(r,c), Index2D(c, r)); }); // CPU answer @@ -148,24 +157,24 @@ void KernelParamReduceTestImpl(const int xdim, const int ydim) VALOPLOC_DATA_TYPE_MIN &_minloc2, VALOPLOC_DATA_TYPE_MAX &_maxloc2 ) { - for (int c = 0; c < xdim; ++c) + for (INDEX_TYPE c(0); c < xdim; ++c) { - _sum += checkarr2D[r][c]; - _min = _min.min(checkarr2D[r][c]); - _max = _max.max(checkarr2D[r][c]); + _sum += CheckView(r,c); + _min = _min.min(CheckView(r,c)); + _max = _max.max(CheckView(r,c)); // loc - _minloc.minloc(checkarr2D[r][c], Index2D(c, r)); - _maxloc.maxloc(checkarr2D[r][c], Index2D(c, r)); - _minloc2.minloc(checkarr2D[r][c], Index2D(c, r)); - _maxloc2.maxloc(checkarr2D[r][c], Index2D(c, r)); + _minloc.minloc(CheckView(r,c), Index2D(c, r)); + _maxloc.maxloc(CheckView(r,c), Index2D(c, r)); + _minloc2.minloc(CheckView(r,c), Index2D(c, r)); + _maxloc2.maxloc(CheckView(r,c), Index2D(c, r)); } }); DATA_TYPE DEBUG_SUM = 0; - for (int r = 0 ; r < ydim; ++r) { - for (int c = 0; c < xdim; ++c) { - DEBUG_SUM += checkarr2D[r][c]; + for (INDEX_TYPE r (0) ; r < ydim; ++r) { + for (INDEX_TYPE c (0); c < xdim; ++c) { + DEBUG_SUM += CheckView(r,c); } } @@ -195,11 +204,6 @@ void KernelParamReduceTestImpl(const int xdim, const int ydim) test_array ); - deallocateForallTestData ( work_res, - workarr2D, - checkarr2D, - testarr2D - ); } diff --git a/test/functional/kernel/region/test-kernel-region-sync.cpp.in b/test/functional/kernel/region/test-kernel-region-sync.cpp.in index 31986a9d77..6e307538a9 100644 --- a/test/functional/kernel/region/test-kernel-region-sync.cpp.in +++ b/test/functional/kernel/region/test-kernel-region-sync.cpp.in @@ -90,7 +90,7 @@ using OpenMPKernelRegionSyncExecPols = // Cartesian product of types used in parameterized tests // using @REGION_BACKEND@KernelRegionSyncTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/region/test-kernel-region.cpp.in b/test/functional/kernel/region/test-kernel-region.cpp.in index 5d9e6b6731..719e06f768 100644 --- a/test/functional/kernel/region/test-kernel-region.cpp.in +++ b/test/functional/kernel/region/test-kernel-region.cpp.in @@ -72,7 +72,7 @@ using OpenMPKernelRegionExecPols = // Cartesian product of types used in parameterized tests // using @REGION_BACKEND@KernelRegionTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/tile-icount-tcount-direct-unchecked/test-kernel-tile-count-direct-unchecked.cpp.in b/test/functional/kernel/tile-icount-tcount-direct-unchecked/test-kernel-tile-count-direct-unchecked.cpp.in index b3bd80e8ca..8e92b0ed89 100644 --- a/test/functional/kernel/tile-icount-tcount-direct-unchecked/test-kernel-tile-count-direct-unchecked.cpp.in +++ b/test/functional/kernel/tile-icount-tcount-direct-unchecked/test-kernel-tile-count-direct-unchecked.cpp.in @@ -109,7 +109,7 @@ using HipKernelTileTCountExecPols = camp::list< // Cartesian product of types used in parameterized tests // using @BACKEND@KernelTile@TESTTYPE@Types = - Test< camp::cartesian_product& _trip_count, RAJA::expt::ValOp& _tile_count) { - _trip_count += 1; + _trip_count += IDX_TYPE(1); if ( i % tsize == t && ii == t ) { - _tile_count += 1; + _tile_count += IDX_TYPE(1); } } ); @@ -62,9 +62,9 @@ CallKernel(IDX_TYPE& _trip_count, RAJA::make_tuple( static_cast(0) ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE ii) { - trip_count += 1; + trip_count += IDX_TYPE(1); if ( i % tsize == t && ii == t ) { - tile_count += 1; + tile_count += IDX_TYPE(1); } } ); @@ -83,7 +83,7 @@ void KernelTileForICountDirectUncheckedTestImpl(IDX_TYPE N, IDX_TYPE tsize) IDX_TYPE tile_expect = N / tsize; if ( t < N % tsize ) { - tile_expect += 1; + tile_expect += IDX_TYPE(1); } ASSERT_EQ(trip_count, (t+1) * N); ASSERT_EQ(tile_count, tile_expect); diff --git a/test/functional/kernel/tile-icount-tcount-direct/test-kernel-tile-count-direct.cpp.in b/test/functional/kernel/tile-icount-tcount-direct/test-kernel-tile-count-direct.cpp.in index b39f1573a5..fa2de78ca2 100644 --- a/test/functional/kernel/tile-icount-tcount-direct/test-kernel-tile-count-direct.cpp.in +++ b/test/functional/kernel/tile-icount-tcount-direct/test-kernel-tile-count-direct.cpp.in @@ -236,7 +236,7 @@ using SyclKernelTileTCountExecPols = camp::list< // Cartesian product of types used in parameterized tests // using @BACKEND@KernelTile@TESTTYPE@Types = - Test< camp::cartesian_product& _trip_count, RAJA::expt::ValOp& _tile_count) { - _trip_count += 1; + _trip_count += IDX_TYPE(1); if ( i % tsize == t && ii == t ) { - _tile_count += 1; + _tile_count += IDX_TYPE(1); } } ); @@ -60,9 +60,9 @@ CallKernel(IDX_TYPE& _trip_count, RAJA::make_tuple( static_cast(0) ), [=] RAJA_HOST_DEVICE(IDX_TYPE i, IDX_TYPE ii) { - trip_count += 1; + trip_count += IDX_TYPE(1); if ( i % tsize == t && ii == t ) { - tile_count += 1; + tile_count += IDX_TYPE(1); } } ); @@ -75,8 +75,8 @@ void KernelTileForICountDirectTestImpl(IDX_TYPE N, IDX_TYPE tsize) { IDX_TYPE trip_count(0); - for (IDX_TYPE t = 0; t < tsize; ++t) { - IDX_TYPE tile_count = 0; + for (IDX_TYPE t(0); t < tsize; ++t) { + IDX_TYPE tile_count(0); CallKernel(trip_count, tile_count, t, N, tsize); IDX_TYPE tile_expect = N / tsize; @@ -105,7 +105,7 @@ TYPED_TEST_P(KernelTileForICountDirectTest, ForICountTileKernel) using REDUCE_POLICY = typename camp::at>::type; using USE_REDUCER_PARAM = typename camp::at>::type; - IDX_TYPE tsize = camp::at_v::value; + IDX_TYPE tsize(camp::at_v::value); KernelTileForICountDirectTestImpl( IDX_TYPE(57), tsize); diff --git a/test/functional/kernel/tile-icount-tcount-loop/test-kernel-tile-count-loop.cpp.in b/test/functional/kernel/tile-icount-tcount-loop/test-kernel-tile-count-loop.cpp.in index 7f9701d4f8..89679d9bdc 100644 --- a/test/functional/kernel/tile-icount-tcount-loop/test-kernel-tile-count-loop.cpp.in +++ b/test/functional/kernel/tile-icount-tcount-loop/test-kernel-tile-count-loop.cpp.in @@ -236,7 +236,7 @@ using SyclKernelTileTCountExecPols = camp::list< // Cartesian product of types used in parameterized tests // using @BACKEND@KernelTile@TESTTYPE@Types = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/tile-variants/test-kernel-tilefixed.cpp.in b/test/functional/kernel/tile-variants/test-kernel-tilefixed.cpp.in index c397aa039b..f12996f0cb 100644 --- a/test/functional/kernel/tile-variants/test-kernel-tilefixed.cpp.in +++ b/test/functional/kernel/tile-variants/test-kernel-tilefixed.cpp.in @@ -270,7 +270,7 @@ using SyclKernelTileExecPols = // Cartesian product of types used in parameterized tests // using @TILE_BACKEND@KernelTileTypes = - Test< camp::cartesian_product, RAJA::statement::CudaSyncThreads, - + RAJA::statement::ForICount<0, RAJA::statement::Param<1>, RAJA::cuda_thread_x_loop, RAJA::statement::ForICount<1, RAJA::statement::Param<0>, RAJA::cuda_thread_y_direct, RAJA::statement::Lambda<1> @@ -146,7 +146,7 @@ using HipKernelTileExecPols = >, RAJA::statement::HipSyncThreads, - + RAJA::statement::ForICount<0, RAJA::statement::Param<1>, RAJA::hip_thread_x_loop, RAJA::statement::ForICount<1, RAJA::statement::Param<0>, RAJA::hip_thread_y_direct, RAJA::statement::Lambda<1> @@ -205,7 +205,7 @@ using HipKernelTileExecPols = // Cartesian product of types used in parameterized tests // using @TILE_BACKEND@KernelTileTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Dynamic2D.hpp b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Dynamic2D.hpp index d13c1a6146..d513fe0a2a 100644 --- a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Dynamic2D.hpp +++ b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Dynamic2D.hpp @@ -12,10 +12,33 @@ #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelTileDynamic2DTestImpl(const int rows, const int cols) +void KernelTileDynamic2DTestImpl(const INDEX_TYPE rows_t, const INDEX_TYPE cols_t) { // This test emulates matrix transposition with tiling. + using raw_index_type = VAL_T; camp::resources::Resource work_res{WORKING_RES::get_default()}; @@ -28,7 +51,7 @@ void KernelTileDynamic2DTestImpl(const int rows, const int cols) DATA_TYPE * check_array_t; DATA_TYPE * test_array_t; - INDEX_TYPE array_length = rows * cols; + INDEX_TYPE array_length = rows_t * cols_t; allocateForallTestData ( array_length, work_res, @@ -44,31 +67,35 @@ void KernelTileDynamic2DTestImpl(const int rows, const int cols) &test_array_t ); - RAJA::View> HostView( test_array, rows, cols ); - RAJA::View> HostTView( test_array_t, cols, rows ); - RAJA::View> WorkView( work_array, rows, cols ); - RAJA::View> WorkTView( work_array_t, cols, rows ); - RAJA::View> CheckTView( check_array_t, cols, rows ); + using LayoutType = + RAJA::TypedLayout>; + using ViewType = RAJA::View; + + ViewType HostView( test_array, rows_t, cols_t ); + ViewType HostTView( test_array_t, cols_t, rows_t ); + ViewType WorkView( work_array, rows_t, cols_t ); + ViewType WorkTView( work_array_t, cols_t, rows_t ); + ViewType CheckTView( check_array_t, cols_t, rows_t ); // initialize arrays - std::iota( test_array, test_array + array_length, 1 ); - std::iota( test_array_t, test_array_t + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); + std::iota( test_array_t, test_array_t + get_val(array_length), 1 ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); - work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); + work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); // transpose test_array on CPU - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - HostTView( cc, rr ) = HostView( rr, cc ); + HostTView( INDEX_TYPE(cc), INDEX_TYPE(rr) ) = HostView( INDEX_TYPE(rr), INDEX_TYPE(cc) ); } } // transpose work_array - RAJA::TypedRangeSegment rowrange( 0, rows ); - RAJA::TypedRangeSegment colrange( 0, cols ); + RAJA::TypedRangeSegment rowrange( 0, rows_t ); + RAJA::TypedRangeSegment colrange( 0, cols_t ); RAJA::kernel_param ( RAJA::make_tuple( colrange, rowrange ), @@ -77,19 +104,20 @@ void KernelTileDynamic2DTestImpl(const int rows, const int cols) WorkTView( cc, rr ) = WorkView( rr, cc ); }); - work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - ASSERT_EQ(CheckTView(cc, rr), HostTView(cc, rr)); + ASSERT_EQ(CheckTView(INDEX_TYPE(cc), INDEX_TYPE(rr)), + HostTView(INDEX_TYPE(cc), INDEX_TYPE(rr))); } } // reset check and work transpose arrays - work_res.memcpy( check_array_t, test_array, sizeof(DATA_TYPE) * array_length ); - work_res.memcpy( work_array_t, test_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array_t, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); + work_res.memcpy( work_array_t, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); // transpose work_array again with different tile sizes RAJA::kernel_param ( @@ -99,13 +127,14 @@ void KernelTileDynamic2DTestImpl(const int rows, const int cols) WorkTView( cc, rr ) = WorkView( rr, cc ); }); - work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - ASSERT_EQ(CheckTView(cc, rr), HostTView(cc, rr)); + ASSERT_EQ(CheckTView(INDEX_TYPE(cc), INDEX_TYPE(rr)), + HostTView(INDEX_TYPE(cc), INDEX_TYPE(rr))); } } @@ -136,9 +165,9 @@ TYPED_TEST_P(KernelTileDynamic2DTest, TileDynamic2DKernel) using WORKING_RES = typename camp::at>::type; using EXEC_POLICY = typename camp::at>::type; - KernelTileDynamic2DTestImpl(10, 10); - KernelTileDynamic2DTestImpl(151, 111); - KernelTileDynamic2DTestImpl(362, 362); + KernelTileDynamic2DTestImpl(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelTileDynamic2DTestImpl(INDEX_TYPE(151), INDEX_TYPE(111)); + KernelTileDynamic2DTestImpl(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelTileDynamic2DTest, diff --git a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2D.hpp b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2D.hpp index 07e5fe361d..b2f105fee2 100644 --- a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2D.hpp +++ b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2D.hpp @@ -12,10 +12,33 @@ #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelTileFixed2DTestImpl(const int rows, const int cols) +void KernelTileFixed2DTestImpl(const INDEX_TYPE rows_t, const INDEX_TYPE cols_t) { // This test emulates matrix transposition with tiling. + using raw_index_type = VAL_T; camp::resources::Resource work_res{WORKING_RES::get_default()}; @@ -28,7 +51,7 @@ void KernelTileFixed2DTestImpl(const int rows, const int cols) DATA_TYPE * check_array_t; DATA_TYPE * test_array_t; - INDEX_TYPE array_length = rows * cols; + INDEX_TYPE array_length = rows_t * cols_t; allocateForallTestData ( array_length, work_res, @@ -44,44 +67,49 @@ void KernelTileFixed2DTestImpl(const int rows, const int cols) &test_array_t ); - RAJA::View> HostView( test_array, rows, cols ); - RAJA::View> HostTView( test_array_t, cols, rows ); - RAJA::View> WorkView( work_array, rows, cols ); - RAJA::View> WorkTView( work_array_t, cols, rows ); - RAJA::View> CheckTView( check_array_t, cols, rows ); + using LayoutType = + RAJA::TypedLayout>; + using ViewType = RAJA::View; + + ViewType HostView( test_array, rows_t, cols_t ); + ViewType HostTView( test_array_t, cols_t, rows_t ); + ViewType WorkView( work_array, rows_t, cols_t ); + ViewType WorkTView( work_array_t, cols_t, rows_t ); + ViewType CheckTView( check_array_t, cols_t, rows_t ); // initialize arrays - std::iota( test_array, test_array + array_length, 1 ); - std::iota( test_array_t, test_array_t + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); + std::iota( test_array_t, test_array_t + get_val(array_length), 1 ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); - work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); + work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); // transpose test_array on CPU - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - HostTView( cc, rr ) = HostView( rr, cc ); + HostTView( INDEX_TYPE(cc), INDEX_TYPE(rr) ) = HostView( INDEX_TYPE(rr), INDEX_TYPE(cc) ); } } // transpose work_array - RAJA::TypedRangeSegment rowrange( 0, rows ); - RAJA::TypedRangeSegment colrange( 0, cols ); + RAJA::TypedRangeSegment rowrange( 0, rows_t ); + RAJA::TypedRangeSegment colrange( 0, cols_t ); RAJA::kernel ( RAJA::make_tuple( colrange, rowrange ), [=] RAJA_HOST_DEVICE ( INDEX_TYPE cc, INDEX_TYPE rr ) { WorkTView( cc, rr ) = WorkView( rr, cc ); }); - work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - ASSERT_EQ(CheckTView(cc, rr), HostTView(cc, rr)); + ASSERT_EQ(CheckTView(INDEX_TYPE(cc), INDEX_TYPE(rr)), + HostTView(INDEX_TYPE(cc), INDEX_TYPE(rr))); } } @@ -112,9 +140,9 @@ TYPED_TEST_P(KernelTileFixed2DTest, TileFixed2DKernel) using WORKING_RES = typename camp::at>::type; using EXEC_POLICY = typename camp::at>::type; - KernelTileFixed2DTestImpl(10, 10); - KernelTileFixed2DTestImpl(151, 111); - KernelTileFixed2DTestImpl(362, 362); + KernelTileFixed2DTestImpl(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelTileFixed2DTestImpl(INDEX_TYPE(151), INDEX_TYPE(111)); + KernelTileFixed2DTestImpl(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelTileFixed2DTest, diff --git a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DMinMax.hpp b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DMinMax.hpp index 02ab058ff8..b935057565 100644 --- a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DMinMax.hpp +++ b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DMinMax.hpp @@ -13,10 +13,33 @@ #include #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelTileFixed2DMinMaxTestImpl(const int rows, const int cols) +void KernelTileFixed2DMinMaxTestImpl(const INDEX_TYPE rows_t, const INDEX_TYPE cols_t) { // This test reduces min and max with tiling. + using raw_index_type = VAL_T; camp::resources::Resource work_res{WORKING_RES::get_default()}; @@ -24,7 +47,7 @@ void KernelTileFixed2DMinMaxTestImpl(const int rows, const int cols) DATA_TYPE * check_array; DATA_TYPE * test_array; - INDEX_TYPE array_length = rows * cols; + INDEX_TYPE array_length = rows_t * cols_t; allocateForallTestData ( array_length, work_res, @@ -34,26 +57,29 @@ void KernelTileFixed2DMinMaxTestImpl(const int rows, const int cols) ); // initialize arrays - std::iota( test_array, test_array + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); // set min and max of the array test_array[4] = -1; - test_array[8] = array_length+2; + test_array[8] = static_cast(get_val(array_length) + 2); - RAJA::View> WorkView( work_array, rows, cols ); + using LayoutType = + RAJA::TypedLayout>; + using ViewType = RAJA::View; + ViewType WorkView( work_array, rows_t, cols_t ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); RAJA::ReduceMin workmin( DATA_TYPE(99999) ); RAJA::ReduceMax workmax( DATA_TYPE(-1) ); // mixed range types - RAJA::TypedRangeSegment rowrange( 0, rows ); + RAJA::TypedRangeSegment rowrange( 0, rows_t ); std::vector colidx; - for (INDEX_TYPE ii = INDEX_TYPE(0); ii < static_cast(cols); ++ii) + for (raw_index_type ii = 0; ii < get_val(cols_t); ++ii) { - colidx.push_back(ii); + colidx.push_back(INDEX_TYPE(ii)); } RAJA::TypedListSegment colrange( &colidx[0], colidx.size(), work_res ); @@ -66,7 +92,7 @@ void KernelTileFixed2DMinMaxTestImpl(const int rows, const int cols) }); ASSERT_EQ(static_cast(-1), static_cast(workmin.get())); - ASSERT_EQ(static_cast(array_length+2), static_cast(workmax.get())); + ASSERT_EQ(static_cast(get_val(array_length) + 2), static_cast(workmax.get())); deallocateForallTestData ( work_res, work_array, @@ -90,9 +116,9 @@ TYPED_TEST_P(KernelTileFixed2DMinMaxTest, TileFixed2DMinMaxKernel) using EXEC_POLICY = typename camp::at>::type; using REDUCE_POLICY = typename camp::at>::type; - KernelTileFixed2DMinMaxTestImpl(10, 10); - KernelTileFixed2DMinMaxTestImpl(151, 111); - KernelTileFixed2DMinMaxTestImpl(362, 362); + KernelTileFixed2DMinMaxTestImpl(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelTileFixed2DMinMaxTestImpl(INDEX_TYPE(151), INDEX_TYPE(111)); + KernelTileFixed2DMinMaxTestImpl(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelTileFixed2DMinMaxTest, diff --git a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DSum.hpp b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DSum.hpp index e9b88cda6f..ec8d93ade5 100644 --- a/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DSum.hpp +++ b/test/functional/kernel/tile-variants/tests/test-kernel-tile-Fixed2DSum.hpp @@ -14,12 +14,35 @@ #include #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelTileFixed2DSumTestImpl(const int rowsin, const int colsin) +void KernelTileFixed2DSumTestImpl(const INDEX_TYPE rowsin, const INDEX_TYPE colsin) { // This test reduces sums with tiling. + using raw_index_type = VAL_T; - int rows, cols; + raw_index_type rows, cols; if ( std::is_same::value ) { // Restrict to a small data size for better float precision. @@ -28,8 +51,8 @@ void KernelTileFixed2DSumTestImpl(const int rowsin, const int colsin) } else { - rows = rowsin; - cols = colsin; + rows = get_val(rowsin); + cols = get_val(colsin); } camp::resources::Resource work_res{WORKING_RES::get_default()}; @@ -39,13 +62,13 @@ void KernelTileFixed2DSumTestImpl(const int rowsin, const int colsin) RAJA::ReduceSum worksum( DATA_TYPE(0) ); // sum on CPU in a tiled manner - for ( int rr = 0; rr < rows; rr += tile_dim_x ) + for ( raw_index_type rr = 0; rr < rows; rr += tile_dim_x ) { - for ( int cc = 0; cc < cols; cc += tile_dim_y ) + for ( raw_index_type cc = 0; cc < cols; cc += tile_dim_y ) { - for ( int r = rr; r < std::min(rr+tile_dim_x, rows); ++r ) + for ( raw_index_type r = rr; r < std::min(rr + tile_dim_x, rows); ++r ) { - for ( int c = cc; c < std::min(cc+tile_dim_y, cols); ++c ) + for ( raw_index_type c = cc; c < std::min(cc + tile_dim_y, cols); ++c ) { hostsum += (DATA_TYPE)(r * 1.1 + c); } @@ -54,12 +77,12 @@ void KernelTileFixed2DSumTestImpl(const int rowsin, const int colsin) } // mixed range types - RAJA::TypedRangeSegment rowrange( 0, rows ); + RAJA::TypedRangeSegment rowrange( 0, INDEX_TYPE(rows) ); std::vector colidx; - for (INDEX_TYPE ii = INDEX_TYPE(0); ii < static_cast(cols); ++ii) + for (raw_index_type ii = 0; ii < cols; ++ii) { - colidx.push_back(ii); + colidx.push_back(INDEX_TYPE(ii)); } RAJA::TypedListSegment colrange( &colidx[0], colidx.size(), work_res ); @@ -67,7 +90,7 @@ void KernelTileFixed2DSumTestImpl(const int rowsin, const int colsin) // sum on target platform RAJA::kernel ( RAJA::make_tuple( colrange, rowrange ), [=] RAJA_HOST_DEVICE ( INDEX_TYPE cc, INDEX_TYPE rr ) { - worksum += (DATA_TYPE)(rr * 1.1 + cc); + worksum += (DATA_TYPE)(get_val(rr) * 1.1 + get_val(cc)); }); ASSERT_FLOAT_EQ(hostsum, (DATA_TYPE)worksum.get()); @@ -88,9 +111,9 @@ TYPED_TEST_P(KernelTileFixed2DSumTest, TileFixed2DSumKernel) using EXEC_POLICY = typename camp::at>::type; using REDUCE_POLICY = typename camp::at>::type; - KernelTileFixed2DSumTestImpl(10, 10); - KernelTileFixed2DSumTestImpl(151, 111); - KernelTileFixed2DSumTestImpl(362, 362); + KernelTileFixed2DSumTestImpl(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelTileFixed2DSumTestImpl(INDEX_TYPE(151), INDEX_TYPE(111)); + KernelTileFixed2DSumTestImpl(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelTileFixed2DSumTest, diff --git a/test/functional/kernel/tile-variants/tests/test-kernel-tile-LocalArray2D.hpp b/test/functional/kernel/tile-variants/tests/test-kernel-tile-LocalArray2D.hpp index bd82821349..4d013d93dd 100644 --- a/test/functional/kernel/tile-variants/tests/test-kernel-tile-LocalArray2D.hpp +++ b/test/functional/kernel/tile-variants/tests/test-kernel-tile-LocalArray2D.hpp @@ -12,10 +12,33 @@ #include +namespace { +template +struct val_t_impl { + using type = T; +}; + +template +struct val_t_impl { + using type = typename T::value_type; +}; + +template +using VAL_T = typename val_t_impl::type; + +template +RAJA_HOST_DEVICE typename T::value_type get_val(T index_val) { return *index_val; } + +template +requires (!RAJA::concepts::IndexValued) +RAJA_HOST_DEVICE auto get_val(T index_val) { return index_val; } +} + template -void KernelTileLocalArray2DTestImpl(const int rows, const int cols) +void KernelTileLocalArray2DTestImpl(const INDEX_TYPE rows_t, const INDEX_TYPE cols_t) { // This test emulates matrix transposition with tiling. + using raw_index_type = VAL_T; camp::resources::Resource work_res{WORKING_RES::get_default()}; @@ -28,7 +51,7 @@ void KernelTileLocalArray2DTestImpl(const int rows, const int cols) DATA_TYPE * check_array_t; DATA_TYPE * test_array_t; - INDEX_TYPE array_length = rows * cols; + INDEX_TYPE array_length = rows_t * cols_t; allocateForallTestData ( array_length, work_res, @@ -44,53 +67,58 @@ void KernelTileLocalArray2DTestImpl(const int rows, const int cols) &test_array_t ); - RAJA::View> HostView( test_array, rows, cols ); - RAJA::View> HostTView( test_array_t, cols, rows ); - RAJA::View> WorkView( work_array, rows, cols ); - RAJA::View> WorkTView( work_array_t, cols, rows ); - RAJA::View> CheckTView( check_array_t, cols, rows ); + using LayoutType = + RAJA::TypedLayout>; + using ViewType = RAJA::View; + + ViewType HostView( test_array, rows_t, cols_t ); + ViewType HostTView( test_array_t, cols_t, rows_t ); + ViewType WorkView( work_array, rows_t, cols_t ); + ViewType WorkTView( work_array_t, cols_t, rows_t ); + ViewType CheckTView( check_array_t, cols_t, rows_t ); // initialize local array (shared mem) using TILE_MEM = RAJA::LocalArray, RAJA::SizeList>; TILE_MEM Tile_Array; // initialize arrays - std::iota( test_array, test_array + array_length, 1 ); - std::iota( test_array_t, test_array_t + array_length, 1 ); + std::iota( test_array, test_array + get_val(array_length), 1 ); + std::iota( test_array_t, test_array_t + get_val(array_length), 1 ); - work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * array_length ); - work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( work_array, test_array, sizeof(DATA_TYPE) * get_val(array_length) ); + work_res.memcpy( work_array_t, test_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); // transpose test_array on CPU - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - HostTView( cc, rr ) = HostView( rr, cc ); + HostTView( INDEX_TYPE(cc), INDEX_TYPE(rr) ) = HostView( INDEX_TYPE(rr), INDEX_TYPE(cc) ); } } // transpose work_array - RAJA::TypedRangeSegment rowrange( 0, rows ); - RAJA::TypedRangeSegment colrange( 0, cols ); + RAJA::TypedRangeSegment rowrange( 0, rows_t ); + RAJA::TypedRangeSegment colrange( 0, cols_t ); - RAJA::kernel_param ( RAJA::make_tuple( colrange, rowrange ), RAJA::make_tuple( (INDEX_TYPE)0, (INDEX_TYPE)0, Tile_Array ), + RAJA::kernel_param ( RAJA::make_tuple( colrange, rowrange ), RAJA::make_tuple( INDEX_TYPE(0), INDEX_TYPE(0), Tile_Array ), [=] RAJA_HOST_DEVICE ( INDEX_TYPE cc, INDEX_TYPE rr, INDEX_TYPE tx, INDEX_TYPE ty, TILE_MEM &_Tile_Array ) { - _Tile_Array( ty, tx ) = WorkView( rr, cc ); + _Tile_Array( get_val(ty), get_val(tx) ) = WorkView( rr, cc ); }, [=] RAJA_HOST_DEVICE ( INDEX_TYPE cc, INDEX_TYPE rr, INDEX_TYPE tx, INDEX_TYPE ty, TILE_MEM &_Tile_Array ) { - WorkTView( cc, rr ) = _Tile_Array( ty, tx ); + WorkTView( cc, rr ) = _Tile_Array( get_val(ty), get_val(tx) ); } ); - work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * array_length ); + work_res.memcpy( check_array_t, work_array_t, sizeof(DATA_TYPE) * get_val(array_length) ); - for ( int rr = 0; rr < rows; ++rr ) + for ( raw_index_type rr = 0; rr < get_val(rows_t); ++rr ) { - for ( int cc = 0; cc < cols; ++cc ) + for ( raw_index_type cc = 0; cc < get_val(cols_t); ++cc ) { - ASSERT_EQ(CheckTView(cc, rr), HostTView(cc, rr)); + ASSERT_EQ(CheckTView(INDEX_TYPE(cc), INDEX_TYPE(rr)), + HostTView(INDEX_TYPE(cc), INDEX_TYPE(rr))); } } @@ -121,9 +149,9 @@ TYPED_TEST_P(KernelTileLocalArray2DTest, TileLocalArray2DKernel) using WORKING_RES = typename camp::at>::type; using EXEC_POLICY = typename camp::at>::type; - KernelTileLocalArray2DTestImpl(10, 10); - KernelTileLocalArray2DTestImpl(151, 111); - KernelTileLocalArray2DTestImpl(362, 362); + KernelTileLocalArray2DTestImpl(INDEX_TYPE(10), INDEX_TYPE(10)); + KernelTileLocalArray2DTestImpl(INDEX_TYPE(151), INDEX_TYPE(111)); + KernelTileLocalArray2DTestImpl(INDEX_TYPE(362), INDEX_TYPE(362)); } REGISTER_TYPED_TEST_SUITE_P(KernelTileLocalArray2DTest, diff --git a/test/functional/launch/nested_direct/test-launch-nested.cpp.in b/test/functional/launch/nested_direct/test-launch-nested.cpp.in index ea92db5cdf..1a78907650 100644 --- a/test/functional/launch/nested_direct/test-launch-nested.cpp.in +++ b/test/functional/launch/nested_direct/test-launch-nested.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/nested_direct_unchecked/test-launch-nested.cpp.in b/test/functional/launch/nested_direct_unchecked/test-launch-nested.cpp.in index 628c7efb08..262924a2d8 100644 --- a/test/functional/launch/nested_direct_unchecked/test-launch-nested.cpp.in +++ b/test/functional/launch/nested_direct_unchecked/test-launch-nested.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/nested_loop/test-launch-nested.cpp.in b/test/functional/launch/nested_loop/test-launch-nested.cpp.in index cd2af8ee26..80d8d7b016 100644 --- a/test/functional/launch/nested_loop/test-launch-nested.cpp.in +++ b/test/functional/launch/nested_loop/test-launch-nested.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/nested_tile_direct/test-launch-nested-tile-direct.cpp.in b/test/functional/launch/nested_tile_direct/test-launch-nested-tile-direct.cpp.in index 59a1ea91a5..425d385bd1 100644 --- a/test/functional/launch/nested_tile_direct/test-launch-nested-tile-direct.cpp.in +++ b/test/functional/launch/nested_tile_direct/test-launch-nested-tile-direct.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/nested_tile_direct_unchecked/test-launch-nested-tile-direct-unchecked.cpp.in b/test/functional/launch/nested_tile_direct_unchecked/test-launch-nested-tile-direct-unchecked.cpp.in index 0238fc01a7..41d263a910 100644 --- a/test/functional/launch/nested_tile_direct_unchecked/test-launch-nested-tile-direct-unchecked.cpp.in +++ b/test/functional/launch/nested_tile_direct_unchecked/test-launch-nested-tile-direct-unchecked.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/nested_tile_loop/test-launch-nested-tile-loop.cpp.in b/test/functional/launch/nested_tile_loop/test-launch-nested-tile-loop.cpp.in index 5d063da3b7..1163dae07d 100644 --- a/test/functional/launch/nested_tile_loop/test-launch-nested-tile-loop.cpp.in +++ b/test/functional/launch/nested_tile_loop/test-launch-nested-tile-loop.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/tile_icount_tcount_direct/test-launch-nested-tile-icount-tcount-direct.cpp.in b/test/functional/launch/tile_icount_tcount_direct/test-launch-nested-tile-icount-tcount-direct.cpp.in index f42482f7b6..199e0f5971 100644 --- a/test/functional/launch/tile_icount_tcount_direct/test-launch-nested-tile-icount-tcount-direct.cpp.in +++ b/test/functional/launch/tile_icount_tcount_direct/test-launch-nested-tile-icount-tcount-direct.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/tile_icount_tcount_direct_unchecked/test-launch-nested-tile-icount-tcount-direct-unchecked.cpp.in b/test/functional/launch/tile_icount_tcount_direct_unchecked/test-launch-nested-tile-icount-tcount-direct-unchecked.cpp.in index 0ca15ad6b0..651ae43baf 100644 --- a/test/functional/launch/tile_icount_tcount_direct_unchecked/test-launch-nested-tile-icount-tcount-direct-unchecked.cpp.in +++ b/test/functional/launch/tile_icount_tcount_direct_unchecked/test-launch-nested-tile-icount-tcount-direct-unchecked.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/functional/launch/tile_icount_tcount_loop/test-launch-nested-tile-icount-tcount-loop.cpp.in b/test/functional/launch/tile_icount_tcount_loop/test-launch-nested-tile-icount-tcount-loop.cpp.in index bee00c999b..38456dfb09 100644 --- a/test/functional/launch/tile_icount_tcount_loop/test-launch-nested-tile-icount-tcount-loop.cpp.in +++ b/test/functional/launch/tile_icount_tcount_loop/test-launch-nested-tile-icount-tcount-loop.cpp.in @@ -29,7 +29,7 @@ // Cartesian product of types used in parameterized tests // using @BACKEND@LaunchNestedTypes = - Test< camp::cartesian_product>::Types; diff --git a/test/include/RAJA_test-forall-data.hpp b/test/include/RAJA_test-forall-data.hpp index 6c45b8f8ab..97c54b8cb1 100644 --- a/test/include/RAJA_test-forall-data.hpp +++ b/test/include/RAJA_test-forall-data.hpp @@ -8,12 +8,13 @@ //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// // -// Utility routines for allocating/deallocating arrays in for forall tests. +// Utility routines for allocating/deallocating arrays in for forall tests. // #ifndef __RAJA_test_forall_data_HPP__ #define __RAJA_test_forall_data_HPP__ +#include "RAJA/index/IndexValue.hpp" #include "camp/resource.hpp" template @@ -33,8 +34,8 @@ void allocateForallTestData(size_t N, // for RAJA strongly typed indices template>::value>::type* = nullptr> -void allocateForallTestData(T N, + RAJA::concepts::IndexValued IdxType> +void allocateForallTestData(IdxType N, camp::resources::Resource work_res, T** work_array, T** check_array, diff --git a/test/include/RAJA_test-index-types.hpp b/test/include/RAJA_test-index-types.hpp index bb8c6acbe9..c861051b66 100644 --- a/test/include/RAJA_test-index-types.hpp +++ b/test/include/RAJA_test-index-types.hpp @@ -27,44 +27,47 @@ // RAJA_INDEX_VALUE(StrongIndexType, "StrongIndexType"); RAJA_INDEX_VALUE_T(StrongInt, int, "StrongIntType"); -RAJA_INDEX_VALUE_T(StrongULL, unsigned long long , "StrongULLType"); +RAJA_INDEX_VALUE_T(StrongUL, unsigned long , "StrongULType"); // -// Standard index types list +// Raw index types list // -// TODO(bowen): do we want to support StrongULL here? -using IdxTypeList = camp::list; + unsigned long long>; // -// Signed index types list +// Raw signed index types list // -// TODO(bowen): do we want to support StrongInt here? -using SignedIdxTypeList = camp::list; +// Use this list for tests that require signed builtin integer semantics. +using RawSignedIdxTypeList = camp::list; // -// Index types w/ Strong types list +// Strong-compatible index types list // +// Use this list for tests that are expected to work with strongly typed +// indices and avoid raw integer interoperability assumptions. using StrongIdxTypeList = camp::list; #endif // __RAJA_test_index_types_HPP__ diff --git a/test/include/RAJA_test-reduceloc-types.hpp b/test/include/RAJA_test-reduceloc-types.hpp index 1109e545f4..c56e9d11e6 100644 --- a/test/include/RAJA_test-reduceloc-types.hpp +++ b/test/include/RAJA_test-reduceloc-types.hpp @@ -15,14 +15,16 @@ #define __RAJA_test_reduceloc_types_HPP__ #include "RAJA/RAJA.hpp" +#include "RAJA/index/IndexValue.hpp" #include "RAJA/util/types.hpp" #include "camp/list.hpp" +template struct Index2D { - RAJA::Index_type idx, idy; + IDX idx, idy; constexpr Index2D() : idx(-1), idy(-1) {} - constexpr Index2D(RAJA::Index_type init) : idx(init), idy(init) {} - constexpr Index2D(RAJA::Index_type ix, RAJA::Index_type iy) : idx(ix), idy(iy) {} + constexpr Index2D(IDX ix) : idx(ix), idy(ix) {} + constexpr Index2D(IDX ix, IDX iy) : idx(ix), idy(iy) {} template RAJA_HOST_DEVICE void operator=(T rhs) { idx = rhs; idx = rhs; } };