CLAUDE: rung C3 - light prepare on frozen conditioning (Andrey's ruling)

Design 3-A4i rung C3 (Andrey 07/17: weights and their normalization once
per scene, not every cycle - re-normalizing on each cycle's re-measurement
makes the solver conditioning wobble with measurement noise while the
information is a per-scene property):
- pose_lma_assemble/pose_lma_prepare_finish gain a 'light' mode: cycle-1
  conditioning (weights/eigen/selection/pull/reg/pure_weight) is FROZEN;
  light cycles write only the fresh measured offsets into the frozen
  structure and add fx (a frozen-selected tile with no fresh valid peak
  abstains: y=fx, zero residual, normalization untouched).
- tp_proc_exec_pose_lma_prepare_resident(int light): guarded by
  pose_prep_frozen_tiles (a full prepare at this tile count must precede);
  light memsets only the measured y region, skips group/pull/reg uploads,
  skips the norm tail and diag/reg/normalize, skips the compact-result D2H.
- Direct test: new LIGHT case (fresh peaks incl. abstains + moved fx) -
  y bit-exact vs serial float oracle, weights/eigen bit-identical frozen.
Gates: direct full+LIGHT bit-exact, NVRTC resident + 19-float, pose_corr
@tol 0, sanitizer 0 errors.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

jna/tp_jna.cpp

@@ -697,6 +697,7 @@ struct TpProc {
   // 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
+  int pose_prep_frozen_tiles;  // rung C3: tile count of the frozen conditioning (0 = no full prepare yet) // By Claude on 07/17/2026
 };
 
 extern "C" {
@@ -749,7 +750,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)
+  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)
   return p;
 }
 
@@ -1452,6 +1453,7 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
                                            float min_confidence, int same_weights,
                                            const float* pull_targets,
                                            const float* param_regweights,
+                                           int light,
                                            float* prep_result,
                                            int capture,
                                            float* out_y,
@@ -1464,16 +1466,23 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
   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||
+  if(!p||!pose_vectors||
+     (!light&&(!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))||
+     (capture&&(light||!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; }
+    seterr("exec_pose_lma_prepare_resident: bad arguments n=%d nc=%d light=%d",num_tiles,num_components,light); 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(light&&(p->pose_prep_frozen_tiles!=num_tiles)){
+    // Rung C3: a light prepare reuses the FROZEN cycle-1 conditioning
+    // (weights/eigen/selection/pull/reg) - a full prepare must have run at
+    // this tile count first. By Claude on 07/17/2026.
+    seterr("exec_pose_lma_prepare_resident: light without frozen conditioning (%d/%d)",
+           p->pose_prep_frozen_tiles,num_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; }
@@ -1605,16 +1614,24 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
     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){
+     (!light&& // rung C3: pull/reg are part of the frozen conditioning // By Claude on 07/17/2026
+      (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; }
-  // Pre-zero the assembled buffers + selection + prep result.
-  if(cudaMemset(p->gpu_pose_lma_y,0,(size_t)num_values*sizeof(float))!=cudaSuccess ||
-     cudaMemset(p->gpu_pose_lma_weights,0,(size_t)num_values*sizeof(float))!=cudaSuccess ||
-     cudaMemset(p->gpu_pose_lma_eigen,0,(size_t)4*num_tiles*sizeof(float))!=cudaSuccess ||
-     cudaMemset(p->gpu_pose_selection,0,(size_t)num_tiles*sizeof(unsigned char))!=cudaSuccess ||
-     cudaMemset(p->gpu_pose_lma_prep,0,POSE_LMA_PREP_RESULT*sizeof(float))!=cudaSuccess){
-    seterr("exec_pose_lma_prepare_resident: memset failed"); return -6; }
+  // Pre-zero: full prepare clears everything; light clears ONLY the measured
+  // y region (pull rows + frozen conditioning survive). By Claude on 07/17/2026 (rung C3).
+  if(light){
+    if(cudaMemset(p->gpu_pose_lma_y,0,(size_t)POSE_NUM_COMPONENTS*num_tiles*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_lma_prepare_resident: light memset failed"); return -6; }
+  } else {
+    p->pose_prep_frozen_tiles=0; // redefining the frozen conditioning
+    if(cudaMemset(p->gpu_pose_lma_y,0,(size_t)num_values*sizeof(float))!=cudaSuccess ||
+       cudaMemset(p->gpu_pose_lma_weights,0,(size_t)num_values*sizeof(float))!=cudaSuccess ||
+       cudaMemset(p->gpu_pose_lma_eigen,0,(size_t)4*num_tiles*sizeof(float))!=cudaSuccess ||
+       cudaMemset(p->gpu_pose_selection,0,(size_t)num_tiles*sizeof(unsigned char))!=cudaSuccess ||
+       cudaMemset(p->gpu_pose_lma_prep,0,POSE_LMA_PREP_RESULT*sizeof(float))!=cudaSuccess){
+      seterr("exec_pose_lma_prepare_resident: memset failed"); return -6; }
+  }
   CUfunction assemble=getfun(p->mod,"pose_lma_assemble");
   CUfunction norm=getfun(p->mod,"pose_lma_prepare_norm");
   CUfunction pose=getfun(p->mod,"pose_fx_jacobian");
@@ -1623,16 +1640,18 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
   int num_corr=p->last_num_corr_tiles;
   int nc=num_components, np=POSE_NUM_PARAMS;
   void* aa[]={ &num_corr,&num_tiles,&corr_slot,&ntile_shift,&nc,&use_eigen,
-               &eig_min_sqrt,&eig_max_sqrt,&min_confidence,&same_weights,
+               &eig_min_sqrt,&eig_max_sqrt,&min_confidence,&same_weights,&light,
                &p->gpu_corr_indices,&p->gpu_peaks,
                &p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_eigen,
                &p->gpu_pose_selection };
   const int a_threads=256;
   if(launch1(assemble,(num_corr+a_threads-1)/a_threads,1,1,a_threads,1,1,aa,
              "pose_lma_assemble")) return -8;
-  void* na[]={ &num_tiles,&nc,&np,&gpu_pull,
-               &p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_prep };
-  if(launch1(norm,1,1,1,1,1,1,na,"pose_lma_prepare_norm")) return -9;
+  if(!light){ // rung C3: sample-weight normalization is frozen from cycle 1
+    void* na[]={ &num_tiles,&nc,&np,&gpu_pull,
+                 &p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_prep };
+    if(launch1(norm,1,1,1,1,1,1,na,"pose_lma_prepare_norm")) return -9;
+  }
   // Raw fx/J at the current pose - resident only, NO readback (that is the
   // whole point of this rung; the oracle path keeps exec_pose_fx_jacobian).
   float* cur_vectors=p->cur_pose_vectors();
@@ -1648,15 +1667,20 @@ int tp_proc_exec_pose_lma_prepare_resident(TpProc* p,
   const int f_threads=256;
   if(launch1(pose,(num_tiles+f_threads-1)/f_threads,1,1,f_threads,1,1,fa,
              "pose_fx_jacobian(prepare)")) return -10;
-  void* za[]={ &num_tiles,&nc,&np,&cur_fx,&cur_jt,&cur_valid,
+  void* za[]={ &num_tiles,&nc,&np,&light,&cur_fx,&cur_jt,&cur_valid,
                &p->gpu_pose_lma_eigen,&gpu_regw,
                &p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_prep };
   if(launch1(finish,1,1,1,1,1,1,za,"pose_lma_prepare_finish")) return -11;
+  if(light){ // rung C3: conditioning + compact result frozen from cycle 1
+    p->pose_raw_num_tiles=num_tiles; p->pose_raw_has_jacobian=1;
+    return 0;
+  }
   if(cudaMemcpy(prep_result,p->gpu_pose_lma_prep,
                 POSE_LMA_PREP_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess){
     seterr("exec_pose_lma_prepare_resident: result DtoH failed"); return -12; }
   if(prep_result[POSE_LMA_PREP_VALID]==0.0f){
     seterr("exec_pose_lma_prepare_resident: invalid prepare (sum weights?)"); return -13; }
+  p->pose_prep_frozen_tiles=num_tiles; // rung C3: conditioning frozen for light cycles // By Claude on 07/17/2026
   if(capture && // one-shot oracle D2H only (debug-gated by the caller) // By Claude on 07/17/2026
      (cudaMemcpy(out_y,p->gpu_pose_lma_y,(size_t)num_values*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess ||
       cudaMemcpy(out_weights,p->gpu_pose_lma_weights,(size_t)num_values*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess ||

src/tests/test_pose_lma_step.cu

@@ -399,11 +399,11 @@ int main(int argc, char ** argv)
 		checkCudaErrors(cudaMemset(gSel.dev(), 0, numTilesP));
 		checkCudaErrors(cudaMemset(gPrep.dev(), 0, POSE_LMA_PREP_RESULT * sizeof(float)));
 		pose_lma_assemble<<<(numCorr + 255) / 256, 256>>>(numCorr, numTilesP, slot, shift,
-				POSE_NUM_COMPONENTS, 1, eigMin, eigMax, minConf, 0,
+				POSE_NUM_COMPONENTS, 1, eigMin, eigMax, minConf, 0, 0,
 				gIdx.dev(), gPeaks.dev(), gY.dev(), gW.dev(), gE.dev(), gSel.dev());
 		pose_lma_prepare_norm<<<1, 1>>>(numTilesP, POSE_NUM_COMPONENTS, POSE_NUM_PARAMS,
 				gPull.dev(), gY.dev(), gW.dev(), gPrep.dev());
-		pose_lma_prepare_finish<<<1, 1>>>(numTilesP, POSE_NUM_COMPONENTS, POSE_NUM_PARAMS,
+		pose_lma_prepare_finish<<<1, 1>>>(numTilesP, POSE_NUM_COMPONENTS, POSE_NUM_PARAMS, 0,
 				gFx.dev(), gJt.dev(), gValid.dev(), gE.dev(), gReg.dev(),
 				gY.dev(), gW.dev(), gPrep.dev());
 		checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize());
@@ -432,6 +432,56 @@ int main(int argc, char ** argv)
 				nMeas, badBits, badSel, gotPrep[POSE_LMA_PREP_NUM_MEAS], nMeas,
 				gotPrep[POSE_LMA_PREP_PURE_WEIGHT], prepOk ? "PASS" : "FAIL");
 		if (!prepOk) bad++;
+
+		// ── Rung C3: LIGHT re-prepare on frozen conditioning ─────────────────
+		// Fresh measurements (some abstaining) + moved fx; weights/eigen/
+		// selection/pull must stay bit-identical, y = fresh offsets + fx (or fx
+		// alone where the fresh peak abstains). By Claude on 07/17/2026.
+		std::vector<float> peaks2(peaks);
+		for (int n = 0; n < numCorr; n++) {
+			float* row = &peaks2[(size_t) n * 8];
+			row[0] = nextf(-2.0f, 2.0f); row[1] = nextf(-2.0f, 2.0f);
+			if ((n % 23) == 11) row[7] = 0.0f; // fresh-cycle reject -> abstain (y = fx)
+		}
+		std::vector<float> rawFx2(rawFxP);
+		for (size_t i = 0; i < rawFx2.size(); i++) if (rawFx2[i] != 0.0f) rawFx2[i] += 0.125f;
+		std::vector<float> expY2(numValuesP, 0.0f);
+		for (int par = 0; par < POSE_NUM_PARAMS; par++) expY2[nsP + par] = pullT[par];
+		for (int n = 0; n < numCorr; n++) {
+			if ((corrIdx[n] & ((1 << shift) - 1)) != slot) continue;
+			const int tile = corrIdx[n] >> shift;
+			const float* row = &peaks2[(size_t) n * 8];
+			if (row[7] == 0.0f) continue;
+			const float str = row[2];
+			if (!(str >= minConf) || !(str > 0.0f) || !expSel[tile]) continue;
+			const int vi = POSE_NUM_COMPONENTS * tile;
+			expY2[vi] = row[0]; expY2[vi + 1] = row[1];
+		}
+		for (int t = 0; t < numTilesP; t++) { const int vi = POSE_NUM_COMPONENTS * t;
+			if ((expW[vi] > 0.0f) && validP[t]) {
+				expY2[vi] = expY2[vi] + rawFx2[vi]; expY2[vi + 1] = expY2[vi + 1] + rawFx2[vi + 1]; } }
+		gPeaks.upload(peaks2.data()); gFx.upload(rawFx2.data());
+		checkCudaErrors(cudaMemset(gY.dev(), 0, (size_t) nsP * sizeof(float))); // measured region only
+		pose_lma_assemble<<<(numCorr + 255) / 256, 256>>>(numCorr, numTilesP, slot, shift,
+				POSE_NUM_COMPONENTS, 1, eigMin, eigMax, minConf, 0, 1,
+				gIdx.dev(), gPeaks.dev(), gY.dev(), gW.dev(), gE.dev(), gSel.dev());
+		pose_lma_prepare_finish<<<1, 1>>>(numTilesP, POSE_NUM_COMPONENTS, POSE_NUM_PARAMS, 1,
+				gFx.dev(), gJt.dev(), gValid.dev(), gE.dev(), gReg.dev(),
+				gY.dev(), gW.dev(), gPrep.dev());
+		checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize());
+		const std::vector<float> gotY2 = gY.download();
+		const std::vector<float> gotW2 = gW.download();
+		const std::vector<float> gotE2 = gE.download();
+		int badLight = 0;
+		for (int i = 0; i < numValuesP; i++) {
+			if (std::memcmp(&gotY2[i], &expY2[i], sizeof(float)) != 0) badLight++;
+			if (std::memcmp(&gotW2[i], &expW[i], sizeof(float)) != 0) badLight++; // frozen
+		}
+		for (size_t i = 0; i < expE.size(); i++)
+			if (std::memcmp(&gotE2[i], &expE[i], sizeof(float)) != 0) badLight++;  // frozen
+		std::printf("pose_lma prepare LIGHT (frozen conditioning): bit mismatches=%d -> %s\n",
+				badLight, (badLight == 0) ? "PASS" : "FAIL");
+		if (badLight) bad++;
 	}
 	return bad ? EXIT_FAILURE : EXIT_SUCCESS;
 }

src/tp_lma.cu

@@ -786,6 +786,7 @@ extern "C" __global__ void pose_lma_assemble(
 		float                 eig_max_sqrt,
 		float                 min_confidence,
 		int                   same_weights,
+		int                   light,          // rung C3: 1 = frozen weights/eigen/selection, write fresh y offsets only // By Claude on 07/17/2026
 		const int           * corr_indices,
 		const float         * peaks,
 		float               * y,
@@ -803,6 +804,18 @@ extern "C" __global__ void pose_lma_assemble(
 	if (peaks[row + 7] == 0.0f) return;               // kernel-rejected peak
 	const float str = peaks[row + 2];
 	if (!(str >= min_confidence)) return;             // NaN strength fails here too
+	if (light) {
+		// Rung C3: cycle-1 conditioning is FROZEN (weights/eigen/selection) so the
+		// objective does not wobble with per-cycle measurement noise; only the
+		// fresh measured offsets enter. A frozen-selected tile with no fresh
+		// valid peak keeps y at the memset 0 and abstains (y becomes fx after
+		// pose_lma_prepare_finish -> zero residual). By Claude on 07/17/2026.
+		if (!(str > 0.0f) || (selection[tile] == 0)) return;
+		const int lvi = num_components * tile;
+		y[lvi]     = peaks[row];
+		y[lvi + 1] = peaks[row + 1];
+		return;
+	}
 	if (use_eigen) {
 		// Java setEigenTransform: a NaN eigenvector leaves transform null and
 		// setSamplesWeights then SKIPS the tile entirely. Non-finite lambdas are
@@ -877,6 +890,7 @@ extern "C" __global__ void pose_lma_prepare_finish(
 		int                   num_tiles,
 		int                   num_components,
 		int                   num_params,
+		int                   light,          // rung C3: 1 = y += fx only, conditioning frozen // By Claude on 07/17/2026
 		const float         * raw_fx,
 		const float         * raw_jt,
 		const unsigned char * valid,
@@ -898,6 +912,7 @@ extern "C" __global__ void pose_lma_prepare_finish(
 			y[vi + 1] = pose_lma_add(y[vi + 1], raw_fx[vi + 1]);
 		}
 	}
+	if (light) return; // rung C3: weights/reg/normalization frozen from cycle 1
 	// H diagonal over the eigen-transformed J, Java getWJtJlambda order (per
 	// parameter the contributions arrive in ascending value order; reg rows
 	// carry zero weight at this point and contribute exactly 0).

src/tp_lma.h

@@ -157,12 +157,13 @@ extern "C" __global__ void pose_lma_assemble(
 		float                 eig_max_sqrt,
 		float                 min_confidence,
 		int                   same_weights,
+		int                   light,          // rung C3: 1 = frozen conditioning, fresh y offsets only // By Claude on 07/17/2026
 		const int           * corr_indices,   // [num_corr] packed (tile<<shift)|slot
 		const float         * peaks,          // [num_corr][8]
-		float               * y,              // [nc*num_tiles+num_params] pre-zeroed
-		float               * weights,        // [nc*num_tiles+num_params] pre-zeroed
-		float               * eigen,          // [4*num_tiles] pre-zeroed
-		unsigned char       * selection);     // [num_tiles] pre-zeroed
+		float               * y,              // [nc*num_tiles+num_params] pre-zeroed (light: measured region pre-zeroed)
+		float               * weights,        // [nc*num_tiles+num_params] pre-zeroed (light: read-only frozen)
+		float               * eigen,          // [4*num_tiles] pre-zeroed (light: untouched frozen)
+		unsigned char       * selection);     // [num_tiles] pre-zeroed (light: read-only frozen)
 
 // Serial tail 1: sample-weight normalization (sum 2w -> scale, duplicate the
 // second component) + pull-target y rows; counts measured tiles.
@@ -183,6 +184,7 @@ extern "C" __global__ void pose_lma_prepare_finish(
 		int                   num_tiles,
 		int                   num_components,
 		int                   num_params,
+		int                   light,          // rung C3: 1 = y+=fx only // By Claude on 07/17/2026
 		const float         * raw_fx,
 		const float         * raw_jt,          // parameter-major [np][nc*num_tiles]
 		const unsigned char * valid,