Skip to content

[Perf] [MI355X] glm53_flash: Score a verify request's index tiles once for its rows - #100

Open
Jacob0226 wants to merge 6 commits into
sgl-project:mainfrom
Jacob0226:glm53-flash-iscore-req
Open

[Perf] [MI355X] glm53_flash: Score a verify request's index tiles once for its rows#100
Jacob0226 wants to merge 6 commits into
sgl-project:mainfrom
Jacob0226:glm53-flash-iscore-req

Conversation

@Jacob0226

@Jacob0226 Jacob0226 commented Nov 4, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Motivation

On long contexts, a verify step's sparse-attention launch slows with every row: 307 µs at 16 rows and 131K tokens. Index scoring gives each CTA one task per (row, tile), so a CTA runs 8 row-tiles back to back, each waiting for its keys. The rows of one verify request share a page table and read the same keys. This PR loads a tile's keys once per request and scores all its rows against them.

Modifications

Only MTP verify launches change. Decode rows and pad rows keep one task a row. Outputs, caches, scores and selections are bitwise those of #98. No new SGLANG_READS, Requirements, env flags or scratch, and the fault-injection site numbers do not move.

File Change
gfx950/kernel.py, index-score stage A task is a (verify request, tile): the tile's keys load once, then each row scores them with its own query and weights, masked to its own pools. The rows' queries and weights arrive in one poll and sit in the stage's LDS.
test/kernels/glm53_flash/time_layers.py --verify L (with --kind dsa) times a verify step: rows / L requests, each request's L rows on its page table at --pos .. --pos + L - 1
spec.py A task reads its request's first page-table row, so an eager verify step checks that the DSA backend's real_page_table repeats each request's row for its rows, and raises otherwise. At 1ced9a9eb0 both the eager and the graph-replay metadata build it with repeat_interleave; a graph replay runs no host code and is not checked.
test/unit/test_glm53_flash_spec.py, test/e2e/test_glm53_flash_server.py The check passes SGLang's layout and raises on one changed row; TestMonokernelMtpEager runs eager MTP verify steps through it

Accuracy Tests

MI355X, rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102, SGLang 1ced9a9eb0. GSM8K: TP4, amd/GLM-5.3-Flash-Quark-MXFP4, TileLang DSA, MTP3, both sides --enable-amd-monokernel; #98 on GPUs 0–3 and this PR on GPUs 4–7 at the same time. GSM8K prompts stay under 2048 tokens, where every pool is selected, so the kernel checks carry the long-context evidence.

Check Result
test/kernels/test_glm53_flash_kernels.py (8 GPUs, TP2/4/8) 10 passed in 30 min, #97's 1M-context test included
test/e2e/test_glm53_flash_server.py 6 passed in 19 min; TestMonokernelMtp (4 requests at a time): GSM8K 0.985, accept length 3.18
TestMonokernelMtpEager: MTP3, --disable-cuda-graph, 4 requests of 2 to 11 pages passed, 8 verify steps a request. Negative control: with the check fed the page table rolled by one row, the server raises the check's error and the test fails
Verify 4 at 12 and 16 rows (lens to 8195), verify 4 at 16 rows and verify 6 at 12 rows past 262,144 tokens (to 650,000), TP4 and TP8 (check_layer.py) 10 passed, bitwise equal to the decode launches
test/unit on SGLang 1ced9a9eb0 every glm53_flash test passes; the other failures are those of main
GSM8K 1319, MTP3, parallel 4 (accuracy / spec accept length / output tok/s) #98 this PR
16 rows in most decode batches 0.920 / 3.31 / 684 0.925 / 3.31 / 687

Long horizon. AIME26 and GPQA Diamond (sgl-eval a231b7a), off (no monokernel) on GPUs 0–3 and this PR on GPUs 4–7 at the same time, the AgentX server flags below with --max-running-requests 8 --cuda-graph-max-bs-decode 8. MTP3 with real acceptance (no SGLANG_SIMULATE_ACC_*), 4 requests at a time, so every verify step of this PR runs the kernels; thinking on, temperature 1.0, top_p 0.95, up to 131,072 tokens. The difference is per question (the mean of its repeats), with a 95% t interval.

Benchmark Samples Off pass@1 This PR pass@1 This PR − off (95% CI) Outputs past 32K tokens, off / this PR Truncated, off / this PR
AIME26 30 × 8 92.50% 94.17% +1.67 [−3.35, +6.68] 28 / 27 6 / 2
GPQA Diamond 198 × 4 91.16% 92.55% +1.39 [−0.27, +3.05] 65 / 57 0 / 2 
python3 -m sgl_eval.cli run aime26 --base-url http://127.0.0.1:30200/v1 --n-repeats 8 --num-threads 4 \
  --max-tokens 131072 --temperature 1.0 --top-p 0.95 --chat-template-kwarg enable_thinking=true
# GPQA Diamond: run gpqa --n-repeats 4, the rest the same

Speed Tests and Profiling

