CLAUDE: DP rung D1 - K chained prepare-light+step cycles from ONE device entry

Self-chaining launcher pose_lma_dp_cycles: each cycle's launcher reads the
LIVE current_set (tail-enqueued behind the previous reduce_rms, so it sees
the accept flip a single unrolled parent could not), resolves cur/cand from
both physical sets in the PoseDpCycles descriptor, tail-enqueues the exact
host per-cycle sequence (clear-y = light memset, assemble(light),
fx_jacobian(current), prepare_finish(light), the six spike step stages),
then pose_lma_dp_commit - the device-side anchor update on accept that the
spike flagged as the multi-step gap - plus a minimal per-cycle 25-float
trace (ruling b), and finally the next launcher. Fixed-K, no device early
exit (ruling a). Interior cycles cross NOTHING host-side.

Wrapper tp_proc_exec_pose_lma_dp_cycles: entry H2D = pose vectors + anchor
+ descriptor; one launch; D2H once = trace + final result + final anchor +
set mirror. Preconditions: frozen conditioning (C3) + resident corr/peaks.

Gate: test_pose_lma_dp_cycles_jna (new, in cases.list as dp_cycles) seeds
resident measurement through the production chain (set_corr_indices_td ->
normalize -> corr2D_peak_eig_f), freezes conditioning, then compares three
chains (K=3 accept, K=2 reject, K=2 post-reject) against host-driven
light-prepare+step cycles: per-cycle traces, final packed, final anchor,
set index ALL BIT-IDENTICAL. run_cases.sh ALL PASS (pose_corr @tol 0,
peak, peak_recenter, avg_td, avg_td_oob, bayer_staged, dp_cycles); DP
spike test still PASS.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

cases.list.tmpl

@@ -65,6 +65,7 @@ case peak_recenter   test_peak             ${REPO}/testdata/synth4/peak_recenter
 case avg_td          test_avg_td           ${REPO}/testdata/synth4/avg_td        --tol 0
 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 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_lma_dp_cycles_jna.cu

