Use concurrent compute command encoders.

Add barriers to synchronize resources.
KhronosGroup · Jan 29, 2025 · 835f85e · 835f85e
1 parent aba9976
commit 835f85e
Show file tree

Hide file tree

Showing 4 changed files with 46 additions and 2 deletions.
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
@@ -243,6 +243,7 @@
                     [mtlTessCtlEncoder setBuffer: vtxOutBuff->_mtlBuffer
                                           offset: vtxOutBuff->_offset
                                          atIndex: cmdEncoder->getDevice()->getMetalBufferIndexForVertexAttributeBinding(kMVKTessCtlInputBufferBinding)];
+					[mtlTessCtlEncoder memoryBarrierWithResources:&vtxOutBuff->_mtlBuffer count:1];
                 }
 
 				NSUInteger sgSize = pipeline->getTessControlStageState().threadExecutionWidth;
@@ -475,6 +476,7 @@
                     [mtlTessCtlEncoder setBuffer: vtxOutBuff->_mtlBuffer
                                           offset: vtxOutBuff->_offset
                                          atIndex: cmdEncoder->getDevice()->getMetalBufferIndexForVertexAttributeBinding(kMVKTessCtlInputBufferBinding)];
+					[mtlTessCtlEncoder memoryBarrierWithResources:&vtxOutBuff->_mtlBuffer count:1];
                 }
 				// The vertex shader produced output in the correct order, so there's no need to use
 				// an index buffer here.
@@ -792,6 +794,8 @@
 					[mtlTessCtlEncoder dispatchThreadgroups: MTLSizeMake(mvkCeilingDivide<NSUInteger>(_drawCount, mtlConvertState.threadExecutionWidth), 1, 1)
 									  threadsPerThreadgroup: MTLSizeMake(mtlConvertState.threadExecutionWidth, 1, 1)];
 				}
+				id<MTLBuffer> barrierBuffers[] = {tempIndirectBuff->_mtlBuffer, tcParamsBuff->_mtlBuffer};
+				[mtlTessCtlEncoder memoryBarrierWithResources:barrierBuffers count:2];
                 // Mark pipelines, resources, and vertex push constants as dirty
                 // so I apply them during the next stage.
                 cmdEncoder->_graphicsPipelineState.beginMetalRenderPass();
@@ -891,6 +895,7 @@
                         [mtlTessCtlEncoder setBuffer: vtxOutBuff->_mtlBuffer
                                               offset: vtxOutBuff->_offset
                                              atIndex: cmdEncoder->getDevice()->getMetalBufferIndexForVertexAttributeBinding(kMVKTessCtlInputBufferBinding)];
