diff --git a/source/lib/common/demangle.cpp b/source/lib/common/demangle.cpp
index cf316deb9f..db1da45c3f 100644
--- a/source/lib/common/demangle.cpp
+++ b/source/lib/common/demangle.cpp
@@ -21,6 +21,7 @@
 // THE SOFTWARE.
 
 #include "lib/common/demangle.hpp"
+#include "lib/common/logging.hpp"
 
 #include <glog/logging.h>
 
@@ -79,14 +80,14 @@ cxa_demangle(std::string_view _mangled_name, int* _status)
         }
         case -1:
         {
-            PLOG(ERROR) << "memory allocation failure occurred demangling " << _demangled_name;
+            ROCP_ERROR << "memory allocation failure occurred demangling " << _demangled_name;
             break;
         }
         case -2: break;
         case -3:
         {
-            PLOG(ERROR) << "Invalid argument in: (\"" << _demangled_name << "\", nullptr, nullptr, "
-                        << _status << ")";
+            ROCP_ERROR << "Invalid argument in: (\"" << _demangled_name << "\", nullptr, nullptr, "
+                       << _status << ")";
             break;
         }
         default: break;
diff --git a/source/lib/common/environment.cpp b/source/lib/common/environment.cpp
index 9065309392..6f464e5b54 100644
--- a/source/lib/common/environment.cpp
+++ b/source/lib/common/environment.cpp
@@ -22,6 +22,7 @@
 
 #include "lib/common/environment.hpp"
 #include "lib/common/demangle.hpp"
+#include "lib/common/logging.hpp"
 
 #include <cctype>
 #include <cstdint>
@@ -103,7 +104,7 @@ get_env(std::string_view env_id, Tp _default, std::enable_if_t<std::is_integral<
                 return static_cast<Tp>(std::stoul(env_var));
         } catch(std::exception& _e)
         {
-            LOG(ERROR) << "[rocprofiler][get_env] Exception thrown converting getenv(\"" << env_id
+            ROCP_ERROR << "[rocprofiler][get_env] Exception thrown converting getenv(\"" << env_id
                        << "\") = " << env_var << " to " << cxx_demangle(typeid(Tp).name())
                        << " :: " << _e.what() << ". Using default value of " << _default << "\n";
         }
diff --git a/source/lib/common/logging.cpp b/source/lib/common/logging.cpp
index be20fd6c06..b1da8529bb 100644
--- a/source/lib/common/logging.cpp
+++ b/source/lib/common/logging.cpp
@@ -25,6 +25,7 @@
 
 #include <fmt/format.h>
 #include <glog/logging.h>
+#include <glog/vlog_is_on.h>
 
 #include <fstream>
 #include <mutex>
@@ -65,25 +66,31 @@ init_logging(std::string_view env_var, logging_config cfg)
         for(auto& itr : loglvl)
             itr = tolower(itr);
         // default to warning
-        auto& loglvl_v = cfg.loglevel;
+        auto& loglvl_v   = cfg.loglevel;
+        auto& vlog_level = cfg.vlog_level;
         if(!loglvl.empty() && loglvl.find_first_not_of("0123456789") == std::string::npos)
         {
-            loglvl_v = std::stoul(loglvl);
+            loglvl_v   = std::stoul(loglvl);
+            vlog_level = loglvl_v;
         }
         else if(!loglvl.empty())
         {
-            const auto opts =
-                std::unordered_map<std::string_view, uint32_t>{{"info", google::INFO},
-                                                               {"warning", google::WARNING},
-                                                               {"error", google::ERROR},
-                                                               {"fatal", google::FATAL}};
+            const auto opts = std::unordered_map<std::string_view, std::pair<uint32_t, uint32_t>>{
+                {"trace", {google::INFO, ROCP_LEVEL_TRACE}},
+                {"info", {google::INFO, ROCP_LEVEL_INFO}},
+                {"warning", {google::WARNING, ROCP_LEVEL_WARNING}},
+                {"error", {google::ERROR, ROCP_NO_VLOG}},
+                {"fatal", {google::ERROR, ROCP_NO_VLOG}}};
             if(opts.find(loglvl) == opts.end())
-                throw std::runtime_error{
-                    fmt::format("invalid specifier for ROCPROFILER_LOG_LEVEL: {}. Supported: info, "
-                                "warning, error, fatal",
-                                loglvl)};
+                throw std::runtime_error{fmt::format(
+                    "invalid specifier for ROCPROFILER_LOG_LEVEL: {}. Supported: trace, info, "
+                    "warning, error, fatal",
+                    loglvl)};
             else
-                loglvl_v = opts.at(loglvl);
+            {
+                loglvl_v   = opts.at(loglvl).first;
+                vlog_level = opts.at(loglvl).second;
+            }
         }
 
         update_logging(cfg, true);
@@ -91,12 +98,24 @@ init_logging(std::string_view env_var, logging_config cfg)
         if(!google::IsGoogleLoggingInitialized())
         {
             static auto argv0 = get_argv0();
+            // Prevent glog from crashing if vmodule is empty
+            if(FLAGS_vmodule.empty())
+            {
+                FLAGS_vmodule = " ";
+            }
+
             google::InitGoogleLogging(argv0.c_str());
+            ROCP_WARNING << "Log Level: " << loglvl << " VLOG Level: " << vlog_level;
+
+            // Swap out memory to avoid leaking the string
+            if(FLAGS_vmodule == " ")
+            {
+                std::string().swap(FLAGS_vmodule);
+            }
         }
 
         update_logging(cfg);
-
-        LOG(INFO) << "logging initialized via " << env_var;
+        ROCP_INFO << "logging initialized via " << env_var;
     });
 }
 
@@ -111,10 +130,13 @@ update_logging(const logging_config& cfg, bool setup_env, int env_override)
     FLAGS_stderrthreshold           = cfg.loglevel;
     FLAGS_logtostderr               = cfg.logtostderr;
     FLAGS_alsologtostderr           = cfg.alsologtostderr;
+    FLAGS_v                         = cfg.vlog_level;
+
     if(cfg.install_failure_handler) install_failure_signal_handler();
 
     if(setup_env)
     {
+        common::set_env("GLOG_v", cfg.vlog_level, env_override);
         common::set_env("GOOGLE_LOG_DIR", get_env("PWD", ""), env_override);
     }
 }
diff --git a/source/lib/common/logging.hpp b/source/lib/common/logging.hpp
index 5ea977be52..33adf983a9 100644
--- a/source/lib/common/logging.hpp
+++ b/source/lib/common/logging.hpp
@@ -27,6 +27,18 @@
 #include <cstdint>
 #include <string_view>
 
+#define ROCP_LEVEL_TRACE   12
+#define ROCP_LEVEL_INFO    11
+#define ROCP_LEVEL_WARNING 10
+#define ROCP_NO_VLOG       -1
+
+#define ROCP_TRACE   VLOG(ROCP_LEVEL_TRACE)
+#define ROCP_INFO    VLOG(ROCP_LEVEL_INFO)
+#define ROCP_WARNING VLOG(ROCP_LEVEL_WARNING)
+#define ROCP_ERROR   LOG(ERROR)
+#define ROCP_FATAL   LOG(FATAL)
+#define ROCP_DFATAL  DLOG(FATAL)
+
 namespace rocprofiler
 {
 namespace common
@@ -36,6 +48,7 @@ struct logging_config
     bool    install_failure_handler = false;
     bool    logtostderr             = true;
     bool    alsologtostderr         = false;
+    int32_t vlog_level              = ROCP_NO_VLOG;
     int32_t loglevel                = google::WARNING;
 };
 
diff --git a/source/lib/common/utility.cpp b/source/lib/common/utility.cpp
index 8ab6bea2dc..945665b19a 100644
--- a/source/lib/common/utility.cpp
+++ b/source/lib/common/utility.cpp
@@ -22,6 +22,7 @@
 //
 
 #include "lib/common/utility.hpp"
+#include "lib/common/logging.hpp"
 
 #include <glog/logging.h>
 
