dtourolle/scene-actor-extraction
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(gemm): the annex is a matrix, not a list — scored by the same GEMM

Browse Source 
The per-film annex was folded in after the gallery multiply by a host-side
cosine loop over a vector of {embedding, actor} structs, justified in-comment
by "tens of embeddings". AR-018/AR-019 retired that assumption: every owned
track promotes, so the annex grows with cast size and film length.

TrackGallery now holds it as a contiguous row-major matrix with a parallel
actor index — the flat_emb_/flat_actor_ shape the baked gallery already uses —
and hands newly promoted rows to the matcher once per frame. The matcher pushes
them into the similarity engine's resident matrix through a new
ISimilarityEngine::append_rows, so one SGEMM covers baked and promoted
references alike and best-of-N is a single pass over one similarity column.
Capacity doubles on overflow, and the GPU backends grow device-to-device, so a
promotion never re-uploads the gallery across the bus.

Absorbing promotions 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 also removes an incidental 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 documented.

OpenBLAS becomes a requirement of the CPU GEMM backend rather than an
opportunistic upgrade. That path is what CI and the cpu builder image run, so
falling back to the scalar loop in silence meant AR-027 could be measured — or
believed — on a kernel no release uses. The loop survives as the correctness
oracle the BLAS backends are diffed against, behind SAE_ALLOW_SCALAR_GEMM.

Call site 3, the deferred TBI pass, is untouched: it does not exist until
AR-020, so AR-026 stays In Progress.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

TRACES: AR-026 | UT-004, UT-005 | SR-001
This commit is contained in:
dtourolle
2026-08-04 21:20:31 +02:00
co-authored by Claude Opus 5
parent f33403fff8
commit c1155cb607
11 changed files with 461 additions and 81 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/backends/gemm_backend.cpp
+79 -14
View File
@@ -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<size_t>(n_rows) * kDim);
n_gallery_ += n_rows;
host_sims_.resize(static_cast<size_t>(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<size_t>(n_gallery_) * kDim;
gpu_malloc(reinterpret_cast<void**>(&d_gallery_), gallery_floats * sizeof(float));
gpu_memcpy_h2d_sync(d_gallery_, gallery_row_major, gallery_floats * sizeof(float));
gpu_malloc(reinterpret_cast<void**>(&d_query_),
static_cast<size_t>(max_faces_) * kDim * sizeof(float));
gpu_malloc(reinterpret_cast<void**>(&d_sims_),
static_cast<size_t>(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<size_t>(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<size_t>(n_gallery_) * kDim,
rows_row_major,
static_cast<size_t>(n_rows) * kDim * sizeof(float));
n_gallery_ = want;
host_sims_.resize(static_cast<size_t>(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<void**>(&d_new_gallery),
static_cast<size_t>(rows) * kDim * sizeof(float));
if (d_gallery_ && n_gallery_ > 0)
gpu_memcpy_d2d_sync(d_new_gallery, d_gallery_,
static_cast<size_t>(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<void**>(&d_new_sims),
static_cast<size_t>(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};