Merge pull request #219 from billhollings/master

MoltenVK 1.0.18

Docs/Whats_New.md

@@ -13,6 +13,19 @@ For best results, use a Markdown reader.*
 
 
 
+MoltenVK 1.0.18
+---------------
+
+Released 2018/08/15
+
+- vkCmdFullBuffer() fills buffer using compute shader.
+- Fix API for updating MVKDeviceConfiguration::synchronousQueueSubmits.
+- vkGetPhysicalDeviceFormatProperties() return VK_FORMAT_FEATURE_VERTEX_BUFFER_BIT 
+ if supported, even if other format properties are not.
+- Support Metal GPU capture scopes.
+- Update to latest SPIRV-Cross, glslang & SPIRV-Tools.
+
+
 MoltenVK 1.0.17
 ---------------
 

ExternalRevisions/SPIRV-Cross_repo_revision

@@ -1 +1 @@
-162eee632599cd077972ee1d88b341eedbcfdf27
+973e2e6e42f0414a304c84a19353cc4719a2bf9f

ExternalRevisions/glslang_repo_revision

@@ -1 +1 @@
-e99a26810f65314183163c07664a40e05647c15f
+ec5c11931bcfc1f438ae4aec46e1783a42cec7e6

MoltenVK/MoltenVK.xcodeproj/project.pbxproj

@@ -13,6 +13,10 @@
  A90C8DEB1F45354D009CB32C /* MVKCommandEncodingPool.h in Headers */ = {isa = PBXBuildFile; fileRef = A90C8DE81F45354D009CB32C /* MVKCommandEncodingPool.h */; };
  A90C8DEC1F45354D009CB32C /* MVKCommandEncodingPool.mm in Sources */ = {isa = PBXBuildFile; fileRef = A90C8DE91F45354D009CB32C /* MVKCommandEncodingPool.mm */; };
  A90C8DED1F45354D009CB32C /* MVKCommandEncodingPool.mm in Sources */ = {isa = PBXBuildFile; fileRef = A90C8DE91F45354D009CB32C /* MVKCommandEncodingPool.mm */; };
+ A93E832F2121C5D4001FEBD4 /* MVKGPUCapture.h in Headers */ = {isa = PBXBuildFile; fileRef = A93E832E2121C5D3001FEBD4 /* MVKGPUCapture.h */; };
+ A93E83302121C5D4001FEBD4 /* MVKGPUCapture.h in Headers */ = {isa = PBXBuildFile; fileRef = A93E832E2121C5D3001FEBD4 /* MVKGPUCapture.h */; };
+ A93E83352121F0C8001FEBD4 /* MVKGPUCapture.mm in Sources */ = {isa = PBXBuildFile; fileRef = A93E83342121F0C8001FEBD4 /* MVKGPUCapture.mm */; };
+ A93E83362121F0C8001FEBD4 /* MVKGPUCapture.mm in Sources */ = {isa = PBXBuildFile; fileRef = A93E83342121F0C8001FEBD4 /* MVKGPUCapture.mm */; };
  A948BB7F1E51642700DE59F2 /* mvk_vulkan.h in Headers */ = {isa = PBXBuildFile; fileRef = A948BB7E1E51642700DE59F2 /* mvk_vulkan.h */; settings = {ATTRIBUTES = (Public, ); }; };
  A948BB801E51642700DE59F2 /* mvk_vulkan.h in Headers */ = {isa = PBXBuildFile; fileRef = A948BB7E1E51642700DE59F2 /* mvk_vulkan.h */; settings = {ATTRIBUTES = (Public, ); }; };
  A94FB7B01C7DFB4800632CA3 /* mvk_datatypes.h in Headers */ = {isa = PBXBuildFile; fileRef = A94FB7671C7DFB4800632CA3 /* mvk_datatypes.h */; settings = {ATTRIBUTES = (Public, ); }; };
