AMD GCN isa と NVIDA SASS でのアドレス演算のチューニング
ちょっとしたきっかけがあって、AMD用のGPUとNVIDIA用のGPU両方で高速化作業を行いました。
そのときに得られた知見を書いておきます。
AMD GCN isa と NVIDIA SASS
GPUでプログラミングを行うときは、OpenCL、CUDA などを使うと思いますが、これらはより低レベルなGPU用の機械語にコンパイルされます。
AMD のGPUでは、これに”ISA”、NVIDIA の GPU では、これを”SASS”という名前が付けられています。ISAは一般的な単語なので、別に正しい名前があるかもしれませんが、ビルド時のオプションなどで指定する場合は、–isa 等を指定しているので、ここでは、GCN isa とします。
通常のGPUプログラミングでこれらを見ることはないかもしれませんが、レジスタがわずかに足りない場合、性能がわずかに足りない場合などは、これを調べると改善できる可能性があります。
今回の作業で得られた経験をもとに、SASS、GCN isaレベルでのチューニング手法に書いていきます。
SASS、GCN isa の確認方法
SASS は、CUDA Toolkit に含まれる、nvcc, nvdisas を使うと確認できます。
(昔似たような文章を書いたことがあるので http://int.main.jp/txt/sass/も参照してください)
- $ nvcc –cubin -arch=sm_20 などで cubin を出す
- $ nvdisasm file.cubin のようにして SASSを出す
GCN isa は AMD CodeXL(http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/)に含まれる、CodeXLAnalyzer を使うと出力できます。
$ CodeXLAnalyzer -s CL <input.cl> -k <kernel名> –isa <output.isa> -c <GPU名>
GPU 名は、一覧を書いた公式ドキュメントが見当たらないのでわからないですが、clGetDeviceInfo(CL_DEVICE_NAME) で確認できる名前を指定すればよさそうです。
共通のチューニング
いくらか見た範囲だと、アドレス計算はどちらのGPUでも無駄な命令が出やすいようです。これには以下の理由が考えられます
- 64bit環境ではアドレスは64bit整数になるが、どちらのGPUも32bit整数命令しか無い
- コンパイラがそれほど良い最適化をしてくれない
- 高性能なCPUと違ってスーパースカラの幅が狭いのでアドレス演算がパイプラインを埋めてしまう
たとえば、次のプログラムを
|
1 2 3 4 5 6 7 |
__global__ void f(int *p,int n) { int i; for (i=0; i<n; i++) { p[i*2] = 100; } } |
nvcc + sm_30 でコンパイルすると ループ部分は ↓ のようなSASSが出ます。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 |
.L_2: /*0070*/ SHL R2, R0, 0x1; /*0078*/ IADD R3, R0, 0x1; /*0088*/ IMAD R6.CC, R2, R5, c[0x0][0x140]; /*0090*/ SHL R3, R3, 0x1; /*0098*/ IMAD.HI.X R7, R2, R5, c[0x0][0x144]; /*00a0*/ IADD R2, R0, 0x2; /*00a8*/ IMAD R8.CC, R3, R5, c[0x0][0x140]; /*00b0*/ SHL R2, R2, 0x1; /*00b8*/ ST.E [R6], R12; /*00c8*/ IMAD.HI.X R9, R3, R5, c[0x0][0x144]; /*00d0*/ IADD R3, R0, 0x3; /*00d8*/ IMAD R10.CC, R2, R5, c[0x0][0x140]; /*00e0*/ SHL R3, R3, 0x1; /*00e8*/ IADD R0, R0, 0x4; /*00f0*/ IMAD.HI.X R11, R2, R5, c[0x0][0x144]; /*00f8*/ ST.E [R8], R12; /*0108*/ IMAD R2.CC, R3, R5, c[0x0][0x140]; /*0110*/ ST.E [R10], R12; /*0118*/ ISETP.LT.AND P0, PT, R0, R4, PT; /*0120*/ IMAD.HI.X R3, R3, R5, c[0x0][0x144]; /*0128*/ ST.E [R2], R12; /*0130*/ @P0 BRA `(.L_2); .L_1: |
ST.E(どこにも説明書いてないですが、おそらくメモリへのストアだと思います) が四回出ていることから、おそらく4回アンロールされているのだと思いますが、4回のストアをするだけのループにしては、22命令は多すぎではないかという気がします
まじめに検証したわけではないですが、NVIDIA、AMDどちらのGPUのコンパイラも、ポインタに関しては、かなりナイーブなコードを出すようです。最適化しやすい形に変形してあげる必要があるでしょう。
上の場合なら、
|
1 2 3 4 5 6 7 8 |
__global__ void g(int *p, int n) { int i; for (i=0; i<n; i++) { p[0] = 100; p+=2; } } |
こんな感じですね。
|
1 2 3 4 5 6 7 8 9 10 |
.L_6: /*0088*/ ST.E [R4], R2; /*0090*/ IADD R0, R0, 0x4; /*0098*/ ST.E [R4+0x8], R2; /*00a0*/ ST.E [R4+0x10], R2; /*00a8*/ ST.E [R4+0x18], R2; /*00b0*/ ISETP.LT.AND P0, PT, R0, R3, PT; /*00b8*/ IADD R4.CC, R4, 0x20; /*00c8*/ IADD.X R5, R5, RZ; /*00d0*/ @P0 BRA `(.L_6); |
9命令になりました。このぐらいなら妥当に見えます。
このように、アドレス演算は配慮するかしないかで、命令量は結構変わります。先日作業していたプログラムは、演算律速の問題だったので 1.5倍くらい変わったと思います。
また、これは、GPUで画像を扱う場合等に、(x,y)とGPUのスレッドを1:1で割り当てるのはあまり良くないということを示しています。
例えば、
|
1 2 3 4 5 6 |
__kernel void f(__global int *p, int step) { int x = get_global_id(0); int y = get_global_id(1); p[y*step + x] = op(p[y*step + x]); } |
こういうコードを書いてしまうと、各画素ごとにy*step + xの計算が入ってしまいます。もう少し大きい粒度でスレッドを割り当てて、
|
1 2 3 4 5 6 7 8 9 10 |
__kernel void f(__global int *p, int w, int step) { int y = get_global_id(0); int *p2 = p + y*step; for (int x=0; x<w; x++) { p[0] = op(p[0]); p++; } } |
このように修正することで性能が改善する可能性があります。
GCN isa のチューニング
GCN isa はwavefront中の64個のワークアイテムが同時に複数のデータに対して命令を実行するSIMD命令と、wavefront1個で実行されるスカラ命令があります。
CodeXLAnalyzer で確認できる命令列のうち、v_ で始まるものが、SIMD命令、s_ で始まるものがスカラ命令になります。
このSIMD命令とスカラ命令は、パイプラインが別になっており、十分な数のワークアイテム(おそらく8wavefront = 512)を投入すれば、2個のwavefrontから64個のSIMD命令と1個のスカラ命令を同時に実行することができます。
全ワークアイテムで共通するアドレス計算をスカラ命令にすることで、データ演算と並行してアドレス計算を実行できるようになります。
どのような演算がスカラ命令になるかはよく調べていないですが、コンパイラは以下の条件が成立する値をスカラ値とみなし、演算の入力が全てスカラ値だった場合にスカラ命令が出るようです。
- カーネル引数
- コンパイル時定数
- __constant からのロードで、アドレスがスカラ値
- スカラ演算の結果
スカラ値ではない値で分岐した中で決まった値もスカラ値にはなりません。
以上を踏まえて、
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 |
__kernel void f(__global int *p, int n) { int lid = get_local_id(0); if (lid < 32) { __global int *p2 = p + get_group_id(0) * n; for (int i=lid; i<n*32; i+=32) { p2[i]++; } } } /* label_0011: v_add_i32 v0, vcc, 32, v0 // 000000000044: 4A0000A0 v_cmp_gt_i32 s[12:13], s3, v0 // 000000000048: D108000C 00020003 v_add_i32 v2, vcc, 0x00000080, v1 // 000000000050: 4A0402FF 00000080 buffer_load_dword v3, v1, s[4:7], 0 offen // 000000000058: E0301000 80010301 s_waitcnt vmcnt(0) // 000000000060: BF8C0F70 v_add_i32 v3, vcc, 1, v3 // 000000000064: 4A060681 buffer_store_dword v3, v1, s[4:7], 0 offen // 000000000068: E0701000 80010301 s_and_b64 exec, exec, s[12:13] // 000000000070: 87FE0C7E s_cbranch_execz label_0020 // 000000000074: BF880002 v_mov_b32 v1, v2 // 000000000078: 7E020302 s_branch label_0011 // 00000000007C: BF82FFF1 v_ : 5個 s_ : 4個 */ |
こういうコードでスカラ命令を使うようにすると、
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 |
__kernel void g(__global int *p, int n) { int lid = get_local_id(0); __global int *p2 = p + get_group_id(0) * n; // アドレス演算をifの外に出す for (int i=0; i<n*32; i+=32) { if (lid < 32) { p2[i+lid]++; // lid は最後に足す } } } /* label_000D: v_cmp_gt_i32 vcc, 32, v0 // 000000000034: 7D0800A0 s_and_saveexec_b64 s[8:9], vcc // 000000000038: BE88246A buffer_load_dword v2, v1, s[4:7], 0 offen // 00000000003C: E0301000 80010201 s_waitcnt vmcnt(0) // 000000000044: BF8C0F70 v_add_i32 v2, vcc, 1, v2 // 000000000048: 4A040481 buffer_store_dword v2, v1, s[4:7], 0 offen // 00000000004C: E0701000 80010201 s_mov_b64 exec, s[8:9] // 000000000054: BEFE0408 s_add_u32 s0, s0, 32 // 000000000058: 8000A000 v_add_i32 v1, vcc, 0x00000080, v1 // 00000000005C: 4A0202FF 00000080 s_cmp_le_i32 s3, s0 // 000000000064: BF050003 s_cbranch_scc0 label_000D // 000000000068: BF84FFF2 v_ : 3個 s_ : 6個 */ |
こんな感じですね。まあ今度は s_ のほうがボトルネックになっているので、この場合はあまり褒められたチューニングではないです。
命令数最小にするなら
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 |
__kernel void h(__global int *p, int n) { int lid = get_local_id(0); if (lid < 32) { __global char *p2 = (__global char*)(p + get_group_id(0) * n + lid); // *sizeof(float) を適切な位置に出すためにchar *にする for (int i=0; i<n*32*sizeof(float); i+=32*sizeof(float)) { __global int *p3 = (__global int*)(p2 + i); p3[0]++; } } } . /* label_0011: v_add_i32 v2, vcc, v0, v1 // 000000000044: 4A040300 v_add_i32 v1, vcc, 0x00000080, v1 // 000000000048: 4A0202FF 00000080 v_cmp_gt_u32 s[10:11], s3, v1 // 000000000050: D188000A 00020203 buffer_load_dword v3, v2, s[4:7], 0 offen // 000000000058: E0301000 80010302 s_waitcnt vmcnt(0) // 000000000060: BF8C0F70 v_add_i32 v3, vcc, 1, v3 // 000000000064: 4A060681 buffer_store_dword v3, v2, s[4:7], 0 offen // 000000000068: E0701000 80010302 s_and_b64 exec, exec, s[10:11] // 000000000070: 87FE0A7E s_cbranch_execnz label_0011 // 000000000074: BF89FFF3 v_ : 4個 s_ : 3個 */ |
こんな感じでしょうか。あまり考えずに色々試して出ただけなので一般化は難しいですが…。
SASSのチューニング
SASS は、アドレスのオペランドが定数オフセットしか使えないようなので、x86やGCN isaと比べるとアドレス演算の問題が表面化しやすいように見えます。
|
1 2 3 4 5 6 7 8 9 |
__global__ void f2(int *p, int n) { __shared__ int a[32]; __shared__ int b[32]; for (int i=0; i<n; i++) { p[i] = a[i] + b[i]; } } |
こういうコードは、
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 |
.L_2: /*0088*/ LDS.64 R4, [R7]; /*0090*/ IADD R0, R0, 0x2; /*0098*/ LDS.64 R8, [R10]; /*00a0*/ IADD R7, R7, 0x8; /*00a8*/ IADD R4, R4, R8; /*00b0*/ IADD R5, R5, R9; /*00b8*/ ST.E [R2], R4; /*00c8*/ ST.E [R2+0x4], R5; /*00d0*/ ISETP.LT.AND P0, PT, R0, R6, PT; /*00d8*/ IADD R10, R10, 0x8; /*00e0*/ IADD R2.CC, R2, 0x8; /*00e8*/ IADD.X R3, R3, RZ; /*00f0*/ @P0 BRA `(.L_2); |
こういう感じに、使っているアドレスごとにIADDを使うループになります。
これは結構どうしようもない問題で、ポインタの数だけ加算が出てしまうのを回避できないのですが、
__shared__ 変数は、一個の配列にすることで、ポインタの数を減らすことができます。
上のプログラムなら、
|
1 2 3 4 5 6 7 8 9 10 |
__global__ void f2(int *p, int n) { __shared__ int ab[64]; int *a = ab; int *b = ab+32; for (int i=0; i<n; i++) { p[i] = a[i] + b[i]; } } |
|
1 2 3 4 5 6 7 8 9 10 11 12 13 |
.L_2: /*0088*/ LDS.64 R2, [R7+0x80]; /*0090*/ IADD R0, R0, 0x2; /*0098*/ LDS.64 R4, [R7]; /*00a0*/ IADD R2, R2, R4; /*00a8*/ IADD R3, R3, R5; /*00b0*/ ST.E [R8], R2; /*00b8*/ ISETP.LT.AND P0, PT, R0, R6, PT; /*00c8*/ ST.E [R8+0x4], R3; /*00d0*/ IADD R7, R7, 0x8; /*00d8*/ IADD R8.CC, R8, 0x8; /*00e0*/ IADD.X R9, R9, RZ; /*00e8*/ @P0 BRA `(.L_2); |
IADD を一個削減できています
About Author
