vLLM TT Plugin

Tenstorrent backend plugin for vLLM.

vllm-tt-plugin integrates Tenstorrent hardware into vLLM using the standard plugin mechanism. Install it alongside vLLM and, when ttnn is importable (see TT-Metal), TT hardware is automatically available as a vLLM platform.

The plugin is self-contained: model registration, platform detection, request validation, scheduling, worker execution, model loading, async decode, single- process standard and multi-lane execution, data-parallel execution all live here. Nothing TT-specific needs to touch vLLM core.

For model authors, the model-plugin interface documents registration, construction, KV storage, execution, and lifecycle, and links to the capability, speculative-decode, and input-reload contracts.

Package Layout

.
+-- src/vllm_tt_plugin/
|   +-- entrypoints.py       # vLLM plugin entry points
|   +-- platform.py          # TTPlatform and config validation
|   +-- model_registry.py    # TT model architecture registration
|   +-- worker.py            # TT worker implementation
|   +-- model_runner.py      # TT model execution bridge
|   +-- scheduler.py         # TT scheduling policy
|   +-- lane_scheduler.py    # Single-process multi-lane (lane-DP) coordinator
|   +-- launcher.py          # retained tt-run / MPI launcher (not hooked by vLLM 0.30.0)
|   +-- loader.py            # TT model loader
|   +-- input_batch.py       # TT input-batch representation
|   +-- async_decode.py      # Decode overlap helpers
|   +-- config.py            # TT plugin config access
|   +-- spec_admission.py    # Speculative configuration and model-plan validation
|   +-- spec_decode.py       # Candidate/verification types and greedy acceptance
|   +-- spec_accept.py       # CPU sampled-acceptance helper (not yet used by the runner)
|   +-- utils/               # Common helpers such as device discovery tools for DP
+-- docs/                    # TT runtime notes
+-- examples/                # Offline and OpenAI-server examples
+-- tests/tt/                # Server-facing TT plugin tests

Requirements

If testing a specific model, check the TT-Metal LLMs table for the appropriate tt-metal and vLLM commits.

vLLM requires Python >=3.10,<3.14. Python 3.10.12 is the default python3 on Ubuntu 22.04.

The installation script builds vLLM 0.30.0 from source with VLLM_TARGET_DEVICE=empty.

To install against older supported vLLM versions X.Y.Z instead (e.g., 0.25.1), check out the compat/vllm-X.Y.Z tag (e.g., compat/vllm-0.25.1), the last plugin state that targets it, and follow the same steps. Nothing is maintained on top of that tag for now. Pair it with the tt-metal commit the LLMs table lists for it.

Environment Setup

Install tt-metal first by following INSTALLING.md. If installing tt-metal from source, build it, create the virtual environment, and set the environment variables needed for tt-metal tests.

Activate the environment where tt-metal is installed, then install vLLM and the TT plugin. Run it from the repository root, which the relative paths inside assume:

source docs/install-vllm-tt.sh

The script installs vLLM with VLLM_TARGET_DEVICE=empty because tt platform is provided by this plugin at runtime. It then installs the plugin with a few dependencies. Most dependencies come from the active tt-metal env.

The script also installs vLLM's dependency list itself, fetched from requirements/common.txt at the pinned vLLM tag, and installs vLLM with --no-deps. Resolving them the usual way pulls vLLM's CUDA dependency set (torch pinned, flashinfer, tilelang, nvidia-*), which would fight the tt-metal env over torch and add several GB. This means the script needs network access to raw.githubusercontent.com beyond the package index. When installing inside a container, also set UV_NO_CACHE=1 to keep the uv cache out of the image layer.

common.txt omits torchvision, which vLLM imports unconditionally from several model and processor modules, so the script installs it separately: the CPU build, with --no-deps, at the version the tt-metal env pins alongside its torch. Environments built from tt-metal's requirements-dev.txt already carry it; the ttnn wheel does not declare it.

To install or refresh only the plugin package:

uv pip install -e .

To run the offline Qwen-VL example, also install its extra:

uv pip install -e ".[examples]"

After the first setup, activate the same environment before running vLLM:

source "$PYTHON_ENV_DIR/bin/activate"

VLLM_TARGET_DEVICE is a build-time variable only and does not need to be set at runtime. The TT platform is detected automatically when ttnn is importable.

Install pre-commit hooks once in your active development environment. After installation, these checks run automatically on every git commit:

uv pip install pre-commit
pre-commit install

Verify Plugin Discovery

The editable install registers two vLLM entry points:

Entry point group	Name	Target
vllm.general_plugins	tt_model_registry	vllm_tt_plugin.entrypoints:register
vllm.platform_plugins	tt	vllm_tt_plugin.entrypoints:platform_plugin

platform_plugin() returns vllm_tt_plugin.platform.TTPlatform only when ttnn is importable. This keeps ordinary vLLM environments from accidentally selecting the TT platform.

Quick checks:

python -c "import vllm_tt_plugin; print(vllm_tt_plugin.__file__)"
python -c "import ttnn; print('ttnn available')"

If VLLM_PLUGINS is set, it must allow both TT entry point names:

export VLLM_PLUGINS=tt,tt_model_registry

Hugging Face Access

To run Meta Llama 3.1 or 3.2 models, request access on Hugging Face:

• Llama 3.1
• Llama 3.2
• Llama 3.2 Vision

