CLAUDE: DP rung D3 - full per-scene parent (K x measure+prepare+step+commit, one entry)

pose_scene_dp: ONE device entry per scene. Self-chaining launchers (D1
pattern): pose_measure_dp (D2, payload unchanged; its tail launcher now
chains onward via a new PoseMeasureChain.scene back-pointer) ->
pose_scene_dp_lma (per-cycle prepare from LIVE corr-row count; full mode
= production per-cycle conditioning re-derivation with clears/assemble/
norm/fx/finish, light mode = the D1 frozen body) -> pose_scene_dp_step
(full-mode pure_weight read LIVE from prep_result - it changes every
cycle; the six validated spike stages) -> pose_scene_dp_commit (anchor
update on accept, 25-float trace + 5-int measure-stats trace rows,
self-chains the next cycle's measure). Fixed-K, no device early exit;
guard/invalid-prepare failures stop the chain fail-safe with
cycles-completed visible to the host.

Wrapper tp_proc_exec_pose_scene_dp: C2 nullable uploads + per-scene
policy/pull/reg/anchor, full provisioning (D2 measure provision + the
prepare/step scratch blocks), B4 bayer fence, ONE launch, ONE
end-of-scene readback (scene status + measure status + traces + final
packed + final anchor + compact prep); peaks/indices stay resident for
the unchanged fetch calls. D2 entry refactored onto shared
pose_measure_provision/pose_mchain_fill (same bytes).

GATE (cases.list scene_dp): test_pose_scene_dp_jna rides the captured
pose_corr case at the production 5120/150 LMA shape (full-grid centers
scattered by task txy - corr indices carry full-grid tile numbers).
Four chains vs the host-driven B3 loop (measure -> prepare_resident ->
resident step, host anchor bookkeeping): full accept K=3, full reject
K=2, full post-reject K=2, light K=2 riding the conditioning frozen by
the preceding full chain. Per-cycle trace rows, measure stats, final
packed/anchor/prep, and last-cycle corr rows keyed by packed index ALL
bit-identical; the measure<->LMA interleave is real (accepts move the
pose the next cycle measures at). run_cases.sh ALL PASS incl. pose_corr
@tol 0 + dp_cycles + measure_dp; DP spike PASS; sanitizer-12.8 memcheck
0 errors. Native-test timing is A==B by construction (no JNA in either
arm); the production win - killing the per-cycle Java/JNA round trips -
lands at rung D4.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

cases.list.tmpl

@@ -67,6 +67,7 @@ case avg_td_oob      test_avg_td_oob       ${REPO}/testdata/synth4/avg_td_oob
 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 scene_dp        test_pose_scene_dp_jna      ${MODEL}/testdata/pose_corr     --tol 0     # rung D3: full per-scene DP parent bit-exact vs host-driven B3 loop (full+light chains) # By Claude on 07/17/2026
 #case 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_scene_dp_jna.cu

+// test_pose_scene_dp_jna.cu - DP rung D3 gate: the FULL per-scene parent
+// (tp_proc_exec_pose_scene_dp - ONE device entry running num_cycles x
+// [D2 measure chain -> prepare -> resident step -> commit]) must be
+// BIT-IDENTICAL to the host-driven rung-B3 per-scene loop
+// (tp_proc_exec_pose_measure_cycle -> tp_proc_exec_pose_lma_prepare_resident
+//  -> tp_proc_exec_pose_lma_resident_step, with host anchor bookkeeping).
+//
+// Two TpProc instances from one NVRTC module (the chain mutates resident CLT/
+// corr/pose state; the D1/D2/DP-spike two-instance reasoning). Real content
+// comes from the captured pose_corr case (images, kernels, geometry, LPF,
+// resident center TD, de-bias); pose/camera-meta/task template are synthetic-
+// but-in-frame with template txy/centers extracted from the captured
+// production stream (the D2 precedent: bit-exactness is between A and B on
+// the SAME bytes). Unlike D1/D2 this exercises the full measure<->LMA
+// interleave: each accepted step moves the pose the NEXT measure cycle runs
+// at, so peaks/conditioning/pure_weight all change per cycle.
+//
+// Chains: full-prepare accept run (production pose_freeze_cycle=0 policy -
+// per-cycle conditioning re-derivation, live pure_weight), reject run
+// (near-infinite lambda; anchor/set retained), post-reject accept run, and a
+// LIGHT run (prepare_light=1, the D1 frozen-conditioning body, conditioning
+// frozen by the preceding full chain). Per chain the per-cycle 25-float trace
+// rows, the 5-int measure status rows, the final packed result, the final
+// anchor, the compact prep result, and the LAST-cycle resident corr rows
+// (raw-TD/PD/peak keyed by packed index - the recorded order-independence
+// rule) must all match bit-exactly. 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_scene_dp_jna.cu src/tests/tp_test_data.cu \
+//     -o tests_bin/test_pose_scene_dp_jna -L jna -ltileproc \
+//     -Xlinker -rpath -Xlinker $PWD/jna
+// Run: tests_bin/test_pose_scene_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_lma_prepare_resident(void*,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,int,int,int,int,int,
+      float,float,float,int,const float*,const float*,int,float*,int,
+      float*,float*,float*);
+  int   tp_proc_exec_pose_lma_resident_step(
+      void*,float,const float*,const float*,const float*,const float*,float,float,
+      int,int,float*,float*,float*);
+  int   tp_proc_exec_pose_scene_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,float,float,float,int,int,int,int,int,float,float,float,int,int,
+      const float*,const float*,const float*,const float*,
+      float*,int*,float*,float*,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, RESULT=25, MSTATS=5, NUM_PARAMS=3, NUM_COMPONENTS=2,
+       PREP_FLOATS=8, PREP_PURE_WEIGHT=3 };
+enum { CORR_SLOT=0xff, NTILE_SHIFT=8, USE_EIGEN=1, SAME_WEIGHTS=0 };
+static const float EIG_MIN_SQRT=0.4f, EIG_MAX_SQRT=4.0f, MIN_CONFIDENCE=0.0f;
+static const float RMS_DIFF=0.001f;
+
+static bool rowEqual(const float* a, const float* b, int n){
+  return std::memcmp(a,b,(size_t)n*sizeof(float))==0;
+}
+
+static int comparePacked(const char* tag,const float* a,const float* b){
+  int diffs=0;
+  for(int i=0;i<RESULT;i++){
+    if(std::memcmp(a+i,b+i,sizeof(float))){
+      unsigned ba,bb; std::memcpy(&ba,a+i,4); std::memcpy(&bb,b+i,4);
+      std::printf("%s packed[%d] host=%.9g (0x%08x) dp=%.9g (0x%08x)\n",
+          tag,i,a[i],ba,b[i],bb);
+      diffs++;
+    }
+  }
+  return diffs;
+}
+
+// LAST-cycle resident corr rows, ORDER-INDEPENDENT KEYED BY PACKED INDEX (the
+// recorded rung-1/B3/D2 rule: correlate2D_inter row order is atomicAdd-
+// nondeterministic; per-index content is the contract).
+static int compareResidentRows(const char* tag, void* pa, void* pb, int n, int corr_rad){
+  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)};
+  void* procs2[2]={pa,pb};
+  for (int k = 0; k < 2; k++){
+    if (tp_proc_get_corr_indices(procs2[k], idx[k].data(), n) < 0 ||
+        tp_proc_get_corr_td(procs2[k], td[k].data()) < 0 ||
+        tp_proc_get_corr2d(procs2[k], pd[k].data(), corr_rad) < 0 ||
+        tp_proc_get_peaks(procs2[k], pk[k].data()) < 0){
+      std::printf("%s: fetch[%d] failed: %s -> FAIL\n", tag, k, tp_last_error()); return 1; }
+  }
+  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 fail = (missing>0)||(dup>0)||(td_mm>0)||(pd_mm>0)||(pk_mm>0)||(nvalid<=0);
+  std::printf("%s: %d last-cycle corr rows by packed index: missing=%ld dup=%ld TD-mm=%ld PD-mm=%ld peak-mm=%ld valid=%d -> %s\n",
+         tag, n, missing, dup, td_mm, pd_mm, pk_mm, nvalid, fail?"FAIL":"PASS");
+  return fail;
+}
+
+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 (D2 rule)
+  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;
+  }
+  // LMA tiles = the FULL grid (the production 5120/150 shape: corr indices
+  // carry full-grid tile numbers; centers are scattered by task txy and the
+  // assemble-built selection masks the unmeasured tiles).
+  const int tilesx = img_width/8, tilesy = img_height/8;
+  const int numTiles = tilesx*tilesy;
+  std::vector<float> lma_centers(3*(size_t)numTiles, 0.0f);
+  for (int i = 0; i < num_tasks; i++){
+    const int tx = task_map[i] & 0xffff, ty = (task_map[i] >> 16) & 0xffff;
+    const int tile = ty*tilesx + tx;
+    if (tile < 0 || tile >= numTiles){ printf("FAIL: task %d txy out of grid\n", i); return EXIT_FAILURE; }
+    lma_centers[3*(size_t)tile+0] = task_centers[3*i+0];
+    lma_centers[3*(size_t)tile+1] = task_centers[3*i+1];
+    lma_centers[3*(size_t)tile+2] = 0.0f;
+  }
+  printf("cams=%d frame %dx%d tasks=%d task_size=%d codes 0x%x/0x%x sel=0x%x lma_tiles=%d\n",
+         num_cams,img_width,img_height,num_tasks,task_size,task_code_main,task_code_sub,sel_sensors,numTiles);
+
+  // synthetic-but-real-dims camera pair (D2 precedent; production-scale rbr)
+  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;
+  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 = host-driven B3 per-scene loop, B = D3 per-scene DP parent
+  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"));
+    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;
+  }
+
+  // LMA policy: initial pose with a real grid-visible error (steps must move),
+  // pull rows near it, uniform regularization (test policy, identical on A/B).
+  const float v0[NUM_PARAMS]={0.0020f,-0.0010f,0.0005f};
+  const float pull[NUM_PARAMS]={v0[0]+0.0004f,v0[1]-0.0002f,v0[2]+0.0001f};
+  const float regw[NUM_PARAMS]={0.08f,0.08f,0.08f};
+
+  const struct { int cycles; float lambda; int light; const char* tag; } chains[4]={
+    {3,0.0005f,0,"chain1(full,accept)"},
+    {2,1.0e6f, 0,"chain2(full,reject)"},
+    {2,0.0005f,0,"chain3(full,post-reject)"},
+    {2,0.0005f,1,"chain4(light)"}};
+  float anchorA[NUM_PARAMS]={v0[0],v0[1],v0[2]};
+  float anchorB[NUM_PARAMS]={v0[0],v0[1],v0[2]};
+  float pwFrozenA=0.0f, pwFrozenB=0.0f; // last FULL prepare's pure_weight (light chains)
+  int bad=0, accepts=0, rejects=0;
+  bool firstA=true, firstB=true;
+  double ms_a=0, ms_b=0; int cyc_timed=0; // chains 2+ (first-call uploads excluded)
+  for (int chain=0; chain<4 && !bad; chain++){
+    const int K=chains[chain].cycles;
+    const float lambda=chains[chain].lambda;
+    const int light=chains[chain].light;
+    // ---- A: host-driven B3 per-scene loop (the production sequence) ----
+    std::vector<float> packedA((size_t)K*RESULT);
+    std::vector<int>   mstatsA((size_t)K*MSTATS);
+    float prepA[PREP_FLOATS]={0};
+    const auto tA0=std::chrono::steady_clock::now();
+    for (int cyc=0; cyc<K && !bad; cyc++){
+      const float poseC[12]={0,0,0, 0,0,0, 0,0,0, anchorA[0],anchorA[1],anchorA[2]};
+      int statsA[3]={-1,-1,-1};
+      const int npk = tp_proc_exec_pose_measure_cycle(procs[0],
+          firstA?meta:nullptr, firstA?radial:nullptr, firstA?rbr.data():nullptr, firstA?rbr_len:0,
+          firstA?ers.data():nullptr, firstA?ers_len:0,
+          firstA?meta:nullptr, firstA?radial:nullptr, firstA?rbr.data():nullptr, firstA?rbr_len:0,
+          firstA?ers.data():nullptr, firstA?ers_len:0,
+          poseC, firstA?task_map.data():nullptr, firstA?task_centers.data():nullptr, statsA);
+      if (npk <= 0){
+        printf("%s A measure cyc%d rc=%d: %s -> FAIL\n", chains[chain].tag, cyc, npk, tp_last_error());
+        return EXIT_FAILURE; }
+      mstatsA[(size_t)cyc*MSTATS+0]=npk;
+      mstatsA[(size_t)cyc*MSTATS+1]=statsA[0];
+      mstatsA[(size_t)cyc*MSTATS+2]=statsA[1];
+      mstatsA[(size_t)cyc*MSTATS+3]=statsA[2];
+      mstatsA[(size_t)cyc*MSTATS+4]=npk;
+      float prep[PREP_FLOATS]={0};
+      int rc = tp_proc_exec_pose_lma_prepare_resident(procs[0],
+          nullptr,nullptr,nullptr,0,nullptr,0,
+          nullptr,nullptr,nullptr,0,nullptr,0,
+          poseC, firstA?lma_centers.data():nullptr, numTiles, NUM_COMPONENTS,
+          CORR_SLOT, NTILE_SHIFT, USE_EIGEN, EIG_MIN_SQRT, EIG_MAX_SQRT,
+          MIN_CONFIDENCE, SAME_WEIGHTS,
+          light?nullptr:pull, light?nullptr:regw, light,
+          light?nullptr:prep, 0, nullptr,nullptr,nullptr);
+      firstA=false;
+      if (rc){
+        printf("%s A prepare cyc%d rc=%d: %s -> FAIL\n", chains[chain].tag, cyc, rc, tp_last_error());
+        return EXIT_FAILURE; }
+      float pw = light ? pwFrozenA : prep[PREP_PURE_WEIGHT];
+      if (!light){ std::memcpy(prepA,prep,sizeof(prepA)); pwFrozenA=pw; }
+      rc = tp_proc_exec_pose_lma_resident_step(procs[0], lambda,
+          nullptr,nullptr,nullptr, anchorA, pw, RMS_DIFF, numTiles, 0,
+          &packedA[(size_t)cyc*RESULT], nullptr, nullptr);
+      if (rc){
+        printf("%s A step cyc%d rc=%d: %s -> FAIL\n", chains[chain].tag, cyc, rc, tp_last_error());
+        return EXIT_FAILURE; }
+      if (packedA[(size_t)cyc*RESULT+23]==1.0f){
+        for(int par=0;par<NUM_PARAMS;par++) anchorA[par]=packedA[(size_t)cyc*RESULT+16+par];
+        accepts++;
+      } else rejects++;
+    }
+    const auto tA1=std::chrono::steady_clock::now();
+    // ---- B: ONE D3 entry for the whole scene ----
+    const float poseK[12]={0,0,0, 0,0,0, 0,0,0, anchorB[0],anchorB[1],anchorB[2]};
+    std::vector<float> trace((size_t)K*RESULT,0.0f);
+    std::vector<int>   mstatsB((size_t)K*MSTATS,-1);
+    float packedB[RESULT], vectorB[NUM_PARAMS], prepB[PREP_FLOATS]={0};
+    int statsB[3]={-1,-1,-1};
+    const auto tB0=std::chrono::steady_clock::now();
+    const int rcB = tp_proc_exec_pose_scene_dp(procs[1],
+        firstB?meta:nullptr, firstB?radial:nullptr, firstB?rbr.data():nullptr, firstB?rbr_len:0,
+        firstB?ers.data():nullptr, firstB?ers_len:0,
+        firstB?meta:nullptr, firstB?radial:nullptr, firstB?rbr.data():nullptr, firstB?rbr_len:0,
+        firstB?ers.data():nullptr, firstB?ers_len:0,
+        poseK, firstB?task_map.data():nullptr, firstB?task_centers.data():nullptr,
+        K, lambda, light?pwFrozenB:0.0f, RMS_DIFF,
+        numTiles, CORR_SLOT, NTILE_SHIFT, NUM_COMPONENTS, USE_EIGEN,
+        EIG_MIN_SQRT, EIG_MAX_SQRT, MIN_CONFIDENCE, SAME_WEIGHTS,
+        light, firstB?lma_centers.data():nullptr,
+        light?nullptr:pull, light?nullptr:regw, anchorB,
+        trace.data(), mstatsB.data(), packedB, vectorB,
+        light?nullptr:prepB, statsB);
+    const auto tB1=std::chrono::steady_clock::now();
+    if (!firstB && chain>0){
+      ms_a += std::chrono::duration<double,std::milli>(tA1-tA0).count();
+      ms_b += std::chrono::duration<double,std::milli>(tB1-tB0).count();
+      cyc_timed += K;
+    }
+    firstB=false;
+    if (rcB <= 0){
+      printf("%s B scene_dp rc=%d: %s -> FAIL\n", chains[chain].tag, rcB, tp_last_error());
+      return EXIT_FAILURE; }
+    if (!light) pwFrozenB=prepB[PREP_PURE_WEIGHT];
+    // ---- compare ----
+    for (int cyc=0; cyc<K; cyc++){
+      char tag[64]; std::snprintf(tag,sizeof(tag),"%s cyc%d",chains[chain].tag,cyc);
+      if (comparePacked(tag,&packedA[(size_t)cyc*RESULT],&trace[(size_t)cyc*RESULT])) bad++;
+      if (std::memcmp(&mstatsA[(size_t)cyc*MSTATS],&mstatsB[(size_t)cyc*MSTATS],MSTATS*sizeof(int))){
+        printf("%s measure stats host={%d,%d,%d,%d,%d} dp={%d,%d,%d,%d,%d}\n",tag,
+            mstatsA[(size_t)cyc*MSTATS+0],mstatsA[(size_t)cyc*MSTATS+1],mstatsA[(size_t)cyc*MSTATS+2],
+            mstatsA[(size_t)cyc*MSTATS+3],mstatsA[(size_t)cyc*MSTATS+4],
+            mstatsB[(size_t)cyc*MSTATS+0],mstatsB[(size_t)cyc*MSTATS+1],mstatsB[(size_t)cyc*MSTATS+2],
+            mstatsB[(size_t)cyc*MSTATS+3],mstatsB[(size_t)cyc*MSTATS+4]); bad++; }
+    }
+    if (comparePacked((std::string(chains[chain].tag)+" final").c_str(),
+        &packedA[(size_t)(K-1)*RESULT],packedB)) bad++;
+    // host bookkeeping already advanced anchorA; the DP arm's device anchor is
+    // vectorB - both must match bit-exactly
+    if (std::memcmp(anchorA,vectorB,sizeof(anchorA))){
+      printf("%s final anchor host={%.9g,%.9g,%.9g} dp={%.9g,%.9g,%.9g}\n",
+          chains[chain].tag,anchorA[0],anchorA[1],anchorA[2],
+          vectorB[0],vectorB[1],vectorB[2]); bad++; }
+    for(int par=0;par<NUM_PARAMS;par++) anchorB[par]=vectorB[par];
+    if (!light && std::memcmp(prepA,prepB,sizeof(prepA))){
+      printf("%s prep result differs (host pw=%.9g dp pw=%.9g)\n",
+          chains[chain].tag,prepA[PREP_PURE_WEIGHT],prepB[PREP_PURE_WEIGHT]); bad++; }
+    if (mstatsA[(size_t)(K-1)*MSTATS+0]!=rcB){
+      printf("%s last-cycle rc host=%d dp=%d\n",chains[chain].tag,
+          mstatsA[(size_t)(K-1)*MSTATS+0],rcB); bad++; }
+    if (mstatsA[(size_t)(K-1)*MSTATS+2]<=0){
+      printf("%s: no surviving tiles - vacuous gate\n",chains[chain].tag); bad++; }
+    char rtag[80]; std::snprintf(rtag,sizeof(rtag),"%s rows",chains[chain].tag);
+    if (compareResidentRows(rtag,procs[0],procs[1],rcB,corr_rad)) bad++;
+    printf("%s K=%d lambda=%g light=%d: last accepted=%g rms %g -> %g : %s\n",
+        chains[chain].tag,K,lambda,light,packedB[23],packedB[19],packedB[21],
+        bad?"DIFF":"bit-identical");
+  }
+  if (!bad && (!accepts || !rejects)){
+    printf("chain coverage broken: accepts=%d rejects=%d (need both)\n",accepts,rejects); bad++;
+  }
+  if (cyc_timed) printf("microbench (chains 2-4, %d cycles, first-call uploads excluded): "
+      "host loop %.3f ms/cycle, D3 scene entry %.3f ms/cycle\n",
+      cyc_timed, ms_a/cyc_timed, ms_b/cyc_timed);
+  printf("D3 per-scene DP parent vs host-driven B3 loop -> %s\n",
+      bad?"FAIL":"PASS (bit-identical)");
+  printf("%s: %s\n", argv[0], bad ? "FAIL" : "PASS");
+  tp_proc_destroy(procs[0]); tp_proc_destroy(procs[1]); tp_destroy_module(m);
+  return bad ? EXIT_FAILURE : EXIT_SUCCESS;
+}

