Avoid calling useResource on resources in argument buffers

MoltenVK/MoltenVK/API/mvk_datatypes.h

@@ -333,6 +333,15 @@ typedef enum {
 	kMVKShaderStageMax = kMVKShaderStageCount	// Public API legacy value
 } MVKShaderStage;
 
+typedef enum {
+	kMVKBarrierStageVertex = 0,
+	kMVKBarrierStageFragment,
+	kMVKBarrierStageCompute,
+	kMVKBarrierStageCopy,
+	kMVKBarrierStageNone,
+	kMVKBarrierStageCount = kMVKBarrierStageNone
+} MVKBarrierStage;
+
 /** Returns the Metal MTLColorWriteMask corresponding to the specified Vulkan VkColorComponentFlags. */
 MTLColorWriteMask mvkMTLColorWriteMaskFromVkChannelFlags(VkColorComponentFlags vkWriteFlags);
 

MoltenVK/MoltenVK/API/mvk_private_api.h

@@ -358,6 +358,7 @@ typedef struct {
 	VkBool32 needsCubeGradWorkaround;				/**< If true, sampling from cube textures with explicit gradients is broken and needs a workaround. */
 	VkBool32 nativeTextureAtomics;                  /**< If true, atomic operations on textures are supported natively. */
 	VkBool32 needsArgumentBufferEncoders;			/**< If true, Metal argument buffer encoders are needed to populate argument buffer content. */
+    VkBool32 residencySets;                         /**< If true, the device supports creating residency sets. */
 } MVKPhysicalDeviceMetalFeatures;
 
 

MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm

@@ -79,6 +79,35 @@
 	return VK_SUCCESS;
 }
 
+static uint64_t mvkPipelineStageFlagsToBarrierStages(VkPipelineStageFlags2 flags) {
+	uint64_t result = 0;
+
+	if (mvkIsAnyFlagEnabled(flags, VK_PIPELINE_STAGE_2_VERTEX_INPUT_BIT | VK_PIPELINE_STAGE_2_VERTEX_SHADER_BIT | VK_PIPELINE_STAGE_2_TESSELLATION_CONTROL_SHADER_BIT |
+							VK_PIPELINE_STAGE_2_TESSELLATION_EVALUATION_SHADER_BIT | VK_PIPELINE_STAGE_2_ALL_GRAPHICS_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT |
+							VK_PIPELINE_STAGE_2_DRAW_INDIRECT_BIT | VK_PIPELINE_STAGE_2_GEOMETRY_SHADER_BIT | VK_PIPELINE_STAGE_2_INDEX_INPUT_BIT |
+							VK_PIPELINE_STAGE_2_VERTEX_ATTRIBUTE_INPUT_BIT | VK_PIPELINE_STAGE_2_PRE_RASTERIZATION_SHADERS_BIT | VK_PIPELINE_STAGE_2_TRANSFORM_FEEDBACK_BIT_EXT |
+							VK_PIPELINE_STAGE_2_TOP_OF_PIPE_BIT | VK_PIPELINE_STAGE_2_BOTTOM_OF_PIPE_BIT))
+		result |= 1 << kMVKBarrierStageVertex;
+
+	if (mvkIsAnyFlagEnabled(flags, VK_PIPELINE_STAGE_2_FRAGMENT_SHADER_BIT | VK_PIPELINE_STAGE_2_ALL_GRAPHICS_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT |
+							VK_PIPELINE_STAGE_2_EARLY_FRAGMENT_TESTS_BIT | VK_PIPELINE_STAGE_2_LATE_FRAGMENT_TESTS_BIT | VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT |
+							VK_PIPELINE_STAGE_2_FRAGMENT_SHADING_RATE_ATTACHMENT_BIT_KHR | VK_PIPELINE_STAGE_2_FRAGMENT_DENSITY_PROCESS_BIT_EXT |
+							VK_PIPELINE_STAGE_2_TOP_OF_PIPE_BIT | VK_PIPELINE_STAGE_2_BOTTOM_OF_PIPE_BIT))
+		result |= 1 << kMVKBarrierStageFragment;
+
+	if (mvkIsAnyFlagEnabled(flags, VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT | VK_PIPELINE_STAGE_2_TOP_OF_PIPE_BIT |
+							VK_PIPELINE_STAGE_2_BOTTOM_OF_PIPE_BIT))
+		result |= 1 << kMVKBarrierStageCompute;
+
+	if (mvkIsAnyFlagEnabled(flags, VK_PIPELINE_STAGE_2_BLIT_BIT | VK_PIPELINE_STAGE_2_COPY_BIT | VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT |
+							VK_PIPELINE_STAGE_2_ALL_TRANSFER_BIT | VK_PIPELINE_STAGE_2_TRANSFER_BIT | VK_PIPELINE_STAGE_2_RESOLVE_BIT |
+							VK_PIPELINE_STAGE_2_CLEAR_BIT | VK_PIPELINE_STAGE_2_ACCELERATION_STRUCTURE_COPY_BIT_KHR | VK_PIPELINE_STAGE_2_TOP_OF_PIPE_BIT |
+							VK_PIPELINE_STAGE_2_BOTTOM_OF_PIPE_BIT))
+		result |= 1 << kMVKBarrierStageCopy;
+
+	return result;
+}
+
 template <size_t N>
 VkResult MVKCmdPipelineBarrier<N>::setContent(MVKCommandBuffer* cmdBuff,
 											  VkPipelineStageFlags srcStageMask,
@@ -157,6 +186,15 @@
 	}
 #endif
 
