CLAUDE: per-kernel test workflow: manifest-driven harness + clt_average_sensors v0

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

.gitignore

@@ -12,3 +12,4 @@ attic
 /clt*
 /clt_big_endian
 /*-old
+tests_bin/

eclipse_setup/nsight_per_kernel_tests.md

+# Per-kernel test binaries: build + NSight/cuda-gdb debugging
+<!-- By Claude on 07/09/2026 -->
+
+Each new kernel gets its own small test program in `src/tests/test_<kernel>.cu`
+(one deterministic kernel launch, data from a Java-oracle-exported directory —
+see `src/tests/testdata_format.md`). The old `test_tp` / Eclipse managed build
+is untouched.
+
+## Build
+
+```bash
+src/tests/build_tests.sh              # all tests, device-debug (-G -g -O0)
+src/tests/build_tests.sh test_avg_td  # one test
+RELEASE=1 src/tests/build_tests.sh    # optimized (-O2 -lineinfo), for timing
+```
+
+Binaries land in `tests_bin/` (gitignored). Uses CUDA 12.8 (`CUDA_HOME` to
+override) and `-arch=native`: the RTX 5060 Ti is Blackwell `sm_120`, which the
+old 12.6 toolchain cannot target natively — 12.8 is required for device
+stepping on this GPU.
+
+## Run (self-checking)
+
+```bash
+tests_bin/test_avg_td --data <case_dir> [--runs N] [--tol X] [--list]
+```
+
+Loads `<case_dir>/manifest.txt`, runs the kernel, writes `<case_dir>/results/`,
+compares against the oracle's `expected_*` buffers and prints PASS/FAIL
+(max|diff|, RMS, worst index). Exit code 0 = pass.
+
+## Debug from Eclipse (NSight EE plugin)
+
+One debug configuration per test binary — create once, reuse forever:
+
+1. Run → Debug Configurations → **CUDA GDB Application** → New.
+2. Name: `test_avg_td`; Application: `<repo>/tests_bin/test_avg_td`;
+   Project: tile_processor_gpu; Working directory: `<repo>`.
+3. Arguments tab: `--data <case_dir> --runs 1`.
+4. Debugger tab: cuda-gdb executable = `/usr/local/cuda-12.8/bin/cuda-gdb`
+   (the default 12.6 one predates this GPU).
+5. Breakpoint inside the kernel (e.g. `tp_consolidate.cu`,
+   `clt_average_sensors`) — step, inspect `threadIdx`/`blockIdx`, switch
+   threads with the CUDA view as usual.
+
+## Debug from the command line
+
+```bash
+/usr/local/cuda-12.8/bin/cuda-gdb --args tests_bin/test_avg_td --data <case_dir>
+(cuda-gdb) break clt_average_sensors
+(cuda-gdb) run
+(cuda-gdb) cuda block (3,0,0) thread (100,0,0)   # jump to a specific tile/coef
+```
+
+## Adding a new kernel test
+
+1. Kernel goes in `src/` (own file, e.g. `tp_<name>.cu` + `.h`).
+2. Copy `src/tests/test_avg_td.cu` → `test_<name>.cu`, adapt buffers/launch.
+3. Java side: export a test case with `com.elphel.imagej.gpu.testdata`
+   (`TestDataWriter`; see `AvgTdExport` as the model — export ALL inputs
+   plus `expected_*` from the CPU oracle).
+4. `build_tests.sh test_<name>`, run, then create the debug configuration
+   (steps above).

src/tests/build_tests.sh

+#!/usr/bin/env bash
+# Build per-kernel standalone test binaries (NSight/cuda-gdb debuggable).
+# By Claude on 07/09/2026.
+#
+# Each src/tests/test_*.cu becomes tests_bin/test_* — device-debug build
+# (-G -g -O0) with native GPU arch so cuda-gdb can step kernel code on
+# Blackwell (RTX 5060 Ti, sm_120 -> needs CUDA >= 12.8, default below).
+#
+# Usage: ./build_tests.sh [test_name ...]     (default: all test_*.cu)
+#        RELEASE=1 ./build_tests.sh           (optimized, no device debug)
+set -e
+cd "$(dirname "$0")"                     # src/tests
+SRC=..                                   # src/
+REPO=../..                               # repo root
+CUDA="${CUDA_HOME:-/usr/local/cuda-12.8}"
+SAMPLES="${CUDA_SAMPLES:-/home/elphel/git/cuda-samples/Common}"  # helper_cuda.h
+OUT="$REPO/tests_bin"
+mkdir -p "$OUT"
+
+if [ -n "$RELEASE" ]; then
+    DBG="-O2 -lineinfo"
+else
+    DBG="-G -g -O0"
+fi
+
+# Kernel + host sources shared by every test ("load all kernels"):
+COMMON="$SRC/TileProcessor.cu $SRC/dtt8x8.cu $SRC/geometry_correction.cu \
+        $SRC/tp_consolidate.cu $SRC/tp_utils.cu $SRC/tp_files.cu tp_test_data.cu"
+
+TESTS="${@:-$(ls test_*.cu | sed 's/\.cu$//')}"
+for t in $TESTS; do
+    echo "== building $t =="
+    "$CUDA/bin/nvcc" -std=c++17 $DBG -arch=native -rdc=true \
+        -I"$SRC" -I"$CUDA/include" -I"$SAMPLES" \
+        "$t.cu" $COMMON \
+        -o "$OUT/$t" \
+        -L"$CUDA/lib64" -lcudart
+    echo "   -> $OUT/$t"
+done
+echo "done: $OUT"

src/tests/test_avg_td.cu

+/*
+ * test_avg_td.cu
+ *
+ * Standalone per-kernel test: clt_average_sensors (16-sensor TD
+ * consolidation, roadmap step 2c). Loads a Java-oracle-exported test case
+ * (see testdata_format.md), runs the kernel once (deterministic — set a
+ * breakpoint inside clt_average_sensors and step in NSight/cuda-gdb),
+ * saves results and compares against the oracle's expected outputs.
+ *
+ * Usage: test_avg_td --data <dir> [--runs N] [--tol X] [--list]
+ *   <dir> must contain manifest.txt with at least:
+ *     prm num_sensors, buf td_sens f32 [nsens][tilesy][tilesx][colors][256]
+ *   optional oracle references: buf expected_td_avg, expected_counts
+ *
+ *  Created on: Jul 9, 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_consolidate.h"
+#include "tests/tp_test_data.h"
+
+int main(int argc, char **argv)
+{
+	std::string data_dir = "testdata/avg_td";
+	int   num_runs  = 1;
+	float tol       = 1e-5f;
+	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);
+	// td_sens dims: [num_sensors][tilesy][tilesx][colors][256] (or any shape
+	// whose product = num_sensors * num_tiles * 256)
+	int num_tiles = (int) (data.elements("td_sens") / ((size_t) num_sensors * TD_TILE_FLOATS));
+	printf("num_sensors=%d num_tiles=%d\n", num_sensors, num_tiles);
+
+	float * gpu_td_sens = data.gpuFloat("td_sens");
+	float * gpu_td_avg  {};
+	float * gpu_counts  {};
+	checkCudaErrors(cudaMalloc(&gpu_td_avg, (size_t) num_tiles * TD_TILE_FLOATS * sizeof(float)));
+	checkCudaErrors(cudaMalloc(&gpu_counts, (size_t) num_tiles * sizeof(float)));
+
+	dim3 blocks(num_tiles);
+	dim3 threads(TD_TILE_FLOATS);
+	for (int run = 0; run < num_runs; run++) {
+		clt_average_sensors<<<blocks, threads>>>(
+				num_sensors,
+				gpu_td_sens,
+				num_tiles,
+				gpu_td_avg,
+				gpu_counts);
+		checkCudaErrors(cudaGetLastError());
+	}
+	checkCudaErrors(cudaDeviceSynchronize());
+	printf("clt_average_sensors: %d run(s) done\n", num_runs);
+
+	data.saveResult("td_avg", gpu_td_avg, {num_tiles, TD_TILE_FLOATS});
+	data.saveResult("counts", gpu_counts, {num_tiles});
+
+	int fail = 0;
+	int r = data.compare("td_avg", gpu_td_avg, tol);
+	if (r > 0) fail = 1;
+	r = data.compare("counts", gpu_counts, 0.0f); // counts must match exactly
+	if (r > 0) fail = 1;
+
+	checkCudaErrors(cudaFree(gpu_td_avg));
+	checkCudaErrors(cudaFree(gpu_counts));
+	printf("%s: %s\n", argv[0], fail ? "FAIL" : "PASS");
+	return fail ? EXIT_FAILURE : EXIT_SUCCESS;
+}

src/tests/testdata_format.md

+# Test-data contract: Java oracle <-> standalone per-kernel GPU tests
+<!-- By Claude on 07/09/2026 -->
+
+One directory per test case. The Java oracle exports **everything** the standalone
+test needs (inputs, parameters, and its own expected outputs) so the test has no
+dependency on old calibration files or on a live Java process.
+
+```
+testdata/<scene>/<kernel>/          e.g. testdata/1773135476_186641/avg_td/
+    manifest.txt                    describes every buffer + scalar parameter
+    <name>.f32 / <name>.i32        raw little-endian arrays, no headers
+    results/                        written by the C++ test (same format back)
+        results_manifest.txt
+        <name>.f32
+```
+
+## manifest.txt
+
+Line-oriented, whitespace-separated, `#` = comment. Two record types:
+
+```
+prm <name> <value>                              scalar parameter (double)
+buf <name> <f32|i32> <ndims> <d0> ... <dN-1> <file>   array buffer
+```
+
+- Dims are row-major, **last dimension fastest** — matches flat Java arrays and C.
+- `<file>` is relative to the manifest's directory.
+- Expected oracle outputs use the `expected_` name prefix; the test compares its
+  GPU result against them and prints PASS/FAIL (max|diff|, RMS, worst index).
+
+Example (16-sensor TD consolidation):
+
+```
+# avg_td test case, exported by Java oracle
+prm num_sensors 16
+prm tilesx     80
+prm tilesy     64
+prm colors      1
+buf td_sens          f32 5 16 64 80 1 256  td_sens.f32
+buf task_xy          f32 3 5120 16 2       task_xy.f32
+buf tile_valid       i32 2 64 80           tile_valid.i32
+buf expected_td_avg  f32 4 64 80 1 256     expected_td_avg.f32
+```
+
+## Round trip
+
+1. Java oracle (CPU) computes reference, exports inputs + `expected_*`.
+2. Standalone test (`src/tests/test_<kernel>.cu`, debuggable in NSight) loads the
+   manifest, uploads, runs ONE kernel, writes `results/`, compares vs `expected_*`.
+3. Java importer reads `results/*.f32` back for ImageJ visual inspection.

src/tests/tp_test_data.cu

+/*
+ * tp_test_data.cu
+ *
+ * Manifest-driven test-data loader for standalone per-kernel GPU tests.
+ * See testdata_format.md for the manifest contract with the Java oracle.
+ *
+ *  Created on: Jul 9, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cmath>
+#include <cstring>
+#include <fstream>
+#include <sstream>
+#include <sys/stat.h>
+
+#include <cuda_runtime.h>
+#include <helper_cuda.h>   // checkCudaErrors
+
+#include "tp_test_data.h"
+
+size_t TpTestData::Buf::elements() const {
+	size_t n = 1;
+	for (int d : dims) n *= (size_t) d;
+	return n;
+}
+
+TpTestData::TpTestData(const std::string & dir) : m_dir(dir) {
+	std::string mpath = m_dir + "/manifest.txt";
+	std::ifstream mf(mpath);
+	if (!mf) {
+		fprintf(stderr, "TpTestData: cannot open %s\n", mpath.c_str());
+		exit(EXIT_FAILURE);
+	}
+	std::string line;
+	int lnum = 0;
+	while (std::getline(mf, line)) {
+		lnum++;
+		size_t hash = line.find('#');
+		if (hash != std::string::npos) line.erase(hash);
+		std::istringstream ls(line);
+		std::string rec;
+		if (!(ls >> rec)) continue; // blank line
+		if (rec == "prm") {
+			std::string name; double val;
+			if (!(ls >> name >> val)) {
+				fprintf(stderr, "TpTestData: bad prm line %d in %s\n", lnum, mpath.c_str());
+				exit(EXIT_FAILURE);
+			}
+			m_params[name] = val;
+		} else if (rec == "buf") {
+			Buf b;
+			int nd;
+			if (!(ls >> b.name >> b.dtype >> nd)) {
+				fprintf(stderr, "TpTestData: bad buf line %d in %s\n", lnum, mpath.c_str());
+				exit(EXIT_FAILURE);
+			}
+			if ((b.dtype != "f32") && (b.dtype != "i32")) {
+				fprintf(stderr, "TpTestData: unsupported dtype '%s' line %d\n", b.dtype.c_str(), lnum);
+				exit(EXIT_FAILURE);
+			}
+			for (int i = 0; i < nd; i++) {
+				int d;
+				if (!(ls >> d)) {
+					fprintf(stderr, "TpTestData: bad dims line %d in %s\n", lnum, mpath.c_str());
+					exit(EXIT_FAILURE);
+				}
+				b.dims.push_back(d);
+			}
+			if (!(ls >> b.file)) {
+				fprintf(stderr, "TpTestData: missing file name line %d in %s\n", lnum, mpath.c_str());
+				exit(EXIT_FAILURE);
+			}
+			m_bufs[b.name] = b;
+		} else {
+			fprintf(stderr, "TpTestData: unknown record '%s' line %d in %s\n", rec.c_str(), lnum, mpath.c_str());
+			exit(EXIT_FAILURE);
+		}
+	}
+	printf("TpTestData: %s — %d buffers, %d parameters\n",
+			mpath.c_str(), (int) m_bufs.size(), (int) m_params.size());
+}
+
+TpTestData::~TpTestData() {
+	for (auto & kv : m_bufs) {
+		if (kv.second.gpu) cudaFree(kv.second.gpu);
+	}
+}
+
+bool TpTestData::has(const std::string & name) const {
+	return m_bufs.count(name) != 0;
+}
+
+void TpTestData::list() const {
+	printf("parameters:\n");
+	for (const auto & kv : m_params) printf("  prm %-24s %g\n", kv.first.c_str(), kv.second);
+	printf("buffers:\n");
+	for (const auto & kv : m_bufs) {
+		const Buf & b = kv.second;
+		printf("  buf %-24s %s [", b.name.c_str(), b.dtype.c_str());
+		for (size_t i = 0; i < b.dims.size(); i++) printf("%s%d", i ? "][" : "", b.dims[i]);
+		printf("] <- %s\n", b.file.c_str());
+	}
+}
+
+TpTestData::Buf & TpTestData::getBuf(const std::string & name) {
+	auto it = m_bufs.find(name);
+	if (it == m_bufs.end()) {
+		fprintf(stderr, "TpTestData: no buffer '%s' in %s/manifest.txt\n", name.c_str(), m_dir.c_str());
+		exit(EXIT_FAILURE);
+	}
+	return it->second;
+}
+
+int TpTestData::ndims(const std::string & name)              { return (int) getBuf(name).dims.size(); }
+int TpTestData::dim  (const std::string & name, int i)       { return getBuf(name).dims.at(i); }
+size_t TpTestData::elements(const std::string & name)        { return getBuf(name).elements(); }
+
+double TpTestData::param(const std::string & name, double dflt) {
+	auto it = m_params.find(name);
+	return (it == m_params.end()) ? dflt : it->second;
+}
+int TpTestData::paramInt(const std::string & name, int dflt) {
+	return (int) lround(param(name, dflt));
+}
+
+void TpTestData::loadHost(Buf & b) {
+	if (!b.host.empty()) return;
+	std::string path = m_dir + "/" + b.file;
+	size_t expect = b.elements() * 4;                 // f32 and i32 are both 4 bytes
+	std::ifstream f(path, std::ios::binary | std::ios::ate);
+	if (!f) {
+		fprintf(stderr, "TpTestData: cannot open %s\n", path.c_str());
+		exit(EXIT_FAILURE);
+	}
+	size_t fsize = (size_t) f.tellg();
+	if (fsize != expect) {
+		fprintf(stderr, "TpTestData: %s is %zu bytes, manifest dims say %zu\n",
+				path.c_str(), fsize, expect);
+		exit(EXIT_FAILURE);
+	}
+	b.host.resize(fsize);
+	f.seekg(0);
+	f.read(b.host.data(), fsize);
+	printf("TpTestData: loaded %-24s %8zu bytes from %s\n", b.name.c_str(), fsize, path.c_str());
+}
+
+float * TpTestData::hostFloat(const std::string & name) {
+	Buf & b = getBuf(name);
+	loadHost(b);
+	return (float *) b.host.data();
+}
+int * TpTestData::hostInt(const std::string & name) {
+	Buf & b = getBuf(name);
+	loadHost(b);
+	return (int *) b.host.data();
+}
+
+float * TpTestData::gpuFloat(const std::string & name) {
+	Buf & b = getBuf(name);
+	if (!b.gpu) {
+		loadHost(b);
+		checkCudaErrors(cudaMalloc(&b.gpu, b.host.size()));
+		checkCudaErrors(cudaMemcpy(b.gpu, b.host.data(), b.host.size(), cudaMemcpyHostToDevice));
+	}
+	return (float *) b.gpu;
+}
+int * TpTestData::gpuInt(const std::string & name) {
+	return (int *) (void *) gpuFloat(name); // same load/upload path, caller-typed
+}
+
+void TpTestData::ensureResultsDir() {
+	if (m_results_dir_ok) return;
+	std::string rdir = m_dir + "/results";
+	mkdir(rdir.c_str(), 0775); // ok if it already exists
+	m_results_dir_ok = true;
+}
+
+void TpTestData::saveResult(const std::string & name, const float * gpu_data,
+		const std::vector<int> & dims) {
+	ensureResultsDir();
+	size_t n = 1;
+	for (int d : dims) n *= (size_t) d;
+	std::vector<float> host(n);
+	checkCudaErrors(cudaMemcpy(host.data(), gpu_data, n * sizeof(float), cudaMemcpyDeviceToHost));
+	std::string path = m_dir + "/results/" + name + ".f32";
+	std::ofstream f(path, std::ios::binary);
+	f.write((const char *) host.data(), n * sizeof(float));
+	std::ofstream rm(m_dir + "/results/results_manifest.txt", std::ios::app);
+	rm << "buf " << name << " f32 " << dims.size();
+	for (int d : dims) rm << " " << d;
+	rm << " " << name << ".f32\n";
+	printf("TpTestData: saved  %-24s %8zu floats to %s\n", name.c_str(), n, path.c_str());
+}
+
+int TpTestData::compare(const std::string & name, const float * gpu_data, float tol) {
+	std::string ename = "expected_" + name;
+	if (!has(ename)) {
+		printf("TpTestData: no '%s' in manifest — comparison skipped\n", ename.c_str());
+		return -1;
+	}
+	float * expected = hostFloat(ename);
+	size_t n = elements(ename);
+	std::vector<float> computed(n);
+	checkCudaErrors(cudaMemcpy(computed.data(), gpu_data, n * sizeof(float), cudaMemcpyDeviceToHost));
+	double sum2 = 0.0, max_abs = 0.0;
+	size_t max_i = 0, num_nan = 0;
+	for (size_t i = 0; i < n; i++) {
+		float e = expected[i], c = computed[i];
+		if (std::isnan(e) || std::isnan(c)) {
+			if (std::isnan(e) != std::isnan(c)) num_nan++; // NaN mismatch counts, NaN==NaN ok
+			continue;
+		}
+		double d = (double) c - (double) e;
+		sum2 += d * d;
+		if (fabs(d) > max_abs) { max_abs = fabs(d); max_i = i; }
+	}
+	double rms = sqrt(sum2 / (double) n);
+	int fail = (max_abs > tol) || (num_nan > 0);
+	printf("compare '%s' vs '%s': %zu elements, max|diff|=%.6g @%zu, rms=%.6g, NaN-mismatch=%zu, tol=%g -> %s\n",
+			name.c_str(), ename.c_str(), n, max_abs, max_i, rms, num_nan, tol,
+			fail ? "FAIL" : "PASS");
+	return fail;
+}

src/tests/tp_test_data.h

+/*
+ * tp_test_data.h
+ *
+ * Manifest-driven test-data loader for standalone per-kernel GPU tests.
+ * Contract: src/tests/testdata_format.md (Java oracle exports inputs +
+ * expected outputs; this class loads/uploads them and checks results).
+ *
+ *  Created on: Jul 9, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#ifndef SRC_TESTS_TP_TEST_DATA_H_
+#define SRC_TESTS_TP_TEST_DATA_H_
+
+#include <string>
+#include <vector>
+#include <map>
+
+class TpTestData {
+public:
+	struct Buf {
+		std::string        name;
+		std::string        dtype;   // "f32" | "i32"
+		std::vector<int>   dims;    // row-major, last fastest
+		std::string        file;    // relative to m_dir
+		std::vector<char>  host;    // lazy-loaded raw bytes
+		void *             gpu {};  // lazy-uploaded device copy
+		size_t             elements() const;
+	};
+
+	explicit TpTestData(const std::string & dir); // parses <dir>/manifest.txt
+	~TpTestData();                                // frees GPU copies
+
+	// inventory
+	bool     has      (const std::string & name) const;
+	void     list     () const;                   // print all buffers/params
+	int      ndims    (const std::string & name);
+	int      dim      (const std::string & name, int i);
+	size_t   elements (const std::string & name);
+
+	// scalar parameters (prm lines)
+	double   param    (const std::string & name, double dflt);
+	int      paramInt (const std::string & name, int dflt);
+
+	// host access (loads file on first use)
+	float *  hostFloat(const std::string & name);
+	int   *  hostInt  (const std::string & name);
+
+	// device access (loads + cudaMalloc + upload on first use)
+	float *  gpuFloat (const std::string & name);
+	int   *  gpuInt   (const std::string & name);
+
+	// download a device buffer and write <dir>/results/<name>.f32 (+ manifest line)
+	void     saveResult(const std::string & name, const float * gpu_data,
+			const std::vector<int> & dims);
+
+	// download computed and compare against buffer "expected_<name>" (if present).
+	// Prints max|diff| (+ its index), RMS, PASS/FAIL vs tol. Returns 0 on pass,
+	// 1 on fail, -1 if no expected buffer exists (prints a note, not a failure).
+	int      compare(const std::string & name, const float * gpu_data, float tol);
+
+private:
+	std::string                 m_dir;
+	std::map<std::string, Buf>  m_bufs;
+	std::map<std::string, double> m_params;
+	Buf &    getBuf (const std::string & name);   // aborts with message if missing
+	void     loadHost(Buf & b);
+	void     ensureResultsDir();
+	bool     m_results_dir_ok {false};
+};
+
+#endif /* SRC_TESTS_TP_TEST_DATA_H_ */

