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

Merge pull request #1847 from billhollings/pipeline-cache-mem-reduction

Browse Source 
VK_EXT_pipeline_creation_cache_control & reduce memory footprint of retained MSL source code
This commit is contained in:
Bill Hollings 2023-03-07 10:15:42 -05:00 committed by GitHub
commit 4421883eeb
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
25 changed files with 766 additions and 456 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
Docs/MoltenVK_Runtime_UserGuide.md
View File

@ -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`

11
Docs/Whats_New.md
View File

@ -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`.

16
MoltenVK/MoltenVK.xcodeproj/project.pbxproj
View File

@ -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 = "<group>"; };
4553AEF62251617100E8EBCD /* MVKBlockObserver.m */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.objc; path = MVKBlockObserver.m; sourceTree = "<group>"; };
4553AEFA2251617100E8EBCD /* MVKBlockObserver.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKBlockObserver.h; sourceTree = "<group>"; };
45557A4D21C9EFF3008868BD /* MVKCodec.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKCodec.cpp; sourceTree = "<group>"; };
45557A4D21C9EFF3008868BD /* MVKCodec.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCodec.mm; sourceTree = "<group>"; };
45557A5121C9EFF3008868BD /* MVKCodec.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCodec.h; sourceTree = "<group>"; };
45557A5721CD83C3008868BD /* MVKDXTnCodec.def */ = {isa = PBXFileReference; explicitFileType = sourcecode.cpp.h; fileEncoding = 4; path = MVKDXTnCodec.def; sourceTree = "<group>"; };
A9096E5C1F81E16300DFBEA6 /* MVKCmdDispatch.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKCmdDispatch.h; sourceTree = "<group>"; };
@ -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 */,

65
MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h
View File

@ -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.

6
MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
View File

@ -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<MTLSamplerState> _defaultMTLSamplerState = nil;
id<MTLBuffer> _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;

78
MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
View File

@ -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];

67
MoltenVK/MoltenVK/GPUObjects/MVKDeviceFeatureStructs.def
View File

@ -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

33
MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
View File

@ -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<MVKShaderModuleKey, MVKShaderLibraryCache*> _shaderCache;
size_t _dataSize = 0;
std::mutex _shaderCacheLock;
bool _isExternallySynchronized = false;
};

172
MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
View File

@ -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<MTLFunction> 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<MTLFunction> 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<MTLFunction> 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<MTLFunction> 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<MTLFunction> 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<mutex> 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<mutex> 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<std::string>& 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<mutex> 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<mutex> 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<std::string> 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<mutex> 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<class Archive>
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<class Archive, class C>
void serialize(Archive & archive, MVKCompressor<C>& 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);
}

60
MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h
View File

@ -20,6 +20,7 @@
#include "MVKDevice.h"
#include "MVKSync.h"
#include "MVKCodec.h"
#include "MVKSmallVector.h"
#include <MoltenVKShaderConverter/SPIRVToMSLConverter.h>
#include <MoltenVKShaderConverter/GLSLToSPIRVConverter.h>
@ -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<MTLFunction> getMTLFunction() { return _mtlFunction; }
MVKMTLFunction(id<MTLFunction> mtlFunc, const SPIRVToMSLConversionResults scRslts, MTLSize tgSize);
MVKMTLFunction(id<MTLFunction> 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<std::string> 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<MTLFunctionConstant*>* mtlFCs, NSUInteger mtlFCID);
void compileLibrary(const std::string& msl);
void compressMSL(const std::string& msl);
void decompressMSL(std::string& msl);
MVKCompressor<std::string>& getCompressedMSL() { return _compressedMSL; }
MVKVulkanAPIDeviceObject* _owner;
id<MTLLibrary> _mtlLibrary;
SPIRVToMSLConversionResults _shaderConversionResults;
std::string _msl;
MVKCompressor<std::string> _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<std::string> 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<uint32_t>& 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<MTLLibrary> newMTLLibrary(NSString* mslSourceCode,
const SPIRVToMSLConversionResults& shaderConversionResults);
const SPIRVToMSLConversionResultInfo& shaderConversionResults);
#pragma mark Construction

138
MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm
View File

