CUDA:Indexing
CUDA Thread Indexing Cheatsheet.
1D grid of 1D blocks
__device__ int getGlobalIdx_1D_1D()
{
return (blockIdx.x * blockDim.x) /*GLOBAL_BLOCK_OFFSET*/
+ threadIdx.x /*LOCAL_THREAD_OFFSET*/;
}
1D grid of 2D blocks
__device__ int getGlobalIdx_1D_2D()
{
return (blockIdx.x * (blockDim.x * blockDim.y)) /*GLOBAL_BLOCK_OFFSET*/
+ (threadIdx.y * blockDim.x) /*LOCAL_THREAD_Y_OFFSET*/
+ threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
}
1D grid of 3D blocks
__device__ int getGlobalIdx_1D_3D()
{
return (blockIdx.x * (blockDim.x * blockDim.y * blockDim.z)) /*GLOBAL_BLOCK_OFFSET*/
+ (threadIdx.z * (blockDim.y * blockDim.x)) /*LOCAL_THREAD_X_OFFSET*/
+ (threadIdx.y * blockDim.x) /*LOCAL_THREAD_Y_OFFSET*/
+ threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
}
2D grid of 1D blocks
__device__ int getGlobalIdx_2D_1D()
{
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
2D grid of 2D blocks
__device__ int getGlobalIdx_2D_2D()
{
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
return threadId;
}
2D grid of 3D blocks
__device__ int getGlobalIdx_2D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
3D grid of 1D blocks
__device__ int getGlobalIdx_3D_1D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
3D grid of 2D blocks
__device__ int getGlobalIdx_3D_2D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y)
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
3D grid of 3D blocks
__device__ int getGlobalIdx_3D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
Warp ID & Lane ID
__forceinline__ __device__ unsigned lane_id()
{
unsigned ret;
asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret));
return ret;
}
__forceinline__ __device__ unsigned warp_id()
{
unsigned ret;
asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
return ret;
}
See also
Favorite site
References
-
Programming_NVIDIA_CUDA_-_1D_array_index.pdf ↩