nVIDIAのOpenCLカーネルをコマンドラインからコンパイルするツールclccを試す

nVIDIAのGPGPUを使ってプログラムを書くために、毎回clCreateProgramWithSourceでコンパイルするのは面倒だ。 それに、逆アセンブルしてどういうカーネルになっているのかを確かめないと納得できない。 それでいろいろ調べていたのだが、nVIDIAのOpenCLカーネルは、オブジェクトファイルを出力できないことが分かってきた。

clCreateProgramWithSource

sygh.hatenadiary.jp

ああーなるほど。OpenCLの対応はnVIDIAのGPGPUではあまり充実していないのか。 その後いろいろ調べると、nVIDIAのOpenCLのカーネルからアセンブリファイルを出力できるclccというプログラムがあることが分かった。

github.com

やっていることは、nVIDIAのOpenCLライブラリを直接叩くことで、プログラムのコンパイルと情報を引き出しているようだ。

...
    strings[0] = source;
    lengths[0] = strlen(source);
    result = NvCliCompileProgram(strings, count, lengths, options, &log, &binary);
...

さっそく試してみよう。環境はVisualStudioとCygwinだ。 まずVisualStudioのソリューションが入っているのでそれを開いて、コンパイルして実行してみる。

$ ./clcc.exe
Usage: clcc ["compiler-options"] input.cl output.ptx

なるほど、コンパイルできている。さて、OpenCLカーネルをコンパイルしている。

• kernel.cl

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

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

まずいろいろ試すと、stat()を使っているものだから、WindowsとUnixの改行コードの種類によってはうまくいかない場合があった。 Linuxでは問題無いが、Cygwinを使っている場合は少し注意しなければならない。ソースコードはLFCRの改行コードを使わなければならない。 さらにコンパイルオプションとして -cl-nv-std=CLX.Xが必要だ。

$ ./clcc.exe -cl-nv-cstd=CL1.1 ../../../add_loop_gpu/add_loop_gpu/Debug/kernel.cl output.ptx
$ less output.ptx

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID:
// Driver
// Based on LLVM 3.4svn
//

.version 4.3
.target sm_35, texmode_independent
.address_size 32

        // .globl       vecAdd

.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;
}

お、アセンブリっぽいものが表示されている！これを使って、いろいろnVIDIAのGPUの勉強をしてみよう。

ちなみに、-cl-nv-cstdをCL1.1/CL1.2/CL2.0に変えても、出力されるログは一緒だった。

$ ./clcc.exe -cl-nv-cstd=CL1.1 ../../../add_loop_gpu/add_loop_gpu/Debug/kernel.cl output.ptx
$ ./clcc.exe -cl-nv-cstd=CL1.2 ../../../add_loop_gpu/add_loop_gpu/Debug/kernel.cl output.ptx
$ ./clcc.exe -cl-nv-cstd=CL2.0 ../../../add_loop_gpu/add_loop_gpu/Debug/kernel.cl output.ptx