Skip to content

M3 perf - #39466

Draft
kevin-mii wants to merge 155 commits into
sgl-project:mainfrom
kevin-mii:M3-perf
Draft

M3 perf#39466
kevin-mii wants to merge 155 commits into
sgl-project:mainfrom
kevin-mii:M3-perf

Conversation

@kevin-mii

@kevin-mii kevin-mii commented Sep 14, 2026 •

Copy link
Copy Markdown
Collaborator

Motivation

Modifications

Accuracy Tests

Speed Tests and Profiling

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS and other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
    • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
  4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

CI States

Latest PR Test (Base): ❌ Missing run-ci label -- add it to run CI tests.
Latest PR Test (Extra): ❌ Blocked -- run-ci is required first.
Latest PR Test (AMD ROCm 10): ➖ No AMD PR run found for this commit.

zcnrex and others added 30 commits September 1, 2026 01:27
Route the MiniMax-M3 sparse prefill main attention (step 3) through AITER's
Gluon paged attention instead of the Triton `flash_prefill_with_gqa_share_sparse`
kernel, behind `SGLANG_OPT_USE_GLUON_PREFILL`. The main KV pool stays NHD
`[max_slots, 1, head_dim]`; each request's context is gathered into persistent
SHUFFLE 5D scratch pages (page size 64, two pages per 128-token sparse block)
and the per-query sparse block table is expanded to a page table, so
`pa_decode_gluon(..., ps=True)` can serve the whole extend batch.

A cheap static gate (`can_use_gluon_prefill`) plus a runtime shape check keep
every unsupported case on the existing Triton path; the fallback logs once.

Also in the prefill kernels this path shares:
- a score-only index-attention kernel for `disable_index_value` layers, with
  per-page K addressing (one base-slot load per page instead of one
  `req_to_token` lookup per token) when a sparse block spans few enough pages;
- CDNA KV sub-tiling in the Triton sparse kernel (`SUB_K`, 0 elsewhere) so each
  QK/PV MFMA is right-sized on gfx942/gfx950.

Co-Authored-By: Kevin Mi <mikevin920@yahoo.com>
Co-Authored-By: Alex Sun <alex.s@amd.com>
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
- matrix_instr_nonkdim/kpack are AMD-only Triton launch kwargs that
  NVIDIA's Triton rejects; only offer those autotune configs on HIP.
- Route the arch detection for KV sub-tiling through is_gfx95_supported /
  is_gfx942_supported instead of parsing gcnArchName.
- The score-only indexer kernel and the Gluon dispatch were only validated
  on gfx950; gate them (_is_hip, and _use_aiter + gfx95 respectively) so
  other platforms keep the shared kernels. The Gluon gate also stops
  non-gfx95 ROCm from allocating scratch pages before the aiter call fails.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01GbK4exGWvsmPdMf3kmBYW4
…_USE_GLUON_PREFILL

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Put the lightning-indexer K cache in `fp8_e4m3fn` on ROCm behind
`SGLANG_OPT_MINIMAX_M3_FP8_INDEX_CACHE`, halving the index pool's bytes.
This is a widening mode -- bf16 q against fp8 k in the score kernels, the
same contract the SM100 fp8 attn-GEMM mode already uses -- and it is
selection-only: the main attention K/V stay in `kv_cache_dtype`.

The Triton qk-norm-rope kernel now stores the index cache in that cache's
own dtype instead of q's, a no-op cast when the two match.

Also `#ifndef USE_ROCM` the device-side assert in `store_kvcache`: HIP
device assertions need hostcall support, which is unavailable during
hipGraph capture, so leaving it in aborts capture with
`hipErrorIllegalState`. ROCm suppresses the invalid write with the same
bounds test moved into the store's own guard instead. The file is
runtime-JIT compiled, so this needs no image rebuild.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Co-Authored-By: Kevin Mi <mikevin920@yahoo.com>
The per-token cell-size estimate hardcoded the lightning-indexer dtype
to model_dtype (bf16) while the constructed pool allocates it in fp8
under the fp8 attn-GEMM mode or SGLANG_OPT_MINIMAX_M3_FP8_INDEX_CACHE,
so the profiled byte budget converted with 2 extra bytes/token per
indexer layer and the pool left the difference unused. Move the dtype
choice into KVCacheConfigurator.minimax_sparse_index_dtype() and use it
at both sites.

