Commit ef36688d authored by Andrey Filippov's avatar Andrey Filippov

CLAUDE: 3-B rung B1 - pose_task_update resident task-build kernel

New NVRTC kernel pose_task_update: one thread per selected task template
entry runs the C1-validated projection chain (pose_world_from_pixel ->
pose_pixel_from_world) and rewrites the 6-float task headers of BOTH
resident ftask slots in place (slot 0 = MB main, slot 1 = MB partner).
Inputs are all resident: per-sequence template (packed txy + reference
centers, configuration-register semantics), per-scene camera blocks
(shared with prepare_resident via pose_prep_ready bits), the
device-authoritative current pose set, and per-scene scalars (uniform-MB
6-float descriptor + margins + pre-encoded |511 task words). Failed
projection / margin-gate tiles become task=0 HOLES in both sets -
index_direct, index_consolidate and the inter correlation all skip them
(the Java condensed stream reached the same end by omission).

JNA API: tp_proc_exec_pose_task_update (nullable groups, rung C2
contract; task_size derived from the uploaded skeleton) +
tp_proc_activate_tasks_slot (make a resident slot the live task stream
with no H2D - geometry/convert consume the kernel-rewritten slots).

Tests (test_avg_td_oob two-tier pairing precedent):
- src/tests/test_pose_task_update (nvcc-direct): grid-stride body makes
  a <<<1,1>>> launch the serial float oracle - parallel launch memcmp
  bit-exact; centerXY == pose_fx_jacobian fx + descriptor at 0 px
  (bit-identical across kernels); host-replayed margin/hole logic exact;
  xy/disp_dist tails untouched. PASS.
- jna/test_pose_task_update_jna (production NVRTC module): fx-consistent
  at 0 px, keep-resident rerun memcmp-identical, no-MB encoding
  (sub=511), activate-slot getter semantics, error paths. PASS.
All regressions PASS: run_cases.sh (pose_corr @tol 0 replay, peak,
peak_recenter, avg_td, avg_td_oob) + pose fx/step/resident/DP/products
standalone tests against the rebuilt module. compute-sanitizer cannot
run on this GPU (sm_120 "Device not supported" - environment, not code).

