後付けピン留め - 東方算程譚

cudaMallocHostを使ってピン留めされたHost-memoryを確保すればcudaMemcpyのスキマを潰すことができ、その分速くなる(てかそうしないと複数streamでのoverlapができん)のですが、あらかじめ(mallocやnewで)確保された領域を後付けでピン留め...できるんですわこれが。

cudaHostRegisterってAPIがありまして、領域の先頭アドレスと大きさ(バイト数)を与えることでピン留めしてくれます。ピンを外すのはcudaHostUnregisterで。

#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 HOSTREGISTER

int main() {

  const unsigned int N = 100000;

// Host-memoryを確保
  float* h_angle;
  float* h_sine;
  h_angle = (float*)malloc(N*sizeof(float));
  h_sine  = (float*)malloc(N*sizeof(float));
#if defined(HOSTREGISTER)
　　// 後付けピン留め
  cudaHostRegister(h_angle, N*sizeof(float), cudaHostRegisterDefault);
  cudaHostRegister(h_sine , N*sizeof(float), cudaHostRegisterDefault);
#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(HOSTREGISTER)
  // ピン留め解除
  cudaHostUnregister(h_angle);
  cudaHostUnregister(h_sine );
#endif
  free(h_angle);
  free(h_sine );
  cudaDeviceReset();
}

timelineはこんな↓、スキマ詰まってますね。

f:id:Episteme:20161019184527p:plain