Merge pull request #191 from billhollings/master

Fix compute shader workgroup size specialization and correct some settings.
KhronosGroup · Jul 4, 2018 · 4dbe302 · 4dbe302
2 parents 5943aae + 6950102
commit 4dbe302
Show file tree

Hide file tree

Showing 10 changed files with 123 additions and 117 deletions.
diff --git a/ExternalRevisions/SPIRV-Cross_repo_revision b/ExternalRevisions/SPIRV-Cross_repo_revision
@@ -1 +1 @@
-d67e586b2e16a46a5cc1515093e8a04bff31c594
+a6814a405abe81545bd3b0a50d374735001173c1
diff --git a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h
@@ -48,7 +48,7 @@ extern "C" {
  */
 #define MVK_VERSION_MAJOR   1
 #define MVK_VERSION_MINOR   0
-#define MVK_VERSION_PATCH   13
+#define MVK_VERSION_PATCH   14
 
 #define MVK_MAKE_VERSION(major, minor, patch)    (((major) * 10000) + ((minor) * 100) + (patch))
 #define MVK_VERSION     MVK_MAKE_VERSION(MVK_VERSION_MAJOR, MVK_VERSION_MINOR, MVK_VERSION_PATCH)

diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -111,33 +111,42 @@
     VkPhysicalDeviceLimits* pLimits = &_properties.limits;
     VkExtent3D maxExt;
     uint32_t maxLayers;
+	uint32_t maxLevels;
     switch (type) {
         case VK_IMAGE_TYPE_1D:
+			// Metal does not allow 1D textures to be used as attachments
+			if (mvkIsAnyFlagEnabled(usage, VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT | VK_IMAGE_USAGE_DEPTH_STENCIL_ATTACHMENT_BIT)) {
+				return VK_ERROR_FORMAT_NOT_SUPPORTED;
+			}
             maxExt.width = pLimits->maxImageDimension1D;
             maxExt.height = 1;
             maxExt.depth = 1;
+			maxLevels = 1;
             maxLayers = pLimits->maxImageArrayLayers;
             break;
         case VK_IMAGE_TYPE_2D:
             maxExt.width = pLimits->maxImageDimension2D;
             maxExt.height = pLimits->maxImageDimension2D;
             maxExt.depth = 1;
+			maxLevels = mvkMipmapLevels3D(maxExt);
             maxLayers = pLimits->maxImageArrayLayers;
             break;
         case VK_IMAGE_TYPE_3D:
             maxExt.width = pLimits->maxImageDimension3D;
             maxExt.height = pLimits->maxImageDimension3D;
             maxExt.depth = pLimits->maxImageDimension3D;
+			maxLevels = mvkMipmapLevels3D(maxExt);
             maxLayers = 1;
             break;
         default:
             maxExt = { 1, 1, 1};
             maxLayers = 1;
+			maxLevels = 1;
             break;
     }
 
     pImageFormatProperties->maxExtent = maxExt;
-    pImageFormatProperties->maxMipLevels = mvkMipmapLevels3D(maxExt);
+    pImageFormatProperties->maxMipLevels = maxLevels;
     pImageFormatProperties->maxArrayLayers = maxLayers;
     pImageFormatProperties->sampleCounts = _metalFeatures.supportedSampleCounts;
     pImageFormatProperties->maxResourceSize = kMVKUndefinedLargeUInt64;
@@ -541,12 +550,12 @@
     _properties.limits.maxVertexInputAttributeOffset = (4 * KIBI);
     _properties.limits.maxVertexInputBindingStride = _properties.limits.maxVertexInputAttributeOffset - 1;
 
+	_properties.limits.maxPerStageDescriptorSamplers = _metalFeatures.maxPerStageSamplerCount;
 	_properties.limits.maxPerStageDescriptorUniformBuffers = _metalFeatures.maxPerStageBufferCount;
 	_properties.limits.maxPerStageDescriptorStorageBuffers = _metalFeatures.maxPerStageBufferCount;
 	_properties.limits.maxPerStageDescriptorSampledImages = _metalFeatures.maxPerStageTextureCount;
 	_properties.limits.maxPerStageDescriptorStorageImages = _metalFeatures.maxPerStageTextureCount;
-    _properties.limits.maxPerStageDescriptorSamplers = _metalFeatures.maxPerStageSamplerCount;
-    _properties.limits.maxDescriptorSetInputAttachments = _metalFeatures.maxPerStageTextureCount;
+	_properties.limits.maxPerStageDescriptorInputAttachments = _metalFeatures.maxPerStageTextureCount;
 
     _properties.limits.maxPerStageResources = (_metalFeatures.maxPerStageBufferCount + _metalFeatures.maxPerStageTextureCount);
     _properties.limits.maxFragmentCombinedOutputResources = _properties.limits.maxPerStageResources;
@@ -558,6 +567,7 @@
 	_properties.limits.maxDescriptorSetStorageBuffersDynamic = (_properties.limits.maxPerStageDescriptorStorageBuffers * 2);
 	_properties.limits.maxDescriptorSetSampledImages = (_properties.limits.maxPerStageDescriptorSampledImages * 2);
 	_properties.limits.maxDescriptorSetStorageImages = (_properties.limits.maxPerStageDescriptorStorageImages * 2);
+	_properties.limits.maxDescriptorSetInputAttachments = (_properties.limits.maxPerStageDescriptorInputAttachments * 2);
 
 	_properties.limits.maxTexelBufferElements = _properties.limits.maxImageDimension2D * _properties.limits.maxImageDimension2D;
 	_properties.limits.maxUniformBufferRange = (uint32_t)_metalFeatures.maxMTLBufferSize;

diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
@@ -485,16 +485,19 @@
 // Ceral archive definitions
 namespace mvk {
 
+	template<class Archive>
+	void serialize(Archive & archive, SPIRVWorkgroupSizeDimension& wsd) {
+		archive(wsd.size,
+				wsd.specializationID,
+				wsd.isSpecialized);
+	}
+
 	template<class Archive>
 	void serialize(Archive & archive, SPIRVEntryPoint& ep) {
 		archive(ep.mtlFunctionName,
 				ep.workgroupSize.width,
 				ep.workgroupSize.height,
-				ep.workgroupSize.depth,
-				ep.workgroupSizeId.width,
-				ep.workgroupSizeId.height,
-				ep.workgroupSizeId.depth,
-				ep.workgroupSizeId.constant);
+				ep.workgroupSize.depth);
 	}
 
 	template<class Archive>

diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.mm
@@ -30,14 +30,19 @@
 #pragma mark -
 #pragma mark MVKShaderLibrary
 
-static uint32_t getOffsetForConstantId(const VkSpecializationInfo* pSpecInfo, uint32_t constantId)
-{
-    for (uint32_t specIdx = 0; specIdx < pSpecInfo->mapEntryCount; specIdx++) {
-        const VkSpecializationMapEntry* pMapEntry = &pSpecInfo->pMapEntries[specIdx];
-        if (pMapEntry->constantID == constantId) { return pMapEntry->offset; }
-    }
+// If the size of the workgroup dimension is specialized, extract it from the
+// specialization info, otherwise use the value specified in the SPIR-V shader code.
+static uint32_t getWorkgroupDimensionSize(const SPIRVWorkgroupSizeDimension& wgDim, const VkSpecializationInfo* pSpecInfo) {
+	if (wgDim.isSpecialized && pSpecInfo) {
+		for (uint32_t specIdx = 0; specIdx < pSpecInfo->mapEntryCount; specIdx++) {
+			const VkSpecializationMapEntry* pMapEntry = &pSpecInfo->pMapEntries[specIdx];
+			if (pMapEntry->constantID == wgDim.specializationID) {
+				return *reinterpret_cast<uint32_t*>((uintptr_t)pSpecInfo->pData + pMapEntry->offset) ;
+			}
+		}
+	}
 
-    return -1;
+	return wgDim.size;
 }
 
 MVKMTLFunction MVKShaderLibrary::getMTLFunction(const VkSpecializationInfo* pSpecializationInfo) {
@@ -88,27 +93,10 @@ static uint32_t getOffsetForConstantId(const VkSpecializationInfo* pSpecInfo, ui
         mvkNotifyErrorWithText(VK_ERROR_INITIALIZATION_FAILED, "Shader module does not contain an entry point named '%s'.", mtlFuncName.UTF8String);
     }
 
-    if (pSpecializationInfo) {
-        // Get the specialization constant values for the work group size
-        if (_entryPoint.workgroupSizeId.constant != 0) {
-            uint32_t widthOffset = getOffsetForConstantId(pSpecializationInfo, _entryPoint.workgroupSizeId.width);
-            if (widthOffset != -1) {
-                _entryPoint.workgroupSize.width = *reinterpret_cast<uint32_t*>((uint8_t*)pSpecializationInfo->pData + widthOffset);
-            }
-
-            uint32_t heightOffset = getOffsetForConstantId(pSpecializationInfo, _entryPoint.workgroupSizeId.height);
-            if (heightOffset != -1) {
-                _entryPoint.workgroupSize.height = *reinterpret_cast<uint32_t*>((uint8_t*)pSpecializationInfo->pData + heightOffset);
-            }
-
-            uint32_t depthOffset = getOffsetForConstantId(pSpecializationInfo, _entryPoint.workgroupSizeId.depth);
-            if (depthOffset != -1) {
-                _entryPoint.workgroupSize.depth = *reinterpret_cast<uint32_t*>((uint8_t*)pSpecializationInfo->pData + depthOffset);
-            }
-        }
-    }
+	return { mtlFunc, MTLSizeMake(getWorkgroupDimensionSize(_entryPoint.workgroupSize.width, pSpecializationInfo),
+								  getWorkgroupDimensionSize(_entryPoint.workgroupSize.height, pSpecializationInfo),
+								  getWorkgroupDimensionSize(_entryPoint.workgroupSize.depth, pSpecializationInfo)) };
 
-    return { mtlFunc, MTLSizeMake(_entryPoint.workgroupSize.width, _entryPoint.workgroupSize.height, _entryPoint.workgroupSize.depth) };
 }
 
 // Returns the MTLFunctionConstant with the specified ID from the specified array of function constants.

diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.cpp
@@ -124,8 +124,8 @@ MVK_PUBLIC_SYMBOL void SPIRVToMSLConverterContext::alignUsageWith(const SPIRVToM
 #pragma mark -
 #pragma mark SPIRVToMSLConverter
 
-/** Populates content extracted from the SPRI-V compiler. */
-void populateFromCompiler(spirv_cross::Compiler* pCompiler, SPIRVEntryPoint& entryPoint, SPIRVToMSLConverterOptions& options);
+// Populates the entry point with info extracted from the SPRI-V compiler.
+void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options);
 
 MVK_PUBLIC_SYMBOL void SPIRVToMSLConverter::setSPIRV(const vector<uint32_t>& spirv) { _spirv = spirv; }
 
@@ -224,7 +224,7 @@ MVK_PUBLIC_SYMBOL bool SPIRVToMSLConverter::convert(SPIRVToMSLConverterContext&
 #endif
 
     // Populate content extracted from the SPRI-V compiler.
-	populateFromCompiler(pMSLCompiler, _entryPoint, context.options);
+	populateEntryPoint(_entryPoint, pMSLCompiler, context.options);
 
     // To check GLSL conversion
     if (shouldLogGLSL) {
@@ -334,7 +334,14 @@ void SPIRVToMSLConverter::logSource(string& src, const char* srcLang, const char
 
 #pragma mark Support functions
 
-void populateFromCompiler(spirv_cross::Compiler* pCompiler, SPIRVEntryPoint& entryPoint, SPIRVToMSLConverterOptions& options) {
+// Populate a workgroup size dimension.
+void populateWorkgroupDimension(SPIRVWorkgroupSizeDimension& wgDim, uint32_t size, spirv_cross::SpecializationConstant& spvSpecConst) {
+	wgDim.size = max(size, 1u);
+	wgDim.isSpecialized = (spvSpecConst.id != 0);
+	wgDim.specializationID = spvSpecConst.constant_id;
+}
+
+void populateEntryPoint(SPIRVEntryPoint& entryPoint, spirv_cross::Compiler* pCompiler, SPIRVToMSLConverterOptions& options) {
 
 	if ( !pCompiler ) { return; }
 
@@ -349,19 +356,13 @@ void populateFromCompiler(spirv_cross::Compiler* pCompiler, SPIRVEntryPoint& ent
 		}
 	}
 
-	uint32_t minDim = 1;
-	auto& wgSize = spvEP.workgroup_size;
+	spirv_cross::SpecializationConstant widthSC, heightSC, depthSC;
+	pCompiler->get_work_group_size_specialization_constants(widthSC, heightSC, depthSC);
 
 	entryPoint.mtlFunctionName = spvEP.name;
-	entryPoint.workgroupSize.width = max(wgSize.x, minDim);
-	entryPoint.workgroupSize.height = max(wgSize.y, minDim);
-	entryPoint.workgroupSize.depth = max(wgSize.z, minDim);
-
-	spirv_cross::SpecializationConstant width, height, depth;
-	entryPoint.workgroupSizeId.constant = pCompiler->get_work_group_size_specialization_constants(width, height, depth);
-	entryPoint.workgroupSizeId.width = width.constant_id;
-	entryPoint.workgroupSizeId.height = height.constant_id;
-	entryPoint.workgroupSizeId.depth = depth.constant_id;
+	populateWorkgroupDimension(entryPoint.workgroupSize.width, spvEP.workgroup_size.x, widthSC);
+	populateWorkgroupDimension(entryPoint.workgroupSize.height, spvEP.workgroup_size.y, heightSC);
+	populateWorkgroupDimension(entryPoint.workgroupSize.depth, spvEP.workgroup_size.z, depthSC);
 }
 
 MVK_PUBLIC_SYMBOL void mvk::logSPIRV(vector<uint32_t>& spirv, string& spvLog) {

diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h
@@ -138,24 +138,29 @@ namespace mvk {
 	} SPIRVToMSLConverterContext;
 
     /**
+     * Describes one dimension of the workgroup size of a SPIR-V entry point, including whether
+	 * it is specialized, and if so, the value of the corresponding specialization ID, which
+	 * is used to map to a value which will be provided when the MSL is compiled into a pipeline.
+     */
+	typedef struct {
+		uint32_t size = 1;
+		uint32_t specializationID = 0;
+		bool isSpecialized = false;
+	} SPIRVWorkgroupSizeDimension;
+
+	/**
      * Describes a SPIRV entry point, including the Metal function name (which may be
      * different than the Vulkan entry point name if the original name was illegal in Metal),
-     * and the number of threads in each workgroup or their specialization constant id, if the shader is a compute shader.
+     * and the size of each workgroup, if the shader is a compute shader.
      */
-    typedef struct {
-        std::string mtlFunctionName = "main0";
-        struct {
-            uint32_t width = 1;
-            uint32_t height = 1;
-            uint32_t depth = 1;
-        } workgroupSize;
-        struct {
-			uint32_t width = 1;
-			uint32_t height = 1;
-			uint32_t depth = 1;
-            uint32_t constant = 0;
-        } workgroupSizeId;
-    } SPIRVEntryPoint;
+	typedef struct {
+		std::string mtlFunctionName = "main0";
+		struct {
+			SPIRVWorkgroupSizeDimension width;
+			SPIRVWorkgroupSizeDimension height;
+			SPIRVWorkgroupSizeDimension depth;
+		} workgroupSize;
+	} SPIRVEntryPoint;
 
 	/** Special constant used in a MSLResourceBinding descriptorSet element to indicate the bindings for the push constants. */
     static const uint32_t kPushConstDescSet = std::numeric_limits<uint32_t>::max();

diff --git a/...MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme b/...MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme
@@ -67,23 +67,27 @@
             isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
-            argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/MoltenVK/External/SPIRV-Cross/shaders-msl"
+            argument = "path-to-shader-directory"
+            isEnabled = "NO">
+         </CommandLineArgument>
+         <CommandLineArgument
+            argument = "-r"
             isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
             argument = "-gi"
             isEnabled = "YES">
          </CommandLineArgument>
          <CommandLineArgument
-            argument = "/Users/bill/Desktop/texture_buffer.vert"
+            argument = "path-to-GLSL-shader-file"
             isEnabled = "YES">
          </CommandLineArgument>
          <CommandLineArgument
             argument = "-si"
             isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
-            argument = "/Users/bill/Documents/Dev/iOSProjects/Molten/Support/2018/MVK_Issue_112/second/vert_bin.spv"
+            argument = "path-to-SPIR-V-shader-file"
             isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
Original file line number	Diff line number	Diff line change
		@@ -1 +1 @@
		d67e586b2e16a46a5cc1515093e8a04bff31c594
		a6814a405abe81545bd3b0a50d374735001173c1