Support resolving attachments with formats that Metal does not natively resolve.

Metal does not support resolving all formats that support MSAA, whereas Vulkan
assumes any MSAA format can be resolved. We fix that by running an optional
post-renderpass compute shader that resolves such textures by simply taking the
first sample as the resolved sample. This works to fix all failing CTS tests,
because such formats are all integer formats, and Vulkan allows an arbitrary
single sample value to be selected.

If we need to resolve, but the Metal format doesn't support it,
cause the Metal renderpass to store the MSAA attachment results.

MVKRenderSubpass don't establish Metal resolve attachment textures if format
is not natively resolvable, and encode Metal renderpass store actions accordingly.
MVKCommandEncodingPool add MTLComputePipelineStates to run simple resolve
compute shaders on attachments that cannot be resolved in Metal renderpass.
Add MVKRenderSubpass::resolveUnresolvableAttachments() and call from
MVKCommandEncoder::endMetalRenderEncoding(), before subpass index is updated.
Rename MVKCommandEncodingPool::getClearStateIndex() to
getRenderpassLoadStoreStateIndex() and remove MVK_MACOS restriction on
clearing shaders to allow compatibility with resolve shader handling.
MVKRenderPassAttachment remove validation of whether a format can be resolved.
MVKPixelFormats::getMTLTextureUsage() add read and write usage as appropriate
to allow compute shader to run to resolve formats not natively resolvable.
MVKPixelFormats remove obsolete unit test code.
MVKImageView clean up access functions and obsolete constructor
use of MVKConfiguration.

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.

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);

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";

MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h

@@ -109,10 +109,11 @@ public:
 	/** Returns a MTLComputePipelineState for filling a buffer. */
 	id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
 
-#if MVK_MACOS
 	/** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */
 	id<MTLComputePipelineState> getCmdClearColorImageMTLComputePipelineState(MVKFormatType type);
-#endif
+
+	/** Returns a MTLComputePipelineState for resolving an image. */
+	id<MTLComputePipelineState> getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type);
 
 	/** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
 	id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
@@ -161,9 +162,8 @@ protected:
     id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState = nil;
     id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
 	id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
-#if MVK_MACOS
 	id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil};
-#endif
+	id<MTLComputePipelineState> _mtlResolveColorImageComputePipelineState[3] = {nil, nil, nil};
 	id<MTLComputePipelineState> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
 	id<MTLComputePipelineState> _mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[2] = {nil, nil};
 	id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil};

MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm

@@ -109,8 +109,7 @@ id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdFillBufferMTLComputePi
 	MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState(_commandPool));
 }
 
-#if MVK_MACOS
+static inline uint32_t getRenderpassLoadStoreStateIndex(MVKFormatType type) {
-static inline uint32_t getClearStateIndex(MVKFormatType type) {
 	switch (type) {
 		case kMVKFormatColorHalf:
 		case kMVKFormatColorFloat:
@@ -129,9 +128,12 @@ static inline uint32_t getClearStateIndex(MVKFormatType type) {
 }
 
 id<MTLComputePipelineState> 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<MTLComputePipelineState> MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type) {
+	MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool));
 }
-#endif
 
 id<MTLComputePipelineState> 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];

MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h

@@ -133,6 +133,24 @@ kernel void cmdClearColorImage2DInt(texture2d<int, access::write> dst [[ texture
     dst.write(clearValue, pos);                                                                                 \n\
 }                                                                                                               \n\
                                                                                                                 \n\
+kernel void cmdResolveColorImage2DFloat(texture2d<float, access::write> dst [[ texture(0) ]],                   \n\
+                                       texture2d_ms<float, access::read> 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<uint, access::write> dst [[ texture(0) ]],                     \n\
+                                       texture2d_ms<uint, access::read> 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<int, access::write> dst [[ texture(0) ]],                       \n\
+                                      texture2d_ms<int, access::read> 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\

MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h

@@ -424,11 +424,13 @@ public:
 	/** Returns a new MTLComputePipelineState for filling a buffer. */
 	id<MTLComputePipelineState> newCmdFillBufferMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner);
 
-#if MVK_MACOS
 	/** Returns a new MTLComputePipelineState for clearing an image. */
 	id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
 																			 MVKVulkanAPIDeviceObject* owner);
-#endif
+
+	/** Returns a new MTLComputePipelineState for resolving an image. */
+	id<MTLComputePipelineState> newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type,
+																			   MVKVulkanAPIDeviceObject* owner);
 
 	/** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */
 	id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,

MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm

@@ -495,7 +495,6 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdFillBufferMTLComput
 	return newMTLComputePipelineState("cmdFillBuffer", owner);
 }
 
