From d0bebe92a70031b6c891626c120c4fc07fd61384 Mon Sep 17 00:00:00 2001
From: Alyan <sah240607@gmail.com>
Date: Tue, 29 Sep 2026 18:27:11 -0400
Subject: [PATCH 1/2] Fixed Bug with assume_aligned not being made for device

---
 tests/memory/cases.hpp |  3 ++-
 tests/memory/cuda.cu   | 46 ++++++++++++++++++++++++++++++++++++++++++
 xpu/memory.hpp         |  6 +++++-
 3 files changed, 53 insertions(+), 2 deletions(-)

diff --git a/tests/memory/cases.hpp b/tests/memory/cases.hpp
index a2577e8..00b44f1 100644
--- a/tests/memory/cases.hpp
+++ b/tests/memory/cases.hpp
@@ -3,6 +3,7 @@
 #include "../support/check.hpp"
 
 #include <xpu/memory.hpp>
+#include <xpu/soa.hpp>
 
 #include <memory>
 
@@ -48,4 +49,4 @@ inline int run_memory_cases() {
   }
 
   return 0;
-}
+}
\ No newline at end of file
diff --git a/tests/memory/cuda.cu b/tests/memory/cuda.cu
index 0b7c87a..721191f 100644
--- a/tests/memory/cuda.cu
+++ b/tests/memory/cuda.cu
@@ -1,5 +1,51 @@
 #include "cases.hpp"
 
+namespace {
+
+__global__ void write_soa_view(xpu::soa_view<int, 2> view) {
+  const auto index{static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x};
+  if (index < view.count()) {
+    view[0][index] = static_cast<int>(index + 1uz);
+    view[1][index] = static_cast<int>(2uz * (index + 1uz));
+  }
+}
+
+__global__ void read_soa_view(const xpu::soa_view<int, 2> input, xpu::soa_view<int, 2> output) {
+  const auto index{static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x};
+  if (index < input.count()) {
+    output[0][index] = (2 * input[0][index]);
+    output[1][index] = input[1][index] + 1;
+  }
+}
+
+}
+
+auto run_device_soa_view_cases() -> int {
+  constexpr auto count{8uz};
+  auto values{xpu::soa<int, 2>{count}};
+  auto output{xpu::soa<int, 2>{count}};
+  write_soa_view<<<1, 32>>>(values.view());
+  xpu::cu_check(cudaGetLastError());
+  read_soa_view<<<1, 32>>>(values.view(), output.view());
+  xpu::cu_check(cudaGetLastError());
+  xpu::cu_check(cudaDeviceSynchronize());
+
+  int result[2][count]{};
+  xpu::copy_n(result[0], output.view()[0], count);
+  xpu::copy_n(result[1], output.view()[1], count);
+  for (auto i{0uz}; i < count; ++i) {
+    if (result[0][i] != static_cast<int>(2uz * (i + 1uz)) 
+    ||  result[1][i] != static_cast<int>((2uz * (i + 1uz)) + 1uz)) {
+      return test::fail("CUDA SoA View Indexing is incorrect");
+    }
+  }
+  return 0;
+}
+
 int main() {
+  if (const auto failure{run_device_soa_view_cases()}; failure) {
+    return failure;
+  }
+
   return run_memory_cases();
 }
diff --git a/xpu/memory.hpp b/xpu/memory.hpp
index 928999c..b3495bc 100644
--- a/xpu/memory.hpp
+++ b/xpu/memory.hpp
@@ -7,6 +7,10 @@
 #include <memory>
 #include <type_traits>
 
+#if defined(XPU_CUDA)
+  #include <cuda/std/memory>
+#endif
+
 namespace xpu {
 
 template <typename T>
@@ -94,7 +98,7 @@ auto make_unique(std::size_t count, T value = T{}) -> unique_ptr<T> {
 template <typename T> [[nodiscard]] XPU_CUDA_CALLABLE
 constexpr auto assume_aligned(T* ptr) noexcept -> T* {
   if constexpr (is_padded<T>) {
-    return std::assume_aligned<default_align<T>>(ptr);
+    return xstd::assume_aligned<default_align<T>>(ptr);
   } else {
     return ptr;
   }

From e3d13f3b2821d6d2e51c6250c9550799fb6d0e71 Mon Sep 17 00:00:00 2001
From: Alyan <sah240607@gmail.com>
Date: Wed, 30 Sep 2026 16:59:54 -0400
Subject: [PATCH 2/2] changed xpu::assume_aligned to be a wrapper

---
 xpu/memory.hpp | 6 +-----
 1 file changed, 1 insertion(+), 5 deletions(-)

diff --git a/xpu/memory.hpp b/xpu/memory.hpp
index b3495bc..4ce95cb 100644
--- a/xpu/memory.hpp
+++ b/xpu/memory.hpp
@@ -97,11 +97,7 @@ auto make_unique(std::size_t count, T value = T{}) -> unique_ptr<T> {
 
 template <typename T> [[nodiscard]] XPU_CUDA_CALLABLE
 constexpr auto assume_aligned(T* ptr) noexcept -> T* {
-  if constexpr (is_padded<T>) {
-    return xstd::assume_aligned<default_align<T>>(ptr);
-  } else {
-    return ptr;
-  }
+  return xstd::assume_aligned<default_align<T>>(ptr);
 }
 
 inline auto memset(