@tanakmura
@tanakmura
id:w_o
F社のほうから来ました
(と言えと言われた気がしたので一応)
実行時にGCNの命令を生成するライブラリ
プログラムのリアルタイム生成(数百usecオーダ)を実現する
ほぼなにもしないプログラム@i7-4700MQ で
約100倍の高速化を実現
単純な処理がいくらかあって、
それを実行時に並びかえたりしたい、みたいなことを言われた
(単純な処理=データを一個拾って少し演算して戻す)
普通の実装だと、 __kernel をいくつか作って、
それを呼び出す順序を変える
R290X
(http://en.wikipedia.org/wiki/Comparison_of_AMD_graphics_processing_unitsより)
bandwidth 320 GB/s = 単精度floatで80G要素/s
read/writeなので転送量半分とすると、40G要素/s
要素あたり100演算ぐらい無いと転送律速になる
各処理ごとにはそんなに演算は無い
環境にもよるけど、Haswell i7で100usec周辺
-
R290X のbandwidth = 320GB/s
1080p SFloatのread/write = 1920x1080x4x2=約16MiB
転送律速だとすると、大体各処理の演算時間は、
16MiB / 320GB = 500usec
処理 500usec + clEnqueueNDRangeKernel呼び出し 100usec = 600 [usec/処理]
FLOPS値とは一体…
複数の処理をいっこのカーネルでやる
レジスタ上で演算できるのでメモリ転送いらない
カーネル呼び出しオーバーヘッドを削減できる
/* before */
__kernel void k0(__global float *out, __global float *in) {
out[i] = f0(in[i]);
}
__kernel void k1(__global float *out, __global float *in) {
out[i] = f1(in[i]);
}
/* after */
__kernel void k01(__global float *out, __global float *in) {
out[i] = f1(f0(in[i]));
}
単純な処理がいくらかあって、
それを実行時に並びかえたりしたい、みたいなことを言われた
OpenCLは、実行時にCをコンパイルしているので
実装はそれほど難しくない
Cプログラム文字列をconcatして実現可能
…だが
clBuildProgramに時間がかかる
プログラム/環境に依存するけど、最低でも30msecぐらい
ちょっと許容できない
(ビルド時に-fbin-exeとか付けると対応するバイナリを出せるので
それで計測 )
全部 : 30msec
C → LLVM IR : 10msec
LLVM IR → AMD IL : 18msec
AMD IL → GCN ISA : 2msec
__kernel void f(__global int *out) { out[0] = 1; }
/* 上のCLと同じ処理をする */
static void
gen_code(struct AIISA_Program *prog, struct AIISA_CodeBuffer *buf)
{
aiisa_code_buffer_reset(buf);
aiisa_s_buffer_load_dword_immoff(buf, S(0), S(8), 0x4);
aiisa_s_waitcnt(buf, LGKMCNT(0));
aiisa_v_mov_b32(buf, V(0), S(0));
aiisa_v_mov_b32(buf, V(1), 129);
aiisa_tbuffer_store_format_x(buf, NFMT_FLOAT, DFMT_32, FLAG_OFFEN, 0,
ZERO, 0, S(4), V(1), V(0));
aiisa_s_endpgm(buf);
aiisa_replace_text(prog, buf);
}
用法・用量を守って正しくお使いください
ELF の .text の中に ELF が入っている。(何故…？)
-fbin-exe を付けてclBuildProgramした場合、
この中のELF(以下innerELF)の.text に
GCN のコードが含まれている。
Linux だと inner ELF はなんかセクションがダブってる。何これ？
[ 0] NULL 00000000 000000 000000 00 0 0 0
[ 1] .shstrtab STRTAB 00000000 0000fc 000028 00 0 0 0
[ 2] .text PROGBITS 00000000 0004d8 0008a8 00 0 0 0
[ 3] .data PROGBITS 00000000 000d80 001280 1280 0 0 0
[ 4] .symtab SYMTAB 00000000 002000 000060 10 5 1 0
[ 5] .strtab STRTAB 00000000 002060 00001b 00 0 0 0
[ 6] .text PROGBITS 00000000 00262f 000068 00 0 0 0 (これ)
[ 7] .data PROGBITS 00000000 002697 001280 1280 0 0 0
[ 8] .symtab SYMTAB 00000000 003917 000060 10 9 1 0
[ 9] .strtab STRTAB 00000000 003977 00001b 00 0 0 0
まあよくわからんのでスーパーダーティーハック…
まあそんな感じでダーティハックにダーティハックを重ねて動くように なった
https://github.com/tanakamura/ai-isa-jit
(ポインタを一個かきかえるだけのデモが動きます)
こういう実装でどうか？と聞いてみたら、
「別にコンパイルはリアルタイムでなくてもいいのだけど」
と言われた
とくにこれ以上の使い途思い付かないのでせっかくだから今日この場で供養してあげようと思って今日はここへ来ました
(Mantleがわかればcalling convetionとかなんかもうちょっとなんとかしようかと思ったがそんなことはなかった)
以下、ダラダラとGCNの話でもします
(まあ別に無理してai-isa-jit使う案件でもないけど)
64way SIMD とは一体…？
性能出すには少なくとも256 item必要
16way SIMD が 4個並んでて、16way SIMDのひとつが4cycleかけて1wave front = 64要素処理する、
みたいなことが書いてある
わかりづらいけど、4wave front いれないとフルで動かない
// vector mov が 128個入ったループを16384回回す
aiisa_s_movk_i32(buf, S(0), 0);
aiisa_s_movk_i32(buf, S(1), 10);
label0 = buf->cur;
aiisa_s_addk_i32(buf, S(0), 1);
for (i=0; i<128; i++) {
aiisa_v_mov_b32(buf, V(0), S(0));
}
aiisa_s_cmpk_le_i32(buf, S(0), 16384);
こういうのを入れてwork item 数を 1, 16, 64, 256 で実行して CodeXL に入れると、
| item数 | VALU Utilization[%] | VALU Busy[%] |
| 1 | 1.56 | 23.52 |
| 16 | 25 | 23.39 |
| 64 | 100 | 23.51 |
| 256 | 100 | 94.04 |
VALU Busy を 100% にするには 256item必要
float 演算はレイテンシ4なのでitem内のスケジューリングは不要
CellのSPUとか、CUDAは、全演算がSIMDなので、ループカウンタのインクリメントとかすると勿体ない気分になる。
// CUDAです
__global void f(int n) {
// この i は、全threadで同じ計算をするのでもったいない
for (int i=0; i<n; i++) { ... }
}
GCNならそんな心配いらない!!
ベクタとスカラで別々のレジスタ/ALUがある。
8つのwave frontから、
4つのベクタ演算、4つのスカラ演算を実行可能
アイテム間で変わらない演算はベクタ演算と並行して実行可能
(8wavefrontなので最低でも512item立ち上げる必要がある)
// コード 1
// ループ中に128個スカラ演算とベクタ演算が交互に出てくる
aiisa_s_addk_i32(buf, S(0), 1);
for (i=0; i<128; i++) {
aiisa_s_movk_i32(buf, S(1), 100); // s_ ではじまるのがスカラ演算
aiisa_v_mov_b32(buf, V(0), S(0)); // v_ ではじまるのがベクタ演算
}
aiisa_s_cmpk_le_i32(buf, S(0), 16384);
// コード 2
// 128個スカラ演算と128個のベクタが連続して出る
aiisa_s_addk_i32(buf, S(0), 1);
for (i=0; i<128; i++) {
aiisa_s_movk_i32(buf, S(1), 100); // s_ ではじまるのがスカラ演算
}
for (i=0; i<128; i++) {
aiisa_v_mov_b32(buf, V(0), S(0)); // v_ ではじまるのがベクタ演算
}
aiisa_s_cmpk_le_i32(buf, S(0), 16384);
| VALU Utilization | VALU Busy | SALU Busy | |
| ループ1 256item | 100 | 47.88 | 48.63 |
| ループ1 512item | 100 | 95.55 | 97.04 |
| ループ2 256item | 100 | 48.09 | 48.84 |
| ループ2 512item | 100 | 64.31 | 65.31 |
→ 8 wave front 詰める + スカラとベクタ演算を詰めることで
ワークアイテム間で変化しない値をベクタ演算とは別に実行可能
Stream Kernel Analyzer で出すと出てくる lgkmcnt とかについて
s_waitcnt lgkmcnt(0)GCNはメモリアクセスが非同期。
メモリアクセスの完了は各自で書かないといけない。
完了を待つのがs_waitcnt命令。3つのカウンタがあって、
メモリアクセスするごとに 1 増える
メモリアクセス完了ごとに 1 減るというようになっている。
このカウントがいくらか減るかまで待つ命令がs_waitcnt
vmcnt(0) とかすると、VMCNT が 0 になるまで待つ。
演算とメモリアクセスをオーバーラップさせやすくするには、
メモリアクセスとs_waitcntの間に演算をたくさん入れるべき
/* ループ 1. ロードのすぐあとに waitcnt */
for (j=0; j<16; j++) {
aiisa_s_buffer_load_dword_immoff(buf, S(3), S(8), 0x4);
aiisa_s_waitcnt(buf, LGKMCNT(0));
for (i=0; i<8; i++) {
aiisa_v_mov_b32(buf, V(0), S(0));
}
}
/* ループ 2. ロード後すこし空けて waitcnt waitcnt */
for (j=0; j<16; j++) {
aiisa_s_buffer_load_dword_immoff(buf, S(3), S(8), 0x4);
for (i=0; i<8; i++) {
aiisa_v_mov_b32(buf, V(0), S(0));
}
aiisa_s_waitcnt(buf, LGKMCNT(0));
}
| VALU Utilization | VALU Busy | |
| ループ1 256item | 100 | 51.16 |
| ループ1 512item | 100 | 78.46 |
| ループ2 256item | 100 | 75.25 |
| ループ2 512item | 100 | 97.89 |
→ 手でスケジュールすることでメモリアクセスを隠蔽できる
他よくわからない点
型変換モードやらオフセットの付けかたが色々ある。
get_global_id(0)を呼んでるかどうかで引数の渡されかたが変わる？
まあ全然理解してない。
ai-isa-jit というのを勢いで作ったけど特に使いみちがなかった
しかたないので普通にアセンブラとして使っていくらか計測した
ご静聴ありがとうございました。