Starting to Implement Acceleration Structure Commands

This commit adds:

* A .h and .mm file for Acceleration Structure commands
* An acceleration structure command encoder into `MVKCommandBuffer`
* An actual acceleration structure handle
* And some other items that are not complete, or need to removed

MoltenVK/MoltenVK.xcodeproj/project.pbxproj

@@ -7,6 +7,9 @@
 	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 */; };
@@ -432,6 +435,8 @@
 /* End PBXContainerItemProxy section */
 
 /* Begin PBXFileReference section */
+		014702702A5855F70040D02D /* MVKCmdAccelerationStructure.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKCmdAccelerationStructure.h; sourceTree = "<group>"; };
+		014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCmdAccelerationStructure.mm; sourceTree = "<group>"; };
 		019795132A5304D600C6CAD0 /* MVKAccelerationStructure.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKAccelerationStructure.h; sourceTree = "<group>"; };
 		0197951A2A56F8AF00C6CAD0 /* MVKAccelerationStructure.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKAccelerationStructure.mm; sourceTree = "<group>"; };
 		2FEA0ABA24902F9F00EEF3AD /* libMoltenVK.a */ = {isa = PBXFileReference; explicitFileType = archive.ar; includeInIndex = 0; path = libMoltenVK.a; sourceTree = BUILT_PRODUCTS_DIR; };
