From 50ff412719b1869927b2f15eed90b101c6d23d5c Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Tue, 8 Sep 2026 19:53:56 +0800 Subject: [PATCH 01/16] [Kernel] Add exact TP2 E4M3 scalar decode fast path Keep the q8 FP32-partial route opt-in while retaining the full-round quality gate and recording measured TP2 limits. Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../kernels/benchmark_sm70_tp2_e4m3_scalar.py | 144 ++++++++++ docs/design/sm70_dflash2_tp2_verifier.md | 156 +++++++++++ docs/design/sm70_v100_migration_control.md | 17 ++ .../flash_attn_v100/flash_attn_interface.py | 17 ++ flash-attention-v100/include/fused_mha.h | 3 + .../kernel/flash_decode_paged.cu | 87 ++++-- flash-attention-v100/kernel/fp8_kv_utils.cuh | 20 +- flash-attention-v100/kernel/fused_mha_api.cpp | 6 + .../test_sm70_tp2_e4m3_scalar_fast.py | 256 ++++++++++++++++++ 9 files changed, 684 insertions(+), 22 deletions(-) create mode 100644 benchmarks/kernels/benchmark_sm70_tp2_e4m3_scalar.py create mode 100644 docs/design/sm70_dflash2_tp2_verifier.md create mode 100644 tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py diff --git a/benchmarks/kernels/benchmark_sm70_tp2_e4m3_scalar.py b/benchmarks/kernels/benchmark_sm70_tp2_e4m3_scalar.py new file mode 100644 index 0000000000..f2954e74df --- /dev/null +++ b/benchmarks/kernels/benchmark_sm70_tp2_e4m3_scalar.py @@ -0,0 +1,144 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Compare TP2 scalar attention over distinct KV layer working sets. + +This measures attention operators, not model verification rounds. Run the +bitwise kernel tests separately before considering a model experiment. +""" + +import argparse +import hashlib +import json +import os +import statistics +from pathlib import Path + +import torch + +FLAG = "VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST" + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--context-lengths", type=int, nargs="+", default=[270, 1100]) + parser.add_argument("--layers", type=int, default=16) + parser.add_argument("--json-out", type=Path, required=True) + args = parser.parse_args() + if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): + raise RuntimeError("requires an owned SM70 GPU") + from flash_attn_v100 import flash_attn_interface as interface + + native = interface.flash_attn_v100_cuda + if getattr(native, "tp2_e4m3_scalar_fast_version", lambda: 0)() < 1: + raise RuntimeError("rebuild Flash-V100 with TP2 scalar fast revision 1") + if args.layers < 1 or any(not 8 <= n <= 262144 for n in args.context_lengths): + raise ValueError("positive layer count and context lengths 8..262144 required") + torch.manual_seed(20260908) + report = { + "measurement": "distinct-KV attention working set, not complete model rounds", + "native_sha256": hashlib.sha256(Path(native.__file__).read_bytes()).hexdigest(), + "gpu": torch.cuda.get_device_name(), + "torch": torch.__version__, + "cuda": torch.version.cuda, + "layers": args.layers, + "rows": [], + } + previous = os.environ.get(FLAG) + try: + for length in args.context_lengths: + report["rows"].append(measure(native, length, args.layers)) + finally: + if previous is None: + os.environ.pop(FLAG, None) + else: + os.environ[FLAG] = previous + args.json_out.parent.mkdir(parents=True, exist_ok=True) + args.json_out.write_text(json.dumps(report, indent=2) + "\n") + print(json.dumps(report, indent=2)) + + +def measure(native, length, layers): + page, parts = 3296, 256 + pages = (length + page - 1) // page + operands = [] + for _ in range(layers): + kv = torch.randn((pages, 2, page, 2, 256), device="cuda", dtype=torch.float16) + k, v = kv.to(torch.float8_e4m3fn).view(torch.uint8).unbind(1) + q = torch.randn((8, 12, 256), device="cuda", dtype=torch.float16) + table = torch.randperm(pages, device="cuda").int()[None].repeat(8, 1) + seq = torch.arange(length - 7, length + 1, device="cuda").int() + operands.append((q, k, v, table, seq)) + out = torch.empty_like(operands[0][0]) + tmp = torch.empty((8, 12, parts, 256), device="cuda") + maxima = torch.empty((8, 12, parts), device="cuda") + sums = torch.empty_like(maxima) + active = torch.full((1,), parts, device="cuda", dtype=torch.int32) + + def call(operand): + q, k, v, table, seq = operand + native.decode_paged_fwd( + q, + k, + v, + out, + table, + seq, + tmp, + maxima, + sums, + active, + 0.0625, + 1024, + parts, + "fp8_e4m3", + 0.5, + 1.25, + -1, + -1, + None, + 0, + ) + + # Compare every layer before capturing a shared-workspace timing graph. + for operand in operands: + os.environ[FLAG] = "0" + call(operand) + expected = out.clone() + os.environ[FLAG] = "1" + call(operand) + if not torch.equal(out.view(torch.int16), expected.view(torch.int16)): + raise AssertionError("candidate output differs from control") + graphs = {} + for enabled in ("0", "1"): + os.environ[FLAG] = enabled + graph = torch.cuda.CUDAGraph() + before = native.tp2_e4m3_scalar_fast_launch_count() + with torch.cuda.graph(graph): + for operand in operands: + call(operand) + count = native.tp2_e4m3_scalar_fast_launch_count() - before + assert count == (layers if enabled == "1" else 0) + graphs[enabled] = graph + samples = {enabled: [] for enabled in graphs} + for trial in range(7): + for enabled in ("0", "1") if trial % 2 == 0 else ("1", "0"): + graph = graphs[enabled] + graph.replay() + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + start.record() + for _ in range(4): + graph.replay() + end.record() + end.synchronize() + samples[enabled].append(start.elapsed_time(end) / 4) + return { + "context_length": length, + "all_layer_outputs_bitwise_equal": True, + "samples_ms": samples, + "median_ms": {k: statistics.median(v) for k, v in samples.items()}, + } + + +if __name__ == "__main__": + main() diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md new file mode 100644 index 0000000000..31faf2514d --- /dev/null +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -0,0 +1,156 @@ +# DFlash2 TP2 verification cost + +## Scope and frozen baseline + +The TP2 campaign targets approximately 25 ms per complete B1/q8 DFlash2 round +on two rear V100-SXM2-32GB GPUs. A round includes target, logits/sampling, +state handling, context work and draft. TP4 optimization is a separate campaign. + +Integration base: `e5d63c51f0fcc1ddf75d229e3df06bf52df206f5`. +Use the QUASAR Qwen3.8-27B NVFP4 checkpoint at +`d8e6fbfa3e3a78899b440222b827430045a05b44` and the DFlash2 checkpoint at +`dedf8df68adfb1afeaf7b7480c0a0243108177b4`. The workload uses Python 3.12.13, +Torch 2.10.0+cu128, CUDA 12.8, TP2 on physical GPUs 4 and 7, FP16 activations, +E4M3 target KV, FP16 draft KV and FP32 logits. Both attention backends are +FLASH_ATTN_V100. Keep V2 runner, target/draft CUDA Graphs, context pipeline and +context KV graph enabled. Maximum context is 262144, batch-token budget 4096, +maximum sequences 4 and memory utilization 0.8. Only one request is active. + +Sampling remains temperature 1, top-k 20, top-p 0.95, xhigh thinking, natural +EOS and at most 1024 output tokens. The release1k fixture uses seed 20260925 +and 1019 input tokens; MBPP28 uses seed 0 and 135 input tokens. Startup and +model preparation are outside decode timing. The original baseline uses +frozen copies of existing native libraries; it is not a rebuild of all main +sources. Retained runtime manifests hash the actual mapped worker libraries. + +One startup, one warmup and five measured requests per fixture gave: + +| Metric | release1k | MBPP28 | +| --- | ---: | ---: | +| Median request-average complete round, ms | 44.973 | 35.119 | +| Median pure decode, tokens/s | 67.832 | 127.662 | +| Median warm TTFT, ms | 575.387 | 147.247 | +| Accepted drafts per round | 2.063291 | 3.500000 | +| Emitted tokens per round | 3.063291 | 4.500000 | +| Output tokens | 242 | 270 | +| Draft rounds | 79 | 60 | + +Outputs repeat within this startup and finish naturally. MBPP28 passes its +three supplied assertions. These are short-context baselines, not a 256K +latency result or the three-startup final acceptance gate. + +## Trace and first optimization + +Ten steady rounds from both ranks show approximately 35.260 ms of target GPU +service, including 14.963 ms of TurboMind projections and 12.838 ms of scalar +attention. Draft GPU service is 6.638 ms; target head/sampling is 2.274 ms. +The profiled critical-rank round interval is 47.057 ms. Service sums and +profiled wall intervals are diagnostic, not unprofiled performance claims. + +TP2's 12 query heads and two KV heads do not enter the existing six-head, +single-KV-head E4M3 grouped route. Its scalar attention uses 1024-token +partitions, FP32 partial output and FP32 partition statistics. The observed +launch is `(8, 12, 256)` CTAs, 256 threads/CTA, 40 registers/thread and +12880 bytes shared memory for the frozen control. + +`VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST=1` selects an experimental specialization +only for q shape `[8,12,256]`, E4M3 KV with two heads, FP32 partial storage, +1024-token partitions and full attention without an anchored window. It is +off by default. Unverified shapes use the original route. A requested matching +route rejects a stale native library instead of silently reporting success. + +The specialization constructs normal E4M3 values directly in FP32 bit fields, +retains the original signed zeros, subnormals and NaN payload, and unrolls the +PV loop by eight. Each output still follows the original ascending-token FMA +chain. It retains partition boundaries, score reductions, FP32 intermediate +storage, output rounding, KV scales and the original final reduction kernel. +The native launch counter proves host dispatch, including capture-time calls; +it does not count CUDA Graph replays or model rounds. + +## Evidence and promotion status + +The initial isolated implementation passes: + +- All 256 E4M3 byte encodings, with bitwise equality against the original + decoder, including both signed zeros and both NaN encodings. +- 33 fixed-operand comparisons across bit conversion alone and PV unroll + factors four/eight. Final outputs and valid partial output/max/sum bits + match the frozen native implementation. Lengths include zero, partition + and page boundaries, 65537 and 262144; live CUDA Graph inputs change between + replays. +- At length 3297, all variants retain the same FP64-reference error: + maximum absolute `3.0444386e-5`, p99 absolute `1.4819749e-5`, relative L2 + `2.0301283e-4`. +- CUDA 12.8 Compute Sanitizer memcheck and racecheck on the winning u8 + partition kernel report zero errors and zero hazards, respectively. +- Live same-call shadow comparison on both model ranks: 27072 attention + calls, 665321472 output elements, zero bit differences and zero nonfinite + outputs. The original result drives generation. These runs contain + diagnostic work and are excluded from speed evidence. + +For sixteen distinct KV layer working sets, the operator median is 12.712 ms +for the frozen scalar implementation and 3.892 ms for exact bit conversion +with PV unroll eight. Conversion alone and unroll four are approximately +6.335/6.314 ms. These are attention operator results, not complete rounds. +The private u8 implementation is pinned by SHA256 +`696545418c6dae261f0bc6a3a530b34464d040de8e404a3069cfd8c2a7762ad3`. +The integrated native build is separately pinned by SHA256 +`f916e9e370eeb8d865b4de9d8b64f6e66d8831c0b4458dbe3087141dbadc1d19`; +its own GPU and runtime gates must pass before substituting it for the +isolated implementation. + +The first contemporaneous control reproduces round cost at 44.986/35.075 ms. +Its release1k trajectory has 349 tokens rather than the original startup's +242, while within-startup repetitions match. The candidate was disabled in +this control. This pre-existing startup variation is not an allowed quality +tolerance; cross-startup token/acceptance comparisons must retain this limit. +The first separate-startup candidate measures 35.921/32.706 ms, with +release1k/MBPP28 outputs of 283/297 tokens. Its corresponding control produces +349/270 tokens. The trajectories and acceptance counts differ, so the +approximately 20.15%/6.75% latency reductions are provisional performance +observations, not accepted quality-preserving gains. A within-startup graph +comparison is used next to keep prefill and projection choices fixed. +The 25 ms target and final promotion remain outstanding. + +The integrated native build passes 14 tests, including exhaustive byte +decoding, 262144-token graph replay, FP64 reference, unsupported-shape fallback +and stale-library rejection. The initial QPN2 TP2 projection screen uses +sixteen real matrices from four adjacent layers. Its best working-set median +is 0.685 ms versus 0.930 ms for TurboMind, but several reference-error metrics +grow. This arithmetic candidate is rejected for model use. Source inspection +identifies early FP16 rounding of the global scale as a separate precision +candidate; it requires fresh operator and model gates. + +## Reproduction and retained negative results + +Build Flash-V100 from this branch with the same CUDA/Torch/compiler flags and +select that module before running the tests. Set `CUDA_VISIBLE_DEVICES` only +to an owned rear GPU, and use private build/compiler caches. + +```bash +TORCH_CUDA_ARCH_LIST=7.0 MAX_JOBS=2 .venv/bin/python -m pytest \ + --confcutdir=tests/kernels/attention \ + tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py \ + tests/kernels/attention/test_sm70_e4m3_scalar_fp32.py -q +``` + +The GPU gate includes an exhaustive decoder comparison, strided output +sentinels, changing page/sequence visibility, graph replay, FP64 reference +and fallback dispatch. The stale-library gate also runs without a GPU. + +Task artifacts are retained under campaign identifier +`v100-quasar-dflash2-tp2-25ms-20260908`. They contain baseline contracts, +worker DSO inventories, Nsight data, raw endpoint responses, operator results, +sanitizer logs, source/build hashes and serial GPU queue records. The baseline +campaign identifier is `v100-quasar-dflash2-tp2-baseline-20260908`. + +A capped partition-grid experiment passed 66 bitwise cases but did not improve +the sixteen-layer working set: 12.701 ms control, 13.641 ms at cap one and +approximately 12.719 ms at caps two through sixteen. It was rejected before +model testing. Do not repeat that path without new bottleneck evidence. + +The first memcheck invocation loaded both experimental u4/u8 DSOs and reported +`CUDA_ERROR_INVALID_HANDLE` in `cuKernelGetFunction` at the second decoder-LUT +launch. The quality gate blocked model work. Running only the winning DSO +passed both sanitizer tools with API error checking retained. The failed +invocation remains recorded rather than counted as a pass. diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index f114bbc08b..e99d400ee4 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -2,6 +2,23 @@ Date: 2026-05-30 +## DFlash2 TP2 verification cost, 2026-09-08 + +[The TP2 worklog](sm70_dflash2_tp2_verifier.md) freezes main at +`e5d63c51f0fcc1ddf75d229e3df06bf52df206f5` and records initial complete-round +costs of 44.973/35.119 ms on release1k/MBPP28. Scalar E4M3 attention accounts +for about 12.84 ms in the release1k trace. An opt-in exact bit-decoder/PV +unroll specialization passes exhaustive encoding, graph/state, FP64-reference +and sanitizer gates; live attention shadow compares 665321472 elements with +zero bit differences. It is restricted to TP2 q8/D256/FP32-partial/P1024. + +The first independent-startup candidate measures 35.921/32.706 ms against +44.986/35.075 ms control. Output trajectories differ between startups, including +between controls, so this is not a quality promotion or a 25 ms result. +Keep the feature off while resolving the comparison and remaining cost. +Capping inactive partition CTAs was numerically exact but had no speed gain; +do not repeat that rejected path without new evidence. + ## DFlash2 E4M3 FP32 default policy, 2026-09-08 [The precision-default change](sm70_dflash2_fp32_defaults.md) is based on diff --git a/flash-attention-v100/flash_attn_v100/flash_attn_interface.py b/flash-attention-v100/flash_attn_v100/flash_attn_interface.py index 30cd71208f..6f362140d1 100644 --- a/flash-attention-v100/flash_attn_v100/flash_attn_interface.py +++ b/flash-attention-v100/flash_attn_v100/flash_attn_interface.py @@ -1012,6 +1012,23 @@ def flash_attn_decode_paged( seq_lens, workspace_seq_capacity_hint=workspace_seq_capacity_hint, ) + if ( + os.getenv("VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST", "0") == "1" + and e4m3_fp32 + and q.dtype == torch.float16 + and tuple(q.shape) == (8, 12, 256) + and k_cache.shape[2:] == (2, 256) + and plan.partition_size == 1024 + and window_size_left == window_size_right == -1 + and anchor_lens is None + ): + version = getattr( + flash_attn_v100_cuda, "tp2_e4m3_scalar_fast_version", None + ) + if not callable(version) or int(version()) < 1: + raise RuntimeError( + "Rebuild Flash-V100 for TP2 E4M3 scalar fast revision 1" + ) tmp_out, max_logits, exp_sums, active_num_partitions = ( _get_decode_workspace_for_plan( q, diff --git a/flash-attention-v100/include/fused_mha.h b/flash-attention-v100/include/fused_mha.h index a29a190dad..291126d61a 100644 --- a/flash-attention-v100/include/fused_mha.h +++ b/flash-attention-v100/include/fused_mha.h @@ -69,6 +69,9 @@ at::Tensor flash_attention_grouped_e4m3_fp32_paged( int64_t flash_attention_grouped_e4m3_fp32_precision_version(); +int64_t flash_attention_tp2_e4m3_scalar_fast_version(); +int64_t flash_attention_tp2_e4m3_scalar_fast_launch_count(); + int64_t flash_attention_grouped_sparse_page4_abi_version(); at::Tensor flash_attention_grouped_sparse_page4( diff --git a/flash-attention-v100/kernel/flash_decode_paged.cu b/flash-attention-v100/kernel/flash_decode_paged.cu index 3d9df0ea23..30a3a00aaf 100644 --- a/flash-attention-v100/kernel/flash_decode_paged.cu +++ b/flash-attention-v100/kernel/flash_decode_paged.cu @@ -21,6 +21,8 @@ namespace { +std::atomic tp2_e4m3_scalar_fast_calls{0}; + int kv_cache_dtype_code_from_string(const std::string& kv_cache_dtype) { if (kv_cache_dtype == "auto" || kv_cache_dtype == "float16" || kv_cache_dtype == "bfloat16") { @@ -989,7 +991,7 @@ __device__ __forceinline__ float dot_qk_half2(const __half* __restrict__ q_ptr, return warp_reduce_sum(acc); } -template +template __device__ __forceinline__ float dot_qk_cache(const __half* __restrict__ q_ptr, const void* __restrict__ k_cache, const int64_t k_index_base, @@ -1017,8 +1019,9 @@ __device__ __forceinline__ float dot_qk_cache(const __half* __restrict__ q_ptr, #pragma unroll for (int d = lane; d < D; d += kWarpSize) { const float qv = __half2float(q_ptr[d]); - const float kv = flash_v100::load_kv_cache_float_unscaled( - k_cache, k_index_base + d); + const float kv = + flash_v100::load_kv_cache_float_unscaled( + k_cache, k_index_base + d); acc = fmaf(qv, kv, acc); } return warp_reduce_sum(acc); @@ -1027,7 +1030,7 @@ __device__ __forceinline__ float dot_qk_cache(const __half* __restrict__ q_ptr, template + typename PARTIAL_T = __half, bool TP2_E4M3_FAST = false> __global__ void flash_attention_decode_partition_kernel( const __half* __restrict__ q, const void* __restrict__ k_cache, const void* __restrict__ v_cache, PARTIAL_T* __restrict__ tmp_out, @@ -1155,7 +1158,8 @@ __global__ void flash_attention_decode_partition_kernel( static_cast(block_offset) * k_token_stride + static_cast(kv_head_idx) * k_head_stride; - float score = dot_qk_cache(q_shared, k_cache, k_index, lane); + float score = dot_qk_cache(q_shared, k_cache, + k_index, lane); if (lane == 0) { if constexpr (ANCHORED_SWA) { const int token_idx = part_start + token_local; @@ -1197,16 +1201,33 @@ __global__ void flash_attention_decode_partition_kernel( for (int d = threadIdx.x; d < D; d += blockDim.x) { float acc = 0.f; - for (int i = 0; i < part_tokens; ++i) { - const int physical_block = block_idx_shared[i]; - const int block_offset = block_offset_shared[i]; - const int64_t v_index = - static_cast(physical_block) * v_block_stride + - static_cast(block_offset) * v_token_stride + - static_cast(kv_head_idx) * v_head_stride + d; - const float vv = - flash_v100::load_kv_cache_float_unscaled(v_cache, v_index); - acc = fmaf(scores_shared[i], vv, acc); + if constexpr (TP2_E4M3_FAST) { + // Expose independent loads while retaining the ascending-token FMA chain. +#pragma unroll 8 + for (int i = 0; i < part_tokens; ++i) { + const int physical_block = block_idx_shared[i]; + const int block_offset = block_offset_shared[i]; + const int64_t v_index = + static_cast(physical_block) * v_block_stride + + static_cast(block_offset) * v_token_stride + + static_cast(kv_head_idx) * v_head_stride + d; + const float vv = + flash_v100::load_kv_cache_float_unscaled(v_cache, + v_index); + acc = fmaf(scores_shared[i], vv, acc); + } + } else { + for (int i = 0; i < part_tokens; ++i) { + const int physical_block = block_idx_shared[i]; + const int block_offset = block_offset_shared[i]; + const int64_t v_index = + static_cast(physical_block) * v_block_stride + + static_cast(block_offset) * v_token_stride + + static_cast(kv_head_idx) * v_head_stride + d; + const float vv = flash_v100::load_kv_cache_float_unscaled( + v_cache, v_index); + acc = fmaf(scores_shared[i], vv, acc); + } } const float out_scale = KV_DTYPE == flash_v100::KV_CACHE_DTYPE_FP16 ? inv_part_sum @@ -3472,10 +3493,11 @@ void launch_flash_attention_decode_paged( // Second kernel version: the anchored decode-window mask is a separate // template instantiation, generated only for the fp16-KV configuration; // the non-anchored instantiations stay untouched. - const auto launch_partition = [&](auto anchored_tag) { + const auto launch_partition = [&](auto anchored_tag, auto fast_tag) { constexpr bool kAnchored = decltype(anchored_tag)::value; - flash_attention_decode_partition_kernel + constexpr bool kFast = decltype(fast_tag)::value; + flash_attention_decode_partition_kernel< + D, PARTITION_SIZE, KV_DTYPE, SEQ_LEN_ROUTE, kAnchored, PARTIAL_T, kFast> <<>>( reinterpret_cast(q.data_ptr()), k_cache.data_ptr(), v_cache.data_ptr(), @@ -3494,14 +3516,29 @@ void launch_flash_attention_decode_paged( }; if constexpr (KV_DTYPE == flash_v100::KV_CACHE_DTYPE_FP16) { if (use_anchored) { - launch_partition(std::true_type{}); + launch_partition(std::true_type{}, std::false_type{}); } else { - launch_partition(std::false_type{}); + launch_partition(std::false_type{}, std::false_type{}); } } else { TORCH_CHECK(!use_anchored, "anchored decode window requires an fp16 KV cache"); - launch_partition(std::false_type{}); + if constexpr (KV_DTYPE == flash_v100::KV_CACHE_DTYPE_FP8_E4M3 && D == 256 && + PARTITION_SIZE == 1024 && std::is_same_v) { + const char* enabled = std::getenv("VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST"); + const bool use_fast = enabled && enabled[0] == '1' && + enabled[1] == '\0' && batch_size == 8 && + num_heads_q == 12 && num_heads_kv == 2 && + window_size_left == -1 && window_size_right == -1; + if (use_fast) { + tp2_e4m3_scalar_fast_calls.fetch_add(1, std::memory_order_relaxed); + launch_partition(std::false_type{}, std::true_type{}); + } else { + launch_partition(std::false_type{}, std::false_type{}); + } + } else { + launch_partition(std::false_type{}, std::false_type{}); + } } if (!launch_reduce) { @@ -4378,6 +4415,14 @@ int64_t flash_attention_grouped_e4m3_fp32_precision_version() { return 4; } +int64_t flash_attention_tp2_e4m3_scalar_fast_version() { return 1; } + +int64_t flash_attention_tp2_e4m3_scalar_fast_launch_count() { + // Includes capture-time launches; CUDA Graph replay does not call this host + // dispatcher again. This counter proves route admission, not round count. + return tp2_e4m3_scalar_fast_calls.load(std::memory_order_relaxed); +} + int64_t flash_attention_grouped_verify_max_query_tokens() { return kGroupedVerifyMaxSupportedQ; } diff --git a/flash-attention-v100/kernel/fp8_kv_utils.cuh b/flash-attention-v100/kernel/fp8_kv_utils.cuh index dfa54a3eae..778a30fd76 100644 --- a/flash-attention-v100/kernel/fp8_kv_utils.cuh +++ b/flash-attention-v100/kernel/fp8_kv_utils.cuh @@ -48,6 +48,21 @@ __device__ __forceinline__ __half fp8_e5m2_to_half(uint8_t raw) { return __ushort_as_half(static_cast(raw) << 8); } +// E4M3 normals map exactly into the IEEE float exponent and mantissa fields. +// Keep the original NaN payload and signed zero, including E4M3 subnormals. +__device__ __forceinline__ float fp8_e4m3fn_to_float_bits(uint8_t raw) { + const uint32_t magnitude = raw & 0x7fu; + const uint32_t sign = static_cast(raw & 0x80u) << 24; + uint32_t bits = (magnitude << 20) + 0x3c000000u; + if (magnitude < 8) { + bits = __float_as_uint(static_cast(magnitude) * 0.001953125f); + } + if (magnitude == 0x7f) { + return quiet_nan_f(); + } + return __uint_as_float(bits | sign); +} + __device__ __forceinline__ __half2 fp8_e5m2_pair_to_half2(uint16_t raw_pair) { const uint32_t half2_bits = (static_cast(raw_pair & 0x00ffu) << 8) | (static_cast(raw_pair & 0xff00u) << 16); @@ -69,7 +84,7 @@ __device__ __forceinline__ float fp8_e5m2_to_float(uint8_t raw) { return __half2float(fp8_e5m2_to_half(raw)); } -template +template __device__ __forceinline__ float load_kv_cache_float_unscaled( const void* __restrict__ cache, const int64_t index) { if constexpr (KV_DTYPE == KV_CACHE_DTYPE_FP16) { @@ -78,6 +93,9 @@ __device__ __forceinline__ float load_kv_cache_float_unscaled( } else { const uint8_t* cache_u8 = reinterpret_cast(cache); const uint8_t raw = cache_u8[index]; + if constexpr (KV_DTYPE == KV_CACHE_DTYPE_FP8_E4M3 && E4M3_BITS) { + return fp8_e4m3fn_to_float_bits(raw); + } const float value = KV_DTYPE == KV_CACHE_DTYPE_FP8_E4M3 ? fp8_e4m3fn_to_float(raw) : fp8_e5m2_to_float(raw); diff --git a/flash-attention-v100/kernel/fused_mha_api.cpp b/flash-attention-v100/kernel/fused_mha_api.cpp index 1eca3b589b..e1104f609a 100644 --- a/flash-attention-v100/kernel/fused_mha_api.cpp +++ b/flash-attention-v100/kernel/fused_mha_api.cpp @@ -30,6 +30,12 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("grouped_e4m3_fp32_precision_version", &flash_attention_grouped_e4m3_fp32_precision_version, "E4M3 grouped FP32 numerical implementation revision"); + m.def("tp2_e4m3_scalar_fast_version", + &flash_attention_tp2_e4m3_scalar_fast_version, + "Capability for the opt-in TP2 E4M3 scalar decoder"); + m.def("tp2_e4m3_scalar_fast_launch_count", + &flash_attention_tp2_e4m3_scalar_fast_launch_count, + "TP2 E4M3 scalar host dispatch count, including graph capture"); m.attr("grouped_verify_e4m3") = true; m.def("grouped_verify_max_query_tokens", &flash_attention_grouped_verify_max_query_tokens, diff --git a/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py b/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py new file mode 100644 index 0000000000..25b3ae3f9c --- /dev/null +++ b/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py @@ -0,0 +1,256 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Bitwise gates for the opt-in TP2 E4M3 scalar decode implementation.""" + +from pathlib import Path +from types import SimpleNamespace + +import pytest +import torch + +FLAG = "VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST" + + +def test_requested_fast_path_rejects_stale_library(monkeypatch): + interface = pytest.importorskip("flash_attn_v100.flash_attn_interface") + monkeypatch.setenv(FLAG, "1") + monkeypatch.setattr( + interface, + "flash_attn_v100_cuda", + SimpleNamespace(grouped_e4m3_fp32_precision_version=lambda: 4), + ) + monkeypatch.setattr( + interface, "flash_attn_grouped_e4m3_fp32_available", lambda: True + ) + monkeypatch.setattr( + interface, + "_get_decode_plan", + lambda *args, **kwargs: SimpleNamespace(partition_size=1024), + ) + monkeypatch.setattr( + interface, "_assert_decode_launch_covers_seq_lens", lambda *args, **kwargs: None + ) + q = torch.empty((8, 12, 256), dtype=torch.float16) + kv = torch.empty((1, 3296, 2, 256), dtype=torch.uint8) + table = torch.zeros((8, 1), dtype=torch.int32) + seq = torch.zeros(8, dtype=torch.int32) + with pytest.raises(RuntimeError, match="TP2 E4M3 scalar fast revision 1"): + interface.flash_attn_decode_paged( + q, kv, kv, table, seq, kv_cache_dtype="fp8_e4m3" + ) + + +@pytest.fixture +def native(): + if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): + pytest.skip("requires SM70") + interface = pytest.importorskip("flash_attn_v100.flash_attn_interface") + extension = interface.flash_attn_v100_cuda + version = getattr(extension, "tp2_e4m3_scalar_fast_version", lambda: 0) + if version() < 1: + pytest.skip("rebuild Flash-V100 with TP2 scalar fast revision 1") + return extension + + +def test_all_256_e4m3_encodings(native, tmp_path): + from torch.utils.cpp_extension import load_inline + + extension = load_inline( + name="tp2_e4m3_codec_test", + cpp_sources="void decode_lut(torch::Tensor output);", + cuda_sources=r""" +#include +#include +#include "fp8_kv_utils.cuh" +__global__ void tp2_e4m3_codec_test_kernel(float* output) { + const uint8_t raw = threadIdx.x; + output[raw] = flash_v100::fp8_e4m3fn_to_float(raw); + output[256 + raw] = flash_v100::fp8_e4m3fn_to_float_bits(raw); +} +void decode_lut(torch::Tensor output) { + tp2_e4m3_codec_test_kernel<<<1, 256, 0, at::cuda::getCurrentCUDAStream()>>>( + output.data_ptr()); +} +""", + functions=["decode_lut"], + extra_include_paths=[ + str(Path(__file__).resolve().parents[3] / "flash-attention-v100/kernel") + ], + extra_cuda_cflags=["-O3", "--use_fast_math"], + build_directory=str(tmp_path), + ) + output = torch.empty((2, 256), device="cuda", dtype=torch.float32) + extension.decode_lut(output) + assert torch.equal(output[0].view(torch.int32), output[1].view(torch.int32)) + expected = torch.arange(256, device="cuda").byte().view(torch.float8_e4m3fn).float() + finite = torch.isfinite(expected) + assert torch.equal( + output[1, finite].view(torch.int32), expected[finite].view(torch.int32) + ) + assert bool(torch.isnan(output[1, ~finite]).all()) + + +def workspace(q, parts): + rows, heads, dim = q.shape + # Check that a noncontiguous output does not overwrite adjacent storage. + storage = torch.full((rows, heads, dim + 16), 123, device=q.device, dtype=q.dtype) + return ( + storage[..., :dim], + torch.empty((rows, heads, parts, dim), device=q.device), + torch.empty((rows, heads, parts), device=q.device), + torch.empty((rows, heads, parts), device=q.device), + storage, + ) + + +@pytest.mark.parametrize("length", [1025, 3297, 262144]) +def test_live_graph_bitwise_state_and_fp64_reference(native, monkeypatch, length): + torch.manual_seed(20260908) + page, parts = 3296, 256 + pages = (length + page - 1) // page + # Interleaved, strided K/V with a nonidentity page table, as in the service. + kv = torch.randn((pages, 2, page, 2, 256), device="cuda", dtype=torch.float16) + kv = kv.to(torch.float8_e4m3fn).view(torch.uint8) + k, v = kv.unbind(1) + order = torch.randperm(pages, device="cuda").int() + table = order[None].repeat(8, 1) + q = torch.randn((8, 12, 256), device="cuda", dtype=torch.float16) + q_initial = q.clone() + seq = torch.zeros(8, device="cuda", dtype=torch.int32) + active = torch.full((1,), parts, device="cuda", dtype=torch.int32) + states = [workspace(q, parts), workspace(q, parts)] + graphs = [] + for enabled, state in enumerate(states): + monkeypatch.setenv(FLAG, str(enabled)) + + def call(state=state): + native.decode_paged_fwd( + q, + k, + v, + state[0], + table, + seq, + *state[1:4], + active, + 0.0625, + 1024, + parts, + "fp8_e4m3", + 0.5, + 1.25, + -1, + -1, + None, + 0, + ) + + before = native.tp2_e4m3_scalar_fast_launch_count() + call() + assert native.tp2_e4m3_scalar_fast_launch_count() - before == enabled + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + call() + graphs.append(graph) + + initial_seq = torch.arange(length - 7, length + 1, device="cuda").int() + for replay in range(4): + seq.copy_(initial_seq) + if replay == 1: + seq[0] = 0 + elif replay == 2: + seq.zero_() + q.copy_(q_initial * (0.5 if replay % 2 else 1.0)) + for state in states: + for tensor in state[:4]: + tensor.fill_(float("nan")) + for graph in graphs: + graph.replay() + assert torch.equal( + states[0][0].view(torch.int16), states[1][0].view(torch.int16) + ) + assert bool(torch.isfinite(states[1][0]).all()) + valid = ( + torch.arange(parts, device="cuda")[None, None] + < ((seq + 1023) // 1024)[:, None, None] + ) + valid = valid.expand(8, 12, parts) + for control, candidate in zip(states[0][1:4], states[1][1:4]): + assert torch.equal( + control[valid].view(torch.int32), candidate[valid].view(torch.int32) + ) + for state in states: + assert bool((state[4][..., 256:] == 123).all()) + + if length == 3297 and replay == 0: + rk = k[order.long()].reshape(-1, 2, 256)[:length] + rv = v[order.long()].reshape(-1, 2, 256)[:length] + rk = rk.view(torch.float8_e4m3fn).double() * 0.5 + rv = rv.view(torch.float8_e4m3fn).double() * 1.25 + expected = torch.empty_like(q, dtype=torch.float64) + for head in range(2): + score = ( + q[:, head * 6 : (head + 1) * 6].transpose(0, 1).double() + @ rk[:, head].T + * 0.0625 + ) + score.masked_fill_( + torch.arange(length, device="cuda")[None, None] + >= seq[None, :, None], + -torch.inf, + ) + expected[:, head * 6 : (head + 1) * 6] = ( + score.softmax(-1) @ rv[:, head] + ).transpose(0, 1) + error = (states[1][0].double() - expected).norm() / expected.norm() + floor = (expected.half().double() - expected).norm() / expected.norm() + assert float(error) <= 1.04 * float(floor) + 2e-6 + + +@pytest.mark.parametrize( + "rows,heads,kv_heads,partition,window", + [ + (4, 12, 2, 1024, -1), + (8, 6, 1, 1024, -1), + (8, 12, 2, 256, -1), + (8, 12, 2, 1024, 127), + ], +) +def test_unverified_shapes_use_original_route( + native, monkeypatch, rows, heads, kv_heads, partition, window +): + q = torch.randn((rows, heads, 256), device="cuda", dtype=torch.float16) + kv = torch.randn((2, 1, 1024, kv_heads, 256), device="cuda", dtype=torch.float16) + k, v = kv.to(torch.float8_e4m3fn).view(torch.uint8).unbind(0) + seq = torch.full((rows,), 65, device="cuda", dtype=torch.int32) + table = torch.zeros((rows, 1), device="cuda", dtype=torch.int32) + parts = 1024 // partition + active = torch.full((1,), parts, device="cuda", dtype=torch.int32) + results = [] + before = native.tp2_e4m3_scalar_fast_launch_count() + for enabled in ("0", "1"): + monkeypatch.setenv(FLAG, enabled) + state = workspace(q, parts) + native.decode_paged_fwd( + q, + k, + v, + state[0], + table, + seq, + *state[1:4], + active, + 0.0625, + partition, + parts, + "fp8_e4m3", + 0.5, + 1.25, + window, + -1, + None, + 0, + ) + results.append(state[0]) + assert native.tp2_e4m3_scalar_fast_launch_count() == before + assert torch.equal(results[0].view(torch.int16), results[1].view(torch.int16)) From c009a7dbfabb86d4b72bd0ff0d75673615a35e98 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Tue, 8 Sep 2026 21:13:16 +0800 Subject: [PATCH 02/16] [Doc] Record paired TP2 attention gains and projection screens Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 103 +++++++++++++++++++-- docs/design/sm70_v100_migration_control.md | 16 +++- 2 files changed, 105 insertions(+), 14 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 31faf2514d..3a82b1de18 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -96,8 +96,10 @@ The private u8 implementation is pinned by SHA256 `696545418c6dae261f0bc6a3a530b34464d040de8e404a3069cfd8c2a7762ad3`. The integrated native build is separately pinned by SHA256 `f916e9e370eeb8d865b4de9d8b64f6e66d8831c0b4458dbe3087141dbadc1d19`; -its own GPU and runtime gates must pass before substituting it for the -isolated implementation. +it passes 14 native tests and supplies the fast partition function in the +paired model comparison below. Its retained build-source snapshot precedes +the final changed-line formatting pass; the manifest hashes the actual +as-built source and library. The first contemporaneous control reproduces round cost at 44.986/35.075 ms. Its release1k trajectory has 349 tokens rather than the original startup's @@ -109,17 +111,98 @@ release1k/MBPP28 outputs of 283/297 tokens. Its corresponding control produces 349/270 tokens. The trajectories and acceptance counts differ, so the approximately 20.15%/6.75% latency reductions are provisional performance observations, not accepted quality-preserving gains. A within-startup graph -comparison is used next to keep prefill and projection choices fixed. -The 25 ms target and final promotion remain outstanding. +comparison now keeps prefill and projection choices fixed and isolates the +attention change. The integrated native build passes 14 tests, including exhaustive byte decoding, 262144-token graph replay, FP64 reference, unsupported-shape fallback -and stale-library rejection. The initial QPN2 TP2 projection screen uses -sixteen real matrices from four adjacent layers. Its best working-set median -is 0.685 ms versus 0.930 ms for TurboMind, but several reference-error metrics -grow. This arithmetic candidate is rejected for model use. Source inspection -identifies early FP16 rounding of the global scale as a separate precision -candidate; it requires fresh operator and model gates. +and stale-library rejection. + +### Three-startup paired attention result + +After warmup, each of three independent services alternates five A/B pairs +per fixture. Between quiescent requests, a diagnostic CUDA driver API helper +changes only the executable graph function for the sixteen scalar partition +nodes on each rank. Node arguments, grid, reducer, buffers, prefill, model +weights and that startup's TurboMind choices stay fixed. The candidate +function comes from the pinned integrated native build. Switching and route +verification happen outside request timing; no profiler or per-round tensor +dump is active. This harness is not a service API change. + +| Startup | release1k control / candidate, ms | MBPP28 control / candidate, ms | +| --- | ---: | ---: | +| 3 | 44.808 / 35.823 | 35.051 / 32.694 | +| 4 | 44.897 / 35.797 | 34.953 / 32.678 | +| 5 | 44.872 / 35.930 | 35.318 / 32.861 | +| Median of startup medians | **44.872 / 35.823** | **35.051 / 32.694** | + +Each cell is the median of five request-average complete-round costs. +All fifteen pairs per fixture have identical token IDs, acceptance counters +and natural EOS. Cross-startup controls still differ; this experiment isolates +the attention optimization without claiming to fix the existing variation. + +Pooled endpoint stream intervals have one interval per round, checked against +the round count. These host-observed intervals include delivery jitter: + +| Fixture / mode | Round p50 / p90 / p99, ms | Median TTFT, ms | Median pure decode, tokens/s | +| --- | ---: | ---: | ---: | +| release1k control | 44.772 / 45.159 / 47.196 | 575.877 | 67.947 | +| release1k candidate | 35.789 / 36.063 / 38.027 | 576.364 | 85.219 | +| MBPP28 control | 35.133 / 36.452 / 37.276 | 148.971 | 128.277 | +| MBPP28 candidate | 32.728 / 33.232 / 33.915 | 149.931 | 137.237 | + +Acceptance is reported separately from emitted tokens: + +| Fixture | Startup | Accepted drafts / round, both modes | Emitted tokens / round, both modes | +| --- | ---: | ---: | ---: | +| release1k | 3 | 2.455446 | 3.455446 | +| release1k | 4 | 2.063291 | 3.063291 | +| release1k | 5 | 2.010638 | 3.010638 | +| MBPP28 | 3, 4 | 3.500000 | 4.500000 | +| MBPP28 | 5 | 3.569231 | 4.569231 | + +These paired results admit the attention component for continued experiments. +The approximately 25 ms target, full context sweep and broader quality suite +remain outstanding. The production flag stays off and the PR stays Draft. +Raw reports are `attention-three-start-pair-summary.json`, +`attention-three-start-secondary-metrics.json`, and +`tp2-attention-within-start-{3,4,5}-switch.json` in the campaign results. + +### Projection screening and rejected paths + +Sixteen real matrices from four adjacent target layers cover all five TP2 +physical projection shapes. Inputs are fixed synthetic M8 operands, so these +are operator screens, not live hidden-state or complete-model evidence. +The first QPN2 screen is faster (0.685 versus 0.930 ms per working set) but +increases several FP64-reference error metrics and is rejected for model use. + +TurboMind first combines each group scale with the global scale in FP32 and +rounds that effective scale to FP16. Matching this order, the actual selected +split-K count, and its K64 chunk boundaries produces bitwise-identical outputs +on all sixteen tested matrices, with identical FP64-reference errors. Observed +split counts vary across startup tuning, including 14 and 15; the experiment +reads the selected kernel rather than assuming a fixed count. This is not yet +proof that tuning explains the model's cross-startup variation. + +The exact prepacked variant measures 0.724 versus 0.891 ms, but duplicates +roughly 5.67 GiB of codes per rank plus scales across the full target. It is +not admitted under the frozen memory/context contract. Reusing the TurboMind +code and effective-scale storage avoids that duplication but is slower: + +| Same-working-set comparison | TurboMind, ms | Candidate, ms | Decision | +| --- | ---: | ---: | --- | +| Shared codes, cached loads | 0.891 | 0.940 | Reject | +| Shared codes, streaming loads | 0.896 | 0.921 | Reject | +| Two / four adjacent N tiles | 0.896 | 1.081 / 1.614 | Reject | +| Vector code load plus lane exchange | 0.887 | 0.907 | Reject | + +All these exact variants match the sixteen outputs bit for bit. No slower +variant advances to model testing. The shared layout reader references PR561; +that memory campaign is separate from this attention PR. Expanding the +attention PV unroll from eight to sixteen retains the operator output and +partial-state bits through length 262144 and changes its sixteen-layer median +from 3.888 to 3.652 ms. This small additional gain is not yet a complete-round +result and is not enabled in the published specialization. ## Reproduction and retained negative results diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index e99d400ee4..3f4563d1b4 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -12,12 +12,20 @@ unroll specialization passes exhaustive encoding, graph/state, FP64-reference and sanitizer gates; live attention shadow compares 665321472 elements with zero bit differences. It is restricted to TP2 q8/D256/FP32-partial/P1024. -The first independent-startup candidate measures 35.921/32.706 ms against -44.986/35.075 ms control. Output trajectories differ between startups, including -between controls, so this is not a quality promotion or a 25 ms result. -Keep the feature off while resolving the comparison and remaining cost. +Three independent startups now each contain five attention-only graph-switch +pairs per fixture, with fixed prefill and projection choices within a startup. +Complete-round medians improve from 44.872/35.051 to 35.823/32.694 ms, and all +fifteen pairs per fixture retain identical tokens, acceptance counts and EOS. +Cross-startup control variation remains unresolved; the first unmatched +comparison is retained as provisional evidence. The approximately 25 ms target +and broader quality/context gates are still pending. Keep PR566 Draft and the +production flag off while reducing the remaining cost. Capping inactive partition CTAs was numerically exact but had no speed gain; do not repeat that rejected path without new evidence. +QPN2 TP2 can match TurboMind's operator bits by matching effective-scale +rounding and the observed K64 split/reduction schedule. A full extra code copy +exceeds the memory budget; shared-code, tiled and vector-reader screens remain +slower and are rejected for model use. The TP2 worklog records the comparisons. ## DFlash2 E4M3 FP32 default policy, 2026-09-08 From 5766763659204a967511053af4b145d9b05f3bfb Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Tue, 8 Sep 2026 21:59:01 +0800 Subject: [PATCH 03/16] [Bugfix] Preserve FP32 beta in packed DFlash2 verification Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 48 +++++++ docs/design/sm70_v100_migration_control.md | 8 ++ .../test_sm70_dflash2_packed_gdn_fp32.py | 119 ++++++++++++++++++ .../layers/mamba/gdn/qwen_gdn_linear_attn.py | 6 +- 4 files changed, 180 insertions(+), 1 deletion(-) create mode 100644 tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 3a82b1de18..503f4b70f5 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -204,6 +204,54 @@ partial-state bits through length 262144 and changes its sixteen-layer median from 3.888 to 3.652 ms. This small additional gain is not yet a complete-round result and is not enabled in the published specialization. +### Selective MLP and packed GDN follow-up + +Only the 128 MLP projections per rank can retain a fast duplicate layout +without duplicating the full target. Their codes and scales cost 4.482422 GiB +per rank. Explicitly disabling the unused TP4-only QPN8 rerank request avoids +FP16 head packing on TP2; both actual target/draft heads still use the original +FP16 parameters and FP32 dense logits. The shadow startup loads 16.82 GiB/rank, +retains 7.09 GiB of KV and reports capacity for 332993 tokens, above the frozen +262144 maximum context. This is capacity evidence, not long-context latency. + +The sixteen-MLP working set across both real TP2 shards improves from 1.203 +to 0.963 ms. Four changing-input graph cases per matrix match bitwise. +Memcheck and racecheck pass for every supported split count 1 through 16 plus +32. A complete live shadow covers all 128 MLP projections on each rank: +248068 calls, 22353903616 output elements, zero bit differences and zero +nonfinite outputs. The TurboMind output drives generation. Both fixtures +finish naturally and repeat within that startup; diagnostic timings are +excluded. A complete-round paired performance result is still pending. + +The first MLP graph-switch startup stops before its first request because the +diagnostic queries edges of an unrelated graph and receives invalid argument. +The helper now skips unmarked graphs and uses the edge-data-aware driver API, +rejecting non-default dependencies inside a marked region. A bounded gate +covers empty, one-node and two-node unmarked graphs plus twelve real-matrix +switches. The failed startup is retained and contributes no speed evidence. + +The packed GDN audit also finds a concrete precision mismatch in its existing +entry: the ordinary speculative path explicitly materializes beta in FP32, +while the packed entry relied on the helper's FP16-input default. On a fixed +TP2 q8 case, the old entry differs from the ordinary FP32-beta path in 5623 +output elements and 3142001 state elements, with maximum absolute differences +of 3.8146973e-6 and 1.9565225e-5. The earlier shared-FP16-beta tests therefore +do not admit the actual packed entry. + +The entry now explicitly requests FP32 beta. With the same FP32 gating, +the packed subchain matches output and every pool-state bit through all eight +acceptance selectors and changing-input graph replays. Sixteen distinct layer +states measure 0.712 ms for QKV materialization, recurrence and output copy, +versus 0.507 ms for direct packed recurrence. Common convolution and gating +are outside this operator timing. The feature remains disabled pending live +state and full-round validation; this finding is not attributed as the cause +of historical text-quality changes while the packed feature was disabled. +The actual-entry GPU regression passes for both TP2 and TP4 head geometry, +with FP32 state, gaps between pool slots, all eight selectors, two changing +replays per selector and untouched retired rows. Reproduce with +`.venv/bin/python -m pytest --confcutdir=tests/kernels tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py -q` +(two tests passed). + ## Reproduction and retained negative results Build Flash-V100 from this branch with the same CUDA/Torch/compiler flags and diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index 3f4563d1b4..da9c7de5cb 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -26,6 +26,14 @@ QPN2 TP2 can match TurboMind's operator bits by matching effective-scale rounding and the observed K64 split/reduction schedule. A full extra code copy exceeds the memory budget; shared-code, tiled and vector-reader screens remain slower and are rejected for model use. The TP2 worklog records the comparisons. +Selective MLP packing fits the frozen context after unused TP2 FP16 head +packing is disabled; its two-rank live shadow has zero differences across +22353903616 output elements. Complete-round MLP performance is still pending. +The packed GDN entry also now requests FP32 beta, matching the ordinary +speculative path. The previous implicit FP16 beta changes output and state +bits in a fixed TP2 reproduction. Actual-entry TP2/TP4 tests pass with strided +state and all acceptance selectors; packed GDN remains opt-in pending live +and full-round gates. ## DFlash2 E4M3 FP32 default policy, 2026-09-08 diff --git a/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py new file mode 100644 index 0000000000..5b46860f20 --- /dev/null +++ b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py @@ -0,0 +1,119 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +from types import SimpleNamespace + +import pytest +import torch + +from vllm.model_executor.layers.fla.ops import fused_recurrent_gated_delta_rule +from vllm.model_executor.layers.mamba.gdn.qwen_gdn_linear_attn import ( + QwenGatedDeltaNetAttention, + fused_gdn_gating, +) + + +@pytest.mark.parametrize("tp_size", [2, 4]) +def test_packed_entry_preserves_fp32_beta_and_strided_state(tp_size: int): + if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): + pytest.skip("The packed DFlash2 verifier requires SM70") + torch.manual_seed(20260911) + q_heads, v_heads, dim, tokens = 16 // tp_size, 48 // tp_size, 128, 8 + mixed = ( + torch.randn( + (tokens, (2 * q_heads + v_heads) * dim), + device="cuda", + dtype=torch.float16, + ) + * 0.1 + ) + a = torch.randn((tokens, v_heads), device="cuda", dtype=torch.float16) + b = torch.randn_like(a) + a_log = torch.randn(v_heads, device="cuda", dtype=torch.float32) + bias = torch.randn(v_heads, device="cuda", dtype=torch.float16) + # Keep a gap between pool slots to exercise the native state stride. + control_storage = ( + torch.randn((12, 2, v_heads, dim, dim), device="cuda", dtype=torch.float32) + * 0.02 + ) + candidate_storage = control_storage.clone() + control_state, candidate_state = control_storage[:, 1], candidate_storage[:, 1] + indices = torch.tensor([[3, 9, 4, 8, 2, 6, 1, 5]], device="cuda", dtype=torch.int32) + retired = torch.tensor([0, 7, 10, 11], device="cuda") + accepted = torch.ones(1, device="cuda", dtype=torch.int32) + cu = torch.tensor([0, tokens], device="cuda", dtype=torch.int32) + expected = torch.empty((tokens, v_heads, dim), device="cuda", dtype=torch.float16) + actual = torch.empty_like(expected) + layer = SimpleNamespace( + A_log=a_log, + dt_bias=bias, + num_k_heads=16, + num_v_heads=48, + tp_size=tp_size, + head_k_dim=dim, + head_v_dim=dim, + ) + + def control(): + q, k, v = torch.split( + mixed, [q_heads * dim, q_heads * dim, v_heads * dim], dim=-1 + ) + g, beta = fused_gdn_gating(a_log, a, b, bias, beta_dtype=torch.float32) + output, _ = fused_recurrent_gated_delta_rule( + q=q.contiguous().view(1, tokens, q_heads, dim), + k=k.contiguous().view(1, tokens, q_heads, dim), + v=v.contiguous().view(1, tokens, v_heads, dim), + g=g, + beta=beta, + initial_state=control_state, + inplace_final_state=True, + cu_seqlens=cu, + ssm_state_indices=indices, + num_accepted_tokens=accepted, + use_qk_l2norm_in_kernel=True, + ) + expected.copy_(output.squeeze(0)) + + def candidate(): + QwenGatedDeltaNetAttention._forward_dflash2_packed_gdn_verify( + layer, + mixed_qkv=mixed, + a=a, + b=b, + core_attn_out=actual, + ssm_state=candidate_state, + spec_query_start_loc=cu, + spec_state_indices_tensor=indices, + spec_state_slot_selectors=accepted, + num_spec_decodes=1, + ) + + control() + candidate() + graphs = [] + for run in (control, candidate): + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + run() + graphs.append(graph) + for selector in range(1, 9): + accepted.fill_(selector) + initial = torch.randn_like(control_storage) * 0.02 + control_storage.copy_(initial) + candidate_storage.copy_(initial) + for _ in range(2): + mixed.copy_(torch.randn_like(mixed) * 0.1) + a.copy_(torch.randn_like(a)) + b.copy_(torch.randn_like(b)) + for graph in graphs: + graph.replay() + torch.accelerator.synchronize() + assert torch.equal(actual.view(torch.int16), expected.view(torch.int16)) + assert torch.equal( + candidate_storage.view(torch.int32), control_storage.view(torch.int32) + ) + assert torch.equal(candidate_storage[:, 0], initial[:, 0]) + assert torch.equal( + candidate_storage.index_select(0, retired), + initial.index_select(0, retired), + ) diff --git a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py index fc99c3a880..c70174ca52 100644 --- a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py +++ b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @@ -5202,7 +5202,11 @@ def _forward_dflash2_packed_gdn_verify( ) -> torch.Tensor: num_tokens = mixed_qkv.shape[0] out = core_attn_out[:num_tokens].unsqueeze(1) - g, beta = fused_gdn_gating(self.A_log, a, b, self.dt_bias) + # Match the ordinary speculative verifier's FP32 beta materialization. + # The gating helper otherwise defaults to the FP16 dtype of b. + g, beta = fused_gdn_gating( + self.A_log, a, b, self.dt_bias, beta_dtype=torch.float32 + ) fused_sigmoid_gating_delta_rule_update_mixed_qkv_out( A_log=self.A_log, a=a, From 1aa45e0a5b48d657c33326a8355617ed8a0ef399 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Tue, 8 Sep 2026 23:32:32 +0800 Subject: [PATCH 04/16] [Doc] Record TP2 projection trace and live GDN coverage Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 81 +++++++++++++++++++++- docs/design/sm70_v100_migration_control.md | 15 +++- 2 files changed, 94 insertions(+), 2 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 503f4b70f5..168f68d6eb 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -195,6 +195,7 @@ code and effective-scale storage avoids that duplication but is slower: | Shared codes, streaming loads | 0.896 | 0.921 | Reject | | Two / four adjacent N tiles | 0.896 | 1.081 / 1.614 | Reject | | Vector code load plus lane exchange | 0.887 | 0.907 | Reject | +| Effective-scale-only repack, shared codes | 0.883 | 0.978 | Reject | All these exact variants match the sixteen outputs bit for bit. No slower variant advances to model testing. The shared layout reader references PR561; @@ -221,7 +222,23 @@ Memcheck and racecheck pass for every supported split count 1 through 16 plus 248068 calls, 22353903616 output elements, zero bit differences and zero nonfinite outputs. The TurboMind output drives generation. Both fixtures finish naturally and repeat within that startup; diagnostic timings are -excluded. A complete-round paired performance result is still pending. +excluded. + +One subsequent startup performs five unprofiled A/B pairs per fixture, +switching only marked MLP graph regions between quiescent requests. Both arms +retain the same prefill, attention u8, allocations, selected TurboMind splits +and sampling. Every pair has identical tokens, acceptance and natural EOS: + +| Fixture | Control / candidate complete round, ms | Accepted drafts / round | Emitted tokens / round | +| --- | ---: | ---: | ---: | +| release1k | 35.903174 / 35.413639 | 2.063291 | 3.063291 | +| MBPP28 | 32.767878 / 32.293540 | 3.500000 | 4.500000 | + +The complete-round benefit is only 0.490/0.474 ms. Do not extrapolate the +approximately 20% projection microbenchmark into a multi-millisecond model +gain. This is one startup, not the final three-startup performance gate. +Raw evidence is `mlp-first-paired-summary.json` and +`tp2-mlp-within-start-2-switch.json`. The first MLP graph-switch startup stops before its first request because the diagnostic queries edges of an unrelated graph and receives invalid argument. @@ -252,6 +269,68 @@ replays per selector and untouched retired rows. Reproduce with `.venv/bin/python -m pytest --confcutdir=tests/kernels tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py -q` (two tests passed). +The first live packed-recurrence shadow has no positive coverage and is not +a pass: it assumes state indices `[1,8]`, while the real q8 graph passes a +padded `[8,8]` index buffer and `[8]` selector buffer. The active sequence count +comes from the two-element cumulative-length tensor. Its metadata report is +also written before capture and therefore misses later dispatches. A second +diagnostic correctly slices the active row and observes zero output/state +differences, but its per-layer gate fails: indexing counters by state-pool base +collapses 48 GDN layers into eight shared pool addresses. These diagnostic +failures remain retained; positive coverage must be attributed to actual +layer identity before model admission. + +The third shadow attributes calls by the active GDN layer prefix and pool +pointer, and resets its GPU counters after graph capture. The final admission +snapshot covers all 48 layers on each rank: 93792 calls, 2305032192 output +elements and 295044120576 state elements, with zero output/state bit +differences, nonfinite values or unsupported active calls. Original output +and state continue to drive generation. Both natural-stop fixtures complete; +these shadow timings are excluded from performance evidence. The fail-closed +client now requires 48 named, positive-coverage layers per rank. +Evidence: `tp2-gdn-shadow-3-admission.json` and the per-rank shadow reports. + +### Updated target and draft attribution + +An actual candidate service with selective MLP and exact u8 attention captures +twelve complete rounds; attribution uses the ten interior rounds, on both +ranks. Both warmup and measured requests finish naturally with 242 tokens, +79 rounds and 163 accepted drafts. The profiler wrapper subsequently exits +137 during shutdown. The completed Nsight capture and exported SQLite are +retained, but the job is not recorded as a clean success. No performance +acceptance claim uses this instrumented service. + +The same critical-rank wall window closes as follows: + +| Diagnostic wall component | Mean, ms | +| --- | ---: | +| Complete round interval | 39.235763 | +| GPU interval union inside that window | 36.186996 | +| Uncovered wall interval | 3.048767 | + +Launch-correlated GPU service identifies the next priorities. These service +sums use both ranks and are separate from the wall-clock closure: + +| GPU work | Mean service per rank and round, ms | +| --- | ---: | +| Target graph, total | 26.222108 | +| Target QPN2 MLP projections | 8.903925 | +| Remaining target TurboMind projections | 5.881736 | +| Target exact scalar attention | 3.845147 | +| GDN recurrent core | 1.577858 | +| Draft proposal, total | 6.831827 | +| Target head and sampling | 2.270963 | + +The QPN2 MLP kernel uses 52 registers/thread with no local-memory allocation +in the native resource dump. Register-cap screens preserve all sixteen real +matrix outputs and their FP64-reference errors, but are slower: matched +0.716544 ms, cap 48 at 0.739584 ms and cap 40 at 0.888576 ms. Both are rejected. +Moving effective-scale conversion to preparation also preserves every tested +output bit but increases the working set from 0.708352 to 0.790528 ms. Reject +it before model work. These are new measured negative results, not evidence +of a changed numerical tolerance. Raw reports are `tp2-mlp-trace.json`, +`tp2-qpn2-register-screen.json` and `tp2-qpn2-effective-scale-screen.json`. + ## Reproduction and retained negative results Build Flash-V100 from this branch with the same CUDA/Torch/compiler flags and diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index da9c7de5cb..ba565ff2fe 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -28,12 +28,25 @@ exceeds the memory budget; shared-code, tiled and vector-reader screens remain slower and are rejected for model use. The TP2 worklog records the comparisons. Selective MLP packing fits the frozen context after unused TP2 FP16 head packing is disabled; its two-rank live shadow has zero differences across -22353903616 output elements. Complete-round MLP performance is still pending. +22353903616 output elements. One five-pair startup measures an additional +35.903/32.768 to 35.414/32.294 ms improvement, with identical paired tokens +and acceptance. The approximately 0.5 ms whole-round gain is much smaller +than the projection microbenchmark; three-startup admission remains pending. The packed GDN entry also now requests FP32 beta, matching the ordinary speculative path. The previous implicit FP16 beta changes output and state bits in a fixed TP2 reproduction. Actual-entry TP2/TP4 tests pass with strided state and all acceptance selectors; packed GDN remains opt-in pending live and full-round gates. +The updated trace still attributes about 14.786 ms to target projections +and 6.832 ms to draft work. Register caps 48/40 and precomputed effective +scales preserve operator bits but lose performance and are rejected. Initial +GDN live diagnostics fail coverage because of padded metadata, then pooled +state-pointer attribution; zero differences without per-layer coverage do +not admit the packed route. The corrected shadow now covers 48 layers per +rank, with zero bit differences across 2305032192 output and 295044120576 +state elements. Graph switching and whole-round benefit remain pending. +The TP2 worklog retains both diagnostic failures and the +profile wrapper's shutdown exit 137 separately from usable captured data. ## DFlash2 E4M3 FP32 default policy, 2026-09-08 From 4e443cdb36d9cd783e68018df3dd901cac51f388 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 00:03:49 +0800 Subject: [PATCH 05/16] [Kernel] Read strided QKV in packed DFlash2 verification Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 54 ++++++++++++++++++- docs/design/sm70_v100_migration_control.md | 9 ++++ .../test_sm70_dflash2_packed_gdn_fp32.py | 40 +++++++++++--- .../layers/fla/ops/fused_sigmoid_gating.py | 11 ++-- .../layers/mamba/gdn/qwen_gdn_linear_attn.py | 3 +- 5 files changed, 102 insertions(+), 15 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 168f68d6eb..e741accb86 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -263,11 +263,11 @@ versus 0.507 ms for direct packed recurrence. Common convolution and gating are outside this operator timing. The feature remains disabled pending live state and full-round validation; this finding is not attributed as the cause of historical text-quality changes while the packed feature was disabled. -The actual-entry GPU regression passes for both TP2 and TP4 head geometry, +The initial actual-entry GPU regression passes for both TP2 and TP4 head geometry, with FP32 state, gaps between pool slots, all eight selectors, two changing replays per selector and untouched retired rows. Reproduce with `.venv/bin/python -m pytest --confcutdir=tests/kernels tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py -q` -(two tests passed). +(two initial tests passed; the extended stride suite below contains four). The first live packed-recurrence shadow has no positive coverage and is not a pass: it assumes state indices `[1,8]`, while the real q8 graph passes a @@ -290,6 +290,26 @@ these shadow timings are excluded from performance evidence. The fail-closed client now requires 48 named, positive-coverage layers per rank. Evidence: `tp2-gdn-shadow-3-admission.json` and the per-rank shadow reports. +The first paired GDN service is blocked before its first request because no +marked packed regions are captured. Source inspection identifies a layout +gate: the Qwen3.5 QKV view shares its row with Z/b/a and the convolution writes +in place, while the packed verifier requires a contiguous QKV matrix. The +native mixed-QKV kernel already accepts a row stride, but its Python wrapper +copies the input and then passes the logical width as that stride. + +The opt-in packed entry now accepts contiguous features with a separate row +stride. Its operator wrapper passes the actual stride, retaining the old +copy fallback for non-unit feature strides. Arithmetic, beta/state precision +and state selectors are unchanged. The extended actual-entry suite passes +four TP2/TP4 contiguous/strided cases, including physical QKV row widths 8256 +and 4128, untouched input padding, every acceptance selector, changing graphs, +all state-pool bits and retired rows. A second gate performs 24 graph-arm +switches with padded metadata, the 8256-element QKV row stride and strided +state pools; every output and pool bit matches. CUDA 12.8 memcheck reports +zero errors and racecheck reports zero hazards for that gate. Live route and +full-round validation remain pending; the failed first paired startup supplies +no speed evidence. + ### Updated target and draft attribution An actual candidate service with selective MLP and exact u8 attention captures @@ -331,6 +351,36 @@ it before model work. These are new measured negative results, not evidence of a changed numerical tolerance. Raw reports are `tp2-mlp-trace.json`, `tp2-qpn2-register-screen.json` and `tp2-qpn2-effective-scale-screen.json`. +Further scheduling screens are also rejected. Compiler unroll one/two/eight +measures 0.886784/0.803584/0.732672 ms against 0.711168 ms for the existing +unroll four. Shortening live input/dequant fragments reduces registers from +52 to 48 without local memory, but measures 0.806400 ms against 0.705792 ms. +All sixteen real matrix outputs and FP64-reference errors remain identical. +Improved occupancy potential alone is not a measured speedup. + +### LM-head width and accumulation order + +The trace spends approximately 4.133 ms across the target and draft dense +FP32 heads. A bounded probe keeps real TP2 target-head weights and input +rows fixed, then selects 256 vocabulary rows for recomputation. Default +`torch.mm` changes its effective split-K grid from two to sixteen. On M8/M7, +2001/1739 FP32 output elements change, with maximum differences of +6.4820051e-7/6.1839819e-7. The cuBLASLt log requests split 19 for the narrower +matrix; its resulting kernel grid uses sixteen partitions. Do not treat +default narrowed GEMM as bit-exact candidate reranking. + +A private probe uses the official cuBLASLt algorithm-selection interface to +retain the complete-head algorithm: ID 21, tile ID 5, stage ID 14, split two, +output-type reduction (scheme 4). It uses the observed Torch workspace limit +of 8519680 bytes. M7 and M8 each pass N64/N256/N1024 with three changing input +amplitudes, with zero FP32 bit differences from the selected complete-head +outputs. The N1024 graph is about 0.032 ms; this excludes candidate selection +and weight gathering and is not a complete-head or model speed result. +Candidate coverage, both actual head inputs, memory safety and full-round +admission remain outstanding. No narrowed LM-head route is enabled. +See `head-cublaslt-probe.json`, `head-lt-plan-probe.json` and the +[CUDA 12.8 cuBLASLt reference](https://docs.nvidia.com/cuda/archive/12.8.0/cublas/index.html). + ## Reproduction and retained negative results Build Flash-V100 from this branch with the same CUDA/Torch/compiler flags and diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index ba565ff2fe..aedbd3de06 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -47,6 +47,15 @@ rank, with zero bit differences across 2305032192 output and 295044120576 state elements. Graph switching and whole-round benefit remain pending. The TP2 worklog retains both diagnostic failures and the profile wrapper's shutdown exit 137 separately from usable captured data. +The packed GDN route also rejected the model's QKV view because its row +shares storage with Z/b/a. The opt-in entry and native wrapper now retain +that row stride directly; four TP2/TP4 actual-entry tests pass, including +input padding and complete state preservation. Full-round validation is +still pending. Additional QPN2 unroll/lifetime screens remain slower. +A private LM-head probe confirms that default reduced vocabulary width +changes split-K and FP32 logits; retaining the complete-head cuBLASLt plan +restores bitwise selected logits. This is a reranking building block, with +candidate coverage and model admission outstanding, not a speed claim. ## DFlash2 E4M3 FP32 default policy, 2026-09-08 diff --git a/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py index 5b46860f20..4afe3d8674 100644 --- a/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py +++ b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py @@ -14,19 +14,22 @@ @pytest.mark.parametrize("tp_size", [2, 4]) -def test_packed_entry_preserves_fp32_beta_and_strided_state(tp_size: int): +@pytest.mark.parametrize("strided_qkv", [False, True]) +def test_packed_entry_preserves_fp32_beta_and_strided_state( + tp_size: int, strided_qkv: bool +): if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): pytest.skip("The packed DFlash2 verifier requires SM70") torch.manual_seed(20260911) q_heads, v_heads, dim, tokens = 16 // tp_size, 48 // tp_size, 128, 8 - mixed = ( - torch.randn( - (tokens, (2 * q_heads + v_heads) * dim), - device="cuda", - dtype=torch.float16, - ) - * 0.1 + width = (2 * q_heads + v_heads) * dim + projection_width = (2 * q_heads + 2 * v_heads) * dim + 2 * v_heads + row_stride = (projection_width + 31) // 32 * 32 if strided_qkv else width + mixed_storage = torch.full( + (tokens, row_stride), -3.0, device="cuda", dtype=torch.float16 ) + mixed = mixed_storage[:, :width] + mixed.copy_(torch.randn_like(mixed) * 0.1) a = torch.randn((tokens, v_heads), device="cuda", dtype=torch.float16) b = torch.randn_like(a) a_log = torch.randn(v_heads, device="cuda", dtype=torch.float32) @@ -52,6 +55,26 @@ def test_packed_entry_preserves_fp32_beta_and_strided_state(tp_size: int): tp_size=tp_size, head_k_dim=dim, head_v_dim=dim, + enable_sm70_dflash2_fused_gdn_verify=True, + ) + metadata = SimpleNamespace( + spec_sequence_masks=torch.ones(1, device="cuda", dtype=torch.bool), + num_spec_decodes=1, + num_prefills=0, + num_decodes=0, + ddtree_parent_ids=None, + spec_query_start_loc=cu, + spec_state_indices_tensor=indices, + spec_state_slot_selectors=accepted, + ) + assert QwenGatedDeltaNetAttention._can_use_dflash2_packed_gdn_verify( + layer, + mixed_qkv=mixed, + a=a, + b=b, + core_attn_out=actual, + ssm_state=candidate_state, + attn_metadata=metadata, ) def control(): @@ -117,3 +140,4 @@ def candidate(): candidate_storage.index_select(0, retired), initial.index_select(0, retired), ) + assert torch.all(mixed_storage[:, width:] == -3.0) diff --git a/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py b/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py index 49f1e3aad1..93c90847e4 100644 --- a/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py +++ b/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py @@ -652,7 +652,9 @@ def fused_sigmoid_gating_delta_rule_update_mixed_qkv_out( """ if mixed_qkv.ndim != 2: raise ValueError("mixed_qkv must have shape [T, qkv_hidden].") - if not mixed_qkv.is_contiguous(): + # Qwen's QKV view shares rows with Z/b/a; convolution preserves that + # row stride. The native mixed-QKV loader already accepts QKV_STRIDE. + if mixed_qkv.stride(1) != 1: mixed_qkv = mixed_qkv.contiguous() if cu_seqlens is None: raise ValueError("cu_seqlens is required for mixed_qkv_out.") @@ -672,11 +674,12 @@ def fused_sigmoid_gating_delta_rule_update_mixed_qkv_out( q_size = H * K k_size = H * K v_size = HV * V - qkv_stride = q_size + k_size + v_size - if mixed_qkv.shape[1] != qkv_stride: + qkv_width = q_size + k_size + v_size + if mixed_qkv.shape[1] != qkv_width: raise ValueError( - f"mixed_qkv width {mixed_qkv.shape[1]} != expected {qkv_stride}." + f"mixed_qkv width {mixed_qkv.shape[1]} != expected {qkv_width}." ) + qkv_stride = mixed_qkv.stride(0) if out.shape != (T, 1, HV, V): raise ValueError(f"out must have shape {(T, 1, HV, V)}, got {out.shape}.") if scale is None: diff --git a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py index c70174ca52..1ebfedcada 100644 --- a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py +++ b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @@ -5181,7 +5181,8 @@ def _can_use_dflash2_packed_gdn_verify( # The supported verifier contract keeps recurrent state in FP32; # an explicit FP16 cache override is also supported. and ssm_state.dtype in (torch.float16, torch.float32) - and mixed_qkv.is_contiguous() + and mixed_qkv.ndim == 2 + and mixed_qkv.stride(1) == 1 and a.is_contiguous() and b.is_contiguous() and core_attn_out.is_contiguous() From ec8e14a42cbe44e7cd611207a398474ee9ebd6ee Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 01:14:42 +0800 Subject: [PATCH 06/16] [Doc] Record paired TP2 GDN gains and rejected fusion Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 72 +++++++++++++++++++++- docs/design/sm70_v100_migration_control.md | 16 ++++- 2 files changed, 82 insertions(+), 6 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index e741accb86..1b124810d2 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -306,9 +306,48 @@ and 4128, untouched input padding, every acceptance selector, changing graphs, all state-pool bits and retired rows. A second gate performs 24 graph-arm switches with padded metadata, the 8256-element QKV row stride and strided state pools; every output and pool bit matches. CUDA 12.8 memcheck reports -zero errors and racecheck reports zero hazards for that gate. Live route and -full-round validation remain pending; the failed first paired startup supplies -no speed evidence. +zero errors and racecheck reports zero hazards for that gate. The failed first +paired startup supplies no speed evidence; the corrected route and three +subsequent paired startups are reported below. + +### Three-startup paired packed GDN result + +The corrected route captures all 48 GDN regions on each rank. The actual q8 +QKV view has shape `[8,5120]` and row stride 8256. Three independent startups +each run five alternating A/B pairs per fixture, holding exact attention u8, +selective MLP, prefill, graph buffers and each startup's projection choices +fixed. Only the packed GDN region changes between quiescent requests. Markers +and the inactive arm are disabled before replay; no profiler or tensor dump +runs during these measurements. + +| Startup | release1k control / candidate, ms | MBPP28 control / candidate, ms | +| --- | ---: | ---: | +| 2 | 36.487586 / 34.678536 | 33.492413 / 31.710558 | +| 3 | 36.305965 / 34.508011 | 33.299356 / 31.508836 | +| 4 | 36.362439 / 34.528673 | 33.312182 / 31.477274 | +| Median of startup medians | **36.362439 / 34.528673** | **33.312182 / 31.508836** | + +Each entry is the median of five request-average complete-round costs. All +15 measured pairs per fixture, and each startup's warmup pair, retain token +IDs, acceptance counters and natural EOS. GDN adds a paired whole-round +benefit of 1.833766/1.803346 ms. The startup controls still vary; this isolates +the GDN change and does not resolve the pre-existing repeatability issue. + +| Fixture / mode | Host round p50 / p90 / p99, ms | Median TTFT, ms | Median pure decode, tokens/s | +| --- | ---: | ---: | ---: | +| release1k control | 36.334 / 36.578 / 38.608 | 575.826 | 82.220 | +| release1k candidate | 34.526 / 34.778 / 36.907 | 575.153 | 86.480 | +| MBPP28 control | 33.313 / 33.883 / 34.619 | 147.617 | 144.241 | +| MBPP28 candidate | 31.526 / 32.105 / 32.897 | 147.968 | 152.346 | + +Host intervals include endpoint delivery jitter and are checked against the +round count. Accepted drafts and emitted tokens are reported separately in +`gdn-three-start-pair-summary.json` and +`tp2-gdn-within-start-{2,3,4}-switch.json`. Both modes share each startup's +acceptance exactly; MBPP28 accepted drafts per round are 3.845070, 3.569231 +and 4.306122, with emitted tokens per round 4.845070, 4.569231 and 5.306122. +The candidate remains above 25 ms on both fixtures. Wider quality/context +gates remain outstanding, so these results do not enable a production default. ### Updated target and draft attribution @@ -358,6 +397,33 @@ unroll four. Shortening live input/dequant fragments reduces registers from All sixteen real matrix outputs and FP64-reference errors remain identical. Improved occupancy potential alone is not a measured speedup. +Changing only weight block placement also loses performance. K-major block +interleaving measures 0.807424 ms and 128-byte N-tile pitch padding 0.728832 ms, +against 0.708608 ms for the existing matched layout. Bounded L2 prefetch eight +groups ahead measures 0.794112 ms against 0.711680 ms. All sixteen tested +matrix outputs and FP64-reference errors remain identical, and padding is +untouched. Reject all three before model testing. These are latency +hypotheses tested by timing; unavailable NCU counters do not establish a +specific stall or cache-bank cause. Reports are +`tp2-qpn2-memory-layout-screen.json` and `tp2-qpn2-prefetch8-screen.json`. + +### Existing communication fusion screening + +The existing TP2 all-reduce/Gemma RMSNorm fusion is compared with the actual +DFlash2 Triton normalization, using real layer-0 norm weights, FP16 `[8,5120]` +rank inputs and FP32 residuals. Five fixed operand amplitudes run on both +ranks through graphs. Residual bits match, but nonzero cases differ in 1 to 4 +normalized FP16 elements. Some FP64 relative-L2 errors also increase slightly; +this is not accepted as a new tolerance. + +The fusion is slower on both ranks: 0.021990/0.022349 ms per graph versus +0.018048/0.018278 ms for the existing collective plus DFlash2 norm. These are +operator timings, not whole-round gains. Reject this fusion for the current +TP2 route. The native CUB reduction and current Triton variance reduction do +not share an established arithmetic order. Do not attribute historical text +quality changes to this disabled candidate. Evidence is +`tp2-fused-comm-norm-rank{0,1}.json`. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index aedbd3de06..b8da908768 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -44,14 +44,24 @@ GDN live diagnostics fail coverage because of padded metadata, then pooled state-pointer attribution; zero differences without per-layer coverage do not admit the packed route. The corrected shadow now covers 48 layers per rank, with zero bit differences across 2305032192 output and 295044120576 -state elements. Graph switching and whole-round benefit remain pending. +state elements. The corrected graph comparison below establishes a +whole-round benefit on the paired short-context fixtures. The TP2 worklog retains both diagnostic failures and the profile wrapper's shutdown exit 137 separately from usable captured data. The packed GDN route also rejected the model's QKV view because its row shares storage with Z/b/a. The opt-in entry and native wrapper now retain that row stride directly; four TP2/TP4 actual-entry tests pass, including -input padding and complete state preservation. Full-round validation is -still pending. Additional QPN2 unroll/lifetime screens remain slower. +input padding and complete state preservation. Three independent startups +now each run five GDN A/B pairs per fixture, with attention and exact MLP +held fixed. Complete-round medians improve from 36.362439/33.312182 to +34.528673/31.508836 ms, with identical paired tokens, acceptance and EOS. +All 48 GDN regions per rank hit the strided route. This establishes a +1.833766/1.803346 ms GDN benefit; 25 ms and wider admission remain pending. +Additional QPN2 unroll/lifetime, block-interleaving, pitch-padding and L2 +prefetch screens preserve operator bits but remain slower. Do not repeat +them without changed evidence. Existing TP2 communication/Gemma fusion also +loses performance and changes 1 to 4 normalized FP16 elements per nonzero +test case versus the actual DFlash2 Triton norm; retain the existing route. A private LM-head probe confirms that default reduced vocabulary width changes split-K and FP32 logits; retaining the complete-head cuBLASLt plan restores bitwise selected logits. This is a reranking building block, with From 1cd8f7a0b97da797e718c1d78ad0c62b4821a22a Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 11:26:25 +0800 Subject: [PATCH 07/16] [Doc] Record TP2 live layout gates and rejected kernels Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 90 ++++++++++++++++++++++ docs/design/sm70_v100_migration_control.md | 12 +++ 2 files changed, 102 insertions(+) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 1b124810d2..cfeea8e876 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -424,6 +424,96 @@ not share an established arithmetic order. Do not attribute historical text quality changes to this disabled candidate. Evidence is `tp2-fused-comm-norm-rank{0,1}.json`. +### Packed-route trace and next projection scope, 2026-09-09 + +A fresh service uses the actual packed GDN route with exact MLP and attention +u8. Both warmup and measured release1k requests finish naturally with 280 +tokens, 96 rounds and 184 accepted drafts. Capture, SQLite export and owned +service shutdown complete with exit zero. Ten interior rounds on both ranks +give a critical-rank interval of 37.306504 ms, GPU union of 35.103202 ms and +uncovered interval of 2.203302 ms. These remain instrumented diagnostics. + +| GPU work | Mean service per rank and round, ms | +| --- | ---: | +| Target graph | 25.422060 | +| Target QPN2 MLP | 8.906423 | +| Remaining target TurboMind projections | 5.881097 | +| Target scalar attention | 3.848959 | +| Draft proposal | 6.699200 | +| Target head and sampling | 2.270655 | +| All-phase gather/scatter/copy | 1.012917 | + +The packed GDN kernel launches 192 CTAs of 32 threads, with 128 registers per +thread and zero static shared memory. Its mean invocation is 32.665 us. A +separate q8 V-tile screen retains one warp and the original K dimension. +BV16/8/4/2/32 all preserve output and every state bit through eight acceptance +selectors and two changing replays, including strided QKV/state and retired +rows. Sixteen distinct layer-state working sets measure respectively +0.525824/0.455936/0.398592/0.392832/0.793984 ms. BV2 passes memcheck, racecheck +and 24 graph-arm switches with padded metadata. Its subsequent live shadow +covers all 48 GDN layers on each rank: 81120 calls compare 1993605120 output +and 255181455360 state elements with zero bit differences, nonfinite values +or unsupported active calls. The ordinary recurrence supplies the outputs +and states used for generation; diagnostic latency is excluded. Complete +rounds against packed BV16 remain pending. Independent multi-request GDN +measurements do not admit this route. Raw evidence is `tp2-packed-trace.json`, +`tp2-gdn-bv-screen.json`, `gdn-bv2-graph-switch-gate.json` and +`tp2-gdn-bv2-shadow-1-admission.json`. + +A separate collective launch-geometry screen keeps the original peer protocol, +rank reduction order and DFlash2 normalization. All tested q7/q8 outputs, +residuals, signed-zero/cancellation inputs and changing graph replays match. +The best q8 median improves by only about 0.33 us per invocation; no model +gain or production setting is established. The first diagnostic stops before +its first replay because the retained graph was not explicitly instantiated; +the corrected run passes and the failed run remains excluded. + +Granular real-weight projection timing identifies a smaller next layout +scope: all 64 target output projections and the 16 attention QKV projections. +Together they require 896532480 additional bytes per rank, compared with the +existing MLP layout's 4812963840 bytes. Their real-shard operator gates pass +four changing-input graph cases and canaries; memcheck and racecheck cover +both TP2 shards and all split counts 1 through 16 plus 32. Its first live +shadow startup fails before generation: available KV is 5.14 GiB, below the +5.58 GiB needed for the unchanged 262144 maximum length. Both ranks prepare +208 projections, but there is no live quality or speed result. The static +layout estimate alone did not establish the full runtime memory budget. + +The next private candidate instead retains 64 MLP down projections and all +128 non-MLP projections; gate/up stays on TurboMind. This uses 3642163200 +layout bytes per rank, less than the previous MLP-only candidate. GDN input +weights/scales retain zero-filled padding from 8240 to 8256 columns, and the +producer output stride is unchanged. Eight real non-MLP shard cases pass +changing-input graphs, full output bits, canaries, memcheck and racecheck +across all supported splits; the MLP down gates are retained. The live shadow +now covers all 192 projections on both ranks: 376904 calls compare +18317493248 output elements with zero bit differences or nonfinite values. +Original outputs drive generation. The observed KV budget is 8870215885 and +8874410189 bytes, about 8.26 GiB per rank, with maximum length 262144 and +memory utilization 0.8 unchanged. The paired speed comparison retains the +same allocation in both arms and switches those 192 projections against +TurboMind; it does not compare separate startups or allocate both complete +layout choices. Complete-round gain remains pending. The source kernels and +layouts remain private experiments. Evidence is +`tp2-balanced-shadow-1-admission.json` and the two rank memory reports. + +An E2M1 register-permutation decoder is also exact on the sixteen-matrix +screen but slower: 0.763648 versus 0.712960 ms. A separate scale-folding probe +initially misreads the half constant `0x5c00` as 64 instead of 256; it changes +1139840 output elements and is rejected before any model use. That diagnostic +failure is retained independently of its corrected probe. Neither decoder +experiment changes repository kernels or production behavior. The corrected +256-factor probe restores all sixteen matrix outputs and FP64-reference +errors, but is slower: 0.774144 versus 0.705024 ms, and is also rejected. + +A separate scheduling probe maps each existing logical split to its own +32-thread CTA, then reduces explicit FP32 partials in the original order. +All sixteen real-matrix outputs and FP64-reference errors match, and scratch +canaries remain intact. The working-set median is 0.886016 ms versus +0.711936 ms for the matched QPN2 kernel, so the extra launch/workspace path +is rejected before model use. The result does not establish an instruction +stall diagnosis; hardware performance counters remain unavailable. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index b8da908768..8e533d8dd3 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -66,6 +66,18 @@ A private LM-head probe confirms that default reduced vocabulary width changes split-K and FP32 logits; retaining the complete-head cuBLASLt plan restores bitwise selected logits. This is a reranking building block, with candidate coverage and model admission outstanding, not a speed claim. +The clean packed-route trace still attributes 14.788 ms to target projections +and 3.849 ms to attention. A smaller GDN V tile (BV2) passes the actual q8 +state/graph and sanitizer gates, then a two-rank live shadow with zero output +and state differences; its full-round comparison against BV16 is pending. +Adding 80 projections to the MLP layout fails the frozen 262144 context +memory budget before generation. A replacement layout packs 64 MLP down and +128 non-MLP projections instead, retaining gate/up on TurboMind. All 192 +layers per rank pass live bitwise comparison, and measured KV budget is +about 8.26 GiB per rank. Complete-round gain is not yet established. +E2M1 decoder changes and external logical-split CTAs are slower than the +matched QPN2 kernel and remain rejected; retain the diagnostic scale-constant +error separately from the corrected but slower implementation. ## DFlash2 E4M3 FP32 default policy, 2026-09-08 From 3bb92fe3eb86ca47a9daf3f8f03edb7e4b0fbf29 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 12:10:38 +0800 Subject: [PATCH 08/16] [Kernel] Add an opt-in TP2 q8 GDN schedule Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 61 ++++++++++++++++++- docs/design/sm70_v100_migration_control.md | 11 +++- .../test_sm70_dflash2_packed_gdn_fp32.py | 27 +++++++- vllm/envs.py | 6 ++ .../layers/fla/ops/fused_sigmoid_gating.py | 19 ++++++ .../layers/mamba/gdn/qwen_gdn_linear_attn.py | 8 +++ 6 files changed, 126 insertions(+), 6 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index cfeea8e876..787b3b5a76 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -454,12 +454,42 @@ and 24 graph-arm switches with padded metadata. Its subsequent live shadow covers all 48 GDN layers on each rank: 81120 calls compare 1993605120 output and 255181455360 state elements with zero bit differences, nonfinite values or unsupported active calls. The ordinary recurrence supplies the outputs -and states used for generation; diagnostic latency is excluded. Complete -rounds against packed BV16 remain pending. Independent multi-request GDN -measurements do not admit this route. Raw evidence is `tp2-packed-trace.json`, +and states used for generation; diagnostic latency is excluded. Independent +multi-request GDN measurements do not admit this route. Raw evidence is +`tp2-packed-trace.json`, `tp2-gdn-bv-screen.json`, `gdn-bv2-graph-switch-gate.json` and `tp2-gdn-bv2-shadow-1-admission.json`. +Three independent startups then each run five BV16/BV2 graph-switch pairs +per fixture, holding the exact MLP projection route, attention u8, prefill +and graph allocations fixed. All fifteen measured pairs and the warmup +pairs retain identical token IDs, acceptance counters and natural EOS. + +| Fixture | Startup | BV16 complete round, ms | BV2 complete round, ms | +| --- | ---: | ---: | ---: | +| release1k | 1 | 34.496555 | 34.132393 | +| release1k | 2 | 34.429623 | 34.174315 | +| release1k | 3 | 34.576248 | 34.201821 | +| MBPP28 | 1 | 31.288474 | 30.914687 | +| MBPP28 | 2 | 31.708845 | 31.277764 | +| MBPP28 | 3 | 31.530112 | 31.126483 | + +The medians of startup medians improve by 0.322239/0.403629 ms to +34.174315/31.126483 ms. Candidate host-observed round p50/p90/p99 are +34.102/34.415/36.301 ms for release1k and 31.060/31.755/33.735 ms for MBPP28. +Median TTFT is 575.230/148.479 ms and pure decode is 87.902/145.594 tokens/s. +Acceptance and emitted counts are reported separately for each startup in +`gdn-bv2-three-start-pair-summary.json`; cross-startup trajectories still +vary. These are unprofiled paired results, not a 25 ms or broad quality gate. +The source now exposes `VLLM_SM70_DFLASH2_TP2_GDN_BV2`, default off and +dependent on the packed verifier flag. It admits only TP2 q8, H8/HV24, +K/V128, FP16 input/output, FP32 state and precomputed gating, retaining the +original recurrent arithmetic and stage count. Eight actual-entry GPU tests +pass, including the original TP4 BV8 fallback with the new flag requested. +The first test revision incorrectly expected TP4 BV16; every output/state +check passed, and only that launch assertion failed. The corrected fixture +and failed evidence are retained. The integrated model A/B remains pending. + A separate collective launch-geometry screen keeps the original peer protocol, rank reduction order and DFlash2 normalization. All tested q7/q8 outputs, residuals, signed-zero/cancellation inputs and changing graph replays match. @@ -496,6 +526,16 @@ TurboMind; it does not compare separate startups or allocate both complete layout choices. Complete-round gain remains pending. The source kernels and layouts remain private experiments. Evidence is `tp2-balanced-shadow-1-admission.json` and the two rank memory reports. +The first paired startup stops before generation because the diagnostic +counts both the 192-region full graph and prefill piecewise graphs. That +failure is retained. The corrected tool selects exactly one full graph, +checks 192 unique prepared weight pointers and layer prefixes, and leaves +all prefill pieces on the control path. The corrected startup runs five +pairs per fixture with exact tokens, acceptance and natural EOS. Complete +rounds improve from 35.178025 to 33.911462 ms on release1k and 32.107532 to +30.847108 ms on MBPP28. Both arms retain packed BV16 and attention u8; the +control uses TurboMind for all target projections. This is one startup and +does not establish a paired comparison with the previous MLP-only layout. An E2M1 register-permutation decoder is also exact on the sixteen-matrix screen but slower: 0.763648 versus 0.712960 ms. A separate scale-folding probe @@ -514,6 +554,21 @@ canaries remain intact. The working-set median is 0.886016 ms versus is rejected before model use. The result does not establish an instruction stall diagnosis; hardware performance counters remain unavailable. +An attention address-reuse screen retains the original shared-memory size, +QK/softmax, ascending PV FMA and FP32 partition/reducer. Forty-four cases +across page sizes 1648/3296 preserve output and valid partial/statistic bits +through 262144 tokens, with unchanged FP64-reference errors. It has no +stable working-set gain: 3.921664 versus u8's 3.859456 ms on the actual +1648-token page, and 3.902720 versus 3.944448 ms on page3296. It is not +promoted. Reusing the existing grouped FP32 kernel separately for both TP2 +KV heads also fails the strict arithmetic gate: all nine tested cases change +some FP16 outputs and some expand FP64 relative-L2 error. Timing and model +admission are skipped. These results do not establish a text-quality cause. +A separate eight-value PV staging kernel also preserves all forty-four +output/partial/statistic comparisons, but increases the page1648 working +set from 3.858176 to 4.857088 ms and page3296 from 3.941120 to 4.908800 ms. +It is rejected before model use. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense diff --git a/docs/design/sm70_v100_migration_control.md b/docs/design/sm70_v100_migration_control.md index 8e533d8dd3..6f089d3856 100644 --- a/docs/design/sm70_v100_migration_control.md +++ b/docs/design/sm70_v100_migration_control.md @@ -69,12 +69,19 @@ candidate coverage and model admission outstanding, not a speed claim. The clean packed-route trace still attributes 14.788 ms to target projections and 3.849 ms to attention. A smaller GDN V tile (BV2) passes the actual q8 state/graph and sanitizer gates, then a two-rank live shadow with zero output -and state differences; its full-round comparison against BV16 is pending. +and state differences. Three independent five-pair startups preserve tokens, +acceptance and natural EOS and improve median complete rounds from +34.496555/31.530112 to 34.174315/31.126483 ms. Its source flag +`VLLM_SM70_DFLASH2_TP2_GDN_BV2` remains off, admits only the FP32-state TP2 q8 +contract, and passes eight integrated-entry GPU tests including TP4 fallback. Adding 80 projections to the MLP layout fails the frozen 262144 context memory budget before generation. A replacement layout packs 64 MLP down and 128 non-MLP projections instead, retaining gate/up on TurboMind. All 192 layers per rank pass live bitwise comparison, and measured KV budget is -about 8.26 GiB per rank. Complete-round gain is not yet established. +about 8.26 GiB per rank. A corrected full-graph-only comparison measures +35.178025/32.107532 to 33.911462/30.847108 ms in one startup, with all five +pairs per fixture exact; two additional startups remain required. The first +diagnostic overcounts prefill piecewise graphs and stops before generation. E2M1 decoder changes and external logical-split CTAs are slower than the matched QPN2 kernel and remain rejected; retain the diagnostic scale-constant error separately from the corrected but slower implementation. diff --git a/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py index 4afe3d8674..1aa9d90c16 100644 --- a/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py +++ b/tests/kernels/test_sm70_dflash2_packed_gdn_fp32.py @@ -1,6 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +import importlib from types import SimpleNamespace import pytest @@ -15,12 +16,32 @@ @pytest.mark.parametrize("tp_size", [2, 4]) @pytest.mark.parametrize("strided_qkv", [False, True]) +@pytest.mark.parametrize("use_bv2", [False, True]) def test_packed_entry_preserves_fp32_beta_and_strided_state( - tp_size: int, strided_qkv: bool + tp_size: int, strided_qkv: bool, use_bv2: bool, monkeypatch: pytest.MonkeyPatch ): if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): pytest.skip("The packed DFlash2 verifier requires SM70") torch.manual_seed(20260911) + fused = importlib.import_module( + "vllm.model_executor.layers.fla.ops.fused_sigmoid_gating" + ) + original_kernel = fused.fused_sigmoid_gating_delta_rule_update_kernel + launches = [] + + class RecordingKernel: + def __getitem__(self, grid): + launch = original_kernel[grid] + + def record(*args, **kwargs): + launches.append((grid, kwargs["BV"], kwargs["num_warps"])) + return launch(*args, **kwargs) + + return record + + monkeypatch.setattr( + fused, "fused_sigmoid_gating_delta_rule_update_kernel", RecordingKernel() + ) q_heads, v_heads, dim, tokens = 16 // tp_size, 48 // tp_size, 128, 8 width = (2 * q_heads + v_heads) * dim projection_width = (2 * q_heads + 2 * v_heads) * dim + 2 * v_heads @@ -56,6 +77,7 @@ def test_packed_entry_preserves_fp32_beta_and_strided_state( head_k_dim=dim, head_v_dim=dim, enable_sm70_dflash2_fused_gdn_verify=True, + enable_sm70_dflash2_tp2_gdn_bv2=use_bv2, ) metadata = SimpleNamespace( spec_sequence_masks=torch.ones(1, device="cuda", dtype=torch.bool), @@ -141,3 +163,6 @@ def candidate(): initial.index_select(0, retired), ) assert torch.all(mixed_storage[:, width:] == -3.0) + expected_bv = (2 if use_bv2 else 16) if tp_size == 2 else 8 + assert launches + assert set(launches) == {((1, dim // expected_bv, v_heads), expected_bv, 1)} diff --git a/vllm/envs.py b/vllm/envs.py index 3bcb9e35b4..b3a8248800 100755 --- a/vllm/envs.py +++ b/vllm/envs.py @@ -241,6 +241,7 @@ VLLM_SM70_DFLASH2_GDN_METADATA_SHADOW: bool = False VLLM_SM70_DFLASH2_GDN_SYNC_ASSERT: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_VERIFY: bool = False + VLLM_SM70_DFLASH2_TP2_GDN_BV2: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_NORM: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT: bool = False VLLM_SM70_DFLASH2_FUSED_SMALLQ_METADATA: bool = False @@ -2223,6 +2224,11 @@ def _resolve_rust_frontend_path() -> str | None: "VLLM_SM70_DFLASH2_FUSED_GDN_VERIFY": lambda: bool( int(os.getenv("VLLM_SM70_DFLASH2_FUSED_GDN_VERIFY", "0")) ), + # Independently gated q8/TP2 packed GDN schedule; other shapes retain the + # accepted recurrent launch geometry. + "VLLM_SM70_DFLASH2_TP2_GDN_BV2": lambda: bool( + int(os.getenv("VLLM_SM70_DFLASH2_TP2_GDN_BV2", "0")) + ), # Route compatible target GDN output gates through the existing one-pass # CUDA RMSNormGated implementation. This remains an explicit opt-in. "VLLM_SM70_DFLASH2_FUSED_GDN_NORM": lambda: bool( diff --git a/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py b/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py index 93c90847e4..0d08424e9b 100644 --- a/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py +++ b/vllm/model_executor/layers/fla/ops/fused_sigmoid_gating.py @@ -643,6 +643,7 @@ def fused_sigmoid_gating_delta_rule_update_mixed_qkv_out( match_recurrent_numerics: bool = False, precomputed_g: torch.Tensor | None = None, precomputed_beta: torch.Tensor | None = None, + sm70_tp2_q8_bv2: bool = False, ) -> tuple[torch.Tensor, torch.Tensor]: """Mixed-QKV update that writes into a caller-provided output buffer. @@ -728,6 +729,24 @@ def fused_sigmoid_gating_delta_rule_update_mixed_qkv_out( mixed_qkv.device, match_recurrent_schedule=match_recurrent_schedule, ) + if ( + sm70_tp2_q8_bv2 + and (N, T, H, HV, K, V) == (1, 8, 8, 24, 128, 128) + and (BV, num_warps) == (16, 1) + and match_recurrent_schedule + and match_recurrent_numerics + and use_precomputed_gating + and kernel_a.dtype == kernel_b.dtype == torch.float32 + and initial_state.dtype == torch.float32 + and mixed_qkv.dtype == out.dtype == torch.float16 + and num_accepted_tokens is not None + and use_qk_l2norm_in_kernel + and not quantize_state_each_step + and _use_sm70_fused_sigmoid_schedule(mixed_qkv.device) + ): + # Same K128 reduction, one warp, gating and recurrent arithmetic. + # Split the independent V columns across more CTAs for this q8 case. + BV = 2 NK, NV = triton.cdiv(K, BK), triton.cdiv(V, BV) assert NK == 1, "NK > 1 is not supported yet" diff --git a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py index 1ebfedcada..74403a80f1 100644 --- a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py +++ b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @@ -2433,6 +2433,11 @@ def __init__( and current_platform.is_device_capability(70) and _is_dflash2_spec_config(vllm_config) ) + self.enable_sm70_dflash2_tp2_gdn_bv2 = bool( + self.enable_sm70_dflash2_fused_gdn_verify + and self.tp_size == 2 + and envs.VLLM_SM70_DFLASH2_TP2_GDN_BV2 + ) self.enable_sm70_dflash2_fused_gdn_norm = bool( envs.VLLM_SM70_DFLASH2_FUSED_GDN_NORM and current_platform.is_device_capability(70) @@ -5234,6 +5239,9 @@ def _forward_dflash2_packed_gdn_verify( quantize_state_each_step=False, match_recurrent_schedule=True, match_recurrent_numerics=True, + sm70_tp2_q8_bv2=getattr( + self, "enable_sm70_dflash2_tp2_gdn_bv2", False + ), ) return out.transpose(0, 1) From a4cbe02912e0365535571ab32a81184736843468 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 14:34:18 +0800 Subject: [PATCH 09/16] [Doc] Record TP2 single-layout performance and corrected gates Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 93 ++++++++++++++++++++++-- 1 file changed, 88 insertions(+), 5 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 787b3b5a76..84681450a5 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -488,7 +488,12 @@ original recurrent arithmetic and stage count. Eight actual-entry GPU tests pass, including the original TP4 BV8 fallback with the new flag requested. The first test revision incorrectly expected TP4 BV16; every output/state check passed, and only that launch assertion failed. The corrected fixture -and failed evidence are retained. The integrated model A/B remains pending. +and failed evidence are retained. A subsequent source-integrated model A/B +verifies the constructor flag on all 48 GDN layers per rank and toggles the +source guard. All five pairs and warmup per fixture retain token IDs, +acceptance and natural EOS. Release1k improves from 34.433621 to 34.063787 ms +and MBPP28 from 31.428823 to 31.008482 ms. This confirms the source integration; +it is a separate startup from the preceding three-start cohort. A separate collective launch-geometry screen keeps the original peer protocol, rank reduction order and DFlash2 normalization. All tested q7/q8 outputs, @@ -523,7 +528,7 @@ Original outputs drive generation. The observed KV budget is 8870215885 and memory utilization 0.8 unchanged. The paired speed comparison retains the same allocation in both arms and switches those 192 projections against TurboMind; it does not compare separate startups or allocate both complete -layout choices. Complete-round gain remains pending. The source kernels and +layout choices. The source kernels and layouts remain private experiments. Evidence is `tp2-balanced-shadow-1-admission.json` and the two rank memory reports. The first paired startup stops before generation because the diagnostic @@ -533,9 +538,57 @@ checks 192 unique prepared weight pointers and layer prefixes, and leaves all prefill pieces on the control path. The corrected startup runs five pairs per fixture with exact tokens, acceptance and natural EOS. Complete rounds improve from 35.178025 to 33.911462 ms on release1k and 32.107532 to -30.847108 ms on MBPP28. Both arms retain packed BV16 and attention u8; the -control uses TurboMind for all target projections. This is one startup and -does not establish a paired comparison with the previous MLP-only layout. +30.847108 ms on MBPP28 in the first corrected startup. Three independent +startups now pass all fifteen measured pairs and warmup per fixture. The +median of startup medians is 35.163780 to 33.911462 ms for release1k and +32.052971 to 30.844085 ms for MBPP28. Candidate round p50/p90/p99 are +33.878/34.173/36.521 and 30.791/31.315/31.913 ms, respectively. Median warm +TTFT is 578.279/147.689 ms and pure decode is 89.938/147.626 tokens/s. +Both arms retain packed BV16 and attention u8; the control uses TurboMind +for all target projections. This does not establish a paired comparison +with the previous MLP-only layout. Raw per-start acceptance and emitted +counts remain separate in `balanced-three-start-pair-summary.json`. + +A subsequent private layout experiment stores one persistent code buffer, +761200640 bytes of extra QPN2 scales and a 44564736-byte shared conversion +workspace per rank. Integer word permutation reconstructs the original +TurboMind prefill layout without changing weights or arithmetic. Live +same-call shadow covers all 256 target projections on each rank: 495818 +calls and 35330540544 output elements total, with zero bit differences or +nonfinite values. Original outputs drive generation. The observed KV +budgets are 11804131533/11808325837 bytes, with the same maximum context +and memory utilization. This is quality and memory evidence, not speed. +The first paired startup stops before generation: the V2 full-graph manager +calls its forward function with runtime mode NONE while capturing, so a +FULL-runtime-mode guard misses the route. The revised private harness uses +the existing SM70 decode-graph capture context and retains unique coverage +and split-K checks for all 256 projections. The corrected startup passes all +five paired requests and warmup per fixture: release1k improves from +35.275544 to 33.318339 ms and MBPP28 from 32.274264 to 30.331608 ms. +Candidate warm TTFT is 591.741/166.197 ms. Both arms materialize original +prefill weights into the fixed workspace, so this new prefill cost must +remain visible in TTFT rather than being attributed to decode. Three +independent startups now pass all fifteen measured pairs and warmup per +fixture. The median of startup medians is 35.281057 to 33.318339 ms for +release1k and 32.180410 to 30.175731 ms for MBPP28. Candidate round +p50/p90/p99 are 33.267/33.602/35.439 and 30.195/30.677/31.347 ms; +warm TTFT is 592.478/165.612 ms and pure decode is 90.140/150.135 tokens/s. +Per-start accepted and emitted counts remain separate in +`single-layout-3-start-pair-summary.json`. A trace of the new projection +combination is pending. This does not establish the 25 ms target. + +The first two conversion sanitizer jobs incorrectly retain a GDN-only +kernel filter. Their zero-error summaries do not establish conversion-kernel +memory or race coverage. Failed admission records and logs remain intact; +the replacement gates explicitly select the conversion/materialization +kernels and retain CUDA API error checking. Both corrected memcheck and +racecheck gates pass twelve real rank/projection cases, conversion in both +directions, original TurboMind M8/M129 consumers, changing graph inputs, +changing layout flags and scratch canaries. The retained native library +SHAs are `7e24b7f014df0060af6ba2e7eb8df88d1839d0cd483e96f1a0990a4954b09657` +for static conversion and +`41c8ad3998f7c826355b5333c17568fa4a617b899b56185debac5c58b7b6f640` +for conversion with dynamic prefill materialization. An E2M1 register-permutation decoder is also exact on the sixteen-matrix screen but slower: 0.763648 versus 0.712960 ms. A separate scale-folding probe @@ -569,6 +622,36 @@ output/partial/statistic comparisons, but increases the page1648 working set from 3.858176 to 4.857088 ms and page3296 from 3.941120 to 4.908800 ms. It is rejected before model use. +An E4M3 decoder probe constructs an exact FP16 bit pattern, converts it to +FP32 and multiplies by 256. It retains signed zeros, reserved-NaN handling, +all original QK/PV arithmetic and FP32 partial/reduction storage. All 256 +encodings and forty-four output/partial/statistic cases through 262144 +tokens match, with the same FP64-reference errors. The sixteen-layer +working-set median is 3.429376 versus u8's 3.859712 ms on page1648, and +3.465472 versus 3.942912 ms on page3296. The winning-library-only memcheck +and racecheck pass; the earlier two-library memcheck reports the previously +observed `cuKernelGetFunction` invalid-handle error and remains excluded. +Twenty-four alternating graph replacements across six changing sequence +lengths retain output and valid partial/statistic bits. Live model shadow +now covers all sixteen logical attention layers on both ranks: 27040 calls +and 664535040 output elements, with zero bit differences or nonfinite values. +The first coverage check incorrectly equates unique KV base pointers with +logical layers and fails after completing generation. Sixteen layers share +eight KV memory pools, with distinct page-table pointers for the paired +layers. The revised diagnostic attributes each call to the backend layer +name and requires positive coverage for every layer. Original u8 outputs +drive generation. Complete-round performance is still pending. This probe changes the +decoder instructions, not the stored E4M3 KV precision. Its library SHA is +`6dc516f8d629f15b578b0c287257fefa70cf4dc58a50878988642b38129b0cd7`. + +Another probe shares KV load/decode between two adjacent query heads in one +CTA while retaining each original scalar head's arithmetic. All forty-four +output/partial/statistic comparisons and FP64-reference errors match, but +the working set regresses from 3.430144 to 4.232704 ms on page1648 and +3.469312 to 4.279040 ms on page3296. It is rejected before sanitizer or +model follow-up. These timings establish a regression, not a measured +hardware-counter explanation of its cause. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense From 7eb145f8c831046bad2b93248f1bb0525dc7f919 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 15:25:28 +0800 Subject: [PATCH 10/16] [Kernel] Add a reproducible TP2 matched QPN2 build Generate the CUDA source used by the admitted TP2 projection experiments without installing a serving route. Record the full-projection trace, three-start decoder results, and rejected activation-layout screen. Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../kernels/build_sm70_tp2_matched_qpn2.py | 155 ++++++++++++++++++ docs/design/sm70_dflash2_tp2_verifier.md | 64 +++++++- 2 files changed, 215 insertions(+), 4 deletions(-) create mode 100644 benchmarks/kernels/build_sm70_tp2_matched_qpn2.py diff --git a/benchmarks/kernels/build_sm70_tp2_matched_qpn2.py b/benchmarks/kernels/build_sm70_tp2_matched_qpn2.py new file mode 100644 index 0000000000..4bf83597c3 --- /dev/null +++ b/benchmarks/kernels/build_sm70_tp2_matched_qpn2.py @@ -0,0 +1,155 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Build an isolated TP2 QPN2 candidate with TurboMind's rounding/split order. + +This reproduces the retained TP2 projection experiment. It does not install a +library or enable a serving route. The caller must use the split count observed +for the matching TurboMind projection; this is not an autotuning interface. +""" + +import argparse +import hashlib +import json +import shutil +from pathlib import Path + + +def replace_exact(source: str, old: str, new: str, count: int) -> str: + if source.count(old) != count: + raise ValueError(f"QPN2 source anchor changed: {old!r}") + return source.replace(old, new) + + +def generate(source: str) -> str: + source = source.replace("nvfp4_qpn2_", "tp2_qpn2_matched_").replace( + "TORCH_LIBRARY_FRAGMENT(_qpn2_candidate,", + "TORCH_LIBRARY_FRAGMENT(_tp2_qpn2_matched,", + ) + source = replace_exact( + source, + " const half2 global_scale2 = __float2half2_rn(global_scale * 16384.0f);\n", + "", + 2, + ) + source = replace_exact( + source, + """ const half2 scale = __hmul2( + fp8e4m3_to_half2(__ldg(scale_ptr + static_cast(group) * 32)), + global_scale2);""", + """ const half2 raw_scale = fp8e4m3_to_half2( + __ldg(scale_ptr + static_cast(group) * 32)); + // Match TurboMind's effective FP16 scale before E2M1 multiplication. + // Do not round the global factor to FP16 before multiplying the group. + const half effective = __float2half_rn(__low2float(raw_scale) * global_scale); + const half2 scale = __hmul2(__halves2half2(effective, effective), + __float2half2_rn(16384.0f));""", + 2, + ) + source = replace_exact( + source, + " const int groups_per_warp = groups_k16 / SplitK;\n" + " const int group_begin = warp * groups_per_warp;", + """ const int chunks = k / 64; + const int chunks_per_warp = chunks / SplitK; + const int extra_begin = SplitK - chunks % SplitK; + const int group_begin = (warp * chunks_per_warp + max(warp - extra_begin, 0)) * 4; + const int groups_per_warp = (chunks_per_warp + (warp >= extra_begin)) * 4;""", + 2, + ) + source = replace_exact( + source, + "split_k == 8 || split_k == 16 || split_k == 32", + "(split_k >= 1 && split_k <= 16) || split_k == 32", + 1, + ) + source = replace_exact( + source, + "(input.size(1) / 16) % split_k == 0", + "input.size(1) / 64 >= split_k", + 2, + ) + pieces = [] + for split in range(1, 17): + if split in (5, 7, 8, 9, 16): + continue + condition = "if" if not pieces else "else if" + pieces.append( + f" {condition} (split_k == {split}) {{\n" + f" VLLM_LAUNCH_QPN2(1, {split}, 1);\n }}" + ) + dispatch = ( + "\n".join(pieces) + + """ else if (split_k == 5) { + VLLM_LAUNCH_QPN2(1, 5, 1); + } else if (split_k == 7) { + VLLM_LAUNCH_QPN2(1, 7, 1); + } else if (split_k == 9) { + VLLM_LAUNCH_QPN2(1, 9, 1); + } else if (native_two_tile && split_k == 8 && accumulator_chains == 1) { + VLLM_LAUNCH_QPN2(2, 8, 1);""" + ) + source = replace_exact( + source, + " if (native_two_tile && split_k == 8 && accumulator_chains == 1) {\n" + " VLLM_LAUNCH_QPN2(2, 8, 1);", + dispatch, + 1, + ) + return source.replace("qpn2_matched", "qpn2_matched_all") + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--output-dir", type=Path, required=True) + parser.add_argument("--build", action="store_true") + args = parser.parse_args() + root = Path(__file__).resolve().parents[2] + parent = root / "csrc/sm70_turbomind/ops/nvfp4_qpn2_sm70.cu" + output = args.output_dir.resolve() + sources = output / "sources" + sources.mkdir(parents=True, exist_ok=True) + path = sources / "qpn2-matched-all.cu" + path.write_text(generate(parent.read_text())) + shutil.copy2(parent.parent / "LICENSE.v100-skinny", sources) + flags = [ + "-O3", + "-std=c++17", + "-DVLLM_NVFP4_QPN2_STANDALONE", + "-DVLLM_NVFP4_QPN2_BENCHMARK_CANDIDATE", + "-U__CUDA_NO_HALF_OPERATORS__", + "-U__CUDA_NO_HALF_CONVERSIONS__", + "-U__CUDA_NO_HALF2_OPERATORS__", + ] + report = { + "parent_source_sha256": hashlib.sha256(parent.read_bytes()).hexdigest(), + "source_sha256": hashlib.sha256(path.read_bytes()).hexdigest(), + "cuda_flags": flags, + "supported_experiment": "TP2 q8, one accumulator, observed TM split", + "serving_route_enabled": False, + } + if args.build: + from torch.utils.cpp_extension import load + + build = output / "build" + build.mkdir(exist_ok=True) + library = Path( + load( + name="tp2_qpn2_matched_all", + sources=[str(path)], + build_directory=str(build), + extra_cflags=["-O3"], + extra_cuda_cflags=flags, + is_python_module=False, + verbose=True, + ) + ) + report.update( + library=str(library), + library_sha256=hashlib.sha256(library.read_bytes()).hexdigest(), + ) + (output / "manifest.json").write_text(json.dumps(report, indent=2) + "\n") + print(json.dumps(report, indent=2), flush=True) + + +if __name__ == "__main__": + main() diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 84681450a5..e6cda2e21c 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -574,8 +574,42 @@ release1k and 32.180410 to 30.175731 ms for MBPP28. Candidate round p50/p90/p99 are 33.267/33.602/35.439 and 30.195/30.677/31.347 ms; warm TTFT is 592.478/165.612 ms and pure decode is 90.140/150.135 tokens/s. Per-start accepted and emitted counts remain separate in -`single-layout-3-start-pair-summary.json`. A trace of the new projection -combination is pending. This does not establish the 25 ms target. +`single-layout-3-start-pair-summary.json`. This does not establish the 25 ms +target. + +The full-target QPN2 trace now covers ten interior rounds on both ranks, +with 256 projection calls per rank per round. Profiled critical-rank wall +is 35.615096 ms, GPU union 33.243090 ms and uncovered time 2.372007 ms. +Target service is 23.626618 ms, draft 6.628412 ms, target head/sampling +2.268724 ms and context/output 0.609346 ms. Projection shapes account for: + +| Projection | N / K / original split | Calls per round | GPU service, ms | +| --- | --- | ---: | ---: | +| MLP gate/up | 17408 / 5120 / 7 | 64 | 6.091905 | +| MLP down | 5120 / 8704 / 8 | 64 | 2.790530 | +| GDN input | 8256 / 5120 / 9 | 48 | 2.008633 | +| Attention/GDN output | 5120 / 3072 / 9 | 64 | 1.278775 | +| Attention QKV | 7168 / 5120 / 5 | 16 | 0.655133 | + +MLP gate/up and down account for approximately 69% of QPN2 service. Draft +dense projections, including its head, cost 5.398522 ms. These are priorities +for further work, not estimates of additive end-to-end savings. Hardware +counters remain unavailable. The first profile attempt fails before +generation because the launcher overrides the requested worker extension. +The corrected capture and request complete with 280 output tokens, 96 rounds +and 184 accepted drafts, matching warmup. Its bounded process cleanup exits +137; the trace is not described as a clean exit-zero benchmark. The retained +report, exported SQLite and interval/route checks admit only diagnostic use. +Evidence is `tp2-single-layout-trace-admission.json`, +`tp2-single-layout-trace.json` and `tp2-single-layout-projection-shapes.json`. + +The reproducible private-kernel generator is now checked in as +`benchmarks/kernels/build_sm70_tp2_matched_qpn2.py`. It preserves the tested +CUDA source byte for byte (SHA256 +`139ff11214d1fb49062efe1e6d9dc824588e5f2f14439435154d30b43915fc62`), +including the original K64 partition boundaries, effective FP16 scale +rounding, single accumulator chain and ordered partial sum. It generates +an isolated library; it does not install a library or enable a serving route. The first two conversion sanitizer jobs incorrectly retain a GDN-only kernel filter. Their zero-error summaries do not establish conversion-kernel @@ -640,8 +674,8 @@ logical layers and fails after completing generation. Sixteen layers share eight KV memory pools, with distinct page-table pointers for the paired layers. The revised diagnostic attributes each call to the backend layer name and requires positive coverage for every layer. Original u8 outputs -drive generation. Complete-round performance is still pending. This probe changes the -decoder instructions, not the stored E4M3 KV precision. Its library SHA is +drive generation. This probe changes the decoder instructions, not the stored +E4M3 KV precision. Its library SHA is `6dc516f8d629f15b578b0c287257fefa70cf4dc58a50878988642b38129b0cd7`. Another probe shares KV load/decode between two adjacent query heads in one @@ -652,6 +686,28 @@ the working set regresses from 3.430144 to 4.232704 ms on page1648 and model follow-up. These timings establish a regression, not a measured hardware-counter explanation of its cause. +The exact FP16-bridge decoder now passes three independent startups with +five alternating pairs after warmup per fixture. Both arms retain balanced +192-projection QPN2, BV16 GDN and original u8 prefill. The median of startup +complete-round medians is 33.895185 to 33.373525 ms on release1k and +30.741620 to 30.641039 ms on MBPP28. All fifteen pairs and warmups match token +IDs, acceptance counters and natural EOS. Candidate round p50/p90/p99 are +33.289/33.683/35.361 and 30.564/31.070/32.148 ms; TTFT is 576.749/148.847 ms +and pure decode 103.063/146.342 tokens/s. Per-start trajectories and accepted +versus emitted counts are retained in `half-bridge-3-start-pair-summary.json`. +This is not a paired comparison with the single-layout campaign. + +A K16-major activation-layout candidate also preserves all sixteen real +matrix outputs, FP64-reference errors and three changing graph replays. +The continuous four-layer working set regresses from matched QPN2's +0.712448 to 0.722688 ms, even before input packing is charged. The tested +candidate is rejected before producer or model integration. Its initial +build omits operator registration and fails before GPU comparison; the +corrected build and failure are retained separately. A separate experiment +changes only the executable grid of the unchanged scalar attention kernel. +It preserves outputs and valid partial/statistic bits but saves only about +0.06–0.07 ms across sixteen KV layers. It is not advanced to a model route. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense From ca0ea462c1877525fb231faf4f817d7929a3a64a Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 16:25:57 +0800 Subject: [PATCH 11/16] [Kernel] Decode TP2 E4M3 through exact half expansion Preserve the scalar q8 reduction and rounding contract while using an exact half bit expansion and FP32 scale. Require native revision 2 when explicitly enabled and reject stale revision 1 modules. Validate with 15 native tests, 24 graph replacements, full-length memcheck, and bounded racecheck. Record combined precision diagnostics, first paired round cost, and rejected QPN2 candidates. Long racecheck timeout remains excluded. Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 97 ++++++++++++++++++- .../flash_attn_v100/flash_attn_interface.py | 4 +- .../kernel/flash_decode_paged.cu | 2 +- flash-attention-v100/kernel/fp8_kv_utils.cuh | 18 ++-- .../test_sm70_tp2_e4m3_scalar_fast.py | 14 ++- 5 files changed, 112 insertions(+), 23 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index e6cda2e21c..34d8a422a5 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -59,9 +59,10 @@ only for q shape `[8,12,256]`, E4M3 KV with two heads, FP32 partial storage, off by default. Unverified shapes use the original route. A requested matching route rejects a stale native library instead of silently reporting success. -The specialization constructs normal E4M3 values directly in FP32 bit fields, -retains the original signed zeros, subnormals and NaN payload, and unrolls the -PV loop by eight. Each output still follows the original ascending-token FMA +Revision 1 constructs normal E4M3 values directly in FP32 bit fields. Revision +2 uses an exact FP16 bit expansion followed by FP32 multiplication by 256. +Both retain the original signed zeros, subnormals and NaN payload and unroll +the PV loop by eight. Each output follows the original ascending-token FMA chain. It retains partition boundaries, score reductions, FP32 intermediate storage, output rounding, KV scales and the original final reduction kernel. The native launch counter proves host dispatch, including capture-time calls; @@ -708,6 +709,83 @@ changes only the executable grid of the unchanged scalar attention kernel. It preserves outputs and valid partial/statistic bits but saves only about 0.06–0.07 ms across sixteen KV layers. It is not advanced to a model route. +The context/probe scheduling experiment retains the original NumPy top-p +cutoff guard. It copies the FP32 probe to a fixed pinned buffer, records an +event, then launches the original context graph on its original stream. +The CPU guard waits for the probe copy while context work continues. Missing +or unsupported probes flush pending context before state updates. Live +shadow compares hidden states, positions and every projected context K/V +before versus after the target head: 966 context and probe checks per rank +pass. Three independent startups retain all fifteen measured pairs and +warmups per fixture. With balanced 192-projection QPN2 and BV16 fixed, +complete rounds improve from 33.778139 to 33.251763 ms on release1k and +30.696557 to 30.018241 ms on MBPP28. Candidate p50/p90/p99 are +33.176/33.809/35.623 and 30.118/30.981/32.607 ms; TTFT is 579.024/148.915 ms +and pure decode 92.123/149.354 tokens/s. Evidence is +`context-probe-3-start-pair-summary.json`. Composition with the full 256 +projection layout, BV2 and the newer decoder requires a separate diagnostic +and unprofiled campaign; these independent gains are not added together. + +The combined diagnostic fixes all 256 QPN2 projections and source-integrated +BV2, then switches the exact decoder and context schedule together. Its first +check incorrectly compares the entire sampled-token allocation. All target +hidden states, full local-vocabulary logits, accepted lengths and final +outputs match, but unused sampled-token tails differ. The original sparse +sampler uses `new_empty` and writes only its valid prefix; downstream output +uses `num_sampled`. The failed report is retained without admission. + +The corrected diagnostic records each raw differing column and confirms it +is outside the valid prefix. It also fills invalid tails with distinct values +in the two arms before downstream consumers run. All later hidden states, +full local-vocabulary logits, valid sampled tokens, acceptance and natural EOS +still match. The comparison covers 1162 complete-vocabulary rows, represented +by 288547840 FP32 elements across the two rank shards. Each arm poisons +476/139 tail elements per rank on release1k/MBPP28; context comparisons cover +95/50 rounds per rank. This tests tail isolation on these requests and does +not resolve the earlier cross-startup variation. Evidence is +`tp2-combined-shadow-2-admission.json` and `combined-shadow-2-summary.json`. + +The first unprofiled combined startup passes all five pairs and warmup per +fixture. With full QPN2/BV2 fixed, adding the private exact decoder and +context overlap changes complete rounds from 33.184675 to 31.896585 ms on +release1k and 30.082788 to 29.146665 ms on MBPP28. Candidate round p50/p90/p99 +are 31.863/32.362/34.516 and 29.078/29.677/30.255 ms. Warm TTFT is +593.983/163.962 ms, and pure decode 108.022/153.820 tokens/s. This is one +startup, not the three-startup gate or the approximately 25 ms goal. It is +also not a performance claim for the newly rebuilt native revision 2. + +A separate MLP two-accumulator-chain screen preserves weights, scale rounding +and logical K64 splits but changes accumulation order. Both TP2 ranks and +three fixed input amplitudes produce twelve checks of the actual gate/up +and down shapes. Eight checks expand independent FP64-reference error, with +122–604 changed FP16 output elements per case. The arithmetic candidate is +rejected before timing or model integration. See +`qpn2-mlp-two-chain-decision.json`. + +Native revision 2 now contains the exact FP16-bridge decoder behind the same +default-off TP2 flag. An explicitly requested matching route rejects revision +1 and older libraries. The isolated build SHA256 is +`9d0fe7186bfe82ccd0b58f0795b9f0b4a70eb7ecf8dc9caf345efb4055752cdf`. +It passes 15 native tests, including both stale-library cases, exhaustive +byte decoding, FP64 reference, changing graphs and unsupported shapes. +Twenty-four replacements of the real kernel function preserve outputs and +valid partial/statistic bits. Memcheck covers 1025, 3297 and 262144 tokens +with zero errors. The 262144-token racecheck reaches the 240-second limit +and is retained as a timeout, not a pass. A bounded racecheck at 1025/3297 +tokens completes with zero hazards. Each successful native invocation +records positive fast-path host dispatch counts. The final source includes +a whitespace-only changed-line formatting pass after the build snapshot. +Whole-model admission of this rebuilt library remains separate from the +private decoder's earlier paired performance evidence. + +A q8-only QPN2 specialization removes unused row predicates and row offsets +while preserving all dot-product arithmetic. All sixteen real projection +outputs, FP64-reference errors, changing replays and output canaries match; +unsupported row counts are rejected. Compiler register use drops from 52 to +48 per thread, but the working-set median worsens from 0.725504 to 0.776960 ms. +It is rejected before model work. Register count alone is not a performance +result; evidence is `qpn2-static-m8-decision.json`. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense @@ -733,6 +811,19 @@ See `head-cublaslt-probe.json`, `head-lt-plan-probe.json` and the ## Reproduction and retained negative results +Generate the isolated TP2 projection candidate without installing it: + +```bash +CUDA_VISIBLE_DEVICES="" TORCH_CUDA_ARCH_LIST=7.0 MAX_JOBS=2 \ + .venv/bin/python benchmarks/kernels/build_sm70_tp2_matched_qpn2.py \ + --output-dir /tmp/tp2-matched-qpn2 --build +``` + +The manifest records source and library hashes. The tested candidate uses +the observed TurboMind split for each real TP2 q8 projection and one +accumulator chain. Other shapes, split choices or new builds require their +own numerical and model admission. + Build Flash-V100 from this branch with the same CUDA/Torch/compiler flags and select that module before running the tests. Set `CUDA_VISIBLE_DEVICES` only to an owned rear GPU, and use private build/compiler caches. diff --git a/flash-attention-v100/flash_attn_v100/flash_attn_interface.py b/flash-attention-v100/flash_attn_v100/flash_attn_interface.py index 6f362140d1..183c853fc0 100644 --- a/flash-attention-v100/flash_attn_v100/flash_attn_interface.py +++ b/flash-attention-v100/flash_attn_v100/flash_attn_interface.py @@ -1025,9 +1025,9 @@ def flash_attn_decode_paged( version = getattr( flash_attn_v100_cuda, "tp2_e4m3_scalar_fast_version", None ) - if not callable(version) or int(version()) < 1: + if not callable(version) or int(version()) < 2: raise RuntimeError( - "Rebuild Flash-V100 for TP2 E4M3 scalar fast revision 1" + "Rebuild Flash-V100 for TP2 E4M3 scalar fast revision 2" ) tmp_out, max_logits, exp_sums, active_num_partitions = ( _get_decode_workspace_for_plan( diff --git a/flash-attention-v100/kernel/flash_decode_paged.cu b/flash-attention-v100/kernel/flash_decode_paged.cu index 30a3a00aaf..a6137cbb32 100644 --- a/flash-attention-v100/kernel/flash_decode_paged.cu +++ b/flash-attention-v100/kernel/flash_decode_paged.cu @@ -4415,7 +4415,7 @@ int64_t flash_attention_grouped_e4m3_fp32_precision_version() { return 4; } -int64_t flash_attention_tp2_e4m3_scalar_fast_version() { return 1; } +int64_t flash_attention_tp2_e4m3_scalar_fast_version() { return 2; } int64_t flash_attention_tp2_e4m3_scalar_fast_launch_count() { // Includes capture-time launches; CUDA Graph replay does not call this host diff --git a/flash-attention-v100/kernel/fp8_kv_utils.cuh b/flash-attention-v100/kernel/fp8_kv_utils.cuh index 778a30fd76..06eadb4bd0 100644 --- a/flash-attention-v100/kernel/fp8_kv_utils.cuh +++ b/flash-attention-v100/kernel/fp8_kv_utils.cuh @@ -48,19 +48,13 @@ __device__ __forceinline__ __half fp8_e5m2_to_half(uint8_t raw) { return __ushort_as_half(static_cast(raw) << 8); } -// E4M3 normals map exactly into the IEEE float exponent and mantissa fields. -// Keep the original NaN payload and signed zero, including E4M3 subnormals. +// Finite E4M3 values map exactly to FP16 bits followed by FP32 scaling. +// Preserve the original NaN payload and signed zeros, including subnormals. __device__ __forceinline__ float fp8_e4m3fn_to_float_bits(uint8_t raw) { - const uint32_t magnitude = raw & 0x7fu; - const uint32_t sign = static_cast(raw & 0x80u) << 24; - uint32_t bits = (magnitude << 20) + 0x3c000000u; - if (magnitude < 8) { - bits = __float_as_uint(static_cast(magnitude) * 0.001953125f); - } - if (magnitude == 0x7f) { - return quiet_nan_f(); - } - return __uint_as_float(bits | sign); + const uint16_t half_bits = ((static_cast(raw) << 7) & 0x3f80u) | + ((static_cast(raw) << 8) & 0x8000u); + const float value = __half2float(__ushort_as_half(half_bits)) * 256.0f; + return (raw & 0x7fu) == 0x7fu ? quiet_nan_f() : value; } __device__ __forceinline__ __half2 fp8_e5m2_pair_to_half2(uint16_t raw_pair) { diff --git a/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py b/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py index 25b3ae3f9c..7ba50ad463 100644 --- a/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py +++ b/tests/kernels/attention/test_sm70_tp2_e4m3_scalar_fast.py @@ -11,13 +11,17 @@ FLAG = "VLLM_FLASH_V100_TP2_E4M3_SCALAR_FAST" -def test_requested_fast_path_rejects_stale_library(monkeypatch): +@pytest.mark.parametrize("native_version", [None, 1]) +def test_requested_fast_path_rejects_stale_library(monkeypatch, native_version): interface = pytest.importorskip("flash_attn_v100.flash_attn_interface") monkeypatch.setenv(FLAG, "1") + stale = SimpleNamespace(grouped_e4m3_fp32_precision_version=lambda: 4) + if native_version is not None: + stale.tp2_e4m3_scalar_fast_version = lambda: native_version monkeypatch.setattr( interface, "flash_attn_v100_cuda", - SimpleNamespace(grouped_e4m3_fp32_precision_version=lambda: 4), + stale, ) monkeypatch.setattr( interface, "flash_attn_grouped_e4m3_fp32_available", lambda: True @@ -34,7 +38,7 @@ def test_requested_fast_path_rejects_stale_library(monkeypatch): kv = torch.empty((1, 3296, 2, 256), dtype=torch.uint8) table = torch.zeros((8, 1), dtype=torch.int32) seq = torch.zeros(8, dtype=torch.int32) - with pytest.raises(RuntimeError, match="TP2 E4M3 scalar fast revision 1"): + with pytest.raises(RuntimeError, match="TP2 E4M3 scalar fast revision 2"): interface.flash_attn_decode_paged( q, kv, kv, table, seq, kv_cache_dtype="fp8_e4m3" ) @@ -47,8 +51,8 @@ def native(): interface = pytest.importorskip("flash_attn_v100.flash_attn_interface") extension = interface.flash_attn_v100_cuda version = getattr(extension, "tp2_e4m3_scalar_fast_version", lambda: 0) - if version() < 1: - pytest.skip("rebuild Flash-V100 with TP2 scalar fast revision 1") + if version() < 2: + pytest.skip("rebuild Flash-V100 with TP2 scalar fast revision 2") return extension From bb333ee528f0d9e4bbe64d65b6a708a0617ea427 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 16:57:39 +0800 Subject: [PATCH 12/16] [Doc] Record the three-start TP2 native combination result Report complete-round latency, TTFT, decode rate and separate accepted/emitted lengths for the native exact decoder plus context overlap, with full QPN2 and BV2 fixed. Retain the unmet 25 ms target and unresolved cross-startup limitation. Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 29 ++++++++++++++++++++++-- 1 file changed, 27 insertions(+), 2 deletions(-) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 34d8a422a5..287c6597a1 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -775,8 +775,33 @@ and is retained as a timeout, not a pass. A bounded racecheck at 1025/3297 tokens completes with zero hazards. Each successful native invocation records positive fast-path host dispatch counts. The final source includes a whitespace-only changed-line formatting pass after the build snapshot. -Whole-model admission of this rebuilt library remains separate from the -private decoder's earlier paired performance evidence. +Whole-model shadow of this rebuilt library now passes on both ranks, including +full local-vocabulary logits, valid acceptance records and distinct invalid-tail +sentinels. It is separate from the private decoder's earlier performance +evidence. + +Three unprofiled native-combination startups now pass all fifteen measured +pairs and warmups per fixture. Both arms keep the single-layout 256-projection +QPN2 path and source-integrated BV2; the candidate adds the rebuilt exact +decoder and context overlap. The median of startup request medians is: + +| Metric | release1k control / candidate | MBPP28 control / candidate | +| --- | ---: | ---: | +| Complete round, ms | 32.934885 / **31.884546** | 29.830193 / **29.279787** | +| Round p50, ms | 32.913 / 31.876 | 29.912 / 29.229 | +| Round p90, ms | 33.361 / 32.370 | 30.454 / 29.913 | +| Round p99, ms | 35.142 / 34.244 | 32.367 / 31.149 | +| Warm TTFT, ms | 591.649 / 591.596 | 165.546 / 164.810 | +| Pure decode, tokens/s | 91.089 / 94.089 | 152.659 / 155.623 | +| Accepted drafts per round, both arms | 2.010638 | 3.569231 | +| Emitted tokens per round, both arms | 3.010638 | 4.569231 | + +All three startups produce 283/297 output tokens and 94/65 draft rounds on +release1k/MBPP28. This does not retroactively resolve the older startup +variation. Source is `ca0ea462c1877525fb231faf4f817d7929a3a64a`; runtime library +and private harness hashes are frozen in `combined-native-three-start-manifest.json`. +Raw evidence is `combined-native-3-start-pair-summary.json`. The approximately +25 ms target remains unmet, so defaults remain off and the PR remains Draft. A q8-only QPN2 specialization removes unused row predicates and row offsets while preserving all dot-product arithmetic. All sixteen real projection From 5ae004b8ce2e7b8eb5d4de0b251f887b5965ae63 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 17:58:51 +0800 Subject: [PATCH 13/16] [Doc] Hold draft arithmetic after same-prefix distribution audit Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 67 ++++++++++++++++++++++++ 1 file changed, 67 insertions(+) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 287c6597a1..668b28ef6b 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -811,6 +811,73 @@ unsupported row counts are rejected. Compiler register use drops from 52 to It is rejected before model work. Register count alone is not a performance result; evidence is `qpn2-static-m8-decision.json`. +### Strict draft GEMM: better local reference error does not preserve proposals + +Actual TP2 draft operands now cover both ranks, twenty projections and five +query steps per rank. The original row-weight GEMM reproduces all 200 retained +outputs. Column-weight GEMM with reduced-precision reduction disabled changes +all 200 outputs, but expands none of the independent FP64 maximum, p99 or +relative-L2 errors. Its twenty-projection working set falls from 2.910515 to +2.362880 ms. This local result does not admit a model optimization. + +A diagnostic captures both projections in the same q8 query graph, selects +the propagated arm with a device flag, and replays control/candidate/control +at each real prefix. The last control replay supplies serving outputs and +query KV. Across 24 prefixes per fixture on both ranks, every retained +projection input/output, FP32 head and selector buffer repeats bytewise in +the two control replays. Query tokens, positions, slot mappings and RNG states +are unchanged. Natural requests before, during and after the audit retain +their token IDs, acceptance counters and EOS. This isolates the candidate +from diagnostic perturbation within this startup; it does not resolve the +earlier cross-startup variation. + +The candidate fails distribution admission despite no top-1 or sampled draft +token changes in the 336 observed rows: + +| Observation | Result | +| --- | ---: | +| Maximum full draft-vocabulary TV | 5.021620% | +| Changed top-20 sets | 26 / 336 | +| Changed diagnostic k20/p0.95 support sets | 18 / 336 | +| Maximum actual selector-proposal TV | 47.719886% | +| Changed actual proposal support sets | 26 / 336 | + +The actual draft selector uses sixteen candidates and proposal top-p 1.0; +the k20/p0.95 row is a separate diagnostic. At release1k prefix step 19, +proposal row 6, token 40718 enters the selector support and receives +47.719886% probability after the selector's edge scores. The first differing +operator is `model.layers.64.self_attn.qkv_proj`, with identical input on both +ranks. Subsequent layer inputs change as the difference propagates. A smaller +local FP64 error and an unchanged sampled token do not establish unchanged +sampling or acceptance. This candidate is rejected before unprofiled model +timing and remains disabled. + +Evidence: `tp2-draft-f16-layout-screen.json`, +`tp2-draft-column-shadow-1-diagnostic.json`, and +`draft-column-tp2-decision.json`. The frozen diagnostic harness and five CPU +checks include candidate-ID permutation invariance and a known TV of 0.5. + +### Trace of the current native combination + +A new release1k trace uses source `bb333ee528f0d9e4bbe64d65b6a708a0617ea427` +and the frozen native revision 2 combination, with the draft arithmetic +candidate disabled. Ten interior rounds cover all 256 QPN2 projections, +48 BV2 GDN kernels and sixteen native attention partitions per rank/round. +Warmup and captured natural requests match at 280 output tokens. Nsight +exits zero; cleanup signals and source/library provenance remain recorded. +The different startup trajectory does not replace the paired unprofiled +31.884546/29.279787-ms result above. + +The profiled critical-rank round is 33.955203 ms, with 32.153534 ms of GPU +activity and 1.801669 ms not covered by GPU activity. Rank-mean service is +12.820451 ms for QPN2, 3.422064 ms for scalar target attention, 1.055777 ms +for recurrent GDN and 1.930007 ms for communication. Draft service is +6.615042 ms; the two full FP32 head GEMMs together take 4.140522 ms and are +already included in the target sampling/draft phases. Context computation +now falls inside the sampling phase, so phase labels must not be read as +independent speedups. These numbers are attribution, not performance +acceptance. See `tp2-combined-native-trace.json` and its admission report. + ### LM-head width and accumulation order The trace spends approximately 4.133 ms across the target and draft dense From 5aa67252674db770eb8b0594963a517456249135 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:56:46 +0800 Subject: [PATCH 14/16] [Kernel] Gate one-copy TP2 combined GDN projection tails Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 34 ++++ .../test_sm70_tp2_combined_gdn_split.py | 169 ++++++++++++++++++ vllm/envs.py | 6 + vllm/model_executor/models/qwen3_5.py | 34 +++- 4 files changed, 240 insertions(+), 3 deletions(-) create mode 100644 tests/kernels/test_sm70_tp2_combined_gdn_split.py diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 668b28ef6b..e0a84b559b 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -901,6 +901,40 @@ admission remain outstanding. No narrowed LM-head route is enabled. See `head-cublaslt-probe.json`, `head-lt-plan-probe.json` and the [CUDA 12.8 cuBLASLt reference](https://docs.nvidia.com/cuda/archive/12.8.0/cublas/index.html). +### Combined TP2 GDN projection copies + +The native-combination trace still contains three tail gathers per GDN layer. +QUASAR uses the combined projection branch, which did not call the existing +one-copy z/b/a helper. A separate, default-off +`VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT` switch now routes the verified TP2 +geometry through that helper. It also requires the existing SM70/DFlash2 split +gate. Other TP sizes, feature dimensions and dtypes retain the old path. +QKV remains a view for the convolution's in-place update; the helper reads +the actual padded row stride and BA view offset. No arithmetic changes. + +The isolated copy screen passes all 65,536 FP16 bit encodings, changing graph +replays, rows 1/7/8/9/32/128/4096, row strides 8240/8256/8320 and storage +offsets 0/17. Input and padding bits are unchanged. Two consecutive 48-layer +working sets take 0.722739 ms per round of three gathers versus 0.100045 ms +for the one-copy helper. This approximately 0.623-ms local saving is not a +complete model-round result. The actual `forward_cuda` entry and existing +split tests pass all twenty cases, including QKV convolution ownership and +tail bits after changing graph replays. Live compiled-model and unprofiled +performance admission are pending. Evidence: `tp2-combined-gdn-split-screen.json` +and serial queue job 309; no new serving default is enabled. + +Two other bounded screens are closed before model work. The N16 QPN2 tile +matches all sixteen retained real outputs and FP64-reference errors, including +changing graph replays and canaries, but worsens the four-adjacent-layer +working set from 0.708352 to 0.795136 ms. The full FP32 head keeps cuBLASLt +algorithm 21, split two, reduction 4 and stage 14; tile IDs 5 and 11 match +all 24 saved real M7/M8 cases across both ranks. Tile 15 is unsupported +(status 15), rather than a numerical failure. Tile 11 saves only about +0.006/0.045 ms across two heads on rank 0/1. Neither screen justifies a model +performance candidate. See `tp2-qpn2-n16-screen.json` and +`head-lt-tiles-screen.json`; the accepted complete-round baseline remains +31.884546/29.279787 ms. + ## Reproduction and retained negative results Generate the isolated TP2 projection candidate without installing it: diff --git a/tests/kernels/test_sm70_tp2_combined_gdn_split.py b/tests/kernels/test_sm70_tp2_combined_gdn_split.py new file mode 100644 index 0000000000..4a36cf98ce --- /dev/null +++ b/tests/kernels/test_sm70_tp2_combined_gdn_split.py @@ -0,0 +1,169 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +from types import SimpleNamespace + +import pytest +import torch + +import vllm.envs as envs +from vllm.model_executor.models import qwen3_5 as model + +pytestmark = pytest.mark.skip_global_cleanup + + +@pytest.mark.parametrize( + "override,expected", + [ + ({}, True), + ({"tp_size": 4}, False), + ({"key_dim": 1024}, False), + ({"value_dim": 3072}, False), + ({"num_v_heads": 24}, False), + ({"enable_sm70_dflash2_fused_gdn_split": False}, False), + ], +) +def test_combined_split_constructor_keeps_unverified_shapes_off( + monkeypatch, override, expected +): + envs.disable_envs_cache() + monkeypatch.setenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "1") + + def init(self): + torch.nn.Module.__init__(self) + for key, value in ( + dict( + quant_config=None, + tp_size=2, + key_dim=2048, + value_dim=6144, + num_v_heads=48, + enable_sm70_dflash2_fused_gdn_split=True, + ) + | override + ).items(): + setattr(self, key, value) + + monkeypatch.setattr(model.QwenGatedDeltaNetAttention, "__init__", init) + monkeypatch.setattr(model, "_uses_split_gdn_input_projections", lambda _: False) + layer = model.Qwen3_5GatedDeltaNet() + assert layer.enable_sm70_dflash2_tp2_combined_gdn_split is expected + monkeypatch.delenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT") + assert not model.Qwen3_5GatedDeltaNet().enable_sm70_dflash2_tp2_combined_gdn_split + monkeypatch.setenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "1") + monkeypatch.setattr(model, "_uses_split_gdn_input_projections", lambda _: True) + assert not model.Qwen3_5GatedDeltaNet().enable_sm70_dflash2_tp2_combined_gdn_split + + +@pytest.mark.parametrize( + "rows,stride,offset", + [ + (1, 8240, 0), + (7, 8256, 17), + (8, 8256, 0), + (8, 8320, 17), + (9, 8256, 0), + (32, 8256, 0), + (4096, 8256, 0), + ], +) +def test_combined_split_forward_preserves_bits_and_convolution_ownership( + monkeypatch, rows, stride, offset +): + if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): + pytest.skip("Requires an owned SM70 GPU") + envs.disable_envs_cache() + monkeypatch.setenv("VLLM_SM70_GDN_MIXED_QKV_CONTIGUOUS", "0") + monkeypatch.setattr(model, "_sm70_gdn_qpn8_ba_dispatch_eligible", lambda *a: False) + monkeypatch.setattr(model, "_sm70_dump_gdn_projection_tensor", lambda *a: a[-1]) + monkeypatch.setattr( + model, "_resolve_qwen_gdn_kv_cache_args", lambda *a: (None, None) + ) + # Force materialized reference slices, as in the compiled q8 control. + monkeypatch.setattr( + model, + "_sm70_compile_graph_slice_dim", + lambda x, dim, start, size: x.index_select( + dim, torch.arange(start, start + size, device=x.device) + ), + ) + arena = torch.full( + (offset + rows * stride + 64,), 16977, device="cuda", dtype=torch.int16 + ) + projection = torch.as_strided( + arena.view(torch.float16), (rows, 8240), (stride, 1), offset + ) + hidden = torch.empty((rows, 5120), device="cuda", dtype=torch.float16) + observed: dict[str, torch.Tensor] = {} + + def recurrent(self, *, mixed_qkv, b, a, core_attn_out, **kwargs): + observed.update(b=b, a=a, qkv=mixed_qkv.clone()) + # The real convolution writes QKV in place. Tail materialization must + # neither detach this view nor let those writes corrupt z/b/a. + mixed_qkv.zero_() + return core_attn_out + + monkeypatch.setattr(model, "_qwen_gdn_run_recurrent_core", recurrent) + layer = SimpleNamespace( + prefix="model.layers.0.linear_attn", + tp_size=2, + key_dim=2048, + value_dim=6144, + num_v_heads=48, + head_v_dim=128, + use_split_input_projections=False, + enable_sm70_dflash2_tp2_combined_gdn_split=False, + in_proj_qkvz=lambda _: (projection, None), + _output_projection=lambda core, z, output, n: ( + z.flatten(1), + observed["b"], + observed["a"], + observed["qkv"], + ), + ) + helper = model._sm70_materialize_qwen35_gdn_splits + hits = [] + + def tracked(qkvz, ba, *sizes): + hits.append((qkvz.stride(), ba.storage_offset() - qkvz.storage_offset())) + return helper(qkvz, ba, *sizes) + + monkeypatch.setattr(model, "_sm70_materialize_qwen35_gdn_splits", tracked) + graphs, outputs = [], [] + for enabled in (False, True): + layer.enable_sm70_dflash2_tp2_combined_gdn_split = enabled + model.Qwen3_5GatedDeltaNet.forward_cuda(layer, hidden, None) + torch.cuda.synchronize() + hits.clear() + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + output = model.Qwen3_5GatedDeltaNet.forward_cuda(layer, hidden, None) + assert hits == ([((stride, 1), 8192)] if enabled else []) + graphs.append(graph) + outputs.append(output) + + # Four changing patterns cover every FP16 payload at q8, including NaNs + # and signed zero, without numerical comparison that would hide bit flips. + for shift in (0, 16384, 32768, 49152): + bits = (torch.arange(rows * 8240, device="cuda") + shift).to(torch.int16) + before = arena.clone() + reference = bits.view(rows, 8240) + expected = [ + reference[:, lo:hi] + for lo, hi in ((5120, 8192), (8192, 8216), (8216, 8240), (0, 5120)) + ] + for graph, output in zip(graphs, outputs, strict=True): + projection.view(torch.int16).copy_(reference) + graph.replay() + assert all( + torch.equal(a.view(torch.int16), b) + for a, b in zip(output, expected, strict=True) + ) + assert torch.count_nonzero(projection[:, :5120].view(torch.int16)) == 0 + assert torch.equal( + projection[:, 5120:].view(torch.int16), reference[:, 5120:] + ) + # Exclude the projection itself when checking padding and canaries. + before_view = torch.as_strided(before, (rows, 8240), (stride, 1), offset) + before_view.copy_(projection.view(torch.int16)) + assert torch.equal(arena, before) diff --git a/vllm/envs.py b/vllm/envs.py index b3a8248800..fdd0104ddd 100755 --- a/vllm/envs.py +++ b/vllm/envs.py @@ -244,6 +244,7 @@ VLLM_SM70_DFLASH2_TP2_GDN_BV2: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_NORM: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT: bool = False + VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT: bool = False VLLM_SM70_DFLASH2_FUSED_SMALLQ_METADATA: bool = False VLLM_SM70_DFLASH2_GROUPED_SMALLQ_METADATA: bool = False VLLM_SM70_DFLASH2_FUSED_QKV_PACK: bool = False @@ -2240,6 +2241,11 @@ def _resolve_rust_frontend_path() -> str | None: "VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT": lambda: bool( int(os.getenv("VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT", "0")) ), + # Opt in separately: the existing split-projection switch is enabled by + # DFlash2 defaults, but combined QUASAR TP2 still needs model admission. + "VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT": lambda: bool( + int(os.getenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "0")) + ), # Build Flash-V100 small-query verifier rows directly in their persistent # graph buffers. This replaces four repeat_interleave scans per KV group. # The matched TP4 trace is token/acceptance exact and cuts the synchronized diff --git a/vllm/model_executor/models/qwen3_5.py b/vllm/model_executor/models/qwen3_5.py index dbc891fee5..94cf5feb4a 100644 --- a/vllm/model_executor/models/qwen3_5.py +++ b/vllm/model_executor/models/qwen3_5.py @@ -345,6 +345,15 @@ def __init__(self, *args, **kwargs): self.use_split_input_projections = _uses_split_gdn_input_projections( self.quant_config ) + self.enable_sm70_dflash2_tp2_combined_gdn_split = bool( + envs.VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT + and self.enable_sm70_dflash2_fused_gdn_split + and not self.use_split_input_projections + and self.tp_size == 2 + and self.key_dim == 2048 + and self.value_dim == 6144 + and self.num_v_heads == 48 + ) def create_qkvz_proj( self, @@ -477,9 +486,28 @@ def forward_cuda( ba_start = z_start + z_size a_start = ba_start + ba_size mixed_qkv = mixed_qkvzba[..., :qkv_size] - z = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, z_start, z_size) - b = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, ba_start, ba_size) - a = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, a_start, ba_size) + if ( + self.enable_sm70_dflash2_tp2_combined_gdn_split + and mixed_qkvzba.is_cuda + and mixed_qkvzba.dtype == torch.float16 + and mixed_qkvzba.ndim == 2 + and num_tokens > 0 + and mixed_qkvzba.stride(1) == 1 + ): + # QUASAR's logical width is 8240, while QPN2 pads rows to + # 8256. Pass views with their actual stride and BA offset. + # The QKV view remains owned by the projection for convolution. + z, b, a = _sm70_materialize_qwen35_gdn_splits( + mixed_qkvzba, + mixed_qkvzba[..., ba_start:], + qkv_size, + z_size, + ba_size, + ) + else: + z = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, z_start, z_size) + b = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, ba_start, ba_size) + a = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, a_start, ba_size) mixed_qkv = _sm70_dump_gdn_projection_tensor( "split_mixed_qkv", layer_name, mixed_qkv From a20f08357ccd88b12e3ffea178f1d31849019973 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Thu, 10 Sep 2026 00:11:36 +0800 Subject: [PATCH 15/16] [Core] Close TP2 tuning at the accepted 31/29 ms endpoint Withdraw the unadmitted combined-copy experiment; retain its audit history. Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 51 ++++-- .../test_sm70_tp2_combined_gdn_split.py | 169 ------------------ vllm/envs.py | 6 - vllm/model_executor/models/qwen3_5.py | 34 +--- 4 files changed, 39 insertions(+), 221 deletions(-) delete mode 100644 tests/kernels/test_sm70_tp2_combined_gdn_split.py diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index e0a84b559b..13552e854d 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -2,9 +2,21 @@ ## Scope and frozen baseline -The TP2 campaign targets approximately 25 ms per complete B1/q8 DFlash2 round -on two rear V100-SXM2-32GB GPUs. A round includes target, logits/sampling, -state handling, context work and draft. TP4 optimization is a separate campaign. +The user closed the optimization campaign on 2026-09-10 at the accepted +31.884546/29.279787-ms release1k/MBPP28 endpoint. The earlier approximately +25-ms target is no longer a merge requirement. These are complete B1/q8 +DFlash2 rounds on two rear V100-SXM2-32GB GPUs, including target, +logits/sampling, state handling, context work and draft. TP4 is a separate +campaign. No later local microbenchmark replaces these accepted measurements. + +The retained endpoint uses three independent paired startups, five measured +pairs per fixture per startup, unchanged token IDs/acceptance/natural EOS, +and a separate full-logits/hidden-state diagnostic. Source integration keeps +the audited native attention, packed GDN repairs/BV2 and matched QPN2 builder. +The unadmitted combined-projection copy experiment has been withdrawn from +this PR's source. New optimization switches remain opt-in; the full measured +combination also uses the retained QPN2/context worker harness described below. +Merging these source components does not make that entire harness a default. Integration base: `e5d63c51f0fcc1ddf75d229e3df06bf52df206f5`. Use the QUASAR Qwen3.8-27B NVFP4 checkpoint at @@ -163,8 +175,9 @@ Acceptance is reported separately from emitted tokens: | MBPP28 | 5 | 3.569231 | 4.569231 | These paired results admit the attention component for continued experiments. -The approximately 25 ms target, full context sweep and broader quality suite -remain outstanding. The production flag stays off and the PR stays Draft. +At this earlier checkpoint, the approximately 25 ms target, full context +sweep and broader quality suite remained outstanding. The production flag +stayed off and the PR remained Draft. Raw reports are `attention-three-start-pair-summary.json`, `attention-three-start-secondary-metrics.json`, and `tp2-attention-within-start-{3,4,5}-switch.json` in the campaign results. @@ -801,7 +814,8 @@ release1k/MBPP28. This does not retroactively resolve the older startup variation. Source is `ca0ea462c1877525fb231faf4f817d7929a3a64a`; runtime library and private harness hashes are frozen in `combined-native-three-start-manifest.json`. Raw evidence is `combined-native-3-start-pair-summary.json`. The approximately -25 ms target remains unmet, so defaults remain off and the PR remains Draft. +25 ms target was not met and was retired by the user at campaign close. +The 31.884546/29.279787-ms endpoint is the accepted scope; defaults remain off. A q8-only QPN2 specialization removes unused row predicates and row offsets while preserving all dot-product arithmetic. All sixteen real projection @@ -905,12 +919,12 @@ See `head-cublaslt-probe.json`, `head-lt-plan-probe.json` and the The native-combination trace still contains three tail gathers per GDN layer. QUASAR uses the combined projection branch, which did not call the existing -one-copy z/b/a helper. A separate, default-off -`VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT` switch now routes the verified TP2 -geometry through that helper. It also requires the existing SM70/DFlash2 split -gate. Other TP sizes, feature dimensions and dtypes retain the old path. -QKV remains a view for the convolution's in-place update; the helper reads -the actual padded row stride and BA view offset. No arithmetic changes. +one-copy z/b/a helper. Commit `5aa67252674db770eb8b0594963a517456249135` +tested a separately gated TP2 integration. QKV remained a view for the +convolution's in-place update; the helper read the actual padded row stride +and BA offset, with no arithmetic changes. This integration was withdrawn +from the final PR because it had not passed complete-round admission when +the user closed optimization. Its source and evidence remain in history. The isolated copy screen passes all 65,536 FP16 bit encodings, changing graph replays, rows 1/7/8/9/32/128/4096, row strides 8240/8256/8320 and storage @@ -919,9 +933,16 @@ working sets take 0.722739 ms per round of three gathers versus 0.100045 ms for the one-copy helper. This approximately 0.623-ms local saving is not a complete model-round result. The actual `forward_cuda` entry and existing split tests pass all twenty cases, including QKV convolution ownership and -tail bits after changing graph replays. Live compiled-model and unprofiled -performance admission are pending. Evidence: `tp2-combined-gdn-split-screen.json` -and serial queue job 309; no new serving default is enabled. +tail bits after changing graph replays. Seven forward cases also pass +memcheck with zero errors. The live source audit covers 48 layers per rank +and 1,462,855,680 FP16 elements with zero bit differences. Its repeated +same-configuration full-model requests preserve hidden states, complete +FP32 logits and valid acceptance records; this is not an original-versus-new +model-performance result. The separate paired model harness fails before +generation because its ctypes CUDA-graph edge type overwrites the QPN2 +reader's binding. Job 313 is excluded, and no later paired performance is +claimed. Evidence: `tp2-combined-gdn-split-screen.json`, +`tp2-combined-split-shadow-1-admission.json` and queue jobs 309--313. Two other bounded screens are closed before model work. The N16 QPN2 tile matches all sixteen retained real outputs and FP64-reference errors, including diff --git a/tests/kernels/test_sm70_tp2_combined_gdn_split.py b/tests/kernels/test_sm70_tp2_combined_gdn_split.py deleted file mode 100644 index 4a36cf98ce..0000000000 --- a/tests/kernels/test_sm70_tp2_combined_gdn_split.py +++ /dev/null @@ -1,169 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project - -from types import SimpleNamespace - -import pytest -import torch - -import vllm.envs as envs -from vllm.model_executor.models import qwen3_5 as model - -pytestmark = pytest.mark.skip_global_cleanup - - -@pytest.mark.parametrize( - "override,expected", - [ - ({}, True), - ({"tp_size": 4}, False), - ({"key_dim": 1024}, False), - ({"value_dim": 3072}, False), - ({"num_v_heads": 24}, False), - ({"enable_sm70_dflash2_fused_gdn_split": False}, False), - ], -) -def test_combined_split_constructor_keeps_unverified_shapes_off( - monkeypatch, override, expected -): - envs.disable_envs_cache() - monkeypatch.setenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "1") - - def init(self): - torch.nn.Module.__init__(self) - for key, value in ( - dict( - quant_config=None, - tp_size=2, - key_dim=2048, - value_dim=6144, - num_v_heads=48, - enable_sm70_dflash2_fused_gdn_split=True, - ) - | override - ).items(): - setattr(self, key, value) - - monkeypatch.setattr(model.QwenGatedDeltaNetAttention, "__init__", init) - monkeypatch.setattr(model, "_uses_split_gdn_input_projections", lambda _: False) - layer = model.Qwen3_5GatedDeltaNet() - assert layer.enable_sm70_dflash2_tp2_combined_gdn_split is expected - monkeypatch.delenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT") - assert not model.Qwen3_5GatedDeltaNet().enable_sm70_dflash2_tp2_combined_gdn_split - monkeypatch.setenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "1") - monkeypatch.setattr(model, "_uses_split_gdn_input_projections", lambda _: True) - assert not model.Qwen3_5GatedDeltaNet().enable_sm70_dflash2_tp2_combined_gdn_split - - -@pytest.mark.parametrize( - "rows,stride,offset", - [ - (1, 8240, 0), - (7, 8256, 17), - (8, 8256, 0), - (8, 8320, 17), - (9, 8256, 0), - (32, 8256, 0), - (4096, 8256, 0), - ], -) -def test_combined_split_forward_preserves_bits_and_convolution_ownership( - monkeypatch, rows, stride, offset -): - if not torch.cuda.is_available() or torch.cuda.get_device_capability() != (7, 0): - pytest.skip("Requires an owned SM70 GPU") - envs.disable_envs_cache() - monkeypatch.setenv("VLLM_SM70_GDN_MIXED_QKV_CONTIGUOUS", "0") - monkeypatch.setattr(model, "_sm70_gdn_qpn8_ba_dispatch_eligible", lambda *a: False) - monkeypatch.setattr(model, "_sm70_dump_gdn_projection_tensor", lambda *a: a[-1]) - monkeypatch.setattr( - model, "_resolve_qwen_gdn_kv_cache_args", lambda *a: (None, None) - ) - # Force materialized reference slices, as in the compiled q8 control. - monkeypatch.setattr( - model, - "_sm70_compile_graph_slice_dim", - lambda x, dim, start, size: x.index_select( - dim, torch.arange(start, start + size, device=x.device) - ), - ) - arena = torch.full( - (offset + rows * stride + 64,), 16977, device="cuda", dtype=torch.int16 - ) - projection = torch.as_strided( - arena.view(torch.float16), (rows, 8240), (stride, 1), offset - ) - hidden = torch.empty((rows, 5120), device="cuda", dtype=torch.float16) - observed: dict[str, torch.Tensor] = {} - - def recurrent(self, *, mixed_qkv, b, a, core_attn_out, **kwargs): - observed.update(b=b, a=a, qkv=mixed_qkv.clone()) - # The real convolution writes QKV in place. Tail materialization must - # neither detach this view nor let those writes corrupt z/b/a. - mixed_qkv.zero_() - return core_attn_out - - monkeypatch.setattr(model, "_qwen_gdn_run_recurrent_core", recurrent) - layer = SimpleNamespace( - prefix="model.layers.0.linear_attn", - tp_size=2, - key_dim=2048, - value_dim=6144, - num_v_heads=48, - head_v_dim=128, - use_split_input_projections=False, - enable_sm70_dflash2_tp2_combined_gdn_split=False, - in_proj_qkvz=lambda _: (projection, None), - _output_projection=lambda core, z, output, n: ( - z.flatten(1), - observed["b"], - observed["a"], - observed["qkv"], - ), - ) - helper = model._sm70_materialize_qwen35_gdn_splits - hits = [] - - def tracked(qkvz, ba, *sizes): - hits.append((qkvz.stride(), ba.storage_offset() - qkvz.storage_offset())) - return helper(qkvz, ba, *sizes) - - monkeypatch.setattr(model, "_sm70_materialize_qwen35_gdn_splits", tracked) - graphs, outputs = [], [] - for enabled in (False, True): - layer.enable_sm70_dflash2_tp2_combined_gdn_split = enabled - model.Qwen3_5GatedDeltaNet.forward_cuda(layer, hidden, None) - torch.cuda.synchronize() - hits.clear() - graph = torch.cuda.CUDAGraph() - with torch.cuda.graph(graph): - output = model.Qwen3_5GatedDeltaNet.forward_cuda(layer, hidden, None) - assert hits == ([((stride, 1), 8192)] if enabled else []) - graphs.append(graph) - outputs.append(output) - - # Four changing patterns cover every FP16 payload at q8, including NaNs - # and signed zero, without numerical comparison that would hide bit flips. - for shift in (0, 16384, 32768, 49152): - bits = (torch.arange(rows * 8240, device="cuda") + shift).to(torch.int16) - before = arena.clone() - reference = bits.view(rows, 8240) - expected = [ - reference[:, lo:hi] - for lo, hi in ((5120, 8192), (8192, 8216), (8216, 8240), (0, 5120)) - ] - for graph, output in zip(graphs, outputs, strict=True): - projection.view(torch.int16).copy_(reference) - graph.replay() - assert all( - torch.equal(a.view(torch.int16), b) - for a, b in zip(output, expected, strict=True) - ) - assert torch.count_nonzero(projection[:, :5120].view(torch.int16)) == 0 - assert torch.equal( - projection[:, 5120:].view(torch.int16), reference[:, 5120:] - ) - # Exclude the projection itself when checking padding and canaries. - before_view = torch.as_strided(before, (rows, 8240), (stride, 1), offset) - before_view.copy_(projection.view(torch.int16)) - assert torch.equal(arena, before) diff --git a/vllm/envs.py b/vllm/envs.py index fdd0104ddd..b3a8248800 100755 --- a/vllm/envs.py +++ b/vllm/envs.py @@ -244,7 +244,6 @@ VLLM_SM70_DFLASH2_TP2_GDN_BV2: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_NORM: bool = False VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT: bool = False - VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT: bool = False VLLM_SM70_DFLASH2_FUSED_SMALLQ_METADATA: bool = False VLLM_SM70_DFLASH2_GROUPED_SMALLQ_METADATA: bool = False VLLM_SM70_DFLASH2_FUSED_QKV_PACK: bool = False @@ -2241,11 +2240,6 @@ def _resolve_rust_frontend_path() -> str | None: "VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT": lambda: bool( int(os.getenv("VLLM_SM70_DFLASH2_FUSED_GDN_SPLIT", "0")) ), - # Opt in separately: the existing split-projection switch is enabled by - # DFlash2 defaults, but combined QUASAR TP2 still needs model admission. - "VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT": lambda: bool( - int(os.getenv("VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT", "0")) - ), # Build Flash-V100 small-query verifier rows directly in their persistent # graph buffers. This replaces four repeat_interleave scans per KV group. # The matched TP4 trace is token/acceptance exact and cuts the synchronized diff --git a/vllm/model_executor/models/qwen3_5.py b/vllm/model_executor/models/qwen3_5.py index 94cf5feb4a..dbc891fee5 100644 --- a/vllm/model_executor/models/qwen3_5.py +++ b/vllm/model_executor/models/qwen3_5.py @@ -345,15 +345,6 @@ def __init__(self, *args, **kwargs): self.use_split_input_projections = _uses_split_gdn_input_projections( self.quant_config ) - self.enable_sm70_dflash2_tp2_combined_gdn_split = bool( - envs.VLLM_SM70_DFLASH2_TP2_COMBINED_GDN_SPLIT - and self.enable_sm70_dflash2_fused_gdn_split - and not self.use_split_input_projections - and self.tp_size == 2 - and self.key_dim == 2048 - and self.value_dim == 6144 - and self.num_v_heads == 48 - ) def create_qkvz_proj( self, @@ -486,28 +477,9 @@ def forward_cuda( ba_start = z_start + z_size a_start = ba_start + ba_size mixed_qkv = mixed_qkvzba[..., :qkv_size] - if ( - self.enable_sm70_dflash2_tp2_combined_gdn_split - and mixed_qkvzba.is_cuda - and mixed_qkvzba.dtype == torch.float16 - and mixed_qkvzba.ndim == 2 - and num_tokens > 0 - and mixed_qkvzba.stride(1) == 1 - ): - # QUASAR's logical width is 8240, while QPN2 pads rows to - # 8256. Pass views with their actual stride and BA offset. - # The QKV view remains owned by the projection for convolution. - z, b, a = _sm70_materialize_qwen35_gdn_splits( - mixed_qkvzba, - mixed_qkvzba[..., ba_start:], - qkv_size, - z_size, - ba_size, - ) - else: - z = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, z_start, z_size) - b = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, ba_start, ba_size) - a = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, a_start, ba_size) + z = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, z_start, z_size) + b = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, ba_start, ba_size) + a = _sm70_compile_graph_slice_dim(mixed_qkvzba, -1, a_start, ba_size) mixed_qkv = _sm70_dump_gdn_projection_tensor( "split_mixed_qkv", layer_name, mixed_qkv From d03edb067c9683c6892972b933bc1bcf52660a5d Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Thu, 10 Sep 2026 00:26:33 +0800 Subject: [PATCH 16/16] [Doc] Pin the accepted TP2 head and integration switches Assisted-by: Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/sm70_dflash2_tp2_verifier.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/docs/design/sm70_dflash2_tp2_verifier.md b/docs/design/sm70_dflash2_tp2_verifier.md index 13552e854d..07940fcb68 100644 --- a/docs/design/sm70_dflash2_tp2_verifier.md +++ b/docs/design/sm70_dflash2_tp2_verifier.md @@ -35,6 +35,15 @@ model preparation are outside decode timing. The original baseline uses frozen copies of existing native libraries; it is not a rebuild of all main sources. Retained runtime manifests hash the actual mapped worker libraries. +The complete-round campaign explicitly holds these four switches at zero: +`VLLM_SM70_DFLASH2_QPN8_RERANK`, +`VLLM_SM70_DFLASH2_QPN8_RERANK_SHADOW`, +`VLLM_SM70_ENABLE_LM_HEAD_FASTPATH`, and `VLLM_SM70_LM_HEAD_TOP1_TC`. +These overrides are part of the measured FP32-logits contract; an automatic +reranking default is not interchangeable with the frozen endpoint. The +final main integration also keeps PR556's combined-copy, direct-output and +fixed-Gemma-norm experiments disabled for this TP2 validation. + One startup, one warmup and five measured requests per fixture gave: | Metric | release1k | MBPP28 |