Skip to content

Add CUDA graph capture/replay support with a graph-safe paged-attention decode path#3733

Open
astorise wants to merge 65 commits into
huggingface:mainfrom
astorise:claude/candle-issue-15-gg7bk2
Open

Add CUDA graph capture/replay support with a graph-safe paged-attention decode path#3733
astorise wants to merge 65 commits into
huggingface:mainfrom
astorise:claude/candle-issue-15-gg7bk2

Conversation

@astorise

Copy link
Copy Markdown

Title: Add CUDA graph capture/replay support with a graph-safe paged-attention decode path

Summary

This PR adds three additive building blocks that together let a decode step
be captured once with a CUDA graph and replayed on every subsequent step,
composed with a vLLM-style paged KV cache:

  1. candle_core::CudaGraph (candle-core/src/cuda_backend/graph.rs): a
    thin capture/replay wrapper around the CUDA stream-capture API, plus an
    htod-upload cache (CudaDevice::enable_cuda_graph_htod_cache) so that
    repeated host-constant uploads inside a captured closure reuse a cached
    allocation instead of failing capture.
  2. PagedKvCache (candle-transformers/src/models/llama.rs):
    caller-owned, block-paged KV storage (vLLM-style) attached per
    transformer layer via Cache::set_paged_kv. When attached, that layer's
    attention routes through a paged/varlen flash-attention kernel against
    caller-managed key_cache/value_cache/block_table instead of the
    contiguous concat-and-narrow cache. Fully additive: layers default to the
    existing contiguous path.
  3. Cache::set_decode_position: a caller-owned, persistent device
    tensor that a layer's rotary-embedding lookup reads via
    Tensor::index_select instead of narrow(0, index_pos, seq_len), so the
    RoPE lookup's device pointer doesn't get baked into a graph capture at a
    position that goes stale on replay.

The bug this PR also fixes (#3732)

With the above three in place, driving the paged decode path through an
actual CudaGraph::capture surfaces a fourth gap: PagedKvCache's
scatter-index computation for writing new K/V is not graph-capturable.

write_new_kv's first operation is a blocking device-to-host readback of
block_table (Tensor::to_vec2) followed by a host-side loop and a fresh
Tensor::from_vec upload of the resulting scatter indices — both of which
are invalid on a stream mid-capture. Left unfixed, a capture either fails
outright or (worse) silently freezes the scatter destination to whatever
block_table contained at capture time, while the K/V values keep
updating correctly on replay — silently misplacing every subsequent decode
step's KV into the position captured at step one.

Fixes this with the same pattern as set_decode_position:

impl Cache {
    /// Attaches a persistent, caller-owned device tensor of flat scatter
    /// indices (shape `(b_sz,)`) for one layer's `PagedKvCache::write_new_kv`
    /// decode step. When attached, `write_new_kv` scatters directly against
    /// it instead of deriving indices from a host readback of `block_table`.
    /// Scoped to the decode step only (`seq_len == 1`).
    pub fn set_paged_kv_decode_slot(&mut self, block_idx: usize, indices: Tensor) -> Result<()>;
    pub fn clear_paged_kv_decode_slot(&mut self, block_idx: usize);
    pub fn paged_kv_decode_slot(&self, block_idx: usize) -> Option<&Tensor>;
}

The caller updates the attached tensor's contents in place (e.g. via
Tensor::slice_set) before each CudaGraph::replay(); its identity/address
stays fixed across replays, which is what makes the capture valid.

Scope / compatibility

Everything here is additive:

  • Cache::new and every existing caller of the contiguous KV path are
    unaffected — no behavior change, no new required arguments.
  • set_paged_kv, set_decode_position, and set_paged_kv_decode_slot all
    default to None per layer; a layer only routes through the new paths
    once a caller explicitly attaches the corresponding tensor.
  • Prefill (seq_len > 1) always uses the existing paths unconditionally;
    the position/decode-slot seams are rejected (return Err) if attached to
    a multi-token call, rather than silently producing wrong output.

