diff --git a/CMakeLists.txt b/CMakeLists.txt index 913000720..99ec81ca4 100644 --- a/CMakeLists.txt +++ b/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 + $) + 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 cdc4c9dd5..97e42b0cc 100644 --- a/docs/npu.md +++ b/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 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..ddad781b7 --- /dev/null +++ b/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 (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("