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..4ce95cb 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>
@@ -93,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 std::assume_aligned<default_align<T>>(ptr);
-  } else {
-    return ptr;
-  }
+  return xstd::assume_aligned<default_align<T>>(ptr);
 }
 
 inline auto memset(