diff --git a/benchmarks/profile_sm70_dflash2_context_cost.py b/benchmarks/profile_sm70_dflash2_context_cost.py new file mode 100644 index 0000000000..29de22def4 --- /dev/null +++ b/benchmarks/profile_sm70_dflash2_context_cost.py @@ -0,0 +1,257 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Fixed-prefix DFlash2 cost windows, separate from natural-output scoring.""" + +import argparse +import hashlib +import json +import time +import urllib.request +from pathlib import Path + +from prometheus_client.parser import text_string_to_metric_families + + +def build_prompt(corpus: dict, length: int) -> list[int]: + remaining = length - len(corpus["prefix"]) - len(corpus["suffix"]) + if remaining < 0: + raise ValueError("Prefix length is smaller than the task template") + repeats, tail = divmod(remaining, len(corpus["filler"])) + return ( + corpus["prefix"] + + corpus["filler"] * repeats + + corpus["filler"][:tail] + + corpus["suffix"] + ) + + +def metrics(base: str) -> dict: + with urllib.request.urlopen(base + "/metrics", timeout=30) as response: + families = text_string_to_metric_families(response.read().decode()) + result = {} + for family in families: + for sample in family.samples: + if "spec_decode" in sample.name or sample.name.endswith( + ( + "_time_seconds_sum", + "_latency_seconds_sum", + "request_prefill_kv_computed_tokens_sum", + "prefix_cache_hits_total", + ) + ): + key = sample.name + if "position" in sample.labels: + key += ":position=" + sample.labels["position"] + result[key] = result.get(key, 0.0) + sample.value + return result + + +def request(base: str, prompt: list[int], output_limit: int, seed: int) -> dict: + if len(prompt) + output_limit > 262144: + raise ValueError("The complete request exceeds the 256K capacity") + payload = { + "model": "quasar-baseline", + "prompt": prompt, + "max_tokens": output_limit, + "temperature": 1.0, + "top_p": 0.95, + "top_k": 20, + "seed": seed, + "stream": True, + "stream_options": {"include_usage": True}, + "return_token_ids": True, + } + before = metrics(base) + start = time.perf_counter() + chunks, token_ids, text = [], [], [] + usage, finish = {}, None + req = urllib.request.Request( + base + "/v1/completions", + data=json.dumps(payload).encode(), + headers={"Content-Type": "application/json"}, + ) + with urllib.request.urlopen(req, timeout=3600) as response: + for raw in response: + now = time.perf_counter() + line = raw.decode().strip() + if not line.startswith("data: "): + continue + if line[6:] == "[DONE]": + break + event = json.loads(line[6:]) + if event.get("error"): + raise RuntimeError(event["error"]) + if event.get("usage"): + usage = event["usage"] + for choice in event.get("choices", []): + ids = choice.get("token_ids") or [] + content = choice.get("text") or "" + if ids or content: + chunks.append({"at_s": now - start, "tokens": len(ids)}) + token_ids.extend(ids) + text.append(content) + if choice.get("finish_reason"): + finish = choice["finish_reason"] + wall = time.perf_counter() - start + after = metrics(base) + delta = {key: after.get(key, 0) - before.get(key, 0) for key in before | after} + assert usage["prompt_tokens"] == len(prompt) + assert usage["completion_tokens"] == len(token_ids) + assert chunks and token_ids + rounds = delta.get("vllm:spec_decode_num_drafts_total", 0) + accepted = delta.get("vllm:spec_decode_num_accepted_tokens_total", 0) + proposed = delta.get("vllm:spec_decode_num_draft_tokens_total", 0) + decode = delta["vllm:request_decode_time_seconds_sum"] + return { + "prompt_tokens": len(prompt), + "prompt_sha256": hashlib.sha256(json.dumps(prompt).encode()).hexdigest(), + "output_limit": output_limit, + "output_tokens": len(token_ids), + "usage": usage, + "token_ids": token_ids, + "text": "".join(text), + "finish_reason": finish, + "wall_s": wall, + "ttft_s": chunks[0]["at_s"], + "engine_prefill_s": delta.get("vllm:request_prefill_time_seconds_sum"), + "prefill_computed_tokens": delta.get( + "vllm:request_prefill_kv_computed_tokens_sum" + ), + "engine_decode_s": decode, + "complete_round_ms": decode * 1000 / rounds if rounds else None, + "pure_decode_tps": (len(token_ids) - 1) / decode if decode else None, + "rounds": rounds, + "accepted_drafts_per_round": accepted / rounds if rounds else None, + "emitted_tokens_per_round": len(token_ids) / rounds if rounds else None, + "accepted_over_proposed": accepted / proposed if proposed else None, + "chunks": chunks, + "stream_intervals_match_round_count": len(chunks) - 1 == rounds, + "metric_deltas": delta, + } + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--corpus", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + parser.add_argument("--base", default="http://127.0.0.1:18215") + parser.add_argument("--lengths", type=int, nargs="+", required=True) + parser.add_argument("--output-tokens", type=int, default=256) + parser.add_argument("--repeats", type=int, default=3) + parser.add_argument("--trace-dir", type=Path) + parser.add_argument("--reset-prefix-cache-before-length", action="store_true") + parser.add_argument("--require-native-prefill", action="store_true") + parser.add_argument("--require-original-gdn-prefill", action="store_true") + args = parser.parse_args() + if args.require_original_gdn_prefill and not args.require_native_prefill: + parser.error("--require-original-gdn-prefill needs --require-native-prefill") + assert not args.output.exists(), args.output + corpus_bytes = args.corpus.read_bytes() + corpus = json.loads(corpus_bytes) + for _ in range(900): + try: + with urllib.request.urlopen(args.base + "/health", timeout=2) as response: + if response.status == 200: + break + except (OSError, TimeoutError): + pass + time.sleep(2) + else: + raise TimeoutError("Context cost server did not become healthy") + report = { + "scope": "Bounded latency diagnostic, not natural-output quality scoring", + "context_capacity": 262144, + "profiler": args.trace_dir is not None, + "reset_prefix_cache_before_length": args.reset_prefix_cache_before_length, + "require_native_prefill": args.require_native_prefill, + "require_original_gdn_prefill": args.require_original_gdn_prefill, + "sampling": {"temperature": 1.0, "top_p": 0.95, "top_k": 20, "seed": 0}, + "corpus_sha256": hashlib.sha256(corpus_bytes).hexdigest(), + "cases": [], + "complete": False, + } + + def save(): + args.output.write_text(json.dumps(report, ensure_ascii=False, indent=2) + "\n") + + def route_snapshot(): + req = urllib.request.Request( + args.base + "/collective_rpc", + data=json.dumps( + {"method": "dflash2_prefill_route_snapshot", "timeout": 60} + ).encode(), + headers={"Content-Type": "application/json"}, + ) + with urllib.request.urlopen(req, timeout=90) as response: + rows = json.load(response)["results"] + assert sorted(row["rank"] for row in rows) == list(range(4)), rows + assert all(row["native_prefill_available"] for row in rows), rows + if args.require_original_gdn_prefill: + assert all(row["gdn_prefill"]["original_tilelang"] for row in rows), rows + return sorted(rows, key=lambda row: row["rank"]) + + if args.require_native_prefill: + report["initial_routes"] = route_snapshot() + save() + + for length in args.lengths: + prompt = build_prompt(corpus, length) + if args.reset_prefix_cache_before_length: + req = urllib.request.Request( + args.base + "/reset_prefix_cache", data=b"", method="POST" + ) + with urllib.request.urlopen(req, timeout=60) as response: + assert response.status == 200 + routes_before = route_snapshot() if args.require_native_prefill else None + for repeat in range(-1, args.repeats): + if args.trace_dir and repeat == 0: + (args.trace_dir / "arm").touch() + row = request(args.base, prompt, args.output_tokens, 0) + row.update(warmup=repeat < 0, repeat=repeat) + report["cases"].append(row) + save() + if args.reset_prefix_cache_before_length and repeat < 0: + # The HTTP reset response alone does not prove cache eviction. + assert row["prefill_computed_tokens"] == length, row + if args.require_native_prefill and repeat < 0: + row["routes_before"] = routes_before + row["routes_after"] = route_snapshot() + save() + if length >= 32768: + for before, after in zip(routes_before, row["routes_after"]): + hits = sum( + count - before["routes"].get(name, 0) + for name, count in after["routes"].items() + if name.startswith("prefill_prefix_fp8_bridge_exact_") + ) + assert hits > 0, (before, after) + print( + json.dumps( + { + key: value + for key, value in row.items() + if key + not in ( + "text", + "token_ids", + "chunks", + "metric_deltas", + "routes_before", + "routes_after", + ) + } + ), + flush=True, + ) + if args.trace_dir: + for rank in range(4): + path = args.trace_dir / f"rank{rank}-rounds.json" + observations = json.loads(path.read_text()) + assert [row["step"] for row in observations] == list(range(21)) + assert all(row["scheduled"] == 8 for row in observations) + report["complete"] = True + save() + + +if __name__ == "__main__": + main() diff --git a/benchmarks/sm70_dflash2_context_trace.py b/benchmarks/sm70_dflash2_context_trace.py new file mode 100644 index 0000000000..b3146797ad --- /dev/null +++ b/benchmarks/sm70_dflash2_context_trace.py @@ -0,0 +1,105 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Opt-in NVTX windows that exclude long-prefill chunks from q8 traces.""" + +import functools +import importlib +import json +import os +from pathlib import Path + + +def is_q8_decode(scheduler_output) -> bool: + scheduled = scheduler_output.num_scheduled_tokens + if len(scheduled) != 1 or scheduler_output.total_num_scheduled_tokens != 8: + return False + request_id = next(iter(scheduled)) + proposals = scheduler_output.scheduled_spec_decode_tokens.get(request_id, ()) + cached = scheduler_output.scheduled_cached_reqs + emitted = dict(zip(cached.req_ids, cached.num_output_tokens)) + return len(proposals) == 7 or emitted.get(request_id, 0) > 0 + + +def install_context_trace(root: Path) -> None: + import torch + + targets = { + "vllm.v1.worker.gpu.model_runner": [ + "GPUModelRunner.execute_model", + "GPUModelRunner.sample_tokens", + "GPUModelRunner.sample", + "GPUModelRunner.postprocess_sampled", + ], + "vllm.v1.worker.gpu.cudagraph_utils": [ + "ModelCudaGraphManager.run_fullgraph", + ], + "vllm.v1.worker.gpu.spec_decode.dflash2.speculator": [ + "DFlash2Speculator.propose", + ], + } + step = -1 + armed = False + finished = False + observations = [] + + def wrap(original, label): + @functools.wraps(original) + def call(*args, **kwargs): + nonlocal step, armed, finished + rank = torch.distributed.get_rank() + if label == "GPUModelRunner.execute_model": + scheduler = args[1] + if not armed and (root / "arm").exists(): + armed = True + if armed and not finished and is_q8_decode(scheduler): + step += 1 + cached = scheduler.scheduled_cached_reqs + observations.append( + { + "step": step, + "scheduled": scheduler.total_num_scheduled_tokens, + "request_ids": list(cached.req_ids), + "computed_tokens": list(cached.num_computed_tokens), + "output_tokens": list(cached.num_output_tokens), + } + ) + if step == 8: + torch.accelerator.synchronize() + torch.distributed.barrier() + torch.cuda.cudart().cudaProfilerStart() + print( + f"CONTEXT_TRACE_START rank={rank} step={step}", flush=True + ) + elif step == 20: + torch.accelerator.synchronize() + torch.distributed.barrier() + torch.cuda.cudart().cudaProfilerStop() + finished = True + (root / f"rank{rank}-rounds.json").write_text( + json.dumps(observations, indent=2) + "\n" + ) + print(f"CONTEXT_TRACE_STOP rank={rank} step={step}", flush=True) + if not armed or finished: + return original(*args, **kwargs) + torch.cuda.nvtx.range_push(f"quasar/rank{rank}/round{step}/{label}") + try: + return original(*args, **kwargs) + finally: + torch.cuda.nvtx.range_pop() + + return call + + for module_name, labels in targets.items(): + module = importlib.import_module(module_name) + for label in labels: + owner_name, method = label.split(".") + owner = getattr(module, owner_name) + setattr(owner, method, wrap(getattr(owner, method), label)) + + +class ContextCostTraceExtension: + """No route changes; explicitly select the performance installer separately.""" + + +if trace_root := os.getenv("VLLM_DFLASH2_CONTEXT_COST_TRACE_DIR"): + install_context_trace(Path(trace_root)) diff --git a/benchmarks/sm70_dflash2_prefill_route_audit.py b/benchmarks/sm70_dflash2_prefill_route_audit.py new file mode 100644 index 0000000000..6dd28b7af8 --- /dev/null +++ b/benchmarks/sm70_dflash2_prefill_route_audit.py @@ -0,0 +1,69 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Inspect native prefill availability and route counts between requests. + +Compose this worker extension with the frozen performance extension. It does +not load missing libraries or instrument model execution; the explicit launch +configuration must resolve native dependencies before model initialization. +""" + +import hashlib +import os +from pathlib import Path + + +class PrefillRouteAuditExtension: + def dflash2_prefill_route_snapshot(self): + import torch + + from vllm.model_executor.layers.mamba.gdn import qwen_gdn_linear_attn as gdn + from vllm.v1.attention.backends import flash_attn_v100 as backend + + required = ( + "sm70_d256_splitd_n32_dense_fwd", + "sm70_d256_splitd_n32_paged_fwd", + "sm70_d256_gqa_v37_fwd", + "sm70_v37_e4m3_bridge", + ) + available = {name: hasattr(torch.ops._vllm_fa2_C, name) for name in required} + flags = { + name: os.getenv(name, default) + for name, default in ( + ("VLLM_FLASH_V100_FP8_PREFILL_BRIDGE", "1"), + ("VLLM_FLASH_V100_PREFILL_D256_GQA_V37", "1"), + ("VLLM_FLASH_V100_FA2_D256_PREFILL", "1"), + ) + } + libraries = {} + for line in Path("/proc/self/maps").read_text().splitlines(): + fields = line.split(maxsplit=5) + if len(fields) == 6 and "_vllm_fa2_C" in fields[5]: + libraries[fields[5]] = None + for name in libraries: + libraries[name] = hashlib.sha256(Path(name).read_bytes()).hexdigest() + return { + "rank": torch.distributed.get_rank(), + "pid": os.getpid(), + "operators": available, + "flags": flags, + "libraries": libraries, + "native_prefill_available": ( + all(available.values()) + and all(value == "1" for value in flags.values()) + and bool(libraries) + ), + "routes": dict(backend._route_counts), + "gdn_prefill": { + "original_tilelang": gdn._sm70_flashqla_original_prefill_enabled(), + "indexed_state": gdn._sm70_flashqla_indexed_prefill_enabled(), + "direct_output": gdn._sm70_flashqla_direct_output_enabled(), + "explicit_original_flag": os.getenv( + "VLLM_SM70_FLASHQLA_ORIGINAL_PREFILL" + ), + "legacy_original_flag": os.getenv( + "FLASH_QLA_SM70_USE_ORIGINAL_TILELANG" + ), + "configuration_only_verify_actual_hit_in_worker_log": True, + }, + "graph_route_counts_are_capture_counts": True, + } diff --git a/benchmarks/sm70_dflash2_prefill_trace.py b/benchmarks/sm70_dflash2_prefill_trace.py new file mode 100644 index 0000000000..856a379ccb --- /dev/null +++ b/benchmarks/sm70_dflash2_prefill_trace.py @@ -0,0 +1,102 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Capture a bounded window of late prefill chunks, excluding q8 decode.""" + +import functools +import importlib +import json +import os +from pathlib import Path + + +def prefill_position(scheduler): + if len(scheduler.num_scheduled_tokens) != 1: + return None + count = scheduler.total_num_scheduled_tokens + if count <= 32 or any(scheduler.scheduled_spec_decode_tokens.values()): + return None + cached = scheduler.scheduled_cached_reqs + if len(cached.req_ids) != 1 or any(cached.num_output_tokens): + return None + return cached.num_computed_tokens[0], count + + +def install_prefill_trace(root: Path, start_position: int, chunks: int): + import torch + + targets = { + "vllm.v1.worker.gpu.model_runner": [ + "GPUModelRunner.execute_model", + "GPUModelRunner.sample_tokens", + "GPUModelRunner.sample", + "GPUModelRunner.postprocess_sampled", + ], + "vllm.v1.worker.gpu.spec_decode.dflash2.speculator": [ + "DFlash2Speculator.propose", + ], + } + active = False + finished = False + step = -1 + observations = [] + + def wrap(original, label): + @functools.wraps(original) + def call(*args, **kwargs): + nonlocal active, finished, step + if finished: + return original(*args, **kwargs) + rank = torch.distributed.get_rank() + if label == "GPUModelRunner.execute_model": + position = prefill_position(args[1]) + if position is not None and position[0] >= start_position: + if not active: + torch.accelerator.synchronize() + torch.distributed.barrier() + torch.cuda.cudart().cudaProfilerStart() + active = True + step += 1 + observations.append( + { + "step": step, + "computed_tokens": position[0], + "scheduled": position[1], + } + ) + elif active: + raise RuntimeError("Prefill capture reached a non-prefill step") + if not active: + return original(*args, **kwargs) + torch.cuda.nvtx.range_push(f"prefill/rank{rank}/chunk{step}/{label}") + try: + result = original(*args, **kwargs) + finally: + torch.cuda.nvtx.range_pop() + if label == "GPUModelRunner.sample_tokens" and step == chunks - 1: + torch.accelerator.synchronize() + torch.distributed.barrier() + torch.cuda.cudart().cudaProfilerStop() + active = False + finished = True + (root / f"rank{rank}-prefill.json").write_text( + json.dumps(observations, indent=2) + "\n" + ) + print(f"PREFILL_TRACE_STOP rank={rank} chunks={chunks}", flush=True) + return result + + return call + + for module_name, labels in targets.items(): + module = importlib.import_module(module_name) + for label in labels: + owner_name, method = label.split(".") + owner = getattr(module, owner_name) + setattr(owner, method, wrap(getattr(owner, method), label)) + + +if trace_root := os.getenv("VLLM_DFLASH2_PREFILL_COST_TRACE_DIR"): + install_prefill_trace( + Path(trace_root), + int(os.getenv("VLLM_DFLASH2_PREFILL_COST_TRACE_START", "115360")), + int(os.getenv("VLLM_DFLASH2_PREFILL_COST_TRACE_CHUNKS", "4")), + ) diff --git a/docs/design/sm70_dflash2_context_cost_20260909.md b/docs/design/sm70_dflash2_context_cost_20260909.md new file mode 100644 index 0000000000..54379397e9 --- /dev/null +++ b/docs/design/sm70_dflash2_context_cost_20260909.md @@ -0,0 +1,268 @@ +# DFlash2 context-length cost audit, 2026-09-09 + +## Scope and frozen contract + +Measure growth in complete verification-round cost at 64K and 128K, with a +same-prompt-template 1K anchor. The user stopped 256K measurements after observing +abnormally slow prefill; queued 256K jobs remain held. This is a latency +diagnostic, separate from the continuing natural-output quality campaign. +No production kernel, arithmetic, weights or serving default changes here. + +The harness branch starts at main `80545c010bbf6f5ed06458d992c189d75d0eff8f`, +which contains optimization PR #556. The serving checkout remains frozen at +`a7cc5ae305149d7a9ffdf42fb224dff34e5606aa`, with the same native libraries as +the approximately 16-ms candidate. This avoids changing the runtime while +measuring context-length scaling. Target QUASAR/draft revisions are +`d8e6fbfa3e3a78899b440222b827430045a05b44` / +`dedf8df68adfb1afeaf7b7480c0a0243108177b4`. + +Use physical GPUs 4–7, four V100-SXM2-32GB GPUs, TP4/B1/q8, CUDA 12.8, +Torch 2.10.0+cu128, Python 3.12.13, E4M3 target KV, FP32 logits/state and +FP16 draft transport. Preserve the fixed Gemma reduction, packed GDN, +BV2 value tile, combined split, QPN2 cap64/publication, grouped attention, +sparse selection, context/probe overlap and CUDA graphs. The server retains +262144 total context capacity, a 4096 prefill chunk budget and four maximum +request slots with only one live request. + +The prompt uses the frozen long-context corpus's prefix, repeated filler and +coding-task suffix. Exact input lengths are 1024, 65536, 131072 and 261888 +tokens. The last point reserves 256 tokens inside the 262144 capacity; it is +not a 262144-token input followed by out-of-capacity generation. Each diagnostic +request has at most 256 output tokens, honors EOS, and retains T1/k20/p.95/seed0 +and the frozen xhigh template. Length stops receive no quality credit. Inputs +are never clipped, and all actual token counts are retained. + +## Measurement and attribution + +`benchmarks/profile_sm70_dflash2_context_cost.py` records an initial/warmup request +and three unprofiled repeats per input length. It separates TTFT, engine decode, +complete-round mean, emitted token throughput and draft acceptance, and retains +per-chunk token counts/times. Client stream intervals are transport evidence, +not instrumented GPU-round latency. No trace or tensor dump runs in this service. + +`benchmarks/sm70_dflash2_context_trace.py` collects twelve q8 rounds in a separate +service after a warmup request and eight preceding q8 rounds. The gate rejects +prefill chunks, initial eight-token prefills, shorter tail queries and multiple +requests. Only capture boundaries synchronize. Every rank retains scheduled +width, computed positions and output positions; missing rank/window evidence +fails the client instead of producing a partial trace. The first and last +captured transitions are excluded from steady attribution. + +The trace uses CUDA Graph node activity and NVTX with Nsight Systems 2025.3.1. +Attribute target graph, target head/sampling, state, draft and host work on the +same critical rank. Retain GPU event union and uncovered wall time separately. +Kernel service, phase envelopes and independent rank maxima must not be summed +as complete-round wall time. Static CTA/register/shared-memory evidence does +not establish achieved occupancy or HBM throughput. + +## Validation and artifacts + +CPU checks cover seven q8/prefill/tail/multiple-request combinations, four exact +prompt lengths and rejection beyond 262144 total tokens. Scoped pre-commit runs +on the benchmark modules. The first GPU observations below exposed a missing +native prefill dependency and do not qualify the intended fast route. + +Artifacts, private launch wrappers, checkpoints and task-local compiler caches +are retained in the task artifact archive, with absolute locations in the +private handoff manifest. Generated libraries and private cache paths are not +part of this source change. +The dataset campaign is checkpointed, preserving completed same-startup pairs +and retaining any interrupted partial case separately. Its queued continuation +remains held during the prefill route investigation. + +## Missing native prefill dependency + +The frozen service selected `FLASH_ATTN_V100`, but its `lib-v4` dependency set +did not contain FA2 and its launch did not set `VLLM_SM70_FA2_D256_LIBRARY`. +The startup log explicitly warned that +`_vllm_fa2_C::sm70_d256_splitd_n32_dense_fwd` was absent and long prefill would +use a slower fallback. The E4M3 bridge is resolved from that same missing +library, so it was unavailable too. The later logs show direct paged E4M3 +prefix prefill, without the D256/v37 bridge route. + +The warning was missed before the long sweep. The partial unprofiled results +are retained as fallback observations, not expected Flash-V100 scaling: + +| Input tokens | Complete round, ms | Pure decode, tokens/s | Accepted drafts/round | Emitted tokens/round | Initial TTFT, s | +| ---: | ---: | ---: | ---: | ---: | ---: | +| 1024 | 16.419 | 254.61 | 3.2295 | 4.1967 | 1.650 | +| 65536 | 34.972 | 142.97 | 4.0784 | 5.0196 | 81.468 | +| 131072 | 53.575 | 80.67 | 3.3390 | 4.3390 | 221.823 | + +Round/decode values are medians of three repeats in one startup. All 12 +completed requests reached the 256-output-token diagnostic cap; none receives +quality credit. Repeated token IDs and acceptance match within each length. +The initial 128K request could reuse the previous 64K prefix, so its TTFT is +not a cold-prefill measurement. Repeat TTFT also includes prefix-cache hits. +No 256K request completed and no graph-node trace was collected in this run. + +PR #548 is already merged at `8d9c3518992059105d89939e8a46d75184505d8e`. +Its CMake-built FA2 library, SHA256 +`ec00745c34b3d146b0200fb9454c1419322072b0ccf0d551d958cbe701e4e15b`, +contains Split-D dense/paged, v37 and the E4M3 bridge. Its eight v37 source +hashes match the frozen serving source. A private copy is frozen under this +audit's `native/fa2-ec00745c` directory; every other native dependency stays +fixed. This is a dependency-loading repair for the diagnostic service, not a +new kernel or a quality promotion. PR #548 disclosed a remaining model token +divergence; operator accuracy alone cannot close that model-quality gate. + +The corrected launch explicitly selects this sidecar. The client checks native +availability and the loaded FA2 SHA on every rank before long requests, then +requires actual exact-bridge route hits. Snapshots occur between requests. +It resets the prefix cache before each new input length and requires the +request's computed-prefill-token counter to equal the full input length; +an HTTP reset response alone does not prove a cold request. Full API usage, +engine prefill time and computed-token metrics are retained. The first focused +repair run stops at 128K; 256K jobs must not resume automatically. + +Missing FA2 explains the prefill fallback. It does not by itself attribute +the q8 decode slope: target verification has a separate grouped E4M3 FP32 +dispatch and still needs a same-route graph-node trace after this repair. + +## Repaired native-prefill probe + +The explicit sidecar launch completed 16 requests through 128K using harness +commit `77ae71b0060acee63881204c4668aae1fe5e3406`. All four ranks mapped the +same FA2 SHA. Every cold request's computed-prefill-token count equals its +full input length. The 32K/64K/128K requests recorded respectively 144/304/624 +v37 and exact E4M3 bridge calls per rank. Existing GPU attention/bridge tests +passed 23 cases, with both 256K cases deselected. The first test launcher +imported the unbuilt harness checkout and failed before GPU validation; +the rerun used frozen serving imports with explicit pytest import isolation. + +| Input tokens | Cold prefill, s | Cold prefill, tokens/s | Complete round, ms | Pure decode, tokens/s | Accepted drafts/round | Emitted tokens/round | +| ---: | ---: | ---: | ---: | ---: | ---: | ---: | +| 1024 | 0.318 | 3222 | 16.429 | 287.43 | 3.7778 | 4.7407 | +| 32768 | 9.670 | 3388 | 25.606 | 168.79 | 3.3729 | 4.3390 | +| 65536 | 21.166 | 3096 | 34.730 | 143.97 | 4.0784 | 5.0196 | +| 131072 | 49.842 | 2630 | 52.981 | 85.95 | 3.5536 | 4.5714 | + +Prefill is one verified cold request per length; round and decode are medians +of three warmed repeats in one startup. These are diagnosis results, not +the plan's three-startup performance acceptance. The 4096 chunk budget, +1648-token page, NVFP4 target and FP32 recurrent state are unchanged. Historical +FP8 target-only/chunk8192 results are a different contract. + +All four requests at each length have identical output-token IDs and acceptance +within that startup. Compared with the missing-library fallback, the 64K +256-token output matches, while 1K and 128K first diverge at zero-based output +positions 114 and 21. None of these capped requests receives quality credit; +there is no promotion or claim that the sampled distribution is unchanged. +The repaired run retains FP32 arithmetic and does not disable compensation to +recover speed. Raw results and route snapshots are in +`results/context-cost-fa2-repaired.json`, with the compact comparison in +`results/prefill-repair-summary.json` under the audit artifact root. + +## Historical prefill regression and precision gate + +The requested historical numbers are real: PR #548's matched release run +records 128000 / 39.778778 = 3218 tokens/s and 256000 / 107.000216 = 2393 +tokens/s. It uses an FP8 target without DFlash, chunk8192 and E4M3 KV. +PR #445 is the closer QUASAR NVFP4 + DFlash2 reference: 63482 input tokens +at 3596.5 tokens/s in the release wheel; same-host source observations at +63488 input tokens are 3611.8 and 3607.2 tokens/s. It uses chunk4096, +max_num_seqs4 and E5M2 KV. The repaired E4M3 run above is not declared to have +recovered historical prefill performance. + +The launch also explicitly sets `FLASH_QLA_SM70_USE_ORIGINAL_TILELANG=0`, +overriding the default original FlashQLA prefill route. This has a recorded +precedent in PR #477's migration worklog, where restoring the original +TileLang route recovered prefill throughput. A bounded late-128K trace +confirms the current native VLK GDN costs 210.862 ms per 3296-token chunk +(48 calls). The actual launch is 96 CTAs, 64 threads/CTA and 79 registers +per thread; an earlier commentary incorrectly inferred 48 CTAs from another +block-group configuration. The trace, not that inference, is authoritative. + +The same trace measures 665.067 ms of attention including its softmax +reduction, about 472.304 ms of remaining dense GEMM/reduction work, +114.835 ms of communication, and 35.103 ms of normalization/residual work +per rank/chunk. These are instrumented GPU service categories, not additive +end-to-end performance predictions. Critical-rank interval is 1562.795 ms, +GPU event union 1549.213 ms and uncovered time 13.582 ms. The four-rank +observations cover four real 3296-token prefills starting at positions +115360, 118656, 121952 and 125248; the middle two chunks are analyzed. + +A native-versus-original GDN screen confirms 4.819 versus 0.832 ms at the +real TP4 Q4/Hv12/D128/T3296 shape. However, an independent FP64 recurrent +reference finds larger original-TileLang errors: at T65 with weak decay, +output relative L2 is 0.04661% versus native 0.02074%, and state relative L2 +is 0.03610% versus native 0.0001018%. The T129 strong-decay case has the same +direction. The initial strict non-worsening reference gate therefore rejected +this arithmetic candidate. The user subsequently accepted output error growth +below one order of magnitude and explicitly requested measuring the historical +GDN fast route first. Output error is about 2x in this operator screen; state +error has a much larger ratio against the very small native FP32 error, so +state and model-quality conclusions remain separate from speed restoration. + +The next candidate changes the number of value columns assigned to each +thread subgroup in the existing FP32 recurrence. It preserves the 16-lane +reductions, recurrence order, FP16 output boundary and FP32 state. The private +screen requires native output and state to be bitwise equal before timing a +candidate; no arithmetic downgrade is enabled. Raw artifacts are +`results/flashqla-original-operator.json` and `results/fp32-gdn-cols.json`. + +The completed 64K q8 trace separately identifies target grouped attention as +the largest decode term: 18.798 ms of GPU service per rank/round, of which +18.471 ms is the compensated partial kernel. Its launch has 240 CTAs, +256 threads, 234 registers/thread and 30464 bytes of shared memory. These +static resources do not establish measured occupancy or memory bandwidth. +The standalone 128K q8 trace is held while prefill recovery takes priority; +the 256K jobs remain held at the user's direction. + +## Restored original-GDN prefill measurements + +The same serving source and FA2 sidecar were rerun with +`VLLM_SM70_FLASHQLA_ORIGINAL_PREFILL=1` and +`FLASH_QLA_SM70_USE_ORIGINAL_TILELANG=1`, removing the task launch's forced +serial-prefill override. The worker log confirms the original TileLang GDN +route with direct output and non-indexed state; FA2 bridge hits are recorded +on all four ranks. The native libraries, TP4/B1/q8, 4096 chunk budget, +262144 total capacity, E4M3 KV, FP32 logits/state, weights and sampling remain +fixed. No profiler or tensor dump is enabled. + +| Input tokens | Previous cold prefill, tokens/s | Restored cold prefill, tokens/s | Restored cold prefill, s | Throughput gain | Complete round, ms | Pure decode, tokens/s | Accepted drafts/round | Emitted tokens/round | +| ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | +| 1024 | 3222 | 3595 | 0.285 | 11.6% | 16.214 | 291.24 | 3.7778 | 4.7407 | +| 32768 | 3388 | 4070 | 8.051 | 20.1% | 26.068 | 148.21 | 2.8939 | 3.8788 | +| 65536 | 3096 | 3651 | 17.949 | 17.9% | 34.633 | 139.84 | 3.8627 | 4.8627 | +| 131072 | 2630 | 3015 | 43.476 | 14.6% | 53.133 | 81.34 | 3.3390 | 4.3390 | + +Each prefill value is a verified cold request with the computed-token counter +equal to the full input length. Round/decode/acceptance values are medians of +three repeats in one startup. This recovers the approximately 3600 tokens/s +64K QUASAR result and 3000+ tokens/s at 128K. It is a speed diagnostic, not +the plan's three-startup acceptance or a prediction for 256K. + +Within each length, all four requests have identical token IDs. Compared with +the serial GDN baseline, 1K is identical; 32K/64K/128K first differ at output +indices 199/126/21. The restored 64K request ends naturally after 248 tokens +with a complete Python function, whereas the control reaches the 256-token +cap inside its return expression. The other restored lengths hit the cap. +There is no claim of general quality parity from these bounded requests. +Accepted drafts/round at 32K/64K/128K change from 3.3729/4.0784/3.5536 to +2.8939/3.8627/3.3390. Different sampled continuations are a confounder, but +these observations do not establish acceptance-length non-inferiority. +Long-context decode is therefore not reported as recovered by this prefill +configuration change. + +The route snapshot now records both original-GDN flags and their resolved +selection, plus indexed-state/direct-output settings. The optional +`--require-original-gdn-prefill` guard, used with `--require-native-prefill`, +rejects a forced serial GDN launch before long measurements. These are +configuration checks; the actual original-GDN hit must still be confirmed in +the worker log. This guard was added after the above run, not retroactively +claimed as part of its launch. + +Retained reports are `results/context-cost-prefill-restored.json`, +`results/prefill-restored-summary.json` and the per-rank runtime-library +manifest. Job `zz3189-prefill-restored` completed successfully and its owned +service was stopped. The explicit fast-route launch is retained for subsequent +verification-cost work; queued 256K requests remain held. + +The precision-preserving alternatives are also retained: value-column +scheduling passes 48/48 native output/state bitwise comparisons and lowers +the T3296 exp-gate operator from 4.842 to 3.271 ms. A four-token input prefetch +preserving recurrence order also passes 48/48 and reaches 2.696 ms. Increasing +prefetch to eight passes 12/12 but regresses to 5.698 ms, so it is rejected +without a model run. These are operator results, not end-to-end gains or +enabled production paths.