CUDA -- 內存分配

 

　　CUDA可以認為是一個由軟件和硬件構成的並行計算系統，其依賴於GPU的並行計算單元，CUDA有類C的API，方便程序編寫。其依賴於CPU和GPU的異構體系，通過在CPU上串行執行環境初始化、內存分配、數據傳輸，然后在GPU上執行並行計算。

內存分配

　　1、一維

int *dev_ans = 0;
cudaMalloc((void**)&dev_ans, d.y * sizeof(int));

　　參數1：顯存中開辟的空間的指針（術語：GPU設備端數據指針）

　　參數2：空間大小，字節為單位

　　2、二維

int *dev_mat = 0;
int pitch;
cudaMallocPitch((void**)&dev_mat, (size_t *)&pitch, d.x * sizeof(int), d.y);

　　參數1：GPU設備端數據指針

　　參數2：一行數據的真實空間大小（字節）【此參數是獲取返回值】，GPU中從256字節對齊的地址（address=0,256,512……）連續訪問最有效率，故每行實際分配的大小要大於需要分配的大小

　　參數3：每行需要分配的空間大小

　　參數4：矩陣行數

 

內存拷貝

　　1、一維

cudaMemcpy(ans, dev_ans, d.y * sizeof(int), cudaMemcpyDeviceToHost);

　　參數1：目標數據地址

　　參數2：源數據地址

　　參數3：數據大小

　　參數4：拷貝類型（主機至主機，主機至設備，設備至主機，設備至設備）

　　2、二維

cudaMemcpy2D(dev_mat, pitch, mat, d.x*sizeof(int), d.x*sizeof(int), d.y, cudaMemcpyHostToDevice);

　　參數1：目標數據地址

　　參數2：pitch，分配空間的行寬（字節單位）

　　參數3：源數據地址

　　參數4：pitch，分配空間的行寬（字節單位）

　　參數5：需要拷貝數據的真實行寬（字節單位）

　　參數6：數據的行數（非字節單位哦！）

　　參數7：數據拷貝類型

　　注：pitch是線性存儲空間的行寬不是數據的行寬，在設備端 pitch大於等於數據行寬，在主機端pitch==數據行寬。

 

內存訪問

　　主機中的內存訪問就是c++的訪存沒什么好說的，現在看看顯存中的訪問方式（也就是在kernel中的訪存）。

__global__ void addKernel(int *mat, int *ans, size_t pitch)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    __shared__ int data[8];
    int *row = (int*)((char*)mat + bid*pitch);
    data[tid] = row[tid];
    __syncthreads();
    for (int i = 4; i > 0; i /= 2) {
        if (tid < i)
            data[tid] = data[tid] + data[tid + i];
        __syncthreads();
    }
    if (tid == 0)
        ans[bid] = data[0];
}

　　一維：

　　　　ans[index]直接訪問

　　二維：

　　　　先計算訪問的行的初始地址　int *row = (int*)((char*)mat + bid*pitch)

　　　　然后訪問此行的對應元素 row[index]

 

內存釋放

cudaFree(dev_mat)