CLAUDE: D5a - k-slot device image-set ring + async conditioning on the copy stream

Overlap-ring design ratified 07/17 (handoffs/2026-07-17_pose_d5_overlap_ring_design.md):
- TP_MAX_IMG_SETS=4 ring; set 0 = the setup allocation, others lazy clones
  (same pitch asserted); images_h/gpu_images always alias the ACTIVE set so
  every consumer (conditioning, granular convert, pose_mchain_fill snapshot)
  follows the ring with zero call-site changes.
- tp_proc_img_set_select / _active / _ready_record; tp_proc_exec_conditioning_async
  = float two-map kernel on the bayer copy stream, ready event after it (ONE
  event = uploaded+conditioned); bayer_fence extended to cover the active set.
- test_img_set_ring_jna (case img_ring): sync-set0 vs async-set1 conditioning
  bit-exact + both set-independence directions; run_cases.sh ALL PASS;
  compute-sanitizer memcheck 0 errors.
Co-authored-by: Claude Opus 4.8 <claude-opus-4-8@anthropic.com>
Co-authored-by: Claude Fable 5 <claude-fable-5@anthropic.com>

cases.list.tmpl

@@ -65,6 +65,7 @@ case peak_recenter   test_peak             ${REPO}/testdata/synth4/peak_recenter
 case avg_td          test_avg_td           ${REPO}/testdata/synth4/avg_td        --tol 0
 case avg_td_oob      test_avg_td_oob       ${REPO}/testdata/synth4/avg_td_oob    --tol 0
 case bayer_staged    test_bayer_staged_jna ${REPO}                               --tol 0     # rung B4: synthetic, --data unused
+case img_ring        test_img_set_ring_jna ${REPO}                               --tol 0     # rung D5a: k-slot image-set ring, sync-vs-async conditioning + set independence bit-exact (synthetic) # By Claude on 07/17/2026
 case dp_cycles       test_pose_lma_dp_cycles_jna ${REPO}                         --tol 0     # rung D1: synthetic, --data unused; DP K-cycle chain bit-exact vs host cycles # By Claude on 07/17/2026
 case measure_dp      test_pose_measure_dp_jna    ${MODEL}/testdata/pose_corr     --tol 0     # rung D2: device measure chain bit-exact vs B3 host-batched entry, keyed by packed index # By Claude on 07/17/2026
 case scene_dp        test_pose_scene_dp_jna      ${MODEL}/testdata/pose_corr     --tol 0     # rung D3: full per-scene DP parent bit-exact vs host-driven B3 loop (full+light chains) # By Claude on 07/17/2026

jna/test_img_set_ring_jna.cu