@@ -599,6 +604,8 @@
 		A94FB76B1C7DFB4800632CA3 /* Commands */ = {
 			isa = PBXGroup;
 			children = (
+				014702702A5855F70040D02D /* MVKCmdAccelerationStructure.h */,
+				014702722A5857600040D02D /* MVKCmdAccelerationStructure.mm */,
 				A99C90EC229455B200A061DA /* MVKCmdDebug.h */,
 				A99C90ED229455B300A061DA /* MVKCmdDebug.mm */,
 				A9096E5C1F81E16300DFBEA6 /* MVKCmdDispatch.h */,
@@ -1406,6 +1413,7 @@
 				2FEA0AB024902F9F00EEF3AD /* MVKFramebuffer.mm in Sources */,
 				2FEA0AB124902F9F00EEF3AD /* MVKMTLBufferAllocation.mm in Sources */,
 				2FEA0AB224902F9F00EEF3AD /* CAMetalLayer+MoltenVK.m in Sources */,
+				014702742A5857600040D02D /* MVKCmdAccelerationStructure.mm in Sources */,
 				2FEA0AB324902F9F00EEF3AD /* MVKCmdDispatch.mm in Sources */,
 				2FEA0AB424902F9F00EEF3AD /* MVKCmdDebug.mm in Sources */,
 			);
@@ -1444,6 +1452,7 @@
 				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 */,
@@ -1505,6 +1514,7 @@
 				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 */,

MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h

@@ -0,0 +1,32 @@
+/*
+ * 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 "MVKCommand.h"
+
+#pragma mark -
+#pragma mark MVKCmdBuildAccelerationStructure
+
+class MVKCmdBuildAccelerationStructure : public MVKCommand {
+
+public:
+    void encode(MVKCommandEncoder* cmdEncoder) override;
+protected:
+    MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
+};

MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm

@@ -0,0 +1,30 @@
+/*
+ * 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"
+
+#import <Metal/Metal.h>
+#import <Metal/MTLAccelerationStructure.h>
+#import <Metal/MTLAccelerationStructureTypes.h>
+
+void MVKCmdBuildAccelerationStructure::encode(MVKCommandEncoder* cmdEncoder) {
+    cmdEncoder->getMTLAccelerationStructureEncoder(kMVKCommandUseNone);
+}

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h

@@ -352,6 +352,8 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
      * the current encoder before beginning BLIT encoding.
 	 */
 	id<MTLBlitCommandEncoder> getMTLBlitEncoder(MVKCommandUse cmdUse);
+
+    id<MTLAccelerationStructureCommandEncoder> getMTLAccelerationStructureEncoder(MVKCommandUse cmdUse); // Write proper comment above
 
 	/**
 	 * Returns the current Metal encoder, which may be any of the Metal render,
@@ -515,6 +517,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
 	MVKSmallVector<MVKSmallVector<MTLSamplePosition>> _subpassSamplePositions;
 	id<MTLComputeCommandEncoder> _mtlComputeEncoder;
 	id<MTLBlitCommandEncoder> _mtlBlitEncoder;
+    id<MTLAccelerationStructureCommandEncoder> _mtlAccelerationStructureEncoder;
 	id<MTLFence> _stageCountersMTLFence;
 	MVKPushConstantsCommandEncoderState _vertexPushConstants;
 	MVKPushConstantsCommandEncoderState _tessCtlPushConstants;
@@ -529,6 +532,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
     uint32_t _flushCount;
 	MVKCommandUse _mtlComputeEncoderUse;
 	MVKCommandUse _mtlBlitEncoderUse;
+    MVKCommandUse _mtlAccelerationStructureUse;
 	bool _isRenderingEntireAttachment;
 };
 

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm

@@ -846,6 +846,10 @@
 	if (_mtlBlitEncoder && _cmdBuffer->_hasStageCounterTimestampCommand) { [_mtlBlitEncoder updateFence: getStageCountersMTLFence()]; }
 	endMetalEncoding(_mtlBlitEncoder);
     _mtlBlitEncoderUse = kMVKCommandUseNone;
+
+    if (_mtlAccelerationStructureEncoder && _cmdBuffer->_hasStageCounterTimestampCommand) { [_mtlAccelerationStructureEncoder updateFence: getStageCountersMTLFence()]; }
+    endMetalEncoding(_mtlAccelerationStructureEncoder);
+    _mtlAccelerationStructureUse = kMVKCommandUseNone;
 
 	encodeTimestampStageCounterSamples();
 }
@@ -883,10 +887,24 @@
 	return _mtlBlitEncoder;
 }
 
+id<MTLAccelerationStructureCommandEncoder> MVKCommandEncoder::getMTLAccelerationStructureEncoder(MVKCommandUse cmdUse) {
+    if ( !_mtlAccelerationStructureEncoder ) {
+        endCurrentMetalEncoding();
+        _mtlAccelerationStructureEncoder = [_mtlCmdBuffer accelerationStructureCommandEncoder];
+        retainIfImmediatelyEncoding(_mtlAccelerationStructureEncoder);
+    }
+    if (_mtlAccelerationStructureUse != cmdUse) {
+        _mtlAccelerationStructureUse = cmdUse;
+        setLabelIfNotNil(_mtlAccelerationStructureEncoder, mvkMTLBlitCommandEncoderLabel(cmdUse));
+    }
+    return _mtlAccelerationStructureEncoder;
+}
+
 id<MTLCommandEncoder> MVKCommandEncoder::getMTLEncoder(){
 	if (_mtlRenderEncoder) { return _mtlRenderEncoder; }
 	if (_mtlComputeEncoder) { return _mtlComputeEncoder; }
 	if (_mtlBlitEncoder) { return _mtlBlitEncoder; }
+    if (_mtlAccelerationStructureEncoder) { return _mtlAccelerationStructureEncoder; }
 	return nil;
 }
 
@@ -1149,6 +1167,8 @@
 			_mtlComputeEncoderUse = kMVKCommandUseNone;
             _mtlBlitEncoder = nil;
             _mtlBlitEncoderUse = kMVKCommandUseNone;
+            _mtlAccelerationStructureEncoder = nil;
+            _mtlAccelerationStructureUse = kMVKCommandUseNone;
 			_pEncodingContext = nullptr;
 			_stageCountersMTLFence = nil;
 			_flushCount = 0;
@@ -1158,6 +1178,7 @@
 	[_mtlRenderEncoder release];
 	[_mtlComputeEncoder release];
 	[_mtlBlitEncoder release];
+    [_mtlAccelerationStructureEncoder release];
 	// _stageCountersMTLFence is released after Metal command buffer completion
 }
 

MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h

@@ -25,8 +25,8 @@
  vkCmdCopyAccelerationStructureToMemoryKHR
  vkCmdCopyMemoryToAccelerationStructureKHR
  vkCmdWriteAccelerationStructuresPropertiesKHR
- vkCreateAccelerationStructureKHR
- vkDestroyAccelerationStructureKHR
+ vkCreateAccelerationStructureKHR - DONE
+ vkDestroyAccelerationStructureKHR - DONE
  vkGetAccelerationStructureBuildSizesKHR - DONE
  vkGetAccelerationStructureDeviceAddressKHR
  vkGetDeviceAccelerationStructureCompatibilityKHR
@@ -37,6 +37,9 @@
 
 #include "MVKVulkanAPIObject.h"
 
+#import <Metal/MTLAccelerationStructure.h>
+#import <Metal/MTLAccelerationStructureTypes.h>
+
 #pragma mark MVKAccelerationStructure
 
 class MVKAccelerationStructure : public MVKVulkanAPIDeviceObject {
@@ -51,9 +54,17 @@ class MVKAccelerationStructure : public MVKVulkanAPIDeviceObject {
     /** Gets the required build sizes for acceleration structure and scratch buffer*/
     static VkAccelerationStructureBuildSizesInfoKHR getBuildSizes();
 
+    /** Gets the device address of the acceleration structure*/
+    void getDeviceAddress();
 
+    /** Builds the acceleration structure as a device command*/
+    void build();
 #pragma mark Construction
     MVKAccelerationStructure(MVKDevice* device) : MVKVulkanAPIDeviceObject(device) {}
 protected:
     void propagateDebugName() override {}
+
+    #if MVK_XCODE_12
+        id<MTLAccelerationStructure> _accelerationStructure;
+    #endif
 };

MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm

@@ -26,10 +26,17 @@
 {
     VkAccelerationStructureBuildSizesInfoKHR vkBuildSizes{};
 
-    MTLAccelerationStructureSizes mtlBuildSizes;
-    vkBuildSizes.accelerationStructureSize = mtlBuildSizes.accelerationStructureSize;
-    vkBuildSizes.buildScratchSize = mtlBuildSizes.buildScratchBufferSize;
-    vkBuildSizes.updateScratchSize = mtlBuildSizes.refitScratchBufferSize;
+    #if MVK_XCODE_12
+        MTLAccelerationStructureSizes mtlBuildSizes;
+        vkBuildSizes.accelerationStructureSize = mtlBuildSizes.accelerationStructureSize;
+        vkBuildSizes.buildScratchSize = mtlBuildSizes.buildScratchBufferSize;
+        vkBuildSizes.updateScratchSize = mtlBuildSizes.refitScratchBufferSize;
+    #endif
 
     return vkBuildSizes;
 }
+
+void MVKAccelerationStructure::getDeviceAddress()
+{
+
+}