Commit e1913933 authored by Andrey Filippov's avatar Andrey Filippov

CLAUDE: M1 - erase-by-task-list in the pose measure chain

Margin rung M1 (ROADMAP 3-A4i D5 tail, item (f) - suspect CONFIRMED by code:
erase_clt_tiles NaN-fills the full 16x5120x256-float CLT buffer (~84 MB)
every cycle while the chain converts/reads only the 150 task tiles):
- new erase_clt_task_tiles(_inner): poison ONLY the task-listed tiles
  (~2.4 MB) - task-tile bytes identical to the full erase (a task tile that
  fails conversion stays NaN), non-task tiles keep stale data no pose-chain
  consumer reads (consolidation/inter-corr/peaks are all task/index-driven).
- pose_measure chain (tp_lma.cu) switches to it; the full-frame
  erase_clt_tiles stays UNCHANGED for every render/debug/legacy path.
- module function list +1 (Java Stage0 41/41 in the main-repo commit).
Gates: run_cases.sh ALL PASS (pose_corr @tol 0 = captured production data
bit-identical through the list-erase chain; scene_dp/scene_dp_split/
measure_dp bit-identical); sanitizer 0 errors. A/B microbench: chain 2.52 ->
~2.35 ms/cycle => ~0.6-0.8 ms/scene off the DP entry.
Co-authored-by: 's avatarClaude Opus 4.8 <claude-opus-4-8@anthropic.com>
Co-authored-by: 's avatarClaude Fable 5 <claude-fable-5@anthropic.com>
parent 25ddfe30
......@@ -49,6 +49,7 @@ static const char* KERNELS[] = {
"convert_direct","imclt_rbg_all","correlate2D","correlate2D_inter","corr2D_combine","corr2D_normalize",
"textures_nonoverlap","generate_RBGA","clear_texture_list","mark_texture_tiles","mark_texture_neighbor_tiles",
"gen_texture_list","clear_texture_rbga","textures_accumulate","create_nonoverlap_list","erase_clt_tiles",
"erase_clt_task_tiles", // M1: erase-by-task-list (pose chain) // By Claude on 07/17/2026
"calculate_tiles_offsets","calc_rot_deriv","calcReverseDistortionTable",
"condition_lwir_f","condition_lwir_d", // production float maps + Java-matched validation // By Codex on 07/15/2026
"clt_average_sensors","index_consolidate","consolidate_oob","clt_average_sensors_list", // By Claude on 07/12/2026
......
......@@ -2842,6 +2842,60 @@ extern "C" __global__ void erase_clt_tiles_inner(
}
}
/**
* Margin rung M1 (pose tail, 07/17/2026): erase-by-task-list - NaN-poison ONLY
* the tiles the task list touches instead of the whole CLT buffer. The pose
* measure chain converts and reads task tiles exclusively (consolidation,
* inter-correlation and peaks are all task/index-list driven), so poisoning
* the 150 task tiles preserves the exact contract (a task tile that fails
* conversion stays NaN and is excluded) while writing ~2.4 MB instead of
* ~84 MB per cycle. Task-tile bytes are IDENTICAL to the full erase; non-task
* tiles keep stale data that no pose-chain consumer reads. The full-frame
* erase_clt_tiles above is UNCHANGED for every render/debug/legacy path that
* does read the whole buffer. By Claude on 07/17/2026.
*/
extern "C" __global__ void erase_clt_task_tiles_inner(
int num_cams, // for get_task_size/get_task_txy
int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
int tiles_x,
float * gpu_ftasks, // flattened task stream (slot 0 = the SET task list)
float ** gpu_clt, // [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE]
float fill_data)
{
int txy = get_task_txy(blockIdx.x, gpu_ftasks, num_cams);
int tile = (txy & 0xffff) + ((txy >> 16) & 0xffff) * tiles_x;
int tile_size = num_colors * (4*DTT_SIZE*DTT_SIZE);
float * data = gpu_clt[blockIdx.y] + tile_size * tile + threadIdx.x;
for (int ncol = 0; ncol < num_colors; ncol++){
#pragma unroll
for (int i = 0; i < (4*DTT_SIZE*DTT_SIZE/NUM_THREADS); i++){
*data = fill_data;
data += NUM_THREADS;
}
}
}
extern "C" __global__ void erase_clt_task_tiles(
int num_cams, // actual number of cameras
int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
int tiles_x,
float * gpu_ftasks, // flattened task stream (slot 0)
int num_tiles, // task-list length
float ** gpu_clt, // [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE]
float fill_data)
{
if (threadIdx.x == 0) { // anyway 1,1,1
dim3 threads_erase(NUM_THREADS, 1, 1); // (32,1,1)
dim3 grid_erase (num_tiles, num_cams, 1);
erase_clt_task_tiles_inner<<<grid_erase,threads_erase>>>(
num_cams, // int num_cams,
num_colors, // int num_colors,
tiles_x, // int tiles_x,
gpu_ftasks, // float * gpu_ftasks,
gpu_clt, // float ** gpu_clt,
fill_data); // float fill_data)
}
}
/**
......
......@@ -232,6 +232,26 @@ extern "C" __global__ void erase_clt_tiles_inner(
float ** gpu_clt, // [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE]
float fill_data);
// M1 (07/17/2026): erase-by-task-list - poison only the task-listed tiles
// (the pose measure chain's exact read set) instead of the full CLT buffer.
// By Claude on 07/17/2026.
extern "C" __global__ void erase_clt_task_tiles(
int num_cams, // actual number of cameras
int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
int tiles_x,
float * gpu_ftasks, // flattened task stream (slot 0 = the SET task list)
int num_tiles, // task-list length
float ** gpu_clt, // [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE]
float fill_data);
extern "C" __global__ void erase_clt_task_tiles_inner(
int num_cams, // for get_task_size/get_task_txy
int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
int tiles_x,
float * gpu_ftasks, // flattened task stream (slot 0)
float ** gpu_clt, // [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE]
float fill_data);
extern "C" __global__ void imclt_rbg(
float * gpu_clt, // [TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float * gpu_rbg, // WIDTH, 3 * HEIGHT
......
......@@ -1412,8 +1412,12 @@ extern "C" __global__ void pose_measure_dp(const PoseMeasureChain * c)
c->uniform_grid, c->num_cams, c->ftasks0, c->num_tasks,
(struct gc *) c->gc, (struct corr_vector *) c->cv,
(float *) c->rbr, (trot_deriv *) c->rot);
erase_clt_tiles<<<1, 1, 0, cudaStreamTailLaunch>>>(
c->num_cams, c->num_colors, c->tilesx, c->tilesy, c->clt, c->erase_fill);
// M1 (07/17/2026): erase-by-task-list - poison only the tiles this chain
// converts/reads (~2.4 MB) instead of the full CLT buffer (~84 MB/cycle);
// task-tile bytes identical, non-task tiles unread here. By Claude on 07/17/2026.
erase_clt_task_tiles<<<1, 1, 0, cudaStreamTailLaunch>>>(
c->num_cams, c->num_colors, c->tilesx, c->ftasks0, c->num_tasks,
c->clt, c->erase_fill);
convert_direct<<<1, 1, 0, cudaStreamTailLaunch>>>(
c->num_cams, c->num_colors, c->kernel_offsets, c->kernels, c->images,
c->ftasks0, c->clt, c->dstride_img, c->num_tasks, 0,
......
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