@@ -77,13 +78,13 @@ get_clock_period_ns_impl(clockid_t _clk_id)
     if(ROCPROFILER_UNLIKELY(ret != 0))
     {
         auto _err = errno;
-        LOG(FATAL) << "error getting clock resolution for " << get_clock_name(_clk_id) << ": "
+        ROCP_FATAL << "error getting clock resolution for " << get_clock_name(_clk_id) << ": "
                    << strerror(_err);
     }
     else if(ROCPROFILER_UNLIKELY(ts.tv_sec != 0 ||
                                  ts.tv_nsec >= std::numeric_limits<uint32_t>::max()))
     {
-        LOG(FATAL) << "clock_getres(" << get_clock_name(_clk_id)
+        ROCP_FATAL << "clock_getres(" << get_clock_name(_clk_id)
                    << ") returned very low frequency (<1Hz)";
     }
 
diff --git a/source/lib/common/utility.hpp b/source/lib/common/utility.hpp
index 500b0e28e3..29b90520b4 100644
--- a/source/lib/common/utility.hpp
+++ b/source/lib/common/utility.hpp
@@ -23,6 +23,7 @@
 #pragma once
 
 #include "lib/common/defines.hpp"
+#include "lib/common/logging.hpp"
 
 #include <glog/logging.h>
 
@@ -69,7 +70,7 @@ get_ticks(clockid_t clk_id_v) noexcept
     if(ROCPROFILER_UNLIKELY(ret != 0))
     {
         auto _err = errno;
-        LOG(FATAL) << "clock_gettime failed: " << strerror(_err);
+        ROCP_FATAL << "clock_gettime failed: " << strerror(_err);
     }
 
     return (static_cast<uint64_t>(ts.tv_sec) * nanosec) + static_cast<uint64_t>(ts.tv_nsec);
diff --git a/source/lib/common/xml.cpp b/source/lib/common/xml.cpp
index 387031b0d7..e62f41d294 100644
--- a/source/lib/common/xml.cpp
+++ b/source/lib/common/xml.cpp
@@ -21,6 +21,7 @@
 // THE SOFTWARE.
 
 #include "lib/common/xml.hpp"
+#include "lib/common/logging.hpp"
 
 #include <glog/logging.h>
 
@@ -304,7 +305,7 @@ Xml::Process()
                 break;
             default:
             {
-                LOG(ERROR) << "XML parser error: wrong state: " << state_;
+                ROCP_ERROR << "XML parser error: wrong state: " << state_;
                 abort();
             }
         }
@@ -418,7 +419,7 @@ void
 Xml::BadFormat(token_t token)
 {
     token.push_back('\0');
-    LOG(ERROR) << "Error: " << file_name_ << ", line " << file_line_ << ", bad XML token '"
+    ROCP_ERROR << "Error: " << file_name_ << ", line " << file_line_ << ", bad XML token '"
                << token.data() << "'";
     abort();
 }
diff --git a/source/lib/rocprofiler-sdk-roctx/roctx.cpp b/source/lib/rocprofiler-sdk-roctx/roctx.cpp
index 0147d351e7..765be8ccfd 100644
--- a/source/lib/rocprofiler-sdk-roctx/roctx.cpp
+++ b/source/lib/rocprofiler-sdk-roctx/roctx.cpp
@@ -211,7 +211,7 @@ get_table_impl()
                                                table_array.size(),
                                                &lib_id);
 
-    LOG(INFO) << "[rocprofiler-sdk-roctx][" << getpid() << "] rocprofiler-register returned code "
+    ROCP_INFO << "[rocprofiler-sdk-roctx][" << getpid() << "] rocprofiler-register returned code "
               << rocp_reg_status << ": " << rocprofiler_register_error_string(rocp_reg_status);
 
     LOG_IF(WARNING, rocp_reg_status != ROCP_REG_SUCCESS && rocp_reg_status != ROCP_REG_NO_TOOLS)
diff --git a/source/lib/rocprofiler-sdk-tool/config.cpp b/source/lib/rocprofiler-sdk-tool/config.cpp
index 1e1b7e67de..2d95319a83 100644
--- a/source/lib/rocprofiler-sdk-tool/config.cpp
+++ b/source/lib/rocprofiler-sdk-tool/config.cpp
@@ -153,12 +153,12 @@ parse_kernel_names(const std::string& line)
     {
         if(has_kernel_name_format(kernel_name))
         {
-            LOG(INFO) << "kernel name " << kernel_names.size() << ": " << kernel_name;
+            ROCP_INFO << "kernel name " << kernel_names.size() << ": " << kernel_name;
             kernel_names.emplace(kernel_name);
         }
         else
         {
-            LOG(ERROR) << "invalid kernel name: " << kernel_name;
+            ROCP_ERROR << "invalid kernel name: " << kernel_name;
         }
     }
 
diff --git a/source/lib/rocprofiler-sdk-tool/output_file.cpp b/source/lib/rocprofiler-sdk-tool/output_file.cpp
index 0b9843844d..9e69339654 100644
--- a/source/lib/rocprofiler-sdk-tool/output_file.cpp
+++ b/source/lib/rocprofiler-sdk-tool/output_file.cpp
@@ -22,6 +22,7 @@
 
 #include "output_file.hpp"
 #include "config.hpp"
+#include "lib/common/logging.hpp"
 
 #include <fmt/format.h>
 
@@ -57,7 +58,7 @@ get_output_stream(const std::string& fname, const std::string& ext)
     if(!_ofs && !*_ofs)
         throw std::runtime_error{fmt::format("Failed to open {} for output", output_file)};
 
-    LOG(ERROR) << "Opened result file: " << output_file;
+    ROCP_ERROR << "Opened result file: " << output_file;
 
     return {_ofs, [](std::ostream*& v) {
                 if(v) dynamic_cast<std::ofstream*>(v)->close();
@@ -69,9 +70,9 @@ get_output_stream(const std::string& fname, const std::string& ext)
 output_file::~output_file()
 {
     if(m_stream)
-        LOG(INFO) << "Closing result file: " << m_name;
+        ROCP_INFO << "Closing result file: " << m_name;
     else
-        LOG(WARNING) << "output_file::~output_file does not have a output stream instance!";
+        ROCP_WARNING << "output_file::~output_file does not have a output stream instance!";
 
     m_dtor(m_stream);
 }
diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp
index 99cdaa2a0a..013a0a1660 100644
--- a/source/lib/rocprofiler-sdk-tool/tool.cpp
+++ b/source/lib/rocprofiler-sdk-tool/tool.cpp
@@ -352,16 +352,16 @@ get_client_ctx()
 void
 flush()
 {
-    LOG(INFO) << "flushing buffers...";
+    ROCP_INFO << "flushing buffers...";
     for(auto itr : get_buffers().as_array())
     {
         if(itr.handle > 0)
         {
-            LOG(INFO) << "flushing buffer " << itr.handle;
+            ROCP_INFO << "flushing buffer " << itr.handle;
             ROCPROFILER_CALL(rocprofiler_flush_buffer(itr), "buffer flush");
         }
     }
-    LOG(INFO) << "Buffers flushed";
+    ROCP_INFO << "Buffers flushed";
 }
 
 void
@@ -633,7 +633,7 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/,
                           void* /*user_data*/,
                           uint64_t /*drop_count*/)
 {
-    LOG(INFO) << "Executing buffered tracing callback for " << num_headers << " headers";
+    ROCP_INFO << "Executing buffered tracing callback for " << num_headers << " headers";
 
     LOG_IF(ERROR, headers == nullptr)
         << "rocprofiler invoked a buffer callback with a null pointer to the array of headers. "
@@ -751,7 +751,7 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/,
             }
             else
             {
-                LOG(FATAL) << fmt::format(
+                ROCP_FATAL << fmt::format(
                     "unsupported category + kind: {} + {}", header->category, header->kind);
             }
         }
@@ -1242,11 +1242,11 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
         {
             auto cb_thread = rocprofiler_callback_thread_t{};
 
-            LOG(INFO) << "creating dedicated callback thread for buffer " << itr.handle;
+            ROCP_INFO << "creating dedicated callback thread for buffer " << itr.handle;
             ROCPROFILER_CALL(rocprofiler_create_callback_thread(&cb_thread),
                              "creating callback thread");
 
-            LOG(INFO) << "assigning buffer " << itr.handle << " to callback thread "
+            ROCP_INFO << "assigning buffer " << itr.handle << " to callback thread "
                       << cb_thread.handle;
             ROCPROFILER_CALL(rocprofiler_assign_callback_thread(itr, cb_thread),
                              "assigning callback thread");
@@ -1346,7 +1346,7 @@ rocprofiler_configure(uint32_t                 version,
             nullptr),
         "Iterate rocporfiler agents")
 
-    LOG(INFO) << id->name << " is using rocprofiler-sdk v" << major << "." << minor << "." << patch
+    ROCP_INFO << id->name << " is using rocprofiler-sdk v" << major << "." << minor << "." << patch
               << " (" << runtime_version << ")";
 
     // create configure data
diff --git a/source/lib/rocprofiler-sdk/agent.cpp b/source/lib/rocprofiler-sdk/agent.cpp
index 52d21cd377..1d666bf17a 100644
--- a/source/lib/rocprofiler-sdk/agent.cpp
+++ b/source/lib/rocprofiler-sdk/agent.cpp
@@ -56,7 +56,7 @@ namespace fs = rocprofiler::common::filesystem;
 
 #if defined(ROCPROFILER_CI)
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__)
-#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(FATAL)
+#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    ROCP_FATAL
 #else
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__)
 #    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(NON_CI_LEVEL)
