1 /* 2 * Copyright (c) Meta Platforms, Inc. and affiliates. 3 * All rights reserved. 4 * 5 * This source code is licensed under the BSD-style license found in the 6 * LICENSE file in the root directory of this source tree. 7 */ 8 9 #pragma once 10 11 // @lint-ignore-every CLANGTIDY facebook-hte-BadMemberName 12 13 #include <executorch/backends/vulkan/runtime/vk_api/vk_api.h> 14 15 #include <executorch/backends/vulkan/runtime/vk_api/Adapter.h> 16 #include <executorch/backends/vulkan/runtime/vk_api/Command.h> 17 #include <executorch/backends/vulkan/runtime/vk_api/Pipeline.h> 18 19 #include <cstdint> 20 #include <functional> 21 22 #ifndef VULKAN_QUERY_POOL_SIZE 23 #define VULKAN_QUERY_POOL_SIZE 4096u 24 #endif 25 26 namespace vkcompute { 27 namespace vkapi { 28 29 struct QueryPoolConfig final { 30 uint32_t max_query_count = VULKAN_QUERY_POOL_SIZE; 31 uint32_t initial_reserve_size = 256u; 32 }; 33 34 struct ShaderDuration final { 35 uint32_t idx; 36 37 // Execution Properties 38 uint32_t dispatch_id; 39 std::string kernel_name; 40 VkExtent3D global_workgroup_size; 41 VkExtent3D local_workgroup_size; 42 43 // Query indexes 44 uint32_t start_query_idx; 45 uint32_t end_query_idx; 46 47 // Timings 48 uint64_t start_time_ns; 49 uint64_t end_time_ns; 50 uint64_t execution_duration_ns; 51 }; 52 53 class QueryPool final { 54 // Configuration 55 QueryPoolConfig config_; 56 uint64_t ns_per_tick_; 57 58 // Vulkan handles 59 VkDevice device_; 60 VkQueryPool querypool_; 61 62 // Internal State 63 uint32_t num_queries_; 64 std::vector<ShaderDuration> shader_durations_; 65 66 std::mutex mutex_; 67 68 public: 69 explicit QueryPool(const QueryPoolConfig&, const Adapter* adapter_p); 70 71 QueryPool(const QueryPool&) = delete; 72 QueryPool& operator=(const QueryPool&) = delete; 73 74 QueryPool(QueryPool&&) = delete; 75 QueryPool& operator=(QueryPool&&) = delete; 76 77 ~QueryPool(); 78 79 void initialize(const Adapter* adapter_p); 80 81 private: 82 size_t write_timestamp(const CommandBuffer&); 83 84 public: 85 void reset_querypool(const CommandBuffer&); 86 87 void reset_state(); 88 89 void shader_profile_begin( 90 const CommandBuffer&, 91 const uint32_t, 92 const std::string&, 93 const VkExtent3D, 94 const VkExtent3D); 95 96 void shader_profile_end(const CommandBuffer&); 97 98 void extract_results(); 99 100 std::vector<std::tuple<std::string, uint32_t, uint64_t, uint64_t>> 101 get_shader_timestamp_data(); 102 std::string generate_string_report(); 103 void print_results(); 104 unsigned long get_total_shader_ns(std::string kernel_name); 105 unsigned long get_mean_shader_ns(std::string kernel_name); 106 107 operator bool() const { 108 return querypool_ != VK_NULL_HANDLE; 109 } 110 }; 111 112 } // namespace vkapi 113 } // namespace vkcompute 114