CUDAのPinnedホストメモリ

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

CUDAを使っていると、

  • GPUとホストとの間の転送が遅く、このオーバーヘッドが高速化分を相殺してしまう。
  • GPUのメモリが少なすぎて、メモリを大量に扱う処理を移植できない。

などの問題点がたびたび発生する。これらの問題点は、Pinnedホストメモリで或る程度解決ができる。

以下の文章でプログラムの実行を行う際の環境は、以下の通りである。

CPU Xeon W3565 (3.2 GHz, 4 cores)
GPU GeForce GTX 580
OS CentOS 5.8
Compiler CUDA 4.2 (-O4 -arch=sm_20)

目次

Pinnedホストメモリとは

CUDAの利用に適した、ページアウトしないホストメモリであり、cudaHostAllocによって新規確保、cudaHostRegisterによって既存のホストメモリをPinned化できる。cudaHostAllocによって確保されたメモリはcudaFreeHostによって解放でき、cudaHostRegisterによってPinned化されたメモリはcudaHostUnregisterによって非Pinnedホストメモリに戻せる。Pinned化したホストメモリは、GPUとの通信が高速に行えるほか、後述するMappedメモリとしても使うことができる。

Pinnedホストメモリの転送速度は、CUDA SDKに含まれる、bandwidthTestサンプルプログラムを実行すると、簡単に確認できる。

通常の非Pinnedホストメモリ使用:

$ NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest --memory=pageable
[bandwidthTest] starting...

NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest Starting...

Running on...

 Device 0: GeForce GTX 580
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     4851.4

 Device to Host Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     4196.7

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     165759.1

[bandwidthTest] test results...
PASSED

> exiting in 3 seconds: 3...2...1...done!

Pinnedホストメモリ使用:

$ NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest --memory=pinned
[bandwidthTest] starting...

NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest Starting...

Running on...

 Device 0: GeForce GTX 580
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Pinned memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     5812.5

 Device to Host Bandwidth, 1 Device(s), Pinned memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     6108.9

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     165766.0

[bandwidthTest] test results...
PASSED

> exiting in 3 seconds: 3...2...1...done!

4 GB/s台だった転送速度が、6 GB/s程度に高速化されているのが分かる。特に、GPUからホストに通信する際の性能の伸びが良い。

なお、Pinnedホストメモリは原則APIを呼び出したContextに於いてのみ利用可能であるが、cudaHostAllocPortable(cudaHostAlloc使用時)またはcudaHostRegisterPortable(cudaHostRegister使用時)をAPIの引数のフラグとして立てることによって、ホスト上の全てのContextで共用できるようになる。

Mappedメモリ

Pinnedホストメモリを割り当てる際に、cudaHostAllocMapped(cudaHostAlloc使用時)またはcudaHostRegisterMapped(cudaHostRegister使用時)をAPIの引数の引数のフラグとして立てると、このPinnedホストメモリはMappedメモリとなる。MappedメモリはGPU上から直接アクセスできるホストメモリであり、cudaHostGetDevicePointerにMappedメモリのポインタを渡すことによって、対応するGPUのメモリ空間上のアドレスを得ることができる。MappedメモリのサイズはGPUのメモリ量に制約されず、GPUのメモリ量を超える大量のメモリを扱うことができる。

以下はベンチマーク用に、Mappedメモリを使うように手を加えたbandwidthTestである。1バイトcharへのメモリアクセスでは性能が出ずに、性能を出すには2バイトshort以上のデータ幅が必要なようであり、そのように書いている。Warpサイズ32と、Xeon W3565のキャッシュライン64バイトとが、何か関係が有るのかもしれないが、調査中である。

--- NVIDIA_GPU_Computing_SDK/C/src/bandwidthTest/bandwidthTest.cu       2012-07-17 10:15:55.000000000 +0900
+++ NVIDIA_GPU_Computing_SDK/C/src/bandwidthTest-Mapped/bandwidthTest.cu        2012-08-13 17:25:04.000000000 +0900
@@ -860,6 +860,17 @@
     return bandwidthInMBs;
 }

+static __global__ void
+myCudaMemcpy(unsigned char *dst, unsigned char *src, size_t n)
+{
+  size_t i = ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) *
+    blockDim.x + threadIdx.x;
+  if (i >= n) {
+    return;
+  }
+  ((short *)dst)[i] = ((short *)src)[i];
+}
+
 ///////////////////////////////////////////////////////////////////////////////
 //! test the bandwidth of a device to device memcopy of a specific size
 ///////////////////////////////////////////////////////////////////////////////
@@ -889,10 +900,14 @@
     }

     //allocate device memory
