#pragma once #include "constant_integral.hip.hpp" template struct static_for_impl { template __host__ __device__ void operator()(F f) const { static_assert(Remaining % Increment == 0, "wrong! Remaining % Increment != 0"); static_assert(Increment <= Remaining, "will go out-of-range"); f(Number{}); static_for_impl{}(f); } }; template struct static_for_impl { template __host__ __device__ void operator()(F) const { // do nothing return; } }; template struct static_for { template __host__ __device__ void operator()(F f) const { static_assert(NBegin < NEnd, "Wrong! we should have NBegin < NEnd"); static_assert((NEnd - NBegin) % Increment == 0, "Wrong! should satisfy (NEnd - NBegin) % Increment == 0"); static_for_impl{}(f); } }; template struct static_const_reduce_n { template __host__ __device__ constexpr auto operator()(F f, Reduce r) const { static_assert(NLoop > 1, "out-of-range"); constexpr auto a = f(Number{}); auto b = static_const_reduce_n{}(f, r); // cannot use constexpr here, weird return r(a, b); } }; template <> struct static_const_reduce_n<1> { template __host__ __device__ constexpr auto operator()(F f, Reduce) const { return f(Number<0>{}); } }; #if 0 template __host__ __device__ constexpr auto unpacker(F f) { return [=](auto xs_array){ f(xs...); }; } #endif namespace mod_conv { template struct multiplies { __host__ __device__ constexpr T operator()(T a, T b) const { return a * b; } }; template struct plus { __host__ __device__ constexpr T operator()(T a, T b) const { return a + b; } }; } // namespace mod_conv