After access is approved, create an access token in Hugging Face settings and log in from Python:

from huggingface_hub import login

login()

Preparing TT-Metal Models

For the target model, follow any setup instructions in the corresponding tt-metal demo. For Llama 3.1, Llama 3.2, and Qwen 2.5 models, follow the tt-transformers demo instructions for weights and environment variables.

Running The Offline Inference Example

Run offline generation with the default Llama 3.1 70B model:

MESH_DEVICE=T3K python examples/offline_inference_tt.py

Measure offline performance for one batch of prompts:

MESH_DEVICE=T3K \
python examples/offline_inference_tt.py --measure_perf

To run a different text model, set MESH_DEVICE to N150, N300, T3K, TG, BH-Galaxy, or a mesh shape such as "(4,8)", then pass --model:

• Llama 3.1 8B: --model "meta-llama/Llama-3.1-8B"
• Llama 3.2 1B: --model "meta-llama/Llama-3.2-1B"
• Llama 3.2 3B: --model "meta-llama/Llama-3.2-3B"
• Qwen 2.5 7B: --model "Qwen/Qwen2.5-7B"
• Qwen 2.5 72B: --model "Qwen/Qwen2.5-72B"
• DeepSeek R1 Distill Llama 70B: --model "deepseek-ai/DeepSeek-R1-Distill-Llama-70B"
• GPT-OSS 20B: --model "openai/gpt-oss-20b"
• GPT-OSS 120B: --model "openai/gpt-oss-120b"

For Llama 3.1 8B on N150, set --max_model_len 32768; see the tt-metal model demo for context-length details.

To run Llama 70B on Wormhole Galaxy:

MESH_DEVICE=TG \
LLAMA_DIR=<path-to-weights> \
TT_LLAMA_TEXT_VER=llama3_70b_galaxy \
python examples/offline_inference_tt.py \
  --model "meta-llama/Llama-3.1-70B-Instruct" \
  --additional-config '{"tt": {"dispatch_core_axis": "col", "sample_on_device_mode": "all", "worker_l1_size": 1344544, "trace_region_size": 216580672}}'

To run GPT-OSS 20B on Wormhole Galaxy:

MESH_DEVICE="(4,8)" \
python examples/offline_inference_tt.py \
  --model "openai/gpt-oss-20b" \
  --max_seqs_in_batch 1

To run Qwen3-32B on Blackhole Galaxy:

MESH_DEVICE=TG \
TT_QWEN3_TEXT_VER=qwen3_32b_galaxy \
python examples/offline_inference_tt.py \
  --model "Qwen/Qwen3-32B" \
  --additional-config '{"tt": {"dispatch_core_axis": "col", "sample_on_device_mode": "all", "worker_l1_size": 1345000, "trace_region_size": 184915840}}'

Wormhole Galaxy defaults to FABRIC_1D_RING and Blackhole Galaxy defaults to FABRIC_2D_TORUS_XY, so those recipes do not need an explicit fabric_config.

Run Llama 3.2 Vision on N300:

MESH_DEVICE=N300 \
python examples/offline_inference_tt.py \
  --model "meta-llama/Llama-3.2-11B-Vision-Instruct" \
  --multi_modal \
  --max_seqs_in_batch 16 \
  --num_repeat_prompts 8

Useful vision-model variants:

• Llama 3.2 11B Vision on QuietBox: set MESH_DEVICE=T3K and --max_seqs_in_batch 32.
• Llama 3.2 90B Vision: set MESH_DEVICE=T3K, --model "meta-llama/Llama-3.2-90B-Vision-Instruct", and --max_seqs_in_batch 4.
• Qwen 2.5-VL 32B: set MESH_DEVICE=T3K, --model "Qwen/Qwen2.5-VL-32B", and --max_seqs_in_batch 32.
• Qwen 2.5-VL 72B: set MESH_DEVICE=T3K, --model "Qwen/Qwen2.5-VL-72B", --max_seqs_in_batch 32, --max_model_len 2048, and --additional-config '{"tt": {"trace_region_size": 28467200}}'.
• Gemma 3 27B: set MESH_DEVICE=T3K, --model "google/gemma-3-27b-it", --max_seqs_in_batch 32, --additional-config '{"tt": {"l1_small_size": 768, "fabric_config": "FABRIC_1D"}}', --multi_modal, --multi_image, and --mm_processor_kwargs '{"use_fast": true, "do_convert_rgb": true}'.

For debugging V1, set VLLM_ENABLE_V1_MULTIPROCESSING=0 to disable multiprocessing. This is useful for stepping through code or making scheduling deterministic, but it is not compatible with DP models.

Running The Server Example

Start the OpenAI-compatible server:

MESH_DEVICE=T3K \
python examples/server_example_tt.py

DiffusionGemma uses a 256-token block-serving contract with stricter launch and request constraints. See DiffusionGemma block serving for the exact command, request validation, and block metrics.

Send a completion request:

curl http://localhost:8000/v1/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "meta-llama/Llama-3.1-70B-Instruct",
    "prompt": "San Francisco is a",
    "max_tokens": 32,
    "temperature": 1,
    "top_p": 0.9,
    "top_k": 10
  }'

Before a forward is submitted, TT selects one sampling path for the whole submitted batch. If any request or batch feature cannot use on-device sampling, the batch uses vLLM's host-side path. Structured outputs have additional rules; see Structured Output Sampling.

