| … | … | @@ -0,0 +1,621 @@ |
| 1 | // ============================================================================ |
| 2 | // TILE GRID GPU GATE — does the paged sampler address an INCOMPLETE world? |
| 3 | // |
| 4 | // THE CLAIM UNDER TEST. custom_data becomes [ INDEX TABLE | per-tile rects ], |
| 5 | // and the emitted GLSL resolves a world position to the tile that owns it and |
| 6 | // then to a texel inside that tile -- returning the STUB where no tile is |
| 7 | // loaded. Contents of custom_data change and the emitted GLSL changes; the |
| 8 | // BINDING SET does not. |
| 9 | // |
| 10 | // WHY THIS INSTRUMENT AND NOT AN EXISTING ONE. |
| 11 | // |
| 12 | // gate_dc_vs_lookdev CANNOT SEE THIS. It walks DCTerrain::sample_density, |
| 13 | // i.e. the CPU path through ProcNodesBridge::evaluate, and this change is |
| 14 | // entirely in the emitted GLSL. It would report identical numbers whether |
| 15 | // the addressing were right or catastrophically wrong -- an instrument that |
| 16 | // cannot fail. |
| 17 | // |
| 18 | // map_dc_diffusion.gd's GPU oracle CAN see it, but it cannot run here: it |
| 19 | // SUBMITS density fills from _process on the live device, and letting it run |
| 20 | // long enough to answer is what interleaved its uniform set with the mesh |
| 21 | // pipeline's and emptied the world on 2026-09-03. Its 20 s deadline is a |
| 22 | // guard, not a defect. |
| 23 | // |
| 24 | // So the subject is compiled and dispatched STANDALONE, on its own VkDevice, |
| 25 | // in its own process, with nothing else on the queue -- the pattern |
| 26 | // test_erosion_parity.cpp established. |
| 27 | // |
| 28 | // ONE COPY, NOT A TRANSCRIPTION. The packing and the emitted sampler both come |
| 29 | // from proc-nodes/src/shared/erosion_wrapper_glsl.h, the same text the engine |
| 30 | // emits. The parity oracle's header names transcription as its own soft spot |
| 31 | // ("wrappers: transcribed ... Hoisting _recompile_glsl's string assembly into |
| 32 | // a godot-free function would close that gap. Not done."). Transcribing the |
| 33 | // packing would put this gate's weakest point exactly on the change it exists |
| 34 | // to check. |
| 35 | // |
| 36 | // THE REFERENCE IS TRUSTED CODE, NOT A SECOND OPINION. Probe A compares the |
| 37 | // paged sampler against the EXISTING single-tile sampler reading that one tile |
| 38 | // alone -- unchanged, shipping, and byte-identical to what it emitted before |
| 39 | // (verified: 2828 bytes, cksum 1189488134). A differential against known-good |
| 40 | // code catches an addressing error that a transcribed twin would reproduce. |
| 41 | // |
| 42 | // CONTROLS, each of which must FAIL for the gate to mean anything: |
| 43 | // B' the single-tile reference at the hole returns real terrain (it can only |
| 44 | // clamp), so "stub" is not trivially equal to "what a clamp would give". |
| 45 | // A' corrupt one index-table entry -> Probe A must FAIL. Proves the shader |
| 46 | // reads the table rather than passing for some other reason. |
| 47 | // C swap tx/ty when packing -> Probe A must FAIL. The axis swap mirrors the |
| 48 | // world and still looks plausible, so it needs its own control. |
| 49 | // |
| 50 | // ⚠ SCOPE, STATED SO IT IS NOT MISTAKEN FOR COVERAGE. This gate checks LEVEL 0 |
| 51 | // addressing. It does NOT cover the mip pyramid (per-tile pyramids are the next |
| 52 | // step and want their own probes), and it does NOT cover points outside the |
| 53 | // grid bounding box: hblk_sample_level clamps world->texel to HBLK_SIDE[0] |
| 54 | // BEFORE the fetch is reached, so a far-outside point resolves to the bbox edge |
| 55 | // rather than to the stub. That divergence from the CPU twin is real and is |
| 56 | // reported below, not asserted away. |
| 57 | // |
| 58 | // Build (no engine, no godot-cpp, no scons): |
| 59 | // g++ -O2 -std=c++17 -I ../proc-nodes/src/shared \ |
| 60 | // -o /tmp/tile_grid_gpu tests/test_tile_grid_gpu.cpp -lvulkan |
| 61 | // |
| 62 | // Run: |
| 63 | // /tmp/tile_grid_gpu [--tiles DIR] |
| 64 | // default DIR = ~/.local/share/godot/app_userdata/HyperTPS/tile_cache/1234_hblk |
| 65 | // |
| 66 | // Exit 0 when every probe and every control behaves. |
| 67 | // ============================================================================ |
| 68 | |
| 69 | #include <vulkan/vulkan.h> |
| 70 | #include "erosion_wrapper_glsl.h" |
| 71 | |
| 72 | #include <cstdio> |
| 73 | #include <cstdlib> |
| 74 | #include <cstring> |
| 75 | #include <cmath> |
| 76 | #include <string> |
| 77 | #include <vector> |
| 78 | #include <map> |
| 79 | #include <algorithm> |
| 80 | #include <dirent.h> |
| 81 | |
| 82 | // The stub is deliberately far outside any real elevation (the whole cache |
| 83 | // spans -4829..5110 m) so "is this the stub?" is never a judgement call. |
| 84 | static const float STUB = -99999.0f; |
| 85 | static int g_fails = 0; |
| 86 | |
| 87 | static void ok(const char* name, bool pass, const char* detail = "") { |
| 88 | printf(" %-4s %-52s %s\n", pass ? "ok" : "FAIL", name, detail); |
| 89 | if (!pass) g_fails++; |
| 90 | } |
| 91 | |
| 92 | // --------------------------------------------------------------------------- |
| 93 | // .hblk, loaded godot-free. Format read from HeightfieldNode::_load_hblk: |
| 94 | // 32-byte header 'HBLK' + ver(u32) + n(u32) + mpp(f32) + ox(f32) + oz(f32), |
| 95 | // then n*n int16 elevations in metres, row = z. |
| 96 | // --------------------------------------------------------------------------- |
| 97 | static bool load_hblk(const std::string& path, HeightTileSrc& out) { |
| 98 | FILE* f = fopen(path.c_str(), "rb"); |
| 99 | if (!f) return false; |
| 100 | unsigned char hdr[32]; |
| 101 | if (fread(hdr, 1, 32, f) != 32 || memcmp(hdr, "HBLK", 4) != 0) { fclose(f); return false; } |
| 102 | uint32_t ver, n; |
| 103 | memcpy(&ver, hdr + 4, 4); memcpy(&n, hdr + 8, 4); |
| 104 | memcpy(&out.mpp, hdr + 12, 4); |
| 105 | float ox, oz; |
| 106 | memcpy(&ox, hdr + 16, 4); memcpy(&oz, hdr + 20, 4); |
| 107 | if (ver != 1u || n == 0u) { fclose(f); return false; } |
| 108 | std::vector<int16_t> raw((size_t)n * n); |
| 109 | if (fread(raw.data(), 2, raw.size(), f) != raw.size()) { fclose(f); return false; } |
| 110 | fclose(f); |
| 111 | out.n = (int)n; |
| 112 | out.apron = 1; |
| 113 | const int payload = out.n - 2 * out.apron; |
| 114 | // Invert stitch_block.py exactly as add_tile does: ox = (ty*payload - apron)*mpp. |
| 115 | out.ty = (int)llround((double)ox / out.mpp + 1.0) / payload; |
| 116 | out.tx = (int)llround((double)oz / out.mpp + 1.0) / payload; |
| 117 | out.h.resize(raw.size()); |
| 118 | for (size_t i = 0; i < raw.size(); i++) out.h[i] = (float)raw[i]; |
| 119 | return true; |
| 120 | } |
| 121 | |
| 122 | // --------------------------------------------------------------------------- |
| 123 | // Vulkan plumbing. A third copy, for the reason test_erosion_parity.cpp gives |
| 124 | // for its second: plumbing drift breaks the instrument loudly, not the terrain |
| 125 | // silently. |
| 126 | // --------------------------------------------------------------------------- |
| 127 | static void die(const char* what, VkResult r) { |
| 128 | fprintf(stderr, "FATAL: %s -> VkResult %d\n", what, (int)r); |
| 129 | exit(1); |
| 130 | } |
| 131 | #define VK(x) do { VkResult _r = (x); if (_r != VK_SUCCESS) die(#x, _r); } while (0) |
| 132 | |
| 133 | struct Ctx { |
| 134 | VkInstance instance = VK_NULL_HANDLE; |
| 135 | VkPhysicalDevice phys = VK_NULL_HANDLE; |
| 136 | VkDevice dev = VK_NULL_HANDLE; |
| 137 | VkQueue queue = VK_NULL_HANDLE; |
| 138 | uint32_t qfam = 0; |
| 139 | VkCommandPool cmdpool = VK_NULL_HANDLE; |
| 140 | VkPhysicalDeviceMemoryProperties memprops{}; |
| 141 | char devname[256] = {0}; |
| 142 | }; |
| 143 | |
| 144 | static uint32_t find_mem(Ctx& c, uint32_t bits, VkMemoryPropertyFlags want) { |
| 145 | for (uint32_t i = 0; i < c.memprops.memoryTypeCount; i++) |
| 146 | if ((bits & (1u << i)) && (c.memprops.memoryTypes[i].propertyFlags & want) == want) |
| 147 | return i; |
| 148 | fprintf(stderr, "FATAL: no memory type with flags 0x%x\n", want); |
| 149 | exit(1); |
| 150 | } |
| 151 | |
| 152 | struct Buf { VkBuffer buf = VK_NULL_HANDLE; VkDeviceMemory mem = VK_NULL_HANDLE; |
| 153 | void* mapped = nullptr; VkDeviceSize size = 0; }; |
| 154 | |
| 155 | static Buf make_host_buf(Ctx& c, VkDeviceSize size) { |
| 156 | if (size == 0) size = 16; |
| 157 | Buf b; b.size = size; |
| 158 | VkBufferCreateInfo bi{VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; |
| 159 | bi.size = size; |
| 160 | bi.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT; |
| 161 | bi.sharingMode = VK_SHARING_MODE_EXCLUSIVE; |
| 162 | VK(vkCreateBuffer(c.dev, &bi, nullptr, &b.buf)); |
| 163 | VkMemoryRequirements mr; |
| 164 | vkGetBufferMemoryRequirements(c.dev, b.buf, &mr); |
| 165 | VkMemoryAllocateInfo ai{VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; |
| 166 | ai.allocationSize = mr.size; |
| 167 | ai.memoryTypeIndex = find_mem(c, mr.memoryTypeBits, |
| 168 | VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); |
| 169 | VK(vkAllocateMemory(c.dev, &ai, nullptr, &b.mem)); |
| 170 | VK(vkBindBufferMemory(c.dev, b.buf, b.mem, 0)); |
| 171 | VK(vkMapMemory(c.dev, b.mem, 0, VK_WHOLE_SIZE, 0, &b.mapped)); |
| 172 | return b; |
| 173 | } |
| 174 | |
| 175 | static void destroy_buf(Ctx& c, Buf& b) { |
| 176 | if (b.mem) { vkUnmapMemory(c.dev, b.mem); vkFreeMemory(c.dev, b.mem, nullptr); } |
| 177 | if (b.buf) vkDestroyBuffer(c.dev, b.buf, nullptr); |
| 178 | b = Buf{}; |
| 179 | } |
| 180 | |
| 181 | static void init_vulkan(Ctx& c) { |
| 182 | VkApplicationInfo app{VK_STRUCTURE_TYPE_APPLICATION_INFO}; |
| 183 | app.pApplicationName = "dc-terrain tile grid gate"; |
| 184 | app.apiVersion = VK_API_VERSION_1_1; |
| 185 | VkInstanceCreateInfo ii{VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO}; |
| 186 | ii.pApplicationInfo = &app; |
| 187 | VK(vkCreateInstance(&ii, nullptr, &c.instance)); |
| 188 | |
| 189 | uint32_t n = 0; |
| 190 | VK(vkEnumeratePhysicalDevices(c.instance, &n, nullptr)); |
| 191 | if (!n) { fprintf(stderr, "FATAL: no Vulkan device\n"); exit(1); } |
| 192 | std::vector<VkPhysicalDevice> devs(n); |
| 193 | VK(vkEnumeratePhysicalDevices(c.instance, &n, devs.data())); |
| 194 | c.phys = devs[0]; |
| 195 | for (auto d : devs) { |
| 196 | VkPhysicalDeviceProperties p; |
| 197 | vkGetPhysicalDeviceProperties(d, &p); |
| 198 | if (p.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU) { c.phys = d; break; } |
| 199 | } |
| 200 | VkPhysicalDeviceProperties props; |
| 201 | vkGetPhysicalDeviceProperties(c.phys, &props); |
| 202 | snprintf(c.devname, sizeof(c.devname), "%s", props.deviceName); |
| 203 | vkGetPhysicalDeviceMemoryProperties(c.phys, &c.memprops); |
| 204 | |
| 205 | uint32_t qn = 0; |
| 206 | vkGetPhysicalDeviceQueueFamilyProperties(c.phys, &qn, nullptr); |
| 207 | std::vector<VkQueueFamilyProperties> qs(qn); |
| 208 | vkGetPhysicalDeviceQueueFamilyProperties(c.phys, &qn, qs.data()); |
| 209 | bool found = false; |
| 210 | for (uint32_t i = 0; i < qn; i++) |
| 211 | if (qs[i].queueFlags & VK_QUEUE_COMPUTE_BIT) { c.qfam = i; found = true; break; } |
| 212 | if (!found) { fprintf(stderr, "FATAL: no compute queue\n"); exit(1); } |
| 213 | |
| 214 | float prio = 1.0f; |
| 215 | VkDeviceQueueCreateInfo qi{VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO}; |
| 216 | qi.queueFamilyIndex = c.qfam; qi.queueCount = 1; qi.pQueuePriorities = &prio; |
| 217 | VkDeviceCreateInfo di{VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO}; |
| 218 | di.queueCreateInfoCount = 1; di.pQueueCreateInfos = &qi; |
| 219 | VK(vkCreateDevice(c.phys, &di, nullptr, &c.dev)); |
| 220 | vkGetDeviceQueue(c.dev, c.qfam, 0, &c.queue); |
| 221 | |
| 222 | VkCommandPoolCreateInfo ci{VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; |
| 223 | ci.queueFamilyIndex = c.qfam; |
| 224 | ci.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT; |
| 225 | VK(vkCreateCommandPool(c.dev, &ci, nullptr, &c.cmdpool)); |
| 226 | } |
| 227 | |
| 228 | static std::vector<uint32_t> compile_glsl(const std::string& src, const char* tag) { |
| 229 | char in_path[] = "/tmp/dc_tile_grid_XXXXXX"; |
| 230 | int fd = mkstemp(in_path); |
| 231 | if (fd < 0) { perror("mkstemp"); exit(1); } |
| 232 | FILE* f = fdopen(fd, "w"); |
| 233 | fwrite(src.data(), 1, src.size(), f); |
| 234 | fclose(f); |
| 235 | std::string spv_path = std::string(in_path) + ".spv"; |
| 236 | std::string cmd = "glslc -fshader-stage=comp -O -o " + spv_path + " " + in_path; |
| 237 | if (system(cmd.c_str()) != 0) { |
| 238 | fprintf(stderr, "FATAL: glslc failed for %s (source left at %s)\n", tag, in_path); |
| 239 | exit(1); |
| 240 | } |
| 241 | FILE* sf = fopen(spv_path.c_str(), "rb"); |
| 242 | if (!sf) { fprintf(stderr, "FATAL: no SPIR-V at %s\n", spv_path.c_str()); exit(1); } |
| 243 | fseek(sf, 0, SEEK_END); long n = ftell(sf); fseek(sf, 0, SEEK_SET); |
| 244 | std::vector<uint32_t> spv(n / 4); |
| 245 | if (fread(spv.data(), 1, n, sf) != (size_t)n) { fprintf(stderr, "FATAL: short read\n"); exit(1); } |
| 246 | fclose(sf); |
| 247 | remove(in_path); remove(spv_path.c_str()); |
| 248 | return spv; |
| 249 | } |
| 250 | |
| 251 | // Dispatch one sampler over probe points. `consts` supplies the storage |
| 252 | // description (HBLK_* and, for the paged variant, HBLK_G_*); `paged` picks the |
| 253 | // texel fetch. Everything else -- the Catmull-Rom and its gradient -- is the |
| 254 | // same shared text either way. |
| 255 | static std::vector<float> run_sampler(Ctx& c, const std::string& consts, bool paged, |
| 256 | const std::vector<float>& data, const std::vector<float>& pts_xz, const char* tag, |
| 257 | const char* expr = "hblk_sample_level(0, pts[i].x, pts[i].y).x") { |
| 258 | std::string src = |
| 259 | "#version 450\n" |
| 260 | "layout(local_size_x = 64) in;\n" |
| 261 | "layout(set=0, binding=0) readonly buffer Hb { float custom_data[]; };\n" |
| 262 | "layout(set=0, binding=1) readonly buffer InP { vec4 pts[]; };\n" |
| 263 | "layout(set=0, binding=2) writeonly buffer OutH { float hs[]; };\n" |
| 264 | "layout(push_constant) uniform PC { uint count; } pc;\n" |
| 265 | // hblk_sample references voxel_scale; this gate probes level 0 directly, |
| 266 | // so it is declared to compile and deliberately not used. |
| 267 | "const float voxel_scale = 1.0;\n"; |
| 268 | src += consts; |
| 269 | src += erosion_heightmap_sampler_glsl(paged); |
| 270 | if (getenv("TG_SRC") && !strcmp(getenv("TG_SRC"), tag)) fputs(src.c_str(), stderr); |
| 271 | src += |
| 272 | "void main() {\n" |
| 273 | " uint i = gl_GlobalInvocationID.x;\n" |
| 274 | " if (i >= pc.count) return;\n" |
| 275 | " hs[i] = " + std::string(expr) + ";\n" |
| 276 | "}\n"; |
| 277 | |
| 278 | std::vector<uint32_t> spv = compile_glsl(src, tag); |
| 279 | const uint32_t n = (uint32_t)(pts_xz.size() / 4); |
| 280 | |
| 281 | Buf hb = make_host_buf(c, data.size() * sizeof(float)); |
| 282 | Buf in = make_host_buf(c, pts_xz.size() * sizeof(float)); |
| 283 | Buf out = make_host_buf(c, n * sizeof(float)); |
| 284 | memcpy(hb.mapped, data.data(), data.size() * sizeof(float)); |
| 285 | memcpy(in.mapped, pts_xz.data(), pts_xz.size() * sizeof(float)); |
| 286 | memset(out.mapped, 0, n * sizeof(float)); |
| 287 | |
| 288 | VkShaderModuleCreateInfo smi{VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO}; |
| 289 | smi.codeSize = spv.size() * 4; smi.pCode = spv.data(); |
| 290 | VkShaderModule mod; |
| 291 | VK(vkCreateShaderModule(c.dev, &smi, nullptr, &mod)); |
| 292 | |
| 293 | VkDescriptorSetLayoutBinding binds[3]{}; |
| 294 | for (int i = 0; i < 3; i++) { |
| 295 | binds[i].binding = i; |
| 296 | binds[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER; |
| 297 | binds[i].descriptorCount = 1; |
| 298 | binds[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT; |
| 299 | } |
| 300 | VkDescriptorSetLayoutCreateInfo dli{VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO}; |
| 301 | dli.bindingCount = 3; dli.pBindings = binds; |
| 302 | VkDescriptorSetLayout dsl; |
| 303 | VK(vkCreateDescriptorSetLayout(c.dev, &dli, nullptr, &dsl)); |
| 304 | |
| 305 | VkPushConstantRange pcr{VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t)}; |
| 306 | VkPipelineLayoutCreateInfo pli{VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO}; |
| 307 | pli.setLayoutCount = 1; pli.pSetLayouts = &dsl; |
| 308 | pli.pushConstantRangeCount = 1; pli.pPushConstantRanges = &pcr; |
| 309 | VkPipelineLayout playout; |
| 310 | VK(vkCreatePipelineLayout(c.dev, &pli, nullptr, &playout)); |
| 311 | |
| 312 | VkComputePipelineCreateInfo cpi{VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO}; |
| 313 | cpi.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO; |
| 314 | cpi.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT; |
| 315 | cpi.stage.module = mod; cpi.stage.pName = "main"; |
| 316 | cpi.layout = playout; |
| 317 | VkPipeline pipe; |
| 318 | VK(vkCreateComputePipelines(c.dev, VK_NULL_HANDLE, 1, &cpi, nullptr, &pipe)); |
| 319 | |
| 320 | VkDescriptorPoolSize ps{VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 3}; |
| 321 | VkDescriptorPoolCreateInfo dpi{VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO}; |
| 322 | dpi.maxSets = 1; dpi.poolSizeCount = 1; dpi.pPoolSizes = &ps; |
| 323 | VkDescriptorPool dpool; |
| 324 | VK(vkCreateDescriptorPool(c.dev, &dpi, nullptr, &dpool)); |
| 325 | VkDescriptorSetAllocateInfo dai{VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO}; |
| 326 | dai.descriptorPool = dpool; dai.descriptorSetCount = 1; dai.pSetLayouts = &dsl; |
| 327 | VkDescriptorSet dset; |
| 328 | VK(vkAllocateDescriptorSets(c.dev, &dai, &dset)); |
| 329 | |
| 330 | VkDescriptorBufferInfo bi[3] = {{hb.buf, 0, VK_WHOLE_SIZE}, |
| 331 | {in.buf, 0, VK_WHOLE_SIZE}, |
| 332 | {out.buf, 0, VK_WHOLE_SIZE}}; |
| 333 | VkWriteDescriptorSet w[3]{}; |
| 334 | for (int i = 0; i < 3; i++) { |
| 335 | w[i].sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET; |
| 336 | w[i].dstSet = dset; w[i].dstBinding = i; w[i].descriptorCount = 1; |
| 337 | w[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER; |
| 338 | w[i].pBufferInfo = &bi[i]; |
| 339 | } |
| 340 | vkUpdateDescriptorSets(c.dev, 3, w, 0, nullptr); |
| 341 | |
| 342 | VkCommandBufferAllocateInfo cbi{VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; |
| 343 | cbi.commandPool = c.cmdpool; cbi.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; |
| 344 | cbi.commandBufferCount = 1; |
| 345 | VkCommandBuffer cmd; |
| 346 | VK(vkAllocateCommandBuffers(c.dev, &cbi, &cmd)); |
| 347 | VkCommandBufferBeginInfo cbb{VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; |
| 348 | cbb.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; |
| 349 | VK(vkBeginCommandBuffer(cmd, &cbb)); |
| 350 | vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipe); |
| 351 | vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, playout, 0, 1, &dset, 0, nullptr); |
| 352 | vkCmdPushConstants(cmd, playout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &n); |
| 353 | vkCmdDispatch(cmd, (n + 63) / 64, 1, 1); |
| 354 | VK(vkEndCommandBuffer(cmd)); |
| 355 | VkSubmitInfo si{VK_STRUCTURE_TYPE_SUBMIT_INFO}; |
| 356 | si.commandBufferCount = 1; si.pCommandBuffers = &cmd; |
| 357 | VK(vkQueueSubmit(c.queue, 1, &si, VK_NULL_HANDLE)); |
| 358 | VK(vkQueueWaitIdle(c.queue)); |
| 359 | |
| 360 | std::vector<float> res(n); |
| 361 | memcpy(res.data(), out.mapped, n * sizeof(float)); |
| 362 | |
| 363 | vkFreeCommandBuffers(c.dev, c.cmdpool, 1, &cmd); |
| 364 | vkDestroyDescriptorPool(c.dev, dpool, nullptr); |
| 365 | vkDestroyPipeline(c.dev, pipe, nullptr); |
| 366 | vkDestroyPipelineLayout(c.dev, playout, nullptr); |
| 367 | vkDestroyDescriptorSetLayout(c.dev, dsl, nullptr); |
| 368 | vkDestroyShaderModule(c.dev, mod, nullptr); |
| 369 | destroy_buf(c, hb); destroy_buf(c, in); destroy_buf(c, out); |
| 370 | return res; |
| 371 | } |
| 372 | |
| 373 | // The constants describing ONE tile, addressed the way the shipping sampler |
| 374 | // addresses it. This is the trusted reference. |
| 375 | static std::string single_tile_consts(const HeightTileSrc& t) { |
| 376 | char b[512]; |
| 377 | const float ox = (float)((double)t.ty * (t.n - 2 * t.apron) - t.apron) * t.mpp; |
| 378 | const float oz = (float)((double)t.tx * (t.n - 2 * t.apron) - t.apron) * t.mpp; |
| 379 | snprintf(b, sizeof(b), |
| 380 | "const int HBLK_N = %d;\n" |
| 381 | "const float HBLK_OX = %.4f;\n" |
| 382 | "const float HBLK_OZ = %.4f;\n" |
| 383 | "const float HBLK_MPP = %.6f;\n" |
| 384 | "const int HBLK_LEVELS = 1;\n" |
| 385 | "const int HBLK_OFF[1] = int[1](0);\n" |
| 386 | "const int HBLK_SIDE[1] = int[1](%d);\n", |
| 387 | t.n, ox, oz, t.mpp, t.n); |
| 388 | return std::string(b); |
| 389 | } |
| 390 | |
| 391 | // The constants describing the GRID. HBLK_OX/OZ move to grid texel (0,0), so |
| 392 | // hblk_sample_level's world->texel mapping is untouched and i,j arrive at the |
| 393 | // fetch as GRID-relative texels. |
| 394 | static std::string grid_consts(const HeightGridPack& g) { |
| 395 | char b[1024]; |
| 396 | const float ox = (float)((double)g.ty0 * g.payload) * g.mpp; |
| 397 | const float oz = (float)((double)g.tx0 * g.payload) * g.mpp; |
| 398 | const int side = std::max(g.nx, g.ny) * g.payload; |
| 399 | snprintf(b, sizeof(b), |
| 400 | "const int HBLK_N = %d;\n" |
| 401 | "const float HBLK_OX = %.4f;\n" |
| 402 | "const float HBLK_OZ = %.4f;\n" |
| 403 | "const float HBLK_MPP = %.6f;\n" |
| 404 | "const int HBLK_LEVELS = 1;\n" |
| 405 | "const int HBLK_OFF[1] = int[1](0);\n" |
| 406 | "const int HBLK_SIDE[1] = int[1](%d);\n", |
| 407 | side, ox, oz, g.mpp, side); |
| 408 | return std::string(b) + erosion_heightmap_grid_glsl(g); |
| 409 | } |
| 410 | |
| 411 | static float tile_centre_x(const HeightTileSrc& t) { |
| 412 | const int payload = t.n - 2 * t.apron; |
| 413 | return (float)((double)t.ty * payload + payload * 0.5) * t.mpp; |
| 414 | } |
| 415 | static float tile_centre_z(const HeightTileSrc& t) { |
| 416 | const int payload = t.n - 2 * t.apron; |
| 417 | return (float)((double)t.tx * payload + payload * 0.5) * t.mpp; |
| 418 | } |
| 419 | |
| 420 | int main(int argc, char** argv) { |
| 421 | std::string dir = std::string(getenv("HOME") ? getenv("HOME") : ".") + |
| 422 | "/.local/share/godot/app_userdata/HyperTPS/tile_cache/1234_hblk"; |
| 423 | for (int i = 1; i < argc; i++) |
| 424 | if (!strcmp(argv[i], "--tiles") && i + 1 < argc) dir = argv[++i]; |
| 425 | |
| 426 | printf("\n=== TILE GRID GPU GATE ===\n"); |
| 427 | printf(" tiles from : %s\n", dir.c_str()); |
| 428 | |
| 429 | DIR* d = opendir(dir.c_str()); |
| 430 | if (!d) { printf("\nFAIL no tile directory. Generate with:\n" |
| 431 | " python3 stitch_block.py --per-tile <dir> --limit 12\n"); return 2; } |
| 432 | std::vector<std::string> names; |
| 433 | while (struct dirent* e = readdir(d)) { |
| 434 | std::string nm = e->d_name; |
| 435 | if (nm.size() > 5 && nm.compare(nm.size() - 5, 5, ".hblk") == 0) names.push_back(nm); |
| 436 | } |
| 437 | closedir(d); |
| 438 | std::sort(names.begin(), names.end()); |
| 439 | |
| 440 | std::vector<HeightTileSrc> tiles; |
| 441 | for (const std::string& nm : names) { |
| 442 | HeightTileSrc t; |
| 443 | if (load_hblk(dir + "/" + nm, t)) tiles.push_back(t); |
| 444 | } |
| 445 | if (tiles.size() < 2) { printf("\nFAIL need >= 2 loadable tiles, got %zu\n", tiles.size()); return 2; } |
| 446 | printf(" loaded : %zu tiles, stored side %d (payload %d), %.1f m/px\n", |
| 447 | tiles.size(), tiles[0].n, tiles[0].n - 2 * tiles[0].apron, tiles[0].mpp); |
| 448 | |
| 449 | std::map<long long, const HeightTileSrc*> by_coord; |
| 450 | for (const auto& t : tiles) by_coord[(long long)t.tx * 100000 + t.ty] = &t; |
| 451 | |
| 452 | // A HOLE WITH LOADED NEIGHBOURS ON BOTH SIDES is what makes the stub |
| 453 | // distinguishable from a clamp. Found in the real set, not constructed. |
| 454 | const HeightTileSrc* A = nullptr; |
| 455 | int hole_tx = 0, hole_ty = 0; |
| 456 | for (const auto& t : tiles) { |
| 457 | if (by_coord.count((long long)t.tx * 100000 + (t.ty + 1))) continue; |
| 458 | if (!by_coord.count((long long)t.tx * 100000 + (t.ty + 2))) continue; |
| 459 | A = &t; hole_tx = t.tx; hole_ty = t.ty + 1; break; |
| 460 | } |
| 461 | if (!A) { printf("\nFAIL no absent cell with loaded neighbours on both sides\n"); return 2; } |
| 462 | printf(" tile A : (%d,%d) hole: (%d,%d) far side: (%d,%d)\n", |
| 463 | A->tx, A->ty, hole_tx, hole_ty, hole_tx, hole_ty + 1); |
| 464 | |
| 465 | Ctx c; |
| 466 | init_vulkan(c); |
| 467 | printf(" device : %s\n\n", c.devname); |
| 468 | |
| 469 | const int payload = A->n - 2 * A->apron; |
| 470 | const float mpp = A->mpp; |
| 471 | |
| 472 | // PROBE POINTS. Tile A's interior, inset 8 texels so all 16 Catmull-Rom taps |
| 473 | // stay inside one tile -- this probe is about addressing, and letting taps |
| 474 | // cross a boundary would fold the neighbour question into it. |
| 475 | std::vector<float> pts; |
| 476 | const int GRID = 8, INSET = 8; |
| 477 | for (int jj = 0; jj < GRID; jj++) |
| 478 | for (int ii = 0; ii < GRID; ii++) { |
| 479 | double fx = INSET + (payload - 2.0 * INSET) * (ii + 0.5) / GRID; |
| 480 | double fz = INSET + (payload - 2.0 * INSET) * (jj + 0.5) / GRID; |
| 481 | pts.push_back((float)(((double)A->ty * payload + fx) * mpp)); |
| 482 | pts.push_back((float)(((double)A->tx * payload + fz) * mpp)); |
| 483 | pts.push_back(0.0f); pts.push_back(0.0f); |
| 484 | } |
| 485 | const int N_A = GRID * GRID; |
| 486 | // The hole's centre. |
| 487 | pts.push_back((float)(((double)hole_ty * payload + payload * 0.5) * mpp)); |
| 488 | pts.push_back((float)(((double)hole_tx * payload + payload * 0.5) * mpp)); |
| 489 | pts.push_back(0.0f); pts.push_back(0.0f); |
| 490 | const int I_HOLE = N_A; |
| 491 | // TEXEL CENTRES IN WORLD METRES. |
| 492 | // |
| 493 | // ⚠ A first draft asserted these are BIT-IDENTICAL in both shaders, on the |
| 494 | // reasoning that Catmull-Rom collapses to p1 at an integer coordinate. That |
| 495 | // claim is UNPROVABLE and the probe never tested it: the world->texel |
| 496 | // division is not exact, so fract(u) came out 0.999973 in the reference and |
| 497 | // 2^-11 in the grid -- neither is an integer texel. The points are kept |
| 498 | // because they are where the coordinate error is largest and cleanest to |
| 499 | // attribute; the claim they carry is now the derived bound below. |
| 500 | const int I_EXACT = I_HOLE + 1; |
| 501 | const int N_EXACT = 16; |
| 502 | for (int k = 0; k < N_EXACT; k++) { |
| 503 | const int tk = INSET + k * 8; |
| 504 | pts.push_back((float)(((double)A->ty * payload + tk) * mpp)); |
| 505 | pts.push_back((float)(((double)A->tx * payload + tk) * mpp)); |
| 506 | pts.push_back(0.0f); pts.push_back(0.0f); |
| 507 | } |
| 508 | |
| 509 | HeightGridPack g = erosion_heightmap_pack_tiles(tiles, STUB); |
| 510 | ok("pack accepted the tile set", g.ok); |
| 511 | if (!g.ok) { printf("\nFAIL cannot continue\n"); return 1; } |
| 512 | printf(" grid %dx%d cells from (%d,%d), %d tiles, %.1f MiB\n\n", |
| 513 | g.nx, g.ny, g.tx0, g.ty0, g.tiles, g.data.size() * 4.0 / 1048576.0); |
| 514 | |
| 515 | std::vector<float> ref = run_sampler(c, single_tile_consts(*A), false, A->h, pts, "reference"); |
| 516 | std::vector<float> got = run_sampler(c, grid_consts(g), true, g.data, pts, "paged"); |
| 517 | |
| 518 | printf("1. THE PAGED SAMPLER AGREES WITH THE SINGLE-TILE SAMPLER\n"); |
| 519 | double worst = 0.0; |
| 520 | for (int i = 0; i < N_A; i++) worst = std::max(worst, (double)fabsf(got[i] - ref[i])); |
| 521 | char det[128]; |
| 522 | snprintf(det, sizeof(det), "max |paged-ref| = %.6f m over %d pts", worst, N_A); |
| 523 | ok("tile A interior matches the trusted reference", worst < 0.01, det); |
| 524 | snprintf(det, sizeof(det), "%.2f m", ref[0]); |
| 525 | ok(" and the reference is real terrain, not a flat buffer", |
| 526 | ref[0] != STUB && fabsf(ref[0]) < 9000.0f, det); |
| 527 | // THE RESIDUAL IS DERIVED, NOT TUNED. |
| 528 | // |
| 529 | // A per-tile origin makes (x - HBLK_OX) a small number: tile A's texels run |
| 530 | // 0..257, so u is exact to well under a ULP. The GRID origin does not |
| 531 | // cancel the large part -- the same point is texel 7176 of the grid -- and |
| 532 | // the world->texel division is NOT exact there. MEASURED: fract(u) = 2^-11 |
| 533 | // = 0.000488281 at a point that is an exact texel centre in world metres, |
| 534 | // which is precisely one ULP of 7176. |
| 535 | // |
| 536 | // One ULP of u, times the local slope in metres per texel, is therefore the |
| 537 | // floor on paged-vs-reference agreement. It is not a defect and it is not a |
| 538 | // tolerance chosen to make the test pass: it is a cost of moving the origin |
| 539 | // to the grid corner, it GROWS with world extent (the full 119x119 world |
| 540 | // reaches texel 30,464, so ~2x this), and a regression that breaks |
| 541 | // ADDRESSING misses by metres, not by ULPs -- controls A' and C below miss |
| 542 | // by 405 m and 102,664 m. |
| 543 | double max_slope = 0.0; // metres per texel, over tile A |
| 544 | for (int j = 1; j < A->n; j++) |
| 545 | for (int i = 1; i < A->n; i++) { |
| 546 | const double h = A->h[(size_t)j * A->n + i]; |
| 547 | max_slope = std::max(max_slope, fabs(h - A->h[(size_t)j * A->n + i - 1])); |
| 548 | max_slope = std::max(max_slope, fabs(h - A->h[(size_t)(j - 1) * A->n + i])); |
| 549 | } |
| 550 | const float u_max = (float)((double)std::max(g.nx, g.ny) * g.payload); |
| 551 | const double ulp_u = (double)(nextafterf(u_max, 2.0f * u_max) - u_max); |
| 552 | const double bound = 4.0 * ulp_u * max_slope; |
| 553 | // Over BOTH probe sets: the fractional interior AND the world-metre texel |
| 554 | // centres, which is where the coordinate error is largest. |
| 555 | double worst_all = worst; |
| 556 | for (int i = I_EXACT; i < I_EXACT + N_EXACT; i++) |
| 557 | worst_all = std::max(worst_all, (double)fabsf(got[i] - ref[i])); |
| 558 | snprintf(det, sizeof(det), "%.6f m <= %.6f = 4 x ULP(%.0f) x %.1f m/texel", |
| 559 | worst_all, bound, u_max, max_slope); |
| 560 | ok(" and the residual is within the coordinate-precision bound", |
| 561 | worst_all <= bound, det); |
| 562 | printf("\n2. AN ABSENT TILE READS AS THE STUB\n"); |
| 563 | snprintf(det, sizeof(det), "%.1f", got[I_HOLE]); |
| 564 | ok("hole samples exactly the stub", got[I_HOLE] == STUB, det); |
| 565 | |
| 566 | printf("\n CONTROL: the stub is not what a clamp would have returned\n"); |
| 567 | snprintf(det, sizeof(det), "reference clamps to %.2f m", ref[I_HOLE]); |
| 568 | ok("single-tile sampler returns terrain at the hole", ref[I_HOLE] != STUB, det); |
| 569 | ok(" so stub != clamp -- the probe above is not vacuous", |
| 570 | got[I_HOLE] != ref[I_HOLE]); |
| 571 | |
| 572 | printf("\n CONTROL A': corrupt one index-table entry -> probe 1 must FAIL\n"); |
| 573 | { |
| 574 | HeightGridPack bad = g; |
| 575 | const size_t cell = (size_t)(A->tx - g.tx0) * (size_t)g.ny + (size_t)(A->ty - g.ty0); |
| 576 | const int was = (int)bad.data[cell]; |
| 577 | bad.data[cell] = (float)((was + 1) % g.tiles); |
| 578 | std::vector<float> bg = run_sampler(c, grid_consts(bad), true, bad.data, pts, "corrupt"); |
| 579 | double w = 0.0; |
| 580 | for (int i = 0; i < N_A; i++) w = std::max(w, (double)fabsf(bg[i] - ref[i])); |
| 581 | snprintf(det, sizeof(det), "max |paged-ref| = %.2f m", w); |
| 582 | ok("a wrong table entry is DETECTED", w > 1.0, det); |
| 583 | } |
| 584 | |
| 585 | printf("\n CONTROL C: swap tx/ty when packing -> probe 1 must FAIL\n"); |
| 586 | { |
| 587 | HeightGridPack sw = erosion_heightmap_pack_tiles(tiles, STUB, true); |
| 588 | if (!sw.ok) { ok("swapped pack built", false); } |
| 589 | else { |
| 590 | std::vector<float> sg = run_sampler(c, grid_consts(sw), true, sw.data, pts, "swapped"); |
| 591 | double w = 0.0; |
| 592 | for (int i = 0; i < N_A; i++) w = std::max(w, (double)fabsf(sg[i] - ref[i])); |
| 593 | snprintf(det, sizeof(det), "max |paged-ref| = %.2f m", w); |
| 594 | ok("a mirrored world is DETECTED", w > 1.0, det); |
| 595 | } |
| 596 | } |
| 597 | |
| 598 | // REPORTED, NOT ASSERTED. hblk_sample_level clamps world->texel to |
| 599 | // HBLK_SIDE[0] before the fetch runs, so a point far outside the grid |
| 600 | // resolves to the bbox edge instead of the stub -- a real divergence from |
| 601 | // HeightfieldNode::sample_height, and work this gate's scope excludes. |
| 602 | printf("\n3. KNOWN GAP (reported, not asserted)\n"); |
| 603 | { |
| 604 | // FAR ON X ONLY, and on tile A's own Z row. Going far on BOTH axes is |
| 605 | // not a test: the second axis's bounds check returns the stub and the |
| 606 | // probe passes for the wrong reason -- which is what a first draft of |
| 607 | // this probe did. |
| 608 | std::vector<float> far_pts; |
| 609 | far_pts.push_back((float)(((double)(g.ty0 + g.ny + 500) * payload) * mpp)); |
| 610 | far_pts.push_back(tile_centre_z(*A)); |
| 611 | far_pts.push_back(0.0f); far_pts.push_back(0.0f); |
| 612 | std::vector<float> fg = run_sampler(c, grid_consts(g), true, g.data, far_pts, "far"); |
| 613 | printf(" far outside the grid returns %.2f m (stub is %.1f)\n", fg[0], STUB); |
| 614 | printf(" hblk_sample_level's clamp runs BEFORE the fetch; the CPU twin\n"); |
| 615 | printf(" returns the stub here. Out of scope for level-0 addressing.\n"); |
| 616 | } |
| 617 | |
| 618 | printf("\n%s (%d failure%s)\n\n", g_fails == 0 ? "PASS" : "FAIL", g_fails, |
| 619 | g_fails == 1 ? "" : "s"); |
| 620 | return g_fails == 0 ? 0 : 1; |
| 621 | } |