From fe32d8ab1f93218bf3d74479e86f7bd490cdbd1a Mon Sep 17 00:00:00 2001 From: Syoyo Fujita Date: Thu, 2 Jul 2026 01:37:01 +0900 Subject: [PATCH 1/3] vk/rtx: large-scene instanced ray-query stack (port from tinyusdz fork) Brings the vendored tinyusdz fork's large-scene RT work back upstream. It scales the VK_KHR_ray_query path from a single-BLAS scene to full-scale instanced assemblies (Moana island, ~42M instances) without faulting the driver, and hardens teardown against device loss. Verified: `make vk_test` passes on an RTX 5060 Ti (BVH8/BVH4 + ray_query + indexed ray_query, 0 mismatches). Included: - vkew: fix the VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_* sType values (they were off -- e.g. GEOMETRY_TRIANGLES_DATA was 1000150000 instead of the correct 1000150005; BUILD_GEOMETRY_INFO is 1000150000, CREATE_INFO 1000150017) and add the AS-build pipeline-stage/access sync flags. - Two-level GPU TLAS: per-prototype BLAS built once and shared across instances (lrt_vk_rtx_scene_build_instanced), instead of one soup BLAS per instance. - Device-local AS build input: stage vertex/index build inputs into device-local memory (NVIDIA requires it) via a one-shot copy. - Pooled per-prototype BLAS allocation (vk_subpool_*): a few 128 MiB blocks for input + AS storage + reused scratch, instead of ~6 vkAllocateMemory + 2 submits per prototype -- the naive path exhausts NVIDIA's allocator on ~100k prototypes. - 64-bit wide hit id: a 5-word/hit encoding (trace_ray_query_wide variant, -DLRT_WIDE_ID) storing instanceId + prim index separately, so scenes past the 32-bit prim_id*instances product still trace. compile_shaders.sh grows a separate source basename + extra-args so one .comp emits both variants. - Multi-TLAS slicing: split instances across K TLASes sharing one BLAS set and CPU-merge the nearest hit, past the single-TLAS instance limit. - Device-lost-safe teardown: a 5 s fence timeout flags device_lost on timeout / VK_ERROR_DEVICE_LOST; every destroy then skips driver calls and only frees host memory (NVIDIA segfaults inside the driver when destroying accels after a fault). - Hardening (pre-push audit): guard the trace descriptor-pool destroy with device_lost; free the suballocation pools on the build-failure path; only flag device_lost on a genuine GPU-build failure (a host-alloc failure keeps the engine usable). Co-Authored-By: Claude Opus 4.8 --- lightrt_c_vk.c | 870 ++++++++++++++++++++++++-- lightrt_c_vk.h | 93 ++- lightrt_vkew.h | 19 +- scripts/compile_shaders.sh | 24 +- vk/shaders/trace_ray_query.comp | 54 +- vk/shaders/trace_ray_query.spv.h | 463 +++++++++----- vk/shaders/trace_ray_query_wide.spv.h | 158 +++++ 7 files changed, 1456 insertions(+), 225 deletions(-) create mode 100644 vk/shaders/trace_ray_query_wide.spv.h diff --git a/lightrt_c_vk.c b/lightrt_c_vk.c index ecb0f8b..935c4d9 100644 --- a/lightrt_c_vk.c +++ b/lightrt_c_vk.c @@ -23,6 +23,7 @@ #include "vk/shaders/trace_bvh.spv.h" /* trace_bvh_spv[] */ #include "vk/shaders/build_morton.spv.h" /* build_morton_spv[] */ #include "vk/shaders/trace_ray_query.spv.h" /* trace_ray_query_spv[] */ +#include "vk/shaders/trace_ray_query_wide.spv.h" /* trace_ray_query_wide_spv[] */ #include "vk/shaders/shade_analytic.spv.h" /* shade_analytic_spv[] */ /* Host-read sync flags not declared in lightrt_vkew.h. */ @@ -99,10 +100,15 @@ struct lrt_vk_engine { uint32_t caps; char device_name[256]; char err[512]; + /* Set once any submit returns VK_ERROR_DEVICE_LOST (e.g. a build ran the GPU + * out of memory). On a lost device the NVIDIA driver crashes inside + * vkDestroy*, so all teardown skips driver calls and just frees host state. */ + int device_lost; vk_pipeline trace_pipes[VK_TRACE_PIPE_CACHE]; vk_pipeline build_pipe; /* build_morton: no spec constants */ vk_pipeline rtx_pipe; /* trace_ray_query: AS + 2 SSBO bindings */ + vk_pipeline rtx_pipe_wide; /* trace_ray_query_wide: 5-word hits (inst,prim) */ vk_pipeline shade_pipe; /* shade_analytic: 3 SSBO bindings + push */ }; @@ -250,8 +256,10 @@ static int vk_buffer_create_staging(lrt_vk_engine *e, VkDeviceSize size, } static void vk_buffer_destroy(lrt_vk_engine *e, vk_buffer *b) { - if (b->buf) vkDestroyBuffer(e->device, b->buf, NULL); - if (b->mem) vkFreeMemory(e->device, b->mem, NULL); + if (!e->device_lost) { /* driver calls crash on a lost device; just free host */ + if (b->buf) vkDestroyBuffer(e->device, b->buf, NULL); + if (b->mem) vkFreeMemory(e->device, b->mem, NULL); + } memset(b, 0, sizeof(*b)); } @@ -650,6 +658,7 @@ lrt_vk_engine *lrt_vk_engine_create(const lrt_vk_engine_options *opts, asf.pNext = &rqf; v12.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_2_FEATURES; v12.bufferDeviceAddress = VK_TRUE; + v12.descriptorIndexing = VK_TRUE; v12.pNext = &asf; f2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2; f2.pNext = &v12; @@ -727,17 +736,22 @@ lrt_vk_engine *lrt_vk_engine_create(const lrt_vk_engine_options *opts, void lrt_vk_engine_destroy(lrt_vk_engine *e) { if (!e) return; - if (e->device) { + /* On a lost/corrupted device (e.g. a build ran the GPU out of memory) the + * NVIDIA driver segfaults inside vkDestroy / vkDeviceWaitIdle, so skip ALL + * driver teardown and just free the host handle -- the OS reclaims the GPU + * objects when the process exits. */ + if (e->device && !e->device_lost) { vkDeviceWaitIdle(e->device); for (int i = 0; i < VK_TRACE_PIPE_CACHE; i++) vk_pipeline_destroy(e, &e->trace_pipes[i]); vk_pipeline_destroy(e, &e->build_pipe); vk_pipeline_destroy(e, &e->rtx_pipe); + vk_pipeline_destroy(e, &e->rtx_pipe_wide); vk_pipeline_destroy(e, &e->shade_pipe); if (e->cmd_pool) vkDestroyCommandPool(e->device, e->cmd_pool, NULL); vkDestroyDevice(e->device, NULL); } - if (e->instance) vkDestroyInstance(e->instance, NULL); + if (e->instance && !e->device_lost) vkDestroyInstance(e->instance, NULL); free(e); } @@ -753,6 +767,59 @@ const char *lrt_vk_engine_last_error(const lrt_vk_engine *e) { return e ? e->err : ""; } +uint64_t lrt_vk_device_local_bytes(int prefer_discrete) { + if (!vkewInit()) return 0; + VkApplicationInfo app; + memset(&app, 0, sizeof(app)); + app.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; + app.pApplicationName = "lightrt"; + app.pEngineName = "lightrt"; + app.apiVersion = VK_API_VERSION_1_2; + VkInstanceCreateInfo ici; + memset(&ici, 0, sizeof(ici)); + ici.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; + ici.pApplicationInfo = &app; + VkInstance inst = VK_NULL_HANDLE; + if (vkCreateInstance(&ici, NULL, &inst) != VK_SUCCESS) return 0; + if (!vkewLoadInstance(inst)) { + vkDestroyInstance(inst, NULL); + return 0; + } + uint32_t ndev = 0; + vkEnumeratePhysicalDevices(inst, &ndev, NULL); + if (ndev == 0) { + vkDestroyInstance(inst, NULL); + return 0; + } + if (ndev > 16) ndev = 16; + VkPhysicalDevice devs[16]; + vkEnumeratePhysicalDevices(inst, &ndev, devs); + uint64_t best = 0; + int best_discrete = 0; + for (uint32_t d = 0; d < ndev; d++) { + VkPhysicalDeviceProperties pp; + vkGetPhysicalDeviceProperties(devs[d], &pp); + int is_discrete = (pp.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU); + VkPhysicalDeviceMemoryProperties mp; + vkGetPhysicalDeviceMemoryProperties(devs[d], &mp); + uint64_t local = 0; + for (uint32_t h = 0; h < mp.memoryHeapCount; h++) { + if ((mp.memoryHeaps[h].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT) && + mp.memoryHeaps[h].size > local) { + local = mp.memoryHeaps[h].size; + } + } + /* Prefer a discrete GPU's VRAM; otherwise take the largest local heap. */ + if ((prefer_discrete && is_discrete && !best_discrete) || + (is_discrete == best_discrete && local > best)) { + best = local; + best_discrete = is_discrete; + } + } + vkDestroyInstance(inst, NULL); + return best; +} + /* ------------------------------------------------------------------------- */ /* Path A: CPU build -> GPU trace. */ /* ------------------------------------------------------------------------- */ @@ -1363,7 +1430,7 @@ int lrt_vk_shade_scene_render(lrt_vk_engine *e, lrt_vk_shade_scene *s, void lrt_vk_shade_scene_free(lrt_vk_engine *e, lrt_vk_shade_scene *s) { if (!s) return; - if (e->device) vkDeviceWaitIdle(e->device); + if (e->device && !e->device_lost) vkDeviceWaitIdle(e->device); if (s->have_cmd) { vkFreeCommandBuffers(e->device, e->cmd_pool, 1, &s->cmd); vkDestroyFence(e->device, s->fence, NULL); @@ -1426,10 +1493,29 @@ static int vk_cmd_end_submit(lrt_vk_engine *e, VkCommandBuffer cb) { si.commandBufferCount = 1; si.pCommandBuffers = &cb; VkResult r = vkQueueSubmit(e->queue, 1, &si, fence); - if (r == VK_SUCCESS) r = vkWaitForFences(e->device, 1, &fence, VK_TRUE, ~0ULL); + if (r == VK_SUCCESS) { + /* 5 second timeout — ~0ULL (infinite) would hang the process on a + * device-lost or misconfigured AS build. */ + r = vkWaitForFences(e->device, 1, &fence, VK_TRUE, 5000000000ULL); + if (r == VK_TIMEOUT) { + vk_set_err(e, "GPU command timed out after 5 s (device lost or " + "misconfigured acceleration-structure build)"); + e->device_lost = 1; + vkDestroyFence(e->device, fence, NULL); + vkFreeCommandBuffers(e->device, e->cmd_pool, 1, &cb); + return 0; + } + } + if (r != VK_SUCCESS) { + vk_set_errf(e, "vkQueueSubmit/vkWaitForFences", r); + if (r == VK_ERROR_DEVICE_LOST || r == VK_TIMEOUT) e->device_lost = 1; + vkDestroyFence(e->device, fence, NULL); + vkFreeCommandBuffers(e->device, e->cmd_pool, 1, &cb); + return 0; + } vkDestroyFence(e->device, fence, NULL); vkFreeCommandBuffers(e->device, e->cmd_pool, 1, &cb); - return r == VK_SUCCESS; + return 1; } typedef struct vk_accel { @@ -1443,7 +1529,8 @@ typedef struct vk_accel { * the build, and resolves the AS device address. */ static int vk_build_accel(lrt_vk_engine *e, VkAccelerationStructureTypeKHR type, const VkAccelerationStructureGeometryKHR *geom, - uint32_t prim_count, vk_accel *out) { + uint32_t prim_count, uint32_t prim_offset, + uint32_t first_vertex, vk_accel *out) { memset(out, 0, sizeof(*out)); VkAccelerationStructureBuildGeometryInfoKHR bgi; memset(&bgi, 0, sizeof(bgi)); @@ -1496,8 +1583,8 @@ static int vk_build_accel(lrt_vk_engine *e, VkAccelerationStructureTypeKHR type, bgi.scratchData.deviceAddress = scratch_addr; VkAccelerationStructureBuildRangeInfoKHR range; range.primitiveCount = prim_count; - range.primitiveOffset = 0; - range.firstVertex = 0; + range.primitiveOffset = prim_offset; + range.firstVertex = first_vertex; range.transformOffset = 0; const VkAccelerationStructureBuildRangeInfoKHR *pranges = ⦥ @@ -1505,6 +1592,16 @@ static int vk_build_accel(lrt_vk_engine *e, VkAccelerationStructureTypeKHR type, int ok = cb != VK_NULL_HANDLE; if (ok) { vkCmdBuildAccelerationStructuresKHR(cb, 1, &bgi, &pranges); + VkMemoryBarrier mb; + memset(&mb, 0, sizeof(mb)); + mb.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER; + mb.srcAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_WRITE_BIT_KHR; + mb.dstAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR; + vkCmdPipelineBarrier(cb, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR | + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + 0, 1, &mb, 0, NULL, 0, NULL); ok = vk_cmd_end_submit(e, cb); } vk_buffer_destroy(e, &scratch); @@ -1525,15 +1622,23 @@ static int vk_build_accel(lrt_vk_engine *e, VkAccelerationStructureTypeKHR type, } static void vk_accel_destroy(lrt_vk_engine *e, vk_accel *a) { - if (a->as) vkDestroyAccelerationStructureKHR(e->device, a->as, NULL); + if (a->as && !e->device_lost) + vkDestroyAccelerationStructureKHR(e->device, a->as, NULL); vk_buffer_destroy(e, &a->storage); memset(a, 0, sizeof(*a)); } typedef struct rtx_push { uint32_t ray_count; + uint32_t tri_chunk; /* prim_id = instanceId*tri_chunk + primitiveIndex */ } rtx_push; +/* Triangles per BLAS chunk. Kept well below the count at which RADV/gfx1201's + * vkGetAccelerationStructureBuildSizesKHR returns a corrupted (under-sized) + * scratch size and the build GPU-faults (empirically OK at 16M, broken by 24M). + * Must be a power of two? No — the trace shader just multiplies by it. */ +#define LRT_VK_BLAS_CHUNK (8u * 1024u * 1024u) + static vk_pipeline *rtx_get_pipeline(lrt_vk_engine *e) { if (e->rtx_pipe.valid) return &e->rtx_pipe; VkDescriptorType types[3] = {VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR, @@ -1545,6 +1650,22 @@ static vk_pipeline *rtx_get_pipeline(lrt_vk_engine *e) { return &e->rtx_pipe; } +/* Wide-id trace pipeline: identical bindings/push to rtx_pipe, but the shader + * writes 5 words/hit ({t,u,v,instanceId,prim}) so instanced scenes past the + * 32-bit prim_id product still trace. Same descriptor layout, so scenes can bind + * against either pipeline. */ +static vk_pipeline *rtx_get_pipeline_wide(lrt_vk_engine *e) { + if (e->rtx_pipe_wide.valid) return &e->rtx_pipe_wide; + VkDescriptorType types[3] = {VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR, + VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + VK_DESCRIPTOR_TYPE_STORAGE_BUFFER}; + if (!vk_pipeline_build(e, trace_ray_query_wide_spv, + sizeof(trace_ray_query_wide_spv), 3, types, + (uint32_t)sizeof(rtx_push), NULL, &e->rtx_pipe_wide)) + return NULL; + return &e->rtx_pipe_wide; +} + /* TLAS + 2 storage buffers (rays, hits). The AS write is chained via pNext. */ static int vk_descriptors_bind_rtx(lrt_vk_engine *e, VkDescriptorSetLayout dsl, VkAccelerationStructureKHR tlas, @@ -1620,15 +1741,73 @@ static int vk_descriptors_bind_rtx(lrt_vk_engine *e, VkDescriptorSetLayout dsl, * staging) is allocated lazily and reused/grown across traces. This is what * makes lrt_vk_rtx_scene_trace measure traversal, not per-call AS build. */ struct lrt_vk_rtx_scene { - vk_accel blas; + /* The geometry is split into chunks of <= LRT_VK_BLAS_CHUNK triangles, one + * BLAS per chunk, to work around drivers (e.g. RADV/gfx1201) whose AS + * build-size query overflows for very large single-BLAS builds and then + * GPU-faults. A single TLAS holds one identity instance per chunk; the trace + * shader recovers the global triangle id as instanceId*tri_chunk + primIndex. */ + vk_accel *blas; /* nblas chunk BLASes */ + uint32_t nblas; + uint32_t tri_chunk; /* triangles per chunk (the prim_id stride) */ vk_accel tlas; vk_buffer rays_dev; /* device-local STORAGE | TRANSFER_DST */ vk_buffer hits_dev; /* device-local STORAGE | TRANSFER_SRC */ vk_buffer rays_stage; /* host-visible TRANSFER_SRC */ vk_buffer hits_stage; /* host-visible TRANSFER_DST */ uint32_t cap; /* trace-buffer capacity in rays */ + uint32_t hit_words; /* words written per hit: 4 (narrow) or 5 (wide id) */ + /* Multi-TLAS: instanced scenes past the device TLAS maxInstanceCount (2^24) + * are split into `ntlas` slices of <= LRT_VK_TLAS_SLICE instances, each its + * own TLAS over the SHARED `blas` set. `tlas` is slice 0; `extra_tlas` holds + * slices 1..ntlas-1. `tlas_base[k]` is the global instance index where slice k + * starts (added back at trace so the reported instanceId is global). ntlas==1 + * is the ordinary single-TLAS scene (extra_tlas NULL). Wide id only. */ + vk_accel *extra_tlas; /* ntlas-1 additional TLAS slices, or NULL */ + uint32_t *tlas_base; /* ntlas global base instance indices */ + uint32_t ntlas; /* number of TLAS slices (>=1) */ + /* BLAS suballocation pools (see vk_subpool). A scene with ~100k prototypes + * would otherwise do ~6 vkAllocateMemory per BLAS -> hundreds of thousands of + * allocations (NVIDIA's allocator is ~O(live allocs); also a per-device count + * limit), exhausting/faulting the driver. These pool the per-BLAS vertex/index + * input and AS storage into a few large blocks each. Live for the scene. */ + struct vk_subpool *blas_input_pool; /* host-visible verts+indices, dev address */ + struct vk_subpool *blas_store_pool; /* device-local AS storage */ }; +/* Bump suballocator: many small "buffers" share a few large VkBuffer blocks (one + * VkDeviceMemory each), handed out as (block buffer, offset, device address). No + * per-item free -- the whole pool is released at scene teardown. Collapses the + * per-BLAS allocation storm into O(total_bytes / block_size) allocations. */ +typedef struct vk_pool_block { + vk_buffer buf; /* the block's VkBuffer + VkDeviceMemory (via vk_buffer) */ + uint64_t base_addr; /* device address of offset 0 (0 if not addressable) */ + void *mapped; /* persistent host pointer (host-visible pools), else NULL */ + VkDeviceSize used; /* bump offset */ +} vk_pool_block; + +typedef struct vk_subpool { + VkBufferUsageFlags usage; + int host_visible; + int device_addr; + VkDeviceSize block_size; + vk_pool_block *blocks; + uint32_t nblocks, cap; +} vk_subpool; + +/* Instances per TLAS slice for the multi-TLAS builder. Kept safely below the + * Vulkan-guaranteed TLAS maxInstanceCount floor (2^24 = 16,777,216). */ +#define LRT_VK_TLAS_SLICE (16u * 1000u * 1000u) + +/* Slice k's TLAS (slice 0 is the inline `tlas`, the rest live in extra_tlas). */ +static vk_accel *rtx_scene_tlas(lrt_vk_rtx_scene *s, uint32_t k) { + return k == 0u ? &s->tlas : &s->extra_tlas[k - 1u]; +} + +/* Bytes the trace shader writes per hit for scene s (narrow lrt_hit = 4 words). */ +static size_t rtx_hit_bytes(const lrt_vk_rtx_scene *s) { + return (size_t)(s->hit_words ? s->hit_words : 4u) * sizeof(uint32_t); +} + /* Shared BLAS+TLAS build behind the soup and indexed entry points. When * `indices` is NULL the BLAS is built from a de-indexed triangle soup * (`vertices` = 9*ntris floats, nverts == 3*ntris, no index buffer). When @@ -1665,32 +1844,54 @@ static lrt_vk_rtx_scene *rtx_scene_build_core(lrt_vk_engine *e, /* The vertex/index/instance build inputs are only needed during the build — * once vkCmdBuildAccelerationStructures completes, the AS is self-contained, - * so they are freed here and the resident scene holds only the BLAS + TLAS. */ + * so they are freed here and the resident scene holds only the BLAS + TLAS. + * The vertex/index build inputs are staged into device-local memory: the AS + * build input must be device-local on many implementations (NVIDIA requires + * it). */ vk_buffer vbuf = {0}, idxbuf = {0}, instbuf = {0}; + vk_buffer vstage = {0}, istage = {0}; int ok = 0; do { VkDeviceSize vbytes = (VkDeviceSize)nverts * 3u * sizeof(float); + VkDeviceSize ibytes = (VkDeviceSize)ntris * 3u * sizeof(uint32_t); + /* Vertex staging (host-visible) + device-local AS build input. */ + if (!vk_buffer_create_ex(e, vbytes, VK_BUFFER_USAGE_TRANSFER_SRC_BIT, 1, 0, + &vstage)) + break; + if (!vk_buffer_write(e, &vstage, vertices, (size_t)vbytes)) + break; if (!vk_buffer_create_ex( e, vbytes, - VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR, - 1, 1, &vbuf)) - break; - if (!vk_buffer_write(e, &vbuf, vertices, (size_t)vbytes)) + VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR | + VK_BUFFER_USAGE_TRANSFER_DST_BIT, + 0, 1, &vbuf)) break; - uint64_t vaddr = vk_device_address(e, &vbuf); - - uint64_t idxaddr = 0; + /* Optional index staging + device-local AS build input. */ if (indices) { - VkDeviceSize ibytes = (VkDeviceSize)ntris * 3u * sizeof(uint32_t); + if (!vk_buffer_create_ex(e, ibytes, VK_BUFFER_USAGE_TRANSFER_SRC_BIT, 1, + 0, &istage)) + break; + if (!vk_buffer_write(e, &istage, indices, (size_t)ibytes)) + break; if (!vk_buffer_create_ex( e, ibytes, - VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR, - 1, 1, &idxbuf)) - break; - if (!vk_buffer_write(e, &idxbuf, indices, (size_t)ibytes)) + VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR | + VK_BUFFER_USAGE_TRANSFER_DST_BIT, + 0, 1, &idxbuf)) break; - idxaddr = vk_device_address(e, &idxbuf); } + /* One command buffer stages both vertex and (optional) index data. */ + VkCommandBuffer cb = vk_cmd_begin(e); + if (!cb) break; + VkBufferCopy vcopy = {0, 0, vbytes}; + vkCmdCopyBuffer(cb, vstage.buf, vbuf.buf, 1, &vcopy); + if (indices) { + VkBufferCopy icopy = {0, 0, ibytes}; + vkCmdCopyBuffer(cb, istage.buf, idxbuf.buf, 1, &icopy); + } + if (!vk_cmd_end_submit(e, cb)) break; + uint64_t vaddr = vk_device_address(e, &vbuf); + uint64_t idxaddr = indices ? vk_device_address(e, &idxbuf) : 0; VkAccelerationStructureGeometryKHR tgeom; memset(&tgeom, 0, sizeof(tgeom)); @@ -1709,23 +1910,56 @@ static lrt_vk_rtx_scene *rtx_scene_build_core(lrt_vk_engine *e, } else { tgeom.geometry.triangles.indexType = VK_INDEX_TYPE_NONE_KHR; } - if (!vk_build_accel(e, VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR, - &tgeom, ntris, &s->blas)) - break; - - VkAccelerationStructureInstanceKHR inst; - memset(&inst, 0, sizeof(inst)); - inst.transform.matrix[0][0] = 1.0f; - inst.transform.matrix[1][1] = 1.0f; - inst.transform.matrix[2][2] = 1.0f; - inst.mask = 0xFFu; - inst.accelerationStructureReference = s->blas.address; - if (!vk_buffer_create_ex( - e, sizeof(inst), + /* Split into chunk BLASes (one TLAS instance each). The build range's + * primitiveOffset is a BYTE offset into the index buffer (indexed) or the + * vertex buffer (soup); the shared tgeom keeps the full buffer addresses + * and maxVertex. */ + const uint32_t chunk = LRT_VK_BLAS_CHUNK; + const uint32_t nblas = (ntris + chunk - 1u) / chunk; + s->blas = (vk_accel *)calloc(nblas, sizeof(vk_accel)); + if (!s->blas) break; + s->nblas = nblas; + s->tri_chunk = chunk; + int blas_ok = 1; + for (uint32_t c = 0; c < nblas; c++) { + uint32_t base = c * chunk; + uint32_t ct = (ntris - base < chunk) ? (ntris - base) : chunk; + uint32_t prim_off = + indices ? base * 3u * (uint32_t)sizeof(uint32_t) + : base * 3u * 3u * (uint32_t)sizeof(float); + if (!vk_build_accel(e, VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR, + &tgeom, ct, prim_off, 0u, &s->blas[c])) { + blas_ok = 0; + break; + } + } + if (!blas_ok) break; + + /* One identity instance per chunk BLAS, in chunk order, so the hit's + * instanceId (build order) recovers the global triangle id in the trace + * shader as instanceId*tri_chunk + primitiveIndex. */ + VkAccelerationStructureInstanceKHR *insts = + (VkAccelerationStructureInstanceKHR *)calloc( + nblas, sizeof(VkAccelerationStructureInstanceKHR)); + if (!insts) break; + for (uint32_t c = 0; c < nblas; c++) { + insts[c].transform.matrix[0][0] = 1.0f; + insts[c].transform.matrix[1][1] = 1.0f; + insts[c].transform.matrix[2][2] = 1.0f; + insts[c].mask = 0xFFu; + insts[c].instanceCustomIndex = c; + insts[c].accelerationStructureReference = s->blas[c].address; + } + VkDeviceSize inst_bytes = + (VkDeviceSize)nblas * sizeof(VkAccelerationStructureInstanceKHR); + int ibuilt = + vk_buffer_create_ex( + e, inst_bytes, VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR, - 1, 1, &instbuf)) - break; - if (!vk_buffer_write(e, &instbuf, &inst, sizeof(inst))) break; + 1, 1, &instbuf) && + vk_buffer_write(e, &instbuf, insts, (size_t)inst_bytes); + free(insts); + if (!ibuilt) break; uint64_t iaddr = vk_device_address(e, &instbuf); VkAccelerationStructureGeometryKHR igeom; @@ -1738,17 +1972,22 @@ static lrt_vk_rtx_scene *rtx_scene_build_core(lrt_vk_engine *e, igeom.geometry.instances.arrayOfPointers = VK_FALSE; igeom.geometry.instances.data.deviceAddress = iaddr; if (!vk_build_accel(e, VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR, &igeom, - 1, &s->tlas)) + nblas, 0u, 0u, &s->tlas)) break; ok = 1; } while (0); + vk_buffer_destroy(e, &vstage); + vk_buffer_destroy(e, &istage); vk_buffer_destroy(e, &vbuf); vk_buffer_destroy(e, &idxbuf); vk_buffer_destroy(e, &instbuf); if (!ok) { vk_accel_destroy(e, &s->tlas); - vk_accel_destroy(e, &s->blas); + if (s->blas) { + for (uint32_t c = 0; c < s->nblas; c++) vk_accel_destroy(e, &s->blas[c]); + free(s->blas); + } free(s); if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return NULL; @@ -1775,6 +2014,435 @@ lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_indexed(lrt_vk_engine *e, return rtx_scene_build_core(e, vertices, nverts, indices, ntris, err); } +/* --- BLAS suballocation pool ------------------------------------------------- */ + +static void vk_subpool_init(vk_subpool *p, VkBufferUsageFlags usage, + int host_visible, int device_addr, + VkDeviceSize block_size) { + memset(p, 0, sizeof(*p)); + p->usage = usage; + p->host_visible = host_visible; + p->device_addr = device_addr; + p->block_size = block_size; +} + +/* Suballocate `size` bytes (aligned to power-of-two `align`) from the pool. Returns + * the containing block's VkBuffer, the byte offset within it, and -- when + * applicable -- the persistent host pointer and device address at that offset. */ +static int vk_subpool_alloc(lrt_vk_engine *e, vk_subpool *p, VkDeviceSize size, + VkDeviceSize align, VkBuffer *out_buf, + VkDeviceSize *out_off, void **out_mapped, + uint64_t *out_addr) { + if (size == 0) size = 16; + if (align == 0) align = 1; + VkDeviceSize off = 0; + int need_new = 1; + if (p->nblocks) { + vk_pool_block *last = &p->blocks[p->nblocks - 1]; + off = (last->used + (align - 1)) & ~(align - 1); + if (off + size <= last->buf.size) need_new = 0; + } + if (need_new) { + VkDeviceSize bs = size > p->block_size ? size : p->block_size; + if (p->nblocks == p->cap) { + uint32_t nc = p->cap ? p->cap * 2u : 8u; + vk_pool_block *nb = + (vk_pool_block *)realloc(p->blocks, nc * sizeof(vk_pool_block)); + if (!nb) return 0; + p->blocks = nb; + p->cap = nc; + } + vk_pool_block *blk = &p->blocks[p->nblocks]; + memset(blk, 0, sizeof(*blk)); + if (!vk_buffer_create_ex(e, bs, p->usage, p->host_visible, p->device_addr, + &blk->buf)) + return 0; + if (p->device_addr) blk->base_addr = vk_device_address(e, &blk->buf); + if (p->host_visible && + vkMapMemory(e->device, blk->buf.mem, 0, VK_WHOLE_SIZE, 0, &blk->mapped) != + VK_SUCCESS) { + vk_buffer_destroy(e, &blk->buf); + return 0; + } + p->nblocks++; + off = 0; + } + vk_pool_block *blk = &p->blocks[p->nblocks - 1]; + blk->used = off + size; + if (out_buf) *out_buf = blk->buf.buf; + if (out_off) *out_off = off; + if (out_mapped) *out_mapped = blk->mapped ? (char *)blk->mapped + off : NULL; + if (out_addr) *out_addr = blk->base_addr ? blk->base_addr + off : 0u; + return 1; +} + +static void vk_subpool_free(lrt_vk_engine *e, vk_subpool *p) { + if (!p) return; + for (uint32_t i = 0; i < p->nblocks; i++) { + if (p->blocks[i].mapped && !e->device_lost) + vkUnmapMemory(e->device, p->blocks[i].buf.mem); + vk_buffer_destroy(e, &p->blocks[i].buf); + } + free(p->blocks); + memset(p, 0, sizeof(*p)); +} + +/* Build all prototype BLAS through two suballocation pools (vertex/index INPUT and + * AS STORAGE) plus one reused scratch buffer, instead of ~6 vkAllocateMemory + 2 + * submits per prototype. On a ~100k-prototype scene (full Moana island) the naive + * path exhausts NVIDIA's allocator and the driver faults; pooling keeps it to a few + * dozen block allocations. Input verts/indices go in a host-visible, device-address + * pool (memcpy'd directly -- no staging, no copy submit); each BLAS is built into a + * pooled device-local storage slice. blas[p].as owns only the AS handle (storage is + * pool-owned; blas[p].storage stays empty). Returns 1 on success. */ +static int build_protos_pooled(lrt_vk_engine *e, const lrt_vk_proto *protos, + uint32_t nprotos, lrt_vk_rtx_scene *s) { + const VkDeviceSize BLOCK = 128u * 1024u * 1024u; // 128 MiB blocks + s->blas_input_pool = (vk_subpool *)calloc(1, sizeof(vk_subpool)); + s->blas_store_pool = (vk_subpool *)calloc(1, sizeof(vk_subpool)); + if (!s->blas_input_pool || !s->blas_store_pool) return 0; + vk_subpool_init( + s->blas_input_pool, + VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR, + /*host_visible=*/1, /*device_addr=*/1, BLOCK); + vk_subpool_init(s->blas_store_pool, + VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_STORAGE_BIT_KHR, + /*host_visible=*/0, /*device_addr=*/0, BLOCK); + + vk_buffer scratch = {0}; // reused across builds, grown on demand + int ok = 1; + for (uint32_t p = 0; p < nprotos && ok; p++) { + const lrt_vk_proto *pr = &protos[p]; + VkDeviceSize vbytes = (VkDeviceSize)pr->nverts * 3u * sizeof(float); + VkDeviceSize ibytes = (VkDeviceSize)pr->ntris * 3u * sizeof(uint32_t); + VkBuffer vbuf, ibuf; + VkDeviceSize voff, ioff; + void *vmap, *imap; + uint64_t vaddr, iaddr; + if (!vk_subpool_alloc(e, s->blas_input_pool, vbytes, 16, &vbuf, &voff, &vmap, + &vaddr) || + !vk_subpool_alloc(e, s->blas_input_pool, ibytes, 16, &ibuf, &ioff, &imap, + &iaddr)) { + ok = 0; + break; + } + memcpy(vmap, pr->vertices, (size_t)vbytes); + memcpy(imap, pr->indices, (size_t)ibytes); + + VkAccelerationStructureGeometryKHR g; + memset(&g, 0, sizeof(g)); + g.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR; + g.geometryType = VK_GEOMETRY_TYPE_TRIANGLES_KHR; + g.flags = VK_GEOMETRY_OPAQUE_BIT_KHR; + g.geometry.triangles.sType = + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_TRIANGLES_DATA_KHR; + g.geometry.triangles.vertexFormat = VK_FORMAT_R32G32B32_SFLOAT; + g.geometry.triangles.vertexData.deviceAddress = vaddr; + g.geometry.triangles.vertexStride = 3u * sizeof(float); + g.geometry.triangles.maxVertex = pr->nverts - 1u; + g.geometry.triangles.indexType = VK_INDEX_TYPE_UINT32; + g.geometry.triangles.indexData.deviceAddress = iaddr; + + VkAccelerationStructureBuildGeometryInfoKHR bgi; + memset(&bgi, 0, sizeof(bgi)); + bgi.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_GEOMETRY_INFO_KHR; + bgi.type = VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR; + bgi.flags = VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_TRACE_BIT_KHR; + bgi.mode = VK_BUILD_ACCELERATION_STRUCTURE_MODE_BUILD_KHR; + bgi.geometryCount = 1; + bgi.pGeometries = &g; + + VkAccelerationStructureBuildSizesInfoKHR sizes; + memset(&sizes, 0, sizeof(sizes)); + sizes.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_SIZES_INFO_KHR; + vkGetAccelerationStructureBuildSizesKHR( + e->device, VK_ACCELERATION_STRUCTURE_BUILD_TYPE_DEVICE_KHR, &bgi, + &pr->ntris, &sizes); + + VkBuffer sbuf; + VkDeviceSize soff; + if (!vk_subpool_alloc(e, s->blas_store_pool, sizes.accelerationStructureSize, + 256, &sbuf, &soff, NULL, NULL)) { + ok = 0; + break; + } + VkAccelerationStructureCreateInfoKHR ci; + memset(&ci, 0, sizeof(ci)); + ci.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_CREATE_INFO_KHR; + ci.buffer = sbuf; + ci.offset = soff; + ci.size = sizes.accelerationStructureSize; + ci.type = VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR; + if (vkCreateAccelerationStructureKHR(e->device, &ci, NULL, &s->blas[p].as) != + VK_SUCCESS) { + ok = 0; + break; + } + + VkDeviceSize need = sizes.buildScratchSize + 256u; + if (scratch.size < need) { + vk_buffer_destroy(e, &scratch); + if (!vk_buffer_create_ex(e, need, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, 0, 1, + &scratch)) { + ok = 0; + break; + } + } + uint64_t saddr = (vk_device_address(e, &scratch) + 255u) & ~(uint64_t)255u; + + bgi.dstAccelerationStructure = s->blas[p].as; + bgi.scratchData.deviceAddress = saddr; + VkAccelerationStructureBuildRangeInfoKHR range = {pr->ntris, 0u, 0u, 0u}; + const VkAccelerationStructureBuildRangeInfoKHR *pranges = ⦥ + + VkCommandBuffer cb = vk_cmd_begin(e); + if (cb == VK_NULL_HANDLE) { + ok = 0; + break; + } + vkCmdBuildAccelerationStructuresKHR(cb, 1, &bgi, &pranges); + VkMemoryBarrier mb; + memset(&mb, 0, sizeof(mb)); + mb.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER; + mb.srcAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_WRITE_BIT_KHR; + mb.dstAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR; + vkCmdPipelineBarrier(cb, VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR, 0, + 1, &mb, 0, NULL, 0, NULL); + if (!vk_cmd_end_submit(e, cb)) { + ok = 0; + break; + } + VkAccelerationStructureDeviceAddressInfoKHR ai; + memset(&ai, 0, sizeof(ai)); + ai.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_DEVICE_ADDRESS_INFO_KHR; + ai.accelerationStructure = s->blas[p].as; + s->blas[p].address = + vkGetAccelerationStructureDeviceAddressKHR(e->device, &ai); + } + vk_buffer_destroy(e, &scratch); + return ok; +} + +/* Build one TLAS over the instance slice insts[off .. off+count) referencing the + * scene's shared BLAS. instanceCustomIndex / build order is the slice-LOCAL index + * (0..count-1); the caller adds tlas_base[k]=off at trace to recover the global + * instance id. Returns 1 on success (out_tlas filled), 0 on failure. */ +static int build_tlas_slice(lrt_vk_engine *e, lrt_vk_rtx_scene *s, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t off, + uint32_t count, vk_accel *out_tlas) { + VkAccelerationStructureInstanceKHR *vi = + (VkAccelerationStructureInstanceKHR *)calloc( + count, sizeof(VkAccelerationStructureInstanceKHR)); + if (!vi) return 0; + vk_buffer instbuf = {0}; + int ok = 0; + do { + int inst_ok = 1; + for (uint32_t j = 0; j < count; j++) { + const lrt_vk_instance *in = &insts[off + j]; + if (in->proto >= nprotos) { + inst_ok = 0; + break; + } + /* transform[12] is row-major 3x4 (world = M*[p;1]), matching + * VkTransformMatrixKHR.matrix[3][4] exactly. */ + memcpy(vi[j].transform.matrix, in->transform, 12u * sizeof(float)); + vi[j].mask = 0xFFu; + vi[j].instanceCustomIndex = j; + vi[j].accelerationStructureReference = s->blas[in->proto].address; + } + if (!inst_ok) break; + VkDeviceSize inst_bytes = + (VkDeviceSize)count * sizeof(VkAccelerationStructureInstanceKHR); + if (!(vk_buffer_create_ex( + e, inst_bytes, + VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR, + 1, 1, &instbuf) && + vk_buffer_write(e, &instbuf, vi, (size_t)inst_bytes))) + break; + VkAccelerationStructureGeometryKHR igeom; + memset(&igeom, 0, sizeof(igeom)); + igeom.sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR; + igeom.geometryType = VK_GEOMETRY_TYPE_INSTANCES_KHR; + igeom.flags = VK_GEOMETRY_OPAQUE_BIT_KHR; + igeom.geometry.instances.sType = + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_INSTANCES_DATA_KHR; + igeom.geometry.instances.arrayOfPointers = VK_FALSE; + igeom.geometry.instances.data.deviceAddress = vk_device_address(e, &instbuf); + if (!vk_build_accel(e, VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR, &igeom, + count, 0u, 0u, out_tlas)) + break; + ok = 1; + } while (0); + vk_buffer_destroy(e, &instbuf); + free(vi); + return ok; +} + +/* Core of all instanced builders. `wide`=0 uses the narrow 4-word trace and + * enforces the 32-bit prim_id product (ninsts*maxPrototypeTris < 2^32); `wide`=1 + * uses the 5-word trace that stores instanceId + prim separately, so there is no + * product to overflow. `max_per_tlas`=0 builds a single TLAS; otherwise the + * instances are split into ceil(ninsts/max_per_tlas) TLAS slices sharing one BLAS + * set (multi-TLAS -- renders scenes past the device TLAS maxInstanceCount). Wide + * is required for multi-TLAS (the CPU merge needs the separate instanceId). */ +static lrt_vk_rtx_scene *rtx_build_instanced_core( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, int wide, + uint32_t max_per_tlas, uint32_t *out_tri_stride, lrt_result *err) { + if (!e || !protos || nprotos == 0 || !insts || ninsts == 0) { + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return NULL; + } + if (!(e->caps & LRT_VK_CAP_RAY_QUERY) || !vkCreateAccelerationStructureKHR) { + vk_set_err(e, "ray_query unavailable (create the engine with " + "want_ray_tracing=1 on an RT-capable device)"); + if (err) *err = LRT_RESULT_NOT_BUILT; + return NULL; + } + if (!(wide ? rtx_get_pipeline_wide(e) : rtx_get_pipeline(e))) { + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + return NULL; + } + /* Narrow: prim_id = instanceId*stride + prim must fit 32 bits and never + * collide with LRT_TRI_NO_HIT (0xFFFFFFFF). stride = max prototype tri count. + * Wide: instanceId and prim are stored separately, so stride is unused for + * encoding (kept only to report *out_tri_stride). */ + uint32_t stride = 1u; + for (uint32_t p = 0; p < nprotos; p++) { + if (!protos[p].vertices || !protos[p].indices || protos[p].ntris == 0 || + protos[p].nverts == 0) { + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return NULL; + } + if (protos[p].ntris > stride) stride = protos[p].ntris; + } + if (!wide && (uint64_t)ninsts * (uint64_t)stride >= (uint64_t)LRT_TRI_NO_HIT) { + vk_set_err(e, "instanced scene exceeds 32-bit prim_id encoding " + "(ninsts*maxPrototypeTris too large; use the wide builder)"); + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return NULL; + } + + lrt_vk_rtx_scene *s = (lrt_vk_rtx_scene *)calloc(1, sizeof(*s)); + if (!s) { + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + return NULL; + } + s->blas = (vk_accel *)calloc(nprotos, sizeof(vk_accel)); + if (!s->blas) { + free(s); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + return NULL; + } + s->nblas = nprotos; + s->tri_chunk = stride; + s->hit_words = wide ? 5u : 4u; + + /* Slice the instances across one or more TLASes (all sharing the BLAS set). + * LRT_VK_TLAS_SLICE (env) shrinks the slice for testing the merge on small + * scenes; it only applies to the multi-TLAS builder (max_per_tlas>0). */ + uint32_t cap = max_per_tlas; + if (cap) { + const char *ev = getenv("LRT_VK_TLAS_SLICE"); + if (ev) { + long v = atol(ev); + if (v > 0) cap = (uint32_t)v; + } + } + const uint32_t slice = (cap && cap < ninsts) ? cap : ninsts; + const uint32_t ntlas = (ninsts + slice - 1u) / slice; + + int ok = 0; + int gpu_fail = 0; /* a GPU build (not a host alloc) failed -> driver may be faulted */ + do { + /* Build every prototype BLAS through the suballocation pools (a few large + * blocks) instead of ~6 vkAllocateMemory + 2 submits per prototype, which + * on ~100k-prototype scenes (full Moana island) exhausts NVIDIA's allocator + * and faults the driver. */ + if (!build_protos_pooled(e, protos, nprotos, s)) { gpu_fail = 1; break; } + + s->tlas_base = (uint32_t *)calloc(ntlas, sizeof(uint32_t)); + if (!s->tlas_base) break; + if (ntlas > 1u) { + s->extra_tlas = (vk_accel *)calloc(ntlas - 1u, sizeof(vk_accel)); + if (!s->extra_tlas) break; + } + int slices_ok = 1; + for (uint32_t k = 0; k < ntlas; k++) { + const uint32_t off = k * slice; + const uint32_t count = (off + slice <= ninsts) ? slice : (ninsts - off); + s->tlas_base[k] = off; + if (!build_tlas_slice(e, s, nprotos, insts, off, count, + rtx_scene_tlas(s, k))) { + slices_ok = 0; + break; + } + } + if (!slices_ok) { gpu_fail = 1; break; } + s->ntlas = ntlas; + ok = 1; + } while (0); + + if (!ok) { + /* A GPU build in the (possibly ~100k-BLAS) storm failed -- typically the GPU + * ran out of memory. NVIDIA's driver is then left in a state where destroying + * the already-built accels segfaults INSIDE the driver, so mark the device + * unusable: vk_accel_destroy / vk_subpool_free / engine teardown then skip all + * driver calls and only free host memory (the OS reclaims the GPU on exit). + * A pure host-allocation failure leaves the device healthy, so tear the + * partial scene down normally and keep the engine usable for a retry. */ + if (gpu_fail) e->device_lost = 1; + vk_accel_destroy(e, &s->tlas); + if (s->extra_tlas) + for (uint32_t k = 1; k < ntlas; k++) + vk_accel_destroy(e, &s->extra_tlas[k - 1u]); + for (uint32_t p = 0; p < s->nblas; p++) vk_accel_destroy(e, &s->blas[p]); + if (s->blas_input_pool) { + vk_subpool_free(e, s->blas_input_pool); + free(s->blas_input_pool); + } + if (s->blas_store_pool) { + vk_subpool_free(e, s->blas_store_pool); + free(s->blas_store_pool); + } + free(s->extra_tlas); + free(s->tlas_base); + free(s->blas); + free(s); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + return NULL; + } + if (out_tri_stride) *out_tri_stride = stride; + if (err) *err = LRT_RESULT_OK; + return s; +} + +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, uint32_t *out_tri_stride, + lrt_result *err) { + return rtx_build_instanced_core(e, protos, nprotos, insts, ninsts, + /*wide=*/0, /*max_per_tlas=*/0, out_tri_stride, + err); +} + +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced_wide( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, lrt_result *err) { + return rtx_build_instanced_core(e, protos, nprotos, insts, ninsts, + /*wide=*/1, /*max_per_tlas=*/0, NULL, err); +} + +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced_multi( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, lrt_result *err) { + return rtx_build_instanced_core(e, protos, nprotos, insts, ninsts, + /*wide=*/1, /*max_per_tlas=*/LRT_VK_TLAS_SLICE, + NULL, err); +} + /* Grow the device-local + staging trace buffers to hold at least n rays. */ static int rtx_scene_ensure(lrt_vk_engine *e, lrt_vk_rtx_scene *s, uint32_t n) { if (s->rays_dev.buf && n <= s->cap) return 1; @@ -1784,7 +2452,7 @@ static int rtx_scene_ensure(lrt_vk_engine *e, lrt_vk_rtx_scene *s, uint32_t n) { vk_buffer_destroy(e, &s->hits_stage); s->cap = 0; VkDeviceSize rb = (VkDeviceSize)n * sizeof(lrt_ray); - VkDeviceSize hb = (VkDeviceSize)n * sizeof(lrt_hit); + VkDeviceSize hb = (VkDeviceSize)n * rtx_hit_bytes(s); if (!vk_buffer_create_ex(e, rb, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, @@ -1812,9 +2480,13 @@ static void vk_buf_barrier(VkCommandBuffer cb, VkFlags src_stage, VkFlags dst_st vkCmdPipelineBarrier(cb, src_stage, dst_stage, 0, 1, &mb, 0, NULL, 0, NULL); } -int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, - const lrt_ray *rays, uint32_t n, lrt_hit *out, - lrt_result *err) { +/* Shared trace body for both the narrow (lrt_hit, 4-word) and wide (lrt_hit_wide, + * 5-word) hit layouts. `pipe` selects the shader variant; `out` must point at n + * hits of rtx_hit_bytes(s). Both layouts store the miss sentinel LRT_TRI_NO_HIT in + * word 3 (narrow prim_id / wide instanceId), so the hit count is one check. */ +static int rtx_trace_impl(lrt_vk_engine *e, lrt_vk_rtx_scene *s, vk_pipeline *pipe, + VkAccelerationStructureKHR tlas, const lrt_ray *rays, + uint32_t n, void *out, lrt_result *err) { if (!e || !s || (n && (!rays || !out))) { if (err) *err = LRT_RESULT_INVALID_ARGUMENT; return -1; @@ -1823,11 +2495,11 @@ int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, if (err) *err = LRT_RESULT_OK; return 0; } - vk_pipeline *pipe = rtx_get_pipeline(e); if (!pipe || !rtx_scene_ensure(e, s, n)) { if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return -1; } + const size_t hit_bytes = rtx_hit_bytes(s); /* Host -> staging (bulk memcpy), then one command buffer does staging -> * device-local rays, dispatch (reads/writes VRAM), device-local hits -> @@ -1839,7 +2511,7 @@ int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, VkDescriptorPool pool; VkDescriptorSet set; - if (!vk_descriptors_bind_rtx(e, pipe->dsl, s->tlas.as, &s->rays_dev, + if (!vk_descriptors_bind_rtx(e, pipe->dsl, tlas, &s->rays_dev, &s->hits_dev, &pool, &set)) { if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return -1; @@ -1855,6 +2527,7 @@ int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, VK_ACCESS_TRANSFER_WRITE_BIT, VK_ACCESS_SHADER_READ_BIT); rtx_push pc; pc.ray_count = n; + pc.tri_chunk = s->tri_chunk; vkCmdBindPipeline(cb, VK_PIPELINE_BIND_POINT_COMPUTE, pipe->pipe); vkCmdBindDescriptorSets(cb, VK_PIPELINE_BIND_POINT_COMPUTE, pipe->layout, 0, 1, &set, 0, NULL); @@ -1864,39 +2537,124 @@ int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, vk_buf_barrier(cb, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_SHADER_WRITE_BIT, VK_ACCESS_TRANSFER_READ_BIT); - VkBufferCopy down = {0, 0, (VkDeviceSize)n * sizeof(lrt_hit)}; + VkBufferCopy down = {0, 0, (VkDeviceSize)n * hit_bytes}; vkCmdCopyBuffer(cb, s->hits_dev.buf, s->hits_stage.buf, 1, &down); vk_buf_barrier(cb, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_HOST_BIT, VK_ACCESS_TRANSFER_WRITE_BIT, VK_ACCESS_HOST_READ_BIT); run_ok = vk_cmd_end_submit(e, cb); } - vkDestroyDescriptorPool(e->device, pool, NULL); + /* vk_cmd_end_submit may have flagged the device lost (trace timeout / loss); on a + * lost device every other site in this file skips driver calls, so match that and + * only leak the pool handle (the OS reclaims the GPU on exit). */ + if (!e->device_lost) vkDestroyDescriptorPool(e->device, pool, NULL); if (!run_ok) { if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return -1; } - if (!vk_buffer_read(e, &s->hits_stage, out, (size_t)n * sizeof(lrt_hit))) { + if (!vk_buffer_read(e, &s->hits_stage, out, (size_t)n * hit_bytes)) { if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return -1; } + const uint32_t hw = s->hit_words ? s->hit_words : 4u; + const uint32_t *w = (const uint32_t *)out; int hits = 0; for (uint32_t i = 0; i < n; i++) - if (out[i].prim_id != LRT_TRI_NO_HIT) hits++; + if (w[(size_t)i * hw + 3u] != LRT_TRI_NO_HIT) hits++; + if (err) *err = LRT_RESULT_OK; + return hits; +} + +int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, + const lrt_ray *rays, uint32_t n, lrt_hit *out, + lrt_result *err) { + if (s && s->hit_words == 5u) { /* built wide: caller must use the wide trace */ + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return -1; + } + return rtx_trace_impl(e, s, rtx_get_pipeline(e), s ? s->tlas.as : VK_NULL_HANDLE, + rays, n, out, err); +} + +int lrt_vk_rtx_scene_trace_wide(lrt_vk_engine *e, lrt_vk_rtx_scene *s, + const lrt_ray *rays, uint32_t n, lrt_hit_wide *out, + lrt_result *err) { + if (!s || s->hit_words != 5u) { /* not built wide */ + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return -1; + } + vk_pipeline *pipe = rtx_get_pipeline_wide(e); + /* Single TLAS: trace it directly into `out`. */ + if (s->ntlas <= 1u) + return rtx_trace_impl(e, s, pipe, s->tlas.as, rays, n, out, err); + + /* Multi-TLAS: trace slice 0 into `out` (its base is 0, so its instanceIds are + * already global), then trace each remaining slice into a scratch buffer and + * keep the nearer hit, offsetting the slice-local instanceId by tlas_base[k]. + * K sequential dispatches + a CPU merge -- offline throughput, no shader/ + * descriptor changes vs the single-TLAS wide path. */ + int r = rtx_trace_impl(e, s, pipe, s->tlas.as, rays, n, out, err); + if (r < 0) return -1; + lrt_hit_wide *tmp = (lrt_hit_wide *)malloc((size_t)n * sizeof(lrt_hit_wide)); + if (!tmp) { + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + return -1; + } + for (uint32_t k = 1; k < s->ntlas; k++) { + if (rtx_trace_impl(e, s, pipe, s->extra_tlas[k - 1u].as, rays, n, tmp, err) < + 0) { + free(tmp); + return -1; + } + const uint32_t base = s->tlas_base[k]; + for (uint32_t i = 0; i < n; i++) { + if (tmp[i].inst == LRT_TRI_NO_HIT) continue; + if (out[i].inst == LRT_TRI_NO_HIT || tmp[i].t < out[i].t) { + out[i] = tmp[i]; + out[i].inst += base; /* slice-local -> global instance id */ + } + } + } + free(tmp); + int hits = 0; + for (uint32_t i = 0; i < n; i++) + if (out[i].inst != LRT_TRI_NO_HIT) hits++; if (err) *err = LRT_RESULT_OK; return hits; } +uint32_t lrt_vk_rtx_scene_ntlas(const lrt_vk_rtx_scene *s) { + return s ? s->ntlas : 0u; +} + void lrt_vk_rtx_scene_free(lrt_vk_engine *e, lrt_vk_rtx_scene *s) { if (!e || !s) return; - vkDeviceWaitIdle(e->device); + if (!e->device_lost) vkDeviceWaitIdle(e->device); vk_buffer_destroy(e, &s->rays_dev); vk_buffer_destroy(e, &s->hits_dev); vk_buffer_destroy(e, &s->rays_stage); vk_buffer_destroy(e, &s->hits_stage); vk_accel_destroy(e, &s->tlas); - vk_accel_destroy(e, &s->blas); + if (s->extra_tlas) { + for (uint32_t k = 1; k < s->ntlas; k++) + vk_accel_destroy(e, &s->extra_tlas[k - 1u]); + free(s->extra_tlas); + } + free(s->tlas_base); + if (s->blas) { + /* Destroys only the AS handles; the underlying storage is pool-owned. */ + for (uint32_t c = 0; c < s->nblas; c++) vk_accel_destroy(e, &s->blas[c]); + free(s->blas); + } + if (s->blas_input_pool) { + vk_subpool_free(e, s->blas_input_pool); + free(s->blas_input_pool); + } + if (s->blas_store_pool) { + vk_subpool_free(e, s->blas_store_pool); + free(s->blas_store_pool); + } free(s); } diff --git a/lightrt_c_vk.h b/lightrt_c_vk.h index 1abc1a4..6992871 100644 --- a/lightrt_c_vk.h +++ b/lightrt_c_vk.h @@ -64,6 +64,12 @@ uint32_t lrt_vk_engine_caps(const lrt_vk_engine *e); /* Selected physical-device name (e.g. "NVIDIA GeForce RTX 3070"). */ const char *lrt_vk_engine_device_name(const lrt_vk_engine *e); +/* Largest DEVICE_LOCAL memory heap (VRAM) in bytes, queried with a throwaway + * Vulkan instance (no engine needed). prefer_discrete!=0 favors a discrete GPU. + * Returns 0 if Vulkan or a suitable device is unavailable. Lets callers size GPU + * memory budgets before building anything. */ +uint64_t lrt_vk_device_local_bytes(int prefer_discrete); + /* Human-readable message for the last failed call on this engine. */ const char *lrt_vk_engine_last_error(const lrt_vk_engine *e); @@ -151,12 +157,97 @@ lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_indexed(lrt_vk_engine *e, const uint32_t *indices, uint32_t ntris, lrt_result *err); +/* --- True two-level (instanced) scene ------------------------------------- + * + * Build a genuine two-level acceleration structure: one BLAS per PROTOTYPE + * (geometry stored on the device ONCE) and one TLAS instance per PLACEMENT (the + * same prototype BLAS referenced under a per-instance transform). This is the + * memory-sharing path the flat builders above cannot express — N copies of a + * prototype cost one BLAS, not N. It reuses the SAME trace pipeline/shader as + * the flat builders (no shader change): the shader already recovers the hit id + * as instanceId*tri_chunk + primitiveIndex, so here tri_chunk is set to the max + * prototype triangle count and the returned lrt_hit.prim_id decodes as: + * instance = prim_id / (*out_tri_stride) + * prototypeLocalTri = prim_id % (*out_tri_stride) + * The caller maps `instance` -> prototype (via its own instance list) and shades + * the prototype-local triangle, transforming object-space attributes by that + * instance's transform. + * + * Each prototype builds as a SINGLE BLAS (this API does no chunk splitting), so a + * prototype larger than the device's one-BLAS build limit must be pre-split by the + * caller into several smaller prototypes (each a triangle slice) sharing the same + * per-instance transform — the encoding is unaffected. Fails (returns NULL, + * LRT_RESULT_INVALID_ARGUMENT) if ninsts*maxPrototypeTris would overflow the 32-bit + * prim_id encoding — the caller should fall back to the flat builder. */ +typedef struct lrt_vk_proto { + const float *vertices; /* 3*nverts floats of unique positions */ + uint32_t nverts; + const uint32_t *indices; /* 3*ntris vertex ids (uint32), prototype-local */ + uint32_t ntris; +} lrt_vk_proto; + +typedef struct lrt_vk_instance { + float transform[12]; /* object->world 3x4 row-major (world = M*[p;1]) */ + uint32_t proto; /* index into the protos[] array */ +} lrt_vk_instance; + +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, uint32_t *out_tri_stride, + lrt_result *err); + +/* Wide-id instanced build: same inputs as lrt_vk_rtx_scene_build_instanced, but + * the hit id is NOT packed into a single 32-bit prim_id. Instead the trace stores + * the TLAS instanceId and the prototype-local triangle index in SEPARATE 32-bit + * words (lrt_hit_wide below), so there is no ninsts*maxPrototypeTris product to + * overflow -- the only remaining ceiling is the device TLAS maxInstanceCount + * (commonly 2^24). Use this for instanced scenes the narrow builder rejects + * (Moana-island scale). A scene built this way MUST be traced with + * lrt_vk_rtx_scene_trace_wide (the narrow trace returns -1 on it). */ +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced_wide( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, lrt_result *err); + +/* Multi-TLAS wide instanced build: as lrt_vk_rtx_scene_build_instanced_wide, but + * splits the instances into ceil(ninsts / ~16M) TLAS slices, each its own TLAS + * over the SAME shared BLAS set, so a scene with MORE than the device TLAS + * maxInstanceCount (2^24) instances renders IN FULL (Moana island's ~42.8M + * instances -> 3 TLASes). Trace with lrt_vk_rtx_scene_trace_wide: it traces the + * slices sequentially and merges the nearest hit on the host, reporting GLOBAL + * instanceIds (0..ninsts-1). Costs K sequential dispatches per trace + K TLAS + * instance buffers of VRAM; the BLAS is stored once regardless of K. */ +lrt_vk_rtx_scene *lrt_vk_rtx_scene_build_instanced_multi( + lrt_vk_engine *e, const lrt_vk_proto *protos, uint32_t nprotos, + const lrt_vk_instance *insts, uint32_t ninsts, lrt_result *err); + /* Trace n rays against the resident AS. Returns #rays that hit, or -1 on error. - * Trace buffers are reused/grown across calls. */ + * Trace buffers are reused/grown across calls. Rejects (-1) a scene built with + * the wide builder -- use lrt_vk_rtx_scene_trace_wide for those. */ int lrt_vk_rtx_scene_trace(lrt_vk_engine *e, lrt_vk_rtx_scene *s, const lrt_ray *rays, uint32_t n, lrt_hit *out, lrt_result *err); +/* Wide-id hit: like lrt_hit but the single 32-bit prim_id is replaced by the two + * fields the wide trace stores separately. `inst` is the TLAS instance (placement) + * id 0..ninsts-1; `local` is the prototype-local triangle index. A miss sets + * inst == LRT_TRI_NO_HIT (as lrt_hit.prim_id does). The caller maps inst -> + * prototype via its own instance list and shades prototype triangle `local`. */ +typedef struct lrt_hit_wide { + float t, u, v; + uint32_t inst; + uint32_t local; +} lrt_hit_wide; + +/* As lrt_vk_rtx_scene_trace, for a scene built with lrt_vk_rtx_scene_build_instanced_wide. + * Writes n lrt_hit_wide. Rejects (-1) a narrow scene. */ +int lrt_vk_rtx_scene_trace_wide(lrt_vk_engine *e, lrt_vk_rtx_scene *s, + const lrt_ray *rays, uint32_t n, lrt_hit_wide *out, + lrt_result *err); + +/* Number of TLAS slices the scene was built with (1 for a single-TLAS scene, >1 + * for a multi-TLAS scene). 0 if s is NULL. Informational (logging / tests). */ +uint32_t lrt_vk_rtx_scene_ntlas(const lrt_vk_rtx_scene *s); + void lrt_vk_rtx_scene_free(lrt_vk_engine *e, lrt_vk_rtx_scene *s); /* --- Analytic-primitive GPU shading (spheres + boxes) --------------------- diff --git a/lightrt_vkew.h b/lightrt_vkew.h index 28045e1..975c624 100644 --- a/lightrt_vkew.h +++ b/lightrt_vkew.h @@ -1224,13 +1224,13 @@ VK_DEFINE_NON_DISPATCHABLE_HANDLE(VkAccelerationStructureKHR) #define VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_BUFFER_DEVICE_ADDRESS_FEATURES 1000257000 #define VK_STRUCTURE_TYPE_BUFFER_DEVICE_ADDRESS_INFO 1000244001 #define VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET_ACCELERATION_STRUCTURE_KHR 1000150007 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_TRIANGLES_DATA_KHR 1000150000 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_AABBS_DATA_KHR 1000150001 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_INSTANCES_DATA_KHR 1000150002 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR 1000150003 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_GEOMETRY_INFO_KHR 1000150004 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_CREATE_INFO_KHR 1000150005 -#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_DEVICE_ADDRESS_INFO_KHR 1000150008 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_GEOMETRY_INFO_KHR 1000150000 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_DEVICE_ADDRESS_INFO_KHR 1000150002 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_AABBS_DATA_KHR 1000150003 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_INSTANCES_DATA_KHR 1000150004 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_TRIANGLES_DATA_KHR 1000150005 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR 1000150006 +#define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_CREATE_INFO_KHR 1000150017 #define VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_SIZES_INFO_KHR 1000150020 #define VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ACCELERATION_STRUCTURE_FEATURES_KHR 1000150013 #define VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_QUERY_FEATURES_KHR 1000348013 @@ -1240,6 +1240,11 @@ VK_DEFINE_NON_DISPATCHABLE_HANDLE(VkAccelerationStructureKHR) #define VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_STORAGE_BIT_KHR 0x00100000 #define VK_BUFFER_USAGE_SHADER_BINDING_TABLE_BIT_KHR 0x00000400 +// New pipeline/access flags for acceleration-structure build synchronization. +#define VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR 0x02000000 +#define VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR 0x00200000 +#define VK_ACCESS_ACCELERATION_STRUCTURE_WRITE_BIT_KHR 0x00400000 + // New VkDescriptorType (for binding a TLAS into a ray_query descriptor set). #define VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR 1000150000 diff --git a/scripts/compile_shaders.sh b/scripts/compile_shaders.sh index 6a79a59..a8f9537 100755 --- a/scripts/compile_shaders.sh +++ b/scripts/compile_shaders.sh @@ -33,17 +33,23 @@ if ! command -v "${GLSLANG}" >/dev/null 2>&1; then exit 1 fi -# gen +# gen [glslang-args...] +# The source basename is separate so one .comp can be compiled into several headers +# under different -D defines (e.g. the wide-id ray-query variant below). gen() { - local base="$1" sym="$2" - "${GLSLANG}" -V --target-env vulkan1.2 --vn "${sym}" \ - -o "${SH}/${base}.spv.h" "${SH}/${base}.comp" - echo "generated ${SH}/${base}.spv.h (${sym})" + local out="$1" sym="$2" src="$3" + shift 3 + "${GLSLANG}" -V --target-env vulkan1.2 --vn "${sym}" "$@" \ + -o "${SH}/${out}.spv.h" "${SH}/${src}.comp" + echo "generated ${SH}/${out}.spv.h (${sym})" } -gen trace_bvh trace_bvh_spv -gen build_morton build_morton_spv -gen trace_ray_query trace_ray_query_spv # needs VK_KHR_ray_query (SPIR-V 1.4) -gen shade_analytic shade_analytic_spv # analytic sphere/box forward shading +gen trace_bvh trace_bvh_spv trace_bvh +gen build_morton build_morton_spv build_morton +gen trace_ray_query trace_ray_query_spv trace_ray_query # needs VK_KHR_ray_query (SPIR-V 1.4) +# Wide-id variant of the ray-query trace: 5 words/hit with instanceId + prim index +# stored separately, so instanced scenes past the 32-bit prim_id product still trace. +gen trace_ray_query_wide trace_ray_query_wide_spv trace_ray_query -DLRT_WIDE_ID +gen shade_analytic shade_analytic_spv shade_analytic # analytic sphere/box forward shading echo "done." diff --git a/vk/shaders/trace_ray_query.comp b/vk/shaders/trace_ray_query.comp index f94026e..1c3addd 100644 --- a/vk/shaders/trace_ray_query.comp +++ b/vk/shaders/trace_ray_query.comp @@ -9,7 +9,18 @@ // read directly. Writes the same lrt_hit layout (t,u,v,prim_id) as the software // trace shader: t is in units of |dir|, (u,v) are the hit barycentrics (which // coincide with Moller-Trumbore u,v), prim_id = triangle index (single geometry, -// no index buffer => primitive index == build order == caller index). +// no index buffer => primitive index == build order == caller index). The +// geometry is split across multiple chunk BLASes (one TLAS instance each, in +// chunk order), so the global triangle id is recovered as +// instanceId*tri_chunk + primitiveIndex. +// +// Two variants are compiled from this one source: +// - default (narrow): 4 words/hit {t,u,v, prim_id=instanceId*tri_chunk+prim}. +// prim_id must fit 32 bits, so ninsts*tri_chunk < 2^32. +// - LRT_WIDE_ID (wide): 5 words/hit {t,u,v, instanceId, primitiveIndex} stored +// SEPARATELY -- no stride multiply, so there is no 32-bit product to overflow. +// Used for instanced scenes whose ninsts*maxPrototypeTris exceeds 32 bits +// (e.g. Moana island); the only remaining ceiling is the TLAS maxInstanceCount. // // SPDX-License-Identifier: Apache-2.0 // @@ -17,10 +28,15 @@ layout(local_size_x = 64) in; layout(set = 0, binding = 0) uniform accelerationStructureEXT tlas; layout(std430, binding = 1) readonly buffer Rays { uint rays[]; }; // 8 words / ray +#ifdef LRT_WIDE_ID +layout(std430, binding = 2) writeonly buffer Hits { uint hits[]; }; // 5 words / hit +#else layout(std430, binding = 2) writeonly buffer Hits { uint hits[]; }; // 4 words / hit +#endif layout(push_constant) uniform PC { uint ray_count; + uint tri_chunk; // triangles per BLAS chunk; global prim_id = instanceId*tri_chunk + primIndex } pc; const uint NO_HIT = 0xFFFFFFFFu; @@ -38,15 +54,46 @@ void main() { float tmax = fu(rays[rb + 7u]); rayQueryEXT rq; - rayQueryInitializeEXT(rq, tlas, gl_RayFlagsOpaqueEXT, 0xFFu, org, tmin, dir, tmax); + // Primary visibility: closest-hit. All geometry is opaque, so the hardware + // commits the nearest intersection during traversal and we read it after the + // proceed loop. Do NOT set TerminateOnFirstHit here — that commits the first + // triangle traversal happens to find, not the closest (correct only for + // shadow/occlusion rays). Cull mask must be 0xFF to match the TLAS instance + // mask (0xFF); a 0x0 mask AND-culls every instance and returns no hits. + rayQueryInitializeEXT(rq, tlas, + gl_RayFlagsOpaqueEXT, + 0xFFu, org, tmin, dir, tmax); while (rayQueryProceedEXT(rq)) {} +#ifdef LRT_WIDE_ID + uint hb = gid * 5u; + if (rayQueryGetIntersectionTypeEXT(rq, true) == + gl_RayQueryCommittedIntersectionTriangleEXT) { + float t = rayQueryGetIntersectionTEXT(rq, true); + vec2 bc = rayQueryGetIntersectionBarycentricsEXT(rq, true); + uint prim = uint(rayQueryGetIntersectionPrimitiveIndexEXT(rq, true)); + uint inst = uint(rayQueryGetIntersectionInstanceIdEXT(rq, true)); + hits[hb + 0u] = floatBitsToUint(t); + hits[hb + 1u] = floatBitsToUint(bc.x); + hits[hb + 2u] = floatBitsToUint(bc.y); + hits[hb + 3u] = inst; // TLAS instance (placement) id, 0..ninsts-1 + hits[hb + 4u] = prim; // prototype-local triangle index + } else { + hits[hb + 0u] = floatBitsToUint(0.0); + hits[hb + 1u] = 0u; + hits[hb + 2u] = 0u; + hits[hb + 3u] = NO_HIT; + hits[hb + 4u] = 0u; + } +#else uint hb = gid * 4u; if (rayQueryGetIntersectionTypeEXT(rq, true) == gl_RayQueryCommittedIntersectionTriangleEXT) { float t = rayQueryGetIntersectionTEXT(rq, true); vec2 bc = rayQueryGetIntersectionBarycentricsEXT(rq, true); - uint pid = uint(rayQueryGetIntersectionPrimitiveIndexEXT(rq, true)); + uint prim = uint(rayQueryGetIntersectionPrimitiveIndexEXT(rq, true)); + uint inst = uint(rayQueryGetIntersectionInstanceIdEXT(rq, true)); + uint pid = inst * pc.tri_chunk + prim; hits[hb + 0u] = floatBitsToUint(t); hits[hb + 1u] = floatBitsToUint(bc.x); hits[hb + 2u] = floatBitsToUint(bc.y); @@ -57,4 +104,5 @@ void main() { hits[hb + 2u] = 0u; hits[hb + 3u] = NO_HIT; } +#endif } diff --git a/vk/shaders/trace_ray_query.spv.h b/vk/shaders/trace_ray_query.spv.h index 4b2d620..7046ff6 100644 --- a/vk/shaders/trace_ray_query.spv.h +++ b/vk/shaders/trace_ray_query.spv.h @@ -1,149 +1,314 @@ - // 1114.3.0 - #pragma once -const uint32_t trace_ray_query_spv[] = { - 0x07230203,0x00010500,0x0008000b,0x000000bf,0x00000000,0x00020011,0x00000001,0x00020011, - 0x00001178,0x0006000a,0x5f565053,0x5f52484b,0x5f796172,0x72657571,0x00000079,0x0006000b, - 0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,0x00000000,0x00000001, - 0x000b000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x00000014,0x0000001c,0x00000031, - 0x00000072,0x00000075,0x00000097,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001, - 0x00000001,0x00030003,0x00000002,0x000001cc,0x00060004,0x455f4c47,0x725f5458,0x715f7961, - 0x79726575,0x00000000,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00040005,0x0000000b, - 0x75287566,0x00003b31,0x00030005,0x0000000a,0x00000069,0x00030005,0x00000011,0x00646967, - 0x00080005,0x00000014,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,0x496e6f69,0x00000044, - 0x00030005,0x0000001a,0x00004350,0x00060006,0x0000001a,0x00000000,0x5f796172,0x6e756f63, - 0x00000074,0x00030005,0x0000001c,0x00006370,0x00030005,0x00000027,0x00006272,0x00030005, - 0x0000002d,0x0067726f,0x00040005,0x0000002f,0x73796152,0x00000000,0x00050006,0x0000002f, - 0x00000000,0x73796172,0x00000000,0x00030005,0x00000031,0x00000000,0x00040005,0x00000034, - 0x61726170,0x0000006d,0x00040005,0x0000003c,0x61726170,0x0000006d,0x00040005,0x00000043, - 0x61726170,0x0000006d,0x00040005,0x00000049,0x6e696d74,0x00000000,0x00040005,0x0000004d, - 0x61726170,0x0000006d,0x00030005,0x00000051,0x00726964,0x00040005,0x00000055,0x61726170, - 0x0000006d,0x00040005,0x0000005c,0x61726170,0x0000006d,0x00040005,0x00000063,0x61726170, - 0x0000006d,0x00040005,0x00000068,0x78616d74,0x00000000,0x00040005,0x0000006c,0x61726170, - 0x0000006d,0x00030005,0x00000072,0x00007172,0x00040005,0x00000075,0x73616c74,0x00000000, - 0x00030005,0x00000082,0x00006268,0x00030005,0x0000008b,0x00000074,0x00030005,0x0000008f, - 0x00006362,0x00030005,0x00000091,0x00646970,0x00040005,0x00000095,0x73746948,0x00000000, - 0x00050006,0x00000095,0x00000000,0x73746968,0x00000000,0x00030005,0x00000097,0x00000000, - 0x00040047,0x00000014,0x0000000b,0x0000001c,0x00050048,0x0000001a,0x00000000,0x00000023, - 0x00000000,0x00030047,0x0000001a,0x00000002,0x00040047,0x0000002e,0x00000006,0x00000004, - 0x00040048,0x0000002f,0x00000000,0x00000018,0x00050048,0x0000002f,0x00000000,0x00000023, - 0x00000000,0x00030047,0x0000002f,0x00000002,0x00040047,0x00000031,0x00000022,0x00000000, - 0x00040047,0x00000031,0x00000021,0x00000001,0x00040047,0x00000075,0x00000022,0x00000000, - 0x00040047,0x00000075,0x00000021,0x00000000,0x00040047,0x00000094,0x00000006,0x00000004, - 0x00040048,0x00000095,0x00000000,0x00000019,0x00050048,0x00000095,0x00000000,0x00000023, - 0x00000000,0x00030047,0x00000095,0x00000002,0x00040047,0x00000097,0x00000022,0x00000000, - 0x00040047,0x00000097,0x00000021,0x00000002,0x00040047,0x000000be,0x0000000b,0x00000019, - 0x00020013,0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020, - 0x00000000,0x00040020,0x00000007,0x00000007,0x00000006,0x00030016,0x00000008,0x00000020, - 0x00040021,0x00000009,0x00000008,0x00000007,0x00040017,0x00000012,0x00000006,0x00000003, - 0x00040020,0x00000013,0x00000001,0x00000012,0x0004003b,0x00000013,0x00000014,0x00000001, - 0x0004002b,0x00000006,0x00000015,0x00000000,0x00040020,0x00000016,0x00000001,0x00000006, - 0x0003001e,0x0000001a,0x00000006,0x00040020,0x0000001b,0x00000009,0x0000001a,0x0004003b, - 0x0000001b,0x0000001c,0x00000009,0x00040015,0x0000001d,0x00000020,0x00000001,0x0004002b, - 0x0000001d,0x0000001e,0x00000000,0x00040020,0x0000001f,0x00000009,0x00000006,0x00020014, - 0x00000022,0x0004002b,0x00000006,0x00000029,0x00000008,0x00040017,0x0000002b,0x00000008, - 0x00000003,0x00040020,0x0000002c,0x00000007,0x0000002b,0x0003001d,0x0000002e,0x00000006, - 0x0003001e,0x0000002f,0x0000002e,0x00040020,0x00000030,0x0000000c,0x0000002f,0x0004003b, - 0x00000030,0x00000031,0x0000000c,0x00040020,0x00000035,0x0000000c,0x00000006,0x0004002b, - 0x00000006,0x0000003a,0x00000001,0x0004002b,0x00000006,0x00000041,0x00000002,0x00040020, - 0x00000048,0x00000007,0x00000008,0x0004002b,0x00000006,0x0000004b,0x00000003,0x0004002b, - 0x00000006,0x00000053,0x00000004,0x0004002b,0x00000006,0x0000005a,0x00000005,0x0004002b, - 0x00000006,0x00000061,0x00000006,0x0004002b,0x00000006,0x0000006a,0x00000007,0x00021178, - 0x00000070,0x00040020,0x00000071,0x00000006,0x00000070,0x0004003b,0x00000071,0x00000072, - 0x00000006,0x000214dd,0x00000073,0x00040020,0x00000074,0x00000000,0x00000073,0x0004003b, - 0x00000074,0x00000075,0x00000000,0x0004002b,0x00000006,0x00000077,0x000000ff,0x00030029, - 0x00000022,0x00000085,0x0004002b,0x0000001d,0x00000086,0x00000001,0x00040017,0x0000008d, - 0x00000008,0x00000002,0x00040020,0x0000008e,0x00000007,0x0000008d,0x0003001d,0x00000094, - 0x00000006,0x0003001e,0x00000095,0x00000094,0x00040020,0x00000096,0x0000000c,0x00000095, - 0x0004003b,0x00000096,0x00000097,0x0000000c,0x0004002b,0x00000008,0x000000b0,0x00000000, - 0x0004002b,0x00000006,0x000000bb,0xffffffff,0x0004002b,0x00000006,0x000000bd,0x00000040, - 0x0006002c,0x00000012,0x000000be,0x000000bd,0x0000003a,0x0000003a,0x00050036,0x00000002, - 0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000007,0x00000011, - 0x00000007,0x0004003b,0x00000007,0x00000027,0x00000007,0x0004003b,0x0000002c,0x0000002d, - 0x00000007,0x0004003b,0x00000007,0x00000034,0x00000007,0x0004003b,0x00000007,0x0000003c, - 0x00000007,0x0004003b,0x00000007,0x00000043,0x00000007,0x0004003b,0x00000048,0x00000049, - 0x00000007,0x0004003b,0x00000007,0x0000004d,0x00000007,0x0004003b,0x0000002c,0x00000051, - 0x00000007,0x0004003b,0x00000007,0x00000055,0x00000007,0x0004003b,0x00000007,0x0000005c, - 0x00000007,0x0004003b,0x00000007,0x00000063,0x00000007,0x0004003b,0x00000048,0x00000068, - 0x00000007,0x0004003b,0x00000007,0x0000006c,0x00000007,0x0004003b,0x00000007,0x00000082, - 0x00000007,0x0004003b,0x00000048,0x0000008b,0x00000007,0x0004003b,0x0000008e,0x0000008f, - 0x00000007,0x0004003b,0x00000007,0x00000091,0x00000007,0x00050041,0x00000016,0x00000017, - 0x00000014,0x00000015,0x0004003d,0x00000006,0x00000018,0x00000017,0x0003003e,0x00000011, - 0x00000018,0x0004003d,0x00000006,0x00000019,0x00000011,0x00050041,0x0000001f,0x00000020, - 0x0000001c,0x0000001e,0x0004003d,0x00000006,0x00000021,0x00000020,0x000500ae,0x00000022, - 0x00000023,0x00000019,0x00000021,0x000300f7,0x00000025,0x00000000,0x000400fa,0x00000023, - 0x00000024,0x00000025,0x000200f8,0x00000024,0x000100fd,0x000200f8,0x00000025,0x0004003d, - 0x00000006,0x00000028,0x00000011,0x00050084,0x00000006,0x0000002a,0x00000028,0x00000029, - 0x0003003e,0x00000027,0x0000002a,0x0004003d,0x00000006,0x00000032,0x00000027,0x00050080, - 0x00000006,0x00000033,0x00000032,0x00000015,0x00060041,0x00000035,0x00000036,0x00000031, - 0x0000001e,0x00000033,0x0004003d,0x00000006,0x00000037,0x00000036,0x0003003e,0x00000034, - 0x00000037,0x00050039,0x00000008,0x00000038,0x0000000b,0x00000034,0x0004003d,0x00000006, - 0x00000039,0x00000027,0x00050080,0x00000006,0x0000003b,0x00000039,0x0000003a,0x00060041, - 0x00000035,0x0000003d,0x00000031,0x0000001e,0x0000003b,0x0004003d,0x00000006,0x0000003e, - 0x0000003d,0x0003003e,0x0000003c,0x0000003e,0x00050039,0x00000008,0x0000003f,0x0000000b, - 0x0000003c,0x0004003d,0x00000006,0x00000040,0x00000027,0x00050080,0x00000006,0x00000042, - 0x00000040,0x00000041,0x00060041,0x00000035,0x00000044,0x00000031,0x0000001e,0x00000042, - 0x0004003d,0x00000006,0x00000045,0x00000044,0x0003003e,0x00000043,0x00000045,0x00050039, - 0x00000008,0x00000046,0x0000000b,0x00000043,0x00060050,0x0000002b,0x00000047,0x00000038, - 0x0000003f,0x00000046,0x0003003e,0x0000002d,0x00000047,0x0004003d,0x00000006,0x0000004a, - 0x00000027,0x00050080,0x00000006,0x0000004c,0x0000004a,0x0000004b,0x00060041,0x00000035, - 0x0000004e,0x00000031,0x0000001e,0x0000004c,0x0004003d,0x00000006,0x0000004f,0x0000004e, - 0x0003003e,0x0000004d,0x0000004f,0x00050039,0x00000008,0x00000050,0x0000000b,0x0000004d, - 0x0003003e,0x00000049,0x00000050,0x0004003d,0x00000006,0x00000052,0x00000027,0x00050080, - 0x00000006,0x00000054,0x00000052,0x00000053,0x00060041,0x00000035,0x00000056,0x00000031, - 0x0000001e,0x00000054,0x0004003d,0x00000006,0x00000057,0x00000056,0x0003003e,0x00000055, - 0x00000057,0x00050039,0x00000008,0x00000058,0x0000000b,0x00000055,0x0004003d,0x00000006, - 0x00000059,0x00000027,0x00050080,0x00000006,0x0000005b,0x00000059,0x0000005a,0x00060041, - 0x00000035,0x0000005d,0x00000031,0x0000001e,0x0000005b,0x0004003d,0x00000006,0x0000005e, - 0x0000005d,0x0003003e,0x0000005c,0x0000005e,0x00050039,0x00000008,0x0000005f,0x0000000b, - 0x0000005c,0x0004003d,0x00000006,0x00000060,0x00000027,0x00050080,0x00000006,0x00000062, - 0x00000060,0x00000061,0x00060041,0x00000035,0x00000064,0x00000031,0x0000001e,0x00000062, - 0x0004003d,0x00000006,0x00000065,0x00000064,0x0003003e,0x00000063,0x00000065,0x00050039, - 0x00000008,0x00000066,0x0000000b,0x00000063,0x00060050,0x0000002b,0x00000067,0x00000058, - 0x0000005f,0x00000066,0x0003003e,0x00000051,0x00000067,0x0004003d,0x00000006,0x00000069, - 0x00000027,0x00050080,0x00000006,0x0000006b,0x00000069,0x0000006a,0x00060041,0x00000035, - 0x0000006d,0x00000031,0x0000001e,0x0000006b,0x0004003d,0x00000006,0x0000006e,0x0000006d, - 0x0003003e,0x0000006c,0x0000006e,0x00050039,0x00000008,0x0000006f,0x0000000b,0x0000006c, - 0x0003003e,0x00000068,0x0000006f,0x0004003d,0x00000073,0x00000076,0x00000075,0x0004003d, - 0x0000002b,0x00000078,0x0000002d,0x0004003d,0x00000008,0x00000079,0x00000049,0x0004003d, - 0x0000002b,0x0000007a,0x00000051,0x0004003d,0x00000008,0x0000007b,0x00000068,0x00091179, - 0x00000072,0x00000076,0x0000003a,0x00000077,0x00000078,0x00000079,0x0000007a,0x0000007b, - 0x000200f9,0x0000007c,0x000200f8,0x0000007c,0x000400f6,0x0000007e,0x0000007f,0x00000000, - 0x000200f9,0x00000080,0x000200f8,0x00000080,0x0004117d,0x00000022,0x00000081,0x00000072, - 0x000400fa,0x00000081,0x0000007d,0x0000007e,0x000200f8,0x0000007d,0x000200f9,0x0000007f, - 0x000200f8,0x0000007f,0x000200f9,0x0000007c,0x000200f8,0x0000007e,0x0004003d,0x00000006, - 0x00000083,0x00000011,0x00050084,0x00000006,0x00000084,0x00000083,0x00000053,0x0003003e, - 0x00000082,0x00000084,0x0005117f,0x00000006,0x00000087,0x00000072,0x00000086,0x000500aa, - 0x00000022,0x00000088,0x00000087,0x0000003a,0x000300f7,0x0000008a,0x00000000,0x000400fa, - 0x00000088,0x00000089,0x000000ad,0x000200f8,0x00000089,0x00051782,0x00000008,0x0000008c, - 0x00000072,0x00000086,0x0003003e,0x0000008b,0x0000008c,0x00051788,0x0000008d,0x00000090, - 0x00000072,0x00000086,0x0003003e,0x0000008f,0x00000090,0x00051787,0x0000001d,0x00000092, - 0x00000072,0x00000086,0x0004007c,0x00000006,0x00000093,0x00000092,0x0003003e,0x00000091, - 0x00000093,0x0004003d,0x00000006,0x00000098,0x00000082,0x00050080,0x00000006,0x00000099, - 0x00000098,0x00000015,0x0004003d,0x00000008,0x0000009a,0x0000008b,0x0004007c,0x00000006, - 0x0000009b,0x0000009a,0x00060041,0x00000035,0x0000009c,0x00000097,0x0000001e,0x00000099, - 0x0003003e,0x0000009c,0x0000009b,0x0004003d,0x00000006,0x0000009d,0x00000082,0x00050080, - 0x00000006,0x0000009e,0x0000009d,0x0000003a,0x00050041,0x00000048,0x0000009f,0x0000008f, - 0x00000015,0x0004003d,0x00000008,0x000000a0,0x0000009f,0x0004007c,0x00000006,0x000000a1, - 0x000000a0,0x00060041,0x00000035,0x000000a2,0x00000097,0x0000001e,0x0000009e,0x0003003e, - 0x000000a2,0x000000a1,0x0004003d,0x00000006,0x000000a3,0x00000082,0x00050080,0x00000006, - 0x000000a4,0x000000a3,0x00000041,0x00050041,0x00000048,0x000000a5,0x0000008f,0x0000003a, - 0x0004003d,0x00000008,0x000000a6,0x000000a5,0x0004007c,0x00000006,0x000000a7,0x000000a6, - 0x00060041,0x00000035,0x000000a8,0x00000097,0x0000001e,0x000000a4,0x0003003e,0x000000a8, - 0x000000a7,0x0004003d,0x00000006,0x000000a9,0x00000082,0x00050080,0x00000006,0x000000aa, - 0x000000a9,0x0000004b,0x0004003d,0x00000006,0x000000ab,0x00000091,0x00060041,0x00000035, - 0x000000ac,0x00000097,0x0000001e,0x000000aa,0x0003003e,0x000000ac,0x000000ab,0x000200f9, - 0x0000008a,0x000200f8,0x000000ad,0x0004003d,0x00000006,0x000000ae,0x00000082,0x00050080, - 0x00000006,0x000000af,0x000000ae,0x00000015,0x0004007c,0x00000006,0x000000b1,0x000000b0, - 0x00060041,0x00000035,0x000000b2,0x00000097,0x0000001e,0x000000af,0x0003003e,0x000000b2, - 0x000000b1,0x0004003d,0x00000006,0x000000b3,0x00000082,0x00050080,0x00000006,0x000000b4, - 0x000000b3,0x0000003a,0x00060041,0x00000035,0x000000b5,0x00000097,0x0000001e,0x000000b4, - 0x0003003e,0x000000b5,0x00000015,0x0004003d,0x00000006,0x000000b6,0x00000082,0x00050080, - 0x00000006,0x000000b7,0x000000b6,0x00000041,0x00060041,0x00000035,0x000000b8,0x00000097, - 0x0000001e,0x000000b7,0x0003003e,0x000000b8,0x00000015,0x0004003d,0x00000006,0x000000b9, - 0x00000082,0x00050080,0x00000006,0x000000ba,0x000000b9,0x0000004b,0x00060041,0x00000035, - 0x000000bc,0x00000097,0x0000001e,0x000000ba,0x0003003e,0x000000bc,0x000000bb,0x000200f9, - 0x0000008a,0x000200f8,0x0000008a,0x000100fd,0x00010038,0x00050036,0x00000008,0x0000000b, - 0x00000000,0x00000009,0x00030037,0x00000007,0x0000000a,0x000200f8,0x0000000c,0x0004003d, - 0x00000006,0x0000000d,0x0000000a,0x0004007c,0x00000008,0x0000000e,0x0000000d,0x000200fe, - 0x0000000e,0x00010038 -}; +#pragma once +const uint32_t trace_ray_query_spv[] = +{0x07230203,0x00010500,0x000d000b,0x000000c9, +0x00000000,0x00020011,0x00000001,0x00020011, +0x00001178,0x0006000a,0x5f565053,0x5f52484b, +0x5f796172,0x72657571,0x00000079,0x0006000b, +0x00000001,0x4c534c47,0x6474732e,0x3035342e, +0x00000000,0x0003000e,0x00000000,0x00000001, +0x000b000f,0x00000005,0x00000004,0x6e69616d, +0x00000000,0x00000014,0x0000001c,0x00000031, +0x00000072,0x00000075,0x000000a1,0x00060010, +0x00000004,0x00000011,0x00000040,0x00000001, +0x00000001,0x00030003,0x00000002,0x000001cc, +0x00060004,0x455f4c47,0x725f5458,0x715f7961, +0x79726575,0x00000000,0x000a0004,0x475f4c47, +0x4c474f4f,0x70635f45,0x74735f70,0x5f656c79, +0x656e696c,0x7269645f,0x69746365,0x00006576, +0x00080004,0x475f4c47,0x4c474f4f,0x6e695f45, +0x64756c63,0x69645f65,0x74636572,0x00657669, +0x00040005,0x00000004,0x6e69616d,0x00000000, +0x00040005,0x0000000b,0x75287566,0x00003b31, +0x00030005,0x0000000a,0x00000069,0x00030005, +0x00000011,0x00646967,0x00080005,0x00000014, +0x475f6c67,0x61626f6c,0x766e496c,0x7461636f, +0x496e6f69,0x00000044,0x00030005,0x0000001a, +0x00004350,0x00060006,0x0000001a,0x00000000, +0x5f796172,0x6e756f63,0x00000074,0x00060006, +0x0000001a,0x00000001,0x5f697274,0x6e756863, +0x0000006b,0x00030005,0x0000001c,0x00006370, +0x00030005,0x00000027,0x00006272,0x00030005, +0x0000002d,0x0067726f,0x00040005,0x0000002f, +0x73796152,0x00000000,0x00050006,0x0000002f, +0x00000000,0x73796172,0x00000000,0x00030005, +0x00000031,0x00000000,0x00040005,0x00000034, +0x61726170,0x0000006d,0x00040005,0x0000003c, +0x61726170,0x0000006d,0x00040005,0x00000043, +0x61726170,0x0000006d,0x00040005,0x00000049, +0x6e696d74,0x00000000,0x00040005,0x0000004d, +0x61726170,0x0000006d,0x00030005,0x00000051, +0x00726964,0x00040005,0x00000055,0x61726170, +0x0000006d,0x00040005,0x0000005c,0x61726170, +0x0000006d,0x00040005,0x00000063,0x61726170, +0x0000006d,0x00040005,0x00000068,0x78616d74, +0x00000000,0x00040005,0x0000006c,0x61726170, +0x0000006d,0x00030005,0x00000072,0x00007172, +0x00040005,0x00000075,0x73616c74,0x00000000, +0x00030005,0x00000082,0x00006268,0x00030005, +0x0000008b,0x00000074,0x00030005,0x0000008f, +0x00006362,0x00040005,0x00000091,0x6d697270, +0x00000000,0x00040005,0x00000094,0x74736e69, +0x00000000,0x00030005,0x00000097,0x00646970, +0x00040005,0x0000009f,0x73746948,0x00000000, +0x00050006,0x0000009f,0x00000000,0x73746968, +0x00000000,0x00030005,0x000000a1,0x00000000, +0x00040047,0x00000014,0x0000000b,0x0000001c, +0x00050048,0x0000001a,0x00000000,0x00000023, +0x00000000,0x00050048,0x0000001a,0x00000001, +0x00000023,0x00000004,0x00030047,0x0000001a, +0x00000002,0x00040047,0x0000002e,0x00000006, +0x00000004,0x00040048,0x0000002f,0x00000000, +0x00000018,0x00050048,0x0000002f,0x00000000, +0x00000023,0x00000000,0x00030047,0x0000002f, +0x00000002,0x00040047,0x00000031,0x00000022, +0x00000000,0x00040047,0x00000031,0x00000021, +0x00000001,0x00040047,0x00000075,0x00000022, +0x00000000,0x00040047,0x00000075,0x00000021, +0x00000000,0x00040047,0x0000009e,0x00000006, +0x00000004,0x00040048,0x0000009f,0x00000000, +0x00000019,0x00050048,0x0000009f,0x00000000, +0x00000023,0x00000000,0x00030047,0x0000009f, +0x00000002,0x00040047,0x000000a1,0x00000022, +0x00000000,0x00040047,0x000000a1,0x00000021, +0x00000002,0x00040047,0x000000c8,0x0000000b, +0x00000019,0x00020013,0x00000002,0x00030021, +0x00000003,0x00000002,0x00040015,0x00000006, +0x00000020,0x00000000,0x00040020,0x00000007, +0x00000007,0x00000006,0x00030016,0x00000008, +0x00000020,0x00040021,0x00000009,0x00000008, +0x00000007,0x00040017,0x00000012,0x00000006, +0x00000003,0x00040020,0x00000013,0x00000001, +0x00000012,0x0004003b,0x00000013,0x00000014, +0x00000001,0x0004002b,0x00000006,0x00000015, +0x00000000,0x00040020,0x00000016,0x00000001, +0x00000006,0x0004001e,0x0000001a,0x00000006, +0x00000006,0x00040020,0x0000001b,0x00000009, +0x0000001a,0x0004003b,0x0000001b,0x0000001c, +0x00000009,0x00040015,0x0000001d,0x00000020, +0x00000001,0x0004002b,0x0000001d,0x0000001e, +0x00000000,0x00040020,0x0000001f,0x00000009, +0x00000006,0x00020014,0x00000022,0x0004002b, +0x00000006,0x00000029,0x00000008,0x00040017, +0x0000002b,0x00000008,0x00000003,0x00040020, +0x0000002c,0x00000007,0x0000002b,0x0003001d, +0x0000002e,0x00000006,0x0003001e,0x0000002f, +0x0000002e,0x00040020,0x00000030,0x0000000c, +0x0000002f,0x0004003b,0x00000030,0x00000031, +0x0000000c,0x00040020,0x00000035,0x0000000c, +0x00000006,0x0004002b,0x00000006,0x0000003a, +0x00000001,0x0004002b,0x00000006,0x00000041, +0x00000002,0x00040020,0x00000048,0x00000007, +0x00000008,0x0004002b,0x00000006,0x0000004b, +0x00000003,0x0004002b,0x00000006,0x00000053, +0x00000004,0x0004002b,0x00000006,0x0000005a, +0x00000005,0x0004002b,0x00000006,0x00000061, +0x00000006,0x0004002b,0x00000006,0x0000006a, +0x00000007,0x00021178,0x00000070,0x00040020, +0x00000071,0x00000006,0x00000070,0x0004003b, +0x00000071,0x00000072,0x00000006,0x000214dd, +0x00000073,0x00040020,0x00000074,0x00000000, +0x00000073,0x0004003b,0x00000074,0x00000075, +0x00000000,0x0004002b,0x00000006,0x00000077, +0x000000ff,0x00030029,0x00000022,0x00000085, +0x0004002b,0x0000001d,0x00000086,0x00000001, +0x00040017,0x0000008d,0x00000008,0x00000002, +0x00040020,0x0000008e,0x00000007,0x0000008d, +0x0003001d,0x0000009e,0x00000006,0x0003001e, +0x0000009f,0x0000009e,0x00040020,0x000000a0, +0x0000000c,0x0000009f,0x0004003b,0x000000a0, +0x000000a1,0x0000000c,0x0004002b,0x00000008, +0x000000ba,0x00000000,0x0004002b,0x00000006, +0x000000c5,0xffffffff,0x0004002b,0x00000006, +0x000000c7,0x00000040,0x0006002c,0x00000012, +0x000000c8,0x000000c7,0x0000003a,0x0000003a, +0x00050036,0x00000002,0x00000004,0x00000000, +0x00000003,0x000200f8,0x00000005,0x0004003b, +0x00000007,0x00000011,0x00000007,0x0004003b, +0x00000007,0x00000027,0x00000007,0x0004003b, +0x0000002c,0x0000002d,0x00000007,0x0004003b, +0x00000007,0x00000034,0x00000007,0x0004003b, +0x00000007,0x0000003c,0x00000007,0x0004003b, +0x00000007,0x00000043,0x00000007,0x0004003b, +0x00000048,0x00000049,0x00000007,0x0004003b, +0x00000007,0x0000004d,0x00000007,0x0004003b, +0x0000002c,0x00000051,0x00000007,0x0004003b, +0x00000007,0x00000055,0x00000007,0x0004003b, +0x00000007,0x0000005c,0x00000007,0x0004003b, +0x00000007,0x00000063,0x00000007,0x0004003b, +0x00000048,0x00000068,0x00000007,0x0004003b, +0x00000007,0x0000006c,0x00000007,0x0004003b, +0x00000007,0x00000082,0x00000007,0x0004003b, +0x00000048,0x0000008b,0x00000007,0x0004003b, +0x0000008e,0x0000008f,0x00000007,0x0004003b, +0x00000007,0x00000091,0x00000007,0x0004003b, +0x00000007,0x00000094,0x00000007,0x0004003b, +0x00000007,0x00000097,0x00000007,0x00050041, +0x00000016,0x00000017,0x00000014,0x00000015, +0x0004003d,0x00000006,0x00000018,0x00000017, +0x0003003e,0x00000011,0x00000018,0x0004003d, +0x00000006,0x00000019,0x00000011,0x00050041, +0x0000001f,0x00000020,0x0000001c,0x0000001e, +0x0004003d,0x00000006,0x00000021,0x00000020, +0x000500ae,0x00000022,0x00000023,0x00000019, +0x00000021,0x000300f7,0x00000025,0x00000000, +0x000400fa,0x00000023,0x00000024,0x00000025, +0x000200f8,0x00000024,0x000100fd,0x000200f8, +0x00000025,0x0004003d,0x00000006,0x00000028, +0x00000011,0x00050084,0x00000006,0x0000002a, +0x00000028,0x00000029,0x0003003e,0x00000027, +0x0000002a,0x0004003d,0x00000006,0x00000032, +0x00000027,0x00050080,0x00000006,0x00000033, +0x00000032,0x00000015,0x00060041,0x00000035, +0x00000036,0x00000031,0x0000001e,0x00000033, +0x0004003d,0x00000006,0x00000037,0x00000036, +0x0003003e,0x00000034,0x00000037,0x00050039, +0x00000008,0x00000038,0x0000000b,0x00000034, +0x0004003d,0x00000006,0x00000039,0x00000027, +0x00050080,0x00000006,0x0000003b,0x00000039, +0x0000003a,0x00060041,0x00000035,0x0000003d, +0x00000031,0x0000001e,0x0000003b,0x0004003d, +0x00000006,0x0000003e,0x0000003d,0x0003003e, +0x0000003c,0x0000003e,0x00050039,0x00000008, +0x0000003f,0x0000000b,0x0000003c,0x0004003d, +0x00000006,0x00000040,0x00000027,0x00050080, +0x00000006,0x00000042,0x00000040,0x00000041, +0x00060041,0x00000035,0x00000044,0x00000031, +0x0000001e,0x00000042,0x0004003d,0x00000006, +0x00000045,0x00000044,0x0003003e,0x00000043, +0x00000045,0x00050039,0x00000008,0x00000046, +0x0000000b,0x00000043,0x00060050,0x0000002b, +0x00000047,0x00000038,0x0000003f,0x00000046, +0x0003003e,0x0000002d,0x00000047,0x0004003d, +0x00000006,0x0000004a,0x00000027,0x00050080, +0x00000006,0x0000004c,0x0000004a,0x0000004b, +0x00060041,0x00000035,0x0000004e,0x00000031, +0x0000001e,0x0000004c,0x0004003d,0x00000006, +0x0000004f,0x0000004e,0x0003003e,0x0000004d, +0x0000004f,0x00050039,0x00000008,0x00000050, +0x0000000b,0x0000004d,0x0003003e,0x00000049, +0x00000050,0x0004003d,0x00000006,0x00000052, +0x00000027,0x00050080,0x00000006,0x00000054, +0x00000052,0x00000053,0x00060041,0x00000035, +0x00000056,0x00000031,0x0000001e,0x00000054, +0x0004003d,0x00000006,0x00000057,0x00000056, +0x0003003e,0x00000055,0x00000057,0x00050039, +0x00000008,0x00000058,0x0000000b,0x00000055, +0x0004003d,0x00000006,0x00000059,0x00000027, +0x00050080,0x00000006,0x0000005b,0x00000059, +0x0000005a,0x00060041,0x00000035,0x0000005d, +0x00000031,0x0000001e,0x0000005b,0x0004003d, +0x00000006,0x0000005e,0x0000005d,0x0003003e, +0x0000005c,0x0000005e,0x00050039,0x00000008, +0x0000005f,0x0000000b,0x0000005c,0x0004003d, +0x00000006,0x00000060,0x00000027,0x00050080, +0x00000006,0x00000062,0x00000060,0x00000061, +0x00060041,0x00000035,0x00000064,0x00000031, +0x0000001e,0x00000062,0x0004003d,0x00000006, +0x00000065,0x00000064,0x0003003e,0x00000063, +0x00000065,0x00050039,0x00000008,0x00000066, +0x0000000b,0x00000063,0x00060050,0x0000002b, +0x00000067,0x00000058,0x0000005f,0x00000066, +0x0003003e,0x00000051,0x00000067,0x0004003d, +0x00000006,0x00000069,0x00000027,0x00050080, +0x00000006,0x0000006b,0x00000069,0x0000006a, +0x00060041,0x00000035,0x0000006d,0x00000031, +0x0000001e,0x0000006b,0x0004003d,0x00000006, +0x0000006e,0x0000006d,0x0003003e,0x0000006c, +0x0000006e,0x00050039,0x00000008,0x0000006f, +0x0000000b,0x0000006c,0x0003003e,0x00000068, +0x0000006f,0x0004003d,0x00000073,0x00000076, +0x00000075,0x0004003d,0x0000002b,0x00000078, +0x0000002d,0x0004003d,0x00000008,0x00000079, +0x00000049,0x0004003d,0x0000002b,0x0000007a, +0x00000051,0x0004003d,0x00000008,0x0000007b, +0x00000068,0x00091179,0x00000072,0x00000076, +0x0000003a,0x00000077,0x00000078,0x00000079, +0x0000007a,0x0000007b,0x000200f9,0x0000007c, +0x000200f8,0x0000007c,0x000400f6,0x0000007e, +0x0000007f,0x00000000,0x000200f9,0x00000080, +0x000200f8,0x00000080,0x0004117d,0x00000022, +0x00000081,0x00000072,0x000400fa,0x00000081, +0x0000007d,0x0000007e,0x000200f8,0x0000007d, +0x000200f9,0x0000007f,0x000200f8,0x0000007f, +0x000200f9,0x0000007c,0x000200f8,0x0000007e, +0x0004003d,0x00000006,0x00000083,0x00000011, +0x00050084,0x00000006,0x00000084,0x00000083, +0x00000053,0x0003003e,0x00000082,0x00000084, +0x0005117f,0x00000006,0x00000087,0x00000072, +0x00000086,0x000500aa,0x00000022,0x00000088, +0x00000087,0x0000003a,0x000300f7,0x0000008a, +0x00000000,0x000400fa,0x00000088,0x00000089, +0x000000b7,0x000200f8,0x00000089,0x00051782, +0x00000008,0x0000008c,0x00000072,0x00000086, +0x0003003e,0x0000008b,0x0000008c,0x00051788, +0x0000008d,0x00000090,0x00000072,0x00000086, +0x0003003e,0x0000008f,0x00000090,0x00051787, +0x0000001d,0x00000092,0x00000072,0x00000086, +0x0004007c,0x00000006,0x00000093,0x00000092, +0x0003003e,0x00000091,0x00000093,0x00051784, +0x0000001d,0x00000095,0x00000072,0x00000086, +0x0004007c,0x00000006,0x00000096,0x00000095, +0x0003003e,0x00000094,0x00000096,0x0004003d, +0x00000006,0x00000098,0x00000094,0x00050041, +0x0000001f,0x00000099,0x0000001c,0x00000086, +0x0004003d,0x00000006,0x0000009a,0x00000099, +0x00050084,0x00000006,0x0000009b,0x00000098, +0x0000009a,0x0004003d,0x00000006,0x0000009c, +0x00000091,0x00050080,0x00000006,0x0000009d, +0x0000009b,0x0000009c,0x0003003e,0x00000097, +0x0000009d,0x0004003d,0x00000006,0x000000a2, +0x00000082,0x00050080,0x00000006,0x000000a3, +0x000000a2,0x00000015,0x0004003d,0x00000008, +0x000000a4,0x0000008b,0x0004007c,0x00000006, +0x000000a5,0x000000a4,0x00060041,0x00000035, +0x000000a6,0x000000a1,0x0000001e,0x000000a3, +0x0003003e,0x000000a6,0x000000a5,0x0004003d, +0x00000006,0x000000a7,0x00000082,0x00050080, +0x00000006,0x000000a8,0x000000a7,0x0000003a, +0x00050041,0x00000048,0x000000a9,0x0000008f, +0x00000015,0x0004003d,0x00000008,0x000000aa, +0x000000a9,0x0004007c,0x00000006,0x000000ab, +0x000000aa,0x00060041,0x00000035,0x000000ac, +0x000000a1,0x0000001e,0x000000a8,0x0003003e, +0x000000ac,0x000000ab,0x0004003d,0x00000006, +0x000000ad,0x00000082,0x00050080,0x00000006, +0x000000ae,0x000000ad,0x00000041,0x00050041, +0x00000048,0x000000af,0x0000008f,0x0000003a, +0x0004003d,0x00000008,0x000000b0,0x000000af, +0x0004007c,0x00000006,0x000000b1,0x000000b0, +0x00060041,0x00000035,0x000000b2,0x000000a1, +0x0000001e,0x000000ae,0x0003003e,0x000000b2, +0x000000b1,0x0004003d,0x00000006,0x000000b3, +0x00000082,0x00050080,0x00000006,0x000000b4, +0x000000b3,0x0000004b,0x0004003d,0x00000006, +0x000000b5,0x00000097,0x00060041,0x00000035, +0x000000b6,0x000000a1,0x0000001e,0x000000b4, +0x0003003e,0x000000b6,0x000000b5,0x000200f9, +0x0000008a,0x000200f8,0x000000b7,0x0004003d, +0x00000006,0x000000b8,0x00000082,0x00050080, +0x00000006,0x000000b9,0x000000b8,0x00000015, +0x0004007c,0x00000006,0x000000bb,0x000000ba, +0x00060041,0x00000035,0x000000bc,0x000000a1, +0x0000001e,0x000000b9,0x0003003e,0x000000bc, +0x000000bb,0x0004003d,0x00000006,0x000000bd, +0x00000082,0x00050080,0x00000006,0x000000be, +0x000000bd,0x0000003a,0x00060041,0x00000035, +0x000000bf,0x000000a1,0x0000001e,0x000000be, +0x0003003e,0x000000bf,0x00000015,0x0004003d, +0x00000006,0x000000c0,0x00000082,0x00050080, +0x00000006,0x000000c1,0x000000c0,0x00000041, +0x00060041,0x00000035,0x000000c2,0x000000a1, +0x0000001e,0x000000c1,0x0003003e,0x000000c2, +0x00000015,0x0004003d,0x00000006,0x000000c3, +0x00000082,0x00050080,0x00000006,0x000000c4, +0x000000c3,0x0000004b,0x00060041,0x00000035, +0x000000c6,0x000000a1,0x0000001e,0x000000c4, +0x0003003e,0x000000c6,0x000000c5,0x000200f9, +0x0000008a,0x000200f8,0x0000008a,0x000100fd, +0x00010038,0x00050036,0x00000008,0x0000000b, +0x00000000,0x00000009,0x00030037,0x00000007, +0x0000000a,0x000200f8,0x0000000c,0x0004003d, +0x00000006,0x0000000d,0x0000000a,0x0004007c, +0x00000008,0x0000000e,0x0000000d,0x000200fe, +0x0000000e,0x00010038} +; diff --git a/vk/shaders/trace_ray_query_wide.spv.h b/vk/shaders/trace_ray_query_wide.spv.h new file mode 100644 index 0000000..9f985ac --- /dev/null +++ b/vk/shaders/trace_ray_query_wide.spv.h @@ -0,0 +1,158 @@ + // 1114.3.0 + #pragma once +const uint32_t trace_ray_query_wide_spv[] = { + 0x07230203,0x00010500,0x0008000b,0x000000c9,0x00000000,0x00020011,0x00000001,0x00020011, + 0x00001178,0x0006000a,0x5f565053,0x5f52484b,0x5f796172,0x72657571,0x00000079,0x0006000b, + 0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,0x00000000,0x00000001, + 0x000b000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x00000014,0x0000001c,0x00000031, + 0x00000072,0x00000075,0x0000009a,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001, + 0x00000001,0x00030003,0x00000002,0x000001cc,0x00060004,0x455f4c47,0x725f5458,0x715f7961, + 0x79726575,0x00000000,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00040005,0x0000000b, + 0x75287566,0x00003b31,0x00030005,0x0000000a,0x00000069,0x00030005,0x00000011,0x00646967, + 0x00080005,0x00000014,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,0x496e6f69,0x00000044, + 0x00030005,0x0000001a,0x00004350,0x00060006,0x0000001a,0x00000000,0x5f796172,0x6e756f63, + 0x00000074,0x00060006,0x0000001a,0x00000001,0x5f697274,0x6e756863,0x0000006b,0x00030005, + 0x0000001c,0x00006370,0x00030005,0x00000027,0x00006272,0x00030005,0x0000002d,0x0067726f, + 0x00040005,0x0000002f,0x73796152,0x00000000,0x00050006,0x0000002f,0x00000000,0x73796172, + 0x00000000,0x00030005,0x00000031,0x00000000,0x00040005,0x00000034,0x61726170,0x0000006d, + 0x00040005,0x0000003c,0x61726170,0x0000006d,0x00040005,0x00000043,0x61726170,0x0000006d, + 0x00040005,0x00000049,0x6e696d74,0x00000000,0x00040005,0x0000004d,0x61726170,0x0000006d, + 0x00030005,0x00000051,0x00726964,0x00040005,0x00000055,0x61726170,0x0000006d,0x00040005, + 0x0000005c,0x61726170,0x0000006d,0x00040005,0x00000063,0x61726170,0x0000006d,0x00040005, + 0x00000068,0x78616d74,0x00000000,0x00040005,0x0000006c,0x61726170,0x0000006d,0x00030005, + 0x00000072,0x00007172,0x00040005,0x00000075,0x73616c74,0x00000000,0x00030005,0x00000082, + 0x00006268,0x00030005,0x0000008b,0x00000074,0x00030005,0x0000008f,0x00006362,0x00040005, + 0x00000091,0x6d697270,0x00000000,0x00040005,0x00000094,0x74736e69,0x00000000,0x00040005, + 0x00000098,0x73746948,0x00000000,0x00050006,0x00000098,0x00000000,0x73746968,0x00000000, + 0x00030005,0x0000009a,0x00000000,0x00040047,0x00000014,0x0000000b,0x0000001c,0x00050048, + 0x0000001a,0x00000000,0x00000023,0x00000000,0x00050048,0x0000001a,0x00000001,0x00000023, + 0x00000004,0x00030047,0x0000001a,0x00000002,0x00040047,0x0000002e,0x00000006,0x00000004, + 0x00040048,0x0000002f,0x00000000,0x00000018,0x00050048,0x0000002f,0x00000000,0x00000023, + 0x00000000,0x00030047,0x0000002f,0x00000002,0x00040047,0x00000031,0x00000022,0x00000000, + 0x00040047,0x00000031,0x00000021,0x00000001,0x00040047,0x00000075,0x00000022,0x00000000, + 0x00040047,0x00000075,0x00000021,0x00000000,0x00040047,0x00000097,0x00000006,0x00000004, + 0x00040048,0x00000098,0x00000000,0x00000019,0x00050048,0x00000098,0x00000000,0x00000023, + 0x00000000,0x00030047,0x00000098,0x00000002,0x00040047,0x0000009a,0x00000022,0x00000000, + 0x00040047,0x0000009a,0x00000021,0x00000002,0x00040047,0x000000c8,0x0000000b,0x00000019, + 0x00020013,0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020, + 0x00000000,0x00040020,0x00000007,0x00000007,0x00000006,0x00030016,0x00000008,0x00000020, + 0x00040021,0x00000009,0x00000008,0x00000007,0x00040017,0x00000012,0x00000006,0x00000003, + 0x00040020,0x00000013,0x00000001,0x00000012,0x0004003b,0x00000013,0x00000014,0x00000001, + 0x0004002b,0x00000006,0x00000015,0x00000000,0x00040020,0x00000016,0x00000001,0x00000006, + 0x0004001e,0x0000001a,0x00000006,0x00000006,0x00040020,0x0000001b,0x00000009,0x0000001a, + 0x0004003b,0x0000001b,0x0000001c,0x00000009,0x00040015,0x0000001d,0x00000020,0x00000001, + 0x0004002b,0x0000001d,0x0000001e,0x00000000,0x00040020,0x0000001f,0x00000009,0x00000006, + 0x00020014,0x00000022,0x0004002b,0x00000006,0x00000029,0x00000008,0x00040017,0x0000002b, + 0x00000008,0x00000003,0x00040020,0x0000002c,0x00000007,0x0000002b,0x0003001d,0x0000002e, + 0x00000006,0x0003001e,0x0000002f,0x0000002e,0x00040020,0x00000030,0x0000000c,0x0000002f, + 0x0004003b,0x00000030,0x00000031,0x0000000c,0x00040020,0x00000035,0x0000000c,0x00000006, + 0x0004002b,0x00000006,0x0000003a,0x00000001,0x0004002b,0x00000006,0x00000041,0x00000002, + 0x00040020,0x00000048,0x00000007,0x00000008,0x0004002b,0x00000006,0x0000004b,0x00000003, + 0x0004002b,0x00000006,0x00000053,0x00000004,0x0004002b,0x00000006,0x0000005a,0x00000005, + 0x0004002b,0x00000006,0x00000061,0x00000006,0x0004002b,0x00000006,0x0000006a,0x00000007, + 0x00021178,0x00000070,0x00040020,0x00000071,0x00000006,0x00000070,0x0004003b,0x00000071, + 0x00000072,0x00000006,0x000214dd,0x00000073,0x00040020,0x00000074,0x00000000,0x00000073, + 0x0004003b,0x00000074,0x00000075,0x00000000,0x0004002b,0x00000006,0x00000077,0x000000ff, + 0x00030029,0x00000022,0x00000085,0x0004002b,0x0000001d,0x00000086,0x00000001,0x00040017, + 0x0000008d,0x00000008,0x00000002,0x00040020,0x0000008e,0x00000007,0x0000008d,0x0003001d, + 0x00000097,0x00000006,0x0003001e,0x00000098,0x00000097,0x00040020,0x00000099,0x0000000c, + 0x00000098,0x0004003b,0x00000099,0x0000009a,0x0000000c,0x0004002b,0x00000008,0x000000b7, + 0x00000000,0x0004002b,0x00000006,0x000000c2,0xffffffff,0x0004002b,0x00000006,0x000000c7, + 0x00000040,0x0006002c,0x00000012,0x000000c8,0x000000c7,0x0000003a,0x0000003a,0x00050036, + 0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000007, + 0x00000011,0x00000007,0x0004003b,0x00000007,0x00000027,0x00000007,0x0004003b,0x0000002c, + 0x0000002d,0x00000007,0x0004003b,0x00000007,0x00000034,0x00000007,0x0004003b,0x00000007, + 0x0000003c,0x00000007,0x0004003b,0x00000007,0x00000043,0x00000007,0x0004003b,0x00000048, + 0x00000049,0x00000007,0x0004003b,0x00000007,0x0000004d,0x00000007,0x0004003b,0x0000002c, + 0x00000051,0x00000007,0x0004003b,0x00000007,0x00000055,0x00000007,0x0004003b,0x00000007, + 0x0000005c,0x00000007,0x0004003b,0x00000007,0x00000063,0x00000007,0x0004003b,0x00000048, + 0x00000068,0x00000007,0x0004003b,0x00000007,0x0000006c,0x00000007,0x0004003b,0x00000007, + 0x00000082,0x00000007,0x0004003b,0x00000048,0x0000008b,0x00000007,0x0004003b,0x0000008e, + 0x0000008f,0x00000007,0x0004003b,0x00000007,0x00000091,0x00000007,0x0004003b,0x00000007, + 0x00000094,0x00000007,0x00050041,0x00000016,0x00000017,0x00000014,0x00000015,0x0004003d, + 0x00000006,0x00000018,0x00000017,0x0003003e,0x00000011,0x00000018,0x0004003d,0x00000006, + 0x00000019,0x00000011,0x00050041,0x0000001f,0x00000020,0x0000001c,0x0000001e,0x0004003d, + 0x00000006,0x00000021,0x00000020,0x000500ae,0x00000022,0x00000023,0x00000019,0x00000021, + 0x000300f7,0x00000025,0x00000000,0x000400fa,0x00000023,0x00000024,0x00000025,0x000200f8, + 0x00000024,0x000100fd,0x000200f8,0x00000025,0x0004003d,0x00000006,0x00000028,0x00000011, + 0x00050084,0x00000006,0x0000002a,0x00000028,0x00000029,0x0003003e,0x00000027,0x0000002a, + 0x0004003d,0x00000006,0x00000032,0x00000027,0x00050080,0x00000006,0x00000033,0x00000032, + 0x00000015,0x00060041,0x00000035,0x00000036,0x00000031,0x0000001e,0x00000033,0x0004003d, + 0x00000006,0x00000037,0x00000036,0x0003003e,0x00000034,0x00000037,0x00050039,0x00000008, + 0x00000038,0x0000000b,0x00000034,0x0004003d,0x00000006,0x00000039,0x00000027,0x00050080, + 0x00000006,0x0000003b,0x00000039,0x0000003a,0x00060041,0x00000035,0x0000003d,0x00000031, + 0x0000001e,0x0000003b,0x0004003d,0x00000006,0x0000003e,0x0000003d,0x0003003e,0x0000003c, + 0x0000003e,0x00050039,0x00000008,0x0000003f,0x0000000b,0x0000003c,0x0004003d,0x00000006, + 0x00000040,0x00000027,0x00050080,0x00000006,0x00000042,0x00000040,0x00000041,0x00060041, + 0x00000035,0x00000044,0x00000031,0x0000001e,0x00000042,0x0004003d,0x00000006,0x00000045, + 0x00000044,0x0003003e,0x00000043,0x00000045,0x00050039,0x00000008,0x00000046,0x0000000b, + 0x00000043,0x00060050,0x0000002b,0x00000047,0x00000038,0x0000003f,0x00000046,0x0003003e, + 0x0000002d,0x00000047,0x0004003d,0x00000006,0x0000004a,0x00000027,0x00050080,0x00000006, + 0x0000004c,0x0000004a,0x0000004b,0x00060041,0x00000035,0x0000004e,0x00000031,0x0000001e, + 0x0000004c,0x0004003d,0x00000006,0x0000004f,0x0000004e,0x0003003e,0x0000004d,0x0000004f, + 0x00050039,0x00000008,0x00000050,0x0000000b,0x0000004d,0x0003003e,0x00000049,0x00000050, + 0x0004003d,0x00000006,0x00000052,0x00000027,0x00050080,0x00000006,0x00000054,0x00000052, + 0x00000053,0x00060041,0x00000035,0x00000056,0x00000031,0x0000001e,0x00000054,0x0004003d, + 0x00000006,0x00000057,0x00000056,0x0003003e,0x00000055,0x00000057,0x00050039,0x00000008, + 0x00000058,0x0000000b,0x00000055,0x0004003d,0x00000006,0x00000059,0x00000027,0x00050080, + 0x00000006,0x0000005b,0x00000059,0x0000005a,0x00060041,0x00000035,0x0000005d,0x00000031, + 0x0000001e,0x0000005b,0x0004003d,0x00000006,0x0000005e,0x0000005d,0x0003003e,0x0000005c, + 0x0000005e,0x00050039,0x00000008,0x0000005f,0x0000000b,0x0000005c,0x0004003d,0x00000006, + 0x00000060,0x00000027,0x00050080,0x00000006,0x00000062,0x00000060,0x00000061,0x00060041, + 0x00000035,0x00000064,0x00000031,0x0000001e,0x00000062,0x0004003d,0x00000006,0x00000065, + 0x00000064,0x0003003e,0x00000063,0x00000065,0x00050039,0x00000008,0x00000066,0x0000000b, + 0x00000063,0x00060050,0x0000002b,0x00000067,0x00000058,0x0000005f,0x00000066,0x0003003e, + 0x00000051,0x00000067,0x0004003d,0x00000006,0x00000069,0x00000027,0x00050080,0x00000006, + 0x0000006b,0x00000069,0x0000006a,0x00060041,0x00000035,0x0000006d,0x00000031,0x0000001e, + 0x0000006b,0x0004003d,0x00000006,0x0000006e,0x0000006d,0x0003003e,0x0000006c,0x0000006e, + 0x00050039,0x00000008,0x0000006f,0x0000000b,0x0000006c,0x0003003e,0x00000068,0x0000006f, + 0x0004003d,0x00000073,0x00000076,0x00000075,0x0004003d,0x0000002b,0x00000078,0x0000002d, + 0x0004003d,0x00000008,0x00000079,0x00000049,0x0004003d,0x0000002b,0x0000007a,0x00000051, + 0x0004003d,0x00000008,0x0000007b,0x00000068,0x00091179,0x00000072,0x00000076,0x0000003a, + 0x00000077,0x00000078,0x00000079,0x0000007a,0x0000007b,0x000200f9,0x0000007c,0x000200f8, + 0x0000007c,0x000400f6,0x0000007e,0x0000007f,0x00000000,0x000200f9,0x00000080,0x000200f8, + 0x00000080,0x0004117d,0x00000022,0x00000081,0x00000072,0x000400fa,0x00000081,0x0000007d, + 0x0000007e,0x000200f8,0x0000007d,0x000200f9,0x0000007f,0x000200f8,0x0000007f,0x000200f9, + 0x0000007c,0x000200f8,0x0000007e,0x0004003d,0x00000006,0x00000083,0x00000011,0x00050084, + 0x00000006,0x00000084,0x00000083,0x0000005a,0x0003003e,0x00000082,0x00000084,0x0005117f, + 0x00000006,0x00000087,0x00000072,0x00000086,0x000500aa,0x00000022,0x00000088,0x00000087, + 0x0000003a,0x000300f7,0x0000008a,0x00000000,0x000400fa,0x00000088,0x00000089,0x000000b4, + 0x000200f8,0x00000089,0x00051782,0x00000008,0x0000008c,0x00000072,0x00000086,0x0003003e, + 0x0000008b,0x0000008c,0x00051788,0x0000008d,0x00000090,0x00000072,0x00000086,0x0003003e, + 0x0000008f,0x00000090,0x00051787,0x0000001d,0x00000092,0x00000072,0x00000086,0x0004007c, + 0x00000006,0x00000093,0x00000092,0x0003003e,0x00000091,0x00000093,0x00051784,0x0000001d, + 0x00000095,0x00000072,0x00000086,0x0004007c,0x00000006,0x00000096,0x00000095,0x0003003e, + 0x00000094,0x00000096,0x0004003d,0x00000006,0x0000009b,0x00000082,0x00050080,0x00000006, + 0x0000009c,0x0000009b,0x00000015,0x0004003d,0x00000008,0x0000009d,0x0000008b,0x0004007c, + 0x00000006,0x0000009e,0x0000009d,0x00060041,0x00000035,0x0000009f,0x0000009a,0x0000001e, + 0x0000009c,0x0003003e,0x0000009f,0x0000009e,0x0004003d,0x00000006,0x000000a0,0x00000082, + 0x00050080,0x00000006,0x000000a1,0x000000a0,0x0000003a,0x00050041,0x00000048,0x000000a2, + 0x0000008f,0x00000015,0x0004003d,0x00000008,0x000000a3,0x000000a2,0x0004007c,0x00000006, + 0x000000a4,0x000000a3,0x00060041,0x00000035,0x000000a5,0x0000009a,0x0000001e,0x000000a1, + 0x0003003e,0x000000a5,0x000000a4,0x0004003d,0x00000006,0x000000a6,0x00000082,0x00050080, + 0x00000006,0x000000a7,0x000000a6,0x00000041,0x00050041,0x00000048,0x000000a8,0x0000008f, + 0x0000003a,0x0004003d,0x00000008,0x000000a9,0x000000a8,0x0004007c,0x00000006,0x000000aa, + 0x000000a9,0x00060041,0x00000035,0x000000ab,0x0000009a,0x0000001e,0x000000a7,0x0003003e, + 0x000000ab,0x000000aa,0x0004003d,0x00000006,0x000000ac,0x00000082,0x00050080,0x00000006, + 0x000000ad,0x000000ac,0x0000004b,0x0004003d,0x00000006,0x000000ae,0x00000094,0x00060041, + 0x00000035,0x000000af,0x0000009a,0x0000001e,0x000000ad,0x0003003e,0x000000af,0x000000ae, + 0x0004003d,0x00000006,0x000000b0,0x00000082,0x00050080,0x00000006,0x000000b1,0x000000b0, + 0x00000053,0x0004003d,0x00000006,0x000000b2,0x00000091,0x00060041,0x00000035,0x000000b3, + 0x0000009a,0x0000001e,0x000000b1,0x0003003e,0x000000b3,0x000000b2,0x000200f9,0x0000008a, + 0x000200f8,0x000000b4,0x0004003d,0x00000006,0x000000b5,0x00000082,0x00050080,0x00000006, + 0x000000b6,0x000000b5,0x00000015,0x0004007c,0x00000006,0x000000b8,0x000000b7,0x00060041, + 0x00000035,0x000000b9,0x0000009a,0x0000001e,0x000000b6,0x0003003e,0x000000b9,0x000000b8, + 0x0004003d,0x00000006,0x000000ba,0x00000082,0x00050080,0x00000006,0x000000bb,0x000000ba, + 0x0000003a,0x00060041,0x00000035,0x000000bc,0x0000009a,0x0000001e,0x000000bb,0x0003003e, + 0x000000bc,0x00000015,0x0004003d,0x00000006,0x000000bd,0x00000082,0x00050080,0x00000006, + 0x000000be,0x000000bd,0x00000041,0x00060041,0x00000035,0x000000bf,0x0000009a,0x0000001e, + 0x000000be,0x0003003e,0x000000bf,0x00000015,0x0004003d,0x00000006,0x000000c0,0x00000082, + 0x00050080,0x00000006,0x000000c1,0x000000c0,0x0000004b,0x00060041,0x00000035,0x000000c3, + 0x0000009a,0x0000001e,0x000000c1,0x0003003e,0x000000c3,0x000000c2,0x0004003d,0x00000006, + 0x000000c4,0x00000082,0x00050080,0x00000006,0x000000c5,0x000000c4,0x00000053,0x00060041, + 0x00000035,0x000000c6,0x0000009a,0x0000001e,0x000000c5,0x0003003e,0x000000c6,0x00000015, + 0x000200f9,0x0000008a,0x000200f8,0x0000008a,0x000100fd,0x00010038,0x00050036,0x00000008, + 0x0000000b,0x00000000,0x00000009,0x00030037,0x00000007,0x0000000a,0x000200f8,0x0000000c, + 0x0004003d,0x00000006,0x0000000d,0x0000000a,0x0004007c,0x00000008,0x0000000e,0x0000000d, + 0x000200fe,0x0000000e,0x00010038 +}; From fc9e650ddb39b3acacb4c97faba45b17a2047b50 Mon Sep 17 00:00:00 2001 From: Syoyo Fujita Date: Thu, 2 Jul 2026 02:42:44 +0900 Subject: [PATCH 2/3] tri: large-scene CPU BVH work (port from tinyusdz fork) Brings the vendored tinyusdz fork's CPU triangle/BVH-core improvements back upstream. All additive relative to upstream (no upstream function or API removed; recent upstream work -- lrt_tri_surface_project, NURBS trim loops -- is preserved). Verified against upstream's own tests: test_lightrt_c_surfaces and test_lightrt_c_quadtetra pass; test_lightrt_c_tri produces the identical 7 pre-existing failures (single-tri layout=10/11, BVH8_Q4) that upstream's own tri.c already fails -- i.e. the port is behavior-neutral. lightrt_c_vk_test also passes on an RTX 5060 Ti with this tri.c. Included: - Indexed triangle build (lrt_tri_scene_build_indexed): callers with shared/ welded vertices pass positions + uint32 indices instead of a 9*ntris soup. - Batched multi-scene build (lrt_tri_scene_build_batch / lrt_batch_build_one): build many small scenes in one parallel pass. - "Drop the soup" leaf-vertex recovery: lrt_tri_get_slot / lrt_tri_get_verts / lrt_tri_scene_has_verts / lrt_tri_slot_count let a caller recover hit-triangle vertices from the BVH leaves instead of retaining the full input soup (large resident-memory saving on 20M+ instance scenes). - Morton-LBVH curve build (parallel curve BVH). - MSVC / llvm-mingw portability: lrt_clzll / lrt_ctz wrap __builtin_clzll / __builtin_ctz so the SIMD leaf/lane code builds on MSVC 2022 and llvm-mingw. Co-Authored-By: Claude Opus 4.8 --- lightrt_c_tri.c | 537 ++++++++++++++++++++++++++++++++++++++++-------- lightrt_c_tri.h | 47 ++++- 2 files changed, 492 insertions(+), 92 deletions(-) diff --git a/lightrt_c_tri.c b/lightrt_c_tri.c index 0064f46..c28a156 100644 --- a/lightrt_c_tri.c +++ b/lightrt_c_tri.c @@ -43,6 +43,42 @@ #define LRT_TRI_HAS_MMAP 0 #endif +#if defined(_WIN32) +#include /* _aligned_malloc / _aligned_free */ +#endif + +/* C11 is optional in practice: some toolchains (notably llvm-mingw) + * do not ship the header even though they leave __STDC_NO_THREADS__ undefined. + * Gate the parallel build/traversal paths on the header actually being present; + * otherwise fall back to the serial path (functionally identical, slower). */ +#if !defined(__STDC_NO_THREADS__) && \ + (!defined(__has_include) || __has_include()) +#define LRT_TRI_HAS_THREADS 1 +#else +#define LRT_TRI_HAS_THREADS 0 +#endif + +/* Count-leading/trailing-zeros: GCC/Clang have __builtin_clzll/__builtin_ctz; + * MSVC exposes the equivalents via bit-scan intrinsics. Both the + * builtins and the intrinsics are undefined for a zero input, so behaviour + * parity is preserved (call sites never pass zero). */ +#if defined(_MSC_VER) +#include +static inline int lrt_clzll(unsigned long long x) { + unsigned long idx; + _BitScanReverse64(&idx, x); + return 63 - (int)idx; +} +static inline int lrt_ctz(unsigned int x) { + unsigned long idx; + _BitScanForward(&idx, x); + return (int)idx; +} +#else +static inline int lrt_clzll(unsigned long long x) { return __builtin_clzll(x); } +static inline int lrt_ctz(unsigned int x) { return __builtin_ctz(x); } +#endif + /* ------------------------------------------------------------------------- */ /* SIMD detection (compile-time; scalar fallback always available). */ /* ------------------------------------------------------------------------- */ @@ -408,6 +444,11 @@ struct lrt_tri_scene { uint32_t node_count; void *blocks; /* lrt_tri4[] when layout==4, lrt_tri8[] when layout==8 */ uint32_t block_count; + /* Optional prim_id -> (block*width + lane) map so callers can recover a + * triangle's object-space vertices from the swizzled leaves (lrt_tri_get_verts) + * and drop their own vertex copy. Built only for plain TRI_PRIM_TRI scenes; + * NULL otherwise. */ + uint32_t *prim2slot; uint32_t block_stride; /* bytes per leaf block (varies for quantized leaves) */ lrt_tri_stats stats; const char *kernel_name; @@ -455,6 +496,73 @@ static inline const float *tri_block_floats(const void *blocks, uint32_t idx, (size_t)idx * tri_block_size(width)); } +/* Build the prim_id -> leaf-slot map for a plain triangle scene, so + * lrt_tri_get_verts can recover object-space vertices from the swizzled leaves. + * Cheap O(ntris) pass over the leaf blocks; leaves s->prim2slot NULL on any + * unsupported layout or OOM (caller then keeps its own vertex copy). */ +static void tri_build_prim2slot(lrt_tri_scene *s) { + if (!s || s->prim_kind != TRI_PRIM_TRI || s->quantized || s->curve) return; + const int w = s->layout; /* 4 or 8 */ + const uint32_t nt = s->original_ntris; + if (nt == 0) return; + uint32_t *p2s = (uint32_t *)malloc((size_t)nt * sizeof(uint32_t)); + if (!p2s) return; + for (uint32_t i = 0; i < nt; i++) p2s[i] = LRT_TRI_NO_HIT; + for (uint32_t blk = 0; blk < s->block_count; blk++) { + const float *f = tri_block_floats(s->blocks, blk, w); + const uint32_t *ids = (const uint32_t *)(const void *)(f + 9 * (size_t)w); + for (int lane = 0; lane < w; lane++) { + const uint32_t prim = ids[lane]; + if (prim < nt) p2s[prim] = blk * (uint32_t)w + (uint32_t)lane; + } + } + s->prim2slot = p2s; +} + +int lrt_tri_scene_has_verts(const lrt_tri_scene *s) { + return (s && s->prim2slot) ? 1 : 0; +} + +int lrt_tri_get_verts(const lrt_tri_scene *s, uint32_t prim_id, float v0[3], + float v1[3], float v2[3]) { + if (!s || !s->prim2slot || prim_id >= s->original_ntris) return 0; + const uint32_t slot = s->prim2slot[prim_id]; + if (slot == LRT_TRI_NO_HIT) return 0; + const int w = s->layout; + const uint32_t blk = slot / (uint32_t)w; + const uint32_t lane = slot % (uint32_t)w; + const float *f = tri_block_floats(s->blocks, blk, w); + /* Leaves store v0 and the precomputed edges e1=v1-v0, e2=v2-v0 (see + * tri_emit_leaf); v0 + e is byte-exact for adjacent mesh coords (Sterbenz). */ + const float v0x = f[0 * w + lane], v0y = f[1 * w + lane], v0z = f[2 * w + lane]; + v0[0] = v0x; + v0[1] = v0y; + v0[2] = v0z; + v1[0] = v0x + f[3 * w + lane]; + v1[1] = v0y + f[4 * w + lane]; + v1[2] = v0z + f[5 * w + lane]; + v2[0] = v0x + f[6 * w + lane]; + v2[1] = v0y + f[7 * w + lane]; + v2[2] = v0z + f[8 * w + lane]; + return 1; +} + +/* Leaf slot (block*width + lane) for a triangle, or LRT_TRI_NO_HIT. Slots run in + * BVH leaf-emission order, so a caller can lay per-triangle shading out in slot + * order for cache-coherent hit-time access (adjacent leaf triangles -> adjacent + * slots). Same lookup lrt_tri_get_verts does internally. */ +uint32_t lrt_tri_get_slot(const lrt_tri_scene *s, uint32_t prim_id) { + if (!s || !s->prim2slot || prim_id >= s->original_ntris) return LRT_TRI_NO_HIT; + return s->prim2slot[prim_id]; +} + +/* Number of leaf slots (block_count * width): the size a slot-indexed array must + * have. Includes padding lanes (their prim_id is LRT_TRI_NO_HIT). */ +uint32_t lrt_tri_slot_count(const lrt_tri_scene *s) { + if (!s) return 0; + return s->block_count * (uint32_t)s->layout; +} + /* Binary build node (arena, freed after collapse). */ typedef struct tri_bnode { float lo[3], hi[3]; @@ -466,7 +574,11 @@ typedef struct tri_bnode { * [node_next, node_end) slice of the shared bnodes arena and a disjoint range * of the shared indices array, so no synchronization is needed. */ typedef struct tri_build_ctx { - const float *verts; /* caller soup, 9*ntris (NULL for curve scenes) */ + const float *verts; /* caller soup, 9*ntris (NULL for curve scenes), OR + * when vidx != NULL: a unique-vertex array (3*nverts) + * gathered through vidx (indexed build, see tri_fetch9). */ + const uint32_t *vidx; /* NULL = flat soup; else 3*ntris vertex indices into + * `verts` (3 floats/vertex) for indexed triangle input */ const tri_subseg *subsegs; /* curve sub-segments (NULL for triangles) */ const tri_rlcseg *rlcsegs; /* round-linear curve segs (TRI_PRIM_RLCURVE) */ size_t ntris; @@ -537,7 +649,9 @@ static void *tri_aligned_alloc(size_t align, size_t size) { } #endif size = (size + align - 1u) & ~(align - 1u); -#if defined(_MSC_VER) +#if defined(_WIN32) + /* mingw (incl. llvm-mingw UCRT) has no C11 aligned_alloc; use the Win32 + * _aligned_malloc/_aligned_free pair (declared in ). */ return _aligned_malloc(size, align); #else return aligned_alloc(align, size); @@ -545,7 +659,7 @@ static void *tri_aligned_alloc(size_t align, size_t size) { } static void tri_aligned_free(void *p) { -#if defined(_MSC_VER) +#if defined(_WIN32) _aligned_free(p); #else free(p); @@ -596,7 +710,7 @@ typedef struct tri_bin { /* ------------------------------------------------------------------------- */ /* Parallel-for helper (C11 threads). */ /* ------------------------------------------------------------------------- */ -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS #include #define TRI_PAR_NODE_MIN (1u << 16) /* intra-node parallelism above this */ @@ -650,7 +764,7 @@ static void tri_parallel_for(uint32_t n, unsigned threads, /* Binary SAH build. */ /* ------------------------------------------------------------------------- */ -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS /* Parallel 3-axis binning over indices[first+begin, first+end) per chunk. */ typedef struct tri_bin_job { const tri_build_ctx *c; @@ -740,6 +854,22 @@ static void tri_part_copyback(void *arg, unsigned chunk, uint32_t begin, /* Per-triangle AABB + centroid precompute. Returns non-zero on non-finite * input. Parallelized by chunks when threads > 1. */ +/* Return a pointer to triangle `i`'s 9 object-space floats (v0 v1 v2). For a + * flat soup (vidx == NULL) this is a zero-copy view into `verts`; for indexed + * input it gathers the three vertices through vidx into `buf` (byte-identical + * floats, just relocated from the caller's de-index into the build). */ +static inline const float *tri_fetch9(const float *verts, const uint32_t *vidx, + size_t i, float buf[9]) { + if (!vidx) return &verts[i * 9]; + const uint32_t a = vidx[i * 3 + 0]; + const uint32_t b = vidx[i * 3 + 1]; + const uint32_t c = vidx[i * 3 + 2]; + memcpy(&buf[0], &verts[(size_t)a * 3], 3 * sizeof(float)); + memcpy(&buf[3], &verts[(size_t)b * 3], 3 * sizeof(float)); + memcpy(&buf[6], &verts[(size_t)c * 3], 3 * sizeof(float)); + return buf; +} + typedef struct tri_precompute_job { tri_build_ctx *c; atomic_int bad; @@ -749,7 +879,8 @@ static void tri_precompute_chunk_impl(tri_build_ctx *c, uint32_t begin, uint32_t end, int *bad_out) { int bad = 0; for (uint32_t i = begin; i < end; i++) { - const float *v = &c->verts[(size_t)i * 9]; + float vbuf[9]; + const float *v = tri_fetch9(c->verts, c->vidx, i, vbuf); for (int a = 0; a < 3; a++) { /* check the vertices, not min/max results: minf/maxf comparisons * are false for NaN and would swallow it */ @@ -767,7 +898,7 @@ static void tri_precompute_chunk_impl(tri_build_ctx *c, uint32_t begin, if (bad) *bad_out = 1; } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS static void tri_precompute_chunk(void *arg, unsigned chunk, uint32_t begin, uint32_t end) { (void)chunk; @@ -779,7 +910,7 @@ static void tri_precompute_chunk(void *arg, unsigned chunk, uint32_t begin, #endif static int tri_precompute(tri_build_ctx *c, unsigned threads) { -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (threads > 1 && c->ntris >= TRI_PAR_NODE_MIN) { tri_precompute_job j; j.c = c; @@ -799,7 +930,7 @@ static int tri_precompute(tri_build_ctx *c, unsigned threads) { * left-side count. Uses the parallel path for large nodes when enabled. */ static uint32_t tri_partition(tri_build_ctx *c, uint32_t first, uint32_t num, int axis, float pos) { -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (c->par_threads > 1 && c->par_scratch && num >= TRI_PAR_NODE_MIN) { tri_part_job j; j.c = c; @@ -909,7 +1040,7 @@ static uint32_t tri_build_node(tri_build_ctx *c, uint32_t first, uint32_t num, } } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (c->par_threads > 1 && num >= TRI_PAR_NODE_MIN) { unsigned threads = c->par_threads; if (threads > TRI_PAR_NODE_THREADS) threads = TRI_PAR_NODE_THREADS; @@ -1103,7 +1234,7 @@ static uint32_t tri_lbvh_find_split(const uint64_t *keys, uint32_t first, uint64_t kl = keys[first + num - 1] >> 32; uint64_t x = kf ^ kl; if (x == 0) return num / 2; /* identical codes: median */ - int bit = 63 - __builtin_clzll(x); + int bit = 63 - lrt_clzll(x); /* sorted: a prefix has the bit clear, the suffix has it set */ uint32_t lo = 0, hi = num - 1; while (hi - lo > 1) { @@ -1191,7 +1322,7 @@ static void tri_morton_chunk_impl(const tri_morton_job *j, uint32_t begin, } } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS static void tri_morton_chunk(void *arg, unsigned chunk, uint32_t begin, uint32_t end) { (void)chunk; @@ -1219,7 +1350,7 @@ static void tri_morton_encode(const tri_build_ctx *c, uint64_t *keys, j.base[a] = clo[a]; j.scale[a] = ext > 0.0f ? 1024.0f / ext : 0.0f; } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (threads > 1 && c->ntris >= TRI_PAR_NODE_MIN) { tri_parallel_for((uint32_t)c->ntris, threads, tri_morton_chunk, &j); return; @@ -1243,7 +1374,7 @@ typedef struct tri_build_task { uint32_t *parent_slot; /* &bnodes[parent].a or .b (unique per task) */ } tri_build_task; -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS typedef struct tri_build_pool { const tri_build_ctx *proto; tri_build_task *tasks; @@ -1286,7 +1417,7 @@ static int tri_build_worker(void *arg) { /* Build the binary tree over all primitives, using up to num_threads workers. * Returns the root node index, or fails via c->failed. */ static uint32_t tri_build_binary(tri_build_ctx *c, unsigned num_threads) { -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && c->ntris >= 4096) { /* Serial frontier expansion: always expand the largest task. */ enum { MAX_TASKS = 256 }; @@ -1407,7 +1538,7 @@ static uint32_t tri_build_binary(tri_build_ctx *c, unsigned num_threads) { * bounds are filled in afterwards from their children (reverse creation * order: children are always created after their parent). */ static uint32_t tri_build_lbvh(tri_build_ctx *c, unsigned num_threads) { -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && c->ntris >= 4096) { enum { MAX_TASKS = 256 }; tri_build_task tasks[MAX_TASKS]; @@ -2700,7 +2831,9 @@ static uint32_t tri_emit_leaf(tri_collapse_ctx *cc, const tri_bnode *bn) { if (k < count) { uint32_t prim = bc->indices[bn->a + k]; prims[lane] = prim; - memcpy(vtx[lane], &bc->verts[(size_t)prim * 9], + float vbuf[9]; + memcpy(vtx[lane], + tri_fetch9(bc->verts, bc->vidx, prim, vbuf), 9 * sizeof(float)); } else { prims[lane] = LRT_TRI_NO_HIT; @@ -2724,7 +2857,8 @@ static uint32_t tri_emit_leaf(tri_collapse_ctx *cc, const tri_bnode *bn) { uint32_t k = b * bw + lane; if (k < count) { uint32_t prim = bc->indices[bn->a + k]; - const float *v = &bc->verts[(size_t)prim * 9]; + float vbuf[9]; + const float *v = tri_fetch9(bc->verts, bc->vidx, prim, vbuf); f[0 * bw + lane] = v[0]; f[1 * bw + lane] = v[1]; f[2 * bw + lane] = v[2]; @@ -5452,7 +5586,7 @@ static inline void tri_block_isect_sse(const lrt_tri4 *blk, _mm_storeu_ps(ua, u); _mm_storeu_ps(va, v); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -5535,7 +5669,7 @@ static inline void tri_mt4_decoded_sse(__m128 v0x, __m128 v0y, __m128 v0z, _mm_storeu_ps(ua, u); _mm_storeu_ps(va, v); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -5715,7 +5849,7 @@ static int tri_occluded_bvh4_sse(const lrt_tri_scene *s, const lrt_ray *ray) { /* Unordered pushes: measured faster for any-hit than tnear-sorted, * octant-ordered, and eager-leaf variants on this workload. */ while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -5811,7 +5945,7 @@ static int tri_curve_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -6050,7 +6184,7 @@ static inline void tri_rlc4_isect_sse(const lrt_rlc4 *blk, _mm_store_ps(ua, u); _mm_store_si128((__m128i *)(void *)pa, primi); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -6147,7 +6281,7 @@ static int tri_rlcurve_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -6227,7 +6361,7 @@ static inline void tri_point4_isect_sse(const lrt_point4 *blk, _mm_store_ps(ta, tcand); _mm_store_si128((__m128i *)(void *)pa, primi); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -6330,7 +6464,7 @@ static inline void tri_flat4_isect_sse(const lrt_flat4 *blk, _mm_store_ps(ua, uu); _mm_store_si128((__m128i *)(void *)pa, primi); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -6419,7 +6553,7 @@ static int tri_point_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -6505,7 +6639,7 @@ static int tri_flatcurve_occluded_bvh4(const lrt_tri_scene *s, int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -6649,7 +6783,7 @@ static int tri_bez_sweep_sse(const float cp_world[16], const tri_ray_ctx *rc, _mm_store_ps(r2a, r2w); int leaf = (depth + 1 >= BEZ_SIMD_DEPTH); while (mask) { - int L = __builtin_ctz((unsigned)mask); + int L = lrt_ctz((unsigned)mask); mask &= mask - 1; if (leaf) { float u0 = usa[L], u1 = uea[L], hL = u1 - u0; @@ -6810,7 +6944,7 @@ static inline void tri_bez4_isect_sse(const lrt_bez4 *blk, const tri_ray_ctx *rc blk->b2x, blk->b2y, blk->b2z, blk->b2r, blk->b3x, blk->b3y, blk->b3z, blk->b3r}; while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; float cp[16]; for (int k = 0; k < 16; k++) cp[k] = src[k][lane]; @@ -6901,7 +7035,7 @@ static int tri_bezcurve_occluded_bvh4(const lrt_tri_scene *s, int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -6954,7 +7088,7 @@ static inline void tri_block_isect_sph_sse(const lrt_sph4 *blk, _mm_storeu_ps(ta, tsel); int mask = _mm_movemask_ps(_mm_cmplt_ps(tsel, bt)); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -7075,7 +7209,7 @@ static int tri_sphere_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -7144,7 +7278,7 @@ static int tri_user_intersect_bvh4(const lrt_tri_scene *s, const lrt_ray *ray, const lrt_user4 *blk = &blocks[blk0 + b]; int m = tri_user_box_mask_sse(blk, &sc, _mm_set1_ps(best_t)); while (m) { - int lane = __builtin_ctz((unsigned)m); + int lane = lrt_ctz((unsigned)m); m &= m - 1; uint32_t pid = blk->prim[lane]; if (pid == LRT_TRI_NO_HIT) continue; @@ -7207,7 +7341,7 @@ static int tri_user_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) { const lrt_user4 *blk = &blocks[blk0 + b]; int m = tri_user_box_mask_sse(blk, &sc, tmax4); while (m) { - int lane = __builtin_ctz((unsigned)m); + int lane = lrt_ctz((unsigned)m); m &= m - 1; uint32_t pid = blk->prim[lane]; if (pid == LRT_TRI_NO_HIT) continue; @@ -7232,7 +7366,7 @@ static int tri_user_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -7329,7 +7463,7 @@ static int tri_qtri_occluded_bvh4(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh4_slab_sse(n, &sc, tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); stack[sp++] = n->child[i]; @@ -7405,7 +7539,7 @@ static inline int tri_pipe4_step(const lrt_tri_scene *s, tri_pipe4_state *st) { float hit_tn[4]; int nhit = 0; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); int j = nhit++; @@ -7514,7 +7648,7 @@ static inline int tri_pipeocc4_step(const lrt_tri_scene *s, int mask = tri_bvh4_slab_sse(n, &st->sc, st->tmax4, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 4); st->stack[st->sp++] = n->child[i]; @@ -7644,7 +7778,7 @@ static inline void tri_block_isect_avx(const lrt_tri8 *blk, _mm256_store_ps(ua, u); _mm256_store_ps(va, v); while (mask) { - int lane = __builtin_ctz((unsigned)mask); + int lane = lrt_ctz((unsigned)mask); mask &= mask - 1; if (ta[lane] < *best_t) { *best_t = ta[lane]; @@ -7756,7 +7890,7 @@ static int tri_intersect_bvh8_avx2(const lrt_tri_scene *s, const lrt_ray *ray, float hit_tn[8]; int nhit = 0; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 8); int j = nhit++; @@ -7813,7 +7947,7 @@ static int tri_occluded_bvh8_avx2(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh8_slab_avx(n, &ac, tmax8, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 8); stack[sp++] = n->child[i]; @@ -7897,7 +8031,7 @@ static int tri_intersect_bvh8q_avx2(const lrt_tri_scene *s, const lrt_ray *ray, float hit_tn[8]; int nhit = 0; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 8); int j = nhit++; @@ -7954,7 +8088,7 @@ static int tri_occluded_bvh8q_avx2(const lrt_tri_scene *s, const lrt_ray *ray) { int mask = tri_bvh8q_slab_avx(n, &rc, ac.tmin, tmax8, tnear); mask &= (1 << n->nchildren) - 1; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, n->child[i], 8); stack[sp++] = n->child[i]; @@ -8077,7 +8211,7 @@ static inline int tri_bvh8q4_slab_avx(const lrt_bvh8q4_node *n, float hit_tn[8]; \ int nhit = 0; \ while (mask) { \ - int i = __builtin_ctz((unsigned)mask); \ + int i = lrt_ctz((unsigned)mask); \ mask &= mask - 1; \ tri_prefetch_ref(s, n->child[i], 8); \ int j = nhit++; \ @@ -8128,7 +8262,7 @@ static inline int tri_bvh8q4_slab_avx(const lrt_bvh8q4_node *n, int mask = SLAB(n, &rc, ac.tmin, tmax8, tnear); \ mask &= (1 << n->nchildren) - 1; \ while (mask) { \ - int i = __builtin_ctz((unsigned)mask); \ + int i = lrt_ctz((unsigned)mask); \ mask &= mask - 1; \ tri_prefetch_ref(s, n->child[i], 8); \ if (sp < TRI_STACK_SIZE) stack[sp++] = n->child[i]; \ @@ -8208,7 +8342,7 @@ static inline int tri_pipe8_step(const lrt_tri_scene *s, tri_pipe8_state *st) { float hit_tn[8]; int nhit = 0; while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, child[i], 8); int j = nhit++; @@ -8329,7 +8463,7 @@ static inline int tri_pipeocc8_step(const lrt_tri_scene *s, child = n->child; } while (mask) { - int i = __builtin_ctz((unsigned)mask); + int i = lrt_ctz((unsigned)mask); mask &= mask - 1; tri_prefetch_ref(s, child[i], 8); st->stack[st->sp++] = child[i]; @@ -8476,6 +8610,7 @@ static int tri_collapse_into(lrt_tri_scene *s, tri_build_ctx *bc, static lrt_tri_scene *tri_scene_build_impl(const float *vertices, size_t ntris, const lrt_tri_build_options *opts, const uint32_t *morton_override, + const uint32_t *vidx, lrt_result *err) { tri_set_err(err, LRT_RESULT_OK); if (!vertices || ntris == 0 || ntris > 0x07FFFFFFu) { @@ -8534,6 +8669,7 @@ static lrt_tri_scene *tri_scene_build_impl(const float *vertices, size_t ntris, tri_build_ctx bc; memset(&bc, 0, sizeof(bc)); bc.verts = vertices; + bc.vidx = vidx; /* NULL => flat soup; else indexed gather (tri_fetch9) */ bc.ntris = ntris; bc.max_leaf = max_leaf; bc.block_shift = layout == 8 ? 3u : 2u; @@ -8567,7 +8703,7 @@ static lrt_tri_scene *tri_scene_build_impl(const float *vertices, size_t ntris, } unsigned num_threads = o.num_threads ? o.num_threads : 1u; -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && ntris >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(ntris * sizeof(uint32_t)); @@ -8741,13 +8877,117 @@ static lrt_tri_scene *tri_scene_build_impl(const float *vertices, size_t ntris, : quantized ? "bvh8q/scalar" : "bvh8/scalar"; #endif + tri_build_prim2slot(s); return s; } lrt_tri_scene *lrt_tri_scene_build(const float *vertices, size_t ntris, const lrt_tri_build_options *opts, lrt_result *err) { - return tri_scene_build_impl(vertices, ntris, opts, NULL, err); + return tri_scene_build_impl(vertices, ntris, opts, NULL, NULL, err); +} + +/* Indexed build: `vertices` is a unique-vertex array (3 floats each, nverts of + * them) and `indices` is 3*ntris vertex ids. Byte-identical to passing the + * de-indexed soup to lrt_tri_scene_build -- the build gathers each triangle's 3 + * vertices through `indices` (tri_fetch9) instead of reading a pre-expanded + * soup, so the caller need not materialize 9*ntris floats. prim_id in the BVH is + * still the triangle index 0..ntris-1, so per-triangle caller data (color, + * material, face map) indexes unchanged. The de-indexed leaf is identical, so + * lrt_tri_get_verts works the same. */ +lrt_tri_scene *lrt_tri_scene_build_indexed(const float *vertices, size_t nverts, + const uint32_t *indices, size_t ntris, + const lrt_tri_build_options *opts, + lrt_result *err) { + if (!vertices || !indices || ntris == 0) { + tri_set_err(err, LRT_RESULT_INVALID_ARGUMENT); + return NULL; + } + (void)nverts; /* bounds are the caller's contract; gather trusts indices */ + return tri_scene_build_impl(vertices, ntris, opts, NULL, indices, err); +} + +/* Batch build: build many independent scenes, parallelizing ACROSS scenes. This + * is the efficient path for two-level (TLAS) builds with many small BLAS — the + * intra-build parallel path only engages at >=4096 tris, so a fleet of small + * prototype BLAS would otherwise build one-at-a-time. Each scene is built + * single-threaded; the workers steal scenes one at a time (nchunks == n) so an + * uneven mix of prototype sizes still balances. out_scenes[i] is the built + * scene or NULL (empty input, i.e. vertices[i]==NULL or ntris[i]==0, or a build + * failure); errs[i] (optional) receives the per-scene result. */ +typedef struct lrt_batch_build_job { + const float *const *verts; + const size_t *ntris; + const lrt_tri_build_options *opts; + lrt_tri_scene **out; + lrt_result *errs; +} lrt_batch_build_job; + +static void lrt_batch_build_one(void *arg, unsigned chunk, uint32_t begin, + uint32_t end) { + lrt_batch_build_job *j = (lrt_batch_build_job *)arg; + uint32_t i; + (void)chunk; + for (i = begin; i < end; i++) { + lrt_result e = LRT_RESULT_OK; + lrt_tri_scene *s = NULL; + if (j->verts[i] && j->ntris[i] > 0) { + lrt_tri_build_options o; + if (j->opts) + o = *j->opts; + else + memset(&o, 0, sizeof(o)); + o.num_threads = 1u; /* parallelism is across scenes, not within */ + s = lrt_tri_scene_build(j->verts[i], j->ntris[i], &o, &e); + } + j->out[i] = s; + if (j->errs) j->errs[i] = e; + } +} + +void lrt_tri_scene_build_batch(const float *const *vertices, + const size_t *ntris, size_t n, + const lrt_tri_build_options *opts, + lrt_tri_scene **out_scenes, lrt_result *errs) { + size_t i; + lrt_batch_build_job job; + unsigned threads; + if (!out_scenes) return; + for (i = 0; i < n; i++) out_scenes[i] = NULL; + if (!vertices || !ntris || n == 0) return; + + job.verts = vertices; + job.ntris = ntris; + job.opts = opts; + job.out = out_scenes; + job.errs = errs; + threads = (opts && opts->num_threads) ? opts->num_threads : 1u; + +#if LRT_TRI_HAS_THREADS + if (threads > 1u && n > 1u && n <= 0xFFFFFFFFu) { + /* One chunk per scene -> the worker loop steals scenes atomically. */ + tri_pfor_job pjob; + thrd_t tids[TRI_PAR_MAX_THREADS]; + unsigned spawned = 0u, k, nw = threads; + if (nw > TRI_PAR_MAX_THREADS) nw = TRI_PAR_MAX_THREADS; + if ((size_t)nw > n) nw = (unsigned)n; + pjob.fn = lrt_batch_build_one; + pjob.arg = &job; + pjob.n = (uint32_t)n; + pjob.nchunks = (uint32_t)n; + atomic_init(&pjob.next, 0u); + for (k = 0; k + 1u < nw; k++) { + if (thrd_create(&tids[k], tri_pfor_worker, &pjob) != thrd_success) + break; + spawned++; + } + tri_pfor_worker(&pjob); /* the calling thread participates */ + for (k = 0; k < spawned; k++) thrd_join(tids[k], NULL); + return; + } +#endif + (void)threads; + lrt_batch_build_one(&job, 0u, 0u, (uint32_t)n); } /* GPU-assisted build hook used by lightrt_c_vk.c (declared extern there; not in @@ -8771,7 +9011,7 @@ lrt_tri_scene *lrt_tri_scene_build_lbvh_morton(const float *vertices, o.layout = layout; o.max_leaf_size = max_leaf_size; o.num_threads = 1; - return tri_scene_build_impl(vertices, ntris, &o, morton, err); + return tri_scene_build_impl(vertices, ntris, &o, morton, NULL, err); } /* Release a memory-mapped scene's backing store. The mmap path is implemented @@ -8798,6 +9038,7 @@ void lrt_tri_scene_free(lrt_tri_scene *s) { tri_aligned_free(s->nodes8q4); tri_aligned_free(s->blocks); } + free(s->prim2slot); /* always heap, even for mmap'd scenes (NULL if absent) */ free(s->owned_user); free(s->trim_loop_off); free(s->trim_pts); @@ -9584,7 +9825,7 @@ lrt_tri_scene *lrt_curve_scene_build(const float *segments, const float *radii, } } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && nsub >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(nsub * sizeof(uint32_t)); @@ -9725,7 +9966,14 @@ lrt_tri_scene *lrt_roundcurve_scene_build(const lrt_hair_strands *strands, bc.ntris = nseg; bc.max_leaf = max_leaf; bc.block_shift = 2u; /* round-linear blocks are 4-wide */ - bc.quality = LRT_TRI_BUILD_DEFAULT; + /* Honor the caller's quality. FAST routes through the parallel Morton LBVH + * below (same path the mesh FAST build uses) instead of the recursive, + * serial-top-level binned SAH -- the SAH build dominated curve-heavy scenes + * (Island isIronwoodA1: 3M hair segments, ~2s, barely scaling past ~2 cores) + * and its 2*nseg bnode arena drove peak RSS. LBVH is fully parallel and uses + * a much smaller frontier. Same closest hits, so renders are byte-identical. */ + bc.quality = (o.quality == LRT_TRI_BUILD_FAST) ? LRT_TRI_BUILD_FAST + : LRT_TRI_BUILD_DEFAULT; bc.emit_kind = TRI_PRIM_RLCURVE; bc.plo = (float *)malloc(nseg * 3 * sizeof(float)); bc.phi = (float *)malloc(nseg * 3 * sizeof(float)); @@ -9810,7 +10058,7 @@ lrt_tri_scene *lrt_roundcurve_scene_build(const lrt_hair_strands *strands, return NULL; } -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && nseg >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(nseg * sizeof(uint32_t)); @@ -9818,7 +10066,44 @@ lrt_tri_scene *lrt_roundcurve_scene_build(const lrt_hair_strands *strands, } #endif - uint32_t b_root = tri_build_binary(&bc, num_threads); + /* LBVH preprocessing (Morton sort) for the FAST path -- mirrors + * tri_scene_build_impl. Centroids (bc.cen) are already filled above. */ + uint64_t *lbvh_keys = NULL, *lbvh_tmp = NULL; + if (bc.quality == LRT_TRI_BUILD_FAST) { + lbvh_keys = (uint64_t *)malloc(nseg * sizeof(uint64_t)); + lbvh_tmp = (uint64_t *)malloc(nseg * sizeof(uint64_t)); + if (!lbvh_keys || !lbvh_tmp) { + free(lbvh_keys); + free(lbvh_tmp); + free(segs); + free(bc.plo); + free(bc.phi); + free(bc.cen); + free(bc.indices); + free(bc.bnodes); + free(bc.par_scratch); + tri_set_err(err, LRT_RESULT_OUT_OF_MEMORY); + return NULL; + } + tri_morton_encode(&bc, lbvh_keys, num_threads); + const uint64_t *sorted = tri_radix_sort_keys(lbvh_keys, lbvh_tmp, nseg); + for (size_t i = 0; i < nseg; i++) bc.indices[i] = (uint32_t)sorted[i]; + bc.lbvh_keys = sorted; + uint32_t bw = 1u << bc.block_shift; + bc.lbvh_leaf = 2u * bw; + if (bc.lbvh_leaf > max_leaf) bc.lbvh_leaf = max_leaf; + /* Centroids are only needed for the Morton encode (the LBVH build uses + * plo/phi + the sorted order); free them now to offset the Morton-key + * buffers (the SAH path keeps cen for binning, so free only here). */ + free(bc.cen); + bc.cen = NULL; + } + + uint32_t b_root = (bc.quality == LRT_TRI_BUILD_FAST) + ? tri_build_lbvh(&bc, num_threads) + : tri_build_binary(&bc, num_threads); + free(lbvh_keys); + free(lbvh_tmp); lrt_tri_scene *s = NULL; if (!bc.failed) { @@ -10004,7 +10289,7 @@ lrt_tri_scene *lrt_flatcurve_scene_build(const lrt_hair_strands *strands, lrt_tri_scene *s = NULL; { /* like tri_finish_aabb_build but with the flat block stride */ -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && nseg >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(nseg * sizeof(uint32_t)); @@ -10126,7 +10411,7 @@ lrt_tri_scene *lrt_bezcurve_scene_build(const float *cps, size_t nseg, } lrt_tri_scene *s = NULL; -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && nsub >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(nsub * sizeof(uint32_t)); @@ -10184,7 +10469,7 @@ lrt_tri_scene *lrt_bezcurve_scene_build(const float *cps, size_t nseg, static lrt_tri_scene *tri_finish_aabb_build(tri_build_ctx *bc, size_t nprims, int prim_kind, unsigned num_threads, lrt_result *err) { -#if !defined(__STDC_NO_THREADS__) +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && nprims >= 4096) { bc->par_threads = num_threads; bc->par_scratch = (uint32_t *)malloc(nprims * sizeof(uint32_t)); @@ -13122,16 +13407,12 @@ static uint32_t tlas_collapse(tlas_collapse_ctx *cc, uint32_t b_idx) { /* (Re)build the TLAS BVH over the instances' current world AABBs. Frees any * existing nodes/inst_refs first. Returns 0 on success. */ -static int tlas_rebuild(lrt_tlas *t, unsigned num_threads) { - size_t ninsts = t->ninsts; - tri_build_ctx bc; - memset(&bc, 0, sizeof(bc)); - bc.ntris = ninsts; - bc.max_leaf = 4u; - bc.block_shift = 2u; - bc.quality = LRT_TRI_BUILD_DEFAULT; - if (tri_alloc_aabb_scratch(&bc, ninsts)) return 1; - for (size_t i = 0; i < ninsts; i++) { +/* Compute world-space AABBs for instances [begin, end) into bc->plo/phi/cen. + * Each instance writes only its own [i*3, i*3+3) slots, so chunks are + * independent. */ +static void tlas_bounds_chunk(const lrt_tlas *t, tri_build_ctx *bc, + uint32_t begin, uint32_t end) { + for (uint32_t i = begin; i < end; i++) { const lrt_instance_int *in = &t->insts[i]; float wlo[3], whi[3]; const lrt_tri_scene *bl = t->blas[in->blas_id]; @@ -13146,13 +13427,49 @@ static int tlas_rebuild(lrt_tlas *t, unsigned num_threads) { tlas_transform_aabb(in->obj2world, bl->root_lo, bl->root_hi, wlo, whi); } for (int a = 0; a < 3; a++) { - bc.plo[i * 3 + a] = wlo[a]; - bc.phi[i * 3 + a] = whi[a]; - bc.cen[i * 3 + a] = 0.5f * (wlo[a] + whi[a]); + bc->plo[(size_t)i * 3 + a] = wlo[a]; + bc->phi[(size_t)i * 3 + a] = whi[a]; + bc->cen[(size_t)i * 3 + a] = 0.5f * (wlo[a] + whi[a]); } - bc.indices[i] = (uint32_t)i; + bc->indices[i] = i; } -#if !defined(__STDC_NO_THREADS__) +} + +#if LRT_TRI_HAS_THREADS +typedef struct tlas_bounds_job { + const lrt_tlas *t; + tri_build_ctx *bc; +} tlas_bounds_job; + +static void tlas_bounds_pfor(void *arg, unsigned chunk, uint32_t begin, + uint32_t end) { + (void)chunk; + tlas_bounds_job *j = (tlas_bounds_job *)arg; + tlas_bounds_chunk(j->t, j->bc, begin, end); +} +#endif + +static int tlas_rebuild(lrt_tlas *t, unsigned num_threads) { + size_t ninsts = t->ninsts; + tri_build_ctx bc; + memset(&bc, 0, sizeof(bc)); + bc.ntris = ninsts; + bc.max_leaf = 4u; + bc.block_shift = 2u; + bc.quality = LRT_TRI_BUILD_DEFAULT; + if (tri_alloc_aabb_scratch(&bc, ninsts)) return 1; +#if LRT_TRI_HAS_THREADS + if (num_threads > 1 && ninsts >= 4096) { + tlas_bounds_job j; + j.t = t; + j.bc = &bc; + tri_parallel_for((uint32_t)ninsts, num_threads, tlas_bounds_pfor, &j); + } else +#endif + { + tlas_bounds_chunk(t, &bc, 0u, (uint32_t)ninsts); + } +#if LRT_TRI_HAS_THREADS if (num_threads > 1 && ninsts >= 4096) { bc.par_threads = num_threads; bc.par_scratch = (uint32_t *)malloc(ninsts * sizeof(uint32_t)); @@ -13195,18 +13512,63 @@ static int tlas_rebuild(lrt_tlas *t, unsigned num_threads) { return failed; } -static int tlas_fill_instances(lrt_tlas *t, const lrt_instance *insts, - size_t ninsts) { - for (size_t i = 0; i < ninsts; i++) { +/* Fill t->insts[begin, end) from insts[]; returns 1 if any instance references + * a missing (out-of-range or NULL) BLAS. Per-instance work (a 3x4 affine copy + + * matrix inversion) is independent and indexed by i, so chunks never collide. */ +static int tlas_fill_chunk(lrt_tlas *t, const lrt_instance *insts, + uint32_t begin, uint32_t end) { + int bad = 0; + for (uint32_t i = begin; i < end; i++) { lrt_instance_int *in = &t->insts[i]; - if (insts[i].blas_id >= t->nblas) return 1; + /* Reject an instance that references a missing (out-of-range or NULL) + * BLAS. NULL entries that no instance references are allowed: the build + * and traversal only ever touch t->blas[in->blas_id] for live + * instances, so callers may pass a sparse blas[] array (e.g. with empty + * prototype slots) without compacting + remapping it first. */ + if (insts[i].blas_id >= t->nblas || !t->blas[insts[i].blas_id]) { + bad = 1; + continue; + } memcpy(in->obj2world, insts[i].obj2world, sizeof(in->obj2world)); in->blas_id = insts[i].blas_id; in->instance_id = insts[i].instance_id; in->mask = insts[i].mask; in->degenerate = tlas_affine_invert(in->obj2world, in->world2obj) ? 0u : 1u; } - return 0; + return bad; +} + +#if LRT_TRI_HAS_THREADS +typedef struct tlas_fill_job { + lrt_tlas *t; + const lrt_instance *insts; + atomic_int bad; +} tlas_fill_job; + +static void tlas_fill_pfor(void *arg, unsigned chunk, uint32_t begin, + uint32_t end) { + (void)chunk; + tlas_fill_job *j = (tlas_fill_job *)arg; + if (tlas_fill_chunk(j->t, j->insts, begin, end)) + atomic_store_explicit(&j->bad, 1, memory_order_relaxed); +} +#endif + +static int tlas_fill_instances(lrt_tlas *t, const lrt_instance *insts, + size_t ninsts, unsigned num_threads) { +#if LRT_TRI_HAS_THREADS + if (num_threads > 1 && ninsts >= 4096) { + tlas_fill_job j; + j.t = t; + j.insts = insts; + atomic_init(&j.bad, 0); + tri_parallel_for((uint32_t)ninsts, num_threads, tlas_fill_pfor, &j); + return atomic_load(&j.bad); + } +#else + (void)num_threads; +#endif + return tlas_fill_chunk(t, insts, 0u, (uint32_t)ninsts); } lrt_tlas *lrt_tlas_build(lrt_tri_scene *const *blas, size_t nblas, @@ -13217,16 +13579,9 @@ lrt_tlas *lrt_tlas_build(lrt_tri_scene *const *blas, size_t nblas, tri_set_err(err, LRT_RESULT_INVALID_ARGUMENT); return NULL; } - /* Validate BLAS pointers are non-NULL. */ - { - size_t i; - for (i = 0; i < nblas; i++) { - if (!blas[i]) { - tri_set_err(err, LRT_RESULT_INVALID_ARGUMENT); - return NULL; - } - } - } + /* Individual blas[] entries may be NULL as long as no instance references + * them (validated per-instance in tlas_fill_instances). This lets callers + * pass a sparse BLAS array without compacting it. */ unsigned num_threads = (opts && opts->num_threads) ? opts->num_threads : 1u; lrt_tlas *t = (lrt_tlas *)calloc(1, sizeof(lrt_tlas)); if (!t) { @@ -13245,7 +13600,7 @@ lrt_tlas *lrt_tlas_build(lrt_tri_scene *const *blas, size_t nblas, return NULL; } memcpy(t->blas, blas, nblas * sizeof(lrt_tri_scene *)); - if (tlas_fill_instances(t, insts, ninsts)) { + if (tlas_fill_instances(t, insts, ninsts, num_threads)) { free(t->insts); free(t); tri_set_err(err, LRT_RESULT_INVALID_ARGUMENT); @@ -13271,7 +13626,7 @@ void lrt_tlas_free(lrt_tlas *t) { lrt_result lrt_tlas_refit(lrt_tlas *t, const lrt_instance *insts, size_t ninsts) { if (!t || !insts || ninsts != t->ninsts) return LRT_RESULT_INVALID_ARGUMENT; - if (tlas_fill_instances(t, insts, ninsts)) return LRT_RESULT_INVALID_ARGUMENT; + if (tlas_fill_instances(t, insts, ninsts, 1u)) return LRT_RESULT_INVALID_ARGUMENT; return tlas_rebuild(t, 1u) ? LRT_RESULT_OUT_OF_MEMORY : LRT_RESULT_OK; } @@ -13571,7 +13926,7 @@ static void tri_intersect4_sse(const lrt_tri_scene *s, const lrt_ray4 *r, best_u = _mm_blendv_ps(best_u, uu, closer); best_v = _mm_blendv_ps(best_v, vv, closer); while (cm) { - int l = __builtin_ctz((unsigned)cm); + int l = lrt_ctz((unsigned)cm); cm &= cm - 1; best_prim[l] = ids[lane]; } diff --git a/lightrt_c_tri.h b/lightrt_c_tri.h index f94d467..1637790 100644 --- a/lightrt_c_tri.h +++ b/lightrt_c_tri.h @@ -88,6 +88,32 @@ lrt_tri_scene *lrt_tri_scene_build(const float *vertices, size_t ntris, const lrt_tri_build_options *opts, lrt_result *err); +/* Indexed build: `vertices` is `nverts` unique vertices (3 floats each) and + * `indices` is 3*ntris vertex ids (3 per triangle). Equivalent to de-indexing + * into a 9*ntris soup and calling lrt_tri_scene_build -- byte-identical result -- + * but the caller avoids materializing the soup (the build gathers each + * triangle's vertices through `indices`). prim_id stays 0..ntris-1 (the triangle + * index), so lrt_tri_get_verts and per-triangle caller data are unchanged. + * Returns NULL on failure (err set when non-NULL). */ +lrt_tri_scene *lrt_tri_scene_build_indexed(const float *vertices, size_t nverts, + const uint32_t *indices, size_t ntris, + const lrt_tri_build_options *opts, + lrt_result *err); + +/* Build `n` independent scenes at once, parallelizing ACROSS scenes (each scene + * is built single-threaded; workers steal scenes one at a time). This is the + * efficient path for two-level/TLAS builds with many small BLAS, where the + * per-scene intra-build threading (engaged only at >=4096 tris) would leave a + * fleet of small prototype scenes building serially. vertices[i]/ntris[i] + * describe scene i exactly as for lrt_tri_scene_build; a NULL vertices[i] or + * zero ntris[i] yields out_scenes[i]=NULL (a permitted empty slot — see + * lrt_tlas_build). opts->num_threads bounds the worker count. out_scenes must + * have room for n entries; errs (optional) receives the n per-scene results. */ +void lrt_tri_scene_build_batch(const float *const *vertices, + const size_t *ntris, size_t n, + const lrt_tri_build_options *opts, + lrt_tri_scene **out_scenes, lrt_result *errs); + void lrt_tri_scene_free(lrt_tri_scene *s); /* Closest hit. Returns 1 and fills *hit on a hit; returns 0 on miss (hit, if @@ -179,6 +205,22 @@ void lrt_tri_scene_stats(const lrt_tri_scene *s, lrt_tri_stats *out); * compiler SIMD flags. */ const char *lrt_tri_kernel_name(const lrt_tri_scene *s); +/* Recover a triangle's object-space vertices from the built BVH leaves, so a + * caller need not retain its own (9-float-per-triangle) vertex copy. Available + * only for plain triangle scenes (lrt_tri_scene_has_verts() != 0); returns 1 and + * fills v0/v1/v2 on success, 0 otherwise. The result is byte-exact with the input + * soup for typical mesh coordinates (leaves store v0 + edges; v0+edge round-trips + * exactly when |vk-v0| < 2|v0|, i.e. adjacent vertices, by Sterbenz's lemma). */ +int lrt_tri_scene_has_verts(const lrt_tri_scene *s); +int lrt_tri_get_verts(const lrt_tri_scene *s, uint32_t prim_id, float v0[3], + float v1[3], float v2[3]); + +/* Leaf slot for a triangle (BVH leaf-emission order), or LRT_TRI_NO_HIT, and the + * total slot count. Lets a caller reorder per-triangle shading into slot order + * for cache-coherent hit-time reads. Plain triangle scenes only (prim2slot). */ +uint32_t lrt_tri_get_slot(const lrt_tri_scene *s, uint32_t prim_id); +uint32_t lrt_tri_slot_count(const lrt_tri_scene *s); + /* --- Hair / curve scenes --------------------------------------------------- * * Axis-aligned boxes around long thin diagonal primitives are mostly empty @@ -825,7 +867,10 @@ typedef struct lrt_tlas_hit { uint32_t inst_id; /* instance_id of the hit instance */ } lrt_tlas_hit; -/* Build a TLAS. Instances with a (near-)singular transform are skipped. */ +/* Build a TLAS. Instances with a (near-)singular transform are skipped. + * Entries of blas[] may be NULL as long as no instance references that index, + * so a sparse BLAS array (e.g. with empty prototype slots) can be passed without + * compacting + remapping it. */ lrt_tlas *lrt_tlas_build(lrt_tri_scene *const *blas, size_t nblas, const lrt_instance *insts, size_t ninsts, const lrt_tri_build_options *opts, lrt_result *err); From 566b462ce7f75100fa58036fe2c2ea228e8590f9 Mon Sep 17 00:00:00 2001 From: Syoyo Fujita Date: Thu, 2 Jul 2026 02:44:46 +0900 Subject: [PATCH 3/3] d3d11: add Direct3D 11 compute trace backend (port from tinyusdz fork) New Windows-only GPU backend, mirror of the Vulkan "Path A" (CPU BVH build -> GPU trace) on a D3D11 compute shader instead of Vulkan. Motivation: on AMD GCN/Polaris the amdvlk Vulkan compute path mis-renders / hangs, while the D3D11 driver is solid; D3D11 is also ubiquitous on Windows and needs no SDK (d3d11 + d3dcompiler + dxgi ship with the OS). Batches a whole ray array into ONE dispatch, so a full-frame trace is a single GPU round-trip. - lightrt_c_d3d11.{cpp,h}: lrt_d3d11_engine_create/destroy + lrt_d3d11_trace_scene, same lrt_tri_scene / lrt_ray / lrt_hit surface as the Vulkan and CPU kernels. The whole .cpp is #if defined(_WIN32) -- an empty translation unit elsewhere. - d3d/shaders/trace_bvh.hlsl (+ gen_trace_bvh_hlsl.sh, trace_bvh_hlsl.h): the trace shader is decompiled from the Vulkan trace_bvh SPIR-V with SPIRV-Cross, so it walks the same serialized BVH4/BVH8 scene and returns identical hits. - CMake: LIGHTRT_BUILD_D3D11 option (OFF by default), builds a lightrt_d3d11 static lib linking d3d11/d3dcompiler/dxgi; on non-Windows it warns and skips. Verified on this Linux host only as far as the platform allows: the guarded TU compiles to an empty object (0 exported symbols), and CMake configures cleanly both by default and with -DLIGHTRT_BUILD_D3D11=ON (warn+skip). The actual D3D11 compile/link/run is UNVERIFIED here -- no Windows toolchain; it was validated in the tinyusdz fork (tusdrender -d3d). Co-Authored-By: Claude Opus 4.8 --- CMakeLists.txt | 18 ++ d3d/shaders/gen_trace_bvh_hlsl.sh | 61 +++++ d3d/shaders/trace_bvh.hlsl | 290 ++++++++++++++++++++++++ d3d/shaders/trace_bvh_hlsl.h | 294 ++++++++++++++++++++++++ lightrt_c_d3d11.cpp | 363 ++++++++++++++++++++++++++++++ lightrt_c_d3d11.h | 57 +++++ 6 files changed, 1083 insertions(+) create mode 100644 d3d/shaders/gen_trace_bvh_hlsl.sh create mode 100644 d3d/shaders/trace_bvh.hlsl create mode 100644 d3d/shaders/trace_bvh_hlsl.h create mode 100644 lightrt_c_d3d11.cpp create mode 100644 lightrt_c_d3d11.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 8454587..c28349b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -423,6 +423,24 @@ if(LIGHTRT_BUILD_CUDA) install(FILES lightrt_c_cuda.h DESTINATION include) endif() +option(LIGHTRT_BUILD_D3D11 "Build the Direct3D 11 compute backend (Windows only)" OFF) +if(LIGHTRT_BUILD_D3D11) + if(NOT WIN32) + message(WARNING "LIGHTRT_BUILD_D3D11 is Windows-only; lightrt_c_d3d11.cpp is " + "an empty translation unit elsewhere. Skipping the target.") + else() + # Static library exposing the D3D11 compute trace backend. Mirrors the + # Vulkan "Path A" (CPU build -> GPU trace) on a D3D11 compute shader + # (d3d/shaders/trace_bvh.hlsl, decompiled from the trace_bvh SPIR-V) and + # walks the same serialized BVH4/BVH8 scene. d3d11 + d3dcompiler + dxgi + # ship with the OS, so no SDK is required. + add_library(lightrt_d3d11 STATIC lightrt_c_d3d11.cpp) + target_include_directories(lightrt_d3d11 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) + target_link_libraries(lightrt_d3d11 PUBLIC d3d11 d3dcompiler dxgi) + install(FILES lightrt_c_d3d11.h DESTINATION include) + endif() +endif() + # Optional: regenerate the checked-in SPIR-V headers (developers with glslang # only; never part of the default build / ALL target). find_program(LIGHTRT_GLSLANG glslangValidator) diff --git a/d3d/shaders/gen_trace_bvh_hlsl.sh b/d3d/shaders/gen_trace_bvh_hlsl.sh new file mode 100644 index 0000000..72b6b8a --- /dev/null +++ b/d3d/shaders/gen_trace_bvh_hlsl.sh @@ -0,0 +1,61 @@ +#!/usr/bin/env bash +# +# Regenerate trace_bvh.hlsl + trace_bvh_hlsl.h from the committed Vulkan SPIR-V +# (../../vk/shaders/trace_bvh.spv.h) by decompiling it with SPIRV-Cross. +# +# The D3D11 backend (lightrt_c_d3d11.cpp) reuses the Vulkan trace_bvh compute +# shader verbatim — decompiling the SPIR-V guarantees the HLSL traverses the +# serialized BVH bit-for-bit like the GLSL/SPIR-V and the scalar CPU kernel, so +# there is no second shader to keep in sync. Run this after editing the Vulkan +# shader and regenerating vk/shaders/trace_bvh.spv.h (compile_shaders.sh). +# +# Needs spirv-cross on PATH or in $SPIRV_CROSS: +# git clone https://github.com/KhronosGroup/SPIRV-Cross && cmake build it, +# or install the Vulkan SDK (ships spirv-cross.exe). +# +# Usage: SPIRV_CROSS=/path/to/spirv-cross ./gen_trace_bvh_hlsl.sh +set -euo pipefail + +HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +SPV_H="${HERE}/../../vk/shaders/trace_bvh.spv.h" +OUT_HLSL="${HERE}/trace_bvh.hlsl" +OUT_H="${HERE}/trace_bvh_hlsl.h" + +SC="${SPIRV_CROSS:-$(command -v spirv-cross || true)}" +[ -z "${SC}" ] && { echo "error: spirv-cross not found (set \$SPIRV_CROSS)" >&2; exit 1; } + +# Reconstruct the raw .spv from the C uint32 array in the header (little-endian). +TMP_SPV="$(mktemp -t trace_bvh.XXXXXX.spv)" +trap 'rm -f "${TMP_SPV}"' EXIT +grep -oE '0x[0-9a-fA-F]+' "${SPV_H}" | \ + perl -ne 'chomp; print pack("V", hex($_));' > "${TMP_SPV}" + +# Decompile to HLSL shader model 5.0 (D3D11). Spec constants W / STACK become +# SPIRV_CROSS_CONSTANT_ID_0 / _1 macros (set per-scene by D3DCompile at runtime). +PROVENANCE="$(sed -n '1,12p' "${OUT_HLSL}" 2>/dev/null | grep -c '^//' || true)" +{ + cat <<'EOF' +// trace_bvh.hlsl — Direct3D 11 (SM5.0) compute BVH traversal for LightRT. +// +// DECOMPILED from vk/shaders/trace_bvh.spv (the committed Vulkan SPIR-V) with +// SPIRV-Cross (--hlsl --shader-model 50). It mirrors the scalar CPU kernel and +// the Vulkan trace_bvh compute shader bit-for-bit. Regenerate after editing the +// GLSL/SPIR-V via d3d/shaders/gen_trace_bvh_hlsl.sh. +// +// Bindings: t0 = BVH nodes, t1 = leaf blocks, t2 = rays, u3 = hits (RW), +// cbuffer PC = {root, node_count, block_count, ray_count}. Spec constants +// W (BVH width 4/8) and STACK are set via SPIRV_CROSS_CONSTANT_ID_0/_1 macros. +EOF + "${SC}" --hlsl --shader-model 50 "${TMP_SPV}" +} > "${OUT_HLSL}" + +# Embed as a raw string literal (the source is < 16 KB, the MSVC literal limit). +{ + echo "// Auto-generated from trace_bvh.hlsl. Do not edit by hand." + echo "#pragma once" + echo "static const char trace_bvh_hlsl[] = R\"HLSLSRC(" + cat "${OUT_HLSL}" + echo ")HLSLSRC\";" +} > "${OUT_H}" + +echo "generated ${OUT_HLSL} and ${OUT_H}" diff --git a/d3d/shaders/trace_bvh.hlsl b/d3d/shaders/trace_bvh.hlsl new file mode 100644 index 0000000..c8b7e71 --- /dev/null +++ b/d3d/shaders/trace_bvh.hlsl @@ -0,0 +1,290 @@ +// trace_bvh.hlsl — Direct3D 11 (SM5.0) compute BVH traversal for LightRT. +// +// DECOMPILED from vk/shaders/trace_bvh.spv (the committed Vulkan SPIR-V) with +// SPIRV-Cross (--hlsl --shader-model 50). It mirrors the scalar CPU kernel and +// the Vulkan trace_bvh compute shader bit-for-bit. Regenerate after editing the +// GLSL/SPIR-V via d3d/shaders/gen_trace_bvh_hlsl.sh. +// +// Bindings: t0 = BVH nodes, t1 = leaf blocks, t2 = rays, u3 = hits (RW), +// cbuffer PC = {root, node_count, block_count, ray_count}. Spec constants +// W (BVH width 4/8) and STACK are set via SPIRV_CROSS_CONSTANT_ID_0/_1 macros. +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 8u +#endif +static const uint W = SPIRV_CROSS_CONSTANT_ID_0; +static const uint _164 = (8u * W); +static const uint _167 = (10u * W); +static const uint _170 = (9u * W); +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 96u +#endif +static const uint STACK = SPIRV_CROSS_CONSTANT_ID_1; +static const uint _264 = (0u * W); +static const uint _274 = (1u * W); +static const uint _284 = (2u * W); +static const uint _294 = (3u * W); +static const uint _304 = (4u * W); +static const uint _314 = (5u * W); +static const uint _324 = (6u * W); +static const uint _334 = (7u * W); +static const uint _344 = (8u * W); +static const uint _542 = (7u * W); +static const uint _558 = (0u * W); +static const uint _568 = (1u * W); +static const uint _578 = (2u * W); +static const uint _588 = (3u * W); +static const uint _598 = (4u * W); +static const uint _608 = (5u * W); +static const uint _753 = (6u * W); +static const uint3 gl_WorkGroupSize = uint3(64u, 1u, 1u); + +ByteAddressBuffer _49 : register(t2); +ByteAddressBuffer _249 : register(t1); +ByteAddressBuffer _540 : register(t0); +RWByteAddressBuffer _801 : register(u3); +cbuffer PC +{ + uint pc_root : packoffset(c0); + uint pc_node_count : packoffset(c0.y); + uint pc_block_count : packoffset(c0.z); + uint pc_ray_count : packoffset(c0.w); +}; + + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +float fu(uint i) +{ + return asfloat(i); +} + +void comp_main() +{ + uint gid = gl_GlobalInvocationID.x; + if (gid >= pc_ray_count) + { + return; + } + uint rb = gid * 8u; + uint param = _49.Load((rb + 0u) * 4 + 0); + uint param_1 = _49.Load((rb + 1u) * 4 + 0); + uint param_2 = _49.Load((rb + 2u) * 4 + 0); + float3 org = float3(fu(param), fu(param_1), fu(param_2)); + uint param_3 = _49.Load((rb + 3u) * 4 + 0); + float tmin = fu(param_3); + uint param_4 = _49.Load((rb + 4u) * 4 + 0); + uint param_5 = _49.Load((rb + 5u) * 4 + 0); + uint param_6 = _49.Load((rb + 6u) * 4 + 0); + float3 dir = float3(fu(param_4), fu(param_5), fu(param_6)); + uint param_7 = _49.Load((rb + 7u) * 4 + 0); + float tmax = fu(param_7); + float3 invd; + for (int k = 0; k < 3; k++) + { + float d = dir[k]; + float inv = 1.0f / d; + if (!((inv >= (-999999984306749440.0f)) && (inv <= 999999984306749440.0f))) + { + float s = (d == 0.0f) ? 1.0f : d; + inv = (s < 0.0f) ? (-999999984306749440.0f) : 999999984306749440.0f; + } + invd[k] = inv; + } + float best_t = tmax; + float best_u = 0.0f; + float best_v = 0.0f; + uint best_prim = 4294967295u; + uint node_stride = _164; + uint block_stride = _167; + uint prim_off = _170; + int sp = 0; + uint stk_ref[STACK]; + stk_ref[0] = pc_root; + float stk_tn[STACK]; + stk_tn[0] = tmin; + sp = 1; + float hit_tn[8]; + uint hit_ref[8]; + while (sp > 0) + { + sp--; + uint ref = stk_ref[sp]; + float tnear_e = stk_tn[sp]; + if (tnear_e >= best_t) + { + continue; + } + if ((ref & 2147483648u) != 0u) + { + uint blk0 = (ref & 2147483647u) >> 4u; + uint nblk = ref & 15u; + for (uint b = 0u; b < nblk; b++) + { + uint bb = (blk0 + b) * block_stride; + for (uint lane = 0u; lane < W; lane++) + { + uint prim = _249.Load(((bb + prim_off) + lane) * 4 + 0); + if (prim == 4294967295u) + { + continue; + } + uint param_8 = _249.Load(((bb + _264) + lane) * 4 + 0); + float v0x = fu(param_8); + uint param_9 = _249.Load(((bb + _274) + lane) * 4 + 0); + float v0y = fu(param_9); + uint param_10 = _249.Load(((bb + _284) + lane) * 4 + 0); + float v0z = fu(param_10); + uint param_11 = _249.Load(((bb + _294) + lane) * 4 + 0); + float e1x = fu(param_11); + uint param_12 = _249.Load(((bb + _304) + lane) * 4 + 0); + float e1y = fu(param_12); + uint param_13 = _249.Load(((bb + _314) + lane) * 4 + 0); + float e1z = fu(param_13); + uint param_14 = _249.Load(((bb + _324) + lane) * 4 + 0); + float e2x = fu(param_14); + uint param_15 = _249.Load(((bb + _334) + lane) * 4 + 0); + float e2y = fu(param_15); + uint param_16 = _249.Load(((bb + _344) + lane) * 4 + 0); + float e2z = fu(param_16); + float px = (dir.y * e2z) - (dir.z * e2y); + float py = (dir.z * e2x) - (dir.x * e2z); + float pz = (dir.x * e2y) - (dir.y * e2x); + float det = ((e1x * px) + (e1y * py)) + (e1z * pz); + if ((det > (-9.9999999600419720025001879548654e-13f)) && (det < 9.9999999600419720025001879548654e-13f)) + { + continue; + } + float inv_det = 1.0f / det; + float tvx = org.x - v0x; + float tvy = org.y - v0y; + float tvz = org.z - v0z; + float uu = (((tvx * px) + (tvy * py)) + (tvz * pz)) * inv_det; + if ((uu < 0.0f) || (uu > 1.0f)) + { + continue; + } + float qx = (tvy * e1z) - (tvz * e1y); + float qy = (tvz * e1x) - (tvx * e1z); + float qz = (tvx * e1y) - (tvy * e1x); + float vv = (((dir.x * qx) + (dir.y * qy)) + (dir.z * qz)) * inv_det; + bool _486 = vv < 0.0f; + bool _494; + if (!_486) + { + _494 = (uu + vv) > 1.0f; + } + else + { + _494 = _486; + } + if (_494) + { + continue; + } + float tt = (((e2x * qx) + (e2y * qy)) + (e2z * qz)) * inv_det; + if ((tt < tmin) || (tt >= best_t)) + { + continue; + } + best_t = tt; + best_u = uu; + best_v = vv; + best_prim = prim; + } + } + continue; + } + uint nb = (ref & 2147483647u) * node_stride; + uint nchildren = _540.Load((nb + _542) * 4 + 0); + int nhit = 0; + for (uint i = 0u; i < nchildren; i++) + { + uint param_17 = _540.Load(((nb + _558) + i) * 4 + 0); + float lo_x = fu(param_17); + uint param_18 = _540.Load(((nb + _568) + i) * 4 + 0); + float lo_y = fu(param_18); + uint param_19 = _540.Load(((nb + _578) + i) * 4 + 0); + float lo_z = fu(param_19); + uint param_20 = _540.Load(((nb + _588) + i) * 4 + 0); + float hi_x = fu(param_20); + uint param_21 = _540.Load(((nb + _598) + i) * 4 + 0); + float hi_y = fu(param_21); + uint param_22 = _540.Load(((nb + _608) + i) * 4 + 0); + float hi_z = fu(param_22); + float tlx = (lo_x - org.x) * invd.x; + float thx = (hi_x - org.x) * invd.x; + float tly = (lo_y - org.y) * invd.y; + float thy = (hi_y - org.y) * invd.y; + float tlz = (lo_z - org.z) * invd.z; + float thz = (hi_z - org.z) * invd.z; + float tnx = min(tlx, thx); + float tfx = max(tlx, thx); + float tny = min(tly, thy); + float tfy = max(tly, thy); + float tnz = min(tlz, thz); + float tfz = max(tlz, thz); + float tnear = max(max(tnx, tny), max(tnz, tmin)); + float tfar = min(min(tfx, tfy), min(tfz, best_t)); + if (tnear <= tfar) + { + int _710 = nhit; + nhit = _710 + 1; + int j = _710; + for (;;) + { + bool _718 = j > 0; + bool _730; + if (_718) + { + _730 = hit_tn[j - 1] > tnear; + } + else + { + _730 = _718; + } + if (_730) + { + hit_tn[j] = hit_tn[j - 1]; + hit_ref[j] = hit_ref[j - 1]; + j--; + continue; + } + else + { + break; + } + } + hit_tn[j] = tnear; + hit_ref[j] = _540.Load(((nb + _753) + i) * 4 + 0); + } + } + if ((uint(sp) + uint(nhit)) > STACK) + { + best_prim = 4294967295u; + break; + } + int _773 = nhit - 1; + for (int i_1 = _773; i_1 >= 0; i_1--) + { + stk_ref[sp] = hit_ref[i_1]; + stk_tn[sp] = hit_tn[i_1]; + sp++; + } + } + uint hb = gid * 4u; + _801.Store((hb + 0u) * 4 + 0, asuint((best_prim != 4294967295u) ? best_t : 0.0f)); + _801.Store((hb + 1u) * 4 + 0, asuint(best_u)); + _801.Store((hb + 2u) * 4 + 0, asuint(best_v)); + _801.Store((hb + 3u) * 4 + 0, best_prim); +} + +[numthreads(64, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/d3d/shaders/trace_bvh_hlsl.h b/d3d/shaders/trace_bvh_hlsl.h new file mode 100644 index 0000000..6b59b08 --- /dev/null +++ b/d3d/shaders/trace_bvh_hlsl.h @@ -0,0 +1,294 @@ +// Auto-generated from trace_bvh.hlsl. Do not edit by hand. +#pragma once +static const char trace_bvh_hlsl[] = R"HLSLSRC( +// trace_bvh.hlsl — Direct3D 11 (SM5.0) compute BVH traversal for LightRT. +// +// DECOMPILED from vk/shaders/trace_bvh.spv (the committed Vulkan SPIR-V) with +// SPIRV-Cross (--hlsl --shader-model 50). It mirrors the scalar CPU kernel and +// the Vulkan trace_bvh compute shader bit-for-bit. Regenerate after editing the +// GLSL/SPIR-V via d3d/shaders/gen_trace_bvh_hlsl.sh. +// +// Bindings: t0 = BVH nodes, t1 = leaf blocks, t2 = rays, u3 = hits (RW), +// cbuffer PC = {root, node_count, block_count, ray_count}. Spec constants +// W (BVH width 4/8) and STACK are set via SPIRV_CROSS_CONSTANT_ID_0/_1 macros. +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 8u +#endif +static const uint W = SPIRV_CROSS_CONSTANT_ID_0; +static const uint _164 = (8u * W); +static const uint _167 = (10u * W); +static const uint _170 = (9u * W); +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 96u +#endif +static const uint STACK = SPIRV_CROSS_CONSTANT_ID_1; +static const uint _264 = (0u * W); +static const uint _274 = (1u * W); +static const uint _284 = (2u * W); +static const uint _294 = (3u * W); +static const uint _304 = (4u * W); +static const uint _314 = (5u * W); +static const uint _324 = (6u * W); +static const uint _334 = (7u * W); +static const uint _344 = (8u * W); +static const uint _542 = (7u * W); +static const uint _558 = (0u * W); +static const uint _568 = (1u * W); +static const uint _578 = (2u * W); +static const uint _588 = (3u * W); +static const uint _598 = (4u * W); +static const uint _608 = (5u * W); +static const uint _753 = (6u * W); +static const uint3 gl_WorkGroupSize = uint3(64u, 1u, 1u); + +ByteAddressBuffer _49 : register(t2); +ByteAddressBuffer _249 : register(t1); +ByteAddressBuffer _540 : register(t0); +RWByteAddressBuffer _801 : register(u3); +cbuffer PC +{ + uint pc_root : packoffset(c0); + uint pc_node_count : packoffset(c0.y); + uint pc_block_count : packoffset(c0.z); + uint pc_ray_count : packoffset(c0.w); +}; + + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +float fu(uint i) +{ + return asfloat(i); +} + +void comp_main() +{ + uint gid = gl_GlobalInvocationID.x; + if (gid >= pc_ray_count) + { + return; + } + uint rb = gid * 8u; + uint param = _49.Load((rb + 0u) * 4 + 0); + uint param_1 = _49.Load((rb + 1u) * 4 + 0); + uint param_2 = _49.Load((rb + 2u) * 4 + 0); + float3 org = float3(fu(param), fu(param_1), fu(param_2)); + uint param_3 = _49.Load((rb + 3u) * 4 + 0); + float tmin = fu(param_3); + uint param_4 = _49.Load((rb + 4u) * 4 + 0); + uint param_5 = _49.Load((rb + 5u) * 4 + 0); + uint param_6 = _49.Load((rb + 6u) * 4 + 0); + float3 dir = float3(fu(param_4), fu(param_5), fu(param_6)); + uint param_7 = _49.Load((rb + 7u) * 4 + 0); + float tmax = fu(param_7); + float3 invd; + for (int k = 0; k < 3; k++) + { + float d = dir[k]; + float inv = 1.0f / d; + if (!((inv >= (-999999984306749440.0f)) && (inv <= 999999984306749440.0f))) + { + float s = (d == 0.0f) ? 1.0f : d; + inv = (s < 0.0f) ? (-999999984306749440.0f) : 999999984306749440.0f; + } + invd[k] = inv; + } + float best_t = tmax; + float best_u = 0.0f; + float best_v = 0.0f; + uint best_prim = 4294967295u; + uint node_stride = _164; + uint block_stride = _167; + uint prim_off = _170; + int sp = 0; + uint stk_ref[STACK]; + stk_ref[0] = pc_root; + float stk_tn[STACK]; + stk_tn[0] = tmin; + sp = 1; + float hit_tn[8]; + uint hit_ref[8]; + while (sp > 0) + { + sp--; + uint ref = stk_ref[sp]; + float tnear_e = stk_tn[sp]; + if (tnear_e >= best_t) + { + continue; + } + if ((ref & 2147483648u) != 0u) + { + uint blk0 = (ref & 2147483647u) >> 4u; + uint nblk = ref & 15u; + for (uint b = 0u; b < nblk; b++) + { + uint bb = (blk0 + b) * block_stride; + for (uint lane = 0u; lane < W; lane++) + { + uint prim = _249.Load(((bb + prim_off) + lane) * 4 + 0); + if (prim == 4294967295u) + { + continue; + } + uint param_8 = _249.Load(((bb + _264) + lane) * 4 + 0); + float v0x = fu(param_8); + uint param_9 = _249.Load(((bb + _274) + lane) * 4 + 0); + float v0y = fu(param_9); + uint param_10 = _249.Load(((bb + _284) + lane) * 4 + 0); + float v0z = fu(param_10); + uint param_11 = _249.Load(((bb + _294) + lane) * 4 + 0); + float e1x = fu(param_11); + uint param_12 = _249.Load(((bb + _304) + lane) * 4 + 0); + float e1y = fu(param_12); + uint param_13 = _249.Load(((bb + _314) + lane) * 4 + 0); + float e1z = fu(param_13); + uint param_14 = _249.Load(((bb + _324) + lane) * 4 + 0); + float e2x = fu(param_14); + uint param_15 = _249.Load(((bb + _334) + lane) * 4 + 0); + float e2y = fu(param_15); + uint param_16 = _249.Load(((bb + _344) + lane) * 4 + 0); + float e2z = fu(param_16); + float px = (dir.y * e2z) - (dir.z * e2y); + float py = (dir.z * e2x) - (dir.x * e2z); + float pz = (dir.x * e2y) - (dir.y * e2x); + float det = ((e1x * px) + (e1y * py)) + (e1z * pz); + if ((det > (-9.9999999600419720025001879548654e-13f)) && (det < 9.9999999600419720025001879548654e-13f)) + { + continue; + } + float inv_det = 1.0f / det; + float tvx = org.x - v0x; + float tvy = org.y - v0y; + float tvz = org.z - v0z; + float uu = (((tvx * px) + (tvy * py)) + (tvz * pz)) * inv_det; + if ((uu < 0.0f) || (uu > 1.0f)) + { + continue; + } + float qx = (tvy * e1z) - (tvz * e1y); + float qy = (tvz * e1x) - (tvx * e1z); + float qz = (tvx * e1y) - (tvy * e1x); + float vv = (((dir.x * qx) + (dir.y * qy)) + (dir.z * qz)) * inv_det; + bool _486 = vv < 0.0f; + bool _494; + if (!_486) + { + _494 = (uu + vv) > 1.0f; + } + else + { + _494 = _486; + } + if (_494) + { + continue; + } + float tt = (((e2x * qx) + (e2y * qy)) + (e2z * qz)) * inv_det; + if ((tt < tmin) || (tt >= best_t)) + { + continue; + } + best_t = tt; + best_u = uu; + best_v = vv; + best_prim = prim; + } + } + continue; + } + uint nb = (ref & 2147483647u) * node_stride; + uint nchildren = _540.Load((nb + _542) * 4 + 0); + int nhit = 0; + for (uint i = 0u; i < nchildren; i++) + { + uint param_17 = _540.Load(((nb + _558) + i) * 4 + 0); + float lo_x = fu(param_17); + uint param_18 = _540.Load(((nb + _568) + i) * 4 + 0); + float lo_y = fu(param_18); + uint param_19 = _540.Load(((nb + _578) + i) * 4 + 0); + float lo_z = fu(param_19); + uint param_20 = _540.Load(((nb + _588) + i) * 4 + 0); + float hi_x = fu(param_20); + uint param_21 = _540.Load(((nb + _598) + i) * 4 + 0); + float hi_y = fu(param_21); + uint param_22 = _540.Load(((nb + _608) + i) * 4 + 0); + float hi_z = fu(param_22); + float tlx = (lo_x - org.x) * invd.x; + float thx = (hi_x - org.x) * invd.x; + float tly = (lo_y - org.y) * invd.y; + float thy = (hi_y - org.y) * invd.y; + float tlz = (lo_z - org.z) * invd.z; + float thz = (hi_z - org.z) * invd.z; + float tnx = min(tlx, thx); + float tfx = max(tlx, thx); + float tny = min(tly, thy); + float tfy = max(tly, thy); + float tnz = min(tlz, thz); + float tfz = max(tlz, thz); + float tnear = max(max(tnx, tny), max(tnz, tmin)); + float tfar = min(min(tfx, tfy), min(tfz, best_t)); + if (tnear <= tfar) + { + int _710 = nhit; + nhit = _710 + 1; + int j = _710; + for (;;) + { + bool _718 = j > 0; + bool _730; + if (_718) + { + _730 = hit_tn[j - 1] > tnear; + } + else + { + _730 = _718; + } + if (_730) + { + hit_tn[j] = hit_tn[j - 1]; + hit_ref[j] = hit_ref[j - 1]; + j--; + continue; + } + else + { + break; + } + } + hit_tn[j] = tnear; + hit_ref[j] = _540.Load(((nb + _753) + i) * 4 + 0); + } + } + if ((uint(sp) + uint(nhit)) > STACK) + { + best_prim = 4294967295u; + break; + } + int _773 = nhit - 1; + for (int i_1 = _773; i_1 >= 0; i_1--) + { + stk_ref[sp] = hit_ref[i_1]; + stk_tn[sp] = hit_tn[i_1]; + sp++; + } + } + uint hb = gid * 4u; + _801.Store((hb + 0u) * 4 + 0, asuint((best_prim != 4294967295u) ? best_t : 0.0f)); + _801.Store((hb + 1u) * 4 + 0, asuint(best_u)); + _801.Store((hb + 2u) * 4 + 0, asuint(best_v)); + _801.Store((hb + 3u) * 4 + 0, best_prim); +} + +[numthreads(64, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} +)HLSLSRC"; diff --git a/lightrt_c_d3d11.cpp b/lightrt_c_d3d11.cpp new file mode 100644 index 0000000..1bf8aa6 --- /dev/null +++ b/lightrt_c_d3d11.cpp @@ -0,0 +1,363 @@ +// SPDX-License-Identifier: Apache-2.0 +// lightrt_c_d3d11.cpp — Direct3D 11 compute trace backend for LightRT. +// See lightrt_c_d3d11.h. Windows-only; empty TU elsewhere. + +#if defined(_WIN32) + +#include "lightrt_c_d3d11.h" + +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "d3d/shaders/trace_bvh_hlsl.h" // const char trace_bvh_hlsl[] + +namespace { + +// LRTS serialization header — must match lightrt_c_tri.c / lightrt_c_vk.c. +struct lrts_header { + char magic[4]; + uint32_t version, endian, flags, layout, prim_kind; + uint32_t node_count, block_count, root, node_stride, block_stride, reserved0; + float root_lo[3], root_hi[3]; + uint64_t node_offset, block_offset, file_size; +}; + +// cbuffer PC { root, node_count, block_count, ray_count } (16 bytes) +struct TracePC { + uint32_t root, node_count, block_count, ray_count; +}; + +// Bucket the traversal stack like the Vulkan path (so shaders are reused). +uint32_t trace_stack_for(uint32_t max_depth, uint32_t w) { + uint32_t need = max_depth * (w - 1u) + w + 1u; + static const uint32_t buckets[] = {32u, 64u, 128u, 256u}; + for (int i = 0; i < 4; i++) + if (need <= buckets[i]) return buckets[i]; + return 0; // too deep for the compute stack +} + +template +void safe_release(T*& p) { if (p) { p->Release(); p = nullptr; } } + +// One cached compute shader keyed by (BVH width, stack depth). +struct ShaderSlot { + uint32_t w = 0, stack = 0; + ID3D11ComputeShader* cs = nullptr; +}; + +} // namespace + +struct lrt_d3d11_engine { + ID3D11Device* dev = nullptr; + ID3D11DeviceContext* ctx = nullptr; + std::string name; + std::string last_error; + ShaderSlot slots[8]; +}; + +static void set_err(lrt_d3d11_engine* e, const char* msg) { + if (e) e->last_error = msg ? msg : ""; +} + +extern "C" lrt_d3d11_engine* lrt_d3d11_engine_create(int prefer_discrete, + lrt_result* err) { + lrt_d3d11_engine* e = new (std::nothrow) lrt_d3d11_engine(); + if (!e) { if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return nullptr; } + + // Pick an adapter (optionally a discrete one). Enumerate via DXGI. + IDXGIFactory* factory = nullptr; + IDXGIAdapter* chosen = nullptr; + if (SUCCEEDED(CreateDXGIFactory(__uuidof(IDXGIFactory), + reinterpret_cast(&factory)))) { + IDXGIAdapter* a = nullptr; + SIZE_T best_vram = 0; + for (UINT i = 0; factory->EnumAdapters(i, &a) != DXGI_ERROR_NOT_FOUND; ++i) { + DXGI_ADAPTER_DESC d{}; + a->GetDesc(&d); + bool sw = (d.VendorId == 0x1414); // Microsoft Basic Render (WARP) + if (!chosen || (prefer_discrete && !sw && + d.DedicatedVideoMemory > best_vram)) { + safe_release(chosen); + chosen = a; // keep (ref held by enum); take ownership + best_vram = d.DedicatedVideoMemory; + continue; + } + a->Release(); + } + factory->Release(); + } + + D3D_FEATURE_LEVEL want[] = {D3D_FEATURE_LEVEL_11_1, D3D_FEATURE_LEVEL_11_0}; + D3D_FEATURE_LEVEL got{}; + UINT flags = 0; + HRESULT hr = D3D11CreateDevice( + chosen, chosen ? D3D_DRIVER_TYPE_UNKNOWN : D3D_DRIVER_TYPE_HARDWARE, + nullptr, flags, want, 2, D3D11_SDK_VERSION, &e->dev, &got, &e->ctx); + if (FAILED(hr)) { + hr = D3D11CreateDevice(nullptr, D3D_DRIVER_TYPE_HARDWARE, nullptr, flags, + want, 2, D3D11_SDK_VERSION, &e->dev, &got, &e->ctx); + } + safe_release(chosen); + if (FAILED(hr) || !e->dev) { + set_err(e, "D3D11CreateDevice failed"); + if (err) *err = LRT_RESULT_NOT_BUILT; + delete e; + return nullptr; + } + + // Device name via the device's own adapter. + IDXGIDevice* dxdev = nullptr; + if (SUCCEEDED(e->dev->QueryInterface(__uuidof(IDXGIDevice), + reinterpret_cast(&dxdev)))) { + IDXGIAdapter* ad = nullptr; + if (SUCCEEDED(dxdev->GetAdapter(&ad))) { + DXGI_ADAPTER_DESC d{}; + ad->GetDesc(&d); + char buf[160]; + WideCharToMultiByte(CP_UTF8, 0, d.Description, -1, buf, sizeof(buf), + nullptr, nullptr); + e->name = buf; + ad->Release(); + } + dxdev->Release(); + } + if (err) *err = LRT_RESULT_OK; + return e; +} + +extern "C" void lrt_d3d11_engine_destroy(lrt_d3d11_engine* e) { + if (!e) return; + for (auto& s : e->slots) safe_release(s.cs); + safe_release(e->ctx); + safe_release(e->dev); + delete e; +} + +extern "C" const char* lrt_d3d11_engine_device_name(const lrt_d3d11_engine* e) { + return e ? e->name.c_str() : ""; +} +extern "C" const char* lrt_d3d11_engine_last_error(const lrt_d3d11_engine* e) { + return e ? e->last_error.c_str() : ""; +} + +// Compile (and cache) the trace compute shader for a given BVH width + stack. +static ID3D11ComputeShader* get_shader(lrt_d3d11_engine* e, uint32_t w, + uint32_t stack) { + for (auto& s : e->slots) + if (s.cs && s.w == w && s.stack == stack) return s.cs; + + char w_s[16], stk_s[16]; + std::snprintf(w_s, sizeof(w_s), "%uu", w); + std::snprintf(stk_s, sizeof(stk_s), "%uu", stack); + D3D_SHADER_MACRO macros[] = { + {"SPIRV_CROSS_CONSTANT_ID_0", w_s}, // W (BVH width) + {"SPIRV_CROSS_CONSTANT_ID_1", stk_s}, // STACK + {nullptr, nullptr}}; + ID3DBlob* code = nullptr; + ID3DBlob* errs = nullptr; + HRESULT hr = D3DCompile(trace_bvh_hlsl, std::strlen(trace_bvh_hlsl), + "trace_bvh.hlsl", macros, nullptr, "main", "cs_5_0", + 0, 0, &code, &errs); + if (FAILED(hr)) { + std::string m = "D3DCompile(trace_bvh) failed"; + if (errs) { m += ": "; m.append((const char*)errs->GetBufferPointer()); } + set_err(e, m.c_str()); + safe_release(errs); + safe_release(code); + return nullptr; + } + safe_release(errs); + ID3D11ComputeShader* cs = nullptr; + hr = e->dev->CreateComputeShader(code->GetBufferPointer(), + code->GetBufferSize(), nullptr, &cs); + safe_release(code); + if (FAILED(hr)) { set_err(e, "CreateComputeShader failed"); return nullptr; } + + // Insert into the cache (evict slot 0 if full). + for (auto& s : e->slots) { + if (!s.cs) { s.cs = cs; s.w = w; s.stack = stack; return cs; } + } + safe_release(e->slots[0].cs); + e->slots[0] = {w, stack, cs}; + return cs; +} + +// Create a raw (ByteAddressBuffer) buffer, optionally with an SRV and/or UAV. +static ID3D11Buffer* make_raw_buffer(ID3D11Device* dev, const void* data, + uint32_t bytes, bool uav, + ID3D11ShaderResourceView** srv, + ID3D11UnorderedAccessView** uav_out) { + D3D11_BUFFER_DESC bd{}; + bd.ByteWidth = (bytes + 3u) & ~3u; // 4-byte aligned + bd.Usage = D3D11_USAGE_DEFAULT; + bd.BindFlags = D3D11_BIND_SHADER_RESOURCE | (uav ? D3D11_BIND_UNORDERED_ACCESS : 0); + bd.MiscFlags = D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS; + D3D11_SUBRESOURCE_DATA init{}; + init.pSysMem = data; + ID3D11Buffer* buf = nullptr; + if (FAILED(dev->CreateBuffer(&bd, data ? &init : nullptr, &buf))) return nullptr; + UINT n4 = bd.ByteWidth / 4u; + if (srv) { + D3D11_SHADER_RESOURCE_VIEW_DESC sd{}; + sd.Format = DXGI_FORMAT_R32_TYPELESS; + sd.ViewDimension = D3D11_SRV_DIMENSION_BUFFEREX; + sd.BufferEx.NumElements = n4; + sd.BufferEx.Flags = D3D11_BUFFEREX_SRV_FLAG_RAW; + if (FAILED(dev->CreateShaderResourceView(buf, &sd, srv))) { buf->Release(); return nullptr; } + } + if (uav && uav_out) { + D3D11_UNORDERED_ACCESS_VIEW_DESC ud{}; + ud.Format = DXGI_FORMAT_R32_TYPELESS; + ud.ViewDimension = D3D11_UAV_DIMENSION_BUFFER; + ud.Buffer.NumElements = n4; + ud.Buffer.Flags = D3D11_BUFFER_UAV_FLAG_RAW; + if (FAILED(dev->CreateUnorderedAccessView(buf, &ud, uav_out))) { buf->Release(); return nullptr; } + } + return buf; +} + +extern "C" int lrt_d3d11_trace_scene(lrt_d3d11_engine* e, const lrt_tri_scene* s, + const lrt_ray* rays, uint32_t n, + lrt_hit* out, lrt_result* err) { + if (!e || !s || (n && (!rays || !out))) { + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return -1; + } + if (n == 0) { if (err) *err = LRT_RESULT_OK; return 0; } + + void* blob = nullptr; + size_t blob_n = 0; + lrt_result sr = lrt_tri_scene_save_to_memory(s, &blob, &blob_n); + if (sr != LRT_RESULT_OK) { + set_err(e, "scene not GPU-traceable (quantized/curve/user?)"); + if (err) *err = sr; + return -1; + } + const lrts_header* h = (const lrts_header*)blob; + uint32_t w = h->layout; // 4 or 8 + if (w != 4 && w != 8) { + set_err(e, "unsupported BVH layout for D3D11 trace (need BVH4/BVH8)"); + free(blob); + if (err) *err = LRT_RESULT_INVALID_ARGUMENT; + return -1; + } + + lrt_tri_stats st; + lrt_tri_scene_stats(s, &st); + uint32_t stack = trace_stack_for(st.max_depth, w); + if (stack == 0) { + set_err(e, "BVH too deep for the GPU compute stack"); + free(blob); + if (err) *err = LRT_RESULT_TRAVERSAL_OVERFLOW; + return -1; + } + + ID3D11ComputeShader* cs = get_shader(e, w, stack); + if (!cs) { free(blob); if (err) *err = LRT_RESULT_OUT_OF_MEMORY; return -1; } + + const uint32_t nodes_bytes = h->node_count * h->node_stride; + const uint32_t blocks_bytes = h->block_count * h->block_stride; + + ID3D11Buffer *b_nodes = nullptr, *b_blocks = nullptr, *b_rays = nullptr, + *b_hits = nullptr, *b_cb = nullptr, *b_stage = nullptr; + ID3D11ShaderResourceView *srv_nodes = nullptr, *srv_blocks = nullptr, + *srv_rays = nullptr; + ID3D11UnorderedAccessView* uav_hits = nullptr; + int result = -1; + + b_nodes = make_raw_buffer(e->dev, (const char*)blob + h->node_offset, + nodes_bytes, false, &srv_nodes, nullptr); + b_blocks = make_raw_buffer(e->dev, (const char*)blob + h->block_offset, + blocks_bytes, false, &srv_blocks, nullptr); + b_rays = make_raw_buffer(e->dev, rays, n * (uint32_t)sizeof(lrt_ray), false, + &srv_rays, nullptr); + b_hits = make_raw_buffer(e->dev, nullptr, n * (uint32_t)sizeof(lrt_hit), true, + nullptr, &uav_hits); + if (!b_nodes || !b_blocks || !b_rays || !b_hits) { + set_err(e, "buffer/view creation failed"); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + goto cleanup; + } + + { + TracePC pc{h->root, h->node_count, h->block_count, n}; + D3D11_BUFFER_DESC cbd{}; + cbd.ByteWidth = sizeof(TracePC); // 16 bytes + cbd.Usage = D3D11_USAGE_DEFAULT; + cbd.BindFlags = D3D11_BIND_CONSTANT_BUFFER; + D3D11_SUBRESOURCE_DATA cinit{}; + cinit.pSysMem = &pc; + if (FAILED(e->dev->CreateBuffer(&cbd, &cinit, &b_cb))) { + set_err(e, "constant buffer creation failed"); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + goto cleanup; + } + } + + { + ID3D11DeviceContext* c = e->ctx; + c->CSSetShader(cs, nullptr, 0); + ID3D11ShaderResourceView* srvs[3] = {srv_nodes, srv_blocks, srv_rays}; + c->CSSetShaderResources(0, 3, srvs); // t0,t1,t2 + UINT init_counts = 0; + c->CSSetUnorderedAccessViews(3, 1, &uav_hits, &init_counts); // u3 + c->CSSetConstantBuffers(0, 1, &b_cb); // b0 + UINT groups = (n + 63u) / 64u; + c->Dispatch(groups, 1, 1); + // Unbind UAV so we can copy it out. + ID3D11UnorderedAccessView* none_uav = nullptr; + c->CSSetUnorderedAccessViews(3, 1, &none_uav, &init_counts); + } + + { // Read hits back via a staging buffer. + D3D11_BUFFER_DESC sd{}; + sd.ByteWidth = n * (uint32_t)sizeof(lrt_hit); + sd.Usage = D3D11_USAGE_STAGING; + sd.CPUAccessFlags = D3D11_CPU_ACCESS_READ; + if (FAILED(e->dev->CreateBuffer(&sd, nullptr, &b_stage))) { + set_err(e, "staging buffer creation failed"); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + goto cleanup; + } + e->ctx->CopyResource(b_stage, b_hits); + D3D11_MAPPED_SUBRESOURCE mapped{}; + if (FAILED(e->ctx->Map(b_stage, 0, D3D11_MAP_READ, 0, &mapped))) { + set_err(e, "map(staging) failed"); + if (err) *err = LRT_RESULT_OUT_OF_MEMORY; + goto cleanup; + } + std::memcpy(out, mapped.pData, n * sizeof(lrt_hit)); + e->ctx->Unmap(b_stage, 0); + } + + { + int hits = 0; + for (uint32_t i = 0; i < n; i++) + if (out[i].prim_id != LRT_TRI_NO_HIT) hits++; + if (err) *err = LRT_RESULT_OK; + result = hits; + } + +cleanup: + safe_release(srv_nodes); + safe_release(srv_blocks); + safe_release(srv_rays); + safe_release(uav_hits); + safe_release(b_nodes); + safe_release(b_blocks); + safe_release(b_rays); + safe_release(b_hits); + safe_release(b_cb); + safe_release(b_stage); + free(blob); + return result; +} + +#endif // _WIN32 diff --git a/lightrt_c_d3d11.h b/lightrt_c_d3d11.h new file mode 100644 index 0000000..bd1b544 --- /dev/null +++ b/lightrt_c_d3d11.h @@ -0,0 +1,57 @@ +/* + * lightrt_c_d3d11.h — Direct3D 11 GPU interop for the LightRT triangle kernel. + * + * Mirror of lightrt_c_vk.h "Path A" (CPU build -> GPU trace), but on a D3D11 + * compute shader instead of Vulkan. The trace shader (d3d/shaders/trace_bvh.hlsl) + * is decompiled from the Vulkan trace_bvh SPIR-V with SPIRV-Cross, so it walks + * the same serialized BVH4/BVH8 scene and returns identical hits. + * + * Motivation: on AMD GCN/Polaris the amdvlk Vulkan compute path mis-renders / + * hangs, while the (much more mature) D3D11 driver is solid. D3D11 is also + * ubiquitous on Windows and needs no SDK to build (d3d11 + d3dcompiler ship with + * the OS). Windows-only; on other platforms this header compiles to nothing. + * + * Unlike the Vulkan helper this batches a whole ray array into ONE dispatch, so + * a full-frame trace is a single GPU round-trip rather than one per pixel. + * + * SPDX-License-Identifier: Apache-2.0 + */ +#ifndef LIGHTRT_C_D3D11_H +#define LIGHTRT_C_D3D11_H + +#include + +#include "lightrt_c.h" /* lrt_result */ +#include "lightrt_c_tri.h" /* lrt_tri_scene, lrt_ray, lrt_hit */ + +#ifdef __cplusplus +extern "C" { +#endif + +typedef struct lrt_d3d11_engine lrt_d3d11_engine; + +/* Create a D3D11 compute engine. prefer_discrete picks a discrete adapter when + * available. Returns NULL on any failure (no D3D11, no compute device); the + * caller should fall back to the CPU kernel. err is set when non-NULL. */ +lrt_d3d11_engine *lrt_d3d11_engine_create(int prefer_discrete, lrt_result *err); +void lrt_d3d11_engine_destroy(lrt_d3d11_engine *e); + +/* Selected adapter name (e.g. "AMD Radeon RX 570"). */ +const char *lrt_d3d11_engine_device_name(const lrt_d3d11_engine *e); + +/* Human-readable message for the last failed call. */ +const char *lrt_d3d11_engine_last_error(const lrt_d3d11_engine *e); + +/* Trace n rays against a CPU-built scene on the GPU. Writes n hits to out + * (prim_id == LRT_TRI_NO_HIT for misses). Returns the number of rays that hit, + * or -1 on error (err set). Only plain BVH4/BVH8 triangle scenes are supported + * (same restriction as lrt_vk_trace_scene). */ +int lrt_d3d11_trace_scene(lrt_d3d11_engine *e, const lrt_tri_scene *s, + const lrt_ray *rays, uint32_t n, lrt_hit *out, + lrt_result *err); + +#ifdef __cplusplus +} +#endif + +#endif /* LIGHTRT_C_D3D11_H */