LLVM Tutorial をやった

TensorFlow XLA が LLVM を使って neural network の最適化をやっている話とか、k0kubun 君の LLVM-based JIT Compiler for Ruby の話を聞いて触発されたので LLVM の勉強をしてみた。Crystal は LLVM を使っているし、LLVM を使えば emscripten のような言語変換処理系を作れるというのも聞いた覚えがあるので雰囲気だけでも掴んでおくと今後良いことがありそう。


https://llvm.org/docs/tutorial/
 kaleidoscope という言語を作りながらLLVMの機能を学んだ。

LLVM IR の生成と、JIT の実装をやった。

LLVM の API はバージョンが変わるとガラッと変わるっぽくて、サンプルコードと手元の LLVM のバージョンがあってないと色々エラーがでて厳しかった。 https://github.com/llvm-mirror/llvm の該当バージョンのブランチ (例: llvm 3.6 ) からサンプルコードを取ってきて、手元のバージョンも http://releases.llvm.org/ から落として合わせると吉。

Parser については、色々クラス (LL(2) など) があるので、もっと深掘りしたほうが良さそう。

SSA については wikipedia の記事を見るとまぁわかる。 https://en.wikipedia.org/wiki/Static_single_assignment_form

IR の構築は、LLVM の API を通して行う

最適化は LLVM Pass というやつを通すと行われる。どういう最適化をさせるかはオプションで選べる。See https://llvm.org/docs/Passes.html

最適化前

ready> def test(x) (1+2+x)*(x+(1+2));
ready> Read function definition:
define double @test(double %x) {
entry:
  %addtmp = fadd double 3.000000e+00, %x
  %addtmp1 = fadd double %x, 3.000000e+00
  %multmp = fmul double %addtmp, %addtmp1
  ret double %multmp
}

最適化後 (1 + 2 が 3 になるのは標準でされてしまうっぽい)

ready>  def test(x) (1+2+x)*(x+(1+2));
ready> Read function definition:
define double @test(double %x) {
entry:
  %addtmp = fadd double %x, 3.000000e+00
  %multmp = fmul double %addtmp, %addtmp
  ret double %multmp
}

JIT した結果の in-memory address pointer を取得して、関数ポインタとして実行すると、JIT 後の関数を呼び出せる。 JIT した結果のアドレスを持っておけば何回も使いまわせる。


所感

LLVM IR を構築しさえすれば、LLVM Passes (Optimizer) や LLVM JIT の恩恵を受けられる。 クロスコンパイルしてプラットフォームごとのバイナリにすることもできる。 動的言語っぽく走らせることも、バイナリを作って走らせることもできるようになるわけだし、これは良いものだなぁ。


Further readings:

Cプログラム高速化研究班を読んだ

とっかかりとしてとても良い本。ただ、薄い本であり、この本だけで十分な知見が得られるかというとそんなことはないので、後は実践あるのみ :muscle: 

Cプログラム高速化研究班 コードを高速化する20の実験と達人の技
片山善夫 
USP研究所 
売り上げランキング: 280,025

1章 CPUとコンパイラについてちょびっと

1章の内容は自分は知っていたので読み飛ばし。詳解システムパフォーマンス輪読会 6章 CPU および perf stat の内容に近かった。

  • キャッシュライン64bytes (8bytes)単位で転送される

  • アセンブリ読むときはレジスタの役割表を見ながらじゃないと読み解けなさそう

767C334A-6211-4A85-B34A-47C22F13C066.JPG

2章 実行コストの感覚を身につける

  • ループ処理を速くするには、単純に if の数を減らすために、中の処理をコピペしてたくさん並べる
    • あとで出て来るが、分岐を減らすだけでなく、CPUに載っている加算器4つを並列で使える効果がある
  • cc -S でアセンブリを生成
b = 0
do {
  a += b;
  a += b;
  a += b;
} while (--n > 0);

b = 0
do {
} while (--n > 0);

のように最適化されてしまうので、実験ではそれを防ぐために b = atoi(argv[0]) としたり苦労している話。

  • 3.2GHz 逆数の0.3ナノ秒がCPUのサイクルタイム
  • CPU内部に複数の加算器があるので、変数を分けて最後にマージしたほうが速い(sequentialにコードを書いても)。部品の数だけ並列にやってくれる。変数が一つだと直列になってしまう
  • 乗算器は加算器ほどの数が用意されてないし、遅い
  • 2の冪数の定数との乗算はコンパイラがシフト演算に最適化する。
    • 2の冪数+1の場合lea命令が使える。変数との乗算は最適化できない
  • メモリ 32kb 8way と言ったら 8本あって並列化されてるということ
    • 合計256kbということであってる?
  • if (b < 0) a++ 条件成立時の加算より、不成立時のジャンプ命令のほうが遅い。!!
  • 分岐命令があると、どちらに飛べばいいかわからなくなるので、先に並列実行しておくことができず、待ちになってしまう。分岐予測して先に実行しておき、間違えたら捨てる。投機的実行
  • 64bitでは関数の引数は6個までレジスタに入れて渡される。それ以上はスタック(メモリ上にある) <= だから gdb で見ると optimized out といつも出てるのか

