Skip to content

Use attention ops from Tokamax#4465

Open
copybara-service[bot] wants to merge 1 commit into
mainfrom
test_945567777
Open

Use attention ops from Tokamax#4465
copybara-service[bot] wants to merge 1 commit into
mainfrom
test_945567777

Conversation

@copybara-service

Copy link
Copy Markdown
Contributor

Use attention ops from Tokamax

Pallas Triton is heading towards deprecation, so these ops will eventually be
removed from JAX.

Closes #4384.

Pallas Triton is heading towards deprecation, so these ops will eventually be
removed from JAX.

Closes #4384.

PiperOrigin-RevId: 945567777
@codecov

codecov Bot commented Jul 14, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 15.38462% with 11 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
src/maxtext/layers/attention_op.py 15.38% 11 Missing ⚠️

📢 Thoughts on this report? Let us know!

out = gpu_pallas_attention.mha(query, key, value, decoder_segment_ids, sm_scale=1.0, causal=True)
mask = tokamax_attention_base.Mask(is_causal=True)
if decoder_segment_ids is not None:
seg_mask = decoder_segment_ids[:, :, None] == decoder_segment_ids[:, None, :]

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.

I needed to add the following to make it work:

seg_mask = seg_mask[:, None, :, :]

Tokamax infers batch dims by rank, so a rank-3 (B, T, t) mask matches its rank-3 (#heads, T, t) annotation exactly and the batch dim gets misread as the head dim — adding an explicit singleton head dim ((B, 1, T, t)) lets it vmap over batch and broadcast across heads.

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.

Deprecate jax.experimental.pallas GPU splash

2 participants