Skip to content

GLM-OCR on HIP: the publisher's BF16 checkpoint, exact against its six model-card tasks - #1940

Merged
ashhart merged 13 commits into
ashhart:mainfrom
laudney:pr-ocr-glm-ocr
Nov 3, 2026
Merged

ashhart merged 13 commits into
ashhart:mainfrom
laudney:pr-ocr-glm-ocr

Conversation

@laudney

@laudney laudney commented Nov 2, 2026 •

Copy link
Copy Markdown
Contributor

Part of #1927.

This serves the publisher's BF16 zai-org/GLM-OCR (revision 2e85a62840ccac27daa451df36c736c4636b8628) natively on HIP. The CogViT image encoder and the GLM decoder run in Zig and HIP, with no Python or libtorch in the served engine. 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. Every reply is checked against the model card's own Transformers example.

Before (main) After (this PR)
zai-org/GLM-OCR (BF16) on HIP not supported served natively on Radeon AI PRO R9700 (gfx1201)
Python or libtorch in the served engine — none; one pinned Transformers 5.5.4 oracle under tools/zig/ for checks
Six model-card tasks vs HF (recognition, description, text, formula, table, extraction) — exact reference IDs and text over HTTP, plain and copy-drafted
Concurrent pairs vs solo — identical IDs
Batched verification rows (widths 1-32) vs single-row projections — 23,519,232 BF16 values, all exact
Strict structured extraction — every schema field, with drafts off and on, identical IDs
Kernel contract on gfx1201 52 kernels pass gates 1-4 70 kernels pass gates 1-4, 5/5 controls red
Inherited limit problems (live gfx1201 / offline, 4 chips) 124 / 503 124 / 503, the same entries (0 new)
Existing models — unchanged: new image budgets, runtime libraries and cache rules are opt-in

Prerequisites #1929 (prompt-cache rules), #1930 (rocBLAS projections) and #1933 (no-reasoning replies) are merged. This branch still carries their earlier copies at its base, so a rebase onto main leaves only the GLM-OCR commits. The GLM-OCR review is the last ten commits, 4e64877cb^..a8594af84.

