Adapt to TVM API and `call_tir` convention changes:
1. Update `DataflowVar.name_hint` to `DataflowVar.name`
2. Change sampler full kernel to accept int64 value before its output buffer and cast to int32
3. Reorder parameters for `tensor_ir_op` PrimFuncs so tensor/scalar arguments precede output buffers (Triton expert-ID generation, MoE indptr generation, vision crop and padding, Phi-3V dynamic repeat)
4. Move crop's `T.reads` and `T.writes` declarations to top of `T.sblock` as required by current TIR parser
A newer TVM bumps tvm-ffi so that `Optional<T>` follows std::optional
semantics: `.defined()` is dropped in favor of `.has_value()`, and
`Optional<Tensor>` no longer implicitly converts to `ObjectRef`. Update
the C++ runtime to call `.has_value()` on the affected `Optional`
receivers (leaving `.defined()` on plain `ObjectRef`/`Function`/`Module`
handles intact) and return `recv.value_or(Tensor(nullptr))` from the
multi-GPU send/recv passthrough.
On the Python side, adapt the compiler passes and ops to the Relax/tirx
API changes. The Relax `Id` indirection is gone, so `PyExprMutator` var
remaps take the `Var` directly instead of `var.vid`. Symbolic size vars
drop `is_size_var`/`SizeVar` for plain `T.int32()`/`tirx.Var`;
`tirx.PrimExpr`/`multiply`/`subtract`/`generic.cast` become
`Expr`/`Mul`/`Sub`/`Cast`; the cross-thread all-reduce idiom uses
`T.int32(0)` with `dtype="void"`; `relax.expr.Call` becomes
`relax.Call`; and handle parameters are detected via
`isinstance(v.ty, PointerType)` now that a var's `.ty` carries a
`PrimType`/`PointerType` rather than a dtype string.
Verified end to end by compiling and chatting with both
Phi-4-mini-instruct and Qwen3-30B-A3B under tensor_parallel_shards=2.
Mainline TVM unified the PrimExpr type mechanism on PrimType (in place of
DataType) and moved the runtime datatype/object accessors onto tvm-ffi.
Update the C++ runtime to match: replace `tvm::Downcast<T>` with the
tvm-ffi accessors `.cast<T>()` / `.as_or_throw<T>()`, and use `DLDataType`
in place of `tvm::runtime::DataType`. In cpp/serve/model.cc the KV-cache
creation function now returns a heterogeneous `Array<Any>`, so its
elements are read as `Array<Any>[i].cast<Tensor>()` to avoid an Array
type error at runtime.
On the Python side, adapt the compiler passes, model definitions, and ops
to the refactored Relax/TIRScript API: parse TIRScript via
`tvm.script.tirx`, follow the StructInfo->Type rename (`.struct_info`->
`.ty`, `sinfo_args`->`ty_args`), use `relax.prim_value` and
`@T.prim_func(s_tir=True)`, and add explicit casts where the stricter
tirx well-formedness rules no longer let `T.let` widen implicitly.
Bump the 3rdparty/tvm submodule to the matching revision.
Also fix the (previously red) Windows CI job: the scikit-build-core wheel
build needs the MSVC toolchain on PATH, so activate the MSVC developer
environment before building; pin the build conda environment to Python
3.11 (core deps are dropping 3.10 wheels); and remove the GNU coreutils
link.exe that conda ships in Library\usr\bin, which otherwise shadows the
MSVC linker and breaks the Rust tokenizers build.
Mainline TVM removed several public headers and re-namespaced their
contents. Update C++ includes to the new tvm/ffi and tvm/ir paths,
rename IntTuple to ffi::Shape, replace tvm::runtime re-exports with
explicit `using tvm::ffi::Object/...; using tvm::Downcast;`, switch
linking from `tvm` to `tvm_runtime`, and vendor two small shims under
cpp/support/ for the removed `parallel_for_with_threading_backend`
helper and the now-private TVM_MODULE_VTABLE_* macros.
Python: ShapeTuple is gone — use tvm_ffi.Shape in the five affected
modules, and drop the now-removed `target` arg from the
tvm_callback_cuda_compile hook (TVM now fetches it via Target.current).
Mirror the TVM setup (apache/tvm#18809): a minimal GHA workflow that
runs `uv sync --group lint` plus `pre-commit/action`. Pre-commit is
rewritten around ruff (check + format), replacing the previous
isort/black/mypy/pylint/cmake-format stack, and gains yamllint,
taplo, check-yaml, and check-toml. The corresponding Jenkins Lint
stage, the ci/task/{isort,black,mypy,pylint,clang-format}.sh scripts,
and .pylintrc are removed. pyproject.toml grows a [tool.ruff] section
(select E/F/I/UP/RUF, line-length 100, py39 target) and a
[dependency-groups] lint entry; the stale isort/black/mypy/pylint
tool tables are dropped.
To get the codebase green under the new ruler, existing UP006/UP035/
UP045/E501/RUF001/RUF002/E402/F821 violations are grandfathered via
inline `# noqa`, so new code in those forms is still flagged. Real
issues surfaced by the new rules are fixed directly: E741 renames
(ambiguous `l`/`O`), F841 unused locals, E722 bare `except`, RUF012
ClassVar annotations for mutable class defaults, RUF013 explicit
`Optional[...]`, and RUF005 iterable unpacking.
With pylint and mypy gone from CI, this also strips all `# pylint:
disable/enable` and `# type: ignore` directives from Python sources
(PEP 484 `# type: <expr>` comments and `# fmt: off/on` are
preserved), and runs the resulting files through ruff-format.
Qwen3's official HF chat template strips `<think>...</think>` blocks
from historical assistant messages (all turns before the last user
message) before rendering the prompt. mlc-llm's `qwen2` template —
which the published Qwen3-*-MLC repos use — does not, so prior
thinking traces get echoed back into context verbatim. On small
Qwen3 variants (e.g. 0.6B) this pushes the model to emit `<|im_end|>`
prematurely inside its next-turn `<think>` block, truncating the
response before `</think>` is ever produced.
This change adds a `strip_reasoning_in_history` flag on the
`Conversation` protocol (both the Python `as_prompt` in
`conversation_protocol.py` and the C++ `CreatePrompt` in
`json_ffi/conv_template.cc`), and registers a new `qwen3` template
that sets it. For Qwen3-*-MLC repos to pick it up, their
`mlc-chat-config.json` on HuggingFace needs the `conv_template`
entry swapped to the new `qwen3` template; no weight changes.
Fixes#3482.
* [WebGPU] Add --enable-subgroups flag for subgroup shuffle support
* Fix _apply_webgpu_subgroups function
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
* Added unit tests for _apply_webgpu_subgroups
---------
Co-authored-by: Akaash Parthasarathy <43900735+akaashrp@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Add dedicated conversation templates for Qwen3.5 models instead of
using the generic ChatML format. Includes two variants:
- qwen3_5: thinking enabled (assistant opens a <think> block)
- qwen3_5_nothink: thinking disabled (empty <think> block prefix)
Sculptor: gabeguralnick
Co-authored-by: Gabriel Guralnick <gabriel@imbue.com>
The post_layernorm was defined in CLIPVisionTransformer.__init__ but
never applied in the forward method. This fix applies it to the final
encoder hidden state, replacing it in the returned tuple. This matches
the HuggingFace CLIPVisionTransformer behavior where last_hidden_state
is post-normalized.
The [-2] indexing in CLIPVisionModel (used by LLaVA and Phi3V to get
penultimate layer features) continues to work correctly since only the
last element of the tuple is replaced.
Fixes#3205
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## Summary
- avoid the Pydantic protected namespace warning for `model_task` in `MLCChatConfig`
- skip `flashinfer-python` for non-Linux editable installs (tested in macos)
## Details
- rename the internal Pydantic field to `field_model_task` and keep `alias="model_task"` so the serialized config format stays unchanged
- add a platform marker to `python/requirements.txt` so editable installs on macOS do not try to resolve `flashinfer-python`
* Add Qwen3.5 model definition and kHybrid KVStateKind to serving runtime
Adds full Qwen3.5 GatedDeltaNet hybrid architecture support to MLC-LLM.
This is a hybrid model: 75% GatedDeltaNet recurrent linear attention layers,
25% standard GQA softmax attention layers (every 4th layer).
C++ changes:
1. KVStateKind enum in cpp/metadata/model.h: added kHybrid = 3, plus
KVStateKindToString/KVStateKindFromString cases for "hybrid".
2. model.cc metadata parsing (cpp/metadata/model.cc): kHybrid is treated
like kKVCache for KVCacheMetadata parsing (reads the "kv_cache" JSON
block), since hybrid models still use PagedKVCache for attention layers.
3. FunctionTable (cpp/serve/function_table.h, function_table.cc): added
create_rnn_state_func_ member. In _InitFunctions, when kv_state_kind
is kHybrid, resolves create_kv_cache_func_ to
"create_tir_paged_kv_cache" and create_rnn_state_func_ to
"create_rnn_state" separately (hybrid needs both).
4. ModelImpl (cpp/serve/model.cc): added rnn_state_ and local_rnn_state_
members. CreateKVCache creates both PagedKVCache and RNNState when
kHybrid. All KV state lifecycle methods (AddNewSequence,
ForkSequence, RemoveSequence, PopNFromKVCache, Reset) mirror
operations to rnn_state_ when kHybrid. BatchPrefill, BatchDecode,
BatchVerify, and their ToLastHidden variants call
kv_cache_begin_forward_func_ and kv_cache_end_forward_func_ on
rnn_state_ when kHybrid. BatchPrefill/BatchDecode/BatchVerify pass
rnn_state_ as an extra argument to the model functions when kHybrid.
GetNumAvailablePages and GetCurrentTotalSequenceLength use
PagedKVCache (not RNNState) for capacity, since PagedKVCache is the
constraining factor.
5. batch_prefill_base.cc: CanPrefill treats kHybrid like kKVCache (uses
PagedKVCache page count for capacity decisions, not unlimited like
kRNNState).
Python changes:
6. compile.py (_infer_kv_state_kind): returns "hybrid" when model_type
contains "qwen3_5".
7. model/model.py: registered "qwen3_5" and "qwen3_5_text" entries in the
MODELS dict, importing from the new qwen35 package. Both map to
Qwen35LMHeadModel and Qwen35Config.
8. model/qwen35/qwen35_model.py: full model definition.
Qwen35Config handles VLM text_config unwrapping and nested
rope_parameters. Qwen35Attention implements GQA with output gate.
Qwen35GatedDeltaNet implements linear attention with causal Conv1D,
L2 normalization, gate computation (exp(-exp(A_log) * softplus(alpha
+ dt_bias))), delta rule recurrence via hand-written TIR kernel
(thread-per-column, fp32 accumulation, GVA support), and per-head
RMSNorm output gating. Two forward paths: forward() with explicit
stacked state tensors, forward_rnn() with RNNState get/set.
create_paged_kv_cache uses num_attention_layers with partial RoPE.
create_rnn_state initializes recurrent (fp32) and conv state per
linear layer.
9. model/qwen35/qwen35_loader.py: HF parameter mapping with VLM prefix
translation (model.language_model.* -> model.*). Fuses QKV for
attention, gate_up for MLP. Maps A_log and dt_bias without .weight
suffix. Adds +1.0 offset to RMSNorm weights where needed.
* Pass rnn_state to ToLastHidden functions for hybrid models, remove unused linear_idx
ToLastHidden variants (prefill, decode, verify, and their single_batch
counterparts) were missing the rnn_state_ argument when kv_state_kind
is kHybrid. These are used for speculative decoding paths. Also removed
dead linear_idx variable from qwen35_loader.py.
* Fixed memory reuse fault that was causing crashes.
TVMs memory planner was reusing the initialization tensor's backing
memory, as it detected they were not being used after the RNN state
create call. In our case, we need to keep them around for every
request. Changing them to constants, and adding the compat fixes
into RNNState.py seem to have solved that problem. Still have a problem
with prefix caching failing.
Also - this likely affects the RWKV series of models as well.
* updated prefix caching logic in engine to work with rnn state semantics
* cleanup
* updated the rnn state builder to always take rx constants
* formatting / ci fixes
* revert chat updates
* cleanup
---------
Co-authored-by: Sidharth N. Babu <sidhartb@andrew.cmu.edu>
Co-authored-by: Sidharth N. Babu <sidharth.n.babu@gmail.com>
This PR is the first staged step toward a first-class embedding
serving/runtime path in MLC LLM.
It introduces an explicit embedding metadata abstraction so
embedding models are described by metadata/config traits instead
of being identified implicitly from exported function names alone.
This PR is scoped to metadata propagation only. It does not change
the `/v1/embeddings` API contract, does not change the normal
chat serving path, and does not yet introduce a dedicated embedding
runtime boundary.
Part of #3451
Fix two issues with `mlc_llm chat` exit behavior:
- Catch `KeyboardInterrupt` and `EOFError` in the input loop so Ctrl+C/Ctrl+D exit cleanly without a traceback.
- Explicitly call `engine.terminate()` after the chat loop ends to stop background threads. Previously the engine relied on `__del__` for cleanup, which never fired due to circular references, causing the process to hang on exit.
NOTE: This PR is only the decoder-side MRoPE/text-model part,
not full end-to-end Qwen2.5-VL support yet.
## What this PR does
- Adds initial Qwen2.5-VL Python skeleton:
- `python/mlc_llm/model/qwen2_5_vl/__init__.py`
- `python/mlc_llm/model/qwen2_5_vl/qwen2_5_vl_model.py`
- Adds MRoPE op implementation and public exports:
- `python/mlc_llm/op/mrope.py`
- `python/mlc_llm/op/__init__.py`
- Makes MRoPE position-id handling layout-compatible by accepting both:
- `(batch, seq, 3)`
- `(3, batch, seq)`
## Why
- Fixes the position-id layout mismatch between `get_mrope_position_ids` output and MRoPE application.
- Unblocks direct use of generated multimodal position IDs in the Qwen2.5-VL path.
A lot of our models have really similar huggingface loaders, so we
should be able to replace with a standardized central loader, and
cut down some redundancy. There are a few models with some
specific quirks that are not yet covered by this.
Explored the codebase, and found that quantization infrastructure
is duplicated on a per-model basis, when it could be centralized.
Tested with qwen3-8b - qualitatively compared the chat outputs of
quantization q4f16_1 across the main branch and this one, and both worked.
Simple PR - noticed we had a "CasualLM" instead of "CausalLM" typo
across the codebase, and just wanted to clean it up. Also removed some
commented out model presets, and removed gen_config support for for
ones that don't have corresponding conv_templates.
Not sure if this last one is necessary / correct, can go back and restore that.
Usage:
```
python python/bench/eval/mmlu.py --model dist/Meta-Llama-3-8B-Instruct-q4f16_1-MLC --dataset /path/to/dataset --device cuda --log-dir debug/mmlu
```
Note that chat mode is problematic, which needs to be fixed.
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
This PR adds API key authentication to the server. Now users can secure
their endpoints with a simple API key.
Key features:
- API key is optional (disabled by default)
- API key is passed via `--api-key` command line argument (not
hardcoded or using environment variables)
- Authentication follows OpenAI API standard (Bearer token in Authorization header)
- No breaking changes - existing functionality remains unchanged
Replace `tir.stmt_functor.renew_defs` with `s_tir.renew_defs` to fix the
pylint E1101 error reporting that `tvm.tir.stmt_functor` has no `renew_defs`
member, aligning with the updated TVM mainline API.
### Problem:
When total tokens in a batch exceed `prefill_chunk_size (2048 in my case)`,
`_embed_decoder()` falls back to sequential processing for **ALL** texts,
causing a throughput cliff (53 → 4 texts/s at batch_size=64)
### Fix:
Greedy sub-batching — pack texts into sub-batches that fit within prefill_chunk,
preserving input order. Oversize single texts fall back to sequential chunked prefill.
## Summary
Add end to end embedding inference support with /v1/embeddings endpoint, also provide related test files.
### Key Changes:
- Add `Qwen3EmbeddingModel` subclass with quantize functions
to fix embedding model compilation
- Add `AsyncEmbeddingEngine` with auto model-type detection,
batch prefill optimization, and chunked prefill fallback
- Add `/v1/embeddings` endpoint similar to openai, has dimension
truncation and base64 encoding support
- Add `bert-bge` model type
- Add comprehensive tests (logit match, engine-level, server test)
Fix missing 1/sqrt(d) scaling in attention fallback path
The _fallback() path in op_ext.attention called
_attention_sequence_prefill with sm_scale=1.0 (default),
but the correct scaling factor is 1/sqrt(head_dim).
This caused attention scores to be head_dim times too large,
making softmax overly peaked and degrading embedding quality
for any model using this fallback (BERT, CLIP Vision, or any
model running on Metal/CPU/non-fp16 configurations).
Decoder models (Llama, Qwen, etc.) were not affected because
they pass sm_scale=self.head_dim**-0.5 through the KV cache
path, bypassing op_ext.attention entirely.
Measured impact on BGE-base-en-v1.5 (BERT encoder, Metal):
Before: TVM vs HuggingFace cosine = 0.862 (avg)
After: TVM vs HuggingFace cosine = 0.999989 (avg)
This was a stale duplicate of
tvm.relax.frontend.nn.llm.position_embedding.
Verified removal is safe by:
Searching for all imports from mlc_llm.op.position_embedding (none
found)
Testing that mlc_llm.op imports correctly after cleanup
Introduces Mac Catalyst (macabi) as a supported device target across build
scripts and Python packaging. Updates ios/prepare_libs.sh to handle --catalyst
builds, adds macabi handling in Python interface and packaging logic, and
extends auto_target.py to recognize and configure macabi targets.
This enables building and packaging static libraries for Mac Catalyst (arm64)
alongside existing iOS and Android targets.
Introduce support for specifying Mac Catalyst architecture (arm64 or x86_64)
in build scripts and Python interface. Update build directories, CMake
arguments, and target triples to handle both architectures. Improves flexibility
for building Mac Catalyst bindings.
Changed tir.Schedule to tvm.s_tir.Schedule in:
- attach_softmax_with_temperature.py
- fuse_dequantize_take.py
Fixes AttributeError during model compilation where tir.Schedule
does not exist. The correct API is tvm.s_tir.Schedule.
Akaash Parthasarathy