memory allocation refactoring

Change-Id: Ic63b4f5ea44f2dc5e009e3e58652a661e957b7d6

test/ctrl/run_kernel.h

@@ -57,22 +57,23 @@ template <class Kernel, class Test> bool RunKernel(int argc, char* argv[], int c
     return false;
   }
 
-  // Run test kernel
+  // Kernel dspatch iterations
   for (int i = 0; i < count; ++i) {
+    // 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;
     }
-  }
 
-  // Verify the results of the execution
-  ret_val = test_aql->VerifyResults();
-  if (ret_val) {
-    std::clog << "Test : Passed" << std::endl;
-  } else {
-    std::clog << "Test : Failed" << std::endl;
+    // Verify the results of the execution
+    ret_val = test_aql->VerifyResults();
+    if (ret_val) {
+      std::clog << "Test : Passed" << std::endl;
+    } else {
+      std::clog << "Test : Failed" << std::endl;
+    }
   }
 
   // Print time taken by sample

test/ctrl/test_hsa.cpp

@@ -115,11 +115,26 @@ bool TestHsa::Setup() {
   mem_map_t& mem_map = test_->GetMemMap();
   for (mem_it_t it = mem_map.begin(); it != mem_map.end(); ++it) {
     mem_descr_t& des = it->second;
-    void* ptr = (des.local) ? hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size)
-                            : hsa_rsrc_->AllocateSysMemory(agent_info_, des.size);
-    des.ptr = ptr;
-    TEST_ASSERT(ptr != NULL);
-    if (ptr == NULL) return false;
+    switch (des.id) {
+      case TestKernel::LOCAL_DES_ID:
+        des.ptr = hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size);
+        break;
+      case TestKernel::KERNARG_DES_ID:
+        des.ptr = hsa_rsrc_->AllocateKernArgMemory(agent_info_, des.size);
+        if (des.ptr) memset(des.ptr, 0, des.size);
+        break;
+      case TestKernel::SYS_DES_ID:
+        des.ptr = hsa_rsrc_->AllocateSysMemory(agent_info_, des.size);
+        if (des.ptr) memset(des.ptr, 0, des.size);
+        break;
+      case TestKernel::NULL_DES_ID:
+        des.ptr = NULL;
+        break;
+      default:
+        break;
+    };
+    TEST_ASSERT(des.ptr != NULL);
+    if (des.ptr == NULL) return false;
   }
   test_->Init();
 
