CLAUDE: clt_average_sensors v1 chain - MB-aware indexing + OOB priority + masked average

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: Claude Fable 5 <noreply@anthropic.com>

src/tests/make_avg_td_oob_case.py

+#!/usr/bin/env python3
+# make_avg_td_oob_case.py - synthetic test cases for the clt_average_sensors
+# kernels (roadmap 2c), independent of any Java run.
+#
+# Generates TWO cases under testdata/synth4/:
+#   avg_td/      - v0 uniform average (verifies the v0 kernel after the
+#                  consolidate_chunk refactor: same tiles, no masks)
+#   avg_td_oob/  - the v1 chain (index_consolidate -> consolidate_oob ->
+#                  clt_average_sensors_list)
+#
+# The scenario is the SAME one the Java oracle self-test passes
+# (imagej-elphel CuasTD.main(), 18/18 on 2026-07-12): 4 sensors, 4x2 tile
+# grid, per-tile per-sensor frame-edge depths chosen to hit every branch of
+# the frozen OOB spec (soft=12 / hard=8):
+#   tile 0: all deep                    -> mask 0b1111
+#   tile 1: deep,deep,soft-zone,<hard   -> mask 0b0011 (soft tier wins)
+#   tile 2: ALL in the soft zone        -> mask 0b1111 (last resort keeps all)
+#   tile 3: all below hard              -> dropped (poisoned NaN, count 0)
+#   tile 4: no MB counterpart (task==0 in set 1)   -> dropped
+#   tile 5: task==0 in set 0                        -> dropped
+#   tile 6: sensor 1 xy = NaN (undefined projection) -> mask 0b1101
+#   tile 7: task==0 in both sets                    -> dropped
+# Expected buffers are computed here in float32 with the ORACLE op order
+# (sequential sum in sensor order, multiply by reciprocal) -> tol 0 holds.
+#
+# Usage: python3 make_avg_td_oob_case.py   (from src/tests/ or anywhere)
+#
+# By Claude on 07/12/2026.
+
+import os
+import struct
+import numpy as np
+
+NS, TX, TY, COLORS = 4, 4, 2, 1
+W, H = 640.0, 512.0
+SOFT, HARD = 12.0, 8.0
+TD = 256                              # TD_TILE_FLOATS
+TASK_SIZE = 6 + 6 * NS                # get_task_size(NS): 6 + 2*NS xy + 4*NS disp_dist
+TASK_AVG_SENS_SHIFT = 11              # informational (masks are checked, not built, here)
+NT = TX * TY
+
+out_root = os.path.join(os.path.dirname(os.path.abspath(__file__)),
+                        "..", "..", "testdata", "synth4")
+
+# ---- per-tile per-sensor depths (distance of the tile center to the LEFT
+# ---- frame edge; y is centered so x IS the min distance to any edge)
+depths0 = [
+    [100, 200, 300, 400],   # 0
+    [100,  50,  10,   5],   # 1
+    [  9,  10,  11, 11.9],  # 2
+    [  3,   5,   0,  7.9],  # 3
+    [100, 100, 100, 100],   # 4
+    [100, 100, 100, 100],   # 5 (task==0 in set 0)
+    [100, 100, 100, 100],   # 6 (sensor 1 xy NaN in set 0)
+    [100, 100, 100, 100],   # 7 (task==0 both sets)
+]
+task0  = [1, 1, 1, 1, 1, 0, 1, 0]     # set-0 task word per tile
+task1  = [1, 1, 1, 1, 0, 1, 1, 0]     # set-1: tile 4 missing, tile 5 active
+
+expected_masks = [0b1111, 0b0011, 0b1111, 0, 0, 0, 0b1101, 0]
+
+def build_task_set(tasks, depth_rows, nan_sensor1_tile6, depth_offset):
+    """Dense full-grid task stream, exact struct tp_task float layout."""
+    ft = np.zeros((NT, TASK_SIZE), dtype=np.float32)
+    for nt in range(NT):
+        ty, tx = divmod(nt, TX)
+        ft[nt, 0] = struct.unpack('<f', struct.pack('<i', tasks[nt]))[0]      # task (int bit-cast)
+        ft[nt, 1] = struct.unpack('<f', struct.pack('<i', tx + (ty << 16)))[0]  # txy
+        # [2] disparity, [3:5] centerXY, [5] scale stay 0
+        for s in range(NS):
+            x = depth_rows[nt][s] + depth_offset
+            y = H / 2.0
+            if nan_sensor1_tile6 and (nt == 6) and (s == 1):
+                x = float('nan')
+            ft[nt, 6 + 2 * s]     = x   # xy[s][0]
+            ft[nt, 6 + 2 * s + 1] = y   # xy[s][1]
+        # disp_dist [6+2*NS ...] stays 0 (host marshalling leaves it 0 too)
+    return ft.reshape(-1)
+
+ftasks0 = build_task_set(task0, depths0, True,  0.0)
+ftasks1 = build_task_set(task1, depths0, False, 1.0)   # MB pair: 1 px deeper
+
+# ---- TD data: sensor s = constant s+1 everywhere; sensor 2 ABSENT on tile 5
+# ---- (whole chunk NaN - exercises the presence test independently of masks)
+td = np.zeros((NS, NT, TD), dtype=np.float32)
+for s in range(NS):
+    td[s, :, :] = np.float32(s + 1)
+td[2, 5, :] = np.nan
+
+def oracle_avg(masks):
+    """Oracle-op-order float32 average; masks[nt]=0 -> NaN tile, count 0."""
+    avg = np.full((NT, TD), np.nan, dtype=np.float32)
+    counts = np.zeros(NT, dtype=np.float32)
+    for nt in range(NT):
+        contrib = [s for s in range(NS)
+                   if ((masks[nt] >> s) & 1) and not np.isnan(td[s, nt, 0])]
+        counts[nt] = len(contrib)
+        if not contrib:
+            continue
+        acc = td[contrib[0], nt, :].copy()          # oracle: arraycopy first
+        for s in contrib[1:]:
+            acc = acc + td[s, nt, :]                # then add, sensor order
+        if len(contrib) > 1:
+            acc = acc * (np.float32(1.0) / np.float32(len(contrib)))  # reciprocal multiply
+        avg[nt, :] = acc
+    return avg, counts
+
+def write_case(path, prms, bufs):
+    os.makedirs(path, exist_ok=True)
+    man = []
+    for k, v in prms:
+        man.append(f"prm {k} {v}")
+    for name, arr, dims in bufs:
+        dtype = 'i32' if arr.dtype == np.int32 else 'f32'
+        fname = f"{name}.{dtype}"
+        arr.astype('<i4' if dtype == 'i32' else '<f4').tofile(os.path.join(path, fname))
+        man.append(f"buf {name} {dtype} {len(dims)} " +
+                   " ".join(str(d) for d in dims) + f" {fname}")
+    with open(os.path.join(path, "manifest.txt"), 'w') as f:
+        f.write("# synthetic case, generator: src/tests/make_avg_td_oob_case.py\n")
+        f.write("\n".join(man) + "\n")
+    print(f"wrote {path}")
+
+# ---- v0 case: plain average, ALL tiles, no masks (NaN-presence only) --------
+avg_v0, counts_v0 = oracle_avg([0b1111] * NT)
+write_case(os.path.join(out_root, "avg_td"),
+    [("num_sensors", NS), ("tilesx", TX), ("tilesy", TY), ("colors", COLORS)],
+    [("td_sens",         td.reshape(-1),      [NS, TY, TX, COLORS, TD]),
+     ("expected_td_avg", avg_v0.reshape(-1),  [TY, TX, COLORS, TD]),
+     ("expected_counts", counts_v0,           [NT])])
+
+# ---- v1 case: the OOB chain -------------------------------------------------
+avg_v1, counts_v1 = oracle_avg(expected_masks)
+write_case(os.path.join(out_root, "avg_td_oob"),
+    [("num_sensors", NS), ("tilesx", TX), ("tilesy", TY), ("colors", COLORS),
+     ("img_width", int(W)), ("img_height", int(H)),
+     ("soft_margin", SOFT), ("hard_margin", HARD),
+     ("has_mb", 1), ("task_size", TASK_SIZE)],
+    [("td_sens",             td.reshape(-1),                        [NS, TY, TX, COLORS, TD]),
+     ("ftasks0",             ftasks0,                               [NT, TASK_SIZE]),
+     ("ftasks1",             ftasks1,                               [NT, TASK_SIZE]),
+     ("expected_sens_masks", np.array(expected_masks, np.int32),    [TY, TX]),
+     ("expected_td_avg",     avg_v1.reshape(-1),                    [TY, TX, COLORS, TD]),
+     ("expected_counts",     counts_v1,                             [NT])])
+
+print("expected masks:", [f"0x{m:x}" for m in expected_masks])

src/tests/test_avg_td_oob.cu

+/*
+ * test_avg_td_oob.cu
+ *
+ * Standalone per-kernel test: the clt_average_sensors v1 chain (roadmap 2c):
+ *   index_consolidate (MB-pair tile selection) ->
+ *   consolidate_oob   (out-of-bounds soft/hard source-coordinate priority) ->
+ *   clt_average_sensors_list (mask-gated consolidation).
+ * Loads a Java-oracle case exported by AvgTdExport.exportOob() (RAW flattened
+ * task streams for both MB sets, margins in the manifest, expected per-tile
+ * sensor masks / TD average / counts from CuasTD.oobSensorMasks +
+ * consolidateSensorsTD). Deterministic single run - set breakpoints in any of
+ * the three kernels and step in NSight/cuda-gdb.
+ *
+ * Usage: test_avg_td_oob --data <dir> [--runs N] [--tol X] [--list]
+ *
+ *  Created on: Jul 12, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <string>
+#include <vector>
+
+#include <cuda_runtime.h>
+#include <helper_cuda.h>      // checkCudaErrors, findCudaDevice
+
+#include "tp_defines.h"            // NUM_CAMS (for geometry_correction.h)
+#include "geometry_correction.h"   // TP_TASK_*_OFFSET task field offsets
+#include "tp_consolidate.h"
+#include "tests/tp_test_data.h"
+#include "tests/tp_gpu_buf.h"
+
+int main(int argc, char **argv)
+{
+	std::string data_dir = "testdata/avg_td_oob";
+	int   num_runs  = 1;
+	float tol       = 0.0f;   // the whole chain is expected bit-exact vs the oracle
+	int   list_only = 0;
+	for (int i = 1; i < argc; i++) {
+		if      (!strcmp(argv[i], "--data") && (i + 1 < argc)) data_dir = argv[++i];
+		else if (!strcmp(argv[i], "--runs") && (i + 1 < argc)) num_runs = atoi(argv[++i]);
+		else if (!strcmp(argv[i], "--tol")  && (i + 1 < argc)) tol = (float) atof(argv[++i]);
+		else if (!strcmp(argv[i], "--list"))                   list_only = 1;
+		else {
+			fprintf(stderr, "Usage: %s --data <dir> [--runs N] [--tol X] [--list]\n", argv[0]);
+			return EXIT_FAILURE;
+		}
+	}
+	printf("%s: data=%s runs=%d tol=%g\n", argv[0], data_dir.c_str(), num_runs, tol);
+
+	findCudaDevice(argc, (const char **) argv);
+	TpTestData data(data_dir);
+	if (list_only) {
+		data.list();
+		return EXIT_SUCCESS;
+	}
+
+	int   num_sensors = data.paramInt("num_sensors", 16);
+	int   tilesx      = data.paramInt("tilesx",  80);
+	int   tilesy      = data.paramInt("tilesy",  64);
+	int   colors      = data.paramInt("colors",   1);
+	int   task_size   = data.paramInt("task_size", 6 + 6 * num_sensors);
+	float img_width   = (float) data.param("img_width",  640);
+	float img_height  = (float) data.param("img_height", 512);
+	float soft_margin = (float) data.param("soft_margin", 12.0);
+	float hard_margin = (float) data.param("hard_margin",  8.0);
+	int   has_mb      = data.paramInt("has_mb", 0);
+	int   num_entries = (int) (data.elements("ftasks0") / task_size); // entries per input set
+	int   num_chunks  = tilesy * tilesx * colors;
+	printf("sensors=%d grid=%dx%dx%d frame=%gx%g soft/hard=%g/%g mb=%d task_size=%d entries=%d\n",
+			num_sensors, tilesy, tilesx, colors, img_width, img_height,
+			soft_margin, hard_margin, has_mb, task_size, num_entries);
+
+	float * gpu_td_sens = data.gpuFloat("td_sens");
+	float * gpu_ftasks0 = data.gpuFloat("ftasks0");
+	float * gpu_ftasks1 = (has_mb && data.has("ftasks1")) ? data.gpuFloat("ftasks1") : nullptr;
+
+	GpuBuf<float> gpu_pairs((size_t) num_entries * 2 * task_size); // worst case: all kept
+	GpuBuf<float> gpu_tasks_out((size_t) num_entries * task_size);
+	GpuBuf<int>   gpu_counters(2);                                 // [0] pairs, [1] out
+	GpuBuf<int>   gpu_misaligned(1);
+	GpuBuf<float> gpu_td_avg((size_t) num_chunks * TD_TILE_FLOATS);
+	GpuBuf<float> gpu_counts((size_t) num_chunks);
+
+	int num_pairs = 0, num_out = 0, num_misaligned = 0;
+	const int IDX_THREADS = 256;
+	for (int run = 0; run < num_runs; run++) {
+		checkCudaErrors(cudaMemset(gpu_counters.dev(),   0, 2 * sizeof(int)));
+		checkCudaErrors(cudaMemset(gpu_misaligned.dev(), 0, sizeof(int)));
+		// host prefill contract of clt_average_sensors_list: tiles absent from
+		// the compacted list stay NaN / count 0 (= the oracle poison semantics);
+		// 0xFF byte fill makes every float the (negative, quiet) NaN 0xFFFFFFFF
+		checkCudaErrors(cudaMemset(gpu_td_avg.dev(), 0xFF,
+				(size_t) num_chunks * TD_TILE_FLOATS * sizeof(float)));
+		checkCudaErrors(cudaMemset(gpu_counts.dev(), 0, (size_t) num_chunks * sizeof(float)));
+
+		// pass 1: MB-pair tile selection -> compacted even/odd pairs
+		index_consolidate<<<(num_entries + IDX_THREADS - 1) / IDX_THREADS, IDX_THREADS>>>(
+				task_size, gpu_ftasks0, gpu_ftasks1, num_entries,
+				gpu_pairs.dev(), gpu_counters.dev() + 0, gpu_misaligned.dev());
+		checkCudaErrors(cudaGetLastError());
+		checkCudaErrors(cudaMemcpy(&num_pairs, gpu_counters.dev(), sizeof(int),
+				cudaMemcpyDeviceToHost));
+
+		// pass 2: OOB soft/hard priority -> one entry per surviving tile,
+		// sensor mask in task bits TASK_AVG_SENS_SHIFT..
+		if (num_pairs > 0) {
+			consolidate_oob<<<(num_pairs + IDX_THREADS - 1) / IDX_THREADS, IDX_THREADS>>>(
+					task_size, num_sensors, gpu_pairs.dev(), num_pairs,
+					img_width, img_height, soft_margin, hard_margin,
+					gpu_tasks_out.dev(), gpu_counters.dev() + 1);
+			checkCudaErrors(cudaGetLastError());
+		}
+		checkCudaErrors(cudaMemcpy(&num_out, gpu_counters.dev() + 1, sizeof(int),
+				cudaMemcpyDeviceToHost));
+
+		// pass 3: mask-gated consolidation over the compacted list
+		if (num_out > 0) {
+			dim3 blocks(num_out, colors);
+			clt_average_sensors_list<<<blocks, TD_TILE_FLOATS>>>(
+					num_sensors, gpu_td_sens, tilesx, tilesy, colors,
+					gpu_tasks_out.dev(), task_size, num_out,
+					gpu_td_avg.dev(), gpu_counts.dev());
+			checkCudaErrors(cudaGetLastError());
+		}
+	}
+	checkCudaErrors(cudaDeviceSynchronize());
+	checkCudaErrors(cudaMemcpy(&num_misaligned, gpu_misaligned.dev(), sizeof(int),
+			cudaMemcpyDeviceToHost));
+	printf("chain done: %d pair(s) -> %d surviving tile(s), %d misaligned, %d run(s)\n",
+			num_pairs, num_out, num_misaligned, num_runs);
+
+	int fail = 0;
+	if (num_misaligned) {
+		printf("FAIL: %d misaligned MB pairs (input sets must be index-aligned)\n",
+				num_misaligned);
+		fail = 1;
+	}
+
+	// masks check (host-side, exact): rebuild the per-GRID-tile sensor-mask
+	// plane from the compacted pass-2 output (0 where no entry survived) and
+	// compare against the oracle CuasTD.oobSensorMasks()
+	{
+		std::vector<float> tasks_out = gpu_tasks_out.download();
+		std::vector<int>   masks((size_t) tilesy * tilesx, 0);
+		unsigned int all_mask = (1u << num_sensors) - 1;
+		for (int i = 0; i < num_out; i++) {
+			const float * t = tasks_out.data() + (size_t) i * task_size;
+			int task = *(const int *) (t + TP_TASK_TASK_OFFSET);
+			int txy  = *(const int *) (t + TP_TASK_TXY_OFFSET);
+			int tx = txy & 0xffff, ty = txy >> 16;
+			if ((tx < tilesx) && (ty < tilesy)) {
+				masks[(size_t) ty * tilesx + tx] =
+						((unsigned int) task >> TASK_AVG_SENS_SHIFT) & all_mask;
+			}
+		}
+		int * expected = data.hostInt("expected_sens_masks");
+		int bad = 0;
+		long first_bad = -1;
+		for (size_t i = 0; i < masks.size(); i++) {
+			if (masks[i] != expected[i]) {
+				if (first_bad < 0) first_bad = (long) i;
+				bad++;
+			}
+		}
+		if (bad) {
+			printf("compare 'sens_masks': %d of %zu tiles mismatch, first @%ld "
+					"(got 0x%x, expected 0x%x) -> FAIL\n", bad, masks.size(), first_bad,
+					masks[first_bad], expected[first_bad]);
+			fail = 1;
+		} else {
+			printf("compare 'sens_masks': %zu tiles, exact match -> PASS\n", masks.size());
+		}
+	}
+
+	data.saveResult("td_avg", gpu_td_avg.dev(), {num_chunks, TD_TILE_FLOATS});
+	data.saveResult("counts", gpu_counts.dev(), {num_chunks});
+
+	int r = data.compare("td_avg", gpu_td_avg.dev(), tol);
+	if (r > 0) fail = 1;
+	r = data.compare("counts", gpu_counts.dev(), 0.0f); // counts must match exactly
+	if (r > 0) fail = 1;
+
+	printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+	return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

src/tp_consolidate.h

@@ -3,23 +3,52 @@
  *
  * Multi-sensor transform-domain (TD) consolidation kernels — the GPU side of
  * roadmap step 2c (16-sensor TD average with per-tile outlier rejection).
- * v0 = uniform average (bring-up of the per-kernel test workflow); OOB
- * source-coordinate priority and median/MAD consensus come next, extending
- * this same entry point.
  *
- *  Created on: Jul 9, 2026
+ * v0: clt_average_sensors — uniform NaN-aware average over the full tile grid
+ *     (bit-exact mirror of the Java oracle CuasTD.consolidateSensorsTD).
+ * v1 chain (Java oracle: CuasTD.oobSensorMasks + masked consolidateSensorsTD;
+ *     testdata cases from AvgTdExport.exportOob):
+ *       1. index_consolidate      — motion-blur-aware tile selection: keep only
+ *          tiles present (task != 0) in BOTH task sets (set 1 may be NULL =
+ *          no-MB mode, same convention as Java GpuQuad.setInterTasksMotionBlur
+ *          returning TpTask[2][] with [1] possibly null). Output = compacted
+ *          even/odd MB PAIRS (Andrey's design 07/12/2026); input may itself be
+ *          full-grid or already-compacted — the passes are not combined.
+ *       2. consolidate_oob        — out-of-bounds source-coordinate PRIORITY
+ *          (frozen spec 07/12/2026): per sensor, depth = min over the MB pair
+ *          of the tile-center distance to the nearest source-frame edge;
+ *          depth < hard (~8 px, 16x16 core overhang) -> sensor always skipped;
+ *          any sensor >= soft (~12 px, 24x24 MCLT support overhang) -> only
+ *          those are used; else the hard survivors serve as the last resort.
+ *          Nobody above hard -> tile dropped (poisoned). The resulting
+ *          per-tile sensor mask is folded into the task word (bits
+ *          TASK_AVG_SENS_SHIFT..+num_sensors-1 — free bits, see the TpTask
+ *          bit-field survey in imagej-elphel-internal handoffs 2026-07-12).
+ *       3. clt_average_sensors_list — the v0 average driven by the compacted
+ *          task list, gated by the per-tile sensor mask; tiles absent from the
+ *          list keep the host-prefilled NaN/0 (= oracle poison semantics).
+ *
+ *  Created on: Jul 9, 2026 (v1 chain added Jul 12, 2026)
  *      Author: Claude (Anthropic), for Elphel
  */
 
 #ifndef SRC_TP_CONSOLIDATE_H_
 #define SRC_TP_CONSOLIDATE_H_
 
-#include <cuda_runtime.h> // __global__ in the prototype below, also for non-nvcc parsers // By Claude on 07/12/2026
+#include <cuda_runtime.h> // __global__ in the prototypes below, also for non-nvcc parsers // By Claude on 07/12/2026
 
 // TD tile = 4 quadrants x 8x8 = 256 floats (see dtt8x8 / TileProcessor docs)
 #define TD_TILE_FLOATS 256
 // shared-memory presence array bound; matches GPUTileProcessor.MAX_NUM_CAMS
 #define CONSOLIDATE_MAX_SENSORS 16
+// tp_task.task bits TASK_AVG_SENS_SHIFT..(+num_sensors-1) = "sensors to average"
+// mask written by consolidate_oob, consumed by clt_average_sensors_list.
+// Bits 11..31 of the task word are unused by ALL other producers/consumers
+// (verified survey 2026-07-12); bits 8/9/10 are the TASK_*_EN bits, 0-7 the
+// texture-neighbor bits. When this becomes a production task bit, add it to
+// the Java GPUTileProcessor constants AND its #define injection block so the
+// two sides cannot diverge.
+#define TASK_AVG_SENS_SHIFT 11
 
 extern "C" __global__ void clt_average_sensors(
 		int           num_sensors, // sensors to consolidate (16 for LWIR)
@@ -28,4 +57,39 @@ extern "C" __global__ void clt_average_sensors(
 		float       * gpu_td_avg,  // [num_tiles][TD_TILE_FLOATS]
 		float       * gpu_counts); // [num_tiles] contributors/tile, may be null
 
+extern "C" __global__ void index_consolidate(
+		int           task_size,        // floats per task entry (get_task_size(num_cams))
+		const float * gpu_ftasks0,      // task set 0, always present (full-grid or compacted)
+		const float * gpu_ftasks1,      // MB pair set (index-aligned with set 0) or NULL = no-MB
+		int           num_tiles,        // entries in EACH input set
+		float       * gpu_ftasks_pairs, // out: even/odd pairs, 2*task_size floats per kept tile
+		                                // (no-MB: odd = copy of even, so pass 2 stays uniform)
+		int         * pnum_pairs,       // atomic pair counter, init 0 (index_direct pattern)
+		int         * pnum_misaligned); // diagnostic: pairs skipped on txy mismatch, may be null
+
+extern "C" __global__ void consolidate_oob(
+		int           task_size,        // floats per task entry
+		int           num_sensors,      // mask width (<= CONSOLIDATE_MAX_SENSORS... <=21 by free bits)
+		const float * gpu_ftasks_pairs, // index_consolidate output (even/odd MB pairs)
+		int           num_pairs,        // *pnum_pairs from pass 1
+		float         width,            // source sensor frame, px
+		float         height,
+		float         soft_margin,      // ~12.0f support overhang; 0 = tier disabled
+		float         hard_margin,      // ~8.0f core overhang; 0 = tier disabled
+		float       * gpu_ftasks_out,   // out: ONE entry per surviving tile = even entry with the
+		                                // sensor mask folded into task bits (dropped tiles absent)
+		int         * pnum_out);        // atomic counter, init 0
+
+extern "C" __global__ void clt_average_sensors_list(
+		int           num_sensors,
+		const float * gpu_td_sens,   // [num_sensors][tilesy*tilesx*colors][TD_TILE_FLOATS]
+		int           tilesx,        // tile grid (for txy -> chunk mapping + bounds)
+		int           tilesy,
+		int           colors,
+		const float * gpu_ftasks,    // consolidate_oob output (sensor mask in task bits)
+		int           task_size,     // floats per task entry
+		int           num_tasks,     // *pnum_out from pass 2
+		float       * gpu_td_avg,    // [tilesy*tilesx*colors][TD_TILE_FLOATS], host-PREFILLED NaN
+		float       * gpu_counts);   // [tilesy*tilesx*colors], host-PREFILLED 0, may be null
+
 #endif /* SRC_TP_CONSOLIDATE_H_ */