Commits

  1. Checkpoint and image preparation: config, the 526-tensor BF16 schema with refusals for missing, extra or wrongly typed tensors, the image processor, and M-RoPE positions, with processor fixtures.
  2. BF16 weights on HIP and matrix products: checked weight ownership and the BF16 projections on HIP: rocBLAS BF16 projections with FP32 sums #1930's rocBLAS module.
  3. The pinned Transformers oracle, tools/zig/glm_ocr_reference.py. It has four commands:
    • capture runs the model card's example and saves the reference evidence.
    • processor writes the committed fixtures.
    • attention writes the SDPA fixtures for the AOTriton ABI test.
    • examples rebuilds the six model-card tasks.
  4. feat(picture): explicit source image budgets. picture.Limits and decodeWithLimits let a caller admit a full linked image before resizing. The default limits are unchanged, and the existing callers keep decode.
  5. The image encoder:
    • CogViT with MIOpen convolutions (hip/convolution.zig), restricted to one GEMM solver family, so its arithmetic does not change between searches.
    • AOTriton flash attention (hip/attention_aotriton.zig).
    • The pointwise kernels in kernels/hip/glm_ocr/pointwise.hip.
  6. Image and text decoding on HIP, with the fixed-algorithm hipBLASLt decode path (hip/blas_fixed.zig).
  7. refactor(hip): share native open and image offers between Qwen3.5 and GLM-OCR. Qwen3.5 behaviour does not change.
  8. Batched exact verification: the vocabulary rows of a window are projected on the GPU with the single-row algorithm.
  9. Serving exact image replies: registration and image offers, and pub const reasoning = false; (Server: one engine flag for models whose replies carry no reasoning block #1933), because GLM-OCR writes no reasoning block. Its prompt still follows enable_thinking.
  10. Prompt-state reuse with exact arithmetic: image repeats are keyed through Rules.repeat_keyed (Prompt cache: an opt-in rule for families to count keyed image prompts as repeats #1929). This commit also adds the recipe docs/recipes/glm-ocr-hip.md.

Runtime libraries (optional, loaded at run time)

Library Used for Notes
rocBLAS (#1930) BF16 projections
hipBLASLt fixed decode algorithm libhipblaslt.so.1
MIOpen image convolutions keep its find database on persistent storage
AOTriton 0.11.1 flash attention MIT; LICENSES/AOTriton-MIT.txt and THIRD_PARTY_NOTICES.md are added; TENSORFOLD_AOTRITON selects it
libjpeg-turbo 3.1.4.1 JPEG decode equal to the reference the repository's pinned decoder

Known limits

Gates

  • Kernel contract on a Radeon AI PRO R9700 (gfx1201): SUMMARY gfx1201 70 kernels: 70 pass gates 1-4, 0 fail, 2 known findings; 5 controls: 5 red as planted, 0 missed; 124 limit problems.
    • The upstream/main base reports 52 kernels ... 124 limit problems on the same GPU, with the same limit entries.
    • All 48 GLM-OCR record combinations of kernel and chip match the code-object metadata.
  • Host: test-hip-blas-host, test-glm-ocr-host, test-glm-ocr-hip-host, test-glm-ocr-native-host and test-lane-api-host pass. zig build test: 470/470 steps. zig build test -Dhip=true: 502/502 steps.
  • Commits: every commit builds native-hip for gfx1201.
  • Lean checker: zero new problems over the base.

@ashhart

ashhart commented Nov 2, 2026

Copy link
Copy Markdown
Owner

Thanks @laudney. We reviewed this independently, and it lands. The shared changes keep today's defaults: Info.reasoning true, the picture.Limits defaults, repeat_keyed false, hip_serve.Options.vision null, and the new refusals only for families that declare them. The host suites pass under -Dhip=true (485 passed, 1 skipped). We have no AMD GPU, so the kernel contract, the gfx1201 reference ids and the six model-card cases are your receipts. It's gated and queued, and we've written the changelog line.

@laudney

laudney commented Nov 2, 2026

Copy link
Copy Markdown
Contributor Author

The #1938 review asked dots.mocr for pinned sources, a device receipt and a per-family reasoning assertion. This PR gets the same now, in two new commits on top (e0546d811, 1aab5dd64), with no force-push.

Item Before Now
Publisher sources the oracle checked the checkpoint revision only the SHA-256 of the 8 non-weight sources at 2e85a6284 (configuration, generation and processor configs, tokenizer, chat template, model card) and of the 5 Transformers 5.5.4 files that hold GLM-OCR's model and processor code, refused on any difference; python3 tools/zig/glm_ocr_reference.py sources <checkpoint> checks them without a GPU or a model load
Reasoning assertion the registry test from #1933 names no family the GLM-OCR flag test asserts !reasoning(glm_ocr.native)
Device receipt the PR body's account below, from a new build of the PR head
capture receipt revision and runtime also the pinned source hashes and runtime versions

GLM-OCR's model code is not in the checkpoint: Transformers 5.5.4 ships it. So the oracle pins those library files by path inside the installed package, which it finds without importing Transformers.

Pinned sources (SHA-256)
File SHA-256
README.md 172334449c3df98650bc9a42781c254c64c9d466b440e2ca8ce31f99829212b2
chat_template.jinja 062e7ee4cc8defa88a5938b5d456dc60366ffd80647f918946ee747bf09ddc7c
config.json ad537817b2becde827cbb23569a1dd83a352ee64cc34ce751a4ef000ecce4b5f
generation_config.json 4e2a3412f65b2a21d315d928986c55723d3e60dfb92c3982a5b9fd56835b0aa5
preprocessor_config.json 02cc50c36240882ae35e8cd4077a25a379664108185d728d261cb785aefeccff
processor_config.json a56fbd559186fefd6c22a3acc6a393a18a9fe52c264692a122a143d24042112b
tokenizer.json aa0fd058c73a5718bb191f6672dc16d122ee0147b20c123d1726514298f9968a
tokenizer_config.json 95abd49942be9cdc6e2653528aed7b4c1cb9e180828421d3488ab38da3057016
transformers/models/glm_ocr/configuration_glm_ocr.py 7015c0aa18811ab058e228a9c343bc07bf266abb2753a66a2c371720b6612c53
transformers/models/glm_ocr/modeling_glm_ocr.py f667126388eacde9ca93983c1fc34488f4c169717a6ff5900c1479d2c6bb7db4
transformers/models/glm46v/image_processing_glm46v.py 41fa88e258a4b291f167d77085eee989dbebff1776c336121e2f4185b6e0eef1
transformers/models/glm46v/image_processing_pil_glm46v.py 3506c3f76284c90d3ed9d4e44e5091cb09ceb4fcd8f382cb2a90aa18923c7598
transformers/models/glm46v/processing_glm46v.py 91ba14d315a2059e4f0f811618394f201401a9858648c1f66c4637f25fe628db

Device receipt

One Radeon AI PRO R9700 (gfx1201), selected by ROCR_VISIBLE_DEVICES and checked by PCI address (0000:c3:00.0); other cards ran unrelated work. HF ids come from the existing Transformers 5.5.4 oracle runs; 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 16 hex digits are shown; we can attach the full values on request.

identity value
candidate 1aab5dd648a6 (PR head); tensorfold-native 1.2.1 sha256 1c09d827c9cef97cd55dd165e183a516be8a77e94bca9de32435ce98152e5a93
base 71932143f95a (upstream/main); tensorfold-native 1.2.1 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 10.1.0 (rocm-gfx120x); hipcc/HIP 7.16.26385; AMD clang 24.0.0git
libraries the server loaded rocBLAS librocblas.so.5.7 75645cbd1f2dd32f; libhipblaslt.so.1.4 3d414626f84ecc43; MIOpen 3.6.1 libMIOpen.so.1.0 0549aa0661734db6; HIP libamdhip64.so.7.16.26385 bddbf3553ee64ae5; AOTriton 0.11.1 libaotriton_v2.so.0.11.1 e0638806efa5d35c + gfx120x images (1888 files, tree f01f1ecb70206d4f); libjpeg-turbo 3.1.4.1 b6a5f4e6066da95c; libtfgrammar (xgrammar v0.2.8 97787376) 0a8a1c3f27c4eef5
checkpoint zai-org/GLM-OCR @ 2e85a62840ccac27daa451df36c736c4636b8628; model.safetensors (2,650,579,464 bytes) a16eb0de98d199293371c560f95f83130d2a2c9612449df16839f08ff9498815 (equal to the HF LFS id)
glm_ocr_reference.py sources pins pass; config ad537817b2becde8, tokenizer.json aa0fd058c73a5718, chat_template.jinja 062e7ee4cc8defa8, preprocessor 02cc50c36240882a; Transformers 5.5.4 modeling_glm_ocr.py f667126388eacde9, processing_glm46v.py 91ba14d315a2059e (config and modeling equal in all 6 HF receipts)
assets car 8cd28b372908bcaa, bee 8b21ba78250f852c (documentation-images @ 9dd8e471, original bytes); text 32b9c78df373c021, formula 7a91aa1b95543f76, table 6b9f3dc2af058e12, extraction 501a86f14f69aaa7 (fictional fixtures for the model-card prompts)
HF reference Transformers 5.5.4, torch 2.10.0+git8514f05, HIP 7.2.53211, MIOpen 3.6.1, SDPA, bfloat16, greedy, same GPU
server serve <ckpt> --host 127.0.0.1 --port 18897 --name glm-ocr --vision --context 8192 --parallel 2 --prompt-cache-gib 1 --mtp-drafts 0 --max-tokens 512 --temperature 0 --no-update-check --idle-unload 0 (recipe flags); a second process with --prompt-cache-gib 0 for the cache-off rows; both SIGINT exit 0
requests model-card prompts, data URLs, return_token_ids: true, temperature 0, max_tokens 512 (bee 128)
check native HF / expected equal
before: base serve, same flags exit 1: the native HIP engine has no backend for glm_ocr checkpoints yet; serve with --engine python refusal before weight load n/a
car: prompt 403 ids; output 7 ids, last 59253 (EOS), stop prompt A 148bff80c05084e5; out A 93aec362c628084a R 02cc86b94f00562d prompt A 148bff80c05084e5; out A 93aec362c628084a R 02cc86b94f00562d yes
bee: prompt 6156 ids; output 6 ids, last 59253 (EOS), stop prompt A 5f659154f6e7b7da; out A c8698f4784ceb1f5 R be2ba128f355f4e1 prompt A 5f659154f6e7b7da; out A c8698f4784ceb1f5 R be2ba128f355f4e1 yes
text: prompt 252 ids; output 26 ids, last 59253 (EOS), stop prompt A 17affe56723c83e3; out A db7213f5ace9e43a R 8eeebd358f0af890 prompt A 17affe56723c83e3; out A db7213f5ace9e43a R 8eeebd358f0af890 yes
formula: prompt 223 ids; output 21 ids, last 59253 (EOS), stop prompt A 54954047df4014b4; out A cec8bf3afa2b36ee R c5a936f8c02793e3 prompt A 54954047df4014b4; out A cec8bf3afa2b36ee R c5a936f8c02793e3 yes
table: prompt 192 ids; output 58 ids, last 59253 (EOS), stop prompt A f6712bfa91c76445; out A b85de97f46b59f52 R 6f2f5b715ffb12cc prompt A f6712bfa91c76445; out A b85de97f46b59f52 R 6f2f5b715ffb12cc yes
extraction: prompt 850 ids; output 155 ids, last 59253 (EOS), stop prompt A ed83816e526add68; out A f12b92635ade64f0 R d081e35d955f873b prompt A ed83816e526add68; out A f12b92635ade64f0 R d081e35d955f873b yes
all 6, cache on: fresh plain (cached 0) ids = HF ×6 HF yes
all 6, cache on: resumed plain / resumed drafted (7/110 accepted); cached car 401, bee 6154, text 250, formula 221, table 190, extraction 848 ids = HF ×12 HF yes
all 6, cache off: fresh plain / fresh drafted (7/110 accepted) ids = HF ×12 HF yes
concurrent pairs table‖extraction, car‖bee, text‖formula: plain / drafted (7/109 accepted); log shared verify: 2 streams ids = HF ×12 HF yes
seeded sampling text, T 0.7, top-k 32, top-p 0.9, seed 9281: solo, repeat, drafted, concurrent plain, concurrent drafted; 26 ids, stop A edeac8e219d80693 R f6d9d06c30c65f41 ×5 equal across 5 cells; differs from greedy yes
seeded sampling formula, T 0.7, top-k 32, top-p 0.9, seed 9281: solo, repeat, drafted, concurrent plain, concurrent drafted; 21 ids, stop A cec8bf3afa2b36ee R c5a936f8c02793e3 ×5 equal across 5 cells; equals greedy yes
seeded sampling table, T 0.7, top-k 32, top-p 0.9, seed 9281: solo, repeat, drafted, concurrent plain, concurrent drafted; 58 ids, stop A b85de97f46b59f52 R 6f2f5b715ffb12cc ×5 equal across 5 cells; equals greedy yes
seeded sampling extraction, T 0.7, top-k 32, top-p 0.9, seed 9281: solo, repeat, drafted, concurrent plain, concurrent drafted; 155 ids, stop A f12b92635ade64f0 R d081e35d955f873b ×5 equal across 5 cells; equals greedy yes
strict json_schema extraction (xgrammar): plain / drafted (26/106 accepted), 151 ids strict JSON, all 11 fields; A aa6e81dd24c5816d / aa6e81dd24c5816d fixture fields; plain = drafted yes
client closes at 0.25 s, before engine submit (body/JPEG decode); then a background car probe and car no ended line, no engine work; probe 0.19 s (before cut 0.19 s); car = HF dropped, no leaked admission count yes
client closes in prefill after the stream opens (fresh bee bytes), then car client-left, tokens 0; car = HF 0 tokens, exact next yes
client closes in decode after 8 answer events, then car client-left, tokens 8; car = HF exact next yes
context edge: 403 prompt + 7,789 max_tokens / + 7,790 200 (car = HF) / 400 context_length_exceeded 200 / 400 yes
invalid image bytes / truncated JPEG 400 image 1 is not PNG, JPEG or WebP; send one of those / 400 image 1 could not be decoded; send PNG, JPEG or WebP 400 / 400 yes
video_url part, then car 400 content parts must be text or image_url; audio and video inputs are unsupported; car = HF 400 yes

In all replies, the OCR text is in message.content, with 0 reasoning tokens. Copy drafts made proposals only for extraction, table; for bee, car, formula, text the drafted = plain result is trivial. Seeded sampling changes only text; the other sampled replies equal greedy, which the peaked OCR distributions explain. The extraction reply keeps the model's Markdown fence, as HF does; only the json_schema request returns strict JSON.

Kernel contract and limits (the two new commits change no kernel)

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

Ordering: #1929, #1930 and #1933 are now on main as your versions, so this branch shows conflicts with main in build.zig, zig/build/hip.zig, zig/src/hip/root.zig, zig/src/native/hip.zig, zig/src/platform/optional_libraries.zig and zig/tests/hip/blas_test.zig. They are the copies of those PRs. A rebase onto main leaves only the GLM-OCR commits. We do not force-push open PRs. If you prefer, we can push a merge of main instead.

ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 2, 2026
ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 2, 2026
ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 2, 2026
ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 3, 2026
ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 3, 2026
@ashhart
ashhart force-pushed the pr-ocr-glm-ocr branch 2 times, most recently from 6010aca to 474ef4a Compare November 3, 2026 11:43
ashhart added a commit to laudney/TensorFold that referenced this pull request Nov 3, 2026
GLM-OCR support starts with the checks that need no GPU. A new
glm_ocr family admits the BF16 checkpoint graph, reads its model and
processor settings, and prepares grouped temporal image patches with
bounded geometry and rotary coordinates before any device memory is
allocated. Host tests compare that preparation with fixtures taken
from the Transformers processor, and test-glm-ocr-host and
test-glm-ocr-fast run them as part of the default test step.
The native engine needs the checkpoint on the GPU in a form the
shared rocBLAS projections can read directly. GLM-OCR now uploads its
tensors unchanged into one admitted HIP allocation, finds them by
name, and refuses a checkpoint that does not fit before it allocates.
Projections reuse hip.blas, so each dot sums in FP32 and rounds to
BF16 once, with bias rows added before that rounding.
test-glm-ocr-hip-host covers ownership and refusal on the host, and
build-glm-ocr-weights-test builds the device upload check.
Native GLM-OCR checks compare against the official Transformers model
at a pinned checkpoint and library version, so the reference must be
reproducible. tools/zig/glm_ocr_reference.py has four commands:
capture runs the model on one HIP GPU and keeps the tensors, tokens
and runtime versions that native comparisons read; processor writes
the processor and RoPE fixtures without a GPU; attention captures
SDPA outputs for the AOTriton checks; and examples prepares the
model-card tasks and linked images. Each command imports its heavy
packages only when it runs.
The image codecs apply one shared source size limit, but GLM-OCR must
decode a full linked page image unchanged before it resizes it. A new
picture.Limits budget and decodeWithLimits entry points let a caller
set its own side and pixel limits, which the PNG, JPEG and WebP
decoders check before they allocate pixels; decode keeps the shared
defaults and the decompression bomb bound still applies. The GLM-OCR
processor passes its own budget, and a new test-picture-host step
covers the caller limits.
GLM-OCR's CogViT vision tower and learned connector now run on HIP
with owned scratch memory: MIOpen runs the patch convolution, rocBLAS
the projections, a standalone AOTriton forward kernel the attention,
and new pointwise kernels the norms, rotary embeddings, activations
and layout steps. The convolution always takes MIOpen's GEMM
solution from the returned candidates, so GPU timing cannot change
its reduction order between requests. The new kernels join the HIP
contract registry with gates and launch records, device tests compare
patch, attention and encoder stages with the reference, and the
AOTriton license joins the third-party notices.
The text decoder now runs on HIP from native image and token rows. A
new hip.blas_fixed wrapper keeps one hipBLASLt single-row algorithm
for each projection shape, so a row reduces in the same order at
every request width, and KV restore, shared decoding and
verification windows keep the reference arithmetic. Device tests
compare decoder states and tokens with the captured references, and
the shared BLAS device test now also checks the fixed path against
exact integer dots.
Native open settings lived in the Qwen3.5 HIP engine, and only a
family's name could tell the lane host that it accepts images. The
settings move into the shared hip_serve layer with an optional vision
field, and the native registry passes a loaded family's image offer
to the lane host, so another native HIP family can open with the same
settings and accept images without a special case.
test-hip-native-options-host and test-hip-family-offers-host cover
both paths without a device.
Draft verification needs logits for several rows at once. The decoder
now keeps up to 32 rows of verification logits on the GPU for the
shared drawer and computes every row with the same single-row
vocabulary product, so batched verification returns the bits that
one-row decoding returns. The operators can also load on a chosen
device, and the device checks compare batched logits with single-row
ones and prompt cuts, including cuts inside image rows, with a
complete prefill.
GLM-OCR now registers as a native HIP family and serves chat requests
through the shared lane core, GPU drawer and copy drafts, with
bounded image preparation and admission against live device memory.
The model never writes a reasoning block although its chat template
has an enable_thinking switch, so the family declares reasoning false
and the server keeps every reply in content while the prompt still
follows enable_thinking. The family refuses MTP drafts and tensor
ranks before it loads, and test-glm-ocr-native-host covers those
refusals.
OCR clients often send the same page image again, and a fresh
prefill of its image rows repeats the costliest prompt work. GLM-OCR
prompt snapshots now keep the full prompt arithmetic shape and one
keyed prepared image row, so a resumed prompt computes the same bytes
as a complete prefill and keeps its image identity. The family opts
in to the shared repeat_keyed cache rule, so a repeated image request
earns a repeat cut, and a device check compares resumed decoding with
a complete prefill. docs/recipes/glm-ocr-hip.md shows how to build,
serve and check the model.
The registry test from ashhart#1933 checks the hook on two anonymous
namespaces, so it does not name GLM-OCR. The GLM-OCR flag test now
asserts that the registered family declares no reasoning block, which
keeps its OCR text in message content. Only this family's assertion
is added, so the other OCR families do not conflict with it.
laudney and others added 2 commits November 3, 2026 12:11
The oracle hashed config.json and the model class file but accepted
any checkpoint or Transformers tree. It now holds the SHA-256 of every
non-weight checkpoint source at 2e85a6284 and of the GLM-OCR model,
configuration and processor files in Transformers 5.5.4, and refuses a
difference before a capture or fixture run. A new sources command
checks both trees and prints the manifest without loading the model.
@ashhart
ashhart merged commit c7952f3 into ashhart:main Nov 3, 2026
0 of 7 checks passed
@ashhart

ashhart commented Nov 3, 2026

Copy link
Copy Markdown
Owner

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

ashhart commented Nov 3, 2026

Copy link
Copy Markdown
Owner

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

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.

2 participants