Ryujinx/moltenvk
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Use device address space when decompressing DXT image data.

Browse Source 
Buffers in constant address space must be aligned to 256 bytes on macOS.

This fixes Metal validation errors:

Compute Function(cmdCopyBufferToImage3DDecompressDXTn): the offset into the buffer src that is bound at buffer index 0 must be a multiple of 256 but was set to 12928.
Compute Function(cmdCopyBufferToImage3DDecompressDXTn): the offset into the buffer src that is bound at buffer index 0 must be a multiple of 256 but was set to 78464.
This commit is contained in:
Józef Kucia 2019-05-27 11:13:42 +02:00
parent 6b556a9625
commit 1f8967a617
2 changed files with 4 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
View File

@ -119,7 +119,7 @@ typedef struct {
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 @@ kernel void cmdCopyBufferToImage3DDecompressDXTn(constant uint8_t* src [[buffer(
} \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\

4
MoltenVK/MoltenVK/Utility/MVKDXTnCodec.def
View File

@ -75,8 +75,8 @@ MVK_DECOMPRESS_CODE(
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];