Skip to content

fold: upstream v0.40.5 — Nemotron 3.5 Super/Bolt; llama.cpp b11669, MLX bebd6422 - #498

Draft
glennneuber wants to merge 57 commits into
mainfrom
task/upstream-sync-0.40.5
Draft

glennneuber wants to merge 57 commits into
mainfrom
task/upstream-sync-0.40.5

Conversation

@glennneuber

@glennneuber glennneuber commented Oct 16, 2026 •

Copy link
Copy Markdown

Folds upstream v0.40.5 into main, taking v0.40.4 on the way. Draft: gates 1–3 are done, and the CUDA image is building.

Why now. v0.40.5 adds NVIDIA Nemotron 3.5 Super (120B-A12B, a hybrid Mamba-2/Transformer MoE with a C-RADIO vision encoder) and Nemotron 3.5 Bolt (74B-A9B). The library refuses both on our current build:

  • Super answers "requires ollama version v0.40.5 or newer";
  • Bolt answers HTTP 412.

The architecture needs v0.40.5's llama.cpp and MLX.

Size: 51 commits, 172 files. All three native pins move:

  • llama.cpp b11351 → b11669;
  • MLX a59cc231 → bebd6422;
  • MLX-C ebc88f10 → a341b492.

So every profile needs a new measurement (ADR 0032). The record is docs/maxusai/tasks/upstream-sync-0.40.5.md.

gate state
1, merge bd36fd015, 5 conflicts resolved (below). go build/go vet clean; go test 63 packages ok. The one failure, cmd/launch TestCodexAppCountsOnlyOllamaRequestsInRegularProfile, also fails 1 of 3 runs on a clean v0.40.5 checkout
2, docs and paths clean
3, patch series 8761533a6: all nine patches apply to a clean b11669. 002 re-anchored, identical changed lines; 903 cut to guard-only, because the pin now contains llama.cpp#29953
GPU-free harness green
4, image building on CUDA (sync-0.40.5-fold, stamp 0.40.3-dynres-12-g8761533)
5, 6 not run

Conflicts, in short:

  • gemma4 (MLX): takes upstream's single SDPA path, which drops the manual CUDA prefill for head_dim > 128, with the fork's bidirectional vision mask. gemma4 prefill numerics on CUDA MLX change.
  • nemotron_h: takes upstream's error-returning patch grid and keeps the fork's per-request floor.
  • pipeline cleanup:
  • sched: takes upstream's LlamaServerBatch and keeps the fork's image-chunk floor.
  • routes:

For each host, once the image exists:

  • gate 5 against new -0-40-5 profiles;
  • 903's GPU retirement gate (CUDA, then gfx1151 for hipcc);
  • 908's loop-rate gate on gemma4:26b think-on (b11669 reworks the FA tile scheduling around it);
  • a Nemotron 3.5 pull, load and vision check.

From ai-server/mlx-cuda via /collab · 2026-10-16 · build 0.40.2-dynres.1-17-g75c6d35 · signed ai-server

🤖 Generated with Claude Code

Anna Rahn and others added 30 commits October 6, 2026 12:28
# Conflicts:
#	docs/api/introduction.mdx
#	docs/cloud.mdx
#	docs/docs.json
#	docs/guide/capabilities/decision.mdx
#	docs/index.mdx
#	docs/integrations/cline.mdx
#	docs/quickstart.mdx
ollama#18947)

* server: reuse verified manifests and loaded models in inference routes

Resolve named manifests and runner selection in one verified cold lookup.
When a runner is loaded, reuse its model after checking that the v2 tag
still points to the same digest. Fully verify regular-file manifests,
including Windows copies.

Share the lookup across generate, chat, embed, and System One. Keep parser
selection request-local so reusing a model does not mutate its config.

Cover tag changes, runner selection, concurrent unloads, and existing
missing-model responses with regression tests and a lookup benchmark.

* address comments
…ama#18955)

* mlx: compact restored recurrent state after speculative rollback

Per-token recurrent captures can leave a restored state pointing into a larger buffer. If a request ends before another model step, the live cache retains that buffer even though it accounts for only the selected state.

Track when a restored state may share that backing buffer. At request close, compact marked states in one batch before retaining them in the prefix cache. Clear the marker when a normal model step replaces the state, and cover restoration and buffer detachment in tests.

* address comments
- Directs users to Github advisories instead of hello email address
* Add Kev decision scoring to MLX

Load Kev models with their FP32 decision head and release metadata.
Support System One scoring with shared-state caching and Kev-specific limits.

* address comments
A failure during decoding can trigger deferred drain and cache close, which may panic again and hide the original error. Skip stateful cleanup while a panic is unwinding so the runner reports the initial failure.

Add regression tests for the masked panic paths and normal cache saving.
Try a little harder instead of giving up on the first failure.
MLX's SDPA already handles head dims > 128 on CUDA, and is faster than our formerly required hand-rolled one.
Anna Rahn and others added 27 commits October 13, 2026 10:05
* Add decision model integration tests and benchmarks

Support external JSONL corpora, quality scoring, concurrent requests, vision, and System One timing metrics.

* Refine decision tests and benchmarks with cache checks, external endpoints, and consistent response validation and timing
* renderers: add Python json.dumps serializer

* decision: render JSON states like the publishers' json.dumps
docs: update README to include cloud and privacy language
* mlx: port the Metal custom kernels to CUDA

mamba2_scan, depthwise_conv_silu and gated_delta (plus their _states/_bias
variants) had a Metal source only, so on CUDA each fell back to graph ops.

Deviations from the Metal sources:

- The scan returns early past D: Metal dispatches an exact thread count,
  CUDA rounds the grid up to whole blocks.
- T is dereferenced; CUDA passes inputs by address, Metal scalarizes
  zero-dimensional ones.
- simd_sum becomes a shfl butterfly, and the threadgroup q/k caches become
  per-warp registers: each warp normalizes the row its own lanes consume,
  costing a redundant norm but no shared memory and no barriers.
- The gates and SiLU call MLX's Sigmoid/LogAddExp/Exp instead of restating
  their formulas. Metal cannot: its custom-kernel preamble is metal::utils()
  alone, while CUDA embeds the op headers for nvrtc.

The parity tests now run on Metal and CUDA hosts.  The fused gated-delta
comparison takes a tolerance: the kernel and the graph path reduce in a
different order, so the bf16 output can land one ULP away.

* address comments
…odels_wording

docs: model recs and language standardization
Co-authored-by: Jeffrey Morgan <jmorganca@gmail.com>
* ci: shorten Windows Vulkan build path

Keep the nested shader-generator compiler probes below Windows path limits in checkouts with longer repository names.

* mlx: fix Nemotron bugs and support global scales

Accept published Nemotron layer names while retaining legacy aliases, and
fix layer-count inference and detached MTP quantization. Stabilize Mamba
softplus, preserve FP32 vision normalization, and compile Mamba gate and
norm scaling.

Support per-expert F32 global scales and preserve activation-scale
companions in ModelOpt and compressed-tensors NVFP4 imports. Normalize
reciprocal globals and reject invalid scales. Keep activation scales
separate from weight-only inference.

* mlx: add Nemotron 3.5 Super

* minor adjustments
…bebd6422

Merges upstream v0.40.4 and v0.40.5 into main (v0.40.3 fold). Five files
conflicted; resolutions:

- mlxrunner/model/gemma4/gemma4.go: upstream ollama#18809 drops the manual CUDA
  prefill attention for head_dim > 128 and uses MLX SDPA everywhere. Taken,
  with the fork's bidirectional vision-chunk mask in place of upstream's
  baseMask parameter. This changes gemma4 prefill numerics on CUDA MLX.
- mlxrunner/model/nemotron_h/vision_prompt.go: upstream's patch grid now
  returns an error; the fork keeps its per-request floor (withMinNumPatches).
  media_budget_test.go follows the new signature.
- mlxrunner/pipeline.go: upstream ollama#18886 (skip stateful cleanup while a panic
  unwinds, so a failed request saves no partial cache) and the fork's
  guardClose solve the same masking bug. Upstream's skipCleanupOnPanic is
  taken at every site; it now re-raises the panic wrapped once as the fork's
  *firstPanic (asFirstPanic in unwind.go), so the first stack still reaches
  the log. pipeline_panic_test.go unwraps it before comparing.
- server/sched.go: upstream folds the batch surcharge into
  llm.LlamaServerBatch (same 768 MiB / 2 GiB steps, plus decision models);
  the fork's image-chunk floor on automaticGenerationBatch stays.
- server/routes.go: upstream ollama#18947 stops mutating m.Config.Parser for
  harmony (models are now reused across requests) and names the parser with
  parserNameForModel. Applied to both of the fork's generate parser blocks
  and to its chat-handler trace logs; upstream's systemOneLlamaServer is
  added beside the fork's structured-outputs helpers.

go build and go vet are clean; go test passes 63 packages. The one failure,
cmd/launch TestCodexAppCountsOnlyOllamaRequestsInRegularProfile, is an
upstream flake: it fails 1 of 3 runs on a pristine v0.40.5 checkout too.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
All nine llama/compat patches apply in order to a clean b11669 with plain
git apply. Two needed work:

- 002 (nemotron dynres): re-anchored only. b11669 adds a separate
  PROJECTOR_TYPE_NEMOTRON_OMNI case right after NEMOTRON_V2_VL in clip.cpp and
  mtmd.cpp, which moved both hunks' context. The changed lines are identical
  (+20/-3). The patch is still needed: the compat layer (upstream's too) maps
  the nemotron_h_omni GGUFs nemotron3 ships to the nemotron_v2_vl projector,
  which b11669 still preprocesses on a fixed canvas.
- 903 (MMQ padding): guard-only. b11669 contains llama.cpp#29953 itself
  (fc9ce6b9d) and #30168, so the backported padding is dropped and 903 keeps
  the y-tile helper and its compile-time static_asserts, re-anchored on
  #30168's template line. The retirement gate's GPU half is still to run.

Upstream changes the other patches sit next to, checked by reading: 908's
config entries and tile sizes apply unchanged around upstream's FA race fix
(#30103) and whole-tile scheduling (#29435), so its loop-rate gate must be
re-measured; clip's build_attn now asserts an f16 mask, and no fork patch or
the nemotron_v2_vl graph passes one; mtmd-image's changes touch the llava-uhd
and longest-edge preprocessors only, not dyn_size.

mlx/compat/mlx-c (upstream's two carry patches, unchanged) applies to MLX-C
a341b49.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… on CUDA

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…_120a and sm_75

The tree (b11669 plus the fold's 805, 903 guard-only, 908) passes every
guard-page, stock and exact-allocation run on both GPUs. The J_best src1
control faults 3/3 on ids16 (both GPUs) and dense321 (sm_75); the no-ids_dst
control faults 3/3 on every MoE case. mmq-0405-gate.sh is mmq-462-gate.sh at
b11669; results beside #462's in mmq-successor-results/.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…red on the fold image

New profiles for the moved payload (llama.cpp b11669 620d2225c, MLX bebd6422),
admitting the v0.40.5-dynres lineage. Measured on sync-0.40.5-fold
(0.40.3-dynres-12-g8761533) in a production-like canary: all six ladders
byte-identical to the 0-40-2 rows they carry; preflight PASS=20 SKIP=9 (cuda)
and PASS=11 SKIP=18 (mlx-cuda), both pins asserted, 0 cudaMemcpy errors. The
interim stamp stays on 0-40-2 by pattern (it cannot be told from v0.40.3), so
the canary used a run-local expectations copy; the test class says why.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… on the fold image

908's loop-rate gate on CUDA: gemma4:26b-a4b think-on NOT CONVERGED 1 with 908,
2 without (b11081: 1 and 6), both arms -b 2048, one library apart. With 908,
gemma4:31b and 26b think-off are byte-identical to production's b11351 cells;
e4b moves in both arms with every check unchanged. Throughput unmeasured
(shared GPU0). Nemotron 3.5 Bolt (IQ2_XXS) and Super (Q4_K_M) load fully on
GPU and answer a text and an image request; a smoke test, not a quality run.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@glennneuber

Copy link
Copy Markdown
Author

CUDA host: gates 4–6 are done on CUDA, and 908 stays. Every result is in the fold record; this is the summary.

  • Gate 4: the image sync-0.40.5-fold (0.40.3-dynres-12-g8761533) is a full native build, 4 h 07 min.
  • Gate 5: new profiles cuda-0-40-5 and mlx-cuda-0-40-5, in cd197ed69.
    • All six ladders are byte-identical to the 0-40-2 rows: nemotron_h_omni, gemma4 and qwen35 on CUDA; gemma4, qwen35moe and qwen35 on MLX.
    • Preflight passes in a production-like canary (--gpus all): PASS=20 SKIP=9 on cuda, PASS=11 SKIP=18 on mlx-cuda.
    • Both payload pins check out (620d2225c, bebd642), with 0 cudaMemcpy errors and 0 OOMs.
  • 903's padding retires (10056c858). On b11669 it passes every guard-page and exact-allocation run on sm_120a and sm_75, and its J_best and no-ids_dst controls fault 3/3. 903 is now the y-tile guard only.
  • 908 stays (f83c5d2b1). gemma4:26b-a4b think-on on the full ladder fails to converge 1 time with 908 and 2 without, against 1 and 6 at b11081. Both arms are at -b 2048, one library apart.
    • With 908, gemma4:31b and 26b think-off are byte-identical to production's b11351 cells on all 27 rows.
    • e4b moves in both arms, with every pass/fail check unchanged.
  • Nemotron 3.5 Bolt and Super load fully on GPU on the fold image and answer one text and one image request each. That's a smoke test, not a quality run.

One thing every host will hit. The interim stamp 0.40.3-dynres-N cannot pull Nemotron 3.5. The model config's requires and the registry's User-Agent check both read the version string, at pull time only. The CUDA host pulled with upstream's stock v0.40.5 client into the same store; a v0.40.5-dynres tag fixes this for fork builds.

For Metal and ROCm:

  • gate 5 against your own -0-40-5 profiles: the 0-40-2 ones claim the interim stamp, so a run-local expectations copy is needed, as canary-0405.sh used;
  • hipcc for 903's guard (gfx1151);
  • your gate 6.

From ai-server/mlx-cuda via /collab · 2026-10-17 · build 0.40.3-dynres-12-g8761533 · signed ai-server

@glennneuber

Copy link
Copy Markdown
Author

CUDA production now serves this fold (deployed untagged, on the maintainer's word)

On 2026-10-21 the maintainer asked for "the latest build on CUDA", so production :11497 on the CUDA host now runs this PR's candidate image. #498 is still an unmerged, untagged draft. Production therefore serves a build that is not on main, and its version reads 0.40.3-dynres-12-g8761533 until v0.40.5-dynres is cut.

before after
container ollama-0.40.2-dynres.1-17-g75c6d35 (stopped, kept) ollama-0.40.3-dynres-12-g8761533
image maxusai/ollama:main-75c6d350 (v0.40.3 fold) maxusai/ollama:sync-0.40.5-fold (built at 8761533a6; later commits touch docs only)
swapped 2026-10-21 15:12 AEST, after 5 idle minutes

The swap mirrored the live container: --gpus all, the same ollama_store_8tb volume, OLLAMA_KV_CACHE_TYPE=f16 and OLLAMA_FORMAT_TWO_PASS=1. It health-checked the new container and was set to roll back on failure; it did not need to.

Post-deploy, against production itself:

check result
preflight cuda (profile cuda-0-40-5) PASS=20 SKIP=9, rc=0
preflight mlx-cuda (profile mlx-cuda-0-40-5) PASS=11 SKIP=18, rc=0
tags in the store 63 → 63
KV allocations f16 only
server config OLLAMA_FORMAT_TWO_PASS:true
cudaMemcpy errors (--gpus all, the #487 case) 0
out-of-memory 0

The preflight read the canary's run-local expectations (expectations-0405-canary-g8761533.toml). Those contain the -0-40-5 profiles widened to admit the interim stamp, because the committed -0-40-5 patterns admit only 0.40.5-dynres and the stamp otherwise resolves to the 0-40-2 profiles. The counts match the device=0 canary's.

Rollback: docker rm -f ollama-0.40.3-dynres-12-g8761533 && docker start ollama-0.40.2-dynres.1-17-g75c6d35

Not changed: the Metal and gfx1151 production hosts still run the v0.40.3 fold. This PR's Metal and ROCm gates are still open. Merging, tagging v0.40.5-dynres and the other hosts' deploys remain the maintainer's call.

From ai-server/mlx-cuda via /collab · 2026-10-21 · build 0.40.3-dynres-12-g8761533 · signed ai-server

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

9 participants