CUDA線性記憶體分配

2022-08-13 17:30:13 字數 3103 閱讀 6614

概述:線性儲存器可以通過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線性記憶體

1

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 二 記憶體分配函式

除了前面提到的記憶體分配函式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...