暫存器
區域性儲存器
共享儲存器
全域性儲存器
主機端記憶體
主機端頁鎖定記憶體
常數儲存器
紋理儲存器
儲存器位置擁有快取
訪問許可權
變數生存週期
register
gpu (芯)片內
n/adevice 可讀/寫
與thread相同
local memory
板載視訊記憶體
無device 可讀/寫
與thread相同
shared memory
gpu 片內
n/adevice 可讀/寫
與block相同
constant memory
板載視訊記憶體
有device 可讀,host可讀/寫
可在程式中保持
texture memory
板載視訊記憶體
有device 可讀,host可讀/寫
可在程式中保持
global memory
板載視訊記憶體
無device 可讀/寫,host可讀/寫
可在程式中保持
host memory
host 記憶體
無host 可讀/寫
可在程式中保持
pinned memory
host 記憶體
無host 可讀/寫
可在程式中保持
共享儲存器
示例:共享儲存器的動態與靜態分配與初始化
int main(int argc, char** argv)
__global__ void testkernel(float* g_idata, float* g_odata)
將共享儲存器中的變數宣告為外部陣列時,陣列的大小將在kernel 啟動時確定,通過其執行引數確定。通過這種方式定義的所有變數都開始於相同的位址,因此陣列中的變數的布局必須通過偏移量顯式管理。例:如果希望在動態分配的共享儲存器內獲得與以下**對應的內容:
short array0[128];
float array1[64];
int array2[256];
應該按照下面的方式定義:
extern __shared__ char array;
__device__ void func()
全域性儲存器
視訊記憶體中的全域性儲存器也稱為線性記憶體。線性記憶體通常使用 cudamalloc() 函式分配, cudafree() 函式釋放,並由 cudamemcpy() 進行主機端與裝置端的資料傳輸。通過cuda api分配的空間未經過初始化,初始化全域性儲存器需要呼叫 cudamemset 函式。
對於二維、三維陣列,我們使用 cudamallocpitch() 和 cudamalloc3d() 分配線性儲存空間。這些函式能夠確保分配滿足對齊要求。
例:分配乙個尺寸為 width * height 的 float 型2d 陣列,以及遍歷陣列元素。
//主機端**
float* devptr;
int pitch;
cudamallocpitch((void**)&devptr, &pitch, width * sizeof(float), height);
mykernel<<<100, 512>>>(devptr, pitch);
//裝置端**
__global__ void mykernel(float* devptr, int pitch)}}
例:分配乙個 width * height * depth 的 float 型3d 陣列,以及遍歷陣列元素。
//主機端**
cudapitchedptr devpitchedptr;
cudaextent extent = make_cudaextent(64, 64, 64);
cudamalloc3d(&devpitchedptr, extent);
mykernel<<<100, 512>>>(devpitchedptr, extent);
//裝置端**
__global__ void mykernel(cudapitchedptr devpitchedptr, cudaextent extent)}}
}例:二維陣列和cuda陣列間的資料拷貝。
cudamemcpy2dtoarray(cuarray, 0, 0, devptr, pitch, width * sizeof(float), height, cudamemcpydevicetodevice);
主機端頁鎖定記憶體
通過 cudahostalloc() 和 cudafreehost() 來分配和釋放 pinned memory。
常數儲存器
定義常數儲存器時,需要將其定義在所有函式之外,作用範圍為整個檔案,並且對主機端和裝置端函式都可見。下面兩段**說明了兩種常數儲存器的使用方法。
第一種方法是直接在定義時直接初始化常熟儲存器,然後再kernel裡面直接使用就可以了。
__constant__ int t_hellocuda[11] = ; __constant__ int num = 11
第二種方法是定義乙個costant 陣列,然後使用函式進行賦值。
__constant__ char p_hellocuda[11];
cuda_safe_call(cudamemcpytosymbol(p_hellocuda, hellocuda, sizeof(char) * 11));
紋理儲存器
在kernel 中訪問紋理儲存器的操作稱為紋理拾取(texture fetching)。紋理拾取使用的座標與資料在視訊記憶體中的位置可以不同,我們通過紋理參照系(texture reference)約定二者的對映方式。將視訊記憶體中的資料與紋理參照系關聯的操作,稱為將資料與紋理繫結(texture binding)。視訊記憶體中可以繫結到紋理的資料有兩種,分別是普通的線性儲存器(linear memory)和 cuda 陣列(cuda array)。
紋理儲存器的使用
例:簡單的紋理使用。
/*宣告紋理參照系:texturetexref;
*///2d float texture
texturetexref;
//裝置端**,乙個簡單的轉換kernel
__global__ void transformkernel(float* output, int width, int height, float theta)
//主機端**
int main()
CSS選擇器種類及使用方法
css選擇器 有萬用字元選擇器書寫格式 聲名塊 並集選擇器 組合選擇器 書寫格式 元素或類或id 元素或類或id 元素或類或id 宣告塊 列 h1,h2,h3 color red 層次選擇器 子集選擇器 格式 父級元素名稱 子級元素名稱 例 div p 後代選擇器 格式 祖先元素名稱 空格 後代元素...
OAuth認證協議原理分析及使用方法
本文 http kejibo.com oauth twitter或豆瓣使用者一定會發現,有時候,在別的 點登入後轉到 twitter登入,之後轉回原 你會發現你已經登入此 了,比如像 feedtwitter rss2twitteroauth是乙個開放的認證協議,讓你可以在web或桌面程式中使用簡單而...
OAuth認證協議原理分析及使用方法
twitter或豆瓣使用者一定會發現,有時候,在別的 點登入後轉到 twitter登入,之後轉回原 你會發現你已經登入此 了,比如像 feedtwitter rss2twitter oauth是乙個開放的認證協議,讓你可以在web或桌面程式中使用簡單而標準的,安全的api認證。網路開放是乙個不變的趨...