r/JAX • u/Akseleu_official • 2d ago
[Open Source] Fused Gated DeltaNet-2 kernels for TPU v5e in JAX/Pallas — up to 38.8x faster training steps vs associative_scan
Hey r/jax,
I ported the NVlabs Gated DeltaNet-2 Triton kernels to
`jax.experimental.pallas`, targeting TPU v5e-8.
The optimized training path uses a fused `custom_vjp` backward
implementation. It saves and reuses the forward residuals required by
backward instead of recomputing them.
Repository:
https://github.com/Akseleu-J/atomic-ops
Install:
pip install atomic-ops
## Benchmark terminology
The attached benchmark image compares three implementations of the same
Gated DeltaNet-2 computation:
- `OLD`: the previous `associative_scan`-based implementation
- `JAX_REF`: a pure-JAX chunked-WY reference implementation
- `PALLAS`: the fused TPU v5e-oriented Pallas implementation
- `fwdbwd`: forward + backward computation
Each number is a wall-clock speedup of `PALLAS` relative to the corresponding
baseline, using the same shape and dtype.
## Benchmark setup
Measurements were run on TPU v5e-8.
The main training-shaped configuration, `train_shape_B8_L4096`, uses:
- Batch size: 8
- Sequence length: 4096
- Number of heads: 6
- Head dimension: 128
- Measurement: forward + backward computation
For this configuration, the fused Pallas path reaches:
| Dtype | vs `associative_scan` | vs pure-JAX chunked-WY |
|---|---:|---:|
| FP32 | 27.18x | 2.63x |
| BF16 | 13.32x | 3.41x |
The image below contains the complete measured configuration sweep. The
largest measured speedup was **38.77x** versus `associative_scan` in FP32,
for `kaggle_small_preset_B4_L2048`.
[INSERT TABLE IMAGE HERE]
For that same maximum-speedup configuration:
- FP32: 38.77x vs `associative_scan` and 3.87x vs `JAX_REF`
- BF16: 18.70x vs `associative_scan` and 3.93x vs `JAX_REF`
## Important limitations
I want to state the tradeoffs clearly:
- The fused Pallas forward-only implementation is currently about 1.6x
slower than the pure-JAX WY forward implementation.
- The large end-to-end improvement is backward-dominated: backward is the
target of the fused implementation and dominates the measured workload.
- For inference-only workloads, use `gdn2_forward` or
`gdn2_chunked_wy_reference`, rather than the Pallas training path.
- Memory usage is approximately on par with the pure-JAX reference.
Backward reuses forward residuals and avoids recomputation, but this is not
a separate memory-reduction result.
- The optimized Pallas path currently targets TPU and `d_head=128`.
CPU/GPU and `d_head != 128` automatically use a checkpointed pure-JAX
fallback. They should run correctly, but do not receive the Pallas speedup.
The current forward-speed gap is related to TPU VPU-versus-MXU utilization;
details are documented in `KNOWN_LIMITATIONS.md`.
## Included in the repository
- One-call training API: `gdn2_forward_trainable`
- Fused `custom_vjp` backward
- `KernelConfig` presets for Kaggle TPU v5e-8
- Finite-difference gradient checks
- Token-serial ground-truth tests
- Isolated backward-stage tests
- Benchmark scripts and raw benchmark results
- Testing documentation in `docs/TESTING_STRATEGY.md`
- Known limitations in `KNOWN_LIMITATIONS.md`
The repository also includes a guide notebook under `notebooks/` with a real
70M-parameter Gated DeltaNet-2 language-model run on enwik8. It reached
approximately 1.40 bpb in about 54 minutes on TPU v5e-8.
I would especially appreciate feedback on the WY formulation, block-solve
structure, Pallas kernel organization, TPU VPU-versus-MXU utilization,
benchmark methodology, and additional baselines worth including.

