fix: enable GLM-5.3-Flash Hopper FP8 KV on supported runtimes

THIRD_PARTY_NOTICES.md

@@ -1285,6 +1285,17 @@ https://github.com/sgl-project/sglang/tree/94602c9c2b7cbdb8efd5c52802dac6a1c1800
 vendored. Upstream licenses are at `LICENSE` under those immutable revisions;
 the repository Apache-2.0 license text applies to these adaptations.
 
+The Hopper SGLang FP8-KV attention precision in
+`crates/core/src/perfmodel/operators/glm53flash.rs` and its arithmetic tests
+is a modified, independently expressed analytical adaptation of
+`python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py` at
+https://github.com/sgl-project/sglang/tree/9336681310c7a7153c4fbd7c2b9ae62011899014
+(Copyright SGLang contributors, Apache-2.0). Changes: express the raw-FP8
+TileLang QK/PV precision as a tensor-core roofline rate on SM90 and verify
+the mathematical work with hand-calculated examples; no kernel source is
+vendored. The upstream license is preserved by the repository and packaged
+Apache-2.0 `LICENSE` text.
+
 ## vLLM DeepSeek V4 cache sizing
 
 `src/aisimulate_core/sdk/memory/deepseek_v4.py` and the indexer cache-dtype

crates/core/src/perfmodel/engine/readiness.rs

@@ -91,13 +91,13 @@ fn validate_compute_capability(db: &PerfDatabase, quant: QuantMapping) -> Result
     Ok(())
 }
 
