diff --git a/Docs/MoltenVK_Runtime_UserGuide.md b/Docs/MoltenVK_Runtime_UserGuide.md
index 117b14e..3dbdf9c 100644
--- a/Docs/MoltenVK_Runtime_UserGuide.md
+++ b/Docs/MoltenVK_Runtime_UserGuide.md
@@ -273,6 +273,7 @@
 - `VK_KHR_get_memory_requirements2`
 - `VK_KHR_get_physical_device_properties2`
 - `VK_KHR_get_surface_capabilities2`
+- `VK_KHR_imageless_framebuffer`
 - `VK_KHR_image_format_list`
 - `VK_KHR_maintenance1`
 - `VK_KHR_maintenance2`
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
index 70c7ba2..8a66785 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
@@ -60,7 +60,7 @@
  * Vulkan command to begin a render pass.
  * Template class to balance vector pre-allocations between very common low counts and fewer larger counts.
  */
-template <size_t N>
+template <size_t N_CV, size_t N_A>
 class MVKCmdBeginRenderPass : public MVKCmdBeginRenderPassBase {
 
 public:
@@ -76,13 +76,26 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKSmallVector<VkClearValue, N> _clearValues;
+	MVKSmallVector<VkClearValue, N_CV> _clearValues;
+    MVKSmallVector<MVKImageView*, N_A> _attachments;
 };
 
 // Concrete template class implementations.
-typedef MVKCmdBeginRenderPass<1> MVKCmdBeginRenderPass1;
-typedef MVKCmdBeginRenderPass<2> MVKCmdBeginRenderPass2;
-typedef MVKCmdBeginRenderPass<9> MVKCmdBeginRenderPassMulti;
+typedef MVKCmdBeginRenderPass<1, 0> MVKCmdBeginRenderPass10;
+typedef MVKCmdBeginRenderPass<2, 0> MVKCmdBeginRenderPass20;
+typedef MVKCmdBeginRenderPass<9, 0> MVKCmdBeginRenderPassMulti0;
+
+typedef MVKCmdBeginRenderPass<1, 1> MVKCmdBeginRenderPass11;
+typedef MVKCmdBeginRenderPass<2, 1> MVKCmdBeginRenderPass21;
+typedef MVKCmdBeginRenderPass<9, 1> MVKCmdBeginRenderPassMulti1;
+
+typedef MVKCmdBeginRenderPass<1, 2> MVKCmdBeginRenderPass12;
+typedef MVKCmdBeginRenderPass<2, 2> MVKCmdBeginRenderPass22;
+typedef MVKCmdBeginRenderPass<9, 2> MVKCmdBeginRenderPassMulti2;
+
+typedef MVKCmdBeginRenderPass<1, 9> MVKCmdBeginRenderPass1Multi;
+typedef MVKCmdBeginRenderPass<2, 9> MVKCmdBeginRenderPass2Multi;
+typedef MVKCmdBeginRenderPass<9, 9> MVKCmdBeginRenderPassMultiMulti;
 
 
 #pragma mark -
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
index 7f523f1..4fc1146 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
@@ -19,6 +19,7 @@
 #include "MVKCmdRenderPass.h"
 #include "MVKCommandBuffer.h"
 #include "MVKCommandPool.h"
+#include "MVKFramebuffer.h"
 #include "MVKRenderPass.h"
 #include "MVKPipeline.h"
 #include "MVKFoundation.h"
@@ -43,10 +44,10 @@
 #pragma mark -
 #pragma mark MVKCmdBeginRenderPass
 
-template <size_t N>
-VkResult MVKCmdBeginRenderPass<N>::setContent(MVKCommandBuffer* cmdBuff,
-											  const VkRenderPassBeginInfo* pRenderPassBegin,
-											  VkSubpassContents contents) {
+template <size_t N_CV, size_t N_A>
+VkResult MVKCmdBeginRenderPass<N_CV, N_A>::setContent(MVKCommandBuffer* cmdBuff,
+													  const VkRenderPassBeginInfo* pRenderPassBegin,
+													  VkSubpassContents contents) {
 	MVKCmdBeginRenderPassBase::setContent(cmdBuff, pRenderPassBegin, contents);
 
 	// Add clear values
@@ -57,26 +58,66 @@
 		_clearValues.push_back(pRenderPassBegin->pClearValues[i]);
 	}
 
+	bool imageless = false;
+	for (auto* next = (const VkBaseInStructure*)pRenderPassBegin->pNext; next; next = next->pNext) {
+		switch (next->sType) {
+		case VK_STRUCTURE_TYPE_RENDER_PASS_ATTACHMENT_BEGIN_INFO: {
+			const auto* pAttachmentBegin = (VkRenderPassAttachmentBeginInfo*)next;
+			for(uint32_t i = 0; i < pAttachmentBegin->attachmentCount; i++) {
+				_attachments.push_back((MVKImageView*)pAttachmentBegin->pAttachments[i]);
+			}
+			imageless = true;
+			break;
+		}
+		default:
+			break;
+		}
+	}
+	
+	if (!imageless) {
+		for(uint32_t i = 0; i < _framebuffer->getAttachmentCount(); i++) {
+			_attachments.push_back((MVKImageView*)_framebuffer->getAttachment(i));
+		}
+	}
+
 	return VK_SUCCESS;
 }
 