3章 遅いのはどこか

  • gprof を使う
    • static link させたものだけプロファイリングされるので、gcc -p -static するな
    • プロファイラ用のライブラリをリンクしないと細かいのは取れない -lc_p
    • gcc に -pg オプションをつけて gmon.out を吐いておかないといけない <= 使いづらいなー

4章 達人の方法論

  • メモリのキャッシュラインを気にして matmulの演算順序を変える話
    • matmul のメモリアクセス順序話はどこでもでてくるなぁ
  • 加算器が4つ載っている
    • ループの展開(unrolling)で演算器4つ並列で使い切る話
    • 単に分岐命令の数を減らす効果ももちろんある
    • 要素が4の倍数じゃないとダメだからゴミ詰め込みしておかないといけないんだよな
  • strcmpは \0 を見つける為に、一文字ずつ比較
    • ワード単位でできないのは、間違えて超えてしまうとセグるから
  • memcmp は8バイトずつ比較できる
    • 8バイトでマッチしなかった場合に1バイトずつの比較にfallbackする
    • コンパイラにmemcmpがインライン展開されて1バイトずつの比較になってしまい、逆に遅くなってしまう事例。-fno-builtin-memcmp で抑制
  • sseなら1クロックで16バイトごと処理できる
    • メモリアクセスは速くならないわけだけど、どうなんだ??
    • sse命令書き方マニュアルどこだ??
  • パイプラインのがバッファは64kb
    • 出力先がファイルではなくパイプの場合は、バッファサイズを64kbにしたほうが速い
  • stdioのバッファは変更可能
    • デフォルト4kb 著者が実験した環境では4mbにしてファイルに書き込むほうが速かったようだ。

5章 コンパイラを骨までしゃぶる

gcc の最適化オプションなど。けっこう一般的な知識なので読み飛ばした。

6章 業務システム向けのヒント

半角文字を全角文字に変換するなどの、具体的な話なので必要になったら改めて読むかも。Unicodeの話は少し勉強になった。

CUDA C プロフェッショナル 第一章 〜 第三章

CUDA C プロフェッショナル第1章から第3章までのメモ。

この本はとても良かった。丁寧に詳細まで解説してある。中級〜上級者向け。Impressのサイトにソースコードがおいてあるので、落として手を動かしながら進めていくのが吉。

CUDA C プロフェッショナル プログラミング (impress top gear)
John Cheng Max Grossman Ty McKercher 
インプレス 
売り上げランキング: 42,483

第1章

nvcc -arch sm_35 hello.cu -o hello

CC(Compute Capability) 3.5 以降 (Kepler 世代以降) に最適なバイナリを生成する。それ以前では動かない。

nvcc -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35

gencode オプションを使うと複数世代に最適なコードを生成し、ワンバイナリに含めることができる

第2章

image.png

  • 1つのカーネル呼び出しによって生成されたすべてのスレッドを、まとめてグリッドと呼ぶ
  • ブロックとは、互いに協調して動作できるスレッドのグループであり、以下の機能を利用できる
    • ブロックに属するスレッド間の同期 (ブロックをまたいだ同期は難しい)
    • ブロック内で共有されるシェアドメモリ
  • kernel<<nBlocks, nThreadsPerBlock>> のように呼び出す
    • nBlocks 1グリッドあたりのブロック数
    • nThreadsPerBlock 1ブロックあたりのスレッド数
  • CUDAでは32本のスレッド(1 Warp)に対して1つのインストラクションを1度に実行する。
    • nThreadsPerBlock は 32の倍数にしておかないと、損をする
    • インストラクションは Warp ごとだが、レジスタとシェアドメモリと同期はブロックごと

Fermi 以降なら、カーネルで printf が使える

CUDA API の実行失敗をチェックするマクロ CHECK()。例えば
CHECK(cudaDeviceSynchronize()); とすればそこまでの処理全部に対して失敗しているものがいないか確認できる。

