Add support for device counter collection ioctl (#46)

Add support for device counter colleciton ioctl Adds support for the device counter collection IOCTL. This IOCTL allows for device wide counters to be collected even if the queue is not intercepted by rocprofiler-sdk (required for system profilers). A test is also included which checks this behavior by creating a queue that does not have profiling enabled on it and checks to see if SQ counters can be read from it. Note: this test will be skipped if the KFD version does not contain this IOCTL. Right now the check is "soft" in that if the IOCTL is present and there is an error with permissions, rocprofiler will continue but will print an error stating that system wide device profiling and collected counter values may be degraded. This is primarily to avoid breaking existing users (like PAPI) who may not need the IOCTL's capability and to give them time to update. Co-authored-by: Benjamin Welton <ben@amd.com> [ROCm/rocprofiler-sdk commit: c574881cdb]
2024-12-19 13:27:35 -08:00
@@ -106,9 +106,9 @@ rocprofiler_configure_device_counting_service(rocprofiler_context_id_t context_i
 * @param [in] context_id context id
 * @param [in] user_data User supplied data, included in records outputted to buffer.
 * @param [in] flags Flags to specify how the counter data should be collected (defaults to sync).
- * @param [in/out] output_records Output records collected via sampling (output is also written to
+ * @param [in] output_records Output records collected via sampling (output is also written to
 * buffer). Must be allocated by caller.
- * @param [in/out] rec_count On entry, this is the maximum number of records rocprof can store in
+ * @param [in] rec_count On entry, this is the maximum number of records rocprof can store in
 * output_records. On exit, contains the number of actual records.
 * @return ::rocprofiler_status_t
 * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID Returned if the context does not exist or
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2024 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
@@ -106,6 +106,7 @@ typedef enum  // NOLINT(performance-enum-size)
                                             ///< status code for more information.
    ROCPROFILER_STATUS_ERROR_EXCEEDS_HW_LIMIT,          ///< Exceeds hardware limits for collection.
    ROCPROFILER_STATUS_ERROR_AGENT_ARCH_NOT_SUPPORTED,  ///< Agent HW architecture not supported.
+    ROCPROFILER_STATUS_ERROR_PERMISSION_DENIED,         ///< Permission denied.
    ROCPROFILER_STATUS_LAST,
 } rocprofiler_status_t;

@@ -7,7 +7,8 @@ set(ROCPROFILER_LIB_COUNTERS_SOURCES
    dispatch_handlers.cpp
    sample_processing.cpp
    controller.cpp
-    device_counting.cpp)
+    device_counting.cpp
+    ioctl.cpp)
 set(ROCPROFILER_LIB_COUNTERS_HEADERS
    metrics.hpp
    dimensions.hpp
@@ -18,10 +19,10 @@ set(ROCPROFILER_LIB_COUNTERS_HEADERS
    sample_processing.hpp
    controller.hpp
    device_counting.hpp
-    sample_consumer.hpp)
+    sample_consumer.hpp
+    ioctl.hpp)
 target_sources(rocprofiler-sdk-object-library PRIVATE ${ROCPROFILER_LIB_COUNTERS_SOURCES}
                                                      ${ROCPROFILER_LIB_COUNTERS_HEADERS})
-
 add_subdirectory(xml)
 add_subdirectory(parser)
 add_subdirectory(yaml)
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2024 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
@@ -29,6 +29,7 @@

 #include "lib/rocprofiler-sdk/buffer.hpp"
 #include "lib/rocprofiler-sdk/context/context.hpp"
+#include "lib/rocprofiler-sdk/counters/ioctl.hpp"

 namespace rocprofiler
 {
@@ -97,6 +98,18 @@ CounterController::configure_agent_collection(rocprofiler_context_id_t context_i
        return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
    }

+    if(counters::counter_collection_has_device_lock())
+    {
+        /**
+         * Note: This should retrun if the lock fails to aquire in the future. However, this
+         * is a change in the required permissions for rocprofiler and needs to be communicated
+         * with partners before strict enforcement. If the required permissions are not obtained,
+         * those profilers will function as they currently do (without any of the benefits of the
+         * IOCTL).
+         */
+        counters::counter_collection_device_lock(rocprofiler::agent::get_agent(agent_id), true);
+    }
+
    ctx.device_counter_collection->agent_data.emplace_back();
    ctx.device_counter_collection->agent_data.back().callback_data =
        rocprofiler_user_data_t{.ptr = user_data};
@@ -1,7 +1,7 @@

 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2024 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
@@ -26,6 +26,7 @@
 #include "lib/common/synchronized.hpp"
 #include "lib/rocprofiler-sdk/aql/packet_construct.hpp"
 #include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp"
+#include "lib/rocprofiler-sdk/counters/ioctl.hpp"
 #include "lib/rocprofiler-sdk/counters/metrics.hpp"

 #include <rocprofiler-sdk/agent.h>
@@ -0,0 +1,119 @@
+// MIT License
+//
+// Copyright (c) 2024 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 "lib/rocprofiler-sdk/counters/ioctl.hpp"
+#include "lib/rocprofiler-sdk/details/kfd_ioctl.h"
+#include "lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp"
+
+#include <sys/ioctl.h>
+#include <cerrno>
+
+namespace rocprofiler
+{
+namespace counters
+{
+bool
+counter_collection_has_device_lock()
+{
+    kfd_ioctl_profiler_args args = {};
+    args.op                      = KFD_IOC_PROFILER_VERSION;
+    int ret = ioctl(pc_sampling::ioctl::get_kfd_fd(), AMDKFD_IOC_PROFILER, &args);
+    if(ret == 0)
+    {
+        return true;
+    }
+    return false;
+}
+
+rocprofiler_status_t
+counter_collection_device_lock(const rocprofiler_agent_t* agent, bool all_queues)
+{
+    CHECK(agent);
+    kfd_ioctl_profiler_args args = {};
+    args.op                      = KFD_IOC_PROFILER_PMC;
+    args.pmc.gpu_id              = agent->gpu_id;
+    args.pmc.lock                = 1;
+    args.pmc.perfcount_enable    = all_queues ? 1 : 0;
+
+    int ret = ioctl(pc_sampling::ioctl::get_kfd_fd(), AMDKFD_IOC_PROFILER, &args);
+    if(ret != 0)
+    {
+        switch(ret)
+        {
+            case -EBUSY:
+                ROCP_WARNING << fmt::format(
+                    "Device {} has a profiler attached to it. PMC Counters may be inaccurate.",
+                    agent->id.handle);
+                return ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES;
+            case -EPERM:
+                ROCP_WARNING << fmt::format(
+                    "Device {} could not be locked for profiling due to lack of permissions "
+                    "(capability SYS_PERFMON). PMC Counters may be inaccurate and System Counter "
+                    "Collection will be degraded.");
+                return ROCPROFILER_STATUS_ERROR_PERMISSION_DENIED;
+            case -EINVAL:
+                ROCP_WARNING << fmt::format(
+                    "Driver/Kernel version does not support locking device {}. PMC Counters may be "
+                    "inaccurate and System Counter Collection will be degraded.",
+                    agent->id.handle);
+                return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI;
+            default:
+                ROCP_WARNING << fmt::format(
+                    "Failed to lock device {}. PMC Counters may be inaccurate and System Counter "
+                    "Collection will be degraded.",
+                    agent->id.handle);
+                return ROCPROFILER_STATUS_ERROR;
+        }
+    }
+
+    return ROCPROFILER_STATUS_SUCCESS;
+}
+
+// Not required now but may be useful in the future.
+// rocprofiler_status_t
+// counter_collection_device_unlock(const rocprofiler_agent_t* agent) {
+//     CHECK(agent);
+//     kfd_ioctl_profiler_args args = {};
+//     args.op = KFD_IOC_PROFILER_PMC;
+//     args.pmc.gpu_id = agent->gpu_id;
+//     args.pmc.lock = 0;
+//     args.pmc.perfcount_enable = 0;
+
+//     int ret = ioctl(pc_sampling::ioctl::get_kfd_fd(), AMDKFD_IOC_PROFILER, &args);
+//     if (ret != 0) {
+//         switch (ret) {
+//             case -EBUSY:
+//             case -EPERM:
+//                 ROCP_WARNING << fmt::format("Could not unlock the device {}", agent->id.handle);
+//                 return ROCPROFILER_STATUS_ERROR;
+//             case -EINVAL:
+//                 return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI;
+//             default:
+//                 ROCP_WARNING << fmt::format("Could not unlock the device {}", agent->id.handle);
+//                 return ROCPROFILER_STATUS_ERROR;
+//         }
+//     }
+
+//     return ROCPROFILER_STATUS_SUCCESS;
+// }
+}  // namespace counters
+}  // namespace rocprofiler
@@ -0,0 +1,37 @@
+// MIT License
+//
+// Copyright (c) 2024 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 <rocprofiler-sdk/rocprofiler.h>
+
+namespace rocprofiler
+{
+namespace counters
+{
+bool
+counter_collection_has_device_lock();
+
+rocprofiler_status_t
+counter_collection_device_lock(const rocprofiler_agent_t* agent, bool all_queues);
+
+}  // namespace counters
+}  // namespace rocprofiler
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2024 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
@@ -77,11 +77,11 @@ findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set<std::st
    std::vector<counters::Metric> ret;
    const auto*                   all_counters = counters::getMetricMap();

-    ROCP_ERROR << "Looking up counters for " << std::string(agent.name());
+    ROCP_WARNING << "Looking up counters for " << std::string(agent.name());
    const auto* gfx_metrics = common::get_val(*all_counters, std::string(agent.name()));
    if(!gfx_metrics)
    {
-        ROCP_ERROR << "No counters found for " << std::string(agent.name());
+        ROCP_WARNING << "No counters found for " << std::string(agent.name());
        return ret;
    }

@@ -145,7 +145,7 @@ check_output_created(rocprofiler_context_id_t,
                break;
            }
            found_value = record->user_data.value;
-            // ROCP_ERROR << fmt::format("Found counter value: {}", record->counter_value);
+            // ROCP_WARNING << fmt::format("Found counter value: {}", record->counter_value);
            global_recs().wlock([&](auto& data) { data.push_back(*record); });
        }
    }
