GPU CUDA之——深入理解threadIdx

http://blog.csdn.net/canhui_wang/article/details/51730264

摘要

本文主要講述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中所有線程同時結束。