Add tested native vLLM TP4 runtime, RTX PRO instructions and validation

.gitattributes

@@ -33,3 +33,5 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zip filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
+model.safetensors.index.json filter=lfs diff=lfs merge=lfs -text
+tokenizer.json filter=lfs diff=lfs merge=lfs -text

LICENSE

@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (c) 2026 Naive AI
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

README.md

@@ -0,0 +1,202 @@
+---
+license: mit
+base_model: NaiveAI/Naive-N0.5-Flash
+base_model_relation: quantized
+library_name: vllm
+pipeline_tag: text-generation
+tags:
+  - nvfp4
+  - modelopt
+  - vllm
+  - mixture-of-experts
+  - custom-code
+  - blackwell
+---
+
+# Naive-N0.5-Flash NVFP4 — vLLM on four RTX PRO 6000 Blackwell GPUs
+
+**Original model by [Naive AI](https://huggingface.co/NaiveAI). This release is built on their work.**
+
+The Naive AI team created and openly released **Naive-N0.5-Flash**, including its sparse-attention adaptation, continued pretraining, model weights and reference implementation. Their research and engineering are the foundation of this repository. Many thanks to the team for sharing the model, code and technical details with the community and making derivatives like this possible.
+
+**Please support and credit the original project:** [original model](https://huggingface.co/NaiveAI/Naive-N0.5-Flash) · [GitHub](https://github.com/NaiveAI-Labs/Naive-N0.5-Flash) · [technical report and research](https://naive.ai/en/research/) · [Naive AI](https://naive.ai/).
+
+Naive's model in turn builds on **[Xiaomi MiMo-V2.5-Base](https://huggingface.co/XiaomiMiMo/MiMo-V2.5-Base)** and the **DeepSeek** team's sparse-attention research. Credit for that foundation belongs to those teams as well. This repository is an **unofficial community quantization and serving adaptation**: its contribution is the expert-only NVFP4 conversion, validation, and native vLLM integration for RTX PRO Blackwell hardware.
+
+NVFP4 expert quantization of [NaiveAI/Naive-N0.5-Flash](https://huggingface.co/NaiveAI/Naive-N0.5-Flash), a sparse-attention derivative of Xiaomi MiMo-V2.5. The model has approximately 309B total parameters and 15.5B active parameters per token. The 49 weight shards occupy **182.57 GB** (170.03 GiB).
+
+**Serve this checkpoint with the included native vLLM model plugin and Docker recipe below.** The tested engine is vLLM 0.28.0, with tensor parallelism across four 96 GB RTX PRO 6000 Blackwell GPUs. vLLM runs scheduling, paged KV-cache allocation, sampling, prefix caching, chunked prefill and the streaming OpenAI-compatible API. The plugin supplies Naive's model implementation and NVFP4 expert kernels.
+
+Naive's architecture requires this plugin; the pinned stock vLLM image does not implement it. Installing the plugin is part of the Docker build. This repository does not provide a tested SGLang implementation. Keep the original `naive_n05_flash` architecture in `config.json`.
+
+## Run with vLLM
+
+The tested machine uses four **NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition 96 GB** GPUs (SM120), a Threadripper PRO 7955WX, 498 GiB host RAM, Linux and driver 595.91.07. Have Docker Compose and NVIDIA Container Toolkit working, with all four GPUs visible to containers. Allow roughly 250 GB of free disk for the checkpoint, container images and compilation caches. This recipe is validated on the four-GPU machine described here; smaller configurations have not been validated.
+
+Download the release using the Hugging Face CLI:
+
+```bash
+hf download piotrkosecki/Naive-N0.5-Flash-NVFP4 \
+  --revision nvfp4-vllm-v1 \
+  --local-dir ./Naive-N0.5-Flash-NVFP4
+cd Naive-N0.5-Flash-NVFP4
+sha256sum --check SHA256SUMS
+```
+
+Build the pinned runtime, check the checkpoint and hardware, and start vLLM:
+
+```bash
+export LOCAL_UID="$(id -u)"
+export LOCAL_GID="$(id -g)"
+mkdir -p runtime-cache
+docker compose build naive-vllm
+docker compose run --rm --no-deps --entrypoint python3 naive-vllm \
+  /opt/naive/scripts/preflight.py
+docker compose up -d naive-vllm
+docker compose logs -f naive-vllm
+```
+
+Free the GPUs from other large models before starting. The recipe uses approximately 81–85 GiB per GPU, rising with workspaces and requests; desktop applications require additional memory. Keep `runtime-cache`: the first start and the first request of a new shape can compile kernels and take substantially longer. Cached starts still need model loading and graph capture. Wait for the API server to become ready:
+
+```bash
+curl --fail http://127.0.0.1:8000/health
+curl --fail http://127.0.0.1:8000/version
+curl --fail http://127.0.0.1:8000/v1/models
+```
+
+`/version` reports `0.28.0`; `/v1/models` lists `Naive-N0.5-Flash-NVFP4`. The API binds to localhost, and container restart is manual. Stop it with `docker compose stop naive-vllm`.
+
+The image entrypoint is `python3 -m vllm.entrypoints.openai.api_server`. The full launch configuration is in [compose.yml](compose.yml); its vLLM arguments are:
+
+```bash
+--model /models/Naive-N0.5-Flash-NVFP4 \
+--served-model-name Naive-N0.5-Flash-NVFP4 \
+--trust-remote-code \
+--tensor-parallel-size 4 --distributed-executor-backend mp \
+--dtype bfloat16 --max-model-len 8192 \
+--max-num-batched-tokens 256 --max-num-seqs 4 \
+--kv-cache-memory-bytes 4294967296 \
+--disable-custom-all-reduce \
+--enable-prefix-caching --enable-chunked-prefill \
+--compilation-config '{"mode":0,"cudagraph_mode":"FULL_DECODE_ONLY","cudagraph_capture_sizes":[1],"cudagraph_num_of_warmups":2}' \
+--enable-auto-tool-choice --tool-call-parser mimo --reasoning-parser mimo \
+--host 127.0.0.1 --port 8000
+```
+
+## Send requests
+
+Streaming chat:
+
+```bash
+curl -N http://127.0.0.1:8000/v1/chat/completions \
+  -H 'Content-Type: application/json' \
+  -d '{
+    "model": "Naive-N0.5-Flash-NVFP4",
+    "messages": [{"role": "user", "content": "Explain how a hash table works."}],
+    "temperature": 0,
+    "max_tokens": 256,
+    "chat_template_kwargs": {"enable_thinking": false},
+    "stream": true,
+    "stream_options": {"include_usage": true}
+  }'
+```
+
+For an OpenAI-compatible client, use base URL `http://127.0.0.1:8000/v1` and model `Naive-N0.5-Flash-NVFP4`. No API key is configured; clients that require a value can use a placeholder.
+
+For thinking, set `chat_template_kwargs` to `{"enable_thinking": true}` and allow enough output tokens. This vLLM version returns separate `message.reasoning` and `message.content` fields. Streaming responses use `delta.reasoning` and `delta.content`.
+
+Tool calls use the normal OpenAI `tools` schema, for example:
+
+```bash
+curl http://127.0.0.1:8000/v1/chat/completions \
+  -H 'Content-Type: application/json' \
+  -d '{
+    "model": "Naive-N0.5-Flash-NVFP4",
+    "messages": [{"role": "user", "content": "Call get_weather for Warsaw."}],
+    "tools": [{"type": "function", "function": {
+      "name": "get_weather",
+      "description": "Get weather for a city",
+      "parameters": {"type": "object", "properties": {
+        "city": {"type": "string"}
+      }, "required": ["city"]}
+    }}],
+    "tool_choice": "auto",
+    "temperature": 0,
+    "max_tokens": 192,
+    "chat_template_kwargs": {"enable_thinking": false}
+  }'
+```
+
+## Measured performance and validation
+
+Measured on September 28, 2026 with the final release build:
+
+| Prompt tokens | Generated tokens | Decode tokens/sec | Time to first token (s) | End-to-end tokens/sec |
+|---:|---:|---:|---:|---:|
+| 76 | 256 | 56.56 | 0.074 | 55.87 |
+| 79 | 256 | 56.56 | 0.076 | 55.85 |
+| 1,685 | 256 | 56.44 | 0.484 | 51.18 |
+| 3,845 | 256 | 56.57 | 0.611 | 50.01 |
+| 7,445 | 256 | 56.44 | 1.113 | 45.46 |
+
+A second run measured **55.00–55.11 tokens/sec** during decode. Both runs exceeded 50 tokens/sec in all five cases. Raw responses, token IDs, timestamps and usage counts are in [run 1](validation/native-vllm-benchmark.json) and [run 2](validation/native-vllm-benchmark-repeat.json).
+
+These measurements use the full model in **native vLLM**, one active request, temperature 0 and 256 generated tokens per response. The client records token-bearing SSE events and checks the exact returned token IDs against vLLM's usage count. Decode rate is the number of tokens after the first event divided by the time from the first to the last token event. Prompt processing and time to first token are reported separately. Prefix caching is enabled; the long prompts share background text, and repeated prompts can reuse cached prefixes. The time-to-first-token column is not a cold-prefill benchmark. Compilation and prompt processing affect end-to-end speed.
+
+CUDA graphs currently cover **one active decode request**. Multiple active requests run through the eager path and are slower per request. Four simultaneous requests passed a functional test; the single-request speeds above do not describe concurrency throughput.
+
+Reproduce the checks while no other client is using the server:
+
+```bash
+python3 runtime/scripts/validate_api.py
+python3 runtime/scripts/benchmark_stream.py \
+  --tokens 256 --cases 5 --output native-vllm-benchmark.json
+```
+
+Validation includes full-model arithmetic, geography, Python code, structured tool calls, reasoning, longer-context generation and four concurrent requests. Tiny-model tests compare 144 greedy tokens exactly with the quantized reference while exercising graph replay, batched scheduling, chunked prefill and prefix caching. A real expert layer was also checked with tensor-parallel slicing. Reports in [validation/](validation/) distinguish these checks from full-model measurements. These are smoke and implementation checks, not an evaluation of accuracy loss versus BF16 on a benchmark suite.
+
+## Runtime limits and implementation
+
+- Tested maximum context: **8,192 total input plus output tokens**, with prompts up to 7,445 tokens in the streamed benchmark. The source model's 1M context is not validated by this deployment.
+- TP4, pipeline parallelism 1; 256-token prefill chunks; up to four scheduled sequences; 4 GiB cache budget per GPU, with 17,472 cache tokens shared across requests (about two full 8K requests). Longer concurrent workloads can require vLLM to preempt and recompute requests. Attention K/V are BF16 values stored losslessly in FP32 cache slots alongside FP8-rounded indexer values with FP32 scales. KV cache is not NVFP4.
+- Single-token execution fuses RMS normalization, Q/K/V projection and rotary embedding operations, while retaining source BF16 rounding points.
+- The plugin preserves the source sliding window of 128, sparse selection of up to 2,048 tokens, stable indexer tie ordering and original attention value scaling.
+- `NCCL_P2P_DISABLE=1` and `--disable-custom-all-reduce` are required by the tested workstation's failed direct peer-copy integrity checks. The recipe keeps those settings and uses NCCL's host transport.
+- FlashInfer 0.7.0's direct-micro expert path returned zeros for this checkpoint's TP4 shape. The plugin disables that path and uses its validated regular MMA micro/static/dynamic kernels. This workaround depends on the pinned FlashInfer version.
+- `sitecustomize.py` isolates FlashInfer 0.7 source-built kernels from the base image's older prebuilt cubins. Keep the supplied dependency pins and persistent cache directories.
+
+The [Dockerfile](runtime/Dockerfile) pins the public vLLM base image by digest. Tested versions: vLLM 0.28.0, PyTorch 2.13.0+cu130, Transformers 5.17.0, FlashInfer 0.7.0, CUTLASS DSL 4.6.2 and Triton 3.7.1. The plugin uses vLLM's [model registration mechanism](https://docs.vllm.ai/en/latest/design/plugin_system/) and is included as source in [runtime/vllm_naive/](runtime/vllm_naive/).
+
+## Quantization and provenance
+
+Source revision: [`0235b3b5ff27422b1f57cdc2acddfaf643e08356`](https://huggingface.co/NaiveAI/Naive-N0.5-Flash/tree/0235b3b5ff27422b1f57cdc2acddfaf643e08356). NVIDIA Model Optimizer revision: [`23355eda90a25c290f9b1fdfb928ad54caae7d10`](https://github.com/NVIDIA/Model-Optimizer/tree/23355eda90a25c290f9b1fdfb928ad54caae7d10).
+
+Only routed experts are quantized: NVFP4 E2M1 weights and activations, block size 16, FP8 block scales and FP32 global scales. Attention, indexer, router, embeddings and output head retain source precision. All 472 non-expert tensors match source shapes, dtypes and payload hashes. There are 36,096 packed expert matrices and 144,856 exported tensors; total tensor payload is 182,549,965,184 bytes.
+
+Layerwise max calibration used 128 sequences of 4,096 tokens (524,288 total), equally divided between Magicoder OSS Instruct and UltraChat, with seed 20260928. Exact dataset revisions and calibration hash are in [quantization_provenance.json](quantization_provenance.json). Layer 47's router was materialized in BF16 during calibration; the final router was restored byte-for-byte from the FP32 source. An exporter naming issue required reconstructing the unchanged embedding, final norm and output-head tail; all 48 calibrated layer shards were retained and the complete export was subsequently audited.
+
+The source model's MIT license is preserved in [LICENSE](LICENSE). Upstream configuration/modeling files retain their Apache-2.0 headers. Additional runtime code is MIT licensed; see [runtime/THIRD_PARTY_NOTICES.md](runtime/THIRD_PARTY_NOTICES.md).
+
+## Acknowledgments and citation
+
+This release depends on substantial work by the original model authors and the open-source inference community:
+
+- **[Naive AI / NaiveAI Team](https://github.com/NaiveAI-Labs/Naive-N0.5-Flash)** — creation and release of Naive-N0.5-Flash, the hybrid SWA/DSA adaptation, continued pretraining, checkpoint, chat template, model implementation and research documentation. Please cite their work when using this derivative.
+- **[Xiaomi MiMo](https://huggingface.co/XiaomiMiMo/MiMo-V2.5-Base)** — the MiMo-V2.5 base model on which Naive-N0.5-Flash builds.
+- **[DeepSeek](https://github.com/deepseek-ai)** — the DeepSeek Sparse Attention research underpinning Naive's sparse-attention design.
+- **[vLLM](https://github.com/vllm-project/vllm)** and **[FlashInfer](https://github.com/flashinfer-ai/flashinfer)** — the serving engine and Blackwell NVFP4 expert kernels used in this deployment.
+- **[NVIDIA Model Optimizer](https://github.com/NVIDIA/Model-Optimizer)** and **[CUTLASS](https://github.com/NVIDIA/cutlass)** — quantization tooling and GPU kernel infrastructure.
+- **[Hugging Face Transformers](https://github.com/huggingface/transformers)**, **[PyTorch](https://github.com/pytorch/pytorch)** and **[Triton](https://github.com/triton-lang/triton)** — model-loading, tensor-computation and kernel-building infrastructure.
+
+The original weights and their capabilities come from the upstream model teams. This quantized release retains Naive AI's model identity, copyright notices and license. Please direct recognition for the model and its research to the original authors.
+
+The following is the citation supplied by Naive AI in the original model card:
+
+```bibtex
+@misc{naiveai2026naiven05flash,
+  title  = {Naive-N0.5-Flash: Building Frontier AI with AI},
+  author = {{NaiveAI Team}},
+  year   = {2026},
+  url    = {https://naive.ai/en/research/}
+}
+```

SHA256SUMS

@@ -0,0 +1,49 @@
+5a5221799458f4f22a969ff5bb06556e74dc214f2168da14dac077e13661c4a5  model-layer-00000.safetensors
+ea02f0c21f403c1b080582f57875e8851a4a77056a2bbb50c1041fd8f2b342e2  model-layer-00001.safetensors
+04c06288df5a794cff43965020272a2e611e52dccc82f85ae4b3a2412b717ff0  model-layer-00002.safetensors
+afde47709532dedae9fc9b3e3b140e31ed2d0027db4461355bc8c00660d95af3  model-layer-00003.safetensors
+13f9ac46ef988032fab0a064a99101a98b0d466da3d743616213aae4d693d607  model-layer-00004.safetensors
+f44cd75dd4060562fcc108614d3104c7c4cda48ef239ee785c50336ecaecb176  model-layer-00005.safetensors
+01910d46a12010345a37136155477ef8eff738af39adf6125254b18f1e7405a8  model-layer-00006.safetensors
+017f899ba6c6890182c334299fb09daa17068e07761c741e1235c7a2ef15075e  model-layer-00007.safetensors
+fa9b19c4633a4bc4e0d83768f5a721c35e7abb049cc0e7a9be0d50d2749038a0  model-layer-00008.safetensors
+1359abf1ed7f385a2cf5d16b3863eb0db5ed34a182d1e99fadbecfd0a0b6370f  model-layer-00009.safetensors
+9f09fd2abf17a4ef2da924785444801edec40a7c52ed84c4c4057f48153d0bf0  model-layer-00010.safetensors
+1dc08baf20ef80bec9c1219f60ba3267b5dc4eaf7dc8cdf4a0c933c05851cd6b  model-layer-00011.safetensors
+654df6d13f2d250fd87a058219350147435207d5e235a74cf5dddefcbffa6afc  model-layer-00012.safetensors
+803c042c1014924977feead399e9d16198dbb63cad43ed2c2340cdcfcb3487ef  model-layer-00013.safetensors
+67ef2e7acc1bf743a8dc43145ad51ac129913fe76060aa3ab19c220b68ad7919  model-layer-00014.safetensors
+053f9b45fa2c66a1e19662327532812ac242a8649dec9c457b1c95030d6e7bdb  model-layer-00015.safetensors
+ee2d0ea2c37f897482e5ebd158a1f7069785962cd48966ebee9d6f1c5e925cb5  model-layer-00016.safetensors
+7964ccac2a882f7ed6c391717837a82a1a94d4b3063348d27abeae121c847105  model-layer-00017.safetensors
+30b2b7bf5cc74bc4564dc820c0e486ebf72815e70b57f01cc5fa1fed512412c1  model-layer-00018.safetensors
+ee74c25f2b33a5a3909083a95bbc25572bfbcc27644600a30cc07a0810c7729a  model-layer-00019.safetensors
+a60ad1d87b423286e9e0c19734021e6da48af9f96bbb101ddde8dbffc5f490bb  model-layer-00020.safetensors
+a36d517f3a1c629afc88d9ed1da8e41ebae34da3bfc28bd4cd5421161203919c  model-layer-00021.safetensors
+c17ef7ede6621cbef4af19e1c1cb90262c5c1a8569e96305314399d9e4743117  model-layer-00022.safetensors
+a3945a6ccd6f8eb49f2fd031c7157c67fdf8be240a5aaf3d264e9d8cef9363ff  model-layer-00023.safetensors
+96de6d3bb2d725f181c8c573ed367db5b3c6bdcffe4664b2ac4b79732fe46928  model-layer-00024.safetensors
+38b03b993fc0d1faa1561d68dadd92d16016d6c8f1ff45bb878f50eac5187d97  model-layer-00025.safetensors
+18a984e0005858c6966cd3583d102cf1449cfb33b4eb956da39fbb7f0201f7b0  model-layer-00026.safetensors
+71b43db7c49d79c9f4e781607368fa427689bc83f04170e07d8f492b004c715b  model-layer-00027.safetensors
+9721a3937de1e68c04b7f27e473c66f027a99db08c183a62f5066a618d204c2f  model-layer-00028.safetensors
+6379291e260ff9c7e6464fb8575bcda918cd2cced4a8aa4c4df25b753fd3c6a8  model-layer-00029.safetensors
+5a8d77fa6ad0d1ade6fc09412bbc024234cf29500e72f275eb7d0f29635396d3  model-layer-00030.safetensors
+266f83df2664128a809e73c5a4a66d52f561e12e490d097e767210076d8ac3c7  model-layer-00031.safetensors
+7fbbe022dee5c5a1abc2f34e4e8233f0889c3a6c136f4a8ade2a30fceb47198f  model-layer-00032.safetensors
+2c9ee12a7ffab178f83ae09b4b32f7d47d0d3f1d66f015b37219684b4fb2fbe5  model-layer-00033.safetensors
+27741fe3a2ce23fd6aec824eed1c6368027d5b4c7ea7aa2f789ad86058cad2a7  model-layer-00034.safetensors
+73de4529e86bad56ea9214e72173f1395e118aa560267eff3089b05d1d895b55  model-layer-00035.safetensors
+e342456af59dd2f1cd9af768495671a55247ebf44005c91d951d8b1fa203022b  model-layer-00036.safetensors
+b1b31e1a4c70987d53af10c6b69da05b4a6089eaac1beec0fdcbcafc639e17eb  model-layer-00037.safetensors
+2f96937c417f671c23184c6b5c0ae7557c3d9647d59b689a10a3ce47eb40fe7c  model-layer-00038.safetensors
+bd1cd8474660df1ae8211ca1ca94b3340727f8f8e35ebe00022837d837952c95  model-layer-00039.safetensors
+5cdcdb39f00310dcce5f0e744a063b83d3fbed1ea30d888827a807302775f2b5  model-layer-00040.safetensors
+b08cdd0830ba4d69de116de6bc84faa92882c8ce44e42c56fa93639067b94225  model-layer-00041.safetensors
+142512bab040f4024047142e4f4da369d0d088b87773c16863716c2b9e56379b  model-layer-00042.safetensors
+ee5c93564e067b630d5f7638711b9d1c01ec6a78beff5714ce749c2c97279b1f  model-layer-00043.safetensors
+59b9a9e0120ba64dd6661c57b2c30522a03a8eee6cf3d7d93e1ea267af72a148  model-layer-00044.safetensors
+c63b3fe88aef75eb0066d7dd5abe30b80d174b9e82be843cd43f0f7e5977b56e  model-layer-00045.safetensors
+8f53dd271c1f77f684d70673817f0f74142c7ca78e1f2fde6af422e201f2ad38  model-layer-00046.safetensors
+5a2dc73e627443f2f95dee391123ef4e069d2ad03bffaf13026b6e012ea22803  model-layer-00047.safetensors
+6394124320964e6dbed2069d11d045c1ee2e8ec4a580c293e2987f1ea1ab6d9c  model-tail.safetensors

chat_template.jinja

@@ -0,0 +1,160 @@
+{%- set effective_reasoning_effort = reasoning_effort if reasoning_effort is defined and reasoning_effort in ['low', 'high'] else 'max' -%}
+{%- if effective_reasoning_effort is not none -%}
+    {{- "<|im_start|>system\nReasoning Effort: " ~ (effective_reasoning_effort | capitalize) ~ "<|im_end|>" }}
+{%- endif -%}
+{%- if not add_generation_prompt is defined -%}
+    {%- set add_generation_prompt = false -%}
+{%- endif -%}
+{%- if not enable_thinking is defined -%}
+    {%- set enable_thinking = true -%}
+{%- endif -%}
+{%- if not keep_all_reasoning is defined -%}
+    {%- set keep_all_reasoning = true -%}
+{%- endif -%}
+{%- macro render_extra_keys(json_dict, handled_keys) -%}
+    {%- if json_dict is mapping %}
+        {%- for json_key in json_dict if json_key not in handled_keys %}
+            {%- if json_dict[json_key] is mapping or (json_dict[json_key] is sequence and json_dict[json_key] is not string) %}
+                {{- '\n<' ~ json_key ~ '>' ~ (json_dict[json_key] | tojson | safe) ~ '</' ~ json_key ~ '>' }}
+            {%- else %}
+                {{-'\n<' ~ json_key ~ '>' ~ (json_dict[json_key] | string) ~ '</' ~ json_key ~ '>' }}
+            {%- endif %}
+        {%- endfor %}
+    {%- endif %}
+{%- endmacro -%}
+{%- macro render_content(message_content) -%}
+    {%- if message_content is string -%}
+        {{- message_content -}}
+    {%- else -%}
+        {%- for content in message_content -%}
+            {%- if 'text' in content -%}
+                {{- content['text'] -}}
+            {%- endif -%}
+        {%- endfor -%}
+    {%- endif -%}
+{%- endmacro -%}
+{%- if messages[0]["role"] == "system" %}
+    {%- set system_message = messages[0]["content"] %}
+    {%- set loop_messages = messages[1:] %}
+{%- else %}
+    {%- set loop_messages = messages %}
+{%- endif %}
+{%- set ns = namespace(last_user_index=-1) %}
+{%- for m in loop_messages %}
+    {%- if m.role == 'user' %}
+        {%- set ns.last_user_index = loop.index0 -%}
+    {%- endif %}
+{%- endfor %}
+{%- if not tools is defined %}
+    {%- set tools = [] %}
+{%- endif %}
+{%- if system_message is defined %}
+    {{- "<|im_start|>system\n" + render_content(system_message) }}
+{%- else %}
+    {{- "<|im_start|>system\nYou are Naive-N0.5-Flash, a helpful AI assistant engineered by Naive AI." }}
+{%- endif %}
+{%- if tools is iterable and tools | length > 0 %}
+    {{- "\n\n# Tools\n\nYou may call one or more functions to assist with the user query.\n\nYou have access to the following functions:\n\n" }}
+    {{- "<tools>" }}
+    {%- for tool in tools %}
+        {%- if tool.function is defined %}
+            {%- set tool = tool.function %}
+        {%- endif %}
+        {{- "\n<function>\n<name>" ~ tool.name ~ "</name>" }}
+        {%- if tool.description is defined %}
+            {{- '\n<description>' ~ (tool.description | trim) ~ '</description>' }}
+        {%- endif %}
+        {{- '\n<parameters>' }}
+        {%- if tool.parameters is defined and tool.parameters is mapping and tool.parameters.properties is defined and tool.parameters.properties is mapping %}
+            {%- for param_name, param_fields in tool.parameters.properties|items %}
+                {{- '\n<parameter>' }}
+                {{- '\n<name>' ~ param_name ~ '</name>' }}
+                {%- if param_fields.type is defined %}
+                    {{- '\n<type>' ~ (param_fields.type | string) ~ '</type>' }}
+                {%- endif %}
+                {%- if param_fields.description is defined %}
+                    {{- '\n<description>' ~ (param_fields.description | trim) ~ '</description>' }}
+                {%- endif %}
+                {%- set handled_keys = ['name', 'type', 'description'] %}
+                {{- render_extra_keys(param_fields, handled_keys) }}
+                {{- '\n</parameter>' }}
+            {%- endfor %}
+        {%- endif %}
+        {%- set handled_keys = ['type', 'properties'] %}
+        {{- render_extra_keys(tool.parameters, handled_keys) }}
+        {{- '\n</parameters>' }}
+        {%- set handled_keys = ['type', 'name', 'description', 'parameters'] %}
+        {{- render_extra_keys(tool, handled_keys) }}
+        {{- '\n</function>' }}
+    {%- endfor %}
+    {{- "\n</tools>" }}
+    {{- '\n\nFor each function call, output the function name and arguments in the following format:\n<tool_call>\n<function=example_function_name>\n<parameter=example_parameter_1>value_1</parameter>\n<parameter=example_parameter_2>This is the value for the second parameter\nthat can span\nmultiple lines</parameter>\n</function>\n</tool_call>\n\n<IMPORTANT>\n- Function calls MUST follow the specified format: an inner <function=...></function> block must be nested within <tool_call></tool_call> XML tags\n- DO NOT use function calls inside <think></think> tags.\n- The value enclosed between parameter tags is preserved exactly as-is, including newlines and spaces.\n</IMPORTANT>' }}
+{%- endif %}
+{{- '<|im_end|>' }}
+{%- for message in loop_messages %}
+    {%- if message.content is string %}
+        {%- set content = message.content %}
+    {%- else %}
+        {%- set content = render_content(message.content) %}
+    {%- endif %}
+    {%- if message.role == "assistant" %}
+        {%- if message.reasoning_content is string %}
+            {%- set reasoning_content = message.reasoning_content %}
+        {%- else %}
+            {%- set reasoning_content = '' %}
+            {%- if '</think>' in content %}
+                {%- set reasoning_content = content.split('</think>')[0].split('<think>')[-1] %}
+                {%- set content = content.split('</think>')[-1] %}
+            {%- endif %}
+        {%- endif %}
+        {%- if (keep_all_reasoning or loop.index0 > ns.last_user_index) and reasoning_content -%}
+            {{- '<|im_start|>' + message.role + '\n<think>' + reasoning_content + '</think>' + content }}
+        {%- else %}
+            {{- '<|im_start|>' + message.role + '\n<think></think>' + content }}
+        {%- endif %}
+        {%- if message.tool_calls is defined and message.tool_calls is iterable and message.tool_calls | length > 0 %}
+            {%- for tool_call in message.tool_calls %}
+                {%- if tool_call.function is defined %}
+                    {%- set tool_call = tool_call.function %}
+                {%- endif %}
+                {{- '<tool_call>\n<function=' + tool_call.name + '>\n' }}
+                {%- if tool_call.arguments is defined %}
+                    {%- for args_name, args_value in tool_call.arguments|items %}
+                        {{- '<parameter=' + args_name + '>' }}
+                        {%- set args_value = args_value | tojson | safe if args_value is mapping or (args_value is sequence and args_value is not string) else args_value | string %}
+                        {{- args_value }}
+                        {{- '</parameter>\n' }}
+                    {%- endfor %}
+                {%- endif %}
+                {{- '</function>\n</tool_call>' }}
+            {%- endfor %}
+        {%- endif %}
+        {{- '<|im_end|>' }}
+    {%- elif message.role == "user" %}
+        {{- '<|im_start|>' + message.role + '\n' + render_content(message.content) + '<|im_end|>' }}
+    {%- elif message.role == "system" %}
+        {{- '<|im_start|>' + message.role + '\n' + render_content(message.content) + '<|im_end|>' }}
+    {%- elif message.role == "tool" %}
+        {%- if loop.previtem and loop.previtem.role != "tool" %}
+            {{- '<|im_start|>tool\n' }}
+        {%- endif %}
+        {{- '<tool_response>\n' }}
+        {{- render_content(message.content) }}
+        {{- '\n</tool_response>\n' }}
+        {%- if not loop.last and loop.nextitem.role != "tool" %}
+            {{- '<|im_end|>' }}
+        {%- elif loop.last %}
+            {{- '<|im_end|>' }}
+        {%- endif %}
+    {%- else %}
+        {{- '<|im_start|>' + message.role + '\n' + render_content(message.content) + '<|im_end|>' }}
+    {%- endif %}
+{%- endfor %}
+{%- if add_generation_prompt %}
+    {{- '<|im_start|>assistant\n' }}
+    {%- if not enable_thinking -%}
+        {{- '<think></think>' -}}
+    {%- else -%}
+        {{- '' -}}
+    {%- endif -%}
+{%- endif %}

compose.yml

@@ -0,0 +1,53 @@
+services:
+  naive-vllm:
+    build: ./runtime
+    image: naive-n05-vllm:local
+    container_name: naive-n05-vllm
+    user: "${LOCAL_UID:-1000}:${LOCAL_GID:-1000}"
+    network_mode: host
+    ipc: host
+    restart: "no"
+    gpus: all
+    environment:
+      CUDA_VISIBLE_DEVICES: "0,1,2,3"
+      MODEL_PATH: /models/Naive-N0.5-Flash-NVFP4
+      NCCL_P2P_DISABLE: "1"
+      OMP_NUM_THREADS: "1"
+      VLLM_FLOAT32_MATMUL_PRECISION: highest
+    volumes:
+      - ./:/models/Naive-N0.5-Flash-NVFP4:ro
+      - ./runtime-cache:/cache
+    command:
+      - --model
+      - /models/Naive-N0.5-Flash-NVFP4
+      - --served-model-name
+      - Naive-N0.5-Flash-NVFP4
+      - --trust-remote-code
+      - --tensor-parallel-size
+      - "4"
+      - --distributed-executor-backend
+      - mp
+      - --dtype
+      - bfloat16
+      - --max-model-len
+      - "8192"
+      - --max-num-batched-tokens
+      - "256"
+      - --max-num-seqs
+      - "4"
+      - --kv-cache-memory-bytes
+      - "4294967296"
+      - --disable-custom-all-reduce
+      - --enable-prefix-caching
+      - --enable-chunked-prefill
+      - --compilation-config
+      - '{"mode":0,"cudagraph_mode":"FULL_DECODE_ONLY","cudagraph_capture_sizes":[1],"cudagraph_num_of_warmups":2}'
+      - --enable-auto-tool-choice
+      - --tool-call-parser
+      - mimo
+      - --reasoning-parser
+      - mimo
+      - --host
+      - "127.0.0.1"
+      - --port
+      - "8000"

config.json

@@ -0,0 +1,342 @@
+{
+  "add_full_attention_sink_bias": false,
+  "add_swa_attention_sink_bias": true,
+  "architectures": [
+    "NaiveN05FlashForCausalLM"
+  ],
+  "attention_bias": false,
+  "attention_chunk_size": 128,
+  "attention_dropout": 0.0,
+  "attention_projection_layout": "split",
+  "attention_value_scale": 0.707,
+  "auto_map": {
+    "AutoConfig": "configuration_naive_n05_flash.NaiveN05FlashConfig",
+    "AutoModel": "modeling_naive_n05_flash.NaiveN05FlashModel",
+    "AutoModelForCausalLM": "modeling_naive_n05_flash.NaiveN05FlashForCausalLM"
+  },
+  "dtype": "bfloat16",
+  "enable_dsa": true,
+  "eos_token_id": 151645,
+  "head_dim": 192,
+  "hidden_act": "silu",
+  "hidden_size": 4096,
+  "hybrid_block_size": null,
+  "hybrid_layer_pattern": [
+    0,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    0
+  ],
+  "index_head_dim": 128,
+  "index_n_heads": 16,
+  "index_n_kv_heads": 1,
+  "index_top_k": 2048,
+  "indexer_activation_dtype": "fp8_e4m3",
+  "initializer_range": 0.02,
+  "intermediate_size": 16384,
+  "layer_types": [
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "sliding_attention",
+    "deepseek_sparse_attention"
+  ],
+  "layernorm_epsilon": 1e-05,
+  "max_position_embeddings": 1048576,
+  "model_type": "naive_n05_flash",
+  "moe_intermediate_size": 2048,
+  "moe_layer_freq": [
+    0,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1,
+    1
+  ],
+  "n_group": 1,
+  "n_routed_experts": 256,
+  "n_shared_experts": null,
+  "norm_topk_prob": true,
+  "num_attention_heads": 64,
+  "num_experts_per_tok": 8,
+  "num_hidden_layers": 48,
+  "num_key_value_heads": 4,
+  "pad_token_id": 151643,
+  "partial_rotary_factor": 0.334,
+  "rope_parameters": {
+    "rope_type": "default",
+    "type": "default"
+  },
+  "rope_theta": 10000000,
+  "routed_scaling_factor": 1.0,
+  "scoring_func": "sigmoid",
+  "sliding_window": 128,
+  "sliding_window_size": 128,
+  "swa_head_dim": 192,
+  "swa_num_attention_heads": 64,
+  "swa_num_key_value_heads": 8,
+  "swa_rope_theta": 10000,
+  "swa_v_head_dim": 128,
+  "tie_word_embeddings": false,
+  "topk_group": 1,
+  "topk_method": "noaux_tc",
+  "transformers_version": "5.17.0",
+  "use_cache": true,
+  "v_head_dim": 128,
+  "vocab_size": 152576,
+  "quantization_config": {
+    "config_groups": {
+      "group_0": {
+        "input_activations": {
+          "dynamic": false,
+          "num_bits": 4,
+          "type": "float",
+          "group_size": 16
+        },
+        "weights": {
+          "dynamic": false,
+          "num_bits": 4,
+          "type": "float",
+          "group_size": 16
+        },
+        "targets": [
+          "Linear"
+        ]
+      }
+    },
+    "ignore": [
+      "lm_head",
+      "model.embed_tokens",
+      "model.layers.0*",
+      "model.layers.1.mlp.gate",
+      "model.layers.1.self_attn*",
+      "model.layers.2.mlp.gate",
+      "model.layers.2.self_attn*",
+      "model.layers.3.mlp.gate",
+      "model.layers.3.self_attn*",
+      "model.layers.4.mlp.gate",
+      "model.layers.4.self_attn*",
+      "model.layers.5.mlp.gate",
+      "model.layers.5.self_attn*",
+      "model.layers.6.mlp.gate",
+      "model.layers.6.self_attn*",
+      "model.layers.7.mlp.gate",
+      "model.layers.7.self_attn*",
+      "model.layers.8.mlp.gate",
+      "model.layers.8.self_attn*",
+      "model.layers.9.mlp.gate",
+      "model.layers.9.self_attn*",
+      "model.layers.10.mlp.gate",
+      "model.layers.10.self_attn*",
+      "model.layers.11.mlp.gate",
+      "model.layers.11.self_attn*",
+      "model.layers.12.mlp.gate",
+      "model.layers.12.self_attn*",
+      "model.layers.13.mlp.gate",
+      "model.layers.13.self_attn*",
+      "model.layers.14.mlp.gate",
+      "model.layers.14.self_attn*",
+      "model.layers.15.mlp.gate",
+      "model.layers.15.self_attn*",
+      "model.layers.16.mlp.gate",
+      "model.layers.16.self_attn*",
+      "model.layers.17.mlp.gate",
+      "model.layers.17.self_attn*",
+      "model.layers.18.mlp.gate",
+      "model.layers.18.self_attn*",
+      "model.layers.19.mlp.gate",
+      "model.layers.19.self_attn*",
+      "model.layers.20.mlp.gate",
+      "model.layers.20.self_attn*",
+      "model.layers.21.mlp.gate",
+      "model.layers.21.self_attn*",
+      "model.layers.22.mlp.gate",
+      "model.layers.22.self_attn*",
+      "model.layers.23.mlp.gate",
+      "model.layers.23.self_attn*",
+      "model.layers.24.mlp.gate",
+      "model.layers.24.self_attn*",
+      "model.layers.25.mlp.gate",
+      "model.layers.25.self_attn*",
+      "model.layers.26.mlp.gate",
+      "model.layers.26.self_attn*",
+      "model.layers.27.mlp.gate",
+      "model.layers.27.self_attn*",
+      "model.layers.28.mlp.gate",
+      "model.layers.28.self_attn*",
+      "model.layers.29.mlp.gate",
+      "model.layers.29.self_attn*",
+      "model.layers.30.mlp.gate",
+      "model.layers.30.self_attn*",
+      "model.layers.31.mlp.gate",
+      "model.layers.31.self_attn*",
+      "model.layers.32.mlp.gate",
+      "model.layers.32.self_attn*",
+      "model.layers.33.mlp.gate",
+      "model.layers.33.self_attn*",
+      "model.layers.34.mlp.gate",
+      "model.layers.34.self_attn*",
+      "model.layers.35.mlp.gate",
+      "model.layers.35.self_attn*",
+      "model.layers.36.mlp.gate",
+      "model.layers.36.self_attn*",
+      "model.layers.37.mlp.gate",
+      "model.layers.37.self_attn*",
+      "model.layers.38.mlp.gate",
+      "model.layers.38.self_attn*",
+      "model.layers.39.mlp.gate",
+      "model.layers.39.self_attn*",
+      "model.layers.40.mlp.gate",
+      "model.layers.40.self_attn*",
+      "model.layers.41.mlp.gate",
+      "model.layers.41.self_attn*",
+      "model.layers.42.mlp.gate",
+      "model.layers.42.self_attn*",
+      "model.layers.43.mlp.gate",
+      "model.layers.43.self_attn*",
+      "model.layers.44.mlp.gate",
+      "model.layers.44.self_attn*",
+      "model.layers.45.mlp.gate",
+      "model.layers.45.self_attn*",
+      "model.layers.46.mlp.gate",
+      "model.layers.46.self_attn*",
+      "model.layers.47.mlp.gate",
+      "model.layers.47.self_attn*"
+    ],
+    "quant_algo": "NVFP4",
+    "producer": {
+      "name": "modelopt",
+      "version": "0.0.1.dev1+g23355eda9"
+    },
+    "quant_method": "modelopt"
+  }
+}

configuration_naive_n05_flash.py

@@ -0,0 +1,155 @@
+# Copyright 2026 Naive AI.
+# Copyright 2026 The HuggingFace Inc. team.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from transformers.configuration_utils import PretrainedConfig
+
+
+class NaiveN05FlashConfig(PretrainedConfig):
+    model_type = "naive_n05_flash"
+    keys_to_ignore_at_inference = ["past_key_values"]
+    attribute_map = {"num_local_experts": "n_routed_experts"}
+
+    def __init__(
+        self,
+        vocab_size: int = 152576,
+        hidden_size: int = 4096,
+        intermediate_size: int = 16384,
+        num_hidden_layers: int = 48,
+        num_attention_heads: int = 64,
+        num_key_value_heads: int = 4,
+        head_dim: int = 192,
+        v_head_dim: int = 128,
+        swa_num_attention_heads: int = 64,
+        swa_num_key_value_heads: int = 8,
+        swa_head_dim: int = 192,
+        swa_v_head_dim: int = 128,
+        hidden_act: str = "silu",
+        max_position_embeddings: int = 1048576,
+        initializer_range: float = 0.02,
+        layernorm_epsilon: float = 1e-5,
+        rope_theta: float = 10000000.0,
+        swa_rope_theta: float = 10000.0,
+        partial_rotary_factor: float = 0.334,
+        sliding_window: int = 128,
+        hybrid_layer_pattern: list[int] | None = None,
+        attention_bias: bool = False,
+        attention_projection_layout: str = "split",
+        attention_value_scale: float | None = 0.707,
+        add_full_attention_sink_bias: bool = False,
+        add_swa_attention_sink_bias: bool = True,
+        n_routed_experts: int = 256,
+        moe_intermediate_size: int = 2048,
+        num_experts_per_tok: int = 8,
+        routed_scaling_factor: float | None = 1.0,
+        scoring_func: str = "sigmoid",
+        topk_method: str = "noaux_tc",
+        n_group: int = 1,
+        topk_group: int = 1,
+        norm_topk_prob: bool = True,
+        moe_layer_freq: list[int] | None = None,
+        enable_dsa: bool = True,
+        index_top_k: int = 2048,
+        index_head_dim: int = 128,
+        index_n_heads: int = 16,
+        index_n_kv_heads: int = 1,
+        indexer_activation_dtype: str = "fp8_e4m3",
+        use_cache: bool = True,
+        tie_word_embeddings: bool = False,
+        pad_token_id: int | None = 151643,
+        eos_token_id: int | list[int] | None = 151645,
+        **kwargs,
+    ) -> None:
+        self.vocab_size = vocab_size
+        self.hidden_size = hidden_size
+        self.intermediate_size = intermediate_size
+        self.num_hidden_layers = num_hidden_layers
+        self.num_attention_heads = num_attention_heads
+        self.num_key_value_heads = num_key_value_heads
+        self.head_dim = head_dim
+        self.v_head_dim = v_head_dim
+        self.swa_num_attention_heads = swa_num_attention_heads
+        self.swa_num_key_value_heads = swa_num_key_value_heads
+        self.swa_head_dim = swa_head_dim
+        self.swa_v_head_dim = swa_v_head_dim
+        self.hidden_act = hidden_act
+        self.max_position_embeddings = max_position_embeddings
+        self.initializer_range = initializer_range
+        self.layernorm_epsilon = layernorm_epsilon
+        self.rope_theta = rope_theta
+        self.swa_rope_theta = swa_rope_theta
+        self.partial_rotary_factor = partial_rotary_factor
+        self.sliding_window = sliding_window
+        self.hybrid_layer_pattern = hybrid_layer_pattern or [
+            int(i != 0 and i % 6 != 5) for i in range(num_hidden_layers)
+        ]
+        self.attention_bias = attention_bias
+        self.attention_projection_layout = attention_projection_layout
+        self.attention_value_scale = attention_value_scale
+        self.add_full_attention_sink_bias = add_full_attention_sink_bias
+        self.add_swa_attention_sink_bias = add_swa_attention_sink_bias
+        self.n_routed_experts = n_routed_experts
+        self.moe_intermediate_size = moe_intermediate_size
+        self.num_experts_per_tok = num_experts_per_tok
+        self.routed_scaling_factor = routed_scaling_factor or 1.0
+        self.scoring_func = scoring_func
+        self.topk_method = topk_method
+        self.n_group = n_group
+        self.topk_group = topk_group
+        self.norm_topk_prob = norm_topk_prob
+        self.moe_layer_freq = moe_layer_freq or ([0] + [1] * (num_hidden_layers - 1))
+        self.enable_dsa = enable_dsa
+        self.index_top_k = index_top_k
+        self.index_head_dim = index_head_dim
+        self.index_n_heads = index_n_heads
+        self.index_n_kv_heads = index_n_kv_heads
+        self.indexer_activation_dtype = indexer_activation_dtype
+        self.use_cache = use_cache
+        super().__init__(
+            tie_word_embeddings=tie_word_embeddings, pad_token_id=pad_token_id, eos_token_id=eos_token_id, **kwargs
+        )
+        self.validate()
+
+    def validate(self):
+        if self.attention_projection_layout != "split" or self.index_n_kv_heads != 1:
+            raise ValueError("NaiveN05Flash requires split Q/K/V and one indexer KV head")
+        if not self.enable_dsa or self.index_top_k <= 0:
+            raise ValueError("NaiveN05Flash requires DSA with a positive index_top_k")
+        if self.indexer_activation_dtype not in ("bf16", "fp8_e4m3"):
+            raise ValueError("indexer_activation_dtype must be bf16 or fp8_e4m3")
+        if self.scoring_func != "sigmoid" or self.topk_method != "noaux_tc":
+            raise ValueError("NaiveN05Flash requires sigmoid routing with correction bias")
+        if self.n_group != 1 or self.topk_group != 1:
+            raise ValueError("NaiveN05Flash uses one expert group")
+        for prefix in ("", "swa_"):
+            heads = getattr(self, prefix + "num_attention_heads")
+            kv_heads = getattr(self, prefix + "num_key_value_heads")
+            head_dim = getattr(self, prefix + "head_dim")
+            rotary_dim = int(head_dim * self.partial_rotary_factor)
+            if kv_heads <= 0 or heads % kv_heads or rotary_dim % 2 or not 0 < rotary_dim <= head_dim:
+                raise ValueError("Invalid attention head or rotary dimensions")
+        if int(self.head_dim * self.partial_rotary_factor) > self.index_head_dim:
+            raise ValueError("Indexer head dimension must cover the rotary dimension")
+        if self._attn_implementation not in (None, "eager"):
+            raise ValueError("NaiveN05Flash supports eager attention only")
+        if len(self.hybrid_layer_pattern) != self.num_hidden_layers:
+            raise ValueError("hybrid_layer_pattern must contain one entry per layer")
+        if len(self.moe_layer_freq) != self.num_hidden_layers:
+            raise ValueError("moe_layer_freq must contain one entry per layer")
+        self.layer_types = [
+            "sliding_attention" if is_swa else "deepseek_sparse_attention" for is_swa in self.hybrid_layer_pattern
+        ]
+
+
+__all__ = ["NaiveN05FlashConfig"]

generation_config.json

@@ -0,0 +1,9 @@
+{
+  "do_sample": true,
+  "temperature": 1.0,
+  "top_p": 0.95,
+  "repetition_penalty": 1.0,
+  "eos_token_id": 151645,
+  "pad_token_id": 151643,
+  "max_new_tokens": 131072
+}

hf_quant_config.json

@@ -0,0 +1,110 @@
+{
+  "producer": {
+    "name": "modelopt",
+    "version": "0.0.1.dev1+g23355eda9"
+  },
+  "quantization": {
+    "quant_algo": "NVFP4",
+    "kv_cache_quant_algo": null,
+    "group_size": 16,
+    "exclude_modules": [
+      "lm_head",
+      "model.embed_tokens",
+      "model.layers.0*",
+      "model.layers.1.mlp.gate",
+      "model.layers.1.self_attn*",
+      "model.layers.2.mlp.gate",
+      "model.layers.2.self_attn*",
+      "model.layers.3.mlp.gate",
+      "model.layers.3.self_attn*",
+      "model.layers.4.mlp.gate",
+      "model.layers.4.self_attn*",
+      "model.layers.5.mlp.gate",
+      "model.layers.5.self_attn*",
+      "model.layers.6.mlp.gate",
+      "model.layers.6.self_attn*",
+      "model.layers.7.mlp.gate",
+      "model.layers.7.self_attn*",
+      "model.layers.8.mlp.gate",
+      "model.layers.8.self_attn*",
+      "model.layers.9.mlp.gate",
+      "model.layers.9.self_attn*",
+      "model.layers.10.mlp.gate",
+      "model.layers.10.self_attn*",
+      "model.layers.11.mlp.gate",
+      "model.layers.11.self_attn*",
+      "model.layers.12.mlp.gate",
+      "model.layers.12.self_attn*",
+      "model.layers.13.mlp.gate",
+      "model.layers.13.self_attn*",
+      "model.layers.14.mlp.gate",
+      "model.layers.14.self_attn*",
+      "model.layers.15.mlp.gate",
+      "model.layers.15.self_attn*",
+      "model.layers.16.mlp.gate",
+      "model.layers.16.self_attn*",
+      "model.layers.17.mlp.gate",
+      "model.layers.17.self_attn*",
+      "model.layers.18.mlp.gate",
+      "model.layers.18.self_attn*",
+      "model.layers.19.mlp.gate",
+      "model.layers.19.self_attn*",
+      "model.layers.20.mlp.gate",
+      "model.layers.20.self_attn*",
+      "model.layers.21.mlp.gate",
+      "model.layers.21.self_attn*",
+      "model.layers.22.mlp.gate",
+      "model.layers.22.self_attn*",
+      "model.layers.23.mlp.gate",
+      "model.layers.23.self_attn*",
+      "model.layers.24.mlp.gate",
+      "model.layers.24.self_attn*",
+      "model.layers.25.mlp.gate",
+      "model.layers.25.self_attn*",
+      "model.layers.26.mlp.gate",
+      "model.layers.26.self_attn*",
+      "model.layers.27.mlp.gate",
+      "model.layers.27.self_attn*",
+      "model.layers.28.mlp.gate",
+      "model.layers.28.self_attn*",
+      "model.layers.29.mlp.gate",
+      "model.layers.29.self_attn*",
+      "model.layers.30.mlp.gate",
+      "model.layers.30.self_attn*",
+      "model.layers.31.mlp.gate",
+      "model.layers.31.self_attn*",
+      "model.layers.32.mlp.gate",
+      "model.layers.32.self_attn*",
+      "model.layers.33.mlp.gate",
+      "model.layers.33.self_attn*",
+      "model.layers.34.mlp.gate",
+      "model.layers.34.self_attn*",
+      "model.layers.35.mlp.gate",
+      "model.layers.35.self_attn*",
+      "model.layers.36.mlp.gate",
+      "model.layers.36.self_attn*",
+      "model.layers.37.mlp.gate",
+      "model.layers.37.self_attn*",
+      "model.layers.38.mlp.gate",
+      "model.layers.38.self_attn*",
+      "model.layers.39.mlp.gate",
+      "model.layers.39.self_attn*",
+      "model.layers.40.mlp.gate",
+      "model.layers.40.self_attn*",
+      "model.layers.41.mlp.gate",
+      "model.layers.41.self_attn*",
+      "model.layers.42.mlp.gate",
+      "model.layers.42.self_attn*",
+      "model.layers.43.mlp.gate",
+      "model.layers.43.self_attn*",
+      "model.layers.44.mlp.gate",
+      "model.layers.44.self_attn*",
+      "model.layers.45.mlp.gate",
+      "model.layers.45.self_attn*",
+      "model.layers.46.mlp.gate",
+      "model.layers.46.self_attn*",
+      "model.layers.47.mlp.gate",
+      "model.layers.47.self_attn*"
+    ]
+  }
+}

model.safetensors.index.json

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b989f8e0c01e2b54f46efcd3d134b53b23c59c8faba37df1d8a6a52a0baccdec
+size 13388642

modeling_naive_n05_flash.py

@@ -0,0 +1,243 @@
+# Copyright 2026 Naive AI.
+# Copyright 2026 The HuggingFace Inc. team.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from copy import copy
+
+import torch
+from torch import nn
+from transformers.cache_utils import DynamicCache
+from transformers.conversion_mapping import get_checkpoint_conversion_mapping, register_checkpoint_conversion_mapping
+from transformers.generation import GenerationMixin
+from transformers.modeling_outputs import CausalLMOutputWithPast
+from transformers.modeling_utils import PreTrainedModel
+from transformers.models.deepseek_v3.modeling_deepseek_v3 import DeepseekV3Experts, DeepseekV3TopkRouter
+from transformers.models.deepseek_v3.modeling_deepseek_v3 import DeepseekV3MLP as NaiveN05FlashMLP
+from transformers.models.gpt_neox.modeling_gpt_neox import GPTNeoXRotaryEmbedding as NaiveN05FlashRotaryEmbedding
+from transformers.models.gpt_neox.modeling_gpt_neox import apply_rotary_pos_emb
+from transformers.models.llama.modeling_llama import LlamaRMSNorm as NaiveN05FlashRMSNorm
+from transformers.models.llama.modeling_llama import repeat_kv
+
+from .configuration_naive_n05_flash import NaiveN05FlashConfig
+
+# Fuse experts weights
+# mlp.experts.{i}.{gate,up,down}_proj.weight[_scale_inv] -> mlp.experts.{gate,up,down}_proj.weight[_scale_inv]
+register_checkpoint_conversion_mapping("naive_n05_flash", get_checkpoint_conversion_mapping("deepseek_v3"), overwrite=True)
+
+
+def round_indexer_fp8(states):
+    states = states.float()
+    scale = states.abs().amax(dim=-1, keepdim=True).clamp_min(1e-4) / 448.0
+    return (states / scale).clamp(-448, 448).to(torch.float8_e4m3fn).float() * scale
+
+
+class NaiveN05FlashMoE(nn.Module):
+    """DeepSeek-V3-style routed experts without a shared expert."""
+
+    def __init__(self, config):
+        super().__init__()
+        self.experts = DeepseekV3Experts(config)
+        self.gate = DeepseekV3TopkRouter(config)
+
+    def forward(self, states):
+        original_shape = states.shape
+        _, topk_weights, topk_indices = self.gate(states)
+        states = states.view(-1, states.shape[-1])
+        states = self.experts(states, topk_indices, topk_weights)
+        return states.view(*original_shape)
+
+
+class NaiveN05FlashIndexer(nn.Module):
+    def __init__(self, config, layer_idx):
+        super().__init__()
+        self.config = config
+        self.layer_idx = layer_idx
+        self.wq = nn.Linear(config.hidden_size, config.index_n_heads * config.index_head_dim, bias=False)
+        self.wk = nn.Linear(config.hidden_size, config.index_head_dim, bias=False)
+        self.k_norm = nn.LayerNorm(config.index_head_dim, eps=1e-5)
+        self.weights_proj = nn.Linear(config.hidden_size, config.index_n_heads, bias=False)
+
+    def forward(self, states, position_embeddings, past_key_values=None):
+        c = self.config
+        batch, length, _ = states.shape
+        query = self.wq(states).view(batch, length, c.index_n_heads, c.index_head_dim).transpose(1, 2)
+        key = self.k_norm(self.wk(states)).unsqueeze(1)
+        query, key = apply_rotary_pos_emb(query, key, *position_embeddings)
+        if c.indexer_activation_dtype == "fp8_e4m3":
+            query, key = round_indexer_fp8(query), round_indexer_fp8(key)
+        if past_key_values is not None:
+            key = past_key_values.update_indexer(key.squeeze(1), self.layer_idx).unsqueeze(1)
+        scores = (query.float() @ key.float().transpose(-1, -2)).relu()
+        weights = self.weights_proj(states) * c.index_n_heads**-0.5
+        return (scores * weights.transpose(1, 2).unsqueeze(-1).float()).sum(dim=1)
+
+
+class NaiveN05FlashAttention(nn.Module):
+    def __init__(self, config, layer_idx):
+        super().__init__()
+        self.config = config
+        self.layer_idx = layer_idx
+        self.is_swa = bool(config.hybrid_layer_pattern[layer_idx])
+        prefix = "swa_" if self.is_swa else ""
+        self.heads = getattr(config, prefix + "num_attention_heads")
+        self.kv_heads = getattr(config, prefix + "num_key_value_heads")
+        self.head_dim = getattr(config, prefix + "head_dim")
+        self.v_dim = getattr(config, prefix + "v_head_dim")
+        rope_config = copy(config)
+        rope_config.head_dim = self.head_dim
+        rope_config.rope_parameters = {
+            "rope_type": "default",
+            "rope_theta": getattr(config, prefix + "rope_theta"),
+            "partial_rotary_factor": config.partial_rotary_factor,
+        }
+        self.rotary_emb = NaiveN05FlashRotaryEmbedding(rope_config)
+        self.q_proj = nn.Linear(config.hidden_size, self.heads * self.head_dim, bias=config.attention_bias)
+        self.k_proj = nn.Linear(
+            config.hidden_size,
+            self.kv_heads * self.head_dim,
+            bias=config.attention_bias,
+        )
+        self.v_proj = nn.Linear(config.hidden_size, self.kv_heads * self.v_dim, bias=config.attention_bias)
+        self.o_proj = nn.Linear(self.heads * self.v_dim, config.hidden_size, bias=False)
+        self.indexer = None if self.is_swa else NaiveN05FlashIndexer(config, layer_idx)
+        sink = config.add_swa_attention_sink_bias if self.is_swa else config.add_full_attention_sink_bias
+        self.attention_sink_bias = nn.Parameter(torch.zeros(self.heads), requires_grad=False) if sink else None
+
+    def forward(self, states, positions, allowed, past_key_values=None):
+        batch, length, _ = states.shape
+        past_length = 0 if past_key_values is None else past_key_values.get_seq_length(self.layer_idx)
+        key_offset = 0 if past_key_values is None else past_key_values.get_mask_sizes(length, self.layer_idx)[1]
+        query = self.q_proj(states).view(batch, length, self.heads, self.head_dim).transpose(1, 2)
+        key = self.k_proj(states).view(batch, length, self.kv_heads, self.head_dim).transpose(1, 2)
+        value = self.v_proj(states).view(batch, length, self.kv_heads, self.v_dim).transpose(1, 2)
+        position_embeddings = self.rotary_emb(states, positions)
+        query, key = apply_rotary_pos_emb(query, key, *position_embeddings)
+        if past_key_values is not None:
+            key, value = past_key_values.update(key, value, self.layer_idx)
+        key_length = key.shape[-2]
+        allowed = allowed[..., key_offset : key_offset + key_length]
+        if self.is_swa:
+            query_offsets = past_length + torch.arange(length, device=states.device)
+            key_offsets = key_offset + torch.arange(key_length, device=states.device)
+            allowed = allowed & (query_offsets[:, None] - key_offsets[None, :] < self.config.sliding_window)
+        else:
+            scores = self.indexer(states, position_embeddings, past_key_values).masked_fill(~allowed, -torch.inf)
+            # Stable ties keep padding from changing which equal-score keys win.
+            selected = scores.argsort(dim=-1, descending=True, stable=True)[..., : self.config.index_top_k]
+            allowed = allowed & torch.zeros_like(scores, dtype=torch.bool).scatter(-1, selected, True)
+        key = repeat_kv(key, self.heads // self.kv_heads)
+        value = repeat_kv(value, self.heads // self.kv_heads)
+        if self.config.attention_value_scale is not None:
+            value = value * self.config.attention_value_scale
+        logits = (query @ key.transpose(-1, -2)) * self.head_dim**-0.5
+        logits = logits.masked_fill(~allowed[:, None], -torch.inf)
+        if self.attention_sink_bias is not None:
+            sink = self.attention_sink_bias[None, :, None, None].expand(batch, -1, length, -1)
+            logits = torch.cat((logits, sink), dim=-1)
+        probabilities = logits.float().softmax(-1).nan_to_num(0.0)[..., :key_length].to(value.dtype)
+        output = (probabilities @ value).transpose(1, 2).reshape(batch, length, -1)
+        return self.o_proj(output)
+
+
+class NaiveN05FlashDecoderLayer(nn.Module):
+    def __init__(self, config, layer_idx):
+        super().__init__()
+        self.self_attn = NaiveN05FlashAttention(config, layer_idx)
+        self.mlp = NaiveN05FlashMoE(config) if config.moe_layer_freq[layer_idx] else NaiveN05FlashMLP(config)
+        self.input_layernorm = NaiveN05FlashRMSNorm(config.hidden_size, config.layernorm_epsilon)
+        self.post_attention_layernorm = NaiveN05FlashRMSNorm(config.hidden_size, config.layernorm_epsilon)
+
+    def forward(self, states, positions, allowed, past_key_values=None):
+        states = states + self.self_attn(self.input_layernorm(states), positions, allowed, past_key_values)
+        return states + self.mlp(self.post_attention_layernorm(states))
+
+
+class NaiveN05FlashModel(PreTrainedModel):
+    config_class = NaiveN05FlashConfig
+    base_model_prefix = "model"
+    _can_set_experts_implementation_cached_value = True
+
+    def __init__(self, config):
+        super().__init__(config)
+        self.embed_tokens = nn.Embedding(config.vocab_size, config.hidden_size)
+        self.layers = nn.ModuleList([NaiveN05FlashDecoderLayer(config, i) for i in range(config.num_hidden_layers)])
+        self.norm = NaiveN05FlashRMSNorm(config.hidden_size, config.layernorm_epsilon)
+
+    def forward(self, states, positions, allowed, past_key_values=None):
+        for layer in self.layers:
+            states = layer(states, positions, allowed, past_key_values)
+        return self.norm(states)
+
+
+class NaiveN05FlashForCausalLM(PreTrainedModel, GenerationMixin):
+    config_class = NaiveN05FlashConfig
+    base_model_prefix = "model"
+    _no_split_modules = ["NaiveN05FlashDecoderLayer"]
+    _skip_keys_device_placement = ["past_key_values"]
+    _keep_in_fp32_modules_strict = ["mlp.gate.weight", "mlp.gate.e_score_correction_bias"]
+    _can_set_experts_implementation_cached_value = True
+
+    def __init__(self, config):
+        config.validate()
+        super().__init__(config)
+        self.model = NaiveN05FlashModel(config)
+        self.lm_head = nn.Linear(config.hidden_size, config.vocab_size, bias=False)
+        self.post_init()
+        self.eval().requires_grad_(False)
+
+    def get_input_embeddings(self):
+        return self.model.embed_tokens
+
+    def get_output_embeddings(self):
+        return self.lm_head
+
+    @torch.inference_mode()
+    def forward(
+        self,
+        input_ids=None,
+        attention_mask=None,
+        position_ids=None,
+        inputs_embeds=None,
+        use_cache=None,
+        past_key_values=None,
+        **kwargs,
+    ):
+        use_cache = self.config.use_cache if use_cache is None else use_cache
+        if past_key_values is not None and not use_cache:
+            raise ValueError("Pass use_cache=True when providing past_key_values")
+        if use_cache and past_key_values is None:
+            past_key_values = DynamicCache(config=self.config)
+        if past_key_values is not None and not isinstance(past_key_values, DynamicCache):
+            raise ValueError("NaiveN05Flash supports DynamicCache only")
+        if (input_ids is None) == (inputs_embeds is None):
+            raise ValueError("Provide exactly one of input_ids or inputs_embeds")
+        states = self.model.embed_tokens(input_ids) if inputs_embeds is None else inputs_embeds
+        batch, length, _ = states.shape
+        past_length = 0 if past_key_values is None else past_key_values.get_seq_length()
+        total_length = past_length + length
+        if attention_mask is None:
+            attention_mask = torch.ones(batch, total_length, device=states.device, dtype=torch.bool)
+        if attention_mask.shape != (batch, total_length):
+            raise ValueError("Expected a 2-D attention_mask covering the cached prefix and new tokens")
+        if position_ids is None:
+            position_ids = (attention_mask.long().cumsum(-1) - 1).clamp_min(0)[:, -length:]
+        query_offsets = past_length + torch.arange(length, device=states.device)
+        key_offsets = torch.arange(total_length, device=states.device)
+        allowed = query_offsets[:, None] >= key_offsets[None, :]
+        allowed = allowed[None] & attention_mask[:, None].bool()
+        logits = self.lm_head(self.model(states, position_ids, allowed, past_key_values))
+        return CausalLMOutputWithPast(logits=logits, past_key_values=past_key_values)
+
+
+__all__ = ["NaiveN05FlashForCausalLM", "NaiveN05FlashModel"]

quantization_provenance.json

@@ -0,0 +1,61 @@
+{
+  "source_repository": "NaiveAI/Naive-N0.5-Flash",
+  "source_revision": "0235b3b5ff27422b1f57cdc2acddfaf643e08356",
+  "modelopt_revision": "23355eda90a25c290f9b1fdfb928ad54caae7d10",
+  "recipe": "NVFP4 experts only; max calibration; original-precision attention, router and KV cache",
+  "calibration": {
+    "sources": [
+      {
+        "id": "ise-uiuc/Magicoder-OSS-Instruct-75K",
+        "revision": "5f839b1f368a76b161028bb9edff055db34022b2",
+        "format": "json",
+        "file": "data-oss_instruct-decontaminated.jsonl"
+      },
+      {
+        "id": "HuggingFaceH4/ultrachat_200k",
+        "revision": "8049631c405ae6576f93f445c6b8166f76f5505a",
+        "format": "parquet",
+        "file": "data/train_sft-00000-of-00003-a3ecf92756993583.parquet"
+      }
+    ],
+    "samples": 128,
+    "sequence_length": 4096,
+    "tokens": 524288,
+    "sha256": "19061f6f772f68b6fe9096a01376e236cf27ba03a1c58f2e36acc5f2f0cab758",
+    "packing": "Chat-formatted records concatenated with EOS, equal coding/chat token counts",
+    "shuffle_seed": 20260928,
+    "shuffle_buffer": 128
+  },
+  "runtime": "Native vLLM 0.28.0 V1 engine, Naive model plugin, TP4, paged KV/indexer cache, FlashInfer 0.7.0 B12x NVFP4 experts, single-request decode CUDA graphs",
+  "export_recovery": {
+    "calibration_layers_reused": 48,
+    "original_precision_tail_tensors": [
+      "lm_head.weight",
+      "model.embed_tokens.weight",
+      "model.norm.weight"
+    ],
+    "tensor_bytes": 182547868032,
+    "reason": "Fused expert source names misidentified as unplaced by generic exporter"
+  },
+  "source_precision_audit": {
+    "passed": true,
+    "tensors_checked": 472,
+    "restored_from_source": {
+      "model-layer-00047.safetensors": [
+        "model.layers.47.mlp.gate.weight"
+      ]
+    },
+    "check": "Exact source dtype, shape and SHA-256 of tensor payload"
+  },
+  "calibration_note": "Layer 47 routing was materialized in BF16 during calibration; its final checkpoint router is restored byte-for-byte from the FP32 source.",
+  "final_export_validation": {
+    "passed": true,
+    "tensor_count": 144856,
+    "packed_expert_weights": 36096,
+    "safetensor_shards": 49,
+    "tensor_bytes": 182549965184,
+    "all_source_tensors_accounted_for": true
+  },
+  "runtime_performance_validation": "validation/native-vllm-benchmark.json",
+  "runtime_build": "runtime/Dockerfile"
+}

runtime-lock.json

@@ -0,0 +1,38 @@
+{
+  "base_image": "vllm/vllm-openai@sha256:61fc8a896b0a4fbbbdc063bc4b0dbc25ce98e02b5050c24aeb7830ac02039b14",
+  "tested_local_image_digest": "naive-n05-vllm@sha256:edc59ac2c49c6d62ee358ffd1ffb4a8b4a59b40a704e54da6286cf8b7f942511",
+  "image_note": "The tested derived image is locally built, not a public registry artifact. Build runtime/Dockerfile from this release.",
+  "runtime_source_sha256": {
+    ".dockerignore": "05cc6ae53cfd6129f1d7010aa874c6cf0a6019b1adfeff3f4fab27f2d04f43ca",
+    "Dockerfile": "201fa4be561e877098945d37897cc0f7a7c9e7dd15eb4779b15830717682f101",
+    "LICENSE": "969b552eb338dc04f165177043e43d2b6daab45accc29021d77c200e3b3ba68a",
+    "LICENSE-APACHE-2.0.txt": "cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30",
+    "THIRD_PARTY_NOTICES.md": "f42c84a8c8a64b743c491d7e1d7ed1e1daa8ea4546f7ec877db30d4379974ece",
+    "pyproject.toml": "90ad59f3b03570a9bb18d5d399a721d5c5daf9e2acba2f5a38d6e44c7cb8e9d3",
+    "requirements.txt": "829ba4fbb1e769da9505c900701f439a3d2126c4c7d2d15e3b539fa37a20aa18",
+    "scripts/benchmark_stream.py": "dab13abe554c6c28f89268ee3a2cec2e6425f73aa91f18692318b76886704157",
+    "scripts/decode_kernels.py": "ea0e71f5c4fd678b5cdf423b62ad07d37bed7de5fdc2af9ee89f08856ee34d09",
+    "scripts/fused_experts.py": "f77eceec0a969eb7ff99e5177bd9d4cda1bfa3f0668ed614c060d28883f46018",
+    "scripts/nvfp4_runtime.py": "be6453ca83d8f9ec528b07f17e78a41bfdcd8859fa9ec2644b6136e7068e9a1a",
+    "scripts/preflight.py": "43ac170da1d7cf614fb5f4d7f0ee70635d2d78bd638bfb3af2322baf91a08600",
+    "scripts/safe_device_transfer.py": "61e931597e59e72ded3525976181e7164933e72f4647423c00ad2be0a57940db",
+    "scripts/validate_api.py": "82e04b0d21e9956da18930e27d35c4fc45f1a4b071bd2a19c6f62271deb74c2d",
+    "sitecustomize.py": "552485e6e5e03c72c1d8b76a43a1ede48f8283a5d1e5ab93761125ec4a2c492b",
+    "vllm_naive/__init__.py": "8a5fddf2b89a4c33b6774a3ed75004a73db990ee711b8411e8772d7f2786f679",
+    "vllm_naive/attention.py": "f3aa650b110495224338b211330f6dcb7598cbeff917e7b1dac77695234f187c",
+    "vllm_naive/model.py": "f0d5769018f8487c267d0ee6c3ba4cbf138e47d65c35c0cc0a00a95dfb69d763"
+  },
+  "versions": {
+    "vllm": "0.28.0",
+    "torch": "2.13.0+cu130",
+    "transformers": "5.17.0",
+    "accelerate": "1.15.0",
+    "tokenizers": "0.23.2",
+    "huggingface-hub": "1.33.0",
+    "safetensors": "0.8.0",
+    "flashinfer-python": "0.7.0",
+    "nvidia-cutlass-dsl": "4.6.2",
+    "triton": "3.7.1"
+  },
+  "tested_at": "2026-09-28T10:42:32.236283+00:00"
+}

runtime/.dockerignore

@@ -0,0 +1,2 @@
+__pycache__
+*.pyc

runtime/Dockerfile

@@ -0,0 +1,29 @@
+FROM vllm/vllm-openai@sha256:61fc8a896b0a4fbbbdc063bc4b0dbc25ce98e02b5050c24aeb7830ac02039b14
+
+# Preserve the pinned vLLM/PyTorch/CUDA/CUTLASS/Triton stack.
+COPY requirements.txt /opt/naive/requirements.txt
+RUN python3 -m pip install --no-cache-dir --no-deps --target /opt/naive-python -r /opt/naive/requirements.txt \
+    && python3 -m pip install --no-cache-dir --no-deps --target /opt/flashinfer-v070 flashinfer-python==0.7.0
+COPY scripts/ /opt/naive/scripts/
+COPY sitecustomize.py /opt/naive-python/sitecustomize.py
+COPY pyproject.toml /opt/vllm-naive/pyproject.toml
+COPY vllm_naive/ /opt/vllm-naive/vllm_naive/
+RUN python3 -m pip install --no-deps --no-build-isolation /opt/vllm-naive
+ENV PYTHONPATH=/opt/flashinfer-v070:/opt/naive-python:/opt/naive/scripts \
+    NAIVE_FLASHINFER_SOURCE_ONLY=1 \
+    NAIVE_EXPERT_BACKEND=fused \
+    FLASHINFER_WORKSPACE_BASE=/cache/flashinfer-v070 \
+    HF_MODULES_CACHE=/cache/hf-modules \
+    TORCH_EXTENSIONS_DIR=/cache/torch-extensions \
+    TRITON_CACHE_DIR=/cache/triton \
+    XDG_CACHE_HOME=/cache/xdg \
+    HF_HUB_OFFLINE=1 \
+    TRANSFORMERS_OFFLINE=1 \
+    OMP_NUM_THREADS=1 \
+    NCCL_P2P_DISABLE=1 \
+    VLLM_PLUGINS=naive_n05 \
+    VLLM_WORKER_MULTIPROC_METHOD=spawn \
+    VLLM_CACHE_ROOT=/cache/vllm \
+    VLLM_FLOAT32_MATMUL_PRECISION=highest
+WORKDIR /opt/naive
+ENTRYPOINT ["python3", "-m", "vllm.entrypoints.openai.api_server"]

runtime/LICENSE

@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (c) 2026 Naive NVFP4 runtime contributors
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

runtime/LICENSE-APACHE-2.0.txt

@@ -0,0 +1,202 @@
+
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

runtime/THIRD_PARTY_NOTICES.md

@@ -0,0 +1,16 @@
+# Licensing and attribution
+
+The quantized model is derived from Naive AI's `Naive-N0.5-Flash` checkpoint.
+Its MIT license and copyright notice are preserved in the repository's root
+`LICENSE` file. Naive AI describes the original model as a derivative of
+Xiaomi MiMo-V2.5; see the linked upstream model card for that lineage.
+
+The upstream `configuration_naive_n05_flash.py` and
+`modeling_naive_n05_flash.py` retain their original Apache-2.0 copyright and
+license headers. The Apache-2.0 license text is in `LICENSE-APACHE-2.0.txt`.
+
+The additional runtime scripts and deployment files in `runtime/`, and the
+root `compose.yml`, are provided under the MIT terms in `runtime/LICENSE`.
+Dependencies downloaded during the Docker build retain their own licenses;
+the dependency image and Python packages are not redistributed in this
+repository.

runtime/pyproject.toml

@@ -0,0 +1,14 @@
+[build-system]
+requires = ["setuptools>=68"]
+build-backend = "setuptools.build_meta"
+
+[project]
+name = "vllm-naive-n05"
+version = "0.1.0"
+requires-python = ">=3.12"
+
+[project.entry-points."vllm.general_plugins"]
+naive_n05 = "vllm_naive:register"
+
+[tool.setuptools.packages.find]
+include = ["vllm_naive*"]

runtime/requirements.txt

@@ -0,0 +1,6 @@
+# Install into an isolated target over the pinned CUDA/PyTorch base image.
+transformers==5.17.0
+accelerate==1.15.0
+tokenizers==0.23.2
+huggingface-hub==1.33.0
+safetensors==0.8.0

runtime/scripts/benchmark_stream.py

@@ -0,0 +1,84 @@
+"""Measure real vLLM streamed token delivery, with exact returned token IDs."""
+import argparse
+import json
+import time
+import urllib.request
+from pathlib import Path
+
+PROMPTS = [
+    'Explain how a hash table works in Python, including collisions, resizing, and time complexity. Write a detailed tutorial of at least 800 words.',
+    'Write a complete Python implementation of an LRU cache with type annotations, a detailed explanation, and examples. Include at least 800 words of explanation after the code.',
+]
+for repetitions in (90, 210, 410):
+    PROMPTS.append(
+        'This is background information for a technical documentation exercise. The system stores records and processes requests.\n' * repetitions
+        + 'Now write a detailed tutorial of at least 800 words explaining database indexes and their tradeoffs.')
+
+
+def request(base_url, model, prompt, maximum):
+    body = {'model':model, 'messages':[{'role':'user','content':prompt}],
+        'temperature':0, 'max_tokens':maximum, 'chat_template_kwargs':{'enable_thinking':False},
+        'stream':True, 'stream_options':{'include_usage':True}, 'return_token_ids':True}
+    req = urllib.request.Request(base_url.rstrip('/')+'/chat/completions',
+        data=json.dumps(body).encode(), headers={'Content-Type':'application/json'})
+    start = time.perf_counter()
+    events, content, reasoning = [], [], []
+    usage = None
+    with urllib.request.urlopen(req, timeout=1800) as response:
+        for line in response:
+            if not line.startswith(b'data: '):
+                continue
+            payload = line[6:].strip()
+            if payload == b'[DONE]':
+                break
+            chunk = json.loads(payload)
+            if 'error' in chunk:
+                raise RuntimeError(chunk['error'])
+            if chunk.get('usage') is not None:
+                usage = chunk['usage']
+            for choice in chunk.get('choices',[]):
+                tokens = choice.get('token_ids') or []
+                if tokens:
+                    events.append({'elapsed_seconds':time.perf_counter()-start,'token_ids':tokens})
+                delta = choice.get('delta',{})
+                content.append(delta.get('content') or '')
+                reasoning.append(delta.get('reasoning') or delta.get('reasoning_content') or '')
+    wall = time.perf_counter()-start
+    count = sum(len(event['token_ids']) for event in events)
+    assert usage and count == usage['completion_tokens'], (count,usage)
+    assert len(events) >= 2, 'At least two token-bearing SSE events are required.'
+    first_count = len(events[0]['token_ids'])
+    seconds = events[-1]['elapsed_seconds']-events[0]['elapsed_seconds']
+    return {'request':body,'usage':usage,'content':''.join(content),'reasoning':''.join(reasoning),
+        'wall_seconds':wall,'time_to_first_token_seconds':events[0]['elapsed_seconds'],
+        'decode_tokens':count-first_count,'decode_seconds':seconds,
+        'decode_tokens_per_second':(count-first_count)/seconds,
+        'end_to_end_tokens_per_second':count/wall,'first_event_token_count':first_count,
+        'token_events':events}
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument('--base-url',default='http://127.0.0.1:8000/v1')
+    parser.add_argument('--model',default='Naive-N0.5-Flash-NVFP4')
+    parser.add_argument('--tokens',type=int,default=256)
+    parser.add_argument('--cases',type=int,default=5)
+    parser.add_argument('--output',type=Path,required=True)
+    args = parser.parse_args()
+    request(args.base_url,args.model,'Explain why the sky is blue in several sentences.',32)
+    results=[]
+    for prompt in PROMPTS[:args.cases]:
+        result=request(args.base_url,args.model,prompt,args.tokens)
+        assert result['usage']['completion_tokens'] >= min(128,args.tokens)
+        results.append(result)
+        print(json.dumps({key:result[key] for key in ('usage','time_to_first_token_seconds',
+            'decode_tokens_per_second','end_to_end_tokens_per_second')}),flush=True)
+        report={'engine':'vLLM','measurement':'Client wall time between token-bearing SSE events, using exact token IDs; first event excluded from decode rate',
+            'single_active_request':True,'cases':results,'complete':len(results)==args.cases,
+            'target_met':len(results)==args.cases and all(r['decode_tokens_per_second']>=50 for r in results)}
+        args.output.parent.mkdir(parents=True,exist_ok=True)
+        args.output.write_text(json.dumps(report,indent=2)+'\n')
+
+
+if __name__ == '__main__':
+    main()

runtime/scripts/decode_kernels.py

@@ -0,0 +1,169 @@
+"""Small fused decode kernels, preserving the reference BF16 rounding points."""
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _rmsnorm(x, weight, output, H: tl.constexpr, EPS: tl.constexpr, BLOCK: tl.constexpr):
+    row = tl.program_id(0)
+    cols = tl.arange(0, BLOCK)
+    data = tl.load(x + row * H + cols, cols < H, other=0).to(tl.float32)
+    scale = tl.rsqrt(tl.sum(data * data, 0) / H + EPS)
+    # LlamaRMSNorm rounds normalized activations before applying its weight.
+    normalized = (data * scale).to(tl.bfloat16).to(tl.float32)
+    w = tl.load(weight + cols, cols < H, other=0).to(tl.float32)
+    tl.store(output + row * H + cols, normalized * w, cols < H)
+
+
+def rmsnorm(states, weight, eps):
+    assert states.is_contiguous() and states.dtype == torch.bfloat16
+    width = states.shape[-1]
+    result = torch.empty_like(states)
+    _rmsnorm[(states.numel() // width,)](states, weight, result, width, eps,
+        triton.next_power_of_2(width), num_warps=4, enable_fp_fusion=False)
+    return result
+
+
+@triton.jit
+def _rotary_pair(Q, K, COS, SIN, POS, OQ, OK, HQ: tl.constexpr, D: tl.constexpr,
+                 R: tl.constexpr, BLOCK: tl.constexpr):
+    head = tl.program_id(0)
+    dims = tl.arange(0, BLOCK)
+    position = tl.load(POS)
+    source = tl.where(head < HQ, Q + head * D, K + (head - HQ) * D)
+    destination = tl.where(head < HQ, OQ + head * D, OK + (head - HQ) * D)
+    values = tl.load(source + dims, dims < D, other=0).to(tl.float32)
+    partner = (dims + R // 2) % R
+    other = tl.load(source + partner, dims < R, other=0).to(tl.float32)
+    other = tl.where(dims < R // 2, -other, other)
+    cos = tl.load(COS + position * R + dims, dims < R, other=1).to(tl.float32)
+    sin = tl.load(SIN + position * R + dims, dims < R, other=0).to(tl.float32)
+    rotated = (values * cos).to(tl.bfloat16).to(tl.float32) + (other * sin).to(tl.bfloat16).to(tl.float32)
+    tl.store(destination + dims, tl.where(dims < R, rotated, values), dims < D)
+
+
+def rotary_pair(query, key, position, cos, sin):
+    dimension = query.shape[-1]
+    assert query.shape[0] == query.shape[2] == 1 and query.stride(1) == dimension
+    oq = torch.empty_like(query, memory_format=torch.contiguous_format)
+    if key is None:
+        key, ok, key_heads = query, oq, 0
+    else:
+        assert key.shape[-1] == dimension and key.stride(1) == dimension
+        ok, key_heads = torch.empty_like(key, memory_format=torch.contiguous_format), key.shape[1]
+    _rotary_pair[(query.shape[1] + key_heads,)](query, key, cos, sin, position, oq, ok,
+        query.shape[1], dimension, cos.shape[-1], triton.next_power_of_2(dimension),
+        num_warps=4, enable_fp_fusion=False)
+    return oq, ok if key_heads else None
+
+
+@triton.jit
+def _small_attention(Q, K, V, MASK, SINK, OUT,
+        L: tl.constexpr, D: tl.constexpr, DV: tl.constexpr, GROUP: tl.constexpr,
+        KS: tl.constexpr, VS: tl.constexpr, SCALE: tl.constexpr, VALUE_SCALE: tl.constexpr,
+        HAS_SINK: tl.constexpr, BN: tl.constexpr, BD: tl.constexpr, BV: tl.constexpr):
+    head = tl.program_id(0)
+    kv = head // GROUP
+    rows, dims = tl.arange(0, BN), tl.arange(0, BD)
+    q = tl.load(Q + head * D + dims, dims < D, other=0).to(tl.float32)
+    k = tl.load(K + kv * KS + rows[:, None] * D + dims[None, :],
+                (rows[:, None] < L) & (dims[None, :] < D), other=0).to(tl.float32)
+    logits = (tl.sum(k * q[None, :], 1).to(tl.bfloat16).to(tl.float32) * SCALE).to(tl.bfloat16).to(tl.float32)
+    valid = tl.load(MASK + rows, rows < L, other=False)
+    logits = tl.where(valid & (rows < L), logits, -float('inf'))
+    maximum = tl.max(logits, 0)
+    sink = -float('inf')
+    if HAS_SINK:
+        sink = tl.load(SINK + head).to(tl.float32)
+        maximum = tl.maximum(maximum, sink)
+    exps = tl.exp(logits - maximum)
+    denominator = tl.sum(exps, 0)
+    if HAS_SINK:
+        denominator += tl.exp(sink - maximum)
+    probabilities = (exps / denominator).to(tl.bfloat16).to(tl.float32)
+    vd = tl.arange(0, BV)
+    values = tl.load(V + kv * VS + rows[:, None] * DV + vd[None, :],
+                    (rows[:, None] < L) & (vd[None, :] < DV), other=0).to(tl.float32)
+    values = (values * VALUE_SCALE).to(tl.bfloat16).to(tl.float32)
+    result = tl.sum(probabilities[:, None] * values, 0)
+    tl.store(OUT + head * DV + vd, result, vd < DV)
+
+
+@triton.jit
+def _attention_logits(Q, K, MASK, LOGITS, L: tl.constexpr, D: tl.constexpr,
+        GROUP: tl.constexpr, KS: tl.constexpr, SCALE: tl.constexpr,
+        BN: tl.constexpr, BD: tl.constexpr):
+    head, block = tl.program_id(0), tl.program_id(1)
+    rows, dims = block * BN + tl.arange(0, BN), tl.arange(0, BD)
+    q = tl.load(Q + head * D + dims, dims < D, other=0).to(tl.float32)
+    k = tl.load(K + (head // GROUP) * KS + rows[:, None] * D + dims[None, :],
+                (rows[:, None] < L) & (dims[None, :] < D), other=0).to(tl.float32)
+    logits = (tl.sum(k * q[None, :], 1).to(tl.bfloat16).to(tl.float32) * SCALE).to(tl.bfloat16)
+    valid = tl.load(MASK + rows, rows < L, other=False)
+    tl.store(LOGITS + head * L + rows, tl.where(valid, logits, -float('inf')), rows < L)
+
+
+@triton.jit
+def _attention_probabilities(LOGITS, SINK, PROBS, L: tl.constexpr, HAS_SINK: tl.constexpr, BN: tl.constexpr):
+    head = tl.program_id(0)
+    rows = tl.arange(0, BN)
+    logits = tl.load(LOGITS + head * L + rows, rows < L, other=-float('inf')).to(tl.float32)
+    maximum = tl.max(logits, 0)
+    sink = -float('inf')
+    if HAS_SINK:
+        sink = tl.load(SINK + head).to(tl.float32)
+        maximum = tl.maximum(maximum, sink)
+    exps = tl.exp(logits - maximum)
+    denominator = tl.sum(exps, 0)
+    if HAS_SINK:
+        denominator += tl.exp(sink - maximum)
+    tl.store(PROBS + head * L + rows, exps / denominator, rows < L)
+
+
+@triton.jit
+def _attention_values(PROBS, V, OUT, L: tl.constexpr, DV: tl.constexpr,
+        GROUP: tl.constexpr, VS: tl.constexpr, VALUE_SCALE: tl.constexpr,
+        BN: tl.constexpr, BV: tl.constexpr):
+    head = tl.program_id(0)
+    dims = tl.arange(0, BV)
+    result = tl.full((BV,), 0, tl.float32)
+    for block in range(tl.cdiv(L, BN)):
+        rows = block * BN + tl.arange(0, BN)
+        probabilities = tl.load(PROBS + head * L + rows, rows < L, other=0).to(tl.float32)
+        values = tl.load(V + (head // GROUP) * VS + rows[:, None] * DV + dims[None, :],
+                        (rows[:, None] < L) & (dims[None, :] < DV), other=0).to(tl.float32)
+        values = (values * VALUE_SCALE).to(tl.bfloat16).to(tl.float32)
+        result += tl.sum(probabilities[:, None] * values, 0)
+    tl.store(OUT + head * DV + dims, result, dims < DV)
+
+
+def decode_attention(query, keys, values, allowed, sink, value_scale):
+    heads, dimension = query.shape[1], query.shape[-1]
+    length, value_dimension = keys.shape[-2], values.shape[-1]
+    assert query.shape[0] == query.shape[2] == 1
+    assert keys.stride(-1) == values.stride(-1) == 1
+    assert keys.stride(-2) == dimension and values.stride(-2) == value_dimension
+    result = torch.empty((1, heads, 1, value_dimension), device=query.device, dtype=query.dtype)
+    sink_pointer = sink if sink is not None else result
+    group = heads // keys.shape[1]
+    scale = dimension ** -.5
+    value_scale = 1. if value_scale is None else value_scale
+    if length <= 128:
+        _small_attention[(heads,)](query, keys, values, allowed, sink_pointer, result,
+            length, dimension, value_dimension, group, keys.stride(1), values.stride(1),
+            scale, value_scale, sink is not None, triton.next_power_of_2(length),
+            triton.next_power_of_2(dimension), triton.next_power_of_2(value_dimension),
+            num_warps=8, enable_fp_fusion=False)
+    else:
+        logits = torch.empty((heads, length), dtype=torch.bfloat16, device=query.device)
+        probabilities = torch.empty_like(logits)
+        _attention_logits[(heads, triton.cdiv(length, 64))](query, keys, allowed, logits,
+            length, dimension, group, keys.stride(1), scale, 64, triton.next_power_of_2(dimension),
+            num_warps=4, enable_fp_fusion=False)
+        _attention_probabilities[(heads,)](logits, sink_pointer, probabilities, length,
+            sink is not None, triton.next_power_of_2(length), num_warps=4)
+        _attention_values[(heads,)](probabilities, values, result, length, value_dimension,
+            group, values.stride(1), value_scale, 64, triton.next_power_of_2(value_dimension),
+            num_warps=4, enable_fp_fusion=False)
+    return result

runtime/scripts/fused_experts.py

@@ -0,0 +1,50 @@
+"""Fused SM120 NVFP4 experts; original checkpoint weights and calibrated scales."""
+import torch
+from torch import nn
+from flashinfer import nvfp4_block_scale_interleave
+from flashinfer.cute_dsl.utils import convert_sf_to_mma_layout
+from flashinfer.fused_moe import B12xMoEWrapper
+
+
+class FusedNVFP4Experts(nn.Module):
+    def __init__(self, reader, prefix, num_experts, hidden, intermediate, device, top_k=8, max_tokens=256):
+        super().__init__()
+        self.num_experts, self.hidden, self.intermediate = num_experts, hidden, intermediate
+        w1 = torch.empty((num_experts, 2 * intermediate, hidden // 2), dtype=torch.uint8, device=device)
+        w2 = torch.empty((num_experts, hidden, intermediate // 2), dtype=torch.uint8, device=device)
+        sf1 = torch.empty((num_experts, 2 * intermediate, hidden // 16), dtype=torch.uint8, device=device)
+        sf2 = torch.empty((num_experts, hidden, intermediate // 16), dtype=torch.uint8, device=device)
+        globals1, globals2, inputs1, inputs2 = [], [], [], []
+        for expert in range(num_experts):
+            name = f'{prefix}{expert}.'
+            values = {p + '.' + k: reader.get_tensor(name + p + '.' + k)
+                for p in ('gate_proj', 'up_proj', 'down_proj')
+                for k in ('weight', 'weight_scale', 'weight_scale_2', 'input_scale')}
+            for kind in ('weight_scale_2', 'input_scale'):
+                assert torch.equal(values['gate_proj.' + kind], values['up_proj.' + kind])
+            # b12x uses [up, gate]; the reference linear uses [gate, up].
+            w1[expert, :intermediate].copy_(values['up_proj.weight'])
+            w1[expert, intermediate:].copy_(values['gate_proj.weight'])
+            w2[expert].copy_(values['down_proj.weight'])
+            logical = torch.cat([values['up_proj.weight_scale'], values['gate_proj.weight_scale']])
+            sf1[expert].copy_(nvfp4_block_scale_interleave(logical.view(torch.uint8)).reshape_as(sf1[expert]))
+            sf2[expert].copy_(nvfp4_block_scale_interleave(values['down_proj.weight_scale'].view(torch.uint8)).reshape_as(sf2[expert]))
+            globals1.append(values['gate_proj.weight_scale_2'].reshape(()))
+            globals2.append(values['down_proj.weight_scale_2'].reshape(()))
+            inputs1.append(values['gate_proj.input_scale'].reshape(()))
+            inputs2.append(values['down_proj.input_scale'].reshape(()))
+        for key, value in dict(w1=w1, w2=w2,
+                sf1=convert_sf_to_mma_layout(sf1.view(torch.float8_e4m3fn), m=2*intermediate, k=hidden, num_groups=num_experts),
+                sf2=convert_sf_to_mma_layout(sf2.view(torch.float8_e4m3fn), m=hidden, k=intermediate, num_groups=num_experts),
+                alpha1=torch.stack(globals1).float(),
+                alpha2=torch.stack(globals2).float() * torch.stack(inputs2).float(),
+                input1=torch.stack(inputs1).float(), input2=torch.stack(inputs2).float()).items():
+            self.register_buffer(key, value)
+        self.wrapper = B12xMoEWrapper(num_experts, top_k, hidden, intermediate,
+            use_cuda_graph=True, max_num_tokens=max_tokens, device=str(device), quant_mode='nvfp4')
+
+    def forward(self, hidden_states, top_k_index, top_k_weights):
+        with torch.cuda.device(hidden_states.device):
+            return self.wrapper.run(hidden_states, self.w1, self.sf1, self.w2, self.sf2,
+                top_k_index, top_k_weights, w1_alpha=self.alpha1, w2_alpha=self.alpha2,
+                fc2_input_scale=self.input2, input_global_scale=self.input1)

runtime/scripts/nvfp4_runtime.py

@@ -0,0 +1,154 @@
+"""Reference local runtime: upstream Naive attention and native Blackwell expert GEMMs.
+
+Loads ModelOpt's standard packed checkpoint without changing the source architecture.
+Expert dispatch is eager PyTorch; the matrix multiplications execute as NVFP4 W4A4.
+"""
+import os
+from pathlib import Path
+
+ROOT = Path(__file__).resolve().parents[1]
+os.environ.setdefault("HF_MODULES_CACHE", str(ROOT / "artifacts/hf-modules"))
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+from accelerate import init_empty_weights, dispatch_model
+from accelerate.utils import set_module_tensor_to_device
+from flashinfer import nvfp4_quantize, nvfp4_block_scale_interleave, mm_fp4, SfLayout
+from safetensors import safe_open
+from transformers import AutoConfig, AutoModelForCausalLM
+from safe_device_transfer import install as install_safe_transfers
+
+class NVFP4Linear(nn.Module):
+    def __init__(self, weight, weight_scale, weight_scale_2, input_scale, backend="cutlass"):
+        super().__init__()
+        assert weight.dtype == torch.uint8
+        assert weight_scale.dtype == torch.float8_e4m3fn
+        assert weight_scale.shape == (weight.shape[0], weight.shape[1] // 8)
+        assert torch.isfinite(weight_scale.float()).all()
+        assert weight_scale_2.numel() == input_scale.numel() == 1
+        assert weight_scale_2.item() > 0 and input_scale.item() > 0
+        self.register_buffer("weight", weight.contiguous())
+        self.register_buffer("scales", nvfp4_block_scale_interleave(weight_scale.contiguous().view(torch.uint8)))
+        self.register_buffer("input_scale_inv", input_scale.float().reciprocal().reshape(1))
+        self.register_buffer("alpha", (weight_scale_2.float() * input_scale.float()).reshape(1))
+        self.backend = backend
+
+    def forward(self, states):
+        shape = states.shape
+        states = states.reshape(-1, shape[-1]).contiguous()
+        # FlashInfer's quantizer calls .cuda() internally, so its current device
+        # must follow the layer's placement in a multi-GPU model.
+        with torch.cuda.device(states.device):
+            packed, scales = nvfp4_quantize(states, self.input_scale_inv,
+                                           sfLayout=SfLayout.layout_128x4, do_shuffle=False)
+            output = mm_fp4(packed, self.weight.T, scales, self.scales,
+                            self.alpha, out_dtype=states.dtype, backend=self.backend)
+        return output.reshape(*shape[:-1], self.weight.shape[0])
+
+
+class NVFP4Expert(nn.Module):
+    def __init__(self, tensors, backend):
+        super().__init__()
+        for suffix in ("weight_scale_2", "input_scale"):
+            if not torch.equal(tensors[f"gate_proj.{suffix}"], tensors[f"up_proj.{suffix}"]):
+                raise ValueError(f"Exported gate/up {suffix} must match before fusion")
+        self.gate_up = NVFP4Linear(
+            torch.cat([tensors["gate_proj.weight"], tensors["up_proj.weight"]]),
+            torch.cat([tensors["gate_proj.weight_scale"], tensors["up_proj.weight_scale"]]),
+            tensors["gate_proj.weight_scale_2"], tensors["gate_proj.input_scale"], backend)
+        self.down = NVFP4Linear(tensors["down_proj.weight"], tensors["down_proj.weight_scale"],
+                               tensors["down_proj.weight_scale_2"], tensors["down_proj.input_scale"], backend)
+
+    def forward(self, states):
+        gate, up = self.gate_up(states).chunk(2, dim=-1)
+        return self.down(F.silu(gate) * up)
+
+
+class NVFP4Experts(nn.Module):
+    def __init__(self, experts):
+        super().__init__()
+        self.experts = nn.ModuleList(experts)
+
+    def forward(self, hidden_states, top_k_index, top_k_weights):
+        output = torch.zeros_like(hidden_states)
+        for expert_id in top_k_index.unique().tolist():
+            tokens, slots = torch.where(top_k_index == expert_id)
+            current = self.experts[expert_id](hidden_states[tokens])
+            weighted = current * top_k_weights[tokens, slots, None].to(current.dtype)
+            output.index_add_(0, tokens, weighted)
+        return output
+
+
+def load_model(checkpoint, devices=None, backend="cutlass"):
+    install_safe_transfers()
+    os.environ.setdefault("HF_MODULES_CACHE", str(ROOT / "artifacts/hf-modules"))
+    checkpoint = Path(checkpoint).resolve()
+    devices = list(range(torch.cuda.device_count())) if devices is None else devices
+    if not devices:
+        raise ValueError("At least one Blackwell GPU is required")
+    for device in devices:
+        assert torch.cuda.get_device_capability(device)[0] >= 10
+    config = AutoConfig.from_pretrained(checkpoint, trust_remote_code=True)
+    if config.model_type != "naive_n05_flash":
+        raise ValueError(f"Unexpected model type: {config.model_type}")
+    if hasattr(config, "quantization_config"):
+        del config.quantization_config
+    config.use_cache = True
+    with init_empty_weights():
+        model = AutoModelForCausalLM.from_config(config, trust_remote_code=True, dtype=torch.bfloat16)
+    device_map = {"model.embed_tokens": devices[0], "model.norm": devices[-1], "lm_head": devices[-1]}
+    for i in range(config.num_hidden_layers):
+        device_map[f"model.layers.{i}"] = devices[min(i * len(devices) // config.num_hidden_layers, len(devices) - 1)]
+    import json
+    index = json.loads((checkpoint / "model.safetensors.index.json").read_text())["weight_map"]
+    loaded = set()
+    # Each ModelOpt layer export is a self-contained shard. Expert tensors are
+    # read by exact name so a missing scale fails instead of silently initializing.
+    for layer_id, layer in enumerate(model.model.layers):
+        prefix = f"model.layers.{layer_id}."
+        gpu = device_map[prefix[:-1]]
+        names = [name for name in index if name.startswith(prefix)]
+        shards = {index[name] for name in names}
+        if len(shards) != 1:
+            raise ValueError("This loader requires one self-contained shard per decoder layer")
+        with torch.cuda.device(gpu), safe_open(checkpoint / next(iter(shards)), framework="pt", device=gpu) as reader:
+            if config.moe_layer_freq[layer_id]:
+                if os.environ.get('NAIVE_EXPERT_BACKEND', 'reference') == 'fused':
+                    from fused_experts import FusedNVFP4Experts
+                    layer.mlp.experts = FusedNVFP4Experts(reader, f'{prefix}mlp.experts.',
+                        config.n_routed_experts, config.hidden_size, config.moe_intermediate_size,
+                        torch.device(f'cuda:{gpu}'), top_k=config.num_experts_per_tok,
+                        max_tokens=int(os.environ.get('PREFILL_CHUNK', '256')))
+                    loaded.update(name for name in names if name.startswith(f'{prefix}mlp.experts.'))
+                else:
+                    experts = []
+                    for expert_id in range(config.n_routed_experts):
+                        ep = f"{prefix}mlp.experts.{expert_id}."
+                        tensors = {projection + "." + kind: reader.get_tensor(ep + projection + "." + kind)
+                                   for projection in ("gate_proj", "up_proj", "down_proj")
+                                   for kind in ("weight", "weight_scale", "weight_scale_2", "input_scale")}
+                        experts.append(NVFP4Expert(tensors, backend))
+                        loaded.update(ep + name for name in tensors)
+                    layer.mlp.experts = NVFP4Experts(experts)
+            for name in names:
+                if name in loaded:
+                    continue
+                tensor = reader.get_tensor(name)
+                set_module_tensor_to_device(model, name, gpu, value=tensor, dtype=tensor.dtype)
+                loaded.add(name)
+        print(f"Loaded layer {layer_id + 1}/{config.num_hidden_layers} on cuda:{gpu}", flush=True)
+    remaining = set(index) - loaded
+    for name in sorted(remaining):
+        module_name = name.rsplit(".", 1)[0]
+        if module_name not in device_map:
+            raise ValueError(f"Unrecognized tensor outside decoder layers: {name}")
+        gpu = device_map[module_name]
+        with safe_open(checkpoint / index[name], framework="pt", device=gpu) as reader:
+            tensor = reader.get_tensor(name)
+            set_module_tensor_to_device(model, name, gpu, value=tensor, dtype=tensor.dtype)
+            loaded.add(name)
+    assert loaded == set(index)
+    assert not any(p.is_meta for p in model.parameters())
+    model = dispatch_model(model, device_map=device_map, force_hooks=True)
+    return model.eval().requires_grad_(False)

runtime/scripts/preflight.py

@@ -0,0 +1,63 @@
+"""Check the tested four-GPU environment and shard headers without loading weights."""
+import importlib.metadata
+import json
+import os
+import struct
+from pathlib import Path
+
+import torch
+
+from flashinfer.fused_moe import B12xMoEWrapper
+from transformers import AutoConfig, AutoTokenizer
+
+
+def main():
+    versions = {name: importlib.metadata.version(name) for name in (
+        'vllm', 'torch', 'transformers', 'accelerate', 'tokenizers', 'huggingface-hub',
+        'safetensors', 'flashinfer-python', 'nvidia-cutlass-dsl', 'triton')}
+    expected = {'vllm': '0.28.0', 'torch': '2.13.0+cu130', 'transformers': '5.17.0', 'accelerate': '1.15.0',
+                'tokenizers': '0.23.2', 'huggingface-hub': '1.33.0', 'safetensors': '0.8.0',
+                'flashinfer-python': '0.7.0', 'nvidia-cutlass-dsl': '4.6.2', 'triton': '3.7.1'}
+    assert versions == expected, {'expected': expected, 'actual': versions}
+    assert torch.cuda.device_count() == 4, 'This deployment recipe requires four visible GPUs.'
+    gpus = []
+    for index in range(4):
+        props = torch.cuda.get_device_properties(index)
+        assert (props.major, props.minor) == (12, 0), 'The fused B12x recipe targets SM120 Blackwell.'
+        assert props.total_memory >= 90 * 1024**3, 'The tested configuration uses four 96 GB GPUs.'
+        gpus.append({'index': index, 'name': props.name, 'memory_bytes': props.total_memory})
+    cache = Path('/cache')
+    cache.mkdir(parents=True, exist_ok=True)
+    probe = cache / '.write-check'
+    probe.write_text('ok\n')
+    probe.unlink()
+    checkpoint = Path(os.environ['MODEL_PATH'])
+    config = AutoConfig.from_pretrained(checkpoint, trust_remote_code=True)
+    assert config.model_type == 'naive_n05_flash'
+    tokenizer = AutoTokenizer.from_pretrained(checkpoint, trust_remote_code=True)
+    assert tokenizer.encode('Hello')
+    weight_map = json.loads((checkpoint / 'model.safetensors.index.json').read_text())['weight_map']
+    shards = sorted(set(weight_map.values()))
+    assert len(shards) == 49
+    discovered = {}
+    for name in shards:
+        shard = checkpoint / name
+        with shard.open('rb') as stream:
+            header_size = struct.unpack('<Q', stream.read(8))[0]
+            assert header_size < 100_000_000
+            header = json.loads(stream.read(header_size))
+        payload_end = 0
+        for key, metadata in header.items():
+            if key == '__metadata__':
+                continue
+            assert key not in discovered
+            discovered[key] = name
+            payload_end = max(payload_end, metadata['data_offsets'][1])
+        assert shard.stat().st_size == 8 + header_size + payload_end, name
+    assert discovered == weight_map, 'Shard headers and weight index differ.'
+    print(json.dumps({'passed': True, 'versions': versions, 'gpus': gpus,
+                      'shards': len(shards), 'tensors': len(discovered)}, indent=2))
+
+
+if __name__ == '__main__':
+    main()

runtime/scripts/safe_device_transfer.py

@@ -0,0 +1,30 @@
+"""Stage Accelerate's cross-GPU copies through host RAM on this workstation.
+
+Direct CUDA peer copies report success but fail equality checks here. The existing
+model-manager profiles also disable NCCL P2P. This adapter changes no driver or
+global machine settings; it affects only this Python process.
+"""
+import torch
+
+
+def install():
+    import accelerate.hooks
+    import accelerate.utils
+    import accelerate.utils.operations as operations
+    if getattr(operations.send_to_device, "_host_staged", False):
+        return
+    original = operations.send_to_device
+
+    def send_to_device(tensor, device, non_blocking=False, skip_keys=None):
+        if isinstance(tensor, torch.Tensor) and tensor.is_cuda:
+            destination = torch.device(f"cuda:{device}" if isinstance(device, int) else device)
+            if destination.type == "cuda":
+                index = torch.cuda.current_device() if destination.index is None else destination.index
+                if tensor.device.index != index:
+                    return tensor.to("cpu", non_blocking=False).to(destination, non_blocking=False)
+        return original(tensor, device, non_blocking=non_blocking, skip_keys=skip_keys)
+
+    send_to_device._host_staged = True
+    operations.send_to_device = send_to_device
+    accelerate.hooks.send_to_device = send_to_device
+    accelerate.utils.send_to_device = send_to_device

runtime/scripts/validate_api.py

@@ -0,0 +1,73 @@
+"""Record actual full-model answers, tool calls, longer context and throughput."""
+import json
+import os
+import re
+import subprocess
+import time
+import urllib.request
+from pathlib import Path
+
+ROOT = Path(__file__).resolve().parents[2]
+BASE = "http://127.0.0.1:8000"
+MODEL = "Naive-N0.5-Flash-NVFP4"
+REPORT_PATH = Path(os.environ.get('NAIVE_VALIDATION_FILE', 'native-vllm-api-validation.json'))
+
+
+def complete(prompt, **extra):
+    body = {"model": MODEL, "messages": [{"role": "user", "content": prompt}],
+            "temperature": 0, "max_tokens": 128,
+            "chat_template_kwargs": {"enable_thinking": False}, **extra}
+    request = urllib.request.Request(BASE + "/v1/chat/completions",
+        data=json.dumps(body).encode(), headers={"Content-Type": "application/json"})
+    print(f'Request: {prompt[-110:]}', flush=True)
+    start = time.monotonic()
+    with urllib.request.urlopen(request, timeout=1200) as response:
+        result = json.load(response)
+    elapsed = time.monotonic() - start
+    result["elapsed_seconds"] = elapsed
+    result["completion_tokens_per_second_including_prefill"] = result["usage"]["completion_tokens"] / elapsed
+    print(f'Completed in {elapsed:.1f}s: {result["usage"]}; {result["choices"][0]["message"]}', flush=True)
+    return result
+
+
+def main():
+    results = {}
+    try:
+        with urllib.request.urlopen(BASE + "/health") as response:
+            assert response.status == 200
+        with urllib.request.urlopen(BASE + "/v1/models") as response:
+            assert MODEL in {entry["id"] for entry in json.load(response)["data"]}
+        results["arithmetic"] = complete("What is 2 + 2? Reply with only the number.", max_tokens=32)
+        assert re.search(r"\b4\b", results["arithmetic"]["choices"][0]["message"].get("content") or "")
+        results["geography"] = complete("What is the capital of France? Reply with only its name.", max_tokens=32)
+        assert "paris" in (results["geography"]["choices"][0]["message"].get("content") or "").lower()
+        results["coding"] = complete("Write a Python function named add that returns the sum of two arguments. Return only the code.", max_tokens=128)
+        assert "def add" in (results["coding"]["choices"][0]["message"].get("content") or "")
+        tools = [{"type": "function", "function": {"name": "get_weather", "description": "Get weather for a city",
+                  "parameters": {"type": "object", "properties": {"city": {"type": "string"}}, "required": ["city"]}}}]
+        results["tools"] = complete("Call the get_weather tool for Warsaw. Do not answer from memory.", tools=tools, max_tokens=192)
+        calls = results["tools"]["choices"][0]["message"].get("tool_calls", [])
+        assert calls and calls[0]["function"]["name"] == "get_weather"
+        assert "warsaw" in calls[0]["function"]["arguments"].lower()
+        results["reasoning"] = complete("Calculate 6 times 7. Think briefly, then give the answer.",
+            chat_template_kwargs={"enable_thinking": True}, reasoning_effort="low", max_tokens=256)
+        reasoning_message = results["reasoning"]["choices"][0]["message"]
+        assert reasoning_message.get("reasoning")
+        assert re.search(r"\b42\b", reasoning_message.get("content") or "")
+        filler = "The following paragraph is background material about testing a language model. It contains no instructions.\n"
+        results["longer_context"] = complete(filler * 220 + "\nFinal question: what is 2 + 2? Answer only the number.", max_tokens=32)
+        assert results["longer_context"]["usage"]["prompt_tokens"] > 2048
+        assert re.search(r"\b4\b", results["longer_context"]["choices"][0]["message"].get("content") or "")
+        results["gpu_memory"] = subprocess.check_output(["nvidia-smi", "--query-gpu=index,memory.used,memory.total", "--format=csv"], text=True)
+        report = {"passed": True, "engine": "vLLM 0.28.0", "model": MODEL, "checks": results,
+                  "scope": "Basic generation, coding, tool use and sparse-context smoke checks; not an accuracy benchmark"}
+    except Exception as error:
+        report = {"passed": False, "model": MODEL, "error": repr(error), "checks": results}
+        REPORT_PATH.write_text(json.dumps(report, indent=2) + "\n")
+        raise
+    REPORT_PATH.write_text(json.dumps(report, indent=2) + "\n")
+    print("FULL_MODEL_API_VALIDATION_PASSED", flush=True)
+
+
+if __name__ == "__main__":
+    main()

runtime/sitecustomize.py

@@ -0,0 +1,7 @@
+"""Isolate FlashInfer source kernels before vLLM worker imports execute."""
+import os
+import sys
+
+if os.environ.get('NAIVE_FLASHINFER_SOURCE_ONLY') == '1':
+    sys.modules['flashinfer_jit_cache'] = None
+    sys.modules['flashinfer_cubin'] = None

runtime/vllm_naive/__init__.py

@@ -0,0 +1,11 @@
+"""Native model registration for vLLM's V1 engine."""
+
+
+def register():
+    import sys
+    # FlashInfer 0.7 kernels must not use the base image's 0.6 prebuilt binaries.
+    sys.modules['flashinfer_jit_cache'] = None
+    sys.modules['flashinfer_cubin'] = None
+    from vllm import ModelRegistry
+    ModelRegistry.register_model('NaiveN05FlashForCausalLM',
+        'vllm_naive.model:NaiveN05FlashForCausalLM')

runtime/vllm_naive/attention.py

@@ -0,0 +1,242 @@
+"""Naive GQA sparse attention using vLLM-managed paged K/V/indexer storage.
+
+Prefill and multiple active requests use eager operations; single-token decode
+supports vLLM CUDA graphs. Both retain upstream indexer FP8 rounding and stable
+sparse-selection ties.
+"""
+from dataclasses import dataclass
+
+import torch
+from torch import nn
+from vllm.distributed import tensor_model_parallel_all_reduce
+from vllm.forward_context import get_forward_context
+from vllm.model_executor.layers.attention_layer_base import AttentionLayerBase
+from vllm.v1.attention.backend import AttentionBackend, AttentionMetadataBuilder, AttentionCGSupport
+from vllm.v1.kv_cache_interface import FullAttentionSpec
+from decode_kernels import decode_attention, rotary_pair
+
+
+@dataclass
+class NaiveMetadata:
+    starts: list[int]
+    lengths: list[int]
+    block_table: torch.Tensor
+    slot_mapping: torch.Tensor
+    num_actual_tokens: int
+    single_decode: bool
+
+
+class NaiveMetadataBuilder(AttentionMetadataBuilder):
+    _cudagraph_support = AttentionCGSupport.UNIFORM_SINGLE_TOKEN_DECODE
+
+    def __init__(self, kv_cache_spec, layer_names, vllm_config, device):
+        super().__init__(kv_cache_spec, layer_names, vllm_config, device)
+
+    def build(self, common_prefix_len, common_attn_metadata, fast_build=False):
+        c = common_attn_metadata
+        single_decode = c.max_query_len == 1 and c.num_actual_tokens == 1
+        lengths = [] if single_decode else c.seq_lens.cpu().tolist()
+        return NaiveMetadata(c.query_start_loc_cpu.tolist(), lengths,
+                             c.block_table_tensor, c.slot_mapping, c.num_actual_tokens, single_decode)
+
+
+class NaiveBackend(AttentionBackend):
+    @staticmethod
+    def get_name():
+        return 'NAIVE_GQA_DSA'
+
+    @staticmethod
+    def get_impl_cls():
+        return NaivePagedAttention
+
+    @staticmethod
+    def get_builder_cls():
+        return NaiveMetadataBuilder
+
+    @staticmethod
+    def get_kv_cache_shape(num_blocks, block_size, num_kv_heads, head_size, cache_dtype_str='auto'):
+        return (num_blocks, 2, block_size, num_kv_heads, head_size)
+
+    @staticmethod
+    def get_kv_cache_stride_order(include_num_layers_dimension=False):
+        return tuple(range(6 if include_num_layers_dimension else 5))
+
+
+def round_fp8(states):
+    values = states.float()
+    scale = values.abs().amax(-1, keepdim=True).clamp_min(1e-4) / 448.
+    return (values / scale).clamp(-448, 448).to(torch.float8_e4m3fn).float() * scale
+
+
+class NaivePagedAttention(nn.Module, AttentionLayerBase):
+    supports_dcp = False
+
+    def __init__(self, source, name, vllm_config, tp_size, tp_rank):
+        super().__init__()
+        self.config = source.config
+        self.name = name
+        self.tp_size = tp_size
+        self.is_swa = source.is_swa
+        self.heads = source.heads // tp_size
+        self.kv_heads = source.kv_heads // tp_size
+        assert source.heads % tp_size == source.kv_heads % tp_size == 0
+        self.head_dim, self.v_dim = source.head_dim, source.v_dim
+        self.num_heads = self.heads
+        self.cache_heads = max(self.config.num_key_value_heads, self.config.swa_num_key_value_heads) // tp_size
+        self.cache_width = max(self.config.head_dim, self.config.swa_head_dim) + self.config.index_head_dim
+        self.block_size = vllm_config.cache_config.block_size
+        self.max_context = vllm_config.model_config.max_model_len
+        for key in ('q_proj', 'k_proj', 'v_proj', 'o_proj', 'rotary_emb', 'indexer'):
+            setattr(self, key, getattr(source, key))
+        self.attention_sink_bias = source.attention_sink_bias
+        for key in ('q_proj', 'k_proj', 'v_proj'):
+            module = getattr(self, key)
+            module.out_features //= tp_size
+            module.weight = nn.Parameter(torch.empty(module.out_features, module.in_features, device='meta', dtype=torch.bfloat16))
+        self.o_proj.in_features //= tp_size
+        self.o_proj.weight = nn.Parameter(torch.empty(self.o_proj.out_features, self.o_proj.in_features, device='meta', dtype=torch.bfloat16))
+        if self.attention_sink_bias is not None:
+            self.attention_sink_bias = nn.Parameter(torch.empty(self.heads, device='meta', dtype=torch.bfloat16), requires_grad=False)
+        self.kv_cache = torch.empty(0)
+        context = vllm_config.compilation_config.static_forward_context
+        assert name not in context
+        context[name] = self
+
+    def get_attn_backend(self):
+        return NaiveBackend
+
+    def get_kv_cache_spec(self, vllm_config):
+        self.block_size = vllm_config.cache_config.block_size
+        # Indexer values after FP8 rounding carry FP32 scales. Storing all cache
+        # elements as FP32 preserves them exactly; K/V are cast back to BF16.
+        return FullAttentionSpec(block_size=self.block_size, num_kv_heads=self.cache_heads,
+                                 head_size=self.cache_width, dtype=torch.float32)
+
+    def initialize_decode(self):
+        """Share packed projection storage and precompute fixed rotary tables."""
+        projections = (self.q_proj, self.k_proj, self.v_proj)
+        assert all(projection.bias is None for projection in projections)
+        sizes = [projection.weight.shape[0] for projection in projections]
+        packed = torch.cat([projection.weight for projection in projections])
+        self.register_buffer('qkv_weight', packed, persistent=False)
+        # Retain the original projection interfaces for prefill without keeping
+        # a second copy of the projection weights on the GPU.
+        for projection, weight in zip(projections, packed.split(sizes)):
+            projection.weight = nn.Parameter(weight, requires_grad=False)
+        self.qkv_sizes = sizes
+        positions = torch.arange(self.max_context, device=packed.device).unsqueeze(0)
+        cos, sin = self.rotary_emb(packed.new_empty((1, 1, 1)), positions)
+        self.register_buffer('decode_cos', cos.contiguous(), persistent=False)
+        self.register_buffer('decode_sin', sin.contiguous(), persistent=False)
+
+    def forward(self, states, positions):
+        from transformers.models.gpt_neox.modeling_gpt_neox import apply_rotary_pos_emb
+        length = states.shape[0]
+        if length == 1:
+            query, key, value = torch.nn.functional.linear(states, self.qkv_weight).split(self.qkv_sizes, dim=-1)
+        else:
+            query, key, value = self.q_proj(states), self.k_proj(states), self.v_proj(states)
+        query = query.view(1, length, self.heads, self.head_dim).transpose(1, 2)
+        key = key.view(1, length, self.kv_heads, self.head_dim).transpose(1, 2)
+        value = value.view(length, self.kv_heads, self.v_dim)
+        if length == 1:
+            query, key = rotary_pair(query, key, positions, self.decode_cos, self.decode_sin)
+        else:
+            embeddings = self.rotary_emb(states.unsqueeze(0), positions.unsqueeze(0))
+            query, key = apply_rotary_pos_emb(query, key, *embeddings)
+        index_query = index_key = index_weights = None
+        if self.indexer is not None:
+            c = self.config
+            index_query = self.indexer.wq(states).view(1, length, c.index_n_heads, c.index_head_dim).transpose(1, 2)
+            index_key = self.indexer.k_norm(self.indexer.wk(states)).view(1, 1, length, c.index_head_dim)
+            if length == 1:
+                index_query, _ = rotary_pair(index_query, None, positions, self.decode_cos, self.decode_sin)
+                index_key, _ = rotary_pair(index_key, None, positions, self.decode_cos, self.decode_sin)
+            else:
+                index_query, index_key = apply_rotary_pos_emb(index_query, index_key, *embeddings)
+            if c.indexer_activation_dtype == 'fp8_e4m3':
+                index_query, index_key = round_fp8(index_query), round_fp8(index_key)
+            index_weights = self.indexer.weights_proj(states) * c.index_n_heads**-.5
+        metadata = get_forward_context().attn_metadata
+        metadata = metadata.get(self.name) if metadata else None
+        if metadata is None:
+            # vLLM's memory profiling pass has no cache metadata.
+            return self.o_proj(torch.zeros(length, self.heads*self.v_dim, dtype=states.dtype, device=states.device))
+        n = metadata.num_actual_tokens
+        # vLLM uses negative slots for graph-capture dummy tokens. Page zero is
+        # its null block; real scheduled tokens always have allocated slots.
+        slots = metadata.slot_mapping[:n].long().clamp_min(0)
+        blocks, offsets = slots // self.block_size, slots % self.block_size
+        self.kv_cache[blocks, 0, offsets, :self.kv_heads, :self.head_dim] = key[0].transpose(0, 1)[:n].float()
+        self.kv_cache[blocks, 1, offsets, :self.kv_heads, :self.v_dim] = value[:n].float()
+        if index_key is not None:
+            self.kv_cache[blocks, 0, offsets, 0, -self.config.index_head_dim:] = index_key[0, 0, :n].float()
+        if metadata.single_decode:
+            return self.decode(query, positions, index_query, index_weights, metadata)
+        output = torch.zeros(length, self.heads*self.v_dim, device=states.device, dtype=states.dtype)
+        for request, (start, end) in enumerate(zip(metadata.starts, metadata.starts[1:])):
+            if start == end:
+                continue
+            total = metadata.lengths[request]
+            first = max(0, total - (end-start) - self.config.sliding_window + 1) if self.is_swa else 0
+            key_positions = torch.arange(first, total, device=states.device)
+            page = metadata.block_table[request, key_positions // self.block_size].long()
+            offset = key_positions % self.block_size
+            keys = self.kv_cache[page, 0, offset, :self.kv_heads, :self.head_dim].to(states.dtype).transpose(0, 1)
+            values = self.kv_cache[page, 1, offset, :self.kv_heads, :self.v_dim].to(states.dtype).transpose(0, 1)
+            allowed = positions[start:end, None] >= key_positions[None]
+            if self.is_swa:
+                allowed = allowed & ((positions[start:end, None] - key_positions[None]) < self.config.sliding_window)
+            else:
+                cached_index = self.kv_cache[page, 0, offset, 0, -self.config.index_head_dim:]
+                scores = (index_query[0, :, start:end].float() @ cached_index.T).relu()
+                scores = (scores * index_weights[start:end].T[:, :, None].float()).sum(0)
+                scores = scores.masked_fill(~allowed, -torch.inf)
+                selected = scores.argsort(dim=-1, descending=True, stable=True)[:, :self.config.index_top_k]
+                allowed = allowed & torch.zeros_like(scores, dtype=torch.bool).scatter(-1, selected, True)
+            keys = keys.repeat_interleave(self.heads//self.kv_heads, dim=0)
+            values = values.repeat_interleave(self.heads//self.kv_heads, dim=0)
+            if self.config.attention_value_scale is not None:
+                values = values * self.config.attention_value_scale
+            logits = (query[0, :, start:end] @ keys.transpose(-1, -2)) * self.head_dim**-.5
+            logits = logits.masked_fill(~allowed[None], -torch.inf)
+            if self.attention_sink_bias is not None:
+                sinks = self.attention_sink_bias[:, None, None].expand(-1, end-start, 1)
+                logits = torch.cat((logits, sinks), -1)
+            probabilities = logits.float().softmax(-1).nan_to_num(0)[..., :len(key_positions)].to(values.dtype)
+            output[start:end] = (probabilities @ values).transpose(0, 1).reshape(end-start, -1)
+        output = self.o_proj(output)
+        return tensor_model_parallel_all_reduce(output) if self.tp_size > 1 else output
+
+    def decode(self, query, positions, index_query, index_weights, metadata):
+        """Fixed-shape single-token path for vLLM's full decode CUDA graphs."""
+        if self.is_swa:
+            key_positions = positions[0] - self.config.sliding_window + 1 + torch.arange(
+                self.config.sliding_window, device=query.device)
+            valid = key_positions >= 0
+            key_positions = key_positions.clamp_min(0)
+        else:
+            key_positions = torch.arange(self.max_context, device=query.device)
+            valid = key_positions <= positions[0]
+            pages = metadata.block_table[0, key_positions // self.block_size].long()
+            pages = torch.where(valid, pages, 0)
+            cached_index = self.kv_cache[pages, 0, key_positions % self.block_size,
+                                         0, -self.config.index_head_dim:]
+            scores = (index_query[0].float() @ cached_index.T).relu()
+            scores = (scores * index_weights.T[:, :, None].float()).sum(0)
+            scores = scores.masked_fill(~valid[None], -torch.inf)
+            selected = scores.argsort(dim=-1, descending=True, stable=True)[0, :self.config.index_top_k]
+            key_positions = selected.sort().values
+            valid = key_positions <= positions[0]
+        pages = metadata.block_table[0, key_positions // self.block_size].long()
+        pages = torch.where(valid, pages, 0)
+        offsets = key_positions % self.block_size
+        keys = self.kv_cache[pages, 0, offsets, :self.kv_heads, :self.head_dim].to(query.dtype).transpose(0, 1).contiguous()
+        values = self.kv_cache[pages, 1, offsets, :self.kv_heads, :self.v_dim].to(query.dtype).transpose(0, 1).contiguous()
+        # Unallocated/padding entries may contain arbitrary bits in the cache.
+        keys = keys.masked_fill(~valid[None, :, None], 0)
+        values = values.masked_fill(~valid[None, :, None], 0)
+        output = decode_attention(query, keys[None], values[None], valid,
+                                  self.attention_sink_bias, self.config.attention_value_scale)
+        output = self.o_proj(output.transpose(1, 2).reshape(1, -1))
+        return tensor_model_parallel_all_reduce(output) if self.tp_size > 1 else output

runtime/vllm_naive/model.py

@@ -0,0 +1,189 @@
+"""Naive-N0.5-Flash NVFP4 model executed by the native vLLM V1 model runner."""
+import copy
+import json
+from pathlib import Path
+
+import torch
+from torch import nn
+from accelerate import init_empty_weights
+from accelerate.utils import set_module_tensor_to_device
+from safetensors import safe_open
+from transformers import AutoModelForCausalLM
+from vllm.distributed import get_tensor_model_parallel_rank, get_tensor_model_parallel_world_size, tensor_model_parallel_all_reduce
+from vllm_naive.attention import NaivePagedAttention
+
+from fused_experts import FusedNVFP4Experts
+from nvfp4_runtime import NVFP4Expert, NVFP4Experts
+from decode_kernels import rmsnorm
+from flashinfer.fused_moe.cute_dsl.blackwell_sm12x import moe_dispatch
+
+# FlashInfer 0.7's direct-micro path returned all-zero output for this model's
+# calibrated NVFP4 scales at TP4 (intermediate=512, one token). Keep the regular
+# MMA micro/static/dynamic paths. Covered by validate_expert_sharding.py.
+moe_dispatch._DIRECT_MICRO_MAX_N = 0
+
+
+class ShardedReader:
+    def __init__(self, reader, rank, size, device):
+        self.reader, self.rank, self.size, self.device = reader, rank, size, device
+
+    def get_tensor(self, name):
+        tensor = self.reader.get_tensor(name)
+        if name.endswith(('.weight', '.weight_scale')):
+            dim = 1 if '.down_proj.' in name else 0
+            assert tensor.shape[dim] % self.size == 0
+            tensor = tensor.chunk(self.size, dim=dim)[self.rank].contiguous()
+        return tensor.to(self.device)
+
+
+class ReduceExperts(nn.Module):
+    def __init__(self, experts, tp_size):
+        super().__init__()
+        self.experts, self.tp_size = experts, tp_size
+
+    def forward(self, hidden_states, top_k_index, top_k_weights):
+        result = self.experts(hidden_states, top_k_index, top_k_weights)
+        return tensor_model_parallel_all_reduce(result) if self.tp_size > 1 else result
+
+
+class StaticReferenceExperts(NVFP4Experts):
+    """Small-fixture reference kernels with no data-dependent Python dispatch."""
+
+    def forward(self, hidden_states, top_k_index, top_k_weights):
+        output = torch.zeros_like(hidden_states)
+        for expert_id, expert in enumerate(self.experts):
+            weight = (top_k_weights * (top_k_index == expert_id)).sum(-1).to(hidden_states.dtype)
+            output = output + expert(hidden_states) * weight[:, None]
+        return output
+
+
+class ShardedDense(nn.Module):
+    def __init__(self, source, tp_size):
+        super().__init__()
+        self.tp_size = tp_size
+        self.act_fn = source.act_fn
+        for name in ('gate_proj', 'up_proj', 'down_proj'):
+            module = getattr(source, name)
+            if name == 'down_proj':
+                module.in_features //= tp_size
+            else:
+                module.out_features //= tp_size
+            module.weight = nn.Parameter(torch.empty(module.out_features, module.in_features,
+                                                     dtype=torch.bfloat16, device='meta'))
+            setattr(self, name, module)
+
+    def forward(self, states):
+        result = self.down_proj(self.act_fn(self.gate_proj(states)) * self.up_proj(states))
+        return tensor_model_parallel_all_reduce(result) if self.tp_size > 1 else result
+
+
+class NaiveN05FlashForCausalLM(nn.Module):
+    packed_modules_mapping = {}
+
+    def __init__(self, *, vllm_config, prefix=''):
+        super().__init__()
+        self.vllm_config = vllm_config
+        assert vllm_config.parallel_config.pipeline_parallel_size == 1
+        self.tp_size = get_tensor_model_parallel_world_size()
+        self.tp_rank = get_tensor_model_parallel_rank()
+        self.config = copy.deepcopy(vllm_config.model_config.hf_config)
+        if hasattr(self.config, 'quantization_config'):
+            del self.config.quantization_config
+        with init_empty_weights():
+            original = AutoModelForCausalLM.from_config(self.config, trust_remote_code=True, dtype=torch.bfloat16)
+        self.model, self.lm_head = original.model, original.lm_head
+        for i, layer in enumerate(self.model.layers):
+            layer.self_attn = NaivePagedAttention(layer.self_attn, f'model.layers.{i}.self_attn',
+                                                 vllm_config, self.tp_size, self.tp_rank)
+            if not self.config.moe_layer_freq[i]:
+                layer.mlp = ShardedDense(layer.mlp, self.tp_size)
+
+    def embed_input_ids(self, input_ids):
+        return self.model.embed_tokens(input_ids)
+
+    def forward(self, input_ids, positions, intermediate_tensors=None, inputs_embeds=None):
+        states = self.embed_input_ids(input_ids) if inputs_embeds is None else inputs_embeds
+        for layer in self.model.layers:
+            norm = layer.input_layernorm
+            states = states + layer.self_attn(rmsnorm(states, norm.weight, norm.variance_epsilon), positions)
+            norm = layer.post_attention_layernorm
+            states = states + layer.mlp(rmsnorm(states, norm.weight, norm.variance_epsilon))
+        return rmsnorm(states, self.model.norm.weight, self.model.norm.variance_epsilon)
+
+    def compute_logits(self, hidden_states):
+        # Vocabulary and embedding/head are replicated across TP ranks.
+        return self.lm_head(hidden_states)
+
+    def load_weights(self, weights):
+        # Layer shards allow streaming only this rank's slices onto its GPU.
+        # The vLLM loader has resolved/downloaded the model path before this call.
+        # Initialize the framework loader's lazy download/timing iterator.
+        next(iter(weights), None)
+        checkpoint = Path(self.vllm_config.model_config.model)
+        if not checkpoint.is_dir():
+            from huggingface_hub import snapshot_download
+            checkpoint = Path(snapshot_download(self.vllm_config.model_config.model,
+                                                 revision=self.vllm_config.model_config.revision))
+        index = json.loads((checkpoint/'model.safetensors.index.json').read_text())['weight_map']
+        device = torch.device('cuda', torch.cuda.current_device())
+        loaded = set()
+        for i, layer in enumerate(self.model.layers):
+            prefix = f'model.layers.{i}.'
+            names = [name for name in index if name.startswith(prefix)]
+            shards = {index[name] for name in names}
+            assert len(shards) == 1
+            with safe_open(checkpoint / next(iter(shards)), framework='pt', device='cpu') as reader:
+                if self.config.moe_layer_freq[i]:
+                    expert_prefix = prefix + 'mlp.experts.'
+                    sliced = ShardedReader(reader, self.tp_rank, self.tp_size, device)
+                    with torch.inference_mode(False):
+                        if self.config.hidden_size >= 4096:
+                            expert_module = FusedNVFP4Experts(sliced, expert_prefix, self.config.n_routed_experts,
+                                self.config.hidden_size, self.config.moe_intermediate_size//self.tp_size,
+                                device, top_k=self.config.num_experts_per_tok,
+                                max_tokens=self.vllm_config.scheduler_config.max_num_batched_tokens)
+                        else:
+                            experts = []
+                            for expert in range(self.config.n_routed_experts):
+                                tensors = {p+'.'+kind:sliced.get_tensor(expert_prefix+str(expert)+'.'+p+'.'+kind).to(device)
+                                    for p in ('gate_proj','up_proj','down_proj')
+                                    for kind in ('weight','weight_scale','weight_scale_2','input_scale')}
+                                experts.append(NVFP4Expert(tensors, 'cutlass'))
+                            expert_module = StaticReferenceExperts(experts)
+                    layer.mlp.experts = ReduceExperts(expert_module, self.tp_size)
+                    loaded.update(name for name in names if name.startswith(expert_prefix))
+                for name in names:
+                    if name in loaded:
+                        continue
+                    tensor = reader.get_tensor(name)
+                    dim = None
+                    if any('.self_attn.'+key+'.' in name for key in ('q_proj','k_proj','v_proj')):
+                        dim = 0
+                    elif '.self_attn.o_proj.weight' in name:
+                        dim = 1
+                    elif name.endswith('.self_attn.attention_sink_bias'):
+                        dim = 0
+                    elif not self.config.moe_layer_freq[i]:
+                        if any('.mlp.'+key+'.' in name for key in ('gate_proj','up_proj')):
+                            dim = 0
+                        elif '.mlp.down_proj.weight' in name:
+                            dim = 1
+                    if dim is not None:
+                        tensor = tensor.chunk(self.tp_size, dim=dim)[self.tp_rank].contiguous()
+                    set_module_tensor_to_device(self, name, device, value=tensor, dtype=tensor.dtype)
+                    loaded.add(name)
+            print(f'NAIVE_VLLM_LOADED_LAYER {i+1}/{self.config.num_hidden_layers} TP={self.tp_rank}', flush=True)
+        for name in sorted(set(index)-loaded):
+            assert name in {'model.embed_tokens.weight','model.norm.weight','lm_head.weight'}
+            with safe_open(checkpoint/index[name], framework='pt', device='cpu') as reader:
+                tensor = reader.get_tensor(name)
+                set_module_tensor_to_device(self, name, device, value=tensor, dtype=tensor.dtype)
+            loaded.add(name)
+        assert loaded == set(index)
+        assert not any(p.is_meta for p in self.parameters())
+        self.eval().requires_grad_(False)
+        for layer in self.model.layers:
+            layer.self_attn.initialize_decode()
+        # Parameter names differ for the fused experts; all source tensors were
+        # explicitly checked above, so use vLLM's optional return convention.
+        return None

tokenizer.json

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9871eef5422b8cd120d4ecb00a08dd2f5c939d6dd46af384233a8c5078b224e2
+size 11422840

tokenizer_config.json

@@ -0,0 +1,32 @@
+{
+  "_comment": "Tokenizer copied from the open-source MiMo 2.5",
+  "add_prefix_space": false,
+  "backend": "tokenizers",
+  "bos_token": null,
+  "clean_up_tokenization_spaces": false,
+  "eos_token": "<|im_end|>",
+  "errors": "replace",
+  "extra_special_tokens": [
+    "<|im_start|>",
+    "<|im_end|>",
+    "<|object_ref_start|>",
+    "<|object_ref_end|>",
+    "<|box_start|>",
+    "<|box_end|>",
+    "<|quad_start|>",
+    "<|quad_end|>",
+    "<|vision_start|>",
+    "<|vision_end|>",
+    "<|vision_pad|>",
+    "<|image_pad|>",
+    "<|video_pad|>",
+    "<|audio_pad|>"
+  ],
+  "is_local": true,
+  "local_files_only": false,
+  "model_max_length": 1048576,
+  "pad_token": "<|endoftext|>",
+  "split_special_tokens": false,
+  "tokenizer_class": "Qwen2Tokenizer",
+  "unk_token": null
+}

validation/expert-sharding-tp4-validation.json

@@ -0,0 +1,17 @@
+{
+  "passed": true,
+  "checks": [
+    {
+      "tokens": 1,
+      "cosine": 0.9999017715454102,
+      "relative_rmse": 0.014022702351212502,
+      "max_error": 0.046875
+    },
+    {
+      "tokens": 16,
+      "cosine": 0.9999954104423523,
+      "relative_rmse": 0.003035087138414383,
+      "max_error": 0.03125
+    }
+  ]
+}

validation/export-validation.json

@@ -0,0 +1,8 @@
+{
+  "passed": true,
+  "tensor_count": 144856,
+  "packed_expert_weights": 36096,
+  "safetensor_shards": 49,
+  "tensor_bytes": 182549965184,
+  "all_source_tensors_accounted_for": true
+}

validation/native-tiny-graphs-validation.json

@@ -0,0 +1,394 @@
+{
+  "passed": true,
+  "decode_graphs": true,
+  "scope": "Real vLLM V1 engine versus quantized upstream reference, tiny architecture fixture",
+  "chunked_prefill": true,
+  "prefix_caching": true,
+  "max_num_seqs": 2,
+  "checks": [
+    {
+      "prompt_tokens": 8,
+      "expected": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "actual": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 35,
+      "expected": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "actual": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 80,
+      "expected": [
+        11,
+        11,
+        11,
+        11,
+        245,
+        11,
+        11,
+        163,
+        11,
+        38,
+        11,
+        245
+      ],
+      "actual": [
+        11,
+        11,
+        11,
+        11,
+        245,
+        11,
+        11,
+        163,
+        11,
+        38,
+        11,
+        245
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 17,
+      "expected": [
+        193,
+        197,
+        11,
+        98,
+        120,
+        46,
+        11,
+        46,
+        11,
+        98,
+        120,
+        46
+      ],
+      "actual": [
+        193,
+        197,
+        11,
+        98,
+        120,
+        46,
+        11,
+        46,
+        11,
+        98,
+        120,
+        46
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 35,
+      "expected": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "actual": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 8,
+      "expected": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "actual": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 8,
+      "expected": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "actual": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 35,
+      "expected": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "actual": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 80,
+      "expected": [
+        11,
+        11,
+        11,
+        11,
+        245,
+        11,
+        11,
+        163,
+        11,
+        38,
+        11,
+        245
+      ],
+      "actual": [
+        11,
+        11,
+        11,
+        11,
+        245,
+        11,
+        11,
+        163,
+        11,
+        38,
+        11,
+        245
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 17,
+      "expected": [
+        193,
+        197,
+        11,
+        98,
+        120,
+        46,
+        11,
+        46,
+        11,
+        98,
+        120,
+        46
+      ],
+      "actual": [
+        193,
+        197,
+        11,
+        98,
+        120,
+        46,
+        11,
+        46,
+        11,
+        98,
+        120,
+        46
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 35,
+      "expected": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "actual": [
+        193,
+        98,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        120,
+        46,
+        98,
+        120
+      ],
+      "exact_token_match": true
+    },
+    {
+      "prompt_tokens": 8,
+      "expected": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "actual": [
+        58,
+        193,
+        46,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58,
+        193,
+        46,
+        58
+      ],
+      "exact_token_match": true
+    }
+  ]
+}

validation/native-vllm-api-validation.json

@@ -0,0 +1,283 @@
+{
+  "passed": true,
+  "engine": "vLLM 0.28.0",
+  "model": "Naive-N0.5-Flash-NVFP4",
+  "checks": {
+    "arithmetic": {
+      "id": "chatcmpl-b57ff21de144f7ca",
+      "object": "chat.completion",
+      "created": 1790592276,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n4",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "stop",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 58,
+        "total_tokens": 61,
+        "completion_tokens": 3,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 0.10958343200036325,
+      "completion_tokens_per_second_including_prefill": 27.376401206252197
+    },
+    "geography": {
+      "id": "chatcmpl-9a328e34a75310ee",
+      "object": "chat.completion",
+      "created": 1790592276,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\nParis",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "stop",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 57,
+        "total_tokens": 60,
+        "completion_tokens": 3,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 0.12031267199927242,
+      "completion_tokens_per_second_including_prefill": 24.93502928783879
+    },
+    "coding": {
+      "id": "chatcmpl-a8ad4110a48efbdb",
+      "object": "chat.completion",
+      "created": 1790592276,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\ndef add(a, b):\n    return a + b",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "stop",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 63,
+        "total_tokens": 76,
+        "completion_tokens": 13,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 0.3035910079997848,
+      "completion_tokens_per_second_including_prefill": 42.82076760326583
+    },
+    "tools": {
+      "id": "chatcmpl-9b641c57e20da9a6",
+      "object": "chat.completion",
+      "created": 1790592277,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": null,
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "tool_calls": [
+              {
+                "id": "chatcmpl-tool-8eeac683e7c34ae0",
+                "type": "function",
+                "function": {
+                  "name": "get_weather",
+                  "arguments": "{\"city\": \"Warsaw\"}"
+                }
+              }
+            ],
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "tool_calls",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 284,
+        "total_tokens": 307,
+        "completion_tokens": 23,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 0.4892891299969051,
+      "completion_tokens_per_second_including_prefill": 47.00697111367563
+    },
+    "reasoning": {
+      "id": "chatcmpl-bebdc22c27f814ba",
+      "object": "chat.completion",
+      "created": 1790592277,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n42",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": "\nWe need answer simple arithmetic. Need think briefly then give answer. We can just answer 42. Need maybe no extra. Ensure final only answer.\n"
+          },
+          "logprobs": null,
+          "finish_reason": "stop",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 57,
+        "total_tokens": 95,
+        "completion_tokens": 38,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 32
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 0.7566590700007509,
+      "completion_tokens_per_second_including_prefill": 50.220768515947725
+    },
+    "longer_context": {
+      "id": "chatcmpl-ab4c6172888cf312",
+      "object": "chat.completion",
+      "created": 1790592278,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n4",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "stop",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 3800,
+        "total_tokens": 3803,
+        "completion_tokens": 3,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 1.0702180820007925,
+      "completion_tokens_per_second_including_prefill": 2.8031669904057726
+    },
+    "gpu_memory": "index, memory.used [MiB], memory.total [MiB]\n0, 86684 MiB, 97887 MiB\n1, 86684 MiB, 97887 MiB\n2, 91002 MiB, 97887 MiB\n3, 86684 MiB, 97887 MiB\n"
+  },
+  "scope": "Basic generation, coding, tool use and sparse-context smoke checks; not an accuracy benchmark"
+}

validation/native-vllm-concurrency-validation.json

@@ -0,0 +1,184 @@
+{
+  "passed": true,
+  "engine": "vLLM 0.28.0",
+  "concurrent_requests": 4,
+  "scope": "Concurrent scheduling/cache/output smoke test, not a throughput or accuracy benchmark",
+  "results": [
+    {
+      "id": "chatcmpl-8f7154ae17b91c78",
+      "object": "chat.completion",
+      "created": 1790592282,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n23\n\nTo add 20 and 3, you can think of it as combining two quantities. Start with the number 20, which represents two tens. Then, add 3 more to it. One simple way is to count forward: begin at 20 and count three steps: 21, 22, 23. Alternatively, you can break it down by place value. The number 20 has a 2 in the tens place and a 0 in the ones place. The number 3 has a 3 in the ones place. When you add them, you add the ones first: 0 + 3 = 3. Then, you add the tens: 2 + 0 = 2. Combining these gives you 2 tens and 3 ones, which is 23. This method works because you are simply combining the values in each place without any carrying",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "length",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 78,
+        "total_tokens": 270,
+        "completion_tokens": 192,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 25.521290875003615,
+      "completion_tokens_per_second_including_prefill": 7.523130430210764
+    },
+    {
+      "id": "chatcmpl-b6950169c10c14bf",
+      "object": "chat.completion",
+      "created": 1790592282,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n25\n\nTo add 20 and 5, you can think of it as combining two quantities. Start with the number 20, which represents two groups of ten. Then, add 5 more to it. One simple way to visualize this is to count forward from 20: 21, 22, 23, 24, 25. After five steps, you land on 25. Alternatively, you can break the numbers down by place value. The number 20 has a 2 in the tens place and a 0 in the ones place. The number 5 has a 5 in the ones place. When you add them, you add the ones digits: 0 + 5 = 5. The tens digit remains 2, since there is no other tens digit to add. So, you get 2 in the tens place and 5 in the",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "length",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 78,
+        "total_tokens": 270,
+        "completion_tokens": 192,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 23.55511791400204,
+      "completion_tokens_per_second_including_prefill": 8.151094836416338
+    },
+    {
+      "id": "chatcmpl-ae5bdf320cff8063",
+      "object": "chat.completion",
+      "created": 1790592282,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n27\n\nTo add 20 and 7, you can think of it as combining two quantities. Start with the number 20, which represents two groups of ten. Then, add 7 more to it. One simple way is to count up from 20: 21, 22, 23, 24, 25, 26, 27. After seven steps, you land on 27. Another method is to break 20 into 10 and 10, then add 7 to one of the tens to get 17, and finally add the other 10 to reach 27. You can also use place value: 20 has a 2 in the tens place and a 0 in the ones place. Adding 7 affects only the ones place, so 0 + 7 = 7, while the",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "length",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 78,
+        "total_tokens": 270,
+        "completion_tokens": 192,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 23.458772466001392,
+      "completion_tokens_per_second_including_prefill": 8.184571476545248
+    },
+    {
+      "id": "chatcmpl-a7f84a9ffb302b3b",
+      "object": "chat.completion",
+      "created": 1790592282,
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "choices": [
+        {
+          "index": 0,
+          "message": {
+            "role": "assistant",
+            "content": "\n\n29\n\nTo add 20 and 9, you can think of it as combining two groups of objects. Imagine you have 20 apples in one basket and 9 apples in another basket. If you pour all the apples together into a single large container, you would count them one by one to find the total. Starting with the 20 apples, you would count 21, 22, 23, 24, 25, 26, 27, 28, and finally 29. This process of counting upward from 20 by nine steps gives you the answer. Alternatively, you can use place value: 20 is two tens and zero ones, while 9 is zero tens and nine ones. Adding the ones together gives 0 + 9 = 9, and adding the tens gives 2 + 0 = 2 tens, which",
+            "refusal": null,
+            "annotations": null,
+            "audio": null,
+            "function_call": null,
+            "reasoning": null
+          },
+          "logprobs": null,
+          "finish_reason": "length",
+          "stop_reason": null,
+          "token_ids": null,
+          "routed_experts": null
+        }
+      ],
+      "service_tier": null,
+      "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b",
+      "usage": {
+        "prompt_tokens": 78,
+        "total_tokens": 270,
+        "completion_tokens": 192,
+        "prompt_tokens_details": null,
+        "completion_tokens_details": {
+          "reasoning_tokens": 0
+        }
+      },
+      "prompt_logprobs": null,
+      "prompt_token_ids": null,
+      "prompt_text": null,
+      "kv_transfer_params": null,
+      "ec_transfer_params": null,
+      "metrics": null,
+      "elapsed_seconds": 25.521558972002822,
+      "completion_tokens_per_second_including_prefill": 7.523051401782478
+    }
+  ]
+}

validation/native-vllm-streaming-validation.json

@@ -0,0 +1,1493 @@
+{
+  "passed": true,
+  "engine": "vLLM 0.28.0",
+  "tool_call": {
+    "request": {
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "messages": [
+        {
+          "role": "user",
+          "content": "Call get_weather for Warsaw. Do not answer from memory."
+        }
+      ],
+      "temperature": 0,
+      "max_tokens": 256,
+      "stream": true,
+      "stream_options": {
+        "include_usage": true
+      },
+      "return_token_ids": true,
+      "chat_template_kwargs": {
+        "enable_thinking": false
+      },
+      "tools": [
+        {
+          "type": "function",
+          "function": {
+            "name": "get_weather",
+            "description": "Get weather for a city",
+            "parameters": {
+              "type": "object",
+              "properties": {
+                "city": {
+                  "type": "string"
+                }
+              },
+              "required": [
+                "city"
+              ]
+            }
+          }
+        }
+      ],
+      "tool_choice": "auto"
+    },
+    "content": "",
+    "reasoning": "",
+    "tool_calls": {
+      "0": {
+        "name": "get_weather",
+        "arguments": "{\"city\": \"Warsaw\"}"
+      }
+    },
+    "usage": {
+      "prompt_tokens": 282,
+      "total_tokens": 305,
+      "completion_tokens": 23,
+      "completion_tokens_details": {
+        "reasoning_tokens": 0
+      }
+    },
+    "events": [
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "role": "assistant",
+              "content": ""
+            },
+            "logprobs": null,
+            "finish_reason": null
+          }
+        ],
+        "prompt_token_ids": [
+          151644,
+          8948,
+          198,
+          25139,
+          287,
+          30022,
+          371,
+          25,
+          7487,
+          151645,
+          151644,
+          8948,
+          198,
+          2610,
+          525,
+          12812,
+          533,
+          11250,
+          15,
+          13,
+          20,
+          12,
+          23876,
+          11,
+          264,
+          10950,
+          15235,
+          17847,
+          44936,
+          553,
+          12812,
+          533,
+          15235,
+          382,
+          2,
+          13852,
+          271,
+          2610,
+          1231,
+          1618,
+          825,
+          476,
+          803,
+          5746,
+          311,
+          7789,
+          448,
+          279,
+          1196,
+          3239,
+          382,
+          2610,
+          614,
+          2615,
+          311,
+          279,
+          2701,
+          5746,
+          1447,
+          27,
+          15918,
+          397,
+          27,
+          1688,
+          397,
+          27,
+          606,
+          29,
+          455,
+          69364,
+          522,
+          606,
+          397,
+          27,
+          4684,
+          29,
+          1949,
+          9104,
+          369,
+          264,
+          3283,
+          522,
+          4684,
+          397,
+          27,
+          13786,
+          397,
+          27,
+          16181,
+          397,
+          27,
+          606,
+          29,
+          8926,
+          522,
+          606,
+          397,
+          27,
+          1313,
+          29,
+          917,
+          522,
+          1313,
+          397,
+          522,
+          16181,
+          397,
+          27,
+          6279,
+          29,
+          1183,
+          8926,
+          1341,
+          522,
+          6279,
+          397,
+          522,
+          13786,
+          397,
+          522,
+          1688,
+          397,
+          522,
+          15918,
+          1339,
+          2461,
+          1817,
+          729,
+          1618,
+          11,
+          2550,
+          279,
+          729,
+          829,
+          323,
+          5977,
+          304,
+          279,
+          2701,
+          3561,
+          510,
+          151657,
+          198,
+          27,
+          1688,
+          28,
+          8687,
+          9174,
+          1269,
+          397,
+          27,
+          16181,
+          28,
+          8687,
+          24899,
+          62,
+          16,
+          29,
+          957,
+          62,
+          16,
+          522,
+          16181,
+          397,
+          27,
+          16181,
+          28,
+          8687,
+          24899,
+          62,
+          17,
+          41993,
+          374,
+          279,
+          897,
+          369,
+          279,
+          2086,
+          5733,
+          198,
+          9033,
+          646,
+          9390,
+          198,
+          35673,
+          5128,
+          522,
+          16181,
+          397,
+          522,
+          1688,
+          397,
+          151658,
+          271,
+          27,
+          98743,
+          397,
+          12,
+          5712,
+          6738,
+          27732,
+          1795,
+          279,
+          5189,
+          3561,
+          25,
+          458,
+          9179,
+          366,
+          1688,
+          28,
+          1112,
+          1472,
+          1688,
+          29,
+          2504,
+          1969,
+          387,
+          24034,
+          2878,
+          220,
+          151657,
+          151658,
+          11874,
+          9492,
+          198,
+          12,
+          9319,
+          4183,
+          990,
+          729,
+          6738,
+          4766,
+          220,
+          151667,
+          151668,
+          9492,
+          624,
+          12,
+          576,
+          897,
+          43810,
+          1948,
+          5733,
+          9492,
+          374,
+          33583,
+          6896,
+          438,
+          30430,
+          11,
+          2670,
+          501,
+          7969,
+          323,
+          12621,
+          624,
+          522,
+          98743,
+          29,
+          151645,
+          151644,
+          872,
+          198,
+          7220,
+          633,
+          69364,
+          369,
+          72176,
+          13,
+          3155,
+          537,
+          4226,
+          504,
+          4938,
+          13,
+          151645,
+          151644,
+          77091,
+          198,
+          151667,
+          151668
+        ],
+        "prompt_text": null
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              271
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              151657
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              198
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              27
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1688
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              28280
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              69364
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "tool_calls": [
+                {
+                  "id": "chatcmpl-tool-b7a4891d78286e7f",
+                  "type": "function",
+                  "index": 0,
+                  "function": {
+                    "name": "get_weather"
+                  }
+                }
+              ]
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              397
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              27
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              16181
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              28
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              8926
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "tool_calls": [
+                {
+                  "index": 0,
+                  "function": {
+                    "arguments": "{\"city\": \""
+                  }
+                }
+              ]
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              29
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "tool_calls": [
+                {
+                  "index": 0,
+                  "function": {
+                    "arguments": "Wars"
+                  }
+                }
+              ]
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              98970
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "tool_calls": [
+                {
+                  "index": 0,
+                  "function": {
+                    "arguments": "aw"
+                  }
+                }
+              ]
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              672
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              522
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              16181
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              397
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              522
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1688
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "tool_calls": [
+                {
+                  "index": 0,
+                  "function": {
+                    "arguments": "\"}"
+                  }
+                }
+              ]
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              397
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              151658
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": "tool_calls",
+            "stop_reason": null,
+            "token_ids": [
+              151645
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a08784d157aa41e3",
+        "object": "chat.completion.chunk",
+        "created": 1790592279,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [],
+        "usage": {
+          "prompt_tokens": 282,
+          "total_tokens": 305,
+          "completion_tokens": 23,
+          "completion_tokens_details": {
+            "reasoning_tokens": 0
+          }
+        },
+        "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b"
+      }
+    ]
+  },
+  "thinking": {
+    "request": {
+      "model": "Naive-N0.5-Flash-NVFP4",
+      "messages": [
+        {
+          "role": "user",
+          "content": "Calculate 6 times 7. Think briefly, then give the answer."
+        }
+      ],
+      "temperature": 0,
+      "max_tokens": 256,
+      "stream": true,
+      "stream_options": {
+        "include_usage": true
+      },
+      "return_token_ids": true,
+      "chat_template_kwargs": {
+        "enable_thinking": true
+      },
+      "reasoning_effort": "low"
+    },
+    "content": "\n\n42",
+    "reasoning": "\nWe need answer simple. Need think briefly then answer. 6*7=42. Final concise.\n",
+    "tool_calls": {},
+    "usage": {
+      "prompt_tokens": 57,
+      "total_tokens": 86,
+      "completion_tokens": 29,
+      "completion_tokens_details": {
+        "reasoning_tokens": 23
+      }
+    },
+    "events": [
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "role": "assistant",
+              "content": ""
+            },
+            "logprobs": null,
+            "finish_reason": null
+          }
+        ],
+        "prompt_token_ids": [
+          151644,
+          8948,
+          198,
+          25139,
+          287,
+          30022,
+          371,
+          25,
+          12041,
+          151645,
+          151644,
+          8948,
+          198,
+          2610,
+          525,
+          12812,
+          533,
+          11250,
+          15,
+          13,
+          20,
+          12,
+          23876,
+          11,
+          264,
+          10950,
+          15235,
+          17847,
+          44936,
+          553,
+          12812,
+          533,
+          15235,
+          13,
+          151645,
+          151644,
+          872,
+          198,
+          47866,
+          220,
+          21,
+          3039,
+          220,
+          22,
+          13,
+          21149,
+          26753,
+          11,
+          1221,
+          2968,
+          279,
+          4226,
+          13,
+          151645,
+          151644,
+          77091,
+          198
+        ],
+        "prompt_text": null
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              151667
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "\n"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              198
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "We"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1654
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " need"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1184
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " answer"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              4226
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " simple"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              4285
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "."
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              13
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " Need"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              14656
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " think"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1744
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " briefly"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              26753
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " then"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              1221
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " answer"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              4226
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "."
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              13
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " "
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              220
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "6"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              21
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "*"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              9
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "7"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              22
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "="
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              28
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "4"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              19
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "2"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              17
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": "."
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              13
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " Final"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              13023
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": " concise"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              63594
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "reasoning": ".\n"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              624
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              151668
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              271
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "content": "\n\n4"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              19
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {
+              "content": "2"
+            },
+            "logprobs": null,
+            "finish_reason": null,
+            "token_ids": [
+              17
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [
+          {
+            "index": 0,
+            "delta": {},
+            "logprobs": null,
+            "finish_reason": "stop",
+            "stop_reason": null,
+            "token_ids": [
+              151645
+            ]
+          }
+        ]
+      },
+      {
+        "id": "chatcmpl-a9f0496c02f6b857",
+        "object": "chat.completion.chunk",
+        "created": 1790592280,
+        "model": "Naive-N0.5-Flash-NVFP4",
+        "choices": [],
+        "usage": {
+          "prompt_tokens": 57,
+          "total_tokens": 86,
+          "completion_tokens": 29,
+          "completion_tokens_details": {
+            "reasoning_tokens": 23
+          }
+        },
+        "system_fingerprint": "vllm-0.28.0-tp4-13be3b5b"
+      }
+    ]
+  }
+}

validation/portable-preflight.json

@@ -0,0 +1,39 @@
+{
+  "passed": true,
+  "versions": {
+    "vllm": "0.28.0",
+    "torch": "2.13.0+cu130",
+    "transformers": "5.17.0",
+    "accelerate": "1.15.0",
+    "tokenizers": "0.23.2",
+    "huggingface-hub": "1.33.0",
+    "safetensors": "0.8.0",
+    "flashinfer-python": "0.7.0",
+    "nvidia-cutlass-dsl": "4.6.2",
+    "triton": "3.7.1"
+  },
+  "gpus": [
+    {
+      "index": 0,
+      "name": "NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition",
+      "memory_bytes": 101973491712
+    },
+    {
+      "index": 1,
+      "name": "NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition",
+      "memory_bytes": 101973491712
+    },
+    {
+      "index": 2,
+      "name": "NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition",
+      "memory_bytes": 101970345984
+    },
+    {
+      "index": 3,
+      "name": "NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition",
+      "memory_bytes": 101973491712
+    }
+  ],
+  "shards": 49,
+  "tensors": 144856
+}

validation/preserved-weights-validation.json

@@ -0,0 +1,10 @@
+{
+  "passed": true,
+  "tensors_checked": 472,
+  "restored_from_source": {
+    "model-layer-00047.safetensors": [
+      "model.layers.47.mlp.gate.weight"
+    ]
+  },
+  "check": "Exact source dtype, shape and SHA-256 of tensor payload"
+}

validation/runtime-source-validation.json

@@ -0,0 +1,17 @@
+{
+  "passed": true,
+  "scope": "Active native vLLM container module bytes match published source files",
+  "files": {
+    "vllm_naive/attention.py": "f3aa650b110495224338b211330f6dcb7598cbeff917e7b1dac77695234f187c",
+    "vllm_naive/__init__.py": "8a5fddf2b89a4c33b6774a3ed75004a73db990ee711b8411e8772d7f2786f679",
+    "vllm_naive/model.py": "f0d5769018f8487c267d0ee6c3ba4cbf138e47d65c35c0cc0a00a95dfb69d763",
+    "scripts/benchmark_stream.py": "dab13abe554c6c28f89268ee3a2cec2e6425f73aa91f18692318b76886704157",
+    "scripts/safe_device_transfer.py": "61e931597e59e72ded3525976181e7164933e72f4647423c00ad2be0a57940db",
+    "scripts/decode_kernels.py": "ea0e71f5c4fd678b5cdf423b62ad07d37bed7de5fdc2af9ee89f08856ee34d09",
+    "scripts/validate_api.py": "82e04b0d21e9956da18930e27d35c4fc45f1a4b071bd2a19c6f62271deb74c2d",
+    "scripts/nvfp4_runtime.py": "be6453ca83d8f9ec528b07f17e78a41bfdcd8859fa9ec2644b6136e7068e9a1a",
+    "scripts/fused_experts.py": "f77eceec0a969eb7ff99e5177bd9d4cda1bfa3f0668ed614c060d28883f46018",
+    "scripts/preflight.py": "43ac170da1d7cf614fb5f4d7f0ee70635d2d78bd638bfb3af2322baf91a08600",
+    "sitecustomize.py": "552485e6e5e03c72c1d8b76a43a1ede48f8283a5d1e5ab93761125ec4a2c492b"
+  }
+}