rust-to-spir.pdf

Transcript

1. Rust to SPIR Rustで快適
 OpenCLプログラミング
 （したい） @iitaku
2. まずは 祝：Rust 1.0 !!!
3. About Me twitter id: @iitaku マルチコアCPU/GPU/FPGAさわってます 最近はストレージファームウェアとか
4. めでたいので、 RustでGPU/MIC/FPGAプログラミングしたい
5. OpenCLバインディングか？いやいや
6. OpenCLカーネル中でRustを使いたい！ __kernel void saxpy(const float alpha, __global const float *x, __global const float *y, __global float *z) { uint i = get_global_id(0); z[i] = alpha * x[i] + y[i]; } #[spir_kernel] fn saxpy<T: Float>(alpha: T, xp: &[T], y: &[T], z: &[T]) { let i = get_global_id(0); pz[i] = alpha * x[i] + y[i]; } Old-School OpenCL C Modern OpenCL Rust
7. The SPIR TM Specification Standard Portable Intermediate Representation Version 1.2 Khronos Group主導で規格化された LLVM IRベースの中間言語仕様 
 OpenCL1.2対応プラットフォームなら、
 SPIRからターゲットへコンパイル・実行可能！
8. Rust Lang Rust AST Rust AST Metadata Rust AST LLVM IR ASM/OBJ Phase1: パース・構文木生成 Phase2: マクロ展開・ユーザー構文拡張 Phase3: 各種静的解析 Phase4: LLVM IR生成 Phase5: リンク・オブジェクト生成 これを
9. Rust Lang Rust AST Rust AST Metadata Rust AST LLVM IR ASM/OBJ SPIR 1.2/2.0 既存のLLVM IR生成系を拡張してSPIR生成 アトリビュートを解析して構文木を拡張
 ビルトイン関数を提供 SPIR生成に必要な メタデータを取得 こうしたい
10. 状況 SPIR向けLLVMターゲット追加 構文木拡張 メタデータ抽出 SPIR生成 まだまともなSPIRでない・・・ ← いまここ
11. 所感 Rustはランタイム層薄いので
 新アーキ追加が楽っぽい python on GPUをやった時は かなり制約が必要だった はやくRustでOpenCLカーネル書きたい！
12. Related Works RustGPU: RustをNVPTXに変換する https://github.com/eholk/RustGPU PyNVVM: PythonをNVVMに変換してGPUで動かす https://github.com/iitaku/pynvvm
13. Congrats ! @iitaku