#ifndef CK_STATIC_BUFFER_HPP #define CK_STATIC_BUFFER_HPP #include "statically_indexed_array.hpp" namespace ck { template struct StaticBuffer : public StaticallyIndexedArray { using type = T; using base = StaticallyIndexedArray; T invalid_element_value_ = T{0}; __host__ __device__ constexpr StaticBuffer() : base{} {} __host__ __device__ constexpr StaticBuffer(T invalid_element_value) : base{}, invalid_element_value_{invalid_element_value} { } __host__ __device__ static constexpr AddressSpaceEnum_t GetAddressSpace() { return BufferAddressSpace; } template __host__ __device__ constexpr auto Get(Number i, bool is_valid_element) const { if constexpr(InvalidElementUseNumericalZeroValue) { return is_valid_element ? At(i) : T{0}; } else { return is_valid_element ? At(i) : invalid_element_value_; } } template __host__ __device__ void Set(Number i, bool is_valid_element, const T& x) { if(is_valid_element) { At(i) = x; } } __host__ __device__ static constexpr bool IsStaticBuffer() { return true; } __host__ __device__ static constexpr bool IsDynamicBuffer() { return false; } }; template struct StaticBufferV2 : public StaticallyIndexedArray { using type = T; using base = StaticallyIndexedArray; using VecBaseType = typename T::d1_t; __host__ __device__ static constexpr index_t GetVectorSize() { return sizeof(typename T::type) / sizeof(VecBaseType); } static constexpr index_t vector_size = GetVectorSize(); VecBaseType invalid_element_value_ = VecBaseType{0}; T invalid_vec_value_ = T{0}; __host__ __device__ constexpr StaticBufferV2() : base{} {} __host__ __device__ constexpr StaticBufferV2(VecBaseType invalid_element_value) : base{}, invalid_vec_value_{invalid_element_value}, invalid_element_value_{invalid_element_value} { } __host__ __device__ static constexpr AddressSpaceEnum_t GetAddressSpace() { return BufferAddressSpace; } template __host__ __device__ constexpr auto& GetVector(Number vec_id) { return this->At(vec_id); } template __host__ __device__ constexpr const auto& GetVector(Number vec_id) const { return this->At(vec_id); } template __host__ __device__ constexpr auto& GetElement(Number i, bool) { constexpr auto vec_id = Number{}; constexpr auto vec_off = Number{}; return this->At(vec_id).template AsType()(vec_off); } template __host__ __device__ constexpr auto GetElement(Number i, bool is_valid_element) const { constexpr auto vec_id = Number{}; constexpr auto vec_off = Number{}; if constexpr(InvalidElementUseNumericalZeroValue) { return is_valid_element ? this->At(vec_id).template AsType()[vec_off] : VecBaseType{0}; } else { return is_valid_element ? this->At(vec_id).template AsType()[vec_off] : invalid_element_value_; } } template __host__ __device__ constexpr auto operator[](Number i) const { return GetElement(i, true); } template __host__ __device__ constexpr auto& operator()(Number i) { return GetElement(i, true); } __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 __host__ __device__ constexpr auto make_static_buffer(Number, T invalid_element_value) { return StaticBuffer{invalid_element_value}; } } // namespace ck #endif