Skip to content

[Cake 9/13] dense_gemm: batched projection and FP32 router GEMMs - #9

Merged
Zhichenzzz merged 98 commits into
radixark:mainfrom
yyihuang:cake-stack/09-dense-gemm
Nov 5, 2026
Merged

Zhichenzzz merged 98 commits into
radixark:mainfrom
yyihuang:cake-stack/09-dense-gemm

Conversation

@yyihuang

@yyihuang yyihuang commented Nov 4, 2026 •

Copy link
Copy Markdown
Contributor

Stack 9/13 · base: main

Summary

This PR adds miles_ext_kernels.cake.dense_gemm on the shared cake runtime. It provides two BF16 GEMMs, each with forward and backward:

  • batched_projection(x, weight, *, weight_layout) is a per-head projection x [T, H, D_in] → [T, H, D_out]. It serves MLA q-absorb ("in_out") and the value and output-group projections ("out_in").
  • router_gating(x, weight) computes FP32 MoE router logits x @ weight.T. It matches Megatron's RouterGatingLinearFunction with --moe-router-dtype fp32 and no bias.

The PR vendors two packages from FlashInfer c5db8d3:

  • dense_projection_det_gemm (det);
  • dense_projection_gemm (fast), which runs its deterministic rows on the det package.

In miles, --batched-projection-gemm-backend cake and --moe-router-gemm-backend cake select the op (radixark/miles#4684). Either one requires --train-backend megatron. The command line offers det only, so every miles process runs det. Direct callers can reach fast through configure(cake_dense_gemm_variant="fast"). The previous version of this layer shipped det only.

Other changes from the previous version:

  • cake.prebuild("dense_gemm", ...) takes model shapes and no token count.
  • The results are documented as not batch-invariant.
  • The runtime owns the variant and the arguments.
  • A new extra, dense_gemm (cuda-bindings>=12.8), is listed in the README Install extras row.

Stack

This is layer 9 of 13, on top of layer 8 (gqa). The stack was rewritten as [Cake 0/13] layout plus 13 op layers of three commits each (vendor / op / tests). Every PR targets main, so this diff contains layers 0 to 9. The previous version of this layer was reviewed in #9 ([Cake 9/11] dense_gemm: det batched projections and FP32 router GEMM).

This layer's commits:

  • ef381a0 vendor: dense_gemm packages
  • 20ccf43 dense_gemm: batched projection and FP32 router GEMMs on the cake runtime
  • 9eead0b tests: dense_gemm

Changes

From the layer 8 head to the layer 9 head: 238 files, +162,232 / −0. Per commit: ef381a0 221 files +159,637; 20ccf43 6 files +640; 9eead0b 11 files +1,955.

Group Files Lines
Vendored cake/_vendor/dense_projection_det_gemm/ 190 +138,713
Vendored cake/_vendor/dense_projection_gemm/ 30 +20,680
cake/_vendor/VENDORED.json (ops.dense_gemm block) 1 +244
Op module and implementation: cake/dense_gemm.py, cake/_impl/dense_gemm/{__init__,ops}.py 3 +633
Tests: tests/cake/dense_gemm/ 11 +1,955
Registry, README (Install extras, Ops and Prebuild rows), pyproject.toml (extra dense_gemm = ["cuda-bindings>=12.8"]) 3 +7

Packages

Package Upstream Role Libraries (sm_100a / sm_103a / sm_107a) Minimum CUDA
dense_projection_det_gemm flashinfer-ai/flashinfer@c5db8d3, flashinfer/experimental/dense_projection_det_gemm (190 files, Cake revision 72d0af54b2ddd96cd9f931508e70999338e18cce) det, and the deterministic sibling that the fast package binds 268 = 87 / 90 / 91 12.8 / 12.9 / 13.4
dense_projection_gemm same commit, flashinfer/experimental/dense_projection_gemm (30 files, same Cake revision) fast: non-deterministic split-K (K1) and stream-K (K2) weight-gradient programs; binds the det package through use_deterministic_backend for every other row 30 = 9 / 10 / 11 12.8 / 12.9 / 13.4

Both packages come from one commit, because the fast package runs its deterministic rows on its sibling. There are no allow entries and no local files. There is no KCP variant, because the op has no context-parallel form.

python tools/vendor_cake.py fi --checkout ../flashinfer --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 \
    --op dense_gemm --package flashinfer/experimental/dense_projection_det_gemm \
    --package flashinfer/experimental/dense_projection_gemm

FlashInfer revision

The message of commit 20ccf43 explains the choice of c5db8d3 from the FlashInfer PR history:

API

OP is name="dense_gemm", variants=("det", "fast"), cli_variants=("det",), selected_by=(("batched_projection_gemm_backend", "cake"), ("moe_router_gemm_backend", "cake")), with packages=("dense_projection_det_gemm", "dense_projection_gemm") and no extra flags. The module is torch-free. It exports OPERATIONS (("forward", "input_grad", "weight_grad")), WEIGHT_LAYOUTS (("in_out", "out_in")), batched_projection, check_projection_dims, fast_projection_programs, fast_router_programs and router_gating.

  • batched_projection:
    • Weight layouts: "in_out" takes weight [H, D_in, D_out] (einsum("thd,hdm->thm")); "out_in" takes weight [H, D_out, D_in] (einsum("thm,hdm->thd")).
    • It is differentiable in x and weight.
    • Each result is an FP32 sum of exact BF16 products, rounded once. These are cuBLAS numerics, not cuBLAS's bits.
    • Views are read in place when they are 16-byte aligned, with unit inner stride and 16-byte outer strides. Other views are copied.
    • D_in and D_out must be positive multiples of 8.
  • router_gating:
    • It takes BF16 x [..., H] and weight [E, H]. E and H must be positive multiples of 8.
    • The forward is a K1 program with FP32 output.
    • The backward computes dx = g @ weight.float() and dweight = g.T @ x.float() in FP32 through the exact BF16×3 split (K2). The results are rounded to the dtypes of x / weight.
  • check_projection_dims(d_in, d_out) raises ValueError unless both are positive multiples of 8.
  • fast_projection_programs(...) / fast_router_programs(...) return the program that fast runs for each operation: a template name, or None for the deterministic program. They need torch, but no device.
  • OP.validate: a selected backend requires --train-backend megatron; otherwise it raises ValueError.
  • An unsupported device raises NotImplementedError. A toolkit too old for the architecture raises RuntimeError (CUDA 13.4 for sm_107a).
  • README Ops row: GPUs sm_100a, sm_103a, sm_107a; prebuild parameters projections, router; tests in tests/cake/dense_gemm/.

Variant routing

Situation Result
Process never configured det
miles CLI with either selector, flag unset, deterministic or not det: --cake-dense-gemm-variant is registered but hidden (one CLI choice, det), and validate_args writes det back
Direct caller, configure(cake_dense_gemm_variant="fast") fast; refused under deterministic_mode=True
At a call every forward reads config.variant("dense_gemm"); the backward runs the variant of its forward
fast under torch.use_deterministic_algorithms(True) runs det, logged once (warn_runs_det)
  • What fast changes. It changes only weight gradients, and only on the rows and token counts its tables select for the device's architecture and SM count:

    • the K1 weight gradients of single-batch rows with d_in 4096–8192 and d_out 32–1024;
    • the router weight gradient of 256 experts at hidden 4096 / 5120 / 6144, inside measured token windows.

    Forwards, input gradients and every other row run the deterministic programs and give det's bits.

  • miles' own calls under fast. In the recorded plan table (token counts 64 to 65536), only one of miles' calls runs a program of its own: the router weight gradient of 256 experts at hidden 4096 / 5120 / 6144. It does so inside token windows, on 152-SM sm_100a / sm_103a parts and on 208–216-SM sm_107a parts. On 148-SM parts and 160-SM sm_103a parts, every miles call runs the deterministic program. fast becomes a CLI choice once it is shown to be reachable and correct for miles' calls.

  • Known accuracy misses of fast. The _impl/dense_gemm/ops.py docstring lists these, taken from perf(cake_dense_projection_gemm): restore the K1 weight-gradient routes that meet the deterministic package's absolute FP64 budgets flashinfer-ai/flashinfer#8743 measurements.

    • On 148-SM sm_103a parts, the two-split K1 weight gradient of the single-head 5120 → 48 and 5120 → 64 BF16 projections missed the det package's BF16 budget (atol = rtol = 1e-2, no element outside) for one of 10 seeds.
    • The misses were at 40961 and 40962 tokens (5120 → 48) and at 47105 tokens (5120 → 64).
    • The deterministic program met the budget on every seed.
    • miles calls neither shape and runs det.
  • Batch invariance. Neither variant is batch-invariant.

    • The plan depends on the token count. A token's outputs, logits and gradients can therefore change in their last bits with the token count of the call, within the accuracy budget.
    • As a result, a router top-k can depend on packing when two logits are that close.
    • det repeats bit for bit for identical inputs of the same shape on the same device.
    • The README Ops row, the op module, ops.py and both entry points' docstrings all say so.

Prebuild

cake.prebuild("dense_gemm", *, device=None, max_workers=None, projections=(), router=()) takes (num_heads, d_in, d_out, weight_layout) tuples in projections and (num_experts, hidden_size) tuples in router. It takes no token count.

  • It builds every program that the two packages register for the device's architecture, for the GEMM kinds given.
  • It checks the shapes first (ValueError), before the device. With both lists empty, it returns [].
  • Afterwards, no call of those GEMM kinds compiles on that device, at any token count, in either variant.

The README Prebuild row reads: dense_gemm | projections: (num_heads, d_in, d_out, weight_layout) tuples; router: (num_experts, hidden_size) tuples. The plan of a GEMM depends on the token count, so it builds every program the two packages register for the architecture for the GEMM kinds given (the projections, the router), not one plan per shape.

Memory

  • The op does not use the runtime workspace pool, so the pool bound (default 0) does not affect it.
  • The vendored entry points plan once per (device, stream, view geometry) and keep a bounded cache of planned launches.
  • The cached launches of a stream share a growable workspace for that device and stream. Calls on two streams never share one.
  • release() (clear_convenience_caches of both packages) drops the plans and workspaces. The README Memory section lists these as out-of-pool per-stream workspaces that grow to the largest call.
  • The op is not CUDA-graph safe. A later growth replaces the workspace that a graph would replay, and an eviction frees it.
  • The packages use the dense_gemm extra (cuda-bindings>=12.8) to call the CUDA driver. They use it to read the L2 size when torch does not report it, and to zero-fill the FP32 router gradients under fast.

Review items

These items come from the review of the previous version: #9 (comment). The measurements are the reviewer's, taken on GB300.

# Reviewer item Resolution Tests
9-1 Mid-step JIT. After a prebuild at T = 33 / 1001 / 2049, 50 random T × 4 ops compiled 2 new libraries at T = 239 (10.1 s and 4.6 s). The reviewer asked for a T-independent prebuild. prebuild now takes model shapes only. It builds every registered program of the given GEMM kinds for the architecture. CPU: test_prebuild_plans.py::test_every_planned_program_at_any_token_count_is_prebuilt[8 parts]; test_package.py::test_prebuild_takes_shapes_only, ::test_prebuild_checks_the_shapes_before_the_device, ::test_every_registered_program_has_a_kind_and_the_kinds_split_the_registry, ::test_the_router_kind_holds_the_k2_programs_the_split_and_the_fp32_forwards, ::test_an_unknown_template_has_no_kind; test_fast_plans.py::test_every_recorded_program_is_built_by_the_package. GPU: the 4 tests of test_prebuild_gpu.py, with T = 239 included; they passed on B200 and GB300 starting from an empty cache
9-2 Not batch-invariant: 2 / 16 router comparisons differed (logits max diff 3.4e-5), and 8 / 32 projection comparisons differed (output-group max diff 9.8e-4). The reviewer asked to document it or fix the split policy. Documented in the README row, the op module, ops.py and the entry-point docstrings, including the consequence for top-k vs packing. The split policy is unchanged. CPU: test_package.py::test_the_docs_state_that_the_results_are_not_batch_invariant. GPU: test_determinism_gpu.py::test_det_is_not_batch_invariant_but_every_token_count_is_within_budget[5 cases]. At T ∈ {1, 33, 257, 2049, 16384}, each T must be repeatable and within the rel-RMS budget against FP64. Router top-k must be equal on rows without a near tie. Cross-T bit differences are recorded, not asserted; the B200 counts are under Validation
9-3 Router backward memory was 1.41 GiB vs 1.32 GiB (Megatron-style) at T = 32768, H = 7168, judged "comparable, not a regression". No change. —
cleared Partial gradients were bit-consistent and FP64-correct, prebuild covered the listed T, and arena growth was 90 MiB. Kept. ROUTES adds input-gradient-only and weight-gradient-only routes. GPU: test_prebuild_gpu.py routes ...-input-gradient-only and ...-weight-gradient-only

Cross-stack items from the review of #1 that touch this layer:

  • The unconfigured default is now det, under one runtime rule.
  • The mid-call nvcc compiles named for this layer are resolved by the shape-only prebuild.
  • Argument ownership moved to the runtime. The op declares OP.validate, and the hidden variant flag is written back.
  • A one-GPU run now deselects the two-device module with -m "not multi_gpu" instead of skipping it at collection. Collecting the module touches no device.

Tests

All tests are in tests/cake/dense_gemm/.

  • CPU, torch-free: test_arguments.py (8).

  • CPU, with torch but no device:

    • test_package.py (7).
    • test_fast_plans.py (9). It covers the fast plans, including the known misses.
    • test_prebuild_plans.py (1). It checks that every planned program is in the prebuilt set, for both variants, over 8 parts: sm_100a with 148 / 152 SMs, sm_103a with 148 / 152 / 160, and sm_107a with 208 / 212 / 216. The token counts are:
      • 1 to 300;
      • every tile / split / window edge ±1 (32 to 65536);
      • 239, 1001, 2049, 10240, 16172, 16231, 20480 and 49152;
      • 24 seeded random counts up to 70000.
  • GPU:

    • test_batched_projection_gpu.py (6) and test_router_gpu.py (5).
    • test_fast_gpu.py (4), including test_fast_runs_det_under_torch_deterministic_algorithms and test_the_known_accuracy_misses_of_fast_are_measured_and_recorded. The latter runs 10 seeds at the 3 documented points. It runs only on 148-SM sm_103a parts and expects a failure where an element is outside the budget.
    • test_determinism_gpu.py (3): det bit for bit in one process and in a fresh one; det on the default stream and a side stream in turn; and the batch-invariance test above.
    • test_multi_gpu.py (2, multi_gpu): det gives the same bits on both devices, and fast keeps det's forward and input-gradient bits on both. The module's skip conditions are marks, so collecting it touches no device.
  • GPU, test_prebuild_gpu.py (4): one prebuild, then 20 ROUTES. Each route runs in det and fast under no_build() at 15 token counts: 1, 33, 64, 128, 239, 1001, 2049, 4097, 10240, 16384, 20480, and 4 seeded random counts in 2 to 20480. The routes are:

    • 6 projection shapes: (16,192,512,in_out), (16,512,256,out_in), (64,256,512,in_out), (8,4096,1024,out_in), (1,4096,1024,out_in) and (1,4096,32,out_in);
    • an input slice of a wider head slot together with a weight split of a parameter (both read in place);
    • a weight split, and a weight that is not 16-byte aligned (copied);
    • input-gradient-only and weight-gradient-only routes, and a side stream;
    • 4 routers (256 × 4096, 256 × 5120, 256 × 6144, 384 × 7168), a router with a [T, H] input, and a router on a side stream;
    • a router and a projection under torch.use_deterministic_algorithms(True).

    The other tests check four things: that one prebuild of the model shapes serves every route and token count; that a fresh process running both variants imports no flashinfer module; that the routes reach this device's fast programs; and that a second prebuild compiles nothing and returns the same libraries.

  • Budgets:

    • BF16 results: rel-RMS 3e-3, plus the det package's elementwise BF16 bar (atol = rtol = 1e-2).
    • FP32 logits: rel-RMS 5e-5, the package's bar.
    • Router FP32 weight gradient under fast: rel-Frobenius 2e-6.
    • Router top-k tie margin: 1e-4.
    • Each test records its measured errors (record_property).

Validation

CPU, at each commit

The image and commands are the same as for layer 8:

  • the full pytest tests run in a torch-free venv and with the image Python;
  • vendor_cake.py --check and pre-commit run --all-files, with CUDA hidden.

At the layer head, three more checks ran:

  • vendor_cake.py --check dense_gemm --upstream;
  • pytest -m 'not multi_gpu' tests/cake/dense_gemm;
  • pytest -m 'not multi_gpu' tests/cake/dense_gemm/test_multi_gpu.py, which must deselect the whole module.

The image is radixark/miles:dev-202611032236 (Python 3.12.16, torch 2.14.1+cu130, apache-tvm-ffi 0.1.11).

Commit Torch-free venv Image Python Other checks Verdict
ef381a0 (vendor) 465 passed, 98 skipped, 1 failed 1024 passed, 249 skipped, 1 failed vendor --check rc 0; pre-commit rc 0, no diff only test_vendored.py::test_every_vendored_op_is_registered fails, as the per-commit rule allows
20ccf43 (op) 471 passed, 100 skipped, 1 failed 1032 passed, 249 skipped, 1 failed vendor --check rc 0; pre-commit rc 0, no diff only test_op_contract.py::test_the_op_has_cpu_tests_and_a_prebuild_test[dense_gemm] fails, as allowed
9eead0b (tests) 487 passed, 123 skipped, 0 failed (rc 0) 1083 passed, 260 skipped, 0 failed (rc 0) vendor --check rc 0; pre-commit rc 0, no diff; --check dense_gemm --upstream rc 0; tests/cake/dense_gemm (image, CPU) 50 passed, 5 skipped, 6 deselected; test_multi_gpu.py with -m 'not multi_gpu': 6 deselected, nothing run (pytest rc 5, as expected) all green

vendor_cake.py --check --upstream at the top of the stack (layer 13, with the same ops.dense_gemm block and package trees) returned rc 0.

GPU

The GPU tests ran on the top of the stack (layer 13). Its dense_gemm files and its ops.dense_gemm block are byte-identical to this layer's. These are cake/dense_gemm.py, _impl/dense_gemm/, both _vendor/dense_projection_* trees and tests/cake/dense_gemm/.

The container was radixark/miles:dev-202611032236 (amd64 on B200, arm64 on GB300), with Python 3.12.16, torch 2.14.1+cu130, nvcc 13.0.88 (CUDA 13.0), apache-tvm-ffi 0.1.11, cuda-bindings 13.4.4 and pytest 9.2.0. Each op had its own JIT cache.

Commands:

  • Cold, then warm: PYTHONPATH=. python -m pytest -rfEs -p no:cacheprovider tests/cake/dense_gemm -m "not multi_gpu".
  • Multi-GPU: -m multi_gpu on 2 GPUs.
  • Prebuild first: as in layer 8.
Site GPU Driver Cold (empty cache) Warm Prebuild first, empty cache multi_gpu (2 GPUs)
B200 NVIDIA B200, sm_100a, 148 SMs 580.82.07 168 passed, 3 skipped, 290.8 s (97 libraries built; 8 ops compiled in parallel) 168 passed, 3 skipped, 160.6 s 4 passed, 115.4 s (96 libraries) 6 passed
GB300 NVIDIA GB300, sm_103a, 152 SMs 580.159.03 168 passed, 3 skipped, 238.7 s (prebuild test first; 101 libraries) 168 passed, 3 skipped, 89.5 s in the cold run: 4 passed 6 passed
R200 sm_107a — pending pending pending pending
  • Skips. The 3 skips on both sites are test_fast_gpu.py::test_the_known_accuracy_misses_of_fast_are_measured_and_recorded at the 3 documented points. These measurements run only on 148-SM sm_103a parts. A 148-SM sm_100a B200 and a 152-SM GB300 are not such parts, so the measurements have not been run in this round.
  • det bitwise. test_determinism_gpu.py::test_det_repeats_bit_for_bit_in_one_process_and_in_a_fresh_one, ::test_det_gives_the_same_bits_on_the_default_stream_and_a_side_stream_in_turn and test_router_gpu.py::test_repeated_router_launches_are_bitwise_identical passed on B200 and GB300, cold and warm.
  • Two devices. test_det_gives_the_same_bits_on_both_devices (3 cases) and test_fast_keeps_det_forward_and_input_gradient_bits_on_both_devices (3 cases) passed on both sites.
  • Prebuild. Starting from an empty cache, one prebuild of the model shapes served all 20 routes at 15 token counts, in det and fast, under no_build(). The routes reached this device's fast programs. A second prebuild compiled nothing. A fresh process ran both variants without importing flashinfer.
  • Deterministic algorithms. fast under torch.use_deterministic_algorithms(True) ran det (test_fast_runs_det_under_torch_deterministic_algorithms).
  • Measured errors against FP64 on B200. These come from the junit properties; cold and warm runs were identical. BF16 outputs and gradients had rel-RMS of at most 1.71e-3 (budget 3e-3). Router FP32 logits had rel-RMS of at most 7.92e-6 (budget 5e-5).
  • Batch invariance on B200. These counts were recorded, not asserted. Each count is the number of elements whose bits differ from the same tokens of the 16384-token call.
    • Projections: the three cases (16,192,512,in_out), (8,4096,1024,out_in) and (1,4096,32,out_in) had 0 differing output and input-gradient elements at T = 1, 33, 257 and 2049.
    • Router logits differed at every T, all within the 5e-5 budget. At T = 1 / 33 / 257 / 2049, 256 experts × 5120 had 252 / 8,371 / 65,182 / 519,702 differing elements, and 384 × 7168 had 382 / 12,597 / 98,058 / 781,731.
    • Router input gradients: 0 differing elements.
  • Pool. The op does not use the workspace pool, so no pool gate applies.
  • sm_107a (R200). Pending: no R200 run of this tree has been recorded yet.

Performance

The performance section does not fit in this description. It is posted in full in the first comment below, Performance (continued). Every number there is copied from FlashInfer pull request bodies; nothing was measured for this PR. It covers:

Notes

  • sm_107a: both packages ship sm_107a libraries (91 det, 11 fast) and need CUDA 13.4 to build them; an older toolkit raises RuntimeError. The CPU planner sweep covers sm_107a parts with 208, 212 and 216 SMs. No GPU run on sm_107a has been recorded for this tree yet.
  • perf(cake_dense_projection_gemm): restore the K1 weight-gradient routes on 216-SM sm_107a parts flashinfer-ai/flashinfer#8782, which restores the weight-gradient rows of 216-SM sm_107a parts, was merged after these packages were vendored. Its merge commit is f8bedda6038a3eff3b6953664c385c6b9eb0f1a3, merged on 2026-11-05. The message of commit 20ccf43 still calls it open.
  • Neither variant is batch-invariant, and the op is not CUDA-graph safe.
  • D_in, D_out, E and H must be multiples of 8.
  • fast is not a CLI choice yet. It has the known accuracy misses listed above, on shapes miles does not call. Those misses were not re-measured in this round, because no 148-SM sm_103a part was in the GPU runs.

🤖 Generated with Claude Code

@Zhichenzzz

Copy link
Copy Markdown
Collaborator

Review of layer 9 (dense_gemm) — suggest approve with nits

Confirmed on GPU

  • Mid-step JIT. After prebuild at T = 33/1001/2049, a sweep of 50 random T × 4 ops compiled 2 new libraries at T=239: glm5_qabs took 10.1 s and dsv4_ogroup 4.6 s. This needs a T-independent prebuild.

  • The ops are not batch-invariant. The same row in batches of T ∈ {1, 33, 257, 2049, 16384} changes bits in:

    • 2 of 16 router comparisons (logits max diff 3.4e-5);
    • 8 of 32 projection comparisons (dsv4 o-group out max diff 9.8e-4).

    Routing can therefore depend on packing, which matters for R3 and true-on-policy. Please document it, or fix the split policy.

  • Router backward memory. Peak extra at T=32768, H=7168 is 1.41 GiB (bf16 dx plus an FP32 [T, H]), vs 1.32 GiB for a Megatron-style torch backward. Comparable, not a regression.

Cleared on GPU

  • Partial grads. 12 cases are bit-consistent and FP64-correct.
  • Prebuild coverage. The listed T values are covered for MLA, DSV4 and the router.
  • Arena growth is 90 MiB over the sweep.

Verified on two GB300 nodes (sm_103a), CUDA 13.0, torch 2.14.1+cu130, at the stack head 2e422c50ca (this layer's files are unchanged there). The full suite passes (35 files, 809 tests). Each item above was checked with a dedicated GPU script; numbers are measured. sm_100a (GB200) paths were not exercised.

averyhNV and others added 26 commits November 4, 2026 23:14
Vendor FlashInfer's two Cake short-conv packages at FlashInfer
c5db8d3ccd3fa3842812280f7f208718dc0b8610 into miles_ext_kernels/cake/_vendor/, one level deep:

- cake_causal_conv1d_det_train_backend (flashinfer/cake_causal_conv1d_det_train_backend,
  87 files): det, fixed-order weight-gradient reduction, with its context-parallel route
- cake_causal_conv1d_train_backend (flashinfer/experimental/cake/cake_causal_conv1d_train_backend,
  75 files): fast, unordered weight-gradient reduction, with its context-parallel route

Both kcp.py modules read one environment variable, a private test hook of the package; the manifest
lists them in allow.env_read. Every file is byte for byte FlashInfer's.

Command (from the repository root, a FlashInfer clone next to it):

    python tools/vendor_cake.py fi --checkout ../flashinfer \
        --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 --op conv1d \
        --package flashinfer/cake_causal_conv1d_det_train_backend \
        --package flashinfer/experimental/cake/cake_causal_conv1d_train_backend \
        --allow-env-read cake_causal_conv1d_det_train_backend:kcp.py \
        --allow-env-read cake_causal_conv1d_train_backend:kcp.py

Co-Authored-By: Claude <noreply@anthropic.com>
Vendor FlashInfer's Cake packed causal GQA training attention package
(head dim 256, SM100a / SM103a / SM107a) from flashinfer-ai/flashinfer
c5db8d3ccd3fa3842812280f7f208718dc0b8610

    flashinfer/experimental/cake_gqa_train

byte for byte into miles_ext_kernels/cake/_vendor/cake_gqa_train/, and
record it as the gqa op's block of _vendor/VENDORED.json, with

    python tools/vendor_cake.py fi --checkout ../flashinfer \
        --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 --op gqa \
        --package flashinfer/experimental/cake_gqa_train

run in this repository with a FlashInfer checkout of that commit next to
it. The package has 21 files; their git blobs and its Cake revision
(7ab5de0c9d4e42a367ba3e9e218d98208f9bd8fa) are the ones it had before
this layout, from the same FlashInfer commit. Nothing outside _vendor/
changes; the op that loads the package comes in the next commit.

Co-Authored-By: Claude <noreply@anthropic.com>
The three FlashInfer Cake mHC training packages at c5db8d3ccd3fa3842812280f7f208718dc0b8610,
byte for byte, each one level under miles_ext_kernels/cake/_vendor/ and in the mhc block of
VENDORED.json:

- cake_mhc_det_train_backend (flashinfer/cake_mhc_det_train_backend): det, the forward and a
  backward whose parameter gradients are fixed-order reductions; 221 files.
- cake_mhc_train_backend (flashinfer/experimental/cake/cake_mhc_train_backend): fast, copies of
  the det forward programs and a backward that adds the dfn partials with FP32 atomics; 181 files.
- cake_mhc_det_train_head_backend (flashinfer/cake_mhc_det_train_head_backend): the model head's
  stream mean and its gradient; 12 files.

The package trees are the ones this op shipped before the layout change (same git trees). No
vendored module reads the environment or imports flashinfer, so the entries list no allow hooks.

Vendored by running, in this repository with a FlashInfer checkout of that commit next to it:

    python tools/vendor_cake.py fi --checkout ../flashinfer \
        --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 --op mhc \
        --package flashinfer/cake_mhc_det_train_backend \
        --package flashinfer/cake_mhc_det_train_head_backend \
        --package flashinfer/experimental/cake/cake_mhc_train_backend

Co-Authored-By: Claude <noreply@anthropic.com>
miles_ext_kernels.cake.mhc declares the op (torch-free): variants det and fast, selected by
--mhc-backend cake, --cake-mhc-variant, the three vendored packages, and the entry points
mhc_prepare, mhc_finish, mhc_head, is_supported, supported_hidden_sizes, HIDDEN_SIZES and STREAMS,
loaded from miles_ext_kernels.cake._impl.mhc on first use. Its validate rejects a --hidden-size
the kernels do not run.

miles_ext_kernels.cake._impl.mhc.ops routes a connection to the det or the fast package by the
variant configured when mhc_prepare runs; the handle carries it to mhc_finish and the backward.

Changes in behaviour:

- A process that never configured mhc runs det, and configure(deterministic_mode=True) without a
  variant runs det (it ran fast). An explicit fast in deterministic mode is rejected.
- The variant comes from the configuration only: torch.use_deterministic_algorithms(True) no longer
  turns a fast call into det; the fast backward then raises torch's alert for an operation without
  a deterministic implementation. The per-call variant argument is gone.
- prebuild("mhc", *, device=None, max_workers=None, hidden, sinkhorn_iters) takes both model
  parameters, without defaults, and builds every library of both variants and the head at that
  shape. A Sinkhorn count other than 20 runs run-time-count programs of the coefficient and
  token-owner passes, which a prebuild at the default count left to compile inside the first call.
  The epsilons are launch arguments and select no library.
- BF16 fn, scale and base (FSDP2 with param_dtype=torch.bfloat16) are cast to FP32 at the call,
  which is exact; their gradients return in BF16. They raised ValueError.
- Workspaces come from the runtime's workspace pool: each launch of the forward, the statistics
  replay and the backward takes a buffer when it runs (roles det-fwd, det-bwd, fast-fwd, fast-bwd,
  sized for the call's power-of-two token-count bucket), so they stay under
  --cake-workspace-pool-bytes and miles_ext_kernels.release() frees them. The packages' own
  per-device default workspaces, which lived for the whole process, are no longer allocated.
- release() drops the launch plans and the loaded libraries; the next call reloads the libraries
  from the cache.

The packages resolve every launch workspace through one function, which the op replaces with the
pooled one. The kernels' work tickets, the first 256 bytes of every workspace layout, must be zero
before a launch, and the content of a workspace is undefined: every launch zeroes them first (256
bytes on the stream of the call), from the pool or not.

Co-Authored-By: Claude <noreply@anthropic.com>
miles_ext_kernels.cake.gqa declares the op (torch-free) and
miles_ext_kernels.cake._impl.gqa runs FlashInfer's cake_gqa_train
package through the runtime's loader: varlen_gqa_attention, packed
causal GQA attention over variable-length segments, BF16 q
[T_q, H_q, 256] and k / v [T_k, H_kv, 256] with an even number of query
heads per key/value head, forward and backward, on SM100a, SM103a and
SM107a. miles selects it with --full-attention-backend cake.

Variants. The op ships one variant, fast: the output, the LSE and dK /
dV repeat bit for bit run to run, dQ is summed with FP32 atomic adds.
There is no det variant, so the runtime rejects the op under
--deterministic-mode, and a process configured with
deterministic_mode=True raises at the call. A process that never
configured the op runs fast.

max_seqlen_k is now a required keyword-only host int and a documented
hard contract: at least every segment's key count, as FlashAttention's.
It sizes the backward grid, so a smaller value leaves key blocks out of
the backward and gives wrong dQ / dK / dV without an error, while the
output is unaffected. The kernels never read the boundaries on the
host; the new option --cake-gqa-check-seqlens (configure
cake_gqa_check_seqlens=True) checks at every call that cu_seqlens_q /
cu_seqlens_k start at 0, never decrease and end at the token counts,
and that no segment has more keys than max_seqlen_k, at the price of
one device synchronization per call. It is off by default.

The implicit None bound is gone: it sized the backward grid by T_k,
which is correct but launches CTA pairs that only retire, and made the
backward slower the more segments a call packs. A caller without the
longest segment can still pass T_k explicitly. The output, the LSE and
dK / dV do not depend on the bound given.

check_supported(num_heads_q, num_heads_kv, head_dim=256) and
is_supported(...) tell, without torch, whether the kernels run a head
layout. validate_args applies them to the heads one tensor-parallel
rank of miles holds (num_attention_heads / TP query heads over
max(num_query_groups, TP) / TP key/value heads, head dim from
--kv-channels or hidden_size / num_attention_heads), so a layout such as
24 query heads over 4 groups at TP 8 (3 query heads over 1 key/value
head per rank) is rejected at startup rather than at the first forward.

Memory. Each backward takes its scratch from the runtime's workspace
pool on the current stream, and each forward zeroes a fresh tile
counter; nothing goes through the package's per-stream default
workspaces, which grow to the largest call and live as long as the
process. release() therefore frees all of it, and calls on two streams
never share a buffer. prebuild("gqa") builds the four libraries of the
forward and the BF16 dK / dV backward for the device's architecture; it
takes no model parameter.

The op's names move from miles_ext_kernels.gqa:
SUPPORTED_CAPABILITIES, device_arch and prebuild are the runtime's
ARCH_BY_CAPABILITY, device_arch and prebuild("gqa"). The README lists
the op in its Ops and Prebuild tables.

Co-Authored-By: Claude <noreply@anthropic.com>
tests/cake/gqa/, CPU:

- test_arguments.py (torch-free): the op registers one variant and no
  variant flag; --cake-gqa-check-seqlens is off by default and needs
  --full-attention-backend cake; --deterministic-mode (also through the
  typed configuration) rejects the op; check_supported / is_supported
  on head layouts; validate_args on the heads of one tensor-parallel
  rank (24 query heads over 4 groups at TP 8 and 8 over 8 at TP 8
  rejected, 16 over 4 at TP 8 accepted, head dim from --kv-channels or
  hidden_size / num_attention_heads); configure and option.
- test_gqa_cpu_plan.py (torch-free): every stage the entry point
  launches has one library per architecture, which the oldest toolkit
  the runtime accepts for it (MIN_CUDA) builds; the grids cover every
  tile; the backward grid covers the key blocks of max_seqlen_k and no
  more; the scratch layout.
- test_gqa_host.py (torch on the CPU): deterministic mode raises at the
  call and configure() lets the next call through; max_seqlen_k is a
  required host int; check_seqlens accepts the contract and rejects a
  bound below the longest segment and malformed boundaries; the option
  runs the check before any launch; with the pool on, the backward
  scratch is one pooled buffer in the package's layout, and at the
  default bound, 0, each backward takes a buffer of its own that the
  pool does not keep; prebuild takes no model parameter,
  raises NotImplementedError on another GPU and RuntimeError when the
  toolkit is older than CUDA 13.4 on SM107a.

GPU (skip without SM100a / SM103a / SM107a):

- test_gqa_gpu.py: FlashInfer's suite layouts and miles' head pairs
  against an FP64 reference (assert_rel_rms) within FlashInfer's budget
  capped at twice the error measured on a B200, noted next to it:
  FlashInfer's budget for out, dQ, dK and dV (1.08 to 1.12 times the
  measured error), twice the measured error for the LSE, whose errors
  are at FP32 rounding;
  out, LSE, dK and dV bit for bit run to run in one process and against
  a new process, for any valid max_seqlen_k (exact, larger, T_k), on
  alternating default and side streams, and pooled against the pool
  turned off;
  dQ within its FP32 reordering; the check rejects max_seqlen_k 1024,
  1920 and 2047 on 2048-key segments and changes no bit when valid;
  document isolation; checkpointing; no device synchronization in
  steady state (and one with the check); with the pool on, the pooled
  scratch and release(); rejected operands; two GPUs give the same bits
  (multi_gpu).
- test_prebuild_gpu.py: ROUTES for every call-time path (exact and T_k
  max_seqlen_k, groups 2 / 4 / 6 / 8, keys fewer than queries, empty
  segments, long keys, one segment, no query, no key, a softmax scale,
  the check option, a side stream, the pool on) after one prebuild
  under no_build(); prebuild returns the four libraries; a fresh
  process loads only those and never imports flashinfer.

Co-Authored-By: Claude <noreply@anthropic.com>
CPU (no GPU): the arguments and the variant rules through the hooks miles calls, deterministic mode
without a variant included; the vendored packages against the routes, constants and hidden sizes
the op relies on; what prebuild builds per Sinkhorn count, checked against the programs the host
rules select at every token count; BF16 parameters reaching the kernels as FP32; the variant read
at each call; with the pool on, every launch workspace taken from the pool, its tickets zeroed at
every launch, and at the default bound, 0, none kept; what release drops.

GPU: det and fast against an FP64 reference at every hidden size the device serves, the fast
outputs against det's bits, another Sinkhorn count and other epsilons, BF16 parameters, det bit for
bit in a new process, on a side stream and on a second GPU (multi_gpu), with the pool on: the bits
under any pool bound, from a buffer a larger call sized and with interleaved connections, release;
the head, and test_prebuild_gpu.py: ROUTES over every call-time switch at the Sinkhorn count 7,
checked from an empty library cache.

Co-Authored-By: Claude <noreply@anthropic.com>
miles_ext_kernels/cake/conv1d.py declares the op: variants det and fast, selected by
--linear-attention-conv-backend cake, the option --cake-conv1d-silu-grad-preactivation, the two
vendored packages and the entry points causal_conv1d, causal_conv1d_kcp, causal_conv1d_kcp_multi and
check_supported. miles_ext_kernels/cake/_impl/conv1d/ implements them on the runtime's loader, config
and prebuild. The op is appended to OPS and to the README's Ops and Prebuild tables.

Changes in behavior:

- A process that never configured the op runs det; it ran fast. miles' command line still defaults
  to fast: validate_args resolves --cake-conv1d-variant to fast and writes it back for configure.
- Each call reads its variant and SiLU' policy when it runs; the backward runs those of its forward.
- The entry points take no variant argument: only configure selects the variant, so a process
  configured with deterministic_mode runs det on every call. A call could pass variant="fast" before.
- --cake-conv1d-silu-grad-preactivation defaults to bf16; it had no default. Without
  --linear-attention-conv-backend cake, an explicit bf16 is now accepted (it raised before); any
  other value still requires the cake backend.
- causal_conv1d without cu_seqlens builds the prefix sums of its one document on the device. It copied
  them from pageable host memory before, which synchronized the host on every call.
- check_supported(..., has_bias=True) raises: the entry points take no bias, so a biased layer
  would lose its bias.
- prebuild("conv1d", weight_dtype=...) builds every library of both variants, both SiLU' policies and
  the context-parallel route. It replaces prebuild(dtype) and prebuild_kcp().
- release() drops the packages' host plans.

Co-Authored-By: Claude <noreply@anthropic.com>
Move the conv1d tests to tests/cake/conv1d/ and run them through the runtime (configure, variant,
prebuild, release).

- CPU: the arguments and the configuration (det unconfigured, fast from miles' command line, an
  explicit default SiLU' policy without the cake backend), the variant and SiLU' policy read at each
  call, no entry point taking a variant (det on every call under deterministic_mode), the
  one-document cu_seqlens built on the device, the rejected bias, prebuild's arguments and its
  libraries against the packages' program selection on sm_100a, sm_103a and sm_107a, release, the
  vendored packages and their loading without flashinfer.
- GPU: the accuracy and determinism tests as before, each variant selected with configure, and new
  ones: det bitwise in a new process, cu_seqlens=None without a host synchronization or a wait for
  queued work, calls on a side stream interleaved with the default stream, an unconfigured process
  running det, and test_prebuild_gpu.py, whose ROUTES cover every call-time path and run under
  no_build() after one prebuild, for each taps dtype in a new process with an empty library cache
  of its own.
- The multi-GPU context-parallel test is marked multi_gpu.

Co-Authored-By: Claude <noreply@anthropic.com>
The four Cake KDA training packages of FlashInfer 7a13d7ba9b75c9bb9c02452df36e3f34417b5d57, byte for byte, one
level deep under miles_ext_kernels/cake/_vendor/ and recorded under ops.kda of VENDORED.json:

- cake_kda_det_train (flashinfer/experimental/cake_kda_det_train): det, single rank.
- cake_kda_det_train_kcp (flashinfer/experimental/cake_kda_det_train_kcp): det, context parallel.
- cake_kda_train (flashinfer/experimental/cake/cake_kda_train): fast (deterministic=False), single rank.
- cake_kda_train_kcp (flashinfer/experimental/cake_kda_train_kcp): fast, context parallel.

One FlashInfer commit for the family, as the context-parallel packages import their single-rank siblings
(..cake_kda_det_train and ..cake.cake_kda_train; the runtime resolves the second to _vendor.cake_kda_train).
7a13d7ba rather than c5db8d3c: it carries FlashInfer #8629, a correctness fix of the det single-rank package (with
the in-kernel q / k L2 norm, its replay backward over a batched arena, KdaDetArena(batch_tokens=...), read up to 63
rows of the next document group before they were written), and the other changes of these packages since c5db8d3c:
#8552, #8620, #8625, #8630 and #8652 (performance; #8625 runs the fused L2 norm's GLM intra stage with another
program on SM103a, #8620 sums the fast packages' raw-gate parameter gradients in FP64 on SM100a / SM103a). The entry
points of this repository pass no arena, so #8629 does not change their results. The later commits up to FlashInfer
main change none of the four directories.

cake_kda_train/cake_backend.py keeps two FlashInfer hooks, listed under allow: it reads CAKE_KDA_CKPT_HANDOFF
(the checkpoint hand-off override of its chunk_kda), and _jit_nvcc_version imports flashinfer.jit.cpp_ext inside
its body to name the toolkit of the libraries; the kda implementation rebinds that function to the runtime loader's
nvcc when it is imported.

Vendored with, from the root of this checkout next to a FlashInfer clone at ../flashinfer:

    python tools/vendor_cake.py fi --checkout ../flashinfer --commit 7a13d7ba9b75c9bb9c02452df36e3f34417b5d57 \
        --op kda \
        --package flashinfer/experimental/cake_kda_det_train \
        --package flashinfer/experimental/cake_kda_det_train_kcp \
        --package flashinfer/experimental/cake/cake_kda_train \
        --package flashinfer/experimental/cake_kda_train_kcp \
        --allow-env-read cake_kda_train:cake_backend.py \
        --allow-flashinfer-import cake_kda_train:cake_backend.py:_jit_nvcc_version

Co-Authored-By: Claude <noreply@anthropic.com>
FlashInfer c5db8d3ccd3fa3842812280f7f208718dc0b8610: the SiLU-gated RMSNorm training packages, deterministic
(cake_rmsnorm_silu_gate_det_train_backend, 14 files) and non-deterministic (cake_rmsnorm_silu_gate_train_backend,
30 files), each one level deep under miles_ext_kernels/cake/_vendor/, and their block ops.gated_rmsnorm in
_vendor/VENDORED.json. Written from a FlashInfer checkout by:

    python tools/vendor_cake.py fi --checkout ../flashinfer --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 \
        --op gated_rmsnorm --package flashinfer/cake_rmsnorm_silu_gate_det_train_backend \
        --package flashinfer/cake_rmsnorm_silu_gate_train_backend

Co-Authored-By: Claude <noreply@anthropic.com>
Vendor FlashInfer's Cake RMSNorm training packages of flashinfer-ai/flashinfer
eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3 into miles_ext_kernels/cake/_vendor/, byte for byte, as the packages of the
op rmsnorm in _vendor/VENDORED.json:

- cake_rmsnorm_det_train_backend (flashinfer/cake_rmsnorm_det_train_backend, 582 files): the deterministic forward
  and backward.
- cake_rmsnorm_train_backend (flashinfer/cake_rmsnorm_train_backend, 581 files): the backward with a non-deterministic
  weight gradient in some token-count windows; it runs its deterministic sibling for deterministic=True, so the two
  packages move together.

eb3eb68d includes flashinfer-ai/flashinfer#8512: on SM100a and SM103a the deterministic backward of the smallest token
counts sums each weight-gradient column in a fixed row order, so its bits no longer depend on the SM count. The files
and their blobs are the ones the previous rmsnorm package vendored from the same commit.

    python tools/vendor_cake.py fi --checkout ../flashinfer --commit eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3 \
        --op rmsnorm --package flashinfer/cake_rmsnorm_det_train_backend \
        --package flashinfer/cake_rmsnorm_train_backend

Co-Authored-By: Claude <noreply@anthropic.com>
Vendor the DSA kernel packages into miles_ext_kernels/cake/_vendor/ with
tools/vendor_cake.py, byte for byte, and record them under ops.dsa of
VENDORED.json:

- cake_dsa_train: FlashInfer's DSA training kernels for 64 local query
  heads (the fast variant: dKV and dk_rope accumulate in FP32 with
  red.global).
- cake_dsa_indexer_det: FlashInfer's lightning-indexer scores and
  deterministic exact top-k.
- cake_dsa_sparse_det: the DSA det version, the generated sparse
  attention forward / backward and lightning-indexer logits / backward
  kernels of radixark/miles#3599 (86 sources and registry.json; no
  floating-point atomics, key-side gradients reduced from FP32 partials
  in a fixed 4-way tree). The tool writes the TVM-FFI host surface in
  place of the upstream torch-extension one and a build.json generated
  from the registry, so these kernels build with nvcc and load with
  tvm-ffi through the same loader as the FlashInfer packages; SM107a
  compiles the SM103a records.

Both FlashInfer packages come from c5db8d3, the commit these kernels
were validated at: cake_dsa_train is byte for byte the same at 7a13d7ba,
and cake_dsa_indexer_det differs there only by a later performance
change (flashinfer-ai/flashinfer#8691).

Commands (run from the repository root, with a FlashInfer checkout at
../flashinfer and a miles checkout at ../miles):

    python tools/vendor_cake.py fi --checkout ../flashinfer \
        --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 --op dsa \
        --package flashinfer/experimental/cake_dsa_train \
        --package flashinfer/experimental/cake_dsa_indexer_det
    python tools/vendor_cake.py miles-pr --checkout ../miles \
        --pr 3599 --commit a1b763c82360afeb00c551aeb7ea93b30870e5bf \
        --op dsa --name cake_dsa_sparse_det \
        --path miles/kernels/attention/dsa/cake \
        --include csrc --include registry.json \
        --exclude csrc/cake_host_shim.h --arch-from sm_107a=sm_103a

Co-Authored-By: Claude <noreply@anthropic.com>
The KDA op of the cake namespace: miles_ext_kernels/cake/kda.py declares it (OpSpec, selected by --kda-backend
cake), miles_ext_kernels/cake/_impl/kda/ implements it on the four vendored packages, OPS gains "kda" and the
README its Ops and Prebuild rows. The entry points are those of the kda family before this layout: chunk_kda,
chunk_kda_fwd, chunk_kda_bwd, and partition_kcp, build_context_kcp, chunk_kda_kcp, chunk_kda_fwd_kcp,
chunk_kda_bwd_kcp, reduce_gate_param_grads_kcp. Each call reads config.variant("kda") when it runs: det unless
configured.

Variants. det (cake_kda_det_train, cake_kda_det_train_kcp) repeats bit for bit. fast (cake_kda_train,
cake_kda_train_kcp, deterministic=False) runs the S1 gate-gradient programs on some backward rows of the activated
gate mode and, in the raw gate mode on SM100a / SM103a, reduces the gate parameters' gradients in FP64 in no fixed
order. fast has no in-kernel q / k L2 norm: such a call runs det and logs it once (config.warn_runs_det). Every
miles call takes that L2 norm, so cli_variants is ("det",): --cake-kda-variant is registered hidden, validate_args
writes det back for every process, and a fast choice from miles, which would change no program, is refused by
argparse. Direct callers configure either variant. A backward reads its saved set's package (a saved set is read by
the package that wrote it); saved=None replays the forward in the package of the call. Every call hands its package
use_qk_l2norm_in_kernel as given, so the backward of a fast saved set asked for the L2 norm, a mode its forward did
not run, is the fast package's ValueError rather than a gradient without the norm. Partitions and contexts hold no
variant.

Operands as miles hands them over. beta may be FP32 (miles gives flash-linear-attention the FP32 sigmoid) and v a
column slice of the packed projection: chunk_kda* round beta to BF16 and make v dense before the packages' checks,
inside the autograd graph, and return dbeta in beta's dtype. Before, both raised ValueError in the packages.

Prebuild. prebuild(num_heads=...) checks 1 <= num_heads <= 64 (an int, not a bool) before anything else and then
builds every library of the four packages for the device's architecture (188 / 186 / 189 on SM100a / SM103a /
SM107a): the record a launch takes depends on its token and document counts, its modes, its path and, under context
parallelism, the rank's shard, so no smaller set holds for every call of a layer. This replaces the family's
route_libraries probe, which planned a fixed list of problems and dropped every problem it could not plan (an
unknown gate mode gave an empty list, not an error). release() drops the packages' per-stream entry pools and the
captured graphs of a loaded context-parallel package.

