#pragma once template struct is_same { static const bool value = false; }; template struct is_same { static const bool value = true; }; __device__ unsigned get_thread_local_1d_id() { return threadIdx.x; } __device__ unsigned get_block_1d_id() { return blockIdx.x; } template struct Constant { static const T mValue = N; __host__ __device__ constexpr T Get() const { return mValue; } }; template using Number = Constant; template struct Sequence { static constexpr unsigned nDim = sizeof...(Is); const unsigned mData[nDim] = {Is...}; template __host__ __device__ constexpr unsigned Get(Number) const { return mData[I]; } template __host__ __device__ constexpr auto Reorder(Number, Number) const { constexpr unsigned IR0 = Get(Number{}); constexpr unsigned IR1 = Get(Number{}); return Sequence{}; } template __host__ __device__ constexpr auto Reorder(Number, Number, Number) const { constexpr unsigned IR0 = Get(Number{}); constexpr unsigned IR1 = Get(Number{}); constexpr unsigned IR2 = Get(Number{}); return Sequence{}; } template __host__ __device__ constexpr auto Reorder(Number, Number, Number, Number) const { constexpr unsigned IR0 = Get(Number{}); constexpr unsigned IR1 = Get(Number{}); constexpr unsigned IR2 = Get(Number{}); constexpr unsigned IR3 = Get(Number{}); return Sequence{}; } template __host__ __device__ constexpr auto Reorder(Sequence) const { constexpr unsigned IR0 = Get(Number{}); constexpr unsigned IR1 = Get(Number{}); constexpr unsigned IR2 = Get(Number{}); constexpr unsigned IR3 = Get(Number{}); return Sequence{}; } };