Tesla

代表的なソート・アルゴリズム「クイックソート」を Tesla K20 によって高速化

 
 

NVIDIA は、この 5 月に開催された GPU Technology Conference にて Kepler アーキテクチャーを採用した NVIDIA Tesla K20 GPU を発表しましたが、このプロセッサが、HPC 業界において史上最高の性能を持つプロセッサとなることを約束します。その理由の一つが、「ダイナミック並列処理」と呼ばれる、広範囲において作業のスピードアップを可能にする機能です。

ダイナミック並列処理は、GPU が CUDA カーネルの実行時に、その実行中のカーネル内部から GPU 自身の新たなタスクを生成するという方法で、CPU からの独立性が高い形でGPUを自律的に動作させることを可能にします。シンプルなコンセプトですが、とてもパワフルな手法で、特に従来 GPU には難しいと言われていた分割統治法などのようなアルゴリズムで GPU プログラミングを容易にし、大きな効果を発揮します。

コンピュータサイエンスの学生が最初に学ぶ題材でもある「クイックソート」を例にして、その能力を見てみましょう。「クイックソート」にダイナミック並列処理を適用すると、コードの行数を半減させた上に、パフォーマンスを倍増することができます。

内部動作について

従来の Fermi アーキテクチャを採用した GPU の場合、実行フローの流れは、ホスト CPU から GPU コアへの一方向に固定されています(下図左側)。

(左) ダイナミック並列処理なし、(右) ダイナミック並列処理あり

ダイナミック並列処理では、CPU を介することなく、GPU が自身のタスクを新しく生成できます。つまり、次になにをするべきかをランタイムで決めるという動的処理が可能で、従来より複雑なアルゴリズムが実現できますし (上図右側)、同時に、CPU に他の作業ができる余裕を生みだします。

このようなダイナミックな処理がおこなえるように、Tesla K20 GPU には、グリッド管理ユニット(GMU) という新しいハードウェアが搭載されました。カーネルの起動、一時停止、再開や、複数ソースの依存性追跡など、動的実行に伴う複雑な処理は、この GMU がハードウェア的に高速処理します。また、GPU 上で走るシステム・ソフトウェアのレイヤーが GMU とやりとりする形で、CUDA のランタイム API(Application Programming Interface) が CUDA カーネルプログラム内から使えるようになります。

クイックソート・アルゴリズムの概要

では、クイックソートのアルゴリズムに話を進めましょう。この例を見れば、ダイナミック並列処理が秘めたパワーを理解していただけるはずです。

まず、その原理をおさらいしましょう。目的は、数字の列をソートすること。そのために、まず、「ピボット」と呼ぶ数値をひとつ選び、ピボット未満の数とピボット以上の数に全体を二分割します。

下に示すダイアグラムでは、単純に、数列の最初 (左端) の数値をピボットにしています。

最初の数列を二分割したら、ふたつの数列、それぞれに対してクイックソートを適用し、4 つの数列に分けます。各数列に含まれる値がひとつになるまで、この操作をくり返し、そのときの数列をひとつにまとめれば、ソートが完了です。問題を順次、小さなものへと分割し、再帰的に解決して行くため、「分割統治法」アルゴリズムと呼ばれます。

クイックソートが簡単に実現-コード行数を半減

では、クイックソートの CUDA コードについて、ダイナミック並列処理がある場合とない場合を比べてみましょう。

ダイナミック並列処理を使ったクイックソート

ダイナミック並列処理を使わないクイックソート

これはプログラミングの専門化でなくても気づくことですが、ダイナミック並列処理を使ったクイックソートは、ダイナミック並列処理なしに比べてコード行数が半分になっています。また、全体の流れも理解しやすくなっています。なぜ、そうなるのでしょうか。

クイックソートの場合、各段階でソートする情報は、その 1 段階前に決定されます。ダイナミック並列処理を使用しない場合は、各段階の起動は、すべて CPU がコントロールしなければならないため、各段階の終了時に、次の段階にすべきことをホスト側に送りかえす必要があります。ここに挙げた例では、単純化して分かりやすくするために、この CPU-GPU 間の通信を CPU/GPU ワークスタックの操作という簡単な形で書いていますが、実際には、アトミック操作やデータ管理などが必要となるとても複雑な処理で、そのコードを書いたとすれば、クイックソートのアルゴリズムと同じくらいの長さになるはずです。

これに対し、ダイナミック並列処理を使用すれば、必要に応じて実行時に GPU が自身でタスクを起動できるため、ある段階のクイックソートが終われば、直ちに次の段階で行う2セットのクイックソートを起動することができます。CPU/GPU のスタック交換といった複雑なオーバーヘッドもなくなりますし、タスクの起動管理するホスト側コードもすべて不要になります。こうして、全体が短く、理解しやすくなりますし、次の項目で見るように、実行も速くなるのです。

ダイナミック並列処理でパフォーマンス向上

上記の 2 種類のコードについて、同じ Tesla K20 GPU を使って性能を比べてみました。結果は、グラフに示すとおりです。ダイナミック並列処理ありのクイックソートは、ダイナミック並列処理なしに比べてスピードが 2 倍になっています。

このようなスピードアップが可能な理由は、タスクの起動方法と密接に関連しています。CPU がコントロールするコードでは、ある段階の処理が終わるまで待ってから次の段階の処理を始めざるを得ません。そのため、各段階ごとに cudaDeviceSynchronize() をコールする必要があります。この操作自体もかなりの負荷になりますが、さらにこのコールによって、ある段階のソートがすべて終了しないと次の段階のソートが始められないという状況になります。つまり、各段階の処理時間は操作に一番時間がかかる部分の処理時間に等しくなるわけで、スピードが一番遅い部分に合わせて全体が進むことになります。

これに対し、ダイナミックな並列処理コードでは、必要に応じて新しいタスクを起動します。GPU と CPU の間でデータをやりとりする必要はありません。また、各段階の処理がすべて終わるのを待たなくても、次の段階を始められるようになります。このため、たくさんの作業を同時に処理できるようになりますし、管理のオーバーヘッドも大幅に小さくなります。

このように、コードが書きやすくなる-読みやすくもなる-だけでなく、実行スピードも格段に向上できるのです。

無限の可能性

ダイナミック並列処理には以下のような 3 つの大きなメリットがあり、これは GPU コンピューティングに大きな変革をもたらす機能です。

  1. GPU のプログラミングがとても簡単になる
  2. いままで GPU による高速化が難しいと考えられてきたアルゴリズムも簡単に高速化できる
  3. CPU に対する GPU の依存性が大きく下がり、CPU、GPU 共に動作効率が向上する

ダイナミック並列処理には無限の可能性があります。今後 2~3 週の間に、パワフルなユースケースをふたつ、記事にする予定です。ひとつは、GPU から並列ライブラリーを直接呼び出して複雑なアルゴリズムを実装する方法で、もうひとつは、小さなジョブをたくさん起動して GPU を最大限に利用する方法です。

もし、ダイナミック並列処理のメリットが活用できそうなコードについてアドバイスがあれば、コメント欄を通してお知らせください。現在我々は、CUDA Toolkit に収録するサンプル・コードを集めている最中なので、いいアイデアがあれば、是非お寄せください。

Kepler やその優れた機能については、@NVIDIAJapan でも情報を提供しています。