canopy

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 (July 18, 2026), Bole (August 3), and TreeWY (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 (September 20) and 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.

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

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 measured an A100 at 295 TFLOPS on zeros and 257 on randn, and Gregersen et al. 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: 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 (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 (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 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: the tree closed form, the finite Neumann solve, and the A100 and DGX Spark numbers.
• TreeWY (Ghantasala, 2026): the same idea as a tree WY transform, with commit by reconstruction.
• SpecLA (Wang et al., 2026): tree masks in the UT transform, and the result that parallel can lose.
• GDN Tree-Scan (Ma, 2026) and STree (Wu et al., NeurIPS 2025): scan and replay in vLLM, and tree speculation for Mamba2 hybrids.
• Gated Delta Networks (Yang et al., ICLR 2025) and flash-linear-attention (MIT): the rule, and its chunked Triton kernels with log-space gates and the WY solve.
• SpecInfer (Miao et al., 2023): token-tree verification for transformers.
• DFlash (Chen, Liang, and Liu, 2026) and DDTree (Ringel and Romano, 2026): a block-diffusion drafter and the draft trees built from it.
• Triton documentation and Colfax's NVFP4 GEMM on RTX PRO Blackwell (sm_12x) 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:

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:

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.