Repository navigation
[Feat] [MI355X] glm53_flash: Serve MTP verify steps on the monokernel - #71
Conversation
|
@Jacob0226 Can you share the accuracy test results |
|
@HaiShaw
Kernel tests: all 9 tests in |
jhinpan
left a comment
There was a problem hiding this comment.
We ran this PR on upstream SGLang, without the extra AMD PRs, and found no correctness problem. Details and requests below.
Setup: MI355X x8; lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261024 with SGLang python/sglang at 5044779f (has sgl-project/sglang#46036, none of #41615, #41871, #42020, #38764, #44415, #44416, #44417); main (a901b72) + this PR; amd/GLM-5.3-Flash-Quark-MXFP4, TP4, bf16 KV, tilelang DSA, MTP3 (EAGLE, 3 steps, topk 1, 4 draft tokens). Off and on run side by side (HIP 0–3 and 4–7) with the same flags; on adds --monokernel-backend amd. The on server logs runs target_verify steps of batch size [1, 2] x 4 tokens.
Tests
test_glm53_flash_kernels.py: 9 passed,test_a_verify_launch_matches_its_draft_tokens_decode_launchesincluded.test_glm53_flash_server.py: 4 passed.TestMonokernelFaultfails, and it fails onmaintoo.- AIME 2025 + 2026 with sgl-eval 0.1.3: 60 problems, 4 samples each, temperature 1.0, top_p 0.95, the default reasoning effort, up to 128K tokens. Two requests at a time, so every verify step has at most 8 rows and runs on the monokernel.
| Off | On | |
|---|---|---|
| Accuracy | 0.958 (230 / 240) | 0.933 (224 / 240) |
| Truncated at 128K | 3 | 3 |
| Accept length (mean of 17,505 / 16,286 log intervals) | 2.757 | 2.764 |
| Output tokens/s, both arms on one node (indicative only) | 504 / 477 | 800 / 738 |
- Paired by problem, on − off is −2.5 points (95% bootstrap interval −6.7 to +1.3), not significant.
- Each wrong answer of the on arm is a coherent solution. Three of them have the right value but no
\boxed{}; the off arm has two such. - The accept length is the same on both arms, so the verify launches agree with SGLang's verify.
Requests
- Speed table. Only the c1 MTP5 row is an A/B on one node. At c4, off ran on another node a day earlier without #46036, the two on runs shared a node, and "+32%" compares MTP3 on with MTP5 off. Please re-measure c4 as a same-configuration A/B on one node, or remove those rows.
- Test stack. The PR was tested with six open AMD PRs and one closed one. Please say whether this PR needs any of them. Our run above suggests it does not.
- Fused accept. A server with
accept_lens_poolset fails at its first verify step. If SGLang exposes this as a server setting, please refuse it at startup with the other config reasons. - One change per PR. The iscore store fix is a decode fix (a fault past 262,144 tokens) that does not need MTP. Please move it to its own PR.
- E2E. Please add an MTP case to
test_glm53_flash_server.py, asTestMonokernelMtpintest_glm53_server.pydoes (GSM8K and an accept-length bound).
Heads-up, not caused by this PR: sgl-project/sglang#46980 removes HybridReqToTokenPool.mamba2_layer_cache, which _bind_kda calls (on main too). On SGLang with that change, for example lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028, the GLM-5.3-Flash monokernel fails at CUDA-graph capture with an AttributeError. #94 reads the cache from the linear-attention backend's state_layer_cache(layer_id). This PR will need a rebase on it.
|
@Jacob0226 pls solve conflict. |
6904f3c to
7f3ccfd
Compare
|
Thanks for the run on upstream SGLang. All five requests are addressed at
Also rebased on #94: |
There was a problem hiding this comment.
The five earlier requests are done. We found no accuracy loss. The decode speed-up is real, but two numbers in the speed table do not reproduce. Four requests are below.
Tested: this PR merged with #97, 8 × MI355X, lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028 (SGLang 5d8e90a0f9), TP4, the flags of test/e2e/glm53_flash_server.py.
Correctness (off → on)
| Check | Result |
|---|---|
| AIME25 + AIME26, MTP3, 2 requests at a time, 3 runs, 720 samples an arm | 0.943 → 0.932, paired −1.11 points (95 % CI −4.44 to +1.94) |
| Verify-versus-prefill logprob KL, MTP3, verify on the kernels | 1 request at a time: 0.0185 → 0.0135; 2 at a time: 0.0755 → 0.0366 |
| RULER2 NIAH, MTP3, 300K / 640K / 900K | 1.000 / 0.964 / 0.963 → 1.000 / 0.964 / 0.990 |
| Few-shot GSM8K, MTP5, 1 request at a time | 0.914 → 0.920 |
The kernel, unit and e2e tests pass on the merge. TestMonokernelFault fails, but it also fails on main. We checked the SGLang side at 5d8e90a0f9. The ring width and slot, the scratch row, the page-table rows and the radix copy agree with the kernel.
Speed: AgentX, 900-s windows, golden acceptance (MTP5 3.61, MTP3 2.99). Both arms ran at the same time on different GPU sets. Round 2 swapped the sets. On / off:
| Point | tok/s/GPU | P90 1/ITL | P90 E2E-norm. interactivity |
|---|---|---|---|
| c1 MTP5 | +4.0 % | +23.4 % | −0.8 % |
| c4 MTP3 | +20.5 % | +5.4 % | +12.3 % |
| c16 MTP3 (control, 1 round) | +3.4 % | −1.1 % | −6.0 % |
At c1, the on arm finished more turns in the window (125-126 vs 110). On the turns that both arms finished, P90 E2E-norm. interactivity is about +8 % at c1 and +14 % at c4.
Requests
- Speed table: two gains in the body do not reproduce. c1 throughput is +18 % in the body and +4.0 % here. c4 P90 1/ITL is +35 % in the body and +5.4 % here. Give the server command, the AgentX client version and the aggregate JSONs. Also add P90 E2E-normalized interactivity, P90 TTFT and a point past the largest width, as the template asks.
TestMonokernelMtpdoes not protect the kernel path. At 12 requests, only about 14 % of the verify work runs on the kernels. If the kernel verify accepts nothing, the mean accept length is still about 2.9, above the 2.5 bound. Run this case at 2 requests, so that every verify step runs on the kernels.- In
_Built.forward, readbackend.linear_attn_backend.accept_lens_pooldirectly. SGLang declares it with the defaultNone, sogetattr(..., None)is not necessary. - Rebase after #97 merges. The only conflict is in
test_glm53_flash_kernels.py.
|
Thanks for the rerun. The four requests, on
Aggregate JSONs: concurrency 1, 4 and 16 (control), off and onConcurrency 1, MTP5, off{
"hw": "mi355x",
"conc": 1,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 287,
"num_requests_successful": 275,
"request_accounting": {
"records_total": 287,
"records_profiled": 275,
"records_dropped_total": 12,
"records_warmup_dropped": 11,
"records_error_dropped": 12,
"error_categories": {
"InvalidInferenceResultError": 12
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3586,
"mean": 0.07641,
"p50": 0.0,
"p75": 0.0,
"p90": 0.0,
"p95": 1,
"std": 0.27188
},
"latency": {
"ttft": {
"mean": 0.2896,
"p50": 0.22105,
"p75": 0.27537,
"p90": 0.39305,
"p95": 0.76438,
"std": 0.27957
},
"e2el": {
"mean": 6.44179,
"p50": 2.92342,
"p75": 8.78648,
"p90": 16.89476,
"p95": 21.26944,
"std": 8.1514
},
"itl": {
"mean": 0.00312,
"p50": 0.0031,
"p75": 0.00319,
"p90": 0.00327,
"p95": 0.00334,
"std": 0.00037
},
"tpot": {
"mean": 0.00312,
"p50": 0.0031,
"p75": 0.00319,
"p90": 0.00327,
"p95": 0.00334,
"std": 0.00037
},
"intvty": {
"mean": 320.73237,
"p50": 322.22571,
"p75": 313.26841,
"p90": 305.54747,
"p95": 299.42202,
"std": 35.98677
},
"e2e_norm_intvty": {
"mean": 237.08256,
"p50": 293.7026,
"p75": 252.75891,
"p90": 193.13917,
"p95": 148.49991,
"std": 60.41264
},
"full_response_itl": {
"mean": 0.00312,
"p50": 0.0031,
"p75": 0.00319,
"p90": 0.00327,
"p95": 0.00334,
"std": 0.00037
},
"full_response_intvty": {
"mean": 320.57678,
"p50": 322.19438,
"p75": 313.22876,
"p90": 305.50454,
"p95": 299.27256,
"std": 35.16414
}
},
"tokens": {
"input": {
"mean": 277301.85818,
"p50": 313689,
"p75": 396036,
"p90": 499689,
"p95": 530162.6,
"std": 162332.58872
},
"output_actual": {
"mean": 1967.56364,
"p50": 801,
"p75": 2715.5,
"p90": 5181.6,
"p95": 6905.1,
"std": 2618.80655
},
"output_expected": {
"mean": 2164.80364,
"p50": 945,
"p75": 3034.5,
"p90": 5559.8,
"p95": 6915.7,
"std": 2883.51305
}
},
"throughput": {
"input": {
"tokens_per_second": 21233.85997
},
"output": {
"tokens_per_second": 150.66243
},
"total": {
"tokens_per_second": 21384.5224
},
"duration_seconds": 3591.34002,
"per_gpu": {
"total_tput_tps": 5346.1306,
"output_tput_tps": 37.66561,
"input_tput_tps": 5308.46499
}
},
"cache": {
"theoretical_cache_hit_rate": 0.98172
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 83,
"cache": {
"gpu_cache_hit_rate": 0.97997,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.97997,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 75755008
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.04,
"gpu_total_tokens": 12586752,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 77303580,
"generation_total": 541029,
"requests_completed": 275,
"prompt_by_source": {
"gpu_cache_hit": 75755008,
"cpu_or_external_cache_hit": null,
"computed": 1884902,
"raw": {
"device": 75755008
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 12586752
}Concurrency 1, MTP5, on{
"hw": "mi355x",
"conc": 1,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 312,
"num_requests_successful": 300,
"request_accounting": {
"records_total": 312,
"records_profiled": 300,
"records_dropped_total": 12,
"records_warmup_dropped": 11,
"records_error_dropped": 12,
"error_categories": {
"InvalidInferenceResultError": 12
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3591,
"mean": 0.08326,
"p50": 0.0,
"p75": 0.0,
"p90": 0.0,
"p95": 1,
"std": 0.28127
},
"latency": {
"ttft": {
"mean": 0.30621,
"p50": 0.23626,
"p75": 0.28379,
"p90": 0.42688,
"p95": 0.77708,
"std": 0.31677
},
"e2el": {
"mean": 5.02879,
"p50": 2.5935,
"p75": 6.79596,
"p90": 12.68529,
"p95": 16.3698,
"std": 6.04598
},
"itl": {
"mean": 0.00236,
"p50": 0.00234,
"p75": 0.00246,
"p90": 0.00257,
"p95": 0.00328,
"std": 0.00042
},
"tpot": {
"mean": 0.00236,
"p50": 0.00234,
"p75": 0.00246,
"p90": 0.00257,
"p95": 0.00328,
"std": 0.00042
},
"intvty": {
"mean": 423.40405,
"p50": 427.96221,
"p75": 406.5724,
"p90": 389.33113,
"p95": 305.08097,
"std": 62.46711
},
"e2e_norm_intvty": {
"mean": 291.32563,
"p50": 382.81547,
"p75": 298.25413,
"p90": 222.7256,
"p95": 168.86507,
"std": 92.65684
},
"full_response_itl": {
"mean": 0.00236,
"p50": 0.00234,
"p75": 0.00246,
"p90": 0.00257,
"p95": 0.00328,
"std": 0.00042
},
"full_response_intvty": {
"mean": 423.14494,
"p50": 427.91301,
"p75": 406.54593,
"p90": 389.21219,
"p95": 305.03044,
"std": 61.81398
}
},
"tokens": {
"input": {
"mean": 301246.48667,
"p50": 336938,
"p75": 438104.5,
"p90": 546812.5,
"p95": 575201.25,
"std": 176555.21982
},
"output_actual": {
"mean": 1993.89667,
"p50": 861.5,
"p75": 2763.25,
"p90": 5108.6,
"p95": 6903.35,
"std": 2568.29819
},
"output_expected": {
"mean": 2122.73,
"p50": 1031.5,
"p75": 2906.5,
"p90": 5422.7,
"p95": 6903.35,
"std": 2788.08961
}
},
"throughput": {
"input": {
"tokens_per_second": 25140.37384
},
"output": {
"tokens_per_second": 166.39964
},
"total": {
"tokens_per_second": 25306.77348
},
"duration_seconds": 3594.77335,
"per_gpu": {
"total_tput_tps": 6326.69337,
"output_tput_tps": 41.59991,
"input_tput_tps": 6285.09346
}
},
"cache": {
"theoretical_cache_hit_rate": 0.98266
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 83,
"cache": {
"gpu_cache_hit_rate": 0.98108,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.98108,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 89363200
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.05,
"gpu_total_tokens": 12577280,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 91086850,
"generation_total": 598130,
"requests_completed": 300,
"prompt_by_source": {
"gpu_cache_hit": 89363200,
"cpu_or_external_cache_hit": null,
"computed": 2062173,
"raw": {
"device": 89363200
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 12577280
}Concurrency 4, MTP3, off{
"hw": "mi355x",
"conc": 4,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 658,
"num_requests_successful": 612,
"request_accounting": {
"records_total": 658,
"records_profiled": 612,
"records_dropped_total": 46,
"records_warmup_dropped": 44,
"records_error_dropped": 42,
"error_categories": {
"InvalidInferenceResultError": 42
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3590,
"mean": 0.17019,
"p50": 0.0,
"p75": 0.0,
"p90": 1,
"p95": 1,
"std": 0.40916
},
"latency": {
"ttft": {
"mean": 0.26828,
"p50": 0.21788,
"p75": 0.28942,
"p90": 0.39464,
"p95": 0.59738,
"std": 0.298
},
"e2el": {
"mean": 4.7896,
"p50": 1.62028,
"p75": 5.77956,
"p90": 13.62804,
"p95": 19.66268,
"std": 7.67236
},
"itl": {
"mean": 0.00349,
"p50": 0.00337,
"p75": 0.00361,
"p90": 0.00389,
"p95": 0.00408,
"std": 0.00129
},
"tpot": {
"mean": 0.00349,
"p50": 0.00337,
"p75": 0.00361,
"p90": 0.00389,
"p95": 0.00408,
"std": 0.00129
},
"intvty": {
"mean": 286.49101,
"p50": 296.91865,
"p75": 276.68714,
"p90": 257.10752,
"p95": 245.13839,
"std": 37.38423
},
"e2e_norm_intvty": {
"mean": 207.07109,
"p50": 248.18072,
"p75": 214.74548,
"p90": 170.25679,
"p95": 135.67705,
"std": 54.66942
},
"full_response_itl": {
"mean": 0.00349,
"p50": 0.00337,
"p75": 0.00361,
"p90": 0.00389,
"p95": 0.00408,
"std": 0.00129
},
"full_response_intvty": {
"mean": 286.78994,
"p50": 296.95611,
"p75": 276.86314,
"p90": 256.99931,
"p95": 245.08109,
"std": 185.74593
}
},
"tokens": {
"input": {
"mean": 183447.80392,
"p50": 114831.5,
"p75": 246051.75,
"p90": 521171.3,
"p95": 580929.7,
"std": 178755.09776
},
"output_actual": {
"mean": 1308.64379,
"p50": 388.5,
"p75": 1549.75,
"p90": 3843.6,
"p95": 5418.35,
"std": 2200.57147
},
"output_expected": {
"mean": 1869.11601,
"p50": 785.5,
"p75": 2719,
"p90": 5299.4,
"p95": 6316,
"std": 2624.25191
}
},
"throughput": {
"input": {
"tokens_per_second": 31193.73498
},
"output": {
"tokens_per_second": 222.52372
},
"total": {
"tokens_per_second": 31416.2587
},
"duration_seconds": 3599.12194,
"per_gpu": {
"total_tput_tps": 7854.06468,
"output_tput_tps": 55.63093,
"input_tput_tps": 7798.43375
}
},
"cache": {
"theoretical_cache_hit_rate": 0.97855
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 83,
"cache": {
"gpu_cache_hit_rate": 0.97614,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.97614,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 111088640
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.09,
"gpu_total_tokens": 11859456,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 113803766,
"generation_total": 800765,
"requests_completed": 612,
"prompt_by_source": {
"gpu_cache_hit": 111088640,
"cpu_or_external_cache_hit": null,
"computed": 3850084,
"raw": {
"device": 111088640
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 11859456
}Concurrency 4, MTP3, on{
"hw": "mi355x",
"conc": 4,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 743,
"num_requests_successful": 697,
"request_accounting": {
"records_total": 743,
"records_profiled": 697,
"records_dropped_total": 46,
"records_warmup_dropped": 44,
"records_error_dropped": 38,
"error_categories": {
"InvalidInferenceResultError": 38
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3588,
"mean": 0.19398,
"p50": 0.0,
"p75": 0.0,
"p90": 1,
"p95": 1,
"std": 0.43438
},
"latency": {
"ttft": {
"mean": 0.27438,
"p50": 0.21694,
"p75": 0.29031,
"p90": 0.39084,
"p95": 0.63419,
"std": 0.31437
},
"e2el": {
"mean": 3.25199,
"p50": 1.24402,
"p75": 3.72078,
"p90": 8.59566,
"p95": 12.11673,
"std": 5.45306
},
"itl": {
"mean": 0.00225,
"p50": 0.00201,
"p75": 0.00236,
"p90": 0.00289,
"p95": 0.00383,
"std": 0.00134
},
"tpot": {
"mean": 0.00225,
"p50": 0.00201,
"p75": 0.00236,
"p90": 0.00289,
"p95": 0.00383,
"std": 0.00134
},
"intvty": {
"mean": 444.49344,
"p50": 498.67062,
"p75": 423.48971,
"p90": 345.8268,
"p95": 261.20668,
"std": 128.03208
},
"e2e_norm_intvty": {
"mean": 269.81446,
"p50": 369.1254,
"p75": 283.30997,
"p90": 199.86559,
"p95": 156.63219,
"std": 115.95015
},
"full_response_itl": {
"mean": 0.00225,
"p50": 0.00201,
"p75": 0.00236,
"p90": 0.00289,
"p95": 0.00383,
"std": 0.00134
},
"full_response_intvty": {
"mean": 445.34782,
"p50": 498.61395,
"p75": 423.82939,
"p90": 346.08897,
"p95": 261.1089,
"std": 266.3816
}
},
"tokens": {
"input": {
"mean": 186892.60258,
"p50": 117893,
"p75": 248646,
"p90": 528988,
"p95": 600750.8,
"std": 181994.68434
},
"output_actual": {
"mean": 1310.5165,
"p50": 439,
"p75": 1505,
"p90": 3785.4,
"p95": 5417.6,
"std": 2258.88879
},
"output_expected": {
"mean": 1802.12769,
"p50": 722,
"p75": 2596,
"p90": 4849.2,
"p95": 6316,
"std": 2631.99162
}
},
"throughput": {
"input": {
"tokens_per_second": 36250.2359
},
"output": {
"tokens_per_second": 254.19161
},
"total": {
"tokens_per_second": 36504.42751
},
"duration_seconds": 3593.47024,
"per_gpu": {
"total_tput_tps": 9126.10688,
"output_tput_tps": 63.5479,
"input_tput_tps": 9062.55897
}
},
"cache": {
"theoretical_cache_hit_rate": 0.9788
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 83,
"cache": {
"gpu_cache_hit_rate": 0.9758,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.9758,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 128559872
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.09,
"gpu_total_tokens": 11852032,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 131748180,
"generation_total": 913295,
"requests_completed": 697,
"prompt_by_source": {
"gpu_cache_hit": 128559872,
"cpu_or_external_cache_hit": null,
"computed": 4328140,
"raw": {
"device": 128559872
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 11852032
}Concurrency 16, MTP3, off (control){
"hw": "mi355x",
"conc": 16,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 2712,
"num_requests_successful": 2533,
"request_accounting": {
"records_total": 2712,
"records_profiled": 2533,
"records_dropped_total": 179,
"records_warmup_dropped": 177,
"records_error_dropped": 128,
"error_categories": {
"InvalidInferenceResultError": 128
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3593,
"mean": 0.7047,
"p50": 0.0,
"p75": 1,
"p90": 2,
"p95": 2,
"std": 0.9036
},
"latency": {
"ttft": {
"mean": 0.3718,
"p50": 0.20113,
"p75": 0.35438,
"p90": 0.72349,
"p95": 0.95393,
"std": 0.58083
},
"e2el": {
"mean": 5.08238,
"p50": 2.12374,
"p75": 4.8533,
"p90": 11.54225,
"p95": 20.2753,
"std": 9.41906
},
"itl": {
"mean": 0.00514,
"p50": 0.00457,
"p75": 0.0054,
"p90": 0.0065,
"p95": 0.00739,
"std": 0.00603
},
"tpot": {
"mean": 0.00514,
"p50": 0.00457,
"p75": 0.0054,
"p90": 0.0065,
"p95": 0.00739,
"std": 0.00603
},
"intvty": {
"mean": 194.71658,
"p50": 218.65511,
"p75": 185.07987,
"p90": 153.86716,
"p95": 135.38621,
"std": 50.45874
},
"e2e_norm_intvty": {
"mean": 144.88623,
"p50": 180.9942,
"p75": 144.81149,
"p90": 110.85913,
"p95": 84.17088,
"std": 52.58548
},
"full_response_itl": {
"mean": 0.00513,
"p50": 0.00457,
"p75": 0.0054,
"p90": 0.0065,
"p95": 0.00739,
"std": 0.00603
},
"full_response_intvty": {
"mean": 194.8846,
"p50": 218.64682,
"p75": 185.15376,
"p90": 153.89042,
"p95": 135.37494,
"std": 2828.20829
}
},
"tokens": {
"input": {
"mean": 152882.1512,
"p50": 108487,
"p75": 178998,
"p90": 278434,
"p95": 524862.8,
"std": 152468.38471
},
"output_actual": {
"mean": 974.77142,
"p50": 354,
"p75": 915,
"p90": 2377.6,
"p95": 4164.2,
"std": 1936.01066
},
"output_expected": {
"mean": 1585.52564,
"p50": 695,
"p75": 1997.5,
"p90": 4034.5,
"p95": 5496.95,
"std": 2463.83183
}
},
"throughput": {
"input": {
"tokens_per_second": 107710.27045
},
"output": {
"tokens_per_second": 686.75704
},
"total": {
"tokens_per_second": 108397.02749
},
"duration_seconds": 3595.29771,
"per_gpu": {
"total_tput_tps": 27099.25687,
"output_tput_tps": 171.68926,
"input_tput_tps": 26927.56761
}
},
"cache": {
"theoretical_cache_hit_rate": 0.97646
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 85,
"cache": {
"gpu_cache_hit_rate": 0.97141,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.97141,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 377209088
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.23,
"gpu_total_tokens": 11580160,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 388311422,
"generation_total": 2468692,
"requests_completed": 2533,
"prompt_by_source": {
"gpu_cache_hit": 377209088,
"cpu_or_external_cache_hit": null,
"computed": 17337739,
"raw": {
"device": 377209088
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 11580160
}Concurrency 16, MTP3, on (control){
"hw": "mi355x",
"conc": 16,
"image": "lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028",
"recipe_fingerprint": "",
"model": "amd/GLM-5.3-Flash-Quark-MXFP4",
"infmax_model_prefix": "glm5.3flash",
"framework": "sglang",
"precision": "fp4",
"spec_decoding": "mtp",
"disagg": false,
"scenario_type": "agentic-coding",
"is_multinode": false,
"num_gpus": 4,
"tp": 4,
"ep": 1,
"dp_attention": "false",
"kv_offloading": "none",
"kv_offload_backend": null,
"allocated_cpu_dram_gb": 0,
"num_requests_total": 2745,
"num_requests_successful": 2566,
"request_accounting": {
"records_total": 2745,
"records_profiled": 2566,
"records_dropped_total": 179,
"records_warmup_dropped": 177,
"records_error_dropped": 127,
"error_categories": {
"InvalidInferenceResultError": 127
}
},
"pp": 1,
"dcp_size": 1,
"pcp_size": 1,
"dataset": {
"source_type": "public_dataset",
"loader": "semianalysis_cc_traces_weka_062126",
"hf_dataset_name": "semianalysisai/cc-traces-weka-062126",
"hf_split": "train",
"num_dataset_entries": 393
},
"request_metrics": {
"qps": {
"window_seconds": 1,
"samples": 3600,
"mean": 0.7125,
"p50": 0.0,
"p75": 1,
"p90": 2,
"p95": 2,
"std": 0.92398
},
"latency": {
"ttft": {
"mean": 0.37323,
"p50": 0.19491,
"p75": 0.34792,
"p90": 0.69994,
"p95": 0.99692,
"std": 0.60507
},
"e2el": {
"mean": 4.6108,
"p50": 1.91142,
"p75": 4.38392,
"p90": 10.72156,
"p95": 18.47175,
"std": 8.5904
},
"itl": {
"mean": 0.00468,
"p50": 0.00437,
"p75": 0.00533,
"p90": 0.00653,
"p95": 0.00756,
"std": 0.00589
},
"tpot": {
"mean": 0.00468,
"p50": 0.00437,
"p75": 0.00533,
"p90": 0.00653,
"p95": 0.00756,
"std": 0.00589
},
"intvty": {
"mean": 213.90309,
"p50": 228.67233,
"p75": 187.63087,
"p90": 153.15686,
"p95": 132.33443,
"std": 135.2905
},
"e2e_norm_intvty": {
"mean": 153.25293,
"p50": 189.34642,
"p75": 146.87591,
"p90": 110.63011,
"p95": 83.33433,
"std": 97.31638
},
"full_response_itl": {
"mean": 0.00467,
"p50": 0.00437,
"p75": 0.00533,
"p90": 0.00653,
"p95": 0.00756,
"std": 0.00589
},
"full_response_intvty": {
"mean": 214.16044,
"p50": 228.63203,
"p75": 187.56532,
"p90": 153.22161,
"p95": 132.33428,
"std": 213.71118
}
},
"tokens": {
"input": {
"mean": 153249.65082,
"p50": 108781.5,
"p75": 181110.5,
"p90": 278159,
"p95": 520922.25,
"std": 151889.8421
},
"output_actual": {
"mean": 977.82307,
"p50": 354,
"p75": 921.5,
"p90": 2385,
"p95": 4165.25,
"std": 1936.41153
},
"output_expected": {
"mean": 1581.42472,
"p50": 693,
"p75": 1982,
"p90": 4033,
"p95": 5515.5,
"std": 2457.8202
}
},
"throughput": {
"input": {
"tokens_per_second": 109206.70879
},
"output": {
"tokens_per_second": 696.80315
},
"total": {
"tokens_per_second": 109903.51194
},
"duration_seconds": 3600.8649,
"per_gpu": {
"total_tput_tps": 27475.87799,
"output_tput_tps": 174.20079,
"input_tput_tps": 27301.6772
}
},
"cache": {
"theoretical_cache_hit_rate": 0.97642
}
},
"server_metrics": {
"present": true,
"adapter": "sglang",
"metric_count": 85,
"cache": {
"gpu_cache_hit_rate": 0.97133,
"cpu_cache_hit_rate": null,
"external_cache_hit_rate": null,
"overall_cache_hit_rate": 0.97133,
"prefix_cache_hits": null,
"prefix_cache_queries": null,
"external_prefix_cache_hits": null,
"external_prefix_cache_queries": null,
"cached_tokens_by_source": {
"device": 382499072
},
"frontend_cache_hit_rate": null,
"router_kv_hit_rate": null,
"router_shared_cache_hit_rate": null,
"frontend_cached_tokens": null,
"frontend_input_tokens": null
},
"kv_cache": {
"gpu_usage_pct": 0.24,
"gpu_total_tokens": 11572736,
"cpu_usage_pct": null,
"cpu_used_tokens": null,
"cpu_total_tokens": null
},
"kv_offload": {
"bytes_gpu_to_cpu": null,
"bytes_cpu_to_gpu": null,
"time_gpu_to_cpu": null,
"time_cpu_to_gpu": null,
"bandwidth_gpu_to_cpu_bytes_per_second": null,
"bandwidth_cpu_to_gpu_bytes_per_second": null
},
"tokens": {
"prompt_total": 393789802,
"generation_total": 2508708,
"requests_completed": 2566,
"prompt_by_source": {
"gpu_cache_hit": 382499072,
"cpu_or_external_cache_hit": null,
"computed": 17596575,
"raw": {
"device": 382499072
}
}
},
"sources": []
},
"kv_cache_pool_tokens": 11572736
} |
… monokernel KDA, DSA, dense, embed and head launches build at widths 12 and 16 (spec WIDTHS 2/4/8/12/16). Above 8 rows the MoE runs expert-major (FlashBuild.grouped): each expert's weights are read once for all the rows that pick it. Steps of 8 rows or fewer keep sgl-project#71's per-row MoE. Thread guards that assumed one pass over 8 rows now loop in passes. Co-authored-by: Cursor <cursoragent@cursor.com>
|
@jhinpan #98 (verify steps of up to 16 rows) and #100 (index tiles scored once per verify request) are ready for review. Both carry this PR and #97 at their current heads; each body names the commits to review. After #97 and this PR merge, I will rebase them onto Could you also approve the workflow runs for the fork? Lint shows |
…ast 262,144 tokens (#97) ## Motivation Past 262,144 tokens, a GLM-5.3-Flash request breaks the monokernel's sparse attention: the launch faults, or silently returns wrong tokens (RULER NIAH at 300K: 0.673 on `main`, 1.000 here). Only there does the score-tile loop run. Since #22 peeled its first tile, FlyDSL's lowering merges the two identical branched index-score stores, giving the loop two entries ([FlyDSL#1358](ROCm/FlyDSL#1358)). The loop's MFMAs then run with the storing lanes off and, writing all 64 lanes, overwrite the store's index. Now every lane stores; masked lanes store past `num_records`, which the buffer drops. Split out of #71; merge this first: both change `kernel.py` and `check_layer.py`. ## Modifications - **Kernel** (`gfx950/kernel.py`): the index-score store in the DSA scoring loop is unconditional. A masked lane stores at `S * MAXP`, one past the buffer's `num_records`, so the store is dropped (as `deepseek_v4`'s `ffn.py` does). The scores of the unmasked lanes are stored as before. The comment links FlyDSL#1358, so the workaround can go once the lowering is fixed. - **Tests**: a kernel test runs decode launches of mixed lengths up to 650,000 tokens, and key selection at 100,003 to 650,000, at TP2/4/8. To reach those rows, `check_layer.py` takes `--max-len` (and `--lens` for `--check-keys`), and `time_layers.py` sizes its caches for `--pos`. `--check-keys` also compares the selected pools with a torch top-512 of the launch's index scores, so a wrong score fails even when the launch does not fault. ## Accuracy Tests 8 × MI355X, `lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028` (SGLang `5d8e90a0f9`), random weights in SGLang's layouts. | Test | `main` | This PR | | --- | --- | --- | | `test_rows_past_65536_pools_of_a_1m_context` | fails: 8 rows at TP4 (262,150 to 650,000 tokens), `hipErrorIllegalAddress` | passed | | the same test's `--check-keys` case (4 rows: 331,001, 262,149, 650,000, 100,003 tokens), TP2/4/8 | fails, no fault: at 262,149, 331,001 and 650,000 tokens, 52, 397 and 507 of the 512 selected pools fall below the torch top-512 | passed | | `test_glm53_flash_kernels.py`, the other 8 tests | passed | passed | Serving: #71's AgentX runs with the monokernel on include this change. 178 of 312 requests (concurrency 1) and 157 of 743 (concurrency 4) passed 262,144 tokens, up to 682,759, with no fault. ## Speed Tests and Profiling Kernel only: `time_layers.py --kind dsa --rows 2,4,8 --pos P`, TP4, 12 launches a graph, 7 trials of 200 replays, median µs a launch. `main`'s kernel and this PR's ran one after the other on the same GPUs. Past 262,144 tokens only this PR runs (`main` faults). | Position | Rows | µs a launch, `main` → this PR | Change | | --- | --- | --- | --- | | 4,095 | 2 | 61.42 → 61.35 | −0.1% | | 4,095 | 4 | 73.48 → 73.92 | +0.6% | | 4,095 | 8 | 100.87 → 101.19 | +0.3% | | 32,767 | 2 | 66.77 → 67.13 | +0.5% | | 32,767 | 4 | 82.30 → 82.64 | +0.4% | | 32,767 | 8 | 116.88 → 117.29 | +0.4% | | 131,071 | 2 | 95.71 → 96.59 | +0.9% | | 131,071 | 4 | 124.15 → 124.02 | −0.1% | | 131,071 | 8 | 179.40 → 180.31 | +0.5% | | 262,000 | 2 | 137.07 → 136.64 | −0.3% | | 262,000 | 4 | 179.66 → 179.26 | −0.2% | | 262,000 | 8 | 262.97 → 264.59 | +0.6% | Within −0.3% to +0.9%; the trials of one setup span up to 0.6%. #71's AgentX runs include this change. ## Checklist - [x] `pre-commit run --all-files` passes. (isort, ruff, ruff-format and codespell at the config's versions, on the changed files.) - [x] Tests under `test/` cover the change. - [x] The SGLang commit or PR this was tested against is named above. - [x] Accuracy and speed results are attached when kernel or forward code changes. (Kernel level; no serving run of this change alone.) - [ ] A new model adds one line to `.github/CODEOWNERS` that names its owner. (No new model.)
With --speculative-algorithm set, SGLang refused --monokernel-backend amd for GLM-5.3-Flash: the spec declared no features. The spec now declares SPECULATIVE_DECODING and builds its launches for the target's verify steps, one draft chain of draft-tokens rows a request (topk 1, draft tokens = steps + 1). A verify step of at most the TP's largest width (8 rows; 4 at TP2) runs the monokernel; the NextN draft, prefill and wider steps keep SGLang's path. KDA: one task a (request, head) runs the request's rows in order from the committed state and leaves it as is. After each draft token it writes the recurrent state and the conv window to SGLang's verify scratch (intermediate_ssm, intermediate_conv_window), which SGLang commits for the accepted token after the step, as its own verify does. DSA: the cache task writes a request's rows in order, into SGLang's kpool ring of verify + 3 slots; a row's pool and its sparse attention wait for the chain's earlier rows, whose keys it attends. Tests: check_layer.py --verify compares a verify launch with its draft tokens' decode launches bitwise (outputs, verify scratch, caches, rings) and checks the committed KDA state is untouched.
The KDA state now carries the verify scratch, which a decode launch leaves unset.
…tup, and test MTP serving With MTP, SGLang's FlashInfer KDA verify skips the KDA state commit that a verify launch reads; the spec now refuses that setup at startup instead of raising on the first verify step. SGLANG_READS drops the unused is_target_verify. TestMonokernelMtp serves GSM8K with MTP3 and checks the score and accept length.
…ery verify step runs the kernels
58578c8 to
464c8a9
Compare
There was a problem hiding this comment.
Follow-up to our review of 11-03. Requests 2, 3 and 4 are done.
Request 1 is open. At c4, P90 1/ITL is +34.5% in the body and +5.4% in our 900-s rerun. Please add one c4 replicate pair with 3,600-s windows, or give the delta on the turns that both arms completed.
Our run, 8 × MI355X, rocm/sgl-dev:v0.5.22-rocm724-mi35x-20261029:
|
For request 1, I split our 3,600-s runs by turn, keyed by
These 85 turns have a lower ITL, so P90 1/ITL over all turns is +34.5%, but +24.7% over the turns that both arms completed. P90 E2E-normalized interactivity over the turns that both arms completed is +19.9%, against +14% in your 900-s run. I've added this split for c1, c4 and c16 to the body. Please let me know if you would still like a replicate pair at c4. |
|
@Jacob0226 Thanks. The split answers request 1. On the turns that both arms completed, c4 is +24.7% P90 1/ITL and +19.9% P90 E2E-normalized interactivity, near our +14%. We do not need a replicate pair. All four requests are done. |
… monokernel KDA, DSA, dense, embed and head launches build at widths 12 and 16 (spec WIDTHS 2/4/8/12/16). Above 8 rows the MoE runs expert-major (FlashBuild.grouped): each expert's weights are read once for all the rows that pick it. Steps of 8 rows or fewer keep sgl-project#71's per-row MoE. Thread guards that assumed one pass over 8 rows now loop in passes. Co-authored-by: Cursor <cursoragent@cursor.com>
Motivation
GLM-5.3-Flash is served with MTP, but the glm53_flash spec declares no features, so SGLang refuses
--monokernel-backend amdonce--speculative-algorithmis set: a server gets MTP or the monokernel, not both. This PR runs MTP verify steps on the monokernel when a step has at most 8 rows (4 at TP2): MTP3 at 1–2 requests, MTP5 at 1. The NextN draft, prefill and wider steps keep SGLang's path. AgentX P90 interactivity, same node: +27% at concurrency 1 (MTP5) and +35% at concurrency 4 (MTP3).SGLang: upstream
5d8e90a0f9(has sgl-project/sglang#46036 and #46980), no other SGLang PRs. Rebased on #94. AgentX contexts pass 262,144 tokens, so the speed runs include #97 (now onmain), which fixes a decode fault there.Modifications
Feature.SPECULATIVE_DECODINGand builds every launch forverify= draft tokens rows a request. A step runs the monokernel only if each request has exactlyverifytokens and the step fits the width; anything else runs SGLang's forward.SGLANG_ENABLE_PP_SPEC, and SGLang's fused KDA accept. With a FlashInfer KDA verify backend, SGLang skips committing the KDA state that a verify launch reads. FlashInfer KDA is CUDA-only, so MI355X never hits this.TestMonokernelMtpserves GSM8K with MTP3, 2 requests at a time, so every verify step runs the kernels. Unit tests cover the startup refusals and the fallback.Every new SGLang name is in
SGLANG_READS.SGLANG_AMD_MONO_CHECKstill checks decode steps only.Accuracy Tests
Setup: 8 × MI355X,
lmsysorg/sglang-rocm:v0.5.22-rocm720-mi35x-20261028(SGLang5d8e90a0f9, no other PRs),amd/GLM-5.3-Flash-Quark-MXFP4, TP4, bf16 KV, tilelang DSA. Off and on differ only in--monokernel-backend amd, and ran at the same time on one node (GPUs 0–3 and 4–7).few_shot_gsm8k, 1,319 questions, parallel 2 (every verify step fits 8 rows)few_shot_gsm8k, 1,319 questions, parallel 1TestMonokernelMtp)Speed Tests and Profiling
AgentX, the setup above with InferenceX's MI355X MTP recipe (golden accept length), 3,600-s windows as InferenceX runs. Off and on ran at the same time on one node.
Off → on. E2E-normalized interactivity is 1 / the P90 of E2E latency per output token, so it includes TTFT. Tok/s/GPU counts prompt tokens, so it follows the requests completed. At concurrency 16 the slow steps pass 8 rows (SGLang's path on both sides); 57% of the on side's logged decode batches had 1–2 requests.
By turn: the on arm completed every turn of the off arm (all but 2 at concurrency 16) and more of the same conversations' later turns (keyed by
conversation_idandturn_index). P90 over the turns that both arms completed, and over the turns that only the on arm completed:The turns that only the on arm completed have a lower ITL, so they raise the on arm's P90 interactivity over all turns (+34.5% at concurrency 4) above that over the turns both arms completed (+24.7%). At concurrency 4 they have longer contexts and a higher TTFT (P90 0.58 s against 0.39 s), so their E2E-normalized interactivity is lower.
Server, concurrency 4 (MTP3). Off drops
--monokernel-backend amd:Client: aiperf 0.12.0 (SemiAnalysisAI/aiperf
754356e9, in InferenceXbee3405ca):export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 aiperf profile --scenario inferencex-agentx-mvp --url http://localhost:29904 --endpoint /v1/chat/completions \ --endpoint-type chat --streaming --model amd/GLM-5.3-Flash-Quark-MXFP4 --tokenizer amd/GLM-5.3-Flash-Quark-MXFP4 \ --tokenizer-trust-remote-code --concurrency 4 --benchmark-duration 3600 --warmup-grace-period 1800 \ --warmup-requests-per-lane 10 --stats-interval 30 --random-seed 42 --failed-request-threshold 0.10 \ --trajectory-start-min-ratio 0.25 --trajectory-start-max-ratio 0.75 --trace-idle-gap-cap-seconds 300 \ --use-server-token-count --no-gpu-telemetry --num-dataset-entries 393 --slice-duration 1.0 \ --public-dataset semianalysis_cc_traces_weka_062126 --server-metrics http://localhost:29904/metrics \ --output-artifact-dir aiperf_artifactsSGLANG_SIMULATE_ACC_LEN--speculative-num-steps/--speculative-num-draft-tokens--max-running-requests/--cuda-graph-max-bs-decode--concurrencyNot covered
SGLANG_AMD_MONO_CHECKdoes not check verify steps.Checklist
pre-commit run --all-filespasses. (isort, ruff and ruff-format at the config's versions, on the changed files.)test/cover the change..github/CODEOWNERSthat names its owner. (No new model.)