-template <size_t N>
-VkResult MVKCmdBeginRenderPass<N>::setContent(MVKCommandBuffer* cmdBuff,
-											  const VkRenderPassBeginInfo* pRenderPassBegin,
-											  const VkSubpassBeginInfo* pSubpassBeginInfo) {
+template <size_t N_CV, size_t N_A>
+VkResult MVKCmdBeginRenderPass<N_CV, N_A>::setContent(MVKCommandBuffer* cmdBuff,
+													  const VkRenderPassBeginInfo* pRenderPassBegin,
+													  const VkSubpassBeginInfo* pSubpassBeginInfo) {
 	return setContent(cmdBuff, pRenderPassBegin, pSubpassBeginInfo->contents);
 }
 
-template <size_t N>
-void MVKCmdBeginRenderPass<N>::encode(MVKCommandEncoder* cmdEncoder) {
+template <size_t N_CV, size_t N_A>
+void MVKCmdBeginRenderPass<N_CV, N_A>::encode(MVKCommandEncoder* cmdEncoder) {
 //	MVKLogDebug("Encoding vkCmdBeginRenderPass(). Elapsed time: %.6f ms.", mvkGetElapsedMilliseconds());
-	cmdEncoder->beginRenderpass(this, _contents, _renderPass, _framebuffer, _renderArea, _clearValues.contents());
+	cmdEncoder->beginRenderpass(this,
+								_contents,
+								_renderPass,
+								_framebuffer->getExtent2D(),
+								_framebuffer->getLayerCount(),
+								_renderArea,
+								_clearValues.contents(),
+								_attachments.contents());
 }
 
-template class MVKCmdBeginRenderPass<1>;
-template class MVKCmdBeginRenderPass<2>;
-template class MVKCmdBeginRenderPass<9>;
+template class MVKCmdBeginRenderPass<1, 0>;
+template class MVKCmdBeginRenderPass<2, 0>;
+template class MVKCmdBeginRenderPass<9, 0>;
 
+template class MVKCmdBeginRenderPass<1, 1>;
+template class MVKCmdBeginRenderPass<2, 1>;
+template class MVKCmdBeginRenderPass<9, 1>;
+
+template class MVKCmdBeginRenderPass<1, 2>;
+template class MVKCmdBeginRenderPass<2, 2>;
+template class MVKCmdBeginRenderPass<9, 2>;
+
+template class MVKCmdBeginRenderPass<1, 9>;
+template class MVKCmdBeginRenderPass<2, 9>;
+template class MVKCmdBeginRenderPass<9, 9>;
 
 #pragma mark -
 #pragma mark MVKCmdNextSubpass
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
index df50a7d..7d1d379 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
@@ -1236,7 +1236,7 @@
 	simd::float4 vertices[vtxCnt];
 	simd::float4 clearColors[kMVKClearAttachmentCount];
 
-	VkExtent2D fbExtent = cmdEncoder->_framebuffer->getExtent2D();
+	VkExtent2D fbExtent = cmdEncoder->_framebufferExtent;
 #if MVK_MACOS_OR_IOS
 	// I need to know if the 'renderTargetWidth' and 'renderTargetHeight' properties
 	// actually do something, but [MTLRenderPassDescriptor instancesRespondToSelector: @selector(renderTargetWidth)]
@@ -1257,7 +1257,7 @@
     // Populate the render pipeline state attachment key with info from the subpass and framebuffer.
 	_rpsKey.mtlSampleCount = mvkSampleCountFromVkSampleCountFlagBits(subpass->getSampleCount());
 	if (cmdEncoder->_canUseLayeredRendering &&
-		(cmdEncoder->_framebuffer->getLayerCount() > 1 || cmdEncoder->getSubpass()->isMultiview())) {
+		(cmdEncoder->_framebufferLayerCount > 1 || cmdEncoder->getSubpass()->isMultiview())) {
 		_rpsKey.enableLayeredRendering();
 	}
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
index 4829b93..952c7ec 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
@@ -273,9 +273,11 @@
 	void beginRenderpass(MVKCommand* passCmd,
 						 VkSubpassContents subpassContents,
 						 MVKRenderPass* renderPass,
-						 MVKFramebuffer* framebuffer,
+						 VkExtent2D framebufferExtent,
+						 uint32_t framebufferLayerCount,
 						 VkRect2D& renderArea,
-						 MVKArrayRef<VkClearValue> clearValues);
+						 MVKArrayRef<VkClearValue> clearValues,
+						 MVKArrayRef<MVKImageView*> attachments);
 
 	/** Begins the next render subpass. */
 	void beginNextSubpass(MVKCommand* subpassCmd, VkSubpassContents renderpassContents);
@@ -420,9 +422,6 @@
 	/** The command buffer whose commands are being encoded. */
 	MVKCommandBuffer* _cmdBuffer;
 
-	/** The framebuffer to which rendering is currently directed. */
-	MVKFramebuffer* _framebuffer;
-
 	/** The current Metal command buffer. */
 	id<MTLCommandBuffer> _mtlCmdBuffer;
 
@@ -474,6 +473,11 @@
 	/** Indicates whether the current draw is an indexed draw. */
 	bool _isIndexedDraw;
 
+	/** The extent of current framebuffer.*/
+	VkExtent2D _framebufferExtent;
+
+	/** The layer count of current framebuffer.*/
+	uint32_t _framebufferLayerCount;
 
 #pragma mark Construction
 
@@ -494,6 +498,7 @@
 	VkRect2D _renderArea;
     MVKActivatedQueries* _pActivatedQueries;
 	MVKSmallVector<VkClearValue, kMVKDefaultAttachmentCount> _clearValues;
+	MVKSmallVector<MVKImageView*, kMVKDefaultAttachmentCount> _attachments;
 	id<MTLComputeCommandEncoder> _mtlComputeEncoder;
 	MVKCommandUse _mtlComputeEncoderUse;
 	id<MTLBlitCommandEncoder> _mtlBlitEncoder;
@@ -523,4 +528,3 @@
 
 /** Returns a name, suitable for use as a MTLComputeCommandEncoder label, based on the MVKCommandUse. */
 NSString* mvkMTLComputeCommandEncoderLabel(MVKCommandUse cmdUse);
-
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
index 352c475..801b3f7 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -20,7 +20,6 @@
 #include "MVKCommandPool.h"
 #include "MVKQueue.h"
 #include "MVKPipeline.h"
-#include "MVKFramebuffer.h"
 #include "MVKQueryPool.h"
 #include "MVKFoundation.h"
 #include "MTLRenderPassDescriptor+MoltenVK.h"
@@ -286,15 +285,19 @@
 void MVKCommandEncoder::beginRenderpass(MVKCommand* passCmd,
 										VkSubpassContents subpassContents,
 										MVKRenderPass* renderPass,
-										MVKFramebuffer* framebuffer,
+										VkExtent2D framebufferExtent,
+										uint32_t framebufferLayerCount,
 										VkRect2D& renderArea,
-										MVKArrayRef<VkClearValue> clearValues) {
+										MVKArrayRef<VkClearValue> clearValues,
+										MVKArrayRef<MVKImageView*> attachments) {
 	_renderPass = renderPass;
-	_framebuffer = framebuffer;
+	_framebufferExtent = framebufferExtent;
+	_framebufferLayerCount = framebufferLayerCount;
 	_renderArea = renderArea;
 	_isRenderingEntireAttachment = (mvkVkOffset2DsAreEqual(_renderArea.offset, {0,0}) &&
-									mvkVkExtent2DsAreEqual(_renderArea.extent, _framebuffer->getExtent2D()));
+									mvkVkExtent2DsAreEqual(_renderArea.extent, _framebufferExtent));
 	_clearValues.assign(clearValues.begin(), clearValues.end());
+	_attachments.assign(attachments.begin(), attachments.end());
 	setSubpass(passCmd, subpassContents, 0);
 }
 
