CLAUDE: T2a - 16-bit feeder native path: u16 staging + one-pass convert/RowCol/condition kernel

condition_lwir_u16: byte-swap (big-endian TIFF strip) + borrowed per-scene
Row/Col in explicit-RN double (the exact applyCalibratedRowColToRaw chain),
ONE double->float cast, then the unchanged production two-map fmaf - all in
one pass from a device u16 scratch into the ACTIVE image set on the copy
stream (D5a ready-event contract kept; serial path works without a prior
img_set_select via lazy ready-event creation).

New entries: tp_proc_bayer_staging_u16 (pinned, 2 slots - the Java feeder
preads raw file bytes straight into it), tp_proc_rowcol_staging (pinned
rows|cols|inv_scales doubles), tp_proc_set_image_staged_u16,
tp_proc_set_rowcol_staged, tp_proc_exec_conditioning_u16. Kernel count 41->42.

Case u16_cond (synthetic, tol 0): host-reference bit-exact, u16-vs-float
path equivalence (the preload byte-identity contract), no-swap/no-rowcol
variant. run_cases.sh ALL PASS (old + new); compute-sanitizer 0 errors.
Co-authored-by: Claude Fable 5 <claude-fable-5@anthropic.com>

cases.list.tmpl

@@ -70,6 +70,7 @@ case dp_cycles       test_pose_lma_dp_cycles_jna ${REPO}
 case measure_dp      test_pose_measure_dp_jna    ${MODEL}/testdata/pose_corr     --tol 0     # rung D2: device measure chain bit-exact vs B3 host-batched entry, keyed by packed index # By Claude on 07/17/2026
 case scene_dp        test_pose_scene_dp_jna      ${MODEL}/testdata/pose_corr     --tol 0     # rung D3: full per-scene DP parent bit-exact vs host-driven B3 loop (full+light chains) # By Claude on 07/17/2026
 case scene_dp_split  test_pose_scene_dp_split_jna ${MODEL}/testdata/pose_corr    --tol 0     # rung D5b: launch/collect split bit-exact vs host loop + concurrent ring-set-1 upload/conditioning stress # By Claude on 07/17/2026
+case u16_cond        test_u16_cond_jna     ${REPO}                               --tol 0     # T2 16-bit feeder: u16 staging + swap/RowCol/two-map kernel bit-exact vs host ref AND vs the float/preload path (synthetic) # By Claude on 07/18/2026
 #case convert_direct test_convert_direct   ${REPO}/testdata/convert_direct_legacy               # local blessed goldens (~200MB, make_convert_direct_case.py)
 
 # gen <name(s)>                      <generator command, run from ${REPO}>

jna/test_u16_cond_jna.cu