#define CHECK(call)                                                  \
{                                                                    \
    const cudaError_t error = call;                                  \
    if (error != cudaSuccess)                                        \
    {                                                                \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                \
        printf("code:%d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                          \
        exit(1);                                                     \
    }                                                                \
}

kernel を送り込むところで CPU で時間がかかる。なので、非同期に送り込みたくなるし、実際 kernel 実行は非同期になっている

Fermi の場合、
グリッドあたりの各次元(x,y,z,w)ごとの最大ブロック数は65536
ブロックあたりのスレッドの数は最大 1024

なお、現行の chainer(cupy) の場合は以下のように x しか使っていない。

134     cpdef linear_launch(self, size_t size, args, size_t shared_mem=0,
135                         size_t block_max_size=128, stream=None):
137         gridx = size // block_max_size + 1
138         if gridx > 65536:
139             gridx = 65536
140         if size > block_max_size:
141             size = block_max_size

スレッド数 128、ブロック数 = 要素数 // スレッド数(128)。
(65536 を超えるようだったら、y を使うようにしても良いかもしれないが、ソフトウェア的にブロック数を増やした所で、ハードウェアのスペックをすでに超えてるから意味ないだろうなぁ)

第3章

SIMDとSIMT

  • SIMD: 複数データにたいして1つのインストラクションを1度に実行する。全体で1つのインストラクション。命令ユニットが少なくてすむ。
  • SIMT: NVIDIA独自(というか提案?)のアーキテクチャ。SIMDの改善。32本のスレッド(1 Warp)に対して1つのインストラクションを1度に実行する。全体としては同時に複数のインストラクションを発行可能(別Streamを指定して発行)。

ワープダイバージェンス

  • 同じワープ内のスレッドが異なる命令を実行することをワープダイバージェンスと呼ぶ
  • ワープ内のスレッドが異なる分岐パスを選択した場合、ワープは各分岐パスを逐次的に実行する(分岐 true を実行してから、分岐 false を実行する)
    • SIMDの場合、そもそも分岐処理できないので、それの改良ではある。
    • ただし、遅くはなる

ダイナミックパラレリズム

  • kepler 以降
  • kernel から kernel 呼び出しができる

nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory

  • メモリ転送速度を帯域幅じゃなくて周波数でみることができる
  • 帯域幅 / 周波数 = bytes / サイクル

3.3 並列性の確保

nvprof --metrics arhieved_occupancy

  • 占有率 = アクティブワープの数 / ワープの最大数 を見る。GPUを有効に使えているかどうかの指標
  • ワープの最大数: SMあたりの最大スレッド数(M2070ではmaxThreadsPerMultiProcessor == 1,536) / 32
  • アクティブワープの数: 有効に使っているワープの数
  • 占有率をあげるには、基本的には nBlocks の数をあげれば良い
    • といって、nThreadsPerBlock を 128 より小さくして nBlocks を大きくしようとすると占有率はあがるものの、1ブロック(ハードウェア的にはSM)あたりのスレッド数が減って並列処理できなくなるので、遅くなる。<= Q. そこもあわせて見る指標はないのか?

nvprof --metrics gld_throughput

  • メモリ読み取りスループットのチェック
  • 補足: cudaMalloc とか cudaMemcpy とかではなくて、GPUプロセッサとGPUメモリ間の転送
  • (CPUでこれを見ようと思うと難しいので簡単に見れるのは嬉しい)
  • Q. ブロック数とスレッド数の変更で、メモリ読み取りスループットが変わる理屈がよくわからない。そもそも説明が書いてない。

nvprof --metrics gld_efficiency

  • メモリ読み取り効率のチェック
  • (さらに%で見れるとは嬉しい)

グリッドのサイズ(nBlock)とブロックのサイズ(nThreadsPerBlock)の決定に関するガイドライン

  • nThreadsPerBlock をワープサイズ(32)の倍数に保つ
  • nThreadsPerBlock は 128 または 256 からチューニングを始める
  • カーネルのリソース要件に従ってブロックサイズを調整する
  • nBlock がSMの数をはるかに上回るようにする (k80では13)
    • 元々のデータサイズが小さいと、SMの数を上回れなくなるので、データをまとめて一気に処理するようにするとか?

3.4, 3.5. リダクション

http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

(日本語) http://gpu-computing.gsic.titech.ac.jp/Japanese/Lecture/2010-06-28/reduction.pdf

うーむ、難しい。あとでまた読む

A Ruby and Fluentd committer working at DeNA. 記事本文および記事中のコード片は引用および特記あるものを除いてすべて修正BSDライセンスとします。 #ruby #fluentd #growthforecast #haikanko #yohoushi #specinfra #serverspec #focuslight
はてぶ人気エントリー

Google