tests: batch_check_state runs in CI, on the runtime stub and a fabricated container (fixes #1946)

CMakeLists.txt

@@ -1512,6 +1512,52 @@ endif()  # ONEBIT_NPU_HOST
 # `llama-server` child (engine #1653 review). Device-free: it drives the helper against a synthetic wrapper.
 add_test(NAME serve_kill_tree COMMAND ${CMAKE_SOURCE_DIR}/tests/serve_kill_tree_test.sh)
 
+# The NPU forward library: one source list and one link interface, built against a runtime target named by the
+# caller. if(ONEBIT_NPU) builds onebit_npu_forward against the private runtime; a configure without it (CI)
+# builds onebit_npu_forward_ci against onnx/ep/npu_runtime_ci_stub.cpp, whose device calls throw -- enough for
+# the host-GEMM forward (init_host), which never opens the device (engine#1946).
+option(ONEBIT_NPU_DEBUG_DUMPS "NPU forward: env-named debug dumps (developer builds)" OFF)
+function(onebit_npu_forward_library name runtime)
+    add_library(${name} STATIC
+        # gemm_q8_raw / gemm_f32: the int8 GEMM the ONNX EP's QDQ leg uses, and now the embedding route's
+        # ternary linears on the array (engine#1311). Its own probe checks the accumulators are exact.
+        npu/ort/onebit_npu_gemm.cpp
+        npu/forward/src/model.c
+        npu/forward/src/q4nx_dequant.c
+        npu/forward/src/q4nx_pack.c
+        # The production weight packer, extracted from I8Ctx so it can be tested without the device runtime (#1416).
+        npu/forward/src/q4nx_pack_blocked.cpp
+        # The lm-head image cache (engine#1716): the file format, and the load-or-pack a serve calls per head tile.
+        npu/forward/src/q4nx_head_image.c
+        npu/forward/src/q4nx_head_tile_cache.cpp
+        npu/forward/src/q4nx_semantics.c
+        npu/forward/src/q4nx_forward.cpp
+        # The NPU target interface and this box's target (engine#799 P4): the model keeps the maths, the
+        # target keeps the runtime, the kernels and the weight residency.
+        npu/target/xdna/xdna_target.cpp)
+    # The Q4NX pack and dequant loops carry `#pragma omp parallel for`; without the
+    # OpenMP flags those compile to serial code. Wire it in (both C and C++ sources)
+    # so they actually use the cores.
+    find_package(OpenMP)
+    if(OpenMP_C_FOUND AND OpenMP_CXX_FOUND)
+        target_link_libraries(${name} PUBLIC OpenMP::OpenMP_C OpenMP::OpenMP_CXX)
+    endif()
+    # NPU_INFER_DUMP_ATTN / _LN / _HEAD write debug dumps to a path taken from the environment;
+    # only debug builds carry them (ONEBIT_NPU_DEBUG_DUMPS)
+    if(ONEBIT_NPU_DEBUG_DUMPS)
+        target_compile_definitions(${name} PRIVATE NPU_INFER_DEBUG_DUMPS)
+    endif()
+    target_include_directories(${name} SYSTEM PUBLIC
+        npu/forward/include npu/forward/engine-src npu ${CMAKE_CURRENT_SOURCE_DIR})
+    # plan.json (npu/plan/plan.h, engine #799 P1) is read with nlohmann's parser, and the header is
+    # part of this library's interface, so the dependency is PUBLIC.
+    # engine#1185: the forward converts an imported q4nx-2 container to the device's 5,120-byte
+    # q4_1 form before model_load (npu/q4nx_device.cpp, onebit_npu), so the host layer is part of
+    # this library's link interface.
+    target_link_libraries(${name} PUBLIC onebit_npu ${runtime} nlohmann_json::nlohmann_json dl
+                                                  pthread m)
+endfunction()
+
 # The lane itself: needs the NPU runtime and an XDNA NPU to run.
 option(ONEBIT_NPU "Build the NPU lane (needs the private NPU runtime: ONEBIT_NPU_RT_LIB or ONEBIT_NPU_RT_DIR)" OFF)
 # The NPU code compiles against one runtime (npu/rt.h, docs/npu.md "The runtime"): our own runtime,
@@ -1637,45 +1683,9 @@ if(ONEBIT_NPU)
     endif()
 endif()
 
-    add_library(onebit_npu_forward STATIC
-        # gemm_q8_raw / gemm_f32: the int8 GEMM the ONNX EP's QDQ leg uses, and now the embedding route's
-        # ternary linears on the array (engine#1311). Its own probe checks the accumulators are exact.
-        npu/ort/onebit_npu_gemm.cpp
-        npu/forward/src/model.c
-        npu/forward/src/q4nx_dequant.c
-        npu/forward/src/q4nx_pack.c
-        # The production weight packer, extracted from I8Ctx so it can be tested without the device runtime (#1416).
-        npu/forward/src/q4nx_pack_blocked.cpp
-        # The lm-head image cache (engine#1716): the file format, and the load-or-pack a serve calls per head tile.
-        npu/forward/src/q4nx_head_image.c
-        npu/forward/src/q4nx_head_tile_cache.cpp
-        npu/forward/src/q4nx_semantics.c
-        npu/forward/src/q4nx_forward.cpp
-        # The NPU target interface and this box's target (engine#799 P4): the model keeps the maths, the
-        # target keeps the runtime, the kernels and the weight residency.
-        npu/target/xdna/xdna_target.cpp)
-    # The Q4NX pack and dequant loops carry `#pragma omp parallel for`; without the
-    # OpenMP flags those compile to serial code. Wire it in (both C and C++ sources)
-    # so they actually use the cores.
-    find_package(OpenMP)
-    if(OpenMP_C_FOUND AND OpenMP_CXX_FOUND)
-        target_link_libraries(onebit_npu_forward PUBLIC OpenMP::OpenMP_C OpenMP::OpenMP_CXX)
-    endif()
-    # NPU_INFER_DUMP_ATTN / _LN / _HEAD write debug dumps to a path taken from the environment;
-    # only debug builds carry them
-    option(ONEBIT_NPU_DEBUG_DUMPS "NPU forward: env-named debug dumps (developer builds)" OFF)
-    if(ONEBIT_NPU_DEBUG_DUMPS)
-        target_compile_definitions(onebit_npu_forward PRIVATE NPU_INFER_DEBUG_DUMPS)
-    endif()
-    target_include_directories(onebit_npu_forward SYSTEM PUBLIC
-        npu/forward/include npu/forward/engine-src npu ${CMAKE_CURRENT_SOURCE_DIR})
-    # plan.json (npu/plan/plan.h, engine #799 P1) is read with nlohmann's parser, and the header is
-    # part of this library's interface, so the dependency is PUBLIC.
-    # engine#1185: the forward converts an imported q4nx-2 container to the device's 5,120-byte
-    # q4_1 form before model_load (npu/q4nx_device.cpp, onebit_npu), so the host layer is part of
-    # this library's link interface.
-    target_link_libraries(onebit_npu_forward PUBLIC onebit_npu onebit_npu_rt nlohmann_json::nlohmann_json dl
-                                                  pthread m)
+    # The library body is onebit_npu_forward_library() (defined above the runtime options), so the CI build of the
+    # same sources against the runtime stub (engine#1946) cannot drift from this one.
+    onebit_npu_forward_library(onebit_npu_forward onebit_npu_rt)
     target_sources(onebit PRIVATE app/forward_serve.cpp)
     # app/embed_serve.cpp's device path includes npu/embed/embed_npu.h (engine#1311): the embedding route's
     # ternary linears on the array, through the same int8 GEMM the forward links.
@@ -1702,8 +1712,8 @@ endif()
     add_test(NAME gdn_batch_order COMMAND gdn_batch_order_test)
     # engine#1905: the NPU_BATCH_CHECK oracle restores the prefix it replays from instead of resetting it, so a chunk
     # at start_pos > 0 is compared (and then run) from the state the batched pass continues. Host-GEMM, no device;
-    # skips (77) unless ONEBIT_TEST_Q4NX names a container. It links onebit_npu_forward, so the row exists only where
-    # ONEBIT_NPU configures; ci.yml does not, so this is a dev-box test (ONEBIT_NPU=ON plus a container), not CI coverage.
+    # skips (77) unless ONEBIT_TEST_Q4NX names a container. This is the dev-box row (the private runtime, a real
+    # container); a configure without ONEBIT_NPU -- CI's -- registers the same test on a fabricated container below.
     add_executable(batch_check_state_test tests/batch_check_state_test.cpp)
     target_include_directories(batch_check_state_test PRIVATE npu/forward/include)
     target_link_libraries(batch_check_state_test PRIVATE onebit_npu_forward)
@@ -1736,6 +1746,29 @@ endif()
     # check, outside this repository (docs/npu.md, "Checking the NPU").
 endif()
 
+# engine#1946: the batch-check state pin in CI. Under if(ONEBIT_NPU) above, batch_check_state_test is a dev-box row (the
+# private runtime, a container named by ONEBIT_TEST_Q4NX), so a regression of #1909 -- the oracle resetting the prefix it
+# replays from -- would land green. Here, in a configure WITHOUT the NPU lane (CI's), the same test is built from the same
+# forward sources linked against the public runtime stub (its device calls throw; the host-GEMM forward never makes one),
+# and tests/batch_check_state_ci.py runs it on a two-layer qwen3 container it writes and packs at test time
+# (tools/safetensors_to_q4nx.py): no device, no model on disk, under a second. Skips (77) without numpy.
+if(ONEBIT_NPU_HOST AND NOT ONEBIT_NPU AND NOT WIN32 AND Python3_Interpreter_FOUND)
+    add_library(onebit_npu_rt_ci_stub STATIC onnx/ep/npu_runtime_ci_stub.cpp)
+    target_include_directories(onebit_npu_rt_ci_stub PUBLIC ${CMAKE_SOURCE_DIR}/npu)
+    target_compile_definitions(onebit_npu_rt_ci_stub PUBLIC ONEBIT_NPU_RUNTIME_OWN)
+    target_link_libraries(onebit_npu_rt_ci_stub PUBLIC Threads::Threads)
+    onebit_npu_forward_library(onebit_npu_forward_ci onebit_npu_rt_ci_stub)
+    add_executable(batch_check_state_test tests/batch_check_state_test.cpp)
+    target_link_libraries(batch_check_state_test PRIVATE onebit_npu_forward_ci)
+    # direction B (tools/ci/ci_targets_check.py): ci.yml's --target list is a workflow file no lane can patch; the
+    # documented hook is the one onebit_npu_forward_syntax uses.
+    add_dependencies(onebit batch_check_state_test)
+    add_test(NAME batch_check_state
+             COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/tests/batch_check_state_ci.py
+                     $<TARGET_FILE:batch_check_state_test>)
+    set_tests_properties(batch_check_state PROPERTIES SKIP_RETURN_CODE 77 TIMEOUT 120)
+endif()
+
 # Private NPU routes (docs/npu.md, "Private routes"): closed-source add-ons in a checkout of
 # the private npu-kernels repository. Its addons/ directory defines the target
 # onebit_npu_private, which holds register_private_addon() (npu/private_route.h) and brings

docs/npu.md

@@ -2476,7 +2476,11 @@ sampled two ways), and `--batchcheck` takes the verify batch's M-curve.
 (`identical=N differing=0`). Both passes continue the same prefix: the oracle snapshots the first `start_pos`
 positions before the replay and restores them afterwards (engine#1905). So a chunk after the first, or the serve
 path's split-off last token, is compared from the state the batched pass really continues, and the flag leaves the
-run's own output unchanged. `tests/batch_check_state_test.cpp` pins both properties.
+run's own output unchanged. `tests/batch_check_state_test.cpp` pins both properties, and it runs in CI (engine#1946):
+a configure without the NPU lane builds the same forward sources against the public runtime stub
+(`onnx/ep/npu_runtime_ci_stub.cpp`, whose device calls throw; the host-GEMM forward makes none), and
+`tests/batch_check_state_ci.py` runs the test on a two-layer qwen3 container it writes and packs itself. With the
+NPU lane on, the row is the dev-box one and takes a real container from `ONEBIT_TEST_Q4NX`.
 
 `tools/mtp_depth_pool.py` pools those tickets. It refuses a pool it cannot read (a repeated tail, a cell scored twice, a
 short read, a mixed budget, a mixed window identity) and, since the tau=0.6 window, it also says when a depth is too

onnx/ep/npu_runtime_ci_stub.cpp

@@ -20,8 +20,9 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-// A link-time stand-in for the private NPU runtime, used ONLY by the CI check that compiles and links
-// npu/ort/onebit_npu_gemm.cpp (engine #874, the ONNX EP's ONEBIT_EP_NPU leg).
+// A link-time stand-in for the private NPU runtime, used ONLY by CI-configure targets: the check that compiles and
+// links npu/ort/onebit_npu_gemm.cpp (engine #874, the ONNX EP's ONEBIT_EP_NPU leg), and onebit_npu_forward_ci, the
+// host-GEMM forward that batch_check_state_test runs on in CI (engine#1946; init_host never opens the device).
 //
 // Why this file exists: that leg is built only when ONEBIT_EP_NPU=ON, which needs ONNX Runtime *and*
 // the private runtime (ONEBIT_NPU_RT_DIR, docs/npu.md), so CI compiled neither it nor anything it
@@ -35,7 +36,7 @@
 // would open, submit to or wait on the device throws, so a build that ever *ran* against it fails
 // immediately instead of silently doing nothing; the accessors that only describe a handle
 // (map/size/address/is_full_elf) return empty values, which is what an empty handle has.
-// It is never linked into libonebit_ep.so, `1bit`, or any probe.
+// It is never linked into libonebit_ep.so, `1bit`, or any probe, and a configure with ONEBIT_NPU never builds it.
 
 #include "npu_rt_api.h"
 

tests/batch_check_state_ci.py

@@ -0,0 +1,144 @@
+#!/usr/bin/env python3
+# Copyright 2026 bong-water-water-bong
+# SPDX-License-Identifier: Apache-2.0
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""The batch-check state pin (tests/batch_check_state_test.cpp, engine#1905/#1946) on a container made here.
+
+usage: tests/batch_check_state_ci.py <batch_check_state_test binary>        (needs numpy)
+
+batch_check_state_test needs a Q4NX container, and CI has none, so on its own it skips (77) there and a regression of
+#1909 -- the NPU_BATCH_CHECK oracle resetting the prefix it replays from -- lands green. This writes a two-layer qwen3
+(hidden 256, two 128-wide query heads on one kv head, ffn 512, vocab 256; the smallest geometry the I8 tiles take whole),
+packs it with tools/safetensors_to_q4nx.py, and runs the test on it in host-GEMM mode: no device, no model on disk.
+
+Measured when it was written (engine#1946): passes on main; with #1909 reverted it fails 3 of 4 arms (the forward
+holds 2 and 1 positions instead of 5 and 6, last-chunk rel-L1 0.78) -- the same failure the real containers show.
+Weights are seeded, so the run is the same every time. Exits 77 (skip) without numpy.
+"""
+import json
+import os
+import struct
+import subprocess
+import sys
+import tempfile
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+TOOL = os.path.join(HERE, "..", "tools", "safetensors_to_q4nx.py")
+
+H, HEADS, KV_HEADS, HEAD_DIM, FFN, VOCAB, LAYERS = 256, 2, 1, 128, 512, 256, 2
+
+
+def write_safetensors(path, arrays):
+    """(name -> float32 numpy array) -> a safetensors file, dense row-major."""
+    header, blobs, off = {}, [], 0
+    for name, a in arrays.items():
+        raw = a.astype("<f4").tobytes()
+        header[name] = {"dtype": "F32", "shape": list(a.shape), "data_offsets": [off, off + len(raw)]}
+        blobs.append(raw)
+        off += len(raw)
+    h = json.dumps(header, separators=(",", ":")).encode()
+    h += b" " * ((-len(h)) % 8)
+    with open(path, "wb") as f:
+        f.write(struct.pack("<Q", len(h)))
+        f.write(h)
+        for b in blobs:
+            f.write(b)
+
+
+def make_model(d, np):
+    rng = np.random.default_rng(1946)
+
+    def w(*shape, sc=0.05):
+        return (rng.standard_normal(shape) * sc).astype(np.float32)
+
+    def norm(n):
+        return (1.0 + 0.1 * rng.standard_normal(n)).astype(np.float32)
+
+    weights = {"model.embed_tokens.weight": w(VOCAB, H, sc=0.5), "model.norm.weight": norm(H),
+               "lm_head.weight": w(VOCAB, H, sc=0.2)}
+    for layer in range(LAYERS):
+        p = f"model.layers.{layer}."
+        weights.update({
+            p + "input_layernorm.weight": norm(H), p + "post_attention_layernorm.weight": norm(H),
+            p + "self_attn.q_proj.weight": w(HEADS * HEAD_DIM, H),
+            p + "self_attn.k_proj.weight": w(KV_HEADS * HEAD_DIM, H),
+            p + "self_attn.v_proj.weight": w(KV_HEADS * HEAD_DIM, H),
+            p + "self_attn.o_proj.weight": w(H, HEADS * HEAD_DIM),
+            p + "self_attn.q_norm.weight": norm(HEAD_DIM), p + "self_attn.k_norm.weight": norm(HEAD_DIM),
+            p + "mlp.gate_proj.weight": w(FFN, H), p + "mlp.up_proj.weight": w(FFN, H),
+            p + "mlp.down_proj.weight": w(H, FFN)})
+    os.makedirs(d)
+    write_safetensors(os.path.join(d, "model.safetensors"), weights)
+    cfg = {"architectures": ["Qwen3ForCausalLM"], "model_type": "qwen3", "hidden_size": H,
+           "num_hidden_layers": LAYERS, "num_attention_heads": HEADS, "num_key_value_heads": KV_HEADS,
+           "head_dim": HEAD_DIM, "intermediate_size": FFN, "vocab_size": VOCAB, "rms_norm_eps": 1e-6,
+           "rope_theta": 1000000.0, "tie_word_embeddings": False, "max_position_embeddings": 64,
+           "hidden_act": "silu", "torch_dtype": "float32"}
+    with open(os.path.join(d, "config.json"), "w") as f:
+        json.dump(cfg, f)
+    with open(os.path.join(d, "tokenizer.json"), "w") as f:
+        f.write("{}")
+    with open(os.path.join(d, "tokenizer_config.json"), "w") as f:
+        json.dump({"add_bos_token": False}, f)
+
+
+def tmp_root():
+    """tests/tmp_root.h's rule (engine#2029): scratch goes on disk, not in tmpfs. ONEBIT_TEST_TMP if set, else
+    ~/.cache/scratch-tests/tmp, else the system temp directory."""
+    e = os.environ.get("ONEBIT_TEST_TMP")
+    if e:
+        return e
+    h = os.environ.get("HOME")
+    if h:
+        p = os.path.join(h, ".cache", "scratch-tests", "tmp")
+        try:
+            os.makedirs(p, exist_ok=True)
+            return p
+        except OSError:
+            pass
+    return None
+
+
+def main():
+    if len(sys.argv) != 2:
+        print(__doc__.strip().splitlines()[2])
+        return 2
+    try:
+        import numpy as np
+    except ImportError:
+        print("SKIP: numpy is not installed")
+        return 77
+    with tempfile.TemporaryDirectory(prefix="batch_check_state_ci.", dir=tmp_root()) as tmp:
+        src, out = os.path.join(tmp, "hf"), os.path.join(tmp, "q4nx")
+        make_model(src, np)
+        r = subprocess.run([sys.executable, TOOL, "pack", src, out], capture_output=True, text=True)
+        if r.returncode != 0:
+            print(f"FAIL  packing the fabricated container: {r.stderr[-600:]}")
+            return 1
+        env = dict(os.environ, ONEBIT_TEST_Q4NX=out)
+        env.pop("NPU_INFER_HOST_GEMM", None)   # init_host is host-GEMM already; nothing else should steer it
+        r = subprocess.run([sys.argv[1]], env=env, capture_output=True, text=True, timeout=110)
+        for line in (r.stdout + r.stderr).splitlines():
+            if line.startswith(("PASS", "FAIL", "SKIP", "ALL PASS", "FAILED", "[batchcheck]")):
+                print(line)
+        if r.returncode == 77:
+            # The container is made here, so a skip would mean the test could not load it: that is a failure of this
+            # arm, not a reason to skip it.
+            print("FAIL  batch_check_state_test skipped on the fabricated container (init_host refused it)")
+            return 1
+        return r.returncode
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tests/batch_check_state_test.cpp

@@ -33,8 +33,9 @@
 //   [2] the last chunk's logits equal the sequential reference's at that position, bit for bit: in host-GEMM mode
 //       both sides are the same arithmetic in the same order, so any tolerance would only hide a one-row prefix leak.
 //
-// Where it runs: the row is registered under if(ONEBIT_NPU) (it links onebit_npu_forward) and ci.yml configures
-// without it, so it runs on a dev box with ONEBIT_NPU=ON and a container, not in CI.
+// Where it runs: in CI (engine#1946), built against the public runtime stub and run by tests/batch_check_state_ci.py
+// on a two-layer container that script fabricates; and on a dev box with ONEBIT_NPU=ON, on the container named by
+// ONEBIT_TEST_Q4NX.
 #include "q4nx_forward.h"
 
 #include <cmath>