From 0452e6b58066a7f9adec4bc9a20fcd0f9ce1b0ae Mon Sep 17 00:00:00 2001 From: shsahiti Date: Thu, 17 Sep 2026 10:10:16 -0400 Subject: [PATCH 1/2] Tidy sincos, memory, random, and soa headers --- xpu/math.hpp | 13 +++++++------ xpu/memory.hpp | 48 +++++++++++++++++++++++------------------------- xpu/random.hpp | 23 +++++++++++++---------- xpu/soa.hpp | 14 +++++++------- 4 files changed, 50 insertions(+), 48 deletions(-) diff --git a/xpu/math.hpp b/xpu/math.hpp index 4c87716..91f27c5 100644 --- a/xpu/math.hpp +++ b/xpu/math.hpp @@ -67,25 +67,26 @@ using xstd::ldexp; using xstd::frexp; using xstd::modf; // Overflow safe version of (num + den - 1) / den template [[nodiscard]] CUDA_CALLABLE -inline constexpr auto ceiling_div(T num, T den) noexcept -> T { +constexpr auto ceiling_div(T num, T den) noexcept -> T { const auto zero_denominator{den == 0}; if (zero_denominator) { xpu::detail::checked_error("zero ceiling division denominator"); } - const auto quotient{num / den + (num % den != 0)}; + const auto quotient{ ( num / den ) + (num % den != 0)}; return quotient; } + template CUDA_CALLABLE -inline auto sincos(T arg, T* RESTRICT s, T* RESTRICT c) noexcept -> void { +inline auto sincos(T arg, T* RESTRICT sin, T* RESTRICT cos) noexcept -> void { #if defined(__CUDA_ARCH__) - if constexpr (std::is_same_v) { ::sincosf(arg, s, c); } - else { ::sincos(arg, s, c); } + if constexpr (std::is_same_v) { ::sincosf(arg, sin, cos); } + else { ::sincos(arg, sin, cos); } #else - *s = xpu::sin(arg); *c = xpu::cos(arg); + *sin = xpu::sin(arg); *cos = xpu::cos(arg); #endif } diff --git a/xpu/memory.hpp b/xpu/memory.hpp index 724a6f9..e9561a2 100644 --- a/xpu/memory.hpp +++ b/xpu/memory.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include #include @@ -10,13 +11,13 @@ namespace xpu { template -inline constexpr auto default_align{(xpu::alignment_bytes > alignof(T)) ? xpu::alignment_bytes : alignof(T)}; +constexpr auto default_align{(xpu::alignment_bytes > alignof(T)) ? xpu::alignment_bytes : alignof(T)}; template -inline constexpr auto is_padded{sizeof(T) < xpu::alignment_bytes}; +constexpr auto is_padded{sizeof(T) < xpu::alignment_bytes}; template [[nodiscard]] CUDA_CALLABLE -inline constexpr auto bytes(std::size_t count) noexcept -> std::size_t { + constexpr auto bytes(std::size_t count) noexcept -> std::size_t { static_assert(std::is_trivially_copyable_v, "ERROR: xpu::bytes requires a trivially copyable type"); const auto byte_count{xpu::detail::checked_bytes(count)}; @@ -24,7 +25,7 @@ inline constexpr auto bytes(std::size_t count) noexcept -> std::size_t { } template [[nodiscard]] CUDA_CALLABLE -inline constexpr auto handle_pad(std::size_t unpadded) noexcept -> std::size_t { + constexpr auto handle_pad(std::size_t unpadded) noexcept -> std::size_t { if constexpr (is_padded) { constexpr auto lanes{xpu::alignment_bytes / sizeof(T)}; const auto padded_count{xpu::detail::checked_round_up(unpadded, lanes)}; @@ -36,24 +37,22 @@ inline constexpr auto handle_pad(std::size_t unpadded) noexcept -> std::size_t { } template [[nodiscard]] -inline auto alloc(std::size_t count) -> T* { +auto alloc(std::size_t count) -> T* { static_assert(std::is_trivially_copyable_v); if (count == 0uz) { return nullptr; } - const auto bytes{xpu::detail::checked_bytes(count)}; - #if defined(XPU_CUDA) auto ptr{static_cast(nullptr)}; - if(cudaMalloc(&ptr, bytes) != cudaSuccess) { ptr = nullptr; } + if(cudaMalloc(&ptr, bytes(count)) != cudaSuccess) { ptr = nullptr; } #else - auto ptr{::operator new(bytes, std::align_val_t{default_align}, std::nothrow)}; + auto ptr{::operator new(bytes(count), std::align_val_t{default_align}, std::nothrow)}; #endif if (!ptr) { std::fprintf( stderr, - "xpu: failed to allocate %zu bytes\n", - bytes + "xpu: failed to allocate {:d} bytes", + bytes(count) ); std::abort(); } @@ -62,7 +61,7 @@ inline auto alloc(std::size_t count) -> T* { } template -inline auto free(T* ptr) noexcept -> void { +auto free(T* ptr) noexcept -> void { if (!ptr) { return; } #if defined(XPU_CUDA) @@ -82,6 +81,11 @@ struct deleter { template using unique_ptr = std::unique_ptr>; +template +auto make_unique(T value = T{}) -> unique_ptr { + return xpu::make_unique(1uz, std::move(value)); +} + template auto make_unique(std::size_t count, T value = T{}) -> unique_ptr { auto* RESTRICT ptr{xpu::alloc(count)}; @@ -93,13 +97,11 @@ auto make_unique(std::size_t count, T value = T{}) -> unique_ptr { template [[nodiscard]] CUDA_CALLABLE constexpr auto assume_aligned(T* ptr) noexcept -> T* { - if constexpr (is_padded) { - return std::assume_aligned>(ptr); - } else { - return ptr; - } + if constexpr (!is_padded) { return ptr; } + return std::assume_aligned>(ptr); } + inline auto memset( void* RESTRICT dst, int value, @@ -129,7 +131,7 @@ inline auto memcpy( } template -inline auto copy_n( +auto copy_n( T* RESTRICT dst, const T* RESTRICT src, std::size_t count @@ -139,13 +141,11 @@ inline auto copy_n( "ERROR: xpu::copy_n requires trivially copyable type" ); - const auto byte_count{xpu::detail::checked_bytes(count)}; - - xpu::memcpy(dst, src, byte_count); + xpu::memcpy(dst, src, bytes(count)); } template -inline auto zero_n( +auto zero_n( T* RESTRICT dst, std::size_t count ) noexcept -> void { @@ -154,9 +154,7 @@ inline auto zero_n( "ERROR: xpu::zero_n requires arithmetic type" ); - const auto byte_count{xpu::detail::checked_bytes(count)}; - - xpu::memset(dst, 0, byte_count); + xpu::memset(dst, 0, bytes(count)); } } // namespace xpu diff --git a/xpu/random.hpp b/xpu/random.hpp index a375ca6..dc6991a 100644 --- a/xpu/random.hpp +++ b/xpu/random.hpp @@ -15,12 +15,15 @@ #include #include -namespace xpu { -namespace random { + +namespace xpu::random { class generator { private: + +static constexpr auto uint32_MAX{std::numeric_limits::digits}; + #if defined(XPU_CUDA) curandStatePhilox4_32_10_t engine_; @@ -51,9 +54,9 @@ class generator { #else auto seed{std::seed_seq{ static_cast(master_seed), - static_cast(master_seed >> 32u), + static_cast(master_seed >> uint32_MAX), static_cast(stream_id), - static_cast(stream_id >> 32u) + static_cast(stream_id >> uint32_MAX) }}; engine_.seed(seed); engine_.discard(offset); @@ -76,7 +79,7 @@ class generator { template [[nodiscard]] DEVICE_ONLY auto uniform(T minimum, T maximum) -> T { assert(minimum < maximum); - const auto value{minimum + (maximum - minimum) * uniform()}; + const auto value{minimum + ((maximum - minimum) * uniform())}; if (value < maximum) { return value; } #if defined(XPU_CUDA) @@ -139,13 +142,13 @@ inline auto seed_n( std::uint64_t offset = 0 ) -> void { const auto range = xpu::range<1uz>{ - {0uz}, - {count}, - {1uz} + .begin={0uz}, + .end={count}, + .step={1uz} }; const auto seed = detail::seed_generators{ - generators, stream_ids, master_seed, offset + .generators=generators, .stream_ids=stream_ids, .master_seed=master_seed, .offset=offset }; xpu::parallel_for(range, seed); @@ -153,4 +156,4 @@ inline auto seed_n( } // namespace xpu::random -} // namespace xpu + diff --git a/xpu/soa.hpp b/xpu/soa.hpp index daab260..ce7eff9 100644 --- a/xpu/soa.hpp +++ b/xpu/soa.hpp @@ -27,7 +27,7 @@ class soa_view { , count_{count} { const auto storage_count{xpu::detail::checked_mul(exposed_arrays, stride())}; - static_cast(xpu::detail::checked_bytes(storage_count)); + static_cast(bytes(storage_count)); } [[nodiscard]] CUDA_CALLABLE @@ -50,7 +50,7 @@ class soa_view { >; return xpu::assume_aligned( - self.base_ + arr_idx * self.stride() + self.base_ + (arr_idx * self.stride()) ); } @@ -118,7 +118,7 @@ class soa_batch_view { >; return xpu::soa_view{ - self.base_ + batch * self.batch_stride(), self.count_ + self.base_ + (batch * self.batch_stride()), self.count_ }; } }; @@ -167,7 +167,7 @@ class soa { >; return xpu::assume_aligned( - self.buffer_.data() + arr_idx * self.stride() + self.buffer_.data() + (arr_idx * self.stride()) ); } @@ -185,7 +185,7 @@ class soa { >; return soa_view{ - self.buffer_.data() + first_array * self.stride(), self.count_ + self.buffer_.data() + (first_array * self.stride()), self.count_ }; } }; @@ -261,7 +261,7 @@ class soa_batch { >; return xpu::soa_batch_view{ - self.storage_.data() + first_array * self.array_stride(), + self.storage_.data() + (first_array * self.array_stride()), self.batches_, self.count_ }; } @@ -283,7 +283,7 @@ class soa_batch { >; return xpu::soa_view{ - self.storage_.data() + batch * self.batch_stride() + first_array * self.array_stride(), + self.storage_.data() + (batch * self.batch_stride()) + (first_array * self.array_stride()), self.element_count() }; } From 9515b956885ad282ca3ba6c9eddcd2cf18d442f6 Mon Sep 17 00:00:00 2001 From: shsahiti Date: Sat, 19 Sep 2026 19:52:14 -0400 Subject: [PATCH 2/2] Fix namespace qualification and drop redundant xpu:: prefixes --- xpu/linear_algebra.hpp | 68 +++++++++++++++++++++--------------------- xpu/random.hpp | 4 +-- xpu/soa.hpp | 40 ++++++++++++------------- 3 files changed, 56 insertions(+), 56 deletions(-) diff --git a/xpu/linear_algebra.hpp b/xpu/linear_algebra.hpp index af246b4..4d1a25b 100644 --- a/xpu/linear_algebra.hpp +++ b/xpu/linear_algebra.hpp @@ -48,7 +48,7 @@ struct build_identity { std::size_t stride; DEVICE_ONLY - auto operator()(const xpu::array& index) const -> void { + auto operator()(const array& index) const -> void { const auto row{index[0]}; const auto column{index[1]}; @@ -66,7 +66,7 @@ struct transpose_square { std::size_t stride; DEVICE_ONLY - auto operator()(const xpu::array& index) const -> void { + auto operator()(const array& index) const -> void { const auto row{index[0]}; const auto column{index[1]}; const auto skip_pair{row >= column}; @@ -87,7 +87,7 @@ struct transpose_square { inline auto create_cusolver_handle() -> cusolverDnHandle_t { auto handle{cusolverDnHandle_t{}}; - xpu::cu_check(cusolverDnCreate(&handle)); + cu_check(cusolverDnCreate(&handle)); return handle; } @@ -103,14 +103,14 @@ inline auto getrf_workspace_size( auto size{0}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrf_bufferSize( + cu_check(cusolverDnSgetrf_bufferSize( handle, vendor_order, vendor_order, nullptr, vendor_stride, &size )); } else { - xpu::cu_check(cusolverDnDgetrf_bufferSize( + cu_check(cusolverDnDgetrf_bufferSize( handle, vendor_order, vendor_order, nullptr, vendor_stride, @@ -135,14 +135,14 @@ inline auto potrf_workspace_size( auto size{0}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrf_bufferSize( + cu_check(cusolverDnSpotrf_bufferSize( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, nullptr, vendor_stride, &size )); } else { - xpu::cu_check(cusolverDnDpotrf_bufferSize( + cu_check(cusolverDnDpotrf_bufferSize( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, nullptr, vendor_stride, @@ -169,14 +169,14 @@ inline auto cusolver_getrf( const auto vendor_stride{xpu::detail::checked_cast(stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrf( + cu_check(cusolverDnSgetrf( handle, vendor_order, vendor_order, matrix, vendor_stride, workspace, pivot, info )); } else { - xpu::cu_check(cusolverDnDgetrf( + cu_check(cusolverDnDgetrf( handle, vendor_order, vendor_order, matrix, vendor_stride, @@ -200,7 +200,7 @@ inline auto cusolver_potrf( const auto vendor_workspace_size{xpu::detail::checked_cast(workspace_size)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrf( + cu_check(cusolverDnSpotrf( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, matrix, vendor_stride, @@ -208,7 +208,7 @@ inline auto cusolver_potrf( info )); } else { - xpu::cu_check(cusolverDnDpotrf( + cu_check(cusolverDnDpotrf( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, matrix, vendor_stride, @@ -237,7 +237,7 @@ inline auto cusolver_getrs( const auto vendor_solution_stride{xpu::detail::checked_cast(solution_stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrs( + cu_check(cusolverDnSgetrs( handle, operation, vendor_order, vendor_right_hand_sides, lower_upper, vendor_lower_upper_stride, @@ -246,7 +246,7 @@ inline auto cusolver_getrs( info )); } else { - xpu::cu_check(cusolverDnDgetrs( + cu_check(cusolverDnDgetrs( handle, operation, vendor_order, vendor_right_hand_sides, lower_upper, vendor_lower_upper_stride, @@ -274,7 +274,7 @@ inline auto cusolver_potrs( const auto vendor_solution_stride{xpu::detail::checked_cast(solution_stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrs( + cu_check(cusolverDnSpotrs( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, vendor_right_hand_sides, factor, vendor_factor_stride, @@ -282,7 +282,7 @@ inline auto cusolver_potrs( info )); } else { - xpu::cu_check(cusolverDnDpotrs( + cu_check(cusolverDnDpotrs( handle, CUBLAS_FILL_MODE_UPPER, vendor_order, vendor_right_hand_sides, factor, vendor_factor_stride, @@ -468,7 +468,7 @@ inline auto transpose_square( stride }; - xpu::parallel_for(range, transpose); + parallel_for(range, transpose); #else for (auto row{0uz}; row < order; ++row) { for (auto column{row + 1uz}; column < order; ++column) { @@ -499,8 +499,8 @@ class cholesky_factorization { #if defined(XPU_CUDA) cusolverDnHandle_t handle_; - xpu::buffer workspace_; - xpu::buffer info_; + buffer workspace_; + buffer info_; #endif [[nodiscard]] @@ -534,7 +534,7 @@ class cholesky_factorization { ~cholesky_factorization() { #if defined(XPU_CUDA) - xpu::cu_check(cusolverDnDestroy(handle_)); + cu_check(cusolverDnDestroy(handle_)); #endif } @@ -559,7 +559,7 @@ class cholesky_factorization { ); auto info{0}; - xpu::copy_n(&info, info_.data(), 1uz); + copy_n(&info, info_.data(), 1uz); if (info < 0) { detail::linalg_error("cuSOLVER potrf received an invalid argument"); } @@ -595,7 +595,7 @@ class cholesky_factorization { detail::linalg_error("rhs and solution must not alias"); } - xpu::copy_n(solution, rhs, order_); + copy_n(solution, rhs, order_); #if defined(XPU_CUDA) detail::cusolver_potrs( @@ -606,7 +606,7 @@ class cholesky_factorization { ); auto info{0}; - xpu::copy_n(&info, info_.data(), 1uz); + copy_n(&info, info_.data(), 1uz); if (info != 0) { detail::linalg_error("cuSOLVER potrs received an invalid argument"); } @@ -640,12 +640,12 @@ class lu_factorization { #endif #if defined(XPU_CUDA) - xpu::buffer pivot_; + buffer pivot_; cusolverDnHandle_t handle_; - xpu::buffer workspace_; - xpu::buffer info_; + buffer workspace_; + buffer info_; #else - xpu::buffer pivot_; + buffer pivot_; #endif [[nodiscard]] @@ -661,7 +661,7 @@ class lu_factorization { const auto matrix_size{xpu::detail::checked_mul(order, stride)}; - static_cast(xpu::detail::checked_bytes(matrix_size)); + static_cast(bytes(matrix_size)); return order; } @@ -682,7 +682,7 @@ class lu_factorization { ~lu_factorization() { #if defined(XPU_CUDA) - xpu::cu_check(cusolverDnDestroy(handle_)); + cu_check(cusolverDnDestroy(handle_)); #endif } @@ -707,7 +707,7 @@ class lu_factorization { ); auto info{0}; - xpu::copy_n(&info, info_.data(), 1uz); + copy_n(&info, info_.data(), 1uz); if (info < 0) { detail::linalg_error("cuSOLVER getrf received an invalid argument"); } @@ -742,7 +742,7 @@ class lu_factorization { detail::linalg_error("rhs and solution must not alias"); } - xpu::copy_n(solution, rhs, order_); + copy_n(solution, rhs, order_); #if defined(XPU_CUDA) detail::cusolver_getrs( @@ -755,7 +755,7 @@ class lu_factorization { ); auto info{0}; - xpu::copy_n(&info, info_.data(), 1uz); + copy_n(&info, info_.data(), 1uz); if (info != 0) { detail::linalg_error("cuSOLVER getrs received an invalid argument"); } @@ -797,7 +797,7 @@ class lu_factorization { stride_ }; - xpu::parallel_for(range, initialize); + parallel_for(range, initialize); detail::cusolver_getrs( handle_, CUBLAS_OP_N, @@ -809,13 +809,13 @@ class lu_factorization { ); auto info{0}; - xpu::copy_n(&info, info_.data(), 1uz); + copy_n(&info, info_.data(), 1uz); if (info != 0) { detail::linalg_error("cuSOLVER getrs received an invalid argument"); } #else for (auto row{0uz}; row < order_; ++row) { - xpu::copy_n( + copy_n( inverse + row * stride_, lower_upper + row * stride_, order_ diff --git a/xpu/random.hpp b/xpu/random.hpp index dc6991a..08d6617 100644 --- a/xpu/random.hpp +++ b/xpu/random.hpp @@ -125,7 +125,7 @@ struct seed_generators { std::uint64_t offset; DEVICE_ONLY - auto operator()(const xpu::array& index) const -> void { + auto operator()(const array& index) const -> void { const auto i{index[0]}; generators[i].seed(master_seed, stream_ids[i], offset); @@ -151,7 +151,7 @@ inline auto seed_n( .generators=generators, .stream_ids=stream_ids, .master_seed=master_seed, .offset=offset }; - xpu::parallel_for(range, seed); + parallel_for(range, seed); } } // namespace xpu::random diff --git a/xpu/soa.hpp b/xpu/soa.hpp index ce7eff9..45e6083 100644 --- a/xpu/soa.hpp +++ b/xpu/soa.hpp @@ -26,7 +26,7 @@ class soa_view { : base_{base} , count_{count} { - const auto storage_count{xpu::detail::checked_mul(exposed_arrays, stride())}; + const auto storage_count{detail::checked_mul(exposed_arrays, stride())}; static_cast(bytes(storage_count)); } @@ -37,7 +37,7 @@ class soa_view { [[nodiscard]] CUDA_CALLABLE constexpr auto stride() const noexcept -> std::size_t { - return xpu::handle_pad(count_); + return handle_pad(count_); } template [[nodiscard]] CUDA_CALLABLE @@ -49,14 +49,14 @@ class soa_view { const T, T >; - return xpu::assume_aligned( + return assume_aligned( self.base_ + (arr_idx * self.stride()) ); } template [[nodiscard]] CUDA_CALLABLE constexpr auto pointers(this Self&& self) noexcept { - auto ptrs{xpu::array{}}; + auto ptrs{array{}}; for (auto i{0uz}; i < exposed_arrays; ++i) { ptrs[i] = self[i]; @@ -100,7 +100,7 @@ class soa_batch_view { [[nodiscard]] CUDA_CALLABLE constexpr auto array_stride() const noexcept -> std::size_t { - return xpu::handle_pad(count_); + return handle_pad(count_); } [[nodiscard]] CUDA_CALLABLE @@ -117,7 +117,7 @@ class soa_batch_view { const T, T >; - return xpu::soa_view{ + return soa_view{ self.base_ + (batch * self.batch_stride()), self.count_ }; } @@ -129,15 +129,15 @@ class soa { private: std::size_t count_; - xpu::buffer buffer_; + buffer buffer_; public: explicit soa(std::size_t count) noexcept : count_{count} , buffer_{ - xpu::detail::checked_mul( + detail::checked_mul( num_arrays, - xpu::handle_pad(count) + handle_pad(count) ) } { } @@ -149,7 +149,7 @@ class soa { [[nodiscard]] constexpr auto stride() const noexcept -> std::size_t { - return xpu::handle_pad(count_); + return handle_pad(count_); } [[nodiscard]] @@ -166,7 +166,7 @@ class soa { const T, T >; - return xpu::assume_aligned( + return assume_aligned( self.buffer_.data() + (arr_idx * self.stride()) ); } @@ -177,7 +177,7 @@ class soa { typename Self > [[nodiscard]] auto view(this Self&& self) noexcept { - static_assert(xpu::detail::checked_add(first_array, exposed_arrays) <= num_arrays); + static_assert(detail::checked_add(first_array, exposed_arrays) <= num_arrays); using element_t = std::conditional_t< std::is_const_v>, @@ -204,11 +204,11 @@ class soa_batch { : batches_{batches} , count_{count} , storage_{ - xpu::detail::checked_mul( + detail::checked_mul( batches, - xpu::detail::checked_mul( + detail::checked_mul( num_arrays, - xpu::handle_pad(count) + handle_pad(count) ) ) } @@ -231,7 +231,7 @@ class soa_batch { [[nodiscard]] CUDA_CALLABLE constexpr auto array_stride() const noexcept -> std::size_t { - return xpu::handle_pad(count_); + return handle_pad(count_); } [[nodiscard]] CUDA_CALLABLE @@ -251,7 +251,7 @@ class soa_batch { > [[nodiscard]] auto view(this Self&& self) noexcept { static_assert( - xpu::detail::checked_add(first_array, exposed_arrays) <= num_arrays, + detail::checked_add(first_array, exposed_arrays) <= num_arrays, "ERROR: number of viewed arrays is too large" ); @@ -260,7 +260,7 @@ class soa_batch { const T, T >; - return xpu::soa_batch_view{ + return soa_batch_view{ self.storage_.data() + (first_array * self.array_stride()), self.batches_, self.count_ }; @@ -273,7 +273,7 @@ class soa_batch { > [[nodiscard]] auto view(this Self&& self, std::size_t batch) noexcept { static_assert(exposed_arrays <= num_arrays, "ERROR: exposed arrays is greater than number of arrays"); - static_assert(xpu::detail::checked_add(first_array, exposed_arrays) <= num_arrays, + static_assert(detail::checked_add(first_array, exposed_arrays) <= num_arrays, "ERROR: number of viewed arrays is too large"); assert(batch < self.batches_); @@ -282,7 +282,7 @@ class soa_batch { const T, T >; - return xpu::soa_view{ + return soa_view{ self.storage_.data() + (batch * self.batch_stride()) + (first_array * self.array_stride()), self.element_count() };