@@ -334,7 +337,14 @@
     endCurrentMetalEncoding();
 
     MTLRenderPassDescriptor* mtlRPDesc = [MTLRenderPassDescriptor renderPassDescriptor];
-    getSubpass()->populateMTLRenderPassDescriptor(mtlRPDesc, _multiviewPassIndex, _framebuffer, _clearValues.contents(), _isRenderingEntireAttachment, loadOverride);
+	getSubpass()->populateMTLRenderPassDescriptor(mtlRPDesc,
+												  _multiviewPassIndex,
+												  _framebufferExtent,
+												  _framebufferLayerCount,
+												  _attachments.contents(),
+												  _clearValues.contents(),
+												  _isRenderingEntireAttachment,
+												  loadOverride);
     if (_cmdBuffer->_needsVisibilityResultMTLBuffer) {
         if (!_visibilityResultMTLBuffer) {
             _visibilityResultMTLBuffer = getTempMTLBuffer(_pDeviceMetalFeatures->maxQueryBufferSize, true, true);
@@ -342,7 +352,7 @@
         mtlRPDesc.visibilityResultBuffer = _visibilityResultMTLBuffer->_mtlBuffer;
     }
 
-    VkExtent2D fbExtent = _framebuffer->getExtent2D();
+	VkExtent2D fbExtent = _framebufferExtent;
     mtlRPDesc.renderTargetWidthMVK = max(min(_renderArea.offset.x + _renderArea.extent.width, fbExtent.width), 1u);
     mtlRPDesc.renderTargetHeightMVK = max(min(_renderArea.offset.y + _renderArea.extent.height, fbExtent.height), 1u);
     if (_canUseLayeredRendering) {
@@ -362,9 +372,9 @@
         if (getSubpass()->isMultiview()) {
             // In the case of a multiview pass, the framebuffer layer count will be one.
             // We need to use the view count for this multiview pass.
-            renderTargetArrayLength = getSubpass()->getViewCountInMetalPass(_multiviewPassIndex);
+			renderTargetArrayLength = getSubpass()->getViewCountInMetalPass(_multiviewPassIndex);
         } else {
-            renderTargetArrayLength = _framebuffer->getLayerCount();
+			renderTargetArrayLength = _framebufferLayerCount;
         }
         // Metal does not allow layered render passes where some RTs are 3D and others are 2D.
         if (!(found3D && found2D) || renderTargetArrayLength > 1) {
@@ -393,7 +403,10 @@
 }
 
 void MVKCommandEncoder::encodeStoreActions(bool storeOverride) {
-	getSubpass()->encodeStoreActions(this, _isRenderingEntireAttachment, storeOverride);
+	getSubpass()->encodeStoreActions(this,
+									 _isRenderingEntireAttachment,
+									 _attachments.contents(),
+									 storeOverride);
 }
 
 MVKRenderSubpass* MVKCommandEncoder::getSubpass() { return _renderPass->getSubpass(_renderSubpassIndex); }
@@ -508,7 +521,7 @@
 		VkClearRect clearRect;
 		clearRect.rect = _renderArea;
 		clearRect.baseArrayLayer = 0;
-		clearRect.layerCount = _framebuffer->getLayerCount();
+		clearRect.layerCount = _framebufferLayerCount;
 
 		// Create and execute a temporary clear attachments command.
 		// To be threadsafe...do NOT acquire and return the command from the pool.
@@ -555,7 +568,9 @@
 	endMetalRenderEncoding();
 
 	_renderPass = nullptr;
-	_framebuffer = nullptr;
+	_framebufferExtent = {};
+	_framebufferLayerCount = 0;
+	_attachments.clear();
 	_renderSubpassIndex = 0;
 }
 
@@ -849,4 +864,3 @@
         default:                                            return @"Unknown Use ComputeEncoder";
     }
 }
