ピン留め(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