contexts complition signals destrotying

[ROCm/rocprofiler commit: b8403823f0]

projects/rocprofiler/src/core/hsa_proxy_queue.h

@@ -29,7 +29,6 @@ class HsaProxyQueue : public ProxyQueue {
                     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_);
     *queue = queue_;

projects/rocprofiler/src/core/profile.h

@@ -82,10 +82,15 @@ class Profile {
     is_legacy_ = (strncmp(agent_info->name, "gfx8", 4) == 0);
   }
   virtual ~Profile() {
-    hsa_memory_free(profile_.command_buffer.ptr);
-    hsa_memory_free(profile_.output_buffer.ptr);
-    free(const_cast<event_t*>(profile_.events));
-    free(const_cast<parameter_t*>(profile_.parameters));
+    if (!info_vector_.empty()) {
+      info_vector_.clear();
+      hsa_memory_free(profile_.command_buffer.ptr);
+      hsa_memory_free(profile_.output_buffer.ptr);
+      free(const_cast<event_t*>(profile_.events));
+      free(const_cast<parameter_t*>(profile_.parameters));
+      hsa_status_t status = hsa_signal_destroy(completion_signal_);
+      if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_destroy " << std::hex << status);
+    }
   }
 
   virtual void Insert(const profile_info_t& info) { info_vector_.push_back(info.rinfo); }
@@ -114,7 +119,7 @@ class Profile {
       start.completion_signal = dummy_signal;
       hsa_signal_t post_signal;
       status = hsa_signal_create(1, 0, NULL, &post_signal);
-      if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "hsa_signal_create");
+      if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
       stop.completion_signal = post_signal;
       completion_signal_ = post_signal;
 

projects/rocprofiler/src/core/simple_proxy_queue.h

@@ -186,11 +186,20 @@ class SimpleProxyQueue : public ProxyQueue {
   }
 
   hsa_status_t Cleanup() const {
-    hsa_status_t status = HSA_STATUS_SUCCESS;
+    hsa_status_t status = HSA_STATUS_ERROR;
+    hsa_signal_t queue_signal = queue_->doorbell_signal;
+
+    // Destroy original HSA queue
     queue_->base_address = base_address_;
     queue_->doorbell_signal = doorbell_signal_;
     status = hsa_queue_destroy_fn(queue_);
+    if (status != HSA_STATUS_SUCCESS) abort();
+
+    // Destroy overloaded virtual queue data and signal
     free(data_array_);
+    status = hsa_signal_destroy(queue_signal);
+    if (status != HSA_STATUS_SUCCESS) abort();
+
     return status;
   }
 

projects/rocprofiler/test/ctrl/tool.cpp

@@ -169,7 +169,7 @@ void output_results(FILE* file, const rocprofiler_feature_t* features, const uns
         break;
       }
       default:
-        std::cout << "Bad result kind (" << p->data.kind << ")" << std::endl;
+        fprintf(file, "undefined data kind(%u)\n", p->data.kind);
     }
   }
 }
@@ -194,7 +194,7 @@ void dump_context(context_entry_t* entry) {
     FILE* file_handle = entry->file_handle;
 
     fprintf(file_handle,
-            "Dispatch[%u], queue_index(%lu), kernel_object(0x%lx), kernel_name(\"%s\"):\n", index,
+            "dispatch[%u], cntx(%p), queue_index(%lu), kernel_object(0x%lx), kernel_name(\"%s\"):\n", index, group.context,
             entry->data.queue_index, entry->data.kernel_object, entry->data.kernel_name);
 
     status = rocprofiler_group_get_data(&group);
@@ -416,9 +416,9 @@ DESTRUCTOR_API void destructor() {
   printf("\nROCPRofiler: %u contexts collected", context_array_count);
   if (result_prefix == NULL) {
     printf("\n");
+    // Dump profiling output data
+    dump_context_array();
   } else {
     printf(", dumping to %s\n", result_prefix);
   }
-  // Dump profiling output data which hasn't yet dumped by completi onhandler
-  dump_context_array();
 }