// SPDX-License-Identifier: MIT // Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once #include "layout.hpp" #include "utils/tensor_utils.hpp" namespace ck { namespace wrapper { template struct Tensor { using ElementSpaceSize = decltype(Layout{ Shape{}, Strides{}}.GetElementSpaceSize()); __host__ __device__ Tensor() = delete; __host__ __device__ Tensor(ElementType* pointer, const Layout& layout) : layout_(layout), dynamic_buffer_( make_dynamic_buffer(pointer, layout.GetElementSpaceSize())) { } template __host__ __device__ index_t operator[](const Tuple& Idx) const { // Padding is not supported, so we can assume that read should be valid. return dynamic_buffer_.template Get(layout_(Idx), true /*is_valid*/); } template __host__ __device__ index_t operator()(const Tuple& Idx) const { return dynamic_buffer_.template Get(layout_(Idx), true /*is_valid*/); } template __host__ __device__ index_t operator()(Idxs... idxs) const { const auto idxs_tuple = make_tuple(idxs...); return dynamic_buffer_.template Get(layout_(idxs_tuple), true /*is_valid*/); } private: const Layout& layout_; DynamicBuffer dynamic_buffer_; }; } // namespace wrapper } // namespace ck