streamで速くする(1) ～イントロ

三角関数表を作りましょう。ひと回り360度を百万等分し、360マイクロ度刻みの sin, cos, tan を求めます。

#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); }

#include <iostream>

using namespace std;

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

void single_stream() {
  cout << "--- single stream ---" << endl;

  // host memories
  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;
  }

  // device memories
  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;

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

  cudaEventRecord(start);
  // copy angle[] from host to device
  cudaMemcpy(dangle, angle, BN, cudaMemcpyHostToDevice);
  cudaEventRecord(mid);

  // compute sine
  kernel_sin<<<grid,block>>>(dsin, dangle, size);
  // write-back to host
  cudaMemcpy(sin, dsin, BN, cudaMemcpyDeviceToHost);

  // compute cosine
  kernel_cos<<<grid,block>>>(dcos, dangle, size);
  // write-back to host
  cudaMemcpy(cos, dcos, BN, cudaMemcpyDeviceToHost);

  // compute tangent
  kernel_tan<<<grid,block>>>(dtan, dangle, size);
  // write-back to host
  cudaMemcpy(tan, dtan, BN, cudaMemcpyDeviceToHost);

  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 host   : { angle,  sin,  cos,  tan  }) 
    free_host(host);
  for ( auto device : { dangle, dsin, dcos, dtan }) 
    free_device(device);

}

int main() {
  single_stream();
}

なんともオーソドックスなコードです。入力(angle)をホストからデバイスに送り、"kernel_うんちゃらで計算して、デバイスからホストに書き戻す"を三回やってるだけ。

CUDAを使う限り、ホスト・デバイス間のデータ転送は避けられません。そのことが大きなボトルネックになるのは事実ですが、だからってこれ以上速くはならんかというと、そうでもない。

上記コードでは、ホスト・デバイス間のデータ転送をやってる間、GPUは仕事せずデータ転送が終わるのを待ってます。

ピザが焼け、バイトの兄ちゃんがバイク飛ばして配達に行って、帰って来るまでピザ窯空けたままぼーっと待ってるあほゥなピザ職人はいないでしょ？窯が空いたら次の注文受けて焼いとけば、兄ちゃんが配達から帰って来次第すぐさま次の配達できるやないですか。

...てなオハナシを次のエントリで。

(original: 2015-04-15)