CUDA線性記憶體分配

概述：線性儲存器可以通過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，const

void* 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線性記憶體

cudaerror_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線性記憶體分配

CUDA 二記憶體分配函式

CUDA之靜態動態共享記憶體分配詳解

CUDA下在Host端分配的幾種記憶體模式

CUDA線性記憶體分配

CUDA 二 記憶體分配函式

CUDA之靜態 動態共享記憶體分配詳解

CUDA下在Host端分配的幾種記憶體模式

相關推薦

CUDA 二記憶體分配函式

CUDA之靜態動態共享記憶體分配詳解