出典: トータル・ディスクロージャ・サイト（事実をありのままに）

先日書いた、Peerアクセスに関する一連の記事で、非同期APIについて、もっと詳しい情報が欲しいという意見があった。そこで、今回は非同期APIについて、簡単に記事にしてみた。

目次 1 CUDAの非同期APIとは 2 簡単な使い方 3 APIの紹介 4 性能評価 4.1 サンプルプログラム 4.2 測定環境 4.3 測定結果

CUDAの非同期APIとは

CUDAの非同期APIとは、CUDA 1.1で追加されたAPIで、データ転送命令などを行う際に、処理の終了を待たないAPIである。このAPIにおいて、カーネル実行やデータ転送命令といった処理を、他の処理の裏で同時実行する枠組みが用意されている。

簡単な使い方

cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, sharedsize, stream>>>(d_a);
cudaMemcpyAsync(h_a, d_a, size, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
cudaStreamDestroy(stream);

cudaStream_t型の変数が表すストリームに対して、カーネルや、名前に「Async」の付いた、一部のメモリ操作命令を実行することができる。同じストリームの処理は直列に処理されるが、違うストリームの処理は並列に実行される可能性がある。ただし、CUDA 4.0では、データ転送を並列実行するためには、ホストメモリはPage-Lockedである必要がある。

非同期APIを使った処理は、CUBLASやCUFFTなどの付属のライブラリでも使用可能なものがある。

APIの紹介

cudaStreamCreate(cudaStream_t *)

引数で渡されたポインタに、新しく作られたストリームを代入する。

cudaStreamQuery(cudaStream_t)

引数で渡されたストリームの状態を、cudaError_t型の返値として返す。ストリームに投入されたタスクが全て完了している場合はcudaSuccessを、未完了のタスクがある場合はcudaErrorNotReadyが返る。

cudaStreamSynchronize(cudaStream_t)

引数で渡されたストリームの、処理の完了を待つ。

cudaStreamDestroy(cudaStream_t)

引数で渡されたストリームを破棄する。

性能評価

非同期APIを使って、カーネルとデータ転送を並列に処理し、性能を向上させるテストを行う。

サンプルプログラム

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#include <curand_kernel.h>

#define N (1000000)
#define M (15)
#define L (100)
#define NTHREADS 512

static __global__ void
kernel(int n, unsigned int *d_a, unsigned long long offset)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    unsigned long long seed = offset + (unsigned long long) i;
    curandState_t state;
    curand_init(seed, 0, 0, &state);
    unsigned int sum = 0;
    for (int j = 0; j < L; ++j) {
      sum ^= curand(&state);
    }
    d_a[i] = sum;
  }
}

void
print(size_t n, unsigned int *a, struct timeval tv[])
{
  size_t i;
  int sec, usec;
  unsigned int sum;

  sec = tv[1].tv_sec - tv[0].tv_sec;
  usec = tv[1].tv_usec - tv[0].tv_usec;
  while (usec < 0) {
    --sec;
    usec += 1000000;
  }
  sum = 0;
  for (i = 0; i < n; ++i) {
    sum ^= a[i];
  }
  printf("%d.%06d\n%08x\n", sec, usec, sum);
}

int
main()
{
  ptrdiff_t i;
  static unsigned int *h_a;
  unsigned int *d_a;
  struct timeval tv[2];
  cudaStream_t stream[M];

  cudaSetDeviceFlags(cudaDeviceScheduleYield);
  cudaHostAlloc((void **)&h_a, M * N * sizeof(unsigned int),
                cudaHostAllocDefault);
  cudaMalloc((void **)&d_a, M * N * sizeof(unsigned int));
  gettimeofday(&tv[0], NULL);
  cudaMemset(d_a, 0, M * N * sizeof(unsigned int));
  cudaMemcpy(h_a, d_a, M * N * sizeof(unsigned int),
             cudaMemcpyDeviceToHost);
  gettimeofday(&tv[1], NULL);
  print(M * N, h_a, tv);
  gettimeofday(&tv[0], NULL);
  kernel<<<(M * N + NTHREADS - 1) / NTHREADS, NTHREADS>>>(M * N, d_a, 0);
  cudaMemcpy(h_a, d_a, M * N * sizeof(unsigned int), cudaMemcpyDeviceToHost);
  gettimeofday(&tv[1], NULL);
  print(M * N, h_a, tv);
  gettimeofday(&tv[0], NULL);
  cudaMemset(d_a, 0, M * N * sizeof(unsigned int));
  cudaMemcpy(h_a, d_a, M * N * sizeof(unsigned int),
             cudaMemcpyDeviceToHost);
  gettimeofday(&tv[1], NULL);
  print(M * N, h_a, tv);
  gettimeofday(&tv[0], NULL);
  for (i = 0; i < M; ++i) {
    cudaStreamCreate(&stream[i]);
  }
  for (i = 0; i < M; ++i) {
    kernel<<<(N + NTHREADS - 1) / NTHREADS, NTHREADS, 0, stream[i]>>>
      (N, &d_a[i * N], (unsigned long long) (i * N));
    cudaMemcpyAsync(&h_a[i * N], &d_a[i * N], N * sizeof(unsigned int),
                    cudaMemcpyDeviceToHost, stream[i]);
  }
  for (i = 0; i < M; ++i) {
    cudaStreamSynchronize(stream[i]);
    cudaStreamDestroy(stream[i]);
  }
  gettimeofday(&tv[1], NULL);
  print(M * N, h_a, tv);
  cudaFree(d_a);
  cudaFreeHost(d_a);

  return 0;
}

CURANDによる、乱数生成処理を、ベンチマーク用のサンプルとして使った。CURANDを使った乱数生成処理のカーネルと、結果をホストに返す（非同期APIを使った）データ転送の処理の組み合わせとを、全データサイズの15分の1に分割し、15個のストリームを作り、それぞれにカーネルと転送の組み合わせのタスクを投げ、最後に同期し、処理を分割しない場合と速度を比較する。あるストリームにおけるデータ転送の裏で、別のストリームがカーネル実行をすることで、性能が向上することが期待される。

測定環境

CPU:	クアッドコア Xeon W3565 3.20 GHz
GPU:	GeForce GTX 580 × 4のうち1枚使用
OS:	CentOS 5.7
CUDA:	CUDA 4.0
GPUドライバ:	バージョン290.10

測定結果

数字は2つペアで、実行時間（秒）とチェックサムである。1つ目と3つ目のペアは、データクリア処理のもので、今回は無視して差し支えない。2つ目のペアが同期APIによる直列処理、4つ目のペアが非同期APIによる並列処理である。

0.009735
00000000
0.046352
57102eea
0.009745
00000000
0.038480
57102eea

同期APIが46ミリ秒に対して、非同期APIが38ミリ秒と、性能向上が確認できる。