CUDA 儲存器種類分析及使用方法指南

2021-05-25 11:48:54 字數 3308 閱讀 2176

暫存器

區域性儲存器

共享儲存器

全域性儲存器

主機端記憶體

主機端頁鎖定記憶體

常數儲存器

紋理儲存器

儲存器位置擁有快取

訪問許可權

變數生存週期

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認證。網路開放是乙個不變的趨...