この記事はGPGPU Advent Calendarの17日目の記事です。ええ、現在時刻は12月17日の27時ですが、何か？

replを作ろう

Read-Eval-Print-Loop、それは闇プログラマーとなりし者に訪れる運命・・・彼は逝ってしまったわ、円環の理に導かれて・・・

ということでそれが世界の選択なのでreplを作ってLLVM PTX Backendで遊んでみましょう。最後まで読めば、この行為が壮大な無駄であることがきっとわかるはず！

おさらい

前回の内容をふりかえると、

LLVM 3.2でLLVM IRからPTXが吐けるようになりました
llc（LLVM System Compiler）を使ってLLVM IRをPTXに変換してみました
PTXに変換できるLLVM IRにはいくつかルールが存在するので、それをざっと見ました

という事でした。

目標

replとか言ったがrpl部分はまあ適当でいいだろ・・・ってことで

S式っぽい何か（整数即値および四則演算 +-*/ のポーランド記法の二項演算）を読んで構文木に変換する
構文木からLLVM IRを生成する
PTXに変換する
CUDAのドライバインタフェースで実行する

という感じを目標にします。llcのソースコードやCUDAのサンプルを組み合わせれば簡単にできそうですが、とはいえ幾つか壁があります。

LLVM IRってどうやって生成すんのよ

なんだかいくつかやり方はあるみたいですが、主な手順は以下のとおりです。

Contextをつくる
Contextを使ってModuleをつくる
Moduleの中にFunctionをつくる
Functionの中にBasicBlockを最低1つはつくる
IRBuilder<>をつくって希望のBasicBlockに向ける
IRBuilderを使って命令列を書き込む

ええ、わかりにくいですね。具体例を。

LLVMContext context;
Module *module = new Module("repl_module", context);

std::vector<Type*> arg_types;
arg_types.push_back(PointerType::get(Type::getInt32Ty(context), 1));

FunctionType *function_type = FunctionType::get( Type::getVoidTy(context), arg_types, false);

Function *function = cast<Function>(module->getOrInsertFunction("repl_function", function_type));
BasicBlock::Create(context, "EntryBlock", function);

Contextを作って、Moduleに名前をつけて、Functionの型をつくって、BasicBlockを作って、ここまでが4の手順です。

さらに、今回はPTX向けのLLVM IRを作らないといけませんから、

kernel関数であることをmetadataを使用して明示する
global memoryのaddrspaceをきちんと指定する

といった手順が必要です。このとおり。

NamedMDNode *annotate = module->getOrInsertNamedMetadata("nvvm.annotations");
std::vector<Value *> vals;
vals.push_back(function);
vals.push_back(MDString::get(context, "kernel"));  vals.push_back(ConstantInt::get(Type::getInt32Ty(context), 1));
annotate->addOperand(MDNode::get(context, vals));

addrspaceは最初の例のarg_typesを作っているところにあります。

arg_types.push_back(PointerType::get(Type::getInt32Ty(context), 1));

PointerType::getのふたつめの引数の1がglobal memoryのaddrspaceです。
余談ですがaddrspaceが定義されているNVPTX.hはLLVMの公開ヘッダでは無いのでこういう時ちょっと微妙ですね・・・

あとは、構文木を手繰って式をつくり、その結果をint *dstの指すポインタにストアする、という命令列を作りましょう。

global::builder = new IRBuilder<>(bb);
AstNode *node = parse(buf);
Argument *dst = function->arg_begin();
Value *value = node->get();
global::builder->CreateStore(value, dst);
global::builder->CreateRetVoid();

AstNodeクラスが構文木を表す基底クラスです。BinOpNodeとValueNodeがこれを継承しています。
AstNodeクラスはllvm::Value get(void)というシグネチャの仮想関数をインタフェースとして定義しており、これを再帰的に呼び出して構文木をLLVM IRに変換します。
続いてIRBBuilder<>::CreateStoreでストア命令をつくり、最後にIRBuilder<>::CreateRetVoidで関数の終了を明示します。

PTX Backendってどう使うのよ

さて、これでLLVM IRはできました。これをPTXに変換するにはどんな手順がいるのでしょうか？

まずはNVPTXターゲットの初期化が必要です。

// Initialize LLVM subsystems
LLVMInitializeNVPTXTargetInfo();
LLVMInitializeNVPTXTarget();
LLVMInitializeNVPTXTargetMC();
LLVMInitializeNVPTXAsmPrinter();

PassRegistry *Registry = PassRegistry::getPassRegistry();
initializeCore(*Registry);
initializeCodeGen(*Registry);
initializeLoopStrengthReducePass(*Registry);
initializeLowerIntrinsicsPass(*Registry);
initializeUnreachableBlockElimPass(*Registry);

LLVMInitializeNVPTX*という関数ですが、ヘッダが公開されていないのでプロトタイプ宣言を適当に宣言して使ってください。llcだとコンパイルしたすべてのターゲットを初期化するのですが、今回はバックエンドはPTXだけでよく、余計なライブラリをリンクするとビルドのたびにえらく時間を食うのでこうなりました。

次に、ターゲットに関する情報を幾つか入力してやります。

