#ifndef CK_TUPLE_HPP #define CK_TUPLE_HPP #include "integral_constant.hpp" #include "sequence.hpp" #include "type.hpp" #include "enable_if.hpp" namespace ck { namespace detail { template struct TupleElementKey { __host__ __device__ constexpr TupleElementKey() = default; }; template struct TupleElement { __host__ __device__ constexpr TupleElement() = default; template >, TupleElement>::value, bool>::type = false> __host__ __device__ constexpr TupleElement(T&& v) : mData(std::forward(v)) { } Data mData; }; template __host__ __device__ constexpr const Data& get_tuple_element(const TupleElement& x) { return static_cast(x.mData); } template __host__ __device__ constexpr Data& get_tuple_element(TupleElement& x) { return x.mData; } // TODO: not sure the use of reference is correct template __host__ __device__ constexpr Data&& get_tuple_element(TupleElement&& x) { return static_cast(x.mData); } template struct TupleImpl; template struct TupleImpl, Xs...> : TupleElement, Xs>... { __host__ __device__ constexpr TupleImpl() = default; template >, TupleImpl>::value, bool>::type = false> __host__ __device__ constexpr TupleImpl(Y&& y) : TupleElement, Xs>(std::forward(y))... { } template = 2, bool>::type = false> __host__ __device__ constexpr TupleImpl(Ys&&... ys) : TupleElement, Xs>(std::forward(ys))... { static_assert(sizeof...(Is) == sizeof...(Xs) && sizeof...(Is) == sizeof...(Ys), "wrong! inconsistent size"); } __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); } template __host__ __device__ constexpr const auto& GetElementByKey(TupleElementKey) const { return get_tuple_element>(*this); } template __host__ __device__ constexpr auto& GetElementByKey(TupleElementKey) { return get_tuple_element>(*this); } }; } // namespace detail template struct Tuple : detail::TupleImpl::type, Xs...> { using base = detail::TupleImpl::type, Xs...>; __host__ __device__ constexpr Tuple() = default; template >, Tuple>::value, bool>::type = false> __host__ __device__ constexpr Tuple(Y&& y) : base(std::forward(y)) { } template = 2, bool>::type = false> __host__ __device__ constexpr Tuple(Ys&&... ys) : base(std::forward(ys)...) { } __host__ __device__ static constexpr index_t Size() { return sizeof...(Xs); } template __host__ __device__ constexpr const auto& At(Number) const { static_assert(I < base::Size(), "wrong! out of range"); return base::GetElementByKey(detail::TupleElementKey{}); } template __host__ __device__ constexpr auto& At(Number) { static_assert(I < base::Size(), "wrong! out of range"); return base::GetElementByKey(detail::TupleElementKey{}); } template __host__ __device__ constexpr const auto& operator[](Number i) const { return At(i); } template __host__ __device__ constexpr auto& operator()(Number i) { return At(i); } template __host__ __device__ constexpr auto operator=(const T& a) { static_assert(T::Size() == Size(), "wrong! size not the same"); static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; }); return *this; } __host__ __device__ static constexpr bool IsStaticBuffer() { return true; } }; template __host__ __device__ constexpr auto make_tuple(Xs&&... xs) { return Tuple...>(std::forward(xs)...); } } // namespace ck #endif