cuda 本地記憶體使用 cuda實戰入門

cuda (compute unified device architecture) 是 nvidia 所推出的一種平行計算平台和平行計算 api。

cuda 在平行計算上可以大顯神威，因此，我們先要找到乙個可並行的問題。乙個很簡單的可並行問題就是計算無窮級數(infinite series)。圓周率 pi 可以通過乙個著名的無窮級數(leibniz formula)進行計算，具體可檢視wiki。我們可以用此公式來逼近圓周率，c **如下。

#define loop_n 8192000000
double pi() 
return pi_qv * 4;
}

下面我們編譯並執行此程式，並用 time 來測量程式執行時間，在我的機器上差不多需要 8s 吧。

gcc pi_number.c -o3 -opi time ./pi

#include#include#include#define block_num 32   //塊數量
#define thread_num 256 // 每個塊中的執行緒數
#define loop_n block_num * thread_num * 1000000
__global__ void leib_pi(double* g_sum) 
g_sum[bid*thread_num+tid] = tmp;
}int main()
printf("calculate %.10fn", pi_v*4);
cudafree(g_sum);
free(h_sum);
}

下面我們編譯並執行**，結果顯示，程式用了 0.07s 就跑完了，就比上面的 c **快了100多倍吧。

nvcc pi_number.cu -o pi time ./pi

1. 上面的**中首先我們注意到函式 leib_pi 前面加上了 __global__ 標誌，這告訴編譯器，此方法是在 gpu 上執行的。而下面的 main 方法則還是在 cpu 上執行。

2. cpu 是不能直接訪問 gpu 記憶體的，反過來也不行。在 cuda 中，記憶體分為 host memory 和 device memory (也稱作 gpu memory)，而指向 host memory 的指標被稱為 host pointer，指向 device memory 的指標被稱為 device memory。

3. 在上面的**中我們首先定義了兩個指標 h_sum 指向 host memory ，而 g_sum 指向 gpu memory，然後用 malloc 系統呼叫給 h_sum 分配記憶體，用 cudamalloc 給 g_sum 分配 gpu memory。

4. 接下來執行計算

leib_pi<<>>(g_sum);

cuda 通過將一組執行緒組成乙個 thread block，並一次 launch 若干個 thread block 的形式來進行平行計算，形成一種網格式的並行結構。在這個例子中，我們 launch 了 32 個 block，每個 block 包含 256 個 thread。你可以理解為有 8192 個執行緒一起被 launch 出去，並行地開始計算他們各自那部分的內容。每個執行緒在執行中，可以拿它其對應的 blockid 和 threadid，比如第乙個 block 的第乙個執行緒對應的兩個 id 就是(0,0)。

const int tid = threadidx.x;
const int bid = blockidx.x;

我們將每個執行緒計算出來的部分結果存入到 gpu memory 中的對應位置。

5. gpu 計算結束後，返回到 cpu 並呼叫 cudamemcpy，將 gpu memory 拷貝到 host memory，然後在 cpu 上完成剩餘操作。

支援 cuda 的 gpu 都被分成若干個流處理器(streaming multiprocessors)，每個流處理器有若干個 32 位暫存器，共享記憶體， a maximum number of active blocks, and a maximum number of active threads。這些資訊可以在wiki上找到。

比如說你正在使用 2080ti 的 gpu，你可以看到你的 cc (compute capability) 是 7.5，而對應的 max number of threads per block 是 1024，等等。

注意，number of threads per block 應該始終是 32 的倍數，因為 gpu kernel 總是以 32 個執行緒一組來啟動執行緒，也就是說，如果你將上面程式中 thread_num 配置為 50，kernel 會啟動 64 個執行緒(per block)。

此外，我們應當盡可能多的 launch 執行緒，這就需要考慮你在使用什麼 gpu，並根據 gpu 的 cc 來決定具體引數了。

除此之外，還需要考慮暫存器和共享記憶體這些資源（這些超出了本文的範圍）。

ps: 原始碼我放在了 github，如下。

cuda 本地記憶體使用 cuda實戰入門

CUDA記憶體使用

CUDA記憶體簡介

CUDA記憶體拷貝

cuda 本地記憶體使用 cuda實戰入門

CUDA記憶體使用

CUDA記憶體簡介

CUDA記憶體拷貝

相關推薦