// ============================================================================
// TILE GRID GPU GATE — does the paged sampler address an INCOMPLETE world?
//
// THE CLAIM UNDER TEST. custom_data becomes [ INDEX TABLE | per-tile rects ],
// and the emitted GLSL resolves a world position to the tile that owns it and
// then to a texel inside that tile -- returning the STUB where no tile is
// loaded. Contents of custom_data change and the emitted GLSL changes; the
// BINDING SET does not.
//
// WHY THIS INSTRUMENT AND NOT AN EXISTING ONE.
//
// gate_dc_vs_lookdev CANNOT SEE THIS. It walks DCTerrain::sample_density,
// i.e. the CPU path through ProcNodesBridge::evaluate, and this change is
// entirely in the emitted GLSL. It would report identical numbers whether
// the addressing were right or catastrophically wrong -- an instrument that
// cannot fail.
//
// map_dc_diffusion.gd's GPU oracle CAN see it, but it cannot run here: it
// SUBMITS density fills from _process on the live device, and letting it run
// long enough to answer is what interleaved its uniform set with the mesh
// pipeline's and emptied the world on 2026-09-03. Its 20 s deadline is a
// guard, not a defect.
//
// So the subject is compiled and dispatched STANDALONE, on its own VkDevice,
// in its own process, with nothing else on the queue -- the pattern
// test_erosion_parity.cpp established.
//
// ONE COPY, NOT A TRANSCRIPTION. The packing and the emitted sampler both come
// from proc-nodes/src/shared/erosion_wrapper_glsl.h, the same text the engine
// emits. The parity oracle's header names transcription as its own soft spot
// ("wrappers: transcribed ... Hoisting _recompile_glsl's string assembly into
// a godot-free function would close that gap. Not done."). Transcribing the
// packing would put this gate's weakest point exactly on the change it exists
// to check.
//
// THE REFERENCE IS TRUSTED CODE, NOT A SECOND OPINION. Probe A compares the
// paged sampler against the EXISTING single-tile sampler reading that one tile
// alone -- unchanged, shipping, and byte-identical to what it emitted before
// (verified: 2828 bytes, cksum 1189488134). A differential against known-good
// code catches an addressing error that a transcribed twin would reproduce.
//
// CONTROLS, each of which must FAIL for the gate to mean anything:
// B' the single-tile reference at the hole returns real terrain (it can only
// clamp), so "stub" is not trivially equal to "what a clamp would give".
// A' corrupt one index-table entry -> Probe A must FAIL. Proves the shader
// reads the table rather than passing for some other reason.
// C swap tx/ty when packing -> Probe A must FAIL. The axis swap mirrors the
// world and still looks plausible, so it needs its own control.
//
// ⚠ SCOPE, STATED SO IT IS NOT MISTAKEN FOR COVERAGE. This gate checks LEVEL 0
// addressing. It does NOT cover the mip pyramid (per-tile pyramids are the next
// step and want their own probes), and it does NOT cover points outside the
// grid bounding box: hblk_sample_level clamps world->texel to HBLK_SIDE[0]
// BEFORE the fetch is reached, so a far-outside point resolves to the bbox edge
// rather than to the stub. That divergence from the CPU twin is real and is
// reported below, not asserted away.
//
// Build (no engine, no godot-cpp, no scons):
// g++ -O2 -std=c++17 -I ../proc-nodes/src/shared \
// -o /tmp/tile_grid_gpu tests/test_tile_grid_gpu.cpp -lvulkan
//
// Run:
// /tmp/tile_grid_gpu [--tiles DIR]
// default DIR = ~/.local/share/godot/app_userdata/HyperTPS/tile_cache/1234_hblk
//
// Exit 0 when every probe and every control behaves.
// ============================================================================
#include <vulkan/vulkan.h>
#include "erosion_wrapper_glsl.h"
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cmath>
#include <string>
#include <vector>
#include <map>
#include <algorithm>
#include <dirent.h>
// The stub is deliberately far outside any real elevation (the whole cache
// spans -4829..5110 m) so "is this the stub?" is never a judgement call.
static const float STUB = -99999.0f;
static int g_fails = 0;
static void ok(const char* name, bool pass, const char* detail = "") {
printf(" %-4s %-52s %s\n", pass ? "ok" : "FAIL", name, detail);
if (!pass) g_fails++;
}
// ---------------------------------------------------------------------------
// .hblk, loaded godot-free. Format read from HeightfieldNode::_load_hblk:
// 32-byte header 'HBLK' + ver(u32) + n(u32) + mpp(f32) + ox(f32) + oz(f32),
// then n*n int16 elevations in metres, row = z.
// ---------------------------------------------------------------------------
static bool load_hblk(const std::string& path, HeightTileSrc& out) {
FILE* f = fopen(path.c_str(), "rb");
if (!f) return false;
unsigned char hdr[32];
if (fread(hdr, 1, 32, f) != 32 || memcmp(hdr, "HBLK", 4) != 0) { fclose(f); return false; }
uint32_t ver, n;
memcpy(&ver, hdr + 4, 4); memcpy(&n, hdr + 8, 4);
memcpy(&out.mpp, hdr + 12, 4);
float ox, oz;
memcpy(&ox, hdr + 16, 4); memcpy(&oz, hdr + 20, 4);
if (ver != 1u || n == 0u) { fclose(f); return false; }
std::vector<int16_t> raw((size_t)n * n);
if (fread(raw.data(), 2, raw.size(), f) != raw.size()) { fclose(f); return false; }
fclose(f);
out.n = (int)n;
out.apron = 1;
const int payload = out.n - 2 * out.apron;
// Invert stitch_block.py exactly as add_tile does: ox = (ty*payload - apron)*mpp.
out.ty = (int)llround((double)ox / out.mpp + 1.0) / payload;
out.tx = (int)llround((double)oz / out.mpp + 1.0) / payload;
out.h.resize(raw.size());
for (size_t i = 0; i < raw.size(); i++) out.h[i] = (float)raw[i];
return true;
}
// ---------------------------------------------------------------------------
// Vulkan plumbing. A third copy, for the reason test_erosion_parity.cpp gives
// for its second: plumbing drift breaks the instrument loudly, not the terrain
// silently.
// ---------------------------------------------------------------------------
static void die(const char* what, VkResult r) {
fprintf(stderr, "FATAL: %s -> VkResult %d\n", what, (int)r);
exit(1);
}
#define VK(x) do { VkResult _r = (x); if (_r != VK_SUCCESS) die(#x, _r); } while (0)
struct Ctx {
VkInstance instance = VK_NULL_HANDLE;
VkPhysicalDevice phys = VK_NULL_HANDLE;
VkDevice dev = VK_NULL_HANDLE;
VkQueue queue = VK_NULL_HANDLE;
uint32_t qfam = 0;
VkCommandPool cmdpool = VK_NULL_HANDLE;
VkPhysicalDeviceMemoryProperties memprops{};
char devname[256] = {0};
};
static uint32_t find_mem(Ctx& c, uint32_t bits, VkMemoryPropertyFlags want) {
for (uint32_t i = 0; i < c.memprops.memoryTypeCount; i++)
if ((bits & (1u << i)) && (c.memprops.memoryTypes[i].propertyFlags & want) == want)
return i;
fprintf(stderr, "FATAL: no memory type with flags 0x%x\n", want);
exit(1);
}
struct Buf { VkBuffer buf = VK_NULL_HANDLE; VkDeviceMemory mem = VK_NULL_HANDLE;
void* mapped = nullptr; VkDeviceSize size = 0; };
static Buf make_host_buf(Ctx& c, VkDeviceSize size) {
if (size == 0) size = 16;
Buf b; b.size = size;
VkBufferCreateInfo bi{VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO};
bi.size = size;
bi.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;
bi.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
VK(vkCreateBuffer(c.dev, &bi, nullptr, &b.buf));
VkMemoryRequirements mr;
vkGetBufferMemoryRequirements(c.dev, b.buf, &mr);
VkMemoryAllocateInfo ai{VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO};
ai.allocationSize = mr.size;
ai.memoryTypeIndex = find_mem(c, mr.memoryTypeBits,
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT);
VK(vkAllocateMemory(c.dev, &ai, nullptr, &b.mem));
VK(vkBindBufferMemory(c.dev, b.buf, b.mem, 0));
VK(vkMapMemory(c.dev, b.mem, 0, VK_WHOLE_SIZE, 0, &b.mapped));
return b;
}
static void destroy_buf(Ctx& c, Buf& b) {
if (b.mem) { vkUnmapMemory(c.dev, b.mem); vkFreeMemory(c.dev, b.mem, nullptr); }
if (b.buf) vkDestroyBuffer(c.dev, b.buf, nullptr);
b = Buf{};
}
static void init_vulkan(Ctx& c) {
VkApplicationInfo app{VK_STRUCTURE_TYPE_APPLICATION_INFO};
app.pApplicationName = "dc-terrain tile grid gate";
app.apiVersion = VK_API_VERSION_1_1;
VkInstanceCreateInfo ii{VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO};
ii.pApplicationInfo = &app;
VK(vkCreateInstance(&ii, nullptr, &c.instance));
uint32_t n = 0;
VK(vkEnumeratePhysicalDevices(c.instance, &n, nullptr));
if (!n) { fprintf(stderr, "FATAL: no Vulkan device\n"); exit(1); }
std::vector<VkPhysicalDevice> devs(n);
VK(vkEnumeratePhysicalDevices(c.instance, &n, devs.data()));
c.phys = devs[0];
for (auto d : devs) {
VkPhysicalDeviceProperties p;
vkGetPhysicalDeviceProperties(d, &p);
if (p.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU) { c.phys = d; break; }
}
VkPhysicalDeviceProperties props;
vkGetPhysicalDeviceProperties(c.phys, &props);
snprintf(c.devname, sizeof(c.devname), "%s", props.deviceName);
vkGetPhysicalDeviceMemoryProperties(c.phys, &c.memprops);
uint32_t qn = 0;
vkGetPhysicalDeviceQueueFamilyProperties(c.phys, &qn, nullptr);
std::vector<VkQueueFamilyProperties> qs(qn);
vkGetPhysicalDeviceQueueFamilyProperties(c.phys, &qn, qs.data());
bool found = false;
for (uint32_t i = 0; i < qn; i++)
if (qs[i].queueFlags & VK_QUEUE_COMPUTE_BIT) { c.qfam = i; found = true; break; }
if (!found) { fprintf(stderr, "FATAL: no compute queue\n"); exit(1); }
float prio = 1.0f;
VkDeviceQueueCreateInfo qi{VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO};
qi.queueFamilyIndex = c.qfam; qi.queueCount = 1; qi.pQueuePriorities = &prio;
VkDeviceCreateInfo di{VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO};
di.queueCreateInfoCount = 1; di.pQueueCreateInfos = &qi;
VK(vkCreateDevice(c.phys, &di, nullptr, &c.dev));
vkGetDeviceQueue(c.dev, c.qfam, 0, &c.queue);
VkCommandPoolCreateInfo ci{VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO};
ci.queueFamilyIndex = c.qfam;
ci.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
VK(vkCreateCommandPool(c.dev, &ci, nullptr, &c.cmdpool));
}
static std::vector<uint32_t> compile_glsl(const std::string& src, const char* tag) {
char in_path[] = "/tmp/dc_tile_grid_XXXXXX";
int fd = mkstemp(in_path);
if (fd < 0) { perror("mkstemp"); exit(1); }
FILE* f = fdopen(fd, "w");
fwrite(src.data(), 1, src.size(), f);
fclose(f);
std::string spv_path = std::string(in_path) + ".spv";
std::string cmd = "glslc -fshader-stage=comp -O -o " + spv_path + " " + in_path;
if (system(cmd.c_str()) != 0) {
fprintf(stderr, "FATAL: glslc failed for %s (source left at %s)\n", tag, in_path);
exit(1);
}
FILE* sf = fopen(spv_path.c_str(), "rb");
if (!sf) { fprintf(stderr, "FATAL: no SPIR-V at %s\n", spv_path.c_str()); exit(1); }
fseek(sf, 0, SEEK_END); long n = ftell(sf); fseek(sf, 0, SEEK_SET);
std::vector<uint32_t> spv(n / 4);
if (fread(spv.data(), 1, n, sf) != (size_t)n) { fprintf(stderr, "FATAL: short read\n"); exit(1); }
fclose(sf);
remove(in_path); remove(spv_path.c_str());
return spv;
}
// Dispatch one sampler over probe points. `consts` supplies the storage
// description (HBLK_* and, for the paged variant, HBLK_G_*); `paged` picks the
// texel fetch. Everything else -- the Catmull-Rom and its gradient -- is the
// same shared text either way.
static std::vector<float> run_sampler(Ctx& c, const std::string& consts, bool paged,
const std::vector<float>& data, const std::vector<float>& pts_xz, const char* tag,
const char* expr = "hblk_sample_level(0, pts[i].x, pts[i].y).x") {
std::string src =
"#version 450\n"
"layout(local_size_x = 64) in;\n"
"layout(set=0, binding=0) readonly buffer Hb { float custom_data[]; };\n"
"layout(set=0, binding=1) readonly buffer InP { vec4 pts[]; };\n"
"layout(set=0, binding=2) writeonly buffer OutH { float hs[]; };\n"
"layout(push_constant) uniform PC { uint count; } pc;\n"
// hblk_sample references voxel_scale; this gate probes level 0 directly,
// so it is declared to compile and deliberately not used.
"const float voxel_scale = 1.0;\n";
src += consts;
src += erosion_heightmap_sampler_glsl(paged);
if (getenv("TG_SRC") && !strcmp(getenv("TG_SRC"), tag)) fputs(src.c_str(), stderr);
src +=
"void main() {\n"
" uint i = gl_GlobalInvocationID.x;\n"
" if (i >= pc.count) return;\n"
" hs[i] = " + std::string(expr) + ";\n"
"}\n";
std::vector<uint32_t> spv = compile_glsl(src, tag);
const uint32_t n = (uint32_t)(pts_xz.size() / 4);
Buf hb = make_host_buf(c, data.size() * sizeof(float));
Buf in = make_host_buf(c, pts_xz.size() * sizeof(float));
Buf out = make_host_buf(c, n * sizeof(float));
memcpy(hb.mapped, data.data(), data.size() * sizeof(float));
memcpy(in.mapped, pts_xz.data(), pts_xz.size() * sizeof(float));
memset(out.mapped, 0, n * sizeof(float));
VkShaderModuleCreateInfo smi{VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO};
smi.codeSize = spv.size() * 4; smi.pCode = spv.data();
VkShaderModule mod;
VK(vkCreateShaderModule(c.dev, &smi, nullptr, &mod));
VkDescriptorSetLayoutBinding binds[3]{};
for (int i = 0; i < 3; i++) {
binds[i].binding = i;
binds[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
binds[i].descriptorCount = 1;
binds[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
}
VkDescriptorSetLayoutCreateInfo dli{VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO};
dli.bindingCount = 3; dli.pBindings = binds;
VkDescriptorSetLayout dsl;
VK(vkCreateDescriptorSetLayout(c.dev, &dli, nullptr, &dsl));
VkPushConstantRange pcr{VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t)};
VkPipelineLayoutCreateInfo pli{VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO};
pli.setLayoutCount = 1; pli.pSetLayouts = &dsl;
pli.pushConstantRangeCount = 1; pli.pPushConstantRanges = &pcr;
VkPipelineLayout playout;
VK(vkCreatePipelineLayout(c.dev, &pli, nullptr, &playout));
VkComputePipelineCreateInfo cpi{VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO};
cpi.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
cpi.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;
cpi.stage.module = mod; cpi.stage.pName = "main";
cpi.layout = playout;
VkPipeline pipe;
VK(vkCreateComputePipelines(c.dev, VK_NULL_HANDLE, 1, &cpi, nullptr, &pipe));
VkDescriptorPoolSize ps{VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 3};
VkDescriptorPoolCreateInfo dpi{VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO};
dpi.maxSets = 1; dpi.poolSizeCount = 1; dpi.pPoolSizes = &ps;
VkDescriptorPool dpool;
VK(vkCreateDescriptorPool(c.dev, &dpi, nullptr, &dpool));
VkDescriptorSetAllocateInfo dai{VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO};
dai.descriptorPool = dpool; dai.descriptorSetCount = 1; dai.pSetLayouts = &dsl;
VkDescriptorSet dset;
VK(vkAllocateDescriptorSets(c.dev, &dai, &dset));
VkDescriptorBufferInfo bi[3] = {{hb.buf, 0, VK_WHOLE_SIZE},
{in.buf, 0, VK_WHOLE_SIZE},
{out.buf, 0, VK_WHOLE_SIZE}};
VkWriteDescriptorSet w[3]{};
for (int i = 0; i < 3; i++) {
w[i].sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;
w[i].dstSet = dset; w[i].dstBinding = i; w[i].descriptorCount = 1;
w[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
w[i].pBufferInfo = &bi[i];
}
vkUpdateDescriptorSets(c.dev, 3, w, 0, nullptr);
VkCommandBufferAllocateInfo cbi{VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO};
cbi.commandPool = c.cmdpool; cbi.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
cbi.commandBufferCount = 1;
VkCommandBuffer cmd;
VK(vkAllocateCommandBuffers(c.dev, &cbi, &cmd));
VkCommandBufferBeginInfo cbb{VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO};
cbb.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
VK(vkBeginCommandBuffer(cmd, &cbb));
vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipe);
vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, playout, 0, 1, &dset, 0, nullptr);
vkCmdPushConstants(cmd, playout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &n);
vkCmdDispatch(cmd, (n + 63) / 64, 1, 1);
VK(vkEndCommandBuffer(cmd));
VkSubmitInfo si{VK_STRUCTURE_TYPE_SUBMIT_INFO};
si.commandBufferCount = 1; si.pCommandBuffers = &cmd;
VK(vkQueueSubmit(c.queue, 1, &si, VK_NULL_HANDLE));
VK(vkQueueWaitIdle(c.queue));
std::vector<float> res(n);
memcpy(res.data(), out.mapped, n * sizeof(float));
vkFreeCommandBuffers(c.dev, c.cmdpool, 1, &cmd);
vkDestroyDescriptorPool(c.dev, dpool, nullptr);
vkDestroyPipeline(c.dev, pipe, nullptr);
vkDestroyPipelineLayout(c.dev, playout, nullptr);
vkDestroyDescriptorSetLayout(c.dev, dsl, nullptr);
vkDestroyShaderModule(c.dev, mod, nullptr);
destroy_buf(c, hb); destroy_buf(c, in); destroy_buf(c, out);
return res;
}
// The constants describing ONE tile, addressed the way the shipping sampler
// addresses it. This is the trusted reference.
static std::string single_tile_consts(const HeightTileSrc& t) {
char b[512];
const float ox = (float)((double)t.ty * (t.n - 2 * t.apron) - t.apron) * t.mpp;
const float oz = (float)((double)t.tx * (t.n - 2 * t.apron) - t.apron) * t.mpp;
snprintf(b, sizeof(b),
"const int HBLK_N = %d;\n"
"const float HBLK_OX = %.4f;\n"
"const float HBLK_OZ = %.4f;\n"
"const float HBLK_MPP = %.6f;\n"
"const int HBLK_LEVELS = 1;\n"
"const int HBLK_OFF[1] = int[1](0);\n"
"const int HBLK_SIDE[1] = int[1](%d);\n",
t.n, ox, oz, t.mpp, t.n);
return std::string(b);
}
// The constants describing the GRID. HBLK_OX/OZ move to grid texel (0,0), so
// hblk_sample_level's world->texel mapping is untouched and i,j arrive at the
// fetch as GRID-relative texels.
static std::string grid_consts(const HeightGridPack& p) {
const HeightGridDesc &g = p.desc;
char b[1024];
const float ox = g.ox();
const float oz = g.oz();
const int side = g.side();
snprintf(b, sizeof(b),
"const int HBLK_N = %d;\n"
"const float HBLK_OX = %.4f;\n"
"const float HBLK_OZ = %.4f;\n"
"const float HBLK_MPP = %.6f;\n"
"const int HBLK_LEVELS = 1;\n"
"const int HBLK_OFF[1] = int[1](0);\n"
"const int HBLK_SIDE[1] = int[1](%d);\n",
side, ox, oz, g.mpp, side);
return std::string(b) + erosion_heightmap_grid_glsl(g);
}
static float tile_centre_x(const HeightTileSrc& t) {
const int payload = t.n - 2 * t.apron;
return (float)((double)t.ty * payload + payload * 0.5) * t.mpp;
}
static float tile_centre_z(const HeightTileSrc& t) {
const int payload = t.n - 2 * t.apron;
return (float)((double)t.tx * payload + payload * 0.5) * t.mpp;
}
int main(int argc, char** argv) {
std::string dir = std::string(getenv("HOME") ? getenv("HOME") : ".") +
"/.local/share/godot/app_userdata/HyperTPS/tile_cache/1234_hblk";
for (int i = 1; i < argc; i++)
if (!strcmp(argv[i], "--tiles") && i + 1 < argc) dir = argv[++i];
printf("\n=== TILE GRID GPU GATE ===\n");
printf(" tiles from : %s\n", dir.c_str());
DIR* d = opendir(dir.c_str());
if (!d) { printf("\nFAIL no tile directory. Generate with:\n"
" python3 stitch_block.py --per-tile <dir> --limit 12\n"); return 2; }
std::vector<std::string> names;
while (struct dirent* e = readdir(d)) {
std::string nm = e->d_name;
if (nm.size() > 5 && nm.compare(nm.size() - 5, 5, ".hblk") == 0) names.push_back(nm);
}
closedir(d);
std::sort(names.begin(), names.end());
std::vector<HeightTileSrc> tiles;
for (const std::string& nm : names) {
HeightTileSrc t;
if (load_hblk(dir + "/" + nm, t)) tiles.push_back(t);
}
if (tiles.size() < 2) { printf("\nFAIL need >= 2 loadable tiles, got %zu\n", tiles.size()); return 2; }
printf(" loaded : %zu tiles, stored side %d (payload %d), %.1f m/px\n",
tiles.size(), tiles[0].n, tiles[0].n - 2 * tiles[0].apron, tiles[0].mpp);
std::map<long long, const HeightTileSrc*> by_coord;
for (const auto& t : tiles) by_coord[(long long)t.tx * 100000 + t.ty] = &t;
// A HOLE WITH LOADED NEIGHBOURS ON BOTH SIDES is what makes the stub
// distinguishable from a clamp. Found in the real set, not constructed.
const HeightTileSrc* A = nullptr;
int hole_tx = 0, hole_ty = 0;
for (const auto& t : tiles) {
if (by_coord.count((long long)t.tx * 100000 + (t.ty + 1))) continue;
if (!by_coord.count((long long)t.tx * 100000 + (t.ty + 2))) continue;
A = &t; hole_tx = t.tx; hole_ty = t.ty + 1; break;
}
if (!A) { printf("\nFAIL no absent cell with loaded neighbours on both sides\n"); return 2; }
printf(" tile A : (%d,%d) hole: (%d,%d) far side: (%d,%d)\n",
A->tx, A->ty, hole_tx, hole_ty, hole_tx, hole_ty + 1);
Ctx c;
init_vulkan(c);
printf(" device : %s\n\n", c.devname);
const int payload = A->n - 2 * A->apron;
const float mpp = A->mpp;
// PROBE POINTS. Tile A's interior, inset 8 texels so all 16 Catmull-Rom taps
// stay inside one tile -- this probe is about addressing, and letting taps
// cross a boundary would fold the neighbour question into it.
std::vector<float> pts;
const int GRID = 8, INSET = 8;
for (int jj = 0; jj < GRID; jj++)
for (int ii = 0; ii < GRID; ii++) {
double fx = INSET + (payload - 2.0 * INSET) * (ii + 0.5) / GRID;
double fz = INSET + (payload - 2.0 * INSET) * (jj + 0.5) / GRID;
pts.push_back((float)(((double)A->ty * payload + fx) * mpp));
pts.push_back((float)(((double)A->tx * payload + fz) * mpp));
pts.push_back(0.0f); pts.push_back(0.0f);
}
const int N_A = GRID * GRID;
// The hole's centre.
pts.push_back((float)(((double)hole_ty * payload + payload * 0.5) * mpp));
pts.push_back((float)(((double)hole_tx * payload + payload * 0.5) * mpp));
pts.push_back(0.0f); pts.push_back(0.0f);
const int I_HOLE = N_A;
// TEXEL CENTRES IN WORLD METRES.
//
// ⚠ A first draft asserted these are BIT-IDENTICAL in both shaders, on the
// reasoning that Catmull-Rom collapses to p1 at an integer coordinate. That
// claim is UNPROVABLE and the probe never tested it: the world->texel
// division is not exact, so fract(u) came out 0.999973 in the reference and
// 2^-11 in the grid -- neither is an integer texel. The points are kept
// because they are where the coordinate error is largest and cleanest to
// attribute; the claim they carry is now the derived bound below.
const int I_EXACT = I_HOLE + 1;
const int N_EXACT = 16;
for (int k = 0; k < N_EXACT; k++) {
const int tk = INSET + k * 8;
pts.push_back((float)(((double)A->ty * payload + tk) * mpp));
pts.push_back((float)(((double)A->tx * payload + tk) * mpp));
pts.push_back(0.0f); pts.push_back(0.0f);
}
// Tile B, on the far side of the hole -- the tile the mutability probe adds.
const HeightTileSrc* B = by_coord[(long long)hole_tx * 100000 + (hole_ty + 1)];
const int I_B = I_EXACT + N_EXACT;
pts.push_back(tile_centre_x(*B));
pts.push_back(tile_centre_z(*B));
pts.push_back(0.0f); pts.push_back(0.0f);
// THE WORLD, from map_dc_diffusion.gd: root span = 28 * 2^(OCTREE_DEPTH-1)
// * VOXEL_SIZE with OCTREE_DEPTH 16, and a tile is 256 px * 30 m.
const HeightWorld world = erosion_heightmap_world(28.0 * (1 << 15) * 1.0,
(double)payload * mpp);
HeightGridPack g = erosion_heightmap_pack_tiles(tiles, STUB, world);
ok("pack accepted the tile set", g.ok);
if (!g.ok) { printf("\nFAIL cannot continue\n"); return 1; }
printf(" world %dx%d cells from (%d,%d) = %d, table %.0f KiB\n",
g.desc.nx, g.desc.ny, g.desc.tx0, g.desc.ty0, g.desc.nx * g.desc.ny, g.desc.nx * g.desc.ny * 4.0 / 1024.0);
printf(" %d tiles packed, %d refused, %.1f MiB total\n\n",
g.tiles, g.rejected, g.data.size() * 4.0 / 1048576.0);
std::vector<float> ref = run_sampler(c, single_tile_consts(*A), false, A->h, pts, "reference");
std::vector<float> got = run_sampler(c, grid_consts(g), true, g.data, pts, "paged");
printf("1. THE PAGED SAMPLER AGREES WITH THE SINGLE-TILE SAMPLER\n");
double worst = 0.0;
for (int i = 0; i < N_A; i++) worst = std::max(worst, (double)fabsf(got[i] - ref[i]));
char det[128];
// REPORTED HERE, ASSERTED AGAINST THE DERIVED BOUND BELOW.
//
// ⚠ This line used to carry `worst < 0.01`, a number picked when the probe
// happened to run on a 12-tile set. On the full 841 it reads 0.012085 --
// because tile A is now a different tile with a steeper local slope (25.0
// against 19.0 m/texel) and the residual is ULP(u) x slope. The magic
// constant was wrong, not the sampler, and a hardcoded tolerance that
// tracks the DATA rather than the mechanism is the thing the bound below
// exists to replace. Removed rather than raised.
snprintf(det, sizeof(det), "max |paged-ref| = %.6f m over %d pts", worst, N_A);
ok("tile A interior matches the trusted reference", true, det);
snprintf(det, sizeof(det), "%.2f m", ref[0]);
ok(" and the reference is real terrain, not a flat buffer",
ref[0] != STUB && fabsf(ref[0]) < 9000.0f, det);
// THE RESIDUAL IS DERIVED, NOT TUNED.
//
// A per-tile origin makes (x - HBLK_OX) a small number: tile A's texels run
// 0..257, so u is exact to well under a ULP. The GRID origin does not
// cancel the large part -- the same point is texel 7176 of the grid -- and
// the world->texel division is NOT exact there. MEASURED: fract(u) = 2^-11
// = 0.000488281 at a point that is an exact texel centre in world metres,
// which is precisely one ULP of 7176.
//
// One ULP of u, times the local slope in metres per texel, is therefore the
// floor on paged-vs-reference agreement. It is not a defect and it is not a
// tolerance chosen to make the test pass: it is a cost of moving the origin
// to the grid corner, it GROWS with world extent (the full 119x119 world
// reaches texel 30,464, so ~2x this), and a regression that breaks
// ADDRESSING misses by metres, not by ULPs -- controls A' and C below miss
// by 405 m and 102,664 m.
double max_slope = 0.0; // metres per texel, over tile A
for (int j = 1; j < A->n; j++)
for (int i = 1; i < A->n; i++) {
const double h = A->h[(size_t)j * A->n + i];
max_slope = std::max(max_slope, fabs(h - A->h[(size_t)j * A->n + i - 1]));
max_slope = std::max(max_slope, fabs(h - A->h[(size_t)(j - 1) * A->n + i]));
}
const float u_max = (float)((double)std::max(g.desc.nx, g.desc.ny) * g.desc.payload);
const double ulp_u = (double)(nextafterf(u_max, 2.0f * u_max) - u_max);
const double bound = 4.0 * ulp_u * max_slope;
// Over BOTH probe sets: the fractional interior AND the world-metre texel
// centres, which is where the coordinate error is largest.
double worst_all = worst;
for (int i = I_EXACT; i < I_EXACT + N_EXACT; i++)
worst_all = std::max(worst_all, (double)fabsf(got[i] - ref[i]));
snprintf(det, sizeof(det), "%.6f m <= %.6f = 4 x ULP(%.0f) x %.1f m/texel",
worst_all, bound, u_max, max_slope);
ok(" and the residual is within the coordinate-precision bound",
worst_all <= bound, det);
printf("\n2. AN ABSENT TILE READS AS THE STUB\n");
snprintf(det, sizeof(det), "%.1f", got[I_HOLE]);
ok("hole samples exactly the stub", got[I_HOLE] == STUB, det);
printf("\n CONTROL: the stub is not what a clamp would have returned\n");
snprintf(det, sizeof(det), "reference clamps to %.2f m", ref[I_HOLE]);
ok("single-tile sampler returns terrain at the hole", ref[I_HOLE] != STUB, det);
ok(" so stub != clamp -- the probe above is not vacuous",
got[I_HOLE] != ref[I_HOLE]);
printf("\n CONTROL A': corrupt one index-table entry -> probe 1 must FAIL\n");
{
HeightGridPack bad = g;
const size_t cell = (size_t)(A->tx - g.desc.tx0) * (size_t)g.desc.ny + (size_t)(A->ty - g.desc.ty0);
const int was = (int)bad.data[cell];
bad.data[cell] = (float)((was + 1) % g.tiles);
std::vector<float> bg = run_sampler(c, grid_consts(bad), true, bad.data, pts, "corrupt");
double w = 0.0;
for (int i = 0; i < N_A; i++) w = std::max(w, (double)fabsf(bg[i] - ref[i]));
snprintf(det, sizeof(det), "max |paged-ref| = %.2f m", w);
ok("a wrong table entry is DETECTED", w > 1.0, det);
}
printf("\n CONTROL C: swap tx/ty when packing -> probe 1 must FAIL\n");
{
HeightGridPack sw = erosion_heightmap_pack_tiles(tiles, STUB, world, true);
if (!sw.ok) { ok("swapped pack built", false); }
else {
std::vector<float> sg = run_sampler(c, grid_consts(sw), true, sw.data, pts, "swapped");
double w = 0.0;
for (int i = 0; i < N_A; i++) w = std::max(w, (double)fabsf(sg[i] - ref[i]));
snprintf(det, sizeof(det), "max |paged-ref| = %.2f m", w);
ok("a mirrored world is DETECTED", w > 1.0, det);
}
}
printf("\n3. A TILE ARRIVES AT RUNTIME -- and the SHADER DOES NOT CHANGE\n");
{
// The same set minus B, which is what the world looks like before B is
// fetched from a peer.
std::vector<HeightTileSrc> without;
for (const auto& t : tiles)
if (!(t.tx == B->tx && t.ty == B->ty)) without.push_back(t);
HeightGridPack g0 = erosion_heightmap_pack_tiles(without, STUB, world);
const std::string c0 = grid_consts(g0);
const std::string c1 = grid_consts(g);
// ★ THE ASSERTION THIS WHOLE STEP EXISTS FOR. If the emitted constants
// differ, adding a tile means recompiling the density shader, and a
// world that recompiles on every fetched tile is not a mutable tileset.
snprintf(det, sizeof(det), "%zu bytes, %d vs %d tiles packed",
c0.size(), g0.tiles, g.tiles);
ok("emitted constants are IDENTICAL with and without B", c0 == c1, det);
std::vector<float> before = run_sampler(c, c0, true, g0.data, pts, "pre-add");
snprintf(det, sizeof(det), "%.1f", before[I_B]);
ok("B reads the stub before it arrives", before[I_B] == STUB, det);
snprintf(det, sizeof(det), "%.2f m", got[I_B]);
ok("B reads real terrain after it arrives", got[I_B] != STUB &&
fabsf(got[I_B]) < 9000.0f, det);
// Removing B shifts every later tile's index in the packed buffer, so
// this also asserts the TABLE absorbed the reshuffle rather than A
// silently following a stale index.
ok(" and A is untouched by the addition", before[0] == got[0]);
}
printf("\n CONTROL: a tile outside the world window is REFUSED, not wrapped\n");
{
std::vector<HeightTileSrc> stray = tiles;
stray[0].tx = world.tx0 + world.nx + 5;
HeightGridPack gs = erosion_heightmap_pack_tiles(stray, STUB, world);
// ⚠ ASSERT THE DELTA, NOT AN ABSOLUTE. A first version required exactly
// one rejection, which held only while the cache had none of its own.
// The real 841-tile set already contains two -- (77,55) and (77,56), at
// tx 77 = 591 km, outside the octree's 458.75 km half-span -- so the
// absolute failed on real data while the mechanism was correct.
snprintf(det, sizeof(det), "%d packed / %d refused, baseline %d / %d",
gs.tiles, gs.rejected, g.tiles, g.rejected);
ok("one MORE tile is dropped than the baseline",
gs.rejected == g.rejected + 1 && gs.tiles == g.tiles - 1, det);
}
printf("\n4. THE C++ TWIN OF THE FETCH AGREES WITH THE GLSL, EXACTLY\n");
{
// The engine meshes on the GPU but answers sample_density from C++, and
// test_erosion_parity.cpp exists because those two drifted once and put
// 4 km of air between what you saw and what you stood on. Compare the
// TEXEL FETCH itself at integer indices -- both sides are pure lookups
// there, no filtering, so anything but bit-identical is a real defect.
std::vector<float> ij;
std::vector<std::pair<int,int>> want;
const int gi0 = (A->ty - g.desc.ty0) * payload; // inside tile A
const int gj0 = (A->tx - g.desc.tx0) * payload;
for (int k = 0; k < 24; k++) {
int i, j;
if (k < 8) { i = gi0 + 3 + k * 31; j = gj0 + 5 + k * 29; }
else if (k < 16) { i = gi0 + payload + 7 * (k - 8); j = gj0 + 11; } // the HOLE
else { i = -900000 + k; j = gj0 + 13; } // outside
want.push_back({i, j});
ij.push_back((float)i); ij.push_back((float)j);
ij.push_back(0.0f); ij.push_back(0.0f);
}
std::vector<float> gpu = run_sampler(c, grid_consts(g), true, g.data, ij,
"texel", "hblk_texel_l(int(pts[i].x), int(pts[i].y), 0, 0)");
int bad = 0, stubs = 0, reals = 0;
for (size_t k = 0; k < want.size(); k++) {
const float cpu = erosion_heightmap_grid_texel(g.data.data(), g.desc,
want[k].first, want[k].second);
if (cpu != gpu[k]) bad++;
if (cpu == STUB) stubs++; else reals++;
}
snprintf(det, sizeof(det), "%d mismatches over %zu fetches", bad, want.size());
ok("GLSL fetch == C++ fetch, bit for bit", bad == 0, det);
// An instrument that only ever saw stubs, or only ever real texels,
// would agree for the wrong reason.
snprintf(det, sizeof(det), "%d real, %d stub", reals, stubs);
ok(" and the probe covered BOTH outcomes", reals > 0 && stubs > 0, det);
}
// REPORTED, NOT ASSERTED, AND NOW LATENT RATHER THAN LIVE.
//
// hblk_sample_level still clamps world->texel to HBLK_SIDE[0] BEFORE the
// fetch runs, so a point outside the world resolves to the edge cell rather
// than missing. With a footprint-sized table that edge cell was loaded and
// the probe returned -3871 m of somebody else's terrain. With the
// WORLD-sized table it is one of 14,161 cells of which 12 are loaded, so
// the clamp almost always lands on an absent cell and the stub comes back
// anyway -- masking, not fixing.
//
// It becomes live again the moment tiles reach the world edge. Left
// reported rather than asserted because this cache cannot demonstrate it:
// no tile sits at tx/ty 59.
printf("\n5. KNOWN GAP (reported, not asserted -- currently MASKED)\n");
{
// FAR ON X ONLY, and on tile A's own Z row. Going far on BOTH axes is
// not a test: the second axis's bounds check returns the stub and the
// probe passes for the wrong reason -- which is what a first draft of
// this probe did.
std::vector<float> far_pts;
far_pts.push_back((float)(((double)(g.desc.ty0 + g.desc.ny + 500) * payload) * mpp));
far_pts.push_back(tile_centre_z(*A));
far_pts.push_back(0.0f); far_pts.push_back(0.0f);
std::vector<float> fg = run_sampler(c, grid_consts(g), true, g.data, far_pts, "far");
printf(" far outside the world returns %.2f m (stub is %.1f)\n", fg[0], STUB);
printf(" The clamp still runs before the fetch; it lands on an ABSENT\n");
printf(" cell here, so the stub masks it. Live again once tiles reach\n");
printf(" the world edge. Out of scope for level-0 addressing.\n");
}
printf("\n%s (%d failure%s)\n\n", g_fails == 0 ? "PASS" : "FAIL", g_fails,
g_fails == 1 ? "" : "s");
return g_fails == 0 ? 0 : 1;
}