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);

参照 : cudaMemcpy2DcudaMemcpy3D

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

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

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

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
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の定義は,
      1
      2
      3
      4
    
    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のいずれかを指定します.


トップ   編集 凍結 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2011-10-27 (木) 15:09:18 (1980d)