Publish Ling 3.0 Flash CIRU Strix-native runtime

.gitattributes

@@ -1,35 +1,6 @@
-*.7z filter=lfs diff=lfs merge=lfs -text
-*.arrow filter=lfs diff=lfs merge=lfs -text
-*.bin filter=lfs diff=lfs merge=lfs -text
-*.bz2 filter=lfs diff=lfs merge=lfs -text
-*.ckpt filter=lfs diff=lfs merge=lfs -text
-*.ftz filter=lfs diff=lfs merge=lfs -text
-*.gz filter=lfs diff=lfs merge=lfs -text
-*.h5 filter=lfs diff=lfs merge=lfs -text
-*.joblib filter=lfs diff=lfs merge=lfs -text
-*.lfs.* filter=lfs diff=lfs merge=lfs -text
-*.mlmodel filter=lfs diff=lfs merge=lfs -text
-*.model filter=lfs diff=lfs merge=lfs -text
-*.msgpack filter=lfs diff=lfs merge=lfs -text
-*.npy filter=lfs diff=lfs merge=lfs -text
-*.npz filter=lfs diff=lfs merge=lfs -text
-*.onnx filter=lfs diff=lfs merge=lfs -text
-*.ot filter=lfs diff=lfs merge=lfs -text
-*.parquet filter=lfs diff=lfs merge=lfs -text
-*.pb filter=lfs diff=lfs merge=lfs -text
-*.pickle filter=lfs diff=lfs merge=lfs -text
-*.pkl filter=lfs diff=lfs merge=lfs -text
-*.pt filter=lfs diff=lfs merge=lfs -text
-*.pth filter=lfs diff=lfs merge=lfs -text
-*.rar filter=lfs diff=lfs merge=lfs -text
-*.safetensors filter=lfs diff=lfs merge=lfs -text
-saved_model/**/* filter=lfs diff=lfs merge=lfs -text
-*.tar.* filter=lfs diff=lfs merge=lfs -text
-*.tar filter=lfs diff=lfs merge=lfs -text
-*.tflite filter=lfs diff=lfs merge=lfs -text
-*.tgz filter=lfs diff=lfs merge=lfs -text
-*.wasm filter=lfs diff=lfs merge=lfs -text
-*.xz 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
+*.bundle binary
+*.patch text eol=lf
+*.sh text eol=lf
+*.md text eol=lf
+*.png filter=lfs diff=lfs merge=lfs -text
+Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz filter=lfs diff=lfs merge=lfs -text

CHANGELOG.md

@@ -0,0 +1,68 @@
+# Runtime change manifest
+
+## Immutable identities
+
+- Original model: `inclusionAI/Ling-3.0-flash-int4`
+- Model revision: `ca3ea63b0255d212c4fe6020db9e0a51ce136006`
+- Model format: 24 safetensors shards, 77,012,299,464 bytes total
+- vLLM upstream base: `d35eb6c44071ea806018841c490f0d2f3219c485`
+- CIRU runtime head: `388d82de2abe418e73380ba250270d3abc62ac12`
+- Net vLLM delta: six files, 365 insertions, 26 deletions
+
+## Weight changes
+
+None. No tensor, quantization scale, tokenizer, template, configuration, or MTP weight was modified. The official packed-INT4 checkpoint is consumed directly.
+
+## Source changes
+
+### `csrc/libtorch_stable/sampler.cu`
+
+Uses 512 rather than 1024 merge threads on ROCm so Wave32 hardware does not exceed the 64 KB LDS limit. This is a general ROCm correctness change.
+
+### `vllm/v1/attention/ops/merge_attn_states.py`
+
+Adds an opt-in Torch implementation selected by `VLLM_ROCM_SAFE_MERGE_ATTN_STATES=1`. It handles token-first and head-first LSE layouts and avoids an observed `gfx1151` HSA fault in the Triton merge kernel.
+
+### `vllm/v1/attention/backends/mla/triton_mla.py`
+
+Declares uniform query-length support and converts causal multi-token verifier blocks into per-token Triton MLA decode rows with correct causal KV-prefix lengths. This is relevant to MLA speculative verification, not W4A16 specifically.
+
+### `vllm/model_executor/layers/fused_moe/fused_moe.py`
+
+Adds a naive small-decode block assignment specialization that bypasses sorting/alignment only under a narrow validated guard.
+
+### `vllm/model_executor/layers/fused_moe/moe_fused_mul_sum.py`
+
+Adds a low-launch-overhead Triton reduction for one to three tokens, top-k 8, and hidden size 2560.
+
+### `vllm/model_executor/layers/fused_moe/experts/triton_moe.py`
+
+Wires the guarded WNA16 paths together and adds the exact small-shape `gfx1151` SiLU-and-multiply kernel. Every adjacent shape or unsupported configuration retains the upstream route.
+
+## Runtime configuration
+
+- Triton MLA attention
+- Triton MoE backend
+- native checkpoint MTP, K1 (`num_speculative_tokens=1`)
+- compile sizes `[1,2]`
+- CUDA graphs disabled
+- chunked prefill and prefix caching enabled
+- AITER and ROCm skinny GEMM disabled
+- safe attention merge enabled
+- native context 262,144 tokens
+- five active sequences for 256K; two for experimental 1M
+
+## Performance context
+
+The runnable public/upstream-compatible starting path measured 0.268944 tok/s in a warmed deterministic target-only decode test. The released runtime measures 21.4407 tok/s target-only and 26.2343 effective tok/s with native K1 MTP: 79.72x and 97.55x the starting throughput, respectively.
+
+The final exact-shape SiLU candidate passed bitwise source fixtures and exact output checks. In the matched final-step comparison against the already-optimized CIRU parent, it improved target-only decode by 7.7457% and K1 effective decode by 4.3743%. Those percentages describe only the last runtime optimization, not the public-to-release gain. Only clean, resource-gated measurements were retained as promotion evidence.
+
+## Public release packaging
+
+- Added the CIRU model-card artwork at `assets/ling30int4.png`.
+- Added distro-aware host dependencies and portable Python 3.12 provisioning through `uv`.
+- Added native-Linux instructions for Ubuntu/Debian, Fedora, Arch, and the validated NixOS boundary.
+- Added an explicitly experimental Windows 11 WSL2/AMD ROCDXG path; native Windows vLLM is not claimed.
+- Changed the public native-256K profile from `gpu_memory_utilization=0.82` to the validated long-context `0.75` envelope. It retains a 2,519,792-token KV pool and five active sequence slots while avoiding the observed 60K memory-pressure failure.
+- Added the retained 60K PP/TG comparison against AtomicChat AD-IQ4_XXS and ROCmFP4 STRIX MTP, with the native row marked for a final-commit refresh.

Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz

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

MANIFEST.json

@@ -0,0 +1,68 @@
+{
+  "release_name": "Ling-3.0-Flash-CIRU-int4-Strix-native",
+  "release_archive": "Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz",
+  "release_type": "runtime-distribution",
+  "weights_changed": false,
+  "base_model": {
+    "repo_id": "inclusionAI/Ling-3.0-flash-int4",
+    "revision": "ca3ea63b0255d212c4fe6020db9e0a51ce136006",
+    "format": "compressed-tensors packed INT4 safetensors",
+    "shards": 24,
+    "shard_bytes": 77012299464,
+    "native_context": 262144
+  },
+  "runtime": {
+    "upstream_repo": "https://github.com/vllm-project/vllm.git",
+    "upstream_base": "d35eb6c44071ea806018841c490f0d2f3219c485",
+    "ciru_head": "388d82de2abe418e73380ba250270d3abc62ac12",
+    "commits_ahead": 15,
+    "net_files_changed": 6,
+    "net_insertions": 365,
+    "net_deletions": 26
+  },
+  "validated_stack": {
+    "python": "3.12.12",
+    "torch": "2.13.0+rocm7.15.0a20260728",
+    "hip": "7.15.0",
+    "triton": "3.8.0+git4cff872c.rocm7.15.0a20260728",
+    "flash_attention": "2.8.3",
+    "gpu": "AMD Radeon 8060S",
+    "architecture": "gfx1151"
+  },
+  "profiles": {
+    "native_256k": {"max_model_len": 262144, "max_num_seqs": 5, "gpu_memory_utilization": 0.75, "kv_pool_tokens_observed": 2519792},
+    "experimental_yarn_1m": {"max_model_len": 1048576, "max_num_seqs": 2, "native_training": false}
+  },
+  "platforms": {
+    "validated": ["NixOS native Linux on gfx1151"],
+    "documented_unvalidated": ["Ubuntu/Debian native Linux", "Fedora native Linux", "Arch native Linux", "Windows 11 WSL2 with AMD ROCDXG"],
+    "unsupported": ["native Windows vLLM"]
+  },
+  "performance": {
+    "host_scope": "single AMD Strix Halo gfx1151 host",
+    "public_starting_target_only_tokens_per_second": 0.268944,
+    "release_target_only_tokens_per_second": 21.4407,
+    "release_native_mtp_k1_effective_tokens_per_second": 26.2343,
+    "release_target_only_multiple": 79.72,
+    "release_native_mtp_k1_multiple": 97.55,
+    "comparison_scope": "full campaign progression; starting TG8 and final TG32, not a single-variable matched A/B",
+    "exact_2k_k1_concurrency": {
+      "c1_aggregate_decode_tokens_per_second": 26.33,
+      "c2_aggregate_decode_tokens_per_second": 38.90,
+      "c4_aggregate_decode_tokens_per_second": 59.03,
+      "c5_aggregate_decode_tokens_per_second": 62.16,
+      "validated_max_concurrency": 5,
+      "c6_status": "engine failure in chunked-prefill MLA accumulation; public profiles cap active sequences below this boundary"
+    },
+    "retained_60k_prompt_512_output": {
+      "native_source": "fabbc7e928dfa64c29dcdf0f795b05246555f4ae",
+      "native_pp_tokens_per_second": 419.334843,
+      "native_tg_tokens_per_second": 19.924752,
+      "atomic_chat_pp_tokens_per_second": 265.63,
+      "atomic_chat_tg_tokens_per_second": 25.11,
+      "rocmfp4_pp_tokens_per_second": 249.23,
+      "rocmfp4_tg_tokens_per_second": 19.00,
+      "refresh_required_for_release_head": true
+    }
+  }
+}

README.md

