indexing: key the shared device tables on their content, not only on an address

The cache returned a device copy for a (device, host address) pair and cast it to
whatever the caller asked for, with nothing checking that the bytes behind that address
were still the same bytes. A host buffer can be mutated in place - PixelMask::LoadMask
does exactly that - or freed and reallocated at the same address, and either hands the
caller a device copy of something else. Nothing would report it: the tables are read-only
geometry, so the engine would simply mask the wrong pixels for the rest of the run while
the azimuthal mapping, the written pixel_mask dataset and the viewer overlay used the new
one. Today that is unreachable, but only because of two guards in unrelated files that
neither state nor assert the requirement.

The byte length and an FNV-1a checksum of the bytes being uploaded are now part of the
key. Both are computed once per engine construction, over a buffer that is about to be
copied to the device anyway, so the cost does not show. Expired entries are pruned on
insert, since distinct content now means distinct entries.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

image_analysis/indexing/CudaSharedTables.h

@@ -6,6 +6,7 @@
 #include <map>
 #include <memory>
 #include <mutex>
+#include <tuple>
 #include <utility>
 
 #include "CUDAMemHelpers.h"
@@ -22,14 +23,35 @@
 // host vector that produced it, which lives in the experiment / integration mapping and therefore
 // outlives every engine.
 //
+// A bare address is not enough on its own to say "same table", though: a host buffer can be mutated
+// in place, or freed and a new one allocated at the same address, and either would hand the caller a
+// device copy of something else - silently, since the data is only ever read. So the byte length and
+// a checksum of the bytes actually uploaded are part of the key too. Both are computed once per
+// engine construction, against an upload of the same buffer, so they cost nothing measurable.
+//
 // Entries are held weakly, so the tables are released once the last engine using them is gone.
 
 namespace jfjoch_cuda_shared_tables {
+    // (device, source address, byte length, checksum of the bytes)
+    using TableKey = std::tuple<int, const void *, size_t, uint64_t>;
+
     struct Registry {
         std::mutex m;
-        std::map<std::pair<int, const void *>, std::weak_ptr<void>> tables;
+        std::map<TableKey, std::weak_ptr<void>> tables;
     };
 
+    // FNV-1a. Not a cryptographic hash and does not need to be - it exists to notice that the bytes
+    // behind a reused address changed, not to resist anyone.
+    inline uint64_t checksum(const void *data, size_t bytes) {
+        const auto *p = static_cast<const unsigned char *>(data);
+        uint64_t h = 1469598103934665603ULL;
+        for (size_t i = 0; i < bytes; i++) {
+            h ^= p[i];
+            h *= 1099511628211ULL;
+        }
+        return h;
+    }
+
     inline Registry &registry() {
         static Registry r;
         return r;
@@ -51,13 +73,22 @@ std::shared_ptr<CudaDevicePtr<T>> SharedDeviceTable(const void *key, size_t coun
     int device = 0;
     jfjoch_cuda_shared_tables::check(cudaGetDevice(&device));
 
+    const size_t bytes = count * sizeof(T);
+    const jfjoch_cuda_shared_tables::TableKey table_key{
+        device, key, bytes, jfjoch_cuda_shared_tables::checksum(host, bytes)};
+
     auto &reg = jfjoch_cuda_shared_tables::registry();
     // The upload happens while the lock is held: another worker must not obtain the pointer before
     // its content is on the device.
     std::lock_guard lock(reg.m);
-    auto &slot = reg.tables[{device, key}];
-    if (auto cached = slot.lock())
-        return std::static_pointer_cast<CudaDevicePtr<T>>(cached);
+    if (auto it = reg.tables.find(table_key); it != reg.tables.end()) {
+        if (auto cached = it->second.lock())
+            return std::static_pointer_cast<CudaDevicePtr<T>>(cached);
+    }
+    // Drop entries whose table is gone before adding one. Without this a long session that reloads
+    // masks or remaps geometry accumulates a dead entry per distinct content, for ever.
+    for (auto it = reg.tables.begin(); it != reg.tables.end();)
+        it = it->second.expired() ? reg.tables.erase(it) : std::next(it);
 
     // Free on the device that allocated it - the last engine to drop the table may well be a worker
     // pinned to a different GPU.
@@ -69,9 +100,9 @@ std::shared_ptr<CudaDevicePtr<T>> SharedDeviceTable(const void *key, size_t coun
         cudaSetDevice(current);
     });
     jfjoch_cuda_shared_tables::check(
-        cudaMemcpyAsync(table->get(), host, count * sizeof(T), cudaMemcpyHostToDevice, stream));
+        cudaMemcpyAsync(table->get(), host, bytes, cudaMemcpyHostToDevice, stream));
     jfjoch_cuda_shared_tables::check(cudaStreamSynchronize(stream));
 
-    slot = std::shared_ptr<void>(table);
+    reg.tables[table_key] = std::shared_ptr<void>(table);
     return table;
 }