CLAUDE: rung C2 - prepare upload groups nullable (configuration registers)

Design 3-A4i rung C2: tp_proc_exec_pose_lma_prepare_resident treats its
uploads as three nullable groups - NULL = keep the resident copy, guarded by
pose_prep_ready bits: reference camera block + centers (per-SEQUENCE), scene
camera block (per-SCENE). Per-cycle payload shrinks to pose_vectors + 6
policy floats + the 8-float result D2H. Tile-capacity regrows clear the
centers ready bit (both entry points). Camera meta now uploads per half.
Gates: full suite PASS (direct step+prepare bit-exact, NVRTC x3, pose_corr
@tol 0, sanitizer 0 errors).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

jna/tp_jna.cpp

@@ -691,6 +691,12 @@ struct TpProc {
   // Rung C1 resident prepare: compact result [8] + pull targets [3] + reg
   // weights [3] in one small device buffer. By Claude on 07/17/2026.
   float *gpu_pose_lma_prep;
+  // Rung C2: which prepare upload groups are resident ("configuration
+  // registers" - per-sequence/per-scene state uploaded once, NULL args on
+  // later calls mean "keep resident"). bit0 = reference camera block,
+  // bit1 = scene camera block, bit2 = centers. By Claude on 07/17/2026.
+  int pose_prep_ready;
+  int pose_prep_centers_tiles; // tile count the resident centers were uploaded for // By Claude on 07/17/2026
 };
 
 extern "C" {
@@ -743,6 +749,7 @@ TpProc* tp_proc_create(TpModule* m){
   p->gpu_pose_lma_candidate_jt=p->gpu_pose_lma_candidate_ymfx=nullptr;
   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; // By Claude on 07/17/2026 (rung C2)
   return p;
 }
 
@@ -1314,6 +1321,7 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p,
     }
   }
   if(num_tiles>p->pose_tile_capacity){
+    p->pose_prep_ready&=~4; p->pose_prep_centers_tiles=0; // regrow frees resident centers (rung C2) // By Claude on 07/17/2026
     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);
