- 14 Jul, 2026 4 commits
-
-
Andrey Filippov authoredThe scene-sequence .list is the single source of truth: it now carries SET casesFile lists/cases.list (relative to rootDirectory, the EyesisCorrectionParameters rule). run_cases.sh discovery order: -c arg -> the scene list's SET casesFile (when -x/-l provides the list; explicit pointer, missing target = hard error) -> convention lookup under rootDirectory -> repo-local gitignored link. Verified: -l, -x and bare modes all resolve to the same manifest/model. Java companion (imagej-elphel): casesFile added to KEY_DIRS (index 12, file-exempt from the auto-mkdirs) so the SET parses cleanly. Co-Authored-By:Claude Fable 5 <noreply@anthropic.com>
-
Andrey Filippov authored
Andrey's rulings 07/13/2026: data directories are not git-controlled; the repo is shared across computers with different data sets, so it carries code, generators and a heavily-commented TEMPLATE (linux-config style) - everything data-specific lives under the data root: - cases.list -> cases.list.tmpl (tracked; install by copying to <dataRoot>/lists/cases.list next to the scene lists; SET listFile may then be a bare filename - resolved relative to the manifest's directory). - /cases.list gitignored: optional repo-local symlink into the data root so a bare ./run_cases.sh works. - run_cases.sh: manifest lookup order -c arg -> ./cases.list link -> via -x/-l through the scene list's rootDirectory (<root>/lists/, <root>/, next to the scene list). -x corr-xml alone now finds everything (list, version, manifest) from the entries the Java run uses. Verified live: installed at /home/elphel/lwir16-proc/LV/lists/cases.list; bare run ALL PASS via the symlink; -x corr-xml discovery ALL resolved with the symlink removed. Co-Authored-By: -
Andrey Filippov authored
CLAUDE: manifest-like tests top - cases.list + run_cases.sh (Andrey's design), de-hardwire test defaults The scene-sequence .list is the human single source of truth (SET header + body lines selected by uncommenting; commented lines document the setup). The tests now follow the SAME convention one level up: - cases.list: SET header (listFile, version - the version deliberately kept out of the scene list, mirroring CORRECTION_PARAMETERS.x3dModelVersion vs .sourceSequencesList) + case lines selected by uncommenting; commented cases stay as documentation. gen lines map synthetic cases to generators. - run_cases.sh: resolves ${MODEL} through the scene list's SET rootDirectory/x3dDirectory + its UNCOMMENTED sequence lines + the version (the *-CENTER dir carries the DERIVED center timestamp -> matched as the smallest center ts >= the sequence's first-scene ts containing <version>/); -x <corr-xml> reads sourceSequencesList/x3dModelVersion directly = the same two sources that drive the Java run. No args = all uncommented cases; named cases run even when commented; missing synthetic data triggers the generator. - test_proc_consolidate: hardwired absolute case/repo paths -> standard --data/--src/--devrt interface (legacy positional kept). - test_pose_corr_jna: srcdir default de-hardwired (repo-relative 'src'). Verified live: ./run_cases.sh ALL PASS (pose_corr real-scene tol-0 resolved through the list+version, 4 synthetic cases); named-commented case avg_td_oob_jna PASS bit-exact via the new interface; gen-on-missing works. Co-Authored-By: -
Andrey Filippov authored
CLAUDE: 3-C rung 2 - corr2D_peak_eig (frozen 15-arg getMaxXYCmEig contract) + checkpoint-2 replay + synthetic numpy gate Andrey's ruling 07/13/2026: the GPU peak contract = the 15-arg superset - closed-form 2x2 eigen, recentering as a runtime param, envelope de-bias as an optional array pointer (NULL = off). - src/tp_peak.{h,cu}: corr2D_peak_eig, one block/tile - parallel load/NaN/ argmax(first-occurrence)/8-connected grow, thread-0 serial DOUBLE centroid/ covariance/eigen/recentering in the Java oracle's exact summation order. - jna/tp_jna.cpp: kernel in the NVRTC module (24 kernels) + tp_proc_exec_corr2d_peak / tp_proc_get_peaks / tp_proc_set_peak_debias. - test_pose_corr + test_pose_corr_jna: checkpoint-2 replay (in-place corr2D_normalize -> corr_pd tol tier, peaks vs the double Java oracle at --peak-tol) when the case carries it; pre-rung-2 cases skip with a note. - make_peak_case.py + test_peak: synthetic gate vs an INDEPENDENT numpy float64 port (8 adversarial tiles + refine/debias case) - PASS at max|diff| 0 / 9.6e-17. - Regressions: test_pose_corr_jna bit-exact @tol 0 on the existing v013 case, avg_td / avg_td_oob / convert_direct PASS. Co-Authored-By:
-
- 13 Jul, 2026 2 commits
-
-
Andrey Filippov authored
CLAUDE: test_pose_corr_jna - the tol-0 verdict binary (NVRTC production module); direct test = stepping companion RUNG 1 VALIDATED on the first real v013 case (scene 1773135467_917220, 4 LMA iterations, 150 tiles/set, 300 corr tiles/iter): test_pose_corr_jna PASS BIT-EXACT @tol 0 - tasks (calculate_tiles_offsets) and raw TD correlation (correlate2D_inter) both max|diff|=0, every iteration. The offline-nvcc direct-launch binary (src/tests/test_pose_corr) cannot be bit-exact vs the NVRTC JIT (FMA/codegen divergence), measured: -G -O0 -> tasks off by 2^-13 px; RELEASE=1 -> tasks BIT-EXACT, corr within ~1 ULP at operand scale (max|diff| 128..256 vs amplitudes ~2e9, float spacing 256). Role split documented in both headers (same pairing as test_avg_td_oob / test_proc_consolidate): direct binary for cuda-gdb stepping, _jna binary for the production-equivalence verdict. Build line in the _jna header. Co-authored-by: -
Andrey Filippov authored
Standalone per-iteration replay of the CuasPoseRT.leanMeasure GPU chain against the PoseCorrExport case (imagej-elphel curt.kernel_test=pose_corr): calcReverseDistortionTable + calc_rot_deriv + calculate_tiles_offsets (NEW under test; GPU-filled tasks compared bit-exact vs the oracle post-offsets streams, header+xy fields - the disp_dist tail is not round-tripped by the Java flatten), erase_clt_tiles(NaN) + convert_direct for both MB sets, the v1 consolidation chain into the cam-0 slot (inputs = the oracle post-offsets streams, the exact bytes production flattens), then correlate2D_inter vs the resident center TD from gpu_clt_ref cam 0 (NEW under test). Raw-TD compare is ORDER-INDEPENDENT (corr slot placement is atomicAdd-nondeterministic): tiles matched by packed index, tol 0 goal; results/corr_td_it<k> saved in oracle order for the Java round trip. All four LPF/HPF __constant__ tables overwritten from the case via extern __constant__ + cudaMemcpyToSymbol (correlate2D_inter reads LoG_corr and lpf_rb_corr). Built by build_tests.sh (auto-discovered), cuda-gdb steppable. Co-authored-by:
-
- 12 Jul, 2026 13 commits
-
-
Andrey Filippov authored
Wire the frozen v1 chain (index_consolidate -> consolidate_oob -> clt_average_sensors_list) into the production JNA surface (roadmap 2c -> step 4 JNA hookup, Andrey's ruling: JNA-only, no JCuda launch code): - tp_jna.cpp: tp_consolidate.h/.cu added to the NVRTC module (23 kernels); new export tp_proc_exec_consolidate = host orchestration of the chain on the RESIDENT per-cam CLT (contiguous DtoD gather, no D2H/H2D of TD data; only the flattened task streams cross the boundary); average lands in the cam-0 slot (what the single conj-multiply correlates); lazy scratch, freed in tp_proc_destroy; stats {pairs, surviving tiles, misaligned}. - tp_consolidate.h/.cu: JCUDA include guards (NVRTC concatenation has no include paths) + NAN/INFINITY fallbacks (NVRTC has no math.h macros). - jna/test_proc_consolidate.cu: standalone validation of the export vs the real-scene avg_td_oob oracle case - PASS bit-exact (max|diff|=0, 5120 pairs -> 5056 tiles == oracle, 0 misaligned). Regressions: tests_bin/test_avg_td, test_avg_td_oob PASS at tol 0 after the guard changes; module smoke = 23/23 kernels resolved. Co-authored-by: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
Three kernels (tp_consolidate.cu, oracle = imagej-elphel CuasTD @eab1660e): - index_consolidate: index_direct successor aware of the motion-blur task pair (gpu_ftasks1 may be NULL = no-MB, mirrors Java TpTask[2][] with [1] null); keeps tiles active in BOTH sets, txy-alignment verified (misaligned counted, skipped); output = compacted even/odd PAIRS (no-MB: odd = copy) so pass 2 is uniform. Activity tested on the task word as int (v0 index_direct float-compare -0.0 quirk noted). - consolidate_oob: frozen 2c spec - per sensor depth = min over the MB pair of tile-center distance to nearest frame edge; <hard(8) always skipped, any >=soft(12) only those, else hard survivors last resort, nobody -> tile dropped; NaN xy excluded (explicit isnan - fminf would silently ignore NaN); mask folded into task bits TASK_AVG_SENS_SHIFT=11 (bits 11+ virgin per the TpTask bit-field survey). - clt_average_sensors_list: v0 core driven by the compacted list, gated by the mask; host prefills NaN/0 = oracle poison semantics. v0 clt_average_sensors refactored onto the shared consolidate_chunk core (all-ones mask), verified bit-exact. test_avg_td_oob: loads AvgTdExport.exportOob cases (raw ftasks0/1 streams); runs the chain; compares masks exactly + td_avg/counts at tol 0. make_avg_td_oob_case.py: synthetic 4-sensor case = the SAME scenario the Java oracle self-test passes (every OOB branch); independent float32 oracle-op-order expectations. VERIFIED: synth avg_td (v0 refactor) max|diff|=0; synth avg_td_oob chain 5 pairs -> 4 tiles, masks exact, avg/counts max|diff|=0 at tol 0; test_convert_direct regression PASS both tiers. Co-Authored-By: -
Andrey Filippov authored
cuda_runtime.h only pulls device_launch_parameters.h under __CUDACC__, which the indexer stopped defining when CDT 2025-09 dropped the Nsight scanner-discovery provider -> 14 false markers (blockIdx/threadIdx/ __syncthreads/NAN/isnan unresolved; F3 still jumped via other TUs). Fix: include cuda_runtime.h + device_launch_parameters.h + math.h explicitly, plus a __CDT_PARSER__-only __syncthreads() decl (its real declaration is __CUDACC__-gated). No-op for nvcc: standalone -dc compile clean, build_tests OK, test_convert_direct PASS both tiers (strict 0). Co-Authored-By: -
Andrey Filippov authored
- launches/test_convert_direct.launch.tmpl: harvested working shared config (C/C++ Application + cuda-gdb 12.8; .tmpl suffix so Eclipse does not discover the template copy as a duplicate; setup.sh strips it on install) - setup.sh: install launches/ - nsight_per_kernel_tests.md: modern Eclipse does not load the Nsight local 'CUDA GDB Application' type - use C/C++ Application with cuda-gdb as GDB; shared-config save recipe; DP child-launch stepping caveat - README: document launches/ templating Co-Authored-By: -
Andrey Filippov authored
CLAUDE: gitignore live /launches (shared Eclipse debug configs; template copy will live in eclipse_setup/) Co-Authored-By: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
-
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
CLAUDE: eclipse_setup + build_tests.sh: ${HOME}-relative cuda-samples path (no absolute /home/elphel) Co-Authored-By: -
Andrey Filippov authored
- .cproject: exclude src/tests|jna|attic from the managed build (tests are many main() executables built by build_tests.sh, not the one-binary managed project); add ${workspace_loc}/src include path (indexer resolves test #includes); sync Release samples include path with live drift - .project: second builder = .externalToolBuilders/build_tests.launch (runs src/tests/build_tests.sh on manual Build; machine-generic via workspace_loc variables) - .settings: adopt live doxygen + UTF-8 encoding prefs - .gitignore: ignore live /.externalToolBuilders (template copies up, then disconnects, per convention) - README.md: the convention + what is workspace-level (CUDA toolkit pref, 12.8 cuda-gdb debug configs) Co-Authored-By:
-
- 10 Jul, 2026 3 commits
-
-
Andrey Filippov authored
- make_convert_direct_case.py packages the legacy clt/ exports + goldens into a manifest case (one-time bridge; new cases come from Java oracle) - test_convert_direct.cu: manifest -> RAII buffers (GpuBuf/GpuPitchedBuf + device pointer arrays) -> host-built tasks from portsxy -> single deterministic launch; the historical (0, kernels_hor) deconvolution-gate quirk is an explicit apply_kernels manifest parameter - two-tier goldens: expected_clt_2025 legacy @ legacy_tol=1.2 (measured 1.107 max, 2.2e-5 relative rms - same class JNA Stage2 accepted <1.0); expected_clt blessed on this machine @ tol=0 for real regression - TpTestData.compareTo(named reference); avg_td re-verified PASS tol=0 Co-Authored-By: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
- src/tests/: TpTestData (manifest.txt contract with the Java oracle, testdata_format.md), test_avg_td main, build_tests.sh (CUDA 12.8, -arch=native, device-debug for NSight on Blackwell sm_120) - src/tp_consolidate.cu: clt_average_sensors v0, bit-exact mirror of CuasTD.consolidateSensorsTD (NaN-absent, count plane, n>1 whole-tile poison, reciprocal-multiply op order); OOB + median/MAD stages next - verified: synthetic + Java-oracle cases PASS at tol=0 on RTX 5060 Ti - pre-existing edits in jna/tp_jna.cpp, src/TileProcessor.cu left alone Co-Authored-By:
-
- 03 Jul, 2026 1 commit
-
-
Andrey Filippov authored
CLAUDE: self-documenting comments: MCLT overview, quad algebra conj/no-conj, tp_task bridge, GPU geometry chain Comment-only. From Andrey's explanation: MCLT perfect-reconstruction overview (TileProcessor.h, with blog ref); convolveTiles = quad product WITHOUT conjugation, kernels DTT-II/III family vs type-IV image tiles; correlateAccumulateTiles = conj(T1) x T2, order-sensitive; shiftTileHor/Vert = lossless fractional-shift phase rotators (differential rectification); struct tp_task bridge role + composed-offset semantics; calc_rot_deriv / calculate_tiles_offsets = GPU side of historic getPortsCoordinates(). Pre-existing dirty jna/tp_jna.cpp intentionally NOT included. Co-Authored-By:
-
- 27 Jun, 2026 3 commits
-
-
Andrey Filippov authored
tpdnn_infer processed the whole `count`-scene request in one shot; at full res (512x640) a 64-scene field tensor is ~10GB (64x124x512x640) + decode p ~10GB -> CUDA OOM on a 16GB card. The 64x80 synthetic smoke didn't expose it. Now loops in GPU_CHUNK-sized sub-batches (== infer_server's GPU_CHUNK), writing each chunk straight to the host output buffers so GPU tensors stay bounded (~3-4GB/chunk at CHUNK=8). L2 hidden/age/sprev carry across chunks; reset only at the first chunk when l2_reset (matches the server). Env TPDNN_GPU_CHUNK (default 8). Validated: parity vs server oracle still EXACT (0.0) at CHUNK=8 (8+4) and CHUNK=4 (4+4+4) - L2 carry across chunk boundaries is correct. Co-Authored-By: -
Andrey Filippov authored
Server-side copies the already-uploaded libtorch zip into mirror.elphel.com/maven-dependencies/org/pytorch/libtorch-cxx11-cu128/2.7.1/ (maven layout: .zip + .pom + .sha1/.md5, no 3.8GB re-upload) so imagej-elphel's -Plibtorch profile can dependency:unpack it. Pairs with fetch_libtorch.sh (direct download). Co-Authored-By: -
Andrey Filippov authored
Piece 2 of the native-JNA DNN path. tp_dnn.cpp is a C-ABI port of infer_server.py's hot path so the Java client can run L1+L2 in-process instead of over TCP: tpdnn_init/upload/infer/free (+ num_levels/level_frames) faithfully reproducing build_pyramid, the 16x shift-and-stitch full-res recovery, decode (ghostbuster + velocity centroid), and the L2 ConvGRU recurrence + track-age. Loads the TorchScript models from imagej_elphel_dnn (export_torchscript / export_l2_torchscript). Disables the TorchScript JIT fuser at init (nvrtc element-wise fusion fails on Blackwell; production wants no runtime nvrtc). Validated: native vs the running Python server (same CUDA) max|diff| offset5=0, roi=0 — bit-for-bit. (Oracle dump_ref.py + driver tpdnn_test.cpp, scratch.) Built standalone via build_dnn.sh (g++ + libtorch 2.7.1+cu128, ABI=1), separate from the nvcc-built libtileproc.so; fetch_libtorch.sh pulls the pinned libtorch. Context unification + zero-copy kernel<->tensor sharing is a later step. Co-Authored-By:
-
- 26 Jun, 2026 5 commits
-
-
Andrey Filippov authored
execImcltRbgAll(ref_scene) was dropped on the JNA side: native imclt was hardcoded to gpu_clt, so rendering the reference scene actually rendered the scene. Add a use_ref arg -> select gpu_clt_ref. Needed so the reference-CLT post-mortem render reflects the real buffer (CORR2D-all-NaN divergence: inter correlation needs BOTH gpu_clt + gpu_clt_ref; gpu_clt proven good via SOURCE). Updated the internal selftest caller to (p,1,0). Co-Authored-By: -
Andrey Filippov authored
Inverse of tp_proc_get_corr_indices/get_corr_td: uploads host-built per-pair TD corr indices + data back to the device (gpu_corr_indices / gpu_num_corr_tiles / gpu_corrs_td, pitched) and sets last_num_corr_tiles so the following corr2d_normalize + get_corr2d use the right count. Backs GpuQuad.setCorrIndicesTdData (TDCorrTile.convertTDtoPD) on the JNA rectilinear CUAS path, which previously fell through to base JCuda and NPE'd on a null device pointer. Buffers are sized num_pairs*ntiles (num_pairs=3 for the rectilinear config), giving ample headroom for the selected-tile count. Co-Authored-By: -
Andrey Filippov authored
Oracle (clt_aberrations_quad_corr_GPU) needs the texture path that RT did not: execTextures (calc_extra -> diff_rgb_combo, calc_textures -> textures), getTextureIndices, getExtra, getFlatTextures. Adds to TpProc: - texture buffers (lazy-allocated in tp_proc_ensure_textures; sizes match GpuQuad: max_texture_size=(num_colors+1+(num_cams+num_colors+1))*256) - tp_proc_exec_textures: textures_nonoverlap<<<1,1>>> mirroring GpuQuad.execTextures_DP. The kernel CDP-builds the index list internally (create_nonoverlap_list) so we do NOT host pre-fill it, and CDP-launches textures_accumulate (dyn-shared attr set on the accumulate fn). linescan_order is taken from the caller (0 in production -> diff_rgb_combo in texture_indices order). This matches PRODUCTION, not the Stage-5 harness convention (linescan=1 + host-prefilled indices + stale golden) that produced the documented diff_rgb_combo mismatch. - tp_proc_get_texture_indices / get_diff_rgb_combo / get_textures readback. Native compiles clean (4 new symbols exported). Not yet deployed: live .so left untouched (run was active); rebuild via jna/build_lib.sh when free. Co-Authored-By: -
Andrey Filippov authored
Native readback of gpu_corr_indices / gpu_corrs_combo_indices / gpu_corrs_td (DtoH) and cudaMemset2D erase of gpu_corrs, for GpuQuadJna's getCorrIndices/getCorrComboIndices/getCorrTdData/eraseGpuCorrs (oracle TD-correlation path). Co-Authored-By: -
Andrey Filippov authored
THE production-mismatch bug (RMSE ~1.7 vs JCUDA, invariant to FPN/row-col/MB). convert_direct gates the deconvolution kernels on `kernels_hor>0` (TileProcessor.cu:2782-2783): with kernels_hor=0 it passes NULL kernels -> NO deconvolution. tp_proc_exec_convert_direct hardcoded kh=0 (copied from the harness, whose golden was itself made with no deconvolution), so JNA skipped aberration deconvolution while production GpuQuad passes (kernels_hor, kernels_vert)=(82,66) and applies it. Fix: add kernels_vert to TpProc (= kern_tiles/(kernels_hor*num_colors)); exec passes (no_kernels?0:kernels_hor, no_kernels?0:kernels_vert). tp_proc_convert_selftest now uses no_kernels=1 to keep matching the NO-deconv harness golden (StageProc still PASS: CLT 0.1085 / RBG 0.0201 / corr 2e-5). Production (GpuQuadJna no_kernels=false) now applies deconvolution = matches JCUDA. .so-only change. Co-Authored-By:
-
- 25 Jun, 2026 9 commits
-
-
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
Co-Authored-By: -
Andrey Filippov authored
Mirrors GpuQuad.setLpfRbg/setLpfCorr (cuModuleGetGlobal + cuMemcpyHtoD) against the native module's lpf_data / lpf_corr / lpf_rb_corr symbols. Fixes the first JNA-mode NPE (setLpfRbg hit the null gpuTileProcessor.module). Co-Authored-By: -
Andrey Filippov authored
CLAUDE: Step 2 corr — split TpProc into granular correlate2D_td/_inter_td/combine/normalize + get_corr2d Match GpuQuad's separate correlation calls: tp_proc_exec_corr2d_td (correlate2D TD, fat_zero=0, corr_radius=0 -> gpu_corrs_td, returns num_corr_tiles), tp_proc_exec_corr2d_inter_td (correlate2D_inter clt vs clt_ref, sel_sensors), tp_proc_exec_corr2d_combine (init|no_transpose<<1, num_pairs, pairs_mask -> num_corr_combo), tp_proc_exec_corr2d_normalize (combo -> corrs_combo / per-pair -> gpu_corrs), tp_proc_get_corr2d (per-pair de-pitch) + tp_proc_num_corr_tiles/combo. Add per-pair pixel buffer gpu_corrs (+dstride_corr). The bundled tp_proc_exec_corr2d/StageProc path is retained. Co-Authored-By: -
Andrey Filippov authored
Add to the persistent TpProc API: tp_proc_setup_rbg_corr (imclt RBG + correlation buffers + corr config), tp_proc_exec_imclt / tp_proc_get_rbg, tp_proc_exec_corr2d (correlate2D TD -> corr2D_combine -> corr2D_normalize) / tp_proc_get_corr2d_combo. launch1() helper. tp_proc_convert_selftest extended to validate imclt vs aux_chnN.rbg and quad corr vs aux_corr-quad.corr (order-independent, stale golden). Validated on RTX 5060 Ti via the persistent API: CLT==golden (0.1085), RBG==golden (0.0201), quad-corr value-err 2.06e-5, no_kernels runs finite. The persistent granular API now covers the full set of GPU ops the LWIR16 CUAS workflow uses (geometry/convert_direct[+no_kernels/use_center_image/ erase_clt/ref_scene]/imclt/correlations) — the surface GpuQuadJna (integration step 2) delegates to. Co-Authored-By: -
Andrey Filippov authored
Add TpProc: the production-facing persistent instance (buffers allocated once in tp_proc_setup, reused across set/exec/get, freed in tp_proc_destroy) — the surface GpuQuadJna will delegate to. API: tp_proc_create/setup, set_geometry/correction_vector/kernels/kernel_offsets/image/center_image/ tasks, exec_geometry (calc_reverse_distortions+rot_derivs+calculate_tiles_offsets), exec_convert_direct (ref_scene/erase_clt/no_kernels), get_clt, destroy. Includes the fragile convert_direct paths the migration must preserve: no_kernels (skip deconvolution -> kernels_hor/vert=0), use_center_image (broadcast one center image to all sensors), erase_clt (erase_clt_tiles), ref_scene (clt_ref buffer). tp_proc_convert_selftest validates end-to-end on RTX 5060 Ti: standard convert CLT == clt/aux_chnN.clt golden (max|CLT-golden|=0.1085, == Stage 2, num_active=5120); no_kernels path runs with finite output. update_image_gpu pitch is in BYTES (the "in floats" comment is misleading). Co-Authored-By: -
Andrey Filippov authored
CLAUDE: Stage 5 — native textures_nonoverlap via JNA (executes on Blackwell; golden mismatch documented) Extend run_convert_pipeline with do_tex: setTextures/setRGBA-equivalent buffers (texture_indices from tasks, gpu_textures, diff_rgb_combo, color_weights, generate_RBGA_params), cuFuncSetAttribute(textures_accumulate, MAX_DYNAMIC_SHARED_SIZE_BYTES, shared_size) for the CDP child, launch textures_nonoverlap <<<1,1>>> (18 args), de-interleave diff_rgb_combo. tp_tex_selftest. RESULT on RTX 5060 Ti: textures_nonoverlap + its CDP child textures_accumulate EXECUTE correctly (no errors, shared 58880, 5120 tiles, output rms within ~1% of golden) => Blackwell compatibility confirmed. BUT diff_rgb_combo does NOT match the Jul-2025 golden numerically (value layers off by constant ~268, diff layers diff_sigma-sensitive). Ruled out input-CLT sensitivity (same error with golden CLT), diff_sigma (10.0 closest), arg/param order. NOT used by the LWIR16 CUAS workflow (cuas/ uses only convert_direct/corr2D_normalize/imclt_rbg_all). Documented known issue (golden staleness / unverified RGB-path drift), to track later via git bisect + the 107 kernel branch switch. See imagej-elphel-internal handoffs/2026-06-25_texture-diff-rgb-combo-mismatch.md. Co-Authored-By: -
Andrey Filippov authored
Extend run_convert_pipeline with do_corr: after convert_direct, allocate corr buffers (gpu_corrs_td/combo_td/combo via alloc_image_gpu + corr_indices/combo_indices/num_corr_tiles), launch correlate2D <<<1,1>>> (TD, CDP; 18 args, generates indices), read num_corr_tiles, corr2D_combine (quad pairs_mask 0x0f), corr2D_normalize (TD->pixel), de-pitch gpu_corrs_combo. tp_corr_selftest wrapper (do_corr=1). Validated on RTX 5060 Ti via JNA: num_pairs=120, num_corr_combo=5120, output stats identical to golden (max 0.6638, rms 0.0717). clt/aux_corr-quad.corr is OLDER (Apr-2025) than the CLT golden (Jul-2025) so the active-tile ORDER differs -> pointwise compare is permutation-dominated (0.66). Order-independent check (sort both, compare distributions): max value error 2.06e-05 == float32 precision => correlate2D/combine/normalize compute the correct values on Blackwell. Co-Authored-By: -
Andrey Filippov authored
Refactor the Stage-2 selftest into run_convert_pipeline(do_imclt) shared helper; tp_convert_direct_selftest is now a thin wrapper (do_imclt=0). Add tp_imclt_selftest (do_imclt=1): after convert_direct, allocate pitched RBG output buffers (alloc_image_gpu, 648x520/cam, mono), launch imclt_rbg_all <<<1,1>>> (gpu_clt -> gpu_corr_images), de-pitch via cudaMemcpy2D, compare to clt/aux_chnN.rbg golden. Validated on RTX 5060 Ti via Java->JNA: max|RBG-golden|=0.0201 over peaks 1535 -> relative ~1.31e-5. convert_direct CLT error unchanged (0.108505) => no regression. Co-Authored-By:
-