From ee32c00cd49130f018d1db91e1d6dacd2e2718d1 Mon Sep 17 00:00:00 2001
From: acd1034 <48613285+acd1034@users.noreply.github.com>
Date: Fri, 13 Feb 2026 15:34:09 +0900
Subject: [PATCH 1/3] =?UTF-8?q?=F0=9F=90=9B=20Disable=20alias=20templates?=
 =?UTF-8?q?=20for=20stride=20views=20in=20HIP?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 include/gpu_array.hpp |  5 ++++-
 test/test.cpp         | 24 ++++++++++++++++--------
 2 files changed, 20 insertions(+), 9 deletions(-)

diff --git a/include/gpu_array.hpp b/include/gpu_array.hpp
index 234f97e..ba7aa09 100644
--- a/include/gpu_array.hpp
+++ b/include/gpu_array.hpp
@@ -2846,13 +2846,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 <std::ranges::random_access_range Range>
     using block_thread_stride_view = detail::stride_view<detail::Stride::BlockThread, Range>;
     template <std::ranges::random_access_range Range>
     using grid_thread_stride_view = detail::stride_view<detail::Stride::GridThread, Range>;
     template <std::ranges::random_access_range Range>
     using grid_block_stride_view = detail::stride_view<detail::Stride::GridBlock, Range>;
-#if !defined(ENABLE_HIP)
+
     template <std::ranges::random_access_range Range>
     using cluster_thread_stride_view = detail::stride_view<detail::Stride::ClusterThread, Range>;
     template <std::ranges::random_access_range Range>
diff --git a/test/test.cpp b/test/test.cpp
index d49b74b..217f445 100644
--- a/test/test.cpp
+++ b/test/test.cpp
@@ -2086,14 +2086,6 @@ __global__ void kernel_stride2(T array)
         for (auto& v : views::block_thread_stride(a)) v = 2;
 }
 
-template <std::ranges::input_range T>
-requires std::ranges::input_range<std::ranges::range_value_t<T>>
-__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<int>(64, 0));
@@ -2108,10 +2100,26 @@ TEST(StrideView, HowToUse)
     api::gpuDeviceSynchronize();
     for (const auto& inner_array : nested_array)
         for (const auto& v : inner_array) EXPECT_EQ(v, 2);
+}
+
+#if !defined(ENABLE_HIP)
+template <std::ranges::input_range T>
+requires std::ranges::input_range<std::ranges::range_value_t<T>>
+__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<int>(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

From a84bea26ffa4e2dd6a91e083a5a942a3c4529aa9 Mon Sep 17 00:00:00 2001
From: acd1034 <48613285+acd1034@users.noreply.github.com>
Date: Fri, 13 Feb 2026 15:44:17 +0900
Subject: [PATCH 2/3] =?UTF-8?q?=F0=9F=94=A5=20Remove=20overloads=20for=20c?=
 =?UTF-8?q?onst=20Range&=20in=20stride=5Fadapter?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 include/gpu_array.hpp | 13 -------------
 1 file changed, 13 deletions(-)

diff --git a/include/gpu_array.hpp b/include/gpu_array.hpp
index ba7aa09..aa56b8a 100644
--- a/include/gpu_array.hpp
+++ b/include/gpu_array.hpp
@@ -2816,12 +2816,6 @@ namespace gpu_array
         template <Stride StrideType>
         struct stride_adapter
         {
-            template <std::ranges::random_access_range Range>
-            requires std::ranges::sized_range<Range>
-            [[nodiscard]] constexpr auto operator()(const Range& r) const noexcept
-            {
-                return stride_view<StrideType, const Range&>(r);
-            }
             template <std::ranges::random_access_range Range>
             requires std::ranges::sized_range<Range>
             [[nodiscard]] constexpr auto operator()(Range& r) const noexcept
@@ -2829,13 +2823,6 @@ namespace gpu_array
                 return stride_view<StrideType, Range&>(r);
             }
 
-            template <std::ranges::random_access_range Range>
-            requires std::ranges::sized_range<Range>
-            [[nodiscard]] friend constexpr std::ranges::view auto operator|(const Range& range,
-                                                                            const stride_adapter& self) noexcept
-            {
-                return self(range);
-            }
             template <std::ranges::random_access_range Range>
             requires std::ranges::sized_range<Range>
             [[nodiscard]] friend constexpr std::ranges::view auto operator|(Range& range,

From 92f90f2030505ebd04aa461e682feb659cf3282e Mon Sep 17 00:00:00 2001
From: acd1034 <48613285+acd1034@users.noreply.github.com>
Date: Fri, 13 Feb 2026 16:58:15 +0900
Subject: [PATCH 3/3] =?UTF-8?q?=F0=9F=9A=A8=20Disable=20kernel=20testing?=
 =?UTF-8?q?=20in=20HIP?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 test/test.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/test/test.cpp b/test/test.cpp
index 217f445..13cbc77 100644
--- a/test/test.cpp
+++ b/test/test.cpp
@@ -2070,6 +2070,7 @@ TEST(JaggedArray, MemoryManagement)
     }
 }
 
+#if !defined(ENABLE_HIP)
 template <std::ranges::input_range T>
 requires std::ranges::input_range<std::ranges::range_value_t<T>>
 __global__ void kernel_stride(T array)
@@ -2102,7 +2103,6 @@ TEST(StrideView, HowToUse)
         for (const auto& v : inner_array) EXPECT_EQ(v, 2);
 }
 
-#if !defined(ENABLE_HIP)
 template <std::ranges::input_range T>
 requires std::ranges::input_range<std::ranges::range_value_t<T>>
 __global__ void kernel_stride3(T array)