From d47aeb0a7702b80973ea511d84afa8a711b61fc5 Mon Sep 17 00:00:00 2001 From: Alexis Duburcq Date: Thu, 23 Jul 2026 17:36:06 +0200 Subject: [PATCH 1/2] Remove orphan ECS system unit, stale header, and dead viewer/internal members - Delete the unused ECS 'System' unit: src/core/system.cpp (listed in no build target) plus include/madrona/system.{hpp,inl} and the device system.{hpp,inl} headers (System/SystemBase referenced nowhere; system.hpp was included only by system.cpp). - Delete src/render/vk/engine_interop.hpp -- a stale duplicate of the EngineInterop struct that actually lives in render_common.hpp (zero includers). - Drop dead BatchRenderer methods left over from the removed viewer: importCudaData (decl only), getComponentBuffer, getLatestWaitSemaphore (all zero callers). - Drop the write-only RenderContext::gpu_input_ member. --- include/madrona/system.hpp | 65 ------- include/madrona/system.inl | 50 ----- src/core/system.cpp | 10 - src/mw/device/include/madrona/system.hpp | 38 ---- src/mw/device/include/madrona/system.inl | 43 ----- src/render/batch_renderer.cpp | 21 --- src/render/batch_renderer.hpp | 6 - src/render/render_ctx.cpp | 3 +- src/render/render_ctx.hpp | 2 - src/render/vk/engine_interop.hpp | 221 ----------------------- 10 files changed, 1 insertion(+), 458 deletions(-) delete mode 100644 include/madrona/system.hpp delete mode 100644 include/madrona/system.inl delete mode 100644 src/core/system.cpp delete mode 100644 src/mw/device/include/madrona/system.hpp delete mode 100644 src/mw/device/include/madrona/system.inl delete mode 100644 src/render/vk/engine_interop.hpp diff --git a/include/madrona/system.hpp b/include/madrona/system.hpp deleted file mode 100644 index b44ec54f..00000000 --- a/include/madrona/system.hpp +++ /dev/null @@ -1,65 +0,0 @@ -#pragma once - -#include -#include -#include -#include - -#include - -namespace madrona { - -class SystemBase { -public: - using EntryFn = void (*)(SystemBase *, void *, uint32_t); - - SystemBase(EntryFn entry_fn); - AtomicU32 numInvocations; -private: - EntryFn entry_fn_; -friend class TaskGraph; -}; - -template -class CustomSystem : public SystemBase { -public: - CustomSystem(); - -private: - static void entry(SystemBase *sys, void *data, - uint32_t invocation_offset); -}; - -template -class ParallelForSystem : public SystemBase { -public: - ParallelForSystem(Context &ctx); - -private: - static void entry(SystemBase *sys, void *data, uint32_t invocation_offset); - - Query query_; -}; - -#if 0 -template -class LambdaParallelForSystem : public ParallelForSystem< - LambdaParallelForSystem, ComponentTs...> { - using ContextT = utils::FirstArgTypeExtractor; -public: - static LambdaParallelForSystem * allocate(Context &ctx); - static void deallocate(Context &ctx, - LambdaParallelForSystem *lambda); - - void run(ContextT &ctx, uint32_t invocation_idx); - -private: - LambdaParallelForSystem(Fn &&fn); - - Fn fn; -}; -#endif - -} - -#include "system.inl" diff --git a/include/madrona/system.inl b/include/madrona/system.inl deleted file mode 100644 index 686d402b..00000000 --- a/include/madrona/system.inl +++ /dev/null @@ -1,50 +0,0 @@ -#pragma once - -namespace madrona { - -#if 0 -template -CustomSystem::CustomSystem() - : SystemBase(&CustomSystem::entry) -{} - -template -void CustomSystem::entry(SystemBase *sys_base, - void *data, - uint32_t invocation_offset) -{ - SystemT *sys = static_cast(sys_base); - sys->run(data, invocation_offset); -} - -template -ParallelForSystem::ParallelForSystem(Context &ctx) - : SystemBase((SystemBase::EntryFn) - &ParallelForSystem::entry), - query_(ctx.query()) -{} - -template -void ParallelForSystem::entry(SystemBase *sys_base, - void *data, uint32_t invocation_offset) -{ - (void)sys_base; - (void)data; - (void)invocation_offset; -} - -template -LambdaParallelFor::LambdaParallelForSystem(Fn &&fn) - : ParallelForSystem, ComponentTs...>(), - fn(std::forward(fn)) -{} - -template -void LambdaParallelForSystem::run(ContextT &ctx, - ComponentTs... &&components) -{ - fn(ctx, std::forward(components)...); -} -#endif - -} diff --git a/src/core/system.cpp b/src/core/system.cpp deleted file mode 100644 index 6aef9f40..00000000 --- a/src/core/system.cpp +++ /dev/null @@ -1,10 +0,0 @@ -#include - -namespace madrona { - -SystemBase::SystemBase(EntryFn fn_ptr) - : numInvocations(0), - entry_fn_(fn_ptr) -{} - -} diff --git a/src/mw/device/include/madrona/system.hpp b/src/mw/device/include/madrona/system.hpp deleted file mode 100644 index 5da72f8f..00000000 --- a/src/mw/device/include/madrona/system.hpp +++ /dev/null @@ -1,38 +0,0 @@ -#pragma once - -#include -#include -#include - -#include - -namespace madrona { - -struct SharedSystemState { - SystemBase **systems; - AtomicU32 numInvocations; - uint32_t sysID; -}; - -class SystemBase { -public: - inline SystemBase(uint32_t sys_id); - -private: - uint32_t sys_id_; - SharedSystemState * shared_; -friend class TaskGraph; -}; - -template -class CustomSystem : public SystemBase { -public: - CustomSystem(); - - static void entry(SystemBase *sys, void *data, - uint32_t invocation_offset); -}; - -} - -#include "system.inl" diff --git a/src/mw/device/include/madrona/system.inl b/src/mw/device/include/madrona/system.inl deleted file mode 100644 index 70157381..00000000 --- a/src/mw/device/include/madrona/system.inl +++ /dev/null @@ -1,43 +0,0 @@ -#pragma once - -namespace madrona { -namespace mwGPU { - -template -__attribute__((used, always_inline)) -inline void systemEntry(SystemBase *sys_base, void *user_data, - uint32_t invocation_offset) -{ - SystemT::entry(sys_base, user_data, invocation_offset); -} - -template -struct SystemIDBase { - static uint32_t id; -}; - -template ) = - systemEntry> -struct SystemID : SystemIDBase {}; - -} - -SystemBase::SystemBase(uint32_t sys_id) - : sys_id_(sys_id) -{} - -template -CustomSystem::CustomSystem() - : SystemBase(mwGPU::SystemID::id) -{} - -template -void CustomSystem::entry(SystemBase *sys_base, void *data, - uint32_t invocation_offset) -{ - SystemT *sys = static_cast(sys_base); - sys->run(data, invocation_offset); -} - -} diff --git a/src/render/batch_renderer.cpp b/src/render/batch_renderer.cpp index 41d4a503..adf2ae76 100644 --- a/src/render/batch_renderer.cpp +++ b/src/render/batch_renderer.cpp @@ -2492,27 +2492,6 @@ BatchImportedBuffers &BatchRenderer::getImportedBuffers(uint32_t frame_id) return impl->batchFrames[frame_id].buffers; } -const vk::LocalBuffer &BatchRenderer::getComponentBuffer(uint32_t frame_id, uint32_t component) const -{ - return impl->batchFrames[frame_id].getComponentOutputBuffer(component); -} - -// Get the semaphore that the viewer renderer has to wait on -VkSemaphore BatchRenderer::getLatestWaitSemaphore() -{ - uint32_t last_frame = (impl->currentFrame + impl->batchFrames.size() - 1) % impl->batchFrames.size(); - assert(impl->batchFrames[last_frame].latestOp != LatestOperation::None); - if (impl->batchFrames[last_frame].latestOp == LatestOperation::RenderPrepare) { - return impl->batchFrames[last_frame].prepareFinished; - } else if (impl->batchFrames[last_frame].latestOp == LatestOperation::RenderViews) { - return impl->batchFrames[last_frame].renderFinished; - } else if (impl->batchFrames[last_frame].latestOp == LatestOperation::Transition) { - return impl->batchFrames[last_frame].layoutTransitionFinished; - } - - return VK_NULL_HANDLE; -} - const void *BatchRenderer::getComponentCUDAPtr(uint32_t frame_id, uint32_t component) const { #ifndef MADRONA_VK_CUDA_SUPPORT diff --git a/src/render/batch_renderer.hpp b/src/render/batch_renderer.hpp index 0b990ece..98b8dfb1 100644 --- a/src/render/batch_renderer.hpp +++ b/src/render/batch_renderer.hpp @@ -183,7 +183,6 @@ struct BatchRenderer { BatchRenderer(const Config& cfg, RenderContext &rctx); ~BatchRenderer(); - void importCudaData(VkCommandBuffer); void setRenderOptions(const render::RenderOptions &render_options); void prepareForRendering(BatchRenderInfo info, EngineInterop *interop); @@ -194,11 +193,6 @@ struct BatchRenderer { BatchImportedBuffers &getImportedBuffers(uint32_t frame_id); - const vk::LocalBuffer &getComponentBuffer(uint32_t frame_id, uint32_t component) const; - // Get the semaphore that the viewer renderer has to wait on. - // This is either going to be the semaphore from prepareForRendering, - // or it's the one from renderViews. - VkSemaphore getLatestWaitSemaphore(); const void *getComponentCUDAPtr(uint32_t frame_id, uint32_t component) const; }; diff --git a/src/render/render_ctx.cpp b/src/render/render_ctx.cpp index ed0d74a3..686df267 100644 --- a/src/render/render_ctx.cpp +++ b/src/render/render_ctx.cpp @@ -1088,8 +1088,7 @@ RenderContext::RenderContext( loaded_assets_(0), sky_(loadSky(dev, alloc, renderQueue)), material_textures_(0), - num_worlds_(cfg.numWorlds), - gpu_input_(cfg.execMode == ExecMode::CUDA) + num_worlds_(cfg.numWorlds) { { VkDescriptorPoolSize pool_sizes[] = { diff --git a/src/render/render_ctx.hpp b/src/render/render_ctx.hpp index 582971f6..118c4a82 100644 --- a/src/render/render_ctx.hpp +++ b/src/render/render_ctx.hpp @@ -83,8 +83,6 @@ struct RenderContext { VkDescriptorSetLayout sky_data_layout_; VkDescriptorSet sky_data_set_; - bool gpu_input_; - VkDescriptorSetLayout aabb_set_layout_; }; diff --git a/src/render/vk/engine_interop.hpp b/src/render/vk/engine_interop.hpp deleted file mode 100644 index f5ca6757..00000000 --- a/src/render/vk/engine_interop.hpp +++ /dev/null @@ -1,221 +0,0 @@ -#pragma once - -#include - -#include "memory.hpp" -#include "cuda_interop.hpp" -#include "utils.hpp" - -namespace madrona::render::vk { - -struct CpuMode {}; -struct CudaMode {}; - -template -struct EngineModeVariant { - using BaseT = EngineModeVariant; - - union { - CudaT cuda; - CpuT cpu; - }; - - bool isCuda; - - template - inline EngineModeVariant(CudaMode, Args && ...args) - : cuda(std::forward(args) ...), - isCuda(true) - {} - - template - inline EngineModeVariant(CpuMode, Args && ...args) - : cpu(std::forward(args) ...), - isCuda(false) - {} - - inline ~EngineModeVariant() - { - if (isCuda) { - cuda.~CudaT(); - } else { - cpu.~CpuT(); - } - } - - inline EngineModeVariant(EngineModeVariant &&o) - { - isCuda = o.isCuda; - - if (isCuda) { - new (&cuda) CudaT(std::move(o.cuda)); - } else { - new (&cpu) CpuT(std::move(o.cpu)); - } - } -}; - -// FIXME: most of the uses of this don't need staging to be kept around -struct HostToEngineBufferCUDA { - HostBuffer staging; - DedicatedBuffer devBuffer; - CudaImportedBuffer cudaImported; - - inline HostToEngineBufferCUDA(const Device &dev, - MemoryAllocator &mem, - uint64_t num_bytes, - int cuda_gpu_id) - : staging(mem.makeStagingBuffer(num_bytes)), - devBuffer(mem.makeDedicatedBuffer(num_bytes, false, true)), - cudaImported(dev, cuda_gpu_id, devBuffer.mem, num_bytes) - {} -}; - -struct HostToEngineBufferCPU { - void *ptr; - uint64_t numBytes; - - inline HostToEngineBufferCPU(uint64_t num_bytes) - : ptr(malloc(num_bytes)), - numBytes(num_bytes) - {} - - inline ~HostToEngineBufferCPU() - { - free(ptr); - } - - inline HostToEngineBufferCPU(HostToEngineBufferCPU &&o) - : ptr(o.ptr), - numBytes(o.numBytes) - { - o.ptr = nullptr; - o.numBytes = 0; - } -}; - -struct HostToEngineBuffer : public EngineModeVariant< - HostToEngineBufferCUDA, HostToEngineBufferCPU> { - using BaseT::BaseT; - - inline void * enginePointer() const - { - if (isCuda) { - return cuda.cudaImported.getDevicePointer(); - } else { - return cpu.ptr; - } - } - - inline void * hostPointer() const - { - if (isCuda) { - return cuda.staging.ptr; - } else { - return cpu.ptr; - } - } - - inline bool needsEngineCopy() const - { - return isCuda; - } - - inline void toEngine(const Device &dev, VkCommandBuffer cmd, - uint32_t offset, uint32_t num_bytes) - { - if (isCuda) { - cuda.staging.flush(dev); - - VkBufferCopy buffer_copy; - buffer_copy.srcOffset = offset; - buffer_copy.dstOffset = offset; - buffer_copy.size = num_bytes; - - dev.dt.cmdCopyBuffer(cmd, cuda.staging.buffer, - cuda.devBuffer.buf.buffer, 1, - &buffer_copy); - } - } -}; - -struct EngineToRendererBufferCUDA { - DedicatedBuffer devBuffer; - CudaImportedBuffer cudaImported; - - inline EngineToRendererBufferCUDA(const Device &dev, - MemoryAllocator &mem, uint64_t num_bytes, int cuda_gpu_id) - : devBuffer(mem.makeDedicatedBuffer(num_bytes, true, true)), - cudaImported(dev, cuda_gpu_id, devBuffer.mem, num_bytes) - {} -}; - -struct EngineToRendererBufferCPU { - HostBuffer staging; - DedicatedBuffer devBuffer; - uint64_t numBytes; - - inline EngineToRendererBufferCPU(MemoryAllocator &mem, uint64_t num_bytes) - : staging(mem.makeStagingBuffer(num_bytes)), - devBuffer(mem.makeDedicatedBuffer(num_bytes, true)), - numBytes(num_bytes) - {} -}; - -struct EngineToRendererBuffer : public EngineModeVariant< - EngineToRendererBufferCUDA, EngineToRendererBufferCPU> { - using BaseT::BaseT; - - inline void * enginePointer() const - { - if (isCuda) { - return cuda.cudaImported.getDevicePointer(); - } else { - return cpu.staging.ptr; - } - } - - inline const LocalBuffer & rendererBuffer() const - { - return isCuda ? cuda.devBuffer.buf : cpu.devBuffer.buf; - } - - // FIXME: this API isn't great, it stops batching pipeline barriers - inline void toRenderer(const Device &dev, VkCommandBuffer cmd, - VkAccessFlagBits pipeline_access, - VkPipelineStageFlagBits pipeline_stage) - { - if (!isCuda) { - cpu.staging.flush(dev); - - VkBufferCopy buffer_copy { - .srcOffset = 0, - .dstOffset = 0, - .size = cpu.numBytes, - }; - - dev.dt.cmdCopyBuffer(cmd, cpu.staging.buffer, - cpu.devBuffer.buf.buffer, 1, &buffer_copy); - - if (pipeline_access != VK_ACCESS_NONE) { - VkBufferMemoryBarrier barrier; - barrier.sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER; - barrier.pNext = nullptr; - barrier.srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT; - barrier.dstAccessMask = pipeline_access; - barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; - barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; - barrier.buffer = cpu.devBuffer.buf.buffer; - barrier.offset = 0; - barrier.size = VK_WHOLE_SIZE; - - dev.dt.cmdPipelineBarrier(cmd, - VK_PIPELINE_STAGE_TRANSFER_BIT, - pipeline_stage, 0, 0, - nullptr, 1, &barrier, 0, nullptr); - } - } - } -}; - -} From b1736396dbb121ba7118574a1eac6247136086e6 Mon Sep 17 00:00:00 2001 From: Alexis Duburcq Date: Thu, 23 Jul 2026 17:43:29 +0200 Subject: [PATCH 2/2] Remove all dead '#if 0' preprocessor blocks Delete the 31 always-false '#if 0' (and '#if 0 && ...') blocks left across the host and nvrtc-device sources -- disabled/reference code, debug prints, and alternate implementations. None had a live '#else' branch (verified), so each is removed whole; pure deletions, no behavior change. Touches render, vk, mw device (bvh/bvh_raycast/sort_archetype), common, bridge, and several ECS headers. --- include/madrona/math.inl | 24 --- include/madrona/mesh_bvh.inl | 153 ------------------ include/madrona/rand.inl | 10 -- include/madrona/state.inl | 6 - src/bridge/sim.cpp | 4 - src/common/table.cpp | 7 - src/mw/cuda_exec.cpp | 7 - src/mw/device/bvh.cpp | 252 ----------------------------- src/mw/device/bvh_raycast.cpp | 46 ------ src/mw/device/sort_archetype.cpp | 24 --- src/render/asset_processor.cpp | 4 - src/render/batch_renderer.cpp | 14 -- src/render/ecs_system.cpp | 57 ------- src/render/render_ctx.cpp | 8 - src/render/shaders/shader_common.h | 14 -- src/render/vk/backend.cpp | 25 --- src/render/vk/memory.cpp | 7 - 17 files changed, 662 deletions(-) diff --git a/include/madrona/math.inl b/include/madrona/math.inl index fef0c8b8..b5f9aa36 100644 --- a/include/madrona/math.inl +++ b/include/madrona/math.inl @@ -756,30 +756,6 @@ Quat Quat::fromAngularVec(Vector3 v) Quat Quat::fromBasis(Vector3 a, Vector3 b, Vector3 c) { //Modified from glm::quat_cast -#if 0 -=============================================================================== -The MIT License -------------------------------------------------------------------------------- -Copyright (c) 2005 - G-Truc Creation - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -#endif float four_x_squared_minus1 = a.x - b.y - c.z; float four_y_squared_minus1 = b.y - a.x - c.z; diff --git a/include/madrona/mesh_bvh.inl b/include/madrona/mesh_bvh.inl index 5ed488b4..4fa7e762 100644 --- a/include/madrona/mesh_bvh.inl +++ b/include/madrona/mesh_bvh.inl @@ -166,153 +166,6 @@ bool MeshBVH::traceRay(math::Vector3 ray_o, return ray_hit; } -#if 0 -bool MeshBVH::traceRay(math::Vector3 ray_o, - math::Vector3 ray_d, - float *out_hit_t, - math::Vector3 *out_hit_normal, - void* shared, - TraversalStack *stack, - float t_max) const -{ - using namespace math; - constexpr float diveps = 0.0000001f; - - Diag3x3 inv_d = Diag3x3::fromVec(ray_d).inv(); - - RayIsectTxfm tri_isect_txfm = computeRayIsectTxfm(ray_o, ray_d, inv_d); - - uint32_t previous_stack_size = stack->size; - - stack->push(0); - -#ifdef SHARED_STACK - const int32_t mwgpu_warp_id = threadIdx.x / 32; - const int32_t mwgpu_warp_lane = threadIdx.x % 32; - const int32_t num_smem_bytes_per_warp = - (mwGPU::SharedMemStorage::numBytesPerWarp()/4)*4; - - auto sharedMem = ((char*)shared) + mwgpu_warp_id * num_smem_bytes_per_warp + - SHARED_STACK_SIZE * sizeof(int32_t) * mwgpu_warp_lane; - int32_t* shared_stack = (int32_t*)sharedMem; - shared_stack[0] = 0; -#endif - - bool ray_hit = false; - Vector3 closest_hit_normal = Vector3{0,0,0}; - - while (stack->size > previous_stack_size) { - int32_t node_idx = stack->pop(); - const QBVHNode &node = nodes[node_idx]; - - float rayXInv = copysignf(ray_d.x == 0 ? 1/diveps : 1/ray_d.x,ray_d.x); - float rayYInv = copysignf(ray_d.y == 0 ? 1/diveps : 1/ray_d.y,ray_d.y); - float rayZInv = copysignf(ray_d.z == 0 ? 1/diveps : 1/ray_d.z,ray_d.z); - //NVIDIA's method, transform for ray plane to quantized space. Shift to IEEE exponent bits. - -#ifdef MADRONA_GPU_MODE - float dirQuantX = __uint_as_float((node.expX + 127) << 23) * rayXInv; - float dirQuantY = __uint_as_float((node.expY + 127) << 23) * rayYInv; - float dirQuantZ = __uint_as_float((node.expZ + 127) << 23) * rayZInv; -#else - float dirQuantX = std::bit_cast((node.expX + 127) << 23) * rayXInv; - float dirQuantY = std::bit_cast((node.expY + 127) << 23) * rayYInv; - float dirQuantZ = std::bit_cast((node.expZ + 127) << 23) * rayZInv; -#endif - - float originQuantX = (node.minX - ray_o.x) * rayXInv; - float originQuantY = (node.minY - ray_o.y) * rayYInv; - float originQuantZ = (node.minZ - ray_o.z) * rayZInv; - - for (CountT i = 0; i < MeshBVH::nodeWidth; i++) { - if (!node.hasChild(i)) { - continue; // Technically this could be break? - }; - - - float t_near_x = node.qMinX[i] * dirQuantX + originQuantX; - float t_near_y = node.qMinY[i] * dirQuantY + originQuantY; - float t_near_z = node.qMinZ[i] * dirQuantZ + originQuantZ; - - float t_far_x = node.qMaxX[i] * dirQuantX + originQuantX; - float t_far_y = node.qMaxY[i] * dirQuantY + originQuantY; - float t_far_z = node.qMaxZ[i] * dirQuantZ + originQuantZ; -/* - madrona::math::AABB child_aabb { - .pMin = { - node.minX + std::bit_cast((node.expX + 127) << 23) * node.qMinX[i], - node.minY + std::bit_cast((node.expY + 127) << 23) * node.qMinY[i], - node.minZ + std::bit_cast((node.expZ + 127) << 23) * node.qMinZ[i], - }, - .pMax = { - node.minX + std::bit_cast((node.expX + 127) << 23) * node.qMaxX[i], - node.minY + std::bit_cast((node.expY + 127) << 23) * node.qMaxY[i], - node.minZ + std::bit_cast((node.expZ + 127) << 23) * node.qMaxZ[i] - }, - }; - - float t_near_x = (child_aabb[tri_isect_txfm.nearX] - - tri_isect_txfm.oNear.x) * - tri_isect_txfm.invDirNear.x; - float t_near_y = (child_aabb[tri_isect_txfm.nearY] - - tri_isect_txfm.oNear.y) * - tri_isect_txfm.invDirNear.y; - float t_near_z = (child_aabb[tri_isect_txfm.nearZ] - - tri_isect_txfm.oNear.z) * - tri_isect_txfm.invDirNear.z; - - float t_far_x = (child_aabb[tri_isect_txfm.farX] - - tri_isect_txfm.oFar.x) * - tri_isect_txfm.invDirFar.x; - float t_far_y = (child_aabb[tri_isect_txfm.farY] - - tri_isect_txfm.oFar.y) * - tri_isect_txfm.invDirFar.y; - float t_far_z = (child_aabb[tri_isect_txfm.farZ] - - tri_isect_txfm.oFar.z) * - tri_isect_txfm.invDirFar.z; - float t_near = fmaxf(t_near_x, fmaxf(t_near_y, - fmaxf(t_near_z, 0.f))); - float t_far = fminf(t_far_x, fminf(t_far_y, - fminf(t_far_z, t_max))); - - */ - - float t_near = fmaxf(fminf(t_near_x,t_far_x), fmaxf(fminf(t_near_y,t_far_y), - fmaxf(fminf(t_near_z,t_far_z), 0.f))); - float t_far = fminf(fmaxf(t_far_x,t_near_x), fminf(fmaxf(t_far_y,t_near_y), - fminf(fmaxf(t_far_z,t_near_z), t_max))); - - if (t_near <= t_far) { - if (node.isLeaf(i)) { - int32_t leaf_idx = node.leafIDX(i); - - float hit_t; - Vector3 leaf_hit_normal; - bool leaf_hit = traceRayLeaf(leaf_idx, node.triSize[i], tri_isect_txfm, - ray_o, t_max, &hit_t, &leaf_hit_normal); - - if (leaf_hit) { - ray_hit = true; - t_max = hit_t; - closest_hit_normal = leaf_hit_normal; - } - } else { - // assert(stack->size < 32); - stack->push(node.childrenIdx[i]); - } - } - } - } - - if (!ray_hit) { - return false; - } - - *out_hit_t = t_max; - *out_hit_normal = closest_hit_normal; - return ray_hit; -} -#endif bool MeshBVH::traceRayLeaf(int32_t leaf_idx, int32_t num_tris, @@ -482,12 +335,6 @@ bool MeshBVH::rayTriangleIntersection( // normalize U, V, W, and T const float rcpDet = 1.0f / det; -#if 0 - hit.u = U * rcpDet; - hit.v = V * rcpDet; - hit.w = W * rcpDet; - hit.t = T * rcpDet; -#endif *out_hit_t = T * rcpDet; *bary_out = Vector3{U,V,W} * rcpDet; diff --git a/include/madrona/rand.inl b/include/madrona/rand.inl index 123ffccb..5aa90a8f 100644 --- a/include/madrona/rand.inl +++ b/include/madrona/rand.inl @@ -209,16 +209,6 @@ constexpr float bitsToFloat01(uint32_t rand_bits) // implementations that generate in the range (0, 1] return (rand_bits >> 8_u32) * 0x1p-24f; -#if 0 - constexpr uint32_t exponent = 0x3f800000; - uint32_t raw = (exponent | (rand_bits >> 9)) - 1; - -#ifdef MADRONA_GPU_MODE - return __uint_as_float(raw); -#else - return std::bit_cast(raw); -#endif -#endif } } diff --git a/include/madrona/state.inl b/include/madrona/state.inl index 73a26d26..2313101c 100644 --- a/include/madrona/state.inl +++ b/include/madrona/state.inl @@ -282,12 +282,6 @@ ComponentID StateManager::componentID() const template ArchetypeID StateManager::archetypeID() const { -#if 0 - uint32_t id = TypeTracker::typeID(); - - assert(id != TypeTracker::unassignedTypeID && - "Trying to access an unregistered archetype!"); -#endif return ArchetypeID { TypeTracker::typeID(), }; diff --git a/src/bridge/sim.cpp b/src/bridge/sim.cpp index cb9e7608..d19b5195 100644 --- a/src/bridge/sim.cpp +++ b/src/bridge/sim.cpp @@ -79,10 +79,6 @@ static void setupRenderTasks(TaskGraphBuilder &builder, Span deps, bool update_mats = false) { -#if 0 - builder.addToGraph>(deps); -#endif RenderingSystem::setupTasks(builder, deps, update_mats); } diff --git a/src/common/table.cpp b/src/common/table.cpp index 6687e8e0..5672a7ab 100644 --- a/src/common/table.cpp +++ b/src/common/table.cpp @@ -28,13 +28,6 @@ Table::Table(const TypeInfo *component_types, CountT num_components, for (int i = 0; i < (int)num_components; i++) { const TypeInfo &type = component_types[i]; -#if 0 - // 3rd argument is offsetting the start from the page aligned boundary - // to avoid everything mapping to the same cache sets. Should revisit - - // maybe add a random offset for each Table as well? - columns_.emplace_back(type.numBytes, type.alignment, - MADRONA_CACHE_LINE * (i + 1), ICfg::maxRowsPerTable); -#endif size_t column_bytes_per_row = (size_t)type.numBytes; columns_[i] = malloc( diff --git a/src/mw/cuda_exec.cpp b/src/mw/cuda_exec.cpp index 256f6fdc..975e8b3b 100644 --- a/src/mw/cuda_exec.cpp +++ b/src/mw/cuda_exec.cpp @@ -646,13 +646,6 @@ static GPUCompileResults compileCode( } }; -#if 0 - // Don't need the device runtime without dynamic parallelism - - checkLinker(cuLinkAddFile(linker, CU_JIT_INPUT_LIBRARY, - MADRONA_CUDADEVRT_PATH, - 0, nullptr, nullptr)); -#endif nvJitLinkInputType linker_input_type; if (opt_mode == CompileConfig::OptMode::LTO) { diff --git a/src/mw/device/bvh.cpp b/src/mw/device/bvh.cpp index ba356810..99ab74a4 100644 --- a/src/mw/device/bvh.cpp +++ b/src/mw/device/bvh.cpp @@ -1206,16 +1206,6 @@ extern "C" __global__ void bvhWidenTree() QBVHNode *current_qbvh_node = &smem->traversalNodes[stored_job.qbvhNodeIndex - 1]; -#if 0 - if (stored_job.qbvhNodeIndex - 1 == 1) { - LOG("CONSTRUCTION LBVH NODE WITH INDEX 1: num_children = {}\n", - num_children); - - for (int i = 0; i < num_children; ++i) { - LOG("child {} is {}\n", i, children_indices[i]); - } - } -#endif *current_qbvh_node = QBVHNode::construct( num_children, @@ -1232,252 +1222,10 @@ extern "C" __global__ void bvhWidenTree() } } -#if 0 -// Phase 1 of the optimization kernel: -// Find the first treelet roots which expand to treelets with at least 7 -// leaf nodes. All these potential treelets will be pushed to a global -// buffer which will then be pulled from in the next stage by each warp -// for processing. -// Takes in the node at which this thread will start searching upwards the tree -static __device__ inline void pushPotentialRoots(uint32_t start_search_idx, - uint32_t total_num_instances, - uint32_t num_resident_threads) -{ - BVHInternalData *internal_data = bvhParams.internalData; - - struct { - uint32_t idx; - uint32_t numInternalNodes; - uint32_t internalNodesOffset; - uint32_t numLeaves; - LBVHNode *leaves; - LBVHNode *internalNodes; - TreeletFormationNode *treeletFormNodes; - } world_info; - - auto update_world_info = [&world_info](uint32_t start_search_idx) { - world_info.idx = bvhParams.instances[start_search_idx].worldID; - world_info.numLeaves = bvhParams.instanceCounts[world_info.idx]; - world_info.numInternalNodes = world_info.numLeaves - 1; - world_info.internalNodesOffset = bvhParams.instanceOffsets[world_info.idx]; - world_info.leaves = internal_data->leaves + world_info.internalNodesOffset; - world_info.internalNodes = internal_data->internalNodes + - world_info.internalNodesOffset; - world_info.treeletFormNodes = internal_data->treeletFormNodes + - world_info.internalNodesOffset; - }; - - update_world_info(start_search_idx); - - // This thread's leaf (tn_offset = thread node offset) - uint32_t tn_offset = start_search_idx - world_info.internalNodesOffset; - LBVHNode *current = &world_info.leaves[tn_offset]; - uint32_t num_leaves = 1; - - // Only the threads which survived push their treelets to shared memory - // (phase 2 shared memory layout). - sm::OptFastBufferTreelets *smem_p2 = - (sm::OptFastBufferTreelets *)sm::buffer; - sm::Treelet *treelets_buffer = (sm::Treelet *) - smem_p2->buffer; - - // TODO: Find break condition here - while (start_search_idx < total_num_instances) { - int32_t parent = current->parent; - - LBVHNode *parent_node = world_info.internalNodes + parent; - TreeletFormationNode *parent_form = world_info.treeletFormNodes + parent; - - parent_form->numLeaves.fetch_add_release(num_leaves); - - if (parent_form->numReached.exchange< - sync::memory_order::relaxed>(1) == 0) { - // Suspend this thread if this is the first thread to reach this node. - // However, if this is the first thread to reach this node but this - // node only has one child, don't suspend. - if (parent_node->numChildren() == 2) { - start_search_idx += num_resident_threads; - update_world_info(start_search_idx); - - tn_offset = start_search_idx - world_info.internalNodesOffset; - current = &world_info.leaves[tn_offset]; - num_leaves = 1; - - continue; - } - } - - // When adding the amount of leaves, exclude what was just - // added by this thread - num_leaves += (parent_form->numLeaves.load_acquire() - num_leaves); - current = &world_info.internalNodes[parent]; - - if (num_leaves > MADRONA_TREELET_SIZE) { - // Push a potential treelet! - sm::InitialTreelet initial_treelet = { - .rootIndex = tn_offset, - .worldIndex = world_info.idx, - .numLeaves = num_leaves - }; - - int32_t treelet_idx = - (int32_t)smem_p2->treeletCounter.fetch_add_relaxed(1); - - treelets_buffer[treelet_idx].initial = initial_treelet; - - // Start a new search - start_search_idx += num_resident_threads; - update_world_info(start_search_idx); - - tn_offset = start_search_idx - world_info.internalNodesOffset; - current = &world_info.leaves[tn_offset]; - num_leaves = 1; - } - } -} - -static __device__ inline void formTreelet(uint32_t treelet_idx) -{ - BVHInternalData *internal_data = bvhParams.internalData; - - uint32_t lane_idx = threadIdx.x % MADRONA_WARP_SIZE; - - sm::OptFastBufferTreelets *smem = (sm::OptFastBufferTreelets *)sm::buffer; - sm::Treelet *treelets = (sm::Treelet *)smem->buffer; - - // Only one thread of the warp actually does the treelet formation - if (lane_idx == 0) { - sm::InitialTreelet *initial = &treelets[treelet_idx].initial; - - uint32_t internal_nodes_offset = - bvhParams.instanceOffsets[initial->worldIndex]; - LBVHNode *internal_nodes = internal_data->internalNodes + - internal_nodes_offset; - LBVHNode *leaf_nodes = internal_data->leaves + internal_nodes_offset; - - sm::FormedTreelet formed_treelet = { - .rootIndex = initial->rootIndex, - .worldIndex = initial->worldIndex - }; - - LBVHNode *current_node = &internal_nodes[formed_treelet.rootIndex]; - uint32_t num_leaves = 2; - formed_treelet.leaves[0] = current_node->left; - formed_treelet.leaves[1] = current_node->right; - - while (num_leaves < MADRONA_TREELET_SIZE) { - int32_t argmax_sah = -1; - float max_sah = -FLT_MAX; - - // Loop through the leaves and figure out which has the largest SAH - // NOTE: you can only replace nodes which are in reality internal - // nodes. Actual leaves cannot be replaced by their children - // because they don't have children. - for (int i = 0; i < num_leaves; ++i) { - bool is_leaf; - int32_t child_idx = LBVHNode::storeIdxToChildIdx( - formed_treelet.leaves[i], is_leaf); - - if (is_leaf) - continue; - - LBVHNode *node = &internal_nodes[child_idx]; - if (node->sah() > max_sah) { - argmax_sah = i; - max_sah = node->sah(); - } - } - - // Now, replace the node with maximum sah with its 2 children. - // This should NEVER fail - // assert(argmax_sah != -1); - - LBVHNode *replaced_node = &internal_nodes[ - formed_treelet.leaves[argmax_sah]]; - - // Normally, after each iteration of the while loop, the number of - // leaves should increase by 1 (until we reach MADRONA_TREELET_SIZE) - formed_treelet.leaves[argmax_sah] = replaced_node->left; - formed_treelet.leaves[num_leaves++] = replaced_node->right; - } - - // Override the bytes in the Treelet struct to reflect the final - // formed treelet. - treelets[treelet_idx].formed = formed_treelet; - } -} - -template -struct WarpRegisterFile -{ - static constexpr uint32_t kNumItemsPerLane = - (N + MADRONA_WARP_SIZE - 1) / MADRONA_WARP_SIZE; - - T items[kNumItemsPerLane]; - - T operator[](uint32_t index) - { - const uint32_t lane_id = index / kNumItemsPerLane; - const uint32_t sub_array_idx = index % kNumItemsPerLane; - return __shlf_sync(0xFFFF'FFFF, items[sub_array_idx], lane_id); - } -}; -#endif // Each warp will maintain a single treelet extern "C" __global__ void bvhOptimizeLBVH() { -#if 0 - BVHInternalData *internal_data = bvhParams.internalData; - - // Phase 1 shared memory layout - sm::OptFastBufferTreelets *smem = (sm::OptFastBufferTreelets *)sm::buffer; - - const uint32_t threads_per_block = blockDim.x; - const uint32_t warps_per_block = threads_per_block / MADRONA_WARP_SIZE; - const uint32_t num_resident_blocks = gridDim.x; - const uint32_t num_resident_threads = threads_per_block * num_resident_blocks; - - if (threadIdx.x == 0) { - uint32_t num_instances = bvhParams.instanceOffsets[bvhParams.numWorlds-1] + - bvhParams.instanceCounts[bvhParams.numWorlds-1]; - smem->totalNumInstances = num_instances; - smem->treeletCounter.store_relaxed(0); - } - - __syncthreads(); - - // For this section, we want all threads who's `thread_inst_offset` is - // beyond the range of allocated instances to lay dormant and wait - // until all the treelets have been formed. - uint32_t thread_inst_offset = blockIdx.x * threads_per_block + threadIdx.x; - uint32_t total_num_instances = smem->totalNumInstances; - - // Push potential roots to shared memory buffer. - pushPotentialRoots(thread_inst_offset, - total_num_instances, - num_resident_threads); - - __syncthreads(); - - // Now, each warp is going to process a single treelet - uint32_t warp_idx = threadIdx.x / MADRONA_WARP_SIZE; - uint32_t lane_idx = threadIdx.x % MADRONA_WARP_SIZE; - uint32_t num_treelets = smem->treeletCounter.load_relaxed(); - - for (uint32_t treelet_idx = warp_idx; - treelet_idx < num_treelets; - treelet_idx += warps_per_block) { - // First, form the treelet of size MADRONA_TREELET_SIZE - formTreelet(treelet_idx); - - // Make the treelet as optimal as possible - - - __syncwarp(); - } - -#endif } extern "C" __global__ void bvhDebug() diff --git a/src/mw/device/bvh_raycast.cpp b/src/mw/device/bvh_raycast.cpp index 1d49ac94..c4e0d1a5 100644 --- a/src/mw/device/bvh_raycast.cpp +++ b/src/mw/device/bvh_raycast.cpp @@ -1,9 +1,5 @@ #define MADRONA_MWGPU_MAX_BLOCKS_PER_SM 4 -#if 0 -#include -#include -#endif #include #include @@ -485,34 +481,7 @@ static TriHitInfo triangleIntersect(int32_t leaf_idx, } } -#if 0 -static void prefetchNode(uint32_t node_idx, - QBVHNode *node_buffer, - QBVHNode *result_smem, - cuda::pipeline &pipe) -{ - pipe.producer_acquire(); - { - cuda::memcpy_async(result_smem, - &node_buffer[node_idx], - sizeof(QBVHNode), - pipe); - } - pipe.producer_commit(); -} -#endif - -#if 0 -static QBVHNode readNode(QBVHNode *node_smem, - cuda::pipeline &pipe) -{ - cuda::pipeline_consumer_wait_prior<0>(pipe); - QBVHNode read_node = *node_smem; - pipe.consumer_release(); - return read_node; -} -#endif static Vector3 hexToRgb(uint32_t hex) { @@ -528,22 +497,7 @@ static __device__ TraceResult traceRay( TraceInfo trace_info, TraceWorldInfo world_info) { -#if 0 - uint8_t *smem_scratch = nullptr; - { - uint32_t linear_tid = - threadIdx.x + - threadIdx.y * blockDim.x + - threadIdx.z * blockDim.x * blockDim.y; - uint32_t bytes_per_thread = smem::kBufSize / - (blockDim.x * blockDim.y * blockDim.z); - smem_scratch = &smem::buffer[linear_tid * bytes_per_thread]; - } -#endif -#if 0 - cuda::pipeline pipe = cuda::make_pipeline(); -#endif // We create these so that we can keep track of the original ray origin, // direction in world space. We will need to transform them when we enter diff --git a/src/mw/device/sort_archetype.cpp b/src/mw/device/sort_archetype.cpp index a8c3b73f..0eb44b79 100644 --- a/src/mw/device/sort_archetype.cpp +++ b/src/mw/device/sort_archetype.cpp @@ -888,30 +888,6 @@ struct SortArchetypeNodeBase::RadixSortOnesweepCustom { } }; -#if 0 && __CUDA_ARCH__ < 800 -static uint32_t __reduce_add_sync(uint32_t mask, uint32_t val) -{ - uint32_t lane_id = threadIdx.x % 32; -#pragma unroll - for (int i = 16; i > 0; i /= 2) { - uint32_t read_lane = lane_id ^ i; - - bool other_active = mask & (1 << read_lane); - - if (!other_active) { - read_lane = lane_id; - } - - uint32_t other = __shfl_sync(mask, val, read_lane); - - if (other_active) { - val += other; - } - } - - return val; -} -#endif SortArchetypeNodeBase::OnesweepNode::OnesweepNode(uint32_t taskgraph_id, ParentNodeT parent, diff --git a/src/render/asset_processor.cpp b/src/render/asset_processor.cpp index 0a4b4c7b..624ba03c 100644 --- a/src/render/asset_processor.cpp +++ b/src/render/asset_processor.cpp @@ -60,10 +60,6 @@ static bool loadCache(const char *location, DynArray nodes{num_nodes}; fread(nodes.data(), sizeof(QBVHNode), num_nodes, ptr); -#if 0 - DynArray leaf_geos{num_leaves}; - fread(leaf_geos.data(), sizeof(MeshBVH::LeafGeometry), num_leaves, ptr); -#endif DynArray vertices{num_verts}; fread(vertices.data(), sizeof(MeshBVH::BVHVertex), num_verts, ptr); diff --git a/src/render/batch_renderer.cpp b/src/render/batch_renderer.cpp index adf2ae76..f9a36e8f 100644 --- a/src/render/batch_renderer.cpp +++ b/src/render/batch_renderer.cpp @@ -2056,20 +2056,6 @@ void BatchRenderer::prepareForRendering(BatchRenderInfo info, 1, &offsets_data_copy); } -#if 0 - { // Import the aabbs for instances - VkDeviceSize num_aabbs_bytes = info.numInstances * sizeof(shader::AABB); - VkBufferCopy aabb_data_copy = { - .srcOffset = 0, - .dstOffset = 0, - .size = num_aabbs_bytes - }; - - impl->dev.dt.cmdCopyBuffer(draw_cmd, interop->aabbHdl, - batch_buffers.aabbs.buffer, - 1, &aabb_data_copy); - } -#endif { // Import the offsets for views VkDeviceSize num_offsets_bytes = info.numWorlds * sizeof(int32_t); diff --git a/src/render/ecs_system.cpp b/src/render/ecs_system.cpp index ee243c03..7a2fc84a 100644 --- a/src/render/ecs_system.cpp +++ b/src/render/ecs_system.cpp @@ -363,34 +363,6 @@ inline void exportCountsGPU(Context &ctx, bvh_internals->numViews = num_views; } -#if 0 - uint32_t *morton_codes = state_mgr->getArchetypeComponent< - RenderableArchetype, MortonCode>(); - - WorldID *world_ids = state_mgr->getArchetypeComponent< - RenderableArchetype, WorldID>(); - - uint32_t current_world = 0; - uint32_t current_world_offset = 0; - - for (int i = 0; - i < state_mgr->getArchetypeNumRows(); - ++i) { - if (world_ids[i].idx != current_world) { - current_world = world_ids[i].idx; - current_world_offset = i; - } - - uint32_t code = morton_codes[i]; - printf(USHORT_TO_BINARY_PATTERN " ", USHORT_TO_BINARY((code>>16))); - printf(USHORT_TO_BINARY_PATTERN " \t", USHORT_TO_BINARY((code))); - - printf("(Leaf node %d)\t %d: (%d)\n", - i - current_world_offset, - world_ids[i].idx, - morton_codes[i]); - } -#endif } #endif @@ -501,11 +473,6 @@ void registerTypes(ECSRegistry ®istry, state_mgr->setArchetypeComponent(bridge->aabbs); } -#if 0 - auto *state_mgr = mwGPU::getStateManager(); - auto instance_ptr = (void *)state_mgr->getArchetypeComponent(); - printf("From rendering system init, instance_ptr=%p\n", instance_ptr); -#endif #else (void)bridge; #endif @@ -597,11 +564,6 @@ TaskGraphNodeID setupTasks(TaskGraphBuilder &builder, builder.addToGraph>( {post_instance_sort_reset_tmp}); -#if 0 - auto sort_views = - builder.addToGraph>( - {post_instance_sort_reset_tmp}); -#endif auto sort_views_world = builder.addToGraph< CompactArchetypeNode>( @@ -656,12 +618,6 @@ void init(Context &ctx, const RenderECSBridge *bridge) system_state.aspectRatio = (float)bridge->renderWidth / (float)bridge->renderHeight; } -#if 0 - bool raycast_enabled = - mwGPU::GPUImplConsts::get().raycastOutputResolution != 0; - - system_state.enableRaycaster = raycast_enabled; -#endif } void makeEntityRenderable(Context &ctx, Entity e) @@ -763,20 +719,7 @@ void makeEntityLightCarrier(Context &ctx, Entity e) }; } -#if 0 -void configureLight(Context &ctx, Entity light, LightDesc desc) -{ - ctx.get(light) = desc; -} -#endif // Add this later when we decide to make the renderer more flexible -#if 0 -void setEntityOutputIndex(Context &ctx, Entity e, uint32_t index) -{ - auto cam_entity = ctx.get(e).cameraEntity; - ctx.get(cam_entity).index = index; -} -#endif } diff --git a/src/render/render_ctx.cpp b/src/render/render_ctx.cpp index 686df267..e21860db 100644 --- a/src/render/render_ctx.cpp +++ b/src/render/render_ctx.cpp @@ -288,14 +288,6 @@ static PipelineShaders makeDrawShaders( shader_path.c_str(), {}, {}, { "frag", ShaderStage::Fragment } ); -#if 0 - {0, 2, repeat_sampler, 1, 0}, - {0, 3, clamp_sampler, 1, 0}, - {1, 1, VK_NULL_HANDLE, - VulkanConfig::max_materials * - VulkanConfig::textures_per_material, - VK_DESCRIPTOR_BINDING_PARTIALLY_BOUND_BIT}, -#endif std::array shaders { std::move(vert_spirv), diff --git a/src/render/shaders/shader_common.h b/src/render/shaders/shader_common.h index 43214a12..d5386f31 100644 --- a/src/render/shaders/shader_common.h +++ b/src/render/shaders/shader_common.h @@ -255,19 +255,5 @@ struct RenderOptions { uint32_t enableAntialiasing; }; -#if 0 -struct PackedDrawInstanceData { - float4 packed[5]; -}; - -struct DrawInstanceData { - float3x3 toViewRot; - float3 toViewTranslation; - float3 objScale; - int32_t viewIdx; - float2 projScale; - float projZNear; -}; -#endif #endif diff --git a/src/render/vk/backend.cpp b/src/render/vk/backend.cpp index 3dd60907..8bb38e9d 100644 --- a/src/render/vk/backend.cpp +++ b/src/render/vk/backend.cpp @@ -842,26 +842,10 @@ Device * Backend::makeDevice( rq_features.pNext = &accel_features; rq_features.rayQuery = true; -#if 0 - VkPhysicalDeviceRobustness2FeaturesEXT robustness_features {}; - robustness_features.sType = - VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ROBUSTNESS_2_FEATURES_EXT; - robustness_features.pNext = &rq_features; - robustness_features.nullDescriptor = true; - - VkPhysicalDeviceLineRasterizationFeaturesEXT line_features {}; - line_features.sType = - VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_LINE_RASTERIZATION_FEATURES_EXT; - line_features.pNext = &robustness_features; - line_features.smoothLines = true; -#endif VkPhysicalDeviceShaderAtomicFloatFeaturesEXT atomic_float_features {}; atomic_float_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_FLOAT_FEATURES_EXT; -#if 0 - atomic_float_features.pNext = &line_features; -#endif if (supports_rt) { atomic_float_features.pNext = &rq_features; } else { @@ -883,15 +867,6 @@ Device * Backend::makeDevice( dyn_render_features.pNext = &subgroup_features; dyn_render_features.dynamicRendering = true; -#if 0 - VkPhysicalDeviceVulkan13Features vk13_features {}; - vk13_features.sType = - VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_3_FEATURES; - vk13_features.pNext = &atomic_float_features; - vk13_features.synchronization2 = true; - vk13_features.computeFullSubgroups = true; - vk13_features.subgroupSizeControl = true; -#endif VkPhysicalDeviceVulkan12Features vk12_features {}; vk12_features.sType = diff --git a/src/render/vk/memory.cpp b/src/render/vk/memory.cpp index 76953596..7867f5e9 100644 --- a/src/render/vk/memory.cpp +++ b/src/render/vk/memory.cpp @@ -62,13 +62,6 @@ static constexpr VkImageUsageFlags depthAttachmentUsage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_SAMPLED_BIT; -#if 0 -static constexpr VkImageUsageFlags rtStorageUsage = - VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_TRANSFER_SRC_BIT; - -static constexpr VkFormatFeatureFlags rtStorageReqs = - VK_FORMAT_FEATURE_STORAGE_IMAGE_BIT | VK_FORMAT_FEATURE_STORAGE_IMAGE_BIT; -#endif }; template