Design: imagej-elphel-internal
handoffs/2026-07-16_3b_measure_chain_residency_design.md (B1; gate D3).
Co-Authored-By: 's avatarClaude Fable 5 <noreply@anthropic.com>
parent 978c5193
This diff is collapsed.
This diff is collapsed.
This diff is collapsed.
......@@ -11,6 +11,8 @@
#ifndef JCUDA
#include <cuda_runtime.h>
#include "tp_defines.h" // NUM_CAMS (for geometry_correction.h) // By Claude on 07/16/2026
#include "geometry_correction.h" // TP_TASK_*_OFFSET task-stream field offsets (pose_task_update) // By Claude on 07/16/2026
#include "tp_lma.h"
#endif
......@@ -473,6 +475,110 @@ extern "C" __global__ void pose_fx_jacobian(
valid[tile] = 1;
}
// 3-B rung B1 (design 2026-07-16): resident task build. One thread per selected
// task template entry runs the SAME C1-validated projection chain as
// pose_fx_jacobian (pose_world_from_pixel -> pose_pixel_from_world) and writes
// the 6-float task headers of BOTH resident ftask slots in place: slot 0 = MB
// main, slot 1 = MB partner. All per-tile variation comes from the projection;
// the uniform-MB contribution is the per-scene 6-float descriptor
// {dx_main, dy_main, dx_sub, dy_sub, scale_main, scale_sub} applied as
// constants (the crank math is tile-invariant under pose_mb_uniform). Task
// codes arrive pre-encoded from the host (already |511, matching
// GpuQuadJna.setTasks): no-MB mode = task_code_sub 511 and a zero descriptor.
// A tile whose projection fails or whose main/sub center falls outside the
// margin gate is written as a HOLE: task=0 in both sets (a value the Java
// upload path never produces) - index_direct skips it in convert_direct,
// index_consolidate drops it, the inter correlation never lists it, so
// downstream it is a missing peak and the per-cycle conditioning abstains
// (the Java build reached the same end by removing the entry from the
// condensed stream). xy/disp_dist tails are untouched - calculate_tiles_offsets
// overwrites them for every entry as before. The body is a grid-stride loop:
// a <<<1,1>>> launch of the SAME kernel is the bit-exact serial oracle
// (pose_lma_prepare_norm/finish precedent). By Claude on 07/16/2026.
extern "C" __global__ void pose_task_update(
int num_tasks,
int task_size, // floats per task entry (get_task_size(num_cams))
const PoseCameraMeta * cameras, // reference, scene
const float * ref_radial, // [7]
const float * ref_rbr,
const float * ref_ers, // [height][14]
const float * scene_radial, // [7]
const float * scene_ers, // [height][14]
const float * pose_vectors, // ref xyz/atr, scene xyz/atr: 4x3 (device current set)
const int * task_map, // [num_tasks] packed txy: tx | (ty<<16) (per-sequence template)
const float * task_centers, // [3*num_tasks] reference px,py,disparity (per-sequence template)
float disparity_corr,
float min_px, // margin gate (host-computed Java bounds)
float max_px,
float min_py,
float max_py,
int task_code_main, // pre-encoded (|511) task word, set 0
int task_code_sub, // pre-encoded task word, set 1 (511 = disabled no-MB partner)
float dx_main, // uniform-MB descriptor: centerXY offsets from the projected point
float dy_main,
float dx_sub,
float dy_sub,
float scale_main, // pre-encoded scale fields (Java setScale semantics)
float scale_sub,
float * ftasks0, // resident task slot 0 (MB main)
float * ftasks1) // resident task slot 1 (MB partner)
{
__shared__ float matrices[2 * 36];
if (threadIdx.x == 0) {
pose_rotation_matrices(pose_vectors + 3, true, matrices); // reference inverse (+derivs)
pose_rotation_matrices(pose_vectors + 9, false, matrices + 36); // scene direct
}
__syncthreads();
for (int task = (int) (blockIdx.x * blockDim.x + threadIdx.x); task < num_tasks;
task += (int) (gridDim.x * blockDim.x)) {
float * t0 = ftasks0 + (size_t) task * task_size;
float * t1 = ftasks1 + (size_t) task * task_size;
const int txy = task_map[task];
*(int *) (t0 + TP_TASK_TXY_OFFSET) = txy;
*(int *) (t1 + TP_TASK_TXY_OFFSET) = txy;
// hole defaults first; overwritten below when the tile survives
*(int *) (t0 + TP_TASK_TASK_OFFSET) = 0;
*(int *) (t1 + TP_TASK_TASK_OFFSET) = 0;
t0[TP_TASK_DISPARITY_OFFSET] = 0.0f;
t1[TP_TASK_DISPARITY_OFFSET] = 0.0f;
t0[TP_TASK_CENTERXY_OFFSET] = 0.0f; t0[TP_TASK_CENTERXY_OFFSET + 1] = 0.0f;
t1[TP_TASK_CENTERXY_OFFSET] = 0.0f; t1[TP_TASK_CENTERXY_OFFSET + 1] = 0.0f;
t0[TP_TASK_SCALE_OFFSET] = 0.0f;
t1[TP_TASK_SCALE_OFFSET] = 0.0f;
const int ci = 3 * task;
const float px = task_centers[ci];
const float py = task_centers[ci + 1];
const float disparity = task_centers[ci + 2];
if (!isfinite(px) || !isfinite(py) || !isfinite(disparity)) continue;
const bool infinity = disparity < POSE_INFINITY_DISPARITY;
float world[3], d_pixel[9], d_atr[9];
if (!pose_world_from_pixel(cameras, ref_radial, ref_rbr, ref_ers,
px, py, disparity, infinity, pose_vectors, matrices,
world, d_pixel, d_atr)) continue;
float scene_pixel[3];
if (!pose_pixel_from_world(cameras + 1, scene_radial, scene_ers,
world, infinity, pose_vectors + 6, matrices + 36, scene_pixel)) continue;
const float cx_main = scene_pixel[0] + dx_main;
const float cy_main = scene_pixel[1] + dy_main;
const float cx_sub = scene_pixel[0] + dx_sub;
const float cy_sub = scene_pixel[1] + dy_sub;
// Java setInterTasksMotionBlur margin gate: main AND sub centers inside
if ((cx_main < min_px) || (cx_main > max_px) || (cy_main < min_py) || (cy_main > max_py) ||
(cx_sub < min_px) || (cx_sub > max_px) || (cy_sub < min_py) || (cy_sub > max_py)) continue;
const float target_disparity = scene_pixel[2] + disparity_corr;
if (!isfinite(target_disparity)) continue;
*(int *) (t0 + TP_TASK_TASK_OFFSET) = task_code_main;
*(int *) (t1 + TP_TASK_TASK_OFFSET) = task_code_sub;
t0[TP_TASK_DISPARITY_OFFSET] = target_disparity;
t1[TP_TASK_DISPARITY_OFFSET] = target_disparity;
t0[TP_TASK_CENTERXY_OFFSET] = cx_main; t0[TP_TASK_CENTERXY_OFFSET + 1] = cy_main;
t1[TP_TASK_CENTERXY_OFFSET] = cx_sub; t1[TP_TASK_CENTERXY_OFFSET + 1] = cy_sub;
t0[TP_TASK_SCALE_OFFSET] = scale_main;
t1[TP_TASK_SCALE_OFFSET] = scale_sub;
}
}
static __device__ __forceinline__ float pose_lma_mul(float a, float b)
{
return __fmul_rn(a, b);
......
......@@ -79,6 +79,38 @@ extern "C" __global__ void pose_fx_jacobian(
float * jt, // parameter-major [3][2*num_tiles], nullable
unsigned char * valid);
// 3-B rung B1: resident task build - the pose_fx_jacobian projection chain
// writes the 6-float task headers of both resident ftask slots from a
// per-sequence template + a per-scene uniform-MB descriptor. Grid-stride body:
// a <<<1,1>>> launch is the bit-exact serial oracle. By Claude on 07/16/2026.
extern "C" __global__ void pose_task_update(
int num_tasks,
int task_size,
const PoseCameraMeta * cameras, // reference, scene
const float * ref_radial, // [7]
const float * ref_rbr,
const float * ref_ers, // [height][14]
const float * scene_radial, // [7]
const float * scene_ers, // [height][14]
const float * pose_vectors, // ref xyz/atr, scene xyz/atr: 4x3
const int * task_map, // [num_tasks] packed txy
const float * task_centers, // [3*num_tasks] reference px,py,disparity
float disparity_corr,
float min_px,
float max_px,
float min_py,
float max_py,
int task_code_main,
int task_code_sub,
float dx_main,
float dy_main,
float dx_sub,
float dy_sub,
float scale_main,
float scale_sub,
float * ftasks0,
float * ftasks1);
// Fixed three-angle candidate calculation in one float thread. Output layout:
// {valid, raw H[9], raw b[3], delta[3], candidate[3]}.
extern "C" __global__ void pose_lma_step(
......
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