@@ -202,7 +202,7 @@ gen_kernel_pkt(uint64_t obj)
    packet.kernel_dispatch.kernel_object    = obj;
    packet.kernel_dispatch.kernarg_address  = nullptr;
    packet.kernel_dispatch.completion_signal.handle = 0;
-    ROCP_ERROR << fmt::format("{:x}", packet.kernel_dispatch.kernel_object);
+    ROCP_WARNING << fmt::format("{:x}", packet.kernel_dispatch.kernel_object);
    return packet;
 }

@@ -247,8 +247,9 @@ protected:
    device_counting_service_test() {}

    static void test_run(rocprofiler_counter_flag_t flags = ROCPROFILER_COUNTER_FLAG_NONE,
-                         const std::unordered_set<std::string>& test_metrics = {},
-                         size_t                                 delay        = 1)
+                         const std::unordered_set<std::string>& test_metrics  = {},
+                         size_t                                 delay         = 1,
+                         bool                                   non_intercept = false)
    {
        hsa_init();
        registration::init_logging();
@@ -282,9 +283,6 @@ protected:
                                      &queue),
                     HSA_STATUS_SUCCESS);

-            // We don't use the queue interceptor, need to enabling profiling manually
-            hsa_amd_profiling_set_profiler_enabled(queue, 1);
-
            hsa_signal_t completion_signal;
            hsa_signal_create(1, 0, nullptr, &completion_signal);

