読者です 読者をやめる 読者になる 読者になる

東方算程譚

επιστημηがヨタをこく、弾幕とは無縁のCUDAなタワゴト

ピン留め(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
f:id:Episteme:20161014233550p:plain

cudaMallocHost
f:id:Episteme:20161014233624p:plain

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が固定されていないとその効果が期待できません。

※ ただし、固定域をたくさん確保するとマシン全体のパフォーマンスを落とすことになりかねません。ご利用は計画的に。