jna/tp_jna.cpp

@@ -711,6 +711,13 @@ struct TpProc {
   // rows - the one end-of-chain D2H). By Claude on 07/17/2026.
   PoseMeasureChain *gpu_pose_mchain;
   int *gpu_pose_mchain_status;
+  // DP rung D3: device copy of the per-scene parent descriptor, its 2-int
+  // status (cycles completed + LMA rc) and the per-cycle measure status trace
+  // ([num_cycles][POSE_MEASURE_STATUS_INTS]); the float result trace reuses
+  // gpu_pose_dp_trace. By Claude on 07/17/2026.
+  PoseSceneDp *gpu_pose_scene;
+  int *gpu_pose_scene_status;
+  int *gpu_pose_scene_mstats; int pose_scene_mstats_capacity;
   // 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,
@@ -805,6 +812,8 @@ TpProc* tp_proc_create(TpModule* m){
   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_scene=nullptr; p->gpu_pose_scene_status=nullptr; // By Claude on 07/17/2026 (rung D3)
+  p->gpu_pose_scene_mstats=nullptr; p->pose_scene_mstats_capacity=0; // By Claude on 07/17/2026 (rung D3)
   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)
@@ -1832,31 +1841,12 @@ int tp_proc_exec_pose_measure_cycle(TpProc* p,
 // 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; }
+// Shared D2/D3 provisioning: the reverse-distortion table, consolidate
+// scratch, worst-case peak rows and the measure-chain descriptor/status
+// buffers - everything the device chain assumes resident (no device-side
+// allocation). Factored from the rung-D2 entry, same bytes/messages.
+// By Claude on 07/17/2026 (rung D3).
+static int pose_measure_provision(TpProc* p){
   // 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){
@@ -1866,7 +1856,7 @@ int tp_proc_exec_pose_measure_cycle_dp(TpProc* p,
     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;
+  int 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
@@ -1882,14 +1872,21 @@ int tp_proc_exec_pose_measure_cycle_dp(TpProc* p,
       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;
+  return 0;
+}
+
+// Shared D2/D3 descriptor fill from provisioned TpProc state (the union of
+// the per-cycle stage arguments + resident pointers). scene defaults to NULL
+// (standalone D2); the D3 entry points it at the per-scene parent descriptor.
+// By Claude on 07/17/2026 (rung D3 refactor of the D2 fill).
+static void pose_mchain_fill(TpProc* p, PoseMeasureChain* ph){
+  PoseMeasureChain& h=*ph;
   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;
@@ -1929,6 +1926,38 @@ int tp_proc_exec_pose_measure_cycle_dp(TpProc* p,
   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;
+  h.scene=nullptr; // standalone D2; the D3 entry overrides // 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; }
+  rc=pose_measure_provision(p); if(rc) return rc;
+  bayer_fence(p); // staged uploads must land before convert reads images (rung B4)
+  PoseMeasureChain h;
+  pose_mchain_fill(p,&h);
   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");
@@ -2380,6 +2409,267 @@ int tp_proc_exec_pose_lma_dp_cycles(TpProc* p,
   return 0;
 }
 
+// DP rung D3 (ladder ratified 2026-07-16): the FULL per-scene solve as ONE
+// device entry - num_cycles x [D2 measure chain -> prepare -> resident step ->
+// commit], self-chained launcher-to-launcher with zero interior host round
+// trips. prepare_light=0 replicates production per-cycle FULL conditioning
+// re-derivation (pose_freeze_cycle=0 after the C3/C3b falsification; the
+// per-cycle pure_weight is read LIVE from prep_result on-device);
+// prepare_light=1 is the D1 frozen-conditioning body (DP-era experiments).
+// Host keeps: the rung-C2 nullable uploads (camera blocks / task template /
+// LMA centers), per-scene policy scalars + pull/reg rows + entry pose/anchor,
+// provisioning, the B4 bayer fence, ONE launch and ONE end-of-scene readback
+// (scene status + measure status + per-cycle trace/mstats + final packed
+// result + final anchor + compact prep). Peaks/corr indices stay resident for
+// the unchanged fetch calls (last-cycle peaks). Returns the LAST cycle's
+// corr-row count (>0) or a negative (host -1xx / device guard codes; a
+// mid-scene device failure reports the failing stage's code with
+// cycles-completed in the message - the D4 FAIL-safe hook).
+// By Claude on 07/17/2026.
+int tp_proc_exec_pose_scene_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 num_cycles,
+                               float lambda,
+                               float pure_weight,   // prepare_light=1 only (frozen)
+                               float rms_diff,
+                               int num_tiles,
+                               int corr_slot, int ntile_shift,
+                               int num_components, int use_eigen,
+                               float eig_min_sqrt, float eig_max_sqrt,
+                               float min_confidence, int same_weights,
+                               int prepare_light,
+                               const float* centers,          // NULL = keep resident (C2)
+                               const float* pull_targets,     // full mode; NULL = keep resident
+                               const float* param_regweights, // full mode; NULL = keep resident
+                               const float* vector,           // entry anchor [3]
+                               float* trace,       // [num_cycles*25] or NULL
+                               int*   mstats,      // [num_cycles*5] or NULL
+                               float* result,      // final packed [25]
+                               float* vector_out,  // final anchor [3] or NULL
+                               float* prep_out,    // last full prep [8] or NULL
+                               int*   stats){      // last-cycle {pairs,surviving,misaligned} or NULL
+  if(!p||!vector||!result||(num_cycles<=0)||(num_tiles<=0)||
+     (num_components!=POSE_NUM_COMPONENTS)||!std::isfinite(rms_diff)||
+     (prepare_light&&(!std::isfinite(pure_weight)||!(pure_weight>0.0f)))){
+    seterr("exec_pose_scene_dp: bad args k=%d nt=%d nc=%d light=%d",
+           num_cycles,num_tiles,num_components,prepare_light); return -1; }
+  if(!p->pcyc_valid){ seterr("exec_pose_scene_dp: descriptor not set (tp_proc_set_pose_cycle)"); return -2; }
+  if(!p->have_corr){ seterr("exec_pose_scene_dp: no corr buffers (call setup_rbg_corr)"); return -2; }
+  if(!p->gc||!p->cv||!p->rot||!p->rbr){ seterr("exec_pose_scene_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_scene_dp: task_size %d != descriptor %d",task_size,p->pcyc_task_size); return -2; }
+  rc=pose_measure_provision(p); if(rc) return rc;
+  // LMA-side provisioning = exec_pose_lma_prepare_resident + resident_step
+  // blocks (same capacity bookkeeping, so mixing granular and DP entries on
+  // one instance never regrows resident state). By Claude on 07/17/2026.
+  #define POSE_ALLOC(field,count,type) do{ if(!(p->field) && \
+      cudaMalloc((void**)&(p->field),(size_t)(count)*sizeof(type))!=cudaSuccess){ \
+      seterr("exec_pose_scene_dp: alloc %s[%d] failed",#field,(int)(count)); return -4; } }while(0)
+  POSE_ALLOC(gpu_pose_lma_prep,POSE_LMA_PREP_RESULT+2*POSE_NUM_PARAMS,float);
+  #undef POSE_ALLOC
+  if(num_tiles>p->pose_tile_capacity){
+    p->pose_prep_ready&=~4; p->pose_prep_centers_tiles=0; // regrow frees resident centers
+    cudaFree(p->gpu_pose_centers); cudaFree(p->gpu_pose_selection); cudaFree(p->gpu_pose_fx);
+    cudaFree(p->gpu_pose_jt); cudaFree(p->gpu_pose_valid); cudaFree(p->gpu_pose_candidate_fx);
+    cudaFree(p->gpu_pose_candidate_jt); cudaFree(p->gpu_pose_candidate_valid);
+    p->gpu_pose_centers=p->gpu_pose_fx=p->gpu_pose_jt=nullptr;
+    p->gpu_pose_candidate_fx=p->gpu_pose_candidate_jt=nullptr;
+    p->gpu_pose_selection=p->gpu_pose_valid=p->gpu_pose_candidate_valid=nullptr;
+    p->pose_tile_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_centers,(size_t)3*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_selection,(size_t)num_tiles*sizeof(unsigned char))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_fx,(size_t)POSE_NUM_COMPONENTS*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_jt,(size_t)POSE_NUM_PARAMS*POSE_NUM_COMPONENTS*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_valid,(size_t)num_tiles*sizeof(unsigned char))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_candidate_fx,(size_t)POSE_NUM_COMPONENTS*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_candidate_jt,(size_t)POSE_NUM_PARAMS*POSE_NUM_COMPONENTS*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_candidate_valid,(size_t)num_tiles*sizeof(unsigned char))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: tile scratch alloc failed n=%d",num_tiles); return -4; }
+    p->pose_tile_capacity=num_tiles;
+  }
+  const int num_values=POSE_NUM_COMPONENTS*num_tiles+POSE_NUM_PARAMS;
+  if(num_values>p->pose_lma_value_capacity){
+    cudaFree(p->gpu_pose_lma_weights); cudaFree(p->gpu_pose_lma_jt);
+    cudaFree(p->gpu_pose_lma_ymfx); cudaFree(p->gpu_pose_lma_y); cudaFree(p->gpu_pose_lma_vector);
+    cudaFree(p->gpu_pose_lma_result); cudaFree(p->gpu_pose_lma_candidate_jt);
+    cudaFree(p->gpu_pose_lma_candidate_ymfx);
+    p->gpu_pose_lma_weights=p->gpu_pose_lma_jt=p->gpu_pose_lma_ymfx=nullptr;
+    p->gpu_pose_lma_y=p->gpu_pose_lma_vector=p->gpu_pose_lma_result=nullptr;
+    p->gpu_pose_lma_candidate_jt=p->gpu_pose_lma_candidate_ymfx=nullptr;
+    p->pose_lma_value_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_lma_weights,(size_t)num_values*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_jt,(size_t)POSE_NUM_PARAMS*num_values*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_ymfx,(size_t)num_values*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_y,(size_t)num_values*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_vector,POSE_NUM_PARAMS*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_result,POSE_LMA_RESIDENT_RESULT*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_candidate_jt,(size_t)POSE_NUM_PARAMS*num_values*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_candidate_ymfx,(size_t)num_values*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: value scratch alloc failed nv=%d",num_values); return -4; }
+    p->pose_lma_value_capacity=num_values;
+  }
+  if(num_tiles>p->pose_lma_tile_capacity){
+    cudaFree(p->gpu_pose_lma_eigen); p->gpu_pose_lma_eigen=nullptr; p->pose_lma_tile_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_lma_eigen,(size_t)4*num_tiles*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: eigen scratch alloc failed"); return -4; }
+    p->pose_lma_tile_capacity=num_tiles;
+  }
+  const int num_partials=(num_tiles+POSE_LMA_REDUCE_THREADS-1)/POSE_LMA_REDUCE_THREADS;
+  if(num_partials>p->pose_lma_partial_capacity){
+    cudaFree(p->gpu_pose_lma_partials); cudaFree(p->gpu_pose_lma_candidate_partials);
+    p->gpu_pose_lma_partials=p->gpu_pose_lma_candidate_partials=nullptr;
+    p->pose_lma_partial_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_lma_partials,
+                  (size_t)POSE_LMA_PARTIAL_VALUES*num_partials*sizeof(float))!=cudaSuccess ||
+       cudaMalloc((void**)&p->gpu_pose_lma_candidate_partials,
+                  (size_t)POSE_LMA_PARTIAL_VALUES*num_partials*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: partial scratch alloc failed np=%d",num_partials); return -4; }
+    p->pose_lma_partial_capacity=num_partials;
+  }
+  // Per-scene uploads (rung-C2 register semantics): LMA centers (per-sequence),
+  // pull/reg rows (full mode; part of the objective), the entry anchor.
+  float* gpu_pull=p->gpu_pose_lma_prep+POSE_LMA_PREP_RESULT;
+  float* gpu_regw=gpu_pull+POSE_NUM_PARAMS;
+  if(centers){
+    if(cudaMemcpy(p->gpu_pose_centers,centers,(size_t)3*num_tiles*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+      seterr("exec_pose_scene_dp: centers HtoD failed"); return -5; }
+    p->pose_prep_ready|=4; p->pose_prep_centers_tiles=num_tiles;
+  } else if(!(p->pose_prep_ready&4)||(p->pose_prep_centers_tiles!=num_tiles)){
+    seterr("exec_pose_scene_dp: resident centers missing (%d/%d)",
+           p->pose_prep_centers_tiles,num_tiles); return -2; }
+  if(!prepare_light){
+    if((pull_targets&&
+        cudaMemcpy(gpu_pull,pull_targets,POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess) ||
+       (param_regweights&&
+        cudaMemcpy(gpu_regw,param_regweights,POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess)){
+      seterr("exec_pose_scene_dp: pull/reg HtoD failed"); return -5; }
+  } else if(p->pose_prep_frozen_tiles!=num_tiles){
+    // light cycles reuse the resident conditioning from the last FULL prepare
+    // (host granular or a previous full-mode scene entry) - rung C3 contract
+    seterr("exec_pose_scene_dp: light without frozen conditioning (%d/%d)",
+           p->pose_prep_frozen_tiles,num_tiles); return -2; }
+  if(cudaMemcpy(p->gpu_pose_lma_vector,vector,POSE_NUM_PARAMS*sizeof(float),
+                cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_scene_dp: anchor HtoD failed"); return -5; }
+  // trace / measure-stats / descriptor / status buffers
+  if(num_cycles>p->pose_dp_trace_capacity){
+    cudaFree(p->gpu_pose_dp_trace); p->gpu_pose_dp_trace=nullptr; p->pose_dp_trace_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_dp_trace,
+                  (size_t)num_cycles*POSE_LMA_RESIDENT_RESULT*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: trace alloc failed k=%d",num_cycles); return -4; }
+    p->pose_dp_trace_capacity=num_cycles;
+  }
+  if(num_cycles>p->pose_scene_mstats_capacity){
+    cudaFree(p->gpu_pose_scene_mstats); p->gpu_pose_scene_mstats=nullptr; p->pose_scene_mstats_capacity=0;
+    if(cudaMalloc((void**)&p->gpu_pose_scene_mstats,
+                  (size_t)num_cycles*POSE_MEASURE_STATUS_INTS*sizeof(int))!=cudaSuccess){
+      seterr("exec_pose_scene_dp: mstats alloc failed k=%d",num_cycles); return -4; }
+    p->pose_scene_mstats_capacity=num_cycles;
+  }
+  if(!p->gpu_pose_scene &&
+     cudaMalloc((void**)&p->gpu_pose_scene,sizeof(PoseSceneDp))!=cudaSuccess){
+    seterr("exec_pose_scene_dp: scene descriptor alloc failed"); return -4; }
+  if(!p->gpu_pose_scene_status &&
+     cudaMalloc((void**)&p->gpu_pose_scene_status,POSE_SCENE_STATUS_INTS*sizeof(int))!=cudaSuccess){
+    seterr("exec_pose_scene_dp: scene status alloc failed"); return -4; }
+  bayer_fence(p); // staged uploads must land before convert reads images (rung B4)
+  PoseMeasureChain hm;
+  pose_mchain_fill(p,&hm);
+  hm.scene=p->gpu_pose_scene; // the tail launcher chains the per-scene LMA continuation
+  if(cudaMemcpy(p->gpu_pose_mchain,&hm,sizeof(PoseMeasureChain),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_scene_dp: measure descriptor HtoD failed"); return -4; }
+  PoseSceneDp hs;
+  hs.num_cycles=num_cycles; hs.cycle=0; hs.prepare_light=prepare_light;
+  hs.num_tiles=num_tiles; hs.num_partials=num_partials;
+  hs.lambda=lambda; hs.pure_weight=pure_weight; hs.rms_diff=rms_diff;
+  hs.corr_slot=corr_slot; hs.ntile_shift=ntile_shift;
+  hs.num_components=num_components; hs.use_eigen=use_eigen;
+  hs.eig_min_sqrt=eig_min_sqrt; hs.eig_max_sqrt=eig_max_sqrt;
+  hs.min_confidence=min_confidence; hs.same_weights=same_weights;
+  hs.mc=p->gpu_pose_mchain;
+  hs.corr_indices=p->gpu_corr_indices; hs.peaks=p->gpu_peaks;
+  hs.num_corr_tiles=p->gpu_num_corr_tiles;
+  hs.cameras=p->gpu_pose_cameras;
+  hs.ref_radial=p->gpu_pose_radial[0];  hs.ref_rbr=p->gpu_pose_rbr[0];   hs.ref_ers=p->gpu_pose_ers[0];
+  hs.scene_radial=p->gpu_pose_radial[1];hs.scene_rbr=p->gpu_pose_rbr[1]; hs.scene_ers=p->gpu_pose_ers[1];
+  hs.centers=p->gpu_pose_centers; hs.selection=p->gpu_pose_selection;
+  hs.set_vectors[0]=p->gpu_pose_vectors;           hs.set_vectors[1]=p->gpu_pose_candidate_vectors;
+  hs.set_fx[0]=p->gpu_pose_fx;                     hs.set_fx[1]=p->gpu_pose_candidate_fx;
+  hs.set_jt[0]=p->gpu_pose_jt;                     hs.set_jt[1]=p->gpu_pose_candidate_jt;
+  hs.set_valid[0]=p->gpu_pose_valid;               hs.set_valid[1]=p->gpu_pose_candidate_valid;
+  hs.y=p->gpu_pose_lma_y; hs.weights=p->gpu_pose_lma_weights; hs.eigen=p->gpu_pose_lma_eigen;
+  hs.vector=p->gpu_pose_lma_vector;
+  hs.pull=gpu_pull; hs.regweights=gpu_regw;
+  hs.jt_out=p->gpu_pose_lma_jt; hs.ymfx_out=p->gpu_pose_lma_ymfx; hs.partials=p->gpu_pose_lma_partials;
+  hs.cand_jt_out=p->gpu_pose_lma_candidate_jt; hs.cand_ymfx_out=p->gpu_pose_lma_candidate_ymfx;
+  hs.cand_partials=p->gpu_pose_lma_candidate_partials;
+  hs.result=p->gpu_pose_lma_result; hs.prep_result=p->gpu_pose_lma_prep;
+  hs.current_set=p->gpu_pose_current_set;
+  hs.trace=trace?p->gpu_pose_dp_trace:nullptr;
+  hs.mstats=mstats?p->gpu_pose_scene_mstats:nullptr;
+  hs.status=p->gpu_pose_scene_status;
+  if(cudaMemcpy(p->gpu_pose_scene,&hs,sizeof(PoseSceneDp),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_scene_dp: scene descriptor HtoD failed"); return -4; }
+  CUfunction dp=getfun(p->mod,"pose_scene_dp");
+  if(!dp){ seterr("pose_scene_dp missing (DP kernel not in module)"); return -6; }
+  void* da[]={ &p->gpu_pose_scene };
+  if(launch1(dp,1,1,1,1,1,1,da,"pose_scene_dp")) return -7;
+  // ONE end-of-scene readback
+  int sst[POSE_SCENE_STATUS_INTS]={0,0};
+  int mst[POSE_MEASURE_STATUS_INTS]={0,0,0,0,0};
+  if(cudaMemcpy(sst,p->gpu_pose_scene_status,sizeof(sst),cudaMemcpyDeviceToHost)!=cudaSuccess ||
+     cudaMemcpy(mst,p->gpu_pose_mchain_status,sizeof(mst),cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_scene_dp: status DtoH failed"); return -8; }
+  if(stats){ stats[0]=mst[1]; stats[1]=mst[2]; stats[2]=mst[3]; }
+  p->last_num_corr_tiles=mst[4]; // fetch calls (indices/TD/peaks) read this mirror
+  if((trace &&
+      cudaMemcpy(trace,p->gpu_pose_dp_trace,
+                 (size_t)num_cycles*POSE_LMA_RESIDENT_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess) ||
+     (mstats &&
+      cudaMemcpy(mstats,p->gpu_pose_scene_mstats,
+                 (size_t)num_cycles*POSE_MEASURE_STATUS_INTS*sizeof(int),cudaMemcpyDeviceToHost)!=cudaSuccess) ||
+     cudaMemcpy(result,p->gpu_pose_lma_result,
+                POSE_LMA_RESIDENT_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess ||
+     (vector_out &&
+      cudaMemcpy(vector_out,p->gpu_pose_lma_vector,
+                 POSE_NUM_PARAMS*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess) ||
+     (prep_out &&
+      cudaMemcpy(prep_out,p->gpu_pose_lma_prep,
+                 POSE_LMA_PREP_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess)){
+    seterr("exec_pose_scene_dp: DtoH failed"); return -8; }
+  if(cudaMemcpy(&p->pose_current_set,p->gpu_pose_current_set,sizeof(int),
+                cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_scene_dp: current_set DtoH failed"); return -8; }
+  if((sst[0]!=num_cycles)||(sst[1]!=0)||(mst[0]<=0)){
+    seterr("exec_pose_scene_dp: chain stopped at cycle %d/%d (lma rc=%d, measure rc=%d pairs=%d out=%d corr=%d)",
+           sst[0],num_cycles,sst[1],mst[0],mst[1],mst[2],mst[4]);
+    if(sst[1]<0) return sst[1];
+    return (mst[0]<0)?mst[0]:-9;
+  }
+  // Resident state now matches a granular prepare+step tail: raw fx/J live at
+  // this tile count; full mode leaves the last cycle's conditioning resident
+  // (light-eligible, the rung-C3 contract).
+  p->pose_raw_num_tiles=num_tiles; p->pose_raw_has_jacobian=1;
+  if(!prepare_light) p->pose_prep_frozen_tiles=num_tiles;
+  return mst[0];
+}
+
 int tp_proc_exec_pose_lma_resident_step(TpProc* p,
                                         float lambda,
                                         const float* weights,
@@ -2803,6 +3093,7 @@ void tp_proc_destroy(TpProc* p){
   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_scene); cudaFree(p->gpu_pose_scene_status); cudaFree(p->gpu_pose_scene_mstats); // By Claude on 07/17/2026 (rung D3)
   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

@@ -1291,6 +1291,9 @@ extern "C" __global__ void pose_dp_fill_bits(
 	}
 }
 
+// fwd (rung D3): the measure tail chains the per-scene LMA continuation.
+extern "C" __global__ void pose_scene_dp_lma(PoseSceneDp * s);
+
 // 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)
@@ -1317,6 +1320,13 @@ extern "C" __global__ void pose_measure_dp_tail(const PoseMeasureChain * c)
 				c->peak_debias, c->peaks);
 	}
 	c->status[0] = nct; // npk == nct (exec_corr2d_peak return convention)
+	// Rung D3: per-scene parent continuation. Enqueued AFTER peak, so the
+	// LMA launcher sees LIVE peaks/counts. Guard failures above returned
+	// without chaining - the scene chain stops there (fail-safe; the host
+	// reads cycles-completed < num_cycles). By Claude on 07/17/2026.
+	if (c->scene != NULL) {
+		pose_scene_dp_lma<<<1, 1, 0, cudaStreamTailLaunch>>>((PoseSceneDp *) c->scene);
+	}
 }
 
 // Average launcher: reads the LIVE surviving-tile count, averages, copies the
@@ -1424,6 +1434,193 @@ extern "C" __global__ void pose_measure_dp(const PoseMeasureChain * c)
 	pose_measure_dp_consolidate<<<1, 1, 0, cudaStreamTailLaunch>>>(c);
 }
 
+// ── DP rung D3: full per-scene parent (ladder ratified 2026-07-16) ───────────
+// See tp_lma.h (PoseSceneDp). Per cycle the launchers replicate the host
+// per-cycle sequence EXACTLY: [the D2 measure chain, unchanged - its tail
+// chains pose_scene_dp_lma] -> [exec_pose_lma_prepare_resident: clears ->
+// assemble -> (norm) -> fx_jacobian(current) -> finish] -> [the six resident
+// step stages; the full-mode pure_weight is read LIVE from prep_result - it
+// changes every cycle under per-cycle conditioning re-derivation (C3/C3b
+// falsified freezing)] -> [commit: anchor update on accept + QC trace + the
+// NEXT cycle's measure chain]. Fixed-K, no device early exit (ruling a);
+// minimal per-cycle trace (ruling b). By Claude on 07/17/2026.
+
+// Byte fill (the cudaMemset of the selection bytes; pose_dp_fill_bits on
+// ceil(n/4) words would overrun a num_tiles%4 buffer).
+extern "C" __global__ void pose_dp_fill_bytes(
+		int             n,
+		unsigned char   value,
+		unsigned char * data)
+{
+	const int i = (int) (blockIdx.x * blockDim.x + threadIdx.x);
+	if (i < n) data[i] = value;
+}
+
+extern "C" __global__ void pose_scene_dp_step(PoseSceneDp * s);
+extern "C" __global__ void pose_scene_dp_commit(PoseSceneDp * s);
+
+// Per-cycle LMA continuation: tail-enqueued by pose_measure_dp_tail AFTER the
+// peak kernel, so the corr-row count and peak rows are LIVE. Snapshots the
+// measure status trace, enqueues the prepare stages (full or light) and
+// chains the step launcher.
+extern "C" __global__ void pose_scene_dp_lma(PoseSceneDp * s)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	const int cycle = s->cycle;
+	if (s->mstats != NULL) {
+		for (int i = 0; i < POSE_MEASURE_STATUS_INTS; i++) {
+			s->mstats[cycle * POSE_MEASURE_STATUS_INTS + i] = s->mc->status[i];
+		}
+	}
+	const int nct = *s->num_corr_tiles; // LIVE (>0: the tail chains only on success)
+	const int set = *s->current_set;
+	float         * cur_vectors = s->set_vectors[set];
+	float         * cur_fx      = s->set_fx[set];
+	float         * cur_jt      = s->set_jt[set];
+	unsigned char * cur_valid   = s->set_valid[set];
+	const int pose_threads = 256;
+	const int pose_blocks = (s->num_tiles + pose_threads - 1) / pose_threads;
+	const int a_threads = 256;                       // host assemble launch shape
+	const int a_blocks = (nct + a_threads - 1) / a_threads;
+	const int light = s->prepare_light;
+	const int nc = s->num_components;
+	const int np = POSE_NUM_PARAMS;
+	if (light) {
+		// the D1 light body: frozen conditioning, fresh measured y offsets only
+		const int clear_n = nc * s->num_tiles;
+		pose_lma_dp_clear<<<(clear_n + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				clear_n, s->y);
+	} else {
+		// exec_pose_lma_prepare_resident(light=0) pre-zero: y/weights over the
+		// full value range, eigen, selection, and the compact prep result (the
+		// resident pull/reg rows behind it survive - the host memset range).
+		const int num_values = nc * s->num_tiles + np;
+		pose_lma_dp_clear<<<(num_values + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				num_values, s->y);
+		pose_lma_dp_clear<<<(num_values + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				num_values, s->weights);
+		pose_lma_dp_clear<<<(4 * s->num_tiles + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				4 * s->num_tiles, s->eigen);
+		pose_dp_fill_bytes<<<(s->num_tiles + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(
+				s->num_tiles, 0, s->selection);
+		pose_lma_dp_clear<<<1, 256, 0, cudaStreamTailLaunch>>>(
+				POSE_LMA_PREP_RESULT, s->prep_result);
+	}
+	pose_lma_assemble<<<a_blocks, a_threads, 0, cudaStreamTailLaunch>>>(
+			nct, s->num_tiles, s->corr_slot, s->ntile_shift,
+			nc, s->use_eigen, s->eig_min_sqrt, s->eig_max_sqrt,
+			s->min_confidence, s->same_weights, light,
+			s->corr_indices, s->peaks,
+			s->y, s->weights, s->eigen, s->selection);
+	if (!light) {
+		pose_lma_prepare_norm<<<1, 1, 0, cudaStreamTailLaunch>>>(
+				s->num_tiles, nc, np, s->pull,
+				s->y, s->weights, s->prep_result);
+	}
+	pose_fx_jacobian<<<pose_blocks, pose_threads, 0, cudaStreamTailLaunch>>>(
+			s->num_tiles, 1, s->cameras,
+			s->ref_radial, s->ref_rbr, s->ref_ers,
+			s->scene_radial, s->scene_rbr, s->scene_ers,
+			cur_vectors, s->centers, s->selection,
+			cur_fx, cur_jt, cur_valid);
+	pose_lma_prepare_finish<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			s->num_tiles, nc, np, light,
+			cur_fx, cur_jt, cur_valid, s->eigen, s->regweights,
+			s->y, s->weights, s->prep_result);
+	pose_scene_dp_step<<<1, 1, 0, cudaStreamTailLaunch>>>(s);
+}
+
+// Step launcher: runs AFTER pose_lma_prepare_finish, so the full-mode
+// pure_weight and prepare validity are LIVE. Enqueues the six validated
+// resident step stages (the spike/D1 sequence) and chains the commit.
+extern "C" __global__ void pose_scene_dp_step(PoseSceneDp * s)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	float pure_weight = s->pure_weight;
+	if (!s->prepare_light) {
+		if (s->prep_result[POSE_LMA_PREP_VALID] == 0.0f) {
+			s->status[1] = -13; // = the host prepare's invalid code; chain stops (fail-safe)
+			return;
+		}
+		pure_weight = s->prep_result[POSE_LMA_PREP_PURE_WEIGHT];
+	}
+	const int set = *s->current_set;
+	float         * cur_vectors = s->set_vectors[set];
+	float         * cur_fx      = s->set_fx[set];
+	float         * cur_jt      = s->set_jt[set];
+	unsigned char * cur_valid   = s->set_valid[set];
+	float         * cand_vectors= s->set_vectors[set ^ 1];
+	float         * cand_fx     = s->set_fx[set ^ 1];
+	float         * cand_jt     = s->set_jt[set ^ 1];
+	unsigned char * cand_valid  = s->set_valid[set ^ 1];
+	const float   * candidate_vector = cand_vectors + 3 * POSE_NUM_PARAMS;
+	const int pose_threads = 256;
+	const int pose_blocks = (s->num_tiles + pose_threads - 1) / pose_threads;
+	pose_lma_prepare_products<<<s->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			s->num_tiles, cur_fx, cur_jt, cur_valid,
+			s->y, s->weights, s->eigen, s->vector,
+			s->jt_out, s->ymfx_out, s->partials);
+	pose_lma_reduce_step<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			s->num_partials, s->lambda, pure_weight,
+			s->partials, s->vector, s->result);
+	pose_lma_dp_candidate_vector<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			cur_vectors, s->result, cand_vectors);
+	pose_fx_jacobian<<<pose_blocks, pose_threads, 0, cudaStreamTailLaunch>>>(
+			s->num_tiles, 1, s->cameras,
+			s->ref_radial, s->ref_rbr, s->ref_ers,
+			s->scene_radial, s->scene_rbr, s->scene_ers,
+			cand_vectors, s->centers, s->selection,
+			cand_fx, cand_jt, cand_valid);
+	pose_lma_prepare_products<<<s->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			s->num_tiles, cand_fx, cand_jt, cand_valid,
+			s->y, s->weights, s->eigen, candidate_vector,
+			s->cand_jt_out, s->cand_ymfx_out, s->cand_partials);
+	pose_lma_reduce_rms<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			s->num_partials, pure_weight, s->rms_diff,
+			s->cand_partials, s->result, s->current_set);
+	pose_scene_dp_commit<<<1, 1, 0, cudaStreamTailLaunch>>>(s);
+}
+
+// Commit: the anchor tracks the ACCEPTED pose (the D1 rule - a stale anchor
+// makes every following candidate overshoot), per-cycle 25-float trace
+// snapshot, then the NEXT cycle's measure chain (self-chaining: this launcher
+// ran tail-enqueued behind reduce_rms, so the accept flip is visible to the
+// next pose_measure_dp).
+extern "C" __global__ void pose_scene_dp_commit(PoseSceneDp * s)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	if (s->result[POSE_LMA_ACCEPTED] != 0.0f) {
+		#pragma unroll
+		for (int par = 0; par < POSE_NUM_PARAMS; par++) {
+			s->vector[par] = s->result[16 + par];
+		}
+	}
+	const int cycle = s->cycle;
+	if (s->trace != NULL) {
+		for (int i = 0; i < POSE_LMA_RESIDENT_RESULT; i++) {
+			s->trace[cycle * POSE_LMA_RESIDENT_RESULT + i] = s->result[i];
+		}
+	}
+	s->status[0] = cycle + 1;
+	if (cycle + 1 < s->num_cycles) {
+		s->cycle = cycle + 1;
+		pose_measure_dp<<<1, 1, 0, cudaStreamTailLaunch>>>(s->mc);
+	}
+}
+
+// The D3 parent: ONE device entry for the whole per-scene solve. The measure
+// pose needs no upload beyond scene entry - after an accept flip the current
+// set already holds the candidate built by pose_lma_dp_candidate_vector
+// (bit-identical to the host re-upload, the D1-proven fact).
+extern "C" __global__ void pose_scene_dp(PoseSceneDp * s)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	s->cycle = 0;
+	s->status[0] = 0;
+	s->status[1] = 0;
+	pose_measure_dp<<<1, 1, 0, cudaStreamTailLaunch>>>(s->mc);
+}
+
 #endif // __CUDACC_RTC__
 
 extern "C" __global__ void lma_normal_products(

src/tp_lma.h

@@ -371,9 +371,75 @@ typedef struct {
 	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
+	// DP rung D3: non-NULL = per-scene parent continuation (PoseSceneDp*). The
+	// peak tail launcher chains pose_scene_dp_lma after a successful measure;
+	// guard failures stop the whole scene chain (fail-safe). NULL = standalone
+	// rung-D2 behavior, unchanged. By Claude on 07/17/2026.
+	void * scene;
 } PoseMeasureChain;
 #define POSE_MEASURE_STATUS_INTS 5
 
+// ── DP rung D3 (ladder ratified 2026-07-16): full per-scene parent ───────────
+// ONE device entry per scene: num_cycles x [measure chain (the D2 stage list,
+// via the resident PoseMeasureChain) -> prepare -> step -> commit], per-cycle
+// QC trace to device buffers, ONE end-of-scene readback. prepare_light selects
+// the per-cycle conditioning policy: 0 = production FULL prepare (per-cycle
+// re-derivation - C3/C3b falsified freezing, pose_freeze_cycle defaults to 0;
+// pure_weight is read LIVE from prep_result each cycle), 1 = the D1 light body
+// (frozen conditioning, static pure_weight; kept for the recorded DP-era
+// experiments, e.g. previous-scene conditioning). The launchers self-chain
+// (D1 pattern): each runs tail-enqueued BEHIND the work it must observe, so
+// LIVE reads of current_set / corr-row count / prep_result are ordered.
+// cycle/status are device-mutated; both descriptors live in device memory.
+// By Claude on 07/17/2026.
+typedef struct {
+	int   num_cycles;
+	int   cycle;               // device-mutable loop index (commit advances it)
+	int   prepare_light;       // 0 = per-cycle full prepare (production), 1 = D1 light body
+	int   num_tiles;
+	int   num_partials;
+	float lambda;              // production policy v1: fixed lambda0, inner=1
+	float pure_weight;         // light mode only; full mode reads prep_result LIVE
+	float rms_diff;
+	// pose_lma_assemble policy (= the host prepare's scalar args)
+	int   corr_slot;
+	int   ntile_shift;
+	int   num_components;
+	int   use_eigen;
+	float eig_min_sqrt;
+	float eig_max_sqrt;
+	float min_confidence;
+	int   same_weights;
+	// the measure chain this parent drives (device copy; its .scene points back)
+	const PoseMeasureChain * mc;
+	// resident correlation / camera / tile state (= PoseDpCycles, live counts)
+	const int   * corr_indices;
+	const float * peaks;
+	const int   * num_corr_tiles; // LIVE corr-row count (assemble grid per cycle)
+	const PoseCameraMeta * cameras;
+	const float * ref_radial;  const float * ref_rbr;  const float * ref_ers;
+	const float * scene_radial;const float * scene_rbr;const float * scene_ers;
+	const float * centers;
+	unsigned char * selection; // full prepare clears+rebuilds it every cycle
+	float         * set_vectors[2];
+	float         * set_fx[2];
+	float         * set_jt[2];
+	unsigned char * set_valid[2];
+	float * y; float * weights; float * eigen;
+	float * vector;            // linearization anchor; device-updated on accept
+	const float * pull;        // pull-target rows (full prepare, resident upload)
+	const float * regweights;
+	float * jt_out; float * ymfx_out; float * partials;
+	float * cand_jt_out; float * cand_ymfx_out; float * cand_partials;
+	float * result;            // [POSE_LMA_RESIDENT_RESULT]
+	float * prep_result;       // [POSE_LMA_PREP_RESULT] (+pull/reg rows behind it)
+	int   * current_set;
+	float * trace;             // [num_cycles][POSE_LMA_RESIDENT_RESULT]
+	int   * mstats;            // [num_cycles][POSE_MEASURE_STATUS_INTS]
+	int   * status;            // [0] cycles completed, [1] LMA rc (0 / -13 invalid prepare)
+} PoseSceneDp;
+#define POSE_SCENE_STATUS_INTS 2
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,