# atlas-gdn

Hand-tuned Gated DeltaNet kernels for the linear-attention path of
Qwen3.6 hybrid models on NVIDIA GB10 (DGX Spark, SM121).

## What's inside

| Op                          | Use                                                    |
|-----------------------------|--------------------------------------------------------|
| `gdn_decode`                | Single-token recurrent decode (FP32 Q/K/V, BF16 out)   |
| `gdn_prefill`               | Multi-token prefill (BF16 throughout)                  |
| `gdn_chunk2` / `gdn_chunk3` | MTP K=2/3 chunkwise verify (Qwen3.6 NVFP4 specialized) |
| `gdn_wy2` / `wy3` / `wy4`   | 2-pass WY-chunkwise verify (general K=2/3/4)           |
| `causal_conv1d_fwd`         | Depthwise causal Conv1d (SSM input projection)         |
| `causal_conv1d_update`      | Single-step Conv1d update (decode)                     |

## Hardware

These kernels target **only** NVIDIA GB10 (compute capability 12.1,
`sm_121f`). They will not load on any other GPU. GB10 has:

- Unified LPDDR5X memory (~273 GB/s) — bandwidth-bound, not occupancy-bound
- No multi-CTA clusters (ClusterShape forced to 1×1×1)
- No `cvt.rn.satfinite.e2m1x2.f32` PTX (software E2M1 conversion path)
- Cooperative-only scheduling (no Pingpong)

`build.toml` pins `cuda-capabilities = ["12.1"]` so the build matrix
yields a single SM121 binary; no fallback binaries are produced.

## Models tested

| Model                                   | Layers using these kernels |
|-----------------------------------------|----------------------------|
| Qwen/Qwen3.6-27B (dense, hybrid)        | 48 GDN layers              |
| Qwen/Qwen3.6-35B-A3B (sparse MoE, hybrid) | 30 GDN layers            |

## Provenance

Sources are extracted from the Atlas inference engine
(<https://github.com/Avarok-Cybersecurity/atlas>, AGPL-3.0). The GDN
NVFP4 variant ships with `__launch_bounds__` annotations specific to
Qwen3.6 hidden dimensions (k_dim=128, v_dim=128, 16/32 K/V heads).

## License

AGPL-3.0-only.