@@ -1449,29 +1457,40 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
                                            float* out_y,
                                            float* out_weights,
                                            float* out_eigen){
-  if(!p||!ref_meta||!ref_radial||!ref_rbr||!ref_ers||
-     !scene_meta||!scene_radial||!scene_rbr||!scene_ers||
-     !pose_vectors||!centers||!pull_targets||!param_regweights||!prep_result||
-     (num_tiles<=0)||(ref_rbr_len<3)||(scene_rbr_len<3)||
-     (num_components!=POSE_NUM_COMPONENTS)|| // contract carries it; this impl is 2-component
-     (capture&&(!out_y||!out_weights||!out_eigen))){
+  // Rung C2 ("configuration registers"): each upload group is nullable - NULL
+  // means "keep the resident copy" (guarded by pose_prep_ready bits). Groups:
+  // reference camera block + centers = per-sequence; scene camera block =
+  // per-scene; only pose_vectors + policy cross per cycle. By Claude on 07/17/2026.
+  const bool up_ref    = (ref_meta!=nullptr)||(ref_radial!=nullptr)||(ref_rbr!=nullptr)||(ref_ers!=nullptr);
+  const bool up_scene  = (scene_meta!=nullptr)||(scene_radial!=nullptr)||(scene_rbr!=nullptr)||(scene_ers!=nullptr);
+  const bool up_centers= (centers!=nullptr);
+  if(!p||!pose_vectors||!pull_targets||!param_regweights||!prep_result||
+     (num_tiles<=0)||(num_components!=POSE_NUM_COMPONENTS)|| // contract carries it; this impl is 2-component
+     (capture&&(!out_y||!out_weights||!out_eigen))||
+     (up_ref&&(!ref_meta||!ref_radial||!ref_rbr||!ref_ers||(ref_rbr_len<3)))||
+     (up_scene&&(!scene_meta||!scene_radial||!scene_rbr||!scene_ers||(scene_rbr_len<3)))){
     seterr("exec_pose_lma_prepare_resident: bad arguments n=%d nc=%d",num_tiles,num_components); return -1; }
+  if((!up_ref&&!(p->pose_prep_ready&1))||(!up_scene&&!(p->pose_prep_ready&2))||
+     (!up_centers&&(!(p->pose_prep_ready&4)||(p->pose_prep_centers_tiles!=num_tiles)))){
+    seterr("exec_pose_lma_prepare_resident: resident group missing (ready=%d ref=%d scene=%d centers=%d n=%d/%d)",
+           p->pose_prep_ready,up_ref,up_scene,up_centers,num_tiles,p->pose_prep_centers_tiles); return -1; }
   if(!p->gpu_corr_indices||!p->gpu_peaks||(p->last_num_corr_tiles<=0)){
     seterr("exec_pose_lma_prepare_resident: no resident correlation/peaks (num_corr=%d)",
            p->last_num_corr_tiles); return -2; }
-  const int ref_width=(int)std::lround(ref_meta[0]);
-  const int ref_height=(int)std::lround(ref_meta[1]);
-  const int scene_width=(int)std::lround(scene_meta[0]);
-  const int scene_height=(int)std::lround(scene_meta[1]);
-  if((ref_width<=0)||(ref_height<=0)||(scene_width<=0)||(scene_height<=0)||
-     (ref_ers_len!=ref_height*POSE_ERS_LINE_STRIDE)||
-     (scene_ers_len!=scene_height*POSE_ERS_LINE_STRIDE)||
-     !std::isfinite(ref_meta[7])||!(ref_meta[7]>0.0f)||
-     !std::isfinite(scene_meta[7])||!(scene_meta[7]>0.0f)){
-    seterr("exec_pose_lma_prepare_resident: dimensions/ERS/step mismatch"); return -3; }
-  const PoseCameraMeta hmeta[2] = {
-    {ref_width,ref_height,ref_rbr_len,ref_meta[2],ref_meta[3],ref_meta[4],ref_meta[5],ref_meta[6],ref_meta[7]},
-    {scene_width,scene_height,scene_rbr_len,scene_meta[2],scene_meta[3],scene_meta[4],scene_meta[5],scene_meta[6],scene_meta[7]}};
+  if(up_ref){
+    const int ref_width=(int)std::lround(ref_meta[0]);
+    const int ref_height=(int)std::lround(ref_meta[1]);
+    if((ref_width<=0)||(ref_height<=0)||(ref_ers_len!=ref_height*POSE_ERS_LINE_STRIDE)||
+       !std::isfinite(ref_meta[7])||!(ref_meta[7]>0.0f)){
+      seterr("exec_pose_lma_prepare_resident: ref dimensions/ERS/step mismatch"); return -3; }
+  }
+  if(up_scene){
+    const int scene_width=(int)std::lround(scene_meta[0]);
+    const int scene_height=(int)std::lround(scene_meta[1]);
+    if((scene_width<=0)||(scene_height<=0)||(scene_ers_len!=scene_height*POSE_ERS_LINE_STRIDE)||
+       !std::isfinite(scene_meta[7])||!(scene_meta[7]>0.0f)){
+      seterr("exec_pose_lma_prepare_resident: scene dimensions/ERS/step mismatch"); return -3; }
+  }
   cuCtxSetCurrent(p->mod->ctx);
   p->pose_raw_num_tiles=0; p->pose_raw_has_jacobian=0; // invalid until this call succeeds
   #define POSE_ALLOC(field,count,type) do{ if(!(p->field) && \
@@ -1490,8 +1509,8 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
       seterr("exec_pose_lma_prepare_resident: alloc current_set failed"); return -4; }
     p->pose_current_set=0;
   }
