nVidiaのGPUのアセンブリ言語を読む

msyksphinz.hatenablog.com

前回は、GPUのアセンブリ言語を出力する方法を勉強した。さて、これを読み解いてみよう。

オリジナルのソースコードは以下だ。

__kernel void vecAdd(__global int *a, __global int *b, __global int *c)
{
        int gid = get_global_id(0);

        c[gid] = a[gid] + b[gid];
}

コンパイル言語は以下だ。

.entry vecAdd(
        .param .u32 .ptr .global .align 4 vecAdd_param_0,
        .param .u32 .ptr .global .align 4 vecAdd_param_1,
        .param .u32 .ptr .global .align 4 vecAdd_param_2
)
{
        .reg .b32       %r<17>;


        ld.param.u32    %r1, [vecAdd_param_0];
        ld.param.u32    %r2, [vecAdd_param_1];
        ld.param.u32    %r3, [vecAdd_param_2];
        mov.b32 %r4, %envreg3;
        mov.u32         %r5, %ntid.x;
        mov.u32         %r6, %ctaid.x;
        mad.lo.s32      %r7, %r6, %r5, %r4;
        mov.u32         %r8, %tid.x;
        add.s32         %r9, %r7, %r8;
        shl.b32         %r10, %r9, 2;
        add.s32         %r11, %r1, %r10;
        ld.global.u32   %r12, [%r11];
        add.s32         %r13, %r2, %r10;
        ld.global.u32   %r14, [%r13];
        add.s32         %r15, %r14, %r12;
        add.s32         %r16, %r3, %r10;
        st.global.u32   [%r16], %r15;
        ret;
}

gidは、get_global_id(0)を使っている。これは、どのようなアセンブリ言語に変換されているのだろうか？ envreg3、ntid、ctaid、tidあたりがキモになりそうだが、よく分からない。 gidは%r10で計算されているが、

%r10 = ((ctaid.x*ntid.x+envreg3) + %tid)

これをベースに、%r1,%r2,%r3をインデックスにしてデータをロードし、加算、結果をストアしている。いろいろググってみたが、実際には上記のパラメータがどのように利用されているのかは良くわからないなあ。。

ちなみに、get_global_id(1), get_global_id(2)に変更してみると、以下のようになる。

        mov.b32 %r10, %envreg4;
        mov.u32         %r11, %ntid.y;
        mov.u32         %r12, %ctaid.y;
        mad.lo.s32      %r13, %r12, %r11, %r10;
        mov.u32         %r14, %tid.y;
        add.s32         %r15, %r13, %r14;
        mov.b32 %r16, %envreg5;
        mov.u32         %r17, %ntid.z;
        mov.u32         %r18, %ctaid.z;
        mad.lo.s32      %r19, %r18, %r17, %r16;
        mov.u32         %r20, %tid.z;