Single-token-output models can explicitly use "seeded_sampling_policy": "host" to sample requests with an API seed on host during both prefill and decode. The default, "auto", preserves the existing routing. With the host policy, one submitted seeded request moves the whole forward to host sampling, including its unseeded companions. Seeded greedy requests also use host sampling. This can reduce throughput and increase latency; batches without submitted seeded requests retain their existing routing. Models that own multi-token sampling and speculative decoding configurations reject this option. It does not promise batch-independent model logits or waive tests.

For vision models, start the server with the correct --model, then send a chat completion request with image content. Qwen 2.5-VL models can use either a base64 data:image/... URL or a real URL such as https://qianwen-res.oss-cn-beijing.aliyuncs.com/Qwen-VL/assets/demo.jpeg.

Configuration

TT options are passed through vLLM's generic additional config namespace:

--additional-config '{"tt": {"sample_on_device_mode": "all"}}'

Plugin code reads this through vllm_tt_plugin.config.get_tt_config(), which returns vllm_config.additional_config["tt"].

Common options:

Key	Purpose
sample_on_device_mode	Select on-device sampling mode, currently all or decode_only when supported by the model.
seeded_sampling_policy	auto (default) or host. Select host sampling for submitted seeded requests during prefill and decode; affects the whole submitted forward. Single-token-output models only; incompatible with speculative decoding.
trace_mode	Control TT tracing: all, decode_only, or none. Default: all.
enable_model_warmup	Warm up the model before the server reports healthy. Default: true.
trace_region_size	Trace region size for TT runtime tracing.
worker_l1_size	Worker L1 size override.
l1_small_size	Small L1 size override.
fabric_config	Fabric config such as DISABLED, FABRIC_1D, FABRIC_2D, FABRIC_1D_RING, FABRIC_2D_TORUS_XY, or CUSTOM. Any ttnn.FabricConfig name is accepted. Overrides model defaults; otherwise defaults to Wormhole Galaxy FABRIC_1D_RING, Blackhole Galaxy FABRIC_2D_TORUS_XY, other multi-device FABRIC_1D.
fabric_reliability_mode	Fabric reliability mode, such as STRICT_INIT or RELAXED_INIT.
dispatch_core_axis	Dispatch core axis, row or col.
always_compat_sampling	Use vLLM's LogitProcessor and sampler path even when not required by the batch. Default: false.
decode_interleave_enabled	Insert decode-only steps into a run of consecutive prefill steps, so a running request's inter-token latency does not scale with another request's prompt length. Default: true.
decode_interleave_prefill_steps	Consecutive prefill steps allowed before one insertion. Raise it to favor time to first token, lower it to favor inter-token latency. Default: 2.
decode_interleave_decode_steps	How many decode-only steps one insertion runs. Reaching this count stops the policy choosing decode; it never forces a prefill step. Does not loosen the latency bound; buys more decode progress per insertion, and more cheaply than a shorter prefill run does. Default: 1.
optimizations	Select model/runtime optimization profile, such as accuracy or performance.
register_test_models	Register non-production TT test models for infrastructure tests. Default: false.

Default top-k for device sampling

The TT device sampler draws only from the 32 most probable tokens. A model class declares this limit as max_device_top_k (see Sampling limits). When a request has a temperature above 0 and a top_k that is unset or above the limit, the plugin samples the whole batch on host for that step. Host sampling reads back the full-vocabulary logits on every decode step, so decode is much slower than with device sampling.

Most requests do not set top_k: the OpenAI Python client has no top_k argument, and vllm bench serve sends no sampling parameters. vLLM then uses the default from the model's generation_config.json. Llama 3.1, Llama 3.3 and GPT-OSS define no top_k default, which means all tokens. Gemma 3 and Gemma 4 define top_k: 64. So with these models, every request that does not set top_k samples on host.

To keep these requests on the device, start the server with a default top_k equal to the device limit:

--override-generation-config '{"top_k": 32}'

vLLM merges this value into the model's default sampling parameters. It applies only to requests that do not set top_k. A request that sets its own top_k keeps it, and a value above 32 still samples on host.

Use this setting when all of these are true:

• sample_on_device_mode is set, and the model supports device sampling on the mesh.
• The model's default temperature is above 0. Greedy requests do not use top_k.
• The model's default top_k is unset or above 32.

Do not use it when the model's default top_k is 32 or less, for example Qwen3 (top_k: 20). The override replaces the model default, so it would widen it. Do not use it when the server samples on host, because it then only changes the distribution and does not make decode faster.

The setting is an approximation: "all tokens", or 64 for Gemma, becomes the 32 most probable tokens. A client that needs the model's full distribution sets top_k in the request (for example extra_body={"top_k": 0} with the OpenAI client) and accepts host sampling for that request.