@@ -208,24 +223,41 @@ bool TestHsa::Run() {
   hsa_signal_wait_acquire(hsa_signal_, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
                           HSA_WAIT_STATE_BLOCKED);
 
+  std::clog << "> DONE, que_idx=" << que_idx << std::endl;
+
   // Stop the timer object
   hsa_timer_.StopTimer(dispatch_timer_idx_);
   dispatch_time_taken_ = hsa_timer_.ReadTimer(dispatch_timer_idx_);
   total_time_taken_ += dispatch_time_taken_;
 
-  // Copy kernel buffers from local memory into system memory
-  const bool suc = hsa_rsrc_->CopyToHost(test_->GetOutputPtr(), test_->GetLocalPtr(), test_->GetOutputSize());
-  if (suc) test_->PrintOutput();
-
-  return suc;
+  return true;
 }
 
 bool TestHsa::VerifyResults() {
-  // Compare the results and see if they match
-  const void* const refout_ptr = test_->GetRefoutPtr();
-  const int32_t cmp_val =
-      (refout_ptr != NULL) ? memcmp(test_->GetOutputPtr(), refout_ptr, test_->GetOutputSize()) : 0;
-  return (cmp_val == 0);
+  bool cmp = false;
+  void* output = NULL;
+  const uint32_t size = test_->GetOutputSize();
+  bool suc = false;
+
+  // Copy local kernel output buffers from local memory into host memory
+  if (test_->IsOutputLocal()) {
+    output = hsa_rsrc_->AllocateSysMemory(agent_info_, size);
+    suc = hsa_rsrc_->Memcpy(agent_info_, output, test_->GetOutputPtr(), size);
+  } else {
+    output = test_->GetOutputPtr();;
+    suc = true;
+  }
+
+  if ((output != NULL) && suc) {
+    // Print the test output
+    test_->PrintOutput(output);
+    // Compare the results and see if they match
+    cmp = (memcmp(output, test_->GetRefOut(), size) == 0);
+  }
+
+  if (test_->IsOutputLocal() && (output != NULL)) hsa_rsrc_->FreeMemory(output);
+
+  return cmp;
 }
 
 void TestHsa::PrintTime() {

test/ctrl/test_kernel.h

@@ -28,22 +28,23 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
 #ifndef TEST_CTRL_TEST_KERNEL_H_
 #define TEST_CTRL_TEST_KERNEL_H_
 
+#include <string.h>
 #include <stdint.h>
 #include <map>
 
-#include "util/hsa_rsrc_factory.h"
-
 // Class implements kernel test
 class TestKernel {
  public:
+  // Exported buffers IDs
+  enum buf_id_t { KERNARG_EXP_ID, OUTPUT_EXP_ID, REFOUT_EXP_ID };
   // Memory descriptors IDs
-  enum { INPUT_DES_ID, OUTPUT_DES_ID, LOCAL_DES_ID, MASK_DES_ID, KERNARG_DES_ID, REFOUT_DES_ID };
+  enum des_id_t { NULL_DES_ID, LOCAL_DES_ID, KERNARG_DES_ID, SYS_DES_ID, REFOUT_DES_ID };
 
   // Memory descriptors vector declaration
   struct mem_descr_t {
+    des_id_t id;
     void* ptr;
     uint32_t size;
-    bool local;
   };
 
   // Memory map declaration
@@ -51,12 +52,7 @@ class TestKernel {
   typedef mem_map_t::iterator mem_it_t;
   typedef mem_map_t::const_iterator mem_const_it_t;
 
-  virtual ~TestKernel() {
-    for (auto& entry : mem_map_) {
-      void* ptr = entry.second.ptr;
-      if (ptr != NULL) HsaRsrcFactory::MemoryFree(ptr);
-    }
-  }
+  virtual ~TestKernel() {}
 
   // Initialize method
   virtual void Init() = 0;
@@ -65,50 +61,79 @@ class TestKernel {
   mem_map_t& GetMemMap() { return mem_map_; }
 
   // Return NULL descriptor
-  static mem_descr_t NullDescriptor() { return {NULL, 0, 0}; }
+  static mem_descr_t NullDescriptor() { return {NULL_DES_ID, NULL, 0}; }
+
+  // Check if decripter is local
+  bool IsLocal(const mem_descr_t& descr) const { return (descr.id == LOCAL_DES_ID); }
 
   // Methods to get the kernel attributes
-  void* GetKernargPtr() const { return GetDescr(KERNARG_DES_ID).ptr; }
-  uint32_t GetKernargSize() const { return GetDescr(KERNARG_DES_ID).size; }
-  void* GetOutputPtr() const { return GetDescr(OUTPUT_DES_ID).ptr; }
-  uint32_t GetOutputSize() const { return GetDescr(OUTPUT_DES_ID).size; }
-  void* GetLocalPtr() const { return GetDescr(LOCAL_DES_ID).ptr; }
-  void* GetRefoutPtr() const { return GetDescr(REFOUT_DES_ID).ptr; }
+  const mem_descr_t& GetKernargDescr() { return *test_map_[KERNARG_EXP_ID]; };
+  const mem_descr_t& GetOutputDescr() { return *test_map_[OUTPUT_EXP_ID]; };
+  void* GetKernargPtr() { return GetKernargDescr().ptr; }
+  uint32_t GetKernargSize() { return GetKernargDescr().size; }
+  void* GetOutputPtr() { return GetOutputDescr().ptr; }
+  uint32_t GetOutputSize() { return GetOutputDescr().size; }
+  bool IsOutputLocal() { return IsLocal(GetOutputDescr()); }
   virtual uint32_t GetGridSize() const = 0;
 
+  // Return reference output
+  void* GetRefOut() { return test_map_[REFOUT_EXP_ID]->ptr; };
+
   // Print output
-  virtual void PrintOutput() const = 0;
+  virtual void PrintOutput(const void* ptr) const = 0;
 
   // Return name
   virtual std::string Name() const = 0;
 
  protected:
-  // Set system memory descriptor
-  bool SetSysDescr(const uint32_t& id, const uint32_t& size) {
-    return SetMemDescr(id, size, false);
+  // Set buffer descriptor
+  bool SetInDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
+    bool suc = SetMemDescr(buf_id, des_id, size);
+    if (des_id == KERNARG_DES_ID) {
+      test_map_[KERNARG_EXP_ID] = &mem_map_[buf_id];
+    }
+    return suc;
   }
 
-  // Set local memory descriptor
-  bool SetLocalDescr(const uint32_t& id, const uint32_t& size) {
-    return SetMemDescr(id, size, true);
+  // Set results descriptor
+  bool SetOutDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
+    bool suc = SetMemDescr(buf_id, des_id, size);
+    test_map_[OUTPUT_EXP_ID] = &mem_map_[buf_id];
+    return suc;
+  }
+
+  // Set host descriptor
+  bool SetHostDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
+    bool suc = SetMemDescr(buf_id, des_id, size);
+    if (suc) {
+      mem_descr_t& descr = mem_map_[buf_id];
+      descr.ptr = malloc(size);
+      if (des_id == REFOUT_DES_ID) {
+        test_map_[REFOUT_EXP_ID] = &descr;
+      }
+      if (descr.ptr == NULL) suc = false;
+    }
+    return suc;
   }
 
   // Get memory descriptor
-  mem_descr_t GetDescr(const uint32_t& id) const {
-    mem_const_it_t it = mem_map_.find(id);
+  mem_descr_t GetDescr(const uint32_t& buf_id) const {
+    mem_const_it_t it = mem_map_.find(buf_id);
     return (it != mem_map_.end()) ? it->second : NullDescriptor();
   }
 
  private:
   // Set memory descriptor
-  bool SetMemDescr(const uint32_t& id, const uint32_t& size, const bool& local) {
-    const mem_descr_t des = {NULL, size, local};
-    auto ret = mem_map_.insert(mem_map_t::value_type(id, des));
+  bool SetMemDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
+    const mem_descr_t des = {des_id, NULL, size};
+    auto ret = mem_map_.insert(mem_map_t::value_type(buf_id, des));
     return ret.second;
   }
 
   // Kernel memory map object
   mem_map_t mem_map_;
+  // Test memory map object
+  std::map<uint32_t, mem_descr_t*> test_map_;
 };
 
 #endif  // TEST_CTRL_TEST_KERNEL_H_

test/ctrl/tool.cpp

@@ -266,16 +266,14 @@ hsa_status_t trace_data_cb(hsa_ven_amd_aqlprofile_info_type_t info_type,
     const uint32_t data_size = info_data->sqtt_data.size;
     const void* data_ptr = info_data->sqtt_data.ptr;
     fprintf(arg->file, "    SE(%u) size(%u)\n", info_data->sample_id, data_size);
-#if 1
-    dump_sqtt_trace(arg->label, info_data->sample_id, data_ptr, data_size);
-#else
-    void* buffer = malloc(data_size);
-    memset(buffer, 0, data_size);
-    const bool suc = HsaRsrcFactory::Instance().CopyToHost(arg->agent, buffer, data_ptr, data_size);
+
+    HsaRsrcFactory* hsa_rsrc = &HsaRsrcFactory::Instance();
+    const AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(arg->agent);
+    void* buffer = hsa_rsrc->AllocateSysMemory(agent_info, data_size);
+    const bool suc = HsaRsrcFactory::Memcpy(arg->agent, buffer, data_ptr, data_size);
     if (suc) dump_sqtt_trace(arg->label, info_data->sample_id, buffer, data_size);
     else fatal("SQTT data memcopy to host failed");
-    free(buffer);
-#endif
+    HsaRsrcFactory::FreeMemory(buffer);
   } else
     status = HSA_STATUS_ERROR;
   return status;
@@ -320,7 +318,7 @@ void output_results(const context_entry_t* entry, const char* label) {
             size += chunk_size;
           }
           fprintf(file, "size(%lu)\n", size);
-          free(p->data.result_bytes.ptr);
+          HsaRsrcFactory::FreeMemory(p->data.result_bytes.ptr);
           const_cast<rocprofiler_feature_t*>(p)->data.result_bytes.size = 0;
         } else {
           fprintf(file, "(\n");