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
KhronosGroup · Jul 7, 2023 · 898e09d · 898e09d
1 parent 0cf8c17
commit 898e09d
Show file tree

Hide file tree

Showing 7 changed files with 123 additions and 6 deletions.
diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/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 */,

diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.h b/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;
+};
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm b/MoltenVK/MoltenVK/Commands/MVKCmdAccelerationStructure.mm
@@ -0,0 +1,32 @@
+/*
+ * 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);
+}
+
+MVKCommandTypePool<MVKCommand>* MVKCmdBuildAccelerationStructure::getTypePool(MVKCommandPool* cmdPool) {}
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/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;
 };
 

diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/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
 }
 

diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.h b/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
 };
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKAccelerationStructure.mm b/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()
+{
+
+}