diff --git a/.github/workflows/pr-agent.yml b/.github/workflows/pr-agent.yml index 6ff716d87..013a9f52b 100644 --- a/.github/workflows/pr-agent.yml +++ b/.github/workflows/pr-agent.yml @@ -38,7 +38,12 @@ jobs: runs-on: [self-hosted, pr-agent] # on the job, not the workflow: a comment that skips the job (a bot's, say) must not cancel a review concurrency: - group: pr-agent-${{ github.event.pull_request.number || github.event.issue.number }} + # A comment and a push are different jobs on one serialised model runner: an `/ask` must not cancel the review of + # the PR it is asked about, and a `synchronize` must cancel the review of the head it replaced. Grouping by event + # name keeps push-cancels-push (right: a newer head makes the older review moot) without letting a busy thread + # starve it. Measured 2026-11-10: PR #2007's review ran 19:21:26 -> cancelled 19:30:31, and `issue_comment: + # [created]` puts every comment on the PR into this same group with `cancel-in-progress: true`. + group: pr-agent-${{ github.event.pull_request.number || github.event.issue.number }}-${{ github.event_name }} cancel-in-progress: true timeout-minutes: 30 # same-repo PRs, and comments by the owner, members or collaborators on a pull request diff --git a/.gitignore b/.gitignore index 7925b4d1d..4820db372 100644 --- a/.gitignore +++ b/.gitignore @@ -41,3 +41,8 @@ pin-report.json # probe - a red that is the tree's, not the repository's (engine#1564). Ignoring them also stops run output being # committed by accident. .hrx123/logs/ + +# npu_pack_import_test's container-genre fixtures (test_q4nx_container): written to /q4nx-container, and to +# ./q4nx-container when the test runs with no golden dir -- so a run from the source root left them where `git add -A` +# picks them up, and #1986 committed five of them (820 KB) and reddened the copyright gate. Run output, not source. +q4nx-container/ diff --git a/CMakeLists.txt b/CMakeLists.txt index 1291e3466..9986d1761 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -197,6 +197,16 @@ add_test(NAME embed_serve COMMAND ${CMAKE_SOURCE_DIR}/tests/embed_serve.sh $) # serve's GET /metrics (docs/serve.md, "Metrics"): a fake backend, no GPU. add_test(NAME serve_hangup COMMAND ${CMAKE_SOURCE_DIR}/tests/serve_hangup_test.sh $) +# engine#1945: scripts/bump-hrx-sync.sh tells AMD's own commit from a line AMD rebased apart from a commit that +# is not AMD's. Local git repositories only: no network, no GPU. +add_test(NAME bump_hrx_sync COMMAND bash ${CMAKE_SOURCE_DIR}/tests/bump_hrx_sync_test.sh ${CMAKE_SOURCE_DIR}/scripts/bump-hrx-sync.sh) +# engine#2102: the sentence a serve owes an operator when no 16-row decode kernel was found (app/npu_decode_note.h). Pure: +# the silent cases are the ones that matter, and they are arms. +add_executable(npu_decode_note_test tests/npu_decode_note_test.cpp) +target_include_directories(npu_decode_note_test PRIVATE app) +add_test(NAME npu_decode_note COMMAND npu_decode_note_test) +add_dependencies(onebit npu_decode_note_test) +add_test(NAME bump_hrx_report COMMAND bash ${CMAKE_SOURCE_DIR}/tests/bump_hrx_report_test.sh ${CMAKE_SOURCE_DIR}/scripts/bump-hrx-report.sh) add_test(NAME metrics_serve COMMAND ${CMAKE_SOURCE_DIR}/tests/metrics_serve.sh $) # serve --print-backend: the llama-server Lemonade starts itself, one hop (docs/serve.md): no GPU. add_test(NAME print_backend COMMAND ${CMAKE_SOURCE_DIR}/tests/print_backend.sh $) @@ -263,9 +273,8 @@ add_test(NAME copyright_gate COMMAND ${CMAKE_SOURCE_DIR}/tests/copyright_gate_te # engine#1701: the NPU hold's queue/token follow ONEBIT_HOLD_DIR, and its queue settings never reach the payload. # Shell only (a stub gpu-hold plus the tree's own), no device, ~1 s. add_test(NAME npu_hold COMMAND bash ${CMAKE_SOURCE_DIR}/tests/npu_hold_test.sh) -# engine#1696: the gpu_hold_* registrations live in their own file, so adding one does not conflict with every -# other test-adding PR through the one contiguous add_test block this used to be part of (see that file's header). -include(cmake/tests-gpu-hold.cmake) +# engine#1696: the gpu_hold_* registrations live in cmake/tests-gpu-hold.cmake. It is no longer named here: every +# cmake/tests-*.cmake is included by the glob near the end of this file (engine#1979). # engine#1485: the two port read-backs the route tests use (tests/serve_port.sh). Script-based, no device, no model -- # it feeds the parsers the exact lines `1bit serve` and `llama-server` print, because a read-back that silently returns # empty turns into "the server did not come up" and is blamed on whatever the row under test checks. @@ -276,6 +285,9 @@ add_test(NAME serve_port COMMAND bash ${CMAKE_SOURCE_DIR}/tests/serve_port_test. # device. Deliberately OUTSIDE if(ONEBIT_SERVE_TEST_GGUF): that row is device-gated, and the whole point is that the # classification is checked in every configure, including the ones that never load a model. add_test(NAME serve_context_limit_classify COMMAND bash ${CMAKE_SOURCE_DIR}/tests/serve_context_limit_test.sh --self-test) +# engine#1785: `serve --flash` names two settings (--ctx-size 1024, --prefix-pin on --device npu) and refuses what +# contradicts them. Every case ends in the option check, so the row needs no model, no device and no port. +add_test(NAME serve_flash_refusals COMMAND bash ${CMAKE_SOURCE_DIR}/tests/serve_flash_test.sh $) # engine#1594: every target ci.yml's `cmake --build --target …` list names must exist in the configuration ci.yml # creates. #1583 added one name for a target defined inside if(ONEBIT_NPU), which CI does not configure -- and ninja # rejects the whole command line before building anything, so `onebit` stopped being compiled for every PR. Reads the @@ -308,7 +320,28 @@ add_test(NAME bench_preflight COMMAND bash ${CMAKE_SOURCE_DIR}/tools/tests/test- # engine#1702: builds a per-layer MoE expert placement table from a moe_trace, and refuses a prefill-only trace because # placement is a decode lever (5.4-6.9 % on batched prefill against dedup's 90-94 %). --self-test drives it on synthetic # traces with a known answer; no model, no device. +# engine#1934: reading ONE tensor out of a safetensors shard needs `data_offsets` measured from 8 + header_len, not +# from byte 8. Getting that wrong returns a *different* tensor of the same size that still looks finite, and it cost +# two false conclusions about a pack in one night. `--self-test` drives the rule on a synthetic shard (no network): +# the same offset from the data section returns the tensor, from byte 8 it does not. +add_python_test(hf_tensor_self ${CMAKE_SOURCE_DIR}/tools/hf_tensor.py --self-test) +# engine#1946: a test that returns 77 is green, and green-because-it-did-not-run is the failure this fleet keeps paying +# for. `tools/ci/ctest_skips.py` reads a ctest log and reports every test that did not execute -- `***Skipped` (77) against +# `tools/ci/ctest_skips.txt`, and `(Not Run)` (a missing binary) as a build-list mistake that can never be baselined. +# Self-test only here: wiring it into ci.yml's Test step needs the log plus a `workflow`-scoped token (engine#1817). +add_python_test(ctest_skips_self ${CMAKE_SOURCE_DIR}/tools/ci/ctest_skips.py --self-test) +# engine#2106 / #2102: the release gate's NPU serve check asserts the fast route, a decode floor, the first token and a +# non-empty answer, because a serve on the slow per-op route still answers. --self-test drives the judging on the shape +# of each failure (the silent fallback at 1.9 tok/s, thinking left on, a degenerate reply); no device, no binary. +add_python_test(npu_serve_smoke_self ${CMAKE_SOURCE_DIR}/tools/npu_serve_smoke.py --self-test) add_python_test(moe_placement ${CMAKE_SOURCE_DIR}/tools/moe_placement.py --self-test) +# engine#1313 item 3: the "after" of "before/after routing accuracy" -- a ridge head over --dump-encoder output, with the +# table's own logits as the before arm. Self-test only, so it needs no dump and no model. +add_python_test(laya_head_finetune ${CMAKE_SOURCE_DIR}/tools/laya_head_finetune.py --self-test) +# engine#856: one HRX process's GTT split into the loader's copy and HRX's packed second copy, from its own -v log. +add_python_test(hrx_gtt_split_self ${CMAKE_SOURCE_DIR}/tools/hrx_gtt_split.py --self-test) +# engine#679: time per HRX dispatch from the backend's serial trace (the id hash and the per-kernel sums). +add_python_test(hrx_dispatch_profile_self ${CMAKE_SOURCE_DIR}/tools/hrx_dispatch_profile.py --self-test) # engine #1279: the version a package is built with (scripts/pack-version.sh) -- a release tag as # it stands, a candidate mapped into the ISO week of its own date with a suffix that sorts below that # release, and the ordering a candidate needs: previous release < candidate < its release. The @@ -319,7 +352,18 @@ add_test(NAME pack_version COMMAND ${CMAKE_SOURCE_DIR}/tests/pack_version_test.s # resolves outside the payload, and a runtime set shipped without the aqlprofile agent libhsa opens. Each # arm asserts the row's message, and a clean tree is the control. No package build, no GPU, runs in CI. add_test(NAME pack_check COMMAND ${CMAKE_SOURCE_DIR}/tests/pack_check_test.sh) +# engine#2106: the hardware gate's `npuserve` phase, end to end against the stand-in serve tools/npu_serve_smoke.py +# writes: a fast-route serve passes the gate, a slow-route one fails it while its answers are fine, a dead one fails, +# and an unconfigured or NPU-less host skips. No device, no 1bit binary, and the fleet's queues are not touched. +add_test(NAME hw_gate_npuserve COMMAND ${CMAKE_SOURCE_DIR}/tests/hw_gate_npuserve_test.sh) +# Both start a loopback serve on a port picked by binding port 0 and releasing it, so they take the lock every +# loopback-server test takes (tests/route_port_lock_test.py): two of them racing for one freed port is a flaky failure. +set_tests_properties(hw_gate_npuserve npu_serve_smoke_self PROPERTIES RESOURCE_LOCK route_ports) add_test(NAME check_ci_runners COMMAND ${CMAKE_SOURCE_DIR}/tests/check_ci_runners_test.sh) +# engine #1504: the release preflight blocks on a runner unit that is down and only warns on a short restart policy. +# The real scripts/ci-runner-health.sh and weekly.sh's preflight_runner_units() against a stub systemctl; no systemd. +add_test(NAME weekly_runner_preflight COMMAND ${CMAKE_SOURCE_DIR}/tests/weekly_runner_preflight_test.sh) +add_test(NAME ci_runner_health_selftest COMMAND bash ${CMAKE_SOURCE_DIR}/scripts/ci-runner-health.sh --self-test) # engine #1750: the carried amdxdna series applies in order to the pinned driver, and the context-restart code it # produces fails a context's jobs (and logs it) when the restart fails, instead of leaving them parked or running # them at close on a half-built firmware context. The driver's own functions in a user-space harness: no NPU, no @@ -384,6 +428,13 @@ add_test(NAME serve_watchdog COMMAND serve_watchdog_test) # engine#1781: the comparator behind NPU_BATCH_CHECK -- pure arithmetic over two float arrays, so it runs with no model, # no device and no runtime. It is a header, so the test includes it directly and needs no link libraries. add_executable(npu_sampling_test tests/npu_sampling_test.cpp) +# The shape rule for a 1.8 external-data initializer (npu/pack_initializer_shape.h, engine#1503's graph-initializer +# reader). Those packs zero every `dims`, so a byte length has to become a shape -- and a shape inferred wrongly imports +# cleanly and serves garbage, so the refusals are the arms worth pinning. Header-only, no parser, no pack, no device. +add_executable(pack_initializer_shape_test tests/pack_initializer_shape_test.cpp) +target_include_directories(pack_initializer_shape_test PRIVATE npu) +add_test(NAME pack_initializer_shape COMMAND pack_initializer_shape_test) +add_dependencies(onebit pack_initializer_shape_test) target_include_directories(npu_sampling_test PRIVATE app) target_link_libraries(npu_sampling_test PRIVATE nlohmann_json::nlohmann_json) # The NPU route's bounded load retry (app/npu_load_retry.h, engine #1010, docs/npu.md "When a load follows an evict"): @@ -418,6 +469,13 @@ if(TARGET OpenMP::OpenMP_CXX) target_link_libraries(mtp_draft_test PRIVATE OpenMP::OpenMP_CXX) endif() add_test(NAME mtp_draft COMMAND mtp_draft_test) +# engine#1781 gate 2: the distribution-following acceptance rule (accept with min(1, p/q), resample the residual at the +# first rejection), with the temperature-0 identity and a deliberately wrong drafter whose output must still follow p. +# Header-only and device-free; pulled into CI through npu_sampling_test for the same reason as mtp_draft_test below. +add_executable(spec_accept_test tests/spec_accept_test.cpp) +target_include_directories(spec_accept_test PRIVATE npu) +add_test(NAME spec_accept COMMAND spec_accept_test) +add_dependencies(npu_sampling_test spec_accept_test) # ci.yml's build step lists every ctest target explicitly, and this branch cannot edit that file (the push # token has no `workflow` scope), so the binary is pulled in as a dependency of a listed test target. When a # workflow-scoped push is available, add `mtp_draft_test` to ci.yml's build list and delete this line. @@ -639,7 +697,7 @@ option(ONEBIT_NPU_HOST "Build the NPU host library and its tests" ${_onebit_npu_ if(ONEBIT_NPU_HOST) find_package(PkgConfig REQUIRED) pkg_check_modules(PCRE2 REQUIRED IMPORTED_TARGET libpcre2-8) -add_library(onebit_npu STATIC npu/file_map.cpp npu/flm_head_import.cpp npu/full_elf.cpp npu/model.cpp +add_library(onebit_npu STATIC npu/file_map.cpp npu/flm_head_import.cpp npu/full_elf.cpp npu/model.cpp npu/vision_pack.cpp npu/vision_tower.cpp npu/vision_image.cpp npu/moe.cpp npu/pack.cpp npu/pack_import.cpp npu/private_route.cpp npu/q4nx.cpp npu/q4nx_device.cpp npu/tokenizer.cpp) target_include_directories(onebit_npu PUBLIC npu) @@ -755,6 +813,12 @@ add_executable(laya_eval_test tests/laya_eval_test.cpp) target_link_libraries(laya_eval_test PRIVATE onebit_laya nlohmann_json::nlohmann_json) add_test(NAME laya_eval COMMAND laya_eval_test) add_dependencies(onebit laya_eval_test) +# engine#1313 item 3: the head-only fine-tune over --dump-encoder dumps, on a synthetic checkpoint and dump (no model, +# no device). It needs torch + numpy, which the engine build does not, so a box without them skips (77). +add_python_test(laya_finetune ${CMAKE_SOURCE_DIR}/tests/laya_finetune_test.py) +if(Python3_Interpreter_FOUND) + set_tests_properties(laya_finetune PROPERTIES SKIP_RETURN_CODE 77) +endif() # Crafted Laya checkpoints the reader and the scorer must refuse (written at run time); runs in CI. add_executable(laya_bounds_test tests/laya_bounds_test.cpp) target_link_libraries(laya_bounds_test PRIVATE onebit_laya nlohmann_json::nlohmann_json) @@ -826,6 +890,12 @@ if(Python3_Interpreter_FOUND) COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/tools/ci/npu_decl_check.py --check --root ${CMAKE_SOURCE_DIR}) add_test(NAME npu_decl_check_self COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/tools/ci/npu_decl_check.py --self-test) + # engine#2057: no fopen(..., "w"/"a") in the tree without a stated reason -- CodeQL's world-writable class, kept fixed + # by a check rather than by the next scan. Source text only: no build, no device. + add_test(NAME file_create_check + COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/tools/ci/file_create_check.py --check --root ${CMAKE_SOURCE_DIR}) + add_test(NAME file_create_check_self + COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/tools/ci/file_create_check.py --self-test) # the new-model watcher's alert contract (offline: synthetic batches against the # committed registry, plus FastFlowLM's org listing and releases from the recorded # payloads under tests/fixtures/fflm/ — engine #588), so a regression cannot leave the @@ -1083,6 +1153,14 @@ target_link_libraries(npu_stop_tokens_test PRIVATE nlohmann_json::nlohmann_json) add_test(NAME npu_stop_tokens COMMAND npu_stop_tokens_test) add_dependencies(npu_sampling_test npu_stop_tokens_test) +# The hand-off between the thread that generates a streamed reply and the one that writes it (app/stream_bridge.h, +# engine#1785: the NPU route sent every frame after the generation ended). Two threads and a clock; no device. +add_executable(stream_bridge_test tests/stream_bridge_test.cpp) +target_include_directories(stream_bridge_test PRIVATE ${CMAKE_SOURCE_DIR}/app) +target_link_libraries(stream_bridge_test PRIVATE Threads::Threads) +add_test(NAME stream_bridge COMMAND stream_bridge_test) +add_dependencies(npu_sampling_test stream_bridge_test) + # The plan.json loader (npu/plan/plan.h, engine #799 P1): schema, refusals and the decision fields that # replaced the filename probes. No model, no container, no device. add_executable(npu_elf_args_test tests/npu_elf_args_test.c npu/forward/src/q4nx_pack.c npu/forward/src/q4nx_dequant.c) @@ -1146,6 +1224,32 @@ if(ONEBIT_NPU_CAPTURED AND ONEBIT_NPU_PDI AND ONEBIT_NPU_GOLDEN) endif() # Q4NX chunk decoders against 1bit-MONSTER's verified decoders, on committed real tiles +# engine#1788: the vision component reader (vision_weight.q4nx). Crafted containers for every refusal; the published +# pack's own numbers when ONEBIT_MINICPM_PACK names a MiniCPM-V 4.7 pack directory. +add_executable(npu_vision_pack_test tests/npu_vision_pack_test.cpp) +target_link_libraries(npu_vision_pack_test PRIVATE onebit_npu) +add_test(NAME npu_vision_pack COMMAND npu_vision_pack_test) +add_dependencies(onebit npu_vision_pack_test) +# The host float32 tower against goldens from the transformers implementation (see the test's header). +add_executable(npu_vision_tower_test tests/npu_vision_tower_test.cpp) +target_link_libraries(npu_vision_tower_test PRIVATE onebit_npu) +add_test(NAME npu_vision_tower COMMAND npu_vision_tower_test) +add_dependencies(onebit npu_vision_tower_test) +# The image plan (sizes, slice grid, text rows) against the reference image processor's answers for 3,481 sizes. +add_executable(npu_vision_plan_test tests/npu_vision_plan_test.cpp) +target_include_directories(npu_vision_plan_test PRIVATE npu) +add_test(NAME npu_vision_plan COMMAND npu_vision_plan_test) +add_dependencies(onebit npu_vision_plan_test) +# The resize (bit for bit against torch's uint8 antialiased bicubic) and the pieces handed to the tower. +add_executable(npu_vision_image_test tests/npu_vision_image_test.cpp) +target_link_libraries(npu_vision_image_test PRIVATE onebit_npu) +add_test(NAME npu_vision_image COMMAND npu_vision_image_test) +add_dependencies(onebit npu_vision_image_test) +# The text side of an image: its tokens and the canvas positions, against the reference model code's own answers. +add_executable(npu_vision_prompt_test tests/npu_vision_prompt_test.cpp) +target_include_directories(npu_vision_prompt_test PRIVATE npu) +add_test(NAME npu_vision_prompt COMMAND npu_vision_prompt_test) +add_dependencies(onebit npu_vision_prompt_test) add_executable(npu_q4nx_test tests/npu_q4nx_test.cpp) target_include_directories(npu_q4nx_test PRIVATE app) target_link_libraries(npu_q4nx_test PRIVATE onebit_npu) @@ -1168,6 +1272,16 @@ endif() # q4nx_pack.c is listed here rather than reached through onebit_npu: the integer host arm's checks # (engine#1808) call q4nx_host_gemm_i8*() directly, and those live in that file. add_executable(npu_pack_test tests/npu_pack_test.cpp npu/forward/src/q4nx_pack.c npu/forward/src/q4nx_dequant.c) +# A DFlash draft package as a model source (engine#1780): config + container -> a checked model whose tensors decode +# one at a time. The loader is compiled here and not into onebit_npu because it needs the q4_1 tile decoder, which +# lives with the forward (npu/forward/src/q4nx_dequant.c); it moves there with the draft's forward. Synthetic package, +# so no model, no device; a real package is an optional argument (-DONEBIT_DFLASH_DRAFT=). +add_executable(dflash_draft_test tests/dflash_draft_test.cpp npu/dflash_draft.cpp npu/forward/src/q4nx_dequant.c) +target_include_directories(dflash_draft_test PRIVATE npu npu/forward/include) +target_link_libraries(dflash_draft_test PRIVATE onebit_npu) +add_dependencies(onebit dflash_draft_test) +set(ONEBIT_DFLASH_DRAFT "" CACHE PATH "A DFlash draft package directory for dflash_draft_test's real-package arm") +add_test(NAME dflash_draft COMMAND dflash_draft_test ${ONEBIT_DFLASH_DRAFT}) add_executable(flm_head_import tools/flm_head_import.cpp) target_link_libraries(flm_head_import PRIVATE onebit_npu) add_executable(npu_flm_head_import_test tests/npu_flm_head_import_test.cpp) @@ -1178,6 +1292,29 @@ target_link_libraries(npu_pack_test PRIVATE onebit_npu) # golden fixture that pins them. target_include_directories(npu_pack_test PRIVATE npu/forward/include ${CMAKE_SOURCE_DIR}/tests) +# NPU_INFER_DUMP_OPS (npu/forward/include/ops-dump.h, engine#1503): the on/off rule, the mean|x|/absmax arithmetic and +# the line's text. Header-only, no device. `add_dependencies(onebit ...)` is what gets it built in CI, whose target +# list is explicit (tools/ci/ci_targets_check.py). +# npu/forward/include/file_create.h (engine#2057): the openers every dump and the head-image cache write through -- 0600 / +# at most 0644 whatever the umask, an existing 0666 file narrowed, a symlink refused with ELOOP. Plain C, no model, no +# device; the call sites only swapped fopen() for an opener, so without this arm a revert to fopen() passes every test. +if(NOT WIN32) + add_executable(file_create_test tests/file_create_test.c) + target_include_directories(file_create_test PRIVATE npu/forward/include) + add_test(NAME file_create COMMAND file_create_test) + add_dependencies(onebit file_create_test) +endif() +add_executable(ops_dump_test tests/ops_dump_test.cpp) +target_include_directories(ops_dump_test PRIVATE npu/forward/include) +add_test(NAME ops_dump COMMAND ops_dump_test) +add_dependencies(onebit ops_dump_test) +# When the dense-tail cache is on by default (npu/forward/include/ffn_cache_policy.h): the size rule and the two +# models it was measured on, one on each side of the old floor. Header-only, no device. +add_executable(ffn_cache_policy_test tests/ffn_cache_policy_test.cpp) +target_include_directories(ffn_cache_policy_test PRIVATE npu/forward/include) +add_test(NAME ffn_cache_policy COMMAND ffn_cache_policy_test) +add_dependencies(onebit ffn_cache_policy_test) + # The device weight layout contract (q4nx_weight_layout): pure index math, no device # and no runtime. CI builds an explicit target list, so the check compiles itself with # the build's compiler at test time (tests/npu_blocked_layout_test.sh) rather than being @@ -1283,6 +1420,11 @@ add_test(NAME quant_i8_row COMMAND quant_i8_row_test) add_executable(attn_math_test tests/attn_math_test.cpp) target_include_directories(attn_math_test PRIVATE npu/forward/include) add_test(NAME attn_math COMMAND attn_math_test) +# engine#1990: the test also holds rows_context() (M rows in one thread team) to the per-row form bit for bit, +# and that comparison is only about threads when the test has them. +if(TARGET OpenMP::OpenMP_CXX) + target_link_libraries(attn_math_test PRIVATE OpenMP::OpenMP_CXX) +endif() # ci.yml's build step names its targets explicitly and needs the `workflow` scope to extend, so bind each new # test binary to a target the list already builds (same reason as npu_moe_test above); drop this line when the # workflow is next touched with that scope and add these to the list itself. `host_gemm_pad_test` is here for the @@ -1356,6 +1498,14 @@ add_dependencies(onebit npu_answer_parser_test) add_executable(batch_check_compare_test tests/batch_check_compare_test.cpp) target_include_directories(batch_check_compare_test PRIVATE npu/forward/include) add_test(NAME batch_check_compare COMMAND batch_check_compare_test) + +# engine#1780: the DFlash draft block's acceptance arithmetic -- tokens = accepted + bonus, the break-even rule, and the +# published OGA-route counts as known answers. Header-only and host-only: no container, no device, no NPU runtime, so it +# runs wherever this file is built. Registered at top level with the same hook as the row above, for the same reason. +add_executable(dflash_accept_test tests/dflash_accept_test.cpp) +target_include_directories(dflash_accept_test PRIVATE npu/forward/include) +add_test(NAME dflash_accept COMMAND dflash_accept_test) +add_dependencies(onebit dflash_accept_test) # direction B (tools/ci/ci_targets_check.py): a target ci.yml's --target list does not name is "never built", and ci.yml # is a workflow file no lane can patch. The documented hook is the one onebit_npu_forward_syntax uses. add_dependencies(onebit batch_check_compare_test) @@ -1445,6 +1595,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, @@ -1523,6 +1719,13 @@ if(ONEBIT_NPU) add_executable(npu_rt_shim_test tests/npu_rt_shim_test.cpp) target_link_libraries(npu_rt_shim_test PRIVATE onebit_npu_lane) add_test(NAME npu_rt_shim COMMAND npu_rt_shim_test) + # engine#753: does the engine EXECUTE on a dma-buf it did not allocate? This one does open the device (it imports a udmabuf and + # hands the exporter's mapping to the runtime), so it takes the kernel ELF as an argument and SKIPS (77) without one -- a box + # with no ELF registers a skip rather than a red row, the same rule as the other device-dependent rows. + add_executable(npu_dmabuf_import_test tests/npu_dmabuf_import_test.cpp) + target_link_libraries(npu_dmabuf_import_test PRIVATE onebit_npu_lane) + add_test(NAME npu_dmabuf_import COMMAND npu_dmabuf_import_test "") + set_tests_properties(npu_dmabuf_import PROPERTIES SKIP_RETURN_CODE 77) # The model-generic forward (npu/forward): the GGUF-on-NPU serve route # (docs/npu.md, "Open"). q4nx_forward.cpp + its Q4NX dequant/pack/semantics @@ -1542,6 +1745,12 @@ if(ONEBIT_NPU) target_include_directories(hybrid_prompt_probe PRIVATE ${CMAKE_SOURCE_DIR}/npu ${CMAKE_SOURCE_DIR}/npu/forward/include ${CMAKE_SOURCE_DIR}/npu/forward/engine-src) target_link_libraries(hybrid_prompt_probe PRIVATE onebit_npu_forward onebit_npu_rt) + # engine#1788: one picture and one question through a MiniCPM-V 4.7 pack with every step the engine's own + # (plan, resize, host tower, prompt tokens, canvas positions, then step_batch_embeds on the text model). + add_executable(minicpm_image_probe npu/forward/minicpm_image_probe.cpp) + target_include_directories(minicpm_image_probe PRIVATE ${CMAKE_SOURCE_DIR}/npu + ${CMAKE_SOURCE_DIR}/npu/forward/include ${CMAKE_SOURCE_DIR}/npu/forward/engine-src) + target_link_libraries(minicpm_image_probe PRIVATE onebit_npu onebit_npu_forward onebit_npu_rt) add_executable(mrope_planes_probe npu/forward/mrope_planes_probe.cpp) target_include_directories(mrope_planes_probe PRIVATE ${CMAKE_SOURCE_DIR}/npu ${CMAKE_SOURCE_DIR}/npu/forward/include ${CMAKE_SOURCE_DIR}/npu/forward/engine-src) @@ -1563,45 +1772,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. @@ -1628,8 +1801,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) @@ -1662,6 +1835,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 @@ -1687,6 +1883,42 @@ if(ONEBIT_NPU_PRIVATE) endif() endif() +# ── Test registrations that live in their own file (engine#1696, engine#1979) ─ +# Every cmake/tests-*.cmake is included here, in name order. A test added as a NEW file there touches no line any +# other pull request touches, so it cannot conflict and its registration cannot go stale while the PR waits: 37 commits +# changed this file in the 24 h before engine#1979 was filed, and three open PRs had a registration hunk that no longer +# applied. The rows above stay where they are; nothing has to move. +# Here, after every target and after find_package(Python3), so a file may use if(TARGET ...), add_python_test() and +# ${Python3_EXECUTABLE}; and before the RESOURCE_LOCK loop below, which asks if(TEST ...). +# CONFIGURE_DEPENDS re-runs the glob at build time, so a file added after the first configure is picked up. +# Readers that follow this shape: tools/ci/ci_targets_check.py (the glob), tools/ci/hrx_build_scope.py and +# .github/workflows/hrx-build.yml (the cmake/tests-*.cmake path), tools/ci/tests_referenced_check.py (all of cmake/). +# The lines between the two markers are run as they stand by tests/registration_glob_test.sh, in a scratch project. +# >>> registration glob +file(GLOB ONEBIT_TEST_REGISTRATION_FILES CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/cmake/tests-*.cmake) +list(SORT ONEBIT_TEST_REGISTRATION_FILES) +# A file this CMakeLists.txt already include()s by name is skipped: that include fixes the file's position (some rows +# need one), and including it twice stops the configure with `add_test given test NAME ... which already exists`. So a +# family file moved out with its own include() line and this glob can land in either order. +# "By name" means a line CMake would execute -- include( at the start of a line, the same rule +# tools/ci/ci_targets_check.py reads by -- and not the same text in a comment: a substring search took +# `# was include(cmake/tests-gpu-hold.cmake)` for an include, and the seven gpu_hold_* tests left the configuration with +# every gate green (review of engine#1983). An include this pattern does not recognise is the loud miss: both includes +# run and the configure stops. +file(READ ${CMAKE_CURRENT_LIST_FILE} _onebit_top_cmake_text) +set(_onebit_include_prefix "(^|\n)[ \t]*include\\([ \t]*\"?(\\$\\{CMAKE_(CURRENT_)?SOURCE_DIR\\}/)?cmake/") +foreach(_onebit_test_registration_file IN LISTS ONEBIT_TEST_REGISTRATION_FILES) + get_filename_component(_onebit_test_registration_name ${_onebit_test_registration_file} NAME) + string(REGEX REPLACE "([.+])" "\\\\\\1" _onebit_test_registration_re "${_onebit_test_registration_name}") + string(REGEX MATCH "${_onebit_include_prefix}${_onebit_test_registration_re}\"?[ \t]*\\)" _onebit_named + "${_onebit_top_cmake_text}") + if(_onebit_named STREQUAL "") + include(${_onebit_test_registration_file}) + endif() +endforeach() +unset(_onebit_top_cmake_text) +# <<< registration glob + # ── Loopback ports: route tests run one at a time (engine#1070) ────────────── # Every test named below starts a server on 127.0.0.1. Most pick a port by binding :0, closing the # socket, then letting the child bind that number; three use fixed 43xxx literals. Under `ctest -j` diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 7d5f58f6d..b86b8a268 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -87,6 +87,17 @@ an RFC or issue. hands the script to the interpreter instead of relying on the file's exec bit; `tests/python_test_order_check.py --check` enforces that it follows `find_package(Python3)`. Two rows may not share a name: CMake stops at configure with `add_test given test NAME "…" which already exists in this directory`. +- **A new test's registration goes in its own file.** Put the `add_test(...)` (or `add_python_test(...)`) row, its + comment, and its `set_tests_properties` in a new `cmake/tests-.cmake` with the Apache notice, and do not touch + `CMakeLists.txt` for it: every `cmake/tests-*.cmake` is included by a glob near the end of `CMakeLists.txt`, after all + targets exist. A new file cannot conflict with another pull request, and a row appended to `CMakeLists.txt` does — + that file changed 37 times in one day and three open pull requests had a registration hunk that no longer applied + (engine [#1979](https://github.com/1bit-MONSTER/engine/issues/1979)). A row that belongs to an existing family goes in + that family's file (`cmake/tests-gpu-hold.cmake`). An `add_executable` for a test binary still goes in + `CMakeLists.txt` next to the sources it builds; only the registration moves. Existing rows stay where they are. + A registration file that needs a fixed position may still be `include()`d by name from `CMakeLists.txt`; the glob + then leaves it alone. `tests/registration_glob_test.sh` (itself registered from `cmake/tests-registration-glob.cmake`) + runs the glob block and holds both rules. - **Pins only move forward.** CI fails a pull request whose submodule pin is behind, or has diverged from, the pin on `main` (`tools/check_pins.py`), because that drops commits the engine already shipped. A rollback on purpose carries the `pin rollback` label. diff --git a/app/forward_serve.cpp b/app/forward_serve.cpp index c70f989e0..db836fc45 100644 --- a/app/forward_serve.cpp +++ b/app/forward_serve.cpp @@ -35,6 +35,9 @@ // resets it from position 0. #include "forward_serve.h" +#include "../npu/forward/include/file_create.h" // dumps are created 0600, never through a symlink +#include "npu_decode_note.h" + #include "../npu/forward/include/q4nx_forward.h" #include "../npu/target/xdna/xdna_target.h" // XdnaTarget::kName for --npu-target #include "npu_device_lock.h" @@ -47,6 +50,7 @@ #include "npu_prompt_ids.h" #include "npu_sampling.h" #include "npu_stop_tokens.h" +#include "stream_bridge.h" #include "npu_systemone.h" #include "tokenizer.h" #include "../npu/decision.h" // the d3 readout head (npu/decision.h, issue #653) @@ -74,6 +78,16 @@ namespace onebit { namespace { +// A prompt that leaves no room for an answer inside the context (the model's, or --ctx-size). +struct ContextOverflow : std::runtime_error { + size_t prompt_tokens; + int context; + ContextOverflow(size_t n, int ctx) + : std::runtime_error("request (" + std::to_string(n) + " tokens) exceeds the available context size (" + + std::to_string(ctx) + " tokens)"), + prompt_tokens(n), context(ctx) {} +}; + using json = nlohmann::json; struct ForwardEngine { @@ -109,6 +123,10 @@ struct ForwardEngine { // slot's attention KV and GDN state. Only ever touched under mu. std::unique_ptr prefix; std::vector prefix_state; + // --prefix-pin (engine#1785): keep the state at the point where a request stops agreeing with the + // previous one, so single-turn requests under one system prompt resume from it (PrefixCache::pin_point). + bool prefix_pin = false; + std::vector prev_ids; // the previous request's prompt, what pin_point() compares against // The d3 decision head (#653), loaded when decision_config.json and decision.q4nx sit next to the // model: it is what turns this server into a decision endpoint as well as a chat one. std::unique_ptr head; @@ -138,6 +156,10 @@ struct CallStats { double prompt_ms = 0; double predicted_ms = 0; npu::PrefixDecision prefix{-1, 0, "off"}; // "off" when --prefix-slots 0 + // engine#2111: each generated token's text, in order, and how many of them were reasoning. -1 = not counted (a + // /v1/completions reply is raw text with no split), so the usage object then carries no details at all. + std::vector pieces; + int reasoning_tokens = -1; }; // One generated token's log-probability entry, with the text that token produced. The two travel @@ -179,7 +201,10 @@ int json_int(const std::string& path, const char* key, int def) { return def; } -void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Response& res, bool chat) { +// `bridge` is where a streamed reply's frames go as they are produced (handle_generate below runs this on its own +// thread for a streaming request); null means the caller wants the whole response in `res`. +void handle_generate_impl(ForwardEngine& e, const httplib::Request& req, httplib::Response& res, bool chat, + npu::StreamBridge* bridge) { // One id per response, taken before any step: the dumps of two prompts in one serve must not share names. const int dump_req = e.dump_logits.empty() ? 0 : ++e.dump_seq; json body; @@ -198,6 +223,12 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // from `ids`. Do not log `prompt` as the request's text without handling this case first. std::vector prompt_ids; bool prompt_is_ids = false; + // --prefix-pin: the leading system messages and a probe user message (npu::pin_probe_messages), rendered. Whatever + // of that the full prompt starts with is the part every question under this system prompt shares, so it can be pinned on + // the FIRST request (PrefixCache::pin_at) -- comparing with the previous request only finds it on the second, + // and can run a few tokens into the question when two questions open with the same words. Empty = the + // request has no system message, or the route is not pinning. + std::string system_prompt; try { if (chat) { if (!e.chat_ok) throw std::runtime_error(e.chat_refusal); @@ -214,8 +245,25 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res extra["enable_thinking"] = npu::thinking_from_body(body); prompt = npu::render_chat_template(e.tmpl, ordered.at("messages"), /*add_generation_prompt=*/true, extra, e.bos_text, e.eos_text); + if (e.prefix_pin) { + // A template that cannot render the probe is no reason to refuse the request: it only means + // there is no early pin, and the log says so. + try { + const nlohmann::ordered_json probe = npu::pin_probe_messages(ordered.at("messages")); + if (!probe.empty()) + system_prompt = npu::render_chat_template(e.tmpl, probe, /*add_generation_prompt=*/false, extra, + e.bos_text, e.eos_text); + } catch (const std::exception& ex) { + system_prompt.clear(); + std::fprintf(stderr, "[prefix] no early pin: the template did not render the system turn (%s)\n", ex.what()); + } + } } else { prompt = npu::chatml(body.at("messages"), npu::thinking_from_body(body)); + if (e.prefix_pin) { + const json probe = npu::pin_probe_messages(body.at("messages")); + if (!probe.empty()) system_prompt = npu::chatml(probe, npu::thinking_from_body(body)); + } } } else { prompt_is_ids = npu::prompt_ids_from_body(body, e.tok->size(), prompt_ids); @@ -273,7 +321,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // Returns the finish_reason the request earned: "stop" for an end-of-sequence token, // a stop sequence or a grammar that allows nothing more, "length" when the token // budget ran out. - auto run = [&e, &req, prompt, prompt_ids, prompt_is_ids, max_tokens, chat, &opts, gram, &stats, dump_req]( + auto run = [&e, &req, prompt, system_prompt, prompt_ids, prompt_is_ids, max_tokens, chat, &opts, gram, &stats, dump_req, bridge]( const std::function&)>& on_text) { std::string finish = "stop"; // The tokenizer is const and keeps no per-call state, so it runs before the lock. A caller @@ -293,7 +341,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // The KV cache grows with every position: the prompt and the answer together stay // inside the model's context. const int room = e.max_context - int(ids.size()); - if (room < 1) throw std::runtime_error("prompt longer than the context"); + if (room < 1) throw ContextOverflow(ids.size(), e.max_context); const int limit = std::clamp(max_tokens, 0, room); std::lock_guard lock(e.mu); @@ -332,11 +380,51 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res std::fprintf(stderr, "[prefix] slot=%d cached=%d of %zu reason=%s\n", stats.prefix.slot, stats.prefix.cached_tokens, ids.size(), stats.prefix.reason.c_str()); } + // (engine#1785) Where this request should leave a pinned snapshot, as a position: the state is saved + // one token short of the shared part for the same reason the whole-prompt snapshot below is saved one + // token short of the prompt. -1 = no pin for this request. pin_point() is 0 when a slot (the pin + // itself, on every request after the one that made it) already covers the shared part, and it is + // always less than the prompt's length, so this cut is never the last-token cut below. + int pin_cut = -1; + size_t pin_tokens = 0; + if (e.prefix_pin && e.prefix && e.prefix->enabled()) { + // The end of the system turn first: it is known on the first request and it is the right place. Only + // a request with no system message falls back to what it shares with the previous request. + if (!system_prompt.empty()) { + std::vector lead = e.tok->encode(system_prompt); + if (chat && e.add_bos && e.bos_id >= 0 && (lead.empty() || lead.front() != e.bos_id)) + lead.insert(lead.begin(), e.bos_id); + pin_tokens = e.prefix->pin_at(ids, npu::shared_prefix_len(lead, ids)); + } else { + pin_tokens = e.prefix->pin_point(e.prev_ids, ids); + } + if (pin_tokens > 0 && int(pin_tokens) - 1 >= resume) pin_cut = int(pin_tokens) - 1; + } { const int chunk_max = std::max(1, e.fwd->max_batch_rows()); const int n = int(ids.size()); for (int off = resume; off < n;) { - const int take = std::min(chunk_max, n - off); + int take = std::min(chunk_max, n - off); + if (pin_cut >= 0 && off == pin_cut) { + Q4nxNpuForward::StateSnapshot snap; + if (e.fwd->snapshot_state(off, snap)) { + std::string why; + const std::vector shared(ids.begin(), ids.begin() + long(pin_tokens)); + const int slot = e.prefix->pin(shared, snap.bytes(), &why); + if (slot >= 0 && size_t(slot) < e.prefix_state.size()) e.prefix_state[size_t(slot)] = std::move(snap); + std::fprintf(stderr, "[prefix] pin slot=%d tokens=%zu state_at=%d %s\n", slot, pin_tokens, off, + why.c_str()); + } + pin_cut = -1; + } else if (pin_cut > off && off + take > pin_cut) { + // Stop this chunk at the pin position; the next pass saves the state there. One more launch. + take = pin_cut - off; + const std::vector head(ids.begin() + off, ids.begin() + off + take); + if (!e.fwd->step_batch(head, off, logits)) + throw std::runtime_error(std::string("forward step failed: ") + e.fwd->last_error()); + off += take; + continue; + } // The state the prefix cache keeps is taken one token short of the prompt, because the answer is // drawn from the last position's logits and a snapshot holds state, not logits -- so the snapshot // is saved just before the last token and that token is run again on a resume. On a GDN model @@ -368,6 +456,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res off += take; } } + if (e.prefix_pin) e.prev_ids = ids; // Prompt processing ends here: every prompt token has been through the model and the first // token's distribution is in `logits`. @@ -439,7 +528,8 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res bool stopped_at_eos = false, stopped_at_stop = false; int produced = 0; for (int n = 0; n < limit; n++) { - if (req.is_connection_closed()) break; // the client is gone: free the NPU + // the client is gone: free the NPU (the bridge knows first when a write to it failed) + if (req.is_connection_closed() || (bridge && bridge->cancelled())) break; if (gram) refresh_allowed(); // Raw distribution for a reference comparison: sample_token() takes logits by const // reference and the copy below happens before it runs, so what lands on disk is what @@ -447,12 +537,12 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res if (!e.dump_logits.empty()) { const std::string path = e.dump_logits + "/step_" + std::to_string(dump_req) + "_" + std::to_string(n) + ".bin"; - if (std::FILE* f = std::fopen(path.c_str(), "wb")) { + if (std::FILE* f = onebit_fopen_private_wb(path.c_str())) { // 0600, no symlink (file_create.h) std::fwrite(logits.data(), sizeof(float), logits.size(), f); std::fclose(f); } } - const int next = npu::sample_token(logits, opts.sample, rng, gram ? &allowed : nullptr); + const int next = npu::sample_token(logits, opts.sample, rng, gram ? &allowed : nullptr, e.tok->size()); // -1: the grammar allows no token here, so the answer ends where it is. An end token is // in the mask only where the grammar can stop (refresh_allowed), so a constrained reply // that completes reports "stop" instead of padding to the budget (engine#1182). @@ -466,6 +556,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // pre-top_k/top_p/temperature and the grammar mask cannot narrow it. if (want_lp) held.emplace_back(Emitted{piece, npu::logprobs_of(logits, next, opts.sample.top_logprobs)}, piece.size()); + stats.pieces.push_back(piece); // engine#2111: reasoning_tokens is counted over these ++produced; const size_t stop_at = npu::find_stop(pending, opts.stop); if (stop_at != std::string::npos) { @@ -514,13 +605,19 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // OpenAI's usage object plus the two fields issue #650 reports: how many prompt tokens came // from a saved prefix, and what the cache decided (which slot, and why). auto usage_json = [](const CallStats& st) { - return json{{"prompt_tokens", st.prompt_tokens}, - {"completion_tokens", st.completion_tokens}, - {"total_tokens", st.prompt_tokens + st.completion_tokens}, - {"prompt_tokens_details", {{"cached_tokens", st.prefix.cached_tokens}}}, - {"flm_prefix_cache", {{"slot", st.prefix.slot}, - {"cached_tokens", st.prefix.cached_tokens}, - {"reason", st.prefix.reason}}}}; + json u{{"prompt_tokens", st.prompt_tokens}, + {"completion_tokens", st.completion_tokens}, + {"total_tokens", st.prompt_tokens + st.completion_tokens}, + {"prompt_tokens_details", {{"cached_tokens", st.prefix.cached_tokens}}}, + {"flm_prefix_cache", {{"slot", st.prefix.slot}, + {"cached_tokens", st.prefix.cached_tokens}, + {"reason", st.prefix.reason}, + {"pinned", st.prefix.pinned}}}}; + // engine#2111: OpenAI's field for "how much of the budget went into reasoning". With content "" and + // finish_reason "length", reasoning_tokens == completion_tokens is the sentence that says the budget ran out + // inside a thinking block, rather than that the model said nothing. Additive: absent where nothing was split. + if (st.reasoning_tokens >= 0) u["completion_tokens_details"] = {{"reasoning_tokens", st.reasoning_tokens}}; + return u; }; // OpenAI's logprobs.content entry: the token's text, its log-probability, and the most likely @@ -541,6 +638,9 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // markup whatever the container is (app/npu_tool_grammar.h), so that reply is split as qwen3. Every // other reply is in the container's own markup. const npu::AnswerFamily family = opts.tool_grammar ? npu::AnswerFamily::qwen3 : e.answer_family; + // A template whose generation prompt already opens the reasoning block leaves the answer inside it + // (npu::prompt_opens_reasoning). Only a rendered chat prompt can: an ids prompt and /v1/completions are raw. + const bool starts_in_reasoning = chat && !prompt_is_ids && npu::prompt_opens_reasoning(prompt, family); auto wrap = [&](const std::string& content, const std::string& finish, const std::vector& lps) { json out; out["id"] = id; @@ -551,7 +651,8 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // (issue #652 item 3) one shared parser splits what the model wrote into OpenAI's three // parts; a reply that is only tool calls says content: null. const npu::ParsedAnswer ans = - npu::parse_answer(content, family, opts.tools_sent); + npu::parse_answer(content, family, opts.tools_sent, starts_in_reasoning); + stats.reasoning_tokens = npu::reasoning_token_count(stats.pieces, family, opts.tools_sent, starts_in_reasoning); json message{{"role", "assistant"}}; message["content"] = ans.content_null ? json(nullptr) : json(ans.content); if (!ans.reasoning.empty()) message["reasoning_content"] = ans.reasoning; @@ -572,7 +673,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res std::vector texts; texts.reserve(lps.size()); for (const auto& x : lps) texts.push_back(x.text); - const std::vector keep = npu::content_mask(texts, family, opts.tools_sent); + const std::vector keep = npu::content_mask(texts, family, opts.tools_sent, starts_in_reasoning); std::vector content_lps; for (size_t i = 0; i < lps.size(); i++) if (i < keep.size() && keep[i]) content_lps.push_back(lps[i]); @@ -610,10 +711,16 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res chunk_tpl["object"] = "chat.completion.chunk"; chunk_tpl["created"] = created; chunk_tpl["model"] = e.id; + // With a bridge, each frame leaves as soon as it exists (the connection thread is already writing them); + // without one they are collected and sent when the handler returns, as this route always did. auto chunks = std::make_shared>(); + auto send = [&](std::string frame) { + if (bridge) bridge->push(std::move(frame)); + else chunks->push_back(std::move(frame)); + }; // The same splitter as the non-streamed path, fed piece by piece, so a streamed reply // and a non-streamed one cannot disagree about what is reasoning and what is content. - npu::AnswerStream stream(family, opts.tools_sent); + npu::AnswerStream stream(family, opts.tools_sent, starts_in_reasoning); auto emit_pieces = [&](const std::string& t, const std::vector& lp) { std::vector> pieces; stream.feed(t, pieces); @@ -628,7 +735,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res // carries no entries. if (!is_reasoning && opts.sample.logprobs && !lp.empty()) chunk["choices"][0]["logprobs"] = logprobs_json(lp); - chunks->push_back("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); + send("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); } return true; }; @@ -640,7 +747,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res json chunk = chunk_tpl; json delta = is_reasoning ? json{{"reasoning_content", text}} : json{{"content", text}}; chunk["choices"] = json::array({{{"index", 0}, {"delta", delta}, {"finish_reason", nullptr}}}); - chunks->push_back("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); + send("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); } if (!stream.calls().empty()) { // The calls arrive whole, in one chunk: OpenAI streams argument fragments, and a @@ -656,7 +763,7 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res json delta{{"tool_calls", calls}}; if (!chunk_tpl.contains("object") || chunk_tpl["object"].is_null()) chunk["object"] = "chat.completion.chunk"; chunk["choices"] = json::array({{{"index", 0}, {"delta", delta}, {"finish_reason", "tool_calls"}}}); - chunks->push_back("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); + send("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); } // OpenAI's stream ends with one choice that carries an empty delta and the reason the // generation ended; without it a streaming client cannot tell a `length` cut from a @@ -668,15 +775,23 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res chunk["choices"] = json::array({{{"index", 0}, {"delta", json::object()}, {"finish_reason", finish}}}); - chunks->push_back("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); + send("data: " + chunk.dump(-1, ' ', false, json::error_handler_t::replace) + "\n\n"); } } { json tail = chunk_tpl; tail["choices"] = json::array(); + // engine#2111: the same count as the non-streamed reply's, over the same pieces. + if (chat) + stats.reasoning_tokens = + npu::reasoning_token_count(stats.pieces, family, opts.tools_sent, starts_in_reasoning); tail["usage"] = usage_json(stats); tail["timings"] = timings_json(stats); - chunks->push_back("data: " + tail.dump() + "\n\n"); + send("data: " + tail.dump() + "\n\n"); + } + if (bridge) { + bridge->push("data: [DONE]\n\n"); + return; } auto cursor = std::make_shared(0); res.set_chunked_content_provider("text/event-stream", @@ -692,12 +807,86 @@ void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Res return true; }); } + } catch (const ContextOverflow& ex) { + // The client's request, not a server fault: 400, with the type and the wording the GGUF route's + // backend uses for the same refusal, so one client-side check covers both routes. + res.status = 400; + res.set_content(json{{"error", {{"message", ex.what()}, {"type", "exceed_context_size_error"}, + {"n_prompt_tokens", ex.prompt_tokens}, {"n_ctx", ex.context}}}}.dump(), + "application/json"); } catch (const std::exception& ex) { res.status = 500; res.set_content(json{{"error", {{"message", ex.what()}}}}.dump(), "application/json"); } } +// The route's entry. A request that does not ask for a stream is answered in place. One that does is generated on +// its own thread, so its frames can be written while later tokens are still being produced (app/stream_bridge.h): +// this thread waits for the first frame or the end, and only then decides between an ordinary response -- the 400s +// and 500s a request can earn before it has produced anything keep their status codes -- and a stream. +void handle_generate(ForwardEngine& e, const httplib::Request& req, httplib::Response& res, bool chat) { + // A non-causal decision pack has no decode to run (npu::decision_only_refusal): say so before any work. + if (e.fwd && e.fwd->noncausal()) { + res.status = 400; + res.set_content(npu::decision_only_refusal(e.id, e.head != nullptr).dump(), "application/json"); + return; + } + bool wants_stream = false; + try { + const json body = json::parse(req.body); + wants_stream = body.is_object() && body.contains("stream") && body["stream"].is_boolean() && body["stream"].get(); + } catch (const std::exception&) { + // not JSON: the implementation answers that itself, without a stream + } + if (!wants_stream) { + handle_generate_impl(e, req, res, chat, nullptr); + return; + } + auto bridge = std::make_shared(); + auto early = std::make_shared(); // what the generation answers when it never starts a stream + // `req` outlives this handler: httplib keeps it until the response, content provider included, has been written, + // and the releaser below joins the thread before that. + auto worker = std::make_shared([&e, &req, chat, bridge, early]() { + try { + handle_generate_impl(e, req, *early, chat, bridge.get()); + } catch (const std::exception& ex) { + early->status = 500; + early->set_content(json{{"error", {{"message", ex.what()}}}}.dump(), "application/json"); + } catch (...) { + early->status = 500; + early->set_content(R"({"error":{"message":"generation failed"}})", "application/json"); + } + // A failure after the stream began cannot change the status any more: say so in the stream and end it. + // One that came before any frame pushes nothing, so the handler below still answers it as an ordinary error. + if (early->status >= 400 && bridge->started()) bridge->push("data: " + early->body + "\n\ndata: [DONE]\n\n"); + bridge->finish(); + }); + if (!bridge->wait_first()) { + worker->join(); + res.status = early->status > 0 ? early->status : 500; + res.set_content(early->body, early->has_header("Content-Type") ? early->get_header_value("Content-Type") + : std::string("application/json")); + return; + } + res.set_chunked_content_provider( + "text/event-stream", + [bridge](size_t, httplib::DataSink& sink) { + std::string frame; + while (bridge->next(frame)) { + if (!sink.write(frame.data(), frame.size())) { + bridge->cancel(); // the client went away: the generation stops at its next token + return false; + } + } + sink.done(); + return true; + }, + [bridge, worker](bool) { + bridge->cancel(); + if (worker->joinable()) worker->join(); + }); +} + // The per-question input limit of the decision route: the text path's own 32,768 tokens, or the // backbone's context when that is smaller. Either overflowing is the same failure to the model, so // the smaller one is the limit the request is checked against. @@ -811,6 +1000,10 @@ int run_forward_ids(const std::string& model_dir, const std::string& kernels, co if (ids.empty()) throw std::runtime_error("--ids: no token ids"); const std::string model = model_dir + "/model.q4nx"; if (!std::filesystem::exists(model)) throw std::runtime_error(model + " does not exist"); + // The tokenizer's size bounds the printed ids (engine#2017); a container without tokenizer.json keeps the whole row. + int tok_vocab = 0; + if (const std::string tj = model_dir + "/tokenizer.json"; std::filesystem::exists(tj)) + tok_vocab = npu::Tokenizer(tj).size(); // Same device discipline as forward-serve: host-GEMM runs take no lock (engine#1353 -- a host-GEMM run that held it // for 4 h 52 m stopped two lanes' device work), a run that opens the device takes the exclusive one. int lock_fd = -1; @@ -841,10 +1034,8 @@ int run_forward_ids(const std::string& model_dir, const std::string& kernels, co " floats for " + std::to_string(m) + " rows)"); const int nv = int(all.size() / size_t(m)); for (int r = 0; r < m; ++r) { - const float* row = all.data() + size_t(r) * size_t(nv); - int best = 0; - for (int v = 1; v < nv; ++v) - if (row[v] > row[best]) best = v; + // Bounded by the tokenizer like the served route (engine#2017): an id past it is a padding position. + const int best = npu::argmax_token(all.data() + size_t(r) * size_t(nv), size_t(nv), tok_vocab); std::printf("%d\n", best); // token ids out, in row order: the id the model chooses at that row } done += m; @@ -868,6 +1059,8 @@ int run_forward_serve(int argc, char** argv) { // instead of being ignored and served by whatever target happened to be linked. std::string npu_target = ::npu::target::xdna::XdnaTarget::kName; npu::PrefixOptions prefix_opts; // --prefix-slots / --prefix-cache-mb / --prefix-min-gain (#650) + bool prefix_pin = false; // --prefix-pin (engine#1785) + int ctx_size = 0; // --ctx-size N: prompt plus answer, at most the model's own (0 = the model's) for (int i = 0; i < argc; i++) { const std::string a = argv[i]; auto next = [&]() -> std::string { @@ -882,6 +1075,8 @@ int run_forward_serve(int argc, char** argv) { else if (a == "--prefix-slots") prefix_opts.slots = std::stoi(next()); else if (a == "--prefix-cache-mb") prefix_opts.cache_mb = std::stoul(next()); else if (a == "--prefix-min-gain") prefix_opts.min_gain = std::stoi(next()); + else if (a == "--prefix-pin") prefix_pin = true; + else if (a == "-c" || a == "--ctx-size") ctx_size = std::stoi(next()); else if (a == "--npu-streams") streams = std::stoi(next()); else if (a == "--npu-target") npu_target = next(); else if (a == "--dump-logits") dump_logits = next(); @@ -920,10 +1115,21 @@ int run_forward_serve(int argc, char** argv) { e.bos_id = json_int(cfg, "bos_token_id", -1); e.add_bos = json_int(cfg, "add_bos_token", 0) != 0; e.eos = npu::eos_token_ids_from_config(cfg); + const size_t eos_from_config = e.eos.size(); e.endoftext = e.tok->token_id("<|endoftext|>"); // q4nx_config_int finds the key at any depth (a text_config's too). e.max_context = q4nx_config_int(cfg.c_str(), "max_position_embeddings", 0); if (e.max_context <= 0) e.max_context = 32768; + // --ctx-size lowers the limit and never raises it: the model's positions are what they are. The same + // guard that refuses a prompt longer than the model's context then refuses one longer than this. + if (ctx_size < 0) throw std::runtime_error("--ctx-size: must be zero or more"); + if (ctx_size > 0) { + if (ctx_size > e.max_context) + std::fprintf(stderr, "1bit forward-serve: --ctx-size %d is above the model's context %d; using %d\n", ctx_size, + e.max_context, e.max_context); + e.max_context = std::min(e.max_context, ctx_size); + } + std::fprintf(stderr, "1bit forward-serve: context %d tokens (prompt plus answer)\n", e.max_context); // #963: the chat route's prompt was ChatML for every container, and a tokenizer with no // <|im_start|>/<|im_end|> added token cannot represent it -- the marker becomes ordinary pieces and @@ -965,6 +1171,38 @@ int run_forward_serve(int argc, char** argv) { (specials.empty() ? std::string() : " (model specials: " + specials + ")"); std::fprintf(stderr, "1bit forward-serve: chat route disabled: %s\n", e.chat_refusal.c_str()); } + // What ends a generation is more than config.json's eos_token_id (npu::stop_token_ids): the tokenizer's own + // eos_token, and, when this route closes its prompt turns with ChatML's marker, that marker. Decided here, + // after the template, because whether the route renders ChatML is known only now. + { + const bool chatml_route = !can_render && e.chat_ok; + // The marker this route's prompts close an assistant turn with. On the ChatML fallback it is <|im_end|>. + // With the container's own template it is read off the template: render a two-message exchange and take + // what follows the assistant's text (npu::turn_end_marker). A template that cannot render that exchange + // contributes nothing, and the load line below shows what the serve stops on either way. + std::string turn_end = chatml_route ? std::string("<|im_end|>") : std::string(); + if (can_render) { + static const char* const kProbe = "\xC2\xA4turn-end-probe\xC2\xA4"; + try { + nlohmann::ordered_json exchange = nlohmann::ordered_json::array(); + exchange.push_back({{"role", "user"}, {"content", "u"}}); + exchange.push_back({{"role", "assistant"}, {"content", kProbe}}); + const std::string rendered = npu::render_chat_template(e.tmpl, exchange, /*add_generation_prompt=*/false, + nlohmann::ordered_json::object(), e.bos_text, e.eos_text); + const size_t at = rendered.find(kProbe); + if (at != std::string::npos) turn_end = npu::turn_end_marker(rendered.substr(at + std::strlen(kProbe))); + } catch (const std::exception& ex) { + std::fprintf(stderr, "1bit forward-serve: the template did not render a two-message exchange (%s); no turn-end marker from it\n", ex.what()); + } + } + e.eos = npu::stop_token_ids(npu::read_json_or_null(cfg), npu::read_json_or_null(dir + "/tokenizer_config.json"), + turn_end, [&](const std::string& text) { return e.tok->token_id(text); }); + std::string named; + for (size_t i = 0; i < e.eos.size(); i++) + named += std::string(i ? ", " : "") + std::to_string(e.eos[i]) + " " + e.tok->decode(e.eos[i]) + + (i < eos_from_config ? "" : " (tokenizer / chat turn end)"); + std::fprintf(stderr, "1bit forward-serve: stops on %s\n", named.empty() ? "<|endoftext|> only" : named.c_str()); + } // The template variables `{{ bos_token }}`/`{{ eos_token }}` render: the strings the tokenizer // actually has (a null bos_token in tokenizer_config.json is "", and the engine's -1 is the same // fact), not a second read of the JSON that could disagree with the tokenizer. @@ -979,6 +1217,8 @@ int run_forward_serve(int argc, char** argv) { std::fprintf(stderr, "1bit forward-serve: prefix cache on: %d slots, %zu MB, min gain %d tokens\n", prefix_opts.slots, prefix_opts.cache_mb, prefix_opts.min_gain); + e.prefix_pin = prefix_pin; + if (prefix_pin) std::fprintf(stderr, "1bit forward-serve: prefix pin on: the part a request shares with the previous one keeps a slot\n"); } e.fwd = std::make_unique(); @@ -1161,6 +1401,13 @@ int run_forward_serve(int argc, char** argv) { // The target that is actually serving, so a log says which one answered rather than leaving it to the build. std::fprintf(stderr, " [npu-target=%s]", npu_target.c_str()); std::fprintf(stderr, "\n"); + // engine#2102: the machine line `[forward] decode family: …` stays byte-for-byte (gates grep it). This is the sentence + // an operator needs at the ready line's level when it says 0 of N -- the route that answers correctly and 50x slower. + { + const std::string note = onebit::npu::decode_family_note(e.fwd->decode_roles(), e.fwd->design_roles(), kernels, + e.fwd->decode_disabled(), e.fwd->decode_image()); + if (!note.empty()) std::fprintf(stderr, "%s\n", note.c_str()); + } while (true) std::this_thread::sleep_for(std::chrono::seconds(3600)); return 0; diff --git a/app/main.cpp b/app/main.cpp index e9249b378..1d7ae3755 100644 --- a/app/main.cpp +++ b/app/main.cpp @@ -43,6 +43,7 @@ int onebit_ort_run_main(int argc, char** argv); #endif #ifdef ONEBIT_NPU_HOST #include "pack_import.h" // 1bit npu-pack: an OGA NPU pack -> a Q4NX model directory +#include "vision_pack.h" // 1bit npu-pack vision: a pack's vision component (engine#1788) #endif #ifdef ONEBIT_NPU #include "forward_serve.h" @@ -94,6 +95,7 @@ void usage(FILE* out) { #endif #ifdef ONEBIT_NPU_HOST " npu-pack import a Ryzen AI OGA NPU pack (npu-pack --help)\n" + " npu-pack vision read a pack's vision component (vision_weight.q4nx)\n" #endif " comfy run a ComfyUI workflow with ComfyUI.cpp (docs/comfyui.md)\n" #ifdef ONEBIT_LAYA @@ -323,6 +325,18 @@ int main(int argc, char** argv) { #ifdef ONEBIT_NPU_HOST if (cmd == "npu-pack") { try { + // `npu-pack vision `: the vision component of a two-component pack (engine#1788), read by + // npu/vision_pack.h rather than the OGA importer, which knows nothing of it. + if (argc >= 3 && std::string(argv[2]) == "vision") { + if (argc != 4 || argv[3][0] == '-') { + std::fprintf(stderr, "usage: 1bit npu-pack vision \n" + " reads the pack's vision component (config.json's vision_model_weight, default\n" + " vision_weight.q4nx): every tensor placed, the tower's geometry from the tensors'\n" + " own shapes, and each config.json field the tensors overrule\n"); + return argc == 4 ? 0 : 2; + } + return onebit::npu::vision::print_info(argv[3]); + } return onebit::npu::pack_import::run(argc - 2, argv + 2); } catch (const std::exception& e) { std::fprintf(stderr, "1bit npu-pack: %s\n", e.what()); diff --git a/app/npu_answer_parser.h b/app/npu_answer_parser.h index 305a6a2c1..f6453fcb0 100644 --- a/app/npu_answer_parser.h +++ b/app/npu_answer_parser.h @@ -142,8 +142,11 @@ inline AnswerFamily answer_family_of_template(const std::string& template_text) // (is_reasoning, text) pieces; finish() flushes what the holdback still held and reports the calls. class AnswerStream { public: - AnswerStream(AnswerFamily family, bool tools_sent) - : fam_(family), tags_(answer_detail::tags_of(family)), tools_sent_(tools_sent) {} + // `starts_in_reasoning`: the PROMPT already opened the reasoning block (prompt_opens_reasoning below), so the + // answer begins inside it and the first thing to look for is the closing tag. + AnswerStream(AnswerFamily family, bool tools_sent, bool starts_in_reasoning = false) + : fam_(family), tags_(answer_detail::tags_of(family)), tools_sent_(tools_sent), + mode_(starts_in_reasoning ? Mode::reasoning : Mode::content) {} void feed(const std::string& chunk, std::vector>& out) { buf_ += chunk; @@ -169,6 +172,9 @@ class AnswerStream { // The calls seen so far, in order. A call needs tools_sent, per the issue's rule. const std::vector& calls() const { return calls_; } bool saw_reasoning() const { return saw_reasoning_; } + // Inside a reasoning block right now (after the opening tag, before the closing one). reasoning_token_count() + // reads it before and after each token, so the tags themselves and held-back partial tags count with the block. + bool in_reasoning() const { return mode_ == Mode::reasoning; } bool saw_text() const { return saw_text_; } const std::string& reasoning() const { return reasoning_; } @@ -177,8 +183,22 @@ class AnswerStream { private: enum class Mode { content, reasoning, call }; - void emit(std::vector>& out, bool is_reasoning, const std::string& text) { - if (text.empty()) return; + // Whitespace around a reasoning block is the model's layout, not part of either text: `\n...\n\n\n` + // then the answer. Left in, the answer to "one word" comes back as "\n\nParis", and a model that thinks nothing + // (`\n\n`) sends a reasoning chunk that is two newlines. So the whitespace that LEADS the reasoning + // is dropped, and so is the whitespace that leads the content when that content follows a reasoning block -- + // the rule app/think_split.h has always applied on the in-process route ("whitespace around the block is + // dropped"), so the two routes return the same strings. Content with no reasoning before it is not touched. + void emit(std::vector>& out, bool is_reasoning, const std::string& raw) { + if (raw.empty()) return; + std::string text = raw; + bool& trim = is_reasoning ? trim_reasoning_ : trim_content_; + if (trim) { + const size_t b = text.find_first_not_of(" \t\r\n"); + if (b == std::string::npos) return; // all whitespace so far: nothing to send yet + text.erase(0, b); + trim = false; + } if (is_reasoning) { reasoning_ += text; saw_reasoning_ = true; @@ -218,6 +238,8 @@ class AnswerStream { } if (starts(tags_.think_close)) { flush_plain(); + // what follows a block is separated from it by layout whitespace -- unless content came first + if (mode_ == Mode::reasoning && !saw_text_) trim_content_ = true; mode_ = Mode::content; i += std::string(tags_.think_close).size(); continue; @@ -374,14 +396,32 @@ class AnswerStream { std::vector calls_; bool saw_reasoning_ = false; bool saw_text_ = false; + bool trim_reasoning_ = true; // until the reasoning's first non-whitespace character + bool trim_content_ = false; // set when a reasoning block closes before any content; cleared by the first text }; +// Does the rendered prompt leave the model INSIDE a reasoning block? Some templates end their generation prompt +// with the opening tag itself -- Qwen3.5's ends `<|im_start|>assistant\n\n` when thinking is on -- so the +// model never writes that tag: its answer starts with the reasoning and the first tag in it is the closing one. +// A splitter that waits to see the block opened then files the whole reasoning under `content` and drops the +// closing tag it never saw opened, so nothing in the reply marks where the answer starts (measured on FastFlowLM's +// Qwen3.5-2B-NPU2, engine#1786). True when the prompt, less trailing whitespace, +// ends with the family's opening tag. A prompt that opens and closes the block (thinking off: +// `\n\n\n\n`) does not end with it, and neither does one that never mentions it. +inline bool prompt_opens_reasoning(const std::string& prompt, AnswerFamily family) { + const std::string open = answer_detail::tags_of(family).think_open; + size_t end = prompt.size(); + while (end > 0 && (prompt[end - 1] == ' ' || prompt[end - 1] == '\n' || prompt[end - 1] == '\r' || prompt[end - 1] == '\t')) --end; + return !open.empty() && end >= open.size() && prompt.compare(end - open.size(), open.size(), open) == 0; +} + // Which of `pieces` (one generated token's text each, in order) belong to `content` -- not to // reasoning and not to call markup. Callers that report a per-token value use this so reasoning and // call markup get no entry (issue #651 item 3). A token that straddles a boundary ("The") // counts as content: its text did reach the caller. -inline std::vector content_mask(const std::vector& pieces, AnswerFamily family, bool tools_sent) { - AnswerStream s(family, tools_sent); +inline std::vector content_mask(const std::vector& pieces, AnswerFamily family, bool tools_sent, + bool starts_in_reasoning = false) { + AnswerStream s(family, tools_sent, starts_in_reasoning); std::vector mask; mask.reserve(pieces.size()); std::vector> out; @@ -405,9 +445,61 @@ inline std::vector content_mask(const std::vector& pieces, An return mask; } +// How many of `pieces` (one generated token's text each, in order) were reasoning: OpenAI's +// usage.completion_tokens_details.reasoning_tokens (engine#2111). An empty `content` with finish_reason "length" reads +// like "the model said nothing" when the whole budget went into a thinking block, and this is the number that says so. +// A token is reasoning when none of its text reached `content` and the block was open before it, after it, or it put +// text into the block: so the opening and closing tags, a partial tag the holdback keeps, and an unclosed block at the +// end of the budget all count, while call markup and the whitespace after a closed block do not. A token that +// straddles the boundary ("The") is content, as content_mask() says. Same splitter, so the count, the split +// and the logprob entries cannot disagree. +inline int reasoning_token_count(const std::vector& pieces, AnswerFamily family, bool tools_sent, + bool starts_in_reasoning = false) { + AnswerStream s(family, tools_sent, starts_in_reasoning); + std::vector> out; + int n = 0; + bool last_counted = false; + // Tokens that produced nothing before anything was produced (leading whitespace, the first half of an opening tag + // the holdback keeps) belong to whatever comes next: to the block if the next token opens it, to nothing otherwise. + int pending = 0; + bool started = false; + for (const auto& piece : pieces) { + const bool before = s.in_reasoning(); + out.clear(); + s.feed(piece, out); + bool content = false, reasoning = false; + for (const auto& [is_reasoning, text] : out) { + (void)text; + (is_reasoning ? reasoning : content) = true; + } + last_counted = !content && (reasoning || before || s.in_reasoning()); + if (last_counted) { + ++n; + if (!started) n += pending; + started = true; + } else if (content) { + started = true; + } else if (!started) { + ++pending; + } + } + // What the holdback still holds is flushed by finish(): if it is content, the last token's text reached `content`. + out.clear(); + s.finish(out); + for (const auto& [is_reasoning, text] : out) { + (void)text; + if (!is_reasoning && last_counted) { + --n; + last_counted = false; + } + } + return n; +} + // The whole answer at once: the same splitter, so the two paths cannot disagree. -inline ParsedAnswer parse_answer(const std::string& text, AnswerFamily family, bool tools_sent) { - AnswerStream s(family, tools_sent); +inline ParsedAnswer parse_answer(const std::string& text, AnswerFamily family, bool tools_sent, + bool starts_in_reasoning = false) { + AnswerStream s(family, tools_sent, starts_in_reasoning); std::vector> pieces; s.feed(text, pieces); s.finish(pieces); diff --git a/app/npu_chatml.h b/app/npu_chatml.h index ed8094357..7ec367fa1 100644 --- a/app/npu_chatml.h +++ b/app/npu_chatml.h @@ -61,6 +61,31 @@ inline std::string chatml(const nlohmann::json& messages, bool thinking = true) return p; } +// The messages a chat route renders to find where a request's system turn ends (--prefix-pin, engine#1785): the +// request's leading system messages, then one user message whose text no question starts with. The rendering of +// these and the rendering of the real request agree up to the first token of the user's text and differ there, so +// their common prefix is the part every question under this system prompt shares. +// +// The user message is there because a template may refuse to render without one. Measured: the Qwen3.5 template +// rendered alone on its system message throws, and the early pin was silently not made on that container. +// +// Empty array when the request does not open with a system message: there is nothing to pin early. +template +Json pin_probe_messages(const Json& messages) { + Json lead = Json::array(); + if (!messages.is_array()) return lead; + for (const auto& m : messages) { + if (!m.is_object() || !m.contains("role") || !m["role"].is_string() || m["role"].template get() != "system") break; + lead.push_back(m); + } + if (lead.empty()) return lead; + Json probe = Json::object(); + probe["role"] = "user"; + probe["content"] = "\xC2\xA4\xC2\xA4 pin probe \xC2\xA4\xC2\xA4"; + lead.push_back(probe); + return lead; +} + // Whether the container's tokenizer can represent the prompt above. `chatml()` is rendered for // every container (#832), and a tokenizer without the anchors spells `<|im_start|>` as ordinary // pieces, so the model is prompted with text its own template never produces (#963: on diff --git a/app/npu_decode_note.h b/app/npu_decode_note.h new file mode 100644 index 000000000..0f208661c --- /dev/null +++ b/app/npu_decode_note.h @@ -0,0 +1,64 @@ +// 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 sentence a serve owes an operator when it is about to run the slow route (engine#2102). +// +// The fact already exists as a machine line — `[forward] decode family: 0 role(s) at M=16, 4 on M=128` — and that line is +// grepped byte-for-byte by the gates, so it must not be reworded. What was missing is a sentence at the ready line's level: +// a serve with no 16-row decode kernels runs every decode step as prefill-width launches, measured at 1.9 tok/s where the +// fused route on the same box does 88-101, and nothing said so. This header is the whole decision, as a pure function, so it +// can be armed without a device and so the "say nothing" case is as testable as the "say something" one. +// +// It deliberately returns "" for the cases it cannot judge (a host-GEMM arm has no family at all; a count of 0 of 0 is a +// different fault, reported elsewhere) — a warning that fires on healthy serves is a warning nobody reads. + +#pragma once + +#include +#include + +namespace onebit::npu { + +// decode_roles: roles that got an M=16 decode kernel. design_roles: roles the model has (0 on the host-GEMM arm, which loads +// no decode family). kernels_dir: what the operator pointed at, echoed so the message can be acted on without reading the +// rest of the log. disabled: NPU_DECODE_M16=0 turned the family off, so nothing is missing. image: a ternary (pb16i/t2s) +// container, whose loader refuses a missing decode kernel; the roles it leaves on M=128 are ones its geometry or the split +// designs cover (named per role in the load log), so for it only 0 of N is this issue. +// +// The way out names the fused route, not the _m16 designs: measured on Qwen3-0.6B (#2102), the per-op forward with all four +// twins decodes at ~2.7 tok/s against ~1.7 without them and 88-101 on the fused route, so the twins are not the demo rate. +inline std::string decode_family_note(int decode_roles, int design_roles, const std::string& kernels_dir, + bool disabled = false, bool image = false) { + if (design_roles <= 0) return {}; // no roles (or the host arm): not this issue's fault, and no number to quote + const char* dir = kernels_dir.empty() ? "(no kernel directory)" : kernels_dir.c_str(); + char buf[640]; + if (disabled) { + std::snprintf(buf, sizeof buf, + "npu: NPU_DECODE_M16=0 turned the 16-row decode family off -- decode runs prefill-width launches on all %d " + "role(s) (engine#2102). Unset it to use the decode kernels in %s.", + design_roles, dir); + return std::string(buf); + } + if (decode_roles >= design_roles) return {}; // a full decode family: say nothing + if (image && decode_roles > 0) return {}; // a ternary family's M=128 roles are its geometry's, already named + std::snprintf(buf, sizeof buf, + "npu: %d of %d role(s) have a 16-row decode kernel in %s -- decode runs prefill-width launches and is far " + "slower than the fused route (engine#2102). The _m16 designs help only part of the way; the demo rate is the " + "fused route's (docs/npu.md, plan.json: decode.dirs).", + decode_roles, design_roles, dir); + return std::string(buf); +} + +} // namespace onebit::npu diff --git a/app/npu_prefix_cache.h b/app/npu_prefix_cache.h index b66562b1c..0412a676d 100644 --- a/app/npu_prefix_cache.h +++ b/app/npu_prefix_cache.h @@ -37,6 +37,7 @@ // What they share is the part before the first difference, so that part gets a slot of its own: // pin_point(prev, ids) how many leading tokens this request shares with the previous one, when that // is worth a slot and no slot covers it already +// pin_at(ids, n) the same decision for a point the caller knows already (the end of the system turn) // pin(prefix, bytes) keep the state at that point. One pinned slot at a time; the budget's eviction // and a longer prompt's replacement both leave it alone, a new pin() replaces it. // The rule reads token ids only, so it needs no knowledge of the template or of what a system prompt is. @@ -186,14 +187,25 @@ class PrefixCache { size_t pin_point(const std::vector& prev, const std::vector& ids) const { if (!enabled()) return 0; const size_t shared = shared_prefix_len(prev, ids); - if (shared == prev.size() || shared == ids.size()) return 0; + if (shared == prev.size()) return 0; + return pin_at(ids, shared); + } + + // The same decision for a point the caller already knows: `n` leading tokens of `ids` that every request of + // this kind will share -- the end of the system turn, which a chat route can find on the FIRST request by + // rendering the system messages alone, instead of waiting for a second request to compare with. Returns `n` + // when the state there is worth pinning, else 0: 0 when `n` is the whole prompt or more (nothing follows it, + // so insert() keeps it), when it adds fewer than prefix-min-gain tokens over the slot that already covers + // this request (the pin itself, on every request after the one that made it), and when the cache is off. + size_t pin_at(const std::vector& ids, size_t n) const { + if (!enabled() || n == 0 || n >= ids.size()) return 0; size_t covered = 0; // the longest live slot that is a prefix of this request for (const PrefixSlot& s : slots_) { if (!s.live || s.ids.size() > ids.size() || s.ids.size() <= covered) continue; if (std::equal(s.ids.begin(), s.ids.end(), ids.begin())) covered = s.ids.size(); } - if (shared <= covered || shared - covered < size_t(opt_.min_gain)) return 0; - return shared; + if (n <= covered || n - covered < size_t(opt_.min_gain)) return 0; + return n; } // Keep `prefix` (the first pin_point() tokens of a request) as the pinned slot. Returns the slot the diff --git a/app/npu_sampling.h b/app/npu_sampling.h index 4e2f41b6f..065005cb6 100644 --- a/app/npu_sampling.h +++ b/app/npu_sampling.h @@ -82,14 +82,25 @@ struct Candidate { int id = 0; }; +// How many leading ids of a logits row are tokens: the tokenizer's vocabulary when it is the +// smaller of the two. The head's width is the container's `vocab_size`, which may exceed the +// tokenizer's (Qwen3-0.6B: 151,936 for 151,669 ids); the ids in between are padding positions with +// no text, and a pick that lands there emits nothing and never ends the reply (engine#1182, +// #2017). `vocab` <= 0 means "no tokenizer to bound by" and keeps the whole row. +inline size_t token_bound(size_t width, int vocab) { + return (vocab > 0 && size_t(vocab) < width) ? size_t(vocab) : width; +} + // Draw one token from `logits`. `allowed`, when given, is a per-token mask (non-zero = // the grammar accepts the token); a mask with nothing set is refused with -1 and the -// caller reports it rather than decoding an unconstrained token. +// caller reports it rather than decoding an unconstrained token. `vocab` is the +// tokenizer's size (token_bound): ids at or past it are never drawn, with or without a mask. inline int sample_token(const std::vector& logits, const SampleParams& p, std::mt19937_64& rng, - const std::vector* allowed = nullptr) { + const std::vector* allowed = nullptr, int vocab = 0) { + const size_t n = token_bound(logits.size(), vocab); std::vector c; - c.reserve(logits.size()); - for (size_t i = 0; i < logits.size(); i++) { + c.reserve(n); + for (size_t i = 0; i < n; i++) { if (allowed && !(*allowed)[i]) continue; c.push_back({logits[i], int(i)}); } @@ -157,6 +168,16 @@ inline int sample_token(const std::vector& logits, const SampleParams& p, return c.back().id; } +// The greedy id of one logits row of `width` floats, bounded like sample_token(): a higher +// logit wins and a tie keeps the lower id. `1bit npu-run` prints this per row. +inline int argmax_token(const float* row, size_t width, int vocab = 0) { + const size_t n = token_bound(width, vocab); + int best = 0; + for (size_t v = 1; v < n; v++) + if (row[v] > row[size_t(best)]) best = int(v); + return best; +} + // One candidate of the raw distribution, as the client reads it: the id and its log-probability. struct Logprob { int id = 0; diff --git a/app/npu_stop_tokens.h b/app/npu_stop_tokens.h index aa910004c..281adf045 100644 --- a/app/npu_stop_tokens.h +++ b/app/npu_stop_tokens.h @@ -56,6 +56,90 @@ inline std::vector eos_token_ids(const nlohmann::json& cfg) { return ids; } +/// tokenizer_config.json's `eos_token` as text: a plain string, or the AddedToken object form +/// `{"content": "<|im_end|>", ...}`. Empty when the key is absent, null or another shape. +inline std::string eos_token_text(const nlohmann::json& tokenizer_config) { + if (!tokenizer_config.is_object() || !tokenizer_config.contains("eos_token")) return {}; + const nlohmann::json& e = tokenizer_config["eos_token"]; + if (e.is_string()) return e.get(); + if (e.is_object() && e.contains("content") && e["content"].is_string()) return e["content"].get(); + return {}; +} + +/// The marker a chat template closes an assistant turn with, read off the template's own output. `after` is the +/// text a template rendered immediately after an assistant message's content; the marker is the `<...>` token it +/// starts with (leading whitespace skipped), or empty when it starts with anything else. +/// +/// Why this exists: a model ends its turn with whatever its template ends an assistant turn with, and neither +/// config.json nor tokenizer_config.json has to name that token. FastFlowLM's Phi4-mini-Instruct-NPU2 names +/// <|endoftext|> in both, and its template closes every turn with <|end|> (200020); the upstream checkpoint +/// lists it in generation_config.json, which the pack does not ship. Served with only the two files' ids, the +/// answer was "Paris<|end|><|assistant|>**Follow-up Question 1:** ..." until the token budget ran out +/// (engine#1786). The template is the one file every chat pack has, and it cannot disagree with itself. +/// +/// "<|end|>" -> "<|end|>" (Phi) +/// "<|im_end|>\n" -> "<|im_end|>" (ChatML) +/// "<|eot_id|>" -> "<|eot_id|>" (Llama 3) +/// "\n" -> "" (Gemma) +/// "\n", "", "x"... -> "" unless it starts with a whole <...> token: plain text after a turn is not a stop. +inline std::string turn_end_marker(const std::string& after) { + size_t i = 0; + while (i < after.size() && (after[i] == ' ' || after[i] == '\n' || after[i] == '\r' || after[i] == '\t')) ++i; + if (i >= after.size() || after[i] != '<') return {}; + const size_t close = after.find('>', i); + if (close == std::string::npos || close - i < 2) return {}; + const std::string marker = after.substr(i, close - i + 1); + for (char c : marker) + if (c == ' ' || c == '\n' || c == '\r' || c == '\t') return {}; + return marker; +} + +/// Every id that ends a generation on a container, in the order they were found and without repeats: +/// +/// 1. config.json's `eos_token_id` (one integer or a list): what eos_token_ids() reads; +/// 2. tokenizer_config.json's `eos_token`, looked up in the tokenizer. The two files are written by +/// different tools and need not agree. FastFlowLM's Qwen3.5-2B container says `eos_token_id: 248044` +/// (<|endoftext|>) in config.json and `eos_token: "<|im_end|>"` (248046) in tokenizer_config.json, and the +/// model ends a chat turn with the second. With only the first, the answer came back as +/// "Item 3 is the 3D printer.<|im_end|>\n": the turn end served as text and one more token generated +/// after it (engine#1785's row on that container); +/// 3. `chat_turn_end`, the marker the chat route's prompt closes an assistant turn with: ChatML's <|im_end|> +/// on the fallback route, or turn_end_marker() of the container's own template when that is rendered. A +/// prompt whose turns end with that marker is answered by a turn that ends with it, whatever the two files +/// say. Pass an empty string when the route has no such marker. +/// +/// `id_of` maps a token's text to its id, or a negative number when the tokenizer has no such token; a name it +/// does not know adds nothing. +template +std::vector stop_token_ids(const nlohmann::json& config, const nlohmann::json& tokenizer_config, + const std::string& chat_turn_end, IdOf&& id_of) { + std::vector ids = eos_token_ids(config); + const auto add = [&](const std::string& text) { + if (text.empty()) return; + const int id = id_of(text); + if (id < 0) return; + for (int have : ids) + if (have == id) return; + ids.push_back(id); + }; + add(eos_token_text(tokenizer_config)); + add(chat_turn_end); + return ids; +} + +/// A JSON file, or null when it is missing or malformed: both stop-id sources are optional. +inline nlohmann::json read_json_or_null(const std::string& path) { + std::ifstream f(path); + if (!f) return nullptr; + nlohmann::json j; + try { + f >> j; + } catch (const std::exception&) { + return nullptr; + } + return j; +} + /// eos_token_ids() for a container's config.json. Unreadable file -> empty. inline std::vector eos_token_ids_from_config(const std::string& cfg_path) { std::ifstream f(cfg_path); diff --git a/app/npu_systemone.h b/app/npu_systemone.h index 22623f7cc..2a6c0bfd4 100644 --- a/app/npu_systemone.h +++ b/app/npu_systemone.h @@ -124,6 +124,22 @@ inline nlohmann::json systemone_question_error(const std::string& kind, const st {"message", message}}; } +// What a text route answers on a pack that cannot generate text. A decision pack that declares +// `attention_mode: noncausal_full_attention` encodes a whole prompt in one pass and reads an answer off the last +// position; there is no token-by-token decode to run, so `/v1/chat/completions` and `/v1/completions` can never +// succeed on it. They used to try, and failed inside the forward with HTTP 500 and the forward's own wording +// ("forward step failed: attention.mode noncausal_full_attention encodes a prompt in one batched pass, but this +// call starts at position 23 ..."; measured on FastFlowLM's D3-Nano-NPU2, engine#1786). That is the client's +// request against the wrong route, so it is a 400 that names the right one. `has_head` says whether the decision +// route is actually up on this serve, because pointing a client at a route that answers 503 is no help. +inline nlohmann::json decision_only_refusal(const std::string& model, bool has_head) { + const std::string where = has_head ? "send decisions to POST /v1/systemone" + : "and this serve has no decision head loaded, so it has no route for it"; + return {{"error", + {{"message", model + " is a decision model (non-causal attention): it does not generate text; " + where}, + {"type", "decision_model_error"}}}}; +} + inline std::string systemone_length_message(size_t tokens, size_t limit) { return "the question prompt has " + std::to_string(tokens) + " tokens, over the maximum context length of " + std::to_string(limit) + " tokens; nothing was truncated"; diff --git a/app/q4nx_cache.cpp b/app/q4nx_cache.cpp index 361974546..3bcb727c6 100644 --- a/app/q4nx_cache.cpp +++ b/app/q4nx_cache.cpp @@ -16,8 +16,10 @@ #include "q4nx_cache.h" #include +#include #include #include +#include #include #include @@ -83,3 +85,22 @@ Q4nxCacheSweep q4nx_cache_evict(const std::string& root, const std::string& keep } return r; } + +int64_t q4nx_cache_bound_bytes(const char* text, int64_t default_gib) { + const int64_t max_gib = INT64_MAX >> 30; // the largest GiB count whose byte count still fits + long long gib = default_gib; + if (text && *text) { + char* end = nullptr; + errno = 0; + const long long v = std::strtoll(text, &end, 10); + if (end != text) gib = v; // a number (ERANGE saturates, which the clamp below handles) + } + if (gib <= 0) return 0; + if (gib > max_gib) gib = max_gib; + return int64_t(gib) << 30; +} + +void q4nx_cache_touch(const std::string& entry) { + std::error_code ec; + if (fs::is_directory(entry, ec)) fs::last_write_time(entry, fs::file_time_type::clock::now(), ec); +} diff --git a/app/q4nx_cache.h b/app/q4nx_cache.h index 53783bac3..f5d987511 100644 --- a/app/q4nx_cache.h +++ b/app/q4nx_cache.h @@ -44,3 +44,16 @@ Q4nxCacheSweep q4nx_cache_evict(const std::string& root, const std::string& keep // Directory bytes under `path` (0 when it does not exist). int64_t q4nx_dir_bytes(const std::string& path); + +// The byte bound for ONEBIT_Q4NX_CACHE_MAX_GIB's value (`text` may be null). Unset, empty or not a number gives the default +// (`default_gib`); zero or negative gives 0, which q4nx_cache_evict() reads as "no sweep". The result never exceeds +// INT64_MAX: the plain `gib << 30` this replaces overflowed from 2^33 GiB up, so a very large number written to mean "no +// limit" could come out as a 1 GiB bound, and a typo ("sixty") silently turned the sweep off. +int64_t q4nx_cache_bound_bytes(const char* text, int64_t default_gib = 60); + +// Mark `entry` (a directory in the cache) as used now. The sweep orders by the entry directory's mtime, which the filesystem +// only moves when a file is created or removed inside it -- so without this the order is by CREATION, and a container +// imported last week and served every day is evicted before yesterday's one-off. Call it on every cache hit. It also puts a +// just-opened entry inside the sweep's age window, which is what makes "a serve may be reading it" true for a reused entry +// and not only for a new one. Failure is ignored (a read-only cache still serves). +void q4nx_cache_touch(const std::string& entry); diff --git a/app/serve.cpp b/app/serve.cpp index 2ea0d3da3..65cc89962 100644 --- a/app/serve.cpp +++ b/app/serve.cpp @@ -185,6 +185,8 @@ struct Options { int prefix_slots = -1; // --prefix-slots N (0 = off) int prefix_cache_mb = -1; // --prefix-cache-mb N int prefix_min_gain = -1; // --prefix-min-gain N: tokens a prompt must add to be worth keeping + bool prefix_pin = false; // --prefix-pin: keep the part single-turn requests share (docs/npu.md) + bool flash = false; // --flash: the single-turn NPU profile = --ctx-size 1024 + --prefix-pin (engine#1785) std::string npu_target; // --npu-target NAME: which NPU target serves the per-op route (empty = the build's default) std::string prefill_provider; // --prefill-provider PATH (or ONEBIT_HYBRID_PREFILL): --device hybrid's prefill (docs/hybrid.md) std::string laya_model; // --laya-model DIR: route each request by the Laya scorer (docs/laya.md) @@ -1149,7 +1151,16 @@ Launch launch_for(const Options& o, const std::string& device, int child_port) { // the experts stay file-backed: without these, llama.cpp copies every one into RAM // (a pinned host buffer, or a CPU repack), which is what streaming avoids // (per_layer_token_embd: Qwen3.8-Flash-Next's 26.8 GiB table, only ever row-gathered) - argv.insert(argv.end(), {"-ot", "exps=CPU,per_layer_token_embd=CPU", "--no-host", "--no-repack", "--load-mode", "mmap"}); + argv.insert(argv.end(), {"-ot", "exps=CPU,per_layer_token_embd=CPU", "--no-host", "--no-repack"}); + // engine#1668: `-lm none` is +17.5 % prefill at -1.6 % decode for ~604 MiB more GTT (docs/bench.md), + // and prefill is what a served model pays for TTFT -- but this site hard-coded `--load-mode mmap`, + // and an explicit flag beats the child's own LLAMA_ARG_LOAD_MODE, so the measured lever could not be + // selected in a server at all. Default unchanged (mmap when the operator has not asked for something + // else); setting LLAMA_ARG_LOAD_MODE=none in the environment now reaches the child, with the GTT cost + // that implies documented beside the measurement. + if (std::getenv("LLAMA_ARG_LOAD_MODE") == nullptr) { + argv.insert(argv.end(), {"--load-mode", "mmap"}); + } #endif } if (device == "hrx") { @@ -2958,12 +2969,12 @@ std::string q4nx_cache_root() { // engine#1956: the cache is regenerable scratch, not user data, and nothing ever removed an old entry -- measured at 104 G in // 29 entries on one box, which closed two lanes' build gates the same day. The bound is a size, not an age, because a busy // day imports more packs than any age window expires; ONEBIT_Q4NX_CACHE_MAX_GIB=0 restores the old unbounded behaviour. +// The order is least-recently-USED: a cache hit touches its entry (q4nx_cache_touch), here and in npu/q4nx_device.cpp. // The entry just written is never a candidate, and neither is one younger than the sweep's age window (a serve may be reading // it) -- see app/q4nx_cache.h for why that is a heuristic rather than a guarantee. static void q4nx_cache_sweep(const std::string& cache, const std::string& keep) { - const char* env = std::getenv("ONEBIT_Q4NX_CACHE_MAX_GIB"); - long long gib = env && *env ? std::atoll(env) : 60; - q4nx_cache_evict(cache, keep, gib > 0 ? gib * (1LL << 30) : 0); + q4nx_cache_touch(keep); // a reused entry counts as used now: the sweep is LRU by this, not by creation + q4nx_cache_evict(cache, keep, q4nx_cache_bound_bytes(std::getenv("ONEBIT_Q4NX_CACHE_MAX_GIB"))); } std::string repack_gguf(const std::string& gguf) { @@ -3041,9 +3052,13 @@ std::string repack_gguf(const std::string& gguf) { "1bit serve: ============================================================\n\n", checker.empty() ? "(unresolved)" : checker.c_str()); } + q4nx_cache_touch(out); // a hit is a use: without this the sweep's order is by creation + return out; + } + if (run_program({ python, checker, gguf, out }, true) == 0) { + q4nx_cache_touch(out); return out; } - if (run_program({ python, checker, gguf, out }, true) == 0) return out; std::fprintf(stderr, "1bit serve: cached Q4NX dir %s FAILED the declared-scale scan; " "discarding and re-repacking\n", out.c_str()); fs::remove_all(out); @@ -3167,6 +3182,8 @@ void usage(FILE* out) { " [--npu-streams N] with --device npu: this serve is one of N sharing the NPU\n" " [--prefix-slots N] [--prefix-cache-mb N] [--prefix-min-gain N] the prompt prefix cache\n" " (docs/npu.md; 0 slots turns it off). Unset keeps the forward's defaults: 8 / 4096 / 128\n" + " [--prefix-pin] keep the part a request shares with the previous one (one system prompt, many questions)\n" + " [--flash] with --device npu: the single-turn profile, --ctx-size 1024 plus --prefix-pin\n" " (default 1, the device exclusive; docs/npu.md)\n" " [--load-queue-ms N] a load refused because another client holds the NPU array\n" " waits at most N ms for it to free before the driver is called dead\n" @@ -3282,6 +3299,8 @@ int run_serve(int argc, char** argv) { else if (a == "--prefix-slots") o.prefix_slots = std::stoi(next()); else if (a == "--prefix-cache-mb") o.prefix_cache_mb = std::stoi(next()); else if (a == "--prefix-min-gain") o.prefix_min_gain = std::stoi(next()); + else if (a == "--prefix-pin") o.prefix_pin = true; + else if (a == "--flash") o.flash = true; else if (a == "--npu-target") o.npu_target = next(); else if (a == "--prefill-provider") o.prefill_provider = next(); else if (a == "--laya-model") o.laya_model = next(); @@ -3318,6 +3337,16 @@ int run_serve(int argc, char** argv) { else if (a == "-h" || a == "--help") { usage(stdout); return 0; } else throw std::runtime_error("unknown option " + a); } + // --flash is a name for two settings, so that the profile is one word and its parts stay visible: a 1024-token + // context and the pinned shared prefix. It refuses what contradicts it instead of quietly winning or losing. + if (o.flash) { + if (o.ctx_size != 0 && o.ctx_size != 1024) + throw std::runtime_error("--flash fixes the context at 1024 tokens; drop --ctx-size " + std::to_string(o.ctx_size)); + if (o.prefix_slots == 0) throw std::runtime_error("--flash needs the prefix cache; drop --prefix-slots 0"); + if (o.device != "npu") throw std::runtime_error("--flash is a profile of --device npu"); + o.ctx_size = 1024; + o.prefix_pin = true; + } if (!o.laya_models.empty()) { #if defined(ONEBIT_LAYA) && !defined(_WIN32) if (o.laya_model.empty() && !o.laya_auto) throw std::runtime_error("--laya-models needs --laya or --laya-model DIR"); @@ -3417,6 +3446,21 @@ int run_serve(int argc, char** argv) { fs::exists(fs::path(o.model) / "config.json") && fs::exists(fs::path(o.model) / "tokenizer.json"); if (model_generic && !kernels.empty() && fs::is_directory(kernels)) { + // engine#1740 (W6): --npu-opt is for a PRIVATE route, and this branch is the per-op forward, which + // cannot consume one. #1868 made the option refuse when no add-on is linked; with an add-on linked the + // option was still accepted here and then bypassed by the CONTAINER TYPE, silently -- a W6 run asked for + // the fused route, got per-op, and produced *identical* numbers, which reads as "no gain" rather than + // "wrong route". Refusing is the only answer that cannot be misread. + if (!o.npu_opts.empty()) { + std::string keys; + for (const auto& kv : o.npu_opts) keys += (keys.empty() ? "" : ", ") + kv.substr(0, kv.find('=')); + std::fprintf(stderr, + "1bit serve: --npu-opt %s was given, but %s is a model-generic container, which serves " + "through the per-op forward: the option has no effect here (it applies to a private " + "route's own model directory). Drop it, or serve the model directory that route " + "expects.\n", keys.c_str(), o.model.c_str()); + return 2; + } std::fprintf(stderr, "1bit serve: %s is a model-generic container: serving it through the per-op forward " "(kernels %s)\n", o.model.c_str(), kernels.c_str()); @@ -3433,6 +3477,10 @@ int run_serve(int argc, char** argv) { if (o.prefix_slots >= 0) { fargs.push_back("--prefix-slots"); fargs.push_back(std::to_string(o.prefix_slots)); } if (o.prefix_cache_mb >= 0) { fargs.push_back("--prefix-cache-mb"); fargs.push_back(std::to_string(o.prefix_cache_mb)); } if (o.prefix_min_gain >= 0) { fargs.push_back("--prefix-min-gain"); fargs.push_back(std::to_string(o.prefix_min_gain)); } + if (o.prefix_pin) fargs.push_back("--prefix-pin"); + // The context limit reaches this route too: it used to stop here, so `serve --ctx-size N` on a + // model-generic container was accepted and the forward kept the model's own limit. + if (o.ctx_size > 0) { fargs.push_back("--ctx-size"); fargs.push_back(std::to_string(o.ctx_size)); } std::vector fav; for (auto& a : fargs) fav.push_back(a.data()); return run_forward_serve(int(fav.size()), fav.data()); diff --git a/app/stream_bridge.h b/app/stream_bridge.h new file mode 100644 index 000000000..cdca732c8 --- /dev/null +++ b/app/stream_bridge.h @@ -0,0 +1,118 @@ +// 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. +// +// stream_bridge.h -- hands a streamed reply from the thread that generates it to the thread that writes it. +// +// The NPU route used to build every server-sent-event frame of a streamed reply and send them all after the +// generation had ended, so a client's time to first token was the whole request (engine#1785: 14 chunks, every +// one arriving at 9.668 s of a 9.668 s reply). httplib writes a response body only after the handler returns, +// from a content provider it calls on the connection's thread, so frames can leave while tokens are still being +// produced only if the generation runs on another thread. This is the piece between the two. +// +// producer (the generation) consumer (the connection) +// push(frame) ... push(frame) wait_first() -> a frame is ready, or it ended without one +// finish() next(frame) -> the next frame in order; false when ended and drained +// cancelled() <- the client went away cancel() +// +// The decision that needs wait_first(): a request that fails before its first frame (a prompt over the context, +// a malformed field) must still be answered with its own status code, and a status cannot change once a stream +// has started. So the connection thread waits for the first frame OR the end, and only then chooses between +// "an ordinary error response" and "a stream". +// +// Frames are delivered in the order pushed, each exactly once. finish() may be called once; push() after it is +// ignored. Nothing here knows about HTTP. + +#pragma once + +#include +#include +#include +#include + +namespace onebit::npu { + +class StreamBridge { +public: + // Producer: queue one frame. Returns false when the consumer has cancelled, so the producer can stop early. + bool push(std::string frame) { + std::lock_guard lock(mu_); + if (finished_) return !cancelled_; + frames_.push_back(std::move(frame)); + ++pushed_; + cv_.notify_all(); + return !cancelled_; + } + + // Producer: no more frames. + void finish() { + std::lock_guard lock(mu_); + finished_ = true; + cv_.notify_all(); + } + + // Producer: has the consumer gone? + bool cancelled() const { + std::lock_guard lock(mu_); + return cancelled_; + } + + // Consumer: block until a frame has been pushed or the producer has finished. True when at least one frame + // was ever pushed (a stream should start), false when it ended with none (answer some other way). + bool wait_first() { + std::unique_lock lock(mu_); + cv_.wait(lock, [&] { return pushed_ > 0 || finished_; }); + return pushed_ > 0; + } + + // Consumer: the next frame, blocking while there is none and the producer is still going. False when the + // producer has finished and every frame has been taken, or after cancel(). + bool next(std::string& frame) { + std::unique_lock lock(mu_); + cv_.wait(lock, [&] { return !frames_.empty() || finished_ || cancelled_; }); + if (cancelled_ || frames_.empty()) return false; + frame = std::move(frames_.front()); + frames_.pop_front(); + return true; + } + + // Consumer: stop. The producer sees it at its next push() or cancelled(). + void cancel() { + std::lock_guard lock(mu_); + cancelled_ = true; + cv_.notify_all(); + } + + // Has any frame ever been pushed? (The producer asks this to decide whether a failure can still be an ordinary + // error response or has to be reported inside the stream.) + bool started() const { + std::lock_guard lock(mu_); + return pushed_ > 0; + } + + bool finished() const { + std::lock_guard lock(mu_); + return finished_; + } + +private: + mutable std::mutex mu_; + std::condition_variable cv_; + std::deque frames_; + size_t pushed_ = 0; + bool finished_ = false; + bool cancelled_ = false; +}; + +} // namespace onebit::npu diff --git a/app/think_split.h b/app/think_split.h index 5ba670400..efb06201f 100644 --- a/app/think_split.h +++ b/app/think_split.h @@ -16,6 +16,7 @@ #include #include +#include namespace onebit { @@ -73,4 +74,35 @@ struct ThinkSplit { } }; +// How many of `pieces` (one generated token's text each, in order) were reasoning, for usage.completion_tokens_details +// .reasoning_tokens (engine#2111): the twin of npu::reasoning_token_count() over this route's splitter. A token is +// reasoning when none of its text reached content and the block was open before it, after it, or it put text into +// it -- so both tags, a held-back partial tag and an unclosed block at the end of the budget count. +inline int think_reasoning_tokens(const std::vector& pieces, bool opened) { + ThinkSplit s(opened); + int n = 0; + bool last_counted = false; + int pending = 0; // leading tokens that produced nothing yet: they join the block if the next one opens it + bool started = false; + for (const auto& piece : pieces) { + const bool before = s.in_think; + bool content = false, reasoning = false; + s.feed(piece, [&](bool r, const std::string&) { (r ? reasoning : content) = true; }); + last_counted = !content && (reasoning || before || s.in_think); + if (last_counted) { + ++n; + if (!started) n += pending; + started = true; + } else if (content) { + started = true; + } else if (!started) { + ++pending; + } + } + bool flushed_content = false; + s.flush([&](bool r, const std::string&) { if (!r) flushed_content = true; }); + if (flushed_content && last_counted) --n; + return n; +} + } // namespace onebit diff --git a/app/unified.cpp b/app/unified.cpp index 102ab340d..c44f2a436 100644 --- a/app/unified.cpp +++ b/app/unified.cpp @@ -237,12 +237,27 @@ void handle_generate(Engine& e, const httplib::Request& req, httplib::Response& return "data: " + c.dump() + "\n\n"; }; - auto usage = [](const npu::GenerateResult& r, size_t n_prompt, size_t n_cached) { + // engine#2111: a chat reply also says how many completion tokens were reasoning (OpenAI's + // completion_tokens_details.reasoning_tokens), so content "" with finish_reason "length" can be told apart from a + // model that said nothing: the budget ran out inside the thinking block. Counted over the generated ids (the end + // token excluded, as completion_tokens excludes it) with this route's own splitter; 0 when nothing is split. + auto usage = [&e, chat, split_reasoning, think_opened](const npu::GenerateResult& r, size_t n_prompt, size_t n_cached) { const size_t completion = r.tokens.size() - (r.stopped_at_eos ? 1 : 0); - return json{{"prompt_tokens", n_prompt}, - {"completion_tokens", completion}, - {"total_tokens", n_prompt + completion}, - {"prompt_tokens_details", {{"cached_tokens", n_cached}}}}; + json u{{"prompt_tokens", n_prompt}, + {"completion_tokens", completion}, + {"total_tokens", n_prompt + completion}, + {"prompt_tokens_details", {{"cached_tokens", n_cached}}}}; + if (chat) { + int reasoning = 0; + if (split_reasoning) { + std::vector pieces; + pieces.reserve(completion); + for (size_t i = 0; i < completion; ++i) pieces.push_back(e.tok->decode(r.tokens[i])); + reasoning = think_reasoning_tokens(pieces, think_opened); + } + u["completion_tokens_details"] = {{"reasoning_tokens", reasoning}}; + } + return u; }; auto timings = [](const npu::GenerateResult& r, const std::string& cache) { json t{{"prompt_ms", r.prefill_ms}, diff --git a/cmake/tests-python-tools.cmake b/cmake/tests-python-tools.cmake index aa973cc8d..909c40d8f 100644 --- a/cmake/tests-python-tools.cmake +++ b/cmake/tests-python-tools.cmake @@ -25,6 +25,16 @@ add_python_test(python_test_order_check ${CMAKE_SOURCE_DIR}/tests/python_test_or add_python_test(python_test_order_check_self ${CMAKE_SOURCE_DIR}/tests/python_test_order_check.py --self-test) add_python_test(ci_targets ${CMAKE_SOURCE_DIR}/tools/ci/ci_targets_check.py --check --root ${CMAKE_SOURCE_DIR}) add_python_test(ci_targets_self ${CMAKE_SOURCE_DIR}/tools/ci/ci_targets_check.py --self-test) +# engine#803 gate, second half: the comparator that turns two --dump-logits directories into the KLD verdict the +# P4 gate asks for. No device and no model, so the gate is one command when a ticket is free; the self-test pins +# the four cases a KLD implementation gets wrong (identical -> exactly 0, a shift is positive and directional, a +# uniform shift stays 0, and an asymmetric file set is refused by name). +add_python_test(logits_kld_self ${CMAKE_SOURCE_DIR}/tools/logits_kld.py --self-test) +# engine#2018: this one was in tools/ci/unreferenced_py_tests.txt as a baseline entry, and it should not have been: the +# file is portable and passes (worst-case |sum| 11.18e6 inside int32 at K=11008, 4 arms), it just never had a runner. A +# baseline records a test as never-run, which for a passing test is the wrong outcome -- so it is registered and the line +# is removed, which is the shrink the baseline's own comment asks for. +add_python_test(moe_expert_chunks_accum_test ${CMAKE_SOURCE_DIR}/tests/moe_expert_chunks_accum_test.py) # engine#1820: the other half of ci_targets_check -- a standalone shell test that nothing references cannot fail. # `need` is every tests/*.sh, `have` is every file that mentions its basename somewhere that can run it, and # tools/ci/unreferenced_tests.txt carries the deliberate exclusions WITH their reasons, so the list can shrink. @@ -37,8 +47,26 @@ add_python_test(tests_referenced_self ${CMAKE_SOURCE_DIR}/tools/ci/tests_referen # Registered through add_python_test (engine#1864), so they run THROUGH the interpreter: my first version used the raw # `${Python3_EXECUTABLE}` form above the find_package and the repository's own ordering guard caught it by name -- # tools/mtp_head_ref.py is the one of the five without an exec bit, and it became BAD_COMMAND under ctest. +# engine#1962: the narrowing-body enumeration, as a ratchet. The corpus is a submodule, so the check +# skips (rc 77) where third_party/llama.cpp is not checked out; the baseline file is in the engine tree, +# so the recorded set is reviewable everywhere. +add_python_test(tool_selftest_hrx_f16_narrowing ${CMAKE_SOURCE_DIR}/tools/ci/hrx_f16_narrowing.py --self-test) +add_test(NAME hrx_f16_narrowing_check COMMAND ${Python3_EXECUTABLE} + ${CMAKE_SOURCE_DIR}/tools/ci/hrx_f16_narrowing.py --check --root ${CMAKE_SOURCE_DIR}) +set_tests_properties(hrx_f16_narrowing_check PROPERTIES SKIP_RETURN_CODE 77) add_python_test(tool_selftest_batched_repeat ${CMAKE_SOURCE_DIR}/tools/batched_repeat.py --selftest) add_python_test(tool_selftest_logprob_agree ${CMAKE_SOURCE_DIR}/tools/logprob_agree.py --selftest) add_python_test(tool_selftest_mtp_head_ref ${CMAKE_SOURCE_DIR}/tools/mtp_head_ref.py --selftest) add_python_test(tool_selftest_spec_decode_acceptance ${CMAKE_SOURCE_DIR}/tools/spec_decode_acceptance.py --selftest) add_python_test(tool_selftest_whisper_encoder_baseline ${CMAKE_SOURCE_DIR}/tools/whisper_encoder_baseline.py --self-test) +# engine#1780: the DFlash draft package's gate before a loader -- tile-grid sizes from the config, tile-header sanity and +# the container's scalars against config.json, on a synthetic package with eight negative controls. Stdlib only. +add_python_test(tool_selftest_dflash_draft_info ${CMAKE_SOURCE_DIR}/tools/dflash_draft_info.py --selftest) +# engine#1750: what a TDR recovery did, counted from the kernel log (failed context restarts among them). The self-test +# parses lines in the driver's own formats for three module generations. Stdlib only, no device. +add_python_test(tool_selftest_amdxdna_tdr_report ${CMAKE_SOURCE_DIR}/tools/amdxdna_tdr_report.py --self-test) +# engine#2013: the qwen35 (gated linear attention) GGUF -> Q4NX converter. Its own selftest (tile round-trips, the +# value-head permutation), and a pack + verify of a small qwen35 GGUF the test writes itself, with the three inputs it +# has to refuse. numpy only: no gguf package, no model, no device. +add_python_test(tool_selftest_qwen35_to_q4nx ${CMAKE_SOURCE_DIR}/tools/qwen35_to_q4nx.py selftest) +add_python_test(qwen35_to_q4nx ${CMAKE_SOURCE_DIR}/tests/qwen35_to_q4nx_test.py) diff --git a/cmake/tests-registration-glob.cmake b/cmake/tests-registration-glob.cmake new file mode 100644 index 000000000..2e4e1b43b --- /dev/null +++ b/cmake/tests-registration-glob.cmake @@ -0,0 +1,23 @@ +# 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 registration glob's own test (engine#1979), registered the way it asks new tests to be: this file is named by +# nothing and is picked up by the cmake/tests-*.cmake glob near the end of CMakeLists.txt. +# +# It configures the glob block of the real CMakeLists.txt in scratch projects: a new file is registered, a file named +# only in a comment is still registered, and a file CMakeLists.txt include()s by name is registered once. Configure +# only (project(... NONE)): no compiler, no device, about a second. +add_test(NAME registration_glob + COMMAND bash ${CMAKE_SOURCE_DIR}/tests/registration_glob_test.sh ${CMAKE_COMMAND} ${CMAKE_SOURCE_DIR}/CMakeLists.txt) diff --git a/cmake/tests-submodule-pin.cmake b/cmake/tests-submodule-pin.cmake new file mode 100644 index 000000000..3a00b8433 --- /dev/null +++ b/cmake/tests-submodule-pin.cmake @@ -0,0 +1,27 @@ +# 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. +# +# engine#2129: every submodule must be checked out at the commit the tree pins. +# +# A build compiles whatever is in the working tree, and every other check in tools/ci reads the tree it is handed - +# so none of them can see that a submodule sits on a different commit than the superproject pins. Measured on +# 2026-11-11: the shared ~/1bit-engine had four submodules off their pins (llama.cpp, hrx-system, tokenizers, +# xdna-driver), and the engine compiles common/jinja from llama.cpp - which was behind the pin on `\xNN` escapes. +# Jinja decides prompt rendering, so that is a difference in what the model sees, not a build failure. +# +# The default run compares only the submodules that ARE checked out (a configure that never inits them is not this +# check's business); `--require-initialised` is the release-gate form, where a missing checkout is a failure too. +add_python_test(submodule_pin_check ${CMAKE_SOURCE_DIR}/tools/ci/submodule_pin_check.py --check --root ${CMAKE_SOURCE_DIR}) +add_python_test(submodule_pin_check_self ${CMAKE_SOURCE_DIR}/tools/ci/submodule_pin_check.py --self-test) diff --git a/cmake/tests-tmp-path.cmake b/cmake/tests-tmp-path.cmake new file mode 100644 index 000000000..588193410 --- /dev/null +++ b/cmake/tests-tmp-path.cmake @@ -0,0 +1,24 @@ +# 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. +# engine#2021: a hardcoded `/tmp` path is RAM on this box, not disk. +# +# `/tmp` is tmpfs here, so a dump, a staging copy or a benchmark log written there is anonymous memory: swap fills +# and every lane on the box stalls (measured 2026-11-10 21:15Z: /tmp 52 G of 62 G, swap 71/71 GiB). engine#1907/#1952 +# bounded the fconst scratch; the literals themselves are still in the tree, 14 of them across 8 files. +# +# Both rows are cheap and device-free: the check is a grep over npu/ app/ tools/ scripts/ against +# tools/ci/tmp_paths.txt, and the ratchet may only shrink. +add_python_test(tmp_path_check ${CMAKE_SOURCE_DIR}/tools/ci/tmp_path_check.py --check --root ${CMAKE_SOURCE_DIR}) +add_python_test(tmp_path_check_self ${CMAKE_SOURCE_DIR}/tools/ci/tmp_path_check.py --self-test) diff --git a/config/llama-mirror-hold.config b/config/llama-mirror-hold.config new file mode 100644 index 000000000..b77bd4d8f --- /dev/null +++ b/config/llama-mirror-hold.config @@ -0,0 +1,34 @@ +# 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 commit of 1bit-MONSTER/llama.cpp `1bit/hrx-vulkan` (our mirror of AMD's HRX line) that the daily +# bump-hrx run must not move away from. While a commit is named here, bump-hrx does not move the mirror, +# does not merge AMD's llama.cpp pin into `1bit/hrx-vulkan-patched`, reports the hold in the run summary and +# succeeds. Delete the commit line to lift the hold. +# +# The hold is for one situation only: AMD rewrote its line, so its pin no longer descends from the mirror, +# and taking it is a hand merge that has not been done yet. The run FAILS, naming this file, when the file +# is out of date (the mirror is no longer the named commit) or no longer needed (AMD's pin contains the +# named commit again), so a hold cannot outlive the state it describes. +# +# HELD 2026-11-11 for the release (orchestrator decision on engine#1945, the owner having delegated technical +# calls): the engine's llama.cpp pin stays where it is, and the merge of AMD's rebased line is post-release work. +# The state this records: +# mirror 1bit/hrx-vulkan 44ee5710d19f7f3af79b99bbcbe5fd13979d1f67 (AMD's line before its rebase) +# AMD pin 9bf5871310a8d7de74f4cb0c57009835fbf91c27 (the rebased line, 1777 upstream commits newer) +# merge into 1bit/hrx-vulkan-patched, against the last AMD commit our line took (43b0ec2c41cc): +# 149 conflicted files, 538 hunks (65 files / 83 hunks in the HRX backend) +# Lift it by deleting the line below, once AMD's line has been merged into ours. +44ee5710d19f7f3af79b99bbcbe5fd13979d1f67 diff --git a/docs/amd-drafters.md b/docs/amd-drafters.md index ee89be046..a7a7235d1 100644 --- a/docs/amd-drafters.md +++ b/docs/amd-drafters.md @@ -88,6 +88,56 @@ So the Gumiho row is *"fails as published, because our verifier is the wrong sha quality (its published numbers are the best in its class: 2.97× mean over Eagle-2's 2.74× on vicuna-7B). Running it would mean implementing a tree verifier first, which is its own project. +## DFlash draft packages on the NPU route: **the container is in forms we already read; the forward is what is missing** (engine#1780) + +FastFlowLM 1.0.12 ships speculative decode as two-package tags: a small DFlash draft beside the main model's package. +A draft package is `config.json` (`model_type: dflash`) plus one Q4NX container. `tools/dflash_draft_info.py ` +checks such a package against its own config and prints what it declares; on +`Nemotron-3.5-Bolt-E4B-DFlash-NPU2` (0.27 GB, sha256 matching the repo's `MANIFEST.sha256`) it passes 19 of 19 +checks in 3.4 s. What that settles, each with the evidence the tool prints: + +| question | answer | evidence | +|---|---|---| +| are the `I8 [tiles, 5120]` projections a new packing? | **No.** They are the ordinary Q4NX q4_1 tile grid (32 x 256 weights per 5120-byte tile), the form every main-model container uses. | all 35 projections have `bytes == (N/32)(K/256) x 5120` for the N x K the config gives the role; 10,240,000 groups with every scale finite and positive; minimum/scale median -7.59 .. -7.33 per tensor (a known-good main model reads -7.41 the same way) | +| is `q_proj` 4096 x 2560 or a half-float 1280 x 2560? | **4096 x 2560**, so `head_dim` is 128 as the config says. | both readings are 6,553,600 bytes, but the half-float one has 25.7 % of its values above 100 in magnitude and 1.8 % non-finite, and 1280 x 2560 as a tile grid would be 400 tiles, not 1280 | +| are the `BF16` labels wrong (really F16)? | **No, they are bf16.** | a norm gain read as half-float looks tight (1.6-2.4) and as bf16 looks wide (0.22-6.3), which reads like a mislabel. The same pattern appears in a main-model container our forward serves correctly with the bf16 reading (1.49-1.88 against 0.12-1.05): half-float folds bf16's exponent range into one octave, so "tight" proves nothing | +| does the container agree with `config.json`? | **Yes.** | `dflash.block_size`, `target_layer_ids`, `sliding_window`, `rope_theta`, `rms_eps`, `embedding_scale` all equal | +| is anything in it unexplained? | **No.** | 64 tensors: 35 projections, 22 norms, `model.fc.weight` (2560 x 12800, the fusion of five target layers' features) and six scalars | + +So importing a DFlash draft is a config-schema translation, not a format project. What our runtime does not have is the +draft's **forward**: fuse the target's hidden states at `target_layer_ids` through `fc`, run five non-causal layers +over a block of `block_size` positions filled with `mask_token_id`, and score the block with the target's own +embedding table and head. `forward-serve` stops on such a package at the first key it expects +(`key 'model' not found`), which is the schema, before any of that. + +The tool is the gate before a loader, and it fails closed: a projection that is a tile short, a `head_dim` that +disagrees with the stored tiles, a NaN scale, a plain half-float matrix of the same size, a scalar that differs from +the config, or a tensor it cannot explain (the second draft family, `DFlash2DraftModel`, has selector and +convolution tensors this tool does not know) each fail by name. `--selftest` runs those controls on a synthetic +package (ctest `tool_selftest_dflash_draft_info`). + +### The loader (`npu/dflash_draft.h`) + +`onebit::npu::dflash::load_draft(dir)` is the C++ half of that gate: it reads the DFlash `config.json`, opens the +container with the engine's own Q4NX reader, checks every tensor against the geometry the config gives, and returns +a model whose tensors are views into the mapped file. Nothing is decoded at load (the Bolt draft's projections are +327,680,000 weights, 1.3 GB as f32); `decode_projection()`, `decode_gain()` and `decode_fusion()` produce one tensor +at a time, through the same q4_1 tile decoder the forward uses. It throws, naming the tensor or key, on a missing or +mis-sized tensor, a dtype label it does not expect, a container scalar that disagrees with the config, a tensor it +cannot explain, another architecture (`DFlash2DraftModel`), or a weight path that leaves the package. + +Two facts about the container that the loader depends on, both measured on the Bolt draft: + +| fact | evidence | +|---|---| +| `model.fc.weight` is **not** row-major. Its declared shape `[80, 12800, 32]` is the element order: 32-row blocks, column by column, the row inside the block last. | Magnitudes that share an input column correlate about five times more strongly than magnitudes that share an output row. Decoded in the declared order the four neighbour figures are consistent (next row 0.107, 32 rows on 0.106; next column 0.022, 32 columns on 0.021). Read as plain row-major they contradict each other (next row 0.019 but 32 rows on 0.107; next column 0.105 but 32 columns on 0.024). `dflash_draft_test` asserts the first and uses the second as its control. | +| A norm gain may be negative. | Three post-attention norms each hold one gain just below zero (-0.0009 to -0.0024) among 2,560. As bf16 that is a gain crossing zero; as half-float the same element would be -0.8 among values of 1.3 to 2.4, which is one more reason the bf16 reading is the right one. | + +`dflash_draft_test` builds a synthetic package (every projection encoded tile by tile from a formula and compared at +every element, the fusion matrix checked for row-major order, sixteen refusals each matched on its message) and needs +no model; given a package directory (`-DONEBIT_DFLASH_DRAFT=` or `ONEBIT_TEST_DFLASH_DRAFT`) it also loads and +decodes the real one. The loader is not wired into `forward-serve` yet: there is no forward for it to feed. + ## PARD-2's repos are weights-only — borrow the target's tokenizer `amd/PARD2-*` ship **`README.md`, `config.json`, `model.safetensors`, `warp_model.bin`** and nothing else: **no tokenizer diff --git a/docs/arch-gaps.md b/docs/arch-gaps.md index 8f7d5fdb5..aa1d2cb0a 100644 --- a/docs/arch-gaps.md +++ b/docs/arch-gaps.md @@ -1075,3 +1075,15 @@ which are retried on every run; and the `!! FLM` lines. Those keep the 2026-11-0 and `smolvla` as above, and `!! NPU FLM` is engine #588. This run's issue window also includes the `gpt-oss` (`GPT-OSS-20B`, `GPT-OSS-Safeguard-20b`) and `lfm2` (`LFM2-1.2B`, `LFM2-2.6B`) Q4NX packs. Their mapping is HRX-only too, so they belong to #588, not to a registry edit. + +## 2026-11-10: `DeepseekV4FlashSuperExpertForCausalLM` — reviewed, and recorded here + +Found by the check this page's contract implies and `tests/census_watch_test.py` now enforces: every class in +`registry/significant.json` has to be cited here, and this one was added there (engine#1922) without a row, so it was +the single undocumented entry out of 37. + +* `DeepseekV4FlashSuperExpertForCausalLM` (`model_type` `deepseek_v4_flash_superexpert`) — DeepSeek V4.1 Flash + SuperExpert, **published as remote code and with AW-SECD-compressed weights** (8.5 GB, claimed 0.962 retention); + MLA-style geometry (`kv_lora_rank` 512, `q_lora_rank` 1536, `qk_rope_head_dim` 64, `v_head_dim` 128, 32 heads / + 16 KV, IM 4096). **Not an alias of `deepseek_v41`**: nothing at the pin decodes AW-SECD weights, so mapping it to + that architecture would name a backend that cannot read the file. Reported, not alerted (`"watch": "report"`). diff --git a/docs/axb35-usb-ports.md b/docs/axb35-usb-ports.md index d369c08ca..06a985d40 100644 --- a/docs/axb35-usb-ports.md +++ b/docs/axb35-usb-ports.md @@ -70,6 +70,34 @@ addressing something a newer kernel had already fixed. Forcing `power/control=on them `D3cold → D0` and their root hubs' ports appear in `/sys/bus/usb/devices`; `--revert` puts them back. So the mechanism works on this kernel; what has not happened is the plug. +## The two halves of this issue have different fixes, and the rule covers one + +Measured on the board 2026-11-10, BIOS 1.10, a dev boot with the rule not installed: + +``` +0000:c5:00.4 0x1587 xhci_hcd D3cold <- rule matches +0000:c7:00.0 0x1588 xhci_hcd D0 <- rule matches +0000:c7:00.3 0x1589 xhci_hcd D3cold <- rule matches +0000:c7:00.4 0x158b xhci_hcd D3cold <- rule matches +0000:c7:00.5 0x158d thunderbolt D3cold <- the rule does NOT match this +0000:c7:00.6 0x158e thunderbolt D3cold <- the rule does NOT match this +``` + +The **rear USB-A** ports hang off the four xHCI functions the rule matches. The **Type-C** ports hang off the two USB4 host routers +(`0x158d`/`0x158e`), which the rule's device list does not contain, and which are in `D3cold` with the `thunderbolt` driver bound and +four live IRQs. So: + +| half of engine#823 | what has to be awake | does the shipped rule cover it | +|---|---|---| +| rear USB-A | one of `0x1587/0x1588/0x1589/0x158b` | **yes** | +| Type-C | `0x158d`/`0x158e` (USB4 host routers) | **no** | + +**A suspended USB4 host router with no partner attached is normal**, so this is not evidence that the routers are broken - it is evidence +that **one plug test cannot test both halves**: a plug into a rear USB-A exercises the rule, a plug into Type-C exercises a path no change +in this repository has touched. Record the router ids with the capture (`tools/xhci_boot_capture.sh` already records `power_state` per +controller; the two `thunderbolt` functions are worth reading the same way) so a "Type-C still dead" result is read as "the router half is +untested" rather than "the keepawake fix failed". + ## What is still owed One device into a rear USB-A or a Type-C port **on a 1bit-OS boot**, checked with `--require-awake` first. If a port is then still diff --git a/docs/bench.md b/docs/bench.md index 9233a739a..9dff89c8a 100644 --- a/docs/bench.md +++ b/docs/bench.md @@ -69,6 +69,13 @@ bytes rather than slack: three more ~306 MiB tensors become device-resident (3 quote a prefill delta measured across it. `-lm none` is not a lever to ship on the strength of one measurement — it costs GTT on a box where GTT is the scarce resource, and the honest statement is the mechanism above, not the percentage. +**How it is selected in a server (engine#1668).** The measurement above was taken with `llama-bench`/`llama-server` directly, and +`1bit serve` hard-coded `--load-mode mmap` for its child, so the lever could not be reached where prefill actually costs +time-to-first-token. It now passes that flag **only when the operator has not chosen a mode**: set `LLAMA_ARG_LOAD_MODE=none` +(the child's own variable, so there is one mechanism rather than two) and the flag is omitted, letting the environment value +ride. The default is unchanged, and the 604 MiB of GTT the `-lm none` row records is the price of asking — a per-model cost on +a box where the pool budget is the scarce resource (engine#1711), which is why this is an override and not a new default. + ## What is measured - **Prompt:** one long prompt (default: the first 8,000 characters of `docs/*.md`, about 1,800 diff --git a/docs/ci-runners.md b/docs/ci-runners.md index f3195c82a..107760817 100644 --- a/docs/ci-runners.md +++ b/docs/ci-runners.md @@ -306,7 +306,9 @@ here: a new workflow changes what the repository's token may do, which is an RFC ## Two rows, one question: are the runners healthy? Since #1682 and #1546 both landed there are two checks, and neither alone is enough — each is blind to what the -other sees. Both run in `scripts/weekly.sh`'s preflight, so a release refuses to start on a runner it cannot trust. +other sees. Both run in `scripts/weekly.sh`'s preflight. A runner that is **not accepting jobs** stops the release +before it starts. A runner that is up but whose **restart policy is short** is a warning, and the release goes on +(owner decision on engine#1504, 2026-11-11): see [the next section](#a-short-restart-policy-warns-it-does-not-block). | | `scripts/check-ci-runners.sh` | `scripts/ci-runner-health.sh` | |---|---|---| @@ -315,7 +317,7 @@ other sees. Both run in `scripts/weekly.sh`'s preflight, so a release refuses to | notices | a registration that was revoked, expired or never made | `Result=oom-kill`, a restart window, a unit that stopped | | cannot see | *why* it died — no `oom-kill`, no restart count | a unit that is running fine while its registration is gone | | speed | after the heartbeat times out | the moment the unit dies | -| exit | 0 online · 1 an offender, named · **2 cannot tell** (unreadable or truncated list) | 0 all accepting · 1 one is not · 2 no such units on this host | +| exit | 0 online · 1 an offender, named · **2 cannot tell** (unreadable or truncated list) | 0 all accepting, policy right · 1 one is not accepting jobs · 2 no such units on this host · **3 all accepting, a restart policy is short** | Run either by hand — both are read-only and neither restarts anything: @@ -335,3 +337,72 @@ rather than report green. One case neither row covers: a runner that is `online` in the API *and* active locally, but wedged so it takes no jobs. That is a hung process, not a registration or unit state; the practical signal for it is a queue that stops moving while both rows stay green. + +## A short restart policy warns, it does not block + +`scripts/ci-runner-health.sh` prints two lines per unit: its state (`PASS`/`FAIL`) and its restart policy +(`PASS`/`WARN`). The exit code keeps them apart, because they call for different things: + +| exit | meaning | `scripts/weekly.sh` preflight | +|---|---|---| +| 0 | every unit is accepting jobs and would come back after an OOM kill | passes | +| 1 | a unit is not accepting jobs now (dead, or in its restart window) | **blocks**: the checks after the release would queue for ever | +| 2 | no `actions.runner.*` units on this host | a note | +| 3 | every unit is accepting jobs, but a `WARN` line names a policy setting | **warns** and goes on | + +A unit that is down outranks a policy line, so 1 wins over 3. The warning in the release log reads +`preflight: WARNING, not blocking: every runner unit is accepting jobs, but a restart policy would leave one down +after the next OOM kill`, with the `WARN` lines above it naming the unit and the setting. + +Why a warning is enough: the policy is a box configuration change, which the owner applies (below), and a release +that cannot start does not make the unit any safer. Meanwhile the heavy builds run under +`scripts/mem-guard.sh --term`, which ends the build before the kernel's OOM killer reaches the runner. **A build step +that ends with exit 143 was stopped by that guard**: the log has a `GUARD: MemAvailable N GB < floor GB` line, the +runner is still alive, and the job is re-run once memory frees. The guard covers only the steps that use it. A server +or a bench started without it can still take the unit down, which is why the drop-ins below are the real fix. + +`tests/weekly_runner_preflight_test.sh` (ctest `weekly_runner_preflight`) runs the real health script and the real +`preflight_runner_units()` from `scripts/weekly.sh` against a stub `systemctl`: the four exit codes, block on a dead +unit, warn on a short policy. + +### The two drop-ins (owner step) + +On strixhalo, `strixhalo-bench` and `strixhalo-org` still run the units that predate `scripts/install-ci-runner.sh`: +`Restart=no`, `OOMPolicy=stop`. One OOM kill of any process in the unit stops it for good. `strixhalo-engine` got +the same drop-in on 2026-11-10. This needs `sudo`, so it is the owner's to run. It does not re-register a runner, +change its user, or stop a job in flight. + +Both files have the same content, [`scripts/ci-runner-restart.conf`](../scripts/ci-runner-restart.conf): + +```ini +# /etc/systemd/system/actions.runner.1bit-MONSTER-engine.strixhalo-bench.service.d/restart.conf +# /etc/systemd/system/actions.runner.1bit-MONSTER.strixhalo-org.service.d/restart.conf +[Unit] +StartLimitIntervalSec=0 + +[Service] +Restart=always +RestartSec=15 +OOMPolicy=continue +``` + +The three commands, from the root of an engine checkout: + +```sh +sudo install -Dm644 scripts/ci-runner-restart.conf /etc/systemd/system/actions.runner.1bit-MONSTER-engine.strixhalo-bench.service.d/restart.conf +sudo install -Dm644 scripts/ci-runner-restart.conf /etc/systemd/system/actions.runner.1bit-MONSTER.strixhalo-org.service.d/restart.conf +sudo systemctl daemon-reload +``` + +Check it took (read-only, no `sudo`): + +```sh +systemctl show actions.runner.1bit-MONSTER-engine.strixhalo-bench.service actions.runner.1bit-MONSTER.strixhalo-org.service -p Id -p Restart -p OOMPolicy +scripts/ci-runner-health.sh; echo "exit $?" +``` + +Want `Restart=always` and `OOMPolicy=continue` on both. The health script still exits 3 afterwards, with one `WARN` +per unit naming only `MemoryMax=infinity(want a ceiling)`: the drop-in sets no memory ceiling. The ceiling this page +documents (`MemoryHigh=43G`, `MemoryMax=48G`) comes with a full `scripts/install-ci-runner.sh` install, or with +`sudo systemctl set-property MemoryHigh=43G MemoryMax=48G` per unit. To undo a drop-in, delete the +`restart.conf` and run `sudo systemctl daemon-reload`. diff --git a/docs/hrx.md b/docs/hrx.md index 8a8b687ba..d9feccefd 100644 --- a/docs/hrx.md +++ b/docs/hrx.md @@ -86,12 +86,31 @@ fails a pin that leaves a gate model's node off HRX. `DRY_RUN=1` and pushes nothing): - syncs the fork (`master` to ggml-org, `1bit/hrx-vulkan` to AMD's pin). When AMD rewrites its history, the old `1bit/hrx-vulkan` tip is tagged `hrx-vulkan-` first, - and the run fails if that branch held a commit that is on no AMD branch; + and the run fails if that branch held a commit that is on no AMD branch. The failure + says which of two things that commit is, because they need opposite actions: + **rebased** (AMD's pin holds a commit with the same author, author date and subject: + AMD rebased its line and dropped the old one, so the commit is AMD's and the pin is merged + by hand) or **foreign** (no such commit: it is not AMD's and moves to + `1bit/hrx-vulkan-patched`). `scripts/bump-hrx-sync.sh off-branch ` + prints the classification alone, from a commits-only clone; - merges AMD's pin into `1bit/hrx-vulkan-patched`, after tagging the previous tip `patched-` so every commit the engine has pinned stays reachable. Conflicts under `ggml/src/ggml-hrx` keep our side; a conflict anywhere else fails the run and names the files; - merges AMD's hrx-system pin into our `1bit/main` the same way; - opens a PR here moving the submodules, listing the llama.cpp commits new to the engine. +- **The mirror hold.** A rewrite of AMD's line that the rules above cannot merge leaves the daily run with + nothing to do but fail, every day, until somebody does the merge by hand. While + [`config/llama-mirror-hold.config`](../config/llama-mirror-hold.config) names the mirror's commit, the run + instead reports the hold in its summary (`mirror 1bit/hrx-vulkan HELD at …; AMD's pin is …, on a line that + does not contain it`), moves nothing and succeeds. The engine's pin does not move in a held run either: moving + it along our own line is a reviewed PR of its own. The hold + describes one state and fails the run when that state ends, so it cannot be forgotten: **stale** when the + mirror is no longer the named commit, **unneeded** when AMD's pin contains the named commit again (an ordinary + sync would work). `scripts/bump-hrx-sync.sh llama-hold ` prints the decision alone + (`none`, `held`, `stale`, `unneeded`). **Held since 2026-11-11** for the release (engine#1945): the file names + the mirror's commit `44ee5710d19f`, and the merge of AMD's rebased line is post-release work. It was added for AMD's 2026-11-09 rebase (engine#1945): taking that pin is a merge of + 149 conflicted files (538 hunks) into `1bit/hrx-vulkan-patched`, measured against the last AMD commit our + line took. - **The hold.** While [`config/hrx-system-hold.config`](../config/hrx-system-hold.config) names a commit, the run does not merge AMD's hrx-system, keeps `third_party/hrx-system` where it is, and takes only AMD's llama.cpp changes outside the HRX backend paths @@ -246,6 +265,163 @@ decode-path logits (`llama-perplexity -c 2049 -b 1 --save-all-logits`) are byte- to the previous pin, the buried code word is exact at 4700 tokens, and no GPU faults were observed. +### Fixed: one-token decode at a head size that is not a power of two ([engine#1926](https://github.com/1bit-MONSTER/engine/issues/1926)) + +A model whose attention head size is not a power of two (stories15M: 48; Sarashina2.2-3B: 160) and whose +Q weights are not Q4_K or Q6_K loaded and prefilled on `HRX0`, then failed at the first decode step: +`HRX Loom JIT Loom compilation failed`, `llama_decode: failed to decode, ret = -3`. The one-token Q +projection kernel wrote each rotated pair at `channel rem head_size`, and the kernel compiler can only +prove that two-wide store in bounds when the remainder is a mask. The kernel now indexes by rotation +pair, which is in bounds at every head size (llama.cpp #282). Q4_K and Q6_K models took a different +kernel at one token and were not affected. + +Measured at pin `a5dc597b8` on gfx1151 (2026-11-10, power mode `performance`): +`test-hrx-attention-q-decode-head-size` compares the fused projection with the CPU backend at head +sizes 48, 96 and 160 (controls 64 and 128) on BF16, F16 and Q8_0 weights: 9 of 15 cases failed before +the fix, 0 of 15 after. `llama-bench -p 512 -n 128 -r 3`, two runs each: Sarashina2.2-3B bf16 decodes +at 26.49 and 25.92 tok/s, stories15M Q8_0 at 317.77 and 338.54 tok/s (the box was shared; the same +binary gave 545.39 tok/s on a quieter run). + +### Fixed: all-NaN logits after a large expert activation ([engine#1539](https://github.com/1bit-MONSTER/engine/issues/1539), [#1962](https://github.com/1bit-MONSTER/engine/issues/1962)) + +`flint-40b-a2.5b` on `HRX0` answered some requests and then failed a later one with +`HRX returned NaN logits: count=150144 of 150144`. The matmul kernels multiply in half precision, and a +half-precision value ends at 65504. An FFN down-projection reads the SwiGLU output, which nothing +normalises: flint's layer 46 reaches about 9.1e4 on an ordinary 22-token prompt. That value became +`inf`, the sum over the routed experts `inf - inf = NaN`, and the NaN was stored as those tokens' +keys and values. The request that overflowed still returned finite logits. The next, shorter request +on the same slot found the stale NaN row inside its masked attention window and returned NaN for the +whole vocabulary. So the failure depended on the prompt (which tokens drive an activation that large) +and on the order of requests, not on memory, load or timing, and the CPU backend, which multiplies in +single precision, was never affected. + +The routed-expert kernel (llama.cpp #284) and the dense kernel for single-precision activations +(llama.cpp #285) now scale the activations by 2^-6 before the half-precision stage and scale the +result back by 2^6. Both factors are powers of two, so values that fit before are rounded as before; +the usable range moves from [6.1e-5, 6.55e4] to [3.9e-3, 4.19e6]. + +Measured on gfx1151 (2026-11-10): + +- flint, the sequence that failed on 9 of 9 fresh loads (a 22-token request, then an 8-token one on + the same slot): 150144 of 150144 NaN before, 0 after; a read-back of all 3056 graph nodes finds no + non-finite value, and the first short request now returns the same logits as its repeat. +- `test-backend-ops -o MUL_MAT_ID`: the new case at flint's shape with activations of 9.1e4 returned + NaN against the CPU's 56214.23 before, and passes after; the whole op is 845 of 845. `-o MUL_MAT`: + the 47 new large-activation cases went from 23 to 47 passing; the whole op is 589 of 589. +- The range the shift gives up at the low end is not used by the models measured: the smallest + activation column maximum is 3.2e-2 on flint's expert down-projections (37278 columns) and 2.4e-2 on + the dense inputs of Qwen3-0.6B, Qwen3-4B, ZAYA1-8B and flint. Cases at those two values are in the + suite, so the shift has a test on each side. + +The prefill kernels followed at pin `7a1f2b81c`. The two small-batch Q6_K kernels were already done +(llama.cpp #288). The Q4_K and Q6_K prefill kernels read a half-precision copy of the activations; +that copy is now stored at 2^-6 and the sums are scaled back (llama.cpp #298). The Q5_K/IQ4_XS +prefill kernel reads the activations as 8-bit blocks, each with a scale and a scaled sum in half +precision. The scale fits up to 8.3e6; the scaled sum is about as large as the largest activation in +the block, so it was the term that overflowed. The kernel now rebuilds a sum that is not finite from +the 8-bit values it already holds, and uses a finite one as stored (llama.cpp #292). + +Measured on gfx1151 (2026-11-11, fork `7a1f2b81c`, full build): + +- `test-backend-ops -o MUL_MAT` with a resident weight (2048 x 4096), Q4_K, Q5_K, Q6_K, IQ4_XS, Q8_0 + and Q4_0 at 1 to 512 tokens, with activation spikes of 3.0, 9.3e4 and 2e6: 180 of 180 pass (run at + llama.cpp #298's head). Before #292 and #298, Q5_K at 512 tokens and Q4_K and Q6_K at 256 and 512 + tokens returned NaN at 9.3e4. The whole op is 617 of 617; `MUL_MAT_ID` is 845 of 845. +- Perplexity on wikitext-2 (`-c 512 -b 512`, 16 chunks): Qwen3-4B-Q4_K_M 17.1836 (17.1864 before + #298), Qwen3-4B-IQ4_XS 17.6304 (the same before and after #292). +- Prefill rate on Qwen3-4B-Q4_K_M across #298, eight alternating runs on a loaded box: 2683 against + 2735 tok/s (ratio 0.98, pairs between 0.88 and 1.14). No regression shown; a cost of 2 to 3 % is + not excluded. + +Still open under #1962: the final projection of the draft head (a shortlist kernel) narrows its +activations too and has no test that reaches it. The other kernels that read the 8-bit blocks (the +fused SwiGLU, the routed kernels and the low-token Q4 paths) use the stored sum as it is and are not +measured above 65504. And the limits are moved, not removed: 4.19e6 for the half-precision copies, +8.3e6 for the 8-bit blocks. + +### Fixed: Qwen3-VL on `HRX0`: it decodes again, and its rotation runs on the device ([engine#1467](https://github.com/1bit-MONSTER/engine/issues/1467)) + +Two faults, both in the text model's graph rather than in the projector: + +- **A refused view (llama.cpp #283).** A deepstack model adds a slice of the wide input embedding to each of + its first layers, taken as a view of a view. With the first view's node in another split, the HRX graph + never holds the wide tensor, the slice could not alias its input, and the graph was refused at + `unsupported HRX node 20: VIEW`. So Qwen3-VL-2B could not decode even text, and every image request failed + with it. Such a view is external to the graph and a layout op writes nothing, so the scheduler now skips it. +- **The rotation on the CPU (llama.cpp #286).** The model's rotation is the interleaved multi-section one + with sections that do not collapse to the plain rotation, so 56 `ROPE` nodes per text graph ran on the CPU + and cut the graph into 114 splits. The two generic rotation kernels gained a mode that reads the position + section each pair selects; the text graph is now 2 or 3 splits with `PAD` (and `GET_ROWS` on token + batches) left on the CPU. + +Measured at pin `50d28600a` on gfx1151 (2026-11-10, power mode `performance`, Qwen3VL-2B-Instruct Q4_K_M with +`mmproj-Qwen3VL-2B-Instruct-F16`): `test-backend-ops -o ROPE -b HRX0` 166/166; 12 greedy text tokens +byte-identical to the CPU backend; `llama-bench -p 512 -n 64 -r 3` at 6,481 and 6,532 tok/s prefill, 133.7 +and 135.5 tok/s decode (2,331 and 57.8 before the rotation moved). `llama-server` with the projector on +`HRX0`: the first request for the 448x448 test image takes 2.99 s (`prompt_ms` 2,782 for 227 prompt tokens) +against the 16.97 s this issue was filed on, and answers "A red rectangle and a green circle." + +The projector's own graph still runs 103 of 626 nodes on the CPU for an image of that size (its 2-D rotation, +the attention score matmul and softmax, `IM2COL`, `UPSCALE`); with it offloaded, encoding the image is no +faster than with `--no-mmproj-offload` (0.5 to 1.1 s against 0.4 to 0.6 s). + +### Fixed: prefill attention lost its rate at long context ([engine#1989](https://github.com/1bit-MONSTER/engine/issues/1989)) + +The cost of prefill attention per key was measured at 0.0331 ms per KV token per 512-token graph on Qwen3-4B +(over 0 to 7.7k keys) and 0.0522 over 0 to 24k. Timed alone on `HRX0`, the attention dispatch costs what its +arithmetic costs up to about 8,192 keys; beyond that, and at particular key counts, it fell into a slower regime. +The cause was not in the attention kernel's arithmetic. From 256 query tokens the dispatch copies V into a +transposed buffer and walks it across channels, and two things about that step were slow at long key counts: + +- the buffer's rows were `2 x keys` bytes apart, which at key counts such as 16,384 and 32,768 puts every channel + of a key block on the same cache sets. A key count that is a multiple of 64 now gets a pitch of `keys + 32`; +- one launch over 512 query tokens was slower per token than two over 256. From 4,096 keys the dispatch now runs + 256-token launches over the one transposed buffer. + +The pitch is derived inside the attention kernel from its key count, so a new key count compiles no new kernel +(llama.cpp #289 and #291; #289 alone made it a compile parameter and must not be pinned without #291). + +Measured at pin `2293f95ea` on gfx1151 (2026-11-10, power mode `performance`, Qwen3-4B Q4_K_M, +`llama-bench -p 24576 -n 0 -b 512 -ub 512 -r 3`, old layout through `GGML_HRX_FA_VALUE_PITCH=0 +GGML_HRX_FA_QUERY_CHUNK=2048`, alternated): warm repetitions 813 to 822 tok/s before, 980 to 1,074 after; +first, cold repetitions 701 to 711 before, 856 to 866 after. At 8,192 tokens 1,545 and 1,600 before, 1,766 and +1,768 after; 2,048 and 512 tokens do not move. A cold run on a 9,230-token prompt takes the same time with either +layout (8.80 s and 8.85 s) and gives the same text. `test-hrx-fa-transposed-value` compares the path with the CPU +backend (8 cases). Why 256-token launches are faster than one 512-token launch at long key counts is measured, +not explained. + +### Fixed: a gated-delta-net layer behind a type cast failed the whole graph ([engine#1925](https://github.com/1bit-MONSTER/engine/issues/1925)) + +On Qwen3.8-Flash-Next, a long prompt failed at the first 512-token batch past 2,048 tokens of context when the +sparse attention path's mask chain ran on the device: + +``` +graph_compute: unsupported HRX node 160: L2_NORM output=f32[128,16,512,1]<-L2_NORM inputs=[f32[128,16,512,1]<-VIEW] consumers=[174:GATED_DELTA_NET] +``` + +The fused gated-delta-net dispatch did match that node. The fault was the order in which the graph was walked. +The scheduler walks a graph in dependency order and takes the matmuls first, which is what puts a layer's gate +and beta projections ahead of the `L2_NORM` its fused pattern starts from. A type cast is an in-place copy that +names its own output as an input. The walk counted that as a dependency, so the copy waited for itself, the walk +stopped there, and every node after it was taken in index order. In index order the `L2_NORM` comes before the +projections, and the pattern was refused because it would have read a projection that had not run. The sparse +path's mask chain contains such casts, so the failure needed both a long context and that chain on the device. + +The walk no longer counts a node's own output as a dependency (llama.cpp #302). The fused matcher also keeps a +match when a node of its pattern that only renames storage was already passed. + +Measured on gfx1151 (2026-11-11, Flash-Next UD-IQ3_XXS, `-c 19456 -np 1`, one 18,782-token request, a build of +the fork with only the two changed objects replaced): `unsupported HRX node` 1 before, 0 after; all 390 fused +gated-delta-net roots accepted; the whole prefill runs. `tests/test-hrx-gdn-passed-alias.cpp` plans the layer +behind a cast on the host and requires the same dispatches as without it. Qwen3.5-4B, which has the same layer +type and no sparse path, gives the same tokens on the old and the new library (128 of 128 in three alternated pairs), uses the same 33 +dispatch registrations in the same order, and runs at the same rate (`llama-bench -p 512 -n 64 -r 5`, old, new, +old, new: decode 57.4, 58.3, 58.1, 58.2 tok/s; prefill 3,516, 3,042, 3,109, 3,213). + +Not shown by this: that the answer at 18.8k is the same with the mask chain on the device as with it on the +CPU. Three of the four long requests in these runs, on the old library and the new, ended in all-NaN logits on +the first sampled token, which is [engine#1536](https://github.com/1bit-MONSTER/engine/issues/1536). + ### MoE decode and the kernel cache (llama.cpp #32, #33) - **MoE decode.** With one to four tokens, a `MUL_MAT_ID` (a MoE expert matmul) runs as a GEMV over @@ -727,6 +903,65 @@ repacked-prefill kernel ([llama.cpp #250](https://github.com/1bit-MONSTER/llama. `GGML_HRX_Q6K_PACKED_PREFILL`) stays **opt-in** for the same reason: it is numerically free and speed-neutral, so it buys nothing on its own. +**Measuring the duplicate on your own build.** `tools/hrx_gtt_split.py` splits one process's GTT into the loader's copy and HRX's +packed second copy, using only lines llama.cpp and HRX already print plus the process's own fdinfo. It needs no root and no +tracefs. Inside your ticket: + +``` +gpu-hold -l -t 5 -- tools/hrx_gtt_split.py run --out split-4b -- \ + llama-cli -m Qwen3-4B-Q4_K_M.gguf -ngl 99 -fit off -c 256 -n 16 --temp 0 -no-cnv -p "The capital of France is" +tools/hrx_gtt_split.py parse split-4b/split.log --samples split-4b/gtt.txt # again, from the files it left +``` + +`run` adds `-v`, sets `GGML_HRX_HOST_WEIGHT_TRACE=1`, and samples `drm-total-gtt` with each DRM client counted once. On the +#1351 pin (llama.cpp `40f58efe1`, the same command): + +| part | Qwen3-4B Q4_K_M | Qwen3.8-27B UD-Q4_K_XL | +|---|---:|---:| +| loader's model buffer (HRX) | 2,375.9 MiB | 15,718.5 MiB | +| HRX's packed copies | **795.8 MiB** (38 weights) | **2,631.3 MiB** (66 weights) | +| KV + recurrent state | 36.0 | 165.6 | +| compute + output buffers | 43.2 | 72.1 | +| not attributed | 674.1 | 1,909.5 | +| process peak GTT | 3,925 | 20,497 | +| share of the loader's copy resident twice | 33.5 % | 16.7 % | + +The packed copies are the duplicate. HRX builds one for each weight that a kernel reads in a re-laid-out form; prefill above +a token threshold still reads the loader's copy. "Not attributed" is mostly runtime module memory and copy staging (the +staging buffer is capped at 64 MiB since [llama.cpp #266](https://github.com/1bit-MONSTER/llama.cpp/pull/266)). It is not +weights. This replaces the `bo_trace.py` reproduction in the issue body, which was never committed. + +**Which kernel a graph's time is in.** `tools/hrx_dispatch_profile.py` gives the time per HRX dispatch, summed per kernel, +on any HRX build (it needs no rebuild). The backend's serial mode (`GGML_HRX_DEBUG_SERIAL_EXECUTION=1`) waits for the device +after every dispatch and prints a line at launch and at completion; the tool sets the switch, time-stamps those lines and +names the kernels from the pinned fork's manifests. Inside your ticket: + +``` +gpu-hold -l -t 10 -- tools/hrx_dispatch_profile.py run --log prof.log --json prof.json -- \ + llama-bench -m Qwen3-4B-Q4_K_M.gguf -ngl 99 -p 8192 -n 16 -r 1 +tools/hrx_dispatch_profile.py parse prof.log --top 20 # again, from the trace it kept +``` + +It prints one table per command program (a graph shape: prefill batch, decode step). Measured on the pin `2293f95ea`, +Qwen3-4B Q4_K_M, the prefill program of a 512-token prompt against the most frequent prefill program of an 8,192-token +prompt (12 of its 16 batches of 512), one run each, 2026-11-11: + +| kernel | 512-token prompt, ms each | share | 8,192-token prompt, ms each | share | +|---|---:|---:|---:|---:| +| `ggml_flash_attention_f32_f16_wmma` | 0.371 | 7.5% | 2.150 | 49.3% | +| `ggml_mul_mat_swiglu_q4_k_f16_wmma_prefill_wave32` | 1.714 | 33.6% | 1.728 | 19.3% | +| `ggml_mul_mat_q4_k_f16_wmma_prefill_wave32` | 0.425 | 34.2% | 0.393 | 18.0% | +| `ggml_mul_mat_q6_k_f16_wmma_prefill_wave32` | 0.635 | 12.4% | 0.632 | 7.0% | + +The matmuls cost the same at any context; the attention kernel is the part that grows. + +What the numbers are and are not: each is the wall time from launch to device-idle with one dispatch in flight, so it holds +one wake-up of the waiting thread and the run is slower than a normal one (the same command: pp512 2,977 -> 2,500 tok/s, +tg16 67 -> 36, pp8192 1,191 -> 1,167). Use it to rank kernels and to see what grows with the context, not as a rate. +Decode is slowed much more than prefill, so compare shares inside one program and never between programs. A dispatch that +was launched and never finished (a ticket that ran out, a cut trace) is listed under the tables and is not in them. A build +with extra kernel directories shows those kernels as numeric ids; name them with a second `--manifests DIR`. + ### Measuring on `HRX0`: what makes a number citable Numbers measured on this device swing far more between *runs* than between *files*, so the harness matters as much as the arm. Everything below is diff --git a/docs/hw-gate.md b/docs/hw-gate.md index 1d7afdeef..ed0dedb4a 100644 --- a/docs/hw-gate.md +++ b/docs/hw-gate.md @@ -72,7 +72,9 @@ hrx-system) through `cmake/hrx.cmake`, then links `onebit` with its HRX wiring. - **Memory.** The box is shared, and on Strix Halo GTT is host RAM. Before it compiles, the job waits for 24 GB of `MemAvailable` (`scripts/mem-wait.sh`, polled once a minute for up to 15 minutes). If memory doesn't free up, the job fails with the reading, and a re-run is the fix. The compile itself runs under `scripts/mem-guard.sh --term 12`, - which stops the whole `thermal-run` scope if `MemAvailable` falls below 12 GB. On 2026-11-08 both self-hosted + which stops the whole `thermal-run` scope if `MemAvailable` falls below 12 GB. **Exit 143 in a build step is + that guard, not a compiler failure**: the step log has a `GUARD: MemAvailable N GB < 12 GB` line just above it, the + runner is still alive, and a re-run once memory frees is the fix. On 2026-11-08 both self-hosted runners were OOM-killed, the second time 76 s after a restart handed this job a compile (engine#1504). - **What it is not.** It is a compile and link check of the public plumbing only. It does not build the private HIP add-on (the gate builds that) and does not dispatch anything. Green means the pinned pair builds into the engine. It @@ -250,6 +252,87 @@ The settings are environment variables at the top of the script: `HW_GATE_STORE` `$HW_GATE_OUT` (default `~/.cache/hw-gate/runs/