Merge pull request #1790 from billhollings/fix-buff-bind-overrides

Fix Metal buffer index binding overrides for push constants and attachment clearing.

Docs/Whats_New.md

@@ -19,6 +19,7 @@ MoltenVK 1.2.1
 Released TBD
 
 - Fix crash on descriptor update with out-of-bounds descriptor count data.
+- Fix Metal buffer index binding overrides for push constants and attachment clearing.
 - Work around `MTLCounterSet` crash on additional Intel Iris Plus Graphics devices.
 - Fix mistaken YCBCR format support indication.
 - Document new linkage model used by *Xcode 14* and later, and how to link **MoltenVK**

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm

@@ -697,7 +697,7 @@ void MVKCommandEncoder::finalizeDrawState(MVKGraphicsStage stage) {
         encodeStoreActions(true);
     }
     _graphicsPipelineState.encode(stage);    // Must do first..it sets others
-    _graphicsResourcesState.encode(stage);
+    _graphicsResourcesState.encode(stage);   // Before push constants, to allow them to override.
     _viewportState.encode(stage);
     _scissorState.encode(stage);
     _depthBiasState.encode(stage);
@@ -763,7 +763,7 @@ void MVKCommandEncoder::beginMetalComputeEncoding(MVKCommandUse cmdUse) {
 
 void MVKCommandEncoder::finalizeDispatchState() {
     _computePipelineState.encode();    // Must do first..it sets others
-    _computeResourcesState.encode();
+    _computeResourcesState.encode();   // Before push constants, to allow them to override.
     _computePushConstants.encode();
 }
 
@@ -882,7 +882,7 @@ void MVKCommandEncoder::setVertexBytes(id<MTLRenderCommandEncoder> mtlEncoder,
     }
 
 	if (descOverride) {
-		_graphicsResourcesState.markBufferIndexDirty(kMVKShaderStageVertex, mtlBuffIndex);
+		_graphicsResourcesState.markBufferIndexOverridden(kMVKShaderStageVertex, mtlBuffIndex);
 	}
 }
 
@@ -899,7 +899,7 @@ void MVKCommandEncoder::setFragmentBytes(id<MTLRenderCommandEncoder> mtlEncoder,
     }
 
 	if (descOverride) {
-		_graphicsResourcesState.markBufferIndexDirty(kMVKShaderStageFragment, mtlBuffIndex);
+		_graphicsResourcesState.markBufferIndexOverridden(kMVKShaderStageFragment, mtlBuffIndex);
 	}
 }
 
@@ -916,7 +916,7 @@ void MVKCommandEncoder::setComputeBytes(id<MTLComputeCommandEncoder> mtlEncoder,
     }
 
 	if (descOverride) {
-		_computeResourcesState.markBufferIndexDirty(mtlBuffIndex);
+		_computeResourcesState.markBufferIndexOverridden(mtlBuffIndex);
 	}
 }
 

MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h

@@ -377,15 +377,25 @@ protected:
         bindingsDirtyFlag = true;
     }
 
-	// Template function to find and mark dirty the binding that uses the index.
+	// Template function to find and mark as overridden the binding that uses the index.
 	template<class T>
-	void markIndexDirty(T& bindings, bool& bindingsDirtyFlag, uint32_t index) {
-		for (auto& b : bindings) {
+	void markBufferIndexOverridden(T& bufferBindings, uint32_t index) {
+		for (auto& b : bufferBindings) {
 			if (b.index == index) {
+				b.isOverridden = true;
+				return;
+			}
+		}
+	}
+
+	// Template function to mark any overridden bindings as dirty.
+	template<class T>
+	void markOverriddenBufferIndexesDirty(T& bufferBindings, bool& bindingsDirtyFlag) {
+		for (auto& b : bufferBindings) {
+			if (b.isOverridden) {
 				b.markDirty();
 				bindingsDirtyFlag = true;
 				MVKCommandEncoderState::markDirty();
-				return;
 			}
 		}
 	}
@@ -548,8 +558,14 @@ public:
 	/** Offset all buffers for vertex attribute bindings with zero divisors by the given number of strides. */
 	void offsetZeroDivisorVertexBuffers(MVKGraphicsStage stage, MVKGraphicsPipeline* pipeline, uint32_t firstInstance);
 
-	/** Marks dirty the buffer binding using the index. */
-	void markBufferIndexDirty(MVKShaderStage stage, uint32_t mtlBufferIndex);
+	/**
+	 * Marks the buffer binding using the index as having been overridden,
+	 * such as by push constants or internal rendering in some transfers.
+	 * */
+	void markBufferIndexOverridden(MVKShaderStage stage, uint32_t mtlBufferIndex);
+
+	/** Marks any overridden buffer indexes as dirty. */
+	void markOverriddenBufferIndexesDirty();
 
 	void markDirty() override;
 
@@ -562,7 +578,7 @@ protected:
     void encodeImpl(uint32_t stage) override;
 	void bindMetalArgumentBuffer(MVKShaderStage stage, MVKMTLBufferBinding& buffBind) override;
 
-    ResourceBindings<8> _shaderStageResourceBindings[4];
+    ResourceBindings<8> _shaderStageResourceBindings[kMVKShaderStageFragment + 1];
 };
 
 
@@ -600,8 +616,14 @@ public:
 										   MTLResourceUsage mtlUsage,
 										   MTLRenderStages mtlStages) override;
 
-	/** Marks dirty the buffer binding using the index. */
-	void markBufferIndexDirty(uint32_t mtlBufferIndex);
+	/**
+	 * Marks the buffer binding using the index as having been overridden,
+	 * such as by push constants or internal rendering in some transfers.
+	 * */
+	void markBufferIndexOverridden(uint32_t mtlBufferIndex);
+
+	/** Marks any overridden buffer indexes as dirty. */
+	void markOverriddenBufferIndexesDirty();
 
     void markDirty() override;
 

MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm

@@ -982,9 +982,15 @@ void MVKGraphicsResourcesCommandEncoderState::encodeArgumentBufferResourceUsage(
 	}
 }
 
-void MVKGraphicsResourcesCommandEncoderState::markBufferIndexDirty(MVKShaderStage stage, uint32_t mtlBufferIndex) {
+void MVKGraphicsResourcesCommandEncoderState::markBufferIndexOverridden(MVKShaderStage stage, uint32_t mtlBufferIndex) {
 	auto& stageRezBinds = _shaderStageResourceBindings[stage];
-	markIndexDirty(stageRezBinds.bufferBindings, stageRezBinds.areBufferBindingsDirty, mtlBufferIndex);
+	MVKResourcesCommandEncoderState::markBufferIndexOverridden(stageRezBinds.bufferBindings, mtlBufferIndex);
+}
+
+void MVKGraphicsResourcesCommandEncoderState::markOverriddenBufferIndexesDirty() {
+	for (auto& stageRezBinds : _shaderStageResourceBindings) {
+		MVKResourcesCommandEncoderState::markOverriddenBufferIndexesDirty(stageRezBinds.bufferBindings, stageRezBinds.areBufferBindingsDirty);
+	}
 }
 
 
@@ -1120,8 +1126,12 @@ void MVKComputeResourcesCommandEncoderState::encodeArgumentBufferResourceUsage(M
 	}
 }
 
-void MVKComputeResourcesCommandEncoderState::markBufferIndexDirty(uint32_t mtlBufferIndex) {
-	markIndexDirty(_resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty, mtlBufferIndex);
+void MVKComputeResourcesCommandEncoderState::markBufferIndexOverridden(uint32_t mtlBufferIndex) {
+	MVKResourcesCommandEncoderState::markBufferIndexOverridden(_resourceBindings.bufferBindings, mtlBufferIndex);
+}
+
+void MVKComputeResourcesCommandEncoderState::markOverriddenBufferIndexesDirty() {
+	MVKResourcesCommandEncoderState::markOverriddenBufferIndexesDirty(_resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty);
 }
 
 

MoltenVK/MoltenVK/Commands/MVKMTLResourceBindings.h

@@ -71,20 +71,23 @@ typedef struct MVKMTLBufferBinding {
     bool justOffset = false;
     bool isDirty = true;
     bool isInline = false;
+	bool isOverridden = false;
 
-    inline void markDirty() { justOffset = false; isDirty = true; }
+	void markDirty() { justOffset = false; isOverridden = false; isDirty = true; }
 
-    inline void update(const MVKMTLBufferBinding &other) {
+    void update(const MVKMTLBufferBinding &other) {
         if (mtlBuffer != other.mtlBuffer || size != other.size || other.isInline) {
             mtlBuffer = other.mtlBuffer;
             size = other.size;
             isInline = other.isInline;
             offset = other.offset;
             justOffset = false;
-            isDirty = true;
+			isOverridden = false;
+			isDirty = true;
         } else if (offset != other.offset) {
             offset = other.offset;
             justOffset = !isDirty || justOffset;
+			isOverridden = false;
             isDirty = true;
         }
     }

MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm

@@ -309,6 +309,8 @@ void MVKGraphicsPipeline::encode(MVKCommandEncoder* cmdEncoder, uint32_t stage)
 
             break;
     }
+
+	cmdEncoder->_graphicsResourcesState.markOverriddenBufferIndexesDirty();
     cmdEncoder->_graphicsResourcesState.bindSwizzleBuffer(_swizzleBufferIndex, _needsVertexSwizzleBuffer, _needsTessCtlSwizzleBuffer, _needsTessEvalSwizzleBuffer, _needsFragmentSwizzleBuffer);
     cmdEncoder->_graphicsResourcesState.bindBufferSizeBuffer(_bufferSizeBufferIndex, _needsVertexBufferSizeBuffer, _needsTessCtlBufferSizeBuffer, _needsTessEvalBufferSizeBuffer, _needsFragmentBufferSizeBuffer);
 	cmdEncoder->_graphicsResourcesState.bindDynamicOffsetBuffer(_dynamicOffsetBufferIndex, _needsVertexDynamicOffsetBuffer, _needsTessCtlDynamicOffsetBuffer, _needsTessEvalDynamicOffsetBuffer, _needsFragmentDynamicOffsetBuffer);
@@ -1814,6 +1816,8 @@ void MVKComputePipeline::encode(MVKCommandEncoder* cmdEncoder, uint32_t) {
 
 	[cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch) setComputePipelineState: _mtlPipelineState];
     cmdEncoder->_mtlThreadgroupSize = _mtlThreadgroupSize;
+
+	cmdEncoder->_computeResourcesState.markOverriddenBufferIndexesDirty();
 	cmdEncoder->_computeResourcesState.bindSwizzleBuffer(_swizzleBufferIndex, _needsSwizzleBuffer);
 	cmdEncoder->_computeResourcesState.bindBufferSizeBuffer(_bufferSizeBufferIndex, _needsBufferSizeBuffer);
 	cmdEncoder->_computeResourcesState.bindDynamicOffsetBuffer(_dynamicOffsetBufferIndex, _needsDynamicOffsetBuffer);