+// test_img_set_ring_jna.cu - validate the D5a k-slot device image-set ring
+// (libtileproc.so): overlap-ring design ratified 2026-07-17.
+//   1. SLOT EQUIVALENCE: the same raw input conditioned on set 0 via the SYNC
+//      float path (tp_proc_exec_conditioning) and on set 1 via the ASYNC
+//      copy-stream path (tp_proc_exec_conditioning_async) must read back
+//      BIT-EXACT equal - same kernel, same input, different buffer only.
+//   2. SET INDEPENDENCE: after set-1 traffic, set 0 still holds its data
+//      bit-exact; after re-uploading DIFFERENT raw into set 0 (staged, no
+//      conditioning), set 1 still holds its conditioned data bit-exact.
+//   3. ACTIVE-SET BOOKKEEPING: tp_proc_img_set_active follows selects.
+// All comparisons tol 0 by design. Fully synthetic (no case data).
+//
+// Build (after jna/build_lib.sh):
+//   cd tile_processor_gpu && /usr/local/cuda-12.8/bin/nvcc -std=c++17 -O2 \
+//     -I src -I $HOME/git/cuda-samples/Common \
+//     jna/test_img_set_ring_jna.cu \
+//     -o tests_bin/test_img_set_ring_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_img_set_ring_jna   (cwd = repo root; no --data needed)
+//
+//  Created on: Jul 17, 2026
+//      Author: Claude (Anthropic), for Elphel
+#include <cstdio>
+#include <cstring>
+#include <cstdlib>
+
+extern "C" {
+  void* tp_create_module(const char*, const char*);
+  const char* tp_last_error();
+  void* tp_proc_create(void*);
+  int   tp_proc_setup(void*, int,int,int,int,int,int);
+  float* tp_proc_bayer_staging(void*, int);
+  int   tp_proc_set_image_staged(void*, int, int);
+  int   tp_proc_set_conditioning_maps(void*, const float*, const float*, int);
+  int   tp_proc_exec_conditioning(void*, int, double);
+  int   tp_proc_img_set_select(void*, int);
+  int   tp_proc_img_set_active(void*);
+  int   tp_proc_img_set_ready_record(void*);
+  int   tp_proc_exec_conditioning_async(void*);
+  int   tp_proc_get_image(void*, int, float*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+static long compare_bits(const float* a, const float* b, size_t n){
+  long bad=0; for(size_t i=0;i<n;i++) if(memcmp(&a[i],&b[i],sizeof(float))) bad++; return bad;
+}
+static void fill_raw(float* dst, int cam, int salt, size_t npix){
+  for(size_t i=0;i<npix;i++) dst[i]=(float)(cam*1000.0+salt*0.001+i*0.000123);
+}
+
+int main(int argc, char** argv){
+  const char* srcd  = "src";
+  const char* devrt = "/usr/local/cuda-12.8/lib64/libcudadevrt.a";
+  for (int i = 1; i < argc; i++){
+    if      (!strcmp(argv[i],"--src")   && (i+1<argc)) srcd  = argv[++i];
+    else if (!strcmp(argv[i],"--devrt") && (i+1<argc)) devrt = argv[++i];
+    else if (!strcmp(argv[i],"--data")  && (i+1<argc)) ++i;  // run_cases.sh uniformity (unused: synthetic)
+    else if (!strcmp(argv[i],"--tol")   && (i+1<argc)) ++i;  // tol-0 by design
+    else { fprintf(stderr,"Usage: %s [--src <kernel src>] [--devrt <libcudadevrt.a>]\n", argv[0]); return 1; }
+  }
+  const int NC=16, W=640, H=512; const size_t NPIX=(size_t)W*H;
+  void* m = tp_create_module(srcd, devrt);
+  if(!m){ printf("FAIL module: %s\n", tp_last_error()); return 1; }
+  void* p = tp_proc_create(m);
+  if(tp_proc_setup(p, NC, 1, W, H, 0, 0)){ printf("FAIL setup: %s\n", tp_last_error()); return 1; }
+
+  // non-trivial per-pixel two-map calibration (production float model)
+  float* scales  = new float[(size_t)NC*NPIX];
+  float* offsets = new float[(size_t)NC*NPIX];
+  for(int c=0;c<NC;c++) for(size_t i=0;i<NPIX;i++){
+    scales [(size_t)c*NPIX+i]=(float)(1.0+c*0.01+i*1.0e-8);
+    offsets[(size_t)c*NPIX+i]=(float)(c*0.5+i*1.0e-7);
+  }
+  if(tp_proc_set_conditioning_maps(p, scales, offsets, (int)NPIX)){
+    printf("FAIL set_conditioning_maps: %s\n", tp_last_error()); return 1; }
+
+  long bad_total = 0;
+  float* got = new float[NPIX];
+  float** ref = new float*[NC];
+  for(int c=0;c<NC;c++) ref[c]=new float[NPIX];
+
+  // ---- 1. sync reference on set 0 ----
+  if(tp_proc_img_set_active(p)!=0){ printf("FAIL initial active set != 0\n"); return 1; }
+  float* staging = tp_proc_bayer_staging(p, 0);
+  if(!staging){ printf("FAIL staging: %s\n", tp_last_error()); return 1; }
+  for(int c=0;c<NC;c++){
+    fill_raw(staging + (size_t)c*NPIX, c, 1, NPIX);
+    if(tp_proc_set_image_staged(p, c, 0)){ printf("FAIL staged cam %d: %s\n", c, tp_last_error()); return 1; }
+  }
+  if(tp_proc_exec_conditioning(p, 0, 0.0)){ printf("FAIL sync conditioning: %s\n", tp_last_error()); return 1; }
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, ref[c])){ printf("FAIL ref readback cam %d\n", c); return 1; }
+  }
+
+  // ---- 2. async path on set 1, SAME raw input ----
+  if(tp_proc_img_set_select(p, 1)){ printf("FAIL select set 1: %s\n", tp_last_error()); return 1; }
+  if(tp_proc_img_set_active(p)!=1){ printf("FAIL active set != 1 after select\n"); return 1; }
+  staging = tp_proc_bayer_staging(p, 1);
+  if(!staging){ printf("FAIL staging slot 1: %s\n", tp_last_error()); return 1; }
+  for(int c=0;c<NC;c++){
+    fill_raw(staging + (size_t)c*NPIX, c, 1, NPIX);   // same salt = same raw
+    if(tp_proc_set_image_staged(p, c, 1)){ printf("FAIL staged set1 cam %d: %s\n", c, tp_last_error()); return 1; }
+  }
+  if(tp_proc_exec_conditioning_async(p)){ printf("FAIL async conditioning: %s\n", tp_last_error()); return 1; }
+  long bad_eq=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL set1 readback cam %d\n", c); return 1; } // fences the ready event
+    bad_eq += compare_bits(got, ref[c], NPIX);
+  }
+  printf("slot equivalence (sync set0 vs async set1): %d sensors x %zu px, bit-mismatches=%ld -> %s\n",
+         NC, NPIX, bad_eq, bad_eq?"FAIL":"PASS");
+  bad_total += bad_eq;
+
+  // ---- 3a. set 0 untouched by set-1 traffic ----
+  if(tp_proc_img_set_select(p, 0)){ printf("FAIL re-select set 0: %s\n", tp_last_error()); return 1; }
+  long bad_s0=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL set0 recheck cam %d\n", c); return 1; }
+    bad_s0 += compare_bits(got, ref[c], NPIX);
+  }
+  printf("set 0 independence after set-1 traffic: bit-mismatches=%ld -> %s\n", bad_s0, bad_s0?"FAIL":"PASS");
+  bad_total += bad_s0;
+
+  // ---- 3b. different raw into set 0 (staged only + ready_record, no kernel);
+  //          set 1 must still hold its conditioned data ----
+  staging = tp_proc_bayer_staging(p, 0);
+  for(int c=0;c<NC;c++){
+    fill_raw(staging + (size_t)c*NPIX, c, 9, NPIX);   // different salt
+    if(tp_proc_set_image_staged(p, c, 0)){ printf("FAIL restage set0 cam %d: %s\n", c, tp_last_error()); return 1; }
+  }
+  if(tp_proc_img_set_ready_record(p)){ printf("FAIL ready_record: %s\n", tp_last_error()); return 1; }
+  if(tp_proc_img_set_select(p, 1)){ printf("FAIL final select set 1: %s\n", tp_last_error()); return 1; }
+  long bad_s1=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL set1 recheck cam %d\n", c); return 1; }
+    bad_s1 += compare_bits(got, ref[c], NPIX);
+  }
+  printf("set 1 independence after set-0 re-upload: bit-mismatches=%ld -> %s\n", bad_s1, bad_s1?"FAIL":"PASS");
+  bad_total += bad_s1;
+
+  tp_proc_destroy(p); tp_destroy_module(m);
+  delete [] got; for(int c=0;c<NC;c++) delete [] ref[c]; delete [] ref;
+  delete [] scales; delete [] offsets;
+  printf("test_img_set_ring_jna: %s\n", bad_total?"FAIL":"PASS (bit-exact)");
+  return bad_total?1:0;
+}

jna/tp_jna.cpp

@@ -594,6 +594,7 @@ int tp_tex_selftest(TpModule* m, int lwir, const char* data_root,
 // erase_clt, ref_scene (separate CLT buffer for inter-scene corr).
 // See handoffs/2026-06-25_convert_direct-fragile-features.md.
 // ============================================================================
+#define TP_MAX_IMG_SETS 4 // D5a device image-set ring capacity (k=2 used now) // By Claude on 07/17/2026
 struct TpProc {
   TpModule* mod;
   int num_cams, num_colors, img_w, img_h, tilesx, tilesy, kernels_hor, kernels_vert, kern_tiles, kern_size, slice;
@@ -756,6 +757,21 @@ struct TpProc {
   cudaStream_t bayer_copy_stream;
   cudaEvent_t  bayer_copy_event;
   bool         bayer_copy_pending;
+  // D5a (overlap-ring design ratified 2026-07-17): k-slot ring of device image
+  // sets. Set 0 = the tp_proc_setup allocation; sets 1..k-1 are lazily
+  // allocated clones (same pitch). images_h/gpu_images always ALIAS the ACTIVE
+  // set, so every existing consumer (conditioning, granular convert, the DP
+  // measure-chain descriptor snapshot in pose_mchain_fill) follows the ring
+  // with zero per-call-site changes, and a scene is consumed from the same set
+  // it was conditioned into (trivial host-driven fallback). img_set_ready[s]
+  // is recorded on the copy stream after the last upload/async-conditioning
+  // launch targeting set s; blocking consumers cover it via bayer_fence, the
+  // D5b launch path turns it into a GPU-side stream-wait. By Claude on 07/17/2026.
+  std::vector<float*> img_sets_h[TP_MAX_IMG_SETS];
+  float**             gpu_img_sets[TP_MAX_IMG_SETS];
+  cudaEvent_t         img_set_ready[TP_MAX_IMG_SETS];
+  bool                img_set_ready_pending[TP_MAX_IMG_SETS];
+  int                 img_set_active;
 };
 
 extern "C" {
@@ -818,6 +834,8 @@ TpProc* tp_proc_create(TpModule* m){
   p->pose_task_capacity=0; p->pose_task_count=0; // By Claude on 07/16/2026 (rung B1)
   p->bayer_staging[0]=p->bayer_staging[1]=nullptr; // By Claude on 07/16/2026 (rung B4)
   p->bayer_copy_stream=nullptr; p->bayer_copy_event=nullptr; p->bayer_copy_pending=false;
+  for(int s=0;s<TP_MAX_IMG_SETS;s++){ p->gpu_img_sets[s]=nullptr; p->img_set_ready[s]=nullptr; p->img_set_ready_pending[s]=false; } // By Claude on 07/17/2026 (rung D5a)
+  p->img_set_active=0; // By Claude on 07/17/2026 (rung D5a)
   p->pcyc_valid=0; // By Claude on 07/16/2026 (rung B3)
   return p;
 }
@@ -844,6 +862,7 @@ int tp_proc_setup(TpProc* p, int num_cams, int num_colors, int img_w, int img_h,
   p->gpu_kernels=copyalloc_pointers_gpu(p->kernels_h.data(),num_cams);
   p->gpu_kernel_offsets=copyalloc_pointers_gpu(p->offs_h.data(),num_cams);
   p->gpu_images=copyalloc_pointers_gpu(p->images_h.data(),num_cams);
+  p->img_sets_h[0]=p->images_h; p->gpu_img_sets[0]=p->gpu_images; p->img_set_active=0; // ring set 0 = the setup allocation // By Claude on 07/17/2026 (rung D5a)
   p->gpu_clt=copyalloc_pointers_gpu(p->clt_h.data(),num_cams);
   p->gpu_clt_ref=copyalloc_pointers_gpu(p->clt_ref_h.data(),num_cams);
   cudaMalloc((void**)&p->active,(size_t)p->tilesx*p->tilesy*sizeof(int));
@@ -872,6 +891,8 @@ int tp_proc_set_kernel_offsets(TpProc* p, int cam, const float* d, int n){ if(!p
 // By Claude on 07/16/2026.
 static void bayer_fence(TpProc* p){
   if(p->bayer_copy_pending){ cudaEventSynchronize(p->bayer_copy_event); p->bayer_copy_pending=false; }
+  const int s=p->img_set_active; // D5a: also cover the active set's async conditioning // By Claude on 07/17/2026
+  if(p->img_set_ready_pending[s]){ cudaEventSynchronize(p->img_set_ready[s]); p->img_set_ready_pending[s]=false; }
 }
 // Pinned host staging for slot (0/1): num_cams*img_w*img_h floats, lazily
 // allocated (non-RT users never pay the pinned memory). Also lazily creates
@@ -907,6 +928,81 @@ int tp_proc_set_image_staged(TpProc* p, int cam, int slot){
   p->bayer_copy_pending=true;
   return 0;
 }
+// ---- D5a (overlap-ring design ratified 2026-07-17): k-slot device image-set
+// ring + async conditioning on the copy stream. By Claude on 07/17/2026. ----
+// Select (lazily allocating) device image set 'set' as the ACTIVE set. All
+// image consumers and the staged uploads follow the active set from now on.
+// Java call-order contract (single-threaded): flip AFTER launching the chain
+// that consumes the previous set - descriptor fills snapshot gpu_images.
+int tp_proc_img_set_select(TpProc* p, int set){
+  if(!p||set<0||set>=TP_MAX_IMG_SETS){ seterr("img_set_select: bad set %d", set); return -1; }
+  cuCtxSetCurrent(p->mod->ctx);
+  if(!p->gpu_img_sets[set]){
+    std::vector<float*> imgs((size_t)p->num_cams,nullptr);
+    size_t dstride=0;
+    for(int c=0;c<p->num_cams;c++){
+      imgs[c]=alloc_image_gpu(&dstride, p->img_w, p->img_h);
+      if(!imgs[c]||(dstride!=p->dstride_img)){
+        for(int k=0;k<=c;k++) if(imgs[k]) cudaFree(imgs[k]);
+        seterr("img_set_select: alloc failed set=%d cam=%d (stride %zu vs %zu)",
+               set,c,dstride,p->dstride_img);
+        return -2; }
+    }
+    float** garr=copyalloc_pointers_gpu(imgs.data(),p->num_cams);
+    if(!garr){
+      for(int c=0;c<p->num_cams;c++) cudaFree(imgs[c]);
+      seterr("img_set_select: pointer-array alloc failed set=%d",set); return -2; }
+    p->img_sets_h[set]=imgs;
+    p->gpu_img_sets[set]=garr;
+  }
+  if(!p->img_set_ready[set] &&
+     cudaEventCreateWithFlags(&p->img_set_ready[set], cudaEventDisableTiming)!=cudaSuccess){
+    p->img_set_ready[set]=nullptr; seterr("img_set_select: ready-event create failed set=%d",set); return -3; }
+  p->images_h=p->img_sets_h[set];
+  p->gpu_images=p->gpu_img_sets[set];
+  p->img_set_active=set;
+  return 0;
+}
+int tp_proc_img_set_active(TpProc* p){ return p?p->img_set_active:-1; }
+// Record the active set's ready event on the copy stream: marks "all staged
+// uploads for this set landed". For the CPU-conditioned fallback (no async
+// kernel); the GPU path's tp_proc_exec_conditioning_async records it itself.
+int tp_proc_img_set_ready_record(TpProc* p){
+  if(!p){ seterr("img_set_ready_record: null"); return -1; }
+  const int s=p->img_set_active;
+  if(!p->img_set_ready[s]||!p->bayer_copy_stream){
+    seterr("img_set_ready_record: set %d not prepared (select + staging first)", s); return -2; }
+  cuCtxSetCurrent(p->mod->ctx);
+  if(cudaEventRecord(p->img_set_ready[s], p->bayer_copy_stream)!=cudaSuccess){
+    seterr("img_set_ready_record: record failed set=%d",s); return -3; }
+  p->img_set_ready_pending[s]=true;
+  return 0;
+}
+// Async in-place conditioning of the ACTIVE image set on the copy stream
+// (production float two-map kernel only - the double oracle stays sync-path).
+// Stream order places the kernel after that set's staged uploads; the set's
+// ready event is recorded after it, so ONE event marks "uploaded+conditioned".
+// No host synchronization anywhere - this is the D5 overlap workhorse (runs
+// while the previous scene's DP chain occupies the legacy stream).
+int tp_proc_exec_conditioning_async(TpProc* p){
+  if(!p){ seterr("exec_conditioning_async: null"); return -1; }
+  if(!p->cond_maps_ready){ seterr("exec_conditioning_async: calibration not uploaded"); return -2; }
+  const int s=p->img_set_active;
+  if(!p->img_set_ready[s]||!p->bayer_copy_stream){
+    seterr("exec_conditioning_async: set %d not prepared (select + staging first)", s); return -3; }
+  cuCtxSetCurrent(p->mod->ctx);
+  CUfunction f=getfun(p->mod,"condition_lwir_f"); if(!f){ seterr("condition_lwir_f missing"); return -4; }
+  int nc=p->num_cams,w=p->img_w,h=p->img_h; size_t stride=p->dstride_img/sizeof(float);
+  void* a[]={&nc,&w,&h,&stride,&p->gpu_images,&p->cond_scale_maps,&p->cond_offset_maps};
+  CUresult cr=cuLaunchKernel(f,(w+31)/32,(h+7)/8,nc,32,8,1,0,
+                             (CUstream)p->bayer_copy_stream,a,nullptr);
+  if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es);
+    seterr("launch condition_lwir_f (async) -> %d (%s)",cr,es); return -5; }
+  if(cudaEventRecord(p->img_set_ready[s], p->bayer_copy_stream)!=cudaSuccess){
+    seterr("exec_conditioning_async: ready record failed set=%d",s); return -6; }
+  p->img_set_ready_pending[s]=true;
+  return 0;
+}
 // pitched image upload (reuses harness update_image_gpu: dstride in floats)
 int tp_proc_set_image(TpProc* p, int cam, const float* d){ if(!p||cam<0||cam>=p->num_cams)return -1; cuCtxSetCurrent(p->mod->ctx);
   bayer_fence(p); // WAW order vs in-flight staged copies // By Claude on 07/16/2026 (rung B4)
@@ -3060,6 +3156,15 @@ void tp_proc_destroy(TpProc* p){
   if(p->bayer_copy_event) cudaEventDestroy(p->bayer_copy_event);
   if(p->bayer_staging[0]) cudaFreeHost(p->bayer_staging[0]);
   if(p->bayer_staging[1]) cudaFreeHost(p->bayer_staging[1]);
+  // D5a: images_h/gpu_images alias the ACTIVE ring set - restore the set-0
+  // alias so the legacy frees below release set 0, then free sets 1..k-1 and
+  // the ready events explicitly. By Claude on 07/17/2026.
+  if(p->gpu_img_sets[0]){ p->images_h=p->img_sets_h[0]; p->gpu_images=p->gpu_img_sets[0]; }
+  for(int s=1;s<TP_MAX_IMG_SETS;s++){
+    for(size_t c=0;c<p->img_sets_h[s].size();c++) cudaFree(p->img_sets_h[s][c]);
+    if(p->gpu_img_sets[s]) cudaFree(p->gpu_img_sets[s]);
+  }
+  for(int s=0;s<TP_MAX_IMG_SETS;s++) if(p->img_set_ready[s]) cudaEventDestroy(p->img_set_ready[s]);
   for(int c=0;c<p->num_cams;c++){ if(c<(int)p->kernels_h.size())cudaFree(p->kernels_h[c]); if(c<(int)p->offs_h.size())cudaFree(p->offs_h[c]);
     if(c<(int)p->images_h.size())cudaFree(p->images_h[c]); if(c<(int)p->clt_h.size())cudaFree(p->clt_h[c]); if(c<(int)p->clt_ref_h.size())cudaFree(p->clt_ref_h[c]); }
   cudaFree(p->gpu_kernels); cudaFree(p->gpu_kernel_offsets); cudaFree(p->gpu_images); cudaFree(p->gpu_clt); cudaFree(p->gpu_clt_ref);