SWDEV-213367 : codeobj event implementation

Change-Id: Ibcaca6869ce96d8802c5fa8ba241f43834d6f2a7 update - codeobj event implementation Change-Id: I4c12f26a19f2b31d9ac2211c3426a0e587a332b3 update2 - codeobj event implementation Change-Id: Ic877549a83542ae00352503471d881e847ebac9c test - codeobj event implementation Change-Id: I0618d3a93de94c3d7467372ba4a3d4ea5520bfc7 URI reference test - codeobj event implementation Change-Id: I6cf7e8a648cf012cb0708058b118a75e58f992b9 adding test/app - codeobj event implementation Change-Id: Idf4c197c7b9116ccde5ec50ff47a26a858bfab32 uri test fix - codeobj event implementation Change-Id: I7c385f82f516d9d8f2cd726366f00be3664006e3 uri test cleanup - codeobj event implementation Change-Id: I542d5baf88c048c8b4717af843b803cd93e8f3bc URI buffer fix - codeobj event implementation Change-Id: Iac65e04c03a0939935c10f53c6b580a2e33878f5 HSA events tests trace-check disabled Change-Id: I0f4d13aeeceb1d1a6e2191673eacbf9c7ae2ae52 [ROCm/roctracer commit: 29c63c5281]
2020-09-03 04:01:28 -05:00
@@ -151,10 +151,10 @@ get_filename_component ( HSA_RUNTIME_LIB_PATH "${HSA_RUNTIME_LIB}" DIRECTORY )

 find_library ( HSA_KMT_LIB "libhsakmt.so" )
 get_filename_component ( HSA_KMT_LIB_PATH "${HSA_KMT_LIB}" DIRECTORY )
-get_filename_component ( ROCM_ROOT_DIR "${HSA_KMT_LIB_PATH}" DIRECTORY )
-
 set ( HSA_KMT_INC_PATH "${HSA_KMT_LIB_PATH}/../include" )
-set ( ROCM_INC_PATH "${HSA_KMT_INC_PATH}" )
+
+get_filename_component ( ROCM_ROOT_DIR "${HSA_KMT_LIB_PATH}" DIRECTORY )
+set ( ROCM_INC_PATH "${ROCM_ROOT_DIR}/include" )

 ## Basic Tool Chain Information
 message ( "----------------NBit: ${NBIT}" )
@@ -167,6 +167,7 @@ message ( "-----HSA-Runtime-Inc: ${HSA_RUNTIME_INC_PATH}" )
 message ( "-----HSA-Runtime-Lib: ${HSA_RUNTIME_LIB_PATH}" )
 message ( "----HSA_KMT_LIB_PATH: ${HSA_KMT_LIB_PATH}" )
 message ( "-------ROCM_ROOT_DIR: ${ROCM_ROOT_DIR}" )
+message ( "-------ROCM_INC_PATH: ${ROCM_INC_PATH}" )
 message ( "-------------KFD-Inc: ${HSA_KMT_INC_PATH}" )
 message ( "-------------HIP-Inc: ${HIP_INC_DIR}" )
 message ( "-------------HIP-VDI: ${HIP_VDI}" )
@@ -27,6 +27,7 @@ THE SOFTWARE.
 #include <hsa_ext_amd.h>

 #include <roctracer.h>
+#include <rocprofiler/activity.h>

 // HSA OP ID enumeration
 enum hsa_op_id_t {
@@ -34,7 +35,7 @@ enum hsa_op_id_t {
  HSA_OP_ID_COPY = 1,
  HSA_OP_ID_BARRIER = 2,
  HSA_OP_ID_RESERVED1 = 3,
-  HSA_OP_ID_NUMBER = 4
+  HSA_OP_ID_NUMBER
 };

 #ifdef __cplusplus
@@ -11,6 +11,9 @@ execute_process ( COMMAND sh -xc "${CMAKE_C_COMPILER} -E ${HIP_PATH}/include/hip
 execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/gen_ostream_ops.py -in ${GEN_INC_DIR}/hip_runtime_api_pp.h -out ${GEN_INC_DIR}/hip_ostream_ops.h" )
 execute_process ( COMMAND sh -xc "${CMAKE_C_COMPILER} -E ${HSA_RUNTIME_INC_PATH}/hsa.h > ${GEN_INC_DIR}/hsa_pp.h" )
 execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/gen_ostream_ops.py -in ${GEN_INC_DIR}/hsa_pp.h -out ${GEN_INC_DIR}/hsa_ostream_ops.h" )
+execute_process ( COMMAND sh -xc "mkdir ${GEN_INC_DIR}/rocprofiler" )
+execute_process ( COMMAND sh -xc "ln -s ${ROOT_DIR}/../rocprofiler/inc/rocprofiler.h ${GEN_INC_DIR}/rocprofiler/rocprofiler.h" )
+execute_process ( COMMAND sh -xc "ln -s ${ROOT_DIR}/../rocprofiler/src/core/activity.h ${GEN_INC_DIR}/rocprofiler/activity.h" )

 # Build dynamic Library object
 set ( TARGET_LIB ${TARGET_NAME} )
@@ -22,14 +25,14 @@ set ( LIB_SRC
  ${LIB_DIR}/util/hsa_rsrc_factory.cpp
 )
 add_library ( ${TARGET_LIB} ${LIBRARY_TYPE} ${LIB_SRC} )
-target_include_directories ( ${TARGET_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HIP_INC_DIR} ${HCC_INC_DIR} ${HSA_KMT_INC_PATH} ${GEN_INC_DIR} )
+target_include_directories ( ${TARGET_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HIP_INC_DIR} ${HSA_KMT_INC_PATH} ${ROCM_INC_PATH} ${GEN_INC_DIR} )
 target_link_libraries( ${TARGET_LIB} PRIVATE ${HSA_RUNTIME_LIB} c stdc++ )

 # Build KFD/Thunk tracing library
 set ( KFD_LIB "kfdwrapper64" )
 set ( KFD_LIB_SRC ${GEN_SRC_DIR}/kfd_wrapper.cpp)
 add_library ( ${KFD_LIB} SHARED ${KFD_LIB_SRC} )
-target_include_directories ( ${KFD_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HSA_KMT_INC_PATH} ${GEN_INC_DIR} )
+target_include_directories ( ${KFD_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_KMT_INC_PATH} ${GEN_INC_DIR} )
 target_link_libraries( ${KFD_LIB} PRIVATE c stdc++ )

 # Build ROCTX tracing library
@@ -39,5 +42,5 @@ set ( ROCTX_LIB_SRC
  ${LIB_DIR}/roctx/roctx_intercept.cpp
 )
 add_library ( ${ROCTX_LIB} SHARED ${ROCTX_LIB_SRC} )
-target_include_directories ( ${ROCTX_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${GEN_INC_DIR} )
+target_include_directories ( ${ROCTX_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${GEN_INC_DIR} )
 target_link_libraries( ${ROCTX_LIB} PRIVATE c stdc++ )
@@ -94,6 +94,10 @@ class RocpApi {
  EnableCallback_t* EnableActivityCallback;
  NameCallback_t* GetOpName;

+  RegisterCallback_t* RegisterEvtCallback;
+  OperateCallback_t* RemoveEvtCallback;
+  NameCallback_t* GetEvtName;
+
  protected:
  void init(Loader* loader) {
    RegisterApiCallback = loader->GetFun<RegisterCallback_t>("RegisterApiCallback");
@@ -101,6 +105,10 @@ class RocpApi {
    InitActivityCallback = loader->GetFun<InitCallback_t>("InitActivityCallback");
    EnableActivityCallback = loader->GetFun<EnableCallback_t>("EnableActivityCallback");
    GetOpName = loader->GetFun<NameCallback_t>("GetOpName");
+
+    RegisterEvtCallback = loader->GetFun<RegisterCallback_t>("RegisterEvtCallback");
+    RemoveEvtCallback = loader->GetFun<OperateCallback_t>("RemoveEvtCallback");
+    GetEvtName = loader->GetFun<NameCallback_t>("GetEvtName");
  }
 };

@@ -290,6 +298,7 @@ typedef HipLoaderShared HipLoader;
  template<class T> bool roctracer::BaseLoader<T>::to_check_open_ = true; \
  template<class T> bool roctracer::BaseLoader<T>::to_check_symb_ = true; \
  template<> const char* roctracer::RocpLoader::lib_name_ = "librocprofiler64.so"; \
+  template<> bool roctracer::RocpLoader::to_load_ = true; \
  template<> const char* roctracer::HccLoader::lib_name_ = "libamdhip64.so"; \
  template<> const char* roctracer::KfdLoader::lib_name_ = "libkfdwrapper64.so"; \
  template<> const char* roctracer::RocTxLoader::lib_name_ = "libroctx64.so"; \
@@ -686,6 +686,8 @@ PUBLIC_API const char* roctracer_op_string(
  switch (domain) {
    case ACTIVITY_DOMAIN_HSA_API:
      return roctracer::hsa_support::GetApiName(op);
+    case ACTIVITY_DOMAIN_HSA_EVT:
+      return roctracer::RocpLoader::Instance().GetEvtName(op);
    case ACTIVITY_DOMAIN_HSA_OPS:
      return roctracer::RocpLoader::Instance().GetOpName(op);
    case ACTIVITY_DOMAIN_HCC_OPS:
@@ -731,6 +733,7 @@ static inline uint32_t get_op_num(const uint32_t& domain) {
  switch (domain) {
    case ACTIVITY_DOMAIN_HSA_OPS: return HSA_OP_ID_NUMBER;
    case ACTIVITY_DOMAIN_HSA_API: return HSA_API_ID_NUMBER;
+    case ACTIVITY_DOMAIN_HSA_EVT: return HSA_EVT_ID_NUMBER;
    case ACTIVITY_DOMAIN_HCC_OPS: return HIP_OP_ID_NUMBER;
    case ACTIVITY_DOMAIN_HIP_API: return HIP_API_ID_NUMBER;
    case ACTIVITY_DOMAIN_KFD_API: return KFD_API_ID_NUMBER;
@@ -760,13 +763,18 @@ static roctracer_status_t roctracer_enable_callback_fun(
 #if 0
      if (op == HSA_API_ID_DISPATCH) {
        const bool succ = roctracer::RocpLoader::Instance().RegisterApiCallback(op, (void*)callback, user_data);
-        if (succ == false) HCC_EXC_RAISING(ROCTRACER_STATUS_HSA_ERR, "HSA::EnableActivityCallback error(" << op << ") failed");
+        if (succ == false) HCC_EXC_RAISING(ROCTRACER_STATUS_HSA_ERR, "HSA::RegisterApiCallback error(" << op << ") failed");
        break;
      }
 #endif
      roctracer::hsa_support::cb_table.set(op, callback, user_data);
      break;
    }
+    case ACTIVITY_DOMAIN_HSA_EVT: {
+      const bool succ = roctracer::RocpLoader::Instance().RegisterEvtCallback(op, (void*)callback, user_data);
+      if (succ == false) HCC_EXC_RAISING(ROCTRACER_STATUS_HSA_ERR, "HSA::RegisterEvtCallback error(" << op << ") failed");
+      break;
+    }
    case ACTIVITY_DOMAIN_HCC_OPS: break;
    case ACTIVITY_DOMAIN_HIP_API: {
      if (roctracer::HipLoader::Instance().Enabled() == false) break;
@@ -875,6 +883,11 @@ static roctracer_status_t roctracer_disable_callback_fun(
      }
      break;
    }
+    case ACTIVITY_DOMAIN_HSA_EVT: {
+      const bool succ = roctracer::RocpLoader::Instance().RemoveEvtCallback(op);
+      if (succ == false) HCC_EXC_RAISING(ROCTRACER_STATUS_HSA_ERR, "HSA::RemoveEvtCallback error(" << op << ") failed");
+      break;
+    }
    case ACTIVITY_DOMAIN_ROCTX: {
      if (roctracer::RocTxLoader::Instance().Enabled()) {
        const bool suc = roctracer::RocTxLoader::Instance().RemoveApiCallback(op);
@@ -984,6 +997,7 @@ static roctracer_status_t roctracer_enable_activity_fun(
      break;
    }
    case ACTIVITY_DOMAIN_HSA_API: break;
+    case ACTIVITY_DOMAIN_HSA_EVT: break;
    case ACTIVITY_DOMAIN_KFD_API: break;
    case ACTIVITY_DOMAIN_HCC_OPS: {
      const bool init_phase = (roctracer::HccLoader::GetRef() == NULL);
@@ -1080,6 +1094,7 @@ static roctracer_status_t roctracer_disable_activity_fun(
      break;
    }
    case ACTIVITY_DOMAIN_HSA_API: break;
+    case ACTIVITY_DOMAIN_HSA_EVT: break;
    case ACTIVITY_DOMAIN_KFD_API: break;
    case ACTIVITY_DOMAIN_HCC_OPS: {
      if (roctracer::HccLoader::Instance().Enabled() == false) break;
@@ -1250,6 +1265,9 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
      roctracer::kfd_support::intercept_KFDApiTable();
      break;
    }
+    case ACTIVITY_DOMAIN_HSA_EVT: {
+      break;
+    }
    case ACTIVITY_DOMAIN_HSA_API: {
      // HSA API properties
      HsaApiTable* table = reinterpret_cast<HsaApiTable*>(properties);
@@ -50,7 +50,8 @@ set ( HSA_REV "19b1191" )
 set ( RUN_SCRIPT "${TEST_DIR}/run.sh" )

 ## build HIP tests
-set ( INC_PATH "${INC_PATH} ${PROJECT_BINARY_DIR}/inc" )
+set ( GEN_INC_DIR ${PROJECT_BINARY_DIR}/inc )
+set ( INC_PATH "${INC_PATH} ${GEN_INC_DIR}" )
 set ( TEST_ENV HIP_VDI=${HIP_VDI} ROCM_PATH=${ROCM_ROOT_DIR} HSA_PATH=${ROCM_ROOT_DIR}/hsa INC_PATH=${INC_PATH} LIB_PATH=${LIB_PATH} HIPCC_VERBOSE=3 )
 add_custom_target( mytest
  COMMAND ${TEST_ENV} make -C "${TEST_DIR}/MatrixTranspose"
@@ -76,10 +77,24 @@ if ( DEFINED ROCTRACER_TARGET )
  set ( TEST_LIB "tracer_tool" )
  set ( TEST_LIB_SRC ${TEST_DIR}/tool/tracer_tool.cpp ${UTIL_SRC} )
  add_library ( ${TEST_LIB} SHARED ${TEST_LIB_SRC} )
-  target_include_directories ( ${TEST_LIB} PRIVATE ${HSA_TEST_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HIP_INC_DIR} ${HCC_INC_DIR} ${HSA_KMT_INC_PATH} ${PROJECT_BINARY_DIR}/inc )
+  target_include_directories ( ${TEST_LIB} PRIVATE ${HSA_TEST_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HIP_INC_DIR} ${HSA_KMT_INC_PATH} ${GEN_INC_DIR} )
  target_link_libraries ( ${TEST_LIB} ${ROCTRACER_TARGET} ${HSA_RUNTIME_LIB} c stdc++ dl pthread rt )
 endif ()

+## Build hsaco_test.cpp referenc test
+set ( CO_LIB_NAME "hsaco_test" )
+set ( CO_LIB_SRC ${TEST_DIR}/app/hsaco_test.cpp )
+add_library ( ${CO_LIB_NAME} SHARED ${CO_LIB_SRC} )
+target_include_directories ( ${CO_LIB_NAME} PRIVATE ${HSA_RUNTIME_INC_PATH} )
+target_link_libraries ( ${CO_LIB_NAME} ${HSA_RUNTIME_LIB} c stdc++ )
+
+## Build codeobj event test
+set ( CO_LIB_NAME "codeobj_test" )
+set ( CO_LIB_SRC ${TEST_DIR}/app/codeobj_test.cpp )
+add_library ( ${CO_LIB_NAME} SHARED ${CO_LIB_SRC} )
+target_include_directories ( ${CO_LIB_NAME} PRIVATE ${TEST_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${GEN_INC_DIR} ${HSA_RUNTIME_INC_PATH} ${ROCM_INC_PATH} )
+target_link_libraries ( ${CO_LIB_NAME} ${ROCTRACER_TARGET} c stdc++ )
+
 ## Build HSA test
 execute_process ( COMMAND sh -xc "if [ ! -e ${TEST_DIR}/hsa ] ; then git clone https://github.com/ROCmSoftwarePlatform/hsa-class.git ${TEST_DIR}/hsa; fi" )
 execute_process ( COMMAND sh -xc "if [ -e ${TEST_DIR}/hsa ] ; then cd ${TEST_DIR}/hsa && git fetch origin && git checkout ${HSA_REV}; fi" )
@@ -0,0 +1,89 @@
+/******************************************************************************
+Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*******************************************************************************/
+
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include "inc/roctracer.h"
+#include "inc/roctracer_hsa.h"
+#include <rocprofiler/rocprofiler.h>
+
+#define PUBLIC_API __attribute__((visibility("default")))
+#define CONSTRUCTOR_API __attribute__((constructor))
+#define DESTRUCTOR_API __attribute__((destructor))
+
+// Check returned HSA API status
+void check_status(roctracer_status_t status) {
+  if (status != ROCTRACER_STATUS_SUCCESS) {
+    const char* error_string = roctracer_error_string();
+    fprintf(stderr, "ERROR: %s\n", error_string);
+    abort();
+  }
+}
+
+// codeobj callback
+void codeobj_callback(uint32_t domain, uint32_t cid, const void* data, void* arg) {
+  const hsa_evt_data_t* evt_data = reinterpret_cast<const hsa_evt_data_t*>(data);
+  const uint32_t uri_length = evt_data->codeobj.uri_length;
+  const char* uri = evt_data->codeobj.uri;
+  printf("codeobj_callback domain(%u) cid(%u): load_delta(0x%lx) load_size(0x%lx) uri_length(%u) uri(\"%s\")\n",
+    domain,
+    cid,
+    evt_data->codeobj.load_delta,
+    evt_data->codeobj.load_size,
+    uri_length,
+    uri);
+  fflush(stdout);
+}
+
+void initialize() {
+  roctracer_status_t status = roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, codeobj_callback, NULL);
+  check_status(status);
+}
+
+void cleanup() {
+  roctracer_status_t status = roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_EVT);
+  check_status(status);
+}
+
+// Tool constructor
+extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) {
+  // Enable HSA events intercepting
+  settings->hsa_intercepting = 1;
+  // Initialize profiling
+  initialize();
+}
+
+// Tool destructor
+extern "C" PUBLIC_API void OnUnloadTool() {
+  // Final resources cleanup
+  cleanup();
+}
+
+extern "C" CONSTRUCTOR_API void constructor() {
+  printf("constructor\n"); fflush(stdout);
+}
+
+extern "C" DESTRUCTOR_API void destructor() {
+  OnUnloadTool();
+}
@@ -0,0 +1,134 @@
+/******************************************************************************
+Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*******************************************************************************/
+
+#include <hsa.h>
+#include <hsa_api_trace.h>
+#include <hsa_ven_amd_loader.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define PUBLIC_API __attribute__((visibility("default")))
+#define CONSTRUCTOR_API __attribute__((constructor))
+#define DESTRUCTOR_API __attribute__((destructor))
+
+#define HSA_RT(call) \
+  do { \
+    const hsa_status_t status = call; \
+    if (status != HSA_STATUS_SUCCESS) { \
+      printf("error \"%s\"\n", #call); fflush(stdout); \
+      abort(); \
+    } \
+  } while(0)
+
+// HSA API intercepting primitives
+decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
+hsa_ven_amd_loader_1_01_pfn_t loader_api_table{};
+
+hsa_status_t code_object_callback(
+  hsa_executable_t executable,
+  hsa_loaded_code_object_t loaded_code_object,
+  void* arg)
+{
+  printf("code_object_callback\n"); fflush(stdout);
+
+  uint64_t load_size = 0;
+  uint64_t load_delta = 0;
+  uint32_t uri_len = 0;
+  char* uri_str = NULL;
+
+  HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
+    loaded_code_object,
+    HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
+    &load_size));
+  HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
+    loaded_code_object,
+    HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
+    &load_delta));
+  HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
+    loaded_code_object,
+    HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
+    &uri_len));
+
+  uri_str = (char*)calloc(uri_len + 1, sizeof(char));
+  if (!uri_str) {
+    perror("calloc");
+    abort();
+  }
+
+  HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
+    loaded_code_object,
+    HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI,
+    uri_str));
+
+  printf("load_size(0x%lx)\n", load_size); fflush(stdout);
+  printf("load_delta(0x%lx)\n", load_delta); fflush(stdout);
+  printf("uri_len(%u)\n", uri_len); fflush(stdout);
+  printf("uri_str(\"%s\")\n", uri_str); fflush(stdout);
+
+  return HSA_STATUS_SUCCESS;
+}
+
+hsa_status_t hsa_executable_freeze_interceptor(
+  hsa_executable_t executable,
+  const char *options)
+{
+  HSA_RT(loader_api_table.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
+    executable,
+    code_object_callback,
+    NULL));
+  HSA_RT(hsa_executable_freeze_fn(
+    executable,
+    options));
+  return HSA_STATUS_SUCCESS;
+}
+
+// HSA-runtime tool on-load method
+extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table,
+                                  uint64_t runtime_version,
+                                  uint64_t failed_tool_count,
+                                  const char* const* failed_tool_names)
+{
+  printf("OnLoad: begin\n"); fflush(stdout);
+  // intercepting hsa_executable_freeze API
+  hsa_executable_freeze_fn = table->core_->hsa_executable_freeze_fn;
+  table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
+  // Fetching AMD Loader HSA extension API
+  HSA_RT(hsa_system_get_major_extension_table(
+    HSA_EXTENSION_AMD_LOADER,
+    1,
+    sizeof(hsa_ven_amd_loader_1_01_pfn_t),
+    &loader_api_table));
+  printf("OnLoad: end\n"); fflush(stdout);
+  return true;
+}
+
+extern "C" PUBLIC_API void OnUnload() {
+  printf("OnUnload\n"); fflush(stdout);
+}
+
+extern "C" CONSTRUCTOR_API void constructor() {
+  printf("constructor\n"); fflush(stdout);
+}
+
+extern "C" DESTRUCTOR_API void destructor() {
+  printf("destructor\n"); fflush(stdout);
+}
@@ -10,3 +10,5 @@ MatrixTranspose_hip_flush_trace --check-order .*
 MatrixTranspose_kfd_trace --check-events .*
 ctrl_hsa_trace --check-event .*
 ctrl_hsa_input_trace --check-event .*
+hsa_co_trace --check-none
+code_obj_trace --check-none
@@ -151,6 +151,13 @@ echo "<trace name=\"HSA\"><parameters api=\"hsa_agent_get_info, hsa_amd_memory_p
 export ROCP_INPUT=input.xml
 eval_test "tool HSA test input" ./test/hsa/ctrl ctrl_hsa_input_trace

+export HSA_TOOLS_LIB=./test/libhsaco_test.so
+eval_test "tool HSA codeobj" ./test/MatrixTranspose hsa_co_trace
+
+export ROCP_TOOL_LIB=./test/libcodeobj_test.so
+export HSA_TOOLS_LIB=librocprofiler64.so
+eval_test "tool tracer codeobj" ./test/MatrixTranspose code_obj_trace
+
 #valgrind --leak-check=full $tbin
 #valgrind --tool=massif $tbin
 #ms_print massif.out.<N>