From 9547b543d17de78ae7ae4a7003639658d62c23c3 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 22 Sep 2025 11:08:26 -0400 Subject: [PATCH 1/3] SWDEV-549518 - Enable logging dynamically through HIP APIS. --- .../include/hip/amd_detail/hip_api_trace.hpp | 12 +- .../include/hip/amd_detail/hip_prof_str.h | 52 +++++- projects/clr/hipamd/src/CMakeLists.txt | 1 + projects/clr/hipamd/src/amdhip.def | 3 + projects/clr/hipamd/src/hip_api_trace.cpp | 14 +- projects/clr/hipamd/src/hip_hcc.map.in | 3 + projects/clr/hipamd/src/hip_log.cpp | 39 +++++ projects/clr/hipamd/src/hip_log.hpp | 59 +++++++ .../clr/hipamd/src/hip_table_interface.cpp | 9 + .../catch/unit/errorHandling/CMakeLists.txt | 1 + .../catch/unit/errorHandling/OutCapture.hh | 101 +++++++++++ .../unit/errorHandling/hipDynamicLogging.cc | 157 ++++++++++++++++++ projects/hip/include/hip/hip_runtime_api.h | 5 + 13 files changed, 451 insertions(+), 5 deletions(-) create mode 100644 projects/clr/hipamd/src/hip_log.cpp create mode 100644 projects/clr/hipamd/src/hip_log.hpp create mode 100644 projects/hip-tests/catch/unit/errorHandling/OutCapture.hh create mode 100644 projects/hip-tests/catch/unit/errorHandling/hipDynamicLogging.cc diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 84c7f3c3a2d..9b621812100 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -63,7 +63,7 @@ #define HIP_API_TABLE_STEP_VERSION 0 #define HIP_COMPILER_API_TABLE_STEP_VERSION 0 #define HIP_TOOLS_API_TABLE_STEP_VERSION 0 -#define HIP_RUNTIME_API_TABLE_STEP_VERSION 15 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 16 // HIP API interface // HIP compiler dispatch functions @@ -1104,6 +1104,9 @@ typedef hipError_t (*t_hipLibraryGetKernel)(hipKernel_t* pKernel, hipLibrary_t l const char* name); typedef hipError_t (*t_hipLibraryGetKernelCount)(unsigned int *count, hipLibrary_t library); +typedef hipError_t (*t_hipExtDisableLogging)(); +typedef hipError_t (*t_hipExtEnableLogging)(); +typedef hipError_t (*t_hipExtSetLoggingParams)(size_t log_level, size_t log_size, size_t log_mask); // HIP Compiler dispatch table struct HipCompilerDispatchTable { @@ -1679,8 +1682,13 @@ struct HipDispatchTable { t_hipLibraryGetKernel hipLibraryGetKernel_fn; t_hipLibraryGetKernelCount hipLibraryGetKernelCount_fn; + // HIP_RUNTIME_API_TABLE_STEP_VERSION = 16 + t_hipExtDisableLogging hipExtDisableLogging_fn; + t_hipExtEnableLogging hipExtEnableLogging_fn; + t_hipExtSetLoggingParams hipExtSetLoggingParams_fn; + // DO NOT EDIT ABOVE! - // HIP_RUNTIME_API_TABLE_STEP_VERSION == 15 + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 16 // ******************************************************************************************* // // diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 8b93feedf06..04aefb2e08d 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -461,7 +461,10 @@ enum hip_api_id_t { HIP_API_ID_hipLibraryUnload = 441, HIP_API_ID_hipLibraryGetKernel = 442, HIP_API_ID_hipLibraryGetKernelCount = 443, - HIP_API_ID_LAST = 443, + HIP_API_ID_hipExtDisableLogging = 444, + HIP_API_ID_hipExtEnableLogging = 445, + HIP_API_ID_hipExtSetLoggingParams = 446, + HIP_API_ID_LAST = 446, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -585,12 +588,15 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipEventRecord: return "hipEventRecord"; case HIP_API_ID_hipEventRecordWithFlags: return "hipEventRecordWithFlags"; case HIP_API_ID_hipEventSynchronize: return "hipEventSynchronize"; + case HIP_API_ID_hipExtDisableLogging: return "hipExtDisableLogging"; + case HIP_API_ID_hipExtEnableLogging: return "hipExtEnableLogging"; case HIP_API_ID_hipExtGetLastError: return "hipExtGetLastError"; case HIP_API_ID_hipExtGetLinkTypeAndHopCount: return "hipExtGetLinkTypeAndHopCount"; case HIP_API_ID_hipExtLaunchKernel: return "hipExtLaunchKernel"; case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: return "hipExtLaunchMultiKernelMultiDevice"; case HIP_API_ID_hipExtMallocWithFlags: return "hipExtMallocWithFlags"; case HIP_API_ID_hipExtModuleLaunchKernel: return "hipExtModuleLaunchKernel"; + case HIP_API_ID_hipExtSetLoggingParams: return "hipExtSetLoggingParams"; case HIP_API_ID_hipExtStreamCreateWithCUMask: return "hipExtStreamCreateWithCUMask"; case HIP_API_ID_hipExtStreamGetCUMask: return "hipExtStreamGetCUMask"; case HIP_API_ID_hipExternalMemoryGetMappedBuffer: return "hipExternalMemoryGetMappedBuffer"; @@ -1023,12 +1029,15 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipEventRecord", name) == 0) return HIP_API_ID_hipEventRecord; if (strcmp("hipEventRecordWithFlags", name) == 0) return HIP_API_ID_hipEventRecordWithFlags; if (strcmp("hipEventSynchronize", name) == 0) return HIP_API_ID_hipEventSynchronize; + if (strcmp("hipExtDisableLogging", name) == 0) return HIP_API_ID_hipExtDisableLogging; + if (strcmp("hipExtEnableLogging", name) == 0) return HIP_API_ID_hipExtEnableLogging; if (strcmp("hipExtGetLastError", name) == 0) return HIP_API_ID_hipExtGetLastError; if (strcmp("hipExtGetLinkTypeAndHopCount", name) == 0) return HIP_API_ID_hipExtGetLinkTypeAndHopCount; if (strcmp("hipExtLaunchKernel", name) == 0) return HIP_API_ID_hipExtLaunchKernel; if (strcmp("hipExtLaunchMultiKernelMultiDevice", name) == 0) return HIP_API_ID_hipExtLaunchMultiKernelMultiDevice; if (strcmp("hipExtMallocWithFlags", name) == 0) return HIP_API_ID_hipExtMallocWithFlags; if (strcmp("hipExtModuleLaunchKernel", name) == 0) return HIP_API_ID_hipExtModuleLaunchKernel; + if (strcmp("hipExtSetLoggingParams", name) == 0) return HIP_API_ID_hipExtSetLoggingParams; if (strcmp("hipExtStreamCreateWithCUMask", name) == 0) return HIP_API_ID_hipExtStreamCreateWithCUMask; if (strcmp("hipExtStreamGetCUMask", name) == 0) return HIP_API_ID_hipExtStreamGetCUMask; if (strcmp("hipExternalMemoryGetMappedBuffer", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedBuffer; @@ -1834,6 +1843,11 @@ typedef struct hip_api_data_s { hipEvent_t stopEvent; unsigned int flags; } hipExtModuleLaunchKernel; + struct { + size_t log_level; + size_t log_size; + size_t log_mask; + } hipExtSetLoggingParams; struct { hipStream_t* stream; hipStream_t stream__val; @@ -4417,6 +4431,12 @@ typedef struct hip_api_data_s { #define INIT_hipEventSynchronize_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipEventSynchronize.event = (hipEvent_t)event; \ }; +// hipExtDisableLogging[] +#define INIT_hipExtDisableLogging_CB_ARGS_DATA(cb_data) { \ +}; +// hipExtEnableLogging[] +#define INIT_hipExtEnableLogging_CB_ARGS_DATA(cb_data) { \ +}; // hipExtGetLastError[] #define INIT_hipExtGetLastError_CB_ARGS_DATA(cb_data) { \ }; @@ -4468,6 +4488,12 @@ typedef struct hip_api_data_s { cb_data.args.hipExtModuleLaunchKernel.stopEvent = (hipEvent_t)stopEvent; \ cb_data.args.hipExtModuleLaunchKernel.flags = (unsigned int)flags; \ }; +// hipExtSetLoggingParams[('size_t', 'log_level'), ('size_t', 'log_size'), ('size_t', 'log_mask')] +#define INIT_hipExtSetLoggingParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipExtSetLoggingParams.log_level = (size_t)log_level; \ + cb_data.args.hipExtSetLoggingParams.log_size = (size_t)log_size; \ + cb_data.args.hipExtSetLoggingParams.log_mask = (size_t)log_mask; \ +}; // hipExtStreamCreateWithCUMask[('hipStream_t*', 'stream'), ('unsigned int', 'cuMaskSize'), ('const unsigned int*', 'cuMask')] #define INIT_hipExtStreamCreateWithCUMask_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipExtStreamCreateWithCUMask.stream = (hipStream_t*)stream; \ @@ -7026,6 +7052,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipEventSynchronize[('hipEvent_t', 'event')] case HIP_API_ID_hipEventSynchronize: break; +// hipExtDisableLogging[] + case HIP_API_ID_hipExtDisableLogging: + break; +// hipExtEnableLogging[] + case HIP_API_ID_hipExtEnableLogging: + break; // hipExtGetLastError[] case HIP_API_ID_hipExtGetLastError: break; @@ -7051,6 +7083,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipExtModuleLaunchKernel.kernelParams) data->args.hipExtModuleLaunchKernel.kernelParams__val = *(data->args.hipExtModuleLaunchKernel.kernelParams); if (data->args.hipExtModuleLaunchKernel.extra) data->args.hipExtModuleLaunchKernel.extra__val = *(data->args.hipExtModuleLaunchKernel.extra); break; +// hipExtSetLoggingParams[('size_t', 'log_level'), ('size_t', 'log_size'), ('size_t', 'log_mask')] + case HIP_API_ID_hipExtSetLoggingParams: + break; // hipExtStreamCreateWithCUMask[('hipStream_t*', 'stream'), ('unsigned int', 'cuMaskSize'), ('const unsigned int*', 'cuMask')] case HIP_API_ID_hipExtStreamCreateWithCUMask: if (data->args.hipExtStreamCreateWithCUMask.stream) data->args.hipExtStreamCreateWithCUMask.stream__val = *(data->args.hipExtStreamCreateWithCUMask.stream); @@ -8990,6 +9025,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << "event="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventSynchronize.event); oss << ")"; break; + case HIP_API_ID_hipExtDisableLogging: + oss << "hipExtDisableLogging("; + oss << ")"; + break; + case HIP_API_ID_hipExtEnableLogging: + oss << "hipExtEnableLogging("; + oss << ")"; + break; case HIP_API_ID_hipExtGetLastError: oss << "hipExtGetLastError("; oss << ")"; @@ -9054,6 +9097,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExtModuleLaunchKernel.flags); oss << ")"; break; + case HIP_API_ID_hipExtSetLoggingParams: + oss << "hipExtSetLoggingParams("; + oss << "log_level="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExtSetLoggingParams.log_level); + oss << ", log_size="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExtSetLoggingParams.log_size); + oss << ", log_mask="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExtSetLoggingParams.log_mask); + oss << ")"; + break; case HIP_API_ID_hipExtStreamCreateWithCUMask: oss << "hipExtStreamCreateWithCUMask("; if (data->args.hipExtStreamCreateWithCUMask.stream == NULL) oss << "stream=NULL"; diff --git a/projects/clr/hipamd/src/CMakeLists.txt b/projects/clr/hipamd/src/CMakeLists.txt index 756dcd7956b..90f4bcfe6f4 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -111,6 +111,7 @@ target_sources(amdhip64 PRIVATE hip_graph.cpp hip_hmm.cpp hip_intercept.cpp + hip_log.cpp hip_memory.cpp hip_mempool.cpp hip_mempool_impl.cpp diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 89db8a4ef5d..56b13f4fc6a 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -516,3 +516,6 @@ hipLibraryLoadFromFile hipLibraryUnload hipLibraryGetKernel hipLibraryGetKernelCount +hipExtDisableLogging +hipExtEnableLogging +hipExtSetLoggingParams \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 16df4d59bd6..c63347ab34a 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -874,6 +874,9 @@ hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName, hipError_t hipLibraryUnload(hipLibrary_t library); hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const char* name); hipError_t hipLibraryGetKernelCount(unsigned int* count, hipLibrary_t library); +hipError_t hipExtDisableLogging(); +hipError_t hipExtEnableLogging(); +hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); } // namespace hip namespace hip { @@ -1414,6 +1417,9 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipLibraryUnload_fn = hip::hipLibraryUnload; ptrDispatchTable->hipLibraryGetKernel_fn = hip::hipLibraryGetKernel; ptrDispatchTable->hipLibraryGetKernelCount_fn = hip::hipLibraryGetKernelCount; + ptrDispatchTable->hipExtDisableLogging_fn = hip::hipExtDisableLogging; + ptrDispatchTable->hipExtEnableLogging_fn = hip::hipExtEnableLogging; + ptrDispatchTable->hipExtSetLoggingParams_fn = hip::hipExtSetLoggingParams; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -2084,15 +2090,19 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryLoadFromFile_fn, 497); HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryUnload_fn, 498); HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernel_fn, 499); HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernelCount_fn, 500); +HIP_ENFORCE_ABI(HipDispatchTable, hipExtDisableLogging_fn, 501); +HIP_ENFORCE_ABI(HipDispatchTable, hipExtEnableLogging_fn, 502); +HIP_ENFORCE_ABI(HipDispatchTable, hipExtSetLoggingParams_fn, 503); + // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.: // // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 501) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 504) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 15, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 16, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " "pointers and then update this check so it is true"); #endif diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index c9fecec49c7..b3fb5b87a93 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -638,6 +638,9 @@ global: hipLibraryUnload; hipLibraryGetKernel; hipLibraryGetKernelCount; + hipExtDisableLogging; + hipExtEnableLogging; + hipExtSetLoggingParams; local: *; } hip_7.1; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_log.cpp b/projects/clr/hipamd/src/hip_log.cpp new file mode 100644 index 00000000000..e968a559b21 --- /dev/null +++ b/projects/clr/hipamd/src/hip_log.cpp @@ -0,0 +1,39 @@ +#include +#include "hip_internal.hpp" +#include "hip_log.hpp" + +namespace hip { + +// Static instance pointer for LoggingInfo singleton +LoggingInfo* LoggingInfo::lginfo_; + +void LoggingInfo::init() { + amd::ScopedLock lock(lg_lock_); + log_level_ = 0; log_size_ = 0; log_mask_ = 0; +} + +hipError_t hipExtEnableLogging() { + HIP_INIT_API(hipExtEnableLogging); + amd::ScopedLock lock(LoggingInfo::instance().lg_lock_); + AMD_LOG_LEVEL = LoggingInfo::instance().log_level_; + AMD_LOG_MASK = LoggingInfo::instance().log_mask_; + HIP_RETURN(hipSuccess); +} + +hipError_t hipExtDisableLogging() { + HIP_INIT_API(hipExtDisableLogging); + amd::ScopedLock lock(LoggingInfo::instance().lg_lock_); + AMD_LOG_LEVEL = 0; + HIP_RETURN(hipSuccess); +} + +hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask) { + HIP_INIT_API(hipExtSetLoggingParams, log_level, log_size, log_mask); + amd::ScopedLock lock(LoggingInfo::instance().lg_lock_); + // Store logging parameters for later activation + LoggingInfo::instance().log_level_ = log_level; + LoggingInfo::instance().log_size_ = log_size; + LoggingInfo::instance().log_mask_ = log_mask; + HIP_RETURN(hipSuccess); +} +} // namespace::hip \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_log.hpp b/projects/clr/hipamd/src/hip_log.hpp new file mode 100644 index 00000000000..0a0d845c7b3 --- /dev/null +++ b/projects/clr/hipamd/src/hip_log.hpp @@ -0,0 +1,59 @@ +/* Copyright (c) 2015 - 2025 Advanced Micro Devices, Inc. + + 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. */ +#pragma once + +#include +#include +#include "vdi_common.hpp" + +namespace hip { + +class LoggingInfo { + +public: + void init(); + + // Singleton instance + static LoggingInfo& instance() { + if (lginfo_ == nullptr) { + // __hipRegisterFatBinary() will call this when app starts, thus + // there is no multiple entry issue here. + lginfo_ = new LoggingInfo(); + } + return *lginfo_; + } + +private: + size_t log_level_; + size_t log_size_; + size_t log_mask_; + + amd::Monitor lg_lock_{true}; + + // Singleton object + static LoggingInfo* lginfo_; + LoggingInfo() {} + ~LoggingInfo() {} + + friend hipError_t hipExtEnableLogging(); + friend hipError_t hipExtDisableLogging(); + friend hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); +}; +} // namespace::hip_impl \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 1160a72745c..b5493875441 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -2037,4 +2037,13 @@ hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const hipError_t hipLibraryGetKernelCount(unsigned int *count, hipLibrary_t library) { return hip::GetHipDispatchTable()->hipLibraryGetKernelCount_fn(count, library); +} +hipError_t hipExtEnableLogging() { + return hip::GetHipDispatchTable()->hipExtEnableLogging_fn(); +} +hipError_t hipExtDisableLogging() { + return hip::GetHipDispatchTable()->hipExtDisableLogging_fn(); +} +hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask) { + return hip::GetHipDispatchTable()->hipExtSetLoggingParams_fn(log_level, log_size, log_mask); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt b/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt index c932a8d3570..8b262175144 100644 --- a/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt @@ -7,6 +7,7 @@ set(TEST_SRC hipDrvGetErrorString.cc hipGetLastError.cc hipPeekAtLastError.cc + hipDynamicLogging.cc ) if(UNIX) diff --git a/projects/hip-tests/catch/unit/errorHandling/OutCapture.hh b/projects/hip-tests/catch/unit/errorHandling/OutCapture.hh new file mode 100644 index 00000000000..3ec60e981cb --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/OutCapture.hh @@ -0,0 +1,101 @@ +/* +Copyright (c) 2023 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. +*/ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +// Class to capture all stderr output (HIP logging uses stderr) +class OutCapture { +private: + std::stringstream captured_stream_; + std::streambuf* cerr_backup_; + int stderr_backup_; + std::string temp_file_; + +public: + OutCapture() : temp_file_("/tmp/hip_stderr_capture.txt") { + // Backup original cerr stream buffer (HIP logging uses stderr) + cerr_backup_ = std::cerr.rdbuf(); + + // Backup original stderr file descriptor + stderr_backup_ = dup(STDERR_FILENO); + } + + void startCapture() { + // Clear any previous content + captured_stream_.str(""); + captured_stream_.clear(); + + // Redirect std::cerr to our stringstream + std::cerr.rdbuf(captured_stream_.rdbuf()); + + // Redirect stderr file descriptor to temp file (for fprintf to stderr) + int temp_fd = open(temp_file_.c_str(), O_WRONLY | O_CREAT | O_TRUNC, 0644); + if (temp_fd != -1) { + dup2(temp_fd, STDERR_FILENO); + close(temp_fd); + } + } + + std::string stopCapture() { + // Restore original cerr stream + std::cerr.rdbuf(cerr_backup_); + + // Restore original stderr file descriptor + dup2(stderr_backup_, STDERR_FILENO); + + // Read from temp file (captures fprintf(stderr) output from HIP logging) + std::ifstream temp_file(temp_file_); + std::string file_content; + if (temp_file.is_open()) { + std::string line; + while (std::getline(temp_file, line)) { + file_content += line + "\n"; + } + temp_file.close(); + } + + // Combine both captures: C++ streams and file descriptor output + std::string stream_content = captured_stream_.str(); + std::string total_output = stream_content + file_content; + + // Clean up temp file + unlink(temp_file_.c_str()); + + return total_output; + } + + ~OutCapture() { + // Ensure everything is restored + std::cerr.rdbuf(cerr_backup_); + dup2(stderr_backup_, STDERR_FILENO); + close(stderr_backup_); + unlink(temp_file_.c_str()); + } +}; diff --git a/projects/hip-tests/catch/unit/errorHandling/hipDynamicLogging.cc b/projects/hip-tests/catch/unit/errorHandling/hipDynamicLogging.cc new file mode 100644 index 00000000000..59ca1bc7ae3 --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/hipDynamicLogging.cc @@ -0,0 +1,157 @@ +/* +Copyright (c) 2023 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 +#include +#include "OutCapture.hh" + +/** + * @addtogroup hipDynamicLogging hipDynamicLogging + * @{ + * @ingroup ErrorTest + * `hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask)` - + * Sets logging parameters for HIP runtime. + * `hipExtEnableLogging()` - + * Enables HIP runtime logging. + * `hipExtDisableLogging()` - + * Disables HIP runtime logging. + */ + +static bool hipDynamicLoggingTest() { + // Create output capture instance + OutCapture capture; + capture.startCapture(); + + // Set Logging params + HIP_CHECK(hipExtSetLoggingParams(4, 0, -1)); + + // Logging is disabled here - allocate memory + int* dptr = nullptr; + HIP_CHECK(hipMalloc(&dptr, sizeof(int))); + + // Stop capture after hipMalloc and check no output (logging disabled) + std::string malloc_output = capture.stopCapture(); + if (malloc_output.size() != 0) { + INFO("Unexpected logging output during hipMalloc (logging should be disabled): " << malloc_output); + return false; + } + + // Start capture before enabling logging + capture.startCapture(); + + // Enable logging and do memset + HIP_CHECK(hipExtEnableLogging()); + HIP_CHECK(hipMemset(dptr, 0x00, sizeof(int))); + + // Disable logging + HIP_CHECK(hipExtDisableLogging()); + + // Stop capture after disabling logging and check for output + std::string logging_output = capture.stopCapture(); + if (logging_output.size() == 0) { + INFO("Expected logging output during enabled logging period, but got none"); + return false; + } + + // Clean up + HIP_CHECK(hipFree(dptr)); + + INFO("Successfully captured HIP logging output (" << logging_output.size() << " bytes)"); + INFO("Logging output: " << logging_output); + + return true; +} + +/** + * Test Description + * ------------------------ + * - Validates that HIP dynamic logging works correctly: + * 1. No output when logging is disabled + * 2. Logging output is captured when logging is enabled + * 3. hipMemset operation produces logging output during enabled period + * Test source + * ------------------------ + * - unit/errorHandling/hipDynamicLogging.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipDynamicLogging_Positive_Basic") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + HipTest::HIP_SKIP_TEST("Skipping hipDynamicLogging test - no devices available"); + return; + } + + REQUIRE(hipDynamicLoggingTest() == true); +} + +/** + * Test Description + * ------------------------ + * - Validates that hipExtSetLoggingParams sets logging parameters correctly + * and that logging can be enabled/disabled multiple times + * Test source + * ------------------------ + * - unit/errorHandling/hipDynamicLogging.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipDynamicLogging_Positive_MultipleEnableDisable") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + HipTest::HIP_SKIP_TEST("Skipping hipDynamicLogging test - no devices available"); + return; + } + + // Test multiple enable/disable cycles + OutCapture capture; + int* dptr = nullptr; + HIP_CHECK(hipMalloc(&dptr, sizeof(int))); + + // Set different logging parameters + HIP_CHECK(hipExtSetLoggingParams(3, 0, -1)); + + for (int i = 0; i < 3; ++i) { + // Start capture and enable logging + capture.startCapture(); + HIP_CHECK(hipExtEnableLogging()); + HIP_CHECK(hipMemset(dptr, 0x42, sizeof(int))); + HIP_CHECK(hipExtDisableLogging()); + + // Check that we captured some output + std::string output = capture.stopCapture(); + REQUIRE(output.size() > 0); + } + + HIP_CHECK(hipFree(dptr)); +} + +/** + * End doxygen group ErrorTest. + * @} + */ diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index d6ac3fc6d38..ffebcd2f744 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -9526,6 +9526,11 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject); /** * @} */ + +hipError_t hipExtEnableLogging(); +hipError_t hipExtDisableLogging(); +hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); + #ifdef __cplusplus } /* extern "c" */ #endif From 24e3d50ad5d0aabc4c8d9b009e20caa6620a90e8 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Thu, 13 Nov 2025 07:10:17 -0500 Subject: [PATCH 2/3] SWDEV-549518 - Adding ROCProfiler related new API changes. --- .../source/include/rocprofiler-sdk/cxx/enum_string.hpp | 7 +++++++ .../source/include/rocprofiler-sdk/hip/api_args.h | 9 +++++++++ .../source/include/rocprofiler-sdk/hip/runtime_api_id.h | 5 +++++ .../source/lib/rocprofiler-sdk/hip/abi.cpp | 8 ++++++++ .../source/lib/rocprofiler-sdk/hip/hip.def.cpp | 5 +++++ 5 files changed, 34 insertions(+) diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp index 514d56e6327..519f0810211 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp @@ -1004,6 +1004,11 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipGetProcAddress_spt) #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo) #endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 +ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipExtDisableLogging) +ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipExtEnableLogging) +ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipExtSetLoggingParams) +#endif #if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0 static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 442); #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1 @@ -1046,6 +1051,8 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 506); static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 507); #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 508); +#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 21 +static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 511); #else # if !defined(ROCPROFILER_UNSAFE_NO_VERSION_CHECK) && \ (defined(ROCPROFILER_CI) && ROCPROFILER_CI > 0) diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h index 17634ce0c84..8aac7ff1e42 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h @@ -3379,4 +3379,13 @@ typedef union rocprofiler_hip_api_args_t #endif } rocprofiler_hip_api_args_t; +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 + struct + { + size_t log_level; + size_t log_size; + size_t log_mask; + } hipExtSetLoggingParams; +#endif + ROCPROFILER_EXTERN_C_FINI diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h index 4c2475ddaf2..b72d4f2cab3 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h @@ -575,6 +575,11 @@ typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size) #endif #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo, +#endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 + ROCPROFILER_HIP_RUNTIME_API_ID_hipExtDisableLogging, + ROCPROFILER_HIP_RUNTIME_API_ID_hipExtEnableLogging, + ROCPROFILER_HIP_RUNTIME_API_ID_hipExtSetLoggingParams, #endif ROCPROFILER_HIP_RUNTIME_API_ID_LAST, } rocprofiler_hip_runtime_api_id_t; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp index 5fa4222ad1c..ac149de1d40 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp @@ -623,6 +623,12 @@ ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipGetProcAddress_spt_fn, 506); ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipKernelGetParamInfo_fn, 507); #endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 +ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipExtDisableLogging_fn, 508); +ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipExtEnableLogging_fn, 509); +ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipExtSetLoggingParams_fn, 510); +#endif + #if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0 ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 442) #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1 @@ -665,6 +671,8 @@ ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 506) ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 507) #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 508) +#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 21 +ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 511) #else INTERNAL_CI_ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 0) #endif diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp index 44cf89d6c1b..5ace9dd6a47 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp @@ -650,6 +650,11 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo, hipKernelGetParamInfo, hipKernelGetParamInfo_fn, kernel, paramIndex, paramOffset, paramSize); #endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 +HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtDisableLogging, hipExtDisableLogging, hipExtDisableLogging_fn); +HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtEnableLogging, hipExtEnableLogging, hipExtEnableLogging_fn); +HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtSetLoggingParams, hipExtSetLoggingParams, hipExtSetLoggingParams_fn, log_level, log_size, log_mask); +#endif // clang-format on #else From e446768becbb7c69cc971825e9d3c8c91dcbf8cb Mon Sep 17 00:00:00 2001 From: Venkateshwar Reddy Kandula Date: Tue, 16 Dec 2025 15:49:48 +0000 Subject: [PATCH 3/3] rocprofiler-sdk changes for hip api additions. --- .../source/include/rocprofiler-sdk/hip/api_args.h | 15 +++++++++++++-- .../source/lib/rocprofiler-sdk/hip/hip.def.cpp | 6 +++--- 2 files changed, 16 insertions(+), 5 deletions(-) diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h index 8aac7ff1e42..1e5608b1374 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h @@ -3377,9 +3377,19 @@ typedef union rocprofiler_hip_api_args_t size_t* paramSize; } hipKernelGetParamInfo; #endif -} rocprofiler_hip_api_args_t; - #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 + struct + { + // Empty struct has a size of 0 in C but size of 1 in C++. + // Add the rocprofiler_hip_api_no_args struct to fix this + rocprofiler_hip_api_no_args no_args; + } hipExtDisableLogging; + struct + { + // Empty struct has a size of 0 in C but size of 1 in C++. + // Add the rocprofiler_hip_api_no_args struct to fix this + rocprofiler_hip_api_no_args no_args; + } hipExtEnableLogging; struct { size_t log_level; @@ -3387,5 +3397,6 @@ typedef union rocprofiler_hip_api_args_t size_t log_mask; } hipExtSetLoggingParams; #endif +} rocprofiler_hip_api_args_t; ROCPROFILER_EXTERN_C_FINI diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp index 5ace9dd6a47..13e4cd1c76e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp @@ -651,10 +651,10 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT #endif #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 21 -HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtDisableLogging, hipExtDisableLogging, hipExtDisableLogging_fn); -HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtEnableLogging, hipExtEnableLogging, hipExtEnableLogging_fn); +HIP_API_INFO_DEFINITION_0(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtDisableLogging, hipExtDisableLogging, hipExtDisableLogging_fn); +HIP_API_INFO_DEFINITION_0(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtEnableLogging, hipExtEnableLogging, hipExtEnableLogging_fn); HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipExtSetLoggingParams, hipExtSetLoggingParams, hipExtSetLoggingParams_fn, log_level, log_size, log_mask); -#endif +#endif // clang-format on #else