-  { const int rbr_lengths[2]={ref_rbr_len,scene_rbr_len};
-    const int ers_lengths[2]={ref_ers_len,scene_ers_len};
+  { const int rbr_lengths[2]={up_ref?ref_rbr_len:0,up_scene?scene_rbr_len:0};
+    const int ers_lengths[2]={up_ref?ref_ers_len:0,up_scene?scene_ers_len:0};
     for(int camera=0;camera<2;camera++){
       if(rbr_lengths[camera]>p->pose_rbr_capacity[camera]){
         cudaFree(p->gpu_pose_rbr[camera]); p->gpu_pose_rbr[camera]=nullptr; p->pose_rbr_capacity[camera]=0;
@@ -1508,6 +1527,7 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
     }
   }
   if(num_tiles>p->pose_tile_capacity){
+    p->pose_prep_ready&=~4; p->pose_prep_centers_tiles=0; // regrow frees resident centers // By Claude on 07/17/2026
     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);
@@ -1555,18 +1575,36 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
       seterr("exec_pose_lma_prepare_resident: eigen scratch alloc failed"); return -4; }
     p->pose_lma_tile_capacity=num_tiles;
   }
-  // Uploads: camera state, pose vector (CURRENT set), centers, policy floats.
+  // Uploads (rung C2): per-group, NULL = keep resident. Per-cycle payload is
+  // only pose_vectors + pull/reg (policy). By Claude on 07/17/2026.
   float* gpu_pull=p->gpu_pose_lma_prep+POSE_LMA_PREP_RESULT;
   float* gpu_regw=gpu_pull+POSE_NUM_PARAMS;
-  if(cudaMemcpy(p->gpu_pose_cameras,hmeta,sizeof(hmeta),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_radial[0],ref_radial,POSE_RADIAL_FLOATS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_radial[1],scene_radial,POSE_RADIAL_FLOATS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_rbr[0],ref_rbr,(size_t)ref_rbr_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_rbr[1],scene_rbr,(size_t)scene_rbr_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_ers[0],ref_ers,(size_t)ref_ers_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_ers[1],scene_ers,(size_t)scene_ers_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->cur_pose_vectors(),pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
-     cudaMemcpy(p->gpu_pose_centers,centers,(size_t)3*num_tiles*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
+  if(up_ref){
+    const PoseCameraMeta href={(int)std::lround(ref_meta[0]),(int)std::lround(ref_meta[1]),ref_rbr_len,
+      ref_meta[2],ref_meta[3],ref_meta[4],ref_meta[5],ref_meta[6],ref_meta[7]};
+    if(cudaMemcpy(p->gpu_pose_cameras,&href,sizeof(href),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_radial[0],ref_radial,POSE_RADIAL_FLOATS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_rbr[0],ref_rbr,(size_t)ref_rbr_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_ers[0],ref_ers,(size_t)ref_ers_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+      seterr("exec_pose_lma_prepare_resident: ref HtoD failed"); return -5; }
+    p->pose_prep_ready|=1;
+  }
+  if(up_scene){
+    const PoseCameraMeta hscene={(int)std::lround(scene_meta[0]),(int)std::lround(scene_meta[1]),scene_rbr_len,
+      scene_meta[2],scene_meta[3],scene_meta[4],scene_meta[5],scene_meta[6],scene_meta[7]};
+    if(cudaMemcpy(p->gpu_pose_cameras+1,&hscene,sizeof(hscene),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_radial[1],scene_radial,POSE_RADIAL_FLOATS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_rbr[1],scene_rbr,(size_t)scene_rbr_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
+       cudaMemcpy(p->gpu_pose_ers[1],scene_ers,(size_t)scene_ers_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+      seterr("exec_pose_lma_prepare_resident: scene HtoD failed"); return -5; }
+    p->pose_prep_ready|=2;
+  }
+  if(up_centers){
+    if(cudaMemcpy(p->gpu_pose_centers,centers,(size_t)3*num_tiles*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
+      seterr("exec_pose_lma_prepare_resident: centers HtoD failed"); return -5; }
+    p->pose_prep_ready|=4; p->pose_prep_centers_tiles=num_tiles;
+  }
+  if(cudaMemcpy(p->cur_pose_vectors(),pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
      cudaMemcpy(gpu_pull,pull_targets,POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
      cudaMemcpy(gpu_regw,param_regweights,POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess){
     seterr("exec_pose_lma_prepare_resident: HtoD failed"); return -5; }