#ifndef CK_BASE_HPP #define CK_BASE_HPP namespace ck { __device__ index_t get_thread_local_1d_id() { return threadIdx.x; } __device__ index_t get_block_1d_id() { return blockIdx.x; } template struct is_same { static constexpr bool value = false; }; template struct is_same { static constexpr bool value = true; }; template __host__ __device__ constexpr bool is_same_type(X, Y) { return is_same::value; } namespace math { // 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::value || is_same::value, "wrong type"); return (a + b - 1) / b; } }; template __host__ __device__ constexpr T integer_divide_ceil(T a, T b) { static_assert(is_same::value || is_same::value, "wrong type"); return (a + b - 1) / b; } template __host__ __device__ constexpr T max(T x, T y) { return x > y ? x : y; } 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::value, "not the same type"); return x > y ? x : y; } template __host__ __device__ constexpr T min(T x, T y) { return x < y ? x : y; } 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::value, "not the same type"); return x < y ? x : y; } // this is wrong // TODO: implement correct least common multiple, instead of calling max() template __host__ __device__ constexpr T lcm(T x, Ts... xs) { return max(x, xs...); } } // namespace math } // namspace ck #endif