profiling controll options

test/util/hsa_rsrc_factory.cpp

@@ -25,6 +25,7 @@ POSSIBILITY OF SUCH DAMAGE.
 #include "util/hsa_rsrc_factory.h"
 
 #include <dlfcn.h>
+#include <fcntl.h>
 #include <hsa.h>
 #include <hsa_ext_amd.h>
 #include <hsa_ext_finalize.h>
@@ -32,6 +33,8 @@ POSSIBILITY OF SUCH DAMAGE.
 #include <stdio.h>
 #include <string.h>
 #include <stdlib.h>
+#include <sys/stat.h>
+#include <sys/types.h>
 
 #include <atomic>
 #include <cassert>
@@ -77,12 +80,13 @@ hsa_status_t HsaRsrcFactory::FindMemRegionsCallback(hsa_region_t region, void* d
 }
 
 // Constructor of the class
-HsaRsrcFactory::HsaRsrcFactory() {
+HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize_hsa) {
+  hsa_status_t status;
   // Initialize the Hsa Runtime
-  printf("HSA init\n");
-  hsa_status_t status = hsa_init();
-  CHECK_STATUS("Error in hsa_init", status);
-
+  if (initialize_hsa_) {
+    status = hsa_init();
+    CHECK_STATUS("Error in hsa_init", status);
+  }
   // Discover the set of Gpu devices available on the platform
   status = hsa_iterate_agents(GetHsaAgentsCallback, this);
   CHECK_STATUS("Error Calling hsa_iterate_agents", status);
@@ -106,10 +110,10 @@ HsaRsrcFactory::HsaRsrcFactory() {
 HsaRsrcFactory::~HsaRsrcFactory() {
   for (auto p : cpu_list_) delete p;
   for (auto p : gpu_list_) delete p;
-
-  printf("HSA shutdown\n");
-  hsa_status_t status = hsa_shut_down();
-  CHECK_STATUS("Error in hsa_shut_down", status);
+  if (initialize_hsa_) {
+    hsa_status_t status = hsa_shut_down();
+    CHECK_STATUS("Error in hsa_shut_down", status);
+  }
 }
 
 hsa_status_t HsaRsrcFactory::LoadAqlProfileLib(aqlprofile_pfn_t* api) {
@@ -365,67 +369,53 @@ bool HsaRsrcFactory::TransferData(void* dest_buff, void* src_buff, uint32_t leng
 //
 // @return bool true if successful, false otherwise
 //
-void* HsaRsrcFactory::LoadAndFinalize(const 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;
+bool HsaRsrcFactory::LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
+                                      const char* kernel_name, hsa_executable_t* executable, hsa_executable_symbol_t* code_desc) {
+  hsa_status_t status = HSA_STATUS_ERROR;
 
   // Build the code object filename
   std::string filename(brig_path);
   std::clog << "Code object filename: " << filename << std::endl;
 
   // Open the file containing code object
-  std::ifstream codeStream(filename.c_str(), std::ios::binary | std::ios::ate);
-  if (!codeStream) {
-    std::cerr << "Error: failed to load " << filename << std::endl;
+  hsa_file_t file_handle = open(filename.c_str(), O_RDONLY);
+  if (file_handle == -1) {
+    std::cerr << "Error: failed to load '" << filename << "'" << std::endl;
     assert(false);
-    return NULL;
+    return false;
   }
 
-  // Allocate memory to read in code object from file
-  size_t size = std::string::size_type(codeStream.tellg());
-  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 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>(), code_buf);
-
-  // De-Serialize the code object that has been read into memory
-  status = hsa_code_object_deserialize(code_buf, size, NULL, &code_object);
+  // Create code object reader
+  hsa_code_object_reader_t code_obj_rdr = {0};
+  status = hsa_code_object_reader_create_from_file(file_handle, &code_obj_rdr);
   if (status != HSA_STATUS_SUCCESS) {
-    std::cerr << "Failed to deserialize code object" << std::endl;
-    if (code_buf) hsa_memory_free(code_buf);
-    return NULL;
+    std::cerr << "Failed to create code object reader '" << filename << "'" << std::endl;
+    return false;
   }
 
   // Create executable.
-  status =
-      hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", hsa_exec);
+  status = hsa_executable_create_alt(HSA_PROFILE_FULL,
+    HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, NULL, executable);
   CHECK_STATUS("Error in creating executable object", status);
 
   // Load code object.
-  status = hsa_executable_load_code_object(*hsa_exec, agent_info->dev_id, code_object, "");
+  status = hsa_executable_load_agent_code_object(*executable, agent_info->dev_id,
+    code_obj_rdr, NULL, NULL);
   CHECK_STATUS("Error in loading executable object", status);
 
   // Freeze executable.
-  status = hsa_executable_freeze(*hsa_exec, "");
+  status = hsa_executable_freeze(*executable, "");
   CHECK_STATUS("Error in freezing executable object", status);
 
   // Get symbol handle.
   hsa_executable_symbol_t kernelSymbol;
-  status = hsa_executable_get_symbol(*hsa_exec, NULL, kernel_name, agent_info->dev_id, 0,
+  status = hsa_executable_get_symbol(*executable, 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 code_buf;
+  return true;
 }
 
 // Print the various fields of Hsa Gpu Agents

test/util/hsa_rsrc_factory.h

@@ -112,20 +112,18 @@ class HsaRsrcFactory {
  public:
   typedef std::recursive_mutex mutex_t;
 
-  static HsaRsrcFactory* Create() { return NULL; }
-
-  static HsaRsrcFactory* CreateInstance() {
+  static HsaRsrcFactory* Create(bool initialize_hsa = true) {
     std::lock_guard<mutex_t> lck(mutex_);
     if (instance_ == NULL) {
-      instance_ = new HsaRsrcFactory();
+      instance_ = new HsaRsrcFactory(initialize_hsa);
     }
     return instance_;
   }
 
   static HsaRsrcFactory& Instance() {
-    CreateInstance();
+    if (instance_ == NULL) instance_ = Create(false);
     hsa_status_t status = (instance_ != NULL) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
-    CHECK_STATUS("HsaRsrcFactory::Instance() is not found", status);
+    CHECK_STATUS("HsaRsrcFactory::Instance() failed", status);
     return *instance_;
   }
 
@@ -229,9 +227,9 @@ class HsaRsrcFactory {
   // @param code_desc Handle of finalized Code Descriptor that could
   // be used to submit for execution
   //
-  // @return code buffer, non NULL if successful, NULL otherwise
+  // @return true if successful, false otherwise
   //
-  void* LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path, const char* kernel_name,
+  bool LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path, const char* kernel_name,
                         hsa_executable_t* hsa_exec, hsa_executable_symbol_t* code_desc);
 
   // Print the various fields of Hsa Gpu Agents
@@ -259,11 +257,14 @@ class HsaRsrcFactory {
 
   // Constructor of the class. Will initialize the Hsa Runtime and
   // query the system topology to get the list of Cpu and Gpu devices
-  HsaRsrcFactory();
+  HsaRsrcFactory(bool initialize_hsa);
 
   // Destructor of the class
   ~HsaRsrcFactory();
 
+  // HSA was initialized
+  const bool initialize_hsa_;
+
   // Add an instance of AgentInfo representing a Hsa Gpu agent
   const AgentInfo* AddAgentInfo(const hsa_agent_t agent);