SWDEV-549518 - Enable logging dynamically through HIP APIS. (#1079)

* SWDEV-549518 - Enable logging dynamically through HIP APIS. * SWDEV-549518 - Adding ROCProfiler related new API changes. * rocprofiler-sdk changes for hip api additions. --------- Co-authored-by: Venkateshwar Reddy Kandula <venkateshwar.kandula1306@gmail.com> Co-authored-by: jainprad <92369414+jainprad@users.noreply.github.com>
2026-01-19 16:16:14 -05:00
commit 99c3a06f4e
@@ -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

  // ******************************************************************************************* //
  //
@@ -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";
@@ -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
@@ -523,3 +523,6 @@ hipKernelGetName
 hipOccupancyAvailableDynamicSMemPerBlock
 hipGetProcAddress_spt
 hipKernelGetParamInfo
+hipExtDisableLogging
+hipExtEnableLogging
+hipExtSetLoggingParams
@@ -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(<table>, <functor>, 8)
 //
 //  HIP_ENFORCE_ABI_VERSIONING(<table>, 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
@@ -645,6 +645,9 @@ global:
    hipOccupancyAvailableDynamicSMemPerBlock;
    hipGetProcAddress_spt;
    hipKernelGetParamInfo;
+    hipExtDisableLogging;
+    hipExtEnableLogging;
+    hipExtSetLoggingParams;
 local:
    *;
 } hip_7.1;
@@ -0,0 +1,31 @@
+#include <hip/hip_runtime.h>
+#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
@@ -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<hipKernel_t, hipLibrary_t> library_functions_;
+
+  // Logging state (moved from LoggingInfo singleton)
+  size_t log_level_;
+  size_t log_size_;
+  size_t log_mask_;
 };
 }  // namespace hip
@@ -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);
 }
@@ -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"
    ]
 }
@@ -7,6 +7,7 @@ set(TEST_SRC
    hipDrvGetErrorString.cc
    hipGetLastError.cc
    hipPeekAtLastError.cc
+    hipDynamicLogging.cc
 )

 if(UNIX)
@@ -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 <iostream>
+#include <sstream>
+#include <streambuf>
+#include <string>
+#include <fstream>
+#include <fcntl.h>
+
+#ifdef _WIN32
+#include <windows.h>
+#include <io.h>
+#include <stdio.h>
+#include <stdlib.h>
+#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 <unistd.h>
+#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());
+    }
+};
@@ -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 <hip_test_common.hh>
+#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.
+ * @}
+ */
@@ -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
@@ -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)
@@ -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
@@ -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;
@@ -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
@@ -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