# canopy: linear attention broke tree speculation, and you write the one-solve fix

---

# Part 1: For you

## The idea in one line

Qwen3.5's linear-attention layers made long context cheap but quietly broke tree speculative
decoding, because every draft branch needs its own copy of the recurrent state. Three papers
in five weeks found the same fix: turn the whole draft tree into one triangular solve. None
of them linked code. You write that kernel for an RTX 5090, a GPU that none of them tested.

## Why it's exciting

- **It's a real bottleneck in a real model.** Qwen3.5-9B has 32 layers, and 24 of them are
  Gated DeltaNet (GDN) linear-attention layers. Each GDN head carries a 128 × 128 state. To
  verify a tree of draft tokens, the standard path walks the tree node by node and forks that
  state at every branch.
- **Three papers, one answer.** [SpecLA](https://arxiv.org/abs/2607.16673) (July 18, 2026),
  [Bole](https://arxiv.org/abs/2608.01651) (August 3), and
  [TreeWY](https://arxiv.org/abs/2608.20961) (August 21) each put an ancestor mask into the
  delta rule's triangular factorization. Every node's output then comes from one solve, and
  only the accepted path's state is rebuilt afterward.
- **The payoff is large on paper.** Bole reports per-layer GDN tree verification on
  Qwen3.5-9B going from 174 µs to 50.6 µs on an A100 at batch 1 (3.4×), and from 3,420 µs to
  443 µs on a DGX Spark at batch 16 (7.7×). Transient state memory drops 82-99×.
- **It's contested, and you get to referee.** SpecLA found that its fully parallel kernel was
  *slower* than root-to-leaf replay on 16-node trees on an H100 (0.52× and 0.58×), and only
  broke even at 32 nodes; a chain-decomposed hybrid won instead. TreeWY says wider trees
  become affordable but are "not yet a throughput win." Your benchmark answers where parallel
  stops winning, on hardware nobody has measured.
- **Nobody has released the closed-form kernel.** None of the three papers links code, and a
  search of SGLang found no merged tree-GDN kernel, although Bole says it integrated there.
  [GDN Tree-Scan](https://arxiv.org/abs/2609.23900) (September 20) and
  [ddtree-mlx](https://github.com/humanrouter/ddtree-mlx) (MIT, Metal) released code, but
  both replay or fork state branch by branch: the serial approach, not the closed form.
- **It's small, checkable, and low-level.** The math fits in six lines, and this plan
  already checked it in NumPy to about 1e-15. You write the kernel in Triton, then in CUDA
  with CuTe, for consumer Blackwell (sm_120), which has warp-level `mma.sync` but no
  `tcgen05` or tensor memory, so datacenter-Blackwell kernels don't carry over.

## What the demo looks like

- **A tree-shape picker.** Choose a chain, a balanced top-k tree, a skewed tree, a random
  tree, or paste your own parent array. The page draws the tree and numbers the nodes in the
  order the kernel sees them.
- **The ancestor mask, live.** A T × T grid lights up each node's ancestors. Hold the pointer
  over a node, and its row and its root path highlight together.
- **The solve, animated.** Step through blocked forward substitution, 64 nodes at a time, or
  through the finite Neumann series, one tree level per term. A counter shows how many
  sequential steps each strategy needs for this tree.
- **Recorded curves and a crossover map.** Serial versus parallel on the RTX 5090 for batch
  1-16 and trees of 3-256 nodes, with Bole's numbers marked, and a depth × width heatmap of
  where parallel wins.
- **A memory bar.** Per-node state snapshots next to canopy's pseudo-value matrix, for the
  tree you picked. If you finish the stretch goal, a tokens-per-second chart joins it.

Nothing runs on a GPU at view time. The page computes the mask client-side and plays back
recorded JSON results.

## How it works

1. **Check the math on your Mac.** Write an fp64 NumPy oracle that runs the recurrence along
   every node's root path, and a closed-form solver that must match it.
2. **Write the slow baseline.** Build a serial parent-pointer verifier in Triton that visits
   one node at a time and forks the state, the way SGLang's tree path does.
3. **Write the fast kernel in Triton.** Build the ancestor-masked, decay-weighted K·Kᵀ
   matrix, solve the unit lower-triangular system for pseudo-values, and read every node's
   output from it.
4. **Compare two solves:** blocked forward substitution in shared memory, and Bole's finite
   Neumann series on tensor cores. Then write the commit kernel for the accepted path.
5. **Port the kernel to CUDA and CuTe** for sm_120, or tune the Triton version if that wins.
6. **Benchmark** at Qwen3.5-9B shapes, and find the crossover.
7. **Build the page and ship it** as a static site on `vm.ifkash.dev`.

## Weekend plan

This is a two-weekend project.

| When | What | Done when |
|---|---|---|
| Weekend 1, Saturday morning | On the Mac: write the tree generators, the fp64 oracle, and the closed form in NumPy | The closed form matches the oracle within 1e-10 on 1,000 random trees |
| Weekend 1, Saturday afternoon | Start the pod; record the GPU's properties; write the serial Triton baseline | The baseline matches the oracle, and FLA matches it on chains |
| Weekend 1, Sunday | Triton tree kernel with forward substitution, then the Neumann variant | Both pass parity on every tree shape up to 256 nodes |
| Weekend 2, Saturday | Commit kernel; CUDA and CuTe port (or tuned Triton); benchmark sweep runs by itself | The commit round-trips over 100 rounds; the sweep finishes |
| Weekend 2, Sunday | Crossover analysis; optional H100 run; demo page, deploy, and write-up | The link works |

## Cost

- **GPU:** 1× RTX 5090 on RunPod, about $0.69 per hour, for about 30 hours: about $21
  (estimates).
- **Optional H100:** about 5 hours on an H100 PCIe at about $1.99 per hour (estimate), so
  about $10, to put your kernel next to Bole's A100 numbers on a datacenter GPU.
- **Total:** about $31 (estimate). The plan sets a hard stop at $40.
- **Free:** your M4 Pro runs the NumPy and PyTorch CPU reference, the web build, and the
  page tests. The page is static on your VM and keeps working after the pod is gone.

## What you have at the end

- Open-source Triton and CUDA kernels that verify a whole draft tree through a GDN layer in
  one solve, and a commit kernel that rebuilds only the accepted state.
- An fp64 oracle and a serial baseline that anyone can use to check a tree-GDN kernel.
- Tree-GDN verification numbers on an RTX 5090, which no paper reports, next to Bole's A100
  and DGX Spark numbers.
- A crossover map that answers where parallel verification stops winning, and a solve
  comparison: forward substitution versus the Neumann series.
- A live page where anyone can draw a tree and watch its mask and its solve.
- A blog post: *"Linear attention broke tree speculation. Here's the one-solve fix, on a
  gaming GPU."*

## What might go wrong

| Problem | What to do |
|---|---|
| Parallel loses on small trees | That's a finding, not a failure. SpecLA saw it on an H100. Report the crossover, and ship a dispatcher that picks the serial kernel below it. |
| Decay ratios overflow or underflow | Never form P_j⁻¹ by itself. Keep cumulative gates in log space, and compute every ratio as exp(ℓ_i − ℓ_j), which is at most 1 on the mask. The NumPy check shows the naive fp32 form producing non-finite values on a 64-node chain. |
| The closed form isn't bitwise equal to the serial path | GDN Tree-Scan rejected whole-tree algebra for this reason. Test against the fp64 oracle with stated tolerances, and measure whether greedy acceptance decisions ever differ from the serial path's. |
| Triton won't compile or runs slowly on sm_120 | Pin a Triton version that targets sm_120, and record it. If `tl.dot` underperforms, try other tile shapes, or move straight to the CUDA version. |
| Shared memory is too small for a 256-node tree | Split the work: one kernel builds and inverts the triangular block per head, and a second applies it per value tile. |
| Consumer-GPU clocks drift | Don't assume locked clocks. Use CUDA events, warmup, and the median of many runs, record clocks and temperature, and interleave the kernels you compare. |
| An RTX 5090 isn't available | Use an RTX PRO 6000, which is also sm_120, at a higher hourly price. Record the change. |
| The stretch goal eats the second weekend | It's optional. The kernels, the benchmark, and the page are the project. |
| Bole's code appears in a release | Good: benchmark against it. canopy stays the open, tested sm_120 version with a public oracle. |
| Paper licenses | Bole is CC BY-NC-SA 4.0: cite it and link it, but don't copy its figures or text. TreeWY is CC BY 4.0. |

## Other ideas the research turned up

- **Bitflip.** Do Blackwell FP4 and FP8 tensor cores slow down or draw more power on random
  data? [Horace He](https://www.thonking.ai/p/strangely-matrix-multiplications) measured an
  A100 at 295 TFLOPS on zeros and 257 on `randn`, and
  [Gregersen et al.](https://arxiv.org/abs/2409.18324) found almost 40% power variation from
  data alone. It lost because it's hardware characterization more than an ML kernel, RunPod
  might not expose reliable power telemetry, and the K-reordering fix is untested.
- **Shrinkage.** Uniform 4-bit training on the RTX 4090's INT4 tensor cores, after
  [arXiv 2606.20381](https://arxiv.org/abs/2606.20381): E2M1 rounding biases values toward
  zero, and its UFP4 recipe removes about 20-22% of the FP4-versus-BF16 loss gap. It lost
  because at a weekend-sized model, the loss difference might sit inside seed noise.
- **MonoMoE on sm_120.** Port the weight-major MoE decode megakernel from
  [arXiv 2609.04244](https://arxiv.org/abs/2609.04244) (1.54× over vLLM's Triton grouped GEMM
  on an H200; code in FlashInfer, datacenter GPUs only). It lost because fusing routing, two
  FP8 projections, and the reduction is too much for two weekends.
- **FlashPairformer on Metal.** Port Anthropic's open-source
  [triangle attention and multiplication kernels](https://github.com/anthropics/uplifting-biomolecular-modeling)
  (Apache-2.0) to Apple Silicon. It lost because the M4 Pro has no matrix units, which caps
  the speedup, and wiring the kernels into Boltz-2 could eat the weekend.
- **Fast Gauss sums.** [arXiv 2609.04910](https://arxiv.org/abs/2609.04910) computes Gaussian
  kernel sums with flash attention and beats PyKeOps for D > 8, but loses to KeOps at
  D ≤ 8. It lost on a quieter hook and a modest expected margin.

## Reading, if you want it

- [Bole](https://arxiv.org/abs/2608.01651): the tree closed form, the finite Neumann solve,
  and the A100 and DGX Spark numbers.
- [TreeWY](https://arxiv.org/abs/2608.20961) (Ghantasala, 2026): the same idea as a tree WY
  transform, with commit by reconstruction.
- [SpecLA](https://arxiv.org/abs/2607.16673) (Wang et al., 2026): tree masks in the UT
  transform, and the result that parallel can lose.
- [GDN Tree-Scan](https://arxiv.org/abs/2609.23900) (Ma, 2026) and
  [STree](https://arxiv.org/abs/2505.14969) (Wu et al., NeurIPS 2025): scan and replay in
  vLLM, and tree speculation for Mamba2 hybrids.
- [Gated Delta Networks](https://arxiv.org/abs/2412.06464) (Yang et al., ICLR 2025) and
  [flash-linear-attention](https://github.com/fla-org/flash-linear-attention) (MIT): the rule,
  and its chunked Triton kernels with log-space gates and the WY solve.
- [SpecInfer](https://arxiv.org/abs/2305.09781) (Miao et al., 2023): token-tree verification
  for transformers.
- [DFlash](https://arxiv.org/abs/2602.06036) (Chen, Liang, and Liu, 2026) and
  [DDTree](https://arxiv.org/abs/2604.12989) (Ringel and Romano, 2026): a block-diffusion
  drafter and the draft trees built from it.
- [Triton documentation](https://triton-lang.org/main/index.html) and Colfax's
  [NVFP4 GEMM on RTX PRO Blackwell (sm_12x)](https://research.colfax-intl.com/cutlass-tutorial-nvfp4-blockscaled-gemm-on-nvidia-rtx-pro-blackwell-gpus-sm12x/)
  tutorial: the kernel language, and what sm_120 does and doesn't have.

---

# Part 2: For the coding agent

## Mission

Build `canopy`, an open-source kernel library and benchmark for verifying a tree of
speculative draft tokens through Gated DeltaNet layers in one triangular solve, on an RTX
5090 (sm_120):

- **Oracle and baseline:** An fp64 reference that runs the gated delta rule along every
  node's root path, and a serial Triton parent-pointer verifier that forks state per node.
- **Tree kernel:** A Triton kernel, then a CUDA and CuTe version (or a tuned Triton version
  if it wins), that takes a draft tree as a parent array of up to 256 topologically ordered
  nodes and outputs every node's o. It has two solve strategies: blocked forward substitution
  and the finite Neumann series.
- **Commit kernel:** Rebuilds only the accepted path's state from the pseudo-values.
- **Benchmarks:** Qwen3.5-9B GDN shapes, batch 1-16, trees of 3-256 nodes, with the
  crossover between serial and parallel.
- **Stretch:** Qwen3.5-9B end to end with a DFlash drafter, in tokens per second.
- **Demo page:** A static site with a tree-shape picker, the ancestor mask, the solve
  schedule, and recorded latency and memory curves.

None of the closed-form papers released code, so you implement everything from the math in
this plan. Record every choice the papers leave open in `NOTES.md`.

Work through milestones M0-M8 in order. Each milestone has acceptance criteria. Don't start a
milestone until the previous one passes, except where a milestone says it runs in parallel.
After each milestone, commit your work and write a short entry in `NOTES.md` with the results
and numbers.

## Hard constraints

- **Secrets:** Read `RUNPOD_API_KEY` and `HF_TOKEN` from environment variables only. Never
  write a secret into any file, log, commit, or echoed command. Commit a `.env.example` that
  has placeholder values only, and add `.env` to `.gitignore`.
- **Budget:** The hard cap is $40 of RunPod spend. Every pod runs a watchdog that stops the pod
  after `MAX_POD_HOURS` hours. The default is `6`. Check `infra/pod.py cost` before you start
  each milestone. At $32, stop and ask the user before you spend more.
- **Pod cleanup:** When a pod isn't running a job, stop it. At the end of the project,
  terminate every pod that you created and report the total spend.
- **Compute split:** The Mac runs the tree generators, the NumPy oracle, the CPU PyTorch
  reference, the web build, and the page tests. Only GPU kernels and benchmarks run on the
  pod.
- **Oracle parity:** Test every kernel against the fp64 oracle on the same inputs:
  - fp32 closed form in NumPy or PyTorch: maximum absolute error 1e-5 or less.
  - fp32 kernels: maximum absolute error 1e-4 or less and relative Frobenius error 1e-5 or
    less, per output tensor.
  - bf16 or TF32 tensor-core kernels: relative Frobenius error 1e-2 or less against the
    oracle on the same bf16-rounded inputs. Report the actual error.
  - These are estimates from the NumPy check (about 3e-7 at small head sizes). Tighten them
    if the data allows; never loosen them without recording why.
- **Log-space gates:** Store cumulative gates as ℓ_i = Σ log α over the root path. Compute
  every decay ratio as exp(ℓ_i − ℓ_j) with j an ancestor of i or i itself. Never materialize
  P⁻¹, and never materialize a T × T diagonal matrix.
- **No unreleased code:** Don't wait for, request, or depend on code from Bole, TreeWY, SpecLA,
  or the TreeWY vLLM fork. Implement from this plan and the papers' math.
- **Licenses:** Use only permissively licensed code: MIT, Apache-2.0, or BSD. Keep the notice
  in any file you adapt from flash-linear-attention (MIT) or ddtree-mlx (MIT). Check each
  dependency's license in M0 and record it. Cite Bole (CC BY-NC-SA 4.0), but don't copy its
  figures or text. State that canopy isn't affiliated with any paper's authors.
- **Reproducible benchmarks:** Time with CUDA events. Run 25 warmup iterations, then report
  the median of 200 timed iterations, with the 10th and 90th percentiles. Flush L2 between
  iterations for cold numbers, and report warm numbers separately. Don't assume locked clocks;
  record the SM clock, memory clock, temperature, and power before and after each sweep with
  `nvidia-smi`, and interleave the kernels you compare. Log the GPU name, driver, CUDA, PyTorch,
  and Triton versions with every result file.
- **Outward actions:** Ask the user before you do any of the following: create the GitHub
  repository, push to the Hugging Face Hub, deploy to a VM, change DNS, or post anything
  publicly. Use the `kashifulhaque` GitHub account (`gh auth switch -u kashifulhaque`).
- **Shared VM:** `vm.ifkash.dev` runs other production apps behind one shared Caddy, which owns
  ports 80 and 443. Its config lives at `~/docs/caddy`.
  - The canopy container lives in `~/docs/canopy` and must not publish any ports. It joins the
    external Docker network `edge` with a stable alias. Check existing aliases with
    `docker network inspect edge` first; use `canopy` if it's free.
  - Add the vhost only by appending to `~/docs/caddy/Caddyfile` with `>>`. Never rewrite,
    rename, or replace that file. It's a single-file bind mount, and a rewrite orphans the
    inode, so `caddy reload` then reports "config is unchanged" while serving the old config.
  - Docker on the VM has no BuildKit for plain `docker build`. Don't use `COPY --chmod`,
    Dockerfile heredocs, or `RUN --mount`. Multi-stage builds and `COPY --from` work.
  - Don't stop, restart, reconfigure, or remove any other container, network, or vhost.

## Science background

### The gated delta rule

Use the convention of flash-linear-attention and Bole. Each head's state S has shape
d_k × d_v. For token t with query q_t, key k_t, value v_t, decay gate α_t ∈ (0, 1), and write
strength β_t ∈ (0, 1):

```
S~_t = α_t S_{t-1}
u_t  = β_t (v_t − S~_tᵀ k_t)
S_t  = S~_t + k_t u_tᵀ            = α_t (I − β_t k_t k_tᵀ) S_{t-1} + β_t k_t v_tᵀ
o_t  = S_tᵀ q_t
```

TreeWY writes the transpose, a d_v × d_k state with S_t = α_t S_{t−1}(I − β_t k_t k_tᵀ) +
β_t v_t k_tᵀ and o_t = S_t q_t; it's the same rule. FLA's naive reference matches the
preceding block. In FLA, g = log α (so g ≤ 0), q and k can be L2-normalized in the kernel,
and the default scale on q is 1/√d_k.

### Qwen3.5-9B shapes

The model's `config.json` sets 32 layers with `full_attention_interval` 4 (three GDN layers,
then one full-attention layer, so 24 GDN layers), `hidden_size` 4096,
`linear_num_key_heads` 16, `linear_num_value_heads` 32, `linear_key_head_dim` 128,
`linear_value_head_dim` 128, and `linear_conv_kernel_dim` 4. Weights are BF16, and the model
card lists Apache-2.0. So each GDN layer holds 32 states of 128 × 128 per sequence, and each
key head serves two value heads (FLA's grouped value attention, HV = 32 and H = 16). In fp32,
the states take 32 × 128 × 128 × 4 bytes = 2 MiB per layer per sequence (computed).

### The tree closed form

Number the T draft nodes 0 to T − 1 so that every parent precedes its child, with
parent(0) = −1. Node 0 is the root, usually the last sampled token, and S₀ is the committed
state before it. Let A(i) be the strict ancestors of i, and let path(i) = A(i) ∪ {i}. Let d
be the tree's depth: the largest |A(i)|, so a single node has d = 0 and a chain of T nodes
has d = T − 1.

Running the recurrence along each root path gives the following, where every sum and mask
follows the tree:

```
ℓ_i   = log α_i + ℓ_parent(i)          (ℓ of the root's parent = 0)
P_i   = exp(ℓ_i)                       cumulative decay over path(i)
Γ_ij  = exp(ℓ_i − ℓ_j)  if j ∈ path(i), else 0          (0 < Γ_ij ≤ 1)

G_ij  = β_i Γ_ij (k_iᵀ k_j)   if j ∈ A(i), else 0       strictly lower triangular
R_i   = β_i (v_i − P_i S₀ᵀ k_i)
(I + G) U = R                                            U: T × d_v pseudo-values
C_ij  = Γ_ij (q_iᵀ k_j)       if j ∈ path(i), else 0
O     = diag(P) Q S₀ + C U                               every node's output
S_a   = P_a S₀ + Σ_{j ∈ path(a)} Γ_aj k_j u_jᵀ           commit to accepted node a
```

In matrix form, O = D_P Q S₀ + C (I + G)⁻¹ D_β (V − D_P K S₀), which is Bole's Theorem 1.
TreeWY writes the same system as (I + diag(β) G′) Ṽ = R with G′ = G without the β, g = P,
and Ṽ = U. Here's why it works: every node's state is S_i = P_i S₀ + Σ_{j ∈ path(i)} Γ_ij k_j
u_jᵀ, so substituting the parent's state into u_i gives one linear equation per node, and
those equations only couple a node to its ancestors.

The facts that matter for the kernel are as follows:

- **Any topological order works.** DFS pre-order (TreeWY's choice) keeps subtrees contiguous.
  Level order makes each level's diagonal block zero, because nodes of equal depth are never
  ancestors of each other.
- **Nilpotency.** G^(d+1) = 0 and G^d is generally non-zero, so
  (I + G)⁻¹ = Σ_{m=0..d} (−G)^m: the Neumann series ends after d matrix products. The NumPy
  check confirmed both on every test tree.
- **Squaring.** (I − G)(I + G²)(I + G⁴)… with k factors equals the first 2^k terms, so
  ⌈log₂(d + 1)⌉ factors are enough. This is this plan's derivation, checked in NumPy.
- **Level-synchronous solve.** For each depth level L in order, U_L = R_L − G_{L,:} U: d + 1
  sequential steps of small matrix products.
- **Depth, not size, sets the sequential cost.** A full top-4 tree reaches 341 nodes at depth
  4 (computed), so Bole's 256-node trees of depth up to 8 must be pruned (an inference). A
  256-node chain has d = 255, so expect the parallel kernel to lose on deep, narrow trees.
- **Memory.** Per-node snapshots cost T × 2 MiB per layer per sequence in fp32, 256 MiB at
  T = 128; the pseudo-values cost 2 MiB at T = 128 (both computed). Bole reports 14.1 GB of
  snapshots versus 151 MB of factors at T = 100 on Qwen3.5-27B.

### Verification of the formulas

This plan verified the closed form in `canopy_check.py` (float64, d_k = 16 or 32, d_v = 12 or
16, L2-normalized q and k, β ∈ (0.05, 0.95), α ∈ (0.5, 0.999)). Five solvers (dense solve,
forward substitution, Neumann, squaring, and level-synchronous) and a 16-node blocked forward
substitution all matched the oracle to 8e-16 or better on six trees: a random 40-node tree in
BFS-like order, a random 40-node tree in DFS pre-order, a 12-node chain, top-4 trees of 128
and 256 nodes, and a 64-node skewed tree of depth 48. The commit formula matched the oracle's
state to 7e-16 or better. The fp32 closed form stayed within 3e-7 of the fp64 oracle. On a
64-node chain with α = 0.05, naive fp32 ratios P_i/P_j produced 630 non-finite values, and the
log-space form stayed within 1.4e-7. No formula needed correction.

### What the papers report

- **Bole** (arXiv 2608.01651, August 3, 2026, CC BY-NC-SA 4.0) solves the closed form with a
  finite Neumann polynomial on tensor cores and applies decay ratios as row-wise scalings. It
  runs H × ⌈d_v / b_v⌉ CTAs per request, each owning one value tile across all T nodes. Its
  serial baseline is SGLang's native tree path, with a state snapshot per node. Table VII,
  per-layer GDN tree verification on Qwen3.5-9B:

  | GPU | Batch | Serial (µs) | Bole (µs) | Speedup | SM occupancy |
  |---|---|---|---|---|---|
  | A100 | 1 | 174 | 50.6 | 3.4× | 1.9% → 15.0% |
  | A100 | 16 | 822 | 110 | 7.5× | 11.4% → 41.4% |
  | GB10 (DGX Spark) | 1 | 333 | 66.6 | 5.0× | 6.6% → 43.1% |
  | GB10 (DGX Spark) | 16 | 3,420 | 443 | 7.7× | 16.4% → 58.9% |

  Trees have depth up to 8 with top-4 branching, at 8, 32, 64, 100, 128, 200, and 256 nodes;
  memory is measured at 128. End to end: up to 4.72× over autoregressive decoding and 2.03×
  over the strongest tree baseline. Colfax lists DGX Spark as sm_121, so the GB10 rows are
  the closest published comparison for sm_120, but the GB10 has only 273 GB/s of LPDDR5x.
- **TreeWY** (arXiv 2608.20961, CC BY 4.0) solves the same system by forward substitution in
  DFS pre-order and matches the recurrence to about 1e-15 in fp64 and 1e-7 in fp32. Only
  chains run fused; general trees run piecewise, outside CUDA graphs. It tests Qwen3.5-35B-A3B
  and -397B-A17B on B200s in a private vLLM fork, with trees up to (3, 3, 3), 39 nodes.
- **SpecLA** (arXiv 2607.16673) uses A_tree = M_tree ⊙ (K Kᵀ), then `solve_tril`. On an H100
  with GDN-1.3B and 16- or 32-node trees, its parallel kernel reaches 0.52-3.47× over
  root-to-leaf replay, and a chain-decomposed hybrid reaches 1.80-7.11×.
- **GDN Tree-Scan** (arXiv 2609.23900) uses branch-local scan and replay in vLLM, with a public
  artifact. It declined whole-tree algebra because its summation order and bf16 boundaries
  differ from the native sequential path.

## Design decisions

- **Triton first.** Triton reaches a working tensor-core kernel with far less code than
  CUDA, it's what FLA and Bole use, and it gives you a baseline to beat. Port to CUDA and CuTe
  only after the Triton kernel passes parity and has a profile.
- **Two solve strategies, one interface.** Both take the same inputs and return U:
  - *Blocked forward substitution:* 64-node blocks; for block b, solve
    (I + G_bb) U_b = R_b − Σ_{c<b} G_bc U_c. Diagonal blocks use row-by-row substitution in
    shared memory; off-diagonal updates are tensor-core matrix products.
  - *Finite Neumann series (Bole):* Z ← −G Z and U ← U + Z for d rounds on tensor cores, plus
    the squaring form. Its cost grows with depth; find where it loses to forward substitution.
- **Tree order and tiling.** The host sorts each tree into DFS pre-order, pads T to a
  multiple of 16, and builds the ancestor mask as a T × T bitmask (256 × 256 bits is 8 KiB,
  computed). The kernel reads the mask, the parent array, and the node depths. Record whether
  level order tiles better than DFS order; both are valid.
- **Batch and heads to CTAs.** Follow Bole: the grid is (batch, value head, value tile). With
  32 value heads and b_v = 32, batch 1 launches 128 CTAs; pick b_v from the SM count you
  record in M0. Each CTA recomputes G for its head. The split alternative, FLA's
  `chunk_scaled_dot_kkt` plus `solve_tril` pattern, builds (I + G)⁻¹ once per (batch, value
  head) in global memory (256 KiB per head at T = 256 in fp32, computed) and applies it per
  value tile. Implement both, and keep the faster one per shape.
- **Memory layout.** q and k are `[B, T, H, 128]` and v is `[B, T, HV, 128]`, in bf16, as in
  FLA. β and g are `[B, T, HV]`, and the kernel computes ℓ in fp32. S₀ is
  `[B, HV, 128, 128]` in fp32. The tree is `parent: int32[B, T]` plus a per-sequence length
  for ragged batches. Accumulate every matrix product in fp32.
- **What the commit kernel stores.** Verification writes U (`[B, T, HV, 128]`, fp32) and ℓ
  (`[B, T, HV]`, fp32); K is already in the layer's activations. Commit takes the accepted
  path as node indices and computes S_new = P_a S₀ + (D_a K_path)ᵀ U_path, one
  (128 × m) · (m × 128) product per head for a path of m nodes. It rebuilds nothing else.
- **Serial baseline.** One Triton program per (batch, value head, value tile) walks the nodes
  in order, reads the parent's state slice from a per-node snapshot buffer, applies one step,
  and writes the node's snapshot and output: SGLang's snapshot path, which Bole compares to.
- **Precision tiers.** Tier A is fp32 everywhere (parity reference on GPU). Tier B takes bf16
  inputs, uses TF32 or bf16 tensor cores, and accumulates in fp32. Report both, and benchmark
  tier B.
- **sm_120 specifics.** Consumer Blackwell uses warp-level `mma.sync` with register operands,
  like sm_80, and has TMA but no `tcgen05` or TMEM. Write the CUDA version with CuTe's SM80
  MMA atoms plus TMA loads. Don't use WGMMA or `tcgen05` code paths.
- **Where parallel stops winning.** Treat this as the main question. Sweep chains, balanced
  top-k trees, skewed trees, Bole's shapes, and SpecLA's shapes (top-2 and top-4, depth 4, 12,
  or 24, 16 or 32 nodes). Report the crossover as a function of T and depth, and ship a
  dispatcher that picks the faster kernel from a small table.

## Tech stack

Pin every version. The stack is as follows:

- **Python:** Python 3.12 with `uv`, `numpy`, `pytest`, and `hypothesis`. On the pod,
  `torch` built for CUDA 12.8 or later, whose `get_device_capability()` returns `(12, 0)`,
  with its bundled Triton; on the Mac, the CPU build. `flash-linear-attention` (MIT) serves
  as the chain reference only.
- **CUDA:** CUDA 12.8 or later with CUTLASS and CuTe headers, built as a PyTorch extension.
- **Benchmarks:** CUDA events or `triton.testing.do_bench`, results as JSONL under
  `results/`, and plots with `matplotlib`. No hosted tracker.
- **Stretch:** `transformers` for Qwen3.5-9B, and `z-lab/Qwen3.5-9B-DFlash` (Apache-2.0, a
  1B-parameter drafter) as the drafter.
- **Web:** TypeScript and `vite`, with hand-written SVG charts and no framework. Test with
  Playwright on Chromium.
- **Pods:** the `runpod` Python SDK, with `runpodctl` inside pods.

## Repository layout

Create the following layout:

```
canopy/
  pyproject.toml  .env.example  .gitignore  README.md  NOTES.md  LICENSE (MIT)
  infra/pod.py  infra/watchdog.sh  infra/bootstrap.sh
  canopy/
    trees.py             # parent arrays: chain, top-k, skewed, random, Bole, SpecLA shapes
    order.py             # DFS pre-order, level order, depths, ancestor bitmask
    oracle.py            # fp64 per-path recurrence; fp64/fp32 closed form (NumPy)
    oracle_torch.py      # the same in PyTorch, CPU or CUDA, for kernel tests
    serial_triton.py     # M2: parent-pointer baseline with per-node snapshots
    tree_triton.py       # M3-M4: fused and split kernels, fwd-sub and Neumann solves
    commit_triton.py     # M5: accepted-path state rebuild
    csrc/                # M6: CUDA + CuTe kernels, PyTorch extension bindings
    dispatch.py          # picks serial or parallel per (B, T, depth) from a table
  tests/                 # parity tests per kernel, tolerances from the hard constraints
  bench/
    shapes.py  sweep.py  crossover.py  memory.py  env.py   # env.py logs GPU, clocks, versions
  e2e/                   # stretch: Qwen3.5-9B + DFlash + DDTree-style tree builder
  web/
    src/tree.ts  src/mask.ts  src/solve.ts  src/charts.ts  src/main.ts  index.html
    public/results/      # recorded JSON from bench/
    tests/               # Playwright smoke test
  deploy/Dockerfile  deploy/compose.yml  deploy/nginx.conf
  results/  post/draft.md
```

## M0: Infrastructure

**Tasks:**

1. Write `infra/pod.py`. It uses the `runpod` SDK and reads `RUNPOD_API_KEY` from the
   environment. It supports the following subcommands:
   - `create`: Creates a pod. The GPU preference order is RTX 5090, then RTX PRO 6000.
     Resolve GPU type IDs at run time. Use an official RunPod image with CUDA 12.8 or later,
     a 50 GB volume at `/workspace`, and SSH. `create --gpu h100` makes the optional H100
     PCIe pod.
   - `status`, `stop`, `terminate`, and `ssh-info`.
   - `cost`: Prints the uptime and spend for every pod whose name has the prefix `canopy-`.
2. Write `infra/watchdog.sh`. It sleeps for `MAX_POD_HOURS`, then runs
   `runpodctl stop pod $RUNPOD_POD_ID`.
3. Write `infra/bootstrap.sh`. It clones the repo, runs `uv sync`, and starts the watchdog.
4. Write `bench/env.py`. It records the GPU name, compute capability, SM count, shared memory
   per block and per SM, L2 size, software versions, clocks, and power limit.
5. Check and record the license of every dependency, and read the Qwen3.5 GDN layer in
   `transformers` to record how it computes β and the log gate from its projections.

**Acceptance criteria:**

- A pod comes up, `bootstrap.sh` finishes without errors, and `bench/env.py` reports compute
  capability 12.0.
- A test run with `MAX_POD_HOURS=0.05` stops the pod within 5 minutes.
- `NOTES.md` records the device properties, the license table, and the gate formula.

## M1: Oracle and tree generators

Do this milestone on the Mac.

**Tasks:**

- Write `canopy/trees.py` and `canopy/order.py`: every tree family in the design decisions,
  DFS pre-order and level order, depths, and the ancestor bitmask.
- Write `canopy/oracle.py` following this spec, which is the logic of the verified
  `canopy_check.py`:

```python
def oracle(parent, q, k, v, beta, alpha, S0):          # fp64, parent[i] < i
    S, O = [None] * T, zeros(T, dv)
    for i in range(T):
        Sp = S0 if parent[i] < 0 else S[parent[i]]
        St = alpha[i] * Sp
        u = beta[i] * (v[i] - St.T @ k[i])
        S[i] = St + outer(k[i], u)
        O[i] = S[i].T @ q[i]
    return O, S

def closed_form(parent, q, k, v, beta, alpha, S0):
    Mm = strict_ancestor_mask(parent); Mp = Mm | eye(T)
    l = cumsum_over_path(log(alpha), parent)          # l[i] = log a[i] + l[parent[i]]
    Gam = where(Mp, exp(where(Mp, l[:, None] - l[None, :], 0)), 0)
    G = beta[:, None] * Gam * (k @ k.T) * Mm
    C = Gam * (q @ k.T) * Mp
    R = beta[:, None] * (v - exp(l)[:, None] * (k @ S0))
    U = forward_substitute(G, R)                      # (I + G) U = R
    return exp(l)[:, None] * (q @ S0) + C @ U, U

def commit(parent, a, k, alpha, U, S0):
    l = cumsum_over_path(log(alpha), parent)
    S = exp(l[a]) * S0
    for j in path(a): S += exp(l[a] - l[j]) * outer(k[j], U[j])
    return S
```

- Add the Neumann, squaring, level-synchronous, and 64-node blocked solvers, and write
  `canopy/oracle_torch.py` with the same functions in PyTorch.
- Write property tests with `hypothesis`: random trees of 1-256 nodes, with random gates
  including α down to 0.01.

**Acceptance criteria:**

- The fp64 closed form matches the oracle within 1e-10 maximum absolute error on 1,000 random
  trees with d_k = d_v = 128, for every solver.
- G^(d+1) is exactly zero on every tree, and the commit matches the oracle's state for every
  leaf within 1e-10.
- The fp32 closed form matches the fp64 oracle within 1e-5. Record the largest error and the
  tree that produced it.
- The PyTorch and NumPy versions agree within 1e-12 in fp64.

## M2: Serial baseline

**Tasks:**

- Write `canopy/serial_triton.py`, the parent-pointer baseline from the design decisions, in
  fp32 with bf16 inputs allowed.
- On chains, compare it with FLA's `fused_recurrent_gated_delta_rule` and
  `chunk_gated_delta_rule` (log-space g, `[B, T, H, K]` q and k, `[B, T, HV, V]` v).
- Record its peak memory for T = 8, 32, 64, 128, and 256 at batch 1 and 16.

**Acceptance criteria:**

- The baseline passes the fp32 tolerances against the oracle on every tree family, batch 1-16,
  with ragged tree sizes in one batch.
- On chains, the baseline and FLA agree within the fp32 tolerances.
- `NOTES.md` has a first latency table at Bole's shape (depth ≤ 8, top-4, T = 128) for
  batch 1 and 16, next to Bole's serial numbers, labeled as a different GPU and baseline.

## M3: Triton tree kernel

**Tasks:**

- Write the fused kernel in `canopy/tree_triton.py`: grid (batch, value head, value tile);
  load K, Q, and the gates for all T nodes; build G and C in fp32 with the ancestor mask and
  log-space ratios; solve by blocked forward substitution with 64-node blocks; write O, U,
  and ℓ.
- Support T up to 256 with padding, ragged batches, and both DFS and level order.
- Add tier A (fp32 `tl.dot` with `input_precision="ieee"`, or `allow_tf32=False` in older
  Triton) and tier B (bf16 inputs, fp32 accumulate).

**Acceptance criteria:**

- Tier A passes the fp32 tolerances on every tree family, T = 3-256, batch 1-16.
- Tier B passes the bf16 tolerance, and `NOTES.md` records the actual errors.
- The kernel writes nothing proportional to T × d_k × d_v; `torch.cuda.max_memory_allocated`
  confirms it.

## M4: Solve strategies

**Tasks:**

- Add the Neumann solve (Z ← −G Z, U ← U + Z, for d rounds, on tensor cores) and the squaring
  form, behind the same interface.
- Add the split design: kernel 1 builds G and its inverse per (batch, value head); kernel 2
  applies it per value tile.
- Autotune tile sizes (b_v, key chunk, `num_warps`, `num_stages`) per (T, batch).
- Optional: add a chain-decomposed hybrid in SpecLA's style (heavy-light decomposition, each
  chain serial, chains in parallel) as a third contender.

**Acceptance criteria:**

- Every variant passes the tolerances of its tier.
- A table in `NOTES.md` compares forward substitution, Neumann, and squaring by latency, and
  by bf16 error, across depths 1-255, at batch 1 and 16.
- `NOTES.md` states which strategy wins at which depth, and whether the split design beats the
  fused one.

## M5: Commit kernel and the full round

**Tasks:**

- Write `canopy/commit_triton.py`: given the accepted path per sequence, rebuild S_new from
  S₀, ℓ, K, and U as one small product per head.
- Write a round test: verify a tree, pick an accepted path, commit, then feed S_new into the
  next round's tree, for 100 rounds with random trees and random accepted paths.
- Compare against the oracle, and against the serial baseline's snapshot of the accepted node.

**Acceptance criteria:**

- One commit matches the oracle's state within the fp32 tolerances.
- After 100 rounds, tier A's state stays within 1e-3 relative Frobenius error of the fp64
  oracle's. Record the error growth curve.
- With tier B, record over 100 rounds how often the argmax of a random projection of o
  differs from the serial baseline's. This stands in for Tree-Scan's acceptance concern.

## M6: CUDA and CuTe kernel

**Tasks:**

- Profile the fastest Triton kernel with Nsight Compute on the 5090: occupancy, shared-memory
  throughput, and the share of time spent in the solve.
- Write the CUDA version in `canopy/csrc/` with CuTe SM80 MMA atoms (`mma.sync`) and TMA
  loads: load the tile, build G in registers and shared memory, run the solve, and write O,
  U, and ℓ. Start with the strategy that won in M4.
- If the CUDA version can't beat Triton within one day of work, tune Triton instead, and
  record the decision.
- Optional: run the fastest kernel and the serial baseline on an H100 PCIe pod, 5 hours at most.

**Acceptance criteria:**

- The CUDA version passes the same parity tests as the Triton version.
- `NOTES.md` has an Nsight comparison of the Triton and CUDA kernels, with SM occupancy next to
  Bole's occupancy numbers, labeled as a different GPU.
- If you ran the H100, `NOTES.md` has its numbers next to Bole's A100 numbers.

## M7: Benchmarks, the crossover, and the stretch goal

Run the sweep in the background while M8's page work starts on the Mac.

**Tasks:**

- Write `bench/sweep.py`: every kernel (serial, fused and split Triton with each solve, CUDA)
  over batch 1, 2, 4, 8, and 16, T = 3, 8, 16, 32, 64, 100, 128, 200, and 256, and chains,
  balanced top-2 and top-4 trees, skewed trees, Bole's shapes, and SpecLA's shapes. Follow the
  benchmark rules in the hard constraints.
- Write `bench/memory.py`: peak transient memory for serial snapshots versus U and ℓ, per
  layer and for all 24 GDN layers.
- Write `bench/crossover.py`: the depth × width map of parallel speedup over serial, and a
  dispatch table for `canopy/dispatch.py`.
- Export `web/public/results/*.json`: latency curves, the crossover map, memory, and the
  environment record.
- **Stretch (optional):** In `e2e/`, run Qwen3.5-9B in BF16 with `z-lab/Qwen3.5-9B-DFlash`,
  a DDTree-style tree builder, a tree attention mask for the 8 full-attention layers, and
  canopy for the 24 GDN layers. The short causal convolution (kernel size 4) before each GDN
  layer must also follow each node's root path, and commit must update the convolution state
  with the accepted tokens. Measure tokens per second against autoregressive decoding and
  against a chain-only DFlash run.

**Acceptance criteria:**

- The sweep finishes with a median, 10th, and 90th percentile for every cell, and the
  environment record for each run.
- `NOTES.md` states the crossover: the depth and T at which parallel stops winning, at batch 1
  and 16, and how that compares with SpecLA's finding.
- The measured memory reduction at T = 128 is in `NOTES.md` next to Bole's 82-99× range.
- Stretch: greedy decoding with canopy produces the same tokens as autoregressive greedy
  decoding on 20 prompts, or `NOTES.md` explains each difference. Tokens per second are
  recorded for each setup.

## M8: The demo page, deploy, and write-up

**Tasks:**

- **Picker and tree view:** chain, balanced top-k (k and depth sliders), skewed, random with
  a seed, and a text field for a parent array of up to 256 nodes. Validate that parents
  precede children, offer to reorder, and draw an SVG tree numbered in DFS or level order.
- **Mask and solve views:** the T × T ancestor mask computed client-side, with the mask row
  and root path highlighted for the node under the pointer; a stepper through blocked forward
  substitution and the Neumann series, with the sequential step count for each.
- **Charts:** latency versus T per batch size, the crossover heatmap, and memory bars from the
  recorded JSON, with Bole's numbers marked as different hardware. Each chart has a text
  table alternative, and the panels stack at phone width.
- **Footer:** links to the four papers and the repository, the environment record, and "not
  affiliated with the authors."
- Write `deploy/Dockerfile` (`nginx:alpine` serving `web/dist`, buildable with the legacy
  builder) and `deploy/compose.yml` (no published ports, the external network `edge`, and the
  alias chosen in the hard constraints).
- Ask the user before deploying. The user confirms a domain such as `canopy.ifkash.dev` and
  adds a non-proxied Cloudflare A record.
- After the user approves, copy the build and `deploy/` to `~/docs/canopy`, and run
  `docker compose up -d` there. Then compare the live and host Caddyfiles with
  `docker exec caddy cat /etc/caddy/Caddyfile | diff - ~/docs/caddy/Caddyfile`. If they
  differ, stop and ask the user. Otherwise, back up, append, validate, and reload:

```
cd ~/docs/caddy
cp Caddyfile "Caddyfile.bak-canopy-$(date +%Y%m%d)"
printf '\ncanopy.ifkash.dev {\n\treverse_proxy canopy:80\n}\n' >> Caddyfile
docker exec -i caddy caddy validate --config - --adapter caddyfile < Caddyfile
docker exec -i caddy caddy reload --config - --adapter caddyfile < Caddyfile
```

  Omit any `tls` block: for a non-proxied A record, the default ACME HTTP-01 challenge works.
- Write `README.md`. It explains the problem, the math in brief, how to run the parity tests
  and the benchmark with one command each, the results tables, the crossover, every open
  choice, the licenses, and the credits.
- Write `post/draft.md`, a blog post of 1,200-1,800 words. Structure it as follows:
  - The hook: linear attention made long context cheap and quietly broke tree speculation.
  - The fix in plain words: one equation per node, coupled only to its ancestors, so the
    whole tree is one triangular solve.
  - Two solves: forward substitution and the Neumann series, and why depth sets the cost.
  - The results on the 5090, and where parallel stops winning.
  - Limitations: one GPU family, synthetic trees unless the stretch goal ran, and numerics
    that aren't bitwise equal to the serial path.
- Ask the user before you push anything. After the user approves, create and push the repo on
  the `kashifulhaque` account. The owner lists the project on projects.dotslasha.me after the
  deploy; that isn't your job.
- Terminate every pod, and write the final spend to `NOTES.md`.

**Acceptance criteria:**

- A Playwright smoke test loads the page, picks each tree family, pastes a parent array, steps
  through both solves, and renders every chart.
- The deployed URL loads over HTTPS in Chrome and Safari, and on one phone.
- Every other site on the VM still responds as it did before the deploy.
- No pods are left running, and the total spend is less than $40.

## Final report to the user

When you finish, report in three short sections, with bullets: what was asked, what changed,
and whether it worked. Include the following:

- **What was asked:** the deliverables, and whether you attempted the stretch goal.
- **What changed:** the repository and deployed URLs, if the user approved them, and the
  kernels that exist (serial, Triton fused and split, each solve, CUDA or tuned Triton,
  commit) with the precision tiers each supports.
- **Did it work:**
  - Parity: the largest error per kernel and tier against the fp64 oracle, from M1-M6.
  - Latency at Bole's shape (T = 128, depth ≤ 8, top-4) at batch 1 and 16, serial versus
    canopy, next to Bole's Table VII, labeled as different hardware.
  - The solve comparison (M4), the crossover (M7), the memory reduction at T = 128 next to
    Bole's 82-99×, and the 100-round commit drift (M5).
  - The stretch result in tokens per second, or that you skipped it.
  - The total spend, with 5090 and H100 hours separately.
  - Everything that you skipped, that failed, or that you couldn't verify, and every choice
    the papers left open.