深入理解threadIdx

摘要

本文主要讲述CUDA的threadIdx。

1. Grid，Block和Thread三者的关系

其中，一个grid包含多个blocks，这些blocks的组织方式可以是一维，二维或者三维。任何一个block包含有多个Threads，这些Threads的组织方式也可以是一维，二维或者三维。举例来讲：比如上图中，任何一个block中有10个Thread，那么，Block(0,0)的第一个Thread的ThreadIdx是0，Block(1,0)的第一个Thread的ThreadIdx是11；Block(2,0)的第一个Thread的ThreadIdx是21，......，依此类推，不难整理出其中的映射公式（表达式已在代码中给出）。

2. GridID，BlockID，ThreadID三者的关系

ThreadID是线性增长的，其目的是用于在硬件和软件上唯一标识每一个线程。CUDA程序中任何一个时刻，每一个线程的ThreadIdx都是特定唯一标识的！grid，block的划分方式不同，比如一维划分，二维划分，或者三维划分。显然，Threads的唯一标识ThreadIdx的表达方式随着grid，block的划分方式（或者说是维度）而不同。下面通过程序给出ThreadIdx的完整的表达式。其中，由于使用的时候会考虑到GPU内存优化等原因，代码可能也会有所不同，但是threadId的计算的表达式是相对固定的。

/**************************************************************/
// ！！！！！！！！！！！！！！注意！！！！！！！！！！！！！！！！
/**************************************************************/
// grid划分成a维，block划分成b维，
// 等价于
// blocks是a维的，Threads是b维的。
// 这里，本人用的是第一中说法。
/**************************************************************/


// 情况1：grid划分成1维，block划分为1维。
__device__ int getGlobalIdx_1D_1D() {
    int threadId = blockIdx.x *blockDim.x + threadIdx.x;
    return threadId;
}

// 情况2：grid划分成1维，block划分为2维。
__device__ int getGlobalIdx_1D_2D() {
    int threadId = blockIdx.x * blockDim.x * blockDim.y
        + threadIdx.y * blockDim.x + threadIdx.x;
    return threadId;
}

// 情况3：grid划分成1维，block划分为3维。
__device__ int getGlobalIdx_1D_3D() {
    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z
        + threadIdx.z * blockDim.y * blockDim.x
        + threadIdx.y * blockDim.x + threadIdx.x;
    return threadId;
}

// 情况4：grid划分成2维，block划分为1维。
__device__ int getGlobalIdx_2D_1D() {
    int blockId = blockIdx.y * gridDim.x + blockIdx.x;
    int threadId = blockId * blockDim.x + threadIdx.x;
    return threadId;
}

// 情况5：grid划分成2维，block划分为2维。
__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;
}

// 情况6：grid划分成2维，block划分为3维。
__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;
}

// 情况7：grid划分成3维，block划分为1维。
__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;
}

// 情况8：grid划分成3维，block划分为2维。
__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;
}

// 情况9：grid划分成3维，block划分为3维。
__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;
}

3. GPU Threads与CPU Threads的比较

GPU Threads的生成代价小，是轻量级的线程；CPU Threads的生成代价大，是重量级的线程。CPU Threads虽然生成的代价高于GPU Threads，但其执行效率高于GPU Threads，所以GPU Threads无法在个体的比较上取胜，只有在数量上取胜。在这个意义上来讲，CPU Threads好比是一头强壮的公牛在耕地，GPU Threads好比是1000头弱小的小牛在耕地。因此，为了保证体现GPU并行计算的优点，线程的数目必须足够多，通常至少得用上1000个GPU线程或者更多才够本，才能很好地体现GPU并行计算的优点！

4. GPU Threads的线程同步

线程同步是针对同一个block中的所有线程而言的，因为只有同一个block中的线程才能在有效的机制中共同访问shared memory。要知道，由于每一个Thread的生命周期长度是不相同的，Thread对Shared Memory的操作可能会导致读写的不一致，因此需要线程的同步，从而保证该block中所有线程同时结束。