CUDAの非同期APIの使い方

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

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

目次

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ミリ秒と、性能向上が確認できる。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:CUDAの非同期APIの使い方

個人用ツール