Commit 1d7235c0 authored by Andrey Filippov's avatar Andrey Filippov

CLAUDE: test_pose_corr_jna - tol-0 gpu_peaks tier (production float capture vs NVRTC float replay)

expected_gpu_peaks_it<k> (exported once the production hookup runs, imagej
f8df8dce) rows align with expected_corr_indices; the float-kernel replay is
compared bit-exactly at --tol, matched by packed index. Absent buffer =
skipped (pre-hookup cases). Current case: ALL PASS unchanged.
Co-Authored-By: 's avatarClaude Fable 5 <noreply@anthropic.com>
parent 4bcc5195
......@@ -351,6 +351,30 @@ int main(int argc, char** argv){
k, n_pd, pk_maxd, pk_maxulp, pk_viol, pk_nan, pk_fail?"FAIL":"PASS", peak_tol, PEAK_TOL_ULP);
printf("it%d: float-vs-double peaks (informational, step 2): max |d dx,dy|=%.3g px, str rel=%.3g, |d eigvec|=%.3g, lambda rel=%.3g, NaN-mismatch=%ld\n",
k, f_dxy, f_str, f_eigv, f_lam, f_nan);
// TOL-0 peak tier (cases exported with the production float kernel active,
// 07/14/2026+): expected_gpu_peaks rows align with expected_corr_indices;
// the NVRTC float replay must reproduce them bit-exactly (rung-1 rule).
snprintf(nm,sizeof(nm),"expected_gpu_peaks_it%d",k);
if (data.has(nm)){
const float* egp = data.hostFloat(nm);
double gp_maxd = 0; long gp_nan = 0, gp_missing = 0;
for (int i = 0; i < n_exp; i++){
std::map<int,int>::const_iterator it = by_index.find(exp_idx[i]);
if (it == by_index.end()){ gp_missing++; continue; }
const float* g = got_pkf.data() + (size_t)it->second*PEAK_ROW_FLOATS;
const float* e = egp + (size_t)i*PEAK_ROW_FLOATS;
for (int j = 0; j < PEAK_ROW_FLOATS; j++){
const bool gn = std::isnan(g[j]), en = std::isnan(e[j]);
if (gn || en){ if(gn!=en) gp_nan++; continue; }
const double d = std::fabs((double)g[j]-(double)e[j]);
if (d > gp_maxd) gp_maxd = d;
}
}
const int gp_fail = (gp_maxd>tol)||(gp_nan>0)||(gp_missing>0);
printf("it%d: compare 'gpu_peaks' (float kernel vs production capture): max|diff|=%g NaN-mismatch=%ld missing=%ld -> %s\n",
k, gp_maxd, gp_nan, gp_missing, gp_fail?"FAIL":"PASS");
if (gp_fail) fail = 1;
}
if (pd_fail || pk_fail) fail = 1;
cudaMalloc((void**)&scratch, res_pd.size()*sizeof(float));
cudaMemcpy(scratch, res_pd.data(), res_pd.size()*sizeof(float), cudaMemcpyHostToDevice);
......
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