cuda 关于线程索引的理解

最近在学习cuda语言，然后发现线程的索引计算貌似有点不一样不太统一，所以写下自己的观点和看法。

只有线程块

kernel_function<<<num_blocks,num_threads>>>(){
    thread_idx=(blockIdx.x*blockDim.x)+threadIdx.x;
}
解释：线程的索引等于所处线程块的Id（blockIdx.x）乘上线程块大小（blockDim.x）再加上当前线程的偏移量。

线程网格

在线程网格下，
dim3 threads_square(16,8);
dim3 threads_square(2,2);
kernel_function<<<threads_square,threads_square>>>(){

    // idx 可以简单的理解为二维空间上x的绝对偏移量， 
    // threadIdx.x则是相对偏移,也就是在一个线程块中的x的偏移量
    const unsighned int idx=(blockIdx.x*blockDim.x)+threadIdx.x;

    // idy 可以简单的理解为二维空间上y的绝对偏移量， 
    // threadIdx.y则是相对偏移,也就是在一个线程块中的y的偏移量
    const unsighned int idy=(blockIdx.y*blockDim.y)+threadIdx.y;
    const unsighned int thread_idx=(gridDim.x*blockDim.x)*idy+idx;
}
解释：其实这个索引有点类似矩阵索引，先确定x方向上的绝对偏移量idx，为什么会有绝对偏移与相对偏移呢，因为整个大的线程网格是由很多网格组成的，同时每个网格里又是二维的，而threadIdx.x是在指该网格内的偏移，所以x的绝对偏移量是网格的x轴方向宽度blockDim.x乘于网格的偏移blockIdx.x再加上threadIdx.x。同理idy也是这么计算。
然后thread_idx也是相似的原理，y的绝对偏移量乘以x轴上线程数目(gridDim.x*blockDim.x)，再加上x轴本身自己的绝对偏移，就可以得到正确的索引了。

两者对比

其实两者是一个原理，是相同的。在没有网格只有线程块时，thread_idx=idx。
因为blockDim.y和gridDim.y 值都是1，所以idy的值只能取0，所以他们公式其实是一致统一的，有兴趣的可以稍微推算一下就会发现。