-    unsigned char *d_idata;
-    checkCudaErrors( cudaMalloc( (void**) &d_idata, memSize));
-    unsigned char *d_odata;
-    checkCudaErrors( cudaMalloc( (void**) &d_odata, memSize));
+    unsigned char* h_d_idata;
+    checkCudaErrors( cudaHostAlloc( (void**) &h_d_idata, memSize, cudaHostAllocMapped | cudaHostAllocWriteCombined ) );
+    unsigned char* d_idata;
+    checkCudaErrors( cudaHostGetDevicePointer( (void**) &d_idata, h_d_idata, 0 ) );
+    unsigned char* h_d_odata;
+    checkCudaErrors( cudaHostAlloc( (void**) &h_d_odata, memSize, cudaHostAllocMapped | cudaHostAllocWriteCombined ) );
+    unsigned char* d_odata;
+    checkCudaErrors( cudaHostGetDevicePointer( (void**) &d_odata, h_d_idata, 0 ) );

     //initialize memory
     checkCudaErrors( cudaMemcpy( d_idata, h_idata, memSize,
@@ -901,10 +916,24 @@
     //run the memcopy
     sdkStartTimer( &timer );
     checkCudaErrors( cudaEventRecord( start, 0 ) );
+    int currentDevice;
+    checkCudaErrors( cudaGetDevice( &currentDevice ) );
+    struct cudaDeviceProp prop;
+    checkCudaErrors( cudaGetDeviceProperties( &prop, currentDevice ) );
+    struct cudaFuncAttributes attr;
+    checkCudaErrors( cudaFuncGetAttributes( &attr, myCudaMemcpy ) );
+    size_t blockx = attr.maxThreadsPerBlock;
+    dim3 grid( (memSize / sizeof(short) + blockx - 1) / blockx %
+               prop.maxGridSize[0],
+               (memSize / sizeof(short) + blockx * prop.maxGridSize[0] - 1) /
+               (blockx * prop.maxGridSize[0]) % prop.maxGridSize[1],
+               (memSize / sizeof(short) +
+                blockx * prop.maxGridSize[0] * prop.maxGridSize[1] - 1) /
+               (blockx * prop.maxGridSize[0] * prop.maxGridSize[1]) );
     for( unsigned int i = 0; i < MEMCOPY_ITERATIONS; i++ )
     {
-        checkCudaErrors( cudaMemcpy( d_odata, d_idata, memSize,
-                                cudaMemcpyDeviceToDevice) );
+      myCudaMemcpy<<<grid, blockx>>>
+        ( d_odata, d_idata, memSize / sizeof(short) );
     }
     checkCudaErrors( cudaEventRecord( stop, 0 ) );

@@ -930,8 +959,8 @@
     free(h_idata);
     checkCudaErrors(cudaEventDestroy(stop));
     checkCudaErrors(cudaEventDestroy(start));
-    checkCudaErrors(cudaFree(d_idata));
-    checkCudaErrors(cudaFree(d_odata));
+    checkCudaErrors(cudaFreeHost(h_d_idata));
+    checkCudaErrors(cudaFreeHost(h_d_odata));

     return bandwidthInMBs;
 }
$ NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest-Mapped --dtod
[bandwidthTest-Mapped] starting...

NVIDIA_GPU_Computing_SDK/C/bin/linux/release/bandwidthTest-Mapped Starting...

Running on...

 Device 0: GeForce GTX 580
 Quick Mode

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     7827.3

[bandwidthTest-Mapped] test results...
PASSED

> exiting in 3 seconds: 3...2...1...done!

PCI Expressの帯域にボトルネックが有るようなベンチマーク結果を示し、通常のデバイスメモリのアクセスに比べては、20分の1程度の速度しか出ていないことが分かる。Mappedメモリへのアクセスは、極力少なくなるようなプログラムを組むのが望ましい。

注意点

Pinnedホストメモリはスワップアウトしない。つまり、ホストの搭載メモリ量ギリギリの量を確保すると、残りの物理メモリでOSやアプリの動作に必要なメモリをやりくりしなければならないために、これらのOSやアプリに必要なメモリのスワップアウトが頻出するようになってしまう。搭載メモリに対して、余裕の有る容量の確保を心掛けたい。

なお、厳密な発生条件は不明であるが、Windows Vista/7 x64に於いて、Pinnedホストメモリが700 MB程度しか確保できない不具合が有るようである。

関連項目

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


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

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

<comments hideform="false" />


Comments

ヴィトン 長財布 メンズ 二つ折り said ...

Tsukihime gives you a standard display to it Arcueid is just are you looking for one specific quite Apostle to deal with given its name Roa. simple fact that Roa doesn thoroughly show themself because of awhile, combined with Shiki truly Arcueid cure tiny his minions, things are all attached with you properly. there a winding sentiment with the demonstration generally considering it slowly and gradually teases out contact lenses and back to you scenario to a great deal, nevertheless,having said that tuning into it from the demonstration sort such as this will show it an extremely tighter piece of work.

--ヴィトン 長財布 メンズ 二つ折り 2014年10月24日 (金) 12:08 (UTC)

個人用ツール