src/tp_consolidate.cu

+/*
+ * tp_consolidate.cu
+ *
+ * Multi-sensor TD consolidation (roadmap step 2c). v0 mirrors the Java CPU
+ * oracle CuasTD.consolidateSensorsTD() exactly (NaN-aware average):
+ *   - a sensor is ABSENT for a tile if the first float of its 256-float TD
+ *     chunk is NaN (whole-tile validity marker, set by the Java side);
+ *   - present sensors are summed and divided by the contributing count;
+ *   - no contributors -> NaN; a stray INTERIOR NaN in a present sensor
+ *     poisons the whole output tile to NaN (fail-visible, same as oracle);
+ *   - counts plane (contributors per tile) is written for weighting/QC.
+ * Planned extensions (same kernel, added inputs):
+ *   1. source-coordinate priority: per-sensor tile XY vs frame edge — MCLT
+ *      24x24 support with replicated OOB rows/cols = defined-but-synthetic,
+ *      center within ~12 px of edge = degraded; weight/exclude accordingly;
+ *   2. statistical consensus: CC(0,0) flux + per-tile energy (sum coeff^2),
+ *      median/MAD across sensors, discard outliers from the average.
+ *
+ *  Created on: Jul 9, 2026
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#include "tp_consolidate.h"
+
+// One block per TD tile (chunk), one thread per TD coefficient (256/block).
+// Thread 0 evaluates per-sensor presence from coefficient 0 into shared
+// memory; then every thread sums its coefficient over present sensors.
+extern "C" __global__ void clt_average_sensors(
+		int           num_sensors, // sensors to consolidate (16 for LWIR)
+		const float * gpu_td_sens, // [num_sensors][num_tiles][TD_TILE_FLOATS]
+		int           num_tiles,   // tilesy * tilesx * colors
+		float       * gpu_td_avg,  // [num_tiles][TD_TILE_FLOATS]
+		float       * gpu_counts)  // [num_tiles] contributors/tile, may be null
+{
+	__shared__ int   present [CONSOLIDATE_MAX_SENSORS];
+	__shared__ int   num_present;
+	__shared__ int   poisoned;
+	int tile = blockIdx.x;
+	int coef = threadIdx.x;                    // 0..TD_TILE_FLOATS-1
+	if ((tile >= num_tiles) || (coef >= TD_TILE_FLOATS)) {
+		return;
+	}
+	size_t sens_stride = (size_t) num_tiles * TD_TILE_FLOATS;
+	size_t tile_offs   = (size_t) tile * TD_TILE_FLOATS;
+	if (coef == 0) {
+		num_present = 0;
+		poisoned = 0;
+		for (int nsens = 0; nsens < num_sensors; nsens++) {
+			// oracle rule: NaN in the chunk's first float == sensor absent
+			int pres = !isnan(gpu_td_sens[nsens * sens_stride + tile_offs]);
+			present[nsens] = pres;
+			num_present += pres;
+		}
+		if (gpu_counts) {
+			gpu_counts[tile] = (float) num_present;
+		}
+	}
+	__syncthreads();
+	float val = NAN;                           // num_present == 0 -> all-NaN tile
+	if (num_present > 0) {
+		float sum = 0.0f;
+		for (int nsens = 0; nsens < num_sensors; nsens++) {
+			if (present[nsens]) {
+				sum += gpu_td_sens[nsens * sens_stride + tile_offs + coef];
+			}
+		}
+		// multiply by precomputed reciprocal, same op order as the Java
+		// oracle (scale = 1.0f/n; avg *= scale) -> bit-identical results
+		val = sum * (1.0f / (float) num_present);
+		// oracle quirk kept: whole-tile poison on interior NaN applies only
+		// when n > 1; a single-contributor tile is copied verbatim
+		if ((num_present > 1) && isnan(val)) {
+			poisoned = 1;
+		}
+	}
+	__syncthreads();
+	if ((num_present > 1) && poisoned) {
+		val = NAN;
+	}
+	gpu_td_avg[tile_offs + coef] = val;
+}

src/tp_consolidate.h

+/*
+ * tp_consolidate.h
+ *
+ * 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
+ *      Author: Claude (Anthropic), for Elphel
+ */
+
+#ifndef SRC_TP_CONSOLIDATE_H_
+#define SRC_TP_CONSOLIDATE_H_
+
+// 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
+
+extern "C" __global__ void clt_average_sensors(
+		int           num_sensors, // sensors to consolidate (16 for LWIR)
+		const float * gpu_td_sens, // [num_sensors][num_tiles][TD_TILE_FLOATS]
+		int           num_tiles,   // tilesy * tilesx * colors
+		float       * gpu_td_avg,  // [num_tiles][TD_TILE_FLOATS]
+		float       * gpu_counts); // [num_tiles] contributors/tile, may be null
+
+#endif /* SRC_TP_CONSOLIDATE_H_ */