coding style formatting

src/core/context.h

@@ -19,7 +19,9 @@ namespace rocprofiler {
 struct rocprofiler_contex_t;
 class Context;
 
-inline unsigned align_size(unsigned size, unsigned alignment) { return ((size + alignment - 1) & ~(alignment - 1)); }
+inline unsigned align_size(unsigned size, unsigned alignment) {
+  return ((size + alignment - 1) & ~(alignment - 1));
+}
 
 // Block descriptor
 struct block_des_t {
@@ -42,9 +44,8 @@ struct block_status_t {
 };
 
 // Metrics arguments
-template <class Map>
-class MetricArgs : public xml::args_cache_t {
-  public:
+template <class Map> class MetricArgs : public xml::args_cache_t {
+ public:
   MetricArgs(const Map& map) : map_(map) {}
   bool Lookup(const std::string& name, uint64_t& result) const {
     rocprofiler_feature_t* info = NULL;
@@ -53,26 +54,29 @@ class MetricArgs : public xml::args_cache_t {
     info = it->second;
     if (info) {
       result = info->data.result_int64;
-      if (info->data.kind == ROCPROFILER_DATA_KIND_UNINIT) EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' is uninitialized");
-      if (info->data.kind != ROCPROFILER_DATA_KIND_INT64) EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' is of incompatible type, not INT64");
-    } else EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' info is NULL");
+      if (info->data.kind == ROCPROFILER_DATA_KIND_UNINIT)
+        EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' is uninitialized");
+      if (info->data.kind != ROCPROFILER_DATA_KIND_INT64)
+        EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' is of incompatible type, not INT64");
+    } else
+      EXC_RAISING(HSA_STATUS_ERROR, "var '" << name << "' info is NULL");
     return (info != NULL);
   }
-  private:
+
+ private:
   const Map& map_;
 };
 
 // Profiling group
 class Group {
-  public:
-  Group(const util::AgentInfo* agent_info, Context *context, const uint32_t& index) :
-    pmc_profile_(agent_info),
-    sqtt_profile_(agent_info),
-    n_profiles_(0),
-    refs_(1),
-    context_(context),
-    index_(index)
-  {}
+ public:
+  Group(const util::AgentInfo* agent_info, Context* context, const uint32_t& index)
+      : pmc_profile_(agent_info),
+        sqtt_profile_(agent_info),
+        n_profiles_(0),
+        refs_(1),
+        context_(context),
+        index_(index) {}
 
   void Insert(const profile_info_t& info) {
     const rocprofiler_feature_kind_t kind = info.rinfo->kind;
@@ -106,9 +110,7 @@ class Group {
     sqtt_profile_.GetProfiles(vec);
   }
 
-  void GetTraceProfiles(profile_vector_t& vec) {
-    sqtt_profile_.GetProfiles(vec);
-  }
+  void GetTraceProfiles(profile_vector_t& vec) { sqtt_profile_.GetProfiles(vec); }
 
   info_vector_t& GetInfoVector() { return info_vector_; }
   const pkt_vector_t& GetStartVector() const { return start_vector_; }
@@ -125,7 +127,7 @@ class Group {
     return refs_;
   }
 
-  private:
+ private:
   PmcProfile pmc_profile_;
   SqttProfile sqtt_profile_;
   info_vector_t info_vector_;
@@ -139,19 +141,19 @@ class Group {
 
 // Profiling context
 class Context {
-  public:
+ public:
   typedef std::mutex mutex_t;
   typedef std::map<std::string, rocprofiler_feature_t*> info_map_t;
 
-  Context(const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info, const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg) :
-    agent_(agent_info->dev_id),
-    agent_info_(agent_info),
-    queue_(queue),
-    hsa_rsrc_(&util::HsaRsrcFactory::Instance()),
-    api_(hsa_rsrc_->AqlProfileApi()),
-    handler_(handler),
-    handler_arg_(handler_arg)
-  {
+  Context(const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info,
+          const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg)
+      : agent_(agent_info->dev_id),
+        agent_info_(agent_info),
+        queue_(queue),
+        hsa_rsrc_(&util::HsaRsrcFactory::Instance()),
+        api_(hsa_rsrc_->AqlProfileApi()),
+        handler_(handler),
+        handler_arg_(handler_arg) {
     metrics_ = MetricsDict::Create(agent_info);
     if (metrics_ == NULL) EXC_RAISING(HSA_STATUS_ERROR, "MetricsDict create failed");
     Initialize(info, info_count);
@@ -163,12 +165,8 @@ class Context {
         const profile_vector_t profile_vector = GetProfiles(group_index);
         for (auto& tuple : profile_vector) {
           // Handler for stop packet completion
-          hsa_amd_signal_async_handler(
-            tuple.completion_signal,
-            HSA_SIGNAL_CONDITION_LT,
-            1,
-            Handler,
-            &set_[group_index]);
+          hsa_amd_signal_async_handler(tuple.completion_signal, HSA_SIGNAL_CONDITION_LT, 1, Handler,
+                                       &set_[group_index]);
         }
       }
     }
@@ -178,7 +176,8 @@ class Context {
     for (const auto& v : info_map_) {
       const std::string& name = v.first;
       const rocprofiler_feature_t* info = v.second;
-      if ((info->kind == ROCPROFILER_FEATURE_KIND_METRIC) && (metrics_map_.find(name) == metrics_map_.end())) {
+      if ((info->kind == ROCPROFILER_FEATURE_KIND_METRIC) &&
+          (metrics_map_.find(name) == metrics_map_.end())) {
         delete info;
       }
     }
@@ -206,14 +205,18 @@ class Context {
       const rocprofiler_feature_kind_t kind = info->kind;
       const char* name = info->name;
 
-      if (kind == ROCPROFILER_FEATURE_KIND_METRIC) { // Processing metrics features
+      if (kind == ROCPROFILER_FEATURE_KIND_METRIC) {  // Processing metrics features
         const Metric* metric = metrics_->Get(name);
-        if (metric == NULL) EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not found");
+        if (metric == NULL)
+          EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not found");
         auto ret = metrics_map_.insert({name, metric});
-        if (!ret.second) EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is registered more then once");
+        if (!ret.second)
+          EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name
+                                                         << "' is registered more then once");
 
         counters_vec_t counters_vec = metric->GetCounters();
-        if (counters_vec.empty()) EXC_RAISING(HSA_STATUS_ERROR, "bad metric '" << name << "' is empty");
+        if (counters_vec.empty())
+          EXC_RAISING(HSA_STATUS_ERROR, "bad metric '" << name << "' is empty");
 
         for (const counter_t* counter : counters_vec) {
           // For metrics expressions checking that there is no the same counter in the input metrics
@@ -238,9 +241,10 @@ class Context {
             query.agent = agent_;
             query.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
             query.events = event;
-  
+
             uint32_t block_counters;
-            hsa_status_t status = api_->hsa_ven_amd_aqlprofile_get_info(&query, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS, &block_counters);
+            hsa_status_t status = api_->hsa_ven_amd_aqlprofile_get_info(
+                &query, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS, &block_counters);
             if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "get block_counters info");
             block_status.max_counters = block_counters;
           }
@@ -254,9 +258,9 @@ class Context {
           const uint32_t group_index = block_status.group_index;
           set_[group_index].Insert(profile_info_t{event, NULL, 0, info});
         }
-      } else if (kind == ROCPROFILER_FEATURE_KIND_TRACE) { // Processing traces features
+      } else if (kind == ROCPROFILER_FEATURE_KIND_TRACE) {  // Processing traces features
         set_[0].Insert(profile_info_t{NULL, info->parameters, info->parameter_count, info});
-      } else { 
+      } else {
         EXC_RAISING(HSA_STATUS_ERROR, "bad rocprofiler feature kind (" << kind << ")");
       }
     }
@@ -269,9 +273,7 @@ class Context {
     }
   }
 
-  void Reset(const uint32_t& group_index) {
-    set_[group_index].ResetRefs();
-  }
+  void Reset(const uint32_t& group_index) { set_[group_index].ResetRefs(); }
 
   uint32_t GetGroupCount() const { return set_.size(); }
 
@@ -285,8 +287,12 @@ class Context {
     return group;
   }
 
-  const pkt_vector_t& StartPackets(const uint32_t& group_index) const { return set_[group_index].GetStartVector(); }
-  const pkt_vector_t& StopPackets(const uint32_t& group_index) const { return set_[group_index].GetStopVector(); }
+  const pkt_vector_t& StartPackets(const uint32_t& group_index) const {
+    return set_[group_index].GetStartVector();
+  }
+  const pkt_vector_t& StopPackets(const uint32_t& group_index) const {
+    return set_[group_index].GetStopVector();
+  }
 
   void Start(const uint32_t& group_index, Queue* const queue = NULL) {
     const pkt_vector_t& start_packets = StartPackets(group_index);
@@ -315,14 +321,11 @@ class Context {
     const profile_vector_t profile_vector = GetProfiles(group_index);
     for (auto& tuple : profile_vector) {
       // Wait for stop packet to complete
-      hsa_signal_wait_scacquire(
-        tuple.completion_signal,
-        HSA_SIGNAL_CONDITION_LT,
-        1,
-        (uint64_t)-1,
-        HSA_WAIT_STATE_BLOCKED);
+      hsa_signal_wait_scacquire(tuple.completion_signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
+                                HSA_WAIT_STATE_BLOCKED);
       callback_data_t callback_data{tuple.info_vector, tuple.info_vector->size(), NULL};
-      const hsa_status_t status = api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, DataCallback, &callback_data);
+      const hsa_status_t status =
+          api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, DataCallback, &callback_data);
       if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "context iterate data failed");
     }
   }
@@ -335,7 +338,8 @@ class Context {
       const xml::Expr* expr = metric->GetExpr();
       if (expr) {
         auto it = info_map_.find(name);
-        if (it == info_map_.end()) EXC_RAISING(HSA_STATUS_ERROR, "metric '" << name << "', rocprofiler info is not found");
+        if (it == info_map_.end())
+          EXC_RAISING(HSA_STATUS_ERROR, "metric '" << name << "', rocprofiler info is not found");
         rocprofiler_feature_t* info = it->second;
         info->data.result_int64 = expr->Eval(args);
         info->data.kind = ROCPROFILER_DATA_KIND_INT64;
@@ -343,16 +347,17 @@ class Context {
     }
   }
 
-  void IterateTraceData(rocprofiler_trace_data_callback_t callback, void *data) {
+  void IterateTraceData(rocprofiler_trace_data_callback_t callback, void* data) {
     profile_vector_t profile_vector;
     set_[0].GetTraceProfiles(profile_vector);
     for (auto& tuple : profile_vector) {
-      const hsa_status_t status = api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, callback, data);
+      const hsa_status_t status =
+          api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, callback, data);
       if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "context iterate data failed");
     }
   }
 
-  private:
+ private:
   // Getting profling packets
   profile_vector_t GetProfiles(const uint32_t& index) {
     profile_vector_t vec;
@@ -374,8 +379,7 @@ class Context {
   }
 
   static hsa_status_t DataCallback(hsa_ven_amd_aqlprofile_info_type_t ainfo_type,
-                                   hsa_ven_amd_aqlprofile_info_data_t* ainfo_data,
-                                   void* data) {
+                                   hsa_ven_amd_aqlprofile_info_data_t* ainfo_data, void* data) {
     hsa_status_t status = HSA_STATUS_SUCCESS;
     callback_data_t* callback_data = reinterpret_cast<callback_data_t*>(data);
     info_vector_t& info_vector = *(callback_data->info_vector);
@@ -403,7 +407,7 @@ class Context {
           char* ptr = (sample_id == 0) ? result_bytes_ptr : callback_data->ptr;
           uint64_t* header = reinterpret_cast<uint64_t*>(ptr);
           char* dest = ptr + sizeof(*header);
-  
+
           if ((dest + size) < end) {
             hsa_status_t status = hsa_memory_copy(dest, src, size);
             if (status == HSA_STATUS_SUCCESS) {
@@ -412,7 +416,8 @@ class Context {
               rinfo->data.result_bytes.instance_count = sample_id + 1;
               callback_data->ptr = dest + align_size(size, sizeof(uint64_t));
             }
-          } else status = HSA_STATUS_ERROR;
+          } else
+            status = HSA_STATUS_ERROR;
         } else {
           if (sample_id == 0) {
             rinfo->data.kind = ROCPROFILER_DATA_KIND_BYTES;
@@ -421,8 +426,10 @@ class Context {
           }
           rinfo->data.result_bytes.instance_count += 1;
         }
-      } else status = HSA_STATUS_ERROR;
-    } else status = HSA_STATUS_ERROR;
+      } else
+        status = HSA_STATUS_ERROR;
+    } else
+      status = HSA_STATUS_ERROR;
 
     return status;
   }
@@ -445,7 +452,7 @@ class Context {
   const pfn_t* api_;
   // Profile group set
   std::vector<Group> set_;
-  // Metrics dictionary 
+  // Metrics dictionary
   MetricsDict* metrics_;
   // Groups map
   std::map<block_des_t, block_status_t, lt_block_des> groups_map_;

src/core/hsa_proxy_queue.h

@@ -15,26 +15,23 @@ extern decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_
 extern decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
 
 class HsaProxyQueue : public ProxyQueue {
-  public:
+ public:
   hsa_status_t SetInterceptCB(on_submit_cb_t on_submit_cb, void* data) {
     return hsa_amd_queue_intercept_register_fn(queue_, on_submit_cb, data);
   }
 
-  void Submit(const packet_t* packet) { EXC_RAISING(HSA_STATUS_ERROR, "HsaProxyQueue::Submit() is not supported"); }
+  void Submit(const packet_t* packet) {
+    EXC_RAISING(HSA_STATUS_ERROR, "HsaProxyQueue::Submit() is not supported");
+  }
 
-  private:
-  hsa_status_t Init(
-    hsa_agent_t agent,
-    uint32_t size,
-    hsa_queue_type32_t type,
-    void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-    void *data,
-    uint32_t private_segment_size,
-    uint32_t group_segment_size,
-    hsa_queue_t **queue)
-  {
+ private:
+  hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
+                    void* data, uint32_t private_segment_size, uint32_t group_segment_size,
+                    hsa_queue_t** queue) {
     printf("HsaProxyQueue::Init()\n");
-    const auto status = hsa_amd_queue_intercept_create_fn(agent, size, type, callback, data, private_segment_size, group_segment_size, &queue_);
+    const auto status = hsa_amd_queue_intercept_create_fn(
+        agent, size, type, callback, data, private_segment_size, group_segment_size, &queue_);
     *queue = queue_;
     return status;
   }
@@ -44,6 +41,6 @@ class HsaProxyQueue : public ProxyQueue {
   hsa_queue_t* queue_;
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_HSA_PROXY_QUEUE_H
+#endif  // _SRC_CORE_HSA_PROXY_QUEUE_H

src/core/hsa_queue.h

@@ -9,7 +9,7 @@
 namespace rocprofiler {
 
 class HsaQueue : public Queue {
-  public:
+ public:
   typedef void (HsaQueue::*submit_fptr_t)(const packet_t* packet);
   enum {
     LEGACY_SLOT_SIZE_W = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_word_t),
@@ -19,9 +19,7 @@ class HsaQueue : public Queue {
     packet_word_t words[LEGACY_SLOT_SIZE_W];
   };
 
-  HsaQueue(const util::AgentInfo* agent_info, hsa_queue_t* queue) :
-    queue_(queue)
- {}
+  HsaQueue(const util::AgentInfo* agent_info, hsa_queue_t* queue) : queue_(queue) {}
 
   void Submit(const packet_t* packet) {
     // Compute the write index of queue and copy Aql packet into it
@@ -51,10 +49,10 @@ class HsaQueue : public Queue {
     hsa_signal_store_relaxed(queue_->doorbell_signal, que_idx);
   }
 
-  private:
+ private:
   hsa_queue_t* queue_;
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_HSA_QUEUE_H
+#endif  // _SRC_CORE_HSA_QUEUE_H

src/core/intercept_queue.cpp

@@ -12,4 +12,4 @@ void* InterceptQueue::on_dispatch_cb_data_ = NULL;
 const char* InterceptQueue::tool_lib_ = NULL;
 void* InterceptQueue::tool_handle_ = NULL;
 InterceptQueue::obj_map_t* InterceptQueue::obj_map_ = NULL;
-} // namespace rocprofiler
+}  // namespace rocprofiler

src/core/intercept_queue.h

@@ -17,7 +17,7 @@ extern decltype(hsa_queue_create)* hsa_queue_create_fn;
 extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
 
 class InterceptQueue {
-  public:
+ public:
   typedef std::recursive_mutex mutex_t;
   typedef std::map<uint64_t, InterceptQueue*> obj_map_t;
 
@@ -25,18 +25,15 @@ class InterceptQueue {
 
   static void SetTool(const char* tool) { tool_lib_ = tool; }
 
-  static void UnloadTool() { if (tool_handle_) dlclose(tool_handle_); }
+  static void UnloadTool() {
+    if (tool_handle_) dlclose(tool_handle_);
+  }
 
-  static hsa_status_t QueueCreate(
-    hsa_agent_t agent,
-    uint32_t size,
-    hsa_queue_type32_t type,
-    void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-    void *data,
-    uint32_t private_segment_size,
-    uint32_t group_segment_size,
-    hsa_queue_t **queue)
-  {
+  static hsa_status_t QueueCreate(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                                  void (*callback)(hsa_status_t status, hsa_queue_t* source,
+                                                   void* data),
+                                  void* data, uint32_t private_segment_size,
+                                  uint32_t group_segment_size, hsa_queue_t** queue) {
     std::lock_guard<mutex_t> lck(mutex_);
 
     hsa_status_t status = HSA_STATUS_ERROR;
@@ -53,7 +50,8 @@ class InterceptQueue {
 
     if (!obj_map_) obj_map_ = new obj_map_t;
 
-    ProxyQueue* proxy = ProxyQueue::Create(agent, size, type, callback, data, private_segment_size, group_segment_size, queue, &status);
+    ProxyQueue* proxy = ProxyQueue::Create(agent, size, type, callback, data, private_segment_size,
+                                           group_segment_size, queue, &status);
     if (status != HSA_STATUS_SUCCESS) {
       InterceptQueue* obj = new InterceptQueue(agent, proxy);
       (*obj_map_)[(uint64_t)(*queue)] = obj;
@@ -63,14 +61,14 @@ class InterceptQueue {
     return status;
   }
 
-  static hsa_status_t QueueDestroy(hsa_queue_t *queue) { 
+  static hsa_status_t QueueDestroy(hsa_queue_t* queue) {
     std::lock_guard<mutex_t> lck(mutex_);
     hsa_status_t status = HSA_STATUS_ERROR;
 
     obj_map_t::iterator it = obj_map_->find((uint64_t)queue);
     if (it != obj_map_->end()) {
       const InterceptQueue* obj = it->second;
-      delete obj; 
+      delete obj;
       obj_map_->erase(it);
       status = HSA_STATUS_SUCCESS;
     }
@@ -78,8 +76,9 @@ class InterceptQueue {
     return status;
   }
 
-  static void OnSubmitCB(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data, hsa_amd_queue_intercept_packet_writer writer) {
-    const packet_t* packets_arr =  reinterpret_cast<const packet_t*>(in_packets);
+  static void OnSubmitCB(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data,
+                         hsa_amd_queue_intercept_packet_writer writer) {
+    const packet_t* packets_arr = reinterpret_cast<const packet_t*>(in_packets);
     InterceptQueue* obj = reinterpret_cast<InterceptQueue*>(data);
     Queue* proxy = obj->proxy_;
 
@@ -89,8 +88,10 @@ class InterceptQueue {
 
       if ((GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) && (on_dispatch_cb_ != NULL)) {
         rocprofiler_group_t group = {};
-        const hsa_kernel_dispatch_packet_t* dispatch_packet = reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
-        rocprofiler_callback_data_t data = {obj->agent_info_->dev_id, dispatch_packet->kernel_object, user_que_idx};
+        const hsa_kernel_dispatch_packet_t* dispatch_packet =
+            reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
+        rocprofiler_callback_data_t data = {obj->agent_info_->dev_id,
+                                            dispatch_packet->kernel_object, user_que_idx};
         hsa_status_t status = on_dispatch_cb_(&data, on_dispatch_cb_data_, &group);
         if (status == HSA_STATUS_SUCCESS) {
           Context* context = reinterpret_cast<Context*>(group.context);
@@ -110,11 +111,11 @@ class InterceptQueue {
       }
 
       if (to_submit) {
-          if (writer != NULL) {
-            writer(packet, 1);
-          } else {
-            proxy->Submit(packet, 1);
-          }
+        if (writer != NULL) {
+          writer(packet, 1);
+        } else {
+          proxy->Submit(packet, 1);
+        }
       }
 
       packet += 1;
@@ -133,12 +134,12 @@ class InterceptQueue {
     on_dispatch_cb_data_ = NULL;
   }
 
-  private:
+ private:
   InterceptQueue(const hsa_agent_t& agent, ProxyQueue* proxy) : proxy_(proxy) {
     agent_info_ = util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
   }
   ~InterceptQueue() { ProxyQueue::Destroy(proxy_); }
-  
+
   static packet_word_t GetHeaderType(const packet_t* packet) {
     const packet_word_t* header = reinterpret_cast<const packet_word_t*>(packet);
     return (*header >> HSA_PACKET_HEADER_TYPE) & header_type_mask;
@@ -156,6 +157,6 @@ class InterceptQueue {
   const util::AgentInfo* agent_info_;
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_INTERCEPT_QUEUE_H
+#endif  // _SRC_CORE_INTERCEPT_QUEUE_H

src/core/metrics.h

@@ -25,61 +25,69 @@ struct counter_t {
 typedef std::vector<const counter_t*> counters_vec_t;
 
 class Metric {
-  public:
+ public:
   Metric(const std::string& name) : name_(name) {}
   std::string GetName() const { return name_; }
-  virtual void GetCounters(counters_vec_t &vec) const = 0;
+  virtual void GetCounters(counters_vec_t& vec) const = 0;
   counters_vec_t GetCounters() const {
     counters_vec_t counters;
     GetCounters(counters);
     return counters;
   }
   virtual const xml::Expr* GetExpr() const = 0;
-  private:
+
+ private:
   std::string name_;
 };
 
 class BaseMetric : public Metric {
-  public:
+ public:
   BaseMetric(const std::string& name, const counter_t& counter) : Metric(name), counter_(counter) {}
-  void GetCounters(counters_vec_t &vec) const { vec.push_back(&counter_); }
+  void GetCounters(counters_vec_t& vec) const { vec.push_back(&counter_); }
   const xml::Expr* GetExpr() const { return NULL; }
-  private:
+
+ private:
   const counter_t counter_;
 };
 
 class ExprMetric : public Metric {
-  public:
-  ExprMetric(const std::string& name, const counters_vec_t& counters, const xml::Expr* expr) : Metric(name), counters_(counters), expr_(expr) {}
-  void GetCounters(counters_vec_t& vec) const { vec.insert(vec.end(), counters_.begin(), counters_.end()); }
+ public:
+  ExprMetric(const std::string& name, const counters_vec_t& counters, const xml::Expr* expr)
+      : Metric(name), counters_(counters), expr_(expr) {}
+  void GetCounters(counters_vec_t& vec) const {
+    vec.insert(vec.end(), counters_.begin(), counters_.end());
+  }
   const xml::Expr* GetExpr() const { return expr_; }
-  private:
+
+ private:
   const counters_vec_t counters_;
   const xml::Expr* expr_;
 };
 
 
 class MetricsDict {
-  public:
+ public:
   typedef std::map<std::string, const Metric*> cache_t;
   typedef cache_t::const_iterator const_iterator_t;
   typedef std::map<std::string, MetricsDict*> map_t;
   typedef std::mutex mutex_t;
 
   class ExprCache : public xml::expr_cache_t {
-    public:
+   public:
     ExprCache(const cache_t* cache) : cache_(cache) {}
     bool Lookup(const std::string& name, std::string& result) const {
       bool ret = false;
       auto it = cache_->find(name);
       if (it != cache_->end()) {
         ret = true;
-        const rocprofiler::ExprMetric* expr_metric = dynamic_cast<const rocprofiler::ExprMetric*>(it->second);
+        const rocprofiler::ExprMetric* expr_metric =
+            dynamic_cast<const rocprofiler::ExprMetric*>(it->second);
         if (expr_metric) result = expr_metric->GetExpr()->GetStr();
       }
       return ret;
     }
-    private:
+
+   private:
     const cache_t* const cache_;
   };
 
@@ -98,7 +106,7 @@ class MetricsDict {
     return metric;
   }
 
-  private:
+ private:
   MetricsDict(const util::AgentInfo* agent_info) : xml_(NULL) {
     const char* xml_name = getenv("ROCP_METRICS");
     if (xml_name != NULL) {
@@ -113,23 +121,26 @@ class MetricsDict {
     auto scope_list = xml_->GetNodes("top." + std::string(scope) + ".metric");
     if (!scope_list.empty()) {
       std::cout << "  " << scope_list.size() << " " << scope << " metrics found" << std::endl;
-  
+
       for (auto node : scope_list) {
         const std::string name = node->opts["name"];
-        if (cache_.find(name) != cache_.end()) EXC_RAISING(HSA_STATUS_ERROR, "ImportMetrics: metrics redefined '" << name << "'");
+        if (cache_.find(name) != cache_.end())
+          EXC_RAISING(HSA_STATUS_ERROR, "ImportMetrics: metrics redefined '" << name << "'");
 
         const std::string expr_str = node->opts["expr"];
         if (expr_str.empty()) {
           const std::string block_name = node->opts["block"];
           const uint32_t event_id = atoi(node->opts["event"].c_str());
-  
+
           hsa_ven_amd_aqlprofile_profile_t profile;
           profile.agent = agent_info->dev_id;
           hsa_ven_amd_aqlprofile_id_query_t query = {block_name.c_str(), 0, 0};
-          hsa_status_t status = util::HsaRsrcFactory::Instance().AqlProfileApi()->
-            hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_ID, &query);
+          hsa_status_t status =
+              util::HsaRsrcFactory::Instance().AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(
+                  &profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_ID, &query);
           if (status == HSA_STATUS_SUCCESS) {
-            const hsa_ven_amd_aqlprofile_block_name_t block_id = (hsa_ven_amd_aqlprofile_block_name_t)query.id;
+            const hsa_ven_amd_aqlprofile_block_name_t block_id =
+                (hsa_ven_amd_aqlprofile_block_name_t)query.id;
             if (query.instance_count > 1) {
               for (unsigned block_index = 0; block_index < query.instance_count; ++block_index) {
                 std::ostringstream os;
@@ -142,14 +153,18 @@ class MetricsDict {
               const counter_t counter = {name, {block_id, 0, event_id}};
               cache_[name] = new BaseMetric(name, counter);
             }
-          } else AQL_EXC_RAISING(HSA_STATUS_ERROR, "ImportMetrics: bad block name '" << block_name << "'");
+          } else
+            AQL_EXC_RAISING(HSA_STATUS_ERROR, "ImportMetrics: bad block name '" << block_name
+                                                                                << "'");
         } else {
           xml::Expr* expr_obj = new xml::Expr(expr_str, new ExprCache(&cache_));
           std::cout << "    " << name << " = " << expr_obj->String() << std::endl;
           counters_vec_t counters_vec;
           for (const std::string var : expr_obj->GetVars()) {
             auto it = cache_.find(var);
-            if (it == cache_.end()) EXC_RAISING(HSA_STATUS_ERROR, "Bad metric '" << name << "', var '" << var << "' is not found");
+            if (it == cache_.end())
+              EXC_RAISING(HSA_STATUS_ERROR, "Bad metric '" << name << "', var '" << var
+                                                           << "' is not found");
             it->second->GetCounters(counters_vec);
           }
           cache_[name] = new ExprMetric(name, counters_vec, expr_obj);

src/core/profile.h

@@ -26,13 +26,13 @@ struct profile_tuple_t {
 };
 typedef std::vector<profile_tuple_t> profile_vector_t;
 
-template<class Item> class ConfigBase {};
+template <class Item> class ConfigBase {};
 
-template<> class ConfigBase<event_t> {
-  public:
-  ConfigBase(profile_t *profile) : profile_(profile) {}
+template <> class ConfigBase<event_t> {
+ public:
+  ConfigBase(profile_t* profile) : profile_(profile) {}
 
-  protected:
+ protected:
   void* Array() { return const_cast<event_t*>(profile_->events); }
   unsigned Count() const { return profile_->event_count; }
   void Set(event_t* events, const unsigned& count) {
@@ -42,11 +42,11 @@ template<> class ConfigBase<event_t> {
   profile_t* profile_;
 };
 
-template<> class ConfigBase<parameter_t> {
-  public:
-  ConfigBase(profile_t *profile) : profile_(profile) {}
+template <> class ConfigBase<parameter_t> {
+ public:
+  ConfigBase(profile_t* profile) : profile_(profile) {}
 
-  protected:
+ protected:
   void* Array() { return const_cast<parameter_t*>(profile_->parameters); }
   unsigned Count() const { return profile_->parameter_count; }
   void Set(parameter_t* parameters, const unsigned& count) {
@@ -56,23 +56,25 @@ template<> class ConfigBase<parameter_t> {
   profile_t* profile_;
 };
 
-template<class Item> 
-class Config : protected ConfigBase<Item> {
+template <class Item> class Config : protected ConfigBase<Item> {
   typedef ConfigBase<Item> Parent;
-  public:
-  Config(profile_t *profile) : Parent(profile) {}
+
+ public:
+  Config(profile_t* profile) : Parent(profile) {}
   void Insert(const Item& item) {
     auto count = Parent::Count();
     count += 1;
-    Item* array = reinterpret_cast<Item*>(realloc(const_cast<void*>(Parent::Array()), count * sizeof(Item)));
+    Item* array =
+        reinterpret_cast<Item*>(realloc(const_cast<void*>(Parent::Array()), count * sizeof(Item)));
     array[count - 1] = item;
     Parent::Set(array, count);
   }
 };
 
 class Profile {
-  public:
-  static const uint32_t LEGACY_SLOT_SIZE_PKT = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t);
+ public:
+  static const uint32_t LEGACY_SLOT_SIZE_PKT =
+      HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t);
 
   Profile(const util::AgentInfo* agent_info) : agent_info_(agent_info) {
     profile_ = {};
@@ -86,9 +88,7 @@ class Profile {
     free(const_cast<parameter_t*>(profile_.parameters));
   }
 
-  virtual void Insert(const profile_info_t& info) {
-    info_vector_.push_back(info.rinfo);
-  }
+  virtual void Insert(const profile_info_t& info) { info_vector_.push_back(info.rinfo); }
 
   hsa_status_t Finalize(pkt_vector_t& start_vector, pkt_vector_t& stop_vector) {
     hsa_status_t status = HSA_STATUS_SUCCESS;
@@ -124,10 +124,14 @@ class Profile {
 
         start_vector.insert(start_vector.end(), LEGACY_SLOT_SIZE_PKT, packet_t{});
         stop_vector.insert(stop_vector.end(), LEGACY_SLOT_SIZE_PKT, packet_t{});
-        status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4(&start, reinterpret_cast<void*>(&start_vector[start_index]));
-        if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
-        status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4(&stop, reinterpret_cast<void*>(&stop_vector[stop_index]));
-        if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
+        status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4(
+            &start, reinterpret_cast<void*>(&start_vector[start_index]));
+        if (status != HSA_STATUS_SUCCESS)
+          AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
+        status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4(
+            &stop, reinterpret_cast<void*>(&stop_vector[stop_index]));
+        if (status != HSA_STATUS_SUCCESS)
+          AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
       } else {
         start_vector.push_back(start);
         stop_vector.push_back(stop);
@@ -145,7 +149,7 @@ class Profile {
 
   bool Empty() const { return info_vector_.empty(); }
 
-  protected:
+ protected:
   virtual hsa_status_t Allocate(util::HsaRsrcFactory* rsrc) = 0;
 
   const util::AgentInfo* const agent_info_;
@@ -156,7 +160,7 @@ class Profile {
 };
 
 class PmcProfile : public Profile {
-  public:
+ public:
   PmcProfile(const util::AgentInfo* agent_info) : Profile(agent_info) {
     profile_.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
   }
@@ -167,14 +171,16 @@ class PmcProfile : public Profile {
   }
 
   hsa_status_t Allocate(util::HsaRsrcFactory* rsrc) {
-    profile_.command_buffer.ptr = rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
+    profile_.command_buffer.ptr =
+        rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
     profile_.output_buffer.ptr = rsrc->AllocateSysMemory(agent_info_, profile_.output_buffer.size);
-    return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
+    return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS
+                                                                       : HSA_STATUS_ERROR;
   }
 };
 
 class SqttProfile : public Profile {
-  public:
+ public:
   static const uint32_t output_buffer_size = 0x2000000;  // 32M
 
   SqttProfile(const util::AgentInfo* agent_info) : Profile(agent_info) {
@@ -197,9 +203,12 @@ class SqttProfile : public Profile {
 
   hsa_status_t Allocate(util::HsaRsrcFactory* rsrc) {
     profile_.output_buffer.size = output_buffer_size;
-    profile_.command_buffer.ptr = rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
-    profile_.output_buffer.ptr = rsrc->AllocateLocalMemory(agent_info_, profile_.output_buffer.size);
-    return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
+    profile_.command_buffer.ptr =
+        rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
+    profile_.output_buffer.ptr =
+        rsrc->AllocateLocalMemory(agent_info_, profile_.output_buffer.size);
+    return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS
+                                                                       : HSA_STATUS_ERROR;
   }
 };
 

src/core/proxy_queue.cpp

@@ -10,25 +10,22 @@ void ProxyQueue::HsaIntercept(HsaApiTable* table) {
   if (rocp_type_) SimpleProxyQueue::HsaIntercept(table);
 }
 
-ProxyQueue* ProxyQueue::Create(
-  hsa_agent_t agent,
-  uint32_t size,
-  hsa_queue_type32_t type,
-  void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-  void *data,
-  uint32_t private_segment_size,
-  uint32_t group_segment_size,
-  hsa_queue_t **queue,
-  hsa_status_t* status)
-{
+ProxyQueue* ProxyQueue::Create(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                               void (*callback)(hsa_status_t status, hsa_queue_t* source,
+                                                void* data),
+                               void* data, uint32_t private_segment_size,
+                               uint32_t group_segment_size, hsa_queue_t** queue,
+                               hsa_status_t* status) {
   hsa_status_t suc = HSA_STATUS_ERROR;
 #ifdef ROCP_HSA_PROXY
-  ProxyQueue* instance = (rocp_type_) ? (ProxyQueue*) new SimpleProxyQueue() : (ProxyQueue*) new HsaProxyQueue();
+  ProxyQueue* instance =
+      (rocp_type_) ? (ProxyQueue*)new SimpleProxyQueue() : (ProxyQueue*)new HsaProxyQueue();
 #else
   ProxyQueue* instance = new SimpleProxyQueue();
 #endif
   if (instance != NULL) {
-    const auto suc = instance->Init(agent, size, type, callback, data, private_segment_size, group_segment_size, queue);
+    const auto suc = instance->Init(agent, size, type, callback, data, private_segment_size,
+                                    group_segment_size, queue);
     if (suc != HSA_STATUS_SUCCESS) {
       delete instance;
       instance = NULL;
@@ -45,4 +42,4 @@ hsa_status_t ProxyQueue::Destroy(const ProxyQueue* obj) {
 }
 
 bool ProxyQueue::rocp_type_ = false;
-} // namespace rocprofiler
+}  // namespace rocprofiler

src/core/proxy_queue.h

@@ -14,10 +14,11 @@ struct HsaApiTable;
 
 namespace rocprofiler {
 typedef void (*hsa_amd_queue_intercept_packet_writer)(const void* packets, uint64_t count);
-typedef void (*on_submit_cb_t)(const void* packet, uint64_t count, uint64_t que_idx, void* data, hsa_amd_queue_intercept_packet_writer writer);
+typedef void (*on_submit_cb_t)(const void* packet, uint64_t count, uint64_t que_idx, void* data,
+                               hsa_amd_queue_intercept_packet_writer writer);
 
 class ProxyQueue : public Queue {
-  public:
+ public:
   static void InitFactory() {
 #ifdef ROCP_HSA_PROXY
     const char* type = getenv("ROCP_PROXY_QUEUE");
@@ -31,40 +32,28 @@ class ProxyQueue : public Queue {
 
   static void HsaIntercept(HsaApiTable* table);
 
-  static ProxyQueue* Create(
-    hsa_agent_t agent,
-    uint32_t size,
-    hsa_queue_type32_t type,
-    void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-    void *data,
-    uint32_t private_segment_size,
-    uint32_t group_segment_size,
-    hsa_queue_t **queue,
-    hsa_status_t* status);
+  static ProxyQueue* Create(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                            void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
+                            void* data, uint32_t private_segment_size, uint32_t group_segment_size,
+                            hsa_queue_t** queue, hsa_status_t* status);
 
   static hsa_status_t Destroy(const ProxyQueue* obj);
 
-  virtual hsa_status_t Init(
-    hsa_agent_t agent,
-    uint32_t size,
-    hsa_queue_type32_t type,
-    void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-    void *data,
-    uint32_t private_segment_size,
-    uint32_t group_segment_size,
-    hsa_queue_t **queue
-  ) = 0;
+  virtual hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                            void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
+                            void* data, uint32_t private_segment_size, uint32_t group_segment_size,
+                            hsa_queue_t** queue) = 0;
   virtual hsa_status_t Cleanup() const = 0;
   virtual hsa_status_t SetInterceptCB(on_submit_cb_t on_submit_cb, void* data) = 0;
   virtual void Submit(const packet_t* packet) = 0;
 
-  protected:
-  virtual ~ProxyQueue() {};
+ protected:
+  virtual ~ProxyQueue(){};
 
-  private:
+ private:
   static bool rocp_type_;
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_PROXY_QUEUE_H
+#endif  // _SRC_CORE_PROXY_QUEUE_H

src/core/queue.h

@@ -6,7 +6,7 @@
 namespace rocprofiler {
 
 class Queue {
-  public:
+ public:
   Queue() {}
   virtual ~Queue() {}
   virtual void Submit(const packet_t* packet) = 0;
@@ -15,6 +15,6 @@ class Queue {
   }
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_QUEUE_H
+#endif  // _SRC_CORE_QUEUE_H

src/core/rocprofiler.cpp

@@ -18,15 +18,15 @@
 #define CONSTRUCTOR_API __attribute__((constructor))
 #define DESTRUCTOR_API __attribute__((destructor))
 
-#define API_METHOD_PREFIX \
-  hsa_status_t status = HSA_STATUS_SUCCESS; \
+#define API_METHOD_PREFIX                                                                          \
+  hsa_status_t status = HSA_STATUS_SUCCESS;                                                        \
   try {
-
-#define API_METHOD_SUFFIX \
-  } catch (std::exception& e) { \
-    ERR_LOGGING(__FUNCTION__ << "(), " << e.what()); \
-    status = rocprofiler::GetExcStatus(e); \
-  } \
+#define API_METHOD_SUFFIX                                                                          \
+  }                                                                                                \
+  catch (std::exception & e) {                                                                     \
+    ERR_LOGGING(__FUNCTION__ << "(), " << e.what());                                               \
+    status = rocprofiler::GetExcStatus(e);                                                         \
+  }                                                                                                \
   return status;
 
 namespace rocprofiler {
@@ -80,12 +80,12 @@ DESTRUCTOR_API void destructor() {
 
 hsa_status_t GetExcStatus(const std::exception& e) {
   const util::exception* rocprofiler_exc_ptr = dynamic_cast<const util::exception*>(&e);
-  return (rocprofiler_exc_ptr) ? static_cast<hsa_status_t>(rocprofiler_exc_ptr->status()) : HSA_STATUS_ERROR;
+  return (rocprofiler_exc_ptr) ? static_cast<hsa_status_t>(rocprofiler_exc_ptr->status())
+                               : HSA_STATUS_ERROR;
 }
 
 util::Logger::mutex_t util::Logger::mutex_;
 util::Logger* util::Logger::instance_ = NULL;
-
 }
 
 extern "C" {
@@ -98,14 +98,9 @@ PUBLIC_API hsa_status_t rocprofiler_error_string(const char** str) {
 }
 
 // Create new profiling context
-PUBLIC_API hsa_status_t rocprofiler_open(
-    hsa_agent_t agent,
-    rocprofiler_feature_t* info,
-    uint32_t info_count,
-    rocprofiler_t** handle,
-    uint32_t mode,
-    rocprofiler_properties_t* properties)
-{
+PUBLIC_API hsa_status_t rocprofiler_open(hsa_agent_t agent, rocprofiler_feature_t* info,
+                                         uint32_t info_count, rocprofiler_t** handle, uint32_t mode,
+                                         rocprofiler_properties_t* properties) {
   API_METHOD_PREFIX
   rocprofiler::util::HsaRsrcFactory* hsa_rsrc = &rocprofiler::util::HsaRsrcFactory::Instance();
   const rocprofiler::util::AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(agent);
@@ -117,7 +112,8 @@ PUBLIC_API hsa_status_t rocprofiler_open(
   if (mode != 0) {
     if (mode & ROCPROFILER_MODE_STANDALONE) {
       if (mode & ROCPROFILER_MODE_CREATEQUEUE) {
-        if (hsa_rsrc->CreateQueue(agent_info, properties->queue_depth, &(properties->queue)) == false) {
+        if (hsa_rsrc->CreateQueue(agent_info, properties->queue_depth, &(properties->queue)) ==
+            false) {
           EXC_RAISING(HSA_STATUS_ERROR, "CreateQueue() failed");
         }
       }
@@ -127,13 +123,13 @@ PUBLIC_API hsa_status_t rocprofiler_open(
     }
   }
 
-  *handle = new rocprofiler::Context(agent_info, queue, info, info_count, properties->handler, properties->handler_arg);
+  *handle = new rocprofiler::Context(agent_info, queue, info, info_count, properties->handler,
+                                     properties->handler_arg);
   API_METHOD_SUFFIX
 }
 
 // Delete profiling info
-PUBLIC_API hsa_status_t rocprofiler_close(rocprofiler_t* handle)
-{
+PUBLIC_API hsa_status_t rocprofiler_close(rocprofiler_t* handle) {
   API_METHOD_PREFIX
   rocprofiler::Context* context = reinterpret_cast<rocprofiler::Context*>(handle);
   if (context) delete context;
@@ -141,8 +137,7 @@ PUBLIC_API hsa_status_t rocprofiler_close(rocprofiler_t* handle)
 }
 
 // Reset context
-PUBLIC_API hsa_status_t rocprofiler_reset(rocprofiler_t* handle, uint32_t group_index)
-{
+PUBLIC_API hsa_status_t rocprofiler_reset(rocprofiler_t* handle, uint32_t group_index) {
   API_METHOD_PREFIX
   rocprofiler::Context* context = reinterpret_cast<rocprofiler::Context*>(handle);
   context->Reset(group_index);
@@ -150,7 +145,8 @@ PUBLIC_API hsa_status_t rocprofiler_reset(rocprofiler_t* handle, uint32_t group_
 }
 
 // Get profiling group count
-PUBLIC_API hsa_status_t rocprofiler_group_count(const rocprofiler_t* handle, uint32_t* group_count) {
+PUBLIC_API hsa_status_t rocprofiler_group_count(const rocprofiler_t* handle,
+                                                uint32_t* group_count) {
   API_METHOD_PREFIX
   const rocprofiler::Context* context = reinterpret_cast<const rocprofiler::Context*>(handle);
   *group_count = context->GetGroupCount();
@@ -158,7 +154,8 @@ PUBLIC_API hsa_status_t rocprofiler_group_count(const rocprofiler_t* handle, uin
 }
 
 // Get profiling group for a given group index
-PUBLIC_API hsa_status_t rocprofiler_get_group(rocprofiler_t* handle, uint32_t group_index, rocprofiler_group_t* group) {
+PUBLIC_API hsa_status_t rocprofiler_get_group(rocprofiler_t* handle, uint32_t group_index,
+                                              rocprofiler_group_t* group) {
   API_METHOD_PREFIX
   rocprofiler::Context* context = reinterpret_cast<rocprofiler::Context*>(handle);
   *group = context->GetGroupInfo(group_index);
@@ -220,7 +217,8 @@ PUBLIC_API hsa_status_t rocprofiler_get_metrics(const rocprofiler_t* handle) {
 }
 
 // Set kernel dispatch observer
-PUBLIC_API hsa_status_t rocprofiler_set_dispatch_callback(rocprofiler_callback_t callback, void* data) {
+PUBLIC_API hsa_status_t rocprofiler_set_dispatch_callback(rocprofiler_callback_t callback,
+                                                          void* data) {
   API_METHOD_PREFIX
   rocprofiler::InterceptQueue::SetDispatchCB(callback, data);
   API_METHOD_SUFFIX
@@ -234,18 +232,16 @@ PUBLIC_API hsa_status_t rocprofiler_remove_dispatch_callback() {
 }
 
 // Method for iterating the events output data
-PUBLIC_API hsa_status_t rocprofiler_iterate_trace_data(rocprofiler_t* handle, hsa_ven_amd_aqlprofile_data_callback_t callback, void* data) {
+PUBLIC_API hsa_status_t rocprofiler_iterate_trace_data(
+    rocprofiler_t* handle, hsa_ven_amd_aqlprofile_data_callback_t callback, void* data) {
   API_METHOD_PREFIX
   rocprofiler::Context* context = reinterpret_cast<rocprofiler::Context*>(handle);
   context->IterateTraceData(callback, data);
   API_METHOD_SUFFIX
 }
 
-PUBLIC_API bool OnLoad(
-    HsaApiTable* table,
-    uint64_t runtime_version,
-    uint64_t failed_tool_count,
-    const char* const * failed_tool_names) {
+PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
+                       const char* const* failed_tool_names) {
   rocprofiler::SaveHsaApi(table);
   rocprofiler::ProxyQueue::InitFactory();
   rocprofiler::InterceptQueue::SetTool(getenv("ROCP_TOOL_LIB"));
@@ -257,8 +253,6 @@ PUBLIC_API bool OnLoad(
   return true;
 }
 
-PUBLIC_API void OnUnload() {
-  rocprofiler::RestoreHsaApi();
-}
+PUBLIC_API void OnUnload() { rocprofiler::RestoreHsaApi(); }
 
-} // extern "C"
+}  // extern "C"

src/core/simple_proxy_queue.cpp

@@ -8,4 +8,4 @@ void SimpleProxyQueue::HsaIntercept(HsaApiTable* table) {
 }
 
 std::map<signal_handle_t, SimpleProxyQueue*> SimpleProxyQueue::queue_map_;
-} // namespace rocprofiler
+}  // namespace rocprofiler

src/core/simple_proxy_queue.h

@@ -20,13 +20,10 @@ typedef decltype(hsa_signal_t::handle) signal_handle_t;
 
 
 class SimpleProxyQueue : public ProxyQueue {
-  public:
+ public:
   static void HsaIntercept(HsaApiTable* table);
- 
-  static void SignalStore(
-    hsa_signal_t signal,
-    hsa_signal_value_t que_idx)
-  {
+
+  static void SignalStore(hsa_signal_t signal, hsa_signal_value_t que_idx) {
     auto it = queue_map_.find(signal.handle);
     if (it != queue_map_.end()) {
       SimpleProxyQueue* instance = it->second;
@@ -37,17 +34,17 @@ class SimpleProxyQueue : public ProxyQueue {
         // Submited packet
         const uint32_t idx = j & instance->queue_mask_;
         packet_t* packet = reinterpret_cast<packet_t*>(instance->queue_->base_address) + idx;
-        if (instance->on_submit_cb_ != NULL) instance->on_submit_cb_(packet, 1, j, instance->on_submit_cb_data_, NULL);
-        else instance->Submit(packet);
+        if (instance->on_submit_cb_ != NULL)
+          instance->on_submit_cb_(packet, 1, j, instance->on_submit_cb_data_, NULL);
+        else
+          instance->Submit(packet);
       }
     } else {
       hsa_signal_store_relaxed_fn(signal, que_idx);
     }
   }
-  
-  static uint64_t LoadIndex(
-    const hsa_queue_t *queue)
-  {
+
+  static uint64_t LoadIndex(const hsa_queue_t* queue) {
     uint64_t index = 0;
     auto it = queue_map_.find(queue->doorbell_signal.handle);
     if (it != queue_map_.end()) {
@@ -60,10 +57,7 @@ class SimpleProxyQueue : public ProxyQueue {
     return index;
   }
 
-  static void StoreIndex(
-    const hsa_queue_t *queue,
-    uint64_t value)
-  {
+  static void StoreIndex(const hsa_queue_t* queue, uint64_t value) {
     auto it = queue_map_.find(queue->doorbell_signal.handle);
     if (it != queue_map_.end()) {
       SimpleProxyQueue* instance = it->second;
@@ -107,31 +101,24 @@ class SimpleProxyQueue : public ProxyQueue {
     hsa_signal_store_relaxed_fn(doorbell_signal_, que_idx);
   }
 
-  SimpleProxyQueue() :
-    agent_info_(NULL),
-    queue_(NULL),
-    base_address_(NULL),
-    doorbell_signal_({}),
-    queue_index_(0),
-    queue_mask_(0),
-    submit_index_(0),
-    on_submit_cb_(0),
-    on_submit_cb_data_(0)
-  {}
+  SimpleProxyQueue()
+      : agent_info_(NULL),
+        queue_(NULL),
+        base_address_(NULL),
+        doorbell_signal_({}),
+        queue_index_(0),
+        queue_mask_(0),
+        submit_index_(0),
+        on_submit_cb_(0),
+        on_submit_cb_data_(0) {}
 
   ~SimpleProxyQueue() {}
 
-  private:
-  hsa_status_t Init(
-    hsa_agent_t agent,
-    uint32_t size,
-    hsa_queue_type32_t type,
-    void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
-    void *data,
-    uint32_t private_segment_size,
-    uint32_t group_segment_size,
-    hsa_queue_t **queue)
-  {
+ private:
+  hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
+                    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
+                    void* data, uint32_t private_segment_size, uint32_t group_segment_size,
+                    hsa_queue_t** queue) {
     auto status = Init(agent, size);
     *queue = queue_;
     return status;
@@ -142,13 +129,14 @@ class SimpleProxyQueue : public ProxyQueue {
     agent_info_ = util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
     if (agent_info_ != NULL) {
       if (agent_info_->dev_type == HSA_DEVICE_TYPE_GPU) {
-        status = hsa_queue_create_fn(agent, size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, UINT32_MAX, UINT32_MAX, &queue_);
+        status = hsa_queue_create_fn(agent, size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, UINT32_MAX,
+                                     UINT32_MAX, &queue_);
         if (status == HSA_STATUS_SUCCESS) {
           base_address_ = reinterpret_cast<packet_t*>(queue_->base_address);
           doorbell_signal_ = queue_->doorbell_signal;
           data_array_ = calloc(size + 1, sizeof(packet_t));
           uintptr_t addr = (uintptr_t)data_array_;
-          queue_->base_address = (void*) ((addr + align_mask_) & ~align_mask_);
+          queue_->base_address = (void*)((addr + align_mask_) & ~align_mask_);
           status = hsa_signal_create(1, 0, NULL, &(queue_->doorbell_signal));
           queue_mask_ = size - 1;
           queue_map_[queue_->doorbell_signal.handle] = this;
@@ -182,6 +170,6 @@ class SimpleProxyQueue : public ProxyQueue {
   void* data_array_;
 };
 
-} // namespace rocprofiler
+}  // namespace rocprofiler
 
-#endif // _SRC_CORE_SIMPLE_PROXY_QUEUE_H
+#endif  // _SRC_CORE_SIMPLE_PROXY_QUEUE_H