#pragma once #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 TupleElementKeyData { #if 0 // workaround compiler complaint about implicitly-deleted default constructor __host__ __device__ constexpr TupleElementKeyData() = default; #else __host__ __device__ constexpr TupleElementKeyData() : mData{} {} #endif template , TupleElementKeyData>::value, bool>::type = false> __host__ __device__ constexpr TupleElementKeyData(T&& v) : mData(std::forward(v)) { } Data mData; }; template __host__ __device__ constexpr const Data& get_tuple_element_data(const TupleElementKeyData& x) { return static_cast(x.mData); } template __host__ __device__ constexpr Data& get_tuple_element_data(TupleElementKeyData& x) { return x.mData; } // TODO: not sure the use of reference is correct template __host__ __device__ constexpr Data&& get_tuple_element_data(TupleElementKeyData&& x) { return static_cast(x.mData); } template struct TupleImpl; template struct TupleImpl, Xs...> : TupleElementKeyData, Xs>... { __host__ __device__ constexpr TupleImpl() = default; template , TupleImpl>::value, bool>::type = false> __host__ __device__ constexpr TupleImpl(Y&& y) : TupleElementKeyData, Xs>(std::forward(y))... { } template = 2, bool>::type = false> __host__ __device__ constexpr TupleImpl(Ys&&... ys) : TupleElementKeyData, 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& GetElementDataByKey(TupleElementKey) const { return get_tuple_element_data>(*this); } template __host__ __device__ constexpr auto& GetElementDataByKey(TupleElementKey) { return get_tuple_element_data>(*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); } // read access template __host__ __device__ constexpr const auto& At(Number) const { static_assert(I < base::Size(), "wrong! out of range"); return base::GetElementDataByKey(detail::TupleElementKey{}); } // write access template __host__ __device__ constexpr auto& At(Number) { static_assert(I < base::Size(), "wrong! out of range"); return base::GetElementDataByKey(detail::TupleElementKey{}); } // read access template __host__ __device__ constexpr const auto& operator[](Number i) const { return At(i); } // write access 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 <> struct Tuple<> { __host__ __device__ constexpr Tuple() = default; __host__ __device__ static constexpr index_t Size() { return 0; } template __host__ __device__ constexpr auto operator=(const T&) { return *this; } __host__ __device__ static constexpr bool IsStaticBuffer() { return true; } }; template struct tuple_element { using type = decltype(TTuple{}.At(Number{})); }; template using tuple_element_t = typename tuple_element::type; template __host__ __device__ constexpr auto make_tuple(Xs&&... xs) { return Tuple...>(std::forward(xs)...); } // https://en.cppreference.com/w/cpp/utility/tuple/tie template constexpr Tuple tie(Args&... args) noexcept { return {args...}; } } // namespace ck