-/// GLM-5.3-Flash sparse MLA on SM9x: the modelled stock runtimes serve the
-/// latent KV only in BF16. SGLang 0.5.20's IndexPool DSA backends cannot read
-/// FP8 KV there, and vLLM 0.31.0 `FLASHINFER_MLA_SPARSE_SM90` hands an FP8
-/// cache to FlashInfer 0.7.0.post1 `plan()` as uint8 storage, which it
-/// rejects. An FP8-KV request (including the FP8 default) fails here in every
-/// database mode instead of pricing a deployment the runtimes cannot serve;
-/// the requested dtype is never switched.
+/// Hopper FP8 KV requires a runtime that implements the GLM IndexPool path.
+/// vLLM 0.32.0 fixes the SM90 FlashInfer plan dtype (#55222, c37f86e5721f),
+/// and SGLang 0.5.22 enables CUDA TileLang FP8 KV (#42957, 9336681310c7).
+/// The latter requires TileLang for both DSA prefill and decode. These are
+/// runtime restrictions, not a hardware prohibition or measured-data coverage.
+/// Sources: https://github.com/vllm-project/vllm/pull/55222
+///          https://github.com/sgl-project/sglang/pull/42957
 fn validate_glm53_kv_capability(
     db: &PerfDatabase,
     op: &crate::operators::glm53flash::Glm53AttentionOp,
@@ -107,12 +107,24 @@ fn validate_glm53_kv_capability(
         && let Some(sm) = db.system_spec.gpu.sm_version
         && (90..100).contains(&sm)
     {
+        let minimum = match op.backend.as_str() {
+            "vllm" => "0.32.0",
+            "sglang" => "0.5.22",
+            _ => return Ok(()),
+        };
+        if db
+            .version
+            .parse::<pep440_rs::Version>()
+            .is_ok_and(|version| version >= minimum.parse::<pep440_rs::Version>().unwrap())
+        {
+            return Ok(());
+        }
         return Err(AicError::InvalidEngineConfig(format!(
             "GLM-5.3-Flash on GPU {} (SM{sm}) cannot use FP8 sparse-MLA KV with {}: the \
-             stock SM90 runtimes (vLLM 0.31.0, SGLang 0.5.20) serve this model with BF16 KV \
-             only. Pass kvcache_quant_mode=\"bfloat16\" (the FP8 default does not apply on \
-             SM90).",
-            db.system, op.backend,
+             requested backend_version={} requires {} >= {}. Select a supported runtime \
+             (SGLang requires TileLang DSA prefill and decode), or pass \
+             kvcache_quant_mode=\"bfloat16\".",
+            db.system, op.backend, db.version, op.backend, minimum,
         )));
     }
     Ok(())
@@ -831,6 +843,65 @@ mod tests {
         }
     }
 
+    #[test]
+    fn glm53_hopper_fp8_kv_admission_tracks_runtime_fixes() {
+        let root = systems();
+        let path = root.path().join("b200_sxm.yaml");
+        let mut spec: serde_yaml::Value =
+            serde_yaml::from_str(&std::fs::read_to_string(&path).unwrap()).unwrap();
+        spec["gpu"]["sm_version"] = serde_yaml::to_value(90).unwrap();
+        std::fs::write(path, serde_yaml::to_string(&spec).unwrap()).unwrap();
+
+        // Stable upstream release boundaries; prereleases of the first fixed
+        // version and unknown labels do not establish that the fix is present.
+        for (backend, version, supported) in [
+            ("vllm", "0.31.0", false),
+            ("vllm", "0.32.0rc1", false),
+            ("vllm", "0.32.0", true),
+            ("vllm", "0.32.0+cu130", true),
+            ("vllm", "0.33.0", true),
+            ("sglang", "0.5.20", false),
+            ("sglang", "0.5.21", false),
+            ("sglang", "0.5.22.dev1", false),
+            ("sglang", "0.5.22", true),
+            ("sglang", "0.5.22.post1", true),
+            ("sglang", "main", false),
+        ] {
+            for mode in [
+                DatabaseMode::Silicon,
+                DatabaseMode::Hybrid,
+                DatabaseMode::Sol,
+            ] {
+                let mut db = PerfDatabase::load(root.path(), "b200_sxm", "vllm", "0.24.0")
+                    .unwrap()
+                    .with_mode(mode, TransferPolicy::ALL);
+                db.tables_mut().version = version.into();
+                let mut sparse = crate::operators::glm53flash::tests::attention("sparse_mla");
+                sparse.backend = backend.into();
+                sparse.checkpoint_format = "fp8".into();
+                let op = Op::TokenScale(crate::operators::op::TokenScaleOp {
+                    op: Box::new(Op::Glm53Attention(sparse.clone())),
+                    numerator: 1,
+                    denominator: 1,
+                });
+                let result = validate_glm53_kv(&db, [&op].into_iter());
+                assert_eq!(result.is_ok(), supported, "{backend} {version}: {result:?}");
+                if supported {
+                    let ready = validate(&db, [&op].into_iter());
+                    if mode == DatabaseMode::Silicon {
+                        // Reopening the runtime path must not supply missing
+                        // measurements or silently switch the FP8 identity.
+                        assert!(matches!(ready, Err(AicError::PerfDatabase(_))), "{ready:?}");
+                    } else {
+                        ready.unwrap();
+                    }
+                }
+                sparse.kv_cache_dtype = KvCacheQuantMode::Bfloat16;
+                validate_glm53_kv_capability(&db, &sparse).unwrap();
+            }
+        }
+    }
+
     #[test]
     fn glm53_fp8_sparse_kv_is_rejected_on_sm90_in_every_mode() {
         let root = systems();
@@ -850,10 +921,15 @@ mod tests {
                 DatabaseMode::Hybrid,
                 DatabaseMode::Sol,
             ] {
-                let db = PerfDatabase::load(root.path(), "b200_sxm", "vllm", "0.24.0")
+                let mut db = PerfDatabase::load(root.path(), "b200_sxm", "vllm", "0.24.0")
                     .unwrap()
                     .with_mode(mode, TransferPolicy::ALL);
                 for backend in ["vllm", "sglang"] {
+                    db.tables_mut().version = match backend {
+                        "vllm" => "0.31.0",
+                        _ => "0.5.20",
+                    }
+                    .into();
                     let mut sparse = crate::operators::glm53flash::tests::attention("sparse_mla");
                     sparse.backend = backend.into();
                     sparse.checkpoint_format = "fp8".into();

crates/core/src/perfmodel/operators/glm53flash.rs

@@ -7,6 +7,9 @@
 //! db9527a46873454610df6dbedf79a36d6bf1a7f6 (v0.31.0), vllm/models/glm5next/common/{model,attention,kda}.py,
 //! and sgl-project/sglang 94602c9c2b7cbdb8efd5c52802dac6a1c180089e,
 //! python/sglang/srt/models/glm5_next.py (Apache-2.0). See THIRD_PARTY_NOTICES.md.
+//! Hopper FP8-KV precision: sgl-project/sglang 9336681310c7a7153c4fbd7c2b9ae62011899014,
+//! python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py (Apache-2.0,
+//! Copyright SGLang contributors); modified analytical adaptation, no kernel copy.
 //! These are payload/roofline bounds, not allocator or kernel-launch predictions.
 
 use crate::common::enums::{DatabaseMode, GemmQuantMode, KvCacheQuantMode, MoeQuantMode};
@@ -488,11 +491,27 @@ impl Glm53AttentionOp {
                 fp32,
             ));
             let latent = self.latent_bytes();
+            // SGLang's raw-FP8 TileLang path casts Q to FP8 and uses FP8
+            // tensor cores for QK/PV. The runtime gate admits this Hopper
+            // combination from 0.5.22; other backend/hardware paths keep
+            // their existing arithmetic contract.
+            // https://github.com/sgl-project/sglang/blob/9336681310c7a7153c4fbd7c2b9ae62011899014/python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py
+            let latent_attention_rate = if self.backend == "sglang"
+                && self.kv_cache_dtype == KvCacheQuantMode::Fp8
+                && spec
+                    .gpu
+                    .sm_version
+                    .is_some_and(|sm| (90..100).contains(&sm))
+            {
+                fp8
+            } else {
+                bf16
+            };
             result = result.plus(leaf(
                 spec,
                 4.0 * selected_pairs * n * k,
                 selected_pairs * k * latent + 4.0 * x * n * k,
-                bf16,
+                latent_attention_rate,
             ));
             result = result.plus(leaf(
                 spec,
@@ -2249,4 +2268,46 @@ pub(crate) mod tests {
             }
         }
     }
+
+    #[test]
+    fn hopper_sglang_fp8_latent_attention_uses_fp8_math() {
+        // With BF16=1 TFLOP/s and FP8=2 TFLOP/s, the two 32-head,
+        // width-512 attention matmuls cost 65,536 FLOPs per selected pair.
+        // Decode at position 4098 selects 2048 pooled tokens + 2 live-tail
+        // tokens: 134,348,800 FLOPs, saving 0.0671744 ms at FP8 precision.
+        // A four-token causal prefill has 1+2+3+4=10 pairs, saving
+        // 0.00032768 ms. All other arithmetic has identical precision.
+        let mut spec = db(DatabaseMode::Sol).system_spec.clone();
+        spec.gpu.bfloat16_tc_flops = Some(1e12);
+        spec.gpu.fp8_tc_flops = Some(2e12);
+        for (backend, sm, uses_fp8_math) in [
+            ("sglang", 90, true),
+            ("sglang", 100, false),
+            ("vllm", 90, false),
+        ] {
+            spec.gpu.sm_version = Some(sm);
+            let mut fp8 = attention("sparse_mla");
+            fp8.backend = backend.into();
+            let mut bf16 = fp8.clone();
+            bf16.kv_cache_dtype = KvCacheQuantMode::Bfloat16;
+            for (is_context, s, expected_delta_ms) in
+                [(false, 4098.0, 0.0671744), (true, 4.0, 0.00032768)]
+            {
+                fp8.is_context = is_context;
+                bf16.is_context = is_context;
+                let a = fp8.sol(&spec, 1.0, s, 0.0).unwrap().sol.unwrap();
+                let b = bf16.sol(&spec, 1.0, s, 0.0).unwrap().sol.unwrap();
+                let expected = if uses_fp8_math {
+                    expected_delta_ms
+                } else {
+                    0.0
+                };
+                assert!(
+                    (b.math_ms - a.math_ms - expected).abs() < 1e-10,
+                    "{backend} SM{sm} context={is_context}: math delta {} vs {expected}",
+                    b.math_ms - a.math_ms
+                );
+            }
+        }
+    }
 }

docs/glm53flash-native-contract.md

@@ -171,8 +171,73 @@ model = RustForwardPassPerfModel.best_available(
 )
 ```
 
-Hopper deployments pass `kvcache_quant_mode="bfloat16"`. The packaged systems
-keep their own `nccl_version`, so the NCCL vocab gather is priced from the
+Hopper deployments using these pinned runtimes pass
+`kvcache_quant_mode="bfloat16"`. FP8 KV is also available for op-level modeling
+on H100/H200 with **vLLM >= 0.32.0** or **SGLang >= 0.5.22**. Select the
+`backend_version` explicitly and retain `kvcache_quant_mode="fp8"`; the engine
+does not silently convert KV precision. Earlier runtimes still reject the
+FP8 request with a version-specific diagnostic.
+
+These newer versions are not yet in the packaged queryable-version slots.
+For an analytical evaluation, opt in with `AIC_ALLOW_UNLISTED_VERSIONS=1`
+and select `database_mode="SOL"`. For example:
+
+```bash
+AIC_ALLOW_UNLISTED_VERSIONS=1 python - <<'PY'
+from aisimulate_core.sdk import ForwardPassPerfModelConfig, RustForwardPassPerfModel
+
+model = RustForwardPassPerfModel.best_available(
+    ForwardPassPerfModelConfig(
+        model="zai-org/GLM-5.3-Flash",
+        system="h200_sxm",
+        backend="sglang",
+        backend_version="0.5.22",  # "0.32.0" for vLLM
+        worker_type="aggregated",
+        tp=8,
+        moe_tp_size=8,
+        moe_ep_size=1,
+        kvcache_quant_mode="fp8",
+        database_mode="SOL",
+        estimation_mode="op_level",
+        fallback_policy="deny",
+    )
+)
+try:
+    print(model.static_phase_latency(
+        batch_size=32, input_tokens=4096, output_tokens=2, prefill=False
+    ))
+finally:
+    model.close()
+PY
+```
+
+SGLang serving requires all three flags below. Its Hopper default backend
+does not select this path automatically:
+
+```bash
+--kv-cache-dtype fp8_e4m3 \
+--dsa-prefill-backend tilelang \
+--dsa-decode-backend tilelang
+```
+
+The SGLang CUDA TileLang path excludes decode context parallelism greater than
+one and HiSparse; the modeled TP-attention baseline uses neither. Support
+comes from [SGLang #42957](https://github.com/sgl-project/sglang/pull/42957)
+(merge `9336681310c7a7153c4fbd7c2b9ae62011899014`, included in v0.5.22).
+The vLLM SM90 plan-dtype fix is
+[#55222](https://github.com/vllm-project/vllm/pull/55222)
+(merge `c37f86e5721f089a27c8e4358c30e435805fcaab`, included in v0.32.0).
+
+Opening this runtime path does not add performance measurements. SOL remains
+an analytical estimate; for Hopper SGLang FP8 KV, its latent attention QK/PV
+terms use FP8 tensor-core throughput to match the TileLang implementation.
+Projection and KDA precision remain unchanged. SILICON still requires
+exact-runtime FP8-KV rows;
+the packaged Hopper BF16-KV data cannot satisfy an FP8 request, and prior
+runtime data cannot establish coverage for the new version. HYBRID retains
+its existing source labels and fallback behavior.
+
+The packaged systems keep their own `nccl_version`, so the NCCL vocab gather is priced from the
 packaged NCCL table, not from the campaign's NCCL 2.30.7 collection.
 
 An external systems root remains optional: a root with its own `<system>.yaml`

docs/perf-model/support-matrix.md

@@ -53,11 +53,17 @@ checkpoint precision, or backends outside vLLM/SGLang are `SDK_UNREPRESENTABLE`.
 The backend rejection describes AISimulate's modeling coverage; it does not
 establish that the backend itself cannot run the model.
 These rows retain their diagnostics and phases,
-with no latency or `PASS` status. Only the recognized `NotImplementedError`
-diagnostics receive these categories. Unexpected build errors and query-time
+with no latency or `PASS` status. Only recognized build-time diagnostics
+receive these categories. Unexpected build errors and query-time
 failures still block nightly qualification. Matching currently uses the exact
-exception type and message. Other restrictions (nextn/decoder replay,
-megamoe/wideep, and non-FP8 KV) and changed diagnostics remain blocking until
+exception type and message. The native build-time GLM-5.3-Flash Hopper FP8-KV
+rejection for vLLM before 0.32.0 or SGLang before 0.5.22 is
+`FRAMEWORK_INCOMPATIBLE`; this is a runtime-version limitation, not a hardware
+prohibition. Newer runtimes remain subject to exact-runtime performance-data
+coverage. See the [KV runtime requirements](../glm53flash-native-contract.md)
+for SGLang's required TileLang launch flags. Other restrictions
+(nextn/decoder replay, megamoe/wideep, and unsupported KV dtypes) and changed
+diagnostics remain blocking until
 explicitly classified; adding those probe axes requires extending coverage.
 
 The site's branch selector identifies the tested main or release snapshot.

python/aisimulate/THIRD_PARTY_NOTICES.md

@@ -1285,6 +1285,17 @@ https://github.com/sgl-project/sglang/tree/94602c9c2b7cbdb8efd5c52802dac6a1c1800
 vendored. Upstream licenses are at `LICENSE` under those immutable revisions;
 the repository Apache-2.0 license text applies to these adaptations.
 
+The Hopper SGLang FP8-KV attention precision in
+`crates/core/src/perfmodel/operators/glm53flash.rs` and its arithmetic tests
+is a modified, independently expressed analytical adaptation of
+`python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py` at
+https://github.com/sgl-project/sglang/tree/9336681310c7a7153c4fbd7c2b9ae62011899014
+(Copyright SGLang contributors, Apache-2.0). Changes: express the raw-FP8
+TileLang QK/PV precision as a tensor-core roofline rate on SM90 and verify
+the mathematical work with hand-calculated examples; no kernel source is
+vendored. The upstream license is preserved by the repository and packaged
+Apache-2.0 `LICENSE` text.
+
 ## vLLM DeepSeek V4 cache sizing
 
 `src/aisimulate_core/sdk/memory/deepseek_v4.py` and the indexer cache-dtype

python/aisimulate/src/aisimulate_core/sdk/models/glm53flash.py

@@ -151,10 +151,9 @@ def __init__(self, info: dict, model_config, backend_name: str):
         expected = self.checkpoint_format if self.checkpoint_format == "nvfp4" else "fp8_block"
         if model_config.gemm_quant_mode.name != expected or model_config.moe_quant_mode.name != expected:
             raise NotImplementedError("GLM-5.3-Flash requires its checkpoint's native FP8/NVFP4 precision partition")
-        # Sparse-MLA latent KV: FP8 or BF16 on both backends. On SM90 both
-        # stock runtimes serve BF16 only: SGLang's IndexPool DSA backends cannot
-        # read FP8 KV, and vLLM 0.31.0 FLASHINFER_MLA_SPARSE_SM90 hands FP8 to
-        # FlashInfer plan() as uint8 storage, which it rejects. KDA state is FP32.
+        # Sparse-MLA latent KV: FP8 or BF16 on both backends. Native readiness
+        # gates SM90 FP8 KV by runtime version: vLLM >= 0.32.0 or SGLang >=
+        # 0.5.22 (TileLang DSA prefill and decode). KDA state remains FP32.
         kv_cache_dtype = model_config.kvcache_quant_mode.name
         if kv_cache_dtype not in {"fp8", "bfloat16"}:
             raise NotImplementedError(