diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj index 91511eeb9..aa52865db 100644 --- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj +++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj @@ -7,6 +7,20 @@ objects = { /* Begin PBXBuildFile section */ + 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 */; }; + 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 */; }; @@ -654,6 +668,12 @@ /* End PBXContainerItemProxy section */ /* 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 = ""; }; + 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 = ""; }; @@ -920,6 +940,8 @@ A94FB76B1C7DFB4800632CA3 /* Commands */ = { isa = PBXGroup; children = ( + 014702702A5855F70040D02D /* MVKCmdAccelerationStructure.h */, + 014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */, A99C90EC229455B200A061DA /* MVKCmdDebug.h */, A99C90ED229455B300A061DA /* MVKCmdDebug.mm */, A9096E5C1F81E16300DFBEA6 /* MVKCmdDispatch.h */, @@ -957,6 +979,8 @@ A94FB77E1C7DFB4800632CA3 /* GPUObjects */ = { isa = PBXGroup; children = ( + 019795132A5304D600C6CAD0 /* MVKAccelerationStructure.h */, + 0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */, A94FB77F1C7DFB4800632CA3 /* MVKBuffer.h */, A94FB7801C7DFB4800632CA3 /* MVKBuffer.mm */, A966A5DC23C535D000BBF9B4 /* MVKDescriptor.h */, @@ -1026,6 +1050,8 @@ A98149401FB6A3F7005F00B4 /* Utility */ = { isa = PBXGroup; children = ( + 1155DEB02C50C1BC009D70F8 /* MVKAddressMap.cpp */, + 1155DEAF2C50C1BC009D70F8 /* MVKAddressMap.h */, A98149421FB6A3F7005F00B4 /* MVKBaseObject.h */, A98149411FB6A3F7005F00B4 /* MVKBaseObject.mm */, A9D7104E25CDE05E00E38106 /* MVKBitArray.h */, @@ -1046,8 +1072,8 @@ A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */, A98149491FB6A3F7005F00B4 /* MVKWatermark.h */, A981494A1FB6A3F7005F00B4 /* MVKWatermark.mm */, - A981494B1FB6A3F7005F00B4 /* MVKWatermarkShaderSource.h */, A981494C1FB6A3F7005F00B4 /* MVKWatermarkTextureContent.h */, + A981494B1FB6A3F7005F00B4 /* MVKWatermarkShaderSource.h */, ); path = Utility; sourceTree = ""; @@ -1202,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 */, @@ -1283,6 +1310,7 @@ A99C91042295FAC600A061DA /* MVKVulkanAPIObject.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 */, @@ -1363,6 +1391,7 @@ A94FB7C11C7DFB4800632CA3 /* MVKCmdQueries.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 */, @@ -1483,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 */, @@ -2073,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 */, @@ -2089,6 +2120,7 @@ 2FEA0AA124902F9F00EEF3AD /* MVKQueue.mm in Sources */, 2FEA0AA224902F9F00EEF3AD /* MTLSamplerDescriptor+MoltenVK.m in Sources */, 2FEA0AA324902F9F00EEF3AD /* MVKRenderPass.mm in Sources */, + 0197951C2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */, 2FEA0AA424902F9F00EEF3AD /* MVKCmdTransfer.mm in Sources */, 2FEA0AA524902F9F00EEF3AD /* MVKCmdQueries.mm in Sources */, 2FEA0AA624902F9F00EEF3AD /* mvk_api.mm in Sources */, @@ -2104,6 +2136,7 @@ 2FEA0AB024902F9F00EEF3AD /* MVKFramebuffer.mm in Sources */, 2FEA0AB124902F9F00EEF3AD /* MVKMTLBufferAllocation.mm in Sources */, 2FEA0AB224902F9F00EEF3AD /* CAMetalLayer+MoltenVK.mm in Sources */, + 014702742A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */, 2FEA0AB324902F9F00EEF3AD /* MVKCmdDispatch.mm in Sources */, 2FEA0AB424902F9F00EEF3AD /* MVKCmdDebug.mm in Sources */, ); @@ -2149,7 +2182,9 @@ A98149551FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */, A9653FBC24129C84005999D7 /* MVKPixelFormats.mm in Sources */, A94FB7E61C7DFB4800632CA3 /* MVKDevice.mm in Sources */, + 014702732A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */, A9E53DF52100B302002781DD /* MTLRenderPassDescriptor+MoltenVK.m in Sources */, + 0197951B2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */, A966A5E123C535D000BBF9B4 /* MVKDescriptor.mm in Sources */, A94FB7FA1C7DFB4800632CA3 /* MVKPipeline.mm in Sources */, A94FB8021C7DFB4800632CA3 /* MVKQueue.mm in Sources */, @@ -2172,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 */, ); @@ -2210,7 +2246,9 @@ A98149561FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */, A9653FBD24129C84005999D7 /* MVKPixelFormats.mm in Sources */, A94FB7E71C7DFB4800632CA3 /* MVKDevice.mm in Sources */, + 014702752A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */, A9E53DF62100B302002781DD /* MTLRenderPassDescriptor+MoltenVK.m in Sources */, + 0197951D2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm in Sources */, A966A5E223C535D000BBF9B4 /* MVKDescriptor.mm in Sources */, A94FB7FB1C7DFB4800632CA3 /* MVKPipeline.mm in Sources */, A94FB8031C7DFB4800632CA3 /* MVKQueue.mm in Sources */, @@ -2233,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 */, ); @@ -2288,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 new file mode 100644 index 000000000..a09154bdc --- /dev/null +++ b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h @@ -0,0 +1,147 @@ +/* + * MVKCmdAccelerationStructure.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 "MVKDevice.h" +#include "MVKCommand.h" +#include "MVKSmallVector.h" + +#import +#import +#import + +#pragma mark - +#pragma mark MVKCmdBuildAccelerationStructure + +class MVKCmdBuildAccelerationStructure : public MVKCommand { + +public: + VkResult setContent(MVKCommandBuffer* cmdBuff, + uint32_t infoCount, + const VkAccelerationStructureBuildGeometryInfoKHR* pInfos, + const VkAccelerationStructureBuildRangeInfoKHR* const* ppBuildRangeInfos); + + void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + struct MVKAccelerationStructureBuildInfo + { + VkAccelerationStructureBuildGeometryInfoKHR info; + MVKSmallVector geometries; + MVKSmallVector ranges; + }; +protected: + MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; + + MVKSmallVector _buildInfos; +}; + +#pragma mark - +#pragma mark MVKCmdCopyAccelerationStructure + +class MVKCmdCopyAccelerationStructure : public MVKCommand { + +public: + VkResult setContent(MVKCommandBuffer* cmdBuff, + VkAccelerationStructureKHR srcAccelerationStructure, + VkAccelerationStructureKHR dstAccelerationStructure, + VkCopyAccelerationStructureModeKHR copyMode); + + void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; + + id _srcAccelerationStructure; + id _dstAccelerationStructure; + + VkCopyAccelerationStructureModeKHR _copyMode; +}; + +#pragma mark - +#pragma mark MVKCmdCopyAccelerationStructureToMemory + +class MVKCmdCopyAccelerationStructureToMemory : public MVKCommand { + +public: + VkResult setContent(MVKCommandBuffer* cmdBuff, + VkAccelerationStructureKHR srcAccelerationStructure, + uint64_t dstAddress, + VkCopyAccelerationStructureModeKHR copyMode); + + void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; + + id _srcAccelerationStructure; + id _srcAccelerationStructureBuffer; + MVKBuffer* _dstBuffer; + uint64_t _copySize; + + uint64_t _dstAddress; + MVKDevice* _mvkDevice; + VkCopyAccelerationStructureModeKHR _copyMode; +}; + +#pragma mark - +#pragma mark MVKCmdCopyMemoryToAccelerationStructure + +class MVKCmdCopyMemoryToAccelerationStructure: public MVKCommand { + +public: + VkResult setContent(MVKCommandBuffer* cmdBuff, + uint64_t srcAddress, + VkAccelerationStructureKHR dstAccelerationStructure, + VkCopyAccelerationStructureModeKHR copyMode); + + void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; + + MVKBuffer* _srcBuffer; + id _dstAccelerationStructure; + id _dstAccelerationStructureBuffer; + uint32_t _copySize; + + uint64_t _srcAddress; + MVKDevice* _mvkDevice; + VkCopyAccelerationStructureModeKHR _copyMode; +}; + +#pragma mark - +#pragma mark MVKCmdWriteAccelerationStructuresProperties + +class MVKCmdWriteAccelerationStructuresProperties: public MVKCommand { + +public: + VkResult setContent(MVKCommandBuffer* cmdBuff, + uint32_t accelerationStructureCount, + const VkAccelerationStructureKHR* pAccelerationStructures, + VkQueryType queryType, + VkQueryPool queryPool, + uint32_t firstQuery); + + void encode(MVKCommandEncoder* cmdEncoder) override; +protected: + MVKCommandTypePool* getTypePool(MVKCommandPool* cmdPool) override; + + uint32_t _accelerationStructureCount; + const MVKAccelerationStructure* _pAccelerationStructures; + VkQueryType _queryType; + VkQueryPool _queryPool; + uint32_t _firstQuery; +}; diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm new file mode 100644 index 000000000..0748c31c3 --- /dev/null +++ b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm @@ -0,0 +1,236 @@ +/* + * MVKCmdAccelerationStructure.mm + * + * 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. + */ + +#include "MVKCmdAccelerationStructure.h" +#include "MVKCmdDebug.h" +#include "MVKCommandBuffer.h" +#include "MVKCommandPool.h" +#include "MVKAccelerationStructure.h" + +#include + +#pragma mark - +#pragma mark MVKCmdBuildAccelerationStructure + +VkResult MVKCmdBuildAccelerationStructure::setContent(MVKCommandBuffer* cmdBuff, + uint32_t infoCount, + const VkAccelerationStructureBuildGeometryInfoKHR* pInfos, + const VkAccelerationStructureBuildRangeInfoKHR* const* 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 (MVKAccelerationStructureBuildInfo& entry : _buildInfos) + { + VkAccelerationStructureBuildGeometryInfoKHR& buildInfo = entry.info; + + MVKAccelerationStructure* mvkSrcAccStruct = (MVKAccelerationStructure*)buildInfo.srcAccelerationStructure; + MVKAccelerationStructure* mvkDstAccStruct = (MVKAccelerationStructure*)buildInfo.dstAccelerationStructure; + + id srcAccStruct = mvkSrcAccStruct->getMTLAccelerationStructure(); + id dstAccStruct = mvkDstAccStruct->getMTLAccelerationStructure(); + + id srcAccStructHeap = mvkSrcAccStruct->getMTLHeap(); + id dstAccStructHeap = mvkDstAccStruct->getMTLHeap(); + + // 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* 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 (buildInfo.mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_BUILD_KHR) + { + MTLAccelerationStructureDescriptor* descriptor = mvkDstAccStruct->populateMTLDescriptor( + mvkDevice, + buildInfo, + entry.ranges.data(), + nullptr + ); + + [accStructEncoder buildAccelerationStructure:dstAccStruct + descriptor:descriptor + scratchBuffer:scratchBuffer + scratchBufferOffset:scratchBufferOffset]; + } + else if (buildInfo.mode == VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR) + { + MTLAccelerationStructureDescriptor* descriptor = [MTLAccelerationStructureDescriptor new]; + + if (mvkIsAnyFlagEnabled(buildInfo.flags, VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR)) + descriptor.usage += MTLAccelerationStructureUsagePreferFastBuild; + + [accStructEncoder refitAccelerationStructure:srcAccStruct + descriptor:descriptor + destination:dstAccStruct + scratchBuffer:scratchBuffer + scratchBufferOffset:scratchBufferOffset]; + } + } +} + +#pragma mark - +#pragma mark MVKCmdCopyAccelerationStructure + +VkResult MVKCmdCopyAccelerationStructure::setContent(MVKCommandBuffer* cmdBuff, + VkAccelerationStructureKHR srcAccelerationStructure, + VkAccelerationStructureKHR dstAccelerationStructure, + VkCopyAccelerationStructureModeKHR copyMode) { + + MVKAccelerationStructure* mvkSrcAccStruct = (MVKAccelerationStructure*)srcAccelerationStructure; + MVKAccelerationStructure* mvkDstAccStruct = (MVKAccelerationStructure*)dstAccelerationStructure; + + _srcAccelerationStructure = mvkSrcAccStruct->getMTLAccelerationStructure(); + _dstAccelerationStructure = mvkDstAccStruct->getMTLAccelerationStructure(); + _copyMode = copyMode; + return VK_SUCCESS; +} + +void MVKCmdCopyAccelerationStructure::encode(MVKCommandEncoder* cmdEncoder) { + id accStructEncoder = cmdEncoder->getMTLAccelerationStructureEncoder(kMVKCommandUseCopyAccelerationStructure); + if(_copyMode == VK_COPY_ACCELERATION_STRUCTURE_MODE_COMPACT_KHR) + { + [accStructEncoder + copyAndCompactAccelerationStructure:_srcAccelerationStructure + toAccelerationStructure:_dstAccelerationStructure]; + + return; + } + + [accStructEncoder + copyAccelerationStructure:_srcAccelerationStructure + toAccelerationStructure:_dstAccelerationStructure]; +} + +#pragma mark - +#pragma mark MVKCmdCopyAccelerationStructureToMemory + +VkResult MVKCmdCopyAccelerationStructureToMemory::setContent(MVKCommandBuffer* cmdBuff, + VkAccelerationStructureKHR srcAccelerationStructure, + uint64_t dstAddress, + VkCopyAccelerationStructureModeKHR copyMode) { + _dstAddress = dstAddress; + _copyMode = copyMode; + + MVKAccelerationStructure* mvkSrcAccStruct = (MVKAccelerationStructure*)srcAccelerationStructure; + _srcAccelerationStructure = mvkSrcAccStruct->getMTLAccelerationStructure(); + + + _dstBuffer = _mvkDevice->getBufferAtAddress(_dstAddress); + return VK_SUCCESS; +} + +void MVKCmdCopyAccelerationStructureToMemory::encode(MVKCommandEncoder* cmdEncoder) { + id blitEncoder = cmdEncoder->getMTLBlitEncoder(kMVKCommandUseCopyAccelerationStructureToMemory); + _mvkDevice = cmdEncoder->getDevice(); + + [blitEncoder copyFromBuffer:_srcAccelerationStructureBuffer sourceOffset:0 toBuffer:_dstBuffer->getMTLBuffer() destinationOffset:0 size:_copySize]; +} + +#pragma mark - +#pragma mark MVKCmdCopyMemoryToAccelerationStructure + +VkResult MVKCmdCopyMemoryToAccelerationStructure::setContent(MVKCommandBuffer* cmdBuff, + uint64_t srcAddress, + VkAccelerationStructureKHR dstAccelerationStructure, + VkCopyAccelerationStructureModeKHR copyMode) { + _srcAddress = srcAddress; + _copyMode = copyMode; + + _srcBuffer = _mvkDevice->getBufferAtAddress(_srcAddress); + + MVKAccelerationStructure* mvkDstAccStruct = (MVKAccelerationStructure*)dstAccelerationStructure; + _dstAccelerationStructure = mvkDstAccStruct->getMTLAccelerationStructure(); + _dstAccelerationStructureBuffer = mvkDstAccStruct->getMTLBuffer(); + return VK_SUCCESS; +} + +void MVKCmdCopyMemoryToAccelerationStructure::encode(MVKCommandEncoder* cmdEncoder) { + id blitEncoder = cmdEncoder->getMTLBlitEncoder(kMVKCommandUseCopyAccelerationStructureToMemory); + _mvkDevice = cmdEncoder->getDevice(); + + [blitEncoder copyFromBuffer:_srcBuffer->getMTLBuffer() sourceOffset:0 toBuffer:_dstAccelerationStructureBuffer destinationOffset:0 size:_copySize]; +} + +#pragma mark - +#pragma mark MVKCmdWriteAccelerationStructuresProperties + +VkResult MVKCmdWriteAccelerationStructuresProperties::setContent(MVKCommandBuffer* cmdBuff, + uint32_t accelerationStructureCount, + const VkAccelerationStructureKHR* pAccelerationStructures, + VkQueryType queryType, + VkQueryPool queryPool, + uint32_t firstQuery) { + + _accelerationStructureCount = accelerationStructureCount; + _pAccelerationStructures = (const MVKAccelerationStructure*)pAccelerationStructures; + _queryType = queryType; + _queryPool = queryPool; + _firstQuery = firstQuery; + return VK_SUCCESS; +} + +void MVKCmdWriteAccelerationStructuresProperties::encode(MVKCommandEncoder* cmdEncoder) { + + for(int i = 0; i < _accelerationStructureCount; i++) + { + if(!_pAccelerationStructures[i].getBuildStatus()) { + return; + } + + // actually finish up the meat of the code here + } + + switch(_queryType) + { + case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR: + break; + case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR: + break; + case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR: + break; + case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR: + break; + default: + break; + } +} diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h index 6bcea53f7..e3bb6094b 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h @@ -350,6 +350,8 @@ class MVKCommandEncoder : public MVKBaseDeviceObject { * the current encoder before beginning BLIT encoding. */ id getMTLBlitEncoder(MVKCommandUse cmdUse); + + id getMTLAccelerationStructureEncoder(MVKCommandUse cmdUse); // Write proper comment above /** * Returns the current Metal encoder, which may be any of the Metal render, @@ -510,6 +512,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject { MVKSmallVector _attachments; id _mtlComputeEncoder; id _mtlBlitEncoder; + id _mtlAccelerationStructureEncoder; id _stageCountersMTLFence; MVKPushConstantsCommandEncoderState _vertexPushConstants; MVKPushConstantsCommandEncoderState _tessCtlPushConstants; @@ -523,6 +526,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject { uint32_t _flushCount; MVKCommandUse _mtlComputeEncoderUse; MVKCommandUse _mtlBlitEncoderUse; + MVKCommandUse _mtlAccelerationStructureEncoderUse; bool _isRenderingEntireAttachment; }; @@ -538,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 031830aa3..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(); } @@ -867,6 +871,19 @@ return _mtlBlitEncoder; } +id MVKCommandEncoder::getMTLAccelerationStructureEncoder(MVKCommandUse cmdUse) { + if ( !_mtlAccelerationStructureEncoder ) { + endCurrentMetalEncoding(); + _mtlAccelerationStructureEncoder = [_mtlCmdBuffer accelerationStructureCommandEncoder]; + retainIfImmediatelyEncoding(_mtlAccelerationStructureEncoder); + } + if (_mtlAccelerationStructureEncoderUse != cmdUse) { + _mtlAccelerationStructureEncoderUse = cmdUse; + setLabelIfNotNil(_mtlAccelerationStructureEncoder, mvkMTLAccelerationStructureCommandEncoderLabel(cmdUse)); + } + return _mtlAccelerationStructureEncoder; +} + id MVKCommandEncoder::getMTLEncoder(){ if (_mtlRenderEncoder) { return _mtlRenderEncoder; } if (_mtlComputeEncoder) { return _mtlComputeEncoder; } @@ -1230,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/Commands/MVKCommandPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h index 3c2768514..4424a5fc5 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h @@ -22,6 +22,7 @@ #include "MVKCommandBuffer.h" #include "MVKCommandEncodingPool.h" #include "MVKCommand.h" +#include "MVKCmdAccelerationStructure.h" #include "MVKCmdPipeline.h" #include "MVKCmdRendering.h" #include "MVKCmdDispatch.h" diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def b/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def index a79fd1992..dea1ade84 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def +++ b/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def @@ -141,7 +141,12 @@ MVK_CMD_TYPE_POOL(DebugMarkerEnd) MVK_CMD_TYPE_POOL(DebugMarkerInsert) MVK_CMD_TYPE_POOLS_FROM_THRESHOLD(WaitEvents, 1) MVK_CMD_TYPE_POOL(SetEvent) -MVK_CMD_TYPE_POOL_LAST(ResetEvent) +MVK_CMD_TYPE_POOL(ResetEvent) +MVK_CMD_TYPE_POOL(BuildAccelerationStructure) +MVK_CMD_TYPE_POOL(CopyAccelerationStructure) +MVK_CMD_TYPE_POOL(CopyAccelerationStructureToMemory) +MVK_CMD_TYPE_POOL(CopyMemoryToAccelerationStructure) +MVK_CMD_TYPE_POOL_LAST(WriteAccelerationStructuresProperties) #undef MVK_CMD_TYPE_POOL #undef MVK_CMD_TYPE_POOL_LAST diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h new file mode 100644 index 000000000..35d166e62 --- /dev/null +++ b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h @@ -0,0 +1,112 @@ +/* + * MVKAccelerationStructure.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, + vkGetDeviceAccelerationStructureCompatibilityKHR * 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. + */ + +/* + Commands that need to be implemented + + vkCmdBuildAccelerationStructuresIndirectKHR + vkCmdBuildAccelerationStructuresKHR + vkCmdCopyAccelerationStructureKHR + vkCmdCopyAccelerationStructureToMemoryKHR + vkCmdCopyMemoryToAccelerationStructureKHR + vkCmdWriteAccelerationStructuresPropertiesKHR + vkCreateAccelerationStructureKHR - Complete + vkDestroyAccelerationStructureKHR - Complete + vkGetAccelerationStructureBuildSizesKHR + vkGetAccelerationStructureDeviceAddressKHR - Complete + vkGetDeviceAccelerationStructureCompatibilityKHR - Complete +*/ + +#pragma once + +#include "MVKDevice.h" + +#import +#import + +#pragma mark - +#pragma mark MVKAccelerationStructure + +class MVKAccelerationStructure : public MVKVulkanAPIDeviceObject { + +public: + VkObjectType getVkObjectType() override { return VK_OBJECT_TYPE_ACCELERATION_STRUCTURE_KHR; } + + VkDebugReportObjectTypeEXT getVkDebugReportObjectType() override { + return VK_DEBUG_REPORT_OBJECT_TYPE_ACCELERATION_STRUCTURE_KHR_EXT; + } + + 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, + const uint32_t* maxPrimitiveCounts); + + /** Gets the actual size of the acceleration structure*/ + uint64_t getMTLSize(); + +#pragma mark - +#pragma mark Getters and Setters + /** Used when building the acceleration structure, to mark whether or not an acceleration structure can be updated, only to be set by MVKCmdBuildAccelerationStructure*/ + void setAllowUpdate(bool value) { _allowUpdate = value; } + + /** Checks if this acceleration structure is allowed to be updated*/ + bool getAllowUpdate() const { return _allowUpdate; } + + /** Only to be called by the MVKCmdBuildAccelerationStructure, and sets the build status*/ + void setBuildStatus(bool value) { _built = value; } + + /** Checks if this acceleration structure has been built*/ + bool getBuildStatus() const { return _built; } + + /** Sets the address of the acceleration structure, only to be used by MVKDevice*/ + void setDeviceAddress(uint64_t address) { _address = address; } + + /** Gets the address of the acceleration structure*/ + uint64_t getDeviceAddress() const { return _address; } + + /** Returns the Metal buffer using the same memory as the acceleration structure*/ + id getMTLBuffer() const { return _buffer; } + + /** Gets the heap allocation that the acceleration structure, and buffer share*/ + id getMTLHeap() const { return _heap; } + + MTLAccelerationStructureTriangleGeometryDescriptor* getTriangleDescriptor(); +#pragma mark - +#pragma mark Construction + MVKAccelerationStructure(MVKDevice* device); + void destroy() override; +protected: + void propagateDebugName() override {} + + id _heap; + id _accelerationStructure; + id _buffer; + + bool _allowUpdate = false; + bool _built = false; + uint64_t _address = 0; +}; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm new file mode 100644 index 000000000..0052623bd --- /dev/null +++ b/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm @@ -0,0 +1,222 @@ +/* + * MVKAccelerationStructure.mm + * + * 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. + */ + +#include "MVKDevice.h" +#include "MVKBuffer.h" +#include "MVKAccelerationStructure.h" + +#include + +#pragma mark - +#pragma mark MVKAcceleration Structure + +id MVKAccelerationStructure::getMTLAccelerationStructure() +{ + return _accelerationStructure; +} + +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: + { + // 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: + { + MTLPrimitiveAccelerationStructureDescriptor* primitive = [MTLPrimitiveAccelerationStructureDescriptor new]; + + NSMutableArray* geoms = [NSMutableArray arrayWithCapacity:buildInfo.geometryCount]; + for (uint32_t i = 0; i < buildInfo.geometryCount; i++) + { + // 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; + } + } + + primitive.geometryDescriptors = geoms; + descriptor = primitive; + } break; + + case VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR: + { + MTLInstanceAccelerationStructureDescriptor* instance = [MTLInstanceAccelerationStructureDescriptor new]; + // add bottom level acceleration structures + + 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{}; + + // TODO: We can't perform host builds, throw an error? + if (type == VK_ACCELERATION_STRUCTURE_BUILD_TYPE_HOST_KHR) + return vkBuildSizes; + + 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; + + return vkBuildSizes; +} + +uint64_t MVKAccelerationStructure::getMTLSize() +{ + if (!_built) { return 0; } + return _accelerationStructure.size; +} + +MVKAccelerationStructure::MVKAccelerationStructure(MVKDevice* device) : MVKVulkanAPIDeviceObject(device) +{ + MTLHeapDescriptor* heapDescriptor = [MTLHeapDescriptor new]; + heapDescriptor.storageMode = MTLStorageModePrivate; +// heapDescriptor.size = getBuildSizes().accelerationStructureSize; + _heap = [getMTLDevice() newHeapWithDescriptor:heapDescriptor]; + +// _accelerationStructure = [_heap newAccelerationStructureWithSize:getBuildSizes().accelerationStructureSize]; +// _buffer = [_heap newBufferWithLength:getBuildSizes().accelerationStructureSize options:MTLResourceOptionCPUCacheModeDefault]; +} + +void MVKAccelerationStructure::destroy() +{ + [_heap release]; + _built = false; +} diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index 0102fa1a9..03f4e2ab5 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -26,6 +26,7 @@ #include "MVKSmallVector.h" #include "MVKPixelFormats.h" #include "MVKOSExtensions.h" +#include "MVKAddressMap.h" #include "mvk_datatypes.hpp" #include #include @@ -69,6 +70,7 @@ class MVKCommandPool; class MVKCommandEncoder; class MVKCommandResourceFactory; class MVKPrivateDataSlot; +class MVKAccelerationStructure; // Not sure where to place, I'll move it there later /** The buffer index to use for vertex content. */ @@ -458,6 +460,8 @@ typedef struct MVKMTLBlitEncoder { id mtlCmdBuffer = nil; } MVKMTLBlitEncoder; +typedef std::pair MVKBufferAddressRange; + /** Represents a Vulkan logical GPU device, associated with a physical device. */ class MVKDevice : public MVKDispatchableVulkanAPIObject { @@ -526,6 +530,10 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { uint64_t* pTimestamps, uint64_t* pMaxDeviation); + /** Returns whether or not the device supports acceleration structures*/ + VkAccelerationStructureCompatibilityKHR getAccelerationStructureCompatibility(const VkAccelerationStructureVersionInfoKHR* pVersionInfo); + + #pragma mark Object lifecycle MVKBuffer* createBuffer(const VkBufferCreateInfo* pCreateInfo, @@ -597,6 +605,11 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { const VkAllocationCallbacks* pAllocator); void destroyPipelineLayout(MVKPipelineLayout* mvkPLL, const VkAllocationCallbacks* pAllocator); + + MVKAccelerationStructure* createAccelerationStructure(const VkAccelerationStructureCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator); + void destroyAccelerationStructure(MVKAccelerationStructure* mvkAccStruct, + const VkAllocationCallbacks* pAllocator); /** * Template function that creates count number of pipelines of type PipelineType, @@ -688,7 +701,13 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { /** Removes the specified timeline semaphore. */ void removeTimelineSemaphore(MVKTimelineSemaphore* sem4, uint64_t value); - + + /** Adds the specified acceleration structure to the address map, so it can be referenced else where*/ + MVKAccelerationStructure* addAccelerationStructure(MVKAccelerationStructure* accStruct); + + /** Removes the specified accelerations from the address map */ + void removeAccelerationStructure(MVKAccelerationStructure* accStruct); + /** Applies the specified global memory barrier to all resource issued by this device. */ void applyMemoryBarrier(MVKPipelineBarrier& barrier, MVKCommandEncoder* cmdEncoder, @@ -746,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 @@ -913,6 +938,9 @@ class MVKDevice : public MVKDispatchableVulkanAPIObject { MVKSmallVector, kMVKQueueFamilyCount> _queuesByQueueFamilyIndex; MVKSmallVector _resources; MVKSmallVector _gpuAddressableBuffers; + MVKAddressMap* _gpuBufferAddressMap; + uint64_t _nextValidAccStructureAddress = 0; + std::unordered_map _gpuAccStructAddressMap; MVKSmallVector _privateDataSlots; MVKSmallVector _privateDataSlotsAvailability; MVKSmallVector _awaitingSemaphores; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index c17740010..cdb875f60 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -33,6 +33,7 @@ #include "MVKFoundation.h" #include "MVKCodec.h" #include "MVKStrings.h" +#include "MVKAccelerationStructure.h" #include #import "CAMetalLayer+MoltenVK.h" @@ -3781,6 +3782,31 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope *pMaxDeviation = cpuEnd - cpuStart; } +MVKBuffer* MVKDevice::getBufferAtAddress(uint64_t address) +{ + void* value = nullptr; + _gpuBufferAddressMap->getValue(address, value); + return (MVKBuffer*)value; +} + +MVKAccelerationStructure* MVKDevice::getAccelerationStructureAtAddress(uint64_t address) +{ + std::unordered_map::iterator accStructIt = _gpuAccStructAddressMap.find(address); + if(accStructIt == _gpuAccStructAddressMap.end()) { return nullptr; } + + return accStructIt->second; +} + +VkAccelerationStructureCompatibilityKHR MVKDevice::getAccelerationStructureCompatibility(const VkAccelerationStructureVersionInfoKHR* pVersionInfo) +{ + if(_enabledAccelerationStructureFeatures.accelerationStructure) + { + return VK_ACCELERATION_STRUCTURE_COMPATIBILITY_COMPATIBLE_KHR; + } + + return VK_ACCELERATION_STRUCTURE_COMPATIBILITY_INCOMPATIBLE_KHR; +} + #pragma mark Object lifecycle uint32_t MVKDevice::getVulkanMemoryTypeIndex(MTLStorageMode mtlStorageMode) { @@ -4035,6 +4061,18 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope if (mvkPLL) { mvkPLL->destroy(); } } +MVKAccelerationStructure* MVKDevice::createAccelerationStructure(const VkAccelerationStructureCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator) { + return addAccelerationStructure(new MVKAccelerationStructure(this)); +} + +void MVKDevice::destroyAccelerationStructure(MVKAccelerationStructure* mvkAccStruct, + const VkAllocationCallbacks* pAllocator) { + if(!mvkAccStruct) { return; } + removeAccelerationStructure(mvkAccStruct); + mvkAccStruct->destroy(); +} + template VkResult MVKDevice::createPipelines(VkPipelineCache pipelineCache, uint32_t count, @@ -4263,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; } @@ -4274,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; } @@ -4328,6 +4376,30 @@ static uint32_t mvkGetEntryProperty(io_registry_entry_t entry, CFStringRef prope mvkRemoveFirstOccurance(_awaitingTimelineSem4s, make_pair(sem4, value)); } +MVKAccelerationStructure* MVKDevice::addAccelerationStructure(MVKAccelerationStructure* accStruct) { + std::pair accStructMemoryPair = std::make_pair(_nextValidAccStructureAddress, accStruct); + _gpuAccStructAddressMap.insert(accStructMemoryPair); + accStruct->setDeviceAddress(_nextValidAccStructureAddress); + _nextValidAccStructureAddress += accStruct->getMTLSize(); + return accStruct; +} + +void MVKDevice::removeAccelerationStructure(MVKAccelerationStructure* accStruct) { + std::unordered_map::iterator accStructIt = _gpuAccStructAddressMap.find(accStruct->getDeviceAddress()); + uint64_t addressOffset = accStructIt->second->getMTLSize(); + _gpuAccStructAddressMap.erase(accStructIt); + + // This can lead to fragmentation over time, so I'll just push all keys after this back + // This, however is also another performance issue + for(auto it = accStructIt; it != _gpuAccStructAddressMap.end(); it++) + { + auto extractedAccStruct = _gpuAccStructAddressMap.extract(it->first); + extractedAccStruct.key() = it->first - addressOffset; + _gpuAccStructAddressMap.insert(std::move(extractedAccStruct)); + _gpuAccStructAddressMap.erase(it->first); + } +} + void MVKDevice::applyMemoryBarrier(MVKPipelineBarrier& barrier, MVKCommandEncoder* cmdEncoder, MVKCommandUse cmdUse) { @@ -4795,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", @@ -5121,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/GPUObjects/MVKDeviceFeatureStructs.def b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def index a8e0ca690..67c1d5193 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def @@ -62,6 +62,7 @@ MVK_DEVICE_FEATURE(TimelineSemaphore, TIMELINE_SEMAPHORE, MVK_DEVICE_FEATURE(UniformBufferStandardLayout, UNIFORM_BUFFER_STANDARD_LAYOUT, 1) MVK_DEVICE_FEATURE(VariablePointer, VARIABLE_POINTER, 2) MVK_DEVICE_FEATURE(VulkanMemoryModel, VULKAN_MEMORY_MODEL, 3) +MVK_DEVICE_FEATURE_EXTN(AccelerationStructure, ACCELERATION_STRUCTURE, KHR, 5) MVK_DEVICE_FEATURE_EXTN(FragmentShaderBarycentric, FRAGMENT_SHADER_BARYCENTRIC, KHR, 1) MVK_DEVICE_FEATURE_EXTN(PortabilitySubset, PORTABILITY_SUBSET, KHR, 15) MVK_DEVICE_FEATURE_EXTN(4444Formats, 4444_FORMATS, EXT, 2) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index 7ec63a9f7..a1364331c 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -712,6 +712,17 @@ // Device extension functions. ADD_DVC_EXT_ENTRY_POINT(vkGetCalibratedTimestampsKHR, KHR_CALIBRATED_TIMESTAMPS); ADD_DVC_EXT_ENTRY_POINT(vkGetPhysicalDeviceCalibrateableTimeDomainsKHR, KHR_CALIBRATED_TIMESTAMPS); + ADD_DVC_EXT_ENTRY_POINT(vkCreateAccelerationStructureKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkDestroyAccelerationStructureKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkGetAccelerationStructureDeviceAddressKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkGetAccelerationStructureBuildSizesKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkGetDeviceAccelerationStructureCompatibilityKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCmdBuildAccelerationStructuresKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCmdCopyAccelerationStructureKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCmdCopyAccelerationStructureToMemoryKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCmdCopyMemoryToAccelerationStructureKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCmdCopyMemoryToAccelerationStructureKHR, KHR_ACCELERATION_STRUCTURE); + ADD_DVC_EXT_ENTRY_POINT(vkCreateDeferredOperationKHR, KHR_DEFERRED_HOST_OPERATIONS); ADD_DVC_EXT_ENTRY_POINT(vkDeferredOperationJoinKHR, KHR_DEFERRED_HOST_OPERATIONS); ADD_DVC_EXT_ENTRY_POINT(vkDestroyDeferredOperationKHR, KHR_DEFERRED_HOST_OPERATIONS); diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def index bbe08cd7b..014fe2a85 100644 --- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def @@ -43,6 +43,7 @@ MVK_EXTENSION(KHR_16bit_storage, KHR_16BIT_STORAGE, DEVICE, 10.11, 8.0, 1.0) MVK_EXTENSION(KHR_8bit_storage, KHR_8BIT_STORAGE, DEVICE, 10.11, 8.0, 1.0) +MVK_EXTENSION(KHR_acceleration_structure, KHR_ACCELERATION_STRUCTURE, DEVICE, 11.0, 14.0, 1.0) MVK_EXTENSION(KHR_bind_memory2, KHR_BIND_MEMORY_2, DEVICE, 10.11, 8.0, 1.0) MVK_EXTENSION(KHR_buffer_device_address, KHR_BUFFER_DEVICE_ADDRESS, DEVICE, 13.0, 16.0, 1.0) MVK_EXTENSION(KHR_calibrated_timestamps, KHR_CALIBRATED_TIMESTAMPS, DEVICE, 10.15, 14.0, 1.0) 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/MVKFoundation.h b/MoltenVK/MoltenVK/Utility/MVKFoundation.h index af7bf71c7..0d1c23cf8 100644 --- a/MoltenVK/MoltenVK/Utility/MVKFoundation.h +++ b/MoltenVK/MoltenVK/Utility/MVKFoundation.h @@ -97,7 +97,12 @@ typedef enum : uint8_t { kMVKCommandUseDrawIndirectConvertBuffers, /**< vkCmdDrawIndirect* convert indirect buffers. */ kMVKCommandUseCopyQueryPoolResults, /**< vkCmdCopyQueryPoolResults. */ kMVKCommandUseAccumOcclusionQuery, /**< Any command terminating a Metal render pass with active visibility buffer. */ - kMVKCommandUseRecordGPUCounterSample /**< Any command triggering the recording of a GPU counter sample. */ + kMVKCommandUseRecordGPUCounterSample, /**< Any command triggering the recording of a GPU counter sample. */ + kMVKCommandUseBuildAccelerationStructure, /**< vkCmdBuiildAccelerationStructure - Builds an acceleration structure */ + kMVKCommandUseCopyAccelerationStructure, /**< vkCmdCopyAccelerationStructure- Copies an acceleration structure to another acceleration structure*/ + kMVKCommandUseCopyAccelerationStructureToMemory,/**< vkCmdCopyAccelerationStructureToMemory - Copies and serializes an acceleration structure to a buffer*/ + kMVKCommandUseCopyMemoryToAccelerationStructure,/**< vkCmdCopyMemoryToAccelerationStructure - Copies and deserializes an acceleration structure from a buffer*/ + kMVKCommandUseWriteAccelerationStructuresProperties, } MVKCommandUse; /** Represents a given stage of a graphics pipeline. */ diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index e28fe7f4c..675f95d98 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -41,6 +41,7 @@ #include "MVKSurface.h" #include "MVKFoundation.h" #include "MVKOSExtensions.h" +#include "MVKAccelerationStructure.h" // I'll reposition this if needed #include @@ -2865,6 +2866,127 @@ MVK_PUBLIC_VULKAN_SYMBOL VkResult vkSetPrivateData( } +#pragma mark - +#pragma mark VK_KHR_acceleration_structure extension + +MVK_PUBLIC_VULKAN_SYMBOL VkResult vkCreateAccelerationStructureKHR( + VkDevice device, + const VkAccelerationStructureCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator, + VkAccelerationStructureKHR* pAccelerationStructure) { + + MVKTraceVulkanCallStart(); + MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); + MVKAccelerationStructure* mvkAccelerationStructure = mvkDev->createAccelerationStructure(pCreateInfo, pAllocator); + *pAccelerationStructure = (VkAccelerationStructureKHR)mvkAccelerationStructure; + VkResult rslt = VK_SUCCESS; + MVKTraceVulkanCallEnd(); + + return rslt; +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkDestroyAccelerationStructureKHR( + VkDevice device, + VkAccelerationStructureKHR accelerationStructure, + const VkAllocationCallbacks* pAllocator) { + + MVKTraceVulkanCallStart(); + MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); + MVKAccelerationStructure* mvkAccelerationStructure = (MVKAccelerationStructure*)accelerationStructure; + mvkDev->destroyAccelerationStructure(mvkAccelerationStructure, pAllocator); + MVKTraceVulkanCallEnd(); +} + + +MVK_PUBLIC_VULKAN_SYMBOL VkDeviceAddress vkGetAccelerationStructureDeviceAddressKHR( + VkDevice device, + const VkAccelerationStructureDeviceAddressInfoKHR* pInfo) { + + MVKTraceVulkanCallStart(); + MVKAccelerationStructure* mvkAccelerationStructure = (MVKAccelerationStructure*)pInfo->accelerationStructure; + uint64_t result = mvkAccelerationStructure->getDeviceAddress(); + MVKTraceVulkanCallEnd(); + + return (VkDeviceAddress)result; +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkGetAccelerationStructureBuildSizesKHR( + VkDevice device, + VkAccelerationStructureBuildTypeKHR buildType, + const VkAccelerationStructureBuildGeometryInfoKHR* pBuildInfo, + const uint32_t* pMaxPrimitiveCounts, + VkAccelerationStructureBuildSizesInfoKHR* pSizeInfo) { + + MVKTraceVulkanCallStart(); + MVKDevice* mvkDev = (MVKDevice*)device; + VkAccelerationStructureBuildSizesInfoKHR buildSizes = MVKAccelerationStructure::getBuildSizes(mvkDev, buildType, pBuildInfo, pMaxPrimitiveCounts); + pSizeInfo = &buildSizes; + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkGetDeviceAccelerationStructureCompatibilityKHR( + VkDevice device, + const VkAccelerationStructureVersionInfoKHR* pVersionInfo, + VkAccelerationStructureCompatibilityKHR* pCompatibility) { + + MVKTraceVulkanCallStart(); + MVKDevice* mvkDev = (MVKDevice*)device; + *pCompatibility = mvkDev->getAccelerationStructureCompatibility(pVersionInfo); + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkCmdBuildAccelerationStructuresKHR( + VkCommandBuffer commandBuffer, + uint32_t infoCount, + const VkAccelerationStructureBuildGeometryInfoKHR* pInfos, + const VkAccelerationStructureBuildRangeInfoKHR* const* ppBuildRangeInfos) { + + MVKTraceVulkanCallStart(); + MVKAddCmd(BuildAccelerationStructure, commandBuffer, infoCount, pInfos, ppBuildRangeInfos); + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkCmdCopyAccelerationStructureKHR( + VkCommandBuffer commandBuffer, + const VkCopyAccelerationStructureInfoKHR* pInfo) { + + MVKTraceVulkanCallStart(); + MVKAddCmd(CopyAccelerationStructure, commandBuffer, pInfo->src, pInfo->dst, pInfo->mode); + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkCmdCopyAccelerationStructureToMemoryKHR( + VkCommandBuffer commandBuffer, + const VkCopyAccelerationStructureToMemoryInfoKHR* pInfo) { + + MVKTraceVulkanCallStart(); + MVKAddCmd(CopyAccelerationStructureToMemory, commandBuffer, pInfo->src, pInfo->dst.deviceAddress, pInfo->mode); + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkCmdCopyMemoryToAccelerationStructureKHR( + VkCommandBuffer commandBuffer, + const VkCopyMemoryToAccelerationStructureInfoKHR* pInfo) { + + MVKTraceVulkanCallStart(); + MVKAddCmd(CopyMemoryToAccelerationStructure, commandBuffer, pInfo->src.deviceAddress, pInfo->dst, pInfo->mode); + MVKTraceVulkanCallEnd(); +} + +MVK_PUBLIC_VULKAN_SYMBOL void vkCmdWriteAccelerationStructuresPropertiesKHR( + VkCommandBuffer commandBuffer, + uint32_t accelerationStructureCount, + const VkAccelerationStructureKHR* pAccelerationStructures, + VkQueryType queryType, + VkQueryPool queryPool, + uint32_t firstQuery) { + + MVKTraceVulkanCallStart(); + MVKAddCmd(WriteAccelerationStructuresProperties, commandBuffer, accelerationStructureCount, pAccelerationStructures, queryType, queryPool, firstQuery); + MVKTraceVulkanCallEnd(); +} + + #pragma mark - #pragma mark VK_KHR_bind_memory2 extension