Skip to content

[Stacked PR 3/5] Add decoupled GDN Pallas backward pass kernel and parity tests - #5153

Open
Rohan-Bierneni wants to merge 1 commit into
rbierneni-gdn-2-fwd-kernelfrom
rbierneni-gdn-3-bwd-kernel
Open

[Stacked PR 3/5] Add decoupled GDN Pallas backward pass kernel and parity tests#5153
Rohan-Bierneni wants to merge 1 commit into
rbierneni-gdn-2-fwd-kernelfrom
rbierneni-gdn-3-bwd-kernel

Conversation

@Rohan-Bierneni

@Rohan-Bierneni Rohan-Bierneni commented Sep 6, 2026

Copy link
Copy Markdown
Collaborator

Stacked PR Chain

Stack Status Branch Base PR
PR 1 🔗 Stack PR 1 rbierneni-gdn-1-ci-hygiene main #5151 - [Stacked PR 1/5] Fix upstream pyink formatting and GMM v2 compatibility for Qwen3.5
PR 2 🔗 Base PR rbierneni-gdn-2-fwd-kernel rbierneni-gdn-1-ci-hygiene #5152 - [Stacked PR 2/5] Import local Tokamax GDN forward kernel with custom remat for backward pass support
PR 3 🚀 This PR rbierneni-gdn-3-bwd-kernel rbierneni-gdn-2-fwd-kernel #5153 - [Stacked PR 3/5] Add decoupled GDN Pallas backward pass kernel and parity tests
PR 4 ⏳ Stacked on PR 3 rbierneni-gdn-4-model-integration rbierneni-gdn-3-bwd-kernel #5154 - [Stacked PR 4/5] Enable GDN backward pass kernel and hybrid precision in Qwen3 model
PR 5 ⏳ Stacked on PR 4 rbierneni-gdnv3-bwd rbierneni-gdn-4-model-integration #5098 - [Stacked PR 5/5] Add GDN backward pass 8k and 64k latency and memory benchmark suite

Description

This is PR 3 of 5 in the stacked series enabling the Pallas Gated Delta Net (GDN) backward pass kernel in MaxText.

This PR introduces the decoupled GDN Pallas backward pass kernel and complete analytical gradient parity testing:

  1. Decoupled Backward Architecture (gdn_bwd_pallas.py):
    • pallas_gdn_bwd_kernel: Executes 40+ GDN adjoint matrix recurrences via pltpu.emit_pipeline with vectorized head processing and zero HBM intermediate spills.
    • conv1d_silu_bwd: Executes immediately after Pallas via unrolled native JAX shift-multiply. Decoupling Conv1D backward avoids binding Conv1D gradient live ranges with GDN adjoint matrix state buffers, eliminates vector register spills, and keeps peak VMEM well within hardware boundaries.
    • gdn_decoupled_conv1d: Custom VJP unifying forward execution (reusing $T^{-1}$ cached in residuals from PR 2) and backward execution. Named gdn_decoupled_conv1d (rather than fused) to accurately convey that Conv1D is decoupled in native JAX rather than fused with GDN recurrence in Pallas.
    • Aliases: Provides gdn_kernel and backward-compatible aliases.
  2. Exhaustive Analytical Parity Test Suite (tests/unit/gdn_bwd_pallas_test.py):
    • Validates analytical gradients against pure JAX autodiff across sequence lengths (128 to 2048), batch sizes, and multi-head settings.
    • Asserts strict numerical parity across $dQ, dK, dV, d�eta, d�lpha, dW_{conv}, db_{conv}$.

Files Changed

  • src/maxtext/models/kernels/gdn/__init__.py: Package export for gdn_decoupled_conv1d, gdn_kernel, pallas_gdn_bwd_kernel.
  • src/maxtext/models/kernels/gdn/gdn_bwd_pallas.py: Decoupled backward pass kernel and custom VJP implementation (gdn_decoupled_conv1d).
  • tests/unit/gdn_bwd_pallas_test.py: Full analytical gradient parity test suite against JAX autodiff.

Tests

  • Verified with pre-commit run --files ... (all hooks passed: codespell, pylint, pyink).
  • Parity test suite: pytest tests/unit/gdn_bwd_pallas_test.py validates analytical gradient equivalence to machine precision (13/13 unit tests passed in 137s).

Checklist

  • I have performed a self-review of my code.
  • I have necessary comments in my code, particularly in hard-to-understand areas.
  • I have run end-to-end tests and verified pre-commit linters pass.

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces a local Gated Delta Net (GDN) backward kernel package using Pallas emit_pipeline with triangular inverse matrix caching, alongside a decoupled Conv1D + SiLU forward and backward implementation in JAX, and comprehensive unit tests. The review feedback identifies two key issues: an epsilon mismatch (1e-12 vs. 1e-6) when use_qk_norm_in_gdn is enabled, which could cause numerical instability, and an unconditional gradient reduction for a_log and dt_bias that fails to preserve batch dimensions for 2D or 3D inputs.

Comment on lines +719 to +730
if use_qk_norm_in_gdn:
d_q_scaled = d_q_proj * scale
r_q = jnp.sqrt(jnp.sum(q_orig**2, axis=-1, keepdims=True) + 1e-12)
q_unit = q_orig / r_q
d_q = (d_q_scaled - q_unit * jnp.sum(d_q_scaled * q_unit, axis=-1, keepdims=True)) / r_q

r_k = jnp.sqrt(jnp.sum(k_orig**2, axis=-1, keepdims=True) + 1e-12)
k_unit = k_orig / r_k
d_k = (d_k_proj - k_unit * jnp.sum(d_k_proj * k_unit, axis=-1, keepdims=True)) / r_k
else:
d_q = d_q_proj * scale
d_k = d_k_proj

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

