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.5

Highlights​

  • The MAX AI kernels library is now the MAX accelerator library, and most of the GPU programming APIs moved out of the Mojo standard library into a new top-level max package: std.algorithm to max.algorithm, std.benchmark to max.benchmark, and std.gpu.{compute,host,memory,sync} to max.gpu.*. The Mojo layout package also now ships with the MAX accelerator library.

  • Support for Apple silicon GPUs now extends back to M1 devices. This release also adds hardware-MMA flash-attention prefill for M5 that lowers TTFT by ~1.4x at sequence length 8192. Image generation is also faster on M5: FLUX.2-klein bf16 checkpoints now default to int8 W8A8 quantization for a ~1.45x end-to-end speedup at near-lossless quality, and NVFP4 checkpoints can opt into an int8 requant that runs ~2.56x faster than the default W4A16 path on FLUX.2-dev.

  • Expanded skills for the model bring-up workflow in our skills repo, including serve-model to bring up an OpenAI-compatible endpoint with max serve, benchmark-model to measure throughput and latency, eval-model to run accuracy evaluations against a served endpoint, and profile-model to profile inference. You can install them in Claude Code from a plugin marketplace with /plugin marketplace add modular/skills, taking either the full skills plugin or the max-skills and mojo-skills subsets. See AI coding assistants for setup.

Documentation​

  • Renamed the MAX AI kernels library to the MAX accelerator library, now including the GPU programming APIs under the new top-level max package; its API reference moved to /api/mojo/ (old links redirect).
  • Added docstrings across the public APIs of the kernel packages (linalg, comm, kv_cache, shmem, state_space, pipeline, structured_kernels, profiling_range, builtin_kernels, and the nvml wrapper), now published in the kernels API reference.
  • Added instructions for installing Modular's agent skills through the Claude Code plugin marketplace on the AI coding assistants page.
  • Brought the OpenAI-compatible REST API reference up to date: each operation is now its own page under a new rest-api/index.mdx overview, and new chat-completion fields and message roles are documented.
  • Restructured the MAX docs site: removed the /max URL prefix and moved deploy/ and inference/ under serve/ and graph/ under develop/.
  • Moved cloud deployment documentation out of the MAX doc set (redirects added).

MAX models​

  • Added GLM-5.2 (GlmMoeDsaForCausalLM) support, extending the GLM-5.1 sparse-attention architecture with cross-layer index sharing.
    • Added multi-token prediction (MTP) speculative decoding for GLM-5.2 (UnifiedMTPGlm5_2ForCausalLM), serving the baked-in NextN layer as a single-layer sparse-MLA draft; enabled automatically for GLM checkpoints that ship a NextN layer with --speculative-method mtp.
    • Added tool-calling, reasoning, and structured-output (response_format) support to GLM-5.1 / GLM-5.2, enabled with --tool-parser glm45 --reasoning-parser glm45 --enable-structured-output.
    • Fixed a GLM-5.1-FP8 crash caused by a shared-experts dtype mismatch.
  • Added Laguna (LagunaForCausalLM) support for poolside/Laguna-M.1-NVFP4, including tool calling.
  • Added DiffusionGemma (DiffusionGemmaForBlockDiffusion) support for google/diffusiongemma-26B-A4B-it (bfloat16) and nvidia/diffusiongemma-26B-A4B-it-NVFP4; text-only for now.
  • Added Nemotron-H (NemotronHForCausalLM) support, NVIDIA's hybrid Mamba-2 + attention decoder, with modelopt per-tensor FP8 and a new Mamba-2 SSD chunked-scan varlen kernel.
    • Extended Nemotron-H with the Nemotron-3-Nano-30B-A3B hybrid MoE variant and enabled the architecture on Apple silicon GPUs in bfloat16.
    • Enabled NVIDIA's official FP8 Nemotron-H checkpoints on Apple silicon (previously crashing or producing all-zero logits) and sped up Nemotron-H decode on Apple M5 by ~41-81%.
    • Added support for serving nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8 on Apple silicon via a tiled simdgroup-MMA grouped-FP8 (W8A16) MoE matmul, decoding faster than bf16 at concurrency with half the weight memory.
  • Fixed the max_batch_size handling for Nemotron-H.
  • Added support for the detail parameter on image and video content parts in chat requests.
  • Added Ideogram 4 (Ideogram4Pipeline) support, a text-to-image flow-matching diffusion transformer; serve via /v1/responses.
    • FP8 checkpoint weights run hot projections on native FP8 GEMMs (~24% faster end-to-end on MI355).
  • Added support for amd/Kimi-K2.7-Code-MXFP4 on AMD GPUs.
  • Expanded Gemma 4 support:
    • Added DSpark speculative decoding for Gemma 4 12B (UnifiedDSparkGemma4ForCausalLM), DeepSeek's block-drafting method: a small draft transformer drafts a 7-token block per step. Enabled with --draft-model-path deepseek-ai/dspark_gemma4_12b_block7 --speculative-method dflash --num-speculative-tokens 7.
    • Sped up Gemma4-12B DSpark decode by up to ~1.3x via a packed wide-N shallow-K GEMV, a single-pass streaming argmax kernel, and device graph capture.
    • Gemma 4 with MTP speculative decoding (UnifiedMTPGemma4ForCausalLM) now supports image and video input; previously the vision encoder output never reached the language model, so image prompts were answered as if the model were blind.
    • MTP speculative decoding now samples recovered tokens from the residual distribution when stochastic acceptance rejects a draft token, preserving the target distribution for argmax draft proposals.
    • Added structured-output and tool-calling support via the xgrammar backend, covering Gemma 4's special tool-call format.
    • Added float16 support, with the logit softcap and vision pooler run in fp32.
    • Added tensor-parallel support for the MoE variant.
    • Video inputs now route through the shared VisionEncoderCache, so a repeated clip is served from cache with no re-encode.
    • Video decoding now runs on a worker thread, so concurrent requests overlap video decode.
    • Improved vision-batch serving latency by concatenating embeddings on-device instead of round-tripping through host numpy.
    • Fixed the MoE expert-router softmax being computed in bfloat16 instead of float32, which degraded MoE quality.
    • Fixed image/video position and scatter indexing desyncs under chunked prefill, which could corrupt vision embeddings on multimodal prompts split across chunks.
    • Fixed crashes in multi-device serving and multi-image batches by making merge_per_device_buffers rank-agnostic.
    • Fixed reasoning being dropped after tool results.
    • Fixed a vision-batch crash caused by constructing a Device() instead of CPU() for host tensors.
  • Expanded DeepSeek-V3 ModuleV3 support:
    • Added NVFP4 (modelopt) weight support, running experts, dense MLPs, and the attention output projection on SM100 block-scaled FP4 matmul kernels.
    • Added data-parallel + expert-parallel (DP-EP) and multi-GPU tensor-parallel + expert-parallel (TP+EP) serving. Note: Tensor.to no longer implicitly calls F.distributed_broadcast; call it explicitly where needed.
    • Fixed the FP8 adapter by casting f32 normalization gammas, resolving a dtype mismatch.
  • Expanded Kimi K2.5 support:
    • Kimi with DFlash speculative decoding (UnifiedDflashKimiK25ForCausalLM) now supports image input; previously the vision encoder was not compiled, so image prompts were answered as if the model were blind.
    • Added support for combining Kimi tool calling with response_format=json_schema on the xgrammar constrained-decoding backend.
  • Expanded FLUX.2 support:
    • FLUX.2-klein bf16 checkpoints on Apple M5 GPUs now default to int8 W8A8 quantization, ~1.45x faster end-to-end than bf16 on FLUX.2-klein-4B at near-lossless quality; set APPLE_FLUX2_INT8_W8A8=0 to opt out.
    • NVFP4 checkpoints can now opt into an int8 W8A8 requant at load on Apple M5 with APPLE_FLUX2_INT8_W8A8=1, ~2.56x faster end-to-end than the default W4A16 path on FLUX.2-dev.
    • Diffusion pipelines now support two denoising-cache backends to skip redundant transformer passes: --taylorseer (recommended default, with balanced and fast presets) and --first-block-caching; the two are mutually exclusive and both off by default.
  • Expanded Qwen support:
    • Added tool-calling and reasoning support to Qwen 3.5 / 3.6.
    • Fixed a Qwen3EmbeddingModel crash.
  • Added per-request LoRA adapter support: LoRALinear and StackedLinearLoRA extend LoRA to standalone and fused-QKV projections, with LoRAManager.apply swapping target layers in a model.
  • Improved Eagle3 speculative-decoding performance by removing a redundant concatenate in the draft path.
  • Fixed Step-3.5-Flash accuracy and performance.
  • Fixed the EAGLE3 MHA draft lm_head all-gather in pure tensor-parallel mode.

MAX framework​

  • The graph compiler now fuses query/key RMSNorm followed by rotate-half RoPE into a single rms_norm_rope GPU kernel even when the RMSNorm upcasts to float32; numerics match the unfused graph.
  • Added a poison-all mode to MODULAR_DEBUG_DEVICE_ALLOCATOR that fills every memory-manager allocation with a configurable NaN-pattern byte (MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_POISON_PATTERN), so uninitialized device-memory reads trip differential tests without kernel instrumentation. Manual debugging aid, not a default.
  • Added conda packages max-benchmark, max-serve, and max-all, plus a max[all] wheel extra, for parity with the existing wheel extras.
  • Multimodal pipelines now compile their vision and language models in parallel via a shared Module container and session.load_all(), cutting compile/load time by up to 1.86x (Qwen3-VL-4B: 614s -> 428s).
  • Made the compiled-model (MEF) cache key relocatable across install paths: absolute-path-valued pipeline options no longer enter the key, so a cache warmed under one install path hits under another.
  • ModuleV3 weights are now sharded and transferred to devices inside the compiled graph rather than via eager ops, reducing per-GPU memory use (about 10 GiB for a DP-EP NVFP4 DeepSeek-V3).
  • The VMM defragmenting allocator is now the default memory manager on NVIDIA GPUs, fixing external-fragmentation OOMs ("plenty free but no contiguous block"); override with MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_VMM=0. Also fixed the earlier opt-in being a silent no-op.
  • Added a HIP-based VMM defragmenting allocator for AMD GPUs (opt-in via MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_VMM=1) on MI300-series hardware.
  • Coalesced consecutive Metal kernel launches into a single shared command buffer with a tunable op cap, reducing per-launch overhead on Apple GPUs; also restored Metal GPU execution aborted by an unimplemented driver-context stub.
  • Improved expert-parallel MoE execution by running the shared expert on a side stream via ops.side_stream, overlapping it with the routed-expert computation.
  • Allowed float16/bfloat16 graphs to load float32 checkpoint weights, with the weight adapter casting at load time.
  • Improved multi-device startup latency by batching replay preface copies into a single submission.

Inference server​

  • Structured output (constrained decoding) and tool calling:
    • Reduced per-iteration latency for structured-output requests: the grammar bitmask is now unpacked and applied on the GPU in a single fused kernel (apply_packed_bitmask), and on speculative-decode models the overlap pipeline enqueues the FSM-advance and bitmask compute in the consuming batch's row order, removing a host synchronization point and a device-side gather across all six speculative-decode architectures.
    • Fixed a structured-output runaway where a response_format schema omitting the root "type" compiled to an unbounded grammar and generated until max_length; object-implying schemas are now normalized to "type": "object" before grammar compilation.
    • Structured-output grammars are now compile-validated at request admission: uncompilable or unsatisfiable schemas return HTTP 400 instead of crashing the model worker, and boolean-valued or non-object root schemas are now accepted per JSON Schema draft-7.
    • Structured-output (response_format) JSON is now emitted compactly, without whitespace around separators, reducing output token count and mitigating runaway generation on Gemma 4.
    • Kimi tool-call arguments are now grammar-enforced via xgrammar structural tags during constrained decoding.
    • Added a pluggable grammar backend for constrained decoding, allowing alternative grammar engines to be selected.
    • Fixed intermittently dropped Kimi K2.5 tool calls under reasoning-enabled tool_choice="auto": the reasoning parser now also ends a reasoning span at <|tool_calls_section_begin|> when the model omits </think>, so the tool call reaches the tool parser instead of being stranded in reasoning.
    • Raised the maximum tool function name length from 64 to 1024 characters; longer client-supplied tool names are now accepted instead of rejected with a 400 error.
  • Extended the chat completions endpoint: accepts the root message role (opt-in via tokenizer extra_chat_roles), a thinking control object, dotted tool function names, and per-image and per-request size limits; reconciles max_tokens with max_completion_tokens; added a schema-aware coerce_arguments tool-parser hook; and added per-video fps, max_frames, and max_long_side_pixel sampling hints on video_url content parts.
  • Speculative decoding:
    • Speculative decoding now runs to max_seq_len instead of early-stopping, allowing requests to reach their full configured length.
    • Greedy speculative-decode acceptance now uses an exact argmax comparison, improving acceptance correctness for greedy requests.
    • Enabled the overlap scheduler for DFlash speculative decoding, reducing per-iteration latency.
  • KV cache and prefix caching:
    • Added opt-in SHA-256 KV-cache block hashing: a new kv_cache_hash_algo field on KVCacheConfig (default ahash64; opt-in sha256 and sha256_64), an optional per-request cache_salt scoping prefix-cache reuse, and support across the host-tier connectors (NullConnector, LocalConnector, TieredConnector) and the dKV external tier (DKVConnector), which advertise accepted algorithms via a supported_hash_algos property validated at startup. Default behavior is unchanged.
    • MAX now refreshes an external KV cache tier's recency when a request's prefix is served from the on-GPU cache, keeping hot shared prefixes warm in an external tier such as dKV; set MODULAR_DKV_DISABLE_G0_TOUCH=1 to disable.
    • Added a Rust-implemented host/disk KV cache connector, selected with --kv-connector rust_tiered (CUDA-only), with asynchronous onload that overlaps compute.
    • Sped up host KV-cache allocation for tiered/local KV caching with a fast page-locked allocator (mmap + parallel page-fault + chunked host registration).
    • Added a per_layer_buffers KV cache option backing each layer (including quantized scale buffers) with its own sub-cap buffer, unlocking larger context lengths by relaxing the per-allocation size cap.
    • Added an enable_dp_cross_replica_prefix_copy KV cache config flag (default on); disabling it serves cross-replica prefix-cache reuse only from the shared host/disk tier instead of device-to-device copies.
    • Added int8 KV-cache storage with fp16 block scales (MXint8) for flash attention, halving KV-cache memory footprint at near-parity performance and accuracy with fp16.
    • Fixed the KV-cache transfer engine (NIXL, behind disaggregated serving) being NVIDIA-only: packages now ship per-vendor UCX plugin flavors, so KV-cache transfer works on AMD GPUs.
    • The dKV connector now supports tensor parallelism (TP > 1) on the multi-tenant path for head-sharded (MHA/GQA), MLA, and GQA head-replicated models, with each GPU handshaking its own per-shard store. Also fixed a cross-shard race where a KV load returning differing block counts could leave a stray in-flight copy that later clobbered a reallocated block.
    • The dKV KV-cache connector (--kv-connector dkv) now supports data-parallel serving and shares its prefix cache across DP replicas, matching the local and tiered connectors.
    • Data-parallel serving now shares the prefix cache across replicas: GPU hits are served by device-to-device copy, and the CPU/disk offload tiers are a single shared pool, so host_kvcache_swap_space_gb now sizes one shared host pool for the whole deployment rather than one per replica.
  • Improved time-to-first-token for multimodal requests: oversized media is now rejected before its bytes are fully materialized, large data: base64 decoding runs on a worker thread instead of blocking the event loop, per-video count and byte limits are enforced up front, and multi-MB data: payloads bypass Pydantic AnyUrl parsing via a lightweight DataUrl wrapper.
  • Made numpy array transport across the API-server-to-model-worker request queue zero-copy: large arrays (notably vision pixel_values) ride out-of-band as their own ZMQ frame and are decoded as views, replacing both the copy and shared-memory transports on this path.
  • Fixed image requests failing with a 400 or 500 across all vision models: data: URIs with unpadded or URL-safe base64 are now decoded tolerantly, and truncated, animated, or content-negotiated images are validated with a full pixel decode that fails fast as a clean 400.
  • Changed VLM image (and Gemma 4 video) prefix-cache/routing keys to hash the raw encoded container bytes plus a resolution size_tier, making keys stable across torch/BLAS/numpy versions and CPU-vs-GPU.
  • Added an opt-in emit_reasoning_content server config that emits a reasoning model's chain-of-thought under reasoning_content instead of reasoning; off by default.
  • Added support for the OpenRouter-style reasoning request field in chat completions.
  • Chat completions now respect max_completion_tokens as the preferred way to cap output length, per the OpenAI API spec; max_tokens is still accepted for backwards compatibility.
  • Added MAX_SERVE_GRACEFUL_SHUTDOWN_TIMEOUT_S (default 5s) controlling how long the server drains in-flight requests after SIGTERM, and fixed SIGTERM cutting off in-flight streaming responses immediately — streaming requests are now bounded by the same timeout.
  • Added request-queue backpressure to MAX serve: --max-queue-size rejects new requests with HTTP 429 once the queue to the model worker is full, and --max-pending-requests stops the worker draining the queue once its prefill queue is that deep. Both default to unbounded; rejections are observable via maxserve.request_count with code="429".
  • Failed request submissions in the OpenAI-compatible endpoints now surface as HTTP 4xx/5xx errors instead of a 200 OK stream carrying an error payload; mid-stream errors are still serialized as error events.
  • Image generation responses on the Open Responses endpoint now report usage: output_tokens/total_tokens carry the total pixel count of the generated images and input_tokens is 0. Previously usage was always null.
  • Added an opt-in MAX_SERVE_STREAM_MIN_CHUNK_TOKENS server config (default 1 = unchanged) that coalesces streamed SSE output into larger chunks without affecting TTFT; streaming now also truncates output at a stop string, and reasoning and content are emitted as separate chunks so a delta at the boundary no longer carries both fields.
  • Added HTTP endpoints to fetch and reset expert-parallel load balancing (EPLB) statistics, plus the MAX_SERVE_EPLB_PROFILE environment variable to enable EPLB stats profiling in max serve.

