GPU向けコンパイラの最適化の紹介と論文のサーベイ
この記事では，私の研究分野であるGPU向けコンパイラの最適化の紹介と論文のサーベイを行う．
以下，随時更新．
分岐発散 （Branch Divergence）
分岐発散とは
Branch Divergence, Warp Divergence, Thread Divergence, Control Flow Divergenceと色々な呼び方がある．分岐発散とは，NVIDIAに代表されるSIMD(SIMT)型の実行モデルのGPUで起きうる実行効率が低下する問題である． NVIDIAのGPUでは，実行単位であるスレッドは32個が１グループとなっていて，そのグループをwarpと呼ぶ．ある分岐命令で，warp内のスレッドが異なる分岐経路に分岐するとき，分岐発散が発生する．例えば，以下のようなコード．(まあ，普通はこんなコードは書かないが...)
if(threadIdx.x % 2 == 0){
//BB1
} else {
//BB2
}
//BB3
１行目の分岐命令によって，スレッドIDが偶数のスレッドはBB1側の命令を，スレッドIDが奇数のスレッドはBB2側の命令を実行しなければならない．このとき，SIMD(SIMT)型の実行モデルのGPUは，分岐先を片方ずつ順に実行する．先にBB1側から実行すると仮定すると，スレッドIDが偶数のスレッドがBB1側を実行し，その間スレッドIDが奇数のスレッドは何もせずに停止している状態になる．次に，スレッドIDが奇数のスレッドがBB2側を実行し，その間スレッドIDが偶数のスレッドは停止している．その後，BB3以降の命令をwarp内の全てのスレッドで実行していく．このとき，スレッドが再収束(reconvergence)するのは，分岐発散した分岐の直後支配節(immediate post-dominator, IPDOM)である．
こんなことが起きてしまうと，当然実行効率は下がってしまう．ループ内で起きてしまったら最悪である．この問題を解決するために，GPU向けコンパイラの最適化が盛んに研究されている．
Independent Thread Scheduling
Volta GV100以降の割と新しいNVIDIAのGPUでは，Independent Thread Schedulingという機能がある． これは，上述のような分岐発散している分岐で，必ず片方の分岐先が終了してからもう片方の分岐先を実行しなければならないという制約を軽くするものである． 以下，NVIDIAの公式webサイトから画像を引用して，説明する．
Independent Thread Schedulingが登場する前は，スレッド達は以下のような実行になる．
分岐条件がtrueになったスレッドがAとBを実行し，その後分岐条件がfalseになったスレッドがXとYを実行する． その後，reconvergenceして，Zを実行する，という流れだ．
しかし，Independent Thread Schedulingを適用すると，次のような実行をする．
図12を見てみると，分岐条件がtrueになったスレッドがAの実行を終えた後，分岐条件がfalseになったスレッドがXの実行をしている．
このように，分岐発散している分岐先が異なるコードの実行の割り込みが発生するようだ．
Zの実行がreconvergenceされていないが，これは保守的に考えてreconvergenceしていないようだ．
もし，Zの実行の後にreconvergenceさせたいなら，CUDA 9以降で使える__syncwarp()を使う．
__syncwarp()を使うと，次のような実行になる．
Independent Thread Schedulingがあっても，図12や図13を見て分かるように，依然として実行はSIMD(SIMT)スタイルのままである． つまり，あるクロックサイクルにおいて，warp内のアクティブスレッドは全て同じ命令を実行している． なので，後述の分岐発散に対する最適化は，依然として重要である．
Independent Thread Schedulingについて，詳細はNVIDIAの公式webサイトを参照してほしい．
分岐発散に対する最適化
Software based approaches
Code Compaction of Matching Single-Entry Multiple-Exit Regions(W.-K. Chen, B. Li and R. Gupta, 2003, SAS)
分岐発散を低減するために提案された手法ではないが，参考までに一応掲載する．tail mergingを提案していて，コードサイズを小さくすることを目的としている．同じコードを含むSEME regionを特定し，同じコードを一箇所に集めて，そこへのジャンプ命令を追加する．分岐発散を低減するために使えなくもない．Reducing Branch Divergence in GPU Programs (T. D. Han and T. S. Abdelrahman, 2011, GPGPU)
分岐発散を低減するために，iteration delayingとbranch distributionを提案している．iteration delayingは，ループ中の条件分岐で，分岐発散しているスレッドの実行タイミングを次回以降の繰返し時まで遅らせ，より多くのスレッドで実行するようにする手法．繰返しを超えて，true側はtrue側で，false側はfalse側でなるべくまとめて実行しようという発想．これにより，計算の並列性が高まる．branch distributionは，分岐発散している分岐の両分岐先に，同じ演算子の命令があったら，それらを分岐の外に括り出すことで分岐発散を低減する手法．このとき，分岐中に代入演算子を挿入することで，同じ演算子だが異なるオペランドを持つ命令を括り出す．Divergence Analysis and Optimizations (B. Coutinho, D. Sampaio, F. M. Q. Pereira and W. Meira Jr., 2011, PACT),
Divergence Analysis (D. Sampaio, R. M. de Souza, S. Collange and F. M. Q. Pereira, 2014, TOPLAS)
分岐発散の有無を静的に解析する分岐発散解析と，分岐発散を低減するための分岐融合 (Branch Fusion) という手法を提案している．分岐発散解析は，optimistic analysis (楽観的解析) なので，実際は分岐発散しない分岐を分岐発散していると見なしてしまう場合がある．逆に，分岐発散している分岐を分岐発散していないと見なすことはない．分岐融合は，選択命令を駆使して，分岐発散している分岐先の命令を括り出し，分岐発散の低減をする手法．しかし，この手法は，両分岐先に同じ順番で演算子が並んでいないと効果を期待できない．On-the-Fly Elimination of Dynamic Irregularities for GPU Computing (E. Z. Zhang, Y. Jiang, Z. Guo, K. Tian and X. Shen, 2011, ASPLOS)
分岐発散とコアレスアクセスになってないメモリアクセス（dynamic irregularitiesと呼ばれている）を解決するG-Streamlineというフレームワークを提案している．スレッドが参照するデータを並べ直したり（data reordering），スレッドが参照するメモリをスレッド間で入れ替えたりする（job swapping）ことで，メモリトランザクションを減らしたり，分岐発散を起こさないようにしている．Characterization and Transformation of Unstructured Control Flow in GPU Applications (H. Wu, G. Diamos, S. Li and S. Yalamanchili, 2011, CACHES)
Unstructured CFG(非構造化CFG)をStructured CFG(構造化CFG)に変換することで，分岐発散の低減と，任意のGPUでプログラムを実行可能にする手法．Unstructured CFGで分岐発散が発生すると，同じコードが２回実行される可能性があるが，それがなくなる．しかし，この手法はUsing Hammock Graphs to Structure Programs (F. Zhang and E. H. D’Hollander, 2004, TSE)の手法に基づいており，コードを複写して変換するので，コードサイズが指数関数的に増加する．このとき，Backward Copyはmultiple entry loopを，Forward Copyはmultiple exit loopを除去するために行う．また，AMDのGPUであるRadeonは，Unstructured CFGのプログラムをサポートしていないので，複数種類のGPUをサポートするコンパイラバックエンドであるOcelotで，NVIDIAのGPU向けに書かれたプログラムをAMDのGPU向けに変換するときに有用だとか．Profiling divergences in GPU applications(B. Coutinho et al., 2012, Concurrency and Computation: Practice and Experience)
分岐発散に対する最適化を提案しているわけではなく，RodiniaやNvidia SDKに対して詳しいプロファイルを取ったよという論文．Reducing Divergence in GPGPU Programs with Loop Merging (T. D. Han and T. S. Abdelrahman, 2013, GPGPU)
分岐発散が生じている多重ループ（ループを脱出するかしないかの分岐で分岐発散が生じている, loop-induced divergence）を，一つのループにまとめることで分岐発散を低減するloop mergingという手法を提案している．loop-induced divergenceがあると，多重ループの内側のループの実行を待つスレッドが存在して無駄になってしまうが，一つにまとめてしまうことで，それがなくなり速くなる．Taming Control Divergence in GPUs through Control Flow Linearization (J. Anantpur and Govindarajan R., 2014, CC)
分岐にガードを付けることで，Unstructured CFGをStructured CFGに変換する手法．4の手法と異なり，この手法ではコードサイズが線形に増加する．結果を見ると，確かにコードサイズの増加は抑えられてると思うが，実行速度については，速くなっているものは少なく，むしろ遅くなっているものの方が多い様子．Loop Optimization for Divergence Reduction on GPUs with SIMT Architecture (R. Novak, 2015, TPDS)
ループの中の分岐で分岐発散が生じているような場合に，スレッドがなるべく同時に実行できるようにスケジューリングすることで，分岐発散を低減しようという手法．スケジューリングの方法として，Frequency schedulingとBalanced schedulingの２つが提案されている．Efficient Control Flow Restructuring for GPUs (N. Reissmann, T. L. Falch, B. A. Bjørnseth, H. Bahmann, J.C. Meyer and M. Jahre, 2016, HPCS)
loop restructuringとbranch restructuringを行い，Unstructured CFGをStructured CFGに変換する手法．loop restructuringは，ループをtail-controlled loopに変換し，branch restructuringは，分岐がきちんと入れ子になっている状態に変換する．head-controlled loopに対しては，loop inversionを適用することで，新しく導入する分岐などによるオーバーヘッドを少なくする工夫も提案されている．この手法は，Perfect Reconstructability of Control Flow from Demand Dependence Graphs (H. Bahmann, N Reissmann, M. Jahre and J. C. Meyer, 2015, TACO)を基にしているらしい．この手法も，4の手法と異なり，コードサイズが指数関数的に増加することはない．実験では，-Xcicc=-O0と-Xptxas=-O0でコンパイルしたプログラムに対して提案手法を適用しているので，何だかなあという印象．On-GPU Thread-Data Remapping for Branch Divergence Reduction(H. Lin, C. Wang and H. Liu, 2018, TACO)
Common Subexpression Convergence: A New Code Optimization for SIMT Processors(S. Damani and V. Sarkar, 2019, LCPC)
Speculative Reconvergence for Improved SIMT Efficiency (S. Damani et al., 2020, CGO)
分岐発散が生じたとき，warp内のスレッドがreconvergenceするのは直後支配節（IPDOM）であるが，それを早めることで無駄にスレッドが停止している状態を無くし．実行効率を高めるSpeculative Reconvergenceを提案している．スレッドブロック内で同期を取るバリア命令をどこに置くかという解析に，データフロー方程式を使っている．reconvergence pointを決定するのに，ユーザガイド版とコンパイラがヒューリスティックにを決定する方法を提案している．実験結果を見るとかなりスピードアップしている．Branch Divergence Reduction Based on Code Motion(J. Fukuhara and M. Takimoto, 2020, JIP)
An Abstract Interpretation for SPMD Divergence on Reducible Control Flow Graphs(J. Rosemann, S. Moll and S. Hack, 2021, POPL)
DARM: Control-Flow Melding for SIMT Thread Divergence Reduction(C. Saumya, K. Sundararaja and M. Kulkarni, 2022, CGO)
分岐発散している分岐のtrue側とfalse側の領域を融合するcontrol-flow meldingという手法を提案している．この手法は，分岐発散している分岐のtrue側とfalse側の部分グラフをSESE regionに変換し，その領域を融合する．分岐融合やCSCよりも一般的な構造のCFGに適用できる．Scalar Replacement Considering Branch Divergence(J. Fukuhara and M. Takimoto, 2022, JIP)
Hardware based approaches
分岐発散が生じないようにするというモチベの手法が多い．とりあえず列挙．
Dynamic Warp Formation and Scheduling for Efficient GPU Control Flow (W. W. L. Fung, I. Sham, G. Yuan and T. M. Aamodt, 2007, MICRO)
Dynamic warp subdivision for integrated branch and memory divergence tolerance (J. Meng, D. Tarjan and K. Skadron, 2010, ISCA)
Thread block compaction for efficient SIMT control flow (W. W. L. Fung and T. M. Aamodt, 2011, HPCA)
SIMD Re-Convergence At Thread Frontiers (G. Diamos, B. Ashbaugh, S. Maiyuran, A. Kerr, H. Wu and S. Yalamanchili, 2011, MICRO)
Simultaneous Branch and Warp Interweaving for Sustained GPU Performance (N. Brunie, S. Collange and G. Diamos, 2012, ISCA)
The Dual-Path Execution Model for Efficient GPU Control Flow (M. Rhu and M. Erez, 2013, HPCA)
A Scalable Multi-Path Microarchitecture for Efficient GPU Control Flow (A. ElTantawy, J. W. Ma, M. O’Connor and T. M. Aamodt, 2014, HPCA)
Efficient warp execution in presence of divergence with collaborative context collection(F. Khorasani, R. Gupta and L. N. Bhuyan, 2015, MICRO)
Dynamic SIMD re-convergence with paired-path comparison(Y. Huang et al, 2016, ISCAS)
Iteration Interleaving--Based SIMD Lane Partition(Y. Wang et al, 2016, TACO)
Control Divergence Optimization through Partial Warp Regrouping in GPGPUs(Y. Yang, S. Zhang and L. Shen, 2018, CSAI)
A Lightweight Method for Handling Control Divergence in GPGPUs(Y. Yang, S. Zhang and L. Shen, 2019, HPC Asia)
その他
- Branch and Data Herding: Reducing Control and Memory Divergence for Error-Tolerant GPU Applications (J. Sartori and R. Kumar, 2013, IEEE Transactions on Multimedia)
分岐発散するwarpで，過半数のスレッドが分岐する経路をwarp中の全てのスレッドが実行するようにすることで，分岐発散を低減する手法．当然，プログラムの意味を変更するので，異なる結果が得られるが，画像処理計算ではそれでもいいらしい．ベンチマークプログラムに適用した結果，結構速くなっている．ソフトウェア方面とハードウェア方面のどちらも提案している．
サーベイ論文
- A Survey of GPGPU Parallel Processing Architecture Performance Optimization(S. Jia et al., 2021, ICIS)
カーネル融合 （Kernel Fusion）
Kernel Fusionとは
CPU側から呼び出してGPU側で実行する関数のことをカーネル（kernel）と呼ぶ． また，複数のカーネルを融合して，１つのカーネルにすることをカーネル融合（kernel fusion）と言う． カーネル融合には，主に次の２つの方法がある．
- 垂直融合（vertical fusion）
- 水平融合（horizontal fusion）
Automatic Horizontal Fusion for GPU Kernelsの図１が分かりやすいので，引用して説明する．
垂直融合（vertical fusion）
上の図１の真ん中にあるのが，垂直融合のパターンを表している． 融合結果は，文字通り２つのカーネルを縦に連ねているような感じになる． 縦に連ねるので，K1とK2のカーネルは同じ次元（スレッドブロックとスレッドの数）を持ち，融合後のカーネルもK1やK2と同じ次元を持つ． 図１では，それぞれのスレッドはK1の分の命令を実行した後にK2の分の命令を実行する． 一般的なカーネル融合といったらこのパターンな気がする． 垂直融合における特徴は以下の通りである．
グローバルメモリへの読み書きが減る．
K1とK2の間にグローバルバリアが必要ないときしか融合できない．1
水平融合（horizontal fusion）
図１の右側にあるのが，水平融合のパターンである．
K1とK2を二つ並べて実行しているような感じになる．
したがって，水平融合では，K1の結果をK2で使うということはできない．
水平融合には，どのレベルで融合しているかでInner Thread BlockとInter Thread Blockの２つの種類に分けられる．
水平融合については，Kernel Fusion : an Effective Method for Better Power Efficiency on Multithreaded GPUの図２が分かりやすいので，引用して説明する．
Inner Thread Block
これは，一つのスレッドブロック内でカーネルを水平融合するパターンである． K1とK2の<スレッドブロックの次元数，スレッドの次元数>を<512, 512>とすると，融合後カーネルの次元数は<512, 1024>となる． 1ブロック内では，1024スレッドのうち，512スレッドがK1を担当し，残り512スレッドがK2を担当する． 図２Bにあるように，イメージでは，以下のような感じ．
if(threadid < 512){
... //instructions of K1
} else {
... //instructions of K2
}
Inner Thread Blockでは，スレッドレベルの並列性が向上し，warpスケジューラなどが命令のレイテンシを隠蔽しやすくする効果がある． しかし，1ブロック内のスレッド数が増えるので，1スレッドが使えるレジスタや共有メモリが減ってしまう．
さらに，K1とK2の両方のカーネル内に，__syncthreads()があると，単純に融合はできない．
なぜなら，__syncthreads()はスレッドブロック内のスレッドの同期を取るので，上のコードのようにif文のthenとelseの両方に__syncthreads()が存在すると，一向に同期を取れないという状態に陥ってしまう．
これを防ぐためには，Automatic Horizontal Fusion for GPU Kernelsが行っているように，PTXのbar.syncを__syncthreads()の代わりに使って，then側だけ，else側だけでそれぞれ同期を取ってあげる必要がある．
Inter Thread Block
これは，スレッドブロックレベルで水平融合するパターンである．
K1とK2の次元を<512, 512>とすると，融合後カーネルの次元数は<1024, 512>となる．
スレッドブロックたちは，1024ブロックのうち，512ブロックがK1を担当し，残りの512ブロックがK2を担当する．
上の図２Cのパターンである．
Inner Thread Blockと違って，スレッドブロックレベルで融合しているので，融合元のカーネルに__syncthreads()があっても問題はない．
しかし，CUDA9からはスレッドブロック間の同期も取れるようになったので，融合元のカーネルにそれがあると，融合はできない．
おそらく，Inner Thread Blockの場合と同じように，then側とelse側のそれぞれでスレッドブロック間の同期を取ってあげれば，実現できそう（CUDAにそれを実現する機能があるかは分からない）．
CUDAでは，一つのスレッドブロックの実行は，一つのSMコアが担当するので，Inter Thread Blockで融合して，SMコア同士でレイテンシなどを隠蔽しあうことで，融合前よりも効率が良くなったりするのだろう．
カーネル融合に関する論文
Kernel Fusion : an Effective Method for Better Power Efficiency on Multithreaded GPU(G. Wang, Y. Lin and W. Yi, 2010, International Conference on Green Computing and Communications & International Conference on Cyber, Physical and Social Computing)
消費電力のことも考えてカーネル融合する手法を提案している．なお，消費電力の測定はGPGPU-Simというシミュレータを使っている．スレッドブロックレベルで水平融合するInter Thread Blockという方法も試している．Optimizing Data Warehousing Applications for GPUs Using Kernel Fusion/Fission(H. Wu et al., 2012, International Parallel and Distributed Processing Symposium Workshops & PhD Forum)
Scalable Kernel Fusion for Memory-Bound GPU Applications(M. Wahib and N. Maruyama, 2014, SC)
Optimizing CUDA code by kernel fusion: application on BLAS(J. Filipovic, M. Madzin, J. Fousek and L. Matyska, 2015, J Supercomput)
mapやreduce，またはそれらの組み合わせといったBLASで使うタイプの関数について特化した垂直カーネル融合をしている．垂直融合はmapなどの関数と相性が良く，速くなっているものが多い．Automatic Kernel Fusion for Image Processing DSLs(B. Qiao, O. Reiche, F. Hannig and J. Teich, 2018, SCOPES)
From Loop Fusion to Kernel Fusion: A Domain-Specific Approach to Locality Optimizationの先行研究．From Loop Fusion to Kernel Fusion: A Domain-Specific Approach to Locality Optimization(B. Qiao, O. Reiche, F. Hannig and J. Teich, 2019, CGO)
ループ融合で使われるグラフの最小カットの手法を，融合できるカーネルを探すために応用している．最小カットするのは重み付き有向グラフだが，カーネル間データ依存が辺の方向，辺の重みはその辺が繋ぐカーネルを融合したときの利益（benefit）を表している．最小カットすることで，融合後に全体で重みが最大化するので，その融合は一番利益が高いということになる．画像処理系のアプリケーションに対して実験しており，なかなかの結果が出ている．Automated OpenCL GPU kernel fusion for Stan Math(T. Ciglaric, R. Cesnovar and E. Strumbelj, 2020, IWOCL)
ベイズ推定のためのオープンソースソフトウェアStanで使われるStan Mathライブラリに着目してカーネル融合している．そのライブラリには，統計に関する数学的な処理が多く含まれているらしい．CUDAではなくOpenCL上で行っている．Automatic Horizontal Fusion for GPU Kernels(A. Li, B. Zheng, G. Pekhimenko and F. Long, 2022, CGO)
Inner Thread Blockで水平融合を行うsource-to-sourceなコンパイラをClang上に構築している．実験は，ディープラーニングで使うカーネルと暗号化で使うカーネルに絞られている．実験結果では，速くなるものは速くなるが遅くなるものは遅くなるといった感じで，差が非常に激しい．
その他のGPU関連の論文
Dimensionally redundant instruction elimination
Microarchitectural Mechanisms to Exploit Value Structure in SIMT Architectures(J. Kim et. al., 2013, ISCA)
Exploiting Uniform Vector Instructions for GPGPU Performance, Energy Efficiency, and Opportunistic Reliability Enhancement(P. Xiang et al., 2013, ICS)
Decoupled Affine Computation for SIMT GPUs(K. Wang and C. Lin, 2017, ISCA)
同じスレッドブロック内で(base, stride)の組で表せるような計算があると，warp間で冗長な計算を行ってしまう．そのような冗長性を除去するDecoupled Affine Computation(DAC)を提案している．DACでは，１つのwarpが代表して(base, stride)の組を計算し，結果をキューに保存し，他のwarpは必要なときにキューから値を取ってくる．GPGPU-simで実験しており，良い結果が出ている．G-Scalar: Cost-Effective Generalized Scalar Execution Architecture for Power-Efficient GPUs(Z. Liu et al., 2017, HPCA)
WIR: Warp Instruction Reuse to Minimize Repeated Computations in GPUs(K. Kim and W. W. Ro, 2018, HPCA)
Dimensionality-Aware Redundant SIMT Instruction Elimination(T. T. Yeh, R. N. Green and T. G. Rogers, 2020, ASPLOS)
スレッドブロック内に多く存在している冗長な命令を除去する手法．スレッドブロックが2次元だと，スレッドIDやブロックの次元，ブロックIDなどについて，それぞれのwarp内のスレッドが結局同じ値を持つことになることがある（論文中図３参照）．特に，ブロックの次元とブロックIDについては，同じブロック内のスレッドは必ず同じ値を持つので，その値をいちいち引っ張ってきたり，それに依存する値を計算するとwarp間で冗長である．そんな冗長な命令に，コンパイル時にアノテーションを付け，実行時にフェッチする前にスキップするようにしている．分岐発散が起きたときは，そのwarpはスキップの対象から外している．GPGPU-sim上で実験し，良い結果が出ている．
Others
Many-Thread Aware Prefetching Mechanisms for GPGPU Applications(J. Lee et al., 2010, MICRO)
Warped-Compression: Enabling Power Efficient GPUs through Register Compression(S. Lee et. al., 2015, ISCA)
Enabling coordinated register allocation and thread-level parallelism optimization for GPUs(X. Xie et al, 2015, MICRO)
CRAT: Enabling Coordinated Register Allocation and Thread-Level Parallelism Optimization for GPUs(X. Xie et al, 2018, IEEE Transactions on Computers)gpucc: An Open-Source GPGPU Compiler(J. Wu et. al., 2016, CGO)
gpuccというLLVMベースでオープンソースなCUDA向けコンパイラを作ったよという論文．フロントエンドはClangベースで，LLVM IR，NVPTXとコンパイルしていくようだ．LLVM IRに対する最適化も作っているが，個々の最適化手法は特に新しいものはない．オープンソースベンチマークを使ってnvccとの比較を行っているが，geomeanで見るとgpuccの方が特段速いコードを生成できるというわけでもないようだ．コンパイル時間の方はgpuccの方がnvccよりも8%速い．Balancing Scalar and Vector Execution on GPU Architectures(Z. Chen and D. Kaeli, 2016, IPDPS)
Orion: A Framework for GPU Occupancy Tuning(A. B. Hayes et al., 2016, Middleware)
Launch-Time Optimization of OpenCL GPU Kernels(A. S. D. Lee and T. S. Abdelrahman, 2017, GPGPU)
カーネルの起動時に定数となるものについて，コンパイル時にアノテーションを付けることで，定数伝播や定数畳み込みなどを行えるようにし，冗長性を除去している．Clangを弄ってアノテーションが付されたPTXを出力するようにしている．実験結果では，そんなに良くなっている印象はない．Compiler Assisted Coalescing(S. Puthoor and M. H. Lipasti, 2018, PACT)
HPVM: Heterogeneous Parallel Virtual Machine(M Kotsifakou et. al., 2018, PPoPP)
Warp Scheduling for Fine-Grained Synchronization(A. ElTantawy and T. M. Aamodt, 2018, HPCA)
IGC: The Open Source Intel Graphics Compiler(A. Chandrasekhar et. al., 2019, CGO)
CORF: Coalescing Operand Register File for GPUs(H. A. Esfeden et. al., 2019, ASPLOS)
ApproxHPVM: a portable compiler IR for accuracy-aware optimizations(H. Sharif et al., 2019, OOPSLA)
Optimizing Occupancy and ILP on the GPU using a Combinatorial Approach(G. Shobaki, A. Kerbow and S. Mekhanoshin, 2020, CGO)
GPUの占有率と命令レベル並列性(ILP)をどちらも向上させようという手法．本来は，ILPを上げようとするとレジスタが多く必要になるので，並列に走れるスレッド数が減り，占有率が下がってしまう．この手法では，最初に最大の占有率となるようなスケジュールを見つけ，次にその占有率を保ちながらILPが最大になるスケジュールを探す．LLVM上に実装し，AMDのGPUに対して実験を行っている．機械学習系のベンチマークに対して実験をしており，かなり速くなっている．Thread-Level Locking for SIMT Architectures(L. Gao et. al., 2020, IEEE Transactions on Parallel and Distributed Systems)
TAFE: Thread Address Footprint Estimation for Capturing Data/Thread Locality in GPU Systems(K. Punniyamurthy and A. Gerstlauer, 2020, PACT)
Fireiron: A Data-Movement-Aware Scheduling Language for GPUs(B. Hagedorn et al., 2020, PACT)
Model-Based Warp Overlapped Tiling for Image Processing Programs on GPUs(A. Jangda and A. Guha, 2020, PACT)
Valkyrie: Leveraging Inter-TLB Locality to Enhance GPU Performance(T. Baruah et al., 2020, PACT)
Modeling and analyzing evaluation cost of CUDA kernels(S. Muller and J. Hoffmann, 2021, POPL)
Are dynamic memory managers on GPUs slow?: a survey and benchmarks(M. Winter et al., 2021, PPoPP)
ApproxTuner: a compiler and runtime system for adaptive approximations(H. Sharif et al., 2021, PPoPP)
Repurposing GPU Microarchitectures with Light-Weight Out-Of-Order Execution(K. Iliakis, S. Xydis and D. Soudris, 2022, TPDS)
PLDI2022参加体験記
言わずと知れたプログラミング関係のトップカンファレンスであるPLDIに参加してきたので，忘れないうちに感想を書いておこうと思う．
PLDIとは
ACM SIGPLAN International Conference on Programming Language Design and Implementation (PLDI) は，プログラミング関係のことを広く扱うトップカンファレンスである．2022年でもう43回目の開催である．名前に「Design and Implementation」と入っているが，基礎理論から応用アプリケーション，パフォーマンス関係のことまで本当に広く扱っている．一般の投稿論文の採択率は，15〜25％あたりで，かなり低めである．PLDI 2022は，アメリカ合衆国カリフォルニア州サンディエゴで，6/13〜6/17にハイブリッド開催された．13日はチュートリアルなど，14日はISMMやLCTESやワークショップなど，15〜17日がPLDIのメイントラックだった．
LCTES
PLDIに参加したのは，併設開催のLCTESに私の論文がめでたく採択され，口頭発表するためであった．LCTESは，ACM SIGPLAN/SIGBED International Conference on Languages, Compilers, and Tools for Embedded Systemsの略で，組み込みシステム関係のことを広く扱った国際会議である．LCTES 2022では，コードサイズや消費電力の削減，FPGA関係のことが多かったように思う．「ように思う」と書いてあるのは，自分の発表を終えるまで緊張しすぎてあまり覚えていないからである．今回の発表が国際会議での初の発表だったので，久しぶりにガチで緊張していた．あまり覚えていないながらも，印象に残っている発表が２つあるので，以下に軽く紹介する．
Tighten Rust’s Belt: Shrinking Embedded Rust Binaries
Rustはバイナリコードのサイズが，Cと比べて大きくなってしまう傾向にあるらしい．ChromeOS向けのCとRustのファームウェア実装のバイナリサイズを比べているが(Table 2, 3)，全体でRust実装の方が1.8倍ほど大きくなってしまっている．この肥大化の理由として，次の４つの理由を挙げている：monomorphization, compiler generated support code, hidden data structures and data, fewer compiler optimizations．それぞれについて，４章と６章で議論されている．ここ数年は，RustをCの代わりに使ってOSなどを開発する流れになっているので，このあたりの話は興味深い内容だった．RollBin: Reducing Code-Size via Loop Rerolling at Binary Level
コードサイズ削減のための最適化は，ソースコードか中間表現レベルでよく行われるが，バイナリコードだけが与えられたときにはその最適化技術の恩恵は受けれない．そこで，バイナリコードのレベルでループをリロール(reroll)することで，コードサイズを削減するRollBinを提案している．ループアンローリング(loop unrolling)は，よく行われる最適化だが，それによって余分な命令が増えてしまうことがある(Figure 1(b))．RollBinでは，この余分に増えている分をなんとかしたらしい．バイナリレベルで，ループと帰納変数の認識，繰り返しの認識，データ依存の認識などをやるなんて大変だ．実験では，有名なSPECベンチマークで，LLVM13のコードサイズ最適化をしたものをベースラインにして，関連手法と比べていて，良い結果が出ている．
PLDI 2022
もうCOVID-19は終息したと言わんばかりの大盛況だった．200〜300人ほどは対面で参加していただろうか．もっと多かったかもしれない． 昼食はパシフィック・ビーチにテーブルが並べられ，ギラギラ照りつける太陽と爽やかな潮風を満喫しながらの食事だった． カリフォルニアにはJune Gloomはあるが，日本の梅雨のような蒸し暑さはなく，基本的にはカラッとした気候だった． Tea Timeには，コーヒーやお菓子が並べられ，みんな立ったままディスカッションや世間話をしていた． 会場には，スポンサーであるJane StreetのTシャツが置かれ，一人一つ取っていってねと書いてあった． デザインはOCamlのラクダが何匹もいるデザインで，個人的には気に入っている．OCaml 5.0のリリースに合わせて作ったのだろうか．
今回のPLDIには，SIGPLANセッションがあり，2020年〜2022年のPOPLやPLDI，OOPSLAなどの発表があった．対面でやりたかったということだろうか． LCTESでもそうだったが，発表は録画され，12時間後に配信される．その配信を見て質問したくなったら，Slackで質問できるという環境で，アジア圏でオンライン参加の人でも参加しやすいようになっていた． 生配信が見たい人は，youtubeのライブ配信を見るか，Gather townに参加して見るという環境だった． 個人的には，15日のCompilationのセッションが印象深かったので，軽く紹介しようと思う．
Finding Typing Compiler Bugs
コンパイラのバグを見つける話．確かdistinguished paperに選ばれていた．著者らが開発したHephaestusが2021年にgroovyc, kotlinc, javacについて調査した結果，156個ものバグがあったらしい．理論的なところは分からなかったが，話自体はとても興味深いものだった．それにしても，型周りの話は沼である．IRDL: An IR Definition Language for SSA Compilers
Intermediate representation(IR)は，昨今では必須と呼ぶべきものだが，新しいIRが欲しくなったときに，それを手動でデザインして修正してツールやら最適化やら作って...というのはコストが高い．そこで，IRを定義するための言語IRDLを提案している．LLVMのMLIRは，28種類もの「方言」(Table 1)(色々な命令セットとか)をカバーしているが，それらに対してIRDLがどれくらい上手く表現できるか実験している．
英語力について
LCTESのところでも述べたが，今回の発表が初めての国際会議での発表だった．それに加えて，日本を出たのも初めてで，まともに英語を話したのも初めてだった．今回の経験から言えることは，圧倒的に自分の英語力が不足していて，それによる機会損失は本当に大きいということだ．研究の裾野を広げるためにも，人脈を広げるためにも，ある程度の英語力はつけていきたい．
今回の経験から言えるその他のこととして，「意外と相手の言いたいことが分かったこと」と，「意外と自分の言いたいことを伝えられた（と思う）こと」がある．これらはおそらく，日本にいながらもリスニングとスピーキングの訓練を地道に続けていたからだと思う．訓練といっても，英会話教室などに通っているわけではない．リスニングの訓練は，PodcastのアプリでBBCのニュースなどを聴いていて，スピーキングの方は，「スラスラ話すための瞬間英作文シャッフルトレーニング」（森沢洋介）という本を何週かこなしているだけである．おそらくこの本でトレーニングしていなかったら，現地で本当に全く言葉が出なかっただろう．「瞬間的に言いたいことを口から出す」トレーニングは実際に有用だったので，スピーキングが苦手な人は是非やってみてほしい．それと，話さなきゃいけないときに，発音や文法の細かいことを考えて黙ってしまうのではなくて，とにかく話すのが大事だと思う．完璧な英語でなくとも意外と伝わるのが分かったのは良い経験だった．
おわりに
総じて，良い経験だった．トップカンファレンスを肌で感じるのは良い．研究のモチベーションに繋がる．また参加したい．
LLVM IRからGPUの実行可能ファイルを出力するコンパイルフローのまとめ
研究の過程で，cudaコードをclangに与えて出力されるLLVM IRを，GPUの実行可能ファイルまでコンパイルする必要があったので，そのコンパイルフローをまとめる．
環境
PCの環境は以下．
- OS: Ubuntu 18.04 LTS
- GPU: Geforce GTX TITAN RTX
- clang,LLVM 11.1.0
- CUDA 11.1
コンパイルフロー
まず，私の目的としては，cudaコード（.cu）をclangに与えて得られるLLVM IR（ホスト側とデバイス側の２つの.llファイル）に，最適化とかしてから，実行可能ファイルにコンパイルしたい． 色々試したところ，結局，「clangが実行可能ファイルを出力する過程をトレースすればいいじゃん」と気づき，おもむろに次のコマンドを試した．
clang++ -v mytest.cu -o mytest --cuda-gpu-arch=sm_75 -L/usr/local/cuda/lib64 -pthread -lcudart_static -ldl -lrt
ここで，mytest.cuにはホスト側とデバイス側のどちらのコードも書いてある．このコマンドによって，mytestというファイルが生成されるわけだが，-vオプションを指定することで，どのように生成されたかが分かる．その結果は，あまりにも長いので全部は貼らないが，LLVM IRを経由して実行可能ファイルまでたどり着くためのコンパイルフローを以下にまとめる．
cu --> ll
まず，cudaコード（.cu）から，LLVM IR（.ll）を得るには，次のコマンドを実行する．
clang++ -S -emit-llvm mytest.cu --cuda-gpu-arch=sm_75
指定するアーキテクチャは適宜，自分の環境に合わせて読み替えてほしい．結果として，mytest.ll（ホスト側）とmytest-cuda-nvptx64-nvidia-cuda-sm_75.ll（デバイス側）のllファイルが得られる．
device.ll --> ptx
次に，デバイス側のllファイルをptxに変換する．LLVMの公式サイトを参考に，次のコマンドを実行する．
llc -mcpu=sm_75 mytest-cuda-nvptx64-nvidia-cuda-sm_75.ll -o mytest-cuda-nvptx64-nvidia-cuda-sm_75.ptx
ptx -> .o
次に，ptxファイルをオブジェクトファイルにコンパイルする．コマンドは以下．-vオプションは，私がいつも指定しているだけで，必要はないだろう．
ptxas -m64 -O3 -v --gpu-name sm_75 -o mytest-cuda-nvptx64-nvidia-cuda-sm_75.o mytest-cuda-nvptx64-nvidia-cuda-sm_75.ptx
ptx, o --> fatbin
次に，ptxファイルとオブジェクトファイルから，デバイス側のバイナリを含むfatbinaryを生成する．
fatbinary -64 --create mytest-cuda-nvptx64-nvidia-cuda-sm_75.fatbin --image=profile=sm_75,file=mytest-cuda-nvptx64-nvidia-cuda-sm_75.o --image=profile=compute_75,file=mytest-cuda-nvptx64-nvidia-cuda-sm_75.ptx
fatbin, host.cu --> host_embedded_fatbin.ll
次に，ホスト側のllファイルに，先ほど作成したfatbinaryを埋め込む必要がある．このステップが重要で，埋め込みをしないでコンパイル-->リンクとやっていくと，カーネル関数が実行できなくなってしまう．実行するコマンドは，clangに-vオプションを渡したときのものを各自参考にしてほしい．clangに-fcuda-include-gpubinaryオプションを渡しているコマンドが，実行すべきコマンドである．-internal-isystemオプションが満載で長すぎるので，適宜省略したものを，以下に記す．（...の部分は省略を表す）
clang++ -cc1 -triple x86_64-pc-linux-gnu -target-sdk-version=11.0 -aux-triple nvptx64-nvidia-cuda -emit-llvm ... -main-file-name mytest.cu ... -fcuda-include-gpubinary mytest-cuda-nvptx64-nvidia-cuda-sm_75.fatbin ... -x cuda mytest.cu -o mytest_embedded_fatbin.ll
ちなみに，clangの内部では，llファイルではなく，オブジェクトファイルを生成し，それをリンクすることで実行可能ファイルを生成している． 私は，ホスト側とデバイス側の両方のllファイルをいじりたいので，llファイルを出力している．
ll --> executable
最後は，↑で得たllファイルをただふつうにコンパイルするだけである．
clang++ mytest_embedded_fatbin.ll -o mytest -L/usr/local/cuda/lib64 -lcudart
以上のステップで，実行可能ファイルmytestが手に入る．
まとめ
LLVM+CUDAの環境では，どうやらllファイルにfatbinaryを埋め込まなければカーネルが実行できないようだ．clang -vでclangが何しているのか確認すれば，clangともお友達になれる気がする．
参考webサイト
Compiling CUDA with clang
User Guide for NVPTX Back-end
NVCC::CUDA Toolkit Documentation
CUDA Binary Utilities::CUDA Toolkit Documentation
各学会の締め切りなどのまとめ
この記事は，各学会の開催時期や締め切りをまとめた備忘録記事である．
国内学会
情報処理学会 プログラミング研究会 PRO
開催時期は，６，７，１０，１，３月．それぞれ約２ヶ月前に発表申し込みを行い，約一ヶ月前に論文を投稿する．投稿論文は発表会後にアクセプトされれば，情報処理学会論文誌 プログラミングに掲載される．論文投稿をしないで，発表だけ行うということもできる．英語論文で投稿しアクセプトされるか，日本語論文をアクセプト後に英語論文にすることで，Journal of Information Processing (JIP)に掲載できる．JIPはオープンアクセスで，dblpにも載る．情報処理学会 全国大会
開催時期は３月．申し込み締め切りは12月．参加したことないので，次は参加したい．日本ソフトウェア科学会 大会
開催時期は，８月下旬〜９月上旬．登壇発表申し込み締め切りは７月で，講演論文やポスターの締め切りは８月．参加したことないので，次は参加したい．プログラミングおよびプログラミング言語ワークショップ PPL
開催時期は３月上旬．カテゴリ1（国内外未発表論文）の発表申し込みは１２月下旬，論文提出は１月上旬．カテゴリ2（国外既発表論文）の発表申し込みは１月下旬．カテゴリ３（ポスターなど）の発表申し込みは２月上旬．例年，合宿形式で温泉に行く．ポスターセッションは，お酒がOKなのでわいわいと盛り上がる．プログラミング・シンポジウム
開催時期は１月上旬．発表申し込み締め切りが１１月で，論文の提出締め切りがその３週間後くらい．これも例年，合宿形式で温泉に行く．
国際学会
とりあえず，列挙する．
Programming Language Design and Implementation, PLDI
開催時期は６月．提出時期は１１月下旬．Principles of Programming Languages, POPL
開催時期は１月．提出時期は７月．Parallel Architectures and Compilation Techniques, PACT
開催時期は１０月．提出時期は４月．Code Generation and Optimization, CGO
開催時期は２月．提出時期は８月．Compiler Construction, CC
開催時期は２月．提出時期は11月．CGOの併設カンファレンス．Asian Symposium on Programming Languages and Systems, APLAS
開催時期は１２月．提出時期は６月．ECOOP
開催時期は７月．提出時期は１月．Languages and Compilers for Parallel Computing, LCPC
開催時期は１０月．提出時期は７月．
CFG，トポロジカルソート，Bit Vector，ワークリストアルゴリズム
はじめに
この記事は，以前に私がQiitaに投稿した記事である． ふと，はてなブログにも残しておきたいと思ったので，移植する．
Control Flow Graph(CFG)
CFGとは、基本ブロック(basic block)をノードとするグラフのこと。 まず、基本ブロックとは何かというところから説明したいと思います。
基本ブロック(basic block)
ある命令 n が唯一の先行節 p を持ち、また p の唯一の後続節が n であるとき、p と n の gen、kill、in、out を合わせて考えることができ、p と n を一つのノードにまとめることができます。 これが、基本ブロックの根底にある考え方です。基本ブロックは、次の定義を満たします。
- 基本ブロック内の最初の命令は、ラベル付きである。
- 基本ブロック内の最後の命令は、無条件ジャンプ、または条件付きジャンプである。
- 基本ブロック内のその他の命令は、ラベル付き、無条件ジャンプ、条件付きジャンプのどれでもない。
つまり、基本ブロックの途中では、どこかに分岐することもなく、どこかから合流することもないということです。また、この定義より、基本ブロック内の最初と最後の命令以外は、常に唯一の先行節と唯一の後続節を持ちます。よって、基本ブロックを１単位として、gen、kill、in、out の計算をすることができます。
簡単な場合として、命令 n が唯一の先行節 p を持ち、p の唯一の後続節が n であるときを考えてみます。
上の図みたいなのを想像してください。 外側の赤い枠が、p と n を合わせてできた基本ブロックだとします。その基本ブロックの in と out を、in[pn]、out[pn] とします。すると、見てわかるように、in[pn] は in[p] と等しく、同様に、out[pn] は out[n] と等しくなります。また、p と n はお互いが唯一の先行節、後続節の関係なので、out[p] と in[n] も等しくなります。 ここで、Reaching Definitionのデータフロー方程式を考えてみましょう。out[n]は以下のように書けます。
out[n] = gen[n] \cup (in[n] - kill[n])
in[n] = out[p] より、
out[n] = gen[n] \cup ((gen[p] \cup (in[p] - kill[p])) - kill[n])
これを少し整理してみると、
out[n] = (gen[n] \cup (gen[p] - kill[n])) \cup (in[p] - (kill[p] \cup kill[n]))
このように書けます。 これはつまり、この p と n を合わせてできた基本ブロックの gen[pn] と kill[pn] が次のように書けることを意味しています。
gen[pn] = gen[n] \cup (gen[p] - kill[n]) kill[pn] = kill[p] \cup kill[n]
よって、out[n] = out[pn]、in[p] = in[pn] なので、次の式が得られます。
out[pn] = gen[pn] \cup (in[pn] - kill[pn])
この out[pn] が、 p と n を合わせてできた基本ブロックの out の値です。これで一つの命令単位でなく、一つの基本ブロック単位で in と out を計算することができます。 このようにして、基本ブロック単位でデータフロー方程式を解くことで、より高速に解を得ることができます。
基本ブロックのトポロジカルソート
トポロジカルソートとは、グラフ理論におけるノードの順序付け方法の一つです。
これをCFGに対して適用すると、あるノードが実行されるとき、そのノードのすべての先行節が実行されているように、ノードを順序付けすることができます。
例えば、次のようなCFGがあったとします。
各ブロックの左側の黒い数字は、このCFG内のそのブロックの固有の番号(basic block id)です。このCFGにトポロジカルソートを適用させると、ブロックを辿る順序は、各ブロックの右側の赤い数字の順番になります。上の図を見ると、あるブロックに到達したときには、そのブロックの全ての先行節がすでに訪問されている状態になっていることがわかると思います。 トポロジカルソートを求めるアルゴリズムは色々あるらしいですが、ここでは深さ優先探索(depth first search)を用いた方法を紹介します。
function DFS(n):
if visited[n] = false
visited[n] ← true;
foreach s in succ[n]
DFS(s);
endfor;
topo_sorted[N] ← n;
N ← N - 1;
endif;
endfunction;
procedure Topological_Sort:
N ← the number of nodes
foreach n : 0 to N
visited[n] ← false;
endfor;
DFS(entry_node);
endprocedure
Topological_Sort手続きが終わったときに、topo_sorted配列に格納されているのがトポロジカルソートされたブロックの順序です。 Reaching DefinitionやAvailable expressionsなどは、先行節のoutの情報を用いてinを計算するforward analysisなので、全ての先行節がすでに訪問され、outが計算されていれば、効率よくデータフロー解析を行うことができます。なので、forward analysisでは、このトポロジカルソートした順序でブロックを訪問します。 また、Liveness analysisなどのbackward analysisは、逆に後続節のinの情報を用いてoutを計算するので、トポロジカルソートの逆順でブロックを辿る必要があります。逆トポロジカルソートは、上記のアルゴリズムで、entry_nodeではなくexit_nodeからDFSを適用し、DFS関数内で後続節の代わりに先行節へのエッジを辿ることで得られます。
Bit Vector
データフロー解析では、フローされているデータは何らかの集合です。それは、一つの命令だったり、変数だったりしますが、そのような有限集合は、Bit Vectorで表すことができます。 Bit Vectorで表すと何がうれしいかというと、集合のAND演算やOR演算を高速に処理することができるようになります。
Reaching Definitionでの例を考えてみます。
1: a ← 1 2: b ← 2 3: if x < 10 goto L1 else goto L2 4: L1: a ← 5 5: goto L3 6: L2: b ← 3 7: c ← a + b 8: goto L3 9: L3: d ← a + b
| n | gen[n] | kill[n] | in[n] | out[n] |
|---|---|---|---|---|
| 1 | 1 | 4 | 1 | |
| 2 | 2 | 6 | 1 | 1,2 |
| 3 | 1,2 | 1,2 | ||
| 4 | 4 | 1 | 1,2 | 2,4 |
| 5 | 2,4 | 2,4 | ||
| 6 | 6 | 2 | 1,2 | 1,6 |
| 7 | 7 | 1,6 | 1,6,7 | |
| 8 | 1,6,7 | 1,6,7 | ||
| 9 | 9 | 1,2,4,6,7 | 1,2,4,6,7,9 |
このような例を考えてみたとき、例えば、９番のノードの in は、５と８番のノードの out の和集合で表されます。これをBit Vectorで表してみると、
| 1 | 2 | 4 | 6 | 7 | |
|---|---|---|---|---|---|
| out[5] | 0 | 1 | 1 | 0 | 0 |
| out[8] | 1 | 0 | 0 | 1 | 1 |
| in[9] | 1 | 1 | 1 | 1 | 1 |
このようになります。out[5]は２と４番の命令に対してビットが立っていることを表し、同様にout[8]は１と６と７番の命令に対してビットが立っていることを表しています。
in[9] = out[5] \cup out[8]
であるので、上の表を見ても分かるように、ビットのOR演算をした結果が、in[9]に格納されます。 逆に、積集合を計算したいときは、ビットのAND演算をすれば良いのです。 このように集合をビット演算で計算できるので、より高速にデータフロー解析をすることができるようになります。
ワークリストアルゴリズム
これは、データフロー方程式を再計算しなければならないところだけ、再計算するというアルゴリズムです。例えば、Reaching Definitionのアルゴリズムはこのように書けるのでした。
foreach n
in[n] ← {};
out[n] ← {};
endfor;
do
foreach n
in2[n] ← in[n];
out2[n] ← out[n];
foreach p in pred[n]
in[n] ← union(in[n],out[p]);
endfor;
out[n] ← union(gen[n],in[n]-kill[n]);
endfor;
while in[n] = in2[n] and out[n] = out2[n] for all n;
このアルゴリズムでは、いずれかのin,out集合が変更されると、また全ての命令についてループを回し、再計算していました。これだと、再計算する必要のないところまで再計算していて、非効率です。そこで、再計算されるべきところだけワークリストに保存して、再計算するというのがこのアルゴリズムです。 Reaching Definitionのワークリストアルゴリズムは、以下のように書くことができます。
worklist ← set of all nodes;
while worklist ≠ empty
pop n from worklist;
old_out ← out[n];
foreach p in pred[n]
in ← union(in,out[p]);
endfor;
out[n] ← union(gen[n],in-kill[n]);
if old_out ≠ out[n]
foreach s in succ[n]
if s ∉ worklist
push s into worklist;
endif;
endfor;
endif;
endwhile;
ループのないCFGならトポロジカルソート順に基本ブロックを訪問すれば、全ブロックの計算は２回で済みますが、ほとんどのプログラムはループを含んでいます。ループがあると、全ブロックの再計算が何回も行われることもあり、非効率的です。しかし、このワークリストアルゴリズムでは、再計算はループの中だけに留まるので、より効率的にデータフロー方程式の解を得ることができます。
参考記事
Reaching Definition
Common Sub-expression Elimination
Dead Code Elimination
Dead Code Elimination(無用コード除去, DCE)
はじめに
この記事は，以前に私がQiitaに投稿した記事である． ふと，はてなブログにも残しておきたいと思ったので，移植する．
Liveness analysis
そのまま日本語に訳すと、生存解析。ある変数の定義がプログラムのどこで使われるのか解析する手法です。 これによって、無用コード除去(Dead Code Elimination, DCE)ができるようになります。
概要
例えば、以下のようなプログラムsample.cがあったとします。
void main(int v){
int a,b,x;
a = 1;
b = 2;
if(v < a){
x = a + b;
} else {
a = 10;
b = 5;
}
a = 100;
x = a * b;
b = 20;
}
このような場合は、if文のfalse側の a = 10; で格納された a の値は使われないです。また、最後の b = 20; の b の値もその後使われません。 そのような使われない値はレジスタに保存しておくのも無駄であるので、命令を削除してしまおうというのがDead Code Eliminationです。
アルゴリズム
それでは、アルゴリズムの説明に入ります。 genとkillの計算表は以下のようになってます。
| Statement s | gen[s] | kill[s] |
|---|---|---|
| t ← a ⊕ b | {a, b} | {t} |
| t ← Mem[a] | {a} | {t} |
| Mem[a] ← b | {b} | {} |
| if a < b goto L1 else goto L2 | {a, b} | {} |
| f(a,b) | {a, b} | {} |
| t ← f(a,b) | {a, b} | {t} |
これは、それぞれ以下のような場合です。
１行目：何らかの計算をして、tに代入する場合。 ２行目：aの場所のメモリの値を、tにロードする場合。 ３行目：aの場所のメモリに、bの値をストアする場合。 ４行目：何らかの条件を判定して、指定した場所へジャンプする条件分岐。 ５行目：関数f()を、引数 a と b で呼び出す場合。 ６行目：関数f()を引数 a と b で呼び出し、その返り値をtに代入する場合。
次に、inとoutの計算式を示します。
in[n] = gen[n] \cup (out[n] - kill[n]) \\
out[n] = \cup_{s \in succ[n]}in[s]
ここで、succ[n]はnの後続節を表します。このLiveness analysisは、Reaching DefinitionやAvailable expressionsと違い、プログラムの実行順とは逆順でデータフロー解析をしていきます。この逆順の解析をbackward dataflow analysisと言います。 上記の式を、すべての命令についてinとoutの変化がなくなるまで繰り返します。以下に、疑似コードを示します。
foreach n
in[n] ← {};
out[n] ← {};
endfor;
do
foreach n
in2[n] ← in[n];
out2[n] ← out[n];
foreach s in succ[n]
out[n] ← union(out[n],in[s]);
endfor;
in[n] ← union(gen[n],out[n]-kill[n]);
endfor;
while in[n] = in2[n] and out[n] = out2[n] for all n;
まず、inとoutは空集合で初期化します。そのあと、データフロー方程式に沿って各inとoutを計算していきます。ここで、気を付けなければならないのが、バックワード解析では、outの方から計算します。また、union関数は引数の和集合を計算する関数です。 では、先のsample.cで例を考えてみましょう。
1: a ← 1 2: b ← 2 3: if v < a goto L1 else goto L2 4: L1: x ← a + b 5 goto L3 6: L2: a ← 10 7: b ← 5 8: goto L3 9: L3: a ← 100 10: x ← a * b 11: b ← 20
次に、gen、kill、in、outを計算すると以下のようになります。
| n | gen[n] | kill[n] | in[n] | out[n] |
|---|---|---|---|---|
| 1 | {a} | {v} | {a, v} | |
| 2 | {b} | {a, v} | {a, b, v} | |
| 3 | {a, v} | {a, b, v} | {a, b} | |
| 4 | {a, b} | {x} | {a, b} | {b} |
| 5 | {b} | {b} | ||
| 6 | {a} | |||
| 7 | {b} | {b} | ||
| 8 | {b} | {b} | ||
| 9 | {a} | {b} | {a, b} | |
| 10 | {a, b} | {x} | {a, b} | |
| 11 | {b} |
バックワード解析なので、11番のノードのoutから計算していきます。 ここで、注目すべきなのは、6番のノードのoutです。6番のノードでは、a ← 10 という操作をしていますが、out[6]に a がないので、その値は今後使われないということがわかります。同様に、11番のノードにも注目してみると、out[11]には b が含まれないので、この b の値は今後使われないということがわかります。 使われないなら無駄なので削除してしまいましょうということで、sample.cをDCEしてみると以下のようになります。
void main(int v){
int a,b,x;
a = 1;
b = 2;
if(v < a){
x = a + b;
} else {
b = 5;
}
a = 100;
x = a * b;
sample2.cでは、無駄な命令が省かれて最適化されました。
DCEの注意点
sample.cは簡単なプログラムだったので、この値は使われない ⇒ よし、消してしまおう！ というようにできたのですが、いつでもこのようにできるとは限りません。 DCEで考慮しなければならない場合は、例外処理を発生させる可能性のある命令がある場合です。 普通の計算でもオーバーフローすると例外が発生する場合がありますし、0除算も例外が発生する場合がありますね。 そういった命令を安易に削除してしまうと、プログラムの意味が変わってしまう場合があります。最適化は、プログラムの意味を変えない範囲で行われるべきであるので、こういった場合は注意が必要です。
参考記事
Reaching Definition
Common Sub-expression Elimination
CFG，トポロジカルソート，Bit Vector，ワークリストアルゴリズム
Lazy Code Motion
Common Sub-expression Elimination(共通部分式削除，CSE)
はじめに
この記事は，以前に私がQiitaに投稿した記事である． ふと，はてなブログにも残しておきたいと思ったので，移植する．
Available expressions
日本語に訳すと、利用可能な式。つまり、プログラム中のあるポイントで利用可能である式があるかどうか解析する手法です。 Available expressionsを解析することによって、Common Sub-expression Elimination(共通部分式削除、CSE)ができるようになります。
概要
例えば、以下のようなプログラムsample.cがあったとします。
void main(int a, int b){
int x,y,z;
x = a + b;
if(a < 5){
y = a + b;
} else {
a = 10;
x = a - b;
}
z = a + b;
return;
}
このような場合は、見てわかるように、a + b の値が２回計算されてしまっていて無駄です。 if文のtrue側のy = a + b; の命令を実行するときには、すでに a + b の値を一回計算していて、しかもその値をそのまま使えるので、使ってしまったほうが計算回数が少なくすみます。 しかし、if文を抜けたあとの z = a + b; では、以前に計算した a + b の値は利用することができません。これは、if文のfalse側で、a の値を変更しているからです。 Available expressionsでは、一時変数 t1,t2 を導入して、次のように変形します。
void main(int a, int b){
int x,y,z;
int t1,t2;
t1 = a + b;
x = t1;
if(a < 5){
t1 = a + b;
y = t1;
} else {
a = 10;
t2 = a - b;
x = t2;
}
t1 = a + b;
z = t1;
return;
}
t1はa + bについての一時変数で、t2はa - bについての一時変数です。 Available expressionsでは、まず最初に、
d ← a + b
という計算があったら、a + b用の一時変数tを導入し、
t ← a + b \\ d ← t
という変換をします。このように変換しておけば、データフロー解析後にa + b の値が使えるとわかったら、t ← a + b の命令を消せばいいだけなので、後々楽です。
アルゴリズム
それでは、アルゴリズムの説明に入ります。 Reaching Definitionと同じように、以下の表のようにgenとkillの計算をします。
| Statement s | gen[s] | kill[s] |
|---|---|---|
| t ← a ⊕ b | {a ⊕ b} - kill[s] | expressions containing t |
| t ← Mem[a] | {Mem[a]} - kill[s] | expressions containing t |
| Mem[a] ← b | {} | expressions of the form Mem[x] |
| f(a,b) | {} | expressions of the form Mem[x] |
| t ← f(a,b) | {} | expressions containing t and expressions of the form Mem[x] |
これは、それぞれ以下のような意味です。
１行目：何らかの計算をして、tに代入する場合。 ２行目：aの場所のメモリの値を、tにロードする場合。 ３行目：aの場所のメモリに、bの値をストアする場合。 ４行目：関数f()を呼び出す場合。 ５行目：関数f()を呼び出し、その返り値をtに代入する場合。
kill[s]にある、expressions containing t とは、t を含む式のことです。つまり、t の値は変更されたので、t をオペランドに持つような式(t + x など)は利用可能でないからkillするということです。 また、expressions of the form Mem[x]とは、Mem[x]の式(任意の x の値について)をキルするということです。メモリアクセスについては、プログラムのどの場所で、メモリのどの位置にアクセスするかといったことを正確に解析するのが難しいです。なので、メモリに値をストアする可能性がある場合は、Mem[x]形式の式はすべてキルします。 また、表中にないような条件分岐文などは、genとkillの計算には関与しません。
次に、inとoutの計算をします。ある命令nについてのinとoutは以下の数式で計算します。
in[n] = \cap_{p \in pred[n]}out[p] \\
out[n] = gen[n] \cup (in[n] - kill[n])
ここで、pred[n]はnの先行節を表します。上記の式を、すべての命令についてinとoutの変化がなくなるまで繰り返します。以下に、疑似コードを示します。
in[entry_node] ← {};
out[entry_node] ← {all expressions};
foreach n
if n is not entry_node then
in[n] ← {all expressions};
out[n] ← {all expressions};
endif;
endfor;
do
foreach n
in2[n] ← in[n];
out2[n] ← out[n];
foreach p in pred[n]
in[n] ← intersection(in[n],out[p]);
endfor;
out[n] ← union(gen[n],in[n]-kill[n]);
endfor;
while in[n] = in2[n] and out[n] = out2[n] for all n;
inとoutは、エントリノードを除いて、全ての式集合で初期化します。エントリノードは空集合で初期化します。また、union()関数は引数の和集合を計算する関数で、intersection()関数は引数の積集合を計算する関数です。 では、先ほどのsample2.cの場合で、例を考えてみます。
1: t1 ← a + b 1*: x ← t1 2: if a < 5 goto L1 else goto L2 3: L1: t1 ← a + b 3*: y ← t1 4: goto L3 5: L2: a ← 10 6: t2 ← a - b 6*: x ← t2 7: goto L3 8: L3: t1 ← a + b 8*: z ← t1
一時変数を導入してみると、このように書けると思います。 次に、gen、kill、in、outを計算すると以下のようになります。
| n | gen[n] | kill[n] | in[n] | out[n] |
|---|---|---|---|---|
| 1 | {a + b} | {a + b} | ||
| 2 | {a + b} | {a + b} | ||
| 3 | {a + b} | {a + b} | {a + b} | |
| 4 | {a + b} | {a + b} | ||
| 5 | {a + b, a - b} | {a + b} | ||
| 6 | {a - b} | {a - b} | ||
| 7 | {a - b} | {a - b} | ||
| 8 | {a + b} | {a + b} |
a + b を計算している1,3,8番のノードでは、{a + b}をgen[n]に加えています。また、a - b を計算している6番のノードでは、{a - b}をgen[n]に加えています。 5番のノードではaの値を変更しているので、a + b と a - b はキルされます。 8番のノードには、4番のノードから{a + b}が、7番のノードから{a - b}がデータフローされてきますが、それらの積集合を取ると空集合となるので、8番のノードのin[n]は空集合です。
ここで、3番のノードのinに注目してみると、a + b があります。つまり、3番のノードの入口でa + b が利用可能であるので、3番のノードは消去することができます。この操作のことをCommon Sub-expression Elimination(共通部分式削除、CSE)と言います。 つまり、sample2.cをCSEすると、以下のようになります。
void main(int a, int b){
int x,y,z;
int t1,t2;
t1 = a + b;
x = t1;
if(a < 5){
y = t1;
} else {
a = 10;
t2 = a - b;
x = t2;
}
t1 = a + b;
z = t1;
return;
}
上記のsample3.cを見るとわかるように、無駄な計算が省かれ最適化されました。
参考記事
Reaching Definition
Dead Code Elimination
CFG，トポロジカルソート，Bit Vector，ワークリストアルゴリズム
Lazy Code Motion
