概述:線性儲存器可以通過cudamalloc()、cudamallocpitch()和cudamalloc3d()分配
1、1d線性記憶體分配
1 cudamalloc(void**,int) //在裝置端分配記憶體
2 cudamemcpy(void* dest,void* source,int size,enum direction) //
資料拷貝
3 cudamemcpytosymbol //
將資料複製到__constant__變數中,或者__device__變數中
4 cudamemcpyfromsynbol //
同上相反
5 cudafree() //
記憶體釋放
6 cudamemset() //
記憶體初始化
注意:主機和裝置間的資料交換會自動同步,而裝置與裝置卻不會,需要使用cudathreadsynchronize()
2、2d線性記憶體分配
2.1 分配
1 cudamallocpitch( void** devptr,size_t* pitch,size_t widthinbytes,size_t height ) //
c語言申請2維記憶體時,一般是連續存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)個位元組。
但在cuda的global memory訪問中,從256位元組對齊的位址(addr=0, 256, 512, ...)開始的連續訪問是最有效率的。這樣,為了提高記憶體訪問的效率,有了cudamallocpitch函式。cudamallocpitch函式分配的記憶體中,陣列的每一行的第乙個元素的開始位址都保證是對齊的。因為每行有多少個資料是不確定的,widthofx*sizeof(元素)不一定是256的倍數。故此,為保證陣列的每一行的第乙個元素的開始位址對齊,cudamallocpitch在分配記憶體時,每行會多分配一些位元組,以保證widthofx*sizeof(元素)+多分配的位元組是256的倍數(對齊)。這樣,上面的y*widthofx*sizeof(元素)+x*sizeof(元素)來計算a[y][x]的位址就不正確了。而應該是y*[widthofx*sizeof(元素)+多分配的位元組]+x*sizeof(元素)。而函式中返回的pitch的值就是widthofx*sizeof(元素)+多分配的位元組。說明:widthinbytes作為輸入引數,應該是widthofx*sizeof(元素);這樣的話,複製內容時也要作相應的修改。
2.2 訪問
1 t* pelement = (t*)((char*)baseaddress + row * pitch) + column; //元素訪問方式
cudamallocpitch()以*pitch的形式返回間距,即所分配儲存器的寬度,以位元組為單位。間距用作儲存器分配的乙個獨立引數,用於在2d陣列內計算位址。
2.3 拷貝
1 cudamemcpy2d( void* dst,size_t dpitch,constvoid* src,size_t spitch,size_t width,size_t height,enum cudamemcpykind kind )
這裡需要特別注意width與pitch的區別,width是實際需要拷貝的資料寬度而pitch是2d線性儲存空間分配時對齊的行寬,而當資料傳遞發生在裝置與主機之間時,主機端pitch==width.
綜上我們可以看到,cuda下對二維線性空間的訪問是不提供多下標支援的,訪問時依然是通過計算偏移量得到,不同的地方在於使用pitch對齊後非常利於實現coalesce訪
問例:下面的**分配了乙個尺寸為width*height的二維浮點陣列,同時演示了怎樣在裝置**中遍歷陣列元素
1//host code
2int width = 64, height = 64;3
float*devptr;
4int
pitch;
5 cudamallocpitch((void**)&devptr, &pitch, width * sizeof(float
), height);
6 mykernel<<<100, 512>>>(devptr, pitch, width, height);7//
device code
8 __global__ void mykernel(float* devptr, int pitch, int width, int
height)14}
15 }
3、3d線性記憶體
1cudaerror_t cudamalloc3d(
2struct cudapitchedptr *pitcheddevptr,
3struct
cudaextent extent
4 )
例:下面的**分配了乙個尺寸為width*height*depth的三維浮點陣列,同時演示了怎樣在裝置**中遍歷陣列元素
1//host code
2cudapitchedptr devpitchedptr;
3 cudaextent extent = make_cudaextent(64, 64, 64
);4 cudamalloc3d(&devpitchedptr, extent);
5 mykernel<<<100, 512>>>(devpitchedptr, extent);6//
device code
7 __global__ void
mykernel(cudapitchedptr devpitchedptr, cudaextent extent) 17}
18 }
CUDA 二 記憶體分配函式
除了前面提到的記憶體分配函式cudamalloc之外,這裡再簡單的介紹幾個常用的記憶體分配函式 cudamallocpitch cudamalloc3d等。cudaerror t cudamallocpitch void devptr,size t pitch,size t width,size t...
CUDA之靜態 動態共享記憶體分配詳解
include global void staticreverse int d,int n global void dynamicreverse int d,int n int main void int d d cudamalloc d d,n sizeof int run version wit...
CUDA下在Host端分配的幾種記憶體模式
覃裡 it168 文件 pageable vs pinned 平時我們使用的記憶體都是pageable 交換頁 的,而另乙個模式就是pinned page locked 實質是強制讓系統在物理記憶體中完成記憶體申請和釋放的工作,不參與頁交換,從而提高系統效率,需要使用cudahostalloc和cu...