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

5036 ワード
CUDA NVPTX clang 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 コードを処理できます.
ただ, そのための環境構築のドキュメントやサンプルはほとんどありません.

ソースコード記述

__device__ などは未定義です.

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

TODO

  • [ ] clang(LLVM) 自体で cubin(ELF?)吐けるかな?
    • 無理そう.
  • nanostl https://github.com/lighttransport/nanostl と組み合わせ, いくらか STL 関数を使えるようにしてみる(特に std::vector)
Hdu 3603
『Java入門第1四半期』のRandomクラスと乱数取得事例