リニアメモリとCUDA配列

Linear Memory†

2次元，3次元配列のメモリ領域を確保したとき，メモリ内でのデータの連続性は保証されません． CUDAではデバイスメモリ内の連続した領域(Linear memoryとよぶ)を確保してくれる命令があります (1次元の場合は，通常のcudaMalloc()でLinear memoryが確保されます)．

cudaError_t cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height);
cudaError_t cudaMalloc3D(struct cudaPitchedPtr* pitchedDevPtr, struct cudaExtent extent);

それぞれ2D,3DのLinear memoryを確保します． cudaMallocPitch, cudaMalloc3D参照． 2Dで(width/sizeof())×height，3Dでextent.width×extent.height×extent.depth(byte)の大きさの多次元配列となります． widthはバイト数で指定(width*sizeof(float)など)し，pitchには確保した領域内のピッチ幅(widthのバイト数)が入ります．つまり，(i,j)要素のアドレスは，

devPtr+j*pitch+i

cudaExtent変数はmake_cudaExtent(size_t w, size_t h, size_t d) 関数で初期化可能です．w,h,dはそれぞれwidth, height, depthをバイト数で指定します． cudaPitchedPtrは pitch : ピッチ幅(size_t,バイト数)，ptr : 確保されたメモリの先頭アドレス(void*)， xsize,ysize : 要素の幅と高さ(size_t) のメンバ変数を持つ構造体です．

Linear memoryのコピーには以下の命令を用います．

cudaError_t cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);
cudaError_t cudaMemcpy3D(const struct cudaMemcpy3DParms *p);

参照 : cudaMemcpy2D， cudaMemcpy3D

2Dの場合の引数はそれぞれ，

dst : コピー先アドレス
dpitch : コピー先配列のピッチ幅(バイト数)
src : コピー元アドレス
spitch : コピー元配列のピッチ幅(バイト数)
width : コピーする領域の幅(バイト数)
height : コピーする領域の高さ(要素数)
kind : コピータイプ : cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice です．

3Dの場合は引数に以下の構造体の変数を指定する．

struct cudaMemcpy3DParms {
  struct cudaArray     *srcArray;
  struct cudaPos        srcPos;
  struct cudaPitchedPtr srcPtr;
  struct cudaArray     *dstArray;
  struct cudaPos        dstPos;
  struct cudaPitchedPtr dstPtr;
  struct cudaExtent     extent;
  enum cudaMemcpyKind   kind;
};

↑

CUDA配列†

CUDA配列(CUDA array)はテクスチャフェッチに最適化されたメモリレイアウトでデータが格納される配列です．テクスチャメモリとして用いる方法は，テクスチャメモリを参照． CUDA配列の確保にはcudaMallocArray を用います．

cudaError_t cudaMallocArray(struct cudaArray ** arrayPtr, const struct cudaChannelFormatDesc * desc, size_t width, size_t height);

引数はそれぞれ，

arrayPtr : 確保されたメモリのアドレス
desc : テクスチャのデータ構造を示す構造体
width,height : 確保幅と高さです． cudaChannelFormatDescの定義は，
```
struct cudaChannelFormatDesc{
  int x, y, z, w;
  enum cudaChannelFormatKind f;
};
```
であり，x,y,z,wはテクスチャの各チャネルのビット数(0,8,16,32)， fは型の種類を示し，cudaChannelFormatKindSigned (符号付整数型の場合)，cudaChannelFormatKindUnsigned (符号なし整数型の場合)， cudaChannelFormatKindFloat (浮動小数点型の場合)のいずれかを指定します．cudaChannelFormatDesc構造体の初期化にはcudaCreateChannelDesc関数を用います．

CUDA配列のコピーを行う関数は，コピーの方向に応じて3種類用意されています．

cudaError_t cudaMemcpyToArray(struct cudaArray * dst, size_t wOffset, size_t hOffset, const void * src, size_t count, enum cudaMemcpyKind kind);
cudaError_t cudaMemcpyFromArray(void * dst, const struct cudaArray * src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind);
cudaError_t cudaMemcpyArrayToArray(struct cudaArray * dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray * src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind);

cudaMemcpyToArray は通常のメモリ領域(アドレスsrc)からCUDA配列(アドレスdstから(wOffset,hOffset)だけオフセット)にcountバイト分コピーします． cudaMemcpyFromArray はCUDA配列(アドレスsrcから(wOffset,hOffset)だけオフセット)から通常のメモリ領域(アドレスdst)にcountバイト分コピーします． cudaMemcpyArrayToArray はCUDA配列(アドレスsrcから(wOffsetSrc,hOffsetSrc)だけオフセット)からCUDA配列(アドレスdstから(wOffsetDst,hOffsetDst)だけオフセット)にcountバイト分コピーします． kindにはコピー元，コピー先の場所に対応して， cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDeviceのいずれかを指定します．

Menu

Linear Memory†

CUDA配列†