Skip to content

Commit 8ebe64f

Browse files
committed
pie 0.4.0: release pipeline, portability fix, and 0.4.0 docs
Per-compute-capability CUDA release matrix plus portable Vulkan / macOS Metal / Windows binaries: node24 CI, manylinux_2_28 (glibc 2.28), self-hosted build removed, standardized artifact names. Rustls/OpenSSL portability fix so the release binaries no longer dynamically link libssl.so.1.1. Per-CC auto-detecting installer (driver + compute capability). Full 0.4.0 website/docs refresh: multimodal guide + a Forward Pass > Media page, drivers (TensorRT-LLM surfaced in sidebar + card; the removed `dev` driver dropped), models (GLM-5.1 / per-CC CUDA), multimodal in the SDK reference, concise CHANGELOG, and the v0.5.0 roadmap. Validated end-to-end on an L40S fleet (cuda12.8 and cuda13.0): documented install, inferlet generation, and perf vs vLLM.
1 parent 0d11b4c commit 8ebe64f

24 files changed

Lines changed: 412 additions & 432 deletions

File tree

.github/workflows/build.yml

Lines changed: 163 additions & 259 deletions
Large diffs are not rendered by default.

.github/workflows/ci.yml

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ jobs:
1919
uses: dtolnay/rust-toolchain@stable
2020

2121
- name: Cache cargo registry + target
22-
uses: actions/cache@v4
22+
uses: actions/cache@v5
2323
with:
2424
path: |
2525
~/.cargo/registry
@@ -59,7 +59,7 @@ jobs:
5959
uses: dtolnay/rust-toolchain@stable
6060

6161
- name: Cache cargo + cmake build
62-
uses: actions/cache@v4
62+
uses: actions/cache@v5
6363
with:
6464
path: |
6565
~/.cargo/registry
@@ -100,7 +100,7 @@ jobs:
100100
uses: dtolnay/rust-toolchain@stable
101101

102102
- name: Cache cargo
103-
uses: actions/cache@v4
103+
uses: actions/cache@v5
104104
with:
105105
path: |
106106
~/.cargo/registry
@@ -137,7 +137,7 @@ jobs:
137137
uses: dtolnay/rust-toolchain@stable
138138

139139
- name: Cache cargo + cmake build
140-
uses: actions/cache@v4
140+
uses: actions/cache@v5
141141
with:
142142
path: |
143143
~/.cargo/registry

CHANGELOG.md

Lines changed: 8 additions & 111 deletions
Original file line numberDiff line numberDiff line change
@@ -6,116 +6,13 @@ adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).
66

77
## [0.4.0] - 2026-06-15
88

