AMD GCN isa と NVIDA SASS でのアドレス演算のチューニング

このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。

2015年6月23日 nakamura

ちょっとしたきっかけがあって、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と違ってスーパースカラの幅が狭いのでアドレス演算がパイプラインを埋めてしまう

たとえば、次のプログラムを

__global__ void f(int *p,int n)
{
    int i;
    for (i=0; i<n; i++) {
        p[i*2] = 100;
    }
}

nvcc + sm_30 でコンパイルするとループ部分は ↓ のようなSASSが出ます。

.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のコンパイラも、ポインタに関しては、かなりナイーブなコードを出すようです。最適化しやすい形に変形してあげる必要があるでしょう。

上の場合なら、

__global__ void g(int *p, int n)
{
    int i;
    for (i=0; i<n; i++) {
        p[0] = 100;
        p+=2;
    }
}

こんな感じですね。

.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で割り当てるのはあまり良くないということを示しています。

例えば、

__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の計算が入ってしまいます。もう少し大きい粒度でスレッドを割り当てて、

__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 からのロードで、アドレスがスカラ値
スカラ演算の結果

スカラ値ではない値で分岐した中で決まった値もスカラ値にはなりません。

以上を踏まえて、

__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個
*/

こういうコードでスカラ命令を使うようにすると、

__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_ のほうがボトルネックになっているので、この場合はあまり褒められたチューニングではないです。

命令数最小にするなら

 __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と比べるとアドレス演算の問題が表面化しやすいように見えます。

__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];
    }
}

こういうコードは、

.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__ 変数は、一個の配列にすることで、ポインタの数を減らすことができます。

上のプログラムなら、

__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];
    }
}

.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

nakamura

Favorite Post

「OpenFOAMスレッド並列化のための基礎検討」を投稿＆発表してきました
2018年2月6日
Livox Mid-360をROS1/ROS2で動かしてみた
2023年1月27日
FFmpeg API の使い方(1): デコードしてみる
2017年8月22日

AMD GCN isa と NVIDA SASS でのアドレス演算のチューニング

AMD GCN isa と NVIDIA SASS

SASS、GCN isa の確認方法

共通のチューニング

GCN isa のチューニング

SASSのチューニング

Tags

About Author

nakamura

Leave a Comment コメントをキャンセル

Tags

Favorite Post

Archives

Categories

keisuke.kimura in Livox Mid-360をROS1/ROS2で動かしてみた

Miya in ウエハースケールエンジン向けSimulated Annealingを複数タイルによる並列化で実装しました

Deivaprakash in Livox Mid-360をROS1/ROS2で動かしてみた

岩崎システム設計　岩崎満 in Alveo U50で10G Ethernetを試してみる

Prabuddhi Wariyapperuma in Livox Mid-360をROS1/ROS2で動かしてみた

Social Media

AMD GCN isa と NVIDA SASS でのアドレス演算のチューニング

AMD GCN isa と NVIDIA SASS

SASS、GCN isa の確認方法

共通のチューニング

GCN isa のチューニング

SASSのチューニング

Tags

About Author

nakamura

Leave a Comment コメントをキャンセル

Tags

Favorite Post

Archives

Categories

keisuke.kimura in Livox Mid-360をROS1/ROS2で動かしてみた

Miya in ウエハースケールエンジン向けSimulated Annealingを複数タイルによる並列化で実装しました

Deivaprakash in Livox Mid-360をROS1/ROS2で動かしてみた

岩崎システム設計 岩崎 満 in Alveo U50で10G Ethernetを試してみる

Prabuddhi Wariyapperuma in Livox Mid-360をROS1/ROS2で動かしてみた

Social Media

岩崎システム設計　岩崎満 in Alveo U50で10G Ethernetを試してみる