CLAUDE: DP spike PASS - CDP2 device-side pose-LMA step bit-identical on Blackwell

pose_lma_dp_step<<<1,1>>> tail-launches all six resident-step stages
(cudaStreamTailLaunch, CDP2) with zero host round trips; the only new
device work is pose_lma_dp_candidate_vector (replaces two host
cudaMemcpyAsync). NVRTC-only (#ifdef __CUDACC_RTC__): the production
module's --extensible-whole-program + libcudadevrt link is exactly the
CDP2 requirement; direct nvcc test builds never see device-side launches.

Feasibility verdict (the spike's deliverable): CDP2 JIT+link+run WORKS in
the production NVRTC module on CC 12.0 - no flag changes needed.

test_pose_lma_dp_jna: two TpProc instances from one module driven to the
identical pre-step resident state (the step mutates state on accept, so
one-instance host-then-DP would compare different starting states); host
prepared-resident path vs DP chain memcmp-equal on packed[0..24] for
accept (device set flip exercised under DP), reject, and post-flip reject;
multi-block child grid (300 tiles -> 2 blocks) + 3-partial reduce.
Wall/step always-reject microbench: host 132 us, DP 118 us (1.12x; the
real win is removing 6 JNA launch/sync round trips per step in production).

Gates: test_pose_lma_dp_jna PASS bit-identical; regressions PASS:
pose_lma_resident, pose_fx_jacobian, lma_products, pose_lma_step (all
NVRTC), pose_corr captured-data @tol 0.

Production notes for the DP milestone: the DP entry reuses the resident
gpu_pose_lma_vector (anchor) - the per-cycle vector update after accept
must move device-side (from result[16..18]) before multi-step DP chains.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

jna/test_pose_lma_dp_jna.cu

+// test_pose_lma_dp_jna.cu - CDP2 feasibility gate: the DP (device-side
+// tail-launch) resident pose-LMA step must be BIT-IDENTICAL to the host-driven
+// resident step. Two fresh TpProc instances are built from the SAME NVRTC
+// module and driven to the identical pre-step resident state (the resident
+// step MUTATES state on accept - current/candidate flip - so "host step then
+// DP step" on one instance would compare different starting states). Instance
+// A then advances with tp_proc_exec_pose_lma_resident_step (prepared-resident
+// path, no per-step H2D of weights/y/eigen), instance B with
+// tp_proc_exec_pose_lma_dp_step. Same kernels, same inputs, same enqueue
+// order - only host-sequenced vs device-tail-launched differs, which changes
+// nothing numeric: packed[0..24] must memcmp equal at every step (accept AND
+// reject paths). Any diff = a real DP bug (chain-descriptor marshalling or a
+// stale pointer after an accept flip). Also reports a host-vs-DP wall-time
+// read on a stateless (always-reject) step. By Claude on 07/16/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_jna.cu -o tests_bin/test_pose_lma_dp_jna \
+//     -L jna -ltileproc -Xlinker -rpath -Xlinker $PWD/jna
+// Run from repo root: ./tests_bin/test_pose_lma_dp_jna src
+
+#include <chrono>
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <vector>
+
+enum { NUM_PARAMS=3, NUM_COMPONENTS=2, RESULT=25 };
+
+extern "C" {
+  void* tp_create_module(const char*, const char*);
+  const char* tp_last_error();
+  int tp_module_num_functions(void*);
+  void* tp_proc_create(void*);
+  int tp_proc_setup(void*,int,int,int,int,int,int);
+  int tp_proc_exec_pose_fx_jacobian(
+      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*,const unsigned char*,int,int,
+      float*,float*,unsigned char*);
+  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_step(void*,float,float,float,int,float*);
+  void tp_proc_destroy(void*);
+  void tp_destroy_module(void*);
+}
+
+// Bit-compare two packed[25] results; print every differing slot.
+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;
+}
+
+int main(int argc,char** argv)
+{
+  const char* srcdir=(argc>1)?argv[1]:"src";
+  const char* devrt=(argc>2)?argv[2]:
+      "/usr/local/cuda-12.8/targets/x86_64-linux/lib/libcudadevrt.a";
+  void* module=tp_create_module(srcdir,devrt);
+  if(!module){ std::fprintf(stderr,"tp_create_module: %s\n",tp_last_error()); return EXIT_FAILURE; }
+  std::printf("module kernels resolved: %d\n",tp_module_num_functions(module));
+
+  // Two instances from one module: A = host-driven oracle, B = DP path.
+  void* procs[2]={nullptr,nullptr};
+  for(int k=0;k<2;k++){
+    procs[k]=tp_proc_create(module);
+    if(!procs[k]||tp_proc_setup(procs[k],1,1,16,16,0,0)){
+      std::fprintf(stderr,"tp_proc setup[%d]: %s\n",k,tp_last_error()); return EXIT_FAILURE; }
+  }
+
+  // Synthetic case sized to exercise MULTI-BLOCK child grids (pose_fx_jacobian
+  // (300+255)/256 = 2 blocks) and multi-partial reduction (300/128 -> 3).
+  const int width=64,height=48,rbrLen=256,ersLen=height*14,numTiles=300;
+  const int numValues=NUM_COMPONENTS*numTiles+NUM_PARAMS;
+  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 vector[NUM_PARAMS]={0.0023f,-0.0017f,0.00091f};
+  const float pose[12]={0,0,0, 0,0,0, 0,0,0, vector[0],vector[1],vector[2]};
+  std::vector<float> centers(3*numTiles);
+  std::vector<unsigned char> selection(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);
+    selection[i]=(i%37)!=0; // a few unselected tiles
+  }
+
+  int bad=0;
+  std::vector<float> rawFx(2*numTiles*2),rawJt((size_t)NUM_PARAMS*2*numTiles*2);
+  std::vector<unsigned char> valid(numTiles*2);
+  for(int k=0;k<2;k++){
+    int rc=tp_proc_exec_pose_fx_jacobian(
+        procs[k],meta,radial,rbr.data(),rbrLen,ers.data(),ersLen,
+        meta,radial,rbr.data(),rbrLen,ers.data(),ersLen,
+        pose,centers.data(),selection.data(),numTiles,1,
+        rawFx.data()+(size_t)k*2*numTiles,
+        rawJt.data()+(size_t)k*NUM_PARAMS*2*numTiles,
+        valid.data()+(size_t)k*numTiles);
+    if(rc){ std::fprintf(stderr,"pose jacobian[%d] rc=%d: %s\n",rc,k,tp_last_error()); bad++; }
+  }
+  if(std::memcmp(rawFx.data(),rawFx.data()+(size_t)2*numTiles,(size_t)2*numTiles*sizeof(float))||
+     std::memcmp(rawJt.data(),rawJt.data()+(size_t)NUM_PARAMS*2*numTiles,
+                 (size_t)NUM_PARAMS*2*numTiles*sizeof(float))||
+     std::memcmp(valid.data(),valid.data()+numTiles,numTiles)){
+    std::printf("instance raw-state divergence before any step\n"); bad++;
+  }
+
+  // Measurement/regularization inputs: candidate must improve on step 1.
+  const float pureWeight=0.98f,rmsDiff=0.001f;
+  std::vector<float> weights(numValues,0.0f),y(numValues,0.0f),eigen(4*numTiles,0.0f);
+  for(int tile=0;tile<numTiles;tile++){
+    const int vi=2*tile,ei=4*tile;
+    if((tile%41)==0) continue; // a few zero-weight tiles (stay 0)
+    weights[vi]=weights[vi+1]=0.245f/numTiles*3.0f; // keep sums near the 3-tile case
+    if(selection[tile]){
+      y[vi]  =rawFx[vi]  +0.031f-0.0001f*(float)tile;
+      y[vi+1]=rawFx[vi+1]-0.022f+0.00007f*(float)tile;
+    }else{
+      // unselected tile kept weighted: exercises the have=false (fx treated as
+      // 0) branch - y must NOT come from rawFx, which is undefined here
+      y[vi]=0.03f; y[vi+1]=-0.02f;
+    }
+    eigen[ei]=0.82f; eigen[ei+1]=0.13f; eigen[ei+2]=0.07f; eigen[ei+3]=-0.61f;
+  }
+  for(int par=0;par<NUM_PARAMS;par++){
+    const int value=2*numTiles+par;
+    weights[value]=0.0066666668f;
+    y[value]=vector[par]+(par-1)*0.0002f;
+  }
+
+  // Step 1 (BOTH instances host-driven, identical): establishes the resident
+  // prepared buffers + advances both to the identical post-accept state.
+  float packed1[2][RESULT];
+  for(int k=0;k<2;k++){
+    int rc=tp_proc_exec_pose_lma_resident_step(procs[k],0.001f,
+        weights.data(),y.data(),eigen.data(),vector,
+        pureWeight,rmsDiff,numTiles,0,packed1[k],nullptr,nullptr);
+    if(rc){ std::fprintf(stderr,"host step1[%d] rc=%d: %s\n",k,rc,tp_last_error()); bad++; }
+  }
+  if(!bad){
+    if(comparePacked("step1(lockstep)",packed1[0],packed1[1])){ bad++; }
+    if(packed1[0][23]!=1.0f){
+      std::printf("step1 not accepted (accepted=%g current=%g candidate=%g) - "
+          "test setup broken\n",packed1[0][23],packed1[0][19],packed1[0][21]); bad++; }
+  }
+
+  // Step 1 consumed the improvement, so from here every same-y step rejects
+  // and the DP accept path (pose_lma_dp_candidate_vector + device set flip)
+  // would never be compared. Shift y and re-upload it through a LOCKSTEP
+  // host-host step on both instances at a near-infinite lambda (candidate ~=
+  // current, state advance identical either way): the next normal-lambda step
+  // then has a real improvement to accept.
+  std::vector<float> y2(y);
+  for(int tile=0;tile<numTiles;tile++){
+    const int vi=2*tile;
+    if(weights[vi]<=0.0f) continue;
+    y2[vi]  +=0.017f-0.00005f*(float)tile;
+    y2[vi+1]-=0.013f-0.00004f*(float)tile;
+  }
+  // The accepted pose from step 1 becomes the linearization anchor: the solved
+  // candidate is vector+delta, so a stale anchor makes every step overshoot.
+  // The non-NULL host path of the y-shift step uploads it to BOTH instances'
+  // resident gpu_pose_lma_vector, which the DP chain then reuses.
+  const float vNew[NUM_PARAMS]={packed1[0][16],packed1[0][17],packed1[0][18]};
+  if(!bad){
+    float packedS[2][RESULT];
+    for(int k=0;k<2;k++){
+      int rc=tp_proc_exec_pose_lma_resident_step(procs[k],1.0e6f,
+          weights.data(),y2.data(),eigen.data(),vNew,
+          pureWeight,rmsDiff,numTiles,0,packedS[k],nullptr,nullptr);
+      if(rc){ std::fprintf(stderr,"host y-shift step[%d] rc=%d: %s\n",k,rc,tp_last_error()); bad++; }
+    }
+    if(!bad&&comparePacked("y-shift(lockstep)",packedS[0],packedS[1])) bad++;
+  }
+
+  // Steps 2..4: A host prepared-resident path vs B DP chain, exercising
+  // accept (set flip) and reject (fallback retained) on both set states.
+  // Vector stays v1 on BOTH paths: the DP entry reuses the resident
+  // gpu_pose_lma_vector (production will need a device-side vector update).
+  const float lambdas[3]={0.0005f,100.0f,0.0005f}; // accept, reject, post-flip
+  int dpAccepts=0;
+  for(int step=0;step<3&&!bad;step++){
+    float packedA[RESULT],packedB[RESULT];
+    int rcA=tp_proc_exec_pose_lma_resident_step(procs[0],lambdas[step],
+        nullptr,nullptr,nullptr,vNew,pureWeight,rmsDiff,numTiles,0,
+        packedA,nullptr,nullptr);
+    int rcB=tp_proc_exec_pose_lma_dp_step(procs[1],lambdas[step],
+        pureWeight,rmsDiff,numTiles,packedB);
+    if(rcA||rcB){
+      std::fprintf(stderr,"step%d rcA=%d rcB=%d: %s\n",step+2,rcA,rcB,tp_last_error());
+      bad++; break; }
+    char tag[32]; std::snprintf(tag,sizeof(tag),"step%d(lambda=%g)",step+2,lambdas[step]);
+    if(comparePacked(tag,packedA,packedB)) bad++;
+    if(packedB[23]==1.0f) dpAccepts++;
+    std::printf("step%d lambda=%g accepted=%g rms %g -> %g : %s\n",
+        step+2,lambdas[step],packedB[23],packedB[19],packedB[21],
+        bad?"DIFF":"bit-identical");
+  }
+  if(!bad&&!dpAccepts){
+    std::printf("no DP step accepted - candidate-vector kernel + device flip "
+        "never exercised under DP\n"); bad++;
+  }
+
+  // Wall-time read (secondary deliverable): always-reject lambda leaves state
+  // untouched, so repeated calls do identical work. Includes each path's own
+  // per-call H2D/D2H (host: 12B vector up; DP: chain descriptor up).
+  if(!bad){
+    const int reps=500; const float bigLambda=1.0e6f;
+    float packedT[RESULT];
+    for(int k=0;k<2;k++){ // warm both paths once
+      if(k) tp_proc_exec_pose_lma_dp_step(procs[1],bigLambda,pureWeight,rmsDiff,numTiles,packedT);
+      else  tp_proc_exec_pose_lma_resident_step(procs[0],bigLambda,nullptr,nullptr,nullptr,
+                vNew,pureWeight,rmsDiff,numTiles,0,packedT,nullptr,nullptr);
+    }
+    auto t0=std::chrono::steady_clock::now();
+    for(int r=0;r<reps;r++)
+      tp_proc_exec_pose_lma_resident_step(procs[0],bigLambda,nullptr,nullptr,nullptr,
+          vNew,pureWeight,rmsDiff,numTiles,0,packedT,nullptr,nullptr);
+    auto t1=std::chrono::steady_clock::now();
+    for(int r=0;r<reps;r++)
+      tp_proc_exec_pose_lma_dp_step(procs[1],bigLambda,pureWeight,rmsDiff,numTiles,packedT);
+    auto t2=std::chrono::steady_clock::now();
+    const double usHost=std::chrono::duration<double,std::micro>(t1-t0).count()/reps;
+    const double usDp  =std::chrono::duration<double,std::micro>(t2-t1).count()/reps;
+    std::printf("wall/step over %d always-reject reps: host-driven %.1f us, DP %.1f us (%.2fx)\n",
+        reps,usHost,usDp,usHost/usDp);
+  }
+
+  std::printf("CDP2 DP pose-LMA step vs host-driven resident step -> %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

@@ -57,6 +57,7 @@ static const char* KERNELS[] = {
   "pose_lma_step", // fixed 3-angle float products + direct solve in one thread // By Codex on 07/15/2026
   "pose_lma_prepare_products","pose_lma_reduce_step","pose_lma_reduce_rms", // resident prepare/solve + candidate RMS
   "pose_lma_assemble","pose_lma_prepare_norm","pose_lma_prepare_finish", // rung C1 resident prepareLMA // By Claude on 07/17/2026
+  "pose_lma_dp_step","pose_lma_dp_candidate_vector", // DP spike: device-side step chain // By Claude on 07/16/2026
   "lma_normal_products"}; // deterministic double normal-equation products // By Codex on 07/14/2026
 static const int N_KERNELS = sizeof(KERNELS)/sizeof(KERNELS[0]);
 
@@ -698,6 +699,7 @@ struct TpProc {
   int pose_prep_ready;
   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
 };
 
 extern "C" {
@@ -751,6 +753,7 @@ TpProc* tp_proc_create(TpModule* m){
   p->gpu_pose_lma_candidate_partials=nullptr; // By Codex on 07/15/2026
   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)
   return p;
 }
 
@@ -1690,6 +1693,65 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
   return 0;
 }
 
+// DP spike (design 2026-07-16): the resident step as ONE device-side launch.
+// Preconditions: a resident prepare (weights/y/eigen/vector) + valid raw pose
+// state at this tile count are live (same as a prepared-resident step with NULL
+// H2D). Fills the chain descriptor host-side, copies it once, launches
+// pose_lma_dp_step<<<1,1>>>; the parent tail-launches all six stages with no
+// host round trip. D2H = the 25-float result + the 4-byte current_set (rung B),
+// exactly as the host-driven step. This is the feasibility spike, not yet the
+// per-scene production entry. By Claude on 07/16/2026.
+int tp_proc_exec_pose_lma_dp_step(TpProc* p,
+                                  float lambda,
+                                  float pure_weight,
+                                  float rms_diff,
+                                  int num_tiles,
+                                  float* result){
+  if(!p||!result||(num_tiles<=0)||!std::isfinite(pure_weight)||!(pure_weight>0.0f)||
+     !std::isfinite(rms_diff)){
+    seterr("exec_pose_lma_dp_step: bad args nt=%d",num_tiles); 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){
+    seterr("exec_pose_lma_dp_step: resident prepared state not ready (nt=%d)",num_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_chain &&
+     cudaMalloc((void**)&p->gpu_pose_dp_chain,sizeof(PoseDpChain))!=cudaSuccess){
+    seterr("exec_pose_lma_dp_step: chain alloc failed"); return -3; }
+  PoseDpChain h;
+  h.num_tiles=num_tiles; h.num_partials=num_partials;
+  h.lambda=lambda; h.pure_weight=pure_weight; h.rms_diff=rms_diff;
+  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.cur_vectors=p->cur_pose_vectors(); h.cur_fx=p->cur_pose_fx();
+  h.cur_jt=p->cur_pose_jt();           h.cur_valid=p->cur_pose_valid();
+  h.cand_vectors=p->cand_pose_vectors();h.cand_fx=p->cand_pose_fx();
+  h.cand_jt=p->cand_pose_jt();          h.cand_valid=p->cand_pose_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.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.current_set=p->gpu_pose_current_set;
+  if(cudaMemcpy(p->gpu_pose_dp_chain,&h,sizeof(PoseDpChain),cudaMemcpyHostToDevice)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_step: chain HtoD failed"); return -4; }
+  CUfunction dp=getfun(p->mod,"pose_lma_dp_step");
+  if(!dp){ seterr("pose_lma_dp_step missing (DP kernel not in module)"); return -5; }
+  void* da[]={ &p->gpu_pose_dp_chain };
+  if(launch1(dp,1,1,1,1,1,1,da,"pose_lma_dp_step")) return -6;
+  if(cudaMemcpy(result,p->gpu_pose_lma_result,
+                POSE_LMA_RESIDENT_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_step: result DtoH failed"); return -7; }
+  if(cudaMemcpy(&p->pose_current_set,p->gpu_pose_current_set,sizeof(int),
+                cudaMemcpyDeviceToHost)!=cudaSuccess){
+    seterr("exec_pose_lma_dp_step: current_set DtoH failed"); return -7; }
+  return 0;
+}
+
 int tp_proc_exec_pose_lma_resident_step(TpProc* p,
                                         float lambda,
                                         const float* weights,
@@ -2105,6 +2167,7 @@ void tp_proc_destroy(TpProc* p){
   cudaFree(p->gpu_pose_lma_candidate_jt); cudaFree(p->gpu_pose_lma_candidate_ymfx);
   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)
   delete p;
 }
 

src/tp_lma.cu

@@ -955,6 +955,68 @@ extern "C" __global__ void pose_lma_prepare_finish(
 	prep_result[POSE_LMA_PREP_VALID] = (isfinite(s) && (full > 0.0f)) ? 1.0f : 0.0f;
 }
 
+// ── DP spike: the resident step as ONE device-side chain ─────────────────────
+// See tp_lma.h. NVRTC-only: the production JIT path (--extensible-whole-program
+// + libcudadevrt link) is CDP2-capable; direct nvcc test builds are not -rdc
+// and must not see device-side launches. By Claude on 07/16/2026.
+#ifdef __CUDACC_RTC__
+
+// Candidate pose-vector build: ref rows copied from the CURRENT set (pose-
+// independent), scene rows overwritten with the solved candidate (result[16..18]).
+// Replaces the two host cudaMemcpyAsync of the host-driven step.
+extern "C" __global__ void pose_lma_dp_candidate_vector(
+		const float * cur_vectors,
+		const float * result,
+		float       * cand_vectors)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	for (int i = 0; i < 4 * POSE_NUM_PARAMS; i++) {
+		cand_vectors[i] = cur_vectors[i];
+	}
+	for (int par = 0; par < POSE_NUM_PARAMS; par++) {
+		cand_vectors[3 * POSE_NUM_PARAMS + par] = result[16 + par];
+	}
+}
+
+#ifndef cudaStreamTailLaunch
+#define cudaStreamTailLaunch ((cudaStream_t)0x3) // CDP2, cuda_device_runtime_api.h
+#endif
+
+extern "C" __global__ void pose_lma_dp_step(const PoseDpChain * c)
+{
+	if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
+	const int pose_threads = 256;
+	const int pose_blocks = (c->num_tiles + pose_threads - 1) / pose_threads;
+	const float * candidate_vector = c->cand_vectors + 3 * POSE_NUM_PARAMS;
+	// All six stages go into THIS grid's tail stream: they start after this
+	// parent exits and run serialized in enqueue order - the whole step with
+	// zero host round trips.
+	pose_lma_prepare_products<<<c->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, c->cur_fx, c->cur_jt, c->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>>>(
+			c->cur_vectors, c->result, c->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,
+			c->cand_vectors, c->centers, c->selection,
+			c->cand_fx, c->cand_jt, c->cand_valid);
+	pose_lma_prepare_products<<<c->num_partials, POSE_LMA_REDUCE_THREADS, 0, cudaStreamTailLaunch>>>(
+			c->num_tiles, c->cand_fx, c->cand_jt, c->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);
+}
+
+#endif // __CUDACC_RTC__
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,

src/tp_lma.h

@@ -194,6 +194,37 @@ extern "C" __global__ void pose_lma_prepare_finish(
 		float               * weights,
 		float               * prep_result);    // [POSE_LMA_PREP_RESULT]
 
+// ── DP spike (design 2026-07-16, the DP endgame seed) ────────────────────────
+// One device-side parent enqueues the WHOLE resident step into its tail-launch
+// stream (CDP2: tail launches from one thread serialize in enqueue order after
+// the parent exits) - no intermediate CPU-GPU synchronization. The payload
+// kernels are the UNMODIFIED validated ones; the only new device work is the
+// 15-float candidate-vector build that used to be two host cudaMemcpyAsync.
+// NVRTC-only (#ifdef __CUDACC_RTC__ in tp_lma.cu): the production module
+// compiles --extensible-whole-program and links libcudadevrt, which is exactly
+// the CDP2 requirement; direct nvcc test builds stay non-rdc and bit-stable.
+// By Claude on 07/16/2026.
+typedef struct {
+	int   num_tiles;
+	int   num_partials;
+	float lambda;
+	float pure_weight;
+	float rms_diff;
+	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;     const unsigned char * selection;
+	const float * cur_vectors; const float * cur_fx;
+	const float * cur_jt;      const unsigned char * cur_valid;
+	float * cand_vectors; float * cand_fx; float * cand_jt; unsigned char * cand_valid;
+	const float * y; const float * weights; const float * eigen;
+	const float * vector;      // current 3-parameter vector (gpu_pose_lma_vector)
+	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]
+	int   * current_set;       // device-authoritative accept flip (rung B)
+} PoseDpChain;
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,