@@ -256,6 +260,8 @@
  A9096E5D1F81E16300DFBEA6 /* MVKCmdDispatch.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCmdDispatch.mm; sourceTree = "<group>"; };
  A90C8DE81F45354D009CB32C /* MVKCommandEncodingPool.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCommandEncodingPool.h; sourceTree = "<group>"; };
  A90C8DE91F45354D009CB32C /* MVKCommandEncodingPool.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCommandEncodingPool.mm; sourceTree = "<group>"; };
+ A93E832E2121C5D3001FEBD4 /* MVKGPUCapture.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKGPUCapture.h; sourceTree = "<group>"; };
+ A93E83342121F0C8001FEBD4 /* MVKGPUCapture.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKGPUCapture.mm; sourceTree = "<group>"; };
  A948BB7E1E51642700DE59F2 /* mvk_vulkan.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = mvk_vulkan.h; sourceTree = "<group>"; };
  A94FB7671C7DFB4800632CA3 /* mvk_datatypes.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = mvk_datatypes.h; sourceTree = "<group>"; };
  A94FB7691C7DFB4800632CA3 /* vk_mvk_moltenvk.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = vk_mvk_moltenvk.h; sourceTree = "<group>"; };
@@ -541,6 +547,8 @@
  A9E53DCD2100B197002781DD /* MTLSamplerDescriptor+MoltenVK.m */,
  A9E53DD02100B197002781DD /* MTLTextureDescriptor+MoltenVK.h */,
  A9E53DD52100B197002781DD /* MTLTextureDescriptor+MoltenVK.m */,
+ A93E832E2121C5D3001FEBD4 /* MVKGPUCapture.h */,
+ A93E83342121F0C8001FEBD4 /* MVKGPUCapture.mm */,
  A9E53DCE2100B197002781DD /* MVKOSExtensions.h */,
  A9E53DCF2100B197002781DD /* MVKOSExtensions.mm */,
  A9E53DD22100B197002781DD /* NSString+MoltenVK.h */,
@@ -619,6 +627,7 @@
  A9C96DD01DDC20C20053187F /* MVKMTLBufferAllocation.h in Headers */,
  A98149571FB6A3F7005F00B4 /* MVKObjectPool.h in Headers */,
  A94FB8141C7DFB4800632CA3 /* MVKSwapchain.h in Headers */,
+ A93E832F2121C5D4001FEBD4 /* MVKGPUCapture.h in Headers */,
  A94FB7DC1C7DFB4800632CA3 /* MVKBuffer.h in Headers */,
  A9F042A41FB4CF83009FCCB8 /* MVKCommonEnvironment.h in Headers */,
  A981495D1FB6A3F7005F00B4 /* MVKWatermark.h in Headers */,
@@ -677,6 +686,7 @@
  A9C96DD11DDC20C20053187F /* MVKMTLBufferAllocation.h in Headers */,
  A98149581FB6A3F7005F00B4 /* MVKObjectPool.h in Headers */,
  A94FB8151C7DFB4800632CA3 /* MVKSwapchain.h in Headers */,
+ A93E83302121C5D4001FEBD4 /* MVKGPUCapture.h in Headers */,
  A94FB7DD1C7DFB4800632CA3 /* MVKBuffer.h in Headers */,
  A9F042A51FB4CF83009FCCB8 /* MVKCommonEnvironment.h in Headers */,
  A981495E1FB6A3F7005F00B4 /* MVKWatermark.h in Headers */,
@@ -891,6 +901,7 @@
  A94FB8301C7DFB4800632CA3 /* vk_mvk_moltenvk.mm in Sources */,
  A94FB8161C7DFB4800632CA3 /* MVKSwapchain.mm in Sources */,
  A95B7D6B1D3EE486003183D3 /* MVKCommandEncoderState.mm in Sources */,
+ A93E83352121F0C8001FEBD4 /* MVKGPUCapture.mm in Sources */,
  A94FB7CE1C7DFB4800632CA3 /* MVKCommand.mm in Sources */,
  A94FB80E1C7DFB4800632CA3 /* MVKShaderModule.mm in Sources */,
  A94FB81A1C7DFB4800632CA3 /* MVKSync.mm in Sources */,
@@ -941,6 +952,7 @@
  A94FB8311C7DFB4800632CA3 /* vk_mvk_moltenvk.mm in Sources */,
  A94FB8171C7DFB4800632CA3 /* MVKSwapchain.mm in Sources */,
  A95B7D6C1D3EE486003183D3 /* MVKCommandEncoderState.mm in Sources */,
+ A93E83362121F0C8001FEBD4 /* MVKGPUCapture.mm in Sources */,
  A94FB7CF1C7DFB4800632CA3 /* MVKCommand.mm in Sources */,
  A94FB80F1C7DFB4800632CA3 /* MVKShaderModule.mm in Sources */,
  A94FB81B1C7DFB4800632CA3 /* MVKSync.mm in Sources */,