AgentX, MTP3, 3,600-s windows, golden accept length 2.99 (SGLANG_SIMULATE_ACC_LEN). The first conc 8 pair and conc 32 ran both sides at once on one node; the replicate ran one server on the node at a time (GPUs 0–3, #98 then this PR). At conc 32 (control), the slow steps pass 16 rows and run SGLang's path on both sides; P90 interactivity and throughput match within 1.1%. P90 E2E-normalized interactivity (1 / P90 E2E latency per output token) moves +0.4% and +10.4% in the two conc 8 pairs. P90 TTFT moves +6.7% and +2.8%. This PR changes only verify launches, and #98's own P90 TTFT moved 0.492 → 0.521 s (+5.9%) between two rounds on one node (#98's AgentX table), so I read it as the setup's noise. Kernel-only: time_layers.py, TP4, 12 launches. One MTP3 request (4 rows) costs up to 2.1 µs a launch (+2.6%) from 4K to 16K tokens: about 23 µs over the 11 sparse-attention layers, 0.4% of its 5.3-ms step. Three requests cost up to 1.8 µs (+1.2%) at 4K. From 32K tokens every width is faster. The tasks are not gated by pool count: a gate would keep two index-score paths for about 2 µs a launch.

Concurrency P90 interactivity, 1 / ITL p90 (tok/s/user) P50 interactivity (tok/s/user) P90 E2E-normalized interactivity (tok/s/user) P90 TTFT (s) Total tok/s/GPU Requests completed
8 258.7 → 273.1 (+5.5%) 392.9 → 438.6 (+11.6%) 185.1 → 185.9 (+0.4%) 0.481 → 0.513 (+6.7%) 15,366 → 15,708 (+2.2%) 1,395 → 1,408
8, replicate 257.3 → 294.0 (+14.3%) 398.6 → 444.4 (+11.5%) 180.6 → 199.5 (+10.4%) 0.485 → 0.499 (+2.8%) 15,488 → 15,572 (+0.5%) 1,397 → 1,400
32 (control) 89.9 → 89.0 (−1.0%) 157.1 → 158.4 (+0.8%) 67.4 → 67.2 (−0.3%) 0.949 → 0.916 (−3.5%) 52,094 → 52,547 (+0.9%) 4,364 → 4,373

Server, this PR, conc 8; #98 differs only in the plugin tree, conc 32 in --max-running-requests 64 --cuda-graph-max-bs-decode 64:

export SGLANG_SIMULATE_ACC_LEN=2.99 SGLANG_SIMULATE_ACC_METHOD=match-expected SGLANG_SIMULATE_ACC_TOKEN_MODE=real-draft-token \
  SGLANG_DSA_MQA_LOGITS_FREE_MEM_FRACTION=0.04 SGLANG_TIMEOUT_KEEP_ALIVE=900
python3 -m sglang.launch_server --model-path amd/GLM-5.3-Flash-Quark-MXFP4 --served-model-name amd/GLM-5.3-Flash-Quark-MXFP4 \
  --host 0.0.0.0 --port 29908 --trust-remote-code --tp 4 --ep-size 1 --kv-cache-dtype bfloat16 \
  --dsa-prefill-backend tilelang --dsa-decode-backend tilelang --moe-runner-backend aiter \
  --mamba-radix-cache-strategy extra_buffer --tool-call-parser glm47 --reasoning-parser glm45 \
  --chunked-prefill-size 16384 --mem-fraction-static 0.85 --max-running-requests 16 --cuda-graph-max-bs-decode 16 \
  --model-loader-extra-config '{"enable_multithread_load": true, "num_threads": 32}' --mamba-full-memory-ratio 0.9 \
  --speculative-algorithm EAGLE --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4 \
  --watchdog-timeout 1800 --enable-metrics --enable-amd-monokernel

Client: aiperf 0.12.0 (SemiAnalysisAI/aiperf 754356e9), --concurrency 8 or 32:

export AIPERF_HTTP_TCP_USER_TIMEOUT=900000
aiperf profile --scenario inferencex-agentx-mvp --url http://localhost:29908 --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 8 --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:29908/metrics \
  --output-artifact-dir aiperf_artifacts
Aggregate JSONs: 6 runs
Conc 8, #98 
{
  "hw": "mi355x",
  "conc": 8,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 1484,
  "num_requests_successful": 1395,
  "request_accounting": {
    "records_total": 1484,
    "records_profiled": 1395,
    "records_dropped_total": 89,
    "records_warmup_dropped": 87,
    "records_error_dropped": 61,
    "error_categories": {
      "InvalidInferenceResultError": 61
    }
  },
  "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": 3605,
      "mean": 0.38669,
      "p50": 0.0,
      "p75": 1,
      "p90": 1,
      "p95": 2,
      "std": 0.64224
    },
    "latency": {
      "ttft": {
        "mean": 0.27648,
        "p50": 0.1691,
        "p75": 0.2852,
        "p90": 0.48056,
        "p95": 0.71676,
        "std": 0.34899
      },
      "e2el": {
        "mean": 3.38279,
        "p50": 1.51652,
        "p75": 3.72127,
        "p90": 7.94662,
        "p95": 13.04511,
        "std": 5.30568
      },
      "itl": {
        "mean": 0.00274,
        "p50": 0.00255,
        "p75": 0.00323,
        "p90": 0.00387,
        "p95": 0.00439,
        "std": 0.00103
      },
      "tpot": {
        "mean": 0.00274,
        "p50": 0.00255,
        "p75": 0.00323,
        "p90": 0.00387,
        "p95": 0.00439,
        "std": 0.00103
      },
      "intvty": {
        "mean": 365.48223,
        "p50": 392.88228,
        "p75": 309.71911,
        "p90": 258.71915,
        "p95": 227.88181,
        "std": 739.12569
      },
      "e2e_norm_intvty": {
        "mean": 252.5847,
        "p50": 319.28832,
        "p75": 244.67984,
        "p90": 185.10495,
        "p95": 143.19212,
        "std": 106.98021
      },
      "full_response_itl": {
        "mean": 0.00273,
        "p50": 0.00255,
        "p75": 0.00322,
        "p90": 0.00386,
        "p95": 0.00439,
        "std": 0.00104
      },
      "full_response_intvty": {
        "mean": 366.2917,
        "p50": 392.8571,
        "p75": 310.24088,
        "p90": 258.81158,
        "p95": 227.89032,
        "std": 434.90007
      }
    },
    "tokens": {
      "input": {
        "mean": 157699.75771,
        "p50": 108206,
        "p75": 174815.5,
        "p90": 375643.6,
        "p95": 491234.3,
        "std": 136545.59901
      },
      "output_actual": {
        "mean": 1148.66882,
        "p50": 439,
        "p75": 1211.5,
        "p90": 2948.8,
        "p95": 4583.4,
        "std": 1889.02694
      },
      "output_expected": {
        "mean": 1943.74942,
        "p50": 1012,
        "p75": 2654,
        "p90": 4796,
        "p95": 6727,
        "std": 2503.14429
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 61018.32035
      },
      "output": {
        "tokens_per_second": 444.45117
      },
      "total": {
        "tokens_per_second": 61462.77152
      },
      "duration_seconds": 3605.32969,
      "per_gpu": {
        "total_tput_tps": 15365.69288,
        "output_tput_tps": 111.11279,
        "input_tput_tps": 15254.58009
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97797
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.97286,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.97286,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 214591232
      },
      "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.13,
      "gpu_total_tokens": 11941632,
      "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": 220577876,
      "generation_total": 1602396,
      "requests_completed": 1395,
      "prompt_by_source": {
        "gpu_cache_hit": 214591232,
        "cpu_or_external_cache_hit": null,
        "computed": 8226655,
        "raw": {
          "device": 214591232
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11941632
}
Conc 8, this PR 
{
  "hw": "mi355x",
  "conc": 8,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 1497,
  "num_requests_successful": 1408,
  "request_accounting": {
    "records_total": 1497,
    "records_profiled": 1408,
    "records_dropped_total": 89,
    "records_warmup_dropped": 87,
    "records_error_dropped": 60,
    "error_categories": {
      "InvalidInferenceResultError": 60
    }
  },
  "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.39083,
      "p50": 0.0,
      "p75": 1,
      "p90": 1,
      "p95": 2,
      "std": 0.65258
    },
    "latency": {
      "ttft": {
        "mean": 0.28035,
        "p50": 0.17531,
        "p75": 0.31917,
        "p90": 0.51292,
        "p95": 0.70566,
        "std": 0.33911
      },
      "e2el": {
        "mean": 3.11904,
        "p50": 1.41704,
        "p75": 3.40461,
        "p90": 7.58669,
        "p95": 11.74724,
        "std": 4.76872
      },
      "itl": {
        "mean": 0.00253,
        "p50": 0.00228,
        "p75": 0.0029,
        "p90": 0.00366,
        "p95": 0.00444,
        "std": 0.00096
      },
      "tpot": {
        "mean": 0.00253,
        "p50": 0.00228,
        "p75": 0.0029,
        "p90": 0.00366,
        "p95": 0.00444,
        "std": 0.00096
      },
      "intvty": {
        "mean": 395.64796,
        "p50": 438.60358,
        "p75": 344.63561,
        "p90": 273.07131,
        "p95": 225.37079,
        "std": 118.98661
      },
      "e2e_norm_intvty": {
        "mean": 266.8831,
        "p50": 349.0965,
        "p75": 266.6213,
        "p90": 185.86914,
        "p95": 144.38071,
        "std": 112.55333
      },
      "full_response_itl": {
        "mean": 0.00253,
        "p50": 0.00228,
        "p75": 0.0029,
        "p90": 0.00366,
        "p95": 0.00444,
        "std": 0.00096
      },
      "full_response_intvty": {
        "mean": 395.94608,
        "p50": 438.73086,
        "p75": 344.90138,
        "p90": 273.04215,
        "p95": 225.26149,
        "std": 211.38688
      }
    },
    "tokens": {
      "input": {
        "mean": 159539.68253,
        "p50": 108280,
        "p75": 178614,
        "p90": 381460.1,
        "p95": 498347.35,
        "std": 138905.5569
      },
      "output_actual": {
        "mean": 1152.16548,
        "p50": 446.5,
        "p75": 1222.5,
        "p90": 2946.4,
        "p95": 4574.85,
        "std": 1883.13003
      },
      "output_expected": {
        "mean": 1940.52227,
        "p50": 1012,
        "p75": 2646.5,
        "p90": 4759.7,
        "p95": 6727,
        "std": 2493.35947
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 62381.68959
      },
      "output": {
        "tokens_per_second": 450.50879
      },
      "total": {
        "tokens_per_second": 62832.19838
      },
      "duration_seconds": 3600.9264,
      "per_gpu": {
        "total_tput_tps": 15708.0496,
        "output_tput_tps": 112.6272,
        "input_tput_tps": 15595.4224
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97827
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.9726,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.9726,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 219047680
      },
      "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.13,
      "gpu_total_tokens": 11941632,
      "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": 225218587,
      "generation_total": 1622252,
      "requests_completed": 1408,
      "prompt_by_source": {
        "gpu_cache_hit": 219047680,
        "cpu_or_external_cache_hit": null,
        "computed": 8435802,
        "raw": {
          "device": 219047680
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11941632
}
Conc 8, replicate, #98 
{
  "hw": "mi355x",
  "conc": 8,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 1486,
  "num_requests_successful": 1397,
  "request_accounting": {
    "records_total": 1486,
    "records_profiled": 1397,
    "records_dropped_total": 89,
    "records_warmup_dropped": 87,
    "records_error_dropped": 62,
    "error_categories": {
      "InvalidInferenceResultError": 62
    }
  },
  "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": 3594,
      "mean": 0.38815,
      "p50": 0.0,
      "p75": 1,
      "p90": 1,
      "p95": 2,
      "std": 0.63857
    },
    "latency": {
      "ttft": {
        "mean": 0.27432,
        "p50": 0.17356,
        "p75": 0.28879,
        "p90": 0.48503,
        "p95": 0.72539,
        "std": 0.33977
      },
      "e2el": {
        "mean": 3.33957,
        "p50": 1.45387,
        "p75": 3.61866,
        "p90": 7.91482,
        "p95": 13.00926,
        "std": 5.16869
      },
      "itl": {
        "mean": 0.00271,
        "p50": 0.00251,
        "p75": 0.00319,
        "p90": 0.00389,
        "p95": 0.00448,
        "std": 0.00106
      },
      "tpot": {
        "mean": 0.00271,
        "p50": 0.00251,
        "p75": 0.00319,
        "p90": 0.00389,
        "p95": 0.00448,
        "std": 0.00106
      },
      "intvty": {
        "mean": 369.23311,
        "p50": 398.61722,
        "p75": 313.65503,
        "p90": 257.31,
        "p95": 223.17066,
        "std": 120.741
      },
      "e2e_norm_intvty": {
        "mean": 254.32865,
        "p50": 329.2181,
        "p75": 250.89462,
        "p90": 180.62697,
        "p95": 140.44627,
        "std": 108.45161
      },
      "full_response_itl": {
        "mean": 0.0027,
        "p50": 0.00251,
        "p75": 0.00319,
        "p90": 0.00388,
        "p95": 0.00448,
        "std": 0.00107
      },
      "full_response_intvty": {
        "mean": 369.85439,
        "p50": 398.77377,
        "p75": 313.70493,
        "p90": 257.43466,
        "p95": 223.32244,
        "std": 354.08388
      }
    },
    "tokens": {
      "input": {
        "mean": 158295.14674,
        "p50": 108215,
        "p75": 175676,
        "p90": 379349.2,
        "p95": 493296.6,
        "std": 137358.01972
      },
      "output_actual": {
        "mean": 1152.14674,
        "p50": 445,
        "p75": 1217,
        "p90": 2970,
        "p95": 4581.6,
        "std": 1889.53536
      },
      "output_expected": {
        "mean": 1946.65221,
        "p50": 1012,
        "p75": 2660,
        "p90": 4796,
        "p95": 6727,
        "std": 2502.19541
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 61503.94137
      },
      "output": {
        "tokens_per_second": 447.6547
      },
      "total": {
        "tokens_per_second": 61951.59607
      },
      "duration_seconds": 3595.51461,
      "per_gpu": {
        "total_tput_tps": 15487.89902,
        "output_tput_tps": 111.91367,
        "input_tput_tps": 15375.98534
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97805
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.97301,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.97301,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 215741184
      },
      "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.13,
      "gpu_total_tokens": 11941632,
      "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": 221725038,
      "generation_total": 1609552,
      "requests_completed": 1397,
      "prompt_by_source": {
        "gpu_cache_hit": 215741184,
        "cpu_or_external_cache_hit": null,
        "computed": 8210469,
        "raw": {
          "device": 215741184
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11941632
}
Conc 8, replicate, this PR 
{
  "hw": "mi355x",
  "conc": 8,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 1489,
  "num_requests_successful": 1400,
  "request_accounting": {
    "records_total": 1489,
    "records_profiled": 1400,
    "records_dropped_total": 89,
    "records_warmup_dropped": 87,
    "records_error_dropped": 55,
    "error_categories": {
      "InvalidInferenceResultError": 55
    }
  },
  "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.38861,
      "p50": 0.0,
      "p75": 1,
      "p90": 1,
      "p95": 2,
      "std": 0.63494
    },
    "latency": {
      "ttft": {
        "mean": 0.28133,
        "p50": 0.17352,
        "p75": 0.29177,
        "p90": 0.49883,
        "p95": 0.74765,
        "std": 0.35798
      },
      "e2el": {
        "mean": 3.08356,
        "p50": 1.38694,
        "p75": 3.35715,
        "p90": 7.52602,
        "p95": 11.49143,
        "std": 4.74295
      },
      "itl": {
        "mean": 0.00246,
        "p50": 0.00225,
        "p75": 0.00286,
        "p90": 0.0034,
        "p95": 0.00392,
        "std": 0.00082
      },
      "tpot": {
        "mean": 0.00246,
        "p50": 0.00225,
        "p75": 0.00286,
        "p90": 0.0034,
        "p95": 0.00392,
        "std": 0.00082
      },
      "intvty": {
        "mean": 405.79662,
        "p50": 444.39619,
        "p75": 349.88594,
        "p90": 293.98086,
        "p95": 255.10684,
        "std": 115.4904
      },
      "e2e_norm_intvty": {
        "mean": 267.32396,
        "p50": 348.33455,
        "p75": 268.26277,
        "p90": 199.49286,
        "p95": 150.07149,
        "std": 111.80236
      },
      "full_response_itl": {
        "mean": 0.00246,
        "p50": 0.00225,
        "p75": 0.00286,
        "p90": 0.0034,
        "p95": 0.00392,
        "std": 0.00083
      },
      "full_response_intvty": {
        "mean": 406.12997,
        "p50": 444.59415,
        "p75": 350.04172,
        "p90": 293.79958,
        "p95": 255.09088,
        "std": 217.62126
      }
    },
    "tokens": {
      "input": {
        "mean": 159051.36786,
        "p50": 108249.5,
        "p75": 178615.75,
        "p90": 380964.9,
        "p95": 496695.45,
        "std": 138160.74373
      },
      "output_actual": {
        "mean": 1149.275,
        "p50": 443,
        "p75": 1217.25,
        "p90": 2944.8,
        "p95": 4578.9,
        "std": 1886.21356
      },
      "output_expected": {
        "mean": 1941.33308,
        "p50": 1012,
        "p75": 2646.5,
        "p90": 4788.2,
        "p95": 6727,
        "std": 2499.04223
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 61842.83704
      },
      "output": {
        "tokens_per_second": 446.8646
      },
      "total": {
        "tokens_per_second": 62289.70165
      },
      "duration_seconds": 3600.60964,
      "per_gpu": {
        "total_tput_tps": 15572.42541,
        "output_tput_tps": 111.71615,
        "input_tput_tps": 15460.70926
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97818
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.97303,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.97303,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 217236736
      },
      "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.13,
      "gpu_total_tokens": 11941632,
      "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": 223258631,
      "generation_total": 1608988,
      "requests_completed": 1400,
      "prompt_by_source": {
        "gpu_cache_hit": 217236736,
        "cpu_or_external_cache_hit": null,
        "computed": 8300380,
        "raw": {
          "device": 217236736
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11941632
}
Conc 32 (control), #98 
{
  "hw": "mi355x",
  "conc": 32,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 4719,
  "num_requests_successful": 4364,
  "request_accounting": {
    "records_total": 4719,
    "records_profiled": 4364,
    "records_dropped_total": 355,
    "records_warmup_dropped": 354,
    "records_error_dropped": 255,
    "error_categories": {
      "InvalidInferenceResultError": 255
    }
  },
  "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": 3613,
      "mean": 1.20758,
      "p50": 1,
      "p75": 2,
      "p90": 3,
      "p95": 4,
      "std": 1.24523
    },
    "latency": {
      "ttft": {
        "mean": 0.63493,
        "p50": 0.31618,
        "p75": 0.51867,
        "p90": 0.94944,
        "p95": 1.65058,
        "std": 1.4367
      },
      "e2el": {
        "mean": 7.99326,
        "p50": 3.50207,
        "p75": 8.60878,
        "p90": 18.8969,
        "p95": 30.97125,
        "std": 13.28627
      },
      "itl": {
        "mean": 0.00774,
        "p50": 0.00637,
        "p75": 0.00778,
        "p90": 0.01112,
        "p95": 0.01341,
        "std": 0.01765
      },
      "tpot": {
        "mean": 0.00774,
        "p50": 0.00637,
        "p75": 0.00778,
        "p90": 0.01112,
        "p95": 0.01341,
        "std": 0.01765
      },
      "intvty": {
        "mean": 129.1376,
        "p50": 157.09178,
        "p75": 128.60362,
        "p90": 89.90402,
        "p95": 74.55645,
        "std": 57.57113
      },
      "e2e_norm_intvty": {
        "mean": 90.63781,
        "p50": 133.47657,
        "p75": 97.75068,
        "p90": 67.43327,
        "p95": 46.64909,
        "std": 51.27983
      },
      "full_response_itl": {
        "mean": 0.00773,
        "p50": 0.00636,
        "p75": 0.00777,
        "p90": 0.01112,
        "p95": 0.01342,
        "std": 0.01763
      },
      "full_response_intvty": {
        "mean": 129.3992,
        "p50": 157.17756,
        "p75": 128.64255,
        "p90": 89.95017,
        "p95": 74.49821,
        "std": 255.42503
      }
    },
    "tokens": {
      "input": {
        "mean": 171632.40445,
        "p50": 126178,
        "p75": 229261.75,
        "p90": 336233,
        "p95": 446562.35,
        "std": 150552.39633
      },
      "output_actual": {
        "mean": 1032.17369,
        "p50": 427,
        "p75": 1102.25,
        "p90": 2629.2,
        "p95": 4088.75,
        "std": 1724.49609
      },
      "output_expected": {
        "mean": 1393.86814,
        "p50": 630,
        "p75": 1617,
        "p90": 3601,
        "p95": 5210,
        "std": 2133.55535
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 207132.00229
      },
      "output": {
        "tokens_per_second": 1245.6634
      },
      "total": {
        "tokens_per_second": 208377.66568
      },
      "duration_seconds": 3616.06997,
      "per_gpu": {
        "total_tput_tps": 52094.41642,
        "output_tput_tps": 311.41585,
        "input_tput_tps": 51783.00057
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97931
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.96995,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.96995,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 728570880
      },
      "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.47,
      "gpu_total_tokens": 11382784,
      "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": 751139001,
      "generation_total": 4504412,
      "requests_completed": 4364,
      "prompt_by_source": {
        "gpu_cache_hit": 728570880,
        "cpu_or_external_cache_hit": null,
        "computed": 32303299,
        "raw": {
          "device": 728570880
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11382784
}
Conc 32 (control), this PR 
{
  "hw": "mi355x",
  "conc": 32,
  "image": "rocm/sgl-dev:v0.5.23-rocm10-mi35x-20261102",
  "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": 4729,
  "num_requests_successful": 4373,
  "request_accounting": {
    "records_total": 4729,
    "records_profiled": 4373,
    "records_dropped_total": 356,
    "records_warmup_dropped": 354,
    "records_error_dropped": 227,
    "error_categories": {
      "InvalidInferenceResultError": 227
    }
  },
  "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": 3605,
      "mean": 1.21276,
      "p50": 1,
      "p75": 2,
      "p90": 3,
      "p95": 4,
      "std": 1.26006
    },
    "latency": {
      "ttft": {
        "mean": 0.58364,
        "p50": 0.31599,
        "p75": 0.5107,
        "p90": 0.91615,
        "p95": 1.52994,
        "std": 1.19281
      },
      "e2el": {
        "mean": 7.88138,
        "p50": 3.49197,
        "p75": 8.43161,
        "p90": 18.95917,
        "p95": 29.64797,
        "std": 13.15991
      },
      "itl": {
        "mean": 0.0075,
        "p50": 0.00631,
        "p75": 0.0078,
        "p90": 0.01124,
        "p95": 0.01365,
        "std": 0.00891
      },
      "tpot": {
        "mean": 0.0075,
        "p50": 0.00631,
        "p75": 0.0078,
        "p90": 0.01124,
        "p95": 0.01365,
        "std": 0.00891
      },
      "intvty": {
        "mean": 133.36892,
        "p50": 158.40058,
        "p75": 128.27245,
        "p90": 88.98291,
        "p95": 73.25823,
        "std": 103.74581
      },
      "e2e_norm_intvty": {
        "mean": 92.69329,
        "p50": 133.53235,
        "p75": 97.90157,
        "p90": 67.21531,
        "p95": 48.25762,
        "std": 56.21257
      },
      "full_response_itl": {
        "mean": 0.00748,
        "p50": 0.00631,
        "p75": 0.00779,
        "p90": 0.01123,
        "p95": 0.01365,
        "std": 0.00891
      },
      "full_response_intvty": {
        "mean": 133.60587,
        "p50": 158.44357,
        "p75": 128.32757,
        "p90": 89.06762,
        "p95": 73.25838,
        "std": 274.93739
      }
    },
    "tokens": {
      "input": {
        "mean": 172410.14887,
        "p50": 126446,
        "p75": 230045,
        "p90": 336755.4,
        "p95": 461986,
        "std": 151646.2072
      },
      "output_actual": {
        "mean": 1032.32426,
        "p50": 427,
        "p75": 1103,
        "p90": 2629.8,
        "p95": 4073,
        "std": 1723.42601
      },
      "output_expected": {
        "mean": 1400.19446,
        "p50": 630.5,
        "p75": 1617.75,
        "p90": 3610,
        "p95": 5242.1,
        "std": 2147.66034
      }
    },
    "throughput": {
      "input": {
        "tokens_per_second": 208937.40524
      },
      "output": {
        "tokens_per_second": 1251.03513
      },
      "total": {
        "tokens_per_second": 210188.44037
      },
      "duration_seconds": 3608.49499,
      "per_gpu": {
        "total_tput_tps": 52547.11009,
        "output_tput_tps": 312.75878,
        "input_tput_tps": 52234.35131
      }
    },
    "cache": {
      "theoretical_cache_hit_rate": 0.97944
    }
  },
  "server_metrics": {
    "present": true,
    "adapter": "sglang",
    "metric_count": 86,
    "cache": {
      "gpu_cache_hit_rate": 0.97021,
      "cpu_cache_hit_rate": null,
      "external_cache_hit_rate": null,
      "overall_cache_hit_rate": 0.97021,
      "prefix_cache_hits": null,
      "prefix_cache_queries": null,
      "external_prefix_cache_hits": null,
      "external_prefix_cache_queries": null,
      "cached_tokens_by_source": {
        "device": 733052672
      },
      "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.47,
      "gpu_total_tokens": 11382784,
      "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": 755564512,
      "generation_total": 4514359,
      "requests_completed": 4373,
      "prompt_by_source": {
        "gpu_cache_hit": 733052672,
        "cpu_or_external_cache_hit": null,
        "computed": 32204895,
        "raw": {
          "device": 733052672
        }
      }
    },
    "sources": []
  },
  "kv_cache_pool_tokens": 11382784
}
Median decode step by requests decoding (server log) rows #98 this PR
1 4 5.60 ms 5.26 ms
2 8 7.63 ms 7.00 ms
3 12 9.48 ms 8.57 ms
4 16 10.70 ms 9.66 ms

pos is the context length (tokens) each verify row sits at; 4 rows is one MTP3 request, 16 rows four. Median of 3 runs in alternating order, nothing else on the node; bold: this PR is faster.

time_layers.py --kind dsa --verify 4, µs a launch, #98 → this PR 4 rows 8 12 16
pos 2047 68.5 → 68.4 (−0.3%) 95.9 → 94.5 (−1.4%) 134.1 → 133.8 (−0.2%) 152.2 → 151.2 (−0.7%)
pos 4095 80.5 → 82.6 (+2.6%) 109.0 → 109.1 (+0.1%) 150.2 → 152.0 (+1.2%) 173.2 → 171.6 (−0.9%)
pos 8191 81.3 → 83.1 (+2.2%) 109.9 → 109.1 (−0.7%) 151.6 → 152.1 (+0.3%) 174.9 → 173.7 (−0.7%)
pos 16383 83.1 → 84.5 (+1.7%) 112.6 → 110.2 (−2.2%) 157.4 → 154.5 (−1.8%) 180.9 → 176.0 (−2.7%)
pos 32767 88.6 → 88.0 (−0.7%) 122.1 → 114.4 (−6.3%) 171.5 → 159.4 (−7.0%) 199.7 → 180.5 (−9.6%)
pos 49151 96.0 → 92.5 (−3.7%) 133.8 → 118.1 (−11.7%) 187.0 → 165.3 (−11.6%) 217.2 → 186.3 (−14.2%)
pos 131071 132.2 → 112.6 (−14.8%) 189.1 → 143.4 (−24.2%) 260.8 → 199.9 (−23.3%) 307.9 → 224.5 (−27.1%)
pos 393215 211.1 → 171.1 (−18.9%) 297.4 → 212.7 (−28.5%) 401.2 → 275.4 (−31.4%) 478.2 → 307.1 (−35.8%)

Checklist

  • pre-commit run --all-files passes.
  • Tests under test/ cover the change.
  • The SGLang commit or PR this was tested against is named above.
  • Accuracy and speed results are attached when kernel or forward code changes.
  • A new model adds one line to .github/CODEOWNERS that names its owner. (No new model.)

@Jacob0226
Jacob0226 force-pushed the glm53-flash-iscore-req branch from 2b1b156 to b1b0651 Compare November 4, 2026 08:07
@Jacob0226
Jacob0226 marked this pull request as ready for review November 4, 2026 11:11
@Jacob0226
Jacob0226 requested a review from jhinpan as a code owner November 4, 2026 11:11
@Jacob0226 Jacob0226 mentioned this pull request Nov 4, 2026
Merged
4 of 5 tasks
@Jacob0226
Jacob0226 force-pushed the glm53-flash-iscore-req branch from b1b0651 to 08768b0 Compare November 5, 2026 01:29
@jhinpan jhinpan mentioned this pull request Nov 5, 2026
Open
4 of 5 tasks
jhinpan
jhinpan approved these changes Nov 5, 2026
View reviewed changes

@jhinpan jhinpan left a comment •
edited
Loading

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approved. The change is correct, and the kernel gain reproduces on our node.

time_layers.py --kind dsa --verify 4, TP4, rocm/sgl-dev:v0.5.22-rocm724-mi35x-20261029, #98 → this PR, median µs per launch, 3 runs an arm:

Pos 4 rows 8 12 16
8191 +2.4% −5.6% −0.7% −1.0%
131071 −13.9% −23.0% −23.3% −28.0% (313.4 → 225.7)
393215 −20.5% −29.2% −30.9% −36.3% (484.5 → 308.6)

test_glm53_flash_kernels.py passes on b1b0651 (10 passed), and verify still equals decode bitwise. 08768b0 keeps the kernel files byte-identical.

Follow-ups, not blocking:

  1. A unit uses its first row's ROW_BT, so it depends on SGLang repeating a verify request's page-table row. Please add a check that fails if SGLang stops this.
  2. One MTP3 request at pos 8191 is 2.4% slower. Please state this cost, or gate the unit tasks by pool count.
  3. At AgentX conc 8, TTFT is 6.7% higher, but this PR changes only verify launches. Please add a replicate pair, with one server on the node at a time.

@jhinpan

jhinpan commented Nov 5, 2026

Copy link
Copy Markdown
Collaborator

Hi @Jacob0226. Can u add more workload accuracy test like GPQA and AIM26 to test its correctness in long horizon for now since CICD is not finalized.

Jacob0226 and others added 3 commits November 6, 2026 01:07
@Jacob0226 @cursoragent
[Feat] [MI355X] glm53_flash: Run verify steps of up to 16 rows on the…
3a473f1 
… 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>
@Jacob0226 @cursoragent
[Test] glm53_flash: Cover widths 12 and 16 in the serving test, and r…
4b8a377 
…un its MTP case 4 requests at a time

The GSM8K decode run takes 24 requests at a time, so its batch still steps
past the largest width; the prefill consistency and SGLANG_AMD_MONO_CHECK
cases add 11 and 16 rows; TestMonokernelMtp runs 4 requests at a time, so
every verify step fits 16 rows and runs the kernels.

Co-authored-by: Cursor <cursoragent@cursor.com>
@Jacob0226
[Test] glm53_flash: Build the grouped MoE from GROUPED_FROM in the ke…
3172459 
…rnel tests

check_layer.py and time_layers.py read the server's threshold instead of restating rows > 8. time_layers.py --per-row-moe times the per-row MoE at every width, the baseline of the grouped one.
@Jacob0226

Jacob0226 commented Nov 6, 2026

Copy link
Copy Markdown
Contributor Author

Sure

Jacob0226 and others added 3 commits November 6, 2026 01:08
@Jacob0226
[Perf] [MI355X] glm53_flash: Score a verify request's index tiles onc…
bd60363 
…e for its rows

The index-score stage ran one task a (row, 256-pool tile): a CTA set up each of
its rows' tiles in turn (kv_ready poll, fence, barrier, key loads, query and
weight polls), so at 16 rows and 131K tokens every CTA ran 8 row-tiles back to
back. A verify request's rows share its page table, so a task is now a
(request, tile): the tile's keys load once and each row scores them with its own
query and weights, masked to its own pools. Decode and pad rows keep a task a
row. The query and the weights arrive in one poll. Outputs, caches, scores and
selections are bitwise those of the row tasks.
@Jacob0226 @cursoragent
[Test] glm53_flash: Time MTP verify steps in time_layers (--verify)
25f51a8 
Co-authored-by: Cursor <cursoragent@cursor.com>
@Jacob0226 @cursoragent
[Test] glm53_flash: Check that a verify request's page-table row repe…
d9f24c1 
…ats for its rows

Co-authored-by: Cursor <cursoragent@cursor.com>
@Jacob0226
Jacob0226 force-pushed the glm53-flash-iscore-req branch from 08768b0 to 25f51a8 Compare November 6, 2026 01:28
@Jacob0226

Jacob0226 commented Nov 6, 2026

Copy link
Copy Markdown
Contributor Author

Thanks, @jhinpan. All three follow-ups are in the body:

  1. Page-table check (the last commit): an eager verify step raises if a request's page-table row is not repeated for its rows. TestMonokernelMtpEager passes, and fails with the table rolled by one row (negative control).
  2. Cost: one MTP3 request costs up to 2.1 µs a launch (+2.6%) from 4K to 16K tokens, about 0.4% of its step; from 32K every width is faster. Not gated; see the time_layers table, now one table from 2K to 393K tokens.
  3. Replicate pair, one server on the node at a time: P90 interactivity +14.3%, P90 TTFT +2.8%. [Feat] [MI355X] glm53_flash: Run verify steps of up to 16 rows on the monokernel #98's own TTFT varies 5.9% between rounds, so the +6.7% reads as noise.

AIME26 and GPQA Diamond (off vs this PR, MTP3, thinking on) are running; I'll post them here when they finish.

@Jacob0226

Jacob0226 commented Nov 6, 2026

Copy link
Copy Markdown
Contributor Author

@jhinpan AIME26 and GPQA Diamond are in Accuracy Tests (Long horizon). Off and this PR ran at the same time, MTP3 with real acceptance, thinking on, up to 131,072 tokens:

  • AIME26 (30 × 8): pass@1 92.50% off, 94.17% this PR; difference +1.67 points, 95% CI [−3.35, +6.68].
  • GPQA Diamond (198 × 4): 91.16% off, 92.55% this PR; +1.39 points, 95% CI [−0.27, +3.05].

Both intervals include zero. 93 outputs off and 84 on this PR ran past 32K tokens.

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

Reviewers

@jhinpan jhinpan jhinpan approved these changes

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@Jacob0226 @jhinpan