MI350X TP8, mem-fraction-static 0.80: max_total_num_tokens
3,189,150 -> 3,686,466 (+15.6%). bs32/il80000 unchanged
(2142/2143/2138 tok/s); bs44/il80000 (3.52M tokens, previously beyond
capacity) now runs at 2574 tok/s.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
torch.float8_e4m3fn is the OCP format; pre-gfx95 CDNA is e4m3fnuz and
was never tested with this cache, so gate the dtype selection (the sole
consumer of SGLANG_OPT_MINIMAX_M3_FP8_INDEX_CACHE) on gfx95.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Replace the decode block top-k's radix-select with a wave64 histogram-select on
ROCm, and raise TopKTrait::kMaxNumBlocks from 4096 to 16384 so the kernel is
usable under CUDA graph capture at long context.

The ROCm path packs (score, ~block_id) into a u64 total order, runs one coarse
histogram round (256 bins up to 1024 rows, else 2048), refines with 11-bit
rounds while more than 128 candidates remain, and finishes with a wave-0
bitonic sort. Bins above the threshold bin are emitted immediately, bins below
die, and the threshold bin survives to the next round, so the row is read from
global memory exactly once. Selection semantics are unchanged; only the
reduction strategy differs. The ascending-order emit at the end of the block-id
kernel now uses a single 64-lane ballot instead of two width-32 reductions.
TopKTrait becomes TopKTraitImpl<kNumThreads> so the ROCm launchers can pick a
1024-thread instantiation when max_seqblock > 1024; CUDA still uses only the
512-thread one.

kMaxNumBlocks 4096 -> 16384 is what unblocks CUDA graphs: capture bounds
max_seqblock by the full context length rather than the batch's real sequence
lengths, so any context above 4096 * block_size (524288 tokens at
block_size 128) tripped the max_seqblock RuntimeCheck at capture time. The
Python wrapper carries the matching _MAX_NUM_BLOCKS assert in both
minimax_decode_topk and minimax_decode_topk_page_table.

Co-Authored-By: Kevin Mi <mikevin920@yahoo.com>
The kMaxNumBlocks bump to 16384 drives kIters (and the multi-element
regime's per-thread key registers) on the unmeasured CUDA path; scope it
to ROCm, mirror the split in _MAX_NUM_BLOCKS and the dispatch bound, and
sweep the registered test to the 2M-token ceiling on HIP.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
aiter's opus MoE sorting kernel is sized for thousands of tokens. At decode
batch sizes it costs ~7us/layer of launch and ramp overhead, and the separate
fused_dynamic_mxfp8_quant_moe_sort stage1 activation quant costs another ~4us.
With P = M*topk <= 64 routing pairs the entire sorting job -- stable
sort-by-expert, per-expert block padding, the expert-id table, num_valid and
the moe_buf zero-fill -- fits in one small Triton program built on P x P rank
compares, and the stage1 mxfp8 quant rides the same launch on dedicated CTAs.

Added as three runtime patch points on aiter.fused_moe's module namespace
(fused_moe, _moe_sorting_impl, fused_dynamic_mxfp8_quant_moe_sort), each
falling back to the original function whenever the fast path does not apply,
so a stock aiter install needs no changes. The patch is installed lazily from
AiterRunnerCore.__init__ and is idempotent, so nothing outside the aiter MoE
runner is touched.

Two supporting changes in the same area:

* sanitize_topk_ids on the native gfx950 MXFP8 MoE. The defensive
  int32 copy + masked_fill_ of out-of-range expert ids is only needed when
  expert parallelism can put a non-local id in topk_ids. The Triton runner now
  passes the same num_experts != num_local_experts predicate it already uses
  for filter_expert on the non-MXFP8 path, so the copy is skipped when every
  expert is local. Default stays True, so out-of-tree callers are unaffected.

* aiter_swiglu_oai_limit. aiter's ActivationType.Swiglu is SwiGLU-OAI with
  alpha=1.702 and the (up + 1.0) bias baked in as compile-time constants. The
  helper recognises exactly that configuration and returns its clamp limit, so
  the bf16 aiter runner can use the kernel instead of silently running plain
  SiLU. The aiter CK bf16 fused-MoE has no SwiGLU-OAI codegen at all, so
  UnquantizedFusedMoEMethod now reports why a layer is unsupported --
  128-alignment or SwiGLU-OAI -- and falls back to the Triton runner; the same
  predicate gates the CK weight shuffle, so the two stay consistent.
…=256

Two variants dispatched by pair count: the proven single-sort-CTA compact
kernel for P <= 64, and a distributed kernel (per-expert-chunk and per-pair
CTAs, one writer per byte) for 64 < P <= 256. The split keeps decode's
smallest batches at the compact kernel's 2.7 us while covering bs <= 51
with the fused sort+quant path.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Lifts the CUDA-only guard on shared-experts fusion so ROCm can use it, gated at
gfx942. minimax_m3_vl.py carries the same gate because the served MXFP8
checkpoint is the VL variant (model_type minimax_m3_vl) -- patching only one of
the two files is a silent no-op.

Deliberately does NOT include the shared-expert alternate-stream overlap that
was originally bundled with this change. The alternate stream is only created
when shared-experts fusion is disabled, so the two are mutually exclusive: with
fusion enabled the overlap is dead code.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The runtime patch previously installed on every aiter build from
AiterRunnerCore.__init__; an aiter without fused_dynamic_mxfp8_quant_moe_sort
would AttributeError and take down every aiter MoE model. Early-return off
gfx950 (the only validated arch), tolerate missing entry points, and defer
the module import so the aiter runner import graph stays Triton-free.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
get_device_capability on ROCm reports the parsed gfx number, so the
(9, 4) tuple compare also admitted RDNA (gfx11xx reports (11, 0)); use
is_gfx942_supported / is_gfx95_supported instead, which also avoids
initializing the device context in the classmethod.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Add a dense-only MXFP8 -> block-fp8 conversion mode
(SGLANG_FORCE_MXFP8_BLOCK_CONVERT_DENSE) so dense linears can run on the
mature aiter block-scale GEMMs while fused-MoE weights stay on the native
MX per-1x32 path. This is needed when converting every MXFP8 weight would
route the MoE to a block-scale kernel that cannot service the model's
activation (SwiGLU-OAI on MiniMax-M3).

Alongside it:
- route MXFP8 dense through torch._scaled_mm Blockwise-1x32 when a runtime
  availability probe says the build supports it, falling back to the
  existing dot_scaled Triton kernel otherwise;
- add a "bf16" --fp8-gemm-backend choice that dequantises MXFP8 weights
  once and uses BF16 F.linear, for A/B and for shapes without a kernel;
- cache a rowwise-fp8 (per-token x per-channel) copy of converted dense
  weights and use it below MXFP8_DENSE_PTPC_DECODE_MAX_M tokens, where the
  ptpc GEMM beats the block-fp8 GEMMs;
- make dispatch_w8a8_mxfp8_linear handle every Mxfp8DenseGemmBackend
  member explicitly instead of falling through to the ROCm import.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
SGLANG_FORCE_MXFP8_BLOCK_CONVERT_DENSE now only applies on HIP,
--fp8-gemm-backend bf16 is rejected at startup off gfx950 instead of at
first dispatch, and the aiter MoE weight reshuffle requires gfx95.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…l-reduce with it

The MXFP8 dense decode path (previous commit) consumes a per-token fp8
(tensor, scale) pair below MXFP8_DENSE_PTPC_DECODE_MAX_M tokens. Producing
that pair with a separate per_token_quant launch re-reads the just-written
normed activation from HBM. Emit it from the producer instead:

- the ROCm Triton Gemma fused-add-RMSNorm gains `emit_fp8`, writing the fp8
  tensor and per-token scale in the same pass and handing them to the
  consumer as `out._fp8_qinput = (fp8, scale)`;
- a fused all-reduce + RMSNorm + per-token-quant path
  (`custom_fused_ar_rms_quant`) folds the previous layer's MoE all-reduce
  into the next layer's norm and qkv quantization, exposed as
  `LayerCommunicator(enable_fused_ar_quant_per_token=True)`. It returns
  None whenever the backend cannot service the request, and every caller
  falls back to the existing unquantized fused-AR-RMSNorm path.

Both are gated on `M <= MXFP8_DENSE_PTPC_DECODE_MAX_M` and gfx950, i.e. the
exact window in which the ptpc decode GEMM -- the only consumer of the
pre-quantized pair -- is selected. Off that window nothing is emitted.

This also passes `is_last_layer` to MiniMaxM3DecoderLayer's
LayerCommunicator, matching DeepSeek-V2. Without it the final layer
published `should_allreduce_fusion`, skipped `postprocess_layer`, and left
its MoE output TP-partial with no next layer to absorb the reduction --
which is what the removed "sparse layers must all-reduce immediately"
workaround was papering over. With the terminator in place that blanket
disable is no longer needed, so the sparse MoE all-reduce can be deferred
into the next layer's fused norm again.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
After vLLM ROCm's wvSplitKQ (vendored in aiter custom_kernels.cu): fp8
operands feed MFMA directly with nontemporal weight streaming. Adapted
to the (16,16)-preshuffled layout -- v_mfma_scale_f32_16x16x128_f8f6f4's
B tile is four consecutive 512B blocks of the shuffle, so each lane
supplies 32 weight bytes with two full-width 16B loads -- and to
per-channel weight scales, which wvSplitKQ (per-tensor scales, row-major
weights) cannot express without a second weight copy. One CTA per
n-tile, four waves splitting K through an LDS reduce; no cross-CTA
traffic.

Cold on the TP8 qkv shapes (1280/1536 x 6144): 9.2/6.9 -> 5.5/5.6 us;
per-output error never worse than the aiter path; e2e bs1/il8192
173.4 -> 181.4 tok/s. Dispatched only for the measured-faster shapes on
gfx950; everything else keeps the aiter path.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…oped

- Sparse-layer MoE all-reduce deferral stays disabled off ROCm; the
  is_last_layer terminator remains on every platform.
- enable_fused_ar_quant_per_token only on HIP, and _use_aiter
  short-circuits before the token-dim size check in the communicator.
- skinny GEMV: arch check first in the dispatch predicate, function-local
  import in fp8_utils, lazy fp8_utils constant lookup in layernorm, and a
  gfx950 skip in the registered benchmark (the AMD lane runs gfx942).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… kernel taxonomy

isort wants is_gfx95_supported before is_gfx1250_supported in layernorm.py.

sgl-project#37436 requires newly added registered tests to live under
test/registered/<kind>/<subsystem>/; both files register *-kernel-* suites,
so kind is kernel.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
sgl-project#36527 landed the shared index top-k, which wraps Step 1 and Step 2 in a
cached_topk_idx branch and threads cu_seqblocks_q/cached_topk_idx through the
backend. Resolutions: keep both parameter sets; thread this PR's page_size into
the Step 1 call inside main's else branch; graft the Gluon Step 3 onto main's
structure so the caching and the Gluon path compose -- Gluon runs after the
top-k is resolved either way, and the MSA/Triton fallback stays behind
`o is None`.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
sgl-project#36527 added a MiniMax-M3 sparse-attention section to environ.py in the same
region. Kept both knobs and moved this PR's fp8 index-cache toggle inside that
section, since it is the same class of ROCm MiniMax-M3 knob.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
kevin-mii and others added 28 commits September 15, 2026 03:18
…_MODEL; DRAFT is the draft-token count

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…ve, bench, summary) with no fixed paths; helpers move to tools/

