天天看點
傳回

CUDA并行程式設計概述CUDA記憶體與顯存栅格結構函數限定詞異常處理主機同步

前往我的首頁以擷取更好的閱讀體驗

CUDA并行程式設計概述 - DearXuan的首頁

CUDA并行程式設計概述CUDA記憶體與顯存栅格結構函數限定詞異常處理主機同步

https://blog.dearxuan.com/2021/11/15/CUDA%E5%B9%B6%E8%A1%8C%E7%BC%96%E7%A8%8B%E6%A6%82%E8%BF%B0/

CUDA

CUDA是英偉達推出的GPU架構平台，通過GPU強大的并行執行效率，為計算密集型應用加速，CUDA檔案以.cu結尾，支援C++語言編寫，在使用CUDA前需要下載下傳 CUDA Toolkit

記憶體與顯存

CPU可以通路記憶體，GPU可以通路顯存，如果需要使用GPU進行計算，必須把資料從記憶體複制到顯存

指向顯存的指針

建立一個指向顯存的指針，下面的代碼可以告訴你為什麼要使用 (void**)類型

int* p; // 這是一個指向int變量的記憶體指針

function(p); // 如果直接把指針傳入函數，那麼它會以參數的形式被帶入計算，函數中的操作無法修改p的值

function(&p); // 隻要把p的位址傳入函數，函數就可以通過位址修改指針的值

void* v; // 但是指針類型很多，為了統一，通常會使用無類型指針

function(&v); // 是以我們應該傳入一個指向無類型指針位址的指針

(void*)p; // 這樣可以把 p 變成無類型指針，但是我們需要的是指向 p 的位址的無類型指針

(void**)&p; // 這樣我們就得到了指向 p 的位址的無類型指針

function((void**)&p); // 這樣function函數就可以直接修改p的數值

void* p;
function(&p); // 如果你的 p 已經是無類型指針，那麼可以直接使用取址符

在GPU中申請顯存，并獲得指向顯存的指針

int length = size * sizeof(int);

int a[size];
int b[size];
int c[size];

int* dev_a;
int* dev_b;
int* dev_c;

cudaMalloc((void**)&dev_a, length);
cudaMalloc((void**)&dev_b, length);
cudaMalloc((void**)&dev_c, length);

 此時的dev_a, dev_b, dev_c已經指向顯存位址，空間大小為 length

記憶體與顯存的資料交換

在使用GPU計算前，需要把資料複制到顯存

cudaMemcpy(dev_a, a, length, cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, length, cudaMemcpyHostToDevice);

dev_a是顯存指針，a是記憶體指針cudaMemcpyHostToDevice表示把長度為length的記憶體資料複制到顯存裡

計算完成後，需要把資料從顯存複制到記憶體以供CPU計算

cudaMemcpy(c, dev_c, length, cudaMemcpyDeviceToHost);

這段代碼的含義是把dev_c指向的顯存位址的資料複制到c指向的記憶體位址

在計算結束後，應該釋放顯存空間

cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);

栅格結構

GPU的結構包含栅格(grid)，塊(block)，線程(thread)，許多線程組成一個“塊”，許多個“塊”組成一個“栅格”，其中grid和block都可以用三維向量表示，假設一個block有1024個線程，如果建立4個block，則總共有4096個線程同時運作

下面的代碼展示了如何擷取block和thread的編号

int i = threadIdx.x;
int j = blockIdx.x;

合理使用這些編号，可以幫助你定位變量位置

__global__ void DoInKernel(int* a, int* b, int* c) {
    int i = blockIdx.x * 1024 + threadIdx.x;
    c[i] = a[i] + b[i];
}

函數限定詞

核函數

核函數使用 __global__ 修飾，它在CPU上調用，在GPU上執行

__global__ void DoInKernel(int* a, int* b, int* c) {
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main(){
    DoInKernel << <1, size >> > (dev_a, dev_b, dev_c);
}

其中 <<< >>>運算符決定了執行核函數的方式，第一個參數是block的數量，即一個grid裡有幾個block，它實際上是一個dim3類型的變量，在處理多元數組時它可以讓你的代碼編寫更加友善，但是這裡不做示範

dim3 dg(10, 10, 10);
DoInKernel << <dg, size >> > (dev_a, dev_b, dev_c);

第二個參數是thread的數量，即一個block裡有幾個線程，它同樣是dim3類型的變量，如果輸入的是int，則預設y和z都是1

後面還有兩個可選參數，分别用來表示共享記憶體大小和流，共享記憶體大小限制了可以動态配置設定的共享記憶體的最大值，流指定使用哪個IO通道在記憶體和顯存之間複制資料，使用不同的流可以防止阻塞

内聯函數

内聯函數使用 __device__ 修飾，它必須在GPU上調用，隻能在GPU上執行

__device__ int add(int a, int b) {
    return a + b;
}

__global__ void DoInKernel(int* a, int* b, int* c) {
    int i = threadIdx.x;
    c[i] = add(a[i], b[i]);
}

主機函數

所有不加修飾的函數都是主機函數，它也可以使用 __host__ 修飾，主機函數隻能在CPU上調用和執行，例如 main 就是一個主機函數

__host__ int main(){
    return 0;
}

異常處理

CUDA代碼極難調試，是以最好在每一步都檢查一次錯誤，一旦發生錯誤，立即轉到錯誤處理

int main()
{
    //無關代碼
    if (cudaMalloc((void**)&dev_a, length) != cudaSuccess) {
        goto OnError;
    }
    if (cudaMalloc((void**)&dev_b, length) != cudaSuccess) {
        goto OnError;
    }
    if (cudaMalloc((void**)&dev_c, length) != cudaSuccess) {
        goto OnError;
    }
    //無關代碼
    return 0;
OnError:
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
}

主機同步

由于GPU的代碼是異步執行的，如果兩個核函數被寫在一起，那麼它們很可能會被同時執行，使用cudaDeviceSynchronize()阻塞主機線程，可以確定所有的核函數或者IO流都已經執行完畢，才會繼續執行下面的代碼

DoInKernel_1 << <1, 10 >> > ();
cudaDeviceSynchronize();
DoInKernel_2 << <1, 10 >> > ();

其中cudaMemcpy()函數必須在所有GPU任務完成後才能執行，是以它已經自帶主機同步，不再需要手動阻塞

cuda gpu c++ 并行計算
上一篇: cuda并行程式設計基礎（二）
下一篇: windows 下 CUDA 并行程式設計環境搭建

繼續閱讀