@@ -0,0 +1,162 @@
+---
+license: mit
+base_model: inclusionAI/Ling-3.0-flash-int4
+pipeline_tag: text-generation
+library_name: vllm
+tags:
+  - rocm
+  - amd
+  - strix-halo
+  - gfx1151
+  - vllm
+  - w4a16
+  - speculative-decoding
+  - mtp
+  - windows
+  - wsl
+---
+
+# Ling-3.0-Flash-CIRU-int4-Strix-native
+
+![Ling 3.0 Flash CIRU INT4 for Strix Halo](assets/ling30int4.png)
+
+`Ling-3.0-Flash-CIRU-int4-Strix-native` is CIRU's native vLLM/ROCm runtime distribution for InclusionAI's official packed-INT4 Ling 3.0 Flash checkpoint on AMD Strix Halo (`gfx1151`).
+
+The model weights were **not** changed, requantized, merged, fine-tuned, or renamed internally. This release combines the original checkpoint at revision `ca3ea63b0255d212c4fe6020db9e0a51ce136006` with a pinned vLLM fork, ROCm 7.15 build recipe, and launch profiles validated on an AMD Radeon 8060S with 128 GB unified memory.
+
+**A 97.5x native ROCm serving campaign:** the first working upstream-compatible build decoded at 0.269 tok/s; this distribution reaches 26.23 tok/s with native MTP K1, while retaining the official vendor checkpoint, OpenAI-compatible serving, continuous batching, and Ling's structured tool parser.
+
+## What this release changes
+
+The runtime is based on vLLM commit `d35eb6c44071ea806018841c490f0d2f3219c485` plus a 15-commit CIRU branch ending at `388d82de2abe418e73380ba250270d3abc62ac12`. Its net source delta is six files, 365 insertions, and 26 deletions.
+
+- Fixes a Wave32 LDS overflow in the ROCm top-k merge launch.
+- Adds an opt-in safe PyTorch attention-state merge for a `gfx1151` Triton/HSA fault, including both supported MLA LSE layouts.
+- Routes causal uniform multi-token MLA verification blocks through Triton decode with correct per-token sequence lengths.
+- Adds narrowly guarded W4A16 MoE decode/verification fast paths for Ling's exact small-token geometry.
+- Adds a `gfx1151` Triton SiLU-and-multiply kernel for Ling's exact eager W4A16 shapes.
+
+See [CHANGELOG.md](CHANGELOG.md) for the public file-level runtime diff.
+
+## Performance: 97.5x over the public starting point
+
+The closest runnable public/upstream starting point was frozen vLLM commit `d35eb6c44071ea806018841c490f0d2f3219c485` plus the publicly proposed five-line ROCm sampler fix required to launch on `gfx1151`. With the unchanged official model, that path produced **0.268944 tok/s** in a warmed deterministic target-only decode test. This release reaches **26.2343 effective tok/s** with the checkpoint's native K1 MTP: **97.55x the starting throughput**, or approximately **100x**.
+
+| Public-to-release progression | Decode throughput | Multiple vs. starting point |
+|---|---:|---:|
+| Upstream-compatible starting path, target only | 0.268944 tok/s | 1.00x |
+| This release, target only | 21.4407 tok/s | 79.72x |
+| This release, native MTP K1 | 26.2343 effective tok/s | 97.55x (~100x) |
+
+The starting row is a warmed deterministic eight-token test; the final rows are medians from deterministic 32-token tests. The multiplier therefore describes the full optimization campaign on the same model and host, not a single-variable matched A/B. The final exact-shape kernel step was also tested separately against the already-optimized CIRU parent: target-only decode rose from 19.8993 to 21.4407 tok/s (+7.7457%), and K1 rose from 25.1348 to 26.2343 tok/s (+4.3743%). That smaller percentage is the final incremental step, not the overall release gain.
+
+MTP acceptance remained 82.35% with 1.8235 mean tokens per step. Retained controls had exact token IDs/text and zero faults. These are local single-host measurements, not claims for other AMD architectures.
+
+### How 0.269 became 26.23 tok/s
+
+| Stage | Decode throughput | Incremental result |
+|---|---:|---:|
+| First working upstream-compatible build | 0.269 tok/s | starting point |
+| Disable pathological ROCm skinny-GEMM dispatch on Wave32 | 7.57 tok/s | 28.16x |
+| vLLM compilation mode 3, decode size 1, graphs off | 11.21 tok/s | +48.1% |
+| Coherent ROCm 7.15 / Torch 2.13 / Triton 3.8 stack | 14.84 tok/s | +32.4% |
+| W4A16 MoE expert assignment and finalize reduction | 19.71 tok/s | +42.4% |
+| Repair optimized multi-token verifier routing; native MTP K1 | 25.13 tok/s | +27.3% |
+| Exact-shape `gfx1151` SiLU-and-multiply kernel | **26.23 tok/s** | +4.37% K1; 97.55x overall |
+
+The largest single problem was silent selection of vLLM's ROCm skinny-GEMM path on Wave32: disabling it alone improved identical-output decode by 28.16x. The remaining gains came from compilation, a coherent AMD software stack, Ling-shaped W4A16 MoE paths, repaired MTP verification routing, and the final exact-shape activation kernel.
+
+### Serving results
+
+On exact 2K prompts with 256 forced output tokens, K1 MTP, and simultaneous bursts:
+
+| Concurrent requests | Aggregate prefill | Aggregate decode | Per-request decode |
+|---:|---:|---:|---:|
+| 1 | 521.06 tok/s | 26.33 tok/s | 26.33 tok/s |
+| 2 | 503.31 tok/s | 38.90 tok/s | 19.73 tok/s |
+| 4 | 508.15 tok/s | 59.03 tok/s | 15.34 tok/s |
+| 5 | 468.30 tok/s | **62.16 tok/s** | 14.40 tok/s |
+
+The C4 row reproduced at 59.15 and 59.03 aggregate decode tok/s across independent loads. For context, CIRU's local llama.cpp Vulkan GGUF reference reached 51.08 tok/s single-stream with MTP: this native vLLM stack is slower single-stream, but exceeds that single-stream reference at C4 and C5 while using the official checkpoint and vLLM serving stack.
+
+The runtime also sustained 540.19, 550.20, and 536.13 prefill tok/s at 2K, 8K, and 16K in clean long-prefill tests on an earlier retained build; treat those rows as a floor rather than a claim measured again on the final commit. The final stack completed CIRU's 20-scenario structured tool-call suite using Ling's native parser.
+
+### 60K PP/TG comparison
+
+This retained comparison uses one frozen 60,000-token prompt followed by 512 forced output tokens:
+
+| Stack | PP | TG | Request wall |
+|---|---:|---:|---:|
+| **CIRU native K1 MTP** | **419.33 tok/s** | 19.92 tok/s | **168.75 s** |
+| AtomicChat AD-IQ4_XXS, llama.cpp Vulkan | 265.63 tok/s | **25.11 tok/s** | 246.28 s |
+| ROCmFP4 STRIX MTP, llama.cpp Vulkan | 249.23 tok/s | 19.00 tok/s | 267.84 s |
+
+Native vLLM gives up 20.66% TG to Atomic at 60K, but its 57.87% PP lead completes the combined request 31.48% sooner. It leads ROCmFP4 by 68.25% PP, 4.88% TG, and 37.00% in combined wall time.
+
+The native row predates the release's final SiLU kernel commit, so this table is a retained baseline rather than a final-commit maximum. See [docs/BENCHMARKS.md](docs/BENCHMARKS.md) for the exact protocol, source identities, and refresh rule.
+
+## Files in this package
+
+- `runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle`: exact Git objects for the CIRU branch, based on the pinned upstream commit.
+- `runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch`: reviewable net source diff.
+- `scripts/download-model.sh`: downloads the unchanged official checkpoint into a directory named `Ling-3.0-Flash-CIRU-int4-Strix-native`.
+- `scripts/install-host-deps.sh`: distro-aware Linux/WSL host dependency installer.
+- `scripts/build-vllm-gfx1151.sh`: reproducible build entry point for the pinned AMD wheel stack.
+- `scripts/run-256k.sh`: native 256K, native-MTP K1 server profile with five active sequence slots.
+- `scripts/run-1m-yarn-experimental.sh`: experimental 4x YaRN 1M profile with two active sequence slots.
+- `MANIFEST.json`: machine-readable identities and scope.
+
+The named release archive is `Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz`. It contains the runtime distribution, image, and documentation—not a second 77 GB copy of unchanged weights.
+
+## Quick start
+
+The exact validated environment is Python 3.12, Torch `2.13.0+rocm7.15.0a20260728`, HIP `7.15.0`, Triton `3.8.0+git4cff872c.rocm7.15.0a20260728`, FlashAttention `2.8.3`, and the CIRU vLLM source at `388d82de2`.
+
+```bash
+bash scripts/install-host-deps.sh
+bash scripts/build-vllm-gfx1151.sh "$HOME/ciru-ling-runtime"
+bash scripts/download-model.sh "$HOME/models"
+
+VENV="$HOME/ciru-ling-runtime/.venv" \
+MODEL_PATH="$HOME/models/Ling-3.0-Flash-CIRU-int4-Strix-native" \
+bash scripts/run-256k.sh
+```
+
+The API advertises model ID `Ling-3.0-Flash-CIRU-int4-Strix-native` on `http://127.0.0.1:18081/v1`.
+
+Choose the platform guide before building:
+
+- [Native Linux: Ubuntu/Debian, Fedora, Arch, and NixOS notes](docs/UBUNTU_BUILD.md)
+- [Windows 11 through WSL2 and AMD ROCDXG](docs/WINDOWS_WSL2.md)
+
+The exact runtime has been validated on CIRU's native-Linux NixOS Strix Halo host. Mainstream native-Linux and Windows/WSL instructions are designed to be reproducible but are not yet CIRU clean-host validations. Native Windows vLLM is not supported; Windows users must use WSL2.
+
+## Context and concurrency
+
+The official checkpoint is natively configured and trained for 256K context. `run-256k.sh` uses `max_model_len=262144`, `max_num_seqs=5`, and the validated `gpu_memory_utilization=0.75` long-context envelope. That envelope reported a 2,519,792-token KV pool, enough for five full 256K sequences.
+
+Five is also the validated serving safety cap. A six-request 2K+256 simultaneous burst terminated the engine in the chunked-prefill MLA accumulation path with an LSE extent mismatch. With `max_num_seqs=5`, additional requests wait in the scheduler instead of entering that unvalidated C6 path.
+
+The 1M profile is an experimental 4x YaRN extension, not native 1M training. It uses two active slots so two full-length requests fit; further requests wait in the vLLM scheduler.
+
+## Does this work for other W4A16 models?
+
+The fork remains a vLLM runtime and safely falls back to upstream paths when its guards do not match, but the measured speedups are mostly Ling-specific:
+
+- The Wave32 top-k fix is broadly relevant to ROCm Wave32 GPUs.
+- Safe attention merging and uniform verifier routing are relevant to compatible ROCm MLA workloads, independent of W4A16.
+- The W4A16 fast assignment path currently requires 512 experts, top-k 8, block size 32, and one to three tokens.
+- The fast reduction requires hidden size 2560 and top-k 8.
+- The SiLU path requires exact `(8,1536)->(8,768)` or `(16,1536)->(16,768)` BF16 shapes on `gfx1151`.
+
+Other W4A16 models will run only if already supported by this vLLM base; they should not be expected to receive these speedups unless their geometry matches. Generalizing those kernels is part of the upstream plan.
+
+## Credits and license
+
+- Model and official INT4 weights: [InclusionAI/Ling-3.0-flash-int4](https://huggingface.co/inclusionAI/Ling-3.0-flash-int4), MIT license.
+- Inference engine: [vLLM](https://github.com/vllm-project/vllm), Apache-2.0.
+- ROCm, Torch, and Triton packages: AMD and their respective upstream projects.
+- Strix Halo kernel, TTM/GTT, and WSL/ROCDXG guidance: [AMD ROCm documentation](https://rocm.docs.amd.com/en/docs-7.2.0/how-to/system-optimization/strixhalo.html) and [ROCDXG](https://github.com/ROCm/librocdxg).
+- Runtime integration, profiling, kernels, and validation: CIRU.
+
+The original model license continues to govern the weights. The vLLM patch and bundle retain vLLM's upstream licensing and file headers.

SHA256SUMS

@@ -0,0 +1,21 @@
+eb208bf38f31d7b068af8e8351378615017b911eb7d0302264d5398cde5d403e  ./.gitattributes
+411ac21b4849cb3bf65f84a55263550dcd658e0623c2b117233ba1fced1c17ff  ./assets/ling30int4.png
+82cf1d95343c44abfbf6daec0d3e0e23240de2caf44da0e2cc797e94fa2f876c  ./CHANGELOG.md
+26c2195a819fe61daa19b1e73b2371d597240aba0dd5f614719209b9dd76f07f  ./config/experimental-yarn-1m.env
+eb4b47674187eaf3035efec1fc0f17a33ac540bf94391cbaae9fec0119ed2fae  ./config/native-256k.env
+788ea98c240e3009cc44b0e1899c87e542f4394922e4ff16a8479c7e1003dc22  ./docs/BENCHMARKS.md
+acdccd00a7605cf06001471cc7cad3f07f8bcd5edafa935cafb521b1d51392d2  ./docs/UBUNTU_BUILD.md
+0d6258e6b4a896feedcf6bf6d6b332a2a6eabde0f7f657828de81c303c540701  ./docs/VALIDATED_ENVIRONMENT.md
+3bfadd289f312c0cd99eea1883c7352828268f360b7b5e53c40f3175d6ed1593  ./docs/WINDOWS_WSL2.md
+88f444c673ba58f9b87d70d664b0876669cf648177358c709dfe180a557ee694  ./MANIFEST.json
+1618febb4927cc20d6dfc237841b1456e782bbcaf5bb16744996972c7a885a5d  ./README.md
+15e48a8df3b65a73c81cf73c820ba12b4ab9474697d6b3f876c26c0b0cd3d455  ./runtime/COMMITS.txt
+6b03af3a1685a1969ff96870ea776537f5f67fd3674fc353880bd088bcef9a9c  ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
+8f307e9f2f16207856ca0b064a765a73bbebc029ad0b0fed61e5975a17e07ab6  ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
+f4caeda80ed751a40fa976aa7853357b88075e6eb6d2d49e68b0bc8599754c34  ./runtime/SHA256SUMS.runtime
+25ab169703ed6c39260f9336f3d484c9a11b1f9b31c7f9072866b158365de983  ./scripts/build-vllm-gfx1151.sh
+6a8530b9b38c52d6fe8efab5dd3a6bbf02a3f24a0b9a800b86a6dde3a3103071  ./scripts/download-model.sh
+d6aa5efee74c850f2ab65efeb4dfb66142c440d66276983282de6aec6a855270  ./scripts/install-host-deps.sh
+27d3cfba9f4178f83cbd2d7e44db098ab69e515f76e5c98c40ac219b7b0597ca  ./scripts/run-1m-yarn-experimental.sh
+296ee94a925db5b97e21afb6ffd4ae3c0cc1c2be67af81e21350f43efaa59ead  ./scripts/run-256k.sh
+aa5f1137c666abe07c9fc37a64547dc535d461958db58f7da80beea4133367f9  ./scripts/run-common.sh

config/experimental-yarn-1m.env

@@ -0,0 +1,6 @@
+MODEL_ID=Ling-3.0-Flash-CIRU-int4-Strix-native
+CONTEXT_LENGTH=1048576
+MAX_NUM_SEQS=2
+GPU_MEMORY_UTILIZATION=0.82
+SPECULATIVE_TOKENS=1
+HF_OVERRIDES='{"max_position_embeddings":1048576,"rope_scaling":{"rope_type":"yarn","factor":4.0,"original_max_position_embeddings":262144}}'

config/native-256k.env

@@ -0,0 +1,5 @@
+MODEL_ID=Ling-3.0-Flash-CIRU-int4-Strix-native
+CONTEXT_LENGTH=262144
+MAX_NUM_SEQS=5
+GPU_MEMORY_UTILIZATION=0.75
+SPECULATIVE_TOKENS=1

docs/BENCHMARKS.md

@@ -0,0 +1,30 @@
+# Benchmark notes
+
+All results below are local measurements on one Ryzen AI MAX+ 395 / Radeon 8060S Strix Halo system with 128 GB unified memory. They compare complete model-format and runtime stacks; they are not model-quality rankings.
+
+## 60K prompt plus 512-token decode
+
+The frozen workload used exactly 60,000 direct prompt tokens, 512 forced output tokens, temperature 0, EOS ignored, a unique cache salt, and no reusable prefix.
+
+| Stack | Format/runtime | PP | TG | Request wall |
+|---|---|---:|---:|---:|
+| **CIRU native K1 MTP** | official InclusionAI packed INT4, vLLM/ROCm | **419.33 tok/s** | 19.92 tok/s | **168.75 s** |
+| [AtomicChat AD-IQ4_XXS](https://huggingface.co/AtomicChat/Ling-3.0-flash-GGUF) | mixed GGUF, llama.cpp Vulkan | 265.63 tok/s | **25.11 tok/s** | 246.28 s |
+| [raulvidis ROCmFP4 STRIX MTP](https://huggingface.co/raulvidis/Ling-3.0-flash-ROCmFP4-STRIX-MTP-GGUF) | ROCmFP4 GGUF, llama.cpp Vulkan MTP | 249.23 tok/s | 19.00 tok/s | 267.84 s |
+
+At this workload, native vLLM trails Atomic's decode by 20.66% but leads its prefill by 57.87%, completing the combined request 31.48% sooner. Against ROCmFP4, native leads prefill by 68.25%, decode by 4.88%, and combined request time by 37.00%.
+
+The native row used the validated `gpu_memory_utilization=0.75`, 262,144-context envelope and K1 MTP. It accepted 254 of 258 draft tokens (98.45%) and completed without cache hits, preemption, JIT during measurement, or device faults.
+
+### Update status
+
+The native 60K row was measured at CIRU runtime commit `fabbc7e928dfa64c29dcdf0f795b05246555f4ae`, immediately before the final exact-shape SiLU kernel commit. The released commit `388d82de2abe418e73380ba250270d3abc62ac12` has not yet been rerun on this 60K protocol. The table is therefore a valid retained comparison, but not a claim that 19.92 tok/s is the final commit's maximum 60K speed. Replace the native row only after a matched rerun of all three stacks or a clearly labeled native-only refresh.
+
+## Short decode and concurrent serving
+
+- Strict final-commit single-request K1 TG32 median: **26.2343 tok/s**.
+- Highest individual strict final-commit TG32 row: **26.2445 tok/s**.
+- Exact-2K C1 serving sweep: **26.33 tok/s** aggregate decode.
+- Exact-2K C5 serving sweep: **62.16 tok/s** aggregate decode, or 14.40 tok/s per request.
+
+Use 26.23 tok/s as the reproducible single-request headline and 62.16 tok/s only when explicitly labeled five-request aggregate throughput.

docs/UBUNTU_BUILD.md

@@ -0,0 +1,125 @@
+# Linux build guide for Strix Halo
+
+This guide covers native Linux on AMD Ryzen AI MAX / MAX+ (`gfx1151`). The release was built and exercised end-to-end on NixOS; Ubuntu/Debian, Fedora, and Arch use the same self-contained ROCm/Torch wheel transaction but remain community portability paths until clean-host reports are retained.
+
+## What is installed where
+
+There are two separate layers:
+
+1. The host must expose the Strix Halo GPU through the Linux AMDGPU/KFD driver as `/dev/kfd` and `/dev/dri`.
+2. `scripts/build-vllm-gfx1151.sh` installs the exact ROCm 7.15, Torch 2.13, Triton 3.8, and development SDK wheels into a private virtual environment. Do not mix a different system PyTorch or ROCm userspace into that environment.
+
+The wheel layout and `device-gfx1151` extra come from [AMD TheRock's Python packaging](https://github.com/ROCm/TheRock/blob/main/RELEASES.md). vLLM lists Ryzen AI MAX / `gfx1151` as an AMD ROCm target, but this release needs the included CIRU source delta rather than a stock vLLM wheel.
+
+## Hardware and storage
+
+- x86-64 Strix Halo / Radeon 8060S (`gfx1151`)
+- 128 GB unified memory strongly recommended; smaller systems are not expected to fit the 77 GB checkpoint plus runtime and KV cache
+- at least 110 GB free disk for weights, the runtime environment, sources, build outputs, and caches
+- a current BIOS and Linux kernel with the Strix Halo KFD fixes
+
+AMD documents the required kernel and memory configuration in its [Strix Halo system-optimization guide](https://rocm.docs.amd.com/en/docs-7.2.0/how-to/system-optimization/strixhalo.html). For generic distributions, Linux 6.18.4 or newer contains the required upstream fixes. AMD also identifies current Fedora, Arch, Ubuntu OEM, and Ubuntu 26.04 paths with those fixes; check the live compatibility table instead of assuming that any kernel carrying an AMDGPU driver is sufficient.
+
+## 1. Verify the host driver
+
+After installing the current AMD-supported driver/kernel path for your distribution, add the user to the GPU access groups and reboot:
+
+```bash
+sudo usermod -a -G render,video "$LOGNAME"
+sudo reboot
+```
+
+Then verify:
+
+```bash
+uname -r
+groups
+ls -l /dev/kfd /dev/dri/renderD*
+```
+
+If `/dev/kfd` is absent, stop here. The Python wheels cannot replace a missing kernel driver.
+
+## 2. Configure unified-memory access
+
+Strix Halo maps system memory dynamically through GTT/TTM. AMD recommends keeping the fixed BIOS GPU reservation small and increasing the shared TTM limit instead. On CIRU's 128 GB host, the validated mapping limit was 124 GiB. That limit is not a permanent reservation, but the model can still leave very little free RAM during long-context requests.
+
+AMD's supported helper is:
+
+```bash
+sudo apt-get install -y pipx
+pipx ensurepath
+pipx install amd-debug-tools
+amd-ttm
+amd-ttm --set 124
+sudo reboot
+```
+
+The first line is the Ubuntu/Debian package name; Fedora and Arch users should install their distribution's `pipx` package instead. Use `124` only on a 128 GB machine. Close other memory-heavy applications before loading the model. A clean 60K+512 run peaked near 117 GB mapped GTT and left about 2 GiB of available RAM.
+
+## 3. Install distro host dependencies
+
+The helper supports Ubuntu/Debian, Fedora/RHEL-family systems, and Arch:
+
+```bash
+bash scripts/install-host-deps.sh
+```
+
+It installs a C/C++ toolchain, Git, `curl`, `pkg-config`, NUMA and DRM headers, and `uv`. `uv` then supplies the exact Python 3.12 interpreter, so the host distribution does not need to package `python3.12` itself.
+
+## 4. Build the pinned runtime
+
+From the root of this Hugging Face repository:
+
+```bash
+bash scripts/build-vllm-gfx1151.sh "$HOME/ciru-ling-runtime"
+```
+
+The script:
+
+- installs the pinned TheRock `device-gfx1151` ROCm/Torch transaction;
+- runs `rocm-sdk init` and `rocm-sdk test`;
+- reconstructs the exact CIRU vLLM commit from the included Git bundle;
+- builds vLLM native extensions for `gfx1151` only;
+- installs the pinned AMD Triton FlashAttention path; and
+- imports every required native vLLM extension.
+
+A successful build ends with:
+
+```text
+CIRU vLLM gfx1151 build: PASS
+```
+
+## 5. Download the unchanged checkpoint and serve
+
+```bash
+export PATH="$HOME/.local/bin:$PATH"
+bash scripts/download-model.sh "$HOME/models"
+
+VENV="$HOME/ciru-ling-runtime/.venv" \
+MODEL_PATH="$HOME/models/Ling-3.0-Flash-CIRU-int4-Strix-native" \
+bash scripts/run-256k.sh
+```
+
+In another shell:
+
+```bash
+curl -fsS http://127.0.0.1:18081/health
+curl -fsS http://127.0.0.1:18081/v1/models
+curl -fsS http://127.0.0.1:18081/v1/chat/completions \
+  -H 'Content-Type: application/json' \
+  -d '{"model":"Ling-3.0-Flash-CIRU-int4-Strix-native","messages":[{"role":"user","content":"Reply with exactly: STRIX READY"}],"temperature":0,"max_tokens":16}'
+```
+
+The native profile uses the validated `gpu_memory_utilization=0.75` envelope, native 262,144 context, K1 MTP, and at most five active requests. Extra requests wait in vLLM's scheduler.
+
+## Portability status
+
+| Platform | Status |
+|---|---|
+| NixOS, native Linux, 128 GB Strix Halo | validated build, load, API, correctness, and performance |
+| Ubuntu/Debian native Linux | documented, not yet clean-host validated by CIRU |
+| Fedora / Arch native Linux | dependency path documented; use a kernel carrying AMD's Strix Halo fixes |
+| Windows 11 through WSL2/ROCDXG | experimental; see [WINDOWS_WSL2.md](WINDOWS_WSL2.md) |
+| Native Windows Python | not supported for this vLLM release |
+
+When reporting a portability result, include the OS, kernel, GPU identity, `rocm-sdk test` result, build tail, model-load result, and one exact API response.

docs/VALIDATED_ENVIRONMENT.md

@@ -0,0 +1,35 @@
+# Validated environment
+
+## Hardware
+
+- AMD Radeon 8060S / Strix Halo
+- GPU target: `gfx1151`, Wave32
+- 128 GB unified memory
+- Linux TTM/GTT mapping limit: 124 GiB
+
+## Software
+
+- NixOS host
+- Python 3.12.12
+- Torch `2.13.0+rocm7.15.0a20260728`
+- HIP `7.15.0`
+- Triton `3.8.0+git4cff872c.rocm7.15.0a20260728`
+- AMD clang 23
+- torchvision `0.28.0+rocm7.15.0a20260728`
+- torchaudio `2.11.0.2+rocm7.15.0a20260728`
+- FlashAttention `2.8.3`, Dao-AILab commit `0e60e39473e8df549a20fb5353760f7a65b30e2d`, Python-only AMD Triton path
+- vLLM upstream base `d35eb6c44071ea806018841c490f0d2f3219c485`
+- CIRU runtime head `388d82de2abe418e73380ba250270d3abc62ac12`
+
+## Build facts
+
+- Fresh Python 3.12 environment.
+- Exact AMD multi-architecture wheel transaction installed in one resolver operation.
+- `rocm-sdk init` and all 27 SDK tests passed.
+- vLLM compiled for `gfx1151` only with 16 jobs.
+- Clean native build completed in 92 seconds on the validation host.
+- Imports passed for `vllm._C`, `vllm._C_stable_libtorch`, `vllm._moe_C_stable_libtorch`, `vllm._rocm_C`, and `vllm.cumem_allocator`.
+
+The Ubuntu/Debian, Fedora, Arch, and Windows/WSL build paths in this package are portability candidates and have not yet been clean-host validated by CIRU.
+
+The published native-256K launcher uses the separately validated `gpu_memory_utilization=0.75` long-context envelope. It exposed a 2,519,792-token KV pool and completed the frozen 60K+512 workload without faults, although system-memory margin was only about 2 GiB at peak.

docs/WINDOWS_WSL2.md

@@ -0,0 +1,87 @@
+# Windows 11 through WSL2 and AMD ROCDXG
+
+This runtime does not build as a native Windows vLLM application. Windows users must run the Linux build inside WSL2 and expose the Strix Halo GPU through AMD's ROCDXG layer.
+
+AMD now supports Strix Halo under WSL2, but CIRU has not yet validated this exact ROCm 7.15 nightly runtime under ROCDXG. Treat this as an experimental portability path and report results before calling it production-ready.
+
+## Requirements
+
+- Windows 11 on a 128 GB Ryzen AI MAX / MAX+ Strix Halo system
+- a current AMD Windows driver compatible with ROCDXG
+- current WSL2 with Ubuntu 24.04 or another distribution listed in AMD's live compatibility matrix
+- at least 110 GB free storage inside WSL or on a fast WSL virtual disk
+
+Use AMD's current [ROCm on Ryzen WSL guide](https://rocm.docs.amd.com/projects/radeon-ryzen/en/latest/docs/install/installryz/wsl/howto_wsl.html) and the [ROCDXG quickstart and compatibility matrix](https://github.com/ROCm/librocdxg). Do not follow older `roc4wsl` instructions.
+
+## 1. Install and size WSL2
+
+In an elevated PowerShell window:
+
+```powershell
+wsl --install -d Ubuntu-24.04
+wsl --update
+wsl --shutdown
+```
+
+For a 128 GB machine, create `%UserProfile%\.wslconfig` with a large WSL memory ceiling:
+
+```ini
+[wsl2]
+memory=120GB
+swap=16GB
+```
+
+Then run `wsl --shutdown` again. This allocation is intentionally large because the checkpoint is 77 GB before KV cache and runtime overhead. It leaves little room for Windows applications; close memory-heavy programs before serving. The 60K benchmark has not been validated under WSL.
+
+## 2. Install the AMD Windows driver and ROCDXG
+
+Follow AMD's live guide in this order:
+
+1. Install the compatible AMD Windows driver.
+2. Install the ROCm packages inside the WSL distribution using AMD's current quickstart.
+3. Install the matching ROCDXG release, either from its `.deb` package or source instructions.
+4. Restart WSL.
+
+Inside Ubuntu, verify that the bridge is active:
+
+```bash
+test -e /dev/dxg
+rocminfo | grep -A4 -B2 gfx1151
+```
+
+Do not proceed until `rocminfo` identifies `gfx1151`. Native-Linux `/dev/kfd` instructions do not apply to the WSL bridge.
+
+## 3. Clone and build the CIRU runtime
+
+Inside WSL:
+
+```bash
+git clone https://huggingface.co/jcbtc/Ling-3.0-Flash-CIRU-int4-Strix-native
+cd Ling-3.0-Flash-CIRU-int4-Strix-native
+
+bash scripts/install-host-deps.sh
+bash scripts/build-vllm-gfx1151.sh "$HOME/ciru-ling-runtime"
+```
+
+The build script detects WSL and refuses to continue when `/dev/dxg` is absent. It installs the release's pinned Linux ROCm/Torch userspace in its own virtual environment; do not install a separate PyTorch into that environment.
+
+## 4. Download and serve
+
+```bash
+export PATH="$HOME/.local/bin:$PATH"
+bash scripts/download-model.sh "$HOME/models"
+
+VENV="$HOME/ciru-ling-runtime/.venv" \
+MODEL_PATH="$HOME/models/Ling-3.0-Flash-CIRU-int4-Strix-native" \
+bash scripts/run-256k.sh
+```
+
+Verify `/health`, `/v1/models`, and an exact short completion using the commands in [UBUNTU_BUILD.md](UBUNTU_BUILD.md).
+
+## Known WSL boundaries
+
+- Native Windows ROCm PyTorch support does not imply native Windows vLLM support.
+- ROCDXG's supported ROCm/driver combinations change independently; use its live matrix.
+- AMD documents limited profiling support under WSL.
+- Very large WSL memory mappings are sensitive to `.wslconfig`, host pressure, and the Windows driver.
+- The published Linux performance numbers must not be presented as WSL performance until measured there.

runtime/COMMITS.txt

@@ -0,0 +1,15 @@
+c3f74efb866edba8892f18506b1c4febc157a799  [Bugfix][Kernel][ROCm] Fix Wave32 LDS overflow in top-k merge launch
+abefcfad484b85af6f57703c815a5f169807a3fd  research(profile): annotate Ling decoder families
+a73488d0243736797f75412ad23549cada767a47  Revert "research(profile): annotate Ling decoder families"
+c7e91b99ce98786dad42f830965759e1ff1f74ac  Reapply "research(profile): annotate Ling decoder families"
+d8481aec96d9a3caca7323f5d6b5bcef3f0fb8ed  Revert "Reapply "research(profile): annotate Ling decoder families""
+8457fad14b028e1c443cb4c5998df377873e3d95  Enable native merge attention states on ROCm
+8743a17f147d899d670af9b808d5bf93a5bf26fb  Revert "Enable native merge attention states on ROCm"
+7deddd0a87a94637e4fbe2d309a484ab9a04e33f  Add opt-in safe ROCm attention merge
+3715c596c1a0df9d42d133a61d14dbed5105ab0a  Handle MLA LSE layouts in safe ROCm merge
+f6773b23215df5aa2a8de5050bd2f9b6fa53b727  ROCm: route uniform MLA verify blocks through causal decode
+15a4190ddf0a363e08d518345a8f3e23e8bc8bd9  Optimize decode WNA16 expert assignment
+abfff126e0698d98f1872f87f501d36a19382240  Optimize single-token MoE reduction
+928e725a9dbcdbea0bf6c914c8b3fe93f08b0d3f  Optimize K2 verification MoE support kernels
+fabbc7e928dfa64c29dcdf0f795b05246555f4ae  Extend MoE support fast path to K1 verification
+388d82de2abe418e73380ba250270d3abc62ac12  ROCm: optimize exact small WNA16 SiLU activation

runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch

@@ -0,0 +1,550 @@
+diff --git a/csrc/libtorch_stable/sampler.cu b/csrc/libtorch_stable/sampler.cu
+index 519e213281cd07a9a4cc81f7bc5fcb843277528d..847d8284afc992c3f75e2d65f240afe3d3d0e0e3 100644
+--- a/csrc/libtorch_stable/sampler.cu
++++ b/csrc/libtorch_stable/sampler.cu
+@@ -711,7 +711,12 @@ void top_k_per_row_decode(const torch::stable::Tensor& logits, int64_t next_n,
+             static_cast<int>(next_n), seqLensIs2D,
+             outLogitsAux.mutable_data_ptr<float>());
+ 
++#ifndef USE_ROCM
+     constexpr int kNumThreadsPerBlockMerge = 1024;
++#else
++    // 1024 threads exceed the 64 KB LDS limit on Wave32 GPUs.
++    constexpr int kNumThreadsPerBlockMerge = 512;
++#endif
+     vllm::topKPerRowDecode<kNumThreadsPerBlockMerge, true, false, true>
+         <<<numRows, kNumThreadsPerBlockMerge, topK * sizeof(int32_t), stream>>>(
+             outLogitsAux.const_data_ptr<float>(), seqLens.const_data_ptr<int>(),
+diff --git a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py b/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
+index 9d93dfe54ac1a81d5c1f512bc266b9732506e31a..dd96f6da4cc3e3655bbab58f64ceed0882f2e7c0 100644
+--- a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
++++ b/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
+@@ -27,6 +27,7 @@ from vllm.model_executor.layers.fused_moe.fused_moe import (
+ from vllm.model_executor.layers.fused_moe.moe_align_block_size import (
+     moe_align_block_size,
+ )
++from vllm.model_executor.layers.fused_moe.moe_fused_mul_sum import moe_sum_decode
+ from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
+     TopKWeightAndReduceNoOP,
+ )
+@@ -55,10 +56,100 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
+     kInt8StaticTensorSym,
+ )
+ from vllm.platforms import current_platform
+-from vllm.triton_utils import tl
++from vllm.triton_utils import tl, triton
+ from vllm.utils.multi_stream_utils import maybe_execute_in_parallel
+ 
+ 
++def _use_wna16_naive_block_assignment(
++    *,
++    num_tokens: int,
++    top_k: int,
++    num_experts: int,
++    global_num_experts: int,
++    block_size_m: int,
++    expert_map: torch.Tensor | None,
++    use_int4_w4a16: bool,
++    block_shape: list[int] | None,
++) -> bool:
++    """Select the narrowly validated gfx1151 WNA16 small-token fast path."""
++    return (
++        num_tokens in (1, 2, 3)
++        and top_k == 8
++        and num_experts == 512
++        and global_num_experts == 512
++        and block_size_m == 16
++        and expert_map is None
++        and use_int4_w4a16
++        and block_shape == [0, 32]
++    )
++
++
++_WNA16_SILU_OUTPUT_SIZE = 768
++_WNA16_SILU_BLOCK_SIZE = 128
++
++
++@triton.jit
++def _wna16_exact_silu_and_mul_kernel(
++    input_ptr,
++    output_ptr,
++    OUTPUT_SIZE: tl.constexpr,
++    BLOCK_SIZE: tl.constexpr,
++):
++    row = tl.program_id(0)
++    block = tl.program_id(1)
++    offsets = block * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
++    mask = offsets < OUTPUT_SIZE
++    row_input = input_ptr + row * (2 * OUTPUT_SIZE)
++    gate = tl.load(row_input + offsets, mask=mask, other=0.0).to(tl.float32)
++    up = tl.load(row_input + OUTPUT_SIZE + offsets, mask=mask, other=0.0).to(
++        tl.float32
++    )
++
++    # Match the native packed HIP kernel's intermediate BF16 rounding.
++    activated = (gate * tl.sigmoid(gate)).to(tl.bfloat16).to(tl.float32)
++    tl.store(
++        output_ptr + row * OUTPUT_SIZE + offsets,
++        activated * up,
++        mask=mask,
++    )
++
++
++def _use_wna16_exact_silu_and_mul(
++    output: torch.Tensor, input: torch.Tensor
++) -> bool:
++    if torch.compiler.is_compiling() or not current_platform.is_rocm():
++        return False
++    from vllm.platforms.rocm import on_gfx1151
++
++    return (
++        on_gfx1151()
++        and input.is_cuda
++        and output.is_cuda
++        and input.dtype == torch.bfloat16
++        and output.dtype == torch.bfloat16
++        and input.is_contiguous()
++        and output.is_contiguous()
++        and input.shape in ((8, 1536), (16, 1536))
++        and output.shape == (input.shape[0], _WNA16_SILU_OUTPUT_SIZE)
++    )
++
++
++def _invoke_wna16_exact_silu_and_mul(
++    output: torch.Tensor, input: torch.Tensor
++) -> None:
++    grid = (
++        input.shape[0],
++        triton.cdiv(_WNA16_SILU_OUTPUT_SIZE, _WNA16_SILU_BLOCK_SIZE),
++    )
++    _wna16_exact_silu_and_mul_kernel[grid](
++        input,
++        output,
++        OUTPUT_SIZE=_WNA16_SILU_OUTPUT_SIZE,
++        BLOCK_SIZE=_WNA16_SILU_BLOCK_SIZE,
++        num_warps=1,
++    )
++
++
+ class TritonExperts(LoRAExpertsMixin, mk.FusedMoEExpertsModular):
+     """Triton-based fused MoE expert implementation."""
+ 
+@@ -603,6 +694,33 @@ class TritonWNA16Experts(TritonExperts):
+             or moe_parallel_config.use_fi_nvl_one_sided_kernels
+         )
+ 
++    def activation(
++        self,
++        activation: MoEActivation,
++        output: torch.Tensor,
++        input: torch.Tensor,
++        **kwargs,
++    ) -> None:
++        # Keep compiled MLA-side MoE activations on Inductor's fused pointwise
++        # route. The exact KDA-side eager shapes otherwise pay the high fixed
++        # latency of the generic HIP activation kernel on gfx1151.
++        if (
++            activation == MoEActivation.SILU
++            and self.activation_config.clamp_limit is None
++            and self.quant_config.use_int4_w4a16
++            and self.block_shape == [0, 32]
++            and _use_wna16_exact_silu_and_mul(output, input)
++        ):
++            _invoke_wna16_exact_silu_and_mul(output, input)
++            return
++
++        super().activation(
++            activation,
++            output,
++            input,
++            **kwargs,
++        )
++
+     def apply(
+         self,
+         output: torch.Tensor,
+@@ -681,9 +799,35 @@ class TritonWNA16Experts(TritonExperts):
+         )
+         intermediate_cache3 = _resize_cache(workspace2, (num_tokens, top_k_num, K))
+ 
+-        sorted_token_ids, expert_ids, num_tokens_post_padded = moe_align_block_size(
+-            topk_ids, config["BLOCK_SIZE_M"], global_num_experts, expert_map
++        use_naive_block_assignment = _use_wna16_naive_block_assignment(
++            num_tokens=num_tokens,
++            top_k=top_k_num,
++            num_experts=E,
++            global_num_experts=global_num_experts,
++            block_size_m=config["BLOCK_SIZE_M"],
++            expert_map=expert_map,
++            use_int4_w4a16=self.quant_config.use_int4_w4a16,
++            block_shape=self.block_shape,
+         )
++        use_naive_block_assignment = (
++            use_naive_block_assignment
++            and topk_ids.shape == (num_tokens, top_k_num)
++            and topk_ids.is_contiguous()
++        )
++        if use_naive_block_assignment:
++            sorted_token_ids = None
++            expert_ids = topk_ids.view(-1)
++            # Compile-time dead in the naive kernel specialization.
++            num_tokens_post_padded = expert_ids
++        else:
++            sorted_token_ids, expert_ids, num_tokens_post_padded = (
++                moe_align_block_size(
++                    topk_ids,
++                    config["BLOCK_SIZE_M"],
++                    global_num_experts,
++                    expert_map,
++                )
++            )
+ 
+         invoke_fused_moe_wna16_triton_kernel(
+             hidden_states,
+@@ -737,5 +881,20 @@ class TritonWNA16Experts(TritonExperts):
+             block_shape=self.block_shape,
+         )
+ 
+-        # separate function is required for MoE + LoRA
+-        self.moe_sum(intermediate_cache3, output)
++        # The native HIP top-k sum has very high fixed latency on gfx1151 for
++        # the exact target-decode and K1/K2 verification shapes. Keep every
++        # adjacent shape on the native implementation.
++        if (
++            use_naive_block_assignment
++            and self._lora_context is None
++            and intermediate_cache3.shape
++            in ((1, 8, 2560), (2, 8, 2560), (3, 8, 2560))
++            and output.shape in ((1, 2560), (2, 2560), (3, 2560))
++            and output.shape[0] == intermediate_cache3.shape[0]
++            and intermediate_cache3.is_contiguous()
++            and output.is_contiguous()
++        ):
++            moe_sum_decode(intermediate_cache3, output)
++        else:
++            # separate function is required for MoE + LoRA
++            self.moe_sum(intermediate_cache3, output)
+diff --git a/vllm/model_executor/layers/fused_moe/fused_moe.py b/vllm/model_executor/layers/fused_moe/fused_moe.py
+index be4930052a9a1c883679f7e7415853b2f9c9f667..034c963b7cee65d253e9c27517463c2e59b2d1c0 100644
+--- a/vllm/model_executor/layers/fused_moe/fused_moe.py
++++ b/vllm/model_executor/layers/fused_moe/fused_moe.py
+@@ -97,6 +97,7 @@ def fused_moe_kernel_gptq_awq(
+     stride_bzn,
+     block_k_diviable: tl.constexpr,
+     group_size: tl.constexpr,
++    naive_block_assignment: tl.constexpr,
+     # Meta-parameters
+     BLOCK_SIZE_M: tl.constexpr,
+     BLOCK_SIZE_N: tl.constexpr,
+@@ -155,12 +156,20 @@ def fused_moe_kernel_gptq_awq(
+     # and accumulate
+     # `a_ptrs` is a block of [BLOCK_SIZE_M, BLOCK_SIZE_K] pointers
+     # `b_ptrs` is a block of [BLOCK_SIZE_K, BLOCK_SIZE_N] pointers
+-    num_tokens_post_padded = tl.load(num_tokens_post_padded_ptr)
+-    if pid_m * BLOCK_SIZE_M >= num_tokens_post_padded:
+-        return
+-    offs_token_id = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
++    offs = tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
++    if not naive_block_assignment:
++        num_tokens_post_padded = tl.load(num_tokens_post_padded_ptr)
++        if pid_m * BLOCK_SIZE_M >= num_tokens_post_padded:
++            return
++        offs_token_id = pid_m * BLOCK_SIZE_M + offs
++        offs_token = tl.load(sorted_token_ids_ptr + offs_token_id)
++    else:
++        # Decode-only fast path: one routed token per program.  The remaining
++        # lanes are masked so no alignment, sorting, or padded-count load is
++        # needed.
++        offs_token = tl.where(offs == 0, pid_m, num_valid_tokens)
+     # Cast to int64 to prevent overflow in stride*offset products
+-    offs_token = tl.load(sorted_token_ids_ptr + offs_token_id).to(tl.int64)
++    offs_token = offs_token.to(tl.int64)
+     token_mask = offs_token < num_valid_tokens
+ 
+     off_experts = tl.load(expert_ids_ptr + pid_m).to(tl.int64)
+@@ -677,7 +686,7 @@ def invoke_fused_moe_wna16_triton_kernel(
+     B_scale: torch.Tensor | None,
+     B_zp: torch.Tensor | None,
+     topk_weights: torch.Tensor | None,
+-    sorted_token_ids: torch.Tensor,
++    sorted_token_ids: torch.Tensor | None,
+     expert_ids: torch.Tensor,
+     num_tokens_post_padded: torch.Tensor,
+     mul_routed_weight: bool,
+@@ -695,8 +704,11 @@ def invoke_fused_moe_wna16_triton_kernel(
+     M = A.size(0)
+     num_tokens = M * top_k
+ 
+-    EM = sorted_token_ids.size(0)
+-    if A.size(0) < config["BLOCK_SIZE_M"]:
++    if sorted_token_ids is None:
++        EM = num_tokens * config["BLOCK_SIZE_M"]
++    else:
++        EM = sorted_token_ids.size(0)
++    if sorted_token_ids is not None and A.size(0) < config["BLOCK_SIZE_M"]:
+         # optimize for small batch_size.
+         # We assume that top_ids of each token is unique,
+         # so num_valid_experts <= batch_size <= BLOCK_SIZE_M,
+@@ -756,6 +768,7 @@ def invoke_fused_moe_wna16_triton_kernel(
+         has_zp=B_zp is not None,
+         use_int4_w4a16=use_int4_w4a16,
+         use_int8_w8a16=use_int8_w8a16,
++        naive_block_assignment=(sorted_token_ids is None),
+         **config,
+     )
+ 
+diff --git a/vllm/model_executor/layers/fused_moe/moe_fused_mul_sum.py b/vllm/model_executor/layers/fused_moe/moe_fused_mul_sum.py
+index 768f41db854e0d532d5faeb3379eab1ab8ad5ef4..827f19fdd2117821861e89cb29571875d56257c0 100644
+--- a/vllm/model_executor/layers/fused_moe/moe_fused_mul_sum.py
++++ b/vllm/model_executor/layers/fused_moe/moe_fused_mul_sum.py
+@@ -7,6 +7,55 @@ from vllm.platforms import current_platform
+ from vllm.triton_utils import tl, triton
+ 
+ 
++@triton.jit
++def moe_sum_decode_kernel(
++    input_ptr,
++    output_ptr,
++    size: tl.constexpr,
++    top_k: tl.constexpr,
++    BLOCK_K: tl.constexpr,
++):
++    """Sum a small token batch's routed outputs in fixed route order."""
++    token = tl.program_id(1)
++    offsets = tl.program_id(0) * BLOCK_K + tl.arange(0, BLOCK_K)
++    mask = offsets < size
++    accumulator = tl.zeros((BLOCK_K,), dtype=tl.float32)
++    for route in tl.static_range(top_k):
++        value = tl.load(
++            input_ptr + (token * top_k + route) * size + offsets,
++            mask=mask,
++            other=0.0,
++        )
++        accumulator += value.to(tl.float32)
++    tl.store(output_ptr + token * size + offsets, accumulator, mask=mask)
++
++
++def moe_sum_decode(
++    inputs: torch.Tensor,
++    output: torch.Tensor,
++) -> torch.Tensor:
++    """Low-launch-overhead sum for exact Ling WNA16 decode/verify shapes."""
++    assert inputs.shape in ((1, 8, 2560), (2, 8, 2560), (3, 8, 2560))
++    assert output.shape in ((1, 2560), (2, 2560), (3, 2560))
++    assert output.shape[0] == inputs.shape[0]
++    assert inputs.dtype == torch.bfloat16
++    assert output.dtype == torch.bfloat16
++    assert inputs.is_contiguous()
++    assert output.is_contiguous()
++
++    if not isinstance(inputs, FakeTensor):
++        moe_sum_decode_kernel[(5, inputs.shape[0])](
++            inputs,
++            output,
++            size=2560,
++            top_k=8,
++            BLOCK_K=512,
++            num_warps=4,
++            num_stages=2,
++        )
++    return output
++
++
+ @triton.jit
+ def moe_fused_mul_sum_kernel(
+     inputs_ptr,
+diff --git a/vllm/v1/attention/backends/mla/triton_mla.py b/vllm/v1/attention/backends/mla/triton_mla.py
+index 71a38e95d0ad49fd6a0511cf6613b5cab1f62af0..641f68b166ff4b4269eb7ec6f07b5cead3849a38 100644
+--- a/vllm/v1/attention/backends/mla/triton_mla.py
++++ b/vllm/v1/attention/backends/mla/triton_mla.py
+@@ -13,6 +13,7 @@ from vllm.model_executor.layers.attention.mla_attention import (
+     MLACommonImpl,
+     MLACommonMetadata,
+     MLACommonMetadataBuilder,
++    QueryLenSupport,
+ )
+ from vllm.platforms import current_platform
+ from vllm.platforms.interface import DeviceCapability
+@@ -51,6 +52,9 @@ class TritonMLAMetadataBuilder(MLACommonMetadataBuilder[MLACommonMetadata]):
+     _cudagraph_support: ClassVar[AttentionCGSupport] = (
+         AttentionCGSupport.UNIFORM_SINGLE_TOKEN_DECODE
+     )
++    # Uniform multi-token groups (for example MTP verification) can use the
++    # decode kernel when each query row is mapped to its causal KV prefix.
++    query_len_support: ClassVar[QueryLenSupport] = QueryLenSupport.UNIFORM
+     # Non-causal DSpark block is flattened to one decode row per query token in
+     # forward_mqa, so no intra-block causal masking is required.
+     supports_non_causal_multi_token_decode: ClassVar[bool] = True
+@@ -74,12 +78,10 @@ class TritonMLAMetadataBuilder(MLACommonMetadataBuilder[MLACommonMetadata]):
+         """
+         if not is_workspace_manager_initialized():
+             return
+-        # Decode reorder threshold is 1, so decode tokens <= max_num_seqs.
++        # Uniform multi-token decode flattens each request's block to query_len
++        # rows. Reserve the largest admitted block before the workspace locks.
+         B = self.vllm_config.scheduler_config.max_num_seqs
+-        # Non-causal DSpark draft flattens each request's block to query_len
+-        # decode rows; cover max_num_seqs * block_len rows.
+-        if getattr(self, "non_causal_multi_token_decode", False):
+-            B *= self.reorder_batch_threshold
++        B *= self.reorder_batch_threshold or 1
+         # DCP all-gathers the query heads before forward_mqa.
+         q_num_heads = self.num_heads * self.dcp_world_size
+         max_splits = _compute_num_kv_splits(
+@@ -285,14 +287,34 @@ class TritonMLAImpl(MLACommonImpl[MLACommonMetadata]):
+ 
+         block_table = attn_metadata.decode.block_table
+         seq_lens = attn_metadata.decode.seq_lens
+-        if not attn_metadata.causal:
+-            # Non-causal DSpark block: flatten to one decode row per query token.
+-            # Each row attends to the same committed KV prefix (per-row seq_lens)
+-            # and never to sibling block tokens = non-causal block semantics.
+-            # Mirrors FlashInferMLA's non-causal path.
+-            query_len = attn_metadata.num_decode_tokens // attn_metadata.num_decodes
+-            if query_len > 1:
+-                block_table = block_table.repeat_interleave(query_len, dim=0)
++        num_decodes = attn_metadata.num_decodes
++        num_decode_tokens = attn_metadata.num_decode_tokens
++        assert num_decodes > 0
++        assert num_decode_tokens == B
++        assert num_decode_tokens % num_decodes == 0
++
++        query_len = num_decode_tokens // num_decodes
++        if query_len > 1:
++            block_table = block_table.repeat_interleave(query_len, dim=0)
++            if attn_metadata.causal:
++                logger.info_once(
++                    "Triton MLA A1 causal spec-as-decode route active: "
++                    "query_len=%d",
++                    query_len,
++                )
++                # seq_lens includes the whole verify block. Row t may attend
++                # only through its own position: context + t + 1.
++                token_offsets = torch.arange(
++                    query_len,
++                    device=seq_lens.device,
++                    dtype=seq_lens.dtype,
++                )
++                seq_lens = (
++                    seq_lens.unsqueeze(1) - (query_len - 1) + token_offsets
++                ).clamp_min(0).flatten()
++            else:
++                # Non-causal DSpark rows share the committed KV prefix and do
++                # not attend to sibling block tokens.
+                 seq_lens = seq_lens.repeat_interleave(query_len)
+ 
+         # Run MQA — always pass layer scales. When KV cache is
+diff --git a/vllm/v1/attention/ops/merge_attn_states.py b/vllm/v1/attention/ops/merge_attn_states.py
+index 20c7503e9d0518a1c9459c6a599708f94a2bae62..cc860e1ab26ca78bf747a00b160314c4d740c5b1 100644
+--- a/vllm/v1/attention/ops/merge_attn_states.py
++++ b/vllm/v1/attention/ops/merge_attn_states.py
+@@ -1,11 +1,87 @@
+ # SPDX-License-Identifier: Apache-2.0
+ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+ 
++import os
++
+ import torch
+ 
+ from vllm.platforms import current_platform
+ 
+ 
++def _merge_attn_states_torch(
++    output: torch.Tensor,
++    prefix_output: torch.Tensor,
++    prefix_lse: torch.Tensor,
++    suffix_output: torch.Tensor,
++    suffix_lse: torch.Tensor,
++    output_lse: torch.Tensor | None,
++    prefill_tokens_with_context: int | None,
++) -> None:
++    """Safe ROCm fallback for the Triton merge kernel.
++
++    This follows the reference implementation in
++    tests/kernels/attention/test_merge_attn_states.py. It is intentionally
++    opt-in because it trades a small amount of prefill performance for avoiding
++    gfx1151 HSA memory faults in merge_attn_states_kernel.
++    """
++    num_context_tokens = (
++        output.shape[0]
++        if prefill_tokens_with_context is None
++        else prefill_tokens_with_context
++    )
++    tokens_first_lse = prefix_lse.shape[:2] == output.shape[:2]
++    if num_context_tokens:
++        if tokens_first_lse:
++            p_lse = prefix_lse[:num_context_tokens, :].clone()
++            s_lse = suffix_lse[:num_context_tokens, :].clone()
++        else:
++            p_lse = prefix_lse[:, :num_context_tokens].clone()
++            s_lse = suffix_lse[:, :num_context_tokens].clone()
++        p_lse[p_lse == torch.inf] = -torch.inf
++        s_lse[s_lse == torch.inf] = -torch.inf
++        max_lse = torch.maximum(p_lse, s_lse)
++        p_exp = torch.nan_to_num(torch.exp(p_lse - max_lse), nan=0.0)
++        s_exp = torch.nan_to_num(torch.exp(s_lse - max_lse), nan=0.0)
++        denom = p_exp + s_exp
++        p_scale = torch.where(denom > 0, p_exp / denom, 0.0)
++        s_scale = torch.where(denom > 0, s_exp / denom, 0.0)
++        p_scale_output_layout = (
++            p_scale.unsqueeze(-1)
++            if tokens_first_lse
++            else p_scale.transpose(0, 1).unsqueeze(-1)
++        )
++        s_scale_output_layout = (
++            s_scale.unsqueeze(-1)
++            if tokens_first_lse
++            else s_scale.transpose(0, 1).unsqueeze(-1)
++        )
++        merged = (
++            prefix_output[:num_context_tokens]
++            * p_scale_output_layout
++            + suffix_output[:num_context_tokens]
++            * s_scale_output_layout
++        )
++        output[:num_context_tokens].copy_(merged.to(output.dtype))
++        if output_lse is not None:
++            merged_lse = torch.logaddexp(p_lse, s_lse)
++            if tokens_first_lse:
++                output_lse[:num_context_tokens, :].copy_(merged_lse)
++            else:
++                output_lse[:, :num_context_tokens].copy_(merged_lse)
++
++    if num_context_tokens < output.shape[0]:
++        output[num_context_tokens:].copy_(suffix_output[num_context_tokens:])
++        if output_lse is not None:
++            if tokens_first_lse:
++                output_lse[num_context_tokens:, :].copy_(
++                    suffix_lse[num_context_tokens:, :]
++                )
++            else:
++                output_lse[:, num_context_tokens:].copy_(
++                    suffix_lse[:, num_context_tokens:]
++                )
++
++
+ def merge_attn_states(
+     output: torch.Tensor,
+     prefix_output: torch.Tensor,
+@@ -46,6 +122,21 @@ def merge_attn_states(
+             When provided, output must be FP8 dtype.
+     """
+ 
++    if (
++        current_platform.is_rocm()
++        and os.getenv("VLLM_ROCM_SAFE_MERGE_ATTN_STATES", "0") == "1"
++        and output_scale is None
++    ):
++        return _merge_attn_states_torch(
++            output,
++            prefix_output,
++            prefix_lse,
++            suffix_output,
++            suffix_lse,
++            output_lse,
++            prefill_tokens_with_context,
++        )
++
+     # Both the CUDA and Triton kernels derive the suffix head stride from
+     # prefix_output, so suffix_output must share the same head stride.
+     assert prefix_output.stride(1) == suffix_output.stride(1), (

runtime/SHA256SUMS.runtime

@@ -0,0 +1,3 @@
+15e48a8df3b65a73c81cf73c820ba12b4ab9474697d6b3f876c26c0b0cd3d455  COMMITS.txt
+6b03af3a1685a1969ff96870ea776537f5f67fd3674fc353880bd088bcef9a9c  Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
+8f307e9f2f16207856ca0b064a765a73bbebc029ad0b0fed61e5975a17e07ab6  Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch

scripts/build-vllm-gfx1151.sh

@@ -0,0 +1,155 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+install_root="${1:-$PWD/ciru-ling-runtime}"
+script_dir="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+package_root="$(cd -- "$script_dir/.." && pwd)"
+venv="$install_root/.venv"
+source_dir="$install_root/vllm"
+index=https://rocm.nightlies.amd.com/whl-multi-arch/
+base=d35eb6c44071ea806018841c490f0d2f3219c485
+head=388d82de2abe418e73380ba250270d3abc62ac12
+bundle="$package_root/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle"
+flash_source="$install_root/flash-attention"
+flash_commit=0e60e39473e8df549a20fb5353760f7a65b30e2d
+
+command -v git >/dev/null
+command -v uv >/dev/null || {
+  echo "Missing uv. Run scripts/install-host-deps.sh first." >&2
+  exit 2
+}
+
+if [ "$(uname -s)" != Linux ]; then
+  echo "This vLLM/ROCm build requires Linux. On Windows, use WSL2 and docs/WINDOWS_WSL2.md." >&2
+  exit 2
+fi
+
+if grep -qi microsoft /proc/sys/kernel/osrelease 2>/dev/null; then
+  test -e /dev/dxg || {
+    echo "WSL2 detected but /dev/dxg is missing. Install AMD ROCDXG before building." >&2
+    exit 2
+  }
+else
+  test -e /dev/kfd || {
+    echo "Missing /dev/kfd. Install/enable the AMDGPU KFD driver for gfx1151 first." >&2
+    exit 2
+  }
+  test -d /dev/dri || {
+    echo "Missing /dev/dri. The AMDGPU render device is not available." >&2
+    exit 2
+  }
+fi
+
+mkdir -p "$install_root"
+if [ ! -x "$venv/bin/python" ]; then
+  UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv python install 3.12
+  UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv venv --python 3.12 "$venv"
+fi
+
+UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" \
+  --index-url "$index" \
+  'torch[device-gfx1151]==2.13.0+rocm7.15.0a20260728' \
+  'torchvision[device-gfx1151]==0.28.0+rocm7.15.0a20260728' \
+  'torchaudio==2.11.0.2+rocm7.15.0a20260728' \
+  'triton==3.8.0+git4cff872c.rocm7.15.0a20260728' \
+  'rocm[libraries,devel,device-gfx1151]==7.15.0a20260728'
+
+"$venv/bin/rocm-sdk" init
+"$venv/bin/rocm-sdk" test
+
+if [ ! -d "$source_dir/.git" ]; then
+  git clone https://github.com/vllm-project/vllm.git "$source_dir"
+fi
+
+git -C "$source_dir" fetch origin "$base"
+git -C "$source_dir" fetch "$bundle" \
+  refs/heads/codex/ling3-gfx1151-rocm715:refs/heads/ciru/ling3-gfx1151-rocm715
+git -C "$source_dir" checkout ciru/ling3-gfx1151-rocm715
+test "$(git -C "$source_dir" rev-parse HEAD)" = "$head"
+
+site="$($venv/bin/python - <<'PY'
+import site
+print(site.getsitepackages()[0])
+PY
+)"
+sdk="$site/_rocm_sdk_devel"
+core="$site/_rocm_sdk_core"
+test -x "$sdk/bin/hipcc"
+
+# Match the compatibility adjustment used by vLLM's ROCm SDK-wheel images.
+rocm_sdk_init="$site/rocm_sdk/__init__.py"
+"$venv/bin/python" - "$rocm_sdk_init" <<'PY'
+from pathlib import Path
+import sys
+p = Path(sys.argv[1])
+s = p.read_text()
+s = s.replace("rtld_global: bool = True", "rtld_global: bool = False")
+p.write_text(s)
+PY
+
+hsakmt_targets="$sdk/lib/cmake/hsakmt/hsakmtTargets.cmake"
+if [ -f "$hsakmt_targets" ]; then
+  sed -i 's|-L/__w/[^;"]*|-L${_IMPORT_PREFIX}/lib/rocm_sysdeps/lib|g' \
+    "$hsakmt_targets"
+  sed -i 's|/usr/lib64/libc.so;||g' "$hsakmt_targets"
+fi
+
+export PATH="$sdk/bin:$sdk/llvm/bin:$venv/bin:$PATH"
+export LD_LIBRARY_PATH="$sdk/lib:$core/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
+export PYTHONPATH="$core/share/amd_smi${PYTHONPATH:+:$PYTHONPATH}"
+export HIP_DEVICE_LIB_PATH="$core/lib/llvm/amdgcn/bitcode"
+export VLLM_TARGET_DEVICE=rocm
+export PYTORCH_ROCM_ARCH=gfx1151
+export ROCM_PATH="$sdk"
+export ROCM_HOME="$sdk"
+export ROCM_SOURCE_DIR="$sdk"
+export HIP_PATH="$sdk"
+export CMAKE_PREFIX_PATH="$sdk/lib/cmake:$site/torch/share/cmake"
+export HIP_FORCE_DEV_KERNARG=1
+export MAX_JOBS="${MAX_JOBS:-16}"
+export CMAKE_BUILD_PARALLEL_LEVEL="$MAX_JOBS"
+export CMAKE_ARGS="-DROCM_PATH=$sdk -DCMAKE_HIP_ARCHITECTURES=gfx1151 ${CMAKE_ARGS:-}"
+
+UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" \
+  'cmake>=3.26.1,<4' ninja 'packaging>=24.2' \
+  'setuptools>=77.0.3,<80' 'setuptools-scm>=8' \
+  'setuptools-rust>=1.9' wheel 'jinja2>=3.1.6'
+UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" "$core/share/amd_smi"
+UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" -r "$source_dir/requirements/rocm.txt"
+
+cd "$source_dir"
+UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" --no-build-isolation --no-deps -e .
+
+if [ ! -d "$flash_source/.git" ]; then
+  git clone https://github.com/Dao-AILab/flash-attention.git "$flash_source"
+fi
+git -C "$flash_source" fetch origin "$flash_commit"
+git -C "$flash_source" checkout --detach "$flash_commit"
+FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE \
+  UV_CACHE_DIR="${UV_CACHE_DIR:-$install_root/uv-cache}" uv pip install \
+  --python "$venv/bin/python" --no-build-isolation --no-deps "$flash_source"
+
+"$venv/bin/python" - <<'PY'
+import importlib
+import importlib.metadata
+import torch
+import triton
+assert torch.__version__ == "2.13.0+rocm7.15.0a20260728"
+assert torch.version.hip == "7.15.0"
+assert triton.__version__ == "3.8.0"
+assert importlib.metadata.version("flash_attn") == "2.8.3"
+for name in (
+    "vllm._C",
+    "vllm._C_stable_libtorch",
+    "vllm._moe_C_stable_libtorch",
+    "vllm._rocm_C",
+    "vllm.cumem_allocator",
+):
+    importlib.import_module(name)
+print("CIRU vLLM gfx1151 build: PASS")
+PY

scripts/download-model.sh

@@ -0,0 +1,23 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+destination_root="${1:-$PWD/models}"
+release_name=Ling-3.0-Flash-CIRU-int4-Strix-native
+revision=ca3ea63b0255d212c4fe6020db9e0a51ce136006
+destination="$destination_root/$release_name"
+
+if command -v hf >/dev/null 2>&1; then
+  hf_cmd=(hf)
+elif command -v uvx >/dev/null 2>&1; then
+  hf_cmd=(uvx --from 'huggingface_hub[cli]' hf)
+else
+  echo "Missing Hugging Face CLI and uvx. Run scripts/install-host-deps.sh first." >&2
+  exit 2
+fi
+
+mkdir -p "$destination_root"
+"${hf_cmd[@]}" download inclusionAI/Ling-3.0-flash-int4 \
+  --revision "$revision" \
+  --local-dir "$destination"
+
+printf 'model_path=%s\nsource_revision=%s\n' "$destination" "$revision"

scripts/install-host-deps.sh

@@ -0,0 +1,47 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+if [ "$(uname -s)" != Linux ]; then
+  echo "Run this inside native Linux or WSL2, not native Windows." >&2
+  exit 2
+fi
+
+if [ ! -r /etc/os-release ]; then
+  echo "Cannot identify the Linux distribution: /etc/os-release is missing." >&2
+  exit 2
+fi
+
+# shellcheck disable=SC1091
+source /etc/os-release
+distro="${ID:-unknown} ${ID_LIKE:-}"
+
+if [[ "$distro" =~ (ubuntu|debian) ]]; then
+  sudo apt-get update
+  sudo apt-get install -y \
+    build-essential git curl ca-certificates pkg-config \
+    libnuma-dev libdrm-dev python3-dev
+elif [[ "$distro" =~ (fedora|rhel|centos) ]]; then
+  sudo dnf install -y \
+    gcc gcc-c++ make git curl ca-certificates pkgconf-pkg-config \
+    numactl-devel libdrm-devel python3-devel
+elif [[ "$distro" =~ (arch) ]]; then
+  sudo pacman -Syu --needed --noconfirm \
+    base-devel git curl ca-certificates pkgconf numactl libdrm python
+else
+  echo "Unsupported automatic dependency install for: $distro" >&2
+  echo "Install a C/C++ toolchain, git, curl, pkg-config, NUMA headers, libdrm headers, and Python 3; then install uv." >&2
+  exit 2
+fi
+
+if ! command -v uv >/dev/null 2>&1; then
+  curl -LsSf https://astral.sh/uv/install.sh | sh
+  export PATH="$HOME/.local/bin:$PATH"
+fi
+
+command -v uv >/dev/null || {
+  echo "uv was installed under ~/.local/bin; add that directory to PATH and rerun." >&2
+  exit 2
+}
+
+printf 'Host dependencies ready on %s.\n' "${PRETTY_NAME:-$distro}"
+printf 'Next: bash scripts/build-vllm-gfx1151.sh "$HOME/ciru-ling-runtime"\n'

scripts/run-1m-yarn-experimental.sh

@@ -0,0 +1,10 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+script_dir="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+package_root="$(cd -- "$script_dir/.." && pwd)"
+# shellcheck disable=SC1091
+set -a
+source "$package_root/config/experimental-yarn-1m.env"
+set +a
+exec bash "$script_dir/run-common.sh"

scripts/run-256k.sh

@@ -0,0 +1,10 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+script_dir="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+package_root="$(cd -- "$script_dir/.." && pwd)"
+# shellcheck disable=SC1091
+set -a
+source "$package_root/config/native-256k.env"
+set +a
+exec bash "$script_dir/run-common.sh"

scripts/run-common.sh

@@ -0,0 +1,77 @@
+#!/usr/bin/env bash
+set -Eeuo pipefail
+
+release_name=Ling-3.0-Flash-CIRU-int4-Strix-native
+venv="${VENV:?Set VENV to the built runtime virtual environment}"
+model_path="${MODEL_PATH:?Set MODEL_PATH to the downloaded checkpoint directory}"
+host="${HOST:-127.0.0.1}"
+port="${PORT:-18081}"
+cache_root="${CACHE_ROOT:-$HOME/.cache/ciru-ling3-vllm}"
+
+test -x "$venv/bin/vllm"
+test -d "$model_path"
+site="$($venv/bin/python - <<'PY'
+import site
+print(site.getsitepackages()[0])
+PY
+)"
+sdk="$site/_rocm_sdk_devel"
+core="$site/_rocm_sdk_core"
+test -x "$sdk/bin/hipcc"
+mkdir -p "$cache_root/vllm" "$cache_root/triton" "$cache_root/huggingface-modules"
+
+export PATH="$sdk/bin:$sdk/llvm/bin:$venv/bin:$PATH"
+export LD_LIBRARY_PATH="$sdk/lib:$core/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
+export PYTHONPATH="$core/share/amd_smi${PYTHONPATH:+:$PYTHONPATH}"
+export HIP_DEVICE_LIB_PATH="$core/lib/llvm/amdgcn/bitcode"
+export ROCM_PATH="$sdk"
+export ROCM_HOME="$sdk"
+export ROCM_SOURCE_DIR="$sdk"
+export HIP_PATH="$sdk"
+export HIP_FORCE_DEV_KERNARG=1
+export HSA_NO_SCRATCH_RECLAIM=1
+export VLLM_ROCM_USE_AITER=0
+export VLLM_ROCM_USE_SKINNY_GEMM=0
+export VLLM_ROCM_SAFE_MERGE_ATTN_STATES=1
+export VLLM_CACHE_ROOT="$cache_root/vllm"
+export TRITON_CACHE_DIR="$cache_root/triton"
+export HF_MODULES_CACHE="$cache_root/huggingface-modules"
+export FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE
+export FLASH_ATTENTION_TRITON_AMD_AUTOTUNE=FALSE
+unset TORCH_BLAS_PREFER_HIPBLASLT
+unset PYTORCH_TUNABLEOP_ENABLED
+unset PYTORCH_TUNABLEOP_TUNING
+unset VLLM_TUNED_CONFIG_FOLDER
+
+args=(
+  serve "$model_path"
+  --host "$host"
+  --port "$port"
+  --served-model-name "$release_name"
+  --trust-remote-code
+  --attention-backend TRITON_MLA
+  --moe-backend triton
+  --enable-chunked-prefill
+  --enable-prefix-caching
+  --mamba-cache-mode align
+  --enable-auto-tool-choice
+  --tool-call-parser ling3
+  --reasoning-parser ling3
+  --max-model-len "$CONTEXT_LENGTH"
+  --max-num-seqs "$MAX_NUM_SEQS"
+  --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION"
+  --compilation-config '{"mode":3,"cudagraph_mode":"NONE","compile_sizes":[1,2]}'
+  --speculative-config "{\"method\":\"mtp\",\"num_speculative_tokens\":$SPECULATIVE_TOKENS}"
+)
+
+if [ -n "${HF_OVERRIDES:-}" ]; then
+  args+=(--hf-overrides "$HF_OVERRIDES")
+fi
+
+if [ "${DRY_RUN:-0}" = 1 ]; then
+  printf '%q ' "$venv/bin/vllm" "${args[@]}"
+  printf '\n'
+  exit 0
+fi
+
+exec "$venv/bin/vllm" "${args[@]}"