CUDA学习笔记
线程的索引和它的线程ID以一种简单的方式相互关联,threadIdx、blockIdx、blockDim是内置变量,用于索引线程位置。线程索引方式对于并行编程至关重要,它制约着内核函数中对循环的改写方式。
线程块需要独立执行:必须能够以任何顺序执行它们,并行执行或串行执行。这种独立性要求允许以任意顺序跨任意数量的内核调度线程块。块中的线程可以通过一些共享内存共享数据,并通过同步它们的执行来协调内存访问进行协作。更准确地说,可以通过调用内部函数来指定内核中的同步点;作为一个屏障,在这个屏障上,块中的所有线程都必须等待,然后才允许任何线程继续。
1.一维线程、一维线程块索引
threadIdx是一个三维矢量,包括threadIdx.x、threadIdx.y、threadIdx.z,对于函数主题设置为一维的线程的索引,仅用到threadIdx.x,下面展示一维线程索引。
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// 1个线程网格上的线程块个数,N:一个线程块上的线程个数
VecAdd<<<1, N>>>(A, B, C);
...
}
2.二维线程、二维线程块索引
线程和线程块均按二维排列的情况如下图,blockIdx为block索引的内置变量,是一个三维矢量;blockDim为一个block块包含的线程规模,也是一个三维矢量(例如下图中一个block中包含4列3行线程,因此blockDim.x大小为4,blockDim.y大小为3)。因此图中Thread(0,0)的线程索引为
i=blockIdx.x * blockDim.x+threadIdx.x;
j=blockIdx.y * blockDim.y + threadIdx.y;
下面是一个二维线程、二维线程块索引的示例。
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}
对于其它维度的线程块和线程索引,可以依照这样的思想类比得到。这里不再赘述,有疑问的可以在评论区交流。