Testing

  • cargo fmt, cargo clippy --lib --tests -- -D warnings (candle-core,
    candle-nn, candle-transformers).
  • cargo test (candle-core, candle-nn, candle-transformers): CPU-only unit
    tests, including numerical equivalence between the new device-tensor
    paths and the existing host-derived paths for the same position/indices,
    and rejection of malformed/prefill input on each new seam.
  • GPU tests (--features cuda,flash-attn, dispatched on a CUDA runner):
    • paged_attention_matches_dense_causal_attention_on_gpu: paged path
      matches the dense causal-attention reference.
    • decode_position_stays_correct_across_cuda_graph_replay: captures a
      decode-step rotary lookup in a real CudaGraph, replays it, updates the
      attached position tensor's contents in place, replays again, and
      asserts the output tracks the new position rather than the one live at
      capture time.
    • paged_kv_decode_slot_stays_correct_across_cuda_graph_replay: same
      property for the paged KV scatter destination — captures
      PagedKvCache::write_new_kv's decode-step scatter, replays, updates the
      decode-slot tensor's contents in place, replays again, and asserts the
      scatter lands at the new slot.

Closes #3732

claude and others added 30 commits June 17, 2026 19:47
…ing CPU

simple_eval ignored the device of its input tensors and always built
initializers, Constant/ConstantOfShape values, RandomUniform/RandomNormal
outputs and Gemm's alpha/beta scalars on Device::Cpu. Mixing those with
GPU inputs raised a device-mismatch error, making the ONNX backend
unusable on CUDA/Metal (huggingface#3491).

simple_eval, simple_eval_, get_tensor and the AttrOwned trait now take an
explicit &Device that is threaded through every previously hard-coded
call site; ops that already had a tensor to read a device from (Range,
Gemm) reuse that tensor's device instead. candle-pyo3's run() infers the
device from the first input tensor to keep its Python-facing signature
unchanged, and the bundled onnx/onnx-llm/onnx_basics/silero-vad examples
are updated for the new parameter.

Adds device-propagation regression tests for initializers, Constant,
ConstantOfShape, RandomUniform, RandomNormal, Range and Gemm, run on
CPU plus cuda/metal-gated variants (mirroring candle-core's test_device!
convention), and a cuda/metal feature section in candle-onnx's Cargo.toml
so those variants can be gated.
HashMap iteration order is unspecified, so picking the "first" input's
device could non-deterministically select a CPU metadata tensor over a
GPU data tensor, reintroducing the device-mismatch bug from huggingface#3491.
Prefer any non-CPU device among the inputs instead.

Addresses review feedback on PR #1.
Fix device mismatch in ONNX evaluation by threading device parameter
…acement

simple_eval runs a whole ONNX graph on one device. This adds an additive,
non-breaking API to spread a single graph across several devices (pipeline
parallelism), so a model that does not fit on one GPU can be split by depth
(e.g. first transformer blocks on cuda:0, the rest on cuda:1).

- New `Placement` type: resolves each node's device from the longest matching
  node-name prefix rule, otherwise the device of the node's first already
  computed input (so a stage's device flows along the graph and only the first
  node of each stage needs an explicit rule), otherwise a fallback device.
- New `simple_eval_with_placement`; `simple_eval` now delegates to it with a
  uniform placement, preserving its exact behavior.
- Initializers are materialized lazily on the device of the node that first
  consumes them (instead of all up front on a single device), which is what
  makes real memory splitting possible. Graph outputs that are unconsumed
  initializers are still materialized before returning.
- Inputs are pre-staged onto each node's device, so cross-device copies are
  inserted automatically at stage boundaries via Tensor::to_device.

Cross-device transfers currently require CUDA (CPU<->CUDA and CUDA<->CUDA);
Metal-to-Metal copies are not implemented in candle-core. Placement is expected
to be monotone along the topological order, which a by-depth pipeline split
satisfies.

Tests: CPU tests cover placement resolution, lazy-initializer handling and
uniform-placement equivalence with simple_eval; feature-gated CUDA tests cover a
real CPU<->cuda:0 round-trip (one GPU) and a cuda:0 -> cuda:1 split (two GPUs,
ignored by default). README documents the new API and its limitations.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VxsGKP59TNS1TDdnvHVA5a
- Placement::uniform (used by plain simple_eval) now always resolves a
  node's device to the fallback, instead of inheriting an already-present
  input's device when no prefix rule matches. With zero rules, input-device
  inheritance was taking priority over the requested device, letting
  simple_eval silently run on the wrong device.

- If subgraphs now inherit their enclosing graph's initializers instead of
  only seeing their own. Lazy initializer materialization only populates
  `values` when a node directly consumes an initializer, so a parent-scope
  initializer referenced solely inside an If branch was never visible to
  that branch and evaluation failed with "cannot find ... for op".

Adds a CPU-runnable regression test for the If/initializer fix and a
cuda-gated regression test for the uniform-placement fix, both verified to
fail without their respective fix.
Add pipeline parallelism support via Placement API
…ing CPU

simple_eval ignored the device of its input tensors and always built
initializers, Constant/ConstantOfShape values, RandomUniform/RandomNormal
outputs and Gemm's alpha/beta scalars on Device::Cpu. Mixing those with
GPU inputs raised a device-mismatch error, making the ONNX backend
unusable on CUDA/Metal (huggingface#3491).

simple_eval, simple_eval_, get_tensor and the AttrOwned trait now take an
explicit &Device that is threaded through every previously hard-coded
call site; ops that already had a tensor to read a device from (Range,
Gemm) reuse that tensor's device instead. candle-pyo3's run() infers the
device from the first input tensor to keep its Python-facing signature
unchanged, and the bundled onnx/onnx-llm/onnx_basics/silero-vad examples
are updated for the new parameter.

Adds device-propagation regression tests for initializers, Constant,
ConstantOfShape, RandomUniform, RandomNormal, Range and Gemm, run on
CPU plus cuda/metal-gated variants (mirroring candle-core's test_device!
convention), and a cuda/metal feature section in candle-onnx's Cargo.toml
so those variants can be gated.
HashMap iteration order is unspecified, so picking the "first" input's
device could non-deterministically select a CPU metadata tensor over a
GPU data tensor, reintroducing the device-mismatch bug from huggingface#3491.
Prefer any non-CPU device among the inputs instead.

Addresses review feedback on PR #1.
…acement

simple_eval runs a whole ONNX graph on one device. This adds an additive,
non-breaking API to spread a single graph across several devices (pipeline
parallelism), so a model that does not fit on one GPU can be split by depth
(e.g. first transformer blocks on cuda:0, the rest on cuda:1).

- New `Placement` type: resolves each node's device from the longest matching
  node-name prefix rule, otherwise the device of the node's first already
  computed input (so a stage's device flows along the graph and only the first
  node of each stage needs an explicit rule), otherwise a fallback device.
- New `simple_eval_with_placement`; `simple_eval` now delegates to it with a
  uniform placement, preserving its exact behavior.
- Initializers are materialized lazily on the device of the node that first
  consumes them (instead of all up front on a single device), which is what
  makes real memory splitting possible. Graph outputs that are unconsumed
  initializers are still materialized before returning.
- Inputs are pre-staged onto each node's device, so cross-device copies are
  inserted automatically at stage boundaries via Tensor::to_device.

Cross-device transfers currently require CUDA (CPU<->CUDA and CUDA<->CUDA);
Metal-to-Metal copies are not implemented in candle-core. Placement is expected
to be monotone along the topological order, which a by-depth pipeline split
satisfies.

Tests: CPU tests cover placement resolution, lazy-initializer handling and
uniform-placement equivalence with simple_eval; feature-gated CUDA tests cover a
real CPU<->cuda:0 round-trip (one GPU) and a cuda:0 -> cuda:1 split (two GPUs,
ignored by default). README documents the new API and its limitations.
- Placement::uniform (used by plain simple_eval) now always resolves a
  node's device to the fallback, instead of inheriting an already-present
  input's device when no prefix rule matches. With zero rules, input-device
  inheritance was taking priority over the requested device, letting
  simple_eval silently run on the wrong device.

- If subgraphs now inherit their enclosing graph's initializers instead of
  only seeing their own. Lazy initializer materialization only populates
  `values` when a node directly consumes an initializer, so a parent-scope
  initializer referenced solely inside an If branch was never visible to
  that branch and evaluation failed with "cannot find ... for op".

Adds a CPU-runnable regression test for the If/initializer fix and a
cuda-gated regression test for the uniform-placement fix, both verified to
fail without their respective fix.
Brings main fully up to date with both onnx fixes (device-aware
simple_eval and pipeline-parallel placement) plus the upstream
huggingface/candle commits already absorbed by that branch, while
the two changes await review as separate pull requests upstream.
Mirrors ci_cuda.yaml: a workflow_dispatch (with optional test_ref) +
pull_request workflow that runs the candle-nn paged-attention tests on
the free GitHub-hosted macos-14 (Apple Silicon) runner with --features
metal and CANDLE_METAL_REQUIRED=1, so the Metal path is genuinely
exercised. Lives on the default branch so it is dispatchable; uses
test_ref to validate upstream-bound branches without adding CI config to
them.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_012k9NFs5twmQSKDumHekARt
candle's Metal backend uses the MTLResidencySet API (Metal 3.2), absent
on the macos-14 image, which made the Metal tests panic with
'MTLResidencySetDescriptor could not be found'. macos-15 provides it.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_012k9NFs5twmQSKDumHekARt
CPU dequantization (quantized_gptq::gptq_linear) plus fused dequant+GEMM
CUDA/Metal kernels (candle-gptq-kernels), including a real tensor-core
(WMMA mma.sync) GEMM path bound via the vendored Marlin kernel for 4-bit
weights. Wired in behind opt-in gptq-cuda/gptq-metal features on
candle-transformers.

Adds candle_transformers::models::gptq_qwen2, a Qwen2 variant that routes
every attention/MLP projection through gptq_linear, and an end-to-end
quantized-gptq-qwen2 example that downloads a real AutoGPTQ/GPTQModel
checkpoint from the Hugging Face Hub and runs generation through it.

This is the GPTQ-only slice of the broader GPTQ/AWQ/FP8 quantization work
for issue huggingface#3650, split out to keep each format's PR independently
reviewable and end-to-end testable in CI.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014pmFaDvgxprkJxEFq6XDFL
…gingface#3650)

CPU dequantization (quantized_awq::awq_linear) plus fused dequant+GEMM
CUDA/Metal kernels (candle-awq-kernels), wired in behind opt-in
awq-cuda/awq-metal features on candle-transformers.

Adds QuantizedLinear: a single enum spanning the GPTQ/AWQ formats and
their dense/CUDA/Metal load paths, analogous to QMatMul for the GGUF
path, so callers no longer match on format and call one of
gptq_linear/awq_linear or four separate *Cuda/*Metal structs themselves.

Generalizes the previous GPTQ-only quantized-gptq-qwen2 example/model
(renamed quantized-qwen2 / quant_linear_qwen2) to read
quantization_config.quant_method from the checkpoint's config.json and
dispatch to GPTQ or AWQ automatically via QuantizedLinear, with
per-checkpoint bias detection since AWQ exports don't put bias on every
projection the way GPTQ exports do. Verified against the real
Qwen/Qwen2-0.5B-Instruct-AWQ checkpoint's safetensors layout.

This is the GPTQ+AWQ slice of the broader GPTQ/AWQ/FP8 quantization work
for issue huggingface#3650, building on the GPTQ-only PR and split out to keep each
format's PR independently reviewable and end-to-end testable in CI.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014pmFaDvgxprkJxEFq6XDFL
…3650)

Adds DeepSeek-V3-style block-wise FP8 (per-128x128-block weight_scale_inv)
support alongside the existing GPTQ/AWQ formats: candle-fp8-kernels provides
the fused dequantize+GEMM CUDA/Metal kernels, and
candle_transformers::quantized_fp8 provides the portable CPU dequantize-at-load
path. quantized_linear::QuantizedLinear gains an Fp8 variant, dispatching to
the fused kernel when the fp8-cuda/fp8-metal feature is enabled and falling
back to the dense path otherwise, exactly mirroring the Gptq/Awq variants.

The quantized-qwen2 example's docs are updated to note that FP8 isn't wired up
end-to-end there: the small public Qwen2 FP8 checkpoints on the Hub use
per-tensor static scales (vLLM/compressed-tensors layout) rather than the
block-wise layout this crate implements, so there's no small checkpoint to
demonstrate it against; quantized_fp8's own unit tests cover it instead.

This is the GPTQ+AWQ+FP8 slice of the broader quantization work for issue
huggingface#3650, building on the GPTQ+AWQ PR (claude/quant-awq) to complete the format
split so each PR's CI independently validates its own format(s).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014pmFaDvgxprkJxEFq6XDFL
…ace#3650)

Brings the GPTQ + AWQ + block-wise FP8 fused dequant+GEMM quantization work
(the three branches proposed upstream as separate PRs) onto the fork's main so
it can be used while those PRs await review.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014pmFaDvgxprkJxEFq6XDFL

# Conflicts:
#	.github/workflows/ci_metal.yaml
Ports the CPU/Metal PagedAttention (candle_nn::attention::paged + its
integration test) from the working branch onto main, alongside the just-merged
GPTQ/AWQ/FP8 quantization work, so the fork's main carries the full feature set
while the upstream PRs await review.

Also restores the candle-{gptq,awq,fp8}-kernels entries in workspace.dependencies
and bumps those crates to 0.11.0: a prior GitHub-side 'merge main into branch'
on the quant branch had dropped them from workspace.dependencies while leaving
candle-transformers referencing them via workspace = true, which broke manifest
parsing for the whole workspace.

Verified: cargo build -p candle-transformers, paged_attention integration test
(6 passed) and GPTQ/AWQ/FP8 dequant roundtrip tests (all passed), clippy
-D warnings and cargo fmt --check all clean.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014pmFaDvgxprkJxEFq6XDFL
…ce#3651)

Adds CudaGraph::capture(&device, || { ... }) and .replay() on top of
the existing CUDA-graph-aware host-to-device upload cache, so a fixed-shape
decode step's kernels can be recorded once and relaunched with a single
cuGraphLaunch call instead of relaunching every kernel per generated token.
A matching dummy stub keeps non-CUDA builds compiling.

FlashInfer backend support from the same issue is not included here; it
needs its own kernel bindings and is left as follow-up work.
candle-flashinfer-kernels is a new optional crate providing
flashinfer_decode_attention(q, k, v, scale): a feature-gated alternative
to candle-flash-attn for the single-new-token decode step referenced in
issue huggingface#3651. It implements the same shape contract FlashInfer's
batch-decode kernels target (one query token per sequence attending over
a KV cache, with GQA support) as a reference streaming-softmax CUDA
kernel; it is not a port of FlashInfer's own tensor-core/split-KV kernels,
so it should not be expected to match their throughput, but it gives
candle a working, swappable decode-attention entry point.

Also points ci_cuda.yaml at the arc-gpu-candle self-hosted GPU runner
(with a workflow_dispatch test_ref input) so this and the CudaGraph
capture/replay work from the previous commit can actually be compiled
and tested against real CUDA hardware, and adds CI steps for both.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
cargo test -p candle-core --features cuda --test cuda_graph_tests was
failing on real CUDA hardware with CUDA_ERROR_STREAM_CAPTURE_INVALIDATED:
the test launched the multiply kernel for the first time directly inside
capture, so the lazy CUDA module load for that kernel happened mid-capture,
which is disallowed and invalidates the capture. Run the op once outside
capture first (and document the requirement on CudaGraph::capture) so
capture only ever records already-loaded kernel launches.
claude and others added 30 commits June 27, 2026 09:36
The decode-attention backend previously had a stub cpu_fwd that bailed and
hard-required the CUDA toolchain to even build. Implement a real reference
cpu_fwd (f32/f16/bf16, GQA, stride/offset aware, numerically-stable softmax
computed in f32) so flashinfer_decode_attention works on CPU.

Make `cuda` an opt-in feature: without it the crate builds CPU-only (no
nvcc, no candle/cuda), with the CUDA kernel and GPU forward pass gated
behind the feature. Drop the candle-nn dev-dependency (inline the softmax
in the test reference) so tests build without CUDA too, and add a
decode_attention_cpu_matches_reference test that runs on Device::Cpu.

CI: build/test the flashinfer step with --features cuda so the GPU kernel
is still compiled and exercised on the runner alongside the CPU test.
… Silicon

CPU: parallelize cpu_decode_attention across (batch, head) with rayon, so the
fallback scales on many-core CPUs including Apple Silicon (per-task softmax
scratch; disjoint output slices via par_chunks_mut).

Metal: add an opt-in `metal` feature with a runtime-compiled decode-attention
shader (kernels/decode_attention.metal) mirroring the CUDA reference path - one
threadgroup per (batch, head), threadgroup-reduced dot products, streaming
softmax - for f32/f16. Wire metal_fwd into the CustomOp3 and cache the compiled
pipeline per (device, function). Add a Metal correctness test and a ci_metal.yaml
workflow that runs it on a macos-15 Apple Silicon runner alongside the CPU test.
The Metal shader failed to compile: MSL requires all [[*_position_*]] /
[[threads_per_threadgroup]] attributes in a kernel to share the same
dimensionality, but the entry points mixed uint2 (threadgroup_position_in_grid)
with uint (thread_position_in_threadgroup, threads_per_threadgroup). Declare all
three as uint2 and take .x for the 1-D threadgroup index/size.
Restore the upstream runner config (group: aws-g5-4xlarge-cache, container
without --gpus all, unscoped `cargo test --features cuda`) instead of the
self-hosted arc-gpu-candle runner and its OOM workarounds, which were specific
to a smaller box. Keep the added FlashInfer decode-attention test step. The
CudaGraph capture/replay integration test is already covered by the main
`cargo test --features cuda` step (candle-core integration test).
… main

# Conflicts:
#	.github/workflows/ci_cuda.yaml
#	.github/workflows/ci_metal.yaml
#	Cargo.toml
Implements candle-nn::lora::LoraLinear, a Linear wrapper supporting
PEFT-style adapter loading via VarBuilder, weight merge/unmerge for
inference, and named multi-adapter registration with active-adapter
selection for hot-swapping.

Closes huggingface#3652

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01N7G4dAaGfkGHtejuq522qW
Implements candle-nn::lora::LoraLinear, a Linear wrapper supporting
PEFT-style adapter loading via VarBuilder, weight merge/unmerge for
inference, and named multi-adapter registration with active-adapter
selection for hot-swapping.

Closes huggingface#3652


Claude-Session: https://claude.ai/code/session_01N7G4dAaGfkGHtejuq522qW

Co-authored-by: Claude <noreply@anthropic.com>
Wires the candle_nn::lora::LoraLinear primitive into candle-transformers'
Llama model. Llama::load_with_config takes a LlamaLoadConfig built via
with_lora_adapter(name, vb, LoraConfig { rank, alpha, target_modules })
to inject adapters into the matching attention/MLP projections, and
Llama::set_active_adapter lets callers hot-swap the active adapter (or
fall back to the frozen base weights) per request without reloading or
duplicating the base model. Llama::load is unchanged and delegates to
load_with_config with no adapters registered.

with_lora_adapter resolves LoRA tensors under the standard PEFT
checkpoint prefix (base_model.model.*) by default;
with_lora_adapter_prefixed lets callers override or disable that
prefix for adapters saved under a different root.

Closes huggingface#3696
Brings in the Llama LoRA loader integration (issue huggingface#3696) on top of
the LoraLinear primitive already merged for huggingface#3652: LlamaLoadConfig,
Llama::load_with_config, Llama::set_active_adapter, and PEFT checkpoint
prefix handling, plus the associated tests. Also carries an unrelated
GELU ONNX test tolerance fix (huggingface#3679), merged as identical in effect to
the version already on main (617d582), keeping main's formatting.
Adds a caller-owned PagedKvCache type and a Cache::new_paged constructor
alongside the existing contiguous Cache::new path, so downstream crates
(e.g. Tachyon-Mesh) can drive candle_flash_attn::flash_attn_varlen_paged_windowed
through Llama::forward without forking the model or reimplementing RoPE/GQA.
Block allocation, eviction and admission stay a caller concern; the model
only writes new K/V into the slots the caller's block table indicates.

Existing Cache::new callers are unaffected. A CPU-only regression test
(no flash-attn feature needed) exercises the dense fallback path, proving
the paged write/gather addressing reproduces the contiguous cache's output
exactly across a prefill + multi-step decode sequence.

Closes #8

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Add Cache::new_paged seam to llama for external paged-attention wiring
This reverts commit 42d6135, reversing
changes made to 60b4ad1.
…wiring

Adds a caller-owned PagedKvCache struct and per-layer Cache::set_paged_kv/
clear_paged_kv/paged_kv methods to candle-transformers::models::llama. When a
layer has paged storage attached, CausalSelfAttention::forward writes the new
K/V into the caller-owned key_cache/value_cache at the slots block_table
designates, then attends via
candle_flash_attn::flash_attn_varlen_paged_windowed instead of the contiguous
concat-and-narrow path.

Layers default to None, so Cache::new, Llama::forward/load, and every existing
field (use_kv_cache, etc.) are unchanged for current callers; the paged path is
purely opt-in and composes with the existing LoRA adapter seam (Proj::Plain /
Proj::Lora) without touching it. This lets a downstream consumer (e.g. a block
allocator/eviction engine) drive paged attention through Llama::forward without
forking the model or reimplementing RoPE/GQA/norm placement to reach the
attention call site.

Includes CPU-only unit tests for the slot-math/write path plus a GPU
correctness test (feature-gated behind flash-attn) comparing the paged path
against the existing dense causal path.

Fixes #8
…dings

paged_test_config()'s base tiny_config() sets max_position_embeddings to 16,
too small for the test's seq_len of 17 - the dense reference path narrows the
RoPE cos/sin cache to max_position_embeddings rows and panicked. Override it
explicitly for this test.
Add additive Cache::Paged seam to llama for external paged-attention wiring
…de-attention wiring

Rebased onto main (which already carries candle-flashinfer-kernels and the
paged-attention Cache::Paged seam from #8) to resolve the merge conflicts
from the original sync/upstream-main-based branch. The crate port is
dropped since main already has it; only the genuinely new delta remains:
wiring candle-flashinfer-kernels into candle-transformers via a new
flashinfer-kernels cargo feature, and threading a use_flashinfer_attention
flag through llama::Config/CausalSelfAttention.

When set and seq_len == 1, the decode step attends via
candle_flashinfer_kernels::flashinfer_decode_attention against the
pre-repeat_kv contiguous cache instead of the dense matmul+softmax or
flash_attn path. A layer with paged KV storage attached always takes the
existing paged path regardless of this flag. Prefill and every other path
are byte-for-byte unchanged when the flag is left at its default (false).

Fixes #11

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Y5Lo9NooyDQgpJkGhaGJcg
Add additive use_flashinfer_attention seam to llama for external decode-attention wiring
…ecode capture

candle_transformers::models::llama's rotary embedding lookup computed
cos/sin via narrow(0, index_pos, seq_len), a host-derived offset baked
into the device pointer at capture time. A CudaGraph capturing that
sequence would replay the same rotary embeddings on every subsequent
decode step instead of re-deriving them, since index_pos increments
each step but the graph does not re-launch with a fresh offset.

Add Cache::set_decode_position/clear_decode_position/decode_position,
mirroring the existing set_paged_kv seam: a caller-owned Tensor((1,))
attached per transformer layer. When attached, apply_rotary_emb reads
cos/sin via Tensor::index_select against that device tensor instead of
narrow, so only the tensor's *contents* (updated in place before each
CudaGraph::replay) change across replays, not the captured kernel
launches or addresses. Scoped to the decode step only (seq_len == 1);
CausalSelfAttention::forward returns an error if a layer with a decode
position attached sees a multi-token (prefill) input.

Layers default to None, so Cache::new and every existing caller are
unaffected; this is purely additive, same pattern as set_paged_kv.

Fixes #12

(cherry picked from commit 82d2118)
The existing tests only checked that index_select against a decode
position tensor matches narrow for a fixed position - they didn't
exercise an actual CudaGraph capture/replay cycle, which is the
scenario issue #12 is about (a graph baking in a stale position).

decode_position_stays_correct_across_cuda_graph_replay captures a
decode step with a position tensor attached, replays it once at the
captured position (0) and checks it matches the dense narrow-based
path, then updates the position tensor's contents in place (not its
identity) to 5, replays again, and checks the output now matches
position 5 instead of staying stuck on 0. Feature-gated behind
flash-attn (implies cuda), same as the existing paged-attention GPU
test in this file; not runnable without a CUDA device.

(cherry picked from commit df72e32)
…e it

The GPU CI run exposed a real gap: decode_position_stays_correct_across_cuda_graph_replay
failed with "h2d param cache miss during CUDA graph capture: warmup did not
populate all required parameter vectors".

CudaGraph::capture only holds CudaDevice::enable_cuda_graph_htod_cache for
its own call, immediately before starting stream capture - it does not
extend that guard to whatever warm-up the caller ran beforehand. Both the
fork's host-to-device upload cache (cuda_backend/device.rs) and upstream's
kernel-launch parameter cache (cuda_backend/mod.rs's params_from_vec, used
for non-contiguous layouts such as the one index_select's gather produces
here) only populate their cache entries while that guard is held and stream
capture is not yet active. A warm-up run outside the guard still JIT-loads
kernels, but never seeds either cache, so capture then fails the first time
it needs an upload neither cache has seen.

Fix the test by wrapping its warm-up call in
cuda_device.enable_cuda_graph_htod_cache(), and make this an explicit part
of CudaGraph::capture's documented contract so the next caller of
Cache::set_decode_position doesn't hit the same silent trap.

(cherry picked from commit f98a8ff)
…ion output

The previous version routed the CudaGraph replay check through the full
CausalSelfAttention::forward and failed its own assert_ne! sanity check:
with a single query attending only to itself and no cached history,
softmax over one key is trivially 1.0, so the attention *output* equals V
regardless of position - it can never distinguish a graph replaying a
stale captured position from one reading the current position, since q/k
rotation doesn't affect the result at all in that degenerate shape.

Capture attn.apply_rotary_emb directly on a q-shaped tensor instead: the
rotated tensor itself does differ between position 0 (identity rotation)
and position 5, so the replay-tracks-updated-position assertions are
actually meaningful now.

(cherry picked from commit d1bedb5)
…elfAttention shape

Rebasing the decode_position seam onto main (which now also carries the
LoRA Proj wrapper and the flashinfer_attention seam from #11) surfaced a
stale test helper: it built CausalSelfAttention with plain Linear fields
and no use_flashinfer_attention, which no longer matches the struct.
Wrap q/k/v/o_proj in Proj::Plain and set use_flashinfer_attention: false,
matching the pattern already used by the paged-attention GPU test in the
same file.
Pre-existing on main, unrelated to the decode-position seam - fixing here
since it currently fails the Clippy check (-D warnings) on any PR against
main, including this one.
Unrelated to the decode-position seam - main's CI clippy job
(cargo clippy --workspace --tests --examples --benches -- -D warnings)
currently fails on a scattering of pre-existing lints across
candle-examples (useless_borrows_in_formatting, iterating a map via
.iter() instead of .keys()/.values()), tensor-tools, and two
manual_filter cases in candle-transformers, none of which this PR
touches. Fixing them here since they block this PR's own CI otherwise;
all mechanical, behavior-preserving changes (cargo clippy --fix plus two
manual equivalents), verified with a clean full-workspace
`cargo clippy --workspace --tests --examples --benches -- -D warnings`
and `cargo fmt --all -- --check`.
Add graph-replay-safe Cache::set_decode_position seam for CudaGraph decode capture
…ture

write_new_kv's scatter-index computation currently starts with a blocking
device-to-host readback of block_table (Tensor::to_vec2) followed by a fresh
Tensor::from_vec upload, both invalid on a stream mid CudaGraph capture.

Adds Cache::set_paged_kv_decode_slot/clear_paged_kv_decode_slot/paged_kv_decode_slot,
mirroring the existing set_decode_position seam: a caller-owned, persistent
device tensor of flat scatter indices, updated in place before each replay.
When attached, write_new_kv scatters directly against it instead of deriving
indices from the host-side block_table readback. Scoped to the decode step
only (seq_len == 1); layers default to no decode slot, so PagedKvCache and
every existing write_new_kv caller are unaffected.

Fixes #15.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

PagedKvCache::write_new_kv's host readback breaks CudaGraph capture for the decode step

2 participants