CLAUDE: 3-C rung 2 - corr2D_peak_eig (frozen 15-arg getMaxXYCmEig contract) +...

CLAUDE: 3-C rung 2 - corr2D_peak_eig (frozen 15-arg getMaxXYCmEig contract) + checkpoint-2 replay + synthetic numpy gate

Andrey's ruling 07/13/2026: the GPU peak contract = the 15-arg superset -
closed-form 2x2 eigen, recentering as a runtime param, envelope de-bias as
an optional array pointer (NULL = off).

- src/tp_peak.{h,cu}: corr2D_peak_eig, one block/tile - parallel load/NaN/
  argmax(first-occurrence)/8-connected grow, thread-0 serial DOUBLE centroid/
  covariance/eigen/recentering in the Java oracle's exact summation order.
- jna/tp_jna.cpp: kernel in the NVRTC module (24 kernels) +
  tp_proc_exec_corr2d_peak / tp_proc_get_peaks / tp_proc_set_peak_debias.
- test_pose_corr + test_pose_corr_jna: checkpoint-2 replay (in-place
  corr2D_normalize -> corr_pd tol tier, peaks vs the double Java oracle at
  --peak-tol) when the case carries it; pre-rung-2 cases skip with a note.
- make_peak_case.py + test_peak: synthetic gate vs an INDEPENDENT numpy
  float64 port (8 adversarial tiles + refine/debias case) - PASS at
  max|diff| 0 / 9.6e-17.
- Regressions: test_pose_corr_jna bit-exact @tol 0 on the existing v013 case,
  avg_td / avg_td_oob / convert_direct PASS.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

jna/test_pose_corr_jna.cu

@@ -8,6 +8,16 @@
 //   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.
 //
+// CHECKPOINT 2 (rung 2, 07/13/2026), when the case carries expected_pd_indices:
+//   exec_corr2d_normalize(combo=0) IN PLACE on the inter buffer (production
+//   path) -> get_corr2d vs expected_corr_pd AT TOL 0 (same NVRTC module ==
+//   the production capture) -> set_peak_debias (if the case has one) +
+//   exec_corr2d_peak (frozen 15-arg getMaxXYCmEig contract, manifest params)
+//   -> get_peaks vs the double-precision Java oracle at --peak-tol (default
+//   1e-6; double kernel math vs Java double - FMA/libm ULPs, not a defect
+//   tier). Once production runs the kernel and re-exports GPU-produced peaks,
+//   THAT compare becomes the tol-0 tier. Old cases skip with a note.
+//
 // 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
@@ -36,6 +46,7 @@
 
 #include <cuda_runtime.h>
 #include "tests/tp_test_data.h"
+#include "tp_peak.h"        // PEAK_ROW_FLOATS (checkpoint 2) // By Claude on 07/13/2026
 
 extern "C" {
   void* tp_create_module(const char*, const char*);
@@ -59,6 +70,12 @@ extern "C" {
   int   tp_proc_get_corr_indices(void*, int*, int);
   int   tp_proc_get_corr_td(void*, float*);
   int   tp_proc_num_corr_tiles(void*);
+  // checkpoint 2 (rung 2) // By Claude on 07/13/2026
+  int   tp_proc_exec_corr2d_normalize(void*, int, double, int);
+  int   tp_proc_get_corr2d(void*, float*, int);
+  int   tp_proc_set_peak_debias(void*, const float*, int);
+  int   tp_proc_exec_corr2d_peak(void*, int, float,float,float,float, int, float,float,float);
+  int   tp_proc_get_peaks(void*, float*);
   void  tp_proc_destroy(void*);
   void  tp_destroy_module(void*);
 }
@@ -91,14 +108,16 @@ int main(int argc, char** argv){
   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;
+  float peak_tol = 1e-6f; // Java-double-oracle tier for the peak kernel (header note) // By Claude on 07/13/2026
   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],"--peak-tol") && (i+1<argc)) peak_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; }
+    else { fprintf(stderr,"Usage: %s --data <dir> [--tol X] [--peak-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);
@@ -153,6 +172,23 @@ int main(int argc, char** argv){
   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);
 
+  // checkpoint 2 (rung 2) statics: FZ-norm in place + frozen peak contract // By Claude on 07/13/2026
+  const int  corr_rad = 7, pd_size = 15*15;
+  const bool have_ckpt2 = data.has("expected_pd_indices_it0");
+  const double fz_inter = data.param("fz_inter", 10000.0);
+  const float p_abs_min   = (float)data.param("eig_min_abs",   0.0);
+  const float p_rel_min   = (float)data.param("eig_min_rel",   0.0);
+  const float p_min_peak  = (float)data.param("min_str_sum",   0.0);
+  const float p_sub_frac  = (float)data.param("eig_sub_frac",  0.0);
+  const int   p_recenter  =        data.paramInt("pose_recenter", 0);
+  const float p_sub_frac1 = (float)data.param("eig_sub_frac1", 0.0);
+  const float p_scale_ax  = (float)data.param("eig_scale_axes",1.2);
+  const float p_inc_ax    = (float)data.param("eig_inc_axes",  1.0);
+  if (have_ckpt2){
+    if (data.has("debias")) tp_proc_set_peak_debias(p, data.hostFloat("debias"), (int)data.elements("debias"));
+    printf("checkpoint 2: fz=%g recenter=%d debias=%s\n", fz_inter, p_recenter, data.has("debias")?"ON":"off");
+  } else printf("NOTE: checkpoint-2 buffers absent (pre-rung-2 case vintage) - FZ-norm/peak replay skipped\n");
+
   int fail = 0;
   char nm[64];
   for (int k = 0; k < num_iters; k++){
@@ -228,6 +264,72 @@ int main(int argc, char** argv){
     snprintf(nm,sizeof(nm),"corr_td_it%d",k);
     data.saveResult(nm, scratch, {n_exp, chunk});
     cudaFree(scratch);
+
+    // ---- CHECKPOINT 2 (rung 2): FZ-norm IN PLACE + peak kernel, production entry points ----
+    // By Claude on 07/13/2026
+    if (have_ckpt2){
+      snprintf(nm,sizeof(nm),"expected_pd_indices_it%d",k);
+      const int* pd_idx = data.hostInt(nm);
+      const int  n_pd   = (int)data.elements(nm);
+      snprintf(nm,sizeof(nm),"expected_corr_pd_it%d",k);
+      const float* exp_pd = data.hostFloat(nm);
+      snprintf(nm,sizeof(nm),"expected_peaks_it%d",k);
+      const float* exp_pk = data.hostFloat(nm);
+      if(tp_proc_exec_corr2d_normalize(p, 0, fz_inter, corr_rad)){
+        printf("FAIL exec_corr2d_normalize: %s\n", tp_last_error()); return EXIT_FAILURE; }
+      std::vector<float> got_pd((size_t)num_corr*pd_size);
+      tp_proc_get_corr2d(p, got_pd.data(), corr_rad);
+      const int n_peaks = tp_proc_exec_corr2d_peak(p, 0xff,
+          p_abs_min, p_rel_min, p_min_peak, p_sub_frac,
+          p_recenter, p_sub_frac1, p_scale_ax, p_inc_ax);
+      if (n_peaks < 0){ printf("FAIL exec_corr2d_peak: %s\n", tp_last_error()); return EXIT_FAILURE; }
+      std::vector<float> got_pk((size_t)n_peaks*PEAK_ROW_FLOATS);
+      tp_proc_get_peaks(p, got_pk.data());
+      // match by packed index (dense tile -> (tile<<8)|0xff sum slot), oracle row order out
+      std::vector<float> res_pd((size_t)n_pd*pd_size, NAN);
+      std::vector<float> res_pk((size_t)n_pd*7, NAN);
+      double pd_maxd = 0, pk_maxd = 0; long pd_nan = 0, pk_nan = 0, pd_missing = 0;
+      for (int i = 0; i < n_pd; i++){
+        const int packed = (pd_idx[i] << 8) | 0xff; // CORR_NTILE_SHIFT
+        std::map<int,int>::const_iterator it = by_index.find(packed);
+        if (it == by_index.end()){ pd_missing++; continue; }
+        const float* gpd = got_pd.data() + (size_t)it->second*pd_size;
+        const float* epd = exp_pd + (size_t)i*pd_size;
+        memcpy(res_pd.data() + (size_t)i*pd_size, gpd, pd_size*sizeof(float));
+        for (int j = 0; j < pd_size; j++){
+          const bool gn = std::isnan(gpd[j]), en = std::isnan(epd[j]);
+          if (gn || en){ if(gn!=en) pd_nan++; continue; }
+          const double d = std::fabs((double)gpd[j]-(double)epd[j]);
+          if (d > pd_maxd) pd_maxd = d;
+        }
+        const float* gpk = got_pk.data() + (size_t)it->second*PEAK_ROW_FLOATS;
+        const float* epk = exp_pk + (size_t)i*7;
+        memcpy(res_pk.data() + (size_t)i*7, gpk, 7*sizeof(float));
+        for (int j = 0; j < 7; j++){
+          const bool gn = std::isnan(gpk[j]), en = std::isnan(epk[j]);
+          if (gn || en){ if(gn!=en) pk_nan++; continue; }
+          const double d = std::fabs((double)gpk[j]-(double)epk[j]);
+          if (d > pk_maxd) pk_maxd = d;
+        }
+      }
+      const int pd_fail = (pd_maxd>tol)      || (pd_nan>0) || (pd_missing>0);
+      const int pk_fail = (pk_maxd>peak_tol) || (pk_nan>0) || (pd_missing>0);
+      printf("it%d: compare 'corr_pd' (%d tiles, in-place FZ-norm): max|diff|=%g NaN-mismatch=%ld missing=%ld -> %s\n",
+             k, n_pd, pd_maxd, pd_nan, pd_missing, pd_fail?"FAIL":"PASS");
+      printf("it%d: compare 'peaks' (%d tiles vs the double Java oracle): max|diff|=%g NaN-mismatch=%ld -> %s (peak_tol %g)\n",
+             k, n_pd, pk_maxd, pk_nan, pk_fail?"FAIL":"PASS", peak_tol);
+      if (pd_fail || pk_fail) fail = 1;
+      cudaMalloc((void**)&scratch, res_pd.size()*sizeof(float));
+      cudaMemcpy(scratch, res_pd.data(), res_pd.size()*sizeof(float), cudaMemcpyHostToDevice);
+      snprintf(nm,sizeof(nm),"corr_pd_it%d",k);
+      data.saveResult(nm, scratch, {n_pd, pd_size});
+      cudaFree(scratch);
+      cudaMalloc((void**)&scratch, res_pk.size()*sizeof(float));
+      cudaMemcpy(scratch, res_pk.data(), res_pk.size()*sizeof(float), cudaMemcpyHostToDevice);
+      snprintf(nm,sizeof(nm),"peaks_it%d",k);
+      data.saveResult(nm, scratch, {n_pd, 7});
+      cudaFree(scratch);
+    }
   }
   printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
   tp_proc_destroy(p); tp_destroy_module(m);

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 23 funcs), // By Claude on 07/12/2026: +4 tp_consolidate kernels
+// extern "C": tp_create_module (NVRTC-compile kernels + cuLink(libcudadevrt) + load 24 funcs), // By Claude on 07/13/2026: +corr2D_peak_eig (rung 2)
 // tp_module_num_functions, tp_last_error, tp_destroy_module. By Claude on 2026-06-25.
 #include <cuda.h>
 #include <nvrtc.h>
@@ -21,6 +21,7 @@
 #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
+#include "tp_peak.h"        // PEAK_ROW_FLOATS/PEAK_MAX_PIX (host use in tp_proc_exec_corr2d_peak) // By Claude on 07/13/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); }
@@ -38,13 +39,15 @@ static const char* TP_DEFINES =
   "#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",
-  "tp_consolidate.h","tp_consolidate.cu"}; // TD consolidation chain (roadmap 2c) // By Claude on 07/12/2026
+  "tp_consolidate.h","tp_consolidate.cu", // TD consolidation chain (roadmap 2c) // By Claude on 07/12/2026
+  "tp_peak.h","tp_peak.cu"}; // peak measurement kernel (roadmap 3-C rung 2) // By Claude on 07/13/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",
-  "clt_average_sensors","index_consolidate","consolidate_oob","clt_average_sensors_list"}; // By Claude on 07/12/2026
+  "clt_average_sensors","index_consolidate","consolidate_oob","clt_average_sensors_list", // By Claude on 07/12/2026
+  "corr2D_peak_eig"}; // By Claude on 07/13/2026
 static const int N_KERNELS = sizeof(KERNELS)/sizeof(KERNELS[0]);
 
 struct TpModule { CUcontext ctx; CUmodule mod; int nfun; };
