由於專案需要用到GPU，所以最近開始學習CUDA程式設計模型，剛開始接觸，先搞清楚執行緒關係和記憶體模型是非常重要的，但是發現書上和許多部落格關於執行緒這些關係沒講明白，所以就著自己的理解，做點筆記，歡迎討論。

        這篇文章針對於已經瞭解過了CUDA執行緒的相關知識，最好已經動手寫過CUDA C的程式碼，而對並行執行緒感到迷惑，不知道怎麼計算執行緒索引的讀者，如果沒接觸過，那麼先看看書，敲兩段程式碼跑跑，如果你理解了那麼恭喜你，如果還有疑惑，那麼再來看看這篇文章，或許有幫助。

        首先，認識一下執行緒。(⊙o⊙)…雖然畫成這樣總是感覺有點怪怪的，但是你應該已經見怪不怪了吧~

下面我們來看一段程式碼，其功能是對兩個陣列求和，並儲存到另一個數組，很簡單吧~

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
 
using namespace std;
 
// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}
 
int main() {
    const int length = 16;                                      // 陣列長度為16
    float a[length], b[length], c[length];                      // host中的陣列
    for (int i = 0; i < length; i++) {                          // 初始賦值
        a[i] = b[i] = i;
    }
    float* a_device, *b_device, *c_device;                      // device中的陣列
 
    cudaMalloc((void**)&a_device, length * sizeof(float));      // 分配記憶體
    cudaMalloc((void**)&b_device, length * sizeof(float));
    cudaMalloc((void**)&c_device, length * sizeof(float));
 
    cudaMemcpy(a_device, a, length * sizeof(float), cudaMemcpyHostToDevice);    // 將host陣列的值拷貝給device陣列
    cudaMemcpy(b_device, b, length * sizeof(float), cudaMemcpyHostToDevice);
 
    // 一：引數配置
    dim3 grid(1, 1, 1), block(length, 1, 1);                    // 設定引數
    vector_add<<<grid,block>>>(a_device, b_device, c_device, length);           // 啟動kernel
 
    cudaMemcpy(c, c_device, length * sizeof(float), cudaMemcpyDeviceToHost);    // 將結果拷貝到host
 
    for (int i = 0; i < length; i++) {                          // 打印出來方便觀察
        cout << c[i] << " ";
    }
    cout << endl;
 
    system("pause");
    return 0;
}
 

執行結果：

結果是對的，也是我們所能預料到的。那麼現在我們來分析程式碼中註釋的一、二處究竟該怎麼來寫。
        首先，我們要明白，上面的程式碼計算的是兩個一維向量的和。由於陣列大小是16，所以我們使用了16個執行緒來計算。

dim3 grid(1, 1, 1), block(length, 1, 1);                    // 設定引數

        先說grid，在這段程式碼中，我們設定引數為執行緒格（grid）中只有一個一維的block，該block的x維度上有16個，這個應該一下就看出來啦。因為grid(x,y,z)中的x=1,y=1,z=1，即各個維度均為1，所以是一維的，數量為x*y*z=1*1*1=1。如果沒明白，再看兩個例子：

dim3 grid1(2, 1, 1); // x=2, y=1, z=1
dim3 grid2(4, 2, 1); // x=4, y=2, z=1
dim3 grid3(2, 3, 4); // x=2, y=3, z=4

        可以知道，grid1是一維的（因為y,z維度是1），grid2是二維的（因為z維度是1），grid3是三維的，且grid1,grid2,grid3中分別有2、8、24個block。
        同理，對於執行緒塊（block），我們知道之前的程式碼中，block中存在16個執行緒，且該執行緒塊維度是一維的，因為block(x,y,z)中x=length=16,y=1,z=1。

我畫個圖來幫助理解，大概就是這樣子的：

dim3 grid(1, 1, 1), block(length, 1, 1);                    // 設定引數

        OK，我想這下應該就清楚了，就是一個一維的block（此處只有x維度上存在16個執行緒）。所以，內建變數只有一個在起作用，就是threadIdx.x，它的範圍是[0,15]。因此，我們在計算執行緒索引是，只用這個內建變數就行了（其他的為0，寫了也不起作用）：

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.x;              // 只使用了threadIdx.x
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

OK，看到這裡，你可能還是不大明白什麼一維二維的，我們再來看一個：

dim3 grid(1, 1, 1), block(8, 2, 1);                    // 設定引數

        根據上面的介紹，我們知道這個執行緒格只有一個一維的執行緒塊，該執行緒塊內的執行緒是二維的，x的維度為8，y的維度為2，共有8*2=16個執行緒，如果要用這16個執行緒來計算陣列的累加，當然是可以的，但是我們這裡需要改動一下執行緒執行程式碼中的索引計算方式了。

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.y * blockDim.x +  threadIdx.x;  // 使用了threadIdx.x, threadIdx.x, blockDim.x
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

        我們一定要有並行思想，這裡有16個執行緒，kernel啟動後，每個執行緒都有自己的索引號，比如某個執行緒位於grid中哪個維度的block（即blockIdx.x,block.y,block.z），又位於該block的哪個維度的執行緒（即threadIdx.x,threadIdx.y,threadIdx.z），利用這些執行緒索引號對映到對應的陣列下標，我們要做的工作就是將保證這些下標不重複（如果重複的話，那就慘了），最初那種一維的計算方式就不行了。因此，通過使用threadIdx，blockDim來進行對映（偏移）。blockDim.x=8,blockDim.y=2，如上程式碼。

        其實，我感覺有些我不能用文字準確、清晰的描述出來，所以咯，我們再來一個例子吧，我相信，多看一看，多想一想就明白了。

dim3 grid(1, 1, 1), block(4, 4, 1);                    // 設定引數

        我們將block改成上面的這樣，其執行緒模型為下圖：

當然，kernel函式的程式碼依然可以不用變動，這個應該想得清楚，還是再寫一下吧。

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.y * blockDim.x +  threadIdx.x;  // 使用了threadIdx.x, threadIdx.x, blockDim.x
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

—————————————————————————————————————————————————————————————————————————————
        以上內容我們分別介紹了用一維和二維執行緒來計算一維陣列的求和，實際上陣列的維度與執行緒格、執行緒塊和執行緒的維度並不是那麼密不可分的，都可以組合實現，只不過在實現時，良好的引數配置對索引的計算很方便，而且由於grid、block、thread維度的限制，還有warpSize的限制，所以對於較大的資料量來說，我們應該做到心中有數，進行有效的塊分解。

        現在來看看二維的block，在整個文章中，我只講解一維、二維的，因為三維的我不知道怎麼畫圖啦，而且不好描述，免得誤導大家。
        還是上面的一維陣列，長度為16。

dim3 grid(16, 1, 1), block(1, 1, 1);                    // 設定引數

        先來個執行緒模型圖，我想大家並不會感到驚訝，綠色的區域表示grid，藍色的區域表示block，圖中有一個grid和16個block，每個block都是一維，而且x維度上只有一個執行緒的：

顯然，我們的執行緒索引程式碼應該為如下：

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = blockIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

或許你會有疑惑，那麼我們再來看一個：
 

dim3 grid(4, 1, 1), block(4, 1, 1);

執行緒索引程式碼應該為如下：

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

        到現在為止，我覺得你應該有所領悟。如果還是不曉得的話，我想你應該認認真真的看圖並動手分析了，圖中的每一個塊，每一個字都是有它的作用的，你不應該就此放棄。
        我依然相信，能用圖解決，就不嗶嗶。就好像你給一個人描述一座宮殿是多麼多麼的巨集偉，富麗堂皇，他並不不會感冒。你就說，嘿大傻，給你瞧瞧我去歐洲玩的教堂，這是照片，不用多說，大傻自己就知道了。
比如，我描述說：

dim3 grid(2, 2, 1), block(2, 2, 1);

這樣肯定不直觀，那我再給你一幅示意圖：

那麼執行程式碼及索引計算如下：

// 二：執行緒執行程式碼
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    // 在第幾個塊中 * 塊的大小 + 塊中的x, y維度（幾行幾列）
    int tid = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.x + threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

        上面的程式碼可能要複雜一點，但是你慢慢的會發現這很有趣。
        到此，我想講的就完了。當然對於二維的陣列或是三維的陣列，我想多看幾個例子也就會有體會了。
        這裡還是忍不住要吐槽一下內建變數threadIdx和blockIdx的命名了，每次看到這些內建變數其最後一個字母是x，就會給我一種誤會是x維度上的，我覺得使用threadId和blockId是多麼的良好。當然，勝利的總是API一方，我也只能吐吐槽。

—————————————————————————————————————————————————————————————————————————————

最後再來一發，我給個圖，我們來倒推其引數及相關執行程式碼，如下：

由於上傳圖片大小限制，由BMP轉成JPG格式的了，有點不清晰，但足夠看了。

顯然引數為：

dim3 grid(8, 4, 1), block(8, 2, 1);

共有8*4*8*2=512個執行緒，當然在CUDA程式設計中，這算很少的了。如果是一幅512x512大小的影象做加或點乘之類的運算，隨隨便便就是幾十萬的執行緒數了。
萬變不離其宗，其一維的計算方式如下：

__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    // 在第幾個塊中 * 塊的大小 + 塊中的x, y維度（幾行幾列）
    int tid = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.x + threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

再給出二維的：

__global__ void vector_add(float** mat1, float** mat2, float** matres, int width) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < width) {
        matres[x][y] = mat1[x][y] + mat2[x][y];
    }
}

--------------------- 
作者：hujingshuang 
來源：CSDN 
原文：https://blog.csdn.net/hujingshuang/article/details/53097222 
版權宣告：本文為博主原創文章，轉載請附上博文連結！