Locations derive from the checkout and M3_WORK (default /scratch if writable, else ~/m3-agentx); preflight checks GPUs,
models, client and imports. Results table: forced c=8 at 3600 s is 13,029 in-window (12,961 reported); ATOM runs every
concurrency for 3600 s, so the low-concurrency points are only comparable at that duration.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…, so the reproduction is clone + one command

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…lts table and chart

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…raph execution; PTPC-FP8 dense becomes opt-in

get_is_capture_mode() is also true while a breakable prefill graph replays, and attention runs there as an
eager break with real shapes, so the cleanup's guard was sending every prefill over a cached prefix to the plain
Triton extend kernel (fresh 100K prefill 2.5 -> 3.0 s, forced c=24 38.2K -> 36.9K tok/s/GPU). The online PTPC-FP8
dense path measured -9% decode at 24 streams on this build (A/B on the same GPUs), so both configs leave it off
unless PTPC_FP8=1.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…nd takes the results root from RESULTS_DIR

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…e it was measured (mirrors PR 36546)

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…nd the follow-up rebase state

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
…nch scripts set both names

The scripts keep exporting the old name too so they still drive builds
predating the rename.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MATMHzpi6K7N2zkpo7mZLc
- MinimaxM3Detector.finish() releases text held back as a possible marker
  prefix (a final "]" was lost on streamed replies; vendor case 12_07).
