Skip to content

perf(gptoss): tune bf16 GEMM shapes to non-hipblaslt kernels - #4352

Merged
valarLip merged 1 commit into
mainfrom
perf/gptoss-bf16-gemm-tune
Jul 23, 2026
Merged

valarLip merged 1 commit into
mainfrom
perf/gptoss-bf16-gemm-tune

Conversation

@valarLip

Copy link
Copy Markdown
Collaborator

Problem

Starting gpt-oss-120b with -tp 2 --enable-dp-attention --enable-expert-parallel --enable-tbo crashes during CUDA/HIP graph capture:

Hip error: 'operation not permitted when stream is capturing'(900) at hipblaslt.cpp:147
Engine Core: load model runner failed

Root cause

Several gpt-oss bf16 GEMM shapes have no tuned config, so they fall back to the torch path (F.linear → hipBLASLt). Under --enable-tbo, decode batches are split into bs/2 micro-batches, producing GEMM shapes (e.g. m=2880 n=128 k=4096, the o/down projection at the TBO ubatch token count) that are executed for the first time during cudagraph capture. hipBLASLt's algorithm/workspace init on first touch of a shape is stream-capture-illegal → HIP error 900.

Confirmed via HIPBLASLT_LOG_MASK=32 trace: the crashing matmul carries a hipBLASLt solution_index, i.e. it went through the hipBLASLt path.

Fix

Tune the affected gpt-oss bf16 GEMM shapes (all dtype=bf16 outdtype=bf16) so they resolve to direct-launch kernels instead of torch/hipBLASLt:

N K role
2880 4096 o / down proj
5120 2880 qkv / gate-up
128 2880 router gate
201088 2880 lm_head

Result: all resolve to asm / flydsl / triton / opus — no torch/hipblaslt. Entries live in the per-model model_configs/gptoss_bf16_tuned_gemm.csv (not the shared bf16_tuned_gemm.csv), sorted by gfx with gfx1250 last.

Tuned with gemm_a16w16_tune.py --libtype asm,opus,flydsl,triton,skinny (hipblaslt and torch excluded).

Verification (gfx950, MI355X)

  • Server starts and passes cudagraph capture — no HIP error 900.
  • GSM8K (chat, num_concurrent=65, num_fewshot=3): flexible-extract 0.872–0.884 across runs (threshold 0.87, baseline 0.9) — accuracy preserved.

Tune gpt-oss-120b bf16 GEMM shapes (o/down-proj N=2880 K=4096, qkv/gate-up
N=5120, router gate N=128, lm_head N=201088; K=2880/4096) to asm/flydsl/
triton/opus kernels so they no longer fall back to torch/hipblaslt. The
hipblaslt algo/workspace init is stream-capture-illegal and crashes
cudagraph capture with HIP error 900 under -tp 2 --enable-dp-attention
--enable-expert-parallel --enable-tbo. Entries live in the per-model
config, sorted by gfx with gfx1250 last.
@valarLip
valarLip requested a review from a team July 23, 2026 10:28
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 4352 --add-label <label>

@valarLip
valarLip merged commit 219d610 into main Jul 23, 2026
40 checks passed
@valarLip
valarLip deleted the perf/gptoss-bf16-gemm-tune branch July 23, 2026 13:52
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.

1 participant