MoltenVK/MoltenVK/API/mvk_datatypes.h

@@ -165,17 +165,18 @@ size_t mvkMTLPixelFormatBytesPerLayer(MTLPixelFormat mtlFormat, size_t bytesPerR
 /** 
  * Returns the default properties for the specified Vulkan format.
  *
- * Not all MTLPixelFormats returned by this function are supported by all GPU's,
- * and, as a result, MoltenVK may return a different value from the 
- * vkGetPhysicalDeviceFormatProperties() function than is returned here.
+ * Not all MTLPixelFormats returned by this function are supported by all GPU's, and, as a
+ * result, MoltenVK may return a different value from the vkGetPhysicalDeviceFormatProperties()
+ * function than is returned here. Use the vkGetPhysicalDeviceFormatProperties() function to
+ * return the properties for a particular GPU.
  *
- * Not all macOS GPU's support the MTLPixelFormatDepth24Unorm_Stencil8 
- * (VK_FORMAT_D24_UNORM_S8_UINT) pixel format. On an macOS device that has more 
- * than one GPU, one of the GPU's may support that format, while another may not.
- * Use the vkGetPhysicalDeviceFormatProperties() function to return the properties
- * for a particular GPU.
+ * Setting assumeGPUSupportsDefault to true allows the default format properties to be returned.
+ * The assumeGPUSupportsDefault flag can be set to false if it is already known that the format
+ * is not supported by a particular GPU for images, in which case all of the returned properties
+ * will be disabled, except possibly VK_FORMAT_FEATURE_VERTEX_BUFFER_BIT, which may be supported
+ * for the format even without image support.
  */
-VkFormatProperties mvkVkFormatProperties(VkFormat vkFormat);
+VkFormatProperties mvkVkFormatProperties(VkFormat vkFormat, bool assumeGPUSupportsDefault = true);
 
 /** Returns the name of the specified Vulkan format. */
 const char* mvkVkFormatName(VkFormat vkFormat);

MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h

@@ -48,7 +48,7 @@ extern "C" {
  */
 #define MVK_VERSION_MAJOR 1
 #define MVK_VERSION_MINOR 0
-#define MVK_VERSION_PATCH 17
+#define MVK_VERSION_PATCH 18
 
 #define MVK_MAKE_VERSION(major, minor, patch) (((major) * 10000) + ((minor) * 100) + (patch))
 #define MVK_VERSION MVK_MAKE_VERSION(MVK_VERSION_MAJOR, MVK_VERSION_MINOR, MVK_VERSION_PATCH)

MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm

@@ -541,6 +541,13 @@
 #pragma mark -
 #pragma mark MVKCmdCopyBuffer
 
+// Matches shader struct.
+typedef struct {
+ uint32_t srcOffset;
+ uint32_t dstOffset;
+ uint32_t size;
+} MVKCmdCopyBufferInfo;
+
 void MVKCmdCopyBuffer::setContent(VkBuffer srcBuffer,
  VkBuffer destBuffer,
  uint32_t regionCount,
@@ -570,17 +577,20 @@
  cpyRgn.dstOffset % buffAlign != 0 ||
  cpyRgn.size % buffAlign != 0);
  if (useComputeCopy) {
- MVKAssert(cpyRgn.srcOffset <= UINT32_MAX || cpyRgn.dstOffset <= UINT32_MAX || cpyRgn.size <= UINT32_MAX,
- "Compute buffer copy region offsets and size must fit into a 32-bit unsigned integer.");
+ MVKAssert(mvkFits<uint32_t>(cpyRgn.srcOffset) && mvkFits<uint32_t>(cpyRgn.dstOffset) && mvkFits<uint32_t>(cpyRgn.size),
+ "Byte-aligned buffer copy region offsets and size must each fit into a 32-bit unsigned integer.");
+
+ MVKCmdCopyBufferInfo copyInfo;
+ copyInfo.srcOffset = (uint32_t)cpyRgn.srcOffset;
+ copyInfo.dstOffset = (uint32_t)cpyRgn.dstOffset;
+ copyInfo.size = (uint32_t)cpyRgn.size;
 
  id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseCopyBuffer);
  [mtlComputeEnc pushDebugGroup: @"vkCmdCopyBuffer"];
- id<MTLComputePipelineState> pipelineState = cmdEncoder->getCommandEncodingPool()->getCopyBufferBytesComputePipelineState();
- [mtlComputeEnc setComputePipelineState:pipelineState];
- [mtlComputeEnc setBuffer:srcMTLBuff offset:srcMTLBuffOffset atIndex:0];
- [mtlComputeEnc setBuffer:dstMTLBuff offset:dstMTLBuffOffset atIndex:1];
- uint32_t copyInfo[3] = { (uint32_t)cpyRgn.srcOffset, (uint32_t)cpyRgn.dstOffset, (uint32_t)cpyRgn.size };
- [mtlComputeEnc setBytes:copyInfo length:sizeof(copyInfo) atIndex:2];
+ [mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdCopyBufferBytesMTLComputePipelineState()];
+ [mtlComputeEnc setBuffer:srcMTLBuff offset: srcMTLBuffOffset atIndex: 0];
+ [mtlComputeEnc setBuffer:dstMTLBuff offset: dstMTLBuffOffset atIndex: 1];
+ [mtlComputeEnc setBytes: &copyInfo length: sizeof(copyInfo) atIndex: 2];
  [mtlComputeEnc dispatchThreads: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)];
  [mtlComputeEnc popDebugGroup];
  } else {
@@ -939,7 +949,7 @@
  ? _image->getLayerCount()
  : (layerStart + layerCnt));
 
