diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj index 445e4deac..aa52865db 100644 --- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj +++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj @@ -10,12 +10,17 @@ 014702732A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */; }; 014702742A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */; }; 014702752A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */; }; - 0147027B2A5B0C010040D02D /* MVKMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 0147027A2A5AF1310040D02D /* MVKMap.h */; }; - 0147027C2A5B0C010040D02D /* MVKMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 0147027A2A5AF1310040D02D /* MVKMap.h */; }; - 0147027D2A5B0C020040D02D /* MVKMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 0147027A2A5AF1310040D02D /* MVKMap.h */; }; 0197951B2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */; }; 0197951C2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */; }; 0197951D2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */ = {isa = PBXBuildFile; fileRef = 0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */; }; + 1155DEB12C50C1BC009D70F8 /* MVKAddressMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */; }; + 1155DEB22C50C1BC009D70F8 /* MVKAddressMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */; }; + 1155DEB32C50C1BC009D70F8 /* MVKAddressMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */; }; + 1155DEB42C50C1BC009D70F8 /* MVKAddressMap.h in Headers */ = {isa = PBXBuildFile; fileRef = 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */; }; + 1155DEB52C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */; }; + 1155DEB62C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */; }; + 1155DEB72C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */; }; + 1155DEB82C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */; }; 2FEA0A4124902F9F00EEF3AD /* MVKExtensions.h in Headers */ = {isa = PBXBuildFile; fileRef = A909F65A213B190600FCD6BE /* MVKExtensions.h */; }; 2FEA0A4224902F9F00EEF3AD /* vk_mvk_moltenvk.h in Headers */ = {isa = PBXBuildFile; fileRef = A94FB7691C7DFB4800632CA3 /* vk_mvk_moltenvk.h */; }; 2FEA0A4324902F9F00EEF3AD /* mvk_datatypes.h in Headers */ = {isa = PBXBuildFile; fileRef = A94FB7671C7DFB4800632CA3 /* mvk_datatypes.h */; }; @@ -665,10 +670,10 @@ /* Begin PBXFileReference section */ 014702702A5855F70040D02D /* MVKCmdAccelerationStructure.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKCmdAccelerationStructure.h; sourceTree = ""; }; 014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCmdAccelerationStructure.mm; sourceTree = ""; }; - 0147027A2A5AF1310040D02D /* MVKMap.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKMap.h; sourceTree = ""; }; - 0147027E2A5B0C9A0040D02D /* MVKMapAllocator.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKMapAllocator.h; sourceTree = ""; }; 019795132A5304D600C6CAD0 /* MVKAccelerationStructure.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKAccelerationStructure.h; sourceTree = ""; }; 0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKAccelerationStructure.mm; sourceTree = ""; }; + 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKAddressMap.h; sourceTree = ""; }; + 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKAddressMap.cpp; sourceTree = ""; }; 2FEA0ABA24902F9F00EEF3AD /* libMoltenVK.a */ = {isa = PBXFileReference; explicitFileType = archive.ar; includeInIndex = 0; path = libMoltenVK.a; sourceTree = BUILT_PRODUCTS_DIR; }; 45003E6F214AD4C900E989CB /* MVKExtensions.def */ = {isa = PBXFileReference; explicitFileType = sourcecode.cpp.h; fileEncoding = 4; path = MVKExtensions.def; sourceTree = ""; }; 4536382D2508A4C6000EFFD3 /* MTLRenderPassStencilAttachmentDescriptor+MoltenVK.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = "MTLRenderPassStencilAttachmentDescriptor+MoltenVK.h"; sourceTree = ""; }; @@ -1045,6 +1050,8 @@ A98149401FB6A3F7005F00B4 /* Utility */ = { isa = PBXGroup; children = ( + 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */, + 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */, A98149421FB6A3F7005F00B4 /* MVKBaseObject.h */, A98149411FB6A3F7005F00B4 /* MVKBaseObject.mm */, A9D7104E25CDE05E00E38106 /* MVKBitArray.h */, @@ -1060,8 +1067,6 @@ A98149441FB6A3F7005F00B4 /* MVKFoundation.h */, A9AB95292B3EDFCC00C4E967 /* MVKInflectionMap.h */, A9F0429E1FB4CF82009FCCB8 /* MVKLogging.h */, - 0147027A2A5AF1310040D02D /* MVKMap.h */, - 0147027E2A5B0C9A0040D02D /* MVKMapAllocator.h */, A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */, A9F3D9DB24732A4D00745190 /* MVKSmallVector.h */, A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */, @@ -1223,6 +1228,7 @@ 2FEA0A6224902F9F00EEF3AD /* MVKMTLBufferAllocation.h in Headers */, 2FEA0A6324902F9F00EEF3AD /* MVKObjectPool.h in Headers */, 2FEA0A6424902F9F00EEF3AD /* MVKSwapchain.h in Headers */, + 1155DEB32C50C1BC009D70F8 /* MVKAddressMap.h in Headers */, 2FEA0A6524902F9F00EEF3AD /* MVKGPUCapture.h in Headers */, 2FEA0A6624902F9F00EEF3AD /* MVKBuffer.h in Headers */, 2FEA0A6724902F9F00EEF3AD /* MVKCommonEnvironment.h in Headers */, @@ -1239,7 +1245,6 @@ 2FEA0A7224902F9F00EEF3AD /* MVKCmdDraw.h in Headers */, A9B3D73C29F9B3B100745CD4 /* mvk_deprecated_api.h in Headers */, 2FEA0A7324902F9F00EEF3AD /* MVKCommandBuffer.h in Headers */, - 0147027C2A5B0C010040D02D /* MVKMap.h in Headers */, 2FEA0A7424902F9F00EEF3AD /* MTLRenderPassDescriptor+MoltenVK.h in Headers */, 2FEA0A7524902F9F00EEF3AD /* MVKCmdDebug.h in Headers */, 2FEA0A7624902F9F00EEF3AD /* MVKWatermarkTextureContent.h in Headers */, @@ -1303,9 +1308,9 @@ A94FB80C1C7DFB4800632CA3 /* MVKShaderModule.h in Headers */, A9AB952B2B3EDFCC00C4E967 /* MVKInflectionMap.h in Headers */, A99C91042295FAC600A061DA /* MVKVulkanAPIObject.h in Headers */, - 0147027B2A5B0C010040D02D /* MVKMap.h in Headers */, A94FB7C01C7DFB4800632CA3 /* MVKCmdQueries.h in Headers */, A9B3D73B29F9B3B100745CD4 /* mvk_deprecated_api.h in Headers */, + 1155DEB22C50C1BC009D70F8 /* MVKAddressMap.h in Headers */, A94FB7CC1C7DFB4800632CA3 /* MVKCommand.h in Headers */, A981494F1FB6A3F7005F00B4 /* MVKBaseObject.h in Headers */, A9C96DD01DDC20C20053187F /* MVKMTLBufferAllocation.h in Headers */, @@ -1384,9 +1389,9 @@ A99C91052295FAC600A061DA /* MVKVulkanAPIObject.h in Headers */, A9AB952D2B3EDFCC00C4E967 /* MVKInflectionMap.h in Headers */, A94FB7C11C7DFB4800632CA3 /* MVKCmdQueries.h in Headers */, - 0147027D2A5B0C020040D02D /* MVKMap.h in Headers */, A94FB7CD1C7DFB4800632CA3 /* MVKCommand.h in Headers */, A9B3D73D29F9B3B100745CD4 /* mvk_deprecated_api.h in Headers */, + 1155DEB12C50C1BC009D70F8 /* MVKAddressMap.h in Headers */, A98149501FB6A3F7005F00B4 /* MVKBaseObject.h in Headers */, A9C96DD11DDC20C20053187F /* MVKMTLBufferAllocation.h in Headers */, A98149581FB6A3F7005F00B4 /* MVKObjectPool.h in Headers */, @@ -1507,6 +1512,7 @@ DCFD7F0F2A45BC6E007BBBF7 /* MVKCmdPipeline.h in Headers */, DCFD7F102A45BC6E007BBBF7 /* MVKSmallVectorAllocator.h in Headers */, DCFD7F112A45BC6E007BBBF7 /* MVKPipeline.h in Headers */, + 1155DEB42C50C1BC009D70F8 /* MVKAddressMap.h in Headers */, DCFD7F122A45BC6E007BBBF7 /* MVKImage.h in Headers */, DCFD7F132A45BC6E007BBBF7 /* MVKBlockObserver.h in Headers */, DCFD7F142A45BC6E007BBBF7 /* MVKCmdTransfer.h in Headers */, @@ -2097,6 +2103,7 @@ 2FEA0A9324902F9F00EEF3AD /* MVKImage.mm in Sources */, 2FEA0A9424902F9F00EEF3AD /* MVKCommandPool.mm in Sources */, 2FEA0A9524902F9F00EEF3AD /* MVKCmdDraw.mm in Sources */, + 1155DEB72C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */, 2FEA0A9624902F9F00EEF3AD /* MVKCommandBuffer.mm in Sources */, 2FEA0A9724902F9F00EEF3AD /* MVKCmdRendering.mm in Sources */, 2FEA0A9824902F9F00EEF3AD /* MVKBuffer.mm in Sources */, @@ -2200,6 +2207,7 @@ A9C96DD21DDC20C20053187F /* MVKMTLBufferAllocation.mm in Sources */, A9E53DE92100B197002781DD /* CAMetalLayer+MoltenVK.mm in Sources */, A9096E5E1F81E16300DFBEA6 /* MVKCmdDispatch.mm in Sources */, + 1155DEB62C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */, A99C90F0229455B300A061DA /* MVKCmdDebug.mm in Sources */, 45E3A40D2166B923005E3E38 /* MTLRenderPipelineColorAttachmentDescriptor+MoltenVK.m in Sources */, ); @@ -2263,6 +2271,7 @@ A9C96DD31DDC20C20053187F /* MVKMTLBufferAllocation.mm in Sources */, A9E53DEA2100B197002781DD /* CAMetalLayer+MoltenVK.mm in Sources */, A9096E5F1F81E16300DFBEA6 /* MVKCmdDispatch.mm in Sources */, + 1155DEB52C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */, A99C90F1229455B300A061DA /* MVKCmdDebug.mm in Sources */, 45E3A40E2166B923005E3E38 /* MTLRenderPipelineColorAttachmentDescriptor+MoltenVK.m in Sources */, ); @@ -2318,6 +2327,7 @@ DCFD7F372A45BC6E007BBBF7 /* MVKInstance.mm in Sources */, DCFD7F382A45BC6E007BBBF7 /* MVKDeviceMemory.mm in Sources */, DCFD7F392A45BC6E007BBBF7 /* MVKImage.mm in Sources */, + 1155DEB82C50C1BC009D70F8 /* MVKAddressMap.cpp in Sources */, DCFD7F3A2A45BC6E007BBBF7 /* MVKCommandPool.mm in Sources */, DCFD7F3B2A45BC6E007BBBF7 /* MVKCmdDraw.mm in Sources */, DCFD7F3C2A45BC6E007BBBF7 /* MVKCommandBuffer.mm in Sources */, diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h index 4971f1efe..a09154bdc 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h +++ b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h @@ -20,6 +20,7 @@ #include "MVKDevice.h" #include "MVKCommand.h" +#include "MVKSmallVector.h" #import #import @@ -37,13 +38,17 @@ class MVKCmdBuildAccelerationStructure : public MVKCommand { const VkAccelerationStructureBuildRangeInfoKHR* const* ppBuildRangeInfos); void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + struct MVKAccelerationStructureBuildInfo + { + VkAccelerationStructureBuildGeometryInfoKHR info; + MVKSmallVector geometries; + MVKSmallVector ranges; + }; protected: MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; - MVKDevice* _mvkDevice; - uint32_t _infoCount; - VkAccelerationStructureBuildGeometryInfoKHR* _geometryInfos; - VkAccelerationStructureBuildRangeInfoKHR const* _buildRangeInfos; + MVKSmallVector _buildInfos; }; #pragma mark - diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm index 2943f718c..0748c31c3 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm @@ -20,9 +20,10 @@ #include "MVKCmdDebug.h" #include "MVKCommandBuffer.h" #include "MVKCommandPool.h" - #include "MVKAccelerationStructure.h" +#include + #pragma mark - #pragma mark MVKCmdBuildAccelerationStructure @@ -30,142 +31,81 @@ uint32_t infoCount, const VkAccelerationStructureBuildGeometryInfoKHR* pInfos, const VkAccelerationStructureBuildRangeInfoKHR* const* ppBuildRangeInfos) { - VkAccelerationStructureBuildGeometryInfoKHR geoInfo = *pInfos; - - _mvkDevice = cmdBuff->getDevice(); - _infoCount = infoCount; - _geometryInfos = &geoInfo; - _buildRangeInfos = *ppBuildRangeInfos; - + _buildInfos.reserve(infoCount); + for (uint32_t i = 0; i < infoCount; i++) + { + MVKAccelerationStructureBuildInfo& info = _buildInfos.emplace_back(); + info.info = pInfos[i]; + + // TODO: ppGeometries + info.geometries.reserve(pInfos[i].geometryCount); + info.ranges.reserve(pInfos[i].geometryCount); + memcpy(info.geometries.data(), pInfos[i].pGeometries, pInfos[i].geometryCount); + memcpy(info.ranges.data(), ppBuildRangeInfos[i], pInfos[i].geometryCount); + + info.info.pGeometries = info.geometries.data(); + } + return VK_SUCCESS; } void MVKCmdBuildAccelerationStructure::encode(MVKCommandEncoder* cmdEncoder) { id accStructEncoder = cmdEncoder->getMTLAccelerationStructureEncoder(kMVKCommandUseBuildAccelerationStructure); - for(int i = 0; i < _infoCount; i++) + for (MVKAccelerationStructureBuildInfo& entry : _buildInfos) { - MVKAccelerationStructure* mvkSrcAccelerationStructure = (MVKAccelerationStructure*)_geometryInfos[i].srcAccelerationStructure; - MVKAccelerationStructure* mvkDstAccelerationStructure = (MVKAccelerationStructure*)_geometryInfos[i].dstAccelerationStructure; - - id srcAccelerationStructure = (id)mvkSrcAccelerationStructure->getMTLAccelerationStructure(); - id dstAccelerationStructure = (id)mvkDstAccelerationStructure->getMTLAccelerationStructure(); + VkAccelerationStructureBuildGeometryInfoKHR& buildInfo = entry.info; + + MVKAccelerationStructure* mvkSrcAccStruct = (MVKAccelerationStructure*)buildInfo.srcAccelerationStructure; + MVKAccelerationStructure* mvkDstAccStruct = (MVKAccelerationStructure*)buildInfo.dstAccelerationStructure; + + id srcAccStruct = mvkSrcAccStruct->getMTLAccelerationStructure(); + id dstAccStruct = mvkDstAccStruct->getMTLAccelerationStructure(); - id srcAccelerationStructureHeap = mvkSrcAccelerationStructure->getMTLHeap(); - id dstAccelerationStructureHeap = mvkDstAccelerationStructure->getMTLHeap(); + id srcAccStructHeap = mvkSrcAccStruct->getMTLHeap(); + id dstAccStructHeap = mvkDstAccStruct->getMTLHeap(); - if(_geometryInfos[i].mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR && !mvkDstAccelerationStructure->getAllowUpdate()) - { + // Should we throw an error here? + // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/vkCmdBuildAccelerationStructuresKHR.html#VUID-vkCmdBuildAccelerationStructuresKHR-pInfos-03667 + if(buildInfo.mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR && !mvkDstAccStruct->getAllowUpdate()) continue; - } - MVKDevice* mvkDvc = cmdEncoder->getDevice(); - MVKBuffer* mvkBuffer = mvkDvc->getBufferAtAddress(_geometryInfos[i].scratchData.deviceAddress); + MVKDevice* mvkDevice = cmdEncoder->getDevice(); + MVKBuffer* mvkBuffer = mvkDevice->getBufferAtAddress(buildInfo.scratchData.deviceAddress); + + // TODO: throw error if mvkBuffer is null? id scratchBuffer = mvkBuffer->getMTLBuffer(); NSInteger scratchBufferOffset = mvkBuffer->getMTLBufferOffset(); - if(_geometryInfos[i].mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_BUILD_KHR) + if (buildInfo.mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_BUILD_KHR) { - if(_geometryInfos[i].type == VK_ACCELERATION_STRUCTURE_TYPE_GENERIC_KHR) - { - MTLAccelerationStructureDescriptor* accStructBuildDescriptor = [MTLAccelerationStructureDescriptor new]; - - if(mvkIsAnyFlagEnabled(_geometryInfos[i].flags, VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_UPDATE_BIT_KHR)){ - accStructBuildDescriptor.usage += MTLAccelerationStructureUsageRefit; - mvkDstAccelerationStructure->setAllowUpdate(true); - }else if(mvkIsAnyFlagEnabled(_geometryInfos[i].flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)){ - accStructBuildDescriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; - }else{ - accStructBuildDescriptor.usage = MTLAccelerationStructureUsageNone; - } - - [dstAccelerationStructureHeap newAccelerationStructureWithDescriptor:accStructBuildDescriptor]; - - [accStructEncoder buildAccelerationStructure:dstAccelerationStructure - descriptor:accStructBuildDescriptor - scratchBuffer:scratchBuffer - scratchBufferOffset:scratchBufferOffset]; - - } - - if(_geometryInfos[i].type == VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR) - { - MTLPrimitiveAccelerationStructureDescriptor* accStructBuildDescriptor = [MTLPrimitiveAccelerationStructureDescriptor new]; - - if(mvkIsAnyFlagEnabled(_geometryInfos[i].flags, VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_UPDATE_BIT_KHR)){ - accStructBuildDescriptor.usage += MTLAccelerationStructureUsageRefit; - mvkDstAccelerationStructure->setAllowUpdate(true); - }else if(mvkIsAnyFlagEnabled(_geometryInfos[i].flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)){ - accStructBuildDescriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; - }else{ - accStructBuildDescriptor.usage = MTLAccelerationStructureUsageNone; - } - - if(_geometryInfos[i].pGeometries->geometryType == VK_GEOMETRY_TYPE_INSTANCES_KHR) { return; } - - if(_geometryInfos[i].pGeometries->geometryType == VK_GEOMETRY_TYPE_TRIANGLES_KHR) - { - VkAccelerationStructureGeometryTrianglesDataKHR triangleGeometryData = _geometryInfos[i].pGeometries->geometry.triangles; - uint64_t vertexBDA = triangleGeometryData.vertexData.deviceAddress; - uint64_t indexBDA = triangleGeometryData.indexData.deviceAddress; - MVKBuffer* mvkVertexBuffer = _mvkDevice->getBufferAtAddress(vertexBDA); - MVKBuffer* mvkIndexBuffer = _mvkDevice->getBufferAtAddress(indexBDA); - - MTLAccelerationStructureTriangleGeometryDescriptor* geometryTriangles = [MTLAccelerationStructureTriangleGeometryDescriptor new]; - geometryTriangles.triangleCount = _geometryInfos[i].geometryCount; - geometryTriangles.vertexBuffer = mvkVertexBuffer->getMTLBuffer(); - geometryTriangles.vertexBufferOffset = mvkVertexBuffer->getMTLBufferOffset(); - - geometryTriangles.indexBuffer = mvkIndexBuffer->getMTLBuffer(); - geometryTriangles.indexBufferOffset = mvkIndexBuffer->getMTLBufferOffset(); - geometryTriangles.indexType = mvkMTLIndexTypeFromVkIndexType(triangleGeometryData.indexType); - accStructBuildDescriptor.geometryDescriptors = @[geometryTriangles]; - - [accStructEncoder buildAccelerationStructure:dstAccelerationStructure - descriptor:accStructBuildDescriptor - scratchBuffer:scratchBuffer - scratchBufferOffset:scratchBufferOffset]; - } - - if(_geometryInfos[i].pGeometries->geometryType == VK_GEOMETRY_TYPE_AABBS_KHR) - { - VkAccelerationStructureGeometryAabbsDataKHR aabbGeometryData = _geometryInfos[i].pGeometries->geometry.aabbs; - uint64_t boundingBoxBDA = aabbGeometryData.data.deviceAddress; - MVKBuffer* mvkBoundingBoxBuffer = _mvkDevice->getBufferAtAddress(boundingBoxBDA); - - MTLAccelerationStructureBoundingBoxGeometryDescriptor* geometryAABBs = [MTLAccelerationStructureBoundingBoxGeometryDescriptor new]; - geometryAABBs.boundingBoxCount = _geometryInfos[i].geometryCount; - geometryAABBs.boundingBoxBuffer = mvkBoundingBoxBuffer->getMTLBuffer(); - geometryAABBs.boundingBoxStride = 0; // Need to get this - geometryAABBs.boundingBoxBufferOffset = mvkBoundingBoxBuffer->getMTLBufferOffset(); - accStructBuildDescriptor.geometryDescriptors = @[geometryAABBs]; - } - } - - if(_geometryInfos[i].type == VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR) - { - MTLInstanceAccelerationStructureDescriptor* accStructInstanceBuildDescriptor = [MTLInstanceAccelerationStructureDescriptor new]; - } + MTLAccelerationStructureDescriptor* descriptor = mvkDstAccStruct->populateMTLDescriptor( + mvkDevice, + buildInfo, + entry.ranges.data(), + nullptr + ); + + [accStructEncoder buildAccelerationStructure:dstAccStruct + descriptor:descriptor + scratchBuffer:scratchBuffer + scratchBufferOffset:scratchBufferOffset]; } - - if(_geometryInfos[i].mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR) + else if (buildInfo.mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR) { - MTLAccelerationStructureDescriptor* accStructRefitDescriptor = [MTLAccelerationStructureDescriptor new]; + MTLAccelerationStructureDescriptor* descriptor = [MTLAccelerationStructureDescriptor new]; - if(mvkIsAnyFlagEnabled(_geometryInfos[i].flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)){ - accStructRefitDescriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; - } + if (mvkIsAnyFlagEnabled(buildInfo.flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)) + descriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; - [accStructEncoder refitAccelerationStructure:srcAccelerationStructure - descriptor:accStructRefitDescriptor - destination:dstAccelerationStructure + [accStructEncoder refitAccelerationStructure:srcAccStruct + descriptor:descriptor + destination:dstAccStruct scratchBuffer:scratchBuffer scratchBufferOffset:scratchBufferOffset]; } } - - return; } #pragma mark - @@ -230,9 +170,9 @@ #pragma mark MVKCmdCopyMemoryToAccelerationStructure VkResult MVKCmdCopyMemoryToAccelerationStructure::setContent(MVKCommandBuffer* cmdBuff, - uint64_t srcAddress, - VkAccelerationStructureKHR dstAccelerationStructure, - VkCopyAccelerationStructureModeKHR copyMode) { + uint64_t srcAddress, + VkAccelerationStructureKHR dstAccelerationStructure, + VkCopyAccelerationStructureModeKHR copyMode) { _srcAddress = srcAddress; _copyMode = copyMode; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h index 9903938a1..e3bb6094b 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h @@ -526,7 +526,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject { uint32_t _flushCount; MVKCommandUse _mtlComputeEncoderUse; MVKCommandUse _mtlBlitEncoderUse; - MVKCommandUse _mtlAccelerationStructureUse; + MVKCommandUse _mtlAccelerationStructureEncoderUse; bool _isRenderingEntireAttachment; }; @@ -542,3 +542,6 @@ NSString* mvkMTLBlitCommandEncoderLabel(MVKCommandUse cmdUse); /** Returns a name, suitable for use as a MTLComputeCommandEncoder label, based on the MVKCommandUse. */ NSString* mvkMTLComputeCommandEncoderLabel(MVKCommandUse cmdUse); + +/** Returns a name, suitable for use as a MTLAccelerationStructureCommandEncoder label, based on the MVKCommandUse. */ +NSString* mvkMTLAccelerationStructureCommandEncoderLabel(MVKCommandUse cmdUse); diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm index 1ce74816b..9ad80d7a0 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm @@ -831,6 +831,10 @@ endMetalEncoding(_mtlBlitEncoder); _mtlBlitEncoderUse = kMVKCommandUseNone; + if (_mtlAccelerationStructureEncoder && _cmdBuffer->_hasStageCounterTimestampCommand) { [_mtlAccelerationStructureEncoder updateFence: getStageCountersMTLFence()]; } + endMetalEncoding(_mtlAccelerationStructureEncoder); + _mtlAccelerationStructureEncoderUse = kMVKCommandUseNone; + encodeTimestampStageCounterSamples(); } @@ -873,9 +877,9 @@ _mtlAccelerationStructureEncoder = [_mtlCmdBuffer accelerationStructureCommandEncoder]; retainIfImmediatelyEncoding(_mtlAccelerationStructureEncoder); } - if (_mtlAccelerationStructureUse != cmdUse) { - _mtlAccelerationStructureUse = cmdUse; - setLabelIfNotNil(_mtlAccelerationStructureEncoder, mvkMTLBlitCommandEncoderLabel(cmdUse)); + if (_mtlAccelerationStructureEncoderUse != cmdUse) { + _mtlAccelerationStructureEncoderUse = cmdUse; + setLabelIfNotNil(_mtlAccelerationStructureEncoder, mvkMTLAccelerationStructureCommandEncoderLabel(cmdUse)); } return _mtlAccelerationStructureEncoder; } @@ -1243,3 +1247,13 @@ default: return @"Unknown Use ComputeEncoder"; } } + +NSString* mvkMTLAccelerationStructureCommandEncoderLabel(MVKCommandUse cmdUse) { + switch (cmdUse) { + case kMVKCommandUseBuildAccelerationStructure: return @"vkCmdBuildAccelerationStructuresKHR AccelerationStructureEncoder"; + case kMVKCommandUseCopyAccelerationStructure: return @"vkCmdCopyAccelerationStructureKHR AccelerationStructureEncoder"; + case kMVKCommandUseCopyAccelerationStructureToMemory: return @"vkCmdCopyAccelerationStructureToMemoryKHR AccelerationStructureEncoder"; + case kMVKCommandUseCopyMemoryToAccelerationStructure: return @"vkCmdCopyMemoryToAccelerationStructureKHR AccelerationStructureEncoder"; + default: return @"Unknown Use AccelerationStructureEncoder"; + } +} diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h index 159a45087..35d166e62 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h @@ -34,7 +34,7 @@ #pragma once -#include "MVKVulkanAPIObject.h" +#include "MVKDevice.h" #import #import @@ -53,8 +53,17 @@ class MVKAccelerationStructure : public MVKVulkanAPIDeviceObject { id getMTLAccelerationStructure(); + /** Populates a MTL acceleration structure descriptor given a vulkan descriptor */ + static MTLAccelerationStructureDescriptor* populateMTLDescriptor(MVKDevice* device, + const VkAccelerationStructureBuildGeometryInfoKHR& buildInfo, + const VkAccelerationStructureBuildRangeInfoKHR* rangeInfos, + const uint32_t* maxPrimitiveCounts); + /** Gets the required build sizes for acceleration structure and scratch buffer*/ - static VkAccelerationStructureBuildSizesInfoKHR getBuildSizes(MVKDevice* device, VkAccelerationStructureBuildTypeKHR buildType, const VkAccelerationStructureBuildGeometryInfoKHR* buildInfo); + static VkAccelerationStructureBuildSizesInfoKHR getBuildSizes(MVKDevice* device, + VkAccelerationStructureBuildTypeKHR buildType, + const VkAccelerationStructureBuildGeometryInfoKHR* buildInfo, + const uint32_t* maxPrimitiveCounts); /** Gets the actual size of the acceleration structure*/ uint64_t getMTLSize(); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm index b94b64800..0052623bd 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm @@ -20,103 +20,177 @@ #include "MVKBuffer.h" #include "MVKAccelerationStructure.h" +#include + #pragma mark - #pragma mark MVKAcceleration Structure -id MVKAccelerationStructure::getMTLAccelerationStructure() { +id MVKAccelerationStructure::getMTLAccelerationStructure() +{ return _accelerationStructure; } - -VkAccelerationStructureBuildSizesInfoKHR MVKAccelerationStructure::getBuildSizes(MVKDevice* device, VkAccelerationStructureBuildTypeKHR type, const VkAccelerationStructureBuildGeometryInfoKHR* info) -{ - VkAccelerationStructureBuildSizesInfoKHR vkBuildSizes{}; - MTLAccelerationStructureDescriptor* accStructDescriptor; - - if(type == VK_ACCELERATION_STRUCTURE_BUILD_TYPE_HOST_KHR) { - // We can't do that, throw an error? - return vkBuildSizes; - } - switch (info->type) +MTLAccelerationStructureDescriptor* MVKAccelerationStructure::populateMTLDescriptor(MVKDevice* device, + const VkAccelerationStructureBuildGeometryInfoKHR& buildInfo, + const VkAccelerationStructureBuildRangeInfoKHR* rangeInfos, + const uint32_t* maxPrimitiveCounts) +{ + MTLAccelerationStructureDescriptor* descriptor = nullptr; + + switch (buildInfo.type) { + default: + break; // TODO: throw error case VK_ACCELERATION_STRUCTURE_TYPE_GENERIC_KHR: { - accStructDescriptor = [MTLAccelerationStructureDescriptor new]; - break; - } - + // TODO: should building generic not be allowed? + // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VkAccelerationStructureTypeKHR.html + } break; + case VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR: { - if(info->pGeometries->geometryType == VK_GEOMETRY_TYPE_AABBS_KHR) - { - accStructDescriptor = [MTLPrimitiveAccelerationStructureDescriptor new]; - MTLPrimitiveAccelerationStructureDescriptor* primitiveAccStructDescriptor = (MTLPrimitiveAccelerationStructureDescriptor*)accStructDescriptor; - - VkAccelerationStructureGeometryAabbsDataKHR aabbGeometryData = info->pGeometries->geometry.aabbs; - uint64_t boundingBoxBDA = aabbGeometryData.data.deviceAddress; - MVKBuffer* mvkBoundingBoxBuffer = device->getBufferAtAddress(boundingBoxBDA); - - MTLAccelerationStructureBoundingBoxGeometryDescriptor* geometryAABBs = [MTLAccelerationStructureBoundingBoxGeometryDescriptor new]; - geometryAABBs.boundingBoxCount = info->geometryCount; - geometryAABBs.boundingBoxBuffer = mvkBoundingBoxBuffer->getMTLBuffer(); - geometryAABBs.boundingBoxStride = 0; // Need to get this - geometryAABBs.boundingBoxBufferOffset = mvkBoundingBoxBuffer->getMTLBufferOffset(); - primitiveAccStructDescriptor.geometryDescriptors = @[geometryAABBs]; - break; - - } - if(info->pGeometries->geometryType == VK_GEOMETRY_TYPE_INSTANCES_KHR) { return vkBuildSizes; } - - if(info->pGeometries->geometryType == VK_GEOMETRY_TYPE_TRIANGLES_KHR) - { - accStructDescriptor = [MTLPrimitiveAccelerationStructureDescriptor new]; - MTLPrimitiveAccelerationStructureDescriptor* primitiveAccStructDescriptor = (MTLPrimitiveAccelerationStructureDescriptor*)accStructDescriptor; - - VkAccelerationStructureGeometryTrianglesDataKHR triangleGeometryData = info->pGeometries->geometry.triangles; - uint64_t vertexBDA = triangleGeometryData.vertexData.deviceAddress; - uint64_t indexBDA = triangleGeometryData.indexData.deviceAddress; - MVKBuffer* mvkVertexBuffer = device->getBufferAtAddress(vertexBDA); - MVKBuffer* mvkIndexBuffer = device->getBufferAtAddress(indexBDA); - - MTLAccelerationStructureTriangleGeometryDescriptor* geometryTriangles = [MTLAccelerationStructureTriangleGeometryDescriptor new]; - geometryTriangles.triangleCount = info->geometryCount; - geometryTriangles.vertexBuffer = mvkVertexBuffer->getMTLBuffer(); - geometryTriangles.vertexBufferOffset = mvkVertexBuffer->getMTLBufferOffset(); - - geometryTriangles.indexBuffer = mvkIndexBuffer->getMTLBuffer(); - geometryTriangles.indexBufferOffset = mvkIndexBuffer->getMTLBufferOffset(); - geometryTriangles.indexType = mvkMTLIndexTypeFromVkIndexType(triangleGeometryData.indexType); - primitiveAccStructDescriptor.geometryDescriptors = @[geometryTriangles]; - break; - } - else + MTLPrimitiveAccelerationStructureDescriptor* primitive = [MTLPrimitiveAccelerationStructureDescriptor new]; + + NSMutableArray* geoms = [NSMutableArray arrayWithCapacity:buildInfo.geometryCount]; + for (uint32_t i = 0; i < buildInfo.geometryCount; i++) { - accStructDescriptor = [MTLPrimitiveAccelerationStructureDescriptor new]; + // TODO: buildInfo.ppGeometries + + const VkAccelerationStructureGeometryKHR& geom = buildInfo.pGeometries[i]; + switch (geom.geometryType) + { + default: + break; + + case VK_GEOMETRY_TYPE_INSTANCES_KHR: + break; + + case VK_GEOMETRY_TYPE_TRIANGLES_KHR: + { + const VkAccelerationStructureGeometryTrianglesDataKHR& triangleData = geom.geometry.triangles; + uint64_t vertexBDA = triangleData.vertexData.deviceAddress; + uint64_t indexBDA = triangleData.indexData.deviceAddress; + uint64_t transformBDA = triangleData.transformData.deviceAddress; + MVKBuffer* mvkVertexBuffer = device->getBufferAtAddress(vertexBDA); + MVKBuffer* mvkIndexBuffer = device->getBufferAtAddress(indexBDA); + MVKBuffer* mvkTransformBuffer = device->getBufferAtAddress(transformBDA); + + // TODO: should validate that buffer->getMTLBufferOffset is a multiple of vertexStride. This could cause issues + NSUInteger vbOffset = (vertexBDA - mvkVertexBuffer->getMTLBufferGPUAddress()) + mvkVertexBuffer->getMTLBufferOffset(); + NSUInteger ibOffset = 0; + NSUInteger tfOffset = 0; + + MTLAccelerationStructureTriangleGeometryDescriptor* geometryTriangles = [MTLAccelerationStructureTriangleGeometryDescriptor new]; + geometryTriangles.vertexBuffer = mvkVertexBuffer->getMTLBuffer(); + geometryTriangles.vertexStride = triangleData.vertexStride; + + if (transformBDA && mvkTransformBuffer) + { + tfOffset = (transformBDA - mvkTransformBuffer->getMTLBufferGPUAddress()) + mvkTransformBuffer->getMTLBufferOffset(); + geometryTriangles.transformationMatrixBuffer = mvkTransformBuffer->getMTLBuffer(); + } + + bool useIndices = indexBDA && mvkIndexBuffer && triangleData.indexType != VK_INDEX_TYPE_NONE_KHR; + if (useIndices) + { + ibOffset = (indexBDA - mvkIndexBuffer->getMTLBufferGPUAddress()) + mvkIndexBuffer->getMTLBufferOffset(); + geometryTriangles.indexBuffer = mvkIndexBuffer->getMTLBuffer(); + geometryTriangles.indexType = mvkMTLIndexTypeFromVkIndexType(triangleData.indexType); + } + + if (rangeInfos) + { + // Utilize range information during build time + + geometryTriangles.triangleCount = rangeInfos[i].primitiveCount; + geometryTriangles.transformationMatrixBufferOffset = tfOffset + rangeInfos[i].transformOffset; + geometryTriangles.vertexBufferOffset = vbOffset; + geometryTriangles.indexBufferOffset = ibOffset + rangeInfos[i].primitiveOffset; + + if (!useIndices) + geometryTriangles.vertexBufferOffset += rangeInfos[i].primitiveOffset + rangeInfos[i].firstVertex * triangleData.vertexStride; + } + else + { + // Less information required when computing size + + geometryTriangles.vertexBufferOffset = vbOffset; + geometryTriangles.triangleCount = maxPrimitiveCounts[i]; + geometryTriangles.indexBufferOffset = ibOffset; + geometryTriangles.transformationMatrixBufferOffset = 0; + } + + [geoms addObject:geometryTriangles]; + } break; + + case VK_GEOMETRY_TYPE_AABBS_KHR: + { + const VkAccelerationStructureGeometryAabbsDataKHR& aabbData = geom.geometry.aabbs; + uint64_t boundingBoxBDA = aabbData.data.deviceAddress; + MVKBuffer* mvkBoundingBoxBuffer = device->getBufferAtAddress(boundingBoxBDA); + + NSUInteger bOffset = (boundingBoxBDA - mvkBoundingBoxBuffer->getMTLBufferGPUAddress()) + mvkBoundingBoxBuffer->getMTLBufferOffset(); + + MTLAccelerationStructureBoundingBoxGeometryDescriptor* geometryAABBs = [MTLAccelerationStructureBoundingBoxGeometryDescriptor new]; + geometryAABBs.boundingBoxStride = aabbData.stride; + geometryAABBs.boundingBoxBuffer = mvkBoundingBoxBuffer->getMTLBuffer(); + geometryAABBs.boundingBoxBufferOffset = bOffset; + + if (rangeInfos) + { + geometryAABBs.boundingBoxCount = rangeInfos[i].primitiveCount; + geometryAABBs.boundingBoxBufferOffset += rangeInfos[i].primitiveOffset; + } + else + geometryAABBs.boundingBoxCount = maxPrimitiveCounts[i]; + + [geoms addObject:geometryAABBs]; + } break; + } } - break; - } + + primitive.geometryDescriptors = geoms; + descriptor = primitive; + } break; + case VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR: { - accStructDescriptor = [MTLInstanceAccelerationStructureDescriptor new]; - MTLInstanceAccelerationStructureDescriptor* instanceAccStructDescriptor = (MTLInstanceAccelerationStructureDescriptor*)accStructDescriptor; + MTLInstanceAccelerationStructureDescriptor* instance = [MTLInstanceAccelerationStructureDescriptor new]; // add bottom level acceleration structures - instanceAccStructDescriptor.instanceDescriptorType = MTLAccelerationStructureInstanceDescriptorTypeDefault; - } - default: - accStructDescriptor = [MTLAccelerationStructureDescriptor new]; - break; + instance.instanceDescriptorType = MTLAccelerationStructureInstanceDescriptorTypeDefault; + + descriptor = instance; + } break; } + + if (!descriptor) + return nullptr; + + if (mvkIsAnyFlagEnabled(buildInfo.flags, VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_UPDATE_BIT_KHR)) + descriptor.usage += MTLAccelerationStructureUsageRefit; + else if (mvkIsAnyFlagEnabled(buildInfo.flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)) + descriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; + else + descriptor.usage = MTLAccelerationStructureUsageNone; + + return descriptor; +} + +VkAccelerationStructureBuildSizesInfoKHR MVKAccelerationStructure::getBuildSizes(MVKDevice* device, + VkAccelerationStructureBuildTypeKHR type, + const VkAccelerationStructureBuildGeometryInfoKHR* info, + const uint32_t* maxPrimitiveCounts) +{ + VkAccelerationStructureBuildSizesInfoKHR vkBuildSizes{}; - if(mvkIsAnyFlagEnabled(info->flags, VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_UPDATE_BIT_KHR)){ - accStructDescriptor.usage += MTLAccelerationStructureUsageRefit; - }else if(mvkIsAnyFlagEnabled(info->flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)){ - accStructDescriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; - }else{ - accStructDescriptor.usage = MTLAccelerationStructureUsageNone; - } + // TODO: We can't perform host builds, throw an error? + if (type == VK_ACCELERATION_STRUCTURE_BUILD_TYPE_HOST_KHR) + return vkBuildSizes; - MTLAccelerationStructureSizes sizes = [device->getMTLDevice() accelerationStructureSizesWithDescriptor: accStructDescriptor]; + MTLAccelerationStructureDescriptor* descriptor = populateMTLDescriptor(device, *info, nullptr, maxPrimitiveCounts); + + MTLAccelerationStructureSizes sizes = [device->getMTLDevice() accelerationStructureSizesWithDescriptor:descriptor]; vkBuildSizes.accelerationStructureSize = sizes.accelerationStructureSize; vkBuildSizes.buildScratchSize = sizes.buildScratchBufferSize; vkBuildSizes.updateScratchSize = sizes.refitScratchBufferSize; @@ -126,7 +200,7 @@ uint64_t MVKAccelerationStructure::getMTLSize() { - if(!_built) { return 0; } + if (!_built) { return 0; } return _accelerationStructure.size; } diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index 4a18f622a..03f4e2ab5 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -26,8 +26,8 @@ #include "MVKSmallVector.h" #include "MVKPixelFormats.h" #include "MVKOSExtensions.h" +#include "MVKAddressMap.h" #include "mvk_datatypes.hpp" -#include "MVKMap.h" #include #include @@ -529,16 +529,11 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { const VkCalibratedTimestampInfoEXT* pTimestampInfos, uint64_t* pTimestamps, uint64_t* pMaxDeviation); - - /** Returns a pointer to the buffer at the provided address*/ - MVKBuffer* getBufferAtAddress(uint64_t address); // Unsure where to place within the file - - /** Returns a pointer to the acceleration structure at the provided address*/ - MVKAccelerationStructure* getAccelerationStructureAtAddress(uint64_t address); - + /** Returns whether or not the device supports acceleration structures*/ VkAccelerationStructureCompatibilityKHR getAccelerationStructureCompatibility(const VkAccelerationStructureVersionInfoKHR* pVersionInfo); + #pragma mark Object lifecycle MVKBuffer* createBuffer(const VkBufferCreateInfo* pCreateInfo, @@ -770,6 +765,12 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { /** Log all performance statistics. */ void logPerformanceSummary(); + + /** Returns a pointer to the buffer at the provided address*/ + MVKBuffer* getBufferAtAddress(uint64_t address); + + /** Returns a pointer to the acceleration structure at the provided address*/ + MVKAccelerationStructure* getAccelerationStructureAtAddress(uint64_t address); #pragma mark Metal @@ -937,9 +938,9 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { MVKSmallVector, kMVKQueueFamilyCount> _queuesByQueueFamilyIndex; MVKSmallVector _resources; MVKSmallVector _gpuAddressableBuffers; - std::unordered_map _gpuBufferAddressMap; + MVKAddressMap* _gpuBufferAddressMap; + uint64_t _nextValidAccStructureAddress = 0; std::unordered_map _gpuAccStructAddressMap; - uint64_t _nextValidAccStructureAddress; MVKSmallVector _privateDataSlots; MVKSmallVector _privateDataSlotsAvailability; MVKSmallVector _awaitingSemaphores; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 0fab842f6..cdb875f60 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -3784,24 +3784,9 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope MVKBuffer* MVKDevice::getBufferAtAddress(uint64_t address) { - lock_guard lock(_rezLock); - - std::unordered_map::iterator it; - // Super inefficent but this can be fixed in the future - for(it = _gpuBufferAddressMap.begin(); it != _gpuBufferAddressMap.end(); it++) - { - // If the beginning address is bigger than, or the ending address is smaller than the passed address, then skip this it - if(it->first.first > address || it->first.second < address) - { - continue; - } - break; - } - - // Couldn't find the buffer at address - if (it == _gpuBufferAddressMap.end()) { return nullptr;} - - return it->second; + void* value = nullptr; + _gpuBufferAddressMap->getValue(address, value); + return (MVKBuffer*)value; } MVKAccelerationStructure* MVKDevice::getAccelerationStructureAtAddress(uint64_t address) @@ -4316,6 +4301,11 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope _resources.push_back(mvkBuff); if (mvkIsAnyFlagEnabled(mvkBuff->getUsage(), VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT)) { _gpuAddressableBuffers.push_back(mvkBuff); + _gpuBufferAddressMap->addEntry({ + mvkBuff->getMTLBufferGPUAddress(), + mvkBuff->getByteCount(), + mvkBuff + }); } return mvkBuff; } @@ -4327,6 +4317,11 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope mvkRemoveFirstOccurance(_resources, mvkBuff); if (mvkIsAnyFlagEnabled(mvkBuff->getUsage(), VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT)) { mvkRemoveFirstOccurance(_gpuAddressableBuffers, mvkBuff); + _gpuBufferAddressMap->removeEntry({ + mvkBuff->getMTLBufferGPUAddress(), + mvkBuff->getByteCount(), + mvkBuff + }); } return mvkBuff; } @@ -4872,6 +4867,8 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope _commandResourceFactory = new MVKCommandResourceFactory(this); + _gpuBufferAddressMap = new MVKAddressMap(); + startAutoGPUCapture(MVK_CONFIG_AUTO_GPU_CAPTURE_SCOPE_DEVICE, getMTLDevice()); MVKLogInfo("Created VkDevice to run on GPU %s with the following %d Vulkan extensions enabled:%s", @@ -5198,6 +5195,8 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope if (_commandResourceFactory) { _commandResourceFactory->destroy(); } + if (_gpuBufferAddressMap) { delete _gpuBufferAddressMap; } + [_globalVisibilityResultMTLBuffer release]; [_defaultMTLSamplerState release]; [_dummyBlitMTLBuffer release]; diff --git a/MoltenVK/MoltenVK/Utility/MVKAddressMap.cpp b/MoltenVK/MoltenVK/Utility/MVKAddressMap.cpp new file mode 100644 index 000000000..3f59b1477 --- /dev/null +++ b/MoltenVK/MoltenVK/Utility/MVKAddressMap.cpp @@ -0,0 +1,211 @@ +/* + * MVKAddressMap.cpp + * + * Copyright (c) 2015-2024 The Brenwill Workshop Ltd. (http://www.brenwill.com) + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "MVKAddressMap.h" +#include + +/** + * Loads the value of an atomic pointer or allocates if it is null in a thread-safe way. + * Returned pointer will never be null. + */ +template +T* loadAtomic(std::atomic& ptr) +{ + T* obj = ptr.load(std::memory_order_acquire); + if (!obj) + { + T* newObj = new T(); + + bool swapped = ptr.compare_exchange_strong(obj, newObj, std::memory_order_release, std::memory_order_acquire); + if (swapped) + obj = newObj; + else + // Someone else allocated, so a new object is no longer needed + delete newObj; + } + + return obj; +} + +MVKAddressMap::~MVKAddressMap() +{ + for (uint64_t i = 0; i < NodeCount; i++) + { + Node* node = _nodes[i].load(std::memory_order_acquire); + if (!node) continue; + + for (uint64_t j = 0; j < BlockCount; j++) + { + SmallStorage* small = node->blocks[j].small.load(std::memory_order_acquire); + if (!small) continue; + + delete small; + } + + delete node; + } +} + +MVKAddressMap::Block* MVKAddressMap::loadBlock(uint64_t addr) +{ + uint64_t blockIdx = getBlockIndex(addr); + uint64_t nodeIdx = getNodeIndex(addr); + + Node* node = loadAtomic(_nodes[nodeIdx]); + + return &node->blocks[blockIdx]; +} + +MVKAddressMap::Block* MVKAddressMap::getBlock(uint64_t addr) const +{ + uint64_t nodeIdx = getNodeIndex(addr); + + Node* node = _nodes[nodeIdx].load(std::memory_order_acquire); + if (!node) + return nullptr; + + uint64_t blockIdx = getBlockIndex(addr); + + return &node->blocks[blockIdx]; +} + +void MVKAddressMap::processEntry(const Entry& entry, bool add) +{ + if (entry.size >= BlockSize) + { + uint64_t low = entry.baseAddress; + uint64_t high = low + entry.size; + + Entry empty{}; + while (low <= high) + { + Block* block = loadBlock(low); + + // If we are adding, insert right only on the first entry, and otherwise + // insert left. If we are removing, we should always reset right and left + // if the value matches. + if (add) + { + if (low == entry.baseAddress) + block->right.store(entry, std::memory_order_relaxed); + else + block->left.store(entry, std::memory_order_relaxed); + } + else + { + if (block->right.load(std::memory_order_relaxed).value == entry.value) + block->right.store(empty, std::memory_order_relaxed); + else if (block->left.load(std::memory_order_relaxed).value == entry.value) + block->left.store(empty, std::memory_order_relaxed); + } + + low += BlockSize; + } + } + else + { + // If the entry is smaller than BlockSize, it is not well-defined to + // mark blocks since one could have multiple small ranges within the same + // block. Thus, these must be stored separately. We will assume that most + // allocations are larger and thus this path is less common. We could optimize + // here and store in a sorted order and binary search later, but that may + // be an unnecessary optimization. + + Block* block = loadBlock(entry.baseAddress); + + SmallStorage* small = loadAtomic(block->small); + + auto lock = std::lock_guard(small->lock); + if (add) + small->entries.emplace_back(entry); + else + { + auto found = std::find_if( + small->entries.begin(), + small->entries.end(), + [&entry](Entry& e) { return e.value == entry.value; } + ); + if (found != small->entries.end()) + small->entries.erase(found); + } + } +} + +void MVKAddressMap::addEntry(const Entry& entry) +{ + processEntry(entry, true); +} + +void MVKAddressMap::removeEntry(const Entry& entry) +{ + processEntry(entry, false); +} + +bool MVKAddressMap::getEntry(uint64_t addr, Entry& outEntry) const +{ + Block* block = getBlock(addr); + + // First check left. This means the address is within the range and the base + // address is to the left. + Entry left = block->left.load(std::memory_order_relaxed); + if (left.baseAddress && addr < left.baseAddress + left.size) + { + outEntry = left; + return true; + } + + // Next check right. This means the base address is within the same block. + Entry right = block->right.load(std::memory_order_relaxed); + if (right.baseAddress && addr >= right.baseAddress) + { + outEntry = right; + return true; + } + + // Otherwise, we need to search for small entries. + SmallStorage* small = block->small.load(std::memory_order_acquire); + if (!small) + return false; + + // Find the small entry where the address is within the range. + auto lock = std::lock_guard(small->lock); + auto found = std::find_if( + small->entries.begin(), + small->entries.end(), + [addr](Entry& e) { return addr >= e.baseAddress && addr < e.baseAddress + e.size; } + ); + if (found != small->entries.end()) + { + outEntry = *found; + return true; + } + + return false; +} + +bool MVKAddressMap::getValue(uint64_t addr, void*& outValue) const +{ + Entry entry; + if (getEntry(addr, entry)) + { + outValue = entry.value; + return true; + } + + return false; +} diff --git a/MoltenVK/MoltenVK/Utility/MVKAddressMap.h b/MoltenVK/MoltenVK/Utility/MVKAddressMap.h new file mode 100644 index 000000000..5402aa801 --- /dev/null +++ b/MoltenVK/MoltenVK/Utility/MVKAddressMap.h @@ -0,0 +1,149 @@ +/* + * MVKAddressMap.h + * + * Copyright (c) 2015-2024 The Brenwill Workshop Ltd. (http://www.brenwill.com) + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "MVKFoundation.h" +#include "MVKSmallVector.h" +#include + +/** + * Maintains a mapping from memory address ranges as keys to arbitrary pointer values. + * + * This data structure is thread-safe. + * + * The map can be queried with any arbitrary address within an inserted range's min and max, + * and they will all map to the same value. + * + * Because not all bits are used in 64-bit memory addresses, this map may not work with + * any arbitrary 64-bit integer range. However, it can always be used with 32-bit integers + * for more generalized use cases. + */ +class MVKAddressMap +{ +public: + + /** + * A key-value entry for the map + */ + struct Entry + { + uint64_t baseAddress; + uint64_t size; + + void* value; + }; + +public: + + /** + * Add an entry to the map. Thread-safe. + * + * The address range must not overlap an existing range, otherwise removal + * and querying are no longer well-defined. + */ + void addEntry(const Entry& entry); + + /** + * Remove an entry to the map. Thread-safe. + * + * The address range must exactly match an existing range, otherwise removal + * and querying are no longer well-defined. + */ + void removeEntry(const Entry& entry); + + /** + * Query the map given an arbitrary address, and return true if it exists. Thread-safe. + * + * Sets outEntry with the queried entry if it exists + */ + bool getEntry(uint64_t addr, Entry& outEntry) const; + + /** + * Query the map given an arbitrary address, and return true if it exists. Thread-safe. + * + * Sets outValue with the queried value if it exists + */ + bool getValue(uint64_t addr, void*& outValue) const; + + ~MVKAddressMap(); + +private: + + static constexpr uint64_t BlockSizeBits = 21; // 2mb + static constexpr uint64_t BlockSize = 1 << BlockSizeBits; + + static constexpr uint64_t BlockCountBits = 18; + static constexpr uint64_t BlockCount = 1 << BlockCountBits; + static constexpr uint64_t BlockCountMask = BlockCount - 1; + + static constexpr uint64_t NodeCountBits = 12; + static constexpr uint64_t NodeCount = 1 << NodeCountBits; + static constexpr uint64_t NodeCountMask = NodeCount - 1; + +private: + + /** Dynamically allocated storage for memory blocks smaller than BlockSize */ + struct SmallStorage + { + std::mutex lock; + MVKSmallVector entries; + }; + + /** Storage for one contiguous memory block of size BlockSize */ + struct Block + { + std::atomic left; + std::atomic right; + + std::atomic small; + }; + + /** Dynamically allocated region with all blocks for that region */ + struct Node + { + Block blocks[BlockCount] = {}; + }; + +private: + + /** + * Load corresponding block where addr is located. Will never return nullptr + * and will allocate if the block was not previously allocated. + */ + Block* loadBlock(uint64_t addr); + + /** + * Get corresponding block where addr is located. Will return nullptr if the + * block was not previously allocated. + */ + Block* getBlock(uint64_t addr) const; + + /** Adds or removes an entry from the map, depending on the value of 'add' */ + void processEntry(const Entry& entry, bool add); + + /** Gets the node index associated with the provided address */ + inline uint64_t getNodeIndex(uint64_t addr) const { return (addr >> (BlockSizeBits + BlockCountBits)) & NodeCountMask; } + + /** Gets the block index associated with the provided address */ + inline uint64_t getBlockIndex(uint64_t addr) const { return (addr >> BlockSizeBits) & BlockCountMask; } + +private: + std::atomic _nodes[NodeCount] = {}; +}; + diff --git a/MoltenVK/MoltenVK/Utility/MVKMap.h b/MoltenVK/MoltenVK/Utility/MVKMap.h deleted file mode 100644 index 19ceaf24b..000000000 --- a/MoltenVK/MoltenVK/Utility/MVKMap.h +++ /dev/null @@ -1,133 +0,0 @@ -/* - * MVKMap.h - * - * Copyright (c) 2015-2023 The Brenwill Workshop Ltd. (http://www.brenwill.com) - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "MVKFoundation.h" -#include "MVKMapAllocator.h" - - -struct MVKHash_uint64_t_pair { // Change the naming style? - size_t operator()(std::pair p) const noexcept { - return size_t(p.first) << 32 | p.second; // Hopefully this is unique enough of a hash, anyway only 1 buffer can ocupy between these 2 memory addreses - } -}; - -template, - class Hash = std::hash> -class MVKMapImpl { - -private: - Allocator alc; - - class iterator - { - const MVKMapImpl* map; - size_t index; - public: - using iterator_category = std::random_access_iterator_tag; - using difference_type = std::ptrdiff_t; - typedef difference_type diff_type; - - iterator() : map{ nullptr }, index{ 0 } { } - iterator(const size_t _index, const MVKMapImpl &_map) : map{ &_map }, index{ _index } { } - - iterator &operator=(const iterator &it) - { - map = it.map; - index = it.index; - return *this; - } - - T *operator->() { return &map->alc.ptr[index]; } - T &operator*() { return map->alc.ptr[index]; } - operator T*() { return &map->alc.ptr[index]; } - - bool operator==( const iterator &it ) const { return map == it.map && index == it.index; } - bool operator!=( const iterator &it ) const { return map != it.map || index != it.index; } - - iterator& operator++() { ++index; return *this; } - iterator operator++( int ) { auto t = *this; ++index; return t; } - iterator& operator--() { --index; return *this; } - iterator operator--( int ) { auto t = *this; --index; return t; } - - iterator operator+ (const diff_type n) { return iterator( index + n, *map ); } - iterator& operator+= (const diff_type n) { index += n; return *this; } - iterator operator- (const diff_type n) { return iterator( index - n, *map ); } - iterator& operator-= (const diff_type n) { index -= n; return *this; } - - diff_type operator- (const iterator& it) { return index - it.index; } - - bool operator< (const iterator& it) { return index < it.index; } - bool operator<= (const iterator& it) { return index <= it.index; } - bool operator> (const iterator& it) { return index > it.index; } - bool operator>= (const iterator& it) { return index >= it.index; } - - const T &operator[]( const diff_type i ) const { return map->alc.ptr[index + i]; } - T &operator[]( const diff_type i ) { return map->alc.ptr[index + i]; } - - bool is_valid() const { return index < map->alc.size(); } - size_t get_position() const { return index; } - }; -protected: - bool empty() { return alc.num_elements_used == 0;} - size_t size() { return alc.size(); } - - T* &at( const size_t i ) { return alc[i]; } - const T* const at(const size_t i) const { return alc[i]; } - - iterator begin() { return iterator(0, this); } - iterator end() { return iterator(size(), this); } - - void erase(const iterator it) - { - if(it.is_valid()) - { - --alc.num_elements_used; - - for(size_t i = it.get_position(); i < alc.num_elements_used; ++i) - { - alc.ptr[i] = alc.ptr[i + 1]; - } - } - } - - void erase(const iterator first, const iterator last) - { - if(first.is_valid()) - { - size_t last_pos = last.is_valid() ? last.get_position() : size(); - size_t n = last_pos - first.get_position(); - alc.num_elements_used -= n; - - for(size_t i = first.get_position(), e = last_pos; i < alc.num_elements_used && e < alc.num_elements_used + n; ++i, ++e) - { - alc.ptr[i] = alc.ptr[e]; - } - } - } - - std::pair insert(const T& value) - { - alc.re_allocate(size() + 1); - alc.ptr[size()] = value; - return std::make_pair(end(), true); - } -}; diff --git a/MoltenVK/MoltenVK/Utility/MVKMapAllocator.h b/MoltenVK/MoltenVK/Utility/MVKMapAllocator.h deleted file mode 100644 index 6387e9d0c..000000000 --- a/MoltenVK/MoltenVK/Utility/MVKMapAllocator.h +++ /dev/null @@ -1,126 +0,0 @@ -/* - * MVKMap.h - * - * Copyright (c) 2015-2023 The Brenwill Workshop Ltd. (http://www.brenwill.com) - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "MVKFoundation.h" -#include - -namespace mvk_map_memory_allocator -{ - inline char *alloc(const size_t num_bytes) - { - return new char[num_bytes]; - } - - inline void free(void *ptr) - { - delete[] (char*)ptr; - } -}; - -template -class mvk_map_allocator final { - -public: - std::pair* ptr; - size_t num_elements_used; -private: - static constexpr size_t CAP_CNT_SIZE = sizeof(size_t); - static constexpr size_t ALIGN_CNT = CAP_CNT_SIZE / sizeof(std::pair); - static constexpr size_t ALIGN_MASK = (ALIGN_CNT> 0) ? (ALIGN_CNT - 1) : 0; - - static constexpr size_t MIN_CNT = M> ALIGN_CNT ? M : ALIGN_CNT; - static constexpr size_t N = (MIN_CNT + ALIGN_MASK) & ~ALIGN_MASK; - - static constexpr size_t MIN_STACK_SIZE = (N * sizeof(std::pair)); - static constexpr size_t STACK_SIZE = MIN_STACK_SIZE> CAP_CNT_SIZE ? MIN_STACK_SIZE : CAP_CNT_SIZE; - alignas(alignof(std::pair)) unsigned char elements_stack[ STACK_SIZE ]; - - void set_num_elements_reserved(const size_t num_elements_reserved) - { - *reinterpret_cast(&elements_stack[0]) = num_elements_reserved; - } -public: - const T &operator[](const size_t i) const { return ptr[i]; } - T &operator[](const size_t i) { return ptr[i]; } - - size_t size() const { return num_elements_used; } - - constexpr T *get_default_ptr() const - { - return reinterpret_cast(const_cast(&elements_stack[0])); - } - - template typename std::enable_if::value>::type - construct(S *_ptr, Args&&... _args) - { - new (_ptr) S(std::forward(_args)...); - } - - template typename std::enable_if::value>::type - construct(S *_ptr, Args&&... _args) - { - *_ptr = S(std::forward(_args)...); - } - - template typename std::enable_if::value>::type - destruct(S *_ptr) - { - _ptr->~S(); - } - - template typename std::enable_if::value>::type - destruct(S *_ptr) {} - - template typename std::enable_if::value>::type - destruct_all() - { - for(size_t i = 0; i < num_elements_used; ++i) - { - ptr[i].~S(); - } - - num_elements_used = 0; - } - - template typename std::enable_if::value>::type - destruct_all() - { - num_elements_used = 0; - } - - void re_allocate(const size_t num_elements_to_reserve) - { - auto *new_ptr = reinterpret_cast(mvk_smallvector_memory_allocator::alloc(num_elements_to_reserve * sizeof(T))); - - for(size_t i = 0; i < num_elements_used; ++i) - { - construct(&new_ptr[i], std::move(ptr[i])); - destruct(&ptr[i]); - } - - if(ptr != get_default_ptr()) - { - mvk_smallvector_memory_allocator::free(ptr); - } - - ptr = new_ptr; - set_num_elements_reserved(num_elements_to_reserve); - } -}; diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index cedb0cc87..675f95d98 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -2919,7 +2919,7 @@ MVK_PUBLIC_VULKAN_SYMBOL void vkGetAccelerationStructureBuildSizesKHR( MVKTraceVulkanCallStart(); MVKDevice* mvkDev = (MVKDevice*)device; - VkAccelerationStructureBuildSizesInfoKHR buildSizes = MVKAccelerationStructure::getBuildSizes(mvkDev, buildType, pBuildInfo); + VkAccelerationStructureBuildSizesInfoKHR buildSizes = MVKAccelerationStructure::getBuildSizes(mvkDev, buildType, pBuildInfo, pMaxPrimitiveCounts); pSizeInfo = &buildSizes; MVKTraceVulkanCallEnd(); }