Zero Copy - 東方算程譚

CUDA C Best Practices Guide 9.1.3 に Zero Copy てのが出てきます...なにコレ? ってんで調べてみました。

サンプル: float列の各要素に対し、その平方根をを求める処理を書きました:

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

// out[i] = √in[i] where i = 0..size-1
__global__ void kernel_square_root(const float* in, float* out, unsigned int size) {
  unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
  if ( i < size ) {
     out[i]  = sqrtf(in[i]);
  }
}

#include <vector>
#include <numeric>
#include <iostream>

void normal() {
  using namespace std;
  const unsigned int N = 10000U;

  // host領域確保
  float* in  = new float[N];
  float* out = new float[N];

  // device領域確保
  size_t before, after, total;
  cudaMemGetInfo(&before, &total);
  cout << "normal   : device free = " << before << " -> ";
  float*  din; cudaMalloc(&din,  2*N*sizeof(float));
  float* dout; cudaMalloc(&dout, N*sizeof(float));
  cudaMemGetInfo(&after, &total);
  cout << after << " (" << (before-after) << " consumed)" << endl;

  // in[] = { 0.0, 1.0, 2.0, 3.0 ... } 
  iota(in, in+N, 0.0f);

  // host -> device
  cudaMemcpy(din, in, N*sizeof(float), cudaMemcpyHostToDevice);
  // kernel-call
  unsigned int block = 256;
  unsigned int grid = (N + block -1U) / block;
  kernel_square_root<<<grid,block>>>(din, dout, N);
  // device -> host
  cudaMemcpy(out, dout, N*sizeof(float), cudaMemcpyDeviceToHost);

  for ( unsigned int i : { 2, 3, 200, 300} ) {
    cout << out[i] << endl;
  }

  // device領域解放
  cudaFree(din);
  cudaFree(dout);

  // host領域解放
  delete[] in;
  delete[] out;
}

いつものダンドリです:

ホスト/デバイスの双方に領域を確保
ホスト領域に入力データを用意
ホストからデバイスへ(入力データを)コピー
カーネル呼び出し
デバイスからホストへ(出力データを)コピー
ホスト/デバイス領域を解放

一方 zero copy ではデバイス側に領域を確保しません。そのかわり、デバイスから読み書きできる領域をホスト側に確保します。 zero copyによる「float列の平方根」はこんなコードになります:

void zero_copy() {
  using namespace std;
  const unsigned int N = 10000U;

  // host領域確保
  size_t before, after, total;
  cudaMemGetInfo(&before, &total);
  cout << "zero-copy: device free = " << before << " -> ";
  float* in;  cudaHostAlloc(&in,  N*sizeof(float), cudaHostAllocMapped);
  float* out; cudaHostAlloc(&out, N*sizeof(float), cudaHostAllocMapped);

  // hostにマップされたdeviceポインタを取得
  float* din;  cudaHostGetDevicePointer(&din,  (void*)in,  0);
  float* dout; cudaHostGetDevicePointer(&dout, (void*)out, 0);

  cudaMemGetInfo(&after, &total);
  cout << after << " (" << (before-after) << " consumed)" << endl;

  // in[] = { 0.0, 1.0, 2.0, 3.0 ... } 
  iota(in, in+N, 0.0f);

  // kernel-call
  unsigned int block = 256;
  unsigned int grid = (N + block -1U) / block;
  kernel_square_root<<<grid,block>>>(din, dout, N);
  // kernel処理完了を待つ
  cudaStreamSynchronize(nullptr);

  for ( unsigned int i : { 2, 3, 200, 300} ) {
    cout << out[i] << endl;
  }

  // host領域解放
  cudaFreeHost(in);
  cudaFreeHost(out);
}

違いは３つ:

ホスト側の領域確保には cudaHostAlloc( ... cudaHostAllocMapped) を用いる。
デバイス側の領域確保は行わず、cudaHostGetDevicePointer でホスト側領域に対応するデバイス・ポインタを手に入れ、kernelに渡す。
kernel呼び出しのあと、kernelの処理が完了するのを待つ(cudaMemcpyは処理完了を暗黙裡に待つが、それが行われないため)。

両者を実行しました:

int main() {
  // zero copy できるか確認
  cudaDeviceProp prop;
  cudaGetDeviceProperties(&prop, 0);
  if ( !prop.canMapHostMemory ) {
    std::cerr << "sorry, not supported." << std::endl;
    return 0;
  }
  // zero copy を有効にする
  cudaSetDeviceFlags(cudaDeviceMapHost);

  normal();
  zero_copy();
}

f:id:Episteme:20161019203327p:plain