CUDA 4.1の新機能(1): Unaligned cudaHostRegister

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

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

目次

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の開発効率が、より向上されることが期待できる。


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

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

<comments hideform="false" />


Comments

ノート:CUDA 4.1の新機能(1): Unaligned cudaHostRegister

個人用ツール