CLAUDE: item-4 R2 native - pixel-LoG + gated subtract-avg kernels on the...

CLAUDE: item-4 R2 native - pixel-LoG + gated subtract-avg kernels on the resident render frame + camera-upload fast path

render_pix_log: NaN-aware LINEAR LoG (convolve2DLReLU form, fmaf throughout
so a same-order CPU std::fma replay is bit-exact) fused with the ruling-3.1
POST-LoG subtract-average (diff = LoG - avg resident detection frame, gated
EMA refresh, NaN-avg bootstrap seed) on the 640x512 sensor window of r_rbg.
render_avg_accum: NaN-aware raw-domain sum/cnt -> the next run's prior-
average file. Entries tp_proc_render_log_setup (ALL allocs at arm - the
no-alloc-class-in-loop rule; re-arm resets = rerun==fresh) and
tp_proc_render_log_get; render_scene enqueues both after render_dp on the
render stream, before render_done. KERNELS 45.

Enqueue shave: tp_proc_render_cameras fast path - scene meta+ERS only
(NULL radial/rbr = keep resident static tables), pinned double-buffered
staging + async copies on the render stream; no stream-sync, no per-scene
rByRDist H2D.

New case log_pipe (registered): L0 log-armed render bit-identical; L1 GPU ==
CPU float replay BIT-exact (raw/diff/avg/accum, 4 frames) + double oracle
rel 1.8e-7 of input mag; NaN substitution exercised; L2 prior base + median
gate, both populations non-empty; L3 rerun==fresh bit-identical; L4 camera
fast-path == full upload bit-exact. render_pipe --bench grows --log (R2
armed): render-8 re-verified 4.365 ms/scene in 8+8 carve, +0.15 with R2.
Full suite ALL PASS; compute-sanitizer 12.8 memcheck 0 errors.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

cases.list.tmpl

@@ -73,6 +73,7 @@ case scene_dp_split  test_pose_scene_dp_split_jna ${MODEL}/testdata/pose_corr
 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 render_full   test_render_full_jna  ${MODEL}/testdata/pose_corr           --tol 0     # R0 bench (ROADMAP item 4): full-frame combined-render chain ms/scene (timing, PASS=sanity only; run by name, pin clocks: perfctl hold measure -- ./run_cases.sh render_full) # By Claude on 07/19/2026
 case render_pipe    test_render_pipe_jna  ${MODEL}/testdata/pose_corr           --tol 0     # item-4 R1a: async full-frame render chain bit-exact vs blocking path + concurrency stress (add --bench 50 for ms/scene) # By Claude on 07/19/2026
+case log_pipe       test_log_pipe_jna     ${MODEL}/testdata/pose_corr           --tol 0     # item-4 R2: pixel-LoG + gated subtract-avg bit-exact vs CPU float replay + rel-tol vs double oracle; rerun==fresh; camera fast-path == full upload # By Claude on 07/19/2026
 #case carve_combo    test_carve_combo_jna  ${MODEL}/testdata/pose_corr           --tol 0     # item-4 step-3: pose DP in its green-ctx carve + render chain in a disjoint carve, one thread; bit-identity + latency percentiles (--pose-green/--render-green/--scenes) # By Claude on 07/19/2026
 #case three_service  test_three_service_jna ${MODEL}/testdata/pose_corr          --tol 0     # item-4 MAP v2 concurrent measure: pose-8 + render-8 + LibTorch DNN thread in rest-20 (TPDNN_GREEN_SM); bit-identity + per-service latency/throughput (--dnn-green/--dnn-period/--no-dnn); needs jna/build_dnn.sh + the DNN models # By Claude on 07/19/2026
 #case convert_direct test_convert_direct   ${REPO}/testdata/convert_direct_legacy               # local blessed goldens (~200MB, make_convert_direct_case.py)

jna/test_log_pipe_jna.cu

