CUDA學習筆記-1: CUDA程式設計概覽

火腿燒豆腐發表於2021-08-04
筆記程式設計

1.GPU程式設計模型及基本步驟

cuda程式的基本步驟如下：

image-20210804164019922

  • 在cpu中初始化資料
  • 將輸入transfer到GPU中
  • 利用分配好的grid和block啟動kernel函式
  • 將計算結果transfer到CPU中
  • 釋放申請的記憶體空間

從上面的步驟可以看出，一個CUDA程式主要包含兩部分，第一部分執行在CPU上，稱作Host code，主要負責完成複雜的指令；第二部分執行在GPU上，稱作Device code，主要負責並行地完成大量的簡單指令（如數值計算）；

2.基本設施

執行在GPU中地函式稱作kernel，該函式有這麼幾個要求：

  • 宣告時在返回型別前需要新增"__globol__"的標識
  • 返回值只能是void
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

這就是一個合規的核函式。

除了宣告時的不同，和函式的呼叫也是不一樣的，需要以 “kernel_name <<< >>>();”的形式呼叫。而在尖括號中間，則是定義了啟用了多少個GPU核，學習這一引數的使用，我們還需要知道下面幾個概念：

  • dim3：一種資料型別，包含x,y,z三個int 型別的成員，在初始化時一個dim3型別的變數時，成員值預設為1
  • grid : 一個grid中包含多個block
  • block: 一個block包含多個thread

我們以一種更抽象的方式來理解GPU中程式的執行方式的話，可以這麼看：

GPU中的每個核可以獨立的執行一個執行緒，那我們就使用thread來代表GPU中的核，但一個GPU中的核數量很多，就需要有更高階的結構對全部用到的核進行約束、管理，這就是block(塊)，一個塊中可以包含多個核，並且這些核在邏輯上的排布可以是三維的，在一個塊中我們可以使用一個dim3型別的量threadIdx來表示每個核所處的位置,threadIdx.x、threadIdx.y、threadIdx.z分別表示在三個維度上的座標；此外，每個塊還帶有一個dim3型別的屬性blockDim，blockDim.x、blockDim.y、blockDim.z分別表示該block三個維度上各有多少個核，這個block中的總核數為blockDim.x * blockDim.y * blockDim.z；

我們一次使用的多個block，最好能使用一個容器把他們都包起來，這就是grid，類比於上文中thread和block的關係，block和grid也有相似的關係。我們使用blockIdx.x、blockIdx.y、blockIdx.z表示每個block在grid中的位置；同樣，grid也具有gridDim.x、gridDim.y和gridDim.z三個屬性以及三者相乘的總block數。

知道了上面這些知識後，我們可以對“kernel_name <<< >>>();”中尖括號中的引數做一個更具體的解釋，它應該被定義為在GPU中執行這一核函式的所有核的組織形式，以"kernel_name <<< number_of_blocks, thread_per_block>>> (arguments)"的形式使用，一個典型的示例如下：

int nx = 16;
int ny = 4;
dim3 block(8, 2); // z預設為1
dim3 grid(nx/8, ny/2);
addKernel << <grid, block >> >(c, a, b);

這一示例中建立了一個有(2*2)個block的grid，每個block中有(8*2)個thread，下圖給出了更直觀的表述：

image-20210804121951952

需要注意的是，對block、grid的尺寸定義並不是沒有限制的，一個GPU中的核的數量同樣是有限制的。對於一個block來說，總的核數不得超過1024，x、y維度都不得超過1024，z維度不得超過64，如下圖

image-20210804121903028

對於整個grid而言，x維度上不得有超過\(2^{32}-1\)個thread，注意這裡是thread而不是block，在其y維度和z維度上thread數量不得超過65536.

image-20210804122008408

在cuda程式設計中我們經常會把陣列的每一個元素分別放到單獨的一個核中處理，我們可以利用核的索引讀取陣列中的資料進行操作，但由於block、grid的存在，索引的獲取需要一定的計算，在exercise2中給出了一個3D模型中取值的訓練，實現如下

__global__ void print_array(int *input)
{
    int tid = (blockDim.x*blockDim.y)*threadIdx.z + blockDim.x*threadIdx.y + threadIdx.x;
    int xoffset = blockDim.x * blockDim.y * blockDim.z;
    int yoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x;
    int zoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x * gridDim.y;
    int gid = zoffset * blockIdx.z + yoffset * blockIdx.y + xoffset * blockIdx.x + tid;
    printf("blockIdx.x : %d, blockIdx.y : %d, blockIdx.z : %d,gid : %d, value: %d\n", blockIdx.x, blockIdx.y, blockIdx.z, gid, input[gid]);
}

