CUDA 4.1の新機能(3): プロセス間Peerアクセス

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

これまでのCUDAでは、デバイスメモリのメモリ空間は、プロセス毎に独立したものであり、あるプロセスが他のプロセスのデバイスメモリにアクセスすることはできなかった。とは言え、CUDA 4.0が出てくるまでは、あえてそういう処理を行う局面は少なかったのだが、CUDA 4.0の新機能、GPUDirectによるPeerアクセスが登場すると、マルチGPUの各々を、別々のプロセスに割り当てる運用において、Peerアクセスが利用不可能、という、デメリットが目立つようになった。

この問題に対して、CUDA 4.1では、プロセス間でデバイスメモリ空間を共有するAPIが提供されることとなった。

APIの紹介

cudaIpcGetMemHandle(cudaIpcMemHandle_t *, void *)
第2引数で指定されたデバイスメモリにアクセスするためのハンドルを、第1引数で指定されたポインタに格納する。このハンドルを他のプロセスと共有することで、プロセス間で間接的にデバイスメモリ空間を共有できる。副作用として、ハンドルを通して共有されたデバイスメモリは、Peerアクセスが可能なデバイスからアクセスできるようになる。
cudaIpcOpenMemHandle(void **, cudaIpcMemHandle_t)
第2引数で指定されたハンドルを、プロセスのデバイスメモリ空間に変換して、第1引数で指定されたポインタに格納する。こうして得たデバイスメモリへのポインタは、プロセスで自由にアクセスできる。
cudaIpcCloseMemHandle(void *)
ハンドルから得たデバイスメモリを、プロセスのメモリ空間から解放する。

サンプルプログラム

#include <cstdio>
#include <cstring>

#include <cuda_runtime.h>
#include <mpi.h>

int
main(int argc, char **argv)
{
  MPI_Init(&argc, &argv);
  int rank;
  MPI_Comm_rank(MPI_COMM_WORLD, &rank);
  cudaSetDevice(rank);
  typedef struct {
    size_t l;
    cudaIpcMemHandle_t h;
  } mpidata_t;
  union {
    mpidata_t d;
    char c[sizeof(mpidata_t)];
  } mpidata;
  if (rank == 0) {
    const char *HELLO = "Hello, World!";
    mpidata.d.l = strlen(HELLO) + 1;
    void *d;
    cudaMalloc(&d, mpidata.d.l);
    cudaMemcpy(d, HELLO, mpidata.d.l, cudaMemcpyHostToDevice);
    cudaIpcGetMemHandle(&mpidata.d.h, d);
    MPI_Send(mpidata.c, sizeof(mpidata_t), MPI_CHAR, 1, 0, MPI_COMM_WORLD);
    MPI_Barrier(MPI_COMM_WORLD);
    cudaFree(d);
  } else /*if (rank == 1)*/ {
    MPI_Status status;
    MPI_Recv(mpidata.c, sizeof(mpidata_t), MPI_CHAR, 0, 0, MPI_COMM_WORLD,
             &status);
    void *d;
    cudaIpcOpenMemHandle(&d, mpidata.d.h);
    char *hello = new char[mpidata.d.l];
    cudaMemcpy(hello, d, mpidata.d.l, cudaMemcpyDeviceToHost);
    cudaIpcCloseMemHandle(d);
    MPI_Barrier(MPI_COMM_WORLD);
    printf("%s\n", hello);
    delete[] hello;
  }
  MPI_Finalize();
}

2プロセスを想定したMPIプログラムで、まずはRankの番号のGPUをSetDeviceする。Rank 0は、デバイスメモリを確保して、"Hello, World!"という文字列をコピーし、cudaIpcGetMemHandleでハンドルを作ったのち、このハンドルをRank 1へMPI_Sendする。Rank 1は、Rank 0からハンドルをMPI_Recvし、cudaIpcGetMemHandleでRank 0で確保したデバイスメモリを得、ホストにコピーして、コピーされた文字列をprintfで表示する。最終的に、デバイスメモリを介して、プロセスを跨いで

Hello, World!

という文字列が表示される。

まとめ

たとえばMPIのような、マルチプロセスを利用したCUDAプログラムにおいて、より効率的なデータ転送が行えるようになった。


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

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

<comments hideform="false" />


Comments

ノート:CUDA 4.1の新機能(3): プロセス間Peerアクセス

個人用ツール