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

Merge pull request #2132 from aitor-lunarg/fix-unresolvable-layered-compute-resolve

Browse Source 
Fix unresolvable layered compute resolve
This commit is contained in:
Bill Hollings 2024-01-25 11:05:10 -05:00 committed by GitHub
commit 3301fdc0f9
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
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. // 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. // Luckily for us, linear images only have one mip and one array layer under Metal.
assert( !isDS ); 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); id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseClearColorImage, true);
[mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"]; [mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"];
[mtlComputeEnc setComputePipelineState: mtlClearState]; [mtlComputeEnc setComputePipelineState: mtlClearState];

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

@ -110,10 +110,10 @@ public:
id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState(); id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
/** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */ /** 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. */ /** 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. */ /** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff); id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
@ -166,8 +166,10 @@ protected:
id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil; id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil; id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
id<MTLComputePipelineState> _mtlDrawIndirectPopulateIndexesComputePipelineState = nil; id<MTLComputePipelineState> _mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil}; // 6 slots, first 3 for non array textures, last 3 for array textures
id<MTLComputePipelineState> _mtlResolveColorImageComputePipelineState[3] = {nil, nil, nil}; 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> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectConvertBuffersComputePipelineState[2] = {nil, nil}; id<MTLComputePipelineState> _mtlDrawIndirectConvertBuffersComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[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)); 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) { switch (type) {
case kMVKFormatColorHalf: case kMVKFormatColorHalf:
case kMVKFormatColorFloat: case kMVKFormatColorFloat:
return 0; return 0 + layeredOffset;
case kMVKFormatColorInt8: case kMVKFormatColorInt8:
case kMVKFormatColorInt16: case kMVKFormatColorInt16:
case kMVKFormatColorInt32: case kMVKFormatColorInt32:
return 1; return 1 + layeredOffset;
case kMVKFormatColorUInt8: case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16: case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32: case kMVKFormatColorUInt32:
return 2; return 2 + layeredOffset;
default: default:
return 0; return 0 + layeredOffset;
} }
} }
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) { id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) {
MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool)); MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool, isTextureArray));
} }
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type) { id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) {
MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool)); MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool, isTextureArray));
} }
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) { id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) {
@ -224,19 +226,13 @@ void MVKCommandEncodingPool::destroyMetalResources() {
[_mtlDrawIndirectPopulateIndexesComputePipelineState release]; [_mtlDrawIndirectPopulateIndexesComputePipelineState release];
_mtlDrawIndirectPopulateIndexesComputePipelineState = nil; _mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
[_mtlClearColorImageComputePipelineState[0] release]; for (uint32_t i = 0; i < kColorImageCount; i++)
[_mtlClearColorImageComputePipelineState[1] release]; {
[_mtlClearColorImageComputePipelineState[2] release]; [_mtlClearColorImageComputePipelineState[i] release];
_mtlClearColorImageComputePipelineState[0] = nil; _mtlClearColorImageComputePipelineState[i] = nil;
_mtlClearColorImageComputePipelineState[1] = nil; [_mtlResolveColorImageComputePipelineState[i] release];
_mtlClearColorImageComputePipelineState[2] = nil; _mtlResolveColorImageComputePipelineState[i] = nil;
}
[_mtlResolveColorImageComputePipelineState[0] release];
[_mtlResolveColorImageComputePipelineState[1] release];
[_mtlResolveColorImageComputePipelineState[2] release];
_mtlResolveColorImageComputePipelineState[0] = nil;
_mtlResolveColorImageComputePipelineState[1] = nil;
_mtlResolveColorImageComputePipelineState[2] = nil;
[_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release]; [_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release];
[_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release]; [_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release];

48
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\ dst.write(clearValue, pos); \n\
} \n\ } \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\ kernel void cmdClearColorImage2DUInt(texture2d<uint, access::write> dst [[ texture(0) ]], \n\
constant uint4& clearValue [[ buffer(0) ]], \n\ constant uint4& clearValue [[ buffer(0) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\ uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(clearValue, pos); \n\ dst.write(clearValue, pos); \n\
} \n\ } \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\ kernel void cmdClearColorImage2DInt(texture2d<int, access::write> dst [[ texture(0) ]], \n\
constant int4& clearValue [[ buffer(0) ]], \n\ constant int4& clearValue [[ buffer(0) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\ uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(clearValue, pos); \n\ dst.write(clearValue, pos); \n\
} \n\ } \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\ kernel void cmdResolveColorImage2DFloat(texture2d<float, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<float, access::read> src [[ texture(1) ]], \n\ texture2d_ms<float, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\ uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\ dst.write(src.read(pos, 0), pos); \n\
} \n\ } \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\ kernel void cmdResolveColorImage2DUInt(texture2d<uint, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<uint, access::read> src [[ texture(1) ]], \n\ texture2d_ms<uint, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\ uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\ dst.write(src.read(pos, 0), pos); \n\
} \n\ } \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\ kernel void cmdResolveColorImage2DInt(texture2d<int, access::write> dst [[ texture(0) ]], \n\
texture2d_ms<int, access::read> src [[ texture(1) ]], \n\ texture2d_ms<int, access::read> src [[ texture(1) ]], \n\
uint2 pos [[thread_position_in_grid]]) { \n\ uint2 pos [[thread_position_in_grid]]) { \n\
dst.write(src.read(pos, 0), pos); \n\ dst.write(src.read(pos, 0), pos); \n\
} \n\ } \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\ typedef struct { \n\
uint32_t srcRowStride; \n\ uint32_t srcRowStride; \n\
uint32_t srcRowStrideHigh; \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. */ /** Returns a new MTLComputePipelineState for clearing an image. */
id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type, id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner); MVKVulkanAPIDeviceObject* owner,
bool isTextureArray);
/** Returns a new MTLComputePipelineState for resolving an image. */ /** Returns a new MTLComputePipelineState for resolving an image. */
id<MTLComputePipelineState> newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, 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. */ /** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */
id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf, 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, id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner) { MVKVulkanAPIDeviceObject* owner,
bool isTextureArray) {
const char* funcName; const char* funcName;
switch (type) { switch (type) {
case kMVKFormatColorHalf: case kMVKFormatColorHalf:
case kMVKFormatColorFloat: case kMVKFormatColorFloat:
funcName = "cmdClearColorImage2DFloat"; funcName = isTextureArray ? "cmdClearColorImage2DFloatArray" : "cmdClearColorImage2DFloat";
break; break;
case kMVKFormatColorInt8: case kMVKFormatColorInt8:
case kMVKFormatColorInt16: case kMVKFormatColorInt16:
case kMVKFormatColorInt32: case kMVKFormatColorInt32:
funcName = "cmdClearColorImage2DInt"; funcName = isTextureArray ? "cmdClearColorImage2DIntArray" : "cmdClearColorImage2DInt";
break; break;
case kMVKFormatColorUInt8: case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16: case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32: case kMVKFormatColorUInt32:
funcName = "cmdClearColorImage2DUInt"; funcName = isTextureArray ? "cmdClearColorImage2DUIntArray" : "cmdClearColorImage2DUInt";
break; break;
default: default:
owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for clearing with a compute shader.", type); 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, id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdResolveColorImageMTLComputePipelineState(MVKFormatType type,
MVKVulkanAPIDeviceObject* owner) { MVKVulkanAPIDeviceObject* owner,
bool isTextureArray) {
const char* funcName; const char* funcName;
switch (type) { switch (type) {
case kMVKFormatColorHalf: case kMVKFormatColorHalf:
case kMVKFormatColorFloat: case kMVKFormatColorFloat:
funcName = "cmdResolveColorImage2DFloat"; funcName = isTextureArray ? "cmdResolveColorImage2DFloatArray" : "cmdResolveColorImage2DFloat";
break; break;
case kMVKFormatColorInt8: case kMVKFormatColorInt8:
case kMVKFormatColorInt16: case kMVKFormatColorInt16:
case kMVKFormatColorInt32: case kMVKFormatColorInt32:
funcName = "cmdResolveColorImage2DInt"; funcName = isTextureArray ? "cmdResolveColorImage2DIntArray" : "cmdResolveColorImage2DInt";
break; break;
case kMVKFormatColorUInt8: case kMVKFormatColorUInt8:
case kMVKFormatColorUInt16: case kMVKFormatColorUInt16:
case kMVKFormatColorUInt32: case kMVKFormatColorUInt32:
funcName = "cmdResolveColorImage2DUInt"; funcName = isTextureArray ? "cmdResolveColorImage2DUIntArray" : "cmdResolveColorImage2DUInt";
break; break;
default: default:
owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED, "Format type %u is not supported for resolving with a compute shader.", type); 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) ) { if ( !mvkAreAllFlagsEnabled(pixFmts->getCapabilities(raImgView->getMTLPixelFormat()), kMVKMTLFmtCapsResolve) ) {
MVKFormatType mvkFmtType = _renderPass->getPixelFormats()->getFormatType(raImgView->getMTLPixelFormat()); 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); id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseResolveImage);
[mtlComputeEnc setComputePipelineState: mtlRslvState]; [mtlComputeEnc setComputePipelineState: mtlRslvState];
[mtlComputeEnc setTexture: raImgView->getMTLTexture() atIndex: 0]; [mtlComputeEnc setTexture: raImgView->getMTLTexture() atIndex: 0];