The fast package's _jit_nvcc_version hook (its checkpoint hand-off keys row records on the toolkit) is rebound to
the runtime loader's nvcc when the implementation is imported: flashinfer is never imported. --kda-backend cake
rejects context parallelism (--context-parallel-size > 1) and a layout other than thd at startup, as before.

Co-Authored-By: Claude <noreply@anthropic.com>
The output norm of miles' GDN layers, RMSNorm(x) * weight * SiLU(z) over rows of 128, becomes the op
miles_ext_kernels.cake.gated_rmsnorm: OP declares its two variants, the miles selector
--linear-attention-output-norm-backend cake, its two vendored packages and the entry points rmsnorm_silu_gate and
check_supported; miles_ext_kernels.cake._impl.gated_rmsnorm implements them.

- A call reads the configured variant when its forward runs, and its backward runs the same package. A process that
  never configured the op runs det (it ran fast); validate_args still resolves miles' command-line default to fast,
  and deterministic mode resolves det and rejects an explicit fast. While torch.use_deterministic_algorithms(True)
  is in effect, which the non-deterministic backward refuses, fast runs det and logs it once.
- prebuild(device=None, max_workers=None), reached through cake.prebuild("gated_rmsnorm"), builds every library of
  both packages on the device's architecture, every row-count route of each, so no call compiles afterwards.
  release() has nothing to drop.
