Add CUDA graph capture/replay support with a graph-safe paged-attention decode path#3733
Open
astorise wants to merge 65 commits into
Open
Add CUDA graph capture/replay support with a graph-safe paged-attention decode path#3733astorise wants to merge 65 commits into
astorise wants to merge 65 commits into
Conversation
…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.
…nto claude/onnx-pipeline-parallel
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.
…attn, huggingface#3643 CPU quant) into fork main
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.
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
…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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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:
candle_core::CudaGraph(candle-core/src/cuda_backend/graph.rs): athin capture/replay wrapper around the CUDA stream-capture API, plus an
htod-upload cache (
CudaDevice::enable_cuda_graph_htod_cache) so thatrepeated host-constant uploads inside a captured closure reuse a cached
allocation instead of failing capture.
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'sattention routes through a paged/varlen flash-attention kernel against
caller-managed
key_cache/value_cache/block_tableinstead of thecontiguous concat-and-narrow cache. Fully additive: layers default to the
existing contiguous path.
Cache::set_decode_position: a caller-owned, persistent devicetensor that a layer's rotary-embedding lookup reads via
Tensor::index_selectinstead ofnarrow(0, index_pos, seq_len), so theRoPE 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::capturesurfaces a fourth gap:PagedKvCache'sscatter-index computation for writing new K/V is not graph-capturable.
write_new_kv's first operation is a blocking device-to-host readback ofblock_table(Tensor::to_vec2) followed by a host-side loop and a freshTensor::from_vecupload of the resulting scatter indices — both of whichare invalid on a stream mid-capture. Left unfixed, a capture either fails
outright or (worse) silently freezes the scatter destination to whatever
block_tablecontained at capture time, while the K/V values keepupdating 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:The caller updates the attached tensor's contents in place (e.g. via
Tensor::slice_set) before eachCudaGraph::replay(); its identity/addressstays fixed across replays, which is what makes the capture valid.
Scope / compatibility
Everything here is additive:
Cache::newand every existing caller of the contiguous KV path areunaffected — no behavior change, no new required arguments.
set_paged_kv,set_decode_position, andset_paged_kv_decode_slotalldefault to
Noneper layer; a layer only routes through the new pathsonce a caller explicitly attaches the corresponding tensor.
seq_len > 1) always uses the existing paths unconditionally;the position/decode-slot seams are rejected (return
Err) if attached toa 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 unittests, 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.
--features cuda,flash-attn, dispatched on a CUDA runner):paged_attention_matches_dense_causal_attention_on_gpu: paged pathmatches the dense causal-attention reference.
decode_position_stays_correct_across_cuda_graph_replay: captures adecode-step rotary lookup in a real
CudaGraph, replays it, updates theattached 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: sameproperty for the paged KV scatter destination — captures
PagedKvCache::write_new_kv's decode-step scatter, replays, updates thedecode-slot tensor's contents in place, replays again, and asserts the
scatter lands at the new slot.
Closes #3732