Repository navigation
GLM-OCR on HIP: the publisher's BF16 checkpoint, exact against its six model-card tasks - #1940
Conversation
|
Thanks @laudney. We reviewed this independently, and it lands. The shared changes keep today's defaults: |
|
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 (
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)
Device receiptOne Radeon AI PRO R9700 (gfx1201), selected by
In all replies, the OCR text is in Kernel contract and limits (the two new commits change no kernel)
Ordering: #1929, #1930 and #1933 are now on main as your versions, so this branch shows conflicts with main in |
1aab5dd to
0a47e4d
Compare
0a47e4d to
6f83312
Compare
6f83312 to
7ac0ae5
Compare
7ac0ae5 to
133348b
Compare
6010aca to
474ef4a
Compare
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.
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.
474ef4a to
216a68e
Compare
|
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. |
|
Thanks @laudney! This is in TensorFold 1.2.2, out now: https://github.com/ashhart/TensorFold/releases/tag/v1.2.2 |
Part of #1927.
This serves the publisher's BF16
zai-org/GLM-OCR(revision2e85a62840ccac27daa451df36c736c4636b8628) 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.zai-org/GLM-OCR(BF16) on HIPtools/zig/for checksPrerequisites #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
tools/zig/glm_ocr_reference.py. It has four commands:captureruns the model card's example and saves the reference evidence.processorwrites the committed fixtures.attentionwrites the SDPA fixtures for the AOTriton ABI test.examplesrebuilds the six model-card tasks.feat(picture): explicit source image budgets.picture.LimitsanddecodeWithLimitslet a caller admit a full linked image before resizing. The default limits are unchanged, and the existing callers keepdecode.hip/convolution.zig), restricted to one GEMM solver family, so its arithmetic does not change between searches.hip/attention_aotriton.zig).kernels/hip/glm_ocr/pointwise.hip.hip/blas_fixed.zig).refactor(hip): share native open and image offers between Qwen3.5 and GLM-OCR. Qwen3.5 behaviour does not change.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 followsenable_thinking.Rules.repeat_keyed(Prompt cache: an opt-in rule for families to count keyed image prompts as repeats #1929). This commit also adds the recipedocs/recipes/glm-ocr-hip.md.Runtime libraries (optional, loaded at run time)
libhipblaslt.so.1LICENSES/AOTriton-MIT.txtandTHIRD_PARTY_NOTICES.mdare added;TENSORFOLD_AOTRITONselects itKnown limits
build.ziggains one line afterconst test_step. core: a GGUF reader and an exact repack of GGUF's quantized blocks (#1920, steps 1-2) #1926 and dots.mocr on HIP: the publisher's BF16 OCR checkpoint, token-exact against its model card #1938 add a line at the same place, so whichever lands later keeps both.Gates
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.upstream/mainbase reports52 kernels ... 124 limit problemson the same GPU, with the same limit entries.test-hip-blas-host,test-glm-ocr-host,test-glm-ocr-hip-host,test-glm-ocr-native-hostandtest-lane-api-hostpass.zig build test: 470/470 steps.zig build test -Dhip=true: 502/502 steps.native-hipfor gfx1201.