Initial commit

src/CMakeLists.txt

@@ -0,0 +1,11 @@
+#
+# Build dynamic Library object
+#
+set ( TARGET_LIB "${TARGET_NAME}" )
+set ( LIB_SRC
+  ${LIB_DIR}/core/roctracer.cpp
+  ${LIB_DIR}/util/hsa_rsrc_factory.cpp
+)
+add_library ( ${TARGET_LIB} SHARED ${LIB_SRC} )
+target_include_directories ( ${TARGET_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${HIP_INC_DIR} )
+target_link_libraries( ${TARGET_LIB} PRIVATE ${HSA_RUNTIME_LIB} c stdc++)

src/core/roctracer.cpp

@@ -0,0 +1,420 @@
+#include "inc/roctracer.h"
+
+#include <atomic>
+#include <hip/hip_runtime.h>
+#include <mutex>
+#include <string.h>
+#include <pthread.h>
+
+#include "util/exception.h"
+#include "util/hsa_rsrc_factory.h"
+#include "util/logger.h"
+
+#define PUBLIC_API __attribute__((visibility("default")))
+#define CONSTRUCTOR_API __attribute__((constructor))
+#define DESTRUCTOR_API __attribute__((destructor))
+
+#define PTHREAD_CALL(call)                                                                         \
+  do {                                                                                             \
+    int err = call;                                                                                \
+    if (err != 0) {                                                                                \
+      errno = err;                                                                                 \
+      perror(#call);                                                                               \
+      abort();                                                                                     \
+    }                                                                                              \
+  } while (0)
+
+#define HSART_CALL(call)                                                                           \
+  do {                                                                                             \
+    hsa_status_t status = call;                                                                    \
+    if (status != HSA_STATUS_SUCCESS) {                                                            \
+      std::cerr << "HSA-rt call '" << #call << "' error(" << std::hex << status << ")"             \
+        << std::dec << std::endl << std::flush;                                                    \
+      abort();                                                                                     \
+    }                                                                                              \
+  } while (0)
+
+#define API_METHOD_PREFIX                                                                          \
+  int err = 0;                                                                                     \
+  try {
+
+#define API_METHOD_SUFFIX                                                                          \
+  }                                                                                                \
+  catch (std::exception & e) {                                                                     \
+    ERR_LOGGING(__FUNCTION__ << "(), " << e.what());                                               \
+    err = roctracer::GetExcStatus(e);                                                              \
+  }                                                                                                \
+  return err;
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+// Internal library methods
+//
+namespace roctracer {
+
+int GetExcStatus(const std::exception& e) {
+  const util::exception* roctracer_exc_ptr = dynamic_cast<const util::exception*>(&e);
+  return (roctracer_exc_ptr) ? static_cast<int>(roctracer_exc_ptr->status()) : 1;
+}
+
+class GlobalCounter {
+  public:
+  typedef std::mutex mutex_t;
+  typedef uint64_t counter_t;
+
+  static counter_t Increment() {
+    std::lock_guard<mutex_t> lock(mutex_);
+    return ++counter_;
+  }
+
+  private:
+  static mutex_t mutex_;
+  static counter_t counter_;
+};
+GlobalCounter::mutex_t GlobalCounter::mutex_;
+GlobalCounter::counter_t GlobalCounter::counter_ = 0;
+
+class MemoryPool {
+  public:
+  typedef std::mutex mutex_t;
+
+  static void allocator_default(char** ptr, size_t size, void* arg) {
+    (void)arg;
+    if (*ptr == NULL) {
+      *ptr = reinterpret_cast<char*>(malloc(size));
+    } else if (size != 0) {
+      *ptr = reinterpret_cast<char*>(realloc(ptr, size));
+    } else {
+      free(*ptr); 
+      *ptr = NULL;
+    }
+  }
+
+  MemoryPool(const roctracer_properties_t& properties) { 
+    // Assigning pool allocator
+    alloc_fun_ = allocator_default;
+    alloc_arg_ = NULL;
+    if (properties.alloc_fun != NULL) {
+      alloc_fun_ = properties.alloc_fun;
+      alloc_arg_ = properties.alloc_arg;
+    }
+
+    // Pool definition
+    buffer_size_ = properties.buffer_size;
+    const size_t pool_size = 2 * buffer_size_;
+    pool_begin_ = NULL;
+    alloc_fun_(&pool_begin_, pool_size, alloc_arg_);
+    if (pool_begin_ == NULL) EXC_ABORT(ROCTRACER_STATUS_ERROR, "pool allocator failed");
+    pool_end_ = pool_begin_ + pool_size;
+    buffer_begin_ = pool_begin_;
+    buffer_end_ = buffer_begin_ + buffer_size_;
+    write_ptr_ = buffer_begin_;
+
+    // Consuming read thread
+    read_callback_fun_ = properties.buffer_callback_fun;
+    read_callback_arg_ = properties.buffer_callback_arg;
+    consumer_arg_ = consumer_arg_t{this, true, NULL, NULL};
+    PTHREAD_CALL(pthread_mutex_init(&read_mutex_, NULL));
+    PTHREAD_CALL(pthread_cond_init(&read_cond_, NULL));
+    PTHREAD_CALL(pthread_create(&consumer_thread_, NULL, reader_fun, &consumer_arg_));
+  }
+
+  ~MemoryPool() {
+    Flush();
+    allocator_default(&pool_begin_, 0, alloc_arg_);
+  }
+
+  template <typename Record>
+  void* Write(const Record& record) {
+    std::lock_guard<mutex_t> lock(write_mutex_);
+    char* next = write_ptr_ + sizeof(Record);
+    if (next > buffer_end_) {
+      if (write_ptr_ == buffer_begin_) EXC_ABORT(ROCTRACER_STATUS_ERROR, "buffer size(" << buffer_size_ << ") is less then the record(" << sizeof(Record) << ")");
+      spawn_reader(buffer_begin_, buffer_end_);
+      buffer_begin_ = (buffer_end_ == pool_end_) ? pool_begin_ : buffer_end_;
+      buffer_end_ = buffer_begin_ + buffer_size_;
+      write_ptr_ = buffer_begin_;
+      next = write_ptr_ + sizeof(Record);
+    }
+    Record* ptr = reinterpret_cast<Record*>(write_ptr_);
+    *ptr = record;
+    write_ptr_ = next;
+    return reinterpret_cast<void*>(ptr);
+  }
+
+  void Flush() {
+    if (write_ptr_ > buffer_begin_) {
+      spawn_reader(buffer_begin_, write_ptr_);
+      sync_reader(&consumer_arg_);
+      buffer_begin_ = write_ptr_;
+    }
+  }
+
+  private:
+  struct consumer_arg_t {
+    MemoryPool* obj;
+    bool valid;
+    const char* begin;
+    const char* end;
+  };
+
+  static void reset_reader(consumer_arg_t* arg) {
+    reinterpret_cast<std::atomic<bool>*>(&(arg->valid))->store(false, std::memory_order_release);
+  }
+
+  static void sync_reader(const consumer_arg_t* arg) {
+    while(arg->valid) PTHREAD_CALL(pthread_yield());
+  }
+
+  static void* reader_fun(void* consumer_arg) {
+    consumer_arg_t* arg = reinterpret_cast<consumer_arg_t*>(consumer_arg);
+    roctracer::MemoryPool* obj = arg->obj;
+
+    reset_reader(arg);
+
+    while (1) {
+      PTHREAD_CALL(pthread_mutex_lock(&(obj->read_mutex_)));
+      while (arg->valid == false) {
+        PTHREAD_CALL(pthread_cond_wait(&(obj->read_cond_), &(obj->read_mutex_)));
+      }
+      obj->read_callback_fun_(arg->begin, arg->end, obj->read_callback_arg_);
+      reset_reader(arg);
+      PTHREAD_CALL(pthread_mutex_unlock(&(obj->read_mutex_)));
+    }
+
+    return NULL;
+  }
+
+  void spawn_reader(const char* data_begin, const char* data_end) {
+    sync_reader(&consumer_arg_);
+    PTHREAD_CALL(pthread_mutex_lock(&read_mutex_));
+    consumer_arg_ = consumer_arg_t{this, true, data_begin, data_end};
+    PTHREAD_CALL(pthread_cond_signal(&read_cond_));
+    PTHREAD_CALL(pthread_mutex_unlock(&read_mutex_));
+  }
+
+  // pool allocator
+  roctracer_allocator_t alloc_fun_;
+  void* alloc_arg_;
+
+  // Pool definition
+  size_t buffer_size_;
+  char* pool_begin_;
+  char* pool_end_;
+  char* buffer_begin_;
+  char* buffer_end_;
+  char* write_ptr_;
+  mutex_t write_mutex_;
+
+  // Consuming read thread
+  roctracer_buffer_callback_t read_callback_fun_;
+  void* read_callback_arg_;
+  consumer_arg_t consumer_arg_;
+  pthread_t consumer_thread_;
+  pthread_mutex_t read_mutex_;
+  pthread_cond_t read_cond_;
+};
+
+class Timer {
+  public:
+  typedef uint64_t timestamp_t;
+  typedef long double freq_t;
+  
+  Timer() {
+    timestamp_t timestamp_hz = 0;
+    HSART_CALL(hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &timestamp_hz));
+    timestamp_factor_ = (freq_t)1000000000 / (freq_t)timestamp_hz;
+  }
+
+  // Return timestamp in 'ns'
+  timestamp_t timestamp_ns() {
+    timestamp_t timestamp;
+    HSART_CALL(hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &timestamp));
+    return timestamp_t((freq_t)timestamp * timestamp_factor_);
+  }
+
+  private:
+  // Timestamp frequency factor
+  freq_t timestamp_factor_;
+};
+
+CONSTRUCTOR_API void constructor() {
+  util::Logger::Create();
+}
+
+DESTRUCTOR_API void destructor() {
+  util::HsaRsrcFactory::Destroy();
+  util::Logger::Destroy();
+}
+
+// Activity callback to generate an activity record
+void ActivityCallback(
+    roctracer_record_t* record,
+    uint32_t activity_kind,
+    const void* callback_data,
+    void* arg)
+{
+  static Timer timer;
+
+  const hip_cb_data_t* data = reinterpret_cast<const hip_cb_data_t*>(callback_data);
+  MemoryPool* pool = reinterpret_cast<MemoryPool*>(arg);
+  if (pool == NULL) EXC_ABORT(ROCTRACER_STATUS_ERROR, "ActivityCallback pool is NULL");
+  if (data->phase == ROCTRACER_API_PHASE_ENTER) {
+    *record = {};
+    record->name = data->name;
+    record->activity_kind = activity_kind;
+    record->begin_ns = timer.timestamp_ns();
+    // Correlation ID generating
+    const auto correlation_id = GlobalCounter::Increment();
+    record->correlation_id = correlation_id;
+    const_cast<hip_cb_data_t*>(data)->correlation_id = correlation_id;
+  } else {
+    record->end_ns = timer.timestamp_ns();
+    pool->Write<roctracer_record_t>(*record);
+  }
+}
+
+util::Logger::mutex_t util::Logger::mutex_;
+util::Logger* util::Logger::instance_ = NULL;
+MemoryPool* memory_pool = NULL;
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+// Public library methods
+//
+extern "C" {
+
+// Returns library vesrion
+PUBLIC_API uint32_t roctracer_version_major() { return ROCTRACER_VERSION_MAJOR; }
+PUBLIC_API uint32_t roctracer_version_minor() { return ROCTRACER_VERSION_MINOR; }
+
+// Returns the last error
+PUBLIC_API const char* roctracer_error_string() {
+  return strdup(roctracer::util::Logger::LastMessage().c_str());
+}
+
+// Return method name by given API domain and call ID
+// NULL returned on the error and the library errno is set
+PUBLIC_API const char* roctracer_get_api_name(roctracer_api_domain_t domain, roctracer_hip_api_cid_t cid) {
+  return NULL;
+}
+
+// Enable runtime API callbacks
+PUBLIC_API int roctracer_enable_api_callback(
+    roctracer_api_domain_t domain,
+    uint32_t cid,
+    roctracer_api_callback_t callback,
+    void* user_data)
+{
+  API_METHOD_PREFIX
+  switch (domain) {
+    case ROCTRACER_API_DOMAIN_HIP: {
+      hipError_t hip_err = hipRegisterApiCallback(cid, callback, user_data);
+      if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRegisterApiCallback error(" << hip_err << ")");
+      break;
+    }
+    default:
+      EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "invalid domain ID(" << domain << ")");
+  }
+  API_METHOD_SUFFIX
+}
+
+// Enable runtime API callbacks
+PUBLIC_API int roctracer_disable_api_callback(
+    roctracer_api_domain_t domain,
+    uint32_t cid)
+{
+  API_METHOD_PREFIX
+  switch (domain) {
+    case ROCTRACER_API_DOMAIN_HIP: {
+      hipError_t hip_err = hipRemoveApiCallback(cid);
+      if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRemoveApiCallback error(" << hip_err << ")");
+      break;
+    }
+    default:
+      EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "invalid domain ID(" << domain << ")");
+  }
+  API_METHOD_SUFFIX
+}
+
+// Return default pool and set new one if parameter pool is not NULL.
+roctracer_pool_t* roctracer_default_pool(roctracer_pool_t* pool) {
+  roctracer_pool_t* p = reinterpret_cast<roctracer_pool_t*>(roctracer::memory_pool);
+  if (pool != NULL) roctracer::memory_pool = reinterpret_cast<roctracer::MemoryPool*>(pool);
+  if (p == NULL) EXC_RAISING(ROCTRACER_STATUS_UNINIT, "default pool is not initialized");
+  return p;
+}
+
+// Open memory pool
+PUBLIC_API int roctracer_open_pool(
+    const roctracer_properties_t* properties,
+    roctracer_pool_t** pool)
+{
+  API_METHOD_PREFIX
+  if ((pool == NULL) && (roctracer::memory_pool != NULL)) {
+    EXC_RAISING(ROCTRACER_STATUS_ERROR, "default pool already set");
+  }
+  roctracer::MemoryPool* p = new roctracer::MemoryPool(*properties);
+  if (p == NULL) EXC_RAISING(ROCTRACER_STATUS_ERROR, "MemoryPool() error");
+  if (pool != NULL) *pool = p;
+  else roctracer::memory_pool = p;
+  API_METHOD_SUFFIX
+}
+
+// Close memory pool
+PUBLIC_API int roctracer_close_pool(roctracer_pool_t* pool) {
+  API_METHOD_PREFIX
+  roctracer_pool_t* ptr = (pool == NULL) ? roctracer_default_pool() : pool;
+  roctracer::MemoryPool* memory_pool = reinterpret_cast<roctracer::MemoryPool*>(ptr);
+  delete(memory_pool);
+  if (pool == NULL) roctracer::memory_pool = NULL;
+  API_METHOD_SUFFIX
+}
+
+// Enable activity records logging
+PUBLIC_API int roctracer_enable_api_activity(
+    roctracer_api_domain_t domain,
+    uint32_t activity_kind,
+    roctracer_pool_t* pool)
+{
+  API_METHOD_PREFIX
+  if (pool == NULL) pool = roctracer_default_pool();
+  switch (domain) {
+    case ROCTRACER_API_DOMAIN_HIP: {
+      const hipError_t hip_err = hipRegisterActivityCallback(activity_kind, roctracer::ActivityCallback, pool);
+      if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRegisterActivityCallback error(" << hip_err << ")");
+      break;
+    }
+    default:
+      EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "invalid domain ID(" << domain << ")");
+  }
+  API_METHOD_SUFFIX
+}
+
+// Disable activity records logging
+PUBLIC_API int roctracer_disable_api_activity(
+    roctracer_api_domain_t domain,
+    uint32_t activity_kind)
+{
+  API_METHOD_PREFIX
+  switch (domain) {
+    case ROCTRACER_API_DOMAIN_HIP: {
+      const hipError_t hip_err = hipRemoveActivityCallback(activity_kind);
+      if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRemoveActivityCallback error(" << hip_err << ")");
+      break;
+    }
+    default:
+      EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "invalid domain ID(" << domain << ")");
+  }
+  API_METHOD_SUFFIX
+}
+
+// Flush available activity records
+PUBLIC_API int roctracer_flush_api_activity(roctracer_pool_t* pool) {
+  API_METHOD_PREFIX
+  if (pool == NULL) pool = roctracer_default_pool();
+  roctracer::MemoryPool* memory_pool = reinterpret_cast<roctracer::MemoryPool*>(pool);
+  memory_pool->Flush();
+  API_METHOD_SUFFIX
+}
+
+}  // extern "C"

