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 f1864eaca4d..7845b69e8a8 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 20 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 21 // HIP API interface // HIP compiler dispatch functions @@ -1113,6 +1113,9 @@ typedef hipError_t (*t_hipKernelGetLibrary)(hipLibrary_t* library, hipKernel_t k typedef hipError_t (*t_hipKernelGetName)(const char** name, hipKernel_t kernel); typedef hipError_t (*t_hipGetProcAddress_spt)(const char* symbol, void** pfn, int hipVersion, uint64_t flags, hipDriverProcAddressQueryResult* symbolStatus); +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); typedef hipError_t (*t_hipKernelGetParamInfo)(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, size_t* paramSize); @@ -1707,8 +1710,13 @@ struct HipDispatchTable { // HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 t_hipKernelGetParamInfo hipKernelGetParamInfo_fn; - // DO NOT EDIT ABOVE! // HIP_RUNTIME_API_TABLE_STEP_VERSION == 21 + t_hipExtDisableLogging hipExtDisableLogging_fn; + t_hipExtEnableLogging hipExtEnableLogging_fn; + t_hipExtSetLoggingParams hipExtSetLoggingParams_fn; + + // DO NOT EDIT ABOVE! + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 22 // ******************************************************************************************* // // 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 7aef3e7f52f..2dc1a0d3b29 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 @@ -468,7 +468,10 @@ enum hip_api_id_t { HIP_API_ID_hipKernelGetName = 448, HIP_API_ID_hipOccupancyAvailableDynamicSMemPerBlock = 449, HIP_API_ID_hipKernelGetParamInfo = 450, - HIP_API_ID_LAST = 450, + HIP_API_ID_hipExtDisableLogging = 451, + HIP_API_ID_hipExtEnableLogging = 452, + HIP_API_ID_hipExtSetLoggingParams = 453, + HIP_API_ID_LAST = 453, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -590,12 +593,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"; @@ -1034,12 +1040,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; @@ -1851,6 +1860,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; @@ -4484,6 +4498,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) { \ }; @@ -4535,6 +4555,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; \ @@ -7125,6 +7151,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; @@ -7150,6 +7182,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); @@ -9124,6 +9159,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 << ")"; @@ -9188,6 +9231,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 aa8d3076886..0f7c537ba42 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -110,6 +110,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 dff046d1acf..9a12cbf0806 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -523,3 +523,6 @@ hipKernelGetName hipOccupancyAvailableDynamicSMemPerBlock hipGetProcAddress_spt hipKernelGetParamInfo +hipExtDisableLogging +hipExtEnableLogging +hipExtSetLoggingParams diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 6399924ec21..564f0bcceae 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -885,6 +885,9 @@ hipError_t hipOccupancyAvailableDynamicSMemPerBlock(size_t* dynamicSmemSize, con int numBlocks, int blockSize); hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, size_t* paramSize); +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 { @@ -1432,6 +1435,9 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipKernelGetName_fn = hip::hipKernelGetName; ptrDispatchTable->hipOccupancyAvailableDynamicSMemPerBlock_fn = hip::hipOccupancyAvailableDynamicSMemPerBlock; ptrDispatchTable->hipKernelGetParamInfo_fn = hip::hipKernelGetParamInfo; + ptrDispatchTable->hipExtDisableLogging_fn = hip::hipExtDisableLogging; + ptrDispatchTable->hipExtEnableLogging_fn = hip::hipExtEnableLogging; + ptrDispatchTable->hipExtSetLoggingParams_fn = hip::hipExtSetLoggingParams; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -2114,15 +2120,19 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipOccupancyAvailableDynamicSMemPerBlock_fn, 5 HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_spt_fn, 506); // HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 HIP_ENFORCE_ABI(HipDispatchTable, hipKernelGetParamInfo_fn, 507); +// HIP_RUNTIME_API_TABLE_STEP_VERSION == 21 +HIP_ENFORCE_ABI(HipDispatchTable, hipExtDisableLogging_fn, 508); +HIP_ENFORCE_ABI(HipDispatchTable, hipExtEnableLogging_fn, 509); +HIP_ENFORCE_ABI(HipDispatchTable, hipExtSetLoggingParams_fn, 510); // 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, 508) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 511) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 20, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 21, "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 37437589e8c..79c15835a5e 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -645,6 +645,9 @@ global: hipOccupancyAvailableDynamicSMemPerBlock; hipGetProcAddress_spt; hipKernelGetParamInfo; + hipExtDisableLogging; + hipExtEnableLogging; + hipExtSetLoggingParams; local: *; } hip_7.1; diff --git a/projects/clr/hipamd/src/hip_log.cpp b/projects/clr/hipamd/src/hip_log.cpp new file mode 100644 index 00000000000..6d43c0dc770 --- /dev/null +++ b/projects/clr/hipamd/src/hip_log.cpp @@ -0,0 +1,31 @@ +#include +#include "hip_internal.hpp" +#include "hip_platform.hpp" + +namespace hip { + +hipError_t hipExtEnableLogging() { + HIP_INIT_API(hipExtEnableLogging); + amd::ScopedLock lock(PlatformState::instance().getLogLock()); + AMD_LOG_LEVEL = PlatformState::instance().log_level_; + AMD_LOG_MASK = PlatformState::instance().log_mask_; + HIP_RETURN(hipSuccess); +} + +hipError_t hipExtDisableLogging() { + HIP_INIT_API(hipExtDisableLogging); + amd::ScopedLock lock(PlatformState::instance().getLogLock()); + 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(PlatformState::instance().getLogLock()); + // Store logging parameters for later activation + PlatformState::instance().log_level_ = log_level; + PlatformState::instance().log_size_ = log_size; + PlatformState::instance().log_mask_ = log_mask; + HIP_RETURN(hipSuccess); +} +} // namespace::hip \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_platform.hpp b/projects/clr/hipamd/src/hip_platform.hpp index bd714206efe..ba8261f3d22 100644 --- a/projects/clr/hipamd/src/hip_platform.hpp +++ b/projects/clr/hipamd/src/hip_platform.hpp @@ -50,9 +50,12 @@ class PlatformState { // Unique FD Store Lock amd::Monitor ufd_lock_{true}; + // Lock for logging operations + amd::Monitor lg_lock_{true}; + // Singleton object static PlatformState* platform_; - PlatformState() {} + PlatformState() : log_level_(0), log_size_(0), log_mask_(0) {} ~PlatformState() {} public: @@ -113,6 +116,14 @@ class PlatformState { size_t UfdMapSize() const { return ufd_map_.size(); } + // Logging lock accessor + amd::Monitor& getLogLock() { return lg_lock_; } + + // Friend functions for logging access + friend hipError_t hipExtEnableLogging(); + friend hipError_t hipExtDisableLogging(); + friend hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); + inline bool RegisterLibraryFunction(const hipKernel_t f, const hipLibrary_t l) { amd::ScopedLock lock(lock_); if (library_functions_.find(f) == library_functions_.end()) { @@ -150,5 +161,10 @@ class PlatformState { void* dynamicLibraryHandle_{nullptr}; std::unordered_map library_functions_; + + // Logging state (moved from LoggingInfo singleton) + size_t log_level_; + size_t log_size_; + size_t log_mask_; }; } // namespace hip diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 443ba81ba5e..d8783c9ee0d 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -2067,4 +2067,13 @@ hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* size_t* paramSize) { return hip::GetHipDispatchTable()->hipKernelGetParamInfo_fn(kernel, paramIndex, paramOffset, paramSize); +} +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/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 85e17cd3810..24e5cefcf5b 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -752,6 +752,8 @@ #endif "=== Following tests disabled as it should be a local perf test", "Performance_hipExtLaunchKernelGGL_QueryGPUFrequency", + "Unit_hipDynamicLogging_Positive_Basic", + "Unit_hipDynamicLogging_Positive_MultipleEnableDisable", "End of json" ] } diff --git a/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt b/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt index b1e9220bfc0..873417baf7f 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..febd7e429f4 --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/OutCapture.hh @@ -0,0 +1,134 @@ +/* +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 + +#ifdef _WIN32 +#include +#include +#include +#include +#define dup _dup +#define dup2 _dup2 +#define fd_close _close +#define unlink _unlink +#define STDERR_FD _fileno(stderr) +#define OPEN_FLAGS (_O_WRONLY | _O_CREAT | _O_TRUNC) +#define OPEN_MODE (_S_IREAD | _S_IWRITE) +#define open _open +#else +#include +#define fd_close close +#define STDERR_FD STDERR_FILENO +#define OPEN_FLAGS (O_WRONLY | O_CREAT | O_TRUNC) +#define OPEN_MODE 0644 +#endif + +// 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_; + + static std::string getTempFilePath() { +#ifdef _WIN32 + char temp_path[MAX_PATH]; + if (GetTempPathA(MAX_PATH, temp_path)) { + return std::string(temp_path) + "hip_stderr_capture.txt"; + } + // Fallback to current directory + return "hip_stderr_capture.txt"; +#else + return "/tmp/hip_stderr_capture.txt"; +#endif + } + +public: + OutCapture() : temp_file_(getTempFilePath()) { + // Backup original cerr stream buffer (HIP logging uses stderr) + cerr_backup_ = std::cerr.rdbuf(); + + // Backup original stderr file descriptor + stderr_backup_ = dup(STDERR_FD); + } + + 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(), OPEN_FLAGS, OPEN_MODE); + if (temp_fd != -1) { + dup2(temp_fd, STDERR_FD); + fd_close(temp_fd); + } + } + + std::string stopCapture() { + // Restore original cerr stream + std::cerr.rdbuf(cerr_backup_); + + // Restore original stderr file descriptor + dup2(stderr_backup_, STDERR_FD); + + // 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_FD); + fd_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..d242b12bb1e --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/hipDynamicLogging.cc @@ -0,0 +1,156 @@ +/* +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 "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 a7bc73deb11..6f6aee86fec 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -9615,6 +9615,45 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject); /** * @} */ + +/** + * @brief Enable HIP runtime logging. + * + * This function enables the HIP runtime logging mechanism, allowing diagnostic + * and trace information to be captured during HIP API execution. + * + * @returns #hipSuccess + * + * @see hipExtDisableLogging, hipExtSetLoggingParams + */ +hipError_t hipExtEnableLogging(); +/** + * @brief Disable HIP runtime logging. + * + * This function disables the HIP runtime logging mechanism, stopping the capture + * of diagnostic and trace information during HIP API execution. + * + * @returns #hipSuccess + * + * @see hipExtEnableLogging, hipExtSetLoggingParams + */ +hipError_t hipExtDisableLogging(); +/** + * @brief Set HIP runtime logging parameters. + * + * This function configures the logging behavior of the HIP runtime, including + * the verbosity level, buffer size, and which components to log. + * + * @param [in] log_level The logging verbosity level. Higher values produce more detailed output. + * @param [in] log_size Reserved for future use. Currently not implemented. + * @param [in] log_mask A bitmask specifying which HIP runtime components to log. + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * @see hipExtEnableLogging, hipExtDisableLogging + */ +hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); + #ifdef __cplusplus } /* extern "c" */ #endif 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..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,6 +3377,26 @@ typedef union rocprofiler_hip_api_args_t size_t* paramSize; } hipKernelGetParamInfo; #endif +#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; + size_t log_size; + size_t log_mask; + } hipExtSetLoggingParams; +#endif } rocprofiler_hip_api_args_t; 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..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 @@ -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_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 // clang-format on #else