Skip to content

[Cake 1/13] gdn: chunked gated delta rule, det and fast, with KCP and FLA-CP - #1

Merged
Zhichenzzz merged 13 commits into
radixark:mainfrom
yyihuang:cake-stack/01-gdn
Nov 5, 2026
Merged

Zhichenzzz merged 13 commits into
radixark:mainfrom
yyihuang:cake-stack/01-gdn

Conversation

@yyihuang

@yyihuang yyihuang commented Nov 4, 2026 •

Copy link
Copy Markdown
Contributor

Stack 1/13 · base: main

Summary

Adds the gdn op, miles_ext_kernels.cake.gdn: the chunked gated delta rule of GDN linear-attention layers, forward and backward, on SM100a, SM103a and SM107a. chunk_gated_delta_rule keeps FLA v0.5.2's signature and autograd contract. kcp_chunk_gated_delta_rule and build_kcp_context expose the packages' own context parallelism (KCP).

The two variants, det and fast, run on five vendored packages:

miles selects the op with --use-miles-ext-kernels --linear-attention-backend cake and imports from miles_ext_kernels.cake.gdn import chunk_gated_delta_rule (paired PR radixark/miles#4676, whose update to this layout follows this stack). Through miles' arguments, an unset --cake-gdn-variant resolves to fast, or to det under --deterministic-mode. An explicit fast under deterministic mode is rejected. A process that never calls configure runs det. Package scratch goes through the runtime pool, which keeps nothing at its default bound of 0.

Stack

Layer 1 of 13. This PR replaces its previous version, "[Cake 1/11] gdn: shared build/vendor infra and GDN kernels (det, fast)" (head c62be504). The shared infrastructure of that version moved to layer 0 ([Cake 0/13] layout); this layer is the gdn op only.

Builds on layer 0 (head e0482323). Every PR targets main, so until layer 0 merges, this diff also contains layer 0's 3 commits. This layer's commits (e0482323..ad8eeef8):

  • e9dd400e vendor: gdn packages
  • 631ab4b2 gdn: chunked gated delta rule on the cake runtime
  • ad8eeef8 tests: gdn

Changes

This layer: 373 files, +316554 / -2 (previous version of this PR: 162 files, +106188 / -4). Per commit: e9dd400e 353 files, +309414 / -1; 631ab4b2 8 files, +3020 / -1; ad8eeef8 12 files, +4120. Almost all of it is vendored:

Group Files Lines
Vendored packages in cake/_vendor/: cake_gdn_det_train 85 (+67823), cake_gdn_det_train_kcp 59 (+64357), cake_gdn_train 80 (+67256), cake_gdn_train_kcp 67 (+72224), cake_gdn_fla_cp_det 61 = 57 upstream + 4 local (+37316) 352 +308976
cake/_vendor/VENDORED.json (ops.gdn block) 1 +438 / -1
Op: cake/gdn.py 65, cake/_impl/gdn/__init__.py 25, ops.py 2535, kcp.py 229, native.py 161; registry.py (OPS = ("gdn",), +1 / -1) 6 +3016 / -1
Tests: tests/cake/gdn/ 12 +4120
README (extras, Ops and Prebuild rows), pyproject.toml (gdn extra) 2 +4

The 56 kernel sources and registry.json of cake_gdn_fla_cp_det keep the blobs of the previous version's gdn/csrc and gdn/registry.json. The previous version vendored only cake_gdn_train, from FlashInfer c5db8d3; the four FlashInfer packages here are at eb3eb68d. The message of e9dd400e lists both vendoring commands (with ../flashinfer and ../miles checkouts).

Packages

Package Upstream Role Libraries sm_100a / sm_103a / sm_107a
cake_gdn_det_train flashinfer-ai/flashinfer@eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3, flashinfer/experimental/cake_gdn_det_train det, no CP 36 / 36 / 36
cake_gdn_train same commit, flashinfer/experimental/cake/cake_gdn_train fast, no CP 35 / 35 / 34
cake_gdn_det_train_kcp same commit, flashinfer/experimental/cake_gdn_det_train_kcp det, KCP 19 / 19 / 19
cake_gdn_train_kcp same commit, flashinfer/experimental/cake_gdn_train_kcp fast, KCP 23 / 23 / 23
cake_gdn_fla_cp_det radixark/miles#3597 head 9d7543812b676a5b50dde107518782933b804b48, miles/kernels/attention/gdn/cake (csrc + registry.json) the det version of the FLA-CP path (FLA cp_context, and KCP calls the KCP packages do not serve), either variant 25 / 25 / 25

Variant routing

The variant is read when the call runs, and a backward runs its forward's choice.

Call det fast
No cp_context cake_gdn_det_train cake_gdn_train
A KcpContext cake_gdn_det_train_kcp cake_gdn_train_kcp
An FLA cp_context, or a KcpContext call the KCP packages do not serve cake_gdn_fla_cp_det cake_gdn_fla_cp_det (the det version of this path), logged once per reason
  • A KcpContext call goes to cake_gdn_fla_cp_det when it uses the in-kernel L2 norm without its records in the package build, or a partition beyond the package's plan caps. That route needs a contiguous partition whose world_size divides the token count; otherwise the call raises.
  • Under fast, these calls run det and say so once through warn_runs_det (the fast variant does not serve an FLA cp_context; running det).
  • An unconfigured process runs det. Through miles' arguments: an unset flag gives fast (written back by validate_args, delivered by configure(args.miles_ext_kernels_args)), --deterministic-mode gives det, and an explicit fast under deterministic mode raises ValueError.
  • Under any cp_context the contract is FLA's: initial_state and output_final_state raise ValueError when the call is made, and the context's cu_seqlens / cu_seqlens_cpu replace the arguments.
  • Scope of det: bits repeat for identical inputs (the whole packed micro-batch; the plan is a pure function of every sequence length and the device) on the same GPU model and SM count, toolkit and build, the same route (no CP / KCP / FLA context, same partition, host or device boundaries) and the same --cake-gdn-recompute-state. Across those, results agree only within the accuracy budget. fast makes no promise about its bits; the tests compare it within the accuracy budgets only. The recomputed det backward is compared within the budgets of the saved one.

API and flags

cake/gdn.py is torch-free: OP = OpSpec(name="gdn", variants=("det", "fast"), selected_by=(("linear_attention_backend", "cake"),), packages=(the 5 above), exports=("KcpContext", "build_kcp_context", "chunk_gated_delta_rule", "kcp_chunk_gated_delta_rule"), flags=(--cake-gdn-recompute-state,)). Entry points load lazily from cake._impl.gdn.

  • chunk_gated_delta_rule(q, k, v, g, beta, scale=None, initial_state=None, output_final_state=False, cu_seqlens=None, cu_seqlens_cpu=None, use_qk_l2norm_in_kernel=False, cp_context=None, chunk_size=64, **unsupported): q/k [B, T, H, K], v [B, T, HV, V], K = V = 128, HV % H == 0, BF16 q/k/v, packed varlen (B = 1), FP32 initial/final state [N, HV, K, V], in-kernel q/k L2 norm. Rejected explicitly: use_gate_in_kernel, beta sigmoid fusion, state_v_first, head_first, chunk sizes other than 64. cu_seqlens_cpu (tuple of ints or CPU tensor) selects the host-table route, with no device-to-host sync. A call without tokens returns an empty output and passes the state through.
  • kcp_chunk_gated_delta_rule(q, k, v, g, beta, kcp_context, scale=None, use_qk_l2norm_in_kernel=False, chunk_size=64) runs only the KCP package of the configured variant; a call it does not serve raises NotImplementedError.
  • build_kcp_context(cu_seqlens_global, group=None, *, device=None, conv_kernel_size=4, layout="contiguous", fragments=None) -> KcpContext builds host metadata only (no tensor read, no collective). Layouts: contiguous (FLA's partition when world_size divides the token count), zigzag, zigzag_per_document or explicit fragments. LRU cache of 64 contexts. KCP takes FP32 g/beta (BF16 beta is cast exactly) and no initial or final state.
  • --cake-gdn-recompute-state (store_true, default off) recomputes w / u / h / v_new in the backward. The single-rank packages take it as save_state, the KCP packages as save_mode="light"; its help states the accuracy and the bits per route. This flag and --cake-gdn-variant are both rejected without --linear-attention-backend cake.

Changes from the previous version of this PR

  • An explicit fast under --deterministic-mode, or through configure(..., deterministic_mode=True), is now rejected: fast does not promise its bits, so deterministic mode runs det. Before, it was accepted.
  • FLA's contract applies under every cp_context. Before, a call with a cp_context used the caller's cu_seqlens (the global boundaries gave non-finite outputs) and accepted output_final_state, returning the local final state with a backward that raised.
  • det without context parallelism runs FlashInfer's cake_gdn_det_train.
  • The FLA-CP det kernels build through the runtime loader (nvcc + tvm-ffi). Removed: _jit.py (torch cpp_extension), the module-level VARIANT / SAVE_STATE / WORKSPACES, and the family's own arguments.py / configure.
  • The scratch buffers of the det and fast packages and of cake_gdn_fla_cp_det go through the runtime's workspace pool instead of per-package pools kept until the process ends. At the default --cake-workspace-pool-bytes of 0, each is a fresh allocation per call. Under a positive bound, the pool keeps one buffer per role and stream, freed by release(). The 0-element state placeholder is one tensor per device, shape and dtype; no call allocates one.
  • New extra gdn = ["flash-linear-attention>=0.5.2"]. The FLA-CP path calls the fla.ops.cp helpers of that release (with their state_v_first keyword) and is checked against it. 0.5.0's helpers lack the keyword, and 0.5.1's helpers module differs from 0.5.2's.

Prebuild

cake.prebuild("gdn", *, device=None, max_workers=None) takes no model parameter. It builds every library of the five packages for the device's architecture (138 / 138 / 137 for sm_100a / sm_103a / sm_107a). It raises NotImplementedError off SM100a / SM103a / SM107a, and RuntimeError when nvcc is too old (sm_107a needs CUDA 13.4). The memo key includes cake_gdn_recompute_state.

One exception, stated in the README Prebuild row and the docstring: FLA's own Triton context-parallel pre-processing (fla.ops.cp) runs only when the caller passes an FLA cp_context with a process group. It is third-party code and compiles in FLA on its first call.

Workspace and memory

  • All package scratch comes from the runtime pool through workspace.take("gdn", role, ...): 14 roles in ops.py (forward f_seg_state, f_seg_flag; backward b_dv_local, b_dh, b_dv2, b_seg_state, b_seg_flag, b_dh_end, b_dq_hv, b_dk_hv, b_dw, b_dg1, b_dk2, b_dg2). The FlashInfer det / fast single-rank packages' scratch hook (cake_chain._POOLED) points at the same pool. At the default bound of 0, every role is a fresh allocation per call and nothing is kept between calls.
  • The 0-element state placeholders (f_zero_state / b_zero_state, for a launch without initial or final state) are one tensor per device, shape and dtype for the process, not pool entries. They hold no memory, so streams share them without a race; release() drops them.
  • Outside the pool: the fast package's self-zeroing tail buffers (at most about 256 MiB per stream), freed by release(). The packages' own small per-layout chunk tables and plans stay.
  • gdn's release() drops the FLA-CP chunk tables, plans and loaded launchers, the cached KcpContext objects, the state placeholders and the single-rank packages' self-zeroing buffers.
  • Measured (one training step, fresh process): at the default bound of 0, the pool holds 0 bytes / 0 entries after the step and release() returns 0 (B200, GB300). With 1 GiB, det keeps 0.232 GiB in 5 entries (B200; GB300 249037952 B, 5 entries), det with FLA context parallelism 0.512 GiB in 11 entries (B200) and fast 0.232 GiB (B200); release() returns exactly the bytes held. det outputs and gradients are bit-identical to bound 0 (B200: det and det with FLA context parallelism; GB300: det, 12/12 tensors over 2 layers, on the miss and the hit run).

Review items

From the review of this PR's previous version. The layout comment and the stack-wide mechanisms are resolved in layer 0.

# Item Resolution Tests
1 det planner crashes on long single sequences (segment_len 171 / 171 / 228 > 128 at HV=48@32k fwd, HV=12@64k bwd, HV=16@128k fwd) det without CP runs FlashInfer cake_gdn_det_train (its own caps); the FLA-CP planner caps segments at PRESCAN_MAX_CHUNKS = 128 and keeps the serial scan when it cannot split within it CPU test_gdn_host.py::test_one_pre_scan_cta_covers_at_most_128_chunks, ::test_without_the_cap_a_long_single_sequence_would_cut_segments_the_pre_scan_cannot_serve, ::test_no_plan_cuts_a_segment_longer_than_the_cap, ::test_a_lower_capacity_bounds_the_split; GPU test_gdn_det_gpu.py::test_a_long_single_sequence_runs_on_the_det_package_and_the_fla_cp_det_kernels (48v_32k, 12v_64k, 16v_128k; host and device), test_gdn_wiring.py::test_a_long_single_sequence_never_plans_a_segment_beyond_the_pre_scan, ::test_segment_planner_only_splits_for_a_modelled_gain; ROUTES *_long_*
2 det has no extent checks (cu_seqlens ending before T, wrong g heads, initial_state N, mismatched cu_seqlens / cu_seqlens_cpu) det without CP: the FlashInfer det package with all its checks; FLA-CP path: _check_inputs (metadata only, no sync), and host boundaries must end at q's tokens, checked before any CUDA call CPU test_gdn_host.py::test_host_boundaries_must_end_at_the_tokens_before_any_cuda_call (tuple, tensor); GPU test_gdn_det_gpu.py::test_the_extents_are_checked_before_a_launch (with and without CP; the well-posed follow-up call uses in-kernel q/k L2 norm, as miles calls it)
3 CP contract differs from FLA (global cu_seqlens with cp_context gives non-finite output; output_final_state under CP returns the local state and the backward raises) under any cp_context, initial/final state are rejected at the call and the context's boundaries replace the arguments; the autograd function no longer materialises a missing incoming gradient (set_materialize_grads(False)) CPU test_gdn_kcp_wiring.py::test_the_context_parallel_contract_is_fla_s, ::test_an_fla_context_runs_the_fla_cp_det_kernels_with_the_context_s_boundaries; GPU test_gdn_fla_cp_gpu.py::test_the_context_s_boundaries_replace_the_call_s, ::test_an_unused_final_state_passes_no_gradient_and_a_final_state_alone_backpropagates, test_gdn_kcp_gpu.py::test_an_initial_or_a_final_state_is_rejected_when_the_call_is_made; multi_gpu ::test_cp_a_final_state_call_is_rejected_when_the_call_is_made (cp2, cp4)
4 workspace pool keeps 6.46 GiB at ~47 % hit rate, little time saved; "consider dropping it, or bounding it and exposing a clear hook for miles' sleep" all scratch through the runtime pool, whose default bound is 0 (keeps nothing; a fresh allocation per call); a positive bound is opt-in, bounded, per stream and freed by release(); out-of-pool tail buffers documented; 0-element placeholders cached instead of allocated GPU test_gdn_wiring.py::test_at_the_default_bound_of_0_a_training_step_retains_no_pooled_bytes, ::test_workspace_pool_hits_after_the_first_call_and_results_match_the_pool_off, ::test_pooled_workspaces_never_alias_saved_returned_or_gradient_tensors, ::test_a_side_stream_gets_its_own_buffers_and_the_same_bits, test_gdn_det_gpu.py::test_the_scratch_comes_from_the_bounded_pool_and_the_pool_does_not_change_the_bits, test_prebuild_gpu.py::test_a_det_call_keeps_its_scratch_in_the_bounded_pool (pool tests configure helpers.POOLED); CPU test_gdn_host.py::test_the_state_placeholder_is_one_tensor_per_device_shape_and_dtype, ::test_the_packages_take_their_scratch_from_the_pool_and_share_the_state_placeholders (det, fast)
5 state the scope of determinism (host vs device tables, SM count, segmented vs serial scan) stated in gdn.py, ops.py and the README (see Variant routing) CPU test_arguments.py::test_the_recompute_help_states_the_accuracy_and_the_bits
6 T=0 raises in det and fast empty output, state passed through (or zeros), correct backward; both variants and an FLA context without a process group CPU test_gdn_host.py::test_a_call_without_tokens_passes_the_state_through, ::test_a_call_without_tokens_and_without_a_final_state, ::test_a_call_without_tokens_checks_the_initial_state; GPU test_gdn_det_gpu.py::test_a_call_without_tokens (det and fast, with and without CP); ROUTES plain_no_token, fla_cp_no_token
code gdn/csrc (58 files) outside every manifest, a second build system (cpp_extension) vendored as cake_gdn_fla_cp_det, drift-checked by --check --upstream; _jit.py / cpp_extension removed, native.py loads loader-built tvm-ffi libraries CPU test_package.py::test_the_fla_cp_package_builds_one_library_per_stage_record_and_architecture, ::test_the_fla_cp_sources_build_with_the_flags_of_the_previous_build, ::test_every_stage_has_the_sm103a_export_sm107a_compiles; GPU test_gdn_fla_cp_gpu.py::test_every_stage_loads_from_the_loader_s_library
code VARIANT, SAVE_STATE, WORKSPACES process-global, set by configure removed; read when the call runs, the backward uses ctx CPU test_arguments.py::test_configure_sets_the_variant_and_the_recompute_switch_the_kernels_read; GPU test_prebuild_gpu.py::test_a_call_runs_the_variant_configured_when_it_is_made
code pyproject.toml declares no dependencies layer 0 adds torch + apache-tvm-ffi; this layer adds the extra gdn = ["flash-linear-attention>=0.5.2"] —
stack unconfigured default, compiles inside a call, argument owner layer 0 runtime; gdn runs det unconfigured, prebuild builds all five packages, flags go through OP CPU test_arguments.py::test_without_the_arguments_the_kernels_keep_the_tape_and_run_det, ::test_an_explicit_fast_is_kept_and_rejected_under_deterministic_mode; GPU test_prebuild_gpu.py::test_after_one_prebuild_no_route_compiles
layout one namespace, one manifest, csrc under _vendor, one runtime layer 0, plus this layer's packages in cake/_vendor/ and its ops.gdn block layer 0 tests; test_package.py::test_the_implementation_loads_exactly_the_declared_packages

Also fixed after GPU runs of earlier versions (not review items): with the shared multi-GPU fixture, a case that raises on every rank before any communication fails only its own tests. The process group has a 5-minute timeout, so a case that raises on some ranks only cannot hang forever, but it is not isolated. Test-only fixes from the GPU rounds of this stack: the extents test's well-posed call normalises q/k in the kernel (unnormalised unit-variance inputs make the state diverge; FLA's Triton kernel also gives non-finite outputs there); planner tests plan at the B200 / GB300 SM counts their expectations were modelled for; the planner-gain test only considers lengths within the 128-chunk cap; the launch-count test skips, naming the device, where torch.profiler records no CUDA API call.

Tests

tests/cake/gdn/ has 119 test functions. CPU: test_arguments.py 13, test_package.py 6, test_gdn_host.py 10, test_gdn_kcp_wiring.py 15. GPU (skipped without CUDA): test_gdn_wiring.py 34, test_gdn_det_gpu.py 12, test_gdn_fast_gpu.py 5, test_gdn_fla_cp_gpu.py 7, test_gdn_kcp_gpu.py 12 (including multi_gpu CP 2 / 4), test_fla_parity.py 1, test_prebuild_gpu.py 4. On a GPU node, 318 are collected without multi_gpu and 30 with it.

  • CPU: arguments, packages and manifests, the planner cap, calls without tokens, host boundaries checked before any CUDA call, the shared state placeholder and package scratch from the pool, and KCP / FLA routing on the packages' torch emulation.
  • GPU: det and fast against FP64, FLA and each other; det bit for bit run to run, in a new process, under deterministic algorithms and on a side stream alternating with the default one; extents; 128k-token sequences; pool tests under helpers.POOLED at the bits of bound 0, and bound 0 keeps no pooled byte after a training step; FLA-CP from the loader's libraries; KCP on one GPU and multi_gpu CP 2 / 4.
  • test_prebuild_gpu.py declares 23 ROUTES: plain_host, plain_device, plain_one_sequence, plain_long_host, plain_long_device, plain_state, plain_state_device, plain_recompute, plain_no_l2norm_fp32_beta, plain_batch, plain_no_token, kcp, kcp_long, kcp_recompute, kcp_no_l2norm_fp32_beta, fla_cp_host, fla_cp_device, fla_cp_long_host, fla_cp_long_device, fla_cp_recompute, fla_cp_no_l2norm_fp32_beta, fla_cp_no_token, side_stream. Layout 300 / 517 / 64 / 7 tokens; long = one 9000-token sequence (141 chunks); 4 key and 12 value heads. check_prebuild runs every route under each variant with no build allowed. Planner expectations use the 148 / 152 SM counts of B200 / GB300.

Validation

CPU: image radixark/miles:dev-202611032236 (Python 3.12.16, torch 2.14.1+cu130, tvm-ffi 0.1.11), x86_64 CPU node, CUDA hidden. The CPU tests compile nothing.

Commit pytest tests, fresh venv without torch pytest tests, image Python vendor_cake.py --check pre-commit run --all-files
e9dd400e (vendor) 178 passed, 5 skipped, 1 failed (allowed by design: test_vendored.py::test_every_vendored_op_is_registered) 221 passed, 2 skipped, same 1 allowed failure rc 0 rc 0, no diff
631ab4b2 (op) 184 passed, 6 skipped, 1 failed (allowed by design: test_op_contract.py::test_the_op_has_cpu_tests_and_a_prebuild_test[gdn]) 229 passed, 1 skipped, same 1 allowed failure rc 0 rc 0, no diff
ad8eeef8 (tests) 221 passed, 15 skipped, 0 failed 388 passed, 8 skipped, 0 failed (skips: GPU tests without CUDA) rc 0; --check gdn --upstream rc 0 rc 0, no diff

tests/cake/gdn alone in the image at ad8eeef8: 158 passed, 7 skipped. --check --upstream also returns 0 at the stack top, whose ops.gdn block, gdn.py, _impl/gdn/, tests/cake/gdn/ and five package trees are identical to this layer's. GitHub CI has not run on this head yet.

GPU: tests/cake/gdn at the stack top (gdn files identical to this layer's), with a private JIT cache per op.

B200 GB300
GPU / arch 8x B200 (148 SMs), sm_100a, driver 580.82.07 4x GB300 (152 SMs), sm_103a, driver 580.159.03
Software torch 2.14.1+cu130, nvcc 13.0.88, flash-linear-attention 0.5.2, tvm-ffi 0.1.11, triton 3.8.0, Python 3.12.16 same versions
Cold (empty cache), not multi_gpu 318 / 318 passed, 0 skipped, 998.8 s (8 ops built in parallel on the node) 318 / 318 passed, 0 skipped, 642.6 s (test_prebuild_gpu.py first)
Warm 318 / 318, 80.3 s 318 / 318, 101.0 s
test_prebuild_gpu.py alone, first, on an empty cache 4 / 4, 145.5 s (builds 138 libraries; every route under no_build()) covered by the cold run (prebuild test first)
multi_gpu (4 GPUs) 30 / 30, 115.3 s: cp2 / cp4, including test_cp_a_final_state_call_is_rejected_when_the_call_is_made (cp2, cp4) 30 / 30, 115.3 s, same
det bit for bit, same process pass (all det repeat tests) pass
det bit for bit, new process: test_gdn_det_gpu.py::test_det_repeats_bit_for_bit_in_a_new_process, test_gdn_fla_cp_gpu.py::test_the_call_repeats_bit_for_bit_in_a_new_process pass pass
default and side stream alternating, bit for bit: test_gdn_det_gpu.py::test_det_repeats_on_a_side_stream_alternating_with_the_default_stream (both cases), test_gdn_fla_cp_gpu.py::test_a_side_stream_alternating_with_the_default_stream_gives_the_same_bits, test_gdn_wiring.py::test_a_side_stream_gets_its_own_buffers_and_the_same_bits pass pass
Pool default (one step: bound 0 keeps nothing; 1 GiB gives the same bits) pass pass

sm_107a (R200): validation of this tree is pending.

Build migration (a one-time check, not a resident test; run on an earlier commit with the identical cake_gdn_fla_cp_det tree): with nvcc 13.0.88, the loader-built FLA-CP libraries have the same SASS, function by function, as the previous torch-extension build for 50/50 libraries (25 per arch, sm_100a and sm_103a). With nvcc 13.4.92, sm_107a and sm_103a match 25/25 each against the previous compile flags. No torch-extension build for CUDA 13.4 was available, so the old side used its flags without torch. The whole-file cuobjdump text differs only by one function-less ELF from the tvm-ffi binding unit. The old build's torch-only flags (-D/-U__CUDA_NO_HALF*, -std=c++20) do not change SASS.

Performance

GDN det A/B: previous layout vs this stack (B200)

Same GPU, node otherwise idle. "old" is the previous layout of this op (per-op runtime, FlashInfer packages at c5db8d3, same cake_gdn_fla_cp_det kernel sources); "new0" is this stack at the default pool bound 0; "new1g" is this stack with a 1 GiB pool. Arms: a = det kernels with FLA's context parallelism called as one rank (same kernel sources on both sides), b = det (FlashInfer det package), c = fast (FlashInfer fast package). Arms b and c also change the FlashInfer commit (c5db8d3 -> eb3eb68d), so their ratios include package changes, not only the layout.

Method: eager, CUDA events, L2 flushed and device synchronized before every rep (host cost inside the window), 5 warm-up + 30 timed reps per process, arms interleaved ABC / CBA, 4 processes per side in rotated order; cell = median of the per-process medians; ratio [min, max] over the 4 rounds. Ops: fwd (with tape), bwd (after an untimed forward), fwdbwd, ckpt (forward under no_grad, forward with tape, backward). 15 shapes: the miles micro-batch layout (3339, 2802, 1340, 263 tokens) with 4x12 and 16x32 heads on host tables and 16x16 on device tables, plus 12 equal-length packed shapes (1/2/4 x 4096/8192 tokens, 8x16 and 16x32 heads). Torch 2.14.1+cu130, CUDA 13.0, driver 580.82.07. In all 12 processes the probes held: b and c bit for bit run to run, b save_state on/off same bits, and a / b / c within budget of an FP64 recurrence.

Summary (median over rows; brackets = min and max over rows; ratios > 1 = this stack slower):

ratio all rows miles shapes equal-length shapes arm a arm b arm c
new0 / old (bound 0) 1.031 [0.971, 1.185] 1.051 1.026 1.027 1.025 1.046
new1g / old (1 GiB) 1.053 [0.926, 1.246] 1.078 1.035 1.032 1.058 1.076
new1g / new0 1.015 [0.867, 1.100] 1.028 1.013 1.008 1.026 1.019

new0 / old by op (median): fwd 1.004, bwd 1.040, fwdbwd 1.038, ckpt 1.043. On the 4x8192-token shapes new0 / old is 0.971-1.046: arms b / c are faster in fwd, fwdbwd and ckpt (0.971-0.999) and slower in bwd (1.014-1.046). A 1 GiB pool is slower than bound 0, which supports the default of 0.

Selected rows (ms; 60 of the 180 rows):

shape op arm old ms new0 ms new1g ms new0/old new1g/old new1g/new0 earlier run: old / new ms (new/old)
miles-4x12-host fwd a 0.1713 0.1847 0.1843 1.078 [1.064, 1.084] 1.076 [1.066, 1.086] 0.998 [0.993, 1.002] 0.1697 / 0.1814 (1.069)
miles-4x12-host fwd b 0.1561 0.1597 0.1595 1.023 [1.012, 1.043] 1.022 [1.013, 1.026] 0.999 [0.983, 1.003] 0.1549 / 0.1567 (1.012)
miles-4x12-host fwd c 0.1547 0.1592 0.1585 1.029 [1.008, 1.050] 1.025 [1.017, 1.028] 0.995 [0.979, 1.009] 0.1528 / 0.1565 (1.025)
miles-4x12-host bwd a 0.5569 0.5960 0.6387 1.070 [1.034, 1.812] 1.147 [1.133, 1.918] 1.072 [1.026, 1.124] 0.5545 / 0.6032 (1.088)
miles-4x12-host bwd b 0.5445 0.5804 0.6154 1.066 [1.031, 2.007] 1.130 [1.059, 2.128] 1.060 [1.028, 1.064] 0.5435 / 0.5789 (1.065)
miles-4x12-host bwd c 0.5261 0.5609 0.5641 1.066 [1.022, 2.225] 1.072 [1.063, 2.205] 1.006 [0.991, 1.063] 0.5376 / 0.5709 (1.062)
miles-4x12-host fwdbwd a 0.7231 0.7989 0.8387 1.105 [1.070, 1.176] 1.160 [1.139, 1.204] 1.050 [1.023, 1.104] 0.7258 / 0.8028 (1.106)
miles-4x12-host fwdbwd b 0.7439 0.8127 0.8573 1.092 [1.068, 1.122] 1.152 [1.108, 1.202] 1.055 [1.038, 1.072] 0.7611 / 0.8099 (1.064)
miles-4x12-host fwdbwd c 0.7292 0.7999 0.8346 1.097 [1.075, 1.110] 1.145 [1.112, 1.231] 1.043 [1.002, 1.110] 0.7376 / 0.7988 (1.083)
miles-4x12-host ckpt a 0.8265 0.9015 0.9565 1.091 [1.063, 1.129] 1.157 [1.118, 1.206] 1.061 [1.051, 1.068] 0.8294 / 0.9079 (1.095)
miles-4x12-host ckpt b 0.9051 0.9364 0.9891 1.035 [0.969, 1.070] 1.093 [1.065, 1.121] 1.056 [0.995, 1.122] 0.8654 / 0.9508 (1.099)
miles-4x12-host ckpt c 0.8505 0.9371 0.9750 1.102 [1.063, 1.139] 1.146 [1.098, 1.188] 1.040 [1.009, 1.056] 0.8364 / 0.8979 (1.074)
miles-16x32-host fwd a 0.3745 0.3848 0.3862 1.028 [1.023, 1.034] 1.031 [1.031, 1.038] 1.004 [0.997, 1.008] 0.3722 / 0.3819 (1.026)
miles-16x32-host fwd b 0.2919 0.2928 0.2937 1.003 [0.996, 1.010] 1.006 [1.001, 1.013] 1.003 [1.000, 1.007] 0.2888 / 0.2883 (0.998)
miles-16x32-host fwd c 0.2897 0.2912 0.2928 1.005 [1.003, 1.011] 1.011 [1.009, 1.013] 1.005 [1.002, 1.007] 0.2872 / 0.2875 (1.001)
miles-16x32-host bwd a 0.8971 0.9116 0.9198 1.016 [1.008, 1.021] 1.025 [1.020, 1.030] 1.009 [1.004, 1.015] 0.8895 / 0.9037 (1.016)
miles-16x32-host bwd b 0.6418 0.6637 0.6709 1.034 [1.026, 1.042] 1.045 [1.039, 1.054] 1.011 [1.005, 1.019] 0.6328 / 0.6571 (1.038)
miles-16x32-host bwd c 0.6304 0.6630 0.6706 1.052 [1.049, 1.064] 1.064 [1.055, 1.081] 1.011 [1.003, 1.017] 0.6242 / 0.6563 (1.051)
miles-16x32-host fwdbwd a 1.0782 1.1056 1.1216 1.025 [1.025, 1.034] 1.040 [1.035, 1.046] 1.015 [1.004, 1.019] 1.0697 / 1.0986 (1.027)
miles-16x32-host fwdbwd b 0.8445 0.8873 0.9103 1.051 [1.049, 1.095] 1.078 [1.068, 1.136] 1.026 [1.012, 1.037] 0.8040 / 0.8696 (1.082)
miles-16x32-host fwdbwd c 0.8057 0.8902 0.9159 1.105 [1.072, 1.127] 1.137 [1.098, 1.159] 1.029 [1.020, 1.046] 0.8228 / 0.8828 (1.073)
miles-16x32-host ckpt a 1.4157 1.4359 1.4348 1.014 [1.013, 1.021] 1.013 [0.992, 1.014] 0.999 [0.972, 1.001] 1.4070 / 1.4267 (1.014)
miles-16x32-host ckpt b 0.9884 1.0313 1.0606 1.043 [1.040, 1.057] 1.073 [0.985, 1.082] 1.028 [0.946, 1.040] 0.9822 / 1.0382 (1.057)
miles-16x32-host ckpt c 0.9845 1.0336 1.0619 1.050 [1.046, 1.054] 1.079 [0.981, 1.085] 1.027 [0.935, 1.037] 0.9767 / 1.0339 (1.059)
miles-16x16-device fwd a 0.2810 0.2917 0.2909 1.038 [1.034, 1.048] 1.035 [1.032, 1.045] 0.997 [0.985, 1.003] 0.2761 / 0.2885 (1.045)
miles-16x16-device fwd b 0.2523 0.2536 0.2526 1.005 [0.990, 1.018] 1.001 [0.989, 1.014] 0.996 [0.985, 1.006] 0.2469 / 0.2488 (1.008)
miles-16x16-device fwd c 0.2498 0.2517 0.2498 1.008 [0.996, 1.014] 1.000 [0.998, 1.002] 0.992 [0.984, 1.005] 0.2456 / 0.2462 (1.003)
miles-16x16-device bwd a 0.6329 0.6586 0.6771 1.041 [1.028, 1.051] 1.070 [0.765, 1.083] 1.028 [0.744, 1.036] 0.6169 / 0.6503 (1.054)
miles-16x16-device bwd b 0.6139 0.6497 0.6741 1.058 [1.049, 1.068] 1.098 [0.754, 1.125] 1.037 [0.719, 1.066] 0.6028 / 0.6418 (1.065)
miles-16x16-device bwd c 0.6057 0.6481 0.6743 1.070 [1.051, 1.103] 1.113 [0.772, 1.145] 1.041 [0.735, 1.068] 0.5943 / 0.6397 (1.076)
miles-16x16-device fwdbwd a 0.8208 0.8762 0.9417 1.067 [0.991, 1.163] 1.147 [0.919, 1.202] 1.075 [0.790, 1.145] 0.8140 / 0.9006 (1.106)
miles-16x16-device fwdbwd b 0.8619 0.8854 0.9649 1.027 [0.968, 1.119] 1.120 [0.844, 1.179] 1.090 [0.755, 1.218] 0.8361 / 0.9153 (1.095)
miles-16x16-device fwdbwd c 0.8250 0.9141 0.9417 1.108 [1.019, 1.183] 1.141 [0.895, 1.192] 1.030 [0.757, 1.171] 0.8023 / 0.9142 (1.140)
miles-16x16-device ckpt a 0.9791 1.0624 1.1084 1.085 [1.080, 1.096] 1.132 [0.948, 1.148] 1.043 [0.876, 1.048] 0.9116 / 1.0662 (1.170)
miles-16x16-device ckpt b 1.0473 1.1009 1.1572 1.051 [0.992, 1.058] 1.105 [0.844, 1.112] 1.051 [0.851, 1.053] 0.9219 / 1.0999 (1.193)
miles-16x16-device ckpt c 1.0183 1.1017 1.1525 1.082 [1.054, 1.161] 1.132 [0.854, 1.227] 1.046 [0.810, 1.057] 0.8953 / 1.0920 (1.220)
fi-1x4096-8x16-host fwd a 0.1382 0.1493 0.1521 1.081 [1.072, 1.097] 1.101 [1.087, 1.124] 1.019 [1.009, 1.025] 0.1364 / 0.1455 (1.067)
fi-1x4096-8x16-host fwd b 0.1455 0.1570 0.1707 1.079 [1.059, 1.108] 1.173 [1.156, 1.202] 1.087 [1.062, 1.106] 0.1440 / 0.1514 (1.052)
fi-1x4096-8x16-host fwd c 0.1457 0.1543 0.1684 1.059 [1.051, 1.088] 1.156 [1.141, 1.202] 1.091 [1.062, 1.125] 0.1434 / 0.1514 (1.056)
fi-1x4096-8x16-host bwd a 0.5602 0.6015 0.6440 1.074 [0.891, 2.234] 1.150 [0.552, 2.377] 1.071 [0.619, 1.081] 0.5117 / 0.6081 (1.188)
fi-1x4096-8x16-host bwd b 0.5548 0.5847 0.6082 1.054 [1.041, 2.606] 1.096 [0.488, 2.722] 1.040 [0.465, 1.048] 0.5526 / 0.5855 (1.060)
fi-1x4096-8x16-host bwd c 0.5446 0.5691 0.5928 1.045 [1.037, 2.801] 1.088 [0.463, 2.885] 1.041 [0.446, 1.053] 0.5325 / 0.5632 (1.058)
fi-1x4096-8x16-host fwdbwd a 0.7152 0.7987 0.8614 1.117 [1.079, 2.080] 1.204 [0.659, 2.235] 1.078 [0.590, 1.119] 0.7055 / 0.8079 (1.145)
fi-1x4096-8x16-host fwdbwd b 0.7599 0.8173 0.8461 1.076 [1.046, 2.145] 1.114 [0.590, 2.288] 1.035 [0.564, 1.067] 0.7327 / 0.8126 (1.109)
fi-1x4096-8x16-host fwdbwd c 0.7513 0.7910 0.8526 1.053 [0.996, 2.170] 1.135 [0.565, 2.400] 1.078 [0.567, 1.106] 0.6909 / 0.7996 (1.157)
fi-1x4096-8x16-host ckpt a 0.7859 0.9314 0.9790 1.185 [1.097, 1.914] 1.246 [0.717, 1.964] 1.051 [0.653, 1.076] 0.8096 / 0.9047 (1.118)
fi-1x4096-8x16-host ckpt b 0.8666 0.9639 0.9982 1.112 [1.049, 1.941] 1.152 [0.683, 2.064] 1.036 [0.646, 1.074] 0.8518 / 0.9673 (1.136)
fi-1x4096-8x16-host ckpt c 0.8306 0.9382 0.9685 1.130 [0.879, 2.006] 1.166 [0.675, 2.185] 1.032 [0.769, 1.089] 0.8276 / 0.9337 (1.128)
fi-4x8192-16x32-host fwd a 1.2470 1.2531 1.2529 1.005 [1.004, 1.006] 1.005 [1.004, 1.006] 1.000 [0.999, 1.000] 1.2486 / 1.2538 (1.004)
fi-4x8192-16x32-host fwd b 0.8535 0.8291 0.8271 0.971 [0.968, 0.973] 0.969 [0.965, 0.971] 0.998 [0.995, 0.999] 0.8570 / 0.8294 (0.968)
fi-4x8192-16x32-host fwd c 0.8533 0.8282 0.8257 0.971 [0.969, 0.972] 0.968 [0.965, 0.969] 0.997 [0.995, 0.999] 0.8569 / 0.8278 (0.966)
fi-4x8192-16x32-host bwd a 2.6219 2.6440 2.6664 1.008 [1.005, 1.012] 1.017 [1.015, 1.020] 1.008 [1.008, 1.009] 2.6313 / 2.6389 (1.003)
fi-4x8192-16x32-host bwd b 1.5517 1.5737 1.5977 1.014 [1.009, 1.024] 1.030 [1.028, 1.034] 1.015 [1.009, 1.024] 1.5531 / 1.5681 (1.010)
fi-4x8192-16x32-host bwd c 1.5383 1.5740 1.5802 1.023 [1.017, 1.030] 1.027 [1.021, 1.032] 1.004 [1.000, 1.005] 1.5422 / 1.5690 (1.017)
fi-4x8192-16x32-host fwdbwd a 3.6738 3.6937 3.6935 1.005 [1.003, 1.007] 1.005 [1.005, 1.007] 1.000 [0.999, 1.001] 3.8241 / 3.8063 (0.995)
fi-4x8192-16x32-host fwdbwd b 2.2050 2.1816 2.1780 0.989 [0.987, 0.992] 0.988 [0.987, 0.991] 0.998 [0.998, 0.999] 2.2841 / 2.2428 (0.982)
fi-4x8192-16x32-host fwdbwd c 2.2008 2.1768 2.1747 0.989 [0.988, 0.990] 0.988 [0.987, 0.989] 0.999 [0.997, 1.001] 2.2759 / 2.2288 (0.979)
fi-4x8192-16x32-host ckpt a 4.8943 4.9182 4.9115 1.005 [1.004, 1.007] 1.004 [1.002, 1.006] 0.999 [0.998, 0.999] 5.1241 / 5.1460 (1.004)
fi-4x8192-16x32-host ckpt b 3.0234 2.9749 2.9640 0.984 [0.983, 0.985] 0.980 [0.974, 0.984] 0.996 [0.990, 0.999] 3.1439 / 3.0842 (0.981)
fi-4x8192-16x32-host ckpt c 3.0121 2.9625 2.9553 0.984 [0.982, 0.985] 0.981 [0.977, 0.984] 0.998 [0.995, 0.999] 3.1399 / 3.0686 (0.977)

The last column is an earlier run on another B200 node with a previous version of layer 0's take (median new/old 1.039 there, miles shapes 1.064). A [max] near 2 comes from bimodal per-process bwd medians in single processes of the old or new1g side (e.g. old 0.22 vs 0.56 ms); the cells use the median over the 4 rounds.

Host time per call under cProfile (miles-4x12-host, us, fwd / bwd; launch-return time, profiler overhead included):

arm old new0 new1g
a 251.3 / 605.4 285.1 / 602.8 285.5 / 677.5
b 296.9 / 555.6 323.9 / 560.2 356.3 / 611.7
c 270.8 / 527.8 311.4 / 575.7 307.5 / 556.6

workspace.take cumulative per op at bound 0: a bwd 12 calls 72.4 us, b bwd 6 calls 45.6 us (1 GiB: 130.6 / 78.0 us). Residual: a small host overhead remains on host-bound small shapes (miles shapes median 1.051x, max 1.185x on one equal-length ckpt row); on the 4x8192-token shapes arm a is within 1.013x, and arms b / c are faster except in bwd (1.014-1.046x).

Vendored kernels (follow-up comments)

The performance of the vendored packages is copied from their source pull requests, with tables and captions unchanged, into five follow-up comments, "Performance (1/5)" to "Performance (5/5)":

  1. Packages and source revisions, the FlashInfer pull requests of the vendored revision, baselines and per-architecture coverage; det single-rank on sm_107a (exact revision, perf(cake_gdn_det): packed-triangular forward WY program at five CTAs per SM on R200 flashinfer-ai/flashinfer#8695) and on sm_100a / sm_103a (perf(cake_gdn_det): packed-triangular forward WY program at five CTAs per SM on B200 and GB300 flashinfer-ai/flashinfer#8615, same programs).
  2. det KCP: fix(cake_gdn_det_kcp): keep M_seg in the full saved set at P > 1 and count the whole plane in the saved byte query flashinfer-ai/flashinfer#8585 (exact revision, host-only fix) and perf(cake_gdn_det_kcp): chain-aware walk widths in the deterministic KCP GDN package flashinfer-ai/flashinfer#8538 (same kernels).
  3. fast single-rank on sm_100a and sm_103a (exact revision, perf(cake_gdn): run the split adjoint scan beside the group-of-three tail on B200 and GB300 flashinfer-ai/flashinfer#8675), against deterministic Cake, the previous fast package and the strongest non-Cake implementation.
  4. fast single-rank on sm_107a (perf(cake_gdn): packed-triangular forward WY program at five CTAs per SM on B200 and GB300 flashinfer-ai/flashinfer#8665), and fast KCP in eager mode with save_mode="full", this layer's default (perf(cake_gdn_kcp): with grad, full-save forward writes the output and the tapes in one walk flashinfer-ai/flashinfer#8644).
  5. fast KCP at the exact revision (perf(cake_gdn_kcp): with grad, graph-replay checkpoint recompute hands its tapes to the backward flashinfer-ai/flashinfer#8638, graph mode, which this layer does not use); the det kernels with FLA's context parallelism (cake_gdn: deterministic chunked GDN training kernels (--linear-attention-backend cake) for the head-sharded Qwen GDN layer miles#3597: whole-model training against FLA on B200 and GB300, kernel gains per optimization round); and the list of missing measurements.

Missing, as listed in comment 5: a det vs non-Cake ratio (single-rank and KCP), det KCP on sm_107a against an earlier version, fast single-rank on sm_107a with BF16 beta, fast KCP at CP > 1 on R200 against a non-Cake implementation, fast KCP with save_mode="light" in eager mode, for the FLA-CP det kernels any sm_107a or kernel-level non-Cake measurement and any measurement of their later WY backward change, and B300 numbers.

Notes

  • sm_107a: 137 libraries; cake_gdn_train lists 34 for sm_107a against 35 for sm_100a and sm_103a. The toolkit floor is CUDA 13.4. cake_gdn_fla_cp_det builds sm_107a from its sm_103a records, and its sm_107a SASS was compared only against the previous compile flags. Validation of this tree on R200 is pending.
  • Limits: K = V = 128, BF16 q/k/v, packed varlen with B = 1, chunk size 64; no initial or final state under any context parallelism. A KCP call routed to the FLA-CP kernels needs a contiguous partition whose world_size divides the token count. prebuild does not cover FLA's Triton CP pre-processing.
  • A small host overhead over the previous layout remains on host-bound small shapes (see the A/B above).
  • Multi-GPU tests: a case that raises on some ranks only is not isolated; the 5-minute process-group timeout keeps it from hanging forever.
  • Shape labels "miles-..." name the miles micro-batch layout, not a model.

🤖 Generated with Claude Code

@Zhichenzzz

Copy link
Copy Markdown
Collaborator

Review of layer 1 (gdn + shared infra) — suggest request changes

Confirmed on GPU

  1. The det planner crashes on long single sequences. _best_split / choose_scan_plan never cap segment_len at the 128 chunks the pre-scan stages check, so fwd_h / dhu raise ValueError: segment_len … exceeds the 128 chunks:

    • HV=48 @ 32k tokens: forward, segment_len 171
    • HV=12 @ 64k: backward, 171
    • HV=16 @ 128k: forward, 228

    Same on the host and device cu_seqlens routes; fast is fine; 16k is fine. det is the only CP variant and the --deterministic-mode default, and the planner tests stop at 16k.

  2. det has no extent checks (fast's _check_inputs has all of them):

    • T=1024 with cu_seqlens ending at 1000 is accepted. With outputs pre-filled with NaN, the last 24 tokens of o, dq, dk, dv, dg and dbeta stay NaN, i.e. they are never written.
    • g of shape [1, T, H] with HV=12 is accepted and launched (out-of-bounds read).
    • initial_state with N=1 for 3 sequences ends in an illegal memory access.
    • Device cu_seqlens with 4 offsets and cu_seqlens_cpu with 2 is accepted silently and runs the host layout.

    Running fast's _check_inputs on the host route would cover all four.

  3. CP contract differs from FLA.

    • Passing the global cu_seqlens together with cp_context (FLA ignores it) gives non-finite o, dq, dk, dg and dbeta; FLA is bit-equal to its local-cu result.
    • output_final_state=True under cp_context: the forward returns a local final state and the backward raises; FLA asserts at entry.
  4. The workspace pool keeps memory. At HV=32, 8 shapes x2, after freeing every tensor the pool still holds 6.46 GiB (fresh allocation: 0). Hit rate is about 47% (72 hits, 154 misses) and step time is 2.6 vs 2.7 ms, so the pool saves little. Consider dropping it, or bounding it and exposing a clear hook for miles' sleep.

  5. The scope of determinism should be stated. det bits differ between the host-table and device-table routes (2 of 4 layouts), and with the SM count (H16/HV32 at 16k: 148 or 160 SMs vs 152 changes the plan and the bits). Segmented vs serial scan also differ in bits, with the same FP64 error (4.15e-3). Suggest documenting it as "same GPU model, same route".

  6. T=0 raises in both det and fast (minor).

Cleared on GPU

  • det CP vs no-CP is within budget (rel_rms ≤ 1.2e-3), and two det CP runs are bit-equal.
  • checkpoint(use_reentrant=False) is bit-exact for det and fast, with save_state on and off.
  • Length-1 and zero-length documents match FP64, including initial and final state.
  • The torch cpp_extension lock does not hang after the builder is SIGKILLed: the second build finished in 51 s. The cake_native lock also recovers (9 s).

Code-level

Across the stack

  • Unconfigured default variant is inconsistent. gdn and dsa default to det; conv1d, gated_rmsnorm, rmsnorm and mhc default to fast, including under deterministic mode.
  • Mid-call nvcc compiles. Several packages compile kernels inside a call (see layers 5, 7, 9 and 11). A single prebuild contract in the README would help.
  • Argument registration will need an owner on the miles side. Each package registers its own argparse flags; miles will need a config owner for them.

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.

@Zhichenzzz

Zhichenzzz commented Nov 5, 2026 •

Copy link
Copy Markdown
Collaborator

Repo layout — please standardize it before splitting into stacked PRs

More kernel sources will come later, so group everything by source under one namespace, with a single layout for every op:

miles_ext_kernels/
  __init__.py              # torch-free, lightweight
  cake/
    _runtime/              # one variant/configure, loader, workspace pool, prebuild
    _vendor/               # one tree, one VENDORED.json
    gdn.py kda.py conv1d.py ... lm_head.py   # one thin wrapper per op
  <future_source>/         # same shape: _runtime/, _vendor/, one wrapper per op
tests/cake/<kernel>/
tools/vendor_cake.py       # the single vendoring entry point

What this fixes at the stack head (2e422c50ca)

  • Vendor paths are inconsistent. Some packages use _vendor/cake/<pkg>/ and others _vendor/<pkg>/. Use one depth everywhere.
  • The manifest is split. There are 11 VENDORED.json files, all pinned to the same flashinfer commit c5db8d3ccd; one is enough.
  • Some sources are unchecked. gdn/csrc and dsa/csrc sit outside every manifest and build through a second system (cpp_extension + _jit.py + registry.json). Move them under _vendor/, list them in the manifest, and build them through the same path as the rest.
  • The runtime is copied per package. Each package has its own arguments.py / configure / module-global VARIANT, and its own prebuild. Move all of this into cake/_runtime/ with one argument surface.

Suggested order: land a "[Cake 0/11] layout" PR first, then put the 11 layers on top of it.

averyhNV and others added 6 commits November 5, 2026 09:01
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>
Move the gdn op onto the cake layout. miles_ext_kernels/cake/gdn.py
declares it (variants det and fast, --linear-attention-backend cake,
--cake-gdn-recompute-state; importing it imports neither torch nor the
kernels); miles_ext_kernels/cake/_impl/gdn/ holds the entry points; the
five vendored packages load through the runtime loader. A call reads the
variant and --cake-gdn-recompute-state when it runs:

- no cp_context: det cake_gdn_det_train, fast cake_gdn_train;
- a KcpContext: cake_gdn_det_train_kcp, cake_gdn_train_kcp;
- an FLA cp_context, or a KcpContext call the KCP packages do not serve:
  cake_gdn_fla_cp_det, the det kernels of FLA's context parallelism,
  under either variant (fast logs it once).

cake_gdn_fla_cp_det builds with nvcc through the runtime loader and loads
with tvm-ffi (native.py); the torch-extension build is gone. prebuild()
builds every library of the five packages for the device's architecture;
a call no longer builds them on its first use. The third-party Triton
kernels that a caller-supplied FLA cp_context runs (FLA's own
context-parallel pre-processing) are not prewarmed: they compile in FLA
on their first call.

Behaviour changes:

- --cake-gdn-variant fast under --deterministic-mode is now rejected,
  and so is configure(..., deterministic_mode=True) with fast: fast does
  not promise its bits, so deterministic mode runs det. Before, fast was
  accepted there.
- Under any cp_context the contract is FLA's: initial_state and
  output_final_state raise ValueError when the call is made, and the
  context's cu_seqlens / cu_seqlens_cpu replace the arguments. Before, a
  KcpContext call with a final state ran the FLA-CP kernels, and a
  cu_seqlens_cpu other than the context's raised.
- The scratch buffers of the det and fast packages and of
  cake_gdn_fla_cp_det go through the runtime's workspace pool instead of
  per-package pools kept until the process ends: at the default
  --cake-workspace-pool-bytes of 0 each is a fresh allocation per call;
  under a positive bound the pool keeps one buffer per role and stream,
  freed by release(). The 0-element state placeholder of a launch
  without an initial or final state holds no memory: it is one tensor
  per device, shape and dtype, not an allocation per call.

Review fixes:

- the FLA-CP state-scan planner never cuts a segment longer than the 128
  chunks one pre-scan CTA covers: a long sequence that cannot be split
  within that keeps the serial scan instead of failing;
- the FLA-CP path checks the extents of every input before a launch, as
  the det package does;
- a call without a token returns an empty output and passes the state
  through, under either variant, also with an FLA context without a
  process group;
- the autograd function of the FLA-CP kernels no longer materialises a
  missing incoming gradient: an unused final state adds no final-state
  gradient (the bits of the call without one), and a loss on the final
  state alone runs with a zero output gradient;
- the scope of det's bits is documented (gdn.py, ops.py).

The gdn extra installs flash-linear-attention>=0.5.2: the FLA-CP path
calls the fla.ops.cp helpers of that release (with their state_v_first
keyword) and is checked against it; 0.5.0's helpers lack the keyword, and
the helpers' module of 0.5.1 differs from 0.5.2's.

Co-Authored-By: Claude <noreply@anthropic.com>
tests/cake/gdn/, on the CPU: the arguments (fast rejected under
deterministic mode), the five packages and their build manifests, the
planner's 128-chunk segment cap, calls without a token, host boundaries
that end before the call's tokens (rejected before any CUDA call), the
state placeholder shared and the packages' other scratch taken from the
workspace pool, and the KCP and FLA context routing on the packages'
torch emulation (FLA's context-parallel contract, a call the KCP package
does not serve logged once).

On a GPU: det and fast against the FP64 recurrence, FLA and each other;
det bit for bit run to run, in a new process, under PyTorch's
deterministic algorithms and on a side stream alternating with the
default one; the extents checked before a launch; one sequence of up to
128k tokens on the det package and on the det kernels with FLA's context
parallelism; the scratch in the workspace pool under a positive bound
(helpers.POOLED, which every pool test configures) at the bits of the
default bound 0, which keeps no pooled byte after a training step; the
det kernels with FLA's context parallelism from the loader's libraries,
where an unused final state passes no gradient (the bits of the call
without one) and a loss on the final state alone backpropagates; the KCP
packages on one GPU and, marked multi_gpu, with context parallelism 2
and 4, where a final state is rejected on every rank when the call is
made. A case that raises on every rank before any communication fails
its own tests with the ranks' tracebacks, and the other cases of the
shared multi-GPU run still report; a case that raises on some ranks only
is not isolated, and the process group's 5-minute timeout bounds the
resulting hang. test_prebuild_gpu.py declares ROUTES, one call per
call-time path, and checks that nothing compiles after one prebuild, and
that a call runs the variant configured when it is made.

fast makes no promise about its bits: its tests compare within the
accuracy budgets only. The recomputed det backward is compared within
the budgets of the saved one (RECOMPUTE_SAME_BITS = False).

The expected scan plans are the cost model's optima at the 148 and 152
SMs of the B200 and GB300 parts: the tests that check them on a device
plan at those SM counts, since a part with another SM count plans other
segment lengths. Where torch.profiler records no CUDA API call (one
known kernel launch shows none either), the launch-count test skips and
names the device.

Co-Authored-By: Claude <noreply@anthropic.com>
@yyihuang

yyihuang commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

Performance (1/5) for [Cake 1/13] gdn, continued from the PR body

Performance

In the copied FlashInfer text, "this PR", "new", "previous" and "main" refer to that FlashInfer pull request and to the FlashInfer package before it, not to this PR. Tables, table captions and quoted paragraphs are copied from the pull request bodies without change; a copied heading is shown in bold.

Packages and variants

This layer vendors four packages from flashinfer-ai/flashinfer eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3 and one package from radixark/miles#3597 at its head 9d7543812b676a5b50dde107518782933b804b48. Every vendored file has the blob of the source tree.

Variant / route Package Source path Source tree Files When this layer calls it
det, single-rank cake_gdn_det_train flashinfer/experimental/cake_gdn_det_train 76c560ea33faf5e78a25dc056344433185903c21 85 det, no cp_context
det, KCP cake_gdn_det_train_kcp flashinfer/experimental/cake_gdn_det_train_kcp d318ca75ffce132e8080208b587dddd0fd1a7262 59 det with a KcpContext; save_mode="full" by default, "light" with --cake-gdn-recompute-state
fast (non-deterministic), single-rank cake_gdn_train flashinfer/experimental/cake/cake_gdn_train 391de6c2718a6773c1b02bffc83704fab25ff75c 80 fast, no cp_context
fast, KCP cake_gdn_train_kcp flashinfer/experimental/cake_gdn_train_kcp d6344ef2be34a8f99efabecf7ccddd22f1208ad3 67 fast with a KcpContext; same save_mode rule. This layer does not pass graph, so calls run the package default graph="off" (eager)
deterministic kernels with FLA's context parallelism cake_gdn_fla_cp_det radixark/miles miles/kernels/attention/gdn/cake (csrc/, registry.json) csrc tree 599cb61345133138d995e9d1744a8bd467a8ee09 57 + 4 layer files an FLA cp_context, or a KcpContext call that the KCP packages do not serve, under either variant

The earlier version of this PR vendored only cake_gdn_train, from FlashInfer c5db8d3ccd3fa3842812280f7f208718dc0b8610 (tree 53fb1746). That tree is not the vendored tree, so its rows are not reused here.

Libraries in build.json: cake_gdn_det_train 36 / 36 / 36 (sm_100a / sm_103a / sm_107a), cake_gdn_det_train_kcp 19 / 19 / 19, cake_gdn_train 35 / 35 / 34, cake_gdn_train_kcp 23 / 23 / 23. min_cuda is 12.8 (sm_100a), 12.9 (sm_103a) and 13.4 (sm_107a) in all four. cake_gdn_fla_cp_det has sm_100a and sm_103a sources; this layer builds its sm_107a libraries from the sm_103a sources.

FlashInfer pull requests of the vendored revision

Every merge commit below is an ancestor of eb3eb68d7deebbe619e05d93ecfd9796e4eaf2e3. Each PR head has the same package tree as its merge commit.

FlashInfer PR Package Merge commit Merged (UTC) Package tree after the merge Relation to the vendored package
flashinfer-ai/flashinfer#8695 det 3b2c21665e6d508c5e4f049ce1f399b92a7bf489 2026-11-04T20:21:30Z 76c560ea (= vendored) exact. It is the last change to the det package before eb3eb68d.
flashinfer-ai/flashinfer#8615 det 35f52e08c412e2490d6e03346da5f23432a1e230 2026-11-04T07:00:22Z 738071dceb53ef047a58d2eed39d6d8c16cc98cd same sm_100a and sm_103a programs. The only later change is flashinfer-ai/flashinfer#8695: it adds #if __CUDA_ARCH__ == 1070 branches to one kernel source and selects that program on sm_107a in the host files. flashinfer-ai/flashinfer#8695 states that the B200 and GB300 compiled programs are identical to FlashInfer main.
flashinfer-ai/flashinfer#8585 det KCP 203cc3c299139b9eea1c0d5417275a779ae38e6e 2026-11-04T04:56:12Z d318ca75 (= vendored) exact. Host-only fix; it is the last change to the det KCP package before eb3eb68d.
flashinfer-ai/flashinfer#8538 det KCP 71b94c30f1f731e0e84d2c75fde26d22e5a1323a 2026-11-03T19:31:55Z 5d9eb839 same kernels. The only later change is flashinfer-ai/flashinfer#8585 (host files cake_kcp.py, cake_bytes.py, README.md).
flashinfer-ai/flashinfer#8675 fast 51f704d0d2ebfd0b73d26c5c2ec68fd18c03c605 2026-11-04T17:45:27Z 391de6c2 (= vendored) exact. It is the last change to the fast package before eb3eb68d. It changes cake_chain.py and README.md only.
flashinfer-ai/flashinfer#8665 fast 8e7be5e3a832fdf5d872cb01b05d6d3c5906b974 2026-11-04T14:23:24Z 572595a4 same kernels. The only later change is flashinfer-ai/flashinfer#8675, which states that every row on sm_107a keeps the previous schedule. Only its R200 rows are used here.
flashinfer-ai/flashinfer#8638 fast KCP 283abf0d4eab8f6ea98d9fbd167fa969561d8aff 2026-11-04T14:49:47Z d6344ef2 (= vendored) exact. Host-only; it changes only the checkpointed step under graph="replay", and states that eager calls launch the same kernels as before.
flashinfer-ai/flashinfer#8644 fast KCP f266f1335f008214bf81639c9b62508ff8952a16 2026-11-04T12:35:07Z 8fc965a2 same kernels. The only later change is flashinfer-ai/flashinfer#8638 (host files). Its tables time eager calls with save_mode="full", the default of this layer.

Baseline identities: the det(main) arm of flashinfer-ai/flashinfer#8644 and flashinfer-ai/flashinfer#8638 is the det KCP package of FlashInfer main f2926f20, whose tree is the vendored d318ca75. The "Deterministic Cake" arm of flashinfer-ai/flashinfer#8675 and flashinfer-ai/flashinfer#8665 is the det package at tree 738071dc: the vendored sm_100a and sm_103a programs, but not the vendored sm_107a forward WY program (flashinfer-ai/flashinfer#8695).

Not vendored: these FlashInfer main changes come after eb3eb68d, and their numbers do not apply to this layer. det KCP: flashinfer-ai/flashinfer#8741 (ab4668c32dd1), flashinfer-ai/flashinfer#8755 (97ce95504d8b). fast single-rank: flashinfer-ai/flashinfer#8776 (6c713b7fcfe8). fast KCP: flashinfer-ai/flashinfer#8754 (8fb7dbf8521c). The det single-rank package has no change after eb3eb68d on FlashInfer main up to 7c774db8c488.

Baselines

Path Previous deterministic Cake version Strongest non-Cake SOTA
det, single-rank flashinfer-ai/flashinfer#8695 (sm_107a) and flashinfer-ai/flashinfer#8615 (sm_100a, sm_103a), each against the det package on FlashInfer main before it. Missing as a ratio. No det pull request reports one. The fast tables of flashinfer-ai/flashinfer#8675 and flashinfer-ai/flashinfer#8665 give det Cake and non-Cake absolute times in the same rows, but print no det / SOTA ratio, and none is computed here.
det, KCP flashinfer-ai/flashinfer#8538 (B/A against the package after flashinfer-ai/flashinfer#8385) and flashinfer-ai/flashinfer#8585 (fix / main). Missing as a ratio. Table B of flashinfer-ai/flashinfer#8644 and of flashinfer-ai/flashinfer#8638 gives det(main) and non-Cake absolute times in the same rows, with no det / SOTA ratio.
fast, single-rank Deterministic Cake (and the previous fast package): flashinfer-ai/flashinfer#8675, flashinfer-ai/flashinfer#8665. Per cell, the fastest accuracy-qualified of cuDNN frontend 1.32.1 FROST (states recomputed, or stored every 64 tokens) and FLA 0.6.0 (Triton): flashinfer-ai/flashinfer#8675, flashinfer-ai/flashinfer#8665.
fast, KCP det(main), the vendored det KCP package: flashinfer-ai/flashinfer#8644, flashinfer-ai/flashinfer#8638. CP1: head split around cuDNN frontend 1.32.1 FROST or FLA, or FLA-CP (FLA 0.6.0), the fastest accuracy-qualified arm. CP > 1: FlashQLA 0.1.3 with its dual context parallelism (B200 CP8, GB300 CP4); not claimed for layout D and on R200 (no FlashQLA sm_107a build). Sources: flashinfer-ai/flashinfer#8644, flashinfer-ai/flashinfer#8638.
det kernels with FLA's context parallelism Kernel-level gains per optimization round against the previous round (radixark/miles#3597). FLA (fla 0.5.2) in the same whole-model training run on the same node: radixark/miles#3597. No kernel-level comparison with a non-Cake implementation is reported.

Per-architecture coverage

Path sm_100a sm_103a sm_107a
det, single-rank measured, same programs: flashinfer-ai/flashinfer#8615 B200 (148 SMs), CUDA 13.4 measured, same programs: flashinfer-ai/flashinfer#8615 GB300 (152 SMs), CUDA 13.4 measured, exact: flashinfer-ai/flashinfer#8695 R200 with 216, 212 and 208 SMs, CUDA 13.4
det, KCP measured: flashinfer-ai/flashinfer#8585 (exact) and flashinfer-ai/flashinfer#8538 on GB200, CUDA 13.4 toolkit; det(main) absolute times on B200 in flashinfer-ai/flashinfer#8644 and flashinfer-ai/flashinfer#8638 measured: flashinfer-ai/flashinfer#8585 and flashinfer-ai/flashinfer#8538 on GB300; det(main) absolute times in flashinfer-ai/flashinfer#8644 and flashinfer-ai/flashinfer#8638 no comparison against an earlier det KCP version (flashinfer-ai/flashinfer#8538: "R200 (sm_107a) rows were not measured"); det(main) absolute times on a 212-SM R200 in flashinfer-ai/flashinfer#8644 and flashinfer-ai/flashinfer#8638
fast, single-rank measured, exact: flashinfer-ai/flashinfer#8675 B200 (148 SMs), CUDA 13.0 measured, exact: flashinfer-ai/flashinfer#8675 GB300 (152 SMs), CUDA 13.0 measured, same programs and schedule: flashinfer-ai/flashinfer#8665 R200 (212 SMs), CUDA 13.0, FP32 beta only
fast, KCP measured: flashinfer-ai/flashinfer#8644 (eager, full save) and flashinfer-ai/flashinfer#8638 (exact; graph replay) on B200 measured: same two PRs on GB300 measured: same two PRs on a 212-SM R200 (earlier revisions of flashinfer-ai/flashinfer#8638 also on a 208-SM R200)
det kernels with FLA's context parallelism measured end to end: radixark/miles#3597, one 8x B200 node measured end to end: radixark/miles#3597, one 4x GB300 node not measured. radixark/miles#3597 targets SM100a / SM103a; this layer builds sm_107a from the sm_103a sources

None of the pull requests above reports B300 numbers.

Deterministic, single-rank, sm_107a (exact revision)

Source: flashinfer-ai/flashinfer#8695. Baseline: the det package on FlashInfer main before it. No non-Cake baseline.

flashinfer.experimental.cake_gdn_det_train is FlashInfer's deterministic Gated DeltaNet chunked training kernel. Its outputs and gradients are bitwise reproducible. flashinfer-ai/flashinfer#8615 added a forward WY program that runs five CTAs per SM on sm_100a (B200) and sm_103a (GB300). That program uses packed lower-triangular L^T / X tiles and 64 tensor-memory columns, which brings each CTA to about 43.6 KB of shared memory and 87 registers. sm_107a (R200) kept the four-CTA program. This PR builds the same program for sm_107a and uses it for every forward WY launch on R200, both in the forward pass and in the backward recompute. On R200 it also uses 87 registers and 43,648 B of shared memory per CTA and runs five CTAs per SM. As a result, every row below changes its launches on R200. Outputs and gradients are bitwise identical to the previous version. The baseline is the previous version of this kernel on current FlashInfer main (speedup = main / this PR, so >1 is faster; median over 56 interleaved clean-process cells). Forward + backward is 1.014x to 1.021x faster on 216 SMs, 1.011x to 1.019x faster on 212 SMs and 1.008x to 1.022x faster on 208 SMs. The forward pass is 1.034x to 1.054x, 1.026x to 1.047x and 1.019x to 1.058x faster on those boards, respectively. Every row is faster on all three boards in both f+b and fwd, and every 95 % confidence interval excludes 1.0.

NVIDIA R200 (sm_107a, 216 SMs), CUDA 13.4 (torch 2.14.1+cu130): BF16 inputs, head dim 128, chunk 64; forward + backward (f+b) and forward (fwd) per call, ms (median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR.

sequences H / HV f+b main ms f+b PR ms f+b speedup fwd main ms fwd PR ms fwd speedup
4 x 4096 8 / 16 0.3726 0.3662 1.017x 0.1548 0.1486 1.042x
3 x ~4.25k (real pack) 8 / 16 0.3144 0.3093 1.016x 0.1319 0.1270 1.039x
8 x 4096 8 / 16 0.6968 0.6859 1.016x 0.2890 0.2779 1.040x
7 x ~4.3k (real pack) 8 / 16 0.6523 0.6428 1.015x 0.2723 0.2633 1.034x
16 x 4096 8 / 16 1.2521 1.2298 1.018x 0.5143 0.4892 1.051x
15 x ~4.3k (real pack) 8 / 16 1.2981 1.2796 1.014x 0.5395 0.5164 1.045x
1 x 16384 8 / 16 0.3876 0.3814 1.016x 0.1714 0.1651 1.038x
4 x 4096 16 / 32 0.7027 0.6921 1.015x 0.2936 0.2830 1.037x
16 x 4096 16 / 32 2.4627 2.4132 1.021x 1.0221 0.9697 1.054x
1 x 16384 16 / 32 0.6763 0.6656 1.016x 0.2811 0.2707 1.038x
64 x 4096 2 / 6 1.8187 1.7838 1.020x 0.6949 0.6600 1.053x
11000 + 12000 + 9768 16 / 48 1.9661 1.9295 1.019x 0.7597 0.7232 1.050x
6 seqs, 241817 tokens 2 / 6 1.7635 1.7303 1.019x 0.6866 0.6536 1.050x
2 seqs, 259174 tokens 2 / 6 1.8588 1.8253 1.018x 0.7187 0.6854 1.049x
131071 + 87531 + 43385 2 / 6 1.8746 1.8403 1.019x 0.7237 0.6897 1.049x

NVIDIA R200 (sm_107a, 212 SMs), CUDA 13.4 (torch 2.14.1+cu130): BF16 inputs, head dim 128, chunk 64; forward + backward (f+b) and forward (fwd) per call, ms (median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR.

sequences H / HV f+b main ms f+b PR ms f+b speedup fwd main ms fwd PR ms fwd speedup
4 x 4096 8 / 16 0.3519 0.3458 1.018x 0.1454 0.1397 1.041x
3 x ~4.25k (real pack) 8 / 16 0.2952 0.2920 1.011x 0.1234 0.1201 1.027x
8 x 4096 8 / 16 0.6584 0.6477 1.017x 0.2690 0.2583 1.041x
7 x ~4.3k (real pack) 8 / 16 0.6037 0.5971 1.011x 0.2480 0.2416 1.026x
16 x 4096 8 / 16 1.1631 1.1410 1.019x 0.4667 0.4458 1.047x
15 x ~4.3k (real pack) 8 / 16 1.2149 1.1932 1.018x 0.4912 0.4702 1.045x
1 x 16384 8 / 16 0.3854 0.3799 1.014x 0.1752 0.1696 1.033x
4 x 4096 16 / 32 0.6661 0.6555 1.016x 0.2740 0.2636 1.039x
16 x 4096 16 / 32 2.2951 2.2550 1.018x 0.9266 0.8881 1.043x
1 x 16384 16 / 32 0.6405 0.6308 1.015x 0.2606 0.2503 1.041x
64 x 4096 2 / 6 1.8575 1.8305 1.015x 0.6417 0.6153 1.043x
11000 + 12000 + 9768 16 / 48 2.0291 1.9937 1.018x 0.7538 0.7235 1.042x
6 seqs, 241817 tokens 2 / 6 1.7852 1.7592 1.015x 0.6301 0.6043 1.043x
2 seqs, 259174 tokens 2 / 6 1.8804 1.8529 1.015x 0.6609 0.6334 1.043x
131071 + 87531 + 43385 2 / 6 1.8940 1.8687 1.014x 0.6627 0.6368 1.041x

NVIDIA R200 (sm_107a, 208 SMs), CUDA 13.4 (torch 2.14.1+cu130): BF16 inputs, head dim 128, chunk 64; forward + backward (f+b) and forward (fwd) per call, ms (median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR.

sequences H / HV f+b main ms f+b PR ms f+b speedup fwd main ms fwd PR ms fwd speedup
4 x 4096 8 / 16 0.3860 0.3806 1.014x 0.1620 0.1566 1.034x
3 x ~4.25k (real pack) 8 / 16 0.3204 0.3177 1.008x 0.1348 0.1323 1.019x
8 x 4096 8 / 16 0.7225 0.7114 1.016x 0.2974 0.2865 1.038x
7 x ~4.3k (real pack) 8 / 16 0.6724 0.6608 1.018x 0.2794 0.2681 1.042x
16 x 4096 8 / 16 1.2955 1.2679 1.022x 0.5331 0.5041 1.058x
15 x ~4.3k (real pack) 8 / 16 1.3215 1.2967 1.019x 0.5404 0.5141 1.051x
1 x 16384 8 / 16 0.4041 0.3986 1.014x 0.1776 0.1723 1.031x
4 x 4096 16 / 32 0.7305 0.7208 1.013x 0.3037 0.2943 1.032x
16 x 4096 16 / 32 2.5679 2.5173 1.020x 1.0662 1.0149 1.051x
1 x 16384 16 / 32 0.7066 0.6968 1.014x 0.2967 0.2870 1.034x
64 x 4096 2 / 6 1.8895 1.8601 1.016x 0.7182 0.6884 1.043x
11000 + 12000 + 9768 16 / 48 2.0701 2.0345 1.017x 0.8312 0.7945 1.046x
6 seqs, 241817 tokens 2 / 6 1.8392 1.8057 1.019x 0.7172 0.6839 1.049x
2 seqs, 259174 tokens 2 / 6 1.9383 1.9054 1.017x 0.7453 0.7123 1.046x
131071 + 87531 + 43385 2 / 6 1.9684 1.9342 1.018x 0.7631 0.7288 1.047x

NVIDIA B200 (sm_100a) and NVIDIA GB300 (sm_103a): unchanged; the compiled programs are identical to FlashInfer main, the launch trace is identical, and outputs are bitwise equal.

Deterministic, single-rank, sm_100a and sm_103a

Source: flashinfer-ai/flashinfer#8615. Baseline: the det package on FlashInfer main before it. No non-Cake baseline. Its sm_100a and sm_103a programs are the vendored ones (see above).

flashinfer.experimental.cake_gdn_det_train is FlashInfer's deterministic Gated DeltaNet chunked training kernel. Its outputs and gradients are bitwise reproducible. The forward WY stage builds the per-chunk WY representation (w and u). The forward pass uses it, and the backward pass recomputes it. This PR adds a new compiled program for that stage on sm_100a (B200) and sm_103a (GB300) that runs five CTAs per SM instead of four. It keeps the triangular L^T / X tiles in a packed lower-triangular layout and uses 64 tensor-memory columns, so each CTA needs about 43.6 KB of shared memory and 87 registers. On B200 and GB300, every launch of this stage uses the new program, both in the forward pass and in the backward recompute. As a result, every row below changes its launches. sm_107a (R200) is unchanged: its compiled programs and launches are identical to FlashInfer main. All other stages are unchanged, and outputs and gradients are bitwise identical to the previous version. The baseline is the previous version of this kernel on current FlashInfer main (speedup = main / this PR, so >1 is faster; median over 56 interleaved clean-process cells). The forward pass is faster on every row on both GPUs: 1.028x to 1.044x on B200 and 1.025x to 1.046x on GB300. For forward + backward, B200 measures 1.007x to 1.015x on 14 rows and 1.000x on 3 x ~4.25k 8 / 16. GB300 measures 1.008x to 1.018x on 14 rows and 0.990x on 1 x 16384 8 / 16. In this run, the 95 % confidence interval of the f+b ratio includes 1.0 on seven B200 rows and two GB300 rows, so those rows are within measurement noise here. The B200 rows are 4 x 4096 8 / 16, 3 x ~4.25k 8 / 16, 15 x ~4.3k 8 / 16, 1 x 16384 8 / 16, 16 x 4096 16 / 32, 11000 + 12000 + 9768 16 / 48, and 2 seqs 259174 tokens 2 / 6. The GB300 rows are 4 x 4096 8 / 16 and 1 x 16384 8 / 16. A repeat measurement of exactly those rows, shown in the last table, finds f+b faster on both GB300 rows and on four of the seven B200 rows. The other three B200 rows remain within noise. No row in either run was slower beyond measurement noise.

NVIDIA B200 (sm_100a, 148 SMs), CUDA 13.4 (torch 2.14.1+cu130): BF16 inputs, head dim 128, chunk 64; forward + backward (f+b) and forward (fwd) per call, ms (median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR.

sequences H / HV f+b main ms f+b PR ms f+b speedup fwd main ms fwd PR ms fwd speedup
4 x 4096 8 / 16 0.6380 0.6301 1.013x 0.2315 0.2253 1.028x
3 x ~4.25k (real pack) 8 / 16 0.7514 0.7516 1.000x 0.1997 0.1926 1.037x
8 x 4096 8 / 16 1.0527 1.0378 1.014x 0.4106 0.3944 1.041x
7 x ~4.3k (real pack) 8 / 16 1.0614 1.0475 1.013x 0.4363 0.4220 1.034x
16 x 4096 8 / 16 2.0858 2.0563 1.014x 0.7876 0.7579 1.039x
15 x ~4.3k (real pack) 8 / 16 2.0979 2.0766 1.010x 0.7876 0.7560 1.042x
1 x 16384 8 / 16 0.6603 0.6554 1.007x 0.2376 0.2294 1.036x
4 x 4096 16 / 32 1.0599 1.0455 1.014x 0.4148 0.4004 1.036x
16 x 4096 16 / 32 4.2880 4.2485 1.009x 1.5572 1.5095 1.032x
1 x 16384 16 / 32 1.0925 1.0773 1.014x 0.4280 0.4136 1.035x
64 x 4096 2 / 6 3.2328 3.1838 1.015x 1.1186 1.0779 1.038x
11000 + 12000 + 9768 16 / 48 3.2952 3.2670 1.009x 1.1499 1.1147 1.032x
6 seqs, 241817 tokens 2 / 6 3.0632 3.0236 1.013x 1.0865 1.0414 1.043x
2 seqs, 259174 tokens 2 / 6 3.2855 3.2379 1.015x 1.1661 1.1167 1.044x
131071 + 87531 + 43385 2 / 6 3.3140 3.2727 1.013x 1.1787 1.1301 1.043x

NVIDIA GB300 (sm_103a, 152 SMs), CUDA 13.4 (torch 2.14.1+cu130; driver 580.159.03): BF16 inputs, head dim 128, chunk 64; forward + backward (f+b) and forward (fwd) per call, ms (median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR.

sequences H / HV f+b main ms f+b PR ms f+b speedup fwd main ms fwd PR ms fwd speedup
4 x 4096 8 / 16 0.5339 0.5298 1.008x 0.2171 0.2109 1.029x
3 x ~4.25k (real pack) 8 / 16 0.4848 0.4792 1.012x 0.1874 0.1818 1.031x
8 x 4096 8 / 16 0.9815 0.9647 1.017x 0.3860 0.3692 1.046x
7 x ~4.3k (real pack) 8 / 16 0.9810 0.9636 1.018x 0.3861 0.3696 1.045x
16 x 4096 8 / 16 1.8945 1.8651 1.016x 0.7304 0.7008 1.042x
15 x ~4.3k (real pack) 8 / 16 1.9216 1.8946 1.014x 0.7325 0.7058 1.038x
1 x 16384 8 / 16 0.5584 0.5640 0.990x 0.2207 0.2140 1.031x
4 x 4096 16 / 32 0.9875 0.9716 1.016x 0.3902 0.3742 1.043x
16 x 4096 16 / 32 3.7229 3.6875 1.010x 1.4288 1.3940 1.025x
1 x 16384 16 / 32 1.0394 1.0242 1.015x 0.4032 0.3871 1.042x
64 x 4096 2 / 6 2.9181 2.8858 1.011x 1.0335 1.0006 1.033x
11000 + 12000 + 9768 16 / 48 2.9767 2.9482 1.010x 1.0637 1.0313 1.031x
6 seqs, 241817 tokens 2 / 6 2.7746 2.7341 1.015x 1.0030 0.9626 1.042x
2 seqs, 259174 tokens 2 / 6 2.9604 2.9174 1.015x 1.0675 1.0245 1.042x
131071 + 87531 + 43385 2 / 6 2.9922 2.9491 1.015x 1.0864 1.0435 1.041x

Repeat measurement of the rows within noise above, on NVIDIA B200 (sm_100a, 148 SMs) and NVIDIA GB300 (sm_103a, 152 SMs), CUDA 13.4 (torch 2.14.1+cu130): fresh run on another node of the same GPU type, same settings (BF16 inputs, head dim 128, chunk 64; f+b per call, ms; median over 56 interleaved clean-process cells); speedup = FlashInfer main / this PR. Rows marked "within noise" have a 95 % confidence interval that includes 1.0.

GPU sequences H / HV f+b main ms f+b PR ms f+b speedup note
B200 4 x 4096 8 / 16 0.6381 0.6308 1.012x
B200 3 x ~4.25k (real pack) 8 / 16 0.7462 0.6929 1.077x within noise
B200 15 x ~4.3k (real pack) 8 / 16 2.0776 2.0430 1.017x
B200 1 x 16384 8 / 16 0.6574 0.6492 1.013x within noise
B200 16 x 4096 16 / 32 4.1558 4.1067 1.012x within noise
B200 11000 + 12000 + 9768 16 / 48 3.2584 3.2253 1.010x
B200 2 seqs, 259174 tokens 2 / 6 3.2587 3.2057 1.017x
GB300 4 x 4096 8 / 16 0.5329 0.5266 1.012x
GB300 1 x 16384 8 / 16 0.5513 0.5442 1.013x

@yyihuang

yyihuang commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

Performance (2/5), continued

As in part 1, "this PR", "new", "previous" and "main" in the copied FlashInfer text refer to that FlashInfer pull request and to the FlashInfer package before it, not to this PR.

Deterministic KCP

Exact revision: flashinfer-ai/flashinfer#8585

This is a correctness fix with no kernel or launch change. Its tables compare the fix with FlashInfer main before it.

With save_mode="full" and save_transfer=False on a rank whose fragments run in P > 1 segments, gradients were wrong; out was not affected. This includes exchange="ring", because the ring requires save_transfer=False. The backward merges read the per-segment transfer plane M_seg, but the forward drops that buffer and the full saved set kept M_seg only under save_transfer, so the merges read memory that no program of the call had written. The fix is host-only, with no kernel or launch change. The full saved set now keeps M_seg at P > 1 for either save_transfer value, and a full saved set without it is refused. gdn_det_saved_bytes_kcp / gdn_det_saved_layout_kcp now count M_seg as P * n_frag states, which is the whole n_frag * P plane that is kept, instead of (P - 1) * n_frag. Defaults (save_mode="light", save_transfer=True) are unaffected, and the light rows are bitwise equal to main in the tables below.

8 GPUs (2 nodes x 4), context parallel 8; four documents of 98,304 / 65,536 / 57,344 / 40,960 tokens, 32,768 tokens per rank; (H_qk, HV) = (16, 48); median ms of the slowest rank over interleaved main/fix visits; CUDA 13.4 toolkit, PyTorch 2.14.1+cu130.

GB300 (sm_103a)

Layout In-kernel L2 norm Scope Form main ms fix ms fix/main
contiguous off forward + checkpoint recompute + backward save_mode="full", save_transfer=False 14.918 14.914 0.9997
contiguous off forward + checkpoint recompute + backward save_mode="light", save_transfer=False 14.916 14.915 0.9999
contiguous on forward shipped defaults 2.709 2.710 1.0004
contiguous on forward save_mode="light", save_transfer=False 2.708 2.712 1.0015
contiguous on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 7.700 7.698 0.9998
contiguous on forward + checkpoint recompute + backward shipped defaults 8.195 8.166 0.9965
contiguous on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.599 7.590 0.9988
per-document zigzag off forward save_mode="light", save_transfer=False 3.973 3.971 0.9995
per-document zigzag off forward + checkpoint recompute + backward save_mode="full", save_transfer=False 14.880 14.874 0.9996
per-document zigzag off forward + checkpoint recompute + backward save_mode="light", save_transfer=False 14.882 14.884 1.0001
per-document zigzag on forward save_mode="light", save_transfer=False 2.748 2.750 1.0007
per-document zigzag on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 7.674 7.674 0.9999
per-document zigzag on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.564 7.568 1.0005
zigzag on forward save_mode="light", save_transfer=False 2.821 2.821 0.9999
zigzag on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 7.935 7.930 0.9994
zigzag on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.824 7.825 1.0002

On the contiguous layout, P = 3 on every rank, so main's gradients are wrong in the save_mode="full", save_transfer=False rows. The main time in those rows is only a cost reference.

GB200 (sm_100a)

Layout In-kernel L2 norm Scope Form main ms fix ms fix/main
contiguous off forward + checkpoint recompute + backward save_mode="full", save_transfer=False 15.164 15.161 0.9998
contiguous off forward + checkpoint recompute + backward save_mode="light", save_transfer=False 15.157 15.161 1.0003
contiguous on forward shipped defaults 2.741 2.741 0.9998
contiguous on forward save_mode="light", save_transfer=False 2.744 2.743 0.9996
contiguous on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 7.840 7.840 1.0001
contiguous on forward + checkpoint recompute + backward shipped defaults 8.165 8.183 1.0021
contiguous on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.731 7.734 1.0004
per-document zigzag off forward save_mode="light", save_transfer=False 4.025 4.028 1.0006
per-document zigzag off forward + checkpoint recompute + backward save_mode="full", save_transfer=False 15.231 15.226 0.9997
per-document zigzag off forward + checkpoint recompute + backward save_mode="light", save_transfer=False 15.223 15.223 1.0000
per-document zigzag on forward save_mode="light", save_transfer=False 2.792 2.795 1.0012
per-document zigzag on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 7.895 7.894 0.9998
per-document zigzag on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.794 7.799 1.0006
zigzag on forward save_mode="light", save_transfer=False 2.855 2.855 1.0001
zigzag on forward + checkpoint recompute + backward save_mode="full", save_transfer=False 8.084 8.086 1.0002
zigzag on forward + checkpoint recompute + backward save_mode="light", save_transfer=False 7.994 7.994 0.9999

Same kernels: flashinfer-ai/flashinfer#8538

A = the det KCP package after flashinfer-ai/flashinfer#8385; B = flashinfer-ai/flashinfer#8538. B/A below 1 means B is faster.

Cake's deterministic context-parallel GDN training kernel has one data-type variant: BF16 q/k/v/out/dout with FP32 scalar log-decay g and beta. The grouped Q/K head counts (H_qk, HV) are runtime parameters with production shape (16, 48). The kernel has three partition layouts: contiguous (C), whole-buffer zigzag (Z) and per-document zigzag + ragged (D). This PR delivers the same variant. It changes the host-side rule that picks the value-head tile width for the cross-rank state-summary walk launches. A walk runs either 128 value columns per CTA or 64, and the 64 form runs two CTAs per segment and head. Forward walks still use the 64 form when the grid fits one wave. They now also use it when an LPT schedule of the walk's segment chain lengths shows that the two-CTA grid finishes within its longest chain. Backward offset walks always use the 64 form. Backward tape walks use it on sm_100a and sm_103a, while sm_107a keeps the previous two-wave rule. One exception applies on sm_100a: a backward walk keeps the previous two-wave rule when its longest chain is at least 400 chunks and its two-CTA schedule would take longer than the one-CTA schedule. The rule reads only host plan integers (segment lengths, head count, SM count and device architecture), never tensor contents. It takes the same decision in the forward, the checkpoint recompute and the recompute-handle replay. The non-deterministic KCP package (PR flashinfer-ai/flashinfer#8521) uses the same thresholds. The change is host-side Python only and the kernels are unchanged. Only launch widths differ, so outputs and gradients are bitwise identical to before on every measured row, and peak memory is unchanged. A is the previous Cake version of this package, as shipped after PR flashinfer-ai/flashinfer#8385. B is this PR on the same base. B/A < 1 means faster. Each GPU was measured on 66 rows: packs P1-P5 x CP 2/4/8 x layouts C/Z/D at B1 and B3. On 33 of these rows the launch form changes, and the other 33 are unchanged controls. On GB300 (sm_103a), changed rows range from 0.984 to 1.004 at B3 (median 0.989) and from 0.976 to 1.001 at B1 (median 0.999), though only 6 B1 rows change. Unchanged GB300 rows range from 0.994 to 1.013 (median 1.000). On GB200 (sm_100a), changed rows range from 0.987 to 1.016 at B3 (median 0.992) and from 0.978 to 1.002 at B1 (median 1.000). Unchanged GB200 rows range from 0.992 to 1.003 (median 1.000). Taking the median over both sessions, changed rows at CP 2 / 4 / 8 come to 0.989 / 0.990 / 0.999 on GB300 and 0.991 / 0.994 / 1.000 on GB200. That is about a 1 % gain at CP 2 and CP 4 and about parity at CP 8. The largest gain is CP 8 P3 C B1 on GB300 at 0.977 / 0.976. The slowest changed row is CP 8 P3 C B3 on GB200, at 1.016 in session 1 and 1.006 in session 2. R200 (sm_107a) rows were not measured.

ms = slowest-rank median; B1 = module forward with autograd recording; B3 = forward + non-reentrant checkpoint replay + backward, including all communication and copies; Walk forms = changed (this PR launches a different width) or unchanged (same launches as A); s1/s2 = two sessions.

Packs: P1 241,817 tokens / 6 documents, P2 261,987 / 3, P3 259,174 / 2, P4 254,976 / 1 (one document spans the whole group), P5 249,856 / 4. Heads (16, 48); causal conv1d width 4 over 10,240 channels.

Software: CUDA 13.4 toolkit, torch 2.14.1+cu130.

NVIDIA GB300 (sm_103a)

CP Pack Layout Walk forms Boundary A ms s1 B ms s1 B/A s1 A ms s2 B ms s2 B/A s2
2 P1 C unchanged B1 7.884 7.898 1.002 7.863 7.865 1.000
2 P1 C changed B3 23.793 23.611 0.992 23.778 23.580 0.992
2 P1 Z unchanged B1 7.753 7.751 1.000 7.757 7.758 1.000
2 P1 Z changed B3 23.728 23.471 0.989 23.740 23.487 0.989
2 P1 D changed B1 7.919 7.913 0.999 7.919 7.916 1.000
2 P1 D changed B3 23.999 23.942 0.998 24.014 23.949 0.997
2 P2 C unchanged B1 8.441 8.449 1.001 8.438 8.443 1.001
2 P2 C changed B3 25.784 25.516 0.990 25.786 25.504 0.989
2 P2 Z unchanged B1 8.383 8.378 0.999 8.398 8.394 1.000
2 P2 Z changed B3 25.590 25.186 0.984 25.609 25.219 0.985
2 P2 D unchanged B1 8.553 8.555 1.000 8.552 8.550 1.000
2 P2 D changed B3 26.260 25.933 0.988 26.249 25.923 0.988
2 P3 C unchanged B1 7.386 7.379 0.999 7.376 7.373 1.000
2 P3 C unchanged B3 25.264 25.261 1.000 25.250 25.240 1.000
2 P3 Z unchanged B1 9.289 9.277 0.999 9.296 9.291 0.999
2 P3 Z changed B3 28.269 27.966 0.989 28.271 27.982 0.990
2 P3 D unchanged B1 8.248 8.237 0.999 8.240 8.245 1.001
2 P3 D changed B3 25.282 24.942 0.987 25.306 24.968 0.987
2 P4 C unchanged B1 8.273 8.284 1.001 8.271 8.253 0.998
2 P4 C changed B3 25.198 24.996 0.992 25.140 24.916 0.991
2 P4 Z unchanged B1 8.463 8.461 1.000 8.454 8.451 1.000
2 P4 Z changed B3 25.718 25.421 0.988 25.734 25.436 0.988
2 P4 D unchanged B1 8.516 8.521 1.001 8.509 8.510 1.000
2 P4 D changed B3 25.838 25.559 0.989 25.839 25.549 0.989
2 P5 C unchanged B1 8.076 8.080 1.000 8.126 8.080 0.994
2 P5 C changed B3 24.852 24.492 0.986 24.835 24.456 0.985
2 P5 Z unchanged B1 8.721 8.712 0.999 8.729 8.724 0.999
2 P5 Z changed B3 26.660 26.375 0.989 26.633 26.332 0.989
2 P5 D unchanged B1 8.199 8.193 0.999 8.233 8.218 0.998
2 P5 D changed B3 25.025 24.737 0.988 25.056 24.753 0.988
4 P1 C unchanged B1 4.360 4.366 1.001 4.352 4.357 1.001
4 P1 C changed B3 12.985 12.987 1.000 12.984 12.985 1.000
4 P1 Z unchanged B1 4.211 4.204 0.998 4.199 4.193 0.999
4 P1 Z changed B3 12.616 12.542 0.994 12.597 12.500 0.992
4 P1 D changed B1 4.325 4.326 1.000 4.321 4.326 1.001
4 P1 D changed B3 13.005 12.996 0.999 13.005 12.999 1.000
4 P2 Z unchanged B1 4.493 4.495 1.000 4.486 4.471 0.997
4 P2 Z unchanged B3 13.538 13.529 0.999 13.518 13.511 0.999
4 P2 D unchanged B1 4.837 4.840 1.000 4.840 4.844 1.001
4 P2 D changed B3 14.439 14.259 0.988 14.464 14.294 0.988
4 P3 C unchanged B1 4.364 4.368 1.001 4.369 4.366 0.999
4 P3 C changed B3 13.225 13.056 0.987 13.213 13.069 0.989
4 P3 Z unchanged B1 4.360 4.356 0.999 4.360 4.352 0.998
4 P3 Z changed B3 13.184 13.018 0.987 13.183 13.016 0.987
4 P3 D unchanged B1 4.691 4.689 1.000 4.685 4.686 1.000
4 P3 D changed B3 14.156 14.080 0.995 14.162 14.082 0.994
4 P4 D unchanged B1 4.427 4.438 1.002 4.426 4.424 1.000
4 P4 D unchanged B3 13.239 13.217 0.998 13.226 13.213 0.999
4 P5 C unchanged B1 4.229 4.218 0.998 4.228 4.229 1.000
4 P5 C changed B3 12.783 12.638 0.989 12.790 12.654 0.989
4 P5 Z unchanged B1 4.253 4.254 1.000 4.257 4.255 0.999
4 P5 Z changed B3 12.862 12.666 0.985 12.859 12.672 0.985
4 P5 D changed B1 4.360 4.358 0.999 4.341 4.341 1.000
4 P5 D changed B3 13.211 13.046 0.988 13.187 13.016 0.987
8 P1 C unchanged B1 2.504 2.536 1.013 2.526 2.527 1.000
8 P1 C unchanged B3 7.880 7.868 0.998 7.880 7.888 1.001
8 P1 D changed B1 2.651 2.649 0.999 2.650 2.647 0.999
8 P1 D changed B3 7.756 7.785 1.004 7.655 7.662 1.001
8 P2 Z unchanged B1 2.692 2.689 0.999 2.695 2.701 1.002
8 P2 Z unchanged B3 7.770 7.767 1.000 7.751 7.751 1.000
8 P3 C changed B1 2.636 2.575 0.977 2.616 2.553 0.976
8 P3 C changed B3 7.973 7.971 1.000 7.830 7.839 1.001
8 P4 D unchanged B1 2.752 2.758 1.002 2.759 2.765 1.002
8 P4 D unchanged B3 8.348 8.380 1.004 8.189 8.210 1.003
8 P5 D changed B1 2.585 2.583 0.999 2.579 2.579 1.000
8 P5 D changed B3 7.550 7.492 0.992 7.533 7.482 0.993

NVIDIA GB200 (sm_100a)

CP Pack Layout Walk forms Boundary A ms s1 B ms s1 B/A s1 A ms s2 B ms s2 B/A s2
2 P1 C unchanged B1 8.016 8.012 0.999 8.032 8.033 1.000
2 P1 C changed B3 24.233 24.104 0.995 24.238 24.106 0.995
2 P1 Z unchanged B1 7.890 7.892 1.000 7.908 7.904 1.000
2 P1 Z changed B3 24.176 23.980 0.992 24.205 24.016 0.992
2 P1 D changed B1 8.072 8.065 0.999 8.073 8.077 1.001
2 P1 D changed B3 24.448 24.438 1.000 24.498 24.517 1.001
2 P2 C unchanged B1 8.600 8.600 1.000 8.606 8.606 1.000
2 P2 C changed B3 26.247 26.035 0.992 26.265 26.046 0.992
2 P2 Z unchanged B1 8.555 8.542 0.998 8.548 8.555 1.001
2 P2 Z changed B3 26.113 25.777 0.987 26.141 25.797 0.987
2 P2 D unchanged B1 8.709 8.716 1.001 8.715 8.720 1.001
2 P2 D changed B3 26.728 26.464 0.990 26.733 26.471 0.990
2 P3 C unchanged B1 7.524 7.525 1.000 7.543 7.545 1.000
2 P3 C unchanged B3 26.088 26.093 1.000 26.127 26.160 1.001
2 P3 Z unchanged B1 9.440 9.433 0.999 9.443 9.445 1.000
2 P3 Z changed B3 28.722 28.493 0.992 28.719 28.500 0.992
2 P3 D unchanged B1 8.386 8.401 1.002 8.397 8.396 1.000
2 P3 D changed B3 25.754 25.473 0.989 25.767 25.503 0.990
2 P4 C unchanged B1 8.418 8.412 0.999 8.415 8.428 1.002
2 P4 C changed B3 25.649 25.585 0.997 25.636 25.567 0.997
2 P4 Z unchanged B1 8.594 8.596 1.000 8.606 8.597 0.999
2 P4 Z changed B3 26.209 25.977 0.991 26.223 25.997 0.991
2 P4 D unchanged B1 8.651 8.656 1.001 8.648 8.654 1.001
2 P4 D changed B3 26.299 26.056 0.991 26.317 26.081 0.991
2 P5 C unchanged B1 8.245 8.230 0.998 8.251 8.251 1.000
2 P5 C changed B3 25.309 24.990 0.987 25.348 25.042 0.988
2 P5 Z unchanged B1 8.840 8.836 1.000 8.847 8.848 1.000
2 P5 Z changed B3 27.040 26.783 0.990 27.068 26.800 0.990
2 P5 D unchanged B1 8.340 8.334 0.999 8.344 8.342 1.000
2 P5 D changed B3 25.501 25.251 0.990 25.529 25.282 0.990
4 P1 C unchanged B1 4.475 4.477 1.000 4.442 4.436 0.999
4 P1 C changed B3 13.335 13.341 1.000 13.274 13.268 1.000
4 P1 Z unchanged B1 4.291 4.296 1.001 4.254 4.252 0.999
4 P1 Z changed B3 12.857 12.799 0.996 12.797 12.765 0.997
4 P1 D changed B1 4.412 4.410 1.000 4.432 4.424 0.998
4 P1 D changed B3 13.268 13.283 1.001 13.272 13.265 0.999
4 P2 Z unchanged B1 4.586 4.590 1.001 4.579 4.581 1.000
4 P2 Z unchanged B3 13.855 13.848 1.000 13.865 13.866 1.000
4 P2 D unchanged B1 4.916 4.915 1.000 4.899 4.899 1.000
4 P2 D changed B3 14.685 14.523 0.989 14.667 14.505 0.989
4 P3 C unchanged B1 4.446 4.454 1.002 4.460 4.472 1.003
4 P3 C changed B3 13.470 13.334 0.990 13.484 13.327 0.988
4 P3 Z unchanged B1 4.451 4.448 0.999 4.457 4.455 1.000
4 P3 Z changed B3 13.428 13.295 0.990 13.431 13.286 0.989
4 P3 D unchanged B1 4.773 4.771 0.999 4.787 4.793 1.001
4 P3 D changed B3 14.415 14.363 0.996 14.446 14.392 0.996
4 P4 D unchanged B1 4.500 4.496 0.999 4.541 4.543 1.000
4 P4 D unchanged B3 13.480 13.479 1.000 13.511 13.511 1.000
4 P5 C unchanged B1 4.296 4.294 1.000 4.310 4.314 1.001
4 P5 C changed B3 13.008 12.890 0.991 13.017 12.908 0.992
4 P5 Z unchanged B1 4.330 4.335 1.001 4.352 4.349 0.999
4 P5 Z changed B3 13.086 12.944 0.989 13.091 12.945 0.989
4 P5 D changed B1 4.428 4.438 1.002 4.467 4.470 1.001
4 P5 D changed B3 13.454 13.315 0.990 13.483 13.354 0.990
8 P1 C unchanged B1 2.538 2.540 1.001 2.599 2.579 0.992
8 P1 C unchanged B3 7.925 7.933 1.001 7.975 7.979 1.001
8 P1 D changed B1 2.671 2.677 1.002 2.685 2.666 0.993
8 P1 D changed B3 7.769 7.776 1.001 7.800 7.789 0.999
8 P2 Z unchanged B1 2.720 2.722 1.001 2.788 2.789 1.000
8 P2 Z unchanged B3 7.933 7.936 1.000 7.923 7.918 0.999
8 P3 C changed B1 2.635 2.576 0.978 2.636 2.585 0.981
8 P3 C changed B3 7.907 8.031 1.016 7.959 8.008 1.006
8 P4 D unchanged B1 2.753 2.753 1.000 2.780 2.770 0.997
8 P4 D unchanged B3 8.311 8.283 0.997 8.200 8.227 1.003
8 P5 D changed B1 2.604 2.604 1.000 2.605 2.608 1.001
8 P5 D changed B3 7.652 7.600 0.993 7.642 7.588 0.993

@yyihuang

yyihuang commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

Performance (3/5), continued

As in part 1, "this PR", "new", "previous" and "main" in the copied FlashInfer text refer to that FlashInfer pull request and to the FlashInfer package before it, not to this PR.

Fast (non-deterministic), single-rank, sm_100a and sm_103a (exact revision)

Source: flashinfer-ai/flashinfer#8675. Its "deterministic Cake" arm has the vendored det sm_100a and sm_103a programs.

Cake ships four GDN training variants in FlashInfer: a non-deterministic and a deterministic form, each as a single-rank package and a context-parallel (KCP) package. This PR changes only the non-deterministic single-rank package flashinfer.experimental.cake.cake_gdn_train (packed variable-length sequences, head dim 128, chunk 64, grouped Q/K heads). On sm_100a (B200) and sm_103a (GB300), group-of-three rows (HV = 3H, shown as H/HV 2/6 in the tables) now use a different backward. The adjoint scan is split into two groups, and the second group runs on a high-priority side stream beside the group-of-three tail. Before this PR, one serial scan ran beside that tail. The deterministic package already uses this backward form on these rows. All other rows keep the previous schedule, and so does every row on sm_107a (R200). On every measured row, the output and all gradients are bitwise identical to the previous non-deterministic package and to the deterministic package.

Ratios are baseline ms / this PR ms, so a value above 1 means this PR is faster. ckpt is the checkpointed step. Before this PR, four cells were slower than deterministic Cake beyond run-to-run noise: B200 FP32-beta 6 seqs, 241817 tokens fwd+bwd; B200 FP32-beta 64 x 4096 fwd+bwd and ckpt; and GB300 FP32-beta 64 x 4096 ckpt. All four are now at parity with deterministic Cake (0.998x-1.004x). Compared with the previous non-deterministic package, they are 1.007x-1.019x. Two of them (B200 64 x 4096 fwd+bwd and GB300 64 x 4096 ckpt) are classed faster, and the other two are at parity.

The bwd rows show where this PR costs time. On the four group-of-three rows, the isolated backward is 0.977x-0.991x of the previous non-deterministic package on both GPUs (0.9-2.4 % more time), which is the same time as the deterministic package (0.999x-1.002x). On those rows, fwd+bwd and ckpt are at parity with the previous package or faster. Two BF16-beta fwd+bwd cells are still slower than deterministic Cake beyond run-to-run noise, both at 7 x ~4.3k (real pack), H/HV 8/16: 0.985x on B200 and 0.992x on GB300. This PR does not change the schedule of those rows, and they are at parity with the previous package.

The tables use three baselines:

  • Deterministic Cake (cake_gdn_det_train at FlashInfer main).
  • The previous non-deterministic Cake package (FlashInfer main before this PR).
  • Per cell, the strongest applicable non-Cake implementation among cuDNN frontend 1.32.1 FROST (chunk 0 and chunk 64) and FLA 0.6.0 (Triton). Where the fastest non-Cake implementation drifted by more than 2 %, the cell is marked "unstable (not published)".

All timings are eager, on NVIDIA B200 (sm_100a, 148 SMs) and NVIDIA GB300 (sm_103a, 152 SMs), with CUDA 13.0 and torch 2.14.1+cu130.

NVIDIA B200 (sm_100a, 148 SMs), CUDA 13.0 (torch 2.14.1+cu130): BF16 q / k / v, FP32 beta, head dim 128, chunk 64, eager; ms per call = mean of two runs in opposite arm orders, each the median of 6 group medians of 16 samples (L2 flushed before each sample); speedup = baseline ms / this PR ms; bwd = the backward pass alone (its forward runs outside the timed region), fwd+bwd = one training step, ckpt = forward, recomputed forward and backward of a checkpointed step.

GPU / architecture workload (H / HV) operation deterministic Cake ms non-Cake SOTA ms previous non-deterministic Cake ms this PR ms speedup vs deterministic Cake speedup vs non-Cake SOTA speedup vs previous
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 fwd 0.2584 unstable (not published) 0.2578 0.2578 1.002x - 1.000x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 bwd 0.6697 1.1385 (cuDNN FROST (chunk 64)) 0.6146 0.6127 1.093x 1.858x 1.003x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 fwd+bwd 0.8622 1.5217 (cuDNN FROST (chunk 64)) 0.8063 0.8105 1.064x 1.877x 0.995x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 ckpt 1.4890 unstable (not published) 1.4227 1.4227 1.047x - 1.000x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 fwd 0.2270 0.2949 (cuDNN FROST) 0.2265 0.2264 1.003x 1.303x 1.000x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 bwd 0.5838 unstable (not published) 0.5575 0.5601 1.042x - 0.995x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 fwd+bwd 0.7781 unstable (not published) 0.7550 0.7502 1.037x - 1.006x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 ckpt 1.3974 2.1785 (cuDNN FROST (chunk 64)) 1.3571 1.3557 1.031x 1.607x 1.001x
NVIDIA B200 sm_100a 8 x 4096, 8/16 fwd 0.4285 0.2788 (cuDNN FROST) 0.4279 0.4285 1.000x 0.651x 0.999x
NVIDIA B200 sm_100a 8 x 4096, 8/16 bwd 0.9232 1.3856 (cuDNN FROST (chunk 64)) 0.9099 0.9074 1.017x 1.527x 1.003x
NVIDIA B200 sm_100a 8 x 4096, 8/16 fwd+bwd 1.1425 1.7861 (cuDNN FROST (chunk 64)) 1.1405 1.1407 1.002x 1.566x 1.000x
NVIDIA B200 sm_100a 8 x 4096, 8/16 ckpt 1.8371 2.5145 (cuDNN FROST (chunk 64)) 1.8266 1.8230 1.008x 1.379x 1.002x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 fwd 0.4550 0.3434 (cuDNN FROST) 0.4540 0.4545 1.001x 0.756x 0.999x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 bwd 0.9502 unstable (not published) 0.9137 0.9192 1.034x - 0.994x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 fwd+bwd 1.1968 unstable (not published) 1.1758 1.1775 1.016x - 0.999x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 ckpt 1.8926 2.6182 (cuDNN FROST (chunk 64)) 1.8733 1.8714 1.011x 1.399x 1.001x
NVIDIA B200 sm_100a 16 x 4096, 8/16 fwd 0.8041 0.4420 (cuDNN FROST) 0.8099 0.8045 1.000x 0.549x 1.007x
NVIDIA B200 sm_100a 16 x 4096, 8/16 bwd 1.5827 2.2769 (cuDNN FROST (chunk 64)) 1.5809 1.5836 0.999x 1.438x 0.998x
NVIDIA B200 sm_100a 16 x 4096, 8/16 fwd+bwd 2.1802 unstable (not published) 2.1878 2.1716 1.004x - 1.007x
NVIDIA B200 sm_100a 16 x 4096, 8/16 ckpt 3.0918 unstable (not published) 3.1370 3.0894 1.001x - 1.015x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 fwd 0.8036 0.4450 (cuDNN FROST) 0.8092 0.8039 1.000x 0.554x 1.007x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 bwd 1.5474 2.3501 (cuDNN FROST (chunk 64)) 1.5248 1.5399 1.005x 1.526x 0.990x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 fwd+bwd 2.2003 unstable (not published) 2.2115 2.1838 1.008x - 1.013x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 ckpt 3.0951 3.5896 (cuDNN FROST (chunk 64)) 3.1514 3.1033 0.997x 1.157x 1.015x
NVIDIA B200 sm_100a 1 x 16384, 8/16 fwd 0.2615 0.3064 (cuDNN FROST) 0.2616 0.2614 1.000x 1.172x 1.001x
NVIDIA B200 sm_100a 1 x 16384, 8/16 bwd 0.6932 1.1321 (cuDNN FROST (chunk 64)) 0.6065 0.6076 1.141x 1.863x 0.998x
NVIDIA B200 sm_100a 1 x 16384, 8/16 fwd+bwd 0.8872 unstable (not published) 0.7963 0.8032 1.105x - 0.991x
NVIDIA B200 sm_100a 1 x 16384, 8/16 ckpt 1.5004 2.3065 (cuDNN FROST (chunk 64)) 1.4157 1.4159 1.060x 1.629x 1.000x
NVIDIA B200 sm_100a 4 x 4096, 16/32 fwd 0.4322 0.2802 (cuDNN FROST) 0.4321 0.4315 1.002x 0.649x 1.001x
NVIDIA B200 sm_100a 4 x 4096, 16/32 bwd 0.9171 1.3992 (cuDNN FROST (chunk 64)) 0.9076 0.9083 1.010x 1.540x 0.999x
NVIDIA B200 sm_100a 4 x 4096, 16/32 fwd+bwd 1.1462 1.7993 (cuDNN FROST (chunk 64)) 1.1459 1.1474 0.999x 1.568x 0.999x
NVIDIA B200 sm_100a 4 x 4096, 16/32 ckpt 1.8436 2.5279 (cuDNN FROST (chunk 64)) 1.8391 1.8369 1.004x 1.376x 1.001x
NVIDIA B200 sm_100a 16 x 4096, 16/32 fwd 1.6243 unstable (not published) 1.6195 1.6038 1.013x - 1.010x
NVIDIA B200 sm_100a 16 x 4096, 16/32 bwd 2.9404 3.9982 (cuDNN FROST (chunk 64)) 2.9541 2.9322 1.003x 1.364x 1.007x
NVIDIA B200 sm_100a 16 x 4096, 16/32 fwd+bwd 4.3998 4.6347 (cuDNN FROST (chunk 64)) 4.4081 4.4036 0.999x 1.052x 1.001x
NVIDIA B200 sm_100a 16 x 4096, 16/32 ckpt 6.1400 5.6580 (cuDNN FROST (chunk 64)) 6.1384 6.1245 1.003x 0.924x 1.002x
NVIDIA B200 sm_100a 1 x 16384, 16/32 fwd 0.4451 0.4289 (cuDNN FROST) 0.4449 0.4451 1.000x 0.964x 1.000x
NVIDIA B200 sm_100a 1 x 16384, 16/32 bwd 1.0850 1.6286 (cuDNN FROST (chunk 64)) 0.9290 0.9296 1.167x 1.752x 0.999x
NVIDIA B200 sm_100a 1 x 16384, 16/32 fwd+bwd 1.3145 2.0147 (cuDNN FROST (chunk 64)) 1.1786 1.1796 1.114x 1.708x 0.999x
NVIDIA B200 sm_100a 1 x 16384, 16/32 ckpt 2.0153 2.7486 (cuDNN FROST (chunk 64)) 1.8771 1.8751 1.075x 1.466x 1.001x
NVIDIA B200 sm_100a 131071 + 87531 + 43385, 2/6 fwd 1.1861 unstable (not published) 1.1717 1.1813 1.004x - 0.992x
NVIDIA B200 sm_100a 131071 + 87531 + 43385, 2/6 bwd 2.4237 3.5866 (cuDNN FROST (chunk 64)) 2.3708 2.4236 1.000x 1.480x 0.978x
NVIDIA B200 sm_100a 131071 + 87531 + 43385, 2/6 fwd+bwd 3.3792 4.3041 (cuDNN FROST (chunk 64)) 3.3901 3.3792 1.000x 1.274x 1.003x
NVIDIA B200 sm_100a 131071 + 87531 + 43385, 2/6 ckpt 4.6529 5.3943 (cuDNN FROST (chunk 64)) 4.6928 4.6890 0.992x 1.150x 1.001x
NVIDIA B200 sm_100a 2 seqs, 259174 tokens, 2/6 fwd 1.1717 0.9097 (cuDNN FROST) 1.1566 1.1675 1.004x 0.779x 0.991x
NVIDIA B200 sm_100a 2 seqs, 259174 tokens, 2/6 bwd 2.4011 unstable (not published) 2.3513 2.4042 0.999x - 0.978x
NVIDIA B200 sm_100a 2 seqs, 259174 tokens, 2/6 fwd+bwd 3.3453 4.2796 (cuDNN FROST (chunk 64)) 3.3914 3.3480 0.999x 1.278x 1.013x
NVIDIA B200 sm_100a 2 seqs, 259174 tokens, 2/6 ckpt 4.6221 5.3806 (cuDNN FROST (chunk 64)) 4.6670 4.6443 0.995x 1.159x 1.005x
NVIDIA B200 sm_100a 6 seqs, 241817 tokens, 2/6 fwd 1.0888 0.9612 (cuDNN FROST) 1.0805 1.0840 1.004x 0.887x 0.997x
NVIDIA B200 sm_100a 6 seqs, 241817 tokens, 2/6 bwd 2.1845 3.4335 (cuDNN FROST (chunk 64)) 2.1680 2.1867 0.999x 1.570x 0.991x
NVIDIA B200 sm_100a 6 seqs, 241817 tokens, 2/6 fwd+bwd 3.1432 4.3035 (cuDNN FROST (chunk 64)) 3.1602 3.1320 1.004x 1.374x 1.009x
NVIDIA B200 sm_100a 6 seqs, 241817 tokens, 2/6 ckpt 4.3387 5.4668 (cuDNN FROST (chunk 64)) 4.3893 4.3356 1.001x 1.261x 1.012x
NVIDIA B200 sm_100a 64 x 4096, 2/6 fwd 1.1462 0.5696 (cuDNN FROST) 1.1435 1.1449 1.001x 0.498x 0.999x
NVIDIA B200 sm_100a 64 x 4096, 2/6 bwd 2.3950 3.0062 (cuDNN FROST (chunk 64)) 2.3630 2.3949 1.000x 1.255x 0.987x
NVIDIA B200 sm_100a 64 x 4096, 2/6 fwd+bwd 3.3236 3.4232 (cuDNN FROST (chunk 64)) 3.3939 3.3316 0.998x 1.027x 1.019x
NVIDIA B200 sm_100a 64 x 4096, 2/6 ckpt 4.6090 4.3206 (cuDNN FROST (chunk 64)) 4.6335 4.6010 1.002x 0.939x 1.007x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 16/48 fwd 1.1780 0.5811 (cuDNN FROST) 1.1747 1.1788 0.999x 0.493x 0.997x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 16/48 bwd 2.3825 3.0242 (cuDNN FROST (chunk 64)) 2.3767 2.3750 1.003x 1.273x 1.001x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 16/48 fwd+bwd 3.3842 3.4538 (cuDNN FROST (chunk 64)) 3.4274 3.3901 0.998x 1.019x 1.011x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 16/48 ckpt 4.7358 4.3091 (cuDNN FROST (chunk 64)) 4.7269 4.7509 0.997x 0.907x 0.995x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 8/24 fwd 0.6149 0.4444 (cuDNN FROST) 0.6142 0.6140 1.001x 0.724x 1.000x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 8/24 bwd 1.5359 unstable (not published) 1.3421 1.3437 1.143x - 0.999x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 8/24 fwd+bwd 1.9482 2.5190 (cuDNN FROST (chunk 64)) 1.7591 1.7592 1.107x 1.432x 1.000x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 8/24 ckpt 2.6600 3.2439 (cuDNN FROST (chunk 64)) 2.4705 2.4709 1.077x 1.313x 1.000x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 4/12 fwd 0.3426 0.3760 (cuDNN FROST) 0.3421 0.3418 1.002x 1.100x 1.001x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 4/12 bwd 0.9398 unstable (not published) 0.8208 0.8224 1.143x - 0.998x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 4/12 fwd+bwd 1.1303 unstable (not published) 0.9997 1.0065 1.123x - 0.993x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 4/12 ckpt 1.7668 2.5279 (cuDNN FROST (chunk 64)) 1.6673 1.6684 1.059x 1.515x 0.999x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 2/6 fwd 0.2481 unstable (not published) 0.2473 0.2475 1.002x - 0.999x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 2/6 bwd 0.5483 0.8953 (cuDNN FROST (chunk 64)) 0.4686 0.4701 1.166x 1.904x 0.997x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 2/6 fwd+bwd 0.8407 unstable (not published) 0.6942 0.7235 1.162x - 0.960x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 2/6 ckpt 1.2972 unstable (not published) 1.2124 1.2144 1.068x - 0.998x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 1/3 fwd 0.1657 unstable (not published) 0.1648 0.1653 1.002x - 0.997x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 1/3 bwd 0.5629 unstable (not published) 0.5428 0.5460 1.031x - 0.994x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 1/3 fwd+bwd 0.5928 unstable (not published) 0.5644 0.5651 1.049x - 0.999x
NVIDIA B200 sm_100a 11000 + 12000 + 9768, 1/3 ckpt 1.1663 1.8673 (cuDNN FROST (chunk 64)) 1.1397 1.1394 1.024x 1.639x 1.000x

NVIDIA B200 (sm_100a, 148 SMs), CUDA 13.0 (torch 2.14.1+cu130): BF16 q / k / v, BF16 beta, head dim 128, chunk 64, eager; ms per call = mean of two runs in opposite arm orders, each the median of 6 group medians of 16 samples (L2 flushed before each sample); speedup = baseline ms / this PR ms; bwd = the backward pass alone (its forward runs outside the timed region), fwd+bwd = one training step, ckpt = forward, recomputed forward and backward of a checkpointed step.

GPU / architecture workload (H / HV) operation deterministic Cake ms non-Cake SOTA ms previous non-deterministic Cake ms this PR ms speedup vs deterministic Cake speedup vs non-Cake SOTA speedup vs previous
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 fwd 0.2594 0.2254 (cuDNN FROST) 0.2592 0.2589 1.002x 0.871x 1.001x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 bwd 0.6623 1.1187 (cuDNN FROST (chunk 64)) 0.6091 0.6104 1.085x 1.833x 0.998x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 fwd+bwd 0.8581 1.5004 (cuDNN FROST (chunk 64)) 0.8044 0.8047 1.066x 1.865x 1.000x
NVIDIA B200 sm_100a 4 x 4096, H/HV 8/16 ckpt 1.4597 2.2197 (cuDNN FROST (chunk 64)) 1.3916 1.3926 1.048x 1.594x 0.999x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 fwd 0.2286 0.2869 (cuDNN FROST) 0.2280 0.2283 1.001x 1.257x 0.999x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 bwd 0.5844 1.0318 (cuDNN FROST (chunk 64)) 0.5606 0.5603 1.043x 1.842x 1.001x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 fwd+bwd 0.7837 unstable (not published) 0.7552 0.7553 1.038x - 1.000x
NVIDIA B200 sm_100a 3 x ~4.25k (real pack), 8/16 ckpt 1.3878 2.1423 (cuDNN FROST (chunk 64)) 1.3566 1.3526 1.026x 1.584x 1.003x
NVIDIA B200 sm_100a 8 x 4096, 8/16 fwd 0.4302 0.2726 (cuDNN FROST) 0.4297 0.4301 1.000x 0.634x 0.999x
NVIDIA B200 sm_100a 8 x 4096, 8/16 bwd 0.9117 1.3515 (cuDNN FROST (chunk 64)) 0.8822 0.8865 1.028x 1.525x 0.995x
NVIDIA B200 sm_100a 8 x 4096, 8/16 fwd+bwd 1.1413 1.7339 (cuDNN FROST (chunk 64)) 1.1416 1.1419 0.999x 1.518x 1.000x
NVIDIA B200 sm_100a 8 x 4096, 8/16 ckpt 1.8014 unstable (not published) 1.7985 1.7975 1.002x - 1.001x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 fwd 0.4592 unstable (not published) 0.4590 0.4589 1.001x - 1.000x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 bwd 0.9703 1.4822 (cuDNN FROST (chunk 64)) 0.9108 0.9119 1.064x 1.625x 0.999x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 fwd+bwd 1.1615 1.8564 (cuDNN FROST (chunk 64)) 1.1797 1.1797 0.985x 1.574x 1.000x
NVIDIA B200 sm_100a 7 x ~4.3k (real pack), 8/16 ckpt 1.8431 2.5824 (cuDNN FROST (chunk 64)) 1.8632 1.8631 0.989x 1.386x 1.000x
NVIDIA B200 sm_100a 16 x 4096, 8/16 fwd 0.8049 0.4318 (cuDNN FROST) 0.8035 0.8047 1.000x 0.537x 0.999x
NVIDIA B200 sm_100a 16 x 4096, 8/16 bwd 1.5879 2.2145 (cuDNN FROST (chunk 64)) 1.5792 1.5855 1.002x 1.397x 0.996x
NVIDIA B200 sm_100a 16 x 4096, 8/16 fwd+bwd 2.1834 unstable (not published) 2.1953 2.1700 1.006x - 1.012x
NVIDIA B200 sm_100a 16 x 4096, 8/16 ckpt 3.0580 3.3284 (cuDNN FROST (chunk 64)) 3.1279 3.0604 0.999x 1.088x 1.022x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 fwd 0.8018 0.4372 (cuDNN FROST) 0.8013 0.8056 0.995x 0.543x 0.995x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 bwd 1.6130 2.3955 (cuDNN FROST (chunk 64)) 1.6121 1.6166 0.998x 1.482x 0.997x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 fwd+bwd 2.1974 2.7736 (cuDNN FROST (chunk 64)) 2.2042 2.1905 1.003x 1.266x 1.006x
NVIDIA B200 sm_100a 15 x ~4.3k (real pack), 8/16 ckpt 3.0900 3.5150 (cuDNN FROST (chunk 64)) 3.1295 3.0811 1.003x 1.141x 1.016x
NVIDIA B200 sm_100a 1 x 16384, 8/16 fwd 0.2636 0.2989 (cuDNN FROST) 0.2633 0.2637 1.000x 1.133x 0.998x
NVIDIA B200 sm_100a 1 x 16384, 8/16 bwd 0.6894 unstable (not published) 0.6063 0.6062 1.137x - 1.000x
NVIDIA B200 sm_100a 1 x 16384, 8/16 fwd+bwd 0.8803 unstable (not published) 0.7870 0.7923 1.111x - 0.993x
NVIDIA B200 sm_100a 1 x 16384, 8/16 ckpt 1.4734 2.2062 (cuDNN FROST (chunk 64)) 1.3910 1.3927 1.058x 1.584x 0.999x
NVIDIA B200 sm_100a 4 x 4096, 16/32 fwd 0.4351 0.2756 (cuDNN FROST) 0.4339 0.4349 1.000x 0.634x 0.998x
NVIDIA B200 sm_100a 4 x 4096, 16/32 bwd 0.9133 1.3389 (cuDNN FROST (chunk 64)) 0.9051 0.8998 1.015x 1.488x 1.006x
NVIDIA B200 sm_100a 4 x 4096, 16/32 fwd+bwd 1.1479 1.7211 (cuDNN FROST (chunk 64)) 1.1479 1.1470 1.001x 1.501x 1.001x
NVIDIA B200 sm_100a 4 x 4096, 16/32 ckpt 1.8006 2.4312 (cuDNN FROST (chunk 64)) 1.7951 1.7931 1.004x 1.356x 1.001x
NVIDIA B200 sm_100a 16 x 4096, 16/32 fwd 1.6195 unstable (not published) 1.6233 1.6167 1.002x - 1.004x
NVIDIA B200 sm_100a 16 x 4096, 16/32 bwd 2.8783 3.8848 (cuDNN FROST (chunk 64)) 2.8739 2.8666 1.004x 1.355x 1.003x
NVIDIA B200 sm_100a 16 x 4096, 16/32 fwd+bwd 4.4182 4.6629 (cuDNN FROST (chunk 64)) 4.4034 4.3883 1.007x 1.063x 1.003x
NVIDIA B200 sm_100a 16 x 4096, 16/32 ckpt 6.1810 5.6679 (cuDNN FROST (chunk 64)) 6.1574 6.1462 1.006x 0.922x 1.002x
NVIDIA B200 sm_100a 1 x 16384, 16/32 fwd 0.4530 0.4306 (cuDNN FROST) 0.4520 0.4522 1.002x 0.952x 1.000x
NVIDIA B200 sm_100a 1 x 16384, 16/32 bwd 1.0372 1.6286 (cuDNN FROST (chunk 64)) 0.9231 0.9233 1.123x 1.764x 1.000x
NVIDIA B200 sm_100a 1 x 16384, 16/32 fwd+bwd 1.2243 2.0081 (cuDNN FROST (chunk 64)) 1.1849 1.1845 1.034x 1.695x 1.000x
NVIDIA B200 sm_100a 1 x 16384, 16/32 ckpt 1.8712 2.7234 (cuDNN FROST (chunk 64)) 1.8675 1.8674 1.002x 1.458x 1.000x

NVIDIA GB300 (sm_103a, 152 SMs), CUDA 13.0 (torch 2.14.1+cu130): BF16 q / k / v, FP32 beta, head dim 128, chunk 64, eager; ms per call = mean of two runs in opposite arm orders, each the median of 6 group medians of 16 samples (L2 flushed before each sample); speedup = baseline ms / this PR ms; bwd = the backward pass alone (its forward runs outside the timed region), fwd+bwd = one training step, ckpt = forward, recomputed forward and backward of a checkpointed step.

GPU / architecture workload (H / HV) operation deterministic Cake ms non-Cake SOTA ms previous non-deterministic Cake ms this PR ms speedup vs deterministic Cake speedup vs non-Cake SOTA speedup vs previous
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 fwd 0.2788 unstable (not published) 0.2779 0.2791 0.999x - 0.996x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 bwd 0.4797 0.8907 (cuDNN FROST (chunk 64)) 0.4304 0.4303 1.115x 2.070x 1.000x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 fwd+bwd 0.6641 unstable (not published) 0.6142 0.6137 1.082x - 1.001x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 ckpt 1.3346 unstable (not published) 1.2784 1.2705 1.050x - 1.006x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 fwd 0.2681 unstable (not published) 0.2665 0.2679 1.001x - 0.995x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 bwd 0.4149 unstable (not published) 0.3932 0.3920 1.058x - 1.003x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 fwd+bwd 0.6179 unstable (not published) 0.5931 0.6008 1.028x - 0.987x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 ckpt 1.2855 unstable (not published) 1.2262 1.2563 1.023x - 0.976x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 fwd 0.4669 unstable (not published) 0.4686 0.4697 0.994x - 0.998x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 bwd 0.7139 unstable (not published) 0.7041 0.7075 1.009x - 0.995x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 fwd+bwd 1.0480 1.5173 (cuDNN FROST (chunk 64)) 1.0469 1.0492 0.999x 1.446x 0.998x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 ckpt 1.5684 unstable (not published) 1.5611 1.5633 1.003x - 0.999x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 fwd 0.4565 unstable (not published) 0.4573 0.4571 0.999x - 1.000x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 bwd 0.7215 1.1727 (cuDNN FROST (chunk 64)) 0.7120 0.7126 1.012x 1.646x 0.999x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 fwd+bwd 1.0551 unstable (not published) 1.0545 1.0542 1.001x - 1.000x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 ckpt 1.5867 unstable (not published) 1.5751 1.5768 1.006x - 0.999x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 fwd 0.8321 unstable (not published) 0.8252 0.8327 0.999x - 0.991x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 bwd 1.3382 2.0155 (cuDNN FROST (chunk 64)) 1.3359 1.3363 1.001x 1.508x 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 fwd+bwd 1.9716 2.4706 (cuDNN FROST (chunk 64)) 1.9706 1.9705 1.001x 1.254x 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 ckpt 2.7897 3.4536 (cuDNN FROST (chunk 64)) 2.7886 2.7916 0.999x 1.237x 0.999x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 fwd 0.8401 unstable (not published) 0.8391 0.8386 1.002x - 1.001x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 bwd 1.3638 2.1502 (cuDNN FROST (chunk 64)) 1.3613 1.3630 1.001x 1.578x 0.999x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 fwd+bwd 2.0015 unstable (not published) 1.9999 2.0005 1.000x - 1.000x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 ckpt 2.8342 unstable (not published) 2.8322 2.8340 1.000x - 0.999x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 fwd 0.3055 unstable (not published) 0.3000 0.3053 1.001x - 0.983x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 bwd 0.5213 unstable (not published) 0.4456 0.4460 1.169x - 0.999x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 fwd+bwd 0.7240 unstable (not published) 0.6431 0.6512 1.112x - 0.988x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 ckpt 1.3908 unstable (not published) 1.3097 1.3038 1.067x - 1.005x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 fwd 0.4496 unstable (not published) 0.4474 0.4502 0.999x - 0.994x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 bwd 0.7044 1.1213 (cuDNN FROST (chunk 64)) 0.6987 0.6986 1.008x 1.605x 1.000x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 fwd+bwd 1.0586 1.5862 (cuDNN FROST (chunk 64)) 1.0586 1.0582 1.000x 1.499x 1.000x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 ckpt 1.5758 unstable (not published) 1.5781 1.5736 1.001x - 1.003x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 fwd 1.4959 0.8347 (cuDNN FROST) 1.4957 1.4979 0.999x 0.557x 0.999x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 bwd 2.4829 3.7228 (cuDNN FROST (chunk 64)) 2.4715 2.4772 1.002x 1.503x 0.998x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 fwd+bwd 3.7800 4.4249 (cuDNN FROST (chunk 64)) 3.7784 3.7794 1.000x 1.171x 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 ckpt 5.2743 5.4027 (cuDNN FROST (chunk 64)) 5.2704 5.2740 1.000x 1.024x 0.999x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 fwd 0.4716 unstable (not published) 0.4718 0.4729 0.997x - 0.998x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 bwd 0.7746 unstable (not published) 0.7284 0.7345 1.055x - 0.992x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 fwd+bwd 1.1324 unstable (not published) 1.0931 1.0958 1.033x - 0.998x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 ckpt 1.7112 unstable (not published) 1.6661 1.6554 1.034x - 1.006x
NVIDIA GB300 sm_103a 131071 + 87531 + 43385, 2/6 fwd 1.1150 0.9006 (cuDNN FROST) 1.1149 1.1153 1.000x 0.807x 1.000x
NVIDIA GB300 sm_103a 131071 + 87531 + 43385, 2/6 bwd 2.0471 3.1925 (cuDNN FROST (chunk 64)) 2.0151 2.0439 1.002x 1.562x 0.986x
NVIDIA GB300 sm_103a 131071 + 87531 + 43385, 2/6 fwd+bwd 3.0266 3.9860 (cuDNN FROST (chunk 64)) 3.0386 3.0277 1.000x 1.317x 1.004x
NVIDIA GB300 sm_103a 131071 + 87531 + 43385, 2/6 ckpt 4.1690 4.9954 (cuDNN FROST (chunk 64)) 4.1820 4.1710 1.000x 1.198x 1.003x
NVIDIA GB300 sm_103a 2 seqs, 259174 tokens, 2/6 fwd 1.0966 0.9024 (cuDNN FROST) 1.0948 1.0973 0.999x 0.822x 0.998x
NVIDIA GB300 sm_103a 2 seqs, 259174 tokens, 2/6 bwd 2.0343 3.1950 (cuDNN FROST (chunk 64)) 1.9898 2.0315 1.001x 1.573x 0.979x
NVIDIA GB300 sm_103a 2 seqs, 259174 tokens, 2/6 fwd+bwd 2.9939 3.9701 (cuDNN FROST (chunk 64)) 2.9927 2.9939 1.000x 1.326x 1.000x
NVIDIA GB300 sm_103a 2 seqs, 259174 tokens, 2/6 ckpt 4.1081 4.9812 (cuDNN FROST (chunk 64)) 4.1132 4.1146 0.998x 1.211x 1.000x
NVIDIA GB300 sm_103a 6 seqs, 241817 tokens, 2/6 fwd 1.0585 unstable (not published) 1.0616 1.0601 0.998x - 1.001x
NVIDIA GB300 sm_103a 6 seqs, 241817 tokens, 2/6 bwd 1.9474 3.1942 (cuDNN FROST (chunk 64)) 1.8990 1.9437 1.002x 1.643x 0.977x
NVIDIA GB300 sm_103a 6 seqs, 241817 tokens, 2/6 fwd+bwd 2.8283 4.0304 (cuDNN FROST (chunk 64)) 2.8393 2.8297 1.000x 1.424x 1.003x
NVIDIA GB300 sm_103a 6 seqs, 241817 tokens, 2/6 ckpt 3.9027 5.1094 (cuDNN FROST (chunk 64)) 3.9167 3.9023 1.000x 1.309x 1.004x
NVIDIA GB300 sm_103a 64 x 4096, 2/6 fwd 1.1353 unstable (not published) 1.1359 1.1376 0.998x - 0.999x
NVIDIA GB300 sm_103a 64 x 4096, 2/6 bwd 2.0537 2.6824 (cuDNN FROST (chunk 64)) 2.0309 2.0524 1.001x 1.307x 0.990x
NVIDIA GB300 sm_103a 64 x 4096, 2/6 fwd+bwd 2.9890 3.1549 (cuDNN FROST (chunk 64)) 3.0176 2.9836 1.002x 1.057x 1.011x
NVIDIA GB300 sm_103a 64 x 4096, 2/6 ckpt 4.0963 4.0172 (cuDNN FROST (chunk 64)) 4.1246 4.0923 1.001x 0.982x 1.008x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 16/48 fwd 1.1287 unstable (not published) 1.1273 1.1286 1.000x - 0.999x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 16/48 bwd 2.0400 2.7102 (cuDNN FROST (chunk 64)) 2.0307 2.0316 1.004x 1.334x 1.000x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 16/48 fwd+bwd 3.0300 unstable (not published) 3.0318 3.0316 0.999x - 1.000x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 16/48 ckpt 4.1712 4.1738 (cuDNN FROST (chunk 64)) 4.1757 4.1705 1.000x 1.001x 1.001x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 8/24 fwd 0.6330 unstable (not published) 0.6335 0.6331 1.000x - 1.001x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 8/24 bwd 1.2683 1.8527 (cuDNN FROST (chunk 64)) 1.0909 1.0908 1.163x 1.698x 1.000x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 8/24 fwd+bwd 1.7784 2.2320 (cuDNN FROST (chunk 64)) 1.6054 1.6036 1.109x 1.392x 1.001x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 8/24 ckpt 2.4225 3.1318 (cuDNN FROST (chunk 64)) 2.2528 2.2487 1.077x 1.393x 1.002x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 4/12 fwd 0.3560 unstable (not published) 0.3566 0.3559 1.000x - 1.002x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 4/12 bwd 0.7222 1.1442 (cuDNN FROST (chunk 64)) 0.6150 0.6161 1.172x 1.857x 0.998x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 4/12 fwd+bwd 0.9718 unstable (not published) 0.8826 0.8830 1.101x - 1.000x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 4/12 ckpt 1.5375 unstable (not published) 1.4136 1.4162 1.086x - 0.998x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 2/6 fwd 0.3140 unstable (not published) 0.3156 0.3142 0.999x - 1.004x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 2/6 bwd 0.4477 unstable (not published) 0.3945 0.3946 1.135x - 1.000x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 2/6 fwd+bwd 0.6623 unstable (not published) 0.6001 0.6049 1.095x - 0.992x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 2/6 ckpt 1.2927 unstable (not published) 1.2290 1.2269 1.054x - 1.002x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 1/3 fwd 0.2408 unstable (not published) 0.2355 0.2364 1.019x - 0.996x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 1/3 bwd 0.3067 unstable (not published) 0.2982 0.2996 1.024x - 0.995x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 1/3 fwd+bwd 0.5173 unstable (not published) 0.4982 0.5152 1.004x - 0.967x
NVIDIA GB300 sm_103a 11000 + 12000 + 9768, 1/3 ckpt 1.2012 unstable (not published) 1.1597 1.1800 1.018x - 0.983x

NVIDIA GB300 (sm_103a, 152 SMs), CUDA 13.0 (torch 2.14.1+cu130): BF16 q / k / v, BF16 beta, head dim 128, chunk 64, eager; ms per call = mean of two runs in opposite arm orders, each the median of 6 group medians of 16 samples (L2 flushed before each sample); speedup = baseline ms / this PR ms; bwd = the backward pass alone (its forward runs outside the timed region), fwd+bwd = one training step, ckpt = forward, recomputed forward and backward of a checkpointed step.

GPU / architecture workload (H / HV) operation deterministic Cake ms non-Cake SOTA ms previous non-deterministic Cake ms this PR ms speedup vs deterministic Cake speedup vs non-Cake SOTA speedup vs previous
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 fwd 0.2774 unstable (not published) 0.2776 0.2773 1.000x - 1.001x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 bwd 0.4965 0.8759 (cuDNN FROST (chunk 64)) 0.4280 0.4276 1.161x 2.048x 1.001x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 fwd+bwd 0.6780 unstable (not published) 0.6146 0.6151 1.102x - 0.999x
NVIDIA GB300 sm_103a 4 x 4096, H/HV 8/16 ckpt 1.2524 unstable (not published) 1.1687 1.1663 1.074x - 1.002x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 fwd 0.2494 unstable (not published) 0.2477 0.2479 1.006x - 0.999x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 bwd 0.4014 unstable (not published) 0.3804 0.3821 1.051x - 0.996x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 fwd+bwd 0.5851 unstable (not published) 0.5582 0.5603 1.044x - 0.996x
NVIDIA GB300 sm_103a 3 x ~4.25k (real pack), 8/16 ckpt 1.1448 unstable (not published) 1.1170 1.1095 1.032x - 1.007x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 fwd 0.4605 unstable (not published) 0.4606 0.4609 0.999x - 0.999x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 bwd 0.7085 1.1097 (cuDNN FROST (chunk 64)) 0.7002 0.7021 1.009x 1.581x 0.997x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 fwd+bwd 1.0509 unstable (not published) 1.0510 1.0520 0.999x - 0.999x
NVIDIA GB300 sm_103a 8 x 4096, 8/16 ckpt 1.5107 2.2732 (cuDNN FROST (chunk 64)) 1.5049 1.5078 1.002x 1.508x 0.998x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 fwd 0.4597 unstable (not published) 0.4595 0.4562 1.008x - 1.007x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 bwd 0.7726 1.1535 (cuDNN FROST (chunk 64)) 0.7126 0.7102 1.088x 1.624x 1.003x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 fwd+bwd 1.0483 unstable (not published) 1.0594 1.0569 0.992x - 1.002x
NVIDIA GB300 sm_103a 7 x ~4.3k (real pack), 8/16 ckpt 1.5595 2.4506 (cuDNN FROST (chunk 64)) 1.5583 1.5573 1.001x 1.574x 1.001x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 fwd 0.7727 unstable (not published) 0.7716 0.7730 1.000x - 0.998x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 bwd 1.3271 unstable (not published) 1.3240 1.3256 1.001x - 0.999x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 fwd+bwd 1.9621 unstable (not published) 1.9596 1.9605 1.001x - 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 8/16 ckpt 2.7414 3.1045 (cuDNN FROST (chunk 64)) 2.7430 2.7452 0.999x 1.131x 0.999x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 fwd 0.7853 unstable (not published) 0.7837 0.7817 1.005x - 1.003x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 bwd 1.3561 2.1072 (cuDNN FROST (chunk 64)) 1.3536 1.3566 1.000x 1.553x 0.998x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 fwd+bwd 1.9861 unstable (not published) 1.9872 1.9840 1.001x - 1.002x
NVIDIA GB300 sm_103a 15 x ~4.3k (real pack), 8/16 ckpt 2.7901 unstable (not published) 2.7902 2.7875 1.001x - 1.001x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 fwd 0.2839 unstable (not published) 0.2833 0.2835 1.001x - 0.999x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 bwd 0.5430 unstable (not published) 0.4378 0.4392 1.236x - 0.997x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 fwd+bwd 0.7488 unstable (not published) 0.6414 0.6453 1.160x - 0.994x
NVIDIA GB300 sm_103a 1 x 16384, 8/16 ckpt 1.3202 unstable (not published) 1.2074 1.2034 1.097x - 1.003x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 fwd 0.4544 unstable (not published) 0.4551 0.4549 0.999x - 1.000x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 bwd 0.7156 unstable (not published) 0.7068 0.7082 1.010x - 0.998x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 fwd+bwd 1.0675 unstable (not published) 1.0696 1.0682 0.999x - 1.001x
NVIDIA GB300 sm_103a 4 x 4096, 16/32 ckpt 1.6127 unstable (not published) 1.6042 1.6048 1.005x - 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 fwd 1.4707 unstable (not published) 1.4684 1.4687 1.001x - 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 bwd 2.4521 3.6307 (cuDNN FROST (chunk 64)) 2.4491 2.4462 1.002x 1.484x 1.001x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 fwd+bwd 3.7847 4.3894 (cuDNN FROST (chunk 64)) 3.7837 3.7833 1.000x 1.160x 1.000x
NVIDIA GB300 sm_103a 16 x 4096, 16/32 ckpt 5.2707 5.3222 (cuDNN FROST (chunk 64)) 5.2734 5.2732 1.000x 1.009x 1.000x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 fwd 0.4557 unstable (not published) 0.4564 0.4540 1.004x - 1.005x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 bwd 0.8066 1.3459 (cuDNN FROST (chunk 64)) 0.7197 0.7174 1.124x 1.876x 1.003x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 fwd+bwd 1.1038 1.7146 (cuDNN FROST (chunk 64)) 1.0846 1.0827 1.019x 1.584x 1.002x
NVIDIA GB300 sm_103a 1 x 16384, 16/32 ckpt 1.6088 unstable (not published) 1.5724 1.5742 1.022x - 0.999x

@yyihuang

yyihuang commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

Performance (4/5), continued

As in part 1, "this PR", "new", "previous" and "main" in the copied FlashInfer text refer to that FlashInfer pull request and to the FlashInfer package before it, not to this PR.

Fast (non-deterministic), single-rank, sm_107a

Source: the R200 table of flashinfer-ai/flashinfer#8665. Its "this PR" column runs the vendored sm_107a programs and schedule. Its "deterministic Cake" column is the det package at FlashInfer main 4046a863 (tree 738071dc), before the sm_107a forward WY program of flashinfer-ai/flashinfer#8695; on R200 that column and the ratio against it do not describe the vendored det package. The non-Cake column and its ratio do apply.

Every ratio is baseline ms / this PR ms, so a value above 1 means this PR is faster. ckpt is the checkpointed step. The tables use three baselines:

  • "Deterministic Cake" is the deterministic package at FlashInfer main 4046a863.
  • "Non-Cake SOTA" is chosen per cell. It is the fastest accuracy-qualified implementation among FLA 0.6.0 (Triton) and the GDN from cuDNN frontend 1.32.1 (FROST) in two modes: states recomputed ("cuDNN FROST" in the tables) or states stored every 64 tokens ("cuDNN FROST (chunk 64)"). FlashQLA 0.1.3, timed earlier on one shape on B200 and GB300, was slower than both cuDNN FROST modes there, so it is not among the candidates. In the published cells, it is cuDNN FROST in fwd and cuDNN FROST (chunk 64) in fwd+bwd and ckpt. Where this comparison was not stable enough to publish, the table shows "unstable (not published)" and no ratio.
  • "Previous" is this package at the same main commit.

R200 was timed with FP32 beta only; this PR does not change the R200 program. On R200 (sm_107a), the package keeps the previous WY program, and every cell is at parity with the previous version (0.994-1.001). Against deterministic Cake, it is 1.000-1.186. Against non-Cake SOTA, the ratio is above 1 in all 19 fwd+bwd cells (1.258-2.527) and all 14 published ckpt cells (1.106-2.066). In fwd, it is below 1 in 12 of 16 published cells (0.452-0.860) and above 1 in the other 4 (1.155-1.291).

NVIDIA R200 (sm_107a, 212 SMs), CUDA 13.0 (torch 2.14.1+cu130): BF16 q / k / v, FP32 beta, head dim 128, chunk 64, eager; ms per call = mean of two runs in opposite arm orders, each the median of 6 group medians of 16 samples (L2 flushed before each sample); speedup = baseline ms / this PR ms.

GPU / architecture workload (H / HV) operation deterministic Cake ms non-Cake SOTA ms previous non-deterministic Cake ms this PR ms speedup vs deterministic Cake speedup vs non-Cake SOTA speedup vs previous
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, H/HV 8/16 fwd 0.1664 0.1960 (cuDNN FROST) 0.1655 0.1661 1.002x 1.180x 0.996x
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, H/HV 8/16 fwd+bwd 0.3765 0.8263 (cuDNN FROST (chunk 64)) 0.3751 0.3756 1.002x 2.200x 0.999x
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, H/HV 8/16 ckpt 0.6396 1.3010 (cuDNN FROST (chunk 64)) 0.6325 0.6347 1.008x 2.050x 0.997x
NVIDIA R200 (212 SMs) sm_107a 3 x ~4.25k (real pack), 8/16 fwd 0.1445 0.1837 (cuDNN FROST) 0.1440 0.1440 1.003x 1.276x 1.000x
NVIDIA R200 (212 SMs) sm_107a 3 x ~4.25k (real pack), 8/16 fwd+bwd 0.3222 0.7535 (cuDNN FROST (chunk 64)) 0.3208 0.3207 1.005x 2.350x 1.000x
NVIDIA R200 (212 SMs) sm_107a 3 x ~4.25k (real pack), 8/16 ckpt 0.6073 1.2363 (cuDNN FROST (chunk 64)) 0.5961 0.5983 1.015x 2.066x 0.996x
NVIDIA R200 (212 SMs) sm_107a 8 x 4096, 8/16 fwd 0.2887 0.1853 (cuDNN FROST) 0.2879 0.2882 1.002x 0.643x 0.999x
NVIDIA R200 (212 SMs) sm_107a 8 x 4096, 8/16 fwd+bwd 0.6833 1.1721 (cuDNN FROST (chunk 64)) 0.6777 0.6777 1.008x 1.730x 1.000x
NVIDIA R200 (212 SMs) sm_107a 8 x 4096, 8/16 ckpt 1.0075 1.6330 (cuDNN FROST (chunk 64)) 1.0016 1.0026 1.005x 1.629x 0.999x
NVIDIA R200 (212 SMs) sm_107a 7 x ~4.3k (real pack), 8/16 fwd 0.2692 0.2310 (cuDNN FROST) 0.2683 0.2687 1.002x 0.860x 0.999x
NVIDIA R200 (212 SMs) sm_107a 7 x ~4.3k (real pack), 8/16 fwd+bwd 0.6316 1.2120 (cuDNN FROST (chunk 64)) 0.6310 0.6318 1.000x 1.918x 0.999x
NVIDIA R200 (212 SMs) sm_107a 7 x ~4.3k (real pack), 8/16 ckpt 0.9352 1.6617 (cuDNN FROST (chunk 64)) 0.9336 0.9351 1.000x 1.777x 0.998x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 8/16 fwd 0.4858 0.3007 (cuDNN FROST) 0.4841 0.4843 1.003x 0.621x 1.000x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 8/16 fwd+bwd 1.1901 2.0595 (cuDNN FROST (chunk 64)) 1.1855 1.1871 1.003x 1.735x 0.999x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 8/16 ckpt 1.7123 2.5015 (cuDNN FROST (chunk 64)) 1.7080 1.7106 1.001x 1.462x 0.998x
NVIDIA R200 (212 SMs) sm_107a 15 x ~4.3k (real pack), 8/16 fwd 0.5115 0.2667 (cuDNN FROST) 0.5087 0.5098 1.003x 0.523x 0.998x
NVIDIA R200 (212 SMs) sm_107a 15 x ~4.3k (real pack), 8/16 fwd+bwd 1.2430 1.9417 (cuDNN FROST (chunk 64)) 1.2364 1.2380 1.004x 1.568x 0.999x
NVIDIA R200 (212 SMs) sm_107a 15 x ~4.3k (real pack), 8/16 ckpt 1.7901 2.3978 (cuDNN FROST (chunk 64)) 1.7833 1.7857 1.002x 1.343x 0.999x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 8/16 fwd 0.1964 0.1987 (cuDNN FROST) 0.1720 0.1720 1.142x 1.155x 1.000x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 8/16 fwd+bwd 0.4337 0.8346 (cuDNN FROST (chunk 64)) 0.3891 0.3895 1.113x 2.143x 0.999x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 8/16 ckpt 0.6932 unstable (not published) 0.6459 0.6463 1.073x - 0.999x
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, 16/32 fwd 0.2915 0.1844 (cuDNN FROST) 0.2904 0.2909 1.002x 0.634x 0.998x
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, 16/32 fwd+bwd 0.6918 1.1791 (cuDNN FROST (chunk 64)) 0.6876 0.6878 1.006x 1.714x 1.000x
NVIDIA R200 (212 SMs) sm_107a 4 x 4096, 16/32 ckpt 1.0188 unstable (not published) 1.0153 1.0164 1.002x - 0.999x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 16/32 fwd 0.9323 0.4212 (cuDNN FROST) 0.9313 0.9324 1.000x 0.452x 0.999x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 16/32 fwd+bwd 2.3143 3.3969 (cuDNN FROST (chunk 64)) 2.3099 2.3107 1.002x 1.470x 1.000x
NVIDIA R200 (212 SMs) sm_107a 16 x 4096, 16/32 ckpt 3.2846 3.9559 (cuDNN FROST (chunk 64)) 3.2808 3.2827 1.001x 1.205x 0.999x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 16/32 fwd 0.2816 unstable (not published) 0.2810 0.2810 1.002x - 1.000x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 16/32 fwd+bwd 0.7252 1.2262 (cuDNN FROST (chunk 64)) 0.6641 0.6654 1.090x 1.843x 0.998x
NVIDIA R200 (212 SMs) sm_107a 1 x 16384, 16/32 ckpt 1.0435 1.6893 (cuDNN FROST (chunk 64)) 0.9817 0.9831 1.061x 1.718x 0.999x
NVIDIA R200 (212 SMs) sm_107a 131071 + 87531 + 43385, 2/6 fwd 0.6818 0.5105 (cuDNN FROST) 0.6811 0.6809 1.001x 0.750x 1.000x
NVIDIA R200 (212 SMs) sm_107a 131071 + 87531 + 43385, 2/6 fwd+bwd 1.9152 2.9953 (cuDNN FROST (chunk 64)) 1.9150 1.9150 1.000x 1.564x 1.000x
NVIDIA R200 (212 SMs) sm_107a 131071 + 87531 + 43385, 2/6 ckpt 2.6295 3.5981 (cuDNN FROST (chunk 64)) 2.6304 2.6289 1.000x 1.369x 1.001x
NVIDIA R200 (212 SMs) sm_107a 2 seqs, 259174 tokens, 2/6 fwd 0.6853 0.5040 (cuDNN FROST) 0.6842 0.6842 1.002x 0.737x 1.000x
NVIDIA R200 (212 SMs) sm_107a 2 seqs, 259174 tokens, 2/6 fwd+bwd 1.9102 2.9655 (cuDNN FROST (chunk 64)) 1.9084 1.9088 1.001x 1.554x 1.000x
NVIDIA R200 (212 SMs) sm_107a 2 seqs, 259174 tokens, 2/6 ckpt 2.6265 3.5593 (cuDNN FROST (chunk 64)) 2.6252 2.6263 1.000x 1.355x 1.000x
NVIDIA R200 (212 SMs) sm_107a 6 seqs, 241817 tokens, 2/6 fwd 0.6532 0.5367 (cuDNN FROST) 0.6532 0.6534 1.000x 0.821x 1.000x
NVIDIA R200 (212 SMs) sm_107a 6 seqs, 241817 tokens, 2/6 fwd+bwd 1.8149 2.9744 (cuDNN FROST (chunk 64)) 1.8051 1.8062 1.005x 1.647x 0.999x
NVIDIA R200 (212 SMs) sm_107a 6 seqs, 241817 tokens, 2/6 ckpt 2.5064 3.6067 (cuDNN FROST (chunk 64)) 2.4973 2.4993 1.003x 1.443x 0.999x
NVIDIA R200 (212 SMs) sm_107a 64 x 4096, 2/6 fwd 0.6639 0.3010 (cuDNN FROST) 0.6626 0.6639 1.000x 0.453x 0.998x
NVIDIA R200 (212 SMs) sm_107a 64 x 4096, 2/6 fwd+bwd 1.8824 2.4361 (cuDNN FROST (chunk 64)) 1.8806 1.8807 1.001x 1.295x 1.000x
NVIDIA R200 (212 SMs) sm_107a 64 x 4096, 2/6 ckpt 2.5845 2.8555 (cuDNN FROST (chunk 64)) 2.5814 2.5821 1.001x 1.106x 1.000x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 16/48 fwd 0.7677 0.3384 (cuDNN FROST) 0.7210 0.7208 1.065x 0.469x 1.000x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 16/48 fwd+bwd 2.0504 2.5312 (cuDNN FROST (chunk 64)) 2.0089 2.0114 1.019x 1.258x 0.999x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 16/48 ckpt 2.8519 unstable (not published) 2.7642 2.7637 1.032x - 1.000x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 8/24 fwd 0.3855 0.3167 (cuDNN FROST) 0.3846 0.3849 1.002x 0.823x 0.999x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 8/24 fwd+bwd 1.0523 1.7575 (cuDNN FROST (chunk 64)) 1.0518 1.0525 1.000x 1.670x 0.999x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 8/24 ckpt 1.4749 2.2035 (cuDNN FROST (chunk 64)) 1.4739 1.4745 1.000x 1.494x 1.000x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 4/12 fwd 0.2165 unstable (not published) 0.2156 0.2159 1.003x - 0.999x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 4/12 fwd+bwd 0.5613 1.0240 (cuDNN FROST (chunk 64)) 0.5607 0.5611 1.000x 1.825x 0.999x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 4/12 ckpt 0.8211 1.4980 (cuDNN FROST (chunk 64)) 0.8169 0.8171 1.005x 1.833x 1.000x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 2/6 fwd 0.1560 0.1811 (cuDNN FROST) 0.1394 0.1403 1.112x 1.291x 0.994x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 2/6 fwd+bwd 0.3691 0.7347 (cuDNN FROST (chunk 64)) 0.3405 0.3414 1.081x 2.152x 0.997x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 2/6 ckpt 0.6391 unstable (not published) 0.6154 0.6182 1.034x - 0.995x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 1/3 fwd 0.1306 unstable (not published) 0.1099 0.1101 1.186x - 0.998x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 1/3 fwd+bwd 0.2815 0.6136 (cuDNN FROST (chunk 64)) 0.2422 0.2428 1.159x 2.527x 0.998x
NVIDIA R200 (212 SMs) sm_107a 11000 + 12000 + 9768, 1/3 ckpt 0.5650 unstable (not published) 0.5407 0.5430 1.041x - 0.996x

Fast (non-deterministic) KCP

The default of this layer is save_mode="full" (eager). flashinfer-ai/flashinfer#8644 times that form. flashinfer-ai/flashinfer#8638, the exact revision, changes only graph="replay" calls, which this layer does not make; its Cake arms in table B run in graph mode.

flashinfer-ai/flashinfer#8644: eager, save_mode="full", two baselines

det(main) is the deterministic KCP package of the same main, and this PR does not change it. In table A this PR is faster than or equal to det(main) in every cell. The det(main) / this PR ratio runs from 1.000 to 1.645 (R200 shape I CP4 C fwd+bwd, 2.913 -> 1.771 ms). In table B, det(main) is faster in one cell: B200 shape H CP8 C fwd, at 1.542 ms vs 1.552 ms for this PR (0.994). In table A the same cell is 1.530 ms on both sides.

Table B also compares this PR with the strongest accuracy-qualified non-Cake arm. Its ratios are non-Cake / this PR, so a value below 1 means the non-Cake arm is faster. At CP1, that arm is a head split around the cuDNN frontend GDN (FROST) in every published cell, and it runs in the same process as this PR. Table B has nine B200 CP1 cells for shapes A, G and M, and seven of them were measured under the sustained protocol. In those nine cells, det(main) / this PR is 1.015-1.123, and FROST is faster than this PR in all nine, at 0.348-0.832. Across the three GPUs, FROST is faster than this PR in 24 of the 26 published CP1 cells, at 0.283-0.933. The widest gap is R200 shape A fwd: 1.885 ms for FROST vs 6.663 ms for this PR. At CP1 this PR is ahead only in R200 fwd+bwd on shapes G and M (1.019 and 1.026).

At CP > 1, the strongest applicable non-Cake arm is FlashQLA 0.1.3 with dual context parallelism (inter- and intra-rank), in the contiguous layout. It was measured on B200 at CP8 and on GB300 at CP4. In the forward at CP > 1, this PR trails FlashQLA in every measured cell. FlashQLA / this PR is 0.736-0.849 on B200 CP8 and 0.738-0.888 on GB300 CP4. The widest gap is at shape H: 1.142 ms vs 1.552 ms on B200 CP8, and 1.883 ms vs 2.552 ms on GB300 CP4. This PR is faster in fwd+bwd, where FlashQLA / this PR is 1.043-1.124 on B200 CP8 and 1.084-1.208 on GB300 CP4. It is also faster in ckpt, at 1.034-1.040 and 1.008-1.128. Layout D (per-document zigzag) cells are not claimed because FlashQLA has no zigzag layout. The R200 CP4 non-Cake cells are not claimed because FlashQLA has no sm_107a build. For both of these groups, table B shows the best in-process arm as a note only.

Workloads: BF16 q / k / v, head dim 128, save_mode="full", in-kernel Q / K L2 norm (use_qk_l2norm_in_kernel=True), eager, GDN layer only (no conv1d); layouts C = contiguous, D = per-document zigzag; CPn = context parallel over n GPUs. Time = median over 5 groups x 12 samples (10 warm-up calls, L2 flushed before each call) of the slowest rank, ms. Operations: fwd = forward with autograd recording; fwd+bwd = forward + backward; ckpt = torch.utils.checkpoint(use_reentrant=False): forward + recompute + backward (2 x fwd + bwd).

shape documents (tokens) tokens heads (q/k, v)
A 4: 98304/65536/57344/40960 262144 (16, 48)
G 12: 3598/4930/6262/7593/8925/10256/11588/12920/14251/15583/16914/18252 131072 (16, 48)
H 4: 49152/32768/28672/20480 131072 (16, 48)
I 4: 24576/16384/14336/10240 65536 (16, 48)
M 5: 12000/9000/6000/3768/2000 32768 (16, 48)

Legend: '*' = ratio below 1.000 (this PR slower than that baseline). Ratios = baseline ms / this PR ms from the shown absolutes. '(sustained)' = a cell whose short run hit the GPU software power cap (>= 5 % of samples), re-measured under the sustained protocol (idle cool-down before each block, 300-sample blocks, 3 block pairs per arm in both orders; value = median of the block medians); 'power-capped' = capped short-run cell without a sustained re-measure: no latency is claimed.

B. Two baselines (det(main), this PR and the in-process non-Cake arms in one run; ‡ = FlashQLA in its own process)

Strongest non-Cake = the fastest arm that passes the accuracy check: FLA-CP = flash-linear-attention 0.6.0's own context parallelism; FLA-HS = head split (all-to-all) around FLA's single-GPU kernels; FROST-ck0 / -ck64 = the same head split around cuDNN frontend 1.32.1's GDN (states recomputed / stored every 64 tokens). A non-Cake cell is published when this PR's arm and the shown arm drift <= 2 % between groups and every arm with drift > 2 % is slower than the shown arm in every group. At CP > 1 the strongest applicable non-Cake arm is FlashQLA 0.1.3 dual-CP (inter + intra rank context parallelism, contiguous layout): it runs in its own process and environment, paired with FLA-CP there and checked against it (relative RMS <= 1e-2, max abs <= 2e-2 of the reference max); '‡' = a FlashQLA value, its ratio divides by this PR's ms of the same table-B run, and the FLA-CP anchor of both processes is listed under the table. A CP > 1 cell without a FlashQLA value (layout D: FlashQLA has no per-document zigzag layout; a GPU without a FlashQLA build) is not claimed; the best in-process arm is shown as a note only.

GPU / arch workload (shape, CP, layout) op det(main) Cake ms strongest non-Cake ms (arm) this PR ms det(main) / this PR non-Cake / this PR
B200 sm_100a A CP1 C fwd (sustained) 11.914 3.866 (FROST-ck0) 11.104 1.073 0.348*
B200 sm_100a A CP1 C fwd+bwd (sustained) 30.379 20.784 (FROST-ck64) 29.592 1.027 0.702*
B200 sm_100a A CP1 C ckpt (sustained) 39.761 25.885 (FROST-ck64) 38.944 1.021 0.665*
B200 sm_100a G CP1 C fwd (sustained) 5.980 1.861 (FROST-ck0) 5.324 1.123 0.350*
B200 sm_100a G CP1 C fwd+bwd (sustained) 15.295 10.642 (FROST-ck64) 14.631 1.045 0.727*
B200 sm_100a G CP1 C ckpt (sustained) 19.735 13.097 (FROST-ck64) 19.437 1.015 0.674*
B200 sm_100a M CP1 C fwd 1.695 0.594 (FROST-ck0) 1.521 1.114 0.391*
B200 sm_100a M CP1 C fwd+bwd (sustained) 4.077 3.244 (FROST-ck64) 3.897 1.046 0.832*
B200 sm_100a M CP1 C ckpt 5.411 4.238 (FROST-ck64) 5.248 1.031 0.808*
B200 sm_100a A CP8 C fwd 2.583 1.828 (FlashQLA 0.1.3 dual-CP‡) 2.384 1.083 0.767*‡
B200 sm_100a A CP8 C fwd+bwd 5.322 5.749 (FlashQLA 0.1.3 dual-CP‡) 5.117 1.040 1.124‡
B200 sm_100a A CP8 C ckpt 7.495 7.525 (FlashQLA 0.1.3 dual-CP‡) 7.277 1.030 1.034‡
B200 sm_100a A CP8 D fwd 2.597 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 1.828 ms); best in-process arm 4.041 (FROST-ck0) 2.424 1.071 not claimed
B200 sm_100a A CP8 D fwd+bwd 5.518 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 5.749 ms); best in-process arm 10.234 (FROST-ck64) 5.335 1.034 not claimed
B200 sm_100a A CP8 D ckpt 7.715 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 7.525 ms); best in-process arm 13.407 (FROST-ck64) 7.512 1.027 not claimed
B200 sm_100a H CP8 C fwd 1.542 1.142 (FlashQLA 0.1.3 dual-CP‡) 1.552 0.994* 0.736*‡
B200 sm_100a H CP8 C fwd+bwd 3.720 3.316 (FlashQLA 0.1.3 dual-CP‡) 3.180 1.170 1.043‡
B200 sm_100a H CP8 C ckpt 5.421 4.645 (FlashQLA 0.1.3 dual-CP‡) 4.475 1.211 1.038‡
B200 sm_100a H CP8 D fwd 1.604 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 1.142 ms); best in-process arm 2.192 (FROST-ck0) 1.491 1.076 not claimed
B200 sm_100a H CP8 D fwd+bwd 3.489 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 3.316 ms); best in-process arm 5.457 (FROST-ck64) 3.117 1.119 not claimed
B200 sm_100a H CP8 D ckpt 5.050 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 4.645 ms); best in-process arm 7.250 (FROST-ck64) 4.428 1.140 not claimed
B200 sm_100a I CP8 C fwd 1.268 0.830 (FlashQLA 0.1.3 dual-CP‡) 0.978 1.297 0.849*‡
B200 sm_100a I CP8 C fwd+bwd 3.118 2.426 (FlashQLA 0.1.3 dual-CP‡) 2.262 1.378 1.073‡
B200 sm_100a I CP8 C ckpt 4.777 3.593 (FlashQLA 0.1.3 dual-CP‡) 3.456 1.382 1.040‡
B200 sm_100a I CP8 D fwd 1.169 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 0.830 ms); best in-process arm 1.248 (FROST-ck0) 1.030 1.135 not claimed
B200 sm_100a I CP8 D fwd+bwd 3.018 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 2.426 ms); best in-process arm 3.036 (FROST-ck64) 2.233 1.352 not claimed
B200 sm_100a I CP8 D ckpt 4.585 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 3.593 ms); best in-process arm 4.186 (FROST-ck64) 3.388 1.353 not claimed
GB300 sm_103a A CP1 C fwd 11.191 3.127 (FROST-ck0) 10.285 1.088 0.304*
GB300 sm_103a A CP1 C fwd+bwd 28.106 18.842 (FROST-ck64) 27.086 1.038 0.696*
GB300 sm_103a A CP1 C ckpt 36.914 23.159 (FROST-ck64) 35.896 1.028 0.645*
GB300 sm_103a G CP1 C fwd 5.659 1.655 (FROST-ck0) 4.961 1.141 0.334*
GB300 sm_103a G CP1 C fwd+bwd 13.868 9.817 (FROST-ck64) 13.186 1.052 0.745*
GB300 sm_103a G CP1 C ckpt 18.246 12.032 (FROST-ck64) 17.569 1.039 0.685*
GB300 sm_103a M CP1 C fwd 1.702 0.651 (FROST-ck0) 1.502 1.133 0.433*
GB300 sm_103a M CP1 C fwd+bwd 3.859 3.101 (FROST-ck64) 3.662 1.054 0.847*
GB300 sm_103a M CP1 C ckpt (re-measure) 5.126 4.169 (FROST-ck64) 4.894 1.047 not published (drift)
GB300 sm_103a A CP4 C fwd 4.274 3.284 (FlashQLA 0.1.3 dual-CP‡) 3.873 1.104 0.848*‡
GB300 sm_103a A CP4 C fwd+bwd 9.049 10.453 (FlashQLA 0.1.3 dual-CP‡) 8.655 1.046 1.208‡
GB300 sm_103a A CP4 C ckpt 12.556 13.710 (FlashQLA 0.1.3 dual-CP‡) 12.151 1.033 1.128‡
GB300 sm_103a A CP4 D fwd 4.331 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 3.284 ms); best in-process arm 7.253 (FROST-ck0) 3.987 1.086 not claimed
GB300 sm_103a A CP4 D fwd+bwd 9.288 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 10.453 ms); best in-process arm 18.667 (FROST-ck64) 8.957 1.037 not claimed
GB300 sm_103a A CP4 D ckpt 12.862 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 13.710 ms); best in-process arm 24.397 (FROST-ck64) 12.525 1.027 not claimed
GB300 sm_103a H CP4 C fwd 2.712 1.883 (FlashQLA 0.1.3 dual-CP‡) 2.552 1.063 0.738*‡
GB300 sm_103a H CP4 C fwd+bwd 5.450 5.584 (FlashQLA 0.1.3 dual-CP‡) 5.150 1.058 1.084‡
GB300 sm_103a H CP4 C ckpt 7.772 7.395 (FlashQLA 0.1.3 dual-CP‡) 7.333 1.060 1.008‡
GB300 sm_103a H CP4 D fwd 2.495 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 1.883 ms); best in-process arm 3.868 (FROST-ck0) 2.276 1.096 not claimed
GB300 sm_103a H CP4 D fwd+bwd 5.110 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 5.584 ms); best in-process arm 9.775 (FROST-ck64) 4.905 1.042 not claimed
GB300 sm_103a H CP4 D ckpt 7.140 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 7.395 ms); best in-process arm 12.852 (FROST-ck64) 6.926 1.031 not claimed
GB300 sm_103a I CP4 C fwd 2.051 1.311 (FlashQLA 0.1.3 dual-CP‡) 1.477 1.389 0.888*‡
GB300 sm_103a I CP4 C fwd+bwd 4.437 3.375 (FlashQLA 0.1.3 dual-CP‡) 2.828 1.569 1.193‡
GB300 sm_103a I CP4 C ckpt 6.873 5.059 (FlashQLA 0.1.3 dual-CP‡) 4.489 1.531 1.127‡
GB300 sm_103a I CP4 D fwd 1.827 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 1.311 ms); best in-process arm 2.085 (FROST-ck0) 1.427 1.280 not claimed
GB300 sm_103a I CP4 D fwd+bwd 4.041 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 3.375 ms); best in-process arm 5.148 (FROST-ck64) 2.889 1.399 not claimed
GB300 sm_103a I CP4 D ckpt 6.256 not claimed (FlashQLA 0.1.3 dual-CP: no per-document zigzag layout; contiguous layout, same documents: 5.059 ms); best in-process arm 6.889 (FROST-ck64) 4.393 1.424 not claimed
R200 sm_107a A CP1 C fwd 6.932 1.885 (FROST-ck0) 6.663 1.040 0.283*
R200 sm_107a A CP1 C fwd+bwd 17.688 16.279 (FROST-ck64) 17.439 1.014 0.933*
R200 sm_107a A CP1 C ckpt 23.345 18.930 (FROST-ck64) 23.113 1.010 0.819*
R200 sm_107a G CP1 C fwd 3.324 0.960 (FROST-ck0) 2.948 1.128 0.326*
R200 sm_107a G CP1 C fwd+bwd 8.642 8.415 (FROST-ck64) 8.261 1.046 1.019
R200 sm_107a G CP1 C ckpt 11.316 9.732 (FROST-ck64) 10.923 1.036 0.891*
R200 sm_107a M CP1 C fwd 1.000 0.353 (FROST-ck0) 0.896 1.116 0.394*
R200 sm_107a M CP1 C fwd+bwd 2.381 2.334 (FROST-ck64) 2.275 1.047 1.026
R200 sm_107a M CP1 C ckpt 3.167 2.847 (FROST-ck64) 3.059 1.035 0.931*
R200 sm_107a A CP4 C fwd 2.765 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 6.460 (FROST-ck0) 2.560 1.080 not claimed
R200 sm_107a A CP4 C fwd+bwd 6.086 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 16.702 (FROST-ck64) 5.885 1.034 not claimed
R200 sm_107a A CP4 C ckpt 8.375 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 21.441 (FROST-ck64) 8.221 1.019 not claimed
R200 sm_107a A CP4 D fwd 2.780 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 6.460 (FROST-ck0) 2.581 1.077 not claimed
R200 sm_107a A CP4 D fwd+bwd 5.957 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 16.707 (FROST-ck64) 5.763 1.034 not claimed
R200 sm_107a A CP4 D ckpt 8.373 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 21.444 (FROST-ck64) 8.170 1.025 not claimed
R200 sm_107a H CP4 C fwd 1.592 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 3.412 (FROST-ck0) 1.422 1.120 not claimed
R200 sm_107a H CP4 C fwd+bwd 3.345 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 8.714 (FROST-ck64) 3.042 1.100 not claimed
R200 sm_107a H CP4 C ckpt 4.723 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 11.248 (FROST-ck64) 4.348 1.086 not claimed
R200 sm_107a H CP4 D fwd 1.578 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 3.414 (FROST-ck0) 1.450 1.088 not claimed
R200 sm_107a H CP4 D fwd+bwd 3.297 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 8.714 (FROST-ck64) 3.171 1.040 not claimed
R200 sm_107a H CP4 D ckpt 4.667 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 11.251 (FROST-ck64) 4.533 1.030 not claimed
R200 sm_107a I CP4 C fwd 0.982 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 1.824 (FROST-ck0) 0.864 1.137 not claimed
R200 sm_107a I CP4 C fwd+bwd 2.246 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 4.570 (FROST-ck64) 1.776 1.265 not claimed
R200 sm_107a I CP4 C ckpt 3.566 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 5.965 (FROST-ck64) 2.545 1.401 not claimed
R200 sm_107a I CP4 D fwd 0.980 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 1.822 (FROST-ck0) 0.903 1.085 not claimed
R200 sm_107a I CP4 D fwd+bwd 1.989 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 4.566 (FROST-ck64) 1.887 1.054 not claimed
R200 sm_107a I CP4 D ckpt 3.152 not claimed (FlashQLA 0.1.3 dual-CP: not available for sm_107a); best in-process arm 5.964 (FROST-ck64) 2.721 1.158 not claimed

