首先要了解什麼是異構架構計算：即gpu和cpu協同工作，cpu端稱為主機端用host表示，gpu端稱為裝置端用device表示。gpu和cpu連線一般協議是pci-e，最新的協議有nvme，延遲更小。
程式執行流程主要分為六個大的部分：
在host端分配記憶體，進行資料初始化。
在device端分配記憶體。
將資料從host拷貝到device。
用cuda核函式在device端完成指定的運算。
將運算結果從device端拷貝回host端。
將host和device上之前分配的記憶體釋放掉。
①，定義執行緒的資料型別
gpu的sm裡面包含很多grid，每個grid包含很多block，每個block包含很多thread。
在分配執行緒時用到的資料型別定義是dim3,dim3就相當於乙個三維的無符號整型。想象一下，sm裡面存在乙個三維的grid，每個grid裡面存在乙個三維的block，每個block裡面存在乙個三維的tread，當然我們可以根據需要定義一維和二維的grid和block，不同的維度執行緒數量計算也會有不一樣。
比如：
dim3 grid(3
,2);
//這裡的3,2分別代表x方向上和y方向上的block數量，即兩行三列
表示我們所分配的每個grid裡面都是如下所示的block索引：（0,
0）（1
,0）（2,0）
（0,1）（1
,1）（2
,1）
dim3 block(4
,3)；對應thread索引：（0,
0） （1
,0） （2
,0） （3,0）
（0,1） （1
,1） （2
,1） （3
,1）
③，執行緒索引計算
在實際使用時，我們所分配的block數量和grid數量都比較多，並且分配的執行緒不止一維，因此需要找到乙個對應關係。類似於在c語言中，我們訪問二維向量時也是從第一行訪問完然後第二行再依次訪問。
比如我們現在分配乙個一維的grid和一維的thread，來計算兩個長度為6的一維向量相加
那麼kernel函式中就應該
有點類似於我們面前有乙個二維矩陣，每一行都是1到10的數字，總共八行。我們需要找到某個數字在全域性中為「第幾個」，就得先計算：「這個數字前面有多少行」×「每一行滿的時候多少個數字」+該數字為第幾列。
再比如我們現在需要知道下圖中紅色位置執行緒在全域性的索引，我們分配的是一維的grid，二維的block
我們首先需要計算該執行緒在「最後乙個塊中的索引」，即「該執行緒前面有多少行」×「每一行滿的時候多少個執行緒」+「該執行緒為第幾列」。，也就是threadidx.y * blockdim.x + threadidx.y
然後用這個值加上前面滿格block中的執行緒數量即：blockdim.x * blockdim.y * blockidx.x
完整**如下：
__global__ void
testblockthread2
(int
*c,const
int*a,
const
int*b)
#include
"cuda_runtime.h"
#include
"device_launch_parameters.h"
#include
__global__ void
addkernel
(int
*c,const
int*a,
const
int*b,
int n)
}int
main()
;// 初始化host端資料
for(
int i =
0; i < n; i++
)// 申請device記憶體
int*d_x,
*d_y,
*d_z;
cudamalloc((
void**
)&d_x, n*
sizeof
(int))
;cudamalloc((
void**
)&d_y, n *
sizeof
(int))
;cudamalloc((
void**
)&d_z, n *
sizeof
(int))
;// 將host資料拷貝到device
cudamemcpy((
void
*)d_x,
(void
*)h_a, n *
sizeof
(int
), cudamemcpyhosttodevice)
;cudamemcpy((
void
*)d_y,
(void
*)h_b, n *
sizeof
(int
), cudamemcpyhosttodevice)
;// 定義kernel的執行配置
dim3 blocksize(5
,3);
dim3 gridsize(3
);// 執行kernel
addkernel <<
< gridsize, blocksize >>
>
(d_z, d_x, d_y, n)
;// 將device得到的結果拷貝到host
cudamemcpy((
void
*)h_c,
(void
*)d_z, n *
sizeof
(int
), cudamemcpydevicetohost)
;//cudadevicesynchronize();
//輸出結果
for(
int i =
0; i < n;i++
)// 釋放device記憶體
cudafree
(d_x)
;cudafree
(d_y)
;cudafree
(d_z)
;// 釋放host記憶體
free
(h_a)
;free
(h_b)
;free
(h_c)
;return0;
}
CUDA學習筆記之 CUDA儲存器模型
gpu片內 register，shared memory 板載視訊記憶體 local memory,constant memory,texture memory,texture memory,global memory host 記憶體 host memory,pinned memory.cuda儲...
CUDA學習筆記之 CUDA儲存器模型
標籤 cuda 儲存binding cache 程式設計api 2010 01 03 20 32 5577人閱讀收藏 舉報 cuda 6 cuda儲存器模型 gpu片內 register，shared memory 板載視訊記憶體 local memory,constant memory,textu...
CUDA學習筆記之 CUDA儲存器模型
標籤 cuda 儲存binding cache 程式設計api 2010 12 14 01 33 1223人閱讀收藏 舉報 cuda 26 gpu片內 register，shared memory 板載視訊記憶體 local memory,constant memory,texture memory...