Skip to content

Add support for Tokamax GMM v2 in MaxText. - #4197

Merged
copybara-service[bot] merged 1 commit into
AI-Hypercomputer:mainfrom
CaptainO5:integrate_v2
Jun 19, 2026
Merged

copybara-service[bot] merged 1 commit into
AI-Hypercomputer:mainfrom
CaptainO5:integrate_v2

Conversation

@CaptainO5

Copy link
Copy Markdown
Collaborator

Description

This change temporarily forks and integrates a new Pallas Mosaic TPU v2 ragged dot implementation from Tokamax. A new configuration option use_gmm_v2 is added to enable this feature. GMM v2 is currently not supported when used with quantization. The forward and backward passes in the grouped matrix multiplication (GMM) operation are updated to use the new v2 implementation when enabled.

BUGS: b/505011787

Tests

I added a unit test in integration/tokamax_test and also verified that the e2e test ran successfully.

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 Bot commented Jun 18, 2026 •

Copy link
Copy Markdown

Comment thread src/maxtext/kernels/megablox/pallas_mosaic_tpu_v2_gmm_kernel.py Outdated
Comment thread src/maxtext/kernels/megablox/pallas_mosaic_tpu_v2_tgmm_kernel.py Outdated
Comment thread tests/integration/tokamax_test.py Outdated
Comment thread tests/integration/tokamax_test.py
Comment thread src/maxtext/layers/moe.py Outdated
Comment thread src/maxtext/kernels/megablox/ops.py

@shuningjin shuningjin left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Thank you for integrating this performant new kernel! I see we are able to use full v2 with bf16 now, which is great.

  • Had a question on transpose_rhs
  • As mentioned by others, let's move the check to types.py (for use_gmm_v2, use_tokamax_gmm, quantization).

Comment thread src/maxtext/kernels/megablox/ops.py Outdated
Comment thread src/maxtext/kernels/megablox/ops.py
Comment thread src/maxtext/layers/moe.py
Comment thread src/maxtext/kernels/megablox/pallas_mosaic_tpu_v2_gmm_kernel.py Outdated
Comment thread src/maxtext/kernels/megablox/pallas_mosaic_tpu_v2_tgmm_kernel.py Outdated
Comment thread src/maxtext/configs/base.yml Outdated

@gobbleturk gobbleturk left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Awesome, thank you!

@shuningjin shuningjin left a comment •

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM. Thank you!

Comment thread src/maxtext/configs/types.py
@copybara-service
copybara-service Bot merged commit 85690da into AI-Hypercomputer:main Jun 19, 2026
29 of 30 checks passed
@shuningjin shuningjin mentioned this pull request Jul 15, 2026
4 tasks done
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants