several fixes; adding direct loading of alprofile library

test/ctrl/test_hsa.cpp

@@ -61,6 +61,7 @@ HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) {
       uint32_t num_pkts = 128;
       if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
         hsa_queue_ = NULL;
+        TEST_ASSERT(false);
       }
     }
   }

test/ctrl/test_pgen_rocp.h

@@ -1,78 +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_ROCP_H_
-#define TEST_CTRL_TEST_PGEN_ROCP_H_
-
-#include <list>
-#include <vector>
-
-#include "ctrl/test_pgen.h"
-#include "util/test_assert.h"
-
-hsa_status_t TestPGenRocpCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
-                                  hsa_ven_amd_aqlprofile_info_data_t* info_data,
-                                  void* callback_data) {
-  hsa_status_t status = HSA_STATUS_SUCCESS;
-  typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> passed_data_t;
-  reinterpret_cast<passed_data_t*>(callback_data)->push_back(*info_data);
-  return status;
-}
-
-// Class implements PMC profiling
-class TestPGenRocp : public TestPGen {
- public:
-  explicit TestPGenRocp(TestAql* t) : TestPGen(t) { std::clog << "Test: PGen ROCP" << std::endl; }
-
-  bool Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
-    status = rocprofiler_on_dispatch(&profile_, PrePacket(), PostPacket());
-    TEST_STATUS(status != HSA_STATUS_SUCCESS);
-    return (status == HSA_STATUS_SUCCESS);
-  }
-
- private:
-  bool BuildPackets() { return true; }
-
-  bool DumpData() {
-    std::clog << "TestPGenRocp::DumpData :" << std::endl;
-
-    typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
-
-    callback_data_t data;
-    api_.hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenRocpCallback, &data);
-    for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
-      std::cout << std::dec << "event(block(" << it->pmc_data.event.block_name << "_"
-                << it->pmc_data.event.block_index << "), id(" << it->pmc_data.event.counter_id
-                << ")), sample(" << it->sample_id << "), result(" << it->pmc_data.result << ")"
-                << std::endl;
-    }
-
-    return true;
-  }
-};
-
-#endif  // TEST_CTRL_TEST_PGEN_ROCP_H_

test/ctrl/test_pmgr.cpp

@@ -62,7 +62,7 @@ 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));
+  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());
@@ -128,16 +128,14 @@ bool TestPMgr::Initialize(int argc, char** argv) {
   hsa_status_t status = HSA_STATUS_ERROR;
   status = hsa_signal_create(1, 0, NULL, &post_signal_);
   TEST_ASSERT(status == HSA_STATUS_SUCCESS);
-  status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &api_);
-  TEST_ASSERT(status == HSA_STATUS_SUCCESS);
+  api_ = HsaRsrcFactory::Instance().AqlProfileApi();;
 
   return true;
 }
 
-TestPMgr::TestPMgr(TestAql* t) : TestAql(t), api_({0}) {
+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_;
-  memset(&api_, 0, sizeof(api_));
 }

test/ctrl/test_pmgr.h

@@ -47,7 +47,7 @@ class TestPMgr : public TestAql {
   hsa_signal_t dummy_signal_;
   hsa_signal_t post_signal_;
 
-  hsa_ven_amd_aqlprofile_1_00_pfn_t api_;
+  HsaRsrcFactory::aqlprofile_pfn_t* api_;
 
   virtual bool BuildPackets() { return false; }
   virtual bool DumpData() { return false; }

test/ctrl/tool.cpp

@@ -320,7 +320,11 @@ CONSTRUCTOR_API void constructor() {
     exit(1);
   }
   printf("ROCProfiler: input from \"%s\"\n", xml_name);
-  xml::Xml* xml = new xml::Xml(xml_name);
+  xml::Xml* xml = xml::Xml::Create(xml_name);
+  if (xml == NULL) {
+    fprintf(stderr, "Input file not found '%s'\n", xml_name);
+    exit(1);
+  }
 
   // Getting metrics
   auto metrics_list = xml->GetNodes("top.metric");

test/util/hsa_rsrc_factory.cpp

@@ -24,6 +24,7 @@ POSSIBILITY OF SUCH DAMAGE.
 
 #include "util/hsa_rsrc_factory.h"
 
+#include <dlfcn.h>
 #include <hsa.h>
 #include <hsa_ext_finalize.h>
 #include <stdint.h>
@@ -116,6 +117,12 @@ HsaRsrcFactory::HsaRsrcFactory() {
   // Discover the set of Gpu devices available on the platform
   status = hsa_iterate_agents(GetHsaAgentsCallback, this);
   CHECK_STATUS("Error Calling hsa_iterate_agents", status);
+
+  // Get AqlProfile API table
+  aqlprofile_api_ = {0};
+  status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &aqlprofile_api_);
+  if (status != HSA_STATUS_SUCCESS) status = LoadAqlProfileLib(&aqlprofile_api_);
+  CHECK_STATUS("aqlprofile API table load failed", status);
 }
 
 // Destructor of the class
@@ -124,6 +131,39 @@ HsaRsrcFactory::~HsaRsrcFactory() {
   CHECK_STATUS("Error in hsa_shut_down", status);
 }
 