@@ -292,21 +290,35 @@ protected:
            CHECK(agent.get_hsa_agent().handle != 0);
            // Set state of the queue to allow profiling (may not be needed since AQL
            // may do this in the future).
-            aql::set_profiler_active_on_queue(
-                agent.cpu_pool(), agent.get_hsa_agent(), [&](hsa::rocprofiler_packet pkt) {
-                    pkt.ext_amd_aql_pm4.completion_signal = completion_signal;
-                    submitPacket(queue, (const void*) &pkt);
+            if(!non_intercept)
+            {
+                // This simulates the presence of us intercepting queues on queue creation.
+                // This is identical to the standard device counting use case where only a single
+                // process is being profiled.
+                hsa_amd_profiling_set_profiler_enabled(queue, 1);
+                aql::set_profiler_active_on_queue(
+                    agent.cpu_pool(), agent.get_hsa_agent(), [&](hsa::rocprofiler_packet pkt) {
+                        pkt.ext_amd_aql_pm4.completion_signal = completion_signal;
+                        submitPacket(queue, (const void*) &pkt);

-                    if(hsa_signal_wait_relaxed(completion_signal,
-                                               HSA_SIGNAL_CONDITION_EQ,
-                                               0,
-                                               20000000,
-                                               HSA_WAIT_STATE_BLOCKED) != 0)
-                    {
-                        ROCP_FATAL << "Failed to set profiling mode on queue";
-                    }
-                    hsa_signal_store_relaxed(completion_signal, 1);
-                });
+                        if(hsa_signal_wait_relaxed(completion_signal,
+                                                   HSA_SIGNAL_CONDITION_EQ,
+                                                   0,
+                                                   20000000,
+                                                   HSA_WAIT_STATE_BLOCKED) != 0)
+                        {
+                            ROCP_FATAL << "Failed to set profiling mode on queue";
+                        }
+                        hsa_signal_store_relaxed(completion_signal, 1);
+                    });
+            }
+            else
+            {
+                // In the non_intercept case, we are simulating queues that are created without
+                // interception on the system. This case is used to test the device counting service
+                // in modes where a system profiler would be present (and we would not have the
+                // ability to intercept queues in order to do the above operations).
+            }

            rocprofiler::hsa::rocprofiler_packet barrier{};

@@ -322,7 +334,7 @@ protected:
                std::vector<rocprofiler_record_counter_t> output_records(10000);
                // global_recs().clear();
                track_metric++;
-                ROCP_ERROR << "Testing metric " << metric.name();
+                ROCP_WARNING << "Testing metric " << metric.name();
                rocprofiler_context_id_t ctx = {.handle = 0};
                ROCPROFILER_CALL(rocprofiler_create_context(&ctx), "context creation failed");
                rocprofiler_buffer_id_t opt_buff_id = {.handle = 0};
@@ -375,8 +387,8 @@ protected:
                auto status = rocprofiler_start_context(ctx);
                if(status == ROCPROFILER_STATUS_ERROR_NO_HARDWARE_COUNTERS)
                {
-                    ROCP_ERROR << fmt::format("No hardware counters for {}, skipping",
-                                              metric.name());
+                    ROCP_WARNING << fmt::format("No hardware counters for {}, skipping",
+                                                metric.name());
                    continue;
                }
                else if(status != ROCPROFILER_STATUS_SUCCESS)
@@ -504,8 +516,8 @@ protected:
            test_kernels kernel_loader(gpu_agent);
            auto         kernel_handle = kernel_loader.load_kernel(gpu_agent, "null_kernel");

-            ROCP_ERROR << fmt::format("Running test on agent {:x}",
-                                      gpu_agent.get_hsa_agent().handle);
+            ROCP_WARNING << fmt::format("Running test on agent {:x}",
+                                        gpu_agent.get_hsa_agent().handle);

            const auto* agent_map = rocprofiler::common::get_val(counters::get_ast_map(),
                                                                 std::string(gpu_agent.name()));
@@ -584,19 +596,19 @@ protected:
                                        UINT32_MAX,
                                        HSA_WAIT_STATE_ACTIVE);

-                ROCP_ERROR << "Processing Next...";
+                ROCP_WARNING << "Processing Next...";
                auto decoded_pkt = counters::EvaluateAST::read_pkt(&pkt_constructor, *inst_pkts);
                CHECK(!decoded_pkt.empty());
-                ROCP_ERROR << "Decoded Packet:";
+                ROCP_WARNING << "Decoded Packet:";
                for(const auto& [id, data_vec] : decoded_pkt)
                {
-                    ROCP_ERROR << fmt::format("\t[{} = {}]", id, fmt::join(data_vec, ","));
+                    ROCP_WARNING << fmt::format("\t[{} = {}]", id, fmt::join(data_vec, ","));
                }

                std::vector<std::unique_ptr<std::vector<rocprofiler_record_counter_t>>> cache;
                auto* ret = counter_ast.evaluate(decoded_pkt, cache);
                CHECK(!ret->empty());
-                ROCP_ERROR << fmt::format(
+                ROCP_WARNING << fmt::format(
                    "Final Decoded Counter Values: {} (iter={})", fmt::join(*ret, ","), i);

                CHECK_EQ(ret->size(), expected_values.size());
@@ -636,7 +648,7 @@ TEST_F(device_counting_service_test, sync_grbm_verify)
 {
    test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"GRBM_COUNT"}, 50000);
    auto local_recs = global_recs().rlock([](const auto& data) { return data; });
