fixing resources releasing

test/ctrl/run_kernel.h

@@ -31,7 +31,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "ctrl/test_hsa.h"
 #include "util/test_assert.h"
 
-template <class Kernel, class Test> bool RunKernel(int argc, char* argv[]) {
+template <class Kernel, class Test> bool RunKernel(int argc, char* argv[], int count = 1) {
   bool ret_val = false;
 
   // Create test kernel object
@@ -58,11 +58,13 @@ template <class Kernel, class Test> bool RunKernel(int argc, char* argv[]) {
   }
 
   // Run test kernel
-  ret_val = test_aql->Run();
-  if (ret_val == false) {
-    std::cerr << "Error in running the test kernel" << std::endl;
-    TEST_ASSERT(ret_val);
-    return false;
+  for (int i = 0; i < count; ++i) {
+    ret_val = test_aql->Run();
+    if (ret_val == false) {
+      std::cerr << "Error in running the test kernel" << std::endl;
+      TEST_ASSERT(ret_val);
+      return false;
+    }
   }
 
   // Verify the results of the execution

test/ctrl/test.cpp

@@ -7,8 +7,12 @@
 #include "simple_convolution/simple_convolution.h"
 
 int main(int argc, char** argv) {
+  const char* kiter_s = getenv("ROCP_KITER");
+  const char* diter_s = getenv("ROCP_DITER");
+  const int kiter = (kiter_s != NULL) ? atol(kiter_s) : 1;
+  const int diter = (diter_s != NULL) ? atol(diter_s) : 1;
   TestHsa::HsaInstantiate();
-  for (int i = 0; i < 3; ++i) RunKernel<SimpleConvolution, TestAql>(argc, argv);
-  TestHsa::HsaShutdown();
+  for (int i = 0; i < kiter; ++i) RunKernel<SimpleConvolution, TestAql>(argc, argv, diter);
+  //TestHsa::HsaShutdown();
   return 0;
 }

test/ctrl/test_hsa.cpp

@@ -69,6 +69,7 @@ HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) {
 }
 
 void TestHsa::HsaShutdown() {
+  if (hsa_queue_ != NULL) hsa_queue_destroy(hsa_queue_);
   if (hsa_rsrc_) hsa_rsrc_->Destroy();
 }
 
@@ -121,11 +122,11 @@ bool TestHsa::Setup() {
 
   // Load and Finalize Kernel Code Descriptor
   char* brig_path = (char*)brig_path_obj_.c_str();
-  const bool ret_val =
-      hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, strdup(name_.c_str()), &kernel_code_desc_);
-  if (ret_val == false) {
+  code_buf_ =
+      hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, name_.c_str(), &hsa_exec_, &kernel_code_desc_);
+  if (code_buf_ == NULL) {
     std::cerr << "Error in loading and finalizing Kernel" << std::endl;
-    return ret_val;
+    return false;
   }
 
   // Stop the timer object
@@ -187,32 +188,16 @@ bool TestHsa::Run() {
   aql.group_segment_size = group_segment_size;
   aql.private_segment_size = private_segment_size;
   // Initialize Aql packet with handle of signal
+  hsa_signal_store_relaxed(hsa_signal_, 1);
   aql.completion_signal = hsa_signal_;
 
-  // Compute the write index of queue and copy Aql packet into it
-  const uint64_t que_idx = hsa_queue_load_write_index_relaxed(hsa_queue_);
-  const uint32_t mask = hsa_queue_->size - 1;
-
   std::clog << "> Executing kernel: \"" << name_ << "\"" << std::endl;
 
   // Start the timer object
   hsa_timer_.StartTimer(dispatch_timer_idx_);
 
-  // Disable packet so that submission to HW is complete
-  const auto header = aql.header;
-  aql.header = HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
-
-  // Copy Aql packet into queue buffer
-  ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask] = aql;
-
-  // After AQL packet is fully copied into queue buffer
-  // update packet header from invalid state to valid state
-  std::atomic_thread_fence(std::memory_order_release);
-  ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask].header = header;
-
-  // Increment the write index and ring the doorbell to dispatch the kernel.
-  hsa_queue_store_write_index_relaxed(hsa_queue_, (que_idx + 1));
-  hsa_signal_store_relaxed(hsa_queue_->doorbell_signal, que_idx);
+  // Submit AQL packet to the queue
+  const uint64_t que_idx = hsa_rsrc_->Submit(hsa_queue_, &aql);
 
   std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl;
 
@@ -251,4 +236,9 @@ void TestHsa::PrintTime() {
             << std::endl;
 }
 
-bool TestHsa::Cleanup() { return true; }
+bool TestHsa::Cleanup() {
+  hsa_executable_destroy(hsa_exec_);
+  hsa_memory_free(code_buf_);
+  hsa_signal_destroy(hsa_signal_);
+  return true;
+}

test/ctrl/test_hsa.h

@@ -47,6 +47,8 @@ class TestHsa : public TestAql {
     total_time_taken_ = 0;
     setup_time_taken_ = 0;
     dispatch_time_taken_ = 0;
+    code_buf_ = NULL;
+    hsa_exec_ = {};
   }
 
   // Get methods for Agent Info, HAS queue, HSA Resourcse Manager
@@ -120,6 +122,10 @@ class TestHsa : public TestAql {
 
   // Test kernel name
   std::string name_;
+
+  // Kernel code buffer
+  void* code_buf_;
+  hsa_executable_t hsa_exec_;
 };
 
 #endif  // TEST_CTRL_TEST_HSA_H_

test/ctrl/test_pgen.h

@@ -1,45 +0,0 @@
-/******************************************************************************
-
-Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
-
-Redistribution and use in source and binary forms, with or without modification,
-are permitted provided that the following conditions are met:
-
-Redistributions of source code must retain the above copyright notice, this list
-of conditions and the following disclaimer.
-
-Redistributions in binary form must reproduce the above copyright notice, this
-list of conditions and the following disclaimer in the documentation and/or
-other materials provided with the distribution.
-
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
-IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
-INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
-BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
-DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
-LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
-OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
-OF THE POSSIBILITY OF SUCH DAMAGE.
-
-*******************************************************************************/
-
-#ifndef TEST_CTRL_TEST_PGEN_H_
-#define TEST_CTRL_TEST_PGEN_H_
-
-#include "ctrl/test_pmgr.h"
-
-// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
-class TestPGen : public TestPMgr {
- protected:
-  typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
-
-  packet_t* PrePacket() { return reinterpret_cast<packet_t*>(&pre_packet_); }
-  packet_t* PostPacket() { return reinterpret_cast<packet_t*>(&post_packet_); }
-
- public:
-  explicit TestPGen(TestAql* t) : TestPMgr(t) {}
-};
-
-#endif  // TEST_CTRL_TEST_PGEN_H_

test/ctrl/test_pmgr.cpp

@@ -1,141 +0,0 @@
-/******************************************************************************
-
-Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
-
-Redistribution and use in source and binary forms, with or without modification,
-are permitted provided that the following conditions are met:
-
-Redistributions of source code must retain the above copyright notice, this list
-of conditions and the following disclaimer.
-
-Redistributions in binary form must reproduce the above copyright notice, this
-list of conditions and the following disclaimer in the documentation and/or
-other materials provided with the distribution.
-
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
-IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
-INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
-BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
-DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
-LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
-OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
-OF THE POSSIBILITY OF SUCH DAMAGE.
-
-*******************************************************************************/
-
-#include "ctrl/test_pmgr.h"
-
-#include <atomic>
-
-#include "ctrl/test_assert.h"
-
-bool TestPMgr::AddPacketGfx9(const packet_t* packet) {
-  packet_t aql_packet = *packet;
-
-  // Compute the write index of queue and copy Aql packet into it
-  uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue());
-  const uint32_t mask = GetQueue()->size - 1;
-  packet_t* slot = (reinterpret_cast<packet_t*>(GetQueue()->base_address)) + (que_idx & mask);
-
-  // Disable packet so that submission to HW is complete
-  const auto header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
-  aql_packet.header &= (~((1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)) << HSA_PACKET_HEADER_TYPE;
-  aql_packet.header |= HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
-
-  // Copy Aql packet into queue buffer
-  *slot = aql_packet;
-  // After AQL packet is fully copied into queue buffer
-  // update packet header from invalid state to valid state
-  auto header_atomic_ptr = reinterpret_cast<std::atomic<uint16_t>*>(&slot->header);
-  header_atomic_ptr->store(header, std::memory_order_release);
-
-  // Increment the write index and ring the doorbell to dispatch the kernel.
-  hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1));
-  hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx);
-
-  return true;
-}
-
-bool TestPMgr::AddPacketGfx8(const packet_t* packet) {
-  // Create legacy devices PM4 data
-  const hsa_ext_amd_aql_pm4_packet_t* aql_packet = (const hsa_ext_amd_aql_pm4_packet_t*)packet;
-  slot_pm4_t data;
-  api_->hsa_ven_amd_aqlprofile_legacy_get_pm4(aql_packet, reinterpret_cast<void*>(data.words));
-
-  // Compute the write index of queue and copy Aql packet into it
-  uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue());
-  const uint32_t mask = GetQueue()->size - 1;
-
-  // Copy Aql/Pm4 blob into queue buffer
-  packet_t* ptr = (reinterpret_cast<packet_t*>(GetQueue()->base_address)) + (que_idx & mask);
-  slot_pm4_t* slot = reinterpret_cast<slot_pm4_t*>(ptr);
-  for (unsigned i = 1; i < SLOT_PM4_SIZE_DW; ++i) {
-    slot->words[i] = data.words[i];
-  }
-  // To maintain global order to ensure the prior copy of the packet contents is made visible
-  // before the header is updated.
-  // With in-order CP it will wait until the first packet in the blob will be valid
-  std::atomic<uint32_t>* header_atomic_ptr =
-      reinterpret_cast<std::atomic<uint32_t>*>(&slot->words[0]);
-  header_atomic_ptr->store(data.words[0], std::memory_order_release);
-
-  // Increment the write index and ring the doorbell to dispatch the kernel.
-  que_idx += SLOT_PM4_SIZE_AQLP - 1;
-  hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1));
-  hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx);
-
-  return true;
-}
-
-bool TestPMgr::AddPacket(const packet_t* packet) {
-  const char* agent_name = GetAgentInfo()->name;
-  return (strncmp(agent_name, "gfx8", 4) == 0) ? AddPacketGfx8(packet) : AddPacketGfx9(packet);
-}
-
-bool TestPMgr::Run() {
-  // Build Aql Pkts
-  const bool active = BuildPackets();
-  if (active) {
-    // Submit Pre-Dispatch Aql packet
-    AddPacket(&pre_packet_);
-  }
-
-  Test()->Run();
-
-  if (active) {
-    // Set post packet completion signal
-    post_packet_.completion_signal = post_signal_;
-
-    // Submit Post-Dispatch Aql packet
-    AddPacket(&post_packet_);
-
-    // Wait for Post-Dispatch packet to complete
-    hsa_signal_wait_acquire(post_signal_, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
-                            HSA_WAIT_STATE_BLOCKED);
-
-    // Dumping profiling data
-    DumpData();
-  }
-
-  return true;
-}
-
-bool TestPMgr::Initialize(int argc, char** argv) {
-  TestAql::Initialize(argc, argv);
-
-  hsa_status_t status = HSA_STATUS_ERROR;
-  status = hsa_signal_create(1, 0, NULL, &post_signal_);
-  TEST_ASSERT(status == HSA_STATUS_SUCCESS);
-  api_ = HsaRsrcFactory::Instance().AqlProfileApi();;
-
-  return true;
-}
-
-TestPMgr::TestPMgr(TestAql* t) : TestAql(t), api_(NULL) {
-  memset(&pre_packet_, 0, sizeof(pre_packet_));
-  memset(&post_packet_, 0, sizeof(post_packet_));
-  dummy_signal_.handle = 0;
-  post_signal_ = dummy_signal_;
-}

