#ifndef CK_BUFFER_HPP #define CK_BUFFER_HPP #include "statically_indexed_array.hpp" namespace ck { template struct StaticBuffer : public StaticallyIndexedArray { using type = T; using base = StaticallyIndexedArray; __host__ __device__ constexpr StaticBuffer() : base{} {} __host__ __device__ static constexpr bool IsStaticBuffer() { return true; } __host__ __device__ static constexpr bool IsDynamicBuffer() { return false; } }; template __host__ __device__ constexpr auto make_static_buffer(Number) { return StaticBuffer{}; } template struct DynamicBuffer { using type = T; T* p_data_; __host__ __device__ constexpr DynamicBuffer(T* p_data) : p_data_{p_data} {} __host__ __device__ constexpr const T& operator[](index_t i) const { return p_data_[i]; } __host__ __device__ constexpr T& operator()(index_t i) { return p_data_[i]; } template >>::type, typename scalar_type>>::type>::value, bool>::type = false> __host__ __device__ constexpr const auto Get(index_t i) const { return *reinterpret_cast(&p_data_[i]); } template >>::type, typename scalar_type>>::type>::value, bool>::type = false> __host__ __device__ void Set(index_t i, const X& x) { *reinterpret_cast(&p_data_[i]) = x; } __host__ __device__ static constexpr bool IsStaticBuffer() { return false; } __host__ __device__ static constexpr bool IsDynamicBuffer() { return true; } }; template __host__ __device__ constexpr auto make_dynamic_buffer(T* p) { return DynamicBuffer{p}; } } // namespace ck #endif