二、一些基礎
上接: 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...