@ -19,12 +19,11 @@
#include "MVKShaderModule.h"
#include "MVKPipeline.h"
#include "MVKFoundation.h"
#include <string>
using namespace std;
MVKMTLFunction::MVKMTLFunction(id<MTLFunction> mtlFunc, const SPIRVToMSLConversionResults scRslts, MTLSize tgSize) {
MVKMTLFunction::MVKMTLFunction(id<MTLFunction> 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<std::string> 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<std::string> 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<mutex> 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<mutex> 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<MTLLibrary> MVKShaderLibraryCompiler::newMTLLibrary(NSString* mslSourceCode,
const SPIRVToMSLConversionResults& shaderConversionResults) {
const SPIRVToMSLConversionResultInfo& shaderConversionResults) {
unique_lock<mutex> lock(_completionLock);
compile(lock, ^{

1
MoltenVK/MoltenVK/Layers/MVKExtensions.def
View File

@ -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)

1
MoltenVK/MoltenVK/Layers/MVKLayers.mm
View File

@ -19,7 +19,6 @@
#include "MVKLayers.h"
#include "MVKEnvironment.h"
#include "MVKFoundation.h"
#include "vk_mvk_moltenvk.h"
#include <mutex>
using namespace std;

88
MoltenVK/MoltenVK/Utility/MVKCodec.h
View File

@ -19,17 +19,18 @@
#pragma once
#include "MVKFoundation.h"
#include "MVKEnvironment.h"
#include <vector>
#include <string>
#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 C>
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<uint8_t> _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<MVKCodec> 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);

75
MoltenVK/MoltenVK/Utility/MVKCodec.cpp → MoltenVK/MoltenVK/Utility/MVKCodec.mm
View File

@ -18,10 +18,13 @@
#include "MVKCodec.h"
#include "MVKBaseObject.h"
#include "MVKFoundation.h"
#include <algorithm>
#include <simd/simd.h>
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<MVKCodec> mvkCreateCodec(VkFormat format) {
#pragma mark -
#pragma mark Support functions
unique_ptr<MVKCodec> 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<MVKCodec> 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<MVKCodec>(new MVKDXTnCodec(format));
return unique_ptr<MVKCodec>(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);
}

14
MoltenVK/MoltenVK/Utility/MVKEnvironment.cpp
View File

@ -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);
}

12
MoltenVK/MoltenVK/Utility/MVKEnvironment.h
View File

@ -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

16
MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLConversion.mm
View File

@ -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;

73
MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.cpp
View File

@ -53,13 +53,12 @@ MVK_PUBLIC_SYMBOL void GLSLToSPIRVConverter::setGLSLs(const std::vector<std::str
}
MVK_PUBLIC_SYMBOL bool GLSLToSPIRVConverter::convert(MVKGLSLConversionShaderStage shaderStage,
GLSLToSPIRVConversionResult& conversionResult,
bool shouldLogGLSL,
bool shouldLogSPIRV) {
_wasConverted = true;
_resultLog.clear();
_spirv.clear();
bool wasConverted = true;
if (shouldLogGLSL) { logGLSL("Converting"); }
if (shouldLogGLSL) { logGLSL(conversionResult.resultLog, "Converting"); }
EShMessages messages = (EShMessages)(EShMsgDefault | EShMsgSpvRules | EShMsgVulkanRules);
@ -79,13 +78,13 @@ MVK_PUBLIC_SYMBOL bool GLSLToSPIRVConverter::convert(MVKGLSLConversionShaderStag
glslShaders.back()->setAutoMapBindings(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<uint32_t> 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";
}

48
MoltenVKShaderConverter/MoltenVKShaderConverter/GLSLToSPIRVConverter.h
View File

@ -27,6 +27,17 @@
namespace mvk {
#pragma mark -
#pragma mark SPIRVToMSLConversionResult
/** The results of a GLSL to SPIRV conversion. */
typedef struct GLSLToSPIRVConversionResult {
std::vector<uint32_t> 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<std::string>& 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<uint32_t>& 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<uint32_t> spirv);
void initGLSLCompilerResources();
std::vector<std::string> _glsls;
std::vector<uint32_t> _spirv;
std::string _resultLog;
bool _wasConverted = false;
};
}

13
MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVConversion.mm
View File

@ -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;

117
MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.cpp
View File

@ -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<char> 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);

76
MoltenVKShaderConverter/MoltenVKShaderConverter/SPIRVToMSLConverter.h
View File

@ -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<uint32_t> _spirv;
std::string _msl;
std::string _resultLog;
SPIRVToMSLConversionResults _shaderConversionResults;
bool _wasConverted = false;
};
}

27
MoltenVKShaderConverter/MoltenVKShaderConverterTool/MoltenVKShaderConverterTool.cpp
View File

@ -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<uint32_t>& 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<uint32_t>& 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<uint32_t>& spv,
}
vector<char> 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);

14
Scripts/runcts
View File

@ -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 --------------------