Some model classes do not declare max_device_top_k, for example the tt-metal models/demos/gemma4 and models/demos/llama31_8b_qb2 classes (tenstorrent/tt-metal#60365). For them, the plugin does not route by top_k, and the device sampler changes any top_k above 32 to 32 without a message. The setting is still recommended for these models: the request then asks for what the device does.

Structured Output Sampling

Structured-output decode can use TT device sampling on models that declare supports_device_grammar. It is off by default. To enable it, set a sample_on_device_mode and turn off async scheduling. The example also sets the default top-k, because requests without top_k to the default Llama model would otherwise sample on host:

MESH_DEVICE=T3K \
python examples/server_example_tt.py \
  --no-async-scheduling \
  --override-generation-config '{"top_k": 32}' \
  --additional-config '{"tt": {"sample_on_device_mode": "decode_only"}}'

vLLM enables async scheduling by default, and the platform disables it only for models without supports_async_decode. On models that declare it, structured requests keep host sampling unless the server is started with --no-async-scheduling.

Step	Sampling path
Structured prefill	Host sampling
Eligible structured decode	TT device sampling after the sample-time grammar mask arrives
Structured decode with logprobs, another host-only sampling option, or no effective device-grammar support	Host sampling

Eligibility also requires a compatible runtime sampler (not row-sharded), model warmup when tracing, and no logprobs (including logprobs=0) or other host-only sampling option. The ordinary device-sampling limits, max_device_top_k and supports_device_penalties, also apply. Eligibility is batch-wide: one ineligible request moves the whole decode step to host sampling. Block-output models reject structured outputs rather than falling back.

Fallback is selected before forward submission. Once an eligible decode is submitted with deferred device sampling, its sample-time grammar mask must be present and valid. A missing mask or a device-sampling failure is terminal for that runner and requires recovery/restart; it is not retried on host, because the device decode and sampler state may already have advanced.

With VLLM_LOGGING_LEVEL=DEBUG, each device-grammar step logs TT device grammar sampling active. See docs/SCHEDULING.md for the execution flow and docs/MODEL_CAPABILITIES.md for the capability contract. Model authors implement the device grammar contract.

Model Fabric Configuration

A model class may declare model_capabilities["fabric_config"] as a dictionary of keyword arguments to ttnn.set_fabric_config, using TTNN enums and config objects directly. For example:

model_capabilities = {
    "fabric_config": {
        "config": ttnn.FabricConfig.FABRIC_1D_RING,
        "num_planes": 2,
    },
}

The worker applies hardware defaults, then the model's dictionary, then explicit fabric_config and fabric_reliability_mode launch overrides. It forwards the result with ttnn.set_fabric_config(**fabric_kwargs) before opening the mesh. Any TTNN fabric argument is supported, including router_config; TTNN validates the arguments. The plugin does not mutate the model's dictionary or store its TTNN objects in the serialized vLLM configuration. Single-device meshes do not initialize fabric. Models without this capability keep the hardware defaults.

max_model_len And KV Cache Capacity

TT does not profile device memory. get_num_available_blocks_tt() instead derives a block count from the model's max_tokens_all_users, a per-model, per-device token budget declared by the model class in tt-metal (get_max_tokens_all_users, default 131072, sometimes exposed as an environment override such as GEMMA4_MAX_TOKENS_ALL_USERS), plus headroom for vLLM's worst-case block allocation and for sliding-window groups on hybrid models. TTWorker.determine_available_memory() publishes that count both as cache_config.num_gpu_blocks_override and as an equivalent byte budget. The latter is what reaches the engine-side KV planner, which in standard DP runs in a different process from the worker that set the override.

That pool is the budget shared by all concurrent requests. It is set independently of max_model_len, the per-request context limit, except that eligible hybrid models raise it to admit one max_model_len request (see below). Since vLLM sizes and validates the KV cache against it, the two must be chosen together:

--max-model-len	Behavior
numeric, fits the pool	Used as given. Maximum concurrency is roughly pool / max_model_len.
numeric, exceeds the pool	Startup fails in get_kv_cache_configs with the estimated maximum servable length. Block-output models fail earlier, in the TT worker: they must fit max_model_len plus one output canvas. Lower --max-model-len or raise max_tokens_all_users. Eligible hybrid models are the exception: those that set _HYBRID_KV_CACHE_GROUPS_ENABLED and whose get_kv_cache_spec hook returns a nonempty spec. For them the TT worker raises the pool to the smallest size upstream admits for one max_model_len request, plus the output reservation, and logs the change. That pool must still fit device memory, which TT does not profile; KV allocation fails at startup if it does not. With chunked prefill off, every sliding-window group is sized at the full context, so the raised pool can be several times max_model_len.
-1	vLLM auto-fits max_model_len to the pool and syncs the result to the workers and the API-server process. For block-output models the TT worker fits max_model_len to the pool minus one output canvas before vLLM's auto-fit runs.
omitted	vLLM uses the HF-derived value. Startup fails if that value exceeds the pool, except for eligible hybrid models, which raise the pool as in the numeric case.

Explicit -1 auto-fit picks the largest length that fits one request, so it lands at roughly the whole pool with maximum concurrency 1.00x: a single long request can occupy the entire KV cache and stall the rest. Prefer an explicit positive --max-model-len for serving, sized around max_tokens_all_users / max_num_seqs, and treat auto-fit as a convenience for one-off runs.

Runtime Architecture

TTPlatform.check_and_update_config() is the main handoff from vLLM into the TT runtime. It validates configuration, registers TT model architectures, and selects the TT-owned runtime classes through vLLM's extension points:

vLLM config field	TT implementation
parallel_config.worker_cls	vllm_tt_plugin.worker.TTWorker
scheduler_config.scheduler_cls	vllm_tt_plugin.scheduler.TTScheduler or vllm_tt_plugin.lane_scheduler.TTLaneCoordinator

A TTWorker subclass names the scheduler it runs with in tt_scheduler_cls, and sets uses_token_kv_cache = False for a stage with no token KV cache. A requested scheduler_cls is kept when it derives from that class, and replaced by it when that class derives from the request. Any other pair raises ValueError. Lane mode also raises unless the worker declares TTScheduler, inherited or explicit, and the request is unset, TTLaneCoordinator, or a class TTScheduler derives from.

The execution model matches TT hardware characteristics:

• A TT step is either prefill-only or decode-only.
• Token-chunked prefill is available to any model whose tt-metal class declares model_capabilities['supports_chunked_prefill']: a long prompt is split across prefill steps and only the chunk that completes the prompt emits a token. Support belongs to the selected tt-metal serving class, including any TT_MODEL_CLASS_OVERRIDES selection. Model-family registration alone does not establish chunked-prefill support.
• Async scheduling overlaps decode submission with host-side scheduling when the model declares support.
• For Galaxy-generator models (Llama3 70B, Qwen3-32B) and GPT-OSS, --data_parallel_size N runs as N in-process TT lanes scheduled by TTLaneCoordinator (one engine, one device mesh); see Single-Process Galaxy Serving.
• For other models, --data_parallel_size N uses standard multi-process DP: each DP rank runs an independent engine core with its own TT submesh, scheduler, and KV cache. Device groups are discovered at startup and assigned via TT_VISIBLE_DEVICES. A per-rank assignment is accepted only when it exactly matches the discovered group; a conflicting GPU-style --device-ids assignment fails before mesh creation. This is upstream vLLM's standard DP mechanism (no gather/scatter; ranks are fully independent).

For a deeper walk-through of the scheduling and execution model, read docs/SCHEDULING.md. The model-facing rules that make resident async decode safe are documented in docs/DECODE_RELOAD_CONTRACT.md. The defaults, validation and consumers of each model_capabilities key are documented in docs/MODEL_CAPABILITIES.md.

Single-Process Galaxy Serving

Galaxy text models served by the single-execute Galaxy generator (Llama 3.3 70B via TT_LLAMA_TEXT_VER=llama3_70b_galaxy, Qwen3-32B via TT_QWEN3_TEXT_VER=qwen3_32b_galaxy) run on a single Galaxy device mesh, so they use single-process TT lanes: one vLLM engine process with internal TT lanes.

Serve them with the familiar --data_parallel_size N --max_num_seqs M flags; the TT backend transparently maps them to N in-process lanes. No config changes are needed. Wormhole Galaxy example (fabric defaults to FABRIC_1D_RING):

MESH_DEVICE=TG \
TT_LLAMA_TEXT_VER=llama3_70b_galaxy \
python examples/server_example_tt.py \
  --model "meta-llama/Llama-3.3-70B-Instruct" \
  --data_parallel_size 4 \
  --max_num_seqs 8 \
  --async-scheduling \
  --override-generation-config '{"top_k": 32}' \
  --additional-config '{"tt": {"dispatch_core_axis": "col", "sample_on_device_mode": "all", "worker_l1_size": 1344544, "trace_region_size": 220000000}}'

--data_parallel_size 4 --max_num_seqs 8 runs 4 TT lanes of 8 requests each (32 concurrent total); --max_num_seqs is the per-lane capacity. This conversion is specific to the Galaxy generators and GPT-OSS; other model families still run --data_parallel_size as multi-process DP. At startup the backend logs that it is running single-process lane-DP.

Supported Model Families

The plugin registers TT-prefixed model architectures backed by tt-metal model implementations. Current families:

• Llama 3.1 / 3.2 / 3.3 text models (TTLlamaForCausalLM)
• Llama 3.2 vision models (TTMllamaForConditionalGeneration)
• Qwen 2.5 and Qwen 3 text models (TTQwen2ForCausalLM, TTQwen3ForCausalLM)
• Qwen 3.5 text models on Blackhole (TTQwen3_5ForConditionalGeneration)
• Qwen 2.5-VL and Qwen 3-VL vision-language models
• Mistral and Mistral 3 multimodal models
• Gemma 3 multimodal models
• Gemma 4 text-only models (TTGemma4ForCausalLM, TTGemma4ForConditionalGeneration, TTGemma4UnifiedForConditionalGeneration)
• DiffusionGemma block-output models (TTDiffusionGemmaForBlockDiffusion, TTDiffusionGemmaForCausalLM)
• DeepSeek V3 (TTDeepseekV3ForCausalLM)
• GPT-OSS 20B / 120B (TTGptOssForCausalLM)

Model availability, supported device shapes, max sequence limits, and required environment variables are documented in the corresponding tt-metal model demos.

Registering models dynamically (EXTRA_MODELS_DIR)

Instead of adding a hard-coded line to platform.py, a model can be registered at startup by dropping a bundle folder under a directory named by the EXTRA_MODELS_DIR environment variable. Each subfolder holds a vllm_metadata.json and the adapter class (plus its dependencies):

$EXTRA_MODELS_DIR/
  my-model/
    vllm_metadata.json      # {"arch": "<HFArch>", "main_class": "module:Class", ...}
    <adapter class + deps>

At import time the plugin scans EXTRA_MODELS_DIR, appends each folder to sys.path (so an installed package of the same name is never shadowed), and registers arch under the plugin's TT-prefixed convention (TT<HFArch>) pointing at main_class. This lets a distribution tool (e.g. tt-kernel) deliver a ready-to-serve model with no source edit to the plugin. The built-in map above stays enabled by default; set TT_VLLM_BUILTIN_MODELS=0 to rely solely on EXTRA_MODELS_DIR.

Selecting a serving class (TT_MODEL_CLASS_OVERRIDES)

A checkpoint whose architecture already has a serving class can be pointed at a DIFFERENT one for a single launch, without editing the plugin or building a bundle. This is how one Gemma 4 checkpoint serves as a plain baseline, as an MTP speculative model, or as a dFlash speculative model.

export TT_MODEL_CLASS_OVERRIDES="TTGemma4ForCausalLM=models.demos.gemma4.tt.generator_vllm:Gemma4DFlashForCausalLM"

Comma-separate several entries. Each is Architecture=module.path:ClassName.

Four things an operator needs to know:

1. Use the TT-prefixed architecture name. TTPlatform.check_and_update_config rewrites every checkpoint architecture in place with a TT prefix before the registry is consulted, so TTGemma4ForCausalLM is the name that resolves. A bare name is accepted and normalised, but the prefixed form is what takes effect.
2. Name the architecture the checkpoint actually resolves to. Several aliases can share one serving class -- Gemma 4 has six -- and overriding one alias leaves the rest on the default class. Check the architectures field of the checkpoint's config.json.
3. Precedence. Overrides register first and unconditionally, so they outrank both EXTRA_MODELS_DIR bundles and the built-in map for the whole process. Every applied override is logged at INFO as Applied TT_MODEL_CLASS_OVERRIDES: <arch> -> <target>; if that line is missing from the server log, the variable was not set in the environment the engine actually started in.
4. The target must be importable in the worker process, i.e. on PYTHONPATH alongside tt-metal's models/ tree.

The registration environment variables, in the order they are consulted: TT_MODEL_CLASS_OVERRIDES, EXTRA_MODELS_DIR, then the built-in map (TT_VLLM_BUILTIN_MODELS=0 disables it). TT_LLAMA_TEXT_VER, TT_QWEN3_TEXT_VER and TT_QWEN35_TEXT_VER select a version WITHIN a built-in family and are unrelated to the above.

Speculative Decoding

The plugin supports greedy speculative decoding through a runner-model contract. TTModelRunner builds candidate blocks, asks the selected tt-metal adapter to verify candidates, accepts matching draft tokens, and reports committed tokens to vLLM. The selected adapter must declare supports_spec_decode=True and return an admissible SpecPlan from spec_plan. Architecture registration alone does not enable speculation.

Two drafting methods execute today:

speculative_config.method	Proposal source	Required model implementation
ngram	vLLM's host NgramProposer	spec_plan and multi-position verification; no model-side proposer
custom_class	The selected adapter's propose_draft_tokens	spec_plan, multi-position verification and device_propose in spec_requirements

Both methods currently use argmax_ids verification and host greedy acceptance. docs/install-vllm-tt.sh installs numba for NgramProposer. The CPU sampled-acceptance helper exists in spec_accept.py, but TTModelRunner does not call it. logits and fused_sample execution, lossless sampled speculation, scheduler-owned paged drafter caches and TTLaneCoordinator execution are not implemented for this contract.

Model selection and launch settings

The following distinctions describe the Gemma serving classes in tt-metal as of 2026-10-01 (Gemma4 source, DiffusionGemma source). The tt-metal model recipe determines supported checkpoints, hardware, batch sizes and context lengths; these declarations alone do not establish validation for every configuration.

tt-metal serving class	Execution path	How the operator selects speculation
Gemma4DFlashContractForCausalLM	Runner-model contract, output_tokens_per_step=1	Select this class and pass the custom_class settings below
Gemma4DFlashForCausalLM	Adaptive block output; the model owns its draft/verify loop	Select this class without speculative_config
Gemma4MTPForCausalLM	Adaptive block output; no MTP contract adapter is present	Select this class without speculative_config
DiffusionGemmaForCausalLM	Fixed block output, not speculative decoding	Follow DiffusionGemma; do not pass speculative_config

New speculative adapters should implement the runner-model contract. Migrating the existing dFlash and MTP speculative block-output classes does not require removing DiffusionGemma's fixed block output.

Gemma4DFlashContractForCausalLM admits max_num_seqs > 1, but the current adapter speculates only for a request that decodes alone. Shared prefill or a peer joining the request's decode step ends speculation for that request; the request continues ordinary decode and does not resume drafting when the peer leaves. Temporary unscheduling alone preserves the request's retained proposal. Admitting concurrent requests does not establish batched dFlash speculation.

For a compatible Gemma4 dFlash model recipe, select the contract class for the architecture the checkpoint uses. For example, a checkpoint that resolves to TTGemma4ForCausalLM uses:

export TT_MODEL_CLASS_OVERRIDES="TTGemma4ForCausalLM=models.demos.gemma4.tt.generator_vllm:Gemma4DFlashContractForCausalLM"
export GEMMA4_DFLASH_VERIFY=5
export GEMMA4_CHUNKED_PREFILL_TRACE=0

Add the following arguments to the compatible model's server command. The model value is a dispatch marker, not a checkpoint or an importable proposer:

--speculative-config '{"method":"custom_class","model":"vllm_tt_plugin.model_owned_drafter","num_speculative_tokens":5}'
--no-async-scheduling

Gemma4DFlashContractForCausalLM.spec_plan uses effective_k=GEMMA4_DFLASH_VERIFY (default 5). spec_plan rejects smaller requested draft lengths, and TTPlatform reduces larger requested draft lengths to effective_k. To request async scheduling, set GEMMA4_CONTRACT_ASYNC=1 before model import and replace --no-async-scheduling with --async-scheduling. Both supports_async_decode and supports_async_spec_decode must be true. Verification steps remain serialized through acceptance and proposal; eligible ordinary decode steps can overlap. Standard multi-process DP is distinct from the unsupported TTLaneCoordinator path. Device validation of contract async execution with DP greater than one remains unverified here.

Two model-side limits require attention in current dFlash recipes:

• Keep GEMMA4_CHUNKED_PREFILL_TRACE=0: traced prefill replay does not capture the dFlash hidden-state taps (tt-metal #58432).
• A bounded sliding KV ring needs enough capacity to preserve the attention window during verification. The contract adapter declines drafting when the bounded sliding KV ring cannot safely hold the candidate block; this can reduce long-context generation to ordinary decode. The contract CI recipe proposed in tt-metal #57855 uses GEMMA4_SPEC_RING_HEADROOM_BLOCKS=16 with a 64-token page size. Extra ring capacity consumes KV memory, so use the model recipe's matching memory budget. Automatic sizing remains tracked by tt-metal #57788. The adaptive block-output dFlash class only warns about insufficient ring capacity and can corrupt in-window KV (tt-metal #57701).

Request controls and sampling limits

Use temperature=0 with no penalties to exercise greedy speculation. TTModelRunner._publish_draft withholds drafts for requests with non-zero temperature or penalties.

A sampled request never enters a verify. A verify in argmax_ids mode commits the target argmax on every row, so it cannot apply a row's temperature or penalties. When the model supports narrow decode, every step that includes a sampled or penalized request runs as the ordinary decode, with the selected model's normal sampling path; the other rows' drafts on that step are dropped, so no request speculates while a sampled one shares its batch. A model without narrow decode verifies every decode step, so TTPlatform.validate_request refuses a request with a non-zero temperature or a penalty on that launch.

TTPlatform.validate_request rejects logprobs, structured output, logit_bias, bad_words, allowed_token_ids and nonzero min_tokens on a speculative launch. Prompt logprobs remain unsupported for all TT models. Only vLLM's Python frontend runs these checks, so a speculative launch with VLLM_USE_RUST_FRONTEND=1 is refused at startup.

See the model contract for candidate shapes, acceptance, K+1 lookahead reservation and lifecycle hooks. See the capability reference for defaults and validation, and the speculative server tests for the dummy model's validation scope.

Operational Constraints

TTPlatform rejects or adjusts unsupported feature combinations early, giving a clear error before anything reaches the device:

• Tensor parallel and pipeline parallel execution are provided by the models internal implementation, not exposed at the vLLM level.
• Speculative decoding supports greedy argmax_ids verification with ngram or the model-owned custom_class drafter on a compatible tt-metal adapter. See Speculative Decoding for launch settings, model selection and sampling limits.
• LoRA is not currently supported.
• Chunked prefill is gated on the model's declared capability, not on a model_type allowlist. vLLM enables it by default; pass --no-enable-chunked-prefill to opt out. When it stays on, max_num_batched_tokens is left as vLLM set it (2048 for vllm serve / server_example_tt.py, 8192 for LLM(), or an explicit --max-num-batched-tokens). When it is disabled, a budget smaller than max_model_len is raised to max_model_len so a full prompt still fits in one step. Resume offsets need an alignment that depends on the model's program config and on the length of each remaining span, and the tt-metal generator corrects them itself.
• Where chunked prefill is active, multimodal inputs are never split across a chunk boundary.
• Prompt logprobs are rejected at request validation time.
• Prefix caching is enabled only for models that declare TT support for it.
• Async decode overlap is enabled only for models that declare the capability.
• Multi-host MPI data parallelism is not supported.
• vLLM's V2 model runner. The plugin implements only the V1 model-runner contract and pins VLLM_USE_V2_MODEL_RUNNER=0; setting it to 1 is refused.

These are TT runtime characteristics, not vLLM plugin API limitations.

Benchmarking

Offline benchmarking is done by passing --measure_perf to offline_inference_tt.py:

MESH_DEVICE=T3K \
python examples/offline_inference_tt.py \
  --model meta-llama/Llama-3.1-70B-Instruct \
  --measure_perf

Client-server benchmarking can be done with vllm bench serve after starting the server:

vllm bench serve --model meta-llama/Llama-3.1-70B-Instruct \
  --dataset-name random \
  --random-input-len 128 \
  --random-output-len 128 \
  --num-prompts 32 \
  --ignore-eos \
  --percentile-metrics ttft,tpot,itl,e2el

For prefix-cache experiments, use prompts with shared prefixes:

python examples/offline_inference_tt.py \
  --prompts_json examples/prompts_overlapping.json

You can also pass --random-prefix-len <N> to vllm bench serve.

Testing

The plugin ships server-facing tests under tests/tt. Start a vLLM server with a TT model, then run:

pytest tests/tt -v \
  --tt-server-url=http://localhost:8000 \
  --tt-model-name=meta-llama/Llama-3.1-8B-Instruct

For reasoning models, pass --tt-reasoning-token-budget=N to add N to the total generation cap in the bad-word and mixed structured/plain chat tests. Their original caps (100 tokens for bad words; 8/16/64/16 for choice/regex/JSON/plain) form the base allowance, and all output assertions remain in place. Reasoning and answer tokens share the resulting total cap. The default is 0, preserving existing runs. This option does not change raw-completion, recall, penalty or token-count tests.

Choose N from matched native-reference completions with a generous discovery ceiling and normal EOS stopping; record reasoning and answer counts separately. A safety allowance of twice the largest completed reasoning count is a useful starting point, not a guarantee for unseen prompts or seeds. Include enough space for the final answer and framing when choosing the extra allowance. A reference that hits its discovery ceiling or runtime limit has not established a usable limit. Keep its failure visible and investigate it before treating a larger test budget as a fix.

For GPT-OSS, run the chunked-prefill recall gate through its supported Harmony chat rendering. The recall-only options preserve the original user prompts, three rounds of four concurrent requests, and complete-identifier matching:

Configure the GPT-OSS server with --structured-outputs-config '{"reasoning_parser":"openai_gptoss"}' so that analysis is separated from the final content returned by chat completions.

"$VIRTUAL_ENV/bin/python" -m pytest tests/tt -v \
  --tt-server-url=http://localhost:8000 \
  --tt-model-name=openai/gpt-oss-120b \
  --tt-chunked-prefill-budget=8192 \
  --tt-recall-reasoning-effort=low \
  --tt-recall-reasoning-token-budget=1000

The recall cap is the original 24 plus the explicit 1,000-token allowance, totaling 1024 for both reasoning and answer tokens. This finite diagnostic ceiling was used for all 14 original GPT-OSS 120B inputs; it is not inferred from a shorter observed answer. The gate requires normal stop and nonempty final content; analysis, tool handoffs and length-truncated answers cannot pass. The chosen effort is sent to chat completions without rewriting the user's text.

Without --tt-recall-reasoning-effort, recall uses the existing legacy requests and 24-token cap, even if a recall allowance is supplied. The recall-only options do not change the bad-word/structured allowance or any other test's request cap. No model name is used to select this behavior. Check the server's effective chunked-prefill settings before running; a recall pass alone cannot prove that the scheduler split the prompt. Supported-input diagnostic evidence does not replace earlier raw-input scores or qualify an untested runtime revision.

Tests cover request isolation, sampling behavior, penalties, logprobs, host-only parameter handling, and TT utility helpers.

Plugin-local unit tests that do not require a running server live directly under tests/, for example:

"$VIRTUAL_ENV/bin/python" -m pytest tests/test_lane_scheduler.py

These need no Tenstorrent hardware, only an importable ttnn. On a host without tt-metal, put the CI stub on the path instead. This is what the host-tests-stubbed workflow job runs:

PYTHONPATH=tests/host-stubs "$VIRTUAL_ENV/bin/python" -m pytest tests/ --ignore=tests/tt

The stub answers only the device-independent ttnn names the plugin touches and raises from anything that would reach a device, so a test that starts depending on real hardware fails loudly rather than passing against a fake.

Hybrid Attention Models

Hybrid attention models combine sliding-window and full-attention layers. The registered TT model class supplies get_kv_cache_spec(vllm_config) so vLLM can plan KV groups. The model adapter allocates physical storage through allocate_kv_cache_per_layer(per_layer_specs) and consumes page_tables_per_layer when the resolved cache has multiple groups.

See model setup and KV storage for specification names, allocation triples, storage sharing, and page-table routing. Model-specific device and DP support must be validated for the selected adapter and mesh; the interface alone does not establish support.

Development Notes

• Normal Python changes under src/vllm_tt_plugin/ take effect after restarting the Python or vLLM process.
• Reinstall the plugin when package metadata or entry points change, such as edits to pyproject.toml.
• Model capability declarations (model_capabilities dict on the model class) are the preferred way to gate features like async decode and prefix caching, rather than hard-coded model-name checks.

Contributing

Contributions are welcome! Bug reports and feature requests should be filed via GitHub Issues. Bug fixes and new functionality are submitted via pull requests. Pull requests are reviewed weekly. See CONTRIBUTING.md for details.

License

• LICENSE: Overall license for this project (Apache 2.0), except where specified
• LICENSE_understanding.txt: Tenstorrent's clarification of how the Apache 2.0 license applies to this repository

Sampling limits

Single-token-output models may declare model_capabilities["max_device_top_k"] to bound stochastic device sampling. Requests submitted in the current step outside 1..max_device_top_k, including unbounded top-k, use the host sampler with the original distribution. Greedy requests are unaffected. Models without the capability keep their existing routing. To keep requests that do not set top_k on the device, see Default top-k for device sampling.

Block-output models (output_tokens_per_step > 1) cannot declare this capability: their model-owned sampler produces a complete canvas, which the host sampler cannot replace. The combination is rejected at startup.

Request seeds

Sampling batches preserve signed 64-bit request seeds, including through slot compaction. Host generators retain the original seed. A model whose device sampler requires a narrower seed remains responsible for conversion at its device boundary. Host and device samplers need not generate identical random streams.

Host sampling readback

Models can declare supports_compact_host_logits to return only the scheduled host sampling rows. supports_selective_host_readback additionally permits selective device-to-host transfers and requires compact-logit support. Adapters must preserve row order, buffer ownership, and asynchronous read completion. Active or custom logits processors can require the full slot layout. See the host sampling output contract for the sample_rows callbacks and fallback requirements.