diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
index 6794ab9..6726212 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
@@ -119,7 +119,7 @@
     VkExtent3D extent;                                                                                          \n\
 } CmdCopyBufferToImageInfo;                                                                                     \n\
                                                                                                                 \n\
-kernel void cmdCopyBufferToImage3DDecompressDXTn(constant uint8_t* src [[buffer(0)]],                           \n\
+kernel void cmdCopyBufferToImage3DDecompressDXTn(const device uint8_t* src [[buffer(0)]],                       \n\
                                                  texture3d<float, access::write> dest [[texture(0)]],           \n\
                                                  constant CmdCopyBufferToImageInfo& info [[buffer(2)]],         \n\
                                                  uint3 pos [[thread_position_in_grid]]) {                       \n\
@@ -144,7 +144,7 @@
     }                                                                                                           \n\
 }                                                                                                               \n\
                                                                                                                 \n\
-kernel void cmdCopyBufferToImage3DDecompressTempBufferDXTn(constant uint8_t* src [[buffer(0)]],                 \n\
+kernel void cmdCopyBufferToImage3DDecompressTempBufferDXTn(const device uint8_t* src [[buffer(0)]],             \n\
                                                            device uint8_t* dest [[buffer(1)]],                  \n\
                                                            constant CmdCopyBufferToImageInfo& info [[buffer(2)]],\n\
                                                            uint3 pos [[thread_position_in_grid]]) {             \n\
diff --git a/MoltenVK/MoltenVK/Utility/MVKDXTnCodec.def b/MoltenVK/MoltenVK/Utility/MVKDXTnCodec.def
index 22cf9da..3448ff4 100644
--- a/MoltenVK/MoltenVK/Utility/MVKDXTnCodec.def
+++ b/MoltenVK/MoltenVK/Utility/MVKDXTnCodec.def
@@ -75,8 +75,8 @@
 		return select(pow((colour + 0.055)/1.055, float3(2.4)), colour/12.92, colour <= 0.04045);
 	}
 
-	static void decompressDXTnBlock(constant void* pSrc, thread void* pDest, VkExtent2D extent, VkDeviceSize destRowPitch, VkFormat format) {
-		constant uint32_t* pSrcBlock = (constant uint32_t *)pSrc;
+	static void decompressDXTnBlock(const device void* pSrc, thread void* pDest, VkExtent2D extent, VkDeviceSize destRowPitch, VkFormat format) {
+		const device uint32_t* pSrcBlock = (const device uint32_t *)pSrc;
 		bool isBC1Alpha = false;
 		float3 colourTable[4];
 		float alphaTable[8];