From 396add581a1dac101bc6998f0daf9f94d8f6c256 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Mon, 27 Aug 2018 12:36:45 -0500 Subject: [PATCH 01/20] Support the VK_KHR_maintenance1 extension. Much of this was already supported, simply by being supported by Metal. Of course, this support is incomplete: Metal doesn't yet allow you to create a 2D texture view from a 3D texture. --- MoltenVK/MoltenVK/Commands/MVKCommandPool.h | 3 +++ MoltenVK/MoltenVK/Commands/MVKCommandPool.mm | 4 ++++ MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm | 2 +- MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm | 1 + MoltenVK/MoltenVK/Loader/MVKLayers.mm | 5 +++++ MoltenVK/MoltenVK/Vulkan/vulkan.mm | 12 ++++++++++++ 6 files changed, 26 insertions(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h index d08045ab..2037bee7 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h @@ -142,6 +142,9 @@ public: void freeCommandBuffers(uint32_t commandBufferCount, const VkCommandBuffer* pCommandBuffers); + /** Release any held but unused memory back to the system. */ + void trimCommandPool(); + #pragma mark Construction diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm index c511f4f9..bf21b894 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm @@ -63,6 +63,10 @@ void MVKCommandPool::freeCommandBuffers(uint32_t commandBufferCount, } } +void MVKCommandPool::trimCommandPool() { + // TODO: Implement. +} + void MVKCommandPool::addCommandBuffer(MVKCommandBuffer* cmdBuffer) { _commandBuffers.insert(cmdBuffer); } diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm index 3ad2aaab..9b514f7d 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm @@ -611,7 +611,7 @@ VkResult MVKDescriptorPool::allocateDescriptorSets(uint32_t count, const VkDescriptorSetLayout* pSetLayouts, VkDescriptorSet* pDescriptorSets) { if (_allocatedSetCount + count > _maxSets) { - return mvkNotifyErrorWithText(VK_ERROR_INITIALIZATION_FAILED, "The maximum number of descriptor sets that can be allocated by this descriptor pool is %d.", _maxSets); + return mvkNotifyErrorWithText(VK_ERROR_OUT_OF_POOL_MEMORY_KHR, "The maximum number of descriptor sets that can be allocated by this descriptor pool is %d.", _maxSets); } for (uint32_t dsIdx = 0; dsIdx < count; dsIdx++) { diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index b50e6447..7ab058d0 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -241,6 +241,7 @@ void MVKInstance::initProcAddrs() { ADD_PROC_ADDR(vkGetSwapchainImagesKHR); ADD_PROC_ADDR(vkAcquireNextImageKHR); ADD_PROC_ADDR(vkQueuePresentKHR); + ADD_PROC_ADDR(vkTrimCommandPoolKHR); ADD_PROC_ADDR(vkGetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkSetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkGetPhysicalDeviceMetalFeaturesMVK); diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.mm b/MoltenVK/MoltenVK/Loader/MVKLayers.mm index 2b514ad5..fed05c0a 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Loader/MVKLayers.mm @@ -92,6 +92,11 @@ MVKLayer::MVKLayer() { extTmplt.specVersion = VK_AMD_NEGATIVE_VIEWPORT_HEIGHT_SPEC_VERSION; _extensions.push_back(extTmplt); + memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); + strcpy(extTmplt.extensionName, VK_KHR_MAINTENANCE1_EXTENSION_NAME); + extTmplt.specVersion = VK_KHR_MAINTENANCE1_SPEC_VERSION; + _extensions.push_back(extTmplt); + #if MVK_IOS memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); strcpy(extTmplt.extensionName, VK_MVK_IOS_SURFACE_EXTENSION_NAME); diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index 28fa8a96..d7cec0e8 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -1601,6 +1601,18 @@ MVK_PUBLIC_SYMBOL VkResult vkCreate_PLATFORM_SurfaceMVK( } +#pragma mark - +#pragma mark VK_KHR_maintenace1 extension + +MVK_PUBLIC_SYMBOL void vkTrimCommandPoolKHR( + VkDevice device, + VkCommandPool commandPool, + VkCommandPoolTrimFlagsKHR flags) { + MVKCommandPool* mvkCmdPool = (MVKCommandPool*)commandPool; + mvkCmdPool->trimCommandPool(); +} + + #pragma mark - #pragma mark Loader and Layer ICD interface extension From ccb1afe42b512f2f7510e2e709c563fea1925f66 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Mon, 27 Aug 2018 12:36:52 -0500 Subject: [PATCH 02/20] Support the VK_KHR_shader_draw_parameters extension. Update SPIRV-Cross to support shaders which use the builtins provided by this extension. --- ExternalRevisions/SPIRV-Cross_repo_revision | 2 +- MoltenVK/MoltenVK/Loader/MVKLayers.mm | 5 +++++ 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/ExternalRevisions/SPIRV-Cross_repo_revision b/ExternalRevisions/SPIRV-Cross_repo_revision index 13a77c96..7e060c34 100644 --- a/ExternalRevisions/SPIRV-Cross_repo_revision +++ b/ExternalRevisions/SPIRV-Cross_repo_revision @@ -1 +1 @@ -6480db7352b154f7decf8df9eb38b4c3c1ec530b +e14bf77b1ac99943aa27c6b9f6446ea2c4a824f7 diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.mm b/MoltenVK/MoltenVK/Loader/MVKLayers.mm index fed05c0a..aa30e7ee 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Loader/MVKLayers.mm @@ -97,6 +97,11 @@ MVKLayer::MVKLayer() { extTmplt.specVersion = VK_KHR_MAINTENANCE1_SPEC_VERSION; _extensions.push_back(extTmplt); + memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); + strcpy(extTmplt.extensionName, VK_KHR_SHADER_DRAW_PARAMETERS_EXTENSION_NAME); + extTmplt.specVersion = VK_KHR_SHADER_DRAW_PARAMETERS_SPEC_VERSION; + _extensions.push_back(extTmplt); + #if MVK_IOS memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); strcpy(extTmplt.extensionName, VK_MVK_IOS_SURFACE_EXTENSION_NAME); From 80bde624cf2f54adb3f951495471f65e057103ca Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Mon, 27 Aug 2018 12:36:57 -0500 Subject: [PATCH 03/20] Support the VK_KHR_get_physical_device_properties2 extension. This extension is a prerequisite for multiple other extensions, the `VK_KHR_push_descriptor` extension in particular. --- MoltenVK/MoltenVK/GPUObjects/MVKDevice.h | 33 +++++++++ MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm | 75 +++++++++++++++++++++ MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm | 7 ++ MoltenVK/MoltenVK/Loader/MVKLayers.mm | 5 ++ MoltenVK/MoltenVK/Vulkan/vulkan.mm | 64 ++++++++++++++++++ 5 files changed, 184 insertions(+) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index 6ca8543d..4eb86f48 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -75,18 +75,27 @@ public: /** Populates the specified structure with the features of this device. */ void getFeatures(VkPhysicalDeviceFeatures* features); + /** Populates the specified structure with the features of this device. */ + void getFeatures(VkPhysicalDeviceFeatures2KHR* features); + /** Populates the specified structure with the Metal-specific features of this device. */ void getMetalFeatures(MVKPhysicalDeviceMetalFeatures* mtlFeatures); /** Populates the specified structure with the properties of this device. */ void getProperties(VkPhysicalDeviceProperties* properties); + /** Populates the specified structure with the properties of this device. */ + void getProperties(VkPhysicalDeviceProperties2KHR* properties); + /** Returns whether the specified format is supported on this device. */ bool getFormatIsSupported(VkFormat format); /** Populates the specified structure with the format properties of this device. */ void getFormatProperties(VkFormat format, VkFormatProperties* pFormatProperties); + /** Populates the specified structure with the format properties of this device. */ + void getFormatProperties(VkFormat format, VkFormatProperties2KHR* pFormatProperties); + /** * Populates the specified structure with the image format properties * supported for the specified image characteristics on this device. @@ -98,6 +107,13 @@ public: VkImageCreateFlags flags, VkImageFormatProperties* pImageFormatProperties); + /** + * Populates the specified structure with the image format properties + * supported for the specified image characteristics on this device. + */ + VkResult getImageFormatProperties(const VkPhysicalDeviceImageFormatInfo2KHR* pImageFormatInfo, + VkImageFormatProperties2KHR* pImageFormatProperties); + #pragma mark Surfaces /** @@ -158,6 +174,20 @@ public: */ VkResult getQueueFamilyProperties(uint32_t* pCount, VkQueueFamilyProperties* properties); + /** + * If properties is null, the value of pCount is updated with the number of + * queue families supported by this instance. + * + * If properties is not null, then pCount queue family properties are copied into the + * array. If the number of available queue families is less than pCount, the value of + * pCount is updated to indicate the number of queue families actually returned in the array. + * + * Returns VK_SUCCESS if successful. Returns VK_INCOMPLETE if the number of queue families + * available in this instance is larger than the specified pCount. Returns other values if + * an error occurs. + */ + VkResult getQueueFamilyProperties(uint32_t* pCount, VkQueueFamilyProperties2KHR* properties); + /** Returns a pointer to the Vulkan instance. */ inline MVKInstance* getInstance() { return _mvkInstance; } @@ -173,6 +203,9 @@ public: /** Populates the specified memory properties with the memory characteristics of this device. */ VkResult getPhysicalDeviceMemoryProperties(VkPhysicalDeviceMemoryProperties* pMemoryProperties); + /** Populates the specified memory properties with the memory characteristics of this device. */ + VkResult getPhysicalDeviceMemoryProperties(VkPhysicalDeviceMemoryProperties2KHR* pMemoryProperties); + /** * Returns a bit mask of all memory type indices. * Each bit [0..31] in the returned bit mask indicates a distinct memory type. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 19bde85f..89411c40 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -58,6 +58,13 @@ void MVKPhysicalDevice::getFeatures(VkPhysicalDeviceFeatures* features) { if (features) { *features = _features; } } +void MVKPhysicalDevice::getFeatures(VkPhysicalDeviceFeatures2KHR* features) { + if (features) { + features->sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2_KHR; + features->features = _features; + } +} + void MVKPhysicalDevice::getMetalFeatures(MVKPhysicalDeviceMetalFeatures* mtlFeatures) { if (mtlFeatures) { *mtlFeatures = _metalFeatures; } } @@ -66,6 +73,13 @@ void MVKPhysicalDevice::getProperties(VkPhysicalDeviceProperties* properties) { if (properties) { *properties = _properties; } } +void MVKPhysicalDevice::getProperties(VkPhysicalDeviceProperties2KHR* properties) { + if (properties) { + properties->sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR; + properties->properties = _properties; + } +} + bool MVKPhysicalDevice::getFormatIsSupported(VkFormat format) { if ( !mvkVkFormatIsSupported(format) ) { return false; } @@ -91,6 +105,15 @@ void MVKPhysicalDevice::getFormatProperties(VkFormat format, VkFormatProperties* } } +void MVKPhysicalDevice::getFormatProperties(VkFormat format, + VkFormatProperties2KHR* pFormatProperties) { + static VkFormatProperties noFmtFeats = { 0, 0, 0 }; + if (pFormatProperties) { + pFormatProperties->sType = VK_STRUCTURE_TYPE_FORMAT_PROPERTIES_2_KHR; + pFormatProperties->formatProperties = getFormatIsSupported(format) ? mvkVkFormatProperties(format) : noFmtFeats; + } +} + VkResult MVKPhysicalDevice::getImageFormatProperties(VkFormat format, VkImageType type, VkImageTiling tiling, @@ -148,6 +171,25 @@ VkResult MVKPhysicalDevice::getImageFormatProperties(VkFormat format, return VK_SUCCESS; } +VkResult MVKPhysicalDevice::getImageFormatProperties(const VkPhysicalDeviceImageFormatInfo2KHR *pImageFormatInfo, + VkImageFormatProperties2KHR* pImageFormatProperties) { + + if ( !pImageFormatInfo || pImageFormatInfo->sType != VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGE_FORMAT_INFO_2_KHR ) { + return VK_ERROR_FORMAT_NOT_SUPPORTED; + } + if ( !getFormatIsSupported(pImageFormatInfo->format) ) { return VK_ERROR_FORMAT_NOT_SUPPORTED; } + + if ( !pImageFormatProperties ) { + return VK_SUCCESS; + } + + pImageFormatProperties->sType = VK_STRUCTURE_TYPE_IMAGE_FORMAT_PROPERTIES_2_KHR; + return getImageFormatProperties(pImageFormatInfo->format, pImageFormatInfo->type, + pImageFormatInfo->tiling, pImageFormatInfo->usage, + pImageFormatInfo->flags, + &pImageFormatProperties->imageFormatProperties); +} + #pragma mark Surfaces @@ -293,6 +335,31 @@ VkResult MVKPhysicalDevice::getQueueFamilyProperties(uint32_t* pCount, return (*pCount <= qfCnt) ? VK_SUCCESS : VK_INCOMPLETE; } +VkResult MVKPhysicalDevice::getQueueFamilyProperties(uint32_t* pCount, + VkQueueFamilyProperties2KHR* queueProperties) { + + uint32_t qfCnt = uint32_t(_queueFamilies.size()); + + // If properties aren't actually being requested yet, simply update the returned count + if ( !queueProperties ) { + *pCount = qfCnt; + return VK_SUCCESS; + } + + // Determine how many families we'll return, and return that number + *pCount = min(*pCount, qfCnt); + + // Now populate the queue families + if (queueProperties) { + for (uint32_t qfIdx = 0; qfIdx < *pCount; qfIdx++) { + queueProperties[qfIdx].sType = VK_STRUCTURE_TYPE_QUEUE_FAMILY_PROPERTIES_2_KHR; + _queueFamilies[qfIdx]->getProperties(&queueProperties[qfIdx].queueFamilyProperties); + } + } + + return (*pCount <= qfCnt) ? VK_SUCCESS : VK_INCOMPLETE; +} + #pragma mark Memory models @@ -302,6 +369,14 @@ VkResult MVKPhysicalDevice::getPhysicalDeviceMemoryProperties(VkPhysicalDeviceMe return VK_SUCCESS; } +VkResult MVKPhysicalDevice::getPhysicalDeviceMemoryProperties(VkPhysicalDeviceMemoryProperties2KHR* pMemoryProperties) { + if (pMemoryProperties) { + pMemoryProperties->sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2_KHR; + pMemoryProperties->memoryProperties = _memoryProperties; + } + return VK_SUCCESS; +} + #pragma mark Construction diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index 7ab058d0..df5b9ff9 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -242,6 +242,13 @@ void MVKInstance::initProcAddrs() { ADD_PROC_ADDR(vkAcquireNextImageKHR); ADD_PROC_ADDR(vkQueuePresentKHR); ADD_PROC_ADDR(vkTrimCommandPoolKHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceFeatures2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceProperties2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceFormatProperties2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceImageFormatProperties2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceQueueFamilyProperties2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceMemoryProperties2KHR); + ADD_PROC_ADDR(vkGetPhysicalDeviceSparseImageFormatProperties2KHR); ADD_PROC_ADDR(vkGetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkSetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkGetPhysicalDeviceMetalFeaturesMVK); diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.mm b/MoltenVK/MoltenVK/Loader/MVKLayers.mm index aa30e7ee..bc3b9e8b 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Loader/MVKLayers.mm @@ -102,6 +102,11 @@ MVKLayer::MVKLayer() { extTmplt.specVersion = VK_KHR_SHADER_DRAW_PARAMETERS_SPEC_VERSION; _extensions.push_back(extTmplt); + memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); + strcpy(extTmplt.extensionName, VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME); + extTmplt.specVersion = VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_SPEC_VERSION; + _extensions.push_back(extTmplt); + #if MVK_IOS memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); strcpy(extTmplt.extensionName, VK_MVK_IOS_SURFACE_EXTENSION_NAME); diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index d7cec0e8..a9434f3a 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -1613,6 +1613,70 @@ MVK_PUBLIC_SYMBOL void vkTrimCommandPoolKHR( } +#pragma mark - +#pragma mark VK_KHR_get_physical_device_properties2 extension + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceFeatures2KHR( + VkPhysicalDevice physicalDevice, + VkPhysicalDeviceFeatures2KHR* pFeatures) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + mvkPD->getFeatures(pFeatures); +} + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceProperties2KHR( + VkPhysicalDevice physicalDevice, + VkPhysicalDeviceProperties2KHR* pProperties) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + mvkPD->getProperties(pProperties); +} + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceFormatProperties2KHR( + VkPhysicalDevice physicalDevice, + VkFormat format, + VkFormatProperties2KHR* pFormatProperties) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + mvkPD->getFormatProperties(format, pFormatProperties); +} + +MVK_PUBLIC_SYMBOL VkResult vkGetPhysicalDeviceImageFormatProperties2KHR( + VkPhysicalDevice physicalDevice, + const VkPhysicalDeviceImageFormatInfo2KHR* pImageFormatInfo, + VkImageFormatProperties2KHR* pImageFormatProperties) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + return mvkPD->getImageFormatProperties(pImageFormatInfo, pImageFormatProperties); +} + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceQueueFamilyProperties2KHR( + VkPhysicalDevice physicalDevice, + uint32_t* pQueueFamilyPropertyCount, + VkQueueFamilyProperties2KHR* pQueueFamilyProperties) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + mvkPD->getQueueFamilyProperties(pQueueFamilyPropertyCount, pQueueFamilyProperties); +} + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceMemoryProperties2KHR( + VkPhysicalDevice physicalDevice, + VkPhysicalDeviceMemoryProperties2KHR* pMemoryProperties) { + + MVKPhysicalDevice* mvkPD = MVKPhysicalDevice::getMVKPhysicalDevice(physicalDevice); + mvkPD->getPhysicalDeviceMemoryProperties(pMemoryProperties); +} + +MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceSparseImageFormatProperties2KHR( + VkPhysicalDevice physicalDevice, + const VkPhysicalDeviceSparseImageFormatInfo2KHR* pFormatInfo, + uint32_t* pPropertyCount, + VkSparseImageFormatProperties2KHR* pProperties) { + + MVKLogUnimplemented("vkGetPhysicalDeviceSparseImageFormatProperties"); +} + + #pragma mark - #pragma mark Loader and Layer ICD interface extension From 54f69c30e7f2e5938dd5b78e26a0785f93a65fb2 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Mon, 27 Aug 2018 12:37:01 -0500 Subject: [PATCH 04/20] Support the VK_KHR_push_descriptor extension. --- MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h | 37 ++++ MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm | 68 +++++++ MoltenVK/MoltenVK/Commands/MVKCommandPool.h | 2 + MoltenVK/MoltenVK/Commands/MVKCommandPool.mm | 3 +- .../MoltenVK/GPUObjects/MVKDescriptorSet.h | 20 ++ .../MoltenVK/GPUObjects/MVKDescriptorSet.mm | 191 ++++++++++++++++++ MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm | 14 ++ MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm | 1 + MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h | 5 + MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm | 11 + MoltenVK/MoltenVK/Loader/MVKLayers.mm | 5 + MoltenVK/MoltenVK/Vulkan/vulkan.mm | 16 ++ 12 files changed, 372 insertions(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h index dfbd7403..f9fd523a 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h +++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h @@ -131,6 +131,35 @@ private: }; +#pragma mark - +#pragma mark MVKCmdPushDescriptorSet + +/** Vulkan command to update a descriptor set. */ +class MVKCmdPushDescriptorSet : public MVKCommand { + +public: + void setContent(VkPipelineBindPoint pipelineBindPoint, + VkPipelineLayout layout, + uint32_t set, + uint32_t descriptorWriteCount, + const VkWriteDescriptorSet* pDescriptorWrites); + + void encode(MVKCommandEncoder* cmdEncoder) override; + + MVKCmdPushDescriptorSet(MVKCommandTypePool* pool); + + ~MVKCmdPushDescriptorSet() override; + +private: + void clearDescriptorWrites(); + + VkPipelineBindPoint _pipelineBindPoint; + MVKPipelineLayout* _pipelineLayout; + std::vector _descriptorWrites; + uint32_t _set; +}; + + #pragma mark - #pragma mark Command creation functions @@ -168,3 +197,11 @@ void mvkCmdPushConstants(MVKCommandBuffer* cmdBuff, uint32_t offset, uint32_t size, const void* pValues); + +/** Adds commands to the specified command buffer that update the specified descriptor set. */ +void mvkCmdPushDescriptorSet(MVKCommandBuffer* cmdBuff, + VkPipelineBindPoint pipelineBindPoint, + VkPipelineLayout layout, + uint32_t set, + uint32_t descriptorWriteCount, + const VkWriteDescriptorSet* pDescriptorWrites); diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm index 8383d653..d21af9ff 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm @@ -179,6 +179,64 @@ MVKCmdPushConstants::MVKCmdPushConstants(MVKCommandTypePool : MVKCommand::MVKCommand((MVKCommandTypePool*)pool) {} +#pragma mark - +#pragma mark MVKCmdPushDescriptorSet + +void MVKCmdPushDescriptorSet::setContent(VkPipelineBindPoint pipelineBindPoint, + VkPipelineLayout layout, + uint32_t set, + uint32_t descriptorWriteCount, + const VkWriteDescriptorSet* pDescriptorWrites) { + _pipelineBindPoint = pipelineBindPoint; + _pipelineLayout = (MVKPipelineLayout*)layout; + _set = set; + + // Add the descriptor writes + clearDescriptorWrites(); // Clear for reuse + _descriptorWrites.reserve(descriptorWriteCount); + for (uint32_t dwIdx = 0; dwIdx < descriptorWriteCount; dwIdx++) { + _descriptorWrites.push_back(pDescriptorWrites[dwIdx]); + VkWriteDescriptorSet& descWrite = _descriptorWrites.back(); + // Make a copy of the associated data. + if (descWrite.pImageInfo) { + auto* pNewImageInfo = new VkDescriptorImageInfo[descWrite.descriptorCount]; + std::copy_n(descWrite.pImageInfo, descWrite.descriptorCount, pNewImageInfo); + descWrite.pImageInfo = pNewImageInfo; + } + if (descWrite.pBufferInfo) { + auto* pNewBufferInfo = new VkDescriptorBufferInfo[descWrite.descriptorCount]; + std::copy_n(descWrite.pBufferInfo, descWrite.descriptorCount, pNewBufferInfo); + descWrite.pBufferInfo = pNewBufferInfo; + } + if (descWrite.pTexelBufferView) { + auto* pNewTexelBufferView = new VkBufferView[descWrite.descriptorCount]; + std::copy_n(descWrite.pTexelBufferView, descWrite.descriptorCount, pNewTexelBufferView); + descWrite.pTexelBufferView = pNewTexelBufferView; + } + } +} + +void MVKCmdPushDescriptorSet::encode(MVKCommandEncoder* cmdEncoder) { + _pipelineLayout->pushDescriptorSet(cmdEncoder, _descriptorWrites, _set); +} + +MVKCmdPushDescriptorSet::MVKCmdPushDescriptorSet(MVKCommandTypePool* pool) + : MVKCommand::MVKCommand((MVKCommandTypePool*)pool) {} + +MVKCmdPushDescriptorSet::~MVKCmdPushDescriptorSet() { + clearDescriptorWrites(); +} + +void MVKCmdPushDescriptorSet::clearDescriptorWrites() { + for (VkWriteDescriptorSet &descWrite : _descriptorWrites) { + if (descWrite.pImageInfo) delete[] descWrite.pImageInfo; + if (descWrite.pBufferInfo) delete[] descWrite.pBufferInfo; + if (descWrite.pTexelBufferView) delete[] descWrite.pTexelBufferView; + } + _descriptorWrites.clear(); +} + + #pragma mark - #pragma mark Command creation functions @@ -232,3 +290,13 @@ void mvkCmdPushConstants(MVKCommandBuffer* cmdBuff, cmdBuff->addCommand(cmd); } +void mvkCmdPushDescriptorSet(MVKCommandBuffer* cmdBuff, + VkPipelineBindPoint pipelineBindPoint, + VkPipelineLayout layout, + uint32_t set, + uint32_t descriptorWriteCount, + const VkWriteDescriptorSet* pDescriptorWrites) { + MVKCmdPushDescriptorSet* cmd = cmdBuff->_commandPool->_cmdPushDescriptorSetPool.acquireObject(); + cmd->setContent(pipelineBindPoint, layout, set, descriptorWriteCount, pDescriptorWrites); + cmdBuff->addCommand(cmd); +} diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h index 2037bee7..52db2654 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h @@ -131,6 +131,8 @@ public: MVKCommandTypePool _cmdDispatchIndirectPool; + MVKCommandTypePool _cmdPushDescriptorSetPool; + #pragma mark Command resources diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm index bf21b894..e12ba783 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm @@ -118,7 +118,8 @@ MVKCommandPool::MVKCommandPool(MVKDevice* device, _cmdCopyQueryPoolResultsPool(this, true), _cmdPushConstantsPool(this, true), _cmdDispatchPool(this, true), - _cmdDispatchIndirectPool(this, true) + _cmdDispatchIndirectPool(this, true), + _cmdPushDescriptorSetPool(this, true) {} // TODO: Destroying a command pool implicitly destroys all command buffers and commands created from it. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h index d03ca573..4e8a4ccf 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h @@ -78,6 +78,16 @@ public: std::vector& dynamicOffsets, uint32_t* pDynamicOffsetIndex); + /** Encodes this binding layout and the specified descriptor binding on the specified command encoder immediately. */ + void push(MVKCommandEncoder* cmdEncoder, + uint32_t& dstArrayElement, + uint32_t& descriptorCount, + VkDescriptorType descriptorType, + const VkDescriptorImageInfo*& pImageInfo, + const VkDescriptorBufferInfo*& pBufferInfo, + const VkBufferView*& pTexelBufferView, + MVKShaderResourceBinding& dslMTLRezIdxOffsets); + /** Populates the specified shader converter context, at the specified descriptor set binding. */ void populateShaderConverterContext(SPIRVToMSLConverterContext& context, MVKShaderResourceBinding& dslMTLRezIdxOffsets, @@ -119,11 +129,20 @@ public: uint32_t* pDynamicOffsetIndex); + /** Encodes this descriptor set layout and the specified descriptor updates on the specified command encoder immediately. */ + void pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + std::vector& descriptorWrites, + MVKShaderResourceBinding& dslMTLRezIdxOffsets); + + /** Populates the specified shader converter context, at the specified DSL index. */ void populateShaderConverterContext(SPIRVToMSLConverterContext& context, MVKShaderResourceBinding& dslMTLRezIdxOffsets, uint32_t dslIndex); + /** Returns true if this layout is for push descriptors only. */ + bool isPushDescriptorLayout() const { return _isPushDescriptorLayout; } + /** Constructs an instance for the specified device. */ MVKDescriptorSetLayout(MVKDevice* device, const VkDescriptorSetLayoutCreateInfo* pCreateInfo); @@ -135,6 +154,7 @@ protected: std::vector _bindings; MVKShaderResourceBinding _mtlResourceCounts; + bool _isPushDescriptorLayout : 1; }; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm index 9b514f7d..d85e29ca 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm @@ -173,6 +173,172 @@ void MVKDescriptorSetLayoutBinding::bind(MVKCommandEncoder* cmdEncoder, } } +void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, + uint32_t& dstArrayElement, + uint32_t& descriptorCount, + VkDescriptorType descriptorType, + const VkDescriptorImageInfo*& pImageInfo, + const VkDescriptorBufferInfo*& pBufferInfo, + const VkBufferView*& pTexelBufferView, + MVKShaderResourceBinding& dslMTLRezIdxOffsets) { + MVKMTLBufferBinding bb; + MVKMTLTextureBinding tb; + MVKMTLSamplerStateBinding sb; + + if (dstArrayElement >= _info.descriptorCount) { + dstArrayElement -= _info.descriptorCount; + return; + } + + if (descriptorType != _info.descriptorType) { + dstArrayElement = 0; + if (_info.descriptorCount > descriptorCount) + descriptorCount = 0; + else { + descriptorCount -= _info.descriptorCount; + pImageInfo += _info.descriptorCount; + pBufferInfo += _info.descriptorCount; + pTexelBufferView += _info.descriptorCount; + } + return; + } + + // Establish the resource indices to use, by combining the offsets of the DSL and this DSL binding. + MVKShaderResourceBinding mtlIdxs = _mtlResourceIndexOffsets + dslMTLRezIdxOffsets; + + for (uint32_t rezIdx = dstArrayElement; + rezIdx < _info.descriptorCount && rezIdx - dstArrayElement < descriptorCount; + rezIdx++) { + switch (_info.descriptorType) { + + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: { + const VkDescriptorBufferInfo& bufferInfo = pBufferInfo[rezIdx - dstArrayElement]; + MVKBuffer* buffer = (MVKBuffer*)bufferInfo.buffer; + bb.mtlBuffer = buffer->getMTLBuffer(); + bb.offset = bufferInfo.offset; + if (_applyToVertexStage) { + bb.index = mtlIdxs.vertexStage.bufferIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexBuffer(bb); + } + if (_applyToFragmentStage) { + bb.index = mtlIdxs.fragmentStage.bufferIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentBuffer(bb); + } + if (_applyToComputeStage) { + bb.index = mtlIdxs.computeStage.bufferIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindBuffer(bb); + } + break; + } + + case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: + case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: + case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: { + const VkDescriptorImageInfo& imageInfo = pImageInfo[rezIdx - dstArrayElement]; + MVKImageView* imageView = (MVKImageView*)imageInfo.imageView; + tb.mtlTexture = imageView->getMTLTexture(); + if (_applyToVertexStage) { + tb.index = mtlIdxs.vertexStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexTexture(tb); + } + if (_applyToFragmentStage) { + tb.index = mtlIdxs.fragmentStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentTexture(tb); + } + if (_applyToComputeStage) { + tb.index = mtlIdxs.computeStage.textureIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindTexture(tb); + } + break; + } + + case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: { + MVKBufferView* bufferView = (MVKBufferView*)pTexelBufferView[rezIdx - dstArrayElement]; + tb.mtlTexture = bufferView->getMTLTexture(); + if (_applyToVertexStage) { + tb.index = mtlIdxs.vertexStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexTexture(tb); + } + if (_applyToFragmentStage) { + tb.index = mtlIdxs.fragmentStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentTexture(tb); + } + if (_applyToComputeStage) { + tb.index = mtlIdxs.computeStage.textureIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindTexture(tb); + } + break; + } + + case VK_DESCRIPTOR_TYPE_SAMPLER: { + MVKSampler* sampler; + if (_immutableSamplers.empty()) + sampler = (MVKSampler*)pImageInfo[rezIdx - dstArrayElement].sampler; + else + sampler = _immutableSamplers[rezIdx]; + sb.mtlSamplerState = sampler->getMTLSamplerState(); + if (_applyToVertexStage) { + sb.index = mtlIdxs.vertexStage.samplerIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexSamplerState(sb); + } + if (_applyToFragmentStage) { + sb.index = mtlIdxs.fragmentStage.samplerIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentSamplerState(sb); + } + if (_applyToComputeStage) { + sb.index = mtlIdxs.computeStage.samplerIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindSamplerState(sb); + } + break; + } + + case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: { + const VkDescriptorImageInfo& imageInfo = pImageInfo[rezIdx - dstArrayElement]; + MVKImageView* imageView = (MVKImageView*)imageInfo.imageView; + MVKSampler* sampler = _immutableSamplers.empty() ? (MVKSampler*)imageInfo.sampler : _immutableSamplers[rezIdx]; + tb.mtlTexture = imageView->getMTLTexture(); + sb.mtlSamplerState = sampler->getMTLSamplerState(); + if (_applyToVertexStage) { + tb.index = mtlIdxs.vertexStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexTexture(tb); + sb.index = mtlIdxs.vertexStage.samplerIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindVertexSamplerState(sb); + } + if (_applyToFragmentStage) { + tb.index = mtlIdxs.fragmentStage.textureIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentTexture(tb); + sb.index = mtlIdxs.fragmentStage.samplerIndex + rezIdx; + cmdEncoder->_graphicsResourcesState.bindFragmentSamplerState(sb); + } + if (_applyToComputeStage) { + tb.index = mtlIdxs.computeStage.textureIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindTexture(tb); + sb.index = mtlIdxs.computeStage.samplerIndex + rezIdx; + cmdEncoder->_computeResourcesState.bindSamplerState(sb); + } + break; + } + + default: + break; + } + } + + dstArrayElement = 0; + if (_info.descriptorCount > descriptorCount) + descriptorCount = 0; + else { + descriptorCount -= _info.descriptorCount; + pImageInfo += _info.descriptorCount; + pBufferInfo += _info.descriptorCount; + pTexelBufferView += _info.descriptorCount; + } +} + void MVKDescriptorSetLayoutBinding::populateShaderConverterContext(SPIRVToMSLConverterContext& context, MVKShaderResourceBinding& dslMTLRezIdxOffsets, uint32_t dslIndex) { @@ -300,6 +466,7 @@ void MVKDescriptorSetLayout::bindDescriptorSet(MVKCommandEncoder* cmdEncoder, vector& dynamicOffsets, uint32_t* pDynamicOffsetIndex) { + if (_isPushDescriptorLayout) return; uint32_t bindCnt = (uint32_t)_bindings.size(); for (uint32_t bindIdx = 0; bindIdx < bindCnt; bindIdx++) { _bindings[bindIdx].bind(cmdEncoder, descSet->_bindings[bindIdx], @@ -308,6 +475,28 @@ void MVKDescriptorSetLayout::bindDescriptorSet(MVKCommandEncoder* cmdEncoder, } } +void MVKDescriptorSetLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + vector& descriptorWrites, + MVKShaderResourceBinding& dslMTLRezIdxOffsets) { + + if (!_isPushDescriptorLayout) return; + for (const VkWriteDescriptorSet& descWrite : descriptorWrites) { + uint32_t bindIdx = descWrite.dstBinding; + uint32_t dstArrayElement = descWrite.dstArrayElement; + uint32_t descriptorCount = descWrite.descriptorCount; + const VkDescriptorImageInfo* pImageInfo = descWrite.pImageInfo; + const VkDescriptorBufferInfo* pBufferInfo = descWrite.pBufferInfo; + const VkBufferView* pTexelBufferView = descWrite.pTexelBufferView; + // Note: This will result in us walking off the end of the array + // in case there are too many updates... but that's ill-defined anyway. + for (; descriptorCount; bindIdx++) { + _bindings[bindIdx].push(cmdEncoder, dstArrayElement, descriptorCount, + descWrite.descriptorType, pImageInfo, pBufferInfo, + pTexelBufferView, dslMTLRezIdxOffsets); + } + } +} + void MVKDescriptorSetLayout::populateShaderConverterContext(SPIRVToMSLConverterContext& context, MVKShaderResourceBinding& dslMTLRezIdxOffsets, uint32_t dslIndex) { @@ -319,6 +508,7 @@ void MVKDescriptorSetLayout::populateShaderConverterContext(SPIRVToMSLConverterC MVKDescriptorSetLayout::MVKDescriptorSetLayout(MVKDevice* device, const VkDescriptorSetLayoutCreateInfo* pCreateInfo) : MVKBaseDeviceObject(device) { + _isPushDescriptorLayout = (pCreateInfo->flags & VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR) != 0; // Create the descriptor bindings _bindings.reserve(pCreateInfo->bindingCount); for (uint32_t i = 0; i < pCreateInfo->bindingCount; i++) { @@ -616,6 +806,7 @@ VkResult MVKDescriptorPool::allocateDescriptorSets(uint32_t count, for (uint32_t dsIdx = 0; dsIdx < count; dsIdx++) { MVKDescriptorSetLayout* mvkDSL = (MVKDescriptorSetLayout*)pSetLayouts[dsIdx]; + if (mvkDSL->isPushDescriptorLayout()) continue; MVKDescriptorSet* mvkDescSet = new MVKDescriptorSet(_device, mvkDSL); _allocatedSets.push_front(mvkDescSet); pDescriptorSets[dsIdx] = (VkDescriptorSet)mvkDescSet; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 89411c40..ee5889c0 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -77,6 +77,20 @@ void MVKPhysicalDevice::getProperties(VkPhysicalDeviceProperties2KHR* properties if (properties) { properties->sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR; properties->properties = _properties; + auto* next = (VkStructureType*)properties->pNext; + while (next) { + switch (*next) { + case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PUSH_DESCRIPTOR_PROPERTIES_KHR: { + auto* pushDescProps = (VkPhysicalDevicePushDescriptorPropertiesKHR*)next; + pushDescProps->maxPushDescriptors = _properties.limits.maxPerStageResources; + next = (VkStructureType*)pushDescProps->pNext; + break; + } + default: + next = *(VkStructureType**)(next+1); + break; + } + } } } diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index df5b9ff9..61df8edc 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -249,6 +249,7 @@ void MVKInstance::initProcAddrs() { ADD_PROC_ADDR(vkGetPhysicalDeviceQueueFamilyProperties2KHR); ADD_PROC_ADDR(vkGetPhysicalDeviceMemoryProperties2KHR); ADD_PROC_ADDR(vkGetPhysicalDeviceSparseImageFormatProperties2KHR); + ADD_PROC_ADDR(vkCmdPushDescriptorSetKHR); ADD_PROC_ADDR(vkGetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkSetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkGetPhysicalDeviceMetalFeaturesMVK); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h index 3772693b..4cd6391c 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h @@ -50,6 +50,11 @@ public: /** Populates the specified shader converter context. */ void populateShaderConverterContext(SPIRVToMSLConverterContext& context); + /** Updates a descriptor set in a command encoder. */ + void pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + std::vector& descriptorWrites, + uint32_t set); + /** Constructs an instance for the specified device. */ MVKPipelineLayout(MVKDevice* device, const VkPipelineLayoutCreateInfo* pCreateInfo); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm index 2efa06f9..151b840d 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm @@ -54,6 +54,17 @@ void MVKPipelineLayout::bindDescriptorSets(MVKCommandEncoder* cmdEncoder, cmdEncoder->getPushConstants(VK_SHADER_STAGE_COMPUTE_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.computeStage.bufferIndex); } +void MVKPipelineLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + vector& descriptorWrites, + uint32_t set) { + + _descriptorSetLayouts[set].pushDescriptorSet(cmdEncoder, descriptorWrites, + _dslMTLResourceIndexOffsets[set]); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.vertexStage.bufferIndex); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_FRAGMENT_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.fragmentStage.bufferIndex); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_COMPUTE_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.computeStage.bufferIndex); +} + void MVKPipelineLayout::populateShaderConverterContext(SPIRVToMSLConverterContext& context) { context.resourceBindings.clear(); diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.mm b/MoltenVK/MoltenVK/Loader/MVKLayers.mm index bc3b9e8b..42f331ea 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Loader/MVKLayers.mm @@ -107,6 +107,11 @@ MVKLayer::MVKLayer() { extTmplt.specVersion = VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_SPEC_VERSION; _extensions.push_back(extTmplt); + memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); + strcpy(extTmplt.extensionName, VK_KHR_PUSH_DESCRIPTOR_EXTENSION_NAME); + extTmplt.specVersion = VK_KHR_PUSH_DESCRIPTOR_SPEC_VERSION; + _extensions.push_back(extTmplt); + #if MVK_IOS memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); strcpy(extTmplt.extensionName, VK_MVK_IOS_SURFACE_EXTENSION_NAME); diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index a9434f3a..db0e6f3a 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -1677,6 +1677,22 @@ MVK_PUBLIC_SYMBOL void vkGetPhysicalDeviceSparseImageFormatProperties2KHR( } +#pragma mark - +#pragma mark VK_KHR_push_descriptor extension + +MVK_PUBLIC_SYMBOL void vkCmdPushDescriptorSetKHR( + VkCommandBuffer commandBuffer, + VkPipelineBindPoint pipelineBindPoint, + VkPipelineLayout layout, + uint32_t set, + uint32_t descriptorWriteCount, + const VkWriteDescriptorSet* pDescriptorWrites) { + + MVKCommandBuffer* cmdBuff = MVKCommandBuffer::getMVKCommandBuffer(commandBuffer); + mvkCmdPushDescriptorSet(cmdBuff, pipelineBindPoint, layout, set, descriptorWriteCount, pDescriptorWrites); +} + + #pragma mark - #pragma mark Loader and Layer ICD interface extension From 63ce4cf1f5773639d4182a9493a7f51afd22ff47 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Sat, 1 Sep 2018 16:42:10 -0400 Subject: [PATCH 05/20] Vulkan extension enhancements. Add MVKExtensionList struct to track supported and enabled Vulkan extensions within MVKInstance & MVKDevice. Log supported and enabled Vulkan extensions for VkInstance & VkDevice. Add error handling for descriptor sets & 3D images in VK_KHR_maintenance1 extension. All header files load mvk_vulkan.h instead of vulkan.h for consistent use of macOS & iOS extensions. --- MoltenVK/MoltenVK.xcodeproj/project.pbxproj | 30 ++- MoltenVK/MoltenVK/API/mvk_datatypes.h | 2 +- MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h | 2 +- MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm | 3 + .../MoltenVK/GPUObjects/MVKDescriptorSet.mm | 6 +- MoltenVK/MoltenVK/GPUObjects/MVKDevice.h | 7 +- MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm | 62 ++++--- MoltenVK/MoltenVK/GPUObjects/MVKImage.h | 1 + MoltenVK/MoltenVK/GPUObjects/MVKImage.mm | 13 ++ MoltenVK/MoltenVK/GPUObjects/MVKInstance.h | 26 +-- MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm | 175 +++++++++--------- MoltenVK/MoltenVK/Layers/MVKExtensions.cpp | 149 +++++++++++++++ MoltenVK/MoltenVK/Layers/MVKExtensions.h | 82 ++++++++ .../MoltenVK/{Loader => Layers}/MVKLayers.h | 15 +- .../MoltenVK/{Loader => Layers}/MVKLayers.mm | 94 +++------- MoltenVK/MoltenVK/OS/MVKOSExtensions.h | 2 +- MoltenVK/MoltenVK/Utility/MVKBaseObject.h | 2 +- .../{MVKFoundation.mm => MVKFoundation.cpp} | 6 +- MoltenVK/MoltenVK/Utility/MVKFoundation.h | 4 +- MoltenVK/MoltenVK/Vulkan/vk_mvk_moltenvk.mm | 8 +- 20 files changed, 455 insertions(+), 234 deletions(-) create mode 100644 MoltenVK/MoltenVK/Layers/MVKExtensions.cpp create mode 100644 MoltenVK/MoltenVK/Layers/MVKExtensions.h rename MoltenVK/MoltenVK/{Loader => Layers}/MVKLayers.h (90%) rename MoltenVK/MoltenVK/{Loader => Layers}/MVKLayers.mm (51%) rename MoltenVK/MoltenVK/Utility/{MVKFoundation.mm => MVKFoundation.cpp} (95%) diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj index c41a0dda..2f87edd4 100644 --- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj +++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj @@ -9,6 +9,10 @@ /* Begin PBXBuildFile section */ A9096E5E1F81E16300DFBEA6 /* MVKCmdDispatch.mm in Sources */ = {isa = PBXBuildFile; fileRef = A9096E5D1F81E16300DFBEA6 /* MVKCmdDispatch.mm */; }; A9096E5F1F81E16300DFBEA6 /* MVKCmdDispatch.mm in Sources */ = {isa = PBXBuildFile; fileRef = A9096E5D1F81E16300DFBEA6 /* MVKCmdDispatch.mm */; }; + A909F65F213B190700FCD6BE /* MVKExtensions.h in Headers */ = {isa = PBXBuildFile; fileRef = A909F65A213B190600FCD6BE /* MVKExtensions.h */; }; + A909F660213B190700FCD6BE /* MVKExtensions.h in Headers */ = {isa = PBXBuildFile; fileRef = A909F65A213B190600FCD6BE /* MVKExtensions.h */; }; + A909F661213B190700FCD6BE /* MVKExtensions.cpp in Sources */ = {isa = PBXBuildFile; fileRef = A909F65E213B190700FCD6BE /* MVKExtensions.cpp */; }; + A909F662213B190700FCD6BE /* MVKExtensions.cpp in Sources */ = {isa = PBXBuildFile; fileRef = A909F65E213B190700FCD6BE /* MVKExtensions.cpp */; }; A90C8DEA1F45354D009CB32C /* MVKCommandEncodingPool.h in Headers */ = {isa = PBXBuildFile; fileRef = A90C8DE81F45354D009CB32C /* MVKCommandEncodingPool.h */; }; A90C8DEB1F45354D009CB32C /* MVKCommandEncodingPool.h in Headers */ = {isa = PBXBuildFile; fileRef = A90C8DE81F45354D009CB32C /* MVKCommandEncodingPool.h */; }; A90C8DEC1F45354D009CB32C /* MVKCommandEncodingPool.mm in Sources */ = {isa = PBXBuildFile; fileRef = A90C8DE91F45354D009CB32C /* MVKCommandEncodingPool.mm */; }; @@ -147,8 +151,8 @@ A98149521FB6A3F7005F00B4 /* MVKEnvironment.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149431FB6A3F7005F00B4 /* MVKEnvironment.h */; }; A98149531FB6A3F7005F00B4 /* MVKFoundation.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149441FB6A3F7005F00B4 /* MVKFoundation.h */; }; A98149541FB6A3F7005F00B4 /* MVKFoundation.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149441FB6A3F7005F00B4 /* MVKFoundation.h */; }; - A98149551FB6A3F7005F00B4 /* MVKFoundation.mm in Sources */ = {isa = PBXBuildFile; fileRef = A98149451FB6A3F7005F00B4 /* MVKFoundation.mm */; }; - A98149561FB6A3F7005F00B4 /* MVKFoundation.mm in Sources */ = {isa = PBXBuildFile; fileRef = A98149451FB6A3F7005F00B4 /* MVKFoundation.mm */; }; + A98149551FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */ = {isa = PBXBuildFile; fileRef = A98149451FB6A3F7005F00B4 /* MVKFoundation.cpp */; }; + A98149561FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */ = {isa = PBXBuildFile; fileRef = A98149451FB6A3F7005F00B4 /* MVKFoundation.cpp */; }; A98149571FB6A3F7005F00B4 /* MVKObjectPool.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */; }; A98149581FB6A3F7005F00B4 /* MVKObjectPool.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */; }; A981495D1FB6A3F7005F00B4 /* MVKWatermark.h in Headers */ = {isa = PBXBuildFile; fileRef = A98149491FB6A3F7005F00B4 /* MVKWatermark.h */; }; @@ -256,6 +260,8 @@ /* Begin PBXFileReference section */ A9096E5C1F81E16300DFBEA6 /* MVKCmdDispatch.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = MVKCmdDispatch.h; sourceTree = ""; }; A9096E5D1F81E16300DFBEA6 /* MVKCmdDispatch.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCmdDispatch.mm; sourceTree = ""; }; + A909F65A213B190600FCD6BE /* MVKExtensions.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKExtensions.h; sourceTree = ""; }; + A909F65E213B190700FCD6BE /* MVKExtensions.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKExtensions.cpp; sourceTree = ""; }; A90C8DE81F45354D009CB32C /* MVKCommandEncodingPool.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCommandEncodingPool.h; sourceTree = ""; }; A90C8DE91F45354D009CB32C /* MVKCommandEncodingPool.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKCommandEncodingPool.mm; sourceTree = ""; }; A93E832E2121C5D3001FEBD4 /* MVKGPUCapture.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKGPUCapture.h; sourceTree = ""; }; @@ -325,7 +331,7 @@ A98149421FB6A3F7005F00B4 /* MVKBaseObject.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKBaseObject.h; sourceTree = ""; }; A98149431FB6A3F7005F00B4 /* MVKEnvironment.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKEnvironment.h; sourceTree = ""; }; A98149441FB6A3F7005F00B4 /* MVKFoundation.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKFoundation.h; sourceTree = ""; }; - A98149451FB6A3F7005F00B4 /* MVKFoundation.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKFoundation.mm; sourceTree = ""; }; + A98149451FB6A3F7005F00B4 /* MVKFoundation.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MVKFoundation.cpp; sourceTree = ""; }; A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKObjectPool.h; sourceTree = ""; }; A98149491FB6A3F7005F00B4 /* MVKWatermark.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKWatermark.h; sourceTree = ""; }; A981494A1FB6A3F7005F00B4 /* MVKWatermark.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = MVKWatermark.mm; sourceTree = ""; }; @@ -365,7 +371,7 @@ A94FB7651C7DFB4800632CA3 /* API */, A94FB76B1C7DFB4800632CA3 /* Commands */, A94FB77E1C7DFB4800632CA3 /* GPUObjects */, - A94FB79F1C7DFB4800632CA3 /* Loader */, + A94FB79F1C7DFB4800632CA3 /* Layers */, A9E53DCC2100B197002781DD /* OS */, A98149401FB6A3F7005F00B4 /* Utility */, A94FB7A81C7DFB4800632CA3 /* Vulkan */, @@ -457,13 +463,15 @@ path = GPUObjects; sourceTree = ""; }; - A94FB79F1C7DFB4800632CA3 /* Loader */ = { + A94FB79F1C7DFB4800632CA3 /* Layers */ = { isa = PBXGroup; children = ( + A909F65A213B190600FCD6BE /* MVKExtensions.h */, + A909F65E213B190700FCD6BE /* MVKExtensions.cpp */, A94FB7A01C7DFB4800632CA3 /* MVKLayers.h */, A94FB7A11C7DFB4800632CA3 /* MVKLayers.mm */, ); - path = Loader; + path = Layers; sourceTree = ""; }; A94FB7A81C7DFB4800632CA3 /* Vulkan */ = { @@ -483,7 +491,7 @@ A98149421FB6A3F7005F00B4 /* MVKBaseObject.h */, A98149431FB6A3F7005F00B4 /* MVKEnvironment.h */, A98149441FB6A3F7005F00B4 /* MVKFoundation.h */, - A98149451FB6A3F7005F00B4 /* MVKFoundation.mm */, + A98149451FB6A3F7005F00B4 /* MVKFoundation.cpp */, A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */, A98149491FB6A3F7005F00B4 /* MVKWatermark.h */, A981494A1FB6A3F7005F00B4 /* MVKWatermark.mm */, @@ -576,6 +584,7 @@ buildActionMask = 2147483647; files = ( A94FB7B41C7DFB4800632CA3 /* vk_mvk_moltenvk.h in Headers */, + A909F65F213B190700FCD6BE /* MVKExtensions.h in Headers */, A94FB7B01C7DFB4800632CA3 /* mvk_datatypes.h in Headers */, A98149511FB6A3F7005F00B4 /* MVKEnvironment.h in Headers */, A948BB7F1E51642700DE59F2 /* mvk_vulkan.h in Headers */, @@ -635,6 +644,7 @@ buildActionMask = 2147483647; files = ( A94FB7B51C7DFB4800632CA3 /* vk_mvk_moltenvk.h in Headers */, + A909F660213B190700FCD6BE /* MVKExtensions.h in Headers */, A94FB7B11C7DFB4800632CA3 /* mvk_datatypes.h in Headers */, A98149521FB6A3F7005F00B4 /* MVKEnvironment.h in Headers */, A948BB801E51642700DE59F2 /* mvk_vulkan.h in Headers */, @@ -866,7 +876,8 @@ A94FB7C61C7DFB4800632CA3 /* MVKCmdRenderPass.mm in Sources */, A94FB7DE1C7DFB4800632CA3 /* MVKBuffer.mm in Sources */, A94FB82A1C7DFB4800632CA3 /* mvk_datatypes.mm in Sources */, - A98149551FB6A3F7005F00B4 /* MVKFoundation.mm in Sources */, + A909F661213B190700FCD6BE /* MVKExtensions.cpp in Sources */, + A98149551FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */, A94FB7E61C7DFB4800632CA3 /* MVKDevice.mm in Sources */, A9E53DF52100B302002781DD /* MTLRenderPassDescriptor+MoltenVK.m in Sources */, A94FB7FA1C7DFB4800632CA3 /* MVKPipeline.mm in Sources */, @@ -917,7 +928,8 @@ A94FB7C71C7DFB4800632CA3 /* MVKCmdRenderPass.mm in Sources */, A94FB7DF1C7DFB4800632CA3 /* MVKBuffer.mm in Sources */, A94FB82B1C7DFB4800632CA3 /* mvk_datatypes.mm in Sources */, - A98149561FB6A3F7005F00B4 /* MVKFoundation.mm in Sources */, + A909F662213B190700FCD6BE /* MVKExtensions.cpp in Sources */, + A98149561FB6A3F7005F00B4 /* MVKFoundation.cpp in Sources */, A94FB7E71C7DFB4800632CA3 /* MVKDevice.mm in Sources */, A9E53DF62100B302002781DD /* MTLRenderPassDescriptor+MoltenVK.m in Sources */, A94FB7FB1C7DFB4800632CA3 /* MVKPipeline.mm in Sources */, diff --git a/MoltenVK/MoltenVK/API/mvk_datatypes.h b/MoltenVK/MoltenVK/API/mvk_datatypes.h index c02676bc..e95f1c6f 100644 --- a/MoltenVK/MoltenVK/API/mvk_datatypes.h +++ b/MoltenVK/MoltenVK/API/mvk_datatypes.h @@ -31,7 +31,7 @@ extern "C" { #endif // __cplusplus -#include +#include #import #import diff --git a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h index 427a3b99..14067b36 100644 --- a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h +++ b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h @@ -26,7 +26,7 @@ extern "C" { #endif // __cplusplus -#include +#include #ifdef __OBJC__ #import diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm index b28c2cd3..b54fc064 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm @@ -62,6 +62,9 @@ void MVKCmdCopyImage::setContent(VkImage srcImage, if (_srcImage->getMTLPixelFormat() != _dstImage->getMTLPixelFormat()) { setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdCopyImage(): The source and destination images must have the same format.")); } + if ((_srcImage->getMTLTextureType() == MTLTextureType3D) || (_dstImage->getMTLTextureType() == MTLTextureType3D)) { + setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdCopyImage(): Metal does not support copying to or from slices of a 3D texture.")); + } } // Adds a Metal copy region structure for each layer in the specified copy region. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm index 9b514f7d..db5013ad 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm @@ -611,7 +611,11 @@ VkResult MVKDescriptorPool::allocateDescriptorSets(uint32_t count, const VkDescriptorSetLayout* pSetLayouts, VkDescriptorSet* pDescriptorSets) { if (_allocatedSetCount + count > _maxSets) { - return mvkNotifyErrorWithText(VK_ERROR_OUT_OF_POOL_MEMORY_KHR, "The maximum number of descriptor sets that can be allocated by this descriptor pool is %d.", _maxSets); + if (_device->_enabledExtensions.vk_KHR_maintenance1.enabled) { + return VK_ERROR_OUT_OF_POOL_MEMORY; // Failure is an acceptable test...don't log as error. + } else { + return mvkNotifyErrorWithText(VK_ERROR_INITIALIZATION_FAILED, "The maximum number of descriptor sets that can be allocated by this descriptor pool is %d.", _maxSets); + } } for (uint32_t dsIdx = 0; dsIdx < count; dsIdx++) { diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index 6ca8543d..7fa5fac4 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -255,7 +255,7 @@ class MVKDevice : public MVKDispatchableObject { public: /** Returns a pointer to the Vulkan instance. */ - inline MVKInstance* getInstance() { return _physicalDevice->getInstance(); } + inline MVKInstance* getInstance() { return _physicalDevice->_mvkInstance; } /** Returns the physical device underlying this logical device. */ inline MVKPhysicalDevice* getPhysicalDevice() { return _physicalDevice; } @@ -471,7 +471,7 @@ public: #pragma mark Properties directly accessible - /** The MoltenVK configuration settings. */ + /** Pointer to the MoltenVK configuration settings. */ const MVKConfiguration* _pMVKConfig; /** Pointer to the feature set of the underlying physical device. */ @@ -486,6 +486,9 @@ public: /** Pointer to the memory properties of the underlying physical device. */ const VkPhysicalDeviceMemoryProperties* _pMemoryProperties; + /** The list of Vulkan extensions, indicating whether each has been enabled by the app for this device. */ + const MVKExtensionList _enabledExtensions; + /** Performance statistics. */ MVKPerformanceStatistics _performanceStatistics; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 19bde85f..1f19e0fe 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -929,39 +929,39 @@ void MVKPhysicalDevice::logGPUInfo() { break; } - string fsMsg = "GPU device:"; - fsMsg += "\n\t\tmodel: %s"; - fsMsg += "\n\t\ttype: %s"; - fsMsg += "\n\t\tvendorID: %#06x"; - fsMsg += "\n\t\tdeviceID: %#06x"; - fsMsg += "\n\t\tpipelineCacheUUID: %s"; - fsMsg += "\n\tsupports the following Metal Feature Sets:"; + string logMsg = "GPU device:"; + logMsg += "\n\t\tmodel: %s"; + logMsg += "\n\t\ttype: %s"; + logMsg += "\n\t\tvendorID: %#06x"; + logMsg += "\n\t\tdeviceID: %#06x"; + logMsg += "\n\t\tpipelineCacheUUID: %s"; + logMsg += "\n\tsupports the following Metal Feature Sets:"; #if MVK_IOS - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily4_v1] ) { fsMsg += "\n\tviOS GPU Family 4 v1"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily4_v1] ) { logMsg += "\n\t\tiOS GPU Family 4 v1"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v3] ) { fsMsg += "\n\t\tiOS GPU Family 3 v3"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v2] ) { fsMsg += "\n\t\tiOS GPU Family 3 v2"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v1] ) { fsMsg += "\n\t\tiOS GPU Family 3 v1"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v3] ) { logMsg += "\n\t\tiOS GPU Family 3 v3"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v2] ) { logMsg += "\n\t\tiOS GPU Family 3 v2"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily3_v1] ) { logMsg += "\n\t\tiOS GPU Family 3 v1"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v4] ) { fsMsg += "\n\t\tiOS GPU Family 2 v4"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v3] ) { fsMsg += "\n\t\tiOS GPU Family 2 v3"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v2] ) { fsMsg += "\n\t\tiOS GPU Family 2 v2"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v1] ) { fsMsg += "\n\t\tiOS GPU Family 2 v1"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v4] ) { logMsg += "\n\t\tiOS GPU Family 2 v4"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v3] ) { logMsg += "\n\t\tiOS GPU Family 2 v3"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v2] ) { logMsg += "\n\t\tiOS GPU Family 2 v2"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily2_v1] ) { logMsg += "\n\t\tiOS GPU Family 2 v1"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v4] ) { fsMsg += "\n\t\tiOS GPU Family 1 v4"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v3] ) { fsMsg += "\n\t\tiOS GPU Family 1 v3"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v2] ) { fsMsg += "\n\t\tiOS GPU Family 1 v2"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v1] ) { fsMsg += "\n\t\tiOS GPU Family 1 v1"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v4] ) { logMsg += "\n\t\tiOS GPU Family 1 v4"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v3] ) { logMsg += "\n\t\tiOS GPU Family 1 v3"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v2] ) { logMsg += "\n\t\tiOS GPU Family 1 v2"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v1] ) { logMsg += "\n\t\tiOS GPU Family 1 v1"; } #endif #if MVK_MACOS - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v3] ) { fsMsg += "\n\t\tmacOS GPU Family 1 v3"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v2] ) { fsMsg += "\n\t\tmacOS GPU Family 1 v2"; } - if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v1] ) { fsMsg += "\n\t\tmacOS GPU Family 1 v1"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v3] ) { logMsg += "\n\t\tmacOS GPU Family 1 v3"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v2] ) { logMsg += "\n\t\tmacOS GPU Family 1 v2"; } + if ( [_mtlDevice supportsFeatureSet: MTLFeatureSet_macOS_GPUFamily1_v1] ) { logMsg += "\n\t\tmacOS GPU Family 1 v1"; } #endif - MVKLogInfo(fsMsg.c_str(), _properties.deviceName, devTypeStr.c_str(), _properties.vendorID, _properties.deviceID, + MVKLogInfo(logMsg.c_str(), _properties.deviceName, devTypeStr.c_str(), _properties.vendorID, _properties.deviceID, [[[NSUUID alloc] initWithUUIDBytes: _properties.pipelineCacheUUID] autorelease].UUIDString.UTF8String); } @@ -1434,7 +1434,7 @@ MVKDevice::MVKDevice(MVKPhysicalDevice* physicalDevice, const VkDeviceCreateInfo initPerformanceTracking(); _physicalDevice = physicalDevice; - _pMVKConfig = &_physicalDevice->getInstance()->_mvkConfig; + _pMVKConfig = _physicalDevice->_mvkInstance->getMoltenVKConfiguration(); _pFeatures = &_physicalDevice->_features; _pMetalFeatures = &_physicalDevice->_metalFeatures; _pProperties = &_physicalDevice->_properties; @@ -1443,11 +1443,13 @@ MVKDevice::MVKDevice(MVKPhysicalDevice* physicalDevice, const VkDeviceCreateInfo _globalVisibilityResultMTLBuffer = nil; _globalVisibilityQueryCount = 0; - // Verify the requested extension names. Should be same as those requested from instance. - setConfigurationResult(_physicalDevice->_mvkInstance->verifyExtensions(pCreateInfo->enabledExtensionCount, - pCreateInfo->ppEnabledExtensionNames)); + _commandResourceFactory = new MVKCommandResourceFactory(this); - _commandResourceFactory = new MVKCommandResourceFactory(this); + // Verify the requested extension names. Should be same as those requested from instance. + MVKExtensionList* pWritableExtns = (MVKExtensionList*)&_enabledExtensions; + setConfigurationResult(pWritableExtns->enable(pCreateInfo->enabledExtensionCount, + pCreateInfo->ppEnabledExtensionNames, + getInstance()->getDriverLayer()->getSupportedExtensions())); // Create the queues uint32_t qrCnt = pCreateInfo->queueCreateInfoCount; @@ -1462,7 +1464,9 @@ MVKDevice::MVKDevice(MVKPhysicalDevice* physicalDevice, const VkDeviceCreateInfo } } - MVKLogInfo("Created VkDevice to run on GPU %s", _pProperties->deviceName); + string logMsg = "Created VkDevice to run on GPU %s with the following Vulkan extensions enabled:"; + logMsg += _enabledExtensions.enabledNamesString("\n\t\t", true); + MVKLogInfo(logMsg.c_str(), _pProperties->deviceName); } void MVKDevice::initPerformanceTracking() { diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h index 087b5f2e..48e0e536 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h @@ -281,6 +281,7 @@ protected: bool matchesSwizzle(VkComponentMapping components, VkComponentMapping pattern); const char* getSwizzleName(VkComponentSwizzle swizzle); void setSwizzleFormatError(VkFormat format, VkComponentMapping components); + void validateImageViewConfig(const VkImageViewCreateInfo* pCreateInfo); MVKImage* _image; VkImageSubresourceRange _subresourceRange; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index 6d27cfd3..6b84a1b2 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -614,6 +614,8 @@ id MVKImageView::newMTLTexture() { MVKImageView::MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCreateInfo) : MVKBaseDeviceObject(device) { + validateImageViewConfig(pCreateInfo); + _image = (MVKImage*)pCreateInfo->image; // Remember the subresource range, and determine the actual number of mip levels and texture slices @@ -631,6 +633,17 @@ MVKImageView::MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCrea initMTLTextureViewSupport(); } +// Validate whether the image view configuration can be supported +void MVKImageView::validateImageViewConfig(const VkImageViewCreateInfo* pCreateInfo) { + VkImageType imgType = ((MVKImage*)pCreateInfo->image)->getImageType(); + VkImageViewType viewType = pCreateInfo->viewType; + + // VK_KHR_maintenance1 supports taking 2D image views of 3D slices. No dice in Metal. + if ((viewType == VK_IMAGE_VIEW_TYPE_2D || viewType == VK_IMAGE_VIEW_TYPE_2D_ARRAY) && (imgType == VK_IMAGE_TYPE_3D)) { + setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): Metal does not support creating a 2D view on a 3D image.")); + } +} + // Returns a MTLPixelFormat, based on the original MTLPixelFormat, as converted from the VkFormat, // but possibly modified by the swizzles defined in the VkComponentMapping of the VkImageViewCreateInfo. // Metal does not support general per-texture swizzles, and so this function relies on a few coincidental diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h index e0222dba..9c6fdc25 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h @@ -18,6 +18,7 @@ #pragma once +#include "MVKLayers.h" #include "MVKSurface.h" #include "MVKBaseObject.h" #include "vk_mvk_moltenvk.h" @@ -53,17 +54,8 @@ public: */ VkResult getPhysicalDevices(uint32_t* pCount, VkPhysicalDevice* pPhysicalDevices); - /** - * Verifies that the list of layers are available, - * and returns VK_SUCCESS or VK_ERROR_LAYER_NOT_PRESENT. - */ - VkResult verifyLayers(uint32_t count, const char* const* names); - - /** - * Verifies that the list of extensions are available, - * and returns VK_SUCCESS or VK_ERROR_EXTENSION_NOT_PRESENT. - */ - VkResult verifyExtensions(uint32_t count, const char* const* names); + /** Returns the driver layer. */ + MVKLayer* getDriverLayer() { return MVKLayerManager::globalManager()->getDriverLayer(); } /** Creates and returns a new object. */ MVKSurface* createSurface(const Vk_PLATFORM_SurfaceCreateInfoMVK* pCreateInfo, @@ -73,8 +65,14 @@ public: void destroySurface(MVKSurface* mvkSrfc, const VkAllocationCallbacks* pAllocator); - /** The MoltenVK configuration settings. */ - MVKConfiguration _mvkConfig; + /** Returns the MoltenVK configuration settings. */ + const MVKConfiguration* getMoltenVKConfiguration() { return &_mvkConfig; } + + /** Returns the MoltenVK configuration settings. */ + void setMoltenVKConfiguration(MVKConfiguration* mvkConfig) { _mvkConfig = *mvkConfig; } + + /** The list of Vulkan extensions, indicating whether each has been enabled by the app. */ + const MVKExtensionList _enabledExtensions; #pragma mark Object Creation @@ -102,7 +100,9 @@ protected: void initProcAddrs(); void initConfig(); void logVersions(); + VkResult verifyLayers(uint32_t count, const char* const* names); + MVKConfiguration _mvkConfig; VkApplicationInfo _appInfo; std::vector _physicalDevices; std::unordered_map _procAddrMap; diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index 7ab058d0..4ab0af7e 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -18,7 +18,6 @@ #include "MVKInstance.h" -#include "MVKLayers.h" #include "MVKDevice.h" #include "MVKFoundation.h" #include "MVKEnvironment.h" @@ -53,27 +52,6 @@ VkResult MVKInstance::getPhysicalDevices(uint32_t* pCount, VkPhysicalDevice* pPh return result; } -VkResult MVKInstance::verifyLayers(uint32_t count, const char* const* names) { - VkResult result = VK_SUCCESS; - for (uint32_t i = 0; i < count; i++) { - if ( !MVKLayerManager::globalManager()->getLayerNamed(names[i]) ) { - result = mvkNotifyErrorWithText(VK_ERROR_LAYER_NOT_PRESENT, "Vulkan layer %s is not supported.", names[i]); - } - } - return result; -} - -VkResult MVKInstance::verifyExtensions(uint32_t count, const char* const* names) { - VkResult result = VK_SUCCESS; - MVKLayer* driverLayer = MVKLayerManager::globalManager()->getDriverLayer(); - for (uint32_t i = 0; i < count; i++) { - if (!driverLayer->hasExtensionNamed(names[i])) { - result = mvkNotifyErrorWithText(VK_ERROR_EXTENSION_NOT_PRESENT, "Vulkan extension %s is not supported.", names[i]); - } - } - return result; -} - MVKSurface* MVKInstance::createSurface(const Vk_PLATFORM_SurfaceCreateInfoMVK* pCreateInfo, const VkAllocationCallbacks* pAllocator) { return new MVKSurface(this, pCreateInfo, pAllocator); @@ -87,6 +65,74 @@ void MVKInstance::destroySurface(MVKSurface* mvkSrfc, #pragma mark Object Creation +// Returns an autoreleased array containing the MTLDevices available on this system, +// sorted according to power, with higher power GPU's at the front of the array. +// This ensures that a lazy app that simply grabs the first GPU will get a high-power one by default. +// If the MVK_FORCE_LOW_POWER_GPU is defined, the returned array will only include low-power devices. +static NSArray>* getAvailableMTLDevices() { +#if MVK_MACOS + NSArray* mtlDevs = [MTLCopyAllDevices() autorelease]; + +#ifdef MVK_FORCE_LOW_POWER_GPU + NSMutableArray* lpDevs = [[NSMutableArray new] autorelease]; + for (id md in mtlDevs) { + if (md.isLowPower) { [lpDevs addObject: md]; } + } + return lpDevs; +#else + return [mtlDevs sortedArrayUsingComparator: ^(id md1, id md2) { + BOOL md1IsLP = md1.isLowPower; + BOOL md2IsLP = md2.isLowPower; + + if (md1IsLP == md2IsLP) { + // If one device is headless and the other one is not, select the + // one that is not headless first. + BOOL md1IsHeadless = md1.isHeadless; + BOOL md2IsHeadless = md2.isHeadless; + if (md1IsHeadless == md2IsHeadless ) { + return NSOrderedSame; + } + return md2IsHeadless ? NSOrderedAscending : NSOrderedDescending; + } + + return md2IsLP ? NSOrderedAscending : NSOrderedDescending; + }]; +#endif // MVK_MACOS + +#endif +#if MVK_IOS + return [NSArray arrayWithObject: MTLCreateSystemDefaultDevice()]; +#endif +} + +MVKInstance::MVKInstance(const VkInstanceCreateInfo* pCreateInfo) { + + _appInfo.apiVersion = MVK_VULKAN_API_VERSION; // Default + mvkSetOrClear(&_appInfo, pCreateInfo->pApplicationInfo); + + initProcAddrs(); // Init function pointers + initConfig(); + + setConfigurationResult(verifyLayers(pCreateInfo->enabledLayerCount, pCreateInfo->ppEnabledLayerNames)); + MVKExtensionList* pWritableExtns = (MVKExtensionList*)&_enabledExtensions; + setConfigurationResult(pWritableExtns->enable(pCreateInfo->enabledExtensionCount, + pCreateInfo->ppEnabledExtensionNames, + getDriverLayer()->getSupportedExtensions())); + logVersions(); // Log the MoltenVK and Vulkan versions + + if (MVK_VULKAN_API_VERSION_CONFORM(MVK_VULKAN_API_VERSION) < + MVK_VULKAN_API_VERSION_CONFORM(_appInfo.apiVersion)) { + setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_INCOMPATIBLE_DRIVER, "Request for driver version %x is not compatible with provided version %x.", _appInfo.apiVersion, MVK_VULKAN_API_VERSION)); + } + + // Populate the array of physical GPU devices + NSArray>* mtlDevices = getAvailableMTLDevices(); + _physicalDevices.reserve(mtlDevices.count); + for (id mtlDev in mtlDevices) { + _physicalDevices.push_back(new MVKPhysicalDevice(this, mtlDev)); + } +} + #define ADD_PROC_ADDR(entrypoint) _procAddrMap[""#entrypoint] = (PFN_vkVoidFunction)&entrypoint; /** Initializes the function pointer map. */ @@ -268,7 +314,6 @@ void MVKInstance::initProcAddrs() { ADD_PROC_ADDR(vkSetMoltenVKDeviceConfigurationMVK); #pragma clang diagnostic pop - } void MVKInstance::logVersions() { @@ -276,74 +321,14 @@ void MVKInstance::logVersions() { char mvkVer[buffLen]; char vkVer[buffLen]; vkGetVersionStringsMVK(mvkVer, buffLen, vkVer, buffLen); - MVKLogInfo("MoltenVK version %s. Vulkan version %s.", mvkVer, vkVer); -} -/** - * Returns an autoreleased array containing the MTLDevices available on this system, - * sorted according to power, with higher power GPU's at the front of the array. - * This ensures that a lazy app that simply grabs the first GPU will get a high-power one by default. - * If the MVK_FORCE_LOW_POWER_GPU is defined, the returned array will only include low-power devices. - */ -static NSArray>* getAvailableMTLDevices() { -#if MVK_MACOS - NSArray* mtlDevs = [MTLCopyAllDevices() autorelease]; - -#ifdef MVK_FORCE_LOW_POWER_GPU - NSMutableArray* lpDevs = [[NSMutableArray new] autorelease]; - for (id md in mtlDevs) { - if (md.isLowPower) { [lpDevs addObject: md]; } - } - return lpDevs; -#else - return [mtlDevs sortedArrayUsingComparator: ^(id md1, id md2) { - BOOL md1IsLP = md1.isLowPower; - BOOL md2IsLP = md2.isLowPower; - - if (md1IsLP == md2IsLP) { - // If one device is headless and the other one is not, select the - // one that is not headless first. - BOOL md1IsHeadless = md1.isHeadless; - BOOL md2IsHeadless = md2.isHeadless; - if (md1IsHeadless == md2IsHeadless ) { - return NSOrderedSame; - } - return md2IsHeadless ? NSOrderedAscending : NSOrderedDescending; - } - - return md2IsLP ? NSOrderedAscending : NSOrderedDescending; - }]; -#endif // MVK_MACOS - -#endif -#if MVK_IOS - return [NSArray arrayWithObject: MTLCreateSystemDefaultDevice()]; -#endif -} - -MVKInstance::MVKInstance(const VkInstanceCreateInfo* pCreateInfo) { - - _appInfo.apiVersion = MVK_VULKAN_API_VERSION; // Default - mvkSetOrClear(&_appInfo, pCreateInfo->pApplicationInfo); - - logVersions(); // Log the MoltenVK and Vulkan versions - initProcAddrs(); // Init function pointers - initConfig(); - - // Populate the array of physical GPU devices - NSArray>* mtlDevices = getAvailableMTLDevices(); - _physicalDevices.reserve(mtlDevices.count); - for (id mtlDev in mtlDevices) { - _physicalDevices.push_back(new MVKPhysicalDevice(this, mtlDev)); - } - - if (MVK_VULKAN_API_VERSION_CONFORM(MVK_VULKAN_API_VERSION) < - MVK_VULKAN_API_VERSION_CONFORM(_appInfo.apiVersion)) { - setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_INCOMPATIBLE_DRIVER, "Request for driver version %x is not compatible with provided version %x.", _appInfo.apiVersion, MVK_VULKAN_API_VERSION)); - } - - setConfigurationResult(verifyLayers(pCreateInfo->enabledLayerCount, pCreateInfo->ppEnabledLayerNames)); - setConfigurationResult(verifyExtensions(pCreateInfo->enabledExtensionCount, pCreateInfo->ppEnabledExtensionNames)); + const char* indent = "\n\t\t"; + string logMsg = "MoltenVK version %s. Vulkan version %s."; + logMsg += "\n\tThe following Vulkan extensions are supported:"; + logMsg += getDriverLayer()->getSupportedExtensions()->enabledNamesString(indent, true); + logMsg += "\n\tCreated VkInstance with the following Vulkan extensions enabled:"; + logMsg += _enabledExtensions.enabledNamesString(indent, true); + MVKLogInfo(logMsg.c_str(), mvkVer, vkVer); } // Init config. @@ -360,6 +345,16 @@ void MVKInstance::initConfig() { _mvkConfig.metalCompileTimeout = MVK_CONFIG_METAL_COMPILE_TIMEOUT; } +VkResult MVKInstance::verifyLayers(uint32_t count, const char* const* names) { + VkResult result = VK_SUCCESS; + for (uint32_t i = 0; i < count; i++) { + if ( !MVKLayerManager::globalManager()->getLayerNamed(names[i]) ) { + result = mvkNotifyErrorWithText(VK_ERROR_LAYER_NOT_PRESENT, "Vulkan layer %s is not supported.", names[i]); + } + } + return result; +} + MVKInstance::~MVKInstance() { mvkDestroyContainerContents(_physicalDevices); } diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp new file mode 100644 index 00000000..f66b76a7 --- /dev/null +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp @@ -0,0 +1,149 @@ +/* + * MVKExtensions.cpp + * + * Copyright (c) 2014-2018 The Brenwill Workshop Ltd. (http://www.brenwill.com) + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "MVKExtensions.h" +#include "MVKFoundation.h" +#include "vk_mvk_moltenvk.h" +#include +#include + +using namespace std; + + +// Returns a VkExtensionProperties struct populated with a name and version +static VkExtensionProperties mvkMakeExtProps(const char* extensionName, uint32_t specVersion) { + VkExtensionProperties extProps; + memset(extProps.extensionName, 0, sizeof(extProps.extensionName)); + if (extensionName) { strcpy(extProps.extensionName, extensionName); } + extProps.specVersion = specVersion; + return extProps; +} + +// Declares and populates a static VkExtensionProperties variable for the extention EXT, +// which should include the unique portion of the extension name, as uppercase. +// For example, for the extension VK_KHR_surface, use KHR_SURFACE. +#define MVK_MAKE_VK_EXT_PROPS(EXT) \ +static VkExtensionProperties kVkExtProps_ ##EXT = mvkMakeExtProps(VK_ ##EXT ##_EXTENSION_NAME, VK_ ##EXT ##_SPEC_VERSION) + +// Extension properties +MVK_MAKE_VK_EXT_PROPS(MVK_MOLTENVK); +MVK_MAKE_VK_EXT_PROPS(MVK_MACOS_SURFACE); +MVK_MAKE_VK_EXT_PROPS(MVK_IOS_SURFACE); +MVK_MAKE_VK_EXT_PROPS(KHR_SURFACE); +MVK_MAKE_VK_EXT_PROPS(KHR_SWAPCHAIN); +MVK_MAKE_VK_EXT_PROPS(KHR_MAINTENANCE1); +MVK_MAKE_VK_EXT_PROPS(IMG_FORMAT_PVRTC); +MVK_MAKE_VK_EXT_PROPS(AMD_NEGATIVE_VIEWPORT_HEIGHT); +MVK_MAKE_VK_EXT_PROPS(KHR_SHADER_DRAW_PARAMETERS); +MVK_MAKE_VK_EXT_PROPS(KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2); +MVK_MAKE_VK_EXT_PROPS(KHR_PUSH_DESCRIPTOR); + +// Calls the constructor for a MVKExtension member variable, using the member name and the +// portion of the extension name, as uppercase, used in the MVK_MAKE_VK_EXT_PROPS() macro above. +// For example, for the memeber variable vk_KHR_surface, use MVKExt_CONSTRUCT(vk_KHR_surface, KHR_SURFACE). +#define MVKExt_CONSTRUCT(var, EXT) var(&kVkExtProps_ ##EXT, enableForPlatform) + +MVKExtensionList::MVKExtensionList(bool enableForPlatform) : + MVKExt_CONSTRUCT(vk_MVK_moltenvk, MVK_MOLTENVK), + MVKExt_CONSTRUCT(vk_MVK_macos_surface, MVK_MACOS_SURFACE), + MVKExt_CONSTRUCT(vk_MVK_ios_surface, MVK_IOS_SURFACE), + MVKExt_CONSTRUCT(vk_KHR_surface, KHR_SURFACE), + MVKExt_CONSTRUCT(vk_KHR_swapchain, KHR_SWAPCHAIN), + MVKExt_CONSTRUCT(vk_KHR_maintenance1, KHR_MAINTENANCE1), + MVKExt_CONSTRUCT(vk_IMG_format_pvrtc, IMG_FORMAT_PVRTC), + MVKExt_CONSTRUCT(vk_AMD_negative_viewport_height, AMD_NEGATIVE_VIEWPORT_HEIGHT), + MVKExt_CONSTRUCT(vk_KHR_shader_draw_parameters, KHR_SHADER_DRAW_PARAMETERS), + MVKExt_CONSTRUCT(vk_KHR_get_physical_device_properties2, KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2), + MVKExt_CONSTRUCT(vk_KHR_push_descriptor, KHR_PUSH_DESCRIPTOR) +{} + +bool MVKExtensionList::isEnabled(const char* extnName) const { + uint32_t extnCnt = getCount(); + const MVKExtension* extnAry = &extensionArray; + for (uint32_t extnIdx = 0; extnIdx < extnCnt; extnIdx++) { + const MVKExtension& extn = extnAry[extnIdx]; + if ( strcmp(extn.pProperties->extensionName, extnName) == 0 ) { + return extn.enabled; + } + } + return false; +} + +void MVKExtensionList::enable(const char* extnName) { + uint32_t extnCnt = getCount(); + MVKExtension* extnAry = &extensionArray; + for (uint32_t extnIdx = 0; extnIdx < extnCnt; extnIdx++) { + MVKExtension& extn = extnAry[extnIdx]; + if ( strcmp(extn.pProperties->extensionName, extnName) == 0 ) { + extn.enabled = true; + return; + } + } +} + +VkResult MVKExtensionList::enable(uint32_t count, const char* const* names, MVKExtensionList* parent) { + VkResult result = VK_SUCCESS; + for (uint32_t i = 0; i < count; i++) { + auto extnName = names[i]; + if (parent && !parent->isEnabled(extnName)) { + result = mvkNotifyErrorWithText(VK_ERROR_EXTENSION_NOT_PRESENT, "Vulkan extension %s is not supported.", extnName); + } else { + enable(extnName); + } + } + return result; +} + +string MVKExtensionList::enabledNamesString(const char* separator, bool prefixFirstWithSeparator) const { + string logMsg; + bool isFirst = true; + uint32_t extnCnt = getCount(); + const MVKExtension* extnAry = &extensionArray; + for (uint32_t extnIdx = 0; extnIdx < extnCnt; extnIdx++) { + const MVKExtension& extn = extnAry[extnIdx]; + if (extn.enabled) { + if ( !isFirst || prefixFirstWithSeparator ) { logMsg += separator; } + logMsg += extn.pProperties->extensionName; + logMsg += " v"; + logMsg += to_string(extn.pProperties->specVersion); + isFirst = false; + } + } + return logMsg; +} + +// Returns whether the specified properties are valid for this platform +static bool mvkIsSupportedOnPlatform(VkExtensionProperties* pProperties) { +#if !(MVK_IOS) + if (pProperties == &kVkExtProps_MVK_IOS_SURFACE) { return false; } + if (pProperties == &kVkExtProps_IMG_FORMAT_PVRTC) { return false; } +#endif +#if !(MVK_MACOS) + if (pProperties == &kVkExtProps_MVK_MACOS_SURFACE) { return false; } +#endif + + if (pProperties == &kVkExtProps_AMD_NEGATIVE_VIEWPORT_HEIGHT) { return false; } + + return true; +} + +// Disable by default unless asked to enable for platform and the extension is valid for this platform +MVKExtension::MVKExtension(VkExtensionProperties* pProperties, bool enableForPlatform) { + this->pProperties = pProperties; + this->enabled = enableForPlatform && mvkIsSupportedOnPlatform(pProperties); +} diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.h b/MoltenVK/MoltenVK/Layers/MVKExtensions.h new file mode 100644 index 00000000..1b2a2e2a --- /dev/null +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.h @@ -0,0 +1,82 @@ +/* + * MVKExtensions.h + * + * Copyright (c) 2014-2018 The Brenwill Workshop Ltd. (http://www.brenwill.com) + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "mvk_vulkan.h" +#include + +/** Describes a Vulkan extension and whether or not it is enabled or supported. */ +struct MVKExtension { + bool enabled = false; + VkExtensionProperties* pProperties; + + MVKExtension(VkExtensionProperties* pProperties, bool enableForPlatform = false); +}; + +/** + * A fixed list of the Vulkan extensions known to MoltenVK, with + * an indication of whether each extension is supported/enabled. + * + * To add support for a Vulkan extension, add a variable to this list. + */ +struct MVKExtensionList { + union { + struct { + MVKExtension vk_MVK_moltenvk; + MVKExtension vk_MVK_macos_surface; + MVKExtension vk_MVK_ios_surface; + MVKExtension vk_KHR_surface; + MVKExtension vk_KHR_swapchain; + MVKExtension vk_KHR_maintenance1; + MVKExtension vk_IMG_format_pvrtc; + MVKExtension vk_AMD_negative_viewport_height; + MVKExtension vk_KHR_shader_draw_parameters; + MVKExtension vk_KHR_get_physical_device_properties2; + MVKExtension vk_KHR_push_descriptor; + }; + MVKExtension extensionArray; + }; + + /** Returns the total number of extensions that are tracked by this object. */ + static uint32_t getCount() { return sizeof(MVKExtensionList) / sizeof(MVKExtension); } + + /** Returns whether the named extension is enabled. */ + bool isEnabled(const char* extnName) const; + + /** Enables the named extension. */ + void enable(const char* extnName); + + /** + * Enables the named extensions. + * + * If parent is non null, the extension must also be enabled in the parent in order + * for it to be enabled here. If it is not enabled in the parent, an error is logged + * and returned. Returns VK_SUCCESS if all requested extensions were able to be enabled. + */ + VkResult enable(uint32_t count, const char* const* names, MVKExtensionList* parent = nullptr); + + /** + * Returns a string containing the names of the enabled extensions, separated by the separator string. + * If prefixFirstWithSeparator is true the separator will also appear before the first extension name. + */ + std::string enabledNamesString(const char* separator = " ", bool prefixFirstWithSeparator = false) const; + + MVKExtensionList(bool enableForPlatform = false); +}; + diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.h b/MoltenVK/MoltenVK/Layers/MVKLayers.h similarity index 90% rename from MoltenVK/MoltenVK/Loader/MVKLayers.h rename to MoltenVK/MoltenVK/Layers/MVKLayers.h index 2fb56dd1..f539a312 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.h +++ b/MoltenVK/MoltenVK/Layers/MVKLayers.h @@ -18,8 +18,8 @@ #pragma once -#include "mvk_vulkan.h" #include "MVKBaseObject.h" +#include "MVKExtensions.h" #include @@ -50,16 +50,16 @@ public: */ VkResult getExtensionProperties(uint32_t* pCount, VkExtensionProperties* pProperties); - /** Returns whether this layer supports the specified extension. */ - bool hasExtensionNamed(const char* extnName); + /** Returns the list of supported extensions. */ + MVKExtensionList* getSupportedExtensions() { return &_supportedExtensions; } /** Default constructor. This represents the driver implementation. */ MVKLayer(); protected: VkLayerProperties _layerProperties; - std::vector _extensions; - + MVKExtensionList _supportedExtensions; + }; @@ -70,12 +70,11 @@ class MVKLayerManager : public MVKConfigurableObject { public: - /** Returns a pointer to the driver layer. */ + /** Returns the driver layer. */ MVKLayer* getDriverLayer(); /** - * Returns a pointe to the layer with the specified name, - * or null if no layer was found with that name. + * Returns the layer with the specified name, or null if no layer was found with that name. * * If pLayerName is null, returns the driver layer, which is * the same layer returned by the getDriverLayer() function. diff --git a/MoltenVK/MoltenVK/Loader/MVKLayers.mm b/MoltenVK/MoltenVK/Layers/MVKLayers.mm similarity index 51% rename from MoltenVK/MoltenVK/Loader/MVKLayers.mm rename to MoltenVK/MoltenVK/Layers/MVKLayers.mm index fed05c0a..d1764cd5 100644 --- a/MoltenVK/MoltenVK/Loader/MVKLayers.mm +++ b/MoltenVK/MoltenVK/Layers/MVKLayers.mm @@ -32,88 +32,43 @@ VkLayerProperties* const MVKLayer::getLayerProperties() { return &_layerProperti VkResult MVKLayer::getExtensionProperties(uint32_t* pCount, VkExtensionProperties* pProperties) { - // If properties aren't actually being requested yet, simply update the returned count - if ( !pProperties ) { - *pCount = (uint32_t)_extensions.size(); - return VK_SUCCESS; + uint32_t enabledCnt = 0; + + // Iterate extensions and handle those that are enabled. Count them, + // and if they are to be returned, and there is room, do so. + uint32_t extnCnt = _supportedExtensions.getCount(); + MVKExtension* extnAry = &_supportedExtensions.extensionArray; + for (uint32_t extnIdx = 0; extnIdx < extnCnt; extnIdx++) { + if (extnAry[extnIdx].enabled) { + if (pProperties) { + if (enabledCnt < *pCount) { + pProperties[enabledCnt] = *(extnAry[extnIdx].pProperties); + } else { + return VK_INCOMPLETE; + } + } + enabledCnt++; + } } - // Othewise, determine how many extensions we'll return, and return that count - uint32_t extCnt = (uint32_t)_extensions.size(); - VkResult result = (*pCount <= extCnt) ? VK_SUCCESS : VK_INCOMPLETE; - *pCount = min(extCnt, *pCount); - - // Now populate the layer properties - for (uint32_t extIdx = 0; extIdx < *pCount; extIdx++) { - pProperties[extIdx] = _extensions[extIdx]; - } - - return result; -} - -bool MVKLayer::hasExtensionNamed(const char* extnName) { - for (auto& extn : _extensions) { - if ( strcmp(extn.extensionName, extnName) == 0 ) { return true; } - } - return false; + // Return the count of enabled extensions. This will either be a + // count of all enabled extensions, or a count of those returned. + *pCount = enabledCnt; + return VK_SUCCESS; } #pragma mark Object Creation -MVKLayer::MVKLayer() { +MVKLayer::MVKLayer() : _supportedExtensions(true) { // The core driver layer + memset(_layerProperties.layerName, 0, sizeof(_layerProperties.layerName)); strcpy(_layerProperties.layerName, "MoltenVK"); + memset(_layerProperties.description, 0, sizeof(_layerProperties.description)); strcpy(_layerProperties.description, "MoltenVK driver layer"); _layerProperties.specVersion = MVK_VULKAN_API_VERSION; _layerProperties.implementationVersion = MVK_VERSION; - - // Extensions - VkExtensionProperties extTmplt; - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_MVK_MOLTENVK_EXTENSION_NAME); - extTmplt.specVersion = VK_MVK_MOLTENVK_SPEC_VERSION; - _extensions.push_back(extTmplt); - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_KHR_SWAPCHAIN_EXTENSION_NAME); - extTmplt.specVersion = VK_KHR_SWAPCHAIN_SPEC_VERSION; - _extensions.push_back(extTmplt); - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_KHR_SURFACE_EXTENSION_NAME); - extTmplt.specVersion = VK_KHR_SURFACE_SPEC_VERSION; - _extensions.push_back(extTmplt); - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_AMD_NEGATIVE_VIEWPORT_HEIGHT_EXTENSION_NAME); - extTmplt.specVersion = VK_AMD_NEGATIVE_VIEWPORT_HEIGHT_SPEC_VERSION; - _extensions.push_back(extTmplt); - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_KHR_MAINTENANCE1_EXTENSION_NAME); - extTmplt.specVersion = VK_KHR_MAINTENANCE1_SPEC_VERSION; - _extensions.push_back(extTmplt); - -#if MVK_IOS - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_MVK_IOS_SURFACE_EXTENSION_NAME); - extTmplt.specVersion = VK_MVK_IOS_SURFACE_SPEC_VERSION; - _extensions.push_back(extTmplt); - - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_IMG_FORMAT_PVRTC_EXTENSION_NAME); - extTmplt.specVersion = VK_IMG_FORMAT_PVRTC_SPEC_VERSION; - _extensions.push_back(extTmplt); -#endif -#if MVK_MACOS - memset(extTmplt.extensionName, 0, sizeof(extTmplt.extensionName)); - strcpy(extTmplt.extensionName, VK_MVK_MACOS_SURFACE_EXTENSION_NAME); - extTmplt.specVersion = VK_MVK_MACOS_SURFACE_SPEC_VERSION; - _extensions.push_back(extTmplt); -#endif } @@ -135,7 +90,6 @@ MVKLayer* MVKLayerManager::getLayerNamed(const char* pLayerName) { return VK_NULL_HANDLE; } - VkResult MVKLayerManager::getLayerProperties(uint32_t* pCount, VkLayerProperties* pProperties) { // If properties aren't actually being requested yet, simply update the returned count diff --git a/MoltenVK/MoltenVK/OS/MVKOSExtensions.h b/MoltenVK/MoltenVK/OS/MVKOSExtensions.h index 4c655907..f2e520b1 100644 --- a/MoltenVK/MoltenVK/OS/MVKOSExtensions.h +++ b/MoltenVK/MoltenVK/OS/MVKOSExtensions.h @@ -18,7 +18,7 @@ #pragma once -#include +#include "mvk_vulkan.h" #import diff --git a/MoltenVK/MoltenVK/Utility/MVKBaseObject.h b/MoltenVK/MoltenVK/Utility/MVKBaseObject.h index 683ba4b5..f15aeac3 100644 --- a/MoltenVK/MoltenVK/Utility/MVKBaseObject.h +++ b/MoltenVK/MoltenVK/Utility/MVKBaseObject.h @@ -18,7 +18,7 @@ #pragma once -#include +#include "mvk_vulkan.h" #include #include diff --git a/MoltenVK/MoltenVK/Utility/MVKFoundation.mm b/MoltenVK/MoltenVK/Utility/MVKFoundation.cpp similarity index 95% rename from MoltenVK/MoltenVK/Utility/MVKFoundation.mm rename to MoltenVK/MoltenVK/Utility/MVKFoundation.cpp index 847c2232..fab2bcc1 100644 --- a/MoltenVK/MoltenVK/Utility/MVKFoundation.mm +++ b/MoltenVK/MoltenVK/Utility/MVKFoundation.cpp @@ -1,5 +1,5 @@ /* - * MVKFoundation.mm + * MVKFoundation.cpp * * Copyright (c) 2014-2018 The Brenwill Workshop Ltd. (http://www.brenwill.com) * @@ -52,8 +52,8 @@ char* mvkResultName(VkResult vkResult, char* name) { CASE_RESULT(VK_ERROR_VALIDATION_FAILED_EXT) CASE_RESULT(VK_ERROR_INVALID_SHADER_NV) - CASE_RESULT(VK_ERROR_OUT_OF_POOL_MEMORY_KHR) - CASE_RESULT(VK_ERROR_INVALID_EXTERNAL_HANDLE_KHR) + CASE_RESULT(VK_ERROR_OUT_OF_POOL_MEMORY) + CASE_RESULT(VK_ERROR_INVALID_EXTERNAL_HANDLE) default: sprintf(name, "UNKNOWN_VkResult(%d)", vkResult); diff --git a/MoltenVK/MoltenVK/Utility/MVKFoundation.h b/MoltenVK/MoltenVK/Utility/MVKFoundation.h index 2e1ae3e7..22539cf5 100644 --- a/MoltenVK/MoltenVK/Utility/MVKFoundation.h +++ b/MoltenVK/MoltenVK/Utility/MVKFoundation.h @@ -20,7 +20,7 @@ #pragma once -#include +#include "mvk_vulkan.h" #include "MVKLogging.h" #include #include @@ -311,11 +311,13 @@ void mvkDestroyContainerContents(C& container) { * Iterates through the contents of the specified Objective-C object pointer * container and releases each object, and clears the container. */ +#ifdef __OBJC__ template void mvkReleaseContainerContents(C& container) { for (auto elem : container) { [elem release]; } container.clear(); } +#endif /** Removes the first occurance of the specified value from the specified container. */ template diff --git a/MoltenVK/MoltenVK/Vulkan/vk_mvk_moltenvk.mm b/MoltenVK/MoltenVK/Vulkan/vk_mvk_moltenvk.mm index fee880f6..cc6bbfb4 100644 --- a/MoltenVK/MoltenVK/Vulkan/vk_mvk_moltenvk.mm +++ b/MoltenVK/MoltenVK/Vulkan/vk_mvk_moltenvk.mm @@ -32,7 +32,7 @@ MVK_PUBLIC_SYMBOL void vkGetMoltenVKConfigurationMVK( MVKConfiguration* pConfiguration) { MVKInstance* mvkInst = MVKInstance::getMVKInstance(instance); - if (pConfiguration) { *pConfiguration = mvkInst->_mvkConfig; } + if (pConfiguration) { *pConfiguration = *(MVKConfiguration*)mvkInst->getMoltenVKConfiguration(); } } MVK_PUBLIC_SYMBOL VkResult vkSetMoltenVKConfigurationMVK( @@ -40,7 +40,7 @@ MVK_PUBLIC_SYMBOL VkResult vkSetMoltenVKConfigurationMVK( MVKConfiguration* pConfiguration) { MVKInstance* mvkInst = MVKInstance::getMVKInstance(instance); - if (pConfiguration) { mvkInst->_mvkConfig = *pConfiguration; } + if (pConfiguration) { mvkInst->setMoltenVKConfiguration(pConfiguration); } return VK_SUCCESS; } @@ -142,7 +142,7 @@ MVK_PUBLIC_SYMBOL void vkGetMoltenVKDeviceConfigurationMVK( MVKDeviceConfiguration* pConfiguration) { MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); - if (pConfiguration) { *pConfiguration = mvkDev->getInstance()->_mvkConfig; } + if (pConfiguration) { *pConfiguration = *(MVKConfiguration*)mvkDev->getInstance()->getMoltenVKConfiguration(); } } MVK_PUBLIC_SYMBOL VkResult vkSetMoltenVKDeviceConfigurationMVK( @@ -150,7 +150,7 @@ MVK_PUBLIC_SYMBOL VkResult vkSetMoltenVKDeviceConfigurationMVK( MVKDeviceConfiguration* pConfiguration) { MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); - if (pConfiguration) { mvkDev->getInstance()->_mvkConfig = *pConfiguration; } + if (pConfiguration) { mvkDev->getInstance()->setMoltenVKConfiguration(pConfiguration); } return VK_SUCCESS; } From c369da1cddee4635b504233b5d9e4b126a00fdf3 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Sat, 1 Sep 2018 18:27:07 -0400 Subject: [PATCH 06/20] Update to What's New document. Update to latest version of SPIRV-Cross. --- Docs/Whats_New.md | 13 +++++++++++++ ExternalRevisions/SPIRV-Cross_repo_revision | 2 +- MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h | 2 +- 3 files changed, 15 insertions(+), 2 deletions(-) diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md index b31f70b9..83969ecd 100644 --- a/Docs/Whats_New.md +++ b/Docs/Whats_New.md @@ -12,6 +12,19 @@ Copyright (c) 2014-2018 [The Brenwill Workshop Ltd.](http://www.brenwill.com) For best results, use a Markdown reader.* +MoltenVK 1.0.20 +--------------- + +Released 2018/09/01 + +- Add support for extensions: + - VK_KHR_maintenance1; + - VK_KHR_shader_draw_parameters; + - VK_KHR_get_physical_device_properties2; + - VK_KHR_push_descriptor; +- Add ability to track and access supported and enabled extensions. +- Update to latest SPIRV-Cross version. + MoltenVK 1.0.19 --------------- diff --git a/ExternalRevisions/SPIRV-Cross_repo_revision b/ExternalRevisions/SPIRV-Cross_repo_revision index 7e060c34..4537bdb3 100644 --- a/ExternalRevisions/SPIRV-Cross_repo_revision +++ b/ExternalRevisions/SPIRV-Cross_repo_revision @@ -1 +1 @@ -e14bf77b1ac99943aa27c6b9f6446ea2c4a824f7 +6fd66664e8bdadd3f6281aad711f771ef9c24bbe diff --git a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h index 14067b36..22b891e7 100644 --- 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 19 +#define MVK_VERSION_PATCH 20 #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) From 63c6e4d8c421f1ba40ec740c131edf2b6a24f0f0 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Thu, 30 Aug 2018 12:06:01 -0500 Subject: [PATCH 07/20] Support the VK_KHR_descriptor_update_template extension. --- MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h | 34 +++ MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm | 71 +++++++ MoltenVK/MoltenVK/Commands/MVKCommandPool.h | 2 + MoltenVK/MoltenVK/Commands/MVKCommandPool.mm | 3 +- .../MoltenVK/GPUObjects/MVKDescriptorSet.h | 58 ++++- .../MoltenVK/GPUObjects/MVKDescriptorSet.mm | 198 ++++++++++++++---- MoltenVK/MoltenVK/GPUObjects/MVKDevice.h | 6 + MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm | 11 + MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm | 4 + MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h | 6 + MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm | 13 ++ MoltenVK/MoltenVK/Layers/MVKExtensions.cpp | 2 + MoltenVK/MoltenVK/Layers/MVKExtensions.h | 1 + MoltenVK/MoltenVK/Vulkan/vulkan.mm | 47 +++++ 14 files changed, 404 insertions(+), 52 deletions(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h index f9fd523a..a6002f7a 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h +++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h @@ -25,6 +25,7 @@ class MVKCommandBuffer; class MVKPipeline; class MVKPipelineLayout; class MVKDescriptorSet; +class MVKDescriptorUpdateTemplate; #pragma mark - @@ -160,6 +161,32 @@ private: }; +#pragma mark - +#pragma mark MVKCmdPushDescriptorSetWithTemplate + +/** Vulkan command to update a descriptor set from a template. */ +class MVKCmdPushDescriptorSetWithTemplate : public MVKCommand { + +public: + void setContent(VkDescriptorUpdateTemplateKHR descUpdateTemplate, + VkPipelineLayout layout, + uint32_t set, + const void* pData); + + void encode(MVKCommandEncoder* cmdEncoder) override; + + MVKCmdPushDescriptorSetWithTemplate(MVKCommandTypePool* pool); + + ~MVKCmdPushDescriptorSetWithTemplate() override; + +private: + MVKDescriptorUpdateTemplate* _descUpdateTemplate; + MVKPipelineLayout* _pipelineLayout; + uint32_t _set; + void* _pData; +}; + + #pragma mark - #pragma mark Command creation functions @@ -205,3 +232,10 @@ void mvkCmdPushDescriptorSet(MVKCommandBuffer* cmdBuff, uint32_t set, uint32_t descriptorWriteCount, const VkWriteDescriptorSet* pDescriptorWrites); + +/** Adds commands to the specified command buffer that update the specified descriptor set from the given template. */ +void mvkCmdPushDescriptorSetWithTemplate(MVKCommandBuffer* cmdBuff, + VkDescriptorUpdateTemplateKHR descUpdateTemplate, + VkPipelineLayout layout, + uint32_t set, + const void* pData); diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm index d21af9ff..57d1ac83 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm @@ -237,6 +237,67 @@ void MVKCmdPushDescriptorSet::clearDescriptorWrites() { } +#pragma mark - +#pragma mark MVKCmdPushDescriptorSetWithTemplate + +void MVKCmdPushDescriptorSetWithTemplate::setContent(VkDescriptorUpdateTemplateKHR descUpdateTemplate, + VkPipelineLayout layout, + uint32_t set, + const void* pData) { + _descUpdateTemplate = (MVKDescriptorUpdateTemplate*)descUpdateTemplate; + _pipelineLayout = (MVKPipelineLayout*)layout; + _set = set; + if (_pData) delete[] (char*)_pData; + // Work out how big the memory block in pData is. + const VkDescriptorUpdateTemplateEntryKHR* pEntry = + _descUpdateTemplate->getEntry(_descUpdateTemplate->getNumberOfEntries()-1); + size_t size = pEntry->offset; + // If we were given a stride, use that; otherwise, assume only one info + // struct of the appropriate type. + if (pEntry->stride) + size += pEntry->stride * pEntry->descriptorCount; + else switch (pEntry->descriptorType) { + + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: + size += sizeof(VkDescriptorBufferInfo); + break; + + case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: + case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: + case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: + case VK_DESCRIPTOR_TYPE_SAMPLER: + case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: + size += sizeof(VkDescriptorImageInfo); + break; + + case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: + size += sizeof(VkBufferView); + break; + + default: + break; + } + _pData = new char[size]; + memcpy(_pData, pData, size); +} + +void MVKCmdPushDescriptorSetWithTemplate::encode(MVKCommandEncoder* cmdEncoder) { + _pipelineLayout->pushDescriptorSet(cmdEncoder, _descUpdateTemplate, _set, _pData); +} + +MVKCmdPushDescriptorSetWithTemplate::MVKCmdPushDescriptorSetWithTemplate( + MVKCommandTypePool* pool) + : MVKCommand::MVKCommand((MVKCommandTypePool*)pool) {} + +MVKCmdPushDescriptorSetWithTemplate::~MVKCmdPushDescriptorSetWithTemplate() { + if (_pData) delete[] (char*)_pData; +} + + #pragma mark - #pragma mark Command creation functions @@ -300,3 +361,13 @@ void mvkCmdPushDescriptorSet(MVKCommandBuffer* cmdBuff, cmd->setContent(pipelineBindPoint, layout, set, descriptorWriteCount, pDescriptorWrites); cmdBuff->addCommand(cmd); } + +void mvkCmdPushDescriptorSetWithTemplate(MVKCommandBuffer* cmdBuff, + VkDescriptorUpdateTemplateKHR descUpdateTemplate, + VkPipelineLayout layout, + uint32_t set, + const void* pData) { + MVKCmdPushDescriptorSetWithTemplate* cmd = cmdBuff->_commandPool->_cmdPushSetWithTemplatePool.acquireObject(); + cmd->setContent(descUpdateTemplate, layout, set, pData); + cmdBuff->addCommand(cmd); +} diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h index 52db2654..3bf9ce75 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.h @@ -133,6 +133,8 @@ public: MVKCommandTypePool _cmdPushDescriptorSetPool; + MVKCommandTypePool _cmdPushSetWithTemplatePool; + #pragma mark Command resources diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm index e12ba783..8623c5c3 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm @@ -119,7 +119,8 @@ MVKCommandPool::MVKCommandPool(MVKDevice* device, _cmdPushConstantsPool(this, true), _cmdDispatchPool(this, true), _cmdDispatchIndirectPool(this, true), - _cmdPushDescriptorSetPool(this, true) + _cmdPushDescriptorSetPool(this, true), + _cmdPushSetWithTemplatePool(this, true) {} // TODO: Destroying a command pool implicitly destroys all command buffers and commands created from it. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h index 4e8a4ccf..4c9e41cb 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h @@ -82,10 +82,10 @@ public: void push(MVKCommandEncoder* cmdEncoder, uint32_t& dstArrayElement, uint32_t& descriptorCount, + uint32_t& descriptorsPushed, VkDescriptorType descriptorType, - const VkDescriptorImageInfo*& pImageInfo, - const VkDescriptorBufferInfo*& pBufferInfo, - const VkBufferView*& pTexelBufferView, + size_t stride, + const void* pData, MVKShaderResourceBinding& dslMTLRezIdxOffsets); /** Populates the specified shader converter context, at the specified descriptor set binding. */ @@ -135,6 +135,13 @@ public: MVKShaderResourceBinding& dslMTLRezIdxOffsets); + /** Encodes this descriptor set layout and the updates from the given template on the specified command encoder immediately. */ + void pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + MVKDescriptorUpdateTemplate* descUpdateTemplates, + const void* pData, + MVKShaderResourceBinding& dslMTLRezIdxOffsets); + + /** Populates the specified shader converter context, at the specified DSL index. */ void populateShaderConverterContext(SPIRVToMSLConverterContext& context, MVKShaderResourceBinding& dslMTLRezIdxOffsets, @@ -189,9 +196,8 @@ public: uint32_t writeBindings(uint32_t srcStartIndex, uint32_t dstStartIndex, uint32_t count, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView); + size_t stride, + const void* pData); /** * Updates the specified content arrays from the internal element bindings. @@ -216,6 +222,7 @@ public: uint32_t readBindings(uint32_t srcStartIndex, uint32_t dstStartIndex, uint32_t count, + VkDescriptorType& descType, VkDescriptorImageInfo* pImageInfo, VkDescriptorBufferInfo* pBufferInfo, VkBufferView* pTexelBufferView); @@ -254,15 +261,15 @@ public: /** Updates the resource bindings in this instance from the specified content. */ template void writeDescriptorSets(const DescriptorAction* pDescriptorAction, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView); + size_t stride, + const void* pData); /** * Reads the resource bindings defined in the specified content * from this instance into the specified collection of bindings. */ void readDescriptorSets(const VkCopyDescriptorSet* pDescriptorCopies, + VkDescriptorType& descType, VkDescriptorImageInfo* pImageInfo, VkDescriptorBufferInfo* pBufferInfo, VkBufferView* pTexelBufferView); @@ -311,6 +318,34 @@ protected: }; +#pragma mark - +#pragma mark MVKDescriptorUpdateTemplate + +/** Represents a Vulkan descriptor update template. */ +class MVKDescriptorUpdateTemplate : public MVKConfigurableObject { + +public: + + /** Get the nth update template entry. */ + const VkDescriptorUpdateTemplateEntryKHR* getEntry(uint32_t n) const; + + /** Get the total number of entries. */ + uint32_t getNumberOfEntries() const; + + /** Get the type of this template. */ + VkDescriptorUpdateTemplateTypeKHR getType() const; + + /** Constructs an instance for the specified device. */ + MVKDescriptorUpdateTemplate(MVKDevice* device, const VkDescriptorUpdateTemplateCreateInfoKHR* pCreateInfo); + + /** Destructor. */ + ~MVKDescriptorUpdateTemplate() override = default; + +private: + VkDescriptorUpdateTemplateTypeKHR _type; + std::vector _entries; +}; + #pragma mark - #pragma mark Support functions @@ -320,6 +355,11 @@ void mvkUpdateDescriptorSets(uint32_t writeCount, uint32_t copyCount, const VkCopyDescriptorSet* pDescriptorCopies); +/** Updates the resource bindings in the given descriptor set from the specified template. */ +void mvkUpdateDescriptorSetWithTemplate(VkDescriptorSet descriptorSet, + VkDescriptorUpdateTemplateKHR updateTemplate, + const void* pData); + /** * If the shader stage binding has a binding defined for the specified stage, populates * the context at the descriptor set binding from the shader stage resource binding. diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm index f47f4ff2..bd431537 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm @@ -173,13 +173,18 @@ void MVKDescriptorSetLayoutBinding::bind(MVKCommandEncoder* cmdEncoder, } } +template +static const T& get(const void* pData, size_t stride, uint32_t index) { + return *(T*)((const char*)pData + stride * index); +} + void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, uint32_t& dstArrayElement, uint32_t& descriptorCount, + uint32_t& descriptorsPushed, VkDescriptorType descriptorType, - const VkDescriptorImageInfo*& pImageInfo, - const VkDescriptorBufferInfo*& pBufferInfo, - const VkBufferView*& pTexelBufferView, + size_t stride, + const void* pData, MVKShaderResourceBinding& dslMTLRezIdxOffsets) { MVKMTLBufferBinding bb; MVKMTLTextureBinding tb; @@ -196,9 +201,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, descriptorCount = 0; else { descriptorCount -= _info.descriptorCount; - pImageInfo += _info.descriptorCount; - pBufferInfo += _info.descriptorCount; - pTexelBufferView += _info.descriptorCount; + descriptorsPushed = _info.descriptorCount; } return; } @@ -215,7 +218,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: { - const VkDescriptorBufferInfo& bufferInfo = pBufferInfo[rezIdx - dstArrayElement]; + const auto& bufferInfo = get(pData, stride, rezIdx - dstArrayElement); MVKBuffer* buffer = (MVKBuffer*)bufferInfo.buffer; bb.mtlBuffer = buffer->getMTLBuffer(); bb.offset = bufferInfo.offset; @@ -237,7 +240,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: { - const VkDescriptorImageInfo& imageInfo = pImageInfo[rezIdx - dstArrayElement]; + const auto& imageInfo = get(pData, stride, rezIdx - dstArrayElement); MVKImageView* imageView = (MVKImageView*)imageInfo.imageView; tb.mtlTexture = imageView->getMTLTexture(); if (_applyToVertexStage) { @@ -257,7 +260,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: { - MVKBufferView* bufferView = (MVKBufferView*)pTexelBufferView[rezIdx - dstArrayElement]; + auto* bufferView = get(pData, stride, rezIdx - dstArrayElement); tb.mtlTexture = bufferView->getMTLTexture(); if (_applyToVertexStage) { tb.index = mtlIdxs.vertexStage.textureIndex + rezIdx; @@ -277,7 +280,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, case VK_DESCRIPTOR_TYPE_SAMPLER: { MVKSampler* sampler; if (_immutableSamplers.empty()) - sampler = (MVKSampler*)pImageInfo[rezIdx - dstArrayElement].sampler; + sampler = (MVKSampler*)get(pData, stride, rezIdx - dstArrayElement).sampler; else sampler = _immutableSamplers[rezIdx]; sb.mtlSamplerState = sampler->getMTLSamplerState(); @@ -297,7 +300,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, } case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: { - const VkDescriptorImageInfo& imageInfo = pImageInfo[rezIdx - dstArrayElement]; + const auto& imageInfo = get(pData, stride, rezIdx - dstArrayElement); MVKImageView* imageView = (MVKImageView*)imageInfo.imageView; MVKSampler* sampler = _immutableSamplers.empty() ? (MVKSampler*)imageInfo.sampler : _immutableSamplers[rezIdx]; tb.mtlTexture = imageView->getMTLTexture(); @@ -333,9 +336,7 @@ void MVKDescriptorSetLayoutBinding::push(MVKCommandEncoder* cmdEncoder, descriptorCount = 0; else { descriptorCount -= _info.descriptorCount; - pImageInfo += _info.descriptorCount; - pBufferInfo += _info.descriptorCount; - pTexelBufferView += _info.descriptorCount; + descriptorsPushed = _info.descriptorCount; } } @@ -475,6 +476,41 @@ void MVKDescriptorSetLayout::bindDescriptorSet(MVKCommandEncoder* cmdEncoder, } } +static const void* getWriteParameters(VkDescriptorType type, const VkDescriptorImageInfo* pImageInfo, + const VkDescriptorBufferInfo* pBufferInfo, const VkBufferView* pTexelBufferView, + size_t& stride) { + const void* pData; + switch (type) { + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: + case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: + pData = pBufferInfo; + stride = sizeof(VkDescriptorBufferInfo); + break; + + case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: + case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: + case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: + case VK_DESCRIPTOR_TYPE_SAMPLER: + case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: + pData = pImageInfo; + stride = sizeof(VkDescriptorImageInfo); + break; + + case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: + case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: + pData = pTexelBufferView; + stride = sizeof(MVKBufferView*); + break; + + default: + pData = nullptr; + stride = 0; + } + return pData; +} + void MVKDescriptorSetLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, vector& descriptorWrites, MVKShaderResourceBinding& dslMTLRezIdxOffsets) { @@ -490,9 +526,42 @@ void MVKDescriptorSetLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, // Note: This will result in us walking off the end of the array // in case there are too many updates... but that's ill-defined anyway. for (; descriptorCount; bindIdx++) { + size_t stride; + const void* pData = getWriteParameters(descWrite.descriptorType, pImageInfo, + pBufferInfo, pTexelBufferView, stride); + uint32_t descriptorsPushed = 0; _bindings[bindIdx].push(cmdEncoder, dstArrayElement, descriptorCount, - descWrite.descriptorType, pImageInfo, pBufferInfo, - pTexelBufferView, dslMTLRezIdxOffsets); + descriptorsPushed, descWrite.descriptorType, + stride, pData, dslMTLRezIdxOffsets); + pBufferInfo += descriptorsPushed; + pImageInfo += descriptorsPushed; + pTexelBufferView += descriptorsPushed; + } + } +} + +void MVKDescriptorSetLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + MVKDescriptorUpdateTemplate* descUpdateTemplate, + const void* pData, + MVKShaderResourceBinding& dslMTLRezIdxOffsets) { + + if (!_isPushDescriptorLayout || + descUpdateTemplate->getType() != VK_DESCRIPTOR_UPDATE_TEMPLATE_TYPE_PUSH_DESCRIPTORS_KHR) + return; + for (uint32_t i = 0; i < descUpdateTemplate->getNumberOfEntries(); i++) { + const VkDescriptorUpdateTemplateEntryKHR* pEntry = descUpdateTemplate->getEntry(i); + uint32_t bindIdx = pEntry->dstBinding; + uint32_t dstArrayElement = pEntry->dstArrayElement; + uint32_t descriptorCount = pEntry->descriptorCount; + const void* pCurData = (const char*)pData + pEntry->offset; + // Note: This will result in us walking off the end of the array + // in case there are too many updates... but that's ill-defined anyway. + for (; descriptorCount; bindIdx++) { + uint32_t descriptorsPushed = 0; + _bindings[bindIdx].push(cmdEncoder, dstArrayElement, descriptorCount, + descriptorsPushed, pEntry->descriptorType, + pEntry->stride, pCurData, dslMTLRezIdxOffsets); + pCurData = (const char*)pCurData + pEntry->stride * descriptorsPushed; } } } @@ -524,9 +593,8 @@ MVKDescriptorSetLayout::MVKDescriptorSetLayout(MVKDevice* device, uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, uint32_t dstStartIndex, uint32_t count, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView) { + size_t stride, + const void* pData) { uint32_t dstCnt = MIN(count, _pBindingLayout->_info.descriptorCount - dstStartIndex); @@ -534,7 +602,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, case VK_DESCRIPTOR_TYPE_SAMPLER: for (uint32_t i = 0; i < dstCnt; i++) { uint32_t dstIdx = dstStartIndex + i; - const VkDescriptorImageInfo* pImgInfo = &pImageInfo[srcStartIndex + i]; + const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; if (_hasDynamicSamplers) { _mtlSamplers[dstIdx] = ((MVKSampler*)pImgInfo->sampler)->getMTLSamplerState(); @@ -545,7 +613,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: for (uint32_t i = 0; i < dstCnt; i++) { uint32_t dstIdx = dstStartIndex + i; - const VkDescriptorImageInfo* pImgInfo = &pImageInfo[srcStartIndex + i]; + const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; _mtlTextures[dstIdx] = ((MVKImageView*)pImgInfo->imageView)->getMTLTexture(); if (_hasDynamicSamplers) { @@ -559,7 +627,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: for (uint32_t i = 0; i < dstCnt; i++) { uint32_t dstIdx = dstStartIndex + i; - const VkDescriptorImageInfo* pImgInfo = &pImageInfo[srcStartIndex + i]; + const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; _mtlTextures[dstIdx] = ((MVKImageView*)pImgInfo->imageView)->getMTLTexture(); } @@ -571,7 +639,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: for (uint32_t i = 0; i < dstCnt; i++) { uint32_t dstIdx = dstStartIndex + i; - const VkDescriptorBufferInfo* pBuffInfo = &pBufferInfo[srcStartIndex + i]; + const auto* pBuffInfo = &get(pData, stride, srcStartIndex + i); _bufferBindings[dstIdx] = *pBuffInfo; MVKBuffer* mtlBuff = (MVKBuffer*)pBuffInfo->buffer; _mtlBuffers[dstIdx] = mtlBuff->getMTLBuffer(); @@ -583,7 +651,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: for (uint32_t i = 0; i < dstCnt; i++) { uint32_t dstIdx = dstStartIndex + i; - const VkBufferView* pBuffView = &pTexelBufferView[srcStartIndex + i]; + const auto* pBuffView = &get(pData, stride, srcStartIndex + i); _texelBufferBindings[dstIdx] = *pBuffView; _mtlTextures[dstIdx] = ((MVKBufferView*)*pBuffView)->getMTLTexture(); } @@ -598,12 +666,14 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, uint32_t MVKDescriptorBinding::readBindings(uint32_t srcStartIndex, uint32_t dstStartIndex, uint32_t count, + VkDescriptorType& descType, VkDescriptorImageInfo* pImageInfo, VkDescriptorBufferInfo* pBufferInfo, VkBufferView* pTexelBufferView) { uint32_t srcCnt = MIN(count, _pBindingLayout->_info.descriptorCount - srcStartIndex); + descType = _pBindingLayout->_info.descriptorType; switch (_pBindingLayout->_info.descriptorType) { case VK_DESCRIPTOR_TYPE_SAMPLER: case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: @@ -714,9 +784,7 @@ void MVKDescriptorBinding::initMTLSamplers(MVKDescriptorSetLayoutBinding* pBindi template void MVKDescriptorSet::writeDescriptorSets(const DescriptorAction* pDescriptorAction, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView) { + size_t stride, const void* pData) { uint32_t dstStartIdx = pDescriptorAction->dstArrayElement; uint32_t binding = pDescriptorAction->dstBinding; uint32_t origCnt = pDescriptorAction->descriptorCount; @@ -731,7 +799,7 @@ void MVKDescriptorSet::writeDescriptorSets(const DescriptorAction* pDescriptorAc uint32_t srcStartIdx = origCnt - remainCnt; remainCnt = mvkDescBind->writeBindings(srcStartIdx, dstStartIdx, remainCnt, - pImageInfo, pBufferInfo, pTexelBufferView); + stride, pData); binding++; // If not consumed, move to next consecutive binding point mvkDescBind = getBinding(binding); @@ -739,17 +807,17 @@ void MVKDescriptorSet::writeDescriptorSets(const DescriptorAction* pDescriptorAc } } -// Create concrete implementations of the two variations of the writeDescriptorSets() function. +// Create concrete implementations of the three variations of the writeDescriptorSets() function. template void MVKDescriptorSet::writeDescriptorSets(const VkWriteDescriptorSet* pDescriptorAction, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView); + size_t stride, const void *pData); template void MVKDescriptorSet::writeDescriptorSets(const VkCopyDescriptorSet* pDescriptorAction, - const VkDescriptorImageInfo* pImageInfo, - const VkDescriptorBufferInfo* pBufferInfo, - const VkBufferView* pTexelBufferView); + size_t stride, const void *pData); +template void MVKDescriptorSet::writeDescriptorSets( + const VkDescriptorUpdateTemplateEntryKHR* pDescriptorAction, + size_t stride, const void *pData); void MVKDescriptorSet::readDescriptorSets(const VkCopyDescriptorSet* pDescriptorCopy, + VkDescriptorType& descType, VkDescriptorImageInfo* pImageInfo, VkDescriptorBufferInfo* pBufferInfo, VkBufferView* pTexelBufferView) { @@ -766,7 +834,7 @@ void MVKDescriptorSet::readDescriptorSets(const VkCopyDescriptorSet* pDescriptor // MVKLogDebug("Reading MVKDescriptorBinding with binding point %d.", binding); uint32_t dstStartIdx = origCnt - remainCnt; - remainCnt = mvkDescBind->readBindings(srcStartIdx, dstStartIdx, remainCnt, + remainCnt = mvkDescBind->readBindings(srcStartIdx, dstStartIdx, remainCnt, descType, pImageInfo, pBufferInfo, pTexelBufferView); binding++; // If not consumed, move to next consecutive binding point @@ -851,6 +919,29 @@ MVKDescriptorPool::~MVKDescriptorPool() { } +#pragma mark - +#pragma mark MVKDescriptorUpdateTemplate + +const VkDescriptorUpdateTemplateEntryKHR* MVKDescriptorUpdateTemplate::getEntry(uint32_t n) const { + return &_entries[n]; +} + +uint32_t MVKDescriptorUpdateTemplate::getNumberOfEntries() const { + return (uint32_t)_entries.size(); +} + +VkDescriptorUpdateTemplateTypeKHR MVKDescriptorUpdateTemplate::getType() const { + return _type; +} + +MVKDescriptorUpdateTemplate::MVKDescriptorUpdateTemplate(MVKDevice* device, const VkDescriptorUpdateTemplateCreateInfoKHR* pCreateInfo) : + MVKConfigurableObject(), _type(pCreateInfo->templateType) { + + for (uint32_t i = 0; i < pCreateInfo->descriptorUpdateEntryCount; i++) + _entries.push_back(pCreateInfo->pDescriptorUpdateEntries[i]); +} + + #pragma mark - #pragma mark Support functions @@ -863,11 +954,12 @@ void mvkUpdateDescriptorSets(uint32_t writeCount, // Perform the write updates for (uint32_t i = 0; i < writeCount; i++) { const VkWriteDescriptorSet* pDescWrite = &pDescriptorWrites[i]; + size_t stride; + const void* pData = getWriteParameters(pDescWrite->descriptorType, pDescWrite->pImageInfo, + pDescWrite->pBufferInfo, pDescWrite->pTexelBufferView, + stride); MVKDescriptorSet* dstSet = (MVKDescriptorSet*)pDescWrite->dstSet; - dstSet->writeDescriptorSets(pDescWrite, - pDescWrite->pImageInfo, - pDescWrite->pBufferInfo, - pDescWrite->pTexelBufferView); + dstSet->writeDescriptorSets(pDescWrite, stride, pData); } // Perform the copy updates by reading bindings from one set and writing to other set. @@ -875,15 +967,37 @@ void mvkUpdateDescriptorSets(uint32_t writeCount, const VkCopyDescriptorSet* pDescCopy = &pDescriptorCopies[i]; uint32_t descCnt = pDescCopy->descriptorCount; + VkDescriptorType descType; VkDescriptorImageInfo imgInfos[descCnt]; VkDescriptorBufferInfo buffInfos[descCnt]; VkBufferView texelBuffInfos[descCnt]; MVKDescriptorSet* srcSet = (MVKDescriptorSet*)pDescCopy->srcSet; - srcSet->readDescriptorSets(pDescCopy, imgInfos, buffInfos, texelBuffInfos); + srcSet->readDescriptorSets(pDescCopy, descType, imgInfos, buffInfos, texelBuffInfos); MVKDescriptorSet* dstSet = (MVKDescriptorSet*)pDescCopy->dstSet; - dstSet->writeDescriptorSets(pDescCopy, imgInfos, buffInfos, texelBuffInfos); + size_t stride; + const void* pData = getWriteParameters(descType, imgInfos, buffInfos, texelBuffInfos, stride); + dstSet->writeDescriptorSets(pDescCopy, stride, pData); + } +} + +/** Updates the resource bindings in the given descriptor set from the specified template. */ +void mvkUpdateDescriptorSetWithTemplate(VkDescriptorSet descriptorSet, + VkDescriptorUpdateTemplateKHR updateTemplate, + const void* pData) { + + MVKDescriptorSet* dstSet = (MVKDescriptorSet*)descriptorSet; + MVKDescriptorUpdateTemplate* pTemplate = (MVKDescriptorUpdateTemplate*)updateTemplate; + + if (pTemplate->getType() != VK_DESCRIPTOR_UPDATE_TEMPLATE_TYPE_DESCRIPTOR_SET_KHR) + return; + + // Perform the updates + for (uint32_t i = 0; i < pTemplate->getNumberOfEntries(); i++) { + const VkDescriptorUpdateTemplateEntryKHR* pEntry = pTemplate->getEntry(i); + const void* pCurData = (const char*)pData + pEntry->offset; + dstSet->writeDescriptorSets(pEntry, pEntry->stride, pCurData); } } diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h index bf19fc84..bc243c0f 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h @@ -53,6 +53,7 @@ class MVKPipeline; class MVKSampler; class MVKDescriptorSetLayout; class MVKDescriptorPool; +class MVKDescriptorUpdateTemplate; class MVKFramebuffer; class MVKRenderPass; class MVKCommandPool; @@ -397,6 +398,11 @@ public: void destroyDescriptorPool(MVKDescriptorPool* mvkDP, const VkAllocationCallbacks* pAllocator); + MVKDescriptorUpdateTemplate* createDescriptorUpdateTemplate(const VkDescriptorUpdateTemplateCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator); + void destroyDescriptorUpdateTemplate(MVKDescriptorUpdateTemplate* mvkDUT, + const VkAllocationCallbacks* pAllocator); + MVKFramebuffer* createFramebuffer(const VkFramebufferCreateInfo* pCreateInfo, const VkAllocationCallbacks* pAllocator); void destroyFramebuffer(MVKFramebuffer* mvkFB, diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index 905a95f0..a2b66c0f 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -1362,6 +1362,17 @@ void MVKDevice::destroyDescriptorPool(MVKDescriptorPool* mvkDP, mvkDP->destroy(); } +MVKDescriptorUpdateTemplate* MVKDevice::createDescriptorUpdateTemplate( + const VkDescriptorUpdateTemplateCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator) { + return new MVKDescriptorUpdateTemplate(this, pCreateInfo); +} + +void MVKDevice::destroyDescriptorUpdateTemplate(MVKDescriptorUpdateTemplate* mvkDUT, + const VkAllocationCallbacks* pAllocator) { + mvkDUT->destroy(); +} + MVKFramebuffer* MVKDevice::createFramebuffer(const VkFramebufferCreateInfo* pCreateInfo, const VkAllocationCallbacks* pAllocator) { return new MVKFramebuffer(this, pCreateInfo); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm index c9761c01..efa4f12f 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.mm @@ -296,6 +296,10 @@ void MVKInstance::initProcAddrs() { ADD_PROC_ADDR(vkGetPhysicalDeviceMemoryProperties2KHR); ADD_PROC_ADDR(vkGetPhysicalDeviceSparseImageFormatProperties2KHR); ADD_PROC_ADDR(vkCmdPushDescriptorSetKHR); + ADD_PROC_ADDR(vkCmdPushDescriptorSetWithTemplateKHR); + ADD_PROC_ADDR(vkCreateDescriptorUpdateTemplateKHR); + ADD_PROC_ADDR(vkDestroyDescriptorUpdateTemplateKHR); + ADD_PROC_ADDR(vkUpdateDescriptorSetWithTemplateKHR); ADD_PROC_ADDR(vkGetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkSetMoltenVKConfigurationMVK); ADD_PROC_ADDR(vkGetPhysicalDeviceMetalFeaturesMVK); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h index 4cd6391c..951b1530 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h @@ -55,6 +55,12 @@ public: std::vector& descriptorWrites, uint32_t set); + /** Updates a descriptor set from a template in a command encoder. */ + void pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + MVKDescriptorUpdateTemplate* descriptorUpdateTemplate, + uint32_t set, + const void* pData); + /** Constructs an instance for the specified device. */ MVKPipelineLayout(MVKDevice* device, const VkPipelineLayoutCreateInfo* pCreateInfo); diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm index 151b840d..e93d2921 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm @@ -65,6 +65,19 @@ void MVKPipelineLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, cmdEncoder->getPushConstants(VK_SHADER_STAGE_COMPUTE_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.computeStage.bufferIndex); } +void MVKPipelineLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder, + MVKDescriptorUpdateTemplate* descUpdateTemplate, + uint32_t set, + const void* pData) { + + _descriptorSetLayouts[set].pushDescriptorSet(cmdEncoder, descUpdateTemplate, + pData, + _dslMTLResourceIndexOffsets[set]); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.vertexStage.bufferIndex); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_FRAGMENT_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.fragmentStage.bufferIndex); + cmdEncoder->getPushConstants(VK_SHADER_STAGE_COMPUTE_BIT)->setMTLBufferIndex(_pushConstantsMTLResourceIndexOffsets.computeStage.bufferIndex); +} + void MVKPipelineLayout::populateShaderConverterContext(SPIRVToMSLConverterContext& context) { context.resourceBindings.clear(); diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp index f66b76a7..213d8af8 100644 --- a/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp @@ -52,6 +52,7 @@ MVK_MAKE_VK_EXT_PROPS(AMD_NEGATIVE_VIEWPORT_HEIGHT); MVK_MAKE_VK_EXT_PROPS(KHR_SHADER_DRAW_PARAMETERS); MVK_MAKE_VK_EXT_PROPS(KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2); MVK_MAKE_VK_EXT_PROPS(KHR_PUSH_DESCRIPTOR); +MVK_MAKE_VK_EXT_PROPS(KHR_DESCRIPTOR_UPDATE_TEMPLATE); // Calls the constructor for a MVKExtension member variable, using the member name and the // portion of the extension name, as uppercase, used in the MVK_MAKE_VK_EXT_PROPS() macro above. @@ -70,6 +71,7 @@ MVKExtensionList::MVKExtensionList(bool enableForPlatform) : MVKExt_CONSTRUCT(vk_KHR_shader_draw_parameters, KHR_SHADER_DRAW_PARAMETERS), MVKExt_CONSTRUCT(vk_KHR_get_physical_device_properties2, KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2), MVKExt_CONSTRUCT(vk_KHR_push_descriptor, KHR_PUSH_DESCRIPTOR) + MVKExt_CONSTRUCT(vk_KHR_descriptor_update_template, KHR_DESCRIPTOR_UPDATE_TEMPLATE) {} bool MVKExtensionList::isEnabled(const char* extnName) const { diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.h b/MoltenVK/MoltenVK/Layers/MVKExtensions.h index 1b2a2e2a..479d71cf 100644 --- a/MoltenVK/MoltenVK/Layers/MVKExtensions.h +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.h @@ -49,6 +49,7 @@ struct MVKExtensionList { MVKExtension vk_KHR_shader_draw_parameters; MVKExtension vk_KHR_get_physical_device_properties2; MVKExtension vk_KHR_push_descriptor; + MVKExtension vk_KHR_descriptor_update_template; }; MVKExtension extensionArray; }; diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm index db0e6f3a..211c92d4 100644 --- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm +++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm @@ -1692,6 +1692,53 @@ MVK_PUBLIC_SYMBOL void vkCmdPushDescriptorSetKHR( mvkCmdPushDescriptorSet(cmdBuff, pipelineBindPoint, layout, set, descriptorWriteCount, pDescriptorWrites); } +MVK_PUBLIC_SYMBOL void vkCmdPushDescriptorSetWithTemplateKHR( + VkCommandBuffer commandBuffer, + VkDescriptorUpdateTemplateKHR descriptorUpdateTemplate, + VkPipelineLayout layout, + uint32_t set, + const void* pData) { + + MVKCommandBuffer* cmdBuff = MVKCommandBuffer::getMVKCommandBuffer(commandBuffer); + mvkCmdPushDescriptorSetWithTemplate(cmdBuff, descriptorUpdateTemplate, layout, set, pData); +} + + +#pragma mark - +#pragma mark VK_KHR_descriptor_update_template extension + +MVK_PUBLIC_SYMBOL VkResult vkCreateDescriptorUpdateTemplateKHR( + VkDevice device, + const VkDescriptorUpdateTemplateCreateInfoKHR* pCreateInfo, + const VkAllocationCallbacks* pAllocator, + VkDescriptorUpdateTemplateKHR* pDescriptorUpdateTemplate) { + + MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); + auto *mvkDUT = mvkDev->createDescriptorUpdateTemplate(pCreateInfo, + pAllocator); + *pDescriptorUpdateTemplate = (VkDescriptorUpdateTemplateKHR)mvkDUT; + return mvkDUT->getConfigurationResult(); +} + +MVK_PUBLIC_SYMBOL void vkDestroyDescriptorUpdateTemplateKHR( + VkDevice device, + VkDescriptorUpdateTemplateKHR descriptorUpdateTemplate, + const VkAllocationCallbacks* pAllocator) { + + if (!descriptorUpdateTemplate) { return; } + MVKDevice* mvkDev = MVKDevice::getMVKDevice(device); + mvkDev->destroyDescriptorUpdateTemplate((MVKDescriptorUpdateTemplate*)descriptorUpdateTemplate, pAllocator); +} + +MVK_PUBLIC_SYMBOL void vkUpdateDescriptorSetWithTemplateKHR( + VkDevice device, + VkDescriptorSet descriptorSet, + VkDescriptorUpdateTemplateKHR descriptorUpdateTemplate, + const void* pData) { + + mvkUpdateDescriptorSetWithTemplate(descriptorSet, descriptorUpdateTemplate, pData); +} + #pragma mark - #pragma mark Loader and Layer ICD interface extension From d557b51433d9621389bf08268fb9dac5ad8b2184 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Sun, 2 Sep 2018 22:23:52 -0500 Subject: [PATCH 08/20] Fix a silly omission that breaks the build. --- MoltenVK/MoltenVK/Layers/MVKExtensions.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp index 213d8af8..d7e2942e 100644 --- a/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp +++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.cpp @@ -70,7 +70,7 @@ MVKExtensionList::MVKExtensionList(bool enableForPlatform) : MVKExt_CONSTRUCT(vk_AMD_negative_viewport_height, AMD_NEGATIVE_VIEWPORT_HEIGHT), MVKExt_CONSTRUCT(vk_KHR_shader_draw_parameters, KHR_SHADER_DRAW_PARAMETERS), MVKExt_CONSTRUCT(vk_KHR_get_physical_device_properties2, KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2), - MVKExt_CONSTRUCT(vk_KHR_push_descriptor, KHR_PUSH_DESCRIPTOR) + MVKExt_CONSTRUCT(vk_KHR_push_descriptor, KHR_PUSH_DESCRIPTOR), MVKExt_CONSTRUCT(vk_KHR_descriptor_update_template, KHR_DESCRIPTOR_UPDATE_TEMPLATE) {} From acf63a16e77d3c0a0ccaddbdb25e0949b03f0072 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 5 Sep 2018 10:05:33 -0500 Subject: [PATCH 09/20] MVKImageView: Create 3D MTLTextureViews for 2D image views of 3D textures. This won't help using these image views from shaders, but it will help attaching them to framebuffers. I've left the warning in `validateImageViewConfig()` in place. --- MoltenVK/MoltenVK/GPUObjects/MVKImage.mm | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index 6b84a1b2..5f738943 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -603,8 +603,13 @@ id MVKImageView::getMTLTexture() { // Creates and returns a retained Metal texture as an // overlay on the Metal texture of the underlying image. id MVKImageView::newMTLTexture() { + MTLTextureType mtlTextureType = _mtlTextureType; + // Fake support for 2D views of 3D textures. + if (_image->getImageType() == VK_IMAGE_TYPE_3D && + (mtlTextureType == MTLTextureType2D || mtlTextureType == MTLTextureType2DArray)) + mtlTextureType = MTLTextureType3D; return [_image->getMTLTexture() newTextureViewWithPixelFormat: _mtlPixelFormat - textureType: _mtlTextureType + textureType: mtlTextureType levels: NSMakeRange(_subresourceRange.baseMipLevel, _subresourceRange.levelCount) slices: NSMakeRange(_subresourceRange.baseArrayLayer, _subresourceRange.layerCount)]; // retained } From 7fd270edb60fe0d2849cba7eaecbb5f994eea63c Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 5 Sep 2018 11:11:40 -0500 Subject: [PATCH 10/20] MVKDevice: Fix a segfault walking unknown extension structs. The problem is that on 64-bit platforms (i.e. every platform we support) there will be padding between the `sType` and `pNext` members of any extensible Vulkan structure, because `sType` is only 4 bytes while `pNext` is 8 (and needs 8-byte alignment). Should fix a segfault running `vulkaninfo`. --- MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm index a2b66c0f..63da27e8 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm @@ -87,7 +87,7 @@ void MVKPhysicalDevice::getProperties(VkPhysicalDeviceProperties2KHR* properties break; } default: - next = *(VkStructureType**)(next+1); + next = (VkStructureType*)((VkPhysicalDeviceProperties2KHR*)next)->pNext; break; } } From f1b62ed297415f90b3c28581780643b4209cdc60 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 5 Sep 2018 10:33:57 -0500 Subject: [PATCH 11/20] Don't fail creation of the ImageView. --- MoltenVK/MoltenVK/GPUObjects/MVKImage.mm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index 5f738943..1fac92f2 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -645,7 +645,7 @@ void MVKImageView::validateImageViewConfig(const VkImageViewCreateInfo* pCreateI // VK_KHR_maintenance1 supports taking 2D image views of 3D slices. No dice in Metal. if ((viewType == VK_IMAGE_VIEW_TYPE_2D || viewType == VK_IMAGE_VIEW_TYPE_2D_ARRAY) && (imgType == VK_IMAGE_TYPE_3D)) { - setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): Metal does not support creating a 2D view on a 3D image.")); + mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): Metal does not support creating a 2D view on a 3D image."); } } From 78a59a3e37b89ab138e9b96bae15293c18cd2bcb Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 5 Sep 2018 11:53:49 -0500 Subject: [PATCH 12/20] Try harder to avoid creating texture views with 3D textures. Also, warn only when the texture may be used for a purpose other than as a color attachment. Fail outright if it won't ever be used for that purpose, or if the view doesn't cover the entirety of the volume. --- MoltenVK/MoltenVK/GPUObjects/MVKImage.mm | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index 1fac92f2..f411483e 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -640,12 +640,19 @@ MVKImageView::MVKImageView(MVKDevice* device, const VkImageViewCreateInfo* pCrea // Validate whether the image view configuration can be supported void MVKImageView::validateImageViewConfig(const VkImageViewCreateInfo* pCreateInfo) { - VkImageType imgType = ((MVKImage*)pCreateInfo->image)->getImageType(); + MVKImage* image = (MVKImage*)pCreateInfo->image; + VkImageType imgType = image->getImageType(); VkImageViewType viewType = pCreateInfo->viewType; // VK_KHR_maintenance1 supports taking 2D image views of 3D slices. No dice in Metal. if ((viewType == VK_IMAGE_VIEW_TYPE_2D || viewType == VK_IMAGE_VIEW_TYPE_2D_ARRAY) && (imgType == VK_IMAGE_TYPE_3D)) { - mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): Metal does not support creating a 2D view on a 3D image."); + if (pCreateInfo->subresourceRange.layerCount != image->_extent.depth) { + setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): Metal does not support views on a subset of a 3D texture.")); + } else if (!(image->_usage & VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT)) { + setConfigurationResult(mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): 2D views on 3D images are only supported for color attachments.")); + } else if (image->_usage & ~VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT) { + mvkNotifyErrorWithText(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImageView(): 2D views on 3D images are only supported for color attachments."); + } } } @@ -768,10 +775,12 @@ void MVKImageView::initMTLTextureViewSupport() { _useMTLTextureView = _image->_canSupportMTLTextureView; // If the view is identical to underlying image, don't bother using a Metal view + bool is3D = _image->_mtlTextureType == MTLTextureType3D; if (_mtlPixelFormat == _image->_mtlPixelFormat && - _mtlTextureType == _image->_mtlTextureType && + (_mtlTextureType == _image->_mtlTextureType || + (_mtlTextureType == MTLTextureType2D || _mtlTextureType == MTLTextureType2DArray) && is3D) && _subresourceRange.levelCount == _image->_mipLevels && - _subresourceRange.layerCount == _image->_arrayLayers) { + _subresourceRange.layerCount == (is3D ? _image->_extent.depth : _image->_arrayLayers)) { _useMTLTextureView = false; } } From a0d567da92349fd6c61cce0032cd7e1676149d58 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Thu, 6 Sep 2018 20:04:41 -0500 Subject: [PATCH 13/20] Fix Clang warning about order of operations. --- MoltenVK/MoltenVK/GPUObjects/MVKImage.mm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm index f411483e..3fa0f96f 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm @@ -778,7 +778,7 @@ void MVKImageView::initMTLTextureViewSupport() { bool is3D = _image->_mtlTextureType == MTLTextureType3D; if (_mtlPixelFormat == _image->_mtlPixelFormat && (_mtlTextureType == _image->_mtlTextureType || - (_mtlTextureType == MTLTextureType2D || _mtlTextureType == MTLTextureType2DArray) && is3D) && + ((_mtlTextureType == MTLTextureType2D || _mtlTextureType == MTLTextureType2DArray) && is3D)) && _subresourceRange.levelCount == _image->_mipLevels && _subresourceRange.layerCount == (is3D ? _image->_extent.depth : _image->_arrayLayers)) { _useMTLTextureView = false; From e721dd6e2cce569ae80d08d137534d32e7c1ff40 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Fri, 7 Sep 2018 11:46:58 -0400 Subject: [PATCH 14/20] Update build process. Allow building and packaging MoltenVK for of only iOS or only macOS. Move packaging scripts out of Xcode projects and into script files. --- MoltenVK/MoltenVK.xcodeproj/project.pbxproj | 4 +- MoltenVK/scripts/create_dylib_ios.sh | 33 +++ MoltenVK/scripts/create_dylib_macos.sh | 24 ++ MoltenVKPackaging.xcodeproj/project.pbxproj | 238 +++++++++++++++--- ...heme => MoltenVK Package (Debug).xcscheme} | 0 ...me => MoltenVK Package (Release).xcscheme} | 0 ...ltenVK Package (iOS only) (Debug).xcscheme | 80 ++++++ ...enVK Package (iOS only) (Release).xcscheme | 80 ++++++ ...enVK Package (macOS only) (Debug).xcscheme | 80 ++++++ ...VK Package (macOS only) (Release).xcscheme | 80 ++++++ Scripts/package_all.sh | 7 + Scripts/package_docs.sh | 12 + Scripts/package_moltenvk.sh | 43 ++++ Scripts/package_shader_converter.sh | 73 ++++++ Scripts/update_latest.sh | 14 ++ 15 files changed, 729 insertions(+), 39 deletions(-) create mode 100755 MoltenVK/scripts/create_dylib_ios.sh create mode 100755 MoltenVK/scripts/create_dylib_macos.sh rename MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/{MoltenVK (Debug).xcscheme => MoltenVK Package (Debug).xcscheme} (100%) rename MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/{MoltenVK (Release).xcscheme => MoltenVK Package (Release).xcscheme} (100%) create mode 100644 MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Debug).xcscheme create mode 100644 MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Release).xcscheme create mode 100644 MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Debug).xcscheme create mode 100644 MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Release).xcscheme create mode 100755 Scripts/package_all.sh create mode 100755 Scripts/package_docs.sh create mode 100755 Scripts/package_moltenvk.sh create mode 100755 Scripts/package_shader_converter.sh create mode 100755 Scripts/update_latest.sh diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj index 2f87edd4..49c83e5e 100644 --- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj +++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj @@ -831,7 +831,7 @@ ); runOnlyForDeploymentPostprocessing = 0; shellPath = /bin/sh; - shellScript = "set -e\n\nexport MVK_PROD_NAME=\"MoltenVK\"\nexport MVK_DYLIB_NAME=\"lib${MVK_PROD_NAME}.dylib\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nexport MVK_SYS_FWK_DIR=\"${SDK_DIR}/System/Library/Frameworks\"\nexport MVK_USR_LIB_DIR=\"${SDK_DIR}/usr/lib\"\n\nclang \\\n-dynamiclib \\\n-arch x86_64 \\\n-mmacosx-version-min=${MACOSX_DEPLOYMENT_TARGET} \\\n-compatibility_version 1.0.0 -current_version 1.0.0 \\\n-install_name \"@rpath/${MVK_DYLIB_NAME}\" \\\n-Wno-incompatible-sysroot \\\n-isysroot ${SDK_DIR} \\\n-iframework ${MVK_SYS_FWK_DIR} \\\n-framework Metal -framework IOSurface -framework IOKit -framework QuartzCore -framework Foundation \\\n--library-directory ${MVK_USR_LIB_DIR} \\\n-lSystem -lc++ \\\n-o \"${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}\" \\\n-force_load \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/${MVK_PROD_NAME}\"\n"; + shellScript = "${SRCROOT}/scripts/create_dylib_macos.sh"; }; A9731FAD1EDDAE39006B7298 /* Create Dynamic Library */ = { isa = PBXShellScriptBuildPhase; @@ -845,7 +845,7 @@ ); runOnlyForDeploymentPostprocessing = 0; shellPath = /bin/sh; - shellScript = "set -e\n\nexport MVK_PROD_NAME=\"MoltenVK\"\nexport MVK_DYLIB_NAME=\"lib${MVK_PROD_NAME}.dylib\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nexport MVK_SYS_FWK_DIR=\"${SDK_DIR}/System/Library/Frameworks\"\nexport MVK_USR_LIB_DIR=\"${SDK_DIR}/usr/lib\"\n\n# Do not link to IOSurface if deploying to iOS versions below 11.0, doing so will\n# link IOSurface as a private framework, which will trigger App Store rejection.\nif [ $(echo \"${IPHONEOS_DEPLOYMENT_TARGET} >= 11.0\" | bc) -eq 1 ]\nthen\n export MVK_IOSURFACE_FWK=\"-framework IOSurface\"\nelse\n export MVK_IOSURFACE_FWK=\"\"\nfi\n\nclang \\\n-dynamiclib \\\n-arch arm64 \\\n-mios-version-min=${IPHONEOS_DEPLOYMENT_TARGET} \\\n-compatibility_version 1.0.0 -current_version 1.0.0 \\\n-install_name \"@rpath/${MVK_DYLIB_NAME}\" \\\n-Wno-incompatible-sysroot \\\n-isysroot ${SDK_DIR} \\\n-iframework ${MVK_SYS_FWK_DIR} \\\n-framework Metal ${MVK_IOSURFACE_FWK} -framework UIKit -framework QuartzCore -framework Foundation \\\n--library-directory ${MVK_USR_LIB_DIR} \\\n-lSystem -lc++ \\\n-o \"${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}\" \\\n-force_load \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/${MVK_PROD_NAME}\"\n"; + shellScript = "${SRCROOT}/scripts/create_dylib_ios.sh"; }; /* End PBXShellScriptBuildPhase section */ diff --git a/MoltenVK/scripts/create_dylib_ios.sh b/MoltenVK/scripts/create_dylib_ios.sh new file mode 100755 index 00000000..544ea793 --- /dev/null +++ b/MoltenVK/scripts/create_dylib_ios.sh @@ -0,0 +1,33 @@ +#!/bin/bash + +set -e + +export MVK_PROD_NAME="MoltenVK" +export MVK_DYLIB_NAME="lib${MVK_PROD_NAME}.dylib" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +export MVK_SYS_FWK_DIR="${SDK_DIR}/System/Library/Frameworks" +export MVK_USR_LIB_DIR="${SDK_DIR}/usr/lib" + +# Do not link to IOSurface if deploying to iOS versions below 11.0, doing so will +# link IOSurface as a private framework, which will trigger App Store rejection. +if [ $(echo "${IPHONEOS_DEPLOYMENT_TARGET} >= 11.0" | bc) -eq 1 ] +then + export MVK_IOSURFACE_FWK="-framework IOSurface" +else + export MVK_IOSURFACE_FWK="" +fi + +clang \ +-dynamiclib \ +-arch arm64 \ +-mios-version-min=${IPHONEOS_DEPLOYMENT_TARGET} \ +-compatibility_version 1.0.0 -current_version 1.0.0 \ +-install_name "@rpath/${MVK_DYLIB_NAME}" \ +-Wno-incompatible-sysroot \ +-isysroot ${SDK_DIR} \ +-iframework ${MVK_SYS_FWK_DIR} \ +-framework Metal ${MVK_IOSURFACE_FWK} -framework UIKit -framework QuartzCore -framework Foundation \ +--library-directory ${MVK_USR_LIB_DIR} \ +-lSystem -lc++ \ +-o "${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}" \ +-force_load "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/${MVK_PROD_NAME}" diff --git a/MoltenVK/scripts/create_dylib_macos.sh b/MoltenVK/scripts/create_dylib_macos.sh new file mode 100755 index 00000000..3e6b33bf --- /dev/null +++ b/MoltenVK/scripts/create_dylib_macos.sh @@ -0,0 +1,24 @@ +#!/bin/bash + +set -e + +export MVK_PROD_NAME="MoltenVK" +export MVK_DYLIB_NAME="lib${MVK_PROD_NAME}.dylib" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +export MVK_SYS_FWK_DIR="${SDK_DIR}/System/Library/Frameworks" +export MVK_USR_LIB_DIR="${SDK_DIR}/usr/lib" + +clang \ +-dynamiclib \ +-arch x86_64 \ +-mmacosx-version-min=${MACOSX_DEPLOYMENT_TARGET} \ +-compatibility_version 1.0.0 -current_version 1.0.0 \ +-install_name "@rpath/${MVK_DYLIB_NAME}" \ +-Wno-incompatible-sysroot \ +-isysroot ${SDK_DIR} \ +-iframework ${MVK_SYS_FWK_DIR} \ +-framework Metal -framework IOSurface -framework IOKit -framework QuartzCore -framework Foundation \ +--library-directory ${MVK_USR_LIB_DIR} \ +-lSystem -lc++ \ +-o "${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}" \ +-force_load "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/${MVK_PROD_NAME}" diff --git a/MoltenVKPackaging.xcodeproj/project.pbxproj b/MoltenVKPackaging.xcodeproj/project.pbxproj index 3fea58c5..e5d0feae 100644 --- a/MoltenVKPackaging.xcodeproj/project.pbxproj +++ b/MoltenVKPackaging.xcodeproj/project.pbxproj @@ -7,14 +7,40 @@ objects = { /* Begin PBXAggregateTarget section */ + A975D5782140585200D4834F /* MoltenVK-iOS */ = { + isa = PBXAggregateTarget; + buildConfigurationList = A975D5882140585200D4834F /* Build configuration list for PBXAggregateTarget "MoltenVK-iOS" */; + buildPhases = ( + A975D5872140585200D4834F /* Package MoltenVK */, + ); + dependencies = ( + A975D5792140585200D4834F /* PBXTargetDependency */, + A975D57D2140585200D4834F /* PBXTargetDependency */, + A975D5812140585200D4834F /* PBXTargetDependency */, + ); + name = "MoltenVK-iOS"; + productName = Package; + }; + A975D58B2140586700D4834F /* MoltenVK-macOS */ = { + isa = PBXAggregateTarget; + buildConfigurationList = A975D59B2140586700D4834F /* Build configuration list for PBXAggregateTarget "MoltenVK-macOS" */; + buildPhases = ( + A975D59A2140586700D4834F /* Package MoltenVK */, + ); + dependencies = ( + A975D58E2140586700D4834F /* PBXTargetDependency */, + A975D5922140586700D4834F /* PBXTargetDependency */, + A975D5962140586700D4834F /* PBXTargetDependency */, + A975D5982140586700D4834F /* PBXTargetDependency */, + ); + name = "MoltenVK-macOS"; + productName = Package; + }; A9FEADBC1F3517480010240E /* MoltenVK */ = { isa = PBXAggregateTarget; buildConfigurationList = A9FEADDC1F3517480010240E /* Build configuration list for PBXAggregateTarget "MoltenVK" */; buildPhases = ( A9FEADD61F3517480010240E /* Package MoltenVK */, - A9FEADD71F3517480010240E /* Package MoltenVKShaderConverter */, - A9A5D8A61FB3AABA00F20475 /* Package Docs */, - A9FEADDB1F3517480010240E /* Update Latest */, ); dependencies = ( A9FEADBD1F3517480010240E /* PBXTargetDependency */, @@ -45,6 +71,55 @@ remoteGlobalIDString = A9CBEE011B6299D800E45FDC; remoteInfo = "MoltenVK-macOS"; }; + A975D57A2140585200D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB3EE1CE0F72500FBC835 /* MoltenVK.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A9B8EE091A98D796009C5A02; + remoteInfo = "MoltenVK-iOS"; + }; + A975D57E2140585200D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A93903B81C57E9D700FE90DC; + remoteInfo = "MoltenVKSPIRVToMSLConverter-iOS"; + }; + A975D5822140585200D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A937472B1A9A8B2900F29B34; + remoteInfo = "MoltenVKGLSLToSPIRVConverter-iOS"; + }; + A975D58F2140586700D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB3EE1CE0F72500FBC835 /* MoltenVK.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A9CBED861B6299D800E45FDC; + remoteInfo = "MoltenVK-macOS"; + }; + A975D5932140586700D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A93903C01C57E9ED00FE90DC; + remoteInfo = "MoltenVKSPIRVToMSLConverter-macOS"; + }; + A975D5972140586700D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A93747701A9A98D000F29B34; + remoteInfo = "MoltenVKGLSLToSPIRVConverter-macOS"; + }; + A975D5992140586700D4834F /* PBXContainerItemProxy */ = { + isa = PBXContainerItemProxy; + containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; + proxyType = 1; + remoteGlobalIDString = A9092A8C1A81717B00051823; + remoteInfo = MoltenVKShaderConverter; + }; A981498A1FB6B566005F00B4 /* PBXContainerItemProxy */ = { isa = PBXContainerItemProxy; containerPortal = A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */; @@ -138,6 +213,13 @@ A92DB3EE1CE0F72500FBC835 /* MoltenVK.xcodeproj */ = {isa = PBXFileReference; lastKnownFileType = "wrapper.pb-project"; name = MoltenVK.xcodeproj; path = MoltenVK/MoltenVK.xcodeproj; sourceTree = ""; }; A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */ = {isa = PBXFileReference; lastKnownFileType = "wrapper.pb-project"; name = MoltenVKShaderConverter.xcodeproj; path = MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj; sourceTree = ""; }; A943100220546CDD00F5CF87 /* fetchDependencies */ = {isa = PBXFileReference; lastKnownFileType = text; name = fetchDependencies; path = ../fetchDependencies; sourceTree = ""; }; + A975D55C213F25D700D4834F /* create_dylib_ios.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; name = create_dylib_ios.sh; path = MoltenVK/scripts/create_dylib_ios.sh; sourceTree = SOURCE_ROOT; }; + A975D55D213F266000D4834F /* create_dylib_macos.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; name = create_dylib_macos.sh; path = MoltenVK/scripts/create_dylib_macos.sh; sourceTree = SOURCE_ROOT; }; + A975D561213F299500D4834F /* update_latest.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = update_latest.sh; sourceTree = ""; }; + A975D562213F2B7700D4834F /* package_docs.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = package_docs.sh; sourceTree = ""; }; + A975D566213F2DAA00D4834F /* package_shader_converter.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = package_shader_converter.sh; sourceTree = ""; }; + A975D573214050AB00D4834F /* package_moltenvk.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = package_moltenvk.sh; sourceTree = ""; }; + A975D5742140567B00D4834F /* package_all.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = package_all.sh; sourceTree = ""; }; A98149E51FB78829005F00B4 /* MoltenVK_Runtime_UserGuide.md */ = {isa = PBXFileReference; lastKnownFileType = net.daringfireball.markdown; name = MoltenVK_Runtime_UserGuide.md; path = Docs/MoltenVK_Runtime_UserGuide.md; sourceTree = ""; }; A9AD67D12054E2D700ED3C08 /* VulkanSamples_repo_revision */ = {isa = PBXFileReference; lastKnownFileType = text; path = VulkanSamples_repo_revision; sourceTree = ""; }; A9AD67D32054E2D700ED3C08 /* SPIRV-Cross_repo_revision */ = {isa = PBXFileReference; lastKnownFileType = text; path = "SPIRV-Cross_repo_revision"; sourceTree = ""; }; @@ -154,6 +236,7 @@ children = ( A92DB3EE1CE0F72500FBC835 /* MoltenVK.xcodeproj */, A92DB40E1CE0F89600FBC835 /* MoltenVKShaderConverter.xcodeproj */, + A975D55B213F25AD00D4834F /* Scripts */, A92DB3E11CE0F34500FBC835 /* Docs */, A939A6FB1F5479D0006ACA0C /* External */, ); @@ -195,6 +278,20 @@ path = ExternalRevisions; sourceTree = ""; }; + A975D55B213F25AD00D4834F /* Scripts */ = { + isa = PBXGroup; + children = ( + A975D55C213F25D700D4834F /* create_dylib_ios.sh */, + A975D55D213F266000D4834F /* create_dylib_macos.sh */, + A975D5742140567B00D4834F /* package_all.sh */, + A975D562213F2B7700D4834F /* package_docs.sh */, + A975D573214050AB00D4834F /* package_moltenvk.sh */, + A975D566213F2DAA00D4834F /* package_shader_converter.sh */, + A975D561213F299500D4834F /* update_latest.sh */, + ); + path = Scripts; + sourceTree = ""; + }; A98149741FB6B565005F00B4 /* Products */ = { isa = PBXGroup; children = ( @@ -242,6 +339,8 @@ projectRoot = ""; targets = ( A9FEADBC1F3517480010240E /* MoltenVK */, + A975D5782140585200D4834F /* MoltenVK-iOS */, + A975D58B2140586700D4834F /* MoltenVK-macOS */, ); }; /* End PBXProject section */ @@ -299,19 +398,33 @@ /* End PBXReferenceProxy section */ /* Begin PBXShellScriptBuildPhase section */ - A9A5D8A61FB3AABA00F20475 /* Package Docs */ = { + A975D5872140585200D4834F /* Package MoltenVK */ = { isa = PBXShellScriptBuildPhase; buildActionMask = 2147483647; files = ( ); inputPaths = ( ); - name = "Package Docs"; + name = "Package MoltenVK"; outputPaths = ( ); runOnlyForDeploymentPostprocessing = 0; shellPath = /bin/sh; - shellScript = "set -e\n\n# Package folder\nexport MVK_WKSPC_PATH=\"${PROJECT_DIR}\"\nexport MVK_PKG_LOCN=\"${MVK_WKSPC_PATH}/Package\"\nexport MVK_PKG_CONFIG_PATH=\"${MVK_PKG_LOCN}/${CONFIGURATION}\"\n\n# Copy the docs. Allow silent fail if a symlinked doc is not built.\ncp -a \"${MVK_WKSPC_PATH}/LICENSE\" \"${MVK_PKG_CONFIG_PATH}\"\ncp -pRLf \"${MVK_WKSPC_PATH}/Docs\" \"${MVK_PKG_CONFIG_PATH}\" 2> /dev/null || true\n"; + shellScript = "${SRCROOT}/Scripts/package_all.sh"; + }; + A975D59A2140586700D4834F /* Package MoltenVK */ = { + isa = PBXShellScriptBuildPhase; + buildActionMask = 2147483647; + files = ( + ); + inputPaths = ( + ); + name = "Package MoltenVK"; + outputPaths = ( + ); + runOnlyForDeploymentPostprocessing = 0; + shellPath = /bin/sh; + shellScript = "${SRCROOT}/Scripts/package_all.sh"; }; A9FEADD61F3517480010240E /* Package MoltenVK */ = { isa = PBXShellScriptBuildPhase; @@ -325,39 +438,46 @@ ); runOnlyForDeploymentPostprocessing = 0; shellPath = /bin/sh; - shellScript = "set -e\n\n# Package folder\nexport MVK_PROD_NAME=\"MoltenVK\"\nexport MVK_DYLIB_NAME=\"lib${MVK_PROD_NAME}.dylib\"\nexport MVK_ICD_NAME=\"${MVK_PROD_NAME}_icd.json\"\nexport MVK_WKSPC_PATH=\"${PROJECT_DIR}\"\nexport MVK_PROD_PROJ_PATH=\"${MVK_WKSPC_PATH}/${MVK_PROD_NAME}\"\nexport MVK_PKG_LOCN=\"${MVK_WKSPC_PATH}/Package\"\nexport MVK_PKG_CONFIG_PATH=\"${MVK_PKG_LOCN}/${CONFIGURATION}\"\nexport MVK_PKG_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/${MVK_PROD_NAME}\"\n\n# Remove the product folder\nrm -rf \"${MVK_PKG_PROD_PATH}\"\n\n# Remove and replace the existing macOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_PROD_PATH}/macOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}\" \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_PROD_PROJ_PATH}/icd/${MVK_ICD_NAME}\" \"${MVK_OS_PROD_PATH}\"\n\n# Remove and replace the existing iOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_PROD_PATH}/iOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}/../${CONFIGURATION}-iphoneos\"\nrm -rf \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}\" \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_PROD_PROJ_PATH}/icd/${MVK_ICD_NAME}\" \"${MVK_OS_PROD_PATH}\"\n\n# Remove and replace header include folder\nrm -rf \"${MVK_PKG_PROD_PATH}/include\"\ncp -pRL \"${MVK_PROD_PROJ_PATH}/include\" \"${MVK_PKG_PROD_PATH}\""; - }; - A9FEADD71F3517480010240E /* Package MoltenVKShaderConverter */ = { - isa = PBXShellScriptBuildPhase; - buildActionMask = 2147483647; - files = ( - ); - inputPaths = ( - ); - name = "Package MoltenVKShaderConverter"; - outputPaths = ( - ); - runOnlyForDeploymentPostprocessing = 0; - shellPath = /bin/sh; - shellScript = "set -e\n\n# Package folder\nexport MVK_PROD_BASE_NAME=\"MoltenVKShaderConverter\"\nexport MVK_WKSPC_PATH=\"${PROJECT_DIR}\"\nexport MVK_PKG_LOCN=\"${MVK_WKSPC_PATH}/Package\"\n\n# Remove the base product folder\nrm -rf \"${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}\"\n\n#-----------------------------------\n# MoltenVKGLSLToSPIRVConverter\nexport MVK_PROD_NAME=\"MoltenVKGLSLToSPIRVConverter\"\nexport MVK_PKG_CONFIG_PATH=\"${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}/${MVK_PROD_NAME}\"\n\n# Remove and replace the existing macOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/macOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\n\n# Remove and replace the existing iOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/iOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}/../${CONFIGURATION}-iphoneos\"\nrm -rf \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\n\n#-----------------------------------\n# MoltenVKSPIRVToMSLConverter\nexport MVK_PROD_NAME=\"MoltenVKSPIRVToMSLConverter\"\nexport MVK_PKG_CONFIG_PATH=\"${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}/${MVK_PROD_NAME}\"\n\n# Remove and replace the existing macOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/macOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\n\n# Remove and replace the existing iOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/iOS\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}/../${CONFIGURATION}-iphoneos\"\nrm -rf \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework\" \"${MVK_OS_PROD_PATH}\"\n\n#-----------------------------------\n# MoltenVKShaderConverter Tool\nexport MVK_PROD_NAME=\"MoltenVKShaderConverter\"\nexport MVK_PKG_CONFIG_PATH=\"${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}\"\n\n# Remove and replace the existing macOS framework folder and copy framework into it\nexport MVK_OS_PROD_PATH=\"${MVK_PKG_CONFIG_PATH}/Tools\"\nexport MVK_BUILT_PROD_PATH=\"${BUILT_PRODUCTS_DIR}\"\nrm -rf \"${MVK_OS_PROD_PATH}\"\nmkdir -p \"${MVK_OS_PROD_PATH}\"\ncp -a \"${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}\" \"${MVK_OS_PROD_PATH}\"\n"; - }; - A9FEADDB1F3517480010240E /* Update Latest */ = { - isa = PBXShellScriptBuildPhase; - buildActionMask = 2147483647; - files = ( - ); - inputPaths = ( - ); - name = "Update Latest"; - outputPaths = ( - ); - runOnlyForDeploymentPostprocessing = 0; - shellPath = /bin/sh; - shellScript = "set -e\n\n# Package folder\nexport MVK_WKSPC_LOCN=\"${PROJECT_DIR}\"\nexport MVK_PKG_LOCN=\"${MVK_WKSPC_LOCN}/Package\"\n\n# Configuration package folder location\nexport MVK_PKG_CONFIG_LOCN=\"${CONFIGURATION}\"\nexport MVK_PKG_LATEST_LOCN=\"Latest\"\n\n# Assign symlink from Latest\nln -sfn \"${MVK_PKG_LOCN}/${MVK_PKG_CONFIG_LOCN}\" \"${MVK_PKG_LOCN}/${MVK_PKG_LATEST_LOCN}\""; + shellScript = "${SRCROOT}/Scripts/package_all.sh"; }; /* End PBXShellScriptBuildPhase section */ /* Begin PBXTargetDependency section */ + A975D5792140585200D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVK-iOS"; + targetProxy = A975D57A2140585200D4834F /* PBXContainerItemProxy */; + }; + A975D57D2140585200D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVKSPIRVToMSLConverter-iOS"; + targetProxy = A975D57E2140585200D4834F /* PBXContainerItemProxy */; + }; + A975D5812140585200D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVKGLSLToSPIRVConverter-iOS"; + targetProxy = A975D5822140585200D4834F /* PBXContainerItemProxy */; + }; + A975D58E2140586700D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVK-macOS"; + targetProxy = A975D58F2140586700D4834F /* PBXContainerItemProxy */; + }; + A975D5922140586700D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVKSPIRVToMSLConverter-macOS"; + targetProxy = A975D5932140586700D4834F /* PBXContainerItemProxy */; + }; + A975D5962140586700D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = "MoltenVKGLSLToSPIRVConverter-macOS"; + targetProxy = A975D5972140586700D4834F /* PBXContainerItemProxy */; + }; + A975D5982140586700D4834F /* PBXTargetDependency */ = { + isa = PBXTargetDependency; + name = MoltenVKShaderConverter; + targetProxy = A975D5992140586700D4834F /* PBXContainerItemProxy */; + }; A98149CB1FB7689D005F00B4 /* PBXTargetDependency */ = { isa = PBXTargetDependency; name = MoltenVKShaderConverter; @@ -408,10 +528,37 @@ }; name = Release; }; + A975D5892140585200D4834F /* Debug */ = { + isa = XCBuildConfiguration; + buildSettings = { + PRODUCT_NAME = "$(TARGET_NAME)"; + }; + name = Debug; + }; + A975D58A2140585200D4834F /* Release */ = { + isa = XCBuildConfiguration; + buildSettings = { + PRODUCT_NAME = "$(TARGET_NAME)"; + }; + name = Release; + }; + A975D59C2140586700D4834F /* Debug */ = { + isa = XCBuildConfiguration; + buildSettings = { + PRODUCT_NAME = "$(TARGET_NAME)"; + }; + name = Debug; + }; + A975D59D2140586700D4834F /* Release */ = { + isa = XCBuildConfiguration; + buildSettings = { + PRODUCT_NAME = "$(TARGET_NAME)"; + }; + name = Release; + }; A9FEADDD1F3517480010240E /* Debug */ = { isa = XCBuildConfiguration; buildSettings = { - GCC_PREPROCESSOR_DEFINITIONS = "$(GCC_PREPROCESSOR_DEFINITIONS)"; PRODUCT_NAME = "$(TARGET_NAME)"; }; name = Debug; @@ -419,7 +566,6 @@ A9FEADDE1F3517480010240E /* Release */ = { isa = XCBuildConfiguration; buildSettings = { - GCC_PREPROCESSOR_DEFINITIONS = "$(GCC_PREPROCESSOR_DEFINITIONS)"; PRODUCT_NAME = "$(TARGET_NAME)"; }; name = Release; @@ -436,6 +582,24 @@ defaultConfigurationIsVisible = 0; defaultConfigurationName = Release; }; + A975D5882140585200D4834F /* Build configuration list for PBXAggregateTarget "MoltenVK-iOS" */ = { + isa = XCConfigurationList; + buildConfigurations = ( + A975D5892140585200D4834F /* Debug */, + A975D58A2140585200D4834F /* Release */, + ); + defaultConfigurationIsVisible = 0; + defaultConfigurationName = Release; + }; + A975D59B2140586700D4834F /* Build configuration list for PBXAggregateTarget "MoltenVK-macOS" */ = { + isa = XCConfigurationList; + buildConfigurations = ( + A975D59C2140586700D4834F /* Debug */, + A975D59D2140586700D4834F /* Release */, + ); + defaultConfigurationIsVisible = 0; + defaultConfigurationName = Release; + }; A9FEADDC1F3517480010240E /* Build configuration list for PBXAggregateTarget "MoltenVK" */ = { isa = XCConfigurationList; buildConfigurations = ( diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK (Debug).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (Debug).xcscheme similarity index 100% rename from MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK (Debug).xcscheme rename to MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (Debug).xcscheme diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK (Release).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (Release).xcscheme similarity index 100% rename from MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK (Release).xcscheme rename to MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (Release).xcscheme diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Debug).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Debug).xcscheme new file mode 100644 index 00000000..eb545a4a --- /dev/null +++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Debug).xcscheme @@ -0,0 +1,80 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Release).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Release).xcscheme new file mode 100644 index 00000000..418cdcc0 --- /dev/null +++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (iOS only) (Release).xcscheme @@ -0,0 +1,80 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Debug).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Debug).xcscheme new file mode 100644 index 00000000..5396b003 --- /dev/null +++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Debug).xcscheme @@ -0,0 +1,80 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Release).xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Release).xcscheme new file mode 100644 index 00000000..9b180e90 --- /dev/null +++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package (macOS only) (Release).xcscheme @@ -0,0 +1,80 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/Scripts/package_all.sh b/Scripts/package_all.sh new file mode 100755 index 00000000..82f4b6d6 --- /dev/null +++ b/Scripts/package_all.sh @@ -0,0 +1,7 @@ +#!/bin/bash + +${SRCROOT}/Scripts/package_moltenvk.sh +${SRCROOT}/Scripts/package_shader_converter.sh +${SRCROOT}/Scripts/package_docs.sh +${SRCROOT}/Scripts/update_latest.sh + diff --git a/Scripts/package_docs.sh b/Scripts/package_docs.sh new file mode 100755 index 00000000..2d634d4e --- /dev/null +++ b/Scripts/package_docs.sh @@ -0,0 +1,12 @@ +#!/bin/bash + +set -e + +# Package folder +export MVK_WKSPC_PATH="${PROJECT_DIR}" +export MVK_PKG_LOCN="${MVK_WKSPC_PATH}/Package" +export MVK_PKG_CONFIG_PATH="${MVK_PKG_LOCN}/${CONFIGURATION}" + +# Copy the docs. Allow silent fail if a symlinked doc is not built. +cp -a "${MVK_WKSPC_PATH}/LICENSE" "${MVK_PKG_CONFIG_PATH}" +cp -pRLf "${MVK_WKSPC_PATH}/Docs" "${MVK_PKG_CONFIG_PATH}" 2> /dev/null || true diff --git a/Scripts/package_moltenvk.sh b/Scripts/package_moltenvk.sh new file mode 100755 index 00000000..c22562ef --- /dev/null +++ b/Scripts/package_moltenvk.sh @@ -0,0 +1,43 @@ +#!/bin/bash + +set -e + +# Package folder +export MVK_PROD_NAME="MoltenVK" +export MVK_DYLIB_NAME="lib${MVK_PROD_NAME}.dylib" +export MVK_ICD_NAME="${MVK_PROD_NAME}_icd.json" +export MVK_WKSPC_PATH="${PROJECT_DIR}" +export MVK_PROD_PROJ_PATH="${MVK_WKSPC_PATH}/${MVK_PROD_NAME}" +export MVK_PKG_LOCN="${MVK_WKSPC_PATH}/Package" +export MVK_PKG_CONFIG_PATH="${MVK_PKG_LOCN}/${CONFIGURATION}" +export MVK_PKG_PROD_PATH="${MVK_PKG_CONFIG_PATH}/${MVK_PROD_NAME}" + +# Remove the product folder +rm -rf "${MVK_PKG_PROD_PATH}" + +# Remove and replace the existing macOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_PROD_PATH}/macOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}" "${MVK_OS_PROD_PATH}" + cp -a "${MVK_PROD_PROJ_PATH}/icd/${MVK_ICD_NAME}" "${MVK_OS_PROD_PATH}" +fi + +# Remove and replace the existing iOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_PROD_PATH}/iOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}-iphoneos" +rm -rf "${MVK_OS_PROD_PATH}" +echo MVK_BUILT_PROD_PATH = "${MVK_BUILT_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + rm -rf "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature" + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_DYLIB_NAME}" "${MVK_OS_PROD_PATH}" + cp -a "${MVK_PROD_PROJ_PATH}/icd/${MVK_ICD_NAME}" "${MVK_OS_PROD_PATH}" +fi +# Remove and replace header include folder +rm -rf "${MVK_PKG_PROD_PATH}/include" +cp -pRL "${MVK_PROD_PROJ_PATH}/include" "${MVK_PKG_PROD_PATH}" diff --git a/Scripts/package_shader_converter.sh b/Scripts/package_shader_converter.sh new file mode 100755 index 00000000..fc001a41 --- /dev/null +++ b/Scripts/package_shader_converter.sh @@ -0,0 +1,73 @@ +#!/bin/bash + +set -e + +# Package folder +export MVK_PROD_BASE_NAME="MoltenVKShaderConverter" +export MVK_WKSPC_PATH="${PROJECT_DIR}" +export MVK_PKG_LOCN="${MVK_WKSPC_PATH}/Package" + +# Remove the base product folder +rm -rf "${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}" + +#----------------------------------- +# MoltenVKGLSLToSPIRVConverter +export MVK_PROD_NAME="MoltenVKGLSLToSPIRVConverter" +export MVK_PKG_CONFIG_PATH="${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}/${MVK_PROD_NAME}" + +# Remove and replace the existing macOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_CONFIG_PATH}/macOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" +fi + +# Remove and replace the existing iOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_CONFIG_PATH}/iOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}-iphoneos" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + rm -rf "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature" + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" +fi + +#----------------------------------- +# MoltenVKSPIRVToMSLConverter +export MVK_PROD_NAME="MoltenVKSPIRVToMSLConverter" +export MVK_PKG_CONFIG_PATH="${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}/${MVK_PROD_NAME}" + +# Remove and replace the existing macOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_CONFIG_PATH}/macOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" +fi + +# Remove and replace the existing iOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_CONFIG_PATH}/iOS" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}-iphoneos" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + rm -rf "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework/_CodeSignature" + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}.framework" "${MVK_OS_PROD_PATH}" +fi + +#----------------------------------- +# MoltenVKShaderConverter Tool +export MVK_PROD_NAME="MoltenVKShaderConverter" +export MVK_PKG_CONFIG_PATH="${MVK_PKG_LOCN}/${CONFIGURATION}/${MVK_PROD_BASE_NAME}" + +# Remove and replace the existing macOS framework folder and copy framework into it +export MVK_OS_PROD_PATH="${MVK_PKG_CONFIG_PATH}/Tools" +export MVK_BUILT_PROD_PATH="${BUILT_PRODUCTS_DIR}" +rm -rf "${MVK_OS_PROD_PATH}" +if [ -e "${MVK_BUILT_PROD_PATH}" ]; then + mkdir -p "${MVK_OS_PROD_PATH}" + cp -a "${MVK_BUILT_PROD_PATH}/${MVK_PROD_NAME}" "${MVK_OS_PROD_PATH}" +fi diff --git a/Scripts/update_latest.sh b/Scripts/update_latest.sh new file mode 100755 index 00000000..8abc75c1 --- /dev/null +++ b/Scripts/update_latest.sh @@ -0,0 +1,14 @@ +#!/bin/bash + +set -e + +# Package folder +export MVK_WKSPC_LOCN="${PROJECT_DIR}" +export MVK_PKG_LOCN="${MVK_WKSPC_LOCN}/Package" + +# Configuration package folder location +export MVK_PKG_CONFIG_LOCN="${CONFIGURATION}" +export MVK_PKG_LATEST_LOCN="Latest" + +# Assign symlink from Latest +ln -sfn "${MVK_PKG_LOCN}/${MVK_PKG_CONFIG_LOCN}" "${MVK_PKG_LOCN}/${MVK_PKG_LATEST_LOCN}" From eb06c9dd5fea6d67481ca8089d02078d7138c013 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 5 Sep 2018 13:51:05 -0500 Subject: [PATCH 15/20] vkCmdFillBuffer: Use the older method to dispatch the compute kernel. We're not even running more than one thread, let alone taking advantage of the new method's automatic threadgroup sizing. --- MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm index b54fc064..68f8759c 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm @@ -1012,7 +1012,7 @@ void MVKCmdFillBuffer::encode(MVKCommandEncoder* cmdEncoder) { [mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdFillBufferMTLComputePipelineState()]; [mtlComputeEnc setBuffer: dstMTLBuff offset: dstMTLBuffOffset atIndex: 0]; [mtlComputeEnc setBytes: &fillInfo length: sizeof(fillInfo) atIndex: 1]; - [mtlComputeEnc dispatchThreads: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)]; + [mtlComputeEnc dispatchThreadgroups: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)]; [mtlComputeEnc popDebugGroup]; } From 38825b960087dd070d24efacde5c07a871279ef4 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Fri, 7 Sep 2018 12:25:40 -0500 Subject: [PATCH 16/20] vkUpdateDescriptorSet: Handle copies of uninitialized descriptors. For a copy, the spec requires only that the source and destination bindings are valid, not that the source has been initialized. If the source binding hasn't been initialized, then we crash attempting to get the Metal resources corresponding to the uninitialized descriptors. So, if the descriptor binding hasn't been initialized, don't try to fetch Metal resources for it. Yes, this causes us to accept writes of NULL descriptors (from templates or otherwise), even though the spec forbids this. I don't know how to solve this without specializing `writeDescriptorSets()` specifically for the `VkCopyDescriptorSet` case (thereby duplicating quite a bit of code). Given the general nature of Vulkan as a framework that does little to no state validation, I wonder if it's even worth it. --- MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm index bd431537..6377ce4e 100644 --- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm +++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm @@ -605,7 +605,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; if (_hasDynamicSamplers) { - _mtlSamplers[dstIdx] = ((MVKSampler*)pImgInfo->sampler)->getMTLSamplerState(); + _mtlSamplers[dstIdx] = pImgInfo->sampler ? ((MVKSampler*)pImgInfo->sampler)->getMTLSamplerState() : nil; } } break; @@ -615,9 +615,9 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, uint32_t dstIdx = dstStartIndex + i; const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; - _mtlTextures[dstIdx] = ((MVKImageView*)pImgInfo->imageView)->getMTLTexture(); + _mtlTextures[dstIdx] = pImgInfo->imageView ? ((MVKImageView*)pImgInfo->imageView)->getMTLTexture() : nil; if (_hasDynamicSamplers) { - _mtlSamplers[dstIdx] = ((MVKSampler*)pImgInfo->sampler)->getMTLSamplerState(); + _mtlSamplers[dstIdx] = pImgInfo->sampler ? ((MVKSampler*)pImgInfo->sampler)->getMTLSamplerState() : nil; } } break; @@ -629,7 +629,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, uint32_t dstIdx = dstStartIndex + i; const auto* pImgInfo = &get(pData, stride, srcStartIndex + i); _imageBindings[dstIdx] = *pImgInfo; - _mtlTextures[dstIdx] = ((MVKImageView*)pImgInfo->imageView)->getMTLTexture(); + _mtlTextures[dstIdx] = pImgInfo->imageView ? ((MVKImageView*)pImgInfo->imageView)->getMTLTexture() : nil; } break; @@ -642,8 +642,8 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, const auto* pBuffInfo = &get(pData, stride, srcStartIndex + i); _bufferBindings[dstIdx] = *pBuffInfo; MVKBuffer* mtlBuff = (MVKBuffer*)pBuffInfo->buffer; - _mtlBuffers[dstIdx] = mtlBuff->getMTLBuffer(); - _mtlBufferOffsets[dstIdx] = mtlBuff->getMTLBufferOffset() + pBuffInfo->offset; + _mtlBuffers[dstIdx] = mtlBuff ? mtlBuff->getMTLBuffer() : nil; + _mtlBufferOffsets[dstIdx] = mtlBuff ? (mtlBuff->getMTLBufferOffset() + pBuffInfo->offset) : 0; } break; @@ -653,7 +653,7 @@ uint32_t MVKDescriptorBinding::writeBindings(uint32_t srcStartIndex, uint32_t dstIdx = dstStartIndex + i; const auto* pBuffView = &get(pData, stride, srcStartIndex + i); _texelBufferBindings[dstIdx] = *pBuffView; - _mtlTextures[dstIdx] = ((MVKBufferView*)*pBuffView)->getMTLTexture(); + _mtlTextures[dstIdx] = *pBuffView ? ((MVKBufferView*)*pBuffView)->getMTLTexture() : nil; } break; default: From e6fe9c093a42c9c334c5edae98281f9a841d7af2 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Fri, 7 Sep 2018 14:19:15 -0500 Subject: [PATCH 17/20] Use the older method for vkCmdCopyBuffers(), too. --- MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm index 68f8759c..d315f367 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm @@ -594,7 +594,7 @@ void MVKCmdCopyBuffer::encode(MVKCommandEncoder* cmdEncoder) { [mtlComputeEnc setBuffer:srcMTLBuff offset: srcMTLBuffOffset atIndex: 0]; [mtlComputeEnc setBuffer:dstMTLBuff offset: dstMTLBuffOffset atIndex: 1]; [mtlComputeEnc setBytes: ©Info length: sizeof(copyInfo) atIndex: 2]; - [mtlComputeEnc dispatchThreads: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)]; + [mtlComputeEnc dispatchThreadgroups: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)]; [mtlComputeEnc popDebugGroup]; } else { id mtlBlitEnc = cmdEncoder->getMTLBlitEncoder(kMVKCommandUseCopyBuffer); From e7cb8b35daddf555e30ac71cc8d4f29ad1999f1b Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Sat, 8 Sep 2018 18:38:32 -0400 Subject: [PATCH 18/20] Update MoltenVK to v1.0.21. Update to latest version of SPRIV-Cross. Update What's New document. --- Docs/Whats_New.md | 30 ++++++++++++++++++--- ExternalRevisions/SPIRV-Cross_repo_revision | 2 +- MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h | 2 +- 3 files changed, 28 insertions(+), 6 deletions(-) diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md index 83969ecd..8dee23b7 100644 --- a/Docs/Whats_New.md +++ b/Docs/Whats_New.md @@ -12,16 +12,38 @@ Copyright (c) 2014-2018 [The Brenwill Workshop Ltd.](http://www.brenwill.com) For best results, use a Markdown reader.* +MoltenVK 1.0.21 +--------------- + +Released 2018/09/08 + +- Add support for extensions: + - VK_KHR_descriptor_update_template +- Create 3D MTLTextureViews for 2D image views of 3D textures. +- Allow building and packaging MoltenVK for of only iOS or only macOS. +- Move packaging scripts out of Xcode projects and into script files. +- vkUpdateDescriptorSet: Handle copies of uninitialized descriptors. +- vkCmdFillBuffer & vkCmdCopyBuffers: Use dispatch call that supports older OS versions. +- Update to latest SPIRV-Cross version: + - MSL: Emit F{Min,Max,Clamp} as fast:: and N{Min,Max,Clamp} as precise + - MSL: Implement multisampled array textures. + - MSL: Emit spvTexelBufferCoord() on ImageWrite to a Buffer. + - MSL: Handle interpolation qualifiers. + - MSL: Account for components when assigning locations to varyings. + - MSL: Do not emit function constants for version < 1.2. + + + MoltenVK 1.0.20 --------------- Released 2018/09/01 - Add support for extensions: - - VK_KHR_maintenance1; - - VK_KHR_shader_draw_parameters; - - VK_KHR_get_physical_device_properties2; - - VK_KHR_push_descriptor; + - VK_KHR_maintenance1 + - VK_KHR_shader_draw_parameters + - VK_KHR_get_physical_device_properties2 + - VK_KHR_push_descriptor - Add ability to track and access supported and enabled extensions. - Update to latest SPIRV-Cross version. diff --git a/ExternalRevisions/SPIRV-Cross_repo_revision b/ExternalRevisions/SPIRV-Cross_repo_revision index 4537bdb3..85beab5d 100644 --- a/ExternalRevisions/SPIRV-Cross_repo_revision +++ b/ExternalRevisions/SPIRV-Cross_repo_revision @@ -1 +1 @@ -6fd66664e8bdadd3f6281aad711f771ef9c24bbe +9ffd4172b46408ab6b03625b0f4f9cfafa5aaa71 diff --git a/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h b/MoltenVK/MoltenVK/API/vk_mvk_moltenvk.h index 22b891e7..a089adf7 100644 --- 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 20 +#define MVK_VERSION_PATCH 21 #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) From 7193ade5ba5186cf37f7b288e027ed6f8cc3d4f5 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Sat, 8 Sep 2018 19:11:27 -0400 Subject: [PATCH 19/20] Update Travis CI to new build process. --- .travis.yml | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/.travis.yml b/.travis.yml index 77d95683..e483e6f2 100644 --- a/.travis.yml +++ b/.travis.yml @@ -1,8 +1,8 @@ language: objective-c # macOS and Xcode Version -# Xcode 9.2 running on macOS 10.12 -osx_image: xcode9.2 +# Xcode 9.4 running on macOS 10.13 +osx_image: xcode9.4 # Build dependencies install: @@ -15,6 +15,6 @@ cache: - External script: - - xcodebuild -scheme "MoltenVK (Debug)" - - xcodebuild -workspace Demos/Demos.xcworkspace -scheme "API-Samples-macOS" + - xcodebuild -scheme "MoltenVK Package (Release)" + - xcodebuild -workspace Demos/Demos.xcworkspace -scheme "Cube-macOS" From 438cc8c382e7a79019014388e66da75b8499e212 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Thu, 6 Sep 2018 20:01:12 -0500 Subject: [PATCH 20/20] Don't pass the offset to the FillBuffer compute kernel. Just add it to the buffer offset when encoding the command. The reason for this is that we were using it to index the buffer--which, according to the C++ spec (on which MSL is based), causes it to be offset by that many 32-bit words instead of bytes. This caused the buffer to be filled incorrectly. While Metal does require the offset to be aligned to the type size (in this case, 4 bytes), Vulkan also requires the offset to `vkCmdFillBuffer()` to be 4-byte aligned, so this shouldn't run into the same problem as `vkCmdCopyBuffer()`. --- MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm | 8 +++----- .../Commands/MVKCommandPipelineStateFactoryShaderSource.h | 3 +-- 2 files changed, 4 insertions(+), 7 deletions(-) diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm index d315f367..0915bb24 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm +++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm @@ -978,7 +978,6 @@ void MVKCmdClearImage::encode(MVKCommandEncoder* cmdEncoder) { // Matches shader struct typedef struct { - uint32_t dstOffset; uint32_t size; uint32_t data; } MVKCmdFillBufferInfo; @@ -999,18 +998,17 @@ void MVKCmdFillBuffer::encode(MVKCommandEncoder* cmdEncoder) { VkDeviceSize byteCnt = (_size == VK_WHOLE_SIZE) ? (_dstBuffer->getByteCount() - (dstMTLBuffOffset + _dstOffset)) : _size; VkDeviceSize wordCnt = byteCnt >> 2; - MVKAssert(mvkFits(_dstOffset) && mvkFits(wordCnt), - "Buffer fill offset and size must each fit into a 32-bit unsigned integer."); + MVKAssert(mvkFits(wordCnt), + "Buffer fill size must fit into a 32-bit unsigned integer."); MVKCmdFillBufferInfo fillInfo; - fillInfo.dstOffset = (uint32_t)_dstOffset; fillInfo.size = (uint32_t)wordCnt; fillInfo.data = _dataValue; id mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseCopyBuffer); [mtlComputeEnc pushDebugGroup: @"vkCmdFillBuffer"]; [mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdFillBufferMTLComputePipelineState()]; - [mtlComputeEnc setBuffer: dstMTLBuff offset: dstMTLBuffOffset atIndex: 0]; + [mtlComputeEnc setBuffer: dstMTLBuff offset: dstMTLBuffOffset+_dstOffset atIndex: 0]; [mtlComputeEnc setBytes: &fillInfo length: sizeof(fillInfo) atIndex: 1]; [mtlComputeEnc dispatchThreadgroups: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)]; [mtlComputeEnc popDebugGroup]; diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h index bda4594b..9da1788d 100644 --- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h +++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h @@ -80,7 +80,6 @@ kernel void cmdCopyBufferBytes(device uint8_t* src [[ buffer(0) ]], }; \n\ \n\ typedef struct { \n\ - uint32_t dstOffset; \n\ uint32_t size; \n\ uint32_t data; \n\ } FillInfo; \n\ @@ -88,7 +87,7 @@ typedef struct { kernel void cmdFillBuffer(device uint32_t* dst [[ buffer(0) ]], \n\ constant FillInfo& info [[ buffer(1) ]]) { \n\ for (uint32_t i = 0; i < info.size; i++) { \n\ - dst[i + info.dstOffset] = info.data; \n\ + dst[i] = info.data; \n\ } \n\ }; \n\ \n\