@@ -200,7 +200,7 @@ parse_cpu_info()
             processor_info.emplace_back(info_v);
         else
         {
-            LOG(ERROR) << "Invalid processor info: "
+            ROCP_ERROR << "Invalid processor info: "
                        << fmt::format("processor={}, vendor={}, family={}, model={}, name={}, "
                                       "physical id={}, core id={}, apicid={}",
                                       info_v.processor,
@@ -326,7 +326,7 @@ read_property(const MapT& data, const std::string& label, Tp& value)
 
         if(data.find(label) == data.end())
         {
-            LOG(ERROR) << "agent properties map missing " << label << " entry";
+            ROCP_ERROR << "agent properties map missing " << label << " entry";
             return;
         }
 
@@ -413,7 +413,7 @@ read_topology()
             gpu_id_prop = read_file(node_path / "gpu_id");
         } catch(std::runtime_error& e)
         {
-            LOG(ERROR) << "Error reading '" << (node_path / "properties").string()
+            ROCP_ERROR << "Error reading '" << (node_path / "properties").string()
                        << "' :: " << e.what();
             continue;
         }
@@ -813,7 +813,7 @@ construct_agent_cache(::HsaApiTable* table)
         }
     }
 
-    LOG(ERROR) << "# agent node maps: " << hsa_agent_node_map.size();
+    ROCP_ERROR << "# agent node maps: " << hsa_agent_node_map.size();
 
     LOG_IF(FATAL, agent_map.size() != hsa_agents.size())
         << "rocprofiler was only able to map " << agent_map.size()
@@ -894,7 +894,7 @@ construct_agent_cache(::HsaApiTable* table)
             {
                 // TODO(aelwazir): To be changed back to use node id once ROCR fixes
                 // the hsa_agents to use the real node id
-                LOG(ERROR) << fmt::format("rocprofiler agent <-> HSA agent mapping failed: {} ({})",
+                ROCP_ERROR << fmt::format("rocprofiler agent <-> HSA agent mapping failed: {} ({})",
                                           rocp_agent->logical_node_id,
                                           err.what());
             }
@@ -905,10 +905,10 @@ construct_agent_cache(::HsaApiTable* table)
 std::optional<hsa_agent_t>
 get_hsa_agent(const rocprofiler_agent_t* agent)
 {
-    LOG(ERROR) << "# of agent mappings: " << get_agent_mapping().size();
+    ROCP_ERROR << "# of agent mappings: " << get_agent_mapping().size();
     for(const auto& itr : get_agent_mapping())
     {
-        LOG(ERROR) << "checking " << itr.rocp_agent->id.handle << " vs. " << agent->id.handle;
+        ROCP_ERROR << "checking " << itr.rocp_agent->id.handle << " vs. " << agent->id.handle;
         if(itr.rocp_agent->id.handle == agent->id.handle) return itr.hsa_agent;
     }
 
@@ -973,14 +973,14 @@ rocprofiler_query_available_agents(rocprofiler_agent_version_t             versi
     {
         if(agent_size > sizeof(rocprofiler_agent_v0_t))
         {
-            LOG(ERROR) << "size of rocprofiler agent struct used by caller is ABI-incompatible "
+            ROCP_ERROR << "size of rocprofiler agent struct used by caller is ABI-incompatible "
                           "with rocprofiler_agent_v0_t in rocprofiler";
             return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI;
         }
     }
     else
     {
-        LOG(FATAL) << "rocprofiler-sdk does not support given agent info version";
+        ROCP_FATAL << "rocprofiler-sdk does not support given agent info version";
     }
 
     auto&& pointers   = rocprofiler::agent::get_agents();
diff --git a/source/lib/rocprofiler-sdk/aql/helpers.cpp b/source/lib/rocprofiler-sdk/aql/helpers.cpp
index bf62d610e8..b5745d3260 100644
--- a/source/lib/rocprofiler-sdk/aql/helpers.cpp
+++ b/source/lib/rocprofiler-sdk/aql/helpers.cpp
@@ -43,7 +43,7 @@ get_query_info(hsa_agent_t agent, const counters::Metric& metric)
     if(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_ID, &query) !=
        HSA_STATUS_SUCCESS)
     {
-        DLOG(FATAL) << fmt::format("AQL failed to query info for counter {}", metric);
+        ROCP_DFATAL << fmt::format("AQL failed to query info for counter {}", metric);
         throw std::runtime_error(fmt::format("AQL failed to query info for counter {}", metric));
     }
     return query;
diff --git a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp
index 0aef048867..0d0f13f00b 100644
--- a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp
+++ b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp
@@ -78,11 +78,11 @@ findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set<std::st
     std::vector<counters::Metric> ret;
     auto                          all_counters = counters::getBaseHardwareMetrics();
 
-    LOG(ERROR) << "Looking up counters for " << std::string(agent.name());
+    ROCP_ERROR << "Looking up counters for " << std::string(agent.name());
     auto gfx_metrics = common::get_val(all_counters, std::string(agent.name()));
     if(!gfx_metrics)
     {
-        LOG(ERROR) << "No counters found for " << std::string(agent.name());
+        ROCP_ERROR << "No counters found for " << std::string(agent.name());
         return ret;
     }
 
diff --git a/source/lib/rocprofiler-sdk/buffer.cpp b/source/lib/rocprofiler-sdk/buffer.cpp
index 0fbe1e3fca..e6f81c5f5d 100644
--- a/source/lib/rocprofiler-sdk/buffer.cpp
+++ b/source/lib/rocprofiler-sdk/buffer.cpp
@@ -139,7 +139,7 @@ flush(rocprofiler_buffer_id_t buffer_id, bool wait)
 {
     if(registration::get_fini_status() > 0)
     {
-        LOG(ERROR) << "ignoring rocprofiler buffer flush (handle=" << buffer_id.handle
+        ROCP_ERROR << "ignoring rocprofiler buffer flush (handle=" << buffer_id.handle
                    << ") request after finalization";
         return ROCPROFILER_STATUS_ERROR_FINALIZED;
     }
@@ -202,14 +202,14 @@ flush(rocprofiler_buffer_id_t buffer_id, bool wait)
                 }
             } catch(std::exception& e)
             {
-                LOG(ERROR) << "buffer callback threw an exception: " << e.what();
+                ROCP_ERROR << "buffer callback threw an exception: " << e.what();
             }
             // clear the buffer
             buff_internal_v.clear();
         }
         else
         {
-            LOG(INFO) << "buffer at " << buffer_id.handle << " is empty...";
+            ROCP_INFO << "buffer at " << buffer_id.handle << " is empty...";
         }
 
         buff_v->syncer.clear();