- rmsnorm_silu_gate documents when a gate that is a column slice of a wider projection is read in place, and that
  the weight gradient sums over the rows of one call: under tensor parallelism the caller all-reduces it over the
  tensor-parallel group.
- registry.OPS lists gated_rmsnorm; the README's Ops and Prebuild tables list its rows.

Gone with the old package: the module-level VARIANT, the op's own arguments module and configure, the per-call
variant keyword, and SUPPORTED_CAPABILITIES (miles_ext_kernels.cake.ARCH_BY_CAPABILITY names the same devices).

Co-Authored-By: Claude <noreply@anthropic.com>
miles_ext_kernels.cake.rmsnorm declares the op: selected by --linear-attention-input-norm cake, variants det and fast
(--cake-rmsnorm-variant), the two vendored packages. miles_ext_kernels.cake._impl.rmsnorm implements it:
ZeroCenteredRMSNorm, zero_centered_rmsnorm, check_supported, supported_hidden_sizes, prebuild(hidden_size=...) and
release. The op is appended to OPS and to the README's Ops and Prebuild tables.

Changes from the previous rmsnorm package:

- A call reads the configured variant when it runs (det in a process that never configured the op), and its backward
  runs the variant of its forward. The module no longer reads or pins a variant at construction; its variant=
  argument is gone, configure is the one switch.
- Constructing the module compiles nothing. prebuild("rmsnorm", hidden_size=H) builds every library of both variants
  for H (forward and backward, every token-count window), so a variant configured after a module was built runs at
  its next call without a build.
- The weight is allocated where torch's factory functions allocate (a with torch.device("meta") block,
  torch.set_default_device), else on the current CUDA device; device="meta" builds without touching CUDA. The device is
  checked when it is a CUDA device; a hidden size the kernels do not serve is refused on any device.
- reset_parameters() zeroes the weight, also after to_empty().
- sequence_parallel= (default None) is set on the weight when given, at construction and by reset_parameters, as
  Transformer Engine's norms set it, so Megatron all-reduces its gradient across the tensor-parallel ranks under
  sequence parallelism. None leaves the attribute to the caller: reset_parameters keeps a mark set on the weight.
- An unsupported device raises NotImplementedError (was ValueError), as the runtime's device_arch does.
- The package's configure, prebuild(hidden, device, variant=) and SUPPORTED_CAPABILITIES are replaced by
  miles_ext_kernels.configure, miles_ext_kernels.cake.prebuild("rmsnorm", hidden_size=...) and
  miles_ext_kernels.cake.ARCH_BY_CAPABILITY.

fast still runs det, logged once, under torch.use_deterministic_algorithms(True), which the non-deterministic package
refuses.

