先行研究
ガチ勢しかいないが、私は今日Hello Worldだけでここに来た。 #rust_jp— 錆びありはぐれベアメタル (@LDScell) November 26, 2018
Accel: GPGPU Framework for Rust
-
- RustでCUDAカーネルを書く
-
- Accel: GPGPU framework for Rust
-
- 2017末に開始~2,3ヶ月で停止
nvptx64-nvidia-cuda ターゲットを追加するためにrustcにパッチ当てたりしないといけなくてつらかった
proc-macroさえ安定化してなかった…
rust-cuda WGが2019/1くらいに発足
libcoreまで公式のrustupで配布されてる
stdsimdの成果として core::arch::nvptx にLLVMバックエンドのintrinsicsがある
2020/1~再開
GitLabに移行 https://gitlab.com/termoshtt/accel
0.3.0開発中…
CUDAのprintf / PTX system call
#include <stdio.h>
__global__ void test() {
printf("Hello world from %d of %d\n", threadIdx.x, blockDim.x);
}
この printf はCPUの方に命令を投げないといけないので、CUDAのコンパイル時にはPTX system callにある vprintf の呼び出しに変換されて、実行時にはドライバによって管理されます。システムコールには以下の４つがあり、それぞれ core::arch::nvptx に対応する命令がある
-
- vprintf
-
- malloc
-
- free
- __assert_fail
print! macro
-
- Rustのprint!などは内部でメモリ確保を行うのでlibcoreにはない
-
- 最近stdからメモリ確保だけを必要とする部分を分離したalloc crateというのがある
Global Allocatorさえあれば動く
malloc/freeがPTX system callにあるので動く
format!マクロはあるので文字列には出来る
しかし標準出力はstdにしか無い
なので vprintf system callに文字列を上げる
#[macro_export]
macro_rules! print {
($($arg:tt)*) => {
let msg = ::alloc::format!($($arg)*);
unsafe {
::core::arch::nvptx::vprintf(msg.as_ptr(), ::core::ptr::null_mut());
}
}
}
Hello World!
use accel::*;
use accel_derive::kernel;
#[kernel]
pub fn print() {
let i = accel_core::index();
accel_core::println!("Hello from {}", i);
}
fn main() -> anyhow::Result<()> {
let grid = Grid::x(1);
let block = Block::x(4);
let device = driver::Device::nth(0)?;
let _ctx = device.create_context_auto()?;
print(grid, block)?;
Ok(())
}
