Update Metal argument buffer resource usage in Metal command encoder.

MVKPipeline track descriptors used by shaders.
Update resources as dirty at start of Metal render pass or compute encoder.
Add MVKCommandEncoder::beginMetalComputeEncoding() to mark a new Metal compute encoder.
MVKResourcesCommandEncoderState track resource usage that needs to be encoded.
Add MVKResourcesCommandEncoderState::encodeArgumentBufferResourceUsage().
Add MVKCommandEncoderState::beginMetalComputeEncoding() to mark compute
state dirty when a MTLComputeEncoder is created.
Add SPIRVToMSLConversionConfiguration::isResourceUsed().
MVKBitArray add ability to retain contents when resizing, and clear bit during getBit().
Add MVKBitArray getBit() option to clear bit.

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h

@@ -298,6 +298,9 @@ public:
 	/** Returns the index of the currently active multiview subpass, or zero if the current render pass is not multiview. */
 	uint32_t getMultiviewPassIndex();
 
+	/** Begins a Metal compute encoding. */
+	void beginMetalComputeEncoding();
+
     /** Binds a pipeline to a bind point. */
     void bindPipeline(VkPipelineBindPoint pipelineBindPoint, MVKPipeline* pipeline);
 

MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm

@@ -415,11 +415,11 @@ NSString* MVKCommandEncoder::getMTLRenderCommandEncoderName() {
 void MVKCommandEncoder::bindPipeline(VkPipelineBindPoint pipelineBindPoint, MVKPipeline* pipeline) {
     switch (pipelineBindPoint) {
         case VK_PIPELINE_BIND_POINT_GRAPHICS:
-            _graphicsPipelineState.setPipeline(pipeline);
+            _graphicsPipelineState.bindPipeline(pipeline);
             break;
 
         case VK_PIPELINE_BIND_POINT_COMPUTE:
-            _computePipelineState.setPipeline(pipeline);
+            _computePipelineState.bindPipeline(pipeline);
             break;
 
         default:
@@ -536,6 +536,12 @@ void MVKCommandEncoder::clearRenderArea() {
 	}
 }
 
+void MVKCommandEncoder::beginMetalComputeEncoding() {
+	_computePipelineState.beginMetalComputeEncoding();
+	_computeResourcesState.beginMetalComputeEncoding();
+	_computePushConstants.beginMetalComputeEncoding();
+}
+
 void MVKCommandEncoder::finalizeDispatchState() {
     _computePipelineState.encode();    // Must do first..it sets others
     _computeResourcesState.encode();
@@ -593,6 +599,7 @@ id<MTLComputeCommandEncoder> MVKCommandEncoder::getMTLComputeEncoder(MVKCommandU
 	if ( !_mtlComputeEncoder ) {
 		endCurrentMetalEncoding();
 		_mtlComputeEncoder = [_mtlCmdBuffer computeCommandEncoder];		// not retained
+		beginMetalComputeEncoding();
 	}
 	if (_mtlComputeEncoderUse != cmdUse) {
 		_mtlComputeEncoderUse = cmdUse;

MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h

@@ -23,6 +23,7 @@
 #include "MVKDevice.h"
 #include "MVKDescriptor.h"
 #include "MVKSmallVector.h"
+#include "MVKBitArray.h"
 #include <unordered_map>
 
 class MVKCommandEncoder;
@@ -68,11 +69,17 @@ public:
      */
 	virtual void beginMetalRenderPass() { if (_isModified) { markDirty(); } }
 
-	/**
+	/** Called automatically when a Metal render pass ends. */
-	 * Called automatically when a Metal render pass ends.
-	 */
 	virtual void endMetalRenderPass() { }
 
+	/**
+	 * Called automatically when a Metal compute pass begins. If the contents have been
+	 * modified from the default values, this instance is marked as dirty, so the contents
+	 * will be encoded to Metal, otherwise it is marked as clean, so the contents will NOT
+	 * be encoded. Default state can be left unencoded on a new Metal encoder.
+	 */
+	virtual void beginMetalComputeEncoding() { if (_isModified) { markDirty(); } }
+
     /**
      * If the content of this instance is dirty, marks this instance as no longer dirty
      * and calls the encodeImpl() function to encode the content onto the Metal encoder.
@@ -105,8 +112,8 @@ class MVKPipelineCommandEncoderState : public MVKCommandEncoderState {
 
 public:
 
-    /** Sets the pipeline during pipeline binding. */
+	/** Binds the pipeline. */
-    void setPipeline(MVKPipeline* pipeline);
+    void bindPipeline(MVKPipeline* pipeline);
 
     /** Returns the currently bound pipeline. */
     MVKPipeline* getPipeline();
@@ -352,19 +359,25 @@ public:
 						   uint32_t& dynamicOffsetIndex);
 
 	/** Returns the dynamic buffer offset for the descriptor in the descriptor set. */
-	inline uint32_t getDynamicBufferOffset(uint32_t descSetIndex, uint32_t descIndex) {
+	uint32_t getDynamicBufferOffset(uint32_t descSetIndex, uint32_t descIndex) {
 		return _dynamicOffsets[getDynamicOffsetKey(descSetIndex, descIndex)];
 	}
 
 	/** Sets the dynamic buffer offset for the descriptor in the descriptor set. */
-	inline void bindDynamicBufferOffset(uint32_t descSetIndex, uint32_t descIndex, uint32_t offset) {
+	void bindDynamicBufferOffset(uint32_t descSetIndex, uint32_t descIndex, uint32_t offset) {
 		_dynamicOffsets[getDynamicOffsetKey(descSetIndex, descIndex)] = offset;
 	}
 
+	/** Encodes the Metal resource to the Metal command encoder. */
+	virtual void encodeArgumentBufferResourceUsage(id<MTLResource> mtlResource,
+										   MTLResourceUsage mtlUsage,
+										   MTLRenderStages mtlStages) = 0;
+
     MVKResourcesCommandEncoderState(MVKCommandEncoder* cmdEncoder) :
 		MVKCommandEncoderState(cmdEncoder), _boundDescriptorSets{} {}
 
 protected:
+	void markDirty() override;
 
     // Template function that marks both the vector and all binding elements in the vector as dirty.
     template<class T>
@@ -453,6 +466,7 @@ protected:
 	};
 
 	MVKDescriptorSet* _boundDescriptorSets[kMVKMaxDescriptorSetCount];
+	MVKBitArray _metalUsageDirtyDescriptors[kMVKMaxDescriptorSetCount];
 
 	std::unordered_map<MVKDescSetDescKey, uint32_t> _dynamicOffsets;
 
@@ -514,6 +528,10 @@ public:
                         std::function<void(MVKCommandEncoder*, MVKMTLTextureBinding&)> bindTexture,
                         std::function<void(MVKCommandEncoder*, MVKMTLSamplerStateBinding&)> bindSampler);
 
+	void encodeArgumentBufferResourceUsage(id<MTLResource> mtlResource,
+										   MTLResourceUsage mtlUsage,
+										   MTLRenderStages mtlStages) override;
+
 	/** 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);
 
@@ -557,6 +575,10 @@ public:
     /** Sets the current buffer size buffer state. */
     void bindBufferSizeBuffer(const MVKShaderImplicitRezBinding& binding, bool needSizeBuffer);
 
+	void encodeArgumentBufferResourceUsage(id<MTLResource> mtlResource,
+										   MTLResourceUsage mtlUsage,
+										   MTLRenderStages mtlStages) override;
+
     void markDirty() override;
 
 #pragma mark Construction

MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm

@@ -35,7 +35,7 @@ MVKVulkanAPIObject* MVKCommandEncoderState::getVulkanAPIObject() { return _cmdEn
 #pragma mark -
 #pragma mark MVKPipelineCommandEncoderState
 
-void MVKPipelineCommandEncoderState::setPipeline(MVKPipeline* pipeline) {
+void MVKPipelineCommandEncoderState::bindPipeline(MVKPipeline* pipeline) {
     _pipeline = pipeline;
     markDirty();
 }
@@ -452,25 +452,92 @@ void MVKBlendColorCommandEncoderState::encodeImpl(uint32_t stage) {
 #pragma mark -
 #pragma mark MVKResourcesCommandEncoderState
 
+// Track the dynamic offsets for later binding, initialize resource usage tracking to match the
+// descriptor set content, and bind the argument buffer MTLBuffer used by the descriptor set as a resource.
 void MVKResourcesCommandEncoderState::bindDescriptorSet(uint32_t descSetIndex,
 														MVKDescriptorSet* descSet,
 														MVKShaderResourceBinding& dslMTLRezIdxOffsets,
 														MVKArrayRef<uint32_t> dynamicOffsets,
 														uint32_t& dynamicOffsetIndex) {
 	_boundDescriptorSets[descSetIndex] = descSet;
-	descSet->bindDynamicOffsets(this, descSetIndex, dynamicOffsets, dynamicOffsetIndex);
 
-	MVKMTLBufferBinding bb;
+	if (descSet->isUsingMetalArgumentBuffers()) {
-	descSet->populateMetalArgumentBufferBinding(bb);
+		descSet->bindDynamicOffsets(this, descSetIndex, dynamicOffsets, dynamicOffsetIndex);
-	bb.index = dslMTLRezIdxOffsets.stages[kMVKShaderStageVertex].bufferIndex;
+
-	bindMetalArgumentBuffer(bb);
+		auto& usageDirty = _metalUsageDirtyDescriptors[descSetIndex];
+		usageDirty.resize(descSet->getDescriptorCount());
+		usageDirty.setAllBits();
+
+		MVKMTLBufferBinding bb;
+		descSet->populateMetalArgumentBufferBinding(bb);
+		bb.index = dslMTLRezIdxOffsets.stages[kMVKShaderStageVertex].bufferIndex;
+		bindMetalArgumentBuffer(bb);
+
+		MVKCommandEncoderState::markDirty();
+	}
 }
 
+// Encode the dirty descriptors to the Metal argument buffer,
+// and set the Metal encoder usage for each resource.
 void MVKResourcesCommandEncoderState::encodeToMetalArgumentBuffer(MVKShaderStage stage) {
-	lock_guard<mutex> lock(getPipeline()->_mtlArgumentEncodingLock);
+	if ( !_cmdEncoder->isUsingMetalArgumentBuffers() ) { return; }
-	for (uint32_t dsIdx = 0; dsIdx < kMVKMaxDescriptorSetCount; dsIdx++) {
+
-		auto& mvkDescSet = _boundDescriptorSets[dsIdx];
+	MVKPipeline* pipeline = getPipeline();
-		if (mvkDescSet) { mvkDescSet->encodeToMetalArgumentBuffer(this, dsIdx, stage); }
+	lock_guard<mutex> lock(pipeline->_mtlArgumentEncodingLock);
+
+	uint32_t dsCnt = pipeline->getDescriptorSetCount();
+	for (uint32_t dsIdx = 0; dsIdx < dsCnt; dsIdx++) {
+		auto* descSet = _boundDescriptorSets[dsIdx];
+		if (descSet) {
+			auto* dsLayout = descSet->getLayout();
+			auto& usedDescriptors = pipeline->getDescriptorUsage(dsIdx);
+			auto& argBuffDirtyDescs = descSet->getMetalArgumentBufferDirtyDescriptors();
+			auto& resourceUsageDirtyDescs = _metalUsageDirtyDescriptors[dsIdx];
+
+			uint32_t elemIdx = 0;
+			uint32_t nextDSLBindDescIdx = 0;
+			MVKDescriptorSetLayoutBinding* mvkDSLBind = nullptr;
+			id<MTLArgumentEncoder> mtlArgEncoder = pipeline->getMTLArgumentEncoder(dsIdx, stage);
+			if (mtlArgEncoder) {
+				[mtlArgEncoder setArgumentBuffer: descSet->getMetalArgumentBuffer()
+										  offset: descSet->getMetalArgumentBufferOffset()];
+
+				// Only update the descriptors that are actually used by the shaders, and only if
+				// the descriptor is dirty relative to the arg buffer or Metal encoder usage setting.
+				usedDescriptors.enumerateEnabledBits(false, [&](size_t descIdx) {
+					bool argBuffDirty = argBuffDirtyDescs.getBit(descIdx, true);
+					bool resourceUsageDirty = resourceUsageDirtyDescs.getBit(descIdx, true);
+					if (argBuffDirty || resourceUsageDirty) {
+						// Get the layout binding associated with this descriptor.
+						// Assume each layout binding will apply to multiple descriptors
+						// and only fetch a new one when necessary, as it is expensive.
+						if (descIdx >= nextDSLBindDescIdx) {
+							mvkDSLBind = dsLayout->getBindingForDescriptorIndex((uint32_t)descIdx);
+							if ( !mvkDSLBind ) { return false; }	// We've run out of layout bindings
+							nextDSLBindDescIdx = mvkDSLBind->getDescriptorIndex(mvkDSLBind->getDescriptorCount(descSet));
+							elemIdx = 0;
+						}
+						auto* mvkDesc = descSet->getDescriptorAt((uint32_t)descIdx);
+						mvkDesc->encodeToMetalArgumentBuffer(this, mtlArgEncoder, dsIdx,
+															 mvkDSLBind, elemIdx++, stage,
+															 argBuffDirty, true);
+					}
+					return true;
+				});
+
+				[mtlArgEncoder setArgumentBuffer: nil offset: 0];
+			}
+		}
+	}
+}
+
+// Mark the resource usage as needing an update for each Metal render encoder.
+void MVKResourcesCommandEncoderState::markDirty() {
+	MVKCommandEncoderState::markDirty();
+	if (_cmdEncoder->isUsingMetalArgumentBuffers()) {
+		for (uint32_t dsIdx = 0; dsIdx < kMVKMaxDescriptorSetCount; dsIdx++) {
+			_metalUsageDirtyDescriptors[dsIdx].setAllBits();
+		}
 	}
 }
 
@@ -617,7 +684,7 @@ void MVKGraphicsResourcesCommandEncoderState::offsetZeroDivisorVertexBuffers(MVK
 
 // Mark everything as dirty
 void MVKGraphicsResourcesCommandEncoderState::markDirty() {
-    MVKCommandEncoderState::markDirty();
+	MVKResourcesCommandEncoderState::markDirty();
     for (uint32_t i = kMVKShaderStageVertex; i <= kMVKShaderStageFragment; i++) {
         MVKResourcesCommandEncoderState::markDirty(_shaderStageResourceBindings[i].bufferBindings, _shaderStageResourceBindings[i].areBufferBindingsDirty);
         MVKResourcesCommandEncoderState::markDirty(_shaderStageResourceBindings[i].textureBindings, _shaderStageResourceBindings[i].areTextureBindingsDirty);
@@ -800,6 +867,19 @@ void MVKGraphicsResourcesCommandEncoderState::bindMetalArgumentBuffer(MVKMTLBuff
 	}
 }
 
+void MVKGraphicsResourcesCommandEncoderState::encodeArgumentBufferResourceUsage(id<MTLResource> mtlResource,
+																				MTLResourceUsage mtlUsage,
+																				MTLRenderStages mtlStages) {
+	auto* mtlRendEnc = _cmdEncoder->_mtlRenderEncoder;
+	if (mtlRendEnc && mtlStages) {
+		if ([mtlRendEnc respondsToSelector: @selector(useResource:usage:stages:)]) {
+			[mtlRendEnc useResource: mtlResource usage: mtlUsage stages: mtlStages];
+		} else {
+			[mtlRendEnc useResource: mtlResource usage: mtlUsage];
+		}
+	}
+}
+
 
 #pragma mark -
 #pragma mark MVKComputeResourcesCommandEncoderState
@@ -830,7 +910,7 @@ void MVKComputeResourcesCommandEncoderState::bindBufferSizeBuffer(const MVKShade
 
 // Mark everything as dirty
 void MVKComputeResourcesCommandEncoderState::markDirty() {
-    MVKCommandEncoderState::markDirty();
+    MVKResourcesCommandEncoderState::markDirty();
     MVKResourcesCommandEncoderState::markDirty(_resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty);
     MVKResourcesCommandEncoderState::markDirty(_resourceBindings.textureBindings, _resourceBindings.areTextureBindingsDirty);
     MVKResourcesCommandEncoderState::markDirty(_resourceBindings.samplerStateBindings, _resourceBindings.areSamplerStateBindingsDirty);
@@ -904,6 +984,13 @@ void MVKComputeResourcesCommandEncoderState::bindMetalArgumentBuffer(MVKMTLBuffe
 	bindBuffer(buffBind);
 }
 
+void MVKComputeResourcesCommandEncoderState::encodeArgumentBufferResourceUsage(id<MTLResource> mtlResource,
+																			   MTLResourceUsage mtlUsage,
+																			   MTLRenderStages mtlStages) {
+	auto* mtlCompEnc = _cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch);
+	[mtlCompEnc useResource: mtlResource usage: mtlUsage];
+}
+
 
 #pragma mark -
 #pragma mark MVKOcclusionQueryCommandEncoderState

MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.h

@@ -120,6 +120,9 @@ public:
 		return _mtlResourceIndexOffsets.stages[stage].resourceIndex + elementIndex;
 	}
 
+	/** Returns a bitwise OR of Metal render stages. */
+	MTLRenderStages getMTLRenderStages();
+
 	MVKDescriptorSetLayoutBinding(MVKDevice* device,
 								  MVKDescriptorSetLayout* layout,
 								  const VkDescriptorSetLayoutBinding* pBinding,
@@ -193,7 +196,9 @@ public:
 											 uint32_t descSetIndex,
 											 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 											 uint32_t elementIndex,
-											 MVKShaderStage stage) = 0;
+											 MVKShaderStage stage,
+											 bool encodeToArgBuffer,
+											 bool encodeUsage) = 0;
 
 	/**
 	 * Updates the internal binding from the specified content. The format of the content depends
@@ -230,6 +235,9 @@ public:
 
 	~MVKDescriptor() { reset(); }
 
+protected:
+	MTLResourceUsage getMTLResourceUsage();
+
 };
 
 
@@ -253,7 +261,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	void write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
 			   MVKDescriptorSet* mvkDescSet,
@@ -340,7 +350,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	void write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
 			   MVKDescriptorSet* mvkDescSet,
@@ -387,7 +399,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	void write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
 			   MVKDescriptorSet* mvkDescSet,
@@ -465,7 +479,8 @@ protected:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage);
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer);
 
 	/**
 	 * Offset to the first sampler index in the argument buffer. Defaults to zero for simple sampler
@@ -517,7 +532,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	void write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
 			   MVKDescriptorSet* mvkDescSet,
@@ -562,7 +579,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	uint32_t getMetalArgumentBufferSamplerIndexOffset(MVKDescriptorSetLayoutBinding* dslBinding) override;
 
@@ -607,7 +626,9 @@ public:
 									 uint32_t descSetIndex,
 									 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 									 uint32_t elementIndex,
-									 MVKShaderStage stage) override;
+									 MVKShaderStage stage,
+									 bool encodeToArgBuffer,
+									 bool encodeUsage) override;
 
 	void write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
 			   MVKDescriptorSet* mvkDescSet,

MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.mm

@@ -453,6 +453,29 @@ bool MVKDescriptorSetLayoutBinding::validate(MVKSampler* mvkSampler) {
 	return true;
 }
 
+MTLRenderStages MVKDescriptorSetLayoutBinding::getMTLRenderStages() {
+	MTLRenderStages mtlStages = 0;
+	for (uint32_t stage = kMVKShaderStageVertex; stage < kMVKShaderStageMax; stage++) {
+		if (_applyToStage[stage]) {
+			switch (stage) {
+				case kMVKShaderStageVertex:
+				case kMVKShaderStageTessCtl:
+				case kMVKShaderStageTessEval:
+					mtlStages |= MTLRenderStageVertex;
+					break;
+
+				case kMVKShaderStageFragment:
+					mtlStages |= MTLRenderStageFragment;
+					break;
+
+				default:
+					break;
+			}
+		}
+	}
+	return mtlStages;
+}
+
 MVKDescriptorSetLayoutBinding::MVKDescriptorSetLayoutBinding(MVKDevice* device,
 															 MVKDescriptorSetLayout* layout,
 															 const VkDescriptorSetLayoutBinding* pBinding,
@@ -606,6 +629,31 @@ void MVKDescriptorSetLayoutBinding::initMetalResourceIndexOffsets(const VkDescri
 }
 
 
+#pragma mark -
+#pragma mark MVKDescriptor
+
+MTLResourceUsage MVKDescriptor::getMTLResourceUsage() {
+	MTLResourceUsage mtlUsage = MTLResourceUsageRead;
+	switch (getDescriptorType()) {
+		case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE:
+		case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER:
+			mtlUsage |= MTLResourceUsageSample;
+			break;
+
+		case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE:
+		case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER:
+		case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER:
+		case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC:
+			mtlUsage |= MTLResourceUsageWrite;
+			break;
+
+		default:
+			break;
+	}
+	return mtlUsage;
+}
+
+
 #pragma mark -
 #pragma mark MVKBufferDescriptor
 
@@ -645,13 +693,22 @@ void MVKBufferDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCommandEncoder
 													  uint32_t descSetIndex,
 													  MVKDescriptorSetLayoutBinding* mvkDSLBind,
 													  uint32_t elementIndex,
-													  MVKShaderStage stage) {
+													  MVKShaderStage stage,
-	NSUInteger bufferDynamicOffset = (usesDynamicBufferOffsets()
+													  bool encodeToArgBuffer,
-									  ? rezEncState->getDynamicBufferOffset(descSetIndex, mvkDSLBind->getDescriptorIndex(elementIndex))
+													  bool encodeUsage) {
-									  : 0);
+	if (encodeToArgBuffer) {
-	[mtlArgEncoder setBuffer: _mvkBuffer ? _mvkBuffer->getMTLBuffer() : nil
+		NSUInteger bufferDynamicOffset = (usesDynamicBufferOffsets()
-					  offset: _mvkBuffer ? _mvkBuffer->getMTLBufferOffset() + _buffOffset + bufferDynamicOffset : 0
+										  ? rezEncState->getDynamicBufferOffset(descSetIndex, mvkDSLBind->getDescriptorIndex(elementIndex))
-					 atIndex: mvkDSLBind->getMetalArgumentBufferIndex(stage, elementIndex)];
+										  : 0);
+		[mtlArgEncoder setBuffer: _mvkBuffer ? _mvkBuffer->getMTLBuffer() : nil
+						  offset: _mvkBuffer ? _mvkBuffer->getMTLBufferOffset() + _buffOffset + bufferDynamicOffset : 0
+						 atIndex: mvkDSLBind->getMetalArgumentBufferIndex(stage, elementIndex)];
+	}
+	if (encodeUsage) {
+		rezEncState->encodeArgumentBufferResourceUsage(_mvkBuffer ? _mvkBuffer->getMTLBuffer() : nil,
+													   getMTLResourceUsage(),
+													   mvkDSLBind->getMTLRenderStages());
+	}
 }
 
 void MVKBufferDescriptor::write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
@@ -727,10 +784,19 @@ void MVKInlineUniformBlockDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCo
 																  uint32_t descSetIndex,
 																  MVKDescriptorSetLayoutBinding* mvkDSLBind,
 																  uint32_t elementIndex,
-																  MVKShaderStage stage) {
+																  MVKShaderStage stage,
-	[mtlArgEncoder setBuffer: _mvkMTLBufferAllocation ? _mvkMTLBufferAllocation->_mtlBuffer : nil
+																  bool encodeToArgBuffer,
-					  offset: _mvkMTLBufferAllocation ? _mvkMTLBufferAllocation->_offset : 0
+																  bool encodeUsage) {
-					 atIndex: mvkDSLBind->getMetalArgumentBufferIndex(stage)];
+	if (encodeToArgBuffer) {
+		[mtlArgEncoder setBuffer: _mvkMTLBufferAllocation ? _mvkMTLBufferAllocation->_mtlBuffer : nil
+						  offset: _mvkMTLBufferAllocation ? _mvkMTLBufferAllocation->_offset : 0
+						 atIndex: mvkDSLBind->getMetalArgumentBufferIndex(stage)];
+	}
+	if (encodeUsage) {
+		rezEncState->encodeArgumentBufferResourceUsage(_mvkMTLBufferAllocation ? _mvkMTLBufferAllocation->_mtlBuffer : nil,
+													   getMTLResourceUsage(),
+													   mvkDSLBind->getMTLRenderStages());
+	}
 }
 
 void MVKInlineUniformBlockDescriptor::write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
@@ -829,7 +895,9 @@ void MVKImageDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCommandEncoderS
 													 uint32_t descSetIndex,
 													 MVKDescriptorSetLayoutBinding* mvkDSLBind,
 													 uint32_t elementIndex,
-													 MVKShaderStage stage) {
+													 MVKShaderStage stage,
+													 bool encodeToArgBuffer,
+													 bool encodeUsage) {
 	VkDescriptorType descType = getDescriptorType();
 	uint8_t planeCount = (_mvkImageView) ? _mvkImageView->getPlaneCount() : 1;
 	uint32_t buffArgIdxOffset = mvkDSLBind->getDescriptorCount() * planeCount;
@@ -839,13 +907,26 @@ void MVKImageDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCommandEncoderS
 
 		id<MTLTexture> mtlTexture = _mvkImageView ? _mvkImageView->getMTLTexture(planeIndex) : nil;
 		uint32_t texArgIdx = mvkDSLBind->getMetalArgumentBufferIndex(stage, planeDescIdx);
-		[mtlArgEncoder setTexture: mtlTexture atIndex: texArgIdx];
+		if (encodeToArgBuffer) {
-
+			[mtlArgEncoder setTexture: mtlTexture atIndex: texArgIdx];
+		}
+		if (encodeUsage) {
+			rezEncState->encodeArgumentBufferResourceUsage(mtlTexture,
+														   getMTLResourceUsage(),
+														   mvkDSLBind->getMTLRenderStages());
+		}
 		if (descType == VK_DESCRIPTOR_TYPE_STORAGE_IMAGE && mtlTexture) {
 			id<MTLTexture> mtlTex = mtlTexture.parentTexture ? mtlTexture.parentTexture : mtlTexture;
-			[mtlArgEncoder setBuffer: mtlTex.buffer
+			if (encodeToArgBuffer) {
-							  offset: mtlTex.bufferOffset
+				[mtlArgEncoder setBuffer: mtlTex.buffer
-							 atIndex: texArgIdx + buffArgIdxOffset];
+								  offset: mtlTex.bufferOffset
+								 atIndex: texArgIdx + buffArgIdxOffset];
+			}
+			if (encodeUsage) {
+				rezEncState->encodeArgumentBufferResourceUsage(mtlTex.buffer,
+															   getMTLResourceUsage(),
+															   mvkDSLBind->getMTLRenderStages());
+			}
 		}
 	}
 }
@@ -923,15 +1004,17 @@ void MVKSamplerDescriptorMixin::encodeToMetalArgumentBuffer(MVKResourcesCommandE
 															uint32_t descSetIndex,
 															MVKDescriptorSetLayoutBinding* mvkDSLBind,
 															uint32_t elementIndex,
-															MVKShaderStage stage) {
+															MVKShaderStage stage,
-
+															bool encodeToArgBuffer) {
-	MVKSampler* imutSamp = mvkDSLBind->getImmutableSampler(elementIndex);
+	if (encodeToArgBuffer) {
-	MVKSampler* mvkSamp = imutSamp ? imutSamp : _mvkSampler;
+		MVKSampler* imutSamp = mvkDSLBind->getImmutableSampler(elementIndex);
-	id<MTLSamplerState> mtlSamp = (mvkSamp
+		MVKSampler* mvkSamp = imutSamp ? imutSamp : _mvkSampler;
-								   ? mvkSamp->getMTLSamplerState()
+		id<MTLSamplerState> mtlSamp = (mvkSamp
-								   : mvkDSLBind->getDevice()->getDefaultMTLSamplerState());
+									   ? mvkSamp->getMTLSamplerState()
-	uint32_t argIdx = mvkDSLBind->getMetalArgumentBufferIndex(stage, getMetalArgumentBufferSamplerIndexOffset(mvkDSLBind) + elementIndex);
+									   : mvkDSLBind->getDevice()->getDefaultMTLSamplerState());
-	[mtlArgEncoder setSamplerState: mtlSamp atIndex: argIdx];
+		uint32_t argIdx = mvkDSLBind->getMetalArgumentBufferIndex(stage, getMetalArgumentBufferSamplerIndexOffset(mvkDSLBind) + elementIndex);
+		[mtlArgEncoder setSamplerState: mtlSamp atIndex: argIdx];
+	}
 }
 
 void MVKSamplerDescriptorMixin::write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
@@ -987,8 +1070,10 @@ void MVKSamplerDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCommandEncode
 													   uint32_t descSetIndex,
 													   MVKDescriptorSetLayoutBinding* mvkDSLBind,
 													   uint32_t elementIndex,
-													   MVKShaderStage stage) {
+													   MVKShaderStage stage,
-	MVKSamplerDescriptorMixin::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage);
+													   bool encodeToArgBuffer,
+													   bool encodeUsage) {
+	MVKSamplerDescriptorMixin::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage, encodeToArgBuffer);
 }
 
 void MVKSamplerDescriptor::write(MVKDescriptorSetLayoutBinding* mvkDSLBind,
@@ -1035,9 +1120,11 @@ void MVKCombinedImageSamplerDescriptor::encodeToMetalArgumentBuffer(MVKResources
 																	uint32_t descSetIndex,
 																	MVKDescriptorSetLayoutBinding* mvkDSLBind,
 																	uint32_t elementIndex,
-																	MVKShaderStage stage) {
+																	MVKShaderStage stage,
-	MVKImageDescriptor::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage);
+																	bool encodeToArgBuffer,
-	MVKSamplerDescriptorMixin::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage);
+																	bool encodeUsage) {
+	MVKImageDescriptor::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage, encodeToArgBuffer, encodeUsage);
+	MVKSamplerDescriptorMixin::encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex, mvkDSLBind, elementIndex, stage, encodeToArgBuffer);
 }
 
 uint32_t MVKCombinedImageSamplerDescriptor::getMetalArgumentBufferSamplerIndexOffset(MVKDescriptorSetLayoutBinding* dslBinding) {
@@ -1118,16 +1205,32 @@ void MVKTexelBufferDescriptor::encodeToMetalArgumentBuffer(MVKResourcesCommandEn
 														   uint32_t descSetIndex,
 														   MVKDescriptorSetLayoutBinding* mvkDSLBind,
 														   uint32_t elementIndex,
-														   MVKShaderStage stage) {
+														   MVKShaderStage stage,
+														   bool encodeToArgBuffer,
+														   bool encodeUsage) {
 	VkDescriptorType descType = getDescriptorType();
 	id<MTLTexture> mtlTexture = _mvkBufferView ? _mvkBufferView->getMTLTexture() : nil;
 	uint32_t texArgIdx = mvkDSLBind->getMetalArgumentBufferIndex(stage, elementIndex);
-	[mtlArgEncoder setTexture: mtlTexture atIndex: texArgIdx];
+	if (encodeToArgBuffer) {
+		[mtlArgEncoder setTexture: mtlTexture atIndex: texArgIdx];
+	}
+	if (encodeUsage) {
+		rezEncState->encodeArgumentBufferResourceUsage(mtlTexture,
+													   getMTLResourceUsage(),
+													   mvkDSLBind->getMTLRenderStages());
+	}
 
 	if (descType == VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER && mtlTexture) {
-		[mtlArgEncoder setBuffer: mtlTexture.buffer
+		if (encodeToArgBuffer) {
-						  offset: mtlTexture.bufferOffset
+			[mtlArgEncoder setBuffer: mtlTexture.buffer
-						 atIndex: texArgIdx + mvkDSLBind->getDescriptorCount()];
+							  offset: mtlTexture.bufferOffset
+							 atIndex: texArgIdx + mvkDSLBind->getDescriptorCount()];
+		}
+		if (encodeUsage) {
+			rezEncState->encodeArgumentBufferResourceUsage(mtlTexture.buffer,
+														   getMTLResourceUsage(),
+														   mvkDSLBind->getMTLRenderStages());
+		}
 	}
 }
 

MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h

@@ -72,6 +72,14 @@ public:
                                         MVKShaderResourceBinding& dslMTLRezIdxOffsets,
                                         uint32_t dslIndex);
 
+	/** Populates the descriptor usage as indicated by the shader converter context. */
+	void populateDescriptorUsage(MVKBitArray& usageArray,
+								 mvk::SPIRVToMSLConversionConfiguration& context,
+								 uint32_t dslIndex);
+
+	/** Returns the binding for the descriptor at the index in a descriptor set. */
+	MVKDescriptorSetLayoutBinding* getBindingForDescriptorIndex(uint32_t descriptorIndex);
+
 	/** Returns true if this layout is for push descriptors only. */
 	bool isPushDescriptorLayout() const { return _isPushDescriptorLayout; }
 
@@ -91,7 +99,6 @@ protected:
 	inline uint32_t getDescriptorCount() { return _descriptorCount; }
 	inline uint32_t getDescriptorIndex(uint32_t binding, uint32_t elementIndex = 0) { return getBinding(binding)->getDescriptorIndex(elementIndex); }
 	inline MVKDescriptorSetLayoutBinding* getBinding(uint32_t binding) { return &_bindings[_bindingToIndex[binding]]; }
-	MVKDescriptorSetLayoutBinding* getBindingForDescriptorIndex(uint32_t descriptorIndex);
 	const VkDescriptorBindingFlags* getBindingFlags(const VkDescriptorSetLayoutCreateInfo* pCreateInfo);
 	inline bool isUsingMetalArgumentBuffer()  { return isUsingMetalArgumentBuffers() && !isPushDescriptorLayout(); };
 
@@ -118,6 +125,9 @@ public:
 	/** Returns the debug report object type of this object. */
 	VkDebugReportObjectTypeEXT getVkDebugReportObjectType() override { return VK_DEBUG_REPORT_OBJECT_TYPE_DESCRIPTOR_SET_EXT; }
 
+	/** Returns the layout that defines this descriptor set. */
+	MVKDescriptorSetLayout* getLayout() { return _layout; }
+
 	/** Returns the descriptor type for the specified binding number. */
 	VkDescriptorType getDescriptorType(uint32_t binding);
 
@@ -141,16 +151,28 @@ public:
 							MVKArrayRef<uint32_t> dynamicOffsets,
 							uint32_t& dynamicOffsetIndex);
 
-	/** Encode any dirty descriptors to the arugment buffer. */
-	void encodeToMetalArgumentBuffer(MVKResourcesCommandEncoderState* rezEncState,
-									 uint32_t descSetIndex,
-									 MVKShaderStage stage);
-
 	/** Populates the buffer binding with the Metal argument buffer and offset. */
 	void populateMetalArgumentBufferBinding(MVKMTLBufferBinding& buffBind);
 
 	/** Returns an MTLBuffer region allocation. */
 	const MVKMTLBufferAllocation* acquireMTLBufferRegion(NSUInteger length);
+	/**
+	 * Returns the Metal argument buffer to which resources are written,
+	 * or return nil if Metal argument buffers are not being used.
+	 */
+	id<MTLBuffer> getMetalArgumentBuffer();
+
+	/** Returns the offset into the Metal argument buffer to which resources are written. */
+	inline NSUInteger getMetalArgumentBufferOffset() { return _metalArgumentBufferOffset; }
+
+	/** Returns an array indicating the descriptors that have changed since the Metal argument buffer was last updated. */
+	MVKBitArray& getMetalArgumentBufferDirtyDescriptors() { return _metalArgumentBufferDirtyDescriptors; }
+
+	/** Returns the descriptor at an index. */
+	MVKDescriptor* getDescriptorAt(uint32_t descIndex) { return _descriptors[descIndex]; }
+
+	/** Returns the number of descriptors in this descriptor set. */
+	uint32_t getDescriptorCount() { return (uint32_t)_descriptors.size(); }
 
 	MVKDescriptorSet(MVKDescriptorPool* pool);
 

MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm

@@ -38,11 +38,10 @@ void MVKDescriptorSetLayout::bindDescriptorSet(MVKCommandEncoder* cmdEncoder,
 	if (!cmdEncoder) { clearConfigurationResult(); }
 	if (_isPushDescriptorLayout ) { return; }
 
-	if (isUsingMetalArgumentBuffer()) {
+	if (cmdEncoder) { cmdEncoder->bindDescriptorSet(pipelineBindPoint, descSetIndex,
-		if (cmdEncoder) { cmdEncoder->bindDescriptorSet(pipelineBindPoint, descSetIndex,
+													descSet, dslMTLRezIdxOffsets,
-														descSet, dslMTLRezIdxOffsets,
+													dynamicOffsets, dynamicOffsetIndex); }
-														dynamicOffsets, dynamicOffsetIndex); }
+	if ( !isUsingMetalArgumentBuffers() ) {
-	} else {
 		for (auto& dslBind : _bindings) {
 			dslBind.bind(cmdEncoder, descSet, dslMTLRezIdxOffsets, dynamicOffsets, dynamicOffsetIndex);
 		}
@@ -184,6 +183,21 @@ void MVKDescriptorSetLayout::populateShaderConverterContext(mvk::SPIRVToMSLConve
 	}
 }
 
+void MVKDescriptorSetLayout::populateDescriptorUsage(MVKBitArray& usageArray,
+													 SPIRVToMSLConversionConfiguration& context,
+													 uint32_t dslIndex) {
+	uint32_t bindCnt = (uint32_t)_bindings.size();
+	for (uint32_t bindIdx = 0; bindIdx < bindCnt; bindIdx++) {
+		auto& dslBind = _bindings[bindIdx];
+		if (context.isResourceUsed(dslIndex, dslBind.getBinding())) {
+			uint32_t elemCnt = dslBind.getDescriptorCount();
+			for (uint32_t elemIdx = 0; elemIdx < elemCnt; elemIdx++) {
+				usageArray.setBit(dslBind.getDescriptorIndex(elemIdx));
+			}
+		}
+	}
+}
+
 id<MTLArgumentEncoder> MVKDescriptorSetLayout::newMTLArgumentEncoder(uint32_t stage,
 																	 mvk::SPIRVToMSLConversionConfiguration& shaderConfig,
 																	 uint32_t descSetIdx) {
@@ -273,6 +287,8 @@ MVKDescriptor* MVKDescriptorSet::getDescriptor(uint32_t binding, uint32_t elemen
 	return _descriptors[_layout->getDescriptorIndex(binding, elementIndex)];
 }
 
+id<MTLBuffer> MVKDescriptorSet::getMetalArgumentBuffer() { return _pool->_metalArgumentBuffer; }
+
 template<typename DescriptorAction>
 void MVKDescriptorSet::write(const DescriptorAction* pDescriptorAction,
 							 size_t stride,
@@ -341,34 +357,6 @@ void MVKDescriptorSet::bindDynamicOffsets(MVKResourcesCommandEncoderState* rezEn
 	});
 }
 
-void MVKDescriptorSet::encodeToMetalArgumentBuffer(MVKResourcesCommandEncoderState* rezEncState,
-												   uint32_t descSetIndex,
-												   MVKShaderStage stage) {
-	uint32_t elemIdx = 0;
-	uint32_t nextDSLBindDescIdx = 0;
-	MVKDescriptorSetLayoutBinding* mvkDSLBind = nullptr;
-	id<MTLArgumentEncoder> mtlArgEncoder = rezEncState->getPipeline()->getMTLArgumentEncoder(descSetIndex, stage);
-	if ( !mtlArgEncoder ) { return; }
-
-	[mtlArgEncoder setArgumentBuffer: _pool->_metalArgumentBuffer offset: _metalArgumentBufferOffset];
-
-	_metalArgumentBufferDirtyDescriptors.enumerateEnabledBits(true, [&](size_t descIdx) {
-		// Get the layout binding associated with this descriptor.
-		// Assume each layout binding will apply to multiple descriptors and only fetch a new one when necessary.
-		if (descIdx >= nextDSLBindDescIdx) {
-			mvkDSLBind = _layout->getBindingForDescriptorIndex((uint32_t)descIdx);
-			if ( !mvkDSLBind ) { return false; }	// We've run out of layout bindings
-			nextDSLBindDescIdx = mvkDSLBind->getDescriptorIndex(mvkDSLBind->getDescriptorCount(this));
-			elemIdx = 0;
-		}
-		_descriptors[descIdx]->encodeToMetalArgumentBuffer(rezEncState, mtlArgEncoder, descSetIndex,
-														   mvkDSLBind, elemIdx++, stage);
-		return true;
-	});
-
-	[mtlArgEncoder setArgumentBuffer: nil offset: 0];
-}
-
 void MVKDescriptorSet::populateMetalArgumentBufferBinding(MVKMTLBufferBinding& buffBind) {
 	buffBind.mtlBuffer = _pool->_metalArgumentBuffer;
 	buffBind.offset = _metalArgumentBufferOffset;

MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h

@@ -23,6 +23,7 @@
 #include "MVKShaderModule.h"
 #include "MVKSync.h"
 #include "MVKSmallVector.h"
+#include "MVKBitArray.h"
 #include <MoltenVKShaderConverter/SPIRVReflection.h>
 #include <MoltenVKShaderConverter/SPIRVToMSLConverter.h>
 #include <unordered_map>
@@ -104,6 +105,9 @@ public:
 	/** Returns the number of descriptor sets in this pipeline layout. */
 	uint32_t getDescriptorSetCount() { return (uint32_t)_descriptorSetLayouts.size(); }
 
+	/** Returns the number of descriptors in the descriptor set layout. */
+	uint32_t getDescriptorCount(uint32_t descSetIndex) { return _descriptorSetLayouts[descSetIndex]->getDescriptorCount(); }
+
 	/** Returns the push constant binding info. */
 	const MVKShaderResourceBinding& getPushConstantBindings() { return _pushConstantsMTLResourceIndexes; }
 
@@ -175,10 +179,16 @@ public:
 	bool hasValidMTLPipelineStates() { return _hasValidMTLPipelineStates; }
 
 	/** Returns the MTLArgumentEncoder for the descriptor set. */
-	inline id<MTLArgumentEncoder> getMTLArgumentEncoder(uint32_t descSetIndex, MVKShaderStage stage) {
+	id<MTLArgumentEncoder> getMTLArgumentEncoder(uint32_t descSetIndex, MVKShaderStage stage) {
 		return _mtlArgumentEncoders[descSetIndex];
 	}
 
+	/** Returns the number of descriptor sets in this pipeline layout. */
+	uint32_t getDescriptorSetCount() { return _descriptorSetCount; }
+
+	/** Returns the descriptor usage array for the descriptor set. */
+	MVKBitArray& getDescriptorUsage(uint32_t descSetIndex) { return _descriptorUsage[descSetIndex]; }
+
 	/** A mutex lock to protect access to the Metal argument encoders. */
 	std::mutex _mtlArgumentEncodingLock;
 
@@ -190,13 +200,16 @@ public:
 protected:
 	void propagateDebugName() override {}
 	void addMTLArgumentEncoders(MVKPipelineLayout* layout, SPIRVToMSLConversionConfiguration& shaderConfig);
+	void initDescriptorUsage(MVKPipelineLayout* layout);
 
 	MVKPipelineCache* _pipelineCache;
+	MVKBitArray _descriptorUsage[kMVKMaxDescriptorSetCount];
 	id<MTLArgumentEncoder> _mtlArgumentEncoders[kMVKMaxDescriptorSetCount];
 	MVKShaderImplicitRezBinding _swizzleBufferIndex;
 	MVKShaderImplicitRezBinding _bufferSizeBufferIndex;
 	MVKShaderImplicitRezBinding _indirectParamsIndex;
 	MVKShaderResourceBinding _pushConstantsMTLResourceIndexes;
+	uint32_t _descriptorSetCount;
 	bool _fullImageViewSwizzle;
 	bool _hasValidMTLPipelineStates = true;
 

MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm

@@ -172,11 +172,16 @@ void MVKPipeline::bindPushConstants(MVKCommandEncoder* cmdEncoder) {
 	}
 }
 
+// For each descriptor set, mark the descriptors that are used by the shaders,
+// and build Metal argument encoders that match the shader requirements.
 void MVKPipeline::addMTLArgumentEncoders(MVKPipelineLayout* layout, SPIRVToMSLConversionConfiguration& shaderConfig) {
-	uint32_t dsCnt = layout->getDescriptorSetCount();
+	if ( !isUsingMetalArgumentBuffers() ) { return; }
+
 	uint32_t stage = kMVKShaderStageVertex;		// Nominal stage. Currently all stages use same encoders.
-	for (uint32_t dsIdx = 0; dsIdx < dsCnt; dsIdx++) {
+	for (uint32_t dsIdx = 0; dsIdx < _descriptorSetCount; dsIdx++) {
-		_mtlArgumentEncoders[dsIdx] = layout->_descriptorSetLayouts[dsIdx]->newMTLArgumentEncoder(stage, shaderConfig, dsIdx);	// retained
+		auto* mvkDSL = layout->_descriptorSetLayouts[dsIdx];
+		mvkDSL->populateDescriptorUsage(_descriptorUsage[dsIdx], shaderConfig, dsIdx);
+		_mtlArgumentEncoders[dsIdx] = mvkDSL->newMTLArgumentEncoder(stage, shaderConfig, dsIdx);	// retained
 	}
 }
 
@@ -185,7 +190,20 @@ MVKPipeline::MVKPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVK
 	_pipelineCache(pipelineCache),
 	_pushConstantsMTLResourceIndexes(layout->getPushConstantBindings()),
 	_fullImageViewSwizzle(mvkConfig()->fullImageViewSwizzle),
-	_mtlArgumentEncoders{} {}
+	_mtlArgumentEncoders{},
+	_descriptorSetCount(layout->getDescriptorSetCount()) {
+		initDescriptorUsage(layout);
+	}
+
+void MVKPipeline::initDescriptorUsage(MVKPipelineLayout* layout) {
+	_descriptorSetCount = layout->getDescriptorSetCount();
+
+	if (isUsingMetalArgumentBuffers() ) {
+		for (uint32_t dsIdx = 0; dsIdx < _descriptorSetCount; dsIdx++) {
+			_descriptorUsage[dsIdx].resize(layout->getDescriptorCount(dsIdx));
+		}
+	}
+}
 
 MVKPipeline::~MVKPipeline() {
 	for (uint32_t dsIdx = 0; dsIdx < kMVKMaxDescriptorSetCount; dsIdx++) {

MoltenVK/MoltenVK/Utility/MVKBitArray.h

@@ -34,9 +34,11 @@ class MVKBitArray {
 
 public:
 
-	/** Returns the value of the bit. */
+	/** Returns the value of the bit, and optionally clears that bit if it was set. */
-	inline bool getBit(size_t bitIndex) {
+	inline bool getBit(size_t bitIndex, bool shouldClear = false) {
-		return mvkIsAnyFlagEnabled(_pSections[getIndexOfSection(bitIndex)], getSectionSetMask(bitIndex));
+		bool val = mvkIsAnyFlagEnabled(_pSections[getIndexOfSection(bitIndex)], getSectionSetMask(bitIndex));
+		if (shouldClear && val) { clearBit(bitIndex); }
+		return val;
 	}
 
 	/** Sets the value of the bit to the val (or to 1 by default). */
@@ -132,22 +134,47 @@ public:
 	/** Returns whether this array is empty. */
 	inline bool empty() { return !_bitCount; }
 
-	/** Resize this array to the specified number of bits, and sets the initial value of all the bits. */
+	/**
+	 * Resize this array to the specified number of bits. The value of existing
+	 * bits that fit within the new size are retained, and any new bits that
+	 * are added to accommodate the new size are set to the given value.
+	 * Consumed memory is retained unless the size is set to zero.
+	 */
 	inline void resize(size_t size = 0, bool val = false) {
-		free(_pSections);
+		size_t oldBitCnt = _bitCount;
+		size_t oldSecCnt = getSectionCount();
 
 		_bitCount = size;
-		_pSections = _bitCount ? (uint64_t*)malloc(getSectionCount() * SectionByteCount) : nullptr;
+		size_t newSecCnt = getSectionCount();
-		if (val) {
+
-			setAllBits();
+		if (newSecCnt > oldSecCnt) {
-		} else {
+			uint64_t* pOldSecs = _pSections;
-			clearAllBits();
+			size_t oldByteCnt = oldSecCnt * SectionByteCount;
+			size_t newByteCnt = newSecCnt * SectionByteCount;
+
+			// Allocate new memory and fill it with the new initial value
+			_pSections = _bitCount ? (uint64_t*)malloc(newByteCnt) : nullptr;
+			if (_pSections) { memset(_pSections, val ? ~0 : 0, newByteCnt); }
+
+			// Copy the old contents to the new memory, and fill any bits in the old
+			// last section that were beyond the old bit count with the new initial value.
+			if (_pSections && pOldSecs) { memcpy(_pSections, pOldSecs, oldByteCnt); }
+			size_t oldEndBitCnt = oldSecCnt << SectionMaskSize;
+			for (size_t bitIdx = oldBitCnt; bitIdx < oldEndBitCnt; bitIdx++) { setBit(bitIdx, val); }
+
+			// If the entire old array and the new array are cleared, move the indicated to the new end.
+			if (_minUnclearedSectionIndex == oldSecCnt && !val) { _minUnclearedSectionIndex = newSecCnt; }
+
+			free(pOldSecs);
+		} else if (newSecCnt == 0) {
+			free(_pSections);
+			_pSections = nullptr;
+			_minUnclearedSectionIndex = 0;
 		}
 	}
 
 	/** Constructs an instance for the specified number of bits, and sets the initial value of all the bits. */
 	MVKBitArray(size_t size = 0, bool val = false) {
-		_pSections = nullptr;
 		resize(size, val);
 	}
 
@@ -195,7 +222,7 @@ protected:
 		_minUnclearedSectionIndex = sectionValue ? 0 : secCnt;
 	}
 
-	uint64_t* _pSections;
+	uint64_t* _pSections = nullptr;
-	size_t _bitCount;
+	size_t _bitCount = 0;
-	size_t _minUnclearedSectionIndex;	// Tracks where to start looking for bits that are set
+	size_t _minUnclearedSectionIndex = 0;	// Tracks where to start looking for bits that are set
 };

MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h

@@ -159,6 +159,9 @@ namespace mvk {
         /** Returns whether the vertex buffer at the specified Vulkan binding is used by the shader. */
 		bool isVertexBufferUsed(uint32_t binding) const { return countShaderInputsAt(binding) > 0; }
 
+		/** Returns whether the resource at the specified descriptor set binding is used by the shader. */
+		bool isResourceUsed(uint32_t descSet, uint32_t binding) const;
+
 		/** Returns the MTLTextureType of the image resource at the descriptor set and binding. */
 		MTLTextureType getMTLTextureType(uint32_t descSet, uint32_t binding) const;
 

MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.mm

@@ -180,6 +180,16 @@ MVK_PUBLIC_SYMBOL uint32_t SPIRVToMSLConversionConfiguration::countShaderInputsA
 	return siCnt;
 }
 
+MVK_PUBLIC_SYMBOL bool SPIRVToMSLConversionConfiguration::isResourceUsed(uint32_t descSet, uint32_t binding) const {
+	for (auto& rb : resourceBindings) {
+		auto& rbb = rb.resourceBinding;
+		if (rbb.desc_set == descSet && rbb.binding == binding) {
+			return rb.isUsedByShader;
+		}
+	}
+	return false;
+}
+
 MVK_PUBLIC_SYMBOL MTLTextureType SPIRVToMSLConversionConfiguration::getMTLTextureType(uint32_t descSet, uint32_t binding) const {
 	for (auto& rb : resourceBindings) {
 		auto& rbb = rb.resourceBinding;