0
0

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

@hiyokomochi

独自GPU言語とRustライブラリでニューラルネットワークの世界に可搬性を！: CommonCL/Rusty Lanternのalpha版を公開しました

0
Last updated at Posted at 2026-05-10

Rust で GPU 計算と機械学習を扱うための実験的なリポジトリを公開しました。

このリポジトリは、大きく分けると CommonCL と Rusty Lantern の2つで構成されています。

CommonCL は、自作の GPU kernel 用言語で書いたコードを CUDA / ROCm / Metal / WGSL / Vulkan / OneAPI / CPU 向けの native source に変換し、同じ runtime から実行するための基盤です。

Rusty Lantern は、CommonCLを使用した Rust 製の機械学習ライブラリです。この記事ではCommonCL/Rusty Lantern の概要と sample_llm_fp16 で Qwen2.5 のサンプルを動かすところまでを紹介します。

この記事で分かること

  • CommonCL と Rusty Lantern の目的・設計の概要
  • CommonCL が移植性と実行速度を両立する仕組み
  • Rusty Lantern ベースのLLM推論デモ（Qwen2.5）を動かす手順

rust、kernel、thread、fp16 などの言葉を見たことがあると読みやすいと思いますが、コマンドのコピペだけで体験できる構成です。

今回のゴールは、手元のマシンで Qwen2.5 0.5B に こんにちは を投げて、generated_text= まで表示することです。実行にはrustの環境が必要です。

CommonCLとRusty Lantern

何を目指しているのか

CommonCL/Rusty Lantern で扱いたいのは、Python系の既存スタックでGPU 計算や ML 実装を扱うときに出てくる次のような問題です。

  1. 移植性が低い
  • 別マシンで動かない
  • 去年のコードが動かない
  • Dockerを持ち出すとマシになるが本質ではない手間と学習が増える
  1. 独自カーネルを書きにくい
  • 独自カーネルを書くとその言語が使えるGPUにロックイン
  • かといってOpenCLを書くとなんだか遅い
  • Tritonとかはあるが対応backendが少ない

これらはだいたい以下が原因です。

  • Pythonのゆるふわさ
  • GPU関連のハードウェア依存
  • GPU向け汎用規格の最適化が往々にして劣る

なのでこうします。

  • Rustでしっかりする
  • GPU関連のハードウェア依存を隠蔽する
  • vendor 独自規格の上にのせて高速に実行する

CommonCL/Rusty Lanternは、この構成を使用して機械学習スタックを1から構築する アホの プロジェクトです。

大雑把に言うと対応関係はこうなっています。

  • CommonCL: CUDA、OpenCLなどのGPU計算APIに相当
  • Rusty Lantern: Pytorchなどの機械学習ライブラリに相当

現行のPython系エコシステムは、現実的により重視すべき他のメリットを優先した結果であり、確かに王道の正解です。しかし、その反対側からのアプローチが可能な選択肢が存在することは、私たちが直面する開発課題に対して、少なくとも一定の意味があるはずです。

実装

CommonCL + Rusty Lanternの呼び出し関係を整理すると以下の図のようになります。

CommonCLは、CommonCL言語で書かれたカーネルを7種のbackend（CUDA / ROCm / Metal / WGSL / Vulkan / OneAPI / CPU）で実行するための基盤です。

Rusty Lanternは、CommonCL をbackendにした Rust 製の ML ライブラリです。拡張性を重視していてCommonCL言語で書いた独自カーネルを使用できます。

まだ alpha 段階なので、API も実装も変わりうる点に注意してください。

Qwen2.5 0.5Bを動かす

ここからはコピペで体感するCommonCL/Rusty Lanternのコーナーです。

Rusty Lanternを使用して実装されたLLM推論サンプルのsample_llm_fp16 を動かします。モデルはちょっと古いですが Qwen2.5 系モデルで、実行はFP16（半精度）です。

対応しているモデルサイズは次の4つです。

  • 0.5b
  • 1.5b
  • 3.0b
  • 7.0b

最初は 0.5b が無難です。モデルは量子化なしのFP16なので、パラメータあたり2 Byte消費し、VRAMはだいたい（モデルサイズ × 2 + 1GB）くらいあれば実行可能です。
必要なものは次のとおりです。

  • Rust / Cargo
  • git
  • 初回モデル取得用のネットワーク
  • 使いたい backend の runtime / driver / device

CommonCL/Rusty Lanternは移植性重視なので非常にシンプルです。

  1. まずclone します。
