Skip to content

T
tile_processor_gpu
Commit 5f473261 authored by Andrey Filippov's avatar Andrey Filippov
Browse files
Options

CLAUDE: rung B - device-visible pose current/candidate set index (no host swap) 

Design 3-A4i rung B (imagej-elphel-internal 2026-07-16 handoff), DP-readiness:
- gpu_pose_* / gpu_pose_candidate_* become FIXED physical slots; the CURRENT
  set is selected by a device-authoritative int flag.
- pose_lma_reduce_rms gains a final int* current_set arg and flips it on
  accept DEVICE-SIDE (NULL = no flip, standalone tests); rejection leaves it
  unchanged so the prior current raw pose remains the fallback.
- tp_proc_exec_pose_lma_resident_step: host std::swap removed; launch
  pointers come from cur_*/cand_* accessors; the host mirror syncs with a
  4-byte D2H after each step. exec_pose_fx_jacobian uploads/reads via the
  CURRENT set. The candidate-base DtoD invariant (ref rows pose-independent)
  is now documented at the copy site.
- test_pose_lma_step: pass + assert the no-flip side (not-accepted case).
C API unchanged. Gates: direct CUDA 5120/150 PASS (products 0.000244141,
candidate 9.69e-08, same as session-40); NVRTC resident accept-carryover +
worsened-rejection PASS (flip proven behaviorally); 19-float bit-exact PASS;
fx_jacobian PASS; compute-sanitizer(12.8) 0 errors; test_pose_corr_jna
bit-exact @tol 0.
Co-Authored-By: 's avatarClaude Fable 5 <noreply@anthropic.com>
parent d9b7afdb
Hide whitespace changes
Inline Side-by-side
Showing with 99 additions and 31 deletions
jna/tp_jna.cpp
View file @ 5f473261
...@@ -658,6 +658,26 @@ struct TpProc { ...@@ -658,6 +658,26 @@ struct TpProc {
float *gpu_pose_candidate_vectors, *gpu_pose_candidate_fx, *gpu_pose_candidate_jt; float *gpu_pose_candidate_vectors, *gpu_pose_candidate_fx, *gpu_pose_candidate_jt;
unsigned char *gpu_pose_selection, *gpu_pose_valid, *gpu_pose_candidate_valid; unsigned char *gpu_pose_selection, *gpu_pose_valid, *gpu_pose_candidate_valid;
int pose_raw_num_tiles, pose_raw_has_jacobian; int pose_raw_num_tiles, pose_raw_has_jacobian;
// Rung B (design 2026-07-16): the gpu_pose_* / gpu_pose_candidate_* pairs are
// FIXED physical slots (set 0 / set 1); which set is CURRENT is selected by a
// DEVICE-authoritative index flipped by pose_lma_reduce_rms on accept (no host
// std::swap - a future DP parent resolves buffers with no host round trip).
// pose_current_set is the host mirror, synced by a 4-byte D2H after each
// resident step; the cur_*/cand_* accessors below select launch pointers.
// The candidate-base DtoD copy (ref rows [0..8] of the pose vectors) is what
// keeps cand vectors valid after any number of flips: ref rows are pose-
// independent and are re-copied from the CURRENT set every step.
// By Claude on 07/16/2026.
int pose_current_set;
int *gpu_pose_current_set;
float* cur_pose_vectors() { return pose_current_set ? gpu_pose_candidate_vectors : gpu_pose_vectors; }
float* cur_pose_fx() { return pose_current_set ? gpu_pose_candidate_fx : gpu_pose_fx; }
float* cur_pose_jt() { return pose_current_set ? gpu_pose_candidate_jt : gpu_pose_jt; }
unsigned char* cur_pose_valid() { return pose_current_set ? gpu_pose_candidate_valid : gpu_pose_valid; }
float* cand_pose_vectors() { return pose_current_set ? gpu_pose_vectors : gpu_pose_candidate_vectors; }
float* cand_pose_fx() { return pose_current_set ? gpu_pose_fx : gpu_pose_candidate_fx; }
float* cand_pose_jt() { return pose_current_set ? gpu_pose_jt : gpu_pose_candidate_jt; }
unsigned char* cand_pose_valid() { return pose_current_set ? gpu_pose_valid : gpu_pose_candidate_valid; }
// One-thread float runLma candidate validation rung. Host-prepared J/W/residual // One-thread float runLma candidate validation rung. Host-prepared J/W/residual
// arrays intentionally cross H2D and the compact result crosses D2H. // arrays intentionally cross H2D and the compact result crosses D2H.
int pose_lma_value_capacity; int pose_lma_value_capacity;
...@@ -710,6 +730,7 @@ TpProc* tp_proc_create(TpModule* m){ ...@@ -710,6 +730,7 @@ TpProc* tp_proc_create(TpModule* m){
p->gpu_pose_candidate_vectors=p->gpu_pose_candidate_fx=p->gpu_pose_candidate_jt=nullptr; p->gpu_pose_candidate_vectors=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; // By Codex on 07/15/2026 p->gpu_pose_selection=p->gpu_pose_valid=p->gpu_pose_candidate_valid=nullptr; // By Codex on 07/15/2026
p->pose_raw_num_tiles=0; p->pose_raw_has_jacobian=0; p->pose_raw_num_tiles=0; p->pose_raw_has_jacobian=0;
p->pose_current_set=0; p->gpu_pose_current_set=nullptr; // By Claude on 07/16/2026
p->pose_lma_value_capacity=0; p->pose_lma_value_capacity=0;
p->gpu_pose_lma_weights=p->gpu_pose_lma_jt=p->gpu_pose_lma_ymfx=nullptr; 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_y=p->gpu_pose_lma_vector=p->gpu_pose_lma_result=nullptr;
...@@ -1261,6 +1282,15 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p, ...@@ -1261,6 +1282,15 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p,
POSE_ALLOC(gpu_pose_vectors,4*POSE_NUM_PARAMS,float); POSE_ALLOC(gpu_pose_vectors,4*POSE_NUM_PARAMS,float);
POSE_ALLOC(gpu_pose_candidate_vectors,4*POSE_NUM_PARAMS,float); POSE_ALLOC(gpu_pose_candidate_vectors,4*POSE_NUM_PARAMS,float);
#undef POSE_ALLOC #undef POSE_ALLOC
// Device-authoritative current/candidate set index (rung B). First allocation
// zeroes it (set 0 = gpu_pose_* buffers current); it persists across scenes -
// uploads always target the CURRENT set via the accessors. By Claude on 07/16/2026.
if(!p->gpu_pose_current_set){
if(cudaMalloc((void**)&p->gpu_pose_current_set,sizeof(int))!=cudaSuccess ||
cudaMemset(p->gpu_pose_current_set,0,sizeof(int))!=cudaSuccess){
seterr("exec_pose_fx_jacobian: alloc current_set failed"); return -3; }
p->pose_current_set=0;
}
const int rbr_lengths[2]={ref_rbr_len,scene_rbr_len}; const int rbr_lengths[2]={ref_rbr_len,scene_rbr_len};
const int ers_lengths[2]={ref_ers_len,scene_ers_len}; const int ers_lengths[2]={ref_ers_len,scene_ers_len};
...@@ -1304,24 +1334,28 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p, ...@@ -1304,24 +1334,28 @@ int tp_proc_exec_pose_fx_jacobian(TpProc* p,
cudaMemcpy(p->gpu_pose_rbr[1],scene_rbr,(size_t)scene_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[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->gpu_pose_ers[1],scene_ers,(size_t)scene_ers_len*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
cudaMemcpy(p->gpu_pose_vectors,pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess || cudaMemcpy(p->cur_pose_vectors(),pose_vectors,(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess || // CURRENT set (rung B) // By Claude on 07/16/2026
cudaMemcpy(p->gpu_pose_centers,centers,(size_t)3*num_tiles*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess || cudaMemcpy(p->gpu_pose_centers,centers,(size_t)3*num_tiles*sizeof(float),cudaMemcpyHostToDevice)!=cudaSuccess ||
cudaMemcpy(p->gpu_pose_selection,selection,(size_t)num_tiles*sizeof(unsigned char),cudaMemcpyHostToDevice)!=cudaSuccess){ cudaMemcpy(p->gpu_pose_selection,selection,(size_t)num_tiles*sizeof(unsigned char),cudaMemcpyHostToDevice)!=cudaSuccess){
seterr("exec_pose_fx_jacobian: HtoD failed"); return -7; } seterr("exec_pose_fx_jacobian: HtoD failed"); return -7; }
CUfunction f=getfun(p->mod,"pose_fx_jacobian"); CUfunction f=getfun(p->mod,"pose_fx_jacobian");
if(!f){ seterr("pose_fx_jacobian missing"); return -8; } if(!f){ seterr("pose_fx_jacobian missing"); return -8; }
float* gpu_jt=calculate_jacobian?p->gpu_pose_jt:nullptr; float* cur_vectors=p->cur_pose_vectors(); // set-selected launch pointers (rung B) // By Claude on 07/16/2026
float* cur_fx=p->cur_pose_fx();
float* cur_jt_buf=p->cur_pose_jt();
unsigned char* cur_valid=p->cur_pose_valid();
float* gpu_jt=calculate_jacobian?cur_jt_buf:nullptr;
void* args[]={ &num_tiles,&calculate_jacobian,&p->gpu_pose_cameras, void* args[]={ &num_tiles,&calculate_jacobian,&p->gpu_pose_cameras,
&p->gpu_pose_radial[0],&p->gpu_pose_rbr[0],&p->gpu_pose_ers[0], &p->gpu_pose_radial[0],&p->gpu_pose_rbr[0],&p->gpu_pose_ers[0],
&p->gpu_pose_radial[1],&p->gpu_pose_rbr[1],&p->gpu_pose_ers[1], &p->gpu_pose_radial[1],&p->gpu_pose_rbr[1],&p->gpu_pose_ers[1],
&p->gpu_pose_vectors,&p->gpu_pose_centers,&p->gpu_pose_selection, &cur_vectors,&p->gpu_pose_centers,&p->gpu_pose_selection,
&p->gpu_pose_fx,&gpu_jt,&p->gpu_pose_valid }; &cur_fx,&gpu_jt,&cur_valid };
const int threads=256; const int threads=256;
if(launch1(f,(num_tiles+threads-1)/threads,1,1,threads,1,1,args,"pose_fx_jacobian")) return -9; if(launch1(f,(num_tiles+threads-1)/threads,1,1,threads,1,1,args,"pose_fx_jacobian")) return -9;
p->pose_raw_num_tiles=num_tiles; p->pose_raw_has_jacobian=calculate_jacobian; p->pose_raw_num_tiles=num_tiles; p->pose_raw_has_jacobian=calculate_jacobian;
if(cudaMemcpy(fx,p->gpu_pose_fx,(size_t)POSE_NUM_COMPONENTS*num_tiles*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess || if(cudaMemcpy(fx,cur_fx,(size_t)POSE_NUM_COMPONENTS*num_tiles*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess ||
(calculate_jacobian && cudaMemcpy(jt,p->gpu_pose_jt,(size_t)POSE_NUM_PARAMS*POSE_NUM_COMPONENTS*num_tiles*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess) || (calculate_jacobian && cudaMemcpy(jt,cur_jt_buf,(size_t)POSE_NUM_PARAMS*POSE_NUM_COMPONENTS*num_tiles*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess) ||
cudaMemcpy(valid,p->gpu_pose_valid,(size_t)num_tiles*sizeof(unsigned char),cudaMemcpyDeviceToHost)!=cudaSuccess){ cudaMemcpy(valid,cur_valid,(size_t)num_tiles*sizeof(unsigned char),cudaMemcpyDeviceToHost)!=cudaSuccess){
seterr("exec_pose_fx_jacobian: DtoH failed"); return -10; } seterr("exec_pose_fx_jacobian: DtoH failed"); return -10; }
return 0; return 0;
} }
...@@ -1452,7 +1486,18 @@ int tp_proc_exec_pose_lma_resident_step(TpProc* p, ...@@ -1452,7 +1486,18 @@ int tp_proc_exec_pose_lma_resident_step(TpProc* p,
CUfunction reduce_rms=getfun(p->mod,"pose_lma_reduce_rms"); CUfunction reduce_rms=getfun(p->mod,"pose_lma_reduce_rms");
CUfunction pose=getfun(p->mod,"pose_fx_jacobian"); CUfunction pose=getfun(p->mod,"pose_fx_jacobian");
if(!prepare||!reduce||!reduce_rms||!pose){ seterr("resident pose LMA kernels missing"); return -7; } if(!prepare||!reduce||!reduce_rms||!pose){ seterr("resident pose LMA kernels missing"); return -7; }
void* pa[]={ &num_tiles,&p->gpu_pose_fx,&p->gpu_pose_jt,&p->gpu_pose_valid, if(!p->gpu_pose_current_set){ seterr("exec_pose_lma_resident_step: current_set not allocated"); return -7; }
// Set-selected launch pointers (rung B): CURRENT/candidate roles are chosen by
// the device-authoritative index, buffers stay in fixed slots. By Claude on 07/16/2026.
float* cur_vectors=p->cur_pose_vectors();
float* cur_fx=p->cur_pose_fx();
float* cur_jt=p->cur_pose_jt();
unsigned char* cur_valid=p->cur_pose_valid();
float* cand_vectors=p->cand_pose_vectors();
float* cand_fx=p->cand_pose_fx();
float* cand_jt=p->cand_pose_jt();
unsigned char* cand_valid=p->cand_pose_valid();
void* pa[]={ &num_tiles,&cur_fx,&cur_jt,&cur_valid,
&p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_eigen, &p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,&p->gpu_pose_lma_eigen,
&p->gpu_pose_lma_vector,&p->gpu_pose_lma_jt,&p->gpu_pose_lma_ymfx, &p->gpu_pose_lma_vector,&p->gpu_pose_lma_jt,&p->gpu_pose_lma_ymfx,
&p->gpu_pose_lma_partials }; &p->gpu_pose_lma_partials };
...@@ -1462,45 +1507,45 @@ int tp_proc_exec_pose_lma_resident_step(TpProc* p, ...@@ -1462,45 +1507,45 @@ int tp_proc_exec_pose_lma_resident_step(TpProc* p,
&p->gpu_pose_lma_vector,&p->gpu_pose_lma_result }; &p->gpu_pose_lma_vector,&p->gpu_pose_lma_result };
if(launch1(reduce,1,1,1,1,1,1,ra,"pose_lma_reduce_step")) return -9; if(launch1(reduce,1,1,1,1,1,1,ra,"pose_lma_reduce_step")) return -9;
// Candidate projection/RMS oracle uses separate buffers, leaving the current // Candidate projection into the candidate set. The base copy from the CURRENT
// raw fx/J state untouched for the existing Java acceptance/fallback path. // set is what keeps the candidate vectors valid after any number of accept
if(cudaMemcpyAsync(p->gpu_pose_candidate_vectors,p->gpu_pose_vectors, // flips: ref rows [0..8] are pose-independent, scene rows [9..11] are then
// overwritten with the solved candidate. By Claude on 07/16/2026 (was implicit).
if(cudaMemcpyAsync(cand_vectors,cur_vectors,
(size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess || (size_t)4*POSE_NUM_PARAMS*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess ||
cudaMemcpyAsync(p->gpu_pose_candidate_vectors+3*POSE_NUM_PARAMS, cudaMemcpyAsync(cand_vectors+3*POSE_NUM_PARAMS,
p->gpu_pose_lma_result+16, p->gpu_pose_lma_result+16,
POSE_NUM_PARAMS*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess){ POSE_NUM_PARAMS*sizeof(float),cudaMemcpyDeviceToDevice)!=cudaSuccess){
seterr("exec_pose_lma_resident_step: candidate vector DtoD failed"); return -10; } seterr("exec_pose_lma_resident_step: candidate vector DtoD failed"); return -10; }
int calculate_jacobian=1; int calculate_jacobian=1;
float* candidate_vector=p->gpu_pose_candidate_vectors+3*POSE_NUM_PARAMS; float* candidate_vector=cand_vectors+3*POSE_NUM_PARAMS;
void* ca[]={ &num_tiles,&calculate_jacobian,&p->gpu_pose_cameras, void* ca[]={ &num_tiles,&calculate_jacobian,&p->gpu_pose_cameras,
&p->gpu_pose_radial[0],&p->gpu_pose_rbr[0],&p->gpu_pose_ers[0], &p->gpu_pose_radial[0],&p->gpu_pose_rbr[0],&p->gpu_pose_ers[0],
&p->gpu_pose_radial[1],&p->gpu_pose_rbr[1],&p->gpu_pose_ers[1], &p->gpu_pose_radial[1],&p->gpu_pose_rbr[1],&p->gpu_pose_ers[1],
&p->gpu_pose_candidate_vectors,&p->gpu_pose_centers,&p->gpu_pose_selection, &cand_vectors,&p->gpu_pose_centers,&p->gpu_pose_selection,
&p->gpu_pose_candidate_fx,&p->gpu_pose_candidate_jt,&p->gpu_pose_candidate_valid }; &cand_fx,&cand_jt,&cand_valid };
const int pose_threads=256; const int pose_threads=256;
if(launch1(pose,(num_tiles+pose_threads-1)/pose_threads,1,1,pose_threads,1,1,ca, if(launch1(pose,(num_tiles+pose_threads-1)/pose_threads,1,1,pose_threads,1,1,ca,
"pose_fx_jacobian(candidate)")) return -11; "pose_fx_jacobian(candidate)")) return -11;
void* cpa[]={ &num_tiles,&p->gpu_pose_candidate_fx,&p->gpu_pose_candidate_jt, void* cpa[]={ &num_tiles,&cand_fx,&cand_jt,
&p->gpu_pose_candidate_valid,&p->gpu_pose_lma_y,&p->gpu_pose_lma_weights, &cand_valid,&p->gpu_pose_lma_y,&p->gpu_pose_lma_weights,
&p->gpu_pose_lma_eigen,&candidate_vector,&p->gpu_pose_lma_candidate_jt, &p->gpu_pose_lma_eigen,&candidate_vector,&p->gpu_pose_lma_candidate_jt,
&p->gpu_pose_lma_candidate_ymfx,&p->gpu_pose_lma_candidate_partials }; &p->gpu_pose_lma_candidate_ymfx,&p->gpu_pose_lma_candidate_partials };
if(launch1(prepare,num_partials,1,1,POSE_LMA_REDUCE_THREADS,1,1,cpa, if(launch1(prepare,num_partials,1,1,POSE_LMA_REDUCE_THREADS,1,1,cpa,
"pose_lma_prepare_products(candidate)")) return -12; "pose_lma_prepare_products(candidate)")) return -12;
void* cr[]={ &num_partials,&pure_weight,&rms_diff,&p->gpu_pose_lma_candidate_partials, void* cr[]={ &num_partials,&pure_weight,&rms_diff,&p->gpu_pose_lma_candidate_partials,
&p->gpu_pose_lma_result }; &p->gpu_pose_lma_result,&p->gpu_pose_current_set };
if(launch1(reduce_rms,1,1,1,1,1,1,cr,"pose_lma_reduce_rms")) return -13; if(launch1(reduce_rms,1,1,1,1,1,1,cr,"pose_lma_reduce_rms")) return -13;
if(cudaMemcpy(result,p->gpu_pose_lma_result, if(cudaMemcpy(result,p->gpu_pose_lma_result,
POSE_LMA_RESIDENT_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess){ POSE_LMA_RESIDENT_RESULT*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess){
seterr("exec_pose_lma_resident_step: result DtoH failed"); return -14; } seterr("exec_pose_lma_resident_step: result DtoH failed"); return -14; }
// The CUDA comparison is the state-transition oracle for this rung. Swapping // Rung B: the accept commit happens DEVICE-SIDE - pose_lma_reduce_rms flipped
// pointers makes the accepted candidate the next current raw pose without a // the current/candidate set index on accept (rejection leaves it unchanged, so
// tile-array copy; rejection leaves the current/fallback buffers untouched. // the prior current raw pose remains the fallback). Sync the host mirror from
if(result[POSE_LMA_ACCEPTED]!=0.0f){ // the device-authoritative flag; no host pointer swap. By Claude on 07/16/2026.
std::swap(p->gpu_pose_vectors,p->gpu_pose_candidate_vectors); if(cudaMemcpy(&p->pose_current_set,p->gpu_pose_current_set,sizeof(int),
std::swap(p->gpu_pose_fx,p->gpu_pose_candidate_fx); cudaMemcpyDeviceToHost)!=cudaSuccess){
std::swap(p->gpu_pose_jt,p->gpu_pose_candidate_jt); seterr("exec_pose_lma_resident_step: current_set DtoH failed"); return -14; }
std::swap(p->gpu_pose_valid,p->gpu_pose_candidate_valid);
}
if(capture_prepared && if(capture_prepared &&
(cudaMemcpy(prepared_jt,p->gpu_pose_lma_jt, (cudaMemcpy(prepared_jt,p->gpu_pose_lma_jt,
(size_t)POSE_NUM_PARAMS*num_values*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess || (size_t)POSE_NUM_PARAMS*num_values*sizeof(float),cudaMemcpyDeviceToHost)!=cudaSuccess ||
...@@ -1766,6 +1811,7 @@ void tp_proc_destroy(TpProc* p){ ...@@ -1766,6 +1811,7 @@ void tp_proc_destroy(TpProc* p){
cudaFree(p->gpu_pose_fx); cudaFree(p->gpu_pose_jt); cudaFree(p->gpu_pose_valid); // By Codex on 07/15/2026 cudaFree(p->gpu_pose_fx); cudaFree(p->gpu_pose_jt); cudaFree(p->gpu_pose_valid); // By Codex on 07/15/2026
cudaFree(p->gpu_pose_candidate_vectors); cudaFree(p->gpu_pose_candidate_fx); cudaFree(p->gpu_pose_candidate_vectors); cudaFree(p->gpu_pose_candidate_fx);
cudaFree(p->gpu_pose_candidate_jt); cudaFree(p->gpu_pose_candidate_valid); cudaFree(p->gpu_pose_candidate_jt); cudaFree(p->gpu_pose_candidate_valid);
cudaFree(p->gpu_pose_current_set); // By Claude on 07/16/2026
cudaFree(p->gpu_pose_lma_weights); cudaFree(p->gpu_pose_lma_jt); cudaFree(p->gpu_pose_lma_ymfx); 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_y); cudaFree(p->gpu_pose_lma_vector); cudaFree(p->gpu_pose_lma_result);
cudaFree(p->gpu_pose_lma_eigen); cudaFree(p->gpu_pose_lma_partials); cudaFree(p->gpu_pose_lma_eigen); cudaFree(p->gpu_pose_lma_partials);
......
src/tests/test_pose_lma_step.cu
View file @ 5f473261
...@@ -253,9 +253,17 @@ int main(int argc, char ** argv) ...@@ -253,9 +253,17 @@ int main(int argc, char ** argv)
checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaGetLastError());
pose_lma_reduce_step<<<1, 1>>>(numPartials, lambda, pureWeight, pose_lma_reduce_step<<<1, 1>>>(numPartials, lambda, pureWeight,
gpuPartials.dev(), gpuVectorResident.dev(), gpuResidentResult.dev()); gpuPartials.dev(), gpuVectorResident.dev(), gpuResidentResult.dev());
// current_set: device set index (rung B) - this case reduces the SAME partials
// as the current step, so candidate RMS == current RMS -> NOT accepted -> the
// index must stay 0 (the no-flip side of the device-side commit).
// By Claude on 07/16/2026.
GpuBuf<int> gpuCurrentSet(1);
const int zeroSet = 0;
gpuCurrentSet.upload(&zeroSet);
pose_lma_reduce_rms<<<1, 1>>>(numPartials, pureWeight, 0.001f, pose_lma_reduce_rms<<<1, 1>>>(numPartials, pureWeight, 0.001f,
gpuPartials.dev(), gpuResidentResult.dev()); gpuPartials.dev(), gpuResidentResult.dev(), gpuCurrentSet.dev());
checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize()); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize());
const std::vector<int> gotCurrentSet = gpuCurrentSet.download();
const std::vector<float> gotJt = gpuPreparedJt.download(); const std::vector<float> gotJt = gpuPreparedJt.download();
const std::vector<float> gotYmfx = gpuPreparedYmfx.download(); const std::vector<float> gotYmfx = gpuPreparedYmfx.download();
const std::vector<float> gotResident = gpuResidentResult.download(); const std::vector<float> gotResident = gpuResidentResult.download();
...@@ -279,7 +287,8 @@ int main(int argc, char ** argv) ...@@ -279,7 +287,8 @@ int main(int argc, char ** argv)
(maxProducts <= 0.02f) && (maxDelta <= 2.0e-5f) && (maxProducts <= 0.02f) && (maxDelta <= 2.0e-5f) &&
(maxCandidate <= 2.0e-5f) && (maxRms <= 2.0e-6f) && (maxCandidate <= 2.0e-5f) && (maxRms <= 2.0e-6f) &&
(gotResident[POSE_LMA_ACCEPTED] == 0.0f) && (gotResident[POSE_LMA_ACCEPTED] == 0.0f) &&
(gotResident[POSE_LMA_STOP] == 0.0f); (gotResident[POSE_LMA_STOP] == 0.0f) &&
(gotCurrentSet[0] == 0); // not accepted -> no device-side flip // By Claude on 07/16/2026
std::printf("pose_lma resident 5120/150: prepared bit mismatches=%d, max products=%g, delta=%g, candidate=%g, RMS=%g -> %s\n", std::printf("pose_lma resident 5120/150: prepared bit mismatches=%d, max products=%g, delta=%g, candidate=%g, RMS=%g -> %s\n",
preparedBad, maxProducts, maxDelta, maxCandidate, maxRms, preparedBad, maxProducts, maxDelta, maxCandidate, maxRms,
residentOk ? "PASS" : "FAIL"); residentOk ? "PASS" : "FAIL");
......
src/tp_lma.cu
View file @ 5f473261
...@@ -737,7 +737,8 @@ extern "C" __global__ void pose_lma_reduce_rms( ...@@ -737,7 +737,8 @@ extern "C" __global__ void pose_lma_reduce_rms(
float pure_weight, float pure_weight,
float rms_diff, float rms_diff,
const float * partials, const float * partials,
float * result) float * result,
int * current_set) // device set index, flipped on accept (NULL = no flip) // By Claude on 07/16/2026
{ {
if ((blockIdx.x != 0) || (threadIdx.x != 0)) return; if ((blockIdx.x != 0) || (threadIdx.x != 0)) return;
float pure_l2 = 0.0f; float pure_l2 = 0.0f;
...@@ -760,6 +761,12 @@ extern "C" __global__ void pose_lma_reduce_rms( ...@@ -760,6 +761,12 @@ extern "C" __global__ void pose_lma_reduce_rms(
pose_lma_sub(1.0f, rms_diff))); pose_lma_sub(1.0f, rms_diff)));
result[POSE_LMA_ACCEPTED] = accepted ? 1.0f : 0.0f; result[POSE_LMA_ACCEPTED] = accepted ? 1.0f : 0.0f;
result[POSE_LMA_STOP] = stop ? 1.0f : 0.0f; result[POSE_LMA_STOP] = stop ? 1.0f : 0.0f;
// Device-authoritative accept commit (rung B): flip the current/candidate
// buffer-set index in device memory. The host (and later a DP parent) reads
// this index instead of swapping pointers. By Claude on 07/16/2026.
if (accepted && (current_set != NULL)) {
*current_set ^= 1;
}
} }
extern "C" __global__ void lma_normal_products( extern "C" __global__ void lma_normal_products(
......
src/tp_lma.h
View file @ 5f473261
...@@ -108,12 +108,18 @@ extern "C" __global__ void pose_lma_reduce_step( ...@@ -108,12 +108,18 @@ extern "C" __global__ void pose_lma_reduce_step(
// Reduce a separately projected candidate pose to RMS only. Results append to // Reduce a separately projected candidate pose to RMS only. Results append to
// the compact resident step record without changing its candidate/normal terms. // the compact resident step record without changing its candidate/normal terms.
// On accept it also flips *current_set (0<->1) DEVICE-SIDE - the authoritative
// current/candidate buffer-set index (rung B, design 2026-07-16): a future DP
// parent resolves the raw pose buffers with no host round trip; today the host
// mirrors the flag with a 4-byte D2H after the step. NULL current_set = no flip
// (standalone kernel tests). By Claude on 07/16/2026.
extern "C" __global__ void pose_lma_reduce_rms( extern "C" __global__ void pose_lma_reduce_rms(
int num_partials, int num_partials,
float pure_weight, float pure_weight,
float rms_diff, float rms_diff,
const float * partials, const float * partials,
float * result); // candidate RMS fields in resident result float * result, // candidate RMS fields in resident result
int * current_set); // device current/candidate set index, flipped on accept
extern "C" __global__ void lma_normal_products( extern "C" __global__ void lma_normal_products(
int num_params, int num_params,
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment