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

東方算程譚

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

streamで速くする(2)

Stream

streamはその名の通り"一連の(処理の)流れ"です。 streamに乗せられたjob(仕事)は並んだ順に実行されます。 複数のstreamにjobを乗せておけば、GPUはその中から実行可能なものを一つ選んで処理します。 ホスト・メモリ間のデータ転送にGPUは関与しないので、データ転送の間でも他のstreamに乗ったjobを処理できます。

ここまでにいくつかサンプル・コードを紹介してきましたが、それらの中には stream を明示的に使ったものはありません。 けども実は暗黙の default stream が一本だけあって、その default stream に cudaMemcpy やら kernel呼び出しやらのjobを乗っけていたんです。

3本のstream(ssin, scos, stan)を用意し、sin,cos,tanの表作成をそれぞれの stream に乗せて処理しましょう。

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

// 時間稼ぎ
__device__ float time_waste() {
  float result = 0.0f;
  const int TIMES = 10U;
  for ( int i = 0; i < TIMES; ++i ) {
    result += 1.0f / TIMES;
  }
  return result;
}

// 度→ラジアン
__device__ __forceinline__ float deg2rad(float angle) {
  return angle / 180.0f * 3.01416f * time_waste();
}

// sine
__global__ void kernel_sin(float* out, const float* angle, unsigned int size) {
  unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
  if ( i < size ) {
    out[i] = sinf(deg2rad(angle[i]));
  }
}

// cosine
__global__ void kernel_cos(float* out, const float* angle, unsigned int size) {
  unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
  if ( i < size ) {
    out[i] = cosf(deg2rad(angle[i]));
  }
}

// tangent
__global__ void kernel_tan(float* out, const float* angle, unsigned int size) {
  unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
  if ( i < size ) {
    out[i] = tanf(deg2rad(angle[i]));
  }
}

// helper funcs
template<typename T> inline T* allocate_host(size_t size)
  { T* p = nullptr; cudaMallocHost(&p, size*sizeof(T)); return p; }

template<typename T> inline void free_host(T* host)
  { cudaFreeHost(host); }

template<typename T> inline T* allocate_device(size_t size) 
  { T* p = nullptr; cudaMalloc(&p, size*sizeof(T)); return p; }

template<typename T> inline void free_device(T* device)
  { cudaFree(device); }

inline cudaEvent_t create_event() 
  { cudaEvent_t e = nullptr; cudaEventCreate(&e); return e; };

inline void destroy_event(cudaEvent_t event)
  { cudaEventDestroy(event); }

inline cudaStream_t create_stream() 
  { cudaStream_t s = nullptr; cudaStreamCreate(&s); return s; };

inline void destroy_stream(cudaStream_t stream)
  { cudaStreamDestroy(stream); }

#include <iostream>

using namespace std;

const unsigned int N = 1000U * 1000U;
const unsigned int BN = N*sizeof(float);

void multiple_stream() {
  cout << "--- multiple stream ---" << endl;

  float* angle = allocate_host<float>(N); // in
  float* sin   = allocate_host<float>(N); // out
  float* cos   = allocate_host<float>(N); // out
  float* tan   = allocate_host<float>(N); // out

  for ( unsigned int i = 0; i < N; ++i ) {
    angle[i] = 360.0f * i / N;
  }

  float* dangle = allocate_device<float>(N);
  float* dsin   = allocate_device<float>(N);
  float* dcos   = allocate_device<float>(N);
  float* dtan   = allocate_device<float>(N);

  unsigned int size  = N;
  unsigned int block = 128U;
  unsigned int grid  = (size + block -1U) / block;

  cudaEvent_t start = create_event();
  cudaEvent_t mid   = create_event();
  cudaEvent_t stop  = create_event();

  cudaStream_t ssin = create_stream();
  cudaStream_t scos = create_stream();
  cudaStream_t stan = create_stream();

  cudaEventRecord(start);
  cudaMemcpy(dangle, angle, BN, cudaMemcpyHostToDevice);
  cudaEventRecord(mid);

  kernel_sin<<<grid,block,0,ssin>>>(dsin, dangle, size);
  cudaMemcpyAsync(sin, dsin, BN, cudaMemcpyDeviceToHost, ssin);

  kernel_cos<<<grid,block,0,scos>>>(dcos, dangle, size);
  cudaMemcpyAsync(cos, dcos, BN, cudaMemcpyDeviceToHost, scos);

  kernel_tan<<<grid,block,0,stan>>>(dtan, dangle, size);
  cudaMemcpyAsync(tan, dtan, BN, cudaMemcpyDeviceToHost, stan);

  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  float elapsed;
  cudaEventElapsedTime(&elapsed, start, mid);
  cout << elapsed << " + ";
  cudaEventElapsedTime(&elapsed, mid, stop);
  cout << elapsed << " = ";
  cudaEventElapsedTime(&elapsed, start, stop);
  cout << elapsed << " [ms]" << endl;
  for ( unsigned int i = 0; i < 5; ++i ) {
    cout << angle[i] << " :"
         << " sin=" << sin[i] 
         << " cos=" << cos[i] 
         << " tan=" << tan[i] 
         << endl; 
  }
  cout << "..." << endl;

  for ( auto event  : { start, mid,  stop }) 
    destroy_event(event);
  for ( auto stream : { ssin,  scos, stan }) 
    destroy_stream(stream);
  for ( auto host   : { angle,  sin,  cos,  tan  }) 
    free_host(host);
  for ( auto device : { dangle, dsin, dcos, dtan }) 
    free_device(device);

}

int main() {
  multiple_stream();
}

default stream 1本で処理した時のタイムラインがコレ。

f:id:Episteme:20161021220448p:plain

対して3本(+default)のstreamでやったのがコチラ。

f:id:Episteme:20161021220455p:plain

ね、データ転送とkernel実行が同時に行われてるでしょ。重なった分だけ速くなるんです。

実装上のキモはふたつ:

  • cudaMemcpyじゃなく cudaMemcpyAsyncを使うべし。 後者はコピー完了を待たずに(job投入後ただちに)返って来るので、すぐさま次のjobを投入できます。

  • ホスト側メモリを malloc/freenew/delete じゃなく cudaMallocHost/cudaFreeHost で確保/解放すべし。 そうしないとデータ転送中に kernelが動いてくれないのですよ。

... streamのおハナシはまだまだ続きます。

(original: 2015-04-16)