diff --git a/Docs/MoltenVK_Runtime_UserGuide.md b/Docs/MoltenVK_Runtime_UserGuide.md
index e4cf172..e735323 100644
--- a/Docs/MoltenVK_Runtime_UserGuide.md
+++ b/Docs/MoltenVK_Runtime_UserGuide.md
@@ -231,6 +231,8 @@
 - `VK_KHR_bind_memory2`
 - `VK_KHR_dedicated_allocation`
 - `VK_KHR_descriptor_update_template`
+- `VK_KHR_device_group`
+- `VK_KHR_device_group_creation`
 - `VK_KHR_get_memory_requirements2`
 - `VK_KHR_get_physical_device_properties2`
 - `VK_KHR_get_surface_capabilities2`
diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md
index 9a5e72f..c3fe0f6 100644
--- a/Docs/Whats_New.md
+++ b/Docs/Whats_New.md
@@ -18,12 +18,16 @@
 
 Released TBD
 
+- Add support for extensions:
+	- `VK_KHR_device_group`
 - Add support for `VkEvent`, using either native `MTLEvent` or emulation when `MTLEvent` not available.
 - Revert to supporting host-coherent memory for linear images on macOS.
 - Ensure Vulkan loader magic number is set every time before returning any dispatchable Vulkan handle.
 - Fix crash when `VkDeviceCreateInfo` specifies queue families out of numerical order.
 - Consolidate the various linkable objects into a `MVKLinkableMixin` template base class.
 - Use `MVKVector` whenever possible in MoltenVK, especially within render loop.
+- No longer prefer dedicated allocations for buffer memory, including buffer-backed images.
+- Handle the `compositeAlpha` member of `VkSwapchainCreateInfoKHR`.
 
 
 MoltenVK 1.0.36
diff --git a/ExternalRevisions/SPIRV-Cross_repo_revision b/ExternalRevisions/SPIRV-Cross_repo_revision
index 337aa79..fc4d9eb 100644
--- a/ExternalRevisions/SPIRV-Cross_repo_revision
+++ b/ExternalRevisions/SPIRV-Cross_repo_revision
@@ -1 +1 @@
-07bb1a53e0cb86cfb9b116623493df974ad9ccee
+4ce04480ec5469fe7ebbdd66c3016090a704d81b
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.h b/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.h
index 036605c..d6b45fa 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.h
@@ -31,14 +31,15 @@
 class MVKCmdDispatch : public MVKCommand {
 
 public:
-	void setContent(uint32_t x, uint32_t y, uint32_t z);
+    void setContent(uint32_t baseGroupX, uint32_t baseGroupY, uint32_t baseGroupZ,
+                    uint32_t groupCountX, uint32_t groupCountY, uint32_t groupCountZ);
 
     void encode(MVKCommandEncoder* cmdEncoder) override;
 
     MVKCmdDispatch(MVKCommandTypePool<MVKCmdDispatch>* pool);
 
 protected:
-    MTLSize  _mtlThreadgroupCount;
+    MTLRegion  _mtlThreadgroupCount;
 };
 
 
@@ -70,3 +71,6 @@
 /** Adds an indirect compute threadgroup dispatch command to the specified command buffer. */
 void mvkCmdDispatchIndirect(MVKCommandBuffer* cmdBuff, VkBuffer buffer, VkDeviceSize offset);
 
+/** Adds a compute threadgroup dispatch command to the specified command buffer, with thread IDs starting from the given base. */
+void mvkCmdDispatchBase(MVKCommandBuffer* cmdBuff, uint32_t baseGroupX, uint32_t baseGroupY, uint32_t baseGroupZ, uint32_t groupCountX, uint32_t groupCountY, uint32_t groupCountZ);
+
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.mm b/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.mm
index d52d36a..f1f4e55 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdDispatch.mm
@@ -20,6 +20,7 @@
 #include "MVKCommandBuffer.h"
 #include "MVKCommandPool.h"
 #include "MVKBuffer.h"
+#include "MVKPipeline.h"
 #include "MVKFoundation.h"
 #include "mvk_datatypes.hpp"
 
@@ -27,18 +28,30 @@
 #pragma mark -
 #pragma mark MVKCmdDispatch
 
-void MVKCmdDispatch::setContent(uint32_t x, uint32_t y, uint32_t z) {
-    _mtlThreadgroupCount.width = x;
-    _mtlThreadgroupCount.height = y;
-    _mtlThreadgroupCount.depth = z;
+void MVKCmdDispatch::setContent(uint32_t baseGroupX, uint32_t baseGroupY, uint32_t baseGroupZ,
+                                uint32_t groupCountX, uint32_t groupCountY, uint32_t groupCountZ) {
+    _mtlThreadgroupCount = MTLRegionMake3D(baseGroupX, baseGroupY, baseGroupZ, groupCountX, groupCountY, groupCountZ);
 }
 
 void MVKCmdDispatch::encode(MVKCommandEncoder* cmdEncoder) {
 //    MVKLogDebug("vkCmdDispatch() dispatching (%d, %d, %d) threadgroups.", _x, _y, _z);
 
 	cmdEncoder->finalizeDispatchState();	// Ensure all updated state has been submitted to Metal
-    [cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch) dispatchThreadgroups: _mtlThreadgroupCount
-															 threadsPerThreadgroup: cmdEncoder->_mtlThreadgroupSize];
+	id<MTLComputeCommandEncoder> mtlEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch);
+	auto* pipeline = (MVKComputePipeline*)cmdEncoder->_computePipelineState.getPipeline();
+	if (pipeline->allowsDispatchBase()) {
+		if ([mtlEncoder respondsToSelector: @selector(setStageInRegion:)]) {
+			// We'll use the stage-input region to pass the base along to the shader.
+			// Hopefully Metal won't complain that we didn't set up a stage-input descriptor.
+			[mtlEncoder setStageInRegion: _mtlThreadgroupCount];
+		} else {
+			// We have to pass the base group in a buffer.
+			unsigned int base[3] = {(uint32_t)_mtlThreadgroupCount.origin.x, (uint32_t)_mtlThreadgroupCount.origin.y, (uint32_t)_mtlThreadgroupCount.origin.z};
+			cmdEncoder->setComputeBytes(mtlEncoder, base, sizeof(base), pipeline->getIndirectParamsIndex().stages[kMVKShaderStageCompute]);
+		}
+	}
+	[mtlEncoder dispatchThreadgroups: _mtlThreadgroupCount.size
+			   threadsPerThreadgroup: cmdEncoder->_mtlThreadgroupSize];
 }
 
 MVKCmdDispatch::MVKCmdDispatch(MVKCommandTypePool<MVKCmdDispatch>* pool)
