Qiita Teams that are logged in
You are not logged in to any team

Log in to Qiita Team
Community
OrganizationAdvent CalendarQiitadon (β)
Service
Qiita JobsQiita ZineQiita Blog
22
Help us understand the problem. What is going on with this article?
@termoshtt

GPUからHello World!

More than 1 year has passed since last update.

GPUからHello World!

by termoshtt
1 / 6

先行研究


Accel: GPGPU Framework for Rust


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 の呼び出しに変換されて、実行時にはドライバによって管理されます。システムコールには以下の4つがあり、それぞれ 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(())
}
22
Help us understand the problem. What is going on with this article?
Why not register and get more from Qiita?
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
ricos
FEMによる構造解析、機械学習の専門家集団。計算資源のクラウド提供もしています。

Comments

No comments
Sign up for free and join this conversation.
Sign Up
If you already have a Qiita account Login
22
Help us understand the problem. What is going on with this article?