#ifndef CK_FUNCTIONAL3_HPP #define CK_FUNCTIONAL3_HPP #include "functional.hpp" #include "functional2.hpp" #include "Sequence.hpp" #include "Array.hpp" namespace ck { namespace detail { // RemainLengths: Sequence<...> // Orders: Sequence<...> template struct static_ford_impl { __host__ __device__ constexpr static_ford_impl() { static_assert(RemainLengths::GetSize() > 0, "wrong! should not get here"); } // F signature: F(Sequence<...>) // CurrentOrderedId: Sequence<...> template __host__ __device__ constexpr void operator()(F f, CurrentOrderedId) const { static_for<0, RemainLengths::Front(), 1>{}([=](auto I) { static_ford_impl{}( f, CurrentOrderedId::PushBack(I)); }); } }; template struct static_ford_impl, Orders> { // F signature: F(Sequence<...>) // OrderedId: Sequence<...> template __host__ __device__ constexpr void operator()(F f, OrderedId) const { // retrive unordered Id f(OrderedId::ReorderGivenOld2New(Orders{})); } }; // RemainLengths: Sequence<...> // Orders: Sequence<...> template struct ford_impl { __host__ __device__ constexpr ford_impl() { static_assert(RemainLengths::GetSize() > 0, "wrong! should not get here"); } // F signature: F(Array<...> multi_id) // CurrentOrderdId: Array<...> template __host__ __device__ constexpr void operator()(F f, CurrentOrderedId current_ordered_id) const { for(index_t i = 0; i < RemainLengths::Front(); ++i) { ford_impl{}( f, current_ordered_id.PushBack(i)); } } }; template struct ford_impl, Orders> { // F signature: F(Array<...> multi_id) // CurrentOrderdId: Array<...> template __host__ __device__ constexpr void operator()(F f, CurrentOrderedId current_ordered_id) const { // retrive unordered Id f(reorder_array_given_old2new(current_ordered_id, Orders{})); } }; } // namespace detail // Lengths is Sequence<...>, it is the length of each dimension for N-dimensional loop // Orders is Sequence<...>, it is the order of dimension in which static_ford will loop over each // dimension template ::type> struct static_ford { __host__ __device__ constexpr static_ford() { static_assert(Lengths::GetSize() > 0, "wrong! Lengths is empty"); static_assert(Lengths::GetSize() == Orders::GetSize(), "wrong! inconsistent size"); } // F signature: F(Sequence<...> multi_id) // multi_id is the unordered multi-index template __host__ __device__ constexpr void operator()(F f) const { constexpr auto ordered_lengths = Lengths::ReorderGivenNew2Old(Orders{}); detail::static_ford_impl{}(f, Sequence<>{}); } }; // Lengths is Sequence<...>, it is the length of each dimension for N-dimensional loop // Orders is Sequence<...>, it is the order of dimension in which ford will loop over each // dimension template ::type> struct ford { __host__ __device__ constexpr ford() { static_assert(Lengths::GetSize() > 0, "wrong! Lengths is empty"); static_assert(Lengths::GetSize() == Orders::GetSize(), "wrong! inconsistent size"); } // F signature: F(Array<...> multi_id) // multi_id is the unordered multi-index template __host__ __device__ constexpr void operator()(F f) const { constexpr auto ordered_lengths = Lengths::ReorderGivenNew2Old(Orders{}); for(index_t i = 0; i < ordered_lengths.Front(); ++i) { detail::ford_impl{}(f, Array{i}); } } }; } // namespace ck #endif