Skip to content

[AMD] Keep the MiniMax sparse decode score tile within gfx942 LDS #41354

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
siliangchen-amd wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[AMD] Keep the MiniMax sparse decode score tile within gfx942 LDS #41354

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
9 changes: 6 additions & 3 deletions python/sglang/kernels/ops/attention/minimax_sparse/decode/flash_with_topk_idx.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@
import triton.language as tl

from sglang.srt.environ import envs
from sglang.srt.utils import is_xpu
from sglang.srt.utils import is_gfx95_supported, is_xpu

from ..common.utils import (
_bitonic_merge,
Expand Down Expand Up @@ -35,11 +35,14 @@ def _prune_decode_configs(configs, named_args, **kwargs):


_ON_HIP = torch.version.hip is not None
_ON_GFX95 = _ON_HIP and is_gfx95_supported()


def _decode_score_block_n(args):
# gfx950 pick: 512 for tiny batches (few CTAs), else 128.
return max(512 if args["batch_size"] <= 4 else 128, args["block_size"])
# gfx950 pick: 512 for tiny batches (few CTAs), else 128. A 512-token K tile
# needs 128 KiB of LDS, which gfx942 (64 KiB) does not have.
tiny_batch = _ON_GFX95 and args["batch_size"] <= 4
return max(512 if tiny_batch else 128, args["block_size"])


# On ROCm the autotune sweep is replaced by a fixed config plus the
Expand Down
Loading