cuda並行程式設計基礎(二)

2021-08-13 19:34:00 字數 3351 閱讀 4072

二、一些基礎

上接: cuda並行程式設計基礎(一)

(二)cuda基礎:gridblock.cu

4.總計算量與block、grid的關係

假設一維陣列總計算量為total_c,怎麼確定block與grid呢?

一維陣列:

int cal_array[xx];

int thd_num=16;

dim3 block(thd_num);//一般block.x * block.y * block.z <=1024,1024這個值根據顯示卡效能確定

dim3 grid((xx + thd_num -1)/thd_num);//這裡為何這麼計算呢?

之所以grid.x = xx + thd_num - 1 /thd_num,原因是要保證block*grid >=xx

二維陣列:

int cal_array[xx][yy];

int thd_x_num=16,thd_y_num=32;

dim3 block(thd_x_num,thd_y_num);//一般block.x * block.y * block.z <=1024,1024這個值根據顯示卡效能確定

dim3 grid((xx + thd_x_num -1)/thd_x_num,(yy + thd_y_num -1)/thd_y_num);//原因同一維陣列

依次類推,讀者可以確定三維陣列的計算方法,至於思維陣列怎麼辦?這個嘿嘿,自己想辦法啦。

問題來了,這樣的話,grid和block肯定有很多分配方法,對計算會有什麼影響嗎?答案是肯定的,別著急,以後再學啦,

5.記憶體管理

當然是gpu的記憶體,跟cpu一樣,gpu的記憶體資源也有限制,申請記憶體太多會爆掉哦,

gpu記憶體分為共享記憶體和全域性記憶體,所謂共享記憶體嘛,就是執行緒可以共享,全域性嘛,所有都能訪問咯

cpu:malloc  memcpy  memset  free

gpu: cudamalloc  cudamemcpy  cudamemset  cudafree

cudaerror_t cudamalloc(void **devptr,size_t size);//只能對gpu資料

cudaerror_t cudamemcpy(void *dst,const void *src,size_t count,cudamemcpykind kind);//可以對gpu也可以對cpu或混合

cudaerror_t cudamemset(void *devptr,int value,size_t count);//只能對gpu資料

cudaerror_t cudafree(void *devptr);//只能對gpu資料

可以看到,gpu的記憶體管理函式相比cpu,無非加了「cuda」,還是很簡單的。

cudaerror_t:正確的話返回cudasuccess,錯誤返回cudaerrormemoryallocation,你可以參看"cuda並行程式設計基礎(一)",或者後面單獨對錯誤處理的章節

cudamemcpykind可以有四個值哦:

cudamemcpyhosttohost //cpu到cpu,相當與memcpy

cudamemcpyhosttodevice //cpu到gpu

cudamemcpydevicetohost //gpu到cpu

cudamemcpydevicetodevice //gpu到gpu  

第三個程式,比較重要但很簡單

和grid的分配方法,一維陣列

第四個程式,很重要哦

記憶體管理

/*authored by alpc40*/

#define num 9//陣列大小

//核心函式定義

__global__ void sumarrayongpu(int *d_arraya, int *d_arrayb, int *d_result)

int main()

; int arrayb[num] = ;

int result[num];//記錄cpu結果

int *d_arraya, *d_arrayb,*d_result;

int nbytes=num * sizeof(int);//注意拷貝的是位元組數,不是個數

check(cudamalloc((int **)&d_arraya, nbytes));//如果覺得check比較煩人,而你不認為這個函式會出錯,那麼刪掉又何妨

check(cudamalloc((int **)&d_arrayb, nbytes));//注意申請gpu記憶體的方法

check(cudamalloc((int **)&d_result, nbytes));

check(cudamemcpy(d_arraya, arraya, nbytes, cudamemcpyhosttodevice));//注意拷貝的方法

check(cudamemcpy(d_arrayb, arrayb, nbytes, cudamemcpyhosttodevice));

int thd_num = 4;

dim3 block(thd_num);//block與grid的定義,結果為(4,1,1),即每個記憶體塊有執行緒4個

dim3 grid((num + thd_num -1)/thd_num);//結果為(3,1,1),即記憶體塊為3個

sumarrayongpu << > > (d_arraya,d_arrayb,d_result);//核心函式

check(cudamemcpy(result, d_result, nbytes, cudamemcpydevicetohost));

printf("計算結果:\n");

for (int i = 0; i < num; i++)

printf(" %d",result[i]);

printf("\n");

cudafree(d_arraya);//**gpu記憶體,一般申請時立馬寫**函式,防止忘記導致記憶體洩露

cudafree(d_arrayb);

cudafree(d_result);

cudadevicereset();//gpu恢復

cuda程式設計 CUDA程式設計入門(四)並行歸約演算法

這一篇我們一起學習一下如何使用cuda實現並行歸約演算法。首先我們要知道什麼是並行歸約。並行歸約 reduction 是一種很基礎的並行演算法,簡單來說,我們有n個輸入資料,使用乙個符合結合律的二元操作符作用其上,最終生成1個結果。這個二元操作符可以是求和 取最大 取最小 平方 邏輯與或等等。我們以...

CUDA程式設計實戰 並行向量求和

多個並行執行緒塊完成兩個向量的求和 如下 使用了10個並行執行緒塊 include include book.h using namespace std define n 10 global void add int a,int b,int c int main void int a n b n c...

CUDA 程式設計 之平行計算思想

思想這個東西,是個人理解的問題。無論是 mpi openmp 等等平行計算的方法,都是使用多執行緒同時併發的執行。所謂併發,就是看起來一起執行罷了,在真正的單核cpu中,是在某段時間內,序列執行,通過執行緒排程來掩蓋其執行的順序。那麼cuda 程式設計中,平行計算的思想是simt,instructi...