Escha Runtime qwen3dense — escha 1.1.0+qwen3dense SGLang wheel, serve.sh, cookbook

.gitattributes

@@ -0,0 +1,36 @@
+*.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
+sglang/escha-1.1.0+qwen3dense-cp312-cp312-manylinux_2_28_x86_64.whl filter=lfs diff=lfs merge=lfs -text

LICENSE

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

README.md

@@ -0,0 +1,164 @@
+---
+license: apache-2.0
+tags:
+- quantization
+- sglang
+- cuda
+- inference
+- escha
+- qwen3
+- dense
+library_name: sglang
+---
+
+# Escha Runtime — `qwen3dense`
+
+By **[Escha Labs Inc.](https://eschalabs.com/)**
+
+The serving runtime for **Escha** 2-/3-bit (`escha`) quantized models of the **`qwen3_5` dense
+architecture** (Qwen3.8-27B and siblings). One repo per model architecture, one directory per
+engine — this architecture currently has **one** engine, [`sglang/`](sglang/).
+
+| | **SGLang** — [`sglang/`](sglang/) |
+|---|---|
+| Best for | everything: single user, teams, agents |
+| Concurrency | continuous batching, paged KV, optional radix prefix cache |
+| Tool calls / JSON schema / thinking parser | yes |
+| Interface | OpenAI-compatible (`/v1/chat/completions`, `/v1/completions`, `/v1/models`) |
+| Install | Python 3.12 venv + CUDA-12 PyTorch, then one wheel |
+
+The engine is a fork of [SGLang](https://github.com/sgl-project/sglang) bundled inside the wheel,
+running the Escha CUDA kernels. No separate `sglang` install is needed, and none should be
+present — the wheel ships its own.
+
+### Compatible models
+
+| Model repo | Bits |
+|---|---|
+| [EschaLabs/Qwen3.8-27B-Escha-W2](https://huggingface.co/EschaLabs/Qwen3.8-27B-Escha-W2) | 2-bit, mixed-rate (`escha`) |
+<!-- add each qwen3_5 dense Escha model you publish here -->
+
+> This runtime targets the **`qwen3_5` dense** architecture. Its wheel also happens to register the
+> `eschamoe` mixture-of-experts method, so a `qwen3_5_moe` model will load too — but the tuning,
+> the defaults in `sglang/serve.sh` and the documentation here are all written for the dense
+> architecture. For a mixture-of-experts model use
+> [`escha-runtime-qwen3moe`](https://huggingface.co/EschaLabs/escha-runtime-qwen3moe), whose
+> defaults are measured on it. A model of a genuinely different architecture will not load — use
+> the matching `escha-runtime-<arch>` repo.
+
+## Quickstart
+
+Full detail, including the per-GPU cookbook and troubleshooting:
+[`sglang/INSTALL.md`](sglang/INSTALL.md).
+
+```bash
+python3.12 -m venv .venv && source .venv/bin/activate
+pip install -U pip wheel
+pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu128   # cu12 torch FIRST
+pip install ./sglang/escha-*.whl   # pulls the bundled sglang fork + its full dep closure
+
+hf download EschaLabs/Qwen3.8-27B-Escha-W2 --local-dir ./Qwen3.8-27B-Escha-W2
+MODEL=./Qwen3.8-27B-Escha-W2 bash sglang/serve.sh
+```
+
+Then check the stack and the endpoint:
+
+```bash
+python -c "import torch, escha, sglang; print(torch.cuda.is_available(), hasattr(torch.ops.escha, 'escham_decode_gemv'), escha.__version__)"
+curl -s http://127.0.0.1:30000/v1/models | python3 -m json.tool
+```
+
+> `pip install "torch==2.9.*"` is a hard pin, not a suggestion. A bare `torch>=2.9` resolves to a
+> newer minor and `import escha` then fails with `undefined symbol: _ZN3c10...` — the compiled
+> extension is ABI-linked to libtorch, and that ABI is not stable across PyTorch minors.
+
+## Thinking, and why you probably want a budget
+
+This is a reasoning model. With thinking on, the reasoning arrives in `reasoning_content` and the
+answer in `content` — **read both**, or you will see half the response.
+
+Two per-request levers, both inside `chat_template_kwargs` (a *top-level* `enable_thinking` field
+is silently ignored):
+
+```jsonc
+{ "chat_template_kwargs": {"enable_thinking": true, "reasoning_effort": "xhigh"} }
+```
+
+`reasoning_effort` is `"xhigh"` (**the default**), `"medium"` or `"low"`; anything else makes the
+template raise, which surfaces as an HTTP 400 rather than a silent fallback. It works by injecting
+one sentence of system instruction — `xhigh` asks the model to validate assumptions and weigh
+alternatives, `low` asks it to keep thinking brief, and **`medium` injects nothing at all**, so
+`medium` is the neutral, unsteered model rather than a midpoint. It therefore *asks* for shorter
+reasoning; it does not bound it. If you are running a benchmark
+or an agent, set a **thinking budget** instead, which forces `</think>` after N reasoning tokens so
+an answer is always produced: see
+[`sglang/INSTALL.md` → Bounded thinking](sglang/INSTALL.md#bounded-thinking-thinking_budget) and
+[`sglang/thinking_budget.py`](sglang/thinking_budget.py). Without one, the usual failure is
+`finish_reason: "length"` with `content: null`, which a harness scores as *wrong* rather than as
+*slow*.
+
+## Requirements
+
+- **NVIDIA GPU, compute capability 8.0–12.0** (Ampere → Blackwell), Linux x86-64 with
+  **glibc ≥ 2.28**. The kernel launch route auto-selects per GPU; you never set it. Per-architecture
+  and per-VRAM launch recipes:
+  [`sglang/INSTALL.md` → Running on your GPU](sglang/INSTALL.md#running-on-your-gpu).
+- **Python 3.12** (the wheel is `cp312`-only) + **CUDA-12 PyTorch 2.9.x**. The wheel handles every
+  other dependency.
+- **A working host C compiler and Python dev headers.** Triton JIT-compiles a small shim at
+  CUDA-graph capture time — this is separate from `ptxas` and from a CUDA toolkit, so "driver only"
+  does not cover it. On slim container images a stripped `libisl` breaks `cc1` while
+  `gcc --version` still succeeds, and the failure surfaces ~40 s in as a `gcc`
+  `CalledProcessError` inside `cuda_graph_runner.py` — which reads like a runtime bug and is not.
+  Preflight in [`sglang/INSTALL.md`](sglang/INSTALL.md#requirements).
+- **24 GB VRAM** for the shipped defaults (65,536-token context, ~8–9 concurrent streams) with a
+  ~10.15 GB model. Note the stream ceiling: `MAMBA_RATIO=0.3` sizes the recurrent-state pool, which
+  clamps `max_running_requests` to 8–9 on a 24 GB card, so the `12`/`16` entries in the default
+  `CUDA_GRAPH_BS` are **dropped and never captured**. To serve more streams raise
+  `MAXREQ`/`MAXMAMBA` with `MEM` — the throughput recipe is in the
+  [model card](https://huggingface.co/EschaLabs/Qwen3.8-27B-Escha-W2#verified-configurations).
+  16 GB should fit at a reduced context — the cookbook has a recipe, but we have not run it.
+
+## Known limitations
+
+- **Serving throughput has been measured on three cards** — RTX 5090 (32 GB, sm_120), RTX 4090
+  (24 GB, sm_89) and RTX 3090 (24 GB, sm_86); the 5090 and 3090 by independent evaluators working
+  only from these docs. The 4090 is also where the model's thinking-on benchmarks (GPQA-Diamond,
+  LiveCodeBench) were produced; its thinking-off commonsense suite ran on an L40. The **16 GB and
+  40 GB+** tiers in the cookbook remain configuration guidance derived from the model size and the
+  wheel's architecture coverage, **not measurements**. Per-GPU numbers live on the
+  [model card](https://huggingface.co/EschaLabs/Qwen3.8-27B-Escha-W2#performance-across-gpus).
+- **On Ampere the auto-selected kernel route is the slower one at batch 1.** `ESCHA_ROUTE`
+  resolves to `lovelace` on sm_80/sm_86, but forcing `ESCHA_ROUTE=blackwell` measured **1.72×
+  faster single-stream on an RTX 3090** (23.6 → 40.7 tok/s, TPOT 42.4 → 24.6 ms) with identical
+  output. The two routes are bit-identical launch geometries, so this is safe to set; the gain is
+  batch-1-only (parity at 2–16). Serving one user on Ampere? Set it.
+- **`DETERMINISTIC=1` fails on consumer Blackwell (sm_120).** The deterministic attention kernel
+  requests 104 KB of shared memory per block, above the sm_120 limit, and the server exits during
+  startup. It works on Ampere, Ada and Hopper.
+- **Greedy output is not bit-reproducible across requests.** Batch composition changes fp16
+  accumulation order, so a near-tie can flip and a long reasoning chain diverges from there. Two
+  identical greedy requests may return different text. Use `DETERMINISTIC=1` when you need
+  reproducibility, and never A/B two configurations by diffing one generation.
+- **CUDA-graph batch sizes are capped at 32**, the decode kernel's maximum M
+  (`torch.ops.escha.escham_decode_gemv_max_m()`). The shipped default list stops at 16 because
+  that is where aggregate throughput peaks on a 4090; capture at `24`/`32` works and is worth it if
+  you serve that many streams. Past 32 a batch falls through to a large-M path meant for prefill,
+  so the runtime refuses to capture it rather than bake in the wrong kernel.
+- **`ATTN_BACKEND=triton` is required on consumer Blackwell (RTX 50-series).** The default
+  flashinfer backend asserts on this hybrid architecture at sm_120. The assertion names three
+  acceptable backends — `triton`, `trtllm_mha`, `fa4` — of which only `triton` has been run on
+  this model. Note that sm_120 shows steeper long-prompt decode decay than sm_89 (88.5% vs 96.3%
+  of short-prompt rate at a 5,000-token prompt); the attention path is the obvious suspect and
+  nobody has run the A/B that would confirm it.
+- **The model's stock chat template raises** on a few malformed conversations (a missing user
+  query, a system message in the wrong position). That surfaces as an HTTP 400 from the template,
+  not as a server fault.
+
+## License
+
+Everything here is released under the **Apache License, Version 2.0** — see [`LICENSE`](LICENSE).
+All bundled third-party code is permissive (Apache-2.0 / MIT / BSD-3-Clause) — **no copyleft**.
+Full texts and the component inventory:
+[`THIRD_PARTY_LICENSES/`](THIRD_PARTY_LICENSES/). Model weights are **not** in this repo and carry
+their own license in the model repository.

THIRD_PARTY_LICENSES/AQLM-NOTICE.txt

@@ -0,0 +1,15 @@
+AQLM — https://github.com/Vahe1994/AQLM
+
+Licensed under the Apache License, Version 2.0.
+Copyright the AQLM authors.
+
+The `escha` runtime's additive-codebook GEMV/HMMA kernels (`escha_aqlm_*`, compiled
+into `escha/_C`) are an independent implementation that follows the AQLM additive
+codebook *numeric format*; they are adapted from the escha runtime's own
+kernels rather than copied from AQLM source. This notice is provided out of an
+abundance of caution.
+
+The full text of the Apache License, Version 2.0 is reproduced in
+`SGLang-LICENSE.txt` in this directory (the license body is identical for every
+Apache-2.0 component in this release; only the per-component copyright line
+differs).

THIRD_PARTY_LICENSES/BSD-3-Clause.txt

@@ -0,0 +1,33 @@
+BSD 3-Clause License
+
+Applies to code within the bundled SGLang fork that is adapted from
+FlashAttention and Mamba:
+  Copyright (c) 2023, Tri Dao.
+  Copyright (c) 2024, Tri Dao, Albert Gu.
+(and other BSD-3-Clause-licensed upstreams whose per-file copyright notices are
+retained verbatim in the corresponding source headers within `sglang/`.)
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice, this
+   list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright notice,
+   this list of conditions and the following disclaimer in the documentation
+   and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its contributors
+   may be used to endorse or promote products derived from this software
+   without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

THIRD_PARTY_LICENSES/MIT.txt

@@ -0,0 +1,26 @@
+MIT License
+
+Applies to code within the bundled SGLang fork that carries an
+`SPDX-License-Identifier: MIT` header and/or is adapted from MIT-licensed
+upstreams (e.g. flash-linear-attention):
+  Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
+(and other MIT-licensed upstreams whose per-file copyright notices are retained
+verbatim in the corresponding source headers within `sglang/`.)
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

THIRD_PARTY_LICENSES/SGLang-LICENSE.txt

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

THIRD_PARTY_LICENSES/THIRD_PARTY_NOTICES.md

@@ -0,0 +1,62 @@
+# Third-Party Notices — Escha Runtime (`escha` 1.1.0)
+
+This repository distributes the **runtime** only (the `escha` wheel: quantization/serving
+kernels + a modified SGLang fork). SGLang is the only engine in this repository. The quantized **model weights** are distributed
+separately and carry their own license — see the model repository
+(**`EschaLabs/Qwen3.8-27B-Escha-W2`**).
+
+All bundled third-party code is under **permissive licenses (Apache-2.0, MIT, BSD-3-Clause)** —
+an audit found **no copyleft (GPL / LGPL / AGPL / MPL / SSPL) code**. The obligations below are
+satisfied by the files in this `THIRD_PARTY_LICENSES/` directory together with the copyright
+notices retained verbatim in the redistributed source (inside the wheel). The same texts are also
+embedded in the wheel metadata (`escha-1.1.0+qwen3dense.dist-info/licenses/LICENSE`).
+
+The `escha` runtime itself is released under the Apache License, Version 2.0 — see the top-level
+`LICENSE` file of this repository.
+
+---
+
+## 1. Acknowledgements — formats referenced, no code bundled
+
+The compiled CUDA runtime (`escha._C`) is **original work**. Its kernels implement, as
+independent (clean-room) reimplementations, numeric *formats* referenced from the two upstream
+projects below. **No upstream source is copied into the binary** and neither project's licence
+obliges us to reproduce it here — the attribution is a courtesy, and we would rather over-credit
+than under-credit work we learned from.
+
+| Upstream | Relationship | License | Text |
+|---|---|---|---|
+| **exllamav3** (Turboderp) | quantization codec *format* referenced by the escha decode kernels | MIT, © 2025 Turboderp | `exllamav3-LICENSE.txt` |
+| **AQLM** (Vahe1994/AQLM) | additive-codebook *format* referenced by `escha_aqlm_*` kernels | Apache-2.0 | `AQLM-NOTICE.txt` (text in `SGLang-LICENSE.txt`) |
+
+`bitblas` (Microsoft, MIT) is an **optional runtime import** only — not bundled or redistributed here.
+
+**Also gratefully acknowledged — the serving framework this engine is built on.**
+Unlike the formats above, its code is genuinely **redistributed** here, so it is covered by real
+licence obligations rather than courtesy: **[SGLang](https://github.com/sgl-project/sglang)**,
+bundled as a modified fork — see §2. This runtime would not exist in this form without it.
+
+## 2. Bundled serving engine — `sglang/` (modified fork)
+
+The wheel bundles a **modified fork of SGLang** as source.
+
+- **License:** Apache-2.0 — © 2023–2026 SGLang Team. Full text: `SGLang-LICENSE.txt`.
+- **Statement of changes (Apache-2.0 §4(b)):** the fork adds quantization-format support
+  (the `escha` dense and `eschamoe` mixture-of-experts methods, plus the int8 dense-layer
+  path) and kernel dispatch.
+- **NOTICE (Apache-2.0 §4(d)):** upstream sgl-project/sglang publishes **no NOTICE file**
+  (verified against the fork's pristine upstream-mirror at the merge base and across the full
+  tree); therefore none is reproduced.
+
+SGLang itself vendors code from other permissive projects; those files are redistributed
+**unmodified with their copyright headers intact**, which preserves their notices:
+
+| Upstream (as vendored in SGLang) | License | Reference text |
+|---|---|---|
+| vLLM project | Apache-2.0 | `SGLang-LICENSE.txt` |
+| IST-DASLab / Marlin (Elias Frantar) | Apache-2.0 | `SGLang-LICENSE.txt` |
+| EleutherAI + HuggingFace (transformers) | Apache-2.0 | `SGLang-LICENSE.txt` |
+| NVIDIA (CUTLASS / kernels) | BSD-3-Clause / Apache-2.0 | `BSD-3-Clause.txt` |
+| FlashAttention, Mamba (Tri Dao, Albert Gu) | BSD-3-Clause | `BSD-3-Clause.txt` |
+| flash-linear-attention (Songlin Yang, Yu Zhang) | MIT | `MIT.txt` |
+| Model configs (Qwen, DeepSeek, Moonshot, LG AI, Black Forest Labs, PaddlePaddle, Antgroup, …) | Apache-2.0 | `SGLang-LICENSE.txt` |

THIRD_PARTY_LICENSES/exllamav3-LICENSE.txt

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

sglang/INSTALL.md

@@ -0,0 +1,328 @@
+# Escha Runtime — install & serve
+
+This repository is the **runtime**: one wheel (`escha-*.whl`) holding the compiled `escha` package +
+a scrubbed `sglang` fork, plus the `serve.sh` launcher. It is **model-agnostic** — point it at any
+Escha model directory of a supported architecture. The quantized weights are a **separate**
+repository: **`EschaLabs/Qwen3.8-27B-Escha-W2`**
+(<https://huggingface.co/EschaLabs/Qwen3.8-27B-Escha-W2>).
+
+## Requirements
+
+**All are hard, and three of them fail quietly:**
+
+- **Linux x86-64 with glibc ≥ 2.28.** That covers RHEL / Rocky / Alma **8** and newer, Ubuntu 20.04+,
+  and Debian 10+. The wheel is tagged `manylinux_2_28_x86_64`, so pip refuses it up front on an
+  older distro instead of failing later at `import escha`.
+- **Python 3.12.** The wheel is built `cp312`; any other interpreter rejects it at install time
+  with "not a supported wheel on this platform". (A compiled-extension wheel is tied to one CPython
+  minor — there is intentionally no 3.11/3.13 build.)
+- **NVIDIA GPU, compute capability 8.0–12.0.** The build is a fat binary carrying SASS for
+  Ampere (`sm_80`/`sm_86`), Ada (`sm_89`, e.g. RTX 4090 / L40S), Hopper (`sm_90`), and Blackwell
+  (`sm_100`/`sm_120`), plus PTX for forward-compat JIT onto newer GPUs. Older cards (`< sm_80`,
+  e.g. Turing) are unsupported. Verify the installed wheel actually covers your card with
+  `cuobjdump --list-elf $(python -c "import escha,os;print(os.path.dirname(escha.__file__))")/_C*.so | grep -oE 'sm_[0-9]+' | sort -u`
+  — you should see all six. If it lists only one, it is a slim build and will fail on other GPUs
+  with "no kernel image is available for execution on the device" at the first launch (install and
+  the step-4 op-registration check still pass — the failure is at runtime).
+- **A working host C compiler and Python dev headers.** Triton JIT-compiles a `cuda_utils.c` shim
+  at CUDA-graph capture time and shells out to `cc`; this is separate from `ptxas` and from a CUDA
+  toolkit, so "driver only, no toolkit" does not cover it. Preflight:
+  ```bash
+  echo '#include <Python.h>' | gcc -xc -I"$(python -c 'import sysconfig;print(sysconfig.get_paths()["include"])')" -c - -o /dev/null
+  ```
+- **CUDA 12.x + matching PyTorch 2.9.x-for-CUDA-12.** `escha._C` is ABI-linked to libtorch and
+  `libcudart.so.12`, so the installed torch must be the CUDA-12 build of 2.9.x (a different torch
+  minor can fail to load `import escha`). Needs a driver new enough for CUDA 12.x.
+- **`transformers >= 5.8`.** Older releases lack the config class for this architecture. The server
+  does **not** abort — it logs a single `WARNING ... config normalize FAILED`, serves with
+  un-normalized architecture parameters, and then produces fluent-looking nonsense. The wheel
+  **pins `transformers>=5.8`**, so a plain `pip install ./escha-*.whl` upgrades an old transformers
+  for you; you only hit this if you *downgrade* transformers afterwards. If generation is garbage,
+  check your transformers version *first*. (Conversely, transformers 5.x logs a scary-looking
+  `WARNING ... you can try downgrading to transformers==4.57.1` about RoPE — that is a generic
+  upstream sglang warning and is **safe to ignore**; do **not** downgrade.)
+
+1. `python3.12 -m venv .venv && source .venv/bin/activate`
+2. `pip install -U pip wheel`
+3. **Install PyTorch from the CUDA-12 index first** (it must be the cu12 build — `escha._C` is
+   ABI-linked to it, and it is the one dependency that cannot come from PyPI). **Pin to 2.9.x**:
+   `pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu128`
+   A bare `torch>=2.9.0` resolves to a newer minor, and **nothing corrects it** — the wheel
+   declares 61 dependencies but *no* `torch` entry, and none of its transitive deps pin one. You
+   stay on the wrong minor and `import escha` fails later with `undefined symbol: _ZN3c10...`,
+   because that is the C++ ABI surface PyTorch does not keep stable across minor versions.
+4. Install the runtime wheel (it ships in this repo):
+   `pip install ./escha-*.whl`
+   — installs the `escha` runtime, the bundled `sglang` fork, **and the fork's full runtime
+   dependency closure** (transformers ≥ 5.8, flashinfer, xgrammar, msgspec, uvloop, … — all pinned
+   in the wheel). No separate `sglang[srt]` install is needed. Sanity-check the whole stack (all
+   three should print `True`):
+   ```bash
+   python -c "import torch, escha, sglang; print(torch.cuda.is_available(), hasattr(torch.ops.escha, 'escham_decode_gemv'), bool(sglang.__version__))"
+   ```
+   (Note `import sglang` in the check — a `torch`+`escha`-only check passes on a box that cannot
+   actually serve.)
+5. Download a model (the weights are a **separate** repo; it is a flat folder — safetensors, config
+   and tokenizer at the root, no nested subdir):
+   ```bash
+   pip install -U "huggingface_hub[cli]"
+   hf download EschaLabs/Qwen3.8-27B-Escha-W2 --local-dir ./Qwen3.8-27B-Escha-W2
+   ```
+6. Serve it — point `MODEL` at the downloaded model directory:
+   ```bash
+   MODEL=./Qwen3.8-27B-Escha-W2 bash serve.sh
+   ```
+   - **Kernel route is auto-selected at runtime by GPU** — no env var needed. The kernels ship a
+     single fatbin with two launch-geometry routes; `escha` queries the device compute capability
+     and picks the measured-optimal one: `blackwell` (barrier-light, SM-adaptive split-K) on
+     **sm_89 Ada, sm_90 Hopper, and sm_100/sm_120 Blackwell**; `lovelace` (original geometry) on
+     **sm_80/sm_86 Ampere**. Both are bit-identical; force one with
+     `ESCHA_ROUTE=lovelace|blackwell` to A/B on your card (Hopper's default is a best-guess —
+     worth A/B'ing). Check the active route with
+     `python -c "import escha,torch; torch.zeros(1,device='cuda'); print(torch.ops.escha.escham_route())"`
+     (`2`=blackwell, `1`=lovelace).
+   - See the top of `serve.sh` for the full knob list (`MEM`, `CTXLEN`, `GRAPHS`, `THINK`, `RADIX`,
+     `INT8`, `DETERMINISTIC`, …). The defaults are the profile every published number for this
+     model was measured on, on a 24 GB card.
+7. OpenAI-compatible endpoint on `:30000` (`http://127.0.0.1:30000/v1`). The served model id is
+   `SERVED_NAME` from `serve.sh` (override with `SERVED_NAME=...`); list it with:
+   ```bash
+   curl -s http://127.0.0.1:30000/v1/models | python3 -m json.tool
+   ```
+
+## Behaviour worth knowing
+
+- **Thinking is per-request.** `chat_template_kwargs: {"enable_thinking": false}` is honoured
+  against the default `THINK=1` server, with or without `tools`; a **top-level `enable_thinking`
+  field is ignored** — it must go inside `chat_template_kwargs`. Set `THINK=0` to serve the whole
+  server with thinking off.
+- **`reasoning_effort` is a second lever, and the default is the expensive one.**
+  `chat_template_kwargs: {"reasoning_effort": "xhigh"|"medium"|"low"}`, default **`"xhigh"`**. Each
+  level is one sentence of injected system instruction: `xhigh` asks the model to validate
+  assumptions and weigh alternatives, `low` asks it to keep thinking brief, and **`medium` injects
+  nothing** — it is the neutral, unsteered model, not a midpoint. If `xhigh` is over-thinking your
+  traffic, `medium` is usually the right first move, `low` for latency-sensitive volume. Any other
+  value makes the template raise, which surfaces as an HTTP 400 — it is a typo, not a server fault.
+  All three only apply while thinking is on, and none of them *bounds* anything: a
+  `thinking_budget` (below) is the only hard guarantee.
+- **Read `reasoning_content` AND `content`.** With thinking on, the reasoning arrives in
+  `reasoning_content` and the answer in `content`. A client that reads only one of them will see
+  half the response — this is the single most common integration mistake.
+- **Structured output is stable.** `serve.sh` defaults `JSON_WS=1`
+  (`--constrained-json-disable-any-whitespace`), which keeps `json_schema` values consistent
+  run-to-run. `JSON_WS=0` opts out.
+- **Over-long prompts are truncated, not rejected.** `serve.sh` defaults `TRUNCATE=1`
+  (`--allow-auto-truncate`). Without it, a prompt above `CTXLEN` is hard-rejected, and a harness
+  that scores the error as an answer reports it as a model failure. Set `TRUNCATE=0` if you would
+  rather see the error.
+- **`usage.reasoning_tokens`** counts the reasoning span on both the streaming and non-streaming
+  paths, and is 0 when no reasoning parser is active. Use it to size a `thinking_budget` from
+  what reasoning actually costs.
+- **`escha.__version__`** — `python -c "import escha; print(escha.__version__)"`. Read from the
+  installed distribution metadata, so it cannot drift from the wheel it shipped in. Paste it
+  into bug reports.
+- **`DETERMINISTIC=1`** sets `--enable-deterministic-inference` — **not available on consumer
+  Blackwell (RTX 50 / sm_120)**, where the deterministic attention kernel exceeds the 104 KB
+  shared-memory-per-block limit and the server exits at startup. `serve.sh` forwards any extra
+  arguments straight to `launch_server`, so a flag it does not model is still reachable.
+
+## Running on your GPU
+
+The weights are **10.15 GB**. What changes across cards is the **compute-capability knob** (one
+attention-backend override on the newest cards) and the **VRAM budget** (context length, the
+CUDA-graph batch list, and whether `lm_head` is kept int8-as-stored). The kernel launch route is
+chosen for you at runtime — you never set it.
+
+> **Measured on three cards: a 32 GB RTX 5090 (sm_120), a 24 GB RTX 4090 (sm_89) and a 24 GB
+> RTX 3090 (sm_86)**; the 4090 is also where the model's thinking-on benchmarks were produced.
+> Per-GPU tokens/s live on the
+> [model card](https://huggingface.co/EschaLabs/Qwen3.8-27B-Escha-W2#performance-across-gpus). The
+> **16 GB and 40 GB+** rows below are still *configuration* guidance derived from the model's size
+> and the wheel's architecture coverage, **not measurements** — we would rather ship a gap than a
+> number we have not run.
+
+### By compute capability (architecture)
+
+| Arch | Compute cap | Cards | What you set |
+|---|---|---|---|
+| Ampere | sm_80 / sm_86 | A100, A6000, RTX 3090 | route auto-selects `lovelace`, but **`ESCHA_ROUTE=blackwell` measured 1.72× faster at batch 1 on an RTX 3090** (23.6 → 40.7 tok/s; parity at batch 2–16) — set it if you serve a single user |
+| Ada | sm_89 | RTX 4090, L40S | nothing — route auto-selects `blackwell` (**the validated card**) |
+| Blackwell (consumer, 32 GB) | sm_120 | RTX 5090 | **`ATTN_BACKEND=triton`** — measured, see *By VRAM* below |
+| Hopper | sm_90 | H100 | nothing (route defaults `blackwell`; A/B `ESCHA_ROUTE` — Hopper is a best-guess) |
+| Blackwell (data-center) | sm_100 | B200 | nothing — route auto-selects `blackwell` |
+| Blackwell (consumer) | sm_120 | RTX 50-series | **`ATTN_BACKEND=triton`** (the default flashinfer backend asserts on this hybrid architecture at sm_120) |
+
+Confirm the fatbin covers your card and the active route with:
+```bash
+python -c "import escha,torch; torch.zeros(1,device='cuda'); print('route', torch.ops.escha.escham_route())"
+# expected: route 2 (blackwell) on sm_89/90/100/120 — incl. RTX 4090/5090; route 1 (lovelace) on sm_80/86
+# Both routes are bit-identical launch geometries — forcing one is a performance choice, not a
+# correctness one. Worth A/B'ing on Ampere (measured 1.72x at bs=1) and on Hopper (best-guess default).
+```
+
+### By VRAM
+
+This checkpoint is text-only, so `SGLANG_VLM_TEXT_ONLY=1` is the `serve.sh` default.
+
+**int8-as-stored `lm_head`.** The runtime can keep `lm_head` int8 in VRAM instead of dequantizing
+it to fp16 at load: ~1.2 GB instead of ~2.4 GB, read on **every** decode step, so it buys both VRAM
+and single-user decode bandwidth. `INT8=auto` (the default) turns it on at ≤ 24 GB VRAM
+(`SGLANG_INT8_AUTO_VRAM_GB`) and off above; `INT8=on` / `INT8=off` force it. Note this is narrower
+than the mixture-of-experts runtime: on this dense model only `lm_head` is affected —
+`embed_tokens` is a row *lookup*, so keeping it int8 would save VRAM but ~no decode bandwidth, and
+it stays on the dequant path.
+
+**CUDA graphs matter on this model** (`GRAPHS=1`, the `serve.sh` default). It is a hybrid
+architecture that launches many small kernels per token, so eager decode is launch-bound. Use
+`GRAPHS=0` only to debug a capture failure.
+
+| VRAM | Example | Suggested launch |
+|---|---|---|
+| **16 GB** | RTX 5060 Ti, RTX 5080 | `ATTN_BACKEND=triton MEM=0.90 CTXLEN=16384 CUDA_GRAPH_BS="1 2 4 8" bash serve.sh` — 10.15 GB of weights leaves a small pool; drop `CTXLEN` first if capture OOMs. Untested. |
+| **24 GB** | RTX 4090, RTX 3090, L40S, A6000 | `MODEL=... bash serve.sh` — the shipped defaults (`MEM=0.72 CTXLEN=65536 MAMBA_RATIO=0.3 RADIX=0 GRAPHS=1 CUDA_GRAPH_BS="1 2 4 8 12 16"`). **Validated on RTX 4090 and RTX 3090.** Single-user profile: the recurrent pool clamps it to **8–9 streams**, so the `12`/`16` capture entries are dropped. For throughput use `MEM=0.86 CTXLEN=32768 MAXREQ=32 MAXMAMBA=32 CUDA_GRAPH_BS="1 2 4 8 12 16 24 32"` (4090: 649 tok/s @ 16; 3090: 383). On a 3090 add `ESCHA_ROUTE=blackwell` for single-user work. |
+| **32 GB** | RTX 5090 | `ATTN_BACKEND=triton MEM=0.85 CTXLEN=65536 MAXREQ=32 MAXMAMBA=32 CUDA_GRAPH_BS="1 2 4 8 12 16 24 32" bash serve.sh` — **measured**: 87.1 tok/s bs=1, 955 tok/s peak at 16 streams, 28.2 GB peak of 32.6. `INT8=on` adds +7.2–7.4% decode at bs=1, paired A/B (and +12.2% KV tokens, not free VRAM — the saved ~1.2 GB is recycled into the pool). |
+| **40 GB+** | A100, H100, L40S-48 | `MEM=0.85 CTXLEN=131072 bash serve.sh`; int8 auto-OFF, raise `CTXLEN`/concurrency freely. Untested. |
+
+**Four settings that are worth getting right:**
+
+1. **`CUDA_GRAPH_BS` must include your maximum batch size — and must not exceed 32.** A batch size
+   that was not captured runs *eager* (the server logs `cuda graph: False` at that
+   `#running-req`), which reads roughly 15% low. 32 is the decode kernel's maximum M
+   (`python -c "import torch, escha; print(torch.ops.escha.escham_decode_gemv_max_m())"`); beyond it
+   a batch falls through to a large-M path meant for prefill, and the runtime declines to capture a
+   graph that would bake in the wrong kernel. The default list stops at 16 because that is where
+   aggregate throughput peaks on a 4090 — extend it to `"1 2 4 8 12 16 24 32"` if you actually serve
+   more than 16 concurrent streams.
+2. **Recurrent state bounds concurrency at short context; the KV pool bounds it at long context.**
+   Every stream holds ~0.15 GB of `ssm_state` regardless of context length — that is the limit with
+   short prompts. With long ones the KV pool binds first: `max_total_num_tokens` (printed at
+   startup) divided by your typical ISL+OSL is your real stream ceiling. **If the server logs a
+   `#running-req` below the batch you requested, the surplus is queueing and your TTFT includes
+   queue time** — measured on a 3090 at 2048/2048: asked 32, ran 18, mean TTFT 140 s, which reads
+   like a prefill catastrophe and is not one.
+3. **`RADIX=0` unless you have measured that prefix caching helps you.** On this hybrid
+   architecture the radix cache is incompatible with the overlap scheduler, so `RADIX=1` *silently
+   disables it* — the server logs `Disabling overlap schedule since mamba no_buffer is not
+   compatible with overlap schedule`. Keep `RADIX=1` only for multi-turn agent traffic that
+   genuinely reuses long prefixes.
+4. **If startup OOMs, step `MEM` DOWN, not up.** `mem-fraction-static` sizes the pool the KV cache
+   is carved from, so *too low* also fails, with the opposite message ("Not enough memory. Please
+   try to increase --mem-fraction-static"). The two failures are: pool-too-small at startup (raise
+   `MEM`) and capture-OOM after the weights load (lower `MEM`, or shorten `CUDA_GRAPH_BS`).
+
+**Concurrency.** This is a **hybrid** model: most layers are linear-attention with a per-request
+recurrent-state pool, so max concurrency is governed by that pool (`MAXMAMBA`), *not* by KV memory.
+The effective `max_running_requests` is clamped to `min(MAXREQ, MAXMAMBA // ratio)`, ratio ≈ 4 with
+the radix cache on, ≈ 1 with `RADIX=0`. A startup log line reports the clamp.
+
+## Bounded thinking (`thinking_budget`)
+
+**Read this before running any benchmark or agent against a thinking-ON server.** With thinking on,
+the model can spend the entire `max_tokens` budget reasoning and never reach its answer. You get:
+
+```
+finish_reason: "length",  content: null      # all tokens went to reasoning_content
+```
+
+An eval harness scores that as **wrong**, not as slow, and an agent framework sees a no-op. Most
+harnesses cap `max_tokens` at 512–4k, which is where this bites. How long the model thinks varies
+run to run, so a cap that works on one prompt can fail on the next — you want a *guarantee*, not a
+bigger number. `reasoning_effort: "low"` *asks* for less; it does not bound it.
+
+A **thinking budget** is that guarantee: after N reasoning tokens the server forces `</think>`, so
+the model must answer within the remaining budget. `serve.sh` enables the processor by default
+(`ENABLE_CLP=1`); pass the budget per request:
+
+```bash
+MODEL=/path/to/model bash serve.sh                                 # ENABLE_CLP=1 is the default
+python thinking_budget.py --model /path/to/model > clp.json        # one-time payload
+```
+
+```jsonc
+{ "messages": [ /* ... */ ], "max_tokens": 32768,
+  "chat_template_kwargs": {"enable_thinking": true, "reasoning_effort": "xhigh"},
+  "custom_logit_processor": "<contents of clp.json>",   // SGLang extension
+  "custom_params": {"thinking_budget": 28672} }
+```
+
+Leave real headroom between `thinking_budget` and `max_tokens` for the answer itself — at least
+~512 tokens, and far more for a task that answers with code. **Every published thinking-on
+benchmark number for this model (GPQA-Diamond, LiveCodeBench) was produced with `thinking_budget`
+28,672 inside `max_tokens` 32,768** (so 4,096 tokens always remain for the answer), at
+`reasoning_effort: "xhigh"`. The commonsense suite ran thinking-off and uses none of this.
+
+`thinking_budget.py` reads `<think>` / `</think>` from your model's own tokenizer, so it works for
+any Escha checkpoint that has a thinking mode. (SGLang's built-in
+`Qwen3ThinkingBudgetLogitProcessor` hardcodes Qwen3's token ids, which are wrong for this model's
+~248k vocabulary — hence the helper.) Verify against a running server with
+`python thinking_budget.py --model <dir> --demo`.
+
+If you would rather not think at all, `THINK=0` patches the model's own chat template to
+default-closed — and also suppresses the reasoning-effort instruction the template would otherwise
+inject into the system prompt — at some cost in reasoning-dependent accuracy.
+
+## Troubleshooting
+
+- **`subprocess.CalledProcessError` from `gcc` during "Capture cuda graph".** Not an escha or
+  sglang fault — your host C compiler cannot build triton's shim, and the traceback's visible
+  frames (`cuda_graph_runner.py`, `triton_backend.py`) all point elsewhere. Run the compiler
+  preflight in *Requirements*. On slim container images a `libisl` registered in dpkg but missing
+  from disk breaks `cc1` while `gcc --version` still works.
+
+- **`finish_reason: "length"` with `content: null` on a thinking-ON server.** The reasoning ran past
+  your `max_tokens`. Raise the cap or, better, set a
+  [thinking budget](#bounded-thinking-thinking_budget) so an answer is always produced.
+
+- **HTTP 400 mentioning "Unexpected reasoning effort".** `reasoning_effort` must be exactly
+  `xhigh`, `medium` or `low`. The chat template raises on anything else.
+
+- **`... has no SGLang implementation and the Transformers implementation is not compatible`.** This
+  error is usually a **red herring** — the fork *does* register this architecture; the real cause is
+  an exception while *importing* the model module (typically an over-new `nvidia-cutlass-dsl`,
+  which the wheel pins `<4.4.0`). Surface the true traceback with:
+  ```bash
+  python -c "import sglang.srt.models.qwen3_5"
+  ```
+
+- **`ValueError: Invalid quantization method ...`.** The wheel's registry ships exactly two
+  methods, `escha` (dense) and `eschamoe` (mixture-of-experts). Check your model's
+  `quantize_config.json` / `config.json` names one of them, and that you installed *this* wheel
+  (`python -c "import escha; print(escha.__version__)"` should print `1.1.0+qwen3dense`) — the
+  1.0.x wheels registered `eschamoe` only and cannot load a dense model at all.
+
+- **`ModuleNotFoundError: pybase64` / `msgspec` / `xgrammar` / `flashinfer` at startup.** The wheel
+  declares the full runtime dependency closure, so a clean `pip install ./escha-*.whl` pulls
+  these automatically. If you see this, you likely installed the wheel with `--no-deps` or into a
+  venv where pip could not resolve them — reinstall without `--no-deps`.
+
+- **`ImportError: undefined symbol: _ZN3c10...` on `import escha`.** Wrong torch. The wheel does
+  not pin it, so an unpinned install leaves you on a newer minor. Check with
+  `python -c "import torch;print(torch.__version__)"` — it must be `2.9.x+cu12`. Fix:
+  `pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu128`.
+
+- **`PermissionError` on a CUDA `bin` directory during startup.** `TRITON_PTXAS_PATH` must point at
+  the ptxas **binary**, not at the directory containing it — `[[ -x <dir> ]]` is true for
+  directories, so a shell profile that sets the directory passes a naive check and then fails
+  inside triton. `serve.sh` detects and repairs this case; if you launch `launch_server` directly,
+  fix the variable yourself (or just unset it — triton ships its own ptxas).
+
+- **A local benchmark client reporting HTTP 504 on `127.0.0.1`.** Check for an `HTTP_PROXY` /
+  `ALL_PROXY` in your environment without `no_proxy=127.0.0.1`. Proxies time out multi-minute
+  requests and synthesize a 504 while the server is still working fine.
+
+- **OOM at startup naming a PID that is not yours.** A previous `sglang::scheduler` child is still
+  holding the GPU — killing the `launch_server` parent does not reap it. Check
+  `nvidia-smi --query-compute-apps=pid,used_memory --format=csv` and kill that PID. (Match on the
+  process *name*; a `pkill -f sglang` pattern also matches the shell running it.)
+
+- **Garbage / fluent-nonsense generations.** Check `transformers` first (must be ≥ 5.8); the
+  transformers 5.x "downgrade to 4.57.1" RoPE warning is benign — do not act on it.
+
+- **Two greedy requests give different text.** Expected: greedy is not bit-reproducible across
+  requests, because fp16 accumulation order depends on batch composition and a long reasoning chain
+  diverges once a near-tie flips. Use `DETERMINISTIC=1` when you need reproducibility.
+
+- **Structured output (`response_format` / JSON schema) on a memory-tight card.** The grammar path
+  needs ~0.5 GB of extra headroom the first time it runs (xgrammar + its sampler bookkeeping). On a
+  card packed to the edge, budget for it — lower `CTXLEN` a notch if a first structured request is
+  tight.

sglang/escha-1.1.0+qwen3dense-cp312-cp312-manylinux_2_28_x86_64.whl

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

sglang/serve.sh

@@ -0,0 +1,313 @@
+#!/usr/bin/env bash
+# Serve a Qwen3.8-27B 2-/3-bit `escha` (dense) export with escha-sglang, OpenAI-compatible,
+# usable by any OpenAI-compatible client. The `escha` quant method is auto-detected from
+# config.json — nothing here needs to name it.
+#
+# Linux x86-64 + NVIDIA CUDA 12.x only (native or WSL2). See ../README.md and INSTALL.md.
+#
+# Required: set MODEL to your downloaded model dir (a flat folder of safetensors +
+# config + tokenizer — the model repo has no nested subfolder). Everything else has
+# safe, MEASURED defaults for a 24 GB card.
+#   MODEL=./Qwen3.8-27B-Escha-W2 bash serve.sh
+set -euo pipefail
+
+# ---- paths ----------------------------------------------------------------------------------
+MODEL=${MODEL:?set MODEL=/path/to/your/escha model dir}
+SERVED_NAME=${SERVED_NAME:-escha-qwen38-27b-w2}   # opencode/LM Studio must use this id verbatim
+VENV=${VENV:-}              # optional: path to your venv; if set, we activate it
+ESCHA_SRC=${ESCHA_SRC:-}    # optional: path to escha *source* tree's src/ (only if not pip-installed)
+
+# ---- server ---------------------------------------------------------------------------------
+HOST=${HOST:-127.0.0.1}     # 0.0.0.0 to serve other machines — do this behind a VPN or tunnel and
+                            #     set API_KEY; never open the raw port to the internet.
+PORT=${PORT:-30000}
+API_KEY=${API_KEY:-}        # REQUIRED if HOST=0.0.0.0 / exposed: clients must send this as the
+                            #     Bearer key. Empty = no auth (fine only for localhost).
+MEM=${MEM:-0.72}            # mem-fraction-static. 0.72 is the MEASURED default for a 24 GB card at
+                            #     CTXLEN 65536 with CUDA graphs on: capture needs headroom, and 0.78
+                            #     left none once the int8 head is resident. Too LOW is the other
+                            #     failure mode — the server exits at startup with "Not enough memory.
+                            #     Please try to increase --mem-fraction-static". If capture OOMs,
+                            #     step DOWN 0.70 / 0.68 rather than up. On a >24 GB card raise it.
+CTXLEN=${CTXLEN:-65536}     # the model's own limit is 262144; 65536 is what fits a 24 GB card
+                            #     alongside the weights. Raise on a bigger card, lower with MEM.
+CHUNK=${CHUNK:-2048}        # chunked-prefill-size. Bounds the activation spike on a long prompt;
+                            #     leave it alone unless you are tuning TTFT on a large card.
+MAXREQ=${MAXREQ:-0}         # max concurrent requests; 0 = auto.
+MAXMAMBA=${MAXMAMBA:-0}     # --max-mamba-cache-size; 0 = auto.
+MAMBA_RATIO=${MAMBA_RATIO:-0.3}  # --mamba-full-memory-ratio. 0.3 (not the 0.9 default) is what
+                            #     leaves the pool room for KV *and* CUDA-graph capture at CTXLEN
+                            #     65536 on 24 GB. Raise it on a card with spare VRAM.
+
+# ---- behavior -------------------------------------------------------------------------------
+GRAPHS=${GRAPHS:-1}        # 1 = CUDA graphs ON (DEFAULT). This is a hybrid model: many small
+                           #     kernels per token, so eager decode is launch-bound. Set GRAPHS=0
+                           #     only to debug a capture failure; expect several-fold slower decode.
+CUDA_GRAPH_BS=${CUDA_GRAPH_BS:-"1 2 4 8 12 16"}  # batch sizes captured when GRAPHS=1.
+                           #     HEADS UP: at the DEFAULT MEM/MAMBA_RATIO below, a 24 GB card's
+                           #     recurrent pool admits only 8-9 streams, so max_running_requests
+                           #     clamps there and the 12/16 entries here are DROPPED (the startup
+                           #     log prints the list actually captured). They start mattering once
+                           #     you raise MAXREQ/MAXMAMBA/MEM. 32 is the decode kernel's max M, so
+                           #     "1 2 4 8 12 16 24 32" captures fine on a throughput profile. Past
+                           #     32 a batch falls through to a large-M path meant for prefill, and
+                           #     the server refuses to capture it.
+THINK=${THINK:-1}          # 1 = thinking-ON + --reasoning-parser qwen3 (simplest, always works).
+                           # 0 = thinking-OFF: patches the model's chat template to default-closed
+                           #     AND drops the reasoning parser.
+                           # A per-request chat_template_kwargs {"enable_thinking": false} is
+                           #     honoured on a THINK=1 server, so THINK=0 is not required for
+                           #     clients that disable thinking per call. With thinking ON, the
+                           #     reasoning arrives in `reasoning_content` and the answer in
+                           #     `content`: read BOTH.
+                           # This model also takes chat_template_kwargs {"reasoning_effort":
+                           #     "xhigh"|"medium"|"low"} (default xhigh). An unrecognised value
+                           #     makes the template raise, which surfaces as a 400.
+RADIX=${RADIX:-0}          # 0 = prefix caching OFF (DEFAULT here). On this hybrid architecture the
+                           #     radix cache without speculative decoding disables the overlap
+                           #     scheduler, which costs more than the prefix reuse returns. Set
+                           #     RADIX=1 only if you have measured a win on your own workload.
+ESCHA_ROUTE=${ESCHA_ROUTE:-}  # kernel launch geometry; empty = auto-select per GPU. The two
+                           #     routes are BIT-IDENTICAL, so this is purely a perf choice. On
+                           #     AMPERE (sm_80/86) auto picks `lovelace` and that is the SLOWER one
+                           #     at batch 1: an RTX 3090 measured 23.6 -> 40.7 tok/s (1.72x) with
+                           #     ESCHA_ROUTE=blackwell. Parity at batch 2-16, so set it for
+                           #     single-user serving and leave it auto for batched.
+if [[ -n "$ESCHA_ROUTE" ]]; then export ESCHA_ROUTE; fi   # NB: `[ .. ] && export` would abort
+                                                        # the script under `set -e` when unset.
+ATTN_BACKEND=${ATTN_BACKEND:-}  # full-attention backend override. Empty = sglang default (fine on
+                           #     Ada/4090). REQUIRED on consumer Blackwell (RTX 50 / sm_120): the
+                           #     default resolves to flashinfer but the fork asserts triton/trtllm_mha
+                           #     /fa4 for hybrid models on Blackwell -> set ATTN_BACKEND=triton.
+TOOL_PARSER=${TOOL_PARSER:-qwen3_coder}   # matches this template's <function=..>/<parameter=..> XML.
+REASONING_PARSER=${REASONING_PARSER:-qwen3}
+ENABLE_CLP=${ENABLE_CLP:-1}   # 1 = --enable-custom-logit-processor (DEFAULT). Needed for the
+                           #     per-request thinking budget (see thinking_budget.py); it is how
+                           #     every published benchmark number for this model was produced.
+TRUNCATE=${TRUNCATE:-1}    # 1 = --allow-auto-truncate (DEFAULT). A prompt longer than CTXLEN is
+                           #     truncated instead of HARD-REJECTED. Without it an over-long prompt
+                           #     returns an error that reads like a model failure.
+JSON_WS=${JSON_WS:-1}      # 1 = --constrained-json-disable-any-whitespace (DEFAULT).
+                           #     With permissive whitespace, a json_schema `integer`/`number`
+                           #     property that follows a string property can take a spurious
+                           #     leading `-`: the grammar allows a "\n    " branch after the
+                           #     colon, and from that off-distribution state the sign/digit choice
+                           #     is near-tied. Compact JSON is what structured-output consumers
+                           #     want anyway. `JSON_WS=0` opts out.
+DETERMINISTIC=${DETERMINISTIC:-0}  # 1 = --enable-deterministic-inference. Greedy is NOT
+                           #     bit-reproducible across requests by default (fp16 accumulation
+                           #     order depends on batch composition): long reasoning chains diverge
+                           #     once a near-tie flips. Set 1 for eval reproducibility; expect a
+                           #     throughput cost. NOT AVAILABLE on consumer Blackwell (RTX 50 /
+                           #     sm_120): the deterministic attention kernel asks for 104 KB of
+                           #     shared memory per block, above the sm_120 limit, and the server
+                           #     exits at startup. Works on Ampere / Ada / Hopper.
+
+# ---- environment the escha stack needs ------------------------------------------------------
+[[ -n "$VENV" && -f "$VENV/bin/activate" ]] && source "$VENV/bin/activate"
+PYBIN=$(command -v python)
+TORCH_LIB=$("$PYBIN" -c 'import torch,os;print(os.path.join(os.path.dirname(torch.__file__),"lib"))')
+export LD_LIBRARY_PATH="$TORCH_LIB${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"   # escha._C links libc10
+[[ -n "$ESCHA_SRC" ]] && export PYTHONPATH="$ESCHA_SRC${PYTHONPATH:+:$PYTHONPATH}"
+export SGLANG_MAMBA_CONV_DTYPE=float16          # must match --dtype float16 or causal_conv1d crashes
+export SGLANG_DISABLE_CUDNN_CHECK=1
+# This checkpoint is TEXT-ONLY (the vision tower is in the quant `ignore` list — the checkpoint
+# has zero visual.* tensors). Default to NOT instantiating the vision tower: it saves VRAM on
+# every card, and prevents an image request from silently decoding through a random-init tower
+# (garbage instead of an error).
+export SGLANG_VLM_TEXT_ONLY=${SGLANG_VLM_TEXT_ONLY:-1}
+export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True   # correct var name — cuts reserved-vs-allocated
+export PYTORCH_ALLOC_CONF=expandable_segments:True        # (alias on newer torch; keep both)
+export HF_HUB_OFFLINE=1
+
+# ---- decode/prefill kernel profile ----------------------------------------------------------
+# This is the exact profile every published number for this model was measured on. All five knobs
+# are pinned rather than left to their defaults, because a stale value inherited from a shell
+# would read as a quality regression rather than as a configuration difference.
+#
+# ESCHA_STRICT=1 — refuse to silently fall back to the reference decode path; raise instead.
+#   The two flags are mutually exclusive AT IMPORT, so the opposite one must not be in the
+#   environment: unset it rather than trusting the caller's shell.
+unset ESCHA_FORCE_REF 2>/dev/null || true
+export ESCHA_STRICT=${ESCHA_STRICT:-1}
+# ESCHA_MULTI=1 — one fused launch per multi-shard projection (q+k+v, gate+up) instead of a
+#   per-shard loop. This is the default; pinned because it is load-bearing for decode latency.
+export ESCHA_MULTI=${ESCHA_MULTI:-1}
+# ESCHA_PREFILL=fused — run prefill through the fused GEMM rather than a dense reconstruction.
+export ESCHA_PREFILL=${ESCHA_PREFILL:-fused}
+# ESCHA_PREFILL_ACC=mixed / ESCHA_PREFILL_FP16ACC=0 — prefill accumulation precision. Under
+#   ESCHA_STRICT every shard's prefill goes through this path, so these govern prefill numerics
+#   for the whole model. "mixed" is the validated setting.
+export ESCHA_PREFILL_ACC=${ESCHA_PREFILL_ACC:-mixed}
+export ESCHA_PREFILL_FP16ACC=${ESCHA_PREFILL_FP16ACC:-0}
+# ESCHA_DENSE_FUSED_EPI=0 — the fused split-K epilogue tail is OFF. It is not a numerics-neutral
+#   speedup on this model (a measured multi-point accuracy regression on MATH-500), and it is the
+#   only process-global mutable state on the decode side: a leaked environment value would look
+#   like a quantization regression. Pinned off.
+export ESCHA_DENSE_FUSED_EPI=${ESCHA_DENSE_FUSED_EPI:-0}
+# ESCHA_ROUTE is now a first-class knob above (see the AMPERE note there) — it is no longer
+# "not set here". Other passthrough knobs, still not set here: ESCHA_KGROUP=0 reverts mixed-rate
+# projections to a per-shard launch (debug only).
+
+# ---- int8-as-stored lm_head -----------------------------------------------------------------
+# The runtime can keep lm_head INT8 in VRAM ("int8-as-stored") instead of dequantizing it to fp16
+# at load. On this dense model the fp16 head is ~2.4 GB of the bytes read EVERY decode step vs
+# ~1.2 GB as stored, so it is worth real single-user throughput plus ~1.2 GB of VRAM. Unlike the
+# mixture-of-experts runtime, only lm_head is affected here: embed_tokens is a row LOOKUP, so
+# keeping it int8 would save VRAM but ~no decode bandwidth, and it stays on the dequant path.
+#   INT8=auto (default) — VRAM-threshold rule: ON at <= SGLANG_INT8_AUTO_VRAM_GB (24), OFF above.
+#   INT8=on   — force ON:  single-user / latency-sensitive serving on ANY card size.
+#   INT8=off  — force OFF: batched / high-concurrency throughput serving on >24 GB cards.
+INT8=${INT8:-auto}
+case "$INT8" in
+  on)  export SGLANG_INT8_LM_HEAD=1 ;;
+  off) export SGLANG_INT8_LM_HEAD=0 ;;
+  auto) : ;;   # leave SGLANG_INT8_LM_HEAD as the caller set it (or unset -> AUTO)
+  *) echo "[serve] ERROR: INT8 must be auto|on|off (got '$INT8')" >&2; exit 1 ;;
+esac
+
+# ---- ptxas discovery ------------------------------------------------------------------------
+#  * Triton — which torch already pulls in — SHIPS a matching ptxas at
+#    site-packages/triton/backends/nvidia/bin/ptxas. Probe that FIRST: a CUDA toolkit is NOT a
+#    requirement of this runtime.
+#  * Validate whatever the caller exported. A common shell profile sets TRITON_PTXAS_PATH to the
+#    CUDA *bin directory*; `[[ -x <dir> ]]` is TRUE for directories, so a naive check accepts it
+#    and triton then dies mid-startup with a permission error on that directory. Require a
+#    regular file (-f) AND executable (-x); if a directory was given, try <dir>/ptxas first.
+_ptxas_ok() { [[ -n "${1:-}" && -f "$1" && -x "$1" ]]; }
+if ! _ptxas_ok "${TRITON_PTXAS_PATH:-}"; then
+  if [[ -n "${TRITON_PTXAS_PATH:-}" ]]; then
+    if _ptxas_ok "${TRITON_PTXAS_PATH%/}/ptxas"; then
+      TRITON_PTXAS_PATH="${TRITON_PTXAS_PATH%/}/ptxas"
+      echo "[serve] TRITON_PTXAS_PATH was a directory; using $TRITON_PTXAS_PATH"
+    else
+      echo "[serve] ignoring TRITON_PTXAS_PATH='$TRITON_PTXAS_PATH' (not an executable file)"
+      TRITON_PTXAS_PATH=""
+    fi
+  fi
+  if [[ -z "${TRITON_PTXAS_PATH:-}" ]]; then
+    for p in \
+      "$("$PYBIN" -c 'import os,triton;print(os.path.join(os.path.dirname(triton.__file__),"backends","nvidia","bin","ptxas"))' 2>/dev/null)" \
+      /usr/local/cuda*/bin/ptxas \
+      "$(command -v ptxas 2>/dev/null)"; do
+      _ptxas_ok "$p" && { TRITON_PTXAS_PATH="$p"; break; }
+    done
+  fi
+  export TRITON_PTXAS_PATH
+fi
+: "${TRITON_PTXAS_PATH:?no usable ptxas found. It normally ships inside the venv with triton (a torch dependency) — check that \`python -c \"import triton\"\` works. Otherwise point TRITON_PTXAS_PATH at a ptxas BINARY (not its directory).}"
+
+# Only if you hit it: the first time triton compiles a kernel (during warmup AND during CUDA-graph
+# capture), the multi-GB scheduler must fork() to spawn ptxas. Under the default overcommit
+# heuristic that fork can fail -> the server dies with a triton JIT memory-allocation failure at
+# kernel-compile time. Setting vm.overcommit_memory=1 fixes it. Many boxes never hit this (a warm
+# triton cache, or enough free RAM), so it is advisory. In a container the sysctl is often
+# read-only — then it is the HOST's job.
+if [[ "$(cat /proc/sys/vm/overcommit_memory 2>/dev/null)" == "0" ]]; then
+  echo "[serve] note: if startup dies with a triton JIT memory-allocation failure at kernel-" >&2
+  echo "[serve]       compile time, set 'sudo sysctl -w vm.overcommit_memory=1' (persist via" >&2
+  echo "[serve]       /etc/sysctl.d/99-escha.conf). In a read-only container this must be set" >&2
+  echo "[serve]       on the host; otherwise it's safe to ignore." >&2
+fi
+
+# Warn when the GPU is already substantially occupied (a foreign server / another model on a
+# shared box is the #1 cause of an immediate, confusing pool-allocation failure at startup).
+_USED_MB=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits 2>/dev/null | head -1 | tr -dc '0-9')
+if [[ -n "$_USED_MB" && "$_USED_MB" -gt 1024 ]]; then
+  echo "[serve] WARNING: GPU already has ${_USED_MB} MiB in use by other processes (nvidia-smi)." >&2
+  echo "[serve]          The static pool sizing assumes a free GPU — expect an allocation failure" >&2
+  echo "[serve]          if another server is resident. Check: nvidia-smi" >&2
+fi
+
+# ---- assemble flags -------------------------------------------------------------------------
+# Piecewise CUDA graphs stay OFF on this model family in every profile (the fork's piecewise path
+# does not handle its positional encoding), so it is not a knob.
+GRAPH_ARGS=(--disable-piecewise-cuda-graph)
+[[ "$GRAPHS" == "1" ]] && GRAPH_ARGS+=(--cuda-graph-bs $CUDA_GRAPH_BS) || GRAPH_ARGS+=(--disable-cuda-graph)
+
+POOL_ARGS=()
+[[ "$MAXREQ"   != "0" ]] && POOL_ARGS+=(--max-running-requests "$MAXREQ")
+[[ "$MAXMAMBA" != "0" ]] && POOL_ARGS+=(--max-mamba-cache-size "$MAXMAMBA")
+[[ -n "$MAMBA_RATIO" ]] && POOL_ARGS+=(--mamba-full-memory-ratio "$MAMBA_RATIO")
+RADIX_ARGS=(); [[ "$RADIX" == "0" ]] && RADIX_ARGS+=(--disable-radix-cache)
+ATTN_ARGS=(); [[ -n "$ATTN_BACKEND" ]] && ATTN_ARGS+=(--attention-backend "$ATTN_BACKEND")
+AUTH_ARGS=(); [[ -n "$API_KEY" ]] && AUTH_ARGS+=(--api-key "$API_KEY")
+CLP_ARGS=(); [[ "$ENABLE_CLP" == "1" ]] && CLP_ARGS+=(--enable-custom-logit-processor)
+TRUNC_ARGS=(); [[ "$TRUNCATE" == "1" ]] && TRUNC_ARGS+=(--allow-auto-truncate)
+JSON_ARGS=(); [[ "$JSON_WS" == "1" ]] && JSON_ARGS+=(--constrained-json-disable-any-whitespace)
+DET_ARGS=(); [[ "$DETERMINISTIC" == "1" ]] && DET_ARGS+=(--enable-deterministic-inference)
+if [[ "$HOST" != "127.0.0.1" && -z "$API_KEY" ]]; then
+  echo "[serve] WARNING: HOST=$HOST is exposed but API_KEY is empty — no auth. Set API_KEY." >&2
+fi
+
+REASON_ARGS=(); TEMPLATE_ARGS=()
+if [[ "$THINK" == "1" ]]; then
+  REASON_ARGS+=(--reasoning-parser "$REASONING_PARSER")
+else
+  # thinking-OFF: generate a patched chat template (default-closed <think>) from the model's own.
+  NOTHINK="${TMPDIR:-/tmp}/escha_nothink_$(basename "$MODEL").jinja"
+  "$PYBIN" - "$MODEL" "$NOTHINK" <<'PY'
+import json, os, sys
+mdir, out = sys.argv[1], sys.argv[2]
+src = os.path.join(mdir, "chat_template.jinja")
+tmpl = open(src).read() if os.path.exists(src) else json.load(open(os.path.join(mdir,"tokenizer_config.json")))["chat_template"]
+# Flip the "enable_thinking" default from ON to OFF. TWO edits are needed on this template and
+# only the first is mandatory:
+#   (1) the generation-prompt block that decides whether <think> is left open or pre-closed;
+#   (2) the reasoning-effort gate near the top, which injects a "think carefully..." system
+#       instruction whenever enable_thinking is UNDEFINED. Without (2) a thinking-off server
+#       still tells the model to reason at length while pre-closing the block it would reason in.
+# Each edit is attempted independently; (1) is required, (2) is best-effort.
+GEN_PROMPT = [
+ ("{%- if enable_thinking is defined and enable_thinking is false %}\n        {{- '<think>\\n\\n</think>\\n\\n' }}\n    {%- else %}\n        {{- '<think>\\n' }}\n    {%- endif %}",
+  "{%- if enable_thinking is defined and enable_thinking is true %}\n        {{- '<think>\\n' }}\n    {%- else %}\n        {{- '<think>\\n\\n</think>\\n\\n' }}\n    {%- endif %}"),
+ ("{%- set enable_thinking = enable_thinking if enable_thinking is defined else true %}",
+  "{%- set enable_thinking = enable_thinking if enable_thinking is defined else false %}"),
+]
+EFFORT_GATE = [
+ ("{%- if enable_thinking is undefined or enable_thinking is true %}",
+  "{%- if enable_thinking is defined and enable_thinking is true %}"),
+]
+patched = tmpl
+for old, new in GEN_PROMPT:
+    if old in patched:
+        patched = patched.replace(old, new); break
+else:
+    sys.exit("THINK=0: could not find the enable_thinking block to patch in this model's chat "
+             "template. Serve with THINK=1, or hand-patch the template and pass it via --chat-template.")
+for old, new in EFFORT_GATE:
+    if old in patched:
+        patched = patched.replace(old, new); break
+open(out, "w").write(patched)
+print(out)
+PY
+  TEMPLATE_ARGS+=(--chat-template "$NOTHINK")   # no reasoning parser when thinking is pre-closed
+fi
+
+echo "[serve] MODEL=$MODEL NAME=$SERVED_NAME $HOST:$PORT MEM=$MEM CTX=$CTXLEN GRAPHS=$GRAPHS THINK=$THINK RADIX=$RADIX INT8=$INT8"
+echo "[serve] tool-call-parser=$TOOL_PARSER  ptxas=$TRITON_PTXAS_PATH"
+
+exec "$PYBIN" -m sglang.launch_server \
+  --model-path "$MODEL" \
+  --served-model-name "$SERVED_NAME" \
+  --host "$HOST" --port "$PORT" \
+  --dtype float16 \
+  --mem-fraction-static "$MEM" \
+  --context-length "$CTXLEN" \
+  --chunked-prefill-size "$CHUNK" \
+  --sampling-backend pytorch \
+  --trust-remote-code \
+  --tool-call-parser "$TOOL_PARSER" \
+  "${REASON_ARGS[@]}" \
+  "${TEMPLATE_ARGS[@]}" \
+  "${POOL_ARGS[@]}" \
+  "${RADIX_ARGS[@]}" \
+  "${ATTN_ARGS[@]}" \
+  "${AUTH_ARGS[@]}" \
+  "${CLP_ARGS[@]}" \
+  "${TRUNC_ARGS[@]}" \
+  "${GRAPH_ARGS[@]}" \
+  "${JSON_ARGS[@]}" \
+  "${DET_ARGS[@]}" \
+  --log-level info \
+  "$@"          # forward any extra sglang flag verbatim

sglang/thinking_budget.py

@@ -0,0 +1,144 @@
+#!/usr/bin/env python3
+"""Bounded thinking for Escha `qwen3_5` models on SGLang.
+
+Problem this solves: with thinking on, the model may spend the whole
+`max_tokens` budget reasoning and never emit an answer. The response then comes
+back as `finish_reason: "length"` with `content: null` (all tokens went to
+`reasoning_content`), which an eval harness scores as *wrong* rather than as
+*slow*, and an agent framework sees as a no-op.
+
+A thinking budget fixes it: once the reasoning has run for N tokens, the server
+forces `</think>` so the model must produce its answer inside the remaining
+budget. This is SGLang's stock `ThinkingBudgetLogitProcessor`; the only reason a
+helper is needed is that SGLang's built-in Qwen3 subclass hardcodes Qwen3's
+token ids (151667/151668), which are wrong for this model's ~248k vocab — so we
+look the ids up from the model's own tokenizer.
+
+Serve with the processor enabled:
+
+    ENABLE_CLP=1 MODEL=/path/to/model bash serve.sh
+
+Emit the payload once, then reuse it for every request:
+
+    python thinking_budget.py --model /path/to/model > clp.json
+
+Send it (any OpenAI client — the two extra fields are SGLang extensions):
+
+    {"messages": [...], "max_tokens": 4096,
+     "chat_template_kwargs": {"enable_thinking": true, "reasoning_effort": "xhigh"},
+     "custom_logit_processor": <contents of clp.json>,
+     "custom_params": {"thinking_budget": 2048}}
+
+Budget guidance: leave enough headroom between `thinking_budget` and
+`max_tokens` for the answer itself — at least ~512 tokens, and far more for a
+task that answers with code. Every published benchmark number for this model
+was produced with `thinking_budget` 28,672 inside `max_tokens` 32,768 (so
+4,096 tokens always remain for the answer), at `reasoning_effort: "xhigh"`.
+
+`reasoning_effort` is a second, independent lever this model exposes through
+`chat_template_kwargs` — `"xhigh"` (default), `"medium"` or `"low"`. It asks
+for shorter reasoning; the budget *enforces* it. They compose.
+
+Self-check (`--demo`) needs a running server and reports, for one prompt, the
+finish_reason / reasoning length / answer length with and without a budget.
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import sys
+
+_THINK_OPEN = "<think>"
+_THINK_CLOSE = "</think>"
+
+
+def token_ids(model_dir: str) -> tuple[int, int, int]:
+    """(start, end, newline) token ids from the model's own tokenizer."""
+    from transformers import AutoTokenizer
+
+    tok = AutoTokenizer.from_pretrained(model_dir, trust_remote_code=False)
+    start, end = tok.convert_tokens_to_ids([_THINK_OPEN, _THINK_CLOSE])
+    unk = getattr(tok, "unk_token_id", None)
+    if start is None or end is None or start == unk or end == unk:
+        raise SystemExit(
+            f"{model_dir}: tokenizer has no {_THINK_OPEN}/{_THINK_CLOSE} tokens — "
+            "this model has no thinking mode, so a thinking budget does not apply."
+        )
+    newline = tok.encode("\n", add_special_tokens=False)
+    if len(newline) != 1:
+        raise SystemExit(f"unexpected newline tokenization: {newline!r}")
+    return int(start), int(end), int(newline[0])
+
+
+def payload(model_dir: str) -> str:
+    """The JSON string to pass as `custom_logit_processor`."""
+    from sglang.srt.sampling.custom_logit_processor import (
+        ThinkingBudgetLogitProcessor,
+    )
+
+    start, end, newline = token_ids(model_dir)
+
+    # dill serializes this subclass by value; the base class is resolved by name
+    # on the server, so the server must run the same SGLang version.
+    class EschaThinkingBudgetLogitProcessor(ThinkingBudgetLogitProcessor):
+        THINKING_START_TOKEN_ID = start
+        THINKING_END_TOKEN_ID = end
+        NEW_LINE_TOKEN_ID = newline
+
+    return EschaThinkingBudgetLogitProcessor.to_str()
+
+
+def demo(base_url: str, model: str, clp: str, budget: int, max_tokens: int) -> None:
+    import urllib.request
+
+    prompt = (
+        "What is the Hubble constant from Planck, and from local distance-ladder "
+        "measurements? Give both numbers with units."
+    )
+    opener = urllib.request.build_opener(urllib.request.ProxyHandler({}))
+    for label, body in (
+        ("no budget ", {}),
+        (f"budget {budget}", {"custom_logit_processor": clp,
+                              "custom_params": {"thinking_budget": budget}}),
+    ):
+        req = {"model": model, "messages": [{"role": "user", "content": prompt}],
+               "max_tokens": max_tokens, "temperature": 0,
+               "chat_template_kwargs": {"enable_thinking": True}, **body}
+        with opener.open(urllib.request.Request(
+                f"{base_url}/chat/completions", json.dumps(req).encode(),
+                {"Content-Type": "application/json"}), timeout=1800) as r:
+            out = json.load(r)
+        choice = out["choices"][0]
+        msg = choice["message"]
+        reasoning = msg.get("reasoning_content") or ""
+        content = msg.get("content") or ""
+        print(f"{label}: finish_reason={choice['finish_reason']:>6} "
+              f"completion_tokens={out['usage']['completion_tokens']:>6} "
+              f"reasoning_chars={len(reasoning):>6} answer_chars={len(content):>5}"
+              f"{'   <-- NO ANSWER' if not content else ''}")
+
+
+def main() -> None:
+    ap = argparse.ArgumentParser(description=__doc__,
+                                 formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--model", required=True, help="model directory (for the tokenizer)")
+    ap.add_argument("--demo", action="store_true", help="probe a running server")
+    ap.add_argument("--base-url", default="http://127.0.0.1:30000/v1")
+    ap.add_argument("--served-name", default="escha-qwen38-27b-w2")
+    ap.add_argument("--budget", type=int, default=2048)
+    ap.add_argument("--max-tokens", type=int, default=4096)
+    args = ap.parse_args()
+
+    clp = payload(args.model)
+    if args.demo:
+        start, end, newline = token_ids(args.model)
+        print(f"# {_THINK_OPEN}={start} {_THINK_CLOSE}={end} newline={newline}",
+              file=sys.stderr)
+        demo(args.base_url, args.served_name, clp, args.budget, args.max_tokens)
+    else:
+        print(clp)
+
+
+if __name__ == "__main__":
+    main()