9-
A large release centered on **multimodal serving**, **native speculative
10-
decoding**, **weight/KV quantization**, a rewritten **scheduler and weight
11-
loader**, and a broad set of **new model architectures**.
12-
13-
### Added
14-
15-
#### Multimodal (vision, video, audio)
16-
- Model-agnostic multimodal API for **vision, video, and audio** — both input
17-
and output. Inferlets hand the host raw encoded bytes; the host owns
18-
decoding and preprocessing, so no model-specific constants leak into inferlet
19-
code.
20-
- `Context` bindings: `append_image`, `append_audio`, `append_video`, plus
21-
`Image`/`Audio`/`Video::from_bytes` with metadata accessors
22-
(`token_count()`, `grid()`).
23-
- Text-to-speech output: `model.speak(text).speaker(id).generate()` returning
24-
self-describing speech (`to_wav()` with `sample_rate`/`channels`).
25-
- `inferlet::http::fetch` for pulling remote media inside an inferlet, and
26-
simplified `media` bindings across the Rust/Python/JS SDKs.
27-
- Example inferlets: `image-qa`, `audio-qa`, `video-qa`, `tts`
28-
(+ `image-qa-bench`), and a multimodal image Pie-vs-vLLM benchmark.
29-
30-
#### Speculative decoding
31-
- **Native Multi-Token-Prediction (MTP) speculation** — lossless, in-driver
32-
draft generation for hybrid GDN models. Opt-in via
33-
`enable_system_speculation` with `mtp_num_drafts` and an optional
34-
`mtp_assistant_snapshot_dir`.
35-
- MTP system speculators wired for **Gemma-4**, **Qwen3.5**, and
36-
**Qwen3.5-MoE**, with fused GEMV+argmax, graph argmax, and per-step position
37-
handling on the draft path.
38-
39-
#### Quantization
40-
- Runtime weight quantization: **fp8, int8, fp4/mxfp4**, with a native MXFP4
41-
loader repack path.
42-
- **KV-cache quantization** formats (fp8_e4m3 / fp8_e5m2 / … / nvfp4),
43-
selectable via `kv_cache_dtype`.
44-
- Dense pre-quantized **block-FP8** weights (e.g. Qwen3-FP8) in the loader.
45-
46-
#### New model architectures (CUDA native)
47-
- **GLM-5.1** (MLA on Blackwell, RoPE, FP4 quant, DSA indexer).
48-
- **DeepSeek V4** and **Kimi** (MLA, fused q_a/kv_a and gate+up projections).
49-
- **Nemotron-H**.
50-
- **Qwen3-MoE** (Qwen3-30B-A3B), **Qwen3.5 / Qwen3.6** (incl. MTP drafters),
51-
and **Qwen3-VL** vision.
52-
- **Gemma-4** (vision + audio in) and **CSM** (audio out).
53-
54-
#### Drivers & runtime
55-
- New **TensorRT-LLM driver** backend.
56-
- `runtime::launch` API for inferlet-to-inferlet invocation.
57-
- Runtime-managed `rs_cache` support and a low-latency IPC profile with a
58-
polling channel.
59-
- Automatic driver **capacity planning** (SM-aware decode capacity, throughput
60-
arena selection, auto admission).
61-
- Portable driver: **sharded GGUF** support, MoE experts from stacked GGUF
62-
tensors, single-`.gguf` `hf_repo` targets, and a tokenizer minted from GGUF
63-
KV metadata.
64-
- Startup banner for `pie serve`.
65-
66-
#### Examples & SDK
67-
- New example inferlets: **AsyncLM** (async function calling), **Reflexion**
68-
reasoning, hierarchical attention + MCTS, modular cache (Rust/Python/JS),
69-
self-correct, and custom-speculator demos.
70-
- E2E tests for raw-completion inferlets and a per-arch SHA smoke gate
71-
(`benches/smoke_deterministic`).
72-
73-
### Changed
74-
75-
#### Performance — scheduler & runtime rewrite
76-
- Main scheduler loop now runs on a **dedicated synchronous OS thread** with
77-
crossbeam channels, firing batches synchronously (closes the batch-coalescing
78-
gap that caused singleton-decode fires under load).
79-
- **Chain-extender worker pool** replaces per-context tasks; pure-decode fast
80-
path extended to chain continuations.
81-
- Switched the global allocator to **mimalloc**; capped tokio worker threads;
82-
`dashmap` for the staged batch; deferred dispatch drops to shrink the
83-
inter-fire gap.
84-
- **Chunked prefill** in the scheduler; numerous CUDA attention, TP1/TP2
85-
throughput, and GDN/FLA kernel-fusion optimizations.
86-
- Result: pie now matches or exceeds vLLM throughput on the measured Qwen3
87-
configurations.
88-
89-
#### Weight loader
90-
- **Rust is now the canonical weight loader** — an algebraic, profile-driven
91-
load-plan pipeline (physical byte-plan layer, streamed checkpoint tensors,
92-
metadata-backed FP8 decode) replacing the legacy C++ storage compiler.
93-
94-
#### Build & packaging
95-
- **Vendored FlashInfer CUTLASS MoE launchers**, removing the build-time Python
96-
codegen and its incidental `tvm_ffi` / `pynvml` / `tqdm` dependencies.
97-
- **CUDA 12.8 / 13 dual-build** support and `sm120` builds. **CUDA 12.8 is the
98-
minimum toolkit** — FlashInfer's Hopper (Mamba SSU) and FP4 kernels require 12.8+.
99-
- Slim CUDA Docker image for `pie-server` (driver-cuda).
100-
- Version bumped to **0.4.0** across all crates, packages, and install scripts.
101-
102-
### Removed
103-
- The **`dev` reference driver** (use `cuda_native` or the portable driver).
104-
- The C++ storage compiler / layout planner and C++ storage-compiler
105-
compatibility shims (superseded by the Rust loader).
106-
107-
### Fixed
108-
- Graceful error instead of a panic on compact wasm import sections.
109-
- Windows portability: `getpid()` in the portable driver entry; explicit
110-
unsupported-driver errors for dev/vllm/sglang on Windows.
111-
- Graph-cache invalidation on page-index growth; lazy-delete of stale
112-
restore-queue entries.
113-
- Paged-attention decode planner `padded_batch_size` / tile-count mismatch.
114-
- XQA `gqa=5` graph-capture failure under TP>1.
115-
- GGUF portable driver build failure; first-shard validation in
116-
`open_sharded`.
117-
- Serialized forked-branch generation; CUDA quantized generation paths.
118-
- GPU-local CPU-affinity parsing in the benchmarks (NIC columns in the
119-
`nvidia-smi` topology), now also applied to `pie_bench`.
9+
- **Multimodal**: vision/video/audio input and audio output via a model-agnostic media API (`append_image`/`append_audio`/`append_video`; TTS via `model.speak`).
10+
- **Native MTP speculative decoding**: lossless, in-driver draft generation (Gemma-4, Qwen3.5, Qwen3.5-MoE).
11+
- **Quantization**: runtime weight quant (fp8/int8/fp4/mxfp4) and KV-cache quant (`kv_cache_dtype`).
12+
- **New model architectures**: GLM-5.1, Nemotron-H, Kimi/DeepSeek, Qwen3-MoE, Qwen3.5/3.6, Qwen3-VL, Gemma-4, CSM.
13+
- **TensorRT-LLM driver**.
14+
- **Scheduler & weight-loader rewrite**: pie now matches or exceeds vLLM throughput on Qwen3; Rust is the canonical weight loader.
15+
- **Removed**: the `dev` reference driver (use `cuda_native` or the portable driver).
16+
- **Build**: CUDA 12.8 is the minimum toolkit; CUDA 12.8/13 dual-build with sm120.
12017