Server metrics​

  • Fixed speculative-decoding metrics being accumulated across all batches since server start; they are now per-batch.
  • Added Prometheus counters for structured-output and tool-call failures: maxserve.tool_call.conformance_errors and maxserve.structured_output.grammar_rejections.
  • Added data-parallel load-balance metrics: dp_active_token_occupancy and dp_context_token_occupancy.
  • Added block-level KV-cache metrics: maxserve.cache.device_blocks_served, maxserve.cache.cross_replica_blocks_copied, and maxserve.cache.cross_replica_bytes_copied.
  • Added MAX_SERVE_OTLP_METRICS_ENDPOINT to push a self-calibrating exponential-histogram shadow (<metric>.exponential) of every histogram metric to an OTLP endpoint, alongside the existing Prometheus /metrics histograms. Unset by default; behavior unchanged.
  • Added video encoder statistics to the scheduler's per-iteration batch log (a Video Encoder clause mirroring Vision Encoder) and as OpenTelemetry metrics under maxserve.video.*. Models managing their own encoder cache implement a SupportsPooledVisionMetrics protocol so their image and video metrics reach the scheduler.
  • Fixed the code label on the maxserve_request_count metric to report the HTTP status actually returned, so requests rejected before generation are counted with their real status code instead of 200 or dropped.
  • Added dKV external-tier health metrics: OpenTelemetry gauges maxserve.dkv.connected_clients, maxserve.dkv.total_clients, and maxserve.dkv.reconnect_attempts, plus a dKV degraded scheduler log clause while any replica client is disconnected.
  • Added maxserve.cache.disk_blocks_read and maxserve.cache.disk_blocks_written counters for KV blocks read from and written to the disk cache tier when tiered KV caching is enabled.
  • Added vision encoder statistics to the scheduler's per-iteration batch log for multimodal models (images encoded vs. served from the vision encoder cache, patches and vision tokens), also exported as OpenTelemetry metrics under the maxserve.vision.* namespace.
  • Fixed maxserve_time_to_first_token_milliseconds under-reporting TTFT by excluding request parsing and media resolution; expect a step up in reported TTFT that reflects measurement, not a regression.
  • Fixed inflated maxserve.cache.h2d_blocks_copied and maxserve.cache.d2h_blocks_copied telemetry on tiered and local KV cache deployments; counters now report per-batch deltas.
  • Fixed the input-token metric being double-counted in MAX Serve.
  • Fixed max-serve telemetry emitting a batch's token counts alongside the previous batch's timing under overlap scheduling.
  • Fixed speculative-decode acceptance metrics leaking from a token-generation batch into a following context-encoding batch.

max CLI​

  • Improved error messages when a subcommand's components are not installed: e.g. max serve from the base package now names the extra to install (max[serve]) instead of a bare ModuleNotFoundError.
  • Trimmed max benchmark dependencies so it no longer requires the full serve dependency set; telemetry is now disabled for max benchmark and max list.
  • Added opt-in token-balanced CE scheduling across data-parallel replicas via --dp-ce-balance-timeout-ms (default off): new context-encoding requests are priced at their post-prefix-cache length and bound to the least-loaded replica, deferring up to the timeout while per-step occupancy stays under --dp-ce-balance-threshold (default 0.8). Dynamic chunk-size reduction (--dp-ce-balance-enable-dynamic-chunk-size, default on) runs below-threshold steps immediately at a reduced chunk size.
  • Added --chunked-prefill-min-chunk-size (default 0 = off) to floor the token size of any chunked-prefill chunk, avoiding degenerate sliver chunks that pay a full step's overhead for almost no progress.
  • Added an experimental --fold-sampler-into-graph option (default off) that folds greedy argmax token selection into the captured forward graph, removing the per-token blocking sampler submission on supported architectures (currently Nemotron-H). Also fixed the serve CLI silently dropping this and related flags on their way to the model worker, which had disabled the folded sampler; with the flags threaded through, decode latency improves substantially.
  • Added a max-pending-futures config (default 1, unchanged behavior); setting it to 2 enables experimental schedule-ahead decoding with two forwards in flight per request.
  • Added --no-enable-tool-call-constrained-decode (default enabled) to keep --tool-parser parsing of generated text while skipping grammar-based constrained decoding for tool calls. With it disabled, tool_choice=required can no longer force a tool call; --enable-structured-output is unaffected.
  • Added a --top-k MAX Serve flag to set a server-wide default top-k when clients omit top_k; request-supplied values still take precedence.
  • Added a max warm-interpreter-cache command that batch-compiles the eager interpreter model matrix into the on-disk cache (parallel workers, --jobs, --check, --force) so later programs start warm. max serve now refuses on-demand eager compiles by default (opt back in with --allow-cold-interpreter-cache or MAX_EAGER_ALLOW_LAZY_COMPILE=1); adopting a warm cache still works.
  • Fixed KV-cache CLI flags (e.g. --kv-connector-config) replacing the entire kv_cache section of a --config-file recipe instead of overriding just the named setting.
  • Fixed speculative-decoding CLI flags (e.g. --num-speculative-tokens) being silently dropped when a --config-file recipe already set a speculative section; flags now merge onto the recipe's values.
  • Fixed default-valued CLI overrides being silently dropped; explicitly passing a flag equal to its default now overrides recipe/config-file settings.

Python API​

  • Added max.graph.ops.reduce_scatter_rms_norm, a distributed op that reduce-scatters a bfloat16 tensor across devices and RMSNorm-normalizes each device's row shard in one collective launch, returning both the normed shard and the residual sum shard; numerically identical to reduce-scatter followed by rms_norm. DeepSeek V3.2 and GLM-5.x now use it in TP-EP non-allreduce mode.
  • Added max.graph.ops.allgather_rms_norm, which fuses an all-gather with the following RMSNorm into one kernel launch, returning both the normed tensor and the bit-identical gathered residual. Dispatches to the fused kernel or falls back by shape; not enabled by default in any shipping pipeline.
  • Added an optional init_value argument to max.graph.ops.buffer_create: the buffer becomes persistent state, allocated and filled once at model load and reused (with mutations preserved) across executions.
  • Added max.graph.ops.floor_div (and F.floor_div), element-wise floor division matching Python //. Integer operands stay in the integer domain, so integer floor division compiles on backends without 64-bit float support (for example, Metal GPUs).
  • Added max.driver.set_virtual_cpu_target() / get_virtual_cpu_target() to fix the CPU codegen target (for example "x86-64-v3" or "generic") before importing max._interpreter_ops, so the eager interpreter's CPU kernel cache is host-independent and shippable across hosts of the same architecture family. Mirrors set_virtual_device_target_arch() for GPUs.
  • Eager execution in max.experimental now routes every realization through the max.experimental.executor.Executor abstraction, selected by MAX_EAGER_EXECUTOR (default composite, the same interpreter-then-cached-compile behavior as before).
  • The eager interpreter now compiles its graph-compiler models lazily, per target on first dispatch, instead of JIT-compiling the full kernel library at import; set MAX_EAGER_OP_PRECOMPILE=1 to precompile at import.
  • The eager interpreter's remaining op families — binary elementwise and comparison, reduce/softmax/argmax/cumsum, shape-rearrange (pad, tile, split, concat, slice), gather/scatter, pooling, conv2d, resize, top_k/bottom_k, where, layer_norm/rms_norm, cast, band_part/arg_nonzero/non_maximum_suppression/random/range/ roi_align, and the last Mojo-bound ops (transpose, broadcast_to, store_slice) — now run through pre-compiled graph-compiler models instead of hand-written Mojo bindings. No Mojo sources ship, removing the multi-minute first-import JIT compile. Structural parameters stay runtime operands, so one compiled graph per (op, device, dtype) serves every shape.
  • Small eager-interpreter behavior fixes from the migration: band_part gained CPU float64 on Apple silicon; arg_nonzero no longer caps input rank; range raises on a non-evenly-divisible interval or zero step; random.uniform samples may differ between eager and compiled graphs from the same seed.
  • Added max.graph.Dim.substitute, which replaces named symbols in a dimension with other dims or static values, folding the result through the compiler's own attribute evaluation. Unmapped symbols are left intact, and a static dimension is always returned unchanged.
  • Added max.graph.KernelLibrary.has_shape_function, reporting whether a kernel in a loaded library registered a shape function.
  • Added max.experimental.nn.subgraphable for Module subgraph compilation: a repeated block lowers to one shared subgraph reused per call (share across modules by passing a name), cutting graph size and compile time — the DeepSeek-V3 ModuleV3 MoE decoder block now compiles as one reused subgraph. Opt out with Module.compile(..., allow_subgraphs=False).
  • max.nn.hooks.PrintHook now supports max.experimental.nn.Module.
  • Added F.print, supporting both single-device and multi-device tensors.
  • Added max.graph.default_custom_extensions() and the default_custom_extensions_scope() context manager: registered paths are merged into the custom_extensions of every new Graph, so a backend's custom-op library is reachable without an explicit custom_extensions=.
  • Added PipelineArgs (max.pipelines) as the user-facing input type for constructing pipelines, consolidating construction arguments into a single explicit type instead of scattered flat kwargs.
  • Added max.experimental.compile_pool.ProcessCompilePool, which compiles graphs in parallel background worker processes; eager execution now compiles out of process, so a first eager call on a warm MEF cache drops from ~5.6s to ~16.7ms.
  • Added InferenceSession.compile_async, which schedules compilation on the runtime worker pool and returns a pending CompiledModel; InferenceSession.compile now wraps it and blocks.
  • Added max.engine.read for loading a compiled-model artifact (.mef) without an InferenceSession; the result can be initialized on any session via InferenceSession.init.
  • Added a LoRA wrapper to max.experimental.nn.common_layers: wrap a Linear or QKVLinear in place to add a ragged multi-adapter delta fed as per-call graph inputs (GPU-only, Mojo SGMV); the QKV path uses a single fused LoRA-B expand op, simplifying weight loading and the StackedLinearLoRA / LoRAMixin interfaces.
  • Added a tile_based_fusion option (default False) to compile / load in the graph API, enabling tile-based codegen at the graph-API level.
  • Added a fused_qk_rms_norm_rope_ragged Python wrapper, exposing the fused QK RMSNorm + RoPE ragged paged kernel to graph building.
  • Added a model_name property to max.driver.Device.
  • Made several previously private modules public so their APIs are officially documented: the GGUFWeights loader wrappers in max.graph.weights and the max.profiler.cpu / max.profiler.gpu modules.
  • Added max.experimental.testing with a shared assert_all_close helper for eager max.experimental.Tensor tests.
  • Added a PipelineModel.sampler_custom_extensions hook so custom extensions are threaded into the token-sampler graph.

C API​

  • Fixed M_borrowTensorInto() copying instead of borrowing a GPU input: device pointers already on the target accelerator are now borrowed in place (zero-copy) on CUDA, so in-place mutation of a BufferType model input is visible to the caller. Host pointers are still staged via copy, as are device pointers on AMD and Apple.

GPU programming​

Most GPU programming APIs formerly packaged with Mojo have been rehomed into the MAX accelerator library, including most of the std.gpu package (now max.gpu), TileTensor, and LayoutTensor. Some low-level APIs remain in the standard library.

Device context and execution​

  • DeviceGraphBuilder.add_function() now covers every live DeviceContext.enqueue_function() form, so any kernel launchable on a device context can also be recorded as a graph node:

    • Added an overload that takes the kernel as a compile-time parameter and compiles it automatically, so callers no longer need a separate DeviceContext.compile_function() step:

      def build(mut builder: DeviceGraphBuilder) raises {read}:
    _ = builder.add_function[kernel](
        42, grid_dim=1, block_dim=1, dependencies=[]
    )

    • Added overloads accepting a DeviceExternalFunction loaded from PTX/SASS via DeviceContext.load_function(), and a capturing kernel as a compile-time parameter with runtime arguments.

    • All add_function() overloads now accept a location argument so wrappers can attribute launch errors to their callers, and the closure overload now accepts (and honors) a func_attribute argument.

  • DeviceGraphBuilder.collect_dependencies() now accepts an optional dependencies argument. The named predecessor handles become ambient predecessors of every node the work closure adds, so the scope's nodes run after those predecessors without the closure threading the handles through to each add_* call. With the default (empty) dependencies the behavior is unchanged. When work adds no nodes, the returned join node falls back to depending on dependencies so it still chains correctly.

    var producers = builder.collect_dependencies(add_producers)
# Every node added by `add_consumers` depends on `producers`:
var consumers = builder.collect_dependencies(
    add_consumers, dependencies=[producers]
)

  • DeviceContext.load_function() now keys its runtime cache on the requested entry-point name as well as the blob. Loading two different entry points (for example, kernel_a and kernel_b) from a single PTX/cubin blob no longer collides—previously the second load silently returned the function resolved by the first. The cache also no longer keys on the entire blob when no module name is supplied: it keys on a short hash of the blob instead, so each call avoids copying, hashing, and byte-comparing the whole blob (and retaining a duplicate of it). The win scales with blob size and matters most for large multi-entry blobs loaded on the per-execution path.

  • The DeviceStream type is now included in the API reference documentation. Returned by DeviceContext.create_stream() and DeviceContext.create_external_stream(), it provides methods for synchronizing and sequencing asynchronous GPU work (for example, synchronize(), record_event(), and enqueue_wait_for()). The type was already public but was previously hidden from the generated docs.

Kernel arguments and address spaces​

  • Int and UInt no longer conform to DevicePassable and can no longer be passed as arguments to GPU kernels (via DeviceContext.enqueue_function() or compile_function()). They are platform-sized index types whose bit width depends on the host, so passing them to an accelerator miscompiles when the host and device disagree on the width (for example, a 64-bit host driving a 32-bit GPU index domain). Use a fixed-width type—Int32, Int64, UInt32, or UInt64—for kernel scalar arguments and parameters, and convert back with Int(...) inside the kernel body if you need a platform Int there. A kernel that still takes a bare Int/UInt argument now fails to compile with: "Int and UInt are not passable to device kernels; use a fixed-width type such as Int32 or Int64 instead".

  • The GPU device-side standard library now uses the unified safe Pointer type throughout std.gpu (memory, compute, intrinsics, sync, and primitives). Public signatures that previously took or returned UnsafePointer are respelled to bare Pointer; since Pointer and UnsafePointer share representation and origin and decay implicitly, this is a type-identical change for callers. One visible difference: external_memory() now returns a safe Pointer instead of an UnsafePointer. Code that performs raw pointer arithmetic on the result can use the unsafe_offset() method.

  • AddressSpace is now target-extensible rather than a fixed, portable enum. The built-in GPU spaces (GENERIC, GLOBAL, SHARED, CONSTANT, LOCAL, SHARED_CLUSTER, BUFFER_RESOURCE) are unchanged, but accessing any other name—for example, an accelerator-specific AddressSpace.SCRATCHPAD—now resolves through the active hardware backend instead of being a hard-coded compile error. The set of valid address-space names is the union of the built-in GPU spaces and whatever the active backend defines, so accelerator backends can provide their own named spaces (with their own values) only where they exist. A name that no backend defines remains a compile-time error.

  • Coord now conforms to DevicePassable, so a Coord embedded in a DevicePassable type (such as a TileTensor's Layout) is encoded to the device through Coord._to_device_type() instead of a raw field bit-copy, the same way IndexList already was.

  • coord is now a comptime expression, and coord[DType]() has been renamed to dyn_coord[DType](). You can now write var my_coord = coord[1, 2, 3] to create a Coord[ComptimeInt[1], ComptimeInt[2], ComptimeInt[3]].

Warp and matrix primitives​

  • Atomic.compare_exchange() now accepts a weak parameter, and requires weak=True to compile on Apple GPU targets: AIR exposes no strong compare-exchange primitive, so Metal only lowers the weak form. This is safe for the common case of a CAS-retry loop, since a spurious failure just costs one extra iteration. Previously any use of compare_exchange(), including helpers built on it like atomic scatter-reduce, failed to compile on Metal.

  • Added warp.match_any() and warp.match_all(): match_any() returns, for each warp lane, the mask of lanes whose value has the same bits, and match_all() returns the warp's active-lane mask if every lane holds the same bits and 0 otherwise. They use NVIDIA's match.any.sync and match.all.sync instructions, a readfirstlane ballot fold on AMD, and a shuffle-based emulation on Apple Silicon GPUs.

  • warp.vote() now works on Apple Silicon GPUs. Metal's AIR backend exposes no usable ballot intrinsic, so it emulates the ballot with an XOR-butterfly OR-reduction over simd_shuffle_xor, returning a 32-bit mask (or a DType.uint64 mask whose upper 32 bits are always zero); NVIDIA and AMD are unchanged.

  • Added an 8x8 simdgroup_matrix matrix multiply-accumulate primitive (_mma_apple_8x8()) with apple_mma_load_8x8() / apple_mma_store_8x8() fragment helpers for Apple Silicon GPUs in std.gpu.compute.arch. Unlike the 16x16 path (Apple M5 only), the 8x8 primitive is available on all Apple GPU generations (M1-M5). It accepts Float16, BFloat16, and Float32 inputs with a Float32 accumulator.

  • Apple M5 simdgroup_matrix MMA now accepts FP8 (float8_e4m3fn, float8_e5m2) inputs with an F32 accumulator, alongside the existing F16/BF16/F32 and 8-bit integer types.

Device targeting and hardware support​

  • GPUInfo.vendor has been removed. It duplicated GPUInfo.api, which identifies the vendor precisely ("cuda", "hip", "metal", or a stdlib plugin's own API name) rather than collapsing every plugin accelerator into one enum value. Compare api instead:

    comptime use_apple_path = ctx.default_device_info.api == "metal"

    Vendor itself remains, as the classifier behind has_amd_gpu_accelerator(), has_nvidia_gpu_accelerator(), and has_apple_gpu_accelerator().

  • A bare --target-accelerator architecture (for example, gfx950 or sm_90) is now handled identically to its vendor-prefixed form (amdgpu:gfx950, nvidia:sm_90). Previously has_amd_gpu_accelerator(), has_nvidia_gpu_accelerator(), and has_apple_gpu_accelerator() only recognized the vendor-prefixed spelling, so code that specialized on them (such as warp-tiling parameters) could silently take the wrong path and fail a downstream comptime constraint. amd:<arch> is also now accepted as an alias for amdgpu:<arch>, mirroring the existing nvidia:<arch> prefix.

  • The GPU Vendor type can now be imported from std.sys (from std.sys import Vendor). It remains importable from std.gpu.host.info for backward compatibility.

  • Added support for the Steam Deck's RDNA2 Van Gogh APU.

Removed​

  • Removed the deprecated DeviceContext.compile_function_experimental() and DeviceContext.enqueue_function_experimental() methods, along with overloads that passed the kernel twice. Use DeviceContext.compile_function[func]() and DeviceContext.enqueue_function[func]() instead.

  • Removed the UInt-returning GPU indexing accessors (thread_idx_uint, block_idx_uint, block_dim_uint, grid_dim_uint, global_idx_uint, lane_id_uint, warp_id_uint). Use the Int-returning thread_idx, block_idx, block_dim, grid_dim, global_idx, lane_id, and warp_id accessors instead.

  • Removed the store_volatile() and load_volatile() intrinsics from std.gpu.intrinsics. Use Pointer.unsafe_store[volatile=True]() and Pointer.unsafe_load[volatile=True]() instead, which work across all supported GPU targets rather than NVIDIA only.

  • Removed the deprecated GPUAddressSpace alias for AddressSpace. Use AddressSpace directly.

MAX kernels​

  • GPU sampling and token selection:
    • Sped up GPU token sampling with top_k >= 10 by 2-4x by fusing softmax, temperature scaling, and min-p masking into the rejection-sampling kernel.
    • Improved the GPU top-k kernel: it now streams past the previous N=2048 cap, with a new single-pass block-wide bitonic sort (persistent_topk_block).
    • Made the top-k + top-p sampler's RNG offset batch-invariant, so sampling results no longer depend on batch composition.
    • Fixed GPU sampling kernel bugs: unwritten output when a per-row top_k is 0, division-by-zero in apply_gumbel_noise_kernel at temperature=0, and NaN-handling bugs in the topk kernel.
  • Sped up GPU softmax with a new split-K path for few-row, wide-column shapes, significantly improving BF16 1024x4096 and 256x4096 cases on both B200 and MI355X.
  • Fixed the streaming GPU argmax failing to launch for reductions over more than 65,535 rows.
  • Attention OrMask and ChunkedMask patterns are now handled natively by the mask iterator, avoiding slower generic fallback paths.
  • Added float16 MMA support on NVIDIA GPUs, enabling float16 matmuls on NVIDIA tensor cores.
  • Improved tiny-M small-N FP32 GEMM performance by routing those shapes to a split-K GEMV kernel.
  • Sped up matmul by checking out-of-bounds only on border tiles.
  • Added fast RMSNorm GPU kernels: a single-pass kernel for small static-shaped rows and a barrier-free warp-per-row kernel.
  • Fixed mis-aligned memory accesses in the rms_norm reduction kernels and in layer_norm_cpu's sum reduction.
  • Added fused_qk_rms_norm_rope_ragged_paged and fused_dual_qk_rms_norm_rope_ragged GPU kernels fusing QK RMSNorm with rotary embeddings over ragged paged inputs; the dual variant also enables scale-free FP8 KV caches.
  • Sped up the Gated DeltaNet recurrence kernel by tiling one CTA per value head.
  • Added an initial functional MegaFFN kernel.
  • Fixed an expert-parallel decode crash when estimated_total_m floored to zero.
  • Added support for non-unit dilation in the CPU direct convolution kernel.
  • SM100 (Blackwell) attention and sparse-attention kernels:
    • Fixed rare crashes in the SM100 attention kernels: an illegal-instruction fault in flash-attention prefill under chunked prefill with tensor parallelism, cta_group=2 TMEM teardown faults, and an out-of-bounds shared-memory write in the 1Q split-K combine kernel.
    • Fixed a hang in the SM100 FA4 attention kernel on all-masked rows (valid_length=0).
    • Improved the SM100 FA4 flash-attention kernel: softmax warp-path optimizations and support for cluster sizes 6, 8, 10, and 16 across more multi-GPU configurations.
    • B200 decode attention partitioning is no longer restricted to powers of 2, improving decode performance for more batch and head configurations.
    • Improved the SM100 MLA prefill kernel: added a 1Q (num_qo=1) mode, Programmatic Dependent Launch for the generic BF16 path (up to ~1.37x on short sequences), and support for v_head_dim != qk_nope_head_dim.
    • Added native FP8 sparse MLA decode and prefill kernels for SM100, and extended prefill head-count coverage with a 64-head path and num_heads=8 support.
    • Added an SM100 tensor-core FP8 MLA sparse-indexer scorer kernel and improved sparse MLA prefill indexer performance via tensor-core scoring over query tiles.
    • Sped up the SM100 MLA sparse decode kernels by ~29% on the bf16 path and ~8% on the fp8 path.
    • Fixed sparse-attention kernel bugs: a crash during CUDA-graph capture in the sparse MLA decode dispatch, wrong results from the sparse MLA indexer with prefix-cached values, a top-k shared-memory overflow, and an MLA indexer fill bug.
  • Added Layout-E and Layout-g MMA paths for MHA prefill, improving attention performance on supported NVIDIA GPUs.
  • Enabled the Mojo SM100 FP32 GEMM in the kernel dispatcher, with optional 1D bias and an allow_tf32 opt-out to force full-FP32 precision.
  • Added weight prefetching to the SM100 block-scaled and structured matmul kernels.
  • B200 collective communication:
    • Enabled the low-latency (Lamport) all-reduce on B200 for small messages (up to 1 MiB at 2, 4, and 8 GPUs), roughly 1.1-1.68x faster than the one-stage path.
    • Added a fused Lamport all-reduce + RMSNorm kernel with PDL early launch, with column-aware dispatch on 8-GPU B200 (up to ~1.14x on the decode-critical-path collective).
    • Added a fused GEMM+ReduceScatter kernel that overlaps computation with communication on B200 (~10% faster than unfused for prefill shapes).
  • Sped up B200 MoE grouped matmul by removing a per-step synchronization.
  • Retuned the SM100 MoE decode kernels from 4-12 to 6 pipeline stages: ~5-7% better decode latency for MXFP8 grouped matmul and up to ~11% throughput for the NVFP4 down-projection at higher active-expert counts, with no change to numerics.
  • Fused SM100 projection and GEMM kernels:
    • Added a fused GEMM+Bias+SwiGLU kernel for SM100 (B200), exposed to graphs as mo.matmul_swiglu (~1.06-1.12x faster than unfused).
    • Added expand_qkv_sm100, a single-launch fused LoRA-B QKV expand kernel for SM100.
    • Added a fused MXFP8 QKV matmul kernel with paged KV-cache write.
    • Added a fused QKV + indexer-QKV projection matmul kernel for sparse-attention MLA models (~1.5-2x faster than the unfused pair).
  • Fixed FP8 MHA silently underflowing P values to 0.0 by scaling them before the softmax accumulation.
  • Fixed the SFB n_group per-lane computation and an N-edge out-of-bounds access in small_bn.
  • Fixed MXFP8 block-scaled matmul dispatch misses by falling back to cuBLAS on unsupported shapes.
  • Fixed LAUNCH_OUT_OF_RESOURCES failures during graph capture on SM100 by declaring the launch block size on the block-scaled small-BN matmul kernel.
  • Fixed an unaligned-access bug in the NVFP4 quantization async kernel.
  • Fixed the cluster_mask_base computation for tall clusters.
  • Fixed GEMM kernels producing incorrect or out-of-bounds C stores for partial tiles and odd sizes (multistage GEMM, SM100 small-MN GEMM).
  • Fixed the vendor-BLAS matmul ignoring the layout row stride, which could compute wrong results for strided inputs.
  • AMD GPU kernels:
    • Sped up GPU RMS norm on AMD CDNA4 (MI355X) by roughly 15-31% on prefill-sized shapes.
    • Extended the AMD MHA prefill kernel with depth-80 support and partial (non-power-of-two) head dimensions.
    • Fixed an AMD MHA decode crash caused by too many split-K partitions.
    • Added speculative-decode query-token folding to the AMD MHA decode kernel for higher multi-token decode throughput.
    • Reworked the AMD sparse-attention decode path: rewrote the sparse-indexer decode scoring and top-k as faster split-K partition kernels (fixing a max_num_blocks double-count), and added multi-token (MTP/speculative) support across the gfx950 FP8 MLA decode and sparse-indexer top-k kernels.
    • Fused MXFP4 expert-parallel operations on AMD MI355X (activation + quantize into down-projection, scale preshuffle into producer/dispatch kernels): +37% total TPM and -15% mean TPOT on an MXFP4 MoE model.
    • Improved AMD MXFP4 matmul and grouped-matmul performance via weight preshuffling, register staging, double buffering, async loads, and shared-memory swizzling.
    • Optimized all-gather and reduce-scatter on AMD CDNA4 (MI355X), bringing 4-GPU bf16 collectives to the PCIe fabric roofline (all-gather ~2.7x its previous default).
    • Sped up expert-parallel dispatch and combine on AMD GPUs via a volatile-load spin loop and reduced atomic contention.
    • Sped up float32 matmul/GEMV on AMD GPUs: enabled the GPU matmul kernel for float32 (fixing graph capture), dispatched float32 GEMV to vectorized/split-K paths (52us to 4us in decode), and enabled split-K for skinny-deep shapes (up to ~10x on small shapes).
    • Sped up bfloat16 matmul on AMD RDNA 3+ GPUs with faster WMMA-based kernels.
    • Fixed a race in the AMD 4-wave split-K matmul kernel.
    • Fixed the GEMV launch grid writing past the end of the output tensor on AMD GPUs.
    • Routed the AMD float32 MoE router through the cached GEMV path.
  • Apple GPU kernels:
    • Extended Apple silicon GPU support for MAX models to M1 and M2 systems by fixing the optimized matrix multiplication kernels on those GPUs.
    • Made the split-K decode attention kernel the default for token-generation attention on Apple GPUs (paged-KV-cache MHA and GQA decode); MODULAR_ENABLE_APPLE_NAIVE_FA_DECODE=0 opts out.
    • Restored the 32-bit atomic scatter_nd reduce path on Apple GPUs.
    • Further optimized the Apple M5 GEMM kernels; neural-accelerator utilization rose from 85% to 95%, within 2-3% of MPS.
    • Added Apple M5 int8 W8A8 matmul optimizations for NT layouts.
    • Added a hardware-MMA flash-attention prefill kernel for Apple M5 (Metal 4) with paged-KV support, giving ~1.4x lower TTFT at seq=8192 (bf16, head dim 128).
    • Added a fused online-im2col bf16 conv2d kernel for Apple M5 (Metal 4) and sped up its gather/MMA path by ~4x.

Breaking changes​

  • Dynamic CE chunk sizing in the data-parallel load-balancer (--dp-ce-balance-enable-dynamic-chunk-size) now defaults to off; re-enable it explicitly where its TTFT win holds for your workload.
  • Removed the deprecated --max-num-steps CLI flag and max_num_steps config field; the pipeline always runs single-step decode and passing the flag now errors.
  • Replaced PipelineTokenizer.eos (scalar) with PipelineTokenizer.eos_token_ids (a set of the declared EOS plus any terminators from the model config's eos_token_id entries).
  • MAX Serve now fails at startup when the device KV cache cannot hold a single request at the configured max sequence length (previously a warning followed by a runtime crash); the error reports the largest --max-length that fits.
  • max.nn.Module.build_subgraph() now takes representative input values (inputs=) instead of input types (input_types=); update call sites from build_subgraph(name, input_types=[v.type for v in values]) to build_subgraph(name, inputs=values). DistributedTransformerBlock now accepts list[PagedCacheValues] directly.
  • Removed the MAX_SERVE_METRIC_LEVEL and MAX_SERVE_DETAILED_METRIC_BUFFER_FACTOR environment variables; MAX Serve now always emits its full metric set. To record no metrics (previously MAX_SERVE_METRIC_LEVEL=NONE), set MAX_SERVE_METRIC_RECORDING_METHOD=NOOP or MAX_SERVE_DISABLE_TELEMETRY=1.
  • Removed InferenceSession.use_old_top_k_kernel() and the USE_OLD_TOP_K_KERNEL environment variable; the two-stage top-k kernel is now used unconditionally.
  • The Input, Output, MutableInput, FusedInput, and FusedOutput IOSpec values are now static members of IOSpec instead of module-level aliases; qualify them as IOSpec.Input etc. (e.g. Tensor[IOSpec.Input, spec]).
  • Removed the compiler Mojo package; import the 4 re-exported symbols from extensibility directly.
  • Renamed the metric maxserve_cache_hit_rate_percent_utilization (OTEL maxserve.cache.hit_rate) to maxserve_cache_request_prefix_coverage_percent (OTEL maxserve.cache.request_prefix_coverage); derive a token-weighted hit rate from maxserve_cache_hits_tokens and maxserve_cache_misses_tokens.
  • Removed the EXAONE model architectures (exaone and exaone_modulev3); EXAONE checkpoints are no longer served.
  • Removed InferenceSession.read; use max.engine.read to load a compiled-model artifact (.mef) without a session, then initialize it on any session via InferenceSession.init.
  • Removed the MAX_USE_EAGER_INTERPRETER environment variable; eager execution is now selected by MAX_EAGER_EXECUTOR (use compile to force compilation), and EagerRealizationContext(use_interpreter=...) is deprecated in favor of executor=....
  • Removed eager-interpreter support that only existed through broken or GPU-only Mojo bindings: conv2d_transpose (cuDNN-dependent, crashed on Apple GPUs), resize_bicubic, and CPU group_norm (now raises NotImplementedError; GPU still works).
  • Narrowed the eager interpreter's CPU coverage to float32/float64 for most float families (the previous Mojo bindings also accepted float16/bfloat16); unsupported dtypes now raise immediately.
  • Changed TileTensor to drop its element_size parameter; element width is now carried by Storage=PointerStorage[element_width=N], or use TileTensor.vectorize().
  • Moved accelerator-programming APIs from the Mojo standard library into the max package: std.algorithm -> max.algorithm, std.benchmark -> max.benchmark, and std.gpu.{compute,host,memory,sync} -> max.gpu.*.
  • The layout package is now bundled with MAX instead of Mojo.
  • Renamed LayoutTensor.get_immutable() to as_imm(); the old name remains as a @deprecated alias.
  • Deprecated the Python local and tiered KV cache connectors in favor of the Rust rust_tiered connector (--kv-connector rust_tiered, CUDA-only); the Python connectors remain for backward compatibility.
  • Made the max.entrypoints package private (max._entrypoints) and deprecated max.entrypoints.LLM; a new offline-inference API will follow in a future release.
  • Renamed the FLUX.2 int8 W8A8 override env var FLUX2_KLEIN_INT8_W8A8 to APPLE_FLUX2_INT8_W8A8; the old name is still honored with a one-time deprecation warning.
  • The dKV connector now requires a non-empty tenant identity (MODULAR_DKV_TENANT_ID); the empty-tenant default path is removed, so every deployment routes through the per-tenant region-sharded store. A single-tenant multi-GPU node must set the dKV server's --fair-share-partitions to its GPU count.
  • Made the tiered (disk) KV cache offload ephemeral: the disk tier starts empty and removes its max_kv_tiered_* directory on shutdown; use dKV for persistent KV across runs.
  • Retuned the Prometheus/OpenTelemetry histogram buckets for MAX metrics: boundaries are now generated instead of hand-written, matched to each histogram's actual range (latencies to 8 hours, token counts to 100M) with a leading 0 boundary. Dashboards and alerts that hardcode le values need updating.
  • Changed maxserve.cache.num_used_blocks and maxserve.cache.num_total_blocks from counters to gauges; the exported series drops the counter _total suffix.

Fixes​

  • Fixed Apple GPU issues:
    • Fixed a per-kernel-launch memory leak on macOS Metal that accumulated autoreleased driver objects for the lifetime of the thread.
    • Fixed three Apple Metal kernel bugs (rope_ragged, concat, split) that silently produced zeros on Metal.
    • Fixed the FlashInfer top-k/top-p samplers returning out-of-top-K tokens on Apple GPUs.
    • Fixed models failing on pre-M5 Apple silicon GPUs (NVFP4 weights, MHA, and 2-D convolution fallbacks).
  • Fixed GPU programming issues:
    • Targeting an MI250X now works. While normalizing the architecture name, gfx90a was rewritten to the nonexistent gfx90aa, so both --target-accelerator and GPUInfo.from_name() reported every spelling of the target (gfx90a, mi250x, amdgpu:gfx90a, and amd:gfx90a) as an unsupported architecture.
    • Fixed print() and debug_assert() emitting garbled output on AMD GPUs when a printed string's byte length was an exact multiple of 8. The AMDGPU hostcall printf interface reads each string up to its nul terminator, and the terminator was being dropped in that case, so the host read past the payload.
    • #6784, #6434 - math.sqrt() on Float64 now works on NVIDIA GPUs. NVIDIA has no approximate F64 sqrt, so it lowers to the IEEE correctly-rounded hardware sqrt (sqrt.rn.f64) instead of being rejected at compile time. The Float32 fast path continues to use sqrt.approx.ftz.f32.
  • Fixed --model-revision and --model-override with an external weights repository or path being resolved against the base model repo, failing weight downloads.
  • Fixed a custom op with a data-dependent output dimension killing the process when its kernel registers no shape function; it now reports a compile error.
  • Fixed MAX Serve containers ignoring SIGTERM during model cold start, leaving pods stuck Terminating until the grace period elapsed; the image now runs under dumb-init, so pods shut down promptly and release their GPUs.
  • Fixed max serve crashing on startup with TypeError: cli_serve() missing 1 required positional argument: 'task'.
  • Fixed max serve failing at startup with a device_graph_capture error for models in the overlap-scheduler auto-disable list (e.g. Gemma 4) when --device-graph-capture is passed explicitly.
  • Fixed a hang during max serve shutdown after failures (e.g. OOM).
  • Fixed max warm-cache crashing in virtual-device (compile-only) mode with --draft-model-path.
  • Fixed a graph-compilation failure on B200 for models with materialized attention masks, such as the padded text encoders in diffusion pipelines (e.g. FLUX.2's Qwen3 text encoder).
  • Fixed structured-output, tool-calling, and reasoning issues:
    • Fixed the structured-output grammar backend silently defaulting to llguidance instead of xgrammar for models launched via max serve.
    • Fixed structured-output responses being truncated when the model sampled an extra eos_token_id mid-response.
    • Fixed an explicitly empty ({}) or boolean-true JSON schema in response_format being treated as "no structured output" instead of "any valid JSON value".
    • Fixed response_format constraints being dropped after an auto tool call for Gemma 4 and Kimi K2.5.
    • Fixed Gemma 4 tool-call structured output enforcing JSON Schema minLength/maxLength as a token count rather than a character count.
    • Fixed grammar enforcement being intermittently dropped during speculative decoding: across tool-call structural tags, mid-span when a committed draft span contained a terminating token, and batch-wide when a single request was preempted.
    • Fixed a grammar FSM off-by-one at the chunked-prefill-to-decode seam that could misalign constrained-decoding state.
    • Fixed streaming tool-call glitches: dangling tool calls with empty arguments, structural namespace tokens leaking into assistant content (MiniMax M2), empty SSE packets while the parser consumed structural tokens, and same-index tool-call entries split across a chunk.
    • Fixed the tool-call conformance check turning successful responses into 500 errors or dropping valid tool calls when a client schema failed validation; parameter-less tools are no longer flagged as unknown_tool.
    • Fixed non-streaming chat completions leaking a literal <tool_call> marker into message.content when max_tokens truncated mid tool-call block; the response now surfaces only the preceding content with finish_reason == "length".
    • Fixed Kimi tool-call arguments decoding unconstrained under tool_choice="auto" when the model opened a tool call without a closing </think>, and the combined tools + response_format=json_schema grammar letting non-conforming output through.
    • Fixed Gemma 4 with MTP speculative decoding crashing the model worker on the first tool-call grammar.
    • Fixed non-streaming Kimi reasoning responses returning null message.content when the model stopped without </think>, and added validation rejecting out-of-range or non-finite frequency_penalty/presence_penalty values.
    • Fixed replayed assistant reasoning being silently dropped when sent back to chat completions; reasoning_content is also accepted.
    • Fixed Nemotron-3-Nano leaking chain-of-thought and a raw </think> delimiter into message.content and emitting zero tool calls with tool_choice="required"; it now defaults --reasoning-parser and --tool-parser to qwen3_5.
  • Fixed ops.scatter_add/scatter_mul/scatter_max/scatter_min and their scatter_nd_* variants silently dropping concurrent updates when indices contains duplicates; reductions are now atomic on CPU and GPU.
  • Fixed scatter/scatter_add writing out of bounds under multi-device sharding in max.experimental, corrupting unrelated allocations.
  • Fixed the compiled-model cache (.max_cache) serving stale compiled models after Mojo kernel-library changes; kernel edits now correctly trigger recompilation.
  • Fixed sparse-attention (MLA) issues:
    • Fixed sparse-attention MLA models (DeepSeek V3.2, GLM 5.1/5.2) with an FP8 latent KV cache falling back to a dense prefill path that cost accuracy and long-context performance.
    • Fixed the SM100 sparse MLA prefill kernel reading the wrong layer's KV cache for every layer above the first, silently corrupting attention for multi-layer sparse-attention models (DeepSeek V3.2, GLM 5.1/5.2); also enabled the kernel for GLM 5.2 tensor-parallel head shards.
    • Fixed sparse MLA top-k dropping the indexer -1 pad sentinels, which could corrupt padding handling in sparse-attention index selection.
    • Fixed a warp-collective deadlock in the sparse MLA decode indexer kernel.
    • Fixed disaggregated prefill_only serving of DeepSeek V3.2 ignoring the sparse indexer's top-k selection.
  • Fixed speculative-decoding issues:
    • Fixed speculative decoding (Eagle) requests stopping short of the per-request max_tokens cap, returning finish_reason="length" with fewer tokens than requested.
    • Fixed the first structured-output request on a speculative-decode (EAGLE) server stalling ~7-9 seconds with the GPU idle.
    • Fixed a crash taking down the decode-only engine on disaggregated (1p1d) deployments running structured output with Eagle3 speculative decoding.
    • Fixed Eagle speculative decoding with FP4-quantized models: Kimi K2.5 with MXFP4 weights produced garbage tokens, and pairing an FP4 target with a bfloat16 Eagle3 draft failed to load.
    • Fixed DeepSeek-V3.1-NVFP4 multi-token prediction failing to load with expert parallelism enabled.
    • Fixed the float32 weight-path fallback firing for any model, breaking architectures without f32 support (e.g. f32-only Eagle3 drafts).
  • Fixed precision loss in rms_norm, layer_norm, group_norm (and fused variants) where a small epsilon such as 1e-6 was silently rounded in bfloat16; epsilon is now carried as float32.
  • Fixed MAX crashing the model worker on the first host KV-cache offload/reload with --kv-connector dkv.
  • Fixed memory-planning and OOM issues:
    • Fixed over-provisioned KV-cache offload budgets (host_kvcache_swap_space_gb, disk_offload_max_gb) OOM-killing the server or filling the disk mid-startup; the local and tiered connectors now fail fast with an actionable error.
    • Fixed the VMM defragmenting allocator failing large contiguous allocations, which could OOM long-context KV caches despite free memory.
    • Fixed steady device-memory drain during serving; model inputs now reuse fixed-capacity backing buffers.
    • Fixed memory planning failing on devices that report zero free memory.
    • Fixed a ValueError during memory planning for diffusion pipelines.
    • Fixed the Gemma 4 vision encoder running one unbounded forward pass over every image regardless of chunked prefill, which could OOM the model worker on many-image requests.
    • Fixed multimodal serving OOMing under load despite free device memory; the vision encoder cache no longer fragments the allocator and releases over-limit entries promptly.
    • Fixed Gemma 4 activation-memory estimation under-estimating with tensor parallelism (TP=2) plus MTP speculative decoding, causing OOMs.
  • Fixed a use-after-free of pinned buffers when multiple streams were in use in the driver.
  • Fixed a token-input staging buffer race under the overlap scheduler that could corrupt inputs.
  • Fixed multimodal serving issues:
    • Fixed a crash in vision encoder cache assembly on a prefix-cache hit covering an evicted image.
    • Fixed Kimi-K2.5 device graph capture failing with an input-count mismatch after the vision cache migration.
    • Fixed two crashes in disaggregated serving of vision-capable models.
  • Fixed tensor-parallel and expert-parallel accuracy and performance issues in DeepSeek-V3 ModuleV3, and added device-graph-capture support.
  • Fixed a GPU memory fault in grouped_matmul_ragged for unquantized bf16 MoE experts under expert parallelism on SM100 (e.g. GLM 5.2 NVFP4's MTP draft).
  • Fixed a crash on data-parallel replicas with no assigned requests.
  • Fixed a FLUX.2 multi-GPU accuracy regression.
  • Fixed failures loading models whose transformers AutoConfig is not registered in the installed transformers version.
  • Fixed Hugging Face weight and dataset downloads failing with FileNotFoundError on a racy .incomplete cache entry; downloads now retry past the transient race.
  • Fixed grouped (groups > 1) ops.conv2d/ops.conv3d on CPU raising grouped conv requires packed filter for non-constant filters.
  • Fixed max.nn.WeightNormConvTranspose1d raising AttributeError when constructed with the default has_bias=False.
  • Fixed max.experimental.nn.Conv2d.forward failing with a device mismatch when the bias started on a different device than the input.
  • Fixed max.experimental.functional.while_loop raising AttributeError at graph-construction time.
  • Fixed slicing and view() on a max.driver.DevicePinnedBuffer silently returning a plain Buffer, causing unexpected device synchronization on to_numpy().
  • Fixed max.driver.set_virtual_device_*() settings having no effect on macOS device creation.
  • Fixed DeviceContext.load_function returning the wrong function when loading two entry points from one PTX/cubin blob.
  • Fixed top_k/bottom_k mis-shaping or rejecting graphs using a negative axis.
  • Fixed sliding_window being silently dropped when combined with a causal mask.
  • Fixed AMD GPU issues:
    • Fixed an AMD MLA decode crash for models with num_heads > 16 on short prefills of 2-8 tokens, common with multi-turn prefix caching.
    • Disabled the v2 attention prefill kernel on AMD GPUs due to a race producing NaNs.
    • Fixed NaNs from the fused SwiGLU kernel on AMD GPUs.
    • Fixed a correctness race in the AMD expert-parallel combine kernel.
    • Fixed diffusion models failing on AMD RDNA GPUs.
  • Fixed usage reporting at zero-token edge cases: prompt_tokens=0 in the final streaming chunk and null usage in non-streaming completions.
  • Fixed image/video URL fetches being rejected by some hosts (now sends a browser User-Agent) and failing with a 500 on slow or large transfers (now a clean 4xx InputError on timeout).
  • Fixed requests containing unpaired UTF-16 surrogates (e.g. a truncated emoji) crashing the tokenizer; they are now replaced with U+FFFD.
  • Fixed the benchmark harness crashing with a 500 when flushing the prefix cache on diffusion models; /reset_prefix_cache now returns a clean 400 for pipelines without a KV cache.
  • Blocked the MagicYUV video codec pending a fix for CVE-2026-8461; video input using that codec is now rejected.

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