Occupancyを可視化する

消えたCUDA関連の旧ブログ記事を復元するひとり Advent Calendar 2024の記事です。

何の話か

CUDAカーネルのプロファイリングをしていると必ず遭遇する言葉「Occupancy」がそもそも何なのか、実際にスレッドが走っている様子を可視化するとどうなるのか、 みたいな話をします。

Occupancyとは？

端的に言えば、あるカーネル関数がStreaming Multiprocessor (SM)の演算器をどれほど使いきれるかという値です。
この値はカーネル関数が使うレジスタ数やSharedメモリ量によって変わってきます。

---

CUDAでは1つのThread blockは1つのSMで実行されますが、1つのSMは同時に1つのThread blockしか実行しないわけではありません。
つまり、SMに搭載されているレジスタやSharedメモリの量的に、同時に2つのThread blockを実行可能な場合は2つ実行します（正確には「する可能性があります」）。
しかし、この同時に走らせられるThread blockの数の制約は、レジスタやSharedメモリなどのプログラマブルな資源によるものだけではありません。
そもそも各SMが同時に走らせられる上限のThread数が決まっているのです（注意：Thread block数ではない）。
最近のGPUですと1,536 threadsで、もしThread sizeを256でカーネルを立ち上げた場合は1,536/256=6 thread blockが1 SMで同時に実行可能であることになります。
この1,536という数字は以下のコードで調べられます。

// nvcc threads_per_sm.cu -arch=sm_80 -lcuda
#include <stdio.h>
#include <cuda.h>

int main() {
  cuInit(0);
  CUdevice device;
  cuDeviceGet(&device, 0);

  int num_threads_per_sm;
  cuDeviceGetAttribute(
      &num_threads_per_sm,
      CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR,
      device);

  printf("Num threads per SM = %d\n", num_threads_per_sm);
}

では、1 SMが同時に実行できるWarp数を考えると、これは1,536/32=48 warpsとなります。

Theoretical occupancyの話

ここでもし、1 threadが使う資源が少なく、SMが同時に48 warp分に資源を割り当てられるカーネル関数があったとします。
これがTheoretical occupancy 100%の状態です。
カーネル関数が使う少し資源が増え、24 warp分の資源しか同時に割当できないカーネル関数であればTheoretical occupancy 50%となります。

Active occupancyの話

Theoretical occupancyはカーネルの立ち上げ時には決定している値です。
一方で、実際に実行してみて1 SM中で同時に48 warp中何Warp走りましたか？という率がActive occupancyです。

Occupancyを時系列で見てみる

Occupancyが何となくなにかわかったところで実際に見てみます。
見たいのは、ある時刻に1 SMで同時にいくつのThread blockが走っているかです。
あるThread blockがどのSMで実行されているかは%smidレジスタを読めばわかります。
ということで、こちらのコードでThread block IDとSMID、カーネルの実行時刻を取得します。
このコードでは1024GClock (約0.6秒くらい)無をさせるカーネル関数を立ち上げます。
立ち上げる際に1 thread blockあたりのSharedメモリのサイズを決定し、Theoretical occupancyを制御します。

#include <iostream>

constexpr unsigned block_size = 256;
constexpr unsigned grid_size = 1u << 10;
constexpr unsigned wait_clock = 1lu << 30;

__global__ void test_kernel() {
  extern __shared__ unsigned smem[];

  const unsigned long t0 = clock64();
  while(clock64() - t0 < wait_clock){}
  const unsigned long t1 = clock64();

  unsigned smid;
  asm(
      R"({mov.u32 %0, %smid;})":"=r"(smid)
      );
  if (threadIdx.x == 0) {
    printf("%u,%u,%lu,%lu\n", blockIdx.x, smid, t0, t1);
  }
}