src/util/exception.h

@@ -0,0 +1,47 @@
+#ifndef SRC_UTIL_EXCEPTION_H_
+#define SRC_UTIL_EXCEPTION_H_
+
+#include <hsa_ven_amd_aqlprofile.h>
+
+#include <exception>
+#include <sstream>
+#include <string>
+
+#define EXC_ABORT(error, stream)                                                                   \
+  {                                                                                                \
+    std::ostringstream oss;                                                                        \
+    oss << __FUNCTION__ << "(), " << stream;                                                       \
+    std::cout << oss.str() << std::endl;                                                           \
+    abort();                                                                                       \
+  }
+
+#define EXC_RAISING(error, stream)                                                                 \
+  {                                                                                                \
+    std::ostringstream oss;                                                                        \
+    oss << __FUNCTION__ << "(), " << stream;                                                       \
+    throw roctracer::util::exception(error, oss.str());                                            \
+  }
+
+#define HIP_EXC_RAISING(error, stream)                                                             \
+  {                                                                                                \
+    EXC_RAISING(error, "HIP error: " << stream);                                            \
+  }
+
+namespace roctracer {
+namespace util {
+
+class exception : public std::exception {
+ public:
+  explicit exception(const uint32_t& status, const std::string& msg) : status_(status), str_(msg) {}
+  const char* what() const throw() { return str_.c_str(); }
+  uint32_t status() const throw() { return status_; }
+
+ protected:
+  const uint32_t status_;
+  const std::string str_;
+};
+
+}  // namespace util
+}  // namespace roctracer
+
+#endif  // SRC_UTIL_EXCEPTION_H_

src/util/hsa_rsrc_factory.cpp

@@ -0,0 +1,572 @@
+/**********************************************************************
+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:
+
+<95>    Redistributions of source code must retain the above copyright notice, this list of
+conditions and the following disclaimer.
+<95>    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 "util/hsa_rsrc_factory.h"
+
+#include <dlfcn.h>
+#include <fcntl.h>
+#include <hsa.h>
+#include <hsa_ext_amd.h>
+#include <hsa_ext_finalize.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <stdlib.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+
+#include <atomic>
+#include <cassert>
+#include <fstream>
+#include <iostream>
+#include <string>
+#include <vector>
+
+#ifndef AQL_PROFILE_READ_API_ENABLE
+#define AQL_PROFILE_READ_API_ENABLE 0
+#endif
+
+namespace roctracer {
+namespace util {
+
+// Callback function to get available in the system agents
+hsa_status_t HsaRsrcFactory::GetHsaAgentsCallback(hsa_agent_t agent, void* data) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  HsaRsrcFactory* hsa_rsrc = reinterpret_cast<HsaRsrcFactory*>(data);
+  const AgentInfo* agent_info = hsa_rsrc->AddAgentInfo(agent);
+  if (agent_info != NULL) status = HSA_STATUS_SUCCESS;
+  return status;
+}
+
+// This function checks to see if the provided
+// pool has the HSA_AMD_SEGMENT_GLOBAL property. If the kern_arg flag is true,
+// the function adds an additional requirement that the pool have the
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT property. If kern_arg is false,
+// pools must NOT have this property.
+// Upon finding a pool that meets these conditions, HSA_STATUS_INFO_BREAK is
+// returned. HSA_STATUS_SUCCESS is returned if no errors were encountered, but
+// no pool was found meeting the requirements. If an error is encountered, we
+// return that error.
+static hsa_status_t
+FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool kern_arg) {
+  hsa_status_t err;
+  hsa_amd_segment_t segment;
+  uint32_t flag;
+
+  if (nullptr == data) {
+    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+  }
+
+  err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+                                     &segment);
+  CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
+  if (HSA_AMD_SEGMENT_GLOBAL != segment) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  err = hsa_amd_memory_pool_get_info(pool,
+                                HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
+  CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
+
+  uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
+
+  if ((karg_st == 0 && kern_arg) ||
+      (karg_st != 0 && !kern_arg)) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  *(reinterpret_cast<hsa_amd_memory_pool_t*>(data)) = pool;
+  return HSA_STATUS_INFO_BREAK;
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that is NOT
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t FindStandardPool(hsa_amd_memory_pool_t pool, void* data) {
+  return FindGlobalPool(pool, data, false);
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that IS
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t FindKernArgPool(hsa_amd_memory_pool_t pool, void* data) {
+  return FindGlobalPool(pool, data, true);
+}
+#if 0
+// Callback function to find and bind kernarg region of an agent
+hsa_status_t HsaRsrcFactory::FindMemRegionsCallback(hsa_region_t region, void* data) {
+  hsa_region_global_flag_t flags;
+  hsa_region_segment_t segment_id;
+
+  hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
+  if (segment_id != HSA_REGION_SEGMENT_GLOBAL) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  AgentInfo* agent_info = (AgentInfo*)data;
+  hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
+  if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
+    agent_info->coarse_region = region;
+  }
+
+  if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) {
+    agent_info->kernarg_region = region;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+#endif
+// Constructor of the class
+HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize_hsa) {
+  hsa_status_t status;
+  // Initialize the Hsa Runtime
+  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);
+
+  // Get AqlProfile API table
+  aqlprofile_api_ = {0};
+#ifdef ROCP_LD_AQLPROFILE
+  status = LoadAqlProfileLib(&aqlprofile_api_);
+#else
+  status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &aqlprofile_api_);
+#endif
+  CHECK_STATUS("aqlprofile API table load failed", status);
+
+  // Get Loader API table
+  loader_api_ = {0};
+  status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_LOADER, 1, 0, &loader_api_);
+  CHECK_STATUS("loader API table query failed", status);
+}
+
+// Destructor of the class
+HsaRsrcFactory::~HsaRsrcFactory() {
+  for (auto p : cpu_list_) delete p;
+  for (auto p : gpu_list_) delete p;
+  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) {
+    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");
+#if AQL_PROFILE_READ_API_ENABLE
+    api->hsa_ven_amd_aqlprofile_read =
+      (decltype(::hsa_ven_amd_aqlprofile_read)*)
+        dlsym(handle, "hsa_ven_amd_aqlprofile_read");
+#endif  // AQL_PROFILE_READ_API_ENABLE
+    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;
+}
+
+// Add system agent info
+const AgentInfo* HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent) {
+  // Determine if device is a Gpu agent
+  hsa_status_t status;
+  AgentInfo* agent_info = NULL;
+
+  hsa_device_type_t type;
+  status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
+  CHECK_STATUS("Error Calling hsa_agent_get_info", status);
+
+  if (type == HSA_DEVICE_TYPE_CPU) {
+    agent_info = new AgentInfo{};
+    agent_info->dev_id = agent;
+    agent_info->dev_type = HSA_DEVICE_TYPE_CPU;
+    agent_info->dev_index = cpu_list_.size();
+
+    status = hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->cpu_pool);
+    CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(cpu pool)", status);
+    status = hsa_amd_agent_iterate_memory_pools(agent, FindKernArgPool, &agent_info->kern_arg_pool);
+    CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(kern arg pool)", status);
+    agent_info->gpu_pool = {};
+
+    cpu_list_.push_back(agent_info);
+    cpu_agents_.push_back(agent);
+  }
+
+  if (type == HSA_DEVICE_TYPE_GPU) {
+    agent_info = new AgentInfo{};
+    agent_info->dev_id = agent;
+    agent_info->dev_type = HSA_DEVICE_TYPE_GPU;
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name);
+    strncpy(agent_info->gfxip, agent_info->name, 4);
+    agent_info->gfxip[4] = '\0';
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE, &agent_info->max_wave_size);
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &agent_info->max_queue_size);
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
+    agent_info->is_apu = (agent_info->profile == HSA_PROFILE_FULL) ? true : false;
+    hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT), &agent_info->cu_num);
+    hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU), &agent_info->waves_per_cu);
+    hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU), &agent_info->simds_per_cu);
+    hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES), &agent_info->se_num);
+    hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE), &agent_info->shader_arrays_per_se);
+
+    agent_info->cpu_pool = {};
+    agent_info->kern_arg_pool = {};
+    status = hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->gpu_pool);
+    CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(gpu pool)", status);
+#if 0
+    // Initialize memory regions to zero
+    agent_info->kernarg_region.handle = 0;
+    agent_info->coarse_region.handle = 0;
+    // Find and Bind Memory regions of the Gpu agent
+    hsa_agent_iterate_regions(agent, FindMemRegionsCallback, agent_info);
+#endif
+
+    // Set GPU index
+    agent_info->dev_index = gpu_list_.size();
+    gpu_list_.push_back(agent_info);
+    gpu_agents_.push_back(agent);
+  }
+
+  if (agent_info) agent_map_[agent.handle] = agent_info;
+
+  return agent_info;
+}
+
+// Return systen agent info
+const AgentInfo* HsaRsrcFactory::GetAgentInfo(const hsa_agent_t agent) {
+  const AgentInfo* agent_info = NULL;
+  auto it = agent_map_.find(agent.handle);
+  if (it != agent_map_.end()) {
+    agent_info = it->second;
+  }
+  return agent_info;
+}
+
+// Get the count of Hsa Gpu Agents available on the platform
+//
+// @return uint32_t Number of Gpu agents on platform
+//
+uint32_t HsaRsrcFactory::GetCountOfGpuAgents() { return uint32_t(gpu_list_.size()); }
+
+// Get the count of Hsa Cpu Agents available on the platform
+//
+// @return uint32_t Number of Cpu agents on platform
+//
+uint32_t HsaRsrcFactory::GetCountOfCpuAgents() { return uint32_t(cpu_list_.size()); }
+
+// Get the AgentInfo handle of a Gpu device
+//
+// @param idx Gpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info) {
+  // Determine if request is valid
+  uint32_t size = uint32_t(gpu_list_.size());
+  if (idx >= size) {
+    return false;
+  }
+
+  // Copy AgentInfo from specified index
+  *agent_info = gpu_list_[idx];
+
+  return true;
+}
+
+// Get the AgentInfo handle of a Cpu device
+//
+// @param idx Cpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info) {
+  // Determine if request is valid
+  uint32_t size = uint32_t(cpu_list_.size());
+  if (idx >= size) {
+    return false;
+  }
+
+  // Copy AgentInfo from specified index
+  *agent_info = cpu_list_[idx];
+  return true;
+}
+
+// Create a Queue object and return its handle. The queue object is expected
+// to support user requested number of Aql dispatch packets.
+//
+// @param agent_info Gpu Agent on which to create a queue object
+//
+// @param num_Pkts Number of packets to be held by queue
+//
+// @param queue Output parameter updated with handle of queue object
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts,
+                                 hsa_queue_t** queue) {
+  hsa_status_t status;
+  status = hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, NULL,
+                            UINT32_MAX, UINT32_MAX, queue);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+// Create a Signal object and return its handle.
+// @param value Initial value of signal object
+// @param signal Output parameter updated with handle of signal object
+// @return bool true if successful, false otherwise
+bool HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t* signal) {
+  hsa_status_t status;
+  status = hsa_signal_create(value, 0, NULL, signal);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+// Allocate memory for use by a kernel of specified size in specified
+// agent's memory region.
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t* HsaRsrcFactory::AllocateLocalMemory(const AgentInfo* agent_info, size_t size) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  uint8_t* buffer = NULL;
+  size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+  status = hsa_amd_memory_pool_allocate(agent_info->gpu_pool, size, 0, (void**)&buffer);
+  // Only GPU can access the memory
+  if (status == HSA_STATUS_SUCCESS) {
+    hsa_agent_t agents_list[1] = {agent_info->dev_id};
+    status = hsa_amd_agents_allow_access(1, agents_list, NULL, buffer);
+  }
+  uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+  printf("AllocateLocalMemory %p\n", ptr);
+  return ptr;
+}
+
+// Allocate memory to pass kernel parameters.
+// Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo parameter.
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t* HsaRsrcFactory::AllocateKernArgMemory(const AgentInfo* agent_info, size_t size) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  uint8_t* buffer = NULL;
+  if (!cpu_agents_.empty()) {
+    size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+    status = hsa_amd_memory_pool_allocate(cpu_list_[0]->kern_arg_pool, size, 0, (void**)&buffer);
+    // Both the CPU and GPU can access the kernel arguments
+    if (status == HSA_STATUS_SUCCESS) {
+      auto agents_vec = cpu_agents_;
+      agents_vec.push_back(agent_info->dev_id);
+      status = hsa_amd_agents_allow_access(agents_vec.size(), &agents_vec[0], NULL, buffer);
+    }
+  }
+  uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+  printf("AllocateKernargMemory %p\n", ptr);
+  return ptr;
+}
+
+// Allocate system memory accessible by both CPU and GPU
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t* HsaRsrcFactory::AllocateSysMemory(const AgentInfo* agent_info, size_t size) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  uint8_t* buffer = NULL;
+  if (!cpu_agents_.empty()) {
+    size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+    status = hsa_amd_memory_pool_allocate(cpu_list_[0]->cpu_pool, size, 0, (void**)&buffer);
+    // Both the CPU and GPU can access the memory
+    if (status == HSA_STATUS_SUCCESS) {
+      auto agents_vec = cpu_agents_;
+      agents_vec.push_back(agent_info->dev_id);
+      status = hsa_amd_agents_allow_access(agents_vec.size(), &agents_vec[0], NULL, buffer);
+    }
+  }
+  uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+  printf("AllocateSysMemory %p\n", ptr);
+  return ptr;
+}
+
+// Copy data from GPU to host memory
+bool HsaRsrcFactory::CopyToHost(const hsa_agent_t& agent, void* dst, const void* src, size_t size) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  if (!cpu_agents_.empty()) {
+    hsa_signal_t s = {};
+    hsa_status_t status = hsa_signal_create(1, 0, NULL, &s);
+    if (status == HSA_STATUS_SUCCESS) {
+      status = hsa_amd_memory_async_copy(dst, cpu_agents_[0], src, agent, size, 0, NULL, s);
+      if (status == HSA_STATUS_SUCCESS) {
+        if (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED) != 0) {
+          status = HSA_STATUS_ERROR;
+        }
+      }
+      status = hsa_signal_destroy(s);
+    }
+  }
+  return (status == HSA_STATUS_SUCCESS);
+}
+bool HsaRsrcFactory::CopyToHost(const AgentInfo* agent_info, void* dst, const void* src, size_t size) {
+  return CopyToHost(agent_info->dev_id, dst, src, size);
+}
+
+// Loads an Assembled Brig file and Finalizes it into Device Isa
+// @param agent_info Gpu device for which to finalize
+// @param brig_path File path of the Assembled Brig file
+// @param kernel_name Name of the kernel to finalize
+// @param code_desc Handle of finalized Code Descriptor that could
+// be used to submit for execution
+// @return bool true if successful, false otherwise
+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
+  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 false;
+  }
+
+  // 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 create code object reader '" << filename << "'" << std::endl;
+    return false;
+  }
+
+  // Create executable.
+  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_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(*executable, "");
+  CHECK_STATUS("Error in freezing executable object", status);
+
+  // Get symbol handle.
+  hsa_executable_symbol_t kernelSymbol;
+  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 true;
+}
+
+// Print the various fields of Hsa Gpu Agents
+bool HsaRsrcFactory::PrintGpuAgents(const std::string& header) {
+  std::clog << header << " :" << std::endl;
+
+  const AgentInfo* agent_info;
+  int size = uint32_t(gpu_list_.size());
+  for (int idx = 0; idx < size; idx++) {
+    agent_info = gpu_list_[idx];
+
+    std::clog << "> agent[" << idx << "] :" << std::endl;
+    std::clog << ">> Name : " << agent_info->name << std::endl;
+    std::clog << ">> APU : " << agent_info->is_apu << std::endl;
+    std::clog << ">> HSAIL profile : " << agent_info->profile << std::endl;
+    std::clog << ">> Max Wave Size : " << agent_info->max_wave_size << std::endl;
+    std::clog << ">> Max Queue Size : " << agent_info->max_queue_size << std::endl;
+//    std::clog << ">> Kernarg Region Id : " << agent_info->coarse_region.handle << std::endl;
+    std::clog << ">> CU number : " << agent_info->cu_num << std::endl;
+    std::clog << ">> Waves per CU : " << agent_info->waves_per_cu << std::endl;
+    std::clog << ">> SIMDs per CU : " << agent_info->simds_per_cu << std::endl;
+    std::clog << ">> SE number : " << agent_info->se_num << std::endl;
+    std::clog << ">> Shader Arrays per SE : " << agent_info->shader_arrays_per_se << std::endl;
+  }
+  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_;
+
+}  // namespace util
+}  // namespace roctracer

src/util/hsa_rsrc_factory.h

@@ -0,0 +1,283 @@
+/**********************************************************************
+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:
+
+<95>    Redistributions of source code must retain the above copyright notice, this list of
+conditions and the following disclaimer.
+<95>    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 _HSA_RSRC_FACTORY_H_
+#define _HSA_RSRC_FACTORY_H_
+
+#include <hsa.h>
+#include <hsa_ext_amd.h>
+#include <hsa_ext_finalize.h>
+#include <hsa_ven_amd_aqlprofile.h>
+#include <hsa_ven_amd_loader.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#include <iostream>
+#include <mutex>
+#include <map>
+#include <string>
+#include <vector>
+
+#define HSA_ARGUMENT_ALIGN_BYTES 16
+#define HSA_QUEUE_ALIGN_BYTES 64
+#define HSA_PACKET_ALIGN_BYTES 64
+
+#define CHECK_STATUS(msg, status)                                                                  \
+  if (status != HSA_STATUS_SUCCESS) {                                                              \
+    const char* emsg = 0;                                                                          \
+    hsa_status_string(status, &emsg);                                                              \
+    printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>");                                      \
+    exit(1);                                                                                       \
+  }
+
+#define CHECK_ITER_STATUS(msg, status)                                                             \
+  if (status != HSA_STATUS_INFO_BREAK) {                                                           \
+    const char* emsg = 0;                                                                          \
+    hsa_status_string(status, &emsg);                                                              \
+    printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>");                                      \
+    exit(1);                                                                                       \
+  }
+
+namespace roctracer {
+namespace util {
+static const unsigned MEM_PAGE_BYTES = 0x1000;
+static const unsigned MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
+typedef decltype(hsa_agent_t::handle) hsa_agent_handle_t;
+
+// Encapsulates information about a Hsa Agent such as its
+// handle, name, max queue size, max wavefront size, etc.
+struct AgentInfo {
+  // Handle of Agent
+  hsa_agent_t dev_id;
+
+  // Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
+  uint32_t dev_type;
+
+  // APU flag
+  bool is_apu;
+
+  // Agent system index
+  uint32_t dev_index;
+
+  // GFXIP name
+  char gfxip[64];
+
+  // Name of Agent whose length is less than 64
+  char name[64];
+
+  // Max size of Wavefront size
+  uint32_t max_wave_size;
+
+  // Max size of Queue buffer
+  uint32_t max_queue_size;
+
+  // Hsail profile supported by agent
+  hsa_profile_t profile;
+#if 0
+  // Memory region supporting kernel parameters
+  hsa_region_t coarse_region;
+
+  // Memory region supporting kernel arguments
+  hsa_region_t kernarg_region;
+#endif
+  // CPU/GPU/kern-arg memory pools
+  hsa_amd_memory_pool_t cpu_pool;
+  hsa_amd_memory_pool_t gpu_pool;
+  hsa_amd_memory_pool_t kern_arg_pool;
+
+  // The number of compute unit available in the agent.
+  uint32_t cu_num;
+
+  // Maximum number of waves possible in a Compute Unit.
+  uint32_t waves_per_cu;
+
+  // Number of SIMD's per compute unit CU
+  uint32_t simds_per_cu;
+
+  // Number of Shader Engines (SE) in Gpu
+  uint32_t se_num;
+
+  // Number of Shader Arrays Per Shader Engines in Gpu
+  uint32_t shader_arrays_per_se;
+};
+
+class HsaRsrcFactory {
+ public:
+  typedef std::recursive_mutex mutex_t;
+
+  static HsaRsrcFactory* Create(bool initialize_hsa = true) {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (instance_ == NULL) {
+      instance_ = new HsaRsrcFactory(initialize_hsa);
+    }
+    return instance_;
+  }
+
+  static HsaRsrcFactory& Instance() {
+    if (instance_ == NULL) instance_ = Create(false);
+    hsa_status_t status = (instance_ != NULL) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
+    CHECK_STATUS("HsaRsrcFactory::Instance() failed", status);
+    return *instance_;
+  }
+
+  static void Destroy() {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (instance_) delete instance_;
+    instance_ = NULL;
+  }
+
+  // Return system agent info
+  const AgentInfo* GetAgentInfo(const hsa_agent_t agent);
+
+  // Get the count of Hsa Gpu Agents available on the platform
+  // @return uint32_t Number of Gpu agents on platform
+  uint32_t GetCountOfGpuAgents();
+
+  // Get the count of Hsa Cpu Agents available on the platform
+  // @return uint32_t Number of Cpu agents on platform
+  uint32_t GetCountOfCpuAgents();
+
+  // Get the AgentInfo handle of a Gpu device
+  // @param idx Gpu Agent at specified index
+  // @param agent_info Output parameter updated with AgentInfo
+  // @return bool true if successful, false otherwise
+  bool GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
+
+  // Get the AgentInfo handle of a Cpu device
+  // @param idx Cpu Agent at specified index
+  // @param agent_info Output parameter updated with AgentInfo
+  // @return bool true if successful, false otherwise
+  bool GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
+
+  // Create a Queue object and return its handle. The queue object is expected
+  // to support user requested number of Aql dispatch packets.
+  // @param agent_info Gpu Agent on which to create a queue object
+  // @param num_Pkts Number of packets to be held by queue
+  // @param queue Output parameter updated with handle of queue object
+  // @return bool true if successful, false otherwise
+  bool CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
+
+  // Create a Signal object and return its handle.
+  // @param value Initial value of signal object
+  // @param signal Output parameter updated with handle of signal object
+  // @return bool true if successful, false otherwise
+  bool CreateSignal(uint32_t value, hsa_signal_t* signal);
+
+  // Allocate local GPU memory
+  // @param agent_info Agent from whose memory region to allocate
+  // @param size Size of memory in terms of bytes
+  // @return uint8_t* Pointer to buffer, null if allocation fails.
+  uint8_t* AllocateLocalMemory(const AgentInfo* agent_info, size_t size);
+
+  // Allocate memory tp pass kernel parameters
+  // Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo parameter.
+  // @param agent_info Agent from whose memory region to allocate
+  // @param size Size of memory in terms of bytes
+  // @return uint8_t* Pointer to buffer, null if allocation fails.
+  uint8_t* AllocateKernArgMemory(const AgentInfo* agent_info, size_t size);
+
+  // Allocate system memory accessible from both CPU and GPU
+  // Memory is alocated accessible to all CPU agents and AgentInfo parameter is ignored.
+  // @param agent_info Agent from whose memory region to allocate
+  // @param size Size of memory in terms of bytes
+  // @return uint8_t* Pointer to buffer, null if allocation fails.
+  uint8_t* AllocateSysMemory(const AgentInfo* agent_info, size_t size);
+
+  // Copy data from GPU to host memory
+  bool CopyToHost(const hsa_agent_t& agent, void* dst, const void* src, size_t size);
+  bool CopyToHost(const AgentInfo* agent_info, void* dst, const void* src, size_t size);
+
+  // Loads an Assembled Brig file and Finalizes it into Device Isa
+  // @param agent_info Gpu device for which to finalize
+  // @param brig_path File path of the Assembled Brig file
+  // @param kernel_name Name of the kernel to finalize
+  // @param code_desc Handle of finalized Code Descriptor that could
+  // be used to submit for execution
+  // @return true if successful, false otherwise
+  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
+  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_; }
+
+  // Return Loader API table
+  const hsa_ven_amd_loader_1_00_pfn_t* LoaderApi() const { return &loader_api_; }
+
+ private:
+  // System agents iterating callback
+  static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data);
+
+  // Callback function to find and bind kernarg region of an agent
+  static hsa_status_t FindMemRegionsCallback(hsa_region_t region, void* data);
+
+  // 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(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);
+
+  static HsaRsrcFactory* instance_;
+  static mutex_t mutex_;
+
+  // Used to maintain a list of Hsa Gpu Agent Info
+  std::vector<const AgentInfo*> gpu_list_;
+  std::vector<hsa_agent_t> gpu_agents_;
+
+  // Used to maintain a list of Hsa Cpu Agent Info
+  std::vector<const AgentInfo*> cpu_list_;
+  std::vector<hsa_agent_t> cpu_agents_;
+
+  // System agents map
+  std::map<hsa_agent_handle_t, const AgentInfo*> agent_map_;
+
+  // AqlProfile API table
+  aqlprofile_pfn_t aqlprofile_api_;
+
+  // Loader API table
+  hsa_ven_amd_loader_1_00_pfn_t loader_api_;
+};
+
+}  // namespace util
+}  // namespace roctracer
+
+#endif  // _HSA_RSRC_FACTORY_H_

src/util/logger.h

@@ -0,0 +1,169 @@
+#ifndef SRC_UTIL_LOGGER_H_
+#define SRC_UTIL_LOGGER_H_
+
+#include <time.h>
+#include <stdio.h>
+#include <unistd.h>
+#include <sys/types.h>
+#include <sys/syscall.h>
+#include <sys/file.h>
+#include <stdarg.h>
+#include <stdlib.h>
+
+#include <string>
+#include <iostream>
+#include <sstream>
+#include <fstream>
+#include <exception>
+#include <mutex>
+#include <map>
+
+namespace roctracer {
+namespace util {
+
+class Logger {
+ public:
+  typedef std::recursive_mutex mutex_t;
+
+  template <typename T> Logger& operator<<(const T& m) {
+    std::ostringstream oss;
+    oss << m;
+    if (!streaming_)
+      Log(oss.str());
+    else
+      Put(oss.str());
+    streaming_ = true;
+    return *this;
+  }
+
+  typedef void (*manip_t)();
+  Logger& operator<<(manip_t f) {
+    f();
+    return *this;
+  }
+
+  static void begm() { Instance().ResetStreaming(true); }
+  static void endl() { Instance().ResetStreaming(false); }
+
+  static const std::string& LastMessage() {
+    Logger& logger = Instance();
+    std::lock_guard<mutex_t> lck(mutex_);
+    return logger.message_[GetTid()];
+  }
+
+  static Logger* Create() {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (instance_ == NULL) instance_ = new Logger();
+    return instance_;
+  }
+
+  static void Destroy() {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (instance_ != NULL) delete instance_;
+    instance_ = NULL;
+  }
+
+  static Logger& Instance() {
+    Create();
+    return *instance_;
+  }
+
+ private:
+  static uint32_t GetPid() { return syscall(__NR_getpid); }
+  static uint32_t GetTid() { return syscall(__NR_gettid); }
+
+  Logger() : file_(NULL), dirty_(false), streaming_(false), messaging_(false) {
+    const char* path = getenv("ROCTRACER_LOG");
+    if (path != NULL) {
+      file_ = fopen("/tmp/roctracer_log.txt", "a");
+    }
+    ResetStreaming(false);
+  }
+
+  ~Logger() {
+    if (file_ != NULL) {
+      if (dirty_) Put("\n");
+      fclose(file_);
+    }
+  }
+
+  void ResetStreaming(const bool messaging) {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (messaging) {
+      message_[GetTid()] = "";
+    } else if (streaming_) {
+      Put("\n");
+      dirty_ = false;
+    }
+    messaging_ = messaging;
+    streaming_ = messaging;
+  }
+
+  void Put(const std::string& m) {
+    std::lock_guard<mutex_t> lck(mutex_);
+    if (messaging_) {
+      message_[GetTid()] += m;
+    }
+    if (file_ != NULL) {
+      dirty_ = true;
+      flock(fileno(file_), LOCK_EX);
+      fprintf(file_, "%s", m.c_str());
+      fflush(file_);
+      flock(fileno(file_), LOCK_UN);
+    }
+  }
+
+  void Log(const std::string& m) {
+    const time_t rawtime = time(NULL);
+    tm tm_info;
+    localtime_r(&rawtime, &tm_info);
+    char tm_str[26];
+    strftime(tm_str, 26, "%Y-%m-%d %H:%M:%S", &tm_info);
+    std::ostringstream oss;
+    oss << "<" << tm_str << std::dec << " pid" << GetPid() << " tid" << GetTid() << "> " << m;
+    Put(oss.str());
+  }
+
+  FILE* file_;
+  bool dirty_;
+  bool streaming_;
+  bool messaging_;
+
+  static mutex_t mutex_;
+  static Logger* instance_;
+  std::map<uint32_t, std::string> message_;
+};
+
+}  // namespace util
+}  // namespace roctracer
+
+#define ERR_LOGGING(stream)                                                                        \
+  {                                                                                                \
+    roctracer::util::Logger::Instance() << "error: " << roctracer::util::Logger::begm          \
+                                          << stream << roctracer::util::Logger::endl;            \
+  }
+
+#define INFO_LOGGING(stream)                                                                       \
+  {                                                                                                \
+    roctracer::util::Logger::Instance() << "info: " << roctracer::util::Logger::begm << stream \
+                                          << roctracer::util::Logger::endl;                      \
+  }
+
+#define WARN_LOGGING(stream)                                                                       \
+  {                                                                                                \
+    std::cerr << "ROCProfiler: " << stream << std::endl;                                                              \
+    roctracer::util::Logger::Instance() << "warning: " << roctracer::util::Logger::begm << stream \
+                                          << roctracer::util::Logger::endl;                      \
+  }
+
+#ifdef DEBUG
+#define DBG_LOGGING(stream)                                                                        \
+  {                                                                                                \
+    roctracer::util::Logger::Instance() << roctracer::util::Logger::begm << "debug: \""        \
+                                          << stream << "\"" < < < <                                \
+        " in " << __FUNCTION__ << " at " << __FILE__ << " line " << __LINE__                       \
+               << roctracer::util::Logger::endl;                                                 \
+  }
+#endif
+
+#endif  // SRC_UTIL_LOGGER_H_