diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm index 9a920d73..a99d4f69 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm @@ -1599,7 +1599,8 @@ void MVKCmdClearImage::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 mtlClearState = cmdEncoder->getCommandEncodingPool()->getCmdClearColorImageMTLComputePipelineState(pixFmts->getFormatType(_image->getVkFormat())); + const bool isTextureArray = _image->getLayerCount() != 1u; + id mtlClearState = cmdEncoder->getCommandEncodingPool()->getCmdClearColorImageMTLComputePipelineState(pixFmts->getFormatType(_image->getVkFormat()), isTextureArray); id mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseClearColorImage, true); [mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"]; [mtlComputeEnc setComputePipelineState: mtlClearState]; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h index 6bab224e..708a3d03 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h @@ -110,10 +110,10 @@ public: id getCmdFillBufferMTLComputePipelineState(); /** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */ - id getCmdClearColorImageMTLComputePipelineState(MVKFormatType type); + id getCmdClearColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray); /** Returns a MTLComputePipelineState for resolving an image. */ - id getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type); + id getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray); /** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */ id getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff); @@ -166,8 +166,10 @@ protected: id _mtlCopyBufferBytesComputePipelineState = nil; id _mtlFillBufferComputePipelineState = nil; id _mtlDrawIndirectPopulateIndexesComputePipelineState = nil; - id _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil}; - id _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 _mtlClearColorImageComputePipelineState[kColorImageCount] = {nil, nil, nil, nil, nil, nil}; + id _mtlResolveColorImageComputePipelineState[kColorImageCount] = {nil, nil, nil, nil, nil, nil}; id _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil}; id _mtlDrawIndirectConvertBuffersComputePipelineState[2] = {nil, nil}; id _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil}; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm index a3308302..b7b9ec65 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm @@ -109,30 +109,32 @@ id 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 MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) { - MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool)); +id MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) { + MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool, isTextureArray)); } -id MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type) { - MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool)); +id MVKCommandEncodingPool::getCmdResolveColorImageMTLComputePipelineState(MVKFormatType type, bool isTextureArray) { + MVK_ENC_REZ_ACCESS(_mtlResolveColorImageComputePipelineState[getRenderpassLoadStoreStateIndex(type, isTextureArray)], newCmdResolveColorImageMTLComputePipelineState(type, _commandPool, isTextureArray)); } id 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]; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h index 7d5a7cee..75d35038 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h @@ -121,36 +121,84 @@ kernel void cmdClearColorImage2DFloat(texture2d dst [[ tex dst.write(clearValue, pos); \n\ } \n\ \n\ +kernel void cmdClearColorImage2DFloatArray(texture2d_array 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 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 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 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 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 dst [[ texture(0) ]], \n\ - texture2d_ms src [[ texture(1) ]], \n\ - uint2 pos [[thread_position_in_grid]]) { \n\ + texture2d_ms 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 dst [[ texture(0) ]], \n\ + texture2d_ms_array 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 dst [[ texture(0) ]], \n\ texture2d_ms 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 dst [[ texture(0) ]], \n\ + texture2d_ms_array 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 dst [[ texture(0) ]], \n\ texture2d_ms 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 dst [[ texture(0) ]], \n\ + texture2d_ms_array 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\ diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h index 1953c20a..4b7c2937 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h @@ -454,11 +454,13 @@ public: /** Returns a new MTLComputePipelineState for clearing an image. */ id newCmdClearColorImageMTLComputePipelineState(MVKFormatType type, - MVKVulkanAPIDeviceObject* owner); + MVKVulkanAPIDeviceObject* owner, + bool isTextureArray); /** Returns a new MTLComputePipelineState for resolving an image. */ id 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 newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf, diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm index d5b72845..74597fdd 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm @@ -523,22 +523,23 @@ id MVKCommandResourceFactory::newCmdFillBufferMTLComput } id 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 MVKCommandResourceFactory::newCmdClearColorImageMTLC } id 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); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm index 3954ce89..6fabf293 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm @@ -406,7 +406,8 @@ void MVKRenderSubpass::resolveUnresolvableAttachments(MVKCommandEncoder* cmdEnco if ( !mvkAreAllFlagsEnabled(pixFmts->getCapabilities(raImgView->getMTLPixelFormat()), kMVKMTLFmtCapsResolve) ) { MVKFormatType mvkFmtType = _renderPass->getPixelFormats()->getFormatType(raImgView->getMTLPixelFormat()); - id mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType); + const bool isTextureArray = raImgView->getImage()->getLayerCount() != 1u; + id mtlRslvState = cmdEncoder->getCommandEncodingPool()->getCmdResolveColorImageMTLComputePipelineState(mvkFmtType, isTextureArray); id mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseResolveImage); [mtlComputeEnc setComputePipelineState: mtlRslvState]; [mtlComputeEnc setTexture: raImgView->getMTLTexture() atIndex: 0];