diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
index 24e840b..5714c1b 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
@@ -1166,8 +1166,8 @@
 
         // Validate
         MVKMTLFmtCaps mtlFmtCaps = cmdBuff->getPixelFormats()->getCapabilities(_image->getMTLPixelFormat(planeIndex));
-        if ((isDS && !mvkAreAllFlagsEnabled(mtlFmtCaps, kMVKMTLFmtCapsDSAtt)) ||
-            ( !isDS && !mvkAreAllFlagsEnabled(mtlFmtCaps, kMVKMTLFmtCapsColorAtt))) {
+		uint32_t reqCap = isDS ? kMVKMTLFmtCapsDSAtt : (_image->getIsLinear() ? kMVKMTLFmtCapsWrite : kMVKMTLFmtCapsColorAtt);
+        if (!mvkAreAllFlagsEnabled(mtlFmtCaps, reqCap)) {
             return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdClear%sImage(): Format %s cannot be cleared on this device.", (isDS ? "DepthStencil" : "Color"), cmdBuff->getPixelFormats()->getName(_image->getVkFormat()));
         }
         
@@ -1196,6 +1196,31 @@
         id<MTLTexture> imgMTLTex = _image->getMTLTexture(MVKImage::getPlaneFromVkImageAspectFlags(srRange.aspectMask));
         if ( !imgMTLTex ) { continue; }
 
+#if MVK_MACOS
+        if ( _image->getIsLinear() ) {
+            // These images cannot be rendered. Instead, use a compute shader.
+            // Luckily for us, linear images only have one mip and one array layer under Metal.
+            assert( !isDS );
+            id<MTLComputePipelineState> mtlClearState = cmdEncoder->getCommandEncodingPool()->getCmdClearColorImageMTLComputePipelineState(pixFmts->getFormatType(_image->getVkFormat()));
+            id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseClearColorImage);
+            [mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"];
+            [mtlComputeEnc setComputePipelineState: mtlClearState];
+            [mtlComputeEnc setTexture: imgMTLTex atIndex: 0];
+            cmdEncoder->setComputeBytes(mtlComputeEnc, &_clearValue, sizeof(_clearValue), 0);
+            MTLSize gridSize = mvkMTLSizeFromVkExtent3D(_image->getExtent3D());
+            MTLSize tgSize = MTLSizeMake(mtlClearState.threadExecutionWidth, 1, 1);
+            if (cmdEncoder->getDevice()->_pMetalFeatures->nonUniformThreadgroups) {
+                [mtlComputeEnc dispatchThreads: gridSize threadsPerThreadgroup: tgSize];
+            } else {
+                MTLSize tgCount = MTLSizeMake(gridSize.width / tgSize.width, gridSize.height, gridSize.depth);
+                if (gridSize.width % tgSize.width) { tgCount.width += 1; }
+                [mtlComputeEnc dispatchThreadgroups: tgCount threadsPerThreadgroup: tgSize];
+            }
+            [mtlComputeEnc popDebugGroup];
+            continue;
+        }
+#endif
+
 		MTLRenderPassDescriptor* mtlRPDesc = [MTLRenderPassDescriptor renderPassDescriptor];
 		MTLRenderPassColorAttachmentDescriptor* mtlRPCADesc = nil;
 		MTLRenderPassDepthAttachmentDescriptor* mtlRPDADesc = nil;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
index 24b65a4..0889577 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -752,6 +752,7 @@
         case kMVKCommandUseCopyBufferToImage:   return @"vkCmdCopyBufferToImage ComputeEncoder";
         case kMVKCommandUseCopyImageToBuffer:   return @"vkCmdCopyImageToBuffer ComputeEncoder";
         case kMVKCommandUseFillBuffer:          return @"vkCmdFillBuffer ComputeEncoder";
+        case kMVKCommandUseClearColorImage:     return @"vkCmdClearColorImage ComputeEncoder";
         case kMVKCommandUseTessellationVertexTessCtl: return @"vkCmdDraw (vertex and tess control stages) ComputeEncoder";
         case kMVKCommandUseMultiviewInstanceCountAdjust: return @"vkCmdDraw (multiview instance count adjustment) ComputeEncoder";
         case kMVKCommandUseCopyQueryPoolResults:return @"vkCmdCopyQueryPoolResults ComputeEncoder";
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
index 52e4704..eac9f3f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
@@ -109,6 +109,11 @@
 	/** Returns a MTLComputePipelineState for filling a buffer. */
 	id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
 
+#if MVK_MACOS
+	/** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */
+	id<MTLComputePipelineState> getCmdClearColorImageMTLComputePipelineState(MVKFormatType type);
+#endif
+
 	/** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
 	id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
 
@@ -151,6 +156,9 @@
     id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState = nil;
     id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
 	id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
+#if MVK_MACOS
+	id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil};
+#endif
 	id<MTLComputePipelineState> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
 	id<MTLComputePipelineState> _mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[2] = {nil, nil};
 	id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil};
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
index da0e661..75c8c48 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
@@ -102,6 +102,30 @@
 	MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState(_commandPool));
 }
 
+#if MVK_MACOS
+static inline uint32_t getClearStateIndex(MVKFormatType type) {
+	switch (type) {
+		case kMVKFormatColorHalf:
+		case kMVKFormatColorFloat:
+			return 0;
+		case kMVKFormatColorInt8:
+		case kMVKFormatColorInt16:
+		case kMVKFormatColorInt32:
+			return 1;
+		case kMVKFormatColorUInt8:
+		case kMVKFormatColorUInt16:
+		case kMVKFormatColorUInt32:
+			return 2;
+		default:
+			return 0;
+	}
+}
+
+id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) {
+	MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getClearStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool));
+}
+#endif
+
 id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) {
 	MVK_ENC_REZ_ACCESS(_mtlCopyBufferToImage3DDecompressComputePipelineState[needsTempBuff ? 1 : 0], newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(needsTempBuff, _commandPool));
 }
@@ -178,6 +202,13 @@
     [_mtlFillBufferComputePipelineState release];
     _mtlFillBufferComputePipelineState = nil;
 
+    [_mtlClearColorImageComputePipelineState[0] release];
+    [_mtlClearColorImageComputePipelineState[1] release];
+    [_mtlClearColorImageComputePipelineState[2] release];
+    _mtlClearColorImageComputePipelineState[0] = nil;
+    _mtlClearColorImageComputePipelineState[1] = nil;
+    _mtlClearColorImageComputePipelineState[2] = nil;
+
     [_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release];
     [_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release];
     _mtlCopyBufferToImage3DDecompressComputePipelineState[0] = nil;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
index 124f6d9..68e60d7 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
@@ -91,13 +91,31 @@
     for (size_t i = 0; i < info.size; i++) {                                                                    \n\
         dst[i + info.dstOffset] = src[i + info.srcOffset];                                                      \n\
     }                                                                                                           \n\
-};                                                                                                              \n\
+}                                                                                                               \n\
                                                                                                                 \n\
 kernel void cmdFillBuffer(device uint32_t* dst [[ buffer(0) ]],                                                 \n\
                           constant uint32_t& fillValue [[ buffer(1) ]],                                         \n\
                           uint pos [[thread_position_in_grid]]) {                                               \n\
     dst[pos] = fillValue;                                                                                       \n\
-};                                                                                                              \n\
+}                                                                                                               \n\
+                                                                                                                \n\
+kernel void cmdClearColorImage2DFloat(texture2d<float, access::write> dst [[ texture(0) ]],                     \n\
+                                      constant float4& clearValue [[ buffer(0) ]],                              \n\
+                                      uint2 pos [[thread_position_in_grid]]) {                                  \n\
+    dst.write(clearValue, pos);                                                                                 \n\
+}                                                                                                               \n\
+                                                                                                                \n\
+kernel void cmdClearColorImage2DUInt(texture2d<uint, access::write> dst [[ texture(0) ]],                       \n\
+                                     constant uint4& clearValue [[ buffer(0) ]],                                \n\
+                                     uint2 pos [[thread_position_in_grid]]) {                                   \n\
+    dst.write(clearValue, pos);                                                                                 \n\
+}                                                                                                               \n\
+                                                                                                                \n\
+kernel void cmdClearColorImage2DInt(texture2d<int, access::write> dst [[ texture(0) ]],                         \n\
+                                    constant int4& clearValue [[ buffer(0) ]],                                  \n\
+                                    uint2 pos [[thread_position_in_grid]]) {                                    \n\
+    dst.write(clearValue, pos);                                                                                 \n\
+}                                                                                                               \n\
                                                                                                                 \n\
 typedef struct {                                                                                                \n\
     uint32_t srcRowStride;                                                                                      \n\
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
index 25327ac..7006fe4 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
@@ -417,6 +417,12 @@
 	/** Returns a new MTLComputePipelineState for filling a buffer. */
 	id<MTLComputePipelineState> newCmdFillBufferMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner);
 
