Repository navigation
Feat/proxy model gfx1250 1023 - #6969
Draft
kkHuang-amd wants to merge 21 commits into
Draft
kkHuang-amd wants to merge 21 commits into
kkHuang-amd wants to merge 21 commits into
Conversation
…columns The fused dynamic MX quant + MoE sort HIP kernel (FUSED_MX_QUANT_MOE_SORT_KERNEL_DISPATCH in csrc/kernels/quant_kernels.cu) supports at most 32 * 256 = 8192 columns. With hidden_size = 16384 the first MoE GEMM's MXFP4 activation quant aborts during decode CUDA-graph capture: quant_kernels.cu:2020 input last dim has exceeded the maximum value 8192 The split path (per_1x32_mx_quant_hip + mxfp4_moe_sort_hip) supports up to 16384 columns and produces the same layout; it is already used above the token-count switch. Route group_size=32 inputs wider than 8192 columns to it. The split path costs about 1% of decode GPU time where it replaces the fused kernel. The fix comes from the Meta proxy MoE SGLang package (patches/aiter_fused_mx_quant_moe_sort_large_hidden.patch), observed on MI350X with amd-aiter 0.1.22.dev120, and reproduces on gfx1250 with amd-aiter 0.1.25.dev178. This version is restructured for black and is logically equivalent to that patch. (cherry picked from commit ee56e61)
….scaled_downcast The gfx1250 Gluon kernels behind dynamic_mxfp4_quant and dynamic_mxfp8_quant call gl.amd.cdna5.scaled_downcast. That API landed on Triton main in triton-lang/triton#11377 and is not in the release/3.8 branch, so with Triton 3.8 (e.g. 3.8.0+git669b31ac in rocm/sgl-dev:v0.5.22-rocm101-mi45x) every MXFP4 linear fails to compile: AttributeError: module 'triton.experimental.gluon.language.amd.cdna5' has no attribute 'scaled_downcast' The gfx950 path already guards the same API with _has_scaled_downcast() and falls back to the Triton kernel. Add the equivalent cdna5 check to both gfx1250 dispatches so they fall back too; backend="gluon" still raises with the missing requirement. Tested on gfx1250 with Triton 3.8.0+git669b31ac: the Triton fallback matches a PyTorch MXFP4 dequant reference for hidden size 16384, M = 1, 16 and 8192, and SGLang serves a Quark MXFP4 W4A4 MoE model. (cherry picked from commit 5dd80d3)
gfx1250 ships no normalization configs, so the first RMSNorm with M > 8192 and small N (e.g. a per-head q-norm over 6000 tokens x 16 heads) raises: FileNotFoundError: Required config file doesn't exist: .../configs/gfx1250/triton/normalization/rmsnorm_large_m_small_n/DEFAULT.json Add the file. gfx942 and gfx950 both use num_warps=8, but gfx1250 runs wave32, so the same value gives half the threads per block; measured on gfx1250 (forward, bf16, N=128), fewer warps is faster: num_warps M=96000 M=131072 2 14.6 us 18.3 us 4 15.1 us 18.4 us 8 15.7 us 19.1 us 16 20.1 us 25.0 us 32 31.9 us 41.2 us All settings match torch.nn.functional.rms_norm; num_stages 1-3 makes no measurable difference. Use num_warps=4, num_stages=2: within 3% of the best at both sizes. (cherry picked from commit d02a891)
Contributor
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
One backend per PR: PR title tags & labels: |
…na bf16x3 router) The default bandwidth_bound BLOCK_M128 config (M_LEQ_8192) hit GPU page faults in the proxy-MoE server when the router GEMM ran at prefill sizes. Use compute_bound 128x128x64 (2 buffers, 4 warps) for M >= 128; M <= 64 keeps DEFAULT. ~20% slower than the faulting config at M=8192, 5x faster than hipBLASLt for this fp32-output shape.
Buckets that spill registers to scratch hit GPU page faults on the scratch stores when dispatches run concurrently (caught with rocgdb in the proxy-MoE TP2 server: SIGSEGV right after scratch_store_b32 in bandwidth_bound 128x128x64 NB4 waves_per_eu=4, 22 VGPR spills, 92 B). - M_LEQ_256: 32x64x256 nb3 wpe4 (132-260 B scratch) -> 64x64x128 nb4 wpe0 - M_LEQ_2048: 128x128x128 nb4 (40 B) -> nb3 - M_LEQ_8192, any: 128x128x64 nb4 wpe4 (92 B) -> nb6 wpe1 (= M_LEQ_4096) All buckets now compile with 0 scratch. On qkv (N=9216 K=16384), o_proj (N=16384 K=8192) and the bf16x3 router (N=1536 K=16384): M=256 -34% / -37% / -8%, M=2048 -3%, M=8192 -4.6% to +1.1%. The N=1536 K=16384 compute_bound override is no longer needed: the router now uses the fixed DEFAULT.
…I size Since 2bafed4 the grouped a8w4 TDM launch passes rank-1 views (out.view(-1), the scales, and quant_scale / ep_row_map which default to out) through the cached FlyDSL launcher, whose C ABI packs a tensor's size as int32. A large prefill makes out = contiguous_m x N exceed 2**31 elements (e.g. 139264 x 16384 = 2,281,701,376) and the launch fails with "struct.error: 'i' format requires -2147483648 <= number <= 2147483647". The launcher reads these tensors only through their base pointer (fx.get_iter), so pass a rank-1 view clamped to 2**31 - 1 elements: same base address, same kernel, no recompile (rank and specialization key are unchanged).
…ace fix, EP-compact gather-reduce - grouped_moe_gfx1250: large-M tile rule for model_dim 16384 when no CSV row matches (inter 4096 / EP8 and inter 512 / EP1, token_num >= 2048; AITER_TDM_PREFILL_TILE_RULE=0 disables). - layout_align_m: expert segments align to max(tile_m, tile_m2); the GEMM with the smaller tile skips tiles that lie wholly in an expert's alignment gap. - mxfp4_preshuffle_gfx1250_tdm: retire LDS reads (s_wait_dscnt) before the stage-free barrier (rare wrong outputs when the compiler moved WMMAs past it); default-off experiment switches AITER_TDM_PW_OVERLAP / AITER_TDM_TILEMAP_LDS / AITER_TDM_SPLIT_FENCE. - moe_gather_reduce: EP compact mode for route maps with dropped routes (skip dropped slots, whole-row block for large T); EP8 0.578 -> 0.227 ms per layer. AITER_GR_EP_COMPACT=0 disables. Verified: op test --scenario verify, a8w4 --no-bias, EP1 (E512 top16 inter512) and EP8 (E64 top2 inter4096) at T=2/64/8192, all pass.
…M<=4, M<=16) Laguna TP8 decode shapes (qkv N=2304 K=16384, o N=16384 K=2048, lm_head N=25256 K=16384) used DEFAULT.json. Swept with the tuning harness (cold weights, CUDA graph); other M fall back to DEFAULT. us, DEFAULT -> tuned: qkv M2 22.4 -> 17.8; o M2 26.0 -> 14.4, M16 31.9 -> 16.4; lm_head M2 190 -> 136, M16 278 -> 159 (qkv M16: DEFAULT already best).
Conflicts in the gfx1250 grouped a8w4 TDM kernel (upstream: IGLP_PARAMS, a_preshuffled, padded/unpadded C store rework): - mxfp4_preshuffle_gfx1250_tdm.py: take main; re-apply layout_align_m (gap-tile skip for tile_m < layout rows) on top. The LDS race fix is already in main. Drop the default-off experiment switches AITER_TDM_PW_OVERLAP / AITER_TDM_TILEMAP_LDS / AITER_TDM_SPLIT_FENCE (no gain measured). - grouped_gemm_mxfp4.py: pass both a_preshuffled and layout_align_m. Verified on gfx1250: op test --scenario verify a8w4 --no-bias, EP1 (E512 top16 inter512, EP1 CSV v4) T=2/64/8192 and EP8 (E64 top2 inter4096) T=2/64/8192, all pass.
EP prefill was re-quantizing a token for every route. Behind AITER_FLYDSL_EP_TOKEN_MULTIDEST, quantize it once and give remote routes a zero-length descriptor so topk 16 stays cheap. The flag stays off by default. Co-authored-by: Cursor <cursoragent@cursor.com>
NOT READY FOR HUMAN REVIEWPR Self-ReviewRevision Waiting for the PR author. The readiness report will appear after the self-review answers are complete for this revision. Questions to answer in the PR description
Edit the Missing or invalid answers
Shadow mode reports what would block but keeps the status successful. Set |
…250-1023 # Conflicts: # aiter/ops/flydsl/grouped_gemm_mxfp4.py # aiter/ops/flydsl/kernels/mxfp4_preshuffle_gfx1250_tdm.py
… GEMM Laguna's shared expert gate/up is GEMM N=1024 K=16384. With no config it uses the 256x256x256 DEFAULT tile. Add a tuned table for the Gluon preshuffle kernel (gfx1250 only): - M<=31: 16x32x1024, M_LEQ_16 uses NUM_BUFFERS 2 - M_LEQ_32 and M_LEQ_64: 32x32x1024, 2 warps, 4 buffers - M_LEQ_128: 32x32x1024, 2 warps, 2 buffers - M_LEQ_256: 64x64x1024; M_LEQ_512: 128x64x512; M_LEQ_1024: 64x128x256 - M_LEQ_2048: 256x128x512; M_LEQ_4096, M_LEQ_8192, any: unchanged 256x256x256 Microbench (10/22, isolated kernel, DEFAULT vs tuned): M=32 48.6 -> 7.9 us, M=64 47.7 -> 8.4 us, M=128 65.8 -> 9.7 us, M=256 66.7 -> 11.1 us, M=512 53.1 -> 22.2 us, M=1024 72.2 -> 41.9 us. M=8192 is 164 -> 142 us and is not used. h13-16, proxy_12l TP8 EP1, bench_serving with the same server args on and off: decode gate/up 54.1 -> 7.9 us per launch, median ITL 1K1K c=64 14.71 -> 14.22 ms, c=32 11.18 -> 10.77 ms. E2E about 3% better. Prefill M=8192 is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
…x1250-1023 Keep both the gfx1250 m_align layout and upstream split-K k_waves on the grouped GEMM launches, and accept gfx1250-strict while still requiring cdna5.scaled_downcast before the Gluon quant kernel. Co-authored-by: Cursor <cursoragent@cursor.com>
…s for 96 CU Laguna TP16 runs three BF16 dense GEMMs per rank with no tuned entry on the 96-CU gfx1250 tree: qkv N=1280 K=16384, o_proj N=16384 K=1024 and lm_head N=12628 K=16384. Every M used the 96cu gemm_a16w16 DEFAULT tile. Add, for gfx1250 96 CU only: - 96cu Gluon gemm_a16w16 tables for the three shapes, with DEFAULT_FALLBACK, so untuned buckets keep DEFAULT.json. o_proj has all 13 buckets, lm_head M<=128 has 7, qkv has the 4 buckets where the best Gluon tile beats DEFAULT by more than 2%. - model_configs/laguna_gfx1250_tp16_bf16_tuned_gemm.csv, 17 rows at power-of-two M, where hipBLASLt or FlyDSL beats the best Gluon tile by more than 3%. The gl=1 padding maps every M of a bucket onto its row. qkv M<=16, 64 and 128 use FlyDSL; qkv M>=256 uses hipBLASLt; o_proj M=512, 1024 and 4096 use hipBLASLt; lm_head M=64 uses FlyDSL. Tuning (h13-16, GPU 0): one CUDA graph over weight copies above 1 GiB, median of rounds. 56 sampled Gluon tiles per bucket at the top M, then the best 6, DEFAULT, hipBLASLt and FlyDSL at three M in the bucket. tgemm.mm on the same bench, A B A B, old -> new us: qkv M=1 10.1 -> 8.2, M=64 18.9 -> 14.1, M=1024 118.7 -> 70.7, M=8192 577 -> 396. o_proj M=1 7.0 -> 5.6, M=32 13.0 -> 7.5, M=256 73.8 -> 26.3, M=1024 153.8 -> 64.6, M=8192 942 -> 466. lm_head M=8 67.8 -> 51.6, M=64 125.9 -> 89.9, M=128 168.3 -> 148.9. All 60 measured M are faster. Error against an fp32 reference is 1.7e-3 in both arms. E2E on h13-16: a 12-layer proxy at TP8 EP8 with 64 heads, 4 KV heads and shared intermediate 2048, so qkv and o_proj per rank match TP16. lm_head stays the TP8 shape there, so it is checked by the kernel bench only. bench_serving random-ids, A B A B, total tok/s: ISL 8192 OSL 1024 c=32 18110 -> 18799 (+3.8%), median TTFT 2275 -> 1942 ms, median ITL 11.70 -> 11.65 ms. ISL 1024 OSL 1024 c=32 5289 -> 5331 (+0.8%), median TTFT 339 -> 304 ms. The two A runs differ by 0.1%. Measured with the CSV rows passed through AITER_CONFIG_GEMM_BF16 on 56fab27. Co-authored-by: Cursor <cursoragent@cursor.com>
…er-of-two M The 96 CU lm_head shape N=12628 K=16384 had a single FlyDSL row at M=64. A FlyDSL plus hipBLASLt sweep on the same tuner writes the measured winner at each power-of-two M: Gluon through M=8, hipBLASLt from M=16 to 8192. Co-authored-by: Cursor <cursoragent@cursor.com>
…unings per M Two TP16 BF16 GEMM tunings existed for the 96 CU gfx1250 tree: this file (88834dd) and a separate aiter tuner run on h13-14 (qkv and o_proj only, hipBLASLt / FlyDSL / torch). Each wins different M, so measure all three dispatch paths on the same GPU and keep the fastest per (N, K, M). Op bench (h13-14, GPU 0, this tree + the Gluon tables of df929c0): one process per AITER_CONFIG_GEMM_BF16 arm (this file / the h13-14 CSV / header only = Gluon tables only), tgemm.mm at every power-of-two M 1..8192, 20 calls in one CUDA graph with weights rotated over more than 1 GiB, median of 7 replays, two rounds in opposite order. A row is kept from the faster arm; no row where the Gluon path is fastest. Changes vs 88834dd (us, this file -> merged): - o_proj N=16384 K=1024: hipBLASLt for M 256..8192, FlyDSL at M 64. M=256 26.9 -> 14.9, M=1024 63.8 -> 46.5, M=8192 455.5 -> 314.5. M <= 32 stays on the Gluon tables (fastest). - qkv N=1280 K=16384: hipBLASLt for M 256, 512, 2048, 4096 (was torch), FlyDSL rows at M 4, 32, 64, 128. M=64 13.7 -> 11.1, M=128 20.3 -> 17.3, M=512 48.3 -> 41.6. torch kept at M 1024 and 8192. - lm_head N=12628 K=16384: unchanged from M 4 up; the M 1 and 2 triton rows are dropped (Gluon path equal within 0.4%). E2E, TP16 across two DPX nodes (24-layer proxy, EP4, carv2), random-ids, TTT/GPU this file -> merged: 70000/300 c1 2800.6 -> 2829.6, TTFT 2044 -> 1998 ms; 8192/1024 c16 1178.9 -> 1186.2, TTFT 1881 -> 1828 ms; 1024/1024 c16 330.9 -> 332.8. TTFT -0.9..-3.2% at all 5 points measured, ITL equal. Single run per point (~3% noise). hipBLASLt solidx values come from image rocm/sgl-dev:v0.5.23-rocm101-mi45x-20261107_8066cd93b__test; retune if the hipBLASLt version differs.
Laguna at TP16 with --ep-size 4 runs MoE as EP4 x MoE-TP4: 128 local experts, inter_dim 1024, fp8 activations x MXFP4 weights. The EP lookup key is (token, 16384, 1024, experts=128, topk=-1) and had no row, so every bucket used the kernel default tile (64 x 256 x 128, 2 buffers, both stages). Add 12 rows for gfx1250 96 CU (DPX partition). Tokens 1024 and 2048 keep the default (it was fastest there). Tuning (h13-14, one DPX partition, other partitions idle): op_tests/flydsl_tests/test_flydsl_grouped_gemm.py a8w4 silu --no-bias, E=128 topk=4 so the per-expert M (T*4/128 = T/32) matches the EP4 rank, whole fused MoE in a CUDA graph. Stage-1 and stage-2 grids, then the finalists in 3 interleaved repeats (median, spread <= 3%). Partitions of one physical GPU running at the same time share power and add up to 20% noise, so the final numbers come from a single busy partition. us per call, default -> row: T=1 56.9 -> 29.5, T=2 70.0 -> 39.4, T=4 86.1 -> 61.7, T=8 117.0 -> 102.0, T=16 220.1 -> 180.1, T=32 304.5 -> 277.0, T=64 397.1 -> 339.9, T=128 441.8 -> 395.3, T=256 463.9 -> 423.7, T=512 487.2 -> 457.5, T=4096 1307.8 -> 1230.0, T=8192 2395.5 -> 2133.4. All rows pass the op test's accuracy gate. E2E, TP16 across two DPX nodes (24-layer proxy, EP4, carv2, merged BF16 GEMM rows of 335b817), random-ids, TTT per GPU without -> with the rows: 70000/300 c1 2829.6 -> 3069.7 (ITL 5.31 -> 4.64 ms), 8192/1024 c2 384.1 -> 428.2, c16 1186.2 -> 1285.7, 1024/1024 c2 93.6 -> 106.1, c16 332.8 -> 366.4. ITL -9..-13% at all 5 points, TTFT -3..-4% on 8192 and 70000. Single run per point (~3% noise). The rows are merged into the default config only when AITER_CONFIG_GROUPED_FMOE is unset (model_configs merge); checked with update_config_files + _find_grouped_config on this tree.
…de M on 96 CU Adds 16 flydsl rows (M 1..128) for the TP16 per-rank shared-expert shapes of the Laguna proxy model on gfx1250 (MXFP4 dequantized to bf16): gate_up N=512 K=16384 and down N=16384 K=256. Tuned with gemm_a16w16_tune.py --with-hipblaslt; only the flydsl winners up to M=128 are kept (the hipBLASLt winners at M >= 256 faulted at other M mapped to the same padded row, e.g. a row tuned at M=2048 at M=1536). Used by sglang SGLANG_OPT_QUARK_BF16_TGEMM=1 (small M only): 15-17 -> 5.3 us (gate_up) and 11 -> 3.1 us (down) per call at M=2 vs F.linear; TP16 two-node 1k/1k c2 TTT/GPU +15.1%, 70000/300 c1 +7.0%.
…l M=8192 Exact-M hipBLASLt solidx 244 for gate_up 8192x512x16384 (176.8 us) and down 8192x16384x256 (147.2 us). h13-15 ran them at 173/148 us solo and 206-209/165-168 us with both partitions of one card, with the same max abs as F.linear.
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist