diff --git a/include/gpu_array.hpp b/include/gpu_array.hpp index 234f97e..aa56b8a 100644 --- a/include/gpu_array.hpp +++ b/include/gpu_array.hpp @@ -2816,12 +2816,6 @@ namespace gpu_array template struct stride_adapter { - template - requires std::ranges::sized_range - [[nodiscard]] constexpr auto operator()(const Range& r) const noexcept - { - return stride_view(r); - } template requires std::ranges::sized_range [[nodiscard]] constexpr auto operator()(Range& r) const noexcept @@ -2829,13 +2823,6 @@ namespace gpu_array return stride_view(r); } - template - requires std::ranges::sized_range - [[nodiscard]] friend constexpr std::ranges::view auto operator|(const Range& range, - const stride_adapter& self) noexcept - { - return self(range); - } template requires std::ranges::sized_range [[nodiscard]] friend constexpr std::ranges::view auto operator|(Range& range, @@ -2846,13 +2833,16 @@ namespace gpu_array }; } // namespace detail +#if !defined(ENABLE_HIP) + // The following three alias templates are also disabled in HIP because HIP does not support alias template argument + // deduction. template using block_thread_stride_view = detail::stride_view; template using grid_thread_stride_view = detail::stride_view; template using grid_block_stride_view = detail::stride_view; -#if !defined(ENABLE_HIP) + template using cluster_thread_stride_view = detail::stride_view; template diff --git a/test/test.cpp b/test/test.cpp index d49b74b..13cbc77 100644 --- a/test/test.cpp +++ b/test/test.cpp @@ -2070,6 +2070,7 @@ TEST(JaggedArray, MemoryManagement) } } +#if !defined(ENABLE_HIP) template requires std::ranges::input_range> __global__ void kernel_stride(T array) @@ -2086,14 +2087,6 @@ __global__ void kernel_stride2(T array) for (auto& v : views::block_thread_stride(a)) v = 2; } -template -requires std::ranges::input_range> -__global__ void kernel_stride3(T array) -{ - for (auto& a : grid_block_stride_view(array)) - for (auto& v : block_thread_stride_view(a)) v = 3; -} - TEST(StrideView, HowToUse) { auto vec_vec = std::vector(32, std::vector(64, 0)); @@ -2108,10 +2101,25 @@ TEST(StrideView, HowToUse) api::gpuDeviceSynchronize(); for (const auto& inner_array : nested_array) for (const auto& v : inner_array) EXPECT_EQ(v, 2); +} + +template +requires std::ranges::input_range> +__global__ void kernel_stride3(T array) +{ + for (auto& a : grid_block_stride_view(array)) + for (auto& v : block_thread_stride_view(a)) v = 3; +} + +TEST(StrideView, AliasTemplate) +{ + auto vec_vec = std::vector(32, std::vector(64, 0)); + auto nested_array = managed_array(vec_vec); kernel_stride3<<<32, 64>>>(nested_array); api::gpuDeviceSynchronize(); for (const auto& inner_array : nested_array) for (const auto& v : inner_array) EXPECT_EQ(v, 3); } +#endif // NOLINTEND