FLA-CP anchor of the ‡ cells (FlashQLA process / table-B run, ms): B200 sm_100a A CP8 fwd 9.561 / 9.583; B200 sm_100a A CP8 fwd+bwd 25.743 / 25.808; B200 sm_100a A CP8 ckpt 35.250 / 35.297; B200 sm_100a H CP8 fwd 5.040 / 5.075; B200 sm_100a H CP8 fwd+bwd 13.189 / 13.209; B200 sm_100a H CP8 ckpt 18.071 / 18.099; B200 sm_100a I CP8 fwd 2.777 / 2.795; B200 sm_100a I CP8 fwd+bwd 7.009 / 7.013; B200 sm_100a I CP8 ckpt 9.603 / 9.615; GB300 sm_103a A CP4 fwd 17.655 / 18.049; GB300 sm_103a A CP4 fwd+bwd 41.095 / 41.809; GB300 sm_103a A CP4 ckpt 58.789 / 59.796; GB300 sm_103a H CP4 fwd 8.959 / 9.113; GB300 sm_103a H CP4 fwd+bwd 20.638 / 20.970; GB300 sm_103a H CP4 ckpt 29.454 / 29.939; GB300 sm_103a I CP4 fwd 4.768 / 4.855; GB300 sm_103a I CP4 fwd+bwd 10.629 / 10.784; GB300 sm_103a I CP4 ckpt 15.192 / 15.420.

