#ifndef CK_MATH_HPP #define CK_MATH_HPP #include "config.hpp" #include "integral_constant.hpp" namespace ck { 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 auto integer_divide_ceil(X x, Y y) { return (x + y - 1) / y; } template __host__ __device__ constexpr auto integer_least_multiple(X x, Y y) { return y * integer_divide_ceil(x, y); } 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