There is an epsilon mismatch between the forward and backward passes when use_qk_norm_in_gdn is enabled. In the forward pass, normalizations.l2norm is called with eps=1e-6. However, in the backward pass, r_q and r_k are computed using 1e-12 as the epsilon.

While this mismatch is negligible in unit tests with random normal inputs (where the sum of squares is large), it can lead to mathematically incorrect gradients, gradient explosion, or NaNs in real training scenarios—especially for padded/masked tokens or during early initialization where values can be extremely small or zero. Changing 1e-12 to 1e-6 ensures mathematical consistency and training stability.

Suggested change
if use_qk_norm_in_gdn:
d_q_scaled = d_q_proj * scale
r_q = jnp.sqrt(jnp.sum(q_orig**2, axis=-1, keepdims=True) + 1e-12)
q_unit = q_orig / r_q
d_q = (d_q_scaled - q_unit * jnp.sum(d_q_scaled * q_unit, axis=-1, keepdims=True)) / r_q
r_k = jnp.sqrt(jnp.sum(k_orig**2, axis=-1, keepdims=True) + 1e-12)
k_unit = k_orig / r_k
d_k = (d_k_proj - k_unit * jnp.sum(d_k_proj * k_unit, axis=-1, keepdims=True)) / r_k
else:
d_q = d_q_proj * scale
d_k = d_k_proj
if use_qk_norm_in_gdn:
d_q_scaled = d_q_proj * scale
r_q = jnp.sqrt(jnp.sum(q_orig**2, axis=-1, keepdims=True) + 1e-6)
q_unit = q_orig / r_q
d_q = (d_q_scaled - q_unit * jnp.sum(d_q_scaled * q_unit, axis=-1, keepdims=True)) / r_q
r_k = jnp.sqrt(jnp.sum(k_orig**2, axis=-1, keepdims=True) + 1e-6)
k_unit = k_orig / r_k
d_k = (d_k_proj - k_unit * jnp.sum(d_k_proj * k_unit, axis=-1, keepdims=True)) / r_k
else:
d_q = d_q_proj * scale
d_k = d_k_proj

Comment on lines +922 to +923
d_a_log_reduced = jnp.sum(d_a_log_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(a_log.dtype)
d_dt_bias_reduced = jnp.sum(d_dt_bias_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(dt_bias.dtype)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

The forward path of _pallas_gdn_bwd_kernel_single_group explicitly supports 1D, 2D, and 3D shapes for a_log and dt_bias. However, the backward path unconditionally reduces the gradients over axis=(0, 1) (batch and chunk dimensions):

d_a_log_reduced = jnp.sum(d_a_log_chunks[..., 0, :num_v_heads], axis=(0, 1))

If a_log or dt_bias is 2D (e.g., batch-dependent), this unconditional reduction over axis=0 will destroy the batch dimension, leading to shape mismatches or incorrect gradient updates in JAX. We should dynamically reduce the gradients based on the input dimensions of a_log and dt_bias.

Suggested change
d_a_log_reduced = jnp.sum(d_a_log_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(a_log.dtype)
d_dt_bias_reduced = jnp.sum(d_dt_bias_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(dt_bias.dtype)
if a_log.ndim == 1:
d_a_log_reduced = jnp.sum(d_a_log_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(a_log.dtype)
elif a_log.ndim == 2:
d_a_log_reduced = jnp.sum(d_a_log_chunks[..., 0, :num_v_heads], axis=1).astype(a_log.dtype)
else:
d_a_log_reduced = d_a_log_chunks[..., 0, :num_v_heads].astype(a_log.dtype)
if dt_bias.ndim == 1:
d_dt_bias_reduced = jnp.sum(d_dt_bias_chunks[..., 0, :num_v_heads], axis=(0, 1)).astype(dt_bias.dtype)
elif dt_bias.ndim == 2:
d_dt_bias_reduced = jnp.sum(d_dt_bias_chunks[..., 0, :num_v_heads], axis=1).astype(dt_bias.dtype)
else:
d_dt_bias_reduced = d_dt_bias_chunks[..., 0, :num_v_heads].astype(dt_bias.dtype)

@Rohan-Bierneni Rohan-Bierneni changed the title [Stacked PR 3/5] Add canonical decoupled GDN Pallas backward kernel and parity tests [Stacked PR 3/5] Add decoupled GDN Pallas backward pass kernel and parity tests Sep 6, 2026
@Rohan-Bierneni
Rohan-Bierneni force-pushed the rbierneni-gdn-3-bwd-kernel branch from 9238216 to d4ca233 Compare September 6, 2026 06:43
…rity tests

- Implement decoupled GDN backward pass using Pallas emit_pipeline (pallas_gdn_bwd_kernel) and custom VJP (gdn_decoupled_conv1d).
- Execute Conv1D backward immediately following Pallas via unrolled native JAX shift-multiply, eliminating live range interference and vector register spills.
- Provide comprehensive analytical gradient parity test suite (gdn_bwd_pallas_test.py) validating exact gradient match against pure JAX autodiff across varying sequence lengths, batch sizes, and head configurations.
@Rohan-Bierneni
Rohan-Bierneni force-pushed the rbierneni-gdn-3-bwd-kernel branch from d4ca233 to c47f1a3 Compare September 6, 2026 07:00
@codecov

codecov Bot commented Sep 6, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 83.28267% with 110 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
src/maxtext/models/kernels/gdn/gdn_bwd_pallas.py 82.91% 82 Missing and 28 partials ⚠️

📢 Thoughts on this report? Let us know!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant