NVIDIA，「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う

　NVIDIAの開発するGPUコンピューティング向け開発環境「CUDA」。その最新版「CUDA 5」は，2012年8月からリリース候補版がダウンロード可能になっていたが，いよいよ正式版がリリースされることになった。
　それに合わせ，アジア太平洋地域の報道関係者を対象とした電話会議が開催され，CUDA 5の概要が語られたので，今回はその内容を簡単にまとめてみたい。

GK110コアの「Tesla K20」に向けた準備が整うCUDA 5

　電話会議で説明を担当したのは，米NVIDIAでGPUコンピューティングソフトウェア部門でシニアプロダクトマネージャを務めるWill Ramey氏だ。同氏が示したCUDA 5のトピックは下に示したスライドのとおりで，これらは第2世代Keplerアーキテクチャを採用した「GK110」コア版Tesla「Tesla K20」の機能として，すでに明らかになっていたものになる（関連記事）。

CUDA 5のトピック。Telsa K20の新要素として紹介されてきたものばかりだ

　そのなかでも最も大きなポイントとなるのは，おそらく「Dynamic Parallelism」（ダイナミック並列処理）だ。
　従来，GPUスレッドはCPU側からしか生成できなかった。それに対してDynamic Parallelismでは，GPUスレッド内から別のGPUスレッドを生成可能になったのだ。これにより，動的にGPU内部でスレッドを生成できるようになるので，より柔軟で効率的なGPUプログラミングが可能になる。

　Ramey氏いわく「利用は非常に簡単」とのこと。CPUからGPUスレッドを生成するコードとあまり変わらないコードでGPU内部からGPUスレッドを生成できるようだ。

青い枠がCPUからGPUスレッドを生成するコード，緑色の枠がCUDA 5からサポートされる「Dynamic Parallelismを用いたGPUスレッド内からGPUスレッドを生成するコード」の例である。「X <<< .... >>> (data);」などとなっている行が，GPUスレッドを生成する命令だが，見てのとおり，CPUとGPUの両方で同じ命令が利用できる

　なお，Dynamic Parallelismを利用できるのは，「GK110ベースのTesla K20と，将来の（開発コードネーム）MaxwellベースのGPUになる」（Ramey氏）。「GK104」ベースの「Tesla K10」や，コンシューマー向けGPUのGeForce GTX 680などでは利用できず，Tesla K20の登場を待つ必要がある機能なのだ。

　2つめの「GPU Callable Libraries」（GPUから呼び出せるライブラリ）だが，これを説明するには従来のCUDAの話が必要だろう。CUDA 4以前は「GPUカーネル内の関数は1つのソースファイルに収めなければならない」という縛りがあり，GPUコードのライブラリ化には限界があった。しかし，CUDA 5からは，関数を異なるファイルに分割し，コンパイル時にリンクできるようになったほか，GPUコードをライブラリ化できるようになった。これがGPU Callable Librariesというわけだ。

これまで1ファイルに収める必要があったGPUカーネルのコードを機能別に分割し，ライブラリ化できるようになる。スライド中の「cuBLAS」というのは，汎用的な演算ライブラリ「BLAS」のCUDA版で，従来はGPUコード側のライブラリ化に制限があったために，こういったライブラリの利用にも制限がつきまとったがCUDA5から利用しやすくなるわけだ

　ソースを分割したりライブラリ化したりするのは，プログラマーからするとごく当たり前の行為なので，「縛りがあった」という歴史を知らないと，何がよくなったのか理解しづらいかもしれない。CUDA 5からは，GPUを利用するライブラリ単独の配布や販売が可能になるため，「CUDAのエコシステムがより大きく広がる」（Ramey氏）という期待が持てるわけだ。
　なお，このライブラリ化に対応できるのは，FermiベースとKeplerベースのGPUすべてだと，Ramey氏は述べている。

　さて，3つめの「GPUDirect」は，GPUから別のGPUのメモリへネットワーク越しにアクセスする機能で，一般に「Remote Direct Memory Access」（RDMA）と呼ばれるものだ。Ramey氏は「CPUを介することなくネットワーク越しにGPUからGPUへメモリアクセスが行えるため，分散コンピューティングの性能が大きく向上する」とし，「OpenMPI」など，既存の分散コンピューティングライブラリからの利用も可能だと説明していた。
　なお，GPUDirectが利用できるのはTesla K10＆K20と，KeplerベースのQuadroになるそうだ。要するにGeForceではサポートされないわけだが，そもそも一般PCユーザーが使う理由もなさそうな機能なので，問題はないだろう。

CUDA 5ではネットワークを経由してあるGPUから別のGPUのローカルメモリへダイレクトにアクセスできるようになる

　最後に「New Nsight Eclipse Edition」は，CUDA向け統合開発ツール「Nsight」（エヌサイト）に，OSを問わずに利用されている統合開発プラットフォーム「Eclipse」向けの「Nsight Eclipse Edition」が登場するという話である。Eclipseというのは，エディタやデバッガなどが一体化された統合開発環境で，プラグインによって開発言語やエディタなどを拡張できるのが特徴だが，要は，Eclipse用のNsightプラグインが登場するというわけなのである。
　これまで，CUDAの統合環境はWindows版「Visual Stduio」のプラグインのみサポートされていた。だが，CUDA 5から新たに追加される「Nsight Eclipse Edition」によって，Windows以外のLinuxやMacでもCUDAアプリケーションの効率的な開発が可能になるのだ。
　なお，Nsight Eclipse Editionに含まれるEclipseプラグインは，専用のコードエディタ「CUDA-Aware Editor」と，CPUとGPUの双方のデバッグが可能な「Nsight Debugger」，性能解析を行う「Nsight Profile」となっている。

Nsight Eclipse Editionの登場で，CUDA向けの統合開発環境がLinuxやMac OS上で利用できるようになる。「CUDAの開発効率が大幅にアップする」とRamey氏は述べているが，LinuxやMacintoshのユーザーには確かに嬉しい新機能といえそうだ

　というわけで，CUDA 5のリリースによって，次世代GPUコアたるGK110の機能がフルサポートされ，第2世代KeplerアーキテクチャベースのTeslaを迎える準備が整った，といったところだろうか。

　ちなみにNVIDIAは，CUDA 5のリリースに合わせて開発者向けサイト「CUDA Zone」をリニューアルし，各種ドキュメントやサポートファイルを提供する「CUDA Resoruce Center」という開発者向けの情報ページをCUDA Zone内で立ち上げる予定予定とのことだ。CUDA Resoruce Centerには，（これまでに蓄積されたものも含めて）CUDA開発に関連するサンプルコードやデータが1500以上用意されるという。
　気になった人は，CUDA Zoneをチェックしてみるのがよさそうだ。

CUDA Zone