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

2011年11月22日現在、リリース候補版が出ている、CUDA 4.1。今回は、その新機能のうちの一つ、アライメントの制限の無くなったcudaHostRegisterを使ってみた。

目次 1 cudaHostRegisterとは 2 実際に使ってみる 2.1 FORTRAN側プログラム 2.2 CUDA側プログラム 2.3 実験環境 2.4 出力結果 3 まとめ

cudaHostRegisterとは

cudaHostRegisterとは、CUDA 4.0で追加されたAPI。すでに確保済みのホストメモリを、GPUからアクセスしやすいPinnedメモリに変更するものである。Pinnedメモリを新規に確保することは、昔からcudaHostAllocで可能であり、このPinnedメモリを使うことで、ホスト-GPU間の高速データ転送や、ゼロコピーのデータ同期などが実現できていた。

cudaHostRegisterを使うことで、ホスト側のプログラムと親和性の高い形でCUDAのプログラムを書くことができ、より効率的なGPGPUコンピューティングが実現できる……、はずであった。ところが、CUDA 4.0のcudaHostRegisterは、OSのページサイズ（x86では、基本的に4096バイト）のアライメントのアドレスのメモリにしか適用できず、実際に使おうとすると、ほとんどのケースで適用不可能な、そんなAPIであった。

実際に使ってみる

このcudaHostRegisterについて、CUDA 4.1ではアライメントの制限が無くなり、いよいよ実用的なAPIになったと考えることができる。そこで、今回はこのAPIを使うテストを行ってみた。

FORTRAN側プログラム

      PROGRAM HOSTREGISTER
      IMPLICIT NONE
      INTEGER I,N
      REAL A,B,C,SUM
      PARAMETER (N=1000000)
      COMMON A(N),B(N)
      DO 10 I=1,N
         A(I)=REAL(I)
 10   CONTINUE
      CALL CUDASQ(N,B,A)
      SUM=0.0
      DO 20 I=1,N
         SUM=SUM+B(I)
 20   CONTINUE
      WRITE(*,*)SUM
      STOP
      END PROGRAM HOSTREGISTER

CUDA側プログラム

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

extern "C" void
cudasq_(int *n_, float *dst, float *src)
{
  size_t n = *n_;
  cudaSetDeviceFlags(cudaDeviceMapHost);
  cudaHostRegister(src, n * sizeof(float),
                   cudaHostRegisterMapped);
  cudaHostRegister(dst, n * sizeof(float),
                   cudaHostRegisterMapped);
  float *d_src, *d_dst;
  cudaHostGetDevicePointer((void **)&d_src, src, 0);
  cudaHostGetDevicePointer((void **)&d_dst, dst, 0);
  dim3 block(512);
  dim3 grid((n + block.x - 1) / block.x);
  cudasq_kernel<<<grid, block>>>(n, d_dst, d_src);
  cudaDeviceSynchronize();
  cudaHostUnregister(dst);
  cudaHostUnregister(src);
}

1 - 1000000の値で初期化された配列を受け取って、その二乗の値を出力配列に書き出すプログラムである。最後に、この出力配列の和をとって、標準出力に出力する。注目すべきは、入力の配列にcudaHostRegisterを適用して、ゼロコピーのデータ同期を行っている点である。

実験環境

CPU:	Core i7 2600K
GPU:	GeForce GTX 580
OS:	Ubuntu 11.04
CUDA:	CUDA 4.1 RC1

出力結果

  3.33381993E+17

FORTRAN側で確保した配列を使って、ゼロコピー通信が実現できているのが分かる。

まとめ

CUDA 4.0で使いにくかったcudaHostRegisterが、CUDA 4.1では非常に使いやすくなっていて、CUDAの開発効率が、より向上されることが期待できる。