@@ -242,7 +242,7 @@ rocprofiler_create_buffer(rocprofiler_context_id_t        context,
     auto* existing_buff = rocprofiler::buffer::get_buffer(*buffer_id);
     if(existing_buff)
     {
-        LOG(ERROR) << "buffer (handle=" << buffer_id->handle
+        ROCP_ERROR << "buffer (handle=" << buffer_id->handle
                    << ") already allocated: handle=" << existing_buff->buffer_id;
         return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED;
     }
diff --git a/source/lib/rocprofiler-sdk/context/context.cpp b/source/lib/rocprofiler-sdk/context/context.cpp
index 842a4ae013..8c55cee2c2 100644
--- a/source/lib/rocprofiler-sdk/context/context.cpp
+++ b/source/lib/rocprofiler-sdk/context/context.cpp
@@ -212,13 +212,13 @@ pop_latest_correlation_id(correlation_id* val)
 {
     if(!val)
     {
-        LOG(ERROR) << "passed nullptr to correlation id";
+        ROCP_ERROR << "passed nullptr to correlation id";
         return nullptr;
     }
 
     if(get_latest_correlation_id_impl().empty())
     {
-        LOG(ERROR) << "empty thread-local correlation id stack";
+        ROCP_ERROR << "empty thread-local correlation id stack";
         return nullptr;
     }
 
diff --git a/source/lib/rocprofiler-sdk/counters/core.cpp b/source/lib/rocprofiler-sdk/counters/core.cpp
index 912fc0c05e..db04529377 100644
--- a/source/lib/rocprofiler-sdk/counters/core.cpp
+++ b/source/lib/rocprofiler-sdk/counters/core.cpp
@@ -173,7 +173,7 @@ counter_callback_info::setup_profile_config(const hsa::AgentCache&           age
 
         if(!req_counters)
         {
-            LOG(ERROR) << fmt::format("Could not find counter {}", metric.name());
+            ROCP_ERROR << fmt::format("Could not find counter {}", metric.name());
             return ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND;
         }
 
@@ -195,14 +195,14 @@ counter_callback_info::setup_profile_config(const hsa::AgentCache&           age
         const auto* agent_map = rocprofiler::common::get_val(asts, agent_name);
         if(!agent_map)
         {
-            LOG(ERROR) << fmt::format("Coult not build AST for {}", agent_name);
+            ROCP_ERROR << fmt::format("Coult not build AST for {}", agent_name);
             return ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED;
         }
 
         const auto* counter_ast = rocprofiler::common::get_val(*agent_map, metric.name());
         if(!counter_ast)
         {
-            LOG(ERROR) << fmt::format("Coult not find AST for {}", metric.name());
+            ROCP_ERROR << fmt::format("Coult not find AST for {}", metric.name());
             return ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND;
         }
         config.asts.push_back(*counter_ast);
@@ -212,7 +212,7 @@ counter_callback_info::setup_profile_config(const hsa::AgentCache&           age
             config.asts.back().set_dimensions();
         } catch(std::runtime_error& e)
         {
-            LOG(ERROR) << metric.name() << " has improper dimensions"
+            ROCP_ERROR << metric.name() << " has improper dimensions"
                        << " " << e.what();
             return ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND;
         }
diff --git a/source/lib/rocprofiler-sdk/counters/dimensions.cpp b/source/lib/rocprofiler-sdk/counters/dimensions.cpp
index db7ce834da..7421200222 100644
--- a/source/lib/rocprofiler-sdk/counters/dimensions.cpp
+++ b/source/lib/rocprofiler-sdk/counters/dimensions.cpp
@@ -80,7 +80,7 @@ getBlockDimensions(std::string_view agent, const Metric& metric)
                     }
                     else
                     {
-                        LOG(ERROR) << "Unknown AQL Profiler Dimension " << id << " " << extent;
+                        ROCP_ERROR << "Unknown AQL Profiler Dimension " << id << " " << extent;
                     }
                 }
             }
@@ -125,7 +125,7 @@ get_dimension_cache()
                             dims.emplace(ast.out_id().handle, ast_copy.set_dimensions());
                         } catch(std::runtime_error& e)
                         {
-                            LOG(ERROR) << metric << " has improper dimensions"
+                            ROCP_ERROR << metric << " has improper dimensions"
                                        << " " << e.what();
                             throw;
                         }
diff --git a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp
index c1a5768f25..8178dd5693 100644
--- a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp
+++ b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp
@@ -140,7 +140,7 @@ get_ast_map()
                 yyparse(&ast);
                 if(!ast)
                 {
-                    LOG(ERROR) << fmt::format("Unable to parse metric {}", metric);
+                    ROCP_ERROR << fmt::format("Unable to parse metric {}", metric);
                     throw std::runtime_error(fmt::format("Unable to parse metric {}", metric));
                 }
                 try
@@ -155,7 +155,7 @@ get_ast_map()
                                    // logic as a Finish() method
                 } catch(std::exception& e)
                 {
-                    LOG(ERROR) << e.what();
+                    ROCP_ERROR << e.what();
                     throw std::runtime_error(
                         fmt::format("AST was not generated for {}:{}", gfx, metric.name()));
                 }
