*Linear Memory [#lc2b8c0b]
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:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_g80d689bc903792f906e49be4a0b6d8db.html]],
[[cudaMalloc3D:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_gc1372614eb614f4689fbb82b4692d30a.html#gc1372614eb614f4689fbb82b4692d30a]]参照．
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):http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/online/group__CUDART__MEMORY_g6ce5d637817b7693f6c2601aef21a294.html]]
関数で初期化可能です．w,h,dはそれぞれwidth, height, depthをバイト数で指定します．
[[cudaPitchedPtr:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/structcudaPitchedPtr.html]]は
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:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_g17f3a55e8c9aef5f90b67cdf22851375.html#g17f3a55e8c9aef5f90b67cdf22851375]]，
[[cudaMemcpy3D:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_gc1372614eb614f4689fbb82b4692d30a.html#gc1372614eb614f4689fbb82b4692d30a]]
2Dの場合の引数はそれぞれ，
-dst : コピー先アドレス
-dpitch : コピー先配列のピッチ幅(バイト数)
-src : コピー元アドレス
-spitch : コピー元配列のピッチ幅(バイト数)
-width : コピーする領域の幅(バイト数)
-height : コピーする領域の高さ(要素数)
-kind : コピータイプ : cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice
です．
3Dの場合は引数に以下の構造体の変数を指定する．
#code(C){{
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配列 [#p197a9fc]
CUDA配列(CUDA array)はテクスチャフェッチに最適化されたメモリレイアウトでデータが格納される配列です．
テクスチャメモリとして用いる方法は，[[テクスチャメモリ]]を参照．
CUDA配列の確保には[[cudaMallocArray:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_ge56101fe6f1ce0b48f163632f6862ae4.html#ge56101fe6f1ce0b48f163632f6862ae4]]
を用います．
cudaError_t cudaMallocArray(struct cudaArray ** arrayPtr, const struct cudaChannelFormatDesc * desc, size_t width, size_t height);
引数はそれぞれ，
-arrayPtr : 確保されたメモリのアドレス
-desc : テクスチャのデータ構造を示す構造体
-width,height : 確保幅と高さ
です．
[[cudaChannelFormatDesc:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/structcudaChannelFormatDesc.html]]の定義は，
#code(C){{
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:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_g06db36948e3ccda65d1adf3529420696.html#g06db36948e3ccda65d1adf3529420696]]
は通常のメモリ領域(アドレスsrc)からCUDA配列(アドレスdstから(wOffset,hOffset)だけオフセット)にcountバイト分コピーします．
[[cudaMemcpyFromArray:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_g1620c76fb3337df8dc7186fd88f40b1a.html#g1620c76fb3337df8dc7186fd88f40b1a]]
はCUDA配列(アドレスsrcから(wOffset,hOffset)だけオフセット)から通常のメモリ領域(アドレスdst)にcountバイト分コピーします．
[[cudaMemcpyArrayToArray:http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_ga37909c2448136db5829682d49bce8d9.html#ga37909c2448136db5829682d49bce8d9]]
はCUDA配列(アドレスsrcから(wOffsetSrc,hOffsetSrc)だけオフセット)からCUDA配列(アドレスdstから(wOffsetDst,hOffsetDst)だけオフセット)にcountバイト分コピーします．
kindにはコピー元，コピー先の場所に対応して，
cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDeviceのいずれかを指定します．
![[PukiWiki]](image/pukiwiki.png "[PukiWiki]")
