#pragma once #define WARPSIZE 32; 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 { using Type = 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 ReorderByGetNewFromOld(Sequence) const { constexpr auto old_sequence = Type{}; constexpr unsigned NR0 = old_sequence.mData[I0]; constexpr unsigned NR1 = old_sequence.mData[I1]; constexpr unsigned NR2 = old_sequence.mData[I2]; constexpr unsigned NR3 = old_sequence.mData[I3]; return Sequence{}; } template __host__ __device__ constexpr auto ReorderByPutOldToNew(Sequence) const { // don't know how to implement this printf("Sequence::ReorderByPutOldToNew not implemented"); assert(false); } };