IMPORTANT: To view this page as Markdown, append `.md` to the URL (e.g. /max/get-started.md). For the complete documentation index, see llms.txt.
Skip to main content

MAX v26.4

Highlights​

  • The new import-model skill converts a Hugging Face model into a high-performance model for production inference on MAX. The new debug-model skill also finds and resolves output token divergences compared to the original model. Get the skills from our skills repo, and check out the guide to bringing up a model with AI agent skills.

  • MAX now serves many common models on M3 and newer Apple silicon GPUs, including the Llama and Qwen model families (among others that can fit into the system memory). Work continues in the following nightly builds to add support for Apple M1 and M2 GPUs.

Documentation​

MAX models​

  • Added support for Tencent Hunyuan Hy3-preview (HYV3ForCausalLM).
  • Added support for LiquidAI LFM2 (LFM2ForCausalLM).
  • Added support for GLM-5 and GLM-5.1 (GlmMoeDsaForCausalLM) with FP8/NVFP4 weights and FP8 sparse MLA decode for GLM-5.1-FP8.
  • Expanded DeepSeek support:
    • Added DeepSeek V2 and V3 (ModuleV3) on a single GPU, with FP8 block-scaled quantization for V3.
    • Added long-context sparse multi-head latent attention (MLA) support for DeepSeek-V3.2.
    • Added tool-call parsers for DeepSeek V3 and V3.1.
  • Expanded MiniMax-M2 support:
    • Added MXFP4 quantization support for MiniMax-M2.
    • Added tensor-parallel attention + expert-parallel MoE (TP+EP) support for MiniMax-M2. Set data_parallel_degree: 1 with runtime.ep_size > 1 to shard attention heads across GPUs while distributing MoE experts via expert parallelism. Both reduce-scatter (default) and allreduce (runtime.ep_use_allreduce: true) collective strategies are supported.
    • Added tensor-parallel support.
    • Added tool-call grammar support and tool-call parsing for MiniMax M2.7.
  • Expanded Gemma 4 support:
    • Added support for native FP8 attention with an FP8 KV cache on B200 (SM100). This roughly matches bf16 accuracy while improving decode throughput and roughly doubling KV cache capacity at the same memory.
    • Added NVFP4 quantization support.
    • Added structured output (response_format) and tool calling with grammar enforcement, and tool-call parsing.
    • Extended tool parameter schemas with JSON Schema anyOf, $ref/$defs, and additionalProperties.
    • Added a reasoning parser.
    • Added MTP speculative decoding (including TP=2) with constrained decoding when MTP is the draft method.
    • Added multi-GPU support for Gemma 4 31B.
  • Kimi K2.5 tool calling now supports interleaved thinking: a single assistant turn may interleave multiple <think>...</think> reasoning blocks with multiple tool-call sections and end with <|im_end|>. The constrained-decoding grammar (used for tool_choice and JSON response_format) admits up to eight tool-call sections with an optional reasoning block before each, and lets the model stop before the cap. This fixes a tool_choice=auto failure where a second tool-call section disabled grammar enforcement for the rest of the request.
  • Expanded speculative decoding model support:
    • Added DFlash speculative decoding for Kimi K2.5 and Llama 3, with CUDA device-graph capture for the Kimi path.
    • Added GQA-based EAGLE drafters for DeepSeek and Kimi.
    • Added Kimi K2.6 Eagle3 speculative decoding and Kimi-K2.6-NVFP4 model support.
  • Expanded pixel-generation model support:
    • Added FLUX.2-Klein-NVFP4 architecture support.
    • Enabled multi-GPU FLUX.2 and FLUX.2-Klein pipelines.
    • Added diffusion pipeline configuration for Qwen-Image.
    • Added TaylorSeer cache support to the Wan T2V pipeline.
    • Added FP8 weight support for Wan 2.2 T2V.
  • Removed FLUX.1 model support.

MAX framework​

Inference server​

  • Improved OpenAI-compatible endpoint responses:

    • Added reasoning_tokens to usage objects when applicable for chat completion responses.
    • Omitted JSON fields set to null from streaming deltas in chat completion responses.
    • /v1/models now includes each model's max_model_len.
    • /v1/completions now reports usage statistics in non-streaming responses.

  • Added temperature and thinking_temperature to CLI/server arguments.

  • Chat completion responses now emit reasoning only under reasoning, aligning with OpenAI's Responses API naming. The reasoning_content alias is no longer included in responses. (Previously we emitted reasoning_content for compatibility with vLLM, SGLang, and the DeepSeek API; vLLM has deprecated reasoning_content in favor of reasoning.) Clients should read chain-of-thought tokens from the reasoning field.

  • response_format JSON schemas with a non-object root are now accepted when the root type is missing or when it's a type union that includes object (for example {"type": ["object", "array", "string"]}); these are valid JSON Schema and compile to a constraining grammar. A root pinned to a single non-object type (for example {"type": "string"}) is still rejected, matching OpenAI's structured-outputs contract.

  • Added a "none" option for runtime.tool_parser and runtime.reasoning_parser in PipelineConfig (CLI flags --tool-parser and --reasoning-parser). Pass none (case-insensitive) to explicitly disable the parser, overriding any architecture-declared default. Leaving the field unset still applies the architecture default as before.

  • Improved the 400 Bad Request response with the underlying message when a prompt is too long for the model, instead of a generic "Value error" response (or, for streaming completions, a 500 Internal Server Error). All architectures now raise a structured PromptTooLongError exposing num_tokens and max_length attributes so callers can handle the failure programmatically. The user-facing message identifies the relevant limit (context window vs sequence length; for example, "Prompt is too long: N tokens exceeds the configured maximum context length of M tokens. Please shorten your prompt").

  • Added support for role: "developer" on /v1/chat/completions, normalizing it to system at the OpenAI-compat route layer. The OpenAI o1/o3 chat-completion spec uses developer in place of system, and recent OpenAI SDKs emit it by default.

  • Added two opt-in server flags for accepting OpenAI-compatible requests that the default behavior rejects:

    • --allow-unsupported-logprobs: when a request asks for logprobs against a runtime that cannot honor them (today, the overlap scheduler), MAX logs a warning and serves the request without logprobs instead of returning a 400.

    • --allow-extra-request-fields: unknown top-level fields on /v1/chat/completions and /v1/completions request bodies are dropped (with a warning) before Pydantic validation, instead of returning a 400. Useful when an upstream proxy sends vendor-specific fields that MAX doesn't need to honor.

    Both flags default to False; the existing strict behavior is unchanged. The corresponding 400 error messages now reference the new flags. As a side effect, the legacy /v1/completions route now surfaces InputError detail strings to the client instead of the generic "Value error." message.

  • Expanded auto-enablement of overlap scheduler and device graph capture:

    • Auto-enabled for OLMo, OLMo2, OLMo3, Qwen2, Qwen3, Qwen3 MoE, and Qwen3.5 text-generation models.
    • Restricted auto-enablement to text-generation pipelines; embedding pipelines (for example Qwen/Qwen3-Embedding-0.6B) are excluded.
    • Enabled HIP device graph capture by default for supported ROCm architectures.

  • Expanded overlap scheduling with structured output:

    • Overlap scheduling now supports speculative decoding together with structured output (response_format and tool-call grammars).
    • Tool-call grammars are no longer gated behind --enable-structured-output.

  • The local and tiered KV connectors now support EAGLE speculative decoding with an MLA target model and an MHA draft model, and can offload the draft model's KV cache to host memory or disk alongside the target model's cache.

  • Extended top_k sampling to support values greater than 255.

  • Benchmark request payloads now forward an OpenAI-style tools=[...] field on chat-completions requests. SampledRequest and RequestFuncInput gained a tools: list[dict] | None = None field; OpenAIChatCompletionsRequestDriver serializes it into the POST body when set. Datasets that supply per-row tool schemas (currently nemotron-opencode) now exercise the server's tool-call grammar / structured-output path end-to-end. Pass enable_tool_calls=False on Nemotron-OpenCode to suppress forwarding.

  • Removed multi-step decode from the text-generation pipelines. The flag --max-num-steps no longer works.

Server metrics​

  • Added a per-phase startup breakdown to the maxserve.model_load_time Prometheus histogram (milliseconds), previously only available in the server logs. In addition to the existing untagged model-load aggregate, the model worker now records each startup phase on the same metric split by a component tag (build, compile, init, graph_capture, pinned_memory, spawn, and total), so a single metric can be plotted and broken down by startup phase to track pod startup time in production.

  • Added a maxserve.time_per_output_token Prometheus histogram (milliseconds). Emitted once per request, it reports the mean decode-phase latency per generated token (decode_time / (num_generated_tokens - 1)), excluding the first token and prefill time. Because the denominator counts the tokens the model actually produced, the metric accounts for speculative decoding.

  • The maxserve.batch_size Prometheus histogram is now labeled by batch_type (CE for prefill, TG for decode), so the token-generation (decode) batch size can be observed separately from prefill. For the prefill token-count view, use maxserve.batch_input_tokens (also labeled by batch_type). Existing aggregate queries over maxserve.batch_size continue to work; selectors that pin a single series now gain the batch_type dimension.

  • Added Prometheus metrics for the API-server ingress backlog: requests accepted by the API server but not yet handed off to the model worker (still API-side, for example in tokenization):

    • maxserve.num_requests_awaiting_admission: an up/down counter with the live value (incremented on arrival, decremented at handoff).
    • maxserve.requests_awaiting_admission: a companion histogram that captures the distribution and tail (p50/p99) over time. A persistently high value points at a backlog in the API server rather than in the scheduler queue (the latter is visible via maxserve.num_requests_queued).

  • Added Prometheus metrics for the egress (response) path, which show whether the API server is shipping tokens back to clients slower than the model produces them:

    • maxserve.num_responses_buffered: a gauge sampling the total model-worker responses received but not yet streamed to clients, with a companion maxserve.responses_buffered distribution histogram.
    • maxserve.response_queue_time: a millisecond histogram of how long a response waits in the API server's per-request output queue before the streaming layer consumes it.

    Together they surface API-side egress bottlenecks (de-tokenization, serialization, slow clients) and the associated unbounded-output-queue memory growth.

  • Added the maxserve.num_requests_queued OTel/Prometheus metric (changed from an UpDownCounter to a synchronous Gauge). The gauge is sampled once per scheduler iteration and reports the depth of the scheduler's CE / prefill queue (the same value as the Pending: N reqs line in scheduler logs). It's published by every text-path scheduler that drives BatchMetrics: TokenGenerationScheduler and PrefillScheduler (via TextBatchConstructor), and DecodeScheduler (via len(pending_reqs) + len(prefill_reqs)). This shows queue buildup during overload conditions.

max CLI​

  • Added --devices=gpu:all to use every visible GPU.

  • Removed the default value for --devices; omit --devices to use the model or config default.

  • The benchmark now defaults --seed to a fixed value instead of drawing a fresh random seed on each run. The seed drives the workload generator (input/output lengths, session structure, content), so a fixed default makes repeated and scheduled runs reproducible and keeps run-to-run deltas reflecting the change under test rather than workload-draw variance. To opt into a fresh seed, pass --seed none on the CLI (or seed: null in a workload/config YAML); the drawn seed is logged and recorded with the results so the run stays reproducible after the fact.

  • Added --profile to max pipelines generate for rudimentary, one-command profiling. With Nsight Systems (nsys) on PATH and an NVIDIA GPU, the timed run is captured into an .nsys-rep file and a ranked top-N GPU kernel summary is printed. Without nsys, a Python/CPU profile is produced from cProfile. The capture window is bounded by cudaProfilerStart/Stop so warmup and graph-compile time are excluded. Use --profile-output to override the report path.

  • Added --profile to max benchmark as a synonym for --trace that also prints a ranked top-N GPU kernel summary at the end of the run. The server still needs to be launched under nsys launch (matching the existing --trace requirement); --profile removes the "now run nsys stats by hand" step.

  • Added architecture YAML recipe support to the CLI:

    • Added reusable recipe files under each architecture's recipes/ directory.
    • Wired recipes to max serve and max benchmark via --config-file.
    • CLI flags override recipe YAML values when both are supplied.

  • Added the nemotron-opencode dataset for max benchmark, backed by nvidia/Nemotron-SFT-OpenCode-v1. Each row is a full Qwen3-Coder OpenCode trace (system prompt, multi-turn user/assistant/tool messages, and tool schemas). Each subset is multi-GB, so the loader streams via datasets.load_dataset(..., streaming=True) and pulls only enough rows to satisfy --num-prompts.

  • The serving benchmark now reports a per-turn KV cache retention percentile metric for multi-turn workloads. For each turn after the first, it compares the server-reported cached prefix against the block-aligned prefix carried over from the previous turn, surfacing when cached tokens are dropped between turns (distinct from the existing cached-token-rate metrics, whose denominator includes new and uncacheable tokens). The KV cache block size used to align the expected prefix is configurable via --kv-block-size (default 128); match it to the server's --kv-cache-page-size.

Python API​

  • Added max.graph.Module for grouping multiple Graph instances into a single compilation unit, replacing the previous alias for the underlying MLIR module. Construct one with Module() and pass it as the module= argument to each Graph; the resulting Module is what you hand to InferenceSession.load_all() to compile every graph together.

    • Removed Graph.empty_module() in favor of max.graph.Module(), and Graph now exposes a module property returning the Module it belongs to.

  • InferenceSession.load_all() now returns a dict[str, Model] keyed by each model's sym_name (the name of its mo.graph op), instead of a list[Model] ordered by MEF position. It also now supports max.graph.Module as a model input, so you can compile a pre-built module containing multiple mo.graph ops directly.

    To migrate positional unpacking call sites, index the returned dict:

    # Before
module = Graph.empty_module()
with Graph("vision", input_types=..., module=module): ...
with Graph("language", input_types=..., module=module): ...
vision_model, language_model = session.load_all(graph, ...)

# After
module = Module()
with Graph("vision", input_types=..., module=module) as vision_graph: ...
with Graph("language", input_types=..., module=module) as language_graph: ...
models = session.load_all(module, ...)
vision_model = models[vision_graph.name]
language_model = models[language_graph.name]

  • Added CompiledModel, which represents the compile artifact. It holds the MEF bytes and one or more sub-models; it is not directly executable. Model continues to represent the runnable, post-init handle.

  • Added CompletionFlag, an 8-byte completion flag in pinned host memory mapped into a device's address space. It lets host code signal a GPU stream (or peer host observer) by writing a 64-bit value to a single location visible to both. Currently CUDA-only; constructing against any other backend raises RuntimeError.

  • Added CompiledModel.export_mef(path) to get the compiled artifact as a MEF file. It serializes the compiled artifact directly, without requiring the model to be initialized on a device, so it works in cross-compilation / virtual-device scenarios where the target device isn't attached. max.experimental.nn.Module.compile() also returns the same artifact.

  • Added Graph.copy() to create a deep copy of a graph's MLIR module (backed by a new Operation.clone() in max._core). The copy shares no MLIR state with the original graph, so you can hand it to another thread (for example, for background compilation) while the original continues to be staged or executed.

  • Added DeviceStream.wait_for_host_value(flag, value) for dispatching a Python callable onto an explicit AsyncRT worker pool from a host-function node and gating the GPU stream on its completion (via the CompletionFlag). The kickoff trampoline returns immediately, letting the GPU stream proceed concurrently with the worker; a downstream wait_for_host_value blocks the stream until the worker stores value. The __unsafe_ prefix marks that the API has no safety net for callbacks that capture state outliving the compiled graph.

  • Added max.nn.kernels.wait_host_value(), which stalls the device stream until a 64-bit host-visible flag reaches a given value. It lets a captured forward graph gate a downstream consumer kernel on CPU-produced data while the rest of the forward body runs concurrently.

  • Increased the default allreduce signal buffer size from 513 MiB to 1025 MiB per GPU (max.nn.comm.allreduce.Signals.NUM_BYTES and the matching constant in max.experimental.realization_context). The previous 512 MiB scratch couldn't hold the per-peer allgather intermediate for models with large hidden dimensions (for example, Kimi-K2.5 at hidden_dim=20480 with max-batch-input-tokens=16384 needs 640 MiB in bf16). This adds ~512 MiB of per-GPU memory use for any multi-GPU model.

  • Added max.experimental.functional.ceil(), an element-wise unary op that rounds each element of a floating-point tensor up toward positive infinity. Complements the existing floor(), round(), and trunc() ops.

  • max.experimental.functional.while_loop() now passes Tensor (not TensorValue) into its predicate and body callbacks. Callbacks can use ordinary Tensor operations directly, without wrapping arguments via Tensor.from_graph_value(...).

  • max.experimental.nn.Module.compile() now emits the same log sequence that pipeline-level CompilationTimer produces, and wraps the compile body in max.profiler.Tracer spans so an nsys capture with MODULAR_ENABLE_PROFILING=1 shows compilation as named ranges. Every ModuleV3 caller, including pixel-generation pipelines that previously compiled silently, now gets this observability for free.

  • max.experimental.nn.Module.load_state_dict() and Module.compile(weights=...) now accept an auto_cast keyword (default False). The framework remains strict by default. When you pass auto_cast=True, loaded weights are automatically cast between float32 and bfloat16 when shapes match, logging a single summary message per load instead of raising. Other dtype mismatches (float16, fp8, fp4, integers, etc.) continue to raise as before. This removes the need for per-adapter astype shims when checkpoint dtypes differ from the module's declared parameter dtype. MAX pipelines opt in via the MODULAR_AUTO_CAST_WEIGHTS environment variable (default true, parsed by max.pipelines.lib.weight_loading.auto_cast_weights_from_env).

  • CPUMetricsCollector in max.diagnostics.cpu is now used as a context manager instead of start/stop, and exposes get_stats() instead of dump_stats(), matching the interface of GPUDiagContext.

  • Added Tensor.from_dim for building symbolic-dimension predicates inside traced graph regions (for example, comparing a tensor dimension against a compile-time constant inside ops.cond). from_dim also works in eager mode.

  • Bound AsyncValue with asyncio interop so InferenceSession.compile() and related compile paths can be awaited from async Python code.

  • Added caching for eager-compiled models for built-in op graphs to speed repeated eager execution.

  • Raised the default MAX_INTERPRETER_MAX_OPS threshold from 30 to 1024 so larger eager graphs stay on the interpreter path without setting the env var.

MAX kernels​

  • The GPU scatter_nd kernel now vectorizes its slice copy with 128-bit accesses when the slice size, row strides, and base pointers are all vector-aligned (scalar fallback otherwise). On B200 this brings the measured row-scatter benchmark (4096 fp32 rows of 1024 columns into a 131072-row table) from 0.25 ms to 0.19 ms end to end, on par with PyTorch index_put_ for the same shape.

  • The GPU scatter_nd kernel (all mo.scatter_nd* variants) now parallelizes over total update elements instead of one thread per index row, removing the serial per-row slice copy. Slice-style scatters (for example row updates into a large table) run roughly 1.7x-2.8x faster end to end on B200 depending on shape; the update kernel itself speeds up far more, but the op's mandatory data-to-output copy now dominates its runtime. Element-style scatters and the CPU path are unchanged.

  • Removed the use_blocking_impl parameter from the foreach custom op helper (and the underlying elementwise primitive), and removed the analogous single_thread_blocking_override parameter from the concat and concat_shape kernels and the reduction-based kernels. Work is always dispatched the same way, with a single worker used automatically when the problem size is small. Removed the dedicated small-tensor concat fast path in favor of the existing serial/parallel dispatch.

  • The axis of the nn.split, nn.cumsum, and fused_concat kernels is now a keyword-only compile-time parameter instead of a runtime argument. Pass it in the parameter list, for example split[..., axis=axis](input, outputs, ctx), cumsum[dtype, exclusive, reverse, axis=axis](output, input), or fused_concat[..., axis=axis](...).

  • Added ManagedTensorSlice.shape_coord() and ManagedTensorSlice.strides_coord(), which return Coord values that preserve the compile-time static shape and stride information that the previous IndexList-returning queries discarded.

  • Improved the short-axis GPU softmax kernel, reaching performance parity on [8, 4096, 1024, 24]-style shapes (~3.1 ms).

  • max.nn.kernels.scatter_set_constant now raises a ValueError at graph construction time when indices doesn't have a statically-known inner dimension of 2. Each indices row is a (row, col) coordinate pair; the kernel previously read the second coordinate out of bounds for [num_indices, 1]-shaped indices, silently scattering the fill value to a heap-content-dependent location (or dropping the write entirely). The Mojo kernel also now raises unconditionally on malformed index shapes instead of only under MODULAR_MAX_DEBUG_ASSERT_LEVEL.

  • Improved SM100 (Blackwell) GEMM kernel performance:

    • Added a cpasync + mma.sync pipelined GEMM kernel for small M and N shapes, providing higher throughput than gemv_splitk for those cases (for example, nearly 2x on 28x384x7168).
    • Added a warp-specialized TMA epilogue warp that asynchronously loads epilogue data into shared memory, improving output-pipeline throughput.
    • Added an experimental 1D residual bias epilogue (matmul + broadcast + add fusion) that uses TMA to asynchronously load the bias tensor.
    • The Mojo SM100 FP4 matmul kernel is now used for small shapes (m <= 128), improving performance and supporting all fusion types on those shapes.
    • Enabled Programmatic Dependent Launch (PDL) for the SM100 (Blackwell) FlashAttention-4 prefill kernel, letting back-to-back attention grids in a stream overlap launch and prologue latency. This reduces per-launch overhead most for shorter sequences (measured ~1.05x–1.5x faster on B200, bf16, head_dim=128 across seq lengths 128–2048). On by default; disable with -D MHA_PDL=false.

  • Improved AMD MI355X GPU kernel performance

    • Added FP8 MLA prefill and HKMhaPrefill long-context MHA prefill
    • Added 4-wave matmul and implicit-GEMM convolution kernels (FP8/BF16/FP16) and native 3D implicit-GEMM convolution
    • Tuned MLA/MXFP4 decode including split-K
    • Overlapped Kimi K2.5 MXFP4 shared-expert compute with expert-parallel communication on MI355
    • Fixed HK MHA prefill correctness for partial-K, odd-N, and depth-64 geometries that corrupted FLUX image generation

  • Improved Blackwell (SM100) GPU kernel performance

    • Added sparse MHA prefill and decode kernels and FP8 MHA prefill/decode support
    • Added sparse MLA prefill kernel (mo.mla.prefill.sparse.paged.fp8)
    • Added fused grouped matmul + SwiGLU + NVFP4/MXFP8 quantization for Kimi K2.5
    • Enabled PDL for split-K MHA decode
    • Fixed an SM100 FP8 MoE decode regression in grouped GEMM
    • Improved FP8 grouped GEMM for non-aligned problem sizes
    • Made LayerNorm ~13% faster at FLUX/DiT-32B tensor shapes

  • Improved Apple silicon support:

    • Added a simdgroup-tiled matmul kernel for the Apple M5 GPU, bringing neural-accelerator-backed matmul to the MAX framework. In-range MAX matmuls (m >= 64, n >= 64, k >= 16; ragged K supported) now use it: fp16/bf16 always, and fp32 a/b by default (accepting the simdgroup MMA's fp19 truncation). Set MODULAR_APPLE_M5_ALLOW_LOSSY_F32_MATMUL=0 for the precise naive fp32 path. An 8x8 simdgroup_matrix matmul path also covers Apple M1-M4 GPUs, roughly 20x faster than the previous naive matmul on those GPUs in local microbenchmarks.
    • Added a naive split-K decode attention kernel for Apple GPUs, opt-in behind MODULAR_ENABLE_APPLE_NAIVE_FA_DECODE (default off). On Apple Metal, token-generation attention otherwise falls back to the unfused mha_gpu_naive; this adds a paged-KV-cache decode path supporting MHA and GQA. It is a correctness-first implementation and not yet performance-optimized.

Breaking changes​

  • Removed the max.interfaces top-level module. Pipeline request/context types and helpers moved into max.pipelines (for example max.pipelines.request, max.pipelines.modeling); SchedulerResult moved to max.serve. No deprecation shim remains at the old paths.

  • KV cache management has moved from max.kv_cache to max.pipelines.kv_cache. Update imports accordingly:

    # Before
from max.kv_cache import PagedKVCacheManager, DummyKVCache

# After
from max.pipelines.kv_cache import PagedKVCacheManager, DummyKVCache

    Deprecation shims with DeprecationWarning remain at the old path.

  • GPU and CPU diagnostic tooling has moved from max.diagnostics to max.profiler: max.diagnostics.gpu → max.profiler.gpu and max.diagnostics.cpu → max.profiler.cpu. Update imports accordingly. Deprecation shims with DeprecationWarning remain at the old paths.

  • Removed support for speculative decoding with a standalone draft model. Use eagle, mtp, or dflash drafters instead.

  • All Mojo custom ops must now declare their ctx parameter as DeviceContext instead of DeviceContextPtr. The DeviceContextPtr type has been removed from the Mojo standard library; multi-device ops should declare their variadic context argument as DeviceContextList[N] (for details, see the corresponding Mojo v1.0.0b2 release notes).

  • max/python/max/benchmark/benchmark_throughput.py, deprecated in v0.26.3, has been removed.

  • Removed the LMCache KV cache connector (use local or tiered connectors).

Fixes​

  • Fixed structured output (response_format: json_schema and grammar-guided tool calling) intermittently emitting raw control characters inside JSON string values on models that use a byte-level BPE (TikToken) tokenizer, producing invalid JSON. The constrained-decoding adapter fed llguidance the tokens' byte->unicode surface bytes (for example, a raw newline rendered as Ċ) instead of their true bytes, so the grammar mask admitted control-char tokens as legal string content. Token bytes are now recovered via the tokenizer's byte_decoder, so raw control characters are correctly excluded. Fast-tokenizer checkpoints were unaffected.

  • Fixed an expert-parallelism dispatch assertion (Cannot dispatch EP kernel with N input tokens when the maximum tokens per rank is N-1) that fired whenever --max-batch-input-tokens wasn't evenly divisible by the tensor-parallel degree. The EP per-rank cap now uses ceiling division to match the ragged binning of reducescatter in TP-attention + EP-MoE mode, so the largest shard fits in the dispatch buffer. Affects DeepSeek-V3, Kimi-K2.5, MiniMax-M2, Qwen3, and Step3.5 deployments configured with non-divisible batch sizes.

  • MODULAR_DEBUG=ir-output-dir=<dir> (and the equivalent [max-debug] ir-output-dir = <dir> config-file entry and InferenceSession.debug.ir_output_dir = <dir> Python setter) now actually dumps per-stage MLIR files to the configured directory. The option was previously parsed but no compiler stage consulted it, so users had to fall back to the legacy MODULAR_MAX_TEMPS_DIR env var. Both spellings are now honored.

  • Fixed page-size=128 support in attention kernels when head depth is ≥ 256, which previously produced incorrect results.

  • Fixed pixel-generation serve configs silently dropping denoising-cache flags (for example TeaCache), so cache tuning options passed through recipes now take effect.

  • Fixed Wan 2.2 image-to-video (I2V) producing a flat green frame when conditioning inputs were misaligned across the dual-transformer MoE path.

  • Fixed a KV-cache host-buffer lifetime bug where cache_lengths and lut_tables could be freed before asynchronous H2D copies finished, causing intermittent corruption under offload or tiered-KV workloads.

  • Fixed an FP8 dynamic-quantization bug that mis-quantized near-zero groups on NVIDIA GPUs (writing NaN into FP8 activations and the FP8 KV cache, surfacing downstream as non-finite logits).

  • Fixed a KV cache offloading correctness bug that corrupted output for multi-cache models (such as Gemma 4's interleaved sliding-window plus global attention) when the local or tiered KV connector was enabled.

  • Fixed JSON response_format and tool-call grammars not being enforced for Kimi K2.5 vision-language checkpoints.

  • Fixed an intermittent constrained-decoding correctness bug under EAGLE speculative decoding.

  • Fixed CreateChatCompletionRequest rejecting explicit null values for optional fields such as tool_choice, tools, and response_format. OpenAI-compatible clients (LangChain, JS SDKs, anything that serializes a dataclass with a None field) that emit "tool_choice": null instead of omitting the key are now accepted, matching the behavior of other OpenAI-compatible inference servers.

  • Fixed miscalculation of scratch space required by some models. The estimate was wrong by a factor of 1/num_devices because each device only needs scratch for its own portion of the collective problem. For example, the fix reduces the footprint for Kimi-K2.5 from 16 GB to 4 GB.

Mojo language​

For all the updates to the Mojo language, standard library, and tools, see the Mojo release notes.

Models

Cookbook

Puzzles

Blog

Community

Contact