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)