+	if (!cmdEncoder->_mtlRenderEncoder && cmdEncoder->isUsingMetalArgumentBuffers() && cmdEncoder->getDevice()->hasResidencySet()) {
+		cmdEncoder->endCurrentMetalEncoding();
+
+		for (auto& b : _barriers) {
+			uint64_t sourceStageMask = mvkPipelineStageFlagsToBarrierStages(b.srcStageMask), destStageMask = mvkPipelineStageFlagsToBarrierStages(b.dstStageMask);
+			cmdEncoder->setBarrier(sourceStageMask, destStageMask);
+		}
+	}
+
 	// Apple GPUs do not support renderpass barriers, and do not support rendering/writing
 	// to an attachment and then reading from that attachment within a single renderpass.
 	// So, in the case where we are inside a Metal renderpass, we need to split those activities

MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm

@@ -617,6 +617,8 @@ static inline MTLSize mvkClampMTLSize(MTLSize size, MTLOrigin origin, MTLSize ma
                 id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPD];
 				cmdEncoder->_cmdBuffer->setMetalObjectLabel(mtlRendEnc, mvkMTLRenderCommandEncoderLabel(commandUse));
 
+				cmdEncoder->barrierWait(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+
                 float zIncr;
                 if (blitKey.srcMTLTextureType == MTLTextureType3D) {
                     // In this case, I need to interpolate along the third dimension manually.
@@ -678,7 +680,10 @@ static inline MTLSize mvkClampMTLSize(MTLSize size, MTLOrigin origin, MTLSize ma
 
                 NSUInteger instanceCount = isLayeredBlit ? mtlRPD.renderTargetArrayLengthMVK : 1;
                 [mtlRendEnc drawPrimitives: MTLPrimitiveTypeTriangleStrip vertexStart: 0 vertexCount: kMVKBlitVertexCount instanceCount: instanceCount];
-                [mtlRendEnc popDebugGroup];
+
+				cmdEncoder->barrierUpdate(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+
+				[mtlRendEnc popDebugGroup];
                 [mtlRendEnc endEncoding];
             }
         }
@@ -897,6 +902,10 @@ static inline MTLSize mvkClampMTLSize(MTLSize size, MTLOrigin origin, MTLSize ma
 
 		[mtlRendEnc pushDebugGroup: @"vkCmdResolveImage"];
 		[mtlRendEnc popDebugGroup];
+
+		cmdEncoder->barrierWait(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+		cmdEncoder->barrierUpdate(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+
 		[mtlRendEnc endEncoding];
 	}
 }
@@ -1700,6 +1709,10 @@ static inline MTLSize mvkClampMTLSize(MTLSize size, MTLOrigin origin, MTLSize ma
 
                 id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPDesc];
 				cmdEncoder->_cmdBuffer->setMetalObjectLabel(mtlRendEnc, mtlRendEncName);
+
+				cmdEncoder->barrierWait(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+				cmdEncoder->barrierUpdate(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+
                 [mtlRendEnc endEncoding];
             } else {
                 for (uint32_t layer = layerStart; layer < layerEnd; layer++) {
@@ -1715,7 +1728,11 @@ static inline MTLSize mvkClampMTLSize(MTLSize size, MTLOrigin origin, MTLSize ma
 
                     id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPDesc];
 					cmdEncoder->_cmdBuffer->setMetalObjectLabel(mtlRendEnc, mtlRendEncName);
-                    [mtlRendEnc endEncoding];
+
+					cmdEncoder->barrierWait(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+					cmdEncoder->barrierUpdate(kMVKBarrierStageCopy, mtlRendEnc, MTLRenderStageFragment);
+
+					[mtlRendEnc endEncoding];
                 }
             }
         }

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h

@@ -48,11 +48,19 @@ typedef uint64_t MVKMTLCommandBufferID;
 #pragma mark -
 #pragma mark MVKCommandEncodingContext
 
+struct BarrierFenceSlots {
+	uint32_t updateDirtyBits = ~0;
+	int update[kMVKBarrierStageCount] = {};
+	int wait[kMVKBarrierStageCount][kMVKBarrierStageCount] = {};
+};
+
 /** Context for tracking information across multiple encodings. */
 typedef struct MVKCommandEncodingContext {
 	NSUInteger mtlVisibilityResultOffset = 0;
 	const MVKMTLBufferAllocation* visibilityResultBuffer = nullptr;
+	BarrierFenceSlots fenceSlots;
 
+	void syncFences(MVKDevice *device, id<MTLCommandBuffer> mtlCommandBuffer);
 	MVKRenderPass* getRenderPass() { return _renderPass; }
 	MVKFramebuffer* getFramebuffer() { return _framebuffer; }
 	void setRenderingContext(MVKRenderPass* renderPass, MVKFramebuffer* framebuffer);
@@ -410,6 +418,27 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
     /** Returns the command encoding pool. */
     MVKCommandEncodingPool* getCommandEncodingPool();
 
+	#pragma mark Barriers
+
+	/** Encode waits in the current command encoder for the stage that corresponds to given use. */
+	void encodeBarrierWaits(MVKCommandUse use);
+
+	/** Update fences for the currently executing pipeline stage. */
+	void encodeBarrierUpdates();
+
+	/** Insert a new execution barrier */
+	void setBarrier(uint64_t sourceStageMask, uint64_t destStageMask);
+
+	/** Encode waits for a specific stage in given encoder. */
+	void barrierWait(MVKBarrierStage stage, id<MTLRenderCommandEncoder> mtlEncoder, MTLRenderStages beforeStages);
+	void barrierWait(MVKBarrierStage stage, id<MTLBlitCommandEncoder> mtlEncoder);
+	void barrierWait(MVKBarrierStage stage, id<MTLComputeCommandEncoder> mtlEncoder);
+
+	/** Encode update for a specific stage in given encoder. */
+	void barrierUpdate(MVKBarrierStage stage, id<MTLRenderCommandEncoder> mtlEncoder, MTLRenderStages afterStages);
+	void barrierUpdate(MVKBarrierStage stage, id<MTLBlitCommandEncoder> mtlEncoder);
+	void barrierUpdate(MVKBarrierStage stage, id<MTLComputeCommandEncoder> mtlEncoder);
+
 #pragma mark Queries
 
     /** Begins an occlusion query. */
@@ -492,6 +521,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
 	NSString* getMTLRenderCommandEncoderName(MVKCommandUse cmdUse);
 	template<typename T> void retainIfImmediatelyEncoding(T& mtlEnc);
 	template<typename T> void endMetalEncoding(T& mtlEnc);
+	id<MTLFence> getBarrierStageFence(MVKBarrierStage stage);
 
 	typedef struct GPUCounterQuery {
 		MVKGPUCounterQueryPool* queryPool = nullptr;