GTC2017 - Tesla V100 と CUDA9.0

今年のGTCはジェンスン・ファンCEOの基調講演が3日目に設定されており、そこでVolta世代の新製品 Tesla V100 と、それを搭載するDGX-1などサーバ製品とGPU クラウドが発表された。
4日間の日程のうち、情報の解禁が3日目の正午なので、その後は慌ただしくVoltaのアーキテクチャや、関連するCUDA9新機能のセッションに参加。
聞くべきセッションが被っちゃったりもしたので、できれば基調講演は初日の方が日程に余裕はあったかなぁ…。

Tesla V100について

5120CUDAコア、16GB(HBM2; 900GB/s)のメモリを搭載するまさに化け物プロセッサ。

デバイス全体で80個のSMを搭載する。SMあたりの構成は64CUDA CoreなのでPascalと同等だが、目新しい基軸としてTensor Coreという新しい計算ユニットが追加されている（SMあたり8個 = 全体で640個）。また、L1 Cacheが128KB x 80 = 10MB、L2 Cacheが6MB搭載されている。

これは4x4行列同士のFMAD演算(←と言っていいの？)を実行する専用のH/Wユニットで、従来であれば、16個のCUDAコアが協調して掛け算/足し算を行うべき箇所を命令一発で実行するというもの。

深層学習のアルゴリズムに関して詳しくないんだけども、ハマる状況であれば綺麗に高速化できるんであろうし、実際、TensorCore機能を使った機械学習ワークロードの実行速度が、P100比で9.3倍に向上したという事はアナウンスされている。

また、6本のリンクを全てフル帯域で使った場合の理論帯域だけども、NVLinkで300GB/sのスループットがあり、NVLink接続のリモートGPUに対するcache coherencyやatomic operationが可能だとのこと。

この辺の発表を聞いていての印象は、15年-20年ほど前からx86 CPUが辿った道のりに似てるなぁ～と。

Intel CPUの場合、マルチメディア処理命令としてMMX/AVXという拡張命令が加わったが、それを自分のプログラムから直接使うという人はあまり多くなく、実態としてはglibcなどのライブラリが内部的にこれらの機能を使うようになったために、それを通じて、ハマるケース限定でハードウェア専用命令を使うようになってきている。
TensorCoreの場合も、cuDNNやcuBLASなどのライブラリ群が内部的にこの機能を使うとのことなので、ユーザプログラムが行列計算や深層学習にこれらのライブラリ群を使っていれば、知らず知らずのうちに恩恵を受けることになる。というか、ラップしてくれないとちょっと使える状況が特殊すぎて辛い。

もう一つ。NVLinkの帯域が向上し、しかもUnified Memoryでアドレス空間の問題が解決し、cache coherencyやatomic operationに対応するとなれば、他のGPUデバイス上のメモリをローカルメモリに近い感覚で使うことができるようになるという事だろう。
そうすると、かつてLinuxがNUMA環境への対応を行ったように、メモリの物理的な距離を意識しなければならない状況が生まれつつあるように思える。例えば Tesla V100 は 16GB(HBM2) のメモリを積んでいるが、アプリケーションが20GBのデータをロードして2台のGPUでぶん回したい場合、データ量は明らかに1台のGPUに収まらないので2台に分散ロードする必要がある。
この時、物理GPUページの割り当てが『前半10GBはGPU0、後半10GBはGPU1』みたいなざっくりで上手く動作するのだろうか？もしかしたら、メモリアクセスの局所性が期待できるタイプのワークロードなのか？　それとも、全体を舐め回すのでインタリーブ型の物理データ配置を行うべきなのか、それほど簡単な問いではないように思える。

SMの構成は図の通りであるが、L1 Cacheの扱いが少し変わっている。
PascalではL1 Cacheが24KBだったのが、128KBあるL1/Shared Memory共通の領域のうち、共有メモリが使用していない領域をL1 Cacheとして利用できるわけである。
Maxwell以前にL1/Shared Memoryの大きさを16KB/48KB、32KB/32KB、48KB/16KBの中から選んでいたのに通ずるが、もう少しフレキシブルに"早いメモリ"を使ってくれているようである。
ただ、GPUのL1キャッシュの場合(少なくともPascalまでは)、キャッシュの下に転がっているGlobal Memoryの内容がGPU kernel実行中に少なくとも変更されない/Read-Onlyであるという制約があったが、これはおそらくVoltaでも変わらないものと思われる（何も言及無かったし）

また、Pascal以前のGPUと大きく変わってくる点として、プログラムカウンタ(PC)がwarp単位ではなくthread単位になるという違いがある。
従来は、if (...) { ... } else { ... } ブロックを実行するときに、warp中の一部スレッドだけが条件が真、他のスレッドが偽である場合には両方のブロックを逐次的に実行することになるため、他方のスレッド群が遊んでしまい、それ故にGPUは条件分岐に弱いと言われてきたが、各スレッドがPCを持つという事はこうした弱点に手を打ってきたという事でもある。
（とは言え、TrueブロックとFalseブロックの分量があまりにアンバランスだとアカンでしょうが）

5月13日追記：
FBでNVIDIA-JPの成瀬さんから教えていただきました。
『プログラムカウンタ(PC)がwarp単位ではなくthread単位になる』というのは必ずしも間違ってはいないのですが、CPUスレッドのように各CUDAコアが全く独立な命令を実行できるという事ではないようです。

より正確に述べると、PC(ここでは次に実行する命令のアドレスとしましょう)は、Voltaではスレッド毎に持ちます。ただし、同じwarp内のスレッドは、PCが同じであるスレッドだけが、同じクロックで実行されます。warp内で、別PCの命令が実行されることは、これまで通り、ありません。

つまり、if (...) {} else {} ブロックで分岐が発生する場合、従来、trueブロックを実行中のCUDAコアがある時はfalseブロックに分岐したCUDAコアは寝ているしかなかった。Voltaにおいては、例えばDRAMアクセスなどでメモリストールが発生しておりtrueブロックを実行中のCUDAコアが寝ている状況では、その間を利用してfalseブロックに分岐したCUDAコアも先に実行を進めることができる、という事のようです。

他にニッチなところとしては、SMが80個にも増えてアプリケーションによっては全てのSMを使いきれないという事を想定しているのか、MPS(Multi Process Service)を使用した時に複数のアプリケーションからのGPU kernel起動要求を調停して、GPU資源を切り分けるという使い方ができますよーとの事。ただ、MPSはcuStreamAddCallback()のようにコールバック関数を呼び出す事ができなかったハズなので、非同期処理が中心のPG-Stromとは今一つ相性がなぁ…。

CUDA 9について

Voltaアーキテクチャ向けのTensorCoreを使うための機能や、cuBLASなどライブラリのTensorCore対応が行われるらしい。

ただ、最も革命的と言ってよい変化は、今までスレッド間の同期は __syncthreads() を使ったSM内の同期のみであったのが、SMを越えて複数のグリッド間で／複数のGPUデバイスで並行動作するカーネル間で同期を取ることが可能になったというもの。

GPU kernelを起動するときは、必要なスレッド数がN個としてN=(ブロックサイズ)×(グリッドサイズ)
という関係になるようスレッド数を与えて起動する。ブロックサイズの上限は 1024 であり、GPU kernelの実行時にはブロックサイズに依らず、特定のブロックは特定のSMで実行される。
（グリッドサイズ > デバイスのSM数な場合は、ブロックを順番に実行する）

で、大雑把に言えば、従来はH/Wに由来するブロックを単位とする同期の他には、GPU kernel全体の実行完了を待つという同期方法しか取れなかったが、CUDA 9で導入される co-operative group という概念を使えば、H/W的な構造と論理的な処理の同期の単位を切り分けることができる。

これがなぜ嬉しいかと言うと、例えばBitonic-Sortingを実行するときに、要素数が1個のSMで処理するには辛い数（1万個とか）になってくると、ステップ事に一度GPU kernelを完了し、そこで同期した上で改めて次のステップを実行するという処理が必要になる。
これは結構しんどい処理で、PG-StromでGPUでのORDER BY処理の実装を断念した要因の一つでもある。
が、SMを越えた単位で処理を同期できるとなれば、ステップ事に同期ポイントを設定してやればよいだけである。

例えば、以下のような感じで処理を記述する。

__global__ void
particleSim(Particle *p, int N)
{
  grid_group g = this_grid();

  // 1. Jobs prior to sync
  for (i = g.thread_rank(); i < N; i += g.size())
    integrate(p[i]);

  g.sync() // 2. Sync whole grid

  // 3. Jobs post sync
  for (i = g.thread_rank(); i < N; i += g.size())
    collide(p[i], p, N);
}

grid_groupというのはco-operative groupを表現するためのクラスで、this_grid()はこのGPU kernelを起動した時の全てのスレッドを含むオブジェクト。
つまり、g.sync()は複数のSMに跨った全てのスレッドで同期を取ることを意味している。

他にも、従来同様に現在のブロックに含まれるスレッドを表現する thread_group や、逆に複数のデバイスで同時に実行中の各GPU kernelに含まれるスレッドを表現する multi_grid_group が定義されている。
なお、SM間/デバイス間での同期機構を使う場合には、GPU kernelを起動する時に用いるAPIが従来のものとは少し異なっており、Runtime APIでは以下のAPIを用いてGPU kernelをキックしなければならない。

grid_groupを使う場合
- cudaLaunchCooperativeKernel()
multi_grid_groupを使う場合
- cudaLaunchCooperativeKernelMultiDevice()

で、更に面白いことに、このco-operative groupというのは更に細かい単位にパーティション化する事ができる*1。

こんな調子で、thread_group gに含まれるスレッドを tiled_partition() を介して4つのパーティションに分割する事ができるらしい。
なので、真ん中の tile.sync() ではパーティション毎の同期という事になる。

  thread_group g = this_thread_label()
  thread_gruup tile = tiled_partition(g, 4);
        :
  tile.sync();
        :

現状、CUDA 9では等分割のパーティショニングのみだが、将来的には labeled partition という形にコンセプトが拡張されるとの事。
"labeled" というのは、以下のように partition() 関数に与えた整数値によってそのスレッドが属するパーティションを決めるというもの。

  int label = foo() % 4;
  thread_group g = partition(this_thread_block(), label);

何が嬉しいかというと、GROUP BY/集約演算をGPUで実装するときに、GROUP BYで指定したカラムの値が同一のスレッド間でのみ同期を取り、縮約演算を行うというロジックが可能になること。これは、現状atomic演算を用いて実装せざるを得ないことを考えると非常に大きなアドバンテージになる。

CUDA 9.0の登場が待ち遠しいところである。

余談

SM間で同期を取るとき、GPU kernelのブロック数がデバイスのSM数よりも多い場合、ハードウェア的にはSMを明け渡さないといけないはずなので、そのタイミングで共有メモリやレジスタをグローバルメモリに退避する（= そんなに軽い処理じゃないのでは？）と思い、セッション後に発表者を捕まえて質問してみた。
曰く、Pascal世代で(H/W的には)対応したプリエンプションの仕組みを使って実装されており、レジスタや共有メモリの使用量がH/Wリソースを超える場合には確かにそうなる。なので、GPU kernelを起動するときのパラメータチューニングが必要で、SMあたり {（ブロック数）/（デバイスのSM数）} のブロックを保持できる程度のレジスタ/共有メモリの消費に抑えておく必要がある。
一応、その辺は Occupancy を計算するAPIを使って効率的なブロックサイズ/グリッドサイズを求めてほしいとの事。