-    ROCP_ERROR << local_recs.size();
+    ROCP_WARNING << local_recs.size();

    for(const auto& val : local_recs)
    {
@@ -644,7 +656,7 @@ TEST_F(device_counting_service_test, sync_grbm_verify)
        rocprofiler_query_record_counter_id(val.id, &id);
        rocprofiler_counter_info_v0_t info;
        rocprofiler_query_counter_info(id, ROCPROFILER_COUNTER_INFO_VERSION_0, &info);
-        ROCP_ERROR << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
+        ROCP_WARNING << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
        EXPECT_GT(val.counter_value, 0.0);
    }
 }
@@ -653,7 +665,7 @@ TEST_F(device_counting_service_test, sync_gpu_util_verify)
 {
    test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"GPU_UTIL"}, 50000);
    auto local_recs = global_recs().rlock([](const auto& data) { return data; });
-    ROCP_ERROR << local_recs.size();
+    ROCP_WARNING << local_recs.size();

    for(const auto& val : local_recs)
    {
@@ -661,7 +673,7 @@ TEST_F(device_counting_service_test, sync_gpu_util_verify)
        rocprofiler_query_record_counter_id(val.id, &id);
        rocprofiler_counter_info_v0_t info;
        rocprofiler_query_counter_info(id, ROCPROFILER_COUNTER_INFO_VERSION_0, &info);
-        ROCP_ERROR << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
+        ROCP_WARNING << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
        EXPECT_GT(val.counter_value, 0.0);
    }
 }
@@ -670,7 +682,7 @@ TEST_F(device_counting_service_test, sync_sq_waves_verify)
 {
    test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"SQ_WAVES_sum"}, 50000);
    auto local_recs = global_recs().rlock([](const auto& data) { return data; });
-    ROCP_ERROR << local_recs.size();
+    ROCP_WARNING << local_recs.size();

    for(const auto& val : local_recs)
    {
@@ -678,7 +690,32 @@ TEST_F(device_counting_service_test, sync_sq_waves_verify)
        rocprofiler_query_record_counter_id(val.id, &id);
        rocprofiler_counter_info_v0_t info;
        rocprofiler_query_counter_info(id, ROCPROFILER_COUNTER_INFO_VERSION_0, &info);
-        ROCP_ERROR << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
+        ROCP_WARNING << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
+        EXPECT_GT(val.counter_value, 0.0);
+    }
+}
+
+TEST_F(device_counting_service_test, sync_sq_waves_verify_non_intercept)
+{
+    // If this test fails, device counters will not be read correctly by a system-wide profiler
+    // deamon.
+    if(!counters::counter_collection_has_device_lock())
+    {
+        ROCP_WARNING << "Unsupported kernel driver version, skipping test";
+        GTEST_SKIP();
+    }
+
+    test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"SQ_WAVES_sum"}, 50000, true);
+    auto local_recs = global_recs().rlock([](const auto& data) { return data; });
+    ROCP_WARNING << local_recs.size();
+
+    for(const auto& val : local_recs)
+    {
+        rocprofiler_counter_id_t id;
+        rocprofiler_query_record_counter_id(val.id, &id);
+        rocprofiler_counter_info_v0_t info;
+        rocprofiler_query_counter_info(id, ROCPROFILER_COUNTER_INFO_VERSION_0, &info);
+        ROCP_WARNING << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value);
        EXPECT_GT(val.counter_value, 0.0);
    }
 }
@@ -23,8 +23,8 @@
 #ifndef KFD_IOCTL_H_INCLUDED
 #define KFD_IOCTL_H_INCLUDED

-#include <drm.h>
 #include <linux/ioctl.h>