test/ctrl/test_pmgr.h

@@ -1,70 +0,0 @@
-/******************************************************************************
-
-Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
-
-Redistribution and use in source and binary forms, with or without modification,
-are permitted provided that the following conditions are met:
-
-Redistributions of source code must retain the above copyright notice, this list
-of conditions and the following disclaimer.
-
-Redistributions in binary form must reproduce the above copyright notice, this
-list of conditions and the following disclaimer in the documentation and/or
-other materials provided with the distribution.
-
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
-IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
-INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
-BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
-DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
-LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
-OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
-OF THE POSSIBILITY OF SUCH DAMAGE.
-
-*******************************************************************************/
-
-#ifndef TEST_CTRL_TEST_PMGR_H_
-#define TEST_CTRL_TEST_PMGR_H_
-
-#include <hsa.h>
-#include <hsa_ven_amd_aqlprofile.h>
-#include <atomic>
-
-#include "ctrl/test_aql.h"
-
-// Class implements profiling manager
-class TestPMgr : public TestAql {
- public:
-  typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
-  explicit TestPMgr(TestAql* t);
-  bool Run();
-
- protected:
-  packet_t pre_packet_;
-  packet_t post_packet_;
-  hsa_signal_t dummy_signal_;
-  hsa_signal_t post_signal_;
-
-  HsaRsrcFactory::aqlprofile_pfn_t* api_;
-
-  virtual bool BuildPackets() { return false; }
-  virtual bool DumpData() { return false; }
-  virtual bool Initialize(int argc, char** argv);
-
- private:
-  enum {
-    SLOT_PM4_SIZE_DW = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t),
-    SLOT_PM4_SIZE_AQLP = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t)
-  };
-  struct slot_pm4_t {
-    uint32_t words[SLOT_PM4_SIZE_DW];
-  };
-
-  bool AddPacket(const packet_t* packet);
-  bool AddPacketGfx8(const packet_t* packet);
-  bool AddPacketGfx9(const packet_t* packet);
-};
-
-#endif  // TEST_CTRL_TEST_PMGR_H_