git clone https://github.com/MochiPiyo/commoncl_and_lantern.git
cd commoncl_and_lantern
  1. cloneしたら--backendを指定して実行します。以上！
cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 0.5b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend metal

sample_llm_fp16 では、backend を次のように指定します。

環境 指定
Apple Silicon / macOS --backend metal
NVIDIA GPU --backend cuda
AMD GPU / ROCm --backend rocm
WebGPU 系 --backend wgsl
Vulkan --backend vulkan
OneAPI --backend oneapi
CPUで小さく確認 --backend cpu

LinuxやWindowsの場合、CUDA / ROCm が使えるなら、まずはそこから試すのが高速です。追加の GPU 開発環境を何も入れていない場合の動作確認には wgsl が試しやすいです。wgsl は WebGPU 系の backend で、グラフィック向け API を使うため、GUI のある OS なら動く可能性が高いです。WGSL 自体は Web ブラウザ向けの shader 言語ですが、CommonCL では Firefox などでも使われている wgpu を内蔵しているため、ブラウザ外でも使えます。

macOSならMetalが最初から使えるはずです。

使えるbackendがわからないときはこのコマンドで確認してください。このコマンドは、CommonCL Runtime から見える backend や device 情報を表示しています。

cargo run -p commoncl_runtime -- caps

実行例

Runtime Capability Check
----------------------------------------

Backends
  backend  runtime     reason
  --------------------------------------------------------
  cpu      available   cpu is always available
  wgsl     available   wgpu found a compatible adapter through Metal
  metal    available   Metal device creation succeeded
  cuda     unavailable runtime_unavailable: nvcc is not available in PATH
  rocm     unavailable runtime_unavailable: hipcc is not available in PATH
  vulkan   unavailable runtime_unavailable: glslangValidator is not available in PATH
  oneapi   unavailable runtime_unavailable: oneapi compiler (dpcpp/icpx) is not available in PATH

...（略）...

Apple Silicon / macOS なら、次のコマンドで試せます。

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 0.5b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend metal

CUDA なら --backend cuda にします。

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 0.5b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend cuda

ROCm なら --backend rocm です。

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 0.5b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend rocm

初回実行時に必要なモデルファイルが無ければ、Hugging Face Hub から取得します。保存先は sample_llm_fp16/model/Qwen_Qwen2.5-0.5B-Instruct/ のようなモデル別ディレクトリです。モデル取得と初回 compile があるので、最初だけ少し時間がかかります。

成功すると、設定値、生成 token id、生成テキストが表示されます。当方の環境では生成文はこんにちは！どのようなお手伝いが必要（生成上限の8トークンで停止）でした。

backend=Metal
model=Qwen/Qwen2.5-0.5B-Instruct
max_new_tokens=8
prompt=こんにちは
generated_token_ids=[151644, 8948, ...(略)...]
generated_text=こんにちは！どのようなお手伝いが必要

まずは generated_text=... まで出れば成功です。生成内容は sampling 設定や backend によって変わることがあります。

0.5Bの次

0.5B が動いたら、--model を変えるだけで大きいモデルも試せます。また、promptやmax-new-tokensを変更することで色々遊べます。

backend指定を変えてみるのも良いでしょう。cpu, wgslは大体の環境で動作するはずです。oneapiやvulkanはコンパイルができる環境が必要なので現状では少し面倒です。

1.5Bで長い生成をする例

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 1.5b \
  --prompt "GPUについて説明してください" \
  --max-new-tokens 64 \
  --backend metal

このプロンプトではこういう結果が得られました。

generated_text=GPU（Graphics Processing Unit）とは、グラフィックス処理ユニットのことを指します。主にゲーム機や高機能なパソコンで使用されるハードウェアアクセサリーです。

GPUは、通常、PC（Personal Computer）やゲーム機、デジタルレンダリングなどの高

3B / 7B も指定できます。

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 3.0b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend metal

cargo run --release --manifest-path sample_llm_fp16/Cargo.toml -- \
  --model 7.0b \
  --prompt "こんにちは" \
  --max-new-tokens 8 \
  --backend metal

中では何が起きているのか

ここまでのサンプルでは、Rust 側から Tensor や LLM の処理を呼び出しています。その下では、CommonCL Runtime が backend を選び、必要な kernel や native library 呼び出しを実行します。

CommonCL の特徴は、CCL で書いた kernel を backend ごとの通常の GPU source に変換するところです。独自 VM で解釈しているわけではありません。

