diff --git a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp index 31196835c8..fc42ac8a16 100644 --- a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp +++ b/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 { diff --git a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h b/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h index e5717b3fae..5109d39761 100644 --- a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h +++ b/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 diff --git a/runtime/hsa-runtime/core/inc/runtime.h b/runtime/hsa-runtime/core/inc/runtime.h index 3d4d889d8c..6bbff2121c 100644 --- a/runtime/hsa-runtime/core/inc/runtime.h +++ b/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& cpu_agents() { return cpu_agents_; } const std::vector& gpu_agents() { return gpu_agents_; } diff --git a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index 9ede1ea301..81b930cc40 100644 --- a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/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() { diff --git a/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp b/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp index 7d38a83b4f..9001b2d628 100644 --- a/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp +++ b/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 diff --git a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index 48dee48585..02d1e6bee9 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/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() { diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index ce8d9256c6..b48c9e4567 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/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 diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index 9ccef65cc6..a64d45c8ed 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/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 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); + else + log_file = stderr; + + return HSA_STATUS_SUCCESS; +} + } // namespace core } // namespace rocr diff --git a/runtime/hsa-runtime/core/util/flag.cpp b/runtime/hsa-runtime/core/util/flag.cpp index f772ea9827..c0ac8bd970 100644 --- a/runtime/hsa-runtime/core/util/flag.cpp +++ b/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 #include @@ -50,6 +51,22 @@ #include 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 split(std::string& str, char sep) { diff --git a/runtime/hsa-runtime/core/util/lnx/os_linux.cpp b/runtime/hsa-runtime/core/util/lnx/os_linux.cpp index d36dc0d5d7..aecca6c0fd 100644 --- a/runtime/hsa-runtime/core/util/lnx/os_linux.cpp +++ b/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()); diff --git a/runtime/hsa-runtime/core/util/os.h b/runtime/hsa-runtime/core/util/os.h index 50a7b0fbd4..2eec51a34e 100644 --- a/runtime/hsa-runtime/core/util/os.h +++ b/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. diff --git a/runtime/hsa-runtime/core/util/utils.h b/runtime/hsa-runtime/core/util/utils.h index 1a454d7bd4..a1479d187e 100644 --- a/runtime/hsa-runtime/core/util/utils.h +++ b/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 #include #include #include +#include +#include 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; \ diff --git a/runtime/hsa-runtime/hsacore.so.def b/runtime/hsa-runtime/hsacore.so.def index dd9b554a18..3a575367f4 100644 --- a/runtime/hsa-runtime/hsacore.so.def +++ b/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: *; }; diff --git a/runtime/hsa-runtime/inc/hsa_api_trace.h b/runtime/hsa-runtime/inc/hsa_api_trace.h index 2a0f59df3b..e0063e6dab 100644 --- a/runtime/hsa-runtime/inc/hsa_api_trace.h +++ b/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 diff --git a/runtime/hsa-runtime/inc/hsa_api_trace_version.h b/runtime/hsa-runtime/inc/hsa_api_trace_version.h index 3393a77620..2033f0bc2e 100644 --- a/runtime/hsa-runtime/inc/hsa_api_trace_version.h +++ b/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 diff --git a/runtime/hsa-runtime/inc/hsa_ext_amd.h b/runtime/hsa-runtime/inc/hsa_ext_amd.h index f9f60edeb9..ae3bcfdb24 100644 --- a/runtime/hsa-runtime/inc/hsa_ext_amd.h +++ b/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