Clang で CUDA コードを NVPTX に変換するメモ #CUDA

背景

CUDA 開発環境とか入れるのめんどい
- PC 環境変えたら毎回開発環境入れるのめんどい
- 自前アプリを CI ビルドするときとか最小限の構成にしたい.
clang から CUDA コードから直接 PTX(as) 吐きたい
なるべく素の機能を使って, C コードと共存したいとか, あとで OpenCL C に移行しやすくしたいとか.
NVRTC(JIT compile) のテストとしてオフラインコンパイルを確認したいとか

Clang CUDA モード?

最近の nvcc(clang/LLVM ベース)は, その通り基本 clang/LLVM のコードを使っている(clang/LLVM に対応コードがコミットされている)ので, clang/LLVM では CUDA 対応(CUDA 構文パースや PTX コード生成)が施されています.

C++ STL も多少は利用できるようです.

CUDA SDK のヘッダファイルと組み合わせることで, CUDA コードを処理できます.
ただ, そのための環境構築のドキュメントやサンプルはほとんどありません.

CUDA 固有の qualifier について

__device__ などは clang cuda モードでは未定義です.

nvcc でも, これは CUDA SDK の crt/host_defines.h を読んで設定しているようです.
crt/host_defines.h を参考に定義します.
clang の場合は __attribute__ にマップします.

#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
#define __host__ __attribute__((host))
#define __shared__ __attribute__((shared))

サンプル

#define __global__ __attribute__((global))

__global__ void add(float a, float b, float *c)
{
	c[0] = a + b;
}

こんな感じの最小のコードを用意します.

clang++ -S -v --cuda-device-only --cuda-gpu-arch=sm_60 -xcuda -nocudainc -nocudalib test.cu

デフォルトでは, CUDA のヘッダやライブラリを見るようになっていますので無効化します. -nocudainc, -nocudalib
-triple=nvptx64-nvidia-cuda を付与するとより確実かもしれません.

//
// Generated by LLVM NVPTX Back-End
//

.version 5.0
.target sm_60
.address_size 64

        // .globl       _Z3addffPf

.visible .entry _Z3addffPf(
        .param .f32 _Z3addffPf_param_0,
        .param .f32 _Z3addffPf_param_1,
        .param .u64 _Z3addffPf_param_2
)
{
        .local .align 8 .b8     __local_depot0[16];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .f32       %f<6>;
        .reg .b64       %rd<5>;

        mov.u64         %SPL, __local_depot0;
        cvta.local.u64  %SP, %SPL;
        ld.param.u64    %rd1, [_Z3addffPf_param_2];
        ld.param.f32    %f2, [_Z3addffPf_param_1];
        ld.param.f32    %f1, [_Z3addffPf_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        cvta.global.u64         %rd3, %rd2;
        st.f32  [%SP+0], %f1;
        st.f32  [%SP+4], %f2;
        st.u64  [%SP+8], %rd3;
        ld.f32  %f3, [%SP+0];
        ld.f32  %f4, [%SP+4];
        add.rn.f32      %f5, %f3, %f4;
        ld.u64  %rd4, [%SP+8];
        st.f32  [%rd4], %f5;
        ret;

}

Voila!

cubin(PTX binary)?

clang では内部では, CUDA SDK の ptxas で cubin(ptx を ELF バイナリにしたもの?)に変換していました.

 "/usr/local/cuda-10.2/bin/ptxas" -m64 -O0 -v --gpu-name sm_60 --output-file test-cuda-nvptx64-nvidia-cuda-sm_60.o /tmp/test-45280d.s

したがって cubin を作りたい場合には CUDA SDK が必要になります.

もしくは, CUDA Driver API で .ptx 読んで elf(cubin)バイナリ取得できるかもしれません.

CUDA SDK からライブラリ抜き出してくる必要がありますが, クライアントサイドで PTX をコンパイルして elf バイナリ得るライブラリもあります.
(実行時には CUDA SDK(Runtime API) 非依存にできる)

PTX Compiler API のメモ
https://qiita.com/syoyo/items/cfaf0f7dd20b67cc734e

Triple bracket(kernel lauch 構文)

一応 clang 単体(cuda モード)でいけるようです.

T.b.W.

TODO

~~[ ] clang(LLVM) 自体で cubin(ELF?)吐けるかな?~~
- 無理そう.
triple bracket 構文動くか試す.
nanostl https://github.com/lighttransport/nanostl と組み合わせ, いくらか STL 関数を使えるようにしてみる(特に std::vector)