ページロックホストメモリ

Page-Locked Host Memory†

ホストメモリにmalloc()やnewで領域を確保する代わりに， cudaMallocHost() や cudaHostAlloc() を使うとPage-locked host memoryを確保できます(解放はcudaFreeHost)． このPage-locked host memoryを用いる利点は，

• page-lockedメモリはホストとデバイス間の転送をより高速に実行できる．
• ホストとデバイスでの同時実行をサポートしているデバイスでは，カーネル実行と同時にホスト-デバイスメモリ間のコピーが可能である．
• デバイスによってはpage-lockedメモリはデバイスのアドレス空間にマップされる． 詳しくはProgramming Guideの3.2.5参照．

Page-lockedメモリは，

cudaError_t cudaMallocHost(void ** ptr, size_t size);
cudaError_t cudaHostAlloc(void ** ptr, size_t size, unsigned int flags);

で確保することができます．

ドライバはこの関数で割り当てられた仮想領域をトラックし， cudaMemcpyなどの命令が呼ばれたときに自動的に高速化してくれます． ただし，ホスト側でそのメモリ領域がページングされることも防いでおり， それによりシステム全体のパフォーマンスが低下する可能性もあるので， ホストとデバイス間で頻繁に転送があるデータのみに適用した方がよいです． それぞれ引数は，

• ptr : 確保されたメモリ領域のポインタ
• size : 確保したいサイズ(バイト数)
• flags : 確保メモリの設定

flagsとしては，

• cudaHostAllocDefault: フラグの値としては0となる．
• cudaHostAllocPortable: 上記のPage-lockedメモリの利点は通常割り当てたホストスレッド以外では得られません．全てのスレッドでPage-lockedメモリの恩恵を得るためにこのフラグを使用する．
• cudaHostAllocMapped : デバイスのアドレス空間にPage-lockedホストメモリブロックをマップする． そのため，このブロックは2つのアドレス(ホストとデバイス)を持つ．ホスト側のアドレスはcudaHostAllocの引数ptrで返され， デバイス側はcudaHostGetDevicePointer()で取得することができる．
• cudaHostAllocWriteCombined : write-combining(WC)としてPage-lockedメモリを確保する．WCによりCPUのL1,L2キャッシュが解放され， 他のアプリケーションが利用可能になる．また，WCメモリはPCI Expressバスを通しての転送中に snoopされないので最大40%の高速化となる．ただし，ホストからのメモリ読込みが遅くなるのでホストは書き込むだけ(ホスト->デバイス転送)といった場合のみに用いた方がよい．

を指定できます． cudaMallocHostはcudaHostAllocの最後の引数にcudaHostAllocDefaultを指定したものです．

cudaHostAllocMappedを指定した場合，カーネルから直接ホストメモリにアクセスできます． これにより，デバイスメモリブロックを明示的に確保する必要がなく， ホストメモリブロックとデバイスメモリブロック間のデータ転送はカーネルにより必要に応じて 暗黙に実行されます．また，ストリームを使わなくてもデータ転送をオーバーラップさせることができます．

デバイス側のアドレスを得るためにはcudaHostGetDevicePointer()を用いますが， そのためには，他のCUDA呼び出しが実行される前にcudaSetDeviceFlags()関数でcudaDeviceMapHostフラグをセットしておかなければなりません． また，cudaHostGetDevicePointer()はデバイスがデバイスがpage-lockedに対応していなくてもエラーを返すので， cudaGetDeviceProperties()でcanMapHostMemoryをチェックしておいた方がよいです．