+#include <linux/types.h>

 /*
 * - 1.1 - initial version
@@ -1789,6 +1789,35 @@ struct kfd_ioctl_pc_sample_args
    __u32 version;
 };

+#define KFD_IOC_PROFILER_VERSION_NUM 1
+enum kfd_profiler_ops
+{
+    KFD_IOC_PROFILER_PMC       = 0,
+    KFD_IOC_PROFILER_PC_SAMPLE = 1,
+    KFD_IOC_PROFILER_VERSION   = 2,
+};
+
+/**
+ * Enables/Disables GPU Specific profiler settings
+ */
+struct kfd_ioctl_pmc_settings
+{
+    __u32 gpu_id;           /* This is the user_gpu_id */
+    __u32 lock;             /* Lock GPU for Profiling */
+    __u32 perfcount_enable; /* Force Perfcount Enable for queues on GPU */
+};
+
+struct kfd_ioctl_profiler_args
+{
+    __u32 op; /* kfd_profiler_op */
+    union
+    {
+        struct kfd_ioctl_pc_sample_args pc_sample;
+        struct kfd_ioctl_pmc_settings   pmc;
+        __u32                           version; /* KFD_IOC_PROFILER_VERSION_NUM */
+    };
+};
+
 #define AMDKFD_IOCTL_BASE     'K'
 #define AMDKFD_IO(nr)         _IO(AMDKFD_IOCTL_BASE, nr)
 #define AMDKFD_IOR(nr, type)  _IOR(AMDKFD_IOCTL_BASE, nr, type)
@@ -1895,7 +1924,9 @@ struct kfd_ioctl_pc_sample_args

 #define AMDKFD_IOC_PC_SAMPLE AMDKFD_IOWR(0x85, struct kfd_ioctl_pc_sample_args)

+#define AMDKFD_IOC_PROFILER AMDKFD_IOWR(0x86, struct kfd_ioctl_profiler_args)
+
 #define AMDKFD_COMMAND_START_2 0x80
-#define AMDKFD_COMMAND_END_2   0x86
+#define AMDKFD_COMMAND_END_2   0x87

 #endif
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 ROCm Developer Tools
+// Copyright (c) 2024 ROCm Developer Tools
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -74,13 +74,6 @@ kfd_open()
    return fd;
 }

-int
-get_kfd_fd()
-{
-    static auto _v = kfd_open();
-    return _v;
-}
-
 /** Call ioctl, restarting if it is interrupted
 * Taken from libhsakmt.c
 */
@@ -283,7 +276,7 @@ ioctl_query_pc_sampling_capabilities(uint32_t  kfd_gpu_id,
            // which is not supported.
            return ROCPROFILER_IOCTL_STATUS_UNAVAILABLE;
        }
-        ROCP_ERROR << "IOCTL failed to query PC sampling configs: " << ret << "\n";
+        ROCP_WARNING << "IOCTL failed to query PC sampling configs: " << ret << "\n";
    }
    *size = args.num_sample_info;

@@ -352,6 +345,13 @@ convert_ioctl_pcs_config_to_rocp(const rocprofiler_ioctl_pc_sampling_info_t& ioc
 }
 }  // namespace

+int
+get_kfd_fd()
+{
+    static auto _v = kfd_open();
+    return _v;
+}
+
 rocprofiler_status_t
 ioctl_query_pcs_configs(const rocprofiler_agent_t* agent, rocp_pcs_cfgs_vec_t& rocp_configs)
 {
@@ -381,7 +381,7 @@ ioctl_query_pcs_configs(const rocprofiler_agent_t* agent, rocp_pcs_cfgs_vec_t& r
    }
    else if(ret != ROCPROFILER_IOCTL_STATUS_SUCCESS)
    {
-        ROCP_ERROR << "......... Failed while iterating over PC sampling configurations\n";
+        ROCP_WARNING << "......... Failed while iterating over PC sampling configurations\n";
        return ROCPROFILER_STATUS_ERROR;
    }

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 ROCm Developer Tools
+// Copyright (c) 2024 ROCm Developer Tools
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -45,6 +45,8 @@ ioctl_pcs_create(const rocprofiler_agent_t*       agent,
                 uint64_t                         interval,
                 uint32_t*                        ioctl_pcs_id);

+int
+get_kfd_fd();
 }  // namespace ioctl
 }  // namespace pc_sampling
 }  // namespace rocprofiler
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2024 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
@@ -116,6 +116,8 @@ ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_EXCEEDS_HW_LIMIT,
                          "Request exceeds the capabilities of the hardware to collect")
 ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_AGENT_ARCH_NOT_SUPPORTED,
                          "Agent HW architecture is not supported, no counter metrics found.")
+ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_PERMISSION_DENIED,
+                          "Required permission (CAP_PERFMON) is not set, permission denied")

 template <size_t Idx, size_t... Tail>
 const char*