#ifndef CK_UTILITY_HPP #define CK_UTILITY_HPP #include #include "config.hpp" namespace ck { template using is_same = std::is_same; namespace math { template struct scales { __host__ __device__ constexpr T operator()(T a) const { return s * a; } }; template struct plus { __host__ __device__ constexpr T operator()(T a, T b) const { return a + b; } }; template struct minus { __host__ __device__ constexpr T operator()(T a, T b) const { return a - b; } }; template struct multiplies { __host__ __device__ constexpr T operator()(T a, T b) const { return a * b; } }; template struct integer_divide_ceiler { __host__ __device__ constexpr T operator()(T a, T b) const { static_assert(is_same{} || is_same{}, "wrong type"); return (a + b - 1) / b; } }; template __host__ __device__ constexpr T integer_divide_ceil(T a, T b) { static_assert(is_same{} || is_same{}, "wrong type"); return (a + b - 1) / b; } template __host__ __device__ constexpr T integer_least_multiple(T a, T b) { static_assert(is_same{} || is_same{}, "wrong type"); return b * integer_divide_ceil(a, b); } template __host__ __device__ constexpr T max(T x) { return x; } template __host__ __device__ constexpr T max(T x, Ts... xs) { static_assert(sizeof...(xs) > 0, "not enough argument"); auto y = max(xs...); static_assert(is_same{}, "not the same type"); return x > y ? x : y; } template __host__ __device__ constexpr T min(T x) { return x; } template __host__ __device__ constexpr T min(T x, Ts... xs) { static_assert(sizeof...(xs) > 0, "not enough argument"); auto y = min(xs...); static_assert(is_same{}, "not the same type"); return x < y ? x : y; } // this is WRONG // TODO: implement least common multiple properly, instead of calling max() template __host__ __device__ constexpr T lcm(T x, Ts... xs) { return max(x, xs...); } } // namespace math } // namspace ck #endif