たとえば、単純なベクトル加算は CCL言語 だとこう書きます。ここでは説明を短くするため、dispatch 範囲と配列長の調整などは省いています。

@versions(commoncl = "0.1.0-alpha.1")

@group(0) @binding(0)
var<storage> lhs: Array<f32>;
@group(0) @binding(1)
var<storage> rhs: Array<f32>;
@group(0) @binding(2)
var<storage> mut out: Array<f32>;

@workgroup_size(64, 1, 1)
fn main() {
    let i: u32 = workgroup_size.x * workgroup_id.x + thread_id.x;
    out[i] = lhs[i] + rhs[i];
}

lhs と rhs が入力、out が出力です。workgroup_id と thread_id から、各 thread が担当する配列 index を作っています。

このCCL言語をCommonCL Compilerに通してCUDAカーネルに変換すると、およそ次のような kernel になります。

extern "C" __global__ void vector_add_kernel(
    const float* __restrict__ lhs,
    const float* __restrict__ rhs,
    float* __restrict__ out
) {
    uint32_t i = blockDim.x * blockIdx.x + threadIdx.x;
    out[i] = lhs[i] + rhs[i];
}

CCL の workgroup_id / thread_id が、CUDA の blockIdx / threadIdx に対応しています。Metal 向けなら Metal の kernel、WGSL 向けなら WGSL の shader になります。

おまけ1: 量子化Llama

⚠️まだ実験段階なので遅いし、動作するのはたぶん metal / cuda / rocm くらいのbackendだけです。

sample_llama_quant では、GGUF 形式の量子化 Llama を読み込んで推論できます。

デフォルトでは unsloth/Llama-3.2-1B-Instruct-GGUF の Q4_K_M モデルを使います。--gguf を指定しない場合は、必要な .gguf ファイルと tokenizer を Hugging Face Hub から取得し、sample_llama_quant/model/ 以下に保存します。

量子化モデルなので fp16 のモデルよりメモリ使用量は小さく、手元では 2GB 程度のメモリがあれば動く想定です。ただし、まだ最適化は十分ではないため、速度を見るサンプルというより「GGUF を読んで CommonCL backend 上で Llama 推論まで通る」ことを確認するためのものです。

cargo run --release -p sample_llama_quant -- \
  --prompt "hello world" \
  --max-new-tokens 8 \
  --backend metal

おまけ2: SDXL

⚠️まだ実験段階なので遅いし、動作するのはたぶん metal / cuda / rocm くらいのbackendだけです。

sample_stable_diffusion では SDXL の画像生成も試せます。

ROCm 環境で 384px 画像を生成する例です。VRAMは 12GB 以上が必要です。

cargo run --release --manifest-path sample_stable_diffusion/Cargo.toml -- \
  --backend rocm \
  --steps 20 \
  --width 384 \
  --height 384 \
  --output sample_stable_diffusion/out_384_blue_sphere.png \
  --prompt "a blue glass sphere on a wooden table, studio lighting, sharp focus" \
  --negative-prompt "red cube, low quality, blurry" \
  --guidance-scale 5.0 \
  --seed 44 \
  --pool-capacity-gb 12

生成例です。一応「青いガラスの球」らしきものは見えます...。

out_384_blue_sphere.png

現時点の制限

現時点では、実際の ML workload を通して設計と実装を検証している段階です。

  • API はまだ変わる可能性があります
  • backend ごとに対応状況や性能差があります
  • fp16、workgroup、barrier などは環境差があります
  • LLM や SDXL の品質・速度はまだ改善中です
  • サンプルによっては初回のモデル取得や kernel compile に時間がかかります

エラーが出る場合

CommonCL runtime/compiler の基本動作は各backendの実機で確認していますが、環境依存のバグはまだ残っていると思われます。
動かしてみてエラーが出た場合は、GitHub Issues で報告していただけるとたいへん助かります。

おわりに

CommonCL と Rusty Lantern は、Rust で GPU 可搬な ML 実行基盤を作るための実験的なプロジェクトです。

まだ alpha 版で、API も性能も実装もこれから変わっていくと思います。それでも、Rust だけで GPU kernel から Tensor、LLM、画像生成までつながる形が見えてきたので、いったん公開することにしました。

手元の環境で試していただいて、動いた backend、動かなかった backend、変なエラーなどがあれば GitHub Issues に投げてもらえると助かります。

機能詳細や設計の話は、別の記事で少しずつ書いていく予定です。

0
0
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
0
0

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?