12118
[0.4.0]: https://github.com/pie-project/pie/compare/0.3.0...0.4.0

driver/cuda/src/loader/weight_copy_engine.hpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -350,11 +350,24 @@ class WeightCopyEngine {
350350
const std::size_t total = batched_dsts_.size();
351351
for (std::size_t off = 0; off < total; off += kChunk) {
352352
const std::size_t n = std::min(kChunk, total - off);
353+
// cudaMemcpyBatchAsync's signature changed between CUDA 12.8
354+
// (preview: non-const ptrs + a trailing `size_t* failIdx`
355+
// out-param, 9 args) and CUDA 13.0 (final: const-qualified
356+
// ptrs, no failIdx, 8 args). Pick the call shape per toolkit.
357+
#if CUDART_VERSION >= 13000
353358
const cudaError_t err = ::cudaMemcpyBatchAsync(
354359
batched_dsts_.data() + off,
355360
const_cast<const void**>(batched_srcs_.data() + off),
356361
batched_sizes_.data() + off,
357362
n, &attr, &attrs_idx, /*numAttrs=*/1, stream);
363+
#else
364+
std::size_t fail_idx = 0;
365+
const cudaError_t err = ::cudaMemcpyBatchAsync(
366+
batched_dsts_.data() + off,
367+
batched_srcs_.data() + off,
368+
batched_sizes_.data() + off,
369+
n, &attr, &attrs_idx, /*numAttrs=*/1, &fail_idx, stream);
370+
#endif
358371
if (err != cudaSuccess) {
359372
throw std::runtime_error(
360373
std::string("cudaMemcpyBatchAsync failed: ") +

driver/cuda/third_party/marlin/scalar_type.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,8 @@
77
#include <cstdint>
88
#include <stdexcept>
99
#include <string>
10+
#include <tuple>
11+
#include <utility>
1012
#include <variant>
1113

1214
#ifndef PIE_MARLIN_CHECK

runtime/Cargo.toml

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -65,7 +65,11 @@ p384 = { version = "0.13", features = ["pkcs8", "ecdsa"] }
6565
sha2 = "0.10"
6666
chrono = "0.4"
6767
ringbuffer = "0.15"
68-
reqwest = { version = "0.12", features = ["rustls-tls"] }
68+
# default-features off so reqwest's `default-tls` (native-tls → openssl) is
69+
# never pulled in: it dynamically links libssl.so.1.1, which doesn't exist on
70+
# OpenSSL-3 distros (Ubuntu 22.04+) and breaks the manylinux_2_28 release
71+
# binary. rustls is statically linked and distro-independent.
72+
reqwest = { version = "0.12", default-features = false, features = ["rustls-tls", "http2", "charset"] }
6973
crossbeam = { version = "0.8.4", features = ["crossbeam-channel"] }
7074
bytemuck = { version = "1.21", features = ["derive"] }
7175
ipc-channel = "0.19"

scripts/install.ps1

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -112,13 +112,13 @@ function Save-Url($Url, $OutFile) {
112112
function Get-AssetName($Flavor) {
113113
switch ($Flavor.ToLowerInvariant()) {
114114
{ $_ -in @("portable", "vulkan", "portable-vulkan") } {
115-
return "pie-x86_64-windows-portable-vulkan.zip"
115+
return "pie-x86_64-windows-vulkan.zip"
116116
}
117117
{ $_ -in @("cuda", "portable-cuda") } {
118-
return "pie-x86_64-windows-portable-cuda.zip"
118+
return "pie-x86_64-windows-cuda.zip"
119119
}
120120
default {
121-
Stop-WithError "no '$Flavor' build for windows/x86_64. Valid flavors: portable-vulkan, portable-cuda."
121+
Stop-WithError "no '$Flavor' build for windows/x86_64. Valid flavors: vulkan, cuda."
122122
}
123123
}
124124
}

scripts/install.sh

Lines changed: 46 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,8 @@
88
# Environment overrides:
99
# PIE_VERSION Release tag (default: 0.4.0).
1010
# PIE_FLAVOR portable | cuda{12.8,13.0}. Auto-detected when unset.
11+
# PIE_CC GPU compute capability for CUDA flavors, e.g. 90, 100
12+
# (auto-detected via nvidia-smi; selects the per-CC binary).
1113
# PIE_INSTALL_DIR Install location for the `pie` binary (default: ~/.local/bin).
1214
# PIE_REPO GitHub owner/name (default: pie-project/pie).
1315
# PIE_DOWNLOAD_BASE Override the asset base URL (default: GitHub releases).
@@ -125,34 +127,65 @@ detect_cuda_flavor() {
125127
fi
126128
}
127129

130+
# Compute capability of the first GPU, normalized to the SM string used in the
131+
# per-CC asset names ("9.0" -> "90", "12.0" -> "120"). Override with PIE_CC.
132+
detect_cuda_cc() {
133+
command -v nvidia-smi >/dev/null 2>&1 || return 0
134+
local cap
135+
cap="$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader 2>/dev/null | head -1 | tr -d ' .' || true)"
136+
case "$cap" in
137+
*[!0-9]*|"") return 0 ;;
138+
esac
139+
PIE_DETECTED_CC="$cap"
140+
}
141+
128142
if [ -z "${PIE_FLAVOR:-}" ]; then
129-
if [ "$os" = linux ] && [ "$arch" = x86_64 ]; then
143+
if [ "$os" = linux ] && { [ "$arch" = x86_64 ] || [ "$arch" = aarch64 ]; }; then
130144
detect_cuda_flavor || true
131145
PIE_FLAVOR="${PIE_DETECTED_FLAVOR:-portable}"
132146
else
133147
PIE_FLAVOR=portable
134148
fi
135149
fi
136150

137-
# Map (os, arch, flavor) -> one or more candidate release asset names, in
138-
# preference order (mirrors what .github/workflows/build.yml uploads). CUDA
139-
# flavors prefer the GH-hosted manylinux_2_28 build (broad glibc compat) and
140-
# fall back to the native linux-x64 build from the self-hosted GPU runners.
151+
# Detect the GPU compute capability so CUDA flavors can pick the per-CC binary.
152+
if [ -z "${PIE_CC:-}" ]; then
153+
detect_cuda_cc || true
154+
PIE_CC="${PIE_DETECTED_CC:-}"
155+
fi
156+
157+
# Map (os, arch, flavor) -> a candidate release asset name (mirrors what
158+
# .github/workflows/build.yml uploads). CUDA flavors select the per-compute-
159+
# capability build (glibc 2.28 floor) for the detected (or PIE_CC) capability.
141160
assets_for() {
142161
case "$os/$arch/$1" in
143-
linux/x86_64/portable) echo "pie-x86_64-manylinux_2_28.tar.gz" ;;
144-
linux/aarch64/portable) echo "pie-aarch64-manylinux_2_28.tar.gz" ;;
145-
darwin/aarch64/portable) echo "pie-aarch64-darwin.tar.gz" ;;
146-
linux/x86_64/cuda12.8 \
147-
| linux/x86_64/cuda13.0)
148-
echo "pie-x86_64-manylinux_2_28-${1}.tar.gz"
149-
echo "pie-x86_64-linux-${1}.tar.gz" ;;
162+
linux/x86_64/portable) echo "pie-x86_64-linux-vulkan.tar.gz" ;;
163+
linux/aarch64/portable) echo "pie-aarch64-linux-vulkan.tar.gz" ;;
164+
darwin/aarch64/portable) echo "pie-aarch64-macos-metal.tar.gz" ;;
165+
linux/x86_64/cuda12.8 | linux/x86_64/cuda13.0 \
166+
| linux/aarch64/cuda12.8 | linux/aarch64/cuda13.0)
167+
# Per-compute-capability binary, selected from the detected (or
168+
# PIE_CC-overridden) compute capability.
169+
if [ -n "${PIE_CC:-}" ]; then
170+
echo "pie-${arch}-linux-${1}-sm${PIE_CC}.tar.gz"
171+
fi
172+
;;
150173
*) return 1 ;;
151174
esac
152175
}
153176

