CLAUDE: JNA export tp_proc_exec_consolidate - on-device TD consolidation chain

Wire the frozen v1 chain (index_consolidate -> consolidate_oob ->
clt_average_sensors_list) into the production JNA surface (roadmap 2c ->
step 4 JNA hookup, Andrey's ruling: JNA-only, no JCuda launch code):
- tp_jna.cpp: tp_consolidate.h/.cu added to the NVRTC module (23 kernels);
  new export tp_proc_exec_consolidate = host orchestration of the chain on
  the RESIDENT per-cam CLT (contiguous DtoD gather, no D2H/H2D of TD data;
  only the flattened task streams cross the boundary); average lands in the
  cam-0 slot (what the single conj-multiply correlates); lazy scratch,
  freed in tp_proc_destroy; stats {pairs, surviving tiles, misaligned}.
- tp_consolidate.h/.cu: JCUDA include guards (NVRTC concatenation has no
  include paths) + NAN/INFINITY fallbacks (NVRTC has no math.h macros).
- jna/test_proc_consolidate.cu: standalone validation of the export vs the
  real-scene avg_td_oob oracle case - PASS bit-exact (max|diff|=0,
  5120 pairs -> 5056 tiles == oracle, 0 misaligned).
Regressions: tests_bin/test_avg_td, test_avg_td_oob PASS at tol 0 after
the guard changes; module smoke = 23/23 kernels resolved.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

jna/test_proc_consolidate.cu

+// test_proc_consolidate.cu - validate the JNA production export
+// tp_proc_exec_consolidate (the granular TpProc surface, libtileproc.so)
+// against a real-scene oracle case exported by AvgTdExport.exportOob (the
+// same case tests_bin/test_avg_td_oob passes at tol 0). Proves the
+// production-facing path end-to-end: set_clt per cam -> exec_consolidate
+// (contiguous DtoD gather + 3-kernel chain + avg->cam0) -> get_clt(0) ==
+// expected_td_avg bit-exactly; stats == oracle (5120 pairs -> 5056 tiles).
+// First PASS 07/12/2026 on the v013 case (max|diff|=0, 0 misaligned).
+//
+// 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_proc_consolidate.cu src/tests/tp_test_data.cu \
+//     -o tests_bin/test_proc_consolidate -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_proc_consolidate [case_dir]
+//
+//  Created on: Jul 12, 2026
+//      Author: Claude (Anthropic), for Elphel
+#include <cstdio>
+#include <cstring>
+#include <cmath>
+#include "tests/tp_test_data.h"
+
+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);
+  int   tp_proc_set_clt(void*, int, const float*, int);
+  int   tp_proc_get_clt(void*, int, int, float*);
+  int   tp_proc_exec_consolidate(void*, const float*, const float*, int, int, float, float, int*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+int main(int argc, char** argv){
+  const char* dir = (argc>1)? argv[1] :
+    "/home/elphel/lwir16-proc/LV/models/models_1773135103-1773135664/1773135476_186641-CENTER/v013/testdata/avg_td_oob";
+  void* m = tp_create_module("/home/elphel/git/tile_processor_gpu/src",
+                             "/usr/local/cuda-12.8/lib64/libcudadevrt.a");
+  if(!m){ printf("FAIL module: %s\n", tp_last_error()); return 1; }
+  TpTestData data(dir);
+  int num_sens  = data.paramInt("num_sensors", 16);
+  int tilesx    = data.paramInt("tilesx", 80);
+  int tilesy    = data.paramInt("tilesy", 64);
+  int colors    = data.paramInt("colors", 1);
+  int task_size = data.paramInt("task_size", 102);
+  int img_w     = data.paramInt("img_width", 640);
+  int img_h     = data.paramInt("img_height", 512);
+  float soft    = (float) data.param("soft_margin", 12.0);
+  float hard    = (float) data.param("hard_margin",  8.0);
+  int num_entries = (int)(data.elements("ftasks0") / task_size);
+  size_t slice = (size_t)tilesy*tilesx*colors*256;
+  printf("case: %d sensors %dx%dx%d frame %dx%d soft/hard %g/%g entries %d\n",
+         num_sens, tilesy, tilesx, colors, img_w, img_h, soft, hard, num_entries);
+
+  void* p = tp_proc_create(m);
+  if(tp_proc_setup(p, num_sens, colors, img_w, img_h, 0, 0)){ printf("FAIL setup: %s\n", tp_last_error()); return 1; }
+  float* td_sens = data.hostFloat("td_sens");
+  for(int c=0;c<num_sens;c++)
+    if(tp_proc_set_clt(p, c, td_sens + (size_t)c*slice, 0)){ printf("FAIL set_clt %d\n", c); return 1; }
+  float* ftasks0 = data.hostFloat("ftasks0");
+  int stats[3] = {-1,-1,-1};
+  int rc = tp_proc_exec_consolidate(p, ftasks0, nullptr, num_entries, task_size, soft, hard, stats);
+  if(rc){ printf("FAIL exec_consolidate rc=%d: %s\n", rc, tp_last_error()); return 1; }
+  printf("exec_consolidate: %d pairs -> %d tiles, %d misaligned\n", stats[0], stats[1], stats[2]);
+
+  float* got = new float[slice];
+  if(tp_proc_get_clt(p, 0, 0, got)){ printf("FAIL get_clt(0)\n"); return 1; }
+  float* exp = data.hostFloat("expected_td_avg");
+  double maxd = 0; long nan_mm = 0; long first = -1;
+  for(size_t i=0;i<slice;i++){
+    bool gn = std::isnan(got[i]), en = std::isnan(exp[i]);
+    if(gn||en){ if(gn!=en){ nan_mm++; if(first<0) first=(long)i; } continue; }
+    double d = std::fabs((double)got[i]-(double)exp[i]);
+    if(d>maxd){ maxd=d; if(d>0 && first<0) first=(long)i; }
+  }
+  int fail = (maxd>0)||(nan_mm>0)||(stats[2]!=0);
+  printf("compare cam0 vs expected_td_avg: max|diff|=%g NaN-mismatch=%ld first@%ld -> %s\n",
+         maxd, nan_mm, first, fail?"FAIL":"PASS (bit-exact)");
+  tp_proc_destroy(p); tp_destroy_module(m);
+  return fail;
+}

jna/tp_jna.cpp

 // Stage-0b: the probe refactored into a JNA-callable shared lib (libtileproc.so).
-// extern "C": tp_create_module (NVRTC-compile kernels + cuLink(libcudadevrt) + load 19 funcs),
+// extern "C": tp_create_module (NVRTC-compile kernels + cuLink(libcudadevrt) + load 23 funcs), // By Claude on 07/12/2026: +4 tp_consolidate kernels
 // tp_module_num_functions, tp_last_error, tp_destroy_module. By Claude on 2026-06-25.
 #include <cuda.h>
 #include <nvrtc.h>
@@ -20,6 +20,7 @@
 #include "tp_paths.h"
 #include "tp_utils.h"   // copyalloc_kernel_gpu / copyalloc_pointers_gpu / copyalloc_image_gpu / alloc_kernel_gpu
 #include "tp_files.h"   // readFloatsFromFile / readAllFloatsFromFile
+#include "tp_consolidate.h" // TD_TILE_FLOATS (host use in tp_proc_exec_consolidate) // By Claude on 07/12/2026
 
 static char g_err[8192] = "";
 static void seterr(const char* fmt, ...){ va_list ap; va_start(ap,fmt); vsnprintf(g_err,sizeof(g_err),fmt,ap); va_end(ap); }
@@ -36,12 +37,14 @@ static const char* TP_DEFINES =
   "#define TASK_TEXT_S_BIT 4\n#define TASK_TEXT_SW_BIT 5\n#define TASK_TEXT_W_BIT 6\n#define TASK_TEXT_NW_BIT 7\n"
   "#define LIST_TEXTURE_BIT 8\n#define TEXT_NTILE_SHIFT 9\n#define FAT_ZERO_WEIGHT 0.0001\n"
   "#define THREADS_DYNAMIC_BITS 5\n#define RBYRDIST_LEN 5001\n#define RBYRDIST_STEP 0.0004\n#define TILES_PER_BLOCK_GEOM 2\n";
-static const char* SRC_FILES[] = {"dtt8x8.h","dtt8x8.cu","geometry_correction.h","geometry_correction.cu","TileProcessor.h","TileProcessor.cu"};
+static const char* SRC_FILES[] = {"dtt8x8.h","dtt8x8.cu","geometry_correction.h","geometry_correction.cu","TileProcessor.h","TileProcessor.cu",
+  "tp_consolidate.h","tp_consolidate.cu"}; // TD consolidation chain (roadmap 2c) // By Claude on 07/12/2026
 static const char* KERNELS[] = {
   "convert_direct","imclt_rbg_all","correlate2D","correlate2D_inter","corr2D_combine","corr2D_normalize",
   "textures_nonoverlap","generate_RBGA","clear_texture_list","mark_texture_tiles","mark_texture_neighbor_tiles",
   "gen_texture_list","clear_texture_rbga","textures_accumulate","create_nonoverlap_list","erase_clt_tiles",
-  "calculate_tiles_offsets","calc_rot_deriv","calcReverseDistortionTable"};
+  "calculate_tiles_offsets","calc_rot_deriv","calcReverseDistortionTable",
+  "clt_average_sensors","index_consolidate","consolidate_oob","clt_average_sensors_list"}; // By Claude on 07/12/2026
 static const int N_KERNELS = sizeof(KERNELS)/sizeof(KERNELS[0]);
 
 struct TpModule { CUcontext ctx; CUmodule mod; int nfun; };
