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
2020-09-03 04:01:28 -05:00
parent 02fcb90a4b
commit 29c63c5281
11 changed files with 1297 additions and 317 deletions
@@ -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);
+}