ピン留め(Page-locked)されたHost-memoryの効果
CUDAによる一連の処理は、
cudaMemcpy
でデータ(input)をHostからDeviceにコピー- kernelを実行: inputを読んで処理してoutputに書く
cudaMemcpy
でデータ(output)をDeviceからHostにコピー
ってゆー一連のダンドリになります。
このとき Host-memoryをmalloc
/new
で確保するのとcudaHostAlloc
/cudaMallocHost
で確保するのとでは動きがかなり異なります。
#include <cuda_runtime.h> #include <device_launch_parameters.h> #include <math_functions.h> __global__ void kernel_sine(const float* angle, float* sine, unsigned int size) { int i = blockDim.x * blockIdx.x + threadIdx.x; if ( i < size ) { sine[i] = sinf(angle[i]); } } #include <cstdlib> #include <iostream> // どちらかひとつを有効にしておくれ #define MALLOC // #define CUDAMALLOCHOST int main() { const unsigned int N = 100000; // Host-memoryを確保 float* h_angle; float* h_sine; #if defined(MALLOC) h_angle = (float*)malloc(N*sizeof(float)); h_sine = (float*)malloc(N*sizeof(float)); #elif defined(CUDAMALLOCHOST) cudaMallocHost(&h_angle, N*sizeof(float)); cudaMallocHost(&h_sine , N*sizeof(float)); #endif // Device-memoryを確保 float* d_angle; float* d_sine; cudaMalloc(&d_angle, N*sizeof(float)); cudaMalloc(&d_sine , N*sizeof(float)); // HtoD, kernel, DtoH を二回実行 cudaMemcpyAsync(d_angle, h_angle, N*sizeof(float), cudaMemcpyDefault); kernel_sine<<<(N+255)/256,256>>>(d_angle, d_sine, N); cudaMemcpyAsync(h_sine , d_sine , N*sizeof(float), cudaMemcpyDefault); cudaMemcpyAsync(d_angle, h_angle, N*sizeof(float), cudaMemcpyDefault); kernel_sine<<<(N+255)/256,256>>>(d_angle, d_sine, N); cudaMemcpyAsync(h_sine , d_sine , N*sizeof(float), cudaMemcpyDefault); cudaStreamSynchronize(nullptr); // defaul-stream上の処理が完了するまで待機 std::cout << "done." << std::endl; // あとしまつ cudaFree(d_angle); cudaFree(d_sine ); #if defined(MALLOC) free(h_angle); free(h_sine ); #elif defined(CUDAMALLOCHOST) cudaFree(h_angle); cudaFree(h_sine ); #endif cudaDeviceReset(); }
実行時のtimelineはそれぞれこんな。
maloc
cudaMallocHost
malloc
だと一回目と二回目との間にスキマができてますが、cudaMallocHost
ではきっちり詰まってますよね。つまりそれだけ速いってことです。
PCI-bsを介してHost-Device間のコピーが行われるとき、DMA(Direct Memory Access)が使われます。CPUの動作とは無関係に(独立して)勝手にメモリをアクセスする機能です。DMAが機能するためにはコピー元/先の物理アドレスが固定されていなくてはなりません。
malloc
/new
で確保された領域は仮想アドレス空間にあり、物理アドレスが定まっていないのでDMAが使えない、そこで一旦 Staging-bufferと呼ばれる固定領域にコピーされた後、DMAが動きます。スキマが空くのはこのため。
一方cudaMallocHost
はハナっから固定された(Page-locked)領域を確保するのでいきなりDMAが動けるってスンポーです。
複数のStreamを使ったmemory-copyとkernelとのoverlapはHost-memoryが固定されていないとその効果が期待できません。
※ ただし、固定域をたくさん確保するとマシン全体のパフォーマンスを落とすことになりかねません。ご利用は計画的に。
cudaMemcpyのコピー方向とUVA
cudaMemcpy の引数は コピー先, コピー元, バイト数, そして コピー方向。 コピー方向は コピー先/元がHostかDeviceのそれぞれに応じて4種類...ともうひとつ cudaMemcpyDefault てのがあります。
cudaMemcpyDefault はコピー元/先に与えたポインタがHost/Deviceのどちらかを判別し善きに計らってくれます。 が cudaMemcpyDefault を使うには「Unified Virtual Addressing をサポートしていること」て但し書きがあります。
UVA(Unified Virtual Addressing) とは、CPUとGPUのメモリ空間を仮想的にひとつの空間に配置できるてゆーよくわかんない機能でして、これが有効じゃないと与えられたポインタがHost/Deviceのどっちか判別できんとのこと。
で、UVAが有効か否かはこんなコードで確認できます:
#include <cuda_runtime.h> #include <device_launch_parameters.h> #include <iostream> using namespace std; int main() { { int device; cudaGetDevice(&device); cudaDeviceProp property; // 現デバイスのプロパティを取得し、 cudaGetDeviceProperties(&property, device); // unifiedAddressing != 0 なら UVA有効 if ( property.unifiedAddressing ) { cout << "UVA enabled, cudaMemcpyDefault can be used." << endl; } else { cout << "sorry, no-UVA" << endl; } } // 試しに cudaMemcpuDefault を使ってみる int* host_ptr; host_ptr = new int[1]; int* device_ptr; cudaMalloc(&device_ptr, sizeof(int)); // host->device *host_ptr = 12345; cudaMemcpy(device_ptr, host_ptr, sizeof(int), cudaMemcpyDefault); // device->host *host_ptr = 0; cudaMemcpy(host_ptr, device_ptr, sizeof(int), cudaMemcpyDefault); if ( *host_ptr == 12345 ) { cout << "ok." << endl; } else { cout << "oops!." << endl; } delete[] host_ptr; cudaFree(device_ptr); cudaDeviceReset(); }
UVAは近頃の大抵のGPUでサポートしてるのですが、32bitだとダメです。32bitではGPU/CPU双方の空間をひとつにするに十分な大きさじゃないんですな。
東方算程譚:只今熱烈引越中
ココからお引越しですー