+// test_log_pipe_jna.cu - imagej ROADMAP item 4 rung R2 verdict case: the GPU
+// pixel-LoG + post-LoG subtract-average stage (render_pix_log) and the raw
+// average accumulator (render_avg_accum) on the resident render frame, plus
+// the R2 enqueue-shave camera fast path (meta+ERS async, static radial/rbr
+// resident).
+//
+// Gates:
+//  L0  non-perturbation: arming the log stage leaves the rendered frame
+//      BIT-identical (the log kernels only read r_rbg).
+//  L1  LoG math: 4 distinct frames (task-center shifts); the GPU output must
+//      be BIT-exact vs a CPU float replay (same loop order, std::fma == GPU
+//      fmaf) for raw LoG / diff / avg / accumulators, and within RELATIVE
+//      tol vs a CPU double oracle (the Java convolve2DLReLU form, ruling 3.3).
+//      NaN coverage is asserted (NaN centers + valid pixels with NaN
+//      neighbors - the center-substitution path must actually run).
+//  L2  prior base + gated EMA: re-arm with avg_init and a data-derived gate;
+//      replay bit-exact; both populations (gated / updated) non-empty.
+//  L3  rerun == fresh (T1 rule): re-arm with L2 params, rerun the scene
+//      sequence, ALL outputs bit-identical to L2.
+//  L4  camera fast path: use_task_update=1 renders, per-scene meta+ERS via
+//      the async fast path == the full blocking upload, BIT-exact frames.
+//
+// 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_log_pipe_jna.cu src/tests/tp_test_data.cu \
+//     -o tests_bin/test_log_pipe_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_log_pipe_jna --data <model>/testdata/pose_corr --tol 0
+//
+//  Created on: Jul 19, 2026
+//      Author: Claude (Anthropic), for Elphel
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cmath>
+#include <string>
+#include <vector>
+#include <algorithm>
+
+#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_const(void*, const char*, const float*, int);
+  int   tp_proc_setup_rbg_corr(void*, int, int,int,int,int, float,float,float, int);
+  // render chain + R2 entries
+  int   tp_proc_render_setup(void*, float, float, float, float, float, float, float);
+  int   tp_proc_render_template(void*, const int*, const float*, int);
+  int   tp_proc_render_cameras(void*, const float*, const float*, const float*, int, const float*, int,
+                               const float*, const float*, const float*, int, const float*, int);
+  int   tp_proc_render_set_tasks(void*, int, const float*, int, int);
+  int   tp_proc_render_scene(void*, const float*, int, float,float,float,float, float,float, int,int, int, int);
+  int   tp_proc_render_status(void*, int, int*);
+  int   tp_proc_render_get(void*, float*);
+  int   tp_proc_render_log_setup(void*, const float*, int, const float*, int, float, float, int, int);
+  int   tp_proc_render_log_get(void*, int, float*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+// bit-exact (NaN==NaN) compare
+static long bitCompare(const float* a, const float* b, size_t n, double* maxd){
+  long bad = 0; *maxd = 0;
+  for (size_t i = 0; i < n; i++){
+    const bool an = std::isnan(a[i]), bn = std::isnan(b[i]);
+    if (an && bn) continue;
+    if (an != bn){ bad++; continue; }
+    const double d = std::fabs((double)a[i]-(double)b[i]);
+    if (d > 0){ bad++; if (d > *maxd) *maxd = d; }
+  }
+  return bad;
+}
+// relative-tol (NaN==NaN) compare vs a double oracle. The LoG of large LWIR
+// counts cancels to a small output, so float accumulation error scales with
+// the INPUT magnitude, not the output: normalize by max(mag, |o|) where mag =
+// max|input| over the window (ruling 3.3: sized to never false-positive; the
+// 49-tap float sum error bound is ~ntaps*eps*mag ~ 3e-6*mag).
+static long relCompare(const float* g, const double* o, size_t n, double tol, double mag, double* maxrel){
+  long bad = 0; *maxrel = 0;
+  for (size_t i = 0; i < n; i++){
+    const bool gn = std::isnan(g[i]), on = std::isnan(o[i]);
+    if (gn && on) continue;
+    if (gn != on){ bad++; continue; }
+    const double den = std::max(std::max(1.0, mag), std::fabs(o[i]));
+    const double r = std::fabs((double)g[i]-o[i]) / den;
+    if (r > *maxrel) *maxrel = r;
+    if (r > tol) bad++;
+  }
+  return bad;
+}
+
+// Java CuasRTUtils.getRadiusLoGKernel + getLoGKenel, verbatim port (double)
+static int logRadius(double psf_radius, double n_sigma){
+  const double sigma = psf_radius / std::sqrt(2.0);
+  return (int)std::ceil(sigma * n_sigma);
+}
+static std::vector<double> logKernel(double psf_radius, int krad){
+  const double sigma = psf_radius / std::sqrt(2.0);
+  const double sigma2 = sigma * sigma;
+  const int width = 2*krad + 1;
+  std::vector<double> kernel((size_t)width*width);
+  double sumPos = 0, sumNeg = 0;
+  for (int y = -krad; y <= krad; y++)
+    for (int x = -krad; x <= krad; x++){
+      const double r2 = (double)x*x + (double)y*y;
+      const double raw = (1.0 - r2/(2.0*sigma2)) * std::exp(-r2/(2.0*sigma2));
+      kernel[(size_t)(y+krad)*width + (x+krad)] = raw;
+      if (raw > 0) sumPos += raw; else sumNeg += raw;
+    }
+  const double absSumNeg = std::fabs(sumNeg);
+  for (auto& k : kernel){ if (k > 0) k /= sumPos; else if (k < 0) k /= absSumNeg; }
+  return kernel;
+}
+
+// CPU float replay of render_pix_log: SAME loop order, std::fma == GPU fmaf.
+// src = padded frame (stride x rows), window (x0,y0)+(w,h). Updates avg/diff/
+// raw in place exactly like the kernel (per-pixel independent).
+static void replayPixLog(const float* src, int stride, int x0, int y0, int w, int h,
+                         const float* taps, int krad,
+                         float* raw, float* avg, float* diff, float ema_k, float gate_abs){
+  const int ksize = 2*krad + 1;
+  for (int y = 0; y < h; y++)
+    for (int x = 0; x < w; x++){
+      const int idx = y*w + x;
+      const float c = src[(size_t)(y0+y)*stride + (x0+x)];
+      float L;
+      if (std::isnan(c)) L = std::nanf("");
+      else {
+        float v = 0.0f;
+        for (int iy = 0; iy < ksize; iy++){
+          const int yy = y + iy - krad;
+          if (yy < 0 || yy >= h) continue;
+          const float* srow = src + (size_t)(y0+yy)*stride + x0;
+          const float* trow = taps + iy*ksize;
+          for (int ix = 0; ix < ksize; ix++){
+            const int xx = x + ix - krad;
+            if (xx < 0 || xx >= w) continue;
+            const float d = srow[xx];
+            v = std::fma(std::isnan(d) ? c : d, trow[ix], v);
+          }
+        }
+        L = v;
+      }
+      if (raw) raw[idx] = L;
+      if (!avg){ diff[idx] = L; continue; }
+      const float A = avg[idx];
+      if (std::isnan(L)) diff[idx] = std::nanf("");
+      else if (std::isnan(A)){
+        if (ema_k > 0.0f){ avg[idx] = L; diff[idx] = 0.0f; }
+        else diff[idx] = std::nanf("");
+      } else {
+        const float d = L - A;
+        diff[idx] = d;
+        if ((ema_k > 0.0f) && ((gate_abs <= 0.0f) || (std::fabs(d) <= gate_abs)))
+          avg[idx] = std::fma(ema_k, d, A);
+      }
+    }
+}
+// CPU double oracle of the LINEAR convolve2DLReLU (Java loop shape, double sum)
+static void oraclePixLog(const float* src, int stride, int x0, int y0, int w, int h,
+                         const double* taps, int krad, double* out){
+  const int ksize = 2*krad + 1;
+  for (int y = 0; y < h; y++)
+    for (int x = 0; x < w; x++){
+      const double c = (double)src[(size_t)(y0+y)*stride + (x0+x)];
+      if (std::isnan(c)){ out[(size_t)y*w+x] = std::nan(""); continue; }
+      double v = 0;
+      for (int iy = 0; iy < ksize; iy++){
+        const int yy = y + iy - krad;
+        if (yy < 0 || yy >= h) continue;
+        for (int ix = 0; ix < ksize; ix++){
+          const int xx = x + ix - krad;
+          if (xx < 0 || xx >= w) continue;
+          const double d = (double)src[(size_t)(y0+yy)*stride + (x0+xx)];
+          v += (std::isnan(d) ? c : d) * taps[(size_t)iy*ksize + ix];
+        }
+      }
+      out[(size_t)y*w+x] = v;
+    }
+}
+
+int main(int argc, char** argv){
+  std::string data_dir = "testdata/pose_corr";
+  std::string srcdir   = "src";
+  std::string devrt    = "/usr/local/cuda-12.8/lib64/libcudadevrt.a";
+  for (int i = 1; i < argc; i++){
+    if      (!strcmp(argv[i],"--data")  && (i+1<argc)) data_dir = 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],"--tol")   && (i+1<argc)) i++; // fixed gates here; accepted for run_cases.sh uniformity
+    else { fprintf(stderr,"Usage: %s --data <pose_corr case> [--src <dir>] [--devrt <a>]\n", argv[0]); return EXIT_FAILURE; }
+  }
+  printf("%s: data=%s (NVRTC module from %s)\n", argv[0], data_dir.c_str(), srcdir.c_str());
+  TpTestData data(data_dir);
+
+  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 float soft_margin = (float)data.param("soft_margin", 12.0);
+  const float hard_margin = (float)data.param("hard_margin",  8.0);
+  const int kern_tiles = kernels_hor*kernels_vert*num_colors;
+  const int ntiles = tilesx*tilesy;
+
+  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; }
+  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_setup_rbg_corr(p, 1, 0,0,0,0, 1.f,1.f,1.f, 7);
+
+  // full-frame tasks with EXPLICIT centers (render_pipe pattern); per-scene
+  // center shift makes 4 distinct frames
+  const float mb_dx = 0.3125f, mb_dy = -0.171875f;
+  const float* tmpl0 = data.hostFloat("ftasks0_it0");
+  auto buildTasks = [&](int nset, float shift, std::vector<float>& out){
+    out.assign((size_t)ntiles*task_size, 0.f);
+    for (int ty = 0; ty < tilesy; ty++)
+      for (int tx = 0; tx < tilesx; tx++){
+        float* t = out.data() + (size_t)(ty*tilesx+tx)*task_size;
+        memcpy(t, tmpl0, 6*sizeof(float));
+        const int txy = (ty << 16) | tx;
+        memcpy(t+1, &txy, sizeof(int));
+        t[3] = tx*8 + 4 + shift + ((nset==1)? mb_dx : 0.f);
+        t[4] = ty*8 + 4 +         ((nset==1)? mb_dy : 0.f);
+      }
+  };
+  const int rbg_w = img_width + 8, rbg_h = num_colors*(img_height + 8);
+  const size_t rbg_n = (size_t)rbg_w*rbg_h;
+  const int x0 = 4, y0 = 4;                    // sensor window inside the padded rbg
+  const size_t npix = (size_t)img_width*img_height;
+  const int NSCENES = 4;
+  const float shifts[NSCENES] = {0.f, 0.25f, -0.375f, 0.5f};
+  int fail = 0;
+
+  if(tp_proc_render_setup(p, 0.f, 0.f,0.f,0.f,0.f, soft_margin, hard_margin)){
+    printf("FAIL render_setup: %s\n", tp_last_error()); return EXIT_FAILURE; }
+
+  auto renderScene = [&](int k, float* frame)->bool{
+    std::vector<float> t0, t1;
+    buildTasks(0, shifts[k], t0); buildTasks(1, shifts[k], t1);
+    if(tp_proc_render_set_tasks(p, 0, t0.data(), ntiles, ntiles*task_size) ||
+       tp_proc_render_set_tasks(p, 1, t1.data(), ntiles, ntiles*task_size)){
+      printf("FAIL render_set_tasks: %s\n", tp_last_error()); return false; }
+    if(tp_proc_render_scene(p, nullptr, 0, 0,0,0,0, 0,0, 0,0, -1, 0)){
+      printf("FAIL render_scene: %s\n", tp_last_error()); return false; }
+    if(tp_proc_render_status(p, 1, nullptr) != 1){
+      printf("FAIL render_status: %s\n", tp_last_error()); return false; }
+    if(frame && tp_proc_render_get(p, frame) != 1){
+      printf("FAIL render_get: %s\n", tp_last_error()); return false; }
+    return true;
+  };
+
+  // ---- taps: the production LoG (psf_radius=1.0, n_sigma=4.0 defaults) ----
+  const double psf_radius = 1.0, n_sigma = 4.0;
+  const int krad = logRadius(psf_radius, n_sigma);
+  const std::vector<double> dtaps = logKernel(psf_radius, krad);
+  std::vector<float> ftaps(dtaps.size());
+  for (size_t i = 0; i < dtaps.size(); i++) ftaps[i] = (float)dtaps[i];
+  printf("LoG taps: psf=%g n_sigma=%g -> krad=%d (%zux%zu)\n", psf_radius, n_sigma, krad,
+         (size_t)(2*krad+1), (size_t)(2*krad+1));
+
+  // ---- L0: non-perturbation - frame before vs after arming the log stage ----
+  std::vector<float> frame_plain(rbg_n), frame_logged(rbg_n);
+  if(!renderScene(0, frame_plain.data())) return EXIT_FAILURE;
+  if(tp_proc_render_log_setup(p, ftaps.data(), krad, nullptr, 0, 0.05f, 0.f, 1, 1)){
+    printf("FAIL render_log_setup: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  if(!renderScene(0, frame_logged.data())) return EXIT_FAILURE;
+  { double maxd=0; const long bad = bitCompare(frame_logged.data(), frame_plain.data(), rbg_n, &maxd);
+    printf("L0 render frame with log armed vs plain: mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS");
+    if (bad) fail = 1; }
+
+  // ---- L1: LoG math over 4 frames, ungated EMA from NaN seed + accumulation ----
+  // re-arm to reset avg/accum (the L0 scene already ran through the stage)
+  if(tp_proc_render_log_setup(p, ftaps.data(), krad, nullptr, 0, 0.05f, 0.f, 1, 1)){
+    printf("FAIL re-arm: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  std::vector<float> rep_avg(npix, std::nanf("")), rep_raw(npix), rep_diff(npix);
+  std::vector<float> rep_sum(npix, 0.f), rep_cnt(npix, 0.f);
+  std::vector<float> g_raw(npix), g_diff(npix), g_avg(npix), g_sum(npix), g_cnt(npix);
+  std::vector<float> frame(rbg_n);
+  std::vector<double> dorc(npix);
+  for (int k = 0; k < NSCENES && !fail; k++){
+    if(!renderScene(k, frame.data())) return EXIT_FAILURE;
+    // replay accumulator BEFORE avg/diff (kernel order irrelevant - per-pixel independent)
+    for (size_t i = 0; i < npix; i++){
+      const float c = frame[(size_t)(y0 + i/img_width)*rbg_w + x0 + (i%img_width)];
+      if (!std::isnan(c)){ rep_sum[i] += c; rep_cnt[i] += 1.f; }
+    }
+    replayPixLog(frame.data(), rbg_w, x0, y0, img_width, img_height,
+                 ftaps.data(), krad, rep_raw.data(), rep_avg.data(), rep_diff.data(), 0.05f, 0.f);
+    if(tp_proc_render_log_get(p, 2, g_raw.data())  != 1 ||
+       tp_proc_render_log_get(p, 0, g_diff.data()) != 1 ||
+       tp_proc_render_log_get(p, 1, g_avg.data())  != 1){
+      printf("FAIL log_get scene %d: %s\n", k, tp_last_error()); return EXIT_FAILURE; }
+    double maxd=0; long bad;
+    bad = bitCompare(g_raw.data(),  rep_raw.data(),  npix, &maxd);
+    printf("L1 scene %d raw  LoG vs float replay: mismatches=%ld -> %s\n", k, bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+    bad = bitCompare(g_diff.data(), rep_diff.data(), npix, &maxd);
+    printf("L1 scene %d diff     vs float replay: mismatches=%ld -> %s\n", k, bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+    bad = bitCompare(g_avg.data(),  rep_avg.data(),  npix, &maxd);
+    printf("L1 scene %d avg      vs float replay: mismatches=%ld -> %s\n", k, bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+    // double oracle, raw LoG only (diff/avg inherit): relative tol (ruling 3.3)
+    oraclePixLog(frame.data(), rbg_w, x0, y0, img_width, img_height, dtaps.data(), krad, dorc.data());
+    double mag = 0;
+    for (int yy = 0; yy < img_height; yy++)
+      for (int xx = 0; xx < img_width; xx++){
+        const float c = frame[(size_t)(y0+yy)*rbg_w + x0+xx];
+        if (!std::isnan(c) && std::fabs((double)c) > mag) mag = std::fabs((double)c);
+      }
+    double maxrel=0;
+    bad = relCompare(g_raw.data(), dorc.data(), npix, 1e-5, mag, &maxrel);
+    printf("L1 scene %d raw LoG vs double oracle: bad=%ld maxrel=%.3g (mag=%.3g) -> %s\n",
+           k, bad, maxrel, mag, bad?"FAIL":"PASS");
+    if (bad) fail=1;
+    if (k == 0){ // NaN coverage: the substitution path must actually run
+      long nnan=0, nvalid=0, nsub=0;
+      for (int yy = 0; yy < img_height; yy++)
+        for (int xx = 0; xx < img_width; xx++){
+          const float c = frame[(size_t)(y0+yy)*rbg_w + x0+xx];
+          if (std::isnan(c)){ nnan++; continue; }
+          nvalid++;
+          bool near_nan = false;
+          for (int dy = -krad; dy <= krad && !near_nan; dy++){
+            const int y2 = yy+dy; if (y2 < 0 || y2 >= img_height) continue;
+            for (int dx = -krad; dx <= krad; dx++){
+              const int x2 = xx+dx; if (x2 < 0 || x2 >= img_width) continue;
+              if (std::isnan(frame[(size_t)(y0+y2)*rbg_w + x0+x2])){ near_nan = true; break; }
+            }
+          }
+          if (near_nan) nsub++;
+        }
+      printf("L1 NaN coverage: nan=%ld valid=%ld valid-with-NaN-neighbor=%ld -> %s\n",
+             nnan, nvalid, nsub, (nnan>0 && nvalid>0 && nsub>0)?"PASS":"FAIL");
+      if (!(nnan>0 && nvalid>0 && nsub>0)) fail = 1;
+    }
+  }
+  if(tp_proc_render_log_get(p, 3, g_sum.data()) != 1 ||
+     tp_proc_render_log_get(p, 4, g_cnt.data()) != 1){
+    printf("FAIL accum get: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  { double maxd=0;
+    long bad = bitCompare(g_sum.data(), rep_sum.data(), npix, &maxd);
+    printf("L1 accum sum vs float replay: mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+    bad = bitCompare(g_cnt.data(), rep_cnt.data(), npix, &maxd);
+    printf("L1 accum cnt vs float replay: mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS"); if (bad) fail=1; }
+
+  // ---- L2: prior-run base + gated EMA (gate derived from L1 data) ----
+  // base = float replay LoG of frame 0; gate = median |L_k - base| of scene 1
+  std::vector<float> base(npix);
+  {
+    if(!renderScene(0, frame.data())) return EXIT_FAILURE;
+    replayPixLog(frame.data(), rbg_w, x0, y0, img_width, img_height,
+                 ftaps.data(), krad, base.data(), nullptr, rep_diff.data(), 0.f, 0.f);
+    if(!renderScene(1, frame.data())) return EXIT_FAILURE;
+    replayPixLog(frame.data(), rbg_w, x0, y0, img_width, img_height,
+                 ftaps.data(), krad, rep_raw.data(), nullptr, rep_diff.data(), 0.f, 0.f);
+  }
+  std::vector<float> absd;
+  absd.reserve(npix);
+  for (size_t i = 0; i < npix; i++)
+    if (!std::isnan(rep_raw[i]) && !std::isnan(base[i])) absd.push_back(std::fabs(rep_raw[i]-base[i]));
+  if (absd.empty()){ printf("L2 FAIL: no valid diff population\n"); return EXIT_FAILURE; }
+  std::nth_element(absd.begin(), absd.begin()+absd.size()/2, absd.end());
+  const float gate = absd[absd.size()/2];
+  const float ema2 = 0.25f;
+  printf("L2 gate = median |LoG - base| of scene 1 = %g\n", (double)gate);
+  std::vector<float> l2_diff[NSCENES], l2_avg_final(npix), l2_sum(npix), l2_cnt(npix);
+  auto runGated = [&](std::vector<float>* diffs, float* avg_out, float* sum_out, float* cnt_out)->bool{
+    if(tp_proc_render_log_setup(p, ftaps.data(), krad, base.data(), (int)npix, ema2, gate, 1, 1)){
+      printf("FAIL gated log_setup: %s\n", tp_last_error()); return false; }
+    std::fill(rep_sum.begin(), rep_sum.end(), 0.f);
+    std::fill(rep_cnt.begin(), rep_cnt.end(), 0.f);
+    std::vector<float> ravg(base);           // replay average starts at the base
+    long gated_total = 0, updated_total = 0;
+    for (int k = 0; k < NSCENES; k++){
+      if(!renderScene(k, frame.data())) return false;
+      for (size_t i = 0; i < npix; i++){
+        const float c = frame[(size_t)(y0 + i/img_width)*rbg_w + x0 + (i%img_width)];
+        if (!std::isnan(c)){ rep_sum[i] += c; rep_cnt[i] += 1.f; }
+      }
+      std::vector<float> ravg_before(ravg);
+      replayPixLog(frame.data(), rbg_w, x0, y0, img_width, img_height,
+                   ftaps.data(), krad, rep_raw.data(), ravg.data(), rep_diff.data(), ema2, gate);
+      for (size_t i = 0; i < npix; i++){
+        if (std::isnan(rep_diff[i]) || rep_diff[i] == 0.f) continue;
+        if (ravg[i] == ravg_before[i]) gated_total++; else updated_total++;
+      }
+      if(tp_proc_render_log_get(p, 0, g_diff.data()) != 1 ||
+         tp_proc_render_log_get(p, 1, g_avg.data())  != 1){
+        printf("FAIL gated log_get: %s\n", tp_last_error()); return false; }
+      double maxd=0;
+      long bad = bitCompare(g_diff.data(), rep_diff.data(), npix, &maxd);
+      bad += bitCompare(g_avg.data(), ravg.data(), npix, &maxd);
+      printf("L2 scene %d gated diff+avg vs float replay: mismatches=%ld -> %s\n", k, bad, bad?"FAIL":"PASS");
+      if (bad) return false;
+      if (diffs) diffs[k] = g_diff;
+    }
+    printf("L2 populations: gated(kept)=%ld updated=%ld -> %s\n", gated_total, updated_total,
+           (gated_total>0 && updated_total>0)?"PASS":"FAIL");
+    if (!(gated_total>0 && updated_total>0)) return false;
+    if(avg_out && tp_proc_render_log_get(p, 1, avg_out) != 1) return false;
+    if(sum_out && tp_proc_render_log_get(p, 3, sum_out) != 1) return false;
+    if(cnt_out && tp_proc_render_log_get(p, 4, cnt_out) != 1) return false;
+    return true;
+  };
+  if(!runGated(l2_diff, l2_avg_final.data(), l2_sum.data(), l2_cnt.data())){ fail = 1; }
+
+  // ---- L3: rerun == fresh (re-arm resets avg + accumulators, T1 rule) ----
+  if (!fail){
+    std::vector<float> l3_diff[NSCENES], l3_avg(npix), l3_sum(npix), l3_cnt(npix);
+    if(!runGated(l3_diff, l3_avg.data(), l3_sum.data(), l3_cnt.data())){ fail = 1; }
+    if (!fail){
+      double maxd=0; long bad = 0;
+      for (int k = 0; k < NSCENES; k++)
+        bad += bitCompare(l3_diff[k].data(), l2_diff[k].data(), npix, &maxd);
+      bad += bitCompare(l3_avg.data(), l2_avg_final.data(), npix, &maxd);
+      bad += bitCompare(l3_sum.data(), l2_sum.data(), npix, &maxd);
+      bad += bitCompare(l3_cnt.data(), l2_cnt.data(), npix, &maxd);
+      printf("L3 rerun-vs-fresh (all diffs + final avg + accum): mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS");
+      if (bad) fail = 1;
+    }
+  }
+
+  // ---- L4: camera fast path (meta+ERS async) == full blocking upload ----
+  // use_task_update=1 with synthetic-but-real-dims cameras (scene_dp precedent);
+  // meta px0/py0 shifts change the projected centers, so the fast path carries
+  // real per-scene content.
+  {
+    const int rbr_len = 5001, ers_len = img_height*14;
+    const float radial[7] = {0,0,0,0,0,0,0};
+    std::vector<float> rbr(rbr_len, 1.0f), ers((size_t)ers_len, 0.0f);
+    for (int line = 0; line < img_height; line++) ers[(size_t)line*14+6] = 1.0f;
+    std::vector<int>   task_map(ntiles);
+    std::vector<float> task_centers((size_t)3*ntiles);
+    for (int ty = 0; ty < tilesy; ty++)
+      for (int tx = 0; tx < tilesx; tx++){
+        const int i = ty*tilesx + tx;
+        task_map[i] = tx | (ty << 16);
+        task_centers[3*(size_t)i+0] = tx*8 + 4;
+        task_centers[3*(size_t)i+1] = ty*8 + 4;
+        task_centers[3*(size_t)i+2] = 0.0f;
+      }
+    // re-arm with REAL pixel bounds - pose_task_update culls by [min_px..max_px]
+    // (the L1/L2 setup used zeros: fine for use_task_update=0, fatal here)
+    if(tp_proc_render_setup(p, 0.f, 8.f, (float)(img_width-9), 8.f, (float)(img_height-9),
+                            soft_margin, hard_margin)){
+      printf("FAIL utu render_setup: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    if(tp_proc_render_template(p, task_map.data(), task_centers.data(), ntiles)){
+      printf("FAIL render_template: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    const int task_code = data.paramInt("task_code_main", 0x203); // scene_dp convention
+    const float pose[12] = {0,0,0, 0,0,0, 0,0,0, 0.0005f,-0.0003f,0.0001f};
+    auto metaFor = [&](int k, float* mv){
+      mv[0]=(float)img_width; mv[1]=(float)img_height;
+      mv[2]=4.0f + 0.5f*k; mv[3]=5.0f - 0.25f*k; mv[4]=2.85f; mv[5]=1.0f; mv[6]=0.0f; mv[7]=0.0004f;
+    };
+    auto renderUtu = [&](float* frame)->bool{
+      if(tp_proc_render_scene(p, pose, 12, 0.1f,-0.05f,0,0, 1.f,0.f, task_code, 511, -1, 1)){
+        printf("FAIL utu render_scene: %s\n", tp_last_error()); return false; }
+      if(tp_proc_render_status(p, 1, nullptr) != 1){
+        printf("FAIL utu render_status: %s\n", tp_last_error()); return false; }
+      return tp_proc_render_get(p, frame) == 1;
+    };
+    // reference camera once (full upload)
+    float meta_ref[8]; metaFor(0, meta_ref);
+    if(tp_proc_render_cameras(p, meta_ref, radial, rbr.data(), rbr_len, ers.data(), ers_len,
+                              nullptr, nullptr, nullptr, 0, nullptr, 0)){
+      printf("FAIL ref cameras: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    // series A: FULL scene upload per scene
+    std::vector<float> A0(rbg_n), A1(rbg_n), B0(rbg_n), B1(rbg_n);
+    float meta_s[8];
+    metaFor(1, meta_s);
+    if(tp_proc_render_cameras(p, nullptr,nullptr,nullptr,0,nullptr,0,
+                              meta_s, radial, rbr.data(), rbr_len, ers.data(), ers_len)){
+      printf("FAIL full scene cameras 0: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    if(!renderUtu(A0.data())) return EXIT_FAILURE;
+    metaFor(2, meta_s);
+    if(tp_proc_render_cameras(p, nullptr,nullptr,nullptr,0,nullptr,0,
+                              meta_s, radial, rbr.data(), rbr_len, ers.data(), ers_len)){
+      printf("FAIL full scene cameras 1: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    if(!renderUtu(A1.data())) return EXIT_FAILURE;
+    { double maxd=0; const long same = bitCompare(A0.data(), A1.data(), rbg_n, &maxd);
+      printf("L4 sanity: meta shift changes the frame (diff pixels=%ld) -> %s\n", same, same>0?"PASS":"FAIL");
+      if (!same) fail = 1; }
+    // series B: FAST path per scene (radial/rbr NULL = resident)
+    metaFor(1, meta_s);
+    if(tp_proc_render_cameras(p, nullptr,nullptr,nullptr,0,nullptr,0,
+                              meta_s, nullptr, nullptr, 0, ers.data(), ers_len)){
+      printf("FAIL fast cameras 0: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    if(!renderUtu(B0.data())) return EXIT_FAILURE;
+    metaFor(2, meta_s);
+    if(tp_proc_render_cameras(p, nullptr,nullptr,nullptr,0,nullptr,0,
+                              meta_s, nullptr, nullptr, 0, ers.data(), ers_len)){
+      printf("FAIL fast cameras 1: %s\n", tp_last_error()); return EXIT_FAILURE; }
+    if(!renderUtu(B1.data())) return EXIT_FAILURE;
+    double maxd=0;
+    long bad = bitCompare(B0.data(), A0.data(), rbg_n, &maxd);
+    printf("L4 fast-vs-full cameras scene 0: mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+    bad = bitCompare(B1.data(), A1.data(), rbg_n, &maxd);
+    printf("L4 fast-vs-full cameras scene 1: mismatches=%ld -> %s\n", bad, bad?"FAIL":"PASS"); if (bad) fail=1;
+  }
+
+  printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+  tp_proc_destroy(p); tp_destroy_module(m);
+  return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

jna/test_render_pipe_jna.cu

@@ -72,6 +72,7 @@ extern "C" {
   int   tp_proc_render_scene(void*, const float*, int, float,float,float,float, float,float, int,int, int, int);
   int   tp_proc_render_status(void*, int, int*);
   int   tp_proc_render_get(void*, float*);
+  int   tp_proc_render_log_setup(void*, const float*, int, const float*, int, float, float, int, int); // item-4 R2 (--log bench) // By Claude on 07/19/2026
   void  tp_proc_destroy(void*);
   void  tp_destroy_module(void*);
 }
@@ -93,14 +94,15 @@ int main(int argc, char** argv){
   std::string data_dir = "testdata/pose_corr";
   std::string srcdir   = "src";
   std::string devrt    = "/usr/local/cuda-12.8/lib64/libcudadevrt.a";
-  int bench = 0;
+  int bench = 0, with_log = 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],"--bench") && (i+1<argc)) bench = atoi(argv[++i]);
+    else if (!strcmp(argv[i],"--log")) with_log = 1;   // item-4 R2: arm LoG+subtract-avg for the bench // By Claude on 07/19/2026
     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],"--tol")   && (i+1<argc)) i++;   // tol is fixed 0 here; accepted for run_cases.sh uniformity
-    else { fprintf(stderr,"Usage: %s --data <pose_corr case> [--src <dir>] [--devrt <a>]\n", argv[0]); return EXIT_FAILURE; }
+    else { fprintf(stderr,"Usage: %s --data <pose_corr case> [--src <dir>] [--devrt <a>] [--bench N [--log]]\n", argv[0]); return EXIT_FAILURE; }
   }
   printf("%s: data=%s (NVRTC module from %s)\n", argv[0], data_dir.c_str(), srcdir.c_str());
   TpTestData data(data_dir);
@@ -239,6 +241,19 @@ int main(int argc, char** argv){
 
   // ---- optional bench: the production async chain, launch -> completion ----
   if (bench > 0){
+    if (with_log){ // item-4 R2: production-shaped taps (psf=1, n_sigma=4 -> krad 3), EMA on, accumulator on // By Claude on 07/19/2026
+      const int krad = 3, ksize = 2*krad+1;
+      std::vector<float> taps((size_t)ksize*ksize);
+      { const double sigma = 1.0/std::sqrt(2.0), s2 = sigma*sigma; double sp=0, sn=0;
+        std::vector<double> d((size_t)ksize*ksize);
+        for (int y=-krad;y<=krad;y++) for (int x=-krad;x<=krad;x++){
+          const double r2=(double)x*x+y*y, raw=(1.0-r2/(2*s2))*std::exp(-r2/(2*s2));
+          d[(size_t)(y+krad)*ksize+(x+krad)]=raw; if (raw>0) sp+=raw; else sn+=raw; }
+        for (size_t i=0;i<d.size();i++) taps[i]=(float)(d[i]>0?d[i]/sp:d[i]/-sn); }
+      if(tp_proc_render_log_setup(p, taps.data(), krad, nullptr, 0, 0.02f, 10.f, 1, 0)){
+        printf("FAIL bench log_setup: %s\n", tp_last_error()); return EXIT_FAILURE; }
+      printf("bench: LoG + subtract-avg + accumulator ARMED (item-4 R2)\n");
+    }
     for (int i = 0; i < 5; i++){ // warmup
       tp_proc_render_scene(p, nullptr, 0, 0,0,0,0, 0,0, 0,0, -1, 0);
       tp_proc_render_status(p, 1, nullptr);

jna/tp_jna.cpp

@@ -53,6 +53,7 @@ static const char* KERNELS[] = {
   "calculate_tiles_offsets","calc_rot_deriv","calcReverseDistortionTable",
   "condition_lwir_f","condition_lwir_d", // production float maps + Java-matched validation // By Codex on 07/15/2026
   "condition_lwir_u16", // T2 16-bit feeder: u16->float + Row/Col + two-map, one pass // By Claude on 07/18/2026
+  "render_pix_log","render_avg_accum", // item-4 R2: pixel LoG + post-LoG subtract-avg / raw-avg accumulation // By Claude on 07/19/2026
   "clt_average_sensors","index_consolidate","consolidate_oob","clt_average_sensors_list", // By Claude on 07/12/2026
   "corr2D_peak_eig","corr2D_peak_eig_f", // double + float instantiations (precision two-step) // By Claude on 07/13-14/2026
   "pose_fx_jacobian", // one float thread per selected tile // By Codex on 07/15/2026
@@ -901,6 +902,20 @@ struct TpProc {
   // guard see the in-flight render. Same-device contexts share the VA space
   // (probe2k + carve_combo --bind gates). By Claude on 07/19/2026.
   TpProc *r_src, *r_client;
+  // item-4 R2 (07/19/2026): pixel-LoG + post-LoG subtract-average stage on the
+  // resident render frame, enqueued on the render stream after render_dp. ALL
+  // buffers allocated at tp_proc_render_log_setup (arm time - the standing
+  // no-alloc-class-in-loop rule). Sensor-window (img_w x img_h) packed layout.
+  bool have_render_log; int log_krad; float log_ema, log_gate;
+  int log_accum, log_want_raw;
+  float *r_log_taps, *r_log_raw, *r_log_avg, *r_log_diff;
+  float *r_acc_sum, *r_acc_cnt;
+  // R2 enqueue shave: per-scene camera update = meta+ERS only (static
+  // radial/rByRDist stay resident after the first full upload), packed into
+  // pinned double-buffered staging and copied ASYNC on the render stream (the
+  // old path stream-synced + issued 4 blocking memcpys per scene).
+  float *r_cam_staging[2]; cudaEvent_t r_cam_staged[2]; int r_cam_slot;
+  int r_cam_staging_floats; int r_rbr_len[2], r_ers_len[2];
 };
 
 extern "C" {
@@ -986,6 +1001,13 @@ TpProc* tp_proc_create(TpModule* m){
   p->r_imclt_clt=p->r_imclt_img=nullptr; p->r_status=nullptr;
   p->gpu_render_dp=nullptr; p->render_staging=nullptr; p->r_vec_staging=nullptr;
   p->r_src=nullptr; p->r_client=nullptr; // By Claude on 07/19/2026 (item-4 step-4 carved-render bind)
+  p->have_render_log=false; p->log_krad=0; p->log_ema=0.0f; p->log_gate=0.0f; // By Claude on 07/19/2026 (item-4 R2)
+  p->log_accum=0; p->log_want_raw=0;
+  p->r_log_taps=p->r_log_raw=p->r_log_avg=p->r_log_diff=nullptr;
+  p->r_acc_sum=p->r_acc_cnt=nullptr;
+  p->r_cam_staging[0]=p->r_cam_staging[1]=nullptr;
+  p->r_cam_staged[0]=p->r_cam_staged[1]=nullptr; p->r_cam_slot=0;
+  p->r_cam_staging_floats=0; p->r_rbr_len[0]=p->r_rbr_len[1]=0; p->r_ers_len[0]=p->r_ers_len[1]=0;
   return p;
 }
 
@@ -1685,6 +1707,16 @@ int tp_proc_render_setup(TpProc* p,
        cudaMalloc((void**)&p->r_radial[0],POSE_RADIAL_FLOATS*sizeof(float))!=cudaSuccess ||
        cudaMalloc((void**)&p->r_radial[1],POSE_RADIAL_FLOATS*sizeof(float))!=cudaSuccess){
       seterr("render_setup: camera alloc failed"); return -3; }
+    // R2 enqueue shave: pinned double-buffered staging for the per-scene
+    // meta+ERS async upload (arm-time alloc; capacity covers any plausible
+    // camera height incl. telemetry rows). By Claude on 07/19/2026.
+    p->r_cam_staging_floats=(int)(sizeof(PoseCameraMeta)/sizeof(float))+(p->img_h+64)*POSE_ERS_LINE_STRIDE;
+    if(cudaMallocHost((void**)&p->r_cam_staging[0],(size_t)p->r_cam_staging_floats*sizeof(float))!=cudaSuccess ||
+       cudaMallocHost((void**)&p->r_cam_staging[1],(size_t)p->r_cam_staging_floats*sizeof(float))!=cudaSuccess ||
+       cudaEventCreateWithFlags(&p->r_cam_staged[0],cudaEventDisableTiming)!=cudaSuccess ||
+       cudaEventCreateWithFlags(&p->r_cam_staged[1],cudaEventDisableTiming)!=cudaSuccess){
+      seterr("render_setup: camera staging alloc failed"); return -3; }
+    p->r_cam_slot=0;
     p->render_ntiles=ntiles; p->render_task_size=ts; p->render_ntasks=ntiles;
     p->have_render=true;
   }
@@ -1748,6 +1780,44 @@ int tp_proc_render_cameras(TpProc* p,
   if(!p||!p->have_render){ seterr("render_cameras: not armed"); return -1; }
   const bool up_ref   = (ref_meta!=nullptr);
   const bool up_scene = (scene_meta!=nullptr);
+  // R2 enqueue shave (07/19/2026): per-scene fast path - scene meta+ERS only,
+  // NULL radial/rbr = "keep the resident static tables" (identical every scene
+  // of a sequence: same camera calibration). Packed into pinned staging and
+  // copied ASYNC on the render stream: stream order already serializes against
+  // the previous render_dp (no sync needed) and precedes the next render_scene
+  // launch. The old path stream-synced + issued 4 blocking memcpys per scene
+  // (~0.4-0.5 ms enqueue). Full uploads (first scene / ref group) keep the
+  // original blocking arm-time behavior below. By Claude on 07/19/2026.
+  if(up_scene&&!up_ref&&!scene_radial&&!scene_rbr){
+    if(!(p->r_cameras_ready&2)){ seterr("render_cameras: static scene tables not resident (full upload first)"); return -2; }
+    if(!scene_ers){ seterr("render_cameras: fast path needs ERS"); return -1; }
+    const int width=(int)std::lround(scene_meta[0]);
+    const int height=(int)std::lround(scene_meta[1]);
+    if((width<=0)||(height<=0)||(scene_ers_len!=height*POSE_ERS_LINE_STRIDE)||
+       !std::isfinite(scene_meta[7])||!(scene_meta[7]>0.0f)){
+      seterr("render_cameras: scene dimensions/ERS/step mismatch (fast path)"); return -3; }
+    if(scene_ers_len>p->r_ers_capacity[1]){
+      seterr("render_cameras: ERS grew %d > resident %d - full upload required",scene_ers_len,p->r_ers_capacity[1]); return -4; }
+    const int meta_floats=(int)(sizeof(PoseCameraMeta)/sizeof(float));
+    if(meta_floats+scene_ers_len>p->r_cam_staging_floats){
+      seterr("render_cameras: staging too small (%d > %d)",meta_floats+scene_ers_len,p->r_cam_staging_floats); return -4; }
+    cuCtxSetCurrent(p->mod->ctx);
+    const int slot=p->r_cam_slot; p->r_cam_slot^=1;
+    // a never-recorded event reports complete - first use passes straight through
+    if(p->r_cam_staged[slot]) cudaEventSynchronize(p->r_cam_staged[slot]);
+    PoseCameraMeta* hm=(PoseCameraMeta*)p->r_cam_staging[slot];
+    *hm=PoseCameraMeta{width,height,p->r_rbr_len[1],
+      scene_meta[2],scene_meta[3],scene_meta[4],scene_meta[5],scene_meta[6],scene_meta[7]};
+    float* ers_stage=p->r_cam_staging[slot]+meta_floats;
+    memcpy(ers_stage,scene_ers,(size_t)scene_ers_len*sizeof(float));
+    if(cudaMemcpyAsync(p->r_cameras+1,hm,sizeof(PoseCameraMeta),cudaMemcpyHostToDevice,p->render_stream)!=cudaSuccess ||
+       cudaMemcpyAsync(p->r_ers[1],ers_stage,(size_t)scene_ers_len*sizeof(float),cudaMemcpyHostToDevice,p->render_stream)!=cudaSuccess){
+      seterr("render_cameras: async meta/ERS H2D failed"); return -5; }
+    if(cudaEventRecord(p->r_cam_staged[slot],p->render_stream)!=cudaSuccess){
+      seterr("render_cameras: staging event record failed"); return -5; }
+    p->r_ers_len[1]=scene_ers_len;
+    return 0;
+  }
   if(up_ref&&(!ref_radial||!ref_rbr||!ref_ers||(ref_rbr_len<3))){ seterr("render_cameras: bad ref group"); return -1; }
   if(up_scene&&(!scene_radial||!scene_rbr||!scene_ers||(scene_rbr_len<3))){ seterr("render_cameras: bad scene group"); return -1; }
   cuCtxSetCurrent(p->mod->ctx);
@@ -1782,11 +1852,78 @@ int tp_proc_render_cameras(TpProc* p,
        cudaMemcpy(p->r_rbr[cam],rbrs[cam],(size_t)rbr_lens[cam]*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
        cudaMemcpy(p->r_ers[cam],erss[cam],(size_t)ers_lens[cam]*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
       seterr("render_cameras: HtoD failed cam %d",cam); return -5; }
+    p->r_rbr_len[cam]=rbr_lens[cam]; p->r_ers_len[cam]=ers_lens[cam]; // fast-path bookkeeping // By Claude on 07/19/2026
     p->r_cameras_ready|=(1<<cam);
   }
   return 0;
 }
 
+// item-4 R2 (07/19/2026): arm the pixel-LoG + post-LoG subtract-average stage.
+// ALL alloc-class work happens HERE (the standing no-alloc-in-loop rule); the
+// per-scene kernels are enqueued by render_scene on the render stream after
+// render_dp. taps = (2*krad+1)^2 LoG taps (Java getLoGKenel, double->float);
+// avg_init = post-LoG-domain base (the prior-run average file LoG'd in Java
+// double), avg_len = img_w*img_h, or NULL/0 = start NaN (EMA bootstrap seeds
+// from the first frame; with ema_k=0 the diff stays NaN = prior-file-only).
+// Re-call per sequence re-uploads taps and RESETS avg + accumulators (rerun ==
+// fresh launch, the T1 rule). By Claude on 07/19/2026.
+int tp_proc_render_log_setup(TpProc* p, const float* taps, int krad,
+                             const float* avg_init, int avg_len,
+                             float ema_k, float gate_abs,
+                             int accum_enable, int want_raw){
+  if(!p||!p->have_render){ seterr("render_log_setup: render not armed (tp_proc_render_setup first)"); return -1; }
+  if(!taps||(krad<0)||(krad>15)){ seterr("render_log_setup: bad taps/krad %d",krad); return -1; }
+  const int npix=p->img_w*p->img_h;
+  if(avg_init&&(avg_len!=npix)){ seterr("render_log_setup: avg_len %d != %d",avg_len,npix); return -2; }
+  cuCtxSetCurrent(p->mod->ctx);
+  const int ntaps=(2*krad+1)*(2*krad+1);
+  if(p->have_render_log&&(krad!=p->log_krad)&&p->r_log_taps){ cudaFree(p->r_log_taps); p->r_log_taps=nullptr; }
+  if(!p->r_log_taps&&cudaMalloc((void**)&p->r_log_taps,(size_t)ntaps*sizeof(float))!=cudaSuccess){
+    seterr("render_log_setup: taps alloc failed"); return -3; }
+  if(!p->r_log_diff&&cudaMalloc((void**)&p->r_log_diff,(size_t)npix*sizeof(float))!=cudaSuccess){
+    seterr("render_log_setup: diff alloc failed"); return -3; }
+  if(!p->r_log_avg&&cudaMalloc((void**)&p->r_log_avg,(size_t)npix*sizeof(float))!=cudaSuccess){
+    seterr("render_log_setup: avg alloc failed"); return -3; }
+  if(want_raw&&!p->r_log_raw&&cudaMalloc((void**)&p->r_log_raw,(size_t)npix*sizeof(float))!=cudaSuccess){
+    seterr("render_log_setup: raw alloc failed"); return -3; }
+  if(accum_enable&&!p->r_acc_sum){
+    if(cudaMalloc((void**)&p->r_acc_sum,(size_t)npix*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->r_acc_cnt,(size_t)npix*sizeof(float))!=cudaSuccess){
+      seterr("render_log_setup: accum alloc failed"); return -3; } }
+  if(cudaMemcpy(p->r_log_taps,taps,(size_t)ntaps*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("render_log_setup: taps H2D failed"); return -4; }
+  if(avg_init){
+    if(cudaMemcpy(p->r_log_avg,avg_init,(size_t)npix*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+      seterr("render_log_setup: avg H2D failed"); return -4; }
+  } else {
+    if(cuMemsetD32((CUdeviceptr)p->r_log_avg,0x7fc00000u,(size_t)npix)!=CUDA_SUCCESS){ // quiet-NaN fill
+      seterr("render_log_setup: avg NaN fill failed"); return -4; } }
+  if(p->r_acc_sum){
+    if(cudaMemset(p->r_acc_sum,0,(size_t)npix*sizeof(float))!=cudaSuccess ||
+       cudaMemset(p->r_acc_cnt,0,(size_t)npix*sizeof(float))!=cudaSuccess){
+      seterr("render_log_setup: accum reset failed"); return -4; } }
+  p->log_krad=krad; p->log_ema=ema_k; p->log_gate=gate_abs;
+  p->log_accum=accum_enable?1:0; p->log_want_raw=want_raw?1:0;
+  p->have_render_log=true;
+  return 0;
+}
+
+// R2 D2H (debug/eval + the end-of-run average save; ruling 3.4: never
+// per-scene in production). which: 0=diff (the detection frame), 1=avg
+// (post-LoG), 2=raw LoG, 3=accum sum, 4=accum count. Blocking (drains the
+// render stream). out = img_w*img_h floats, sensor-window packed.
+// By Claude on 07/19/2026.
+int tp_proc_render_log_get(TpProc* p, int which, float* out){
+  if(!p||!p->have_render_log||!out){ seterr("render_log_get: bad args / not armed"); return -1; }
+  cuCtxSetCurrent(p->mod->ctx);
+  if(cudaStreamSynchronize(p->render_stream)!=cudaSuccess){ seterr("render_log_get: stream sync failed"); return -2; }
+  const float* buf = (which==0)?p->r_log_diff : (which==1)?p->r_log_avg : (which==2)?p->r_log_raw :
+                     (which==3)?p->r_acc_sum : (which==4)?p->r_acc_cnt : nullptr;
+  if(!buf){ seterr("render_log_get: which=%d not armed",which); return -2; }
+  const int npix=p->img_w*p->img_h;
+  return cudaMemcpy(out,buf,(size_t)npix*sizeof(float),cudaMemcpyDeviceToHost)==cudaSuccess?1:-3;
+}
+
 // Direct task-slot seed (tests + Java-centers A/B): bypasses pose_task_update.
 int tp_proc_render_set_tasks(TpProc* p, int slot, const float* ftasks, int ntiles, int total_floats){
   if(!p||!p->have_render||!ftasks||slot<0||slot>1){ seterr("render_set_tasks: bad args / not armed"); return -1; }
@@ -1878,6 +2015,33 @@ int tp_proc_render_scene(TpProc* p, const float* pose_vec, int vec_floats,
     CUresult cr=cuLaunchKernel(f,1,1,1,1,1,1,0,(CUstream)p->render_stream,a,nullptr);
     if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es);
       seterr("render_scene: launch render_dp -> %d (%s)",cr,es); return -5; } }
+  // item-4 R2: pixel-LoG + subtract-average (+ optional raw-average
+  // accumulation) on the finished frame - stream order after the CDP parent
+  // (a grid is not complete until its children are), before render_done.
+  // Window = sensor image at (+4,+4) inside the padded r_rbg (the renderGet
+  // de-pitch convention). By Claude on 07/19/2026.
+  if(p->have_render_log){
+    CUfunction fl=getfun(p->mod,"render_pix_log");
+    if(!fl){ seterr("render_scene: render_pix_log missing"); return -5; }
+    int sstride=(int)(p->r_dstride_rslt/sizeof(float));
+    int xo=4, yo=4; // DTT_SIZE/2 - matches the img_w+8 x img_h+8 rbg padding
+    int w=p->img_w, h=p->img_h;
+    void* al[]={ &p->r_rbg,&sstride,&xo,&yo,&w,&h,&p->r_log_taps,&p->log_krad,
+                 &p->r_log_raw,&p->r_log_avg,&p->r_log_diff,&p->log_ema,&p->log_gate };
+    CUresult cr=cuLaunchKernel(fl,(unsigned)((w+15)/16),(unsigned)((h+15)/16),1,16,16,1,0,
+                               (CUstream)p->render_stream,al,nullptr);
+    if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es);
+      seterr("render_scene: launch render_pix_log -> %d (%s)",cr,es); return -5; }
+    if(p->log_accum){
+      CUfunction fa=getfun(p->mod,"render_avg_accum");
+      if(!fa){ seterr("render_scene: render_avg_accum missing"); return -5; }
+      void* aa[]={ &p->r_rbg,&sstride,&xo,&yo,&w,&h,&p->r_acc_sum,&p->r_acc_cnt };
+      cr=cuLaunchKernel(fa,(unsigned)((w+15)/16),(unsigned)((h+15)/16),1,16,16,1,0,
+                        (CUstream)p->render_stream,aa,nullptr);
+      if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es);
+        seterr("render_scene: launch render_avg_accum -> %d (%s)",cr,es); return -5; }
+    }
+  }
   if(cudaEventRecord(p->render_done,p->render_stream)!=cudaSuccess){
     seterr("render_scene: event record failed"); return -5; }
   p->render_pending=true; p->render_img_set=img_set;
@@ -3904,6 +4068,13 @@ void tp_proc_destroy(TpProc* p){
   cudaFree(p->r_active); cudaFree(p->r_num_active); cudaFree(p->r_rbg);
   cudaFree(p->r_imclt_clt); cudaFree(p->r_imclt_img); cudaFree(p->r_status);
   cudaFree(p->gpu_render_dp);
+  // item-4 R2 buffers + camera staging // By Claude on 07/19/2026
+  cudaFree(p->r_log_taps); cudaFree(p->r_log_raw); cudaFree(p->r_log_avg); cudaFree(p->r_log_diff);
+  cudaFree(p->r_acc_sum); cudaFree(p->r_acc_cnt);
+  if(p->r_cam_staging[0]) cudaFreeHost(p->r_cam_staging[0]);
+  if(p->r_cam_staging[1]) cudaFreeHost(p->r_cam_staging[1]);
+  if(p->r_cam_staged[0]) cudaEventDestroy(p->r_cam_staged[0]);
+  if(p->r_cam_staged[1]) cudaEventDestroy(p->r_cam_staged[1]);
   // rung B4: drain the copy stream before freeing pinned staging // By Claude on 07/16/2026
   if(p->bayer_copy_stream){ cudaStreamSynchronize(p->bayer_copy_stream); cudaStreamDestroy(p->bayer_copy_stream); }
   if(p->bayer_copy_event) cudaEventDestroy(p->bayer_copy_event);

src/tp_condition.cu

@@ -109,3 +109,89 @@ extern "C" __global__ void condition_lwir_u16(
 			scale_maps[map_index],
 			-offset_maps[map_index]);
 }
+
+// imagej ROADMAP item 4 rung R2: pixel LoG + post-LoG subtract-average on the
+// resident render frame. See tp_condition.h for the frozen contract (Java
+// oracle = CuasRTUtils.convolve2DLReLU LINEAR form on the sensor window; all
+// accumulation through fmaf so a same-order CPU std::fma replay is bit-exact).
+// By Claude on 07/19/2026.
+extern "C" __global__ void render_pix_log(
+		const float * src,
+		int           src_stride,
+		int           x0,
+		int           y0,
+		int           width,
+		int           height,
+		const float * taps,
+		int           krad,
+		float       * log_raw,
+		float       * avg,
+		float       * diff,
+		float         ema_k,
+		float         gate_abs)
+{
+	const int x = (int) (blockIdx.x * blockDim.x + threadIdx.x);
+	const int y = (int) (blockIdx.y * blockDim.y + threadIdx.y);
+	if ((x >= width) || (y >= height)) return;
+	const int idx = y * width + x;
+	const float c = src[(size_t) (y0 + y) * src_stride + (x0 + x)];
+	float L;
+	if (isnan(c)) {
+		L = nanf("");
+	} else {
+		// same loop order as the Java oracle (iy outer, ix inner, top-left first)
+		float v = 0.0f;
+		const int ksize = 2 * krad + 1;
+		for (int iy = 0; iy < ksize; iy++) {
+			const int yy = y + iy - krad;
+			if ((yy < 0) || (yy >= height)) continue; // out-of-window taps contribute 0
+			const float * srow = src + (size_t) (y0 + yy) * src_stride + x0;
+			const float * trow = taps + iy * ksize;
+			for (int ix = 0; ix < ksize; ix++) {
+				const int xx = x + ix - krad;
+				if ((xx < 0) || (xx >= width)) continue;
+				const float d = srow[xx];
+				v = fmaf(isnan(d) ? c : d, trow[ix], v);
+			}
+		}
+		L = v;
+	}
+	if (log_raw) log_raw[idx] = L;
+	if (!avg) { diff[idx] = L; return; }
+	const float A = avg[idx];
+	if (isnan(L)) {
+		diff[idx] = nanf("");            // average keeps its last state
+	} else if (isnan(A)) {
+		if (ema_k > 0.0f) { avg[idx] = L; diff[idx] = 0.0f; } // bootstrap seed
+		else diff[idx] = nanf("");       // prior-file-only mode: never seen
+	} else {
+		const float d = L - A;
+		diff[idx] = d;
+		if ((ema_k > 0.0f) && ((gate_abs <= 0.0f) || (fabsf(d) <= gate_abs))) {
+			avg[idx] = fmaf(ema_k, d, A);    // gated EMA refresh (fused axpy)
+		}
+	}
+}
+
+// R2 companion: raw-domain running average accumulation (next run's prior file).
+// By Claude on 07/19/2026.
+extern "C" __global__ void render_avg_accum(
+		const float * src,
+		int           src_stride,
+		int           x0,
+		int           y0,
+		int           width,
+		int           height,
+		float       * acc_sum,
+		float       * acc_cnt)
+{
+	const int x = (int) (blockIdx.x * blockDim.x + threadIdx.x);
+	const int y = (int) (blockIdx.y * blockDim.y + threadIdx.y);
+	if ((x >= width) || (y >= height)) return;
+	const int idx = y * width + x;
+	const float c = src[(size_t) (y0 + y) * src_stride + (x0 + x)];
+	if (!isnan(c)) {
+		acc_sum[idx] += c;
+		acc_cnt[idx] += 1.0f;
+	}
+}

src/tp_condition.h

@@ -66,4 +66,51 @@ extern "C" __global__ void condition_lwir_u16(
 		const float          * scale_maps,
 		const float          * offset_maps);
 
+/* imagej ROADMAP item 4 rung R2 (07/19/2026): pixel-domain LoG + subtract-
+ * average on the GPU-resident detection-input render (r_rbg), enqueued on the
+ * render stream right after the render_dp chain. The math is the LINEAR
+ * (alpha=1) form of CuasRTUtils.convolve2DLReLU restricted to the sensor
+ * window (width x height at (x0,y0) inside the padded imclt buffer):
+ *   - center NaN -> output NaN (a NaN neighbor never blooms into a valid pixel);
+ *   - NaN tap    -> substitute the center value (edge replication);
+ *   - out-of-window tap -> contributes 0 (the Java TileNeibs margin behavior).
+ * Every product accumulates through fmaf() so a CPU float replay using
+ * std::fma in the same loop order is BIT-exact (the log_pipe case gate).
+ *
+ * Subtract-average runs in the POST-LoG domain (Andrey ruling 3.1: LoG is
+ * linear, so maintaining the average post-LoG is the same math, fused axpy):
+ *   diff = LoG - avg; gated EMA refresh avg += ema_k*diff only where
+ *   |diff| <= gate_abs, so slow-apparent-motion targets are never absorbed.
+ * avg NaN + LoG valid seeds avg = LoG (diff 0) when EMA is on - the no-prior-
+ * file bootstrap; with EMA off the diff stays NaN there (prior file only).
+ * By Claude on 07/19/2026. */
+extern "C" __global__ void render_pix_log(
+		const float * src,        // padded render buffer (r_rbg)
+		int           src_stride, // floats per row
+		int           x0,         // window origin in src (DTT_SIZE/2)
+		int           y0,
+		int           width,      // sensor window (640 x 512)
+		int           height,
+		const float * taps,       // (2*krad+1)^2 LoG taps, row-major
+		int           krad,
+		float       * log_raw,    // may be NULL: raw LoG (debug/eval)
+		float       * avg,        // may be NULL: post-LoG average, updated in place
+		float       * diff,       // detection frame out: LoG - avg (or LoG when avg NULL)
+		float         ema_k,      // 0 = no EMA refresh
+		float         gate_abs);  // >0: update only where |diff| <= gate_abs
+
+/* R2 companion: NaN-aware raw-domain accumulation of the rendered frames
+ * (sum += pixel, cnt += 1 where valid) over the run - D2H once after the loop
+ * builds the NEXT run's prior-average file (ruling 3.1: the base is a prior-
+ * run average, first rotations are unrepresentative). By Claude on 07/19/2026. */
+extern "C" __global__ void render_avg_accum(
+		const float * src,        // padded render buffer (r_rbg)
+		int           src_stride,
+		int           x0,
+		int           y0,
+		int           width,
+		int           height,
+		float       * acc_sum,    // [height*width] running sum
+		float       * acc_cnt);   // [height*width] valid-sample count
+
 #endif /* SRC_TP_CONDITION_H_ */