Skip to content

[ROCm] Enable gluon paged MQA logits on gfx950 (MI355X) #42062

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
frida-andersson wants to merge 2 commits into
base: main
Choose a base branch
from
+3 −2
Open

[ROCm] Enable gluon paged MQA logits on gfx950 (MI355X) #42062

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
7bd872b
[ROCm] Enable gluon paged MQA logits on gfx950 (MI355X)
frida-andersson May 8, 2026
d614a81
Merge branch 'main' into rocm/gluon-gfx950-dispatch
tjtanaa May 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 3 additions & 2 deletions vllm/v1/attention/ops/rocm_aiter_mla_sparse.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,9 +16,10 @@
from vllm.v1.attention.ops.common import pack_seq_triton, unpack_seq_triton

if current_platform.is_rocm():
from vllm.platforms.rocm import _ON_GFX942
from vllm.platforms.rocm import _ON_GFX942, _ON_GFX950
else:
_ON_GFX942 = False
_ON_GFX950 = False
Comment on lines 18 to +22
Copy link
Copy Markdown
Contributor

@gemini-code-assist gemini-code-assist Bot May 8, 2026

Choose a reason for hiding this comment

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

high

For better maintainability and to avoid listing individual GPU architectures, consider using the _ON_MI3XX flag which is a boolean that is true for both gfx942 and gfx950. This will make the code cleaner and easier to extend for future MI300-series GPUs.

Suggested change
if current_platform.is_rocm():
from vllm.platforms.rocm import _ON_GFX942
from vllm.platforms.rocm import _ON_GFX942, _ON_GFX950
else:
_ON_GFX942 = False
_ON_GFX950 = False
if current_platform.is_rocm():
from vllm.platforms.rocm import _ON_MI3XX
else:
_ON_MI3XX = False

Copy link
Copy Markdown
Contributor

@maeehart maeehart May 8, 2026

Choose a reason for hiding this comment

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

In this case we have to be more specific as Mi308 might have some exceptions.



@triton.jit
Expand Down Expand Up @@ -385,7 +386,7 @@ def rocm_fp8_paged_mqa_logits(
aiter_paged_mqa_logits_module = paged_mqa_logits_module()

if aiter_paged_mqa_logits_module is not None:
if _ON_GFX942:
if _ON_GFX942 or _ON_GFX950:
Copy link
Copy Markdown
Contributor

@gemini-code-assist gemini-code-assist Bot May 8, 2026

Choose a reason for hiding this comment

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

high

Following the suggestion to use _ON_MI3XX, this condition can be simplified.

Suggested change
if _ON_GFX942 or _ON_GFX950:
if _ON_MI3XX:

Copy link
Copy Markdown
Contributor

@maeehart maeehart May 8, 2026

Choose a reason for hiding this comment

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

We need both _on_gfx942 and _on_gfx950 elsewhere, so it does not also make sense to apply above recommendation. Also see the comment above.

deepgemm_fp8_paged_mqa_logits = (
aiter_paged_mqa_logits_module.deepgemm_fp8_paged_mqa_logits
)
Expand Down
Loading