B200 CP1 cells marked (sustained) in table B: one process per shape (A; G; M), each holding det(main), this PR, previous and the four in-process non-Cake arms, at the PR head.

Table A of flashinfer-ai/flashinfer#8644 (this PR against the previous fast KCP package, 81 cells, with a det(main) column from the same runs) is not copied here. Its summary:

"Previous" is this package at FlashInfer main f2926f20. Ratios are previous / this PR, so a value above 1 means this PR is faster. In table A, no cell is slower than previous on any GPU. On B200 and GB300, 23 of 27 cells are faster and 4 are at parity. On a 212-SM R200, all 27 cells are faster. The parity cells are all contiguous-layout cells of shapes H and I, at CP8 on B200 and CP4 on GB300. The forward gains the most: 1.004-1.104 on B200, 1.011-1.111 on GB300 and 1.024-1.103 on R200. On every GPU the best cell is shape G at CP1: 5.899 -> 5.344 ms on B200, 5.522 -> 4.972 ms on GB300 and 3.252 -> 2.947 ms on R200. fwd+bwd gains 1.003-1.042 on B200, 1.002-1.043 on GB300 and 1.011-1.038 on R200. The checkpointed step (ckpt) gains 1.004-1.041, 1.001-1.032 and 1.007-1.030 on the same three GPUs.

