CLAUDE: DP rung D2 - device measure chain from ONE parent entry

PoseMeasureChain descriptor (tp_lma.h) = the rung-B3 per-scene pcyc_*
union + the resident buffer pointers the granular stage functions launch
with - the DP parent's marshalling, consumed verbatim. Geometry buffers
stay opaque (void*) so the header remains host-compilable; tp_lma.cu
casts at the launch sites (one NVRTC translation unit).

pose_measure_dp tail-enqueues the EXACT B3 host stage sequence:
task_update -> [slot0 rot-deriv + tile offsets + erase(NaN) + convert
SET] -> [slot1 rot-deriv + tile offsets + convert SUBTRACT] ->
consolidate -> inter-corr -> normalize -> peak. Payload kernels are the
UNMODIFIED production ones (most already single-thread CDP parents -
the pre-JNA top-DP pattern), so the chain is byte-identical BY
CONSTRUCTION; only the interior host syncs and the mid-chain 4-byte
count readbacks disappear. Mid-chain counts (MB pairs / surviving tiles
/ corr rows) are read by self-chaining launchers (the D1 pattern) that
run tail-enqueued BEHIND the counting kernel and check the LIVE count
against the descriptor's provisioned capacity (device count guards;
static over-provisioned shapes per the ratified D2 rule). New helper
kernels only move bytes: gather/scatter of the per-cam CLT slices and
the 0xFF NaN-poison fill (the host DtoD memcpys/memsets).

Wrapper tp_proc_exec_pose_measure_cycle_dp (B3-signature drop-in):
host keeps the rung-C2 nullable uploads (pose_task_upload split out of
exec_pose_task_update, same bytes), the per-sequence reverse-distortion
build, consolidate/peak capacity provisioning (no device-side malloc),
the B4 bayer fence, ONE parent launch, ONE end-of-chain 5-int status
D2H (rc/pairs/surviving/misaligned/corr rows). Peaks/corr indices stay
resident for the unchanged fetch calls.

Gate: test_pose_measure_dp_jna (cases.list 'measure_dp', rides the
captured pose_corr case for real images/kernels/geometry/center-TD;
task template txy/centers extracted from the captured production task
stream; camera pair synthetic per the B1/D1 two-instance rule). Three
poses, B3 entry vs DP chain, ORDER-INDEPENDENT KEYED BY PACKED INDEX:
counts/stats equal, index sets equal, raw-TD (256) + FZ-normalized PD
(225) + peak (8) rows ALL memcmp-equal (300 corr rows, 150/150 tiles).
run_cases.sh ALL PASS (pose_corr @tol 0, peak, peak_recenter, avg_td,
avg_td_oob, bayer_staged, dp_cycles, measure_dp); DP spike test still
PASS; compute-sanitizer 12.8 memcheck 0 errors.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

cases.list.tmpl

@@ -66,6 +66,7 @@ case avg_td          test_avg_td           ${REPO}/testdata/synth4/avg_td
 case avg_td_oob      test_avg_td_oob       ${REPO}/testdata/synth4/avg_td_oob    --tol 0
 case bayer_staged    test_bayer_staged_jna ${REPO}                               --tol 0     # rung B4: synthetic, --data unused
 case dp_cycles       test_pose_lma_dp_cycles_jna ${REPO}                         --tol 0     # rung D1: synthetic, --data unused; DP K-cycle chain bit-exact vs host cycles # By Claude on 07/17/2026
+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 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_pose_measure_dp_jna.cu

+// test_pose_measure_dp_jna.cu - DP rung D2 gate: the device measure chain
+// (tp_proc_exec_pose_measure_cycle_dp - ONE parent launch tail-running
+// task_update -> geometry/convert x2 -> consolidate -> inter-corr ->
+// normalize -> peak) must be BIT-IDENTICAL to the rung-B3 host-batched entry
+// (tp_proc_exec_pose_measure_cycle) that drives production today.
+//
+// Two TpProc instances from one NVRTC module (the chain mutates resident CLT/
+// corr state; the D1/DP-spike two-instance reasoning). Real content comes from
+// the captured pose_corr case (images, deconvolution kernels, geometry,
+// correction vector, LPF constants, resident center TD, de-bias); the pose/
+// camera-meta/task template are synthetic-but-in-frame (B1/D1 test precedent:
+// bit-exactness is between A and B on the SAME bytes, not against an analytic
+// truth). The task template txy/centers are extracted from the captured
+// production task stream so tiles land exactly where production put them.
+//
+// Compare per iteration, ORDER-INDEPENDENT KEYED BY PACKED INDEX (the recorded
+// rung-1/B3 rule: correlate2D_inter row order is atomicAdd-nondeterministic;
+// per-index content is the contract): corr-row counts + consolidate stats
+// equal, index SETS equal, and per matched packed index the raw-TD row (256),
+// the FZ-normalized PD row (225) and the peak row (8) all memcmp-equal
+// (NaN==NaN by construction - same kernels, same bytes).
+// By Claude on 07/17/2026.
+//
+// Build after jna/build_lib.sh:
+//   cd tile_processor_gpu && /usr/local/cuda-12.8/bin/nvcc -std=c++17 -O2 \
+//     -I src -I $HOME/git/cuda-samples/Common \
+//     jna/test_pose_measure_dp_jna.cu src/tests/tp_test_data.cu \
+//     -o tests_bin/test_pose_measure_dp_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_pose_measure_dp_jna --data <pose_corr case> [--src src]
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cmath>
+#include <chrono>
+#include <map>
+#include <string>
+#include <vector>
+
+#include <cuda_runtime.h>
+#include "tests/tp_test_data.h"
+#include "tp_peak.h"                 // PEAK_ROW_FLOATS
+#include "geometry_correction.h"     // TP_TASK_*_OFFSET
+
+extern "C" {
+  void* tp_create_module(const char*, const char*);
+  const char* tp_last_error();
+  void* tp_proc_create(void*);
+  int   tp_proc_setup(void*, int,int,int,int,int,int);
+  int   tp_proc_set_geometry(void*, const float*, int);
+  int   tp_proc_set_correction_vector(void*, const float*, int);
+  int   tp_proc_set_kernels(void*, int, const float*, int);
+  int   tp_proc_set_kernel_offsets(void*, int, const float*, int);
+  int   tp_proc_set_image(void*, int, const float*);
+  int   tp_proc_set_tasks_slot(void*, const float*, int, int, int);
+  int   tp_proc_set_const(void*, const char*, const float*, int);
+  int   tp_proc_set_clt(void*, int, const float*, int);
+  int   tp_proc_setup_rbg_corr(void*, int, int,int,int,int, float,float,float, int);
+  int   tp_proc_set_peak_debias(void*, const float*, int);
+  int   tp_proc_set_pose_cycle(void*, int,int, float, float,float,float,float,
+                               int,int, float,float,float,float, float,float,
+                               float,float, int, float,float,float, double,int,
+                               int, float,float,float,float, int, float,float,float, int);
+  int   tp_proc_exec_pose_measure_cycle(void*,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const int*,const float*,int*);
+  int   tp_proc_exec_pose_measure_cycle_dp(void*,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const int*,const float*,int*);
+  int   tp_proc_get_corr_indices(void*, int*, int);
+  int   tp_proc_get_corr_td(void*, float*);
+  int   tp_proc_get_corr2d(void*, float*, int);
+  int   tp_proc_get_peaks(void*, float*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+enum { CHUNK=256, PD_SIZE=15*15 };
+
+// bit-exact row compare (NaN==NaN via memcmp) of n floats
+static bool rowEqual(const float* a, const float* b, int n){
+  return std::memcmp(a,b,(size_t)n*sizeof(float))==0;
+}
+
+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++; // manifest form; gate is bit-exact
+    else { fprintf(stderr,"Usage: %s --data <pose_corr case> [--src <dir>] [--devrt <lib>]\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 img_width    = data.paramInt("img_width",   640);
+  const int img_height   = data.paramInt("img_height",  512);
+  const int kernels_hor  = data.paramInt("kernels_hor",  82);
+  const int kernels_vert = data.paramInt("kernels_vert", 66);
+  const int task_size    = data.paramInt("task_size", 6 + 6*num_cams);
+  const int sel_sensors  = data.paramInt("sel_sensors",   1);
+  const float soft_margin = (float)data.param("soft_margin", 12.0);
+  const float hard_margin = (float)data.param("hard_margin",  8.0);
+  const float s0 = (float)data.param("scale0",1.0), s1 = (float)data.param("scale1",0.0), s2 = (float)data.param("scale2",0.0);
+  const double fz_inter   = data.param("fz_inter", 10000.0);
+  const int   corr_rad    = 7;
+  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);
+  const int kern_tiles = kernels_hor*kernels_vert*num_colors;
+
+  // per-sequence task template from the captured production stream: txy +
+  // center px,py (disparity 0 = the CUAS infinity branch); task codes verbatim
+  const float* ft0 = data.hostFloat("ftasks0_it0");
+  const float* ft1 = data.hostFloat("ftasks1_it0");
+  const int num_tasks = (int)(data.elements("ftasks0_it0")/task_size);
+  if (num_tasks <= 0 || (int)(data.elements("ftasks1_it0")/task_size) != num_tasks){
+    printf("FAIL: captured task streams missing/mismatched\n"); return EXIT_FAILURE; }
+  std::vector<int>   task_map(num_tasks);
+  std::vector<float> task_centers(3*num_tasks);
+  int task_code_main = 0, task_code_sub = 0;
+  std::memcpy(&task_code_main, ft0 + TP_TASK_TASK_OFFSET, sizeof(int));
+  std::memcpy(&task_code_sub,  ft1 + TP_TASK_TASK_OFFSET, sizeof(int));
+  for (int i = 0; i < num_tasks; i++){
+    std::memcpy(&task_map[i], ft0 + (size_t)i*task_size + TP_TASK_TXY_OFFSET, sizeof(int));
+    task_centers[3*i+0] = ft0[(size_t)i*task_size + TP_TASK_CENTERXY_OFFSET];
+    task_centers[3*i+1] = ft0[(size_t)i*task_size + TP_TASK_CENTERXY_OFFSET+1];
+    task_centers[3*i+2] = 0.0f;
+  }
+  printf("cams=%d frame %dx%d tasks=%d task_size=%d codes 0x%x/0x%x sel=0x%x\n",
+         num_cams,img_width,img_height,num_tasks,task_size,task_code_main,task_code_sub,sel_sensors);
+
+  // synthetic-but-real-dims camera pair (B1/D1 precedent: identical on A and B).
+  // rbr must cover the full-frame radius at step_r (the 64x48 D1 test got away
+  // with 256 entries; 640x512 needs the production RBYRDIST_LEN-scale table).
+  const int rbr_len=5001, ers_len=img_height*14;
+  const float meta[8]={(float)img_width,(float)img_height,4.0f,5.0f,2.85f,1.0f,0.0f,0.0004f};
+  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;
+  // uniform-MB descriptor (plausible magnitudes; identical bytes on A and B)
+  const float dx_main=0.35f, dy_main=-0.25f, dx_sub=-0.35f, dy_sub=0.25f;
+  const float scale_main=1.0f, scale_sub=-1.0f;
+  const float min_px=8.0f, max_px=(float)(img_width-8), min_py=8.0f, max_py=(float)(img_height-8);
+
+  void* m = tp_create_module(srcdir.c_str(), devrt.c_str());
+  if(!m){ printf("FAIL module: %s\n", tp_last_error()); return EXIT_FAILURE; }
+  // A = rung-B3 host-batched oracle, B = DP device chain
+  void* procs[2]={nullptr,nullptr};
+  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);
+  const int num_pairs = num_cams*(num_cams-1)/2;
+  for (int k = 0; k < 2; k++){
+    void* p = tp_proc_create(m);
+    if(!p || tp_proc_setup(p, num_cams, num_colors, img_width, img_height, kernels_hor, kern_tiles)){
+      printf("FAIL setup[%d]: %s\n", k, 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"));
+    for (int c = 0; c < num_cams; c++){
+      tp_proc_set_kernels       (p, c, h_kern + (size_t)c*kern_elems, kern_elems);
+      tp_proc_set_kernel_offsets(p, c, h_offs + (size_t)c*offs_elems, offs_elems);
+      tp_proc_set_image         (p, c, h_imgs + (size_t)c*img_width*img_height);
+    }
+    tp_proc_set_const(p, "lpf_data",    data.hostFloat("lpf_data"),    4*64);
+    tp_proc_set_const(p, "lpf_corr",    data.hostFloat("lpf_corr"),    64);
+    tp_proc_set_const(p, "lpf_rb_corr", data.hostFloat("lpf_rb_corr"), 64);
+    tp_proc_set_const(p, "LoG_corr",    data.hostFloat("log_corr"),    64);
+    tp_proc_set_clt(p, 0, data.hostFloat("clt_ref_cam0"), 1); // resident center TD
+    tp_proc_setup_rbg_corr(p, num_pairs, 0,0,0,0, 1.f,1.f,1.f, corr_rad);
+    if (data.has("debias")) tp_proc_set_peak_debias(p, data.hostFloat("debias"), (int)data.elements("debias"));
+    // skeleton task slots (shape source for the resident task build)
+    if(tp_proc_set_tasks_slot(p, ft0, num_tasks, num_tasks*task_size, 0) ||
+       tp_proc_set_tasks_slot(p, ft1, num_tasks, num_tasks*task_size, 1)){
+      printf("FAIL set_tasks_slot[%d]: %s\n", k, tp_last_error()); return EXIT_FAILURE; }
+    if(tp_proc_set_pose_cycle(p, num_tasks, task_size, 0.0f,
+        min_px, max_px, min_py, max_py, task_code_main, task_code_sub,
+        dx_main, dy_main, dx_sub, dy_sub, scale_main, scale_sub,
+        soft_margin, hard_margin, sel_sensors, s0, s1, s2,
+        fz_inter, corr_rad, 0xff,
+        p_abs_min, p_rel_min, p_min_peak, p_sub_frac,
+        p_recenter, p_sub_frac1, p_scale_ax, p_inc_ax, 1)){
+      printf("FAIL set_pose_cycle[%d]: %s\n", k, tp_last_error()); return EXIT_FAILURE; }
+    procs[k]=p;
+  }
+
+  // three poses: identity, small +, small - (grid-visible shifts, tiles stay in frame)
+  const float poses[3][12]={
+    {0,0,0, 0,0,0, 0,0,0,  0.0f,     0.0f,    0.0f},
+    {0,0,0, 0,0,0, 0,0,0,  0.0020f, -0.0010f, 0.0005f},
+    {0,0,0, 0,0,0, 0,0,0, -0.0015f,  0.0010f,-0.0008f}};
+  int fail = 0;
+  double ms_a = 0, ms_b = 0; int timed = 0;
+  for (int it = 0; it < 3 && !fail; it++){
+    const bool first = (it == 0);
+    int rc[2]={0,0}; int stats[2][3]={{-1,-1,-1},{-1,-1,-1}};
+    for (int k = 0; k < 2; k++){
+      const auto t0=std::chrono::steady_clock::now();
+      const int r = (k==0)
+        ? tp_proc_exec_pose_measure_cycle(procs[k],
+            first?meta:nullptr, first?radial:nullptr, first?rbr.data():nullptr, first?rbr_len:0,
+            first?ers.data():nullptr, first?ers_len:0,
+            first?meta:nullptr, first?radial:nullptr, first?rbr.data():nullptr, first?rbr_len:0,
+            first?ers.data():nullptr, first?ers_len:0,
+            poses[it], first?task_map.data():nullptr, first?task_centers.data():nullptr, stats[k])
+        : tp_proc_exec_pose_measure_cycle_dp(procs[k],
+            first?meta:nullptr, first?radial:nullptr, first?rbr.data():nullptr, first?rbr_len:0,
+            first?ers.data():nullptr, first?ers_len:0,
+            first?meta:nullptr, first?radial:nullptr, first?rbr.data():nullptr, first?rbr_len:0,
+            first?ers.data():nullptr, first?ers_len:0,
+            poses[it], first?task_map.data():nullptr, first?task_centers.data():nullptr, stats[k]);
+      const auto t1=std::chrono::steady_clock::now();
+      if (!first){ (k==0?ms_a:ms_b) += std::chrono::duration<double,std::milli>(t1-t0).count(); timed += (k==0); }
+      if (r <= 0){
+        printf("it%d: %s entry rc=%d: %s -> FAIL\n", it, k?"DP":"B3", r, tp_last_error());
+        return EXIT_FAILURE; }
+      rc[k]=r;
+    }
+    printf("it%d: B3 rc=%d stats={%d,%d,%d}  DP rc=%d stats={%d,%d,%d}\n", it,
+           rc[0],stats[0][0],stats[0][1],stats[0][2], rc[1],stats[1][0],stats[1][1],stats[1][2]);
+    if (rc[0]!=rc[1] || std::memcmp(stats[0],stats[1],sizeof(stats[0]))){
+      printf("it%d: count/stats mismatch -> FAIL\n", it); fail=1; break; }
+    if (stats[0][1] <= 0){ printf("it%d: no surviving tiles - vacuous gate -> FAIL\n", it); fail=1; break; }
+    const int n = rc[0];
+    std::vector<int>   idx[2] = {std::vector<int>(n), std::vector<int>(n)};
+    std::vector<float> td [2] = {std::vector<float>((size_t)n*CHUNK),  std::vector<float>((size_t)n*CHUNK)};
+    std::vector<float> pd [2] = {std::vector<float>((size_t)n*PD_SIZE),std::vector<float>((size_t)n*PD_SIZE)};
+    std::vector<float> pk [2] = {std::vector<float>((size_t)n*PEAK_ROW_FLOATS),std::vector<float>((size_t)n*PEAK_ROW_FLOATS)};
+    for (int k = 0; k < 2; k++){
+      if (tp_proc_get_corr_indices(procs[k], idx[k].data(), n) < 0 ||
+          tp_proc_get_corr_td(procs[k], td[k].data()) < 0 ||
+          tp_proc_get_corr2d(procs[k], pd[k].data(), corr_rad) < 0 ||
+          tp_proc_get_peaks(procs[k], pk[k].data()) < 0){
+        printf("it%d: fetch[%d] failed: %s -> FAIL\n", it, k, tp_last_error()); return EXIT_FAILURE; }
+    }
+    // ORDER-INDEPENDENT compare keyed by packed index (the recorded rule)
+    std::map<int,int> byB;
+    for (int i = 0; i < n; i++) byB[idx[1][i]] = i;
+    long missing=0, dup=(long)(n - (int)byB.size());
+    long td_mm=0, pd_mm=0, pk_mm=0; int nvalid=0;
+    for (int i = 0; i < n; i++){
+      std::map<int,int>::const_iterator it2 = byB.find(idx[0][i]);
+      if (it2 == byB.end()){ missing++; continue; }
+      const int j = it2->second;
+      if (!rowEqual(&td[0][(size_t)i*CHUNK],  &td[1][(size_t)j*CHUNK],  CHUNK))  td_mm++;
+      if (!rowEqual(&pd[0][(size_t)i*PD_SIZE],&pd[1][(size_t)j*PD_SIZE],PD_SIZE)) pd_mm++;
+      if (!rowEqual(&pk[0][(size_t)i*PEAK_ROW_FLOATS],&pk[1][(size_t)j*PEAK_ROW_FLOATS],PEAK_ROW_FLOATS)) pk_mm++;
+      if (pk[0][(size_t)i*PEAK_ROW_FLOATS+7]!=0.0f) nvalid++;
+    }
+    const int it_fail = (missing>0)||(dup>0)||(td_mm>0)||(pd_mm>0)||(pk_mm>0)||(nvalid<=0);
+    printf("it%d: %d corr rows by packed index: missing=%ld dup=%ld TD-mm=%ld PD-mm=%ld peak-mm=%ld valid-peaks=%d -> %s\n",
+           it, n, missing, dup, td_mm, pd_mm, pk_mm, nvalid, it_fail?"FAIL":"PASS (bit-identical)");
+    if (it_fail) fail = 1;
+  }
+  if (timed) printf("microbench (it1..2, incl. per-call H2D/D2H + fetches excluded): B3 %.3f ms/cycle, DP %.3f ms/cycle\n",
+                    ms_a/timed, ms_b/timed);
+  printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+  tp_proc_destroy(procs[0]); tp_proc_destroy(procs[1]); tp_destroy_module(m);
+  return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

jna/tp_jna.cpp

@@ -59,6 +59,7 @@ static const char* KERNELS[] = {
   "pose_lma_prepare_products","pose_lma_reduce_step","pose_lma_reduce_rms", // resident prepare/solve + candidate RMS
   "pose_lma_assemble","pose_lma_prepare_norm","pose_lma_prepare_finish", // rung C1 resident prepareLMA // By Claude on 07/17/2026
   "pose_lma_dp_step","pose_lma_dp_candidate_vector", // DP spike: device-side step chain // By Claude on 07/16/2026
+  "pose_measure_dp", // DP rung D2: device measure-chain parent // By Claude on 07/17/2026
   "lma_normal_products"}; // deterministic double normal-equation products // By Codex on 07/14/2026
 static const int N_KERNELS = sizeof(KERNELS)/sizeof(KERNELS[0]);
 
@@ -705,6 +706,11 @@ struct TpProc {
   // trace (num_cycles x POSE_LMA_RESIDENT_RESULT floats). By Claude on 07/17/2026.
   PoseDpCycles *gpu_pose_dp_cycles;
   float *gpu_pose_dp_trace; int pose_dp_trace_capacity;
+  // DP rung D2: device copy of the measure-chain descriptor + the small status
+  // buffer the launchers report through (rc/pairs/surviving/misaligned/corr
+  // rows - the one end-of-chain D2H). By Claude on 07/17/2026.
+  PoseMeasureChain *gpu_pose_mchain;
+  int *gpu_pose_mchain_status;
   // 3-B rung B1: per-sequence task template (packed txy + reference centers per
   // selected task entry) for the resident task-build kernel. NULL template args
   // on later calls mean "keep resident" (configuration-register semantics,
@@ -798,6 +804,7 @@ TpProc* tp_proc_create(TpModule* m){
   p->pose_prep_ready=0; p->pose_prep_centers_tiles=0; p->pose_prep_frozen_tiles=0; // By Claude on 07/17/2026 (rungs C2/C3)
   p->gpu_pose_dp_chain=nullptr; // By Claude on 07/16/2026 (DP spike)
   p->gpu_pose_dp_cycles=nullptr; p->gpu_pose_dp_trace=nullptr; p->pose_dp_trace_capacity=0; // By Claude on 07/17/2026 (rung D1)
+  p->gpu_pose_mchain=nullptr; p->gpu_pose_mchain_status=nullptr; // By Claude on 07/17/2026 (rung D2)
   p->gpu_pose_task_map=nullptr; p->gpu_pose_task_centers=nullptr; // By Claude on 07/16/2026 (rung B1)
   p->pose_task_capacity=0; p->pose_task_count=0; // By Claude on 07/16/2026 (rung B1)
   p->bayer_staging[0]=p->bayer_staging[1]=nullptr; // By Claude on 07/16/2026 (rung B4)
@@ -1496,24 +1503,24 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p,
 // production caller uploads the Java-tracked measure pose, 48 B). The task
 // slots must hold the per-sequence skeleton (tp_proc_set_tasks_slot) with
 // exactly num_tasks entries. By Claude on 07/16/2026.
-int tp_proc_exec_pose_task_update(TpProc* p,
-                                  const float* ref_meta,
-                                  const float* ref_radial,
-                                  const float* ref_rbr, int ref_rbr_len,
-                                  const float* ref_ers, int ref_ers_len,
-                                  const float* scene_meta,
-                                  const float* scene_radial,
-                                  const float* scene_rbr, int scene_rbr_len,
-                                  const float* scene_ers, int scene_ers_len,
-                                  const float* pose_vectors,
-                                  const int*   task_map,
-                                  const float* task_centers,
-                                  int num_tasks, int num_cams,
-                                  float disparity_corr,
-                                  float min_px, float max_px, float min_py, float max_py,
-                                  int task_code_main, int task_code_sub,
-                                  float dx_main, float dy_main, float dx_sub, float dy_sub,
-                                  float scale_main, float scale_sub){
+// Upload/validation half of tp_proc_exec_pose_task_update, shared with the DP
+// rung-D2 measure-chain entry (same bytes reach the device either way; only
+// who launches pose_task_update differs). Returns 0 and the skeleton-derived
+// task_size on success. By Claude on 07/17/2026 (rung D2 refactor).
+static int pose_task_upload(TpProc* p,
+                            const float* ref_meta,
+                            const float* ref_radial,
+                            const float* ref_rbr, int ref_rbr_len,
+                            const float* ref_ers, int ref_ers_len,
+                            const float* scene_meta,
+                            const float* scene_radial,
+                            const float* scene_rbr, int scene_rbr_len,
+                            const float* scene_ers, int scene_ers_len,
+                            const float* pose_vectors,
+                            const int*   task_map,
+                            const float* task_centers,
+                            int num_tasks, int num_cams,
+                            int* out_task_size){
   const bool up_ref   = (ref_meta!=nullptr)||(ref_radial!=nullptr)||(ref_rbr!=nullptr)||(ref_ers!=nullptr);
   const bool up_scene = (scene_meta!=nullptr)||(scene_radial!=nullptr)||(scene_rbr!=nullptr)||(scene_ers!=nullptr);
   const bool up_templ = (task_map!=nullptr)||(task_centers!=nullptr);
@@ -1623,6 +1630,33 @@ int tp_proc_exec_pose_task_update(TpProc* p,
   if(pose_vectors&&
      cudaMemcpy(p->cur_pose_vectors(),pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
     seterr("exec_pose_task_update: pose HtoD failed"); return -5; }
+  if(out_task_size) *out_task_size=task_size;
+  return 0;
+}
+
+int tp_proc_exec_pose_task_update(TpProc* p,
+                                  const float* ref_meta,
+                                  const float* ref_radial,
+                                  const float* ref_rbr, int ref_rbr_len,
+                                  const float* ref_ers, int ref_ers_len,
+                                  const float* scene_meta,
+                                  const float* scene_radial,
+                                  const float* scene_rbr, int scene_rbr_len,
+                                  const float* scene_ers, int scene_ers_len,
+                                  const float* pose_vectors,
+                                  const int*   task_map,
+                                  const float* task_centers,
+                                  int num_tasks, int num_cams,
+                                  float disparity_corr,
+                                  float min_px, float max_px, float min_py, float max_py,
+                                  int task_code_main, int task_code_sub,
+                                  float dx_main, float dy_main, float dx_sub, float dy_sub,
+                                  float scale_main, float scale_sub){
+  int task_size=0;
+  int rc=pose_task_upload(p,ref_meta,ref_radial,ref_rbr,ref_rbr_len,ref_ers,ref_ers_len,
+                          scene_meta,scene_radial,scene_rbr,scene_rbr_len,scene_ers,scene_ers_len,
+                          pose_vectors,task_map,task_centers,num_tasks,num_cams,&task_size);
+  if(rc) return rc;
   CUfunction f=getfun(p->mod,"pose_task_update");
   if(!f){ seterr("pose_task_update missing"); return -6; }
   float* cur_vectors=p->cur_pose_vectors(); // device-authoritative current set (rung B)
@@ -1787,6 +1821,133 @@ int tp_proc_exec_pose_measure_cycle(TpProc* p,
   return npk;
 }
 
+// DP rung D2 (ladder ratified 2026-07-16): the SAME measure cycle as
+// tp_proc_exec_pose_measure_cycle, but the whole stage chain runs DEVICE-SIDE
+// from ONE parent launch (pose_measure_dp tail-enqueues the exact host
+// sequence; mid-chain counts are read by self-chaining launchers instead of
+// 4-byte D2H round trips). Host keeps: the rung-C2 nullable uploads, the
+// per-sequence reverse-distortion table build, scratch/capacity provisioning
+// (no device-side allocation), the bayer fence, and ONE end-of-chain status
+// D2H. Signature = the B3 entry; returns the corr-row count (>0), a
+// device-reported stage code (-18 no corr tiles, -24 capacity guard), or a
+// host-side negative. Peaks/corr indices stay resident; the existing fetch
+// calls read them exactly as after the B3 entry. By Claude on 07/17/2026.
+int tp_proc_exec_pose_measure_cycle_dp(TpProc* p,
+                                       const float* ref_meta,
+                                       const float* ref_radial,
+                                       const float* ref_rbr, int ref_rbr_len,
+                                       const float* ref_ers, int ref_ers_len,
+                                       const float* scene_meta,
+                                       const float* scene_radial,
+                                       const float* scene_rbr, int scene_rbr_len,
+                                       const float* scene_ers, int scene_ers_len,
+                                       const float* pose_vectors,
+                                       const int*   task_map,
+                                       const float* task_centers,
+                                       int* stats){
+  if(!p){ seterr("exec_pose_measure_cycle_dp: null"); return -1; }
+  if(!p->pcyc_valid){ seterr("exec_pose_measure_cycle_dp: descriptor not set (tp_proc_set_pose_cycle)"); return -2; }
+  if(!p->have_corr){ seterr("exec_pose_measure_cycle_dp: no corr buffers (call setup_rbg_corr)"); return -2; }
+  if(!p->gc||!p->cv||!p->rot||!p->rbr){ seterr("exec_pose_measure_cycle_dp: geometry not set"); return -2; }
+  cuCtxSetCurrent(p->mod->ctx);
+  int task_size=0;
+  int rc=pose_task_upload(p,ref_meta,ref_radial,ref_rbr,ref_rbr_len,ref_ers,ref_ers_len,
+                          scene_meta,scene_radial,scene_rbr,scene_rbr_len,scene_ers,scene_ers_len,
+                          pose_vectors,task_map,task_centers,p->pcyc_num_tasks,p->num_cams,&task_size);
+  if(rc) return -10;
+  if(task_size!=p->pcyc_task_size){
+    seterr("exec_pose_measure_cycle_dp: task_size %d != descriptor %d",task_size,p->pcyc_task_size); return -2; }
+  // per-sequence reverse-distortion table (exec_geometry builds it lazily; the
+  // device chain assumes it resident - same bytes, built once here if missing)
+  if(!p->rbr_ready){
+    CUfunction f_rbr=getfun(p->mod,"calcReverseDistortionTable");
+    if(!f_rbr){ seterr("calcReverseDistortionTable missing"); return -5; }
+    void* a[]={ &p->gc,&p->rbr };
+    if(launch1(f_rbr,16,1,1,3,3,3,a,"calcReverseDistortionTable")) return -6;
+    p->rbr_ready=true;
+  }
+  rc=prepare_consolidate(p,p->pcyc_num_tasks,p->pcyc_task_size); if(rc) return rc;
+  // Provision the peak rows for the worst-case inter-corr emission (per task
+  // tile: selected sensors + the sum row), bounded by the corr-TD row capacity.
+  // Static over-provisioning is the ratified D2 shape rule - the launcher's
+  // count guard checks the LIVE count against this. By Claude on 07/17/2026.
+  {
+    int cap=(p->num_cams+1)*p->pcyc_num_tasks;
+    const int cap_td=p->num_pairs*p->tilesx*p->tilesy;
+    if(cap>cap_td) cap=cap_td;
+    if(cap>p->peak_capacity){
+      cudaFree(p->gpu_peaks); p->gpu_peaks=nullptr;
+      if(cudaMalloc((void**)&p->gpu_peaks,(size_t)cap*PEAK_ROW_FLOATS*sizeof(float))!=cudaSuccess){
+        p->peak_capacity=0; seterr("exec_pose_measure_cycle_dp: peaks alloc failed (%d rows)",cap); return -4; }
+      p->peak_capacity=cap;
+    }
+  }
+  bayer_fence(p); // staged uploads must land before convert reads images (rung B4)
+  if(!p->gpu_pose_mchain &&
+     cudaMalloc((void**)&p->gpu_pose_mchain,sizeof(PoseMeasureChain))!=cudaSuccess){
+    seterr("exec_pose_measure_cycle_dp: descriptor alloc failed"); return -4; }
+  if(!p->gpu_pose_mchain_status &&
+     cudaMalloc((void**)&p->gpu_pose_mchain_status,POSE_MEASURE_STATUS_INTS*sizeof(int))!=cudaSuccess){
+    seterr("exec_pose_measure_cycle_dp: status alloc failed"); return -4; }
+  PoseMeasureChain h;
+  h.num_cams=p->num_cams; h.num_colors=p->num_colors; h.tilesx=p->tilesx; h.tilesy=p->tilesy;
+  h.num_tasks=p->pcyc_num_tasks; h.task_size=p->pcyc_task_size;
+  h.img_w=p->img_w; h.img_h=p->img_h; h.kernels_hor=p->kernels_hor; h.kernels_vert=p->kernels_vert;
+  h.uniform_grid=0; h.corr_capacity=p->peak_capacity;
+  h.disparity_corr=p->pcyc_disparity_corr;
+  h.min_px=p->pcyc_min_px; h.max_px=p->pcyc_max_px; h.min_py=p->pcyc_min_py; h.max_py=p->pcyc_max_py;
+  h.task_code_main=p->pcyc_task_code_main; h.task_code_sub=p->pcyc_task_code_sub;
+  h.dx_main=p->pcyc_dx_main; h.dy_main=p->pcyc_dy_main; h.dx_sub=p->pcyc_dx_sub; h.dy_sub=p->pcyc_dy_sub;
+  h.scale_main=p->pcyc_scale_main; h.scale_sub=p->pcyc_scale_sub;
+  h.erase_fill=NAN; // slot-0 SET semantics (exec_convert_direct erase_clt=1)
+  h.soft_margin=p->pcyc_soft_margin; h.hard_margin=p->pcyc_hard_margin;
+  h.sel_sensors=p->pcyc_sel_sensors;
+  h.sc[0]=p->pcyc_sc[0]; h.sc[1]=p->pcyc_sc[1]; h.sc[2]=p->pcyc_sc[2];
+  h.fat_zero2=(float)(p->pcyc_fat_zero*p->pcyc_fat_zero); // = exec_corr2d_normalize
+  h.corr_radius=p->pcyc_corr_radius;
+  h.pair_select=p->pcyc_pair_select;
+  h.abs_min=p->pcyc_abs_min; h.rel_min=p->pcyc_rel_min; h.min_peak=p->pcyc_min_peak;
+  h.eig_sub_frac=p->pcyc_eig_sub_frac; h.refine=p->pcyc_refine; h.eig_sub_frac1=p->pcyc_eig_sub_frac1;
+  h.scale_axes=p->pcyc_scale_axes; h.inc_axes=p->pcyc_inc_axes; h.use_float=p->pcyc_use_float;
+  h.cameras=p->gpu_pose_cameras;
+  h.ref_radial=p->gpu_pose_radial[0]; h.ref_rbr=p->gpu_pose_rbr[0]; h.ref_ers=p->gpu_pose_ers[0];
+  h.scene_radial=p->gpu_pose_radial[1]; h.scene_ers=p->gpu_pose_ers[1];
+  h.task_map=p->gpu_pose_task_map; h.task_centers=p->gpu_pose_task_centers;
+  h.set_vectors[0]=p->gpu_pose_vectors; h.set_vectors[1]=p->gpu_pose_candidate_vectors;
+  h.current_set=p->gpu_pose_current_set;
+  h.gc=p->gc; h.cv=p->cv; h.rbr=p->rbr; h.rot=p->rot;
+  h.ftasks0=p->task_slots[0]; h.ftasks1=p->task_slots[1];
+  h.kernel_offsets=p->gpu_kernel_offsets; h.kernels=p->gpu_kernels; h.images=p->gpu_images;
+  h.clt=p->gpu_clt; h.clt_ref=p->gpu_clt_ref;
+  h.dstride_img=p->dstride_img/sizeof(float);
+  h.active=p->active; h.num_active=p->num_active;
+  h.cons_td_sens=p->cons_td_sens; h.cons_td_avg=p->cons_td_avg;
+  h.cons_pairs=p->cons_pairs; h.cons_tasks_out=p->cons_tasks_out;
+  h.cons_counters=p->cons_counters; h.slice=p->slice;
+  h.corr_indices=p->gpu_corr_indices; h.num_corr_tiles=p->gpu_num_corr_tiles;
+  h.dstride_corr_td=p->dstride_corr_td/sizeof(float); h.corrs_td=p->gpu_corrs_td;
+  h.dstride_corr=p->dstride_corr/sizeof(float); h.corrs=p->gpu_corrs;
+  h.peak_debias=p->gpu_peak_debias; h.peaks=p->gpu_peaks;
+  h.status=p->gpu_pose_mchain_status;
+  if(cudaMemcpy(p->gpu_pose_mchain,&h,sizeof(PoseMeasureChain),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_measure_cycle_dp: descriptor HtoD failed"); return -4; }
+  CUfunction dp=getfun(p->mod,"pose_measure_dp");
+  if(!dp){ seterr("pose_measure_dp missing (DP kernel not in module)"); return -5; }
+  void* da[]={ &p->gpu_pose_mchain };
+  if(launch1(dp,1,1,1,1,1,1,da,"pose_measure_dp")) return -6;
+  int st[POSE_MEASURE_STATUS_INTS]={0,0,0,0,0};
+  if(cudaMemcpy(st,p->gpu_pose_mchain_status,sizeof(st),cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_measure_cycle_dp: status DtoH failed"); return -7; }
+  if(stats){ stats[0]=st[1]; stats[1]=st[2]; stats[2]=st[3]; }
+  p->last_num_corr_tiles=st[4]; // fetch calls (indices/TD/peaks) read this mirror
+  if(st[0]<=0){
+    seterr("exec_pose_measure_cycle_dp: device chain rc=%d (pairs=%d out=%d corr=%d)",
+           st[0],st[1],st[2],st[4]);
+    return (st[0]<0)?st[0]:-6;
+  }
+  return st[0];
+}
+
 // Resident validation boundary: raw fx/J are consumed directly from the most
 // recent pose_fx_jacobian launch. Only y/weights/eigen/current vector cross H2D;
 // transformed tile arrays cross D2H only for the explicitly requested oracle.
@@ -2641,6 +2802,7 @@ void tp_proc_destroy(TpProc* p){
   cudaFree(p->gpu_pose_lma_prep); // By Claude on 07/17/2026
   cudaFree(p->gpu_pose_dp_chain); // By Claude on 07/16/2026 (DP spike)
   cudaFree(p->gpu_pose_dp_cycles); cudaFree(p->gpu_pose_dp_trace); // By Claude on 07/17/2026 (rung D1)
+  cudaFree(p->gpu_pose_mchain); cudaFree(p->gpu_pose_mchain_status); // By Claude on 07/17/2026 (rung D2)
   cudaFree(p->gpu_pose_task_map); cudaFree(p->gpu_pose_task_centers); // By Claude on 07/16/2026 (rung B1)
   delete p;
 }

src/tp_lma.cu

@@ -1237,6 +1237,193 @@ extern "C" __global__ void pose_lma_dp_cycles(const PoseDpCycles * c, int cycle)
 	}
 }
 
+// ── DP rung D2: device measure chain (ladder ratified 2026-07-16) ────────────
+// pose_measure_dp tail-enqueues the EXACT host stage sequence of the rung-B3
+// entry tp_proc_exec_pose_measure_cycle: task_update -> [slot0 rot-deriv +
+// tile offsets + erase(NaN) + convert SET] -> [slot1 rot-deriv + tile offsets
+// + convert SUBTRACT] -> consolidate -> inter-corr -> normalize -> peak.
+// The payload kernels are the UNMODIFIED production ones; most are themselves
+// single-thread CDP parents (Andrey's pre-JNA top-DP pattern), so this nests
+// them one level deeper - byte-identical BY CONSTRUCTION, only the interior
+// host syncs and the mid-chain 4-byte count readbacks disappear. Counters
+// written mid-chain are read by the NEXT launcher (it runs tail-enqueued
+// BEHIND the counting kernel) and checked against the descriptor's
+// provisioned capacities (device count guards). By Claude on 07/17/2026.
+
+// Contiguous [num_cams][slice] gather of the per-cam CLT slices (the device
+// equivalent of exec_consolidate_device's per-cam DtoD memcpy loop).
+extern "C" __global__ void pose_dp_gather_slices(
+		float * const * src,      // per-cam CLT slice pointers
+		size_t          slice,    // floats per slice
+		float         * dst)      // contiguous [num_cams][slice]
+{
+	const float * s = src[blockIdx.y];
+	float * d = dst + (size_t) blockIdx.y * slice;
+	for (size_t i = (size_t) blockIdx.x * blockDim.x + threadIdx.x; i < slice;
+			i += (size_t) gridDim.x * blockDim.x) {
+		d[i] = s[i];
+	}
+}
+
+// The consolidated average -> the cam-0 CLT slice (the setCltData(0,..) copy).
+extern "C" __global__ void pose_dp_scatter_slice(
+		float * const * dst,      // per-cam CLT slice pointers (cam 0 target)
+		const float   * src,
+		size_t          slice)
+{
+	float * d = dst[0];
+	for (size_t i = (size_t) blockIdx.x * blockDim.x + threadIdx.x; i < slice;
+			i += (size_t) gridDim.x * blockDim.x) {
+		d[i] = src[i];
+	}
+}
+
+// Bit-pattern fill (the cudaMemset(0xFF) NaN poison of cons_td_avg).
+extern "C" __global__ void pose_dp_fill_bits(
+		size_t          n,        // floats
+		unsigned int    bits,
+		float         * data)
+{
+	unsigned int * d = (unsigned int *) data;
+	for (size_t i = (size_t) blockIdx.x * blockDim.x + threadIdx.x; i < n;
+			i += (size_t) gridDim.x * blockDim.x) {
+		d[i] = bits;
+	}
+}
+
+// Tail launcher: runs AFTER correlate2D_inter, reads the LIVE corr-row count,
+// guards it, then enqueues normalize + peak with the exact host arguments.
+extern "C" __global__ void pose_measure_dp_tail(const PoseMeasureChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	const int nct = *c->num_corr_tiles;
+	c->status[4] = nct;
+	if (nct <= 0)               { c->status[0] = -18; return; } // = B3 -18 (no corr tiles)
+	if (nct > c->corr_capacity) { c->status[0] = -24; return; } // device count guard
+	corr2D_normalize<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			nct, c->dstride_corr_td, c->corrs_td, (float *) 0,
+			c->dstride_corr, c->corrs, c->fat_zero2, c->corr_radius);
+	if (c->use_float) {
+		corr2D_peak_eig_f<<<nct, 256, 0, cudaStreamTailLaunch>>>(
+				nct, c->corr_indices, c->pair_select, c->dstride_corr, c->corrs,
+				c->corr_radius, c->abs_min, c->rel_min, c->min_peak, c->eig_sub_frac,
+				c->refine, c->eig_sub_frac1, c->scale_axes, c->inc_axes,
+				c->peak_debias, c->peaks);
+	} else {
+		corr2D_peak_eig<<<nct, 256, 0, cudaStreamTailLaunch>>>(
+				nct, c->corr_indices, c->pair_select, c->dstride_corr, c->corrs,
+				c->corr_radius, c->abs_min, c->rel_min, c->min_peak, c->eig_sub_frac,
+				c->refine, c->eig_sub_frac1, c->scale_axes, c->inc_axes,
+				c->peak_debias, c->peaks);
+	}
+	c->status[0] = nct; // npk == nct (exec_corr2d_peak return convention)
+}
+
+// Average launcher: reads the LIVE surviving-tile count, averages, copies the
+// average into cam 0 and starts the inter correlation (which resets its own
+// counter device-side), then chains the tail.
+extern "C" __global__ void pose_measure_dp_avg(const PoseMeasureChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	const int num_out = c->cons_counters[1];
+	c->status[2] = num_out;
+	c->status[3] = c->cons_counters[2];
+	if (num_out > 0) {
+		clt_average_sensors_list<<<dim3((unsigned) num_out, (unsigned) c->num_colors, 1),
+				TD_TILE_FLOATS, 0, cudaStreamTailLaunch>>>(
+				c->num_cams, c->cons_td_sens, c->tilesx, c->tilesy, c->num_colors,
+				c->cons_tasks_out, c->task_size, num_out, c->cons_td_avg, (float *) 0);
+	}
+	pose_dp_scatter_slice<<<1024, 256, 0, cudaStreamTailLaunch>>>(
+			c->clt, c->cons_td_avg, c->slice);
+	correlate2D_inter<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_cams, c->sel_sensors, c->clt, c->clt_ref, c->num_colors,
+			c->sc[0], c->sc[1], c->sc[2], c->ftasks1, c->num_tasks, c->tilesx,
+			c->corr_indices, c->num_corr_tiles, c->dstride_corr_td, c->corrs_td);
+	pose_measure_dp_tail<<<1, 1, 0, cudaStreamTailLaunch>>>(c);
+}
+
+// OOB launcher: reads the LIVE MB-pair count and runs the margin/priority pass.
+extern "C" __global__ void pose_measure_dp_oob(const PoseMeasureChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	const int num_pairs = c->cons_counters[0];
+	c->status[1] = num_pairs;
+	if (num_pairs > 0) {
+		consolidate_oob<<<(num_pairs + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				c->task_size, c->num_cams, c->cons_pairs, num_pairs,
+				(float) c->img_w, (float) c->img_h, c->soft_margin, c->hard_margin,
+				c->cons_tasks_out, c->cons_counters + 1);
+	}
+	pose_measure_dp_avg<<<1, 1, 0, cudaStreamTailLaunch>>>(c);
+}
+
+// Consolidate launcher: runs AFTER both converts; zeroes the counters (the
+// host cudaMemset), gathers/poisons the scratch and indexes the MB pairs.
+extern "C" __global__ void pose_measure_dp_consolidate(const PoseMeasureChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	c->cons_counters[0] = 0;
+	c->cons_counters[1] = 0;
+	c->cons_counters[2] = 0;
+	pose_dp_gather_slices<<<dim3(1024, (unsigned) c->num_cams, 1), 256, 0, cudaStreamTailLaunch>>>(
+			c->clt, c->slice, c->cons_td_sens);
+	pose_dp_fill_bits<<<1024, 256, 0, cudaStreamTailLaunch>>>(
+			c->slice, 0xFFFFFFFFu, c->cons_td_avg);
+	index_consolidate<<<(c->num_tasks + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+			c->task_size, c->ftasks0, c->ftasks1, c->num_tasks,
+			c->cons_pairs, c->cons_counters + 0, c->cons_counters + 2);
+	pose_measure_dp_oob<<<1, 1, 0, cudaStreamTailLaunch>>>(c);
+}
+
+// The D2 parent: ONE device entry for the whole measure cycle. Resolves the
+// CURRENT pose set LIVE (rung-B device-authoritative index; equals the host
+// mirror at entry today, and stays correct when D3 interleaves accept flips).
+extern "C" __global__ void pose_measure_dp(const PoseMeasureChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	c->status[0] = 0; c->status[1] = 0; c->status[2] = 0;
+	c->status[3] = 0; c->status[4] = 0;
+	const float * pose_vectors = c->set_vectors[*c->current_set];
+	pose_task_update<<<(c->num_tasks + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+			c->num_tasks, c->task_size, c->cameras,
+			c->ref_radial, c->ref_rbr, c->ref_ers,
+			c->scene_radial, c->scene_ers,
+			pose_vectors, c->task_map, c->task_centers,
+			c->disparity_corr, c->min_px, c->max_px, c->min_py, c->max_py,
+			c->task_code_main, c->task_code_sub,
+			c->dx_main, c->dy_main, c->dx_sub, c->dy_sub,
+			c->scale_main, c->scale_sub,
+			c->ftasks0, c->ftasks1);
+	// slot 0 (MB main): geometry -> erase(NaN) -> convert SET
+	calc_rot_deriv<<<c->num_cams, dim3(3, 3, 3), 0, cudaStreamTailLaunch>>>(
+			c->num_cams, (struct corr_vector *) c->cv, (trot_deriv *) c->rot);
+	calculate_tiles_offsets<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->uniform_grid, c->num_cams, c->ftasks0, c->num_tasks,
+			(struct gc *) c->gc, (struct corr_vector *) c->cv,
+			(float *) c->rbr, (trot_deriv *) c->rot);
+	erase_clt_tiles<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_cams, c->num_colors, c->tilesx, c->tilesy, c->clt, c->erase_fill);
+	convert_direct<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_cams, c->num_colors, c->kernel_offsets, c->kernels, c->images,
+			c->ftasks0, c->clt, c->dstride_img, c->num_tasks, 0,
+			c->img_w, c->img_h, c->kernels_hor, c->kernels_vert,
+			c->active, c->num_active, c->tilesx);
+	// slot 1 (MB partner): geometry -> convert SUBTRACT (negative scale in tasks)
+	calc_rot_deriv<<<c->num_cams, dim3(3, 3, 3), 0, cudaStreamTailLaunch>>>(
+			c->num_cams, (struct corr_vector *) c->cv, (trot_deriv *) c->rot);
+	calculate_tiles_offsets<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->uniform_grid, c->num_cams, c->ftasks1, c->num_tasks,
+			(struct gc *) c->gc, (struct corr_vector *) c->cv,
+			(float *) c->rbr, (trot_deriv *) c->rot);
+	convert_direct<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_cams, c->num_colors, c->kernel_offsets, c->kernels, c->images,
+			c->ftasks1, c->clt, c->dstride_img, c->num_tasks, 0,
+			c->img_w, c->img_h, c->kernels_hor, c->kernels_vert,
+			c->active, c->num_active, c->tilesx);
+	pose_measure_dp_consolidate<<<1, 1, 0, cudaStreamTailLaunch>>>(c);
+}
+
 #endif // __CUDACC_RTC__
 
 extern "C" __global__ void lma_normal_products(

src/tp_lma.h

@@ -311,6 +311,69 @@ typedef struct {
 	float * trace;             // [num_cycles][POSE_LMA_RESIDENT_RESULT], NULL = off
 } PoseDpCycles;
 
+// ── DP rung D2 (ladder ratified 2026-07-16): device measure chain ────────────
+// PoseMeasureChain = the rung-B3 per-scene descriptor (pcyc_*) plus the
+// resident buffer pointers the granular stage functions launch with - the DP
+// parent's marshalling, consumed verbatim. Geometry buffers are OPAQUE (void*)
+// so this header stays host-compilable without the geometry structs; tp_lma.cu
+// casts at the launch sites (same NVRTC translation unit). Counts produced
+// MID-CHAIN (MB pairs / surviving tiles / corr rows) are read by small
+// self-chaining launchers (the D1 pattern: each runs tail-enqueued BEHIND the
+// counting kernel, so it sees the LIVE counter) and are bounded by the
+// descriptor's provisioned capacities - the device count guards; the corr-row
+// and tile counts are per-sequence constants, so the static shapes hold.
+// By Claude on 07/17/2026.
+typedef struct {
+	// shapes
+	int num_cams, num_colors, tilesx, tilesy;
+	int num_tasks, task_size;
+	int img_w, img_h, kernels_hor, kernels_vert;
+	int uniform_grid;          // 0 = use provided centers (the lean chain)
+	int corr_capacity;         // corr rows peaks/status are provisioned for (count guard)
+	// task-update policy (per-scene uniform-MB descriptor, = pcyc_*)
+	float disparity_corr, min_px, max_px, min_py, max_py;
+	int   task_code_main, task_code_sub;
+	float dx_main, dy_main, dx_sub, dy_sub, scale_main, scale_sub;
+	// consolidate / correlate / normalize / peak policy (= pcyc_*)
+	float erase_fill;          // host NAN (slot-0 SET semantics, erase_clt=1)
+	float soft_margin, hard_margin;
+	int   sel_sensors; float sc[3];
+	float fat_zero2;           // host float((double)fz*fz), = exec_corr2d_normalize
+	int   corr_radius;
+	int   pair_select; float abs_min, rel_min, min_peak, eig_sub_frac;
+	int   refine; float eig_sub_frac1, scale_axes, inc_axes; int use_float;
+	// camera/template state (resident, rung-C2 registers)
+	const PoseCameraMeta * cameras;
+	const float * ref_radial; const float * ref_rbr; const float * ref_ers;
+	const float * scene_radial; const float * scene_ers;
+	const int   * task_map; const float * task_centers;
+	// pose: fixed physical slots + live set index (rung B; D3-readiness - the
+	// parent resolves the CURRENT set exactly like pose_lma_dp_cycles)
+	float * set_vectors[2];
+	int   * current_set;
+	// geometry buffers (struct gc / corr_vector / trot_deriv in the TU)
+	void * gc; void * cv; void * rbr; void * rot;
+	// resident task slots + convert inputs
+	float * ftasks0; float * ftasks1;
+	float ** kernel_offsets; float ** kernels; float ** images;
+	float ** clt; float ** clt_ref;
+	size_t dstride_img;        // floats
+	int * active; int * num_active;
+	// consolidate scratch (host prepare_consolidate provisioned)
+	float * cons_td_sens; float * cons_td_avg;
+	float * cons_pairs; float * cons_tasks_out;
+	int * cons_counters;       // [0] pairs, [1] surviving tiles, [2] misaligned
+	size_t slice;              // floats per cam CLT slice
+	// correlation / peak (resident)
+	int * corr_indices; int * num_corr_tiles;
+	size_t dstride_corr_td; float * corrs_td;
+	size_t dstride_corr; float * corrs;
+	const float * peak_debias; float * peaks;
+	int * status;              // [0] rc (corr rows or stage-tagged negative),
+	                           // [1] pairs, [2] surviving, [3] misaligned, [4] corr rows
+} PoseMeasureChain;
+#define POSE_MEASURE_STATUS_INTS 5
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,