-
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def b/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def
index 26009aa..a6966f4 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandTypePools.def
@@ -56,11 +56,26 @@
 	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##threshold3)									\
 	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##Multi)
 
+#define MVK_CMD_TYPE_POOLS_FROM_5_THRESHOLDS(cmdType,											\
+											 arg1Threshold1, arg1Threshold2,					\
+											 arg2Threshold1, arg2Threshold2, arg2Threshold3)	\
+    MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold1 ##arg2Threshold1)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold1 ##arg2Threshold2)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold1 ##arg2Threshold3)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold1 ##Multi)         					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold2 ##arg2Threshold1)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold2 ##arg2Threshold2)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold2 ##arg2Threshold3)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##arg1Threshold2 ##Multi)         					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##Multi          ##arg2Threshold1)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##Multi          ##arg2Threshold2)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##Multi          ##arg2Threshold3)					\
+	MVK_TMPLT_DECL MVK_CMD_TYPE_POOL(cmdType ##Multi          ##Multi)
 
 MVK_CMD_TYPE_POOLS_FROM_2_THRESHOLDS(PipelineBarrier, 1, 4)
 MVK_CMD_TYPE_POOL(BindGraphicsPipeline)
 MVK_CMD_TYPE_POOL(BindComputePipeline)
-MVK_CMD_TYPE_POOLS_FROM_2_THRESHOLDS(BeginRenderPass, 1, 2)
+MVK_CMD_TYPE_POOLS_FROM_5_THRESHOLDS(BeginRenderPass, 1, 2, 0, 1, 2)
 MVK_CMD_TYPE_POOL(NextSubpass)
 MVK_CMD_TYPE_POOL(EndRenderPass)
 MVK_CMD_TYPE_POOLS_FROM_THRESHOLD(ExecuteCommands, 1)
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
index b1bd357..8eef871 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
@@ -738,6 +738,7 @@
 	const VkPhysicalDeviceTexelBufferAlignmentFeaturesEXT _enabledTexelBuffAlignFeatures;
 	const VkPhysicalDeviceVertexAttributeDivisorFeaturesEXT _enabledVtxAttrDivFeatures;
 	const VkPhysicalDevicePortabilitySubsetFeaturesKHR _enabledPortabilityFeatures;
+	const VkPhysicalDeviceImagelessFramebufferFeaturesKHR _enabledImagelessFramebufferFeatures;
 
 	/** The list of Vulkan extensions, indicating whether each has been enabled by the app for this device. */
 	const MVKExtensionList _enabledExtensions;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
index 40ee051..6ca2971 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -267,6 +267,11 @@
 				inlineUniformBlockFeatures->descriptorBindingInlineUniformBlockUpdateAfterBind = true;
 				break;
 			}
