// SPDX-License-Identifier: MIT // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once namespace ck { template struct MetaDataBuffer { __host__ __device__ constexpr MetaDataBuffer() : buffer_{}, size_{0} {} template __host__ __device__ constexpr MetaDataBuffer(const X& x, const Xs&... xs) : buffer_{}, size_{0} { Push(x, xs...); } template __host__ __device__ constexpr void Push(const T& data) { if constexpr(!is_empty_v) { constexpr index_t size = sizeof(T); auto tmp = bit_cast>(data); for(int i = 0; i < size; i++) { buffer_(size_) = tmp[i]; size_++; } } } template __host__ __device__ constexpr void Push(const X& x, const Xs&... xs) { Push(x); Push(xs...); } template __host__ __device__ constexpr T Pop(index_t& pos) const { T data; if constexpr(!is_empty_v) { constexpr index_t size = sizeof(T); Array tmp; for(int i = 0; i < size; i++) { tmp(i) = buffer_[pos]; pos++; } data = bit_cast(tmp); } return data; } template __host__ __device__ constexpr T Get(index_t pos) const { constexpr index_t size = sizeof(T); Array tmp; for(int i = 0; i < size; i++) { tmp(i) = buffer_[pos]; pos++; } auto data = bit_cast(tmp); return data; } // Array buffer_; index_t size_ = 0; }; } // namespace ck