[ET-VK] Print op breakdown in tsv format

backends/vulkan/runtime/VulkanBackend.cpp

@@ -334,9 +334,6 @@ class GraphBuilder {
     }
 
     // Parse the operators
-    uint32_t last_prepack_node_ct = 0;
-    uint32_t last_execute_node_ct = 0;
-
     for (OpCallPtr op_call : *(flatbuffer_->chain())) {
       std::string op_name = op_call->name()->str();
       ET_CHECK_MSG(VK_HAS_OP(op_name), "Missing operator: %s", op_name.c_str());
@@ -351,22 +348,6 @@ class GraphBuilder {
 
       auto vkFn = VK_GET_OP_FN(op_name);
       vkFn(*compute_graph_, args);
-      if (compute_graph_->graphconfig().enable_querypool) {
-        for (uint32_t idx_prepack = last_prepack_node_ct;
-             idx_prepack < compute_graph_->prepack_nodes().size();
-             idx_prepack++) {
-          compute_graph_->prepack_nodes()[idx_prepack]->set_node_id(
-              op_call->node_id());
-        }
-        for (uint32_t idx_execute = last_execute_node_ct;
-             idx_execute < compute_graph_->execute_nodes().size();
-             idx_execute++) {
-          compute_graph_->execute_nodes()[idx_execute]->set_node_id(
-              op_call->node_id());
-        }
-        last_prepack_node_ct = compute_graph_->prepack_nodes().size();
-        last_execute_node_ct = compute_graph_->execute_nodes().size();
-      }
     }
 
     // Parse the outputs, which will be mostly tensors.  For some reason,
@@ -379,6 +360,15 @@ class GraphBuilder {
         compute_graph_->set_output_tensor(ref);
       }
     }
+
+    if (compute_graph_->graphconfig().enable_querypool) {
+      for (uint32_t i = 0; i < compute_graph_->prepack_nodes().size(); ++i) {
+        compute_graph_->prepack_nodes()[i]->set_node_id(i);
+      }
+      for (uint32_t i = 0; i < compute_graph_->execute_nodes().size(); ++i) {
+        compute_graph_->execute_nodes()[i]->set_node_id(i);
+      }
+    }
   }
 };
 

backends/vulkan/runtime/vk_api/QueryPool.cpp

@@ -233,9 +233,44 @@ std::string QueryPool::generate_string_report() {
   return ss.str();
 }
 
-void QueryPool::print_results() {
+std::string QueryPool::generate_tsv_string_report() {
+  std::lock_guard<std::mutex> lock(mutex_);
+
+  std::stringstream ss;
+
+  ss << "Kernel Name\t";
+  ss << "Global Workgroup Size\t";
+  ss << "Local Workgroup Size\t";
+  ss << "Duration (ns)\t";
+  ss << std::endl;
+
+  ss << "===========\t";
+  ss << "=====================\t";
+  ss << "====================\t";
+  ss << "=============\t";
+  ss << std::endl;
+
+  for (ShaderDuration& entry : shader_durations_) {
+    std::chrono::duration<size_t, std::nano> exec_duration_ns(
+        entry.execution_duration_ns);
+
+    ss << entry.kernel_name << "\t";
+    ss << stringize(entry.global_workgroup_size) << "\t";
+    ss << stringize(entry.local_workgroup_size) << "\t";
+    ss << exec_duration_ns.count() << "\t";
+    ss << std::endl;
+  }
+
+  return ss.str();
+}
+
+void QueryPool::print_results(const bool tsv_format) {
   EARLY_RETURN_IF_UNINITIALIZED();
-  std::cout << generate_string_report() << std::endl;
+  if (tsv_format) {
+    std::cout << generate_tsv_string_report() << std::endl;
+  } else {
+    std::cout << generate_string_report() << std::endl;
+  }
 }
 
 unsigned long QueryPool::get_total_shader_ns(std::string kernel_name) {

backends/vulkan/runtime/vk_api/QueryPool.h

@@ -99,14 +99,17 @@ class QueryPool final {
 
   std::vector<std::tuple<std::string, uint32_t, uint64_t, uint64_t>>
   get_shader_timestamp_data();
-  std::string generate_string_report();
-  void print_results();
+  void print_results(const bool tsv_format = false);
   unsigned long get_total_shader_ns(std::string kernel_name);
   unsigned long get_mean_shader_ns(std::string kernel_name);
 
   operator bool() const {
     return querypool_ != VK_NULL_HANDLE;
   }
+
+ private:
+  std::string generate_string_report();
+  std::string generate_tsv_string_report();
 };
 
 } // namespace vkapi