@@ -75,7 +88,7 @@
 
 void mvkCmdDispatch(MVKCommandBuffer* cmdBuff, uint32_t x, uint32_t y, uint32_t z) {
 	MVKCmdDispatch* cmd = cmdBuff->_commandPool->_cmdDispatchPool.acquireObject();
-	cmd->setContent(x, y, z);
+	cmd->setContent(0, 0, 0, x, y, z);
 	cmdBuff->addCommand(cmd);
 }
 
@@ -85,4 +98,11 @@
 	cmdBuff->addCommand(cmd);
 }
 
+void mvkCmdDispatchBase(MVKCommandBuffer* cmdBuff, uint32_t baseGroupX, uint32_t baseGroupY, uint32_t baseGroupZ,
+						uint32_t groupCountX, uint32_t groupCountY, uint32_t groupCountZ) {
+	MVKCmdDispatch* cmd = cmdBuff->_commandPool->_cmdDispatchPool.acquireObject();
+	cmd->setContent(baseGroupX, baseGroupY, baseGroupZ, groupCountX, groupCountY, groupCountZ);
+	cmdBuff->addCommand(cmd);
+}
+
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
index 1c543aa..3812739 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
@@ -313,3 +313,6 @@
 					  const VkBufferMemoryBarrier* pBufferMemoryBarriers,
 					  uint32_t imageMemoryBarrierCount,
 					  const VkImageMemoryBarrier* pImageMemoryBarriers);
+
+/** Indicates that following commands are to be recorded only for the devices in the given device mask. */
+void mvkCmdSetDeviceMask(MVKCommandBuffer* cmdBuff, uint32_t deviceMask);
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
index 7cd6e09..afab76a 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
@@ -503,5 +503,8 @@
 					bufferMemoryBarrierCount, pBufferMemoryBarriers,
 					imageMemoryBarrierCount, pImageMemoryBarriers);
 	cmdBuff->addCommand(cmd);
+}
 
+void mvkCmdSetDeviceMask(MVKCommandBuffer* cmdBuff, uint32_t deviceMask) {
+	// No-op for now...
 }
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
index 17bdb02..5f267df 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
@@ -402,6 +402,9 @@
 	/** Populates the device group surface presentation modes. */
 	VkResult getDeviceGroupSurfacePresentModes(MVKSurface* surface, VkDeviceGroupPresentModeFlagsKHR* pModes);
 
+	/** Populates the device group peer memory features. */
+	void getPeerMemoryFeatures(uint32_t heapIndex, uint32_t localDevice, uint32_t remoteDevice, VkPeerMemoryFeatureFlags* pPeerMemoryFeatures);
+
 
 #pragma mark Object lifecycle
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
index d3d66a8..2f73064 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -242,7 +242,8 @@
 	if ( !pImageFormatProperties ) { return VK_SUCCESS; }
 
 	// Metal does not support creating uncompressed views of compressed formats.
-	if (mvkIsAnyFlagEnabled(flags, VK_IMAGE_CREATE_BLOCK_TEXEL_VIEW_COMPATIBLE_BIT)) {
+	// Metal does not support split-instance images.
+	if (mvkIsAnyFlagEnabled(flags, VK_IMAGE_CREATE_BLOCK_TEXEL_VIEW_COMPATIBLE_BIT | VK_IMAGE_CREATE_SPLIT_INSTANCE_BIND_REGIONS_BIT)) {
 		return VK_ERROR_FORMAT_NOT_SUPPORTED;
 	}
 
@@ -1780,6 +1781,10 @@
 	return VK_SUCCESS;
 }
 
+void MVKDevice::getPeerMemoryFeatures(uint32_t heapIndex, uint32_t localDevice, uint32_t remoteDevice, VkPeerMemoryFeatureFlags* pPeerMemoryFeatures) {
+	*pPeerMemoryFeatures = VK_PEER_MEMORY_FEATURE_COPY_SRC_BIT | VK_PEER_MEMORY_FEATURE_COPY_DST_BIT;
+}
+
 
 #pragma mark Object lifecycle
 
@@ -1837,6 +1842,20 @@
 
 MVKImage* MVKDevice::createImage(const VkImageCreateInfo* pCreateInfo,
 								 const VkAllocationCallbacks* pAllocator) {
+	// If there's a VkImageSwapchainCreateInfoKHR, then we need to create a swapchain image.
+	const VkImageSwapchainCreateInfoKHR* swapchainInfo = nullptr;
+	for (const auto* next = (const VkBaseInStructure*)pCreateInfo->pNext; next; next = next->pNext) {
+		switch (next->sType) {
+		case VK_STRUCTURE_TYPE_IMAGE_SWAPCHAIN_CREATE_INFO_KHR:
+			swapchainInfo = (const VkImageSwapchainCreateInfoKHR*)next;
+			break;
+		default:
+			break;
+		}
+	}
+	if (swapchainInfo) {
+		return (MVKImage*)addResource(new MVKSwapchainImage(this, pCreateInfo, (MVKSwapchain*)swapchainInfo->swapchain));
+	}
 	return (MVKImage*)addResource(new MVKImage(this, pCreateInfo));
 }
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
index 68fdd23..0358457 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
@@ -393,9 +393,6 @@
 #pragma mark -
 #pragma mark MVKSwapchainImage
 
-/** Tracks a semaphore and fence for later signaling. */
-typedef std::pair<MVKSemaphore*, MVKFence*> MVKSwapchainSignaler;
-
 /** Indicates the relative availability of each image in the swapchain. */
 typedef struct MVKSwapchainImageAvailability_t {
 	uint64_t acquisitionID;			/**< When this image was last made available, relative to the other images in the swapchain. Smaller value is earlier. */
@@ -405,11 +402,18 @@
 	bool operator< (const MVKSwapchainImageAvailability_t& rhs) const;
 } MVKSwapchainImageAvailability;
 