+#if MVK_MACOS
+	/** Returns a new MTLComputePipelineState for clearing an image. */
+	id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
+																			 MVKVulkanAPIDeviceObject* owner);
+#endif
+
 	/** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */
 	id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
 																						   MVKVulkanAPIDeviceObject* owner);
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
index a616a64..751fb18 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
@@ -410,6 +410,34 @@
 	return newMTLComputePipelineState("cmdFillBuffer", owner);
 }
 
+#if MVK_MACOS
+id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
+					MVKVulkanAPIDeviceObject* owner) {
+	const char* funcName;
+	switch (type) {
+		case kMVKFormatColorHalf:
+		case kMVKFormatColorFloat:
+			funcName = "cmdClearColorImage2DFloat";
+			break;
+		case kMVKFormatColorInt8:
+		case kMVKFormatColorInt16:
+		case kMVKFormatColorInt32:
+			funcName = "cmdClearColorImage2DInt";
+			break;
+		case kMVKFormatColorUInt8:
+		case kMVKFormatColorUInt16:
+		case kMVKFormatColorUInt32:
+			funcName = "cmdClearColorImage2DUInt";
+			break;
+		default:
+			owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED,
+							   "Format type %u is not supported for clearing with a compute shader.", type);
+			return nil;
+	}
+	return newMTLComputePipelineState(funcName, owner);
+}
+#endif
+
 id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
 																												  MVKVulkanAPIDeviceObject* owner) {
 	return newMTLComputePipelineState(needTempBuf
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
index 65caf00..9fb584e 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
@@ -477,6 +477,14 @@
 
 		mvkEnableFlags(mtlUsage, MTLTextureUsageShaderWrite);
 	}
+#if MVK_MACOS
+    // Clearing a linear image may use shader writes.
+    if (mvkIsAnyFlagEnabled(vkImageUsageFlags, VK_IMAGE_USAGE_TRANSFER_DST_BIT) &&
+        mvkIsAnyFlagEnabled(mtlFmtCaps, kMVKMTLFmtCapsWrite) && isLinear) {
+
+		mvkEnableFlags(mtlUsage, MTLTextureUsageShaderWrite);
+    }
+#endif
 
 	// Render to but only if format supports rendering...
 	if (mvkIsAnyFlagEnabled(vkImageUsageFlags, (VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT |