From 170b45f657cde37cda186ffb0f96949f6df07597 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Tue, 10 Nov 2026 18:05:27 -0400 Subject: [PATCH 1/2] tests: batch_check_state runs in CI, on the runtime stub and a fabricated container (engine#1946) batch_check_state_test pins #1909 (the NPU_BATCH_CHECK oracle must restore the prefix it replays from), but it links the NPU forward, whose runtime is private, and it needs a Q4NX container. So it existed only under ONEBIT_NPU and a regression of #1909 could land green. - onebit_npu_forward_library(name runtime): the forward's one source list and link interface, used by onebit_npu_forward (private runtime) and by onebit_npu_forward_ci (onnx/ep/npu_runtime_ci_stub.cpp, whose device calls throw; the host-GEMM forward makes none). - In a configure without ONEBIT_NPU (CI's), batch_check_state_test links the stub-built forward and is built through add_dependencies(onebit ...). tests/batch_check_state_ci.py writes a seeded two-layer qwen3 (hidden 256), packs it with tools/safetensors_to_q4nx.py and runs the test on it: no device, no model, ~3 s. Skips (77) only without numpy. Measured: passes on main; with #1909 reverted it fails 3 of 4 (holds 2 and 1 positions instead of 5 and 6, rel-L1 0.78). The NPU-lane configure still builds its three container rows. ci_targets_check: the by-design bucket goes 5 -> 4 and nothing new is unbuilt. Co-Authored-By: Claude Opus 5.5 --- CMakeLists.txt | 115 ++++++++++++++++++---------- docs/npu.md | 6 +- onnx/ep/npu_runtime_ci_stub.cpp | 7 +- tests/batch_check_state_ci.py | 127 +++++++++++++++++++++++++++++++ tests/batch_check_state_test.cpp | 5 +- 5 files changed, 213 insertions(+), 47 deletions(-) create mode 100755 tests/batch_check_state_ci.py diff --git a/CMakeLists.txt b/CMakeLists.txt index fb9feacfd..d4bc2b01a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1503,6 +1503,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, @@ -1621,45 +1667,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. @@ -1686,8 +1696,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) @@ -1720,6 +1730,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 + $) + 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 diff --git a/docs/npu.md b/docs/npu.md index 2448c851c..194b9302a 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -2389,7 +2389,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 diff --git a/onnx/ep/npu_runtime_ci_stub.cpp b/onnx/ep/npu_runtime_ci_stub.cpp index 2697740c7..e2ebca31c 100644 --- a/onnx/ep/npu_runtime_ci_stub.cpp +++ b/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" diff --git a/tests/batch_check_state_ci.py b/tests/batch_check_state_ci.py new file mode 100755 index 000000000..4c3c9fa88 --- /dev/null +++ b/tests/batch_check_state_ci.py @@ -0,0 +1,127 @@ +#!/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 (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(" From 1336b50a947442ee398476f2b6012b798a9e4d6f Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Tue, 10 Nov 2026 18:51:09 -0400 Subject: [PATCH 2/2] tests: batch_check_state_ci's container goes under tmp_root(), not /tmp (engine#2029 rule, #2034) The fabricated container and its HF source are about 2 MB, but tests/tmp_root.h's rule now applies to every test: ONEBIT_TEST_TMP, else ~/.cache/scratch-tests/tmp, else the system temp directory. Re-run in the CI-style configure: Passed, directory removed. Co-Authored-By: Claude Opus 5.5 --- tests/batch_check_state_ci.py | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) diff --git a/tests/batch_check_state_ci.py b/tests/batch_check_state_ci.py index 4c3c9fa88..ddad781b7 100755 --- a/tests/batch_check_state_ci.py +++ b/tests/batch_check_state_ci.py @@ -93,6 +93,23 @@ def norm(n): 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]) @@ -102,7 +119,7 @@ def main(): except ImportError: print("SKIP: numpy is not installed") return 77 - with tempfile.TemporaryDirectory() as tmp: + 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)