Upgrade to Pro — share decks privately, control downloads, hide ads and more …

GPUの実行モデルを理解してうまく使いたい

ainozaki
September 14, 2024
500

 GPUの実行モデルを理解してうまく使いたい

ainozaki

September 14, 2024
Tweet

Transcript

  1. n各SMには… u実⾏ユニット(Lane) n INT32/FP32/FP64 : 64/64/32個 n Tensor Core n

    Load/Store Unit n Special Function Unit u実⾏可能な命令を選ぶスケジューラ uL1キャッシュ: 192KB n ⼀部をプログラムから明⽰的に使える n Shared Memory uレジスタファイル など 8 ハードウェアの話:SM内 この図はSMの1/4だけ
  2. n4. サイクル毎にReadyなWarpを1つ選ぶ! uWarp schedulerが選ぶ n スケジューリングポリシー:LRR(loose round robin), GTO(greedy-then-oldest) uその後命令を実⾏ユニットに割り当て、実⾏する

    n 同じくWarp単位で同じ動き(SIMD)をする、条件分岐も時分割で実⾏ 14 実⾏モデル SM0 実⾏ユニット Warp スケジューラ I-Buffer w1 Ld.global %r4, [%r10] Ready w2 Add %r3, %r1, %r2 Wait
  3. nしかし計算機を埋めるのは難しい n例えばNVIDIA A100で演算器を休みなく使うためには… uDRAMバンド幅1.5TByte/s uINT32ピーク性能:19.5TOPS より u1Byteのデータを取ってきて、19.5/1.5=13程度演算をする必要がある n キャッシュなどは考慮していない n

    ベクトル加算なぞではメモリアクセスがボトルネック u⾏列積のような計算がたくさん発⽣するカーネルばかりではない 19 GPUの計算資源を有効に使いたいが…
  4. nHfusion uソースコードレベルで2つのカーネルを⽔平にfusion n ⼀般的なfusionでは中間結果のメモリ読み書きを減らすべく垂直にfusion 21 Concurrent Kernel Executionの例 Ao Li,

    Bojian Zheng, Gennady Pekhimenko, and Fan Long. 2022. Automatic horizontal fusion for GPU kernels. In Proceedings of the 20th IEEE/ACM International Symposium on Code Generation and Optimization (CGO '22) __global__ void kernel_fused(){ if (threadIdx.x < THREAD1){ // memory-bound }else { // compute-bound } } __global__ void kernel1(){ // memory-bound } __global__ void kenrel2(){ // compute-bound }
  5. n Matthew D. Sinclair, “CS 758: Advanced Topics in Computer

    Architecture”, https://pages.cs.wisc.edu/~sinclair/courses/cs758/fall2019/handouts/lecture/cs758-fall19- gpu_uarch2.pdf n M. Lee et al., "Improving GPGPU resource utilization through alternative thread block scheduling," 2014 IEEE 20th International Symposium on High Performance Computer Architecture (HPCA), Orlando, FL, USA, 2014, n Xiaodan Serina Tan, Pavel Golikov, Nandita Vijaykumar, and Gennady Pekhimenko. 2023. GPUPool: A Holistic Approach to Fine-Grained GPU Sharing in the Cloud. In Proceedings of the International Conference on Parallel Architectures and Compilation Techniques (PACT ʻ22) n H. Zhao et al., "Tacker: Tensor-CUDA Core Kernel Fusion for Improving the GPU Utilization while Ensuring QoS," 2022 IEEE International Symposium on High-Performance Computer Architecture (HPCA), 23 参考⽂献