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: bbf2e5d533
Branches Tags
${ item.name }
Create branch ${ searchTerm }
from 'bbf2e5d533'
${ noResults }
ncnn/src/gpu.cpp

4551 lines
202 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 <stdlib.h>

  20. #include <string.h>

  21. 

  22. #include "glslang/SPIRV/GlslangToSpv.h"

  23. #if NCNN_SYSTEM_GLSLANG

  24. #include "glslang/Public/ShaderLang.h"

  25. #else

  26. #include "glslang/glslang/Public/ShaderLang.h"

  27. #endif

  28. 

  29. #include "layer/vulkan/shader/vulkan_activation.comp.hex.h"

  30. 

  31. #include "command.h"

  32. #include "layer.h"

  33. #include "layer/vulkan/packing_vulkan.h"

  34. #include "layer_type.h"

  35. #include "mat.h"

  36. #include "pipelinecache.h"

  37. 

  38. // There is known issue that vkDestroyDebugUtilsMessengerEXT crash on exit when vulkan validation layer enabled

  39. // upstream fix https://github.com/KhronosGroup/Vulkan-Loader/pull/539

  40. #define ENABLE_VALIDATION_LAYER 0

  41. 

  42. namespace ncnn {

  43. 

  44. // global

  45. static Mutex g_instance_lock;

  46. 

  47. class __ncnn_vulkan_instance_holder

  48. {

  49. public:

  50.     __ncnn_vulkan_instance_holder()

  51.     {

  52.         instance = 0;

  53.         created = 0;

  54.         glslang_initialized = false;

  55. 

  56. #if NCNN_VULKAN_LOADER

  57.         libvulkan = 0;

  58. #if defined __ANDROID__

  59.         hvkdi = 0;

  60. #endif

  61. #endif // NCNN_VULKAN_LOADER

  62. 

  63. #if ENABLE_VALIDATION_LAYER

  64.         callback = 0;

  65. #endif

  66.     }

  67. 

  68.     ~__ncnn_vulkan_instance_holder()

  69.     {

  70.         destroy_gpu_instance();

  71.     }

  72. 

  73.     operator VkInstance()

  74.     {

  75.         return instance;

  76.     }

  77. 

  78.     VkInstance instance;

  79.     int created;

  80.     bool glslang_initialized;

  81. 

  82. #if ENABLE_VALIDATION_LAYER

  83.     VkDebugUtilsMessengerEXT callback;

  84. #endif

  85. };

  86. static __ncnn_vulkan_instance_holder g_instance;

  87. 

  88. static int g_gpu_count = 0;

  89. static int g_default_gpu_index = -1;

  90. 

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

  92. #define NCNN_MAX_GPU_COUNT 8

  93. static GpuInfo* g_gpu_infos[NCNN_MAX_GPU_COUNT] = {0};

  94. 

  95. // default vulkan device

  96. static Mutex g_default_vkdev_lock;

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

  98. 

  99. struct layer_shader_registry_entry

  100. {

  101.     const char* comp_data;

  102.     int comp_data_size;

  103. };

  104. 

  105. #include "layer_shader_spv_data.h"

  106. 

  107. static const layer_shader_registry_entry layer_shader_registry[] = {

  108. #include "layer_shader_registry.h"

  109. };

  110. 

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

  112. 

  113. // vulkan core

  114. PFN_vkAllocateCommandBuffers vkAllocateCommandBuffers = 0;

  115. PFN_vkAllocateDescriptorSets vkAllocateDescriptorSets = 0;

  116. PFN_vkAllocateMemory vkAllocateMemory = 0;

  117. PFN_vkBeginCommandBuffer vkBeginCommandBuffer = 0;

  118. PFN_vkBindBufferMemory vkBindBufferMemory = 0;

  119. PFN_vkBindImageMemory vkBindImageMemory = 0;

  120. PFN_vkCmdBeginQuery vkCmdBeginQuery = 0;

  121. PFN_vkCmdBindDescriptorSets vkCmdBindDescriptorSets = 0;

  122. PFN_vkCmdBindIndexBuffer vkCmdBindIndexBuffer = 0;

  123. PFN_vkCmdBindPipeline vkCmdBindPipeline = 0;

  124. PFN_vkCmdCopyBuffer vkCmdCopyBuffer = 0;

  125. PFN_vkCmdCopyBufferToImage vkCmdCopyBufferToImage = 0;

  126. PFN_vkCmdCopyImage vkCmdCopyImage = 0;

  127. PFN_vkCmdCopyImageToBuffer vkCmdCopyImageToBuffer = 0;

  128. PFN_vkCmdCopyQueryPoolResults vkCmdCopyQueryPoolResults = 0;

  129. PFN_vkCmdDispatch vkCmdDispatch = 0;

  130. PFN_vkCmdDispatchIndirect vkCmdDispatchIndirect = 0;

  131. PFN_vkCmdEndQuery vkCmdEndQuery = 0;

  132. PFN_vkCmdExecuteCommands vkCmdExecuteCommands = 0;

  133. PFN_vkCmdFillBuffer vkCmdFillBuffer = 0;

  134. PFN_vkCmdPipelineBarrier vkCmdPipelineBarrier = 0;

  135. PFN_vkCmdPushConstants vkCmdPushConstants = 0;

  136. PFN_vkCmdResetQueryPool vkCmdResetQueryPool = 0;

  137. PFN_vkCmdResolveImage vkCmdResolveImage = 0;

  138. PFN_vkCmdUpdateBuffer vkCmdUpdateBuffer = 0;

  139. PFN_vkCmdWriteTimestamp vkCmdWriteTimestamp = 0;

  140. PFN_vkCreateBuffer vkCreateBuffer = 0;

  141. PFN_vkCreateBufferView vkCreateBufferView = 0;

  142. PFN_vkCreateCommandPool vkCreateCommandPool = 0;

  143. PFN_vkCreateComputePipelines vkCreateComputePipelines = 0;

  144. PFN_vkCreateDescriptorPool vkCreateDescriptorPool = 0;

  145. PFN_vkCreateDescriptorSetLayout vkCreateDescriptorSetLayout = 0;

  146. PFN_vkCreateDevice vkCreateDevice = 0;

  147. PFN_vkCreateFence vkCreateFence = 0;

  148. PFN_vkCreateImage vkCreateImage = 0;

  149. PFN_vkCreateImageView vkCreateImageView = 0;

  150. PFN_vkCreatePipelineCache vkCreatePipelineCache = 0;

  151. PFN_vkCreatePipelineLayout vkCreatePipelineLayout = 0;

  152. PFN_vkCreateQueryPool vkCreateQueryPool = 0;

  153. PFN_vkCreateSampler vkCreateSampler = 0;

  154. PFN_vkCreateSemaphore vkCreateSemaphore = 0;

  155. PFN_vkCreateShaderModule vkCreateShaderModule = 0;

  156. PFN_vkDestroyBuffer vkDestroyBuffer = 0;

  157. PFN_vkDestroyBufferView vkDestroyBufferView = 0;

  158. PFN_vkDestroyCommandPool vkDestroyCommandPool = 0;

  159. PFN_vkDestroyDescriptorPool vkDestroyDescriptorPool = 0;

  160. PFN_vkDestroyDescriptorSetLayout vkDestroyDescriptorSetLayout = 0;

  161. PFN_vkDestroyDevice vkDestroyDevice = 0;

  162. PFN_vkDestroyFence vkDestroyFence = 0;

  163. PFN_vkDestroyImage vkDestroyImage = 0;

  164. PFN_vkDestroyImageView vkDestroyImageView = 0;

  165. PFN_vkDestroyInstance vkDestroyInstance = 0;

  166. PFN_vkDestroyPipeline vkDestroyPipeline = 0;

  167. PFN_vkDestroyPipelineCache vkDestroyPipelineCache = 0;

  168. PFN_vkDestroyPipelineLayout vkDestroyPipelineLayout = 0;

  169. PFN_vkDestroyQueryPool vkDestroyQueryPool = 0;

  170. PFN_vkDestroySampler vkDestroySampler = 0;

  171. PFN_vkDestroySemaphore vkDestroySemaphore = 0;

  172. PFN_vkDestroyShaderModule vkDestroyShaderModule = 0;

  173. PFN_vkDeviceWaitIdle vkDeviceWaitIdle = 0;

  174. PFN_vkEndCommandBuffer vkEndCommandBuffer = 0;

  175. PFN_vkEnumerateDeviceExtensionProperties vkEnumerateDeviceExtensionProperties = 0;

  176. PFN_vkEnumerateDeviceLayerProperties vkEnumerateDeviceLayerProperties = 0;

  177. PFN_vkEnumeratePhysicalDevices vkEnumeratePhysicalDevices = 0;

  178. PFN_vkFlushMappedMemoryRanges vkFlushMappedMemoryRanges = 0;

  179. PFN_vkFreeCommandBuffers vkFreeCommandBuffers = 0;

  180. PFN_vkFreeDescriptorSets vkFreeDescriptorSets = 0;

  181. PFN_vkFreeMemory vkFreeMemory = 0;

  182. PFN_vkGetBufferMemoryRequirements vkGetBufferMemoryRequirements = 0;

  183. PFN_vkGetDeviceMemoryCommitment vkGetDeviceMemoryCommitment = 0;

  184. PFN_vkGetDeviceProcAddr vkGetDeviceProcAddr = 0;

  185. PFN_vkGetDeviceQueue vkGetDeviceQueue = 0;

  186. PFN_vkGetFenceStatus vkGetFenceStatus = 0;

  187. PFN_vkGetImageMemoryRequirements vkGetImageMemoryRequirements = 0;

  188. PFN_vkGetImageSubresourceLayout vkGetImageSubresourceLayout = 0;

  189. PFN_vkGetPhysicalDeviceFeatures vkGetPhysicalDeviceFeatures = 0;

  190. PFN_vkGetPhysicalDeviceFormatProperties vkGetPhysicalDeviceFormatProperties = 0;

  191. PFN_vkGetPhysicalDeviceImageFormatProperties vkGetPhysicalDeviceImageFormatProperties = 0;

  192. PFN_vkGetPhysicalDeviceMemoryProperties vkGetPhysicalDeviceMemoryProperties = 0;

  193. PFN_vkGetPhysicalDeviceProperties vkGetPhysicalDeviceProperties = 0;

  194. PFN_vkGetPhysicalDeviceQueueFamilyProperties vkGetPhysicalDeviceQueueFamilyProperties = 0;

  195. PFN_vkGetPipelineCacheData vkGetPipelineCacheData = 0;

  196. PFN_vkGetQueryPoolResults vkGetQueryPoolResults = 0;

  197. PFN_vkInvalidateMappedMemoryRanges vkInvalidateMappedMemoryRanges = 0;

  198. PFN_vkMapMemory vkMapMemory = 0;

  199. PFN_vkMergePipelineCaches vkMergePipelineCaches = 0;

  200. PFN_vkQueueSubmit vkQueueSubmit = 0;

  201. PFN_vkQueueWaitIdle vkQueueWaitIdle = 0;

  202. PFN_vkResetCommandBuffer vkResetCommandBuffer = 0;

  203. PFN_vkResetCommandPool vkResetCommandPool = 0;

  204. PFN_vkResetDescriptorPool vkResetDescriptorPool = 0;

  205. PFN_vkResetFences vkResetFences = 0;

  206. PFN_vkUnmapMemory vkUnmapMemory = 0;

  207. PFN_vkUpdateDescriptorSets vkUpdateDescriptorSets = 0;

  208. PFN_vkWaitForFences vkWaitForFences = 0;

  209. 

  210. int support_VK_KHR_external_memory_capabilities = 0;

  211. int support_VK_KHR_get_physical_device_properties2 = 0;

  212. int support_VK_KHR_get_surface_capabilities2 = 0;

  213. int support_VK_KHR_portability_enumeration = 0;

  214. int support_VK_KHR_surface = 0;

  215. int support_VK_EXT_debug_utils = 0;

  216. int support_VK_EXT_validation_features = 0;

  217. int support_VK_EXT_validation_flags = 0;

  218. #if __ANDROID_API__ >= 26

  219. int support_VK_KHR_android_surface = 0;

  220. #endif // __ANDROID_API__ >= 26

  221. 

  222. // VK_KHR_cooperative_matrix

  223. PFN_vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR = 0;

  224. 

  225. // VK_KHR_external_memory_capabilities

  226. PFN_vkGetPhysicalDeviceExternalBufferPropertiesKHR vkGetPhysicalDeviceExternalBufferPropertiesKHR = 0;

  227. 

  228. // VK_KHR_get_physical_device_properties2

  229. PFN_vkGetPhysicalDeviceFeatures2KHR vkGetPhysicalDeviceFeatures2KHR = 0;

  230. PFN_vkGetPhysicalDeviceProperties2KHR vkGetPhysicalDeviceProperties2KHR = 0;

  231. PFN_vkGetPhysicalDeviceFormatProperties2KHR vkGetPhysicalDeviceFormatProperties2KHR = 0;

  232. PFN_vkGetPhysicalDeviceImageFormatProperties2KHR vkGetPhysicalDeviceImageFormatProperties2KHR = 0;

  233. PFN_vkGetPhysicalDeviceQueueFamilyProperties2KHR vkGetPhysicalDeviceQueueFamilyProperties2KHR = 0;

  234. PFN_vkGetPhysicalDeviceMemoryProperties2KHR vkGetPhysicalDeviceMemoryProperties2KHR = 0;

  235. 

  236. // VK_KHR_get_surface_capabilities2

  237. PFN_vkGetPhysicalDeviceSurfaceCapabilities2KHR vkGetPhysicalDeviceSurfaceCapabilities2KHR = 0;

  238. PFN_vkGetPhysicalDeviceSurfaceFormats2KHR vkGetPhysicalDeviceSurfaceFormats2KHR = 0;

  239. 

  240. // VK_KHR_surface

  241. PFN_vkDestroySurfaceKHR vkDestroySurfaceKHR = 0;

  242. PFN_vkGetPhysicalDeviceSurfaceSupportKHR vkGetPhysicalDeviceSurfaceSupportKHR = 0;

  243. PFN_vkGetPhysicalDeviceSurfaceCapabilitiesKHR vkGetPhysicalDeviceSurfaceCapabilitiesKHR = 0;

  244. PFN_vkGetPhysicalDeviceSurfaceFormatsKHR vkGetPhysicalDeviceSurfaceFormatsKHR = 0;

  245. PFN_vkGetPhysicalDeviceSurfacePresentModesKHR vkGetPhysicalDeviceSurfacePresentModesKHR = 0;

  246. 

  247. #if __ANDROID_API__ >= 26

  248. // VK_KHR_android_surface

  249. PFN_vkCreateAndroidSurfaceKHR vkCreateAndroidSurfaceKHR = 0;

  250. #endif // __ANDROID_API__ >= 26

  251. 

  252. // VK_NV_cooperative_matrix

  253. PFN_vkGetPhysicalDeviceCooperativeMatrixPropertiesNV vkGetPhysicalDeviceCooperativeMatrixPropertiesNV = 0;

  254. 

  255. class GpuInfoPrivate

  256. {

  257. public:

  258.     // vulkan physical device

  259.     VkPhysicalDevice physical_device;

  260. 

  261.     // memory properties

  262.     VkPhysicalDeviceMemoryProperties physical_device_memory_properties;

  263. 

  264.     // info

  265.     uint32_t api_version;

  266.     uint32_t driver_version;

  267.     uint32_t vendor_id;

  268.     uint32_t device_id;

  269.     char device_name[VK_MAX_PHYSICAL_DEVICE_NAME_SIZE];

  270.     uint8_t pipeline_cache_uuid[VK_UUID_SIZE];

  271. 

  272.     // 0 = discrete gpu

  273.     // 1 = integrated gpu

  274.     // 2 = virtual gpu

  275.     // 3 = cpu

  276.     int type;

  277. 

  278.     // hardware limit

  279.     uint32_t max_shared_memory_size;

  280.     uint32_t max_workgroup_count_x;

  281.     uint32_t max_workgroup_count_y;

  282.     uint32_t max_workgroup_count_z;

  283.     uint32_t max_workgroup_invocations;

  284.     uint32_t max_workgroup_size_x;

  285.     uint32_t max_workgroup_size_y;

  286.     uint32_t max_workgroup_size_z;

  287.     size_t memory_map_alignment;

  288.     size_t buffer_offset_alignment;

  289.     size_t non_coherent_atom_size;

  290.     size_t buffer_image_granularity;

  291.     uint32_t max_image_dimension_1d;

  292.     uint32_t max_image_dimension_2d;

  293.     uint32_t max_image_dimension_3d;

  294.     float timestamp_period;

  295. 

  296.     // runtime

  297.     uint32_t compute_queue_family_index;

  298.     uint32_t graphics_queue_family_index;

  299.     uint32_t transfer_queue_family_index;

  300. 

  301.     uint32_t compute_queue_count;

  302.     uint32_t graphics_queue_count;

  303.     uint32_t transfer_queue_count;

  304. 

  305.     // property

  306.     bool unified_compute_transfer_queue;

  307. 

  308.     // subgroup

  309.     uint32_t subgroup_size;

  310.     bool support_subgroup_basic;

  311.     bool support_subgroup_vote;

  312.     bool support_subgroup_ballot;

  313.     bool support_subgroup_shuffle;

  314. 

  315.     // bug is not feature

  316.     bool bug_storage_buffer_no_l1;

  317.     bool bug_corrupted_online_pipeline_cache;

  318.     bool bug_buffer_image_load_zero;

  319. 

  320.     // but sometimes bug is a feature

  321.     bool bug_implicit_fp16_arithmetic;

  322. 

  323.     // fp16 and int8 feature

  324.     bool support_fp16_packed;

  325.     bool support_fp16_storage;

  326.     bool support_fp16_uniform;

  327.     bool support_fp16_arithmetic;

  328.     bool support_int8_packed;

  329.     bool support_int8_storage;

  330.     bool support_int8_uniform;

  331.     bool support_int8_arithmetic;

  332. 

  333.     // ycbcr conversion feature

  334.     bool support_ycbcr_conversion;

  335. 

  336.     // cooperative matrix

  337.     bool support_cooperative_matrix;

  338.     bool support_cooperative_matrix_8_8_16;

  339.     bool support_cooperative_matrix_16_8_8;

  340.     bool support_cooperative_matrix_16_8_16;

  341.     bool support_cooperative_matrix_16_16_16;

  342. 

  343.     // extension capability

  344.     int support_VK_KHR_8bit_storage;

  345.     int support_VK_KHR_16bit_storage;

  346.     int support_VK_KHR_bind_memory2;

  347.     int support_VK_KHR_buffer_device_address;

  348.     int support_VK_KHR_create_renderpass2;

  349.     int support_VK_KHR_cooperative_matrix;

  350.     int support_VK_KHR_dedicated_allocation;

  351.     int support_VK_KHR_descriptor_update_template;

  352.     int support_VK_KHR_external_memory;

  353.     int support_VK_KHR_get_memory_requirements2;

  354.     int support_VK_KHR_maintenance1;

  355.     int support_VK_KHR_maintenance2;

  356.     int support_VK_KHR_maintenance3;

  357.     int support_VK_KHR_multiview;

  358.     int support_VK_KHR_portability_subset;

  359.     int support_VK_KHR_push_descriptor;

  360.     int support_VK_KHR_sampler_ycbcr_conversion;

  361.     int support_VK_KHR_shader_float16_int8;

  362.     int support_VK_KHR_shader_float_controls;

  363.     int support_VK_KHR_storage_buffer_storage_class;

  364.     int support_VK_KHR_swapchain;

  365.     int support_VK_EXT_buffer_device_address;

  366.     int support_VK_EXT_descriptor_indexing;

  367.     int support_VK_EXT_memory_budget;

  368.     int support_VK_EXT_memory_priority;

  369.     int support_VK_EXT_queue_family_foreign;

  370.     int support_VK_AMD_device_coherent_memory;

  371. #if __ANDROID_API__ >= 26

  372.     int support_VK_ANDROID_external_memory_android_hardware_buffer;

  373. #endif // __ANDROID_API__ >= 26

  374.     int support_VK_NV_cooperative_matrix;

  375. };

  376. 

  377. GpuInfo::GpuInfo()

  378.     : d(new GpuInfoPrivate)

  379. {

  380. }

  381. 

  382. GpuInfo::~GpuInfo()

  383. {

  384.     delete d;

  385. }

  386. 

  387. GpuInfo::GpuInfo(const GpuInfo&)

  388.     : d(0)

  389. {

  390. }

  391. 

  392. GpuInfo& GpuInfo::operator=(const GpuInfo&)

  393. {

  394.     return *this;

  395. }

  396. 

  397. VkPhysicalDevice GpuInfo::physical_device() const

  398. {

  399.     return d->physical_device;

  400. }

  401. 

  402. const VkPhysicalDeviceMemoryProperties& GpuInfo::physical_device_memory_properties() const

  403. {

  404.     return d->physical_device_memory_properties;

  405. }

  406. 

  407. uint32_t GpuInfo::api_version() const

  408. {

  409.     return d->api_version;

  410. }

  411. 

  412. uint32_t GpuInfo::driver_version() const

  413. {

  414.     return d->driver_version;

  415. }

  416. 

  417. uint32_t GpuInfo::vendor_id() const

  418. {

  419.     return d->vendor_id;

  420. }

  421. 

  422. uint32_t GpuInfo::device_id() const

  423. {

  424.     return d->device_id;

  425. }

  426. 

  427. const char* GpuInfo::device_name() const

  428. {

  429.     return d->device_name;

  430. }

  431. 

  432. uint8_t* GpuInfo::pipeline_cache_uuid() const

  433. {

  434.     return d->pipeline_cache_uuid;

  435. }

  436. 

  437. int GpuInfo::type() const

  438. {

  439.     return d->type;

  440. }

  441. 

  442. uint32_t GpuInfo::max_shared_memory_size() const

  443. {

  444.     return d->max_shared_memory_size;

  445. }

  446. 

  447. uint32_t GpuInfo::max_workgroup_count_x() const

  448. {

  449.     return d->max_workgroup_count_x;

  450. }

  451. 

  452. uint32_t GpuInfo::max_workgroup_count_y() const

  453. {

  454.     return d->max_workgroup_count_y;

  455. }

  456. 

  457. uint32_t GpuInfo::max_workgroup_count_z() const

  458. {

  459.     return d->max_workgroup_count_z;

  460. }

  461. 

  462. uint32_t GpuInfo::max_workgroup_invocations() const

  463. {

  464.     return d->max_workgroup_invocations;

  465. }

  466. 

  467. uint32_t GpuInfo::max_workgroup_size_x() const

  468. {

  469.     return d->max_workgroup_size_x;

  470. }

  471. 

  472. uint32_t GpuInfo::max_workgroup_size_y() const

  473. {

  474.     return d->max_workgroup_size_y;

  475. }

  476. 

  477. uint32_t GpuInfo::max_workgroup_size_z() const

  478. {

  479.     return d->max_workgroup_size_z;

  480. }

  481. 

  482. size_t GpuInfo::memory_map_alignment() const

  483. {

  484.     return d->memory_map_alignment;

  485. }

  486. 

  487. size_t GpuInfo::buffer_offset_alignment() const

  488. {

  489.     return d->buffer_offset_alignment;

  490. }

  491. 

  492. size_t GpuInfo::non_coherent_atom_size() const

  493. {

  494.     return d->non_coherent_atom_size;

  495. }

  496. 

  497. size_t GpuInfo::buffer_image_granularity() const

  498. {

  499.     return d->buffer_image_granularity;

  500. }

  501. 

  502. uint32_t GpuInfo::max_image_dimension_1d() const

  503. {

  504.     return d->max_image_dimension_1d;

  505. }

  506. 

  507. uint32_t GpuInfo::max_image_dimension_2d() const

  508. {

  509.     return d->max_image_dimension_2d;

  510. }

  511. 

  512. uint32_t GpuInfo::max_image_dimension_3d() const

  513. {

  514.     return d->max_image_dimension_3d;

  515. }

  516. 

  517. float GpuInfo::timestamp_period() const

  518. {

  519.     return d->timestamp_period;

  520. }

  521. 

  522. uint32_t GpuInfo::compute_queue_family_index() const

  523. {

  524.     return d->compute_queue_family_index;

  525. }

  526. 

  527. uint32_t GpuInfo::graphics_queue_family_index() const

  528. {

  529.     return d->graphics_queue_family_index;

  530. }

  531. 

  532. uint32_t GpuInfo::transfer_queue_family_index() const

  533. {

  534.     return d->transfer_queue_family_index;

  535. }

  536. 

  537. uint32_t GpuInfo::compute_queue_count() const

  538. {

  539.     return d->compute_queue_count;

  540. }

  541. 

  542. uint32_t GpuInfo::graphics_queue_count() const

  543. {

  544.     return d->graphics_queue_count;

  545. }

  546. 

  547. uint32_t GpuInfo::transfer_queue_count() const

  548. {

  549.     return d->transfer_queue_count;

  550. }

  551. 

  552. bool GpuInfo::unified_compute_transfer_queue() const

  553. {

  554.     return d->unified_compute_transfer_queue;

  555. }

  556. 

  557. uint32_t GpuInfo::subgroup_size() const

  558. {

  559.     return d->subgroup_size;

  560. }

  561. 

  562. bool GpuInfo::support_subgroup_basic() const

  563. {

  564.     return d->support_subgroup_basic;

  565. }

  566. 

  567. bool GpuInfo::support_subgroup_vote() const

  568. {

  569.     return d->support_subgroup_vote;

  570. }

  571. 

  572. bool GpuInfo::support_subgroup_ballot() const

  573. {

  574.     return d->support_subgroup_ballot;

  575. }

  576. 

  577. bool GpuInfo::support_subgroup_shuffle() const

  578. {

  579.     return d->support_subgroup_shuffle;

  580. }

  581. 

  582. bool GpuInfo::bug_storage_buffer_no_l1() const

  583. {

  584.     return d->bug_storage_buffer_no_l1;

  585. }

  586. 

  587. bool GpuInfo::bug_corrupted_online_pipeline_cache() const

  588. {

  589.     return d->bug_corrupted_online_pipeline_cache;

  590. }

  591. 

  592. bool GpuInfo::bug_buffer_image_load_zero() const

  593. {

  594.     return d->bug_buffer_image_load_zero;

  595. }

  596. 

  597. bool GpuInfo::bug_implicit_fp16_arithmetic() const

  598. {

  599.     return d->bug_implicit_fp16_arithmetic;

  600. }

  601. 

  602. bool GpuInfo::support_fp16_packed() const

  603. {

  604.     return d->support_fp16_packed;

  605. }

  606. 

  607. bool GpuInfo::support_fp16_storage() const

  608. {

  609.     return d->support_fp16_storage;

  610. }

  611. 

  612. bool GpuInfo::support_fp16_uniform() const

  613. {

  614.     return d->support_fp16_uniform;

  615. }

  616. 

  617. bool GpuInfo::support_fp16_arithmetic() const

  618. {

  619.     return d->support_fp16_arithmetic;

  620. }

  621. 

  622. bool GpuInfo::support_int8_packed() const

  623. {

  624.     return d->support_int8_packed;

  625. }

  626. 

  627. bool GpuInfo::support_int8_storage() const

  628. {

  629.     return d->support_int8_storage;

  630. }

  631. 

  632. bool GpuInfo::support_int8_uniform() const

  633. {

  634.     return d->support_int8_uniform;

  635. }

  636. 

  637. bool GpuInfo::support_int8_arithmetic() const

  638. {

  639.     return d->support_int8_arithmetic;

  640. }

  641. 

  642. bool GpuInfo::support_ycbcr_conversion() const

  643. {

  644.     return d->support_ycbcr_conversion;

  645. }

  646. 

  647. bool GpuInfo::support_cooperative_matrix() const

  648. {

  649.     return d->support_cooperative_matrix;

  650. }

  651. 

  652. bool GpuInfo::support_cooperative_matrix_8_8_16() const

  653. {

  654.     return d->support_cooperative_matrix_8_8_16;

  655. }

  656. 

  657. bool GpuInfo::support_cooperative_matrix_16_8_8() const

  658. {

  659.     return d->support_cooperative_matrix_16_8_8;

  660. }

  661. 

  662. bool GpuInfo::support_cooperative_matrix_16_8_16() const

  663. {

  664.     return d->support_cooperative_matrix_16_8_16;

  665. }

  666. 

  667. bool GpuInfo::support_cooperative_matrix_16_16_16() const

  668. {

  669.     return d->support_cooperative_matrix_16_16_16;

  670. }

  671. 

  672. int GpuInfo::support_VK_KHR_8bit_storage() const

  673. {

  674.     return d->support_VK_KHR_8bit_storage;

  675. }

  676. 

  677. int GpuInfo::support_VK_KHR_16bit_storage() const

  678. {

  679.     return d->support_VK_KHR_16bit_storage;

  680. }

  681. 

  682. int GpuInfo::support_VK_KHR_bind_memory2() const

  683. {

  684.     return d->support_VK_KHR_bind_memory2;

  685. }

  686. 

  687. int GpuInfo::support_VK_KHR_buffer_device_address() const

  688. {

  689.     return d->support_VK_KHR_buffer_device_address;

  690. }

  691. 

  692. int GpuInfo::support_VK_KHR_create_renderpass2() const

  693. {

  694.     return d->support_VK_KHR_create_renderpass2;

  695. }

  696. 

  697. int GpuInfo::support_VK_KHR_cooperative_matrix() const

  698. {

  699.     return d->support_VK_KHR_cooperative_matrix;

  700. }

  701. 

  702. int GpuInfo::support_VK_KHR_dedicated_allocation() const

  703. {

  704.     return d->support_VK_KHR_dedicated_allocation;

  705. }

  706. 

  707. int GpuInfo::support_VK_KHR_descriptor_update_template() const

  708. {

  709.     return d->support_VK_KHR_descriptor_update_template;

  710. }

  711. 

  712. int GpuInfo::support_VK_KHR_external_memory() const

  713. {

  714.     return d->support_VK_KHR_external_memory;

  715. }

  716. 

  717. int GpuInfo::support_VK_KHR_get_memory_requirements2() const

  718. {

  719.     return d->support_VK_KHR_get_memory_requirements2;

  720. }

  721. 

  722. int GpuInfo::support_VK_KHR_maintenance1() const

  723. {

  724.     return d->support_VK_KHR_maintenance1;

  725. }

  726. 

  727. int GpuInfo::support_VK_KHR_maintenance2() const

  728. {

  729.     return d->support_VK_KHR_maintenance2;

  730. }

  731. 

  732. int GpuInfo::support_VK_KHR_maintenance3() const

  733. {

  734.     return d->support_VK_KHR_maintenance3;

  735. }

  736. 

  737. int GpuInfo::support_VK_KHR_multiview() const

  738. {

  739.     return d->support_VK_KHR_multiview;

  740. }

  741. 

  742. int GpuInfo::support_VK_KHR_portability_subset() const

  743. {

  744.     return d->support_VK_KHR_portability_subset;

  745. }

  746. 

  747. int GpuInfo::support_VK_KHR_push_descriptor() const

  748. {

  749.     return d->support_VK_KHR_push_descriptor;

  750. }

  751. 

  752. int GpuInfo::support_VK_KHR_sampler_ycbcr_conversion() const

  753. {

  754.     return d->support_VK_KHR_sampler_ycbcr_conversion;

  755. }

  756. 

  757. int GpuInfo::support_VK_KHR_shader_float16_int8() const

  758. {

  759.     return d->support_VK_KHR_shader_float16_int8;

  760. }

  761. 

  762. int GpuInfo::support_VK_KHR_shader_float_controls() const

  763. {

  764.     return d->support_VK_KHR_shader_float_controls;

  765. }

  766. 

  767. int GpuInfo::support_VK_KHR_storage_buffer_storage_class() const

  768. {

  769.     return d->support_VK_KHR_storage_buffer_storage_class;

  770. }

  771. 

  772. int GpuInfo::support_VK_KHR_swapchain() const

  773. {

  774.     return d->support_VK_KHR_swapchain;

  775. }

  776. 

  777. int GpuInfo::support_VK_EXT_buffer_device_address() const

  778. {

  779.     return d->support_VK_EXT_buffer_device_address;

  780. }

  781. 

  782. int GpuInfo::support_VK_EXT_descriptor_indexing() const

  783. {

  784.     return d->support_VK_EXT_descriptor_indexing;

  785. }

  786. 

  787. int GpuInfo::support_VK_EXT_memory_budget() const

  788. {

  789.     return d->support_VK_EXT_memory_budget;

  790. }

  791. 

  792. int GpuInfo::support_VK_EXT_memory_priority() const

  793. {

  794.     return d->support_VK_EXT_memory_priority;

  795. }

  796. 

  797. int GpuInfo::support_VK_EXT_queue_family_foreign() const

  798. {

  799.     return d->support_VK_EXT_queue_family_foreign;

  800. }

  801. 

  802. int GpuInfo::support_VK_AMD_device_coherent_memory() const

  803. {

  804.     return d->support_VK_AMD_device_coherent_memory;

  805. }

  806. 

  807. #if __ANDROID_API__ >= 26

  808. int GpuInfo::support_VK_ANDROID_external_memory_android_hardware_buffer() const

  809. {

  810.     return d->support_VK_ANDROID_external_memory_android_hardware_buffer;

  811. }

  812. #endif // __ANDROID_API__ >= 26

  813. 

  814. int GpuInfo::support_VK_NV_cooperative_matrix() const

  815. {

  816.     return d->support_VK_NV_cooperative_matrix;

  817. }

  818. 

  819. static int init_instance_core()

  820. {

  821.     vkAllocateCommandBuffers = (PFN_vkAllocateCommandBuffers)vkGetInstanceProcAddr(g_instance, "vkAllocateCommandBuffers");

  822.     vkAllocateDescriptorSets = (PFN_vkAllocateDescriptorSets)vkGetInstanceProcAddr(g_instance, "vkAllocateDescriptorSets");

  823.     vkAllocateMemory = (PFN_vkAllocateMemory)vkGetInstanceProcAddr(g_instance, "vkAllocateMemory");

  824.     vkBeginCommandBuffer = (PFN_vkBeginCommandBuffer)vkGetInstanceProcAddr(g_instance, "vkBeginCommandBuffer");

  825.     vkBindBufferMemory = (PFN_vkBindBufferMemory)vkGetInstanceProcAddr(g_instance, "vkBindBufferMemory");

  826.     vkBindImageMemory = (PFN_vkBindImageMemory)vkGetInstanceProcAddr(g_instance, "vkBindImageMemory");

  827.     vkCmdBeginQuery = (PFN_vkCmdBeginQuery)vkGetInstanceProcAddr(g_instance, "vkCmdBeginQuery");

  828.     vkCmdBindDescriptorSets = (PFN_vkCmdBindDescriptorSets)vkGetInstanceProcAddr(g_instance, "vkCmdBindDescriptorSets");

  829.     vkCmdBindIndexBuffer = (PFN_vkCmdBindIndexBuffer)vkGetInstanceProcAddr(g_instance, "vkCmdBindIndexBuffer");

  830.     vkCmdBindPipeline = (PFN_vkCmdBindPipeline)vkGetInstanceProcAddr(g_instance, "vkCmdBindPipeline");

  831.     vkCmdCopyBuffer = (PFN_vkCmdCopyBuffer)vkGetInstanceProcAddr(g_instance, "vkCmdCopyBuffer");

  832.     vkCmdCopyBufferToImage = (PFN_vkCmdCopyBufferToImage)vkGetInstanceProcAddr(g_instance, "vkCmdCopyBufferToImage");

  833.     vkCmdCopyImage = (PFN_vkCmdCopyImage)vkGetInstanceProcAddr(g_instance, "vkCmdCopyImage");

  834.     vkCmdCopyImageToBuffer = (PFN_vkCmdCopyImageToBuffer)vkGetInstanceProcAddr(g_instance, "vkCmdCopyImageToBuffer");

  835.     vkCmdCopyQueryPoolResults = (PFN_vkCmdCopyQueryPoolResults)vkGetInstanceProcAddr(g_instance, "vkCmdCopyQueryPoolResults");

  836.     vkCmdDispatch = (PFN_vkCmdDispatch)vkGetInstanceProcAddr(g_instance, "vkCmdDispatch");

  837.     vkCmdDispatchIndirect = (PFN_vkCmdDispatchIndirect)vkGetInstanceProcAddr(g_instance, "vkCmdDispatchIndirect");

  838.     vkCmdEndQuery = (PFN_vkCmdEndQuery)vkGetInstanceProcAddr(g_instance, "vkCmdEndQuery");

  839.     vkCmdExecuteCommands = (PFN_vkCmdExecuteCommands)vkGetInstanceProcAddr(g_instance, "vkCmdExecuteCommands");

  840.     vkCmdFillBuffer = (PFN_vkCmdFillBuffer)vkGetInstanceProcAddr(g_instance, "vkCmdFillBuffer");

  841.     vkCmdPipelineBarrier = (PFN_vkCmdPipelineBarrier)vkGetInstanceProcAddr(g_instance, "vkCmdPipelineBarrier");

  842.     vkCmdPushConstants = (PFN_vkCmdPushConstants)vkGetInstanceProcAddr(g_instance, "vkCmdPushConstants");

  843.     vkCmdResetQueryPool = (PFN_vkCmdResetQueryPool)vkGetInstanceProcAddr(g_instance, "vkCmdResetQueryPool");

  844.     vkCmdResolveImage = (PFN_vkCmdResolveImage)vkGetInstanceProcAddr(g_instance, "vkCmdResolveImage");

  845.     vkCmdUpdateBuffer = (PFN_vkCmdUpdateBuffer)vkGetInstanceProcAddr(g_instance, "vkCmdUpdateBuffer");

  846.     vkCmdWriteTimestamp = (PFN_vkCmdWriteTimestamp)vkGetInstanceProcAddr(g_instance, "vkCmdWriteTimestamp");

  847.     vkCreateBuffer = (PFN_vkCreateBuffer)vkGetInstanceProcAddr(g_instance, "vkCreateBuffer");

  848.     vkCreateBufferView = (PFN_vkCreateBufferView)vkGetInstanceProcAddr(g_instance, "vkCreateBufferView");

  849.     vkCreateCommandPool = (PFN_vkCreateCommandPool)vkGetInstanceProcAddr(g_instance, "vkCreateCommandPool");

  850.     vkCreateComputePipelines = (PFN_vkCreateComputePipelines)vkGetInstanceProcAddr(g_instance, "vkCreateComputePipelines");

  851.     vkCreateDescriptorPool = (PFN_vkCreateDescriptorPool)vkGetInstanceProcAddr(g_instance, "vkCreateDescriptorPool");

  852.     vkCreateDescriptorSetLayout = (PFN_vkCreateDescriptorSetLayout)vkGetInstanceProcAddr(g_instance, "vkCreateDescriptorSetLayout");

  853.     vkCreateDevice = (PFN_vkCreateDevice)vkGetInstanceProcAddr(g_instance, "vkCreateDevice");

  854.     vkCreateFence = (PFN_vkCreateFence)vkGetInstanceProcAddr(g_instance, "vkCreateFence");

  855.     vkCreateImage = (PFN_vkCreateImage)vkGetInstanceProcAddr(g_instance, "vkCreateImage");

  856.     vkCreateImageView = (PFN_vkCreateImageView)vkGetInstanceProcAddr(g_instance, "vkCreateImageView");

  857.     vkCreatePipelineCache = (PFN_vkCreatePipelineCache)vkGetInstanceProcAddr(g_instance, "vkCreatePipelineCache");

  858.     vkCreatePipelineLayout = (PFN_vkCreatePipelineLayout)vkGetInstanceProcAddr(g_instance, "vkCreatePipelineLayout");

  859.     vkCreateQueryPool = (PFN_vkCreateQueryPool)vkGetInstanceProcAddr(g_instance, "vkCreateQueryPool");

  860.     vkCreateSampler = (PFN_vkCreateSampler)vkGetInstanceProcAddr(g_instance, "vkCreateSampler");

  861.     vkCreateSemaphore = (PFN_vkCreateSemaphore)vkGetInstanceProcAddr(g_instance, "vkCreateSemaphore");

  862.     vkCreateShaderModule = (PFN_vkCreateShaderModule)vkGetInstanceProcAddr(g_instance, "vkCreateShaderModule");

  863.     vkDestroyBuffer = (PFN_vkDestroyBuffer)vkGetInstanceProcAddr(g_instance, "vkDestroyBuffer");

  864.     vkDestroyBufferView = (PFN_vkDestroyBufferView)vkGetInstanceProcAddr(g_instance, "vkDestroyBufferView");

  865.     vkDestroyCommandPool = (PFN_vkDestroyCommandPool)vkGetInstanceProcAddr(g_instance, "vkDestroyCommandPool");

  866.     vkDestroyDescriptorPool = (PFN_vkDestroyDescriptorPool)vkGetInstanceProcAddr(g_instance, "vkDestroyDescriptorPool");

  867.     vkDestroyDescriptorSetLayout = (PFN_vkDestroyDescriptorSetLayout)vkGetInstanceProcAddr(g_instance, "vkDestroyDescriptorSetLayout");

  868.     vkDestroyDevice = (PFN_vkDestroyDevice)vkGetInstanceProcAddr(g_instance, "vkDestroyDevice");

  869.     vkDestroyFence = (PFN_vkDestroyFence)vkGetInstanceProcAddr(g_instance, "vkDestroyFence");

  870.     vkDestroyImage = (PFN_vkDestroyImage)vkGetInstanceProcAddr(g_instance, "vkDestroyImage");

  871.     vkDestroyImageView = (PFN_vkDestroyImageView)vkGetInstanceProcAddr(g_instance, "vkDestroyImageView");

  872.     vkDestroyInstance = (PFN_vkDestroyInstance)vkGetInstanceProcAddr(g_instance, "vkDestroyInstance");

  873.     vkDestroyPipeline = (PFN_vkDestroyPipeline)vkGetInstanceProcAddr(g_instance, "vkDestroyPipeline");

  874.     vkDestroyPipelineCache = (PFN_vkDestroyPipelineCache)vkGetInstanceProcAddr(g_instance, "vkDestroyPipelineCache");

  875.     vkDestroyPipelineLayout = (PFN_vkDestroyPipelineLayout)vkGetInstanceProcAddr(g_instance, "vkDestroyPipelineLayout");

  876.     vkDestroyQueryPool = (PFN_vkDestroyQueryPool)vkGetInstanceProcAddr(g_instance, "vkDestroyQueryPool");

  877.     vkDestroySampler = (PFN_vkDestroySampler)vkGetInstanceProcAddr(g_instance, "vkDestroySampler");

  878.     vkDestroySemaphore = (PFN_vkDestroySemaphore)vkGetInstanceProcAddr(g_instance, "vkDestroySemaphore");

  879.     vkDestroyShaderModule = (PFN_vkDestroyShaderModule)vkGetInstanceProcAddr(g_instance, "vkDestroyShaderModule");

  880.     vkDeviceWaitIdle = (PFN_vkDeviceWaitIdle)vkGetInstanceProcAddr(g_instance, "vkDeviceWaitIdle");

  881.     vkEndCommandBuffer = (PFN_vkEndCommandBuffer)vkGetInstanceProcAddr(g_instance, "vkEndCommandBuffer");

  882.     vkEnumerateDeviceExtensionProperties = (PFN_vkEnumerateDeviceExtensionProperties)vkGetInstanceProcAddr(g_instance, "vkEnumerateDeviceExtensionProperties");

  883.     vkEnumerateDeviceLayerProperties = (PFN_vkEnumerateDeviceLayerProperties)vkGetInstanceProcAddr(g_instance, "vkEnumerateDeviceLayerProperties");

  884.     vkEnumeratePhysicalDevices = (PFN_vkEnumeratePhysicalDevices)vkGetInstanceProcAddr(g_instance, "vkEnumeratePhysicalDevices");

  885.     vkFlushMappedMemoryRanges = (PFN_vkFlushMappedMemoryRanges)vkGetInstanceProcAddr(g_instance, "vkFlushMappedMemoryRanges");

  886.     vkFreeCommandBuffers = (PFN_vkFreeCommandBuffers)vkGetInstanceProcAddr(g_instance, "vkFreeCommandBuffers");

  887.     vkFreeDescriptorSets = (PFN_vkFreeDescriptorSets)vkGetInstanceProcAddr(g_instance, "vkFreeDescriptorSets");

  888.     vkFreeMemory = (PFN_vkFreeMemory)vkGetInstanceProcAddr(g_instance, "vkFreeMemory");

  889.     vkGetBufferMemoryRequirements = (PFN_vkGetBufferMemoryRequirements)vkGetInstanceProcAddr(g_instance, "vkGetBufferMemoryRequirements");

  890.     vkGetDeviceMemoryCommitment = (PFN_vkGetDeviceMemoryCommitment)vkGetInstanceProcAddr(g_instance, "vkGetDeviceMemoryCommitment");

  891.     vkGetDeviceProcAddr = (PFN_vkGetDeviceProcAddr)vkGetInstanceProcAddr(g_instance, "vkGetDeviceProcAddr");

  892.     vkGetDeviceQueue = (PFN_vkGetDeviceQueue)vkGetInstanceProcAddr(g_instance, "vkGetDeviceQueue");

  893.     vkGetFenceStatus = (PFN_vkGetFenceStatus)vkGetInstanceProcAddr(g_instance, "vkGetFenceStatus");

  894.     vkGetImageMemoryRequirements = (PFN_vkGetImageMemoryRequirements)vkGetInstanceProcAddr(g_instance, "vkGetImageMemoryRequirements");

  895.     vkGetImageSubresourceLayout = (PFN_vkGetImageSubresourceLayout)vkGetInstanceProcAddr(g_instance, "vkGetImageSubresourceLayout");

  896.     vkGetPhysicalDeviceFeatures = (PFN_vkGetPhysicalDeviceFeatures)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceFeatures");

  897.     vkGetPhysicalDeviceFormatProperties = (PFN_vkGetPhysicalDeviceFormatProperties)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceFormatProperties");

  898.     vkGetPhysicalDeviceImageFormatProperties = (PFN_vkGetPhysicalDeviceImageFormatProperties)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceImageFormatProperties");

  899.     vkGetPhysicalDeviceMemoryProperties = (PFN_vkGetPhysicalDeviceMemoryProperties)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceMemoryProperties");

  900.     vkGetPhysicalDeviceProperties = (PFN_vkGetPhysicalDeviceProperties)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceProperties");

  901.     vkGetPhysicalDeviceQueueFamilyProperties = (PFN_vkGetPhysicalDeviceQueueFamilyProperties)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceQueueFamilyProperties");

  902.     vkGetPipelineCacheData = (PFN_vkGetPipelineCacheData)vkGetInstanceProcAddr(g_instance, "vkGetPipelineCacheData");

  903.     vkGetQueryPoolResults = (PFN_vkGetQueryPoolResults)vkGetInstanceProcAddr(g_instance, "vkGetQueryPoolResults");

  904.     vkInvalidateMappedMemoryRanges = (PFN_vkInvalidateMappedMemoryRanges)vkGetInstanceProcAddr(g_instance, "vkInvalidateMappedMemoryRanges");

  905.     vkMapMemory = (PFN_vkMapMemory)vkGetInstanceProcAddr(g_instance, "vkMapMemory");

  906.     vkMergePipelineCaches = (PFN_vkMergePipelineCaches)vkGetInstanceProcAddr(g_instance, "vkMergePipelineCaches");

  907.     vkQueueSubmit = (PFN_vkQueueSubmit)vkGetInstanceProcAddr(g_instance, "vkQueueSubmit");

  908.     vkQueueWaitIdle = (PFN_vkQueueWaitIdle)vkGetInstanceProcAddr(g_instance, "vkQueueWaitIdle");

  909.     vkResetCommandBuffer = (PFN_vkResetCommandBuffer)vkGetInstanceProcAddr(g_instance, "vkResetCommandBuffer");

  910.     vkResetCommandPool = (PFN_vkResetCommandPool)vkGetInstanceProcAddr(g_instance, "vkResetCommandPool");

  911.     vkResetDescriptorPool = (PFN_vkResetDescriptorPool)vkGetInstanceProcAddr(g_instance, "vkResetDescriptorPool");

  912.     vkResetFences = (PFN_vkResetFences)vkGetInstanceProcAddr(g_instance, "vkResetFences");

  913.     vkUnmapMemory = (PFN_vkUnmapMemory)vkGetInstanceProcAddr(g_instance, "vkUnmapMemory");

  914.     vkUpdateDescriptorSets = (PFN_vkUpdateDescriptorSets)vkGetInstanceProcAddr(g_instance, "vkUpdateDescriptorSets");

  915.     vkWaitForFences = (PFN_vkWaitForFences)vkGetInstanceProcAddr(g_instance, "vkWaitForFences");

  916. 

  917.     return 0;

  918. }

  919. 

  920. static int init_instance_extension()

  921. {

  922.     if (support_VK_KHR_external_memory_capabilities)

  923.     {

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

  925.     }

  926. 

  927.     if (support_VK_KHR_get_physical_device_properties2)

  928.     {

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

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

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

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

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

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

  935.     }

  936. 

  937.     if (support_VK_KHR_get_surface_capabilities2)

  938.     {

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

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

  941.     }

  942. 

  943.     if (support_VK_KHR_surface)

  944.     {

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

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

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

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

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

  950.     }

  951. 

  952. #if __ANDROID_API__ >= 26

  953.     if (support_VK_KHR_android_surface)

  954.     {

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

  956.     }

  957. #endif // __ANDROID_API__ >= 26

  958. 

  959.     // VK_KHR_cooperative_matrix

  960.     {

  961.         vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR = (PFN_vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR");

  962.     }

  963. 

  964.     // VK_NV_cooperative_matrix

  965.     {

  966.         vkGetPhysicalDeviceCooperativeMatrixPropertiesNV = (PFN_vkGetPhysicalDeviceCooperativeMatrixPropertiesNV)vkGetInstanceProcAddr(g_instance, "vkGetPhysicalDeviceCooperativeMatrixPropertiesNV");

  967.     }

  968. 

  969.     return 0;

  970. }

  971. 

  972. #if ENABLE_VALIDATION_LAYER

  973. static VKAPI_ATTR VkBool32 VKAPI_CALL debugCallback(

  974.     VkDebugUtilsMessageSeverityFlagBitsEXT /*messageSeverity*/,

  975.     VkDebugUtilsMessageTypeFlagsEXT /*messageType*/,

  976.     const VkDebugUtilsMessengerCallbackDataEXT* pCallbackData,

  977.     void* /*pUserData*/)

  978. {

  979.     NCNN_LOGE("validation layer: %s", pCallbackData->pMessage);

  980. 

  981.     return VK_FALSE;

  982. }

  983. 

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

  985. {

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

  987.     if (func)

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

  989. 

  990.     return VK_ERROR_EXTENSION_NOT_PRESENT;

  991. }

  992. 

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

  994. {

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

  996.     if (func)

  997.         func(instance, callback, pAllocator);

  998. }

  999. #endif // ENABLE_VALIDATION_LAYER

  1000. 

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

  1002. {

  1003.     // first try, compute only queue

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

  1005.     {

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

  1007. 

  1008.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

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

  1010.         {

  1011.             return i;

  1012.         }

  1013.     }

  1014. 

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

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

  1017.     {

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

  1019. 

  1020.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  1021.                 && (queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT))

  1022.         {

  1023.             return i;

  1024.         }

  1025.     }

  1026. 

  1027.     // third try, any queue with compute

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

  1029.     {

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

  1031. 

  1032.         if (queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT)

  1033.         {

  1034.             return i;

  1035.         }

  1036.     }

  1037. 

  1038.     //     NCNN_LOGE("no compute queue");

  1039.     return -1;

  1040. }

  1041. 

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

  1043. {

  1044.     // first try, graphics only queue

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

  1046.     {

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

  1048. 

  1049.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

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

  1051.         {

  1052.             return i;

  1053.         }

  1054.     }

  1055. 

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

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

  1058.     {

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

  1060. 

  1061.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

  1062.                 && (queueFamilyProperty.queueFlags & VK_QUEUE_COMPUTE_BIT))

  1063.         {

  1064.             return i;

  1065.         }

  1066.     }

  1067. 

  1068.     // third try, any queue with graphics

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

  1070.     {

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

  1072. 

  1073.         if (queueFamilyProperty.queueFlags & VK_QUEUE_GRAPHICS_BIT)

  1074.         {

  1075.             return i;

  1076.         }

  1077.     }

  1078. 

  1079.     //     NCNN_LOGE("no graphics queue");

  1080.     return -1;

  1081. }

  1082. 

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

  1084. {

  1085.     // first try, transfer only queue

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

  1087.     {

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

  1089. 

  1090.         if ((queueFamilyProperty.queueFlags & VK_QUEUE_TRANSFER_BIT)

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

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

  1093.         {

  1094.             return i;

  1095.         }

  1096.     }

  1097. 

  1098.     // second try, any queue with transfer

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

  1100.     {

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

  1102. 

  1103.         if (queueFamilyProperty.queueFlags & VK_QUEUE_TRANSFER_BIT)

  1104.         {

  1105.             return i;

  1106.         }

  1107.     }

  1108. 

  1109.     // third try, use compute queue

  1110.     uint32_t compute_queue_index = find_device_compute_queue(queueFamilyProperties);

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

  1112.     {

  1113.         return compute_queue_index;

  1114.     }

  1115. 

  1116.     // fourth try, use graphics queue

  1117.     uint32_t graphics_queue_index = find_device_graphics_queue(queueFamilyProperties);

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

  1119.     {

  1120.         return graphics_queue_index;

  1121.     }

  1122. 

  1123.     //     NCNN_LOGE("no transfer queue");

  1124.     return -1;

  1125. }

  1126. 

  1127. static int find_default_vulkan_device_index()

  1128. {

  1129.     // first try, discrete gpu

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

  1131.     {

  1132.         if (g_gpu_infos[i]->type() == 0)

  1133.             return i;

  1134.     }

  1135. 

  1136.     // second try, integrated gpu

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

  1138.     {

  1139.         if (g_gpu_infos[i]->type() == 1)

  1140.             return i;

  1141.     }

  1142. 

  1143.     // third try, any probed device

  1144.     if (g_gpu_count > 0)

  1145.         return 0;

  1146. 

  1147.     NCNN_LOGE("no vulkan device");

  1148.     return -1;

  1149. }

  1150. 

  1151. int create_gpu_instance(const char* driver_path)

  1152. {

  1153.     MutexLockGuard lock(g_instance_lock);

  1154. 

  1155.     if (g_instance.created != 0)

  1156.         return g_instance.instance ? 0 : -1;

  1157. 

  1158.     g_instance.created = 1;

  1159. 

  1160.     // NCNN_LOGE("create_gpu_instance");

  1161. 

  1162. #if NCNN_SIMPLEVK

  1163.     // load vulkan driver

  1164.     {

  1165.         int ret = load_vulkan_driver(driver_path);

  1166.         if (ret != 0)

  1167.         {

  1168.             NCNN_LOGE("load vulkan driver failed");

  1169.             return -1;

  1170.         }

  1171.     }

  1172. #else

  1173.     if (driver_path)

  1174.     {

  1175.         NCNN_LOGE("custom vulkan driver is not supported when NCNN_SIMPLEVK is off");

  1176.         NCNN_LOGE("will always use the system vulkan driver");

  1177.     }

  1178. #endif // NCNN_SIMPLEVK

  1179. 

  1180.     VkResult ret;

  1181. 

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

  1183. 

  1184. #if ENABLE_VALIDATION_LAYER

  1185.     uint32_t instanceLayerPropertyCount;

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

  1187.     if (ret != VK_SUCCESS)

  1188.     {

  1189.         NCNN_LOGE("vkEnumerateInstanceLayerProperties failed %d", ret);

  1190.         return -1;

  1191.     }

  1192. 

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

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

  1195.     if (ret != VK_SUCCESS)

  1196.     {

  1197.         NCNN_LOGE("vkEnumerateInstanceLayerProperties failed %d", ret);

  1198.         return -1;

  1199.     }

  1200. 

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

  1202.     {

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

  1204.         //         NCNN_LOGE("instance layer %s = %u", lp.layerName, lp.implementationVersion);

  1205. 

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

  1207.         {

  1208.             enabledLayers.push_back("VK_LAYER_LUNARG_standard_validation");

  1209.         }

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

  1211.         {

  1212.             enabledLayers.push_back("VK_LAYER_LUNARG_parameter_validation");

  1213.         }

  1214.         if (strcmp(lp.layerName, "VK_LAYER_KHRONOS_validation") == 0)

  1215.         {

  1216.             enabledLayers.push_back("VK_LAYER_KHRONOS_validation");

  1217.         }

  1218.     }

  1219. #endif // ENABLE_VALIDATION_LAYER

  1220. 

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

  1222. 

  1223.     uint32_t instanceExtensionPropertyCount;

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

  1225.     if (ret != VK_SUCCESS)

  1226.     {

  1227.         NCNN_LOGE("vkEnumerateInstanceExtensionProperties failed %d", ret);

  1228.         return -1;

  1229.     }

  1230. 

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

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

  1233.     if (ret != VK_SUCCESS)

  1234.     {

  1235.         NCNN_LOGE("vkEnumerateInstanceExtensionProperties failed %d", ret);

  1236.         return -1;

  1237.     }

  1238. 

  1239.     support_VK_KHR_get_physical_device_properties2 = 0;

  1240.     support_VK_KHR_get_surface_capabilities2 = 0;

  1241.     support_VK_KHR_portability_enumeration = 0;

  1242.     support_VK_KHR_surface = 0;

  1243.     support_VK_EXT_debug_utils = 0;

  1244.     support_VK_EXT_validation_features = 0;

  1245.     support_VK_EXT_validation_flags = 0;

  1246. #if __ANDROID_API__ >= 26

  1247.     support_VK_KHR_android_surface = 0;

  1248. #endif // __ANDROID_API__ >= 26

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

  1250.     {

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

  1252.         //         NCNN_LOGE("instance extension %s = %u", exp.extensionName, exp.specVersion);

  1253. 

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

  1255.             support_VK_KHR_external_memory_capabilities = exp.specVersion;

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

  1257.             support_VK_KHR_get_physical_device_properties2 = exp.specVersion;

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

  1259.             support_VK_KHR_get_surface_capabilities2 = exp.specVersion;

  1260.         else if (strcmp(exp.extensionName, "VK_KHR_portability_enumeration") == 0)

  1261.             support_VK_KHR_portability_enumeration = exp.specVersion;

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

  1263.             support_VK_KHR_surface = exp.specVersion;

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

  1265.             support_VK_EXT_debug_utils = exp.specVersion;

  1266.         else if (strcmp(exp.extensionName, "VK_EXT_validation_features") == 0)

  1267.             support_VK_EXT_validation_features = exp.specVersion;

  1268.         else if (strcmp(exp.extensionName, "VK_EXT_validation_flags") == 0)

  1269.             support_VK_EXT_validation_flags = exp.specVersion;

  1270. #if __ANDROID_API__ >= 26

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

  1272.             support_VK_KHR_android_surface = exp.specVersion;

  1273. #endif // __ANDROID_API__ >= 26

  1274.     }

  1275. 

  1276.     if (support_VK_EXT_validation_features)

  1277.     {

  1278.         // we prefer the modern one

  1279.         support_VK_EXT_validation_flags = 0;

  1280.     }

  1281. 

  1282.     if (support_VK_KHR_external_memory_capabilities)

  1283.         enabledExtensions.push_back("VK_KHR_external_memory_capabilities");

  1284.     if (support_VK_KHR_get_physical_device_properties2)

  1285.         enabledExtensions.push_back("VK_KHR_get_physical_device_properties2");

  1286.     if (support_VK_KHR_get_surface_capabilities2)

  1287.         enabledExtensions.push_back("VK_KHR_get_surface_capabilities2");

  1288.     if (support_VK_KHR_portability_enumeration)

  1289.         enabledExtensions.push_back("VK_KHR_portability_enumeration");

  1290.     if (support_VK_KHR_surface)

  1291.         enabledExtensions.push_back("VK_KHR_surface");

  1292. #if ENABLE_VALIDATION_LAYER

  1293.     if (support_VK_EXT_debug_utils)

  1294.         enabledExtensions.push_back("VK_EXT_debug_utils");

  1295.     if (support_VK_EXT_validation_features)

  1296.         enabledExtensions.push_back("VK_EXT_validation_features");

  1297.     if (support_VK_EXT_validation_flags)

  1298.         enabledExtensions.push_back("VK_EXT_validation_flags");

  1299. #endif // ENABLE_VALIDATION_LAYER

  1300. #if __ANDROID_API__ >= 26

  1301.     if (support_VK_KHR_android_surface)

  1302.         enabledExtensions.push_back("VK_KHR_android_surface");

  1303. #endif // __ANDROID_API__ >= 26

  1304. 

  1305.     uint32_t instance_api_version = VK_MAKE_VERSION(1, 0, 0);

  1306.     typedef VkResult(VKAPI_PTR * PFN_vkEnumerateInstanceVersion)(uint32_t * pApiVersion);

  1307.     PFN_vkEnumerateInstanceVersion vkEnumerateInstanceVersion = (PFN_vkEnumerateInstanceVersion)vkGetInstanceProcAddr(0, "vkEnumerateInstanceVersion");

  1308.     if (vkEnumerateInstanceVersion)

  1309.     {

  1310.         ret = vkEnumerateInstanceVersion(&instance_api_version);

  1311.         if (ret != VK_SUCCESS)

  1312.         {

  1313.             NCNN_LOGE("vkEnumerateInstanceVersion failed %d", ret);

  1314.             return -1;

  1315.         }

  1316.     }

  1317. 

  1318.     // NCNN_LOGE("instance apiVersion = %u.%u.%u", VK_VERSION_MAJOR(instance_api_version), VK_VERSION_MINOR(instance_api_version), VK_VERSION_PATCH(instance_api_version));

  1319. 

  1320.     VkApplicationInfo applicationInfo;

  1321.     applicationInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;

  1322.     applicationInfo.pNext = 0;

  1323.     applicationInfo.pApplicationName = "ncnn";

  1324.     applicationInfo.applicationVersion = 0;

  1325.     applicationInfo.pEngineName = "ncnn";

  1326.     applicationInfo.engineVersion = 20201010;

  1327.     applicationInfo.apiVersion = instance_api_version;

  1328. 

  1329.     void* enabledExtensionFeatures = 0;

  1330. 

  1331. #if ENABLE_VALIDATION_LAYER

  1332.     std::vector<VkValidationFeatureEnableEXT> enabledValidationFeature;

  1333.     enabledValidationFeature.push_back(VK_VALIDATION_FEATURE_ENABLE_GPU_ASSISTED_EXT);

  1334.     enabledValidationFeature.push_back(VK_VALIDATION_FEATURE_ENABLE_GPU_ASSISTED_RESERVE_BINDING_SLOT_EXT);

  1335.     enabledValidationFeature.push_back(VK_VALIDATION_FEATURE_ENABLE_BEST_PRACTICES_EXT);

  1336.     enabledValidationFeature.push_back(VK_VALIDATION_FEATURE_ENABLE_DEBUG_PRINTF_EXT);

  1337.     enabledValidationFeature.push_back(VK_VALIDATION_FEATURE_ENABLE_SYNCHRONIZATION_VALIDATION_EXT);

  1338. 

  1339.     VkValidationFeaturesEXT validationFeatures;

  1340.     validationFeatures.sType = VK_STRUCTURE_TYPE_VALIDATION_FEATURES_EXT;

  1341.     validationFeatures.pNext = 0;

  1342.     validationFeatures.enabledValidationFeatureCount = enabledValidationFeature.size();

  1343.     validationFeatures.pEnabledValidationFeatures = enabledValidationFeature.data();

  1344.     validationFeatures.disabledValidationFeatureCount = 0;

  1345.     validationFeatures.pDisabledValidationFeatures = 0;

  1346.     if (support_VK_EXT_validation_features)

  1347.     {

  1348.         validationFeatures.pNext = enabledExtensionFeatures;

  1349.         enabledExtensionFeatures = &validationFeatures;

  1350.     }

  1351. 

  1352.     VkValidationFlagsEXT validationFlags;

  1353.     validationFlags.sType = VK_STRUCTURE_TYPE_VALIDATION_FLAGS_EXT;

  1354.     validationFlags.pNext = 0;

  1355.     validationFlags.disabledValidationCheckCount = 0;

  1356.     validationFlags.pDisabledValidationChecks = 0;

  1357.     if (support_VK_EXT_validation_flags)

  1358.     {

  1359.         validationFlags.pNext = enabledExtensionFeatures;

  1360.         enabledExtensionFeatures = &validationFlags;

  1361.     }

  1362. #endif // ENABLE_VALIDATION_LAYER

  1363. 

  1364.     VkInstanceCreateInfo instanceCreateInfo;

  1365.     instanceCreateInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;

  1366.     instanceCreateInfo.pNext = enabledExtensionFeatures;

  1367.     instanceCreateInfo.flags = 0;

  1368.     if (support_VK_KHR_portability_enumeration)

  1369.         instanceCreateInfo.flags |= VK_INSTANCE_CREATE_ENUMERATE_PORTABILITY_BIT_KHR;

  1370.     instanceCreateInfo.pApplicationInfo = &applicationInfo;

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

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

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

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

  1375. 

  1376.     VkInstance instance = 0;

  1377.     ret = vkCreateInstance(&instanceCreateInfo, 0, &instance);

  1378.     if (ret != VK_SUCCESS)

  1379.     {

  1380.         NCNN_LOGE("vkCreateInstance failed %d", ret);

  1381.         return -1;

  1382.     }

  1383. 

  1384.     g_instance.instance = instance;

  1385. 

  1386.     init_instance_core();

  1387. 

  1388. #if ENABLE_VALIDATION_LAYER

  1389.     if (support_VK_EXT_debug_utils)

  1390.     {

  1391.         VkDebugUtilsMessengerCreateInfoEXT createInfo = {};

  1392.         createInfo.sType = VK_STRUCTURE_TYPE_DEBUG_UTILS_MESSENGER_CREATE_INFO_EXT;

  1393.         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;

  1394.         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;

  1395.         createInfo.pfnUserCallback = debugCallback;

  1396.         createInfo.pUserData = 0;

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

  1398.         if (ret != VK_SUCCESS)

  1399.         {

  1400.             NCNN_LOGE("CreateDebugUtilsMessengerEXT failed %d", ret);

  1401.             return -1;

  1402.         }

  1403.     }

  1404. #endif // ENABLE_VALIDATION_LAYER

  1405. 

  1406.     init_instance_extension();

  1407. 

  1408.     uint32_t physicalDeviceCount = 0;

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

  1410.     if (ret != VK_SUCCESS)

  1411.     {

  1412.         NCNN_LOGE("vkEnumeratePhysicalDevices failed %d", ret);

  1413.         return -1;

  1414.     }

  1415. 

  1416.     if (physicalDeviceCount > NCNN_MAX_GPU_COUNT)

  1417.         physicalDeviceCount = NCNN_MAX_GPU_COUNT;

  1418. 

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

  1420. 

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

  1422.     if (ret != VK_SUCCESS)

  1423.     {

  1424.         NCNN_LOGE("vkEnumeratePhysicalDevices failed %d", ret);

  1425.         return -1;

  1426.     }

  1427. 

  1428.     // find proper device and queue

  1429.     int gpu_info_index = 0;

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

  1431.     {

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

  1433.         delete g_gpu_infos[gpu_info_index];

  1434.         g_gpu_infos[gpu_info_index] = new GpuInfo;

  1435.         GpuInfoPrivate& gpu_info = *(g_gpu_infos[gpu_info_index]->d);

  1436. 

  1437.         // device type

  1438.         VkPhysicalDeviceProperties physicalDeviceProperties;

  1439.         vkGetPhysicalDeviceProperties(physicalDevice, &physicalDeviceProperties);

  1440. 

  1441.         //         NCNN_LOGE("[%u] apiVersion = %u.%u.%u", i, VK_VERSION_MAJOR(physicalDeviceProperties.apiVersion),

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

  1443.         //         NCNN_LOGE("[%u] driverVersion = %u.%u.%u", i, VK_VERSION_MAJOR(physicalDeviceProperties.driverVersion),

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

  1445.         //         NCNN_LOGE("[%u] vendorID = %x", i, physicalDeviceProperties.vendorID);

  1446.         //         NCNN_LOGE("[%u] deviceID = %x", i, physicalDeviceProperties.deviceID);

  1447.         //         NCNN_LOGE("[%u] deviceType = %x", i, physicalDeviceProperties.deviceType);

  1448.         //         NCNN_LOGE("[%u] deviceName = %s", i, physicalDeviceProperties.deviceName);

  1449.         //         NCNN_LOGE("[%u] pipelineCacheUUID = %u", i, physicalDeviceProperties.pipelineCacheUUID);

  1450. 

  1451.         // mali

  1452.         // t760 = 0x13b5 0x7500001 / 0x7501000

  1453.         // t860 = 0x13b5 0x8602000

  1454.         // t880 = 0x13b5 0x8800020

  1455.         // g31  = 0x13b5 0x70930000

  1456.         // g51  = 0x13b5 0x70901010

  1457.         // g52  = 0x13b5 0x74021000 / 0x72120000

  1458.         // g71  = 0x13b5 0x60a00002

  1459.         // g72  = 0x13b5 0x62210001

  1460.         // g76  = 0x13b5 0x72110000

  1461.         // g77  = 0x13b5 0x90800011

  1462. 

  1463.         // adreno

  1464.         // 506 = 0x5143 0x5000600

  1465.         // 510 = 0x5143 0x5010000

  1466.         // 512 = 0x5143 0x5010200

  1467.         // 530 = 0x5143 0x5030004

  1468.         // 540 = 0x5143 0x5040001

  1469.         // 616 = 0x5143 0x6010600

  1470.         // 630 = 0x5143 0x6030001

  1471.         // 640 = 0x5143 0x6040001

  1472.         // 650 = 0x5143 0x6050002

  1473. 

  1474.         gpu_info.bug_storage_buffer_no_l1 = false;

  1475.         gpu_info.bug_corrupted_online_pipeline_cache = false;

  1476.         gpu_info.bug_implicit_fp16_arithmetic = false;

  1477.         gpu_info.bug_buffer_image_load_zero = false;

  1478. 

  1479.         if (physicalDeviceProperties.vendorID == 0x5143 && physicalDeviceProperties.apiVersion < VK_MAKE_VERSION(1, 0, 66))

  1480.         {

  1481.             // qcom adreno with old buggy driver cannot share created pipeline properly

  1482.             gpu_info.bug_corrupted_online_pipeline_cache = true;

  1483.         }

  1484. 

  1485.         if (physicalDeviceProperties.vendorID == 0x5143 && !(physicalDeviceProperties.deviceID == 0x6040001 || physicalDeviceProperties.deviceID == 0x6050002))

  1486.         {

  1487.             // NOTE but qcom855/qcom855plus/qcom865 are known exceptions

  1488.             // qcom adreno storage buffer without L1 cache

  1489.             gpu_info.bug_storage_buffer_no_l1 = true;

  1490.         }

  1491. 

  1492.         if (physicalDeviceProperties.vendorID == 0x5143 && physicalDeviceProperties.apiVersion < VK_MAKE_VERSION(1, 1, 87))

  1493.         {

  1494.             // HACK buffer2image before image-read dependency does not work properly

  1495.             // even promised with full image memory barrier on old adreno driver

  1496.             // TODO figure out a proper workaround without hurt speed too much

  1497.             // TODO only for old drivers

  1498.             gpu_info.bug_buffer_image_load_zero = true;

  1499.         }

  1500. 

  1501.         if (physicalDeviceProperties.vendorID == 0x13b5

  1502.                 && (physicalDeviceProperties.deviceID == 0x7500001

  1503.                     || physicalDeviceProperties.deviceID == 0x7501000

  1504.                     || physicalDeviceProperties.deviceID == 0x8602000

  1505.                     || physicalDeviceProperties.deviceID == 0x8800020

  1506.                     || physicalDeviceProperties.deviceID == 0x70930000

  1507.                     || physicalDeviceProperties.deviceID == 0x70901010

  1508.                     || physicalDeviceProperties.deviceID == 0x72120000

  1509.                     || physicalDeviceProperties.deviceID == 0x74021000

  1510.                     || physicalDeviceProperties.deviceID == 0x60a00002

  1511.                     || physicalDeviceProperties.deviceID == 0x62210001))

  1512.         {

  1513.             // NOTE rk3288/rk3399/t880/g31/g51/g52/g71/g72

  1514.             // however, g76/g77 has explicit fp16 arithmetic

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

  1516.             gpu_info.bug_implicit_fp16_arithmetic = true;

  1517.         }

  1518. 

  1519.         if (physicalDeviceProperties.vendorID == 0x5143

  1520.                 && (physicalDeviceProperties.deviceID == 0x6030001

  1521.                     || physicalDeviceProperties.deviceID == 0x6040001

  1522.                     || physicalDeviceProperties.deviceID == 0x6050002))

  1523.         {

  1524.             // TODO enable devices other than qcom845/qcom855/qcom855plus/qcom865

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

  1526.             gpu_info.bug_implicit_fp16_arithmetic = true;

  1527.         }

  1528. 

  1529.         gpu_info.physical_device = physicalDevice;

  1530. 

  1531.         // info

  1532.         gpu_info.api_version = physicalDeviceProperties.apiVersion;

  1533.         gpu_info.driver_version = physicalDeviceProperties.driverVersion;

  1534.         gpu_info.vendor_id = physicalDeviceProperties.vendorID;

  1535.         gpu_info.device_id = physicalDeviceProperties.deviceID;

  1536.         memcpy(gpu_info.device_name, physicalDeviceProperties.deviceName, VK_MAX_PHYSICAL_DEVICE_NAME_SIZE);

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

  1538. 

  1539.         if (physicalDeviceProperties.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU)

  1540.             gpu_info.type = 0;

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

  1542.             gpu_info.type = 1;

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

  1544.             gpu_info.type = 2;

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

  1546.             gpu_info.type = 3;

  1547.         else

  1548.             gpu_info.type = -1;

  1549. 

  1550.         // device capability

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

  1552. 

  1553.         gpu_info.max_workgroup_count_x = physicalDeviceProperties.limits.maxComputeWorkGroupCount[0];

  1554.         gpu_info.max_workgroup_count_y = physicalDeviceProperties.limits.maxComputeWorkGroupCount[1];

  1555.         gpu_info.max_workgroup_count_z = physicalDeviceProperties.limits.maxComputeWorkGroupCount[2];

  1556. 

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

  1558. 

  1559.         gpu_info.max_workgroup_size_x = physicalDeviceProperties.limits.maxComputeWorkGroupSize[0];

  1560.         gpu_info.max_workgroup_size_y = physicalDeviceProperties.limits.maxComputeWorkGroupSize[1];

  1561.         gpu_info.max_workgroup_size_z = physicalDeviceProperties.limits.maxComputeWorkGroupSize[2];

  1562. 

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

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

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

  1566.         gpu_info.buffer_image_granularity = physicalDeviceProperties.limits.bufferImageGranularity;

  1567.         gpu_info.max_image_dimension_1d = physicalDeviceProperties.limits.maxImageDimension1D;

  1568.         gpu_info.max_image_dimension_2d = physicalDeviceProperties.limits.maxImageDimension2D;

  1569.         gpu_info.max_image_dimension_3d = physicalDeviceProperties.limits.maxImageDimension3D;

  1570. 

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

  1572. 

  1573.         //         NCNN_LOGE("[%u] max_shared_memory_size = %u", i, gpu_info.max_shared_memory_size);

  1574.         //         NCNN_LOGE("[%u] max_workgroup_count = %u %u %u", i, gpu_info.max_workgroup_count[0], gpu_info.max_workgroup_count[1], gpu_info.max_workgroup_count[2]);

  1575.         //         NCNN_LOGE("[%u] max_workgroup_invocations = %u", i, gpu_info.max_workgroup_invocations);

  1576.         //         NCNN_LOGE("[%u] max_workgroup_size = %u %u %u", i, gpu_info.max_workgroup_size[0], gpu_info.max_workgroup_size[1], gpu_info.max_workgroup_size[2]);

  1577.         //         NCNN_LOGE("[%u] memory_map_alignment = %lu", i, gpu_info.memory_map_alignment);

  1578.         //         NCNN_LOGE("[%u] buffer_offset_alignment = %lu", i, gpu_info.buffer_offset_alignment);

  1579. 

  1580.         // find compute queue

  1581.         uint32_t queueFamilyPropertiesCount;

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

  1583. 

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

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

  1586. 

  1587.         gpu_info.compute_queue_family_index = find_device_compute_queue(queueFamilyProperties);

  1588.         gpu_info.graphics_queue_family_index = find_device_graphics_queue(queueFamilyProperties);

  1589.         gpu_info.transfer_queue_family_index = find_device_transfer_queue(queueFamilyProperties);

  1590. 

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

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

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

  1594. 

  1595.         gpu_info.unified_compute_transfer_queue = gpu_info.compute_queue_family_index == gpu_info.transfer_queue_family_index;

  1596. 

  1597.         // additional device properties

  1598.         gpu_info.subgroup_size = 64;

  1599.         gpu_info.support_subgroup_basic = false;

  1600.         gpu_info.support_subgroup_vote = false;

  1601.         gpu_info.support_subgroup_ballot = false;

  1602.         gpu_info.support_subgroup_shuffle = false;

  1603.         if (support_VK_KHR_get_physical_device_properties2)

  1604.         {

  1605.             void* queryDeviceProperties = 0;

  1606. 

  1607.             // query subgroup

  1608.             VkPhysicalDeviceSubgroupProperties physicalDeviceSubgroupProperties;

  1609.             physicalDeviceSubgroupProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SUBGROUP_PROPERTIES;

  1610.             physicalDeviceSubgroupProperties.pNext = queryDeviceProperties;

  1611.             if (VK_VERSION_MAJOR(instance_api_version) >= 1 && VK_VERSION_MINOR(instance_api_version) >= 1)

  1612.             {

  1613.                 queryDeviceProperties = &physicalDeviceSubgroupProperties;

  1614.             }

  1615. 

  1616.             VkPhysicalDeviceProperties2KHR queryProperties;

  1617.             queryProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR;

  1618.             queryProperties.pNext = queryDeviceProperties;

  1619. 

  1620.             vkGetPhysicalDeviceProperties2KHR(physicalDevice, &queryProperties);

  1621. 

  1622.             if (VK_VERSION_MAJOR(instance_api_version) >= 1 && VK_VERSION_MINOR(instance_api_version) >= 1)

  1623.             {

  1624.                 gpu_info.subgroup_size = physicalDeviceSubgroupProperties.subgroupSize;

  1625.                 if (physicalDeviceSubgroupProperties.supportedStages & VK_SHADER_STAGE_COMPUTE_BIT)

  1626.                 {

  1627.                     gpu_info.support_subgroup_basic = physicalDeviceSubgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_BASIC_BIT;

  1628.                     gpu_info.support_subgroup_vote = physicalDeviceSubgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_VOTE_BIT;

  1629.                     gpu_info.support_subgroup_ballot = physicalDeviceSubgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_BALLOT_BIT;

  1630.                     gpu_info.support_subgroup_shuffle = physicalDeviceSubgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_SHUFFLE_BIT;

  1631.                 }

  1632.             }

  1633.             else

  1634.             {

  1635.                 if (physicalDeviceProperties.vendorID == 0x5143) // qcom adreno prefer very large workgroup :P

  1636.                     gpu_info.subgroup_size = 128;

  1637.                 if (physicalDeviceProperties.vendorID == 0x13b5) // arm mali

  1638.                     gpu_info.subgroup_size = 16;

  1639.                 if (physicalDeviceProperties.vendorID == 0x1010) // imgtec powervr

  1640.                     gpu_info.subgroup_size = 32;

  1641.                 if (physicalDeviceProperties.vendorID == 0x1002) // amd

  1642.                     gpu_info.subgroup_size = 64;

  1643.                 if (physicalDeviceProperties.vendorID == 0x10de) // nvidia

  1644.                     gpu_info.subgroup_size = 32;

  1645.                 if (physicalDeviceProperties.vendorID == 0x8086) // intel

  1646.                     gpu_info.subgroup_size = 32;

  1647.             }

  1648.         }

  1649. 

  1650.         // cache memory properties

  1651.         vkGetPhysicalDeviceMemoryProperties(physicalDevice, &gpu_info.physical_device_memory_properties);

  1652. 

  1653.         // get device extension

  1654.         uint32_t deviceExtensionPropertyCount = 0;

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

  1656.         if (ret != VK_SUCCESS)

  1657.         {

  1658.             NCNN_LOGE("vkEnumerateDeviceExtensionProperties failed %d", ret);

  1659.             return -1;

  1660.         }

  1661. 

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

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

  1664.         if (ret != VK_SUCCESS)

  1665.         {

  1666.             NCNN_LOGE("vkEnumerateDeviceExtensionProperties failed %d", ret);

  1667.             return -1;

  1668.         }

  1669. 

  1670.         // extension capability

  1671.         gpu_info.support_VK_KHR_8bit_storage = 0;

  1672.         gpu_info.support_VK_KHR_16bit_storage = 0;

  1673.         gpu_info.support_VK_KHR_bind_memory2 = 0;

  1674.         gpu_info.support_VK_KHR_buffer_device_address = 0;

  1675.         gpu_info.support_VK_KHR_create_renderpass2 = 0;

  1676.         gpu_info.support_VK_KHR_cooperative_matrix = 0;

  1677.         gpu_info.support_VK_KHR_dedicated_allocation = 0;

  1678.         gpu_info.support_VK_KHR_descriptor_update_template = 0;

  1679.         gpu_info.support_VK_KHR_external_memory = 0;

  1680.         gpu_info.support_VK_KHR_get_memory_requirements2 = 0;

  1681.         gpu_info.support_VK_KHR_maintenance1 = 0;

  1682.         gpu_info.support_VK_KHR_maintenance2 = 0;

  1683.         gpu_info.support_VK_KHR_maintenance3 = 0;

  1684.         gpu_info.support_VK_KHR_multiview = 0;

  1685.         gpu_info.support_VK_KHR_portability_subset = 0;

  1686.         gpu_info.support_VK_KHR_push_descriptor = 0;

  1687.         gpu_info.support_VK_KHR_sampler_ycbcr_conversion = 0;

  1688.         gpu_info.support_VK_KHR_shader_float16_int8 = 0;

  1689.         gpu_info.support_VK_KHR_shader_float_controls = 0;

  1690.         gpu_info.support_VK_KHR_storage_buffer_storage_class = 0;

  1691.         gpu_info.support_VK_KHR_swapchain = 0;

  1692.         gpu_info.support_VK_EXT_buffer_device_address = 0;

  1693.         gpu_info.support_VK_EXT_descriptor_indexing = 0;

  1694.         gpu_info.support_VK_EXT_memory_budget = 0;

  1695.         gpu_info.support_VK_EXT_memory_priority = 0;

  1696.         gpu_info.support_VK_EXT_queue_family_foreign = 0;

  1697.         gpu_info.support_VK_AMD_device_coherent_memory = 0;

  1698. #if __ANDROID_API__ >= 26

  1699.         gpu_info.support_VK_ANDROID_external_memory_android_hardware_buffer = 0;

  1700. #endif // __ANDROID_API__ >= 26

  1701.         gpu_info.support_VK_NV_cooperative_matrix = 0;

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

  1703.         {

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

  1705.             // NCNN_LOGE("device extension %s = %u", exp.extensionName, exp.specVersion);

  1706. 

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

  1708.                 gpu_info.support_VK_KHR_8bit_storage = exp.specVersion;

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

  1710.                 gpu_info.support_VK_KHR_16bit_storage = exp.specVersion;

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

  1712.                 gpu_info.support_VK_KHR_bind_memory2 = exp.specVersion;

  1713.             else if (strcmp(exp.extensionName, "VK_KHR_buffer_device_address") == 0)

  1714.                 gpu_info.support_VK_KHR_buffer_device_address = exp.specVersion;

  1715.             else if (strcmp(exp.extensionName, "VK_KHR_create_renderpass2") == 0)

  1716.                 gpu_info.support_VK_KHR_create_renderpass2 = exp.specVersion;

  1717.             else if (strcmp(exp.extensionName, "VK_KHR_cooperative_matrix") == 0)

  1718.                 gpu_info.support_VK_KHR_cooperative_matrix = exp.specVersion;

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

  1720.                 gpu_info.support_VK_KHR_dedicated_allocation = exp.specVersion;

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

  1722.                 gpu_info.support_VK_KHR_descriptor_update_template = exp.specVersion;

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

  1724.                 gpu_info.support_VK_KHR_external_memory = exp.specVersion;

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

  1726.                 gpu_info.support_VK_KHR_get_memory_requirements2 = exp.specVersion;

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

  1728.                 gpu_info.support_VK_KHR_maintenance1 = exp.specVersion;

  1729.             else if (strcmp(exp.extensionName, "VK_KHR_maintenance2") == 0)

  1730.                 gpu_info.support_VK_KHR_maintenance2 = exp.specVersion;

  1731.             else if (strcmp(exp.extensionName, "VK_KHR_maintenance3") == 0)

  1732.                 gpu_info.support_VK_KHR_maintenance3 = exp.specVersion;

  1733.             else if (strcmp(exp.extensionName, "VK_KHR_multiview") == 0)

  1734.                 gpu_info.support_VK_KHR_multiview = exp.specVersion;

  1735.             else if (strcmp(exp.extensionName, "VK_KHR_portability_subset") == 0)

  1736.                 gpu_info.support_VK_KHR_portability_subset = exp.specVersion;

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

  1738.                 gpu_info.support_VK_KHR_push_descriptor = exp.specVersion;

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

  1740.                 gpu_info.support_VK_KHR_sampler_ycbcr_conversion = exp.specVersion;

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

  1742.                 gpu_info.support_VK_KHR_shader_float16_int8 = exp.specVersion;

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

  1744.                 gpu_info.support_VK_KHR_shader_float_controls = exp.specVersion;

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

  1746.                 gpu_info.support_VK_KHR_storage_buffer_storage_class = exp.specVersion;

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

  1748.                 gpu_info.support_VK_KHR_swapchain = exp.specVersion;

  1749.             else if (strcmp(exp.extensionName, "VK_EXT_buffer_device_address") == 0)

  1750.                 gpu_info.support_VK_EXT_buffer_device_address = exp.specVersion;

  1751.             else if (strcmp(exp.extensionName, "VK_EXT_descriptor_indexing") == 0)

  1752.                 gpu_info.support_VK_EXT_descriptor_indexing = exp.specVersion;

  1753.             else if (strcmp(exp.extensionName, "VK_EXT_memory_budget") == 0)

  1754.                 gpu_info.support_VK_EXT_memory_budget = exp.specVersion;

  1755.             else if (strcmp(exp.extensionName, "VK_EXT_memory_priority") == 0)

  1756.                 gpu_info.support_VK_EXT_memory_priority = exp.specVersion;

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

  1758.                 gpu_info.support_VK_EXT_queue_family_foreign = exp.specVersion;

  1759.             else if (strcmp(exp.extensionName, "VK_AMD_device_coherent_memory") == 0)

  1760.                 gpu_info.support_VK_AMD_device_coherent_memory = exp.specVersion;

  1761. #if __ANDROID_API__ >= 26

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

  1763.                 gpu_info.support_VK_ANDROID_external_memory_android_hardware_buffer = exp.specVersion;

  1764. #endif // __ANDROID_API__ >= 26

  1765.             else if (strcmp(exp.extensionName, "VK_NV_cooperative_matrix") == 0)

  1766.                 gpu_info.support_VK_NV_cooperative_matrix = exp.specVersion;

  1767.         }

  1768. 

  1769.         if (gpu_info.support_VK_KHR_buffer_device_address)

  1770.         {

  1771.             // we prefer khr extension

  1772.             gpu_info.support_VK_EXT_buffer_device_address = 0;

  1773.         }

  1774. 

  1775.         if (gpu_info.support_VK_KHR_cooperative_matrix)

  1776.         {

  1777.             // we prefer khr extension

  1778.             gpu_info.support_VK_NV_cooperative_matrix = 0;

  1779.         }

  1780. 

  1781.         // check features

  1782.         gpu_info.support_fp16_packed = true;

  1783.         gpu_info.support_fp16_storage = false;

  1784.         gpu_info.support_fp16_uniform = false;

  1785.         gpu_info.support_fp16_arithmetic = false;

  1786.         gpu_info.support_int8_packed = true;

  1787.         gpu_info.support_int8_storage = false;

  1788.         gpu_info.support_int8_uniform = false;

  1789.         gpu_info.support_int8_arithmetic = false;

  1790.         gpu_info.support_ycbcr_conversion = false;

  1791.         gpu_info.support_cooperative_matrix = false;

  1792.         gpu_info.support_cooperative_matrix_8_8_16 = false;

  1793.         gpu_info.support_cooperative_matrix_16_8_8 = false;

  1794.         gpu_info.support_cooperative_matrix_16_8_16 = false;

  1795.         gpu_info.support_cooperative_matrix_16_16_16 = false;

  1796.         if (support_VK_KHR_get_physical_device_properties2)

  1797.         {

  1798.             void* queryExtensionFeatures = 0;

  1799. 

  1800.             // query int8 storage

  1801.             VkPhysicalDevice8BitStorageFeaturesKHR query8BitStorageFeatures;

  1802.             query8BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_8BIT_STORAGE_FEATURES_KHR;

  1803.             query8BitStorageFeatures.pNext = 0;

  1804.             if (gpu_info.support_VK_KHR_8bit_storage)

  1805.             {

  1806.                 query8BitStorageFeatures.pNext = queryExtensionFeatures;

  1807.                 queryExtensionFeatures = &query8BitStorageFeatures;

  1808.             }

  1809. 

  1810.             // query fp16/int16 storage

  1811.             VkPhysicalDevice16BitStorageFeaturesKHR query16BitStorageFeatures;

  1812.             query16BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR;

  1813.             query16BitStorageFeatures.pNext = 0;

  1814.             if (gpu_info.support_VK_KHR_16bit_storage)

  1815.             {

  1816.                 query16BitStorageFeatures.pNext = queryExtensionFeatures;

  1817.                 queryExtensionFeatures = &query16BitStorageFeatures;

  1818.             }

  1819. 

  1820.             // query fp16/int8 arithmetic

  1821.             VkPhysicalDeviceFloat16Int8FeaturesKHR queryFloat16Int8Features;

  1822.             queryFloat16Int8Features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR;

  1823.             queryFloat16Int8Features.pNext = 0;

  1824.             if (gpu_info.support_VK_KHR_shader_float16_int8)

  1825.             {

  1826.                 queryFloat16Int8Features.pNext = queryExtensionFeatures;

  1827.                 queryExtensionFeatures = &queryFloat16Int8Features;

  1828.             }

  1829. 

  1830.             // query ycbcr_conversion

  1831.             VkPhysicalDeviceSamplerYcbcrConversionFeaturesKHR querySamplerYcbcrConversionFeatures;

  1832.             querySamplerYcbcrConversionFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SAMPLER_YCBCR_CONVERSION_FEATURES_KHR;

  1833.             querySamplerYcbcrConversionFeatures.pNext = 0;

  1834.             if (gpu_info.support_VK_KHR_sampler_ycbcr_conversion)

  1835.             {

  1836.                 querySamplerYcbcrConversionFeatures.pNext = queryExtensionFeatures;

  1837.                 queryExtensionFeatures = &querySamplerYcbcrConversionFeatures;

  1838.             }

  1839. 

  1840.             // query cooperative_matrix

  1841.             VkPhysicalDeviceCooperativeMatrixFeaturesKHR queryCooperativeMatrixFeatures;

  1842.             queryCooperativeMatrixFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_FEATURES_KHR;

  1843.             queryCooperativeMatrixFeatures.pNext = 0;

  1844.             VkPhysicalDeviceCooperativeMatrixFeaturesNV queryCooperativeMatrixFeaturesNV;

  1845.             queryCooperativeMatrixFeaturesNV.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_FEATURES_NV;

  1846.             queryCooperativeMatrixFeaturesNV.pNext = 0;

  1847.             if (gpu_info.support_VK_KHR_cooperative_matrix)

  1848.             {

  1849.                 queryCooperativeMatrixFeatures.pNext = queryExtensionFeatures;

  1850.                 queryExtensionFeatures = &queryCooperativeMatrixFeatures;

  1851.             }

  1852.             else if (gpu_info.support_VK_NV_cooperative_matrix)

  1853.             {

  1854.                 queryCooperativeMatrixFeaturesNV.pNext = queryExtensionFeatures;

  1855.                 queryExtensionFeatures = &queryCooperativeMatrixFeaturesNV;

  1856.             }

  1857. 

  1858.             VkPhysicalDeviceFeatures2KHR queryFeatures;

  1859.             queryFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2_KHR;

  1860.             queryFeatures.pNext = queryExtensionFeatures;

  1861. 

  1862.             vkGetPhysicalDeviceFeatures2KHR(physicalDevice, &queryFeatures);

  1863. 

  1864.             if (gpu_info.support_VK_KHR_8bit_storage)

  1865.             {

  1866.                 gpu_info.support_int8_storage = query8BitStorageFeatures.storageBuffer8BitAccess;

  1867.                 gpu_info.support_int8_uniform = query8BitStorageFeatures.uniformAndStorageBuffer8BitAccess;

  1868.             }

  1869.             if (gpu_info.support_VK_KHR_16bit_storage && queryFeatures.features.shaderStorageImageExtendedFormats)

  1870.             {

  1871.                 // shaderStorageImageExtendedFormats enables r16f format in storage image

  1872.                 gpu_info.support_fp16_storage = query16BitStorageFeatures.storageBuffer16BitAccess;

  1873.                 gpu_info.support_fp16_uniform = query16BitStorageFeatures.uniformAndStorageBuffer16BitAccess;

  1874.             }

  1875.             if (gpu_info.support_VK_KHR_shader_float16_int8)

  1876.             {

  1877.                 gpu_info.support_fp16_arithmetic = queryFloat16Int8Features.shaderFloat16;

  1878.                 gpu_info.support_int8_arithmetic = queryFloat16Int8Features.shaderInt8;

  1879.             }

  1880.             if (gpu_info.support_VK_KHR_sampler_ycbcr_conversion)

  1881.             {

  1882.                 gpu_info.support_ycbcr_conversion = querySamplerYcbcrConversionFeatures.samplerYcbcrConversion;

  1883.             }

  1884.             if (gpu_info.support_VK_KHR_cooperative_matrix)

  1885.             {

  1886.                 gpu_info.support_cooperative_matrix = queryCooperativeMatrixFeatures.cooperativeMatrix;

  1887.             }

  1888.             else if (gpu_info.support_VK_NV_cooperative_matrix)

  1889.             {

  1890.                 gpu_info.support_cooperative_matrix = queryCooperativeMatrixFeaturesNV.cooperativeMatrix;

  1891.             }

  1892.         }

  1893.         else

  1894.         {

  1895.             //             // TODO

  1896.             //             VkPhysicalDeviceFeatures features;

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

  1898.         }

  1899. 

  1900.         if (physicalDeviceProperties.vendorID == 0x13b5 && physicalDeviceProperties.apiVersion < VK_MAKE_VERSION(1, 0, 82))

  1901.         {

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

  1903.             gpu_info.support_fp16_storage = false;

  1904.         }

  1905. 

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

  1907.         {

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

  1909.             gpu_info.support_fp16_storage = false;

  1910.         }

  1911. 

  1912.         if (gpu_info.bug_implicit_fp16_arithmetic)

  1913.         {

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

  1915.             gpu_info.support_fp16_arithmetic = true;

  1916.         }

  1917. 

  1918.         if (physicalDeviceProperties.vendorID == 0x5143 && !gpu_info.support_fp16_storage)

  1919.         {

  1920.             // fp16 arithmetic yields wrong result on old adreno drivers :(

  1921.             gpu_info.support_fp16_arithmetic = false;

  1922.         }

  1923. 

  1924.         if (gpu_info.support_cooperative_matrix)

  1925.         {

  1926.             // query supported cooperative matrix types and operations

  1927.             if (gpu_info.support_VK_KHR_cooperative_matrix)

  1928.             {

  1929.                 uint32_t propertyCount = 0;

  1930.                 ret = vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR(physicalDevice, &propertyCount, 0);

  1931.                 if (ret != VK_SUCCESS)

  1932.                 {

  1933.                     NCNN_LOGE("vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR failed %d", ret);

  1934.                 }

  1935. 

  1936.                 std::vector<VkCooperativeMatrixPropertiesKHR> properties(propertyCount);

  1937.                 ret = vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR(physicalDevice, &propertyCount, properties.data());

  1938.                 if (ret != VK_SUCCESS)

  1939.                 {

  1940.                     NCNN_LOGE("vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR failed %d", ret);

  1941.                 }

  1942. 

  1943.                 for (uint32_t j = 0; j < properties.size(); j++)

  1944.                 {

  1945.                     const VkCooperativeMatrixPropertiesKHR& cmp = properties[j];

  1946.                     // NCNN_LOGE("cpm %2d %2d %2d  %d %d %d %d  %d", cmp.MSize, cmp.NSize, cmp.KSize, cmp.AType, cmp.BType, cmp.CType, cmp.ResultType, cmp.scope);

  1947. 

  1948.                     if (cmp.MSize == 8 && cmp.NSize == 8 && cmp.KSize == 16

  1949.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_KHR && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_KHR

  1950.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_KHR && cmp.ResultType == VK_COMPONENT_TYPE_FLOAT32_KHR

  1951.                             && cmp.scope == VK_SCOPE_SUBGROUP_KHR)

  1952.                     {

  1953.                         gpu_info.support_cooperative_matrix_8_8_16 = true;

  1954.                     }

  1955.                     if (cmp.MSize == 16 && cmp.NSize == 8 && cmp.KSize == 8

  1956.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_KHR && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_KHR

  1957.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_KHR && cmp.ResultType == VK_COMPONENT_TYPE_FLOAT32_KHR

  1958.                             && cmp.scope == VK_SCOPE_SUBGROUP_KHR)

  1959.                     {

  1960.                         gpu_info.support_cooperative_matrix_16_8_8 = true;

  1961.                     }

  1962.                     if (cmp.MSize == 16 && cmp.NSize == 8 && cmp.KSize == 16

  1963.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_KHR && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_KHR

  1964.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_KHR && cmp.ResultType == VK_COMPONENT_TYPE_FLOAT32_KHR

  1965.                             && cmp.scope == VK_SCOPE_SUBGROUP_KHR)

  1966.                     {

  1967.                         gpu_info.support_cooperative_matrix_16_8_16 = true;

  1968.                     }

  1969.                     if (cmp.MSize == 16 && cmp.NSize == 16 && cmp.KSize == 16

  1970.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_KHR && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_KHR

  1971.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_KHR && cmp.ResultType == VK_COMPONENT_TYPE_FLOAT32_KHR

  1972.                             && cmp.scope == VK_SCOPE_SUBGROUP_KHR)

  1973.                     {

  1974.                         gpu_info.support_cooperative_matrix_16_16_16 = true;

  1975.                     }

  1976.                 }

  1977.             }

  1978.             else

  1979.             {

  1980.                 uint32_t propertyCount = 0;

  1981.                 ret = vkGetPhysicalDeviceCooperativeMatrixPropertiesNV(physicalDevice, &propertyCount, 0);

  1982.                 if (ret != VK_SUCCESS)

  1983.                 {

  1984.                     NCNN_LOGE("vkGetPhysicalDeviceCooperativeMatrixPropertiesNV failed %d", ret);

  1985.                 }

  1986. 

  1987.                 std::vector<VkCooperativeMatrixPropertiesNV> properties(propertyCount);

  1988.                 for (uint32_t j = 0; j < properties.size(); j++)

  1989.                 {

  1990.                     properties[j].sType = VK_STRUCTURE_TYPE_COOPERATIVE_MATRIX_PROPERTIES_NV;

  1991.                     properties[j].pNext = 0;

  1992.                 }

  1993.                 ret = vkGetPhysicalDeviceCooperativeMatrixPropertiesNV(physicalDevice, &propertyCount, properties.data());

  1994.                 if (ret != VK_SUCCESS)

  1995.                 {

  1996.                     NCNN_LOGE("vkGetPhysicalDeviceCooperativeMatrixPropertiesNV failed %d", ret);

  1997.                 }

  1998. 

  1999.                 for (uint32_t j = 0; j < properties.size(); j++)

  2000.                 {

  2001.                     const VkCooperativeMatrixPropertiesNV& cmp = properties[j];

  2002.                     // NCNN_LOGE("cpm %2d %2d %2d  %d %d %d %d  %d", cmp.MSize, cmp.NSize, cmp.KSize, cmp.AType, cmp.BType, cmp.CType, cmp.DType, cmp.scope);

  2003. 

  2004.                     if (cmp.MSize == 8 && cmp.NSize == 8 && cmp.KSize == 16

  2005.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_NV && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_NV

  2006.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_NV && cmp.DType == VK_COMPONENT_TYPE_FLOAT32_NV

  2007.                             && cmp.scope == VK_SCOPE_SUBGROUP_NV)

  2008.                     {

  2009.                         gpu_info.support_cooperative_matrix_8_8_16 = true;

  2010.                     }

  2011.                     if (cmp.MSize == 16 && cmp.NSize == 8 && cmp.KSize == 8

  2012.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_NV && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_NV

  2013.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_NV && cmp.DType == VK_COMPONENT_TYPE_FLOAT32_NV

  2014.                             && cmp.scope == VK_SCOPE_SUBGROUP_NV)

  2015.                     {

  2016.                         gpu_info.support_cooperative_matrix_16_8_8 = true;

  2017.                     }

  2018.                     if (cmp.MSize == 16 && cmp.NSize == 8 && cmp.KSize == 16

  2019.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_NV && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_NV

  2020.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_NV && cmp.DType == VK_COMPONENT_TYPE_FLOAT32_NV

  2021.                             && cmp.scope == VK_SCOPE_SUBGROUP_NV)

  2022.                     {

  2023.                         gpu_info.support_cooperative_matrix_16_8_16 = true;

  2024.                     }

  2025.                     if (cmp.MSize == 16 && cmp.NSize == 16 && cmp.KSize == 16

  2026.                             && cmp.AType == VK_COMPONENT_TYPE_FLOAT16_NV && cmp.BType == VK_COMPONENT_TYPE_FLOAT16_NV

  2027.                             && cmp.CType == VK_COMPONENT_TYPE_FLOAT32_NV && cmp.DType == VK_COMPONENT_TYPE_FLOAT32_NV

  2028.                             && cmp.scope == VK_SCOPE_SUBGROUP_NV)

  2029.                     {

  2030.                         gpu_info.support_cooperative_matrix_16_16_16 = true;

  2031.                     }

  2032.                 }

  2033.             }

  2034.         }

  2035. 

  2036.         NCNN_LOGE("[%u %s]  queueC=%u[%u]  queueG=%u[%u]  queueT=%u[%u]", i, physicalDeviceProperties.deviceName,

  2037.                   gpu_info.compute_queue_family_index, gpu_info.compute_queue_count,

  2038.                   gpu_info.graphics_queue_family_index, gpu_info.graphics_queue_count,

  2039.                   gpu_info.transfer_queue_family_index, gpu_info.transfer_queue_count);

  2040. 

  2041.         NCNN_LOGE("[%u %s]  bugsbn1=%d  bugbilz=%d  bugcopc=%d  bugihfa=%d", i, physicalDeviceProperties.deviceName,

  2042.                   gpu_info.bug_storage_buffer_no_l1, gpu_info.bug_buffer_image_load_zero, gpu_info.bug_corrupted_online_pipeline_cache, gpu_info.bug_implicit_fp16_arithmetic);

  2043. 

  2044.         NCNN_LOGE("[%u %s]  fp16-p/s/u/a=%d/%d/%d/%d  int8-p/s/u/a=%d/%d/%d/%d", i, physicalDeviceProperties.deviceName,

  2045.                   gpu_info.support_fp16_packed, gpu_info.support_fp16_storage, gpu_info.support_fp16_uniform, gpu_info.support_fp16_arithmetic,

  2046.                   gpu_info.support_int8_packed, gpu_info.support_int8_storage, gpu_info.support_int8_uniform, gpu_info.support_int8_arithmetic);

  2047. 

  2048.         NCNN_LOGE("[%u %s]  subgroup=%u  basic/vote/ballot/shuffle=%d/%d/%d/%d", i, physicalDeviceProperties.deviceName,

  2049.                   gpu_info.subgroup_size, gpu_info.support_subgroup_basic, gpu_info.support_subgroup_vote,

  2050.                   gpu_info.support_subgroup_ballot, gpu_info.support_subgroup_shuffle);

  2051. 

  2052.         NCNN_LOGE("[%u %s]  fp16-8x8x16/16x8x8/16x8x16/16x16x16=%d/%d/%d/%d", i, physicalDeviceProperties.deviceName,

  2053.                   gpu_info.support_cooperative_matrix_8_8_16, gpu_info.support_cooperative_matrix_16_8_8,

  2054.                   gpu_info.support_cooperative_matrix_16_8_16, gpu_info.support_cooperative_matrix_16_16_16);

  2055. 

  2056.         gpu_info_index++;

  2057.     }

  2058. 

  2059.     g_gpu_count = gpu_info_index;

  2060. 

  2061.     // the default gpu device

  2062.     g_default_gpu_index = find_default_vulkan_device_index();

  2063. 

  2064.     g_instance.glslang_initialized = glslang::InitializeProcess();

  2065. 

  2066.     // the global __ncnn_vulkan_instance_holder destructor will call destroy_gpu_instance() on exit

  2067.     // but it seems to be too late for nvidia driver :(

  2068.     // driver's internal data structure has been destroyed when called, causing segfault

  2069.     // atexit() seems to be helpful for calling it earlier    --- nihui

  2070.     static int destroy_gpu_instance_atexit_registered = 0;

  2071.     if (!destroy_gpu_instance_atexit_registered)

  2072.     {

  2073.         atexit(destroy_gpu_instance);

  2074.         destroy_gpu_instance_atexit_registered = 1;

  2075.     }

  2076. 

  2077.     return 0;

  2078. }

  2079. 

  2080. VkInstance get_gpu_instance()

  2081. {

  2082.     return (VkInstance)g_instance;

  2083. }

  2084. 

  2085. void destroy_gpu_instance()

  2086. {

  2087.     MutexLockGuard lock(g_instance_lock);

  2088. 

  2089.     if (g_instance.created == 0)

  2090.         return;

  2091. 

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

  2093.     {

  2094.         VulkanDevice* vulkan_device = g_default_vkdev[i];

  2095.         if (vulkan_device)

  2096.         {

  2097.             VkDevice vkdev = g_default_vkdev[i]->vkdevice();

  2098.             if (vkdev)

  2099.             {

  2100.                 vkDeviceWaitIdle(vkdev);

  2101.             }

  2102.         }

  2103.     }

  2104. 

  2105.     // NCNN_LOGE("destroy_gpu_instance");

  2106. 

  2107.     if (g_instance.glslang_initialized)

  2108.     {

  2109.         glslang::FinalizeProcess();

  2110.         g_instance.glslang_initialized = false;

  2111.     }

  2112. 

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

  2114.     {

  2115.         delete g_default_vkdev[i];

  2116.         g_default_vkdev[i] = 0;

  2117. 

  2118.         delete g_gpu_infos[i];

  2119.         g_gpu_infos[i] = 0;

  2120.     }

  2121. 

  2122. #if ENABLE_VALIDATION_LAYER

  2123.     if (support_VK_EXT_debug_utils && g_instance.callback)

  2124.     {

  2125.         DestroyDebugUtilsMessengerEXT(g_instance, g_instance.callback, NULL);

  2126.         g_instance.callback = 0;

  2127.     }

  2128. #endif // ENABLE_VALIDATION_LAYER

  2129. 

  2130.     if (vkDestroyInstance)

  2131.     {

  2132.         vkDestroyInstance(g_instance, 0);

  2133.         vkDestroyInstance = 0;

  2134.     }

  2135. 

  2136.     g_instance.instance = 0;

  2137. 

  2138. #if NCNN_SIMPLEVK

  2139.     unload_vulkan_driver();

  2140. #endif

  2141. 

  2142.     g_instance.created = 0;

  2143. }

  2144. 

  2145. static void try_create_gpu_instance()

  2146. {

  2147.     {

  2148.         MutexLockGuard lock(g_instance_lock);

  2149. 

  2150.         if (g_instance.created != 0)

  2151.             return;

  2152.     }

  2153. 

  2154.     create_gpu_instance();

  2155. }

  2156. 

  2157. int get_gpu_count()

  2158. {

  2159.     try_create_gpu_instance();

  2160. 

  2161.     return g_gpu_count;

  2162. }

  2163. 

  2164. int get_default_gpu_index()

  2165. {

  2166.     try_create_gpu_instance();

  2167. 

  2168.     return g_default_gpu_index;

  2169. }

  2170. 

  2171. const GpuInfo& get_gpu_info(int device_index)

  2172. {

  2173.     try_create_gpu_instance();

  2174. 

  2175.     return *g_gpu_infos[device_index];

  2176. }

  2177. 

  2178. class VkDummyAllocator : public VkBlobAllocator

  2179. {

  2180. public:

  2181.     // NOTE 16k is large enough I think ...

  2182.     VkDummyAllocator(const VulkanDevice* _vkdev)

  2183.         : VkBlobAllocator(_vkdev, 16 * 1024)

  2184.     {

  2185.     }

  2186. };

  2187. 

  2188. class VkDummyCompute : public VkCompute

  2189. {

  2190. public:

  2191.     VkDummyCompute(const VulkanDevice* _vkdev)

  2192.         : VkCompute(_vkdev)

  2193.     {

  2194.     }

  2195. 

  2196.     void record_dummy(const VkMat& buffer)

  2197.     {

  2198.         barrier_readwrite(buffer);

  2199.     }

  2200. 

  2201.     void record_dummy(const VkImageMat& image)

  2202.     {

  2203.         barrier_readwrite(image);

  2204.     }

  2205. 

  2206.     void record_dummy_readonly(const VkImageMat& image)

  2207.     {

  2208.         barrier_readonly(image);

  2209.     }

  2210. };

  2211. 

  2212. class VulkanDevicePrivate

  2213. {

  2214. public:

  2215.     VulkanDevicePrivate(VulkanDevice* _vkdev)

  2216.         : vkdev(_vkdev)

  2217.     {

  2218.     }

  2219.     VulkanDevice* const vkdev;

  2220. 

  2221.     // dummy buffer and image

  2222.     int create_dummy_buffer_image();

  2223.     void destroy_dummy_buffer_image();

  2224. 

  2225.     // utility operator

  2226.     const ncnn::Packing_vulkan* get_utility_operator(int storage_type_from, int storage_type_to, int cast_type_from_index, int cast_type_to_index, int packing_type_to_index) const;

  2227.     void destroy_utility_operator();

  2228. 

  2229.     VkDevice device;

  2230. 

  2231.     // hardware queue

  2232.     mutable std::vector<VkQueue> compute_queues;

  2233.     mutable std::vector<VkQueue> graphics_queues;

  2234.     mutable std::vector<VkQueue> transfer_queues;

  2235.     mutable int free_compute_queue_count;

  2236.     mutable int free_graphics_queue_count;

  2237.     mutable int free_transfer_queue_count;

  2238.     mutable Mutex compute_queue_lock;

  2239.     mutable Mutex graphics_queue_lock;

  2240.     mutable Mutex transfer_queue_lock;

  2241.     mutable ConditionVariable compute_queue_condition;

  2242.     mutable ConditionVariable graphics_queue_condition;

  2243.     mutable ConditionVariable transfer_queue_condition;

  2244. 

  2245.     // default blob allocator for each queue

  2246.     mutable std::vector<VkAllocator*> blob_allocators;

  2247.     mutable Mutex blob_allocator_lock;

  2248. 

  2249.     // default staging allocator for each queue

  2250.     mutable std::vector<VkAllocator*> staging_allocators;

  2251.     mutable Mutex staging_allocator_lock;

  2252. 

  2253.     // nearest sampler for texelfetch

  2254.     VkSampler texelfetch_sampler;

  2255. 

  2256.     // dummy buffer and image

  2257.     VkAllocator* dummy_allocator;

  2258.     VkMat dummy_buffer;

  2259.     VkImageMat dummy_image;

  2260.     VkImageMat dummy_image_readonly;

  2261. 

  2262.     // device-wide pipeline cache

  2263.     PipelineCache* pipeline_cache;

  2264. 

  2265.     // utility operator

  2266.     // from buffer | image

  2267.     // to buffer | image

  2268.     // from fp32-b/i | fp16p-b/i | fp16s-b/i

  2269.     // to fp32-b/i | fp16p-b/i | fp16s-b/i

  2270.     // to pack1 | pack4 | pack8

  2271.     mutable ncnn::Packing_vulkan* uop_packing[2][2][3][3][3];

  2272.     mutable Mutex uop_lock;

  2273. };

  2274. 

  2275. int VulkanDevicePrivate::create_dummy_buffer_image()

  2276. {

  2277.     dummy_allocator = new VkDummyAllocator(vkdev);

  2278. 

  2279.     dummy_buffer.create(1, 4u, dummy_allocator);

  2280.     dummy_image.create(1, 4u, dummy_allocator);

  2281. #if __APPLE__

  2282.     if (vkdev->info.type() == 0)

  2283.         dummy_image_readonly.create(1, 4u, dummy_allocator);

  2284. #else

  2285.     dummy_image_readonly.create(1, 4u, dummy_allocator);

  2286. #endif

  2287. 

  2288.     VkDummyCompute cmd(vkdev);

  2289. 

  2290.     cmd.record_dummy(dummy_buffer);

  2291.     cmd.record_dummy(dummy_image);

  2292. #if __APPLE__

  2293.     if (vkdev->info.type() == 0)

  2294.         cmd.record_dummy_readonly(dummy_image_readonly);

  2295. #else

  2296.     cmd.record_dummy_readonly(dummy_image_readonly);

  2297. #endif

  2298. 

  2299.     return cmd.submit_and_wait();

  2300. }

  2301. 

  2302. void VulkanDevicePrivate::destroy_dummy_buffer_image()

  2303. {

  2304.     dummy_buffer.release();

  2305.     dummy_image.release();

  2306. #if __APPLE__

  2307.     if (vkdev->info.type() == 0)

  2308.         dummy_image_readonly.release();

  2309. #else

  2310.     dummy_image_readonly.release();

  2311. #endif

  2312. 

  2313.     delete dummy_allocator;

  2314. }

  2315. 

  2316. const ncnn::Packing_vulkan* VulkanDevicePrivate::get_utility_operator(int storage_type_from, int storage_type_to, int cast_type_from_index, int cast_type_to_index, int packing_type_to_index) const

  2317. {

  2318.     MutexLockGuard lock(uop_lock);

  2319. 

  2320.     const ncnn::Packing_vulkan* cached_uop = uop_packing[storage_type_from][storage_type_to][cast_type_from_index][cast_type_to_index][packing_type_to_index];

  2321.     if (cached_uop)

  2322.         return cached_uop;

  2323. 

  2324.     if ((cast_type_from_index == 1 && cast_type_to_index == 2) || (cast_type_from_index == 2 && cast_type_to_index == 1))

  2325.     {

  2326.         NCNN_LOGE("no fp16p to/from fp16s conversion");

  2327.         return 0;

  2328.     }

  2329. 

  2330.     // create uop

  2331.     Option opt;

  2332.     opt.use_image_storage = (storage_type_from == 1 || storage_type_to == 1);

  2333.     opt.use_fp16_packed = (cast_type_from_index == 1 || cast_type_to_index == 1);

  2334.     opt.use_fp16_storage = (cast_type_from_index == 2 || cast_type_to_index == 2);

  2335. 

  2336.     if (!vkdev->info.support_fp16_packed() && opt.use_fp16_packed)

  2337.     {

  2338.         NCNN_LOGE("cannot create uop with use_fp16_packed if not support_fp16_packed");

  2339.         return 0;

  2340.     }

  2341. 

  2342.     if (!vkdev->info.support_fp16_storage() && opt.use_fp16_storage)

  2343.     {

  2344.         NCNN_LOGE("cannot create uop with use_fp16_storage if not support_fp16_storage");

  2345.         return 0;

  2346.     }

  2347. 

  2348.     // fp16/int8 arithmetic are not necessary for packing

  2349.     // and may conflict with storage options

  2350.     opt.use_fp16_arithmetic = false;

  2351.     opt.use_int8_arithmetic = false;

  2352. 

  2353.     // enable pack8 for pack8to1/pack8to4

  2354.     opt.use_shader_pack8 = true;

  2355. 

  2356.     // do not enable spirv-1.3 from cooperative matrix

  2357.     opt.use_cooperative_matrix = false;

  2358. 

  2359.     opt.use_vulkan_compute = true;

  2360. 

  2361.     // cache uop pipeline as device member explicitly

  2362.     opt.pipeline_cache = 0;

  2363. 

  2364.     ncnn::Packing_vulkan* uop = new ncnn::Packing_vulkan;

  2365.     uop->vkdev = vkdev;

  2366. 

  2367.     ncnn::ParamDict pd;

  2368.     pd.set(0, packing_type_to_index == 0 ? 1 : packing_type_to_index == 1 ? 4 : 8); // out_elempack

  2369.     pd.set(2, cast_type_from_index + 1);                                            // 0=auto 1=fp32 2=fp16p 3=fp16s

  2370.     pd.set(3, cast_type_to_index + 1);

  2371.     pd.set(4, storage_type_from); // 0=buffer 1=image

  2372.     pd.set(5, storage_type_to);

  2373. 

  2374.     uop->load_param(pd);

  2375. 

  2376.     uop->create_pipeline(opt);

  2377. 

  2378.     uop_packing[storage_type_from][storage_type_to][cast_type_from_index][cast_type_to_index][packing_type_to_index] = uop;

  2379. 

  2380.     return uop;

  2381. }

  2382. 

  2383. void VulkanDevicePrivate::destroy_utility_operator()

  2384. {

  2385.     Option opt;

  2386.     opt.use_vulkan_compute = true;

  2387.     opt.use_fp16_arithmetic = false;

  2388.     opt.use_int8_arithmetic = false;

  2389.     opt.use_cooperative_matrix = false;

  2390.     opt.pipeline_cache = 0;

  2391. 

  2392.     // from buffer | image

  2393.     // to buffer | image

  2394.     for (int i0 = 0; i0 < 2; i0++)

  2395.     {

  2396.         for (int i1 = 0; i1 < 2; i1++)

  2397.         {

  2398.             opt.use_image_storage = (i0 == 1 || i1 == 1);

  2399. 

  2400.             // from fp32-b/i | fp16p-b/i | fp16s-b/i

  2401.             // to fp32-b/i | fp16p-b/i | fp16s-b/i

  2402.             for (int j0 = 0; j0 < 3; j0++)

  2403.             {

  2404.                 for (int j1 = 0; j1 < 3; j1++)

  2405.                 {

  2406.                     if ((j0 == 1 && j1 == 2) || (j0 == 2 && j1 == 1))

  2407.                     {

  2408.                         // no fp16p to/from fp16s conversion

  2409.                         continue;

  2410.                     }

  2411. 

  2412.                     opt.use_fp16_packed = (j0 == 1 || j1 == 1);

  2413.                     opt.use_fp16_storage = (j0 == 2 || j1 == 2);

  2414. 

  2415.                     if (!vkdev->info.support_fp16_packed() && opt.use_fp16_packed)

  2416.                         continue;

  2417. 

  2418.                     if (!vkdev->info.support_fp16_storage() && opt.use_fp16_storage)

  2419.                         continue;

  2420. 

  2421.                     // to pack1 | pack4 | pack8

  2422.                     for (int k = 0; k < 3; k++)

  2423.                     {

  2424.                         // enable pack8 for pack8to1/pack8to4

  2425.                         opt.use_shader_pack8 = true;

  2426. 

  2427.                         ncnn::Layer* uop = uop_packing[i0][i1][j0][j1][k];

  2428.                         if (!uop)

  2429.                             continue;

  2430. 

  2431.                         uop->destroy_pipeline(opt);

  2432. 

  2433.                         delete uop;

  2434. 

  2435.                         uop_packing[i0][i1][j0][j1][k] = 0;

  2436.                     }

  2437.                 }

  2438.             }

  2439.         }

  2440.     }

  2441. }

  2442. 

  2443. VulkanDevice::VulkanDevice(int device_index)

  2444.     : info(get_gpu_info(device_index)), d(new VulkanDevicePrivate(this))

  2445. {

  2446.     try_create_gpu_instance();

  2447. 

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

  2449.     if (info.support_VK_KHR_8bit_storage())

  2450.         enabledExtensions.push_back("VK_KHR_8bit_storage");

  2451.     if (info.support_VK_KHR_16bit_storage())

  2452.         enabledExtensions.push_back("VK_KHR_16bit_storage");

  2453.     if (info.support_VK_KHR_bind_memory2())

  2454.         enabledExtensions.push_back("VK_KHR_bind_memory2");

  2455.     if (info.support_VK_KHR_buffer_device_address())

  2456.         enabledExtensions.push_back("VK_KHR_buffer_device_address");

  2457.     if (info.support_VK_KHR_create_renderpass2())

  2458.         enabledExtensions.push_back("VK_KHR_create_renderpass2");

  2459.     if (info.support_VK_KHR_cooperative_matrix())

  2460.         enabledExtensions.push_back("VK_KHR_cooperative_matrix");

  2461.     if (info.support_VK_KHR_dedicated_allocation())

  2462.         enabledExtensions.push_back("VK_KHR_dedicated_allocation");

  2463.     if (info.support_VK_KHR_descriptor_update_template())

  2464.         enabledExtensions.push_back("VK_KHR_descriptor_update_template");

  2465.     if (info.support_VK_KHR_external_memory())

  2466.         enabledExtensions.push_back("VK_KHR_external_memory");

  2467.     if (info.support_VK_KHR_get_memory_requirements2())

  2468.         enabledExtensions.push_back("VK_KHR_get_memory_requirements2");

  2469.     if (info.support_VK_KHR_maintenance1())

  2470.         enabledExtensions.push_back("VK_KHR_maintenance1");

  2471.     if (info.support_VK_KHR_maintenance2())

  2472.         enabledExtensions.push_back("VK_KHR_maintenance2");

  2473.     if (info.support_VK_KHR_maintenance3())

  2474.         enabledExtensions.push_back("VK_KHR_maintenance3");

  2475.     if (info.support_VK_KHR_multiview())

  2476.         enabledExtensions.push_back("VK_KHR_multiview");

  2477.     if (info.support_VK_KHR_portability_subset())

  2478.         enabledExtensions.push_back("VK_KHR_portability_subset");

  2479.     if (info.support_VK_KHR_push_descriptor())

  2480.         enabledExtensions.push_back("VK_KHR_push_descriptor");

  2481.     if (info.support_VK_KHR_sampler_ycbcr_conversion())

  2482.         enabledExtensions.push_back("VK_KHR_sampler_ycbcr_conversion");

  2483.     if (info.support_VK_KHR_shader_float16_int8())

  2484.         enabledExtensions.push_back("VK_KHR_shader_float16_int8");

  2485.     if (info.support_VK_KHR_shader_float_controls())

  2486.         enabledExtensions.push_back("VK_KHR_shader_float_controls");

  2487.     if (info.support_VK_KHR_storage_buffer_storage_class())

  2488.         enabledExtensions.push_back("VK_KHR_storage_buffer_storage_class");

  2489.     if (info.support_VK_KHR_swapchain())

  2490.         enabledExtensions.push_back("VK_KHR_swapchain");

  2491.     if (info.support_VK_EXT_buffer_device_address())

  2492.         enabledExtensions.push_back("VK_EXT_buffer_device_address");

  2493.     if (info.support_VK_EXT_descriptor_indexing())

  2494.         enabledExtensions.push_back("VK_EXT_descriptor_indexing");

  2495.     if (info.support_VK_EXT_memory_budget())

  2496.         enabledExtensions.push_back("VK_EXT_memory_budget");

  2497.     if (info.support_VK_EXT_memory_priority())

  2498.         enabledExtensions.push_back("VK_EXT_memory_priority");

  2499.     if (info.support_VK_EXT_queue_family_foreign())

  2500.         enabledExtensions.push_back("VK_EXT_queue_family_foreign");

  2501.     if (info.support_VK_AMD_device_coherent_memory())

  2502.         enabledExtensions.push_back("VK_AMD_device_coherent_memory");

  2503. #if __ANDROID_API__ >= 26

  2504.     if (info.support_VK_ANDROID_external_memory_android_hardware_buffer())

  2505.         enabledExtensions.push_back("VK_ANDROID_external_memory_android_hardware_buffer");

  2506. #endif // __ANDROID_API__ >= 26

  2507.     if (info.support_VK_NV_cooperative_matrix())

  2508.         enabledExtensions.push_back("VK_NV_cooperative_matrix");

  2509. 

  2510.     void* enabledExtensionFeatures = 0;

  2511. 

  2512.     // enable int8 storage

  2513.     VkPhysicalDevice8BitStorageFeaturesKHR enabled8BitStorageFeatures;

  2514.     enabled8BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_8BIT_STORAGE_FEATURES_KHR;

  2515.     enabled8BitStorageFeatures.pNext = 0;

  2516.     enabled8BitStorageFeatures.storageBuffer8BitAccess = info.support_int8_storage();

  2517.     enabled8BitStorageFeatures.uniformAndStorageBuffer8BitAccess = info.support_int8_uniform();

  2518.     enabled8BitStorageFeatures.storagePushConstant8 = VK_FALSE;

  2519.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_8bit_storage())

  2520.     {

  2521.         enabled8BitStorageFeatures.pNext = enabledExtensionFeatures;

  2522.         enabledExtensionFeatures = &enabled8BitStorageFeatures;

  2523.     }

  2524. 

  2525.     // enable fp16/int16 storage

  2526.     VkPhysicalDevice16BitStorageFeaturesKHR enabled16BitStorageFeatures;

  2527.     enabled16BitStorageFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR;

  2528.     enabled16BitStorageFeatures.pNext = 0;

  2529.     enabled16BitStorageFeatures.storageBuffer16BitAccess = info.support_fp16_storage();

  2530.     enabled16BitStorageFeatures.uniformAndStorageBuffer16BitAccess = info.support_fp16_uniform();

  2531.     enabled16BitStorageFeatures.storagePushConstant16 = VK_FALSE;

  2532.     enabled16BitStorageFeatures.storageInputOutput16 = VK_FALSE;

  2533.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_16bit_storage())

  2534.     {

  2535.         enabled16BitStorageFeatures.pNext = enabledExtensionFeatures;

  2536.         enabledExtensionFeatures = &enabled16BitStorageFeatures;

  2537.     }

  2538. 

  2539.     // enable fp16/int8 arithmetic

  2540.     VkPhysicalDeviceFloat16Int8FeaturesKHR enabledFloat16Int8Features;

  2541.     enabledFloat16Int8Features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR;

  2542.     enabledFloat16Int8Features.pNext = 0;

  2543.     enabledFloat16Int8Features.shaderFloat16 = info.support_fp16_arithmetic();

  2544.     enabledFloat16Int8Features.shaderInt8 = info.support_int8_arithmetic();

  2545.     if (support_VK_KHR_get_physical_device_properties2 && info.support_VK_KHR_shader_float16_int8())

  2546.     {

  2547.         enabledFloat16Int8Features.pNext = enabledExtensionFeatures;

  2548.         enabledExtensionFeatures = &enabledFloat16Int8Features;

  2549.     }

  2550. 

  2551.     // enable ycbcr conversion

  2552.     VkPhysicalDeviceSamplerYcbcrConversionFeaturesKHR querySamplerYcbcrConversionFeatures;

  2553.     querySamplerYcbcrConversionFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SAMPLER_YCBCR_CONVERSION_FEATURES_KHR;

  2554.     querySamplerYcbcrConversionFeatures.pNext = 0;

  2555.     querySamplerYcbcrConversionFeatures.samplerYcbcrConversion = info.support_ycbcr_conversion();

  2556.     if (support_VK_KHR_get_physical_device_properties2 && info.support_ycbcr_conversion())

  2557.     {

  2558.         querySamplerYcbcrConversionFeatures.pNext = enabledExtensionFeatures;

  2559.         enabledExtensionFeatures = &querySamplerYcbcrConversionFeatures;

  2560.     }

  2561. 

  2562.     // enable cooperative matrix

  2563.     VkPhysicalDeviceCooperativeMatrixFeaturesKHR queryCooperativeMatrixFeatures;

  2564.     queryCooperativeMatrixFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_FEATURES_KHR;

  2565.     queryCooperativeMatrixFeatures.pNext = 0;

  2566.     queryCooperativeMatrixFeatures.cooperativeMatrix = info.support_cooperative_matrix();

  2567.     queryCooperativeMatrixFeatures.cooperativeMatrixRobustBufferAccess = VK_FALSE;

  2568.     VkPhysicalDeviceCooperativeMatrixFeaturesNV queryCooperativeMatrixFeaturesNV;

  2569.     queryCooperativeMatrixFeaturesNV.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_FEATURES_NV;

  2570.     queryCooperativeMatrixFeaturesNV.pNext = 0;

  2571.     queryCooperativeMatrixFeaturesNV.cooperativeMatrix = info.support_cooperative_matrix();

  2572.     queryCooperativeMatrixFeaturesNV.cooperativeMatrixRobustBufferAccess = VK_FALSE;

  2573.     if (support_VK_KHR_get_physical_device_properties2 && info.support_cooperative_matrix())

  2574.     {

  2575.         if (info.support_VK_KHR_cooperative_matrix())

  2576.         {

  2577.             queryCooperativeMatrixFeatures.pNext = enabledExtensionFeatures;

  2578.             enabledExtensionFeatures = &queryCooperativeMatrixFeatures;

  2579.         }

  2580.         else

  2581.         {

  2582.             queryCooperativeMatrixFeaturesNV.pNext = enabledExtensionFeatures;

  2583.             enabledExtensionFeatures = &queryCooperativeMatrixFeaturesNV;

  2584.         }

  2585.     }

  2586. 

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

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

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

  2590. 

  2591.     VkDeviceQueueCreateInfo deviceQueueCreateInfos[3];

  2592. 

  2593.     VkDeviceQueueCreateInfo deviceComputeQueueCreateInfo;

  2594.     deviceComputeQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  2595.     deviceComputeQueueCreateInfo.pNext = 0;

  2596.     deviceComputeQueueCreateInfo.flags = 0;

  2597.     deviceComputeQueueCreateInfo.queueFamilyIndex = info.compute_queue_family_index();

  2598.     deviceComputeQueueCreateInfo.queueCount = info.compute_queue_count();

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

  2600. 

  2601.     VkDeviceQueueCreateInfo deviceGraphicsQueueCreateInfo;

  2602.     deviceGraphicsQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  2603.     deviceGraphicsQueueCreateInfo.pNext = 0;

  2604.     deviceGraphicsQueueCreateInfo.flags = 0;

  2605.     deviceGraphicsQueueCreateInfo.queueFamilyIndex = info.graphics_queue_family_index();

  2606.     deviceGraphicsQueueCreateInfo.queueCount = info.graphics_queue_count();

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

  2608. 

  2609.     VkDeviceQueueCreateInfo deviceTransferQueueCreateInfo;

  2610.     deviceTransferQueueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;

  2611.     deviceTransferQueueCreateInfo.pNext = 0;

  2612.     deviceTransferQueueCreateInfo.flags = 0;

  2613.     deviceTransferQueueCreateInfo.queueFamilyIndex = info.transfer_queue_family_index();

  2614.     deviceTransferQueueCreateInfo.queueCount = info.transfer_queue_count();

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

  2616. 

  2617.     VkDeviceCreateInfo deviceCreateInfo;

  2618.     deviceCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;

  2619.     deviceCreateInfo.pNext = enabledExtensionFeatures;

  2620.     deviceCreateInfo.flags = 0;

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

  2622.     {

  2623.         deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  2624.         deviceCreateInfo.queueCreateInfoCount = 1;

  2625.     }

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

  2627.     {

  2628.         deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  2629.         deviceQueueCreateInfos[1] = deviceTransferQueueCreateInfo;

  2630.         deviceCreateInfo.queueCreateInfoCount = 2;

  2631.     }

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

  2633.     {

  2634.         deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  2635.         deviceQueueCreateInfos[1] = deviceGraphicsQueueCreateInfo;

  2636.         deviceCreateInfo.queueCreateInfoCount = 2;

  2637.     }

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

  2639.     {

  2640.         deviceQueueCreateInfos[0] = deviceComputeQueueCreateInfo;

  2641.         deviceQueueCreateInfos[1] = deviceGraphicsQueueCreateInfo;

  2642.         deviceQueueCreateInfos[2] = deviceTransferQueueCreateInfo;

  2643.         deviceCreateInfo.queueCreateInfoCount = 3;

  2644.     }

  2645.     deviceCreateInfo.pQueueCreateInfos = deviceQueueCreateInfos;

  2646.     deviceCreateInfo.enabledLayerCount = 0;

  2647.     deviceCreateInfo.ppEnabledLayerNames = 0;

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

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

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

  2651. 

  2652.     VkResult ret = vkCreateDevice(info.physical_device(), &deviceCreateInfo, 0, &d->device);

  2653.     if (ret != VK_SUCCESS)

  2654.     {

  2655.         NCNN_LOGE("vkCreateDevice failed %d", ret);

  2656.     }

  2657. 

  2658.     init_device_extension();

  2659. 

  2660.     d->free_compute_queue_count = 0;

  2661.     d->free_graphics_queue_count = 0;

  2662.     d->free_transfer_queue_count = 0;

  2663. 

  2664.     d->free_compute_queue_count = info.compute_queue_count();

  2665.     d->compute_queues.resize(info.compute_queue_count());

  2666.     d->blob_allocators.resize(info.compute_queue_count());

  2667.     d->staging_allocators.resize(info.compute_queue_count());

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

  2669.     {

  2670.         vkGetDeviceQueue(d->device, info.compute_queue_family_index(), i, &d->compute_queues[i]);

  2671.         d->blob_allocators[i] = new VkBlobAllocator(this);

  2672.         d->staging_allocators[i] = new VkStagingAllocator(this);

  2673.     }

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

  2675.     {

  2676.         d->free_graphics_queue_count = info.graphics_queue_count();

  2677.         d->graphics_queues.resize(info.graphics_queue_count());

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

  2679.         {

  2680.             vkGetDeviceQueue(d->device, info.graphics_queue_family_index(), i, &d->graphics_queues[i]);

  2681.         }

  2682.     }

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

  2684.     {

  2685.         d->free_transfer_queue_count = info.transfer_queue_count();

  2686.         d->transfer_queues.resize(info.transfer_queue_count());

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

  2688.         {

  2689.             vkGetDeviceQueue(d->device, info.transfer_queue_family_index(), i, &d->transfer_queues[i]);

  2690.         }

  2691.     }

  2692. 

  2693.     // prepare immutable texelfetch sampler

  2694.     {

  2695.         VkSamplerCreateInfo samplerCreateInfo;

  2696.         samplerCreateInfo.sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO;

  2697.         samplerCreateInfo.pNext = 0;

  2698.         samplerCreateInfo.flags = 0;

  2699.         samplerCreateInfo.magFilter = VK_FILTER_NEAREST;

  2700.         samplerCreateInfo.minFilter = VK_FILTER_NEAREST;

  2701.         samplerCreateInfo.mipmapMode = VK_SAMPLER_MIPMAP_MODE_NEAREST;

  2702.         samplerCreateInfo.addressModeU = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE;

  2703.         samplerCreateInfo.addressModeV = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE;

  2704.         samplerCreateInfo.addressModeW = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE;

  2705.         samplerCreateInfo.mipLodBias = 0.0f;

  2706.         samplerCreateInfo.anisotropyEnable = VK_FALSE;

  2707.         samplerCreateInfo.maxAnisotropy = 1;

  2708.         samplerCreateInfo.compareEnable = VK_FALSE;

  2709.         samplerCreateInfo.compareOp = VK_COMPARE_OP_NEVER;

  2710.         samplerCreateInfo.minLod = 0.0f;

  2711.         samplerCreateInfo.maxLod = 0.0f;

  2712.         samplerCreateInfo.borderColor = VK_BORDER_COLOR_FLOAT_TRANSPARENT_BLACK;

  2713.         samplerCreateInfo.unnormalizedCoordinates = VK_TRUE;

  2714. 

  2715.         d->texelfetch_sampler = 0;

  2716.         ret = vkCreateSampler(d->device, &samplerCreateInfo, 0, &d->texelfetch_sampler);

  2717.         if (ret != VK_SUCCESS)

  2718.         {

  2719.             NCNN_LOGE("vkCreateSampler failed %d", ret);

  2720.         }

  2721.     }

  2722. 

  2723.     int cret = d->create_dummy_buffer_image();

  2724.     if (cret != 0)

  2725.     {

  2726.         NCNN_LOGE("VulkanDevice create_dummy_buffer_image failed %d", cret);

  2727.     }

  2728. 

  2729.     d->pipeline_cache = new PipelineCache(this);

  2730. 

  2731.     memset(d->uop_packing, 0, sizeof(d->uop_packing));

  2732. }

  2733. 

  2734. VulkanDevice::~VulkanDevice()

  2735. {

  2736.     d->destroy_utility_operator();

  2737. 

  2738.     d->destroy_dummy_buffer_image();

  2739. 

  2740.     if (d->texelfetch_sampler)

  2741.     {

  2742.         vkDestroySampler(d->device, d->texelfetch_sampler, 0);

  2743.     }

  2744. 

  2745.     for (size_t i = 0; i < d->blob_allocators.size(); i++)

  2746.     {

  2747.         delete d->blob_allocators[i];

  2748.     }

  2749.     d->blob_allocators.clear();

  2750.     for (size_t i = 0; i < d->staging_allocators.size(); i++)

  2751.     {

  2752.         delete d->staging_allocators[i];

  2753.     }

  2754.     d->staging_allocators.clear();

  2755. 

  2756.     delete d->pipeline_cache;

  2757. 

  2758.     vkDestroyDevice(d->device, 0);

  2759. 

  2760.     delete d;

  2761. }

  2762. 

  2763. VulkanDevice::VulkanDevice(const VulkanDevice&)

  2764.     : info(get_gpu_info(0)), d(0)

  2765. {

  2766. }

  2767. 

  2768. VulkanDevice& VulkanDevice::operator=(const VulkanDevice&)

  2769. {

  2770.     return *this;

  2771. }

  2772. 

  2773. VkDevice VulkanDevice::vkdevice() const

  2774. {

  2775.     return d->device;

  2776. }

  2777. 

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

  2779. {

  2780.     VkShaderModuleCreateInfo shaderModuleCreateInfo;

  2781.     shaderModuleCreateInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;

  2782.     shaderModuleCreateInfo.pNext = 0;

  2783.     shaderModuleCreateInfo.flags = 0;

  2784.     shaderModuleCreateInfo.codeSize = spv_data_size;

  2785.     shaderModuleCreateInfo.pCode = spv_data;

  2786. 

  2787.     VkShaderModule shader_module;

  2788.     VkResult ret = vkCreateShaderModule(d->device, &shaderModuleCreateInfo, 0, &shader_module);

  2789.     if (ret != VK_SUCCESS)

  2790.     {

  2791.         NCNN_LOGE("vkCreateShaderModule failed %d", ret);

  2792.         return 0;

  2793.     }

  2794. 

  2795.     return shader_module;

  2796. }

  2797. 

  2798. 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)

  2799. {

  2800.     uint32_t local_size_x_id = -1;

  2801.     uint32_t local_size_y_id = -1;

  2802.     uint32_t local_size_z_id = -1;

  2803.     uint32_t gl_WorkGroupSize_id = -1;

  2804. 

  2805.     const uint32_t* p = code;

  2806.     uint32_t* dp = dstcode;

  2807. 

  2808.     // skip magic version generator bound schema

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

  2810.     p += 5;

  2811.     dp += 5;

  2812. 

  2813.     // foreach op

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

  2815.     {

  2816.         uint32_t opcode = p[0];

  2817. 

  2818.         uint16_t wordcount = opcode >> 16;

  2819.         uint16_t op = opcode & 0xffff;

  2820. 

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

  2822.         {

  2823.             uint32_t mode = p[2];

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

  2825.             {

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

  2827. 

  2828.                 // set local_size_xyz

  2829.                 dp[3] = local_size_x;

  2830.                 dp[4] = local_size_y;

  2831.                 dp[5] = local_size_z;

  2832. 

  2833.                 p += wordcount;

  2834.                 dp += wordcount;

  2835.                 continue;

  2836.             }

  2837.         }

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

  2839.         {

  2840.             uint32_t id = p[2];

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

  2842.             {

  2843.                 p += wordcount;

  2844.                 continue;

  2845.             }

  2846.         }

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

  2848.         {

  2849.             uint32_t id = p[2];

  2850.             if (id == gl_WorkGroupSize_id)

  2851.             {

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

  2853.                 {

  2854.                     p += wordcount;

  2855.                     continue;

  2856.                 }

  2857.             }

  2858.         }

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

  2860.         {

  2861.             uint32_t id = p[1];

  2862.             uint32_t decoration = p[2];

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

  2864.             {

  2865.                 uint32_t specid = p[3];

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

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

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

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

  2870.                 {

  2871.                     p += wordcount;

  2872.                     continue;

  2873.                 }

  2874.             }

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

  2876.             {

  2877.                 uint32_t builtin = p[3];

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

  2879.                 {

  2880.                     gl_WorkGroupSize_id = id;

  2881.                     p += wordcount;

  2882.                     continue;

  2883.                 }

  2884.             }

  2885.         }

  2886. 

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

  2888.         p += wordcount;

  2889.         dp += wordcount;

  2890.     }

  2891. 

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

  2893. }

  2894. 

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

  2896. {

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

  2898.     size_t spv_data_size_modified = spv_data_size;

  2899.     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);

  2900. 

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

  2902. 

  2903.     free(spv_data_modified);

  2904. 

  2905.     return shader_module;

  2906. }

  2907. 

  2908. int VulkanDevice::create_descriptorset_layout(int binding_count, const int* binding_types, VkDescriptorSetLayout* descriptorset_layout) const

  2909. {

  2910.     if (binding_count == 0)

  2911.     {

  2912.         *descriptorset_layout = 0;

  2913.         return 0;

  2914.     }

  2915. 

  2916.     std::vector<VkDescriptorSetLayoutBinding> descriptorSetLayoutBindings(binding_count);

  2917.     for (int i = 0; i < binding_count; i++)

  2918.     {

  2919.         int binding_type = binding_types[i];

  2920. 

  2921.         descriptorSetLayoutBindings[i].binding = i;

  2922.         descriptorSetLayoutBindings[i].descriptorCount = 1;

  2923.         descriptorSetLayoutBindings[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;

  2924. 

  2925.         if (binding_type == 1)

  2926.         {

  2927.             descriptorSetLayoutBindings[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;

  2928.             descriptorSetLayoutBindings[i].pImmutableSamplers = 0;

  2929.         }

  2930.         else if (binding_type == 2)

  2931.         {

  2932.             descriptorSetLayoutBindings[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;

  2933.             descriptorSetLayoutBindings[i].pImmutableSamplers = 0;

  2934.         }

  2935.         else // if (binding_type == 3)

  2936.         {

  2937.             descriptorSetLayoutBindings[i].descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER;

  2938.             descriptorSetLayoutBindings[i].pImmutableSamplers = immutable_texelfetch_sampler(); // we always use texelfetch

  2939.         }

  2940.     }

  2941. 

  2942.     VkDescriptorSetLayoutCreateInfo descriptorSetLayoutCreateInfo;

  2943.     descriptorSetLayoutCreateInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;

  2944.     descriptorSetLayoutCreateInfo.pNext = 0;

  2945.     descriptorSetLayoutCreateInfo.flags = 0;

  2946.     descriptorSetLayoutCreateInfo.bindingCount = binding_count;

  2947.     descriptorSetLayoutCreateInfo.pBindings = descriptorSetLayoutBindings.data();

  2948. 

  2949.     if (info.support_VK_KHR_push_descriptor())

  2950.     {

  2951.         descriptorSetLayoutCreateInfo.flags |= VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR;

  2952.     }

  2953. 

  2954.     VkResult ret = vkCreateDescriptorSetLayout(d->device, &descriptorSetLayoutCreateInfo, 0, descriptorset_layout);

  2955.     if (ret != VK_SUCCESS)

  2956.     {

  2957.         NCNN_LOGE("vkCreateDescriptorSetLayout failed %d", ret);

  2958.         return -1;

  2959.     }

  2960. 

  2961.     return 0;

  2962. }

  2963. 

  2964. int VulkanDevice::create_pipeline_layout(int push_constant_count, VkDescriptorSetLayout descriptorset_layout, VkPipelineLayout* pipeline_layout) const

  2965. {

  2966.     VkPushConstantRange pushConstantRange;

  2967.     pushConstantRange.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;

  2968.     pushConstantRange.offset = 0;

  2969.     pushConstantRange.size = sizeof(vk_constant_type) * push_constant_count;

  2970. 

  2971.     VkPipelineLayoutCreateInfo pipelineLayoutCreateInfo;

  2972.     pipelineLayoutCreateInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;

  2973.     pipelineLayoutCreateInfo.pNext = 0;

  2974.     pipelineLayoutCreateInfo.flags = 0;

  2975. 

  2976.     if (descriptorset_layout)

  2977.     {

  2978.         pipelineLayoutCreateInfo.setLayoutCount = 1;

  2979.         pipelineLayoutCreateInfo.pSetLayouts = &descriptorset_layout;

  2980.     }

  2981.     else

  2982.     {

  2983.         pipelineLayoutCreateInfo.setLayoutCount = 0;

  2984.         pipelineLayoutCreateInfo.pSetLayouts = 0;

  2985.     }

  2986. 

  2987.     if (push_constant_count > 0)

  2988.     {

  2989.         pipelineLayoutCreateInfo.pushConstantRangeCount = 1;

  2990.         pipelineLayoutCreateInfo.pPushConstantRanges = &pushConstantRange;

  2991.     }

  2992.     else

  2993.     {

  2994.         pipelineLayoutCreateInfo.pushConstantRangeCount = 0;

  2995.         pipelineLayoutCreateInfo.pPushConstantRanges = 0;

  2996.     }

  2997. 

  2998.     VkResult ret = vkCreatePipelineLayout(d->device, &pipelineLayoutCreateInfo, 0, pipeline_layout);

  2999.     if (ret != VK_SUCCESS)

  3000.     {

  3001.         NCNN_LOGE("vkCreatePipelineLayout failed %d", ret);

  3002.         return -1;

  3003.     }

  3004. 

  3005.     return 0;

  3006. }

  3007. 

  3008. int VulkanDevice::create_pipeline(VkShaderModule shader_module, VkPipelineLayout pipeline_layout, const std::vector<vk_specialization_type>& specializations, VkPipeline* pipeline) const

  3009. {

  3010.     const int specialization_count = specializations.size();

  3011. 

  3012.     std::vector<VkSpecializationMapEntry> specializationMapEntries(specialization_count);

  3013.     for (int i = 0; i < specialization_count; i++)

  3014.     {

  3015.         specializationMapEntries[i].constantID = i;

  3016.         specializationMapEntries[i].offset = i * sizeof(vk_specialization_type);

  3017.         specializationMapEntries[i].size = sizeof(vk_specialization_type);

  3018.     }

  3019. 

  3020.     VkSpecializationInfo specializationInfo;

  3021.     specializationInfo.mapEntryCount = specializationMapEntries.size();

  3022.     specializationInfo.pMapEntries = specializationMapEntries.data();

  3023.     specializationInfo.dataSize = specializations.size() * sizeof(vk_specialization_type);

  3024.     specializationInfo.pData = specializations.data();

  3025. 

  3026.     VkPipelineShaderStageCreateInfo pipelineShaderStageCreateInfo;

  3027.     pipelineShaderStageCreateInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;

  3028.     pipelineShaderStageCreateInfo.pNext = 0;

  3029.     pipelineShaderStageCreateInfo.flags = 0;

  3030.     pipelineShaderStageCreateInfo.stage = VK_SHADER_STAGE_COMPUTE_BIT;

  3031.     pipelineShaderStageCreateInfo.module = shader_module;

  3032.     pipelineShaderStageCreateInfo.pName = "main";

  3033.     pipelineShaderStageCreateInfo.pSpecializationInfo = &specializationInfo;

  3034. 

  3035.     VkComputePipelineCreateInfo computePipelineCreateInfo;

  3036.     computePipelineCreateInfo.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;

  3037.     computePipelineCreateInfo.pNext = 0;

  3038.     computePipelineCreateInfo.flags = 0;

  3039.     computePipelineCreateInfo.stage = pipelineShaderStageCreateInfo;

  3040.     computePipelineCreateInfo.layout = pipeline_layout;

  3041.     computePipelineCreateInfo.basePipelineHandle = 0;

  3042.     computePipelineCreateInfo.basePipelineIndex = 0;

  3043. 

  3044.     VkResult ret = vkCreateComputePipelines(d->device, 0, 1, &computePipelineCreateInfo, 0, pipeline);

  3045.     if (ret != VK_SUCCESS)

  3046.     {

  3047.         NCNN_LOGE("vkCreateComputePipelines failed %d", ret);

  3048.         return -1;

  3049.     }

  3050. 

  3051.     return 0;

  3052. }

  3053. 

  3054. int VulkanDevice::create_descriptor_update_template(int binding_count, const int* binding_types, VkDescriptorSetLayout descriptorset_layout, VkPipelineLayout pipeline_layout, VkDescriptorUpdateTemplateKHR* descriptor_update_template) const

  3055. {

  3056.     if (binding_count == 0)

  3057.     {

  3058.         *descriptor_update_template = 0;

  3059.         return 0;

  3060.     }

  3061. 

  3062.     std::vector<VkDescriptorUpdateTemplateEntryKHR> descriptorUpdateTemplateEntries(binding_count);

  3063.     size_t offset = 0;

  3064.     for (int i = 0; i < binding_count; i++) // TODO do not update weights

  3065.     {

  3066.         int binding_type = binding_types[i];

  3067. 

  3068.         descriptorUpdateTemplateEntries[i].dstBinding = i;

  3069.         descriptorUpdateTemplateEntries[i].dstArrayElement = 0;

  3070.         descriptorUpdateTemplateEntries[i].descriptorCount = 1;

  3071.         descriptorUpdateTemplateEntries[i].offset = offset;

  3072. 

  3073.         if (binding_type == 1)

  3074.         {

  3075.             descriptorUpdateTemplateEntries[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;

  3076.             descriptorUpdateTemplateEntries[i].stride = sizeof(VkDescriptorBufferInfo);

  3077.         }

  3078.         else if (binding_type == 2)

  3079.         {

  3080.             descriptorUpdateTemplateEntries[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;

  3081.             descriptorUpdateTemplateEntries[i].stride = sizeof(VkDescriptorImageInfo);

  3082.         }

  3083.         else // if (binding_type == 3)

  3084.         {

  3085.             descriptorUpdateTemplateEntries[i].descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER;

  3086.             descriptorUpdateTemplateEntries[i].stride = sizeof(VkDescriptorImageInfo);

  3087.         }

  3088. 

  3089.         offset += descriptorUpdateTemplateEntries[i].stride;

  3090.     }

  3091. 

  3092.     VkDescriptorUpdateTemplateCreateInfoKHR descriptorUpdateTemplateCreateInfo;

  3093.     descriptorUpdateTemplateCreateInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_UPDATE_TEMPLATE_CREATE_INFO_KHR;

  3094.     descriptorUpdateTemplateCreateInfo.pNext = 0;

  3095.     descriptorUpdateTemplateCreateInfo.flags = 0;

  3096.     descriptorUpdateTemplateCreateInfo.descriptorUpdateEntryCount = binding_count; // TODO do not update weights

  3097.     descriptorUpdateTemplateCreateInfo.pDescriptorUpdateEntries = descriptorUpdateTemplateEntries.data();

  3098.     if (info.support_VK_KHR_push_descriptor())

  3099.     {

  3100.         descriptorUpdateTemplateCreateInfo.templateType = VK_DESCRIPTOR_UPDATE_TEMPLATE_TYPE_PUSH_DESCRIPTORS_KHR;

  3101.     }

  3102.     else

  3103.     {

  3104.         descriptorUpdateTemplateCreateInfo.templateType = VK_DESCRIPTOR_UPDATE_TEMPLATE_TYPE_DESCRIPTOR_SET_KHR;

  3105.     }

  3106.     // descriptorSetLayout should be ignored if VK_DESCRIPTOR_UPDATE_TEMPLATE_TYPE_PUSH_DESCRIPTORS_KHR

  3107.     // FIXME HACK WARNING TODO NOTE but crash on radv if set NULL  :(

  3108.     descriptorUpdateTemplateCreateInfo.descriptorSetLayout = descriptorset_layout;

  3109.     descriptorUpdateTemplateCreateInfo.pipelineBindPoint = VK_PIPELINE_BIND_POINT_COMPUTE;

  3110.     descriptorUpdateTemplateCreateInfo.pipelineLayout = pipeline_layout;

  3111.     descriptorUpdateTemplateCreateInfo.set = 0;

  3112. 

  3113.     VkResult ret = vkCreateDescriptorUpdateTemplateKHR(d->device, &descriptorUpdateTemplateCreateInfo, 0, descriptor_update_template);

  3114.     if (ret != VK_SUCCESS)

  3115.     {

  3116.         NCNN_LOGE("vkCreateDescriptorUpdateTemplateKHR failed %d", ret);

  3117.         return -1;

  3118.     }

  3119. 

  3120.     return 0;

  3121. }

  3122. 

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

  3124. {

  3125.     const VkPhysicalDeviceMemoryProperties& memory_properties = info.physical_device_memory_properties();

  3126. 

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

  3128.     for (uint32_t i = 0; i < memory_properties.memoryTypeCount; i++)

  3129.     {

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

  3131.         if (is_required)

  3132.         {

  3133.             const VkMemoryType& memoryType = memory_properties.memoryTypes[i];

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

  3135.                     && (preferred && (memoryType.propertyFlags & preferred))

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

  3137.             {

  3138.                 return i;

  3139.             }

  3140.         }

  3141.     }

  3142. 

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

  3144.     for (uint32_t i = 0; i < memory_properties.memoryTypeCount; i++)

  3145.     {

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

  3147.         if (is_required)

  3148.         {

  3149.             const VkMemoryType& memoryType = memory_properties.memoryTypes[i];

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

  3151.                     && (preferred && (memoryType.propertyFlags & preferred)))

  3152.             {

  3153.                 return i;

  3154.             }

  3155.         }

  3156.     }

  3157. 

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

  3159.     for (uint32_t i = 0; i < memory_properties.memoryTypeCount; i++)

  3160.     {

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

  3162.         if (is_required)

  3163.         {

  3164.             const VkMemoryType& memoryType = memory_properties.memoryTypes[i];

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

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

  3167.             {

  3168.                 return i;

  3169.             }

  3170.         }

  3171.     }

  3172. 

  3173.     // fourth try, find any required

  3174.     for (uint32_t i = 0; i < memory_properties.memoryTypeCount; i++)

  3175.     {

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

  3177.         if (is_required)

  3178.         {

  3179.             const VkMemoryType& memoryType = memory_properties.memoryTypes[i];

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

  3181.             {

  3182.                 return i;

  3183.             }

  3184.         }

  3185.     }

  3186. 

  3187.     NCNN_LOGE("no such memory type %u %u %u %u", memory_type_bits, required, preferred, preferred_not);

  3188.     return -1;

  3189. }

  3190. 

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

  3192. {

  3193.     const VkMemoryType& memoryType = info.physical_device_memory_properties().memoryTypes[memory_type_index];

  3194. 

  3195.     return memoryType.propertyFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT;

  3196. }

  3197. 

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

  3199. {

  3200.     const VkMemoryType& memoryType = info.physical_device_memory_properties().memoryTypes[memory_type_index];

  3201. 

  3202.     return memoryType.propertyFlags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT;

  3203. }

  3204. 

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

  3206. {

  3207.     if (queue_family_index != info.compute_queue_family_index()

  3208.             && queue_family_index != info.graphics_queue_family_index()

  3209.             && queue_family_index != info.transfer_queue_family_index())

  3210.     {

  3211.         NCNN_LOGE("invalid queue_family_index %u", queue_family_index);

  3212.         return 0;

  3213.     }

  3214. 

  3215.     Mutex& queue_lock = queue_family_index == info.compute_queue_family_index() ? d->compute_queue_lock

  3216.                         : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queue_lock

  3217.                         : d->transfer_queue_lock;

  3218. 

  3219.     queue_lock.lock();

  3220. 

  3221.     ConditionVariable& queue_condition = queue_family_index == info.compute_queue_family_index() ? d->compute_queue_condition

  3222.                                          : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queue_condition

  3223.                                          : d->transfer_queue_condition;

  3224. 

  3225.     int& free_queue_count = queue_family_index == info.compute_queue_family_index() ? d->free_compute_queue_count

  3226.                             : queue_family_index == info.graphics_queue_family_index() ? d->free_graphics_queue_count

  3227.                             : d->free_transfer_queue_count;

  3228. 

  3229.     while (free_queue_count == 0)

  3230.     {

  3231.         // no free queues, wait for recleams from other threads

  3232.         queue_condition.wait(queue_lock);

  3233.     }

  3234. 

  3235.     std::vector<VkQueue>& queues = queue_family_index == info.compute_queue_family_index() ? d->compute_queues

  3236.                                    : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queues

  3237.                                    : d->transfer_queues;

  3238. 

  3239.     VkQueue queue = 0;

  3240.     for (size_t i = 0; i < queues.size(); i++)

  3241.     {

  3242.         if (queues[i])

  3243.         {

  3244.             queue = queues[i];

  3245.             queues[i] = 0;

  3246.             break;

  3247.         }

  3248.     }

  3249. 

  3250.     if (!queue)

  3251.     {

  3252.         NCNN_LOGE("FATAL ERROR! out of hardware queue %u", queue_family_index);

  3253.     }

  3254. 

  3255.     free_queue_count -= 1;

  3256. 

  3257.     queue_lock.unlock();

  3258. 

  3259.     queue_condition.signal();

  3260. 

  3261.     return queue;

  3262. }

  3263. 

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

  3265. {

  3266.     if (queue_family_index != info.compute_queue_family_index()

  3267.             && queue_family_index != info.graphics_queue_family_index()

  3268.             && queue_family_index != info.transfer_queue_family_index())

  3269.     {

  3270.         NCNN_LOGE("invalid queue_family_index %u", queue_family_index);

  3271.         return;

  3272.     }

  3273. 

  3274.     Mutex& queue_lock = queue_family_index == info.compute_queue_family_index() ? d->compute_queue_lock

  3275.                         : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queue_lock

  3276.                         : d->transfer_queue_lock;

  3277. 

  3278.     queue_lock.lock();

  3279. 

  3280.     ConditionVariable& queue_condition = queue_family_index == info.compute_queue_family_index() ? d->compute_queue_condition

  3281.                                          : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queue_condition

  3282.                                          : d->transfer_queue_condition;

  3283. 

  3284.     int& free_queue_count = queue_family_index == info.compute_queue_family_index() ? d->free_compute_queue_count

  3285.                             : queue_family_index == info.graphics_queue_family_index() ? d->free_graphics_queue_count

  3286.                             : d->free_transfer_queue_count;

  3287. 

  3288.     std::vector<VkQueue>& queues = queue_family_index == info.compute_queue_family_index() ? d->compute_queues

  3289.                                    : queue_family_index == info.graphics_queue_family_index() ? d->graphics_queues

  3290.                                    : d->transfer_queues;

  3291. 

  3292.     size_t i = 0;

  3293.     for (; i < queues.size(); i++)

  3294.     {

  3295.         if (!queues[i])

  3296.         {

  3297.             queues[i] = queue;

  3298.             break;

  3299.         }

  3300.     }

  3301. 

  3302.     if (i == queues.size())

  3303.     {

  3304.         NCNN_LOGE("FATAL ERROR! reclaim_queue get wild queue %u %p", queue_family_index, queue);

  3305.     }

  3306. 

  3307.     free_queue_count += 1;

  3308. 

  3309.     queue_lock.unlock();

  3310. 

  3311.     queue_condition.signal();

  3312. }

  3313. 

  3314. VkAllocator* VulkanDevice::acquire_blob_allocator() const

  3315. {

  3316.     MutexLockGuard lock(d->blob_allocator_lock);

  3317. 

  3318.     for (int i = 0; i < (int)d->blob_allocators.size(); i++)

  3319.     {

  3320.         VkAllocator* allocator = d->blob_allocators[i];

  3321.         if (allocator)

  3322.         {

  3323.             d->blob_allocators[i] = 0;

  3324.             return allocator;

  3325.         }

  3326.     }

  3327. 

  3328.     // pre-allocated allcator exhausted, create new

  3329.     VkAllocator* allocator = new VkBlobAllocator(this);

  3330.     d->blob_allocators.push_back(allocator);

  3331.     d->blob_allocators[d->blob_allocators.size() - 1] = 0;

  3332.     return allocator;

  3333. }

  3334. 

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

  3336. {

  3337.     MutexLockGuard lock(d->blob_allocator_lock);

  3338. 

  3339.     for (int i = 0; i < (int)d->blob_allocators.size(); i++)

  3340.     {

  3341.         if (!d->blob_allocators[i])

  3342.         {

  3343.             d->blob_allocators[i] = allocator;

  3344.             return;

  3345.         }

  3346.     }

  3347. 

  3348.     NCNN_LOGE("FATAL ERROR! reclaim_blob_allocator get wild allocator %p", allocator);

  3349. }

  3350. 

  3351. VkAllocator* VulkanDevice::acquire_staging_allocator() const

  3352. {

  3353.     MutexLockGuard lock(d->staging_allocator_lock);

  3354. 

  3355.     for (int i = 0; i < (int)d->staging_allocators.size(); i++)

  3356.     {

  3357.         VkAllocator* allocator = d->staging_allocators[i];

  3358.         if (allocator)

  3359.         {

  3360.             d->staging_allocators[i] = 0;

  3361.             return allocator;

  3362.         }

  3363.     }

  3364. 

  3365.     // pre-allocated allcator exhausted, create new

  3366.     VkAllocator* allocator = new VkStagingAllocator(this);

  3367.     d->staging_allocators.push_back(allocator);

  3368.     d->staging_allocators[d->staging_allocators.size() - 1] = 0;

  3369.     return allocator;

  3370. }

  3371. 

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

  3373. {

  3374.     MutexLockGuard lock(d->staging_allocator_lock);

  3375. 

  3376.     for (int i = 0; i < (int)d->staging_allocators.size(); i++)

  3377.     {

  3378.         if (!d->staging_allocators[i])

  3379.         {

  3380.             d->staging_allocators[i] = allocator;

  3381.             return;

  3382.         }

  3383.     }

  3384. 

  3385.     NCNN_LOGE("FATAL ERROR! reclaim_staging_allocator get wild allocator %p", allocator);

  3386. }

  3387. 

  3388. const VkSampler* VulkanDevice::immutable_texelfetch_sampler() const

  3389. {

  3390.     return &d->texelfetch_sampler;

  3391. }

  3392. 

  3393. VkMat VulkanDevice::get_dummy_buffer() const

  3394. {

  3395.     return d->dummy_buffer;

  3396. }

  3397. 

  3398. VkImageMat VulkanDevice::get_dummy_image() const

  3399. {

  3400.     return d->dummy_image;

  3401. }

  3402. 

  3403. VkImageMat VulkanDevice::get_dummy_image_readonly() const

  3404. {

  3405. #if __APPLE__

  3406.     if (info.type() != 0)

  3407.         return d->dummy_image;

  3408. #endif

  3409.     return d->dummy_image_readonly;

  3410. }

  3411. 

  3412. const PipelineCache* VulkanDevice::get_pipeline_cache() const

  3413. {

  3414.     return d->pipeline_cache;

  3415. }

  3416. 

  3417. bool VulkanDevice::shape_support_image_storage(const Mat& shape) const

  3418. {

  3419.     int dims = shape.dims;

  3420.     int width = shape.w;

  3421.     int height = shape.h;

  3422.     int depth = shape.c;

  3423.     int elempack = shape.elempack;

  3424. 

  3425.     // large elempack spills on image w

  3426.     if (elempack == 8) width *= 2;

  3427.     if (elempack == 16) width *= 4;

  3428.     if (elempack == 32) width *= 8;

  3429.     if (elempack == 64) width *= 16;

  3430. 

  3431.     if (dims == 1)

  3432.     {

  3433.         if (width > (int)info.max_image_dimension_1d())

  3434.         {

  3435.             return false;

  3436.         }

  3437.     }

  3438.     else if (dims == 2)

  3439.     {

  3440.         if (width > (int)info.max_image_dimension_2d() || height > (int)info.max_image_dimension_2d())

  3441.         {

  3442.             return false;

  3443.         }

  3444.     }

  3445.     else // if (dims == 3)

  3446.     {

  3447.         if (width > (int)info.max_image_dimension_3d() || height > (int)info.max_image_dimension_3d() || depth > (int)info.max_image_dimension_3d())

  3448.         {

  3449.             return false;

  3450.         }

  3451.     }

  3452. 

  3453.     return true;

  3454. }

  3455. 

  3456. uint32_t VulkanDevice::get_heap_budget() const

  3457. {

  3458.     const VkPhysicalDeviceMemoryProperties& memory_properties = info.physical_device_memory_properties();

  3459. 

  3460.     uint32_t buffer_memory_type_index = d->dummy_allocator->buffer_memory_type_index;

  3461.     uint32_t buffer_heap_index = memory_properties.memoryTypes[buffer_memory_type_index].heapIndex;

  3462. 

  3463.     if (!info.support_VK_EXT_memory_budget())

  3464.     {

  3465.         //         NCNN_LOGE("heap budget from assumption\n");

  3466.         uint32_t device_local_heap_size = memory_properties.memoryHeaps[buffer_heap_index].size / 1024 / 1024;

  3467. 

  3468.         // we usually cannot use all heap

  3469.         // 70% for 4G+

  3470.         // 50% for 4G-

  3471.         return device_local_heap_size >= 4000 ? device_local_heap_size * 0.7 : device_local_heap_size * 0.5;

  3472.     }

  3473. 

  3474.     VkPhysicalDeviceMemoryBudgetPropertiesEXT memoryBudgetProperties;

  3475.     memoryBudgetProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT;

  3476.     memoryBudgetProperties.pNext = 0;

  3477. 

  3478.     VkPhysicalDeviceMemoryProperties2KHR memoryProperties;

  3479.     memoryProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_PROPERTIES_2_KHR;

  3480.     memoryProperties.pNext = &memoryBudgetProperties;

  3481. 

  3482.     vkGetPhysicalDeviceMemoryProperties2KHR(info.physical_device(), &memoryProperties);

  3483. 

  3484.     return memoryBudgetProperties.heapBudget[buffer_heap_index] / 1024 / 1024;

  3485. }

  3486. 

  3487. void VulkanDevice::convert_packing(const VkMat& src, VkMat& dst, int dst_elempack, VkCompute& cmd, const Option& _opt) const

  3488. {

  3489.     // buffer2buffer uop is created with use_image_storage disabled

  3490.     Option opt = _opt;

  3491.     opt.use_image_storage = false;

  3492. 

  3493.     int cast_type_to_index = opt.use_fp16_storage ? 2 : opt.use_fp16_packed ? 1 : 0;

  3494.     int packing_type_to_index = dst_elempack == 1 ? 0 : dst_elempack == 4 ? 1 : 2;

  3495. 

  3496.     int cast_type_from_index;

  3497.     if (src.elembits() == 32)

  3498.     {

  3499.         cast_type_from_index = 0;

  3500.     }

  3501.     else // if (src.elembits() == 16)

  3502.     {

  3503.         if (cast_type_to_index != 0)

  3504.         {

  3505.             cast_type_from_index = cast_type_to_index;

  3506.         }

  3507.         else if (info.support_fp16_storage())

  3508.         {

  3509.             cast_type_from_index = 2;

  3510.         }

  3511.         else // if (info.support_fp16_packed())

  3512.         {

  3513.             cast_type_from_index = 1;

  3514.         }

  3515.     }

  3516. 

  3517.     // NCNN_LOGE("convert_packing b2b %d %d %d", cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3518. 

  3519.     const ncnn::Packing_vulkan* uop = d->get_utility_operator(0, 0, cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3520.     uop->forward(src, dst, cmd, opt);

  3521. }

  3522. 

  3523. void VulkanDevice::convert_packing(const VkImageMat& src, VkImageMat& dst, int dst_elempack, VkCompute& cmd, const Option& opt) const

  3524. {

  3525.     int cast_type_to_index = opt.use_fp16_storage ? 2 : opt.use_fp16_packed ? 1 : 0;

  3526.     int packing_type_to_index = dst_elempack == 1 ? 0 : dst_elempack == 4 ? 1 : 2;

  3527. 

  3528.     int cast_type_from_index;

  3529.     if (src.elembits() == 32)

  3530.     {

  3531.         cast_type_from_index = 0;

  3532.     }

  3533.     else // if (src.elembits() == 16)

  3534.     {

  3535.         if (cast_type_to_index != 0)

  3536.         {

  3537.             cast_type_from_index = cast_type_to_index;

  3538.         }

  3539.         else if (info.support_fp16_storage())

  3540.         {

  3541.             cast_type_from_index = 2;

  3542.         }

  3543.         else // if (info.support_fp16_packed())

  3544.         {

  3545.             cast_type_from_index = 1;

  3546.         }

  3547.     }

  3548. 

  3549.     // NCNN_LOGE("convert_packing i2i %d %d %d", cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3550. 

  3551.     const ncnn::Packing_vulkan* uop = d->get_utility_operator(1, 1, cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3552.     uop->forward(src, dst, cmd, opt);

  3553. }

  3554. 

  3555. void VulkanDevice::convert_packing(const VkMat& src, VkImageMat& dst, int dst_elempack, VkCompute& cmd, const Option& opt) const

  3556. {

  3557.     int cast_type_to_index = opt.use_fp16_storage ? 2 : opt.use_fp16_packed ? 1 : 0;

  3558.     int packing_type_to_index = dst_elempack == 1 ? 0 : dst_elempack == 4 ? 1 : 2;

  3559. 

  3560.     int cast_type_from_index;

  3561.     if (src.elembits() == 32)

  3562.     {

  3563.         cast_type_from_index = 0;

  3564.     }

  3565.     else // if (src.elembits() == 16)

  3566.     {

  3567.         if (cast_type_to_index != 0)

  3568.         {

  3569.             cast_type_from_index = cast_type_to_index;

  3570.         }

  3571.         else if (info.support_fp16_storage())

  3572.         {

  3573.             cast_type_from_index = 2;

  3574.         }

  3575.         else // if (info.support_fp16_packed())

  3576.         {

  3577.             cast_type_from_index = 1;

  3578.         }

  3579.     }

  3580. 

  3581.     // NCNN_LOGE("convert_packing b2i %d %d %d", cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3582. 

  3583.     const ncnn::Packing_vulkan* uop = d->get_utility_operator(0, 1, cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3584.     uop->forward(src, dst, cmd, opt);

  3585. }

  3586. 

  3587. void VulkanDevice::convert_packing(const VkImageMat& src, VkMat& dst, int dst_elempack, VkCompute& cmd, const Option& opt) const

  3588. {

  3589.     int cast_type_to_index = opt.use_fp16_storage ? 2 : opt.use_fp16_packed ? 1 : 0;

  3590.     int packing_type_to_index = dst_elempack == 1 ? 0 : dst_elempack == 4 ? 1 : 2;

  3591. 

  3592.     int cast_type_from_index;

  3593.     if (src.elembits() == 32)

  3594.     {

  3595.         cast_type_from_index = 0;

  3596.     }

  3597.     else // if (src.elembits() == 16)

  3598.     {

  3599.         if (cast_type_to_index != 0)

  3600.         {

  3601.             cast_type_from_index = cast_type_to_index;

  3602.         }

  3603.         else if (info.support_fp16_storage())

  3604.         {

  3605.             cast_type_from_index = 2;

  3606.         }

  3607.         else // if (info.support_fp16_packed())

  3608.         {

  3609.             cast_type_from_index = 1;

  3610.         }

  3611.     }

  3612. 

  3613.     // NCNN_LOGE("convert_packing i2b %d %d %d", cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3614. 

  3615.     const ncnn::Packing_vulkan* uop = d->get_utility_operator(1, 0, cast_type_from_index, cast_type_to_index, packing_type_to_index);

  3616.     uop->forward(src, dst, cmd, opt);

  3617. }

  3618. 

  3619. int VulkanDevice::init_device_extension()

  3620. {

  3621.     if (info.support_VK_KHR_bind_memory2())

  3622.     {

  3623.         vkBindBufferMemory2KHR = (PFN_vkBindBufferMemory2KHR)vkGetDeviceProcAddr(d->device, "vkBindBufferMemory2KHR");

  3624.         vkBindImageMemory2KHR = (PFN_vkBindImageMemory2KHR)vkGetDeviceProcAddr(d->device, "vkBindImageMemory2KHR");

  3625.     }

  3626. 

  3627.     if (info.support_VK_KHR_buffer_device_address())

  3628.     {

  3629.         vkGetBufferDeviceAddressKHR = (PFN_vkGetBufferDeviceAddressKHR)vkGetDeviceProcAddr(d->device, "vkGetBufferDeviceAddressKHR");

  3630.         vkGetBufferOpaqueCaptureAddressKHR = (PFN_vkGetBufferOpaqueCaptureAddressKHR)vkGetDeviceProcAddr(d->device, "vkGetBufferOpaqueCaptureAddressKHR");

  3631.         vkGetDeviceMemoryOpaqueCaptureAddressKHR = (PFN_vkGetDeviceMemoryOpaqueCaptureAddressKHR)vkGetDeviceProcAddr(d->device, "vkGetDeviceMemoryOpaqueCaptureAddressKHR");

  3632.     }

  3633. 

  3634.     if (info.support_VK_KHR_descriptor_update_template())

  3635.     {

  3636.         vkCreateDescriptorUpdateTemplateKHR = (PFN_vkCreateDescriptorUpdateTemplateKHR)vkGetDeviceProcAddr(d->device, "vkCreateDescriptorUpdateTemplateKHR");

  3637.         vkDestroyDescriptorUpdateTemplateKHR = (PFN_vkDestroyDescriptorUpdateTemplateKHR)vkGetDeviceProcAddr(d->device, "vkDestroyDescriptorUpdateTemplateKHR");

  3638.         vkUpdateDescriptorSetWithTemplateKHR = (PFN_vkUpdateDescriptorSetWithTemplateKHR)vkGetDeviceProcAddr(d->device, "vkUpdateDescriptorSetWithTemplateKHR");

  3639.     }

  3640. 

  3641.     if (info.support_VK_KHR_get_memory_requirements2())

  3642.     {

  3643.         vkGetImageMemoryRequirements2KHR = (PFN_vkGetImageMemoryRequirements2KHR)vkGetDeviceProcAddr(d->device, "vkGetImageMemoryRequirements2KHR");

  3644.         vkGetBufferMemoryRequirements2KHR = (PFN_vkGetBufferMemoryRequirements2KHR)vkGetDeviceProcAddr(d->device, "vkGetBufferMemoryRequirements2KHR");

  3645.     }

  3646. 

  3647.     if (info.support_VK_KHR_maintenance1())

  3648.     {

  3649.         vkTrimCommandPoolKHR = (PFN_vkTrimCommandPoolKHR)vkGetDeviceProcAddr(d->device, "vkTrimCommandPoolKHR");

  3650.     }

  3651. 

  3652.     if (info.support_VK_KHR_maintenance3())

  3653.     {

  3654.         vkGetDescriptorSetLayoutSupportKHR = (PFN_vkGetDescriptorSetLayoutSupportKHR)vkGetDeviceProcAddr(d->device, "vkGetDescriptorSetLayoutSupportKHR");

  3655.     }

  3656. 

  3657.     if (info.support_VK_KHR_push_descriptor())

  3658.     {

  3659.         if (info.support_VK_KHR_descriptor_update_template())

  3660.         {

  3661.             vkCmdPushDescriptorSetWithTemplateKHR = (PFN_vkCmdPushDescriptorSetWithTemplateKHR)vkGetDeviceProcAddr(d->device, "vkCmdPushDescriptorSetWithTemplateKHR");

  3662.         }

  3663. 

  3664.         vkCmdPushDescriptorSetKHR = (PFN_vkCmdPushDescriptorSetKHR)vkGetDeviceProcAddr(d->device, "vkCmdPushDescriptorSetKHR");

  3665.     }

  3666. 

  3667.     if (info.support_VK_KHR_sampler_ycbcr_conversion())

  3668.     {

  3669.         vkCreateSamplerYcbcrConversionKHR = (PFN_vkCreateSamplerYcbcrConversionKHR)vkGetDeviceProcAddr(d->device, "vkCreateSamplerYcbcrConversionKHR");

  3670.         vkDestroySamplerYcbcrConversionKHR = (PFN_vkDestroySamplerYcbcrConversionKHR)vkGetDeviceProcAddr(d->device, "vkDestroySamplerYcbcrConversionKHR");

  3671.     }

  3672. 

  3673.     if (info.support_VK_KHR_swapchain())

  3674.     {

  3675.         vkCreateSwapchainKHR = (PFN_vkCreateSwapchainKHR)vkGetDeviceProcAddr(d->device, "vkCreateSwapchainKHR");

  3676.         vkDestroySwapchainKHR = (PFN_vkDestroySwapchainKHR)vkGetDeviceProcAddr(d->device, "vkDestroySwapchainKHR");

  3677.         vkGetSwapchainImagesKHR = (PFN_vkGetSwapchainImagesKHR)vkGetDeviceProcAddr(d->device, "vkGetSwapchainImagesKHR");

  3678.         vkAcquireNextImageKHR = (PFN_vkAcquireNextImageKHR)vkGetDeviceProcAddr(d->device, "vkAcquireNextImageKHR");

  3679.         vkQueuePresentKHR = (PFN_vkQueuePresentKHR)vkGetDeviceProcAddr(d->device, "vkQueuePresentKHR");

  3680.     }

  3681. 

  3682.     if (info.support_VK_EXT_buffer_device_address())

  3683.     {

  3684.         vkGetBufferDeviceAddressEXT = (PFN_vkGetBufferDeviceAddressEXT)vkGetDeviceProcAddr(d->device, "vkGetBufferDeviceAddressEXT");

  3685.     }

  3686. 

  3687. #if __ANDROID_API__ >= 26

  3688.     if (info.support_VK_ANDROID_external_memory_android_hardware_buffer())

  3689.     {

  3690.         vkGetAndroidHardwareBufferPropertiesANDROID = (PFN_vkGetAndroidHardwareBufferPropertiesANDROID)vkGetDeviceProcAddr(d->device, "vkGetAndroidHardwareBufferPropertiesANDROID");

  3691.         vkGetMemoryAndroidHardwareBufferANDROID = (PFN_vkGetMemoryAndroidHardwareBufferANDROID)vkGetDeviceProcAddr(d->device, "vkGetMemoryAndroidHardwareBufferANDROID");

  3692.     }

  3693. #endif // __ANDROID_API__ >= 26

  3694. 

  3695.     return 0;

  3696. }

  3697. 

  3698. VulkanDevice* get_gpu_device(int device_index)

  3699. {

  3700.     try_create_gpu_instance();

  3701. 

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

  3703.         return 0;

  3704. 

  3705.     MutexLockGuard lock(g_default_vkdev_lock);

  3706. 

  3707.     if (!g_default_vkdev[device_index])

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

  3709. 

  3710.     return g_default_vkdev[device_index];

  3711. }

  3712. 

  3713. static TBuiltInResource get_default_TBuiltInResource()

  3714. {

  3715.     TBuiltInResource resource;

  3716. 

  3717.     resource.maxLights = 32;

  3718.     resource.maxClipPlanes = 6;

  3719.     resource.maxTextureUnits = 32;

  3720.     resource.maxTextureCoords = 32;

  3721.     resource.maxVertexAttribs = 64;

  3722.     resource.maxVertexUniformComponents = 4096;

  3723.     resource.maxVaryingFloats = 64;

  3724.     resource.maxVertexTextureImageUnits = 32;

  3725.     resource.maxCombinedTextureImageUnits = 80;

  3726.     resource.maxTextureImageUnits = 32;

  3727.     resource.maxFragmentUniformComponents = 4096;

  3728.     resource.maxDrawBuffers = 32;

  3729.     resource.maxVertexUniformVectors = 128;

  3730.     resource.maxVaryingVectors = 8;

  3731.     resource.maxFragmentUniformVectors = 16;

  3732.     resource.maxVertexOutputVectors = 16;

  3733.     resource.maxFragmentInputVectors = 15;

  3734.     resource.minProgramTexelOffset = -8;

  3735.     resource.maxProgramTexelOffset = 7;

  3736.     resource.maxClipDistances = 8;

  3737.     resource.maxComputeWorkGroupCountX = 65535;

  3738.     resource.maxComputeWorkGroupCountY = 65535;

  3739.     resource.maxComputeWorkGroupCountZ = 65535;

  3740.     resource.maxComputeWorkGroupSizeX = 1024;

  3741.     resource.maxComputeWorkGroupSizeY = 1024;

  3742.     resource.maxComputeWorkGroupSizeZ = 64;

  3743.     resource.maxComputeUniformComponents = 1024;

  3744.     resource.maxComputeTextureImageUnits = 16;

  3745.     resource.maxComputeImageUniforms = 8;

  3746.     resource.maxComputeAtomicCounters = 8;

  3747.     resource.maxComputeAtomicCounterBuffers = 1;

  3748.     resource.maxVaryingComponents = 60;

  3749.     resource.maxVertexOutputComponents = 64;

  3750.     resource.maxGeometryInputComponents = 64;

  3751.     resource.maxGeometryOutputComponents = 128;

  3752.     resource.maxFragmentInputComponents = 128;

  3753.     resource.maxImageUnits = 8;

  3754.     resource.maxCombinedImageUnitsAndFragmentOutputs = 8;

  3755.     resource.maxCombinedShaderOutputResources = 8;

  3756.     resource.maxImageSamples = 0;

  3757.     resource.maxVertexImageUniforms = 0;

  3758.     resource.maxTessControlImageUniforms = 0;

  3759.     resource.maxTessEvaluationImageUniforms = 0;

  3760.     resource.maxGeometryImageUniforms = 0;

  3761.     resource.maxFragmentImageUniforms = 8;

  3762.     resource.maxCombinedImageUniforms = 8;

  3763.     resource.maxGeometryTextureImageUnits = 16;

  3764.     resource.maxGeometryOutputVertices = 256;

  3765.     resource.maxGeometryTotalOutputComponents = 1024;

  3766.     resource.maxGeometryUniformComponents = 1024;

  3767.     resource.maxGeometryVaryingComponents = 64;

  3768.     resource.maxTessControlInputComponents = 128;

  3769.     resource.maxTessControlOutputComponents = 128;

  3770.     resource.maxTessControlTextureImageUnits = 16;

  3771.     resource.maxTessControlUniformComponents = 1024;

  3772.     resource.maxTessControlTotalOutputComponents = 4096;

  3773.     resource.maxTessEvaluationInputComponents = 128;

  3774.     resource.maxTessEvaluationOutputComponents = 128;

  3775.     resource.maxTessEvaluationTextureImageUnits = 16;

  3776.     resource.maxTessEvaluationUniformComponents = 1024;

  3777.     resource.maxTessPatchComponents = 120;

  3778.     resource.maxPatchVertices = 32;

  3779.     resource.maxTessGenLevel = 64;

  3780.     resource.maxViewports = 16;

  3781.     resource.maxVertexAtomicCounters = 0;

  3782.     resource.maxTessControlAtomicCounters = 0;

  3783.     resource.maxTessEvaluationAtomicCounters = 0;

  3784.     resource.maxGeometryAtomicCounters = 0;

  3785.     resource.maxFragmentAtomicCounters = 8;

  3786.     resource.maxCombinedAtomicCounters = 8;

  3787.     resource.maxAtomicCounterBindings = 1;

  3788.     resource.maxVertexAtomicCounterBuffers = 0;

  3789.     resource.maxTessControlAtomicCounterBuffers = 0;

  3790.     resource.maxTessEvaluationAtomicCounterBuffers = 0;

  3791.     resource.maxGeometryAtomicCounterBuffers = 0;

  3792.     resource.maxFragmentAtomicCounterBuffers = 1;

  3793.     resource.maxCombinedAtomicCounterBuffers = 1;

  3794.     resource.maxAtomicCounterBufferSize = 16384;

  3795.     resource.maxTransformFeedbackBuffers = 4;

  3796.     resource.maxTransformFeedbackInterleavedComponents = 64;

  3797.     resource.maxCullDistances = 8;

  3798.     resource.maxCombinedClipAndCullDistances = 8;

  3799.     resource.maxSamples = 4;

  3800.     resource.maxMeshOutputVerticesNV = 256;

  3801.     resource.maxMeshOutputPrimitivesNV = 512;

  3802.     resource.maxMeshWorkGroupSizeX_NV = 32;

  3803.     resource.maxMeshWorkGroupSizeY_NV = 1;

  3804.     resource.maxMeshWorkGroupSizeZ_NV = 1;

  3805.     resource.maxTaskWorkGroupSizeX_NV = 32;

  3806.     resource.maxTaskWorkGroupSizeY_NV = 1;

  3807.     resource.maxTaskWorkGroupSizeZ_NV = 1;

  3808.     resource.maxMeshViewCountNV = 4;

  3809. 

  3810.     // TODO compile-time glslang version check

  3811.     // resource.maxDualSourceDrawBuffersEXT = 1;

  3812. 

  3813.     resource.limits.nonInductiveForLoops = 1;

  3814.     resource.limits.whileLoops = 1;

  3815.     resource.limits.doWhileLoops = 1;

  3816.     resource.limits.generalUniformIndexing = 1;

  3817.     resource.limits.generalAttributeMatrixVectorIndexing = 1;

  3818.     resource.limits.generalVaryingIndexing = 1;

  3819.     resource.limits.generalSamplerIndexing = 1;

  3820.     resource.limits.generalVariableIndexing = 1;

  3821.     resource.limits.generalConstantMatrixVectorIndexing = 1;

  3822. 

  3823.     return resource;

  3824. }

  3825. 

  3826. class VulkanShaderIncluder : public glslang::TShader::Includer

  3827. {

  3828. public:

  3829.     virtual glslang::TShader::Includer::IncludeResult* includeLocal(const char* headerName, const char* /*includerName*/, size_t /*inclusionDepth*/)

  3830.     {

  3831.         if (strcmp(headerName, "vulkan_activation.comp") == 0)

  3832.         {

  3833.             const char* const headerData = vulkan_activation_comp_data;

  3834.             const size_t headerLength = sizeof(vulkan_activation_comp_data);

  3835.             glslang::TShader::Includer::IncludeResult* r = new glslang::TShader::Includer::IncludeResult(headerName, headerData, headerLength, 0);

  3836.             return r;

  3837.         }

  3838. 

  3839.         return 0;

  3840.     }

  3841. 

  3842.     virtual void releaseInclude(glslang::TShader::Includer::IncludeResult* r)

  3843.     {

  3844.         delete r;

  3845.     }

  3846. };

  3847. 

  3848. int compile_spirv_module(const char* comp_string, const Option& opt, std::vector<uint32_t>& spirv)

  3849. {

  3850.     // -1 for omitting the tail '\0'

  3851.     int length = strlen(comp_string) - 1;

  3852.     return compile_spirv_module(comp_string, length, opt, spirv);

  3853. }

  3854. 

  3855. int compile_spirv_module(const char* comp_data, int comp_data_size, const Option& opt, std::vector<uint32_t>& spirv)

  3856. {

  3857.     std::vector<std::pair<const char*, const char*> > custom_defines;

  3858. 

  3859.     if (opt.use_fp16_storage)

  3860.     {

  3861.         custom_defines.push_back(std::make_pair("sfp", "float16_t"));

  3862.         custom_defines.push_back(std::make_pair("sfpvec2", "f16vec2"));

  3863.         custom_defines.push_back(std::make_pair("sfpvec4", "f16vec4"));

  3864. 

  3865.         if (opt.use_fp16_arithmetic)

  3866.         {

  3867.             custom_defines.push_back(std::make_pair("sfpvec8", "f16mat2x4"));

  3868.             custom_defines.push_back(std::make_pair("sfpmat4", "f16mat4"));

  3869.         }

  3870.     }

  3871.     else if (opt.use_fp16_packed)

  3872.     {

  3873.         custom_defines.push_back(std::make_pair("sfp", "float"));

  3874.         custom_defines.push_back(std::make_pair("sfpvec2", "uint"));

  3875.         custom_defines.push_back(std::make_pair("sfpvec4", "uvec2"));

  3876.         custom_defines.push_back(std::make_pair("sfpvec8", "uvec4"));

  3877.     }

  3878.     else

  3879.     {

  3880.         custom_defines.push_back(std::make_pair("sfp", "float"));

  3881.         custom_defines.push_back(std::make_pair("sfpvec2", "vec2"));

  3882.         custom_defines.push_back(std::make_pair("sfpvec4", "vec4"));

  3883.         custom_defines.push_back(std::make_pair("sfpvec8", "mat2x4"));

  3884.         custom_defines.push_back(std::make_pair("sfpmat4", "mat4"));

  3885.     }

  3886. 

  3887.     if (opt.use_fp16_arithmetic)

  3888.     {

  3889.         custom_defines.push_back(std::make_pair("afp", "float16_t"));

  3890.         custom_defines.push_back(std::make_pair("afpvec2", "f16vec2"));

  3891.         custom_defines.push_back(std::make_pair("afpvec4", "f16vec4"));

  3892.         custom_defines.push_back(std::make_pair("afpvec8", "f16mat2x4"));

  3893.         custom_defines.push_back(std::make_pair("afpmat4", "f16mat4"));

  3894.     }

  3895.     else

  3896.     {

  3897.         custom_defines.push_back(std::make_pair("afp", "float"));

  3898.         custom_defines.push_back(std::make_pair("afpvec2", "vec2"));

  3899.         custom_defines.push_back(std::make_pair("afpvec4", "vec4"));

  3900.         custom_defines.push_back(std::make_pair("afpvec8", "mat2x4"));

  3901.         custom_defines.push_back(std::make_pair("afpmat4", "mat4"));

  3902.     }

  3903. 

  3904.     if (opt.use_fp16_storage && opt.use_fp16_uniform && opt.use_fp16_arithmetic)

  3905.     {

  3906.         custom_defines.push_back(std::make_pair("lfp", "float16_t"));

  3907.         custom_defines.push_back(std::make_pair("lfpvec4", "f16vec4"));

  3908.     }

  3909.     else if (opt.use_fp16_storage && opt.use_fp16_arithmetic)

  3910.     {

  3911.         custom_defines.push_back(std::make_pair("lfp", "float"));

  3912.         custom_defines.push_back(std::make_pair("lfpvec4", "uint64_t"));

  3913.     }

  3914.     else if (opt.use_fp16_storage || opt.use_fp16_packed)

  3915.     {

  3916.         custom_defines.push_back(std::make_pair("lfp", "float"));

  3917.         custom_defines.push_back(std::make_pair("lfpvec4", "uvec2"));

  3918.     }

  3919.     else

  3920.     {

  3921.         custom_defines.push_back(std::make_pair("lfp", "float"));

  3922.         custom_defines.push_back(std::make_pair("lfpvec4", "vec4"));

  3923.     }

  3924. 

  3925.     if (opt.use_fp16_storage && opt.use_fp16_uniform && opt.use_fp16_arithmetic)

  3926.     {

  3927.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "v"));

  3928.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "v"));

  3929. 

  3930.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "v"));

  3931.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "v"));

  3932.     }

  3933.     else if (opt.use_fp16_storage && opt.use_fp16_arithmetic)

  3934.     {

  3935.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "float(v)"));

  3936.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "pack64(halfBitsToUInt16(v))"));

  3937. 

  3938.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "float16_t(v)"));

  3939.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "int16BitsToHalf(unpack16(v))"));

  3940.     }

  3941.     else if (opt.use_fp16_packed && opt.use_fp16_arithmetic)

  3942.     {

  3943.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "v"));

  3944.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "v"));

  3945. 

  3946.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "float16_t(v)"));

  3947.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "f16vec4(unpackFloat2x16(v.x),unpackFloat2x16(v.y))"));

  3948.     }

  3949.     else if (opt.use_fp16_storage)

  3950.     {

  3951.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "float(v)"));

  3952.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "uvec2(packHalf2x16(vec4(v).rg),packHalf2x16(vec4(v).ba))"));

  3953. 

  3954.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "v"));

  3955.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "vec4(unpackHalf2x16(v.x),unpackHalf2x16(v.y))"));

  3956.     }

  3957.     else if (opt.use_fp16_packed)

  3958.     {

  3959.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "v"));

  3960.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "v"));

  3961. 

  3962.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "v"));

  3963.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "vec4(unpackHalf2x16(v.x),unpackHalf2x16(v.y))"));

  3964.     }

  3965.     else

  3966.     {

  3967.         custom_defines.push_back(std::make_pair("sfp2lfp(v)", "v"));

  3968.         custom_defines.push_back(std::make_pair("sfp2lfpvec4(v)", "v"));

  3969. 

  3970.         custom_defines.push_back(std::make_pair("lfp2afp(v)", "v"));

  3971.         custom_defines.push_back(std::make_pair("lfp2afpvec4(v)", "v"));

  3972.     }

  3973. 

  3974.     if (opt.use_fp16_storage && opt.use_fp16_arithmetic)

  3975.     {

  3976.         custom_defines.push_back(std::make_pair("buffer_ld1(buf,i)", "buf[i]"));

  3977.         custom_defines.push_back(std::make_pair("buffer_st1(buf,i,v)", "{buf[i]=v;}"));

  3978.         custom_defines.push_back(std::make_pair("buffer_cp1(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  3979.         custom_defines.push_back(std::make_pair("buffer_cp1to4(buf,i,sbuf,si4)", "{buf[i]=f16vec4(sbuf[si4.r],sbuf[si4.g],sbuf[si4.b],sbuf[si4.a]);}"));

  3980.         custom_defines.push_back(std::make_pair("buffer_cp1to8(buf,i,sbuf,si4,sii4)", "{buf[i]=f16mat2x4(sbuf[si4.r],sbuf[si4.g],sbuf[si4.b],sbuf[si4.a],sbuf[sii4.r],sbuf[sii4.g],sbuf[sii4.b],sbuf[sii4.a]);}"));

  3981.         custom_defines.push_back(std::make_pair("buffer_ld2(buf,i)", "buf[i]"));

  3982.         custom_defines.push_back(std::make_pair("buffer_st2(buf,i,v)", "{buf[i]=v;}"));

  3983.         custom_defines.push_back(std::make_pair("buffer_cp2(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  3984.         custom_defines.push_back(std::make_pair("buffer_ld4(buf,i)", "buf[i]"));

  3985.         custom_defines.push_back(std::make_pair("buffer_st4(buf,i,v)", "{buf[i]=v;}"));

  3986.         custom_defines.push_back(std::make_pair("buffer_cp4(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  3987.         custom_defines.push_back(std::make_pair("buffer_cp4to1(buf,i4,sbuf,si)", "{buf[i4.r]=sbuf[si].r;buf[i4.g]=sbuf[si].g;buf[i4.b]=sbuf[si].b;buf[i4.a]=sbuf[si].a;}"));

  3988.         custom_defines.push_back(std::make_pair("buffer_cp4to8(buf,i,sbuf,si2)", "{buf[i]=f16mat2x4(sbuf[si2.r],sbuf[si2.g]);}"));

  3989.         custom_defines.push_back(std::make_pair("buffer_ld8(buf,i)", "buf[i]"));

  3990.         custom_defines.push_back(std::make_pair("buffer_st8(buf,i,v)", "{buf[i]=v;}"));

  3991.         custom_defines.push_back(std::make_pair("buffer_cp8(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  3992.         custom_defines.push_back(std::make_pair("buffer_cp8to1(buf,i4,ii4,sbuf,si)", "{f16mat2x4 _v=sbuf[si]; buf[i4.r]=_v[0].r;buf[i4.g]=_v[0].g;buf[i4.b]=_v[0].b;buf[i4.a]=_v[0].a; buf[ii4.r]=_v[1].r;buf[ii4.g]=_v[1].g;buf[ii4.b]=_v[1].b;buf[ii4.a]=_v[1].a;}"));

  3993.         custom_defines.push_back(std::make_pair("buffer_cp8to4(buf,i2,sbuf,si)", "{f16mat2x4 _v=sbuf[si]; buf[i2.r]=_v[0];buf[i2.g]=_v[1];}"));

  3994.         custom_defines.push_back(std::make_pair("sfp2afpmat4(v)", "v"));

  3995.         custom_defines.push_back(std::make_pair("afp2sfpmat4(v)", "v"));

  3996.     }

  3997.     else if (opt.use_fp16_packed && opt.use_fp16_arithmetic)

  3998.     {

  3999.         custom_defines.push_back(std::make_pair("buffer_ld1(buf,i)", "float16_t(buf[i])"));

  4000.         custom_defines.push_back(std::make_pair("buffer_st1(buf,i,v)", "{buf[i]=float(v);}"));

  4001.         custom_defines.push_back(std::make_pair("buffer_cp1(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4002.         custom_defines.push_back(std::make_pair("buffer_cp1to4(buf,i,sbuf,si4)", "{buf[i]=uvec2(packFloat2x16(f16vec2(sbuf[si4.r],sbuf[si4.g])),packFloat2x16(f16vec2(sbuf[si4.b],sbuf[si4.a])));}"));

  4003.         custom_defines.push_back(std::make_pair("buffer_cp1to8(buf,i,sbuf,si4,sii4)", "{buf[i]=uvec4(packFloat2x16(f16vec2(sbuf[si4.r],sbuf[si4.g])),packFloat2x16(f16vec2(sbuf[si4.b],sbuf[si4.a])),packFloat2x16(f16vec2(sbuf[sii4.r],sbuf[sii4.g])),packFloat2x16(f16vec2(sbuf[sii4.b],sbuf[sii4.a])));}"));

  4004.         custom_defines.push_back(std::make_pair("buffer_ld2(buf,i)", "unpackFloat2x16(buf[i])"));

  4005.         custom_defines.push_back(std::make_pair("buffer_st2(buf,i,v)", "{buf[i]=packFloat2x16(v)}"));

  4006.         custom_defines.push_back(std::make_pair("buffer_cp2(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4007.         custom_defines.push_back(std::make_pair("buffer_ld4(buf,i)", "f16vec4(unpackFloat2x16(buf[i].x),unpackFloat2x16(buf[i].y))"));

  4008.         custom_defines.push_back(std::make_pair("buffer_st4(buf,i,v)", "{buf[i]=uvec2(packFloat2x16(v.rg),packFloat2x16(v.ba));}"));

  4009.         custom_defines.push_back(std::make_pair("buffer_cp4(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4010.         custom_defines.push_back(std::make_pair("buffer_cp4to1(buf,i4,sbuf,si)", "{uvec2 _v=sbuf[si]; f16vec2 _v0=unpackFloat2x16(_v.x);f16vec2 _v1=unpackFloat2x16(_v.y); buf[i4.r]=_v0.r;buf[i4.g]=_v0.g;buf[i4.b]=_v1.r;buf[i4.a]=_v1.g;}"));

  4011.         custom_defines.push_back(std::make_pair("buffer_cp4to8(buf,i,sbuf,si2)", "{buf[i]=uvec4(sbuf[si2.r],sbuf[si2.g]);}"));

  4012.         custom_defines.push_back(std::make_pair("buffer_ld8(buf,i)", "f16mat2x4(f16vec4(unpackFloat2x16(buf[i].r),unpackFloat2x16(buf[i].g)),f16vec4(unpackFloat2x16(buf[i].b),unpackFloat2x16(buf[i].a)))"));

  4013.         custom_defines.push_back(std::make_pair("buffer_st8(buf,i,v)", "{buf[i]=uvec4(uvec2(packFloat2x16(v[0].rg),packFloat2x16(v[0].ba)),uvec2(packFloat2x16(v[1].rg),packFloat2x16(v[1].ba)));}"));

  4014.         custom_defines.push_back(std::make_pair("buffer_cp8(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4015.         custom_defines.push_back(std::make_pair("buffer_cp8to1(buf,i4,ii4,sbuf,si)", "{uvec4 _v=sbuf[si]; f16vec2 _v0=unpackFloat2x16(_v.r);f16vec2 _v1=unpackFloat2x16(_v.g);f16vec2 _v2=unpackFloat2x16(_v.b);f16vec2 _v3=unpackFloat2x16(_v.a); buf[i4.r]=_v0.r;buf[i4.g]=_v0.g;buf[i4.b]=_v1.r;buf[i4.a]=_v1.g; buf[ii4.r]=_v2.r;buf[ii4.g]=_v2.g;buf[ii4.b]=_v3.r;buf[ii4.a]=_v3.g;}"));

  4016.         custom_defines.push_back(std::make_pair("buffer_cp8to4(buf,i2,sbuf,si)", "{uvec4 _v=sbuf[si]; buf[i2.r]=_v.rg;buf[i2.g]=_v.ba;}"));

  4017.     }

  4018.     else if (opt.use_fp16_storage)

  4019.     {

  4020.         custom_defines.push_back(std::make_pair("buffer_ld1(buf,i)", "float(buf[i])"));

  4021.         custom_defines.push_back(std::make_pair("buffer_st1(buf,i,v)", "{buf[i]=float16_t(v);}"));

  4022.         custom_defines.push_back(std::make_pair("buffer_cp1(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4023.         custom_defines.push_back(std::make_pair("buffer_cp1to4(buf,i,sbuf,si4)", "{buf[i].r=sbuf[si4.r];buf[i].g=sbuf[si4.g];buf[i].b=sbuf[si4.b];buf[i].a=sbuf[si4.a];}"));

  4024.         custom_defines.push_back(std::make_pair("buffer_cp1to8(buf,i,sbuf,si4,sii4)", "{buf[i].abcd.r=sbuf[si4.r];buf[i].abcd.g=sbuf[si4.g];buf[i].abcd.b=sbuf[si4.b];buf[i].abcd.a=sbuf[si4.a];buf[i].efgh.r=sbuf[sii4.r];buf[i].efgh.g=sbuf[sii4.g];buf[i].efgh.b=sbuf[sii4.b];buf[i].efgh.a=sbuf[sii4.a];}"));

  4025.         custom_defines.push_back(std::make_pair("buffer_ld2(buf,i)", "vec2(buf[i])"));

  4026.         custom_defines.push_back(std::make_pair("buffer_st2(buf,i,v)", "{buf[i]=f16vec2(v);}"));

  4027.         custom_defines.push_back(std::make_pair("buffer_cp2(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4028.         custom_defines.push_back(std::make_pair("buffer_ld4(buf,i)", "vec4(buf[i])"));

  4029.         custom_defines.push_back(std::make_pair("buffer_st4(buf,i,v)", "{buf[i]=f16vec4(v);}"));

  4030.         custom_defines.push_back(std::make_pair("buffer_cp4(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4031.         custom_defines.push_back(std::make_pair("buffer_cp4to1(buf,i4,sbuf,si)", "{buf[i4.r]=sbuf[si].r;buf[i4.g]=sbuf[si].g;buf[i4.b]=sbuf[si].b;buf[i4.a]=sbuf[si].a;}"));

  4032.         custom_defines.push_back(std::make_pair("buffer_cp4to8(buf,i,sbuf,si2)", "{buf[i].abcd=sbuf[si2.r];buf[i].efgh=sbuf[si2.g];}"));

  4033.         custom_defines.push_back(std::make_pair("buffer_ld8(buf,i)", "mat2x4(vec4(buf[i].abcd),vec4(buf[i].efgh))"));

  4034.         custom_defines.push_back(std::make_pair("buffer_st8(buf,i,v)", "{buf[i].abcd=f16vec4(v[0]);buf[i].efgh=f16vec4(v[1]);}"));

  4035.         custom_defines.push_back(std::make_pair("buffer_cp8(buf,i,sbuf,si)", "{buf[i].abcd=sbuf[si].abcd;buf[i].efgh=sbuf[si].efgh;}"));

  4036.         custom_defines.push_back(std::make_pair("buffer_cp8to1(buf,i4,ii4,sbuf,si)", "{buf[i4.r]=sbuf[si].abcd.r;buf[i4.g]=sbuf[si].abcd.g;buf[i4.b]=sbuf[si].abcd.b;buf[i4.a]=sbuf[si].abcd.a; buf[ii4.r]=sbuf[si].efgh.r;buf[ii4.g]=sbuf[si].efgh.g;buf[ii4.b]=sbuf[si].efgh.b;buf[ii4.a]=sbuf[si].efgh.a;}"));

  4037.         custom_defines.push_back(std::make_pair("buffer_cp8to4(buf,i2,sbuf,si)", "{buf[i2.r]=sbuf[si].abcd;buf[i2.g]=sbuf[si].efgh;}"));

  4038.     }

  4039.     else if (opt.use_fp16_packed)

  4040.     {

  4041.         custom_defines.push_back(std::make_pair("buffer_ld1(buf,i)", "buf[i]"));

  4042.         custom_defines.push_back(std::make_pair("buffer_st1(buf,i,v)", "{buf[i]=v;}"));

  4043.         custom_defines.push_back(std::make_pair("buffer_cp1(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4044.         custom_defines.push_back(std::make_pair("buffer_cp1to4(buf,i,sbuf,si4)", "{buf[i]=uvec2(packHalf2x16(vec2(sbuf[si4.r],sbuf[si4.g])),packHalf2x16(vec2(sbuf[si4.b],sbuf[si4.a])));}"));

  4045.         custom_defines.push_back(std::make_pair("buffer_cp1to8(buf,i,sbuf,si4,sii4)", "{buf[i]=uvec4(packHalf2x16(vec2(sbuf[si4.r],sbuf[si4.g])),packHalf2x16(vec2(sbuf[si4.b],sbuf[si4.a])),packHalf2x16(vec2(sbuf[sii4.r],sbuf[sii4.g])),packHalf2x16(vec2(sbuf[sii4.b],sbuf[sii4.a])));}"));

  4046.         custom_defines.push_back(std::make_pair("buffer_ld2(buf,i)", "unpackHalf2x16(buf[i])"));

  4047.         custom_defines.push_back(std::make_pair("buffer_st2(buf,i,v)", "{buf[i]=packHalf2x16(v)}"));

  4048.         custom_defines.push_back(std::make_pair("buffer_cp2(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4049.         custom_defines.push_back(std::make_pair("buffer_ld4(buf,i)", "vec4(unpackHalf2x16(buf[i].x),unpackHalf2x16(buf[i].y))"));

  4050.         custom_defines.push_back(std::make_pair("buffer_st4(buf,i,v)", "{buf[i]=uvec2(packHalf2x16(v.rg),packHalf2x16(v.ba));}"));

  4051.         custom_defines.push_back(std::make_pair("buffer_cp4(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4052.         custom_defines.push_back(std::make_pair("buffer_cp4to1(buf,i4,sbuf,si)", "{uvec2 _v=sbuf[si]; vec2 _v0=unpackHalf2x16(_v.x);vec2 _v1=unpackHalf2x16(_v.y); buf[i4.r]=_v0.r;buf[i4.g]=_v0.g;buf[i4.b]=_v1.r;buf[i4.a]=_v1.g;}"));

  4053.         custom_defines.push_back(std::make_pair("buffer_cp4to8(buf,i,sbuf,si2)", "{buf[i]=uvec4(sbuf[si2.r],sbuf[si2.g]);}"));

  4054.         custom_defines.push_back(std::make_pair("buffer_ld8(buf,i)", "mat2x4(vec4(unpackHalf2x16(buf[i].r),unpackHalf2x16(buf[i].g)),vec4(unpackHalf2x16(buf[i].b),unpackHalf2x16(buf[i].a)))"));

  4055.         custom_defines.push_back(std::make_pair("buffer_st8(buf,i,v)", "{buf[i]=uvec4(uvec2(packHalf2x16(v[0].rg),packHalf2x16(v[0].ba)),uvec2(packHalf2x16(v[1].rg),packHalf2x16(v[1].ba)));}"));

  4056.         custom_defines.push_back(std::make_pair("buffer_cp8(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4057.         custom_defines.push_back(std::make_pair("buffer_cp8to1(buf,i4,ii4,sbuf,si)", "{uvec4 _v=sbuf[si]; vec2 _v0=unpackHalf2x16(_v.r);vec2 _v1=unpackHalf2x16(_v.g);vec2 _v2=unpackHalf2x16(_v.b);vec2 _v3=unpackHalf2x16(_v.a); buf[i4.r]=_v0.r;buf[i4.g]=_v0.g;buf[i4.b]=_v1.r;buf[i4.a]=_v1.g; buf[ii4.r]=_v2.r;buf[ii4.g]=_v2.g;buf[ii4.b]=_v3.r;buf[ii4.a]=_v3.g;}"));

  4058.         custom_defines.push_back(std::make_pair("buffer_cp8to4(buf,i2,sbuf,si)", "{uvec4 _v=sbuf[si]; buf[i2.r]=_v.rg;buf[i2.g]=_v.ba;}"));

  4059.     }

  4060.     else

  4061.     {

  4062.         custom_defines.push_back(std::make_pair("buffer_ld1(buf,i)", "buf[i]"));

  4063.         custom_defines.push_back(std::make_pair("buffer_st1(buf,i,v)", "{buf[i]=v;}"));

  4064.         custom_defines.push_back(std::make_pair("buffer_cp1(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4065.         custom_defines.push_back(std::make_pair("buffer_cp1to4(buf,i,sbuf,si4)", "{buf[i]=vec4(sbuf[si4.r],sbuf[si4.g],sbuf[si4.b],sbuf[si4.a]);}"));

  4066.         custom_defines.push_back(std::make_pair("buffer_cp1to8(buf,i,sbuf,si4,sii4)", "{buf[i]=mat2x4(sbuf[si4.r],sbuf[si4.g],sbuf[si4.b],sbuf[si4.a],sbuf[sii4.r],sbuf[sii4.g],sbuf[sii4.b],sbuf[sii4.a]);}"));

  4067.         custom_defines.push_back(std::make_pair("buffer_ld2(buf,i)", "buf[i]"));

  4068.         custom_defines.push_back(std::make_pair("buffer_st2(buf,i,v)", "{buf[i]=v;}"));

  4069.         custom_defines.push_back(std::make_pair("buffer_cp2(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4070.         custom_defines.push_back(std::make_pair("buffer_ld4(buf,i)", "buf[i]"));

  4071.         custom_defines.push_back(std::make_pair("buffer_st4(buf,i,v)", "{buf[i]=v;}"));

  4072.         custom_defines.push_back(std::make_pair("buffer_cp4(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4073.         custom_defines.push_back(std::make_pair("buffer_cp4to1(buf,i4,sbuf,si)", "{vec4 _v=sbuf[si]; buf[i4.r]=_v.r;buf[i4.g]=_v.g;buf[i4.b]=_v.b;buf[i4.a]=_v.a;}"));

  4074.         custom_defines.push_back(std::make_pair("buffer_cp4to8(buf,i,sbuf,si2)", "{buf[i]=mat2x4(sbuf[si2.r],sbuf[si2.g]);}"));

  4075.         custom_defines.push_back(std::make_pair("buffer_ld8(buf,i)", "buf[i]"));

  4076.         custom_defines.push_back(std::make_pair("buffer_st8(buf,i,v)", "{buf[i]=v;}"));

  4077.         custom_defines.push_back(std::make_pair("buffer_cp8(buf,i,sbuf,si)", "{buf[i]=sbuf[si];}"));

  4078.         custom_defines.push_back(std::make_pair("buffer_cp8to1(buf,i4,ii4,sbuf,si)", "{mat2x4 _v=sbuf[si]; buf[i4.r]=_v[0].r;buf[i4.g]=_v[0].g;buf[i4.b]=_v[0].b;buf[i4.a]=_v[0].a; buf[ii4.r]=_v[1].r;buf[ii4.g]=_v[1].g;buf[ii4.b]=_v[1].b;buf[ii4.a]=_v[1].a;}"));

  4079.         custom_defines.push_back(std::make_pair("buffer_cp8to4(buf,i2,sbuf,si)", "{mat2x4 _v=sbuf[si]; buf[i2.r]=_v[0];buf[i2.g]=_v[1];}"));

  4080.         custom_defines.push_back(std::make_pair("sfp2afpmat4(v)", "v"));

  4081.         custom_defines.push_back(std::make_pair("afp2sfpmat4(v)", "v"));

  4082.     }

  4083. 

  4084.     if (opt.use_image_storage)

  4085.     {

  4086.         if (opt.use_fp16_storage)

  4087.         {

  4088.             custom_defines.push_back(std::make_pair("imfmtc1", "r16f"));

  4089.             custom_defines.push_back(std::make_pair("imfmtc4", "rgba16f"));

  4090.             custom_defines.push_back(std::make_pair("unfp", "mediump"));

  4091.         }

  4092.         else if (opt.use_fp16_packed)

  4093.         {

  4094.             custom_defines.push_back(std::make_pair("imfmtc1", "r32f"));

  4095.             custom_defines.push_back(std::make_pair("imfmtc4", "rgba16f"));

  4096.             custom_defines.push_back(std::make_pair("unfp", "mediump"));

  4097.         }

  4098.         else

  4099.         {

  4100.             custom_defines.push_back(std::make_pair("imfmtc1", "r32f"));

  4101.             custom_defines.push_back(std::make_pair("imfmtc4", "rgba32f"));

  4102.             custom_defines.push_back(std::make_pair("unfp", "highp"));

  4103.         }

  4104. 

  4105.         if (opt.use_fp16_storage && opt.use_fp16_arithmetic)

  4106.         {

  4107.             custom_defines.push_back(std::make_pair("image1d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4108.             custom_defines.push_back(std::make_pair("image2d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4109.             custom_defines.push_back(std::make_pair("image3d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4110.             custom_defines.push_back(std::make_pair("image1d_st1(img,p,v)", "{vec4 _v;_v.r=float(v);imageStore(img,p,_v);}"));

  4111.             custom_defines.push_back(std::make_pair("image2d_st1(img,p,v)", "{vec4 _v;_v.r=float(v);imageStore(img,p,_v);}"));

  4112.             custom_defines.push_back(std::make_pair("image3d_st1(img,p,v)", "{vec4 _v;_v.r=float(v);imageStore(img,p,_v);}"));

  4113.             custom_defines.push_back(std::make_pair("image1d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4114.             custom_defines.push_back(std::make_pair("image2d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4115.             custom_defines.push_back(std::make_pair("image3d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4116.             custom_defines.push_back(std::make_pair("image1d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4117.             custom_defines.push_back(std::make_pair("image2d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4118.             custom_defines.push_back(std::make_pair("image3d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4119.             custom_defines.push_back(std::make_pair("image1d_st4(img,p,v)", "{imageStore(img,p,vec4(v));}"));

  4120.             custom_defines.push_back(std::make_pair("image2d_st4(img,p,v)", "{imageStore(img,p,vec4(v));}"));

  4121.             custom_defines.push_back(std::make_pair("image3d_st4(img,p,v)", "{imageStore(img,p,vec4(v));}"));

  4122.             custom_defines.push_back(std::make_pair("image1d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4123.             custom_defines.push_back(std::make_pair("image2d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4124.             custom_defines.push_back(std::make_pair("image3d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4125.             custom_defines.push_back(std::make_pair("image1d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,(p)*2,0),texelFetch(tex,(p)*2+1,0))"));

  4126.             custom_defines.push_back(std::make_pair("image2d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,ivec2(p.x*2,p.y),0),texelFetch(tex,ivec2(p.x*2+1,p.y),0))"));

  4127.             custom_defines.push_back(std::make_pair("image3d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,ivec3(p.x*2,p.y,p.z),0),texelFetch(tex,ivec3(p.x*2+1,p.y,p.z),0))"));

  4128.             custom_defines.push_back(std::make_pair("image1d_st8(img,p,v)", "{imageStore(img,(p)*2,vec4(v[0]));imageStore(img,(p)*2+1,vec4(v[1]));}"));

  4129.             custom_defines.push_back(std::make_pair("image2d_st8(img,p,v)", "{imageStore(img,ivec2(p.x*2,p.y),vec4(v[0]));imageStore(img,ivec2(p.x*2+1,p.y),vec4(v[1]));}"));

  4130.             custom_defines.push_back(std::make_pair("image3d_st8(img,p,v)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),vec4(v[0]));imageStore(img,ivec3(p.x*2+1,p.y,p.z),vec4(v[1]));}"));

  4131.             custom_defines.push_back(std::make_pair("image1d_cp8(img,p,tex,sp)", "{imageStore(img,(p)*2,texelFetch(tex,sp*2,0));imageStore(img,(p)*2+1,texelFetch(tex,sp*2+1,0));}"));

  4132.             custom_defines.push_back(std::make_pair("image2d_cp8(img,p,tex,sp)", "{imageStore(img,ivec2(p.x*2,p.y),texelFetch(tex,ivec2(sp.x*2,sp.y),0));imageStore(img,ivec2(p.x*2+1,p.y),texelFetch(tex,ivec2(sp.x*2+1,sp.y),0));}"));

  4133.             custom_defines.push_back(std::make_pair("image3d_cp8(img,p,tex,sp)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),texelFetch(tex,ivec3(sp.x*2,sp.y,sp.z),0));imageStore(img,ivec3(p.x*2+1,p.y,p.z),texelFetch(tex,ivec3(sp.x*2+1,sp.y,sp.z),0));}"));

  4134.         }

  4135.         else if (opt.use_fp16_packed && opt.use_fp16_arithmetic)

  4136.         {

  4137.             custom_defines.push_back(std::make_pair("image1d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4138.             custom_defines.push_back(std::make_pair("image2d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4139.             custom_defines.push_back(std::make_pair("image3d_ld1(tex,p)", "float16_t(texelFetch(tex,p,0).r)"));

  4140.             custom_defines.push_back(std::make_pair("image1d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4141.             custom_defines.push_back(std::make_pair("image2d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4142.             custom_defines.push_back(std::make_pair("image3d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4143.             custom_defines.push_back(std::make_pair("image1d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4144.             custom_defines.push_back(std::make_pair("image2d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4145.             custom_defines.push_back(std::make_pair("image3d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4146.             custom_defines.push_back(std::make_pair("image1d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4147.             custom_defines.push_back(std::make_pair("image2d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4148.             custom_defines.push_back(std::make_pair("image3d_ld4(tex,p)", "f16vec4(texelFetch(tex,p,0))"));

  4149.             custom_defines.push_back(std::make_pair("image1d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4150.             custom_defines.push_back(std::make_pair("image2d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4151.             custom_defines.push_back(std::make_pair("image3d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4152.             custom_defines.push_back(std::make_pair("image1d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4153.             custom_defines.push_back(std::make_pair("image2d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4154.             custom_defines.push_back(std::make_pair("image3d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4155.             custom_defines.push_back(std::make_pair("image1d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,(p)*2,0),texelFetch(tex,(p)*2+1,0))"));

  4156.             custom_defines.push_back(std::make_pair("image2d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,ivec2(p.x*2,p.y),0),texelFetch(tex,ivec2(p.x*2+1,p.y),0))"));

  4157.             custom_defines.push_back(std::make_pair("image3d_ld8(tex,p)", "f16mat2x4(texelFetch(tex,ivec3(p.x*2,p.y,p.z),0),texelFetch(tex,ivec3(p.x*2+1,p.y,p.z),0))"));

  4158.             custom_defines.push_back(std::make_pair("image1d_st8(img,p,v)", "{imageStore(img,(p)*2,v[0]);imageStore(img,(p)*2+1,v[1]);}"));

  4159.             custom_defines.push_back(std::make_pair("image2d_st8(img,p,v)", "{imageStore(img,ivec2(p.x*2,p.y),v[0]);imageStore(img,ivec2(p.x*2+1,p.y),v[1]);}"));

  4160.             custom_defines.push_back(std::make_pair("image3d_st8(img,p,v)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),v[0]);imageStore(img,ivec3(p.x*2+1,p.y,p.z),v[1]);}"));

  4161.             custom_defines.push_back(std::make_pair("image1d_cp8(img,p,tex,sp)", "{imageStore(img,(p)*2,texelFetch(tex,sp*2,0));imageStore(img,(p)*2+1,texelFetch(tex,sp*2+1,0));}"));

  4162.             custom_defines.push_back(std::make_pair("image2d_cp8(img,p,tex,sp)", "{imageStore(img,ivec2(p.x*2,p.y),texelFetch(tex,ivec2(sp.x*2,sp.y),0));imageStore(img,ivec2(p.x*2+1,p.y),texelFetch(tex,ivec2(sp.x*2+1,sp.y),0));}"));

  4163.             custom_defines.push_back(std::make_pair("image3d_cp8(img,p,tex,sp)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),texelFetch(tex,ivec3(sp.x*2,sp.y,sp.z),0));imageStore(img,ivec3(p.x*2+1,p.y,p.z),texelFetch(tex,ivec3(sp.x*2+1,sp.y,sp.z),0));}"));

  4164.         }

  4165.         else if (opt.use_fp16_storage)

  4166.         {

  4167.             custom_defines.push_back(std::make_pair("image1d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4168.             custom_defines.push_back(std::make_pair("image2d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4169.             custom_defines.push_back(std::make_pair("image3d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4170.             custom_defines.push_back(std::make_pair("image1d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4171.             custom_defines.push_back(std::make_pair("image2d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4172.             custom_defines.push_back(std::make_pair("image3d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4173.             custom_defines.push_back(std::make_pair("image1d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4174.             custom_defines.push_back(std::make_pair("image2d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4175.             custom_defines.push_back(std::make_pair("image3d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4176.             custom_defines.push_back(std::make_pair("image1d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4177.             custom_defines.push_back(std::make_pair("image2d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4178.             custom_defines.push_back(std::make_pair("image3d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4179.             custom_defines.push_back(std::make_pair("image1d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4180.             custom_defines.push_back(std::make_pair("image2d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4181.             custom_defines.push_back(std::make_pair("image3d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4182.             custom_defines.push_back(std::make_pair("image1d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4183.             custom_defines.push_back(std::make_pair("image2d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4184.             custom_defines.push_back(std::make_pair("image3d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4185.             custom_defines.push_back(std::make_pair("image1d_ld8(tex,p)", "mat2x4(texelFetch(tex,(p)*2,0),texelFetch(tex,(p)*2+1,0))"));

  4186.             custom_defines.push_back(std::make_pair("image2d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec2(p.x*2,p.y),0),texelFetch(tex,ivec2(p.x*2+1,p.y),0))"));

  4187.             custom_defines.push_back(std::make_pair("image3d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec3(p.x*2,p.y,p.z),0),texelFetch(tex,ivec3(p.x*2+1,p.y,p.z),0))"));

  4188.             custom_defines.push_back(std::make_pair("image1d_st8(img,p,v)", "{imageStore(img,(p)*2,v[0]);imageStore(img,(p)*2+1,v[1]);}"));

  4189.             custom_defines.push_back(std::make_pair("image2d_st8(img,p,v)", "{imageStore(img,ivec2(p.x*2,p.y),v[0]);imageStore(img,ivec2(p.x*2+1,p.y),v[1]);}"));

  4190.             custom_defines.push_back(std::make_pair("image3d_st8(img,p,v)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),v[0]);imageStore(img,ivec3(p.x*2+1,p.y,p.z),v[1]);}"));

  4191.             custom_defines.push_back(std::make_pair("image1d_cp8(img,p,tex,sp)", "{imageStore(img,(p)*2,texelFetch(tex,sp*2,0));imageStore(img,(p)*2+1,texelFetch(tex,sp*2+1,0));}"));

  4192.             custom_defines.push_back(std::make_pair("image2d_cp8(img,p,tex,sp)", "{imageStore(img,ivec2(p.x*2,p.y),texelFetch(tex,ivec2(sp.x*2,sp.y),0));imageStore(img,ivec2(p.x*2+1,p.y),texelFetch(tex,ivec2(sp.x*2+1,sp.y),0));}"));

  4193.             custom_defines.push_back(std::make_pair("image3d_cp8(img,p,tex,sp)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),texelFetch(tex,ivec3(sp.x*2,sp.y,sp.z),0));imageStore(img,ivec3(p.x*2+1,p.y,p.z),texelFetch(tex,ivec3(sp.x*2+1,sp.y,sp.z),0));}"));

  4194.         }

  4195.         else if (opt.use_fp16_packed)

  4196.         {

  4197.             custom_defines.push_back(std::make_pair("image1d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4198.             custom_defines.push_back(std::make_pair("image2d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4199.             custom_defines.push_back(std::make_pair("image3d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4200.             custom_defines.push_back(std::make_pair("image1d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4201.             custom_defines.push_back(std::make_pair("image2d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4202.             custom_defines.push_back(std::make_pair("image3d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4203.             custom_defines.push_back(std::make_pair("image1d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4204.             custom_defines.push_back(std::make_pair("image2d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4205.             custom_defines.push_back(std::make_pair("image3d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4206.             custom_defines.push_back(std::make_pair("image1d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4207.             custom_defines.push_back(std::make_pair("image2d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4208.             custom_defines.push_back(std::make_pair("image3d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4209.             custom_defines.push_back(std::make_pair("image1d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4210.             custom_defines.push_back(std::make_pair("image2d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4211.             custom_defines.push_back(std::make_pair("image3d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4212.             custom_defines.push_back(std::make_pair("image1d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4213.             custom_defines.push_back(std::make_pair("image2d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4214.             custom_defines.push_back(std::make_pair("image3d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4215.             custom_defines.push_back(std::make_pair("image1d_ld8(tex,p)", "mat2x4(texelFetch(tex,(p)*2,0),texelFetch(tex,(p)*2+1,0))"));

  4216.             custom_defines.push_back(std::make_pair("image2d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec2(p.x*2,p.y),0),texelFetch(tex,ivec2(p.x*2+1,p.y),0))"));

  4217.             custom_defines.push_back(std::make_pair("image3d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec3(p.x*2,p.y,p.z),0),texelFetch(tex,ivec3(p.x*2+1,p.y,p.z),0))"));

  4218.             custom_defines.push_back(std::make_pair("image1d_st8(img,p,v)", "{imageStore(img,(p)*2,v[0]);imageStore(img,(p)*2+1,v[1]);}"));

  4219.             custom_defines.push_back(std::make_pair("image2d_st8(img,p,v)", "{imageStore(img,ivec2(p.x*2,p.y),v[0]);imageStore(img,ivec2(p.x*2+1,p.y),v[1]);}"));

  4220.             custom_defines.push_back(std::make_pair("image3d_st8(img,p,v)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),v[0]);imageStore(img,ivec3(p.x*2+1,p.y,p.z),v[1]);}"));

  4221.             custom_defines.push_back(std::make_pair("image1d_cp8(img,p,tex,sp)", "{imageStore(img,(p)*2,texelFetch(tex,sp*2,0));imageStore(img,(p)*2+1,texelFetch(tex,sp*2+1,0));}"));

  4222.             custom_defines.push_back(std::make_pair("image2d_cp8(img,p,tex,sp)", "{imageStore(img,ivec2(p.x*2,p.y),texelFetch(tex,ivec2(sp.x*2,sp.y),0));imageStore(img,ivec2(p.x*2+1,p.y),texelFetch(tex,ivec2(sp.x*2+1,sp.y),0));}"));

  4223.             custom_defines.push_back(std::make_pair("image3d_cp8(img,p,tex,sp)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),texelFetch(tex,ivec3(sp.x*2,sp.y,sp.z),0));imageStore(img,ivec3(p.x*2+1,p.y,p.z),texelFetch(tex,ivec3(sp.x*2+1,sp.y,sp.z),0));}"));

  4224.         }

  4225.         else

  4226.         {

  4227.             custom_defines.push_back(std::make_pair("image1d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4228.             custom_defines.push_back(std::make_pair("image2d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4229.             custom_defines.push_back(std::make_pair("image3d_ld1(tex,p)", "texelFetch(tex,p,0).r"));

  4230.             custom_defines.push_back(std::make_pair("image1d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4231.             custom_defines.push_back(std::make_pair("image2d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4232.             custom_defines.push_back(std::make_pair("image3d_st1(img,p,v)", "{vec4 _v;_v.r=v;imageStore(img,p,_v);}"));

  4233.             custom_defines.push_back(std::make_pair("image1d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4234.             custom_defines.push_back(std::make_pair("image2d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4235.             custom_defines.push_back(std::make_pair("image3d_cp1(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4236.             custom_defines.push_back(std::make_pair("image1d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4237.             custom_defines.push_back(std::make_pair("image2d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4238.             custom_defines.push_back(std::make_pair("image3d_ld4(tex,p)", "texelFetch(tex,p,0)"));

  4239.             custom_defines.push_back(std::make_pair("image1d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4240.             custom_defines.push_back(std::make_pair("image2d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4241.             custom_defines.push_back(std::make_pair("image3d_st4(img,p,v)", "{imageStore(img,p,v);}"));

  4242.             custom_defines.push_back(std::make_pair("image1d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4243.             custom_defines.push_back(std::make_pair("image2d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4244.             custom_defines.push_back(std::make_pair("image3d_cp4(img,p,tex,sp)", "{imageStore(img,p,texelFetch(tex,sp,0));}"));

  4245.             custom_defines.push_back(std::make_pair("image1d_ld8(tex,p)", "mat2x4(texelFetch(tex,(p)*2,0),texelFetch(tex,(p)*2+1,0))"));

  4246.             custom_defines.push_back(std::make_pair("image2d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec2(p.x*2,p.y),0),texelFetch(tex,ivec2(p.x*2+1,p.y),0))"));

  4247.             custom_defines.push_back(std::make_pair("image3d_ld8(tex,p)", "mat2x4(texelFetch(tex,ivec3(p.x*2,p.y,p.z),0),texelFetch(tex,ivec3(p.x*2+1,p.y,p.z),0))"));

  4248.             custom_defines.push_back(std::make_pair("image1d_st8(img,p,v)", "{imageStore(img,(p)*2,v[0]);imageStore(img,(p)*2+1,v[1]);}"));

  4249.             custom_defines.push_back(std::make_pair("image2d_st8(img,p,v)", "{imageStore(img,ivec2(p.x*2,p.y),v[0]);imageStore(img,ivec2(p.x*2+1,p.y),v[1]);}"));

  4250.             custom_defines.push_back(std::make_pair("image3d_st8(img,p,v)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),v[0]);imageStore(img,ivec3(p.x*2+1,p.y,p.z),v[1]);}"));

  4251.             custom_defines.push_back(std::make_pair("image1d_cp8(img,p,tex,sp)", "{imageStore(img,(p)*2,texelFetch(tex,sp*2,0));imageStore(img,(p)*2+1,texelFetch(tex,sp*2+1,0));}"));

  4252.             custom_defines.push_back(std::make_pair("image2d_cp8(img,p,tex,sp)", "{imageStore(img,ivec2(p.x*2,p.y),texelFetch(tex,ivec2(sp.x*2,sp.y),0));imageStore(img,ivec2(p.x*2+1,p.y),texelFetch(tex,ivec2(sp.x*2+1,sp.y),0));}"));

  4253.             custom_defines.push_back(std::make_pair("image3d_cp8(img,p,tex,sp)", "{imageStore(img,ivec3(p.x*2,p.y,p.z),texelFetch(tex,ivec3(sp.x*2,sp.y,sp.z),0));imageStore(img,ivec3(p.x*2+1,p.y,p.z),texelFetch(tex,ivec3(sp.x*2+1,sp.y,sp.z),0));}"));

  4254.         }

  4255.     }

  4256. 

  4257.     custom_defines.push_back(std::make_pair("psc(x)", "(x==0?p.x:x)"));

  4258. 

  4259.     if (opt.use_fp16_storage)

  4260.     {

  4261.         custom_defines.push_back(std::make_pair("NCNN_fp16_storage", "1"));

  4262.     }

  4263.     else if (opt.use_fp16_packed)

  4264.     {

  4265.         custom_defines.push_back(std::make_pair("NCNN_fp16_packed", "1"));

  4266.     }

  4267. 

  4268.     if (opt.use_fp16_uniform)

  4269.     {

  4270.         custom_defines.push_back(std::make_pair("NCNN_fp16_uniform", "1"));

  4271.     }

  4272. 

  4273.     if (opt.use_fp16_arithmetic)

  4274.     {

  4275.         custom_defines.push_back(std::make_pair("NCNN_fp16_arithmetic", "1"));

  4276.     }

  4277. 

  4278.     if (opt.use_int8_storage)

  4279.     {

  4280.         custom_defines.push_back(std::make_pair("NCNN_int8_storage", "1"));

  4281.     }

  4282.     else if (opt.use_int8_packed)

  4283.     {

  4284.         custom_defines.push_back(std::make_pair("NCNN_int8_packed", "1"));

  4285.     }

  4286. 

  4287.     if (opt.use_int8_uniform)

  4288.     {

  4289.         custom_defines.push_back(std::make_pair("NCNN_int8_uniform", "1"));

  4290.     }

  4291. 

  4292.     if (opt.use_int8_arithmetic)

  4293.     {

  4294.         custom_defines.push_back(std::make_pair("NCNN_int8_arithmetic", "1"));

  4295.     }

  4296. 

  4297.     if (opt.use_image_storage)

  4298.     {

  4299.         custom_defines.push_back(std::make_pair("NCNN_image_shader", "1"));

  4300.     }

  4301. 

  4302.     if (opt.use_subgroup_basic)

  4303.     {

  4304.         custom_defines.push_back(std::make_pair("NCNN_subgroup_basic", "1"));

  4305. 

  4306.         if (opt.use_subgroup_vote)

  4307.         {

  4308.             custom_defines.push_back(std::make_pair("NCNN_subgroup_vote", "1"));

  4309.         }

  4310.         if (opt.use_subgroup_ballot)

  4311.         {

  4312.             custom_defines.push_back(std::make_pair("NCNN_subgroup_ballot", "1"));

  4313.         }

  4314.         if (opt.use_subgroup_shuffle)

  4315.         {

  4316.             custom_defines.push_back(std::make_pair("NCNN_subgroup_shuffle", "1"));

  4317.         }

  4318.     }

  4319. 

  4320.     if (opt.use_shader_local_memory)

  4321.     {

  4322.         custom_defines.push_back(std::make_pair("NCNN_shader_local_memory", "1"));

  4323.     }

  4324. 

  4325. #if __APPLE__

  4326.     custom_defines.push_back(std::make_pair("NCNN_moltenvk", "1"));

  4327. #endif

  4328. 

  4329.     std::string preamble;

  4330.     std::vector<std::string> processes;

  4331. 

  4332.     processes.resize(custom_defines.size());

  4333.     for (size_t i = 0; i < custom_defines.size(); i++)

  4334.     {

  4335.         const char* key = custom_defines[i].first;

  4336.         const char* def = custom_defines[i].second;

  4337. 

  4338.         preamble += std::string("#define ") + key + " " + def + "\n";

  4339.         processes[i] = std::string("define-macro ") + key + "=" + def;

  4340.     }

  4341. 

  4342.     bool compile_success = true;

  4343. 

  4344.     {

  4345.         glslang::TShader s(EShLangCompute);

  4346. 

  4347.         s.setStringsWithLengths(&comp_data, &comp_data_size, 1);

  4348. 

  4349.         s.setPreamble(preamble.c_str());

  4350.         s.addProcesses(processes);

  4351.         s.setEntryPoint("main");

  4352.         s.setSourceEntryPoint("main");

  4353. 

  4354.         s.setEnvInput(glslang::EShSourceGlsl, EShLangCompute, glslang::EShClientVulkan, 1);

  4355. 

  4356.         if (opt.use_subgroup_basic || opt.use_cooperative_matrix)

  4357.         {

  4358.             // subgroup / cooperative_matrix need vulkan-1.1 and spirv-1.3

  4359.             s.setEnvClient(glslang::EShClientVulkan, glslang::EShTargetVulkan_1_1);

  4360.             s.setEnvTarget(glslang::EshTargetSpv, glslang::EShTargetSpv_1_3);

  4361.         }

  4362.         else

  4363.         {

  4364.             s.setEnvClient(glslang::EShClientVulkan, glslang::EShTargetVulkan_1_0);

  4365.             s.setEnvTarget(glslang::EshTargetSpv, glslang::EShTargetSpv_1_0);

  4366.         }

  4367. 

  4368.         TBuiltInResource resources = get_default_TBuiltInResource();

  4369. 

  4370.         VulkanShaderIncluder includer;

  4371. 

  4372.         bool pr = s.parse(&resources, 100, ENoProfile, false, false, EShMsgDefault, includer);

  4373.         if (!pr)

  4374.         {

  4375.             NCNN_LOGE("compile spir-v module failed");

  4376.             NCNN_LOGE("%s", s.getInfoLog());

  4377.             NCNN_LOGE("%s", s.getInfoDebugLog());

  4378. 

  4379.             compile_success = false;

  4380.         }

  4381.         else

  4382.         {

  4383.             glslang::TIntermediate* ir = s.getIntermediate();

  4384.             glslang::GlslangToSpv(*ir, spirv);

  4385.         }

  4386.     }

  4387. 

  4388.     return compile_success ? 0 : -1;

  4389. }

  4390. 

  4391. int compile_spirv_module(int shader_type_index, const Option& opt, std::vector<uint32_t>& spirv)

  4392. {

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

  4394.     {

  4395.         NCNN_LOGE("no such shader module %d", shader_type_index);

  4396.         return -1;

  4397.     }

  4398. 

  4399.     const char* comp_data = layer_shader_registry[shader_type_index].comp_data;

  4400.     int comp_data_size = layer_shader_registry[shader_type_index].comp_data_size;

  4401. 

  4402.     return compile_spirv_module(comp_data, comp_data_size, opt, spirv);

  4403. }

  4404. 

  4405. int resolve_shader_info(const uint32_t* spv_data, size_t spv_data_size, ShaderInfo& shader_info)

  4406. {

  4407.     shader_info.specialization_count = 0;

  4408.     shader_info.binding_count = 0;

  4409.     shader_info.push_constant_count = 0;

  4410. 

  4411.     uint32_t parameter_id = -233;

  4412. 

  4413.     int specialization_count = 0;

  4414.     int binding_count = 0;

  4415.     int push_constant_count = 0;

  4416. 

  4417.     // id -> binding_type

  4418.     std::vector<int> id_types;

  4419. 

  4420.     // binding_id -> binding_type

  4421.     std::vector<int> binding_types;

  4422. 

  4423.     const uint32_t* p = spv_data;

  4424. 

  4425.     int bound = p[3];

  4426. 

  4427.     id_types.resize(bound);

  4428. 

  4429.     // skip magic version generator bound schema

  4430.     p += 5;

  4431. 

  4432.     // foreach op

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

  4434.     {

  4435.         uint32_t opcode = p[0];

  4436. 

  4437.         uint16_t wordcount = opcode >> 16;

  4438.         uint16_t op = opcode & 0xffff;

  4439. 

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

  4441.         {

  4442.             uint32_t id = p[1];

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

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

  4445.             {

  4446.                 parameter_id = id;

  4447.             }

  4448.         }

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

  4450.         {

  4451.             uint32_t id = p[1];

  4452.             if (id == parameter_id)

  4453.             {

  4454.                 push_constant_count++;

  4455.             }

  4456.         }

  4457.         else if (op == 25) // OpTypeImage

  4458.         {

  4459.             uint32_t id = p[1];

  4460.             id_types[id] = 2;

  4461.         }

  4462.         else if (op == 27) // OpTypeSampledImage

  4463.         {

  4464.             uint32_t id = p[1];

  4465.             id_types[id] = 3;

  4466.         }

  4467.         else if (op == 32) // OpTypePointer

  4468.         {

  4469.             uint32_t id = p[1];

  4470.             uint32_t storage_class = p[2];

  4471.             uint32_t type = p[3];

  4472.             if (storage_class == 0) // UniformConstant

  4473.             {

  4474.                 id_types[id] = id_types[type];

  4475.             }

  4476.             if (storage_class == 2) // Uniform

  4477.             {

  4478.                 id_types[id] = id_types[type];

  4479.             }

  4480.             if (storage_class == 12) // StorageBuffer

  4481.             {

  4482.                 id_types[type] = 1;

  4483.                 id_types[id] = id_types[type];

  4484.             }

  4485.         }

  4486.         else if (op == 59) // OpVariable

  4487.         {

  4488.             uint32_t id = p[1];

  4489.             uint32_t var_id = p[2];

  4490.             uint32_t storage_class = p[3];

  4491.             if (storage_class == 0) // UniformConstant

  4492.             {

  4493.                 id_types[var_id] = id_types[id];

  4494.             }

  4495.             if (storage_class == 2) // Uniform

  4496.             {

  4497.                 id_types[var_id] = id_types[id];

  4498.             }

  4499.             if (storage_class == 12) // StorageBuffer

  4500.             {

  4501.                 id_types[var_id] = id_types[id];

  4502.             }

  4503.         }

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

  4505.         {

  4506.             uint32_t id = p[1];

  4507.             uint32_t decoration = p[2];

  4508.             uint32_t binding_id = p[3];

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

  4510.             {

  4511.                 specialization_count++;

  4512.             }

  4513.             if (decoration == 3) // BufferBlock

  4514.             {

  4515.                 id_types[id] = 1;

  4516.             }

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

  4518.             {

  4519.                 binding_count = std::max(binding_count, (int)binding_id + 1);

  4520. 

  4521.                 binding_types.resize(binding_count);

  4522.                 binding_types[binding_id] = id;

  4523.             }

  4524.         }

  4525. 

  4526.         p += wordcount;

  4527.     }

  4528. 

  4529.     if (binding_count > 16)

  4530.     {

  4531.         NCNN_LOGE("too many binding %d", binding_count);

  4532.         return -1;

  4533.     }

  4534. 

  4535.     shader_info.specialization_count = specialization_count;

  4536.     shader_info.binding_count = binding_count;

  4537.     shader_info.push_constant_count = push_constant_count;

  4538. 

  4539.     // resolve binding_types

  4540.     for (int i = 0; i < binding_count; i++)

  4541.     {

  4542.         shader_info.binding_types[i] = id_types[binding_types[i]];

  4543.     }

  4544. 

  4545.     return 0;

  4546. }

  4547. 

  4548. } // namespace ncnn

  4549. 

  4550. #endif // NCNN_VULKAN