test/ctrl/tool.cpp

@@ -184,6 +184,8 @@ void output_results(FILE* file, const rocprofiler_feature_t* features, const uns
             ptr = chunk_data + off;
             size += chunk_size;
           }
+          free(p->data.result_bytes.ptr);
+
           fprintf(file, "size(%lu)\n", size);
           if (size > p->data.result_bytes.size) {
             fprintf(stderr, "SQTT data size is out of the result buffer size\n");

test/util/hsa_rsrc_factory.cpp

@@ -32,6 +32,7 @@ POSSIBILITY OF SUCH DAMAGE.
 #include <string.h>
 #include <stdlib.h>
 
+#include <atomic>
 #include <cassert>
 #include <fstream>
 #include <iostream>
@@ -111,6 +112,7 @@ static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data) {
 // Constructor of the class
 HsaRsrcFactory::HsaRsrcFactory() {
   // Initialize the Hsa Runtime
+  printf("HSA init\n");
   hsa_status_t status = hsa_init();
   CHECK_STATUS("Error in hsa_init", status);
 
@@ -130,6 +132,10 @@ HsaRsrcFactory::HsaRsrcFactory() {
 
 // Destructor of the class
 HsaRsrcFactory::~HsaRsrcFactory() {
+  for (auto p : cpu_list_) free(p);
+  for (auto p : gpu_list_) free(p);
+
+  printf("HSA shutdown\n");
   hsa_status_t status = hsa_shut_down();
   CHECK_STATUS("Error in hsa_shut_down", status);
 }
@@ -318,8 +324,8 @@ bool HsaRsrcFactory::TransferData(void* dest_buff, void* src_buff, uint32_t leng
 //
 // @return bool true if successful, false otherwise
 //
-bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path,
-                                     char* kernel_name, hsa_executable_symbol_t* code_desc) {
+void* HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path,
+                                      const char* kernel_name, hsa_executable_t* hsa_exec, hsa_executable_symbol_t* code_desc) {
   // Finalize the Hsail object into code object
   hsa_status_t status;
   hsa_code_object_t code_object;
@@ -333,52 +339,52 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_pat
   if (!codeStream) {
     std::cerr << "Error: failed to load " << filename << std::endl;
     assert(false);
-    return false;
+    return NULL;
   }
 
   // Allocate memory to read in code object from file
   size_t size = std::string::size_type(codeStream.tellg());
-  char* codeBuff = (char*)AllocateSysMemory(agent_info, size);
-  if (!codeBuff) {
+  char* code_buf = (char*)AllocateSysMemory(agent_info, size);
+  if (!code_buf) {
     std::cerr << "Error: failed to allocate memory for code object." << std::endl;
     assert(false);
-    return false;
+    return NULL;
   }
 
   // Read the code object into allocated memory
   codeStream.seekg(0, std::ios::beg);
-  std::copy(std::istreambuf_iterator<char>(codeStream), std::istreambuf_iterator<char>(), codeBuff);
+  std::copy(std::istreambuf_iterator<char>(codeStream), std::istreambuf_iterator<char>(), code_buf);
 
   // De-Serialize the code object that has been read into memory
-  status = hsa_code_object_deserialize(codeBuff, size, NULL, &code_object);
+  status = hsa_code_object_deserialize(code_buf, size, NULL, &code_object);
   if (status != HSA_STATUS_SUCCESS) {
     std::cerr << "Failed to deserialize code object" << std::endl;
-    return false;
+    if (code_buf) hsa_memory_free(code_buf);
+    return NULL;
   }
 
   // Create executable.
-  hsa_executable_t hsaExecutable;
   status =
-      hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable);
+      hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", hsa_exec);
   CHECK_STATUS("Error in creating executable object", status);
 
   // Load code object.
-  status = hsa_executable_load_code_object(hsaExecutable, agent_info->dev_id, code_object, "");
+  status = hsa_executable_load_code_object(*hsa_exec, agent_info->dev_id, code_object, "");
   CHECK_STATUS("Error in loading executable object", status);
 
   // Freeze executable.
-  status = hsa_executable_freeze(hsaExecutable, "");
+  status = hsa_executable_freeze(*hsa_exec, "");
   CHECK_STATUS("Error in freezing executable object", status);
 
   // Get symbol handle.
   hsa_executable_symbol_t kernelSymbol;
-  status = hsa_executable_get_symbol(hsaExecutable, NULL, kernel_name, agent_info->dev_id, 0,
+  status = hsa_executable_get_symbol(*hsa_exec, NULL, kernel_name, agent_info->dev_id, 0,
                                      &kernelSymbol);
   CHECK_STATUS("Error in looking up kernel symbol", status);
 
   // Update output parameter
   *code_desc = kernelSymbol;
-  return true;
+  return code_buf;
 }
 
 // Add an instance of AgentInfo representing a Hsa Gpu agent
@@ -411,5 +417,32 @@ bool HsaRsrcFactory::PrintGpuAgents(const std::string& header) {
   return true;
 }
 
+uint64_t HsaRsrcFactory::Submit(hsa_queue_t* queue, void* packet) {
+  const uint32_t slot_size_b = 0x40;
+
+  // adevance command queue
+  const uint64_t write_idx = hsa_queue_load_write_index_relaxed(queue);
+  hsa_queue_store_write_index_relaxed(queue, write_idx + 1);
+  while ((write_idx - hsa_queue_load_read_index_relaxed(queue)) >= queue->size) {
+    sched_yield();
+  }
+
+  uint32_t slot_idx = (uint32_t)(write_idx % queue->size);
+  uint32_t* queue_slot = (uint32_t*)((uintptr_t)(queue->base_address) + (slot_idx * slot_size_b));
+  uint32_t* slot_data = (uint32_t*)packet;
+
+  // Copy buffered commands into the queue slot.
+  // Overwrite the AQL invalid header (first dword) last.
+  // This prevents the slot from being read until it's fully written.
+  memcpy(&queue_slot[1], &slot_data[1], slot_size_b - sizeof(uint32_t));
+  std::atomic<uint32_t>* header_atomic_ptr = reinterpret_cast<std::atomic<uint32_t>*>(&queue_slot[0]);
+  header_atomic_ptr->store(slot_data[0], std::memory_order_release);
+
+  // ringdoor bell
+  hsa_signal_store_relaxed(queue->doorbell_signal, write_idx);
+
+  return write_idx;
+}
+
 HsaRsrcFactory* HsaRsrcFactory::instance_ = NULL;
 HsaRsrcFactory::mutex_t HsaRsrcFactory::mutex_;

test/util/hsa_rsrc_factory.h

@@ -196,10 +196,10 @@ class HsaRsrcFactory {
   // @param code_desc Handle of finalized Code Descriptor that could
   // be used to submit for execution
   //
-  // @return bool true if successful, false otherwise
+  // @return code buffer, non NULL if successful, NULL otherwise
   //
-  bool LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, char* kernel_name,
-                       hsa_executable_symbol_t* code_desc);
+  void* LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, const char* kernel_name,
+                        hsa_executable_t* hsa_exec, hsa_executable_symbol_t* code_desc);
 
   // Add an instance of AgentInfo representing a Hsa Gpu agent
   void AddAgentInfo(AgentInfo* agent_info, bool gpu);
@@ -207,6 +207,9 @@ class HsaRsrcFactory {
   // Print the various fields of Hsa Gpu Agents
   bool PrintGpuAgents(const std::string& header);
 
+  // Submit AQL packet to given queue
+  static uint64_t Submit(hsa_queue_t* queue, void* packet);
+
   // Return AqlProfile API table
   typedef hsa_ven_amd_aqlprofile_1_00_pfn_t aqlprofile_pfn_t;
   const aqlprofile_pfn_t* AqlProfileApi() const { return &aqlprofile_api_; }
@@ -225,12 +228,6 @@ class HsaRsrcFactory {
   static HsaRsrcFactory* instance_;
   static mutex_t mutex_;
 
-  // Used to maintain a list of Hsa Queue handles
-  std::vector<hsa_queue_t*> queue_list_;
-
-  // Used to maintain a list of Hsa Signal handles
-  std::vector<hsa_signal_t*> signal_list_;
-
   // Used to maintain a list of Hsa Gpu Agent Info
   std::vector<AgentInfo*> gpu_list_;