読書活動 PMPP Ch. 3, 4
ひきつづき Programming Massively Parallel Processors – 4th Edition. だらだらインターネットしてたら出遅れたが読んでいきます。
CHAPTER 3 Multidimensional grids and data.
- 三次元まである構造に対しどういうハードウェアの支援があるのかはわからなかった・・・・。ただデータの数 (ex. 画像のピクセル数) ぶんだけスレッドがあるという抽象は面白いというか、スレッドプール脳の自分はメンタルモデルを補正しなければいけないと思いました。
CHAPTER 4 Compute architecture and scheduling
- Stream Multiprocessor (SM) は複数の CUDA Core を持つ.
- SM には速いメモリ(キャッシュみたいの?) がついている
- Thread block = SM! こういうのが知りたかった!
- SM には数十しかコアがないのに 1024 per block のスレッドを許すのだから、SM は 1024 スレッドぶんのローカルなステート(stack, pc, etc.) はコアの外のどこかに保持するのだね。この shared memory に入っているのだろうか。
- __syncthreads()!
- WARP はスケジューリングの単位で、つまりほんとに並列に動く。現状は 32 thread = 1 WARP らしい。つまり最低でも各 SM は 32 cuda core. 64 core ある SM は 2 WARP を並列に動かさせるということか。それは誰得なのだろうな・・・。
- とおもいきや SM の Core が WARP より小さい (ex. 16 core) "processing block" に分割され、これが SIMD の単位らしい。なんでやねん！WARP = SIMD ならわかるが、これが違うなら WARM の存在意義とは...
- というとこれは歴史的経緯で、昔は 1 SM = 1 WARP (= 32 Core) みたいな構成だったらしい。メンタルモデル的には WARP = SIMD と思っておけばよいもよう。
- そして Volta (2017) 以降では WARP 内で control divergence が起きても同一 WARP 内のそれぞれのコードパスが同時に動けるらしい。SIMD とはなんだったのか。わけわかんねーそりゃ消費電力も増えるわけだわ。
- なお、厳密には SM は 1 つ "以上" つまりハイエンドだと 2 つとかの thread block を持てる模様。これも SM = WARP というマッピングを破壊しているが、アーキテクチャの進化の帰結なのだろうね。
- Latency Tolerance. これすごいよな。OS がスケジューリングに関与せずソフトウェアからスレッドのスケジューリングが見えないとかプロプリエタリにも程があるしプロファイリングとかも困難でクソだなとおもっていたが、DRAM アクセスの瞬間にスレッドを切り替えるんだからハードウェア側でやらざるを得ない。仕方ないですね・・・と納得した。というか CPU もメモリアクセスでストールするとかはプロファイラでは見えない（なんらかのカウンタはあると思うけど）ので、可視性は大差ないとも言える。
- そして latency を上手に隠したいなら thread block のサイズはなるべくデカくしておいた方が良いということなんですね・・・とおもいきや SM は複数の thread blocks を同時にホストできるので、かならずしも block size を 1024 にしなくてもよい。とはい同時にホストできる block 数にも上限があるので、多いに越したことはないそうな。更にスレッドが使うレジスタ数でも utilization が制限される！大変だなー・・・。あるカーネルが必要とするレジスタ数とかクエリできるのだろうか。デバイスの性能がわかっても、コードの性質がわからないとな。
以上。Chapter 4 は面白かったね。
疑問があるとすれば、モバイルの GPU の "コア" たとえば M1 が 10 GPU Core というときそのコアは CUDA の語彙で何に相当するのだろうか。SM? CUDA Core? 常識的に考えたら SM だが、ならその "Core / SM" の中にはいくつの CUDA Core 相当があるの？まあそれは CUDA の教科書には出てこないよなあ。