File size: 2,008 Bytes
3f6ddf9 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 | # 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. |