- // Iterate across mipmap levels and layers, and render to clear each
+ // Iterate across mipmap levels and layers, and perform and empty render to clear each
  for (uint32_t mipLvl = mipLvlStart; mipLvl < mipLvlEnd; mipLvl++) {
  mtlRPCADesc.level = mipLvl;
  mtlRPDADesc.level = mipLvl;
@@ -963,6 +973,13 @@
 #pragma mark -
 #pragma mark MVKCmdFillBuffer
 
+// Matches shader struct
+typedef struct {
+ uint32_t dstOffset;
+ uint32_t size;
+ uint32_t data;
+} MVKCmdFillBufferInfo;
+
 void MVKCmdFillBuffer::setContent(VkBuffer dstBuffer,
  VkDeviceSize dstOffset,
  VkDeviceSize size,
@@ -974,18 +991,26 @@
 }
 
 void MVKCmdFillBuffer::encode(MVKCommandEncoder* cmdEncoder) {
-
- id<MTLBlitCommandEncoder> mtlBlitEnc = cmdEncoder->getMTLBlitEncoder(kMVKCommandUseFillBuffer);
-
  id<MTLBuffer> dstMTLBuff = _dstBuffer->getMTLBuffer();
- NSUInteger dstMTLBuffOffset = _dstBuffer->getMTLBufferOffset() + _dstOffset;
- VkDeviceSize byteCnt = (_size == VK_WHOLE_SIZE) ? (_dstBuffer->getByteCount() - _dstOffset) : _size;
-
- // Metal only supports filling with a single byte value, so each byte in the
- // buffer will be filled with the lower 8 bits of the Vulkan 32-bit data value.
- [mtlBlitEnc fillBuffer: dstMTLBuff
- range: NSMakeRange(dstMTLBuffOffset, byteCnt)
- value: (uint8_t)_dataValue];
+ VkDeviceSize dstMTLBuffOffset = _dstBuffer->getMTLBufferOffset();
+ VkDeviceSize byteCnt = (_size == VK_WHOLE_SIZE) ? (_dstBuffer->getByteCount() - (dstMTLBuffOffset + _dstOffset)) : _size;
+ VkDeviceSize wordCnt = byteCnt >> 2;
+
+ MVKAssert(mvkFits<uint32_t>(_dstOffset) && mvkFits<uint32_t>(wordCnt),
+ "Buffer fill offset and size must each fit into a 32-bit unsigned integer.");
+
+ MVKCmdFillBufferInfo fillInfo;
+ fillInfo.dstOffset = (uint32_t)_dstOffset;
+ fillInfo.size = (uint32_t)wordCnt;
+ fillInfo.data = _dataValue;
+
+ id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseCopyBuffer);
+ [mtlComputeEnc pushDebugGroup: @"vkCmdFillBuffer"];
+ [mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdFillBufferMTLComputePipelineState()];
+ [mtlComputeEnc setBuffer: dstMTLBuff offset: dstMTLBuffOffset atIndex: 0];
+ [mtlComputeEnc setBytes: &fillInfo length: sizeof(fillInfo) atIndex: 1];
+ [mtlComputeEnc dispatchThreads: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)];
+ [mtlComputeEnc popDebugGroup];
 }
 
 

MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h

@@ -91,11 +91,11 @@ class MVKCommandEncodingPool : public MVKBaseDeviceObject {
  */
  MVKImage* getTransferMVKImage(MVKImageDescriptorData& imgData);
 
- /**
-  * Returns an MTLComputePipelineState dedicated to copying bytes between two buffers
- * with unaligned copy regions.
-  */
- id<MTLComputePipelineState> getCopyBufferBytesComputePipelineState();
+ /** Returns a MTLComputePipelineState for copying between two buffers with byte-aligned copy regions. */
+ id<MTLComputePipelineState> getCmdCopyBufferBytesMTLComputePipelineState();
+
+ /** Returns a MTLComputePipelineState for filling a buffer. */
+ id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
 
 #pragma mark Construction
 
@@ -113,12 +113,13 @@ class MVKCommandEncodingPool : public MVKBaseDeviceObject {
  std::unordered_map<MVKImageDescriptorData, MVKImage*> _transferImages;
  MVKDeviceMemory* _transferImageMemory;
  MVKMTLBufferAllocator _mtlBufferAllocator;
- id<MTLSamplerState> _cmdBlitImageLinearMTLSamplerState;
- id<MTLSamplerState> _cmdBlitImageNearestMTLSamplerState;
- id<MTLDepthStencilState> _cmdClearDepthOnlyDepthStencilState;
- id<MTLDepthStencilState> _cmdClearStencilOnlyDepthStencilState;
- id<MTLDepthStencilState> _cmdClearDepthAndStencilDepthStencilState;
- id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState;
- id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState;
+ id<MTLSamplerState> _cmdBlitImageLinearMTLSamplerState = nil;
+ id<MTLSamplerState> _cmdBlitImageNearestMTLSamplerState = nil;
+ id<MTLDepthStencilState> _cmdClearDepthOnlyDepthStencilState = nil;
+ id<MTLDepthStencilState> _cmdClearStencilOnlyDepthStencilState = nil;
+ id<MTLDepthStencilState> _cmdClearDepthAndStencilDepthStencilState = nil;
+ id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState = nil;
+ id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
+ id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
 };
 

MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm

@@ -111,26 +111,25 @@
  return mvkImg;
 }
 
-id<MTLComputePipelineState> MVKCommandEncodingPool::getCopyBufferBytesComputePipelineState() {
+id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferBytesMTLComputePipelineState() {
  if (_mtlCopyBufferBytesComputePipelineState == nil) {
- _mtlCopyBufferBytesComputePipelineState = _device->getCommandResourceFactory()->newCopyBytesMTLComputePipelineState();
+ _mtlCopyBufferBytesComputePipelineState = _device->getCommandResourceFactory()->newCmdCopyBufferBytesMTLComputePipelineState();
  }
  return _mtlCopyBufferBytesComputePipelineState;
 }
 
+id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdFillBufferMTLComputePipelineState() {
+ if (_mtlFillBufferComputePipelineState == nil) {
+ _mtlFillBufferComputePipelineState = _device->getCommandResourceFactory()->newCmdFillBufferMTLComputePipelineState();
+ }
+ return _mtlFillBufferComputePipelineState;
+}
+
 #pragma mark Construction
 
 MVKCommandEncodingPool::MVKCommandEncodingPool(MVKDevice* device) : MVKBaseDeviceObject(device),
  _mtlBufferAllocator(device, device->_pMetalFeatures->maxMTLBufferSize) {
 
- _cmdBlitImageLinearMTLSamplerState = nil;
- _cmdBlitImageNearestMTLSamplerState = nil;
- _cmdClearDepthAndStencilDepthStencilState = nil;
- _cmdClearDepthOnlyDepthStencilState = nil;
- _cmdClearStencilOnlyDepthStencilState = nil;
- _cmdClearDefaultDepthStencilState = nil;
- _mtlCopyBufferBytesComputePipelineState = nil;
-
  initTextureDeviceMemory();
 }