+			case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGELESS_FRAMEBUFFER_FEATURES: {
+				auto* imagelessFramebufferFeatures = (VkPhysicalDeviceImagelessFramebufferFeaturesKHR*)next;
+				imagelessFramebufferFeatures->imagelessFramebuffer = true;
+				break;
+			}
 			default:
 				break;
 		}
@@ -3789,6 +3794,7 @@
 	_enabledVtxAttrDivFeatures(),
 	_enabledPrivateDataFeatures(),
 	_enabledPortabilityFeatures(),
+	_enabledImagelessFramebufferFeatures(),
 	_enabledExtensions(this),
 	_isCurrentlyAutoGPUCapturing(false)
 {
@@ -3899,11 +3905,16 @@
 	mvkClear(&_enabledTexelBuffAlignFeatures);
 	mvkClear(&_enabledVtxAttrDivFeatures);
 	mvkClear(&_enabledPortabilityFeatures);
+	mvkClear(&_enabledImagelessFramebufferFeatures);
 
+	VkPhysicalDeviceImagelessFramebufferFeaturesKHR pdImagelessFramebufferFeatures;
+	pdImagelessFramebufferFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGELESS_FRAMEBUFFER_FEATURES;
+	pdImagelessFramebufferFeatures.pNext = NULL;
+    
 	// Fetch the available physical device features.
 	VkPhysicalDevicePortabilitySubsetFeaturesKHR pdPortabilityFeatures;
 	pdPortabilityFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_KHR;
-	pdPortabilityFeatures.pNext = NULL;
+	pdPortabilityFeatures.pNext = &pdImagelessFramebufferFeatures;
 
 	VkPhysicalDeviceVertexAttributeDivisorFeaturesEXT pdVtxAttrDivFeatures;
 	pdVtxAttrDivFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VERTEX_ATTRIBUTE_DIVISOR_FEATURES_EXT;
@@ -4088,6 +4099,13 @@
 							   &pdPortabilityFeatures.constantAlphaColorBlendFactors, 15);
 				break;
 			}
+			case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGELESS_FRAMEBUFFER_FEATURES: {
+				auto* requestedFeatures = (VkPhysicalDeviceImagelessFramebufferFeaturesKHR*)next;
+				enableFeatures(&_enabledImagelessFramebufferFeatures.imagelessFramebuffer,
+							   &requestedFeatures->imagelessFramebuffer,
+							   &pdImagelessFramebufferFeatures.imagelessFramebuffer, 1);
+				break;
+			}
 			default:
 				break;
 		}
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
index 16a2a22..32cd5dd 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
@@ -44,7 +44,8 @@
 
 	/** Returns the attachment at the specified index.  */
 	inline MVKImageView* getAttachment(uint32_t index) { return _attachments[index]; }
-
+	
+	inline size_t getAttachmentCount() {return _attachments.size(); }
 
 #pragma mark Construction
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.mm b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.mm
index 66af186..91d3bde 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.mm
@@ -28,10 +28,11 @@
     _extent = { .width = pCreateInfo->width, .height = pCreateInfo->height };
 	_layerCount = pCreateInfo->layers;
 
-	// Add attachments
-	_attachments.reserve(pCreateInfo->attachmentCount);
-	for (uint32_t i = 0; i < pCreateInfo->attachmentCount; i++) {
-		_attachments.push_back((MVKImageView*)pCreateInfo->pAttachments[i]);
+	if (!(pCreateInfo->flags & VK_FRAMEBUFFER_CREATE_IMAGELESS_BIT_KHR)) {
+		// Add attachments
+		_attachments.reserve(pCreateInfo->attachmentCount);
+		for (uint32_t i = 0; i < pCreateInfo->attachmentCount; i++) {
+			_attachments.push_back((MVKImageView*)pCreateInfo->pAttachments[i]);
+		}
 	}
 }
-
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
index fdb9077..c813fa6 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
@@ -95,7 +95,9 @@
 	 */
 	void populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* mtlRPDesc,
 										 uint32_t passIdx,
-										 MVKFramebuffer* framebuffer,
+										 VkExtent2D framebufferExtent,
+										 uint32_t framebufferLayerCount,
+										 const MVKArrayRef<MVKImageView*>& attachments,
 										 const MVKArrayRef<VkClearValue>& clearValues,
 										 bool isRenderingEntireAttachment,
                                          bool loadOverride = false);
