Merge pull request #613 from jozefkucia/fix-dxt-decompression
Use device address space when decompressing DXT image data.
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];