diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/device_counting_service.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/device_counting_service.h index 7ba72cc4d5..cc203062bd 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/device_counting_service.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/device_counting_service.h @@ -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 diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h index 952e80267f..b4c632eafa 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h @@ -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; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/CMakeLists.txt index ef8985e07e..ab71747ee1 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/CMakeLists.txt @@ -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) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.cpp index 843262b5e2..583d0c4396 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.cpp @@ -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}; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.hpp index 4acd904e81..f5e78ab0b4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/controller.hpp @@ -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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.cpp new file mode 100644 index 0000000000..069a07f41e --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.cpp @@ -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 +#include + +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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.hpp new file mode 100644 index 0000000000..6d54117bfe --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/ioctl.hpp @@ -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 + +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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp index c86b009107..332e8adbc1 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp @@ -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 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& test_metrics = {}, - size_t delay = 1) + const std::unordered_set& 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 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>> 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); } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/details/kfd_ioctl.h b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/details/kfd_ioctl.h index 32ca4b76a6..7acc4f88f7 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/details/kfd_ioctl.h +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/details/kfd_ioctl.h @@ -23,8 +23,8 @@ #ifndef KFD_IOCTL_H_INCLUDED #define KFD_IOCTL_H_INCLUDED -#include #include +#include /* * - 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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp index cdd108fe6e..aae235f501 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp @@ -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; } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp index 2a0e91fadc..23be8c1584 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp @@ -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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp index 6977aaf5f3..a7e33d7518 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp @@ -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 const char*