• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  G-DEPの高速演算記  >  高速演算記 第12回 「Tesla M2090 と CUDA 4.0 速報」

高速演算記 第12回 「Tesla M2090 と CUDA 4.0 速報」

以前よりGeForceシリーズでは512個のFermiコアを搭載した GeForce GTX 580 が販売されていますが, この度, Teslaシリーズでも512個のFermiコアが搭載された Tesla M2090 が発売されることになりました.

今 回, NVIDIAの協力を頂き, Tesla M2090の販売開始にあたって, テクニカルサンプルを貸与頂き, 評価する機会を得ましたので, 今回の高速演算記では, このサンプルを使用したベンチマーク結果をご紹介いたします. また, NVIDIAからGPUプログラム開発環境の最新バージョンであるCUDA4.0が発表されましたので, こちらも合わせてレポート致します.

Tesla 2090のベンチマーク

Tesla 2090の特徴ですが, 6GBのメモリと512コアを搭載しています. CUDAのSDKに添付されているバンド幅テストを実行してみたところ以下のような結果がでました. 参考のため, C2050の測定した結果も掲載されています.

  Tesla C2050 (ecc on) Tesla 2090相当 (ecc on)
Host-to-device [MB/s] 5059 5107
Device-to-host [MB/s] 4116 4007
Device-to-device [MB/s] 84259 105398
Memory clock [Mhz] 1500 1848
GPU Clock [GHz] 1.15 1.3

Host- to-deviceはCPUからGPUへ, Device-to-hostはGPUからCPUへ, Device-to-deviceはデバイス内でのデータ転送時の結果です. 上記を見る限り、新たにデバイス間データ転送が高速になっているのが伺えます. また、nbodyで倍精度演算能力をみますと314 GFLOP/sと, Tesla C2050の242 GFLOP/s よりもさらに演算能力が高まっていることが伺えます.

なお, 今回ベンチマークを使用したGPUのスペック・演算性能等はTesla M2090と同等との報告を受けております. Tesla M2090は, ブレードサーバーへの搭載を目的としたGPUで, 冷却ファンが GPUカード上には搭載されず, ヒートシンクのみが搭載されています. 冷却のための空気の循環は, カードが装着される筐体の方で提供することを前提にしている設計です. そのため, Tesla M2090 を単体で購入してPCへ装着することはできません. また残念ながら, カード単体で一般に販売することは予定されていないようです.

したがって, Tesla M2090 はサーバーを利用するハイエンドユース向けで, ワークステーション向けには, しばらく Tesla C2050 か Tesla C2070を利用することになりそうです.
 

(G-DEP副理事長 藤澤智光, CUDA エンジニア 田原哲雄)

CUDA 4.0の新機能

Fermiカードの機能を利用しやすくする新しいCUDAツールキット4.0がダウンロード可能となっています. 現在, 正式リリース前のRC2がNVIDIAのダウンロードサイトから入手することが可能となっていますのでご存知の方もいるかもしれません.

CUDA 4.0では何が新しくなったのでしょうか. 一つは1スレッドで複数GPUを扱えるようになったこと, 64ビット環境ではUnified Virtual Address空間が利用可能となったことなどが挙げられます. また, プロファイラも機能強化され, 最適化の箇所について, ヒントを得ることが可能となりました.

複数GPUカードの利用


CUDA4.0 から, cudaSetDevice()をプログラム中なんどでも呼び出すことが可能となりました. cudaSetDevice()によって選択されたデバイスが以降の関数呼び出しの対象になります. cudaSetDevice()を頻繁に呼び出してもオーバーヘッドはそれほどないようです.
1スレッドが複数のデバイスを使用できるようになり, 今まで存在していたいくつかの関数の名称が動作と合わなくなったため, 名称が変更された関数がいくつか存在します. ただし, いままでに存在していた関数はなくなったわけではないので, 以前のコードをそのままコンパイルすることも可能となっています.

変更された関数の例として

  • cudaThreadSynchronize -> cudaDeviceSynchronize
  • cudaThreadExit -> cudaDeviceExit

などです.

GPU間のデータ転送


複 数のGPUカードが同一PCIバス上に存在する場合, 片方がPCIマスタとなって他方のカードをアクセスすることができるようになりました. これを実現するための関数がいくつか追加されています. メモリ転送を実現するのが cudaMemcpyPeer() です. この関数は転送先アドレス, 転送先デバイス番号, 転送元アドレス, 転送元デバイス番号を引数として必要とし, データをGPUカード間で転送します.

すべてのカードの組み合わせで動作するわけではなく, Fermiアーキテクチャのカード,同一アーキテクチャ同士が通信可能となっているようです. 例としてはTeslaC2050が2枚あった場合, カード間でデータ転送が可能となります.  特定のカードから特定のカードをアクセス可能にするためにcudaDeviceEnablePeerAccess() という関数が用意されています. デフォルトではすべてOFFになっているので, 相互にカード間転送を行いたい場合は2度関数を呼び出す必要があります.

GPU間転送に関しては, サポートされていない組み合わせであっても, 関数内部で自動的に一旦ホストを経由してデータ転送を行うようになっているようです.

簡略化されたコードを添付しますと

...
cudaSetDevice(dev0);
cudaDeviceEnablePeerAccess(dev1, 0);
cudaSetDevice(dev1);
cudaDeviceEnablePeerAccess(dev0, 0);
...
cudaMemcpyPeer(mem1, dev1, mem0, dev0, memSz);

と いった流れになります. 試しに2枚のTeslaC2050上でそれぞれ64MBずつメモリを確保し、GPU間のコピーを試した結果を下に示しました. 下の表でホスト経由と記載のあるのはGPU0とGPU1間でメモリ転送する際に一旦ホストを経由した場合の結果です. GPU間とあるのは, 両GPUカード間で直接転送が実現できるようにcudaDeviceEnablePeerAccess()を設定した場合となります. また、最後の欄ではcudaMemcpyPeer()を使用せずに以前のSDKで提供されていたcudaMemcpy()を利用, ホスト側にPinnedメモリを用意してGPU0から転送, 次にGPU1へ転送した場合の結果を示しました.

  時間 [ms] 転送レート [MB/s]
ホスト経由 156 4124
GPU間 140 4583
cudaMemcpy 226 2831

この結果から, 特別にバッファ処理用コードを記述することなくGPU間転送を実現することが可能となりました.

Unified Virtual Address の利用

64 bit のlinux環境で Fermiカードを使用している場合, あるいは 64 bit Windows 環境で FermiカードをTCCモードで使用している場合は Unified Virtual Address を使用することが可能となっています. このモードが有効になっているカードでは, GPUとcudaHostAlloc()で確保したCPUのメモリアドレスが統一空間内に配置されます.  以前は各GPUで個別にアドレス割り当てが行われているため、カードで確保したメモリアドレスが他カードで確保したメモリアドレスと同一の番地に置かれる 可能性がありましたが, 今回の機能によりGPUメモリ, CPUメモリを区別して扱う必要が少なくなります.  


UVA モードが使用可能であるかは cudaGetDeviceProperties() 関数から取得できます. また, UVA空間を指すポインタは cudaPointerGetAttributes() を使用することによってポインタがどのメモリを指すかを特定することが可能となりました.

pinnedメモリの指定

cudaMemHostAlloc() 関数によってpinnedメモリを確保することは以前から可能でしたが, 4.0ではcudaHostRegister()関数によって通常のmalloc()によって確保したメモリを後からpinnedメモリに指定することが 可能となり、プログラム中のメモリ確保関数をcudaのものと置き換えなくてもすむようになります.

プロファイラの強化

プ ロファイラを実行した後,  それぞれのカーネル名をダブルクリックすれば, 最適化に貢献するヒントを表示させることが可能となりました. また, GUIが更新され, 新たに設けられたカーネル解析ウィンドーから, 各カーネルが計算あるいはメモリ転送能力によって束縛されているかどうかを各項目の数値から確認することが可能になります.  カーネルごとにL1/L2キャッシュのヒット率, 共有メモリのバンクコンフリクト情報など, 確認することが可能です.

ま た、プログラムからGPUプログラムのプロファイル情報やイベントを取得することができるようなライブラリが用意されました. これを利用するとCUDAカーネルの起動と終了時に呼ばれるコールバックを登録することが可能になります. イベントは, visual profilerと同様のイベント情報を取得することが可能なようです. これを利用することによってアプリケーションに特化したプロファイラツールを作成することが可能となります.

この他にも、Tesla, Quadroカードの状態を取得するためのライブラリが提供されています。nvidia-smiツールと同様にeccエラー情報, GPU使用率, ファン情報などがアプリケーションから取得できるようになります.

最後にVisual Studio 2010からでもCUDAが利用可能となっているようで、GPUプログラミングの利便性が向上しているのではないでしょうか.

(CUDA エンジニア 田原哲雄)