diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
index fa3bb7f..7d9c92f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
@@ -52,13 +52,7 @@
 /** Context for tracking information across multiple encodings. */
 typedef struct MVKCommandEncodingContext {
 	NSUInteger mtlVisibilityResultOffset = 0;
-
-	void incrementMTLVisibilityResultOffset(MVKCommandEncoder* cmdEncoder);
-	const MVKMTLBufferAllocation* getVisibilityResultBuffer(MVKCommandEncoder* cmdEncoder);
-
-private:
-	const MVKMTLBufferAllocation* _visibilityResultBuffer = nullptr;
-
+	const MVKMTLBufferAllocation* visibilityResultBuffer = nullptr;
 } MVKCommandEncodingContext;
 
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
index 7225d93..18d2441 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -350,7 +350,10 @@
 												  _isRenderingEntireAttachment,
 												  loadOverride);
 	if (_cmdBuffer->_needsVisibilityResultMTLBuffer) {
-		mtlRPDesc.visibilityResultBuffer = _pEncodingContext->getVisibilityResultBuffer(this)->_mtlBuffer;
+		if ( !_pEncodingContext->visibilityResultBuffer ) {
+			_pEncodingContext->visibilityResultBuffer = getTempMTLBuffer(_pDeviceMetalFeatures->maxQueryBufferSize, true, true);
+		}
+		mtlRPDesc.visibilityResultBuffer = _pEncodingContext->visibilityResultBuffer->_mtlBuffer;
 	}
 
 	VkExtent2D fbExtent = _framebufferExtent;
@@ -804,32 +807,6 @@
 
 
 #pragma mark -