+
 /** Represents a Vulkan image used as a rendering destination within a swapchain. */
 class MVKSwapchainImage : public MVKImage {
 
 public:
 
+	/** Binds this resource to the specified offset within the specified memory allocation. */
+	VkResult bindDeviceMemory(MVKDeviceMemory* mvkMem, VkDeviceSize memOffset) override;
+
+	/** Binds this resource according to the specified bind information. */
+	VkResult bindDeviceMemory2(const void* pBindInfo) override;
+
 	/** Returns the encompassing swapchain. */
 	inline MVKSwapchain* getSwapchain() { return _swapchain; }
 
@@ -431,7 +435,7 @@
 
 	/**
 	 * Presents the contained drawable to the OS, releases the Metal drawable and its 
-	 * texture back to the Metal layer's pool, and makes this image available for new use.
+	 * texture back to the Metal layer's pool, and makes the image memory available for new use.
 	 *
 	 * If mtlCmdBuff is not nil, the contained drawable is scheduled for presentation using
 	 * the presentDrawable: method of the command buffer. If mtlCmdBuff is nil, the contained
@@ -448,25 +452,18 @@
 					  MVKSwapchain* swapchain,
 					  uint32_t swapchainIndex);
 
-	~MVKSwapchainImage() override;
+	/** Constructs an instance for the specified device and swapchain, without binding to a particular swapchain image index. */
+	MVKSwapchainImage(MVKDevice* device,
+					  const VkImageCreateInfo* pCreateInfo,
+					  MVKSwapchain* swapchain);
 
 protected:
 	id<MTLTexture> newMTLTexture() override;
 	id<CAMetalDrawable> getCAMetalDrawable();
-	void resetCAMetalDrawable();
     void resetMetalSurface();
-	void signal(MVKSwapchainSignaler& signaler);
-	void markAsTracked(MVKSwapchainSignaler& signaler);
-	void unmarkAsTracked(MVKSwapchainSignaler& signaler);
-	void makeAvailable();
     void renderWatermark(id<MTLCommandBuffer> mtlCmdBuff);
 
 	MVKSwapchain* _swapchain;
 	uint32_t _swapchainIndex;
-	id<CAMetalDrawable> _mtlDrawable;
-	std::mutex _availabilityLock;
-	MVKVectorInline<MVKSwapchainSignaler, 4> _availabilitySignalers;
-	MVKSwapchainSignaler _preSignaled;
-	MVKSwapchainImageAvailability _availability;
 };
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
index 9cf0cce..2a20417 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
@@ -662,6 +662,9 @@
 	if (mvkIsAnyFlagEnabled(pCreateInfo->flags, VK_IMAGE_CREATE_BLOCK_TEXEL_VIEW_COMPATIBLE_BIT)) {
 		setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImage() : Metal does not allow uncompressed views of compressed images."));
 	}