@@ -117,7 +119,10 @@
 									 uint32_t caIdx, VkImageAspectFlags aspectMask);
 
 	/** If a render encoder is active, sets the store actions for all attachments to it. */
-	void encodeStoreActions(MVKCommandEncoder* cmdEncoder, bool isRenderingEntireAttachment, bool storeOverride = false);
+	void encodeStoreActions(MVKCommandEncoder* cmdEncoder,
+							bool isRenderingEntireAttachment,
+							const MVKArrayRef<MVKImageView*>& attachments,
+							bool storeOverride = false);
 
 	/** Constructs an instance for the specified parent renderpass. */
 	MVKRenderSubpass(MVKRenderPass* renderPass, const VkSubpassDescription* pCreateInfo,
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
index 4a09569..0443597 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
@@ -175,7 +175,9 @@
 
 void MVKRenderSubpass::populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* mtlRPDesc,
 													   uint32_t passIdx,
-													   MVKFramebuffer* framebuffer,
+													   VkExtent2D framebufferExtent,
+													   uint32_t framebufferLayerCount,
+													   const MVKArrayRef<MVKImageView*>& attachments,
 													   const MVKArrayRef<VkClearValue>& clearValues,
 													   bool isRenderingEntireAttachment,
 													   bool loadOverride) {
@@ -195,7 +197,8 @@
             uint32_t rslvRPAttIdx = _resolveAttachments.empty() ? VK_ATTACHMENT_UNUSED : _resolveAttachments[caIdx].attachment;
             bool hasResolveAttachment = (rslvRPAttIdx != VK_ATTACHMENT_UNUSED);
             if (hasResolveAttachment) {
-                framebuffer->getAttachment(rslvRPAttIdx)->populateMTLRenderPassAttachmentDescriptorResolve(mtlColorAttDesc);
+				attachments[rslvRPAttIdx]->populateMTLRenderPassAttachmentDescriptorResolve(mtlColorAttDesc);
+
 				// In a multiview render pass, we need to override the starting layer to ensure
 				// only the enabled views are loaded.
 				if (isMultiview()) {
@@ -209,10 +212,10 @@
 
             // Configure the color attachment
             MVKRenderPassAttachment* clrMVKRPAtt = &_renderPass->_attachments[clrRPAttIdx];
-			framebuffer->getAttachment(clrRPAttIdx)->populateMTLRenderPassAttachmentDescriptor(mtlColorAttDesc);
+			attachments[clrRPAttIdx]->populateMTLRenderPassAttachmentDescriptor(mtlColorAttDesc);
 			bool isMemorylessAttachment = false;
 #if MVK_APPLE_SILICON
-			isMemorylessAttachment = framebuffer->getAttachment(clrRPAttIdx)->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
+			isMemorylessAttachment = attachments[clrRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
 #endif
 			if (clrMVKRPAtt->populateMTLRenderPassAttachmentDescriptor(mtlColorAttDesc, this,
                                                                        isRenderingEntireAttachment,
@@ -236,12 +239,12 @@
 	uint32_t dsRslvRPAttIdx = _depthStencilResolveAttachment.attachment;
 	if (dsRPAttIdx != VK_ATTACHMENT_UNUSED) {
 		MVKRenderPassAttachment* dsMVKRPAtt = &_renderPass->_attachments[dsRPAttIdx];
-		MVKImageView* dsImage = framebuffer->getAttachment(dsRPAttIdx);
+		MVKImageView* dsImage = attachments[dsRPAttIdx];
 		MVKImageView* dsRslvImage = nullptr;
 		MTLPixelFormat mtlDSFormat = dsImage->getMTLPixelFormat(0);
 
 		if (dsRslvRPAttIdx != VK_ATTACHMENT_UNUSED) {
-			dsRslvImage = framebuffer->getAttachment(dsRslvRPAttIdx);
+			dsRslvImage = attachments[dsRslvRPAttIdx];
 		}
 
 		if (pixFmts->isDepthFormat(mtlDSFormat)) {
@@ -312,7 +315,7 @@
         }
 
 		// Add a dummy attachment so this passes validation.
-		VkExtent2D fbExtent = framebuffer->getExtent2D();
+		VkExtent2D fbExtent = framebufferExtent;
 		MTLTextureDescriptor* mtlTexDesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat: MTLPixelFormatR8Unorm width: fbExtent.width height: fbExtent.height mipmapped: NO];
 		if (isMultiview()) {
 #if MVK_MACOS_OR_IOS
@@ -326,7 +329,7 @@
 			mtlTexDesc.textureType = MTLTextureType2DArray;
 #endif
 			mtlTexDesc.arrayLength = getViewCountInMetalPass(passIdx);
-		} else if (framebuffer->getLayerCount() > 1) {
+		} else if (framebufferLayerCount > 1) {
 #if MVK_MACOS
 			if (sampleCount > 1 && _renderPass->getDevice()->_pMetalFeatures->multisampleLayeredRendering) {
 				mtlTexDesc.textureType = MTLTextureType2DMultisampleArray;
@@ -337,7 +340,7 @@
 #else
 			mtlTexDesc.textureType = MTLTextureType2DArray;
 #endif
-			mtlTexDesc.arrayLength = framebuffer->getLayerCount();
+			mtlTexDesc.arrayLength = framebufferLayerCount;
 		} else if (sampleCount > 1) {
 			mtlTexDesc.textureType = MTLTextureType2DMultisample;
 			mtlTexDesc.sampleCount = sampleCount;
@@ -366,6 +369,7 @@
 
 void MVKRenderSubpass::encodeStoreActions(MVKCommandEncoder* cmdEncoder,
                                           bool isRenderingEntireAttachment,
+										  const MVKArrayRef<MVKImageView*>& attachments,
                                           bool storeOverride) {
     if (!cmdEncoder->_mtlRenderEncoder) { return; }
 	if (!_renderPass->getDevice()->_pMetalFeatures->deferredStoreActions) { return; }
@@ -377,7 +381,7 @@
             bool hasResolveAttachment = _resolveAttachments.empty() ? false : _resolveAttachments[caIdx].attachment != VK_ATTACHMENT_UNUSED;
 			bool isMemorylessAttachment = false;
 #if MVK_APPLE_SILICON
-			isMemorylessAttachment = cmdEncoder->_framebuffer->getAttachment(clrRPAttIdx)->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
+			isMemorylessAttachment = attachments[clrRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
 #endif
             _renderPass->_attachments[clrRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasResolveAttachment, caIdx, false, storeOverride);
         }
@@ -389,7 +393,7 @@
         bool hasStencilResolveAttachment = hasResolveAttachment && _stencilResolveMode != VK_RESOLVE_MODE_NONE;
 		bool isMemorylessAttachment = false;
 #if MVK_APPLE_SILICON
-		isMemorylessAttachment = cmdEncoder->_framebuffer->getAttachment(dsRPAttIdx)->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
+		isMemorylessAttachment = attachments[dsRPAttIdx]->getMTLTexture(0).storageMode == MTLStorageModeMemoryless;
 #endif
         _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasDepthResolveAttachment, 0, false, storeOverride);
         _renderPass->_attachments[dsRPAttIdx].encodeStoreAction(cmdEncoder, this, isRenderingEntireAttachment, isMemorylessAttachment, hasStencilResolveAttachment, 0, true, storeOverride);
diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
index 5fb1c18..a80de24 100644
--- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def
+++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
@@ -58,6 +58,7 @@
 MVK_EXTENSION(KHR_get_memory_requirements2, KHR_GET_MEMORY_REQUIREMENTS_2, DEVICE)
 MVK_EXTENSION(KHR_get_physical_device_properties2, KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2, INSTANCE)
 MVK_EXTENSION(KHR_get_surface_capabilities2, KHR_GET_SURFACE_CAPABILITIES_2, INSTANCE)
+MVK_EXTENSION(KHR_imageless_framebuffer, KHR_IMAGELESS_FRAMEBUFFER, DEVICE)
 MVK_EXTENSION(KHR_image_format_list, KHR_IMAGE_FORMAT_LIST, DEVICE)
 MVK_EXTENSION(KHR_maintenance1, KHR_MAINTENANCE1, DEVICE)
 MVK_EXTENSION(KHR_maintenance2, KHR_MAINTENANCE2, DEVICE)
diff --git a/MoltenVK/MoltenVK/Utility/MVKSmallVector.h b/MoltenVK/MoltenVK/Utility/MVKSmallVector.h
index 4648d19..e20e7fa 100755
--- a/MoltenVK/MoltenVK/Utility/MVKSmallVector.h
+++ b/MoltenVK/MoltenVK/Utility/MVKSmallVector.h
@@ -781,6 +781,18 @@
     alc.num_elements_used = new_size;
   }
 
+  template <class InputIterator>
+  void assign( InputIterator first, InputIterator last )
+  {
+    clear();
+
+    while( first != last )
+    {
+      push_back( *first );
+      ++first;
+    }
+  }
+
   void resize( const size_t new_size, const Type *t = nullptr )
   {
     if ( new_size == alc.num_elements_used )
diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
index 25eb9d6..e595e7e 100644
--- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm
+++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
@@ -131,6 +131,36 @@
 		MVKAddCmd(baseCmdType ##Multi, vkCmdBuff, ##__VA_ARGS__);											\
 	}
 
+// Add one of nine commands, based on comparing a command parameter against four threshold values
+#define MVKAddCmdFrom5Thresholds(baseCmdType, value1, arg1Threshold1, arg1Threshold2,			\
+								 value2, arg2Threshold1, arg2Threshold2, arg2Threshold3,		\
+								 vkCmdBuff, ...)												\
+	if (value1 <= arg1Threshold1 && value2 <= arg2Threshold1) {									\
+		MVKAddCmd(baseCmdType ##arg1Threshold1 ##arg2Threshold1, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 <= arg1Threshold2 && value2 <= arg2Threshold1) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold1 ##arg2Threshold1, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 > arg1Threshold2 && value2 <= arg2Threshold1) {							\
+		MVKAddCmd(baseCmdType ##Multi ##arg2Threshold1, vkCmdBuff, ##__VA_ARGS__);				\
+	} else if (value1 <= arg1Threshold1 && value2 <= arg2Threshold2) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold1 ##arg2Threshold2, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 <= arg1Threshold2 && value2 <= arg2Threshold2) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold2 ##arg2Threshold2, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 > arg1Threshold2 && value2 <= arg2Threshold2) {							\
+		MVKAddCmd(baseCmdType ##Multi ##arg2Threshold2, vkCmdBuff, ##__VA_ARGS__);				\
+	} else if (value1 <= arg1Threshold1 && value2 <= arg2Threshold3) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold1 ##arg2Threshold3, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 <= arg1Threshold2 && value2 <= arg2Threshold3) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold2 ##arg2Threshold3, vkCmdBuff, ##__VA_ARGS__);		\
+	} else if (value1 > arg1Threshold2 && value2 <= arg2Threshold3) {							\
+		MVKAddCmd(baseCmdType ##Multi ##arg2Threshold3, vkCmdBuff, ##__VA_ARGS__);				\
+	} else if (value1 <= arg1Threshold1 && value2 > arg2Threshold3) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold1 ##Multi, vkCmdBuff, ##__VA_ARGS__);				\
+	} else if (value1 <= arg1Threshold2 && value2 > arg2Threshold3) {							\
+		MVKAddCmd(baseCmdType ##arg1Threshold2 ##Multi, vkCmdBuff, ##__VA_ARGS__);				\
+	} else {																					\
+		MVKAddCmd(baseCmdType ##Multi ##Multi, vkCmdBuff, ##__VA_ARGS__);						\
+	}
+
 // Define an extension call as an alias of a core call
 #define MVK_PUBLIC_CORE_ALIAS(vkf)	MVK_PUBLIC_ALIAS(vkf##KHR, vkf)
 
@@ -1869,7 +1899,24 @@
     VkSubpassContents							contents) {
 	
 	MVKTraceVulkanCallStart();
-	MVKAddCmdFrom2Thresholds(BeginRenderPass, pRenderPassBegin->clearValueCount, 1, 2, commandBuffer,pRenderPassBegin, contents);
+	uint32_t attachmentCount = 0;
+	for (const auto* next = (VkBaseInStructure*)pRenderPassBegin->pNext; next; next = next->pNext) {
+		switch(next->sType) {
+			case VK_STRUCTURE_TYPE_RENDER_PASS_ATTACHMENT_BEGIN_INFO: {
+				auto* pAttachmentBegin = (VkRenderPassAttachmentBeginInfo*)next;
+				attachmentCount = pAttachmentBegin->attachmentCount;
+				break;
+			}
+			default:
+				break;
+		}
+	}
+	MVKAddCmdFrom5Thresholds(BeginRenderPass,
+							 pRenderPassBegin->clearValueCount, 1, 2,
+							 attachmentCount, 0, 1, 2,
+							 commandBuffer,
+							 pRenderPassBegin,
+							 contents);
 	MVKTraceVulkanCallEnd();
 }
 
@@ -2283,7 +2330,24 @@
 	const VkSubpassBeginInfo*					pSubpassBeginInfo) {
 
 	MVKTraceVulkanCallStart();
-	MVKAddCmdFrom2Thresholds(BeginRenderPass, pRenderPassBegin->clearValueCount, 1, 2, commandBuffer, pRenderPassBegin, pSubpassBeginInfo);
+	uint32_t attachmentCount = 0;
+	for (const auto* next = (VkBaseInStructure*)pRenderPassBegin->pNext; next; next = next->pNext) {
+		switch(next->sType) {
+			case VK_STRUCTURE_TYPE_RENDER_PASS_ATTACHMENT_BEGIN_INFO: {
+				auto* pAttachmentBegin = (VkRenderPassAttachmentBeginInfo*)next;
+				attachmentCount = pAttachmentBegin->attachmentCount;
+				break;
+			}
+			default:
+				break;
+		}
+	}
+	MVKAddCmdFrom5Thresholds(BeginRenderPass,
+							 pRenderPassBegin->clearValueCount, 1, 2,
+							 attachmentCount, 0, 1, 2,
+							 commandBuffer,
+							 pRenderPassBegin,
+							 pSubpassBeginInfo);
 	MVKTraceVulkanCallEnd();
 }