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

Fix unresolvable layered compute resolve

Browse Source 
Implementation did not correctly handle array textures when resolving
unresolvable textures like r8uint formats. Add kernel modules to
correctly resolve through the compute pass.

Fixes CTS failures in test families:
dEQP-VK.api.image_clearing.core.clear_color_attachment.*
dEQP-VK.renderpass.suballocation.multisample_resolve.*
This commit is contained in:
Aitor Camacho 2024-01-19 13:53:19 +01:00
parent c4f90e84b2
commit 18f06de878
7 changed files with 92 additions and 40 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
View File

@ -1599,7 +1599,8 @@ void MVKCmdClearImage<N>::encode(MVKCommandEncoder* cmdEncoder) {
// 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()));
const bool isTextureArray = _image->getLayerCount() != 1u;
id<MTLComputePipelineState> mtlClearState = cmdEncoder->getCommandEncodingPool()->getCmdClearColorImageMTLComputePipelineState(pixFmts->getFormatType(_image->getVkFormat()), isTextureArray);
id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseClearColorImage, true);
[mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"];
[mtlComputeEnc setComputePipelineState: mtlClearState];

10
MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
View File

@ -110,10 +110,10 @@ public:
id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
/** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */
id<MTLComputePipelineState> getCmdClearColorImageMTLComputePipelineState(MVKFormatType type);
id<MTLComputePipelineState> getCmdClearColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray);
/** Returns a MTLComputePipelineState for resolving an image. */
id<MTLComputePipelineState> getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type);
id<MTLComputePipelineState> getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray);
/** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
@ -166,8 +166,10 @@ protected:
id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
id<MTLComputePipelineState> _mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil};
id<MTLComputePipelineState> _mtlResolveColorImageComputePipelineState[3] = {nil, nil, nil};
// 6 slots, first 3 for non array textures, last 3 for array textures
static constexpr uint32_t kColorImageCount = 6u;
id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[kColorImageCount] = {nil, nil, nil, nil, nil, nil};
id<MTLComputePipelineState> _mtlResolveColorImageComputePipelineState[kColorImageCount] = {nil, nil, nil, nil, nil, nil};
id<MTLComputePipelineState> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectConvertBuffersComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil};

40
MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
View File

@ -109,30 +109,32 @@ id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdFillBufferMTLComputePi
MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState(_commandPool));
}
static constexpr uint32_t getRenderpassLoadStoreStateIndex(MVKFormatType type) {
static constexpr uint32_t getRenderpassLoadStoreStateIndex(MVKFormatType type, bool isTextureArray) {
// Kernels for array textures are stored from slot 3 onwards
uint32_t layeredOffset = isTextureArray ? 3u : 0u;
switch (type) {
case kMVKFormatColorHalf:
case kMVKFormatColorFloat:
return 0;
return 0 + layeredOffset;
case kMVKFormatColorInt8:
case kMVKFormatColorInt16:
case kMVKFormatColorInt32:
return 1;
return 1 + layeredOffset;
case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32:
return 2;
return 2 + layeredOffset;
default:
return 0;
return 0 + layeredOffset;
}
}
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) {
MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool));
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) {
MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool, isTextureArray));
}
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type) {
MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool));
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) {
MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool, isTextureArray));
}
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) {
@ -224,19 +226,13 @@ void MVKCommandEncodingPool::destroyMetalResources() {
[_mtlDrawIndirectPopulateIndexesComputePipelineState release];
_mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
[_mtlClearColorImageComputePipelineState[0] release];
[_mtlClearColorImageComputePipelineState[1] release];
[_mtlClearColorImageComputePipelineState[2] release];
_mtlClearColorImageComputePipelineState[0] = nil;
_mtlClearColorImageComputePipelineState[1] = nil;
_mtlClearColorImageComputePipelineState[2] = nil;
[_mtlResolveColorImageComputePipelineState[0] release];
[_mtlResolveColorImageComputePipelineState[1] release];
[_mtlResolveColorImageComputePipelineState[2] release];
_mtlResolveColorImageComputePipelineState[0] = nil;
_mtlResolveColorImageComputePipelineState[1] = nil;
_mtlResolveColorImageComputePipelineState[2] = nil;
for (uint32_t i = 0; i < kColorImageCount; i++)
{
[_mtlClearColorImageComputePipelineState[i] release];
_mtlClearColorImageComputePipelineState[i] = nil;
[_mtlResolveColorImageComputePipelineState[i] release];
_mtlResolveColorImageComputePipelineState[i] = nil;
}
[_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release];
[_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release];

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

@ -121,36 +121,84 @@ kernel void cmdClearColorImage2DFloat(texture2d<float, access::write> dst [[ tex
dst.write(clearValue, pos); \n\
} \n\
\n\
kernel void cmdClearColorImage2DFloatArray(texture2d_array<float, access::write> dst [[ texture(0) ]], \n\
constant float4& clearValue [[ buffer(0) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < dst.get_array_size(); ++i) { \n\
dst.write(clearValue, pos, i); \n\
} \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 cmdClearColorImage2DUIntArray(texture2d_array<uint, access::write> dst [[ texture(0) ]], \n\
constant uint4& clearValue [[ buffer(0) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < dst.get_array_size(); ++i) { \n\
dst.write(clearValue, pos, i); \n\
} \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\
kernel void cmdClearColorImage2DIntArray(texture2d_array<int, access::write> dst [[ texture(0) ]], \n\
constant int4& clearValue [[ buffer(0) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < dst.get_array_size(); ++i) { \n\
dst.write(clearValue, pos, i); \n\
} \n\
} \n\
\n\
kernel void cmdResolveColorImage2DFloat(texture2d<float, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<float, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
texture2d_ms<float, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\
} \n\
\n\
kernel void cmdResolveColorImage2DFloatArray(texture2d_array<float, access::write> dst [[ texture(0) ]], \n\
texture2d_ms_array<float, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < src.get_array_size(); ++i) { \n\
dst.write(src.read(pos, i, 0), pos, i); \n\
} \n\
} \n\
\n\
kernel void cmdResolveColorImage2DUInt(texture2d<uint, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<uint, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\
} \n\
\n\
kernel void cmdResolveColorImage2DUIntArray(texture2d_array<uint, access::write> dst [[ texture(0) ]], \n\
texture2d_ms_array<uint, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < src.get_array_size(); ++i) { \n\
dst.write(src.read(pos, i, 0), pos, i); \n\
} \n\
} \n\
\n\
kernel void cmdResolveColorImage2DInt(texture2d<int, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<int, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\
} \n\
\n\
kernel void cmdResolveColorImage2DIntArray(texture2d_array<int, access::write> dst [[ texture(0) ]], \n\
texture2d_ms_array<int, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\
for (uint i = 0u; i < src.get_array_size(); ++i) { \n\
dst.write(src.read(pos, i, 0), pos, i); \n\
} \n\
} \n\
\n\
typedef struct { \n\
uint32_t srcRowStride; \n\
uint32_t srcRowStrideHigh; \n\

6
MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
View File

@ -454,11 +454,13 @@ public:
/** Returns a new MTLComputePipelineState for clearing an image. */
id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner);
MVKVulkanAPIDeviceObject* owner,
bool isTextureArray);
/** Returns a new MTLComputePipelineState for resolving an image. */
id<MTLComputePipelineState> newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner);
MVKVulkanAPIDeviceObject* owner,
bool isTextureArray);
/** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */
id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,

18
MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
View File

@ -523,22 +523,23 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdFillBufferMTLComput
}
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner) {
MVKVulkanAPIDeviceObject* owner,
bool isTextureArray) {
const char* funcName;
switch (type) {
case kMVKFormatColorHalf:
case kMVKFormatColorFloat:
funcName = "cmdClearColorImage2DFloat";
funcName = isTextureArray ? "cmdClearColorImage2DFloatArray" : "cmdClearColorImage2DFloat";
break;
case kMVKFormatColorInt8:
case kMVKFormatColorInt16:
case kMVKFormatColorInt32:
funcName = "cmdClearColorImage2DInt";
funcName = isTextureArray ? "cmdClearColorImage2DIntArray" : "cmdClearColorImage2DInt";
break;
case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32:
funcName = "cmdClearColorImage2DUInt";
funcName = isTextureArray ? "cmdClearColorImage2DUIntArray" : "cmdClearColorImage2DUInt";
break;
default:
owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for clearing with a compute shader.", type);
@ -548,22 +549,23 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLC
}
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner) {
MVKVulkanAPIDeviceObject* owner,
bool isTextureArray) {
const char* funcName;
switch (type) {
case kMVKFormatColorHalf:
case kMVKFormatColorFloat:
funcName = "cmdResolveColorImage2DFloat";
funcName = isTextureArray ? "cmdResolveColorImage2DFloatArray" : "cmdResolveColorImage2DFloat";
break;
case kMVKFormatColorInt8:
case kMVKFormatColorInt16:
case kMVKFormatColorInt32:
funcName = "cmdResolveColorImage2DInt";
funcName = isTextureArray ? "cmdResolveColorImage2DIntArray" : "cmdResolveColorImage2DInt";
break;
case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32:
funcName = "cmdResolveColorImage2DUInt";
funcName = isTextureArray ? "cmdResolveColorImage2DUIntArray" : "cmdResolveColorImage2DUInt";
break;
default:
owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for resolving with a compute shader.", type);

3
MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
View File

@ -406,7 +406,8 @@ void MVKRenderSubpass::resolveUnresolvableAttachments(MVKCommandEncoder* cmdEnco
if ( !mvkAreAllFlagsEnabled(pixFmts->getCapabilities(raImgView->getMTLPixelFormat()), kMVKMTLFmtCapsResolve) ) {
MVKFormatType mvkFmtType = _renderPass->getPixelFormats()->getFormatType(raImgView->getMTLPixelFormat());
id<MTLComputePipelineState> mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType);
const bool isTextureArray = raImgView->getImage()->getLayerCount() != 1u;
id<MTLComputePipelineState> mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType, isTextureArray);
id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseResolveImage);
[mtlComputeEnc setComputePipelineState: mtlRslvState];
[mtlComputeEnc setTexture: raImgView->getMTLTexture() atIndex: 0];