3.資料在host和device之間的遷移

我們前邊提到，cuda的程式設計步驟是將資料移入GPU，待計算完成後將其取出，官方對可能涉及到的記憶體操作類的操作都給出了介面。

首先是cudaMemCpy函式，其定義為

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

該函式是將資料從CPU移入到GPU或者從GPU移出到CPU中，引數0指向目標區域的地址，引數1指向資料的源地址，引數2表示要移動的資料的位元組數，最後一個參數列示資料的移動方向（cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost或cudaMemcpyDeviceToDevice）

此外，對應C語言的記憶體空間操作，cuda也推出了CudaMalloc, CudaMemset, CudaFree三個介面

cudaError_t  cudaMalloc ( void** devPtr, size_t size );
cudaError_t  cudaMemset ( void* devPtr, int  value, size_t count );
cudaError_t  cudaFree ( void* devPtr );

這裡需要注意的一個點是cudaMalloc的第一引數的資料型別為void**，這一點怎麼理解呢？

這裡我們結合一個示例進行解釋:

int *d_input;
cudaMalloc((void **) &d_input, bytesize);

之所以使用void，是因為這一步只管分配記憶體，不考慮如何解釋指標，所以只需要傳入待分配記憶體的地址，不需要傳入具體的型別，其他API中的 void* 也是同理。為什麼是兩個*呢，這是因為我們在定義d_input時是定義了主存中的一個指標，它指向主存中的一個地址；而&d_input則是取得了儲存該指標值的地址，cudaMalloc利用這一地址將在GPU中分配給該緩衝區的首地址賦值給d_input。

利用上述的幾個介面函式，我們就可以實現一個基本的cuda程式的主函式：

int main()
{
    const int arraySize = 64;
    const int byteSize = arraySize * sizeof(int);

    int *h_input,*d_input;
    h_input = (int*)malloc(byteSize);
    cudaMalloc((void **)&d_input,byteSize);

    srand((unsigned)time(NULL));
    for (int i = 0; i < 64; ++i)
    {
        if(h_input[i] != NULL)h_input[i] = (int)rand()& 0xff;
    }

    cudaMemcpy(d_input, h_input, byteSize, cudaMemcpyHostToDevice);

    int nx = 4, ny = 4, nz = 4;
    dim3 block(2, 2, 2);
    dim3 grid(nx/2, ny/2, nz/2);
    print_array << < grid, block >> > (d_input);
    cudaDeviceSynchronize();

    cudaFree(d_input);
    free(h_input);

    return 0;
}

其中 cudaDeviceSynchronize(); 的作用是在此處等待GPU中計算完成後再繼續執行後續的程式碼。

4 錯誤處理

在C++中，可以使用異常機制處理執行時錯誤，而cuda程式設計中由於Host和Device共同使用，難以利用異常機制，因此，cuda提供了檢測執行時錯誤的機制。

看上面的API時會發現，每個函式的返回值型別都是 cudaError_t ，這正是cuda提供的錯誤檢測機制，如果返回值是cudaSuccess則說明執行正確，否則就是出現了錯誤。可以使用 cudaGetErrorString( error )獲取返回值的代表的錯誤的文字。前面的程式碼中沒有使用這一機制主要是為了便於閱讀，但實際的使用中這一機制是必不可少的，也會看到VS生成的demo程式碼中就包含著大量的錯誤檢測程式碼

	cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

	cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
...
...

5 其他

  1. 不同的block_size計算耗時會不同，可以多嘗試後選擇計算的更快的引數（學DL的調參是吧，這也搞黑盒？）；考慮GPU的計算時間時要考慮資料移入移出GPU的時間。

  2. 不同的GPU有不同的性質，裝置中也可能存在多個GPU，在設計程式時需要考慮這些問題，cuda也提供了訪問這些資訊的介面

    // 獲取裝置數量
int deviceCount = 0;
cudaGetDeviceCount(&deviceCount);

//獲取第一個裝置的各項性質
int devNo = 0;
cudaDeviceProp iProp;
cudaGetDeviceProperties(&iprop, devNo);

相關文章