jcbtc commited on
Commit
9d29a15
·
verified ·
1 Parent(s): f504162

Release validated Ling runtime and installer fixes

Browse files
.gitattributes CHANGED
@@ -1,4 +1,4 @@
1
- *.bundle binary
2
  *.patch text eol=lf
3
  *.sh text eol=lf
4
  *.md text eol=lf
 
1
+ *.bundle filter=lfs diff=lfs merge=lfs -text
2
  *.patch text eol=lf
3
  *.sh text eol=lf
4
  *.md text eol=lf
CHANGELOG.md CHANGED
@@ -6,13 +6,40 @@
6
  - Model revision: `ca3ea63b0255d212c4fe6020db9e0a51ce136006`
7
  - Model format: 24 safetensors shards, 77,012,299,464 bytes total
8
  - vLLM upstream base: `d35eb6c44071ea806018841c490f0d2f3219c485`
9
- - CIRU runtime head: `91fab4e6dcaa683911655630dae8a5b140a7cb91`
10
- - Net vLLM delta: eight files, 459 insertions, 28 deletions, including regression tests
11
 
12
  ## Weight changes
13
 
14
  None. No tensor, quantization scale, tokenizer, template, configuration, or MTP weight was modified. The official packed-INT4 checkpoint is consumed directly.
15
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
16
  ## Installer release-gate hardening — August 10, 2026
17
 
18
  - Reconstruct the pinned vLLM commit through a dedicated remote ref and a
 
6
  - Model revision: `ca3ea63b0255d212c4fe6020db9e0a51ce136006`
7
  - Model format: 24 safetensors shards, 77,012,299,464 bytes total
8
  - vLLM upstream base: `d35eb6c44071ea806018841c490f0d2f3219c485`
9
+ - CIRU runtime head: `838616875c5dd4913d9f753221bf05f64cb7a7ed`
10
+ - Net vLLM delta: 10 files, 477 insertions, 28 deletions, including regression tests
11
 
12
  ## Weight changes
13
 
14
  None. No tensor, quantization scale, tokenizer, template, configuration, or MTP weight was modified. The official packed-INT4 checkpoint is consumed directly.
15
 
16
+ ## Cold-start diagnostics and environment registration — August 10, 2026
17
+
18
+ - Register the required `VLLM_ROCM_SAFE_MERGE_ATTN_STATES` switch in vLLM's
19
+ environment registry and read it through that registry. The fork no longer
20
+ warns that its own packaged setting is unknown.
21
+ - Make the Inductor max-autotune defaults explicit in the launcher and reject
22
+ malformed values before model loading.
23
+ - Explain before launch that the pinned Triton `make_block_ptr`, PyTorch
24
+ `TypedStorage`, and Inductor `AUTOTUNE` output is expected during a cold
25
+ compile. Readiness remains `Application startup complete` plus `/health`.
26
+ - Print the package's expected fork head beside the installed editable-source
27
+ head. A mismatch now produces the exact model-preserving CIRU rebuild command;
28
+ the ROCr-only `install.sh --upgrade` path is no longer easy to mistake for a
29
+ vLLM source update.
30
+
31
+ ## Clean NixOS installer correction — August 10, 2026
32
+
33
+ - Restore the host-toolchain bridge used by the validated NixOS build. The
34
+ installer now derives the active GCC, glibc, libstdc++, and NUMA paths instead
35
+ of embedding machine-specific Nix store hashes.
36
+ - Supply the pinned ROCr rebuild's `pkg-config`, `xxd`, DRM, ELF, NUMA, and
37
+ OpenSSL inputs through an ephemeral `nix-shell`, then translate its search
38
+ flags for CMake. No package is added to the user's Nix profile.
39
+ - Make the existing-install command update both the CIRU vLLM fork and the
40
+ corrected ROCr runtime while preserving the 77 GB checkpoint. Keep
41
+ `install.sh --upgrade` explicitly ROCr-only.
42
+
43
  ## Installer release-gate hardening — August 10, 2026
44
 
45
  - Reconstruct the pinned vLLM commit through a dedicated remote ref and a
Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz CHANGED
@@ -1,3 +1,3 @@
1
  version https://git-lfs.github.com/spec/v1
2
- oid sha256:ebda74c970ce2794942c259d52eea4666de000a54788575849f8a2dd2dd37d1f
3
- size 114586870
 
1
  version https://git-lfs.github.com/spec/v1