diff --git a/source/lib/rocprofiler-sdk/counters/metrics.cpp b/source/lib/rocprofiler-sdk/counters/metrics.cpp
index 680a3d53b3..7c4bd43be8 100644
--- a/source/lib/rocprofiler-sdk/counters/metrics.cpp
+++ b/source/lib/rocprofiler-sdk/counters/metrics.cpp
@@ -90,7 +90,7 @@ MetricMap
 loadXml(const std::string& filename, bool load_constants = false)
 {
     MetricMap ret;
-    DLOG(INFO) << "Loading Counter Config: " << filename;
+    ROCP_INFO << "Loading Counter Config: " << filename;
     // todo: return unique_ptr....
     auto xml = common::Xml::Create(filename);
     LOG_IF(FATAL, !xml)
@@ -142,7 +142,7 @@ std::string
 findViaInstallPath(const std::string& filename)
 {
     Dl_info dl_info = {};
-    DLOG(INFO) << filename << " is being looked up via install path";
+    ROCP_INFO << filename << " is being looked up via install path";
     if(dladdr(reinterpret_cast<const void*>(rocprofiler_query_available_agents), &dl_info) != 0)
     {
         return common::filesystem::path{dl_info.dli_fname}.parent_path().parent_path() /
@@ -156,7 +156,7 @@ findViaEnvironment(const std::string& filename)
 {
     if(const char* metrics_path = nullptr; (metrics_path = getenv("ROCPROFILER_METRICS_PATH")))
     {
-        DLOG(INFO) << filename << " is being looked up via env variable ROCPROFILER_METRICS_PATH";
+        ROCP_INFO << filename << " is being looked up via env variable ROCPROFILER_METRICS_PATH";
         return common::filesystem::path{std::string{metrics_path}} / filename;
     }
     // No environment variable, lookup via install path
diff --git a/source/lib/rocprofiler-sdk/counters/parser/parser.cpp b/source/lib/rocprofiler-sdk/counters/parser/parser.cpp
index 6c36e1b69d..c4038cb257 100644
--- a/source/lib/rocprofiler-sdk/counters/parser/parser.cpp
+++ b/source/lib/rocprofiler-sdk/counters/parser/parser.cpp
@@ -81,7 +81,7 @@ yylex(void);
 void
 yyerror(rocprofiler::counters::RawAST**, const char* s)
 {
-    LOG(ERROR) << s;
+    ROCP_ERROR << s;
 }
 
 #line 85 "parser.cpp"
diff --git a/source/lib/rocprofiler-sdk/counters/parser/parser.y b/source/lib/rocprofiler-sdk/counters/parser/parser.y
index 12d24274e7..2cf483cdef 100644
--- a/source/lib/rocprofiler-sdk/counters/parser/parser.y
+++ b/source/lib/rocprofiler-sdk/counters/parser/parser.y
@@ -16,7 +16,7 @@ using namespace rocprofiler::counters;
 
 int yyparse(rocprofiler::counters::RawAST** result);
 int yylex(void);
-void yyerror(rocprofiler::counters::RawAST**, const char *s) { LOG(ERROR) << s; }
+void yyerror(rocprofiler::counters::RawAST**, const char *s) { ROCP_ERROR << s; }
 %}
 
 /* declare tokens */
diff --git a/source/lib/rocprofiler-sdk/counters/parser/raw_ast.hpp b/source/lib/rocprofiler-sdk/counters/parser/raw_ast.hpp
index f1e24d69b8..3bfb67241d 100644
--- a/source/lib/rocprofiler-sdk/counters/parser/raw_ast.hpp
+++ b/source/lib/rocprofiler-sdk/counters/parser/raw_ast.hpp
@@ -194,7 +194,7 @@ struct RawAST
         }
         else
         {
-            LOG(ERROR) << "select_dimension_set creation failed.";
+            ROCP_ERROR << "select_dimension_set creation failed.";
         }
     }
 
diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp
index d5ac27d88e..45f98d678d 100644
--- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp
+++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp
@@ -113,11 +113,11 @@ findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set<std::st
     std::vector<counters::Metric> ret;
     auto                          all_counters = counters::getMetricMap();
 
-    LOG(ERROR) << "Looking up counters for " << std::string(agent.name());
+    ROCP_ERROR << "Looking up counters for " << std::string(agent.name());
     auto gfx_metrics = common::get_val(*all_counters, std::string(agent.name()));
     if(!gfx_metrics)
     {
-        LOG(ERROR) << "No counters found for " << std::string(agent.name());
+        ROCP_ERROR << "No counters found for " << std::string(agent.name());
         return ret;
     }
 
diff --git a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
index 647bf1ce47..da57745df2 100644
--- a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
+++ b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
@@ -187,11 +187,11 @@ findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set<std::st
     std::vector<counters::Metric> ret;
     auto                          all_counters = counters::getMetricMap();
 
-    LOG(ERROR) << "Looking up counters for " << std::string(agent.name());
+    ROCP_ERROR << "Looking up counters for " << std::string(agent.name());
     auto gfx_metrics = common::get_val(*all_counters, std::string(agent.name()));
     if(!gfx_metrics)
     {
-        LOG(ERROR) << "No counters found for " << std::string(agent.name());
+        ROCP_ERROR << "No counters found for " << std::string(agent.name());
         return ret;
     }
 
@@ -239,7 +239,7 @@ TEST(dimension, block_dim_test)
              */
             std::unordered_map<counters::rocprofiler_profile_counter_instance_types, uint64_t>
                 rocp_dims;
-            LOG(ERROR) << metric.name() << " " << metric.special();
+            ROCP_ERROR << metric.name() << " " << metric.special();
             if(!metric.special().empty())
             {
                 rocp_dims[counters::rocprofiler_profile_counter_instance_types::
diff --git a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp
index 32c0e60fec..089897f5c3 100644
--- a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp
+++ b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp
@@ -219,7 +219,7 @@ add_constants(std::unordered_map<std::string, Metric>& metrics, uint64_t start_i
     rocprofiler::agent::get_agents();
     for(const auto& prop : rocprofiler::agent::get_agent_available_properties())
     {
-        LOG(ERROR) << prop;
+        ROCP_ERROR << prop;
         metrics[prop] = {"constant",
                          prop,
                          "",
@@ -299,7 +299,7 @@ TEST(evaluate_ast, counter_constants)
     {
         auto eval_counters =
             rocprofiler::counters::get_required_hardware_counters(asts, "gfx9", metrics[name]);
-        LOG(INFO) << name;
+        ROCP_INFO << name;
         ASSERT_TRUE(eval_counters);
         EXPECT_EQ(eval_counters->size(), expected.size());
         for(const auto& c : *eval_counters)
@@ -505,7 +505,7 @@ TEST(evaluate_ast, evaluate_simple_counters)
 
     for(auto& [name, expected, eval_count] : derived_counters)
     {
-        LOG(INFO) << name;
+        ROCP_INFO << name;
         auto eval_counters =
             rocprofiler::counters::get_required_hardware_counters(asts, "gfx9", metrics[name]);
         ASSERT_TRUE(eval_counters);
@@ -558,7 +558,7 @@ run_reduce_test(
 
     for(auto& [name, expected, eval_count] : derived_counters)
     {
-        LOG(INFO) << name;
+        ROCP_INFO << name;
         auto eval_counters =
             rocprofiler::counters::get_required_hardware_counters(asts, "gfx9", metrics[name]);
         ASSERT_TRUE(eval_counters);
@@ -1001,7 +1001,7 @@ TEST(evaluate_ast, evaluate_mixed_counters)
 
     for(auto& [name, expected, eval_count] : derived_counters)
     {
-        LOG(INFO) << name;
+        ROCP_INFO << name;
         auto eval_counters =
             rocprofiler::counters::get_required_hardware_counters(asts, "gfx9", metrics[name]);
         ASSERT_TRUE(eval_counters);
diff --git a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp
index 4b842bb5d5..b38c37e346 100644
--- a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp
+++ b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp
@@ -74,7 +74,7 @@ TEST(metrics, base_load)
     auto find = [&rocp_data_v](const auto& v) -> std::optional<counters::Metric> {
         for(const auto& ditr : rocp_data_v)
         {
-            LOG(ERROR) << fmt::format("{}", ditr);
+            ROCP_ERROR << fmt::format("{}", ditr);
             if(ditr.name() == v.name()) return ditr;
         }
         return std::nullopt;
diff --git a/source/lib/rocprofiler-sdk/hip/hip.cpp b/source/lib/rocprofiler-sdk/hip/hip.cpp
index ab61007167..6f370ff25d 100644
--- a/source/lib/rocprofiler-sdk/hip/hip.cpp
+++ b/source/lib/rocprofiler-sdk/hip/hip.cpp
@@ -173,7 +173,7 @@ hip_api_impl<TableIdx, OpIdx>::exec(FuncT&& _func, Args&&... args)
     }
 
     using info_type = hip_api_info<TableIdx, OpIdx>;
-    LOG(ERROR) << "nullptr to next hip function for " << info_type::name << " ("
+    ROCP_ERROR << "nullptr to next hip function for " << info_type::name << " ("
                << info_type::operation_idx << ")";
 
     return get_default_retval<return_type>();
@@ -533,12 +533,12 @@ copy_table(Tp* _orig, uint64_t _tbl_instance, std::integral_constant<size_t, OpI
 
         if(!_copy_func)
         {
-            LOG(INFO) << "copying table entry for " << _info.name;
+            ROCP_INFO << "copying table entry for " << _info.name;
             _copy_func = _orig_func;
         }
         else
         {
-            LOG(INFO) << "skipping copying table entry for " << _info.name
+            ROCP_INFO << "skipping copying table entry for " << _info.name
                       << " from table instance " << _tbl_instance;
         }
     }
@@ -562,7 +562,7 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
                _info.callback_domain_idx, _info.buffered_domain_idx, _info.operation_idx))
             return;
 
-        LOG(INFO) << "updating table entry for " << _info.name;
+        ROCP_INFO << "updating table entry for " << _info.name;
 
         // 1. get the sub-table containing the function pointer in original table
         // 2. get reference to function pointer in sub-table in original table
diff --git a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp
index fc6cd694a2..f8dfb8f489 100644
--- a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp
@@ -50,7 +50,7 @@
 
 #if defined(ROCPROFILER_CI)
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__)
-#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(FATAL)
+#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    ROCP_FATAL
 #else
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__)
 #    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(NON_CI_LEVEL)
@@ -491,7 +491,7 @@ async_copy_impl(Args... args)
 
         if(_status != HSA_STATUS_SUCCESS)
         {
-            LOG(ERROR) << "hsa_signal_create returned non-zero error code " << _status;
+            ROCP_ERROR << "hsa_signal_create returned non-zero error code " << _status;
 
             delete _data;
             return invoke(get_next_dispatch<TableIdx, OpIdx>(),
@@ -513,7 +513,7 @@ async_copy_impl(Args... args)
 
         if(_status != HSA_STATUS_SUCCESS)
         {
-            LOG(ERROR) << "hsa_amd_signal_async_handler returned non-zero error code " << _status;
+            ROCP_ERROR << "hsa_amd_signal_async_handler returned non-zero error code " << _status;
 
             ROCP_HSA_TABLE_CALL(ERROR, get_core_table()->hsa_signal_destroy_fn(_data->rocp_signal))
                 << ":: failed to destroy signal after async handler failed";
@@ -576,12 +576,12 @@ async_copy_save(hsa_amd_ext_table_t* _orig, uint64_t _tbl_instance)
 
     if(!_copy_func)
     {
-        LOG(INFO) << "copying table entry for " << _meta.name;
+        ROCP_INFO << "copying table entry for " << _meta.name;
         _copy_func = _orig_func;
     }
     else
     {
-        LOG(INFO) << "skipping copying table entry for " << _meta.name << " from table instance "
+        ROCP_INFO << "skipping copying table entry for " << _meta.name << " from table instance "
                   << _tbl_instance;
     }
 }
diff --git a/source/lib/rocprofiler-sdk/hsa/code_object.cpp b/source/lib/rocprofiler-sdk/hsa/code_object.cpp
index c36cacfa49..bc076c8352 100644
--- a/source/lib/rocprofiler-sdk/hsa/code_object.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/code_object.cpp
@@ -773,8 +773,8 @@ code_object_load_callback(hsa_executable_t         executable,
     const auto* _rocp_agent = agent::get_rocprofiler_agent(data.hsa_agent);
     if(!_rocp_agent)
     {
-        ROCP_CI_LOG(ERROR) << "hsa agent (handle=" << _hsa_agent.handle
-                           << ") did not map to a rocprofiler agent";
+        ROCP_ERROR << "hsa agent (handle=" << _hsa_agent.handle
+                   << ") did not map to a rocprofiler agent";
         return HSA_STATUS_ERROR_INVALID_AGENT;
     }
     data.rocp_agent = _rocp_agent->id;
@@ -792,7 +792,7 @@ code_object_load_callback(hsa_executable_t         executable,
     }
     else
     {
-        LOG(ERROR) << "hsa_executable_iterate_agent_symbols failed for " << data.uri;
+        ROCP_ERROR << "hsa_executable_iterate_agent_symbols failed for " << data.uri;
     }
 
     return _status;
@@ -814,13 +814,13 @@ code_object_unload_callback(hsa_executable_t         executable,
     CHECK_NOTNULL(code_obj_arr);
 
     // auto _size = get_code_objects().rlock([](const auto& data) { return data.size(); });
-    // LOG(INFO) << "[inp] executable=" << executable.handle
+    // ROCP_INFO << "[inp] executable=" << executable.handle
     //            << ", code_object=" << loaded_code_object.handle << " vs. " << _size;
 
     get_code_objects().rlock([&](const code_object_array_t& arr) {
         for(const auto& itr : arr)
         {
-            // LOG(INFO) << "[cmp] executable=" << itr->hsa_executable.handle
+            // ROCP_INFO << "[cmp] executable=" << itr->hsa_executable.handle
             //            << ", code_object=" << itr->hsa_code_object.handle;
             if(itr->hsa_executable.handle == executable.handle &&
                itr->hsa_code_object.handle == loaded_code_object.handle)
@@ -858,7 +858,7 @@ executable_freeze(hsa_executable_t executable, const char* options)
     hsa_status_t status = CHECK_NOTNULL(get_freeze_function())(executable, options);
     if(status != HSA_STATUS_SUCCESS) return status;
 
-    LOG(INFO) << "running " << __FUNCTION__ << " (executable=" << executable.handle << ")...";
+    ROCP_INFO << "running " << __FUNCTION__ << " (executable=" << executable.handle << ")...";
 
     get_executables().wlock(
         [executable](executable_array_t& data) { data.emplace_back(executable); });
@@ -999,7 +999,7 @@ executable_destroy(hsa_executable_t executable)
 std::vector<code_object_unload>
 shutdown(hsa_executable_t executable)
 {
-    LOG(INFO) << "running " << __FUNCTION__ << " (executable=" << executable.handle << ")...";
+    ROCP_INFO << "running " << __FUNCTION__ << " (executable=" << executable.handle << ")...";
 
     auto _unloaded = std::vector<code_object_unload>{};
     hsa::get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
diff --git a/source/lib/rocprofiler-sdk/hsa/hsa.cpp b/source/lib/rocprofiler-sdk/hsa/hsa.cpp
index 1b822dbe80..c60fe6e0dc 100644
--- a/source/lib/rocprofiler-sdk/hsa/hsa.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/hsa.cpp
@@ -605,12 +605,12 @@ copy_table(Tp* _orig, uint64_t _tbl_instance, std::integral_constant<size_t, OpI
 
         if(!_copy_func)
         {
-            LOG(INFO) << "copying table entry for " << _info.name;
+            ROCP_INFO << "copying table entry for " << _info.name;
             _copy_func = _orig_func;
         }
         else
         {
-            LOG(INFO) << "skipping copying table entry for " << _info.name
+            ROCP_INFO << "skipping copying table entry for " << _info.name
                       << " from table instance " << _tbl_instance;
         }
     }
@@ -639,7 +639,7 @@ update_table(const context::context_array_t& _contexts,
                                 _info.operation_idx))
             return;
 
-        LOG(INFO) << "updating table entry for " << _info.name;
+        ROCP_INFO << "updating table entry for " << _info.name;
 
         // 1. get the sub-table containing the function pointer in original table
         // 2. get reference to function pointer in sub-table in original table
diff --git a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp
index a359141445..3de9de9e89 100644
--- a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp
@@ -85,7 +85,7 @@ hsa_barrier::enqueue_packet(const Queue* queue)
     rocprofiler_packet barrier{};
     barrier.barrier_and.header        = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
     barrier.barrier_and.dep_signal[0] = _barrier_signal;
-    LOG(ERROR) << "Barrier Added: " << _barrier_signal.handle;
+    ROCP_ERROR << "Barrier Added: " << _barrier_signal.handle;
     return barrier;
 }
 
diff --git a/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp
index db9684e272..41110c460e 100644
--- a/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp
@@ -72,7 +72,7 @@ profiler_serializer::add_queue(hsa_queue_t** hsa_queues, const Queue& queue)
                                              -1,
                                              profiler_serializer_ready_signal_handler,
                                              *hsa_queues);
-    if(status != HSA_STATUS_SUCCESS) LOG(FATAL) << "hsa_amd_signal_async_handler failed";
+    if(status != HSA_STATUS_SUCCESS) ROCP_FATAL << "hsa_amd_signal_async_handler failed";
 }
 
 void
@@ -124,25 +124,25 @@ void
 profiler_serializer::queue_ready(hsa_queue_t* hsa_queue, const Queue& queue)
 {
     {
-        LOG(INFO) << "Obtaining queue mutex lock...";
+        ROCP_TRACE << "Obtaining queue mutex lock...";
         std::lock_guard<std::mutex> cv_lock(queue.cv_mutex);
-        LOG(INFO) << "Queue mutex lock obtained";
+        ROCP_TRACE << "Queue mutex lock obtained";
         if(queue.get_state() == queue_state::to_destroy)
         {
-            LOG(INFO) << "Setting queue state to done_destroy...";
+            ROCP_TRACE << "Setting queue state to done_destroy...";
             CHECK_NOTNULL(get_queue_controller())
                 ->set_queue_state(queue_state::done_destroy, hsa_queue);
-            LOG(INFO) << "Destroying ready signal...";
+            ROCP_TRACE << "Destroying ready signal...";
             CHECK_NOTNULL(get_queue_controller())
                 ->get_core_table()
                 .hsa_signal_destroy_fn(queue.ready_signal);
-            LOG(INFO) << "Notifying queue condition variable...";
+            ROCP_TRACE << "Notifying queue condition variable...";
             queue.cv_ready_signal.notify_one();
             return;
         }
     }
 
-    LOG(INFO) << "setting queue ready signal to 1...";
+    ROCP_TRACE << "setting queue ready signal to 1...";
     CHECK_NOTNULL(get_queue_controller())
         ->get_core_table()
         .hsa_signal_store_screlease_fn(queue.ready_signal, 1);
@@ -204,7 +204,7 @@ profiler_serializer::kernel_dispatch(const Queue& queue) const
 void
 profiler_serializer::destroy_queue(hsa_queue_t* id, const Queue& queue)
 {
-    LOG(INFO) << "destroying queue...";
+    ROCP_INFO << "destroying queue...";
 
     /*Deletes the queue to be destructed from the dispatch ready.*/
     for(auto& barriers : _barrier)
@@ -224,7 +224,7 @@ profiler_serializer::destroy_queue(hsa_queue_t* id, const Queue& queue)
                     {
                         // insert fatal condition here
                         // ToDO [srnagara]: Need to find a solution rather than abort.
-                        LOG(FATAL)
+                        ROCP_FATAL
                             << "Queue is being destroyed while kernel launch is still active";
                     }
                     return true;
@@ -237,7 +237,7 @@ profiler_serializer::destroy_queue(hsa_queue_t* id, const Queue& queue)
         ->get_core_table()
         .hsa_signal_store_screlease_fn(queue.ready_signal, 0);
 
-    LOG(INFO) << "queue destroyed";
+    ROCP_INFO << "queue destroyed";
 }
 
 // Enable the serializer
@@ -246,7 +246,7 @@ profiler_serializer::enable(const queue_map_t& queues)
 {
     if(_serializer_status == Status::ENABLED) return;
 
-    LOG(INFO) << "Enabling profiler serialization...";
+    ROCP_INFO << "Enabling profiler serialization...";
 
     _serializer_status = Status::ENABLED;
     if(queues.empty()) return;
@@ -259,7 +259,7 @@ profiler_serializer::enable(const queue_map_t& queues)
     _serializer_status = Status::ENABLED;
     _barrier.back().barrier->set_barrier(queues);
 
-    LOG(INFO) << "Profiler serialization enabled";
+    ROCP_INFO << "Profiler serialization enabled";
 }
 
 // Disable the serializer
@@ -268,7 +268,7 @@ profiler_serializer::disable(const queue_map_t& queues)
 {
     if(_serializer_status == Status::DISABLED) return;
 
-    LOG(INFO) << "Disabling profiler serialization...";
+    ROCP_INFO << "Disabling profiler serialization...";
 
     _serializer_status = Status::DISABLED;
     if(queues.empty()) return;
@@ -281,7 +281,7 @@ profiler_serializer::disable(const queue_map_t& queues)
     _serializer_status = Status::DISABLED;
     _barrier.back().barrier->set_barrier(queues);
 
-    LOG(INFO) << "Profiler serialization disabled";
+    ROCP_INFO << "Profiler serialization disabled";
 }
 
 }  // namespace hsa
diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp
index 190311f7d9..5af1f9d225 100644
--- a/source/lib/rocprofiler-sdk/hsa/queue.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp
@@ -56,7 +56,7 @@ static_assert(offsetof(hsa_ext_amd_aql_pm4_packet_t, completion_signal) ==
 
 #if defined(ROCPROFILER_CI)
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__)
-#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(FATAL)
+#    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    ROCP_FATAL
 #else
 #    define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__)
 #    define ROCP_CI_LOG(NON_CI_LEVEL, ...)    LOG(NON_CI_LEVEL)
diff --git a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
index e6e45e52fb..51fc03c4e2 100644
--- a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
+++ b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
@@ -71,7 +71,7 @@ create_queue(hsa_agent_t        agent,
             return HSA_STATUS_SUCCESS;
         }
     }
-    LOG(FATAL) << "Could not find agent - " << agent.handle;
+    ROCP_FATAL << "Could not find agent - " << agent.handle;
     return HSA_STATUS_ERROR_FATAL;
 }
 
@@ -117,13 +117,13 @@ QueueController::destroy_queue(hsa_queue_t* id)
     // return if queue does not exist
     if(!queue) return;
 
-    LOG(INFO) << "destroying queue...";
+    ROCP_INFO << "destroying queue...";
 
     queue->sync();
     if(queue->block_signal.handle != 0) get_core_table().hsa_signal_destroy_fn(queue->block_signal);
     _queues.wlock([&](auto& map) { map.erase(id); });
 
-    LOG(INFO) << "queue destroyed";
+    ROCP_INFO << "queue destroyed";
 }
 
 ClientID
@@ -262,7 +262,7 @@ QueueController::print_debug_signals() const
     _debug_signals.rlock([&](const auto& signals) {
         for(const auto& [id, signal] : signals)
         {
-            LOG(ERROR) << "Signal " << signal.handle << " "
+            ROCP_ERROR << "Signal " << signal.handle << " "
                        << get_core_table().hsa_signal_load_scacquire_fn(signal);
         }
     });
@@ -271,7 +271,7 @@ QueueController::print_debug_signals() const
     _queues.rlock([&](const auto& queues) {
         for(const auto& [_, queue] : queues)
         {
-            LOG(ERROR) << "Queue " << queue->get_id().handle << " " << queue->ready_signal.handle
+            ROCP_ERROR << "Queue " << queue->get_id().handle << " " << queue->ready_signal.handle
                        << ":" << get_core_table().hsa_signal_load_scacquire_fn(queue->ready_signal)
                        << " " << queue->block_signal.handle << ":"
                        << get_core_table().hsa_signal_load_scacquire_fn(queue->block_signal);
diff --git a/source/lib/rocprofiler-sdk/marker/marker.cpp b/source/lib/rocprofiler-sdk/marker/marker.cpp
index 3de208cb87..63e36e44a2 100644
--- a/source/lib/rocprofiler-sdk/marker/marker.cpp
+++ b/source/lib/rocprofiler-sdk/marker/marker.cpp
@@ -121,7 +121,7 @@ roctx_api_impl<TableIdx, OpIdx>::exec(FuncT&& _func, Args&&... args)
     }
 
     using info_type = roctx_api_info<TableIdx, OpIdx>;
-    LOG(ERROR) << "nullptr to next roctx function for " << info_type::name << " ("
+    ROCP_ERROR << "nullptr to next roctx function for " << info_type::name << " ("
                << info_type::operation_idx << ")";
 
     if constexpr(std::is_void<return_type>::value)
@@ -484,12 +484,12 @@ copy_table(Tp* _orig, uint64_t _tbl_instance, std::integral_constant<size_t, OpI
 
         if(!_copy_func)
         {
-            LOG(INFO) << "copying table entry for " << _info.name;
+            ROCP_INFO << "copying table entry for " << _info.name;
             _copy_func = _orig_func;
         }
         else
         {
-            LOG(INFO) << "skipping copying table entry for " << _info.name
+            ROCP_INFO << "skipping copying table entry for " << _info.name
                       << " from table instance " << _tbl_instance;
         }
     }
@@ -514,7 +514,7 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
                _info.callback_domain_idx, _info.buffered_domain_idx, _info.operation_idx))
             return;
 
-        LOG(INFO) << "updating table entry for " << _info.name;
+        ROCP_INFO << "updating table entry for " << _info.name;
 
         // 1. get the sub-table containing the function pointer in original table
         // 2. get reference to function pointer in sub-table in original table
diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp
index dc165d03ea..69ee673f97 100644
--- a/source/lib/rocprofiler-sdk/registration.cpp
+++ b/source/lib/rocprofiler-sdk/registration.cpp
@@ -198,7 +198,7 @@ find_clients()
 
     if(get_forced_configure() && is_unique_configure_func(get_forced_configure()))
     {
-        LOG(ERROR) << "adding forced configure";
+        ROCP_ERROR << "adding forced configure";
         emplace_client("(forced)", nullptr, get_forced_configure());
     }
 
@@ -235,7 +235,7 @@ find_clients()
     {
         for(const auto& itr : env)
         {
-            LOG(INFO) << "[env] searching " << itr << " for rocprofiler_configure";
+            ROCP_INFO << "[env] searching " << itr << " for rocprofiler_configure";
 
             void* handle = dlopen(itr.c_str(), RTLD_NOLOAD | RTLD_LAZY);
 
@@ -248,7 +248,7 @@ find_clients()
 
             if(!handle)
             {
-                LOG(ERROR) << "error dlopening " << itr;
+                ROCP_ERROR << "error dlopening " << itr;
                 continue;
             }
 
@@ -290,7 +290,7 @@ find_clients()
     {
         for(const auto& itr : get_link_map())
         {
-            LOG(INFO) << "searching " << itr << " for rocprofiler_configure";
+            ROCP_INFO << "searching " << itr << " for rocprofiler_configure";
 
             void* handle = dlopen(itr.c_str(), RTLD_LAZY | RTLD_NOLOAD);
             LOG_IF(ERROR, handle == nullptr) << "error dlopening " << itr;
@@ -300,7 +300,7 @@ find_clients()
             // symbol not found
             if(!_sym)
             {
-                LOG(INFO) << "|_" << itr << " did not contain rocprofiler_configure symbol";
+                ROCP_INFO << "|_" << itr << " did not contain rocprofiler_configure symbol";
                 continue;
             }
 
@@ -327,7 +327,7 @@ find_clients()
         }
     }
 
-    LOG(ERROR) << __FUNCTION__ << " found " << data.size() << " clients";
+    ROCP_ERROR << __FUNCTION__ << " found " << data.size() << " clients";
 
     return data;
 }
@@ -356,7 +356,7 @@ invoke_client_configures()
 
     auto _lk = scoped_lock_t{get_registration_mutex()};
 
-    LOG(ERROR) << __FUNCTION__;
+    ROCP_ERROR << __FUNCTION__;
 
     if(!get_clients()) return false;
 
@@ -366,7 +366,7 @@ invoke_client_configures()
 
         if(!itr->configure_func)
         {
-            LOG(ERROR) << "rocprofiler::registration::invoke_client_configures() attempted to "
+            ROCP_ERROR << "rocprofiler::registration::invoke_client_configures() attempted to "
                           "invoke configure function from "
                        << itr->name << " that had no configuration function";
             continue;
@@ -374,7 +374,7 @@ invoke_client_configures()
 
         if(get_invoked_configures().find(itr->configure_func) != get_invoked_configures().end())
         {
-            LOG(ERROR) << "rocprofiler::registration::invoke_client_configures() attempted to "
+            ROCP_ERROR << "rocprofiler::registration::invoke_client_configures() attempted to "
                           "invoke configure function from "
                        << itr->name << " (addr="
                        << fmt::format("{:#018x}", reinterpret_cast<uint64_t>(itr->configure_func))
@@ -383,7 +383,7 @@ invoke_client_configures()
         }
         else
         {
-            LOG(INFO) << "rocprofiler::registration::invoke_client_configures() invoking configure "
+            ROCP_INFO << "rocprofiler::registration::invoke_client_configures() invoking configure "
                          "function from "
                       << itr->name << " (addr="
                       << fmt::format("{:#018x}", reinterpret_cast<uint64_t>(itr->configure_func))
@@ -418,7 +418,7 @@ invoke_client_initializers()
 
     auto _lk = scoped_lock_t{get_registration_mutex()};
 
-    LOG(ERROR) << __FUNCTION__;
+    ROCP_ERROR << __FUNCTION__;
 
     if(!get_clients()) return false;
 
@@ -460,7 +460,7 @@ invoke_client_finalizers()
 void
 invoke_client_finalizer(rocprofiler_client_id_t client_id)
 {
-    LOG(ERROR) << __FUNCTION__ << "(client_id=" << client_id.handle << ")";
+    ROCP_ERROR << __FUNCTION__ << "(client_id=" << client_id.handle << ")";
 
     auto _lk = scoped_lock_t{get_registration_mutex()};
 
@@ -539,17 +539,17 @@ set_fini_status(int v)
 void
 initialize()
 {
-    LOG(INFO) << "rocprofiler initialize called...";
+    ROCP_INFO << "rocprofiler initialize called...";
 
     if(get_init_status() != 0)
     {
-        LOG(INFO) << "rocprofiler initialize ignored...";
+        ROCP_INFO << "rocprofiler initialize ignored...";
         return;
     }
 
     static auto _once = std::once_flag{};
     std::call_once(_once, []() {
-        LOG(INFO) << "rocprofiler initialize started...";
+        ROCP_INFO << "rocprofiler initialize started...";
         // initialization is in process
         set_init_status(-1);
         std::atexit([]() {
@@ -570,20 +570,20 @@ finalize()
 {
     if(get_fini_status() != 0)
     {
-        LOG(INFO) << "ignoring finalization request (value=" << get_fini_status() << ")";
+        ROCP_INFO << "ignoring finalization request (value=" << get_fini_status() << ")";
         return;
     }
 
     static auto _sync = std::atomic_flag{};
     if(_sync.test_and_set())
     {
-        LOG(INFO) << "ignoring finalization request [already finalized] (value="
+        ROCP_INFO << "ignoring finalization request [already finalized] (value="
                   << get_fini_status() << ")";
         return;
     }
     // above returns true for all invocations after the first one
 
-    LOG(INFO) << "finalizing rocprofiler (value=" << get_fini_status() << ")";
+    ROCP_INFO << "finalizing rocprofiler (value=" << get_fini_status() << ")";
 
     static auto _once = std::once_flag{};
     std::call_once(_once, []() {
@@ -620,7 +620,7 @@ rocprofiler_is_finalized(int* status)
 rocprofiler_status_t
 rocprofiler_force_configure(rocprofiler_configure_func_t configure_func)
 {
-    LOG(INFO) << "forcing rocprofiler configuration";
+    ROCP_INFO << "forcing rocprofiler configuration";
 
     auto& forced_config = rocprofiler::registration::get_forced_configure();
 
@@ -650,7 +650,7 @@ rocprofiler_set_api_table(const char* name,
     // implementation has a call once
     rocprofiler::registration::init_logging();
 
-    LOG(ERROR) << __FUNCTION__ << "(\"" << name << "\", " << lib_version << ", " << lib_instance
+    ROCP_ERROR << __FUNCTION__ << "(\"" << name << "\", " << lib_version << ", " << lib_instance
                << ", ..., " << num_tables << ")";
 
     static auto _once = std::once_flag{};
@@ -780,7 +780,7 @@ rocprofiler_set_api_table(const char* name,
     }
     else
     {
-        LOG(ERROR) << "rocprofiler does not accept API tables from " << name;
+        ROCP_ERROR << "rocprofiler does not accept API tables from " << name;
 
         return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
     }
@@ -816,8 +816,8 @@ OnLoad(HsaApiTable*       table,
 void
 OnUnload()
 {
-    LOG(INFO) << "Unloading hsa-runtime...";
+    ROCP_INFO << "Unloading hsa-runtime...";
     ::rocprofiler::registration::finalize();
-    LOG(INFO) << "Finalization complete.";
+    ROCP_INFO << "Finalization complete.";
 }
 }
diff --git a/source/lib/rocprofiler-sdk/shared_library.cpp b/source/lib/rocprofiler-sdk/shared_library.cpp
index 9d85b42958..9883a99bd6 100644
--- a/source/lib/rocprofiler-sdk/shared_library.cpp
+++ b/source/lib/rocprofiler-sdk/shared_library.cpp
@@ -46,9 +46,9 @@ lifetime::lifetime()
 
     if(common::get_env("ROCPROFILER_LIBRARY_CTOR", false))
     {
-        LOG(INFO) << "Initializing rocprofiler-sdk library...";
+        ROCP_INFO << "Initializing rocprofiler-sdk library...";
         registration::initialize();
-        LOG(INFO) << "rocprofiler-sdk library initialized";
+        ROCP_INFO << "rocprofiler-sdk library initialized";
     }
 }
 
@@ -56,9 +56,9 @@ lifetime::~lifetime()
 {
     if(common::get_env("ROCPROFILER_LIBRARY_DTOR", false))
     {
-        LOG(INFO) << "Finalizing rocprofiler-sdk library...";
+        ROCP_INFO << "Finalizing rocprofiler-sdk library...";
         registration::finalize();
-        LOG(INFO) << "rocprofiler-sdk library finalized";
+        ROCP_INFO << "rocprofiler-sdk library finalized";
     }
 }