マルチGPUプログラミングにおけるコンテキスト切り替え処理の改善について

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

目次

概要

比較的古い資料には、「CUDAでマルチGPUをプログラミングするためには、ホストで複数スレッドを起動し、各ホストスレッドにそれぞれGPUを割り当てる必要がある。」というような記述がされているものがある。昔のCUDAでは、確かにそうなのであるが、今のCUDAは、もはやそういう仕様になっていない。

CUDAのプログラミングモデルは、GPUを指定してコンテキストを初期化し、そのコンテキストを通して各種処理を行う、というものである。この点に関しては、今も昔も変わらないのであるが、3.1よりも昔のCUDAは、同じスレッドが複数のコンテキストを持つことができなかった。無理やりシングルスレッドでマルチGPUを扱おうとすれば、コンテキストを破棄して、新しいコンテキストを初期化する、という手間をとる必要があり、この「破棄→初期化」の処理に大きなオーバーヘッドが起き、現実的な速度でマルチGPU計算を行うことは、不可能であった。

CUDA 3.1で、以下のAPIがドライバAPIに提供される。

スレッドで複数コンテキストの初期化ができるようになり、スレッドがコンテキストをスタックの形で保持できるようになった。スタックの先頭にあるコンテキストが、処理に使われるコンテキストとなり、cuCtxPushCurrentで、指定したコンテキストをスタックの先頭に追加し、cuCtxPopCurrentで、スタックの先頭のコンテキストを切り離すことで、コンテキストの切り替えを行う。このAPIを使えば、小さなオーバーヘッドで、シングルスレッドによるマルチGPU計算を行うことができる。

ただし、この時点では、ドライバAPIよりも広く使われているであろうランタイムAPIに、対応する処理が実装されなかった。記述量の多いドライバAPIで、さらにリソース管理の大変なスタックを扱ってマルチGPU向けの処理を書くよりも、ランタイムAPIにマルチスレッドの処理を追加したほうが、簡潔な記述になることが多いため、この時点ではあまりメリットは大きくなかった。

CUDA 4.0において、より汎用的なAPI関数が、2種類提供された。

もはやスタックとは関係無く、処理に使われるコンテキストを直接読み書きできるようになった。さらにランタイムAPIのcudaSetDeviceも、このドライバAPIを呼び出すようになったことで、ランタイムAPIを普通に使うだけで、シングルスレッドによるマルチGPU処理が、現実的な速度で実現できるようになった。

プログラム例

#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <cblas.h>

#define N (16384)
#define NB (16)

#define NGPU (4)

static __global__ void
square_matrix(float (*dst)[N], const float (*src)[N], ptrdiff_t i0)
{
  ptrdiff_t i = blockIdx.y * NB;
  ptrdiff_t j = blockIdx.x * NB;
  float dstij = 0.0f;
  for (ptrdiff_t k = 0; k < N; k += NB) {
    __shared__ float s_a[NB][NB], s_b[NB][NB];
    s_a[threadIdx.y][threadIdx.x] = src[i + i0 + threadIdx.y][k + threadIdx.x];
    s_b[threadIdx.y][threadIdx.x] = src[k + threadIdx.y][j + threadIdx.x];
    __syncthreads();
    for (ptrdiff_t kk = 0; kk < NB; ++kk) {
      dstij += s_a[threadIdx.y][kk] * s_b[kk][threadIdx.x];
    }
    __syncthreads();
  }
  dst[i + threadIdx.y][j + threadIdx.x] = dstij;
}

int
main()
{
  float (*h_src)[N], (*h_dst)[N], (*h_gold)[N];
  h_src = new float[N][N];
  for (ptrdiff_t i = 0; i < N; ++i) {
    for (ptrdiff_t j = 0; j < N; ++j) {
      h_src[i][j] = std::rand() / (float) RAND_MAX;
    }
  }
  h_dst = new float[N][N];
  h_gold = new float[N][N];

  std::time_t start, finish;
  std::time(&start);
  float (*d_src[NGPU])[N], (*d_dst[NGPU])[N];
  for (ptrdiff_t i = 0; i < NGPU; ++i) {
    ptrdiff_t begin = i * (N / NB) / NGPU;
    ptrdiff_t end = (i + 1) * (N / NB) / NGPU;
    ptrdiff_t width = end - begin;
    cudaSetDevice(i);
    cudaMalloc((void **)&d_src[i], N * N * sizeof(float));
    cudaMalloc((void **)&d_dst[i], width * NB * N * sizeof(float));
    cudaMemcpy(d_src[i], h_src, N * N * sizeof(float),
               cudaMemcpyHostToDevice);
  }
  for (ptrdiff_t i = 0; i < NGPU; ++i) {
    ptrdiff_t begin = i * (N / NB) / NGPU;
    ptrdiff_t end = (i + 1) * (N / NB) / NGPU;
    ptrdiff_t width = end - begin;
    cudaSetDevice(i);
    square_matrix<<<dim3(N / NB, width), dim3(NB, NB)>>>
      (d_dst[i], d_src[i], begin * NB);
  }
  for (ptrdiff_t i = 0; i < NGPU; ++i) {
    ptrdiff_t begin = i * (N / NB) / NGPU;
    ptrdiff_t end = (i + 1) * (N / NB) / NGPU;
    ptrdiff_t width = end - begin;
    cudaSetDevice(i);
    cudaMemcpy(h_dst[begin * NB], d_dst[i], width * NB * N * sizeof(float),
               cudaMemcpyDeviceToHost);
    cudaFree(d_dst[i]);
    cudaFree(d_src[i]);
  }
  std::time(&finish);
  std::printf("Elapsed time: %.6f seconds.\n", std::difftime(finish, start));

  cblas_sgemm(CblasRowMajor, CblasNoTrans, CblasNoTrans,
              N, N, N, 1.0f, h_src[0], N, h_src[0], N, 0.0f, h_gold[0], N);
  double error = 0.0;
  for (ptrdiff_t i = 0; i < N; ++i) {
    for (ptrdiff_t j = 0; j < N; ++j) {
      error += std::fabs((double) h_dst[i][j] - h_gold[i][j]);
    }
  }
  std::printf("Error: %.6f.\n", error / (N * N));

  return 0;
}

ランタイムAPIを使った、簡単な行列積プログラムであるが、カーネル呼び出し(CUDAのデフォルトでは非同期実行)を単純なループで回した、マルチGPUのプログラムである。CUDA 4.1で動くTesla M2090×4を積んだ、Xeon X5690搭載のWindows 7 Professional機での実行結果は、以下のようになった。

ファイル:20120223gpubmt.png

1 GPU

Elapsed time: 41.000000 seconds.
Error: 0.005488.

2 GPUs

Elapsed time: 20.000000 seconds.
Error: 0.005488.

3 GPUs

Elapsed time: 14.000000 seconds.
Error: 0.005488.

4 GPUs

Elapsed time: 13.000000 seconds.
Error: 0.005488.

おおむね、理想的な効率で並列化ができていることがわかる。


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

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

<comments hideform="false" />


Comments

ノート:マルチGPUプログラミングにおけるコンテキスト切り替え処理の改善について

個人用ツール