+// test_pose_lma_dp_cycles_jna.cu - DP rung D1 gate: K chained prepare-light +
+// step cycles as ONE device-side entry (tp_proc_exec_pose_lma_dp_cycles) must
+// be BIT-IDENTICAL to the host-driven per-cycle pair
+// (tp_proc_exec_pose_lma_prepare_resident(light) +
+//  tp_proc_exec_pose_lma_resident_step(prepared)).
+//
+// Two TpProc instances from one NVRTC module (the step MUTATES state on
+// accept, same reasoning as the DP spike test). Resident measurement is seeded
+// through the PRODUCTION chain - synthetic TD correlation tiles uploaded with
+// tp_proc_set_corr_indices_td, then exec_corr2d_normalize + exec_corr2d_peak
+// leave resident corr indices + peaks exactly as a lean measure cycle would.
+// One FULL resident prepare freezes the conditioning (rung C3), one lockstep
+// host step advances both instances; then three chains exercise the launcher:
+//   chain 1: K=3, production lambda   -> accepts (device anchor update + flips)
+//   chain 2: K=2, near-infinite lambda -> rejects (anchor + set retained)
+//   chain 3: K=2, production lambda   -> post-reject accepts
+// Instance A mirrors every chain with host-driven cycles; per-cycle trace rows,
+// the final packed result, the final anchor and the set index must all memcmp
+// equal. Any diff = a real D1 bug (live-set resolution, anchor commit, or the
+// device-side light prepare diverging from the host sequence).
+// 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 \
+//     jna/test_pose_lma_dp_cycles_jna.cu -o tests_bin/test_pose_lma_dp_cycles_jna \
+//     -L jna -ltileproc -Xlinker -rpath -Xlinker $PWD/jna
+// Run from repo root: ./tests_bin/test_pose_lma_dp_cycles_jna src
+
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <string>
+#include <vector>
+
+enum { NUM_PARAMS=3, NUM_COMPONENTS=2, RESULT=25, PEAK_FLOATS=8, TD_FLOATS=256 };
+
+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_setup_rbg_corr(void*,int,int,int,int,int,float,float,float,int);
+  int tp_proc_set_corr_indices_td(void*,const int*,const float*,int);
+  int tp_proc_exec_corr2d_normalize(void*,int,double,int);
+  int tp_proc_exec_corr2d_peak(void*,int,float,float,float,float,int,float,float,float,int);
+  int tp_proc_get_peaks(void*,float*);
+  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_lma_dp_cycles(void*,int,float,float,float,int,
+      int,int,int,int,float,float,float,int,
+      const float*,const float*,float*,float*,float*);
+  void tp_proc_destroy(void*);
+  void tp_destroy_module(void*);
+}
+
+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;
+}
+
+// Peak/assemble policy - shared by the peak kernel call, the full prepare and
+// the DP chain descriptor (identical on both instances; values are test policy,
+// not production constants).
+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 PURE_WEIGHT_FALLBACK=0.98f; // replaced by prep_result below
+static const float RMS_DIFF=0.001f;
+
+int main(int argc,char** argv)
+{
+  // Accept both the bare positional form (srcdir [devrt]) and the run_cases.sh
+  // manifest form (--data <unused> --tol <unused> --src <dir> --devrt <lib>).
+  const char* srcdir="src";
+  const char* devrt="/usr/local/cuda-12.8/targets/x86_64-linux/lib/libcudadevrt.a";
+  int npos=0;
+  for(int i=1;i<argc;i++){
+    if(!std::strcmp(argv[i],"--src")&&(i+1<argc)) srcdir=argv[++i];
+    else if(!std::strcmp(argv[i],"--devrt")&&(i+1<argc)) devrt=argv[++i];
+    else if((!std::strcmp(argv[i],"--data")||!std::strcmp(argv[i],"--tol"))&&(i+1<argc)) i++;
+    else if(argv[i][0]!='-'){ if(npos==0) srcdir=argv[i]; else if(npos==1) devrt=argv[i]; npos++; }
+  }
+  void* module=tp_create_module(srcdir,devrt);
+  if(!module){ std::fprintf(stderr,"tp_create_module: %s\n",tp_last_error()); return EXIT_FAILURE; }
+
+  // Two instances: A = host-driven oracle, B = DP chained cycles.
+  void* procs[2]={nullptr,nullptr};
+  for(int k=0;k<2;k++){
+    procs[k]=tp_proc_create(module);
+    // img 160x128 -> 20x16 = 320 corr-capable tiles >= the 280 seeded rows
+    // (setup dims size the corr buffers only; pose projection uses the meta
+    // camera dims below).
+    if(!procs[k]||tp_proc_setup(procs[k],1,1,160,128,0,0)||
+       tp_proc_setup_rbg_corr(procs[k],1,0xffff,0,0,0,1.0f,0.0f,0.0f,7)){
+      std::fprintf(stderr,"tp_proc setup[%d]: %s\n",k,tp_last_error()); return EXIT_FAILURE; }
+  }
+
+  // Tile-grid geometry sized like the spike test: multi-block child grids
+  // ((300+255)/256 = 2) and multi-partial reduce (300/128 -> 3).
+  const int width=64,height=48,rbrLen=256,ersLen=height*14,numTiles=300;
+  const float meta[8]={(float)width,(float)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(rbrLen,1.0f),ers(ersLen,0.0f);
+  for(int line=0;line<height;line++) ers[(size_t)line*14+6]=1.0f;
+  const float v0[NUM_PARAMS]={0.0023f,-0.0017f,0.00091f};
+  std::vector<float> centers(3*numTiles);
+  for(int i=0;i<numTiles;i++){
+    centers[3*i+0]=4.0f+(float)((i*7)%56);
+    centers[3*i+1]=4.0f+(float)((i*5)%40);
+    centers[3*i+2]=(float)(i%6);
+  }
+
+  // Synthetic resident measurement through the production chain: per measured
+  // tile a TD tile whose 4 DTT quadrants encode a windowed bump shifted by
+  // (dx,dy) - normalize converts TD->PD, corr2D_peak_eig_f measures it. Exact
+  // DTT sign conventions do not matter here: any stable finite sub-pixel peak
+  // with nonzero offsets gives the LMA real residuals, and both instances see
+  // the SAME bytes (bit-exactness is between A and B, not against an analytic
+  // truth). A few tiles are left unmeasured to exercise selection gaps.
+  const int numCorr=280;
+  std::vector<int> corrIndices(numCorr);
+  std::vector<float> td((size_t)numCorr*TD_FLOATS,0.0f);
+  {
+    int row=0;
+    for(int tile=0;tile<numTiles&&row<numCorr;tile++){
+      if((tile%15)==7) continue; // unmeasured tiles
+      corrIndices[row]=(tile<<NTILE_SHIFT)|CORR_SLOT;
+      const float dx=0.6f*std::sin(0.13f*(float)tile);
+      const float dy=0.5f*std::cos(0.11f*(float)tile+0.4f);
+      const float amp=200.0f*(1.0f+0.3f*std::sin(0.05f*(float)tile));
+      float* t=&td[(size_t)row*TD_FLOATS];
+      for(int i=0;i<8;i++){
+        const float a=(float)M_PI*((float)i+0.5f)/8.0f;
+        const float gi=std::exp(-0.02f*(float)(i*i));
+        const float cx=std::cos(a*dx)*gi, sx=std::sin(a*dx)*gi;
+        for(int j=0;j<8;j++){
+          const float b=(float)M_PI*((float)j+0.5f)/8.0f;
+          const float gj=std::exp(-0.02f*(float)(j*j));
+          const float cy=std::cos(b*dy)*gj, sy=std::sin(b*dy)*gj;
+          t[0*64+i*8+j]=amp*cx*cy;  // CC
+          t[1*64+i*8+j]=amp*sx*cy;  // SC
+          t[2*64+i*8+j]=amp*cx*sy;  // CS
+          t[3*64+i*8+j]=amp*sx*sy;  // SS
+        }
+      }
+      row++;
+    }
+    if(row<numCorr){ std::fprintf(stderr,"internal: only %d corr rows\n",row); return EXIT_FAILURE; }
+  }
+
+  int bad=0;
+  std::vector<float> peaks((size_t)numCorr*PEAK_FLOATS);
+  for(int k=0;k<2;k++){
+    int rcs=tp_proc_set_corr_indices_td(procs[k],corrIndices.data(),td.data(),numCorr);
+    if(rcs){
+      std::fprintf(stderr,"set_corr_indices_td[%d] rc=%d: %s\n",k,rcs,tp_last_error()); return EXIT_FAILURE; }
+    if(tp_proc_exec_corr2d_normalize(procs[k],0,30.0,7)){
+      std::fprintf(stderr,"corr2d_normalize[%d]: %s\n",k,tp_last_error()); return EXIT_FAILURE; }
+    const int npk=tp_proc_exec_corr2d_peak(procs[k],CORR_SLOT,
+        1.0e-6f,0.0f,1.0e-6f,0.0f,1,0.0f,1.0f,0.0f,1); // use_float=1 (lean path)
+    if(npk!=numCorr){
+      std::fprintf(stderr,"corr2d_peak[%d] rc=%d: %s\n",k,npk,tp_last_error()); return EXIT_FAILURE; }
+  }
+  { // sanity + instance identity of the seeded measurement
+    std::vector<float> peaksB((size_t)numCorr*PEAK_FLOATS);
+    if(tp_proc_get_peaks(procs[0],peaks.data())!=numCorr||
+       tp_proc_get_peaks(procs[1],peaksB.data())!=numCorr||
+       std::memcmp(peaks.data(),peaksB.data(),peaks.size()*sizeof(float))){
+      std::printf("seeded peaks diverge between instances\n"); bad++; }
+    int nvalid=0; float adx=0.0f;
+    for(int r=0;r<numCorr;r++){
+      if(peaks[(size_t)r*PEAK_FLOATS+7]!=0.0f){
+        nvalid++; adx+=std::fabs(peaks[(size_t)r*PEAK_FLOATS]); }
+    }
+    std::printf("seeded peaks: %d/%d valid, mean|dx|=%.4f\n",nvalid,numCorr,
+        nvalid?adx/(float)nvalid:0.0f);
+    if(nvalid<numCorr/2){ std::printf("too few valid peaks - synthetic TD too weak\n"); bad++; }
+    if(nvalid&&(adx/(float)nvalid)<1.0e-3f){
+      std::printf("peaks have ~zero offsets - steps would never accept\n"); bad++; }
+  }
+  if(bad) return EXIT_FAILURE;
+
+  // One FULL resident prepare per instance: freezes conditioning (rung C3) and
+  // projects fx/J at the initial pose. pure_weight comes from prep_result.
+  const float pose0[12]={0,0,0, 0,0,0, 0,0,0, v0[0],v0[1],v0[2]};
+  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};
+  float pureWeight=PURE_WEIGHT_FALLBACK;
+  for(int k=0;k<2;k++){
+    float prep[8]={0};
+    int rc=tp_proc_exec_pose_lma_prepare_resident(procs[k],
+        meta,radial,rbr.data(),rbrLen,ers.data(),ersLen,
+        meta,radial,rbr.data(),rbrLen,ers.data(),ersLen,
+        pose0,centers.data(),numTiles,NUM_COMPONENTS,CORR_SLOT,NTILE_SHIFT,
+        USE_EIGEN,EIG_MIN_SQRT,EIG_MAX_SQRT,MIN_CONFIDENCE,SAME_WEIGHTS,
+        pull,regw,0,prep,0,nullptr,nullptr,nullptr);
+    if(rc){ std::fprintf(stderr,"full prepare[%d] rc=%d: %s\n",k,rc,tp_last_error()); return EXIT_FAILURE; }
+    if(k==0){
+      std::printf("full prepare: valid=%g num_meas=%g pure_weight=%g\n",
+          prep[0],prep[1],prep[3]);
+      pureWeight=prep[3];
+    }
+  }
+
+  // Lockstep host step 1 on both instances (must accept - the measured offsets
+  // are real residuals at v0).
+  float anchor[NUM_PARAMS]={v0[0],v0[1],v0[2]};
+  float packed1[2][RESULT];
+  for(int k=0;k<2;k++){
+    int rc=tp_proc_exec_pose_lma_resident_step(procs[k],0.001f,
+        nullptr,nullptr,nullptr,anchor,pureWeight,RMS_DIFF,numTiles,0,
+        packed1[k],nullptr,nullptr);
+    if(rc){ std::fprintf(stderr,"host step1[%d] rc=%d: %s\n",k,rc,tp_last_error()); return EXIT_FAILURE; }
+  }
+  if(comparePacked("step1(lockstep)",packed1[0],packed1[1])) bad++;
+  if(packed1[0][23]!=1.0f){
+    std::printf("step1 not accepted (current=%g candidate=%g) - seed too weak\n",
+        packed1[0][19],packed1[0][21]); bad++; }
+  if(bad) return EXIT_FAILURE;
+  for(int par=0;par<NUM_PARAMS;par++) anchor[par]=packed1[0][16+par]; // accepted pose
+
+  // Three chains: accept run, reject run, post-reject accept run. Instance A
+  // mirrors each chain with host light-prepare + step cycles at the SAME fixed
+  // lambda; the anchor/pose bookkeeping below is exactly the production host
+  // policy (update on accept only).
+  const struct { int cycles; float lambda; const char* tag; } chains[3]={
+    {3,0.0005f,"chain1(accept)"},{2,1.0e6f,"chain2(reject)"},{2,0.0005f,"chain3(post-reject)"}};
+  int dpAccepts=0,dpRejects=0;
+  for(int chain=0;chain<3&&!bad;chain++){
+    const int K=chains[chain].cycles;
+    const float lambda=chains[chain].lambda;
+    float poseK[12]={0,0,0, 0,0,0, 0,0,0, anchor[0],anchor[1],anchor[2]};
+    // A: host-driven cycles
+    std::vector<float> packedA((size_t)K*RESULT);
+    float anchorA[NUM_PARAMS]={anchor[0],anchor[1],anchor[2]};
+    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 rc=tp_proc_exec_pose_lma_prepare_resident(procs[0],
+          nullptr,nullptr,nullptr,0,nullptr,0,
+          nullptr,nullptr,nullptr,0,nullptr,0,
+          poseC,nullptr,numTiles,NUM_COMPONENTS,CORR_SLOT,NTILE_SHIFT,
+          USE_EIGEN,EIG_MIN_SQRT,EIG_MAX_SQRT,MIN_CONFIDENCE,SAME_WEIGHTS,
+          nullptr,nullptr,1,nullptr,0,nullptr,nullptr,nullptr);
+      if(rc){ std::fprintf(stderr,"%s host light prepare cyc%d rc=%d: %s\n",
+          chains[chain].tag,cyc,rc,tp_last_error()); bad++; break; }
+      rc=tp_proc_exec_pose_lma_resident_step(procs[0],lambda,
+          nullptr,nullptr,nullptr,anchorA,pureWeight,RMS_DIFF,numTiles,0,
+          &packedA[(size_t)cyc*RESULT],nullptr,nullptr);
+      if(rc){ std::fprintf(stderr,"%s host step cyc%d rc=%d: %s\n",
+          chains[chain].tag,cyc,rc,tp_last_error()); bad++; break; }
+      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];
+    }
+    if(bad) break;
+    // B: ONE DP entry for the whole chain
+    std::vector<float> trace((size_t)K*RESULT,0.0f);
+    float packedB[RESULT],vectorB[NUM_PARAMS];
+    int rc=tp_proc_exec_pose_lma_dp_cycles(procs[1],K,lambda,pureWeight,RMS_DIFF,
+        numTiles,CORR_SLOT,NTILE_SHIFT,NUM_COMPONENTS,USE_EIGEN,
+        EIG_MIN_SQRT,EIG_MAX_SQRT,MIN_CONFIDENCE,SAME_WEIGHTS,
+        poseK,anchor,trace.data(),packedB,vectorB);
+    if(rc){ std::fprintf(stderr,"%s dp_cycles rc=%d: %s\n",
+        chains[chain].tag,rc,tp_last_error()); bad++; break; }
+    // Compare: per-cycle trace rows, final packed, final anchor.
+    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(trace[(size_t)cyc*RESULT+23]==1.0f) dpAccepts++; else dpRejects++;
+    }
+    if(comparePacked((std::string(chains[chain].tag)+" final").c_str()
+        ,&packedA[(size_t)(K-1)*RESULT],packedB)) bad++;
+    if(std::memcmp(anchorA,vectorB,sizeof(anchorA))){
+      std::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++; }
+    std::printf("%s K=%d lambda=%g: last accepted=%g rms %g -> %g : %s\n",
+        chains[chain].tag,K,lambda,packedB[23],packedB[19],packedB[21],
+        bad?"DIFF":"bit-identical");
+    for(int par=0;par<NUM_PARAMS;par++) anchor[par]=anchorA[par]; // both advanced equally
+  }
+  if(!bad&&(!dpAccepts||!dpRejects)){
+    std::printf("chain coverage broken: accepts=%d rejects=%d (need both)\n",
+        dpAccepts,dpRejects); bad++;
+  }
+
+  std::printf("DP K-cycle chain vs host-driven light-prepare+step cycles -> %s\n",
+      bad?"FAIL":"PASS (bit-identical)");
+  tp_proc_destroy(procs[0]); tp_proc_destroy(procs[1]); tp_destroy_module(module);
+  return bad?EXIT_FAILURE:EXIT_SUCCESS;
+}

jna/tp_jna.cpp

@@ -701,6 +701,10 @@ struct TpProc {
   int pose_prep_centers_tiles; // tile count the resident centers were uploaded for // By Claude on 07/17/2026
   int pose_prep_frozen_tiles;  // rung C3: tile count of the frozen conditioning (0 = no full prepare yet) // By Claude on 07/17/2026
   PoseDpChain *gpu_pose_dp_chain; // DP spike: device copy of the chain descriptor // By Claude on 07/16/2026
+  // DP rung D1: device copy of the K-cycle chain descriptor + per-cycle result
+  // 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;
   // 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,
@@ -793,6 +797,7 @@ TpProc* tp_proc_create(TpModule* m){
   p->gpu_pose_lma_prep=nullptr; // By Claude on 07/17/2026
   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_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)
@@ -2108,6 +2113,112 @@ int tp_proc_exec_pose_lma_dp_step(TpProc* p,
   return 0;
 }
 
+// DP rung D1: K chained prepare-light + step cycles as ONE device-side entry
+// (fixed measurement - the resident corr indices/peaks are re-assembled each
+// cycle, exactly the host light-prepare + resident-step pair). Preconditions:
+// a FULL resident prepare froze the conditioning at this tile count (rung C3)
+// and the raw pose state is live. Per-call H2D = the 12-float pose vectors +
+// the 3-float anchor + the chain descriptor; interior cycles cross NOTHING.
+// D2H once at the end: per-cycle trace (num_cycles x 25), final result[25],
+// final anchor, current_set mirror. By Claude on 07/17/2026.
+int tp_proc_exec_pose_lma_dp_cycles(TpProc* p,
+                                    int num_cycles,
+                                    float lambda,
+                                    float pure_weight,
+                                    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,
+                                    const float* pose_vectors,
+                                    const float* vector,
+                                    float* trace,
+                                    float* result,
+                                    float* vector_out){
+  if(!p||!result||!pose_vectors||!vector||(num_cycles<=0)||(num_tiles<=0)||
+     (num_components!=POSE_NUM_COMPONENTS)||
+     !std::isfinite(pure_weight)||!(pure_weight>0.0f)||!std::isfinite(rms_diff)){
+    seterr("exec_pose_lma_dp_cycles: bad args k=%d nt=%d nc=%d",num_cycles,num_tiles,num_components); return -1; }
+  if((p->pose_raw_num_tiles!=num_tiles)||!p->pose_raw_has_jacobian||
+     (p->pose_lma_value_capacity<POSE_NUM_COMPONENTS*num_tiles+POSE_NUM_PARAMS)||
+     !p->gpu_pose_lma_weights||!p->gpu_pose_lma_y||!p->gpu_pose_lma_eigen||
+     !p->gpu_pose_lma_vector||!p->gpu_pose_current_set||!p->gpu_pose_lma_prep){
+    seterr("exec_pose_lma_dp_cycles: resident prepared state not ready (nt=%d)",num_tiles); return -2; }
+  if(p->pose_prep_frozen_tiles!=num_tiles){ // light cycles need the frozen conditioning (rung C3)
+    seterr("exec_pose_lma_dp_cycles: conditioning not frozen (%d/%d)",
+           p->pose_prep_frozen_tiles,num_tiles); return -2; }
+  if(!p->gpu_corr_indices||!p->gpu_peaks||(p->last_num_corr_tiles<=0)){
+    seterr("exec_pose_lma_dp_cycles: no resident correlation/peaks (num_corr=%d)",
+           p->last_num_corr_tiles); return -2; }
+  cuCtxSetCurrent(p->mod->ctx);
+  const int num_partials=(num_tiles+POSE_LMA_REDUCE_THREADS-1)/POSE_LMA_REDUCE_THREADS;
+  if(!p->gpu_pose_dp_cycles &&
+     cudaMalloc((void**)&p->gpu_pose_dp_cycles,sizeof(PoseDpCycles))!=cudaSuccess){
+    seterr("exec_pose_lma_dp_cycles: chain alloc failed"); return -3; }
+  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_lma_dp_cycles: trace alloc failed k=%d",num_cycles); return -3; }
+    p->pose_dp_trace_capacity=num_cycles;
+  }
+  // Chain-entry state: pose vectors into the CURRENT set + the anchor - the
+  // same two uploads a host cycle would make; the chained cycles then keep
+  // both device-resident. By Claude on 07/17/2026.
+  if(cudaMemcpy(p->cur_pose_vectors(),pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),
+                cudaMemcpyHostToDevice)!=cudaSuccess ||
+     cudaMemcpy(p->gpu_pose_lma_vector,vector,POSE_NUM_PARAMS*sizeof(float),
+                cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_cycles: entry HtoD failed"); return -4; }
+  PoseDpCycles h;
+  h.num_cycles=num_cycles; h.num_tiles=num_tiles;
+  h.num_corr=p->last_num_corr_tiles; h.num_partials=num_partials;
+  h.lambda=lambda; h.pure_weight=pure_weight; h.rms_diff=rms_diff;
+  h.corr_slot=corr_slot; h.ntile_shift=ntile_shift;
+  h.num_components=num_components; h.use_eigen=use_eigen;
+  h.eig_min_sqrt=eig_min_sqrt; h.eig_max_sqrt=eig_max_sqrt;
+  h.min_confidence=min_confidence; h.same_weights=same_weights;
+  h.corr_indices=p->gpu_corr_indices; h.peaks=p->gpu_peaks;
+  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_rbr=p->gpu_pose_rbr[1]; h.scene_ers=p->gpu_pose_ers[1];
+  h.centers=p->gpu_pose_centers; h.selection=p->gpu_pose_selection;
+  h.set_vectors[0]=p->gpu_pose_vectors;           h.set_vectors[1]=p->gpu_pose_candidate_vectors;
+  h.set_fx[0]=p->gpu_pose_fx;                     h.set_fx[1]=p->gpu_pose_candidate_fx;
+  h.set_jt[0]=p->gpu_pose_jt;                     h.set_jt[1]=p->gpu_pose_candidate_jt;
+  h.set_valid[0]=p->gpu_pose_valid;               h.set_valid[1]=p->gpu_pose_candidate_valid;
+  h.y=p->gpu_pose_lma_y; h.weights=p->gpu_pose_lma_weights; h.eigen=p->gpu_pose_lma_eigen;
+  h.vector=p->gpu_pose_lma_vector;
+  h.regweights=p->gpu_pose_lma_prep+POSE_LMA_PREP_RESULT+POSE_NUM_PARAMS;
+  h.jt_out=p->gpu_pose_lma_jt; h.ymfx_out=p->gpu_pose_lma_ymfx; h.partials=p->gpu_pose_lma_partials;
+  h.cand_jt_out=p->gpu_pose_lma_candidate_jt; h.cand_ymfx_out=p->gpu_pose_lma_candidate_ymfx;
+  h.cand_partials=p->gpu_pose_lma_candidate_partials;
+  h.result=p->gpu_pose_lma_result; h.prep_result=p->gpu_pose_lma_prep;
+  h.current_set=p->gpu_pose_current_set;
+  h.trace=trace?p->gpu_pose_dp_trace:nullptr;
+  if(cudaMemcpy(p->gpu_pose_dp_cycles,&h,sizeof(PoseDpCycles),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_cycles: chain HtoD failed"); return -4; }
+  CUfunction dp=getfun(p->mod,"pose_lma_dp_cycles");
+  if(!dp){ seterr("pose_lma_dp_cycles missing (DP kernel not in module)"); return -5; }
+  int cycle0=0;
+  void* da[]={ &p->gpu_pose_dp_cycles, &cycle0 };
+  if(launch1(dp,1,1,1,1,1,1,da,"pose_lma_dp_cycles")) return -6;
+  if((trace &&
+      cudaMemcpy(trace,p->gpu_pose_dp_trace,
+                 (size_t)num_cycles*POSE_LMA_RESIDENT_RESULT*sizeof(float),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)){
+    seterr("exec_pose_lma_dp_cycles: DtoH failed"); return -7; }
+  if(cudaMemcpy(&p->pose_current_set,p->gpu_pose_current_set,sizeof(int),
+                cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_cycles: current_set DtoH failed"); return -7; }
+  return 0;
+}
+
 int tp_proc_exec_pose_lma_resident_step(TpProc* p,
                                         float lambda,
                                         const float* weights,
@@ -2529,6 +2640,7 @@ void tp_proc_destroy(TpProc* p){
   cudaFree(p->gpu_pose_lma_candidate_partials); // By Codex on 07/15/2026
   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_task_map); cudaFree(p->gpu_pose_task_centers); // By Claude on 07/16/2026 (rung B1)
   delete p;
 }

src/tp_lma.cu

@@ -1121,6 +1121,122 @@ extern "C" __global__ void pose_lma_dp_step(const PoseDpChain * c)
 			c->cand_partials, c->result, c->current_set);
 }
 
+// ── DP rung D1: K chained prepare-light + step cycles, fixed measurement ─────
+// See tp_lma.h (PoseDpCycles). The launcher replicates the host per-cycle
+// sequence EXACTLY: [clear measured y region = the light cudaMemset] ->
+// pose_lma_assemble(light) -> pose_fx_jacobian(current) ->
+// pose_lma_prepare_finish(light) -> the six spike step stages ->
+// pose_lma_dp_commit (anchor update on accept - the spike's noted gap - plus
+// the per-cycle trace snapshot) -> the NEXT launcher. The pose vectors need no
+// per-cycle upload: after an accept flip the current set already holds the
+// candidate built by pose_lma_dp_candidate_vector (bit-identical to the host's
+// re-upload); after a reject the current set is untouched. By Claude on 07/17/2026.
+
+// Device equivalent of the light-prepare cudaMemset on the measured y region.
+extern "C" __global__ void pose_lma_dp_clear(
+		int     n,
+		float * data)
+{
+	const int i = (int) (blockIdx.x * blockDim.x + threadIdx.x);
+	if (i < n) data[i] = 0.0f;
+}
+
+// Post-step commit: the linearization anchor must track the ACCEPTED pose
+// (a stale anchor makes every following candidate overshoot - the spike's
+// test-authoring lesson). Reject keeps the anchor, exactly like the host
+// re-uploading an unchanged parameters_vector. Optional 25-float snapshot
+// into the per-cycle trace (ruling b: minimal convenient trace).
+extern "C" __global__ void pose_lma_dp_commit(
+		const float * result,
+		float       * vector,
+		float       * trace)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	if (result[POSE_LMA_ACCEPTED] != 0.0f) {
+		#pragma unroll
+		for (int par = 0; par < POSE_NUM_PARAMS; par++) {
+			vector[par] = result[16 + par];
+		}
+	}
+	if (trace != NULL) {
+		for (int i = 0; i < POSE_LMA_RESIDENT_RESULT; i++) {
+			trace[i] = result[i];
+		}
+	}
+}
+
+extern "C" __global__ void pose_lma_dp_cycles(const PoseDpCycles * c, int cycle)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	// LIVE set read: this launcher runs AFTER the previous cycle's
+	// pose_lma_reduce_rms (it was tail-enqueued behind it), so it sees the flip.
+	const int set = *c->current_set;
+	float         * cur_vectors = c->set_vectors[set];
+	float         * cur_fx      = c->set_fx[set];
+	float         * cur_jt      = c->set_jt[set];
+	unsigned char * cur_valid   = c->set_valid[set];
+	float         * cand_vectors= c->set_vectors[set ^ 1];
+	float         * cand_fx     = c->set_fx[set ^ 1];
+	float         * cand_jt     = c->set_jt[set ^ 1];
+	unsigned char * cand_valid  = c->set_valid[set ^ 1];
+	const float   * candidate_vector = cand_vectors + 3 * POSE_NUM_PARAMS;
+	const int pose_threads = 256;
+	const int pose_blocks = (c->num_tiles + pose_threads - 1) / pose_threads;
+	const int a_threads = 256;                       // host assemble launch shape
+	const int a_blocks = (c->num_corr + a_threads - 1) / a_threads;
+	const int clear_n = c->num_components * c->num_tiles;
+	const int clear_blocks = (clear_n + 255) / 256;
+	const int light = 1;
+	// prepare-light (exec_pose_lma_prepare_resident light=1, same grids):
+	pose_lma_dp_clear<<<clear_blocks, 256, 0, cudaStreamTailLaunch>>>(
+			clear_n, c->y);
+	pose_lma_assemble<<<a_blocks, a_threads, 0, cudaStreamTailLaunch>>>(
+			c->num_corr, c->num_tiles, c->corr_slot, c->ntile_shift,
+			c->num_components, c->use_eigen, c->eig_min_sqrt, c->eig_max_sqrt,
+			c->min_confidence, c->same_weights, light,
+			c->corr_indices, c->peaks,
+			c->y, c->weights, c->eigen, c->selection);
+	pose_fx_jacobian<<<pose_blocks, pose_threads, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, 1, c->cameras,
+			c->ref_radial, c->ref_rbr, c->ref_ers,
+			c->scene_radial, c->scene_rbr, c->scene_ers,
+			cur_vectors, c->centers, c->selection,
+			cur_fx, cur_jt, cur_valid);
+	pose_lma_prepare_finish<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, c->num_components, POSE_NUM_PARAMS, light,
+			cur_fx, cur_jt, cur_valid, c->eigen, c->regweights,
+			c->y, c->weights, c->prep_result);
+	// the resident step (the six validated spike stages):
+	pose_lma_prepare_products<<<c->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, cur_fx, cur_jt, cur_valid,
+			c->y, c->weights, c->eigen, c->vector,
+			c->jt_out, c->ymfx_out, c->partials);
+	pose_lma_reduce_step<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_partials, c->lambda, c->pure_weight,
+			c->partials, c->vector, c->result);
+	pose_lma_dp_candidate_vector<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			cur_vectors, c->result, cand_vectors);
+	pose_fx_jacobian<<<pose_blocks, pose_threads, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, 1, c->cameras,
+			c->ref_radial, c->ref_rbr, c->ref_ers,
+			c->scene_radial, c->scene_rbr, c->scene_ers,
+			cand_vectors, c->centers, c->selection,
+			cand_fx, cand_jt, cand_valid);
+	pose_lma_prepare_products<<<c->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, cand_fx, cand_jt, cand_valid,
+			c->y, c->weights, c->eigen, candidate_vector,
+			c->cand_jt_out, c->cand_ymfx_out, c->cand_partials);
+	pose_lma_reduce_rms<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->num_partials, c->pure_weight, c->rms_diff,
+			c->cand_partials, c->result, c->current_set);
+	pose_lma_dp_commit<<<1, 1, 0, cudaStreamTailLaunch>>>(
+			c->result, c->vector,
+			(c->trace != NULL) ? (c->trace + (size_t) cycle * POSE_LMA_RESIDENT_RESULT) : NULL);
+	if (cycle + 1 < c->num_cycles) {
+		pose_lma_dp_cycles<<<1, 1, 0, cudaStreamTailLaunch>>>(c, cycle + 1);
+	}
+}
+
 #endif // __CUDACC_RTC__
 
 extern "C" __global__ void lma_normal_products(

src/tp_lma.h

@@ -257,6 +257,60 @@ typedef struct {
 	int   * current_set;       // device-authoritative accept flip (rung B)
 } PoseDpChain;
 
+// ── DP rung D1 (ladder ratified 2026-07-16): K chained prepare-light + step
+// cycles from ONE device-side entry, fixed measurement. The launcher kernel
+// pose_lma_dp_cycles resolves CURRENT/candidate from the LIVE *current_set at
+// each cycle start (a single unrolled parent cannot: tail-launch children run
+// after the parent exits, so cycle k+1's pointer choice depends on a flip the
+// parent has not seen - hence the SELF-CHAINING launcher, one per cycle, each
+// tail-enqueueing its stages and then the next launcher). Both physical buffer
+// sets cross in the struct; per-cycle H2D is ZERO. Fixed-K, no device early
+// exit (ruling a); per-cycle result snapshots to a minimal trace (ruling b).
+// By Claude on 07/17/2026.
+typedef struct {
+	int   num_cycles;
+	int   num_tiles;
+	int   num_corr;        // resident correlation rows (assemble grid)
+	int   num_partials;
+	float lambda;          // production policy: fixed lambda0, inner=1
+	float pure_weight;
+	float rms_diff;
+	// pose_lma_assemble(light) policy args (frozen conditioning, rung C3)
+	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;
+	const int   * corr_indices;
+	const float * peaks;
+	// camera/geometry (shared, per-sequence/per-scene resident)
+	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; // read-only in light cycles (frozen)
+	// FIXED physical slots (rung B): [0] = gpu_pose_*, [1] = gpu_pose_candidate_*;
+	// the launcher picks cur = set[*current_set], cand = the other, LIVE per cycle.
+	float         * set_vectors[2];
+	float         * set_fx[2];
+	float         * set_jt[2];
+	unsigned char * set_valid[2];
+	// shared LMA buffers (resident prepared state)
+	float * y; float * weights; float * eigen;
+	float * vector;            // linearization anchor (gpu_pose_lma_vector);
+	                           // updated DEVICE-SIDE on accept (the spike's noted gap)
+	const float * regweights;  // frozen; unused by light finish but part of its contract
+	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;       // light finish contract arg (untouched in light mode)
+	int   * current_set;
+	float * trace;             // [num_cycles][POSE_LMA_RESIDENT_RESULT], NULL = off
+} PoseDpCycles;
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,