diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md index ea3fd4ff..899f6217 100644 --- a/Docs/Whats_New.md +++ b/Docs/Whats_New.md @@ -19,6 +19,7 @@ MoltenVK 1.1.5 Released TBD - Vulkan timestamp query pools use Metal GPU counters when available. +- Support resolving attachments with formats that Metal does not natively resolve. - Fix issue where swapchain images were acquired out of order under heavy load. - Fix issue with `vkCmdBlitImage()` from compressed textures. - Fix incorrect translation of clear color values on Apple Silicon. diff --git a/MoltenVK/MoltenVK/API/mvk_datatypes.h b/MoltenVK/MoltenVK/API/mvk_datatypes.h index 186a4597..275a048d 100644 --- a/MoltenVK/MoltenVK/API/mvk_datatypes.h +++ b/MoltenVK/MoltenVK/API/mvk_datatypes.h @@ -358,7 +358,7 @@ MTLTriangleFillMode mvkMTLTriangleFillModeFromVkPolygonMode(VkPolygonMode vkFill MTLLoadAction mvkMTLLoadActionFromVkAttachmentLoadOp(VkAttachmentLoadOp vkLoadOp); /** Returns the Metal MTLStoreAction corresponding to the specified Vulkan VkAttachmentStoreOp. */ -MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOp(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment); +MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOp(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat = true); /** Returns the Metal MTLMultisampleDepthResolveFilter corresponding to the specified Vulkan VkResolveModeFlagBits. */ MTLMultisampleDepthResolveFilter mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBits(VkResolveModeFlagBits vkResolveMode); diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm index d3a351c1..411e17cb 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm @@ -310,10 +310,12 @@ void MVKCommandEncoder::beginNextSubpass(MVKCommand* subpassCmd, VkSubpassConten } // Sets the current render subpass to the subpass with the specified index. +// End current Metal renderpass before udpating subpass index. void MVKCommandEncoder::setSubpass(MVKCommand* subpassCmd, VkSubpassContents subpassContents, uint32_t subpassIndex) { encodeStoreActions(); + endMetalRenderEncoding(); _lastMultiviewPassCmd = subpassCmd; _subpassContents = subpassContents; @@ -579,12 +581,13 @@ void MVKCommandEncoder::endRenderpass() { } void MVKCommandEncoder::endMetalRenderEncoding() { -// MVKLogDebugIf(_mtlRenderEncoder, "Render subpass end MTLRenderCommandEncoder."); if (_mtlRenderEncoder == nil) { return; } [_mtlRenderEncoder endEncoding]; _mtlRenderEncoder = nil; // not retained + getSubpass()->resolveUnresolvableAttachments(this, _attachments.contents()); + _graphicsPipelineState.endMetalRenderPass(); _graphicsResourcesState.endMetalRenderPass(); _viewportState.endMetalRenderPass(); @@ -932,6 +935,7 @@ NSString* mvkMTLComputeCommandEncoderLabel(MVKCommandUse cmdUse) { case kMVKCommandUseCopyImageToBuffer: return @"vkCmdCopyImageToBuffer ComputeEncoder"; case kMVKCommandUseFillBuffer: return @"vkCmdFillBuffer ComputeEncoder"; case kMVKCommandUseClearColorImage: return @"vkCmdClearColorImage ComputeEncoder"; + case kMVKCommandUseResolveImage: return @"Resolve Subpass Attachment ComputeEncoder"; case kMVKCommandUseTessellationVertexTessCtl: return @"vkCmdDraw (vertex and tess control stages) ComputeEncoder"; case kMVKCommandUseMultiviewInstanceCountAdjust: return @"vkCmdDraw (multiview instance count adjustment) ComputeEncoder"; case kMVKCommandUseCopyQueryPoolResults: return @"vkCmdCopyQueryPoolResults ComputeEncoder"; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h index 6318f640..a57ee88b 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h @@ -109,10 +109,11 @@ public: /** Returns a MTLComputePipelineState for filling a buffer. */ id getCmdFillBufferMTLComputePipelineState(); -#if MVK_MACOS /** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */ id getCmdClearColorImageMTLComputePipelineState(MVKFormatType type); -#endif + + /** Returns a MTLComputePipelineState for resolving an image. */ + id getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type); /** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */ id getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff); @@ -161,9 +162,8 @@ protected: id _cmdClearDefaultDepthStencilState = nil; id _mtlCopyBufferBytesComputePipelineState = nil; id _mtlFillBufferComputePipelineState = nil; -#if MVK_MACOS id _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil}; -#endif + id _mtlResolveColorImageComputePipelineState[3] = {nil, nil, nil}; id _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil}; id _mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[2] = {nil, nil}; id _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil}; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm index 78f8bbc5..00fd9ee8 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm @@ -109,8 +109,7 @@ id MVKCommandEncodingPool::getCmdFillBufferMTLComputePi MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState(_commandPool)); } -#if MVK_MACOS -static inline uint32_t getClearStateIndex(MVKFormatType type) { +static inline uint32_t getRenderpassLoadStoreStateIndex(MVKFormatType type) { switch (type) { case kMVKFormatColorHalf: case kMVKFormatColorFloat: @@ -129,9 +128,12 @@ static inline uint32_t getClearStateIndex(MVKFormatType type) { } id MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) { - MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getClearStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool)); + MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool)); +} + +id MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type) { + MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool)); } -#endif id MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) { MVK_ENC_REZ_ACCESS(_mtlCopyBufferToImage3DDecompressComputePipelineState[needsTempBuff ? 1 : 0], newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(needsTempBuff, _commandPool)); @@ -215,14 +217,19 @@ void MVKCommandEncodingPool::destroyMetalResources() { [_mtlFillBufferComputePipelineState release]; _mtlFillBufferComputePipelineState = nil; -#if MVK_MACOS [_mtlClearColorImageComputePipelineState[0] release]; [_mtlClearColorImageComputePipelineState[1] release]; [_mtlClearColorImageComputePipelineState[2] release]; _mtlClearColorImageComputePipelineState[0] = nil; _mtlClearColorImageComputePipelineState[1] = nil; _mtlClearColorImageComputePipelineState[2] = nil; -#endif + + [_mtlResolveColorImageComputePipelineState[0] release]; + [_mtlResolveColorImageComputePipelineState[1] release]; + [_mtlResolveColorImageComputePipelineState[2] release]; + _mtlResolveColorImageComputePipelineState[0] = nil; + _mtlResolveColorImageComputePipelineState[1] = nil; + _mtlResolveColorImageComputePipelineState[2] = nil; [_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release]; [_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release]; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h index 7b715f34..b26b35b7 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h @@ -133,6 +133,24 @@ kernel void cmdClearColorImage2DInt(texture2d dst [[ texture dst.write(clearValue, pos); \n\ } \n\ \n\ +kernel void cmdResolveColorImage2DFloat(texture2d dst [[ texture(0) ]], \n\ + texture2d_ms src [[ texture(1) ]], \n\ + uint2 pos [[thread_position_in_grid]]) { \n\ + dst.write(src.read(pos, 0), pos); \n\ +} \n\ + \n\ +kernel void cmdResolveColorImage2DUInt(texture2d dst [[ texture(0) ]], \n\ + texture2d_ms src [[ texture(1) ]], \n\ + uint2 pos [[thread_position_in_grid]]) { \n\ + dst.write(src.read(pos, 0), pos); \n\ +} \n\ + \n\ +kernel void cmdResolveColorImage2DInt(texture2d dst [[ texture(0) ]], \n\ + texture2d_ms src [[ texture(1) ]], \n\ + uint2 pos [[thread_position_in_grid]]) { \n\ + dst.write(src.read(pos, 0), pos); \n\ +} \n\ + \n\ typedef struct { \n\ uint32_t srcRowStride; \n\ uint32_t srcRowStrideHigh; \n\ diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h index e036c45a..7dd04b4b 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h @@ -424,11 +424,13 @@ public: /** Returns a new MTLComputePipelineState for filling a buffer. */ id newCmdFillBufferMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner); -#if MVK_MACOS /** Returns a new MTLComputePipelineState for clearing an image. */ id newCmdClearColorImageMTLComputePipelineState(MVKFormatType type, MVKVulkanAPIDeviceObject* owner); -#endif + + /** Returns a new MTLComputePipelineState for resolving an image. */ + id newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, + MVKVulkanAPIDeviceObject* owner); /** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */ id newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf, diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm index 2a3b02a4..e37bd06f 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm @@ -495,7 +495,6 @@ id MVKCommandResourceFactory::newCmdFillBufferMTLComput return newMTLComputePipelineState("cmdFillBuffer", owner); } -#if MVK_MACOS id MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type, MVKVulkanAPIDeviceObject* owner) { const char* funcName; @@ -515,13 +514,36 @@ id MVKCommandResourceFactory::newCmdClearColorImageMTLC funcName = "cmdClearColorImage2DUInt"; break; default: - owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, - "Format type %u is not supported for clearing with a compute shader.", type); + owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for clearing with a compute shader.", type); + return nil; + } + return newMTLComputePipelineState(funcName, owner); +} + +id MVKCommandResourceFactory::newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, + MVKVulkanAPIDeviceObject* owner) { + const char* funcName; + switch (type) { + case kMVKFormatColorHalf: + case kMVKFormatColorFloat: + funcName = "cmdResolveColorImage2DFloat"; + break; + case kMVKFormatColorInt8: + case kMVKFormatColorInt16: + case kMVKFormatColorInt32: + funcName = "cmdResolveColorImage2DInt"; + break; + case kMVKFormatColorUInt8: + case kMVKFormatColorUInt16: + case kMVKFormatColorUInt32: + funcName = "cmdResolveColorImage2DUInt"; + break; + default: + owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for resolving with a compute shader.", type); return nil; } return newMTLComputePipelineState(funcName, owner); } -#endif id MVKCommandResourceFactory::newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf, MVKVulkanAPIDeviceObject* owner) { diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h index 4817b139..c6490697 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h @@ -190,29 +190,20 @@ public: /** Returns whether this image is allowed to alias another image. */ bool getIsAliasable() { return _isAliasable; } - /** - * Returns the 3D extent of this image at the base mipmap level. - * For 2D or cube images, the Z component will be 1. - */ - inline VkExtent3D getExtent3D() { return _extent; } - - /** - * Returns the 3D extent of this image at the specified mipmap level. - * For 2D or cube images, the Z component will be 1. - */ - VkExtent3D getExtent3D(uint8_t planeIndex, uint32_t mipLevel); + /** Returns the 3D extent of this image at the specified mipmap level. */ + VkExtent3D getExtent3D(uint8_t planeIndex = 0, uint32_t mipLevel = 0); /** Returns the number of mipmap levels in this image. */ - inline uint32_t getMipLevelCount() { return _mipLevels; } + uint32_t getMipLevelCount() { return _mipLevels; } /** * Returns the number of layers at each mipmap level. For an array image type, this is * the number of elements in the array. For cube image type, this is a multiple of 6. */ - inline uint32_t getLayerCount() { return _arrayLayers; } + uint32_t getLayerCount() { return _arrayLayers; } /** Returns the number of samples for each pixel of this image. */ - inline VkSampleCountFlagBits getSampleCount() { return _samples; } + VkSampleCountFlagBits getSampleCount() { return _samples; } /** * Returns the number of bytes per image row at the specified zero-based mip level. @@ -231,7 +222,7 @@ public: VkDeviceSize getBytesPerLayer(uint8_t planeIndex, uint32_t mipLevel); /** Returns the number of planes of this image view. */ - inline uint8_t getPlaneCount() { return _planes.size(); } + uint8_t getPlaneCount() { return _planes.size(); } /** Populates the specified layout for the specified sub-resource. */ VkResult getSubresourceLayout(const VkImageSubresource* pSubresource, @@ -268,7 +259,7 @@ public: #pragma mark Metal /** Returns the Metal texture underlying this image. */ - virtual id getMTLTexture(uint8_t planeIndex); + virtual id getMTLTexture(uint8_t planeIndex = 0); /** Returns a Metal texture that interprets the pixels in the specified format. */ id getMTLTexture(uint8_t planeIndex, MTLPixelFormat mtlPixFmt); @@ -307,7 +298,7 @@ public: IOSurfaceRef getIOSurface(); /** Returns the Metal pixel format of this image. */ - inline MTLPixelFormat getMTLPixelFormat(uint8_t planeIndex) { return _planes[planeIndex]->_mtlPixFmt; } + inline MTLPixelFormat getMTLPixelFormat(uint8_t planeIndex = 0) { return _planes[planeIndex]->_mtlPixFmt; } /** Returns the Metal texture type of this image. */ inline MTLTextureType getMTLTextureType() { return _mtlTextureType; } @@ -570,13 +561,16 @@ public: /** Returns the debug report object type of this object. */ VkDebugReportObjectTypeEXT getVkDebugReportObjectType() override { return VK_DEBUG_REPORT_OBJECT_TYPE_IMAGE_VIEW_EXT; } + /** Returns the 3D extent of this image at the specified mipmap level. */ + VkExtent3D getExtent3D(uint8_t planeIndex = 0, uint32_t mipLevel = 0) { return _image->getExtent3D(planeIndex, mipLevel); } + #pragma mark Metal /** Returns the Metal texture underlying this image view. */ - id getMTLTexture(uint8_t planeIndex) { return _planes[planeIndex]->getMTLTexture(); } + id getMTLTexture(uint8_t planeIndex = 0) { return _planes[planeIndex]->getMTLTexture(); } /** Returns the Metal pixel format of this image view. */ - MTLPixelFormat getMTLPixelFormat(uint8_t planeIndex) { return _planes[planeIndex]->_mtlPixFmt; } + MTLPixelFormat getMTLPixelFormat(uint8_t planeIndex = 0) { return _planes[planeIndex]->_mtlPixFmt; } /** Returns the packed component swizzle of this image view. */ uint32_t getPackedSwizzle() { return _planes[0]->getPackedSwizzle(); } @@ -602,9 +596,7 @@ public: #pragma mark Construction - MVKImageView(MVKDevice* device, - const VkImageViewCreateInfo* pCreateInfo, - const MVKConfiguration* pAltMVKConfig = nullptr); + MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCreateInfo); ~MVKImageView(); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index 1a4f78ed..e24a1957 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -846,7 +846,7 @@ MTLTextureUsage MVKImage::getMTLTextureUsage(MTLPixelFormat mtlPixFmt) { needsReinterpretation = needsReinterpretation || !pixFmts->compatibleAsLinearOrSRGB(mtlPixFmt, viewFmt); } - MTLTextureUsage mtlUsage = pixFmts->getMTLTextureUsage(_usage, mtlPixFmt, _isLinear, needsReinterpretation, _hasExtendedUsage); + MTLTextureUsage mtlUsage = pixFmts->getMTLTextureUsage(_usage, mtlPixFmt, _samples, _isLinear, needsReinterpretation, _hasExtendedUsage); // Metal before 3.0 doesn't support 3D compressed textures, so we'll // decompress the texture ourselves, and we need to be able to write to it. @@ -1731,9 +1731,7 @@ void MVKImageView::populateMTLRenderPassAttachmentDescriptorResolve(MTLRenderPas #pragma mark Construction -MVKImageView::MVKImageView(MVKDevice* device, - const VkImageViewCreateInfo* pCreateInfo, - const MVKConfiguration* pAltMVKConfig) : MVKVulkanAPIDeviceObject(device) { +MVKImageView::MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCreateInfo) : MVKVulkanAPIDeviceObject(device) { _image = (MVKImage*)pCreateInfo->image; // Transfer commands don't use image views. _usage = _image->_usage; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.h b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.h index eb1304c8..fe55655e 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.h @@ -369,6 +369,7 @@ public: */ MTLTextureUsage getMTLTextureUsage(VkImageUsageFlags vkImageUsageFlags, MTLPixelFormat mtlFormat, + VkSampleCountFlagBits samples = VK_SAMPLE_COUNT_1_BIT, bool isLinear = false, bool needsReinterpretation = true, bool isExtended = false); @@ -422,11 +423,6 @@ protected: MTLVertexFormat mtlVtxFmt, MVKMTLFmtCaps mtlFmtCaps); - template - void testFmt(const T v1, const T v2, const char* fmtName, const char* funcName); - void testProps(const VkFormatProperties p1, const VkFormatProperties p2, const char* fmtName); - void test(); - MVKPhysicalDevice* _physicalDevice; MVKVkFormatDesc _vkFormatDescriptions[_vkFormatCount]; MVKMTLFormatDesc _mtlPixelFormatDescriptions[_mtlPixelFormatCount]; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm index fb305ba2..3192c8ca 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm @@ -630,6 +630,7 @@ VkImageUsageFlags MVKPixelFormats::getVkImageUsageFlags(MTLTextureUsage mtlUsage MTLTextureUsage MVKPixelFormats::getMTLTextureUsage(VkImageUsageFlags vkImageUsageFlags, MTLPixelFormat mtlFormat, + VkSampleCountFlagBits samples, bool isLinear, bool needsReinterpretation, bool isExtended) { @@ -680,6 +681,14 @@ MTLTextureUsage MVKPixelFormats::getMTLTextureUsage(VkImageUsageFlags vkImageUsa #endif } + // Resolving an MSAA color attachment whose format Metal cannot resolve natively, may use a compute shader + // to perform theh resolve, by reading from the multisample texture and writing to the single-sample texture. + if (mvkIsAnyFlagEnabled(vkImageUsageFlags, (VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT)) && + !mvkIsAnyFlagEnabled(mtlFmtCaps, kMVKMTLFmtCapsResolve)) { + + mvkEnableFlags(mtlUsage, samples == VK_SAMPLE_COUNT_1_BIT ? MTLTextureUsageShaderWrite : MTLTextureUsageShaderRead); + } + // Create view on, but only on color formats, or combined depth-stencil formats if supported by the GPU... if ((mvkIsAnyFlagEnabled(vkImageUsageFlags, (VK_IMAGE_USAGE_TRANSFER_SRC_BIT)) || // May use temp view if transfer involves format change (needsReinterpretation && @@ -744,8 +753,6 @@ MVKPixelFormats::MVKPixelFormats(MVKPhysicalDevice* physicalDevice) : _physicalD // Build the Vulkan formats and link them to the Metal formats initVkFormatCapabilities(); buildVkFormatMaps(); - -// test(); } #define addVkFormatDescFull(VK_FMT, MTL_FMT, MTL_FMT_ALT, MTL_VTX_FMT, MTL_VTX_FMT_ALT, CSPC, CSCB, BLK_W, BLK_H, BLK_BYTE_CNT, MVK_FMT_TYPE) \ @@ -2101,103 +2108,3 @@ void MVKPixelFormats::setFormatProperties(MVKVkFormatDesc& vkDesc) { enableFormatFeatures(Vertex, Buf, getMTLVertexFormatDesc(vkDesc.mtlVertexFormat).mtlFmtCaps, vkProps.bufferFeatures); } } - - -#pragma mark - -#pragma mark Unit Testing - -template -void MVKPixelFormats::testFmt(const T v1, const T v2, const char* fmtName, const char* funcName) { - MVKAssert(mvkAreEqual(&v1,&v2), "Results not equal for format %s on test %s.", fmtName, funcName); -} - -void MVKPixelFormats::testProps(const VkFormatProperties p1, const VkFormatProperties p2, const char* fmtName) { - MVKLogErrorIf(!mvkAreEqual(&p1, &p2), - "Properties not equal for format %s. " - "\n\tgetVkFormatProperties() linear %d, optimal %d, buffer %d. " - "\n\tmvkVkFormatProperties(): linear %d, optimal %d, buffer %d" - "\n\tdifference: linear %d, optimal %d, buffer %d", fmtName, - p1.linearTilingFeatures, p1.optimalTilingFeatures, p1.bufferFeatures, - p2.linearTilingFeatures, p2.optimalTilingFeatures, p2.bufferFeatures, - std::abs((int)p2.linearTilingFeatures - (int)p1.linearTilingFeatures), - std::abs((int)p2.optimalTilingFeatures - (int)p1.optimalTilingFeatures), - std::abs((int)p2.bufferFeatures - (int)p1.bufferFeatures)); -} - -// Validate the functionality of this class against the previous format data within MoltenVK. -// This is a temporary function to confirm that converting to using this class matches existing behaviour at first. -#define testFmt(V1, V2) testFmt(V1, V2, fd.name, #V1) -#define testProps(V1, V2) testProps(V1, V2, fd.name) -void MVKPixelFormats::test() { - if ( !_physicalDevice ) { return; } // Don't test a static instance not associated with a physical device - - // If more than one GPU, only test the system default MTLDevice. - // Can release system MTLDevice immediates because we are just comparing it's address. - id sysMTLDvc = MTLCreateSystemDefaultDevice(); // temp retained - [sysMTLDvc release]; // release temp instance - if ( _physicalDevice->getMTLDevice() != sysMTLDvc ) { return; } - - MVKLogInfo("Starting testing formats"); - for (uint32_t fmtIdx = 0; fmtIdx < _vkFormatCount; fmtIdx++) { - auto& fd = _vkFormatDescriptions[fmtIdx]; - VkFormat vkFmt = fd.vkFormat; - MTLPixelFormat mtlFmt = fd.mtlPixelFormat; - - if (fd.vkFormat) { - if (fd.isSupportedOrSubstitutable()) { - MVKLogInfo("Testing %s", fd.name); - - testFmt(isSupported(vkFmt), mvkVkFormatIsSupported(vkFmt)); - testFmt(isSupported(mtlFmt), mvkMTLPixelFormatIsSupported(mtlFmt)); - testFmt(isDepthFormat(mtlFmt), mvkMTLPixelFormatIsDepthFormat(mtlFmt)); - testFmt(isStencilFormat(mtlFmt), mvkMTLPixelFormatIsStencilFormat(mtlFmt)); - testFmt(isPVRTCFormat(mtlFmt), mvkMTLPixelFormatIsPVRTCFormat(mtlFmt)); - testFmt(getFormatType(vkFmt), mvkFormatTypeFromVkFormat(vkFmt)); - testFmt(getFormatType(mtlFmt), mvkFormatTypeFromMTLPixelFormat(mtlFmt)); - testFmt(getMTLPixelFormat(vkFmt), mvkMTLPixelFormatFromVkFormat(vkFmt)); - testFmt(getVkFormat(mtlFmt), mvkVkFormatFromMTLPixelFormat(mtlFmt)); - testFmt(getBytesPerBlock(vkFmt), mvkVkFormatBytesPerBlock(vkFmt)); - testFmt(getBytesPerBlock(mtlFmt), mvkMTLPixelFormatBytesPerBlock(mtlFmt)); - testFmt(getBlockTexelSize(vkFmt), mvkVkFormatBlockTexelSize(vkFmt)); - testFmt(getBlockTexelSize(mtlFmt), mvkMTLPixelFormatBlockTexelSize(mtlFmt)); - testFmt(getBytesPerTexel(vkFmt), mvkVkFormatBytesPerTexel(vkFmt)); - testFmt(getBytesPerTexel(mtlFmt), mvkMTLPixelFormatBytesPerTexel(mtlFmt)); - testFmt(getBytesPerRow(vkFmt, 4), mvkVkFormatBytesPerRow(vkFmt, 4)); - testFmt(getBytesPerRow(mtlFmt, 4), mvkMTLPixelFormatBytesPerRow(mtlFmt, 4)); - testFmt(getBytesPerLayer(vkFmt, 256, 4), mvkVkFormatBytesPerLayer(vkFmt, 256, 4)); - testFmt(getBytesPerLayer(mtlFmt, 256, 4), mvkMTLPixelFormatBytesPerLayer(mtlFmt, 256, 4)); - testProps(getVkFormatProperties(vkFmt), mvkVkFormatProperties(vkFmt)); - testFmt(strcmp(getName(vkFmt), mvkVkFormatName(vkFmt)), 0); - testFmt(strcmp(getName(mtlFmt), mvkMTLPixelFormatName(mtlFmt)), 0); - testFmt(getMTLClearColor(VkClearValue(), vkFmt), - mvkMTLClearColorFromVkClearValue(VkClearValue(), vkFmt)); - - testFmt(getVkImageUsageFlags(MTLTextureUsageUnknown, mtlFmt), - mvkVkImageUsageFlagsFromMTLTextureUsage(MTLTextureUsageUnknown, mtlFmt)); - testFmt(getVkImageUsageFlags(MTLTextureUsageShaderRead, mtlFmt), - mvkVkImageUsageFlagsFromMTLTextureUsage(MTLTextureUsageShaderRead, mtlFmt)); - testFmt(getVkImageUsageFlags(MTLTextureUsageShaderWrite, mtlFmt), - mvkVkImageUsageFlagsFromMTLTextureUsage(MTLTextureUsageShaderWrite, mtlFmt)); - testFmt(getVkImageUsageFlags(MTLTextureUsageRenderTarget, mtlFmt), - mvkVkImageUsageFlagsFromMTLTextureUsage(MTLTextureUsageRenderTarget, mtlFmt)); - testFmt(getVkImageUsageFlags(MTLTextureUsagePixelFormatView, mtlFmt), - mvkVkImageUsageFlagsFromMTLTextureUsage(MTLTextureUsagePixelFormatView, mtlFmt)); - - VkImageUsageFlags vkUsage; - vkUsage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_SAMPLED_BIT | VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_INPUT_ATTACHMENT_BIT | VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT | VK_IMAGE_USAGE_DEPTH_STENCIL_ATTACHMENT_BIT; - testFmt(getMTLTextureUsage(vkUsage, mtlFmt), mvkMTLTextureUsageFromVkImageUsageFlags(vkUsage, mtlFmt)); - - vkUsage = VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT | VK_IMAGE_USAGE_DEPTH_STENCIL_ATTACHMENT_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT | VK_IMAGE_USAGE_STORAGE_BIT; - testFmt(getMTLTextureUsage(vkUsage, mtlFmt), mvkMTLTextureUsageFromVkImageUsageFlags(vkUsage, mtlFmt)); - - testFmt(getMTLVertexFormat(vkFmt), mvkMTLVertexFormatFromVkFormat(vkFmt)); - - } else { - MVKLogInfo("%s not supported or substitutable on this device.", fd.name); - } - } - } - MVKLogInfo("Finished testing formats.\n"); -} -#undef testFmt -#undef testProps diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h index 76f83a6a..0119d2a8 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h @@ -130,12 +130,13 @@ public: const MVKArrayRef& attachments, bool storeOverride = false); - /** Constructs an instance for the specified parent renderpass. */ + /** Resolves any resolve attachments that cannot be handled by native Metal subpass resolve behavior. */ + void resolveUnresolvableAttachments(MVKCommandEncoder* cmdEncoder, const MVKArrayRef& attachments); + MVKRenderSubpass(MVKRenderPass* renderPass, const VkSubpassDescription* pCreateInfo, const VkRenderPassInputAttachmentAspectCreateInfo* pInputAspects, uint32_t viewMask); - /** Constructs an instance for the specified parent renderpass. */ MVKRenderSubpass(MVKRenderPass* renderPass, const VkSubpassDescription2* pCreateInfo); private: @@ -187,7 +188,8 @@ public: MVKRenderSubpass* subpass, bool isRenderingEntireAttachment, bool isMemorylessAttachment, - bool hasResolveAttachment, + bool hasResolveAttachment, + bool canResolveFormat, bool isStencil, bool loadOverride = false); @@ -197,6 +199,7 @@ public: bool isRenderingEntireAttachment, bool isMemorylessAttachment, bool hasResolveAttachment, + bool canResolveFormat, uint32_t caIdx, bool isStencil, bool storeOverride = false); @@ -225,6 +228,7 @@ protected: bool isRenderingEntireAttachment, bool isMemorylessAttachment, bool hasResolveAttachment, + bool canResolveFormat, bool isStencil, bool storeOverride); void validateFormat(); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm index bced6485..870dfc77 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm @@ -19,6 +19,7 @@ #include "MVKRenderPass.h" #include "MVKFramebuffer.h" #include "MVKCommandBuffer.h" +#include "MVKCommandEncodingPool.h" #include "MVKFoundation.h" #include "mvk_datatypes.hpp" #include "MTLRenderPassDepthAttachmentDescriptor+MoltenVK.h" @@ -200,22 +201,27 @@ void MVKRenderSubpass::populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* MTLRenderPassColorAttachmentDescriptor* mtlColorAttDesc = mtlRPDesc.colorAttachments[caIdx]; // If it exists, configure the resolve attachment first, - // as it affects how the store action of the color attachment. + // as it affects the store action of the color attachment. uint32_t rslvRPAttIdx = _resolveAttachments.empty() ? VK_ATTACHMENT_UNUSED : _resolveAttachments[caIdx].attachment; bool hasResolveAttachment = (rslvRPAttIdx != VK_ATTACHMENT_UNUSED); - if (hasResolveAttachment) { - attachments[rslvRPAttIdx]->populateMTLRenderPassAttachmentDescriptorResolve(mtlColorAttDesc); + bool canResolveFormat = true; + if (hasResolveAttachment) { + MVKImageView* raImgView = attachments[rslvRPAttIdx]; + canResolveFormat = mvkAreAllFlagsEnabled(pixFmts->getCapabilities(raImgView->getMTLPixelFormat()), kMVKMTLFmtCapsResolve); + if (canResolveFormat) { + raImgView->populateMTLRenderPassAttachmentDescriptorResolve(mtlColorAttDesc); - // In a multiview render pass, we need to override the starting layer to ensure - // only the enabled views are loaded. - if (isMultiview()) { - uint32_t startView = getFirstViewIndexInMetalPass(passIdx); - if (mtlColorAttDesc.resolveTexture.textureType == MTLTextureType3D) - mtlColorAttDesc.resolveDepthPlane += startView; - else - mtlColorAttDesc.resolveSlice += startView; + // In a multiview render pass, we need to override the starting layer to ensure + // only the enabled views are loaded. + if (isMultiview()) { + uint32_t startView = getFirstViewIndexInMetalPass(passIdx); + if (mtlColorAttDesc.resolveTexture.textureType == MTLTextureType3D) + mtlColorAttDesc.resolveDepthPlane += startView; + else + mtlColorAttDesc.resolveSlice += startView; + } } - } + } // Configure the color attachment MVKRenderPassAttachment* clrMVKRPAtt = &_renderPass->_attachments[clrRPAttIdx]; @@ -225,10 +231,9 @@ void MVKRenderSubpass::populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* isMemorylessAttachment = attachments[clrRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless; #endif if (clrMVKRPAtt->populateMTLRenderPassAttachmentDescriptor(mtlColorAttDesc, this, - isRenderingEntireAttachment, - isMemorylessAttachment, - hasResolveAttachment, false, - loadOverride)) { + isRenderingEntireAttachment, isMemorylessAttachment, + hasResolveAttachment, canResolveFormat, + false, loadOverride)) { mtlColorAttDesc.clearColor = pixFmts->getMTLClearColor(clearValues[clrRPAttIdx], clrMVKRPAtt->getFormat()); } if (isMultiview()) { @@ -381,16 +386,19 @@ void MVKRenderSubpass::encodeStoreActions(MVKCommandEncoder* cmdEncoder, if (!cmdEncoder->_mtlRenderEncoder) { return; } if (!_renderPass->getDevice()->_pMetalFeatures->deferredStoreActions) { return; } + MVKPixelFormats* pixFmts = _renderPass->getPixelFormats(); uint32_t caCnt = getColorAttachmentCount(); for (uint32_t caIdx = 0; caIdx < caCnt; ++caIdx) { uint32_t clrRPAttIdx = _colorAttachments[caIdx].attachment; if (clrRPAttIdx != VK_ATTACHMENT_UNUSED) { - bool hasResolveAttachment = _resolveAttachments.empty() ? false : _resolveAttachments[caIdx].attachment != VK_ATTACHMENT_UNUSED; + uint32_t rslvRPAttIdx = _resolveAttachments.empty() ? VK_ATTACHMENT_UNUSED : _resolveAttachments[caIdx].attachment; + bool hasResolveAttachment = (rslvRPAttIdx != VK_ATTACHMENT_UNUSED); + bool canResolveFormat = hasResolveAttachment && mvkAreAllFlagsEnabled(pixFmts->getCapabilities(attachments[rslvRPAttIdx]->getMTLPixelFormat()), kMVKMTLFmtCapsResolve); bool isMemorylessAttachment = false; #if MVK_APPLE_SILICON isMemorylessAttachment = attachments[clrRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless; #endif - _renderPass->_attachments[clrRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, caIdx, false, storeOverride); + _renderPass->_attachments[clrRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, canResolveFormat, caIdx, false, storeOverride); } } uint32_t dsRPAttIdx = _depthStencilAttachment.attachment; @@ -398,12 +406,13 @@ void MVKRenderSubpass::encodeStoreActions(MVKCommandEncoder* cmdEncoder, bool hasResolveAttachment = _depthStencilResolveAttachment.attachment != VK_ATTACHMENT_UNUSED; bool hasDepthResolveAttachment = hasResolveAttachment && _depthResolveMode != VK_RESOLVE_MODE_NONE; bool hasStencilResolveAttachment = hasResolveAttachment && _stencilResolveMode != VK_RESOLVE_MODE_NONE; + bool canResolveFormat = true; bool isMemorylessAttachment = false; #if MVK_APPLE_SILICON isMemorylessAttachment = attachments[dsRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless; #endif - _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasDepthResolveAttachment, 0, false, storeOverride); - _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasStencilResolveAttachment, 0, true, storeOverride); + _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasDepthResolveAttachment, canResolveFormat, 0, false, storeOverride); + _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasStencilResolveAttachment, canResolveFormat, 0, true, storeOverride); } } @@ -489,6 +498,37 @@ MVKMTLFmtCaps MVKRenderSubpass::getRequiredFormatCapabilitiesForAttachmentAt(uin return caps; } +void MVKRenderSubpass::resolveUnresolvableAttachments(MVKCommandEncoder* cmdEncoder, const MVKArrayRef& attachments) { + MVKPixelFormats* pixFmts = cmdEncoder->getPixelFormats(); + size_t raCnt = _resolveAttachments.size(); + for (uint32_t raIdx = 0; raIdx < raCnt; raIdx++) { + auto& ra = _resolveAttachments[raIdx]; + auto& ca = _colorAttachments[raIdx]; + if (ra.attachment != VK_ATTACHMENT_UNUSED && ca.attachment != VK_ATTACHMENT_UNUSED) { + MVKImageView* raImgView = attachments[ra.attachment]; + MVKImageView* caImgView = attachments[ca.attachment]; + + if ( !mvkAreAllFlagsEnabled(pixFmts->getCapabilities(raImgView->getMTLPixelFormat()), kMVKMTLFmtCapsResolve) ) { + MVKFormatType mvkFmtType = _renderPass->getPixelFormats()->getFormatType(raImgView->getMTLPixelFormat()); + id mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType); + id mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseResolveImage); + [mtlComputeEnc setComputePipelineState: mtlRslvState]; + [mtlComputeEnc setTexture: raImgView->getMTLTexture() atIndex: 0]; + [mtlComputeEnc setTexture: caImgView->getMTLTexture() atIndex: 1]; + MTLSize gridSize = mvkMTLSizeFromVkExtent3D(raImgView->getExtent3D()); + MTLSize tgSize = MTLSizeMake(mtlRslvState.threadExecutionWidth, 1, 1); + if (cmdEncoder->getDevice()->_pMetalFeatures->nonUniformThreadgroups) { + [mtlComputeEnc dispatchThreads: gridSize threadsPerThreadgroup: tgSize]; + } else { + MTLSize tgCount = MTLSizeMake(gridSize.width / tgSize.width, gridSize.height, gridSize.depth); + if (gridSize.width % tgSize.width) { tgCount.width += 1; } + [mtlComputeEnc dispatchThreadgroups: tgCount threadsPerThreadgroup: tgSize]; + } + } + } + } +} + MVKRenderSubpass::MVKRenderSubpass(MVKRenderPass* renderPass, const VkSubpassDescription* pCreateInfo, const VkRenderPassInputAttachmentAspectCreateInfo* pInputAspects, @@ -611,7 +651,8 @@ bool MVKRenderPassAttachment::populateMTLRenderPassAttachmentDescriptor(MTLRende MVKRenderSubpass* subpass, bool isRenderingEntireAttachment, bool isMemorylessAttachment, - bool hasResolveAttachment, + bool hasResolveAttachment, + bool canResolveFormat, bool isStencil, bool loadOverride) { // Only allow clearing of entire attachment if we're actually @@ -635,7 +676,7 @@ bool MVKRenderPassAttachment::populateMTLRenderPassAttachmentDescriptor(MTLRende if ( _renderPass->getDevice()->_pMetalFeatures->deferredStoreActions ) { mtlAttDesc.storeAction = MTLStoreActionUnknown; } else { - mtlAttDesc.storeAction = getMTLStoreAction(subpass, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, isStencil, false); + mtlAttDesc.storeAction = getMTLStoreAction(subpass, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, canResolveFormat, isStencil, false); } return (mtlLA == MTLLoadActionClear); } @@ -644,11 +685,12 @@ void MVKRenderPassAttachment::encodeStoreAction(MVKCommandEncoder* cmdEncoder, MVKRenderSubpass* subpass, bool isRenderingEntireAttachment, bool isMemorylessAttachment, - bool hasResolveAttachment, + bool hasResolveAttachment, + bool canResolveFormat, uint32_t caIdx, bool isStencil, bool storeOverride) { - MTLStoreAction storeAction = getMTLStoreAction(subpass, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, isStencil, storeOverride); + MTLStoreAction storeAction = getMTLStoreAction(subpass, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, canResolveFormat, isStencil, storeOverride); MVKPixelFormats* pixFmts = _renderPass->getPixelFormats(); MTLPixelFormat mtlFmt = pixFmts->getMTLPixelFormat(_info.format); @@ -698,10 +740,11 @@ MTLStoreAction MVKRenderPassAttachment::getMTLStoreAction(MVKRenderSubpass* subp bool isRenderingEntireAttachment, bool isMemorylessAttachment, bool hasResolveAttachment, + bool canResolveFormat, bool isStencil, bool storeOverride) { // If a resolve attachment exists, this attachment must resolve once complete. - if (hasResolveAttachment && !_renderPass->getDevice()->_pMetalFeatures->combinedStoreResolveAction) { + if (hasResolveAttachment && canResolveFormat && !_renderPass->getDevice()->_pMetalFeatures->combinedStoreResolveAction) { return MTLStoreActionMultisampleResolve; } // Memoryless can't be stored. @@ -712,10 +755,10 @@ MTLStoreAction MVKRenderPassAttachment::getMTLStoreAction(MVKRenderSubpass* subp // Only allow the attachment to be discarded if we're actually // rendering to the entire attachment and we're in the last subpass. if (storeOverride || !isRenderingEntireAttachment || !isLastUseOfAttachment(subpass)) { - return hasResolveAttachment ? MTLStoreActionStoreAndMultisampleResolve : MTLStoreActionStore; + return hasResolveAttachment && canResolveFormat ? MTLStoreActionStoreAndMultisampleResolve : MTLStoreActionStore; } VkAttachmentStoreOp storeOp = isStencil ? _info.stencilStoreOp : _info.storeOp; - return mvkMTLStoreActionFromVkAttachmentStoreOp(storeOp, hasResolveAttachment); + return mvkMTLStoreActionFromVkAttachmentStoreOp(storeOp, hasResolveAttachment, canResolveFormat); } bool MVKRenderPassAttachment::shouldUseClearAttachment(MVKRenderSubpass* subpass) { @@ -765,7 +808,10 @@ void MVKRenderPassAttachment::validateFormat() { // Validate that the attachment pixel format supports the capabilities required by the subpass. // Use MTLPixelFormat to look up capabilities to permit Metal format substitution. - if ( !mvkAreAllFlagsEnabled(pixFmts->getCapabilities(pixFmts->getMTLPixelFormat(_info.format)), reqCaps) ) { + // It's okay if the format does not support the resolve capability, as this can be handled via a compute shader. + MVKMTLFmtCaps availCaps = pixFmts->getCapabilities(pixFmts->getMTLPixelFormat(_info.format)); + mvkEnableFlags(availCaps, kMVKMTLFmtCapsResolve); + if ( !mvkAreAllFlagsEnabled(availCaps, reqCaps) ) { _renderPass->setConfigurationResult(reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "vkCreateRenderPass(): Attachment format %s on this device does not support the VkFormat attachment capabilities required by the subpass at index %d.", _renderPass->getPixelFormats()->getName(_info.format), spIdx)); } } diff --git a/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.hpp b/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.hpp index 8d408f56..fb5f7cc0 100644 --- a/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.hpp +++ b/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.hpp @@ -59,8 +59,8 @@ MTLTriangleFillMode mvkMTLTriangleFillModeFromVkPolygonModeInObj(VkPolygonMode v MTLLoadAction mvkMTLLoadActionFromVkAttachmentLoadOpInObj(VkAttachmentLoadOp vkLoadOp, MVKBaseObject* mvkObj); #define mvkMTLLoadActionFromVkAttachmentLoadOp(vkLoadOp) mvkMTLLoadActionFromVkAttachmentLoadOpInObj(vkLoadOp, this) -MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOpInObj(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, MVKBaseObject* mvkObj); -#define mvkMTLStoreActionFromVkAttachmentStoreOp(vkStoreOp, hasResolveAttachment) mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, this) +MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOpInObj(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat, MVKBaseObject* mvkObj); +#define mvkMTLStoreActionFromVkAttachmentStoreOp(vkStoreOp, hasResolveAttachment, canResolveFormat) mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, canResolveFormat, this) MTLMultisampleDepthResolveFilter mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBitsInObj(VkResolveModeFlagBits vkResolveMode, MVKBaseObject* mvkObj); #define mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBits(vkResolveMode) mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBitsInObj(vkResolveMode, this) diff --git a/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.mm b/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.mm index 9953aa2e..7ce43a17 100644 --- a/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.mm +++ b/MoltenVK/MoltenVK/Vulkan/mvk_datatypes.mm @@ -494,14 +494,15 @@ MTLLoadAction mvkMTLLoadActionFromVkAttachmentLoadOpInObj(VkAttachmentLoadOp vkL } #undef mvkMTLStoreActionFromVkAttachmentStoreOp -MVK_PUBLIC_SYMBOL MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOp(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment) { - return mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, nullptr); +MVK_PUBLIC_SYMBOL MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOp(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat) { + return mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, canResolveFormat, nullptr); } -MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOpInObj(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, MVKBaseObject* mvkObj) { +// If we need to resolve, but the format doesn't support it, we must store the attachment so we can run a post-renderpass compute shader to perform the resolve. +MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOpInObj(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat, MVKBaseObject* mvkObj) { switch (vkStoreOp) { - case VK_ATTACHMENT_STORE_OP_STORE: return hasResolveAttachment ? MTLStoreActionStoreAndMultisampleResolve : MTLStoreActionStore; - case VK_ATTACHMENT_STORE_OP_DONT_CARE: return hasResolveAttachment ? MTLStoreActionMultisampleResolve : MTLStoreActionDontCare; + case VK_ATTACHMENT_STORE_OP_STORE: return hasResolveAttachment && canResolveFormat ? MTLStoreActionStoreAndMultisampleResolve : MTLStoreActionStore; + case VK_ATTACHMENT_STORE_OP_DONT_CARE: return hasResolveAttachment ? (canResolveFormat ? MTLStoreActionMultisampleResolve : MTLStoreActionStore) : MTLStoreActionDontCare; default: MVKBaseObject::reportError(mvkObj, VK_ERROR_FORMAT_NOT_SUPPORTED, "VkAttachmentStoreOp value %d is not supported.", vkStoreOp);