@@ -599,6 +602,15 @@ struct TpProc {
   float *gpu_tex_textures, *gpu_tex_diff_rgb_combo, *gpu_tex_color_weights, *gpu_tex_rbga_params;
   int *gpu_tex_indices, *gpu_tex_pnum;
   int last_num_tex_tiles;
+  // TD consolidation chain (tp_consolidate.cu, roadmap 2c) — lazily allocated in
+  // tp_proc_exec_consolidate; cons_capacity = task entries the per-entry scratch is
+  // sized for (grown on demand). cons_td_sens = contiguous [num_cams][slice] gather
+  // of the per-cam clt_h slices (the kernels take one contiguous block, clt_h are 16
+  // separate allocations); cons_td_avg = full-grid average, D2D-copied into clt_h[0]
+  // (the slot the single conj-multiply correlates). By Claude on 07/12/2026.
+  int cons_capacity;
+  float *cons_ftasks0, *cons_ftasks1, *cons_pairs, *cons_tasks_out, *cons_td_sens, *cons_td_avg;
+  int   *cons_counters; // [0] pairs, [1] surviving tiles, [2] misaligned
 };
 
 extern "C" {
@@ -616,6 +628,8 @@ TpProc* tp_proc_create(TpModule* m){
   p->last_num_corr_combo=0; p->last_num_corr_tiles=0;
   p->have_tex=false; p->gpu_tex_textures=p->gpu_tex_diff_rgb_combo=p->gpu_tex_color_weights=p->gpu_tex_rbga_params=nullptr;
   p->gpu_tex_indices=p->gpu_tex_pnum=nullptr; p->last_num_tex_tiles=0;
+  p->cons_capacity=0; p->cons_ftasks0=p->cons_ftasks1=p->cons_pairs=p->cons_tasks_out=p->cons_td_sens=p->cons_td_avg=nullptr;
+  p->cons_counters=nullptr; // By Claude on 07/12/2026
   return p;
 }
 
@@ -735,6 +749,85 @@ int tp_proc_set_clt(TpProc* p, int cam, const float* data, int ref_scene){ if(!p
   float* dst = (ref_scene? p->clt_ref_h:p->clt_h)[cam];
   return cudaMemcpy(dst,data,(size_t)p->slice*sizeof(float),cudaMemcpyHostToDevice)==cudaSuccess?0:-2; }
 
+// On-device 16-sensor TD consolidation — the v1 chain of tp_consolidate.cu (roadmap 2c):
+// index_consolidate (MB-pair tile selection) -> consolidate_oob (soft/hard source-frame
+// margin priority, per-tile sensor mask folded into task bits TASK_AVG_SENS_SHIFT..) ->
+// clt_average_sensors_list (mask-gated NaN-aware average). Replaces the Java CPU bridge
+// getCltData -> CuasTD.consolidateSensorsTD -> setCltData(0): the per-sensor TD never
+// leaves the device; only the small task streams cross the JNA boundary. The average
+// lands in the cam-0 CLT slice (the slot correlate2D_inter reads with sensor mask 1);
+// tiles absent from the surviving list stay NaN (erase_clt poison semantics — includes
+// the everything-poisoned num_out==0 case, which still returns 0 with stats[1]==0).
+// ftasks1 = the MB pair task stream (index-aligned with ftasks0) or NULL = no-MB.
+// stats (may be NULL) = {num_pairs, surviving tiles, misaligned pairs}. Margins 0 = v0.
+// Host orchestration mirrors src/tests/test_avg_td_oob.cu (the NSight-debuggable test);
+// Java oracle: CuasTD.oobSensorMasks + masked consolidateSensorsTD. By Claude on 07/12/2026.
+int tp_proc_exec_consolidate(TpProc* p, const float* ftasks0, const float* ftasks1,
+                             int num_entries, int task_size,
+                             float soft_margin, float hard_margin, int* stats){
+  if(!p||!ftasks0||(num_entries<=0)||(task_size<=0)){ seterr("exec_consolidate: bad args"); return -1; }
+  cuCtxSetCurrent(p->mod->ctx);
+  CUfunction f_idx=getfun(p->mod,"index_consolidate"), f_oob=getfun(p->mod,"consolidate_oob"),
+             f_avg=getfun(p->mod,"clt_average_sensors_list");
+  if(!f_idx||!f_oob||!f_avg){ seterr("exec_consolidate: consolidate kernels missing"); return -2; }
+  // fixed-size scratch (first call): contiguous per-sensor gather + full-grid average + counters
+  if(!p->cons_td_sens){
+    if(cudaMalloc((void**)&p->cons_td_sens,(size_t)p->num_cams*p->slice*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->cons_td_avg,(size_t)p->slice*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->cons_counters,3*sizeof(int))!=cudaSuccess){
+      seterr("exec_consolidate: scratch alloc failed"); return -3; }
+  }
+  // per-entry scratch, grown on demand (gpu_pairs worst case: every entry survives)
+  if(num_entries>p->cons_capacity){
+    cudaFree(p->cons_ftasks0); cudaFree(p->cons_ftasks1); cudaFree(p->cons_pairs); cudaFree(p->cons_tasks_out);
+    size_t ts=(size_t)num_entries*task_size*sizeof(float);
+    if(cudaMalloc((void**)&p->cons_ftasks0,ts)!=cudaSuccess ||
+       cudaMalloc((void**)&p->cons_ftasks1,ts)!=cudaSuccess ||
+       cudaMalloc((void**)&p->cons_pairs,2*ts)!=cudaSuccess ||
+       cudaMalloc((void**)&p->cons_tasks_out,ts)!=cudaSuccess){
+      p->cons_capacity=0; seterr("exec_consolidate: task scratch alloc failed (%d entries)",num_entries); return -3; }
+    p->cons_capacity=num_entries;
+  }
+  size_t tbytes=(size_t)num_entries*task_size*sizeof(float);
+  if(cudaMemcpy(p->cons_ftasks0,ftasks0,tbytes,cudaMemcpyHostToDevice)!=cudaSuccess){ seterr("exec_consolidate: HtoD ftasks0"); return -4; }
+  float* d_ftasks1=nullptr;
+  if(ftasks1){
+    if(cudaMemcpy(p->cons_ftasks1,ftasks1,tbytes,cudaMemcpyHostToDevice)!=cudaSuccess){ seterr("exec_consolidate: HtoD ftasks1"); return -4; }
+    d_ftasks1=p->cons_ftasks1;
+  }
+  // gather the per-cam CLT slices into the contiguous [num_cams][slice] block the kernels
+  // index (clt_h are separate allocations); DtoD — no host round-trip
+  for(int c=0;c<p->num_cams;c++){
+    if(cudaMemcpy(p->cons_td_sens+(size_t)c*p->slice,p->clt_h[c],(size_t)p->slice*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess){
+      seterr("exec_consolidate: DtoD gather cam %d",c); return -4; } }
+  // prefill contract (test_avg_td_oob.cu): counters 0; td_avg all-NaN (0xFF bytes) = poison for absent tiles
+  cudaMemset(p->cons_counters,0,3*sizeof(int));
+  cudaMemset(p->cons_td_avg,0xFF,(size_t)p->slice*sizeof(float));
+  const int THREADS=256;
+  int* pnum_pairs=p->cons_counters+0; int* pnum_out=p->cons_counters+1; int* pmis=p->cons_counters+2;
+  void* a1[]={ &task_size,&p->cons_ftasks0,&d_ftasks1,&num_entries,&p->cons_pairs,&pnum_pairs,&pmis };
+  if(launch1(f_idx,(num_entries+THREADS-1)/THREADS,1,1,THREADS,1,1,a1,"index_consolidate")) return -5;
+  int num_pairs=0; cudaMemcpy(&num_pairs,p->cons_counters,sizeof(int),cudaMemcpyDeviceToHost);
+  int num_out=0;
+  if(num_pairs>0){
+    float w=(float)p->img_w, h=(float)p->img_h;
+    void* a2[]={ &task_size,&p->num_cams,&p->cons_pairs,&num_pairs,&w,&h,&soft_margin,&hard_margin,&p->cons_tasks_out,&pnum_out };
+    if(launch1(f_oob,(num_pairs+THREADS-1)/THREADS,1,1,THREADS,1,1,a2,"consolidate_oob")) return -5;
+    cudaMemcpy(&num_out,p->cons_counters+1,sizeof(int),cudaMemcpyDeviceToHost);
+  }
+  if(num_out>0){
+    float* counts=nullptr; // per-tile contributor counts not needed in the production bridge
+    void* a3[]={ &p->num_cams,&p->cons_td_sens,&p->tilesx,&p->tilesy,&p->num_colors,&p->cons_tasks_out,&task_size,&num_out,&p->cons_td_avg,&counts };
+    if(launch1(f_avg,num_out,p->num_colors,1,TD_TILE_FLOATS,1,1,a3,"clt_average_sensors_list")) return -5;
+  }
+  // the average -> the cam-0 slot (exactly what the CPU bridge's setCltData(0,...) did)
+  if(cudaMemcpy(p->clt_h[0],p->cons_td_avg,(size_t)p->slice*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess){
+    seterr("exec_consolidate: DtoD avg -> cam0"); return -4; }
+  int mis=0; cudaMemcpy(&mis,p->cons_counters+2,sizeof(int),cudaMemcpyDeviceToHost);
+  if(stats){ stats[0]=num_pairs; stats[1]=num_out; stats[2]=mis; }
+  return 0;
+}
+
 // Allocate the imclt (RBG) + correlation buffers and store the corr config params.
 int tp_proc_setup_rbg_corr(TpProc* p, int num_pairs, int s0,int s1,int s2,int s3,
                            float cw0,float cw1,float cw2, int corr_out_rad){
@@ -1066,6 +1159,8 @@ void tp_proc_destroy(TpProc* p){
     cudaFree(p->gpu_corr_indices); cudaFree(p->gpu_corrs_combo_indices); cudaFree(p->gpu_num_corr_tiles); }
   if(p->have_tex){ cudaFree(p->gpu_tex_textures); cudaFree(p->gpu_tex_diff_rgb_combo); cudaFree(p->gpu_tex_indices);
     cudaFree(p->gpu_tex_pnum); cudaFree(p->gpu_tex_color_weights); cudaFree(p->gpu_tex_rbga_params); }
+  cudaFree(p->cons_ftasks0); cudaFree(p->cons_ftasks1); cudaFree(p->cons_pairs); cudaFree(p->cons_tasks_out); // By Claude on 07/12/2026
+  cudaFree(p->cons_td_sens); cudaFree(p->cons_td_avg); cudaFree(p->cons_counters); // cudaFree(nullptr) is a no-op // By Claude on 07/12/2026
   delete p;
 }
 

src/tp_consolidate.cu

@@ -13,12 +13,20 @@
  *      Author: Claude (Anthropic), for Elphel
  */
 
+#ifndef JCUDA // NVRTC concatenation (tp_jna.cpp TP_DEFINES, headers merged in order) has no include paths // By Claude on 07/12/2026
 #include <cuda_runtime.h>              // __global__/__shared__ also for non-nvcc parsers // By Claude on 07/12/2026
 #include <device_launch_parameters.h>  // threadIdx/blockIdx: cuda_runtime.h gates this on __CUDACC__ // By Claude on 07/12/2026
 #include <math.h>                      // NAN, isnan // By Claude on 07/12/2026
 #include "tp_defines.h"                // NUM_CAMS (for geometry_correction.h) // By Claude on 07/12/2026
 #include "geometry_correction.h"       // struct tp_task field offsets TP_TASK_*_OFFSET // By Claude on 07/12/2026
 #include "tp_consolidate.h"
+#endif // By Claude on 07/12/2026
+#ifndef NAN // NVRTC has isnan built in but no math.h NAN/INFINITY macros (device-only use is fine) // By Claude on 07/12/2026
+#define NAN __int_as_float(0x7fc00000)
+#endif // By Claude on 07/12/2026
+#ifndef INFINITY // By Claude on 07/12/2026
+#define INFINITY __int_as_float(0x7f800000)
+#endif // By Claude on 07/12/2026
 #ifdef __CDT_PARSER__                  // Eclipse CDT indexer only, nvcc never defines this // By Claude on 07/12/2026
 extern "C" void __syncthreads();       // real decl is __CUDACC__-gated in crt/device_functions.h // By Claude on 07/12/2026
 #endif

src/tp_consolidate.h

@@ -35,7 +35,9 @@
 #ifndef SRC_TP_CONSOLIDATE_H_
 #define SRC_TP_CONSOLIDATE_H_
 
+#ifndef JCUDA // NVRTC concatenation (tp_jna.cpp TP_DEFINES) has no include paths // By Claude on 07/12/2026
 #include <cuda_runtime.h> // __global__ in the prototypes below, also for non-nvcc parsers // By Claude on 07/12/2026
+#endif // By Claude on 07/12/2026
 
 // TD tile = 4 quadrants x 8x8 = 256 floats (see dtt8x8 / TileProcessor docs)
 #define TD_TILE_FLOATS 256