-#if MVK_MACOS
 id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
 					MVKVulkanAPIDeviceObject* owner) {
 	const char* funcName;
@@ -515,13 +514,36 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLC
 			funcName = "cmdClearColorImage2DUInt";
 			break;
 		default:
-			owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED,
+			owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for clearing with a compute shader.", type);
-							   "Format type %u is not supported for clearing with a compute shader.", type);
+			return nil;
+	}
+	return newMTLComputePipelineState(funcName, owner);
+}
+
+id<MTLComputePipelineState> 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<MTLComputePipelineState> MVKCommandResourceFactory::newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
 																												  MVKVulkanAPIDeviceObject* owner) {

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 specified mipmap level. */
-	 * Returns the 3D extent of this image at the base mipmap level.
+	VkExtent3D getExtent3D(uint8_t planeIndex = 0, uint32_t mipLevel = 0);
-	 * 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 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<MTLTexture> getMTLTexture(uint8_t planeIndex);
+	virtual id<MTLTexture> getMTLTexture(uint8_t planeIndex = 0);
 
 	/** Returns a Metal texture that interprets the pixels in the specified format. */
 	id<MTLTexture> 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<MTLTexture> getMTLTexture(uint8_t planeIndex) { return _planes[planeIndex]->getMTLTexture(); }
+	id<MTLTexture> 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,
+	MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCreateInfo);
-				 const VkImageViewCreateInfo* pCreateInfo,
-				 const MVKConfiguration* pAltMVKConfig = nullptr);
 
 	~MVKImageView();
 

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,
+MVKImageView::MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCreateInfo) : MVKVulkanAPIDeviceObject(device) {
-						   const VkImageViewCreateInfo* pCreateInfo,
-						   const MVKConfiguration* pAltMVKConfig) : MVKVulkanAPIDeviceObject(device) {
 	_image = (MVKImage*)pCreateInfo->image;
 	// Transfer commands don't use image views.
 	_usage = _image->_usage;

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<typename T>
-	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];

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<typename T>
-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<MTLDevice> 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

MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h

@@ -130,12 +130,13 @@ public:
 							const MVKArrayRef<MVKImageView*>& 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<MVKImageView*>& 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:
@@ -188,6 +189,7 @@ public:
                                                    bool isRenderingEntireAttachment,
 												   bool isMemorylessAttachment,
 												   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();

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,11 +201,15 @@ 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);
+			bool canResolveFormat = true;
 			if (hasResolveAttachment) {
-				attachments[rslvRPAttIdx]->populateMTLRenderPassAttachmentDescriptorResolve(mtlColorAttDesc);
+				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.
@@ -216,6 +221,7 @@ void MVKRenderSubpass::populateMTLRenderPassDescriptor(MTLRenderPassDescriptor*
 							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,
+                                                                       isRenderingEntireAttachment, isMemorylessAttachment,
-																	   isMemorylessAttachment,
+                                                                       hasResolveAttachment, canResolveFormat,
-                                                                       hasResolveAttachment, false,
+																	   false, loadOverride)) {
-                                                                       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, hasDepthResolveAttachment, canResolveFormat, 0, false, storeOverride);
-        _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasStencilResolveAttachment, 0, true, 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<MVKImageView*>& 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<MTLComputePipelineState> mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType);
+				id<MTLComputeCommandEncoder> 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,
@@ -612,6 +652,7 @@ bool MVKRenderPassAttachment::populateMTLRenderPassAttachmentDescriptor(MTLRende
                                                                         bool isRenderingEntireAttachment,
 																		bool isMemorylessAttachment,
 																		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);
 }
@@ -645,10 +686,11 @@ void MVKRenderPassAttachment::encodeStoreAction(MVKCommandEncoder* cmdEncoder,
                                                 bool isRenderingEntireAttachment,
 												bool isMemorylessAttachment,
 												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));
 			}
 		}

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);
+MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOpInObj(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat, MVKBaseObject* mvkObj);
-#define mvkMTLStoreActionFromVkAttachmentStoreOp(vkStoreOp, hasResolveAttachment) mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, this)
+#define mvkMTLStoreActionFromVkAttachmentStoreOp(vkStoreOp, hasResolveAttachment, canResolveFormat) mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, canResolveFormat, this)
 
 MTLMultisampleDepthResolveFilter mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBitsInObj(VkResolveModeFlagBits vkResolveMode, MVKBaseObject* mvkObj);
 #define mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBits(vkResolveMode) mvkMTLMultisampleDepthResolveFilterFromVkResolveModeFlagBitsInObj(vkResolveMode, this)

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) {
+MVK_PUBLIC_SYMBOL MTLStoreAction mvkMTLStoreActionFromVkAttachmentStoreOp(VkAttachmentStoreOp vkStoreOp, bool hasResolveAttachment, bool canResolveFormat) {
-	return mvkMTLStoreActionFromVkAttachmentStoreOpInObj(vkStoreOp, hasResolveAttachment, nullptr);
+	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_STORE:		return hasResolveAttachment && canResolveFormat ? MTLStoreActionStoreAndMultisampleResolve : MTLStoreActionStore;
-		case VK_ATTACHMENT_STORE_OP_DONT_CARE:	return hasResolveAttachment ? MTLStoreActionMultisampleResolve : MTLStoreActionDontCare;
+		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);