void launch(const unsigned smem_size) {
  int num_blocks_per_sm;
  cudaOccupancyMaxActiveBlocksPerMultiprocessor(
      &num_blocks_per_sm,
      test_kernel,
      block_size,
      smem_size);
  std::printf(
    "smem_size = %u,num_blocks / SM = %d\n",
    smem_size,
    num_blocks_per_sm);

  cudaFuncSetAttribute(
    &test_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
      smem_size);
  test_kernel<<<grid_size, block_size, smem_size>>>();
  cudaDeviceSynchronize();
}

int main() {
  launch(10 * 1024);
  launch(40 * 1024);
}

見やすいよう開始時刻をシフトし、SM 0番で実行されているThread blockをグラフで表します。
横軸はカーネルを立ち上げてからの経過Clock、縦軸は実行されているThread block IDです。
Theoretical occupancy = 100%の場合、1 SMあたり6 thread blocks (=1,536/256(thread size))が実行されます。

---

Sharedメモリの使用量を少なめに設定し、Theoretical occupancy = 100%(6/6)としたのが下のグラフです。
重なっていて少し見にくいですが、確かに1つのclock区間に6つのThread blockが走っていることがわかります。

▲ SM 0番で動作するThread blockの時系列表示。

---

次はSharedメモリの使用量を増やし、Theoretical occupancy = 33.3%(2/6)としました。
確かに1つのclock区間に2つのThread blockが走っていることがわかります。

▲ SM 0番で動作するThread blockの時系列表示。

---

この1 clock区間に走っているThread blockの数（の最大値(=1,536/thread_size)に対する割合）こそがOccupancyです。

Theoretical occupancyを2倍にできれば理論計算速度も2倍になる？

という疑問は当然出てくると思いますが、残念ながら2倍にはなりません。
1 SMが複数のThread blockを実行できるのは、命令の遅延を隠すための機能です。
同時に複数のThread blockが実行されると言っても、演算器の個数はThread数と比較して少ないため、同じタイミングで全Thread blockが同じ命令を実行されるわけではありません。
Threadはタイミングが微妙にずれながら動いています。
（タイミングが揃っていることが保証されているのはsyncされているWarpのみです。）
このズレが演算器を衝突なく利用するための鍵で、Occupancyを高めるということはこのズレたThreadをたくさん生やせるようにするということです。
この思想はCUDAでの命令の遅延隠蔽の根幹となっています。
では、逆にOccupancyが低いと隠蔽を隠せないかと言うと必ずしもそうではなく、Thread数はThread sizeを大きくすることでも増やせるのでこちらで隠蔽していくという手もあります。
遅延隠蔽は他にも使用する回路の異なる命令を同時に動かす命令重ね合わせ（Instruction-level overlap）や、AmpereからはAsynchronous global-to-shared data copyなどがあります。

おまけ：SM割当

では、あるthread blockがどのSMで実行されるのかには規則があるのでしょうか？
例えばblockIdx.x % num_smのようにblockIdx.xから簡単に計算可能なものなのでしょうか？
ということで、横軸に%smid、縦軸にblockIdx.xをとったものが下のグラフたちです。

• 1 SMあたり6 thread blocksが実行される場合（Theoretical occupancy=100%の場合）
  
• 1 SMあたり2 thread blocksが実行される場合（Theoretical occupancy=33.3%の場合）
  </ul>

  RTX 3080を用いており、SMは64基搭載されています。
  このためSMIDは64で折り返されます。
  1周目の割当時のSMIDはblockIdx.x % num_smで決まり、2周め以降は適当に、という感じに見えますね。

  おわり

  あまり日本語でOccupancyの説明が書かれたWebページが見当たらなかったので書いてみましたが、読んでいただき何か得るものがあったのであれば幸いです。
  特に可視化やsmid割当に関しては英語などでも説明されているものを見たことがないので、少しはオリジナリティがあるかなと思います。
  SMIDとSIMDはアナグラムになっていてパット見見分けがつかないので、どこかでtypoしていたらごめんなさい。