+// test_u16_cond_jna.cu - validate the T2 16-bit feeder path (libtileproc.so):
+// pinned u16 staging -> async H2D -> condition_lwir_u16 (byte-swap + borrowed
+// Row/Col in explicit-RN double + ONE double->float cast + the production
+// two-map fmaf) into the ACTIVE image set. Fully synthetic (no case data):
+//   1. HOST REFERENCE: per-pixel expected value computed with the exact Java
+//      chain (double row/col subtraction, single float cast, fmaf) -> readback
+//      must be BIT-EXACT;
+//   2. FLOAT-PATH EQUIVALENCE (the production gate): feed the float staging
+//      path with the host-computed raw floats (what CuasScenePreload built)
+//      and run condition_lwir_f -> both paths' conditioned images must be
+//      BIT-IDENTICAL. This is the preload-vs-feeder byte-identity contract.
+//   3. ROWCOL-DISABLED + NO-SWAP variant of (1) (FLIR-corrected production).
+// All comparisons tol 0 by design.
+//
+// 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_u16_cond_jna.cu \
+//     -o tests_bin/test_u16_cond_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_u16_cond_jna   (cwd = repo root; no --data needed)
+//
+//  Created on: Jul 18, 2026
+//      Author: Claude (Anthropic), for Elphel
+#include <cstdio>
+#include <cstring>
+#include <cstdlib>
+#include <cmath>
+
+extern "C" {
+  void* tp_create_module(const char*, const char*);
+  const char* tp_last_error();
+  void* tp_proc_create(void*);
+  int   tp_proc_setup(void*, int,int,int,int,int,int);
+  float* tp_proc_bayer_staging(void*, int);
+  unsigned short* tp_proc_bayer_staging_u16(void*, int);
+  double* tp_proc_rowcol_staging(void*, int);
+  int   tp_proc_set_image_staged(void*, int, int);
+  int   tp_proc_set_image_staged_u16(void*, int, int);
+  int   tp_proc_set_rowcol_staged(void*, int);
+  int   tp_proc_set_conditioning_maps(void*, const float*, const float*, int);
+  int   tp_proc_exec_conditioning_u16(void*, int, int);
+  int   tp_proc_exec_conditioning(void*, int, double);
+  int   tp_proc_get_image(void*, int, float*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+static long compare_bits(const float* a, const float* b, size_t n){
+  long bad=0; for(size_t i=0;i<n;i++) if(memcmp(&a[i],&b[i],sizeof(float))) bad++; return bad;
+}
+// Exact Java raw chain: raw = (double)u16 - (row+col)*inv, one float cast.
+// volatile intermediates forbid host FP contraction (the kernel uses __d*_rn).
+static float raw_float(unsigned short v, double row, double col, double inv, int rowcol){
+  volatile double raw = (double) v;
+  if (rowcol) {
+    volatile double t = row + col;
+    volatile double u = t * inv;
+    raw = raw - u;
+  }
+  return (float) raw;
+}
+
+int main(int argc, char** argv){
+  const char* srcd  = "src";
+  const char* devrt = "/usr/local/cuda-12.8/lib64/libcudadevrt.a";
+  for (int i = 1; i < argc; i++){
+    if      (!strcmp(argv[i],"--src")   && (i+1<argc)) srcd  = argv[++i];
+    else if (!strcmp(argv[i],"--devrt") && (i+1<argc)) devrt = argv[++i];
+    else if (!strcmp(argv[i],"--data")  && (i+1<argc)) ++i;  // run_cases.sh uniformity (unused: synthetic)
+    else if (!strcmp(argv[i],"--tol")   && (i+1<argc)) ++i;  // tol-0 by design
+    else { fprintf(stderr,"Usage: %s [--src <kernel src>] [--devrt <libcudadevrt.a>]\n", argv[0]); return 1; }
+  }
+  const int NC=16, W=640, H=512; const size_t NPIX=(size_t)W*H;
+  void* m = tp_create_module(srcd, devrt);
+  if(!m){ printf("FAIL module: %s\n", tp_last_error()); return 1; }
+  void* p = tp_proc_create(m);
+  if(tp_proc_setup(p, NC, 1, W, H, 0, 0)){ printf("FAIL setup: %s\n", tp_last_error()); return 1; }
+
+  // resident two-map calibration (nontrivial, sensor+pixel dependent)
+  float* scale_maps  = new float[(size_t)NC*NPIX];
+  float* offset_maps = new float[(size_t)NC*NPIX];
+  for(int c=0;c<NC;c++) for(size_t i=0;i<NPIX;i++){
+    scale_maps [(size_t)c*NPIX+i] = 1.0f + 0.001f*c + 1e-8f*(float)i;
+    offset_maps[(size_t)c*NPIX+i] = -3500.0f + 0.5f*c + 0.0001f*(float)(i%997);
+  }
+  if(tp_proc_set_conditioning_maps(p, scale_maps, offset_maps, (int)NPIX)){
+    printf("FAIL set_conditioning_maps: %s\n", tp_last_error()); return 1; }
+
+  // synthetic u16 scene (stored BIG-ENDIAN like the TIFF strip) + rowcol block
+  unsigned short* u16 = tp_proc_bayer_staging_u16(p, 0);
+  if(!u16){ printf("FAIL u16 staging: %s\n", tp_last_error()); return 1; }
+  double* rc = tp_proc_rowcol_staging(p, 0);
+  if(!rc){ printf("FAIL rowcol staging: %s\n", tp_last_error()); return 1; }
+  unsigned short* value = new unsigned short[(size_t)NC*NPIX]; // logical (post-swap) values
+  for(int c=0;c<NC;c++) for(size_t i=0;i<NPIX;i++){
+    unsigned short v = (unsigned short)(17000 + 37*c + (int)((i*2654435761u)%9000));
+    value[(size_t)c*NPIX+i]=v;
+    u16  [(size_t)c*NPIX+i]=(unsigned short)((v>>8)|(v<<8)); // store big-endian
+  }
+  double* rows = rc;                       // [cam][W]
+  double* cols = rc + (size_t)NC*W;        // [cam][H]
+  double* invs = rc + (size_t)NC*(W+H);    // [cam]
+  for(int c=0;c<NC;c++){
+    for(int x=0;x<W;x++) rows[(size_t)c*W+x] = 0.37*c - 11.25 + 0.013*x + 0.7*((x*13)%29);
+    for(int y=0;y<H;y++) cols[(size_t)c*H+y] = -0.21*c + 5.5 - 0.007*y + 0.3*((y*7)%23);
+    invs[c] = 1.0/(0.985 + 0.002*c); // borrowed prior-scale inverse
+  }
+
+  long bad_total = 0;
+  float* got = new float[NPIX];
+  float* expected = new float[NPIX];
+  float* u16_result = new float[(size_t)NC*NPIX];
+
+  // ---- 1. host reference, swap+rowcol ON ----
+  for(int c=0;c<NC;c++)
+    if(tp_proc_set_image_staged_u16(p, c, 0)){ printf("FAIL set_image_staged_u16 cam %d: %s\n", c, tp_last_error()); return 1; }
+  if(tp_proc_set_rowcol_staged(p, 0)){ printf("FAIL set_rowcol_staged: %s\n", tp_last_error()); return 1; }
+  if(tp_proc_exec_conditioning_u16(p, 1, 1)){ printf("FAIL exec_conditioning_u16: %s\n", tp_last_error()); return 1; }
+  long bad1=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL get_image cam %d\n", c); return 1; }
+    memcpy(u16_result+(size_t)c*NPIX, got, NPIX*sizeof(float));
+    for(size_t i=0;i<NPIX;i++){
+      const int x=(int)(i%W), y=(int)(i/W);
+      const float rawf = raw_float(value[(size_t)c*NPIX+i],
+          rows[(size_t)c*W+x], cols[(size_t)c*H+y], invs[c], 1);
+      expected[i] = fmaf(rawf, scale_maps[(size_t)c*NPIX+i], -offset_maps[(size_t)c*NPIX+i]);
+    }
+    bad1 += compare_bits(got, expected, NPIX);
+  }
+  printf("u16 host-reference (swap+rowcol): %d sensors x %zu px, bit-mismatches=%ld -> %s\n",
+         NC, NPIX, bad1, bad1?"FAIL":"PASS");
+  bad_total += bad1;
+
+  // ---- 2. float-path equivalence: preload-style raw floats + condition_lwir_f ----
+  float* fstage = tp_proc_bayer_staging(p, 1);
+  if(!fstage){ printf("FAIL float staging: %s\n", tp_last_error()); return 1; }
+  for(int c=0;c<NC;c++){
+    for(size_t i=0;i<NPIX;i++){
+      const int x=(int)(i%W), y=(int)(i/W);
+      fstage[(size_t)c*NPIX+i] = raw_float(value[(size_t)c*NPIX+i],
+          rows[(size_t)c*W+x], cols[(size_t)c*H+y], invs[c], 1);
+    }
+    if(tp_proc_set_image_staged(p, c, 1)){ printf("FAIL set_image_staged cam %d: %s\n", c, tp_last_error()); return 1; }
+  }
+  if(tp_proc_exec_conditioning(p, 0, 0.0)){ printf("FAIL exec_conditioning (float): %s\n", tp_last_error()); return 1; }
+  long bad2=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL get_image cam %d\n", c); return 1; }
+    bad2 += compare_bits(got, u16_result+(size_t)c*NPIX, NPIX);
+  }
+  printf("u16-vs-float path equivalence (preload contract): bit-mismatches=%ld -> %s\n",
+         bad2, bad2?"FAIL":"PASS");
+  bad_total += bad2;
+
+  // ---- 3. rowcol OFF, swap OFF (native-endian source, FLIR-corrected) ----
+  for(int c=0;c<NC;c++) for(size_t i=0;i<NPIX;i++)
+    u16[(size_t)c*NPIX+i]=value[(size_t)c*NPIX+i]; // native order this time
+  for(int c=0;c<NC;c++)
+    if(tp_proc_set_image_staged_u16(p, c, 0)){ printf("FAIL set_image_staged_u16(3) cam %d: %s\n", c, tp_last_error()); return 1; }
+  if(tp_proc_exec_conditioning_u16(p, 0, 0)){ printf("FAIL exec_conditioning_u16(3): %s\n", tp_last_error()); return 1; }
+  long bad3=0;
+  for(int c=0;c<NC;c++){
+    if(tp_proc_get_image(p, c, got)){ printf("FAIL get_image cam %d\n", c); return 1; }
+    for(size_t i=0;i<NPIX;i++){
+      const float rawf = (float) value[(size_t)c*NPIX+i];
+      expected[i] = fmaf(rawf, scale_maps[(size_t)c*NPIX+i], -offset_maps[(size_t)c*NPIX+i]);
+    }
+    bad3 += compare_bits(got, expected, NPIX);
+  }
+  printf("u16 no-swap/no-rowcol: bit-mismatches=%ld -> %s\n", bad3, bad3?"FAIL":"PASS");
+  bad_total += bad3;
+
+  printf("test_u16_cond_jna: %s\n", bad_total?"FAIL":"ALL PASS");
+  tp_proc_destroy(p);
+  tp_destroy_module(m);
+  return bad_total?1:0;
+}

src/tp_condition.cu

@@ -68,3 +68,44 @@ extern "C" __global__ void condition_lwir_d(
 	value = __dsub_rn(value, dc_level);
 	image[(size_t) y * image_stride + x] = (float) value;
 }
+
+// T2 16-bit feeder (see tp_condition.h). The double chain replicates
+// CuasConditioning.applyCalibratedRowColToRaw byte-for-byte:
+//   raw -= (row[x] + col[y]) * inv_scale       (three RN double ops, no FMA)
+// then ONE double->float cast (CuasScenePreload's fscene build) and the
+// unchanged production fmaf of condition_lwir_f. By Claude on 07/18/2026.
+extern "C" __global__ void condition_lwir_u16(
+		int                    num_cams,
+		int                    width,
+		int                    height,
+		size_t                 image_stride,
+		float               ** gpu_images,
+		const unsigned short * u16_src,
+		int                    swap_bytes,
+		int                    rowcol_enable,
+		const double         * rowcol,
+		const float          * scale_maps,
+		const float          * offset_maps)
+{
+	const int x = (int) (blockIdx.x * blockDim.x + threadIdx.x);
+	const int y = (int) (blockIdx.y * blockDim.y + threadIdx.y);
+	const int c = (int) blockIdx.z;
+	if ((c >= num_cams) || (x >= width) || (y >= height)) return;
+	const size_t pixel = (size_t) y * width + x;
+	const size_t map_index = (size_t) c * width * height + pixel;
+	unsigned short v = u16_src[map_index];
+	if (swap_bytes) v = (unsigned short) ((v >> 8) | (v << 8));
+	double raw = (double) v;
+	if (rowcol_enable) {
+		const double row = rowcol[(size_t) c * width + x];
+		const double col = rowcol[(size_t) num_cams * width + (size_t) c * height + y];
+		const double inv = rowcol[(size_t) num_cams * (width + height) + c];
+		raw = __dsub_rn(raw, __dmul_rn(__dadd_rn(row, col), inv));
+	}
+	const float raw_f = (float) raw;
+	float * image = gpu_images[c];
+	image[(size_t) y * image_stride + x] = fmaf(
+			raw_f,
+			scale_maps[map_index],
+			-offset_maps[map_index]);
+}

src/tp_condition.h

@@ -44,4 +44,26 @@ extern "C" __global__ void condition_lwir_d(
 		const double * fpn,
 		double         dc_level);
 
+/* T2 16-bit feeder (07/18/2026): convert the resident device uint16 scratch
+ * ([cam][height][width], raw file bytes) into the ACTIVE image set and apply
+ * the two-map conditioning in ONE pass. swap_bytes: source file is big-endian
+ * (TIFF 'MM'). rowcol_enable: subtract the borrowed per-scene Row/Col profiles
+ * exactly as the Java preload did on raw doubles - explicit round-to-nearest
+ * double ops (no FMA fusion), ONE double->float cast, then the identical
+ * production fmaf. rowcol layout: rows [cam][width], cols [cam][height],
+ * inv_scales [cam], one contiguous double block.
+ * By Claude on 07/18/2026, from Andrey's T2 design. */
+extern "C" __global__ void condition_lwir_u16(
+		int                    num_cams,
+		int                    width,
+		int                    height,
+		size_t                 image_stride,
+		float               ** gpu_images,
+		const unsigned short * u16_src,
+		int                    swap_bytes,
+		int                    rowcol_enable,
+		const double         * rowcol,     // rows | cols | inv_scales (may be NULL when disabled)
+		const float          * scale_maps,
+		const float          * offset_maps);
+
 #endif /* SRC_TP_CONDITION_H_ */