2
+ oid sha256:c4c88fb110c892da8069c59b370c2994f59ce968ce2baf8632b18b0aee9bdb2f
3
+ size 114330579
MANIFEST.json CHANGED
@@ -22,12 +22,12 @@
22
  "runtime": {
23
  "upstream_repo": "https://github.com/vllm-project/vllm.git",
24
  "upstream_base": "d35eb6c44071ea806018841c490f0d2f3219c485",
25
- "ciru_head": "91fab4e6dcaa683911655630dae8a5b140a7cb91",
26
- "commits_ahead": 16,
27
- "net_files_changed": 8,
28
- "net_insertions": 459,
29
  "net_deletions": 28,
30
- "engine_revision": "2026-08-10-installer-release-gate",
31
  "rocr_idle_fix": {
32
  "upstream_repo": "https://github.com/ROCm/rocm-systems.git",
33
  "upstream_commit": "44be71b52284948e58c93f65f46910399773fdcd",
 
22
  "runtime": {
23
  "upstream_repo": "https://github.com/vllm-project/vllm.git",
24
  "upstream_base": "d35eb6c44071ea806018841c490f0d2f3219c485",
25
+ "ciru_head": "838616875c5dd4913d9f753221bf05f64cb7a7ed",
26
+ "commits_ahead": 17,
27
+ "net_files_changed": 10,
28
+ "net_insertions": 477,
29
  "net_deletions": 28,
30
+ "engine_revision": "2026-08-10-startup-and-installer-hardening",
31
  "rocr_idle_fix": {
32
  "upstream_repo": "https://github.com/ROCm/rocm-systems.git",
33
  "upstream_commit": "44be71b52284948e58c93f65f46910399773fdcd",
README.md CHANGED
@@ -75,11 +75,21 @@ quoted. Because Linux `LD_PRELOAD` uses whitespace and colons as separators,
75
  the runtime install root must not contain either character. The installer now
76
  rejects such a root immediately instead of producing a broken launch later.
77
 
 
 
 
 
 
 
 
 
 
 
78
  **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.
79
 
80
  ## What this release changes
81
 
82
- The runtime is based on vLLM commit `d35eb6c44071ea806018841c490f0d2f3219c485` plus a 16-commit CIRU branch ending at `91fab4e6dcaa683911655630dae8a5b140a7cb91`. Its net source delta is eight files, 459 insertions, and 28 deletions, including regression tests.
83
 
84
  - Includes [`liminfei-amd`'s Wave32 LDS fix](https://github.com/vllm-project/vllm/pull/46012) for the ROCm top-k merge launch, preserving the AMD author's commit credit.
85
  - Adds an opt-in safe PyTorch attention-state merge for a `gfx1151` Triton/HSA fault, including both supported MLA LSE layouts.
@@ -167,7 +177,7 @@ The named release archive is `Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz`. It
167
 
168
  ## Quick start
169
 
170
- 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 `91fab4e6d`.
171
 
172
  For a clean runtime installation, download the 115 MB release archive instead
173
  of cloning this model repository. A normal Git LFS clone can also download the
@@ -217,24 +227,27 @@ installed. Add `--download-model` only when you want this installer to fetch
217
  the checkpoint; omit it for a runtime-only install. Use
218
  `bash install.sh --help` for the complete agent interface.
219
 
220
- ### Upgrade an existing CIRU engine
221
 
222
- Users who already installed this release do not need to download the model or
223
- rebuild vLLM. From the updated package directory, run:
224
 
225
  ```bash
226
- bash install.sh --upgrade \
227
  --install-root "$HOME/ciru-ling-runtime" \
 
228
  --install-host-deps
229
  ```
230
 
231
  Replace the install root with the directory that contains the existing
232
- `.venv`. The upgrade compiles the pinned ROCr userspace library with the host
233
- GCC, records and verifies its SHA-256 checksum, and makes the launchers preload
234
- it automatically. It does not modify model weights, the model directory, or
235
- `ling3-runtime.env`. Restart the server after the command completes. If an
236
- emergency rollback is needed, launch once with
237
- `CIRU_DISABLE_ROCR_IDLE_FIX=1`.
 
 
238
 
239
  The API advertises model ID `Ling-3.0-Flash-CIRU-int4-Strix-native` on `http://127.0.0.1:18081/v1`.
240
 
@@ -251,6 +264,35 @@ bash scripts/run-256k.sh
251
 
252
  Lower values reduce KV-cache capacity. Even at the packaged 0.72 setting, six requests cannot all occupy the full 256K limit simultaneously; verify the reported KV-cache capacity at startup for your workload and host.
253
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
254
  For the experimental 1M profile instead:
255
 
256
  ```bash
 
75
  the runtime install root must not contain either character. The installer now
76
  rejects such a root immediately instead of producing a broken launch later.
77
 
78
+ The current launcher also identifies cold-cache compilation before vLLM starts.
79
+ Messages containing `tl.make_block_ptr is deprecated`, `TypedStorage is
80
+ deprecated`, `Autotune Choices Stats`, or `AUTOTUNE mm(...)` are emitted by the
81
+ pinned Triton/PyTorch compiler path and are not startup failures. Wait for
82
+ `Application startup complete`, then verify `/health`; a traceback,
83
+ `EngineCore failed`, a nonzero process exit, or a failed health check is an
84
+ actual failure. The required `VLLM_ROCM_SAFE_MERGE_ATTN_STATES` setting is now
85
+ registered by the CIRU fork, so vLLM no longer mislabels it as an unknown
86
+ environment variable.
87
+
88
  **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.
89
 
90
  ## What this release changes
91
 
92
+ The runtime is based on vLLM commit `d35eb6c44071ea806018841c490f0d2f3219c485` plus a 17-commit CIRU branch ending at `838616875c5dd4913d9f753221bf05f64cb7a7ed`. Its net source delta is 10 files, 477 insertions, and 28 deletions, including regression tests.
93
 
94
  - Includes [`liminfei-amd`'s Wave32 LDS fix](https://github.com/vllm-project/vllm/pull/46012) for the ROCm top-k merge launch, preserving the AMD author's commit credit.
95
  - Adds an opt-in safe PyTorch attention-state merge for a `gfx1151` Triton/HSA fault, including both supported MLA LSE layouts.
 
177
 
178
  ## Quick start
179
 
180
+ 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 `838616875`.
181
 
182
  For a clean runtime installation, download the 115 MB release archive instead
183
  of cloning this model repository. A normal Git LFS clone can also download the
 
227
  the checkpoint; omit it for a runtime-only install. Use
228
  `bash install.sh --help` for the complete agent interface.
229
 
230
+ ### Update an existing CIRU installation
231
 
232
+ From the updated package directory, rebuild the packaged fork and the corrected
233
+ ROCr runtime while preserving the existing model:
234
 
235
  ```bash
236
+ bash install.sh \
237
  --install-root "$HOME/ciru-ling-runtime" \
238
+ --model-path "/path/to/Ling-checkpoint" \
239
  --install-host-deps
240
  ```
241
 
242
  Replace the install root with the directory that contains the existing
243
+ `.venv` and use the existing checkpoint directory. This does not download,
244
+ move, or modify the weights. It installs runtime head `838616875`, builds the
245
+ pinned ROCr userspace override that fixes the idle 100% CPU spin, and rewrites
246
+ `ling3-runtime.env` with those same paths. Restart the server afterward.
247
+
248
+ `install.sh --upgrade` is intentionally ROCr-only. Use it only when the vLLM
249
+ source already matches the package and just the idle-CPU override needs repair.
250
+ For emergency ROCr rollback, launch once with `CIRU_DISABLE_ROCR_IDLE_FIX=1`.
251
 
252
  The API advertises model ID `Ling-3.0-Flash-CIRU-int4-Strix-native` on `http://127.0.0.1:18081/v1`.
253
 
 
264
 
265
  Lower values reduce KV-cache capacity. Even at the packaged 0.72 setting, six requests cannot all occupy the full 256K limit simultaneously; verify the reported KV-cache capacity at startup for your workload and host.
266
 
267
+ ### Cold-start compiler output
268
+
269
+ The optimized launcher deliberately uses vLLM compilation mode 3 with static
270
+ decode sizes `[1,2]` and Inductor max autotuning. On an empty cache, model
271
+ loading plus compilation can take several minutes and prints kernel comparisons
272
+ such as `AUTOTUNE mm(...)`. The two deprecation warnings named above are also
273
+ expected with the pinned dependency versions. They occurred in the validated
274
+ run before the API became healthy and returned HTTP 200.
275
+
276
+ Do not remove launcher flags merely to make these messages disappear. First
277
+ run the packaged command unchanged and wait for `Application startup complete`:
278
+
279
+ ```bash
280
+ source ling3-runtime.env
281
+ bash scripts/run-256k.sh
282
+ ```
283
+
284
+ The launcher prints both the runtime head expected by the package and the head
285
+ of the installed editable vLLM source. If they differ, update the CIRU fork by
286
+ running the exact rebuild command printed there. That command reuses the model
287
+ path and does not download weights; `install.sh --upgrade` is intentionally
288
+ ROCr-only and is not a vLLM source update.
289
+
290
+ If a modified command is required on another distribution, retain the complete
291
+ command, the first traceback or `EngineCore failed` block, and the final 100 log
292
+ lines. Performance from a command with MTP, compilation, max autotuning, prefix
293
+ caching, or the 8,192-token scheduler budget disabled is not comparable to the
294
+ published benchmark profile.
295
+
296
  For the experimental 1M profile instead:
297
 
298
  ```bash
SHA256SUMS CHANGED
@@ -1,23 +1,25 @@
1
- 7b77dda00cc6bd6d0ce15f7c055f9e997f56e1cf1709c61cca5a57283ce3f008 ./.gitattributes
2
- b3a12fee0f8e5187dcb389bbb8fc944b84e132b395a43edd8feee1c9e725a057 ./CHANGELOG.md
3
- 95d15a4b46e2c38c1c2af5e098166d43bdd192653a12cc53ad7159420f81e711 ./MANIFEST.json
4
- 9151f96f0a7322d759e1c2f9aa95530fa621f4a78a10f8e85e2fa04bf59e61bb ./README.md
5
  411ac21b4849cb3bf65f84a55263550dcd658e0623c2b117233ba1fced1c17ff ./assets/ling30int4.png
6
  dbf3f912d175e69d30b9d352bb01446f231d72112181be6ba3d5b34be741a869 ./config/experimental-yarn-1m.env
7
  725bb450e9df2f76bc34604ec6c08c55aa7e19034f244e76a6cc57d0b9ebcbaa ./config/native-256k.env
8
- c9fb6ed33fa8f98e1c4969f5e06118836da733f719b24f7a95a5a48bcb85b3f6 ./docs/BENCHMARKS.md
9
- 801506edabd1e1f3e15c41a3890084d8e15946123652d8677379699009866ace ./docs/UBUNTU_BUILD.md
10
- cb071718556c99c46e7168e02a6ed645daaad0d078f96d9d948217e261a30e4d ./docs/VALIDATED_ENVIRONMENT.md
11
  eb927bfb99e64911dbc0f6bc0b41f692a07f8d9aed4b551de2de3b1a1acbfd64 ./docs/WINDOWS_WSL2.md
12
  06c78a1ef6f5ee0c3571bd5d35a0ea0f32e2620d999d02f2c9e9f175141851e1 ./install.sh
13
- a14fcd3747829be98a73a453057bd451a492d87a236c2affcf57db1cbbf18386 ./runtime/COMMITS.txt
14
- 9cbcbbef5e113e22b753801ee141bd877b20c84f25f15cc159d0ae508aaf4324 ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
15
- c0c1446ea2c7334ef8b888f3104ec660d2ab41d2cb08bb9d90a3f589dba040e8 ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
16
- ecab16a7d189d48ff85942ebefd49f2c206445a4e4a85de440ebde19b6555551 ./runtime/SHA256SUMS.runtime
17
- 0831d1f5c296525333323250227b20be8cee450c33a398d5f2629893b634e1ce ./scripts/build-vllm-gfx1151.sh
18
  e41ff92b71b146a745161f6b6b09f1c85d4b828e68aa1e2a151311656c19a7d7 ./scripts/download-model.sh
19
- 566b45b97266bdf215c2da2765b52465c515808223098819ac4086335e6759e4 ./scripts/install-host-deps.sh
20
- 2ace002dc64d1ccd20df9ad02b8d3d0212c05c50ced8f26be9c24cf889c0f925 ./scripts/install-rocr-idle-fix.sh
 
 
21
  72f4024af2ad19ee58af201d4fc29f84481b521a80e36c88ae1ab0b6edf7e0e1 ./scripts/run-1m-yarn-experimental.sh
22
  f3fa5364ce194b6676bbe3566604f73d2df9ac381fe1937bf7abc465dde8fd2c ./scripts/run-256k.sh
23
- 581792ad19506e9db61fddcf1c077c8b81dc72e8e5a5ac72eba722597052dfb6 ./scripts/run-common.sh
 
1
+ 46cf1b9e12bdc3a2efbcd762a98b55b2a2b0e30ade3eb03134e0b8c45e84049a ./.gitattributes
2
+ 90b6bea03f33e6c0a9830ea4e43eb0c1ab9477241af0c42e2a7eb7da689e4a38 ./CHANGELOG.md
3
+ c641ecff9c6ef45d00e2bda4d7995d4e25522b213fc838100babd8fac86121f0 ./MANIFEST.json
4
+ 49062288d0b8895baca666fc5d76559ae5ae1bedad7114e62000dc5b1538cbcb ./README.md
5
  411ac21b4849cb3bf65f84a55263550dcd658e0623c2b117233ba1fced1c17ff ./assets/ling30int4.png
6
  dbf3f912d175e69d30b9d352bb01446f231d72112181be6ba3d5b34be741a869 ./config/experimental-yarn-1m.env
7
  725bb450e9df2f76bc34604ec6c08c55aa7e19034f244e76a6cc57d0b9ebcbaa ./config/native-256k.env
8
+ 3f67af3f3acaf08815ece0d373f2a9edf7abfc4e3910bb7061e7b8f01070152c ./docs/BENCHMARKS.md
9
+ 28fb88619cb4b51790dd267506410bf81530d9ad17c1cee290866d3c23c4c708 ./docs/UBUNTU_BUILD.md
10
+ 86501c16acd668d0e25613c60c61a395c31731836c09c37fafe11490d0eb2370 ./docs/VALIDATED_ENVIRONMENT.md
11
  eb927bfb99e64911dbc0f6bc0b41f692a07f8d9aed4b551de2de3b1a1acbfd64 ./docs/WINDOWS_WSL2.md
12
  06c78a1ef6f5ee0c3571bd5d35a0ea0f32e2620d999d02f2c9e9f175141851e1 ./install.sh
13
+ 1bec28ea9c6bfa28aea0e0c7adf3d46b56eb399156680278fba0e5b5e3cc87b7 ./runtime/COMMITS.txt
14
+ 58dde46a772bedfaa7a2bd806a774db15352d5b53ad7c5c681b6857c273561b5 ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
15
+ c6839a15d8250d7e9541f6a12b8cef0f9ccfbfdd825f095ee313048c56bdb0f4 ./runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
16
+ fc47eb458b4d8e4db52b5cb8d9dbd912ee16dbd63ade61861d196e300374c4ca ./runtime/SHA256SUMS.runtime
17
+ 8541e877203c3954bf0df44b2a4027188ef694a0ee644e29ddcdfa0d4d2c21c0 ./scripts/build-vllm-gfx1151.sh
18
  e41ff92b71b146a745161f6b6b09f1c85d4b828e68aa1e2a151311656c19a7d7 ./scripts/download-model.sh
19
+ 37e5e506ce70b0388926456c59ed14b3565a6067ecdb19242e73252fb91b0792 ./scripts/hipcc-nixos-linker-helper.sh
20
+ a16abc24cfcb590b191bf122aaf3df01549ea3baf6b5eaa8d089da064b15f043 ./scripts/install-host-deps.sh
21
+ d2793fc18ee5438efa55b77ff0f4cfe090c52ea1541b82fa94cd8da5d128a7d6 ./scripts/install-rocr-idle-fix.sh
22
+ 881b565217b0881b78e0e9bec074ffd27ba013ec21f63d595de76b0693289067 ./scripts/rocm715-nixos-numa-target.cmake
23
  72f4024af2ad19ee58af201d4fc29f84481b521a80e36c88ae1ab0b6edf7e0e1 ./scripts/run-1m-yarn-experimental.sh
24
  f3fa5364ce194b6676bbe3566604f73d2df9ac381fe1937bf7abc465dde8fd2c ./scripts/run-256k.sh
25
+ e705d68ea20fc4dd07ead15111ace08f29b6ea7c60f3613690a111bdfd31b437 ./scripts/run-common.sh
SHA256SUMS.release CHANGED
@@ -1,25 +1,27 @@
1
- ebda74c970ce2794942c259d52eea4666de000a54788575849f8a2dd2dd37d1f ./Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz
2
- 7b77dda00cc6bd6d0ce15f7c055f9e997f56e1cf1709c61cca5a57283ce3f008 ./Ling-3.0-Flash-CIRU-int4-Strix-native/.gitattributes
3
- b3a12fee0f8e5187dcb389bbb8fc944b84e132b395a43edd8feee1c9e725a057 ./Ling-3.0-Flash-CIRU-int4-Strix-native/CHANGELOG.md
4
- 95d15a4b46e2c38c1c2af5e098166d43bdd192653a12cc53ad7159420f81e711 ./Ling-3.0-Flash-CIRU-int4-Strix-native/MANIFEST.json
5
- 9151f96f0a7322d759e1c2f9aa95530fa621f4a78a10f8e85e2fa04bf59e61bb ./Ling-3.0-Flash-CIRU-int4-Strix-native/README.md
6
- 5a7af736e6c1e1d51717f32cdb9b0331261f3b5d0c2491c474905ee8b1c6639c ./Ling-3.0-Flash-CIRU-int4-Strix-native/SHA256SUMS
7
- 411ac21b4849cb3bf65f84a55263550dcd658e0623c2b117233ba1fced1c17ff ./Ling-3.0-Flash-CIRU-int4-Strix-native/assets/ling30int4.png
8
- dbf3f912d175e69d30b9d352bb01446f231d72112181be6ba3d5b34be741a869 ./Ling-3.0-Flash-CIRU-int4-Strix-native/config/experimental-yarn-1m.env
9
- 725bb450e9df2f76bc34604ec6c08c55aa7e19034f244e76a6cc57d0b9ebcbaa ./Ling-3.0-Flash-CIRU-int4-Strix-native/config/native-256k.env
10
- c9fb6ed33fa8f98e1c4969f5e06118836da733f719b24f7a95a5a48bcb85b3f6 ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/BENCHMARKS.md
11
- 801506edabd1e1f3e15c41a3890084d8e15946123652d8677379699009866ace ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/UBUNTU_BUILD.md
12
- cb071718556c99c46e7168e02a6ed645daaad0d078f96d9d948217e261a30e4d ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/VALIDATED_ENVIRONMENT.md
13
- eb927bfb99e64911dbc0f6bc0b41f692a07f8d9aed4b551de2de3b1a1acbfd64 ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/WINDOWS_WSL2.md
14
- 06c78a1ef6f5ee0c3571bd5d35a0ea0f32e2620d999d02f2c9e9f175141851e1 ./Ling-3.0-Flash-CIRU-int4-Strix-native/install.sh
15
- a14fcd3747829be98a73a453057bd451a492d87a236c2affcf57db1cbbf18386 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/COMMITS.txt
16
- 9cbcbbef5e113e22b753801ee141bd877b20c84f25f15cc159d0ae508aaf4324 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
17
- c0c1446ea2c7334ef8b888f3104ec660d2ab41d2cb08bb9d90a3f589dba040e8 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
18
- ecab16a7d189d48ff85942ebefd49f2c206445a4e4a85de440ebde19b6555551 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/SHA256SUMS.runtime
19
- 0831d1f5c296525333323250227b20be8cee450c33a398d5f2629893b634e1ce ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/build-vllm-gfx1151.sh
20
- e41ff92b71b146a745161f6b6b09f1c85d4b828e68aa1e2a151311656c19a7d7 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/download-model.sh
21
- 566b45b97266bdf215c2da2765b52465c515808223098819ac4086335e6759e4 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/install-host-deps.sh
22
- 2ace002dc64d1ccd20df9ad02b8d3d0212c05c50ced8f26be9c24cf889c0f925 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/install-rocr-idle-fix.sh
23
- 72f4024af2ad19ee58af201d4fc29f84481b521a80e36c88ae1ab0b6edf7e0e1 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-1m-yarn-experimental.sh
24
- f3fa5364ce194b6676bbe3566604f73d2df9ac381fe1937bf7abc465dde8fd2c ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-256k.sh
25
- 581792ad19506e9db61fddcf1c077c8b81dc72e8e5a5ac72eba722597052dfb6 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-common.sh
 
 
 
1
+ c4c88fb110c892da8069c59b370c2994f59ce968ce2baf8632b18b0aee9bdb2f ./Ling-3.0-Flash-CIRU-int4-Strix-native.tar.gz
2
+ 46cf1b9e12bdc3a2efbcd762a98b55b2a2b0e30ade3eb03134e0b8c45e84049a ./Ling-3.0-Flash-CIRU-int4-Strix-native/.gitattributes
3
+ 411ac21b4849cb3bf65f84a55263550dcd658e0623c2b117233ba1fced1c17ff ./Ling-3.0-Flash-CIRU-int4-Strix-native/assets/ling30int4.png
4
+ 90b6bea03f33e6c0a9830ea4e43eb0c1ab9477241af0c42e2a7eb7da689e4a38 ./Ling-3.0-Flash-CIRU-int4-Strix-native/CHANGELOG.md
5
+ dbf3f912d175e69d30b9d352bb01446f231d72112181be6ba3d5b34be741a869 ./Ling-3.0-Flash-CIRU-int4-Strix-native/config/experimental-yarn-1m.env
6
+ 725bb450e9df2f76bc34604ec6c08c55aa7e19034f244e76a6cc57d0b9ebcbaa ./Ling-3.0-Flash-CIRU-int4-Strix-native/config/native-256k.env
7
+ 3f67af3f3acaf08815ece0d373f2a9edf7abfc4e3910bb7061e7b8f01070152c ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/BENCHMARKS.md
8
+ 28fb88619cb4b51790dd267506410bf81530d9ad17c1cee290866d3c23c4c708 ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/UBUNTU_BUILD.md
9
+ 86501c16acd668d0e25613c60c61a395c31731836c09c37fafe11490d0eb2370 ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/VALIDATED_ENVIRONMENT.md
10
+ eb927bfb99e64911dbc0f6bc0b41f692a07f8d9aed4b551de2de3b1a1acbfd64 ./Ling-3.0-Flash-CIRU-int4-Strix-native/docs/WINDOWS_WSL2.md
11
+ 06c78a1ef6f5ee0c3571bd5d35a0ea0f32e2620d999d02f2c9e9f175141851e1 ./Ling-3.0-Flash-CIRU-int4-Strix-native/install.sh
12
+ c641ecff9c6ef45d00e2bda4d7995d4e25522b213fc838100babd8fac86121f0 ./Ling-3.0-Flash-CIRU-int4-Strix-native/MANIFEST.json
13
+ 49062288d0b8895baca666fc5d76559ae5ae1bedad7114e62000dc5b1538cbcb ./Ling-3.0-Flash-CIRU-int4-Strix-native/README.md
14
+ 1bec28ea9c6bfa28aea0e0c7adf3d46b56eb399156680278fba0e5b5e3cc87b7 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/COMMITS.txt
15
+ 58dde46a772bedfaa7a2bd806a774db15352d5b53ad7c5c681b6857c273561b5 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
16
+ c6839a15d8250d7e9541f6a12b8cef0f9ccfbfdd825f095ee313048c56bdb0f4 ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
17
+ fc47eb458b4d8e4db52b5cb8d9dbd912ee16dbd63ade61861d196e300374c4ca ./Ling-3.0-Flash-CIRU-int4-Strix-native/runtime/SHA256SUMS.runtime
18
+ 8541e877203c3954bf0df44b2a4027188ef694a0ee644e29ddcdfa0d4d2c21c0 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/build-vllm-gfx1151.sh
19
+ e41ff92b71b146a745161f6b6b09f1c85d4b828e68aa1e2a151311656c19a7d7 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/download-model.sh
20
+ 37e5e506ce70b0388926456c59ed14b3565a6067ecdb19242e73252fb91b0792 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/hipcc-nixos-linker-helper.sh
21
+ a16abc24cfcb590b191bf122aaf3df01549ea3baf6b5eaa8d089da064b15f043 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/install-host-deps.sh
22
+ d2793fc18ee5438efa55b77ff0f4cfe090c52ea1541b82fa94cd8da5d128a7d6 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/install-rocr-idle-fix.sh
23
+ 881b565217b0881b78e0e9bec074ffd27ba013ec21f63d595de76b0693289067 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/rocm715-nixos-numa-target.cmake
24
+ 72f4024af2ad19ee58af201d4fc29f84481b521a80e36c88ae1ab0b6edf7e0e1 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-1m-yarn-experimental.sh
25
+ f3fa5364ce194b6676bbe3566604f73d2df9ac381fe1937bf7abc465dde8fd2c ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-256k.sh
26
+ e705d68ea20fc4dd07ead15111ace08f29b6ea7c60f3613690a111bdfd31b437 ./Ling-3.0-Flash-CIRU-int4-Strix-native/scripts/run-common.sh
27
+ 9b462200154d5951ced0bf2ea2e11c7aeeb76e4ad54769cfa1ac6af81de8038a ./Ling-3.0-Flash-CIRU-int4-Strix-native/SHA256SUMS
docs/BENCHMARKS.md CHANGED
@@ -18,7 +18,7 @@ The native row used the then-current `gpu_memory_utilization=0.75`, 262,144-cont
18
 
19
  ### Update status
20
 
21
- The native 60K row was measured at CIRU runtime commit `fabbc7e928dfa64c29dcdf0f795b05246555f4ae`, before the final exact-shape SiLU kernel and ROCm MLA LSE-layout commits. The released commit `91fab4e6dcaa683911655630dae8a5b140a7cb91` 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.
22
 
23
  ## Short decode and concurrent serving
24
 
 
18
 
19
  ### Update status
20
 
21
+ The native 60K row was measured at CIRU runtime commit `fabbc7e928dfa64c29dcdf0f795b05246555f4ae`, before the final exact-shape SiLU kernel and ROCm MLA LSE-layout commits. The release branch now ends at `838616875c5dd4913d9f753221bf05f64cb7a7ed`; its last performance-affecting runtime commit is `91fab4e6dcaa683911655630dae8a5b140a7cb91`. That branch 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 release branch's maximum 60K speed. Replace the native row only after a matched rerun of all three stacks or a clearly labeled native-only refresh.
22
 
23
  ## Short decode and concurrent serving
24
 
docs/UBUNTU_BUILD.md CHANGED
@@ -65,22 +65,29 @@ spaces.
65
  For non-interactive agents, see `bash install.sh --help`. The numbered steps
66
  below remain available for manual diagnosis and custom automation.
67
 
68
- ### Existing installations: engine-only upgrade
69
 
70
- After updating this package, upgrade an existing runtime in place with:
 
71
 
72
  ```bash
73
- bash install.sh --upgrade \
74
  --install-root "$HOME/ciru-ling-runtime" \
 
75
  --install-host-deps
76
  ```
77
 
78
- The install root is the parent of the existing `.venv`. This operation does
79
- not download, move, or alter the model, and it leaves `ling3-runtime.env`
80
- unchanged. It builds only the pinned ROCr userspace library, verifies the
81
- result, and enables it in the packaged launchers. Restart the running service
82
- afterward. `CIRU_DISABLE_ROCR_IDLE_FIX=1` temporarily disables the preload for
83
- rollback or diagnosis.
 
 
 
 
 
84
 
85
  ## 2. Configure unified-memory access
86
 
@@ -154,6 +161,15 @@ MODEL_PATH="$HOME/models/Ling-3.0-Flash-CIRU-int4-Strix-native" \
154
  bash scripts/run-256k.sh
155
  ```
156
 
 
 
 
 
 
 
 
 
 
157
  In another shell:
158
 
159
  ```bash
 
65
  For non-interactive agents, see `bash install.sh --help`. The numbered steps
66
  below remain available for manual diagnosis and custom automation.
67
 
68
+ ### Existing installations: full runtime update, model preserved
69
 
70
+ After updating this package, rebuild the packaged fork and corrected ROCr
71
+ runtime in place with:
72
 
73
  ```bash
74
+ bash install.sh \
75
  --install-root "$HOME/ciru-ling-runtime" \
76
+ --model-path "/path/to/Ling-checkpoint" \
77
  --install-host-deps
78
  ```
79
 
80
+ The install root is the parent of the existing `.venv`; the model path points
81
+ to the existing checkpoint. This operation does not download, move, or alter
82
+ the model. It updates both the CIRU vLLM fork and the pinned ROCr userspace
83
+ override that fixes the idle 100% CPU spin. Restart the service afterward.
84
+ `install.sh --upgrade` remains a ROCr-only repair path for an installation whose
85
+ vLLM source already matches the package. `CIRU_DISABLE_ROCR_IDLE_FIX=1`
86
+ temporarily disables the preload for rollback or diagnosis.
87
+
88
+ On NixOS, the build derives the active GCC/glibc store paths instead of using
89
+ host-specific hashes. The ROCr step enters an ephemeral `nix-shell` for its
90
+ host build dependencies; it does not add packages to the user's profile.
91
 
92
  ## 2. Configure unified-memory access
93
 
 
161
  bash scripts/run-256k.sh
162
  ```
163
 
164
+ On a cold compiler cache, keep the launch terminal open for several minutes.
165
+ The pinned stack can print `tl.make_block_ptr is deprecated`, `TypedStorage is
166
+ deprecated`, `Autotune Choices Stats`, and `AUTOTUNE mm(...)` while selecting
167
+ kernels. Those lines are nonfatal; readiness is the later `Application startup
168
+ complete` line and a successful `/health` response. A traceback, `EngineCore
169
+ failed`, process exit, or failed health check is an actual failure. Test the
170
+ packaged launcher unchanged before removing optimization flags, because doing
171
+ so changes both behavior and the performance profile.
172
+
173
  In another shell:
174
 
175
  ```bash
docs/VALIDATED_ENVIRONMENT.md CHANGED
@@ -19,7 +19,7 @@
19
  - torchaudio `2.11.0.2+rocm7.15.0a20260728`
20
  - FlashAttention `2.8.3`, Dao-AILab commit `0e60e39473e8df549a20fb5353760f7a65b30e2d`, Python-only AMD Triton path
21
  - vLLM upstream base `d35eb6c44071ea806018841c490f0d2f3219c485`
22
- - CIRU runtime head `91fab4e6dcaa683911655630dae8a5b140a7cb91`
23
 
24
  ## Build facts
25
 
@@ -27,6 +27,8 @@
27
  - Exact AMD multi-architecture wheel transaction installed in one resolver operation.
28
  - `rocm-sdk init` and all 27 SDK tests passed.
29
  - vLLM compiled for `gfx1151` only with 16 jobs.
 
 
30
  - Clean native build completed in 92 seconds on the validation host.
31
  - Imports passed for `vllm._C`, `vllm._C_stable_libtorch`, `vllm._moe_C_stable_libtorch`, `vllm._rocm_C`, and `vllm.cumem_allocator`.
32
 
 
19
  - torchaudio `2.11.0.2+rocm7.15.0a20260728`
20
  - FlashAttention `2.8.3`, Dao-AILab commit `0e60e39473e8df549a20fb5353760f7a65b30e2d`, Python-only AMD Triton path
21
  - vLLM upstream base `d35eb6c44071ea806018841c490f0d2f3219c485`
22
+ - CIRU runtime head `838616875c5dd4913d9f753221bf05f64cb7a7ed`
23
 
24
  ## Build facts
25
 
 
27
  - Exact AMD multi-architecture wheel transaction installed in one resolver operation.
28
  - `rocm-sdk init` and all 27 SDK tests passed.
29
  - vLLM compiled for `gfx1151` only with 16 jobs.
30
+ - The clean package derived the NixOS GCC/glibc bridge without fixed store
31
+ hashes and supplied ROCr host dependencies through an ephemeral `nix-shell`.
32
  - Clean native build completed in 92 seconds on the validation host.
33
  - Imports passed for `vllm._C`, `vllm._C_stable_libtorch`, `vllm._moe_C_stable_libtorch`, `vllm._rocm_C`, and `vllm.cumem_allocator`.
34
 
runtime/COMMITS.txt CHANGED
@@ -14,3 +14,4 @@ abfff126e0698d98f1872f87f501d36a19382240 Optimize single-token MoE reduction
14
  fabbc7e928dfa64c29dcdf0f795b05246555f4ae Extend MoE support fast path to K1 verification
15
  388d82de2abe418e73380ba250270d3abc62ac12 ROCm: optimize exact small WNA16 SiLU activation
16
  91fab4e6dcaa683911655630dae8a5b140a7cb91 [Bugfix][ROCm][MLA] Normalize upstream FlashAttention LSE layout
 
 
14
  fabbc7e928dfa64c29dcdf0f795b05246555f4ae Extend MoE support fast path to K1 verification
15
  388d82de2abe418e73380ba250270d3abc62ac12 ROCm: optimize exact small WNA16 SiLU activation
16
  91fab4e6dcaa683911655630dae8a5b140a7cb91 [Bugfix][ROCm][MLA] Normalize upstream FlashAttention LSE layout
17
+ 838616875c5dd4913d9f753221bf05f64cb7a7ed ROCm: register safe attention merge environment flag
runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle CHANGED
@@ -1,3 +1,3 @@
1
  version https://git-lfs.github.com/spec/v1
2
- oid sha256:9cbcbbef5e113e22b753801ee141bd877b20c84f25f15cc159d0ae508aaf4324
3
- size 117890929
 
1
  version https://git-lfs.github.com/spec/v1
2
+ oid sha256:58dde46a772bedfaa7a2bd806a774db15352d5b53ad7c5c681b6857c273561b5
3
+ size 117557885
runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch CHANGED
@@ -15,6 +15,26 @@ index 519e21328..847d8284a 100644
15
  vllm::topKPerRowDecode<kNumThreadsPerBlockMerge, true, false, true>
16
  <<<numRows, kNumThreadsPerBlockMerge, topK * sizeof(int32_t), stream>>>(
17
  outLogitsAux.const_data_ptr<float>(), seqLens.const_data_ptr<int>(),
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
18
  diff --git a/tests/v1/attention/test_mla_prefill_quant_output.py b/tests/v1/attention/test_mla_prefill_quant_output.py
19
  index d7659485a..93df32d38 100644
20
  --- a/tests/v1/attention/test_mla_prefill_quant_output.py
@@ -136,6 +156,31 @@ index d7659485a..93df32d38 100644
136
  @pytest.mark.parametrize(
137
  ("version", "is_vllm_fa", "dc_major", "quant_key", "expected"),
138
  [
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
139
  diff --git a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py b/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
140
  index 9d93dfe54..dd96f6da4 100644
141
  --- a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
@@ -575,17 +620,14 @@ index 71a38e95d..641f68b16 100644
575
 
576
  # Run MQA — always pass layer scales. When KV cache is
577
  diff --git a/vllm/v1/attention/ops/merge_attn_states.py b/vllm/v1/attention/ops/merge_attn_states.py
578
- index 20c7503e9..cc860e1ab 100644
579
  --- a/vllm/v1/attention/ops/merge_attn_states.py
580
  +++ b/vllm/v1/attention/ops/merge_attn_states.py
581
- @@ -1,11 +1,87 @@
582
- # SPDX-License-Identifier: Apache-2.0
583
- # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
584
 
585
- +import os
586
- +
587
  import torch
588
 
 
589
  from vllm.platforms import current_platform
590
 
591
 
@@ -666,13 +708,13 @@ index 20c7503e9..cc860e1ab 100644
666
  def merge_attn_states(
667
  output: torch.Tensor,
668
  prefix_output: torch.Tensor,
669
- @@ -46,6 +122,21 @@ def merge_attn_states(
670
  When provided, output must be FP8 dtype.
671
  """
672
 
673
  + if (
674
  + current_platform.is_rocm()
675
- + and os.getenv("VLLM_ROCM_SAFE_MERGE_ATTN_STATES", "0") == "1"
676
  + and output_scale is None
677
  + ):
678
  + return _merge_attn_states_torch(
 
15
  vllm::topKPerRowDecode<kNumThreadsPerBlockMerge, true, false, true>
16
  <<<numRows, kNumThreadsPerBlockMerge, topK * sizeof(int32_t), stream>>>(
17
  outLogitsAux.const_data_ptr<float>(), seqLens.const_data_ptr<int>(),
18
+ diff --git a/tests/config/test_config_generation.py b/tests/config/test_config_generation.py
19
+ index c7edf2b97..7c9b21c24 100644
20
+ --- a/tests/config/test_config_generation.py
21
+ +++ b/tests/config/test_config_generation.py
22
+ @@ -109,3 +109,15 @@ def test_unrecognized_env(monkeypatch):
23
+ fail_on_environ_validation=True,
24
+ )
25
+ engine_args.create_engine_config()
26
+ +
27
+ +
28
+ +def test_rocm_safe_merge_env_is_registered(monkeypatch):
29
+ + from vllm.envs import environment_variables
30
+ +
31
+ + name = "VLLM_ROCM_SAFE_MERGE_ATTN_STATES"
32
+ + assert name in environment_variables
33
+ +
34
+ + monkeypatch.setenv(name, "1")
35
+ + assert environment_variables[name]() is True
36
+ + monkeypatch.setenv(name, "0")
37
+ + assert environment_variables[name]() is False
38
  diff --git a/tests/v1/attention/test_mla_prefill_quant_output.py b/tests/v1/attention/test_mla_prefill_quant_output.py
39
  index d7659485a..93df32d38 100644
40
  --- a/tests/v1/attention/test_mla_prefill_quant_output.py
 
156
  @pytest.mark.parametrize(
157
  ("version", "is_vllm_fa", "dc_major", "quant_key", "expected"),
158
  [
159
+ diff --git a/vllm/envs.py b/vllm/envs.py
160
+ index 3f23fb615..3a15c5c93 100755
161
+ --- a/vllm/envs.py
162
+ +++ b/vllm/envs.py
163
+ @@ -143,6 +143,7 @@ if TYPE_CHECKING:
164
+ VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS: bool = False
165
+ VLLM_ROCM_USE_AITER_TRITON_GEMM: bool = True
166
+ VLLM_ROCM_USE_SKINNY_GEMM: bool = True
167
+ + VLLM_ROCM_SAFE_MERGE_ATTN_STATES: bool = False
168
+ VLLM_ROCM_FP8_PADDING: bool = True
169
+ VLLM_ROCM_MOE_PADDING: bool = True
170
+ VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT: bool = False
171
+ @@ -1297,6 +1298,12 @@ environment_variables: dict[str, Callable[[], Any]] = {
172
+ "VLLM_ROCM_USE_SKINNY_GEMM": lambda: (
173
+ os.getenv("VLLM_ROCM_USE_SKINNY_GEMM", "True").lower() in ("true", "1")
174
+ ),
175
+ + # Use the safe Torch attention-state merge on ROCm. This is an opt-in
176
+ + # fallback for platforms where the Triton merge kernel is unstable.
177
+ + "VLLM_ROCM_SAFE_MERGE_ATTN_STATES": lambda: (
178
+ + os.getenv("VLLM_ROCM_SAFE_MERGE_ATTN_STATES", "False").lower()
179
+ + in ("true", "1")
180
+ + ),
181
+ # Pad the fp8 weights to 256 bytes for ROCm
182
+ "VLLM_ROCM_FP8_PADDING": lambda: bool(int(os.getenv("VLLM_ROCM_FP8_PADDING", "1"))),
183
+ # Pad the weights for the moe kernel
184
  diff --git a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py b/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
185
  index 9d93dfe54..dd96f6da4 100644
186
  --- a/vllm/model_executor/layers/fused_moe/experts/triton_moe.py
 
620
 
621
  # Run MQA — always pass layer scales. When KV cache is
622
  diff --git a/vllm/v1/attention/ops/merge_attn_states.py b/vllm/v1/attention/ops/merge_attn_states.py
623
+ index 20c7503e9..dc3ee0f97 100644
624
  --- a/vllm/v1/attention/ops/merge_attn_states.py
625
  +++ b/vllm/v1/attention/ops/merge_attn_states.py
626
+ @@ -3,9 +3,84 @@
 
 
627
 
 
 
628
  import torch
629
 
630
+ +from vllm import envs
631
  from vllm.platforms import current_platform
632
 
633
 
 
708
  def merge_attn_states(
709
  output: torch.Tensor,
710
  prefix_output: torch.Tensor,
711
+ @@ -46,6 +121,21 @@ def merge_attn_states(
712
  When provided, output must be FP8 dtype.
713
  """
714
 
715
  + if (
716
  + current_platform.is_rocm()
717
+ + and envs.VLLM_ROCM_SAFE_MERGE_ATTN_STATES
718
  + and output_scale is None
719
  + ):
720
  + return _merge_attn_states_torch(
runtime/SHA256SUMS.runtime CHANGED
@@ -1,3 +1,3 @@
1
- a14fcd3747829be98a73a453057bd451a492d87a236c2affcf57db1cbbf18386 COMMITS.txt
2
- 9cbcbbef5e113e22b753801ee141bd877b20c84f25f15cc159d0ae508aaf4324 Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
3
- c0c1446ea2c7334ef8b888f3104ec660d2ab41d2cb08bb9d90a3f589dba040e8 Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
 
1
+ 1bec28ea9c6bfa28aea0e0c7adf3d46b56eb399156680278fba0e5b5e3cc87b7 COMMITS.txt
2
+ 58dde46a772bedfaa7a2bd806a774db15352d5b53ad7c5c681b6857c273561b5 Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle
3
+ c6839a15d8250d7e9541f6a12b8cef0f9ccfbfdd825f095ee313048c56bdb0f4 Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.patch
scripts/build-vllm-gfx1151.sh CHANGED
@@ -52,7 +52,7 @@ venv="$install_root/.venv"
52
  source_dir="$install_root/vllm"
53
  index=https://rocm.nightlies.amd.com/whl-multi-arch/
54
  base=d35eb6c44071ea806018841c490f0d2f3219c485
55
- head=91fab4e6dcaa683911655630dae8a5b140a7cb91
56
  bundle="$package_root/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle"
57
  flash_source="$install_root/flash-attention"
58
  flash_commit=0e60e39473e8df549a20fb5353760f7a65b30e2d
@@ -160,7 +160,63 @@ export CMAKE_PREFIX_PATH="$sdk/lib/cmake:$site/torch/share/cmake"
160
  export HIP_FORCE_DEV_KERNARG=1
161
  export MAX_JOBS="${MAX_JOBS:-16}"
162
  export CMAKE_BUILD_PARALLEL_LEVEL="$MAX_JOBS"
163
- export CMAKE_ARGS="-DROCM_PATH=$sdk -DCMAKE_HIP_ARCHITECTURES=gfx1151 ${CMAKE_ARGS:-}"
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
164
 
165
  # Run the SDK checks against this virtual environment's development root, not
166
  # an inherited system ROCm installation.
 
52
  source_dir="$install_root/vllm"
53
  index=https://rocm.nightlies.amd.com/whl-multi-arch/
54
  base=d35eb6c44071ea806018841c490f0d2f3219c485
55
+ head=838616875c5dd4913d9f753221bf05f64cb7a7ed
56
  bundle="$package_root/runtime/Ling-3.0-Flash-CIRU-int4-Strix-native-vllm.bundle"
57
  flash_source="$install_root/flash-attention"
58
  flash_commit=0e60e39473e8df549a20fb5353760f7a65b30e2d
 
160
  export HIP_FORCE_DEV_KERNARG=1
161
  export MAX_JOBS="${MAX_JOBS:-16}"
162
  export CMAKE_BUILD_PARALLEL_LEVEL="$MAX_JOBS"
163
+ nixos_cmake_args=
164
+ host_id="$(
165
+ # Keep distro variables out of the installer environment.
166
+ # shellcheck disable=SC1091
167
+ source /etc/os-release 2>/dev/null || true
168
+ printf '%s\n' "${ID:-unknown}"
169
+ )"
170
+ if [ "$host_id" = nixos ]; then
171
+ # The SDK-wheel HIP compiler is intentionally self-contained and cannot
172
+ # discover Nix store C++ headers, CRT objects, libstdc++, or libnuma by
173
+ # following conventional /usr paths. Derive the active system toolchain
174
+ # paths instead of baking Ciru's store hashes into the public package.
175
+ command -v c++ >/dev/null || die 'NixOS build requires a C++ compiler in PATH'
176
+ command -v nix-store >/dev/null || die 'NixOS build requires nix-store in PATH'
177
+
178
+ crtbegin="$(c++ -print-file-name=crtbegin.o)"
179
+ libstdcpp="$(c++ -print-file-name=libstdc++.so)"
180
+ crt1="$(c++ -print-file-name=crt1.o)"
181
+ [ -f "$crtbegin" ] || die "cannot resolve the NixOS GCC toolchain: $crtbegin"
182
+ [ -f "$libstdcpp" ] || die "cannot resolve the NixOS libstdc++: $libstdcpp"
183
+ [ -f "$crt1" ] || die "cannot resolve the NixOS glibc runtime: $crt1"
184
+
185
+ gcc_toolchain="${crtbegin%%/lib/gcc/*}"
186
+ gcc_runtime_lib="$(dirname -- "$libstdcpp")"
187
+ gcc_crt_lib="$(dirname -- "$crtbegin")"
188
+ glibc_runtime_lib="$(dirname -- "$crt1")"
189
+ glibc_headers="$(
190
+ c++ -E -x c++ -v /dev/null 2>&1 |
191
+ awk '/^ \/nix\/store\/.*-glibc-.*-dev\/include\/?$/ && !found {
192
+ sub(/^ /, ""); found=$0
193
+ } END { print found }'
194
+ )"
195
+ [ -d "$glibc_headers" ] || \
196
+ die 'cannot resolve the NixOS glibc development headers from c++'
197
+
198
+ numactl_store="$(
199
+ nix-store -q --requisites /run/current-system |
200
+ awk '/\/nix\/store\/.*-numactl-[^/]+$/ && !found { found=$0 }
201
+ END { print found }'
202
+ )"
203
+ numa_lib="$numactl_store/lib"
204
+ [ -f "$numa_lib/libnuma.so" ] || \
205
+ die 'NixOS system closure does not provide libnuma; add pkgs.numactl'
206
+ dynamic_linker="$glibc_runtime_lib/ld-linux-x86-64.so.2"
207
+ [ -f "$dynamic_linker" ] || \
208
+ die "cannot resolve the NixOS dynamic linker: $dynamic_linker"
209
+
210
+ export LD_LIBRARY_PATH="$gcc_runtime_lib:$numa_lib:$LD_LIBRARY_PATH"
211
+ export CMAKE_LIBRARY_PATH="$numa_lib${CMAKE_LIBRARY_PATH:+:$CMAKE_LIBRARY_PATH}"
212
+ export HIPFLAGS="--rocm-path=$sdk --gcc-toolchain=$gcc_toolchain -idirafter $glibc_headers ${HIPFLAGS:-}"
213
+ export HIPCC_COMPILE_FLAGS_APPEND="--gcc-toolchain=$gcc_toolchain -idirafter $glibc_headers ${HIPCC_COMPILE_FLAGS_APPEND:-}"
214
+ export LDFLAGS="-B$glibc_runtime_lib -B$gcc_runtime_lib -L$glibc_runtime_lib -L$gcc_crt_lib -L$gcc_runtime_lib -Wl,--dynamic-linker=$dynamic_linker -Wl,-rpath,$gcc_runtime_lib -Wl,-rpath,$glibc_runtime_lib ${LDFLAGS:-}"
215
+ chmod +x "$script_dir/hipcc-nixos-linker-helper.sh"
216
+ nixos_cmake_args="-DHIP_HIPCC_CMAKE_LINKER_HELPER=$script_dir/hipcc-nixos-linker-helper.sh -DCMAKE_PROJECT_TOP_LEVEL_INCLUDES=$script_dir/rocm715-nixos-numa-target.cmake"
217
+ printf 'Enabled the derived NixOS HIP host-toolchain bridge.\n'
218
+ fi
219
+ export CMAKE_ARGS="-DROCM_PATH=$sdk -DCMAKE_HIP_ARCHITECTURES=gfx1151 $nixos_cmake_args ${CMAKE_ARGS:-}"
220
 
221
  # Run the SDK checks against this virtual environment's development root, not
222
  # an inherited system ROCm installation.
scripts/hipcc-nixos-linker-helper.sh ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ #!/usr/bin/env bash
2
+ set -Eeuo pipefail
3
+
4
+ # TheRock's generated helper uses /bin/bash, which NixOS intentionally omits.
5
+ # Invoke it through the bash selected by the package instead; BASH_SOURCE in
6
+ # the SDK helper still resolves HIP_PATH relative to the SDK.
7
+ exec bash \
8
+ "${ROCM_PATH:?ROCM_PATH must identify the TheRock SDK}/bin/hipcc_cmake_linker_helper" \
9
+ "$@"
scripts/install-host-deps.sh CHANGED
@@ -31,7 +31,14 @@ run_as_root() {
31
  fi
32
  }
33
 
34
- if [[ "$distro" =~ (ubuntu|debian) ]]; then
 
 
 
 
 
 
 
35
  run_as_root apt-get update
36
  run_as_root apt-get install -y \
37
  build-essential git curl ca-certificates pkg-config \
@@ -60,4 +67,9 @@ command -v uv >/dev/null || {
60
  exit 2
61
  }
62
 
63
- printf 'Host dependencies ready on %s.\n' "${PRETTY_NAME:-$distro}"
 
 
 
 
 
 
31
  fi
32
  }
33
 
34
+ if [[ "$distro" =~ (nixos) ]]; then
35
+ command -v nix-shell >/dev/null 2>&1 || {
36
+ echo 'NixOS requires nix-shell for the package build dependencies.' >&2
37
+ exit 2
38
+ }
39
+ nix-shell -p pkg-config xxd libdrm elfutils numactl openssl --run \
40
+ 'command -v pkg-config >/dev/null && command -v xxd >/dev/null'
41
+ elif [[ "$distro" =~ (ubuntu|debian) ]]; then
42
  run_as_root apt-get update
43
  run_as_root apt-get install -y \
44
  build-essential git curl ca-certificates pkg-config \
 
67
  exit 2
68
  }
69
 
70
+ if [[ "$distro" =~ (nixos) ]]; then
71
+ printf 'Host dependencies ready on demand through nix-shell on %s.\n' \
72
+ "${PRETTY_NAME:-$distro}"
73
+ else
74
+ printf 'Host dependencies ready on %s.\n' "${PRETTY_NAME:-$distro}"
75
+ fi
scripts/install-rocr-idle-fix.sh CHANGED
@@ -84,6 +84,29 @@ if [ "$force" != "--force" ] && [ -r "$runtime_path" ] && \
84
  exit 0
85
  fi
86
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
87
  for command_name in git gcc g++ pkg-config sha256sum xxd; do
88
  command -v "$command_name" >/dev/null 2>&1 || \
89
  die "missing $command_name; run scripts/install-host-deps.sh"
@@ -164,6 +187,33 @@ host_cmake_prefix="${CMAKE_PREFIX_PATH:-}"
164
  host_cmake_prefix="${host_cmake_prefix//:/;}"
165
  cmake_prefix="$sdk;$core${host_cmake_prefix:+;$host_cmake_prefix}"
166
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
167
  printf 'Building pinned ROCr idle-runtime override\n'
168
  printf ' source: %s @ %s\n compiler: %s / %s\n' \
169
  "$rocm_systems_repo" "$rocm_systems_commit" "$cc" "$cxx"
 
84
  exit 0
85
  fi
86
 
87
+ host_id="$(
88
+ # Keep distro variables out of the installer environment.
89
+ # shellcheck disable=SC1091
90
+ source /etc/os-release 2>/dev/null || true
91
+ printf '%s\n' "${ID:-unknown}"
92
+ )"
93
+ if [ "$host_id" = nixos ] && \
94
+ { ! command -v pkg-config >/dev/null 2>&1 || \
95
+ ! command -v xxd >/dev/null 2>&1; }; then
96
+ command -v nix-shell >/dev/null 2>&1 || \
97
+ die 'NixOS build needs nix-shell to provide pkg-config and xxd'
98
+ [ "${CIRU_NIX_ROCR_DEPS_ACTIVE:-0}" != 1 ] || \
99
+ die 'NixOS dependency shell did not provide pkg-config and xxd'
100
+ printf -v rerun_command '%q ' env CIRU_NIX_ROCR_DEPS_ACTIVE=1 \
101
+ bash "${BASH_SOURCE[0]}" "$install_root_arg"
102
+ if [ "$force" = --force ]; then
103
+ printf -v rerun_command '%s%q ' "$rerun_command" --force
104
+ fi
105
+ printf 'Entering an ephemeral Nix shell for ROCr host dependencies.\n'
106
+ exec nix-shell -p pkg-config xxd libdrm elfutils numactl openssl \
107
+ --run "$rerun_command"
108
+ fi
109
+
110
  for command_name in git gcc g++ pkg-config sha256sum xxd; do
111
  command -v "$command_name" >/dev/null 2>&1 || \
112
  die "missing $command_name; run scripts/install-host-deps.sh"
 
187
  host_cmake_prefix="${host_cmake_prefix//:/;}"
188
  cmake_prefix="$sdk;$core${host_cmake_prefix:+;$host_cmake_prefix}"
189
 
190
+ if [ "$host_id" = nixos ]; then
191
+ # CMake's source-tree find_library/find_path calls do not understand the
192
+ # Nix compiler wrapper flags. Convert the ephemeral shell's -L/-isystem
193
+ # inputs into ordinary search variables and add glibc explicitly.
194
+ glibc_runtime_lib="$(dirname -- "$(gcc -print-file-name=crt1.o)")"
195
+ [ -f "$glibc_runtime_lib/libc.so" ] || \
196
+ die "cannot resolve the NixOS libc linker input: $glibc_runtime_lib/libc.so"
197
+ nix_library_path="$(
198
+ printf '%s\n' "${NIX_LDFLAGS:-}" | tr ' ' '\n' |
199
+ sed -n 's/^-L\(\/nix\/store\/.*\)$/\1/p' | paste -sd: -
200
+ )"
201
+ nix_include_path="$(
202
+ printf '%s\n' "${NIX_CFLAGS_COMPILE:-}" | tr ' ' '\n' |
203
+ awk 'previous == "-isystem" && /^\/nix\/store\// { print }
204
+ { previous=$0 }' | paste -sd: -
205
+ )"
206
+ [ -n "$nix_library_path" ] || \
207
+ die 'NixOS dependency shell did not expose library search paths'
208
+ [ -n "$nix_include_path" ] || \
209
+ die 'NixOS dependency shell did not expose include search paths'
210
+ export CMAKE_LIBRARY_PATH="$glibc_runtime_lib:$nix_library_path${CMAKE_LIBRARY_PATH:+:$CMAKE_LIBRARY_PATH}"
211
+ export CMAKE_INCLUDE_PATH="$nix_include_path${CMAKE_INCLUDE_PATH:+:$CMAKE_INCLUDE_PATH}"
212
+ export LIBRARY_PATH="$glibc_runtime_lib:$nix_library_path${LIBRARY_PATH:+:$LIBRARY_PATH}"
213
+ export CPATH="$nix_include_path${CPATH:+:$CPATH}"
214
+ export LD_LIBRARY_PATH="$glibc_runtime_lib:$nix_library_path:$LD_LIBRARY_PATH"
215
+ fi
216
+
217
  printf 'Building pinned ROCr idle-runtime override\n'
218
  printf ' source: %s @ %s\n compiler: %s / %s\n' \
219
  "$rocm_systems_repo" "$rocm_systems_commit" "$cc" "$cxx"
scripts/rocm715-nixos-numa-target.cmake ADDED
@@ -0,0 +1,6 @@
 
 
 
 
 
 
 
1
+ if(NOT TARGET numa::numa)
2
+ find_library(NUMA_LIBRARY NAMES numa REQUIRED)
3
+ add_library(numa::numa UNKNOWN IMPORTED)
4
+ set_target_properties(numa::numa PROPERTIES
5
+ IMPORTED_LOCATION "${NUMA_LIBRARY}")
6
+ endif()
scripts/run-common.sh CHANGED
@@ -6,6 +6,8 @@ if [ "$#" -ne 0 ]; then
6
  exit 2
7
  fi
8
 
 
 
9
  release_name=Ling-3.0-Flash-CIRU-int4-Strix-native
10
  venv="${VENV:?Set VENV to the built runtime virtual environment}"
11
  model_path="${MODEL_PATH:?Set MODEL_PATH to the downloaded checkpoint directory}"
@@ -13,6 +15,21 @@ host="${HOST:-127.0.0.1}"
13
  port="${PORT:-18081}"
14
  cache_root="${CACHE_ROOT:-$HOME/.cache/ciru-ling3-vllm}"
15
  max_num_batched_tokens="${MAX_NUM_BATCHED_TOKENS:-8192}"
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
16
 
17
  case "$venv" in
18
  *[[:space:]:]*)
@@ -60,6 +77,14 @@ if [ ! -x "$sdk/bin/hipcc" ]; then
60
  exit 2
61
  fi
62
  install_root="$(cd -- "$(dirname -- "$venv")" && pwd -P)"
 
 
 
 
 
 
 
 
63
  rocr_override_dir="${ROCR_IDLE_RUNTIME_DIR:-$install_root/runtime-overrides/rocr-7.15-gcc-idle}"
64
  case "$rocr_override_dir" in
65
  *[[:space:]:]*)
@@ -100,6 +125,8 @@ export HSA_NO_SCRATCH_RECLAIM=1
100
  export VLLM_ROCM_USE_AITER=0
101
  export VLLM_ROCM_USE_SKINNY_GEMM=0
102
  export VLLM_ROCM_SAFE_MERGE_ATTN_STATES=1
 
 
103
  export VLLM_CACHE_ROOT="$cache_root/vllm"
104
  export TRITON_CACHE_DIR="$cache_root/triton"
105
  export HF_MODULES_CACHE="$cache_root/huggingface-modules"
@@ -149,4 +176,22 @@ fi
149
  printf 'Launching %s\n' "$release_name"
150
  printf ' vLLM: %s\n model: %s\n listen: %s:%s\n' \
151
  "$venv/bin/vllm" "$model_path" "$host" "$port"
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
152
  exec "$venv/bin/vllm" "${args[@]}"
 
6
  exit 2
7
  fi
8
 
9
+ script_dir="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
10
+ package_root="$(cd -- "$script_dir/.." && pwd)"
11
  release_name=Ling-3.0-Flash-CIRU-int4-Strix-native
12
  venv="${VENV:?Set VENV to the built runtime virtual environment}"
13
  model_path="${MODEL_PATH:?Set MODEL_PATH to the downloaded checkpoint directory}"
 
15
  port="${PORT:-18081}"
16
  cache_root="${CACHE_ROOT:-$HOME/.cache/ciru-ling3-vllm}"
17
  max_num_batched_tokens="${MAX_NUM_BATCHED_TOKENS:-8192}"
18
+ inductor_max_autotune="${VLLM_ENABLE_INDUCTOR_MAX_AUTOTUNE:-1}"
19
+ inductor_coordinate_descent="${VLLM_ENABLE_INDUCTOR_COORDINATE_DESCENT_TUNING:-1}"
20
+
21
+ for setting in \
22
+ "VLLM_ENABLE_INDUCTOR_MAX_AUTOTUNE:$inductor_max_autotune" \
23
+ "VLLM_ENABLE_INDUCTOR_COORDINATE_DESCENT_TUNING:$inductor_coordinate_descent"; do
24
+ case "${setting#*:}" in
25
+ 0|1) ;;
26
+ *)
27
+ printf '%s must be 0 or 1, got: %s\n' \
28
+ "${setting%%:*}" "${setting#*:}" >&2
29
+ exit 2
30
+ ;;
31
+ esac
32
+ done
33
 
34
  case "$venv" in
35
  *[[:space:]:]*)
 
77
  exit 2
78
  fi
79
  install_root="$(cd -- "$(dirname -- "$venv")" && pwd -P)"
80
+ package_head=unknown
81
+ if [ -r "$package_root/runtime/COMMITS.txt" ]; then
82
+ package_head="$(tail -n 1 "$package_root/runtime/COMMITS.txt" | cut -d ' ' -f 1)"
83
+ fi
84
+ installed_head=unavailable
85
+ if [ -d "$install_root/vllm/.git" ] && command -v git >/dev/null 2>&1; then
86
+ installed_head="$(git -C "$install_root/vllm" rev-parse HEAD 2>/dev/null || printf unavailable)"
87
+ fi
88
  rocr_override_dir="${ROCR_IDLE_RUNTIME_DIR:-$install_root/runtime-overrides/rocr-7.15-gcc-idle}"
89
  case "$rocr_override_dir" in
90
  *[[:space:]:]*)
 
125
  export VLLM_ROCM_USE_AITER=0
126
  export VLLM_ROCM_USE_SKINNY_GEMM=0
127
  export VLLM_ROCM_SAFE_MERGE_ATTN_STATES=1
128
+ export VLLM_ENABLE_INDUCTOR_MAX_AUTOTUNE="$inductor_max_autotune"
129
+ export VLLM_ENABLE_INDUCTOR_COORDINATE_DESCENT_TUNING="$inductor_coordinate_descent"
130
  export VLLM_CACHE_ROOT="$cache_root/vllm"
131
  export TRITON_CACHE_DIR="$cache_root/triton"
132
  export HF_MODULES_CACHE="$cache_root/huggingface-modules"
 
176
  printf 'Launching %s\n' "$release_name"
177
  printf ' vLLM: %s\n model: %s\n listen: %s:%s\n' \
178
  "$venv/bin/vllm" "$model_path" "$host" "$port"
179
+ printf ' package runtime: %s\n installed source: %s\n max autotune: %s\n' \
180
+ "$package_head" "$installed_head" "$inductor_max_autotune"
181
+ if [ "$package_head" != unknown ] && [ "$installed_head" != unavailable ] && \
182
+ [ "$package_head" != "$installed_head" ]; then
183
+ printf 'WARNING: the installed CIRU fork does not match this package.\n' >&2
184
+ printf 'Rebuild it without downloading the model:\n' >&2
185
+ printf ' bash %q --install-root %q --model-path %q\n' \
186
+ "$package_root/install.sh" "$install_root" "$model_path" >&2
187
+ fi
188
+ if [ "$inductor_max_autotune" = 1 ]; then
189
+ printf '%s\n' \
190
+ 'Cold-cache startup compiles and benchmarks kernels for several minutes.' \
191
+ 'The make_block_ptr/TypedStorage deprecations and AUTOTUNE tables are nonfatal.' \
192
+ 'The server is ready only after: Application startup complete.'
193
+ else
194
+ printf '%s\n' \
195
+ 'NOTICE: Inductor max autotuning is disabled; this is not the benchmarked profile.'
196
+ fi
197
  exec "$venv/bin/vllm" "${args[@]}"