diff --git a/simd/src/Kokkos_SIMD_AVX512.hpp b/simd/src/Kokkos_SIMD_AVX512.hpp index 7d1acbc35ab..89c297bb0fe 100644 --- a/simd/src/Kokkos_SIMD_AVX512.hpp +++ b/simd/src/Kokkos_SIMD_AVX512.hpp @@ -3231,6 +3231,7 @@ class where_expression>, } }; +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int32_t hmax(const_where_expression< @@ -3240,6 +3241,7 @@ class where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int32_t reduce_max( const_where_expression< @@ -3250,6 +3252,7 @@ class where_expression>, _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int32_t hmin(const_where_expression< @@ -3259,6 +3262,7 @@ class where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int32_t reduce_min( const_where_expression< @@ -3285,6 +3289,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint32_t hmax(const_where_expression< @@ -3294,6 +3299,7 @@ class where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint32_t reduce_max( const_where_expression< @@ -3304,6 +3310,7 @@ class where_expression>, _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint32_t hmin(const_where_expression< @@ -3313,6 +3320,7 @@ class where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castsi256_si512(static_cast<__m256i>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint32_t reduce_min( const_where_expression< @@ -3339,6 +3347,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int64_t hmax(const_where_expression< @@ -3347,6 +3356,7 @@ class where_expression>, return _mm512_mask_reduce_max_epi64(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512i>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int64_t reduce_max( const_where_expression< @@ -3356,6 +3366,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int64_t hmin(const_where_expression< @@ -3364,6 +3375,7 @@ class where_expression>, return _mm512_mask_reduce_min_epi64(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512i>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::int64_t reduce_min( const_where_expression< @@ -3373,6 +3385,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint64_t hmax(const_where_expression< @@ -3381,6 +3394,7 @@ class where_expression>, return _mm512_mask_reduce_max_epu64(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512i>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint64_t reduce_max( const_where_expression< @@ -3390,6 +3404,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint64_t hmin(const_where_expression< @@ -3398,6 +3413,7 @@ class where_expression>, return _mm512_mask_reduce_min_epu64(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512i>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION std::uint64_t reduce_min( const_where_expression< @@ -3407,6 +3423,7 @@ class where_expression>, static_cast<__m512i>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION double hmax(const_where_expression>, simd>> const& @@ -3414,6 +3431,7 @@ hmax(const_where_expression>, return _mm512_mask_reduce_max_pd(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512d>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION double reduce_max( const_where_expression>, @@ -3423,6 +3441,7 @@ hmax(const_where_expression>, static_cast<__m512d>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION double hmin(const_where_expression>, simd>> const& @@ -3430,6 +3449,7 @@ hmin(const_where_expression>, return _mm512_mask_reduce_min_pd(static_cast<__mmask8>(x.impl_get_mask()), static_cast<__m512d>(x.impl_get_value())); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION double reduce_min( const_where_expression>, @@ -3439,6 +3459,7 @@ hmin(const_where_expression>, static_cast<__m512d>(x.impl_get_value())); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION float hmax(const_where_expression>, simd>> const& @@ -3447,6 +3468,7 @@ hmax(const_where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castps256_ps512(static_cast<__m256>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION float reduce_max( const_where_expression>, @@ -3457,6 +3479,7 @@ hmax(const_where_expression>, _mm512_castps256_ps512(static_cast<__m256>(x.impl_get_value()))); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION float hmin(const_where_expression>, simd>> const& @@ -3465,6 +3488,7 @@ hmin(const_where_expression>, static_cast<__mmask8>(x.impl_get_mask()), _mm512_castps256_ps512(static_cast<__m256>(x.impl_get_value()))); } +#endif [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION float reduce_min( const_where_expression>, diff --git a/simd/src/Kokkos_SIMD_Common_Math.hpp b/simd/src/Kokkos_SIMD_Common_Math.hpp index 2e78e2700fb..8b96bc1dd9d 100644 --- a/simd/src/Kokkos_SIMD_Common_Math.hpp +++ b/simd/src/Kokkos_SIMD_Common_Math.hpp @@ -32,6 +32,7 @@ class simd_mask; template class const_where_expression; +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 template [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION T hmin(const_where_expression, simd> const& x) { @@ -55,6 +56,7 @@ hmax(const_where_expression, simd> const& x) { } return result; } +#endif template [[nodiscard]] KOKKOS_IMPL_HOST_FORCEINLINE_FUNCTION T diff --git a/simd/src/Kokkos_SIMD_Scalar.hpp b/simd/src/Kokkos_SIMD_Scalar.hpp index 62b2e14e994..e2c82f0eb8c 100644 --- a/simd/src/Kokkos_SIMD_Scalar.hpp +++ b/simd/src/Kokkos_SIMD_Scalar.hpp @@ -426,6 +426,7 @@ reduce(const_where_expression, : identity_element; } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 template [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_FORCEINLINE_FUNCTION T hmax(const_where_expression, @@ -434,6 +435,7 @@ hmax(const_where_expression, ? static_cast(x.impl_get_value()) : Kokkos::reduction_identity::max(); } +#endif template [[nodiscard]] KOKKOS_FORCEINLINE_FUNCTION T @@ -444,6 +446,7 @@ reduce_max(const_where_expression, : Kokkos::reduction_identity::max(); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4 template [[nodiscard]] KOKKOS_DEPRECATED KOKKOS_FORCEINLINE_FUNCTION T hmin(const_where_expression, @@ -452,6 +455,7 @@ hmin(const_where_expression, ? static_cast(x.impl_get_value()) : Kokkos::reduction_identity::min(); } +#endif template [[nodiscard]] KOKKOS_FORCEINLINE_FUNCTION T diff --git a/simd/unit_tests/include/SIMDTesting_Ops.hpp b/simd/unit_tests/include/SIMDTesting_Ops.hpp index 6103efd90c2..3ce00ddefff 100644 --- a/simd/unit_tests/include/SIMDTesting_Ops.hpp +++ b/simd/unit_tests/include/SIMDTesting_Ops.hpp @@ -335,12 +335,12 @@ template > class reduce_where_expr { public: template - KOKKOS_INLINE_FUNCTION auto on_host(T const& a, MaskType mask) const { + auto on_host(T const& a, MaskType mask) const { auto w = Kokkos::Experimental::where(mask, a); return Kokkos::Experimental::reduce(w, BinaryOperation()); } template - KOKKOS_INLINE_FUNCTION auto on_host_serial(T const& a, MaskType mask) const { + auto on_host_serial(T const& a, MaskType mask) const { auto w = Kokkos::Experimental::where(mask, a); auto const& v = w.impl_get_value(); auto const& m = w.impl_get_mask(); @@ -373,11 +373,11 @@ class reduce_where_expr { class reduce_min { public: template - KOKKOS_INLINE_FUNCTION auto on_host(T const& a, MaskType mask) const { + auto on_host(T const& a, MaskType mask) const { return Kokkos::Experimental::reduce_min(a, mask); } template - KOKKOS_INLINE_FUNCTION auto on_host_serial(T const& a, MaskType mask) const { + auto on_host_serial(T const& a, MaskType mask) const { auto w = Kokkos::Experimental::where(mask, a); auto const& v = w.impl_get_value(); auto const& m = w.impl_get_mask(); @@ -409,11 +409,11 @@ class reduce_min { class reduce_max { public: template - KOKKOS_INLINE_FUNCTION auto on_host(T const& a, MaskType mask) const { + auto on_host(T const& a, MaskType mask) const { return Kokkos::Experimental::reduce_max(a, mask); } template - KOKKOS_INLINE_FUNCTION auto on_host_serial(T const& a, MaskType mask) const { + auto on_host_serial(T const& a, MaskType mask) const { auto w = Kokkos::Experimental::where(mask, a); auto const& v = w.impl_get_value(); auto const& m = w.impl_get_mask(); @@ -446,11 +446,11 @@ template > class reduce { public: template - KOKKOS_INLINE_FUNCTION auto on_host(T const& a, MaskType mask) const { + auto on_host(T const& a, MaskType mask) const { return Kokkos::Experimental::reduce(a, mask, BinaryOperation()); } template - KOKKOS_INLINE_FUNCTION auto on_host_serial(T const& a, MaskType mask) const { + auto on_host_serial(T const& a, MaskType mask) const { return reduce_where_expr().on_host_serial(a, mask); }