From 39938a677c2f48ed0da9696eaa03250e86533d0c Mon Sep 17 00:00:00 2001 From: njayj <124410198+njayj@users.noreply.github.com> Date: Mon, 22 Jun 2026 22:49:20 -0400 Subject: [PATCH 01/16] feat(build): SAM3_CUDA ggml-CUDA backend option (NVIDIA fast-path) --- CMakeLists.txt | 25 ++++++++++++++++++++----- 1 file changed, 20 insertions(+), 5 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index dc0029e..62c822c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,6 +26,20 @@ if(NOT APPLE AND SAM3_VULKAN) set(GGML_VULKAN ON CACHE BOOL "" FORCE) endif() +# RFD 0011 (NVIDIA fast-path): ggml-CUDA runs the same EdgeTAM graph on NVIDIA GPUs +# with tensor-core FP16 kernels + CUDA graphs + FlashAttention. It is the NVIDIA +# sibling of the Vulkan (cross-vendor) and Metal/CoreML (Apple) backends. The cgo +# C-ABI (coreml/edgetam_capi.cpp) is reused below — its CoreML stage calls are +# #ifdef SAM3_COREML-guarded, so a CUDA build runs the pure-ggml path on the CUDA +# device. Setting GGML_CUDA=ON makes ggml define GGML_USE_CUDA PUBLIC (ggml_add_backend +# CUDA), which the model loader's #ifdef GGML_USE_CUDA branch keys on. Requires the +# CUDA Toolkit (nvcc) at build. SAM3_CUDA and SAM3_VULKAN may both be ON (fat build): +# the loader prefers CUDA then falls back to Vulkan then CPU. +option(SAM3_CUDA "Enable the ggml-CUDA backend (NVIDIA GPUs)" OFF) +if(NOT APPLE AND SAM3_CUDA) + set(GGML_CUDA ON CACHE BOOL "" FORCE) +endif() + add_subdirectory(ggml) # sam3 static library @@ -75,11 +89,12 @@ if(APPLE) endif() endif() -# RFD 0011 (Windows/Vulkan): reuse the cgo tracker C-ABI on the non-Apple Vulkan -# path WITHOUT the CoreML bridge. edgetam_capi.cpp's CoreML stage calls + producer- -# encoder pool are #ifdef SAM3_COREML-guarded, so here the tracker runs the pure- -# ggml path on the Vulkan device. SAM3_COREML is NOT defined, no .mm, no frameworks. -if(NOT APPLE AND SAM3_VULKAN) +# RFD 0011 (Windows/Vulkan + NVIDIA/CUDA): reuse the cgo tracker C-ABI on the +# non-Apple GPU paths WITHOUT the CoreML bridge. edgetam_capi.cpp's CoreML stage +# calls + producer-encoder pool are #ifdef SAM3_COREML-guarded, so here the tracker +# runs the pure-ggml path on the Vulkan or CUDA device. SAM3_COREML is NOT defined, +# no .mm, no frameworks. +if(NOT APPLE AND (SAM3_VULKAN OR SAM3_CUDA)) target_sources(sam3 PRIVATE coreml/edgetam_capi.cpp) target_include_directories(sam3 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/coreml) endif() From fa06b8519be35b21e0156ed085bfbebba38fc56f Mon Sep 17 00:00:00 2001 From: njayj <124410198+njayj@users.noreply.github.com> Date: Mon, 22 Jun 2026 22:50:40 -0400 Subject: [PATCH 02/16] feat(cuda): select the ggml-CUDA backend in the model loader (CUDA->Vulkan->CPU) --- sam3.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/sam3.cpp b/sam3.cpp index 8947d16..38d4c92 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -18,6 +18,13 @@ #include "ggml-vulkan.h" #endif +// RFD 0011 (NVIDIA fast-path): ggml-CUDA backend for NVIDIA GPUs (SAM3_CUDA → +// GGML_CUDA → GGML_USE_CUDA). The model loader selects it below, preferred over +// Vulkan when both are compiled (a fat build). +#ifdef GGML_USE_CUDA +#include "ggml-cuda.h" +#endif + /* stb (implementation compiled here -- order is pinned) */ #define STB_IMAGE_IMPLEMENTATION #include "stb_image.h" @@ -3431,6 +3438,18 @@ std::shared_ptr sam3_load_model(const sam3_params& params) { model->backend = ggml_backend_metal_init(); } #endif +#ifdef GGML_USE_CUDA + // NVIDIA fast-path. ggml-CUDA runs the EdgeTAM graph with tensor-core FP16 + // kernels, CUDA graphs, and FlashAttention. Preferred over Vulkan when both are + // compiled in (a fat NVIDIA build). device 0 = first CUDA GPU. Returns NULL if + // there is no CUDA device/driver (or a CUDA 12/13 lib mismatch) — the chain then + // falls through to Vulkan, then CPU. We log the selection so a silent CPU + // fallback is visible (cf. Egor part-2: "always check which provider ran"). + if (params.use_gpu && !model->backend) { + fprintf(stderr, "%s: using CUDA backend\n", __func__); + model->backend = ggml_backend_cuda_init(0); + } +#endif #ifdef GGML_USE_VULKAN // RFD 0011: cross-vendor GPU. ggml-Vulkan runs the same EdgeTAM graph on any // Vulkan device (AMD/NVIDIA/Intel; Apple via MoltenVK). device 0 = first GPU. From 127e61b518e31bdab1331e10acc60ab0485009f9 Mon Sep 17 00:00:00 2001 From: njayj <124410198+njayj@users.noreply.github.com> Date: Mon, 22 Jun 2026 22:51:57 -0400 Subject: [PATCH 03/16] feat(capi): expose active ggml backend name (anti silent-CPU-fallback) --- coreml/edgetam_capi.cpp | 7 +++++++ coreml/edgetam_capi.h | 5 +++++ sam3.cpp | 7 +++++++ sam3.h | 5 +++++ 4 files changed, 24 insertions(+) diff --git a/coreml/edgetam_capi.cpp b/coreml/edgetam_capi.cpp index 1c99dcd..9ca18a1 100644 --- a/coreml/edgetam_capi.cpp +++ b/coreml/edgetam_capi.cpp @@ -245,3 +245,10 @@ extern "C" et_result edgetam_capi_track_slot(edgetam_tracker_t h, int slot, extern "C" const char* edgetam_capi_version(void) { return SAM3_VERSION; } + +extern "C" const char* edgetam_capi_backend(edgetam_tracker_t h) { + auto* t = static_cast(h); + // *t->model matches the existing capi pattern (sam3_encode_image(*t->state, + // *t->model, ...)); sam3_backend_name reads the model's active ggml backend. + return (t && t->model) ? sam3_backend_name(*t->model) : "none"; +} diff --git a/coreml/edgetam_capi.h b/coreml/edgetam_capi.h index 4849ecd..9bf9957 100644 --- a/coreml/edgetam_capi.h +++ b/coreml/edgetam_capi.h @@ -65,6 +65,11 @@ void edgetam_capi_reset(edgetam_tracker_t t); void edgetam_capi_destroy(edgetam_tracker_t t); const char* edgetam_capi_version(void); // SAM3_VERSION, for cgo link sanity +// Active ggml backend name (e.g. "CUDA0", "Vulkan0", "Metal", "CPU", or "none"). +// Lets the Go side assert the GPU backend really initialized — catching a silent +// CPU fallback that would be "fast and wrong" (cf. Egor part-2 CUDA 12/13 mismatch). +const char* edgetam_capi_backend(edgetam_tracker_t t); + #ifdef __cplusplus } #endif diff --git a/sam3.cpp b/sam3.cpp index 38d4c92..1690a49 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -3538,6 +3538,13 @@ void sam3_free_model(sam3_model& model) { } } +const char* sam3_backend_name(const sam3_model& model) { + // ggml_backend_name returns a static string owned by ggml; safe to hand across + // the C-ABI without copying. model.backend is set by sam3_load_model's + // preference chain (CUDA/Vulkan/Metal/CPU). + return model.backend ? ggml_backend_name(model.backend) : "none"; +} + bool sam3_is_visual_only(const sam3_model& model) { return model.hparams.visual_only != 0 || model.hparams.is_sam2() || model.hparams.is_edgetam(); } diff --git a/sam3.h b/sam3.h index 809d4de..7c624c6 100644 --- a/sam3.h +++ b/sam3.h @@ -222,6 +222,11 @@ std::shared_ptr sam3_load_model(const sam3_params & params); /* Free all resources held by a loaded model. */ void sam3_free_model(sam3_model & model); +// Active ggml backend name for the loaded model (e.g. "CUDA0", "Vulkan0", "Metal", +// "CPU"). Lets callers verify the GPU accelerator actually initialized rather than +// silently falling back to CPU. Returns "none" if no backend is set. +const char* sam3_backend_name(const sam3_model& model); + /* Returns true if the model was loaded as visual-only (no text/detector path). ** SAM2 models are always considered visual-only. */ bool sam3_is_visual_only(const sam3_model & model); From 260454437c188d59cc4e3c298b9f62516d09ab91 Mon Sep 17 00:00:00 2001 From: njayj <124410198+njayj@users.noreply.github.com> Date: Sat, 4 Jul 2026 10:11:11 +0200 Subject: [PATCH 04/16] build(capi): SAM3_CAPI_SHARED C-ABI DLL target + force CUDA graphs on SAM3_CUDA MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - SAM3_CUDA now forces GGML_CUDA_GRAPHS=ON: standalone ggml defaults it OFF (llama.cpp-only flag), so the previously advertised 'CUDA graphs' path was never compiled. ggml auto-falls back at runtime for uncapturable graphs. - New SAM3_CAPI_SHARED option builds the cgo C-ABI as its own shared library (sam3capi.dll). This is the Windows/CUDA toolchain bridge: nvcc requires MSVC as host on Windows, and MSVC-built C++ static libs cannot link into a mingw cgo binary — a C-ABI DLL is toolchain-neutral. ET_API in edgetam_capi.h carries the dllexport (extern "C" alone does not export from an MSVC DLL). Co-Authored-By: Claude Fable 5 --- CMakeLists.txt | 37 ++++++++++++++++++++++++++++++++----- coreml/edgetam_capi.h | 33 ++++++++++++++++++++++----------- 2 files changed, 54 insertions(+), 16 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 62c822c..a9b5a6b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -38,6 +38,13 @@ endif() option(SAM3_CUDA "Enable the ggml-CUDA backend (NVIDIA GPUs)" OFF) if(NOT APPLE AND SAM3_CUDA) set(GGML_CUDA ON CACHE BOOL "" FORCE) + # CUDA graphs: standalone ggml defaults GGML_CUDA_GRAPHS to OFF (llama.cpp + # flips it on in its own build). Without it, every frame pays per-kernel + # launch overhead across the EdgeTAM graph's hundreds of small kernels. + # Force it ON so the advertised "tensor-core FP16 + CUDA graphs" path is + # actually compiled; ggml auto-falls back to normal kernel launches at + # runtime for graphs it cannot capture, so this is safe. + set(GGML_CUDA_GRAPHS ON CACHE BOOL "" FORCE) endif() add_subdirectory(ggml) @@ -91,12 +98,32 @@ endif() # RFD 0011 (Windows/Vulkan + NVIDIA/CUDA): reuse the cgo tracker C-ABI on the # non-Apple GPU paths WITHOUT the CoreML bridge. edgetam_capi.cpp's CoreML stage -# calls + producer-encoder pool are #ifdef SAM3_COREML-guarded, so here the tracker -# runs the pure-ggml path on the Vulkan or CUDA device. SAM3_COREML is NOT defined, -# no .mm, no frameworks. +# calls are #ifdef SAM3_COREML-guarded, so here the tracker runs the pure-ggml +# path on the Vulkan or CUDA device (with the opt-in ggml encoder-ahead producer +# on CUDA). SAM3_COREML is NOT defined, no .mm, no frameworks. +# +# SAM3_CAPI_SHARED: build the C-ABI as its own SHARED library (sam3capi.dll / +# libsam3capi.so) instead of folding it into the static libsam3. This is the +# Windows/CUDA cgo boundary: nvcc requires MSVC as its host compiler on Windows, +# and an MSVC-built C++ static lib cannot link into a mingw cgo binary (C++ ABI / +# runtime mismatch). A C-ABI DLL is toolchain-neutral — all C++ (sam3 + ggml + +# ggml-cuda + ggml-vulkan) stays inside the MSVC-built DLL and Go links only the +# 15 extern "C" entry points. Keep BUILD_SHARED_LIBS=OFF with this so ggml links +# statically INTO the DLL (one artifact + the CUDA runtime DLLs). +option(SAM3_CAPI_SHARED "Build the cgo C-ABI tracker as a shared library (Windows/MSVC cgo boundary)" OFF) if(NOT APPLE AND (SAM3_VULKAN OR SAM3_CUDA)) - target_sources(sam3 PRIVATE coreml/edgetam_capi.cpp) - target_include_directories(sam3 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/coreml) + if(SAM3_CAPI_SHARED) + add_library(sam3capi SHARED coreml/edgetam_capi.cpp) + target_include_directories(sam3capi PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/coreml) + # ET_CAPI_BUILD_DLL turns on __declspec(dllexport) for the capi entry + # points (see ET_API in edgetam_capi.h) — extern "C" alone does not + # export symbols from an MSVC DLL. + target_compile_definitions(sam3capi PRIVATE ET_CAPI_BUILD_DLL) + target_link_libraries(sam3capi PRIVATE sam3) + else() + target_sources(sam3 PRIVATE coreml/edgetam_capi.cpp) + target_include_directories(sam3 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/coreml) + endif() endif() # Examples (only when top-level) diff --git a/coreml/edgetam_capi.h b/coreml/edgetam_capi.h index 9bf9957..2ca1113 100644 --- a/coreml/edgetam_capi.h +++ b/coreml/edgetam_capi.h @@ -9,6 +9,17 @@ #include +// ET_API: symbol visibility for the SAM3_CAPI_SHARED build (sam3capi.dll — the +// Windows/MSVC cgo boundary; see CMakeLists.txt). extern "C" alone does not +// export symbols from an MSVC DLL, so the DLL build defines ET_CAPI_BUILD_DLL +// and every entry point below carries __declspec(dllexport). Static builds +// (Mac CoreML, Linux CUDA, mingw Vulkan) expand ET_API to nothing. +#if defined(_WIN32) && defined(ET_CAPI_BUILD_DLL) + #define ET_API __declspec(dllexport) +#else + #define ET_API +#endif + #ifdef __cplusplus extern "C" { #endif @@ -33,23 +44,23 @@ typedef struct { // path + ggml fallback). models_dir = dir holding the 4 CoreML .mlpackages; if // non-empty, the CoreML encoder/mem-attn/decoder stages are enabled. Returns // NULL on failure. -edgetam_tracker_t edgetam_capi_create(const char* models_dir, const char* ggml_model, +ET_API edgetam_tracker_t edgetam_capi_create(const char* models_dir, const char* ggml_model, int use_samurai, int use_lifecycle); // Seed the target from a box on frame 0. rgb = contiguous RGB24 (w*h*3 bytes). // seed box in PIXELS. Returns 1 on success. -int edgetam_capi_seed(edgetam_tracker_t t, const uint8_t* rgb, int w, int h, et_box seed_px); +ET_API int edgetam_capi_seed(edgetam_tracker_t t, const uint8_t* rgb, int w, int h, et_box seed_px); // Track one frame. rgb = contiguous RGB24 (w*h*3). Returns the per-frame result // (box normalized [0,1]); valid==0 when the target is lost this frame. -et_result edgetam_capi_track(edgetam_tracker_t t, const uint8_t* rgb, int w, int h); +ET_API et_result edgetam_capi_track(edgetam_tracker_t t, const uint8_t* rgb, int w, int h); // RFD 0011 U8 mask export — copy the LAST track's binary mask (0/255, row-major, // w*h bytes, original-frame resolution) into `out` (capacity `cap`). Writes dims // to *w,*h. Returns the byte count (w*h); if cap < w*h nothing is copied and w*h // is returned so the caller can size its buffer. 0 => no mask (lost / pre-track). // Valid after edgetam_capi_track / edgetam_capi_track_slot on the same handle. -int edgetam_capi_last_mask(edgetam_tracker_t t, uint8_t* out, int cap, int* w, int* h); +ET_API int edgetam_capi_last_mask(edgetam_tracker_t t, uint8_t* out, int cap, int* w, int* h); // RFD 0011 U8 Wave 5 — encoder-ahead threading split. pool_size returns the slot // count (0 if CoreML/threading is unavailable). encode_slot (PRODUCER thread) @@ -57,18 +68,18 @@ int edgetam_capi_last_mask(edgetam_tracker_t t, uint8_t* out, int cap, int* w, i // handle. track_slot (CONSUMER thread) consumes slot's prefetched neck + propagates. // Go owns slot free/ready coordination (channels) — one writer then one reader per // slot, so these are safe to overlap across the two threads on DIFFERENT slots. -int edgetam_capi_pool_size(edgetam_tracker_t t); -int edgetam_capi_encode_slot(edgetam_tracker_t t, int slot, const uint8_t* rgb, int w, int h); -et_result edgetam_capi_track_slot(edgetam_tracker_t t, int slot, const uint8_t* rgb, int w, int h); +ET_API int edgetam_capi_pool_size(edgetam_tracker_t t); +ET_API int edgetam_capi_encode_slot(edgetam_tracker_t t, int slot, const uint8_t* rgb, int w, int h); +ET_API et_result edgetam_capi_track_slot(edgetam_tracker_t t, int slot, const uint8_t* rgb, int w, int h); -void edgetam_capi_reset(edgetam_tracker_t t); -void edgetam_capi_destroy(edgetam_tracker_t t); -const char* edgetam_capi_version(void); // SAM3_VERSION, for cgo link sanity +ET_API void edgetam_capi_reset(edgetam_tracker_t t); +ET_API void edgetam_capi_destroy(edgetam_tracker_t t); +ET_API const char* edgetam_capi_version(void); // SAM3_VERSION, for cgo link sanity // Active ggml backend name (e.g. "CUDA0", "Vulkan0", "Metal", "CPU", or "none"). // Lets the Go side assert the GPU backend really initialized — catching a silent // CPU fallback that would be "fast and wrong" (cf. Egor part-2 CUDA 12/13 mismatch). -const char* edgetam_capi_backend(edgetam_tracker_t t); +ET_API const char* edgetam_capi_backend(edgetam_tracker_t t); #ifdef __cplusplus } From a3194c7a095cb5e1ef22b6b40bc64d2fd136e7f6 Mon Sep 17 00:00:00 2001 From: njayj <124410198+njayj@users.noreply.github.com> Date: Sat, 4 Jul 2026 10:11:11 +0200 Subject: [PATCH 05/16] feat(cuda): guarded runtime fallback chain + ggml encoder-ahead producer MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Loader (fix): the CUDA/Vulkan selection logged 'using X backend' BEFORE init and unconditionally — a machine without the device logged the wrong backend. Now: device-count guard, log AFTER successful init, explicit fallback warning. ggml-Vulkan THROWS (vk::SystemError) when the Vulkan loader/driver is absent — the probe+init are now wrapped so a GPU-less machine falls through to the CPU backend instead of failing the whole model load (fat-build CUDA->Vulkan->CPU chain is now safe end to end). Encoder-ahead (feat): the CoreML-threading equivalent for NVIDIA. The prefetched-neck seam moves out of #ifdef SAM3_COREML (it is raw floats — no CoreML dependency); a factored edgetam_load_neck_from_floats() is shared by the CoreML helper and the new ggml consume path in edgetam_encode_image. sam3_encoder_ahead_{create,encode,destroy} run the EdgeTAM RepViT+FPN encoder graph on a SECOND ggml-CUDA backend instance (own stream; shared read-only weights) so frame N+1's encode overlaps frame N's mem-attn/decoder. The capi wires it behind SAM3_GGML_ENCODER_AHEAD=1 (default OFF: same-GPU producer/consumer contention must be A/B-measured on real NVIDIA hardware — the M4 sweep showed co-location can lose; pool_size()==0 keeps the synchronous path otherwise). Verified: -fsyntax-only across {plain, SAM3_COREML, GGML_USE_CUDA, GGML_USE_VULKAN, CUDA+VULKAN, ET_CAPI_BUILD_DLL}; full Mac cmake build (SAM3_COREML=ON) + sam3_edgetam_bench link green. Co-Authored-By: Claude Fable 5 --- coreml/edgetam_capi.cpp | 72 +++++++--- sam3.cpp | 290 ++++++++++++++++++++++++++++++---------- sam3.h | 32 ++++- 3 files changed, 300 insertions(+), 94 deletions(-) diff --git a/coreml/edgetam_capi.cpp b/coreml/edgetam_capi.cpp index 9ca18a1..a4ecc64 100644 --- a/coreml/edgetam_capi.cpp +++ b/coreml/edgetam_capi.cpp @@ -20,12 +20,18 @@ struct EtTracker { sam3_tracker_ptr tracker; // RFD 0011 U8 Wave 5 (encoder-ahead threading): a producer encoder handle - // (its OWN MLModel, per-thread ownership) + a fixed neck pool. Go owns slot - // free/ready coordination via channels, so NO lock here — each slot is written - // once by encode_slot (producer thread) then read once by track_slot (consumer - // thread). track_slot is the ONLY caller of sam3_coreml_set_prefetched_neck - // (global, consumer-thread-only). prod_enc is non-null only when CoreML is on. + // (its OWN engine instance, per-thread ownership) + a fixed neck pool. Go owns + // slot free/ready coordination via channels, so NO lock here — each slot is + // written once by encode_slot (producer thread) then read once by track_slot + // (consumer thread). track_slot is the ONLY caller of + // sam3_coreml_set_prefetched_neck (global, consumer-thread-only). + // At most ONE producer engine is ever non-null: + // prod_enc — CoreML MLModel (SAM3_COREML builds) + // ggml_ea — second ggml-CUDA backend instance (non-CoreML builds, opt-in + // via SAM3_GGML_ENCODER_AHEAD=1; the CoreML-threading + // equivalent for NVIDIA) edgetam_coreml_handle prod_enc = nullptr; + sam3_encoder_ahead* ggml_ea = nullptr; static const int POOL = 3; std::vector neck[POOL][3]; // n0 256*256*256, n1 128*128*256, n2 64*64*256 @@ -89,17 +95,33 @@ extern "C" edgetam_tracker_t edgetam_capi_create(const char* models_dir, t->tracker = sam3_create_visual_tracker(*t->model, vp); #ifdef SAM3_COREML // Wave 5: producer encoder handle + neck pool (CoreML encoder-ahead threading). - // Absent on non-CoreML builds → prod_enc stays null → pool_size()==0 → - // the cgo session uses the synchronous (single-thread) Track path. if (models_dir && *models_dir) { std::string d = models_dir; t->prod_enc = edgetam_coreml_create((d + "/edgetam_encoder_neck_nhwc.mlpackage").c_str(), /*ANE*/ 1); + } +#else + // RFD 0011 (CUDA): ggml encoder-ahead producer — the CoreML-threading + // equivalent for NVIDIA. A second ggml-CUDA backend instance encodes + // frame N+1 on its own stream while the consumer propagates frame N. + // OPT-IN (SAM3_GGML_ENCODER_AHEAD=1) and default OFF: same-GPU + // producer/consumer contention is hardware-dependent (the M4 sweep + // showed co-location can LOSE) — A/B on the target GPU first. + // sam3_encoder_ahead_create itself returns NULL unless the active + // backend is CUDA and the model is EdgeTAM, so a Vulkan/CPU fallback + // run cleanly keeps pool_size()==0 → the synchronous Track path. + const char* ea_env = getenv("SAM3_GGML_ENCODER_AHEAD"); + if (ea_env && *ea_env && strcmp(ea_env, "0") != 0) { + t->ggml_ea = sam3_encoder_ahead_create(*t->model); + } +#endif + // Either producer engine present → allocate the fixed neck pool; neither + // → pool_size()==0 → the cgo session uses the synchronous Track path. + if (t->prod_enc || t->ggml_ea) { const int Wd[3] = {256, 128, 64}; for (int s = 0; s < EtTracker::POOL; ++s) for (int i = 0; i < 3; ++i) t->neck[s][i].assign((size_t)256 * Wd[i] * Wd[i], 0.f); } -#endif return t.release(); } catch (...) { return nullptr; @@ -174,6 +196,7 @@ extern "C" void edgetam_capi_destroy(edgetam_tracker_t h) { #ifdef SAM3_COREML if (t && t->prod_enc) edgetam_coreml_destroy(t->prod_enc); #endif + if (t && t->ggml_ea) sam3_encoder_ahead_destroy(t->ggml_ea); delete t; } @@ -186,28 +209,33 @@ extern "C" void edgetam_capi_destroy(edgetam_tracker_t h) { extern "C" int edgetam_capi_pool_size(edgetam_tracker_t h) { auto* t = static_cast(h); - return (t && t->prod_enc) ? EtTracker::POOL : 0; // 0 => threading unavailable + return (t && (t->prod_enc || t->ggml_ea)) ? EtTracker::POOL : 0; // 0 => threading unavailable } extern "C" int edgetam_capi_encode_slot(edgetam_tracker_t h, int slot, const uint8_t* rgb, int w, int hgt) { auto* t = static_cast(h); - if (!t || !t->prod_enc || !rgb || slot < 0 || slot >= EtTracker::POOL) return 0; -#ifdef SAM3_COREML + if (!t || (!t->prod_enc && !t->ggml_ea) || !rgb || slot < 0 || slot >= EtTracker::POOL) return 0; try { +#ifdef SAM3_COREML sam3_image img = make_image(rgb, w, hgt); std::vector norm = sam3_coreml_preprocess_image(img, 1024); return edgetam_coreml_encode(t->prod_enc, norm.data(), t->neck[slot][0].data(), t->neck[slot][1].data(), t->neck[slot][2].data()) ? 1 : 0; +#else + // ggml encoder-ahead (CUDA): preprocess + encode on the producer's own + // backend instance; write the 3 neck levels into this slot's buffers. + sam3_image img = make_image(rgb, w, hgt); + return sam3_encoder_ahead_encode(t->ggml_ea, *t->model, img, + t->neck[slot][0].data(), + t->neck[slot][1].data(), + t->neck[slot][2].data()) ? 1 : 0; +#endif } catch (...) { return 0; } -#else - (void)rgb; (void)w; (void)hgt; // encoder-ahead threading is CoreML-only - return 0; -#endif } extern "C" et_result edgetam_capi_track_slot(edgetam_tracker_t h, int slot, @@ -217,13 +245,15 @@ extern "C" et_result edgetam_capi_track_slot(edgetam_tracker_t h, int slot, auto* t = static_cast(h); if (!t || !rgb || slot < 0 || slot >= EtTracker::POOL) return r; try { -#ifdef SAM3_COREML // Consumer thread only: hand the producer-encoded neck to the library so - // sam3_propagate_frame's encode step is skipped (no ggml/CoreML encode here). - sam3_coreml_set_prefetched_neck(t->neck[slot][0].data(), - t->neck[slot][1].data(), - t->neck[slot][2].data()); -#endif + // sam3_propagate_frame's encode step is skipped (no in-line encode here). + // The seam is backend-agnostic — the slot was filled by the CoreML + // producer OR the ggml-CUDA encoder-ahead producer. + if (t->prod_enc || t->ggml_ea) { + sam3_coreml_set_prefetched_neck(t->neck[slot][0].data(), + t->neck[slot][1].data(), + t->neck[slot][2].data()); + } sam3_image img = make_image(rgb, w, hgt); sam3_result res = sam3_propagate_frame(*t->tracker, *t->state, *t->model, img); if (res.detections.empty()) { t->last_mask = sam3_mask{}; return r; } diff --git a/sam3.cpp b/sam3.cpp index 1690a49..93e79b8 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -3441,21 +3441,44 @@ std::shared_ptr sam3_load_model(const sam3_params& params) { #ifdef GGML_USE_CUDA // NVIDIA fast-path. ggml-CUDA runs the EdgeTAM graph with tensor-core FP16 // kernels, CUDA graphs, and FlashAttention. Preferred over Vulkan when both are - // compiled in (a fat NVIDIA build). device 0 = first CUDA GPU. Returns NULL if - // there is no CUDA device/driver (or a CUDA 12/13 lib mismatch) — the chain then - // falls through to Vulkan, then CPU. We log the selection so a silent CPU - // fallback is visible (cf. Egor part-2: "always check which provider ran"). + // compiled in (a fat NVIDIA build). device 0 = first CUDA GPU. The device-count + // guard + log-AFTER-init make the runtime fallback chain (CUDA → Vulkan → CPU) + // observable and honest: a machine without an NVIDIA device/driver (or with a + // CUDA 12/13 lib mismatch) logs the fallback instead of claiming CUDA ran + // (cf. Egor part-2: "always check which provider ran"). if (params.use_gpu && !model->backend) { - fprintf(stderr, "%s: using CUDA backend\n", __func__); - model->backend = ggml_backend_cuda_init(0); + if (ggml_backend_cuda_get_device_count() > 0) { + model->backend = ggml_backend_cuda_init(0); + } + if (model->backend) { + fprintf(stderr, "%s: using CUDA backend (%s)\n", __func__, ggml_backend_name(model->backend)); + } else { + fprintf(stderr, "%s: CUDA backend unavailable (no NVIDIA device/driver) — falling back\n", __func__); + } } #endif #ifdef GGML_USE_VULKAN // RFD 0011: cross-vendor GPU. ggml-Vulkan runs the same EdgeTAM graph on any // Vulkan device (AMD/NVIDIA/Intel; Apple via MoltenVK). device 0 = first GPU. + // ggml-Vulkan THROWS (vk::SystemError) when the Vulkan loader/driver is absent + // or pre-1.2 — both the device-count probe and init are wrapped so a GPU-less + // machine falls through to the CPU backend instead of failing the model load. if (params.use_gpu && !model->backend) { - fprintf(stderr, "%s: using Vulkan backend\n", __func__); - model->backend = ggml_backend_vk_init(0); + try { + if (ggml_backend_vk_get_device_count() > 0) { + model->backend = ggml_backend_vk_init(0); + } + } catch (const std::exception& e) { + fprintf(stderr, "%s: Vulkan probe failed (%s)\n", __func__, e.what()); + model->backend = nullptr; + } catch (...) { + model->backend = nullptr; + } + if (model->backend) { + fprintf(stderr, "%s: using Vulkan backend (%s)\n", __func__, ggml_backend_name(model->backend)); + } else { + fprintf(stderr, "%s: Vulkan backend unavailable (no Vulkan device/driver) — falling back\n", __func__); + } } #endif if (!model->backend) { @@ -5012,30 +5035,22 @@ static void edgetam_build_repvit_graph(struct ggml_context* ctx, } } -#ifdef SAM3_COREML -#include "edgetam_coreml.h" - -// RFD 0011 U8 (threaded encoder-ahead): the producer thread runs preprocess + -// CoreML encode for frame N+1 into s_prefetch_neck while the main thread runs the -// consumer (assembly+memattn+decode+memencode+bank) for frame N. When a prefetched -// neck is ready, the main thread's encode step copies it (skipping BOTH preprocess -// and the CoreML encode), so the ~17 ms encoder leg overlaps the ~52 ms consumer. -// SPSC: the bench serializes set→consume per frame, so no lock is needed here. +// ── RFD 0011: prefetched-neck seam (backend-agnostic) ──────────────────────── +// Threaded encoder-ahead: a PRODUCER thread runs preprocess + encode for frame +// N+1 into s_prefetch_neck while the main thread runs the consumer +// (assembly+memattn+decode+memencode+bank) for frame N. When a prefetched neck +// is ready, the main thread's encode step copies it (skipping BOTH preprocess +// and the encode), so the encoder leg overlaps the consumer. The producer engine +// is CoreML on Apple (edgetam_coreml_encode) and a SECOND ggml backend instance +// on CUDA (sam3_encoder_ahead_encode) — the seam itself is raw float arrays with +// no backend dependency. Function names keep the historic sam3_coreml_ prefix +// for source compatibility with the existing capi + threaded benches. +// SPSC: the consumer thread is the only reader AND the only caller of +// sam3_coreml_set_prefetched_neck (per the capi contract), so no lock is needed. +// Fixed EdgeTAM@1024 contract: levels 256/128/64 x 256ch. static std::vector s_prefetch_neck[3]; static bool s_prefetch_ready = false; -// RFD 0011 U8 — per-chip CoreML compute-unit profile (0=ALL 1=ANE 2=GPU 3=CPU). -// M4 Pro sweep (sam3_coreml_pipeline, isolated 20-rep median): encoder ANE 11ms -// (GPU 14, CPU 38), mem-attn GPU 17ms (ANE 26, CPU 24), decoder ANE 7.8ms (GPU 8.2), -// memenc ANE 5.0 / GPU 3.9. Defaults = the validated M4 optimum (enc ANE, mem-attn -// GPU, decoder ANE, memenc ANE — memenc's isolated GPU edge washes out in thermal -// noise end-to-end). NOT inverted vs M1 like the ORT/EfficientTAM stack. Override -// per chip via env (e.g. SAM3_COREML_MA_UNIT=1) if M1/M2/M3 profiles differ. -static int sam3_coreml_unit(const char* env, int def) { - const char* v = getenv(env); - return (v && *v) ? atoi(v) : def; -} - void sam3_coreml_set_prefetched_neck(const float* n0, const float* n1, const float* n2) { const int Wd[3] = {256, 128, 64}, D = 256; const float* src[3] = {n0, n1, n2}; @@ -5048,6 +5063,69 @@ std::vector sam3_coreml_preprocess_image(const sam3_image& image, int img return sam2_preprocess_image(image, img_size); } +// Load 3 neck levels (contiguous [D,W,H] floats, EdgeTAM@1024: 256/128/64 x 256ch) +// into the persistent state tensors, (re)allocating state + the sinusoidal PE on +// the first frame after create/reset. Factored from the CoreML encode helper so +// the ggml prefetch-consume path (encoder-ahead on CUDA) loads necks identically. +// Per rule reference-rules-in-comments: think-in-systems — one load path, two +// producers (CoreML / second ggml backend). +static bool edgetam_load_neck_from_floats(sam3_state& state, const sam3_model& model, + std::vector nbuf[3]) { + const int D = model.hparams.neck_dim; // 256 + const int Wd[3] = {256, 128, 64}; + const int Hd[3] = {256, 128, 64}; + SAM3_TIME_SCOPE(state_update_ms); + bool reuse = state.buffer && state.pe_buf && state.neck_trk[0] + && state.neck_trk[0]->ne[0] == D && state.neck_trk[0]->ne[1] == Wd[0]; + if (!reuse) { + if (state.buffer) { ggml_backend_buffer_free(state.buffer); state.buffer = nullptr; } + if (state.pe_buf) { ggml_backend_buffer_free(state.pe_buf); state.pe_buf = nullptr; } + if (state.pe_ctx) { ggml_free(state.pe_ctx); state.pe_ctx = nullptr; } + if (state.ctx) { ggml_free(state.ctx); state.ctx = nullptr; } + struct ggml_init_params sp = {ggml_tensor_overhead() * 32, nullptr, true}; + state.ctx = ggml_init(sp); + for (int i = 0; i < 3; ++i) { + state.neck_trk[i] = ggml_new_tensor_4d(state.ctx, GGML_TYPE_F32, D, Wd[i], Hd[i], 1); + char nm[32]; snprintf(nm, sizeof(nm), "neck_trk_%d", i); ggml_set_name(state.neck_trk[i], nm); + } + state.neck_trk[3] = nullptr; + state.buffer = ggml_backend_alloc_ctx_tensors(state.ctx, model.backend); + // Positional encoding (same as the ggml path: sinusoidal_pe_2d(H, W, D)). + struct ggml_init_params pp = {ggml_tensor_overhead() * 16, nullptr, true}; + state.pe_ctx = ggml_init(pp); + for (int i = 0; i < 3; ++i) { + state.neck_trk_pe[i] = ggml_new_tensor_4d(state.pe_ctx, GGML_TYPE_F32, D, Wd[i], Hd[i], 1); + char nm[32]; snprintf(nm, sizeof(nm), "neck_trk_pe_%d", i); ggml_set_name(state.neck_trk_pe[i], nm); + } + state.pe_buf = ggml_backend_alloc_ctx_tensors(state.pe_ctx, model.backend); + for (int i = 0; i < 3; ++i) { + auto pe = sam3_sinusoidal_pe_2d(Hd[i], Wd[i], D); + ggml_backend_tensor_set(state.neck_trk_pe[i], pe.data(), 0, pe.size() * sizeof(float)); + } + } + // Contiguous [D,W,H] floats load with a DIRECT copy — no per-frame transpose + // (CoreML exports channels-last [1,H,W,D] whose bytes are identical; a ggml + // producer's ggml_backend_tensor_get output is already in this layout). + for (int i = 0; i < 3; ++i) + ggml_backend_tensor_set(state.neck_trk[i], nbuf[i].data(), 0, nbuf[i].size() * sizeof(float)); + return true; +} + +#ifdef SAM3_COREML +#include "edgetam_coreml.h" + +// RFD 0011 U8 — per-chip CoreML compute-unit profile (0=ALL 1=ANE 2=GPU 3=CPU). +// M4 Pro sweep (sam3_coreml_pipeline, isolated 20-rep median): encoder ANE 11ms +// (GPU 14, CPU 38), mem-attn GPU 17ms (ANE 26, CPU 24), decoder ANE 7.8ms (GPU 8.2), +// memenc ANE 5.0 / GPU 3.9. Defaults = the validated M4 optimum (enc ANE, mem-attn +// GPU, decoder ANE, memenc ANE — memenc's isolated GPU edge washes out in thermal +// noise end-to-end). NOT inverted vs M1 like the ORT/EfficientTAM stack. Override +// per chip via env (e.g. SAM3_COREML_MA_UNIT=1) if M1/M2/M3 profiles differ. +static int sam3_coreml_unit(const char* env, int def) { + const char* v = getenv(env); + return (v && *v) ? atoi(v) : def; +} + // RFD 0011 (CoreML/ANE hybrid): run the EdgeTAM encoder on CoreML/ANE (~12 ms) // instead of the ggml/Metal RepViT (~122 ms) and load the ggml-matching 256-ch // neck features into state.neck_trk. The CoreML model exports neck(trunk(x))[0:3] @@ -5087,45 +5165,10 @@ static bool edgetam_encode_image_coreml(sam3_state& state, const sam3_model& mod return false; } - // Build/refresh the persistent state (U1 reuse) and permute CoreML's - // [D,H,W] (NCHW, W innermost) into ggml neck_trk [D,W,H] (D innermost). - { - SAM3_TIME_SCOPE(state_update_ms); - bool reuse = state.buffer && state.pe_buf && state.neck_trk[0] - && state.neck_trk[0]->ne[0] == D && state.neck_trk[0]->ne[1] == Wd[0]; - if (!reuse) { - if (state.buffer) { ggml_backend_buffer_free(state.buffer); state.buffer = nullptr; } - if (state.pe_buf) { ggml_backend_buffer_free(state.pe_buf); state.pe_buf = nullptr; } - if (state.pe_ctx) { ggml_free(state.pe_ctx); state.pe_ctx = nullptr; } - if (state.ctx) { ggml_free(state.ctx); state.ctx = nullptr; } - struct ggml_init_params sp = {ggml_tensor_overhead() * 32, nullptr, true}; - state.ctx = ggml_init(sp); - for (int i = 0; i < 3; ++i) { - state.neck_trk[i] = ggml_new_tensor_4d(state.ctx, GGML_TYPE_F32, D, Wd[i], Hd[i], 1); - char nm[32]; snprintf(nm, sizeof(nm), "neck_trk_%d", i); ggml_set_name(state.neck_trk[i], nm); - } - state.neck_trk[3] = nullptr; - state.buffer = ggml_backend_alloc_ctx_tensors(state.ctx, model.backend); - // Positional encoding (same as the ggml path: sinusoidal_pe_2d(H, W, D)). - struct ggml_init_params pp = {ggml_tensor_overhead() * 16, nullptr, true}; - state.pe_ctx = ggml_init(pp); - for (int i = 0; i < 3; ++i) { - state.neck_trk_pe[i] = ggml_new_tensor_4d(state.pe_ctx, GGML_TYPE_F32, D, Wd[i], Hd[i], 1); - char nm[32]; snprintf(nm, sizeof(nm), "neck_trk_pe_%d", i); ggml_set_name(state.neck_trk_pe[i], nm); - } - state.pe_buf = ggml_backend_alloc_ctx_tensors(state.pe_ctx, model.backend); - for (int i = 0; i < 3; ++i) { - auto pe = sam3_sinusoidal_pe_2d(Hd[i], Wd[i], D); - ggml_backend_tensor_set(state.neck_trk_pe[i], pe.data(), 0, pe.size() * sizeof(float)); - } - } - // The CoreML model exports channels-last [1,H,W,D], whose contiguous - // bytes are byte-identical to ggml neck_trk [D,W,H] (d innermost, then W, - // then H). So each level loads with a DIRECT copy — no per-frame transpose. - for (int i = 0; i < 3; ++i) - ggml_backend_tensor_set(state.neck_trk[i], nbuf[i].data(), 0, nbuf[i].size() * sizeof(float)); - } - return true; + // Build/refresh the persistent state (U1 reuse) and load CoreML's channels- + // last [1,H,W,D] output — byte-identical to ggml neck_trk [D,W,H] — via the + // shared (backend-agnostic) neck loader. + return edgetam_load_neck_from_floats(state, model, nbuf); } #endif // SAM3_COREML @@ -5158,6 +5201,18 @@ static bool edgetam_encode_image(sam3_state& state, } #endif + // RFD 0011 (ggml encoder-ahead): a producer thread (sam3_encoder_ahead_encode, + // a SECOND ggml backend instance — own CUDA stream) already encoded THIS + // frame's neck into the prefetch seam. Load it into state and skip the whole + // in-line encoder graph, mirroring the CoreML prefetch path above. Guarded to + // the seam's fixed EdgeTAM@1024 contract (levels 256/128/64). + if (s_prefetch_ready && hp.is_edgetam() && img_size == 1024) { + static std::vector nbuf[3]; // consumer-thread-only (SPSC seam) + for (int i = 0; i < 3; ++i) nbuf[i].swap(s_prefetch_neck[i]); + s_prefetch_ready = false; + return edgetam_load_neck_from_floats(state, model, nbuf); + } + // ── Preprocess (same ImageNet normalization as SAM2) ───────────────── // RFD 0011 U0: preprocess timed separately from the graph build/alloc/compute. std::vector img_data; @@ -5362,6 +5417,105 @@ static bool edgetam_encode_image(sam3_state& state, return true; } +/***************************************************************************** +** RFD 0011 — ggml encoder-ahead producer (CUDA) +** +** The NVIDIA counterpart of the CoreML producer-encoder pool: a SECOND ggml +** backend instance on the same CUDA device gives the producer thread its own +** stream + cublas handles, so the frame-N+1 encoder graph overlaps the +** frame-N consumer (mem-attn/decoder/mem-enc) on the GPU's own scheduler. +** The weight tensors stay shared (same-device memory, read-only after load). +** +** CUDA-only for now: cross-instance buffer reads are unverified on Vulkan, +** and the M4 compute-unit sweep showed same-accelerator producer/consumer +** co-location can LOSE (the Mac win came from ANE parallel to GPU) — so this +** must be A/B-measured on real NVIDIA hardware (the platform encoder-ahead +** bench) before it can default on. Callers opt in via the capi +** (SAM3_GGML_ENCODER_AHEAD=1). +*****************************************************************************/ + +struct sam3_encoder_ahead { + ggml_backend_t backend = nullptr; // producer-owned second backend instance + ggml_gallocr_t galloc = nullptr; // producer-owned graph allocator +}; + +sam3_encoder_ahead* sam3_encoder_ahead_create(const sam3_model& model) { +#ifdef GGML_USE_CUDA + // A ggml_backend_t is not safe for concurrent graph_compute from two + // threads, so the producer gets its own CUDA context (own stream). Weights + // live in device memory on the same GPU and are read-only after load, so + // both instances can read them. Only offered when the consumer actually + // runs on CUDA (a CPU/Vulkan consumer cannot read another instance's + // buffers safely) and the model is EdgeTAM (the seam's fixed contract). + if (!model.backend || !ggml_backend_is_cuda(model.backend)) return nullptr; + if (!model.hparams.is_edgetam()) return nullptr; + ggml_backend_t b = ggml_backend_cuda_init(0); + if (!b) return nullptr; + auto* ea = new sam3_encoder_ahead(); + ea->backend = b; + ea->galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(b)); + if (!ea->galloc) { ggml_backend_free(b); delete ea; return nullptr; } + fprintf(stderr, "%s: ggml encoder-ahead producer on %s\n", __func__, ggml_backend_name(b)); + return ea; +#else + (void)model; // producer pool unavailable on non-CUDA ggml builds + return nullptr; +#endif +} + +// Preprocess + run the EdgeTAM RepViT+FPN encoder graph on the PRODUCER backend +// and extract the 3 neck levels as contiguous [D,W,H] floats (the prefetch-seam +// layout). n0/n1/n2 must hold 256*256*256, 256*128*128, 256*64*64 floats. +bool sam3_encoder_ahead_encode(sam3_encoder_ahead* ea, const sam3_model& model, + const sam3_image& image, + float* n0, float* n1, float* n2) { + if (!ea || !ea->backend || !n0 || !n1 || !n2) return false; + const int img_size = 1024; // the prefetch seam's fixed EdgeTAM contract + + std::vector img_data = sam2_preprocess_image(image, img_size); + + // Build the same encoder graph edgetam_encode_image builds (build cost is + // ~0.3 ms per U0; the gallocr reuses its reservation across frames). + const size_t buf_size = ggml_tensor_overhead() * 16384 + ggml_graph_overhead() * 2; + struct ggml_init_params gparams = {buf_size, nullptr, true}; + struct ggml_context* ctx0 = ggml_init(gparams); + if (!ctx0) return false; + + struct ggml_tensor* inp = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, img_size, img_size, 3, 1); + ggml_set_name(inp, "input_image"); + ggml_set_input(inp); + + struct ggml_tensor* stage_outs[4] = {}; + edgetam_build_repvit_graph(ctx0, inp, model, stage_outs); + struct ggml_tensor* fpn_outs[4] = {}; + edgetam_build_fpn_neck_graph(ctx0, stage_outs, model, fpn_outs); + for (int i = 0; i < 3; ++i) ggml_set_output(fpn_outs[i]); + + struct ggml_cgraph* graph = ggml_new_graph_custom(ctx0, 32768, false); + for (int i = 0; i < 3; ++i) ggml_build_forward_expand(graph, fpn_outs[i]); + + bool ok = ggml_gallocr_alloc_graph(ea->galloc, graph); + if (ok) { + ggml_backend_tensor_set(inp, img_data.data(), 0, img_data.size() * sizeof(float)); + // n_threads only applies to a CPU backend; the producer is CUDA here. + ok = sam3_graph_compute(ea->backend, graph, 4); + } + if (ok) { + float* dst[3] = {n0, n1, n2}; + for (int i = 0; i < 3; ++i) + ggml_backend_tensor_get(fpn_outs[i], dst[i], 0, ggml_nbytes(fpn_outs[i])); + } + ggml_free(ctx0); + return ok; +} + +void sam3_encoder_ahead_destroy(sam3_encoder_ahead* ea) { + if (!ea) return; + if (ea->galloc) ggml_gallocr_free(ea->galloc); + if (ea->backend) ggml_backend_free(ea->backend); + delete ea; +} + /***************************************************************************** ** EdgeTAM Profiling ** diff --git a/sam3.h b/sam3.h index 7c624c6..629428d 100644 --- a/sam3.h +++ b/sam3.h @@ -365,17 +365,39 @@ sam3_image sam3_decode_video_frame(const std::string & video_path, int fram sam3_video_info sam3_get_video_info(const std::string & video_path); /* -** ── RFD 0011 U8: threaded encoder-ahead (CoreML) ───────────────────────── -** The producer thread preprocesses + CoreML-encodes frame N+1 off the main -** thread, then hands the 3 neck levels here; the next sam3_propagate_frame's -** encode step consumes them (skipping preprocess + encode) so the encoder leg -** overlaps the consumer. Built only with SAM3_COREML; no-ops otherwise. +** ── RFD 0011 U8: threaded encoder-ahead (backend-agnostic seam) ────────── +** The producer thread preprocesses + encodes frame N+1 off the main thread, +** then hands the 3 neck levels here; the next sam3_propagate_frame's encode +** step consumes them (skipping preprocess + encode) so the encoder leg +** overlaps the consumer. The producer engine is CoreML on Apple or a second +** ggml-CUDA backend instance on NVIDIA (sam3_encoder_ahead_* below); the +** seam itself is raw floats — available on every build. Names keep the +** historic sam3_coreml_ prefix for source compatibility. */ void sam3_coreml_set_prefetched_neck(const float* neck0, const float* neck1, const float* neck2); std::vector sam3_coreml_preprocess_image(const sam3_image& image, int img_size); +/* +** ── RFD 0011: ggml encoder-ahead producer (CUDA) ───────────────────────── +** NVIDIA counterpart of the CoreML producer encoder: a SECOND ggml backend +** instance (own CUDA stream) runs the EdgeTAM encoder graph for frame N+1 +** concurrently with the consumer's mem-attn/decoder work for frame N. +** create returns NULL unless the model's active backend is CUDA and the +** model is EdgeTAM. encode writes the 3 neck levels in the prefetch-seam +** layout ([D,W,H] contiguous floats; 256*256*256 / 256*128*128 / 256*64*64). +** Same-GPU producer/consumer contention is hardware-dependent — A/B on the +** target GPU before defaulting on (the capi gates this behind +** SAM3_GGML_ENCODER_AHEAD=1). +*/ +struct sam3_encoder_ahead; +sam3_encoder_ahead* sam3_encoder_ahead_create(const sam3_model& model); +bool sam3_encoder_ahead_encode(sam3_encoder_ahead* ea, const sam3_model& model, + const sam3_image& image, + float* neck0, float* neck1, float* neck2); +void sam3_encoder_ahead_destroy(sam3_encoder_ahead* ea); + /***************************************************************************** ** Test and Debug API ** From d0d2c4565fd8d3389e859cfea0798f9f72582e0e Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 16:01:47 -0600 Subject: [PATCH 06/16] cuda: fall back to manual attention for head dims ggml-cuda flash-attn lacks MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ggml-CUDA's fused flash-attention kernels only cover head sizes {40,64,72,80,96,112,128,256,576} (ggml-cuda/fattn.cu). EdgeTAM's SAM mask-decoder attention uses head_dim=32, so ggml_cuda_get_best_fattn_kernel returns BEST_FATTN_KERNEL_NONE and GGML_ABORTs on the first Track() — while model load and the RepViT encoder run fine on CUDA0. Metal and Vulkan accept head_dim=32, so this gap is CUDA-specific and never surfaced on those backends. Add sam3_fa_ext(), a wrapper around ggml_flash_attn_ext. On the CUDA build only (#ifdef GGML_USE_CUDA) it routes head dims outside the supported set through a mathematically equivalent manual attention — the exact QK^T -> soft_max_ext -> .V -> permute(0,2,1,3) idiom already used interchangeably with flash in sam3_ddec_layer_forward, so the [HD,NH,N_q,B] output layout is identical and no call site changes. Flash is still used on CUDA for supported head dims (encoder, memory attention), preserving perf there. On non-CUDA builds sam3_fa_ext forwards verbatim to ggml_flash_attn_ext, so the Metal/CoreML graph and numerics are byte-for-byte unchanged. All 17 flash-attn call sites route through the wrapper. Enables the Windows x64 + CUDA EdgeTAM 512 hold to run per-frame inference on NVIDIA GPUs (RTX 4060 / Ada sm_89, CUDA 12.9). Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 71 ++++++++++++++++++++++++++++++++++++++++++-------------- 1 file changed, 54 insertions(+), 17 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 93e79b8..86326d2 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -3915,6 +3915,43 @@ static struct ggml_tensor* sam3_apply_rope(struct ggml_context* ctx, return ggml_reshape_3d(ctx, ggml_cont(ctx, out), head_dim, N, nheads_B); } +// ── CUDA flash-attn head-dim fallback (Windows x64 + CUDA port) ────────────── +// ggml-CUDA's fused flash-attention kernels only cover head sizes +// {40,64,72,80,96,112,128,256,576} (ggml-cuda/fattn.cu). EdgeTAM's SAM +// mask-decoder attention uses head_dim=32, for which the CUDA selector returns +// BEST_FATTN_KERNEL_NONE and GGML_ABORTs. Metal and Vulkan accept head_dim=32, +// so this is a CUDA-only gap — NOT a model/graph change. On the CUDA build ONLY +// we route the unsupported head dims through a mathematically equivalent manual +// attention (the exact idiom the fork already uses for the fp32-mask path in +// sam3_ddec_layer_forward): QK^T → soft_max_ext(scale,+mask) → ·V, laid out to +// match ggml_flash_attn_ext's [HD, NH, N_q, B] output so every call site is +// unchanged. On non-CUDA builds this forwards verbatim to ggml_flash_attn_ext, +// so the Metal/CoreML path and its numerics are byte-for-byte untouched. +static struct ggml_tensor* sam3_fa_ext(struct ggml_context* ctx, + struct ggml_tensor* q, // [HD, N_q, NH, B] + struct ggml_tensor* k, // [HD, N_kv, NH_kv, B] + struct ggml_tensor* v, // [HD, N_kv, NH_kv, B] + struct ggml_tensor* mask, // additive [N_kv, N_q, ...] or null + float scale, float max_bias, float logit_softcap) { +#ifdef GGML_USE_CUDA + const int64_t hd = q->ne[0]; + const bool fa_ok = + (hd == 40 || hd == 64 || hd == 72 || hd == 80 || hd == 96 || + hd == 112 || hd == 128 || hd == 256 || hd == 576) && + k->ne[0] == hd && v->ne[0] == hd && + max_bias == 0.0f && logit_softcap == 0.0f; + if (!fa_ok) { + struct ggml_tensor* kq = ggml_mul_mat(ctx, k, q); // [N_kv, N_q, NH, B] + kq = ggml_soft_max_ext(ctx, kq, mask, scale, max_bias); // softmax(scale*kq + mask) + struct ggml_tensor* v_t = ggml_cont(ctx, ggml_transpose(ctx, v)); // [N_kv, HD, NH_kv, B] + struct ggml_tensor* kqv = ggml_mul_mat(ctx, v_t, kq); // [HD, N_q, NH, B] + kqv = ggml_cont(ctx, ggml_permute(ctx, kqv, 0, 2, 1, 3)); // [HD, NH, N_q, B] + return kqv; + } +#endif + return ggml_flash_attn_ext(ctx, q, k, v, mask, scale, max_bias, logit_softcap); +} + // Single ViT block forward: pre-norm → attn (window or global, with RoPE) → residual → pre-norm → MLP → residual // x: [E, W, H, B] in ggml layout (following sam.cpp convention) static struct ggml_tensor* sam3_vit_block_forward(struct ggml_context* ctx, @@ -3981,7 +4018,7 @@ static struct ggml_tensor* sam3_vit_block_forward(struct ggml_context* ctx, K = ggml_reshape_4d(ctx, K, HD, W_cur * H_cur, NH, B_cur); float scale = 1.0f / sqrtf((float)HD); - auto* attn_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* attn_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); // flash_attn_ext returns [HD, NH, N, B_cur] — HD and NH adjacent, // so reshaping directly to [E, W, H, B] is correct. x = ggml_reshape_4d(ctx, attn_out, E, W_cur, H_cur, B_cur); @@ -4201,7 +4238,7 @@ static struct ggml_tensor* sam3_text_block_forward(struct ggml_context* ctx, V = ggml_permute(ctx, V, 0, 2, 1, 3); // non-contiguous; flash_attn uses strides float scale = 1.0f / sqrtf((float)HD); - auto* attn_out = ggml_flash_attn_ext(ctx, Q, K, V, causal_mask, scale, 0.0f, 0.0f); + auto* attn_out = sam3_fa_ext(ctx, Q, K, V, causal_mask, scale, 0.0f, 0.0f); x = ggml_reshape_2d(ctx, attn_out, E, L); x = ggml_mul_mat(ctx, blk.attn_out_proj_w, x); @@ -4602,7 +4639,7 @@ static struct ggml_tensor* sam2_hiera_block_forward(struct ggml_context* ctx, V = ggml_permute(ctx, V, 0, 2, 1, 3); // non-contiguous OK for flash_attn float scale = 1.0f / sqrtf((float)head_dim); - auto* attn_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0, 0); + auto* attn_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0, 0); // Recombine: flash_attn output is [HD, N_q, NH, B_win] // → reshape to [C_out, N_q, B_win] @@ -6013,7 +6050,7 @@ static struct ggml_tensor* edgetam_perceiver_layer_forward( k = ggml_reshape_4d(ctx, k, D, N_x, 1, batch); v = ggml_reshape_4d(ctx, v, D, N_x, 1, batch); - auto* attn = ggml_flash_attn_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); + auto* attn = sam3_fa_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); // Output: [D, 1, N_lat, batch] (permuted) → reshape to [D, N_lat, batch] attn = ggml_reshape_3d(ctx, attn, D, N_lat, batch); @@ -6051,7 +6088,7 @@ static struct ggml_tensor* edgetam_perceiver_layer_forward( k = ggml_reshape_4d(ctx, k, D, N_lat, 1, batch); v = ggml_reshape_4d(ctx, v, D, N_lat, 1, batch); - auto* attn = ggml_flash_attn_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); + auto* attn = sam3_fa_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); attn = ggml_reshape_3d(ctx, attn, D, N_lat, batch); auto* sa_out = ggml_mul_mat(ctx, layer.sa_out_w, attn); @@ -7353,7 +7390,7 @@ static struct ggml_tensor * sam3_build_vit_attn_core_from_qkv(struct ggml_contex K = ggml_reshape_4d(ctx, K, HD, W_cur * H_cur, NH, B_cur); const float scale = 1.0f / sqrtf((float) HD); - struct ggml_tensor * attn_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + struct ggml_tensor * attn_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); return ggml_cont(ctx, ggml_reshape_4d(ctx, attn_out, E, W_cur, H_cur, B_cur)); } @@ -8070,7 +8107,7 @@ static struct ggml_tensor* sam3_multihead_attn_fused( V = ggml_permute(ctx, V, 0, 2, 1, 3); // [HD, N_kv, NH, B] non-contiguous; flash_attn uses strides float scale = 1.0f / sqrtf((float)HD); - auto* attn_out = ggml_flash_attn_ext(ctx, Q, K, V, attn_mask, scale, 0.0f, 0.0f); + auto* attn_out = sam3_fa_ext(ctx, Q, K, V, attn_mask, scale, 0.0f, 0.0f); auto* merged = ggml_reshape_3d(ctx, attn_out, D, N_q, B); merged = ggml_mul_mat(ctx, out_proj_w, merged); @@ -8253,7 +8290,7 @@ static sam3_geom_result sam3_build_geom_enc_graph( V = ggml_permute(ctx, V, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)HD); - auto* sa_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* sa_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); sa_out = ggml_reshape_3d(ctx, sa_out, D, S, 1); sa_out = ggml_mul_mat(ctx, ly.sa_out_proj_w, sa_out); @@ -8295,7 +8332,7 @@ static sam3_geom_result sam3_build_geom_enc_graph( V = ggml_permute(ctx, V, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)HD); - auto* ca_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* ca_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); ca_out = ggml_reshape_3d(ctx, ca_out, D, S_q, 1); ca_out = ggml_mul_mat(ctx, ly.ca_out_w, ca_out); @@ -8603,7 +8640,7 @@ static struct ggml_tensor* sam3_fenc_layer_forward( V = ggml_permute(ctx, V, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)HD); - auto* sa_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* sa_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); sa_out = ggml_reshape_3d(ctx, sa_out, D, N, B); sa_out = ggml_mul_mat(ctx, ly.sa_out_proj_w, sa_out); @@ -8645,7 +8682,7 @@ static struct ggml_tensor* sam3_fenc_layer_forward( auto* ca_mask = sam3_expand_token_attn_bias(ctx, prompt_attn_bias, N_q, n_heads, B); float scale = 1.0f / sqrtf((float)HD); - auto* ca_out = ggml_flash_attn_ext(ctx, Q, K, V, ca_mask, scale, 0.0f, 0.0f); + auto* ca_out = sam3_fa_ext(ctx, Q, K, V, ca_mask, scale, 0.0f, 0.0f); ca_out = ggml_reshape_3d(ctx, ca_out, D, N_q, B); ca_out = ggml_mul_mat(ctx, ly.ca_out_w, ca_out); @@ -9021,7 +9058,7 @@ static struct ggml_tensor* sam3_ddec_layer_forward( V = ggml_permute(ctx, V, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)HD); - auto* sa_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* sa_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); sa_out = ggml_reshape_3d(ctx, sa_out, D, N, B); sa_out = ggml_mul_mat(ctx, ly.sa_out_proj_w, sa_out); sa_out = ggml_add(ctx, sa_out, ly.sa_out_proj_b); @@ -9064,7 +9101,7 @@ static struct ggml_tensor* sam3_ddec_layer_forward( auto* text_mask = sam3_expand_token_attn_bias(ctx, text_attn_bias, N_q, n_heads, B); float scale = 1.0f / sqrtf((float)HD); - auto* ca_out = ggml_flash_attn_ext(ctx, Q, K, V, text_mask, scale, 0.0f, 0.0f); + auto* ca_out = sam3_fa_ext(ctx, Q, K, V, text_mask, scale, 0.0f, 0.0f); ca_out = ggml_reshape_3d(ctx, ca_out, D, N_q, B); ca_out = ggml_mul_mat(ctx, ly.ca_text_out_w, ca_out); ca_out = ggml_add(ctx, ca_out, ly.ca_text_out_b); @@ -9120,7 +9157,7 @@ static struct ggml_tensor* sam3_ddec_layer_forward( ca_out = ggml_mul_mat(ctx, v_t, kq); // [HD, N_q, NH, B] ca_out = ggml_cont(ctx, ggml_permute(ctx, ca_out, 0, 2, 1, 3)); } else { - ca_out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + ca_out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); } ca_out = ggml_reshape_3d(ctx, ca_out, D, N_q, B); @@ -9800,7 +9837,7 @@ static struct ggml_tensor* sam3_build_mem_attn_graph( v = ggml_permute(ctx, v, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)D); - auto* sa_out = ggml_flash_attn_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); + auto* sa_out = sam3_fa_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); sa_out = ggml_reshape_3d(ctx, sa_out, D, N, 1); sa_out = ggml_add(ctx, ggml_mul_mat(ctx, ly.sa_out_w, sa_out), ly.sa_out_b); x = ggml_add(ctx, x, sa_out); @@ -9845,7 +9882,7 @@ static struct ggml_tensor* sam3_build_mem_attn_graph( v = ggml_permute(ctx, v, 0, 2, 1, 3); float scale = 1.0f / sqrtf((float)D); - auto* ca_out = ggml_flash_attn_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); + auto* ca_out = sam3_fa_ext(ctx, q, k, v, nullptr, scale, 0.0f, 0.0f); ca_out = ggml_reshape_3d(ctx, ca_out, D, N, 1); ca_out = ggml_add(ctx, ggml_mul_mat(ctx, ly.ca_out_w, ca_out), ly.ca_out_b); x = ggml_add(ctx, x, ca_out); @@ -10819,7 +10856,7 @@ static struct ggml_tensor* sam3_sam_attention( // Attention float scale = 1.0f / sqrtf((float)HD); - auto* out = ggml_flash_attn_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); + auto* out = sam3_fa_ext(ctx, Q, K, V, nullptr, scale, 0.0f, 0.0f); // out: [HD, NH, N_q, B] (flash_attn_ext swaps dims 1,2 vs input) #if 0 // Manual SDPA (for debugging only) From 356475eb17a94373407eab28a96d0ff9ccc82c17 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 17:50:42 -0600 Subject: [PATCH 07/16] =?UTF-8?q?cuda:=20stage-transport=20driver=20increm?= =?UTF-8?q?ent=201=20=E2=80=94=20sam3=5Fstage=5Ffeed=20+=20the=204=20neck?= =?UTF-8?q?=5Ftrk=20edges?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit report/CUDA-RESIDENCY-DESIGN.md §5 increment 1. Adds the sam3_transport driver (host default on every backend = today's exact tensor_get/tensor_set sequence; CUDA device opt-in via SAM3_STAGE_TRANSPORT=device = one same-backend D2D ggml_backend_tensor_copy) and converts the four heaviest inter-stage edges: mem-attn curr (with the 4D-leaf + in-graph-reshape fix so D2D layouts match; host graph bit-for-bit unchanged), decoder trk_s0/trk_s1, memenc pix_in_raw. Removes 4 D2H + 4 H2D = ~176.7 MB/frame (76% of PCIe bytes). Gates: host-mode output byte-identical to pre-change baseline; host-vs-device A/B byte-identical over 39 synthetic frames incl. masks; 7.3 -> 9.5-9.7 fps. Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 131 ++++++++++++++++++++++++++++++++++++++++++++++--------- 1 file changed, 110 insertions(+), 21 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 86326d2..5e5a123 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -934,6 +934,87 @@ struct edgetam_perceiver { std::vector layers; }; +/***************************************************************************** +** Stage-transport driver (CUDA device residency) +** +** report/CUDA-RESIDENCY-DESIGN.md (trackbench Windows port). The CLAUDE.md +** stage-isolation pattern moves inter-stage tensors through CPU-side +** std::vector buffers — free on Apple unified memory, ~231 MB/frame over +** PCIe on a discrete NVIDIA card (~95% of GPU time measured as memcpy). +** This driver changes ONLY the transport between stages: +** HOST (default on EVERY backend): today's exact instruction sequence — +** ggml_backend_tensor_get into a temp vector + tensor_set. The mac +** Metal/CoreML path is byte-identical by construction. +** DEVICE (CUDA only, opt-in via SAM3_STAGE_TRANSPORT=device): one +** same-backend D2D ggml_backend_tensor_copy — byte-identical +** destination bytes, zero PCIe. In-tree precedent: the encoder's +** steady-state FPN → state.neck_trk copy. +** Graph isolation (CLAUDE.md rules 1-2) is untouched: per-stage graphs and +** fresh input leaves stay; only HOW bytes reach those leaves changes. +*****************************************************************************/ + +enum sam3_transport_kind { SAM3_TRANSPORT_HOST = 0, SAM3_TRANSPORT_DEVICE = 1 }; + +struct sam3_transport { + sam3_transport_kind kind = SAM3_TRANSPORT_HOST; + ggml_backend_t backend = nullptr; +}; + +static bool sam3_same_layout(const struct ggml_tensor* a, const struct ggml_tensor* b) { + if (a->type != b->type) return false; + for (int i = 0; i < GGML_MAX_DIMS; ++i) + if (a->ne[i] != b->ne[i] || a->nb[i] != b->nb[i]) return false; + return true; +} + +// Persistent tensor -> fresh graph-input tensor. Device mode pre-checks and +// DEGRADES to the verbatim host path with a once-per-site stderr line instead +// of aborting — a silently degraded site is then caught by the per-increment +// nsys transfer-count assert as count drift, never as wrongness. +static void sam3_stage_feed(const sam3_transport& tr, struct ggml_tensor* src, + struct ggml_tensor* dst, const char* site) { + GGML_ASSERT(ggml_nbytes(src) == ggml_nbytes(dst)); + if (tr.kind == SAM3_TRANSPORT_DEVICE) { + if (sam3_same_layout(src, dst) && + !ggml_backend_buffer_is_host(src->buffer) && + !ggml_backend_buffer_is_host(dst->buffer)) { + ggml_backend_tensor_copy(src, dst); // same-backend D2D, self-fencing + return; + } + static std::map warned; + if (!warned[site]) { + warned[site] = true; + fprintf(stderr, "sam3_stage_feed: %s: layout/buffer mismatch — using host path\n", site); + } + } + std::vector tmp(ggml_nbytes(src) / sizeof(float)); + ggml_backend_tensor_get(src, tmp.data(), 0, ggml_nbytes(src)); + ggml_backend_tensor_set(dst, tmp.data(), 0, ggml_nbytes(dst)); +} + +// Resolved once at model load, after the backend chain (Metal>CUDA>Vulkan>CPU). +// Default host on every backend; device only when explicitly requested AND the +// answered backend is CUDA — Metal can never pass this gate, so the mac cannot +// select the device driver under any configuration. +static void sam3_transport_init(sam3_transport& tr, ggml_backend_t backend) { + tr.backend = backend; + tr.kind = SAM3_TRANSPORT_HOST; + const char* env = getenv("SAM3_STAGE_TRANSPORT"); + const bool force_host = env && strcmp(env, "host") == 0; + const bool want_device = env && strcmp(env, "device") == 0; + (void)force_host; +#ifdef GGML_USE_CUDA + if (want_device && !force_host && ggml_backend_is_cuda(backend)) { + tr.kind = SAM3_TRANSPORT_DEVICE; + fprintf(stderr, "%s: stage transport = device (CUDA-resident)\n", __func__); + if (getenv("SAM3_GGML_ENCODER_AHEAD")) + fprintf(stderr, "%s: note: encoder-ahead handoff stays host-float until its copy_async clause lands\n", __func__); + } +#endif + if (want_device && tr.kind != SAM3_TRANSPORT_DEVICE) + fprintf(stderr, "%s: SAM3_STAGE_TRANSPORT=device requested but backend is not CUDA; using host\n", __func__); +} + /***************************************************************************** ** Top-Level Opaque Types (defined here, forward-declared in sam3.h) *****************************************************************************/ @@ -988,6 +1069,9 @@ struct sam3_model { ggml_backend_t backend = nullptr; ggml_backend_buffer_t buffer = nullptr; + // stage-transport driver (host default everywhere; CUDA device opt-in) + sam3_transport transport; + // tensor lookup std::map tensors; @@ -3489,6 +3573,7 @@ std::shared_ptr sam3_load_model(const sam3_params& params) { fprintf(stderr, "%s: failed to init backend\n", __func__); return nullptr; } + sam3_transport_init(model->transport, model->backend); // ── Create ggml context (no_alloc — we use backend_alloc_ctx_tensors) // Estimate: ~3000 tensors, generous overhead @@ -12047,7 +12132,7 @@ static sam3_prop_output sam3_propagate_single( // Memory-attention inputs (ggml path only). In the CoreML mem-attn path these // stay null and cond_spatial is an input tensor filled from the CoreML output. - struct ggml_tensor* curr = nullptr, * src_pos_t = nullptr; + struct ggml_tensor* curr = nullptr, * curr_in = nullptr, * src_pos_t = nullptr; struct ggml_tensor* prompt_t = nullptr, * prompt_pos_t = nullptr; struct ggml_tensor* rope_q_t = nullptr, * rope_k_t = nullptr; struct ggml_tensor* cond_spatial = nullptr; @@ -12061,8 +12146,23 @@ static sam3_prop_output sam3_propagate_single( // CRITICAL: create fresh input tensors for state features. // Using state.neck_trk[*] directly as ggml_reshape operands pulls in // the entire ViT+neck recomputation from the image encoder graph. - curr = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N, 1); - ggml_set_name(curr, "prop_curr"); ggml_set_input(curr); + if (model.transport.kind == SAM3_TRANSPORT_DEVICE) { + // Device transport: the fresh input LEAF mirrors state.neck_trk[2]'s + // 4D layout so sam3_stage_feed can whole-tensor D2D into it; the + // mem-attn builder still sees [D, N, 1] via an in-graph reshape of + // the fresh leaf (no src[] ancestry — CLAUDE.md rule 2 cannot fire). + // Host mode keeps today's 3D tensor and graph bit-for-bit. + curr_in = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, + state.neck_trk[2]->ne[0], + state.neck_trk[2]->ne[1], + state.neck_trk[2]->ne[2], 1); + ggml_set_name(curr_in, "prop_curr"); ggml_set_input(curr_in); + curr = ggml_reshape_3d(ctx0, curr_in, D, N, 1); + } else { + curr = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N, 1); + ggml_set_name(curr, "prop_curr"); ggml_set_input(curr); + curr_in = curr; + } // src_pos (sinusoidal PE 256-dim for 72×72) src_pos_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N, 1); @@ -12167,21 +12267,14 @@ static sam3_prop_output sam3_propagate_single( ggml_backend_tensor_set(dense_emb, state.dense_nomask_cache.data(), 0, state.dense_nomask_cache.size() * sizeof(float)); - // Copy tracker features from state to fresh input tensors + // Copy tracker features from state to fresh input tensors (transport + // driver: host = today's round-trip, device = same-backend D2D). { if (curr) { // ggml mem-attn path only (CoreML reads neck_trk[2] itself) - std::vector c2(D * N); - ggml_backend_tensor_get(state.neck_trk[2], c2.data(), 0, D * N * sizeof(float)); - ggml_backend_tensor_set(curr, c2.data(), 0, D * N * sizeof(float)); + sam3_stage_feed(model.transport, state.neck_trk[2], curr_in, "prop_curr"); } - - std::vector s0(D * H0 * H0); - ggml_backend_tensor_get(state.neck_trk[0], s0.data(), 0, D * H0 * H0 * sizeof(float)); - ggml_backend_tensor_set(trk_s0, s0.data(), 0, D * H0 * H0 * sizeof(float)); - - std::vector s1(D * H1 * H1); - ggml_backend_tensor_get(state.neck_trk[1], s1.data(), 0, D * H1 * H1 * sizeof(float)); - ggml_backend_tensor_set(trk_s1, s1.data(), 0, D * H1 * H1 * sizeof(float)); + sam3_stage_feed(model.transport, state.neck_trk[0], trk_s0, "prop_trk_s0"); + sam3_stage_feed(model.transport, state.neck_trk[1], trk_s1, "prop_trk_s1"); } // RFD 0011 U0: backend compute of the shared mem-attn + mask-decoder graph @@ -12558,12 +12651,8 @@ static bool sam3_encode_memory( SAM3_TIME_END(mem_encoder_alloc_ms, _t_mem_alloc); ggml_backend_tensor_set(mask_in, m_interp.data(), 0, m_interp.size() * sizeof(float)); - // Copy pixel features from state tensor to fresh input - { - std::vector pix_data(D * H * H); - ggml_backend_tensor_get(state.neck_trk[2], pix_data.data(), 0, D * H * H * sizeof(float)); - ggml_backend_tensor_set(pix_in_raw, pix_data.data(), 0, D * H * H * sizeof(float)); - } + // Copy pixel features from state tensor to fresh input (transport driver) + sam3_stage_feed(model.transport, state.neck_trk[2], pix_in_raw, "mem_pix_feat"); // RFD 0011 U0: backend compute of the memory encoder (mem_encoder_compute_ms). SAM3_TIME_BEGIN(_t_mem_compute); if (!sam3_graph_compute(model.backend, g, 4)) { From d35fab437087ef325b9f73d93840e4634b13d3a0 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 17:59:20 -0600 Subject: [PATCH 08/16] =?UTF-8?q?cuda:=20stage-transport=20increment=202?= =?UTF-8?q?=20=E2=80=94=20constant=20device=20twins=20+=20weight-read=20ca?= =?UTF-8?q?ches?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit report/CUDA-RESIDENCY-DESIGN.md §5 increment 2. sam3_stagebuf residency handle: the existing host vector stays the value; device mode adds a persistent twin uploaded once per content generation, then D2D into the fresh graph input. Twins: rope_q/rope_k/src_pos (tracker, keyed on pe_gen), sparse/image_pe/ dense_emb (state, keyed on pe_cache_gen). Load-time host caches (device driver only) for the per-frame weight re-reads: obj-ptr MLP W/b, no_obj_ptr, perceiver latents_1d/2d, maskmem tpos. Host driver byte-identical: pushes are verbatim tensor_set, caches stay empty. Gates: host-mode byte-identical to pre-change baseline; host-vs-device A/B byte-identical; 9.6-9.8 fps. Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 222 ++++++++++++++++++++++++++++++++++++++++++++++++------- 1 file changed, 195 insertions(+), 27 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 5e5a123..fcd15ff 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -992,6 +992,66 @@ static void sam3_stage_feed(const sam3_transport& tr, struct ggml_tensor* src, ggml_backend_tensor_set(dst, tmp.data(), 0, ggml_nbytes(dst)); } +// Residency handle for values that today live in host vectors (increment 2). +// The EXISTING host vector remains the value and keeps its name/type/owner; +// the stagebuf only ADDS a persistent device twin beside it. HOST mode never +// touches the twin — the push is verbatim today's ggml_backend_tensor_set. +// DEVICE mode uploads the host value into the twin only when its content +// generation changes (constants: once), then D2D-copies twin -> fresh input. +struct sam3_stagebuf { + struct ggml_context* ctx = nullptr; // owns dev's metadata (no_alloc) + struct ggml_tensor* dev = nullptr; // persistent device twin + ggml_backend_buffer_t buf = nullptr; // owns dev's storage + size_t nbytes = 0; + uint64_t gen = (uint64_t)-1; // content generation last uploaded +}; + +static void sam3_stagebuf_release(sam3_stagebuf& sb) { + if (sb.buf) ggml_backend_buffer_free(sb.buf); + if (sb.ctx) ggml_free(sb.ctx); + sb.ctx = nullptr; sb.dev = nullptr; sb.buf = nullptr; + sb.nbytes = 0; sb.gen = (uint64_t)-1; +} + +// Feed a graph-input leaf from a host-side value. `gen` keys the twin's +// content: pass a counter that increments whenever the host value is +// recomputed (PE-cache generations); size changes (e.g. rope_k during the +// memory-bank ramp) re-shape the twin automatically. +static void sam3_stagebuf_push(const sam3_transport& tr, sam3_stagebuf& sb, + const void* host, size_t nbytes, + struct ggml_tensor* dst, uint64_t gen, const char* site) { + if (tr.kind != SAM3_TRANSPORT_DEVICE) { + ggml_backend_tensor_set(dst, host, 0, nbytes); // verbatim host path + return; + } + GGML_ASSERT(nbytes == ggml_nbytes(dst)); + if (sb.dev && (sb.nbytes != nbytes || !sam3_same_layout(sb.dev, dst))) + sam3_stagebuf_release(sb); + if (!sb.dev) { + struct ggml_init_params ip = { ggml_tensor_overhead() * 2, nullptr, true }; + sb.ctx = ggml_init(ip); + sb.dev = sb.ctx ? ggml_dup_tensor(sb.ctx, dst) : nullptr; + sb.buf = sb.dev ? ggml_backend_alloc_ctx_tensors(sb.ctx, tr.backend) : nullptr; + if (!sb.buf || ggml_backend_buffer_is_host(sb.dev->buffer)) { + static std::map warned; + if (!warned[site]) { + warned[site] = true; + fprintf(stderr, "sam3_stagebuf_push: %s: twin alloc failed — using host path\n", site); + } + sam3_stagebuf_release(sb); + ggml_backend_tensor_set(dst, host, 0, nbytes); + return; + } + sb.nbytes = nbytes; + sb.gen = (uint64_t)-1; + } + if (sb.gen != gen) { + ggml_backend_tensor_set(sb.dev, host, 0, nbytes); // one upload per content change + sb.gen = gen; + } + ggml_backend_tensor_copy(sb.dev, dst); // same-backend D2D +} + // Resolved once at model load, after the backend chain (Metal>CUDA>Vulkan>CPU). // Default host on every backend; device only when explicitly requested AND the // answered backend is CUDA — Metal can never pass this gate, so the mac cannot @@ -1072,6 +1132,19 @@ struct sam3_model { // stage-transport driver (host default everywhere; CUDA device opt-in) sam3_transport transport; + // Device-transport weight caches (increment 2): host copies of the small + // immutable weights that per-frame CPU logic re-reads from the device + // every frame (obj-ptr MLP + no_obj_ptr, perceiver latents, maskmem tpos). + // Populated ONCE at load and ONLY when transport == device; empty on the + // host driver, so every per-frame read stays verbatim today's sequence. + // Each cache replicates its consuming site's exact read (sam3_read_f32 + // for converted weights, raw tensor_get for no_obj_ptr). + struct { + std::vector objptr_w[3], objptr_b[3], no_obj_ptr; + std::vector perc_latents_1d, perc_latents_2d; + std::vector tpos; + } wcache; + // tensor lookup std::map tensors; @@ -1111,6 +1184,12 @@ struct sam3_state { float no_mask_emb_cache[256] = {}; std::vector dense_pe_cache; // [D * H * H] -- PE grid std::vector dense_nomask_cache; // [D * H * H] -- no-mask tiled + + // Stage-transport device twins for the PE-cache constants (increment 2). + // Inert on the host driver; content keyed on pe_cache_gen (bumped whenever + // sam3_populate_pe_cache actually recomputes). Freed in sam3_free_state. + uint64_t pe_cache_gen = 0; + sam3_stagebuf tw_sparse, tw_image_pe, tw_dense; }; /* @@ -1221,6 +1300,14 @@ struct sam3_tracker { // EdgeTAM-specific: RoPE for 16x16 grid (cross-attn K on perceiver 2D latents) std::vector cached_axial_cis_k16_reord; // [2, 128, 256] for 16x16 grid + + // Stage-transport device twins for the propagate constants (increment 2). + // Inert on the host driver; content keyed on pe_gen (bumped whenever + // sam3_ensure_tracker_pe_caches recomputes). rope_k's content is a pure + // function of (k16 cache, M_spatial) — the twin auto-reshapes on size + // change, so pe_gen keys it too. Freed with owned_buffers on reset. + uint64_t pe_gen = 0; + sam3_stagebuf tw_rope_q, tw_rope_k, tw_src_pos; }; // Resolve effective img_size / feat_size from state (which may override hp defaults). @@ -3627,6 +3714,41 @@ std::shared_ptr sam3_load_model(const sam3_params& params) { fprintf(stderr, "%s: visual-only model — skipping tokenizer\n", __func__); } + // Device-transport weight caches (increment 2): read each small immutable + // weight ONCE now instead of every frame. Device driver only — on the + // host driver the caches stay empty and the per-frame reads are verbatim. + if (model->transport.kind == SAM3_TRANSPORT_DEVICE) { + const int D = model->hparams.neck_dim; + const int MD = model->hparams.mem_out_dim; + for (int j = 0; j < 3; ++j) { + if (model->obj_ptr_proj_w[j] && model->obj_ptr_proj_b[j]) { + const int nel_w = (int)(model->obj_ptr_proj_w[j]->ne[0] * model->obj_ptr_proj_w[j]->ne[1]); + model->wcache.objptr_w[j].resize(nel_w); + sam3_read_f32(model->obj_ptr_proj_w[j], model->wcache.objptr_w[j].data(), nel_w); + model->wcache.objptr_b[j].resize(D); + sam3_read_f32(model->obj_ptr_proj_b[j], model->wcache.objptr_b[j].data(), D); + } + } + if (model->no_obj_ptr) { + model->wcache.no_obj_ptr.resize(D); + ggml_backend_tensor_get(model->no_obj_ptr, model->wcache.no_obj_ptr.data(), 0, D * sizeof(float)); + } + if (model->hparams.has_perceiver && model->perceiver.latents_1d && model->perceiver.latents_2d) { + const int n1 = (int)(model->perceiver.latents_1d->ne[0] * model->perceiver.latents_1d->ne[1]); + const int n2 = (int)(model->perceiver.latents_2d->ne[0] * model->perceiver.latents_2d->ne[1]); + model->wcache.perc_latents_1d.resize(n1); + sam3_read_f32(model->perceiver.latents_1d, model->wcache.perc_latents_1d.data(), n1); + model->wcache.perc_latents_2d.resize(n2); + sam3_read_f32(model->perceiver.latents_2d, model->wcache.perc_latents_2d.data(), n2); + } + if (model->mem_enc.tpos[0]) { + model->wcache.tpos.resize(MD * model->hparams.num_maskmem); + sam3_read_f32(model->mem_enc.tpos[0], model->wcache.tpos.data(), + MD * model->hparams.num_maskmem); + } + fprintf(stderr, "%s: device-transport weight caches populated\n", __func__); + } + fprintf(stderr, "%s: model loaded successfully\n", __func__); return model; } @@ -3706,6 +3828,10 @@ void sam3_state_set_orig_dims(sam3_state& state, int w, int h) { } void sam3_free_state(sam3_state& state) { + // Stage-transport twins (increment 2) + sam3_stagebuf_release(state.tw_sparse); + sam3_stagebuf_release(state.tw_image_pe); + sam3_stagebuf_release(state.tw_dense); if (state.galloc) { ggml_gallocr_free(state.galloc); state.galloc = nullptr; @@ -6292,11 +6418,20 @@ static bool edgetam_perceiver_forward( // ── Read learnable latent tokens from model weights ───────────────── // latents_1d: ggml shape [64, 256] → ne[0]=64(D), ne[1]=256(N) + // (device-transport weight cache when populated — increment 2) std::vector latents_1d_data(D * N_1d); - sam3_read_f32(perc.latents_1d, latents_1d_data.data(), D * N_1d); + if (!model.wcache.perc_latents_1d.empty()) + std::copy(model.wcache.perc_latents_1d.begin(), model.wcache.perc_latents_1d.end(), + latents_1d_data.begin()); + else + sam3_read_f32(perc.latents_1d, latents_1d_data.data(), D * N_1d); std::vector latents_2d_data(D * N_2d); - sam3_read_f32(perc.latents_2d, latents_2d_data.data(), D * N_2d); + if (!model.wcache.perc_latents_2d.empty()) + std::copy(model.wcache.perc_latents_2d.begin(), model.wcache.perc_latents_2d.end(), + latents_2d_data.begin()); + else + sam3_read_f32(perc.latents_2d, latents_2d_data.data(), D * N_2d); // ── Prepare 2D windowed features on CPU ───────────────────────────── // mem_features layout: [D, H*W] (ggml: ne[0]=D, ne[1]=H*W, stored as @@ -9798,8 +9933,11 @@ static sam3_prompt_data sam3_build_prompt_and_pos( pd.num_obj_ptr_tokens = 0; // Read maskmem_tpos_enc from model (one tensor [MD, 1, 1, 7]) + // (device-transport weight cache when populated — increment 2) std::vector tpos_all(MD * hp.num_maskmem); - if (model.mem_enc.tpos[0]) { + if (!model.wcache.tpos.empty()) { + std::copy(model.wcache.tpos.begin(), model.wcache.tpos.end(), tpos_all.begin()); + } else if (model.mem_enc.tpos[0]) { sam3_read_f32(model.mem_enc.tpos[0], tpos_all.data(), MD * hp.num_maskmem); } @@ -10015,19 +10153,25 @@ static void sam3_extract_obj_ptr_cpu( // λ = (obj_score > 0) ? 1.0 : 0.0 (hard threshold, not sigmoid) float lambda = (obj_score > 0.0f) ? 1.0f : 0.0f; - // Project token through MLP + // Project token through MLP (weights from the device-transport cache + // when populated — increment 2; host driver reads verbatim as today) std::vector h(D), tmp(D); std::copy(sam_token_data, sam_token_data + D, h.data()); for (int j = 0; j < 3; ++j) { auto* w = model.obj_ptr_proj_w[j]; auto* b = model.obj_ptr_proj_b[j]; int nel_w = (int)(w->ne[0] * w->ne[1]); - std::vector w_data(nel_w), b_data(D); - sam3_read_f32(w, w_data.data(), nel_w); - sam3_read_f32(b, b_data.data(), D); + const bool cached = !model.wcache.objptr_w[j].empty(); + std::vector w_data, b_data; + if (!cached) { + w_data.resize(nel_w); sam3_read_f32(w, w_data.data(), nel_w); + b_data.resize(D); sam3_read_f32(b, b_data.data(), D); + } + const float* W = cached ? model.wcache.objptr_w[j].data() : w_data.data(); + const float* B = cached ? model.wcache.objptr_b[j].data() : b_data.data(); for (int o = 0; o < D; ++o) { - float sum = b_data[o]; - for (int i = 0; i < D; ++i) sum += w_data[o * D + i] * h[i]; + float sum = B[o]; + for (int i = 0; i < D; ++i) sum += W[o * D + i] * h[i]; tmp[o] = (j < 2) ? std::max(0.0f, sum) : sum; } std::swap(h, tmp); @@ -10035,7 +10179,10 @@ static void sam3_extract_obj_ptr_cpu( // Blend: obj_ptr = λ * projected + (1-λ) * no_obj_ptr std::vector no_ptr(D); - ggml_backend_tensor_get(model.no_obj_ptr, no_ptr.data(), 0, D * sizeof(float)); + if (!model.wcache.no_obj_ptr.empty()) + std::copy(model.wcache.no_obj_ptr.begin(), model.wcache.no_obj_ptr.end(), no_ptr.begin()); + else + ggml_backend_tensor_get(model.no_obj_ptr, no_ptr.data(), 0, D * sizeof(float)); for (int i = 0; i < D; ++i) out_ptr[i] = lambda * h[i] + (1.0f - lambda) * no_ptr[i]; return; @@ -10043,7 +10190,10 @@ static void sam3_extract_obj_ptr_cpu( // SAM3 / default path: binary threshold if (obj_score <= 0.0f) { - ggml_backend_tensor_get(model.no_obj_ptr, out_ptr, 0, D * sizeof(float)); + if (!model.wcache.no_obj_ptr.empty()) + std::copy(model.wcache.no_obj_ptr.begin(), model.wcache.no_obj_ptr.end(), out_ptr); + else + ggml_backend_tensor_get(model.no_obj_ptr, out_ptr, 0, D * sizeof(float)); return; } @@ -10055,16 +10205,19 @@ static void sam3_extract_obj_ptr_cpu( auto* b = model.obj_ptr_proj_b[j]; int nel_w = (int)(w->ne[0] * w->ne[1]); - std::vector w_data(nel_w); - sam3_read_f32(w, w_data.data(), nel_w); - - std::vector b_data(D); - sam3_read_f32(b, b_data.data(), D); + const bool cached = !model.wcache.objptr_w[j].empty(); + std::vector w_data, b_data; + if (!cached) { + w_data.resize(nel_w); sam3_read_f32(w, w_data.data(), nel_w); + b_data.resize(D); sam3_read_f32(b, b_data.data(), D); + } + const float* W = cached ? model.wcache.objptr_w[j].data() : w_data.data(); + const float* B = cached ? model.wcache.objptr_b[j].data() : b_data.data(); for (int o = 0; o < D; ++o) { - float sum = b_data[o]; + float sum = B[o]; for (int i = 0; i < D; ++i) { - sum += w_data[o * D + i] * h[i]; + sum += W[o * D + i] * h[i]; } tmp[o] = (j < 2) ? std::max(0.0f, sum) : sum; } @@ -11088,6 +11241,7 @@ static void sam3_populate_pe_cache(sam3_state& state, const sam3_model& model) { } state.pe_cache_valid = true; + state.pe_cache_gen++; // invalidate stage-transport twins keyed on this cache fprintf(stderr, "%s: PE cache populated (%d embeddings, %.1f KB dense grids)\n", __func__, pe_nel, 2.0f * D * H * H * sizeof(float) / 1024.0f); } @@ -11901,6 +12055,7 @@ static void sam3_ensure_tracker_pe_caches(sam3_tracker& tracker, const sam3_hpar } tracker.pe_caches_valid = true; + tracker.pe_gen++; // invalidate stage-transport twins keyed on these caches SAM3_LOG(2, "%s: tracker PE caches populated (%.1f KB)\n", __func__, (tracker.cached_sinpe_256.size() + tracker.cached_sinpe_64.size() + tracker.cached_axial_cis_reord.size()) * sizeof(float) / 1024.0f); @@ -12243,9 +12398,13 @@ static sam3_prop_output sam3_propagate_single( // Upload prompt data ggml_backend_tensor_set(prompt_t, pd.prompt.data(), 0, pd.prompt.size() * sizeof(float)); ggml_backend_tensor_set(prompt_pos_t, pd.prompt_pos.data(), 0, pd.prompt_pos.size() * sizeof(float)); - ggml_backend_tensor_set(rope_q_t, rope_q_reord.data(), 0, rope_q_reord.size() * sizeof(float)); + sam3_stagebuf_push(model.transport, tracker.tw_rope_q, rope_q_reord.data(), + rope_q_reord.size() * sizeof(float), rope_q_t, + tracker.pe_gen, "rope_q"); if (rope_k_t && !rope_k_data.empty()) - ggml_backend_tensor_set(rope_k_t, rope_k_data.data(), 0, rope_k_data.size() * sizeof(float)); + sam3_stagebuf_push(model.transport, tracker.tw_rope_k, rope_k_data.data(), + rope_k_data.size() * sizeof(float), rope_k_t, + tracker.pe_gen, "rope_k"); } // Set default obj_score when pred_obj_scores=False (older SAM2 models) @@ -12256,16 +12415,20 @@ static sam3_prop_output sam3_propagate_single( // Upload src_pos (sinusoidal PE 256-dim) — ggml mem-attn path only. if (src_pos_t) - ggml_backend_tensor_set(src_pos_t, tracker.cached_sinpe_256.data(), 0, - tracker.cached_sinpe_256.size() * sizeof(float)); + sam3_stagebuf_push(model.transport, tracker.tw_src_pos, tracker.cached_sinpe_256.data(), + tracker.cached_sinpe_256.size() * sizeof(float), src_pos_t, + tracker.pe_gen, "src_pos"); // Upload not_a_point_embed, image_pe, dense_emb from state PE cache sam3_populate_pe_cache(state, model); - ggml_backend_tensor_set(sparse_in, state.not_a_point_cache, 0, D * sizeof(float)); - ggml_backend_tensor_set(image_pe, state.dense_pe_cache.data(), 0, - state.dense_pe_cache.size() * sizeof(float)); - ggml_backend_tensor_set(dense_emb, state.dense_nomask_cache.data(), 0, - state.dense_nomask_cache.size() * sizeof(float)); + sam3_stagebuf_push(model.transport, state.tw_sparse, state.not_a_point_cache, + D * sizeof(float), sparse_in, state.pe_cache_gen, "prop_sparse"); + sam3_stagebuf_push(model.transport, state.tw_image_pe, state.dense_pe_cache.data(), + state.dense_pe_cache.size() * sizeof(float), image_pe, + state.pe_cache_gen, "prop_pe"); + sam3_stagebuf_push(model.transport, state.tw_dense, state.dense_nomask_cache.data(), + state.dense_nomask_cache.size() * sizeof(float), dense_emb, + state.pe_cache_gen, "prop_dense"); // Copy tracker features from state to fresh input tensors (transport // driver: host = today's round-trip, device = same-backend D2D). @@ -13227,6 +13390,11 @@ void sam3_tracker_reset(sam3_tracker& tracker) { for (auto* b : tracker.owned_buffers) if (b) ggml_backend_buffer_free(b); tracker.owned_buffers.clear(); + // Stage-transport twins (increment 2): release with the other device + // resources; they repopulate lazily on the next device-mode propagate. + sam3_stagebuf_release(tracker.tw_rope_q); + sam3_stagebuf_release(tracker.tw_rope_k); + sam3_stagebuf_release(tracker.tw_src_pos); if (tracker.ctx) { ggml_free(tracker.ctx); tracker.ctx = nullptr; From 092b3136d13023d55ecc723b709e31ef68b71775 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 18:15:03 -0600 Subject: [PATCH 09/16] =?UTF-8?q?cuda:=20stage-transport=20increment=203?= =?UTF-8?q?=20=E2=80=94=20fixed=20bank=20ring=20+=20device=20prompt=20asse?= =?UTF-8?q?mbly?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit report/CUDA-RESIDENCY-DESIGN.md §5 increment 3. Ring: bank slots + obj pointers now recycle their device tensors on eviction (sam3_acquire_slot_pair / recycle free-lists) instead of leaking buffers until reset (~263 KB VRAM per credible frame) and accreting tensor metadata in tracker.ctx (exhaustion at ~1365 credible frames). Whole-bank instance eviction recycles too. Host-driver bytes unchanged (same tensor_set order); only the allocation pattern differs. Device prompt assembly (CUDA device driver + EdgeTAM + full steady-state capacity only; ramp frames keep the verbatim host path — no padding): the prompt's spatial region is concatenated IN-GRAPH from fresh slot leaves fed by D2D from the bank ring, with the tpos broadcast add done on device (single fp32 add — bit-exact vs the CPU add). The pointer region is built by the exact host code path (sam3_build_prompt_and_pos with empty slots) and uploaded (~34 KB). Removes the per-frame slot downloads (1.8 MB D2H), the 917 KB prompt/pos uploads, and 16 pointer D2H gets + their syncs. Obj-ptr tpos-proj weights added to the load-time cache; host mirror of pointers maintained alongside ptr_banks. Gates: host-mode byte-identical to pre-change baseline; host-vs-device A/B byte-identical at 40 frames AND a 300-frame soak (19.6 MB, ring wraparound; "device prompt assembly engaged (7 slots x 512 tokens + 64 ptr tokens)" asserted in stderr). Steady-state device fps 10.46 (from 9.6). Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 377 +++++++++++++++++++++++++++++++++++++++++++++---------- 1 file changed, 312 insertions(+), 65 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index fcd15ff..8aae358 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -1143,6 +1143,7 @@ struct sam3_model { std::vector objptr_w[3], objptr_b[3], no_obj_ptr; std::vector perc_latents_1d, perc_latents_2d; std::vector tpos; + std::vector tpos_proj_w, tpos_proj_b; // obj_ptr_tpos_proj (inc 3) } wcache; // tensor lookup @@ -1227,6 +1228,12 @@ struct sam3_masklet { struct sam3_memory_slot { struct ggml_tensor* spatial_feats = nullptr; // [64, 72, 72] struct ggml_tensor* spatial_pe = nullptr; // [64, 72, 72] + // Increment 3 (fixed ring): the backend buffer holding BOTH tensors above, + // recorded so eviction can recycle the pair instead of leaking it until + // reset (~263 KB VRAM per credible frame) and so tensor metadata stops + // accreting in tracker.ctx. Null for slots stored by paths that predate + // the ring (CoreML store) — those simply aren't recycled. + ggml_backend_buffer_t buf = nullptr; int frame_index = -1; bool is_cond_frame = false; // RFD 0011 U3: per-slot bbox + quality recorded at the time this memory @@ -1308,8 +1315,68 @@ struct sam3_tracker { // change, so pe_gen keys it too. Freed with owned_buffers on reset. uint64_t pe_gen = 0; sam3_stagebuf tw_rope_q, tw_rope_k, tw_src_pos; + + // Increment 3: recycling free-lists for bank-slot / obj-ptr device tensors. + // Evicted slots return their tensor pair + buffer here; the next store + // reuses them in place (fixes the eviction VRAM leak, stops per-frame + // cudaMalloc churn, and keeps leaf data pointers frame-invariant for + // CUDA-graph capture). Buffers are ALSO in owned_buffers, which remains + // the single point of freeing on reset — these lists never free. + struct sam3_slot_rec { struct ggml_tensor* feats; struct ggml_tensor* pe; ggml_backend_buffer_t buf; }; + struct sam3_ptr_rec { struct ggml_tensor* pt; ggml_backend_buffer_t buf; }; + std::vector slot_free; + std::vector ptr_free; + + // Increment 3: host mirror of stored object pointers, strictly parallel to + // ptr_banks (same push/trim/erase). The device prompt assembly reads these + // instead of doing 16 per-frame 1 KB D2H gets. The values pass through the + // host at store time anyway, so this costs one memcpy per credible frame. + std::map>> ptr_host; }; +// Increment 3: acquire a bank-slot tensor pair (feats + pe sharing one backend +// buffer), reusing a recycled pair when the dims match; otherwise allocate and +// track the new buffer in owned_buffers exactly as the old per-frame path did. +// Returns false only on allocation failure. +static bool sam3_acquire_slot_pair(sam3_tracker& tracker, const sam3_model& model, + int64_t ne0, int64_t ne1, int64_t ne2, + struct ggml_tensor** out_feats, + struct ggml_tensor** out_pe, + ggml_backend_buffer_t* out_buf) { + for (size_t i = 0; i < tracker.slot_free.size(); ++i) { + auto& r = tracker.slot_free[i]; + if (r.feats && r.feats->ne[0] == ne0 && r.feats->ne[1] == ne1 && + r.feats->ne[2] == ne2 && r.feats->ne[3] == 1) { + *out_feats = r.feats; *out_pe = r.pe; *out_buf = r.buf; + tracker.slot_free.erase(tracker.slot_free.begin() + i); + return true; + } + } + if (!tracker.ctx) { + struct ggml_init_params tp = {ggml_tensor_overhead() * 4096, nullptr, true}; + tracker.ctx = ggml_init(tp); + } + const size_t nbytes = (size_t)ne0 * ne1 * ne2 * sizeof(float); + auto* feats = ggml_new_tensor_3d(tracker.ctx, GGML_TYPE_F32, ne0, ne1, ne2); + auto* pe = ggml_new_tensor_3d(tracker.ctx, GGML_TYPE_F32, ne0, ne1, ne2); + auto* buf = ggml_backend_alloc_buffer(model.backend, 2 * nbytes + 512); + if (!feats || !pe || !buf) { + if (buf) ggml_backend_buffer_free(buf); + return false; + } + struct ggml_tallocr ta = ggml_tallocr_new(buf); + ggml_tallocr_alloc(&ta, feats); + ggml_tallocr_alloc(&ta, pe); + tracker.owned_buffers.push_back(buf); + *out_feats = feats; *out_pe = pe; *out_buf = buf; + return true; +} + +static void sam3_recycle_slot(sam3_tracker& tracker, const sam3_memory_slot& slot) { + if (slot.buf && slot.spatial_feats && slot.spatial_pe) + tracker.slot_free.push_back({slot.spatial_feats, slot.spatial_pe, slot.buf}); +} + // Resolve effective img_size / feat_size from state (which may override hp defaults). static int sam3_eff_img_size(const sam3_state& s, const sam3_hparams& hp) { return (s.encode_img_size > 0) ? s.encode_img_size : hp.img_size; @@ -3746,6 +3813,12 @@ std::shared_ptr sam3_load_model(const sam3_params& params) { sam3_read_f32(model->mem_enc.tpos[0], model->wcache.tpos.data(), MD * model->hparams.num_maskmem); } + if (model->obj_ptr_tpos_w && model->obj_ptr_tpos_b) { + model->wcache.tpos_proj_w.resize(D * MD); + sam3_read_f32(model->obj_ptr_tpos_w, model->wcache.tpos_proj_w.data(), D * MD); + model->wcache.tpos_proj_b.resize(MD); + sam3_read_f32(model->obj_ptr_tpos_b, model->wcache.tpos_proj_b.data(), MD); + } fprintf(stderr, "%s: device-transport weight caches populated\n", __func__); } @@ -9964,8 +10037,12 @@ static sam3_prompt_data sam3_build_prompt_and_pos( const int split = D / MD; // 4 // Read obj_ptr_tpos_proj weights for CPU-side matmul + // (device-transport weight cache when populated — increment 3) std::vector tpos_w(D * MD), tpos_b(MD); - if (model.obj_ptr_tpos_w) { + if (!model.wcache.tpos_proj_w.empty()) { + std::copy(model.wcache.tpos_proj_w.begin(), model.wcache.tpos_proj_w.end(), tpos_w.begin()); + std::copy(model.wcache.tpos_proj_b.begin(), model.wcache.tpos_proj_b.end(), tpos_b.begin()); + } else if (model.obj_ptr_tpos_w) { sam3_read_f32(model.obj_ptr_tpos_w, tpos_w.data(), D * MD); sam3_read_f32(model.obj_ptr_tpos_b, tpos_b.data(), MD); } @@ -12086,39 +12163,71 @@ static sam3_prop_output sam3_propagate_single( const int N_per_slot = use_perceiver ? (hp.perceiver_n_latents_1d + hp.perceiver_n_latents_2d) : N; int n_sel = (int)sel.size(); + int P = std::min((int)ptr_bank.size(), hp.max_obj_ptrs); + + // Increment 3: device prompt assembly. At full steady-state capacity on + // the CUDA device driver (EdgeTAM only) the prompt's spatial region is + // assembled ON DEVICE — in-graph concat of D2D-fed slot leaves plus the + // bit-exact tpos broadcast add — so the per-frame slot downloads and the + // ~917 KB prompt/prompt_pos uploads disappear. Ramp frames (bank not yet + // full) take the verbatim host path: no padding — padding would append + // real attention tokens and change softmax denominators (design §3.4-3). + bool device_prompt = + model.transport.kind == SAM3_TRANSPORT_DEVICE && + use_perceiver && hp.is_edgetam() && !model.wcache.tpos.empty() && + n_sel == hp.num_maskmem && n_sel <= 16 && + P == hp.max_obj_ptrs && (int)ptr_bank.size() == P; + if (device_prompt) { + for (int s = 0; s < n_sel; ++s) + if (!mem_bank[sel[s]].spatial_feats || !mem_bank[sel[s]].spatial_pe) { + device_prompt = false; + break; + } + auto hm = tracker.ptr_host.find(masklet.instance_id); + if (hm == tracker.ptr_host.end() || (int)hm->second.size() != (int)ptr_bank.size()) + device_prompt = false; + } + std::vector> slot_feats(n_sel), slot_pes(n_sel); std::vector spatial_tpos(n_sel, 1); // default t_pos=1 for non-cond for (int s = 0; s < n_sel; ++s) { - slot_feats[s].resize(MD * N_per_slot); - ggml_backend_tensor_get(mem_bank[sel[s]].spatial_feats, - slot_feats[s].data(), 0, MD * N_per_slot * sizeof(float)); - slot_pes[s].resize(MD * N_per_slot); - if (mem_bank[sel[s]].spatial_pe) { - ggml_backend_tensor_get(mem_bank[sel[s]].spatial_pe, - slot_pes[s].data(), 0, MD * N_per_slot * sizeof(float)); - } else { - sam3_ensure_tracker_pe_caches(tracker, hp, H); - if (use_perceiver) { - // For perceiver: zeros for 1D tokens, sinusoidal for 2D tokens - const int N_1d = hp.perceiver_n_latents_1d; - const int N_2d = hp.perceiver_n_latents_2d; - memset(slot_pes[s].data(), 0, MD * N_1d * sizeof(float)); - auto pe_2d = sam3_sinusoidal_pe_2d(16, 16, MD); - memcpy(slot_pes[s].data() + MD * N_1d, pe_2d.data(), MD * N_2d * sizeof(float)); + if (!device_prompt) { + slot_feats[s].resize(MD * N_per_slot); + ggml_backend_tensor_get(mem_bank[sel[s]].spatial_feats, + slot_feats[s].data(), 0, MD * N_per_slot * sizeof(float)); + slot_pes[s].resize(MD * N_per_slot); + if (mem_bank[sel[s]].spatial_pe) { + ggml_backend_tensor_get(mem_bank[sel[s]].spatial_pe, + slot_pes[s].data(), 0, MD * N_per_slot * sizeof(float)); } else { - slot_pes[s] = tracker.cached_sinpe_64; + sam3_ensure_tracker_pe_caches(tracker, hp, H); + if (use_perceiver) { + // For perceiver: zeros for 1D tokens, sinusoidal for 2D tokens + const int N_1d = hp.perceiver_n_latents_1d; + const int N_2d = hp.perceiver_n_latents_2d; + memset(slot_pes[s].data(), 0, MD * N_1d * sizeof(float)); + auto pe_2d = sam3_sinusoidal_pe_2d(16, 16, MD); + memcpy(slot_pes[s].data() + MD * N_1d, pe_2d.data(), MD * N_2d * sizeof(float)); + } else { + slot_pes[s] = tracker.cached_sinpe_64; + } } } spatial_tpos[s] = mem_bank[sel[s]].is_cond_frame ? 0 : (n_sel - s); } - int P = std::min((int)ptr_bank.size(), hp.max_obj_ptrs); std::vector> obj_ptrs(P); std::vector ptr_tpos(P); int cur_frame = tracker.frame_index; for (int p = 0; p < P; ++p) { - obj_ptrs[p].resize(D); - ggml_backend_tensor_get(ptr_bank[p].second, obj_ptrs[p].data(), 0, D * sizeof(float)); + if (device_prompt) { + // host mirror maintained by sam3_store_obj_ptr — same values the + // D2H get would return, without the 16 per-frame syncs + obj_ptrs[p] = tracker.ptr_host[masklet.instance_id][p]; + } else { + obj_ptrs[p].resize(D); + ggml_backend_tensor_get(ptr_bank[p].second, obj_ptrs[p].data(), 0, D * sizeof(float)); + } // Use actual frame distance (matches Python: abs(frame_idx - t)) ptr_tpos[p] = std::abs(cur_frame - ptr_bank[p].first); if (ptr_tpos[p] < 1) ptr_tpos[p] = 1; // minimum distance of 1 @@ -12147,7 +12256,35 @@ static sam3_prop_output sam3_propagate_single( } } #endif - auto pd = sam3_build_prompt_and_pos(model, slot_feats, slot_pes, spatial_tpos, obj_ptrs, ptr_tpos, H); + // Device prompt assembly builds ONLY the pointer region on the host — the + // exact code path the host driver runs, so its bytes are identical by + // construction — and accounts for the device-assembled spatial region in + // the M fields (rope_k sizing and the graph read them). + sam3_prompt_data pd; + if (device_prompt) { + const std::vector> no_slots; + const std::vector no_tpos; + pd = sam3_build_prompt_and_pos(model, no_slots, no_slots, no_tpos, obj_ptrs, ptr_tpos, H); + pd.M_spatial = n_sel * N_per_slot; + pd.M_total = pd.M_spatial + pd.num_obj_ptr_tokens; + } else { + pd = sam3_build_prompt_and_pos(model, slot_feats, slot_pes, spatial_tpos, obj_ptrs, ptr_tpos, H); + } + + // tpos rows for the selected slots (device path; ~1.8 KB upload). Matches + // the host semantics exactly: rows outside [0, num_maskmem) add nothing. + std::vector tpos_sel; + if (device_prompt) { + tpos_sel.assign((size_t)MD * n_sel, 0.0f); + for (int s = 0; s < n_sel; ++s) { + const int tpos_idx = spatial_tpos[s]; + if (tpos_idx >= 0 && tpos_idx < hp.num_maskmem) { + const int enc_idx = hp.num_maskmem - tpos_idx - 1; + memcpy(&tpos_sel[(size_t)s * MD], &model.wcache.tpos[(size_t)enc_idx * MD], + MD * sizeof(float)); + } + } + } // ── RoPE frequencies (cached) ────────────────────────────────────── sam3_ensure_tracker_pe_caches(tracker, hp, H); @@ -12291,6 +12428,12 @@ static sam3_prop_output sam3_propagate_single( struct ggml_tensor* prompt_t = nullptr, * prompt_pos_t = nullptr; struct ggml_tensor* rope_q_t = nullptr, * rope_k_t = nullptr; struct ggml_tensor* cond_spatial = nullptr; + // Device prompt assembly leaves (increment 3; null on the host path) + struct ggml_tensor* slot_feat_leaf[16] = {}; + struct ggml_tensor* slot_pe_leaf[16] = {}; + struct ggml_tensor* tpos_leaf = nullptr; + struct ggml_tensor* ptr_leaf = nullptr, * ptr_pos_leaf = nullptr; + struct ggml_tensor* prompt_node = nullptr, * prompt_pos_node = nullptr; if (use_cml_ma) { // CoreML produced the attended features; inject them as a decoder input. @@ -12323,11 +12466,43 @@ static sam3_prop_output sam3_propagate_single( src_pos_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N, 1); ggml_set_name(src_pos_t, "src_pos"); ggml_set_input(src_pos_t); - // Prompt and prompt_pos - prompt_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, MD, pd.M_total, 1); - ggml_set_name(prompt_t, "prompt"); ggml_set_input(prompt_t); - prompt_pos_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, MD, pd.M_total, 1); - ggml_set_name(prompt_pos_t, "prompt_pos"); ggml_set_input(prompt_pos_t); + // Prompt and prompt_pos. Device prompt assembly (increment 3): the + // spatial region is concatenated IN-GRAPH from fresh slot leaves fed + // by D2D (bank slots never round-trip the bus); the pointer region is + // a small host upload built by the exact host code path. All operands + // are fresh leaves — no state-tensor ancestry (CLAUDE.md rule 2). + if (device_prompt) { + struct ggml_tensor* pr = nullptr; + struct ggml_tensor* pp = nullptr; + tpos_leaf = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, MD, n_sel); + ggml_set_name(tpos_leaf, "prompt_tpos"); ggml_set_input(tpos_leaf); + for (int s = 0; s < n_sel; ++s) { + slot_feat_leaf[s] = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, MD, N_per_slot); + ggml_set_name(slot_feat_leaf[s], "prompt_slot_feat"); ggml_set_input(slot_feat_leaf[s]); + slot_pe_leaf[s] = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, MD, N_per_slot); + ggml_set_name(slot_pe_leaf[s], "prompt_slot_pe"); ggml_set_input(slot_pe_leaf[s]); + // pos_s = slot_pe + tpos[enc_idx] broadcast over the slot's + // tokens — a single element-wise fp32 add, bit-identical to + // the host's CPU add of the same operands. + auto* tp_s = ggml_cont(ctx0, ggml_view_2d(ctx0, tpos_leaf, MD, 1, + tpos_leaf->nb[1], + (size_t)s * tpos_leaf->nb[1])); + auto* pos_s = ggml_add(ctx0, slot_pe_leaf[s], tp_s); + pr = pr ? ggml_concat(ctx0, pr, slot_feat_leaf[s], 1) : slot_feat_leaf[s]; + pp = pp ? ggml_concat(ctx0, pp, pos_s, 1) : pos_s; + } + ptr_leaf = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, MD, pd.num_obj_ptr_tokens); + ggml_set_name(ptr_leaf, "prompt_ptr"); ggml_set_input(ptr_leaf); + ptr_pos_leaf = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, MD, pd.num_obj_ptr_tokens); + ggml_set_name(ptr_pos_leaf, "prompt_ptr_pos"); ggml_set_input(ptr_pos_leaf); + prompt_node = ggml_concat(ctx0, pr, ptr_leaf, 1); + prompt_pos_node = ggml_concat(ctx0, pp, ptr_pos_leaf, 1); + } else { + prompt_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, MD, pd.M_total, 1); + ggml_set_name(prompt_t, "prompt"); ggml_set_input(prompt_t); + prompt_pos_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, MD, pd.M_total, 1); + ggml_set_name(prompt_pos_t, "prompt_pos"); ggml_set_input(prompt_pos_t); + } // RoPE frequencies rope_q_t = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, 2, half_d, N); @@ -12338,7 +12513,8 @@ static sam3_prop_output sam3_propagate_single( } auto* conditioned = sam3_build_mem_attn_graph(ctx0, model, curr, src_pos_t, - prompt_t, prompt_pos_t, + device_prompt ? prompt_node : prompt_t, + device_prompt ? prompt_pos_node : prompt_pos_t, rope_q_t, rope_k_t, pd.num_obj_ptr_tokens); cond_spatial = ggml_reshape_4d(ctx0, conditioned, D, H, H, 1); @@ -12396,8 +12572,29 @@ static sam3_prop_output sam3_propagate_single( ggml_backend_tensor_set(cond_spatial, cml_cond.data(), 0, cml_cond.size() * sizeof(float)); } else { // Upload prompt data - ggml_backend_tensor_set(prompt_t, pd.prompt.data(), 0, pd.prompt.size() * sizeof(float)); - ggml_backend_tensor_set(prompt_pos_t, pd.prompt_pos.data(), 0, pd.prompt_pos.size() * sizeof(float)); + if (prompt_t) { + ggml_backend_tensor_set(prompt_t, pd.prompt.data(), 0, pd.prompt.size() * sizeof(float)); + ggml_backend_tensor_set(prompt_pos_t, pd.prompt_pos.data(), 0, pd.prompt_pos.size() * sizeof(float)); + } else if (device_prompt) { + // Increment 3: spatial region D2D from the bank slots; the small + // tpos + pointer blocks are the only prompt uploads left (~34 KB). + static bool logged = false; + if (!logged) { + logged = true; + fprintf(stderr, "sam3_propagate_single: device prompt assembly engaged " + "(%d slots x %d tokens + %d ptr tokens)\n", + n_sel, N_per_slot, pd.num_obj_ptr_tokens); + } + ggml_backend_tensor_set(tpos_leaf, tpos_sel.data(), 0, tpos_sel.size() * sizeof(float)); + ggml_backend_tensor_set(ptr_leaf, pd.prompt.data(), 0, pd.prompt.size() * sizeof(float)); + ggml_backend_tensor_set(ptr_pos_leaf, pd.prompt_pos.data(), 0, pd.prompt_pos.size() * sizeof(float)); + for (int s = 0; s < n_sel; ++s) { + sam3_stage_feed(model.transport, mem_bank[sel[s]].spatial_feats, + slot_feat_leaf[s], "prompt_slot_feat"); + sam3_stage_feed(model.transport, mem_bank[sel[s]].spatial_pe, + slot_pe_leaf[s], "prompt_slot_pe"); + } + } sam3_stagebuf_push(model.transport, tracker.tw_rope_q, rope_q_reord.data(), rope_q_reord.size() * sizeof(float), rope_q_t, tracker.pe_gen, "rope_q"); @@ -12627,6 +12824,19 @@ static void sam3_update_tracker(sam3_tracker& tracker, int frame_idx) { } for (auto it = tracker.masklets.begin(); it != tracker.masklets.end();) { if (frame_idx - it->last_seen > tracker.params.max_keep_alive) { + // Increment 3: recycle the evicted instance's slot/ptr tensors so + // whole-bank eviction doesn't leak until reset either. + { + auto mb = tracker.mem_banks.find(it->instance_id); + if (mb != tracker.mem_banks.end()) + for (const auto& s : mb->second) sam3_recycle_slot(tracker, s); + auto pb = tracker.ptr_banks.find(it->instance_id); + if (pb != tracker.ptr_banks.end()) + for (const auto& p : pb->second) + if (p.second && p.second->buffer) + tracker.ptr_free.push_back({p.second, nullptr}); + tracker.ptr_host.erase(it->instance_id); + } tracker.mem_banks.erase(it->instance_id); tracker.ptr_banks.erase(it->instance_id); // RFD 0011 U3: evict the per-instance motion model alongside its @@ -12911,25 +13121,25 @@ static bool sam3_encode_memory( } } #endif - // Store perceiver output: [MD=64, N_perc] (N_1d + N_2d latents) - auto* st = ggml_new_tensor_2d(tracker.ctx, GGML_TYPE_F32, MD, N_perc); - auto* sb = ggml_backend_alloc_buffer(model.backend, MD * N_perc * sizeof(float)); - struct ggml_tallocr ta_perc = ggml_tallocr_new(sb); - ggml_tallocr_alloc(&ta_perc, st); - tracker.owned_buffers.push_back(sb); + // Store perceiver output + PE (increment 3: fixed-ring pair — reuses a + // recycled tensor pair in place instead of allocating per frame) + struct ggml_tensor* st = nullptr; + struct ggml_tensor* spe = nullptr; + ggml_backend_buffer_t slot_buf = nullptr; + if (!sam3_acquire_slot_pair(tracker, model, MD, N_perc, 1, &st, &spe, &slot_buf)) { + fprintf(stderr, "%s: slot pair alloc failed\n", __func__); + SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); + ggml_gallocr_free(ga); + ggml_free(ctx0); + return false; + } ggml_backend_tensor_set(st, perc_latents.data(), 0, MD * N_perc * sizeof(float)); - - // Store perceiver PE: [MD=64, 512] (first 256 zeros, last 256 sinusoidal) - auto* spe = ggml_new_tensor_2d(tracker.ctx, GGML_TYPE_F32, MD, N_perc); - auto* speb = ggml_backend_alloc_buffer(model.backend, MD * N_perc * sizeof(float)); - struct ggml_tallocr ta_perc2 = ggml_tallocr_new(speb); - ggml_tallocr_alloc(&ta_perc2, spe); - tracker.owned_buffers.push_back(speb); ggml_backend_tensor_set(spe, perc_pos.data(), 0, MD * N_perc * sizeof(float)); sam3_memory_slot slot; slot.spatial_feats = st; slot.spatial_pe = spe; + slot.buf = slot_buf; slot.frame_index = frame_idx; slot.is_cond_frame = is_cond; auto& bk = tracker.mem_banks[inst_id]; @@ -12937,15 +13147,20 @@ static bool sam3_encode_memory( // RFD 0011 U3: trim to the memory POOL cap (>= num_maskmem) rather than // exactly num_maskmem, so sam3_select_memory_frames has a real choice of // motion-consistent slots. Still drop the oldest non-cond slot first. + // Increment 3: evicted slots recycle their tensors instead of leaking. while ((int)bk.size() > sam3_mem_pool_cap(hp.num_maskmem)) { bool removed = false; for (auto it = bk.begin(); it != bk.end(); ++it) if (!it->is_cond_frame) { + sam3_recycle_slot(tracker, *it); bk.erase(it); removed = true; break; } - if (!removed) bk.erase(bk.begin() + 1); + if (!removed) { + sam3_recycle_slot(tracker, bk[1]); + bk.erase(bk.begin() + 1); + } } SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); ggml_gallocr_free(ga); @@ -12954,27 +13169,28 @@ static bool sam3_encode_memory( } // ── Standard path (SAM2/SAM3): store raw [MD, H, H] features ──────── - // Store spatial features - auto* st = ggml_new_tensor_4d(tracker.ctx, GGML_TYPE_F32, MD, H, H, 1); - auto* sb = ggml_backend_alloc_buffer(model.backend, MD * H * H * sizeof(float)); - struct ggml_tallocr ta = ggml_tallocr_new(sb); - ggml_tallocr_alloc(&ta, st); - tracker.owned_buffers.push_back(sb); + // (increment 3: fixed-ring pair — reuses a recycled tensor pair in place) + struct ggml_tensor* st = nullptr; + struct ggml_tensor* spe = nullptr; + ggml_backend_buffer_t slot_buf = nullptr; + if (!sam3_acquire_slot_pair(tracker, model, MD, H, H, &st, &spe, &slot_buf)) { + fprintf(stderr, "%s: slot pair alloc failed\n", __func__); + SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); + ggml_gallocr_free(ga); + ggml_free(ctx0); + return false; + } ggml_backend_tensor_set(st, md.data(), 0, md.size() * sizeof(float)); // Compute and store sinusoidal spatial PE sam3_ensure_tracker_pe_caches(tracker, hp, H); const auto& pe_data = tracker.cached_sinpe_64; - auto* spe = ggml_new_tensor_4d(tracker.ctx, GGML_TYPE_F32, MD, H, H, 1); - auto* speb = ggml_backend_alloc_buffer(model.backend, MD * H * H * sizeof(float)); - struct ggml_tallocr ta2 = ggml_tallocr_new(speb); - ggml_tallocr_alloc(&ta2, spe); - tracker.owned_buffers.push_back(speb); ggml_backend_tensor_set(spe, pe_data.data(), 0, pe_data.size() * sizeof(float)); sam3_memory_slot slot; slot.spatial_feats = st; slot.spatial_pe = spe; + slot.buf = slot_buf; slot.frame_index = frame_idx; slot.is_cond_frame = is_cond; auto& bk = tracker.mem_banks[inst_id]; @@ -12982,15 +13198,20 @@ static bool sam3_encode_memory( // RFD 0011 U3: trim to the memory POOL cap (see perceiver path above and // sam3_mem_pool_cap). Selection of the attended num_maskmem subset is the // motion-aware sam3_select_memory_frames; eviction here only bounds storage. + // Increment 3: evicted slots recycle their tensors instead of leaking. while ((int)bk.size() > sam3_mem_pool_cap(hp.num_maskmem)) { bool removed = false; for (auto it = bk.begin(); it != bk.end(); ++it) if (!it->is_cond_frame) { + sam3_recycle_slot(tracker, *it); bk.erase(it); removed = true; break; } - if (!removed) bk.erase(bk.begin() + 1); + if (!removed) { + sam3_recycle_slot(tracker, bk[1]); + bk.erase(bk.begin() + 1); + } } SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); ggml_gallocr_free(ga); @@ -13002,19 +13223,40 @@ static void sam3_store_obj_ptr( sam3_tracker& tracker, const sam3_model& model, int inst_id, const float* pd, int frame_idx) { const int D = model.hparams.neck_dim; - if (!tracker.ctx) { - struct ggml_init_params tp = {ggml_tensor_overhead() * 4096, nullptr, true}; - tracker.ctx = ggml_init(tp); + // Increment 3: reuse a recycled pointer tensor when available (fixes the + // 1 KB/credible-frame buffer leak and metadata accretion in tracker.ctx). + struct ggml_tensor* pt = nullptr; + for (size_t i = 0; i < tracker.ptr_free.size(); ++i) { + auto& r = tracker.ptr_free[i]; + if (r.pt && r.pt->ne[0] == D && r.pt->ne[1] == 1) { + pt = r.pt; + tracker.ptr_free.erase(tracker.ptr_free.begin() + i); + break; + } + } + if (!pt) { + if (!tracker.ctx) { + struct ggml_init_params tp = {ggml_tensor_overhead() * 4096, nullptr, true}; + tracker.ctx = ggml_init(tp); + } + pt = ggml_new_tensor_2d(tracker.ctx, GGML_TYPE_F32, D, 1); + auto* pb = ggml_backend_alloc_buffer(model.backend, D * sizeof(float)); + struct ggml_tallocr ta = ggml_tallocr_new(pb); + ggml_tallocr_alloc(&ta, pt); + tracker.owned_buffers.push_back(pb); } - auto* pt = ggml_new_tensor_2d(tracker.ctx, GGML_TYPE_F32, D, 1); - auto* pb = ggml_backend_alloc_buffer(model.backend, D * sizeof(float)); - struct ggml_tallocr ta = ggml_tallocr_new(pb); - ggml_tallocr_alloc(&ta, pt); - tracker.owned_buffers.push_back(pb); ggml_backend_tensor_set(pt, pd, 0, D * sizeof(float)); auto& bk = tracker.ptr_banks[inst_id]; bk.push_back({frame_idx, pt}); - while ((int)bk.size() > model.hparams.max_obj_ptrs) bk.erase(bk.begin()); + // Host mirror (increment 3): strictly parallel to the bank — the device + // prompt assembly reads pointer values from here instead of D2H gets. + auto& hm = tracker.ptr_host[inst_id]; + hm.emplace_back(pd, pd + D); + while ((int)bk.size() > model.hparams.max_obj_ptrs) { + if (bk.front().second) tracker.ptr_free.push_back({bk.front().second, nullptr}); + bk.erase(bk.begin()); + if (!hm.empty()) hm.erase(hm.begin()); + } } sam3_tracker_ptr sam3_create_tracker(const sam3_model& model, @@ -13395,6 +13637,11 @@ void sam3_tracker_reset(sam3_tracker& tracker) { sam3_stagebuf_release(tracker.tw_rope_q); sam3_stagebuf_release(tracker.tw_rope_k); sam3_stagebuf_release(tracker.tw_src_pos); + // Increment 3: the recycling lists hold pointers into owned_buffers' + // storage (freed just above) — clear them, never free through them. + tracker.slot_free.clear(); + tracker.ptr_free.clear(); + tracker.ptr_host.clear(); if (tracker.ctx) { ggml_free(tracker.ctx); tracker.ctx = nullptr; From e1b5665a195b7867392245d7af2108847dad3426 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 18:21:24 -0600 Subject: [PATCH 10/16] =?UTF-8?q?cuda:=20stage-transport=20increment=204?= =?UTF-8?q?=20=E2=80=94=20device=20memenc=E2=86=92perceiver=20chain?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit report/CUDA-RESIDENCY-DESIGN.md §5 increment 4 (device driver + EdgeTAM only; SAM2.1 keeps the verbatim host path — its no_obj_embed_spatial branch mutates the memenc output on the host). The perceiver now feeds its two sub-graphs by same-backend D2D from the LIVE memenc output tensor (mo) while its graph memory is still allocated — the 1 MB per-frame download of md disappears. The 1D path uses the 4D-leaf + in-graph-reshape pattern; the 2D path performs the window partition IN-GRAPH (reshape/permute/cont/reshape — a pure byte gather, bit-identical to edgetam_window_partition_cpu, verified by the byte-diff gate). Perceiver pos (cached_sinpe_64) and latents ride persistent twins. Gates: host-mode byte-identical to pre-change baseline; host-vs-device A/B byte-identical at 40 frames and a 300-frame soak; 10.0 fps smoke / 10.45 fps steady-state. Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 137 ++++++++++++++++++++++++++++++++++++++++++++----------- 1 file changed, 111 insertions(+), 26 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 8aae358..3bc27e1 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -1315,6 +1315,9 @@ struct sam3_tracker { // change, so pe_gen keys it too. Freed with owned_buffers on reset. uint64_t pe_gen = 0; sam3_stagebuf tw_rope_q, tw_rope_k, tw_src_pos; + // Increment 4: perceiver-input twins (pos = cached_sinpe_64 keyed on + // pe_gen; latents are immutable weights, gen fixed at 1) + sam3_stagebuf tw_perc_pos, tw_perc_lat1, tw_perc_lat2; // Increment 3: recycling free-lists for bank-slot / obj-ptr device tensors. // Evicted slots return their tensor pair + buffer here; the next store @@ -6462,7 +6465,14 @@ static bool edgetam_perceiver_forward( const std::vector& mem_pos, // [64 * H * W] int H, int W, std::vector& out_latents, // output: [512 * 64] - std::vector& out_pos) { // output: [512 * 64] + std::vector& out_pos, // output: [512 * 64] + // Increment 4 (device transport): when set, feed the memory features + // by same-backend D2D from this live memenc output tensor instead of + // uploading `mem_features` (which may then be empty), and route the 2D + // window partition IN-GRAPH (a pure byte gather — bit-identical to + // edgetam_window_partition_cpu). trk supplies the input twins. + struct ggml_tensor* mem_features_dev = nullptr, + sam3_tracker* trk = nullptr) { auto t_start = std::chrono::high_resolution_clock::now(); const auto& perc = model.perceiver; @@ -6515,8 +6525,12 @@ static bool edgetam_perceiver_forward( // pos = w + h * W. // Window partition: [D, H*W] → [D, ws2, N_2d] i.e. [64, 16, 256] - std::vector feat_windowed(D * ws2 * N_2d); - edgetam_window_partition_cpu(mem_features.data(), D, H, W, ws, feat_windowed.data()); + // (host path only — the device path partitions in-graph from the leaf) + std::vector feat_windowed; + if (!mem_features_dev) { + feat_windowed.resize((size_t)D * ws2 * N_2d); + edgetam_window_partition_cpu(mem_features.data(), D, H, W, ws, feat_windowed.data()); + } // Reshape latents_2d for windowed processing: [D, N_2d] → [D, 1, N_2d] // Each window has exactly 1 latent token. @@ -6543,9 +6557,22 @@ static bool edgetam_perceiver_forward( ggml_set_name(lat_in, "lat_1d_in"); ggml_set_input(lat_in); - auto* x_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, HW, 1); - ggml_set_name(x_in, "feat_1d_in"); - ggml_set_input(x_in); + // Device transport: the fresh leaf mirrors the memenc output's 4D + // layout so it can be fed by whole-tensor D2D; the graph still sees + // [D, HW, 1] via an in-graph reshape of the leaf (curr-fix pattern). + struct ggml_tensor* x4 = nullptr; + struct ggml_tensor* x_in = nullptr; + if (mem_features_dev) { + x4 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, mem_features_dev->ne[0], + mem_features_dev->ne[1], mem_features_dev->ne[2], 1); + ggml_set_name(x4, "feat_1d_in"); ggml_set_input(x4); + x_in = ggml_reshape_3d(ctx0, x4, D, HW, 1); + } else { + x_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, HW, 1); + ggml_set_name(x_in, "feat_1d_in"); + ggml_set_input(x_in); + x4 = x_in; + } auto* pos_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, HW, 1); ggml_set_name(pos_in, "pos_1d_in"); @@ -6577,13 +6604,25 @@ static bool edgetam_perceiver_forward( return false; } - // Set inputs - ggml_backend_tensor_set(lat_in, latents_1d_data.data(), 0, - D * N_1d * sizeof(float)); - ggml_backend_tensor_set(x_in, mem_features.data(), 0, - D * HW * sizeof(float)); - ggml_backend_tensor_set(pos_in, mem_pos.data(), 0, - D * HW * sizeof(float)); + // Set inputs (device transport: latents/pos via persistent twins, + // features by D2D from the live memenc output — increment 4) + if (trk) + sam3_stagebuf_push(model.transport, trk->tw_perc_lat1, latents_1d_data.data(), + D * N_1d * sizeof(float), lat_in, 1, "perc_lat1"); + else + ggml_backend_tensor_set(lat_in, latents_1d_data.data(), 0, + D * N_1d * sizeof(float)); + if (mem_features_dev) + sam3_stage_feed(model.transport, mem_features_dev, x4, "perc_feat_1d"); + else + ggml_backend_tensor_set(x_in, mem_features.data(), 0, + D * HW * sizeof(float)); + if (trk) + sam3_stagebuf_push(model.transport, trk->tw_perc_pos, mem_pos.data(), + D * HW * sizeof(float), pos_in, trk->pe_gen, "perc_pos"); + else + ggml_backend_tensor_set(pos_in, mem_pos.data(), 0, + D * HW * sizeof(float)); if (!sam3_graph_compute(model.backend, graph, 4)) { fprintf(stderr, "%s: 1D graph compute failed\n", __func__); @@ -6612,10 +6651,33 @@ static bool edgetam_perceiver_forward( return false; } - // Input: windowed features [D, ws2, N_2d] where batch dim = N_2d = 256 - auto* x_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, ws2, N_2d); - ggml_set_name(x_in, "feat_2d_in"); - ggml_set_input(x_in); + // Input: windowed features [D, ws2, N_2d] where batch dim = N_2d = 256. + // Device transport (increment 4): the leaf mirrors the memenc output's + // 4D layout and the window partition happens IN-GRAPH — a pure byte + // gather (reshape/permute/cont), bit-identical to the CPU partition: + // [D, W, H] x = wx*ws+ix, y = wy*ws+iy + // → [D, ws, nw, H] split x into (ix, wx) + // → perm(0,1,3,2) → [D, ws, H, nw] (d, ix, y, wx) + // → [D*ws, ws, nw, nw] split y into (iy, wy) + // → perm(0,1,3,2) → [D*ws, ws, nw, nw] (d·ix, iy, wx, wy) + // → [D, ws2, N_2d] token = iy*ws+ix, window = wy*nw+wx + struct ggml_tensor* x4 = nullptr; + struct ggml_tensor* x_in = nullptr; + if (mem_features_dev) { + x4 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, mem_features_dev->ne[0], + mem_features_dev->ne[1], mem_features_dev->ne[2], 1); + ggml_set_name(x4, "feat_2d_in"); ggml_set_input(x4); + auto* a = ggml_reshape_4d(ctx0, x4, D, ws, nw, H); + auto* b = ggml_cont(ctx0, ggml_permute(ctx0, a, 0, 1, 3, 2)); + auto* c = ggml_reshape_4d(ctx0, b, D * ws, ws, nw, nw); + auto* dperm = ggml_cont(ctx0, ggml_permute(ctx0, c, 0, 1, 3, 2)); + x_in = ggml_reshape_3d(ctx0, dperm, D, ws2, N_2d); + } else { + x_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, ws2, N_2d); + ggml_set_name(x_in, "feat_2d_in"); + ggml_set_input(x_in); + x4 = x_in; + } // Input: latents [D, 1, N_2d] — 1 latent per window auto* lat_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, 1, N_2d); @@ -6650,13 +6712,21 @@ static bool edgetam_perceiver_forward( return false; } - // Set windowed features - ggml_backend_tensor_set(x_in, feat_windowed.data(), 0, - D * ws2 * N_2d * sizeof(float)); + // Set windowed features (device: D2D from the live memenc output; the + // partition itself runs in-graph) + if (mem_features_dev) + sam3_stage_feed(model.transport, mem_features_dev, x4, "perc_feat_2d"); + else + ggml_backend_tensor_set(x_in, feat_windowed.data(), 0, + D * ws2 * N_2d * sizeof(float)); // Set latents: [D, 256] → interpret as [D, 1, 256] - ggml_backend_tensor_set(lat_in, latents_2d_data.data(), 0, - D * N_2d * sizeof(float)); + if (trk) + sam3_stagebuf_push(model.transport, trk->tw_perc_lat2, latents_2d_data.data(), + D * N_2d * sizeof(float), lat_in, 1, "perc_lat2"); + else + ggml_backend_tensor_set(lat_in, latents_2d_data.data(), 0, + D * N_2d * sizeof(float)); if (!sam3_graph_compute(model.backend, graph, 4)) { fprintf(stderr, "%s: 2D graph compute failed\n", __func__); @@ -13041,11 +13111,21 @@ static bool sam3_encode_memory( // and pushing/evicting slots in tracker.mem_banks. SAM3_TIME_BEGIN(_t_mem_bank); - std::vector md(MD * H * H); - ggml_backend_tensor_get(mo, md.data(), 0, md.size() * sizeof(float)); + // Increment 4: on the CUDA device driver the EdgeTAM chain keeps the + // memenc output on device (the perceiver feeds from `mo` by D2D while its + // graph memory is still live), so the 1 MB download disappears. Gated on + // is_edgetam(): SAM2.1's no_obj_embed_spatial branch below mutates md on + // the host and must keep the verbatim path (design §3.4-4). + const bool device_chain = model.transport.kind == SAM3_TRANSPORT_DEVICE && + hp.is_edgetam() && hp.has_perceiver; + std::vector md; + if (!device_chain) { + md.resize((size_t)MD * H * H); + ggml_backend_tensor_get(mo, md.data(), 0, md.size() * sizeof(float)); + } // Apply no_obj_embed_spatial if occluded (SAM2.1 only — EdgeTAM does not have this) - if (obj_score <= 0.0f && model.no_obj_embed_spatial) { + if (!device_chain && obj_score <= 0.0f && model.no_obj_embed_spatial) { std::vector no_obj_emb(MD); auto* noe = model.no_obj_embed_spatial; if (noe->type == GGML_TYPE_F16) { @@ -13072,7 +13152,9 @@ static bool sam3_encode_memory( std::vector perc_latents, perc_pos; if (!edgetam_perceiver_forward(model, md, mem_pos, H, H, - perc_latents, perc_pos)) { + perc_latents, perc_pos, + device_chain ? mo : nullptr, + device_chain ? &tracker : nullptr)) { fprintf(stderr, "%s: perceiver forward failed\n", __func__); SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); ggml_gallocr_free(ga); @@ -13637,6 +13719,9 @@ void sam3_tracker_reset(sam3_tracker& tracker) { sam3_stagebuf_release(tracker.tw_rope_q); sam3_stagebuf_release(tracker.tw_rope_k); sam3_stagebuf_release(tracker.tw_src_pos); + sam3_stagebuf_release(tracker.tw_perc_pos); + sam3_stagebuf_release(tracker.tw_perc_lat1); + sam3_stagebuf_release(tracker.tw_perc_lat2); // Increment 3: the recycling lists hold pointers into owned_buffers' // storage (freed just above) — clear them, never free through them. tracker.slot_free.clear(); From edb60a5f3281369ed19103cd6ba2b99466d1e9a9 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 18:26:25 -0600 Subject: [PATCH 11/16] =?UTF-8?q?cuda:=20stage-transport=20increment=205?= =?UTF-8?q?=20=E2=80=94=20device=20default=20on=20CUDA=20+=20async=20D2D?= =?UTF-8?q?=20feeds?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit report/CUDA-RESIDENCY-DESIGN.md §5 increment 5 (the severable core). The CUDA backend now DEFAULTS to the device transport; SAM3_STAGE_TRANSPORT=host is the runtime escape hatch. The flip sits inside ggml_backend_is_cuda — Metal/ Vulkan/CPU can never take it, so every non-CUDA backend keeps today's host path unconditionally. All device-path D2D feeds (stage_feed + stagebuf twin copies) switch to ggml_backend_tensor_copy_async on the backend's compute stream: same-stream FIFO orders each copy before the graph_compute that consumes it, removing the per-copy stream fence. Every feed source is long-lived (state/ring/twins, or a graph output whose gallocr outlives the consuming compute), so no copy can read recycled memory. Twin UPLOADS stay synchronous (host-side sources may be transient). Deferred (severable, per design §5-inc5): per-stage ctx/graph/gallocr reuse for CUDA-graph replay (needs the CLAUDE.md rule-1 addendum + owner sign-off), pinned host staging for the frame/mask pair, and the encoder-ahead cross-instance copy_async handoff (the transport logs a note when SAM3_GGML_ENCODER_AHEAD is combined with device transport). Gates: host-mode byte-identical to pre-change baseline; host-vs-device A/B byte-identical at 40 frames and a 300-frame soak; default engagement asserted ("stage transport = device (CUDA-resident, default)"); 10.0-10.5 fps. Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 22 +++++++++++++++++----- 1 file changed, 17 insertions(+), 5 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 3bc27e1..bc762c9 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -978,7 +978,13 @@ static void sam3_stage_feed(const sam3_transport& tr, struct ggml_tensor* src, if (sam3_same_layout(src, dst) && !ggml_backend_buffer_is_host(src->buffer) && !ggml_backend_buffer_is_host(dst->buffer)) { - ggml_backend_tensor_copy(src, dst); // same-backend D2D, self-fencing + // Async same-backend D2D (increment 5): queued on the backend's + // compute stream, so FIFO ordering serializes it before the + // graph_compute that consumes dst — no per-copy fence needed. + // Every feed source is long-lived (persistent state / ring / + // twins, or a graph output whose gallocr outlives the consuming + // compute), so the async copy can never read recycled memory. + ggml_backend_tensor_copy_async(tr.backend, tr.backend, src, dst); return; } static std::map warned; @@ -1049,7 +1055,9 @@ static void sam3_stagebuf_push(const sam3_transport& tr, sam3_stagebuf& sb, ggml_backend_tensor_set(sb.dev, host, 0, nbytes); // one upload per content change sb.gen = gen; } - ggml_backend_tensor_copy(sb.dev, dst); // same-backend D2D + // Async D2D from the persistent twin (increment 5) — same-stream FIFO + // orders it before the consuming graph_compute; the twin never moves. + ggml_backend_tensor_copy_async(tr.backend, tr.backend, sb.dev, dst); } // Resolved once at model load, after the backend chain (Metal>CUDA>Vulkan>CPU). @@ -1062,11 +1070,15 @@ static void sam3_transport_init(sam3_transport& tr, ggml_backend_t backend) { const char* env = getenv("SAM3_STAGE_TRANSPORT"); const bool force_host = env && strcmp(env, "host") == 0; const bool want_device = env && strcmp(env, "device") == 0; - (void)force_host; + (void)force_host; (void)want_device; #ifdef GGML_USE_CUDA - if (want_device && !force_host && ggml_backend_is_cuda(backend)) { + // Increment 5: device is the DEFAULT on CUDA (SAM3_STAGE_TRANSPORT=host is + // the escape hatch). This flip lives inside ggml_backend_is_cuda, which + // Metal/Vulkan/CPU can never pass — every other backend stays host. + if (!force_host && ggml_backend_is_cuda(backend)) { tr.kind = SAM3_TRANSPORT_DEVICE; - fprintf(stderr, "%s: stage transport = device (CUDA-resident)\n", __func__); + fprintf(stderr, "%s: stage transport = device (CUDA-resident%s)\n", __func__, + want_device ? "" : ", default"); if (getenv("SAM3_GGML_ENCODER_AHEAD")) fprintf(stderr, "%s: note: encoder-ahead handoff stays host-float until its copy_async clause lands\n", __func__); } From f5875dbc7b524e28dc172285e1e80715df959879 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 18:49:22 -0600 Subject: [PATCH 12/16] =?UTF-8?q?cuda:=20increment=206a=20=E2=80=94=20cach?= =?UTF-8?q?ed=20propagate=20stage=20graph=20(CUDA-graph=20replay=20path)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CLAUDE.md graph-isolation addendum 4 (owner-approved): on the CUDA device driver a stage may keep its ggml_context/cgraph/gallocr alive across frames and recompute in place, keyed on every topology-affecting shape/config and released on any key change / reset / compute failure. Applied to the fused mem-attn + mask-decoder graph in sam3_propagate_single (sam3_prop_gcache on the tracker): steady-state frames skip the graph rebuild AND the gallocr reserve+alloc entirely — no per-frame cudaMalloc/free, stable leaf data pointers (the precondition for ggml-CUDA graph capture to reach replay). Ramp frames (growing M_total) miss the key and rebuild as before. Host driver keeps the verbatim build→compute→free per call. Gates: host byte-identical to pre-program baseline; host-vs-device A/B byte-identical at 40 frames + 300-frame soak; 10.45 fps smoke, 11.11 fps steady state (from 10.0/10.5). Co-Authored-By: Claude Opus 4.8 (1M context) --- CLAUDE.md | 15 +++++ sam3.cpp | 167 +++++++++++++++++++++++++++++++++++++++++++++--------- 2 files changed, 156 insertions(+), 26 deletions(-) diff --git a/CLAUDE.md b/CLAUDE.md index c471d85..d70f398 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -39,6 +39,21 @@ sam3.cpp — a C++14 port of Meta's SAM 3 (Segment Anything Model 3) using ggml **Functions that follow this pattern:** `sam3_segment_pcs` (5 sub-graphs), `sam3_segment_pvs`, `sam3_propagate_single`, `sam3_encode_memory`. +4. **Device-transport addendum (CUDA stage-graph reuse).** On the CUDA + device-transport driver (`sam3_transport`; see `report/CUDA-RESIDENCY-DESIGN.md` + in the trackbench repo) a stage MAY keep its own `ggml_context` + + `ggml_cgraph` + `ggml_gallocr` alive across frames and recompute it in place + (`sam3_stage_cache`), instead of rule 1's build→compute→free-per-call, + PROVIDED: (a) the cached graph is keyed on every shape/config input that + affects its topology and is released + rebuilt on any key change; (b) stages + are never merged — one cached graph per stage, so rule 1's buffer-aliasing + rationale still holds within each compute; (c) all per-frame data still + enters through fresh input leaves (rule 2 untouched). This is what lets + ggml-CUDA's graph capture reach steady replay (stable node/leaf data + pointers) and removes the per-frame cudaMalloc/free churn. Precedent: the + encoder's persistent `state.neck_trk` buffers (steady-state fast path). The + host driver (mac/Metal/CPU) keeps the verbatim per-call pattern. + ## Implementation plan All work follows the phased plan in `PLAN.md`. Read it before starting any phase. Each phase has concrete steps, verification criteria, and the exact structs/functions to implement. diff --git a/sam3.cpp b/sam3.cpp index bc762c9..6d9ddd7 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -1237,6 +1237,42 @@ struct sam3_masklet { int occl_count = 0; // consecutive absent frames in OCCLUDED }; +// Increment 6 (CLAUDE.md graph-isolation addendum 4): a stage's cached +// ggml_context + cgraph + gallocr, kept alive across frames on the CUDA +// device driver and recomputed in place. Keyed on every topology-affecting +// shape/config; released + rebuilt on any key change. Stable node/leaf data +// pointers are what let ggml-CUDA's graph capture reach steady replay, and +// keeping the gallocr kills the per-frame cudaMalloc/free churn. The host +// driver never touches these — it keeps the verbatim build→compute→free. +struct sam3_prop_gcache { + struct ggml_context* ctx = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + uint64_t key = 0; + // input leaves (subset used depends on the cached configuration) + struct ggml_tensor *curr_in = nullptr, *src_pos = nullptr, + *prompt = nullptr, *prompt_pos = nullptr, + *rope_q = nullptr, *rope_k = nullptr, + *sparse = nullptr, *image_pe = nullptr, *dense = nullptr, + *trk_s0 = nullptr, *trk_s1 = nullptr, + *tpos = nullptr, *ptr = nullptr, *ptr_pos = nullptr, + *slot_feat[16] = {}, *slot_pe[16] = {}; + // outputs + struct ggml_tensor *masks = nullptr, *iou = nullptr, *obj = nullptr, + *sam_token = nullptr, *mask_tokens = nullptr; +}; + +static void sam3_prop_gcache_release(sam3_prop_gcache& gc) { + if (gc.galloc) ggml_gallocr_free(gc.galloc); + if (gc.ctx) ggml_free(gc.ctx); + gc = sam3_prop_gcache(); +} + +// FNV-style mix for graph-cache keys. +static inline void sam3_key_mix(uint64_t& k, uint64_t v) { + k ^= v + 0x9e3779b97f4a7c15ull + (k << 6) + (k >> 2); +} + struct sam3_memory_slot { struct ggml_tensor* spatial_feats = nullptr; // [64, 72, 72] struct ggml_tensor* spatial_pe = nullptr; // [64, 72, 72] @@ -1347,6 +1383,9 @@ struct sam3_tracker { // instead of doing 16 per-frame 1 KB D2H gets. The values pass through the // host at store time anyway, so this costs one memcpy per credible frame. std::map>> ptr_host; + + // Increment 6: cached propagate (mem-attn + decoder) stage graph. + sam3_prop_gcache gc_prop; }; // Increment 3: acquire a bank-slot tensor pair (feats + pe sharing one backend @@ -12499,10 +12538,33 @@ static sam3_prop_output sam3_propagate_single( // stays 0 and the combined compute lands in mem_attn_compute_ms. U1 reuses // this graph across frames, which should drive build_ms+alloc_ms → 0. SAM3_TIME_BEGIN(_t_prop_build); - const size_t buf_size = ggml_tensor_overhead() * 32768 + ggml_graph_overhead() * 2; - struct ggml_init_params gparams = {buf_size, nullptr, true}; - auto* ctx0 = ggml_init(gparams); - if (!ctx0) return output; + // Increment 6 (CLAUDE.md addendum 4): on the CUDA device driver, reuse the + // stage graph across frames when every topology-affecting input matches. + // Ramp frames (growing M_total) miss the key and rebuild; steady state + // hits every frame — no rebuild, no gallocr realloc, stable leaf pointers + // (the precondition for ggml-CUDA graph replay). + const bool use_gcache = model.transport.kind == SAM3_TRANSPORT_DEVICE && !use_cml_ma; + uint64_t gkey = 1469598103934665603ull; + sam3_key_mix(gkey, (uint64_t)H); + sam3_key_mix(gkey, (uint64_t)pd.M_total); + sam3_key_mix(gkey, (uint64_t)pd.num_obj_ptr_tokens); + sam3_key_mix(gkey, (uint64_t)pd.M_spatial); + sam3_key_mix(gkey, device_prompt ? (uint64_t)n_sel : 0ull); + sam3_key_mix(gkey, device_prompt ? 2ull : 1ull); + sam3_key_mix(gkey, (uint64_t)N_per_slot); + auto& gc = tracker.gc_prop; + const bool gcache_hit = use_gcache && gc.ctx && gc.key == gkey; + if (use_gcache && !gcache_hit && gc.ctx) sam3_prop_gcache_release(gc); + + struct ggml_context* ctx0 = nullptr; + if (gcache_hit) { + ctx0 = gc.ctx; + } else { + const size_t buf_size = ggml_tensor_overhead() * 32768 + ggml_graph_overhead() * 2; + struct ggml_init_params gparams = {buf_size, nullptr, true}; + ctx0 = ggml_init(gparams); + if (!ctx0) return output; + } // Memory-attention inputs (ggml path only). In the CoreML mem-attn path these // stay null and cond_spatial is an input tensor filled from the CoreML output. @@ -12516,7 +12578,31 @@ static sam3_prop_output sam3_propagate_single( struct ggml_tensor* tpos_leaf = nullptr; struct ggml_tensor* ptr_leaf = nullptr, * ptr_pos_leaf = nullptr; struct ggml_tensor* prompt_node = nullptr, * prompt_pos_node = nullptr; + // Hoisted so both the build path and the cache-hit path can set them + struct ggml_tensor* sparse_in = nullptr, * image_pe = nullptr, * dense_emb = nullptr; + struct ggml_tensor* trk_s0 = nullptr, * trk_s1 = nullptr; + struct ggml_tensor* o_masks = nullptr, * o_iou = nullptr, * o_obj = nullptr; + struct ggml_tensor* o_sam = nullptr, * o_mtok = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + const int H0 = H * 4, H1 = H * 2; + if (gcache_hit) { + // Reuse the cached stage graph in place; only the inputs change. + graph = gc.graph; galloc = gc.galloc; + curr_in = gc.curr_in; curr = gc.curr_in; src_pos_t = gc.src_pos; + prompt_t = gc.prompt; prompt_pos_t = gc.prompt_pos; + rope_q_t = gc.rope_q; rope_k_t = gc.rope_k; + sparse_in = gc.sparse; image_pe = gc.image_pe; dense_emb = gc.dense; + trk_s0 = gc.trk_s0; trk_s1 = gc.trk_s1; + tpos_leaf = gc.tpos; ptr_leaf = gc.ptr; ptr_pos_leaf = gc.ptr_pos; + for (int i = 0; i < 16; ++i) { + slot_feat_leaf[i] = gc.slot_feat[i]; + slot_pe_leaf[i] = gc.slot_pe[i]; + } + o_masks = gc.masks; o_iou = gc.iou; o_obj = gc.obj; + o_sam = gc.sam_token; o_mtok = gc.mask_tokens; + } else { if (use_cml_ma) { // CoreML produced the attended features; inject them as a decoder input. cond_spatial = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); @@ -12603,22 +12689,21 @@ static sam3_prop_output sam3_propagate_single( } // Bug 3 fix: single not_a_point_embed token instead of empty sparse - auto* sparse_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, 1, 1); + sparse_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, 1, 1); ggml_set_name(sparse_in, "prop_sparse"); ggml_set_input(sparse_in); - auto* image_pe = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); + image_pe = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); ggml_set_name(image_pe, "prop_pe"); ggml_set_input(image_pe); - auto* dense_emb = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); + dense_emb = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); ggml_set_name(dense_emb, "prop_dense"); ggml_set_input(dense_emb); - const int H0 = H * 4, H1 = H * 2; - auto* trk_s0 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H0, H0, 1); + trk_s0 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H0, H0, 1); ggml_set_name(trk_s0, "prop_trk_s0"); ggml_set_input(trk_s0); - auto* trk_s1 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H1, H1, 1); + trk_s1 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H1, H1, 1); ggml_set_name(trk_s1, "prop_trk_s1"); ggml_set_input(trk_s1); @@ -12631,7 +12716,7 @@ static sam3_prop_output sam3_propagate_single( ggml_set_output(dec.sam_token); if (dec.mask_tokens) ggml_set_output(dec.mask_tokens); - auto* graph = ggml_new_graph_custom(ctx0, 32768, false); + graph = ggml_new_graph_custom(ctx0, 32768, false); ggml_build_forward_expand(graph, dec.masks); ggml_build_forward_expand(graph, dec.iou_pred); ggml_build_forward_expand(graph, dec.obj_score); @@ -12641,7 +12726,7 @@ static sam3_prop_output sam3_propagate_single( // RFD 0011 U0: gallocr reserve+alloc region (mem_attn_alloc_ms). SAM3_TIME_BEGIN(_t_prop_alloc); - auto* galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); + galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); if (!ggml_gallocr_reserve(galloc, graph) || !ggml_gallocr_alloc_graph(galloc, graph)) { ggml_gallocr_free(galloc); ggml_free(ctx0); @@ -12649,6 +12734,27 @@ static sam3_prop_output sam3_propagate_single( } SAM3_TIME_END(mem_attn_alloc_ms, _t_prop_alloc); + o_masks = dec.masks; o_iou = dec.iou_pred; o_obj = dec.obj_score; + o_sam = dec.sam_token; o_mtok = dec.mask_tokens; + if (use_gcache) { + // Stash the stage graph for reuse (CLAUDE.md addendum 4). Released on + // any key change, tracker reset, or compute failure. + gc.ctx = ctx0; gc.graph = graph; gc.galloc = galloc; gc.key = gkey; + gc.curr_in = curr_in; gc.src_pos = src_pos_t; + gc.prompt = prompt_t; gc.prompt_pos = prompt_pos_t; + gc.rope_q = rope_q_t; gc.rope_k = rope_k_t; + gc.sparse = sparse_in; gc.image_pe = image_pe; gc.dense = dense_emb; + gc.trk_s0 = trk_s0; gc.trk_s1 = trk_s1; + gc.tpos = tpos_leaf; gc.ptr = ptr_leaf; gc.ptr_pos = ptr_pos_leaf; + for (int i = 0; i < 16; ++i) { + gc.slot_feat[i] = slot_feat_leaf[i]; + gc.slot_pe[i] = slot_pe_leaf[i]; + } + gc.masks = o_masks; gc.iou = o_iou; gc.obj = o_obj; + gc.sam_token = o_sam; gc.mask_tokens = o_mtok; + } + } // end build path (gcache_hit skips straight to the input uploads) + if (use_cml_ma) { // Inject the CoreML mem-attn output as the decoder's conditioned input. ggml_backend_tensor_set(cond_spatial, cml_cond.data(), 0, cml_cond.size() * sizeof(float)); @@ -12724,8 +12830,12 @@ static sam3_prop_output sam3_propagate_single( // because the decoder is fused into this same graph on the EdgeTAM path). SAM3_TIME_BEGIN(_t_prop_compute); if (!sam3_graph_compute(model.backend, graph, 4)) { - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (use_gcache) { + sam3_prop_gcache_release(gc); // never keep a wedged graph + } else { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } return output; } SAM3_TIME_END(mem_attn_compute_ms, _t_prop_compute); @@ -12740,7 +12850,7 @@ static sam3_prop_output sam3_propagate_single( if (use_multimask) { // Read all 4 IoU predictions std::vector all_ious(num_mask_tokens); - ggml_backend_tensor_get(dec.iou_pred, all_ious.data(), 0, num_mask_tokens * sizeof(float)); + ggml_backend_tensor_get(o_iou, all_ious.data(), 0, num_mask_tokens * sizeof(float)); // Multimask uses mask tokens 1-3 (skip token 0 which is the single-mask output) int best_idx = 1; @@ -12754,35 +12864,37 @@ static sam3_prop_output sam3_propagate_single( output.mask_w = mhw; output.mask_logits.resize(mhw * mhw); // Read best mask (offset by best_idx * mhw * mhw) - ggml_backend_tensor_get(dec.masks, output.mask_logits.data(), + ggml_backend_tensor_get(o_masks, output.mask_logits.data(), best_idx * mhw * mhw * sizeof(float), mhw * mhw * sizeof(float)); output.iou_scores.resize(1); output.iou_scores[0] = best_iou; - ggml_backend_tensor_get(dec.obj_score, &output.obj_score, 0, sizeof(float)); + ggml_backend_tensor_get(o_obj, &output.obj_score, 0, sizeof(float)); // Object pointer token: use best multimask token if use_multimask_token_for_obj_ptr output.sam_token.resize(D); - if (hp.use_multimask_token_for_obj_ptr && dec.mask_tokens) { - ggml_backend_tensor_get(dec.mask_tokens, output.sam_token.data(), + if (hp.use_multimask_token_for_obj_ptr && o_mtok) { + ggml_backend_tensor_get(o_mtok, output.sam_token.data(), best_idx * D * sizeof(float), D * sizeof(float)); } else { - ggml_backend_tensor_get(dec.sam_token, output.sam_token.data(), 0, D * sizeof(float)); + ggml_backend_tensor_get(o_sam, output.sam_token.data(), 0, D * sizeof(float)); } } else { output.n_masks = 1; output.mask_h = mhw; output.mask_w = mhw; output.mask_logits.resize(mhw * mhw); - ggml_backend_tensor_get(dec.masks, output.mask_logits.data(), 0, mhw * mhw * sizeof(float)); + ggml_backend_tensor_get(o_masks, output.mask_logits.data(), 0, mhw * mhw * sizeof(float)); output.iou_scores.resize(1); - ggml_backend_tensor_get(dec.iou_pred, output.iou_scores.data(), 0, sizeof(float)); - ggml_backend_tensor_get(dec.obj_score, &output.obj_score, 0, sizeof(float)); + ggml_backend_tensor_get(o_iou, output.iou_scores.data(), 0, sizeof(float)); + ggml_backend_tensor_get(o_obj, &output.obj_score, 0, sizeof(float)); output.sam_token.resize(D); - ggml_backend_tensor_get(dec.sam_token, output.sam_token.data(), 0, D * sizeof(float)); + ggml_backend_tensor_get(o_sam, output.sam_token.data(), 0, D * sizeof(float)); } - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (!use_gcache) { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } return output; } @@ -13739,6 +13851,9 @@ void sam3_tracker_reset(sam3_tracker& tracker) { tracker.slot_free.clear(); tracker.ptr_free.clear(); tracker.ptr_host.clear(); + // Increment 6: drop the cached stage graph (its leaves may reference + // dims tied to the old session; it rebuilds on the next device frame). + sam3_prop_gcache_release(tracker.gc_prop); if (tracker.ctx) { ggml_free(tracker.ctx); tracker.ctx = nullptr; From b8e7ecb7beb1e1b8e886564980f844a1ad0acba4 Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 18:57:26 -0600 Subject: [PATCH 13/16] =?UTF-8?q?cuda:=20increment=206b=20=E2=80=94=20cach?= =?UTF-8?q?ed=20memory-encoder=20+=20perceiver=20stage=20graphs?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Same CLAUDE.md addendum-4 contract as 6a, applied to the three remaining per-credible-frame graphs: the memory encoder (sam3_stage_gcache on the tracker; keyed on H/INTERPOL/D) and the perceiver 1D/2D sub-graphs (keyed on their fixed dims + device flag). Steady state skips build+reserve+alloc for all of them; the cached memenc output tensor (mo) keeps a frame-invariant address, which also stabilizes the perceiver's D2D feed source. Compute failure releases the affected cache. Host driver unchanged. Gates: byte-identity A/B 40f + 300-frame soak; 10.81 fps smoke, 11.28 fps steady state (from 10.45/11.11). Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 195 +++++++++++++++++++++++++++++++++++++++++++------------ 1 file changed, 152 insertions(+), 43 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 6d9ddd7..465be92 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -1268,6 +1268,23 @@ static void sam3_prop_gcache_release(sam3_prop_gcache& gc) { gc = sam3_prop_gcache(); } +// Generic small-stage cache (same addendum-4 contract as sam3_prop_gcache) +// for stages with few leaves: memory encoder, perceiver 1D/2D sub-graphs. +struct sam3_stage_gcache { + struct ggml_context* ctx = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + uint64_t key = 0; + struct ggml_tensor* in[3] = {}; + struct ggml_tensor* out[1] = {}; +}; + +static void sam3_stage_gcache_release(sam3_stage_gcache& gc) { + if (gc.galloc) ggml_gallocr_free(gc.galloc); + if (gc.ctx) ggml_free(gc.ctx); + gc = sam3_stage_gcache(); +} + // FNV-style mix for graph-cache keys. static inline void sam3_key_mix(uint64_t& k, uint64_t v) { k ^= v + 0x9e3779b97f4a7c15ull + (k << 6) + (k >> 2); @@ -1386,6 +1403,8 @@ struct sam3_tracker { // Increment 6: cached propagate (mem-attn + decoder) stage graph. sam3_prop_gcache gc_prop; + // Increment 6b: cached memory-encoder + perceiver sub-graphs. + sam3_stage_gcache gc_memenc, gc_perc1, gc_perc2; }; // Increment 3: acquire a bank-slot tensor pair (feats + pe sharing one backend @@ -6595,24 +6614,42 @@ static bool edgetam_perceiver_forward( // ══════════════════════════════════════════════════════════════════════ std::vector result_1d(D * N_1d); { + // Increment 6b: cached 1D sub-graph (fixed shapes; device driver only) + const bool use_pgc = trk && model.transport.kind == SAM3_TRANSPORT_DEVICE; + uint64_t pkey = 1469598103934665603ull; + sam3_key_mix(pkey, (uint64_t)D); sam3_key_mix(pkey, (uint64_t)HW); + sam3_key_mix(pkey, (uint64_t)N_1d); + sam3_key_mix(pkey, mem_features_dev ? 2ull : 1ull); + sam3_stage_gcache* pgc = use_pgc ? &trk->gc_perc1 : nullptr; + const bool pgc_hit = pgc && pgc->ctx && pgc->key == pkey; + if (pgc && !pgc_hit && pgc->ctx) sam3_stage_gcache_release(*pgc); + + struct ggml_context* ctx0 = nullptr; + struct ggml_tensor* lat_in = nullptr, * x4 = nullptr, * x_in = nullptr; + struct ggml_tensor* pos_in = nullptr, * latents = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + if (pgc_hit) { + ctx0 = pgc->ctx; graph = pgc->graph; galloc = pgc->galloc; + lat_in = pgc->in[0]; x4 = pgc->in[1]; pos_in = pgc->in[2]; + latents = pgc->out[0]; + } else { const size_t buf_size = ggml_tensor_overhead() * 4096 + ggml_graph_overhead(); struct ggml_init_params gparams = {buf_size, nullptr, true}; - auto* ctx0 = ggml_init(gparams); + ctx0 = ggml_init(gparams); if (!ctx0) { fprintf(stderr, "%s: failed to init 1D context\n", __func__); return false; } // Input tensors (fresh, no dependency chains) - auto* lat_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N_1d, 1); + lat_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, N_1d, 1); ggml_set_name(lat_in, "lat_1d_in"); ggml_set_input(lat_in); // Device transport: the fresh leaf mirrors the memenc output's 4D // layout so it can be fed by whole-tensor D2D; the graph still sees // [D, HW, 1] via an in-graph reshape of the leaf (curr-fix pattern). - struct ggml_tensor* x4 = nullptr; - struct ggml_tensor* x_in = nullptr; if (mem_features_dev) { x4 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, mem_features_dev->ne[0], mem_features_dev->ne[1], mem_features_dev->ne[2], 1); @@ -6625,12 +6662,12 @@ static bool edgetam_perceiver_forward( x4 = x_in; } - auto* pos_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, HW, 1); + pos_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, HW, 1); ggml_set_name(pos_in, "pos_1d_in"); ggml_set_input(pos_in); // Run perceiver layers - auto* latents = lat_in; + latents = lat_in; for (int l = 0; l < n_layers; ++l) { latents = edgetam_perceiver_layer_forward(ctx0, latents, x_in, pos_in, perc.layers[l]); @@ -6642,10 +6679,10 @@ static bool edgetam_perceiver_forward( ggml_set_output(latents); // Build + allocate + compute - auto* graph = ggml_new_graph_custom(ctx0, 16384, false); + graph = ggml_new_graph_custom(ctx0, 16384, false); ggml_build_forward_expand(graph, latents); - auto* galloc = ggml_gallocr_new( + galloc = ggml_gallocr_new( ggml_backend_get_default_buffer_type(model.backend)); if (!ggml_gallocr_reserve(galloc, graph) || !ggml_gallocr_alloc_graph(galloc, graph)) { @@ -6655,6 +6692,13 @@ static bool edgetam_perceiver_forward( return false; } + if (use_pgc) { + pgc->ctx = ctx0; pgc->graph = graph; pgc->galloc = galloc; pgc->key = pkey; + pgc->in[0] = lat_in; pgc->in[1] = x4; pgc->in[2] = pos_in; + pgc->out[0] = latents; + } + } // end build path + // Set inputs (device transport: latents/pos via persistent twins, // features by D2D from the live memenc output — increment 4) if (trk) @@ -6677,16 +6721,22 @@ static bool edgetam_perceiver_forward( if (!sam3_graph_compute(model.backend, graph, 4)) { fprintf(stderr, "%s: 1D graph compute failed\n", __func__); - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (pgc && (pgc_hit || pgc->ctx)) { + sam3_stage_gcache_release(*pgc); + } else { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } return false; } ggml_backend_tensor_get(latents, result_1d.data(), 0, D * N_1d * sizeof(float)); - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (!use_pgc) { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } } // ══════════════════════════════════════════════════════════════════════ @@ -6694,9 +6744,28 @@ static bool edgetam_perceiver_forward( // ══════════════════════════════════════════════════════════════════════ std::vector result_2d(D * N_2d); { + // Increment 6b: cached 2D sub-graph (fixed shapes; device driver only) + const bool use_pgc = trk && model.transport.kind == SAM3_TRANSPORT_DEVICE; + uint64_t pkey = 1469598103934665603ull; + sam3_key_mix(pkey, (uint64_t)D); sam3_key_mix(pkey, (uint64_t)ws2); + sam3_key_mix(pkey, (uint64_t)N_2d); + sam3_key_mix(pkey, mem_features_dev ? 2ull : 1ull); + sam3_stage_gcache* pgc = use_pgc ? &trk->gc_perc2 : nullptr; + const bool pgc_hit = pgc && pgc->ctx && pgc->key == pkey; + if (pgc && !pgc_hit && pgc->ctx) sam3_stage_gcache_release(*pgc); + + struct ggml_context* ctx0 = nullptr; + struct ggml_tensor* x4 = nullptr, * x_in = nullptr, * lat_in = nullptr; + struct ggml_tensor* lat_out = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + if (pgc_hit) { + ctx0 = pgc->ctx; graph = pgc->graph; galloc = pgc->galloc; + x4 = pgc->in[0]; lat_in = pgc->in[1]; lat_out = pgc->out[0]; + } else { const size_t buf_size = ggml_tensor_overhead() * 4096 + ggml_graph_overhead(); struct ggml_init_params gparams = {buf_size, nullptr, true}; - auto* ctx0 = ggml_init(gparams); + ctx0 = ggml_init(gparams); if (!ctx0) { fprintf(stderr, "%s: failed to init 2D context\n", __func__); return false; @@ -6712,8 +6781,6 @@ static bool edgetam_perceiver_forward( // → [D*ws, ws, nw, nw] split y into (iy, wy) // → perm(0,1,3,2) → [D*ws, ws, nw, nw] (d·ix, iy, wx, wy) // → [D, ws2, N_2d] token = iy*ws+ix, window = wy*nw+wx - struct ggml_tensor* x4 = nullptr; - struct ggml_tensor* x_in = nullptr; if (mem_features_dev) { x4 = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, mem_features_dev->ne[0], mem_features_dev->ne[1], mem_features_dev->ne[2], 1); @@ -6731,7 +6798,7 @@ static bool edgetam_perceiver_forward( } // Input: latents [D, 1, N_2d] — 1 latent per window - auto* lat_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, 1, N_2d); + lat_in = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, D, 1, N_2d); ggml_set_name(lat_in, "lat_2d_in"); ggml_set_input(lat_in); @@ -6743,17 +6810,17 @@ static bool edgetam_perceiver_forward( } // The 2D latents: [D, 1, N_2d] → squeeze to [D, N_2d] for output - auto* lat_out = ggml_reshape_2d(ctx0, latents, D, N_2d); + lat_out = ggml_reshape_2d(ctx0, latents, D, N_2d); // Final LayerNorm (shared norm_w/norm_b) lat_out = sam3_layer_norm(ctx0, lat_out, perc.norm_w, perc.norm_b); ggml_set_name(lat_out, "lat_2d_out"); ggml_set_output(lat_out); - auto* graph = ggml_new_graph_custom(ctx0, 16384, false); + graph = ggml_new_graph_custom(ctx0, 16384, false); ggml_build_forward_expand(graph, lat_out); - auto* galloc = ggml_gallocr_new( + galloc = ggml_gallocr_new( ggml_backend_get_default_buffer_type(model.backend)); if (!ggml_gallocr_reserve(galloc, graph) || !ggml_gallocr_alloc_graph(galloc, graph)) { @@ -6763,6 +6830,13 @@ static bool edgetam_perceiver_forward( return false; } + if (use_pgc) { + pgc->ctx = ctx0; pgc->graph = graph; pgc->galloc = galloc; pgc->key = pkey; + pgc->in[0] = x4; pgc->in[1] = lat_in; + pgc->out[0] = lat_out; + } + } // end build path + // Set windowed features (device: D2D from the live memenc output; the // partition itself runs in-graph) if (mem_features_dev) @@ -6781,16 +6855,22 @@ static bool edgetam_perceiver_forward( if (!sam3_graph_compute(model.backend, graph, 4)) { fprintf(stderr, "%s: 2D graph compute failed\n", __func__); - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (pgc && (pgc_hit || pgc->ctx)) { + sam3_stage_gcache_release(*pgc); + } else { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } return false; } ggml_backend_tensor_get(lat_out, result_2d.data(), 0, D * N_2d * sizeof(float)); - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (!use_pgc) { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } } // ══════════════════════════════════════════════════════════════════════ @@ -13151,12 +13231,33 @@ static bool sam3_encode_memory( // RFD 0011 U0: memory-encoder graph CONSTRUCTION region (mem_encoder_build_ms). // U1 will reuse this graph across frames; build_ms+alloc_ms should drop to ~0. SAM3_TIME_BEGIN(_t_mem_build); + // Increment 6b: cached memory-encoder graph (CLAUDE.md addendum 4) — fixed + // shapes at a given feat size, so steady state hits every credible frame. + const bool use_mgc = model.transport.kind == SAM3_TRANSPORT_DEVICE; + uint64_t mkey = 1469598103934665603ull; + sam3_key_mix(mkey, (uint64_t)H); + sam3_key_mix(mkey, (uint64_t)INTERPOL); + sam3_key_mix(mkey, (uint64_t)D); + auto& mgc = tracker.gc_memenc; + const bool mgc_hit = use_mgc && mgc.ctx && mgc.key == mkey; + if (use_mgc && !mgc_hit && mgc.ctx) sam3_stage_gcache_release(mgc); + + struct ggml_context* ctx0 = nullptr; + struct ggml_tensor* mask_in = nullptr; + struct ggml_tensor* pix_in_raw = nullptr; + struct ggml_tensor* mo = nullptr; + struct ggml_cgraph* g = nullptr; + struct ggml_gallocr* ga = nullptr; + if (mgc_hit) { + ctx0 = mgc.ctx; g = mgc.graph; ga = mgc.galloc; + mask_in = mgc.in[0]; pix_in_raw = mgc.in[1]; mo = mgc.out[0]; + } else { const size_t bs = ggml_tensor_overhead() * 16384 + ggml_graph_overhead(); struct ggml_init_params gp = {bs, nullptr, true}; - auto* ctx0 = ggml_init(gp); + ctx0 = ggml_init(gp); if (!ctx0) return false; - auto* mask_in = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, INTERPOL, INTERPOL, 1, 1); + mask_in = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, INTERPOL, INTERPOL, 1, 1); ggml_set_name(mask_in, "mem_mask"); ggml_set_input(mask_in); @@ -13179,7 +13280,7 @@ static bool sam3_encode_memory( // Pixel projection — use fresh input tensor to avoid pulling in the // entire ViT+neck recomputation from state.neck_trk[2]'s dependency tree - auto* pix_in_raw = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); + pix_in_raw = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, D, H, H, 1); ggml_set_name(pix_in_raw, "mem_pix_feat"); ggml_set_input(pix_in_raw); auto* pix_in = ggml_cont(ctx0, ggml_permute(ctx0, pix_in_raw, 2, 0, 1, 3)); @@ -13197,19 +13298,19 @@ static bool sam3_encode_memory( model.mem_enc.fuser_fc2_w[i], model.mem_enc.fuser_fc2_b[i], model.mem_enc.fuser_gamma[i]); auto* fused_out = ggml_cont(ctx0, ggml_permute(ctx0, fused, 2, 0, 1, 3)); - auto* mo = ggml_conv_2d(ctx0, model.mem_enc.out_proj_w, fused_out, 1, 1, 0, 0, 1, 1); + mo = ggml_conv_2d(ctx0, model.mem_enc.out_proj_w, fused_out, 1, 1, 0, 0, 1, 1); mo = ggml_add(ctx0, mo, ggml_reshape_4d(ctx0, model.mem_enc.out_proj_b, 1, 1, MD, 1)); mo = ggml_cont(ctx0, ggml_permute(ctx0, mo, 1, 2, 0, 3)); ggml_set_name(mo, "mem_out"); ggml_set_output(mo); - auto* g = ggml_new_graph_custom(ctx0, 16384, false); + g = ggml_new_graph_custom(ctx0, 16384, false); ggml_build_forward_expand(g, mo); SAM3_TIME_END(mem_encoder_build_ms, _t_mem_build); // RFD 0011 U0: gallocr reserve+alloc region (mem_encoder_alloc_ms). SAM3_TIME_BEGIN(_t_mem_alloc); - auto* ga = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); + ga = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); if (!ggml_gallocr_reserve(ga, g) || !ggml_gallocr_alloc_graph(ga, g)) { ggml_gallocr_free(ga); ggml_free(ctx0); @@ -13217,14 +13318,24 @@ static bool sam3_encode_memory( } SAM3_TIME_END(mem_encoder_alloc_ms, _t_mem_alloc); + if (use_mgc) { + mgc.ctx = ctx0; mgc.graph = g; mgc.galloc = ga; mgc.key = mkey; + mgc.in[0] = mask_in; mgc.in[1] = pix_in_raw; mgc.out[0] = mo; + } + } // end build path (mgc_hit skips straight to the input uploads) + ggml_backend_tensor_set(mask_in, m_interp.data(), 0, m_interp.size() * sizeof(float)); // Copy pixel features from state tensor to fresh input (transport driver) sam3_stage_feed(model.transport, state.neck_trk[2], pix_in_raw, "mem_pix_feat"); // RFD 0011 U0: backend compute of the memory encoder (mem_encoder_compute_ms). SAM3_TIME_BEGIN(_t_mem_compute); if (!sam3_graph_compute(model.backend, g, 4)) { - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (use_mgc) { + sam3_stage_gcache_release(mgc); // never keep a wedged graph + } else { + ggml_gallocr_free(ga); + ggml_free(ctx0); + } return false; } SAM3_TIME_END(mem_encoder_compute_ms, _t_mem_compute); @@ -13281,8 +13392,7 @@ static bool sam3_encode_memory( device_chain ? &tracker : nullptr)) { fprintf(stderr, "%s: perceiver forward failed\n", __func__); SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (!use_mgc) { ggml_gallocr_free(ga); ggml_free(ctx0); } return false; } @@ -13335,8 +13445,7 @@ static bool sam3_encode_memory( if (!sam3_acquire_slot_pair(tracker, model, MD, N_perc, 1, &st, &spe, &slot_buf)) { fprintf(stderr, "%s: slot pair alloc failed\n", __func__); SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (!use_mgc) { ggml_gallocr_free(ga); ggml_free(ctx0); } return false; } ggml_backend_tensor_set(st, perc_latents.data(), 0, MD * N_perc * sizeof(float)); @@ -13369,8 +13478,7 @@ static bool sam3_encode_memory( } } SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (!use_mgc) { ggml_gallocr_free(ga); ggml_free(ctx0); } return true; } @@ -13382,8 +13490,7 @@ static bool sam3_encode_memory( if (!sam3_acquire_slot_pair(tracker, model, MD, H, H, &st, &spe, &slot_buf)) { fprintf(stderr, "%s: slot pair alloc failed\n", __func__); SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (!use_mgc) { ggml_gallocr_free(ga); ggml_free(ctx0); } return false; } ggml_backend_tensor_set(st, md.data(), 0, md.size() * sizeof(float)); @@ -13420,8 +13527,7 @@ static bool sam3_encode_memory( } } SAM3_TIME_END(memory_bank_update_ms, _t_mem_bank); - ggml_gallocr_free(ga); - ggml_free(ctx0); + if (!use_mgc) { ggml_gallocr_free(ga); ggml_free(ctx0); } return true; } @@ -13851,9 +13957,12 @@ void sam3_tracker_reset(sam3_tracker& tracker) { tracker.slot_free.clear(); tracker.ptr_free.clear(); tracker.ptr_host.clear(); - // Increment 6: drop the cached stage graph (its leaves may reference - // dims tied to the old session; it rebuilds on the next device frame). + // Increment 6: drop the cached stage graphs (their leaves may reference + // dims tied to the old session; they rebuild on the next device frame). sam3_prop_gcache_release(tracker.gc_prop); + sam3_stage_gcache_release(tracker.gc_memenc); + sam3_stage_gcache_release(tracker.gc_perc1); + sam3_stage_gcache_release(tracker.gc_perc2); if (tracker.ctx) { ggml_free(tracker.ctx); tracker.ctx = nullptr; From e865052a26e16762fe96aba4a842084843c08bab Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 19:03:02 -0600 Subject: [PATCH 14/16] =?UTF-8?q?cuda:=20increment=206c=20=E2=80=94=20cach?= =?UTF-8?q?ed=20image-encoder=20stage=20graph?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Same addendum-4 contract applied to the largest per-frame graph (RepViT + FPN neck, runs every frame): sam3_stage_gcache on sam3_state, keyed on img_size/n_fpn, released in sam3_free_state / on compute failure. Steady state skips the encoder's build+reserve+alloc entirely; the FPN output tensors keep frame-invariant addresses feeding the existing D2D fpn->neck_trk fast path. Pinned host staging for the 12.6 MB frame upload evaluated and deferred: preprocess returns a fresh vector per frame, so pinning needs a preprocess-into-persistent-buffer restructure for ~1 ms — logged as backlog. Gates: byte-identity A/B 40f + 300-frame soak; 11.49 fps smoke, 12.05 fps steady state (from 10.81/11.28). Co-Authored-By: Claude Opus 4.8 (1M context) --- sam3.cpp | 153 ++++++++++++++++++++++++++++++++++--------------------- 1 file changed, 94 insertions(+), 59 deletions(-) diff --git a/sam3.cpp b/sam3.cpp index 465be92..c169325 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -1060,6 +1060,59 @@ static void sam3_stagebuf_push(const sam3_transport& tr, sam3_stagebuf& sb, ggml_backend_tensor_copy_async(tr.backend, tr.backend, sb.dev, dst); } +// Increment 6 (CLAUDE.md graph-isolation addendum 4): a stage's cached +// ggml_context + cgraph + gallocr, kept alive across frames on the CUDA +// device driver and recomputed in place. Keyed on every topology-affecting +// shape/config; released + rebuilt on any key change. Stable node/leaf data +// pointers are what let ggml-CUDA's graph capture reach steady replay, and +// keeping the gallocr kills the per-frame cudaMalloc/free churn. The host +// driver never touches these — it keeps the verbatim build→compute→free. +struct sam3_prop_gcache { + struct ggml_context* ctx = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + uint64_t key = 0; + // input leaves (subset used depends on the cached configuration) + struct ggml_tensor *curr_in = nullptr, *src_pos = nullptr, + *prompt = nullptr, *prompt_pos = nullptr, + *rope_q = nullptr, *rope_k = nullptr, + *sparse = nullptr, *image_pe = nullptr, *dense = nullptr, + *trk_s0 = nullptr, *trk_s1 = nullptr, + *tpos = nullptr, *ptr = nullptr, *ptr_pos = nullptr, + *slot_feat[16] = {}, *slot_pe[16] = {}; + // outputs + struct ggml_tensor *masks = nullptr, *iou = nullptr, *obj = nullptr, + *sam_token = nullptr, *mask_tokens = nullptr; +}; + +static void sam3_prop_gcache_release(sam3_prop_gcache& gc) { + if (gc.galloc) ggml_gallocr_free(gc.galloc); + if (gc.ctx) ggml_free(gc.ctx); + gc = sam3_prop_gcache(); +} + +// Generic small-stage cache (same addendum-4 contract as sam3_prop_gcache) +// for stages with few leaves: image encoder, memory encoder, perceiver 1D/2D. +struct sam3_stage_gcache { + struct ggml_context* ctx = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_gallocr* galloc = nullptr; + uint64_t key = 0; + struct ggml_tensor* in[3] = {}; + struct ggml_tensor* out[4] = {}; // encoder has 4 FPN outputs +}; + +static void sam3_stage_gcache_release(sam3_stage_gcache& gc) { + if (gc.galloc) ggml_gallocr_free(gc.galloc); + if (gc.ctx) ggml_free(gc.ctx); + gc = sam3_stage_gcache(); +} + +// FNV-style mix for graph-cache keys. +static inline void sam3_key_mix(uint64_t& k, uint64_t v) { + k ^= v + 0x9e3779b97f4a7c15ull + (k << 6) + (k >> 2); +} + // Resolved once at model load, after the backend chain (Metal>CUDA>Vulkan>CPU). // Default host on every backend; device only when explicitly requested AND the // answered backend is CUDA — Metal can never pass this gate, so the mac cannot @@ -1203,6 +1256,9 @@ struct sam3_state { // sam3_populate_pe_cache actually recomputes). Freed in sam3_free_state. uint64_t pe_cache_gen = 0; sam3_stagebuf tw_sparse, tw_image_pe, tw_dense; + + // Increment 6c: cached image-encoder stage graph (device driver only). + sam3_stage_gcache gc_encoder; }; /* @@ -1237,59 +1293,6 @@ struct sam3_masklet { int occl_count = 0; // consecutive absent frames in OCCLUDED }; -// Increment 6 (CLAUDE.md graph-isolation addendum 4): a stage's cached -// ggml_context + cgraph + gallocr, kept alive across frames on the CUDA -// device driver and recomputed in place. Keyed on every topology-affecting -// shape/config; released + rebuilt on any key change. Stable node/leaf data -// pointers are what let ggml-CUDA's graph capture reach steady replay, and -// keeping the gallocr kills the per-frame cudaMalloc/free churn. The host -// driver never touches these — it keeps the verbatim build→compute→free. -struct sam3_prop_gcache { - struct ggml_context* ctx = nullptr; - struct ggml_cgraph* graph = nullptr; - struct ggml_gallocr* galloc = nullptr; - uint64_t key = 0; - // input leaves (subset used depends on the cached configuration) - struct ggml_tensor *curr_in = nullptr, *src_pos = nullptr, - *prompt = nullptr, *prompt_pos = nullptr, - *rope_q = nullptr, *rope_k = nullptr, - *sparse = nullptr, *image_pe = nullptr, *dense = nullptr, - *trk_s0 = nullptr, *trk_s1 = nullptr, - *tpos = nullptr, *ptr = nullptr, *ptr_pos = nullptr, - *slot_feat[16] = {}, *slot_pe[16] = {}; - // outputs - struct ggml_tensor *masks = nullptr, *iou = nullptr, *obj = nullptr, - *sam_token = nullptr, *mask_tokens = nullptr; -}; - -static void sam3_prop_gcache_release(sam3_prop_gcache& gc) { - if (gc.galloc) ggml_gallocr_free(gc.galloc); - if (gc.ctx) ggml_free(gc.ctx); - gc = sam3_prop_gcache(); -} - -// Generic small-stage cache (same addendum-4 contract as sam3_prop_gcache) -// for stages with few leaves: memory encoder, perceiver 1D/2D sub-graphs. -struct sam3_stage_gcache { - struct ggml_context* ctx = nullptr; - struct ggml_cgraph* graph = nullptr; - struct ggml_gallocr* galloc = nullptr; - uint64_t key = 0; - struct ggml_tensor* in[3] = {}; - struct ggml_tensor* out[1] = {}; -}; - -static void sam3_stage_gcache_release(sam3_stage_gcache& gc) { - if (gc.galloc) ggml_gallocr_free(gc.galloc); - if (gc.ctx) ggml_free(gc.ctx); - gc = sam3_stage_gcache(); -} - -// FNV-style mix for graph-cache keys. -static inline void sam3_key_mix(uint64_t& k, uint64_t v) { - k ^= v + 0x9e3779b97f4a7c15ull + (k << 6) + (k >> 2); -} - struct sam3_memory_slot { struct ggml_tensor* spatial_feats = nullptr; // [64, 72, 72] struct ggml_tensor* spatial_pe = nullptr; // [64, 72, 72] @@ -3974,10 +3977,11 @@ void sam3_state_set_orig_dims(sam3_state& state, int w, int h) { } void sam3_free_state(sam3_state& state) { - // Stage-transport twins (increment 2) + // Stage-transport twins (increment 2) + cached encoder graph (increment 6c) sam3_stagebuf_release(state.tw_sparse); sam3_stagebuf_release(state.tw_image_pe); sam3_stagebuf_release(state.tw_dense); + sam3_stage_gcache_release(state.gc_encoder); if (state.galloc) { ggml_gallocr_free(state.galloc); state.galloc = nullptr; @@ -5624,6 +5628,24 @@ static bool edgetam_encode_image(sam3_state& state, struct ggml_tensor* inp = nullptr; int n_fpn = 4 - hp.scalp; struct ggml_tensor* fpn_outs[4] = {}; + struct ggml_gallocr* galloc = nullptr; + + // Increment 6c: cached encoder graph (CLAUDE.md addendum 4) — the biggest + // per-frame graph; fixed shapes at a given img_size, so steady state hits + // every frame and ggml-CUDA graph capture sees stable pointers. + const bool use_egc = model.transport.kind == SAM3_TRANSPORT_DEVICE; + uint64_t ekey = 1469598103934665603ull; + sam3_key_mix(ekey, (uint64_t)img_size); + sam3_key_mix(ekey, (uint64_t)n_fpn); + auto& egc = state.gc_encoder; + const bool egc_hit = use_egc && egc.ctx && egc.key == ekey; + if (use_egc && !egc_hit && egc.ctx) sam3_stage_gcache_release(egc); + + if (egc_hit) { + ctx0 = egc.ctx; graph = egc.graph; galloc = egc.galloc; + inp = egc.in[0]; + for (int i = 0; i < 4; ++i) fpn_outs[i] = egc.out[i]; + } else { { SAM3_TIME_SCOPE(image_encoder_build_ms); const size_t buf_size = ggml_tensor_overhead() * 16384 + ggml_graph_overhead() * 2; @@ -5667,7 +5689,7 @@ static bool edgetam_encode_image(sam3_state& state, // ── Allocate ───────────────────────────────────────────────────────── // RFD 0011 U0: gallocr reserve+alloc region (image_encoder_alloc_ms), // separate from build and compute. U1 reuses a pre-reserved allocator here. - auto* galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); + galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); { SAM3_TIME_SCOPE(image_encoder_alloc_ms); if (!ggml_gallocr_reserve(galloc, graph)) { @@ -5684,6 +5706,13 @@ static bool edgetam_encode_image(sam3_state& state, } } + if (use_egc) { + egc.ctx = ctx0; egc.graph = graph; egc.galloc = galloc; egc.key = ekey; + egc.in[0] = inp; + for (int i = 0; i < 4; ++i) egc.out[i] = fpn_outs[i]; + } + } // end build path (egc_hit skips straight to the input upload) + // Set input image ggml_backend_tensor_set(inp, img_data.data(), 0, img_data.size() * sizeof(float)); @@ -5694,8 +5723,12 @@ static bool edgetam_encode_image(sam3_state& state, SAM3_TIME_SCOPE(image_encoder_compute_ms); if (!sam3_graph_compute(model.backend, graph, state.n_threads)) { fprintf(stderr, "%s: graph compute failed\n", __func__); - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (use_egc) { + sam3_stage_gcache_release(egc); // never keep a wedged graph + } else { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } return false; } } @@ -5796,8 +5829,10 @@ static bool edgetam_encode_image(sam3_state& state, } } - ggml_gallocr_free(galloc); - ggml_free(ctx0); + if (!use_egc) { + ggml_gallocr_free(galloc); + ggml_free(ctx0); + } auto t_end = std::chrono::high_resolution_clock::now(); auto ms = std::chrono::duration_cast(t_end - t_start).count(); From 56a27ece74d229b3c03a2de4c078c6883e06ebfa Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Tue, 7 Jul 2026 19:30:12 -0600 Subject: [PATCH 15/16] =?UTF-8?q?cuda:=20increment=207=20=E2=80=94=20devic?= =?UTF-8?q?e-resident=20encoder-ahead=20seam=20(18=20fps=20threaded)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The encoder-ahead handoff moved 2×84 MB of neck floats through the host per frame (producer tensor_get -> Go slot buffers -> consumer tensor_set), which capped the overlap win at 1.04x. The device seam keeps everything on-GPU: the producer encodes with a CACHED graph (addendum-4; producer-stream CUDA graph replay) and same-instance-D2D's the three FPN levels into per-slot persistent device tensors (+ one producer sync so writes are complete before the Go channel signals — happens-before across threads); the consumer's new device-prefetch branch in edgetam_encode_image D2D-feeds them straight into state.neck_trk. Zero PCIe on the seam. Falls back to the host-float seam for host transport / CoreML (that path is untouched). capi: encode_slot tries the device seam first; track_slot prefers it per slot. Measured (RTX 4060, device transport): serial 79.2 ms/frame -> threaded 55.7 ms/frame = 18.0 fps (speedup 1.42x, was 1.04x with the host seam). J2-13 gate: threaded-vs-serial outputs BYTE-IDENTICAL over a 299-frame moving-subject A/B (boxes, scores, states, masks) — the second CUDA instance produces bit-identical necks, so encoder-ahead costs zero accuracy here. Co-Authored-By: Claude Opus 4.8 (1M context) --- coreml/edgetam_capi.cpp | 15 ++++-- sam3.cpp | 117 ++++++++++++++++++++++++++++++++++++++++ sam3.h | 9 ++++ 3 files changed, 137 insertions(+), 4 deletions(-) diff --git a/coreml/edgetam_capi.cpp b/coreml/edgetam_capi.cpp index a4ecc64..ff31f0a 100644 --- a/coreml/edgetam_capi.cpp +++ b/coreml/edgetam_capi.cpp @@ -226,8 +226,11 @@ extern "C" int edgetam_capi_encode_slot(edgetam_tracker_t h, int slot, t->neck[slot][2].data()) ? 1 : 0; #else // ggml encoder-ahead (CUDA): preprocess + encode on the producer's own - // backend instance; write the 3 neck levels into this slot's buffers. + // backend instance. Device transport keeps the neck in per-slot DEVICE + // tensors (zero PCIe); host transport writes this slot's float buffers. sam3_image img = make_image(rgb, w, hgt); + if (sam3_encoder_ahead_encode_dev(t->ggml_ea, *t->model, img, slot)) + return 1; return sam3_encoder_ahead_encode(t->ggml_ea, *t->model, img, t->neck[slot][0].data(), t->neck[slot][1].data(), @@ -250,9 +253,13 @@ extern "C" et_result edgetam_capi_track_slot(edgetam_tracker_t h, int slot, // The seam is backend-agnostic — the slot was filled by the CoreML // producer OR the ggml-CUDA encoder-ahead producer. if (t->prod_enc || t->ggml_ea) { - sam3_coreml_set_prefetched_neck(t->neck[slot][0].data(), - t->neck[slot][1].data(), - t->neck[slot][2].data()); + // Prefer the device-resident seam (slot filled by encode_dev); + // fall back to the host-float seam for CoreML / host transport. + if (!(t->ggml_ea && sam3_set_prefetched_neck_dev(t->ggml_ea, slot))) { + sam3_coreml_set_prefetched_neck(t->neck[slot][0].data(), + t->neck[slot][1].data(), + t->neck[slot][2].data()); + } } sam3_image img = make_image(rgb, w, hgt); sam3_result res = sam3_propagate_frame(*t->tracker, *t->state, *t->model, img); diff --git a/sam3.cpp b/sam3.cpp index c169325..b652d98 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -5449,6 +5449,12 @@ static void edgetam_build_repvit_graph(struct ggml_context* ctx, static std::vector s_prefetch_neck[3]; static bool s_prefetch_ready = false; +// Device-resident prefetch seam (CUDA device transport): the producer's +// per-slot persistent DEVICE tensors, handed to the consumer as pointers — +// the consumer D2D-copies them into state.neck_trk instead of round-tripping +// 2×84 MB of floats through the host. Same SPSC discipline as the host seam. +static struct { struct ggml_tensor* n[3]; bool ready; } s_prefetch_dev = {{nullptr, nullptr, nullptr}, false}; + void sam3_coreml_set_prefetched_neck(const float* n0, const float* n1, const float* n2) { const int Wd[3] = {256, 128, 64}, D = 256; const float* src[3] = {n0, n1, n2}; @@ -5604,6 +5610,28 @@ static bool edgetam_encode_image(sam3_state& state, // frame's neck into the prefetch seam. Load it into state and skip the whole // in-line encoder graph, mirroring the CoreML prefetch path above. Guarded to // the seam's fixed EdgeTAM@1024 contract (levels 256/128/64). + // Device-resident prefetch (increment 7): the producer already encoded + // this frame's neck into per-slot DEVICE tensors — three same-GPU D2D + // copies into state and the whole in-line encoder is skipped, with zero + // PCIe traffic. Requires the state buffers to exist (always true after + // the synchronous seed frame); otherwise falls through to the sync path. + if (s_prefetch_dev.ready && hp.is_edgetam() && img_size == 1024 && + model.transport.kind == SAM3_TRANSPORT_DEVICE) { + s_prefetch_dev.ready = false; + bool ok = state.buffer && state.pe_buf; + for (int i = 0; ok && i < 3; ++i) + ok = state.neck_trk[i] && s_prefetch_dev.n[i] && + ggml_nbytes(state.neck_trk[i]) == ggml_nbytes(s_prefetch_dev.n[i]); + if (ok) { + SAM3_TIME_SCOPE(state_update_ms); + for (int i = 0; i < 3; ++i) + sam3_stage_feed(model.transport, s_prefetch_dev.n[i], + state.neck_trk[i], "ea_neck"); + return true; + } + fprintf(stderr, "%s: device prefetch present but state not primed — sync encode\n", __func__); + } + if (s_prefetch_ready && hp.is_edgetam() && img_size == 1024) { static std::vector nbuf[3]; // consumer-thread-only (SPSC seam) for (int i = 0; i < 3; ++i) nbuf[i].swap(s_prefetch_neck[i]); @@ -5866,6 +5894,19 @@ static bool edgetam_encode_image(sam3_state& state, struct sam3_encoder_ahead { ggml_backend_t backend = nullptr; // producer-owned second backend instance ggml_gallocr_t galloc = nullptr; // producer-owned graph allocator + + // Device seam (increment 7): cached producer graph (addendum-4 contract — + // built once, recomputed in place; stable pointers let the producer's own + // stream reach CUDA-graph replay) + per-slot persistent neck tensors. + static const int SLOTS = 3; // == EtTracker::POOL + struct ggml_context* gctx = nullptr; + struct ggml_cgraph* graph = nullptr; + struct ggml_tensor* inp = nullptr; + struct ggml_tensor* fpn[3] = {}; + struct ggml_context* slot_ctx = nullptr; + ggml_backend_buffer_t slot_buf = nullptr; + struct ggml_tensor* slot_n[SLOTS][3] = {}; + bool slot_filled[SLOTS] = {}; }; sam3_encoder_ahead* sam3_encoder_ahead_create(const sam3_model& model) { @@ -5938,8 +5979,84 @@ bool sam3_encoder_ahead_encode(sam3_encoder_ahead* ea, const sam3_model& model, return ok; } +// Device-seam producer (increment 7): encode on the producer instance with a +// CACHED graph, then same-instance async D2D into this slot's persistent neck +// tensors + one producer-stream sync so the writes are complete before the Go +// channel signals the consumer (happens-before across threads). +bool sam3_encoder_ahead_encode_dev(sam3_encoder_ahead* ea, const sam3_model& model, + const sam3_image& image, int slot) { +#ifdef GGML_USE_CUDA + if (!ea || !ea->backend || slot < 0 || slot >= sam3_encoder_ahead::SLOTS) return false; + if (model.transport.kind != SAM3_TRANSPORT_DEVICE) return false; + const int img_size = 1024; // the prefetch seam's fixed EdgeTAM contract + + std::vector img_data = sam2_preprocess_image(image, img_size); + + if (!ea->graph) { + const size_t buf_size = ggml_tensor_overhead() * 16384 + ggml_graph_overhead() * 2; + struct ggml_init_params gparams = {buf_size, nullptr, true}; + ea->gctx = ggml_init(gparams); + if (!ea->gctx) return false; + ea->inp = ggml_new_tensor_4d(ea->gctx, GGML_TYPE_F32, img_size, img_size, 3, 1); + ggml_set_name(ea->inp, "input_image"); + ggml_set_input(ea->inp); + struct ggml_tensor* stage_outs[4] = {}; + edgetam_build_repvit_graph(ea->gctx, ea->inp, model, stage_outs); + struct ggml_tensor* fpn_outs[4] = {}; + edgetam_build_fpn_neck_graph(ea->gctx, stage_outs, model, fpn_outs); + for (int i = 0; i < 3; ++i) { ggml_set_output(fpn_outs[i]); ea->fpn[i] = fpn_outs[i]; } + ea->graph = ggml_new_graph_custom(ea->gctx, 32768, false); + for (int i = 0; i < 3; ++i) ggml_build_forward_expand(ea->graph, ea->fpn[i]); + if (!ggml_gallocr_alloc_graph(ea->galloc, ea->graph)) { + ggml_free(ea->gctx); ea->gctx = nullptr; ea->graph = nullptr; + return false; + } + // Per-slot persistent neck tensors: one ctx + one buffer for all slots. + struct ggml_init_params sp = {ggml_tensor_overhead() * 16, nullptr, true}; + ea->slot_ctx = ggml_init(sp); + for (int s = 0; s < sam3_encoder_ahead::SLOTS; ++s) + for (int i = 0; i < 3; ++i) + ea->slot_n[s][i] = ggml_dup_tensor(ea->slot_ctx, ea->fpn[i]); + ea->slot_buf = ggml_backend_alloc_ctx_tensors(ea->slot_ctx, ea->backend); + if (!ea->slot_buf) { + if (ea->slot_ctx) { ggml_free(ea->slot_ctx); ea->slot_ctx = nullptr; } + ggml_free(ea->gctx); ea->gctx = nullptr; ea->graph = nullptr; + return false; + } + fprintf(stderr, "%s: device-resident encoder-ahead seam active (%d slots)\n", + __func__, sam3_encoder_ahead::SLOTS); + } + + ggml_backend_tensor_set(ea->inp, img_data.data(), 0, img_data.size() * sizeof(float)); + if (!sam3_graph_compute(ea->backend, ea->graph, 4)) return false; + for (int i = 0; i < 3; ++i) + ggml_backend_tensor_copy_async(ea->backend, ea->backend, ea->fpn[i], ea->slot_n[slot][i]); + ggml_backend_synchronize(ea->backend); + ea->slot_filled[slot] = true; + return true; +#else + (void)ea; (void)model; (void)image; (void)slot; + return false; +#endif +} + +// Consumer side of the device seam: stash this slot's device tensors for +// edgetam_encode_image's device-prefetch branch (consumer thread only, SPSC). +bool sam3_set_prefetched_neck_dev(sam3_encoder_ahead* ea, int slot) { + if (!ea || slot < 0 || slot >= sam3_encoder_ahead::SLOTS || !ea->slot_filled[slot]) + return false; + ea->slot_filled[slot] = false; + for (int i = 0; i < 3; ++i) s_prefetch_dev.n[i] = ea->slot_n[slot][i]; + s_prefetch_dev.ready = true; + return true; +} + void sam3_encoder_ahead_destroy(sam3_encoder_ahead* ea) { if (!ea) return; + s_prefetch_dev.ready = false; + if (ea->slot_buf) ggml_backend_buffer_free(ea->slot_buf); + if (ea->slot_ctx) ggml_free(ea->slot_ctx); + if (ea->gctx) ggml_free(ea->gctx); if (ea->galloc) ggml_gallocr_free(ea->galloc); if (ea->backend) ggml_backend_free(ea->backend); delete ea; diff --git a/sam3.h b/sam3.h index 629428d..ef83540 100644 --- a/sam3.h +++ b/sam3.h @@ -398,6 +398,15 @@ bool sam3_encoder_ahead_encode(sam3_encoder_ahead* ea, const sam3 float* neck0, float* neck1, float* neck2); void sam3_encoder_ahead_destroy(sam3_encoder_ahead* ea); +// Device-resident encoder-ahead seam (CUDA device transport only): the +// producer encodes into per-slot persistent DEVICE tensors and the consumer +// D2D-copies them into state — no 2×84 MB host round trip per frame. Both +// return false when unavailable (host transport, non-CUDA build, bad slot, +// slot not device-filled); callers then fall back to the host-float seam. +bool sam3_encoder_ahead_encode_dev(sam3_encoder_ahead* ea, const sam3_model& model, + const sam3_image& image, int slot); +bool sam3_set_prefetched_neck_dev(sam3_encoder_ahead* ea, int slot); + /***************************************************************************** ** Test and Debug API ** From 3772b3537fac055c921a1366b47f9ea9bef35dfa Mon Sep 17 00:00:00 2001 From: Hunter Dolan Date: Fri, 10 Jul 2026 19:43:51 -0600 Subject: [PATCH 16/16] cuda: device frame preprocess (plan v1 step 5) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit coreml/sam3_preproc.cu: bit-exact CUDA twin of sam3_resize_bilinear (double chain, -fmad=false — load-bearing for the byte gate) + ImageNet normalize, writing f32 CHW straight into the encoder input tensor's device memory. Replaces ~14ms single-threaded CPU preprocess + 12.6MB f32 H2D with a ~2.7MB u8 upload on the device-transport path. SAM3_PREPROC_DEVICE: unset/1 = device (when transport==DEVICE); 0 = shipped CPU path byte-identical; 2 = VERIFY (both paths, memcmp, abort on mismatch). Fail-soft CPU fallback on any CUDA error. Encoder-ahead producer path unchanged (CPU preprocess; default-off). --- CMakeLists.txt | 10 ++++ coreml/sam3_preproc.cu | 124 +++++++++++++++++++++++++++++++++++++++++ sam3.cpp | 60 ++++++++++++++++++++ 3 files changed, 194 insertions(+) create mode 100644 coreml/sam3_preproc.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index a9b5a6b..e2d762f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -51,6 +51,16 @@ add_subdirectory(ggml) # sam3 static library add_library(sam3 STATIC sam3.cpp sam3.h) +if(NOT APPLE AND SAM3_CUDA) + # STEP-5 device frame preprocess (Windows parity plan v1): bit-exact + # CUDA twin of sam3_resize_bilinear + ImageNet normalize. -fmad=false is + # LOAD-BEARING for the byte gate (no FMA contraction in the double chain). + enable_language(CUDA) + target_sources(sam3 PRIVATE coreml/sam3_preproc.cu) + set_source_files_properties(coreml/sam3_preproc.cu PROPERTIES LANGUAGE CUDA) + set_property(SOURCE coreml/sam3_preproc.cu APPEND PROPERTY COMPILE_OPTIONS "-fmad=false") + target_compile_definitions(sam3 PRIVATE SAM3_PREPROC_CUDA) +endif() target_include_directories(sam3 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/stb diff --git a/coreml/sam3_preproc.cu b/coreml/sam3_preproc.cu new file mode 100644 index 0000000..524eb21 --- /dev/null +++ b/coreml/sam3_preproc.cu @@ -0,0 +1,124 @@ +// sam3_preproc.cu — device frame preprocess (Windows parity plan v1 STEP 5). +// +// Replicates sam3.cpp's CPU reference BIT-EXACTLY: +// sam3_resize_bilinear (torch bilinear align_corners=False, ALL arithmetic +// in double, round-to-nearest to uint8) → sam2_preprocess_image's +// ImageNet normalize in float, CHW. +// Bit-exactness contract: double precision per output pixel in the SAME +// operation order as the CPU loop; this file is compiled with -fmad=false so +// no FMA contraction can perturb the double chain (IEEE round-to-nearest per +// op on both MSVC x64 and sm_86/89 then yields identical bits); the float +// normalize is two float ops in the same order. Verified at runtime by +// SAM3_PREPROC_DEVICE=2 (compute both paths, memcmp, abort on mismatch). +// +// The kernel writes the f32 CHW result DIRECTLY into the encoder input +// tensor's device memory (ggml-CUDA tensor->data is a raw device pointer), +// replacing a 12.6 MB f32 H2D with a ~2.7 MB u8 upload and deleting the +// ~14ms single-threaded CPU preprocess from the hold critical path. + +#include +#include + +static __global__ void sam3_preproc_kernel(const uint8_t* __restrict__ src, + int src_w, int src_h, + float* __restrict__ dst, + int img_size, + float mean0, float mean1, float mean2, + float std0, float std1, float std2) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= img_size || y >= img_size) return; + + const double sx = (double)src_w / img_size; + const double sy = (double)src_h / img_size; + + double fy = (y + 0.5) * sy - 0.5; + if (fy < 0.0) fy = 0.0; + const int y0 = (int)fy; + const int y1 = (y0 < src_h - 1) ? y0 + 1 : y0; + const double wy = fy - y0; + const double wy0 = 1.0 - wy; + + double fx = (x + 0.5) * sx - 0.5; + if (fx < 0.0) fx = 0.0; + const int x0 = (int)fx; + const int x1 = (x0 < src_w - 1) ? x0 + 1 : x0; + const double wx = fx - x0; + const double wx0 = 1.0 - wx; + + const float mean[3] = {mean0, mean1, mean2}; + const float std_d[3] = {std0, std1, std2}; + const int plane = img_size * img_size; + + for (int c = 0; c < 3; ++c) { + const double p00 = src[(y0 * src_w + x0) * 3 + c]; + const double p01 = src[(y0 * src_w + x1) * 3 + c]; + const double p10 = src[(y1 * src_w + x0) * 3 + c]; + const double p11 = src[(y1 * src_w + x1) * 3 + c]; + // Same association order as the CPU reference. + double v = wy0 * (wx0 * p00 + wx * p01) + + wy * (wx0 * p10 + wx * p11); + int iv = (int)(v + 0.5); + if (iv < 0) iv = 0; + if (iv > 255) iv = 255; + // CPU reference: float division by 255.0f, then (v-mean)/std in float. + float fv = (float)(uint8_t)iv / 255.0f; + dst[c * plane + y * img_size + x] = (fv - mean[c]) / std_d[c]; + } +} + +// Identity fast path (src already img_size×img_size): normalize only, +// matching the CPU reference's no-resize branch. +static __global__ void sam3_preproc_norm_kernel(const uint8_t* __restrict__ src, + float* __restrict__ dst, + int img_size, + float mean0, float mean1, float mean2, + float std0, float std1, float std2) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= img_size || y >= img_size) return; + const float mean[3] = {mean0, mean1, mean2}; + const float std_d[3] = {std0, std1, std2}; + const int plane = img_size * img_size; + for (int c = 0; c < 3; ++c) { + float fv = (float)src[(y * img_size + x) * 3 + c] / 255.0f; + dst[c * plane + y * img_size + x] = (fv - mean[c]) / std_d[c]; + } +} + +// C entry point (called from sam3.cpp under GGML_USE_CUDA + device +// transport). dst_dev = the encoder input tensor's device pointer (f32 CHW, +// img_size²·3). Uses an owned staging buffer + stream; host-synchronizes +// before returning so the subsequent graph launch (any stream) observes the +// writes. Returns false on any CUDA error — caller falls back to the CPU +// path (fail-soft: preprocess must never kill a frame). +extern "C" bool sam3_cuda_preprocess_imagenet(const uint8_t* rgb, int src_w, int src_h, + float* dst_dev, int img_size) { + static uint8_t* s_stage = nullptr; + static size_t s_stage_cap = 0; + const size_t need = (size_t)src_w * src_h * 3; + if (need > s_stage_cap) { + if (s_stage) cudaFree(s_stage); + if (cudaMalloc(&s_stage, need) != cudaSuccess) { s_stage = nullptr; s_stage_cap = 0; return false; } + s_stage_cap = need; + } + if (cudaMemcpy(s_stage, rgb, need, cudaMemcpyHostToDevice) != cudaSuccess) return false; + + dim3 blk(16, 16); + dim3 grd((img_size + 15) / 16, (img_size + 15) / 16); + const float mean[3] = {0.485f, 0.456f, 0.406f}; + const float std_d[3] = {0.229f, 0.224f, 0.225f}; + if (src_w == img_size && src_h == img_size) { + sam3_preproc_norm_kernel<<>>(s_stage, dst_dev, img_size, + mean[0], mean[1], mean[2], + std_d[0], std_d[1], std_d[2]); + } else { + sam3_preproc_kernel<<>>(s_stage, src_w, src_h, dst_dev, img_size, + mean[0], mean[1], mean[2], + std_d[0], std_d[1], std_d[2]); + } + if (cudaGetLastError() != cudaSuccess) return false; + // Host sync: the caller launches the encoder graph after this returns; + // host ordering guarantees visibility across streams. + return cudaDeviceSynchronize() == cudaSuccess; +} diff --git a/sam3.cpp b/sam3.cpp index b652d98..b65350d 100644 --- a/sam3.cpp +++ b/sam3.cpp @@ -955,6 +955,14 @@ struct edgetam_perceiver { enum sam3_transport_kind { SAM3_TRANSPORT_HOST = 0, SAM3_TRANSPORT_DEVICE = 1 }; +#if defined(GGML_USE_CUDA) && defined(SAM3_PREPROC_CUDA) +// STEP-5 device frame preprocess (coreml/sam3_preproc.cu): bit-exact CUDA +// twin of sam3_resize_bilinear + ImageNet normalize, writing f32 CHW into +// the encoder input tensor's device memory. Host-synchronizes internally. +extern "C" bool sam3_cuda_preprocess_imagenet(const uint8_t* rgb, int src_w, int src_h, + float* dst_dev, int img_size); +#endif + struct sam3_transport { sam3_transport_kind kind = SAM3_TRANSPORT_HOST; ggml_backend_t backend = nullptr; @@ -5641,11 +5649,32 @@ static bool edgetam_encode_image(sam3_state& state, // ── Preprocess (same ImageNet normalization as SAM2) ───────────────── // RFD 0011 U0: preprocess timed separately from the graph build/alloc/compute. + // STEP-5 (plan v1): on the CUDA device-transport path the resize+normalize + // moves onto the GPU (coreml/sam3_preproc.cu — bit-exact double-chain twin + // of sam3_resize_bilinear, -fmad=false) and the 12.6MB f32 upload becomes a + // ~2.7MB u8 upload. SAM3_PREPROC_DEVICE: unset/1 = device path when + // available; 0 = CPU path (byte-identical shipped behavior); 2 = VERIFY + // (both paths, memcmp, abort on mismatch — the byte gate). std::vector img_data; + bool preproc_on_device = false; +#if defined(GGML_USE_CUDA) && defined(SAM3_PREPROC_CUDA) { + const char* pd = getenv("SAM3_PREPROC_DEVICE"); + const int pd_mode = pd ? atoi(pd) : 1; + if (pd_mode != 0 && model.transport.kind == SAM3_TRANSPORT_DEVICE) { + preproc_on_device = true; + if (pd_mode == 2) { + SAM3_TIME_SCOPE(preprocess_ms); + img_data = sam2_preprocess_image(image, img_size); // reference for VERIFY + } + } + } +#endif + if (!preproc_on_device) { SAM3_TIME_SCOPE(preprocess_ms); img_data = sam2_preprocess_image(image, img_size); } + (void)preproc_on_device; // ── Build graph ────────────────────────────────────────────────────── // RFD 0011 U0: graph CONSTRUCTION region (image_encoder_build_ms). U1 will @@ -5742,7 +5771,38 @@ static bool edgetam_encode_image(sam3_state& state, } // end build path (egc_hit skips straight to the input upload) // Set input image +#if defined(GGML_USE_CUDA) && defined(SAM3_PREPROC_CUDA) + if (preproc_on_device) { + SAM3_TIME_SCOPE(preprocess_ms); + bool pd_ok = sam3_cuda_preprocess_imagenet(image.data.data(), image.width, + image.height, (float*)inp->data, + img_size); + if (!pd_ok) { + // Fail-soft: fall back to the CPU path for this frame (log once). + static bool pd_warned = false; + if (!pd_warned) { fprintf(stderr, "%s: device preprocess failed - CPU fallback\n", __func__); pd_warned = true; } + if (img_data.empty()) img_data = sam2_preprocess_image(image, img_size); + ggml_backend_tensor_set(inp, img_data.data(), 0, img_data.size() * sizeof(float)); + } else if (!img_data.empty()) { + // VERIFY mode: download the kernel's output and byte-compare. + std::vector dev_out(img_data.size()); + ggml_backend_tensor_get(inp, dev_out.data(), 0, dev_out.size() * sizeof(float)); + if (memcmp(dev_out.data(), img_data.data(), img_data.size() * sizeof(float)) != 0) { + size_t bad = 0, first = (size_t)-1; + for (size_t i = 0; i < img_data.size(); ++i) + if (dev_out[i] != img_data[i]) { ++bad; if (first == (size_t)-1) first = i; } + fprintf(stderr, "%s: SAM3_PREPROC_DEVICE=2 BYTE GATE FAIL: %zu/%zu floats differ (first %zu: dev %.9g cpu %.9g)\n", + __func__, bad, img_data.size(), first, (double)dev_out[first], (double)img_data[first]); + abort(); + } + fprintf(stderr, "%s: SAM3_PREPROC_DEVICE=2 byte gate PASS (%zu floats)\n", __func__, img_data.size()); + } + } else { + ggml_backend_tensor_set(inp, img_data.data(), 0, img_data.size() * sizeof(float)); + } +#else ggml_backend_tensor_set(inp, img_data.data(), 0, img_data.size() * sizeof(float)); +#endif // ── Compute ────────────────────────────────────────────────────────── // RFD 0011 U0: backend compute region (image_encoder_compute_ms). This is