diff --git a/Docs/MoltenVK_Runtime_UserGuide.md b/Docs/MoltenVK_Runtime_UserGuide.md index 805671a3..d61496e7 100644 --- a/Docs/MoltenVK_Runtime_UserGuide.md +++ b/Docs/MoltenVK_Runtime_UserGuide.md @@ -329,6 +329,7 @@ In addition to core *Vulkan* functionality, **MoltenVK** also supports the foll - `VK_EXT_memory_budget` *(requires Metal 2.0)* - `VK_EXT_metal_objects` - `VK_EXT_metal_surface` +- `VK_EXT_pipeline_creation_cache_control` - `VK_EXT_post_depth_coverage` *(iOS and macOS, requires family 4 (A11) or better Apple GPU)* - `VK_EXT_private_data ` - `VK_EXT_robustness2` diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md index 32ff22e1..5c895704 100644 --- a/Docs/Whats_New.md +++ b/Docs/Whats_New.md @@ -19,6 +19,7 @@ MoltenVK 1.2.3 Released TBA - Add support for extensions: + - `VK_EXT_pipeline_creation_cache_control` - `VK_EXT_swapchain_maintenance1` - `VK_EXT_surface_maintenance1` - Fix crash when `VkCommandBufferInheritanceInfo::renderPass` is `VK_NULL_HANDLE` during dynamic rendering. @@ -33,6 +34,16 @@ Released TBA - Queue submissions retain wait semaphores until `MTLCommandBuffer` finishes. - Use a different visibility buffer for each `MTLCommandBuffer` in a queue submit. - Work around problems with using explicit LoD with arrayed depth images on Apple Silicon. +- Reduce memory footprint of MSL source code retained in pipeline cache. +- Add `MVKConfiguration::shaderSourceCompressionAlgorithm` and + env var `MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM` to support + compressing MSL shader source code held in a pipeline cache. +- Add `MVKShaderCompilationPerformance::mslCompress` and `mslDecompress` + to allow performance of MSL compression to be tracked and queried. +- Add support for logging performance stats accumulated in a `VkDevice`, when it is destroyed. +- Change `MVKConfiguration::logActivityPerformanceInline` boolean to `activityPerformanceLoggingStyle` enumeration value. +- Add `MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE` environment variable and + build setting to set `MVKConfiguration::activityPerformanceLoggingStyle` value. - Update `VK_MVK_MOLTENVK_SPEC_VERSION` to version `37`. diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj index 88783912..2c4b87eb 100644 --- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj +++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj @@ -112,7 +112,7 @@ 2FEA0AAA24902F9F00EEF3AD /* MVKOSExtensions.mm in Sources */ = {isa = PBXBuildFile; fileRef = A9B51BD2225E986A00AC74D2 /* MVKOSExtensions.mm */; }; 2FEA0AAB24902F9F00EEF3AD /* MVKShaderModule.mm in Sources */ = {isa = PBXBuildFile; fileRef = A94FB7981C7DFB4800632CA3 /* MVKShaderModule.mm */; }; 2FEA0AAC24902F9F00EEF3AD /* MVKSync.mm in Sources */ = {isa = PBXBuildFile; fileRef = A94FB79E1C7DFB4800632CA3 /* MVKSync.mm */; }; - 2FEA0AAD24902F9F00EEF3AD /* MVKCodec.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.cpp */; }; + 2FEA0AAD24902F9F00EEF3AD /* MVKCodec.mm in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.mm */; }; 2FEA0AAE24902F9F00EEF3AD /* MVKCmdPipeline.mm in Sources */ = {isa = PBXBuildFile; fileRef = A94FB76F1C7DFB4800632CA3 /* MVKCmdPipeline.mm */; }; 2FEA0AAF24902F9F00EEF3AD /* MVKLayers.mm in Sources */ = {isa = PBXBuildFile; fileRef = A94FB7A11C7DFB4800632CA3 /* MVKLayers.mm */; }; 2FEA0AB024902F9F00EEF3AD /* MVKFramebuffer.mm in Sources */ = {isa = PBXBuildFile; fileRef = A94FB7881C7DFB4800632CA3 /* MVKFramebuffer.mm */; }; @@ -136,8 +136,8 @@ 4553AEFC2251617100E8EBCD /* MVKBlockObserver.m in Sources */ = {isa = PBXBuildFile; fileRef = 4553AEF62251617100E8EBCD /* MVKBlockObserver.m */; }; 4553AEFD2251617100E8EBCD /* MVKBlockObserver.h in Headers */ = {isa = PBXBuildFile; fileRef = 4553AEFA2251617100E8EBCD /* MVKBlockObserver.h */; }; 4553AEFE2251617100E8EBCD /* MVKBlockObserver.h in Headers */ = {isa = PBXBuildFile; fileRef = 4553AEFA2251617100E8EBCD /* MVKBlockObserver.h */; }; - 45557A5221C9EFF3008868BD /* MVKCodec.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.cpp */; }; - 45557A5321C9EFF3008868BD /* MVKCodec.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.cpp */; }; + 45557A5221C9EFF3008868BD /* MVKCodec.mm in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.mm */; }; + 45557A5321C9EFF3008868BD /* MVKCodec.mm in Sources */ = {isa = PBXBuildFile; fileRef = 45557A4D21C9EFF3008868BD /* MVKCodec.mm */; }; 45557A5421C9EFF3008868BD /* MVKCodec.h in Headers */ = {isa = PBXBuildFile; fileRef = 45557A5121C9EFF3008868BD /* MVKCodec.h */; }; 45557A5521C9EFF3008868BD /* MVKCodec.h in Headers */ = {isa = PBXBuildFile; fileRef = 45557A5121C9EFF3008868BD /* MVKCodec.h */; }; A9096E5E1F81E16300DFBEA6 /* MVKCmdDispatch.mm in Sources */ = {isa = PBXBuildFile; fileRef = A9096E5D1F81E16300DFBEA6 /* MVKCmdDispatch.mm */; }; @@ -428,7 +428,7 @@ 453638312508A4C7000EFFD3 /* MTLRenderPassDepthAttachmentDescriptor+MoltenVK.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = "MTLRenderPassDepthAttachmentDescriptor+MoltenVK.h"; sourceTree = ""; }; 4553AEF62251617100E8EBCD /* MVKBlockObserver.m */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.objc; path = MVKBlockObserver.m; sourceTree = ""; }; 4553AEFA2251617100E8EBCD /* MVKBlockObserver.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKBlockObserver.h; sourceTree = ""; }; - 45557A4D21C9EFF3008868BD /* MVKCodec.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKCodec.cpp; sourceTree = ""; }; + 45557A4D21C9EFF3008868BD /* MVKCodec.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCodec.mm; sourceTree = ""; }; 45557A5121C9EFF3008868BD /* MVKCodec.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCodec.h; sourceTree = ""; }; 45557A5721CD83C3008868BD /* MVKDXTnCodec.def */ = {isa = PBXFileReference; explicitFileType = sourcecode.cpp.h; fileEncoding = 4; path = MVKDXTnCodec.def; sourceTree = ""; }; A9096E5C1F81E16300DFBEA6 /* MVKCmdDispatch.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKCmdDispatch.h; sourceTree = ""; }; @@ -691,8 +691,8 @@ A9D7104E25CDE05E00E38106 /* MVKBitArray.h */, 4553AEFA2251617100E8EBCD /* MVKBlockObserver.h */, 4553AEF62251617100E8EBCD /* MVKBlockObserver.m */, - 45557A4D21C9EFF3008868BD /* MVKCodec.cpp */, 45557A5121C9EFF3008868BD /* MVKCodec.h */, + 45557A4D21C9EFF3008868BD /* MVKCodec.mm */, 45557A5721CD83C3008868BD /* MVKDXTnCodec.def */, A9A5E9C525C0822700E9085E /* MVKEnvironment.cpp */, A98149431FB6A3F7005F00B4 /* MVKEnvironment.h */, @@ -1368,7 +1368,7 @@ 2FEA0AAA24902F9F00EEF3AD /* MVKOSExtensions.mm in Sources */, 2FEA0AAB24902F9F00EEF3AD /* MVKShaderModule.mm in Sources */, 2FEA0AAC24902F9F00EEF3AD /* MVKSync.mm in Sources */, - 2FEA0AAD24902F9F00EEF3AD /* MVKCodec.cpp in Sources */, + 2FEA0AAD24902F9F00EEF3AD /* MVKCodec.mm in Sources */, 2FEA0AAE24902F9F00EEF3AD /* MVKCmdPipeline.mm in Sources */, 2FEA0AAF24902F9F00EEF3AD /* MVKLayers.mm in Sources */, 2FEA0AB024902F9F00EEF3AD /* MVKFramebuffer.mm in Sources */, @@ -1427,7 +1427,7 @@ A9B51BD7225E986A00AC74D2 /* MVKOSExtensions.mm in Sources */, A94FB80E1C7DFB4800632CA3 /* MVKShaderModule.mm in Sources */, A94FB81A1C7DFB4800632CA3 /* MVKSync.mm in Sources */, - 45557A5221C9EFF3008868BD /* MVKCodec.cpp in Sources */, + 45557A5221C9EFF3008868BD /* MVKCodec.mm in Sources */, A94FB7BE1C7DFB4800632CA3 /* MVKCmdPipeline.mm in Sources */, A94FB81E1C7DFB4800632CA3 /* MVKLayers.mm in Sources */, A94FB7EE1C7DFB4800632CA3 /* MVKFramebuffer.mm in Sources */, @@ -1487,7 +1487,7 @@ A9B51BD8225E986A00AC74D2 /* MVKOSExtensions.mm in Sources */, A94FB80F1C7DFB4800632CA3 /* MVKShaderModule.mm in Sources */, A94FB81B1C7DFB4800632CA3 /* MVKSync.mm in Sources */, - 45557A5321C9EFF3008868BD /* MVKCodec.cpp in Sources */, + 45557A5321C9EFF3008868BD /* MVKCodec.mm in Sources */, A94FB7BF1C7DFB4800632CA3 /* MVKCmdPipeline.mm in Sources */, A94FB81F1C7DFB4800632CA3 /* MVKLayers.mm in Sources */, A94FB7EF1C7DFB4800632CA3 /* MVKFramebuffer.mm in Sources */, diff --git a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h index 662cea11..7377f05f 100644 --- a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h +++ b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h @@ -130,6 +130,24 @@ typedef enum MVKConfigFastMath { MVK_CONFIG_FAST_MATH_MAX_ENUM = 0x7FFFFFFF } MVKConfigFastMath; +/** Identifies available system data compression algorithms. */ +typedef enum MVKConfigCompressionAlgorithm { + MVK_CONFIG_COMPRESSION_ALGORITHM_NONE = 0, /**< No compression. */ + MVK_CONFIG_COMPRESSION_ALGORITHM_LZFSE = 1, /**< Apple proprietary. Good balance of high performance and small compression size, particularly for larger data content. */ + MVK_CONFIG_COMPRESSION_ALGORITHM_ZLIB = 2, /**< Open cross-platform ZLib format. For smaller data content, has better performance and smaller size than LZFSE. */ + MVK_CONFIG_COMPRESSION_ALGORITHM_LZ4 = 3, /**< Fastest performance. Largest compression size. */ + MVK_CONFIG_COMPRESSION_ALGORITHM_LZMA = 4, /**< Slowest performance. Smallest compression size, particular with larger content. */ + MVK_CONFIG_COMPRESSION_ALGORITHM_MAX_ENUM = 0x7FFFFFFF, +} MVKConfigCompressionAlgorithm; + +/** Identifies the style of activity performance logging to use. */ +typedef enum MVKConfigActivityPerformanceLoggingStyle { + MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_FRAME_COUNT = 0, /**< Repeatedly log performance after a configured number of frames. */ + MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_IMMEDIATE = 1, /**< Log immediately after each performance measurement. */ + MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_DEVICE_LIFETIME = 2, /**< Log at the end of the VkDevice lifetime. This is useful for one-shot apps such as testing frameworks. */ + MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_MAX_ENUM = 0x7FFFFFFF, +} MVKConfigActivityPerformanceLoggingStyle; + /** * MoltenVK configuration settings. * @@ -361,8 +379,8 @@ typedef struct { * If enabled, performance statistics, as defined by the MVKPerformanceStatistics structure, * are collected, and can be retrieved via the vkGetPerformanceStatisticsMVK() function. * - * You can also use the performanceLoggingFrameCount or logActivityPerformanceInline - * parameters to automatically log the performance statistics collected by this parameter. + * You can also use the activityPerformanceLoggingStyle and performanceLoggingFrameCount + * parameters to configure when to log the performance statistics collected by this parameter. * * The value of this parameter must be changed before creating a VkDevice, * for the change to take effect. @@ -770,21 +788,20 @@ typedef struct { VkBool32 useMTLHeap; /** - * Controls whether MoltenVK should log the performance of individual activities as they happen. - * If this setting is enabled, activity performance will be logged when each activity happens. - * If this setting is disabled, activity performance will be logged when frame peformance is - * logged as determined by the performanceLoggingFrameCount value. + * Controls when MoltenVK should log activity performance events. * * The value of this parameter must be changed before creating a VkDevice, * for the change to take effect. * * The initial value or this parameter is set by the - * MVK_CONFIG_PERFORMANCE_LOGGING_INLINE + * MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE * runtime environment variable or MoltenVK compile-time build setting. - * If neither is set, this setting is disabled by default, and activity - * performance will be logged only when frame activity is logged. + * If neither is set, this setting is set to + * MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_FRAME_COUNT by default, + * and activity performance will be logged when frame activity is logged. */ - VkBool32 logActivityPerformanceInline; + MVKConfigActivityPerformanceLoggingStyle activityPerformanceLoggingStyle; +#define logActivityPerformanceInline activityPerformanceLoggingStyle /** * Controls the Vulkan API version that MoltenVK should advertise in vkEnumerateInstanceVersion(). @@ -877,6 +894,27 @@ typedef struct { */ MVKUseMetalArgumentBuffers useMetalArgumentBuffers; + /** + * Controls the type of compression to use on the MSL source code that is stored in memory + * for use in a pipeline cache. After being converted from SPIR-V, or loaded directly into + * a VkShaderModule, and then compiled into a MTLLibrary, the MSL source code is no longer + * needed for operation, but it is retained so it can be written out as part of a pipeline + * cache export. When a large number of shaders are loaded, this can consume significant + * memory. In such a case, this parameter can be used to compress the MSL source code that + * is awaiting export as part of a pipeline cache. + * + * The value of this parameter can be changed at any time, and will affect the size of + * the cached MSL from subsequent shader compilations. + * + * The initial value or this parameter is set by the + * MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM + * runtime environment variable or MoltenVK compile-time build setting. + * If neither is set, this setting is set to + * MVK_CONFIG_COMPRESSION_ALGORITHM_NONE by default, + * and MoltenVK will not compress the MSL source code after compilation into a MTLLibrary. + */ + MVKConfigCompressionAlgorithm shaderSourceCompressionAlgorithm; + } MVKConfiguration; /** Identifies the type of rounding Metal uses for float to integer conversions in particular calculatons. */ @@ -999,6 +1037,8 @@ typedef struct { MVKPerformanceTracker spirvToMSL; /** Convert SPIR-V to MSL source code. */ MVKPerformanceTracker mslCompile; /** Compile MSL source code into a MTLLibrary. */ MVKPerformanceTracker mslLoad; /** Load pre-compiled MSL code into a MTLLibrary. */ + MVKPerformanceTracker mslCompress; /** Compress MSL source code after compiling a MTLLibrary, to hold it in a pipeline cache. */ + MVKPerformanceTracker mslDecompress; /** Decompress MSL source code to write the MSL when serializing a pipeline cache. */ MVKPerformanceTracker shaderLibraryFromCache; /** Retrieve a shader library from the cache, lazily creating it if needed. */ MVKPerformanceTracker functionRetrieval; /** Retrieve a MTLFunction from a MTLLibrary. */ MVKPerformanceTracker functionSpecialization; /** Specialize a retrieved MTLFunction. */ @@ -1220,9 +1260,8 @@ VKAPI_ATTR void VKAPI_CALL vkGetVersionStringsMVK( /** * Sets the number of threads in a workgroup for a compute kernel. * - * This needs to be called if you are creating compute shader modules from MSL - * source code or MSL compiled code. Workgroup size is determined automatically - * if you're using SPIR-V. + * This needs to be called if you are creating compute shader modules from MSL source code + * or MSL compiled code. If you are using SPIR-V, workgroup size is determined automatically. * * This function is not supported by the Vulkan SDK Loader and Layers framework * and is unavailable when using the Vulkan SDK Loader and Layers framework. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index f58d9ee3..c9b0e97e 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -686,7 +686,9 @@ public: // Log call not locked. Very minor chance that the tracker data will be updated during log call, // resulting in an inconsistent report. Not worth taking lock perf hit for rare inline reporting. - if (_logActivityPerformanceInline) { logActivityPerformance(activityTracker, _performanceStatistics, true); } + if (_activityPerformanceLoggingStyle == MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_IMMEDIATE) { + logActivityPerformance(activityTracker, _performanceStatistics, true); + } } }; @@ -891,7 +893,7 @@ protected: id _defaultMTLSamplerState = nil; id _dummyBlitMTLBuffer = nil; uint32_t _globalVisibilityQueryCount = 0; - bool _logActivityPerformanceInline = false; + MVKConfigActivityPerformanceLoggingStyle _activityPerformanceLoggingStyle = MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_FRAME_COUNT; bool _isPerformanceTracking = false; bool _isCurrentlyAutoGPUCapturing = false; bool _isUsingMetalArgumentBuffers = false; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 11ab9443..efa11fea 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -390,6 +390,11 @@ void MVKPhysicalDevice::getFeatures(VkPhysicalDeviceFeatures2* features) { swapchainMaintenance1Features->swapchainMaintenance1 = true; break; } + case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PIPELINE_CREATION_CACHE_CONTROL_FEATURES_EXT: { + auto* pipelineCreationCacheControlFeatures = (VkPhysicalDevicePipelineCreationCacheControlFeaturesEXT*)next; + pipelineCreationCacheControlFeatures->pipelineCreationCacheControl = true; + break; + } case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_TEXEL_BUFFER_ALIGNMENT_FEATURES_EXT: { auto* texelBuffAlignFeatures = (VkPhysicalDeviceTexelBufferAlignmentFeaturesEXT*)next; texelBuffAlignFeatures->texelBufferAlignment = _metalFeatures.texelBuffers && [_mtlDevice respondsToSelector: @selector(minimumLinearTextureAlignmentForPixelFormat:)]; @@ -3726,10 +3731,16 @@ VkResult MVKDevice::createPipelines(VkPipelineCache pipelineCache, const PipelineInfoType* pCreateInfos, const VkAllocationCallbacks* pAllocator, VkPipeline* pPipelines) { + bool ignoreFurtherPipelines = false; VkResult rslt = VK_SUCCESS; MVKPipelineCache* mvkPLC = (MVKPipelineCache*)pipelineCache; for (uint32_t plIdx = 0; plIdx < count; plIdx++) { + + // Ensure all slots are purposefully set. + pPipelines[plIdx] = VK_NULL_HANDLE; + if (ignoreFurtherPipelines) { continue; } + const PipelineInfoType* pCreateInfo = &pCreateInfos[plIdx]; // See if this pipeline has a parent. This can come either directly @@ -3742,18 +3753,19 @@ VkResult MVKDevice::createPipelines(VkPipelineCache pipelineCache, parentPL = vkParentPL ? (MVKPipeline*)vkParentPL : VK_NULL_HANDLE; } - // Create the pipeline and if creation was successful, insert the new pipeline - // in the return array and add it to the pipeline cache (if the cache was specified). - // If creation was unsuccessful, insert NULL into the return array, change the - // result code of this function, and destroy the broken pipeline. + // Create the pipeline and if creation was successful, insert the new pipeline in the return array. MVKPipeline* mvkPL = new PipelineType(this, mvkPLC, parentPL, pCreateInfo); VkResult plRslt = mvkPL->getConfigurationResult(); if (plRslt == VK_SUCCESS) { pPipelines[plIdx] = (VkPipeline)mvkPL; } else { - rslt = plRslt; - pPipelines[plIdx] = VK_NULL_HANDLE; - mvkPL->destroy(); + // If creation was unsuccessful, destroy the broken pipeline, change the result + // code of this function, and if the VK_PIPELINE_CREATE_EARLY_RETURN_ON_FAILURE_BIT + // flag is set, don't build any further pipelines. + mvkPL->destroy(); + if (rslt == VK_SUCCESS) { rslt = plRslt; } + ignoreFurtherPipelines = (_enabledPipelineCreationCacheControlFeatures.pipelineCreationCacheControl && + mvkIsAnyFlagEnabled(pCreateInfo->flags, VK_PIPELINE_CREATE_EARLY_RETURN_ON_FAILURE_BIT)); } } @@ -4000,7 +4012,7 @@ void MVKDevice::logActivityPerformance(MVKPerformanceTracker& activity, MVKPerfo } void MVKDevice::logPerformanceSummary() { - if (_logActivityPerformanceInline) { return; } + if (_activityPerformanceLoggingStyle == MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_IMMEDIATE) { return; } // Get a copy to minimize time under lock MVKPerformanceStatistics perfStats; @@ -4014,6 +4026,8 @@ void MVKDevice::logPerformanceSummary() { logActivityPerformance(perfStats.shaderCompilation.spirvToMSL, perfStats); logActivityPerformance(perfStats.shaderCompilation.mslCompile, perfStats); logActivityPerformance(perfStats.shaderCompilation.mslLoad, perfStats); + logActivityPerformance(perfStats.shaderCompilation.mslCompress, perfStats); + logActivityPerformance(perfStats.shaderCompilation.mslDecompress, perfStats); logActivityPerformance(perfStats.shaderCompilation.shaderLibraryFromCache, perfStats); logActivityPerformance(perfStats.shaderCompilation.functionRetrieval, perfStats); logActivityPerformance(perfStats.shaderCompilation.functionSpecialization, perfStats); @@ -4028,6 +4042,8 @@ const char* MVKDevice::getActivityPerformanceDescription(MVKPerformanceTracker& if (&activity == &perfStats.shaderCompilation.spirvToMSL) { return "Convert SPIR-V to MSL source code"; } if (&activity == &perfStats.shaderCompilation.mslCompile) { return "Compile MSL source code into a MTLLibrary"; } if (&activity == &perfStats.shaderCompilation.mslLoad) { return "Load pre-compiled MSL code into a MTLLibrary"; } + if (&activity == &perfStats.shaderCompilation.mslCompress) { return "Compress MSL source code after compiling a MTLLibrary"; } + if (&activity == &perfStats.shaderCompilation.mslDecompress) { return "Decompress MSL source code during pipeline cache write"; } if (&activity == &perfStats.shaderCompilation.shaderLibraryFromCache) { return "Retrieve shader library from the cache"; } if (&activity == &perfStats.shaderCompilation.functionRetrieval) { return "Retrieve a MTLFunction from a MTLLibrary"; } if (&activity == &perfStats.shaderCompilation.functionSpecialization) { return "Specialize a retrieved MTLFunction"; } @@ -4377,29 +4393,25 @@ MVKDevice::MVKDevice(MVKPhysicalDevice* physicalDevice, const VkDeviceCreateInfo void MVKDevice::initPerformanceTracking() { _isPerformanceTracking = mvkConfig().performanceTracking; - _logActivityPerformanceInline = mvkConfig().logActivityPerformanceInline; + _activityPerformanceLoggingStyle = mvkConfig().activityPerformanceLoggingStyle; - MVKPerformanceTracker initPerf; - initPerf.count = 0; - initPerf.averageDuration = 0.0; - initPerf.minimumDuration = 0.0; - initPerf.maximumDuration = 0.0; - - _performanceStatistics.shaderCompilation.hashShaderCode = initPerf; - _performanceStatistics.shaderCompilation.spirvToMSL = initPerf; - _performanceStatistics.shaderCompilation.mslCompile = initPerf; - _performanceStatistics.shaderCompilation.mslLoad = initPerf; - _performanceStatistics.shaderCompilation.shaderLibraryFromCache = initPerf; - _performanceStatistics.shaderCompilation.functionRetrieval = initPerf; - _performanceStatistics.shaderCompilation.functionSpecialization = initPerf; - _performanceStatistics.shaderCompilation.pipelineCompile = initPerf; - _performanceStatistics.pipelineCache.sizePipelineCache = initPerf; - _performanceStatistics.pipelineCache.writePipelineCache = initPerf; - _performanceStatistics.pipelineCache.readPipelineCache = initPerf; - _performanceStatistics.queue.mtlQueueAccess = initPerf; - _performanceStatistics.queue.mtlCommandBufferCompletion = initPerf; - _performanceStatistics.queue.nextCAMetalDrawable = initPerf; - _performanceStatistics.queue.frameInterval = initPerf; + _performanceStatistics.shaderCompilation.hashShaderCode = {}; + _performanceStatistics.shaderCompilation.spirvToMSL = {}; + _performanceStatistics.shaderCompilation.mslCompile = {}; + _performanceStatistics.shaderCompilation.mslLoad = {}; + _performanceStatistics.shaderCompilation.mslCompress = {}; + _performanceStatistics.shaderCompilation.mslDecompress = {}; + _performanceStatistics.shaderCompilation.shaderLibraryFromCache = {}; + _performanceStatistics.shaderCompilation.functionRetrieval = {}; + _performanceStatistics.shaderCompilation.functionSpecialization = {}; + _performanceStatistics.shaderCompilation.pipelineCompile = {}; + _performanceStatistics.pipelineCache.sizePipelineCache = {}; + _performanceStatistics.pipelineCache.writePipelineCache = {}; + _performanceStatistics.pipelineCache.readPipelineCache = {}; + _performanceStatistics.queue.mtlQueueAccess = {}; + _performanceStatistics.queue.mtlCommandBufferCompletion = {}; + _performanceStatistics.queue.nextCAMetalDrawable = {}; + _performanceStatistics.queue.frameInterval = {}; } void MVKDevice::initPhysicalDevice(MVKPhysicalDevice* physicalDevice, const VkDeviceCreateInfo* pCreateInfo) { @@ -4666,9 +4678,15 @@ void MVKDevice::reservePrivateData(const VkDeviceCreateInfo* pCreateInfo) { } MVKDevice::~MVKDevice() { + if (_activityPerformanceLoggingStyle == MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_DEVICE_LIFETIME) { + MVKLogInfo("Device activity performance summary:"); + logPerformanceSummary(); + } + for (auto& queues : _queuesByQueueFamilyIndex) { mvkDestroyContainerContents(queues); } + if (_commandResourceFactory) { _commandResourceFactory->destroy(); } [_globalVisibilityResultMTLBuffer release]; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def index c6de213e..05086327 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def @@ -35,39 +35,40 @@ #pragma push_macro("INTEL") #undef INTEL -MVK_DEVICE_FEATURE(16BitStorage, 16BIT_STORAGE, 4) -MVK_DEVICE_FEATURE(8BitStorage, 8BIT_STORAGE, 3) -MVK_DEVICE_FEATURE(BufferDeviceAddress, BUFFER_DEVICE_ADDRESS, 3) -MVK_DEVICE_FEATURE(DescriptorIndexing, DESCRIPTOR_INDEXING, 20) -MVK_DEVICE_FEATURE(DynamicRendering, DYNAMIC_RENDERING, 1) -MVK_DEVICE_FEATURE(HostQueryReset, HOST_QUERY_RESET, 1) -MVK_DEVICE_FEATURE(ImagelessFramebuffer, IMAGELESS_FRAMEBUFFER, 1) -MVK_DEVICE_FEATURE(ImageRobustness, IMAGE_ROBUSTNESS, 1) -MVK_DEVICE_FEATURE(InlineUniformBlock, INLINE_UNIFORM_BLOCK, 2) -MVK_DEVICE_FEATURE(Multiview, MULTIVIEW, 3) -MVK_DEVICE_FEATURE(PrivateData, PRIVATE_DATA, 1) -MVK_DEVICE_FEATURE(ProtectedMemory, PROTECTED_MEMORY, 1) -MVK_DEVICE_FEATURE(SamplerYcbcrConversion, SAMPLER_YCBCR_CONVERSION, 1) -MVK_DEVICE_FEATURE(ScalarBlockLayout, SCALAR_BLOCK_LAYOUT, 1) -MVK_DEVICE_FEATURE(SeparateDepthStencilLayouts, SEPARATE_DEPTH_STENCIL_LAYOUTS, 1) -MVK_DEVICE_FEATURE(ShaderDrawParameters, SHADER_DRAW_PARAMETERS, 1) -MVK_DEVICE_FEATURE(ShaderAtomicInt64, SHADER_ATOMIC_INT64, 2) -MVK_DEVICE_FEATURE(ShaderFloat16Int8, SHADER_FLOAT16_INT8, 2) -MVK_DEVICE_FEATURE(ShaderSubgroupExtendedTypes, SHADER_SUBGROUP_EXTENDED_TYPES, 1) -MVK_DEVICE_FEATURE(SubgroupSizeControl, SUBGROUP_SIZE_CONTROL, 2) -MVK_DEVICE_FEATURE(TextureCompressionASTCHDR, TEXTURE_COMPRESSION_ASTC_HDR, 1) -MVK_DEVICE_FEATURE(TimelineSemaphore, TIMELINE_SEMAPHORE, 1) -MVK_DEVICE_FEATURE(UniformBufferStandardLayout, UNIFORM_BUFFER_STANDARD_LAYOUT, 1) -MVK_DEVICE_FEATURE(VariablePointer, VARIABLE_POINTER, 2) -MVK_DEVICE_FEATURE(VulkanMemoryModel, VULKAN_MEMORY_MODEL, 3) -MVK_DEVICE_FEATURE_EXTN(FragmentShaderBarycentric, FRAGMENT_SHADER_BARYCENTRIC, KHR, 1) -MVK_DEVICE_FEATURE_EXTN(PortabilitySubset, PORTABILITY_SUBSET, KHR, 15) -MVK_DEVICE_FEATURE_EXTN(FragmentShaderInterlock, FRAGMENT_SHADER_INTERLOCK, EXT, 3) -MVK_DEVICE_FEATURE_EXTN(Robustness2, ROBUSTNESS_2, EXT, 3) -MVK_DEVICE_FEATURE_EXTN(SwapchainMaintenance1, SWAPCHAIN_MAINTENANCE_1, EXT, 1) -MVK_DEVICE_FEATURE_EXTN(TexelBufferAlignment, TEXEL_BUFFER_ALIGNMENT, EXT, 1) -MVK_DEVICE_FEATURE_EXTN(VertexAttributeDivisor, VERTEX_ATTRIBUTE_DIVISOR, EXT, 2) -MVK_DEVICE_FEATURE_EXTN(ShaderIntegerFunctions2, SHADER_INTEGER_FUNCTIONS_2, INTEL, 1) +MVK_DEVICE_FEATURE(16BitStorage, 16BIT_STORAGE, 4) +MVK_DEVICE_FEATURE(8BitStorage, 8BIT_STORAGE, 3) +MVK_DEVICE_FEATURE(BufferDeviceAddress, BUFFER_DEVICE_ADDRESS, 3) +MVK_DEVICE_FEATURE(DescriptorIndexing, DESCRIPTOR_INDEXING, 20) +MVK_DEVICE_FEATURE(DynamicRendering, DYNAMIC_RENDERING, 1) +MVK_DEVICE_FEATURE(HostQueryReset, HOST_QUERY_RESET, 1) +MVK_DEVICE_FEATURE(ImagelessFramebuffer, IMAGELESS_FRAMEBUFFER, 1) +MVK_DEVICE_FEATURE(ImageRobustness, IMAGE_ROBUSTNESS, 1) +MVK_DEVICE_FEATURE(InlineUniformBlock, INLINE_UNIFORM_BLOCK, 2) +MVK_DEVICE_FEATURE(Multiview, MULTIVIEW, 3) +MVK_DEVICE_FEATURE(PrivateData, PRIVATE_DATA, 1) +MVK_DEVICE_FEATURE(ProtectedMemory, PROTECTED_MEMORY, 1) +MVK_DEVICE_FEATURE(SamplerYcbcrConversion, SAMPLER_YCBCR_CONVERSION, 1) +MVK_DEVICE_FEATURE(ScalarBlockLayout, SCALAR_BLOCK_LAYOUT, 1) +MVK_DEVICE_FEATURE(SeparateDepthStencilLayouts, SEPARATE_DEPTH_STENCIL_LAYOUTS, 1) +MVK_DEVICE_FEATURE(ShaderDrawParameters, SHADER_DRAW_PARAMETERS, 1) +MVK_DEVICE_FEATURE(ShaderAtomicInt64, SHADER_ATOMIC_INT64, 2) +MVK_DEVICE_FEATURE(ShaderFloat16Int8, SHADER_FLOAT16_INT8, 2) +MVK_DEVICE_FEATURE(ShaderSubgroupExtendedTypes, SHADER_SUBGROUP_EXTENDED_TYPES, 1) +MVK_DEVICE_FEATURE(SubgroupSizeControl, SUBGROUP_SIZE_CONTROL, 2) +MVK_DEVICE_FEATURE(TextureCompressionASTCHDR, TEXTURE_COMPRESSION_ASTC_HDR, 1) +MVK_DEVICE_FEATURE(TimelineSemaphore, TIMELINE_SEMAPHORE, 1) +MVK_DEVICE_FEATURE(UniformBufferStandardLayout, UNIFORM_BUFFER_STANDARD_LAYOUT, 1) +MVK_DEVICE_FEATURE(VariablePointer, VARIABLE_POINTER, 2) +MVK_DEVICE_FEATURE(VulkanMemoryModel, VULKAN_MEMORY_MODEL, 3) +MVK_DEVICE_FEATURE_EXTN(FragmentShaderBarycentric, FRAGMENT_SHADER_BARYCENTRIC, KHR, 1) +MVK_DEVICE_FEATURE_EXTN(PortabilitySubset, PORTABILITY_SUBSET, KHR, 15) +MVK_DEVICE_FEATURE_EXTN(FragmentShaderInterlock, FRAGMENT_SHADER_INTERLOCK, EXT, 3) +MVK_DEVICE_FEATURE_EXTN(Robustness2, ROBUSTNESS_2, EXT, 3) +MVK_DEVICE_FEATURE_EXTN(PipelineCreationCacheControl, PIPELINE_CREATION_CACHE_CONTROL, EXT, 1) +MVK_DEVICE_FEATURE_EXTN(SwapchainMaintenance1, SWAPCHAIN_MAINTENANCE_1, EXT, 1) +MVK_DEVICE_FEATURE_EXTN(TexelBufferAlignment, TEXEL_BUFFER_ALIGNMENT, EXT, 1) +MVK_DEVICE_FEATURE_EXTN(VertexAttributeDivisor, VERTEX_ATTRIBUTE_DIVISOR, EXT, 2) +MVK_DEVICE_FEATURE_EXTN(ShaderIntegerFunctions2, SHADER_INTEGER_FUNCTIONS_2, INTEL, 1) #pragma pop_macro("INTEL") #undef MVK_DEVICE_FEATURE_EXTN diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h index f08eb11c..caffd5c2 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h @@ -155,8 +155,18 @@ public: /** Returns the number of descriptor sets in this pipeline layout. */ uint32_t getDescriptorSetCount() { return _descriptorSetCount; } + /** Returns the pipeline cache used by this pipeline. */ + MVKPipelineCache* getPipelineCache() { return _pipelineCache; } + + /** Returns whether the pipeline creation fail if a pipeline compile is required. */ + bool shouldFailOnPipelineCompileRequired() { + return (_device->_enabledPipelineCreationCacheControlFeatures.pipelineCreationCacheControl && + mvkIsAnyFlagEnabled(_flags, VK_PIPELINE_CREATE_FAIL_ON_PIPELINE_COMPILE_REQUIRED_BIT)); + } + /** Constructs an instance for the device. layout, and parent (which may be NULL). */ - MVKPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipelineLayout* layout, MVKPipeline* parent); + MVKPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipelineLayout* layout, + VkPipelineCreateFlags flags, MVKPipeline* parent); protected: void propagateDebugName() override {} @@ -172,6 +182,7 @@ protected: MVKShaderImplicitRezBinding _dynamicOffsetBufferIndex; MVKShaderImplicitRezBinding _indirectParamsIndex; MVKShaderImplicitRezBinding _pushConstantsBufferIndex; + VkPipelineCreateFlags _flags; uint32_t _descriptorSetCount; bool _stageUsesPushConstants[kMVKShaderStageCount]; bool _fullImageViewSwizzle; @@ -324,6 +335,9 @@ protected: bool verifyImplicitBuffer(bool needsBuffer, MVKShaderImplicitRezBinding& index, MVKShaderStage stage, const char* name); uint32_t getTranslatedVertexBinding(uint32_t binding, uint32_t translationOffset, uint32_t maxBinding); uint32_t getImplicitBufferIndex(MVKShaderStage stage, uint32_t bufferIndexOffset); + MVKMTLFunction getMTLFunction(SPIRVToMSLConversionConfiguration& shaderConfig, + const VkPipelineShaderStageCreateInfo* pShaderStage, + const char* pStageName); const VkPipelineShaderStageCreateInfo* _pVertexSS = nullptr; const VkPipelineShaderStageCreateInfo* _pTessCtlSS = nullptr; @@ -456,8 +470,14 @@ public: */ VkResult writeData(size_t* pDataSize, void* pData); - /** Return a shader library from the shader conversion configuration and sourced from the specified shader module. */ - MVKShaderLibrary* getShaderLibrary(SPIRVToMSLConversionConfiguration* pContext, MVKShaderModule* shaderModule); + /** + * Return a shader library for the shader conversion configuration, from the + * pipeline's pipeline cache, or compiled from source in the shader module. + */ + MVKShaderLibrary* getShaderLibrary(SPIRVToMSLConversionConfiguration* pContext, + MVKShaderModule* shaderModule, + MVKPipeline* pipeline, + uint64_t startTime = 0); /** Merges the contents of the specified number of pipeline caches into this cache. */ VkResult mergePipelineCaches(uint32_t srcCacheCount, const VkPipelineCache* pSrcCaches); @@ -474,11 +494,18 @@ protected: MVKShaderLibraryCache* getShaderLibraryCache(MVKShaderModuleKey smKey); void readData(const VkPipelineCacheCreateInfo* pCreateInfo); void writeData(std::ostream& outstream, bool isCounting = false); + MVKShaderLibrary* getShaderLibraryImpl(SPIRVToMSLConversionConfiguration* pContext, + MVKShaderModule* shaderModule, + MVKPipeline* pipeline, + uint64_t startTime); + VkResult writeDataImpl(size_t* pDataSize, void* pData); + VkResult mergePipelineCachesImpl(uint32_t srcCacheCount, const VkPipelineCache* pSrcCaches); void markDirty(); std::unordered_map _shaderCache; size_t _dataSize = 0; std::mutex _shaderCacheLock; + bool _isExternallySynchronized = false; }; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm index 30acafc1..c254df95 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm @@ -208,9 +208,11 @@ void MVKPipeline::addMTLArgumentEncoders(MVKMTLFunction& mvkMTLFunc, } } -MVKPipeline::MVKPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipelineLayout* layout, MVKPipeline* parent) : +MVKPipeline::MVKPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipelineLayout* layout, + VkPipelineCreateFlags flags, MVKPipeline* parent) : MVKVulkanAPIDeviceObject(device), _pipelineCache(pipelineCache), + _flags(flags), _descriptorSetCount(layout->getDescriptorSetCount()), _fullImageViewSwizzle(mvkConfig().fullImageViewSwizzle) { @@ -393,7 +395,7 @@ MVKGraphicsPipeline::MVKGraphicsPipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipeline* parent, const VkGraphicsPipelineCreateInfo* pCreateInfo) : - MVKPipeline(device, pipelineCache, (MVKPipelineLayout*)pCreateInfo->layout, parent) { + MVKPipeline(device, pipelineCache, (MVKPipelineLayout*)pCreateInfo->layout, pCreateInfo->flags, parent) { // Determine rasterization early, as various other structs are validated and interpreted in this context. const VkPipelineRenderingCreateInfo* pRendInfo = getRenderingCreateInfo(pCreateInfo); @@ -431,6 +433,14 @@ MVKGraphicsPipeline::MVKGraphicsPipeline(MVKDevice* device, } } + // Tessellation - must ignore allowed bad pTessellationState pointer if not tess pipeline + _outputControlPointCount = reflectData.numControlPoints; + mvkSetOrClear(&_tessInfo, (_pTessCtlSS && _pTessEvalSS) ? pCreateInfo->pTessellationState : nullptr); + + // Render pipeline state. Do this as early as possible, to fail fast if pipeline requires a fail on cache-miss. + initMTLRenderPipelineState(pCreateInfo, reflectData); + if ( !_hasValidMTLPipelineStates ) { return; } + // Track dynamic state const VkPipelineDynamicStateCreateInfo* pDS = pCreateInfo->pDynamicState; if (pDS) { @@ -455,10 +465,6 @@ MVKGraphicsPipeline::MVKGraphicsPipeline(MVKDevice* device, } } - // Tessellation - must ignore allowed bad pTessellationState pointer if not tess pipeline - _outputControlPointCount = reflectData.numControlPoints; - mvkSetOrClear(&_tessInfo, (_pTessCtlSS && _pTessEvalSS) ? pCreateInfo->pTessellationState : nullptr); - // Rasterization _mtlCullMode = MTLCullModeNone; _mtlFrontWinding = MTLWindingCounterClockwise; @@ -481,9 +487,6 @@ MVKGraphicsPipeline::MVKGraphicsPipeline(MVKDevice* device, // Must run after _isRasterizing and _dynamicState are populated initCustomSamplePositions(pCreateInfo); - // Render pipeline state - initMTLRenderPipelineState(pCreateInfo, reflectData); - // Depth stencil content - clearing will disable depth and stencil testing // Must ignore allowed bad pDepthStencilState pointer if rasterization disabled or no depth attachment mvkSetOrClear(&_depthStencilInfo, _isRasterizingDepthStencil ? pCreateInfo->pDepthStencilState : nullptr); @@ -605,8 +608,10 @@ void MVKGraphicsPipeline::initMTLRenderPipelineState(const VkGraphicsPipelineCre } else { getOrCompilePipeline(plDesc, _mtlPipelineState); } + [plDesc release]; // temp release + } else { + _hasValidMTLPipelineStates = false; } - [plDesc release]; // temp release } else { // In this case, we need to create three render pipelines. But, the way Metal handles // index buffers for compute stage-in means we have to defer creation of stage 1 until @@ -621,6 +626,8 @@ void MVKGraphicsPipeline::initMTLRenderPipelineState(const VkGraphicsPipelineCre if (getOrCompilePipeline(tcPLDesc, _mtlTessControlStageState, "Tessellation control")) { getOrCompilePipeline(rastPLDesc, _mtlPipelineState); } + } else { + _hasValidMTLPipelineStates = false; } [tcPLDesc release]; // temp release [rastPLDesc release]; // temp release @@ -910,13 +917,10 @@ bool MVKGraphicsPipeline::addVertexShaderToPipeline(MTLRenderPipelineDescriptor* shaderConfig.options.mslOptions.disable_rasterization = !_isRasterizing; addVertexInputToShaderConversionConfig(shaderConfig, pCreateInfo); - MVKMTLFunction func = ((MVKShaderModule*)_pVertexSS->module)->getMTLFunction(&shaderConfig, _pVertexSS->pSpecializationInfo, _pipelineCache); + MVKMTLFunction func = getMTLFunction(shaderConfig, _pVertexSS, "Vertex"); id mtlFunc = func.getMTLFunction(); - if ( !mtlFunc ) { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Vertex shader function could not be compiled into pipeline. See previous logged error.")); - return false; - } plDesc.vertexFunction = mtlFunc; + if ( !mtlFunc ) { return false; } auto& funcRslts = func.shaderConversionResults; plDesc.rasterizationEnabled = !funcRslts.isRasterizationDisabled; @@ -975,22 +979,19 @@ bool MVKGraphicsPipeline::addVertexShaderToPipeline(MTLComputePipelineDescriptor addVertexInputToShaderConversionConfig(shaderConfig, pCreateInfo); addNextStageInputToShaderConversionConfig(shaderConfig, tcInputs); + // We need to compile this function three times, with no indexing, 16-bit indices, and 32-bit indices. static const CompilerMSL::Options::IndexType indexTypes[] = { CompilerMSL::Options::IndexType::None, CompilerMSL::Options::IndexType::UInt16, CompilerMSL::Options::IndexType::UInt32, }; - // We need to compile this function three times, with no indexing, 16-bit indices, and 32-bit indices. MVKMTLFunction func; for (uint32_t i = 0; i < sizeof(indexTypes)/sizeof(indexTypes[0]); i++) { shaderConfig.options.mslOptions.vertex_index_type = indexTypes[i]; - func = ((MVKShaderModule*)_pVertexSS->module)->getMTLFunction(&shaderConfig, _pVertexSS->pSpecializationInfo, _pipelineCache); + func = getMTLFunction(shaderConfig, _pVertexSS, "Vertex"); id mtlFunc = func.getMTLFunction(); - if ( !mtlFunc ) { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Vertex shader function could not be compiled into pipeline. See previous logged error.")); - return false; - } _mtlTessVertexFunctions[i] = [mtlFunc retain]; + if ( !mtlFunc ) { return false; } auto& funcRslts = func.shaderConversionResults; _needsVertexSwizzleBuffer = funcRslts.needsSwizzleBuffer; @@ -1044,13 +1045,10 @@ bool MVKGraphicsPipeline::addTessCtlShaderToPipeline(MTLComputePipelineDescripto addPrevStageOutputToShaderConversionConfig(shaderConfig, vtxOutputs); addNextStageInputToShaderConversionConfig(shaderConfig, teInputs); - MVKMTLFunction func = ((MVKShaderModule*)_pTessCtlSS->module)->getMTLFunction(&shaderConfig, _pTessCtlSS->pSpecializationInfo, _pipelineCache); + MVKMTLFunction func = getMTLFunction(shaderConfig, _pTessCtlSS, "Tessellation control"); id mtlFunc = func.getMTLFunction(); - if ( !mtlFunc ) { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Tessellation control shader function could not be compiled into pipeline. See previous logged error.")); - return false; - } plDesc.computeFunction = mtlFunc; + if ( !mtlFunc ) { return false; } auto& funcRslts = func.shaderConversionResults; _needsTessCtlSwizzleBuffer = funcRslts.needsSwizzleBuffer; @@ -1105,14 +1103,10 @@ bool MVKGraphicsPipeline::addTessEvalShaderToPipeline(MTLRenderPipelineDescripto shaderConfig.options.mslOptions.disable_rasterization = !_isRasterizing; addPrevStageOutputToShaderConversionConfig(shaderConfig, tcOutputs); - MVKMTLFunction func = ((MVKShaderModule*)_pTessEvalSS->module)->getMTLFunction(&shaderConfig, _pTessEvalSS->pSpecializationInfo, _pipelineCache); + MVKMTLFunction func = getMTLFunction(shaderConfig, _pTessEvalSS, "Tessellation evaluation"); id mtlFunc = func.getMTLFunction(); - if ( !mtlFunc ) { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Tessellation evaluation shader function could not be compiled into pipeline. See previous logged error.")); - return false; - } - // Yeah, you read that right. Tess. eval functions are a kind of vertex function in Metal. - plDesc.vertexFunction = mtlFunc; + plDesc.vertexFunction = mtlFunc; // Yeah, you read that right. Tess. eval functions are a kind of vertex function in Metal. + if ( !mtlFunc ) { return false; } auto& funcRslts = func.shaderConversionResults; plDesc.rasterizationEnabled = !funcRslts.isRasterizationDisabled; @@ -1166,13 +1160,10 @@ bool MVKGraphicsPipeline::addFragmentShaderToPipeline(MTLRenderPipelineDescripto } addPrevStageOutputToShaderConversionConfig(shaderConfig, shaderOutputs); - MVKMTLFunction func = ((MVKShaderModule*)_pFragmentSS->module)->getMTLFunction(&shaderConfig, _pFragmentSS->pSpecializationInfo, _pipelineCache); + MVKMTLFunction func = getMTLFunction(shaderConfig, _pFragmentSS, "Fragment"); id mtlFunc = func.getMTLFunction(); - if ( !mtlFunc ) { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Fragment shader function could not be compiled into pipeline. See previous logged error.")); - return false; - } plDesc.fragmentFunction = mtlFunc; + if ( !mtlFunc ) { return false; } auto& funcRslts = func.shaderConversionResults; _needsFragmentSwizzleBuffer = funcRslts.needsSwizzleBuffer; @@ -1796,6 +1787,23 @@ bool MVKGraphicsPipeline::isRasterizationDisabled(const VkGraphicsPipelineCreate (mvkMTLPrimitiveTopologyClassFromVkPrimitiveTopology(pCreateInfo->pInputAssemblyState->topology) == MTLPrimitiveTopologyClassTriangle)))); } +MVKMTLFunction MVKGraphicsPipeline::getMTLFunction(SPIRVToMSLConversionConfiguration& shaderConfig, + const VkPipelineShaderStageCreateInfo* pShaderStage, + const char* pStageName) { + MVKShaderModule* shaderModule = (MVKShaderModule*)pShaderStage->module; + MVKMTLFunction func = shaderModule->getMTLFunction(&shaderConfig, + pShaderStage->pSpecializationInfo, + this); + if ( !func.getMTLFunction() ) { + if (shouldFailOnPipelineCompileRequired()) { + setConfigurationResult(VK_PIPELINE_COMPILE_REQUIRED); + } else { + setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "%s shader function could not be compiled into pipeline. See previous logged error.", pStageName)); + } + } + return func; +} + MVKGraphicsPipeline::~MVKGraphicsPipeline() { @synchronized (getMTLDevice()) { [_mtlTessVertexStageDesc release]; @@ -1830,7 +1838,7 @@ MVKComputePipeline::MVKComputePipeline(MVKDevice* device, MVKPipelineCache* pipelineCache, MVKPipeline* parent, const VkComputePipelineCreateInfo* pCreateInfo) : - MVKPipeline(device, pipelineCache, (MVKPipelineLayout*)pCreateInfo->layout, parent) { + MVKPipeline(device, pipelineCache, (MVKPipelineLayout*)pCreateInfo->layout, pCreateInfo->flags, parent) { _allowsDispatchBase = mvkAreAllFlagsEnabled(pCreateInfo->flags, VK_PIPELINE_CREATE_DISPATCH_BASE_BIT); @@ -1863,7 +1871,7 @@ MVKComputePipeline::MVKComputePipeline(MVKDevice* device, if ( !_mtlPipelineState ) { _hasValidMTLPipelineStates = false; } } else { - setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Compute shader function could not be compiled into pipeline. See previous logged error.")); + _hasValidMTLPipelineStates = false; } if (_needsSwizzleBuffer && _swizzleBufferIndex.stages[kMVKShaderStageCompute] > _device->_pMetalFeatures->maxPerStageBufferCount) { @@ -1931,8 +1939,14 @@ MVKMTLFunction MVKComputePipeline::getMTLFunction(const VkComputePipelineCreateI shaderConfig.options.mslOptions.dynamic_offsets_buffer_index = _dynamicOffsetBufferIndex.stages[kMVKShaderStageCompute]; shaderConfig.options.mslOptions.indirect_params_buffer_index = _indirectParamsIndex.stages[kMVKShaderStageCompute]; - MVKMTLFunction func = ((MVKShaderModule*)pSS->module)->getMTLFunction(&shaderConfig, pSS->pSpecializationInfo, _pipelineCache); - + MVKMTLFunction func = ((MVKShaderModule*)pSS->module)->getMTLFunction(&shaderConfig, pSS->pSpecializationInfo, this); + if ( !func.getMTLFunction() ) { + if (shouldFailOnPipelineCompileRequired()) { + setConfigurationResult(VK_PIPELINE_COMPILE_REQUIRED); + } else { + setConfigurationResult(reportError(VK_ERROR_INVALID_SHADER_NV, "Compute shader function could not be compiled into pipeline. See previous logged error.")); + } + } auto& funcRslts = func.shaderConversionResults; _needsSwizzleBuffer = funcRslts.needsSwizzleBuffer; _needsBufferSizeBuffer = funcRslts.needsBufferSizeBuffer; @@ -1959,12 +1973,25 @@ MVKComputePipeline::~MVKComputePipeline() { #pragma mark MVKPipelineCache // Return a shader library from the specified shader conversion configuration sourced from the specified shader module. -MVKShaderLibrary* MVKPipelineCache::getShaderLibrary(SPIRVToMSLConversionConfiguration* pContext, MVKShaderModule* shaderModule) { - lock_guard lock(_shaderCacheLock); +MVKShaderLibrary* MVKPipelineCache::getShaderLibrary(SPIRVToMSLConversionConfiguration* pContext, + MVKShaderModule* shaderModule, + MVKPipeline* pipeline, + uint64_t startTime) { + if (_isExternallySynchronized) { + return getShaderLibraryImpl(pContext, shaderModule, pipeline, startTime); + } else { + lock_guard lock(_shaderCacheLock); + return getShaderLibraryImpl(pContext, shaderModule, pipeline, startTime); + } +} +MVKShaderLibrary* MVKPipelineCache::getShaderLibraryImpl(SPIRVToMSLConversionConfiguration* pContext, + MVKShaderModule* shaderModule, + MVKPipeline* pipeline, + uint64_t startTime) { bool wasAdded = false; MVKShaderLibraryCache* slCache = getShaderLibraryCache(shaderModule->getKey()); - MVKShaderLibrary* shLib = slCache->getShaderLibrary(pContext, shaderModule, &wasAdded); + MVKShaderLibrary* shLib = slCache->getShaderLibrary(pContext, shaderModule, pipeline, &wasAdded, startTime); if (wasAdded) { markDirty(); } return shLib; } @@ -2004,8 +2031,8 @@ protected: bool next() { return (++_index < (_pSLCache ? _pSLCache->_shaderLibraries.size() : 0)); } SPIRVToMSLConversionConfiguration& getShaderConversionConfig() { return _pSLCache->_shaderLibraries[_index].first; } - std::string& getMSL() { return _pSLCache->_shaderLibraries[_index].second->_msl; } - SPIRVToMSLConversionResults& getShaderConversionResults() { return _pSLCache->_shaderLibraries[_index].second->_shaderConversionResults; } + MVKCompressor& getCompressedMSL() { return _pSLCache->_shaderLibraries[_index].second->getCompressedMSL(); } + SPIRVToMSLConversionResultInfo& getShaderConversionResultInfo() { return _pSLCache->_shaderLibraries[_index].second->_shaderConversionResultInfo; } MVKShaderCacheIterator(MVKShaderLibraryCache* pSLCache) : _pSLCache(pSLCache) {} MVKShaderLibraryCache* _pSLCache; @@ -2013,14 +2040,21 @@ protected: int32_t _index = -1; }; +VkResult MVKPipelineCache::writeData(size_t* pDataSize, void* pData) { + if (_isExternallySynchronized) { + return writeDataImpl(pDataSize, pData); + } else { + lock_guard lock(_shaderCacheLock); + return writeDataImpl(pDataSize, pData); + } +} + // If pData is not null, serializes at most pDataSize bytes of the contents of the cache into that // memory location, and returns the number of bytes serialized in pDataSize. If pData is null, // returns the number of bytes required to serialize the contents of this pipeline cache. // This is the compliment of the readData() function. The two must be kept aligned. -VkResult MVKPipelineCache::writeData(size_t* pDataSize, void* pData) { +VkResult MVKPipelineCache::writeDataImpl(size_t* pDataSize, void* pData) { #if MVK_USE_CEREAL - lock_guard lock(_shaderCacheLock); - try { if ( !pDataSize ) { return VK_SUCCESS; } @@ -2086,8 +2120,8 @@ void MVKPipelineCache::writeData(ostream& outstream, bool isCounting) { writer(cacheEntryType); writer(smKey); writer(cacheIter.getShaderConversionConfig()); - writer(cacheIter.getShaderConversionResults()); - writer(cacheIter.getMSL()); + writer(cacheIter.getShaderConversionResultInfo()); + writer(cacheIter.getCompressedMSL()); _device->addActivityPerformance(activityTracker, startTime); } } @@ -2149,16 +2183,16 @@ void MVKPipelineCache::readData(const VkPipelineCacheCreateInfo* pCreateInfo) { SPIRVToMSLConversionConfiguration shaderConversionConfig; reader(shaderConversionConfig); - SPIRVToMSLConversionResults shaderConversionResults; - reader(shaderConversionResults); + SPIRVToMSLConversionResultInfo resultInfo; + reader(resultInfo); - string msl; - reader(msl); + MVKCompressor compressedMSL; + reader(compressedMSL); // Add the shader library to the staging cache. MVKShaderLibraryCache* slCache = getShaderLibraryCache(smKey); _device->addActivityPerformance(_device->_performanceStatistics.pipelineCache.readPipelineCache, startTime); - slCache->addShaderLibrary(&shaderConversionConfig, msl, shaderConversionResults); + slCache->addShaderLibrary(&shaderConversionConfig, resultInfo, compressedMSL); break; } @@ -2184,6 +2218,15 @@ void MVKPipelineCache::markDirty() { } VkResult MVKPipelineCache::mergePipelineCaches(uint32_t srcCacheCount, const VkPipelineCache* pSrcCaches) { + if (_isExternallySynchronized) { + return mergePipelineCachesImpl(srcCacheCount, pSrcCaches); + } else { + lock_guard lock(_shaderCacheLock); + return mergePipelineCachesImpl(srcCacheCount, pSrcCaches); + } +} + +VkResult MVKPipelineCache::mergePipelineCachesImpl(uint32_t srcCacheCount, const VkPipelineCache* pSrcCaches) { for (uint32_t srcIdx = 0; srcIdx < srcCacheCount; srcIdx++) { MVKPipelineCache* srcPLC = (MVKPipelineCache*)pSrcCaches[srcIdx]; for (auto& srcPair : srcPLC->_shaderCache) { @@ -2374,7 +2417,7 @@ namespace mvk { } template - void serialize(Archive & archive, SPIRVToMSLConversionResults& scr) { + void serialize(Archive & archive, SPIRVToMSLConversionResultInfo& scr) { archive(scr.entryPoint, scr.isRasterizationDisabled, scr.isPositionInvariant, @@ -2396,10 +2439,21 @@ void serialize(Archive & archive, MVKShaderModuleKey& k) { k.codeHash); } +template +void serialize(Archive & archive, MVKCompressor& comp) { + archive(comp._compressed, + comp._uncompressedSize, + comp._algorithm); +} + #pragma mark Construction -MVKPipelineCache::MVKPipelineCache(MVKDevice* device, const VkPipelineCacheCreateInfo* pCreateInfo) : MVKVulkanAPIDeviceObject(device) { +MVKPipelineCache::MVKPipelineCache(MVKDevice* device, const VkPipelineCacheCreateInfo* pCreateInfo) : + MVKVulkanAPIDeviceObject(device), + _isExternallySynchronized(device->_enabledPipelineCreationCacheControlFeatures.pipelineCreationCacheControl && + mvkIsAnyFlagEnabled(pCreateInfo->flags, VK_PIPELINE_CACHE_CREATE_EXTERNALLY_SYNCHRONIZED_BIT)) { + readData(pCreateInfo); } diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h index 2db3f3c7..a7e3417f 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h @@ -20,6 +20,7 @@ #include "MVKDevice.h" #include "MVKSync.h" +#include "MVKCodec.h" #include "MVKSmallVector.h" #include #include @@ -40,11 +41,11 @@ using namespace mvk; /** A MTLFunction and corresponding result information resulting from a shader conversion. */ typedef struct MVKMTLFunction { - SPIRVToMSLConversionResults shaderConversionResults; + SPIRVToMSLConversionResultInfo shaderConversionResults; MTLSize threadGroupSize; inline id getMTLFunction() { return _mtlFunction; } - MVKMTLFunction(id mtlFunc, const SPIRVToMSLConversionResults scRslts, MTLSize tgSize); + MVKMTLFunction(id mtlFunc, const SPIRVToMSLConversionResultInfo scRslts, MTLSize tgSize); MVKMTLFunction(const MVKMTLFunction& other); MVKMTLFunction& operator=(const MVKMTLFunction& other); MVKMTLFunction() {} @@ -56,7 +57,7 @@ private: } MVKMTLFunction; /** A MVKMTLFunction indicating an invalid MTLFunction. The mtlFunction member is nil. */ -const MVKMTLFunction MVKMTLFunctionNull(nil, SPIRVToMSLConversionResults(), MTLSizeMake(1, 1, 1)); +const MVKMTLFunction MVKMTLFunctionNull(nil, SPIRVToMSLConversionResultInfo(), MTLSizeMake(1, 1, 1)); /** Wraps a single MTLLibrary. */ class MVKShaderLibrary : public MVKBaseObject { @@ -84,12 +85,13 @@ public: */ void setWorkgroupSize(uint32_t x, uint32_t y, uint32_t z); - /** Constructs an instance from the specified MSL source code. */ MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, - const std::string& mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults); + const SPIRVToMSLConversionResult& conversionResult); + + MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, + const SPIRVToMSLConversionResultInfo& resultInfo, + const MVKCompressor compressedMSL); - /** Constructs an instance from the specified compiled MSL code data. */ MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, const void* mslCompiledCodeData, size_t mslCompiledCodeLength); @@ -108,11 +110,15 @@ protected: MVKMTLFunction getMTLFunction(const VkSpecializationInfo* pSpecializationInfo, MVKShaderModule* shaderModule); void handleCompilationError(NSError* err, const char* opDesc); MTLFunctionConstant* getFunctionConstant(NSArray* mtlFCs, NSUInteger mtlFCID); + void compileLibrary(const std::string& msl); + void compressMSL(const std::string& msl); + void decompressMSL(std::string& msl); + MVKCompressor& getCompressedMSL() { return _compressedMSL; } MVKVulkanAPIDeviceObject* _owner; id _mtlLibrary; - SPIRVToMSLConversionResults _shaderConversionResults; - std::string _msl; + MVKCompressor _compressedMSL; + SPIRVToMSLConversionResultInfo _shaderConversionResultInfo; }; @@ -128,15 +134,17 @@ public: MVKVulkanAPIObject* getVulkanAPIObject() override { return _owner->getVulkanAPIObject(); }; /** - * Returns a shader library from the shader conversion configuration sourced from the shader module, - * lazily creating the shader library from source code in the shader module, if needed. + * Returns a shader library from the shader conversion configuration sourced from the + * shader module, lazily creating the shader library from source code in the shader + * module, if needed, and if the pipeline is not configured to fail if a pipeline compile + * is required. In that case, the new shader library is not created, and nil is returned. * * If pWasAdded is not nil, this function will set it to true if a new shader library was created, * and to false if an existing shader library was found and returned. */ MVKShaderLibrary* getShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, - MVKShaderModule* shaderModule, - bool* pWasAdded = nullptr); + MVKShaderModule* shaderModule, MVKPipeline* pipeline, + bool* pWasAdded, uint64_t startTime = 0); MVKShaderLibraryCache(MVKVulkanAPIDeviceObject* owner) : _owner(owner) {}; @@ -147,10 +155,12 @@ protected: friend MVKPipelineCache; friend MVKShaderModule; - MVKShaderLibrary* findShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig); - MVKShaderLibrary* addShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, - const std::string& mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults); + MVKShaderLibrary* findShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, uint64_t startTime = 0); + MVKShaderLibrary* addShaderLibrary(const SPIRVToMSLConversionConfiguration* pShaderConfig, + const SPIRVToMSLConversionResult& conversionResult); + MVKShaderLibrary* addShaderLibrary(const SPIRVToMSLConversionConfiguration* pShaderConfig, + const SPIRVToMSLConversionResultInfo& resultInfo, + const MVKCompressor compressedMSL); void merge(MVKShaderLibraryCache* other); MVKVulkanAPIDeviceObject* _owner; @@ -197,23 +207,15 @@ public: /** Returns the Metal shader function, possibly specialized. */ MVKMTLFunction getMTLFunction(SPIRVToMSLConversionConfiguration* pShaderConfig, const VkSpecializationInfo* pSpecializationInfo, - MVKPipelineCache* pipelineCache); + MVKPipeline* pipeline); /** Convert the SPIR-V to MSL, using the specified shader conversion configuration. */ - bool convert(SPIRVToMSLConversionConfiguration* pShaderConfig); + bool convert(SPIRVToMSLConversionConfiguration* pShaderConfig, + SPIRVToMSLConversionResult& conversionResult); /** Returns the original SPIR-V code that was specified when this object was created. */ const std::vector& getSPIRV() { return _spvConverter.getSPIRV(); } - /** - * Returns the Metal Shading Language source code as converted by the most recent - * call to convert() function, or set directly using the setMSL() function. - */ - const std::string& getMSL() { return _spvConverter.getMSL(); } - - /** Returns information about the shader conversion results. */ - const SPIRVToMSLConversionResults& getConversionResults() { return _spvConverter.getConversionResults(); } - /** Sets the number of threads in a single compute kernel workgroup, per dimension. */ void setWorkgroupSize(uint32_t x, uint32_t y, uint32_t z); @@ -258,7 +260,7 @@ public: * nanoseconds, an error will be generated and logged, and nil will be returned. */ id newMTLLibrary(NSString* mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults); + const SPIRVToMSLConversionResultInfo& shaderConversionResults); #pragma mark Construction diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm index 3de7df9e..2eb57ebb 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm @@ -19,12 +19,11 @@ #include "MVKShaderModule.h" #include "MVKPipeline.h" #include "MVKFoundation.h" -#include using namespace std; -MVKMTLFunction::MVKMTLFunction(id mtlFunc, const SPIRVToMSLConversionResults scRslts, MTLSize tgSize) { +MVKMTLFunction::MVKMTLFunction(id mtlFunc, const SPIRVToMSLConversionResultInfo scRslts, MTLSize tgSize) { _mtlFunction = [mtlFunc retain]; // retained shaderConversionResults = scRslts; threadGroupSize = tgSize; @@ -74,7 +73,7 @@ MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkSpecializationInfo* pSpe @synchronized (_owner->getMTLDevice()) { @autoreleasepool { - NSString* mtlFuncName = @(_shaderConversionResults.entryPoint.mtlFunctionName.c_str()); + NSString* mtlFuncName = @(_shaderConversionResultInfo.entryPoint.mtlFunctionName.c_str()); MVKDevice* mvkDev = _owner->getDevice(); uint64_t startTime = mvkDev->getPerformanceTimestamp(); @@ -121,8 +120,8 @@ MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkSpecializationInfo* pSpe if ( !dbName ) { dbName = _owner-> getDebugName(); } setLabelIfNotNil(mtlFunc, dbName); - auto& wgSize = _shaderConversionResults.entryPoint.workgroupSize; - return MVKMTLFunction(mtlFunc, _shaderConversionResults, MTLSizeMake(getWorkgroupDimensionSize(wgSize.width, pSpecializationInfo), + auto& wgSize = _shaderConversionResultInfo.entryPoint.workgroupSize; + return MVKMTLFunction(mtlFunc, _shaderConversionResultInfo, MTLSizeMake(getWorkgroupDimensionSize(wgSize.width, pSpecializationInfo), getWorkgroupDimensionSize(wgSize.height, pSpecializationInfo), getWorkgroupDimensionSize(wgSize.depth, pSpecializationInfo))); } @@ -130,29 +129,55 @@ MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkSpecializationInfo* pSpe } void MVKShaderLibrary::setEntryPointName(string& funcName) { - _shaderConversionResults.entryPoint.mtlFunctionName = funcName; + _shaderConversionResultInfo.entryPoint.mtlFunctionName = funcName; } void MVKShaderLibrary::setWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { - auto& wgSize = _shaderConversionResults.entryPoint.workgroupSize; + auto& wgSize = _shaderConversionResultInfo.entryPoint.workgroupSize; wgSize.width.size = x; wgSize.height.size = y; wgSize.depth.size = z; } +// Sets the cached MSL source code, after first compressing it. +void MVKShaderLibrary::compressMSL(const string& msl) { + MVKDevice* mvkDev = _owner->getDevice(); + uint64_t startTime = mvkDev->getPerformanceTimestamp(); + _compressedMSL.compress(msl, mvkConfig().shaderSourceCompressionAlgorithm); + mvkDev->addActivityPerformance(mvkDev->_performanceStatistics.shaderCompilation.mslCompress, startTime); +} + +// Decompresses the cached MSL into the string. +void MVKShaderLibrary::decompressMSL(string& msl) { + MVKDevice* mvkDev = _owner->getDevice(); + uint64_t startTime = mvkDev->getPerformanceTimestamp(); + _compressedMSL.decompress(msl); + mvkDev->addActivityPerformance(mvkDev->_performanceStatistics.shaderCompilation.mslDecompress, startTime); +} + MVKShaderLibrary::MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, - const string& mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults) : _owner(owner) { + const SPIRVToMSLConversionResult& conversionResult) : _owner(owner) { + _shaderConversionResultInfo = conversionResult.resultInfo; + compressMSL(conversionResult.msl); + compileLibrary(conversionResult.msl); +} + +MVKShaderLibrary::MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, + const SPIRVToMSLConversionResultInfo& resultInfo, + const MVKCompressor compressedMSL) : _owner(owner) { + _shaderConversionResultInfo = resultInfo; + _compressedMSL = compressedMSL; + string msl; + decompressMSL(msl); + compileLibrary(msl); +} + +void MVKShaderLibrary::compileLibrary(const string& msl) { MVKShaderLibraryCompiler* slc = new MVKShaderLibraryCompiler(_owner); - - NSString* nsSrc = [[NSString alloc] initWithUTF8String: mslSourceCode.c_str()]; // temp retained - _mtlLibrary = slc->newMTLLibrary(nsSrc, shaderConversionResults); // retained - [nsSrc release]; // release temp string - + NSString* nsSrc = [[NSString alloc] initWithUTF8String: msl.c_str()]; // temp retained + _mtlLibrary = slc->newMTLLibrary(nsSrc, _shaderConversionResultInfo); // retained + [nsSrc release]; // release temp string slc->destroy(); - - _shaderConversionResults = shaderConversionResults; - _msl = mslSourceCode; } MVKShaderLibrary::MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, @@ -176,8 +201,8 @@ MVKShaderLibrary::MVKShaderLibrary(MVKVulkanAPIDeviceObject* owner, MVKShaderLibrary::MVKShaderLibrary(const MVKShaderLibrary& other) { _owner = other._owner; _mtlLibrary = [other._mtlLibrary retain]; - _shaderConversionResults = other._shaderConversionResults; - _msl = other._msl; + _shaderConversionResultInfo = other._shaderConversionResultInfo; + _compressedMSL = other._compressedMSL; } MVKShaderLibrary& MVKShaderLibrary::operator=(const MVKShaderLibrary& other) { @@ -186,8 +211,8 @@ MVKShaderLibrary& MVKShaderLibrary::operator=(const MVKShaderLibrary& other) { _mtlLibrary = [other._mtlLibrary retain]; } _owner = other._owner; - _shaderConversionResults = other._shaderConversionResults; - _msl = other._msl; + _shaderConversionResultInfo = other._shaderConversionResultInfo; + _compressedMSL = other._compressedMSL; return *this; } @@ -216,13 +241,14 @@ MVKShaderLibrary::~MVKShaderLibrary() { #pragma mark MVKShaderLibraryCache MVKShaderLibrary* MVKShaderLibraryCache::getShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, - MVKShaderModule* shaderModule, - bool* pWasAdded) { + MVKShaderModule* shaderModule, MVKPipeline* pipeline, + bool* pWasAdded, uint64_t startTime) { bool wasAdded = false; - MVKShaderLibrary* shLib = findShaderLibrary(pShaderConfig); - if ( !shLib ) { - if (shaderModule->convert(pShaderConfig)) { - shLib = addShaderLibrary(pShaderConfig, shaderModule->getMSL(), shaderModule->getConversionResults()); + MVKShaderLibrary* shLib = findShaderLibrary(pShaderConfig, startTime); + if ( !shLib && !pipeline->shouldFailOnPipelineCompileRequired() ) { + SPIRVToMSLConversionResult conversionResult; + if (shaderModule->convert(pShaderConfig, conversionResult)) { + shLib = addShaderLibrary(pShaderConfig, conversionResult); wasAdded = true; } } @@ -234,10 +260,13 @@ MVKShaderLibrary* MVKShaderLibraryCache::getShaderLibrary(SPIRVToMSLConversionCo // Finds and returns a shader library matching the shader config, or returns nullptr if it doesn't exist. // If a match is found, the shader config is aligned with the shader config of the matching library. -MVKShaderLibrary* MVKShaderLibraryCache::findShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig) { +MVKShaderLibrary* MVKShaderLibraryCache::findShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, + uint64_t startTime) { for (auto& slPair : _shaderLibraries) { if (slPair.first.matches(*pShaderConfig)) { pShaderConfig->alignWith(slPair.first); + MVKDevice* mvkDev = _owner->getDevice(); + mvkDev->addActivityPerformance(mvkDev->_performanceStatistics.shaderCompilation.shaderLibraryFromCache, startTime); return slPair.second; } } @@ -245,10 +274,18 @@ MVKShaderLibrary* MVKShaderLibraryCache::findShaderLibrary(SPIRVToMSLConversionC } // Adds and returns a new shader library configured from the specified conversion configuration. -MVKShaderLibrary* MVKShaderLibraryCache::addShaderLibrary(SPIRVToMSLConversionConfiguration* pShaderConfig, - const string& mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults) { - MVKShaderLibrary* shLib = new MVKShaderLibrary(_owner, mslSourceCode, shaderConversionResults); +MVKShaderLibrary* MVKShaderLibraryCache::addShaderLibrary(const SPIRVToMSLConversionConfiguration* pShaderConfig, + const SPIRVToMSLConversionResult& conversionResult) { + MVKShaderLibrary* shLib = new MVKShaderLibrary(_owner, conversionResult); + _shaderLibraries.emplace_back(*pShaderConfig, shLib); + return shLib; +} + +// Adds and returns a new shader library configured from contents read from a pipeline cache. +MVKShaderLibrary* MVKShaderLibraryCache::addShaderLibrary(const SPIRVToMSLConversionConfiguration* pShaderConfig, + const SPIRVToMSLConversionResultInfo& resultInfo, + const MVKCompressor compressedMSL) { + MVKShaderLibrary* shLib = new MVKShaderLibrary(_owner, resultInfo, compressedMSL); _shaderLibraries.emplace_back(*pShaderConfig, shLib); return shLib; } @@ -274,18 +311,17 @@ MVKShaderLibraryCache::~MVKShaderLibraryCache() { MVKMTLFunction MVKShaderModule::getMTLFunction(SPIRVToMSLConversionConfiguration* pShaderConfig, const VkSpecializationInfo* pSpecializationInfo, - MVKPipelineCache* pipelineCache) { - lock_guard lock(_accessLock); - + MVKPipeline* pipeline) { MVKShaderLibrary* mvkLib = _directMSLLibrary; if ( !mvkLib ) { uint64_t startTime = _device->getPerformanceTimestamp(); + MVKPipelineCache* pipelineCache = pipeline->getPipelineCache(); if (pipelineCache) { - mvkLib = pipelineCache->getShaderLibrary(pShaderConfig, this); + mvkLib = pipelineCache->getShaderLibrary(pShaderConfig, this, pipeline, startTime); } else { - mvkLib = _shaderLibraryCache.getShaderLibrary(pShaderConfig, this); + lock_guard lock(_accessLock); + mvkLib = _shaderLibraryCache.getShaderLibrary(pShaderConfig, this, pipeline, nullptr, startTime); } - _device->addActivityPerformance(_device->_performanceStatistics.shaderCompilation.shaderLibraryFromCache, startTime); } else { mvkLib->setEntryPointName(pShaderConfig->options.entryPointName); pShaderConfig->markAllInterfaceVarsAndResourcesUsed(); @@ -294,7 +330,8 @@ MVKMTLFunction MVKShaderModule::getMTLFunction(SPIRVToMSLConversionConfiguration return mvkLib ? mvkLib->getMTLFunction(pSpecializationInfo, this) : MVKMTLFunctionNull; } -bool MVKShaderModule::convert(SPIRVToMSLConversionConfiguration* pShaderConfig) { +bool MVKShaderModule::convert(SPIRVToMSLConversionConfiguration* pShaderConfig, + SPIRVToMSLConversionResult& conversionResult) { bool shouldLogCode = mvkConfig().debugMode; bool shouldLogEstimatedGLSL = shouldLogCode; @@ -302,27 +339,28 @@ bool MVKShaderModule::convert(SPIRVToMSLConversionConfiguration* pShaderConfig) // convert the GLSL code to SPIR-V and set it into the SPIR-V conveter. if ( !_spvConverter.hasSPIRV() && _glslConverter.hasGLSL() ) { + GLSLToSPIRVConversionResult glslConversionResult; uint64_t startTime = _device->getPerformanceTimestamp(); - bool wasConverted = _glslConverter.convert(getMVKGLSLConversionShaderStage(pShaderConfig), shouldLogCode, false); + bool wasConverted = _glslConverter.convert(getMVKGLSLConversionShaderStage(pShaderConfig), glslConversionResult, shouldLogCode, false); _device->addActivityPerformance(_device->_performanceStatistics.shaderCompilation.glslToSPRIV, startTime); if (wasConverted) { - if (shouldLogCode) { MVKLogInfo("%s", _glslConverter.getResultLog().c_str()); } - _spvConverter.setSPIRV(_glslConverter.getSPIRV()); + if (shouldLogCode) { MVKLogInfo("%s", glslConversionResult.resultLog.c_str()); } + _spvConverter.setSPIRV(glslConversionResult.spirv); } else { - reportError(VK_ERROR_INVALID_SHADER_NV, "Unable to convert GLSL to SPIR-V:\n%s", _glslConverter.getResultLog().c_str()); + reportError(VK_ERROR_INVALID_SHADER_NV, "Unable to convert GLSL to SPIR-V:\n%s", glslConversionResult.resultLog.c_str()); } shouldLogEstimatedGLSL = false; } uint64_t startTime = _device->getPerformanceTimestamp(); - bool wasConverted = _spvConverter.convert(*pShaderConfig, shouldLogCode, shouldLogCode, shouldLogEstimatedGLSL); + bool wasConverted = _spvConverter.convert(*pShaderConfig, conversionResult, shouldLogCode, shouldLogCode, shouldLogEstimatedGLSL); _device->addActivityPerformance(_device->_performanceStatistics.shaderCompilation.spirvToMSL, startTime); if (wasConverted) { - if (shouldLogCode) { MVKLogInfo("%s", _spvConverter.getResultLog().c_str()); } + if (shouldLogCode) { MVKLogInfo("%s", conversionResult.resultLog.c_str()); } } else { - reportError(VK_ERROR_INVALID_SHADER_NV, "Unable to convert SPIR-V to MSL:\n%s", _spvConverter.getResultLog().c_str()); + reportError(VK_ERROR_INVALID_SHADER_NV, "Unable to convert SPIR-V to MSL:\n%s", conversionResult.resultLog.c_str()); } return wasConverted; } @@ -345,7 +383,6 @@ MVKGLSLConversionShaderStage MVKShaderModule::getMVKGLSLConversionShaderStage(SP } void MVKShaderModule::setWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { - _spvConverter.setWorkgroupSize(x, y, z); if(_directMSLLibrary) { _directMSLLibrary->setWorkgroupSize(x, y, z); } } @@ -394,8 +431,9 @@ MVKShaderModule::MVKShaderModule(MVKDevice* device, codeHash = mvkHash(pMSLCode, mslCodeLen, codeHash); _device->addActivityPerformance(_device->_performanceStatistics.shaderCompilation.hashShaderCode, startTime); - _spvConverter.setMSL(pMSLCode, nullptr); - _directMSLLibrary = new MVKShaderLibrary(this, _spvConverter.getMSL().c_str(), _spvConverter.getConversionResults()); + SPIRVToMSLConversionResult conversionResult; + conversionResult.msl = pMSLCode; + _directMSLLibrary = new MVKShaderLibrary(this, conversionResult); break; } @@ -441,7 +479,7 @@ MVKShaderModule::~MVKShaderModule() { #pragma mark MVKShaderLibraryCompiler id MVKShaderLibraryCompiler::newMTLLibrary(NSString* mslSourceCode, - const SPIRVToMSLConversionResults& shaderConversionResults) { + const SPIRVToMSLConversionResultInfo& shaderConversionResults) { unique_lock lock(_completionLock); compile(lock, ^{ diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def index 598757c2..c619f4b4 100644 --- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def @@ -104,6 +104,7 @@ MVK_EXTENSION(EXT_inline_uniform_block, EXT_INLINE_UNIFORM_BLOCK, MVK_EXTENSION(EXT_memory_budget, EXT_MEMORY_BUDGET, DEVICE, 10.13, 11.0) MVK_EXTENSION(EXT_metal_objects, EXT_METAL_OBJECTS, DEVICE, 10.11, 8.0) MVK_EXTENSION(EXT_metal_surface, EXT_METAL_SURFACE, INSTANCE, 10.11, 8.0) +MVK_EXTENSION(EXT_pipeline_creation_cache_control, EXT_PIPELINE_CREATION_CACHE_CONTROL, DEVICE, 10.11, 8.0) MVK_EXTENSION(EXT_post_depth_coverage, EXT_POST_DEPTH_COVERAGE, DEVICE, 11.0, 11.0) MVK_EXTENSION(EXT_private_data, EXT_PRIVATE_DATA, DEVICE, 10.11, 8.0) MVK_EXTENSION(EXT_robustness2, EXT_ROBUSTNESS_2, DEVICE, 10.11, 8.0) diff --git a/MoltenVK/MoltenVK/Layers/MVKLayers.mm b/MoltenVK/MoltenVK/Layers/MVKLayers.mm index 3405a56e..040712b3 100644 --- a/MoltenVK/MoltenVK/Layers/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Layers/MVKLayers.mm @@ -19,7 +19,6 @@ #include "MVKLayers.h" #include "MVKEnvironment.h" #include "MVKFoundation.h" -#include "vk_mvk_moltenvk.h" #include using namespace std; diff --git a/MoltenVK/MoltenVK/Utility/MVKCodec.h b/MoltenVK/MoltenVK/Utility/MVKCodec.h index 136f4204..595a50fa 100644 --- a/MoltenVK/MoltenVK/Utility/MVKCodec.h +++ b/MoltenVK/MoltenVK/Utility/MVKCodec.h @@ -19,17 +19,18 @@ #pragma once -#include "MVKFoundation.h" +#include "MVKEnvironment.h" +#include #include +#pragma mark - #pragma mark Texture data codecs /** * This is the base class implemented by all codecs supported by MoltenVK. - * Objects of this class are used to decompress texture data for upload to a - * 3D texture. + * Objects of this class are used to decompress texture data for upload to a 3D texture. */ class MVKCodec { @@ -43,8 +44,89 @@ public: }; + +#pragma mark - +#pragma mark General data compressor + +/** + * Holds compressed data, along with information allowing it to be decompressed again. + * The template class C must support the basic data container methods data(), size() and resize(). + * + * THIS CLASS IS STREAMED OUT AS PART OF THE PIEPLINE CACHE. + * STURCTURAL CHANGES TO THIS CLASS MUST BE CAPTURED IN THE STREAMING LOGIC OF THE PIPELINE CACHE. + */ +template +class MVKCompressor { + +public: + + /** + * Compresses the content in the data container using the algorithm, and retains + * the compressed content. If an error occurs, or if the compressed data is actually + * larger (which can happen with some compression algorithms if the source is small), + * the uncompressed content is retained. Returns true if the content was successfully + * compressed, or returns false if the content was retained as uncompressed, + */ + bool compress(const C& uncompressed, MVKConfigCompressionAlgorithm algorithm) { + + _uncompressedSize = uncompressed.size(); + _compressed.resize(_uncompressedSize); + _algorithm = algorithm; + size_t compSize = mvkCompress((uint8_t*)uncompressed.data(), uncompressed.size(), + _compressed.data(), _compressed.size(), + _algorithm); + + bool wasCompressed = (compSize > 0); + if ( !wasCompressed ) { + _algorithm = MVK_CONFIG_COMPRESSION_ALGORITHM_NONE; + compSize = mvkCompress((uint8_t*)uncompressed.data(), uncompressed.size(), + _compressed.data(), _compressed.size(), + _algorithm); + } + + _compressed.resize(compSize); + _compressed.shrink_to_fit(); + + return wasCompressed; + } + + /** Decompress the retained compressed content into the data container. */ + void decompress(C& uncompressed) { + uncompressed.resize(_uncompressedSize); + mvkDecompress(_compressed.data(), _compressed.size(), + (uint8_t*)uncompressed.data(), uncompressed.size(), + _algorithm); + } + + std::vector _compressed; + size_t _uncompressedSize = 0; + MVKConfigCompressionAlgorithm _algorithm = MVK_CONFIG_COMPRESSION_ALGORITHM_NONE; +}; + + +#pragma mark - +#pragma mark Support functions + /** Returns an appropriate codec for the given format, or nullptr if the format is not supported. */ std::unique_ptr mvkCreateCodec(VkFormat format); /** Returns whether or not the given format can be decompressed. */ bool mvkCanDecodeFormat(VkFormat format); + +/** + * Compresses the source bytes into the destination bytes using a compression algorithm, + * and returns the number of bytes written to dstBytes. If an error occurs, or the compressed + * data is larger than dstSize, no data is copied to dstBytes, and zero is returned. + */ +size_t mvkCompress(const uint8_t* srcBytes, size_t srcSize, + uint8_t* dstBytes, size_t dstSize, + MVKConfigCompressionAlgorithm compAlgo); + +/** + * Decompresses the source bytes into the destination bytes using a compression algorithm, + * and returns the number of bytes written to dstBytes. If an error occurs, or the decompressed + * data is larger than dstSize, no data is copied to dstBytes, and zero is returned. + */ +size_t mvkDecompress(const uint8_t* srcBytes, size_t srcSize, + uint8_t* dstBytes, size_t dstSize, + MVKConfigCompressionAlgorithm compAlgo); diff --git a/MoltenVK/MoltenVK/Utility/MVKCodec.cpp b/MoltenVK/MoltenVK/Utility/MVKCodec.mm similarity index 53% rename from MoltenVK/MoltenVK/Utility/MVKCodec.cpp rename to MoltenVK/MoltenVK/Utility/MVKCodec.mm index 09669e49..7c07d6af 100644 --- a/MoltenVK/MoltenVK/Utility/MVKCodec.cpp +++ b/MoltenVK/MoltenVK/Utility/MVKCodec.mm @@ -18,10 +18,13 @@ #include "MVKCodec.h" +#include "MVKBaseObject.h" +#include "MVKFoundation.h" #include #include +using namespace std; using simd::float3; using simd::float4; @@ -62,8 +65,8 @@ public: for (uint32_t y = 0; y < extent.height; y += 4) { for (uint32_t x = 0; x < extent.width; x += 4) { VkExtent2D blockExtent; - blockExtent.width = std::min(extent.width - x, 4u); - blockExtent.height = std::min(extent.height - y, 4u); + blockExtent.width = min(extent.width - x, 4u); + blockExtent.height = min(extent.height - y, 4u); decompressDXTnBlock(pSrcRow + x * (blockByteCount / 4), pDestRow + x * 4, blockExtent, destLayout.rowPitch, _format); } @@ -90,7 +93,11 @@ private: VkFormat _format; }; -std::unique_ptr mvkCreateCodec(VkFormat format) { + +#pragma mark - +#pragma mark Support functions + +unique_ptr mvkCreateCodec(VkFormat format) { switch (format) { case VK_FORMAT_BC1_RGB_UNORM_BLOCK: case VK_FORMAT_BC1_RGB_SRGB_BLOCK: @@ -100,7 +107,7 @@ std::unique_ptr mvkCreateCodec(VkFormat format) { case VK_FORMAT_BC2_SRGB_BLOCK: case VK_FORMAT_BC3_UNORM_BLOCK: case VK_FORMAT_BC3_SRGB_BLOCK: - return std::unique_ptr(new MVKDXTnCodec(format)); + return unique_ptr(new MVKDXTnCodec(format)); default: return nullptr; @@ -123,3 +130,63 @@ bool mvkCanDecodeFormat(VkFormat format) { return false; } } + +static NSDataCompressionAlgorithm getSystemCompressionAlgo(MVKConfigCompressionAlgorithm compAlgo) { + switch (compAlgo) { + case MVK_CONFIG_COMPRESSION_ALGORITHM_NONE: return NSDataCompressionAlgorithmLZFSE; + case MVK_CONFIG_COMPRESSION_ALGORITHM_LZFSE: return NSDataCompressionAlgorithmLZFSE; + case MVK_CONFIG_COMPRESSION_ALGORITHM_LZ4: return NSDataCompressionAlgorithmLZ4; + case MVK_CONFIG_COMPRESSION_ALGORITHM_LZMA: return NSDataCompressionAlgorithmLZMA; + case MVK_CONFIG_COMPRESSION_ALGORITHM_ZLIB: return NSDataCompressionAlgorithmZlib; + default: return NSDataCompressionAlgorithmLZFSE; + } +} + +// Only copy into the dstBytes if it can fit, otherwise the data will be corrupted +static size_t mvkCompressDecompress(const uint8_t* srcBytes, size_t srcSize, + uint8_t* dstBytes, size_t dstSize, + MVKConfigCompressionAlgorithm compAlgo, + bool isCompressing) { + size_t dstByteCount = 0; + if (compAlgo != MVK_CONFIG_COMPRESSION_ALGORITHM_NONE) { + @autoreleasepool { + NSDataCompressionAlgorithm sysCompAlgo = getSystemCompressionAlgo(compAlgo); + NSData* srcData = [NSData dataWithBytesNoCopy: (void*)srcBytes length: srcSize freeWhenDone: NO]; + + NSError* err = nil; + NSData* dstData = (isCompressing + ? [srcData compressedDataUsingAlgorithm: sysCompAlgo error: &err] + : [srcData decompressedDataUsingAlgorithm: sysCompAlgo error: &err]); + if ( !err ) { + size_t dataLen = dstData.length; + if (dstSize >= dataLen) { + [dstData getBytes: dstBytes length: dstSize]; + dstByteCount = dataLen; + } + } else { + MVKBaseObject::reportError(nullptr, VK_ERROR_FORMAT_NOT_SUPPORTED, + "Could not %scompress data (Error code %li):\n%s", + (isCompressing ? "" : "de"), + (long)err.code, err.localizedDescription.UTF8String); + } + } + } else if (dstSize >= srcSize) { + mvkCopy(dstBytes, srcBytes, srcSize); + dstByteCount = srcSize; + } + return dstByteCount; +} + +size_t mvkCompress(const uint8_t* srcBytes, size_t srcSize, + uint8_t* dstBytes, size_t dstSize, + MVKConfigCompressionAlgorithm compAlgo) { + + return mvkCompressDecompress(srcBytes, srcSize, dstBytes, dstSize, compAlgo, true); +} + +size_t mvkDecompress(const uint8_t* srcBytes, size_t srcSize, + uint8_t* dstBytes, size_t dstSize, + MVKConfigCompressionAlgorithm compAlgo) { + + return mvkCompressDecompress(srcBytes, srcSize, dstBytes, dstSize, compAlgo, false); +} diff --git a/MoltenVK/MoltenVK/Utility/MVKEnvironment.cpp b/MoltenVK/MoltenVK/Utility/MVKEnvironment.cpp index 9861a359..1381c616 100644 --- a/MoltenVK/MoltenVK/Utility/MVKEnvironment.cpp +++ b/MoltenVK/MoltenVK/Utility/MVKEnvironment.cpp @@ -39,7 +39,7 @@ static void mvkInitConfigFromEnvVars() { MVK_SET_FROM_ENV_OR_BUILD_INT64 (evCfg.metalCompileTimeout, MVK_CONFIG_METAL_COMPILE_TIMEOUT); MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.performanceTracking, MVK_CONFIG_PERFORMANCE_TRACKING); MVK_SET_FROM_ENV_OR_BUILD_INT32 (evCfg.performanceLoggingFrameCount, MVK_CONFIG_PERFORMANCE_LOGGING_FRAME_COUNT); - MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.logActivityPerformanceInline, MVK_CONFIG_PERFORMANCE_LOGGING_INLINE); + MVK_SET_FROM_ENV_OR_BUILD_INT32 (evCfg.activityPerformanceLoggingStyle, MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE); MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.displayWatermark, MVK_CONFIG_DISPLAY_WATERMARK); MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.specializedQueueFamilies, MVK_CONFIG_SPECIALIZED_QUEUE_FAMILIES); MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.switchSystemGPU, MVK_CONFIG_SWITCH_SYSTEM_GPU); @@ -62,6 +62,7 @@ static void mvkInitConfigFromEnvVars() { MVK_SET_FROM_ENV_OR_BUILD_INT32 (evCfg.advertiseExtensions, MVK_CONFIG_ADVERTISE_EXTENSIONS); MVK_SET_FROM_ENV_OR_BUILD_BOOL (evCfg.resumeLostDevice, MVK_CONFIG_RESUME_LOST_DEVICE); MVK_SET_FROM_ENV_OR_BUILD_INT32 (evCfg.useMetalArgumentBuffers, MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS); + MVK_SET_FROM_ENV_OR_BUILD_INT32 (evCfg.shaderSourceCompressionAlgorithm, MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM); // Deprected legacy VkSemaphore MVK_ALLOW_METAL_FENCES and MVK_ALLOW_METAL_EVENTS config. // Legacy MVK_ALLOW_METAL_EVENTS is covered by MVK_CONFIG_VK_SEMAPHORE_SUPPORT_STYLE, @@ -75,6 +76,17 @@ static void mvkInitConfigFromEnvVars() { evCfg.semaphoreUseMTLEvent = (MVKVkSemaphoreSupportStyle)false; // Disabled. Also semaphoreSupportStyle MVK_CONFIG_VK_SEMAPHORE_SUPPORT_STYLE_SINGLE_QUEUE. } + // Deprecated legacy env var MVK_CONFIG_PERFORMANCE_LOGGING_INLINE config. If legacy + // MVK_CONFIG_PERFORMANCE_LOGGING_INLINE env var was used, and activityPerformanceLoggingStyle + // was not already set by MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE, set + // activityPerformanceLoggingStyle to MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_IMMEDIATE. + bool logPerfInline; + MVK_SET_FROM_ENV_OR_BUILD_BOOL(logPerfInline, MVK_CONFIG_PERFORMANCE_LOGGING_INLINE); + if (logPerfInline && evCfg.activityPerformanceLoggingStyle == MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_FRAME_COUNT) { + evCfg.activityPerformanceLoggingStyle = MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_IMMEDIATE; + } + + mvkSetConfig(evCfg); } diff --git a/MoltenVK/MoltenVK/Utility/MVKEnvironment.h b/MoltenVK/MoltenVK/Utility/MVKEnvironment.h index 16833d08..cec1740e 100644 --- a/MoltenVK/MoltenVK/Utility/MVKEnvironment.h +++ b/MoltenVK/MoltenVK/Utility/MVKEnvironment.h @@ -151,8 +151,11 @@ void mvkSetConfig(const MVKConfiguration& mvkConfig); # define MVK_CONFIG_PERFORMANCE_LOGGING_FRAME_COUNT 0 #endif -/** Log activity performance every time an activity occurs. Disabled by default. */ -# ifndef MVK_CONFIG_PERFORMANCE_LOGGING_INLINE +/** Activity performance logging style. Default is to log after a configured number of frames. */ +# ifndef MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE +# define MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE_FRAME_COUNT +# endif +# ifndef MVK_CONFIG_PERFORMANCE_LOGGING_INLINE // Deprecated # define MVK_CONFIG_PERFORMANCE_LOGGING_INLINE 0 # endif @@ -286,3 +289,8 @@ void mvkSetConfig(const MVKConfiguration& mvkConfig); #ifndef MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS # define MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS_NEVER #endif + +/** Compress MSL shader source code in a pipeline cache. Defaults to no compression. */ +#ifndef MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM +# define MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM MVK_CONFIG_COMPRESSION_ALGORITHM_NONE +#endif diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLConversion.mm b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLConversion.mm index 128ffb50..8ff31cf4 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLConversion.mm +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLConversion.mm @@ -32,27 +32,27 @@ MVK_PUBLIC_SYMBOL bool mvkConvertGLSLToSPIRV(const char* glslSource, char** pResultLog, bool shouldLogGLSL, bool shouldLogSPIRV) { - GLSLToSPIRVConverter glslConverter; + + GLSLToSPIRVConversionResult conversionResult; + GLSLToSPIRVConverter glslConverter; glslConverter.setGLSL(glslSource); - bool wasConverted = glslConverter.convert(shaderStage, shouldLogGLSL, shouldLogSPIRV); + bool wasConverted = glslConverter.convert(shaderStage, conversionResult, shouldLogGLSL, shouldLogSPIRV); size_t spvLen = 0; if (pSPIRVCode) { uint32_t* spvCode = NULL; if (wasConverted) { - auto spv = glslConverter.getSPIRV(); - spvLen = spv.size() * sizeof(uint32_t); + spvLen = conversionResult.spirv.size() * sizeof(uint32_t); spvCode = (uint32_t*)malloc(spvLen); - memcpy(spvCode, spv.data(), spvLen); + memcpy(spvCode, conversionResult.spirv.data(), spvLen); } *pSPIRVCode = spvCode; } if (pSPIRVLength) { *pSPIRVLength = spvLen; } if (pResultLog) { - auto log = glslConverter.getResultLog(); - *pResultLog = (char*)malloc(log.size() + 1); - strcpy(*pResultLog, log.data()); + *pResultLog = (char*)malloc(conversionResult.resultLog.size() + 1); + strcpy(*pResultLog, conversionResult.resultLog.data()); } return wasConverted; diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.cpp b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.cpp index bb85a8e4..136aa5c1 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.cpp +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.cpp @@ -53,13 +53,12 @@ MVK_PUBLIC_SYMBOL void GLSLToSPIRVConverter::setGLSLs(const std::vectorsetAutoMapBindings(true); if (glslShaders.back()->parse(&glslCompilerResources, 100, false, messages)) { if (shouldLogGLSL) { - logMsg(glslShaders.back()->getInfoLog()); - logMsg(glslShaders.back()->getInfoDebugLog()); + logMsg(conversionResult.resultLog, glslShaders.back()->getInfoLog()); + logMsg(conversionResult.resultLog, glslShaders.back()->getInfoDebugLog()); } } else { - logError(glslShaders.back()->getInfoLog()); - logError(glslShaders.back()->getInfoDebugLog()); - return logError("Error compiling GLSL when converting GLSL to SPIR-V."); + logError(conversionResult.resultLog, glslShaders.back()->getInfoLog()); + logError(conversionResult.resultLog, glslShaders.back()->getInfoDebugLog()); + return logError(conversionResult.resultLog, "Error compiling GLSL when converting GLSL to SPIR-V."); } // Add a shader to the program. Each shader added will be linked together. glslProgram.addShader(glslShaders.back().get()); @@ -93,61 +92,57 @@ MVK_PUBLIC_SYMBOL bool GLSLToSPIRVConverter::convert(MVKGLSLConversionShaderStag // Create and link a shader program if ( !glslProgram.link(messages) ) { - logError(glslProgram.getInfoLog()); - logError(glslProgram.getInfoDebugLog()); - return logError("Error creating GLSL program when converting GLSL to SPIR-V."); + logError(conversionResult.resultLog, glslProgram.getInfoLog()); + logError(conversionResult.resultLog, glslProgram.getInfoDebugLog()); + return logError(conversionResult.resultLog, "Error creating GLSL program when converting GLSL to SPIR-V."); } // Output the SPIR-V code from the shader program - glslang::GlslangToSpv(*glslProgram.getIntermediate(stage), _spirv); + glslang::GlslangToSpv(*glslProgram.getIntermediate(stage), conversionResult.spirv); - if (shouldLogSPIRV) { logSPIRV("Converted"); } + if (shouldLogSPIRV) { logSPIRV(conversionResult, "Converted"); } - return _wasConverted; + return wasConverted; } /** Appends the message text to the result log. */ -void GLSLToSPIRVConverter::logMsg(const char* logMsg) { +void GLSLToSPIRVConverter::logMsg(string& log, const char* logMsg) { string trimMsg = trim(logMsg); if ( !trimMsg.empty() ) { - _resultLog += trimMsg; - _resultLog += "\n\n"; + log += trimMsg; + log += "\n\n"; } } /** Appends the error text to the result log, sets the wasConverted property to false, and returns it. */ -bool GLSLToSPIRVConverter::logError(const char* errMsg) { - logMsg(errMsg); - _wasConverted = false; - return _wasConverted; +bool GLSLToSPIRVConverter::logError(string& log, const char* errMsg) { + logMsg(log, errMsg); + return false; } /** Appends the SPIR-V to the result log, indicating whether it is being converted or was converted. */ -void GLSLToSPIRVConverter::logSPIRV(const char* opDesc) { +void GLSLToSPIRVConverter::logSPIRV(GLSLToSPIRVConversionResult& conversionResult, const char* opDesc) { - string spvLog; - mvk::logSPIRV(_spirv, spvLog); - - _resultLog += opDesc; - _resultLog += " SPIR-V:\n"; - _resultLog += spvLog; - _resultLog += "\nEnd SPIR-V\n\n"; + conversionResult.resultLog += opDesc; + conversionResult.resultLog += " SPIR-V:\n"; + mvk::logSPIRV(conversionResult.spirv, conversionResult.resultLog); + conversionResult.resultLog += "\nEnd SPIR-V\n\n"; } /** Validates that the SPIR-V code will disassemble during logging. */ -bool GLSLToSPIRVConverter::validateSPIRV() { - if (_spirv.size() < 5) { return false; } - if (_spirv[0] != spv::MagicNumber) { return false; } - if (_spirv[4] != 0) { return false; } +bool GLSLToSPIRVConverter::validateSPIRV(vector spirv) { + if (spirv.size() < 5) { return false; } + if (spirv[0] != spv::MagicNumber) { return false; } + if (spirv[4] != 0) { return false; } return true; } /** Appends the GLSL to the result log, indicating whether it is being converted or was converted. */ -void GLSLToSPIRVConverter::logGLSL(const char* opDesc) { - _resultLog += opDesc; - _resultLog += " GLSL:\n"; - for (const auto& glsl : _glsls) { _resultLog += glsl + "\n"; } - _resultLog += "End GLSL\n\n"; +void GLSLToSPIRVConverter::logGLSL(string& log, const char* opDesc) { + log += opDesc; + log += " GLSL:\n"; + for (const auto& glsl : _glsls) { log += glsl + "\n"; } + log += "End GLSL\n\n"; } diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.h b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.h index c83443da..53ba9b2c 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.h +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.h @@ -27,6 +27,17 @@ namespace mvk { + +#pragma mark - +#pragma mark SPIRVToMSLConversionResult + + /** The results of a GLSL to SPIRV conversion. */ + typedef struct GLSLToSPIRVConversionResult { + std::vector spirv; + std::string resultLog; + } GLSLToSPIRVConversionResult; + + #pragma mark - #pragma mark GLSLToSPIRVConverter @@ -50,8 +61,7 @@ namespace mvk { /** * Sets the GLSL source code that is to be converted to the specified strings. * - * A separate shader will be compiled for each source and linked together into a single - * program. + * A separate shader will be compiled for each source and linked together into a single program. */ void setGLSLs(const std::vector& glslSrcs); @@ -67,36 +77,20 @@ namespace mvk { * The boolean flags indicate whether the original GLSL code and resulting SPIR-V code should * be logged to the result log of this converter. This can be useful during shader debugging. */ - bool convert(MVKGLSLConversionShaderStage shaderStage, bool shouldLogGLSL, bool shouldLogSPIRV); - - /** - * Returns whether the most recent conversion was successful. - * - * The initial value of this property is NO. It is set to YES upon successful conversion. - */ - bool wasConverted() { return _wasConverted; } - - /** Returns the SPIRV code most recently converted by the convert() function. */ - const std::vector& getSPIRV() { return _spirv; } - - /** - * Returns a human-readable log of the most recent conversion activity. - * This may be empty if the conversion was successful. - */ - const std::string& getResultLog() { return _resultLog; } + bool convert(MVKGLSLConversionShaderStage shaderStage, + GLSLToSPIRVConversionResult& conversionResult, + bool shouldLogGLSL, + bool shouldLogSPIRV); protected: - void logMsg(const char* logMsg); - bool logError(const char* errMsg); - void logGLSL(const char* opDesc); - void logSPIRV(const char* opDesc); - bool validateSPIRV(); + void logMsg(std::string& log, const char* logMsg); + bool logError(std::string& log, const char* errMsg); + void logGLSL(std::string& log, const char* opDesc); + void logSPIRV(GLSLToSPIRVConversionResult& conversionResult, const char* opDesc); + bool validateSPIRV(std::vector spirv); void initGLSLCompilerResources(); std::vector _glsls; - std::vector _spirv; - std::string _resultLog; - bool _wasConverted = false; }; } diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVConversion.mm b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVConversion.mm index 3d44b4bd..257018bf 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVConversion.mm +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVConversion.mm @@ -32,20 +32,19 @@ MVK_PUBLIC_SYMBOL bool mvkConvertSPIRVToMSL(uint32_t* spvCode, bool shouldLogSPIRV, bool shouldLogMSL) { SPIRVToMSLConversionConfiguration spvCtx; + SPIRVToMSLConversionResult conversionResult; SPIRVToMSLConverter spvConverter; spvConverter.setSPIRV(spvCode, spvLength); - bool wasConverted = spvConverter.convert(spvCtx, shouldLogSPIRV, shouldLogMSL); + bool wasConverted = spvConverter.convert(spvCtx, conversionResult, shouldLogSPIRV, shouldLogMSL); if (pMSL) { - auto& msl = spvConverter.getMSL(); - *pMSL = (char*)malloc(msl.size() + 1); - strcpy(*pMSL, msl.data()); + *pMSL = (char*)malloc(conversionResult.msl.size() + 1); + strcpy(*pMSL, conversionResult.msl.data()); } if (pResultLog) { - auto log = spvConverter.getResultLog(); - *pResultLog = (char*)malloc(log.size() + 1); - strcpy(*pResultLog, log.data()); + *pResultLog = (char*)malloc(conversionResult.resultLog.size() + 1); + strcpy(*pResultLog, conversionResult.resultLog.data()); } return wasConverted; diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.cpp b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.cpp index 3a7f53f4..f1672e2b 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.cpp +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.cpp @@ -261,6 +261,7 @@ MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const uint32_t* spirvCode, } MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConversionConfiguration& shaderConfig, + SPIRVToMSLConversionResult& conversionResult, bool shouldLogSPIRV, bool shouldLogMSL, bool shouldLogGLSL) { @@ -270,14 +271,10 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConversionConfigur // spvFile.write((char*)_spirv.data(), _spirv.size() << 2); // spvFile.close(); - _wasConverted = true; - _resultLog.clear(); - _msl.clear(); - _shaderConversionResults.reset(); - - if (shouldLogSPIRV) { logSPIRV("Converting"); } + if (shouldLogSPIRV) { logSPIRV(conversionResult.resultLog, "Converting"); } CompilerMSL* pMSLCompiler = nullptr; + bool wasConverted = true; #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS try { @@ -341,42 +338,42 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConversionConfigur } } } - _msl = pMSLCompiler->compile(); + conversionResult.msl = pMSLCompiler->compile(); - if (shouldLogMSL) { logSource(_msl, "MSL", "Converted"); } + if (shouldLogMSL) { logSource(conversionResult.resultLog, conversionResult.msl, "MSL", "Converted"); } #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS } catch (CompilerError& ex) { string errMsg("MSL conversion error: "); errMsg += ex.what(); - logError(errMsg.data()); + logError(conversionResult.resultLog, errMsg.data()); if (shouldLogMSL && pMSLCompiler) { - _msl = pMSLCompiler->get_partial_source(); - logSource(_msl, "MSL", "Partially converted"); + auto partialMSL = pMSLCompiler->get_partial_source(); + logSource(conversionResult.resultLog, partialMSL, "MSL", "Partially converted"); } } #endif // Populate the shader conversion results with info from the compilation run, // and mark which vertex attributes and resource bindings are used by the shader - populateEntryPoint(pMSLCompiler, shaderConfig.options); - _shaderConversionResults.isRasterizationDisabled = pMSLCompiler && pMSLCompiler->get_is_rasterization_disabled(); - _shaderConversionResults.isPositionInvariant = pMSLCompiler && pMSLCompiler->is_position_invariant(); - _shaderConversionResults.needsSwizzleBuffer = pMSLCompiler && pMSLCompiler->needs_swizzle_buffer(); - _shaderConversionResults.needsOutputBuffer = pMSLCompiler && pMSLCompiler->needs_output_buffer(); - _shaderConversionResults.needsPatchOutputBuffer = pMSLCompiler && pMSLCompiler->needs_patch_output_buffer(); - _shaderConversionResults.needsBufferSizeBuffer = pMSLCompiler && pMSLCompiler->needs_buffer_size_buffer(); - _shaderConversionResults.needsInputThreadgroupMem = pMSLCompiler && pMSLCompiler->needs_input_threadgroup_mem(); - _shaderConversionResults.needsDispatchBaseBuffer = pMSLCompiler && pMSLCompiler->needs_dispatch_base_buffer(); - _shaderConversionResults.needsViewRangeBuffer = pMSLCompiler && pMSLCompiler->needs_view_mask_buffer(); + populateEntryPoint(pMSLCompiler, shaderConfig.options, conversionResult.resultInfo.entryPoint); + conversionResult.resultInfo.isRasterizationDisabled = pMSLCompiler && pMSLCompiler->get_is_rasterization_disabled(); + conversionResult.resultInfo.isPositionInvariant = pMSLCompiler && pMSLCompiler->is_position_invariant(); + conversionResult.resultInfo.needsSwizzleBuffer = pMSLCompiler && pMSLCompiler->needs_swizzle_buffer(); + conversionResult.resultInfo.needsOutputBuffer = pMSLCompiler && pMSLCompiler->needs_output_buffer(); + conversionResult.resultInfo.needsPatchOutputBuffer = pMSLCompiler && pMSLCompiler->needs_patch_output_buffer(); + conversionResult.resultInfo.needsBufferSizeBuffer = pMSLCompiler && pMSLCompiler->needs_buffer_size_buffer(); + conversionResult.resultInfo.needsInputThreadgroupMem = pMSLCompiler && pMSLCompiler->needs_input_threadgroup_mem(); + conversionResult.resultInfo.needsDispatchBaseBuffer = pMSLCompiler && pMSLCompiler->needs_dispatch_base_buffer(); + conversionResult.resultInfo.needsViewRangeBuffer = pMSLCompiler && pMSLCompiler->needs_view_mask_buffer(); // When using Metal argument buffers, if the shader is provided with dynamic buffer offsets, // then it needs a buffer to hold these dynamic offsets. - _shaderConversionResults.needsDynamicOffsetBuffer = false; + conversionResult.resultInfo.needsDynamicOffsetBuffer = false; if (shaderConfig.options.mslOptions.argument_buffers) { for (auto& db : shaderConfig.dynamicBufferDescriptors) { if (db.stage == shaderConfig.options.entryPointStage) { - _shaderConversionResults.needsDynamicOffsetBuffer = true; + conversionResult.resultInfo.needsDynamicOffsetBuffer = true; } } } @@ -418,68 +415,64 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConversionConfigur options.separate_shader_objects = true; pGLSLCompiler->set_common_options(options); string glsl = pGLSLCompiler->compile(); - logSource(glsl, "GLSL", "Estimated original"); + logSource(conversionResult.resultLog, glsl, "GLSL", "Estimated original"); #ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS } catch (CompilerError& ex) { string errMsg("Original GLSL extraction error: "); errMsg += ex.what(); - logMsg(errMsg.data()); + logMsg(conversionResult.resultLog, errMsg.data()); if (pGLSLCompiler) { string glsl = pGLSLCompiler->get_partial_source(); - logSource(glsl, "GLSL", "Partially converted"); + logSource(conversionResult.resultLog, glsl, "GLSL", "Partially converted"); } } #endif delete pGLSLCompiler; } - return _wasConverted; + return wasConverted; } // Appends the message text to the result log. -void SPIRVToMSLConverter::logMsg(const char* logMsg) { +void SPIRVToMSLConverter::logMsg(string& log, const char* logMsg) { string trimMsg = trim(logMsg); if ( !trimMsg.empty() ) { - _resultLog += trimMsg; - _resultLog += "\n\n"; + log += trimMsg; + log += "\n\n"; } } -// Appends the error text to the result log, sets the wasConverted property to false, and returns it. -bool SPIRVToMSLConverter::logError(const char* errMsg) { - logMsg(errMsg); - _wasConverted = false; - return _wasConverted; +// Appends the error text to the result log, and returns false to indicate an error. +bool SPIRVToMSLConverter::logError(string& log, const char* errMsg) { + logMsg(log, errMsg); + return false; } // Appends the SPIR-V to the result log, indicating whether it is being converted or was converted. -void SPIRVToMSLConverter::logSPIRV(const char* opDesc) { +void SPIRVToMSLConverter::logSPIRV(string& log, const char* opDesc) { - string spvLog; - mvk::logSPIRV(_spirv, spvLog); - - _resultLog += opDesc; - _resultLog += " SPIR-V:\n"; - _resultLog += spvLog; - _resultLog += "\nEnd SPIR-V\n\n"; + log += opDesc; + log += " SPIR-V:\n"; + mvk::logSPIRV(_spirv, log); + log += "\nEnd SPIR-V\n\n"; // Uncomment one or both of the following lines to get additional debugging and tracability capabilities. // The SPIR-V can be written in binary form to a file, and/or logged in human readable form to the console. // These can be helpful if errors occur during conversion of SPIR-V to MSL. -// writeSPIRVToFile("spvout.spv"); -// printf("\n%s\n", getResultLog().c_str()); +// writeSPIRVToFile("spvout.spv", log); +// printf("\n%s\n", log.c_str()); } // Writes the SPIR-V code to a file. This can be useful for debugging // when the SPRIR-V did not originally come from a known file -void SPIRVToMSLConverter::writeSPIRVToFile(string spvFilepath) { +void SPIRVToMSLConverter::writeSPIRVToFile(string spvFilepath, string& log) { vector fileContents; spirvToBytes(_spirv, fileContents); string errMsg; if (writeFile(spvFilepath, fileContents, errMsg)) { - _resultLog += "Saved SPIR-V to file: " + absolutePath(spvFilepath) + "\n\n"; + log += "Saved SPIR-V to file: " + absolutePath(spvFilepath) + "\n\n"; } else { - _resultLog += "Could not write SPIR-V file. " + errMsg + "\n\n"; + log += "Could not write SPIR-V file. " + errMsg + "\n\n"; } } @@ -492,15 +485,15 @@ bool SPIRVToMSLConverter::validateSPIRV() { } // Appends the source to the result log, prepending with the operation. -void SPIRVToMSLConverter::logSource(string& src, const char* srcLang, const char* opDesc) { - _resultLog += opDesc; - _resultLog += " "; - _resultLog += srcLang; - _resultLog += ":\n"; - _resultLog += src; - _resultLog += "\nEnd "; - _resultLog += srcLang; - _resultLog += "\n\n"; +void SPIRVToMSLConverter::logSource(string& log, string& src, const char* srcLang, const char* opDesc) { + log += opDesc; + log += " "; + log += srcLang; + log += ":\n"; + log += src; + log += "\nEnd "; + log += srcLang; + log += "\n\n"; } void SPIRVToMSLConverter::populateWorkgroupDimension(SPIRVWorkgroupSizeDimension& wgDim, @@ -513,7 +506,8 @@ void SPIRVToMSLConverter::populateWorkgroupDimension(SPIRVWorkgroupSizeDimension // Populates the entry point with info extracted from the SPRI-V compiler. void SPIRVToMSLConverter::populateEntryPoint(Compiler* pCompiler, - SPIRVToMSLConversionOptions& options) { + SPIRVToMSLConversionOptions& options, + SPIRVEntryPoint& entryPoint) { if ( !pCompiler ) { return; } @@ -528,14 +522,13 @@ void SPIRVToMSLConverter::populateEntryPoint(Compiler* pCompiler, } } - auto& ep = _shaderConversionResults.entryPoint; - ep.mtlFunctionName = spvEP.name; - ep.supportsFastMath = !spvEP.flags.get(ExecutionModeSignedZeroInfNanPreserve); + entryPoint.mtlFunctionName = spvEP.name; + entryPoint.supportsFastMath = !spvEP.flags.get(ExecutionModeSignedZeroInfNanPreserve); SpecializationConstant widthSC, heightSC, depthSC; pCompiler->get_work_group_size_specialization_constants(widthSC, heightSC, depthSC); - auto& wgSize = ep.workgroupSize; + auto& wgSize = entryPoint.workgroupSize; populateWorkgroupDimension(wgSize.width, spvEP.workgroup_size.x, widthSC); populateWorkgroupDimension(wgSize.height, spvEP.workgroup_size.y, heightSC); populateWorkgroupDimension(wgSize.depth, spvEP.workgroup_size.z, depthSC); diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h index 1acc2bec..765e110f 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h +++ b/MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h @@ -192,7 +192,7 @@ namespace mvk { #pragma mark - -#pragma mark SPIRVToMSLConversionResults +#pragma mark SPIRVToMSLConversionResult /** * Describes one dimension of the workgroup size of a SPIR-V entry point, including whether @@ -227,12 +227,12 @@ namespace mvk { } SPIRVEntryPoint; /** - * Contains the results of the shader conversion that can be used to populate a pipeline. + * Contains information about a shader conversion that can be used to populate a pipeline. * * THIS STRUCT IS STREAMED OUT AS PART OF THE PIEPLINE CACHE. * CHANGES TO THIS STRUCT SHOULD BE CAPTURED IN THE STREAMING LOGIC OF THE PIPELINE CACHE. */ - typedef struct SPIRVToMSLConversionResults { + typedef struct SPIRVToMSLConversionResultInfo { SPIRVEntryPoint entryPoint; bool isRasterizationDisabled = false; bool isPositionInvariant = false; @@ -245,9 +245,14 @@ namespace mvk { bool needsDispatchBaseBuffer = false; bool needsViewRangeBuffer = false; - void reset() { *this = SPIRVToMSLConversionResults(); } + } SPIRVToMSLConversionResultInfo; - } SPIRVToMSLConversionResults; + /** The results of a SPIRV to MSL conversion. */ + typedef struct SPIRVToMSLConversionResult { + SPIRVToMSLConversionResultInfo resultInfo = {}; + std::string msl; + std::string resultLog; + } SPIRVToMSLConversionResult; #pragma mark - @@ -274,68 +279,29 @@ namespace mvk { bool hasSPIRV() { return !_spirv.empty(); } /** - * Converts SPIR-V code, set using setSPIRV() to MSL code, which can be retrieved using getMSL(). + * Converts SPIR-V code, set using setSPIRV() to MSL code. * * The boolean flags indicate whether the original SPIR-V code, the resulting MSL code, * and optionally, the original GLSL (as converted from the SPIR_V), should be logged * to the result log of this converter. This can be useful during shader debugging. */ bool convert(SPIRVToMSLConversionConfiguration& shaderConfig, - bool shouldLogSPIRV = false, - bool shouldLogMSL = false, - bool shouldLogGLSL = false); - - /** - * Returns whether the most recent conversion was successful. - * - * The initial value of this property is NO. It is set to YES upon successful conversion. - */ - bool wasConverted() { return _wasConverted; } - - /** - * Returns the Metal Shading Language source code most recently converted - * by the convert() function, or set directly using the setMSL() function. - */ - const std::string& getMSL() { return _msl; } - - /** Returns information about the shader conversion. */ - const SPIRVToMSLConversionResults& getConversionResults() { return _shaderConversionResults; } - - /** Sets the number of threads in a single compute kernel workgroup, per dimension. */ - void setWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { - auto& wgSize = _shaderConversionResults.entryPoint.workgroupSize; - wgSize.width.size = x; - wgSize.height.size = y; - wgSize.depth.size = z; - } - - /** - * Returns a human-readable log of the most recent conversion activity. - * This may be empty if the conversion was successful. - */ - const std::string& getResultLog() { return _resultLog; } - - /** Sets MSL source code. This can be used when MSL is supplied directly. */ - void setMSL(const std::string& msl, const SPIRVToMSLConversionResults* pShaderConversionResults) { - _msl = msl; - if (pShaderConversionResults) { _shaderConversionResults = *pShaderConversionResults; } - } + SPIRVToMSLConversionResult& conversionResult, + bool shouldLogSPIRV = false, + bool shouldLogMSL = false, + bool shouldLogGLSL = false); protected: - void logMsg(const char* logMsg); - bool logError(const char* errMsg); - void logSPIRV(const char* opDesc); + void logMsg(std::string& log, const char* logMsg); + bool logError(std::string& log, const char* errMsg); + void logSPIRV(std::string& log, const char* opDesc); + void logSource(std::string& log, std::string& src, const char* srcLang, const char* opDesc); bool validateSPIRV(); - void writeSPIRVToFile(std::string spvFilepath); - void logSource(std::string& src, const char* srcLang, const char* opDesc); + void writeSPIRVToFile(std::string spvFilepath, std::string& log); void populateWorkgroupDimension(SPIRVWorkgroupSizeDimension& wgDim, uint32_t size, SPIRV_CROSS_NAMESPACE::SpecializationConstant& spvSpecConst); - void populateEntryPoint(SPIRV_CROSS_NAMESPACE::Compiler* pCompiler, SPIRVToMSLConversionOptions& options); + void populateEntryPoint(SPIRV_CROSS_NAMESPACE::Compiler* pCompiler, SPIRVToMSLConversionOptions& options, SPIRVEntryPoint& entryPoint); std::vector _spirv; - std::string _msl; - std::string _resultLog; - SPIRVToMSLConversionResults _shaderConversionResults; - bool _wasConverted = false; }; } diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverterTool/MoltenVKShaderConverterTool.cpp b/MoltenVKShaderConverter/MoltenVKShaderConverterTool/MoltenVKShaderConverterTool.cpp index e38033b9..9a911e3d 100644 --- a/MoltenVKShaderConverter/MoltenVKShaderConverterTool/MoltenVKShaderConverterTool.cpp +++ b/MoltenVKShaderConverter/MoltenVKShaderConverterTool/MoltenVKShaderConverterTool.cpp @@ -137,24 +137,23 @@ bool MoltenVKShaderConverterTool::convertGLSL(string& glslInFile, } // Convert GLSL to SPIR-V + GLSLToSPIRVConversionResult conversionResult; GLSLToSPIRVConverter glslConverter; glslConverter.setGLSL(glslCode); uint64_t startTime = _glslConversionPerformance.getTimestamp(); - bool wasConverted = glslConverter.convert(shaderStage, _shouldLogConversions, _shouldLogConversions); + bool wasConverted = glslConverter.convert(shaderStage, conversionResult, _shouldLogConversions, _shouldLogConversions); _glslConversionPerformance.accumulate(startTime); if (wasConverted) { - if (_shouldLogConversions) { log(glslConverter.getResultLog().data()); } + if (_shouldLogConversions) { log(conversionResult.resultLog.data()); } } else { string logMsg = "Could not convert GLSL in file: " + absolutePath(path); log(logMsg.data()); - log(glslConverter.getResultLog().data()); + log(conversionResult.resultLog.data()); return false; } - const vector& spv = glslConverter.getSPIRV(); - // Write the SPIR-V code to a file. // If no file has been supplied, create one from the GLSL file name. if (_shouldWriteSPIRV) { @@ -162,9 +161,9 @@ bool MoltenVKShaderConverterTool::convertGLSL(string& glslInFile, if (path.empty()) { path = pathWithExtension(glslInFile, _shouldOutputAsHeaders ? "h" : "spv", _shouldIncludeOrigPathExtn, _origPathExtnSep); } if (_shouldOutputAsHeaders) { - spirvToHeaderBytes(spv, fileContents, fileName(path, false)); + spirvToHeaderBytes(conversionResult.spirv, fileContents, fileName(path, false)); } else { - spirvToBytes(spv, fileContents); + spirvToBytes(conversionResult.spirv, fileContents); } if (writeFile(path, fileContents, errMsg)) { @@ -177,7 +176,7 @@ bool MoltenVKShaderConverterTool::convertGLSL(string& glslInFile, } } - return convertSPIRV(spv, glslInFile, mslOutFile, false); + return convertSPIRV(conversionResult.spirv, glslInFile, mslOutFile, false); } // Read SPIR-V code from a SPIR-V file, convert to MSL, and write the MSL code to files. @@ -227,25 +226,25 @@ bool MoltenVKShaderConverterTool::convertSPIRV(const vector& spv, spvConverter.setSPIRV(spv); uint64_t startTime = _spvConversionPerformance.getTimestamp(); - bool wasConverted = spvConverter.convert(mslContext, shouldLogSPV, _shouldLogConversions, (_shouldLogConversions && shouldLogSPV)); + SPIRVToMSLConversionResult conversionResult; + bool wasConverted = spvConverter.convert(mslContext, conversionResult, shouldLogSPV, _shouldLogConversions, (_shouldLogConversions && shouldLogSPV)); _spvConversionPerformance.accumulate(startTime); if (wasConverted) { - if (_shouldLogConversions) { log(spvConverter.getResultLog().data()); } + if (_shouldLogConversions) { log(conversionResult.resultLog.data()); } } else { string errMsg = "Could not convert SPIR-V in file: " + absolutePath(inFile); log(errMsg.data()); - log(spvConverter.getResultLog().data()); + log(conversionResult.resultLog.data()); return false; } // Write the MSL to file string path = mslOutFile; if (mslOutFile.empty()) { path = pathWithExtension(inFile, "metal", _shouldIncludeOrigPathExtn, _origPathExtnSep); } - const string& msl = spvConverter.getMSL(); string compileErrMsg; - bool wasCompiled = compile(msl, compileErrMsg, _mslVersionMajor, _mslVersionMinor, _mslVersionPatch); + bool wasCompiled = compile(conversionResult.msl, compileErrMsg, _mslVersionMajor, _mslVersionMinor, _mslVersionPatch); if (compileErrMsg.size() > 0) { string preamble = wasCompiled ? "is valid but the validation compilation produced warnings: " : "failed a validation compilation: "; compileErrMsg = "Generated MSL " + preamble + compileErrMsg; @@ -255,7 +254,7 @@ bool MoltenVKShaderConverterTool::convertSPIRV(const vector& spv, } vector fileContents; - fileContents.insert(fileContents.end(), msl.begin(), msl.end()); + fileContents.insert(fileContents.end(), conversionResult.msl.begin(), conversionResult.msl.end()); string writeErrMsg; if (writeFile(path, fileContents, writeErrMsg)) { string logMsg = "Saved MSL to file: " + fileName(path); diff --git a/Scripts/runcts b/Scripts/runcts index 4cc33e78..2349e75d 100755 --- a/Scripts/runcts +++ b/Scripts/runcts @@ -94,22 +94,24 @@ if [ "${is_portability}" != "" ]; then export MVK_CONFIG_ADVERTISE_EXTENSIONS=0xA fi +# ----- Metal validation settings ------ export METAL_DEVICE_WRAPPER_TYPE=1 export METAL_ERROR_MODE=3 export METAL_DEBUG_ERROR_MODE=3 # ----- MoltenVK config settings ------ -export MVK_CONFIG_LOG_LEVEL=1 +export MVK_CONFIG_LOG_LEVEL=1 #(1 = Errors only, 3 = Info) export MVK_DEBUG=0 -# Additional MoltenVK configuration can be set here by -# editing below, or can be set before calling this script. +# Additional MoltenVK configuration can be set here by editing below. export MVK_CONFIG_RESUME_LOST_DEVICE=1 export MVK_CONFIG_FAST_MATH_ENABLED=1 -export MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS=0 #(2 = VK_EXT_descriptor_indexing enabled) export MVK_CONFIG_FORCE_LOW_POWER_GPU=0 -export MVK_CONFIG_VK_SEMAPHORE_SUPPORT_STYLE=2 #(2 = MTLEvents always) - +export MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS=0 #(2 = VK_EXT_descriptor_indexing enabled) +export MVK_CONFIG_VK_SEMAPHORE_SUPPORT_STYLE=2 #(2 = MTLEvents always) +export MVK_CONFIG_SHADER_COMPRESSION_ALGORITHM=0 #(2 = ZLIB, 3 = LZ4) +export MVK_CONFIG_PERFORMANCE_TRACKING=0 +export MVK_CONFIG_ACTIVITY_PERFORMANCE_LOGGING_STYLE=2 #(2 = Device lifetime) # -------------- Operation --------------------