Initial external logging API

New API to accept a file stream for logging

Co-authored-by: David Yat Sin <David.YatSin@amd.com>

Change-Id: Ie09c35ae14ca86a97eb25f61251be287c55d7169
Signed-off-by: Chris Freehill <cfreehil@amd.com>

runtime/hsa-runtime/core/common/hsa_table_interface.cpp

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -1308,6 +1308,10 @@ hsa_status_t HSA_API hsa_amd_queue_get_info(hsa_queue_t* queue,
   return amdExtTable->hsa_amd_queue_get_info_fn(queue, attribute, value);
 }
 
+hsa_status_t HSA_API hsa_amd_enable_logging(uint8_t* flags, void* file) {
+  return amdExtTable->hsa_amd_enable_logging_fn(flags, file);
+}
+
 // Tools only table interfaces.
 namespace rocr {
 

runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -357,6 +357,9 @@ hsa_status_t HSA_API hsa_amd_agent_set_async_scratch_limit(hsa_agent_t agent, si
 hsa_status_t hsa_amd_queue_get_info(hsa_queue_t* queue, hsa_queue_info_attribute_t attribute,
                                     void* value);
 
+// Mirrors Amd Extension Apis
+hsa_status_t HSA_API hsa_amd_enable_logging(uint8_t* flags, void* file);
+
 }  // namespace amd
 }  // namespace rocr
 

runtime/hsa-runtime/core/inc/runtime.h

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -393,6 +393,8 @@ class Runtime {
                                                    const core::MemoryRegion** mem_region,
                                                    hsa_amd_memory_type_t* type);
 
+  hsa_status_t EnableLogging(uint8_t* flags, void* file);
+
   const std::vector<Agent*>& cpu_agents() { return cpu_agents_; }
 
   const std::vector<Agent*>& gpu_agents() { return gpu_agents_; }

runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp

@@ -2,24 +2,24 @@
 //
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
-// 
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
-// 
+//
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
+//
 // Developed by:
-// 
+//
 //                 AMD Research and AMD HSA Software Development
-// 
+//
 //                 Advanced Micro Devices, Inc.
-// 
+//
 //                 www.amd.com
-// 
+//
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to
 // deal with 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:
-// 
+//
 //  - Redistributions of source code must retain the above copyright notice,
 //    this list of conditions and the following disclaimers.
 //  - Redistributions in binary form must reproduce the above copyright
@@ -29,7 +29,7 @@
 //    nor the names of its contributors may be used to endorse or promote
 //    products derived from this Software without specific prior written
 //    permission.
-// 
+//
 // 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
@@ -898,6 +898,10 @@ static int kCopyMisalignedUnroll = GetKernelSourceParam("kCopyMisalignedUnroll")
 static int kFillVecWidth = GetKernelSourceParam("kFillVecWidth");
 static int kFillUnroll = GetKernelSourceParam("kFillUnroll");
 
+static unsigned extractAqlBits(unsigned v, unsigned pos, unsigned width) {
+  return (v >> pos) & ((1 << width) - 1);
+};
+
 BlitKernel::BlitKernel(core::Queue* queue)
     : core::Blit(),
       queue_(queue),
@@ -1041,6 +1045,25 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(
       std::atomic_thread_fence(std::memory_order_release);
       queue_buffer[(write_index)&queue_bitmask_].header = kBarrierPacketHeader;
 
+      LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS,
+      "HWq=%p, id=%d, Barrier Header = "
+      "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
+      "dep_signal=[0x%zx 0x%zx 0x%zx 0x%zx 0x%zx], completion_signal=0x%zx "
+      "rptr=%u, wptr=%u",
+      queue_->public_handle()->base_address, queue_->public_handle()->id,
+      kBarrierPacketHeader,
+      extractAqlBits(kBarrierPacketHeader,
+                    HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
+      extractAqlBits(kBarrierPacketHeader,
+                    HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
+      extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
+                    HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
+      extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
+                    HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
+      barrier_packet.dep_signal[0], barrier_packet.dep_signal[1], barrier_packet.dep_signal[2],
+      barrier_packet.dep_signal[3], barrier_packet.dep_signal[4],
+      barrier_packet.completion_signal, queue_->LoadReadIndexRelaxed(), write_index);
+
       ++write_index;
 
       memset(&barrier_packet, 0, sizeof(hsa_barrier_and_packet_t));
@@ -1234,6 +1257,28 @@ void BlitKernel::PopulateQueue(uint64_t index, uint64_t code_handle, void* args,
   queue_buffer[index & queue_bitmask_] = packet;
   std::atomic_thread_fence(std::memory_order_release);
   queue_buffer[index & queue_bitmask_].header = kDispatchPacketHeader;
+
+  LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS,
+    "HWq=%p, id=%d, Dispatch Header = "
+    "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
+    "setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, "
+    "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx "
+    "rptr=%u, wptr=%u",
+    queue_->public_handle()->base_address, queue_->public_handle()->id,
+    kDispatchPacketHeader,
+    extractAqlBits(kDispatchPacketHeader,
+                   HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
+    extractAqlBits(kDispatchPacketHeader,
+                   HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
+    extractAqlBits(kDispatchPacketHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
+                   HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
+    extractAqlBits(kDispatchPacketHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
+                   HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
+    packet.setup, packet.grid_size_x, packet.grid_size_y, packet.grid_size_z,
+    packet.workgroup_size_x, packet.workgroup_size_y, packet.workgroup_size_z,
+    packet.private_segment_size, packet.group_segment_size,
+    packet.kernel_object,packet.kernarg_address,
+    completion_signal, queue_->LoadReadIndexRelaxed(), index);
 }
 
 BlitKernel::KernelArgs* BlitKernel::ObtainAsyncKernelCopyArg() {

runtime/hsa-runtime/core/runtime/amd_memory_region.cpp

@@ -2,24 +2,24 @@
 //
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
-// 
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
-// 
+//
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
+//
 // Developed by:
-// 
+//
 //                 AMD Research and AMD HSA Software Development
-// 
+//
 //                 Advanced Micro Devices, Inc.
-// 
+//
 //                 www.amd.com
-// 
+//
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to
 // deal with 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:
-// 
+//
 //  - Redistributions of source code must retain the above copyright notice,
 //    this list of conditions and the following disclaimers.
 //  - Redistributions in binary form must reproduce the above copyright
@@ -29,7 +29,7 @@
 //    nor the names of its contributors may be used to endorse or promote
 //    products derived from this Software without specific prior written
 //    permission.
-// 
+//
 // 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

runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -80,7 +80,7 @@ void HsaApiTable::Init() {
   // they can add preprocessor macros on the new functions
 
   constexpr size_t expected_core_api_table_size = 1016;
-  constexpr size_t expected_amd_ext_table_size = 576;
+  constexpr size_t expected_amd_ext_table_size = 584;
   constexpr size_t expected_image_ext_table_size = 120;
   constexpr size_t expected_finalizer_ext_table_size = 64;
   constexpr size_t expected_tools_table_size = 64;
@@ -466,6 +466,7 @@ void HsaApiTable::UpdateAmdExts() {
       AMD::hsa_amd_vmem_get_alloc_properties_from_handle;
   amd_ext_api.hsa_amd_agent_set_async_scratch_limit_fn = AMD::hsa_amd_agent_set_async_scratch_limit;
   amd_ext_api.hsa_amd_queue_get_info_fn = AMD::hsa_amd_queue_get_info;
+  amd_ext_api.hsa_amd_enable_logging_fn = AMD::hsa_amd_enable_logging;
 }
 
 void HsaApiTable::UpdateTools() {

runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -1415,5 +1415,11 @@ hsa_status_t HSA_API hsa_amd_queue_get_info(hsa_queue_t* _queue,
   CATCH;
 }
 
+hsa_status_t hsa_amd_enable_logging(uint8_t* flags, void *file) {
+  TRY;
+  return core::Runtime::runtime_singleton_->EnableLogging(flags, file);
+  CATCH;
+}
+
 }   //  namespace amd
 }   //  namespace rocr

runtime/hsa-runtime/core/runtime/runtime.cpp

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -125,6 +125,7 @@ hsa_status_t Runtime::Acquire() {
   ScopedAcquire<KernelMutex> boot(&bootstrap_lock_);
 
   if (runtime_singleton_ == NULL) {
+    memset(log_flags, 0, sizeof(log_flags));
     runtime_singleton_ = new Runtime();
   }
 
@@ -3504,5 +3505,16 @@ hsa_status_t Runtime::VMemoryGetAllocPropertiesFromHandle(hsa_amd_vmem_alloc_han
   return HSA_STATUS_SUCCESS;
 }
 
+hsa_status_t Runtime::EnableLogging(uint8_t* flags, void* file) {
+  memcpy(log_flags, flags, sizeof(log_flags));
+
+  if (file)
+    log_file = reinterpret_cast<FILE*>(file);
+  else
+    log_file = stderr;
+
+  return HSA_STATUS_SUCCESS;
+}
+
 }  // namespace core
 }  // namespace rocr

runtime/hsa-runtime/core/util/flag.cpp

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2021-2021, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2021-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -42,6 +42,7 @@
 
 #include "core/util/flag.h"
 #include "core/util/utils.h"
+#include "core/util/os.h"
 
 #include <vector>
 #include <map>
@@ -50,6 +51,22 @@
 #include <locale>
 
 namespace rocr {
+FILE* log_file = stderr;
+uint8_t log_flags[8];
+
+void log_printf(const char* file, int line, const char* format, ...) {
+    va_list ap;
+    std::stringstream str_thrd_id;
+    str_thrd_id << std::hex << std::this_thread::get_id();
+    va_start(ap, format);
+    char message[4096];
+    vsnprintf(message, sizeof(message), format, ap);
+    va_end(ap);
+    fprintf(log_file, ":%-25s:%-4d: %010lld us: [pid:%-5d tid:0x%s] [***rocr***] %s\n",
+            file, line, os::ReadAccurateClock()/1000ULL, os::GetProcessId(),
+            str_thrd_id.str().c_str(), message);
+    fflush(log_file);
+}
 
 // split at separators
 static std::vector<std::string> split(std::string& str, char sep) {

runtime/hsa-runtime/core/util/lnx/os_linux.cpp

@@ -2,24 +2,24 @@
 //
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
-// 
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
-// 
+//
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
+//
 // Developed by:
-// 
+//
 //                 AMD Research and AMD HSA Software Development
-// 
+//
 //                 Advanced Micro Devices, Inc.
-// 
+//
 //                 www.amd.com
-// 
+//
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to
 // deal with 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:
-// 
+//
 //  - Redistributions of source code must retain the above copyright notice,
 //    this list of conditions and the following disclaimers.
 //  - Redistributions in binary form must reproduce the above copyright
@@ -29,7 +29,7 @@
 //    nor the names of its contributors may be used to endorse or promote
 //    products derived from this Software without specific prior written
 //    permission.
-// 
+//
 // 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
@@ -435,6 +435,10 @@ void SetEnvVar(std::string env_var_name, std::string env_var_value) {
   setenv(env_var_name.c_str(), env_var_value.c_str(), 1);
 }
 
+int GetProcessId() {
+  return ::getpid();
+}
+
 std::string GetEnvVar(std::string env_var_name) {
   char* buff;
   buff = getenv(env_var_name.c_str());

runtime/hsa-runtime/core/util/os.h

@@ -2,24 +2,24 @@
 //
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
-// 
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
-// 
+//
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
+//
 // Developed by:
-// 
+//
 //                 AMD Research and AMD HSA Software Development
-// 
+//
 //                 Advanced Micro Devices, Inc.
-// 
+//
 //                 www.amd.com
-// 
+//
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to
 // deal with 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:
-// 
+//
 //  - Redistributions of source code must retain the above copyright notice,
 //    this list of conditions and the following disclaimers.
 //  - Redistributions in binary form must reproduce the above copyright
@@ -29,7 +29,7 @@
 //    nor the names of its contributors may be used to endorse or promote
 //    products derived from this Software without specific prior written
 //    permission.
-// 
+//
 // 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
@@ -246,6 +246,11 @@ void SetEnvVar(std::string env_var_name, std::string env_var_value);
 /// @return: std::string, value of the environment value, returned as string.
 std::string GetEnvVar(std::string env_var_name);
 
+/// @brief: Gets the process ID.
+/// @param: void
+/// @return: int, process ID returned as int.
+int GetProcessId();
+
 /// @brief: Gets the max virtual memory size accessible to the application.
 /// @param: void.
 /// @return: size_t, size of the accessible memory to the application.

runtime/hsa-runtime/core/util/utils.h

@@ -2,24 +2,24 @@
 //
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
-// 
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
-// 
+//
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
+//
 // Developed by:
-// 
+//
 //                 AMD Research and AMD HSA Software Development
-// 
+//
 //                 Advanced Micro Devices, Inc.
-// 
+//
 //                 www.amd.com
-// 
+//
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to
 // deal with 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:
-// 
+//
 //  - Redistributions of source code must retain the above copyright notice,
 //    this list of conditions and the following disclaimers.
 //  - Redistributions in binary form must reproduce the above copyright
@@ -29,7 +29,7 @@
 //    nor the names of its contributors may be used to endorse or promote
 //    products derived from this Software without specific prior written
 //    permission.
-// 
+//
 // 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
@@ -48,12 +48,18 @@
 #include "stdint.h"
 #include "stddef.h"
 #include "stdlib.h"
+#include "stdarg.h"
+#include "unistd.h"
 #include <assert.h>
 #include <iostream>
 #include <string>
 #include <algorithm>
+#include <sstream>
+#include <thread>
 
 namespace rocr {
+extern FILE* log_file;
+extern uint8_t log_flags[8];
 
 typedef unsigned int uint;
 typedef uint64_t uint64;
@@ -69,6 +75,8 @@ typedef uint64_t uint64;
 #define __stdcall  // __attribute__((__stdcall__))
 #define __ALIGNED__(x) __attribute__((aligned(x)))
 
+void log_printf(const char* file, int line, const char* format, ...);
+
 static __forceinline void* _aligned_malloc(size_t size, size_t alignment) {
 #ifdef _ISOC11_SOURCE
   return aligned_alloc(alignment, size);
@@ -135,6 +143,15 @@ static __forceinline unsigned long long int strtoull(const char* str,
 #define ifdebug if (true)
 #endif
 
+#define __FILENAME__ (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__)
+
+#define LogPrint(flag, format, ...)                                                                \
+  do {                                                                                             \
+    if (hsa_flag_isset64(log_flags, flag))                                                         \
+      rocr::log_printf(__FILENAME__, __LINE__, format, ##__VA_ARGS__);                             \
+  } while (false);
+
+
 // A macro to disallow the copy and move constructor and operator= functions
 #define DISALLOW_COPY_AND_ASSIGN(TypeName)                                                         \
   TypeName(const TypeName&) = delete;                                                              \

runtime/hsa-runtime/hsacore.so.def

@@ -261,6 +261,7 @@ global:
 	hsa_ven_amd_pcs_stop;
 	hsa_ven_amd_pcs_flush;
 	hsa_amd_queue_get_info;
+	hsa_amd_enable_logging;
 local:
     *;
 };

runtime/hsa-runtime/inc/hsa_api_trace.h

@@ -265,6 +265,7 @@ struct AmdExtTable {
   decltype(hsa_amd_agent_set_async_scratch_limit)* hsa_amd_agent_set_async_scratch_limit_fn;
   decltype(hsa_amd_queue_get_info)* hsa_amd_queue_get_info_fn;
   decltype(hsa_amd_vmem_address_reserve_align)* hsa_amd_vmem_address_reserve_align_fn;
+  decltype(hsa_amd_enable_logging)* hsa_amd_enable_logging_fn;
 };
 
 // Table to export HSA Core Runtime Apis

runtime/hsa-runtime/inc/hsa_api_trace_version.h

@@ -58,7 +58,7 @@
 // Step Ids of the Api tables exported by Hsa Core Runtime
 #define HSA_API_TABLE_STEP_VERSION                  0x01
 #define HSA_CORE_API_TABLE_STEP_VERSION             0x00
-#define HSA_AMD_EXT_API_TABLE_STEP_VERSION          0x03
+#define HSA_AMD_EXT_API_TABLE_STEP_VERSION          0x04
 #define HSA_FINALIZER_API_TABLE_STEP_VERSION        0x00
 #define HSA_IMAGE_API_TABLE_STEP_VERSION            0x00
 #define HSA_AQLPROFILE_API_TABLE_STEP_VERSION       0x00

runtime/hsa-runtime/inc/hsa_ext_amd.h

@@ -3,7 +3,7 @@
 // The University of Illinois/NCSA
 // Open Source License (NCSA)
 //
-// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
 //
 // Developed by:
 //
@@ -70,7 +70,18 @@ extern "C" {
  */
 
 /**
- * @brief Macro to use to determine that a  flag is set when querying flags within uint8_t[8]
+ * @brief Macro to set a flag within uint8_t[8] types
+ * types
+ */
+static __inline__ __attribute__((always_inline)) void hsa_flag_set64(uint8_t* value,
+                                                                       uint32_t bit) {
+  unsigned int index = bit / 8;
+  unsigned int subBit = bit % 8;
+  (((uint8_t*)value)[index]) |= (1 << subBit);
+}
+
+/**
+ * @brief Macro to use to determine that a flag is set when querying flags within uint8_t[8]
  * types
  */
 static __inline__ __attribute__((always_inline)) bool hsa_flag_isset64(uint8_t* value,
@@ -3140,6 +3151,30 @@ typedef enum {
 hsa_status_t hsa_amd_queue_get_info(hsa_queue_t* queue, hsa_queue_info_attribute_t attribute,
                                     void* value);
 
+/**
+ * @brief logging types
+ */
+typedef enum hsa_amd_log_flag_s {
+   /* Log AQL packets internally enqueued by HSA for Blit Kernels */
+  HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS = 0,
+} hsa_amd_log_flag_t;
+
+/**
+ * @brief Enable logging via external file
+ * If this function is called multiple times, the last call to this function will overwrite the
+ * previous @p flags and @p file.
+ *
+ * @param[in] flags is used to filter types of logging. Type is uint8_t[8].
+ * Can be set using the hsa_flag_set64 macro. Setting @p flags to 0 will disable logging.
+ * @param[in] file file stream to output logging. If file is NULL, prints are sent to stderr.
+ *
+ * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
+ *
+ * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
+ * initialized.
+ */
+hsa_status_t hsa_amd_enable_logging(uint8_t* flags, void* file);
+
 #ifdef __cplusplus
 }  // end extern "C" block
 #endif