@@ -611,6 +614,12 @@ struct TpProc {
   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
+  // peak measurement (tp_peak.cu, roadmap 3-C rung 2) — lazily allocated in
+  // tp_proc_exec_corr2d_peak; gpu_peak_debias uploaded once via
+  // tp_proc_set_peak_debias (len 0 = off -> kernel gets NULL). By Claude on 07/13/2026.
+  int    peak_capacity;   // corr tiles gpu_peaks is sized for (PEAK_ROW_FLOATS each)
+  float *gpu_peaks;
+  float *gpu_peak_debias; int peak_debias_len;
 };
 
 extern "C" {
@@ -630,6 +639,7 @@ TpProc* tp_proc_create(TpModule* m){
   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
+  p->peak_capacity=0; p->gpu_peaks=nullptr; p->gpu_peak_debias=nullptr; p->peak_debias_len=0; // By Claude on 07/13/2026
   return p;
 }
 
@@ -955,6 +965,56 @@ int tp_proc_exec_corr2d_normalize(TpProc* p, int combo, double fat_zero, int cor
     return launch1(f,1,1,1,1,1,1,a,"corr2D_normalize_pair")?-3:0;
   }
 }
+
+// --- peak measurement (tp_peak.cu corr2D_peak_eig, roadmap 3-C rung 2) --- // By Claude on 07/13/2026
+// Optional envelope de-bias array (CuasPoseRT.windowAutocorrDebias, (2*corr_rad+1)^2
+// floats): upload once per config; len 0 (or null data) clears it -> kernel gets NULL.
+int tp_proc_set_peak_debias(TpProc* p, const float* data, int len){
+  if(!p){ seterr("set_peak_debias: null"); return -1; } cuCtxSetCurrent(p->mod->ctx);
+  if(!data || (len<=0)){ cudaFree(p->gpu_peak_debias); p->gpu_peak_debias=nullptr; p->peak_debias_len=0; return 0; }
+  if(len>PEAK_MAX_PIX){ seterr("set_peak_debias: len %d > %d",len,PEAK_MAX_PIX); return -2; }
+  if(len!=p->peak_debias_len){ cudaFree(p->gpu_peak_debias); p->gpu_peak_debias=nullptr;
+    if(cudaMalloc((void**)&p->gpu_peak_debias,(size_t)len*sizeof(float))!=cudaSuccess){
+      p->peak_debias_len=0; seterr("set_peak_debias: alloc failed"); return -3; } }
+  if(cudaMemcpy(p->gpu_peak_debias,data,(size_t)len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("set_peak_debias: HtoD failed"); return -3; }
+  p->peak_debias_len=len;
+  return 0;
+}
+// Run the frozen 15-arg getMaxXYCmEig contract on the pixel-domain correlation buffer
+// (gpu_corrs, i.e. AFTER tp_proc_exec_corr2d_normalize combo=0) - one row of
+// PEAK_ROW_FLOATS {dx,dy,str,e0x,e0y,l0,l1,valid} per resident corr tile.
+// pair_select: process only slots with this pair code (lean sum = 0xff; -1 = all).
+// Returns the number of corr tiles (rows) or <0 on error.
+int tp_proc_exec_corr2d_peak(TpProc* p, int pair_select,
+                             float abs_min, float rel_min, float min_peak, float eig_sub_frac,
+                             int refine, float eig_sub_frac1, float scale_axes, float inc_axes){
+  if(!p||!p->have_corr){ seterr("exec_corr2d_peak: no corr buffers"); return -1; } cuCtxSetCurrent(p->mod->ctx);
+  CUfunction f=getfun(p->mod,"corr2D_peak_eig"); if(!f){ seterr("corr2D_peak_eig missing"); return -2; }
+  int nct=p->last_num_corr_tiles;
+  if(nct<=0){ seterr("exec_corr2d_peak: no corr tiles (%d)",nct); return -3; }
+  if(nct>p->peak_capacity){
+    cudaFree(p->gpu_peaks); p->gpu_peaks=nullptr;
+    if(cudaMalloc((void**)&p->gpu_peaks,(size_t)nct*PEAK_ROW_FLOATS*sizeof(float))!=cudaSuccess){
+      p->peak_capacity=0; seterr("exec_corr2d_peak: peaks alloc failed (%d tiles)",nct); return -4; }
+    p->peak_capacity=nct;
+  }
+  size_t sp=p->dstride_corr/sizeof(float); int crad=p->corr_out_rad;
+  void* a[]={ &nct,&p->gpu_corr_indices,&pair_select,&sp,&p->gpu_corrs,&crad,
+              &abs_min,&rel_min,&min_peak,&eig_sub_frac,
+              &refine,&eig_sub_frac1,&scale_axes,&inc_axes,
+              &p->gpu_peak_debias,&p->gpu_peaks }; // gpu_peak_debias nullptr = de-bias off
+  if(launch1(f,nct,1,1,256,1,1,a,"corr2D_peak_eig")) return -5; // 256 = PEAK_THREADS (tp_peak.cu)
+  return nct;
+}
+// D2H the peak rows (last_num_corr_tiles x PEAK_ROW_FLOATS floats). Returns the count.
+int tp_proc_get_peaks(TpProc* p, float* out){
+  if(!p||!p->gpu_peaks){ seterr("get_peaks: no peaks buffer"); return -1; } cuCtxSetCurrent(p->mod->ctx);
+  int nct=p->last_num_corr_tiles;
+  if(cudaMemcpy(out,p->gpu_peaks,(size_t)nct*PEAK_ROW_FLOATS*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("get_peaks: DtoH failed"); return -2; }
+  return nct;
+}
 // de-pitch per-pair gpu_corrs (last_num_corr_tiles x (2*corr_rad+1)^2). returns count.
 int tp_proc_get_corr2d(TpProc* p, float* out, int corr_rad){ if(!p||!p->have_corr)return -1; cuCtxSetCurrent(p->mod->ctx);
   int per=(2*corr_rad+1)*(2*corr_rad+1);
@@ -1161,6 +1221,7 @@ void tp_proc_destroy(TpProc* p){
     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
+  cudaFree(p->gpu_peaks); cudaFree(p->gpu_peak_debias); // By Claude on 07/13/2026
   delete p;
 }
 

src/tests/build_tests.sh

@@ -25,7 +25,7 @@ fi
 
 # Kernel + host sources shared by every test ("load all kernels"):
 COMMON="$SRC/TileProcessor.cu $SRC/dtt8x8.cu $SRC/geometry_correction.cu \
-        $SRC/tp_consolidate.cu $SRC/tp_utils.cu $SRC/tp_files.cu tp_test_data.cu"
+        $SRC/tp_consolidate.cu $SRC/tp_peak.cu $SRC/tp_utils.cu $SRC/tp_files.cu tp_test_data.cu"
 
 TESTS="${@:-$(ls test_*.cu | sed 's/\.cu$//')}"
 for t in $TESTS; do

src/tests/make_peak_case.py

+#!/usr/bin/env python3
+# make_peak_case.py - synthetic test case for corr2D_peak_eig (tp_peak.cu,
+# roadmap 3-C rung 2): NO Java/GPU needed. The expected values come from an
+# INDEPENDENT numpy float64 port of the frozen 15-arg getMaxXYCmEig contract
+# (Correlation2d.java, eig_fast2x2 branch) - same math, same summation order,
+# so the double-math kernel must agree to rounding (test_peak default tol 1e-9).
+#
+# Tiles exercise the contract edges:
+#   0: centered isotropic Gaussian            (baseline)
+#   1: offset anisotropic rotated Gaussian    (eigen axes + subpixel centroid)
+#   2: wide peak near the border              (region clipped by the tile edge)
+#   3: weak peak below min_peak               (reject -> NaN row)
+#   4: NaN poisoned tile                      (reject -> NaN row)
+#   5: main peak + DISCONNECTED second blob above threshold (8-connectivity:
+#      the far blob must NOT bias the centroid)
+#   6: two peaks TOUCHING diagonally          (8-connectivity: must merge)
+#   7: negative-pedestal peak                 (threshold min(abs_min, rel_min*mx))
+# Cases written:
+#   testdata/synth4/peak          - refine=0, no debias (the validated baseline shape)
+#   testdata/synth4/peak_recenter - refine=2 + debias   (both runtime hooks ON)
+#
+# Usage: python3 src/tests/make_peak_case.py
+# Then:  tests_bin/test_peak --data testdata/synth4/peak --tol 1e-9
+#        tests_bin/test_peak --data testdata/synth4/peak_recenter --tol 1e-9
+#
+# By Claude on 07/13/2026.
+import os
+import numpy as np
+
+CORR_RAD = 7
+W = 2 * CORR_RAD + 1          # 15
+CENTER = CORR_RAD
+SHIFT = 8                     # CORR_NTILE_SHIFT
+
+# ---- the frozen contract, independent float64 port (Correlation2d.getMaxXYCmEig 15-arg) ----
+
+def eigen2x2(cxx, cxy, cyy):
+    """Correlation2d.getEigen2x2 port: {e0x,e0y,l0,l1} or None."""
+    e = 1e-12
+    hapd = (cxx + cyy) / 2
+    hamd = (cxx - cyy) / 2
+    bc = cxy * cxy
+    d = np.sqrt(hamd * hamd + bc)
+    with np.errstate(divide='ignore', invalid='ignore'):
+        if d / abs(hapd) < e:
+            return (1.0, 0.0, hapd, hapd)
+    lam0, lam1 = hapd - d, hapd + d
+    v0x, v0y = cxx - lam1, cxy
+    l = np.sqrt(v0x * v0x + v0y * v0y)
+    if l == 0:
+        return None
+    k = 1.0 / l
+    if v0x < 0:
+        k = -k
+    return (k * v0x, k * v0y, lam0, lam1)
+
+
+def connected8(above, seed):
+    """TileNeibs.getConnected (dirs=8) as a BFS set."""
+    en = np.zeros_like(above, dtype=bool)
+    if not above.flat[seed]:
+        return en
+    front = [seed]
+    en.flat[seed] = True
+    while front:
+        i = front.pop(0)
+        iy, ix = divmod(i, W)
+        for dy in (-1, 0, 1):
+            for dx in (-1, 0, 1):
+                if dx == 0 and dy == 0:
+                    continue
+                y, x = iy + dy, ix + dx
+                if 0 <= y < W and 0 <= x < W:
+                    j = y * W + x
+                    if above.flat[j] and not en.flat[j]:
+                        en.flat[j] = True
+                        front.append(j)
+    return en
+
+
+def get_max_xy_cm_eig(data, abs_min, rel_min, min_peak, eig_sub_frac,
+                      refine, eig_sub_frac1, scale_axes, inc_axes):
+    """15-arg getMaxXYCmEig, eig_fast2x2 branch. data float64 [W*W]. None = oracle null."""
+    if np.any(np.isnan(data)):
+        return None
+    imax = int(np.argmax(data))            # first occurrence, as the Java forward scan
+    mx = data[imax]
+    if mx < min_peak:
+        return None
+    ix0, iy0 = imax % W, imax // W
+    min_d = min(abs_min, rel_min * mx)
+    sub_pedestal = min_d * eig_sub_frac
+    sub_pedestal1 = min_d * eig_sub_frac1
+    en = connected8(data.reshape(W, W) >= min_d, imax)
+    s0 = sx = sy = sx2 = sy2 = sxy = 0.0
+    for iy in range(W):                    # exact oracle summation order
+        y = float(iy - iy0)
+        for ix in range(W):
+            if en[iy, ix]:
+                x = float(ix - ix0)
+                d = data[iy * W + ix] - sub_pedestal
+                s0 += d
+                sx += d * x
+                sy += d * y
+                sx2 += d * x * x
+                sy2 += d * y * y
+                sxy += d * x * y
+    x0 = ix0 + sx / s0
+    y0 = iy0 + sy / s0
+    cxx = sx2 - sx * sx / s0
+    cyy = sy2 - sy * sy / s0
+    cxy = sxy - sx * sy / s0
+    if (sx2 == 0) or (sy2 == 0) or (cxx == 0) or (cyy == 0):
+        return None
+    eig = eigen2x2(cxx, cxy, cyy)
+    if eig is None:
+        return None
+    xy = [x0 - CENTER, y0 - CENTER]
+    if refine > 0:
+        ha0 = scale_axes * np.sqrt(eig[2]) + inc_axes
+        ha1 = scale_axes * np.sqrt(eig[3]) + inc_axes
+        t00, t01 = eig[0] / ha0, eig[1] / ha0
+        t10, t11 = eig[1] / ha1, -eig[0] / ha1
+        for _ in range(refine):
+            r0 = rx = ry = 0.0
+            for iy in range(W):
+                y = iy - xy[1] - CENTER
+                for ix in range(W):
+                    x = ix - xy[0] - CENTER
+                    de0 = t00 * x + t01 * y
+                    de1 = t10 * x + t11 * y
+                    er2 = de0 * de0 + de1 * de1
+                    if er2 < 1.0:
+                        d = data[iy * W + ix] - sub_pedestal1
+                        if d > 0:
+                            d *= np.cos(0.5 * np.pi * np.sqrt(er2))
+                            r0 += d
+                            rx += d * x
+                            ry += d * y
+            if r0 > 0:
+                xy[0] += rx / r0
+                xy[1] += ry / r0
+    return (xy[0], xy[1], mx, eig[0], eig[1], eig[2], eig[3])
+
+
+# ---- synthetic tiles ----
+
+def gauss(dx, dy, sx, sy, rot=0.0, amp=1.0):
+    yy, xx = np.mgrid[0:W, 0:W].astype(np.float64)
+    x = xx - (CENTER + dx)
+    y = yy - (CENTER + dy)
+    c, s = np.cos(rot), np.sin(rot)
+    u, v = c * x + s * y, -s * x + c * y
+    return amp * np.exp(-0.5 * ((u / sx) ** 2 + (v / sy) ** 2))
+
+
+def build_tiles():
+    tiles = []
+    tiles.append(gauss(0, 0, 1.5, 1.5))                          # 0 baseline
+    tiles.append(gauss(1.3, -0.7, 2.8, 1.2, rot=0.5))            # 1 offset anisotropic rotated
+    tiles.append(gauss(5.4, 4.9, 3.0, 2.5, rot=-0.3))            # 2 wide, near border (clipped region)
+    tiles.append(gauss(0.5, 0.5, 1.5, 1.5, amp=0.005))           # 3 weak -> reject
+    t = gauss(0, 0, 1.5, 1.5); t[3, 3] = np.nan; tiles.append(t) # 4 NaN -> reject
+    t = gauss(-1.0, 0.5, 1.4, 1.4) + gauss(6.0, -6.0, 1.0, 1.0, amp=0.6)
+    tiles.append(t)                                              # 5 disconnected far blob
+    t = gauss(-1.5, -1.5, 1.2, 1.2) + gauss(1.5, 1.5, 1.2, 1.2, amp=0.8)
+    tiles.append(t)                                              # 6 diagonal-touching pair
+    tiles.append(gauss(0.8, -1.9, 2.0, 1.6, rot=1.1) - 0.02)     # 7 negative pedestal
+    return tiles
+
+
+def window_autocorr_debias(frac):
+    """CuasPoseRT.windowAutocorrDebias port (MCLT half-sine window autocorr)."""
+    ts = CORR_RAD + 1                     # 8
+    wl = 2 * ts
+    wnd = np.sin(np.pi * (np.arange(wl) + 0.5) / wl)
+    ac = np.array([np.sum(wnd[max(0, -d):min(wl, wl - d)] * wnd[max(0, -d) + d:min(wl, wl - d) + d])
+                   for d in range(-(ts - 1), ts)])
+    ac = ac / ac[ts - 1]
+    env = np.outer(ac, ac)
+    return env ** (-frac)
+
+
+def write_case(dirname, tiles, params, debias):
+    os.makedirs(dirname, exist_ok=True)
+    n = len(tiles)
+    # packed indices: fake dense tile numbers 10,20,... with a non-sum slot
+    # interleaved before each sum slot (pair 0) - the kernel must skip those
+    idx = []
+    pd = []
+    for i, t in enumerate(tiles):
+        tile_no = 10 * (i + 1)
+        idx.append((tile_no << SHIFT) | 0x00)   # per-sensor slot - skipped by pair_select
+        pd.append(np.zeros(W * W))              # content irrelevant (skipped)
+        idx.append((tile_no << SHIFT) | 0xff)   # the sum slot the kernel measures
+        pd.append(t.reshape(-1))
+    pd = np.array(pd, dtype=np.float64)
+    # expected: on the DEBIASED tile when debias is present (kernel multiplies float
+    # debias in double - reference uses the same float-cast array)
+    dz = debias.astype(np.float32).astype(np.float64).reshape(-1) if debias is not None else None
+    peaks = []
+    pd_indices = []
+    for i in range(n):
+        data = pd[2 * i + 1].copy()
+        if dz is not None and not np.any(np.isnan(data)):
+            data = data * dz
+        r = get_max_xy_cm_eig(data, **params)
+        peaks.append([np.nan] * 7 if r is None else list(r))
+        pd_indices.append(10 * (i + 1))
+    peaks = np.array(peaks, dtype=np.float32)
+    manifest = ["# synthetic corr2D_peak_eig case (make_peak_case.py) - numpy float64 reference",
+                "prm num_corr_tiles %d" % (2 * n),
+                "prm corr_radius %d" % CORR_RAD,
+                "prm eig_min_abs %.17g" % params['abs_min'],
+                "prm eig_min_rel %.17g" % params['rel_min'],
+                "prm min_str_sum %.17g" % params['min_peak'],
+                "prm eig_sub_frac %.17g" % params['eig_sub_frac'],
+                "prm pose_recenter %d" % params['refine'],
+                "prm eig_sub_frac1 %.17g" % params['eig_sub_frac1'],
+                "prm eig_scale_axes %.17g" % params['scale_axes'],
+                "prm eig_inc_axes %.17g" % params['inc_axes']]
+
+    def buf(name, arr, dtype, dims):
+        fn = name + ('.i32' if dtype == 'i32' else '.f32')
+        arr.astype(np.int32 if dtype == 'i32' else np.float32).tofile(os.path.join(dirname, fn))
+        # manifest buf line: buf <name> <dtype> <ndims> <dim...> <file>
+        manifest.append("buf %s %s %d %s %s" % (name, dtype, len(dims), ' '.join(map(str, dims)), fn))
+
+    buf("corr_indices", np.array(idx), 'i32', [2 * n])
+    buf("corr_pd", pd, 'f32', [2 * n, W * W])
+    buf("expected_pd_indices", np.array(pd_indices), 'i32', [n])
+    buf("expected_peaks", peaks, 'f32', [n, 7])
+    if debias is not None:
+        buf("debias", debias.reshape(-1), 'f32', [W * W])
+    with open(os.path.join(dirname, "manifest.txt"), "w") as f:
+        f.write("\n".join(manifest) + "\n")
+    print("wrote %s (%d tiles, refine=%d, debias=%s)" %
+          (dirname, n, params['refine'], "ON" if debias is not None else "off"))
+    # NOTE: expected_peaks computed from float32-CAST pd tiles (the kernel input),
+    # promoted back to float64 - exactly what the double-math kernel sees.
+
+
+def main():
+    root = os.path.join(os.path.dirname(os.path.abspath(__file__)), "..", "..", "testdata", "synth4")
+    tiles32 = [t.astype(np.float32).astype(np.float64) for t in build_tiles()]  # kernel sees float32
+    base = dict(abs_min=0.03, rel_min=0.15, min_peak=0.02, eig_sub_frac=1.0,
+                refine=0, eig_sub_frac1=0.0, scale_axes=1.2, inc_axes=1.0)
+    write_case(os.path.join(root, "peak"), tiles32, base, None)
+    rec = dict(base, refine=2)
+    write_case(os.path.join(root, "peak_recenter"), tiles32, rec, window_autocorr_debias(0.85))
+
+
+if __name__ == "__main__":
+    main()

src/tests/test_peak.cu

+/*
+ * test_peak.cu
+ *
+ * Standalone test of corr2D_peak_eig (tp_peak.cu, roadmap 3-C rung 2) against
+ * the SYNTHETIC numpy reference case (src/tests/make_peak_case.py): float64
+ * port of the frozen 15-arg getMaxXYCmEig contract, same summation order, so
+ * the double-math kernel must agree to rounding. NaN rows (rejects) compare
+ * NaN==NaN; interleaved non-sum slots verify the pair_select skip.
+ *
+ * This is the KERNEL-LOGIC gate that runs with no Java/GPU-export dependency
+ * (flood-fill connectivity, argmax tie-break, eigen, recentering, de-bias);
+ * the REAL-DATA gates are test_pose_corr / test_pose_corr_jna checkpoint 2
+ * on a rung-2 pose_corr case export.
+ *
+ * Usage: test_peak --data testdata/synth4/peak [--tol 1e-9] [--list]
+ *
+ *  Created on: Jul 13, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cmath>
+#include <string>
+#include <vector>
+
+#include <cuda_runtime.h>
+#include <helper_cuda.h>
+
+#include "tp_defines.h"            // CORR_NTILE_SHIFT
+#include "tp_peak.h"
+#include "tests/tp_test_data.h"
+#include "tests/tp_gpu_buf.h"
+
+int main(int argc, char **argv)
+{
+	std::string data_dir = "testdata/synth4/peak";
+	float tol = 1e-9f;             // double-math kernel vs the float64 numpy reference
+	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], "--list"))                   list_only = 1;
+		else {
+			fprintf(stderr, "Usage: %s --data <dir> [--tol X] [--list]\n", argv[0]);
+			return EXIT_FAILURE;
+		}
+	}
+	printf("%s: data=%s tol=%g\n", argv[0], data_dir.c_str(), tol);
+	findCudaDevice(argc, (const char **) argv);
+	TpTestData data(data_dir);
+	if (list_only) { data.list(); return EXIT_SUCCESS; }
+
+	const int num_corr  = data.paramInt("num_corr_tiles", 0);
+	const int corr_rad  = data.paramInt("corr_radius", 7);
+	const int pd_size   = (2 * corr_rad + 1) * (2 * corr_rad + 1);
+	const float p_abs_min   = (float) data.param("eig_min_abs",   0.0);
+	const float p_rel_min   = (float) data.param("eig_min_rel",   0.0);
+	const float p_min_peak  = (float) data.param("min_str_sum",   0.0);
+	const float p_sub_frac  = (float) data.param("eig_sub_frac",  0.0);
+	const int   p_recenter  =         data.paramInt("pose_recenter", 0);
+	const float p_sub_frac1 = (float) data.param("eig_sub_frac1", 0.0);
+	const float p_scale_ax  = (float) data.param("eig_scale_axes",1.2);
+	const float p_inc_ax    = (float) data.param("eig_inc_axes",  1.0);
+	printf("tiles=%d rad=%d abs/rel/min=%g/%g/%g sub=%g/%g recenter=%d axes=%gx+%g debias=%s\n",
+			num_corr, corr_rad, p_abs_min, p_rel_min, p_min_peak, p_sub_frac, p_sub_frac1,
+			p_recenter, p_scale_ax, p_inc_ax, data.has("debias") ? "ON" : "off");
+
+	// PD tiles into a pitched buffer (the kernel takes a stride, as production)
+	GpuPitchedBuf<float> gpu_pd(pd_size, num_corr);
+	gpu_pd.upload(data.hostFloat("corr_pd"));
+	GpuBuf<int> gpu_idx(num_corr);
+	gpu_idx.upload(data.hostInt("corr_indices"));
+	GpuBuf<float> gpu_peaks((size_t) num_corr * PEAK_ROW_FLOATS);
+	GpuBuf<float> gpu_debias(pd_size);
+	float * d_debias = nullptr;
+	if (data.has("debias")) {
+		gpu_debias.upload(data.hostFloat("debias"));
+		d_debias = gpu_debias.dev();
+	}
+
+	corr2D_peak_eig<<<num_corr, 256>>>(
+			num_corr, gpu_idx.dev(),
+			0xff,                                   // sum slots only (interleaved 0x00 slots skipped)
+			gpu_pd.pitchFloats(), gpu_pd.dev(), corr_rad,
+			p_abs_min, p_rel_min, p_min_peak, p_sub_frac,
+			p_recenter, p_sub_frac1, p_scale_ax, p_inc_ax,
+			d_debias, gpu_peaks.dev());
+	checkCudaErrors(cudaGetLastError());
+	checkCudaErrors(cudaDeviceSynchronize());
+	std::vector<float> got = gpu_peaks.download();
+
+	const int * pd_idx = data.hostInt("expected_pd_indices");
+	const int   n_exp  = (int) data.elements("expected_pd_indices");
+	const float * exp_pk = data.hostFloat("expected_peaks");
+	const int * idx = data.hostInt("corr_indices");
+	int fail = 0;
+	// non-sum slots must be skipped (NaN + valid 0)
+	for (int i = 0; i < num_corr; i++) {
+		if ((idx[i] & ((1 << CORR_NTILE_SHIFT) - 1)) != 0xff) {
+			const float * g = got.data() + (size_t) i * PEAK_ROW_FLOATS;
+			if (!std::isnan(g[0]) || (g[PEAK_ROW_FLOATS - 1] != 0.0f)) {
+				printf("slot %d (pair 0x%02x): NOT skipped -> FAIL\n", i, idx[i] & 0xff);
+				fail = 1;
+			}
+		}
+	}
+	double maxd = 0;
+	long nan_mm = 0;
+	for (int i = 0; i < n_exp; i++) {
+		// find the sum slot of this dense tile
+		const int packed = (pd_idx[i] << CORR_NTILE_SHIFT) | 0xff;
+		int row = -1;
+		for (int j = 0; j < num_corr; j++) if (idx[j] == packed) { row = j; break; }
+		if (row < 0) { printf("tile %d: sum slot missing -> FAIL\n", pd_idx[i]); fail = 1; continue; }
+		const float * g = got.data() + (size_t) row * PEAK_ROW_FLOATS;
+		const float * e = exp_pk + (size_t) i * 7;
+		double tile_maxd = 0;
+		for (int j = 0; j < 7; j++) {
+			const bool gn = std::isnan(g[j]), en = std::isnan(e[j]);
+			if (gn || en) { if (gn != en) nan_mm++; continue; }
+			const double d = std::fabs((double) g[j] - (double) e[j]);
+			if (d > tile_maxd) tile_maxd = d;
+		}
+		if (tile_maxd > maxd) maxd = tile_maxd;
+		printf("tile %2d: got {%9.6f,%9.6f,%8.5f | %8.5f,%8.5f,%9.4f,%9.4f} v=%g max|diff|=%g%s\n",
+				pd_idx[i], g[0], g[1], g[2], g[3], g[4], g[5], g[6], g[7], tile_maxd,
+				(tile_maxd > tol) ? " <- OVER TOL" : "");
+	}
+	if ((maxd > tol) || (nan_mm > 0)) fail = 1;
+	printf("compare 'peaks' (%d tiles): max|diff|=%g NaN-mismatch=%ld tol=%g -> %s\n",
+			n_exp, maxd, nan_mm, tol, ((maxd > tol) || nan_mm) ? "FAIL" : "PASS");
+	printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+	return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

src/tests/test_pose_corr.cu

@@ -45,7 +45,19 @@
  * --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]
+ * CHECKPOINT 2 (rung 2, added 07/13/2026), when the case carries the
+ * expected_pd_indices_it<k> buffers (re-export with the rung-2 Java):
+ *   corr2D_normalize IN PLACE on the inter TD buffer (weights NULL == the
+ *   production uniform 1.0) -> compare pixel-domain tiles vs expected_corr_pd
+ *   (dense-tile-indexed, matched via packed (tile<<8)|0xff, --tol tier) ->
+ *   corr2D_peak_eig (the frozen 15-arg getMaxXYCmEig contract, params from
+ *   the manifest incl. optional debias array) -> compare {dx,dy,str,eigen}
+ *   vs the double-precision Java oracle rows at --peak-tol (default 1e-6:
+ *   the kernel computes in double but offline-nvcc FMA/codegen and the float
+ *   PD input make bit-parity with Java unattainable - the tol-0 verdict is
+ *   the NVRTC replay, test_pose_corr_jna). Old cases skip with a note.
+ *
+ * Usage: test_pose_corr --data <dir> [--tol X] [--peak-tol X] [--list]
  *
  *  Created on: Jul 13, 2026
  *      Author: Claude (Anthropic), for Elphel
@@ -67,6 +79,7 @@
 #include "geometry_correction.h"   // struct gc/corr_vector/trot_deriv, geometry kernels, TP_TASK_*_OFFSET
 #include "TileProcessor.h"         // convert_direct, correlate2D_inter, erase_clt_tiles
 #include "tp_consolidate.h"        // index_consolidate, consolidate_oob, clt_average_sensors_list
+#include "tp_peak.h"               // corr2D_peak_eig (checkpoint 2, rung 2) // By Claude on 07/13/2026
 #include "tests/tp_test_data.h"
 #include "tests/tp_gpu_buf.h"
 
@@ -121,17 +134,19 @@ int main(int argc, char **argv)
 {
 	std::string data_dir = "testdata/pose_corr";
 	float tol       = 0.0f;   // the whole chain is expected bit-exact vs the oracle
+	float peak_tol  = 1e-6f;  // Java-double-oracle tier for the peak kernel (header note)
 	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], "--peak-tol") && (i + 1 < argc)) peak_tol = (float) atof(argv[++i]);
 		else if (!strcmp(argv[i], "--list"))                   list_only = 1;
 		else {
-			fprintf(stderr, "Usage: %s --data <dir> [--tol X] [--list]\n", argv[0]);
+			fprintf(stderr, "Usage: %s --data <dir> [--tol X] [--peak-tol X] [--list]\n", argv[0]);
 			return EXIT_FAILURE;
 		}
 	}
-	printf("%s: data=%s tol=%g\n", argv[0], data_dir.c_str(), tol);
+	printf("%s: data=%s tol=%g peak_tol=%g\n", argv[0], data_dir.c_str(), tol, peak_tol);
 
 	findCudaDevice(argc, (const char **) argv);
 	TpTestData data(data_dir);
@@ -252,6 +267,37 @@ int main(int argc, char **argv)
 	GpuBuf<float> gpu_result_scratch(max_corr * chunk);      // reordered result upload for saveResult
 	const size_t corr_stride = gpu_corrs_td.pitchFloats();
 
+	// ---- checkpoint 2 (rung 2) statics: FZ-norm in place + frozen peak contract ----
+	// By Claude on 07/13/2026
+	const int  corr_rad  = DTT_SIZE - 1;                                 // 7 -> 15x15
+	const int  pd_size   = (2 * corr_rad + 1) * (2 * corr_rad + 1);      // 225
+	const bool have_ckpt2 = data.has("expected_pd_indices_it0");
+	const float fz_inter = (float) data.param("fz_inter", 10000.0);
+	const float fz2 = (float) ((double) fz_inter * (double) fz_inter);   // GpuQuad: (float)(fat_zero*fat_zero)
+	const float p_abs_min   = (float) data.param("eig_min_abs",   0.0);
+	const float p_rel_min   = (float) data.param("eig_min_rel",   0.0);
+	const float p_min_peak  = (float) data.param("min_str_sum",   0.0);
+	const float p_sub_frac  = (float) data.param("eig_sub_frac",  0.0);
+	const int   p_recenter  =         data.paramInt("pose_recenter", 0);
+	const float p_sub_frac1 = (float) data.param("eig_sub_frac1", 0.0);
+	const float p_scale_ax  = (float) data.param("eig_scale_axes",1.2);
+	const float p_inc_ax    = (float) data.param("eig_inc_axes",  1.0);
+	GpuPitchedBuf<float> gpu_corrs_pd(pd_size, (int) max_corr);
+	GpuBuf<float> gpu_peaks(max_corr * PEAK_ROW_FLOATS);
+	GpuBuf<float> gpu_debias(pd_size);
+	float * d_debias = nullptr;                                          // NULL = de-bias off (the contract)
+	if (have_ckpt2 && data.has("debias")) {
+		gpu_debias.upload(data.hostFloat("debias"));
+		d_debias = gpu_debias.dev();
+	}
+	if (have_ckpt2) {
+		printf("checkpoint 2: fz=%g recenter=%d debias=%s peak params abs/rel/min=%g/%g/%g sub=%g/%g axes=%gx+%g\n",
+				fz_inter, p_recenter, d_debias ? "ON" : "off",
+				p_abs_min, p_rel_min, p_min_peak, p_sub_frac, p_sub_frac1, p_scale_ax, p_inc_ax);
+	} else {
+		printf("NOTE: checkpoint-2 buffers absent (pre-rung-2 case vintage) - FZ-norm/peak replay skipped\n");
+	}
+
 	int fail = 0;
 	const int IDX_THREADS = 256;
 	for (int k = 0; k < num_iters; k++) {
@@ -420,6 +466,88 @@ int main(int argc, char **argv)
 				reordered.size() * sizeof(float), cudaMemcpyHostToDevice));
 		snprintf(nm, sizeof(nm), "corr_td_it%d", k);
 		data.saveResult(nm, gpu_result_scratch.dev(), {n_exp, chunk});
+
+		// ---- CHECKPOINT 2 (rung 2): FZ-norm IN PLACE + peak kernel ----
+		// By Claude on 07/13/2026
+		if (have_ckpt2) {
+			snprintf(nm, sizeof(nm), "expected_pd_indices_it%d", k);
+			const int * pd_idx = data.hostInt(nm);
+			const int   n_pd   = (int) data.elements(nm);
+			snprintf(nm, sizeof(nm), "expected_corr_pd_it%d", k);
+			const float * exp_pd = data.hostFloat(nm);
+			snprintf(nm, sizeof(nm), "expected_peaks_it%d", k);
+			const float * exp_pk = data.hostFloat(nm);
+			// normalize ALL resident slots where they sit (per-tile independent math;
+			// NULL weights == the production uniform 1.0 - see TDCorrTile.convertTDtoPDInPlace)
+			corr2D_normalize<<<1, 1>>>(
+					num_corr, corr_stride, gpu_corrs_td.dev(),
+					(float *) NULL,
+					gpu_corrs_pd.pitchFloats(), gpu_corrs_pd.dev(),
+					fz2, corr_rad);
+			checkCudaErrors(cudaGetLastError());
+			checkCudaErrors(cudaDeviceSynchronize());
+			corr2D_peak_eig<<<num_corr, 256>>>(
+					num_corr, gpu_corr_indices.dev(),
+					0xff,                                   // lean processes the sum slot only
+					gpu_corrs_pd.pitchFloats(), gpu_corrs_pd.dev(), corr_rad,
+					p_abs_min, p_rel_min, p_min_peak, p_sub_frac,
+					p_recenter, p_sub_frac1, p_scale_ax, p_inc_ax,
+					d_debias, gpu_peaks.dev());
+			checkCudaErrors(cudaGetLastError());
+			checkCudaErrors(cudaDeviceSynchronize());
+			std::vector<float> got_pd((size_t) num_corr * pd_size);
+			checkCudaErrors(cudaMemcpy2D(got_pd.data(), (size_t) pd_size * sizeof(float),
+					gpu_corrs_pd.dev(), gpu_corrs_pd.pitchBytes(),
+					(size_t) pd_size * sizeof(float), num_corr, cudaMemcpyDeviceToHost));
+			std::vector<float> got_pk((size_t) num_corr * PEAK_ROW_FLOATS);
+			checkCudaErrors(cudaMemcpy(got_pk.data(), gpu_peaks.dev(),
+					got_pk.size() * sizeof(float), cudaMemcpyDeviceToHost));
+			// match by packed index (dense tile -> (tile<<shift)|0xff sum slot), reusing
+			// the checkpoint-1 map; save results in ORACLE row order for the Java round trip
+			std::vector<float> res_pd((size_t) n_pd * pd_size, NAN);
+			std::vector<float> res_pk((size_t) n_pd * 7, NAN);
+			double pd_maxd = 0, pk_maxd = 0;
+			long pd_nan = 0, pk_nan = 0, pd_missing = 0;
+			for (int i = 0; i < n_pd; i++) {
+				const int packed = (pd_idx[i] << CORR_NTILE_SHIFT) | 0xff;
+				const std::map<int, int>::const_iterator it = got_by_index.find(packed);
+				if (it == got_by_index.end()) { pd_missing++; continue; }
+				const float * gpd = got_pd.data() + (size_t) it->second * pd_size;
+				const float * epd = exp_pd + (size_t) i * pd_size;
+				memcpy(res_pd.data() + (size_t) i * pd_size, gpd, pd_size * sizeof(float));
+				for (int j = 0; j < pd_size; j++) {
+					const bool gn = std::isnan(gpd[j]), en = std::isnan(epd[j]);
+					if (gn || en) { if (gn != en) pd_nan++; continue; }
+					const double d = std::fabs((double) gpd[j] - (double) epd[j]);
+					if (d > pd_maxd) pd_maxd = d;
+				}
+				const float * gpk = got_pk.data() + (size_t) it->second * PEAK_ROW_FLOATS;
+				const float * epk = exp_pk + (size_t) i * 7;
+				memcpy(res_pk.data() + (size_t) i * 7, gpk, 7 * sizeof(float));
+				for (int j = 0; j < 7; j++) {
+					const bool gn = std::isnan(gpk[j]), en = std::isnan(epk[j]);
+					if (gn || en) { if (gn != en) pk_nan++; continue; }
+					const double d = std::fabs((double) gpk[j] - (double) epk[j]);
+					if (d > pk_maxd) pk_maxd = d;
+				}
+			}
+			const int pd_fail = (pd_maxd > tol)      || (pd_nan > 0) || (pd_missing > 0);
+			const int pk_fail = (pk_maxd > peak_tol) || (pk_nan > 0) || (pd_missing > 0);
+			printf("it%d: compare 'corr_pd' (%d tiles, in-place FZ-norm): max|diff|=%g NaN-mismatch=%ld "
+					"missing=%ld -> %s\n", k, n_pd, pd_maxd, pd_nan, pd_missing, pd_fail ? "FAIL" : "PASS");
+			printf("it%d: compare 'peaks' (%d tiles vs the double Java oracle): max|diff|=%g "
+					"NaN-mismatch=%ld -> %s (peak_tol %g)\n",
+					k, n_pd, pk_maxd, pk_nan, pk_fail ? "FAIL" : "PASS", peak_tol);
+			if (pd_fail || pk_fail) fail = 1;
+			checkCudaErrors(cudaMemcpy(gpu_result_scratch.dev(), res_pd.data(),
+					res_pd.size() * sizeof(float), cudaMemcpyHostToDevice));
+			snprintf(nm, sizeof(nm), "corr_pd_it%d", k);
+			data.saveResult(nm, gpu_result_scratch.dev(), {n_pd, pd_size});
+			checkCudaErrors(cudaMemcpy(gpu_result_scratch.dev(), res_pk.data(),
+					res_pk.size() * sizeof(float), cudaMemcpyHostToDevice));
+			snprintf(nm, sizeof(nm), "peaks_it%d", k);
+			data.saveResult(nm, gpu_result_scratch.dev(), {n_pd, 7});
+		}
 	}
 	printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
 	return fail ? EXIT_FAILURE : EXIT_SUCCESS;

src/tp_peak.cu

+/*
+ * tp_peak.cu
+ *
+ * corr2D_peak_eig — per-tile argmax/centroid/eigen peak measurement, the
+ * frozen 15-arg getMaxXYCmEig contract (see tp_peak.h for the contract and
+ * precision notes). One block per correlation tile: the parallel threads do
+ * the load/de-bias, NaN scan, first-max argmax reduction, thresholding and
+ * the 8-connected region grow; thread 0 then runs the centroid/covariance/
+ * eigen/recentering math SERIALLY in double, in the exact linescan summation
+ * order of the Java oracle (Correlation2d.getMaxXYCmEig, eig_fast2x2 branch) —
+ * 225 pixels/tile make the serial tail negligible and keep the result
+ * order-deterministic.
+ *
+ *  Created on: Jul 13, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#ifndef JCUDA // NVRTC concatenation (tp_jna.cpp TP_DEFINES, headers merged in order) has no include paths
+#include <cuda_runtime.h>              // __global__/__shared__ also for non-nvcc parsers
+#include <device_launch_parameters.h>  // threadIdx/blockIdx: cuda_runtime.h gates this on __CUDACC__
+#include <math.h>                      // NAN, isnan, sqrt, cos
+#include "tp_defines.h"                // CORR_NTILE_SHIFT
+#include "tp_peak.h"
+#endif
+#ifndef NAN // NVRTC has isnan built in but no math.h NAN/INFINITY macros (device-only use is fine)
+#define NAN __int_as_float(0x7fc00000)
+#endif
+#ifdef __CDT_PARSER__                  // Eclipse CDT indexer only, nvcc never defines this
+extern "C" void __syncthreads();       // real decl is __CUDACC__-gated in crt/device_functions.h
+#endif
+
+#define PEAK_THREADS 256               // block size (>= PEAK_MAX_PIX)
+#define PEAK_HALF_PI 1.5707963267948966 // 0.5 * Java Math.PI, for the recentering cosine mask
+
+/**
+ * getEigen2x2 port (Correlation2d.getEigen2x2, double, verbatim semantics):
+ * eigenvalues/eigenvectors of the symmetric 2x2 covariance {{cxx,cxy},{cxy,cyy}}.
+ * Returns 0 on the oracle's null (l == 0), else fills
+ * eig = {e0x, e0y, lambda0, lambda1} (lambda0 <= lambda1, vector of lambda0,
+ * sign-normalized so e0x >= 0 — "to be compatible with standard").
+ */
+__device__ static int peak_eigen2x2(double cxx, double cxy, double cyy, double * eig)
+{
+	const double e = 1E-12;
+	double hapd = (cxx + cyy) / 2;
+	double hamd = (cxx - cyy) / 2;
+	double bc   = cxy * cxy;                 // A[1][0]*A[0][1], symmetric
+	double d    = sqrt(hamd * hamd + bc);
+	if (d / fabs(hapd) < e) {                // near-isotropic: any direction
+		eig[0] = 1.0; eig[1] = 0.0; eig[2] = hapd; eig[3] = hapd;
+		return 1;
+	}
+	double lambda0 = hapd - d, lambda1 = hapd + d;
+	double v0x = cxx - lambda1, v0y = cxy;
+	double l = sqrt(v0x * v0x + v0y * v0y);
+	if (l == 0.0) {
+		return 0;                            // oracle returns null
+	}
+	double k = 1.0 / l;
+	if (v0x < 0) {
+		k = -k;
+	}
+	eig[0] = k * v0x; eig[1] = k * v0y; eig[2] = lambda0; eig[3] = lambda1;
+	return 1;
+}
+
+extern "C" __global__ void corr2D_peak_eig(
+		int           num_corr_tiles,
+		const int   * gpu_corr_indices,
+		int           pair_select,
+		const size_t  corr_stride,
+		const float * gpu_corrs,
+		int           corr_radius,
+		float         abs_min,
+		float         rel_min,
+		float         min_peak,
+		float         eig_sub_frac,
+		int           refine,
+		float         eig_sub_frac1,
+		float         scale_axes,
+		float         inc_axes,
+		const float * gpu_debias,
+		float       * gpu_peaks)
+{
+	const int slot = blockIdx.x;
+	const int tid  = threadIdx.x;
+	if (slot >= num_corr_tiles) return;
+	const int width  = 2 * corr_radius + 1;
+	const int npix   = width * width;
+	const int center = corr_radius;                       // (width-1)/2 = transform_size-1
+	float * out = gpu_peaks + (size_t) slot * PEAK_ROW_FLOATS;
+
+	__shared__ double sd  [PEAK_MAX_PIX];                 // (de-biased) tile, double = the oracle's input
+	__shared__ char   above[PEAK_MAX_PIX];                // >= min_d threshold mask
+	__shared__ char   en   [PEAK_MAX_PIX];                // 8-connected region grown from the argmax
+	__shared__ double red_v[PEAK_THREADS];                // argmax reduction: values
+	__shared__ int    red_i[PEAK_THREADS];                //                   first-max indices
+	__shared__ int    sh_flag;                            // NaN found / fill-pass changed
+	__shared__ int    sh_reject;                          // block-wide early-out (syncthreads-safe)
+
+	// skip non-selected slots (lean processes the 0xff sum slot only)
+	if ((pair_select >= 0) && ((gpu_corr_indices[slot] & ((1 << CORR_NTILE_SHIFT) - 1)) != pair_select)) {
+		if (tid == 0) {
+			for (int i = 0; i < PEAK_ROW_FLOATS - 1; i++) out[i] = NAN;
+			out[PEAK_ROW_FLOATS - 1] = 0.0f;
+		}
+		return;
+	}
+
+	// ---- load (+ optional de-bias), NaN scan ----
+	// float -> double is exact, so sd matches the Java oracle's input bit-for-bit;
+	// with de-bias the oracle multiplies the double envelope while the kernel gets
+	// the float-cast array — an epsilon-tier difference, default OFF (contract note).
+	if (tid == 0) sh_flag = 0;
+	__syncthreads();
+	const float * src = gpu_corrs + (size_t) slot * corr_stride;
+	for (int i = tid; i < npix; i += PEAK_THREADS) {
+		float v = src[i];
+		if (isnan(v)) sh_flag = 1;                        // benign race: any writer sets 1
+		sd[i] = gpu_debias ? ((double) v * (double) gpu_debias[i]) : (double) v;
+	}
+	__syncthreads();
+	if (sh_flag) {                                        // oracle: NaN anywhere -> null
+		if (tid == 0) {
+			for (int i = 0; i < PEAK_ROW_FLOATS - 1; i++) out[i] = NAN;
+			out[PEAK_ROW_FLOATS - 1] = 0.0f;
+		}
+		return;
+	}
+
+	// ---- argmax, FIRST occurrence in linescan order (oracle: strict > scan) ----
+	{
+		double bv = sd[0];                                // every thread covers <= 1 extra pixel
+		int    bi = 0;
+		for (int i = tid; i < npix; i += PEAK_THREADS) {
+			if (sd[i] > bv) { bv = sd[i]; bi = i; }
+		}
+		red_v[tid] = bv; red_i[tid] = bi;
+		__syncthreads();
+		for (int s = PEAK_THREADS / 2; s > 0; s >>= 1) {
+			if (tid < s) {
+				if ((red_v[tid + s] > red_v[tid]) ||
+						((red_v[tid + s] == red_v[tid]) && (red_i[tid + s] < red_i[tid]))) {
+					red_v[tid] = red_v[tid + s];
+					red_i[tid] = red_i[tid + s];
+				}
+			}
+			__syncthreads();
+		}
+	}
+	const int    imax = red_i[0];
+	const double mx   = red_v[0];
+	if (tid == 0) sh_reject = (mx < (double) min_peak);   // oracle: too weak -> null
+	__syncthreads();
+	if (sh_reject) {
+		if (tid == 0) {
+			for (int i = 0; i < PEAK_ROW_FLOATS - 1; i++) out[i] = NAN;
+			out[PEAK_ROW_FLOATS - 1] = 0.0f;
+		}
+		return;
+	}
+	const int ix0 = imax % width;
+	const int iy0 = imax / width;
+
+	// ---- threshold + 8-connected region grow from the argmax seed ----
+	// the region is a SET (order-independent), so the parallel grow-until-stable
+	// loop yields exactly the oracle's TileNeibs.getConnected (dirs = 8) result
+	const double min_d = fmin((double) abs_min, (double) rel_min * mx);
+	for (int i = tid; i < npix; i += PEAK_THREADS) {
+		above[i] = (sd[i] >= min_d);
+		en[i] = 0;
+	}
+	__syncthreads();
+	if (tid == 0) en[imax] = above[imax];                 // oracle: seed not above -> empty region
+	__syncthreads();
+	for (;;) {
+		if (tid == 0) sh_flag = 0;
+		__syncthreads();
+		for (int i = tid; i < npix; i += PEAK_THREADS) {
+			if (above[i] && !en[i]) {
+				const int ix = i % width, iy = i / width;
+				char hit = 0;
+				for (int dy = -1; dy <= 1 && !hit; dy++) {
+					const int y = iy + dy;
+					if ((y < 0) || (y >= width)) continue;
+					for (int dx = -1; dx <= 1; dx++) {
+						const int x = ix + dx;
+						if ((x < 0) || (x >= width) || ((dx == 0) && (dy == 0))) continue;
+						if (en[y * width + x]) { hit = 1; break; }
+					}
+				}
+				if (hit) { en[i] = 1; sh_flag = 1; }
+			}
+		}
+		__syncthreads();
+		if (!sh_flag) break;
+		__syncthreads();                                  // writers of the next round wait for readers
+	}
+
+	// ---- serial double tail: centroid/covariance/eigen/recentering (thread 0) ----
+	// exact oracle summation order (linescan, sequential) - see tp_peak.h precision notes
+	if (tid != 0) return;
+	const double sub_pedestal  = min_d * (double) eig_sub_frac;
+	const double sub_pedestal1 = min_d * (double) eig_sub_frac1;
+	double s0 = 0, sx = 0, sy = 0, sx2 = 0, sy2 = 0, sxy = 0;
+	for (int iy = 0; iy < width; iy++) {
+		const double y = iy - (double) iy0;
+		for (int ix = 0; ix < width; ix++) {
+			const int indx = iy * width + ix;
+			if (en[indx]) {
+				const double x = ix - (double) ix0;
+				const double d = sd[indx] - sub_pedestal;
+				s0  += d;
+				sx  += d * x;
+				sy  += d * y;
+				sx2 += d * x * x;
+				sy2 += d * y * y;
+				sxy += d * x * y;
+			}
+		}
+	}
+	const double x0 = (double) ix0 + sx / s0;             // relative to top-left
+	const double y0 = (double) iy0 + sy / s0;
+	const double cxx = sx2 - sx * sx / s0;
+	const double cyy = sy2 - sy * sy / s0;
+	const double cxy = sxy - sx * sy / s0;
+	double eig[4];
+	if ((sx2 == 0) || (sy2 == 0) || (cxx == 0) || (cyy == 0) ||    // oracle degenerate checks
+			!peak_eigen2x2(cxx, cxy, cyy, eig)) {
+		for (int i = 0; i < PEAK_ROW_FLOATS - 1; i++) out[i] = NAN;
+		out[PEAK_ROW_FLOATS - 1] = 0.0f;
+		return;
+	}
+	double xy0 = x0 - center, xy1 = y0 - center;          // relative to the tile center
+
+	if (refine > 0) {                                     // recentering passes (runtime contract param)
+		const double ha0 = (double) scale_axes * sqrt(eig[2]) + (double) inc_axes;
+		const double ha1 = (double) scale_axes * sqrt(eig[3]) + (double) inc_axes;
+		const double t00 = eig[0] / ha0, t01 = eig[1] / ha0;    // ellipse -> unit-disk transform
+		const double t10 = eig[1] / ha1, t11 = -eig[0] / ha1;   // (fixed across passes, as the oracle)
+		for (int nref = 1; nref <= refine; nref++) {
+			double r0 = 0, rx = 0, ry = 0;
+			for (int iy = 0; iy < width; iy++) {
+				const double y = iy - xy1 - center;
+				for (int ix = 0; ix < width; ix++) {
+					const double x  = ix - xy0 - center;
+					const double de0 = t00 * x + t01 * y;
+					const double de1 = t10 * x + t11 * y;
+					const double er2 = de0 * de0 + de1 * de1;
+					if (er2 < 1.0) {
+						double d = sd[iy * width + ix] - sub_pedestal1;
+						if (d > 0) {                      // ignore negative
+							d *= cos(PEAK_HALF_PI * sqrt(er2));
+							r0 += d;
+							rx += d * x;
+							ry += d * y;
+						}
+					}
+				}
+			}
+			if (r0 > 0) {
+				xy0 += rx / r0;
+				xy1 += ry / r0;
+			}
+		}
+	}
+	out[0] = (float) xy0;
+	out[1] = (float) xy1;
+	out[2] = (float) mx;
+	out[3] = (float) eig[0];
+	out[4] = (float) eig[1];
+	out[5] = (float) eig[2];
+	out[6] = (float) eig[3];
+	out[7] = 1.0f;
+}

src/tp_peak.h

+/*
+ * tp_peak.h
+ *
+ * Per-tile correlation peak measurement kernel — the GPU side of the lean
+ * pose-chain CHECKPOINT 2 second half (roadmap 3-C rung 2).
+ *
+ * FROZEN CONTRACT (Andrey's ruling 07/13/2026): the 15-arg
+ * Correlation2d.getMaxXYCmEig superset —
+ *   argmax (first occurrence in linescan order) -> weak-peak reject
+ *   -> 8-connected region above min(abs_min, rel_min*max) grown from the
+ *      argmax seed -> pedestal-subtracted centroid + covariance
+ *   -> closed-form 2x2 eigen (getEigen2x2 port, eig_fast2x2 == the contract)
+ *   -> OPTIONAL recentering passes (refine > 0): cosine-mask window over the
+ *      eigen-ellipse (scale_axes*sqrt(lambda)+inc_axes half-axes),
+ *      re-centroid each pass
+ *   -> OPTIONAL envelope de-bias: gpu_debias (NULL = off) is multiplied into
+ *      the tile before ANY peak math (the window-autocorr envelope
+ *      under-reports offsets by ~0.75-0.9x; array from
+ *      CuasPoseRT.windowAutocorrDebias, single Java source).
+ * NOT in the contract: the oracle's fpn_mask/ignore_border branch — the lean
+ * chain always passes null (input is FPN-subtracted); host-side policy.
+ *
+ * All accumulation and the eigen/refine math run in DOUBLE with the SAME
+ * summation order as the Java oracle (serial linescan), so the float-input /
+ * double-math results track the double oracle to rounding (offline-nvcc vs
+ * NVRTC FMA/codegen caveats per the rung-1 lesson: the tol-0 verdict is
+ * NVRTC-replay-vs-production-capture, Java-oracle compare is an epsilon tier).
+ *
+ * Java oracle: Correlation2d.getMaxXYCmEig (15-arg, eig_fast2x2=true);
+ * case export: PoseCorrExport (curt.kernel_test = pose_corr, checkpoint 2).
+ *
+ *  Created on: Jul 13, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#ifndef SRC_TP_PEAK_H_
+#define SRC_TP_PEAK_H_
+
+#ifndef JCUDA // NVRTC concatenation (tp_jna.cpp TP_DEFINES) has no include paths
+#include <cuda_runtime.h> // __global__ in the prototype below, also for non-nvcc parsers
+#endif
+
+// correlation tile side bound: corr_radius <= DTT_SIZE-1 = 7 -> 15x15
+#define PEAK_MAX_WIDTH  15
+#define PEAK_MAX_PIX    (PEAK_MAX_WIDTH * PEAK_MAX_WIDTH)
+// output row stride (floats): {dx, dy, strength, e0x, e0y, lambda0, lambda1, valid}
+#define PEAK_ROW_FLOATS 8
+
+extern "C" __global__ void corr2D_peak_eig(
+		int           num_corr_tiles,   // resident correlation tiles (one block each)
+		const int   * gpu_corr_indices, // packed (tile << CORR_NTILE_SHIFT) | pair (0xff = sum)
+		int           pair_select,      // process only slots with this pair code (lean: 0xff); -1 = all
+		const size_t  corr_stride,      // in floats: pixel-domain correlation buffer stride
+		const float * gpu_corrs,        // FZ-normalized pixel-domain correlation tiles
+		int           corr_radius,      // 7 -> 15x15
+		float         abs_min,          // imp.eig_min_abs
+		float         rel_min,          // imp.eig_min_rel
+		float         min_peak,         // imp.min_str_sum
+		float         eig_sub_frac,     // pedestal subtract fraction (centroid/covariance)
+		int           refine,           // recentering passes (curt.pose_recenter; 0 = single-pass)
+		float         eig_sub_frac1,    // pedestal subtract during recentering
+		float         scale_axes,       // eigen-ellipse half-axis scale (imp.eig_scale_axes)
+		float         inc_axes,         // eigen-ellipse half-axis increment, pix (imp.eig_inc_axes)
+		const float * gpu_debias,       // NULL or [width*width] multiplicative envelope de-bias
+		float       * gpu_peaks);       // [num_corr_tiles][PEAK_ROW_FLOATS]; rejected/skipped ->
+		                                // 7x NaN + valid 0.0f; measured -> values + valid 1.0f
+
+#endif /* SRC_TP_PEAK_H_ */