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

これまでの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プログラムにおいて、より効率的なデータ転送が行えるようになった。