Fuse partial sum accumulation into TGMM kernel

[Shuwen-Fang]

Description

This is a follow up to PR, see previous PR for more context on motivation of this change.

Changes in this PR

1. Backward TGMM Kernel
   • Added optional partial_sum input tensor and block specs to tgmm_v2() .
   • Updated VMEM budget calculations ( calculate_tgmm_tiling ) to allocate double-
   buffered space for incoming partial sum tiles.
   • Updated tgmm_inner_kernel() to accumulate incoming partial sum tiles in MXU
   registers ( acc += tiled_ps_ref ) on expert group transitions.
   • Added zero-token expert group handling ( group_size == 0 ) via jnp.where to
   ensure empty experts preserve incoming gradient partial sums rather than being
   overwritten with zero.
2. Kernel & Integration Test Suite
   • Unit Tests ( tests/unit/pallas_mosaic_tpu_v2_kernel_test.py ): Migrated Tokamax
   test suite into MaxText and expanded coverage for partial sum feature

FIXES: b/522351030
FIXES: b/512156666

Tests

Ran pytest src/maxtext/kernels/megablox/pallas_mosaic_tpu_v2_gmm_kernel_test.py

Checklist

Before submitting this PR, please make sure (put X in square brackets):

• I have performed a self-review of my code. For an optional AI review, add the gemini-review label.
• I have necessary comments in my code, particularly in hard-to-understand areas.
• I have run end-to-end tests tests and provided workload links above if applicable.
• I have made or will make corresponding changes to the doc if needed, including adding new documentation pages to the relevant Table of Contents (toctree directive) as explained in our documentation.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

📢 Thoughts on this report? Let us know!