diff --git a/projects/rocprofiler/src/core/hsa_proxy_queue.h b/projects/rocprofiler/src/core/hsa_proxy_queue.h index e1c33da0fe..030eb8b65f 100644 --- a/projects/rocprofiler/src/core/hsa_proxy_queue.h +++ b/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_; diff --git a/projects/rocprofiler/src/core/profile.h b/projects/rocprofiler/src/core/profile.h index 3847cb565d..eb169dbe14 100644 --- a/projects/rocprofiler/src/core/profile.h +++ b/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(profile_.events)); - free(const_cast(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(profile_.events)); + free(const_cast(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; diff --git a/projects/rocprofiler/src/core/simple_proxy_queue.h b/projects/rocprofiler/src/core/simple_proxy_queue.h index 0695fb4596..6b3f73b7d1 100644 --- a/projects/rocprofiler/src/core/simple_proxy_queue.h +++ b/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; } diff --git a/projects/rocprofiler/test/ctrl/tool.cpp b/projects/rocprofiler/test/ctrl/tool.cpp index 9077c4718a..dd68bf7005 100644 --- a/projects/rocprofiler/test/ctrl/tool.cpp +++ b/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(); }