-#pragma mark MVKCommandEncodingContext
-
-// Increment to the next query slot offset. If we reach the size of the visibility buffer,
-// reset to retrieve and start filling another visibility buffer. This approach may still
-// cause Metal validation errors if the platform does not permit offsets to be reused
-// witin a MTLCommandBuffer, even when a different visibility buffer is used.
-// We don't test against the size of the visibility buffer itself, because this call may
-// arrive before the visibiltiy buffer in the case of a query that ends before the renderpass.
-void MVKCommandEncodingContext::incrementMTLVisibilityResultOffset(MVKCommandEncoder* cmdEncoder) {
-	mtlVisibilityResultOffset += kMVKQuerySlotSizeInBytes;
-
-	if (mtlVisibilityResultOffset + kMVKQuerySlotSizeInBytes > cmdEncoder->_pDeviceMetalFeatures->maxQueryBufferSize) {
-		_visibilityResultBuffer = nullptr;
-		mtlVisibilityResultOffset = 0;
-	}
-}
-
-const MVKMTLBufferAllocation* MVKCommandEncodingContext::getVisibilityResultBuffer(MVKCommandEncoder* cmdEncoder) {
-	if ( !_visibilityResultBuffer ) {
-		_visibilityResultBuffer = cmdEncoder->getTempMTLBuffer(cmdEncoder->_pDeviceMetalFeatures->maxQueryBufferSize, true, true);
-	}
-	return _visibilityResultBuffer;
-}
-
-
-#pragma mark -
 #pragma mark Support functions
 
 NSString* mvkMTLCommandBufferLabel(MVKCommandUse cmdUse) {
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
index 2f9312b..8093cab 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
@@ -1073,10 +1073,9 @@
 #pragma mark MVKOcclusionQueryCommandEncoderState
 
 void MVKOcclusionQueryCommandEncoderState::endMetalRenderPass() {
+	const MVKMTLBufferAllocation* vizBuff = _cmdEncoder->_pEncodingContext->visibilityResultBuffer;
+    if ( !vizBuff || _mtlRenderPassQueries.empty() ) { return; }  // Nothing to do.
 
-    if (_mtlRenderPassQueries.empty()) { return; }  // Nothing to do.
-
-	const MVKMTLBufferAllocation* vizResultBuffer = _cmdEncoder->_pEncodingContext->getVisibilityResultBuffer(_cmdEncoder);
 	id<MTLComputePipelineState> mtlAccumState = _cmdEncoder->getCommandEncodingPool()->getAccumulateOcclusionQueryResultsMTLComputePipelineState();
     id<MTLComputeCommandEncoder> mtlAccumEncoder = _cmdEncoder->getMTLComputeEncoder(kMVKCommandUseAccumOcclusionQuery);
     [mtlAccumEncoder setComputePipelineState: mtlAccumState];
@@ -1085,8 +1084,8 @@
         [mtlAccumEncoder setBuffer: qryLoc.queryPool->getVisibilityResultMTLBuffer()
                             offset: qryLoc.queryPool->getVisibilityResultOffset(qryLoc.query)
                            atIndex: 0];
-        [mtlAccumEncoder setBuffer: vizResultBuffer->_mtlBuffer
-                            offset: vizResultBuffer->_offset + qryLoc.visibilityBufferOffset
+        [mtlAccumEncoder setBuffer: vizBuff->_mtlBuffer
+                            offset: vizBuff->_offset + qryLoc.visibilityBufferOffset
                            atIndex: 1];
         [mtlAccumEncoder dispatchThreadgroups: MTLSizeMake(1, 1, 1)
                         threadsPerThreadgroup: MTLSizeMake(1, 1, 1)];
@@ -1095,16 +1094,28 @@
     _mtlRenderPassQueries.clear();
 }
 
+// The Metal visibility buffer has a finite size, and on some Metal platforms (looking at you M1),
+// query offsets cannnot be reused with the same MTLCommandBuffer. If enough occlusion queries are
+// begun within a single MTLCommandBuffer, it may exhaust the visibility buffer. If that occurs,
+// report an error and disable further visibility tracking for the remainder of the MTLCommandBuffer.
+// In most cases, a MTLCommandBuffer corresponds to a Vulkan command submit (VkSubmitInfo),
+// and so the error text is framed in terms of the Vulkan submit.
 void MVKOcclusionQueryCommandEncoderState::beginOcclusionQuery(MVKOcclusionQueryPool* pQueryPool, uint32_t query, VkQueryControlFlags flags) {
-    bool shouldCount = _cmdEncoder->_pDeviceFeatures->occlusionQueryPrecise && mvkAreAllFlagsEnabled(flags, VK_QUERY_CONTROL_PRECISE_BIT);
-    _mtlVisibilityResultMode = shouldCount ? MTLVisibilityResultModeCounting : MTLVisibilityResultModeBoolean;
-	_mtlRenderPassQueries.emplace_back(pQueryPool, query, _cmdEncoder->_pEncodingContext->mtlVisibilityResultOffset);
+	if (_cmdEncoder->_pEncodingContext->mtlVisibilityResultOffset + kMVKQuerySlotSizeInBytes <= _cmdEncoder->_pDeviceMetalFeatures->maxQueryBufferSize) {
+		bool shouldCount = _cmdEncoder->_pDeviceFeatures->occlusionQueryPrecise && mvkAreAllFlagsEnabled(flags, VK_QUERY_CONTROL_PRECISE_BIT);
+		_mtlVisibilityResultMode = shouldCount ? MTLVisibilityResultModeCounting : MTLVisibilityResultModeBoolean;
+		_mtlRenderPassQueries.emplace_back(pQueryPool, query, _cmdEncoder->_pEncodingContext->mtlVisibilityResultOffset);
+	} else {
+		reportError(VK_ERROR_OUT_OF_DEVICE_MEMORY, "vkCmdBeginQuery(): The maximum number of queries in a single Vulkan command submission is %llu.", _cmdEncoder->_pDeviceMetalFeatures->maxQueryBufferSize / kMVKQuerySlotSizeInBytes);
+		_mtlVisibilityResultMode = MTLVisibilityResultModeDisabled;
+		_cmdEncoder->_pEncodingContext->mtlVisibilityResultOffset -= kMVKQuerySlotSizeInBytes;
+	}
     markDirty();
 }
 
 void MVKOcclusionQueryCommandEncoderState::endOcclusionQuery(MVKOcclusionQueryPool* pQueryPool, uint32_t query) {
 	_mtlVisibilityResultMode = MTLVisibilityResultModeDisabled;
-	_cmdEncoder->_pEncodingContext->incrementMTLVisibilityResultOffset(_cmdEncoder);
+	_cmdEncoder->_pEncodingContext->mtlVisibilityResultOffset += kMVKQuerySlotSizeInBytes;
 	markDirty();
 }