+hsa_status_t HsaRsrcFactory::LoadAqlProfileLib(aqlprofile_pfn_t* api) {
+    void* handle = dlopen(kAqlProfileLib, RTLD_NOW);
+    if (handle == NULL) {
+      fprintf(stderr, "Loading '%s' failed, %s\n", kAqlProfileLib, dlerror());
+      return HSA_STATUS_ERROR;
+    }
+    dlerror(); /* Clear any existing error */
+
+    api->hsa_ven_amd_aqlprofile_error_string =
+      (decltype(::hsa_ven_amd_aqlprofile_error_string)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_error_string");
+    api->hsa_ven_amd_aqlprofile_validate_event =
+      (decltype(::hsa_ven_amd_aqlprofile_validate_event)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_validate_event");
+    api->hsa_ven_amd_aqlprofile_start =
+      (decltype(::hsa_ven_amd_aqlprofile_start)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_start");
+    api->hsa_ven_amd_aqlprofile_stop =
+      (decltype(::hsa_ven_amd_aqlprofile_stop)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_stop");
+    api->hsa_ven_amd_aqlprofile_legacy_get_pm4 =
+      (decltype(::hsa_ven_amd_aqlprofile_legacy_get_pm4)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
+    api->hsa_ven_amd_aqlprofile_get_info =
+      (decltype(::hsa_ven_amd_aqlprofile_get_info)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_get_info");
+    api->hsa_ven_amd_aqlprofile_iterate_data =
+      (decltype(::hsa_ven_amd_aqlprofile_iterate_data)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_iterate_data");
+
+  return HSA_STATUS_SUCCESS;
+}
+
 // Get the count of Hsa Gpu Agents available on the platform
 //
 // @return uint32_t Number of Gpu agents on platform

test/util/hsa_rsrc_factory.h

@@ -27,6 +27,7 @@ POSSIBILITY OF SUCH DAMAGE.
 
 #include <hsa.h>
 #include <hsa_ext_finalize.h>
+#include <hsa_ven_amd_aqlprofile.h>
 #include <stdint.h>
 #include <stdio.h>
 #include <stdlib.h>
@@ -54,7 +55,7 @@ static const unsigned MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
 
 // Encapsulates information about a Hsa Agent such as its
 // handle, name, max queue size, max wavefront size, etc.
-typedef struct {
+struct AgentInfo {
   // Handle of Agent
   hsa_agent_t dev_id;
 
@@ -78,8 +79,7 @@ typedef struct {
 
   // Memory region supporting kernel arguments
   hsa_region_t kernarg_region;
-
-} AgentInfo;
+};
 
 class HsaRsrcFactory {
  public:
@@ -207,7 +207,14 @@ class HsaRsrcFactory {
   // Print the various fields of Hsa Gpu Agents
   bool PrintGpuAgents(const std::string& header);
 
+  // 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_; }
+
  private:
+  // Load AQL profile HSA extension library directly
+  static hsa_status_t LoadAqlProfileLib(aqlprofile_pfn_t* api);
+
   // Constructor of the class. Will initialize the Hsa Runtime and
   // query the system topology to get the list of Cpu and Gpu devices
   HsaRsrcFactory();
@@ -229,6 +236,9 @@ class HsaRsrcFactory {
 
   // Used to maintain a list of Hsa Cpu Agent Info
   std::vector<AgentInfo*> cpu_list_;
+
+  // AqlProfile API table
+  aqlprofile_pfn_t aqlprofile_api_;
 };
 
 #endif  // TEST_UTIL_HSA_RSRC_FACTORY_H_

test/util/xml.h

@@ -27,6 +27,33 @@ class Xml {
 
   enum { DECL_STATE, BODY_STATE };
 
+  static Xml* Create(const char* file_name) {
+    Xml* xml = new Xml(file_name);
+    if (xml->fd_ == -1) {
+      delete xml;
+      xml = NULL;
+    }
+    return xml;
+  }
+
+  static void Destroy(Xml *xml) { delete xml; }
+
+  std::vector<level_t*> GetNodes(std::string global_tag) { return map_[global_tag]; }
+
+  void Print() const {
+    for (auto& elem : map_) {
+      for (auto node : elem.second) {
+        if (node->opts.size()) {
+          std::cout << elem.first << ":" << std::endl;
+          for (auto& opt : node->opts) {
+            std::cout << "  " << opt.first << " = " << opt.second << std::endl;
+          }
+        }
+      }
+    }
+  }
+
+ private:
   Xml(const char* file_name)
       : file_name_(file_name),
         file_line_(0),
@@ -39,7 +66,7 @@ class Xml {
 
     fd_ = open(file_name, O_RDONLY);
     if (fd_ == -1) {
-      std::cout << "XML file not found: " << file_name << std::endl;
+      perror("open XML file");
       return;
     }
 
@@ -117,22 +144,8 @@ class Xml {
     }
   }
 
-  std::vector<level_t*> GetNodes(std::string global_tag) { return map_[global_tag]; }
+  ~Xml() {}
 
-  void Print() const {
-    for (auto& elem : map_) {
-      for (auto node : elem.second) {
-        if (node->opts.size()) {
-          std::cout << elem.first << ":" << std::endl;
-          for (auto& opt : node->opts) {
-            std::cout << "  " << opt.first << " = " << opt.second << std::endl;
-          }
-        }
-      }
-    }
-  }
-
- private:
   bool LineEndCheck() {
     bool found = false;
     if (buffer_[index_] == '\n') {