154177
candidates="$(assets_for "$PIE_FLAVOR")" || err \
155-
"no '$PIE_FLAVOR' build for $os/$arch. Valid flavors: portable; or (Linux x86_64 only) cuda12.8, cuda13.0."
178+
"no '$PIE_FLAVOR' build for $os/$arch. Valid flavors: portable; or (Linux) cuda12.8, cuda13.0."
179+
180+
# CUDA flavors select per-compute-capability binaries; without a CC we have
181+
# nothing to download. Guide the user to set PIE_CC.
182+
case "$PIE_FLAVOR" in
183+
cuda*)
184+
[ -n "${PIE_CC:-}" ] || err \
185+
"CUDA flavor '${PIE_FLAVOR}' needs a GPU compute capability, but none was detected. \
186+
Set PIE_CC explicitly, e.g. PIE_CC=90 (H100), 100 (B200), 89 (L40/RTX40), 86 (A10), 80 (A100)."
187+
;;
188+
esac
156189

157190
heading "Installing Pie"
158191
detail "Version: ${PIE_VERSION}"
Lines changed: 65 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,65 @@
1+
---
2+
title: Media (multimodal)
3+
description: Splicing image, video, and audio spans into a forward pass
4+
sidebar_position: 3
5+
---
6+
7+
# Media (multimodal)
8+
9+
A forward pass accepts **encoded media spans** (images, video frames, audio clips) alongside ordinary tokens. The inferlet hands the host raw encoded bytes; the driver decodes them, runs the model's vision/audio encoder, and **scatters the projected embedding rows into the hidden state** at the positions you choose. Surrounding text attends to the span through the normal causal mask, so generation after a media span is ordinary text decoding. The API is **model-agnostic**: no model constants leak into inferlet code.
10+
11+
Two layers sit over this, mirroring tokens vs. `Forward`:
12+
13+
- **`Context::append_*`**: the high-level wrapper that splices a span at the context's current position and advances the cursor for you. Most inferlets use this; see [Multimodal generation](../model/multimodal) for the full guide and the supported-model matrix.
14+
- **`Forward::input_*`**: the per-pass builder, covered here. You pick the anchor position. Reach for it when you build passes directly: custom masks, scoring, or a hand-written decoder.
15+
16+
## Encoding media
17+
18+
Fetch the bytes and build a media handle. Decoding is synchronous and exposes layout metadata *before* any forward pass:
19+
20+
```rust
21+
use inferlet::media::{Image, Audio, Video};
22+
23+
let bytes = inferlet::http::fetch(&url).await?;
24+
let image = Image::from_bytes(&model, &bytes).map_err(|e| e.to_string())?;
25+
26+
image.token_count(); // soft tokens the span occupies in the sequence
27+
image.grid(); // (t, h, w) merged-token layout
28+
image.position_span(); // sequence positions to advance past the span
29+
```
30+
31+
`Audio::from_bytes(&model, &bytes)` and `Video::from_bytes(&model, &bytes, max_frames)` follow the same shape. For video, `max_frames` is your KV-budget knob: the host demuxes and uniformly samples up to that many frames, preprocessing each as an image span.
32+
33+
## Splicing into a forward pass
34+
35+
`Forward::input_image` / `input_audio` attach an encoded span at sequence position `anchor`. At `execute()`, the driver runs the encoder and overwrites `hidden[anchor .. anchor + token_count]` with the projected rows. Advance your position cursor by `position_span()` for whatever follows.
36+
37+
| Method | What it does |
38+
|---|---|
39+
| `fwd.input_image(&image, anchor)` | Splice a visual span (image or video frame) at `anchor`. |
40+
| `fwd.input_audio(&audio, anchor)` | Splice an audio clip at `anchor`. |
41+
42+
```rust
43+
let mut fwd = ctx.forward();
44+
let anchor = fwd.start_position(); // first free sequence position
45+
fwd.input_image(&image, anchor); // driver encodes + scatters rows at execute()
46+
let out = fwd.execute().await?;
47+
```
48+
49+
Spans are attached by reference and emitted in declaration order, so several can be spliced into one pass. The positions a span occupies are reserved like any other input; `Context::append_image` (below) does that bookkeeping for you, and at the raw level you manage pages through `ctx.inner()`.
50+
51+
## The high-level wrapper
52+
53+
`Context::append_image` is exactly `input_image` at the context's own position, plus the bookkeeping to reserve the span's pages and advance the cursor:
54+
55+
```rust
56+
ctx.append_image(&image).await?; // input_image(&image, ctx.seq_len), reserve, advance
57+
ctx.append_audio(&clip).await?;
58+
ctx.append_video(&video).await?;
59+
```
60+
61+
Use these unless you specifically need to control the anchor.
62+
63+
## Audio output
64+
65+
Audio **output** (text-to-speech) is a separate generation path via `model.speak(text).speaker(id).generate()`, not a forward-pass input. See [Multimodal generation](../model/multimodal#audio-output-text-to-speech).

website/docs/guide/forward/overview.mdx

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -114,6 +114,7 @@ The handles are typed: `out.token(probe_handle)` does not compile. Mismatched ac
114114
## What this section covers
115115

116116
- [Tokens, positions, and masks](./inputs). Building the input side of a forward pass.
117+
- [Media (multimodal)](./media). Splicing image, video, and audio spans into a pass.
117118
- [Samplers and probabilities](./samplers). Every `Sampler` kind and every `Probe` kind.
118119
- [Constrained generation](./constrained). Logit masks driven by JSON Schema, regex, EBNF, and custom grammars.
119120
- [Adapters](./adapters). LoRA adapters scoped to a model.

0 commit comments

Comments
 (0)