【CUDA学习笔记】定位一个线程id

1、查看当前线程的序号

• 线程格里面总的线程个数N即可通过下面的公式算出：

N = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

• 当前线程位于线程格中的哪一个线程块blockId

blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

• 当前线程位于线程块中的哪一个线程threadId

 threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

• 计算一个线程块中一共有多少个线程M

M = blockDim.x*blockDim.y*blockDim.z

• 求得当前的线程序列号idx

idx = threadId + M*blockId;

2、例子： add2.cu,CUDA函数实现

先放代码，这里实现的功能是两个长度为的tensor相加，每个block有1024个线程，一共有个block。

__global__ void add2_kernel(float* c,
                            const float* a,
                            const float* b,
                            int n) {
    // 若n=20，而只有3个blcok块，每个blcok块5个线程，一共15个线程，则第0号线程需要处理数组中idx=0和idx=15中的数，因此i+=gridDim.x * blockDim.x -> i += 15，也就是这个线程需要执行的第二个数在数组中的idx
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
            i < n; i += gridDim.x * blockDim.x) {
        c[i] = a[i] + b[i];
    }
}

void launch_add2(float* c,
                 const float* a,
                 const float* b,
                 int n) {
    dim3 grid((n + 1023) / 1024); #+1023是为了向上取整
    dim3 block(1024);
    add2_kernel<<<grid, block>>>(c, a, b, n);
}