On SM107a, hidden size 4096 launches two programs above the 232448 B shared-memory opt-in (271360 B, CUDA 13.4's
oversized shared-memory mode): det's backward reduce above 4096 tokens and its copy in fast's backward from 4097 to
24064 tokens. A process whose CUDA runtime is older than 13.4, such as the 13.0 runtime wheel a cu130 torch installs,
rejects that mode at the launch ("cudaFuncSetAttribute failed: invalid argument"), and the error it leaves pending
fails the next CUDA call too. The op now checks the loaded runtime (libcudart of the toolkit's major) and the
driver's oversized ceiling for every program of the width that sets the mode, at construction on a CUDA device,
check_supported, prebuild and the first call of a width on a device, and raises RuntimeError naming the programs,
the runtime and the remedy before any launch. Other widths and SM100a / SM103a launch no such program. The README's
Ops and Prebuild rows state the limit.

Co-Authored-By: Claude <noreply@anthropic.com>
Add the DSA op: miles_ext_kernels/cake/dsa.py (its OpSpec, variants det
and fast, --cake-dsa-bwd-workspace-bytes and --cake-dsa-lazy-build) and
its implementation under miles_ext_kernels/cake/_impl/dsa/:

- ops.py: sparse_attention, the lightning indexer, the det forward and
  backward entry points, prebuild and release. The variant is read when
  a call runs; a call the fast variant does not serve (other than 64
  heads, a sink, another d_qk) runs det and logs it once per reason.
- native.py: the DSA det version (cake_dsa_sparse_det) built and loaded
  through the runtime loader, one library per stage and architecture.
  Every argument binds by its exported name; a TMA source binds as a
  tensor whose descriptor the stage's launcher encodes on the host.
- kernels.py, reduction.py, layouts.py, indexer_topk.py: the det host
  side, the fixed-order reduction, the layout adapters and the det
  indexer top-k (cake_dsa_indexer_det).

Register the op in OPS and add its rows to the README's Ops and
Prebuild tables.

Behaviour, compared with the previous layout:

- sm_scale is required: the d_qk ** -0.5 default is not the scale of
  the absorbed MLA these kernels serve.
- A call takes no variant argument; configure() selects it.
- A batch-local bshd index outside its batch element's keys is an
  empty slot; it no longer reads the next element's keys. So is an
  indexer index past the keys.
- --cake-dsa-bwd-workspace-bytes bounds the det backward's partials
  and the index scratch of their sort (one sort of every chunk when it
  fits, else one per chunk); the chunking, and with it the det bits,
  follows from it.
- The fast backward receives the packed segment count (num_segments;
  the batch for bshd) to plan its key passes; the bits do not change.
- A call without query rows returns empty outputs and zero gradients,
  forward and backward, attention and indexer (and the indexer's top-k
  selection, with or without key rows).
- The indexer serves 8, 16, 32 or 64 heads of dimension 128 in its
  logits, its backward and the selected scores alike; 24 is rejected
  before any launch.
- A selected key with a -inf score (outside its query's window) gets
  no gradient, even a NaN one: the backward sees it as an empty slot.

Co-Authored-By: Claude <noreply@anthropic.com>
tests/cake/kda/: the kda family's tests on the cake layout, with per-tensor accuracy budgets.

- test_arguments.py (torch-free): the spec, the hidden --cake-kda-variant (det written back, fast refused), the
  settings validate_args rejects, configure for direct callers, an import without torch.
- test_kda_variant.py (CPU, torch): the package of every entry point under det and fast, the in-kernel L2 norm call
  (logged once), the backward of a saved set of either package (a fast set asked for the in-kernel L2 norm is the
  fast package's ValueError), the beta / v seam and its gradients, the head-count
  check before any package or build runs, the fast toolkit hook, that fast plans programs det does not ship (S1 in
  the activated gate mode on every architecture, R1c in the raw gate mode on SM100a / SM103a), that prebuild builds
  the same 188 / 186 / 189 libraries for every head count, and release().
- test_kda_gpu.py: det in miles' call (q / k / v column slices, raw gate, in-kernel L2 norm) against the exact FP64
  recurrence, every output and gradient within twice the largest error measured (_budgets.py, one B200 and one GB300):
  the token tensors in relative RMS error (a 1.02 scale of out, dq, dk, dv or dbeta fails it), dA_log and ddt_bias
  in error over the root sum of squares of the terms of their exact sums, free of the shape's cancellation;
  bit-identical repeats, packed views and an FP32 beta against contiguous BF16 operands, autograd / saved set /
  replay / checkpoint bits, default and side stream bits, a fresh process's bits, and flash-linear-attention parity
  on every shape (each tensor within twice fla's error on the same inputs).
- test_kda_fast_gpu.py: fast against the same reference within twice det's error on the same inputs (no bit checks:
  fast promises none), its forward / backward split and checkpointed node, saved sets of either package, and the
  in-kernel L2 norm call that runs det giving det's bits.
- test_kda_kcp_gpu.py: the context-parallel entry points on 2 and 4 GPUs, marked multi_gpu, so a one-GPU run
  deselects it (-m "not multi_gpu") instead of reporting 20 skips; with fewer GPUs than a group needs it skips with
  that count.
- test_prebuild_gpu.py: ROUTES over every call-time switch (gate and L2 norm modes, packed views and FP32 beta,
  explicit forward and backward, replay, both checkpoints, the segmented route with saved states, 1 to 64 heads,
  more than 256 documents, a token sweep with empty documents, the context-parallel entry points on a one-rank
  partition with host cu_seqlens), checked after one prebuild with no build allowed.

The elementwise gate the family used (atol = rtol = 1e-2 on every tensor) let a 2 % error of out, dv or dbeta pass;
the relative RMS budgets do not, and the elementwise gate is not carried over. dg is exempt from the 1.02 detection
check, as its own error is about 1 %.

dA_log and ddt_bias are sums over every token of a head. Their relative error is set by how far the terms cancel,
which varies with the shape (on one B200 and one GB300, 0.8 % to 29 % for dA_log and 4.1 % to 6.3 % for ddt_bias),
and dA_log has one value per head, so a relative budget per shape is loose at one head. Over the root sum of squares
of the terms (_reference.gate_param_terms) the error is 0.6 % to 3.3 % for dA_log and 4.7 % to 6.3 % for ddt_bias,
its largest at one head; fla's on the same inputs is 1.2 % to 2.8 % and 4.6 % to 5.6 % (the fla parity rows print
both). These budgets are twice the largest.

Co-Authored-By: Claude <noreply@anthropic.com>
tests/cake/gated_rmsnorm/:

- test_arguments.py (CPU): the declaration and its variant flag; an unconfigured process and an unconfigured call
  run det; miles' resolution (fast by default, det in deterministic mode, explicit values kept, fast rejected in
  deterministic mode, the flag rejected without the backend); the variant read at every call and det under torch's
  deterministic algorithms, with the packages replaced by recorders; check_supported's rejections before any device
  query; import and configuration without torch.
- inputs.py: miles' GDN layout. The gate is a column slice of the fused [q | k | v | z | b | a] projection, whose
  width (2 * Hk * 128 + 2 * Hv * 128 + 2 * Hv, padded to a multiple of 8) is the gate's token stride: the kernels
  read the gate in place when it is a multiple of 16 elements and copy it otherwise. Every forward + backward
  checks the documented result types: y BF16 of x's shape, every gradient BF16, dw of shape [128].
- test_gated_rmsnorm_gpu.py: every value-head count (12, 16, 32, 48) with the gate read in place and copied, the
  path asserted, among them 16 key / 32 value heads (width 12352, in place) and its twin with 8 more columns (12360,
  copied), at 1 to 32771 tokens: y, dx and dz within one BF16 ulp of FP64 and dw within det's bound, every gate
  failing a 2 % error; det bitwise three times in one process and once in a new process; a contiguous gate the same
  bits; fast's y, dx and dz det's bits. Also calls alternating between the default and a side stream,
  checkpointing, accumulation, a tensor-parallel head split (each rank's dw covers its own heads, their sum is the
  layer's) and the toolkit check.
- test_prebuild_gpu.py: ROUTES (both gate paths at every head count, a 2-D input, a token sweep that reaches every
  row-count module of the non-deterministic package, checked against the device's SM count, a weight without a
  gradient, torch's deterministic algorithms) under check_prebuild; prebuild's libraries and idempotence; configure
  after a call switches the next call without a build; no flashinfer module is imported.

Co-Authored-By: Claude <noreply@anthropic.com>
tests/cake/rmsnorm/:

- test_arguments.py (torch-free): registration, variant resolution with and without miles' deterministic mode,
  rejections, configure, det in an unconfigured process.
- test_module.py (CPU and meta tensors, the kernels replaced by a recorder): construction on the meta device,
  reset_parameters after to_empty, the sequence_parallel attribute (none by default; set from the argument; a mark the
  caller set kept by reset_parameters) and Megatron's selection rule, refusal of widths the kernels do not serve,
  construction without a build, the variant each call reads, fast running det under torch's deterministic algorithms,
  the arguments of prebuild (its refusal without a CUDA device in a child process that sees none, so it runs on GPU
  nodes too); the programs that launch in the oversized shared-memory mode (SM107a at 4096 only) and
  the refusal of their width, with the CUDA runtime and driver facts replaced (runtime 13.0, none loaded, no oversized
  ceiling, a ceiling too small).
- test_rmsnorm_gpu.py: accuracy against FP64 at FlashInfer's budgets; det bit for bit three times in one process,
  once in a fresh process, at the smallest token counts and on a side stream (hidden sizes 4096 and 5120); fast's y and
  dx against det; strided rows, checkpointing, no tokens, a module built on the meta device, the toolkit check, a
  backward that runs its forward's variant. A width this process cannot launch (SM107a at 4096 without a CUDA 13.4
  runtime) is refused instead of run (_refusal.py); the refusal holds exactly where a fresh process with the check
  disabled fails to launch the width, and with the runtime taken as 13.0 SM107a refuses 4096 at check_supported,
  prebuild, construction and the first call, while every other width and architecture runs.
- test_prebuild_gpu.py: ROUTES (every token-count window of both packages, [T, H] input, strided rows, FP32 weight, no
  tokens, checkpointing, torch's deterministic algorithms) for every width after one prebuild, with no build allowed
  (prebuild of a width this process cannot launch raises its refusal);
  a variant configured after construction runs without a build; no flashinfer module is imported.
- test_sequence_parallel_gpu.py (multi_gpu, 2 GPUs): under tensor and sequence parallelism the all-reduce rule selects
  the weight, and the ranks keep one bit-identical weight that follows the single-GPU run.

Co-Authored-By: Claude <noreply@anthropic.com>
Add tests/cake/dsa/:

- test_arguments.py, test_package.py, test_dispatch.py and
  test_indexer_topk.py (CPU): the arguments and configuration, the
  three packages and the CUDA 13.4 floor of SM107a whatever a package
  records, the stage selection, the argument binding of every det
  stage (TMA sources included), the calls fast serves, the required
  sm_scale, the segment count, the index ranges, the det backward's
  memory plan, calls without query rows (the indexer forward and
  backward and its top-k selection included), the indexer head counts
  and the gradient mask of the selected scores.
- test_dsa_gpu.py and test_indexer_topk_gpu.py: accuracy against FP32
  and FP64 references at every head count from 4 to 128, det bits
  repeated in one process, in a fresh process, on a side stream and on
  a second GPU, the det backward within its workspace budget, and the
  reviewed limits on a GPU (batch-local and packed indices outside the
  keys are empty slots). The two-GPU tests are marked multi_gpu.
- test_prebuild_gpu.py: ROUTES over every call-time path (each head
  count from 4 to 128, 24 included; RoPE and NoPE; with and without a
  sink; thd, bshd and sbhd; no query rows; the indexer and its det
  top-k), each group checked after one prebuild.

Co-Authored-By: Claude <noreply@anthropic.com>
FlashInfer c5db8d3ccd3fa3842812280f7f208718dc0b8610: the dense projection GEMM training packages, deterministic
(dense_projection_det_gemm, 190 files) and non-deterministic (dense_projection_gemm, 30 files), each one level deep
under miles_ext_kernels/cake/_vendor/, and their block ops.dense_gemm in _vendor/VENDORED.json. Both come from one
commit: the non-deterministic package runs its deterministic rows on its sibling. Written from a FlashInfer checkout
by:

    python tools/vendor_cake.py fi --checkout ../flashinfer --commit c5db8d3ccd3fa3842812280f7f208718dc0b8610 \
        --op dense_gemm --package flashinfer/experimental/dense_projection_det_gemm \
        --package flashinfer/experimental/dense_projection_gemm

Co-Authored-By: Claude <noreply@anthropic.com>
Vendor the five packages of the gdn op into miles_ext_kernels/cake/_vendor/
with tools/vendor_cake.py, one level deep, recorded under ops.gdn of
_vendor/VENDORED.json:

- cake_gdn_det_train, cake_gdn_det_train_kcp, cake_gdn_train and
  cake_gdn_train_kcp: the FlashInfer Cake packages of
  flashinfer-ai/flashinfer eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3, byte
  for byte (85, 59, 80 and 67 files). This commit carries the fixes of
  the context-parallel packages, flashinfer-ai/flashinfer#8585 and #8623.
- cake_gdn_fla_cp_det: the det kernels with FLA's context-parallel state
  passing generated for radixark/miles#3597, at
  9d7543812b676a5b50dde107518782933b804b48: registry.json and the 56
  kernel and launcher sources of miles/kernels/attention/gdn/cake/csrc.
  The tool adds an empty __init__.py, FlashInfer's loader cake_build.py,
  the TVM-FFI host surface csrc/cake_host_shim.h in place of the
  torch-extension one (excluded, its blob recorded) and a build.json
  generated from the registry, so these sources build with nvcc and load
  with tvm-ffi like the FlashInfer packages; sm_107a compiles the sm_103a
  records.

The vendored files have the blobs the gdn packages, kernel sources and
registry had in the previous layout.

Commands, from the repository root, with a FlashInfer checkout at
../flashinfer and a miles checkout at ../miles:

    python tools/vendor_cake.py fi --checkout ../flashinfer \
        --commit eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3 --op gdn \
        --package flashinfer/experimental/cake_gdn_det_train \
        --package flashinfer/experimental/cake_gdn_det_train_kcp \
        --package flashinfer/experimental/cake/cake_gdn_train \
        --package flashinfer/experimental/cake_gdn_train_kcp
    python tools/vendor_cake.py miles-pr --checkout ../miles --pr 3597 \
        --commit 9d7543812b676a5b50dde107518782933b804b48 --op gdn \
        --name cake_gdn_fla_cp_det --path miles/kernels/attention/gdn/cake \
        --include csrc --include registry.json \
        --exclude csrc/cake_host_shim.h --arch-from sm_107a=sm_103a

Co-Authored-By: Claude <noreply@anthropic.com>
Lay out one namespace per kernel source, miles_ext_kernels.cake, with one shared runtime instead of a copy per
package:

- _runtime/spec.py: OpSpec and Flag, the whole declaration of an op (variants, the variants miles may choose, the
  miles selector, vendored packages, lazy entry points, options, its own argument checks).
- _runtime/registry.py: OPS in layer order (empty here), op_spec, prebuild(op, **model_params) and release(op=None);
  prebuild remembers its paths per op, CUDA device (index and architecture), toolkit (nvcc and its CUDA version),
  cache directory, model parameters and op options until release().
- _runtime/config.py: the one argument surface (add_arguments, validate_args, configure) and the variant each op
  runs. An unconfigured process runs det; validate_args keeps miles' command-line default fast for a selected op and
  writes the resolved variant back, also through a hidden flag when miles may choose only one variant, so it reaches
  every process; deterministic mode resolves det, rejects an explicit fast and rejects an op without det, also for
  direct configure callers. configure replaces the whole configuration and the implementations read the variant when
  a call runs; that read looks the op's declaration up in sys.modules instead of importing it again, and variant and
  option, which run at every call, remember what they found per configuration and op registry.
- _runtime/loader.py: the nvcc + tvm-ffi loader of the vendored packages, unchanged in behavior, plus install_op,
  load_package, prebuild_libraries, no_build (a build inside it raises, for the prebuild tests), device_arch, and a
  finder that resolves FlashInfer's ..cake.<name> imports to the one copy of _vendor.<name>, name by name.
- _runtime/workspace.py: one workspace pool per process, device and stream, bounded by --cake-workspace-pool-bytes
  on every request, LRU eviction, release() and stats(). The default bound is 0: nothing is kept and a request is a
  plain torch.empty on the current stream (measured GDN training steps were no faster with a 1 GiB pool, which kept
  up to 0.75 GiB per device); a positive bound reuses buffers. A stream that captures a CUDA graph allocates outside
  the pool; a lock keeps the pool consistent across threads. take runs at every call of an op, so its host work is
  small: once a configuration's bound was read as 0 and no device holds a buffer, take counts the request and calls
  torch.empty on the resolved device; with a positive bound it reads the current device, its stream and their capture
  state with one CUDA runtime query each (no device switch for the current device, no stream object), each device
  keeps a running byte total, and a hit evicts nothing unless the bound was lowered.
- miles_ext_kernels: SOURCES and the four hooks miles calls (add_arguments, validate_args, configure, release),
  torch-free and importing no op.

Tests under tests/cake/_runtime/ run every check on a synthetic op, so they do not pass vacuously while no op is
registered; tests/cake/helpers.py holds the shared accuracy and bitwise gates of the op tests, check_prebuild, which
runs every route an op's test_prebuild_gpu.py declares (ROUTES) with every variant after one prebuild, with no build
allowed, and POOLED, the 1 GiB bound a test of pooled buffers configures (the fixture pooled applies it). The op
contract also forbids call-time switches as prebuild parameters and checks that importing an implementation rebinds
every FlashInfer hook its vendored packages are allowed to keep. tests/cake/_runtime/fake_cuda.py is a fake CUDA
device layer on which the pool's CUDA path (current device, streams, capture) runs on the CPU; bench_workspace.py
next to it times take and the per-call lookups (a script, not a test).

Co-Authored-By: Claude <noreply@anthropic.com>
tools/vendor_cake.py is the one way sources enter miles_ext_kernels/cake/_vendor/, one level deep, recorded in one
_vendor/VENDORED.json grouped by op, with the source repository and commit of every package:

- fi: FlashInfer Cake packages, byte for byte.
- miles-pr: the generated kernel sources of a miles pull request (registry and csrc). The tool adds FlashInfer's
  loader, the TVM-FFI host surface tools/templates/cake_host_shim.h in place of the excluded torch-extension one, and
  a cake-build/1 build.json generated from the registry, so they build through the same loader as every other
  package instead of torch.utils.cpp_extension.
- remove, and --check [OP ...]: the exact file set and blob of every package, each build.json, one FlashInfer commit
  per op, that a vendored module imports packages of its own op only, and no environment read or flashinfer import in
  vendored modules unless allowed (a flashinfer import only inside a listed function, a hook the op rebinds).
  --check --upstream compares the recorded files with the upstream commits (trees only) for both kinds of source.
- merge: the git merge driver of VENDORED.json (.gitattributes), a three-way merge package by package, so layers
  that each add or change their own op merge and rebase without conflicts.

Co-Authored-By: Claude <noreply@anthropic.com>
- pyproject.toml: depend on torch and apache-tvm-ffi>=0.1.11,<0.2; a test extra; the multi_gpu marker.
- CI: the CPU tests with CPU torch, tools/vendor_cake.py --check, and a separate --check --upstream job.
- pre-commit: leave the vendored tree and FlashInfer's loader template byte-identical.
- README: install, the hooks miles calls, the Ops and Prebuild tables the op layers fill, variants and the scope of
  det, the prebuild contract, its routes and its limits, the workspace pool and the buffers packages keep outside it,
  the layout and the modules an op owns, vendoring, adding an op, the merge driver of the vendoring manifest, and how
  tests configure the pool.

Co-Authored-By: Claude <noreply@anthropic.com>
… not the README

The README no longer lists ops in a table; the op module and implementation
docstrings already state that neither variant is batch-invariant.
@Zhichenzzz
Zhichenzzz merged commit 80e21f2 into radixark:main Nov 5, 2026
4 checks passed
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.

3 participants