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

MAX v26.7

Highlights​

  • New models: Added Muse Glimmer (meta-models/Muse-Glimmer-30B) and MiMo-V2.6-Flash (XiaomiMiMo/MiMo-V2.6-Flash-RL), and added support for video input with Qwen3.5 and Qwen3-VL.

  • New Decisions API: /v1/decisions returns probabilities for classification, scoring, and yes/no questions without generating text. Available on every text-generation model. See the decisions guide.

  • More capable speculative decoding: DFlash, DFlash2, and DSpark can now sample draft tokens using each request's sampling settings. MAX can also dynamically choose how many draft tokens to verify based on which width delivers the highest measured decode throughput. Learn more in the speculative decoding guide.

  • Faster startup and GPU execution: Kimi K2.5 MXFP4 warm starts on eight AMD MI355 GPUs drop from about 350 seconds to 48 seconds, while peak host memory during loading drops from about 590 GiB to 13 GiB. New GPU paths for graph ops deliver compute-kernel speedups up to 152x, and Qwen3.5 and Qwen3.8 recurrent prefill kernels improve by up to about 38x.

  • Expanded caching: Jenga becomes the default KV cache. Qwen3.5-family and Inkling models gain prefix caching, and recurrent-state models can use tiered KV connectors.

Documentation​

  • Added a chat templates guide that explains how to use the --chat-template flag with the max serve CLI command.

  • Added a prefill-decode disaggregation guide that describes how to run prefill and decode on separate workers.

  • Added a decisions guide that describes how to use the /v1/decisions endpoint to answer typed questions about an input with a probability for every option, such as to classify text or route tickets.

Models​

  • Added the Muse Glimmer architecture (MuseGlimmerForConditionalGeneration_ModuleV3) for meta-models/Muse-Glimmer-30B, with text and image inputs, reasoning and tool-call parsing, decode graph capture, and a serving recipe for one B200 GPU (muse_glimmer_30b.yaml).

  • Speculative decoding for Llama 3 (EAGLE and DFlash drafts) and Gemma 4 (DFlash and both DSpark drafts) now builds its fused graph from the target model and the draft model, instead of using a separate architecture per pairing. The startup banner names the fused architecture after the method and the target, for example LlamaForCausalLM+eagle instead of UnifiedEagleLlama3ForCausalLM.

  • Added the MiMo-V2.6-Flash architecture (MiMoV2ForCausalLM) for Xiaomi's checkpoint, XiaomiMiMo/MiMo-V2.6-Flash-RL. Its 39 sliding-window layers use learned attention sinks, and its Q/K and V head dims (192 and 128) are padded to 256 for paged attention. The MXFP4 routed experts run with MXFP4 activations, and the dense projections run as stored, in FP8 with 128x128 block scales. The architecture needs two or more B200 (SM100) GPUs, since its weights (about 158 GiB) leave no room for the KV cache on one B200. It supports speculative decoding with the DFlash drafter that its checkpoint ships, with greedy and per-row sampled acceptance.

  • Expanded support for Qwen3.5:

    • The fused Qwen3.5 speculative-decoding graph (qwen3_5_with_mtp_graph) now accepts M-RoPE positions, so speculative decoding composes with the vision path instead of excluding it. The input covers the merged [real, draft_1..draft_k] window and is appended after the existing spec-decode signature, so earlier inputs keep their positions. It is declared only when the target runs M-RoPE, leaving text-only exports byte-identical, and a graph exported without it still serves text.
    • Qwen3.5-family models now support serving with DFlash2 speculative decoding (unified_dflash2_qwen3_5), with a Qwen3.8-27B NVFP4 recipe (qwen38_27b_nvfp4_dflash2.yaml).
    • Added video input to the Qwen3.5 and Qwen3-VL architectures, which previously rejected any request carrying a video part. Clips are sampled at 2 fps, sized against a whole-clip pixel budget rather than a per-frame one, and patchified so that each patch row spans two consecutive frames. A clip becomes one vision encoder grid item but several placeholder runs in the prompt, each preceded by a <T seconds> timestamp label, so the model reads the temporal ordering from the prompt text. Send the clip's bytes in videos with a video content part in messages. An OpenAI-compatible request may instead carry a video_url part, which the server resolves and rewrites into that form.
    • Qwen3.5-family and Inkling models now support prefix caching. Their recurrent and convolution state is checkpointed in the KV cache at block boundaries, so a prefix hit resumes from the saved state instead of recomputing the prompt, including with the overlap scheduler and device graph capture.

  • Added a ModuleV3 implementation of the Nemotron-H architecture (NemotronHForCausalLM), selected with --prefer-module-v3. It serves Nemotron-3.5-Lightning-30B-A3B in BF16 and NVFP4, with prefix caching and tensor parallelism. It supports DSpark speculative decoding: the BF16 Nemotron-3.5-Lightning-30B-A3B checkpoint can serve with the NVFP4 DSpark draft (--speculative-method dspark) on one or two B200 GPUs.

  • GLM-5.3 FP8 (zai-org/GLM-5.3) now runs on 8x AMD MI355X GPUs, with a new serving recipe (glm_5_3_fp8_tp_ep_8x_mi355.yaml).

  • Improved startup time for Kimi K2.5 MXFP4 checkpoints, such as amd/Kimi-K2.5-MXFP4, on AMD MI355 GPUs. A warm start on eight GPUs dropped from about 350 seconds to 48 seconds, and peak host memory during loading from about 590 GiB to 13 GiB.

  • Every unified speculative-decoding architecture now passes its full [batch_size] per-row seed tensor to acceptance sampling. A row's tokens no longer depend on its position in the batch, on how many requests share it, or on the batch's verify width. The committed distribution is unchanged, so this buys reproducibility rather than accuracy, and a single-row batch is unchanged.

  • Added the GLM-5.3-Flash architecture (Glm5NextForConditionalGeneration) for zai-org/GLM-5.3-Flash. This model is still in early development.

  • Added the DeepSeek-V4-Flash architecture (DeepseekV4ForCausalLM) for deepseek-ai/DeepSeek-V4-Flash-0731. This model is still in early development.

Inference server​

  • Removed the --rejection-sampling-strategy flag, the SpeculativeConfig field it set, and the RejectionSamplingStrategy type. The flag had no effect because each architecture's acceptance sampler chooses the acceptance rule.

  • Removed the llguidance structured-output backend. Structured and constrained output now uses xgrammar exclusively, and --structured-output-backend no longer accepts llguidance.

  • Removed the mojo_tiered KV connector type. Use rust_tiered instead. A --kv-connector-config that sets mojo_tiered now fails with an error.

  • Removed the sha256_64 value of kv_cache_hash_algo, and a config that sets it now fails validation at startup. It kept the first 8 bytes of each block's SHA-256 digest but chained the next block onto the full digest, so a request whose hash chain grew after its first block (decode past a block boundary, chunked prefill) raised an error. Use sha256 instead. The dKV connector keys sha256 blocks by the first 8 bytes of the digest, so dKV keys are unchanged.

  • Added support for the Decisions API:

    • Added POST /v1/decisions, which answers typed questions (choice, score, yes_no) about an input with a probability for every option, read from the model's next-token scores at the answer position. No text is generated, and usage.completion_tokens is 0. The wire format matches SGLang's /v1/decisions. It needs no server flag: any text generation model serves it, and a model with a decider_config.json is served in the decider prompt format. See Decisions.
    • Added POST /v1/systemone, a System One compatible route over the same scoring path as /v1/decisions, so System One clients such as typesafe-sdk can use a MAX server.
    • Decision models such as Mapika/decider-2b, which read a plain state-first prompt and pick an option letter, are served in that format automatically. The Qwen3.5 text-only decider checkpoints load through the existing qwen3_5 architecture. See Decisions.

  • Improved speculative decoding:

    • --draft-proposal sampled now works with block speculative decoding (DFlash, DFlash2, and DSpark drafts). The draft samples each proposal at the request's temperature, top-k, and top-p and hands the verifier the distribution it sampled from, so acceptance runs true speculative sampling instead of typical acceptance.

    • A speculative architecture that can't sample its draft now refuses --draft-proposal sampled at startup rather than silently drafting by argmax.

    • A num_speculative_tokens_per_batch_size range whose count is 0 now skips the drafter as well as the verification, on the DFlash and EAGLE Llama 3 architectures, on DFlash Gemma 4 31B, on Gemma 4 MTP, and on DeepSeek-V3 MTP, including MTP served tensor-parallel or, for DeepSeek-V3, data-parallel, where every replica skips at once. DFlash skips its whole draft block, while EAGLE and MTP keep the first draft step and skip the rest. Every other architecture is unchanged and still drafts at its configured depth. Behavior without a schedule is unchanged. A draft built with --draft-proposal sampled skips as well.

    • DSpark speculative decoding now has its own method: serve a DSpark draft with --speculative-method dspark instead of dflash. DSpark drafts a block at a time like DFlash but runs a different draft step, so a DSpark draft under any other method now fails at startup with a message naming the method to use.

  • /v1/completions now honors reasoning_split. By default the server's reasoning parser still hides the reasoning span from text and counts it in usage.completion_tokens_details.reasoning_tokens. A request that sends reasoning_split: false skips the parser, so text and logprobs cover every generated token, reasoning span included, as in vLLM.

  • Added the return_token_ids request parameter to /v1/chat/completions and /v1/completions. When it's set, each choice includes token_ids for the generated tokens and prompt_token_ids for the templated prompt, and each streamed chunk includes the token IDs it adds.

  • Text generation now stops on every eos_token_id in a model's generation_config.json, as Hugging Face generate does, in addition to the tokenizer's EOS token and the eos_token_id in config.json. This applies to TextTokenizer, TextAndVisionTokenizer, and the architecture tokenizers built on them. Before, those extra ids were only a default for stop_token_ids, so a request that set its own stop_token_ids, or that bypassed the model's sampling defaults, could generate past them. Models whose generation_config.json lists more EOS ids than config.json (Llama 3.1 Instruct, for example) may now end some responses sooner. stop_token_ids still adds to the set, and ignore_eos still disables all of it.

  • Hardened handling of client-supplied media:

    • Hardened decoding of client-supplied images. Image.open is now restricted to an explicit format allowlist (PNG, JPEG, WEBP, GIF, BMP, PPM, TIFF, TGA, and AVIF where the platform provides it), shrinking the native-decoder attack surface and keeping image bytes away from decoders such as EPS/Ghostscript.
    • Media memory is now governed by a single knob, MAX_SERVE_MAX_MEDIA_BYTES (default 100 MiB). The fetched size of all resolved media (http(s)://, data:, and file:// images and videos) for a request is bounded in aggregate by it, rather than each item being capped independently. No single image may decode to more than it either. That decoded size is estimated from the image header and rejected before the pixel buffer is allocated, which is what catches a decompression bomb (a small on-the-wire image that expands to hundreds of MB). This replaces both the former per-item MAX_SERVE_MAX_BYTES server cap and a separate decoded-pixel limit. It is deliberately separate from MAX_SERVE_MAX_REQUEST_BYTES, which bounds only the request body. A small body can name URLs the server then fetches, so raising one limit should not silently widen the other.

  • Changed the tiered KV connector:

    • The tiered KV connector can allocate its pinned host staging region from the GPU driver instead of registering ordinary host memory, with "host_staging_alloc": "driver" in --kv-connector-config. The default is unchanged. On AMD, registered host memory stays movable: when the kernel compacts memory it migrates those pages, which makes the driver evict the process's GPU queues and revalidate the whole registration before they resume. At a large host tier that takes tens of seconds, long enough to fail a health check. Driver-allocated memory cannot be migrated. It needs HSA_USERPTR_FOR_PAGED_MEM=0 in the environment before the server starts, because the HIP runtime otherwise backs the allocation with movable pages anyway; the server refuses to start rather than run unprotected. The memory counts against the GPU driver's system memory limit, which every process on the node shares, and it is not charged to the server's memory cgroup.
    • A model with recurrent state, such as Qwen3.5 and the other hybrid architectures on the recurrent cache group, can now use a KV connector. The host and disk tiers hold each state checkpoint beside the KV blocks of the boundary it was taken at, and a prompt that has left the device cache resumes from the deepest checkpoint the tiers still hold. Configuring --kv-connector-config beside such a model used to be refused at startup.

  • Changed the default KV cache:

    • The Jenga KV cache is now the default for every model, including custom architecture names and local model paths. Set MODULAR_USE_LEGACY_KV_CACHE=1 to opt out.
    • The legacy KV cache manager (MODULAR_USE_LEGACY_KV_CACHE=1) no longer supports KV connectors or disaggregated inference, and requesting either now fails at startup. Disaggregated inference with a cache that has non-full-attention layers also fails instead of falling back to the legacy manager.

  • MAX now builds against NIXL 1.5.0, whose libfabric (EFA) transport adds a mandatory connection handshake that a NIXL 1.3.0 peer never answers. A MAX built on 1.5.0 cannot exchange KV cache transfers over libfabric with one built on 1.3.0, so disaggregated prefill and decode have to be upgraded together.

Server metrics​

  • Expanded OpenTelemetry traces, metrics, and logs:

    • MAX now honors OTEL_SDK_DISABLED and, for metrics, OTEL_EXPORTER_OTLP_METRICS_ENDPOINT and OTEL_EXPORTER_OTLP_ENDPOINT, so a deployment that set the latter for traces now sends metrics there too, and loses them if that endpoint is gRPC.

    • MAX Serve can now export spans over OTLP/gRPC. Set OTEL_EXPORTER_OTLP_TRACES_PROTOCOL (or the generic OTEL_EXPORTER_OTLP_PROTOCOL) to grpc, and set OTEL_EXPORTER_OTLP_TRACES_ENDPOINT to a gRPC receiver such as http://<host>:4317. Keep the http:// scheme for a plaintext collector: without it the OTel SDK dials TLS. The protocol defaults to http/protobuf, so existing deployments are unchanged, and metrics and logs still export over HTTP.

    • The scheduler's max.phase.prefill and max.phase.decode spans now nest under the request's max.request span on chat and completions requests, instead of beside it or, without a traceparent, in a trace of their own.

    • Added correlation IDs to structured log records on every route: request_id, previously always empty, and, while tracing is enabled with OTEL_EXPORTER_OTLP_TRACES_ENDPOINT set, dd.trace_id, except on a probe that arrives without a traceparent. Structured logging is on by default in the MAX container images; elsewhere, set MODULAR_STRUCTURED_LOGGING=1.

    • Populated batch_id in structured log records that a text-generation model worker logs during a forward pass while tracing is enabled. It was previously always empty, and still is on disaggregated prefill and decode workers.

    • With OTEL_EXPORTER_OTLP_TRACES_ENDPOINT set, MAX Serve starts an HTTP server span for each request except health, version, ping and metrics probes, such as /v1/health. It continues any inbound traceparent and parents the request's max.request span, so the request's logs carry a trace ID without a traceparent. It ends when the response body finishes, so a streamed response's span covers the whole stream.

    • A request can now ask for a kernel trace of the forward passes it runs in by setting the x-max-trace-level header to batch, op, kernel-sampled, kernel or full. This needs MAX_SERVE_KERNEL_TRACE_HEADERS=true, span export through OTEL_EXPORTER_OTLP_TRACES_ENDPOINT, and a MAX build that ships the profiler plugin. Each traced pass exports a max.batch span with GPU spans beneath it, and max.batch spans now link to their requests' spans: every request's at MAX_SERVE_KERNEL_TRACE_LEVEL=batch or higher, otherwise only traced ones. The kernel_trace_max_passes, kernel_trace_max_spans, kernel_trace_max_capture_bytes and kernel_trace_max_batch_links fields in the profiling section of a max serve --config-file set its limits of 64 passes, 20,000 spans, a 64 MiB capture file and 128 links per span.

    • MAX Serve now exports spans only when OTEL_EXPORTER_OTLP_TRACES_ENDPOINT is set. It no longer sends them to Modular's collector by default, and OTEL_EXPORTER_OTLP_ENDPOINT alone no longer turns tracing on. To keep exporting spans, set OTEL_EXPORTER_OTLP_TRACES_ENDPOINT, for example to http://collector:4318/v1/traces.

    • MAX_SERVE_TRACE_PREFILL_BATCH is renamed MAX_SERVE_TRACE_BATCH and now brackets the model-execute call on the decode scheduler as well as prefill. A dispatch line with no matching completion identifies which worker is stuck inside one call, which the prefill-only form could not do for a decode-side stall. Log lines name the role, as Dispatching decode batch: ..., and an execute that raises closes its bracket with Failed ... batch so it is distinguishable from one that never returned.

  • Added tool-calling and structured-output metrics:

    • Added counters for how much traffic uses tool calling and structured output: maxserve.tool_call.requests (the request declared tools, tagged choice), maxserve.tool_call.responses (its response actually contained a tool call), and maxserve.structured_output.requests (tagged kind). The first two together show how often a declared tool inventory is used.
    • Added maxserve.tool_call.tools_per_request, a histogram of how many tools a request declared. Tool schemas are rendered into the prompt, so this is the explanatory variable behind a client's prompt length and grammar compile cost.

  • Added media and vision preprocess metrics:

    • Added the maxserve.media.* family, which accounts for the media work a multimodal request pays for before it reaches the model. Volume and origin come from maxserve.media.items (tagged media_kind and source), maxserve.media.items_per_request and maxserve.media.item_size; the cost of turning a reference into bytes from maxserve.media.resolve_time; the codec cost from maxserve.media.image_size and the maxserve.media.image_decodes / maxserve.media.image_decode_ms counter pair, both tagged format, whose quotient is the mean decode cost for a format; and refusals from maxserve.media.rejections (tagged reason). Cache pressure comes from maxserve.media.preprocess_cache_evictions, maxserve.media.preprocess_cache_size and maxserve.media.preprocess_cache_capacity (all tagged media_kind), which separate a low preprocess-cache hit rate caused by a workload with no repeats from one caused by a cache thrashing at its byte budget. The resolve and decode time is inside maxserve.time_to_first_token, where it was previously unattributable.
    • maxserve.vision.preprocess_cache_hits and maxserve.vision.preprocess_cache_misses are now recorded on every architecture with a preprocess cache, not only on those the API server can ask at admission whether an image is already preprocessed. Elsewhere the pair sat at zero. On those other architectures it comes from the tokenizer's own cache lookup after tokenization rather than from the admission peek, so the two windows differ and the rate isn't comparable across architectures.

  • Changed KV cache connector metrics:

    • Added maxserve.cache.connector_loads_refused and maxserve.cache.connector_offload_blocks_dropped for the tiered KV connector. The first counts loads its host and disk tiers refused, each served by recomputing the blocks instead. The second counts offloaded blocks the host pool had no room for, which rises when blocks pinned for in-flight transfers starve the pool, before the hit rate falls.
    • KVCacheMetrics drops nixl_read_blocks_local, nixl_read_blocks_remote, and the remote_read_ratio property computed over them. No code path ever populated either field, so the ratio returned 0.0 for every caller. Six populated counters are added in their place: dkv_peer_attaches, dkv_peer_attach_failures, dkv_peers_dropped, dkv_peer_loads, dkv_peer_load_failures, and dkv_hints_rejected. They are zero unless the external KV-cache connector is in use.

max CLI​

  • Removed --enable-tool-call-constrained-decode. Use --tool-call-policy=force_unconstrained instead of setting it to false.

  • Added --tool-call-policy, which sets how tool-call arguments are constrained:

    • force_unconstrained: no tool-call grammar.
    • force_strict_false: only the tool-call envelope is constrained; arguments are free-form, ignoring each tool's strict field.
    • force_strict_true_*: every tool is constrained to its argument schema.
    • default_strict_true_*: tools are constrained to their argument schema unless the request sends strict: false.
    • default_strict_false_* (default): tools are constrained to their argument schema only when the request sends strict: true.

    The * suffix is one of and_best_effort or and_reject_unsupported, which controls what happens to a schema that the grammar cannot enforce.

    The default, default_strict_false_and_best_effort, matches OpenAI's Chat Completions API.

  • Changed expert-parallel load balancing:

    • Renamed the --eplb-replicas-per-gpu option to --eplb-redundant-experts-per-gpu (eplb_redundant_experts_per_gpu in config files). The value and its effect are unchanged.
    • Added the --eplb-planner option to choose how expert-parallel load balancing (EPLB) places experts. The new default, per_layer, balances each MoE layer against its own load, while aggregated keeps the previous behavior of one placement shared by every layer.

  • Changes to max generate:

    • max generate with speculative decoding now verifies drafts on decode steps, as serving does.
    • max generate now runs with device graph capture off unless you pass --device-graph-capture, which it then warms up so decode steps replay. It also plans for a batch of one unless you pass --max-batch-size, since it runs one prompt.
    • max generate now raises an error that gives the max serve command to use when you pass --device-graph-capture with --data-parallel-degree above 1, or when the prompt is over a replica's prefill token budget.
    • On AMD GPUs, max generate and AutoModel now turn off the vendor BLAS matmul fallback, as max serve does when it captures device graphs, so they compile the same matmul kernels. Pass --no-disable-vendor-blas-fallback or disable_vendor_blas_fallback=False to keep the fallback.

  • Added --use-dummy-weights to max warm-cache. Compilation only needs tensor names, shapes and dtypes, so with the flag set each safetensors shard is stood in by a sparse, zero-filled stub built from the remote file's header instead of being downloaded, making a compile-only run independent of checkpoint size. It requires --target (compile-only mode) and a safetensors repo on Hugging Face.

  • Added --cascade to max serve, which routes the request to the experimental Cascade server (max.experimental.cascade.serve.main.serve) instead of the standard API server + model worker. The resolved PipelineArgs is forwarded to the Cascade entrypoint, so max serve --cascade --model ... runs the Cascade server with the same model-selection flags. The Cascade deployment context is configurable from the same command: --cascade-transport (http/grpc), --cascade-local-cpu-workers, --cascade-local-gpu-workers, and --cascade-remote-cpu-workers / --cascade-remote-gpu-workers map onto ContextConfig. --host is now exposed for both paths and defaults max serve --cascade to 0.0.0.0, matching max serve, instead of the Cascade entrypoint's localhost default.

  • max serve --disable-vendor-blas-fallback disables the vendor-BLAS matmul fallback on any GPU. Setting MODULAR_DISABLE_VENDOR_FALLBACK=1 in the environment of max serve sets its default. Previously the variable had no effect under max serve.

  • Changes to overlap scheduling:

    • max serve --disable-overlap keeps the overlap pipeline but synchronizes after every model execution, so device graph capture can run without CPU/GPU overlap. MAX_DISABLE_OVERLAP=1 sets its default.
    • max serve --no-enable-overlap-scheduler (or enable_overlap_scheduler=False) now turns the overlap scheduler off without --force. Device graph capture, which needs the overlap scheduler, then stays off unless you pass --device-graph-capture. Leaving the flag unset still lets the architecture turn the overlap scheduler on.

  • The max CLI now allocates through jemalloc on Linux instead of glibc malloc. Set MODULAR_MAX_ALLOCATOR=system to keep glibc malloc; sanitizer builds and non-Linux platforms are unaffected. Scripts that use InferenceSession directly are not affected either; to run one on jemalloc, add the installed modular/lib/libjemalloc_preload.so to LD_PRELOAD before starting Python.

  • Added prefill-admission controls:

    • Added --prefill-coalesce-min-pending (default 0, off): under in-flight batching, hold pending fresh prefills until that many can share one mixed step instead of admitting them one by one. With data parallelism the count is the total across all replicas, not one replica's queue. Mixed steps forfeit device graph capture for their decode rows, so coalescing admissions keeps more decode steps on the captured fast path, trading a bounded prefill-admission delay for lower decode latency at high concurrency. That delay is at most the same number of decode steps, or --prefill-coalesce-max-held-steps when set. Mid-prefill chunked continuations are never held.
    • Added --prefill-coalesce-max-held-steps (default 0): a cap on how many decode steps in a row a held prefill waits before it is admitted, no matter how few are queued. Without it, --prefill-coalesce-min-pending does both jobs: it sets the queue depth that releases a prefill, and the number of steps after which one is released anyway. So the queue depth could not be raised without also making prefills wait longer. This flag splits the two. At 0 it falls back to --prefill-coalesce-min-pending, and it does nothing while that is 0.
    • Added --max-request-input-tokens (default 0, off): a ceiling on how many prefill tokens one request may draw from the batch's context-encoding budget in a single step, applied by chunking. Without it a long prefill can claim the whole --max-batch-input-tokens budget step after step while short requests queue behind it, so the cap trades the long request's time to first token for the interactivity of the short ones it no longer blocks. It requires chunked prefill and is inert without it, and --chunked-prefill-min-chunk-size must not exceed it.
    • Added --prefill-schedule-interval (default 1, every step): admit prefill work only on every Nth scheduler step, leaving the steps in between entirely to decode. Data-parallel ranks advance in lockstep, so prefill on any one rank stalls the whole group; scattering it across steps pays that stall repeatedly, while concentrating it onto a shared cadence pays it once. A step with no decode work on any replica admits prefill regardless, rather than run an empty batch. The cost is delayed prefill admission, bounded at N-1 steps. Unlike prefill coalescing, mid-prefill chunked continuations are held too: the budget emits one chunk per step, so exempting them defeats the cadence.

  • Added adaptive speculative decoding:

    • Added --adaptive-speculative-widths to pick how many drafts speculative decoding verifies by measured decode tokens per second, instead of a fixed width. Takes a list such as 1,3,5, or all. Each width adds graph capture time at startup.
    • Added --adaptive-speculative-min-batch-size to capture narrower adaptive verify widths only from that decode batch size up. Smaller batches verify the widest width, which cuts startup time and graph memory.

  • Expanded max benchmark image mixing:

    • Added dataset-agnostic image mixing to max benchmark (--image-fraction, --image-count, --image-long-side, --image-aspect-ratio, --image-turn), so any dataset (sonnet, sharegpt, instruct-coder, and so on) can have generated images mixed into a fraction of its requests or chat turns, not just the random dataset's existing fixed --random-image-count/--random-image-size. --dry-run reports an image-count distribution table (and a decoded image-long-side-pixel table) when a workload has images. See the image mixing quick-start guide.
    • max benchmark's image mixing now draws its selection from a private seeded RNG rather than the global one, so adding another augmentation ahead of it no longer changes which requests or sessions carry images. It also logs, over the turns it sampled, the newly-encoded images per request and the share of turns carrying one, the two shares a multimodal workload is calibrated against. augment_samples_with_images's image_fraction and image_turn arguments are renamed fraction and turn to match augment_samples_with_response_format; the --image-fraction and --image-turn flags are unchanged.
    • max benchmark's --image-fraction now picks exactly that share of requests or chat sessions, rounding up or down at random when the share isn't whole, instead of drawing each one independently. A run picks no images only when that share is under one item or too few items can carry them.
    • Added --image-num-turns to max benchmark, which cuts chat sessions that get images to a drawn number of turns, so image-bearing requests can carry as many resent image parts as production's. The image-mixing log now also reports image parts per image-bearing request.

  • Expanded max benchmark constrained requests and tool calls:

    • max benchmark's --response-format now reaches multi-turn workloads, which previously logged a warning and dropped it. A constrained request or turn runs without ignore_eos: a schema-shaped response ends where its schema is satisfied, and generating past that point makes the server drop enforcement for the remainder of the request. A chat session's running prompt length now charges what each turn actually generated rather than the length it drew.
    • max benchmark's --response-format now applies to a share of traffic rather than all of it: --response-format-fraction (default 1.0) sets the fraction of requests, or of eligible user turns, that are constrained, and --response-format-turn (every, first, last) chooses which turns of a chat session are eligible. The fraction is drawn per request and per turn, so at the default every it lands directly on the share of requests that set response_format; first and last narrow eligibility to one turn per session, so the realized request share is correspondingly lower.
    • Added --tools and --tools-fraction to max benchmark, which send OpenAI tool definitions on any workload, multi-turn chat sessions included. Previously only datasets that carry their own tools sent any, and only on single-turn requests. --tools takes a tools list (inline JSON or a file path); --tools-fraction (default 1.0) sets the fraction of requests, or of chat sessions, that carry it, and a selected session sends its tools on every turn. The rendered definitions are carved out of a session's first turn or a single-turn string prompt, so the input-length distribution still describes the whole prompt. Results report tool_request_rate and tool_call_response_rate: the share of requests that offered tools, and the share of those whose response called one.

  • max benchmark now reports its realized request mix in its own "Request Mix" section and result_groups.request_mix, rather than among the headline metrics in result_groups.summary. The group holds the structured-output and tool-calling rates, plus two new ones: image_request_rate, the share of requests whose payload carried an image (including images resent with a session's earlier turns), and lora_request_rate, the share routed to a LoRA adapter. Existing top-level keys in the result JSON are unchanged.

  • Added --workload-stats-json to max benchmark, which writes the statistics of the requests a workload would send as JSON. With --dry-run it needs no server, only the tokenizer. Input length is split into first and later turns and reports the share of requests above 128k, 192k and 256k tokens. The tool-calling, image and structured-output shares are included, as are any turns a session drops at the context limit. Warm-up prefix turns, which are never sent, are left out. --print-workload-stats and --dry-run print the same figures in a new "Requests Sent" section.

  • Added a Cat(v1:w1, v2:w2, ...) categorical distribution for every max benchmark config field that accepts a distribution string (for example --image-long-side, --image-count, --random-input-len), so an explicit, empirically-measured distribution (such as image sizes measured from real production traffic) can be reproduced exactly instead of approximated with a parametric shape like N/U/LN. The :weight suffix is optional per entry (uniform when omitted), and weights don't need to sum to 1.

  • max benchmark no longer strips a leading @ from a file path, whether it is passed to --response-format, --tools, --extra-body or --agentic-tool-profiles or named in a --workload-config YAML. Pass the bare path instead: --tools path/to/tools.json.

Python API​

  • Removed the deprecated max.kv_cache module. It only re-exported max.pipelines.kv_cache and emitted a DeprecationWarning on import. Import PagedKVCacheManager, DummyKVCache, load_kv_manager, and the other names it re-exported from max.pipelines.kv_cache instead.

  • Changes to driver buffers:

    • Removed the DeviceQueue overloads of max.driver.Buffer.to() and Buffer.copy(). Pass a Device (or, for to(), a sequence of devices) instead: buffer.to(device) and buffer.copy(device=device).
    • Removed max.driver.DevicePinnedBuffer. Allocate Buffer(dtype, shape, device=gpu, usage=Usage.STAGING | Usage.UNTRACKED) instead; UNTRACKED keeps the old behavior, where reads never wait on device copies. Replace isinstance(b, DevicePinnedBuffer) with b.pinned to check for page-locked memory, or Usage.UNTRACKED in b.usage to check whether reads skip the hazard wait.

  • Expanded custom op support:

    • Added reusable custom op declarations, so one input and output signature can drive the same operation in eager code and graph construction, using max.experimental.custom.declare(), which returns an immutable custom.CustomOp. Eager calls reuse a cached binding across shapes of the same rank, avoiding shape-specific recompilation. Output dims may be symbolic, parameterized, or left to the kernel's shape function to determine at run time. declare() raises ValueError if the op's custom_extensions or the process-wide defaults don't register the kernel.
    • ops.custom() and ops.inplace_custom() now accept float values in their parameters dictionaries, lowering to a rank-0 f32 attribute that a kernel declares as a Float32 parameter.

  • Added module trace and compile:

    • Added max.nn.Module.trace() and max.nn.Module.compile(), which trace a module's __call__ into a graph and compile it in a given InferenceSession, matching the methods of the same names on max.experimental.nn.Module. Pass module= to trace() to compile several modules' graphs into one artifact.
    • max.experimental.nn.Module.compile(), max.experimental.compilation.compile(), and CompiledCallable take a session to compile in, so a compile uses that session's precompiled or exported MEFs and kernel options rather than the process-wide session's. Module.trace() and Module.compile() also take a graph name.

  • Stabilized pytrees in max.tree:

    • Removed max.experimental.nn.CompiledModel. Module.compile() now returns a max.experimental.compilation.CompiledCallable, which also accepts a Buffer for a single-device input.
    • Removed the Module parameter helpers apply_to_parameters, map_parameters, and load_state. Use max.tree instead.
    • Removed the SubgraphInput alias. Use the public Tree type alias in max.tree for nested pytree values; layer subgraph inputs now use Tree[Any].
    • Moved pytree utilities out of experimental to stable max.tree. max.experimental.nn.Module is now a pytree: its attributes are its children, so max.tree functions such as tree.map() and tree.flatten() walk a module directly.
    • Added the max.tree.dataclass decorator, which registers a dataclass as a pytree node so max.tree functions such as tree.map() and tree.flatten() walk its fields. Fields that are None are dropped on flatten and restored on unflatten.

  • Expanded mixture-of-experts layers:

    • max.nn.moe.MoE and MoEQuantized take the expert activation as one activation= argument, a GatedActivation from max.nn.activation, in place of use_swigluoai, swiglu_alpha, swiglu_limit, and gated_activation_fn. Replace use_swigluoai=True, swiglu_alpha=a, swiglu_limit=l with activation=GatedActivation("swigluoai", alpha=a, limit=l), and gated_activation_fn=fn with activation=GatedActivation(gated_activation_fn=fn). The fused SwiGLU kernel wrappers max.nn.kernels.grouped_matmul_blocked_swiglu and grouped_dynamic_block_scaled_matmul_amd_swiglu_quant, and max.nn.comm.ep.ep_kernels.fused_silu and fused_silu_quantized, take the same activation= argument.
    • max.nn.moe.StackedMoE can now run MXFP4 experts W4A8 on SM100 GPUs, including under tensor parallelism: pass mxfp8_activations=True with an MXFP4 quant_config. The expert scales are then declared in the grouped matmul's interleaved layout (see max.nn.moe.interleaved_block_scales_shape). The new router_dtype and combine_dtype options run the router and the weighted combine in float32.
    • Added max.nn.moe.SigmoidTopKRouter, the sigmoid top-k router with an expert score correction bias (noaux_tc with one expert group) that MiniMax-M2, HY-V3, and MiMo-V2 share.
    • max.nn.moe.MoE takes interleaved_gate_up=True for gate/up projections stored as [gate_0, up_0, gate_1, up_1, ...], and runs its fused activation on them; max.nn.comm.ep.ep_kernels.fused_silu takes interleaved=True to read that layout. Under expert parallelism, an interleaved MoEQuantized layer needs the fused NVFP4 or MXFP8 grouped-matmul SwiGLU path, since the fused SiLU quantize kernels read only [gate | up].

  • Expanded quantization support:

    • max.experimental.nn.Linear takes a quant_config, like max.nn.Linear. It supports static per-tensor FP8: a float8_e4m3fn weight with a weight_scale and an input_scale, the layout modelopt FP8 checkpoints use.
    • Added the QuantConfig.blockscaled_fp8(), QuantConfig.mxfp4(), and QuantConfig.nvfp4() constructors in max.nn.quant_config, which build the scale specs of block-scaled FP8 (N x N weight blocks, dynamic 1 x N activation blocks), MXFP4 (E8M0 per 32, dynamic activations), and modelopt NVFP4 checkpoints.
    • Added max.pipelines.weights.read_modelopt_quantized_layers(), which reads the per-module quantized_layers map of a modelopt MIXED_PRECISION checkpoint and refuses any module, algorithm, or group size the caller did not declare it can load.

  • Expanded GPU execution for graph and eager ops. The following ops now execute on GPU instead of copying tensor data to the CPU and back:

    • ops.cumsum: Float32 inputs accumulate in float32 rather than float64. max.experimental.functional.cumsum also has its own sharding rule, cumsum_rule.
    • ops.scatter, ops.scatter_add, ops.scatter_max, ops.scatter_min, and ops.scatter_mul: These ops now work under strict_device_placement=DevicePlacementPolicy.Error. As on CPU, the reductions apply duplicate indices atomically in an unspecified order, and ops.scatter keeps an unspecified winner among duplicates. On Apple silicon GPUs, the reductions support only 32-bit dtypes.
    • ops.tile: This op now works under strict_device_placement=DevicePlacementPolicy.Error and accepts tensors of any rank on every device, up from a maximum of 4 dimensions. Ranks beyond 5 compile on first use rather than being precompiled. Repeats remain host metadata, and a data-dependent output size still prevents device-graph capture.
    • mo.non_maximum_suppression: On one B200, compute-kernel time is 60x faster than the host path at 1024 boxes (0.7 ms versus 42 ms for 1 batch and 1 class) and 152x faster at 4096 boxes (1.3 ms versus 203 ms).
    • mo.gather_sum: The fused gather on axis 0 and sum on axis 1 is now single-device. The gather-reduce fusion that produces this op in compiled graphs is also restored. On an M-series GPU, the fused op runs 24-37x faster than the previous round-trip for a batch of 1024, 64 multi-hot IDs, embedding dimensions of 128 or 512, and float32 data.
    • mo.roi_align / ops.roi_align: ROI align is now single-device with host-side scalar operands and preserves the CPU implementation's results.
    • ops.repeat_interleave: Scalar and per-element repeats now run on the input's device. Per-element repeats build their output-to-input offset map on the device, so no repeat value is read on the host. FLUX.2 and Z-Image RoPE embeddings now use this op instead of stack-and-reshape workarounds.
    • mo.arg_nonzero / ops.nonzero: The GPU implementation preserves stable row-major ordering and is now single-device. Because the output row count is data-dependent, only that count is still computed on the host.

  • The // operator on TensorValue and max.experimental.tensor.Tensor now matches ops.floor_div: floor division of two integer tensors returns an integer tensor of the operand dtype instead of float64. Float operands are unaffected. Code that relied on the float result should cast explicitly, for example (x // n).cast(DType.float64).

  • Improvements for in-process text-generation:

    • Removed max.pipelines.GenerateMixin and the generate() and generate_async() methods it gave text generation pipelines. Build contexts with pipeline.tokenizer.new_context() and run them with GenerateIterator.from_contexts(), or pass token ids to max.pipelines.lib.generate.generate(). TextGenerationPipelineInterface now declares pipeline_config and tokenizer as abstract properties, so custom subclasses must implement them.
    • Added AutoModel and AutoTokenizer to max.pipelines for generating with a model in-process. model.generate(**tokenizer(prompts), max_new_tokens=16) runs exactly the batch you pass, takes the same sampling arguments as max generate, and can return each token's logits and log probabilities. model.generate_iter() returns a GenerateIterator from max.pipelines.lib.generate that runs one forward pass per step. Asking for a feature that only serving supports raises an error with the matching max serve command.
    • Added AutoPipeline to max.pipelines, which uses AutoModel and AutoTokenizer to take text and return text. Each output starts with its prompt, and the prompt runs without the chat template.

  • Reworked experimental sharding:

    • Removed max.experimental.sharding.PlacementMapping and DeviceMapping.to_placements(). Instead use DeviceMapping and its placements attribute, which they aliased.
    • Removed DeviceMesh.default(), DeviceMesh.is_single, DeviceMesh.is_simulated and DeviceMapping.to_mesh(). Instead use DeviceMesh.single(CPU()), mesh.num_devices == 1 and a new DeviceMapping on the target mesh.
    • Removed max.experimental.sharding.mesh_context(). Pass a DeviceMesh to max.experimental.tensor.default_device() instead. defaults() now returns the device as a DeviceMesh.
    • Removed max.experimental.nn.common_layers.functional_kernels.local_map(). Instead wrap the per-device function in F.functional() and claim its output's placement with Tensor.rebind_mapping().
    • Removed max.experimental.functional.per_shard_dispatch(). Use max.experimental.functional.call_on_mesh() instead to run a function on each device of a mesh, or on each group of devices along given mesh axes, and return its outputs as distributed tensors. Use it for code that a sharding rule cannot express, such as a collective kernel. to_tensors(), map_tensors(), tensor_to_layout() and any_distributed() are also no longer exported from max.experimental.functional; use Tensor.layout for a tensor's layout.
    • max.experimental.sharding no longer re-exports P, R, Action, PerShard, PerShardDim, Collective, ReduceOp, get_active_mesh, as_device_mapping, as_layout or the *_rule functions.
    • Removed ActionSet, its extras and finalize hooks, build_action_set, and force_replicated_action_set. Sharding rules now return a plain list of AxisAssignment rows, and replicated_rows and match_operand_placement build the common cases.
    • Removed ReduceOp. max.experimental.sharding.Partial now always means a pending sum because no other reduction could be resolved.
    • Removed mode(), isolated_solver(), Solver, ReshardBehavior, GreedyReshard, NoReshard and PartialsOnly; for example, mode(NoReshard()) becomes auto_reshard(mode="raise"). Use max.experimental.sharding.auto_reshard to control automatic resharding. It now raises by default instead of silently inserting collectives when an op can't accept its inputs' placements; use auto_reshard(mode="silent") or mode="warn" to allow it.
    • Added max.experimental.sharding.visualize, which draws a DeviceMesh, TensorLayout or Tensor twice over: as a grid of colored blocks sized in proportion to the slice each device holds, and as a tree with one level per mesh axis. A slice that is a partial sum, or whose relation to the others is Unknown, is marked and drawn see-through. Pass blocks=False or tree=False for one of the two, or a dict of values to draw them side by side. It renders as SVG in a notebook and as Unicode text anywhere else.
    • max.experimental.random.uniform() and gaussian() now accept a DeviceMapping.
    • max.experimental.nn.Module.to() now only moves a module to one device. Build a multi-device module inside default_device(mesh) instead.
    • Added max.experimental.sharding.Unknown, the placement for per-device values with no known relation, and Tensor.rebind_mapping(), which relabels a tensor's placements without moving data. Ops on Unknown inputs run on each device's own shard and return Unknown results.
    • max.experimental.functional.allreduce_sum(), allgather() and reduce_scatter() accept several mesh axes, by index or name, and run one collective over all the devices those axes span. A tensor axis sharded over several mesh axes is now split once over all those devices, the way a reduce-scatter over them splits it: 10 rows on a 2x2 mesh become 3, 3, 2 and 2 rows instead of 3, 2, 3 and 2.
    • max.experimental.functional.stack() takes a device mapping that places the stack of host tensors directly, so each device receives only its part.
    • A sharding rule in max.experimental.sharding now receives the op's own arguments with each tensor replaced by its layout, at any depth, and returns a plain list of AxisAssignment rows: one placement per tensor operand and, in AxisAssignment.outputs (formerly output), one per result. Rules no longer rewrite non-tensor arguments.
    • reshape, broadcast_to and rebind in max.experimental.functional take global shapes on a distributed tensor, as on one device. A sharded axis stays sharded on the first axis it splits into, or on the axis it merges into when it's the first of the merged axes, and is gathered otherwise. split takes global sizes: along a sharded axis, each device splits its own piece by its share of every size when every size divides evenly, and the axis is gathered otherwise.

  • OverlapTextGenerationPipeline has a new graph_capture_ready property.

  • Added max.nn.Module.named_modules(), which yields every module in the tree with the prefix its weights carry in raw_state_dict(), so code can find modules by type and address their weights without parsing names.

  • Added max.profiler.oneshot.cuda_profiler_region(), a context manager that brackets a region with cudaProfilerStart/cudaProfilerStop so nsys/ncu capture only the wrapped region.

  • Added an eager usage validator for ModuleV3 and eager Tensor code. It can be enabled with a new --eager-usage-validator flag in the max CLI.

  • Added max.nn.HyperConnection and max.experimental.nn.HyperConnection, the ModuleV2 and ModuleV3 forms of a manifold-constrained hyper-connection (mHC) site. The layer generalizes the residual connection: hc_mult residual streams run in parallel, and a learned gate collapses them into the sublayer's input and decides how the sublayer's output is written back across them. It owns the hc_fn, hc_base, and hc_scale weights, and returns (post, comb, collapsed) so the caller drives the residual update. Float32, GPU-only.

Kernels and GPU programming​

  • Deprecated LayoutTensor and LayoutTensorIter in the layout package. Use TileTensor and TileTensorIter instead.

  • Removed the LayoutTensor overloads of sum() and max() from layout.math and of create_tensor_tile() from layout.tma_async. Use the TileTensor overloads instead.

  • Removed max.nn.kernels.mhc_split_sinkhorn. Use max.nn.kernels.hyper_connection_gates instead. The fused kernel computes the Manifold-Constrained Hyper-Connections (mHC) gates between the stream projection and stream collapse. It replaces the sigmoids, softmax, and Sinkhorn-Knopp projection with one launch that keeps the whole hc_mult x hc_mult mixer in one warp's registers, instead of roughly 40 elementwise and reduction launches at 20 Sinkhorn iterations. It is float32 and GPU-only.

  • Changed capturing closures to runtime arguments:

    • Removed the extensibility.foreach[func, target=...](output, ctx) overloads that took a capturing closure as a compile-time parameter. Pass the body as a closure in the first runtime argument instead: foreach[target=target](body, output, ctx).

    • max.gpu.primitives.warp.reduce() and lane_group_reduce() now take the reduction function as a runtime closure argument instead of a compile-time capturing parameter. The shuffle function stays a compile-time parameter:

      def add[dtype: DType, width: SIMDLength](
    x: SIMD[dtype, width], y: SIMD[dtype, width]
) -> SIMD[dtype, width]:
    return x + y

warp.reduce[shuffle_down](val, add)  # was: warp.reduce[shuffle_down, add](val)

  • Replaced index lists with coordinates:

    • TMATensorTile, TMATensorTileArray, TMATensorTileIm2col, and SplitLastDimTMATensorTile now take their shapes as flat Coord values instead of IndexList[rank], and no longer take a separate rank parameter.

    • The reduction callbacks in the max.algorithm package now take a Coord instead of an IndexList[rank], and no longer carry a rank parameter. This changes the input_fn and output_fn contracts of sum(), min(), max(), mean() and product(), along with row_reduce() and reduce_launch() in the GPU backend. The coordinate holds the same values as before, so results are unchanged:

      # was: def input_fn[width: Int, rank: Int](idx: IndexList[rank]) -> ...
def input_fn[width: Int](idx: Coord) -> SIMD[dtype, width]:
    return input.load[width=width](idx)

  • Expanded TileTensor:

    • TensorEngine no longer requires load() and store(), and DefaultEngine, DevicePointerEngine, and StaticOffsetEngine no longer implement them. Loads and stores go through the raw pointer that unsafe_ptr() returns: Engine.load[width=w, alignment=a](storage, offset) becomes Engine.unsafe_ptr(storage).load[width=w, alignment=a](offset), and TileTensor.unsafe_ptr() exposes the same pointer at the tensor level.
    • Deprecated TileTensor.as_immut() in favor of TileTensor.as_imm(), which returns the same immutable view and matches the naming of Pointer.as_imm().
    • TileTensor's two runtime slicing methods are replaced by subscript syntax, and the All marker they used is removed along with them. t.slice(batch, All, All) becomes t[batch, :, :] and t.slice((0, h), (0, w)) becomes t[0:h, 0:w]. All was a CoordLike that stood for a dimension rather than a coordinate, which is why Coord carried a contains_slices member to reject it; both are gone. Where the indices are compile-time values, prefer slice[](), which keeps the view fully static.
    • TileTensor gained as_span(), returning a Span over the tensor's elements.
    • TileTensor gained unsafe_ptr(), returning the raw scalar Pointer its engine exposes for the tensor's storage. Element loads and stores on the tensor go through that pointer.
    • TileTensor.slice now accepts an Int in place of a slice literal, which fixes that dimension to the given index and drops it from the result, the way squeeze would. The output rank is the number of slice arguments, so t.slice[1:3, 2, 0:4]() returns a rank-2 view of a rank-3 tensor. Any axis may be dropped, not only trailing ones, and each surviving axis keeps its own stride. Bounds stay compile-time and are now range-checked against the parent's static shape at compile time. A view of a fully static tensor remains fully static: its extents and strides are ComptimeInt, and the base offset is carried in the view's engine as a ComptimeInt rather than computed at runtime.
    • TileTensor now slices through subscript syntax: t[batch, 0:n, :] fixes the batch axis, narrows the next one to a runtime subrange, and keeps the last whole. Every dimension takes exactly one argument. A subscript is a view whenever one of its arguments is a slice, and an element load otherwise, so existing indexing is unchanged. Compile-time and runtime indices share the path: an Idx[n] index on an axis with a compile-time stride is folded into a ComptimeInt component of the view's offset, and only the rest is computed at runtime. Slice bounds are runtime values, so a sliced extent is runtime too -- : means 0:dim, not a marker. Strides are always inherited whole.
    • Added TileTensorIter in layout.tile_iterator, the TileTensor counterpart of LayoutTensorIter. It steps a fixed tile layout through memory, dereferences to a TileTensor view of the current tile, supports circular multi-stage shared-memory buffers, and can carry a runtime valid-row count for masked global-memory tiles.

  • Changes to custom-kernel tensor arguments:

    • Added TensorTuple to the extensibility package, so a custom kernel can take a variadic output tensor argument and a variadic input tensor argument together. Each is a TensorTuple[*Ts] over an infer-only TypeList bound by tensor-argument traits, and its elements may differ in dtype and shape. A signature takes at most one tuple per operand list. The tuple holds its elements by value and, when every element is DevicePassable, crosses a kernel launch as one argument. Kernel import now rejects a variadic pack of tensor-trait arguments (*xs: *Ts) with a diagnostic that points at TensorTuple.
    • Changed shape functions registered with @extensibility.register_shape_function() to require tensor-trait arguments. Use Some[Tensor] for shapes, Some[TileTensorable] for element values, or TensorTuple for variadic inputs. Legacy tensor aliases remain supported on execute().

  • The rms_norm functions and GPU kernels in nn.normalization, the KV cache rms_norm ops, and gemv_and_partial_norm now take epsilon and weight_offset as Float32 compile-time parameters instead of runtime arguments.

  • Expanded DeviceContext:

    • DeviceContext is now an alias for DeviceContextBase[], a new struct parameterized on the CompilationTarget it runs on, which defaults to the default accelerator. Code that uses DeviceContext doesn't change, but its API reference is now the DeviceContextBase page.
    • Added DeviceContext.CPU, a DeviceContext type whose target is the host CPU, along with the DeviceContext.cpu() factory and DeviceContext.assert_cpu_target(). parallelize(), sync_parallelize(), and parallelism_level() in max.algorithm now take an Optional[DeviceContext.CPU] instead of an Optional[DeviceContext].
    • DeviceContext.stream() now works on Apple GPUs, where each stream maps to a new Metal command queue, so stream-based event recording (record_event(), enqueue_wait_for()) behaves the same as on CUDA and HIP.
    • Added DeviceContext.wrap_host_memory(): makes a caller-owned host range device-accessible for as long as the returned DeviceBuffer lives. It grants access, not ownership, and the range must be addressed through that buffer rather than through the pointer passed in. CUDA, HIP and Metal only; Metal also requires a page-aligned base and a page-multiple length.

  • Added GPU resize:

    • Added GPU support for ops.resize_linear / mo.resize.linear in graphs and eager mode. The op was host-only; resize_linear_gpu now runs one pass per resized axis through device scratch buffers with the same per-point interpolation as the CPU kernel. Interpolation accumulates in float32 (float64 for float64), so bfloat16 and float16 inputs no longer round at every tap. The eager resize family sweeps accelerators for every resize op.
    • Added GPU support for ops.resize_nearest / mo.resize.nearest. The op was host-only; its kernel is a pure gather, so it now launches on the input's device (the op is single-device with a host size operand, like bicubic). The eager resize family sweeps accelerators for all three modes; new Metal-verified GPU kernel tests cover up/downsampling and both coordinate modes, plus an eager GPU-residency test.
    • Trimmed the eager resize family's accelerator warm to rank 4. The GPU warm action (warm_accel_resize) swept all five ranks for resize.nearest/resize.linear across every dtype and variant, which alone made the action compile 363 graphs per virtual accelerator and put it among the slowest warm fragments. The warm now builds only the image-shaped rank-4 graphs per accelerator (75 per slot, a 4.8x drop in compile count and artifact size); CPU keeps the full 1..5 sweep. A rank outside the warm is not unsupported — it lazy-compiles on first dispatch and is cached, now covered by a GPU test.
    • mo.resize.nearest skips the coordinate arithmetic when every axis scales by an exact integer of at least 1 and the modes reduce to floor division (asymmetric + floor, half_pixel + half_down, half_pixel + half_up), gathering from out_coord // scale (a shift when the scale is a power of two) instead. On the benchmark's vision-encoder shape (float32, 224x224 -> 448x448) this is 1.66x faster than the generic path (22.9us vs 38.2us median). The VAE and diffusion 2x-upsample helpers in max.pipelines.architectures (autoencoders, flux, qwen_image and their ModuleV3 ports) now delegate to ops.resize_nearest instead of hand-rolling reshape + broadcast / concat expansions; output is bit-identical at 2x.

  • Added Blackwell tensor memory:

    • Added layout.tmem_engine.TMemEngine, a TensorEngine that lets a TileTensor view Blackwell Tensor Memory (TMEM). A TMEM tile's layout places elements on the 128-lane by 512-column grid lane-first: a lane stride of 1 and a column stride of TMEM_NUM_LANES (128), so the whole accumulator is (128, 512):(1, 128) and nested layouts express the lane placement of other MMA shapes. The engine encodes grid positions into the hardware's lane-and-column addresses itself. TMEM has no pointer, so the engine's data path is TileTensor.copy_from in either direction: a copy between a TMEM row and a register or shared-memory tile moves consecutive columns of the lane the calling thread owns through the warp-collective tcgen05.ld or tcgen05.st in the 32x32b shape, one instruction per power-of-two chunk of at most 64 registers, and one wait per 64-column slice. The copy_from_async and copy_to_async engine methods and the tile-level tmem_copy_async issue the same instructions without waiting, so several copies can share one TMemEngine.wait_store or wait_load; an async row is capped at 64 columns, so a wider row is tiled into one copy per slice. A thread's view of a warp's (32, N) tile is its row 0, since the hardware adds the lane to the warp base address. The engine supports 4-byte element types and requires an SM100 target.
    • Added TensorEngine.copy_to, the source side of a copy. TileTensor.copy_from now calls it on the source tensor's engine, and its default forwards to the destination engine's copy_from, so existing engines are unaffected. An engine whose storage has no pointer, such as TMemEngine, overrides it to run its own load loop.

  • Improved KDA and gated-DeltaNet kernels:

    • Reworked the gated_delta_conv1d_fwd GPU kernel (Gated DeltaNet causal conv, Qwen3.5/3.8 and GLM-5/Kimi-Delta prefill pass 1) from one thread per channel walking the sequence serially to one thread per token tile. A 2048-token prefill chunk at the Qwen3.8 TP2 conv width drops from ~650 us to ~32 us per layer, and 8192 tokens from ~4.3 ms to ~112 us, bit-identically, with decode shapes unchanged.
    • The KDA / gated-DeltaNet chunk prefill now routes the production configuration (fp32 output, bf16 K_FIRST state pool, original softplus gate, probability beta, GVA head grouping on SM100) to the two-kernel bt16 pair instead of the fused kernel: a prep pass materializes gate-rebased tiles per chunk group, then a chain kernel walks chunks per (sequence, value-head) with tcgen05 MMA and drains fp32 output directly. At the Qwen3.5 prefill shape (16 key / 48 value heads, 128 dims, 2048 tokens per forward) the op drops from 0.59 ms to 0.23 ms per layer, and the fused kernel remains the fallback for other configurations.
    • Added max.nn.state_space.kda_chunk, the chunk-blocked form of the KDA / gated-DeltaNet recurrence, alongside the existing kda_decode. Its launcher can now reach the fused Blackwell kernel. The fused path is opt-in through a compile-time parameter and off by default. At Qwen3.5 geometry the launcher runs 1024 tokens in 0.315 ms against the three-stage pipeline's 1.853 ms.

  • Added max.nn.kernels.keyed_uniform, which draws one uniform value in [0, 1) per row of a seed tensor. Every row is its own Philox key, so a row's value is a function of its seed alone, not of the row's position or of what else shares the launch. ops.random.uniform keys the whole tensor off index 0 of the graph seed and walks the flat element index as its Philox counter, so it cannot express that. Speculative decoding's accept coin uses it to draw one uniform per (request, draft position).

  • max.nn.kernels.grouped_matmul_block_scaled and max.nn.kernels.grouped_matmul_blocked_swiglu accept an optional a_row_scales: one bfloat16 scale per activation row, multiplied into that row's output together with its expert scale (before the SwiGLU in the fused op). It dequantizes NVFP4 activations that were quantized with a per-token tensor scale. NVFP4 on SM100 only.

  • Upgraded the bundled libnvptxcompiler from CUDA 13.3 to CUDA 13.4 Update 1 (13.4.92), bringing the latest bug fixes and performance improvements from NVIDIA's PTX compiler.

Tools​

  • Added new debug options:
    • Added the pre-jit debug option (MODULAR_DEBUG=pre-jit, [max-debug] pre-jit, or InferenceSession.debug.pre_jit). It stops graph compilation once the Mojo for the graph has been emitted into ir-output-dir, skipping the kernel JIT, and InferenceSession.compile() raises max.engine.CompilationStopped in place of returning a model. max warm-cache treats that as success, so MODULAR_DEBUG=ir-output-dir=<dir>,pre-jit max warm-cache ... is a quick way to get the Mojo a model compiles to.
    • Added the uninitialized-read-mode debug option (MODULAR_DEBUG=uninitialized-read-mode=report, [max-debug] uninitialized-read-mode, or InferenceSession.debug.uninitialized_read_mode). With uninitialized-read-check on, report prints each poisoned read and keeps running, so one run lists every offending load instead of stopping at the first. The default, abort, is unchanged. On Apple GPUs, report currently aborts without printing.

Fixes​

  • Device graph capture no longer aborts the model worker on AMD GPUs.

  • Kimi, MiniMax-M2, GLM-4.7, and Gemma 4 no longer fail a request whose strict: true tool schema has unenforceable keywords; they compile it best-effort, unless the tool-call policy selects different behavior.

  • Fixed the paged KV cache overflowing its page budget when memory estimation caps max_length to what fits. The budget now reserves the null block's page, and a budget too small for that page plus one request fails startup with a clear error.

  • Fixed the tiered KV cache connector leaving its max_kv_tiered_* disk offload directory behind after shutdown. Startup now reclaims an auto-created directory that no live process still holds, including one abandoned by SIGTERM, SIGKILL, or an OOM kill. Directories left by earlier versions are still only reported, not deleted.

  • Fixed a decode-engine crash in disaggregated serving when KV allocation failed for a request that partly hit the prefix cache. A failed allocation now leaves the request unchanged, in both the paged and the Jenga cache managers.

  • Fixed a prefill worker crash on response_format requests under disaggregated serving. Prefill no longer enforces the grammar; the decode worker samples the first token, so the schema is still honored.

  • Fixed the scheduler refusing to admit queued prefill on models with a recurrent state cache while KV cache memory was still free.

  • Fixed indexing a buffer, such as a paged KV cache gather, copying the entire source on every execution. The gather now reads only the requested rows.

  • Fixed eager attention with a paged KV cache through the functional kernel wrappers in max.experimental.nn.common_layers.

  • Fixed KV cache transfers over NIXL on InfiniBand hosts aborting the process when a NIXL agent shut down.

  • Fixed ops.floor_div() on mixed integer dtypes such as uint8 // int16. 7 // -2 now returns -4.

  • Fixed tracing and telemetry export:

    • Fixed MAX sending the headers in OTEL_EXPORTER_OTLP_HEADERS, or its signal-specific forms, to Modular's telemetry collector.
    • Fixed max serve with MAX_SERVE_KERNEL_TRACE_LEVEL=kernel never writing its kernel trace. The trace is written when the server stops, if the model worker shuts down within its 5 second grace period.
    • Fixed max.phase.prefill and max.phase.decode spans not being exported for a completed request.

  • Fixed max.experimental.functional.relu() so it runs eagerly on float Tensor inputs.

  • Fixed a crash serving Gemma 4, Gemma 3 multimodal, Pixtral, and Qwen3 embedding with a batch larger than one.

  • Fixed max generate decode on vision-language models:

    • Fixed max generate crashing after the first token for Gemma 4 and Idefics3 models, whose tokenizers rejected the CLI's token list on decode.
    • Fixed the same max generate decode crash for the other vision-language models: Gemma 3 multimodal, InternVL, Kimi K2.5, Pixtral, Qwen2.5-VL and Qwen3-VL.

  • Fixed streaming /v1/completions sending an unreadable frame, or ending the response abruptly, when a request failed part-way through a stream. The error now arrives as an error object the OpenAI client surfaces as an APIError, matching /v1/chat/completions.

  • Fixed sampling reproducibility:

    • Fixed sampling_params.seed not reproducing across requests. Requests that pin the same seed and have generated the same number of tokens now draw the same tokens.
    • Fixed sampled tokens changing with the batch slot a request occupied, including after preemption. Samples now depend only on the request's seed and how many tokens it has generated. The distribution is unchanged, but the exact token for a given seed can differ from a previous build.
    • Fixed identical requests producing different tokens when they occupied different batch slots on multiple GPUs. Reduce-scatter now sums peers in a fixed order. Outputs of existing multi-GPU batches may shift within rounding noise.

  • Fixed a negative-bound slice of a concatenated value returning the value rotated instead of the requested region, as in concat(x[..., :-rope_dim], x[..., -rope_dim:]). Slice bounds are now applied against the axis length.

  • Fixed tool-call argument schemas:

    • Fixed tool schemas that use const or enum with an object or array value failing to compile on models that frame arguments in XML (GLM, MiniMax, Qwen, DeepSeek). The literal now keeps its JSON form. Tag-keyed formats reject it with an explicit message.
    • Fixed GLM and MiniMax-M2 tool calls with free-form JSON arguments running to max_tokens on /v1/chat/completions. A string nested in such a value is now a quoted JSON string, so the call can finish.
    • Fixed a GLM tool-call argument pinned by const or enum to a string that reads as a number or a boolean ("2.0", "true") being returned as that number or boolean. The value stays a string.

  • Fixed structured-output integer fields rejecting numbers with an exponent, such as 1e80. The grammar now allows an exponent of + or no sign, and still rejects a negative exponent such as 1e-5. This applies to response_format schemas and tool calls.

  • Fixed Qwen3.5 returning fluent nonsense on every prompt, including text. Each linear-attention layer now keeps its own recurrent state.

  • Fixed top_k on Apple GPUs returning stale values past the 32nd output when k is large. Ties in rows of up to 2048 elements now come back smallest index first, as on other GPUs.

  • Fixed PyTorch raising HIP error: peer access is already enabled on AMD GPUs after MAX was set up across several GPUs more than once, such as a second multi-GPU InferenceSession.

  • Fixed DeviceContext.execution_time() and execution_time_iter() on Apple GPUs reporting host enqueue time. They now return GPU execution time, as on NVIDIA and AMD GPUs.

  • Fixed grouped collectives failing to compile when device groups had different shapes. ops.allreduce.sum(), ops.reducescatter.sum(), ops.allgather_rms_norm() (including its MXFP8 and MXFP6 variants), and ops.reduce_scatter_rms_norm() now accept a different row count per group. The fused RMSNorm ops reject groups whose shapes differ outside axis 0.

  • Fixed weight loading treating nested checkpoint-<step>/ directories as extra shards, which could load an arbitrary training step. Those directories are excluded unless requested through subfolder.

Mojo language​

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

Models

Cookbook

Puzzles

Blog

Community