diff --git a/CMakeLists.txt b/CMakeLists.txt index 79f7d6e..2009689 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -55,6 +55,13 @@ set_property(CACHE SAE_GEMM_BACKEND PROPERTY STRINGS ROCM CUDA CPU) # default so ROCm/CPU builds don't reference unavailable EPs. option(SAE_ORT_TRT_EP "ORT backend: enable TensorRT/CUDA execution providers" OFF) +# AR-026/AR-027: the CPU GEMM path is backed by OpenBLAS, and its absence is a +# configure error rather than a silent downgrade to the scalar loop. Declared at +# top level because the unit-test target compiles the CPU kernel regardless of +# which backend the main build selected, and both must make the same choice. +option(SAE_ALLOW_SCALAR_GEMM + "Permit the scalar-loop GEMM fallback when OpenBLAS is absent" OFF) + # Back-compat: a legacy -DSAE_WITH_TRT=ON/OFF seeds the new vars (ON⇒TRT+CUDA, # OFF⇒ORT+ROCM) unless the user set them explicitly. if(DEFINED SAE_WITH_TRT) @@ -153,10 +160,16 @@ if(SAE_GEMM_BACKEND STREQUAL "CPU") target_include_directories(gemm_backend PRIVATE src) target_compile_definitions(gemm_backend PRIVATE SAE_GEMM_CPU) - # AR-026/AR-027: back the CPU path with OpenBLAS when present. Optional, so - # the build gains no hard dependency — but without it the fallback is a - # scalar loop, which does not hold up against a library-scale gallery, and - # the CPU path is exactly what CI (no GPU) and the cpu builder image use. + # AR-026/AR-027: the CPU path is backed by OpenBLAS, and that is REQUIRED + # rather than opportunistic. The CPU backend is what CI (no GPU) and the cpu + # builder image actually run, so a silent fall back to the scalar loop means + # AR-027 is measured — or worse, believed — on a path no release uses. A + # missing dependency should stop the build and name itself, not degrade into + # a slower answer nobody notices. + # + # The scalar loop survives as the correctness oracle the two backends are + # diffed against; -DSAE_ALLOW_SCALAR_GEMM=ON is how you ask for it, which + # keeps that an explicit, visible choice. find_package(PkgConfig QUIET) if(PkgConfig_FOUND) pkg_check_modules(OPENBLAS QUIET openblas) @@ -166,9 +179,16 @@ if(SAE_GEMM_BACKEND STREQUAL "CPU") target_compile_definitions(gemm_backend PRIVATE SAE_GEMM_CBLAS) target_include_directories(gemm_backend PRIVATE ${OPENBLAS_INCLUDE_DIRS}) target_link_libraries(gemm_backend PRIVATE ${OPENBLAS_LINK_LIBRARIES}) + elseif(SAE_ALLOW_SCALAR_GEMM) + message(WARNING "GEMM backend: CPU scalar fallback (SAE_ALLOW_SCALAR_GEMM=ON) — " + "correct, but slow on a large gallery. Do not measure AR-027 here.") else() - message(WARNING "GEMM backend: CPU scalar fallback — OpenBLAS not found. " - "Correct, but slow on a large gallery (AR-027).") + message(FATAL_ERROR + "OpenBLAS not found, and the CPU GEMM backend requires it (AR-026/AR-027).\n" + " Install it: Fedora dnf install openblas-devel\n" + " Arch pacman -S openblas\n" + " Debian apt install libopenblas-dev\n" + " Or build the scalar fallback deliberately: -DSAE_ALLOW_SCALAR_GEMM=ON") endif() elseif(SAE_GEMM_BACKEND STREQUAL "CUDA") find_library(CUBLAS_LIB cublas diff --git a/docs/SPEC.md b/docs/SPEC.md index 3a6d27a..6f8fc46 100644 --- a/docs/SPEC.md +++ b/docs/SPEC.md @@ -958,21 +958,45 @@ is the only viable formulation — a per-pair loop is orders of magnitude off. **All similarity computation goes through the GEMM path**, with no exception justified by "this set is small". Three call sites: -1. **Baked gallery** — already GEMM (`sim_engine_->compute()`, - `identity_matcher_node.hpp:143`, backend from `SAE_GEMM_BACKEND`). ✓ -2. **Per-film annex** — currently a **CPU loop** - (`identity_matcher_node.hpp:159-162`), justified in-comment by "tens of - embeddings". AR-018…AR-021 invalidates that assumption: every owned track now - contributes, so the annex grows with cast size and film length. It must move - into the GEMM path — appended to the gallery matrix, or a second multiply. +1. **Baked gallery** — GEMM (`sim_engine_->compute()`, backend from + `SAE_GEMM_BACKEND`). ✓ +2. **Per-film annex** — was a **CPU loop**, justified in-comment by "tens of + embeddings". AR-018…AR-021 invalidated that assumption: every owned track + contributes, so the annex grows with cast size and film length. Now appended + to the gallery matrix rather than scored separately — promotions are pushed + into the engine's resident matrix (`ISimilarityEngine::append_rows`, + capacity doubling, device-to-device on the GPU backends) and `flat_actor_` + grows in lockstep, so one multiply covers baked and promoted references and + best-of-N is a single pass over one similarity column. ✓ 3. **Deferred TBI pass (AR-020)** — the most GEMM-friendly operation in the pipeline: all TBI embeddings against the full gallery-plus-annex, offline, operands resident, no streaming. One large multiply, not a loop over entries. + Not yet built; AR-020 owns it. + +**Current:** 1 and 2 done. `TrackGallery` holds the annex as a contiguous +row-major matrix plus a parallel actor index, and hands newly promoted rows to +the matcher once per frame (`drain_promotions`), which is what call site 3 will +score against. + +**Gap:** call site 3, gated on AR-020 existing at all. This constrains AR-018…AR-021's implementation: the annex must be a **contiguous matrix** with promotions appended, plus a parallel actor-index mapping — exactly the `flat_emb_`/`flat_actor_` arrangement the baked gallery already uses. +**Ordering note.** Absorbing promotions is a once-per-frame step that runs after +every face in the frame has been scored, not mid-frame. Appending mid-frame would +invalidate the similarity pointer the matcher is still reading, and it also +removes an accidental dependence on face order within a frame: a promotion helps +subsequent frames, never the one that produced it, which is the semantics the +expansion store already documents. + +**The CPU GEMM path requires OpenBLAS.** It is what CI and the cpu builder image +run, so a silent fall back to the scalar loop would mean AR-027 is measured — or +believed — on a path no release uses. Absence is a configure error; the scalar +loop survives as the correctness oracle, reachable only via +`-DSAE_ALLOW_SCALAR_GEMM=ON`. + ### Scaling characteristics that must be known, not assumed - **Throughput versus gallery size must be measured** (VR-008) and published. The diff --git a/docs/plan.md b/docs/plan.md index b1a606b..c0623b3 100644 --- a/docs/plan.md +++ b/docs/plan.md @@ -255,8 +255,11 @@ Context crops opt-in behind `--dump-unidentified-crops`. ## AR-026, AR-027 — GEMM and scale -**Depends on:** nothing to start. The annex CPU loop -(`identity_matcher_node.hpp:159-162`) moves into the GEMM path. +**Depends on:** nothing to start. The annex CPU loop has moved into the GEMM +path: the annex is a contiguous matrix, promotions are appended to the engine's +resident gallery, and the CPU backend now requires OpenBLAS. What is left of +AR-026 is call site 3, the deferred pass — so the rest of AR-026 lands *with* +AR-020 rather than before it. --- diff --git a/docs/requirements.md b/docs/requirements.md index 5d49c83..d1a2e2a 100644 --- a/docs/requirements.md +++ b/docs/requirements.md @@ -53,7 +53,7 @@ Status: `Done` · `In Progress` · `Planned` · `TBD` · `Withdrawn` | AR-023 | Fit sigmoid calibration from intra/inter similarity distributions | SR-002 | High | Done | | AR-024 | **Always the calibrated probability, never a raw cosine** — exceptions recorded | SR-002 | High | **Done** — association, accumulation and expansion all in probability space; `track_max_embed_dist`, `cut_revive_sim`, `expand_novelty_sim`, `expand_track_spread_max` retired | | AR-025 | Per-track Bayesian accumulation in log-odds, with correlated-observation discounting | SR-002 | High | **Done** — log-odds accumulation with correlation discounting owned by the registry, `src/evidence_discount.hpp` | -| AR-026 | All similarity computed as GEMM, including annex and deferred pass | SR-001 | High | In Progress | +| AR-026 | All similarity computed as GEMM, including annex and deferred pass | SR-001 | High | **In Progress** — two of the three call sites done. Baked gallery was already GEMM; the annex now is too — it is a contiguous row-major matrix (`track_gallery.hpp`) whose promoted rows are appended to the engine's resident matrix (`ISimilarityEngine::append_rows`), so one multiply covers baked and promoted references and the host-side cosine loop is gone. CPU path requires OpenBLAS (scalar fallback now opt-in behind `SAE_ALLOW_SCALAR_GEMM`). Remaining: the deferred pass, which does not exist until AR-020 | | AR-027 | Throughput acceptable for **arbitrary** gallery size | SR-001 | High | Planned | | AR-028 | **Embedding input quality assessed and carried** — every face scored on size, sharpness and visibility before its embedding is used as identity evidence; the vector travels with the face and reaches the VR-001 dump | SR-002 | High | Planned | | AR-029 | Sharpness measure on the **aligned crop** (scale-normalised, so it cannot re-measure size) | SR-002 | Medium | Planned | diff --git a/src/backends/gemm_backend.cpp b/src/backends/gemm_backend.cpp index d1d09c4..af10550 100644 --- a/src/backends/gemm_backend.cpp +++ b/src/backends/gemm_backend.cpp @@ -42,17 +42,20 @@ constexpr int kDim = 512; // Used for CI and as the correctness oracle for the GPU backends. // // TRACES: AR-026, AR-027 | SR-001 -// Backed by CBLAS (OpenBLAS) when available, falling back to a scalar loop when -// not. The fallback is portable but scales badly: scoring one face against a +// Backed by CBLAS (OpenBLAS), which CMake now REQUIRES for this backend. The +// scalar loop below is portable but scales badly: scoring one face against a // 5000-embedding gallery is 2.6 MFLOP, and a crowded frame multiplies that by // the face count. Since AR-003 removed the per-frame face cap and CI has no GPU, -// the CPU path is now the one that has to hold up under a library-scale gallery -// (AR-027) rather than merely be correct. +// the CPU path is the one that has to hold up under a library-scale gallery +// (AR-027) rather than merely be correct — so falling back to it silently would +// mean measuring AR-027 on a path no release runs. // -// The fallback is kept rather than made mandatory so the build has no hard new -// dependency, and so the two can be diffed when a similarity looks wrong. The gallery is L2-normalised (as are the queries), so each -// similarity is a plain dot product. S is stored column-major to match the GPU -// backends: the gallery similarities for face fi start at result + fi*n_gallery. +// The fallback is kept as the correctness oracle the two BLAS backends are +// diffed against when a similarity looks wrong, and is reachable only via +// -DSAE_ALLOW_SCALAR_GEMM=ON. The gallery is L2-normalised (as are the queries), +// so each similarity is a plain dot product. S is stored column-major to match +// the GPU backends: the gallery similarities for face fi start at +// result + fi*n_gallery(). class SimilarityEngine final : public ISimilarityEngine { public: SimilarityEngine(const float* gallery_row_major, int n_gallery, int max_faces) @@ -72,6 +75,19 @@ public: } int max_faces() const override { return max_faces_; } + int n_gallery() const override { return n_gallery_; } + + /// TRACES: AR-026 | SR-001 + /// Promotions join the resident matrix, so the annex is scored by the same + /// SGEMM as the baked references. std::vector already grows geometrically, + /// so this is amortised O(1) per row. + void append_rows(const float* rows_row_major, int n_rows) override { + if (n_rows <= 0) return; + gallery_.insert(gallery_.end(), rows_row_major, + rows_row_major + static_cast(n_rows) * kDim); + n_gallery_ += n_rows; + host_sims_.resize(static_cast(max_faces_) * n_gallery_); + } const float* compute(const float* query_row_major, int n_faces) override { if (n_faces <= 0) return host_sims_.data(); @@ -143,6 +159,7 @@ inline void gpu_free(void* p) { cudaFree(p) inline void gpu_memcpy_h2d(void* dst, const void* src, size_t n, stream_t s) { check_gpu(cudaMemcpyAsync(dst, src, n, cudaMemcpyHostToDevice, s), "H2D"); } inline void gpu_memcpy_d2h(void* dst, const void* src, size_t n, stream_t s) { check_gpu(cudaMemcpyAsync(dst, src, n, cudaMemcpyDeviceToHost, s), "D2H"); } inline void gpu_memcpy_h2d_sync(void* dst, const void* src, size_t n) { check_gpu(cudaMemcpy(dst, src, n, cudaMemcpyHostToDevice), "H2D_sync"); } +inline void gpu_memcpy_d2d_sync(void* dst, const void* src, size_t n) { check_gpu(cudaMemcpy(dst, src, n, cudaMemcpyDeviceToDevice), "D2D_sync"); } inline void stream_create(stream_t* s) { check_gpu(cudaStreamCreate(s), "cudaStreamCreate"); } inline void stream_destroy(stream_t s) { cudaStreamDestroy(s); } inline void stream_sync(stream_t s) { check_gpu(cudaStreamSynchronize(s), "cudaStreamSync"); } @@ -177,6 +194,7 @@ inline void gpu_free(void* p) { (void)hipFr inline void gpu_memcpy_h2d(void* dst, const void* src, size_t n, stream_t s) { check_gpu(hipMemcpyAsync(dst, src, n, hipMemcpyHostToDevice, s), "H2D"); } inline void gpu_memcpy_d2h(void* dst, const void* src, size_t n, stream_t s) { check_gpu(hipMemcpyAsync(dst, src, n, hipMemcpyDeviceToHost, s), "D2H"); } inline void gpu_memcpy_h2d_sync(void* dst, const void* src, size_t n) { check_gpu(hipMemcpy(dst, src, n, hipMemcpyHostToDevice), "H2D_sync"); } +inline void gpu_memcpy_d2d_sync(void* dst, const void* src, size_t n) { check_gpu(hipMemcpy(dst, src, n, hipMemcpyDeviceToDevice), "D2D_sync"); } inline void stream_create(stream_t* s) { check_gpu(hipStreamCreate(s), "hipStreamCreate"); } inline void stream_destroy(stream_t s) { (void)hipStreamDestroy(s); } inline void stream_sync(stream_t s) { check_gpu(hipStreamSynchronize(s), "hipStreamSync"); } @@ -201,14 +219,15 @@ public: SimilarityEngine(const float* gallery_row_major, int n_gallery, int max_faces) : n_gallery_(n_gallery), max_faces_(max_faces) { - const size_t gallery_floats = static_cast(n_gallery_) * kDim; - gpu_malloc(reinterpret_cast(&d_gallery_), gallery_floats * sizeof(float)); - gpu_memcpy_h2d_sync(d_gallery_, gallery_row_major, gallery_floats * sizeof(float)); - gpu_malloc(reinterpret_cast(&d_query_), static_cast(max_faces_) * kDim * sizeof(float)); - gpu_malloc(reinterpret_cast(&d_sims_), - static_cast(max_faces_) * n_gallery_ * sizeof(float)); + + // Allocates d_gallery_/d_sims_ at the initial row count; append_rows() + // grows them geometrically from here. + reserve_rows(std::max(n_gallery_, 1)); + const size_t gallery_floats = static_cast(n_gallery_) * kDim; + if (gallery_floats) + gpu_memcpy_h2d_sync(d_gallery_, gallery_row_major, gallery_floats * sizeof(float)); stream_create(&stream_); blas_create(&handle_); @@ -232,6 +251,24 @@ public: SimilarityEngine& operator=(const SimilarityEngine&) = delete; int max_faces() const override { return max_faces_; } + int n_gallery() const override { return n_gallery_; } + + /// TRACES: AR-026 | SR-001 + /// Promotions join the GPU-resident matrix, so the annex is scored by the + /// same SGEMM as the baked references rather than by a host-side loop. + /// Capacity doubles on overflow, so the gallery is re-uploaded O(log n) + /// times over a film rather than once per promotion. + void append_rows(const float* rows_row_major, int n_rows) override { + if (n_rows <= 0) return; + const int want = n_gallery_ + n_rows; + if (want > capacity_) reserve_rows(std::max(want, capacity_ * 2)); + + gpu_memcpy_h2d_sync(d_gallery_ + static_cast(n_gallery_) * kDim, + rows_row_major, + static_cast(n_rows) * kDim * sizeof(float)); + n_gallery_ = want; + host_sims_.resize(static_cast(max_faces_) * n_gallery_); + } const float* compute(const float* query_row_major, int n_faces) override { if (n_faces <= 0) return host_sims_.data(); @@ -251,7 +288,35 @@ public: } private: + // Grow the resident gallery (and the similarity output sized against it) to + // `rows` capacity, preserving the n_gallery_ rows already there. The copy is + // device-to-device, so a promotion never re-uploads the baked gallery across + // the bus. + void reserve_rows(int rows) { + if (rows <= capacity_) return; + + float* d_new_gallery = nullptr; + gpu_malloc(reinterpret_cast(&d_new_gallery), + static_cast(rows) * kDim * sizeof(float)); + if (d_gallery_ && n_gallery_ > 0) + gpu_memcpy_d2d_sync(d_new_gallery, d_gallery_, + static_cast(n_gallery_) * kDim * sizeof(float)); + if (d_gallery_) gpu_free(d_gallery_); + d_gallery_ = d_new_gallery; + + // S is (capacity × n_faces); its contents are rewritten by every + // compute(), so this one is a plain reallocation with nothing to keep. + float* d_new_sims = nullptr; + gpu_malloc(reinterpret_cast(&d_new_sims), + static_cast(max_faces_) * rows * sizeof(float)); + if (d_sims_) gpu_free(d_sims_); + d_sims_ = d_new_sims; + + capacity_ = rows; + } + int n_gallery_{0}; + int capacity_{0}; int max_faces_{0}; float* d_gallery_{nullptr}; float* d_query_{nullptr}; diff --git a/src/gallery/track_gallery.hpp b/src/gallery/track_gallery.hpp index db6cebe..dfd82e0 100644 --- a/src/gallery/track_gallery.hpp +++ b/src/gallery/track_gallery.hpp @@ -46,18 +46,19 @@ // cosines this replaces, expand_novelty_sim and expand_track_spread_max, are // retired under AR-024. // -// The annex is CPU-side and in-memory: it is small (tens of embeddings) so the -// matcher scans it with a scalar loop, and it is discarded when the process -// exits. Promoted embeddings only help SUBSEQUENT frames and later tracks of A — -// the pipeline stays streaming, no emitted output is buffered or relabelled. +// TRACES: AR-026 | SR-001 +// The annex is in-memory and discarded when the process exits, but it is NOT +// small: every owned track contributes, so it grows with cast size and film +// length. It is therefore held as a contiguous row-major matrix with a parallel +// actor index — the same flat_emb_/flat_actor_ shape the baked gallery uses — +// and the matcher hands promoted rows to the similarity engine rather than +// scanning them with a host-side loop. The deferred pass (AR-020) needs the same +// contiguous operand to score the TBI queue against in one multiply. +// +// Promoted embeddings only help SUBSEQUENT frames and later tracks of A — the +// pipeline stays streaming, no emitted output is buffered or relabelled. struct TrackGallery { - // One promoted reference view held in the per-actor annex. - struct AnnexEntry { - Embedding emb; - int actor_idx{-1}; - }; - explicit TrackGallery(const Config& cfg) : enabled_(cfg.expand_gallery) , buffer_size_(std::max(1, cfg.expand_buffer_size)) @@ -80,10 +81,39 @@ struct TrackGallery { bool enabled() const { return enabled_; } - // Current annex contents (empty when disabled). The matcher scans these - // alongside the baked gallery so a promoted view can win best-of-N for its - // actor. Returned by const-ref; only grows, never reordered. - const std::vector& annex() const { return annex_; } + /// TRACES: AR-026 | SR-001 + /// The annex as a contiguous row-major matrix (annex_size() × 512) plus the + /// parallel actor index. Only ever grows, never reordered, so a row index is + /// stable for the life of the film — which is what lets the similarity + /// engine hold the same rows and the actor mapping stay a plain vector. + int annex_size() const { return static_cast(annex_actor_.size()); } + const float* annex_data() const { return annex_emb_.data(); } + const std::vector& annex_actors() const { return annex_actor_; } + + /// One annex row (512 floats). The deferred pass (AR-020) scores the whole + /// matrix at once via annex_data(); this is for inspecting a single view. + const float* annex_row(int i) const { + return annex_emb_.data() + static_cast(i) * kEmbDim; + } + + /// TRACES: AR-026 | SR-001 + /// Hand the caller every row promoted since the previous call, appending to + /// its buffers, and return how many. The matcher pushes these into the + /// similarity engine so the next frame's single GEMM covers the annex — + /// draining rather than re-reading the whole matrix keeps that O(promoted), + /// not O(annex), per frame. + int drain_promotions(std::vector& emb_out, std::vector& actor_out) { + const int pending = annex_size() - drained_; + if (pending <= 0) return 0; + + emb_out.insert(emb_out.end(), + annex_emb_.begin() + static_cast(drained_) * kEmbDim, + annex_emb_.end()); + actor_out.insert(actor_out.end(), + annex_actor_.begin() + drained_, annex_actor_.end()); + drained_ = annex_size(); + return pending; + } // Offer one observed face to its track's diversity buffer. // track_id : face_tracker track (−1 = untracked, ignored) @@ -245,7 +275,10 @@ private: int added = 0; for (const auto& be : ts.buf) { - annex_.push_back({be.emb, actor}); + // Row-major append: the matrix stays contiguous so the matcher can + // hand whole blocks of new rows to the GEMM path (AR-026). + annex_emb_.insert(annex_emb_.end(), be.emb.begin(), be.emb.end()); + annex_actor_.push_back(actor); if (!debug_dir_.empty() && !be.crop.empty()) dump_mugshot(track_id, actor, added, be); ++added; @@ -255,7 +288,7 @@ private: << " confirmed actor " << actor << " (" << ts.accepted_frames << " accepted frames, worst " << "pairwise P=" << worst << ") — promoted " << added - << " views; annex now " << annex_.size() << "\n"; + << " views; annex now " << annex_size() << "\n"; } /// Prefer the registry's verdict; fall back to the local tally only when no @@ -320,5 +353,12 @@ private: std::string debug_dir_; std::map tracks_; - std::vector annex_; + + /// TRACES: AR-026 | SR-001 + /// Contiguous annex matrix and its parallel actor index. `drained_` marks + /// how much of it the similarity engine already holds. + static constexpr int kEmbDim = 512; + std::vector annex_emb_; ///< annex_size() × 512, row-major + std::vector annex_actor_; ///< actor index per annex row + int drained_{0}; }; diff --git a/src/inference/similarity.hpp b/src/inference/similarity.hpp index c1efe5a..eeee39e 100644 --- a/src/inference/similarity.hpp +++ b/src/inference/similarity.hpp @@ -15,6 +15,15 @@ // by CMake (SAE_GEMM_BACKEND); backends/gemm_backend.cpp provides // make_similarity_engine(). The core matcher node sees only this interface and // holds no CUDA/HIP/BLAS headers. +// +// TRACES: AR-026 | SR-001 +// The resident matrix GROWS. Per-film expansion (AR-018/AR-019) promotes new +// reference views mid-film, and those have to be scored by the same multiply as +// the baked references rather than by a side loop — "this set is small" is not +// an exception, because the annex grows with cast size and film length. Rows are +// therefore appended to the resident matrix and the next compute() covers baked +// and promoted references alike, in one GEMM. The deferred pass (AR-020) then +// inherits a single contiguous operand to score the TBI queue against. struct ISimilarityEngine { virtual ~ISimilarityEngine() = default; @@ -22,11 +31,23 @@ struct ISimilarityEngine { // Largest n_faces accepted by compute() per call (bounds GPU buffer sizes). virtual int max_faces() const = 0; + // Rows currently resident: the baked gallery plus every appended promotion. + // This is compute()'s column stride, and it changes as rows are appended — + // read it per call rather than caching it across frames. + virtual int n_gallery() const = 0; + + /// TRACES: AR-026 | SR-001 + /// Append n_rows unit-norm embeddings (row-major, 512 floats each) to the + /// resident matrix. Amortised O(1) per row: capacity grows geometrically, so + /// a promotion does not re-upload the gallery. Invalidates any pointer + /// previously returned by compute(). + virtual void append_rows(const float* rows_row_major, int n_rows) = 0; + // Compute similarities for n_faces query embeddings. // query_row_major: n_faces × 512, row fi at query + fi*512. // Returns a pointer to host memory holding S column-major: the gallery - // similarities for face fi start at result + fi*n_gallery. The pointer is - // owned by the engine and valid until the next compute() call. + // similarities for face fi start at result + fi*n_gallery(). The pointer is + // owned by the engine and valid until the next compute() or append_rows(). virtual const float* compute(const float* query_row_major, int n_faces) = 0; }; diff --git a/src/nodes/identity_matcher_node.hpp b/src/nodes/identity_matcher_node.hpp index 776e658..a281688 100644 --- a/src/nodes/identity_matcher_node.hpp +++ b/src/nodes/identity_matcher_node.hpp @@ -39,6 +39,13 @@ // uploaded and a single SGEMM computes the full similarity matrix in well under // a millisecond. The GPU math backend (cuBLAS or rocBLAS) lives behind // ISimilarityEngine (backends/gemm_backend.cpp) and is selected at compile time. +// +// TRACES: AR-026 | SR-001 +// That resident matrix grows during a film: per-film expansion (AR-019) promotes +// pose-varied views, and they are APPENDED to it rather than scored separately, +// so one multiply covers baked and promoted references alike and best-of-N is a +// single pass over one similarity column. There is no second similarity path in +// this node to fall out of step with the first. struct IdentityMatcherFunc { static constexpr std::string_view label() { return "identity_matcher"; } @@ -183,29 +190,30 @@ struct IdentityMatcherFunc { tf.embeddings[base + k].data(), 512 * sizeof(float)); } - // S (N_gallery × chunk) col-major: face k's gallery sims at sims + k*n_gallery. + /// TRACES: AR-026 | SR-001 + // One GEMM now covers baked references AND the per-film annex: promoted + // rows were appended to the engine's resident matrix, so they are just + // more gallery rows with an entry in flat_actor_. The annex used to be + // folded in afterwards by a host-side cosine loop, justified by "tens of + // embeddings" — an assumption AR-018/AR-019 retired, since every owned + // track promotes and the annex grows with cast size and film length. + // + // n_gallery() is read per frame, not cached: it grows as promotions land. + const int n_gal = sim_engine_->n_gallery(); const float* host_sims = sim_engine_->compute(host_query.data(), chunk); for (int ci = 0; ci < chunk; ++ci) { const int fi = base + ci; - const float* sims = host_sims + static_cast(ci) * n_gallery_; + const float* sims = host_sims + static_cast(ci) * n_gal; std::vector best_sim(gallery_.actors.size(), -std::numeric_limits::max()); - for (int ei = 0; ei < n_gallery_; ++ei) { + for (int ei = 0; ei < n_gal; ++ei) { float sim = sims[ei]; int ai = flat_actor_[ei]; if (sim > best_sim[ai]) best_sim[ai] = sim; } - // Fold in the per-film annex (CPU-side, tens of embeddings). Promoted - // pose-varied views compete for best-of-N exactly like baked refs, so - // a face at a pose the gallery lacked can now win its true actor. - for (const auto& ae : track_gallery_.annex()) { - float sim = cosine_similarity(tf.embeddings[fi], ae.emb); - if (sim > best_sim[ae.actor_idx]) best_sim[ae.actor_idx] = sim; - } - int best_actor = -1; int second_actor = -1; float best_s = -std::numeric_limits::max(); @@ -299,10 +307,34 @@ struct IdentityMatcherFunc { } } // chunk loop + absorb_promotions(); + return {std::move(tf.source), std::move(actors)}; } private: + /// TRACES: AR-026 | SR-001 + /// Move rows promoted during this frame into the resident gallery matrix, + /// extending the actor mapping in lockstep so row i keeps naming the actor + /// at flat_actor_[i]. Runs once per frame, after every face has been scored: + /// appending mid-frame would invalidate the similarity pointer the chunk + /// loop is still reading, and it is also the semantics the expansion store + /// documents — a promotion helps SUBSEQUENT frames, never the one that + /// produced it, so identification cannot depend on face order within a frame. + void absorb_promotions() { + if (!track_gallery_.enabled()) return; + + pending_emb_.clear(); + pending_actor_.clear(); + const int n = track_gallery_.drain_promotions(pending_emb_, pending_actor_); + if (n == 0) return; + + sim_engine_->append_rows(pending_emb_.data(), n); + flat_actor_.insert(flat_actor_.end(), + pending_actor_.begin(), pending_actor_.end()); + n_gallery_ = sim_engine_->n_gallery(); + } + ActorGallery gallery_; GalleryCalibration cal_; float prob_threshold_; @@ -310,10 +342,18 @@ private: float threshold_; float ratio_; float ratio_ceil_; + /// flat_emb_ is the BAKED reference set only — it is the calibration fit's + /// input (AR-023) and is not touched again after construction. flat_actor_, + /// by contrast, is the actor mapping parallel to the *engine's* rows, so it + /// grows with every promotion absorbed (AR-026) and is the longer of the two. std::vector flat_emb_; std::vector flat_actor_; int n_gallery_{0}; + // Reused across frames so absorbing a promotion allocates nothing. + std::vector pending_emb_; + std::vector pending_actor_; + std::unique_ptr sim_engine_; TrackGallery track_gallery_; std::shared_ptr registry_; diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index dfb2587..fe8a4f2 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -46,6 +46,14 @@ endif() if(OPENBLAS_T_FOUND) target_include_directories(sae_tests PRIVATE ${OPENBLAS_T_INCLUDE_DIRS}) target_link_libraries(sae_tests PRIVATE ${OPENBLAS_T_LINK_LIBRARIES}) +elseif(NOT SAE_ALLOW_SCALAR_GEMM) + # Same rule as the CPU backend itself: testing the scalar loop while the + # shipped CPU path is OpenBLAS means the suite is not evidence about the + # kernel that runs. + message(FATAL_ERROR + "OpenBLAS not found, and the unit tests compile the CPU GEMM kernel " + "(AR-026). Install openblas-devel, or pass -DSAE_ALLOW_SCALAR_GEMM=ON " + "to test the scalar fallback deliberately.") endif() target_compile_definitions(sae_tests PRIVATE diff --git a/tests/test_similarity.cpp b/tests/test_similarity.cpp index b508077..595448a 100644 --- a/tests/test_similarity.cpp +++ b/tests/test_similarity.cpp @@ -1,4 +1,4 @@ -// TRACES: AR-026 | SR-001 +// TRACES: UT-004 | AR-026 | SR-001 // // Unit tests for the CPU reference similarity engine (backends/gemm_backend.cpp, // SAE_GEMM_CPU) and the l2_normalise helper. All pure, GPU-free, model-free. @@ -82,6 +82,107 @@ TEST_CASE("CPU similarity engine matches hand-computed dot products", "[similari CHECK_THAT(S[2 + 1 * n_gallery], WithinAbs(s, 1e-6f)); } +// TRACES: UT-004 | AR-026 | SR-001 +// The annex half of AR-026: promoted rows are appended to the resident matrix +// and scored by the same GEMM as the baked references. What used to be a +// host-side cosine loop over TrackGallery::annex() is now these extra columns, +// so the equivalence that matters is "an appended row scores exactly what the +// reference dot product says", and "appending changes nothing about the rows +// already there". +TEST_CASE("appended rows are scored by the same GEMM as the baked gallery", + "[similarity][AR-026]") { + std::vector gallery; + for (int slot : {0, 1}) { + auto e = one_hot(slot); + gallery.insert(gallery.end(), e.begin(), e.end()); + } + auto engine = make_similarity_engine(gallery.data(), /*n_gallery=*/2, /*max_faces=*/2); + REQUIRE(engine->n_gallery() == 2); + + // One query face at 45° between slots 1 and 2. Slot 2 is not in the baked + // gallery yet, so the face is currently "unrecognised at that pose". + const float s = std::sqrt(0.5f); + std::vector query(static_cast(2) * 512, 0.0f); + query[1] = s; + query[2] = s; + + const float* before = engine->compute(query.data(), 1); + CHECK_THAT(before[0], WithinAbs(0.0f, 1e-6f)); // vs slot 0 + CHECK_THAT(before[1], WithinAbs(s, 1e-6f)); // vs slot 1 + + // Promote the missing view — the annex row an owned track would contribute. + auto promoted = one_hot(2); + engine->append_rows(promoted.data(), 1); + REQUIRE(engine->n_gallery() == 3); + + const float* after = engine->compute(query.data(), 1); + CHECK_THAT(after[0], WithinAbs(0.0f, 1e-6f)); // baked rows unchanged + CHECK_THAT(after[1], WithinAbs(s, 1e-6f)); + CHECK_THAT(after[2], WithinAbs(s, 1e-6f)); // appended row, same multiply +} + +TEST_CASE("appending many rows keeps every similarity exact", "[similarity][AR-026]") { + // Start from a one-row gallery and append past the initial capacity several + // times over — the growth path has to preserve what is already resident, and + // a film promotes far more rows than the gallery starts with. + auto seed = one_hot(0); + auto engine = make_similarity_engine(seed.data(), /*n_gallery=*/1, /*max_faces=*/1); + + constexpr int kAppended = 40; + for (int i = 1; i <= kAppended; ++i) { + auto e = one_hot(i); + engine->append_rows(e.data(), 1); + } + REQUIRE(engine->n_gallery() == kAppended + 1); + + // Query one-hot slot k: similarity is 1 against row k and 0 against all others. + for (int k : {0, 1, 17, kAppended}) { + auto q = one_hot(k); + const float* S = engine->compute(q.data(), 1); + for (int g = 0; g <= kAppended; ++g) + CHECK_THAT(S[g], WithinAbs(g == k ? 1.0f : 0.0f, 1e-6f)); + } +} + +TEST_CASE("appending a block of rows matches appending them one at a time", + "[similarity][AR-026]") { + // A promotion hands over a whole diversity buffer at once; that must be + // indistinguishable from the same rows arriving singly. + auto seed = one_hot(0); + + std::vector block; + for (int slot : {1, 2, 3}) { + auto e = one_hot(slot); + block.insert(block.end(), e.begin(), e.end()); + } + + auto bulk = make_similarity_engine(seed.data(), 1, 1); + bulk->append_rows(block.data(), 3); + + auto singly = make_similarity_engine(seed.data(), 1, 1); + for (int i = 0; i < 3; ++i) singly->append_rows(block.data() + i * 512, 1); + + REQUIRE(bulk->n_gallery() == singly->n_gallery()); + + const float t = std::sqrt(1.0f / 3.0f); + std::array q{}; + q[1] = t; q[2] = t; q[3] = t; + + const float* a = bulk->compute(q.data(), 1); + std::vector a_copy(a, a + bulk->n_gallery()); + const float* b = singly->compute(q.data(), 1); + + for (int g = 0; g < bulk->n_gallery(); ++g) + CHECK_THAT(a_copy[g], WithinAbs(b[g], 1e-6f)); +} + +TEST_CASE("appending zero rows is a no-op", "[similarity][AR-026]") { + auto e = one_hot(0); + auto engine = make_similarity_engine(e.data(), 1, 1); + CHECK_NOTHROW(engine->append_rows(nullptr, 0)); + CHECK(engine->n_gallery() == 1); +} + TEST_CASE("CPU similarity engine rejects too many faces", "[similarity]") { auto e = one_hot(0); auto engine = make_similarity_engine(e.data(), /*n_gallery=*/1, /*max_faces=*/1); diff --git a/tests/test_track_gallery.cpp b/tests/test_track_gallery.cpp index 46422b1..7c47ae1 100644 --- a/tests/test_track_gallery.cpp +++ b/tests/test_track_gallery.cpp @@ -1,4 +1,4 @@ -// TRACES: AR-018, AR-019, AR-024 | SR-005 +// TRACES: UT-005 | AR-018, AR-019, AR-024, AR-026 | SR-005, SR-001 // // Unit tests for TrackGallery (gallery/track_gallery.hpp): per-film gallery // expansion driven by track continuity. Pure, GPU-free, model-free — exercises @@ -18,8 +18,10 @@ #include "gallery/track_gallery.hpp" #include "types.hpp" +#include #include #include +#include namespace { @@ -43,6 +45,17 @@ Embedding one_hot(int slot) { return e; } +// TRACES: AR-026 | SR-001 +// The annex is a contiguous row-major matrix, not a vector of structs, so that +// the matcher can hand whole blocks of new rows to the GEMM path. Tests that +// want to compare one promoted view read it back through this. +Embedding annex_view(const TrackGallery& tg, int row) { + const float* p = tg.annex_row(row); + Embedding e{}; + std::copy(p, p + 512, e.begin()); + return e; +} + // Unit-norm embedding in the plane of axes i,j at cosine `cos_t` from axis i. // Cosine sim to one_hot(i) is exactly cos_t. Embedding at_sim(int i, int j, float cos_t) { @@ -95,7 +108,7 @@ TEST_CASE("disabled: no annex growth when expand_gallery is off", "[track_galler REQUIRE_FALSE(tg.enabled()); for (int f = 0; f < 10; ++f) tg.observe(1, one_hot(1), /*actor*/ 0, /*sim*/ 0.2f, /*accept*/ true, kNoCrop); - CHECK(tg.annex().empty()); + CHECK(tg.annex_size() == 0); } // ── AR-018: the band ───────────────────────────────────────────────────────── @@ -199,9 +212,9 @@ TEST_CASE("a two-person track never poisons the annex", "[track_gallery][AR-018] tg.observe(3, one_hot(400), 0, 0.30f, true, kNoCrop); // orthogonal outlier CHECK(tg.band_rejected() == 1); // refused at the door - REQUIRE_FALSE(tg.annex().empty()); // the legitimate views still promote - for (const auto& e : tg.annex()) - CHECK(cosine_similarity(e.emb, one_hot(400)) < 0.5f); + REQUIRE(tg.annex_size() > 0); // the legitimate views still promote + for (int i = 0; i < tg.annex_size(); ++i) + CHECK(cosine_similarity(annex_view(tg, i), one_hot(400)) < 0.5f); } TEST_CASE("a track that drifts through the band is refused at promotion", @@ -219,7 +232,7 @@ TEST_CASE("a track that drifts through the band is refused at promotion", CHECK(tg.band_rejected() == 0); // every step passed the door... // ...but 0° and 50° are P=0.643 apart, below the floor: the whole track goes. - CHECK(tg.annex().empty()); + CHECK(tg.annex_size() == 0); } // ── AR-019: ownership and promotion ────────────────────────────────────────── @@ -235,8 +248,8 @@ TEST_CASE("confirmed track promotes its store", "[track_gallery][AR-019]") { tg.observe(7, spoke(k, kSpokeCos), 0, 0.30f, true, kNoCrop); // 3 accepted frames == min_anchor_frames → confirmed and promoted. - CHECK(tg.annex().size() == 3); - for (const auto& ae : tg.annex()) CHECK(ae.actor_idx == 0); + CHECK(tg.annex_size() == 3); + for (int actor : tg.annex_actors()) CHECK(actor == 0); } TEST_CASE("registry ownership overrides the local tally", "[track_gallery][AR-019]") { @@ -247,8 +260,8 @@ TEST_CASE("registry ownership overrides the local tally", "[track_gallery][AR-01 tg.set_owner(11, 9); for (int k = 1; k <= 3; ++k) tg.observe(11, spoke(k, kSpokeCos), 5, 0.30f, true, kNoCrop); - REQUIRE_FALSE(tg.annex().empty()); - for (const auto& ae : tg.annex()) CHECK(ae.actor_idx == 9); + REQUIRE(tg.annex_size() > 0); + for (int actor : tg.annex_actors()) CHECK(actor == 9); } TEST_CASE("unconfirmed track (too few accepts) does not promote", "[track_gallery][AR-019]") { @@ -259,7 +272,7 @@ TEST_CASE("unconfirmed track (too few accepts) does not promote", "[track_galler tg.observe(4, spoke(1, kSpokeCos), 0, 0.30f, true, kNoCrop); tg.observe(4, spoke(2, kSpokeCos), 0, 0.30f, true, kNoCrop); tg.observe(4, spoke(3, kSpokeCos), 0, 0.30f, false, kNoCrop); - CHECK(tg.annex().empty()); + CHECK(tg.annex_size() == 0); } TEST_CASE("plurality actor wins a mixed-vote track", "[track_gallery][AR-019]") { @@ -270,8 +283,8 @@ TEST_CASE("plurality actor wins a mixed-vote track", "[track_gallery][AR-019]") tg.observe(8, spoke(1, kSpokeCos), 5, 0.30f, true, kNoCrop); tg.observe(8, spoke(2, kSpokeCos), 5, 0.30f, true, kNoCrop); tg.observe(8, spoke(3, kSpokeCos), 6, 0.30f, true, kNoCrop); - REQUIRE_FALSE(tg.annex().empty()); - for (const auto& ae : tg.annex()) CHECK(ae.actor_idx == 5); + REQUIRE(tg.annex_size() > 0); + for (int actor : tg.annex_actors()) CHECK(actor == 5); } TEST_CASE("promotion is idempotent across a long track", "[track_gallery][AR-019]") { @@ -279,12 +292,12 @@ TEST_CASE("promotion is idempotent across a long track", "[track_gallery][AR-019 tg.set_calibration(identity_cal); for (int k = 1; k <= 3; ++k) tg.observe(9, spoke(k, kSpokeCos), 0, 0.30f, true, kNoCrop); - size_t after_confirm = tg.annex().size(); + const int after_confirm = tg.annex_size(); REQUIRE(after_confirm > 0); // Keep feeding the confirmed track: annex must not grow again. for (int f = 0; f < 10; ++f) tg.observe(9, spoke(4 + f, kSpokeCos), 0, 0.30f, true, kNoCrop); - CHECK(tg.annex().size() == after_confirm); + CHECK(tg.annex_size() == after_confirm); } TEST_CASE("clear_tracks drops buffers before confirmation", "[track_gallery][AR-019]") { @@ -296,7 +309,7 @@ TEST_CASE("clear_tracks drops buffers before confirmation", "[track_gallery][AR- tg.observe(1, spoke(2, kSpokeCos), 0, 0.30f, true, kNoCrop); tg.clear_tracks(); tg.observe(1, spoke(3, kSpokeCos), 0, 0.30f, true, kNoCrop); - CHECK(tg.annex().empty()); + CHECK(tg.annex_size() == 0); } TEST_CASE("eviction keeps the gallery-far views", "[track_gallery][AR-018]") { @@ -315,11 +328,56 @@ TEST_CASE("eviction keeps the gallery-far views", "[track_gallery][AR-018]") { tg.observe(6, spoke(3, kSpokeCos), 0, 0.20f, true, kNoCrop); // novel: evicts the 0.80 tg.observe(6, spoke(4, kSpokeCos), 0, 0.90f, true, kNoCrop); // least novel: dropped - REQUIRE(tg.annex().size() == 2); + REQUIRE(tg.annex_size() == 2); // Survivors are the two most gallery-far views: spokes 2 and 3. - for (const auto& ae : tg.annex()) { - const bool is_2 = cosine_similarity(ae.emb, spoke(2, kSpokeCos)) > 0.99f; - const bool is_3 = cosine_similarity(ae.emb, spoke(3, kSpokeCos)) > 0.99f; + for (int i = 0; i < tg.annex_size(); ++i) { + const Embedding view = annex_view(tg, i); + const bool is_2 = cosine_similarity(view, spoke(2, kSpokeCos)) > 0.99f; + const bool is_3 = cosine_similarity(view, spoke(3, kSpokeCos)) > 0.99f; CHECK((is_2 || is_3)); } } + +// TRACES: UT-005 | AR-026 | SR-001 +// The annex reaches the GEMM path by being drained, not re-read: the matcher +// pushes newly promoted rows into the similarity engine once per frame. Draining +// must therefore be exactly-once — a row handed over twice becomes a duplicate +// gallery entry that quietly doubles an actor's best-of-N chances, and a row +// never handed over is a promotion that silently does nothing. +TEST_CASE("promotions drain exactly once, in matrix order", + "[track_gallery][AR-026]") { + TrackGallery tg(expand_cfg()); + tg.set_calibration(identity_cal); + + std::vector emb; + std::vector actor; + + CHECK(tg.drain_promotions(emb, actor) == 0); // nothing promoted yet + + for (int k = 1; k <= 3; ++k) + tg.observe(7, spoke(k, kSpokeCos), 0, 0.30f, true, kNoCrop); + REQUIRE(tg.annex_size() == 3); + + const int drained = tg.drain_promotions(emb, actor); + CHECK(drained == 3); + CHECK(actor.size() == 3); + CHECK(emb.size() == 3 * 512); + for (int a : actor) CHECK(a == 0); + + // Drained rows are the annex rows, in the same order — the engine's row i + // and flat_actor_[i] have to keep naming the same view. + for (int i = 0; i < drained; ++i) + for (int d = 0; d < 512; ++d) + CHECK(emb[static_cast(i) * 512 + d] == tg.annex_row(i)[d]); + + // Draining again yields nothing: the engine already holds these. + CHECK(tg.drain_promotions(emb, actor) == 0); + CHECK(actor.size() == 3); + + // A second track promotes, and only its rows are handed over. + for (int k = 1; k <= 3; ++k) + tg.observe(8, spoke(k, kSpokeCos), 4, 0.30f, true, kNoCrop); + CHECK(tg.drain_promotions(emb, actor) == 3); + CHECK(actor.size() == 6); + CHECK(actor[5] == 4); +}