+	if (mvkIsAnyFlagEnabled(pCreateInfo->flags, VK_IMAGE_CREATE_SPLIT_INSTANCE_BIND_REGIONS_BIT)) {
+		setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImage() : Metal does not support split-instance memory binding."));
+	}
 }
 
 VkSampleCountFlagBits MVKImage::validateSamples(const VkImageCreateInfo* pCreateInfo, bool isAttachment) {
@@ -1232,101 +1235,36 @@
 #pragma mark -
 #pragma mark MVKSwapchainImage
 
-bool MVKSwapchainImageAvailability_t::operator< (const MVKSwapchainImageAvailability_t& rhs) const {
-	if (  isAvailable && !rhs.isAvailable) { return true; }
-	if ( !isAvailable &&  rhs.isAvailable) { return false; }
-
-	if (waitCount < rhs.waitCount) { return true; }
-	if (waitCount > rhs.waitCount) { return false; }
-
-	return acquisitionID < rhs.acquisitionID;
+VkResult MVKSwapchainImage::bindDeviceMemory(MVKDeviceMemory*, VkDeviceSize) {
+	return VK_ERROR_OUT_OF_DEVICE_MEMORY;
 }
 
-// Makes this image available for acquisition by the app.
-// If any semaphores are waiting to be signaled when this image becomes available, the
-// earliest semaphore is signaled, and this image remains unavailable for other uses.
-void MVKSwapchainImage::makeAvailable() {
-	lock_guard<mutex> lock(_availabilityLock);
-
-	// Mark when this event happened, relative to that of other images
-	_availability.acquisitionID = _swapchain->getNextAcquisitionID();
-
-	// Mark this image as available if no semaphores or fences are waiting to be signaled.
-	_availability.isAvailable = _availabilitySignalers.empty();
-
-	MVKSwapchainSignaler signaler;
-	if (_availability.isAvailable) {
-		// If this image is now available, signal the semaphore and fence that were associated
-		// with the last time this image was acquired while available. This is a workaround for
-		// when an app uses a single semaphore or fence for more than one swapchain image.
-		// Becuase the semaphore or fence will be signaled by more than one image, it will
-		// get out of sync, and the final use of the image would not be signaled as a result.
-
-		signaler = _preSignaled;
-	} else {
-		// If this image is not yet available, extract and signal the first semaphore and fence.
-
-		signaler = _availabilitySignalers.front();
-		_availabilitySignalers.erase( _availabilitySignalers.begin() );
+VkResult MVKSwapchainImage::bindDeviceMemory2(const void* pBindInfo) {
+	const auto* imageInfo = (const VkBindImageMemoryInfo*)pBindInfo;
+	const VkBindImageMemorySwapchainInfoKHR* swapchainInfo = nullptr;
+	for (const auto* next = (const VkBaseInStructure*)imageInfo->pNext; next; next = next->pNext) {
+		switch (next->sType) {
+		case VK_STRUCTURE_TYPE_BIND_IMAGE_MEMORY_SWAPCHAIN_INFO_KHR:
+			swapchainInfo = (const VkBindImageMemorySwapchainInfoKHR*)next;
+			break;
+		default:
+			break;
+		}
+		if (swapchainInfo) { break; }
 	}
-
-	// Signal the semaphore and fence, and let them know they are no longer being tracked.
-	signal(signaler);
-	unmarkAsTracked(signaler);
-
-//	MVKLogDebug("Signaling%s swapchain image %p semaphore %p from present, with %lu remaining semaphores.", (_availability.isAvailable ? " pre-signaled" : ""), this, signaler.first, _availabilitySignalers.size());
+	if (!swapchainInfo) {
+		return VK_ERROR_OUT_OF_DEVICE_MEMORY;
+	}
+	_swapchainIndex = swapchainInfo->imageIndex;
+	return VK_SUCCESS;
 }
 
 void MVKSwapchainImage::signalWhenAvailable(MVKSemaphore* semaphore, MVKFence* fence) {
-	lock_guard<mutex> lock(_availabilityLock);
-	auto signaler = make_pair(semaphore, fence);
-	if (_availability.isAvailable) {
-		_availability.isAvailable = false;
-		signal(signaler);
-		if (_device->_useMTLEventsForSemaphores) {
-			// Unfortunately, we can't assume we have an MTLSharedEvent here.
-			// This means we need to execute a command on the device to signal
-			// the semaphore. Alternatively, we could always use an MTLSharedEvent,
-			// but that might impose unacceptable performance costs just to handle
-			// this one case.
-			MVKQueue* queue = _device->getQueue(0, 0);	
-			id<MTLCommandQueue> mtlQ = queue->getMTLCommandQueue();
-			id<MTLCommandBuffer> mtlCmdBuff = [mtlQ commandBufferWithUnretainedReferences];
-			[mtlCmdBuff enqueue];
-			signaler.first->encodeSignal(mtlCmdBuff);
-			[mtlCmdBuff commit];
-		}
-		_preSignaled = signaler;
-	} else {
-		_availabilitySignalers.push_back(signaler);
-	}
-	markAsTracked(signaler);
-
-//	MVKLogDebug("%s swapchain image %p semaphore %p in acquire with %lu other semaphores.", (_availability.isAvailable ? "Signaling" : "Tracking"), this, semaphore, _availabilitySignalers.size());
-}
-
-// Signal either or both of the semaphore and fence in the specified tracker pair.
-void MVKSwapchainImage::signal(MVKSwapchainSignaler& signaler) {
-	if (signaler.first && !_device->_useMTLEventsForSemaphores) { signaler.first->signal(); }
-	if (signaler.second) { signaler.second->signal(); }
-}
-
-// Tell the semaphore and fence that they are being tracked for future signaling.
-void MVKSwapchainImage::markAsTracked(MVKSwapchainSignaler& signaler) {
-	if (signaler.first) { signaler.first->retain(); }
-	if (signaler.second) { signaler.second->retain(); }
-}
-
-// Tell the semaphore and fence that they are no longer being tracked for future signaling.
-void MVKSwapchainImage::unmarkAsTracked(MVKSwapchainSignaler& signaler) {
-	if (signaler.first) { signaler.first->release(); }
-	if (signaler.second) { signaler.second->release(); }
+	_swapchain->signalWhenAvailable( _swapchainIndex, semaphore, fence );
 }
 
 const MVKSwapchainImageAvailability* MVKSwapchainImage::getAvailability() {
-	lock_guard<mutex> lock(_availabilityLock);
-	_availability.waitCount = (uint32_t)_availabilitySignalers.size();
-	return &_availability;
+	return _swapchain->getAvailability( _swapchainIndex );
 }
 
 
@@ -1339,13 +1277,9 @@
 }
 
 id<CAMetalDrawable> MVKSwapchainImage::getCAMetalDrawable() {
-	if ( !_mtlDrawable ) {
-		@autoreleasepool {		// Allow auto-released drawable object to be reclaimed before end of loop
-			_mtlDrawable = [_swapchain->getNextCAMetalDrawable() retain];	// retained
-		}
-		MVKAssert(_mtlDrawable, "Could not aquire an available CAMetalDrawable from the CAMetalLayer in MVKSwapchain image: %p.", this);
-	}
-	return _mtlDrawable;
+	id<CAMetalDrawable> mtlDrawable = _swapchain->getCAMetalDrawable(_swapchainIndex);
+	MVKAssert(mtlDrawable, "Could not acquire an available CAMetalDrawable from the CAMetalLayer in MVKSwapchain image: %p.", this);
+	return mtlDrawable;
 }
 
 void MVKSwapchainImage::presentCAMetalDrawable(id<MTLCommandBuffer> mtlCmdBuff) {
@@ -1364,33 +1298,26 @@
 		if (scName) { [mtlCmdBuff popDebugGroup]; }
 
 		resetMetalSurface();
-        if (_device->_useMTLEventsForSemaphores && !_availabilitySignalers.empty()) {
-            // Signal the semaphore device-side.
-            _availabilitySignalers.front().first->encodeSignal(mtlCmdBuff);
+        if (_device->_useMTLEventsForSemaphores) {
+            _swapchain->signalOnDevice(_swapchainIndex, mtlCmdBuff);
         }
 
 		retain();	// Ensure this image is not destroyed while awaiting MTLCommandBuffer completion
         [mtlCmdBuff addCompletedHandler: ^(id<MTLCommandBuffer> mcb) {
-			makeAvailable();
+			_swapchain->makeAvailable(_swapchainIndex);
 			release();
 		}];
     } else {
         [mtlDrawable present];
         resetMetalSurface();
-        makeAvailable();
+        _swapchain->makeAvailable(_swapchainIndex);
     }
 }
 
-// Removes and releases the Metal drawable object, so that it can be lazily created by getCAMetalDrawable().
-void MVKSwapchainImage::resetCAMetalDrawable() {
-	[_mtlDrawable release];
-	_mtlDrawable = nil;
-}
-
 // Resets the MTLTexture and CAMetalDrawable underlying this image.
 void MVKSwapchainImage::resetMetalSurface() {
     resetMTLTexture();			// Release texture first so drawable will be last to release it
-    resetCAMetalDrawable();
+    _swapchain->resetCAMetalDrawable(_swapchainIndex);
 }
 
 
@@ -1402,14 +1329,13 @@
 									 uint32_t swapchainIndex) : MVKImage(device, pCreateInfo) {
 	_swapchain = swapchain;
 	_swapchainIndex = swapchainIndex;
-	_availability.acquisitionID = _swapchain->getNextAcquisitionID();
-	_availability.isAvailable = true;
-	_preSignaled = make_pair(nullptr, nullptr);
-	_mtlDrawable = nil;
 }
 
-MVKSwapchainImage::~MVKSwapchainImage() {
-	resetCAMetalDrawable();
+MVKSwapchainImage::MVKSwapchainImage(MVKDevice* device,
+									 const VkImageCreateInfo* pCreateInfo,
+									 MVKSwapchain* swapchain) : MVKImage(device, pCreateInfo) {
+	_swapchain = swapchain;
+	_swapchainIndex = uint32_t(-1);
 }
 
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm
index 31558bb..7a32e7a 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm
@@ -609,6 +609,9 @@
 	ADD_DVC_EXT_ENTRY_POINT(vkCreateDescriptorUpdateTemplateKHR, KHR_DESCRIPTOR_UPDATE_TEMPLATE);
 	ADD_DVC_EXT_ENTRY_POINT(vkDestroyDescriptorUpdateTemplateKHR, KHR_DESCRIPTOR_UPDATE_TEMPLATE);
 	ADD_DVC_EXT_ENTRY_POINT(vkUpdateDescriptorSetWithTemplateKHR, KHR_DESCRIPTOR_UPDATE_TEMPLATE);
+	ADD_DVC_EXT_ENTRY_POINT(vkGetDeviceGroupPeerMemoryFeaturesKHR, KHR_DEVICE_GROUP);
+	ADD_DVC_EXT_ENTRY_POINT(vkCmdSetDeviceMaskKHR, KHR_DEVICE_GROUP);
+	ADD_DVC_EXT_ENTRY_POINT(vkCmdDispatchBaseKHR, KHR_DEVICE_GROUP);
 	ADD_DVC_EXT_ENTRY_POINT(vkGetBufferMemoryRequirements2KHR, KHR_GET_MEMORY_REQUIREMENTS_2);
 	ADD_DVC_EXT_ENTRY_POINT(vkGetImageMemoryRequirements2KHR, KHR_GET_MEMORY_REQUIREMENTS_2);
 	ADD_DVC_EXT_ENTRY_POINT(vkGetImageSparseMemoryRequirements2KHR, KHR_GET_MEMORY_REQUIREMENTS_2);
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
index 99bb888..30d48c2 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
@@ -156,6 +156,9 @@
 	/** Returns the current buffer size buffer bindings. */
 	const MVKShaderImplicitRezBinding& getBufferSizeBufferIndex() { return _bufferSizeBufferIndex; }
 
+	/** Returns the current indirect parameter buffer bindings. */
+	const MVKShaderImplicitRezBinding& getIndirectParamsIndex() { return _indirectParamsIndex; }
+
 	/** Returns whether or not full image view swizzling is enabled for this pipeline. */
 	bool fullImageViewSwizzle() const { return _fullImageViewSwizzle; }
 
@@ -171,6 +174,7 @@
 	MVKPipelineCache* _pipelineCache;
 	MVKShaderImplicitRezBinding _swizzleBufferIndex;
 	MVKShaderImplicitRezBinding _bufferSizeBufferIndex;
+	MVKShaderImplicitRezBinding _indirectParamsIndex;
 	MVKShaderResourceBinding _pushConstantsMTLResourceIndexes;
 	bool _fullImageViewSwizzle;
 	bool _hasValidMTLPipelineStates = true;
@@ -204,9 +208,6 @@
     /** Returns the number of output tessellation patch control points. */
     uint32_t getOutputControlPointCount() { return _outputControlPointCount; }
 
-	/** Returns the current indirect parameter buffer bindings. */
-	const MVKShaderImplicitRezBinding& getIndirectParamsIndex() { return _indirectParamsIndex; }
-
 	/** Returns the current captured output buffer bindings. */
 	const MVKShaderImplicitRezBinding& getOutputBufferIndex() { return _outputBufferIndex; }
 
@@ -281,7 +282,6 @@
 
     float _blendConstants[4] = { 0.0, 0.0, 0.0, 1.0 };
     uint32_t _outputControlPointCount;
-	MVKShaderImplicitRezBinding _indirectParamsIndex;
 	MVKShaderImplicitRezBinding _outputBufferIndex;
 	uint32_t _tessCtlPatchOutputBufferIndex = 0;
 	uint32_t _tessCtlLevelBufferIndex = 0;
@@ -317,6 +317,9 @@
 	/** Binds this pipeline to the specified command encoder. */
 	void encode(MVKCommandEncoder* cmdEncoder, uint32_t = 0) override;
 
+	/** Returns if this pipeline allows non-zero dispatch bases in vkCmdDispatchBase(). */
+	bool allowsDispatchBase() { return _allowsDispatchBase; }
+
 	/** Constructs an instance for the device and parent (which may be NULL). */
 	MVKComputePipeline(MVKDevice* device,
 					   MVKPipelineCache* pipelineCache,
@@ -332,6 +335,8 @@
     MTLSize _mtlThreadgroupSize;
     bool _needsSwizzleBuffer = false;
     bool _needsBufferSizeBuffer = false;
+    bool _needsDispatchBaseBuffer = false;
+    bool _allowsDispatchBase = false;
 };
 
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
index c21340d..265d922 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
@@ -1294,6 +1294,8 @@
 									   const VkComputePipelineCreateInfo* pCreateInfo) :
 	MVKPipeline(device, pipelineCache, (MVKPipelineLayout*)pCreateInfo->layout, parent) {
 
+	_allowsDispatchBase = mvkAreAllFlagsEnabled(pCreateInfo->flags, VK_PIPELINE_CREATE_DISPATCH_BASE);	// sic; drafters forgot the 'BIT' suffix
+
 	MVKMTLFunction func = getMTLFunction(pCreateInfo);
 	_mtlThreadgroupSize = func.threadGroupSize;
 	_mtlPipelineState = nil;
@@ -1324,6 +1326,9 @@
 	if (_needsBufferSizeBuffer && _bufferSizeBufferIndex.stages[kMVKShaderStageCompute] > _device->_pMetalFeatures->maxPerStageBufferCount) {
 		setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Compute shader requires buffer size buffer, but there is no free slot to pass it."));
 	}
+	if (_needsDispatchBaseBuffer && _indirectParamsIndex.stages[kMVKShaderStageCompute] > _device->_pMetalFeatures->maxPerStageBufferCount) {
+		setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Compute shader requires dispatch base buffer, but there is no free slot to pass it."));
+	}
 }
 
 // Returns a MTLFunction to use when creating the MTLComputePipelineState.
@@ -1339,6 +1344,7 @@
     shaderContext.options.mslOptions.texel_buffer_texture_width = _device->_pMetalFeatures->maxTextureDimension;
 	shaderContext.options.mslOptions.swizzle_texture_samples = _fullImageViewSwizzle;
 	shaderContext.options.mslOptions.texture_buffer_native = _device->_pMetalFeatures->textureBuffers;
+	shaderContext.options.mslOptions.dispatch_base = _allowsDispatchBase;
 
     MVKPipelineLayout* layout = (MVKPipelineLayout*)pCreateInfo->layout;
     layout->populateShaderConverterContext(shaderContext);
@@ -1346,12 +1352,14 @@
     _bufferSizeBufferIndex = layout->getBufferSizeBufferIndex();
     shaderContext.options.mslOptions.swizzle_buffer_index = _swizzleBufferIndex.stages[kMVKShaderStageCompute];
     shaderContext.options.mslOptions.buffer_size_buffer_index = _bufferSizeBufferIndex.stages[kMVKShaderStageCompute];
+    shaderContext.options.mslOptions.indirect_params_buffer_index = _indirectParamsIndex.stages[kMVKShaderStageCompute];
 
     MVKMTLFunction func = ((MVKShaderModule*)pSS->module)->getMTLFunction(&shaderContext, pSS->pSpecializationInfo, _pipelineCache);
 
 	auto& funcRslts = func.shaderConversionResults;
 	_needsSwizzleBuffer = funcRslts.needsSwizzleBuffer;
     _needsBufferSizeBuffer = funcRslts.needsBufferSizeBuffer;
+    _needsDispatchBaseBuffer = funcRslts.needsDispatchBaseBuffer;
 
 	return func;
 }
@@ -1710,7 +1718,8 @@
 				scr.needsOutputBuffer,
 				scr.needsPatchOutputBuffer,
 				scr.needsBufferSizeBuffer,
-				scr.needsInputThreadgroupMem);
+				scr.needsInputThreadgroupMem,
+				scr.needsDispatchBaseBuffer);
 	}
 
 }
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
index 802f2ad..24421f4 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
@@ -19,9 +19,9 @@
 #pragma once
 
 #include "MVKDevice.h"
+#include "MVKImage.h"
 #include "MVKVector.h"
 
-class MVKSwapchainImage;
 class MVKWatermark;
 
 @class MVKBlockObserver;
@@ -29,6 +29,9 @@
 
 #pragma mark MVKSwapchain
 
+/** Tracks a semaphore and fence for later signaling. */
+typedef std::pair<MVKSemaphore*, MVKFence*> MVKSwapchainSignaler;
+
 /** Represents a Vulkan swapchain. */
 class MVKSwapchain : public MVKVulkanAPIDeviceObject {
 
@@ -77,17 +80,27 @@
 	/** Returns the specified performance stats structure. */
 	const MVKSwapchainPerformance* getPerformanceStatistics() { return &_performanceStatistics; }
 
+	/**
+	 * Registers a semaphore and/or fence that will be signaled when the image at the given index becomes available.
+	 * This function accepts both a semaphore and a fence, and either none, one, or both may be provided.
+	 * If this image is available already, the semaphore and fence are immediately signaled.
+	 */
+	void signalWhenAvailable(uint32_t imageIndex, MVKSemaphore* semaphore, MVKFence* fence);
+
+	/** Returns the availability status of the image at the given index, relative to other images in the swapchain. */
+	const MVKSwapchainImageAvailability* getAvailability(uint32_t imageIndex);
 
 #pragma mark Metal
 
 	/** 
-	 * Returns the next Metal drawable available to provide backing for 
-	 * an image in this swapchain. The returned object is autoreleased.
+	 * Returns the Metal drawable providing backing for the image at the given
+	 * index in this swapchain. If none is established, the next available
+	 * drawable is acquired and returned.
 	 *
 	 * This function may block until the next drawable is available, 
 	 * and may return nil if no drawable is available at all.
 	 */
-	id<CAMetalDrawable> getNextCAMetalDrawable();
+	id<CAMetalDrawable> getCAMetalDrawable(uint32_t imgIdx);
 
 
 #pragma mark Construction
@@ -99,6 +112,12 @@
 protected:
 	friend class MVKSwapchainImage;
 
+	struct Availability {
+		MVKSwapchainImageAvailability status;
+		MVKVectorInline<MVKSwapchainSignaler, 4> signalers;
+		MVKSwapchainSignaler preSignaled;
+	};
+
 	void propogateDebugName() override;
 	void initCAMetalLayer(const VkSwapchainCreateInfoKHR* pCreateInfo, uint32_t imgCnt);
 	void initSurfaceImages(const VkSwapchainCreateInfoKHR* pCreateInfo, uint32_t imgCnt);
@@ -108,10 +127,19 @@
     void willPresentSurface(id<MTLTexture> mtlTexture, id<MTLCommandBuffer> mtlCmdBuff);
     void renderWatermark(id<MTLTexture> mtlTexture, id<MTLCommandBuffer> mtlCmdBuff);
     void markFrameInterval();
+	void resetCAMetalDrawable(uint32_t imgIdx);
+	void signal(MVKSwapchainSignaler& signaler);
+	void signalOnDevice(uint32_t imgIdx, id<MTLCommandBuffer> mtlCmdBuff);
+	static void markAsTracked(MVKSwapchainSignaler& signaler);
+	static void unmarkAsTracked(MVKSwapchainSignaler& signaler);
+	void makeAvailable(uint32_t imgIdx);
 
 	CAMetalLayer* _mtlLayer;
     MVKWatermark* _licenseWatermark;
 	MVKVectorInline<MVKSwapchainImage*, kMVKMaxSwapchainImageCount> _surfaceImages;
+	MVKVectorInline<id<CAMetalDrawable>, kMVKMaxSwapchainImageCount> _mtlDrawables;
+	MVKVectorInline<Availability, kMVKMaxSwapchainImageCount> _imageAvailability;
+	std::mutex _availabilityLock;
 	std::atomic<uint64_t> _currentAcquisitionID;
     CGSize _mtlLayerOrigDrawSize;
     MVKSwapchainPerformance _performanceStatistics;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.mm b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.mm
index cae47c8..e1e6467 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.mm
@@ -19,6 +19,7 @@
 #include "MVKSurface.h"
 #include "MVKSwapchain.h"
 #include "MVKImage.h"
+#include "MVKQueue.h"
 #include "MVKFoundation.h"
 #include "MVKOSExtensions.h"
 #include "MVKWatermark.h"
@@ -34,6 +35,16 @@
 
 #pragma mark MVKSwapchain
 
+bool MVKSwapchainImageAvailability_t::operator< (const MVKSwapchainImageAvailability_t& rhs) const {
+	if (  isAvailable && !rhs.isAvailable) { return true; }
+	if ( !isAvailable &&  rhs.isAvailable) { return false; }
+
+	if (waitCount < rhs.waitCount) { return true; }
+	if (waitCount > rhs.waitCount) { return false; }
+
+	return acquisitionID < rhs.acquisitionID;
+}
+
 void MVKSwapchain::propogateDebugName() {
 	if (_debugName) {
 		size_t imgCnt = _surfaceImages.size();
@@ -45,7 +56,7 @@
 	}
 }
 
-uint32_t MVKSwapchain::getImageCount() { return (uint32_t)_surfaceImages.size(); }
+uint32_t MVKSwapchain::getImageCount() { return (uint32_t)_imageAvailability.size(); }
 
 MVKSwapchainImage* MVKSwapchain::getImage(uint32_t index) { return _surfaceImages[index]; }
 
@@ -85,16 +96,16 @@
     MVKSwapchainImageAvailability minAvailability = { .acquisitionID = kMVKUndefinedLargeUInt64,
 													  .waitCount = kMVKUndefinedLargeUInt32,
 													  .isAvailable = false };
-    for (MVKSwapchainImage* mvkSCImg : _surfaceImages) {
-        const MVKSwapchainImageAvailability* currAvailability = mvkSCImg->getAvailability();
-        if (*currAvailability < minAvailability) {
-            minAvailability = *currAvailability;
-            minWaitIndex = mvkSCImg->getSwapchainIndex();
+    for (uint32_t imgIdx = 0; imgIdx < _imageAvailability.size(); imgIdx++) {
+        const Availability& avail = _imageAvailability[imgIdx];
+        if (avail.status < minAvailability) {
+            minAvailability = avail.status;
+            minWaitIndex = imgIdx;
         }
     }
 
     *pImageIndex = minWaitIndex;	// Return the index of the image with the shortest wait
-    _surfaceImages[minWaitIndex]->signalWhenAvailable((MVKSemaphore*)semaphore, (MVKFence*)fence);
+    signalWhenAvailable(minWaitIndex, (MVKSemaphore*)semaphore, (MVKFence*)fence);
     return getHasSurfaceSizeChanged() ? VK_ERROR_OUT_OF_DATE_KHR : VK_SUCCESS;
 }
 
@@ -110,6 +121,103 @@
  */
 void MVKSwapchain::releaseUndisplayedSurfaces() {}
 
+// Makes an image available for acquisition by the app.
+// If any semaphores are waiting to be signaled when this image becomes available, the
+// earliest semaphore is signaled, and this image remains unavailable for other uses.
+void MVKSwapchain::makeAvailable(uint32_t imgIdx) {
+	lock_guard<mutex> lock(_availabilityLock);
+	auto& availability = _imageAvailability[imgIdx].status;
+
+	// Mark when this event happened, relative to that of other images
+	availability.acquisitionID = getNextAcquisitionID();
+
+	// Mark this image as available if no semaphores or fences are waiting to be signaled.
+	availability.isAvailable = _imageAvailability[imgIdx].signalers.empty();
+
+	MVKSwapchainSignaler signaler;
+	if (availability.isAvailable) {
+		// If this image is now available, signal the semaphore and fence that were associated
+		// with the last time this image was acquired while available. This is a workaround for
+		// when an app uses a single semaphore or fence for more than one swapchain image.
+		// Becuase the semaphore or fence will be signaled by more than one image, it will
+		// get out of sync, and the final use of the image would not be signaled as a result.
+
+		signaler = _imageAvailability[imgIdx].preSignaled;
+	} else {
+		// If this image is not yet available, extract and signal the first semaphore and fence.
+
+		signaler = _imageAvailability[imgIdx].signalers.front();
+		_imageAvailability[imgIdx].signalers.erase( _imageAvailability[imgIdx].signalers.begin() );
+	}
+
+	// Signal the semaphore and fence, and let them know they are no longer being tracked.
+	signal(signaler);
+	unmarkAsTracked(signaler);
+
+//	MVKLogDebug("Signaling%s swapchain image %p semaphore %p from present, with %lu remaining semaphores.", (_availability.isAvailable ? " pre-signaled" : ""), this, signaler.first, _availabilitySignalers.size());
+}
+
+void MVKSwapchain::signalWhenAvailable(uint32_t imageIndex, MVKSemaphore* semaphore, MVKFence* fence) {
+	lock_guard<mutex> lock(_availabilityLock);
+	auto signaler = make_pair(semaphore, fence);
+	auto& availability = _imageAvailability[imageIndex].status; 
+	if (availability.isAvailable) {
+		availability.isAvailable = false;
+		signal(signaler);
+		if (_device->_useMTLEventsForSemaphores) {
+			// Unfortunately, we can't assume we have an MTLSharedEvent here.
+			// This means we need to execute a command on the device to signal
+			// the semaphore. Alternatively, we could always use an MTLSharedEvent,
+			// but that might impose unacceptable performance costs just to handle
+			// this one case.
+			MVKQueue* queue = _device->getQueue(0, 0);	
+			id<MTLCommandQueue> mtlQ = queue->getMTLCommandQueue();
+			id<MTLCommandBuffer> mtlCmdBuff = [mtlQ commandBufferWithUnretainedReferences];
+			[mtlCmdBuff enqueue];
+			signaler.first->encodeSignal(mtlCmdBuff);
+			[mtlCmdBuff commit];
+		}
+		_imageAvailability[imageIndex].preSignaled = signaler;
+	} else {
+		_imageAvailability[imageIndex].signalers.push_back(signaler);
+	}
+	markAsTracked(signaler);
+
+//	MVKLogDebug("%s swapchain image %p semaphore %p in acquire with %lu other semaphores.", (_availability.isAvailable ? "Signaling" : "Tracking"), this, semaphore, _availabilitySignalers.size());
+}
+
+// Signal either or both of the semaphore and fence in the specified tracker pair.
+void MVKSwapchain::signal(MVKSwapchainSignaler& signaler) {
+	if (signaler.first && !_device->_useMTLEventsForSemaphores) { signaler.first->signal(); }
+	if (signaler.second) { signaler.second->signal(); }
+}
+
+// If present, signal the semaphore for the first waiter for the given image.
+void MVKSwapchain::signalOnDevice(uint32_t imgIdx, id<MTLCommandBuffer> mtlCmdBuff) {
+	lock_guard<mutex> lock(_availabilityLock);
+	MVKSemaphore* mvkSem = _imageAvailability[imgIdx].signalers.front().first;
+	if (mvkSem) { mvkSem->encodeSignal(mtlCmdBuff); }
+}
+
+// Tell the semaphore and fence that they are being tracked for future signaling.
+void MVKSwapchain::markAsTracked(MVKSwapchainSignaler& signaler) {
+	if (signaler.first) { signaler.first->retain(); }
+	if (signaler.second) { signaler.second->retain(); }
+}
+
+// Tell the semaphore and fence that they are no longer being tracked for future signaling.
+void MVKSwapchain::unmarkAsTracked(MVKSwapchainSignaler& signaler) {
+	if (signaler.first) { signaler.first->release(); }
+	if (signaler.second) { signaler.second->release(); }
+}
+
+const MVKSwapchainImageAvailability* MVKSwapchain::getAvailability(uint32_t imageIndex) {
+	lock_guard<mutex> lock(_availabilityLock);
+	auto& availability = _imageAvailability[imageIndex].status;
+	availability.waitCount = (uint32_t)_imageAvailability[imageIndex].signalers.size();
+	return &availability;
+}
+
 
 #pragma mark Rendering
 
@@ -176,12 +284,22 @@
 
 #pragma mark Metal
 
-id<CAMetalDrawable> MVKSwapchain::getNextCAMetalDrawable() {
-    id<CAMetalDrawable> nextDrwbl = nil;
-    while ( !(nextDrwbl = [_mtlLayer nextDrawable]) ) {
-        MVKLogError("Drawable could not be retrieved! Elapsed time: %.6f ms.", mvkGetElapsedMilliseconds());
+id<CAMetalDrawable> MVKSwapchain::getCAMetalDrawable(uint32_t imageIndex) {
+    if ( _mtlDrawables[imageIndex] ) { return _mtlDrawables[imageIndex]; }
+    @autoreleasepool {      // Allow auto-released drawable object to be reclaimed before end of loop
+        id<CAMetalDrawable> nextDrwbl = nil;
+        while ( !(nextDrwbl = [_mtlLayer nextDrawable]) ) {
+            MVKLogError("Drawable could not be retrieved! Elapsed time: %.6f ms.", mvkGetElapsedMilliseconds());
+        }
+        _mtlDrawables[imageIndex] = [nextDrwbl retain];
     }
-    return nextDrwbl;
+    return _mtlDrawables[imageIndex];
+}
+
+// Removes and releases a Metal drawable object, so that it can be lazily created by getCAMetalDrawable().
+void MVKSwapchain::resetCAMetalDrawable(uint32_t imgIdx) {
+	[_mtlDrawables[imgIdx] release];
+	_mtlDrawables[imgIdx] = nil;
 }
 
 
@@ -304,10 +422,20 @@
 	if (mvkAreAllFlagsEnabled(pCreateInfo->flags, VK_SWAPCHAIN_CREATE_MUTABLE_FORMAT_BIT_KHR)) {
 		mvkEnableFlag(imgInfo.flags, VK_IMAGE_CREATE_MUTABLE_FORMAT_BIT | VK_IMAGE_CREATE_EXTENDED_USAGE_BIT);
 	}
+	if (mvkAreAllFlagsEnabled(pCreateInfo->flags, VK_SWAPCHAIN_CREATE_SPLIT_INSTANCE_BIND_REGIONS_BIT_KHR)) {
+		// We don't really support this, but set the flag anyway.
+		mvkEnableFlag(imgInfo.flags, VK_IMAGE_CREATE_SPLIT_INSTANCE_BIND_REGIONS_BIT);
+	}
 
 	_surfaceImages.reserve(imgCnt);
+	_mtlDrawables.resize(imgCnt);
+	_imageAvailability.resize(imgCnt);
     for (uint32_t imgIdx = 0; imgIdx < imgCnt; imgIdx++) {
         _surfaceImages.push_back(_device->createSwapchainImage(&imgInfo, this, imgIdx, NULL));
+        _imageAvailability[imgIdx].status.acquisitionID = getNextAcquisitionID();
+        _imageAvailability[imgIdx].status.isAvailable = true;
+        _imageAvailability[imgIdx].preSignaled = make_pair(nullptr, nullptr);
+        _mtlDrawables[imgIdx] = nil;
     }
 
     MVKLogInfo("Created %d swapchain images with initial size (%d, %d).", imgCnt, imgExtent.width, imgExtent.height);
diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
index 44589cb..a7edcb7 100644
--- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def
+++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
@@ -35,6 +35,7 @@
 MVK_EXTENSION(KHR_bind_memory2, KHR_BIND_MEMORY_2)
 MVK_EXTENSION(KHR_dedicated_allocation, KHR_DEDICATED_ALLOCATION)
 MVK_EXTENSION(KHR_descriptor_update_template, KHR_DESCRIPTOR_UPDATE_TEMPLATE)
+MVK_EXTENSION(KHR_device_group, KHR_DEVICE_GROUP)
 MVK_EXTENSION(KHR_device_group_creation, KHR_DEVICE_GROUP_CREATION)
 MVK_EXTENSION(KHR_get_memory_requirements2, KHR_GET_MEMORY_REQUIREMENTS_2)
 MVK_EXTENSION(KHR_get_physical_device_properties2, KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2)
diff --git a/MoltenVK/MoltenVK/Utility/MVKVector.h b/MoltenVK/MoltenVK/Utility/MVKVector.h
index e7843a8..4d4adee 100755
--- a/MoltenVK/MoltenVK/Utility/MVKVector.h
+++ b/MoltenVK/MoltenVK/Utility/MVKVector.h
@@ -191,7 +191,7 @@
   virtual void                reset()                                        = 0;
   virtual void                reserve( const size_t new_size )               = 0;
   virtual void                assign( const size_t new_size, const Type *t ) = 0;
-  virtual void                resize( const size_t new_size, const Type *t ) = 0;
+  virtual void                resize( const size_t new_size, const Type *t = nullptr ) = 0;
   virtual void                shrink_to_fit()                                = 0;
   virtual void                push_back( const Type *t )                     = 0;
 };
@@ -880,7 +880,7 @@
     alc.num_elements_used = new_size;
   }
 
-  void resize( const size_t new_size, const Type *t ) override
+  void resize( const size_t new_size, const Type *t = nullptr ) override
   {
     if ( new_size == alc.num_elements_used )
     {
diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
index 1604d24..e44f368 100644
--- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm
+++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
@@ -1969,6 +1969,48 @@
 
 
 #pragma mark -
+#pragma mark VK_KHR_device_group extension
+
+MVK_PUBLIC_SYMBOL void vkGetDeviceGroupPeerMemoryFeaturesKHR(
+    VkDevice                                    device,
+    uint32_t                                    heapIndex,
+    uint32_t                                    localDeviceIndex,
+    uint32_t                                    remoteDeviceIndex,
+    VkPeerMemoryFeatureFlagsKHR*                pPeerMemoryFeatures) {
+
+    MVKTraceVulkanCallStart();
+    MVKDevice* mvkDev = MVKDevice::getMVKDevice(device);
+    mvkDev->getPeerMemoryFeatures(heapIndex, localDeviceIndex, remoteDeviceIndex, pPeerMemoryFeatures);
+    MVKTraceVulkanCallEnd();
+}
+
+MVK_PUBLIC_SYMBOL void vkCmdSetDeviceMaskKHR(
+    VkCommandBuffer                             commandBuffer,
+    uint32_t                                    deviceMask) {
+
+    MVKTraceVulkanCallStart();
+    MVKCommandBuffer* cmdBuff = MVKCommandBuffer::getMVKCommandBuffer(commandBuffer);
+    mvkCmdSetDeviceMask(cmdBuff, deviceMask);
+    MVKTraceVulkanCallEnd();
+}
+
+MVK_PUBLIC_SYMBOL void vkCmdDispatchBaseKHR(
+    VkCommandBuffer                             commandBuffer,
+    uint32_t                                    baseGroupX,
+    uint32_t                                    baseGroupY,
+    uint32_t                                    baseGroupZ,
+    uint32_t                                    groupCountX,
+    uint32_t                                    groupCountY,
+    uint32_t                                    groupCountZ) {
+	
+	MVKTraceVulkanCallStart();
+    MVKCommandBuffer* cmdBuff = MVKCommandBuffer::getMVKCommandBuffer(commandBuffer);
+	mvkCmdDispatchBase(cmdBuff, baseGroupX, baseGroupY, baseGroupZ, groupCountX, groupCountY, groupCountZ);
+	MVKTraceVulkanCallEnd();
+}
+
+
+#pragma mark -
 #pragma mark VK_KHR_device_group_creation extension
 
 MVK_PUBLIC_SYMBOL VkResult vkEnumeratePhysicalDeviceGroupsKHR(
diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp
index 885bb79..e331486 100644
--- a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp
+++ b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp
@@ -308,6 +308,7 @@
 	_shaderConversionResults.needsPatchOutputBuffer = pMSLCompiler && pMSLCompiler->needs_patch_output_buffer();
 	_shaderConversionResults.needsBufferSizeBuffer = pMSLCompiler && pMSLCompiler->needs_buffer_size_buffer();
 	_shaderConversionResults.needsInputThreadgroupMem = pMSLCompiler && pMSLCompiler->needs_input_threadgroup_mem();
+	_shaderConversionResults.needsDispatchBaseBuffer = pMSLCompiler && pMSLCompiler->needs_dispatch_base_buffer();
 
 	if (context.stageSupportsVertexAttributes()) {
 		for (auto& ctxVA : context.vertexAttributes) {
diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h
index 82024cd..0b7f463 100644
--- a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h
+++ b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h
@@ -204,6 +204,7 @@
 		bool needsPatchOutputBuffer = false;
 		bool needsBufferSizeBuffer = false;
 		bool needsInputThreadgroupMem = false;
+		bool needsDispatchBaseBuffer = false;
 
 		void reset() { *this = SPIRVToMSLConversionResults(); }