CLAUDE: test_pose_corr_jna - the tol-0 verdict binary (NVRTC production...

CLAUDE: test_pose_corr_jna - the tol-0 verdict binary (NVRTC production module); direct test = stepping companion

RUNG 1 VALIDATED on the first real v013 case (scene 1773135467_917220,
4 LMA iterations, 150 tiles/set, 300 corr tiles/iter): test_pose_corr_jna
PASS BIT-EXACT @tol 0 - tasks (calculate_tiles_offsets) and raw TD
correlation (correlate2D_inter) both max|diff|=0, every iteration.

The offline-nvcc direct-launch binary (src/tests/test_pose_corr) cannot be
bit-exact vs the NVRTC JIT (FMA/codegen divergence), measured: -G -O0 ->
tasks off by 2^-13 px; RELEASE=1 -> tasks BIT-EXACT, corr within ~1 ULP at
operand scale (max|diff| 128..256 vs amplitudes ~2e9, float spacing 256).
Role split documented in both headers (same pairing as test_avg_td_oob /
test_proc_consolidate): direct binary for cuda-gdb stepping, _jna binary
for the production-equivalence verdict. Build line in the _jna header.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

jna/test_pose_corr_jna.cu

+// test_pose_corr_jna.cu - the TOL-0 verdict binary for the lean pose-chain
+// CHECKPOINT 1 (roadmap 3-C rung 1): replay the per-LMA-iteration chain of
+// CuasPoseRT.leanMeasure through the SAME NVRTC-compiled module and the SAME
+// granular TpProc functions the production JNA backend (GpuQuadJna) calls:
+//   set_tasks -> exec_geometry(0) -> [get_tasks vs oracle, bit-exact] ->
+//   exec_convert_direct (erase NaN on the first MB set only) ->
+//   exec_consolidate (oracle post-offsets streams = production bytes) ->
+//   exec_corr2d_inter_td vs the resident center TD (set_clt ref cam 0) ->
+//   get_corr_indices/get_corr_td vs oracle, ORDER-INDEPENDENT by packed index.
+//
+// Why this exists next to src/tests/test_pose_corr.cu (direct launches, -G,
+// cuda-gdb steppable): an offline nvcc build cannot be bit-exact against the
+// NVRTC-JIT production module (FMA/codegen divergence; measured: tasks match
+// bit-exactly only in the RELEASE build, corr TD stays within ~1 ULP at
+// operand scale). Identical to the test_avg_td_oob / test_proc_consolidate
+// pairing: the direct binary is for stepping into kernels, THIS binary is the
+// bit-exact production-equivalence check.
+//
+// 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_pose_corr_jna.cu src/tests/tp_test_data.cu \
+//     -o tests_bin/test_pose_corr_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_pose_corr_jna --data <case> --tol 0
+//
+//  Created on: Jul 13, 2026
+//      Author: Claude (Anthropic), for Elphel
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cmath>
+#include <map>
+#include <string>
+#include <vector>
+
+#include <cuda_runtime.h>
+#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_geometry(void*, const float*, int);
+  int   tp_proc_set_correction_vector(void*, const float*, int);
+  int   tp_proc_set_kernels(void*, int, const float*, int);
+  int   tp_proc_set_kernel_offsets(void*, int, const float*, int);
+  int   tp_proc_set_image(void*, int, const float*);
+  int   tp_proc_set_tasks(void*, const float*, int, int);
+  int   tp_proc_get_tasks(void*, float*, int);
+  int   tp_proc_set_const(void*, const char*, const float*, int);
+  int   tp_proc_exec_geometry(void*, int);
+  int   tp_proc_exec_convert_direct(void*, int, int, int);
+  int   tp_proc_set_clt(void*, int, const float*, int);
+  int   tp_proc_exec_consolidate(void*, const float*, const float*, int, int, float, float, int*);
+  int   tp_proc_setup_rbg_corr(void*, int, int,int,int,int, float,float,float, int);
+  int   tp_proc_exec_corr2d_inter_td(void*, int, float,float,float);
+  int   tp_proc_get_corr_indices(void*, int*, int);
+  int   tp_proc_get_corr_td(void*, float*);
+  int   tp_proc_num_corr_tiles(void*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+// bit-exact (NaN==NaN) compare of the task fields the GPU/oracle fill: 6 header
+// floats + 2*num_cams xy; the disp_dist tail is never round-tripped by the Java
+// flatten (TpTask.asFloatArray) and is excluded. Same policy as test_pose_corr.cu.
+static int compareTasks(const float* got, const float* expected,
+                        int n_entries, int task_size, int num_cams, const char* label){
+  const int cmp_floats = 6 + 2*num_cams;
+  int bad = 0; long first_bad = -1; double maxd = 0;
+  for (int i = 0; i < n_entries; i++)
+    for (int j = 0; j < cmp_floats; j++){
+      const size_t idx = (size_t)i*task_size + j;
+      const float g = got[idx], e = expected[idx];
+      const bool gn = std::isnan(g), en = std::isnan(e);
+      if (gn && en) continue;
+      if (gn != en){ bad++; if(first_bad<0) first_bad=(long)idx; continue; }
+      const double d = std::fabs((double)g-(double)e);
+      if (d > 0){ bad++; if(d>maxd) maxd=d; if(first_bad<0) first_bad=(long)idx; }
+    }
+  if (bad) printf("compare '%s': %d of %dx%d floats mismatch, max|diff|=%g, first @%ld -> FAIL\n",
+                  label, bad, n_entries, cmp_floats, maxd, first_bad);
+  else     printf("compare '%s': %d entries, bit-exact -> PASS\n", label, n_entries);
+  return bad;
+}
+
+int main(int argc, char** argv){
+  std::string data_dir = "testdata/pose_corr";
+  std::string srcdir   = "/home/elphel/git/tile_processor_gpu/src";
+  std::string devrt    = "/usr/local/cuda-12.8/lib64/libcudadevrt.a";
+  float tol = 0.0f;
+  int list_only = 0;
+  for (int i = 1; i < argc; i++){
+    if      (!strcmp(argv[i],"--data")  && (i+1<argc)) data_dir = argv[++i];
+    else if (!strcmp(argv[i],"--tol")   && (i+1<argc)) tol = (float)atof(argv[++i]);
+    else if (!strcmp(argv[i],"--src")   && (i+1<argc)) srcdir = argv[++i];
+    else if (!strcmp(argv[i],"--devrt") && (i+1<argc)) devrt = argv[++i];
+    else if (!strcmp(argv[i],"--list"))                list_only = 1;
+    else { fprintf(stderr,"Usage: %s --data <dir> [--tol X] [--src <kernel src>] [--devrt <libcudadevrt.a>] [--list]\n", argv[0]); return EXIT_FAILURE; }
+  }
+  printf("%s: data=%s tol=%g (NVRTC module from %s)\n", argv[0], data_dir.c_str(), tol, srcdir.c_str());
+  TpTestData data(data_dir);
+  if (list_only){ data.list(); return EXIT_SUCCESS; }
+
+  const int num_cams     = data.paramInt("num_sensors",  16);
+  const int num_colors   = data.paramInt("colors",        1);
+  const int tilesx       = data.paramInt("tilesx",       80);
+  const int tilesy       = data.paramInt("tilesy",       64);
+  const int img_width    = data.paramInt("img_width",   640);
+  const int img_height   = data.paramInt("img_height",  512);
+  const int kernels_hor  = data.paramInt("kernels_hor",  82);
+  const int kernels_vert = data.paramInt("kernels_vert", 66);
+  const int task_size    = data.paramInt("task_size", 6 + 6*num_cams);
+  const int num_iters    = data.paramInt("num_iters",     1);
+  const int has_mb       = data.paramInt("has_mb",        1);
+  const int sel_sensors  = data.paramInt("sel_sensors",   1);
+  const float soft_margin = (float)data.param("soft_margin", 12.0);
+  const float hard_margin = (float)data.param("hard_margin",  8.0);
+  const float s0 = (float)data.param("scale0",1.0), s1 = (float)data.param("scale1",0.0), s2 = (float)data.param("scale2",0.0);
+  const int chunk = 256;
+  const size_t slice = (size_t)tilesy*tilesx*num_colors*chunk;
+  const int kern_tiles = kernels_hor*kernels_vert*num_colors;
+  printf("cams=%d colors=%d frame %dx%d tiles %dx%d kernels %dx%d iters=%d mb=%d sel=0x%x soft/hard=%g/%g\n",
+         num_cams,num_colors,img_width,img_height,tilesx,tilesy,kernels_hor,kernels_vert,num_iters,has_mb,sel_sensors,soft_margin,hard_margin);
+
+  void* m = tp_create_module(srcdir.c_str(), devrt.c_str());
+  if(!m){ printf("FAIL module: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  void* p = tp_proc_create(m);
+  if(tp_proc_setup(p, num_cams, num_colors, img_width, img_height, kernels_hor, kern_tiles)){
+    printf("FAIL setup: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  // statics: geometry, kernels, images, LPF constants, resident center TD (ref cam 0)
+  tp_proc_set_geometry(p, data.hostFloat("geometry_correction"), (int)data.elements("geometry_correction"));
+  tp_proc_set_correction_vector(p, data.hostFloat("correction_vector"), (int)data.elements("correction_vector"));
+  const float* h_kern = data.hostFloat("kernels");
+  const float* h_offs = data.hostFloat("kernel_offsets");
+  const float* h_imgs = data.hostFloat("images");
+  const int kern_elems = (int)(data.elements("kernels")/num_cams);
+  const int offs_elems = (int)(data.elements("kernel_offsets")/num_cams);
+  for (int c = 0; c < num_cams; c++){
+    tp_proc_set_kernels       (p, c, h_kern + (size_t)c*kern_elems, kern_elems);
+    tp_proc_set_kernel_offsets(p, c, h_offs + (size_t)c*offs_elems, offs_elems);
+    tp_proc_set_image         (p, c, h_imgs + (size_t)c*img_width*img_height);
+  }
+  tp_proc_set_const(p, "lpf_data",    data.hostFloat("lpf_data"),    4*64);
+  tp_proc_set_const(p, "lpf_corr",    data.hostFloat("lpf_corr"),    64);
+  tp_proc_set_const(p, "lpf_rb_corr", data.hostFloat("lpf_rb_corr"), 64);
+  tp_proc_set_const(p, "LoG_corr",    data.hostFloat("log_corr"),    64);
+  tp_proc_set_clt(p, 0, data.hostFloat("clt_ref_cam0"), 1); // resident center TD, ref buffer cam 0
+  // correlation buffers: same sizing call ensureRbgCorr makes (num_pairs = C(16,2)=120,
+  // sel_pairs placeholders, corr_out_rad 7)
+  const int num_pairs = num_cams*(num_cams-1)/2;
+  tp_proc_setup_rbg_corr(p, num_pairs, 0,0,0,0, 1.f,1.f,1.f, 7);
+
+  int fail = 0;
+  char nm[64];
+  for (int k = 0; k < num_iters; k++){
+    int n_set[2] = {0,0};
+    const int num_sets = has_mb ? 2 : 1;
+    for (int nset = 0; nset < num_sets; nset++){
+      snprintf(nm,sizeof(nm),"ftasks%d_it%d",nset,k);
+      const int n = (int)(data.elements(nm)/task_size);
+      n_set[nset] = n;
+      tp_proc_set_tasks(p, data.hostFloat(nm), n, n*task_size);
+      if(tp_proc_exec_geometry(p, 0)){ printf("FAIL exec_geometry: %s\n", tp_last_error()); return EXIT_FAILURE; }
+      { // checkpoint: GPU-filled tasks vs the oracle post-offsets stream (bit-exact)
+        std::vector<float> got((size_t)n*task_size);
+        tp_proc_get_tasks(p, got.data(), n*task_size);
+        char lbl[64]; snprintf(lbl,sizeof(lbl),"tasks%d_it%d",nset,k);
+        snprintf(nm,sizeof(nm),"expected_tasks%d_it%d",nset,k);
+        if (compareTasks(got.data(), data.hostFloat(nm), n, task_size, num_cams, lbl)) fail = 1;
+      }
+      // erase NaN before the FIRST (SET) convert only, as interCorrTDMotionBlur
+      if(tp_proc_exec_convert_direct(p, 0, (nset==0)?1:-1, 0)){
+        printf("FAIL exec_convert_direct: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    }
+    // consolidation: the oracle post-offsets streams = the exact bytes production flattens
+    snprintf(nm,sizeof(nm),"expected_tasks0_it%d",k);
+    const float* cons0 = data.hostFloat(nm);
+    const float* cons1 = nullptr;
+    if (has_mb){ snprintf(nm,sizeof(nm),"expected_tasks1_it%d",k); cons1 = data.hostFloat(nm); }
+    int stats[3] = {-1,-1,-1};
+    if(tp_proc_exec_consolidate(p, cons0, cons1, n_set[0], task_size, soft_margin, hard_margin, stats)){
+      printf("FAIL exec_consolidate: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    // the single conj-multiply vs the resident center TD; device tasks = the LAST
+    // uploaded set (set 1 in MB mode), exactly the state execCorr2D_inter_TD sees
+    const int num_corr = tp_proc_exec_corr2d_inter_td(p, sel_sensors, s0, s1, s2);
+    printf("it%d: sets %d+%d -> %d pairs -> %d tiles (%d misaligned) -> %d corr tiles\n",
+           k, n_set[0], n_set[1], stats[0], stats[1], stats[2], num_corr);
+    if (stats[2]) fail = 1;
+    // checkpoint 1 compare: raw TD correlation, ORDER-INDEPENDENT by packed index
+    snprintf(nm,sizeof(nm),"expected_corr_indices_it%d",k);
+    const int* exp_idx = data.hostInt(nm);
+    const int  n_exp   = (int)data.elements(nm);
+    snprintf(nm,sizeof(nm),"expected_corr_td_it%d",k);
+    const float* exp_td = data.hostFloat(nm);
+    std::vector<int> got_idx(num_corr);
+    tp_proc_get_corr_indices(p, got_idx.data(), num_corr);
+    std::vector<float> got_td((size_t)num_corr*chunk);
+    tp_proc_get_corr_td(p, got_td.data());
+    if (num_corr != n_exp){ printf("it%d: num_corr %d != expected %d -> FAIL\n", k, num_corr, n_exp); fail = 1; }
+    std::map<int,int> by_index;
+    for (int i = 0; i < num_corr; i++) by_index[got_idx[i]] = i;
+    std::vector<float> reordered((size_t)n_exp*chunk, NAN);
+    double maxd = 0; long nan_mm = 0, missing = 0, first_bad = -1;
+    for (int i = 0; i < n_exp; i++){
+      std::map<int,int>::const_iterator it = by_index.find(exp_idx[i]);
+      if (it == by_index.end()){ missing++; continue; }
+      const float* g = got_td.data() + (size_t)it->second*chunk;
+      const float* e = exp_td + (size_t)i*chunk;
+      memcpy(reordered.data() + (size_t)i*chunk, g, chunk*sizeof(float));
+      for (int j = 0; j < chunk; j++){
+        const bool gn = std::isnan(g[j]), en = std::isnan(e[j]);
+        if (gn || en){ if(gn!=en){ nan_mm++; if(first_bad<0) first_bad=(size_t)i*chunk+j; } continue; }
+        const double d = std::fabs((double)g[j]-(double)e[j]);
+        if (d > maxd){ maxd = d; if((d>tol)&&(first_bad<0)) first_bad=(size_t)i*chunk+j; }
+      }
+    }
+    const int it_fail = (maxd>tol)||(nan_mm>0)||(missing>0)||(num_corr!=n_exp);
+    printf("it%d: compare 'corr_td' (%d tiles, by packed index): max|diff|=%g NaN-mismatch=%ld missing=%ld first@%ld -> %s\n",
+           k, n_exp, maxd, nan_mm, missing, first_bad, it_fail?"FAIL":"PASS");
+    if (it_fail) fail = 1;
+    // save the reordered result (oracle index order) for the Java round trip
+    float* scratch = nullptr;
+    cudaMalloc((void**)&scratch, reordered.size()*sizeof(float));
+    cudaMemcpy(scratch, reordered.data(), reordered.size()*sizeof(float), cudaMemcpyHostToDevice);
+    snprintf(nm,sizeof(nm),"corr_td_it%d",k);
+    data.saveResult(nm, scratch, {n_exp, chunk});
+    cudaFree(scratch);
+  }
+  printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+  tp_proc_destroy(p); tp_destroy_module(m);
+  return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

src/tests/test_pose_corr.cu

@@ -33,6 +33,18 @@
  * in calculate_tiles_offsets / correlate2D_inter_inner and step in
  * NSight/cuda-gdb.
  *
+ * ROLE (measured 07/13/2026 on the first real case): this DIRECT-LAUNCH binary
+ * is the DEBUG/STEPPING companion; the TOL-0 verdict binary is
+ * jna/test_pose_corr_jna (same replay through the NVRTC-compiled production
+ * module - PASSed bit-exact). An offline nvcc build cannot be bit-exact vs the
+ * NVRTC JIT (FMA/codegen divergence): default -G -O0 build -> tasks off by
+ * ~2^-13 px (ULP of px-scale values) and corr off ~ULP-cascade; RELEASE=1
+ * build -> tasks BIT-EXACT, corr within ~1 ULP at operand scale (max|diff|
+ * 128..256 vs amplitudes ~2e9, spacing 256). Same two-tier situation as
+ * test_convert_direct's goldens. So: expect PASS here only with a nonzero
+ * --tol (~1 ULP at scale) or against goldens blessed from THIS binary;
+ * use test_pose_corr_jna for the production-equivalence check.
+ *
  * Usage: test_pose_corr --data <dir> [--tol X] [--list]
  *
  *  Created on: Jul 13, 2026