GPUの実⾏モデルを理解してうまく使いたい Ai Nozaki (@ainno) 2024/09/14 情報科学若⼿の会＠軽井沢

nAi Nozaki uX: @ainno321 uGitHub: @ainozaki u…⾃⼰紹介タイムがあったのでスキップ！ 2 ⾃⼰紹介 ←変えたい ↑変えたい

nみなさんGPUを使っていますか？ n研究で暗号（準同型暗号）のGPUアクセラレーションをしているが… 3 はじめに ハイエンドなGPUを全 然使いきれてない 既存研究より遅すぎる

n普段使っているGPUがどう動いているのか、ハードウェア/プログラ ミングモデルの⾯から少し理解する nGPUの利⽤率を⾼めるために⾏われていることを少し紹介する u 今回の発表ではNVIDIA GPUを扱い、以降のスライドでGPUと略します！ 4 発表の概要

nGPUの実⾏モデルを理解する編 nGPUの利⽤率を⾼めるために編 5 ⽬次

nCUDAのハードウェアとソフトウェアの概念がややこしい！ u1. ハードウェアの整理（NVIDIA A100を例に) u2. CUDAプログラムの整理 u3. プログラムがどうハードウェアで実⾏されるか、の流れで話す 6 GPUの実⾏モデルを理解するにあたって CUDAなんて 書かないよ..

n⼤量の演算器がSMという単位で束ねられた構成 uSM = Streaming Multiprocessor, A100だと108個 7 ハードウェアの話 ：全体像 https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/ SM L2 Cache Global Memory

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だけ

nどう並列な処理を記述するか？→Thread, Block, Gridという概念 u処理の最⼩単位がThread n Laneで処理する単位 uThreadの集合がBlock uBlockの集合がGrid u物理コア数よりも圧倒的に多くのThreadを起動可能 9 プログラムの話 Thread Blockサイズ（1024, 1, 1） Gridサイズ（4, 2, 1）

10 プログラムの話：ベクター加算の例 GPUで実⾏するコード（カーネル） ←128Threadを持つBlockを256個⽴ち上げる

n1. CPU （ホスト）はGPU（デバイス）上でカーネルを⽴ち上げる！ 11 実⾏モデル CPU GPU cudaMalloc() cudaMemcpy() add<<>>()

GPU n2. BlockをどのSM上で実⾏するか割り当てる！ u1Blockは1SMに割り当てられ、1SMは複数Blockを処理する n 最⼤Block数/SMは、HWリソースとプログラムの使⽤量から決まる n レジスタ、Shared Memory、Threadの制約 n Blockを詰めるほど利⽤率は上がるが、リソースの競合が発⽣しうる 12 実⾏モデル SM0 SM107 ・・・ SM1 B B B B B B Block スケジューラ

n3. 割り当てられたThreadが実⾏する命令を読み実⾏の準備！ u命令が格納されているメモリから命令を読んでくる（Fetch） u命令を解釈する（Decode） u命令がREADYかどうか（依存が解決されているか）確認する uこの時、32ThreadずつWarpという単位でまとめて処理される 13 実⾏モデル SM0 命令 キャッシュ w1 Ld.global %r4, [%r10] Ready w2 Add %r3, %r1, %r2 Wait 命令 デコード レジスタの管理テブル（Score Board）を⽤ いて確認した結果 I-Buffer

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

n4ʼ. 実⾏可能なWarpが1つもない場合はストール 15 実⾏モデル SM0 実⾏ユニット Warp スケジューラ w1 Ld.global %r4, [%r10] Wait w2 Add %r3, %r1, %r2 Wait

nまとめると… nCUDAプログラムは、Block単位でSM（演算ユニットの塊）に割り当 てられ、Warp = 32Threadの単位でまとめて命令実⾏される nGPUは実⾏可能なWarpをサイクル毎に次々と切り替えて実⾏すること で命令実⾏のレイテンシを隠蔽している！ uGlobal Memoryアクセス命令のレイテンシ：500cycleくらい u算術命令のレイテンシ：数⼗cycleくらい 16 実⾏モデル

nGPUの実⾏モデルを理解する編 nGPUの利⽤率を⾼めるために編 17 ⽬次

nベクター加算のプログラムをプロファイル u計算ユニットの利⽤率は16%程度 n メモリアクセスを待つ間、演算器が空いている n …もったいない！ 18 GPUの計算資源を有効に使いたいが… Warpスケジューラの ストール発⽣理由 Memory stall ←SM利⽤率 16%

nしかし計算機を埋めるのは難しい n例えばNVIDIA A100で演算器を休みなく使うためには… uDRAMバンド幅1.5TByte/s uINT32ピーク性能：19.5TOPS より u1Byteのデータを取ってきて、19.5/1.5=13程度演算をする必要がある n キャッシュなどは考慮していない n ベクトル加算なぞではメモリアクセスがボトルネック u⾏列積のような計算がたくさん発⽣するカーネルばかりではない 19 GPUの計算資源を有効に使いたいが…

nメモリボトルネックなカーネルでは、演算器が暇になってしまい資源 の利⽤率が低い>< ↓ 様々なアプローチがあるが…アプローチの1つとして nConcurrent Kernel Execution u名の通り同時に複数のカーネルを実⾏する u特に要求するリソースの異なる2つのカーネルを、1つのSMで同時に実⾏ n リソースが競合しない命令を詰め込むので、Warpスケジューラが実⾏可能な命令 が増えるはず 20 1つのアプローチ

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 }

n仕組みを知るのは楽しい n少し速くなると嬉しい nかなり使われているけれどまだやることがありそうな感じが楽しい ありがとうございました！ 22 まとめ

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 参考⽂献