+						[mtlTessCtlEncoder memoryBarrierWithResources:&vtxOutBuff->_mtlBuffer count:1];
                     }
                     [mtlTessCtlEncoder dispatchThreadgroupsWithIndirectBuffer: mtlIndBuff
                                                          indirectBufferOffset: mtlIndBuffOfst
@@ -1237,6 +1242,7 @@
 					[mtlTessCtlEncoder setBuffer: vtxIndexBuff->_mtlBuffer
 										  offset: vtxIndexBuff->_offset
 										 atIndex: pipeline->getIndirectParamsIndex().stages[kMVKShaderStageVertex]];
+					[mtlTessCtlEncoder memoryBarrierWithResources:&vtxIndexBuff->_mtlBuffer count:1];
 					[mtlTessCtlEncoder setStageInRegion: MTLRegionMake2D(0, 0, vertexCount, vertexCount)];
 					if ([mtlTessCtlEncoder respondsToSelector: @selector(setStageInRegionWithIndirectBuffer:indirectBufferOffset:)]) {
 						[mtlTessCtlEncoder setStageInRegionWithIndirectBuffer: mtlIndBuff
@@ -1275,11 +1281,13 @@
 					[mtlTessCtlEncoder setBuffer: tcParamsBuff->_mtlBuffer
 										  offset: mtlParmBuffOfst
 										 atIndex: pipeline->getIndirectParamsIndex().stages[kMVKShaderStageTessCtl]];
+					if (drawIdx == 0) [mtlTessCtlEncoder memoryBarrierWithResources:&tcParamsBuff->_mtlBuffer count:1];
 					mtlParmBuffOfst += paramsIncr;
                     if (pipeline->needsVertexOutputBuffer()) {
                         [mtlTessCtlEncoder setBuffer: vtxOutBuff->_mtlBuffer
                                               offset: vtxOutBuff->_offset
                                              atIndex: cmdEncoder->getDevice()->getMetalBufferIndexForVertexAttributeBinding(kMVKTessCtlInputBufferBinding)];
+						[mtlTessCtlEncoder memoryBarrierWithResources:&vtxOutBuff->_mtlBuffer count:1];
                     }
                     [mtlTessCtlEncoder dispatchThreadgroupsWithIndirectBuffer: mtlIndBuff
                                                          indirectBufferOffset: mtlTempIndBuffOfst

diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
@@ -113,6 +113,40 @@
 
 	auto& mtlFeats = cmdEncoder->getMetalFeatures();
 
+	if (cmdEncoder->_mtlComputeEncoder) {
+		for (auto& b : _barriers) {
+			const uint64_t stages = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT|VK_PIPELINE_STAGE_2_DRAW_INDIRECT_BIT
+				|VK_PIPELINE_STAGE_2_TRANSFER_BIT|VK_PIPELINE_STAGE_2_COPY_BIT|VK_PIPELINE_STAGE_2_BLIT_BIT|VK_PIPELINE_STAGE_2_CLEAR_BIT;
+
+			if (!(b.srcStageMask & stages) || !(b.dstStageMask & stages))
+				continue;
+
+			switch (b.type) {
+			case MVKPipelineBarrier::Memory: {
+				[cmdEncoder->_mtlComputeEncoder memoryBarrierWithScope:MTLBarrierScopeBuffers|MTLBarrierScopeTextures];
+				break;
+			}
+
+			case MVKPipelineBarrier::Buffer: {
+				id<MTLResource> mtlRez = b.mvkBuffer->getMTLBuffer();
+				[cmdEncoder->_mtlComputeEncoder memoryBarrierWithResources:&mtlRez count:1];
+				break;
+			}
+			case MVKPipelineBarrier::Image: {
+				uint32_t plnCnt = b.mvkImage->getPlaneCount();
+				id<MTLResource> mtlRezs[plnCnt];
+				for (uint8_t plnIdx = 0; plnIdx < plnCnt; plnIdx++) {
+					mtlRezs[plnIdx] = b.mvkImage->getMTLTexture(plnIdx);
+				}
+				[cmdEncoder->_mtlComputeEncoder memoryBarrierWithResources:mtlRezs count:plnCnt];
+				break;
+			}
+			default:
+				break;
+			}
+		}
+	}
+
 #if MVK_MACOS
 	// Calls below invoke MTLBlitCommandEncoder so must apply this first.
 	// Check if pipeline barriers are available and we are in a renderpass.

diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
@@ -438,6 +438,9 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
 	/** The current Metal render encoder. */
 	id<MTLRenderCommandEncoder> _mtlRenderEncoder;
 
+	/** The current Metal compute encoder. */
+	id<MTLComputeCommandEncoder> _mtlComputeEncoder;
+
     /** Tracks the current graphics pipeline bound to the encoder. */
 	MVKPipelineCommandEncoderState _graphicsPipelineState;
 
@@ -504,7 +507,6 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
 	MVKSmallVector<GPUCounterQuery, 16> _timestampStageCounterQueries;
 	MVKSmallVector<VkClearValue, kMVKDefaultAttachmentCount> _clearValues;
 	MVKSmallVector<MVKImageView*, kMVKDefaultAttachmentCount> _attachments;
-	id<MTLComputeCommandEncoder> _mtlComputeEncoder;
 	id<MTLBlitCommandEncoder> _mtlBlitEncoder;
 	id<MTLFence> _stageCountersMTLFence;
 	MVKPushConstantsCommandEncoderState _vertexPushConstants;

diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -835,7 +835,7 @@
 id<MTLComputeCommandEncoder> MVKCommandEncoder::getMTLComputeEncoder(MVKCommandUse cmdUse, bool markCurrentComputeStateDirty) {
 	if ( !_mtlComputeEncoder ) {
 		endCurrentMetalEncoding();
-		_mtlComputeEncoder = [_mtlCmdBuffer computeCommandEncoder];
+		_mtlComputeEncoder = [_mtlCmdBuffer computeCommandEncoderWithDispatchType:MTLDispatchTypeConcurrent];
 		retainIfImmediatelyEncoding(_mtlComputeEncoder);
 		beginMetalComputeEncoding(cmdUse);
 		markCurrentComputeStateDirty = false;	// Already marked dirty above in endCurrentMetalEncoding()