GPUDirectによるPeerアクセスについて

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

前回、CUDA 4.1で強化されたPeerアクセスの話をしたところ、そもそも従来のGPUDirectによるPeerアクセスの方法が、よくわからない、という指摘を受けた。そこで、今回は、このGPUDirectによるPeerアクセスの方法について、説明をする。

GPUDirectとは

GPUDirectとは、CUDAを使って、GPUと同じバスに繋がっている(他のGPUを含む)デバイスに対して、GPUから直接アクセスする技術である。CPUを経由しない、高速なデータ転送が実現でき、GPU間のデータ同期や、ネットワークデバイスを使った他のノードとの通信を、効率的に行うことができる。この技術を使った、GPU間の直接データ転送を、Peerアクセスと呼ぶ。

Peerアクセスを使うことで、他のGPUで割り当てたデバイスメモリに対して、カーネルやcudaMemcpy系の関数で、直接アクセスすることができるようになる。

注意が必要なのは、「同じバス」という点であり、たとえばマルチソケットのCPUを使ったPCの中には、複数本のバスを持ち、GPUと別のバスに他のデバイスが繋がっているような構成があるため、気を付けていただきたい。

NVIDIA社によると、この機能はTeslaとQuadroのプレミアム機能のようであるが、弊社のCentOS 5.7 + CUDA 4.0 + ドライバVer. 290.10の環境において、GeForce GTX 580でGPUDirectが可能であった。どういう状況で機能が有効になるのか、弊社では問題を切り分けているところである。

APIの紹介

cudaDeviceCanAccessPeer(int *, int, int)
第2引数と第3引数とで指定されたデバイス番号のGPUの間で、Peerアクセスが可能であるかを、第1引数のポインタに返す。可能であれば1を、不可能であれば0を返す。
cudaDeviceEnablePeerAccess(int, int)
cudaSetDeviceで指定された現在のGPUから、第1引数で指定されたデバイス番号のGPUへの、Peerアクセスを有効化する。第2引数は、CUDA 4.0では必ず0を指定する。
cudaDeviceDisablePeerAccess(int)
cudaSetDeviceで指定された現在のGPUから、引数で指定されたデバイス番号のGPUへの、Peerアクセスを無効化する。
cudaMemcpyPeer(void *, int, const void *, int, size_t)
第4引数で指定されたデバイス番号のGPU上の、第3引数の領域にある、第5引数バイトのデータを、第2引数で指定されたデバイス番号のGPU上の、第1引数の領域にコピーする。この関数は、明示的にPeerアクセスを有効化しなくても、成功するようである。同期APIであるが、CPUは転送の終了を待たず、転送に関わった2つのGPUと、cudaSetDeviceで指定された現在のGPUとの、3デバイスのみが転送の終了を待つ。
cudaMemcpyPeerAsync(void *, int, const void *, int, size_t, cudaStream_t)
同期APIのcudaMemcpyPeerに対応する、非同期API。第6引数でストリームを指定する。

サンプルプログラム

#include <stdio.h>
#include <stdlib.h>

#define N (1000)
#define NTHREADS (32)

__global__ void
kernel(int n, float *d_a)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    d_a[i] = i;
  }
}

int
main()
{
  /* GPU 1からGPU 0へのPeerアクセスを許すか調べる */
  int canAccessPeer;
  cudaDeviceCanAccessPeer(&canAccessPeer, 1, 0);
  if (!canAccessPeer) {
    printf("Peer Access Disabled!\n");
    return -1;
  }
  /* GPU 0でメモリを割り当てる */
  cudaSetDevice(0);
  float *d_a;
  cudaMalloc((void **)&d_a, N * sizeof(float));
  /* GPU 1からGPU 0へのPeerアクセスを有効化する */
  cudaSetDevice(1);
  cudaDeviceEnablePeerAccess(0, 0);
  /* GPU 1のKernelで、GPU 0で確保したメモリにアクセスする */
  dim3 grid((N + NTHREADS - 1) / NTHREADS);
  kernel<<<grid, NTHREADS>>>(N, d_a);
  /* GPU 1のコンテキストのままで、GPU 0のメモリをホストにコピーする */
  float *h_a = (float *)malloc(N * sizeof(float));
  cudaMemcpy(h_a, d_a, N * sizeof(float), cudaMemcpyDeviceToHost);
  /* 結果(和は499500)を得る */
  float sum = 0.0f;
  for (ptrdiff_t i = 0; i < N; ++i) {
    sum += h_a[i];
  }
  printf("%f\n", sum);
  return 0;
}

このプログラムを、GPUDirectが使えるマルチGPUの環境で実行すると、

499500.000000

という正しい結果が返り、GPUDirectによるPeerアクセスが正しく実行されていることがわかる。


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

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

<comments hideform="false" />


Comments

ノート:GPUDirectによるPeerアクセスについて

個人用ツール