Triple triple;
std::string err;
const Target *target = TargetRegistry::lookupTarget("nvptx64", triple, err);

std::auto_ptr<TargetMachine> target_machine(target->createTargetMachine(triple.getTriple(), "sm_20", "", TargetOptions()));

TargetMachineはターゲットに関する情報を保持するクラスです。
TargetRegistry::lookupTargetを使用すると、アーキテクチャ名から適切なTriple（ターゲットの情報を表す3つ組の文字列）を導出できるので、これを使用してTargetMachineを作ります。
また、GPUのCompute Capability "sm_20"等もここで文字列として与えます。

さて、そしていよいよPTXへの変換部分です。

PassManager pm;
pm.add(new TargetLibraryInfo(triple));
pm.add(new TargetTransformInfo(target_machine->getScalarTargetTransformInfo(), target_machine->getVectorTargetTransformInfo()));
pm.add(new DataLayout(*(target_machine->getDataLayout())));

target_machine->setAsmVerbosityDefault(true);

std::string ptxcode;
raw_string_ostream ros(ptxcode);
formatted_raw_ostream fos(ros);
target_machine->addPassesToEmitFile(pm, fos, TargetMachine::CGFT_AssemblyFile);

PassManagerとは何でしょうか？
すこしだけLLVMの内部実装の話になりますが、LLVMのバックエンドはパスという仕組みを利用して実装されています。
パスとは、簡単に言ってしまえばModuleやFunctionを取って、それを改変したり外にダンプしたりする仕組みです。
バックエンドだけでなくLLVM IRの最適化もこのパスを使用して実装されており、llcのコンパイルフェーズは何十から何百という様々なパスの組み合わせから構成されています。

ここではPassManagerに対して、ptxcodeというstd::stringをバッファにもつstreamにアセンブリ形式で結果を出力するというパスを追加しました。replのループ中で、これらのパスを以下のようにして実行します。

// Codegen PTX using NVPTX
pm.run(*module);
fos.flush();

fos.flush()を忘れないでください。raw_string_ostreamは内部バッファを持っており、書き込まれたストリームの大きさが一定以上の大きさになるまで、外部バッファにデータを出力しません。

これでめでたくptxcodeにPTXが得られました。

PTXをどうやってGPUで実行させんのよ

CUDAのドライバインタフェース(cuda*ではなくcu*というインタフェース)を使用すると、動的にPTXをロードしてGPU上で実行することができます。

// Load PTX and Exec

// CUmoduleをつくります
CUmodule cu_module;
cuModuleLoadDataEx(&cu_module, ptxcode.c_str(), 0, 0, 0);

// CUfunctionをつくります
CUfunction cu_function;
cuModuleGetFunction(&cu_function, cu_module, "repl_function");

// Blockのサイズを設定します
cuFuncSetBlockShape(cu_function, 1, 1, 1);

// カーネルの引数を設定します
cuParamSetv(cu_function, 0, &d_dstptr, sizeof(d_dstptr));
cuParamSetSize(cu_function, sizeof(d_dstptr));

// いけ！
cuLaunchGrid(cu_function, 1, 1);

// 結果くれ
cudaMemcpy(&h_dst, d_dstptr, sizeof(int32_t), cudaMemcpyDeviceToHost);

簡単ですね。

結果

$ rlwrap ./gpu-repl
>>> (+ 1 2)
3
>>> (* (+ 1 2) 4)
12
>>>

ふむ、なにやら動いているみたいですね。でもほんとに動いているのか心配なのでPTXをダンプしてみましょう。

>>> (* (+ 1 2) 4)
//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_20, texmode_independent
.address_size 64


    // .globl   repl_function
.entry repl_function(
    .param .u64 .ptr .global .align 4 repl_function_param_0
)                                       // @repl_function
{
    .reg .pred %p<396>;
    .reg .s16 %rc<396>;
    .reg .s16 %rs<396>;
    .reg .s32 %r<396>;
    .reg .s64 %rl<396>;
    .reg .f32 %f<396>;
    .reg .f64 %fl<396>;

// BB#0:                                // %EntryBlock
    mov.u32     %r0, 12;
    ld.param.u64    %rl0, [repl_function_param_0];
    st.global.u32   [%rl0], %r0;
    ret;
}


12
>>>

12をストア・・・だと・・・LLVM IRはどうなのさ

>>> (* (+ 1 2) 4)
; ModuleID = 'repl_module'

define void @repl_function(i32 addrspace(1)*) {
EntryBlock:
  store i32 12, i32 addrspace(1)* %0
  ret void
}

declare void @abort()

!nvvm.annotations = !{!0}

!0 = metadata !{void (i32 addrspace(1)*)* @repl_function, metadata !"kernel", i32 1}
12
>>>

・・・まあ、そーですよねー

今回のreplの入力が受け付けるのは整数定数なので、IRの構築時に全てConstant Foldingされて定数ストアになってしまうのでした。

というわけで！GPUの有り余る計算パワーを使って！定数でメモリを埋めるという新たなプログラムがここに爆誕しました！おめでとう！ありがとう！

サンタさんからのプレゼントだよぉ

ソースコード

それでは皆さん良いクリスマスを。

GPGPU Advent Calendar #17