@yyihuang

yyihuang commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

Performance (5/5), continued (fast KCP, continued; det kernels with FLA's context parallelism; missing measurements)

As in part 1, "this PR", "new", "previous" and "main" in the copied FlashInfer text refer to that FlashInfer pull request and to the FlashInfer package before it, not to this PR.

flashinfer-ai/flashinfer#8638 (exact revision): PR head, graph mode

det(main) is the deterministic KCP package of the same main. Against it, this PR is faster or at parity in every published A4 cell:

  • fwd+bwd: 1.029-1.059
  • M ckpt: 1.344-1.361 (B200 6.551 -> 4.845 ms)
  • N ckpt: 1.007-1.021 on B200, 1.073-1.075 on GB300 and 1.119-1.126 on R200

Table B also compares this PR with the strongest accuracy-qualified non-Cake arm. Its ratios are non-Cake ms / this PR ms, so above 1 means this PR is faster. At CP1, the cuDNN frontend GDN arms (FROST) are faster than this PR on shapes G, A and M. There non-Cake / this PR is 0.396-0.866 on B200, 0.341-0.954 on GB300 and 0.316-0.938 on R200. The one exception is R200 shape M ckpt, where the ratio is 1.003. On shape N, this PR is faster where a cell is published: 1.888-1.902 on B200 and 1.562 on R200. At CP > 1, the strongest accuracy-qualified non-Cake arm is FlashQLA 0.1.3 with its own dual context parallelism. FlashQLA needs its own environment, so it runs in a separate process. FLA-CP runs in both processes and is the anchor between them. On B200 at CP8, FlashQLA / this PR is 0.887 in fwd, 0.928 in fwd+bwd and 1.102 in ckpt. On GB300 at CP4, it is 0.951, 0.983 and 1.208. FlashQLA is faster in fwd and fwd+bwd, and this PR is faster in ckpt. Table B has no CP > 1 rows on R200.

Shapes (head dim 128, BF16 q / k / v, FP32 g / beta, in-kernel q / k L2 norm, layout contiguous, GDN layer only): A = 262144 tokens, 4 documents (98304 / 65536 / 57344 / 40960), heads (16, 48) (q/k heads, v heads); G = 131072 tokens, 12 documents (3598 ... 18252), heads (16, 48); M = 32768 tokens, 5 documents (12000 / 9000 / 6000 / 3768 / 2000), heads (16, 48); N = 8192 tokens, 3 documents (2049 / 3001 / 3142), heads (2, 6). CPn = context parallel over n GPUs. The single-rank proxy forces two segments per fragment (the segment walks and carry merges a CP > 1 fragment runs) on one GPU. Non-Cake arms: FLA-CP = flash-linear-attention 0.6.0's own context parallelism; FLA-HS = head split (all-to-all) around FLA's single-GPU kernels; FROST-ck0 / -ck64 = the same head split around cuDNN frontend 1.32.1's GDN (states recomputed / stored every 64 tokens).

Marks used in these tables, from the legend of flashinfer-ai/flashinfer#8638: * = ratio below 1.000 (this PR slower than that baseline); § = a power-capped cell taken from the sustained protocol (median of the block medians [min, max], each arm's share of samples at the power limit and its median SM clock); ‡ = FlashQLA 0.1.3 with its own dual context parallelism, timed in a separate process; † = a run published by the paired-ratio test only; n/p = not published.

B200 (sm_100a):

B. Two baselines at the PR head (graph mode for the Cake arms; the non-Cake arms have no graph form here and run eager; det(main) = the deterministic KCP package of main f2926f20)

workload operation det(main) ms this PR ms det(main) / this PR strongest non-Cake ms (arm) this PR ms (same process) non-Cake / this PR
G CP1 fwd 5.728 [5.726, 5.730] cap 10 % 1965 MHz 4.639 [4.633, 4.642] cap 0 % 1965 MHz 1.235§ 1.838 [1.834, 1.845] cap 75 % 1800 MHz (FROST-ck0) 4.639 [4.633, 4.642] cap 0 % 1965 MHz 0.396*§
G CP1 fwd+bwd 19.026 [18.837, 19.100] cap 87 % 1920 MHz 17.768 [17.744, 17.788] cap 83 % 1935 MHz 1.071§ 10.630 [10.618, 10.647] cap 0 % 1965 MHz (FROST-ck64) 17.768 [17.744, 17.788] cap 83 % 1935 MHz 0.598*§
G CP1 ckpt 24.784 [24.746, 24.809] cap 93 % 1935 MHz 17.880 [17.876, 17.884] cap 83 % 1942 MHz 1.386§ 12.981 [12.962, 13.002] cap 33 % 1965 MHz (FROST-ck64) 17.880 [17.876, 17.884] cap 83 % 1942 MHz 0.726*§
A CP1 fwd 11.627 [11.607, 11.637] cap 63 % 1957 MHz 9.871 [9.869, 9.872] cap 0 % 1965 MHz 1.178§ 3.894 [3.884, 3.910] cap 86 % 1717 MHz (FROST-ck0) 9.824 [9.819, 9.830] cap 15 % 1965 MHz 0.396*§
A CP1 fwd+bwd 38.827 [38.573, 38.902] cap 97 % 1897 MHz 36.873 [36.788, 36.930] cap 94 % 1920 MHz 1.053§ 20.717 [20.678, 20.733] cap 57 % 1954 MHz (FROST-ck64) 37.165 [37.133, 37.211] cap 98 % 1905 MHz 0.557*§
A CP1 ckpt 50.521 [50.499, 50.574] cap 97 % 1920 MHz 36.989 [36.978, 37.009] cap 93 % 1920 MHz 1.366§ 26.005 [25.977, 26.095] cap 96 % 1905 MHz (FROST-ck64) 37.277 [37.127, 37.306] cap 98 % 1905 MHz 0.698*§
A CP8 fwd 2.342 2.050 1.142 1.818 (FlashQLA 0.1.3 dual-CP‡) 2.050 0.887*‡
A CP8 fwd+bwd 6.473 [6.470, 6.474] cap 1 % 1965 MHz 6.187 [6.175, 6.193] cap 16 % 1965 MHz 1.046§ 5.744 [5.733, 5.747] cap 0 % 1965 MHz (FlashQLA 0.1.3 dual-CP‡) 6.187 [6.175, 6.193] cap 16 % 1965 MHz 0.928*‡§
A CP8 ckpt 8.903 [8.899, 8.905] cap 3 % 1965 MHz 6.831 [6.830, 6.835] cap 0 % 1965 MHz 1.303§ 7.528 [7.525, 7.532] cap 0 % 1965 MHz (FlashQLA 0.1.3 dual-CP‡) 6.831 [6.830, 6.835] cap 0 % 1965 MHz 1.102‡§
M CP1 fwd 1.509 1.258 1.200 0.592 (FROST-ck0) 1.258 0.471*
M CP1 fwd+bwd 4.940 4.692 1.053 3.216 (FROST-ck64) 4.692 0.685*
M CP1 ckpt 6.558 4.842 1.354 4.194 (FROST-ck64) 4.842 0.866*
N CP1 fwd 0.153 0.140 1.093 n/p -- --
N CP1 fwd+bwd 0.674 0.653 1.032 1.246 (FROST-ck0) 0.655 1.902
N CP1 ckpt 1.172 1.158 1.012 2.186 (FROST-ck0) 1.158 1.888

FLA-CP anchor of the FlashQLA cells (FlashQLA process / two-baseline process, ms): A CP8 fwd 9.575 / 9.636; A CP8 fwd+bwd (sustained) 25.765 / 25.736; A CP8 ckpt (sustained) 35.246 / 35.191.

A4. PR head (on FlashInfer main f2926f20, which holds flashinfer-ai/flashinfer#8600 and flashinfer-ai/flashinfer#8623): graph N / M fwd+bwd and checkpointed step (this PR vs previous = main f2926f20, same process, A/A arm, both arm orders twice)

workload fwd: previous ms -> this PR ms (previous / this PR per run, class) fwd+bwd ckpt (2 x fwd + bwd)
M CP1 not measured 4.644 -> 4.653 (0.998*/0.998*/0.998* REGRESSION) 6.050 -> 4.845 (1.249/1.248/1.247/1.247/1.247/1.249/1.248/1.249 win)
N CP1 not measured 0.643 -> 0.648 (0.992*/0.997*†/0.997*/0.997* parity) 1.130 -> 1.130 (1.000/0.992*/1.001/0.985*/0.991* parity)

Same runs, det(main) (the deterministic KCP package of main f2926f20) vs this PR:

workload fwd+bwd: det(main) ms -> this PR ms (det(main) / this PR per run, class) ckpt (2 x fwd + bwd)
M CP1 4.900 -> 4.653 (1.053/1.052/1.052 win) 6.551 -> 4.845 (1.352/1.354/1.352/1.354/1.352/1.356/1.353/1.356 win)
N CP1 0.664 -> 0.643 (1.033/1.034 win) 1.154 -> 1.130 (1.021/1.013/1.007/1.010/1.014 win/win/parity/parity/win)

GB300 (sm_103a):

B. Two baselines at the PR head (graph mode for the Cake arms; the non-Cake arms have no graph form here and run eager; det(main) = the deterministic KCP package of main f2926f20)

workload operation det(main) ms this PR ms det(main) / this PR strongest non-Cake ms (arm) this PR ms (same process) non-Cake / this PR
G CP1 fwd 5.394 4.351 1.240 1.714 (FROST-ck0) 4.351 0.394*
G CP1 fwd+bwd 17.451 16.386 1.065 9.949 (FROST-ck64) 16.386 0.607*
G CP1 ckpt 22.902 16.539 1.385 12.170 (FROST-ck64) 16.539 0.736*
A CP1 fwd 10.865 9.165 1.185 3.125 (FROST-ck0) 9.165 0.341*
A CP1 fwd+bwd 35.628 33.898 1.051 19.041 (FROST-ck64) 33.898 0.562*
A CP1 ckpt 46.605 34.087 1.367 23.351 (FROST-ck64) 34.087 0.685*
A CP4 fwd 4.006 3.452 1.160 3.283 (FlashQLA 0.1.3 dual-CP‡) 3.452 0.951*‡
A CP4 fwd+bwd 11.194 10.658 1.050 10.474 (FlashQLA 0.1.3 dual-CP‡) 10.658 0.983*‡
A CP4 ckpt 15.248 11.548 1.320 13.948 (FlashQLA 0.1.3 dual-CP‡) 11.548 1.208‡
M CP1 fwd 1.471 1.230 1.196 n/p -- --
M CP1 fwd+bwd 4.628 4.388 1.055 3.117 (FROST-ck64) 4.388 0.710*
M CP1 ckpt 6.146 4.528 1.357 4.320 (FROST-ck64) 4.528 0.954*
N CP1 fwd n/p unstable -- -- n/p -- --
N CP1 fwd+bwd 0.499 0.485 1.029 n/p -- --
N CP1 ckpt 1.015 0.963 1.054 n/p -- --

FLA-CP anchor of the FlashQLA cells (FlashQLA process / two-baseline process, ms): A CP4 fwd 17.784 / 17.875; A CP4 fwd+bwd 41.132 / 41.349; A CP4 ckpt 59.121 / 59.003.

A4. PR head (on FlashInfer main f2926f20, which holds flashinfer-ai/flashinfer#8600 and flashinfer-ai/flashinfer#8623): graph N / M fwd+bwd and checkpointed step (this PR vs previous = main f2926f20, same process, A/A arm, both arm orders twice)

workload fwd: previous ms -> this PR ms (previous / this PR per run, class) fwd+bwd ckpt (2 x fwd + bwd)
M CP1 not measured 4.352 -> 4.364 (0.997*/1.000/0.997*/1.000 parity/parity/REGRESSION/parity) 5.598 -> 4.495 (1.245/1.248/1.246/1.248 win)
N CP1 not measured 0.452 -> 0.457 (0.989*/0.993* parity) 0.909 -> 0.866 (1.050†/1.046†/1.049† win)

Same runs, det(main) (the deterministic KCP package of main f2926f20) vs this PR:

workload fwd+bwd: det(main) ms -> this PR ms (det(main) / this PR per run, class) ckpt (2 x fwd + bwd)
M CP1 4.607 -> 4.364 (1.056/1.059/1.056/1.058 win) 6.100 -> 4.495 (1.357/1.361/1.358/1.359 win)
N CP1 0.473 -> 0.458 (1.033†/1.029†/1.037/1.031 win) 0.936 -> 0.871 (1.075†/1.073† win)

R200 (sm_107a, 212 SMs):

B. Two baselines at the PR head (graph mode for the Cake arms; the non-Cake arms have no graph form here and run eager; det(main) = the deterministic KCP package of main f2926f20), on a 212-SM board

workload operation det(main) ms this PR ms det(main) / this PR strongest non-Cake ms (arm) this PR ms (same process) non-Cake / this PR
G CP1 fwd 3.179 2.609 1.218 0.959 (FROST-ck0) 2.609 0.368*
G CP1 fwd+bwd 10.827 10.270 1.054 8.363 (FROST-ck64) 10.270 0.814*
G CP1 ckpt 14.119 10.366 1.362 9.720 (FROST-ck64) 10.366 0.938*
A CP1 fwd 6.750 5.939 1.137 1.878 (FROST-ck0) 5.939 0.316*
A CP1 fwd+bwd 22.589 21.766 1.038 16.460 (FROST-ck64) 21.766 0.756*
A CP1 ckpt 29.409 21.895 1.343 18.897 (FROST-ck0) 21.895 0.863*
M CP1 fwd 0.886 0.757 1.170 0.349 (FROST-ck0) 0.757 0.461*
M CP1 fwd+bwd 2.901 2.769 1.048 2.337 (FROST-ck64) 2.769 0.844*
M CP1 ckpt 3.826 2.845 1.345 2.853 (FROST-ck64) 2.845 1.003
N CP1 fwd 0.115 0.108 1.065 n/p -- --
N CP1 fwd+bwd 0.333 0.323 1.031 0.506 (FROST-ck64) 0.324 1.562
N CP1 ckpt 0.545 0.491 1.110 n/p -- --

A4. PR head on a 212-SM board (on FlashInfer main f2926f20, which holds flashinfer-ai/flashinfer#8600 and flashinfer-ai/flashinfer#8623): graph N / M fwd+bwd and checkpointed step (this PR vs previous = main f2926f20, same process, A/A arm, both arm orders twice)

workload fwd: previous ms -> this PR ms (previous / this PR per run, class) fwd+bwd ckpt (2 x fwd + bwd)
M CP1 not measured 2.762 -> 2.760 (1.001/0.999*/1.001/1.000 parity) 3.554 -> 2.828 (1.257/1.253/1.257/1.253 win)
N CP1 not measured 0.325 -> 0.324 (1.003/1.000/1.000/1.000 parity) 0.518 -> 0.472 (1.097/1.105/1.096/1.101 win)

Same runs, det(main) (the deterministic KCP package of main f2926f20) vs this PR:

workload fwd+bwd: det(main) ms -> this PR ms (det(main) / this PR per run, class) ckpt (2 x fwd + bwd)
M CP1 2.891 -> 2.760 (1.047/1.046/1.048/1.046 win) 3.815 -> 2.828 (1.349/1.345/1.348/1.344 win)
N CP1 0.334 -> 0.324 (1.031/1.037/1.031/1.037 win) 0.528 -> 0.472 (1.119/1.126/1.120/1.120 win)

Tables B and A4 above time the PR head. Tables A, A2 and A3 of flashinfer-ai/flashinfer#8638 time earlier revisions of that PR and are not copied here.

Deterministic kernels with FLA's context parallelism (radixark/miles#3597)

These are the deterministic chunked GDN training kernels of radixark/miles#3597 (open; head 9d7543812b676a5b50dde107518782933b804b48, the vendored source commit). This layer runs them for an FLA cp_context under either variant. The PR body reports whole-model training times against FLA and kernel-level gains per optimization round. It reports no per-kernel latency against a non-Cake implementation.

Revision note: the PR body's whole-model runs and kernel gates are at branch head 09933356e. One later branch commit changes the kernels: 45d998d7b5 replaces the WY backward kernels (wy_bwd, wy_bwd_bf16) with tcgen05 kernels on sm_100a and sm_103a (the other 23 stages are unchanged, per its commit message). The PR body has no measurement of that commit. The other branch commit that touches the kernel sources, 68a77b4d30, only moves them to one source pair per stage: every source blob after it is a blob of the measured tree.

Setup from the PR body: a 27B dense GDN model, GRPO, --linear-attention-backend cake against an FLA control arm on the same node and container image (torch 2.13.0+cu130, fla 0.5.2). B200: one 8x B200 node (sm_100a), TP4. GB300: one 4x GB300 node (sm_103a), TP2 x DP2. Ratios are cake / fla, so a value below 1 means the cake arm takes less time.

Sampled arms. Means over steps >= 1. Times are in seconds.

metric B200 cake B200 fla B200 cake/fla GB300 cake GB300 fla GB300 cake/fla
step time 293.0 310.6 0.943 339.6 345.8 0.982
rollout 197.6 201.7 0.980 212.6 216.2 0.983
train time (ref + actor log-probs + actor train) 81.74 94.22 0.868 114.6 117.5 0.976
ref log-probs 13.21 16.06 0.822 18.09 19.19 0.943
actor log-probs 11.84 14.29 0.829 17.72 18.84 0.940
actor train 55.46 62.43 0.888 74.96 78.04 0.961
weight sync 3.633 3.644 0.997 2.72 2.72 —
actor-train tokens/s, length-normalized 23,403 23,228 1.008 18,600 18,510 1.00
mean sampled response length (tokens) 4882 5471 — 5219 5449 —
train–rollout logprob abs diff 0.01029 0.01055 — 0.01040 0.01048 —
train–rollout KL 0.0003791 0.0003862 — 0.000381 0.000384 —
AIME-2024 (480 samples per eval) 0.5417 → 0.6792 after 10 steps not evaluated — 0.5417 → 0.5729 after 5 steps not evaluated —

Three rollouts were recorded with cake. They were then replayed train-only twice with cake (a, b) and once with fla (f). The cake value is the mean of a and b. Means are over all 3 steps, in seconds; train_time includes the train-phase waits.

metric B200 cake B200 fla B200 cake/fla GB300 cake GB300 fla GB300 cake/fla
train_time 93.42 97.27 0.960 112.8 118.1 0.955
ref log-probs 15.88 17.17 0.925 19.41 20.71 0.937
actor log-probs 12.58 13.63 0.923 17.59 17.82 0.987
actor train 64.29 65.86 0.976 74.09 77.53 0.956

With step 0 excluded (steps 1–2), the cake/fla ratios on B200 are 0.975 for train_time and 0.982 for actor train. On GB300 they are 0.974 and 0.969. On both machines, train_time and actor train are lower with cake than with fla at every replayed step, in both cake replays.

The table below lists kernel-level gains per round. Each gain is the CUPTI GPU time of the GDN chunk kernels (forward, and the forward+backward union) on the best row of the 16-row kernel table, from a same-node interleaved A/B against the previous round's kernels. Round 12 kept no candidate.

round change best-row gain vs previous round (sm_103a / sm_100a)
9 WY pass tiles streamed through per-warp staging fwd +5.65 % / +5.35 %
10 fwd_h split entry path scoped to the 64-row scan fwd +7.49 % / +8.80 %, fwd+bwd +2.71 % / +3.25 %
11 eight-warp WY pullback on a padded-pitch pool, stmatrix.x4 tiles, packed f32x2 epilogue fwd+bwd +3.36 % / +3.32 %
13 forward WY on a padded-pitch pool with trace-time address folding and a line-aligned cp.async fwd +5.1…+5.8 % / +4.9…+5.5 % on the 1 × 4096-token TP4 row (both launch orders); fwd+bwd +1.3 % / +1.1 % on that row

Missing facts

  • det against a non-Cake implementation, single-rank and KCP: no FlashInfer body prints the ratio. None is computed here from absolute times.
  • det single-rank on sm_107a in the same run as a non-Cake implementation: the only such R200 det times (the R200 table of flashinfer-ai/flashinfer#8665) are of the det package before flashinfer-ai/flashinfer#8695.
  • det KCP on sm_107a against an earlier det KCP version: not measured (flashinfer-ai/flashinfer#8538).
  • fast single-rank on sm_107a with BF16 beta: not timed (flashinfer-ai/flashinfer#8665: "R200 was timed with FP32 beta only").
  • fast KCP at CP > 1 on R200 against a non-Cake implementation: not claimed (no FlashQLA sm_107a build; flashinfer-ai/flashinfer#8644).
  • fast KCP with save_mode="light" (--cake-gdn-recompute-state) in eager mode: not in the copied tables. flashinfer-ai/flashinfer#8644 times the full save mode only, and the PR-head tables of flashinfer-ai/flashinfer#8638 are graph mode. Eager light-mode rows exist only for earlier package trees and are not used.
  • Deterministic kernels with FLA's context parallelism: no measurement of the WY backward change 45d998d7b5; no kernel-level latency against FLA or another non-Cake implementation; no sm_107a measurement.
  • B300: none of the pull requests above reports B300 numbers.

The README repeated the docstrings of _runtime/config.py, workspace.py,
loader.py, registry.prebuild and tools/vendor_cake.py, which would drift
from the code. Keep the install, miles contract, layout and adding-an-op
sections, shorten the rest to pointers into those modules, and keep the
Ops and Prebuild tables (test_op_contract reads the Ops table; later
layers append rows) with their surrounding lines unchanged so the layers
merge without conflicts.

The version bump is not needed for this PR.
Install, the miles contract in one paragraph, the Ops and Prebuild tables
(test_op_contract reads the Ops table; later layers append rows), the
layout and the development commands. Everything else is documented in
the runtime modules and tools/vendor_cake.py.
workspace.take read the current device, stream and capture state through
private torch._C functions and kept several identity-keyed fast paths
(_direct, _bound, _devices, a running byte total) to shave host time.
The pool is off by default and measured no faster with it on, so the
extra state is not worth its cost: read torch.cuda's public API, compute
held bytes from the pool, and keep the same take / release / stats
behaviour (one buffer per op, role and stream; LRU eviction under the
bound; capture bypass). workspace.py goes from 248 to 130 lines.

config.variant / config.option no longer memoise per configuration; a
dict lookup and a scan over an op's few flags is enough.

Tests: drop the cases that pinned the removed internals and the
micro-benchmark bench_workspace.py; the fake CUDA layer patches the
public torch.cuda functions only. On GB300 (torch 2.14.1+cu130) the
runtime suite passes: 214 passed, 1 skipped (was 223, minus the 9
internal-only cases).

CI: the upstream vendoring check fetches every recorded commit, so it
moves to its own workflow that runs when VENDORED.json or the tool
changes, weekly, and on demand.
main's tree equals cake-stack/00-layout at d67ad25, which this branch already contains, so this records main as merged without changing files.
@Zhichenzzz
Zhichenzzz merged commit 4bb96a7 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