- X-Request-Id response header on every response, errors included
  (vendor case 06_08 wants a trace id on 4xx replies).
- SGLANG_ENABLE_STRICT_MODEL_NAME: 404 for an unknown `model` (vendor
  case 20_02); off by default.
- The two Kimi-K3 serving tests answer their tool call, as the new
  tool-message pairing validator requires.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Same layer-entry capture as the EAGLE3 hook, driven by the draft
checkpoint's target_layer_ids (nvidia/MiniMax-M3-DSpark).

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…stead of 503

Opt-in; MiniMax's provider check expects overload to surface as HTTP 429.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…ockerfile, build runbook)

serve_endpoint.sh is the production launch (EAGLE3 3 steps / 4 draft tokens, mem 0.9,
strict model name, 429 on a full queue, tool calls parsed with no inventory).
Dockerfile overlays this branch's python tree and the aiter FlyDSL swizzle fix on
lmsysorg/sglang-rocm:v0.5.19-rocm724-mi35x-20260911, the compiled stack the branch
was validated on. build_on_cpu_vm.sh follows rdxa-skills/build-sglang-images-cpu-vm.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…me= glue

At 60-80K tokens of context the model emits <invoke name Read">, <invoke nameRead">,
<invoke name name="Read"> or splices a namespace token into the tag; the strict
prefix match dropped every such call (Provider-Verifier ToolCalls-Match-Rate 84.8%
against the official 98.8%). Parse the text between <invoke and the real tag close,
take the last identifier and prefer a declared tool name (case-insensitive), in both
the block and the streaming parser. The handbook asks providers for exactly this
leniency.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…ACHE, GLUON_PREFILL); handoff doc draft

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
SGLANG_MINIMAX_M3_INDEX_TOPK_FREQ=4 (ATOM's value) makes the model drop the '="'
of its tool-call tags past ~60K tokens of context (Provider-Verifier raw invokes
32/42 malformed at 4, 5/44 at 2, 0/54 at 1; bisected against quick-reduce INT4,
fp8 KV/index caches and the Gluon prefill, which were not the cause). At 1 the
EAGLE3 accept length also rises from 2.65 to 3.07.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…k bisection and the recommendation

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
The base image lmsysorg/sglang-rocm:v0.5.19-rocm724-mi35x-20260911 already
ships 2 of the aiter patch's 4 files pre-modified in its working tree
(pa_mqa_logits_fp4_prefill.py and torch_utils.py, byte-identical to our
post-image blobs). A whole-patch `git apply --check` therefore fails, and
because it sits on the LEFT of an `&&` list its failure does not trip
`set -e` -- so the build silently applied nothing at all, including the
two files that carry the XCD-swizzle fix itself
(mixed_moe_gemm_2stage_common.py, moe_kernels.py). The build still exited
0 and the image looked fine; only the documented "expect 4 files changed"
check caught it.

Apply the patch one file at a time, classifying each as APPLIED /
ALREADY-IN-BASE / FAILED, then assert the end state by content: both XCD
remap guards, the split-K and xcd_swizzle candidate expansion, the two
already-in-base hunks, a compileall over all four, and the 4-file
`git diff --stat`. Not --3way: it implies --index, which would stage the
result and empty the very `git diff` the verification reads.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The wrapper script carried ~20 exported env vars, half of which did
nothing. Replace it with the literal command in ENDPOINT.md, matching the
cookbook shape, after checking every variable against environ.py:

- dead names, read nowhere in the tree: SGLANG_MINIMAX_OPT_USE_GLUON_PREFILL,
  SGLANG_ENABLE_TRITON_EXTEND_LONG_PREFIX
- already the image's ENV: SGLANG_USE_AITER, HIP_FORCE_DEV_KERNARG
- already the code default True: SGLANG_OPT_MINIMAX_M3_FP8_INDEX_CACHE,
  SGLANG_OPT_USE_MINIMAX_GLUON_PREFILL
- redundant when the container gets exactly 4 GPUs: HIP_VISIBLE_DEVICES,
  CUDA_VISIBLE_DEVICES

The 12 that remain are each non-default and load-bearing, and are now
documented in a table with their defaults. Note
SGLANG_MINIMAX_M3_INDEX_TOPK_FREQ's real default is 2, not the 4 the old
knob text claimed -- either way it must be set to 1 explicitly, which is
the tool-call-tag fix.

Also `--tp` -> `--tp-size` (`--tp` is not a declared alias; it resolved
only by argparse prefix matching) and `python -m sglang.launch_server` ->
`sglang serve`, which launch_server.py itself now recommends in a
UserWarning and which is a strict superset: same load_plugins(), same
kill_process_tree cleanup, plus backend detection.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
It has no row in OPTIMIZATIONS.md section 3, the section the old
serve_endpoint.sh comment pointed at; its only other appearance is a
boilerplate export line in reproduce.sh. This config also routes the hot
all-reduces around RCCL entirely (custom AR owns decode sizes at 17us vs
NCCL 48us, quick-reduce INT4 owns >= 64 MB prefill messages), so the knob
has little left to act on. Keep it -- every measured run had it set, so
dropping it would decouple the numbers from the config -- but label it
honestly as inherited and unverified rather than measured.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The base image's aiter carries its own MiniMax-M3 table,
aiter/configs/model_configs/minimax_m3_fp4_tuned_fmoe.csv (32 rows, all at
model_dim=6144 expert=129, all FP4). Our tuned_fmoe_m3_gfx950.csv duplicated
11 of those shapes.

aiter merges configs/tuned_fmoe.csv with every model_configs/*tuned_fmoe*.csv
and keys duplicates on the untuned columns + cu_num + gfx + _tag. The 11
overlapping FP4 shapes therefore collide, and aiter aborts the FIRST start of
every fresh container with

  RuntimeError: Found 22 duplicate shape entries during merge of 'tuned_fmoe'.
  Auto-resolved by keeping best performing (lowest 'us') ... Please re-run.

after destructively rewriting our file 32 -> 21 rows. The second start then
works, so in Kubernetes this reads as a pod that crash-loops once and
"recovers" -- while silently running on a table 11 rows short.

Ship the de-conflicted table instead: the 16 FP8 rows plus the 5 FP4 shapes
the base image does not already cover. Verified 0 collisions against the base
table and 0 internal duplicates. This is byte-for-byte what aiter itself
resolves to, and the configuration the endpoint was validated on
(GSM8K-500 0.890, accept len 2.67-2.76).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0136r14rmgkS24vpRShiJ5FJ

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

amd documentation Improvements or additions to documentation jit-kernel memory-pool quant LLM Quantization

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants