[laudney]

Part of #1927.

This serves the publisher's BF16 dots-studio/dots.mocr (revision e539fbb52280393adc081b289ec597430a0f9031) natively on HIP. Its template, tokenizer, image processor, weights and three EOS IDs come from the checkpoint unchanged. It is the checkpoint other engines publish numbers for and the one users download first, so serving the publisher's weights matters for benchmarking and adoption. The image tower and the Qwen2 decoder run in Zig and HIP, with no Python in the served engine, behind /v1/chat/completions with image inputs, shared lanes, drafted verification and the prompt cache.

	Before (main)	After (this PR)
dots-studio/dots.mocr (BF16) on HIP	not supported	served natively on Radeon AI PRO R9700 (gfx1201)
Python or libtorch in the served engine	—	none; one pinned HF oracle under tools/zig/ for checks
Model-card demo vs HF	—	all 4,634 output IDs match through EOS
Two-image case vs HF	—	all 108 IDs match
Kernel contract on gfx1201	52 kernels pass gates 1-4	63 kernels pass gates 1-4, 9/9 controls red
Inherited limit problems (live gfx1201 / offline, 4 chips)	124 / 503	124 / 503, the same entries (0 new)
Existing models	—	unchanged: causal prompt tiles compile to the same code; new cache and lane rules are opt-in

Prerequisites #1929, #1930 and #1933 are merged. This branch still carries their earlier copies at its base, so a rebase onto main leaves only the dots.mocr commits, de07fc7a9^..459512e31.

Commits

1. Checkpoint and image preparation: config, weights and image contracts against the pinned checkpoint, and the Qwen2-VL image processor with its fixtures. This commit also adds the HF oracle, tools/zig/dots_mocr_reference.py:
   • oracle runs the model card's own example and saves IDs, scores and seams.
   • processor regenerates the fixtures.
   • compare checks a native greedy trace against the oracle.
2. BF16 weights on HIP: checked ownership, layout and a budget refusal before any upload.
3. The native graph:
   • The image tower: patch convolution, 2D vision RoPE, full attention for each image (as the publisher's cu_seqlens spans), and the merger LayerNorm/GELU.
   • The Qwen2 decoder.
   • New HIP ops (columns, dense_text, vision) with kernel-contract cases and launch records.
   • Full image attention reuses the prompt tile in attention/prefill.hip as a non-causal instantiation. A NaN score makes its row NaN, as in Torch. The existing causal instantiations compile to the same code and keep their contract results.
4. Live KV prefix copies across cache capacities.
5. Serving buffers sized inside the GPU's free memory, refused before load when they do not fit.
6. Window and partial-keep equivalence tests.
7. Serving through native HIP lanes: registration, image offers, flag refusals (--tp, MTP drafts, unsupported policies), warm passes that yield to foreground requests, and the recipe docs/recipes/dots-mocr-hip.md.
8. Review follow-ups: the dots refusal test asserts !reasoning(dots.native), and the oracle pins the SHA-256 of all 14 non-weight sources at e539fbb52, with a sources command that checks a checkpoint without a GPU.

Shared changes (other models unchanged)

• The family declares pub const reasoning = false; (Server: one engine flag for models whose replies carry no reasoning block #1933), so OCR text is always answer content. Its chat template does not read enable_thinking, so the prompt is the same either way.
• native/hip.zig: three optional family hooks (refusal, openVision, attach). A family without them takes the existing path. The dots refusal test asserts that dots.mocr declares no reasoning block; Server: one engine flag for models whose replies carry no reasoning block #1933's test checks the hook itself.
• Rules.warm_usable (default false): with it, a warm pass keeps its end only where Store.usable resumes it. With it off, both call sites in lane_host_warm.zig reduce to the existing condition. lane_host_warm_test.zig covers both settings.
• build.zig gains one line after const test_step. core: a GGUF reader and an exact repack of GGUF's quantized blocks (#1920, steps 1-2) #1926 and Native OCR models on HIP: dots.mocr, GLM-OCR, LightOnOCR-3 and Nemotron Parse, CUDA to follow (tracker) #1927's GLM-OCR PR add a line at the same place, so whichever lands later keeps both lines.

Gates

• Against HF:
  • The complete model-card demo returns all 4,634 HF output IDs through EOS over HTTP, with the full 5,096-token prompt and the native tower's own features.
  • A two-image case matches all 108 HF IDs.
  • A square synthetic case reads the same text but differs in four box-coordinate tokens, by one pixel per edge. The tower activations are not byte-identical to HF's, and the recipe keeps this limit visible.
  • Plain and drafted, fresh and resumed, and solo and concurrent replies match for the checked cases, including seeded sampling. Cancellation, cache image identity, eviction, idle reload and failure cleanup are checked.
• Kernel contract on a Radeon AI PRO R9700 (gfx1201): SUMMARY gfx1201 63 kernels: 63 pass gates 1-4, 0 fail, 2 known findings; 9 controls: 9 red as planted, 0 missed; 124 limit problems.
  • The upstream/main base reports 52 kernels ... 124 limit problems on the same GPU, and the same limit problems appear on both. So the 124 problems are inherited, and none is a new launch.
  • New full image attention: 31040 non-finite and 460480 finite outputs match the float64 classes, five byte-identical runs, 8 shapes within the float64 bound, and exact sums at k 57600.
  • Offline records: an offline limits audit for gfx1030, gfx1100, gfx1151 and gfx1201 lists the same problems as upstream/main, and no new problem.
• Host: test-hip-blas-host, test-dots-mocr-host, test-dots-hip-host, test-dots-hip-weights-host, test-dots-native-host and test-lane-api-host pass. zig build test: 469/469 steps; zig build test -Dhip=true: 502/502 steps.
• Commits: every commit builds native-hip for gfx1201.
• Lean checker: zero new problems over the 425 baseline.
• Not qualified: CUDA, Metal, video, the MTP head and tensor parallelism. Native OCR models on HIP: dots.mocr, GLM-OCR, LightOnOCR-3 and Nemotron Parse, CUDA to follow (tracker) #1927 tracks CUDA.

[ashhart]

Thanks @laudney. We've reviewed this independently, the code against your account, since we have no AMD GPU.

What the review established:

• Nothing shared changes behaviour for another family. The HIP registry hooks are optional, and existing families declare none. Info.reasoning defaults to true. The new prompt-cache rules default off. Existing attention instantiations keep CAUSAL=true, and the NaN-row handling applies only to non-causal attention. rocBLAS loads only when its owner opens.
• The checks pass on host: 495 tests (4 checkpoint or image skips) and 261 HTTP and admission checks.
• The checkpoint is validated first: all 643 tensor specs and 1,056 shape axes are checked before allocation, and the EOS ids follow generation_config.
• The image processor (Qwen2VL resize, patch 14, merge 2, bicubic, rescale and checkpoint mean and std) reproduces your four synthetic pixel hashes.

Before it lands:

1. The publisher's source, pinned. Please share modeling_dots_ocr.py, modeling_dots_vision.py, the configuration and the processor/template assets at revision e539fbb52, with hashes. Your oracle can record that source manifest. Two questions belong in that comparison:
   • The native tower calls full attention at every block, but the body says "window and full attention" (families/dots_mocr/hip_vision.zig:26,58). Which does the publisher's tower use?
   • The patch projection rounds to BF16 before adding the patch bias, while the decoder and merger projections add the bias first. Does the publisher round the same way?
2. Compact device receipts. We need candidate and base binary identities, checkpoint and asset hashes, versions and request settings. We also need the native and HF prompt and output ids with EOS termination, plus hashes for plain against drafted, fresh against resumed, and solo against concurrent. Add the seeded cells, cancellation and rejection, and the contract and limit outputs.
3. One small correction: the body says the registry test checks every family's reasoning capability, but it checks two anonymous namespaces. Please narrow the wording or add the per-family assertion.

Ordering: the body says the last seven commits depend on #1929, #1930 and #1933. #1929 and #1930 are on main now, and #1933 lands shortly. Once it's in, a rebase leaves only this family's commits.

[laudney]

Thanks @ashhart. All three points are answered below, and the code changes are two new commits on top (a3f797c49, 459512e31), with no force-push.

Request	Answer
Pinned publisher source	The oracle now holds the SHA-256 of all 14 non-weight sources at e539fbb52 and refuses a checkpoint that differs. python3 tools/zig/dots_mocr_reference.py sources <checkpoint> checks a checkpoint and prints the manifest without a GPU or model load.
Window or full attention?	Full attention for each image. The publisher's tower has no window attention: every block calls attention over cu_seqlens with one span per image and is_causal: false (modeling_dots_vision.py, DotsVisionTransformer.forward). The native tower does the same. "Window and full attention" in the body was wrong. The body now says "full attention for each image". The "windows" in the gates are decoder verification windows, not the tower.
Patch bias rounding	The publisher's run rounds the same way. DotsPatchEmbed.proj is an nn.Conv2d with bias. In the pinned HF run on PyTorch ROCm, the patch output matches a BF16 convolution output with the bias added afterwards. nn.Linear adds the bias in the GEMM epilogue before rounding. We measured both orders on the demo image: with the bias added before rounding, about 7.9 million of 29,982,720 patch values differed from HF's; with HF's order, 994 differ, the largest by 0.015625. A projection device test has a cancellation case that tells the two orders apart.
Device receipts	Below, from a new build of the PR head on one R9700.
Registry test wording	a3f797c49 adds try std.testing.expect(!reasoning(dots.native)) to the dots refusal test. It asserts only dots.mocr, so GLM-OCR and Parse do not conflict with it. The body now says this.

Pinned sources at e539fbb52 (SHA-256)

File	SHA-256
README.md	381645b2bd63c2adff9d0c266f79e86b25657394c48c35e68551f1da73178259
chat_template.json	c8d629ba2b3133376f1e64b980f82438945741fba5334db768afb281af93c4a9
config.json	41d5c9aad72c730bb97062cab7a59cdb1d427eeec0b7a6e2f95e1ee72ad3b498
configuration_dots.py	aa85add6cdaedfd947d25b850410de4506fc299dc55140d0cc5352c9737151fe
generation_config.json	d9ca5bd51bbf3b3b2b418f2b2d296cb4286ef2f02dc76dd7361d82ffd9d4e7ca
merges.txt	599bab54075088774b1733fde865d5bd747cbcc7a547c5bc12610e874e26f5e3
model.safetensors.index.json	0bb4faedeb3cc8b43c71fe1db6a761ea6557d4ab8ee3f54519e641078ded1c8a
modeling_dots_ocr.py	274b5bfa35f624f15001da70414e733b89f5cb44b5453d80023bff26a68d0135
modeling_dots_vision.py	f22d6251a05a03c7510d329ba1be9e959438cd54b32204eb51ccb1b56392a96f
preprocessor_config.json	fec8b3187ab0ae340f6d0b1f6f892a53c63b73b328836c83dd7bff1d5bc4c79a
special_tokens_map.json	92b060f39417514eed2f7652f454a3d5afb0e4a57b9f4b51dd4d1dfd8fdb9736
tokenizer.json	386545eb05f08c51352cde2fcc2c867f1592bb330f305efd1c6a57a93b1244cd
tokenizer_config.json	4f09a26dd6b20c82e9678d96b7363a58a778a948a21261d3315a949d47ed4cf4
vocab.json	ca10d7e9fb3ed18575dd1e277a2579c16d108e32f27439684afa0e10b1440910

Device receipt (PR head 459512e31, 2026-11-02)

One Radeon AI PRO R9700 (gfx1201), selected by ROCR_VISIBLE_DEVICES and checked by PCI address; other cards were busy with unrelated services. HF ids come from the existing oracle runs for the demo, square and two-image cases; the oracle was not re-run. A = sha256 of comma-joined ASCII ids (the server's token_sha256). R = sha256 of little-endian u32 ids. Only the first 16 hex digits are shown; full values are in receipt.json.

identity	value
candidate	459512e31e51 (PR head); tensorfold-native sha256 89a6d579d596cfd173245f0714ee25b20465803b427eeed7dbe3e223c47bdc3e
base	71932143f95a (upstream/main); tensorfold-native sha256 2e17004a34f084c3e2d1f84557cdb0c5f2ede94ff806913e94acf974b9806e33
build (both)	zig build native-hip -Dhip=true -Dhip-arch=gfx1201 -Dhipcc=/opt/rocm/core/bin/hipcc -Dhip-device-lib=/opt/rocm/core/lib/llvm/amdgcn/bitcode -Doptimize=ReleaseFast -j8
toolchain	Zig 0.17.0; ROCm core 10.1.0; hipcc/HIP 7.16.26385; AMD clang 24.0.0git
libraries the server loaded	rocBLAS librocblas.so.5.7 75645cbd1f2dd32f; HIP libamdhip64.so.7.16.26385 bddbf3553ee64ae5; libhipblaslt.so.1.4 3d414626f84ecc43; libhsa-runtime64.so.1.21.0 0457249a68424615; libjpeg-turbo 3.1.4.1 (official .deb 3.1.4.1-20260327, via TENSORFOLD_LIBTURBOJPEG) b6a5f4e6066da95c; libwebp 1.6.0 decoder a9d75f4997d79878
checkpoint	dots-studio/dots.mocr @ e539fbb52280393adc081b289ec597430a0f9031; shard 1 d8f4bc75340279da003609fe93f2eb02cc1a77087f5dfb6ba46c0980e1b4da81, shard 2 0101b36fe6620ba135d4bb8efbbf275fc27b4363b10b7f632c058c0955f3dc4d (equal to the HF LFS ids)
dots_mocr_reference.py sources	14 files, e.g. config 41d5c9aad72c730b, modeling_dots_ocr.py 274b5bfa35f624f1, modeling_dots_vision.py f22d6251a05a03c7, tokenizer.json 386545eb05f08c51, chat_template.json c8d629ba2b313337; all 9 files in the HF run manifests are equal
assets	demo_image1.jpg (rednote-hilab/dots.mocr @ 23f3e561) 90345584ccc2c4a8; square.png 56f28596c8b92fcc; wide.png e8f59ece1a61661b
HF reference	torch 2.10.0+git8514f05, transformers 4.57.6, flash-attn 2.8.3 (Triton AMD backend), qwen-vl-utils 0.0.10, HIP 7.2.53211; greedy decoding, flash_attention_2; max_new_tokens 24000 (demo) and 128 (square, two-image)
server	serve <ckpt> --host 127.0.0.1 --port 18896 --name dots-mocr --vision --context 16384 --parallel 2 --max-tokens 8192 --temperature 0 --no-update-check (recipe flags); v1.2.1, 10,693,404,224 tracked device bytes; SIGINT exit 0
requests	model-card prompt, data URLs, temperature 0, return_token_ids: true; max_tokens 64 / 8192 (demo), 128 (square, two-image)

check	native	HF / expected	equal
before: base serve, same flags	exit 1: the native HIP engine has no backend for dots_ocr checkpoints yet; serve with --engine python	refusal before weight load (device memory unchanged)	n/a
demo prompt /tokenize, 5096 ids	A e523420cba1e851e R feef1012bc620281	same	yes
square prompt, 280 ids / two-image prompt, 330 ids	A c1386e1fe8f59690 / A b4dee8bffad44e22	same / same	yes
demo, full, plain, 4634 ids, last 151673 (EOS), stop	A baddbe15d181fcad R 25ea393e611c83c5	same	yes
demo, full, drafted (1288/3280 accepted)	A baddbe15d181fcad R 25ea393e611c83c5	same	yes
two-image, plain, 108 ids, last 151673 (EOS)	A 8e9c4f1aa6a208d6 R a402ec9a718f471d	same	yes
square, plain, 40 ids, last 151673 (EOS)	A f6acbea11c225fc6, bbox [15,44,131,64]	A c2b79c3c8d0da748, bbox [16,45,130,63]	no: 4 bbox ids differ, text ALPHA 123 is equal (a known difference that the PR discloses)
demo, 64 ids: fresh (cached 0) / resumed (cached 4864) / drafted	A d8c09779b3ce16b8 ×3	HF first 64: A d8c09779b3ce16b8	yes
two-image: fresh (0) / resumed (256) / drafted (9/28 accepted)	A 8e9c4f1aa6a208d6 ×3	solo plain	yes
square: fresh (0) / resumed (256) / drafted	A f6acbea11c225fc6 ×3	solo plain	yes
concurrent: demo 64 plain ‖ two-image drafted	A d8c09779b3ce16b8 / A 8e9c4f1aa6a208d6	solo	yes
seeded sampling, two-image, T 0.7, top-k 32, top-p 0.9, seed 9281, 128 ids: solo plain, solo drafted (36/41), concurrent plain ‖ concurrent drafted	A 2aa5a1927fd3f35d R 9824c9bc8d76d17b ×4	equal across all 4 cells; differs from greedy	yes
client closes in prefill (0 output tokens, client-left), then square	A f6acbea11c225fc6	solo square	yes
client closes in decode (1 output token, client-left), then square	A f6acbea11c225fc6	solo square	yes
invalid image bytes	400 image 1 is not PNG, JPEG or WebP; send one of those	400	yes
excess context (max_tokens 50000)	400 context_length_exceeded	400	yes
video_url part	400 content parts must be text or image_url; audio and video inputs are unsupported	400	yes

In all replies, the OCR text is in message.content and there is no reasoning_content. The 64-id demo and square drafted requests made zero copy proposals, so their drafted = plain result is trivial. The full demo, two-image and seeded rows do use drafts.

Kernel contract and limits (PR head kernels; the two new commits change no kernel)

	upstream/main	this PR
gfx1201 contract	52 kernels pass gates 1-4, 5/5 controls red	63 kernels pass gates 1-4, 9/9 controls red
Limit problems (live gfx1201 / offline, 4 chips)	124 / 503	124 / 503, the same entries

Ordering: #1929 and #1930 are now on main as your versions, so this branch shows conflicts with main in build.zig, zig/build/hip.zig, zig/src/core/prompt_rules.zig and zig/src/hip/root.zig. They are the copies of those two PRs. #1933 is merged too, so a rebase onto main now leaves only the dots.mocr commits. We do not force-push open PRs. If you prefer, we can push a merge of main instead.

[ashhart]

Thanks @laudney: all three asks are answered (the oracle pins the revision and the SHA-256 of all 14 non-weight sources, refuses a differing checkpoint, and has a GPU-free sources command). Please go ahead with the rebase you offered. Its four conflicting files are the copies of #1929/#1930, which are on main now, so dropping those commits should leave only the dots.mocr ones. We'll land it then. #1940 is landing ahead of it, so rebase after that if you can.

[laudney]

Thanks, Ash. I will rebase this branch onto main as soon as #1940 lands: the copies of #1929/#1930 drop out, only the dots.mocr commits stay, and I re-run the dots host suites and a device smoke before I push. I will post here when the rebased head is up.

[laudney]

Rebased on main c7952f398 (after #1940), as you asked. Only the dots.mocr commits are left, and the device record on the rebased head equals the earlier receipt hash for hash.

	before	after
head	459512e31	7e5b5c4f0
base	71932143f plus copies of #1929, #1930, #1933	main c7952f398
commits	14	9, all dots
files	75	64
mergeable	conflicting	yes

Conflicts. Each one is a file where #1940 and this PR both add an entry, and both entries are kept:

file	resolution
build.zig, zig/build/hip.zig	GLM-OCR's steps and modules, then dots.mocr's
hip_records.zig, hip/root.zig, tests/hip/contract_main.zig	GLM-OCR's records, ops and contract kernels, then dots.mocr's
native/hip.zig	registry …, glm_ocr.native, dots.native; flagRefused keeps GLM's tp_refusal/mtp_refusal and dots' refusal hook; openWith keeps #1940's loaded.cache block and dots' attach hook (dots' Loaded has no vision or cache field, so #1940's path gives it nothing and attach sets its offers as before)
core/lane_host_warm.zig	#1962's keepBornFrame (exact image identities) is now gated by dots' keepsEnd

Checks on 7e5b5c4f0: zig fmt --check; zig build test without and with -Dhip=true -Dhip-arch=gfx1201; host suites test-dots-native-host, test-glm-ocr-native-host, test-hip-family-offers-host, test-hip-native-options-host, test-draft-hip-host, test-lane-api-host, test-qwen35-hip-host, test-hip-blas-host; zig build native-hip … -Doptimize=ReleaseFast; lean 424 problems, equal to main.

Device record (gfx1201, Radeon AI PRO R9700): the same driver and server flags as the receipt in this thread, on binary sha256 480519f455952411f7863af5142fafedd6e30235a512a1b32a540c966ead477d.

check	result
prompt ids: demo 5096, square 280, two-image 330	equal to HF
demo full, plain and drafted, 4634 ids to EOS	A baddbe15d181fcad, equal to HF
demo 64: fresh / resumed (cached 4864) / drafted	A d8c09779b3ce16b8 ×3, equal to HF
two-image: plain 108 ids to EOS / resumed (256) / drafted	A 8e9c4f1aa6a208d6 ×3, equal to HF
square: fresh / resumed (256) / drafted	A f6acbea11c225fc6 ×3; vs HF the same 4 bbox ids differ as before (disclosed)
concurrent demo 64 plain ‖ two-image drafted	equal to solo
seeded (T 0.7, top-k 32, top-p 0.9, seed 9281): solo/concurrent × plain/drafted	A 2aa5a1927fd3f35d ×4
cancel in prefill / in decode, then square	A f6acbea11c225fc6
invalid image, excess context, video_url	400 each

25 of 26 checks pass; the one failure is the known square bbox difference, with the same positions and hashes as before the rebase.

[ashhart]

Thanks @laudney. The rebase is clean: 9 commits, all dots, on current main, and each resolution keeps both GLM-OCR's
entry and yours. We checked that it applies on our main with the lean check at 424. Your gfx1201 record on the
rebased head matches the earlier receipt hash for hash:

• the demo's 4,634 ids to EOS equal Transformers', plain and drafted;
• the two-image case's 108 ids equal Transformers';
• resumed, concurrent and seeded runs equal their solo plain runs;
• the only difference is the disclosed four box ids on the square page.

With the pinned sources, the attention and bias-order answers, and the registry assertion already in, everything we asked for is done, so this is ready to land as it is.

One heads-up for #1941: both PRs add ops/vision.hpp and hip_records_vision.zig, so #1941 will need the rebase you offered once this is in. When the two vision sets merge, one LayerNorm kernel may be enough. If the two round differently, keep both and note why.

[ashhart]

Thank you @laudney. This is merged into main with your commits under your own name, so you're listed as a TensorFold contributor. It ships in the next release, credited to you in the changelog and the release notes.

[ashhart]

Thanks @laudney! This is in TensorFold 1.2.3, out now: https://github.com/ashhart/TensorFold/releases/tag/v1.2.3