nihui
/
ncnn
Not watched
1
0
Fork 0
Code
Releases 46 Wiki evaluate Activity
Issues 0 Pull Requests 0 Datasets Model Cloudbrain HPC
You can not select more than 25 topics Topics must start with a chinese character,a letter or number, can include dashes ('-') and can be up to 35 characters long.
Tree: 1469bc8b19
Branches Tags
${ item.name }
Create branch ${ searchTerm }
from '1469bc8b19'
${ noResults }
ncnn/src/gpu.cpp

1696 lines
69 kB
Raw Blame History 

  1. // Tencent is pleased to support the open source community by making ncnn available.

  2. //

  3. // Copyright (C) 2018 THL A29 Limited, a Tencent company. All rights reserved.

  4. //

  5. // Licensed under the BSD 3-Clause License (the "License"); you may not use this file except

  6. // in compliance with the License. You may obtain a copy of the License at

  7. //

  8. // https://opensource.org/licenses/BSD-3-Clause

  9. //

  10. // Unless required by applicable law or agreed to in writing, software distributed

  11. // under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR

  12. // CONDITIONS OF ANY KIND, either express or implied. See the License for the

  13. // specific language governing permissions and limitations under the License.

  14. 

  15. #include "gpu.h"

  16. 

  17. #if NCNN_VULKAN

  18. 

  19. #include <vulkan/vulkan.h>

  20. 

  21. #include <math.h>

  22. #include <stdio.h>

  23. #include <string.h>

  24. 

  25. #include <algorithm>

  26. #include <string>

  27. #include <vector>

  28. 

  29. #include "mat.h"

  30. 

  31. #if __ANDROID__

  32. #define ENABLE_VALIDATION_LAYER 0

  33. #else

  34. #define ENABLE_VALIDATION_LAYER 0

  35. #endif

  36. 

  37. namespace ncnn {

  38. 

  39. // global

  40. static VkInstance g_instance = 0;

  41. static int g_gpu_count = 0;

  42. static int g_default_gpu_index = -1;

  43. 

  44. // NOTE 8 is large enough i think ...

  45. #define NCNN_MAX_GPU_COUNT 8

  46. static GpuInfo g_gpu_infos[NCNN_MAX_GPU_COUNT];

  47. 

  48. // default vulkan device

  49. static Mutex g_default_vkdev_lock;

  50. static VulkanDevice* g_default_vkdev[NCNN_MAX_GPU_COUNT] = {0};

  51. 

  52. // precompiled spirv

  53. struct layer_shader_registry_entry

  54. {

  55.     const uint32_t* spv_data;

  56.     size_t spv_data_size;

  57. };

  58. 

  59. #include "layer_shader_spv_data.h"

  60. 

  61. static const layer_shader_registry_entry layer_shader_registry[] =

  62. {

  63. #include "layer_shader_registry.h"

  64. };

  65. 

  66. static ShaderInfo layer_shader_infos[sizeof(layer_shader_registry) / sizeof(layer_shader_registry_entry)];

  67. 

  68. static const int layer_shader_registry_entry_count = sizeof(layer_shader_registry) / sizeof(layer_shader_registry_entry);

  69. 

  70. int support_VK_KHR_external_memory_capabilities = 0;

  71. int support_VK_KHR_get_physical_device_properties2 = 0;

  72. int support_VK_KHR_get_surface_capabilities2 = 0;

  73. int support_VK_KHR_surface = 0;

  74. int support_VK_EXT_debug_utils = 0;

  75. #if __ANDROID_API__ >= 26

  76. int support_VK_KHR_android_surface = 0;

  77. #endif // __ANDROID_API__ >= 26

  78. 

  79. // VK_KHR_external_memory_capabilities

  80. PFN_vkGetPhysicalDeviceExternalBufferPropertiesKHR vkGetPhysicalDeviceExternalBufferPropertiesKHR = 0;

  81. 

  82. // VK_KHR_get_physical_device_properties2

  83. PFN_vkGetPhysicalDeviceFeatures2KHR vkGetPhysicalDeviceFeatures2KHR = 0;

  84. PFN_vkGetPhysicalDeviceProperties2KHR vkGetPhysicalDeviceProperties2KHR = 0;

  85. PFN_vkGetPhysicalDeviceFormatProperties2KHR vkGetPhysicalDeviceFormatProperties2KHR = 0;

  86. PFN_vkGetPhysicalDeviceImageFormatProperties2KHR vkGetPhysicalDeviceImageFormatProperties2KHR = 0;

  87. PFN_vkGetPhysicalDeviceQueueFamilyProperties2KHR vkGetPhysicalDeviceQueueFamilyProperties2KHR = 0;

  88. PFN_vkGetPhysicalDeviceMemoryProperties2KHR vkGetPhysicalDeviceMemoryProperties2KHR = 0;

  89. PFN_vkGetPhysicalDeviceSparseImageFormatProperties2KHR vkGetPhysicalDeviceSparseImageFormatProperties2KHR = 0;

  90. 

  91. // VK_KHR_get_surface_capabilities2

  92. PFN_vkGetPhysicalDeviceSurfaceCapabilities2KHR vkGetPhysicalDeviceSurfaceCapabilities2KHR = 0;

  93. PFN_vkGetPhysicalDeviceSurfaceFormats2KHR vkGetPhysicalDeviceSurfaceFormats2KHR = 0;

  94. 

  95. // VK_KHR_surface

  96. PFN_vkDestroySurfaceKHR vkDestroySurfaceKHR = 0;

  97. PFN_vkGetPhysicalDeviceSurfaceSupportKHR vkGetPhysicalDeviceSurfaceSupportKHR = 0;

  98. PFN_vkGetPhysicalDeviceSurfaceCapabilitiesKHR vkGetPhysicalDeviceSurfaceCapabilitiesKHR = 0;

  99. PFN_vkGetPhysicalDeviceSurfaceFormatsKHR vkGetPhysicalDeviceSurfaceFormatsKHR = 0;

  100. PFN_vkGetPhysicalDeviceSurfacePresentModesKHR vkGetPhysicalDeviceSurfacePresentModesKHR = 0;

  101. 

  102. #if __ANDROID_API__ >= 26

  103. // VK_KHR_android_surface

  104. PFN_vkCreateAndroidSurfaceKHR vkCreateAndroidSurfaceKHR = 0;

  105. #endif // __ANDROID_API__ >= 26

  106. 

  107. // compile with old vulkan sdk

  108. #if VK_HEADER_VERSION < 80

  109. #define VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_8BIT_STORAGE_FEATURES_KHR (VkStructureType)1000177000

  110. typedef struct VkPhysicalDevice8BitStorageFeaturesKHR {

  111.     VkStructureType    sType;

  112.     void*              pNext;

  113.     VkBool32           storageBuffer8BitAccess;

  114.     VkBool32           uniformAndStorageBuffer8BitAccess;

  115.     VkBool32           storagePushConstant8;

  116. } VkPhysicalDevice8BitStorageFeaturesKHR;

  117. #endif // VK_HEADER_VERSION < 80

  118. #if VK_HEADER_VERSION < 95

  119. #define VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR (VkStructureType)1000082000

  120. typedef struct VkPhysicalDeviceFloat16Int8FeaturesKHR {

  121.     VkStructureType    sType;

  122.     void*              pNext;

  123.     VkBool32           shaderFloat16;

  124.     VkBool32           shaderInt8;

  125. } VkPhysicalDeviceFloat16Int8FeaturesKHR;

  126. #endif // VK_HEADER_VERSION < 95

  127. 

  128. static int init_instance_extension()

  129. {

  130.     if (support_VK_KHR_external_memory_capabilities)

  131.     {

  132.         vkGetPhysicalDeviceExternalBufferPropertiesKHR = (PFN_vkGetPhysicalDeviceExternalBufferPropertiesKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceExternalBufferPropertiesKHR");

  133.     }

  134. 

  135.     if (support_VK_KHR_get_physical_device_properties2)

  136.     {

  137.         vkGetPhysicalDeviceFeatures2KHR = (PFN_vkGetPhysicalDeviceFeatures2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceFeatures2KHR");

  138.         vkGetPhysicalDeviceProperties2KHR = (PFN_vkGetPhysicalDeviceProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceProperties2KHR");

  139.         vkGetPhysicalDeviceFormatProperties2KHR = (PFN_vkGetPhysicalDeviceFormatProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceFormatProperties2KHR");

  140.         vkGetPhysicalDeviceImageFormatProperties2KHR = (PFN_vkGetPhysicalDeviceImageFormatProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceImageFormatProperties2KHR");

  141.         vkGetPhysicalDeviceQueueFamilyProperties2KHR = (PFN_vkGetPhysicalDeviceQueueFamilyProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceQueueFamilyProperties2KHR");

  142.         vkGetPhysicalDeviceMemoryProperties2KHR = (PFN_vkGetPhysicalDeviceMemoryProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceMemoryProperties2KHR");

  143.         vkGetPhysicalDeviceSparseImageFormatProperties2KHR = (PFN_vkGetPhysicalDeviceSparseImageFormatProperties2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSparseImageFormatProperties2KHR");

  144.     }

  145. 

  146.     if (support_VK_KHR_get_surface_capabilities2)

  147.     {

  148.         vkGetPhysicalDeviceSurfaceCapabilities2KHR = (PFN_vkGetPhysicalDeviceSurfaceCapabilities2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfaceCapabilities2KHR");;

  149.         vkGetPhysicalDeviceSurfaceFormats2KHR = (PFN_vkGetPhysicalDeviceSurfaceFormats2KHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfaceFormats2KHR");;

  150.     }

  151. 

  152.     if (support_VK_KHR_surface)

  153.     {

  154.         vkDestroySurfaceKHR = (PFN_vkDestroySurfaceKHR)vkGetInstanceProcAddr(g_instance, "vkDestroySurfaceKHR");;

  155.         vkGetPhysicalDeviceSurfaceSupportKHR = (PFN_vkGetPhysicalDeviceSurfaceSupportKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfaceSupportKHR");;

  156.         vkGetPhysicalDeviceSurfaceCapabilitiesKHR = (PFN_vkGetPhysicalDeviceSurfaceCapabilitiesKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfaceCapabilitiesKHR");;

  157.         vkGetPhysicalDeviceSurfaceFormatsKHR = (PFN_vkGetPhysicalDeviceSurfaceFormatsKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfaceFormatsKHR");;

  158.         vkGetPhysicalDeviceSurfacePresentModesKHR = (PFN_vkGetPhysicalDeviceSurfacePresentModesKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceSurfacePresentModesKHR");;

  159.     }

  160. 

  161. #if __ANDROID_API__ >= 26

  162.     if (support_VK_KHR_android_surface)

  163.     {

  164.         vkCreateAndroidSurfaceKHR = (PFN_vkCreateAndroidSurfaceKHR)vkGetInstanceProcAddr(g_instance, "vkCreateAndroidSurfaceKHR");

  165.     }

  166. #endif // __ANDROID_API__ >= 26

  167. 

  168.     return 0;

  169. }

  170. 

  171. #if ENABLE_VALIDATION_LAYER

  172. static VkDebugUtilsMessengerEXT callback;

  173. 

  174. static VKAPI_ATTR VkBool32 VKAPI_CALL debugCallback(

  175.     VkDebugUtilsMessageSeverityFlagBitsEXT /*messageSeverity*/,

  176.     VkDebugUtilsMessageTypeFlagsEXT /*messageType*/,

  177.     const VkDebugUtilsMessengerCallbackDataEXT* pCallbackData,

  178.     void* /*pUserData*/)

  179. {

  180.     fprintf(stderr, "validation layer: %s\n", pCallbackData->pMessage);

  181. 

  182.     return VK_FALSE;

  183. }

  184. 

  185. VkResult CreateDebugUtilsMessengerEXT(VkInstance instance, const VkDebugUtilsMessengerCreateInfoEXT* pCreateInfo, const VkAllocationCallbacks* pAllocator, VkDebugUtilsMessengerEXT* pCallback)

  186. {

  187.     PFN_vkCreateDebugUtilsMessengerEXT func = (PFN_vkCreateDebugUtilsMessengerEXT)vkGetInstanceProcAddr(instance, "vkCreateDebugUtilsMessengerEXT");

  188.     if (func)

  189.         return func(instance, pCreateInfo, pAllocator, pCallback);

  190. 

  191.     return VK_ERROR_EXTENSION_NOT_PRESENT;

  192. }

  193. 

  194. void DestroyDebugUtilsMessengerEXT(VkInstance instance, VkDebugUtilsMessengerEXT callback, const VkAllocationCallbacks* pAllocator)

  195. {

  196.     PFN_vkDestroyDebugUtilsMessengerEXT func = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr(instance, "vkDestroyDebugUtilsMessengerEXT");

  197.     if (func)

  198.         func(instance, callback, pAllocator);

  199. }

  200. #endif // ENABLE_VALIDATION_LAYER

  201. 

  202. static uint32_t find_device_compute_queue(const std::vector<VkQueueFamilyProperties>& queueFamilyProperties)

  203. {

  204.     // first try, compute only queue

  205.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  206.     {

  207.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  208. 

  209.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  210.             && !(queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT))

  211.         {

  212.             return i;

  213.         }

  214.     }

  215. 

  216.     // second try, any queue with compute and graphics

  217.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  218.     {

  219.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  220. 

  221.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  222.             && (queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT))

  223.         {

  224.             return i;

  225.         }

  226.     }

  227. 

  228.     // third try, any queue with compute

  229.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  230.     {

  231.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  232. 

  233.         if (queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  234.         {

  235.             return i;

  236.         }

  237.     }

  238. 

  239. //     fprintf(stderr, "no compute queue\n");

  240.     return -1;

  241. }

  242. 

  243. static uint32_t find_device_graphics_queue(const std::vector<VkQueueFamilyProperties>& queueFamilyProperties)

  244. {

  245.     // first try, graphics only queue

  246.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  247.     {

  248.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  249. 

  250.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

  251.             && !(queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT))

  252.         {

  253.             return i;

  254.         }

  255.     }

  256. 

  257.     // second try, any queue with graphics and compute

  258.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  259.     {

  260.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  261. 

  262.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

  263.             && (queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT))

  264.         {

  265.             return i;

  266.         }

  267.     }

  268. 

  269.     // third try, any queue with graphics

  270.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  271.     {

  272.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  273. 

  274.         if (queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

  275.         {

  276.             return i;

  277.         }

  278.     }

  279. 

  280. //     fprintf(stderr, "no graphics queue\n");

  281.     return -1;

  282. }

  283. 

  284. static uint32_t find_device_transfer_queue(const std::vector<VkQueueFamilyProperties>& queueFamilyProperties)

  285. {

  286.     // first try, transfer only queue

  287.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  288.     {

  289.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  290. 

  291.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_TRANSFER_BIT)

  292.             && !(queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  293.             && !(queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT))

  294.         {

  295.             return i;

  296.         }

  297.     }

  298. 

  299.     // second try, any queue with transfer

  300.     for (uint32_t i=0; i<queueFamilyProperties.size(); i++)

  301.     {

  302.         const VkQueueFamilyProperties& queueFamilyProperty = queueFamilyProperties[i];

  303. 

  304.         if (queueFamilyProperty.queueFlags & VK_QUEUE_TRANSFER_BIT)

  305.         {

  306.             return i;

  307.         }

  308.     }

  309. 

  310.     // third try, use compute queue

  311.     uint32_t compute_queue_index = find_device_compute_queue(queueFamilyProperties);

  312.     if (compute_queue_index != (uint32_t)-1)

  313.     {

  314.         return compute_queue_index;

  315.     }

  316. 

  317.     // fourth try, use graphics queue

  318.     uint32_t graphics_queue_index = find_device_graphics_queue(queueFamilyProperties);

  319.     if (graphics_queue_index != (uint32_t)-1)

  320.     {

  321.         return graphics_queue_index;

  322.     }

  323. 

  324. //     fprintf(stderr, "no transfer queue\n");

  325.     return -1;

  326. }

  327. 

  328. static int find_default_vulkan_device_index()

  329. {

  330.     // first try, discrete gpu

  331.     for (int i=0; i<g_gpu_count; i++)

  332.     {

  333.         if (g_gpu_infos[i].type == 0)

  334.             return i;

  335.     }

  336. 

  337.     // second try, integrated gpu

  338.     for (int i=0; i<g_gpu_count; i++)

  339.     {

  340.         if (g_gpu_infos[i].type == 1)

  341.             return i;

  342.     }

  343. 

  344.     // third try, any probed device

  345.     if (g_gpu_count > 0)

  346.         return 0;

  347. 

  348.     fprintf(stderr, "no vulkan device\n");

  349.     return -1;

  350. }

  351. 

  352. int create_gpu_instance()

  353. {

  354.     VkResult ret;

  355. 

  356.     std::vector<const char*> enabledLayers;

  357. 

  358. #if ENABLE_VALIDATION_LAYER

  359.     uint32_t instanceLayerPropertyCount;

  360.     ret = vkEnumerateInstanceLayerProperties(&instanceLayerPropertyCount, NULL);

  361.     if (ret != VK_SUCCESS)

  362.     {

  363.         fprintf(stderr, "vkEnumerateInstanceLayerProperties failed %d\n", ret);

  364.         return -1;

  365.     }

  366. 

  367.     std::vector<VkLayerProperties> instanceLayerProperties(instanceLayerPropertyCount);

  368.     ret = vkEnumerateInstanceLayerProperties(&instanceLayerPropertyCount, instanceLayerProperties.data());

  369.     if (ret != VK_SUCCESS)

  370.     {

  371.         fprintf(stderr, "vkEnumerateInstanceLayerProperties failed %d\n", ret);

  372.         return -1;

  373.     }

  374. 

  375.     for (uint32_t i=0; i<instanceLayerPropertyCount; i++)

  376.     {

  377.         const VkLayerProperties& lp = instanceLayerProperties[i];

  378. //         fprintf(stderr, "instance layer %s = %u\n", lp.layerName, lp.implementationVersion);

  379. 

  380.         if (strcmp(lp.layerName, "VK_LAYER_LUNARG_standard_validation") == 0)

  381.         {

  382.             enabledLayers.push_back("VK_LAYER_LUNARG_standard_validation");

  383.         }

  384.         if (strcmp(lp.layerName, "VK_LAYER_LUNARG_parameter_validation") == 0)

  385.         {

  386.             enabledLayers.push_back("VK_LAYER_LUNARG_parameter_validation");

  387.         }

  388.     }

  389. #endif // ENABLE_VALIDATION_LAYER

  390. 

  391.     std::vector<const char*> enabledExtensions;

  392. 

  393.     uint32_t instanceExtensionPropertyCount;

  394.     ret = vkEnumerateInstanceExtensionProperties(NULL, &instanceExtensionPropertyCount, NULL);

  395.     if (ret != VK_SUCCESS)

  396.     {

  397.         fprintf(stderr, "vkEnumerateInstanceExtensionProperties failed %d\n", ret);

  398.         return -1;

  399.     }

  400. 

  401.     std::vector<VkExtensionProperties> instanceExtensionProperties(instanceExtensionPropertyCount);

  402.     ret = vkEnumerateInstanceExtensionProperties(NULL, &instanceExtensionPropertyCount, instanceExtensionProperties.data());

  403.     if (ret != VK_SUCCESS)

  404.     {

  405.         fprintf(stderr, "vkEnumerateInstanceExtensionProperties failed %d\n", ret);

  406.         return -1;

  407.     }

  408. 

  409.     support_VK_KHR_get_physical_device_properties2 = 0;

  410.     support_VK_KHR_get_surface_capabilities2 = 0;

  411.     support_VK_KHR_surface = 0;

  412.     support_VK_EXT_debug_utils = 0;

  413. #if __ANDROID_API__ >= 26

  414.     support_VK_KHR_android_surface = 0;

  415. #endif // __ANDROID_API__ >= 26

  416.     for (uint32_t j=0; j<instanceExtensionPropertyCount; j++)

  417.     {

  418.         const VkExtensionProperties& exp = instanceExtensionProperties[j];

  419. //         fprintf(stderr, "instance extension %s = %u\n", exp.extensionName, exp.specVersion);

  420. 

  421.         if (strcmp(exp.extensionName, "VK_KHR_external_memory_capabilities") == 0)

  422.             support_VK_KHR_external_memory_capabilities = exp.specVersion;

  423.         else if (strcmp(exp.extensionName, "VK_KHR_get_physical_device_properties2") == 0)

  424.             support_VK_KHR_get_physical_device_properties2 = exp.specVersion;

  425.         else if (strcmp(exp.extensionName, "VK_KHR_get_surface_capabilities2") == 0)

  426.             support_VK_KHR_get_surface_capabilities2 = exp.specVersion;

  427.         else if (strcmp(exp.extensionName, "VK_KHR_surface") == 0)

  428.             support_VK_KHR_surface = exp.specVersion;

  429.         else if (strcmp(exp.extensionName, "VK_EXT_debug_utils") == 0)

  430.             support_VK_EXT_debug_utils = exp.specVersion;

  431. #if __ANDROID_API__ >= 26

  432.         else if (strcmp(exp.extensionName, "VK_KHR_android_surface") == 0)

  433.             support_VK_KHR_android_surface = exp.specVersion;

  434. #endif // __ANDROID_API__ >= 26

  435.     }

  436. 

  437.     if (support_VK_KHR_external_memory_capabilities)

  438.         enabledExtensions.push_back("VK_KHR_external_memory_capabilities");

  439.     if (support_VK_KHR_get_physical_device_properties2)

  440.         enabledExtensions.push_back("VK_KHR_get_physical_device_properties2");

  441.     if (support_VK_KHR_get_surface_capabilities2)

  442.         enabledExtensions.push_back("VK_KHR_get_surface_capabilities2");

  443.     if (support_VK_KHR_surface)

  444.         enabledExtensions.push_back("VK_KHR_surface");

  445. #if ENABLE_VALIDATION_LAYER

  446.     if (support_VK_EXT_debug_utils)

  447.         enabledExtensions.push_back("VK_EXT_debug_utils");

  448. #endif // ENABLE_VALIDATION_LAYER

  449. #if __ANDROID_API__ >= 26

  450.     if (support_VK_KHR_android_surface)

  451.         enabledExtensions.push_back("VK_KHR_android_surface");

  452. #endif // __ANDROID_API__ >= 26

  453. 

  454.     VkApplicationInfo applicationInfo;

  455.     applicationInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;

  456.     applicationInfo.pNext = 0;

  457.     applicationInfo.pApplicationName = "ncnn";

  458.     applicationInfo.applicationVersion = 0;

  459.     applicationInfo.pEngineName = "ncnn";

  460.     applicationInfo.engineVersion = 20200222;

  461.     applicationInfo.apiVersion = VK_MAKE_VERSION(1, 0, 0);

  462. 

  463.     VkInstanceCreateInfo instanceCreateInfo;

  464.     instanceCreateInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;

  465.     instanceCreateInfo.pNext = 0;

  466.     instanceCreateInfo.flags = 0;

  467.     instanceCreateInfo.pApplicationInfo = &applicationInfo;

  468.     instanceCreateInfo.enabledLayerCount = enabledLayers.size();

  469.     instanceCreateInfo.ppEnabledLayerNames = enabledLayers.data();

  470.     instanceCreateInfo.enabledExtensionCount = enabledExtensions.size();

  471.     instanceCreateInfo.ppEnabledExtensionNames = enabledExtensions.data();

  472. 

  473.     ret = vkCreateInstance(&instanceCreateInfo, 0, &g_instance);

  474.     if (ret != VK_SUCCESS)

  475.     {

  476.         fprintf(stderr, "vkCreateInstance failed %d\n", ret);

  477.         return -1;

  478.     }

  479. 

  480. #if ENABLE_VALIDATION_LAYER

  481.     if (support_VK_EXT_debug_utils)

  482.     {

  483.         VkDebugUtilsMessengerCreateInfoEXT createInfo = {};

  484.         createInfo.sType = VK_STRUCTURE_TYPE_DEBUG_UTILS_MESSENGER_CREATE_INFO_EXT;

  485.         createInfo.messageSeverity = VK_DEBUG_UTILS_MESSAGE_SEVERITY_VERBOSE_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_SEVERITY_WARNING_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_SEVERITY_ERROR_BIT_EXT;

  486.         createInfo.messageType = VK_DEBUG_UTILS_MESSAGE_TYPE_GENERAL_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_TYPE_VALIDATION_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_TYPE_PERFORMANCE_BIT_EXT;

  487.         createInfo.pfnUserCallback = debugCallback;

  488.         createInfo.pUserData = 0;

  489.         ret = CreateDebugUtilsMessengerEXT(g_instance, &createInfo, NULL, &callback);

  490.         if (ret != VK_SUCCESS)

  491.         {

  492.             fprintf(stderr, "CreateDebugUtilsMessengerEXT failed %d\n", ret);

  493.             return -1;

  494.         }

  495.     }

  496. #endif // ENABLE_VALIDATION_LAYER

  497. 

  498.     init_instance_extension();

  499. 

  500.     uint32_t physicalDeviceCount = 0;

  501.     ret = vkEnumeratePhysicalDevices(g_instance, &physicalDeviceCount, 0);

  502.     if (ret != VK_SUCCESS)

  503.     {

  504.         fprintf(stderr, "vkEnumeratePhysicalDevices failed %d\n", ret);

  505.         return -1;

  506.     }

  507. 

  508.     if (physicalDeviceCount > NCNN_MAX_GPU_COUNT)

  509.         physicalDeviceCount = NCNN_MAX_GPU_COUNT;

  510. 

  511.     std::vector<VkPhysicalDevice> physicalDevices(physicalDeviceCount);

  512. 

  513.     ret = vkEnumeratePhysicalDevices(g_instance, &physicalDeviceCount, physicalDevices.data());

  514.     if (ret != VK_SUCCESS)

  515.     {

  516.         fprintf(stderr, "vkEnumeratePhysicalDevices failed %d\n", ret);

  517.         return -1;

  518.     }

  519. 

  520.     // find proper device and queue

  521.     int gpu_info_index = 0;

  522.     for (uint32_t i=0; i<physicalDeviceCount; i++)

  523.     {

  524.         const VkPhysicalDevice& physicalDevice = physicalDevices[i];

  525.         GpuInfo& gpu_info = g_gpu_infos[gpu_info_index];

  526. 

  527.         // device type

  528.         VkPhysicalDeviceProperties physicalDeviceProperties;

  529.         vkGetPhysicalDeviceProperties(physicalDevice, &physicalDeviceProperties);

  530. 

  531. //         fprintf(stderr, "[%u] apiVersion = %u.%u.%u\n", i, VK_VERSION_MAJOR(physicalDeviceProperties.apiVersion),

  532. //             VK_VERSION_MINOR(physicalDeviceProperties.apiVersion), VK_VERSION_PATCH(physicalDeviceProperties.apiVersion));

  533. //         fprintf(stderr, "[%u] driverVersion = %u.%u.%u\n", i, VK_VERSION_MAJOR(physicalDeviceProperties.driverVersion),

  534. //             VK_VERSION_MINOR(physicalDeviceProperties.driverVersion), VK_VERSION_PATCH(physicalDeviceProperties.driverVersion));

  535. //         fprintf(stderr, "[%u] vendorID = %x\n", i, physicalDeviceProperties.vendorID);

  536. //         fprintf(stderr, "[%u] deviceID = %x\n", i, physicalDeviceProperties.deviceID);

  537. //         fprintf(stderr, "[%u] deviceType = %x\n", i, physicalDeviceProperties.deviceType);

  538. //         fprintf(stderr, "[%u] deviceName = %s\n", i, physicalDeviceProperties.deviceName);

  539. //         fprintf(stderr, "[%u] pipelineCacheUUID = %u\n", i, physicalDeviceProperties.pipelineCacheUUID);

  540. 

  541.         gpu_info.bug_local_size_spec_const = false;

  542.         gpu_info.bug_implicit_fp16_arithmetic = false;

  543. 

  544.         if (physicalDeviceProperties.vendorID == 0x13b5 && physicalDeviceProperties.apiVersion < VK_MAKE_VERSION(1, 0, 66))

  545.         {

  546.             // arm mali with old buggy driver

  547.             gpu_info.bug_local_size_spec_const = true;

  548.         }

  549. 

  550.         if (physicalDeviceProperties.vendorID == 0x5143 && physicalDeviceProperties.apiVersion < VK_MAKE_VERSION(1, 0, 49))

  551.         {

  552.             // qcom adreno with old buggy driver

  553.             gpu_info.bug_local_size_spec_const = true;

  554.         }

  555. 

  556.         if (physicalDeviceProperties.vendorID == 0x13b5 && (physicalDeviceProperties.deviceID == 0x7500001 || physicalDeviceProperties.deviceID == 0x8602000))

  557.         {

  558.             // TODO enable devices other than rk3288/rk3399

  559.             // arm mali driver accept spirv with fp16 arithmetic

  560.             gpu_info.bug_implicit_fp16_arithmetic = true;

  561.         }

  562. 

  563.         if (physicalDeviceProperties.vendorID == 0x5143 && (physicalDeviceProperties.deviceID == 0x6030001 || physicalDeviceProperties.deviceID == 0x6040001))

  564.         {

  565.             // TODO enable devices other than qcom855/qcom855plus

  566.             // qcom adreno driver accept spirv with fp16 arithmetic

  567.             gpu_info.bug_implicit_fp16_arithmetic = true;

  568.         }

  569. 

  570.         gpu_info.physical_device = physicalDevice;

  571. 

  572.         // info

  573.         gpu_info.api_version = physicalDeviceProperties.apiVersion;

  574.         gpu_info.driver_version = physicalDeviceProperties.driverVersion;

  575.         gpu_info.vendor_id = physicalDeviceProperties.vendorID;

  576.         gpu_info.device_id = physicalDeviceProperties.deviceID;

  577.         memcpy(gpu_info.pipeline_cache_uuid, physicalDeviceProperties.pipelineCacheUUID, VK_UUID_SIZE);

  578. 

  579.         if (physicalDeviceProperties.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU)

  580.             gpu_info.type = 0;

  581.         else if (physicalDeviceProperties.deviceType == VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU)

  582.             gpu_info.type = 1;

  583.         else if (physicalDeviceProperties.deviceType == VK_PHYSICAL_DEVICE_TYPE_VIRTUAL_GPU)

  584.             gpu_info.type = 2;

  585.         else if (physicalDeviceProperties.deviceType == VK_PHYSICAL_DEVICE_TYPE_CPU)

  586.             gpu_info.type = 3;

  587.         else

  588.             gpu_info.type = -1;

  589. 

  590.         // device capability

  591.         gpu_info.max_shared_memory_size = physicalDeviceProperties.limits.maxComputeSharedMemorySize;

  592. 

  593.         gpu_info.max_workgroup_count[0] = physicalDeviceProperties.limits.maxComputeWorkGroupCount[0];

  594.         gpu_info.max_workgroup_count[1] = physicalDeviceProperties.limits.maxComputeWorkGroupCount[1];

  595.         gpu_info.max_workgroup_count[2] = physicalDeviceProperties.limits.maxComputeWorkGroupCount[2];

  596. 

  597.         gpu_info.max_workgroup_invocations = physicalDeviceProperties.limits.maxComputeWorkGroupInvocations;

  598. 

  599.         gpu_info.max_workgroup_size[0] = physicalDeviceProperties.limits.maxComputeWorkGroupSize[0];

  600.         gpu_info.max_workgroup_size[1] = physicalDeviceProperties.limits.maxComputeWorkGroupSize[1];

  601.         gpu_info.max_workgroup_size[2] = physicalDeviceProperties.limits.maxComputeWorkGroupSize[2];

  602. 

  603.         gpu_info.memory_map_alignment = physicalDeviceProperties.limits.minMemoryMapAlignment;

  604.         gpu_info.buffer_offset_alignment = physicalDeviceProperties.limits.minStorageBufferOffsetAlignment;

  605.         gpu_info.non_coherent_atom_size = physicalDeviceProperties.limits.nonCoherentAtomSize;

  606. 

  607.         gpu_info.timestamp_period = physicalDeviceProperties.limits.timestampPeriod;

  608. 

  609. //         fprintf(stderr, "[%u] max_shared_memory_size = %u\n", i, gpu_info.max_shared_memory_size);

  610. //         fprintf(stderr, "[%u] max_workgroup_count = %u %u %u\n", i, gpu_info.max_workgroup_count[0], gpu_info.max_workgroup_count[1], gpu_info.max_workgroup_count[2]);

  611. //         fprintf(stderr, "[%u] max_workgroup_invocations = %u\n", i, gpu_info.max_workgroup_invocations);

  612. //         fprintf(stderr, "[%u] max_workgroup_size = %u %u %u\n", i, gpu_info.max_workgroup_size[0], gpu_info.max_workgroup_size[1], gpu_info.max_workgroup_size[2]);

  613. //         fprintf(stderr, "[%u] memory_map_alignment = %lu\n", i, gpu_info.memory_map_alignment);

  614. //         fprintf(stderr, "[%u] buffer_offset_alignment = %lu\n", i, gpu_info.buffer_offset_alignment);

  615. 

  616.         // find compute queue

  617.         uint32_t queueFamilyPropertiesCount;

  618.         vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &queueFamilyPropertiesCount, 0);

  619. 

  620.         std::vector<VkQueueFamilyProperties> queueFamilyProperties(queueFamilyPropertiesCount);

  621.         vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &queueFamilyPropertiesCount, queueFamilyProperties.data());

  622. 

  623.         gpu_info.compute_queue_family_index = find_device_compute_queue(queueFamilyProperties);

  624.         gpu_info.graphics_queue_family_index = find_device_graphics_queue(queueFamilyProperties);

  625.         gpu_info.transfer_queue_family_index = find_device_transfer_queue(queueFamilyProperties);

  626. 

  627.         gpu_info.compute_queue_count = queueFamilyProperties[gpu_info.compute_queue_family_index].queueCount;

  628.         gpu_info.graphics_queue_count = queueFamilyProperties[gpu_info.graphics_queue_family_index].queueCount;

  629.         gpu_info.transfer_queue_count = queueFamilyProperties[gpu_info.transfer_queue_family_index].queueCount;

  630. 

  631.         // cache memory properties

  632.         vkGetPhysicalDeviceMemoryProperties(physicalDevice, &gpu_info.physicalDeviceMemoryProperties);

  633. 

  634.         // get device extension

  635.         uint32_t deviceExtensionPropertyCount = 0;

  636.         ret = vkEnumerateDeviceExtensionProperties(physicalDevice, NULL, &deviceExtensionPropertyCount, NULL);

  637.         if (ret != VK_SUCCESS)

  638.         {

  639.             fprintf(stderr, "vkEnumerateDeviceExtensionProperties failed %d\n", ret);

  640.             return -1;

  641.         }

  642. 

  643.         std::vector<VkExtensionProperties> deviceExtensionProperties(deviceExtensionPropertyCount);

  644.         ret = vkEnumerateDeviceExtensionProperties(physicalDevice, NULL, &deviceExtensionPropertyCount, deviceExtensionProperties.data());

  645.         if (ret != VK_SUCCESS)

  646.         {

  647.             fprintf(stderr, "vkEnumerateDeviceExtensionProperties failed %d\n", ret);

  648.             return -1;

  649.         }

  650. 

  651.         // extension capability

  652.         gpu_info.support_VK_KHR_8bit_storage = 0;

  653.         gpu_info.support_VK_KHR_16bit_storage = 0;

  654.         gpu_info.support_VK_KHR_bind_memory2 = 0;

  655.         gpu_info.support_VK_KHR_dedicated_allocation = 0;

  656.         gpu_info.support_VK_KHR_descriptor_update_template = 0;

  657.         gpu_info.support_VK_KHR_external_memory = 0;

  658.         gpu_info.support_VK_KHR_get_memory_requirements2 = 0;

  659.         gpu_info.support_VK_KHR_maintenance1 = 0;

  660.         gpu_info.support_VK_KHR_push_descriptor = 0;

  661.         gpu_info.support_VK_KHR_sampler_ycbcr_conversion = 0;

  662.         gpu_info.support_VK_KHR_shader_float16_int8 = 0;

  663.         gpu_info.support_VK_KHR_shader_float_controls = 0;

  664.         gpu_info.support_VK_KHR_storage_buffer_storage_class = 0;

  665.         gpu_info.support_VK_KHR_swapchain = 0;

  666.         gpu_info.support_VK_EXT_queue_family_foreign = 0;

  667. #if __ANDROID_API__ >= 26

  668.         gpu_info.support_VK_ANDROID_external_memory_android_hardware_buffer = 0;

  669. #endif // __ANDROID_API__ >= 26

  670.         for (uint32_t j=0; j<deviceExtensionPropertyCount; j++)

  671.         {

  672.             const VkExtensionProperties& exp = deviceExtensionProperties[j];

  673. //             fprintf(stderr, "device extension %s = %u\n", exp.extensionName, exp.specVersion);

  674. 

  675.             if (strcmp(exp.extensionName, "VK_KHR_8bit_storage") == 0)

  676.                 gpu_info.support_VK_KHR_8bit_storage = exp.specVersion;

  677.             else if (strcmp(exp.extensionName, "VK_KHR_16bit_storage") == 0)

  678.                 gpu_info.support_VK_KHR_16bit_storage = exp.specVersion;

  679.             else if (strcmp(exp.extensionName, "VK_KHR_bind_memory2") == 0)

  680.                 gpu_info.support_VK_KHR_bind_memory2 = exp.specVersion;

  681.             else if (strcmp(exp.extensionName, "VK_KHR_dedicated_allocation") == 0)

  682.                 gpu_info.support_VK_KHR_dedicated_allocation = exp.specVersion;

  683.             else if (strcmp(exp.extensionName, "VK_KHR_descriptor_update_template") == 0)

  684.                 gpu_info.support_VK_KHR_descriptor_update_template = exp.specVersion;

  685.             else if (strcmp(exp.extensionName, "VK_KHR_external_memory") == 0)

  686.                 gpu_info.support_VK_KHR_external_memory = exp.specVersion;

  687.             else if (strcmp(exp.extensionName, "VK_KHR_get_memory_requirements2") == 0)

  688.                 gpu_info.support_VK_KHR_get_memory_requirements2 = exp.specVersion;

  689.             else if (strcmp(exp.extensionName, "VK_KHR_maintenance1") == 0)

  690.                 gpu_info.support_VK_KHR_maintenance1 = exp.specVersion;

  691.             else if (strcmp(exp.extensionName, "VK_KHR_push_descriptor") == 0)

  692.                 gpu_info.support_VK_KHR_push_descriptor = exp.specVersion;

  693.             else if (strcmp(exp.extensionName, "VK_KHR_sampler_ycbcr_conversion") == 0)

  694.                 gpu_info.support_VK_KHR_sampler_ycbcr_conversion = exp.specVersion;

  695.             else if (strcmp(exp.extensionName, "VK_KHR_shader_float16_int8") == 0)

  696.                 gpu_info.support_VK_KHR_shader_float16_int8 = exp.specVersion;

  697.             else if (strcmp(exp.extensionName, "VK_KHR_shader_float_controls") == 0)

  698.                 gpu_info.support_VK_KHR_shader_float_controls = exp.specVersion;

  699.             else if (strcmp(exp.extensionName, "VK_KHR_storage_buffer_storage_class") == 0)

  700.                 gpu_info.support_VK_KHR_storage_buffer_storage_class = exp.specVersion;

  701.             else if (strcmp(exp.extensionName, "VK_KHR_swapchain") == 0)

  702.                 gpu_info.support_VK_KHR_swapchain = exp.specVersion;

  703.             else if (strcmp(exp.extensionName, "VK_EXT_queue_family_foreign") == 0)

  704.                 gpu_info.support_VK_EXT_queue_family_foreign = exp.specVersion;

  705. #if __ANDROID_API__ >= 26

  706.             else if (strcmp(exp.extensionName, "VK_ANDROID_external_memory_android_hardware_buffer") == 0)

  707.                 gpu_info.support_VK_ANDROID_external_memory_android_hardware_buffer = exp.specVersion;

  708. #endif // __ANDROID_API__ >= 26

  709.         }

  710. 

  711.         // check features

  712.         gpu_info.support_fp16_packed = true;

  713.         gpu_info.support_fp16_storage = false;

  714.         gpu_info.support_fp16_arithmetic = false;

  715.         gpu_info.support_int8_storage = false;

  716.         gpu_info.support_int8_arithmetic = false;

  717.         gpu_info.support_ycbcr_conversion = false;

  718.         if (support_VK_KHR_get_physical_device_properties2)

  719.         {

  720.             void* queryExtensionFeatures = 0;

  721. 

  722.             // query int8 storage

  723.             VkPhysicalDevice8BitStorageFeaturesKHR query8BitStorageFeatures;

  724.             query8BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_8BIT_STORAGE_FEATURES_KHR;

  725.             query8BitStorageFeatures.pNext = 0;

  726.             if (gpu_info.support_VK_KHR_8bit_storage)

  727.             {

  728.                 query8BitStorageFeatures.pNext = queryExtensionFeatures;

  729.                 queryExtensionFeatures = &query8BitStorageFeatures;

  730.             }

  731. 

  732.             // query fp16/int16 storage

  733.             VkPhysicalDevice16BitStorageFeaturesKHR query16BitStorageFeatures;

  734.             query16BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR;

  735.             query16BitStorageFeatures.pNext = 0;

  736.             if (gpu_info.support_VK_KHR_16bit_storage)

  737.             {

  738.                 query16BitStorageFeatures.pNext = queryExtensionFeatures;

  739.                 queryExtensionFeatures = &query16BitStorageFeatures;

  740.             }

  741. 

  742.             // query fp16/int8 arithmetic

  743.             VkPhysicalDeviceFloat16Int8FeaturesKHR queryFloat16Int8Features;

  744.             queryFloat16Int8Features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR;

  745.             queryFloat16Int8Features.pNext = 0;

  746.             if (gpu_info.support_VK_KHR_shader_float16_int8)

  747.             {

  748.                 queryFloat16Int8Features.pNext = queryExtensionFeatures;

  749.                 queryExtensionFeatures = &queryFloat16Int8Features;

  750.             }

  751. 

  752.             // query ycbcr_conversion

  753.             VkPhysicalDeviceSamplerYcbcrConversionFeaturesKHR querySamplerYcbcrConversionFeatures;

  754.             querySamplerYcbcrConversionFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SAMPLER_YCBCR_CONVERSION_FEATURES_KHR;

  755.             querySamplerYcbcrConversionFeatures.pNext = 0;

  756.             if (gpu_info.support_VK_KHR_sampler_ycbcr_conversion)

  757.             {

  758.                 querySamplerYcbcrConversionFeatures.pNext = queryExtensionFeatures;

  759.                 queryExtensionFeatures = &querySamplerYcbcrConversionFeatures;

  760.             }

  761. 

  762.             VkPhysicalDeviceFeatures2KHR queryFeatures;

  763.             queryFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2_KHR,

  764.             queryFeatures.pNext = queryExtensionFeatures;

  765. 

  766.             vkGetPhysicalDeviceFeatures2KHR(physicalDevice, &queryFeatures);

  767. 

  768.             if (gpu_info.support_VK_KHR_8bit_storage)

  769.             {

  770.                 gpu_info.support_int8_storage = query8BitStorageFeatures.storageBuffer8BitAccess && query8BitStorageFeatures.uniformAndStorageBuffer8BitAccess;

  771.             }

  772.             if (gpu_info.support_VK_KHR_16bit_storage)

  773.             {

  774.                 gpu_info.support_fp16_storage = query16BitStorageFeatures.storageBuffer16BitAccess && query16BitStorageFeatures.uniformAndStorageBuffer16BitAccess;

  775.             }

  776.             if (gpu_info.support_VK_KHR_shader_float16_int8)

  777.             {

  778.                 gpu_info.support_fp16_arithmetic = queryFloat16Int8Features.shaderFloat16;

  779.                 gpu_info.support_int8_arithmetic = queryFloat16Int8Features.shaderInt8;

  780.             }

  781.             if (gpu_info.support_VK_KHR_sampler_ycbcr_conversion)

  782.             {

  783.                 gpu_info.support_ycbcr_conversion = querySamplerYcbcrConversionFeatures.samplerYcbcrConversion;

  784.             }

  785.         }

  786.         else

  787.         {

  788. //             // TODO

  789. //             VkPhysicalDeviceFeatures features;

  790. //             vkGetPhysicalDeviceFeatures(physicalDevice, &features);

  791.         }

  792. 

  793.         if (physicalDeviceProperties.vendorID == 0x13b5)

  794.         {

  795.             // the 16bit_storage implementation of arm mali driver is buggy :[

  796.             gpu_info.support_fp16_storage = false;

  797.         }

  798. 

  799.         if (physicalDeviceProperties.vendorID == 0x10002 && physicalDeviceProperties.deviceID == 0x70006214 && physicalDeviceProperties.apiVersion == VK_MAKE_VERSION(1, 1, 82))

  800.         {

  801.             // the 16bit_storage implementation of vivante gc1700 driver is buggy :[

  802.             gpu_info.support_fp16_storage = false;

  803.         }

  804. 

  805.         if (gpu_info.bug_implicit_fp16_arithmetic)

  806.         {

  807.             // force capability on as long as the driver accept spirv with fp16 arithmetic :D

  808.             gpu_info.support_fp16_arithmetic = true;

  809.         }

  810. 

  811.         fprintf(stderr, "[%u %s]  queueC=%u[%u]  queueG=%u[%u]  queueT=%u[%u]\n", i, physicalDeviceProperties.deviceName,

  812.                 gpu_info.compute_queue_family_index, gpu_info.compute_queue_count,

  813.                 gpu_info.graphics_queue_family_index, gpu_info.graphics_queue_count,

  814.                 gpu_info.transfer_queue_family_index, gpu_info.transfer_queue_count);

  815. 

  816.         fprintf(stderr, "[%u %s]  buglssc=%d  bugihfa=%d\n", i, physicalDeviceProperties.deviceName,

  817.                 gpu_info.bug_local_size_spec_const, gpu_info.bug_implicit_fp16_arithmetic);

  818. 

  819.         fprintf(stderr, "[%u %s]  fp16p=%d  fp16s=%d  fp16a=%d  int8s=%d  int8a=%d\n", i, physicalDeviceProperties.deviceName,

  820.                 gpu_info.support_fp16_packed, gpu_info.support_fp16_storage, gpu_info.support_fp16_arithmetic,

  821.                 gpu_info.support_int8_storage, gpu_info.support_int8_arithmetic);

  822. 

  823.         gpu_info_index++;

  824.     }

  825. 

  826.     g_gpu_count = gpu_info_index;

  827. 

  828.     // the default gpu device

  829.     g_default_gpu_index = find_default_vulkan_device_index();

  830. 

  831.     // resolve shader info

  832.     for (int i=0; i<layer_shader_registry_entry_count; i++)

  833.     {

  834.         layer_shader_infos[i] = resolve_shader_info(layer_shader_registry[i].spv_data, layer_shader_registry[i].spv_data_size);

  835.     }

  836. 

  837.     return 0;

  838. }

  839. 

  840. void destroy_gpu_instance()

  841. {

  842.     for (int i=0; i<NCNN_MAX_GPU_COUNT; i++)

  843.     {

  844.         delete g_default_vkdev[i];

  845.         g_default_vkdev[i] = 0;

  846.     }

  847. 

  848. #if ENABLE_VALIDATION_LAYER

  849.     if (support_VK_EXT_debug_utils)

  850.     {

  851.         DestroyDebugUtilsMessengerEXT(g_instance, callback, NULL);

  852.     }

  853. #endif // ENABLE_VALIDATION_LAYER

  854. 

  855.     vkDestroyInstance(g_instance, 0);

  856. }

  857. 

  858. int get_gpu_count()

  859. {

  860.     return g_gpu_count;

  861. }

  862. 

  863. int get_default_gpu_index()

  864. {

  865.     return g_default_gpu_index;

  866. }

  867. 

  868. const GpuInfo& get_gpu_info(int device_index)

  869. {

  870.     return g_gpu_infos[device_index];

  871. }

  872. 

  873. VulkanDevice::VulkanDevice(int device_index) : info(g_gpu_infos[device_index])

  874. {

  875.     std::vector<const char*> enabledExtensions;

  876.     if (info.support_VK_KHR_8bit_storage)

  877.         enabledExtensions.push_back("VK_KHR_8bit_storage");

  878.     if (info.support_VK_KHR_16bit_storage)

  879.         enabledExtensions.push_back("VK_KHR_16bit_storage");

  880.     if (info.support_VK_KHR_bind_memory2)

  881.         enabledExtensions.push_back("VK_KHR_bind_memory2");

  882.     if (info.support_VK_KHR_dedicated_allocation)

  883.         enabledExtensions.push_back("VK_KHR_dedicated_allocation");

  884.     if (info.support_VK_KHR_descriptor_update_template)

  885.         enabledExtensions.push_back("VK_KHR_descriptor_update_template");

  886.     if (info.support_VK_KHR_external_memory)

  887.         enabledExtensions.push_back("VK_KHR_external_memory");

  888.     if (info.support_VK_KHR_get_memory_requirements2)

  889.         enabledExtensions.push_back("VK_KHR_get_memory_requirements2");

  890.     if (info.support_VK_KHR_maintenance1)

  891.         enabledExtensions.push_back("VK_KHR_maintenance1");

  892.     if (info.support_VK_KHR_push_descriptor)

  893.         enabledExtensions.push_back("VK_KHR_push_descriptor");

  894.     if (info.support_VK_KHR_sampler_ycbcr_conversion)

  895.         enabledExtensions.push_back("VK_KHR_sampler_ycbcr_conversion");

  896.     if (info.support_VK_KHR_shader_float16_int8)

  897.         enabledExtensions.push_back("VK_KHR_shader_float16_int8");

  898.     if (info.support_VK_KHR_shader_float_controls)

  899.         enabledExtensions.push_back("VK_KHR_shader_float_controls");

  900.     if (info.support_VK_KHR_storage_buffer_storage_class)

  901.         enabledExtensions.push_back("VK_KHR_storage_buffer_storage_class");

  902.     if (info.support_VK_KHR_swapchain)

  903.         enabledExtensions.push_back("VK_KHR_swapchain");

  904.     if (info.support_VK_EXT_queue_family_foreign)

  905.         enabledExtensions.push_back("VK_EXT_queue_family_foreign");

  906. #if __ANDROID_API__ >= 26

  907.     if (info.support_VK_ANDROID_external_memory_android_hardware_buffer)

  908.         enabledExtensions.push_back("VK_ANDROID_external_memory_android_hardware_buffer");

  909. #endif // __ANDROID_API__ >= 26

  910. 

  911.     void* enabledExtensionFeatures = 0;

  912. 

  913.     // enable int8 storage

  914.     VkPhysicalDevice8BitStorageFeaturesKHR enabled8BitStorageFeatures;

  915.     enabled8BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_8BIT_STORAGE_FEATURES_KHR;

  916.     enabled8BitStorageFeatures.pNext = 0;

  917.     enabled8BitStorageFeatures.storageBuffer8BitAccess = info.support_int8_storage;

  918.     enabled8BitStorageFeatures.uniformAndStorageBuffer8BitAccess = info.support_int8_storage;

  919.     enabled8BitStorageFeatures.storagePushConstant8 = VK_FALSE;

  920.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_8bit_storage)

  921.     {

  922.         enabled8BitStorageFeatures.pNext = enabledExtensionFeatures;

  923.         enabledExtensionFeatures = &enabled8BitStorageFeatures;

  924.     }

  925. 

  926.     // enable fp16/int16 storage

  927.     VkPhysicalDevice16BitStorageFeaturesKHR enabled16BitStorageFeatures;

  928.     enabled16BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR;

  929.     enabled16BitStorageFeatures.pNext = 0;

  930.     enabled16BitStorageFeatures.storageBuffer16BitAccess = info.support_fp16_storage;

  931.     enabled16BitStorageFeatures.uniformAndStorageBuffer16BitAccess = info.support_fp16_storage;

  932.     enabled16BitStorageFeatures.storagePushConstant16 = VK_FALSE;

  933.     enabled16BitStorageFeatures.storageInputOutput16 = VK_FALSE;

  934.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_16bit_storage)

  935.     {

  936.         enabled16BitStorageFeatures.pNext = enabledExtensionFeatures;

  937.         enabledExtensionFeatures = &enabled16BitStorageFeatures;

  938.     }

  939. 

  940.     // enable fp16/int8 arithmetic

  941.     VkPhysicalDeviceFloat16Int8FeaturesKHR enabledFloat16Int8Features;

  942.     enabledFloat16Int8Features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR;

  943.     enabledFloat16Int8Features.pNext = 0;

  944.     enabledFloat16Int8Features.shaderFloat16 = info.support_fp16_arithmetic;

  945.     enabledFloat16Int8Features.shaderInt8 = info.support_int8_arithmetic;

  946.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_shader_float16_int8)

  947.     {

  948.         enabledFloat16Int8Features.pNext = enabledExtensionFeatures;

  949.         enabledExtensionFeatures = &enabledFloat16Int8Features;

  950.     }

  951. 

  952.     // enable ycbcr conversion

  953.     VkPhysicalDeviceSamplerYcbcrConversionFeaturesKHR querySamplerYcbcrConversionFeatures;

  954.     querySamplerYcbcrConversionFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SAMPLER_YCBCR_CONVERSION_FEATURES_KHR;

  955.     querySamplerYcbcrConversionFeatures.pNext = 0;

  956.     querySamplerYcbcrConversionFeatures.samplerYcbcrConversion = info.support_ycbcr_conversion;

  957.     if (support_VK_KHR_get_physical_device_properties2 && info.support_ycbcr_conversion)

  958.     {

  959.         querySamplerYcbcrConversionFeatures.pNext = enabledExtensionFeatures;

  960.         enabledExtensionFeatures = &querySamplerYcbcrConversionFeatures;

  961.     }

  962. 

  963.     std::vector<float> compute_queue_priorities(info.compute_queue_count, 1.f);// 0.f ~ 1.f

  964.     std::vector<float> graphics_queue_priorities(info.graphics_queue_count, 1.f);// 0.f ~ 1.f

  965.     std::vector<float> transfer_queue_priorities(info.transfer_queue_count, 1.f);// 0.f ~ 1.f

  966. 

  967.     VkDeviceQueueCreateInfo deviceQueueCreateInfos[3];

  968. 

  969.     VkDeviceQueueCreateInfo deviceComputeQueueCreateInfo;

  970.     deviceComputeQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  971.     deviceComputeQueueCreateInfo.pNext = 0;

  972.     deviceComputeQueueCreateInfo.flags = 0;

  973.     deviceComputeQueueCreateInfo.queueFamilyIndex = info.compute_queue_family_index;

  974.     deviceComputeQueueCreateInfo.queueCount = info.compute_queue_count;

  975.     deviceComputeQueueCreateInfo.pQueuePriorities = compute_queue_priorities.data();

  976. 

  977.     VkDeviceQueueCreateInfo deviceGraphicsQueueCreateInfo;

  978.     deviceGraphicsQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  979.     deviceGraphicsQueueCreateInfo.pNext = 0;

  980.     deviceGraphicsQueueCreateInfo.flags = 0;

  981.     deviceGraphicsQueueCreateInfo.queueFamilyIndex = info.graphics_queue_family_index;

  982.     deviceGraphicsQueueCreateInfo.queueCount = info.graphics_queue_count;

  983.     deviceGraphicsQueueCreateInfo.pQueuePriorities = graphics_queue_priorities.data();

  984. 

  985.     VkDeviceQueueCreateInfo deviceTransferQueueCreateInfo;

  986.     deviceTransferQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  987.     deviceTransferQueueCreateInfo.pNext = 0;

  988.     deviceTransferQueueCreateInfo.flags = 0;

  989.     deviceTransferQueueCreateInfo.queueFamilyIndex = info.transfer_queue_family_index;

  990.     deviceTransferQueueCreateInfo.queueCount = info.transfer_queue_count;

  991.     deviceTransferQueueCreateInfo.pQueuePriorities = transfer_queue_priorities.data();

  992. 

  993.     VkDeviceCreateInfo deviceCreateInfo;

  994.     deviceCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;

  995.     deviceCreateInfo.pNext = enabledExtensionFeatures;

  996.     deviceCreateInfo.flags = 0;

  997.     if (info.compute_queue_family_index == info.graphics_queue_family_index && info.compute_queue_family_index == info.transfer_queue_family_index)

  998.     {

  999.     deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  1000.     deviceCreateInfo.queueCreateInfoCount = 1;

  1001.     }

  1002.     else if (info.compute_queue_family_index == info.graphics_queue_family_index && info.compute_queue_family_index != info.transfer_queue_family_index)

  1003.     {

  1004.     deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  1005.     deviceQueueCreateInfos[1] = deviceTransferQueueCreateInfo;

  1006.     deviceCreateInfo.queueCreateInfoCount = 2;

  1007.     }

  1008.     else if (info.compute_queue_family_index != info.graphics_queue_family_index && info.graphics_queue_family_index == info.transfer_queue_family_index)

  1009.     {

  1010.     deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  1011.     deviceQueueCreateInfos[1] = deviceGraphicsQueueCreateInfo;

  1012.     deviceCreateInfo.queueCreateInfoCount = 2;

  1013.     }

  1014.     else // if (info.compute_queue_family_index != info.graphics_queue_family_index && info.graphics_queue_family_index != info.transfer_queue_family_index)

  1015.     {

  1016.     deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  1017.     deviceQueueCreateInfos[1] = deviceGraphicsQueueCreateInfo;

  1018.     deviceQueueCreateInfos[2] = deviceTransferQueueCreateInfo;

  1019.     deviceCreateInfo.queueCreateInfoCount = 3;

  1020.     }

  1021.     deviceCreateInfo.pQueueCreateInfos = deviceQueueCreateInfos;

  1022.     deviceCreateInfo.enabledLayerCount = 0;

  1023.     deviceCreateInfo.ppEnabledLayerNames = 0;

  1024.     deviceCreateInfo.enabledExtensionCount = enabledExtensions.size();

  1025.     deviceCreateInfo.ppEnabledExtensionNames = enabledExtensions.data();

  1026.     deviceCreateInfo.pEnabledFeatures = 0;// VkPhysicalDeviceFeatures pointer

  1027. 

  1028.     VkResult ret = vkCreateDevice(info.physical_device, &deviceCreateInfo, 0, &device);

  1029.     if (ret != VK_SUCCESS)

  1030.     {

  1031.         fprintf(stderr, "vkCreateDevice failed %d\n", ret);

  1032.     }

  1033. 

  1034.     init_device_extension();

  1035. 

  1036.     create_shader_module();

  1037. 

  1038.     compute_queues.resize(info.compute_queue_count);

  1039.     blob_allocators.resize(info.compute_queue_count);

  1040.     staging_allocators.resize(info.compute_queue_count);

  1041.     for (uint32_t i = 0; i < info.compute_queue_count; i++)

  1042.     {

  1043.         vkGetDeviceQueue(device, info.compute_queue_family_index, i, &compute_queues[i]);

  1044.         blob_allocators[i] = new VkBlobBufferAllocator(this);

  1045.         staging_allocators[i] = new VkStagingBufferAllocator(this);

  1046.     }

  1047.     if (info.compute_queue_family_index != info.graphics_queue_family_index)

  1048.     {

  1049.         graphics_queues.resize(info.graphics_queue_count);

  1050.         for (uint32_t i = 0; i < info.graphics_queue_count; i++)

  1051.         {

  1052.             vkGetDeviceQueue(device, info.graphics_queue_family_index, i, &graphics_queues[i]);

  1053.         }

  1054.     }

  1055.     if (info.compute_queue_family_index != info.transfer_queue_family_index && info.graphics_queue_family_index != info.transfer_queue_family_index)

  1056.     {

  1057.         transfer_queues.resize(info.transfer_queue_count);

  1058.         for (uint32_t i = 0; i < info.transfer_queue_count; i++)

  1059.         {

  1060.             vkGetDeviceQueue(device, info.transfer_queue_family_index, i, &transfer_queues[i]);

  1061.         }

  1062.     }

  1063. }

  1064. 

  1065. VulkanDevice::~VulkanDevice()

  1066. {

  1067.     for (uint32_t i = 0; i < info.compute_queue_count; i++)

  1068.     {

  1069.         delete blob_allocators[i];

  1070.         delete staging_allocators[i];

  1071.     }

  1072.     blob_allocators.clear();

  1073.     staging_allocators.clear();

  1074. 

  1075.     destroy_shader_module();

  1076. 

  1077.     vkDestroyDevice(device, 0);

  1078. }

  1079. 

  1080. VkShaderModule VulkanDevice::get_shader_module(int shader_type_index) const

  1081. {

  1082.     if (shader_type_index < 0 || shader_type_index >= layer_shader_registry_entry_count)

  1083.     {

  1084.         fprintf(stderr, "no such shader module %d\n", shader_type_index);

  1085.         return 0;

  1086.     }

  1087. 

  1088.     return shader_modules[shader_type_index];

  1089. }

  1090. 

  1091. VkShaderModule VulkanDevice::create_shader_module(int shader_type_index, uint32_t local_size_x, uint32_t local_size_y, uint32_t local_size_z) const

  1092. {

  1093.     if (shader_type_index < 0 || shader_type_index >= layer_shader_registry_entry_count)

  1094.     {

  1095.         fprintf(stderr, "no such shader module %d\n", shader_type_index);

  1096.         return 0;

  1097.     }

  1098. 

  1099.     const uint32_t* spv_data = layer_shader_registry[shader_type_index].spv_data;

  1100.     size_t spv_data_size = layer_shader_registry[shader_type_index].spv_data_size;

  1101. 

  1102.     return compile_shader_module(spv_data, spv_data_size, local_size_x, local_size_y, local_size_z);

  1103. }

  1104. 

  1105. VkShaderModule VulkanDevice::compile_shader_module(const uint32_t* spv_data, size_t spv_data_size) const

  1106. {

  1107.     VkShaderModuleCreateInfo shaderModuleCreateInfo;

  1108.     shaderModuleCreateInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;

  1109.     shaderModuleCreateInfo.pNext = 0;

  1110.     shaderModuleCreateInfo.flags = 0;

  1111.     shaderModuleCreateInfo.codeSize = spv_data_size;

  1112.     shaderModuleCreateInfo.pCode = spv_data;

  1113. 

  1114.     VkShaderModule shader_module;

  1115.     VkResult ret = vkCreateShaderModule(device, &shaderModuleCreateInfo, 0, &shader_module);

  1116.     if (ret != VK_SUCCESS)

  1117.     {

  1118.         fprintf(stderr, "vkCreateShaderModule failed %d\n", ret);

  1119.         return 0;

  1120.     }

  1121. 

  1122.     return shader_module;

  1123. }

  1124. 

  1125. static void inject_local_size_xyz(const uint32_t* code, size_t size, uint32_t local_size_x, uint32_t local_size_y, uint32_t local_size_z, uint32_t* dstcode, size_t* dstsize)

  1126. {

  1127.     uint32_t local_size_x_id = -1;

  1128.     uint32_t local_size_y_id = -1;

  1129.     uint32_t local_size_z_id = -1;

  1130.     uint32_t gl_WorkGroupSize_id = -1;

  1131. 

  1132.     const uint32_t* p = code;

  1133.     uint32_t* dp = dstcode;

  1134. 

  1135.     // skip magic version generator bound schema

  1136.     memcpy(dp, p, 5 * sizeof(uint32_t));

  1137.     p += 5;

  1138.     dp += 5;

  1139. 

  1140.     // foreach op

  1141.     while ((const unsigned char*)p < (const unsigned char*)code + size)

  1142.     {

  1143.         uint32_t opcode = p[0];

  1144. 

  1145.         uint16_t wordcount = opcode >> 16;

  1146.         uint16_t op = opcode & 0xffff;

  1147. 

  1148.         if (op == 16) // OpExecutionMode

  1149.         {

  1150.             uint32_t mode = p[2];

  1151.             if (mode == 17) // LocalSize

  1152.             {

  1153.                 memcpy(dp, p, wordcount * sizeof(uint32_t));

  1154. 

  1155.                 // set local_size_xyz

  1156.                 dp[3] = local_size_x;

  1157.                 dp[4] = local_size_y;

  1158.                 dp[5] = local_size_z;

  1159. 

  1160.                 p += wordcount;

  1161.                 dp += wordcount;

  1162.                 continue;

  1163.             }

  1164.         }

  1165.         else if (op == 50) // OpSpecConstant

  1166.         {

  1167.             uint32_t id = p[2];

  1168.             if (id == local_size_x_id || id == local_size_y_id || id == local_size_z_id)

  1169.             {

  1170.                 p += wordcount;

  1171.                 continue;

  1172.             }

  1173.         }

  1174.         else if (op == 51) // OpSpecConstantComposite

  1175.         {

  1176.             uint32_t id = p[2];

  1177.             if (id == gl_WorkGroupSize_id)

  1178.             {

  1179.                 if (wordcount == 6 && (p[3] == local_size_x_id || p[4] == local_size_y_id || p[5] == local_size_z_id))

  1180.                 {

  1181.                     p += wordcount;

  1182.                     continue;

  1183.                 }

  1184.             }

  1185.         }

  1186.         else if (op == 71) // OpDecorate

  1187.         {

  1188.             uint32_t id = p[1];

  1189.             uint32_t decoration = p[2];

  1190.             if (decoration == 1) // SpecId

  1191.             {

  1192.                 uint32_t specid = p[3];

  1193.                 if (specid == 233) local_size_x_id = id;

  1194.                 if (specid == 234) local_size_y_id = id;

  1195.                 if (specid == 235) local_size_z_id = id;

  1196.                 if (specid == 233 || specid == 234 || specid == 235)

  1197.                 {

  1198.                     p += wordcount;

  1199.                     continue;

  1200.                 }

  1201.             }

  1202.             else if (decoration == 11) // BuiltIn

  1203.             {

  1204.                 uint32_t builtin = p[3];

  1205.                 if (builtin == 25) // WorkgroupSize

  1206.                 {

  1207.                     gl_WorkGroupSize_id = id;

  1208.                     p += wordcount;

  1209.                     continue;

  1210.                 }

  1211.             }

  1212.         }

  1213. 

  1214.         memcpy(dp, p, wordcount * sizeof(uint32_t));

  1215.         p += wordcount;

  1216.         dp += wordcount;

  1217.     }

  1218. 

  1219.     *dstsize = (unsigned char*)dp - (unsigned char*)dstcode;

  1220. }

  1221. 

  1222. VkShaderModule VulkanDevice::compile_shader_module(const uint32_t* spv_data, size_t spv_data_size, uint32_t local_size_x, uint32_t local_size_y, uint32_t local_size_z) const

  1223. {

  1224.     uint32_t* spv_data_modified = (uint32_t*)malloc(spv_data_size);

  1225.     size_t spv_data_size_modified = spv_data_size;

  1226.     inject_local_size_xyz(spv_data, spv_data_size, local_size_x, local_size_y, local_size_z, spv_data_modified, &spv_data_size_modified);

  1227. 

  1228.     VkShaderModule shader_module = compile_shader_module(spv_data_modified, spv_data_size_modified);

  1229. 

  1230.     free(spv_data_modified);

  1231. 

  1232.     return shader_module;

  1233. }

  1234. 

  1235. uint32_t VulkanDevice::find_memory_index(uint32_t memory_type_bits, VkFlags required, VkFlags preferred, VkFlags preferred_not) const

  1236. {

  1237.     // first try, find required and with preferred and without preferred_not

  1238.     for (uint32_t i=0; i<info.physicalDeviceMemoryProperties.memoryTypeCount; i++)

  1239.     {

  1240.         bool is_required = (1 << i) & memory_type_bits;

  1241.         if (is_required)

  1242.         {

  1243.             const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[i];

  1244.             if ((memoryType.propertyFlags & required) == required

  1245.                 && (preferred && (memoryType.propertyFlags & preferred))

  1246.                 && (preferred_not && !(memoryType.propertyFlags & preferred_not)))

  1247.             {

  1248.                 return i;

  1249.             }

  1250.         }

  1251.     }

  1252. 

  1253.     // second try, find required and with preferred

  1254.     for (uint32_t i=0; i<info.physicalDeviceMemoryProperties.memoryTypeCount; i++)

  1255.     {

  1256.         bool is_required = (1 << i) & memory_type_bits;

  1257.         if (is_required)

  1258.         {

  1259.             const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[i];

  1260.             if ((memoryType.propertyFlags & required) == required

  1261.                 && (preferred && (memoryType.propertyFlags & preferred)))

  1262.             {

  1263.                 return i;

  1264.             }

  1265.         }

  1266.     }

  1267. 

  1268.     // third try, find required and without preferred_not

  1269.     for (uint32_t i=0; i<info.physicalDeviceMemoryProperties.memoryTypeCount; i++)

  1270.     {

  1271.         bool is_required = (1 << i) & memory_type_bits;

  1272.         if (is_required)

  1273.         {

  1274.             const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[i];

  1275.             if ((memoryType.propertyFlags & required) == required

  1276.                 && (preferred_not && !(memoryType.propertyFlags & preferred_not)))

  1277.             {

  1278.                 return i;

  1279.             }

  1280.         }

  1281.     }

  1282. 

  1283.     // fourth try, find any required

  1284.     for (uint32_t i=0; i<info.physicalDeviceMemoryProperties.memoryTypeCount; i++)

  1285.     {

  1286.         bool is_required = (1 << i) & memory_type_bits;

  1287.         if (is_required)

  1288.         {

  1289.             const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[i];

  1290.             if ((memoryType.propertyFlags & required) == required)

  1291.             {

  1292.                 return i;

  1293.             }

  1294.         }

  1295.     }

  1296. 

  1297.     fprintf(stderr, "no such memory type %u %u %u %u\n", memory_type_bits, required, preferred, preferred_not);

  1298.     return -1;

  1299. }

  1300. 

  1301. bool VulkanDevice::is_mappable(uint32_t memory_type_index) const

  1302. {

  1303.     const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[memory_type_index];

  1304. 

  1305.     return memoryType.propertyFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT;

  1306. }

  1307. 

  1308. bool VulkanDevice::is_coherent(uint32_t memory_type_index) const

  1309. {

  1310.     const VkMemoryType& memoryType = info.physicalDeviceMemoryProperties.memoryTypes[memory_type_index];

  1311. 

  1312.     return memoryType.propertyFlags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT;

  1313. }

  1314. 

  1315. VkQueue VulkanDevice::acquire_queue(uint32_t queue_family_index) const

  1316. {

  1317.     if (queue_family_index != info.compute_queue_family_index

  1318.         && queue_family_index != info.graphics_queue_family_index

  1319.         && queue_family_index != info.transfer_queue_family_index)

  1320.     {

  1321.         fprintf(stderr, "invalid queue_family_index %u\n", queue_family_index);

  1322.         return 0;

  1323.     }

  1324. 

  1325.     MutexLockGuard lock(queue_lock);

  1326. 

  1327.     std::vector<VkQueue>& queues = queue_family_index == info.compute_queue_family_index ? compute_queues

  1328.                                  : queue_family_index == info.graphics_queue_family_index ? graphics_queues : transfer_queues;

  1329.     for (int i=0; i<(int)queues.size(); i++)

  1330.     {

  1331.         VkQueue queue = queues[i];

  1332.         if (queue)

  1333.         {

  1334.             queues[i] = 0;

  1335.             return queue;

  1336.         }

  1337.     }

  1338. 

  1339.     // out of hardware queue

  1340.     return 0;

  1341. }

  1342. 

  1343. void VulkanDevice::reclaim_queue(uint32_t queue_family_index, VkQueue queue) const

  1344. {

  1345.     if (queue_family_index != info.compute_queue_family_index

  1346.         && queue_family_index != info.graphics_queue_family_index

  1347.         && queue_family_index != info.transfer_queue_family_index)

  1348.     {

  1349.         fprintf(stderr, "invalid queue_family_index %u\n", queue_family_index);

  1350.         return;

  1351.     }

  1352. 

  1353.     MutexLockGuard lock(queue_lock);

  1354. 

  1355.     std::vector<VkQueue>& queues = queue_family_index == info.compute_queue_family_index ? compute_queues

  1356.                                  : queue_family_index == info.graphics_queue_family_index ? graphics_queues : transfer_queues;

  1357.     for (int i=0; i<(int)queues.size(); i++)

  1358.     {

  1359.         if (!queues[i])

  1360.         {

  1361.             queues[i] = queue;

  1362.             return;

  1363.         }

  1364.     }

  1365. 

  1366.     fprintf(stderr, "FATAL ERROR! reclaim_queue get wild queue %u %p\n", queue_family_index, queue);

  1367. }

  1368. 

  1369. VkAllocator* VulkanDevice::acquire_blob_allocator() const

  1370. {

  1371.     MutexLockGuard lock(blob_allocator_lock);

  1372. 

  1373.     for (int i=0; i<(int)blob_allocators.size(); i++)

  1374.     {

  1375.         VkAllocator* allocator = blob_allocators[i];

  1376.         if (allocator)

  1377.         {

  1378.             blob_allocators[i] = 0;

  1379.             return allocator;

  1380.         }

  1381.     }

  1382. 

  1383.     // out of blob allocator

  1384.     return 0;

  1385. }

  1386. 

  1387. void VulkanDevice::reclaim_blob_allocator(VkAllocator* allocator) const

  1388. {

  1389.     MutexLockGuard lock(blob_allocator_lock);

  1390. 

  1391.     for (int i=0; i<(int)blob_allocators.size(); i++)

  1392.     {

  1393.         if (!blob_allocators[i])

  1394.         {

  1395.             blob_allocators[i] = allocator;

  1396.             return;

  1397.         }

  1398.     }

  1399. 

  1400.     fprintf(stderr, "FATAL ERROR! reclaim_blob_allocator get wild allocator %p\n", allocator);

  1401. }

  1402. 

  1403. VkAllocator* VulkanDevice::acquire_staging_allocator() const

  1404. {

  1405.     MutexLockGuard lock(staging_allocator_lock);

  1406. 

  1407.     for (int i=0; i<(int)staging_allocators.size(); i++)

  1408.     {

  1409.         VkAllocator* allocator = staging_allocators[i];

  1410.         if (allocator)

  1411.         {

  1412.             staging_allocators[i] = 0;

  1413.             return allocator;

  1414.         }

  1415.     }

  1416. 

  1417.     // out of staging allocator

  1418.     return 0;

  1419. }

  1420. 

  1421. void VulkanDevice::reclaim_staging_allocator(VkAllocator* allocator) const

  1422. {

  1423.     MutexLockGuard lock(staging_allocator_lock);

  1424. 

  1425.     for (int i=0; i<(int)staging_allocators.size(); i++)

  1426.     {

  1427.         if (!staging_allocators[i])

  1428.         {

  1429.             staging_allocators[i] = allocator;

  1430.             return;

  1431.         }

  1432.     }

  1433. 

  1434.     fprintf(stderr, "FATAL ERROR! reclaim_staging_allocator get wild allocator %p\n", allocator);

  1435. }

  1436. 

  1437. static inline bool string_ends_with_fp16p(const char* name)

  1438. {

  1439.     int len = strlen(name);

  1440.     if (len < 6)

  1441.         return false;

  1442. 

  1443.     return memcmp(name + len - 6, "_fp16p", 6) == 0;

  1444. }

  1445. 

  1446. static inline bool string_ends_with_fp16pa(const char* name)

  1447. {

  1448.     int len = strlen(name);

  1449.     if (len < 7)

  1450.         return false;

  1451. 

  1452.     return memcmp(name + len - 7, "_fp16pa", 7) == 0;

  1453. }

  1454. 

  1455. static inline bool string_ends_with_fp16s(const char* name)

  1456. {

  1457.     int len = strlen(name);

  1458.     if (len < 6)

  1459.         return false;

  1460. 

  1461.     return memcmp(name + len - 6, "_fp16s", 6) == 0;

  1462. }

  1463. 

  1464. static inline bool string_ends_with_fp16sa(const char* name)

  1465. {

  1466.     int len = strlen(name);

  1467.     if (len < 7)

  1468.         return false;

  1469. 

  1470.     return memcmp(name + len - 7, "_fp16sa", 7) == 0;

  1471. }

  1472. 

  1473. int VulkanDevice::create_shader_module()

  1474. {

  1475.     if (info.bug_local_size_spec_const)

  1476.     {

  1477.         // do not cache shader module

  1478.         return 0;

  1479.     }

  1480. 

  1481.     shader_modules.resize(layer_shader_registry_entry_count, VK_NULL_HANDLE);

  1482. 

  1483.     for (int i=0; i<layer_shader_registry_entry_count; i++)

  1484.     {

  1485.         // ncnn_add_shader cmake macro

  1486.         // 0 = fp32

  1487.         // 1 = fp16p

  1488.         // 2 = fp16pa

  1489.         // 3 = fp16s

  1490.         // 4 = fp16sa

  1491. 

  1492.         if (!info.support_fp16_packed)

  1493.         {

  1494.             if (i % 5 == 1)

  1495.                 continue;

  1496.         }

  1497. 

  1498.         if (!info.support_fp16_packed || !info.support_fp16_arithmetic)

  1499.         {

  1500.             if (i % 5 == 2)

  1501.                 continue;

  1502.         }

  1503. 

  1504.         if (!info.support_fp16_storage)

  1505.         {

  1506.             if (i % 5 == 3)

  1507.                 continue;

  1508.         }

  1509. 

  1510.         if (!info.support_fp16_storage || !info.support_fp16_arithmetic)

  1511.         {

  1512.             if (i % 5 == 4)

  1513.                 continue;

  1514.         }

  1515. 

  1516.         const uint32_t* spv_data = layer_shader_registry[i].spv_data;

  1517.         size_t spv_data_size = layer_shader_registry[i].spv_data_size;

  1518. 

  1519.         VkShaderModule shader_module = compile_shader_module(spv_data, spv_data_size);

  1520.         if (shader_module == 0)

  1521.         {

  1522.             fprintf(stderr, "compile_shader_module %d failed\n", i);

  1523.             return -1;

  1524.         }

  1525. 

  1526.         shader_modules[i] = shader_module;

  1527. 

  1528. //         fprintf(stderr, "shader_module %d created\n", i);

  1529.     }

  1530. 

  1531.     return 0;

  1532. }

  1533. 

  1534. void VulkanDevice::destroy_shader_module()

  1535. {

  1536.     for (int i=0; i<(int)shader_modules.size(); i++)

  1537.     {

  1538.         vkDestroyShaderModule(device, shader_modules[i], 0);

  1539.     }

  1540. 

  1541.     shader_modules.clear();

  1542. }

  1543. 

  1544. int VulkanDevice::init_device_extension()

  1545. {

  1546.     if (info.support_VK_KHR_bind_memory2)

  1547.     {

  1548.         vkBindBufferMemory2KHR = (PFN_vkBindBufferMemory2KHR)vkGetDeviceProcAddr(device, "vkBindBufferMemory2KHR");

  1549.         vkBindImageMemory2KHR = (PFN_vkBindImageMemory2KHR)vkGetDeviceProcAddr(device, "vkBindImageMemory2KHR");

  1550.     }

  1551. 

  1552.     if (info.support_VK_KHR_descriptor_update_template)

  1553.     {

  1554.         vkCreateDescriptorUpdateTemplateKHR = (PFN_vkCreateDescriptorUpdateTemplateKHR)vkGetDeviceProcAddr(device, "vkCreateDescriptorUpdateTemplateKHR");

  1555.         vkDestroyDescriptorUpdateTemplateKHR = (PFN_vkDestroyDescriptorUpdateTemplateKHR)vkGetDeviceProcAddr(device, "vkDestroyDescriptorUpdateTemplateKHR");

  1556.         vkUpdateDescriptorSetWithTemplateKHR = (PFN_vkUpdateDescriptorSetWithTemplateKHR)vkGetDeviceProcAddr(device, "vkUpdateDescriptorSetWithTemplateKHR");

  1557.     }

  1558. 

  1559.     if (info.support_VK_KHR_get_memory_requirements2)

  1560.     {

  1561.         vkGetImageMemoryRequirements2KHR = (PFN_vkGetImageMemoryRequirements2KHR)vkGetDeviceProcAddr(device, "vkGetImageMemoryRequirements2KHR");

  1562.         vkGetBufferMemoryRequirements2KHR = (PFN_vkGetBufferMemoryRequirements2KHR)vkGetDeviceProcAddr(device, "vkGetBufferMemoryRequirements2KHR");

  1563.         vkGetImageSparseMemoryRequirements2KHR = (PFN_vkGetImageSparseMemoryRequirements2KHR)vkGetDeviceProcAddr(device, "vkGetImageSparseMemoryRequirements2KHR");

  1564.     }

  1565. 

  1566.     if (info.support_VK_KHR_maintenance1)

  1567.     {

  1568.         vkTrimCommandPoolKHR = (PFN_vkTrimCommandPoolKHR)vkGetDeviceProcAddr(device, "vkTrimCommandPoolKHR");

  1569.     }

  1570. 

  1571.     if (info.support_VK_KHR_push_descriptor)

  1572.     {

  1573.         if (info.support_VK_KHR_descriptor_update_template)

  1574.         {

  1575.             vkCmdPushDescriptorSetWithTemplateKHR = (PFN_vkCmdPushDescriptorSetWithTemplateKHR)vkGetDeviceProcAddr(device, "vkCmdPushDescriptorSetWithTemplateKHR");

  1576.         }

  1577. 

  1578.         vkCmdPushDescriptorSetKHR = (PFN_vkCmdPushDescriptorSetKHR)vkGetDeviceProcAddr(device, "vkCmdPushDescriptorSetKHR");

  1579.     }

  1580. 

  1581.     if (info.support_VK_KHR_sampler_ycbcr_conversion)

  1582.     {

  1583.         vkCreateSamplerYcbcrConversionKHR = (PFN_vkCreateSamplerYcbcrConversionKHR)vkGetDeviceProcAddr(device, "vkCreateSamplerYcbcrConversionKHR");

  1584.         vkDestroySamplerYcbcrConversionKHR = (PFN_vkDestroySamplerYcbcrConversionKHR)vkGetDeviceProcAddr(device, "vkDestroySamplerYcbcrConversionKHR");

  1585.     }

  1586. 

  1587.     if (info.support_VK_KHR_swapchain)

  1588.     {

  1589.         vkCreateSwapchainKHR = (PFN_vkCreateSwapchainKHR)vkGetDeviceProcAddr(device, "vkCreateSwapchainKHR");

  1590.         vkDestroySwapchainKHR = (PFN_vkDestroySwapchainKHR)vkGetDeviceProcAddr(device, "vkDestroySwapchainKHR");

  1591.         vkGetSwapchainImagesKHR = (PFN_vkGetSwapchainImagesKHR)vkGetDeviceProcAddr(device, "vkGetSwapchainImagesKHR");

  1592.         vkAcquireNextImageKHR = (PFN_vkAcquireNextImageKHR)vkGetDeviceProcAddr(device, "vkAcquireNextImageKHR");

  1593.         vkQueuePresentKHR = (PFN_vkQueuePresentKHR)vkGetDeviceProcAddr(device, "vkQueuePresentKHR");

  1594.     }

  1595. 

  1596. #if __ANDROID_API__ >= 26

  1597.     if (info.support_VK_ANDROID_external_memory_android_hardware_buffer)

  1598.     {

  1599.         vkGetAndroidHardwareBufferPropertiesANDROID = (PFN_vkGetAndroidHardwareBufferPropertiesANDROID)vkGetDeviceProcAddr(device, "vkGetAndroidHardwareBufferPropertiesANDROID");

  1600.         vkGetMemoryAndroidHardwareBufferANDROID = (PFN_vkGetMemoryAndroidHardwareBufferANDROID)vkGetDeviceProcAddr(device, "vkGetMemoryAndroidHardwareBufferANDROID");

  1601.     }

  1602. #endif // __ANDROID_API__ >= 26

  1603. 

  1604.     return 0;

  1605. }

  1606. 

  1607. VulkanDevice* get_gpu_device(int device_index)

  1608. {

  1609.     if (device_index < 0 || device_index >= g_gpu_count)

  1610.         return 0;

  1611. 

  1612.     MutexLockGuard lock(g_default_vkdev_lock);

  1613. 

  1614.     if (!g_default_vkdev[device_index])

  1615.         g_default_vkdev[device_index] = new VulkanDevice(device_index);

  1616. 

  1617.     return g_default_vkdev[device_index];

  1618. }

  1619. 

  1620. const ShaderInfo& get_shader_info(int shader_type_index)

  1621. {

  1622.     if (shader_type_index < 0 || shader_type_index >= layer_shader_registry_entry_count)

  1623.     {

  1624.         fprintf(stderr, "no such shader module %d\n", shader_type_index);

  1625.         return layer_shader_infos[0];

  1626.     }

  1627. 

  1628.     return layer_shader_infos[shader_type_index];

  1629. }

  1630. 

  1631. ShaderInfo resolve_shader_info(const uint32_t* spv_data, size_t spv_data_size)

  1632. {

  1633.     uint32_t parameter_id = -233;

  1634. 

  1635.     int specialization_count = 0;

  1636.     int binding_count = 0;

  1637.     int push_constant_count = 0;

  1638. 

  1639.     const uint32_t* p = spv_data;

  1640. 

  1641.     // skip magic version generator bound schema

  1642.     p += 5;

  1643. 

  1644.     // foreach op

  1645.     while ((const unsigned char*)p < (const unsigned char*)spv_data + spv_data_size)

  1646.     {

  1647.         uint32_t opcode = p[0];

  1648. 

  1649.         uint16_t wordcount = opcode >> 16;

  1650.         uint16_t op = opcode & 0xffff;

  1651. 

  1652.         if (op == 5) // OpName

  1653.         {

  1654.             uint32_t id = p[1];

  1655.             const char* name = (const char*)&p[2];

  1656.             if (strcmp(name, "parameter") == 0)

  1657.             {

  1658.                 parameter_id = id;

  1659.             }

  1660.         }

  1661.         else if (op == 6) // OpMemberName

  1662.         {

  1663.             uint32_t id = p[1];

  1664.             if (id == parameter_id)

  1665.             {

  1666.                 push_constant_count++;

  1667.             }

  1668.         }

  1669.         else if (op == 71) // OpDecorate

  1670.         {

  1671.             uint32_t decoration = p[2];

  1672.             if (decoration == 1) // SpecId

  1673.             {

  1674.                 specialization_count++;

  1675.             }

  1676.             else if (decoration == 33) // Binding

  1677.             {

  1678.                 binding_count++;

  1679.             }

  1680.         }

  1681. 

  1682.         p += wordcount;

  1683.     }

  1684. 

  1685.     ShaderInfo si;

  1686.     si.specialization_count = specialization_count;

  1687.     si.binding_count = binding_count;

  1688.     si.push_constant_count = push_constant_count;

  1689. 

  1690.     return si;

  1691. }

  1692. 

  1693. } // namespace ncnn

  1694. 

  1695. #endif // NCNN_VULKAN