GPUからHello World!

先行研究

ガチ勢しかいないが、私は今日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(())
}