diff --git a/projects/aqlprofile/.github/workflows/codeql.yml b/projects/aqlprofile/.github/workflows/codeql.yml index 5451f4d2cc..89660c9543 100644 --- a/projects/aqlprofile/.github/workflows/codeql.yml +++ b/projects/aqlprofile/.github/workflows/codeql.yml @@ -53,7 +53,7 @@ jobs: run: | git config --global --add safe.directory '*' apt-get update - apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev rocm-llvm-dev + apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev rocm-llvm-dev libgtest-dev libgmock-dev update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11 --slave /usr/bin/gcov gcov /usr/bin/gcov-11 update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12 --slave /usr/bin/gcov gcov /usr/bin/gcov-12 diff --git a/projects/aqlprofile/.github/workflows/continuous_integration.yml b/projects/aqlprofile/.github/workflows/continuous_integration.yml index 6d079489c9..5e22c859a5 100644 --- a/projects/aqlprofile/.github/workflows/continuous_integration.yml +++ b/projects/aqlprofile/.github/workflows/continuous_integration.yml @@ -66,7 +66,7 @@ jobs: run: | git config --global --add safe.directory '*' apt-get update - apt-get install -y build-essential cmake g++-11 g++-12 python3-pip + apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libgtest-dev libgmock-dev update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11 --slave /usr/bin/gcov gcov /usr/bin/gcov-11 update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12 --slave /usr/bin/gcov gcov /usr/bin/gcov-12 @@ -89,6 +89,7 @@ jobs: -DCTEST_BINARY_DIRECTORY="$(pwd)/build" -DAQLPROFILE_BUILD_NUM_JOBS="16" -DCTEST_SITE="${RUNNER_HOSTNAME}" -DCTEST_BUILD_NAME=PR_${{ github.ref_name }}_${{ github.repository }}-${{ matrix.os }}-${{ matrix.runner }}-core -DCMAKE_CTEST_ARGUMENTS="" + -DAQLPROFILE_BUILD_TESTS=ON -DAQLPROFILE_EXTRA_CONFIGURE_ARGS="" -S ./dashboard.cmake @@ -142,5 +143,6 @@ jobs: -DCTEST_BINARY_DIRECTORY="$(pwd)/build" -DAQLPROFILE_BUILD_NUM_JOBS="16" -DCTEST_SITE="${RUNNER_HOSTNAME}" -DCTEST_BUILD_NAME=PR_${{ github.ref_name }}_${{ github.repository }}-${{ matrix.os }}-${{ matrix.runner }}-core -DCMAKE_CTEST_ARGUMENTS="" + -DAQLPROFILE_BUILD_TESTS=ON -DAQLPROFILE_EXTRA_CONFIGURE_ARGS="" -S ./dashboard.cmake diff --git a/projects/aqlprofile/CMakeLists.txt b/projects/aqlprofile/CMakeLists.txt index ce7c8ebf9c..9f8bc3a678 100644 --- a/projects/aqlprofile/CMakeLists.txt +++ b/projects/aqlprofile/CMakeLists.txt @@ -65,6 +65,9 @@ set ( DEST_NAME ${AQLPROFILE_NAME} ) install ( TARGETS ${AQLPROFILE_TARGET} LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT runtime ) install ( TARGETS ${AQLPROFILE_TARGET} LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT asan ) +# Enable/disable test +option(AQLPROFILE_BUILD_TESTS "Build tests for AQLProfile" OFF) + ## Add the packaging directives for the runtime library. if ( ENABLE_ASAN_PACKAGING ) set ( CPACK_PACKAGE_NAME ${AQLPROFILE_NAME}-asan ) diff --git a/projects/aqlprofile/build.sh b/projects/aqlprofile/build.sh index 8164d4ce1e..75996f7d35 100755 --- a/projects/aqlprofile/build.sh +++ b/projects/aqlprofile/build.sh @@ -65,6 +65,7 @@ cmake \ -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=FALSE \ -DCPACK_GENERATOR="STGZ" \ -DGPU_TARGETS="$GPU_LIST" \ + -DAQLPROFILE_BUILD_TESTS=OFF \ -DCPACK_OBJCOPY_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-objcopy" \ -DCPACK_READELF_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-readelf" \ -DCPACK_STRIP_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-strip" \ diff --git a/projects/aqlprofile/dashboard.cmake b/projects/aqlprofile/dashboard.cmake index 04a8a0b65c..05716f1de4 100644 --- a/projects/aqlprofile/dashboard.cmake +++ b/projects/aqlprofile/dashboard.cmake @@ -42,7 +42,7 @@ if(NOT DEFINED AQLPROFILE_BUILD_NUM_JOBS) set(AQLPROFILE_BUILD_NUM_JOBS "16") endif() -set(CTEST_CONFIGURE_COMMAND "cmake -B ${CTEST_BINARY_DIRECTORY} -DCMAKE_BUILD_TYPE='RelWithDebInfo' -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_INSTALL_PREFIX=/opt/rocm -DCPACK_PACKAGING_INSTALL_PREFIX=/opt/rocm -DCPACK_GENERATOR='DEB;RPM;STGZ' -DGPU_TARGETS='gfx906,gfx90a,gfx942,gfx1101,gfx1201' ${AQLPROFILE_EXTRA_CONFIGURE_ARGS} ${CTEST_SOURCE_DIRECTORY}") +set(CTEST_CONFIGURE_COMMAND "cmake -B ${CTEST_BINARY_DIRECTORY} -DCMAKE_BUILD_TYPE='RelWithDebInfo' -DCMAKE_PREFIX_PATH=/opt/rocm -DAQLPROFILE_BUILD_TESTS=ON -DCMAKE_INSTALL_PREFIX=/opt/rocm -DCPACK_PACKAGING_INSTALL_PREFIX=/opt/rocm -DCPACK_GENERATOR='DEB;RPM;STGZ' -DGPU_TARGETS='gfx906,gfx90a,gfx942,gfx1101,gfx1201' ${AQLPROFILE_EXTRA_CONFIGURE_ARGS} ${CTEST_SOURCE_DIRECTORY}") set(CTEST_BUILD_COMMAND "cmake --build \"${CTEST_BINARY_DIRECTORY}\" -- -j ${AQLPROFILE_BUILD_NUM_JOBS} all mytest") if(NOT DEFINED CTEST_SITE) diff --git a/projects/aqlprofile/src/CMakeLists.txt b/projects/aqlprofile/src/CMakeLists.txt index 320a5bd8d8..3f7bfee02c 100644 --- a/projects/aqlprofile/src/CMakeLists.txt +++ b/projects/aqlprofile/src/CMakeLists.txt @@ -73,4 +73,4 @@ add_custom_target( mygen COMMAND sh -xc "sed 's/_GPU_BLOCKINFO_H_/SRC_DEF_GPU_BLOCK_INFO_H_/' ${BINFO_DEF} >>${BINFO_HEADER}" ) -add_subdirectory(src/core) \ No newline at end of file +add_subdirectory(src/core) diff --git a/projects/aqlprofile/src/core/CMakeLists.txt b/projects/aqlprofile/src/core/CMakeLists.txt index 2ae4097206..2845d75296 100644 --- a/projects/aqlprofile/src/core/CMakeLists.txt +++ b/projects/aqlprofile/src/core/CMakeLists.txt @@ -1 +1,5 @@ -add_subdirectory(include) \ No newline at end of file +add_subdirectory(include) +if(AQLPROFILE_BUILD_TESTS) + enable_testing() + add_subdirectory(tests) +endif() \ No newline at end of file diff --git a/projects/aqlprofile/src/core/commandbuffermgr.hpp b/projects/aqlprofile/src/core/commandbuffermgr.hpp index cb5d7dc1ee..8456e676b7 100644 --- a/projects/aqlprofile/src/core/commandbuffermgr.hpp +++ b/projects/aqlprofile/src/core/commandbuffermgr.hpp @@ -23,6 +23,7 @@ #pragma once #include +#include #include #include #include diff --git a/projects/aqlprofile/src/core/tests/CMakeLists.txt b/projects/aqlprofile/src/core/tests/CMakeLists.txt new file mode 100644 index 0000000000..eaabc56536 --- /dev/null +++ b/projects/aqlprofile/src/core/tests/CMakeLists.txt @@ -0,0 +1,119 @@ +# write a CMakeLists.txt file for the gfx9_factory_test +# that includes the necessary libraries and sets the properties for the test +include(GoogleTest) +find_package(GTest REQUIRED) +include_directories(${GTEST_INCLUDE_DIRS}) + + +find_library( + hsa-amd-aqlprofile64 + REQUIRED + NAMES hsa-amd-aqlprofile64 hsa-amd-aqlprofile + HINTS /opt/rocm/lib /opt/rocm + PATHS /opt/rocm/lib /opt/rocm) + +# Add test for memory manager +add_executable(gfx9-memory-manager-test) +SET(AQLPROFILE_MEMORYMANAGER_SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/memorymanager_tests.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/../memorymanager.cpp +) + +target_sources(gfx9-memory-manager-test PRIVATE ${AQLPROFILE_MEMORYMANAGER_SOURCES}) +target_include_directories(gfx9-memory-manager-test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${LIB_DIR}) +target_link_libraries( + gfx9-memory-manager-test + PRIVATE + hsa-runtime64::hsa-runtime64 + GTest::gtest + GTest::gtest_main) + + +gtest_add_tests( + TARGET gfx9-memory-manager-test + SOURCES ${AQLPROFILE_MEMORYMANAGER_SOURCES} + TEST_LIST gfx9-memory-manager_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + +set_tests_properties( + ${gfx9-memory-manager_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION + "${AQLPROFILE_DEFAULT_FAIL_REGEX}") + +# Add test for aql profile +add_executable(aqlprofile-test) +SET(AQLPROFILE_TEST_SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/aql_profile_tests.cpp +) + +target_sources(aqlprofile-test PRIVATE ${AQLPROFILE_TEST_SOURCES}) +target_include_directories(aqlprofile-test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${LIB_DIR}) +target_link_libraries( + aqlprofile-test + PRIVATE + hsa-runtime64::hsa-runtime64 + GTest::gtest + GTest::gtest_main + GTest::gmock + GTest::gmock_main) + + +gtest_add_tests( + TARGET aqlprofile-test + SOURCES ${AQLPROFILE_TEST_SOURCES} + TEST_LIST aqlprofile-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + +set_tests_properties( + ${aqlprofile-test_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION + "${AQLPROFILE_DEFAULT_FAIL_REGEX}") + + +# Add tests for command buffer +add_executable(command-buffer-test) +SET(AQLPROFILE_COMMAND_BUFFER_SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/command_buffer_tests.cpp +) +target_sources(command-buffer-test PRIVATE ${AQLPROFILE_COMMAND_BUFFER_SOURCES}) +target_include_directories(command-buffer-test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${LIB_DIR}) +target_link_libraries( + command-buffer-test + PRIVATE + hsa-runtime64::hsa-runtime64 + GTest::gtest + GTest::gtest_main + GTest::gmock + GTest::gmock_main) + +gtest_add_tests( + TARGET command-buffer-test + SOURCES ${AQLPROFILE_COMMAND_BUFFER_SOURCES} + TEST_LIST command-buffer-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) +set_tests_properties( + ${command-buffer-test_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION + "${AQLPROFILE_DEFAULT_FAIL_REGEX}") + +# Add tests for counters +add_executable(counters-test) +SET(AQLPROFILE_COUNTERS_SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/counter_tests.cpp +) +target_sources(counters-test PRIVATE ${AQLPROFILE_COUNTERS_SOURCES}) +target_include_directories(counters-test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${LIB_DIR}) +target_link_libraries( + counters-test + PRIVATE + hsa-runtime64::hsa-runtime64 + ${hsa-amd-aqlprofile64} + GTest::gtest + GTest::gtest_main + GTest::gmock + GTest::gmock_main) +gtest_add_tests( + TARGET counters-test + SOURCES ${AQLPROFILE_COUNTERS_SOURCES} + TEST_LIST counters-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) +set_tests_properties( + ${counters-test_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION + "${AQLPROFILE_DEFAULT_FAIL_REGEX}") \ No newline at end of file diff --git a/projects/aqlprofile/src/core/tests/aql_profile_tests.cpp b/projects/aqlprofile/src/core/tests/aql_profile_tests.cpp new file mode 100644 index 0000000000..ba932e64c4 --- /dev/null +++ b/projects/aqlprofile/src/core/tests/aql_profile_tests.cpp @@ -0,0 +1,245 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include + +#include "core/aql_profile.hpp" +#include "core/pm4_factory.h" +// header for memcpy +#include + + +//#include "core/counter_dimensions.hpp" +//#include "core/include/aql_profile_v2.h" + +using namespace aql_profile; +using namespace testing; + +namespace aql_profile { +bool Pm4Factory::concurrent_create_mode_ = false; +bool Pm4Factory::spm_kfd_mode_ = false; +Pm4Factory::mutex_t Pm4Factory::mutex_; +Pm4Factory::instances_t* Pm4Factory::instances_ = nullptr; +} + +// Mock classes to simulate Pm4Factory and related functionality +class MockPm4Factory : public Pm4Factory { +public: + MockPm4Factory() : Pm4Factory(BlockInfoMap(nullptr, 0)) {} + MOCK_METHOD(const GpuBlockInfo*, GetBlockInfo, (const hsa_ven_amd_aqlprofile_event_t*), (const)); + MOCK_METHOD(bool, IsGFX9, (), (const)); + MOCK_METHOD(bool, IsConcurrent, (), (const)); + MOCK_METHOD(int, GetNumWGPs, (), (const)); + MOCK_METHOD(bool, SpmKfdMode, (), (const)); + MOCK_METHOD(bool, SPISkip, (uint32_t, uint32_t), (const)); +}; + + +// Helper to create a mock GpuBlockInfo +GpuBlockInfo* CreateBlockInfo(uint32_t id, uint32_t counter_count, uint32_t attr = 0) { + auto* info = new GpuBlockInfo(); + info->id = id; + info->counter_count = counter_count; + info->attr = attr; + info->instance_count = 1; + return info; +} + +// Helper to create a profile with specified events +hsa_ven_amd_aqlprofile_profile_t* CreateProfile(const std::vector& events) { + auto* profile = new hsa_ven_amd_aqlprofile_profile_t(); + profile->event_count = events.size(); + if (!events.empty()) { + memcpy(reinterpret_cast(&profile->events), &events, sizeof(hsa_ven_amd_aqlprofile_event_t)); + } else { + profile->events = nullptr; + } + return profile; +} + +void DeleteProfile(hsa_ven_amd_aqlprofile_profile_t* profile) { + if (profile) { + delete[] profile->events; + delete profile; + } +} + +hsa_status_t DefaultTracedataCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, + hsa_ven_amd_aqlprofile_info_data_t* info_data, + void* callback_data) { + hsa_status_t status = HSA_STATUS_SUCCESS; + hsa_ven_amd_aqlprofile_info_data_t* passed_data = + reinterpret_cast(callback_data); + + if (info_type == HSA_VEN_AMD_AQLPROFILE_INFO_TRACE_DATA) { + if (info_data->sample_id == passed_data->sample_id) { + passed_data->trace_data = info_data->trace_data; + status = HSA_STATUS_INFO_BREAK; + } + } + + return status; +} + +// Test fixture for CountersVec tests +class CountersVecTest : public Test { +protected: + void SetUp() override { + pm4_factory = new NiceMock(); + ON_CALL(*pm4_factory, IsGFX9()).WillByDefault(Return(true)); + } + void TearDown() override { + delete pm4_factory; + } + NiceMock* pm4_factory; + + pm4_builder::counters_vector CountersVec(const profile_t* profile, const Pm4Factory* pm4_factory); +}; + +pm4_builder::counters_vector CountersVecTest::CountersVec(const profile_t* profile, + const Pm4Factory* pm4_factory) { + pm4_builder::counters_vector vec; + std::map index_map; + for (const hsa_ven_amd_aqlprofile_event_t* p = profile->events; + p < profile->events + profile->event_count; ++p) { + const GpuBlockInfo* block_info = pm4_factory->GetBlockInfo(p); + const block_des_t block_des = {pm4_factory->GetBlockInfo(p)->id, p->block_index}; + + // Counting counter register index per block + const auto ret = index_map.insert({block_des, 0}); + uint32_t& reg_index = ret.first->second; + + if (pm4_builder::SPISkip(block_info->attr, p->counter_id)) + { + vec.push_back({p->counter_id, reg_index, block_des, block_info}); + continue; + } + + if (reg_index >= block_info->counter_count) { + throw event_exception("Event is out of block counter registers number limit, ", *p); + } + + vec.push_back({p->counter_id, reg_index, block_des, block_info}); + + ++reg_index; + } + return vec; +} + +//Test case: Empty profile (no events) +TEST_F(CountersVecTest, EmptyProfile) { + auto profile = CreateProfile({}); + auto counters = CountersVec(profile, pm4_factory); + EXPECT_TRUE(counters.empty()); + DeleteProfile(profile); +} + +// Test case: Profile with regular events +TEST_F(CountersVecTest, RegularEvents) { + GpuBlockInfo* block_info1 = CreateBlockInfo(1, 4); + GpuBlockInfo* block_info2 = CreateBlockInfo(2, 2); + + std::vector events = { + {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4} + }; + + auto profile = CreateProfile(events); + EXPECT_NE(profile,nullptr); + EXPECT_EQ(profile->events->block_index, 0); + EXPECT_EQ(profile->events->counter_id, 4); + + pm4_factory->GetBlockInfo(profile->events); + bool is_gfx9 = pm4_factory->IsGFX9(); + EXPECT_TRUE(is_gfx9); + +} + +// Test fixture for the DefaultTracedataCallback function +class DefaultTracedataCallbackTest : public Test { +protected: + hsa_ven_amd_aqlprofile_info_data_t CreateInfoData(uint32_t sample_id) { + hsa_ven_amd_aqlprofile_info_data_t data{}; + data.sample_id = sample_id; + data.trace_data.ptr = reinterpret_cast(0x1000 + sample_id); + data.trace_data.size = 0x100 + sample_id; + return data; + } +}; + +//Test case: DefaultTracedataCallback with matching sample IDs +TEST_F(DefaultTracedataCallbackTest, MatchingSampleId) { + auto info_data = CreateInfoData(42); + hsa_ven_amd_aqlprofile_info_data_t callback_data{}; + callback_data.sample_id = 42; + callback_data.trace_data.ptr = nullptr; + callback_data.trace_data.size = 0; + + hsa_status_t status = DefaultTracedataCallback( + HSA_VEN_AMD_AQLPROFILE_INFO_TRACE_DATA, + &info_data, + &callback_data); + + EXPECT_EQ(status, HSA_STATUS_INFO_BREAK); + EXPECT_EQ(callback_data.trace_data.ptr, info_data.trace_data.ptr); + EXPECT_EQ(callback_data.trace_data.size, info_data.trace_data.size); +} + +// Test case: DefaultTracedataCallback with non-matching sample IDs +TEST_F(DefaultTracedataCallbackTest, NonMatchingSampleId) { + auto info_data = CreateInfoData(42); + hsa_ven_amd_aqlprofile_info_data_t callback_data{}; + callback_data.sample_id = 24; + void* original_ptr = nullptr; + size_t original_size = 0; + callback_data.trace_data.ptr = original_ptr; + callback_data.trace_data.size = original_size; + + hsa_status_t status = DefaultTracedataCallback( + HSA_VEN_AMD_AQLPROFILE_INFO_TRACE_DATA, + &info_data, + &callback_data); + + EXPECT_EQ(status, HSA_STATUS_SUCCESS); + EXPECT_EQ(callback_data.trace_data.ptr, original_ptr); + EXPECT_EQ(callback_data.trace_data.size, original_size); +} + +// Test case: DefaultTracedataCallback with non-trace info type +TEST_F(DefaultTracedataCallbackTest, NonTraceInfoType) { + auto info_data = CreateInfoData(42); + hsa_ven_amd_aqlprofile_info_data_t callback_data{}; + callback_data.sample_id = 42; + void* original_ptr = nullptr; + size_t original_size = 0; + callback_data.trace_data.ptr = original_ptr; + callback_data.trace_data.size = original_size; + + hsa_status_t status = DefaultTracedataCallback( + HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA, + &info_data, + &callback_data); + + EXPECT_EQ(status, HSA_STATUS_SUCCESS); + EXPECT_EQ(callback_data.trace_data.ptr, original_ptr); + EXPECT_EQ(callback_data.trace_data.size, original_size); +} diff --git a/projects/aqlprofile/src/core/tests/command_buffer_tests.cpp b/projects/aqlprofile/src/core/tests/command_buffer_tests.cpp new file mode 100644 index 0000000000..f86566da53 --- /dev/null +++ b/projects/aqlprofile/src/core/tests/command_buffer_tests.cpp @@ -0,0 +1,51 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include "core/commandbuffermgr.hpp" +#include + + +using namespace aql_profile; + +namespace { + +struct DummyBuffer { + std::array data; +}; + +TEST(CommandBufferMgrTest, BasicPrefixAndSize) { + DummyBuffer buf; + CommandBufferMgr mgr(buf.data.data(), sizeof(buf.data)); + // Should not throw and should have a nonzero size + EXPECT_GT(mgr.GetSize(), 0u); +} + +TEST(CommandBufferMgrTest, FinalizeThrowsOnSmallData) { + DummyBuffer buf; + CommandBufferMgr mgr(buf.data.data(), sizeof(buf.data)); + mgr.SetPreSize(128); + // Finalize with data_size <= precmds_size should throw + EXPECT_THROW(mgr.Finalize(64), aql_profile_exc_msg); +} + +} // namespace diff --git a/projects/aqlprofile/src/core/tests/counter_tests.cpp b/projects/aqlprofile/src/core/tests/counter_tests.cpp new file mode 100644 index 0000000000..9c4d6518f9 --- /dev/null +++ b/projects/aqlprofile/src/core/tests/counter_tests.cpp @@ -0,0 +1,118 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include +#include "core/include/aql_profile_v2.h" + +// Mocks and helpers +namespace { + +struct MockMemory { + std::vector data; + void* alloc(size_t size) { + data.resize(size); + return data.data(); + } + void dealloc(void* /*ptr*/) { + // No-op for vector-backed memory + } + void copy(void* dst, const void* src, size_t size) { + memcpy(dst, src, size); + } +}; + +// Corrected mock_alloc matching aqlprofile_memory_alloc_callback_t +hsa_status_t mock_alloc(void** ptr, uint64_t size, aqlprofile_buffer_desc_flags_t /*flags*/, void* userdata) { + auto* mem = static_cast(userdata); + *ptr = mem->alloc(size); + return HSA_STATUS_SUCCESS; +} + +// Corrected mock_dealloc matching aqlprofile_memory_dealloc_callback_t +void mock_dealloc(void* /*ptr*/, void* /*userdata*/) { + // No-op +} + +// Corrected mock_memcpy matching aqlprofile_memory_copy_t +hsa_status_t mock_memcpy(void* dst, const void* src, size_t size, void* /*userdata*/) { + memcpy(dst, src, size); + return HSA_STATUS_SUCCESS; +} + +} // namespace + +// Test: CreatePacketsSuccess +TEST(CountersTest, CreatePacketsSuccess) { + MockMemory mem; + aqlprofile_pmc_profile_t profile = {}; + // Fill profile with minimal valid data + profile.agent.handle = 0; // Use a valid agent handle in real test + profile.events = nullptr; + profile.event_count = 0; + + aqlprofile_handle_t handle = {}; + aqlprofile_pmc_aql_packets_t packets = {}; + + hsa_status_t status = aqlprofile_pmc_create_packets( + &handle, &packets, profile, mock_alloc, mock_dealloc, mock_memcpy, &mem); + + // Accept HSA_STATUS_ERROR if agent/event is not valid in this test context + EXPECT_TRUE(status == HSA_STATUS_SUCCESS || status == HSA_STATUS_ERROR); +} + +// Test: DeletePackets +TEST(CountersTest, DeletePackets) { + MockMemory mem; + aqlprofile_pmc_profile_t profile = {}; + profile.agent.handle = 0; + profile.events = nullptr; + profile.event_count = 0; + + aqlprofile_handle_t handle = {}; + aqlprofile_pmc_aql_packets_t packets = {}; + + hsa_status_t status = aqlprofile_pmc_create_packets( + &handle, &packets, profile, mock_alloc, mock_dealloc, mock_memcpy, &mem); + + // Only proceed if creation succeeded + if (status == HSA_STATUS_SUCCESS) { + // This should not crash or throw + aqlprofile_pmc_delete_packets(handle); + } +} + +// Test: ValidateEvent +TEST(CountersTest, ValidateEvent) { + aqlprofile_agent_handle_t agent = {}; + agent.handle = 0; + + aqlprofile_pmc_event_t event = {}; + event.block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GRBM; + + bool result = true; + hsa_status_t status = aqlprofile_validate_pmc_event(agent, &event, &result); + + // In a mock environment, we can't guarantee validation, but we can check that it runs + EXPECT_TRUE(status == HSA_STATUS_SUCCESS || status == HSA_STATUS_ERROR); +} \ No newline at end of file diff --git a/projects/aqlprofile/src/core/tests/memorymanager_tests.cpp b/projects/aqlprofile/src/core/tests/memorymanager_tests.cpp new file mode 100644 index 0000000000..c6fb97ec97 --- /dev/null +++ b/projects/aqlprofile/src/core/tests/memorymanager_tests.cpp @@ -0,0 +1,124 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "core/memorymanager.hpp" +#include "gtest/gtest.h" +#include + +// Dummy alloc/dealloc functions for testing +hsa_status_t dummy_alloc(void** ptr, size_t size, aqlprofile_buffer_desc_flags_t, void*) { + *ptr = malloc(size); + return *ptr ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR; +} +void dummy_dealloc(void* ptr, void*) { + free(ptr); +} +hsa_status_t dummy_alloc_fail(void**, size_t, aqlprofile_buffer_desc_flags_t, void*) { + return HSA_STATUS_ERROR; +} +hsa_status_t dummy_copy(void* dst, const void* src, size_t size, void*) { + memcpy(dst, src, size); + return HSA_STATUS_SUCCESS; +} + +// Helper subclass to expose protected AllocMemory for testing +class TestCounterMemoryManager : public CounterMemoryManager { +public: + using CounterMemoryManager::CounterMemoryManager; + using CounterMemoryManager::AllocMemory; +}; + +TEST(CounterMemoryManagerTest, AllocMemory_Success) { + hsa_agent_t agent = {.handle = 1}; + TestCounterMemoryManager mgr(agent, dummy_alloc, dummy_dealloc, nullptr); + aqlprofile_buffer_desc_flags_t flags{}; + auto mem = mgr.AllocMemory(64, flags); + ASSERT_NE(mem.get(), nullptr); +} + +TEST(CounterMemoryManagerTest, AllocMemory_FailureThrows) { + hsa_agent_t agent = {.handle = 1}; + TestCounterMemoryManager mgr(agent, dummy_alloc_fail, dummy_dealloc, nullptr); + aqlprofile_buffer_desc_flags_t flags{}; + EXPECT_THROW(mgr.AllocMemory(64, flags), hsa_status_t); +} + +TEST(CounterMemoryManagerTest, CmdBufAndOutputBuf) { + hsa_agent_t agent = {.handle = 1}; + CounterMemoryManager mgr(agent, dummy_alloc, dummy_dealloc, nullptr); + mgr.CreateCmdBuf(128); + ASSERT_NE(mgr.GetCmdBuf(), nullptr); + mgr.CreateOutputBuf(256); + ASSERT_NE(mgr.GetOutputBuf(), nullptr); + ASSERT_EQ(mgr.GetOutputBufSize(), 256u); +} + +TEST(CounterMemoryManagerTest, RegisterAndGetManager) { + hsa_agent_t agent = {.handle = 1}; + auto mgr = std::make_shared(agent, dummy_alloc, dummy_dealloc, nullptr); + size_t handle = mgr->GetHandler(); + CounterMemoryManager::RegisterManager(mgr); + auto found = CounterMemoryManager::GetManager(handle); + ASSERT_EQ(found.get(), mgr.get()); + CounterMemoryManager::DeleteManager(handle); + ASSERT_EQ(CounterMemoryManager::GetManager(handle), nullptr); +} + +TEST(CounterMemoryManagerTest, CopyEvents) { + hsa_agent_t agent = {.handle = 1}; + CounterMemoryManager mgr(agent, dummy_alloc, dummy_dealloc, nullptr); + + aqlprofile_pmc_event_t events[2]{}; + events[0].event_id = 1; + events[0].flags.raw = 0; + events[1].event_id = 2; + events[1].flags.raw = 1; + + mgr.CopyEvents(events, 2); + auto& ev = mgr.GetEvents(); + ASSERT_GE(ev.size(), 2u); +} + +TEST(TraceMemoryManagerTest, TraceControlBufAndATTParams) { + hsa_agent_t agent = {.handle = 1}; + TraceMemoryManager mgr(agent, dummy_alloc, dummy_dealloc, dummy_copy, nullptr); + mgr.CreateTraceControlBuf(64); + auto buf = mgr.GetTraceControlBuf(); + ASSERT_NE(buf, nullptr); + + hsa_ven_amd_aqlprofile_parameter_t params[2]{}; + params[0].parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET; + params[0].value = 42; + params[1].parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SIMD_SELECTION; + params[1].value = 0xF; + + mgr.CopyATTParams(params, 2); + ASSERT_EQ(mgr.GetSimdMask(), 0xF); + const auto& att_params = mgr.GetATTParams(); + ASSERT_EQ(att_params.size(), 2u); +} + +TEST(CodeobjMemoryManagerTest, CmdBufferAlloc) { + hsa_agent_t agent = {.handle = 1}; + CodeobjMemoryManager mgr(agent, dummy_alloc, dummy_dealloc, 128, nullptr); + ASSERT_NE(mgr.cmd_buffer.get(), nullptr); +} \ No newline at end of file diff --git a/projects/aqlprofile/src/pm4/tests/CMakeLists.txt b/projects/aqlprofile/src/pm4/tests/CMakeLists.txt new file mode 100644 index 0000000000..afd41700c8 --- /dev/null +++ b/projects/aqlprofile/src/pm4/tests/CMakeLists.txt @@ -0,0 +1,28 @@ +include(GoogleTest) +find_package(GTest REQUIRED) +include_directories(${GTEST_INCLUDE_DIRS}) + + +# Add a test for gfx9 command builder +add_executable(gfx9-command-builder-test) +SET(AQLPROFILE_COMMAND_BUILDER_SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/cmd_builder_tests.cpp +) +target_sources(gfx9-command-builder-test PRIVATE ${AQLPROFILE_COMMAND_BUILDER_SOURCES}) +target_include_directories(gfx9-command-builder-test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${LIB_DIR}) +target_link_libraries( + gfx9-command-builder-test + PRIVATE + hsa-runtime64::hsa-runtime64 + GTest::gtest + GTest::gtest_main + GTest::gmock + GTest::gmock_main) +gtest_add_tests( + TARGET gfx9-command-builder-test + SOURCES ${AQLPROFILE_COMMAND_BUILDER_SOURCES} + TEST_LIST gfx9-command-builder-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) +set_tests_properties( + ${gfx9-command-builder-test_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION + "${AQLPROFILE_DEFAULT_FAIL_REGEX}") \ No newline at end of file diff --git a/projects/aqlprofile/src/pm4/tests/cmd_builder_tests.cpp b/projects/aqlprofile/src/pm4/tests/cmd_builder_tests.cpp new file mode 100644 index 0000000000..0741aacfc8 --- /dev/null +++ b/projects/aqlprofile/src/pm4/tests/cmd_builder_tests.cpp @@ -0,0 +1,71 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include "pm4/cmd_builder.h" + +using pm4_builder::CmdBuffer; + +struct DummyPacket { + uint32_t a; + uint32_t b; +}; + +TEST(CmdBufferTest, AppendSinglePacket) { + CmdBuffer buf; + DummyPacket pkt{0x12345678, 0x9abcdef0}; + buf.Append(pkt); + + // Should have 2 dwords + EXPECT_EQ(buf.DwSize(), 2u); + + // Data should match what we appended + const uint32_t* data = static_cast(buf.Data()); + EXPECT_EQ(data[0], 0x12345678u); + EXPECT_EQ(data[1], 0x9abcdef0u); +} + +TEST(CmdBufferTest, AppendMultiplePackets) { + CmdBuffer buf; + DummyPacket pkt1{1, 2}; + DummyPacket pkt2{3, 4}; + buf.Append(pkt1, pkt2); + + EXPECT_EQ(buf.DwSize(), 4u); + const uint32_t* data = static_cast(buf.Data()); + EXPECT_EQ(data[0], 1u); + EXPECT_EQ(data[1], 2u); + EXPECT_EQ(data[2], 3u); + EXPECT_EQ(data[3], 4u); +} + +TEST(CmdBufferTest, AppendRawData) { + CmdBuffer buf; + uint32_t raw[3] = {10, 20, 30}; + buf.Append(raw, 3); + + EXPECT_EQ(buf.DwSize(), 4u); + const uint32_t* data = static_cast(buf.Data()); + EXPECT_EQ(data[0], 10u); + EXPECT_EQ(data[1], 20u); + EXPECT_EQ(data[2], 30u); +} \ No newline at end of file diff --git a/projects/aqlprofile/test/CMakeLists.txt b/projects/aqlprofile/test/CMakeLists.txt index 70f3001e58..26a431eda3 100644 --- a/projects/aqlprofile/test/CMakeLists.txt +++ b/projects/aqlprofile/test/CMakeLists.txt @@ -82,3 +82,8 @@ install(FILES ${TEST_DIR}/run_install.sh RENAME run_tests.sh #add_subdirectory(parser) add_test(NAME legacy-tests COMMAND "${TEST_BINARY_DIR}/run.sh") + +if(AQLPROFILE_BUILD_TESTS) + add_subdirectory(integration) + enable_testing() +endif() diff --git a/projects/aqlprofile/test/integration/CMakeLists.txt b/projects/aqlprofile/test/integration/CMakeLists.txt new file mode 100644 index 0000000000..2df0ecbc93 --- /dev/null +++ b/projects/aqlprofile/test/integration/CMakeLists.txt @@ -0,0 +1,67 @@ +cmake_minimum_required(VERSION 3.21 FATAL_ERROR) +project(Aqlprofile_v2_tests LANGUAGES C CXX HIP) +set(CMAKE_CXX_STANDARD 17) + + +find_library( + hsa-amd-aqlprofile64 + REQUIRED + NAMES hsa-amd-aqlprofile64 hsa-amd-aqlprofile64 + HINTS /opt/rocm/lib /opt/rocm + PATHS /opt/rocm/lib /opt/rocm) + +find_package( + hsa-runtime64 + REQUIRED + CONFIG + HINTS /opt/rocm/ + PATHS /opt/rocm/) + +get_property( + HSA_RUNTIME_INCLUDE_DIRECTORIES + TARGET hsa-runtime64::hsa-runtime64 + PROPERTY INTERFACE_INCLUDE_DIRECTORIES) +find_file( + HSA_H hsa.h + PATHS ${HSA_RUNTIME_INCLUDE_DIRECTORIES} + PATH_SUFFIXES hsa + NO_DEFAULT_PATH REQUIRED) +get_filename_component(HSA_RUNTIME_INC_PATH ${HSA_H} DIRECTORY) +include_directories(${HSA_RUNTIME_INC_PATH}) + +if(NOT CMAKE_HIP_COMPILER) + find_program( + amdclangpp_EXECUTABLE + NAMES amdclang++ + HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATH_SUFFIXES bin llvm/bin NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP) + +set(CMAKE_BUILD_TYPE Debug) +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -O1" ) + +add_library(intercept SHARED) +target_include_directories(intercept PRIVATE ${HSA_RUNTIME_INC_PATH} ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/src/core/include/ /opt/rocm/include) +target_sources(intercept PRIVATE intercept.cpp) +target_link_libraries(intercept PRIVATE hsa-runtime64::hsa-runtime64 ${hsa-amd-aqlprofile64}) +target_link_options(intercept PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exportmap -Wl,--no-undefined) +target_compile_definitions(intercept PUBLIC AMD_INTERNAL_BUILD) +add_executable(testv2) +target_sources(testv2 PRIVATE main.cpp workload.cpp counter.cpp agent.cpp) +target_include_directories(testv2 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/src/core/include/ ${HSA_RUNTIME_INC_PATH} /opt/rocm/include) +target_link_libraries(testv2 PRIVATE hsa-runtime64::hsa-runtime64 ${hsa-amd-aqlprofile64}) +target_compile_definitions(testv2 PUBLIC AMD_INTERNAL_BUILD) + +# Add a PRELOAD environment with libintercept +set(ENV{LD_PRELOAD} "$ENV{LD_PRELOAD}:${CMAKE_CURRENT_BINARY_DIR}/libintercept.so") + +add_test(NAME testv2 COMMAND testv2) +set_tests_properties(testv2 PROPERTIES ENVIRONMENT "${LD_PRELOAD}" TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION "${AQLPROFILE_DEFAULT_FAIL_REGEX}") diff --git a/projects/aqlprofile/test/integration/agent.cpp b/projects/aqlprofile/test/integration/agent.cpp new file mode 100644 index 0000000000..d149c3f333 --- /dev/null +++ b/projects/aqlprofile/test/integration/agent.cpp @@ -0,0 +1,235 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "agent.hpp" +#include + +#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); } + +std::vector> AgentInfo::gpu_agents{}; +hsa_agent_t AgentInfo::cpu_agent{0}; + +hsa_amd_memory_pool_t AgentInfo::cpu_pool; +hsa_amd_memory_pool_t AgentInfo::kernarg_pool; + +void AgentInfo::add_event(aqlprofile_pmc_event_t block, const std::string& counter, int block_cnt, int event_id) +{ + block.event_id = event_id; + std::vector cnt{}; + for (int i=0; i info = std::make_shared(); + info->hsa_agent = agent; + CHECK_HSA(hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, info->gfxip.data())); + CHECK_HSA(hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_NUM_XCC), &info->info.xcc_num)); + CHECK_HSA(hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES), &info->info.se_num)); + CHECK_HSA(hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT), &info->info.cu_num)); + CHECK_HSA(hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE), &info->info.shader_arrays_per_se)); + + info->info.agent_gfxip = info->gfxip.data(); + CHECK_HSA(aqlprofile_register_agent(&info->handle, &info->info)); + + aqlprofile_pmc_event_flags_t flags{.raw = 0}; + aqlprofile_pmc_event_t grbm {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GRBM}; + aqlprofile_pmc_event_t sq {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ}; + aqlprofile_pmc_event_t ta {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TA}; + aqlprofile_pmc_event_t tcp {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP}; + aqlprofile_pmc_event_t tcc {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC}; + aqlprofile_pmc_event_t gl2c {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2C}; + + info->add_event(grbm, "GRBM_COUNT", 1, 0); + info->add_event(grbm, "GRBM_GUI_ACTIVE", 1, 2); + info->add_event(sq, "SQ_WAVES", 1, 4); + info->add_event(sq, "SQ_BUSY_CYCLES", 1, 3); + info->add_event(sq, "SQ_INSTS_VALU", 1, (info->gfxip.find("gfx1")==0) ? 62 : 26); + + info->add_event(ta, "TA_BUSY", 16, (info->gfxip.find("gfx94") != 0 || info->gfxip.find("gfx95") != 0) ? 13 : 15); + + if (info->gfxip.find("gfx1") == 0) + { + info->add_event(gl2c, "GL2C_REQ", 32, 3); + info->add_event(gl2c, "GL2C_READ", 32, 6); + info->add_event(gl2c, "GL2C_WRITE", 32, 7); + } + else if (info->gfxip.find("gfx95") == 0) + { + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP16", 10, 81); + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP32", 10, 82); + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP64", 10, 83); + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP16_TRANS", 10, 84); + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP32_TRANS", 10, 85); + info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP64_TRANS", 10, 86); + + info->add_event(tcp, "TCP_READ", 10, 28); + info->add_event(tcp, "TCP_WRITE", 10, 30); + info->add_event(tcp, "TCP_CACHE_ACCESS", 10, 58); + info->add_event(tcp, "TCP_CACHE_MISS_TG0", 10, 59); + info->add_event(tcp, "TCP_CACHE_MISS_TG1", 10, 60); + info->add_event(tcp, "TCP_CACHE_MISS_TG2", 10, 61); + info->add_event(tcp, "TCP_CACHE_MISS_TG3", 10, 62); + info->add_event(tcp, "TCP_CACHE_MISS", 10, 63); + + info->add_event(tcc, "TCC_EA0_RDREQ", 16, 42); + info->add_event(tcc, "TCC_EA0_RDREQ_DRAM", 16, 108); + info->add_event(tcc, "TCC_EA0_WRREQ_DRAM", 16, 109); + info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_DRAM", 16, 110); + info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_DRAM", 16, 111); + info->add_event(tcc, "TCC_EA0_RDREQ_DRAM_32B", 16, 112); + info->add_event(tcc, "TCC_EA0_RDREQ_GMI_32B", 16, 113); + info->add_event(tcc, "TCC_EA0_RDREQ_IO_32B", 16, 114); + info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_DRAM_32B", 16, 115); + info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_DRAM_32B", 16, 116); + info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_GMI_32B", 16, 117); + info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_GMI_32B", 16, 118); + info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_IO_32B", 16, 119); + info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_IO_32B", 16, 119); + } + else if (info->gfxip.find("gfx94") == 0) + { + info->add_event(tcc, "TCC_REQ", 16, 3); + info->add_event(tcc, "TCC_ATOMIC", 16, 14); + info->add_event(tcc, "TCC_EA0_ATOMIC", 16, 36); + + info->add_event(tcc, "TCC_EA0_WRREQ_CREDIT_STALL", 16, 30); + info->add_event(tcc, "TCC_EA0_WRREQ_IO_CREDIT_STALL", 16, 31); + info->add_event(tcc, "TCC_EA0_WRREQ_GMI_CREDIT_STALL", 16, 32); + info->add_event(tcc, "TCC_EA0_WRREQ_DRAM_CREDIT_STALL", 16, 33); + info->add_event(tcc, "TCC_EA0_RDREQ", 16, 38); + info->add_event(tcc, "TCC_EA0_RDREQ_IO_CREDIT_STALL", 16, 41); + info->add_event(tcc, "TCC_EA0_RDREQ_GMI_CREDIT_STALL", 16, 42); + info->add_event(tcc, "TCC_EA0_RDREQ_DRAM_CREDIT_STALL", 16, 43); + + info->add_event(tcp, "TCP_READ", 10, 28); + info->add_event(tcp, "TCP_WRITE", 10, 30); + info->add_event(tcp, "TCP_CACHE_ACCESS", 10, 60); + info->add_event(tcp, "TCP_CACHE_MISS_TG0", 10, 61); + info->add_event(tcp, "TCP_CACHE_MISS_TG1", 10, 62); + info->add_event(tcp, "TCP_CACHE_MISS_TG2", 10, 63); + info->add_event(tcp, "TCP_CACHE_MISS_TG3", 10, 64); + } + else if (info->gfxip.find("gfx90a") == 0) + { + info->add_event(tcp, "TCP_READ", 16, 30); + info->add_event(tcp, "TCP_WRITE", 16, 32); + } + else if (info->gfxip.find("gfx900") == 0) + { + info->add_event(tcp, "TCP_READ", 16, 30); + info->add_event(tcp, "TCP_WRITE", 16, 32); + } + else + { + assert(false); + } + + gpu_agents.push_back(info); + return HSA_STATUS_SUCCESS; +} + +hsa_status_t +FindGlobalPool(hsa_amd_memory_pool_t pool, void* data) +{ + hsa_amd_segment_t segment; + CHECK_HSA(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment)); + + if(HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS; + + uint32_t flag; + CHECK_HSA(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag)); + uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT; + + if (karg_st) + AgentInfo::kernarg_pool = pool; + else + AgentInfo::cpu_pool = pool; + + return HSA_STATUS_SUCCESS; +} + +void AgentInfo::iterate_agents() +{ + CHECK_HSA(hsa_iterate_agents(get_agent_handle_cb, nullptr)); + CHECK_HSA(hsa_amd_agent_iterate_memory_pools(cpu_agent, FindGlobalPool, nullptr)); +} + +bool +Queue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet) +{ + const uint64_t write_idx = hsa_queue_add_write_index_relaxed(queue, 1); + + size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t); + auto* queue_slot = reinterpret_cast(size_t(queue->base_address) + index); // NOLINT + + const auto* slot_data = reinterpret_cast(packet); + + std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t)); + auto* header = reinterpret_cast*>(queue_slot); + + header->store(slot_data[0], std::memory_order_release); + hsa_signal_store_screlease(queue->doorbell_signal, write_idx); + + int loops = 0; + while(hsa_queue_load_read_index_relaxed(queue) <= write_idx) + { + loops++; + usleep(1); + if(loops > 10000) + { + std::cerr << "Codeobj packet submission failed!" << std::endl; + return false; + } + } + return true; +} + +Queue::Queue(std::shared_ptr& _agent): agent(_agent) +{ + CHECK_HSA(hsa_queue_create(agent->hsa_agent, 64, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &this->queue)); +} + +void Queue::flush() +{ + return; + hsa_barrier_and_packet_t barrier{}; + barrier.header = HSA_PACKET_TYPE_BARRIER_OR | (1< +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +//#include "/opt/rocm/include/aqlprofile-sdk/aql_profile_v2.h" +#include "aql_profile_v2.h" + +#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); } + +inline bool operator==(const aqlprofile_pmc_event_t& a, const aqlprofile_pmc_event_t& b) { + return a.event_id == b.event_id && a.block_name == b.block_name; +} + +inline bool operator<(const aqlprofile_pmc_event_t& a, const aqlprofile_pmc_event_t& b) { + return (a.block_name == b.block_name) ? (a.event_id < b.event_id) : (a.block_name < b.block_name); +} + +struct AgentInfo +{ + AgentInfo() { gfxip.resize(64); } + + std::string gfxip; + hsa_agent_t hsa_agent; + aqlprofile_agent_info_t info; + aqlprofile_agent_handle_t handle; + std::map> counters; + + static void iterate_agents(); + + static std::vector> gpu_agents; + static hsa_agent_t cpu_agent; + static hsa_amd_memory_pool_t cpu_pool; + static hsa_amd_memory_pool_t kernarg_pool; + +private: + void add_event(aqlprofile_pmc_event_t block, const std::string& counter, int block_cnt, int event_id); + + static hsa_status_t get_agent_handle_cb(hsa_agent_t agent, void* userdata); +}; + +class Queue +{ +public: + Queue(std::shared_ptr& _agent); + ~Queue() { if (queue) hsa_queue_destroy(queue); } + + bool Submit(hsa_ext_amd_aql_pm4_packet_t* packet); + void flush(); + + hsa_queue_t* queue = nullptr; + std::shared_ptr agent; +}; \ No newline at end of file diff --git a/projects/aqlprofile/test/integration/counter.cpp b/projects/aqlprofile/test/integration/counter.cpp new file mode 100644 index 0000000000..03f32a8835 --- /dev/null +++ b/projects/aqlprofile/test/integration/counter.cpp @@ -0,0 +1,127 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include "counter.hpp" +#include + +#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); } + +hsa_status_t data_callback( + aqlprofile_pmc_event_t event, + uint64_t counter_id, + uint64_t counter_value, + void* userdata +) { + auto* packet = static_cast(userdata); + try { + packet->results.at(event) += counter_value; + } catch(...) { abort(); } + return HSA_STATUS_SUCCESS; +} + +hsa_status_t +AQLPacket::Alloc(void** ptr, size_t size, desc_t flags, void* data) +{ + auto* packet = reinterpret_cast(data); + assert(packet && "Invalid aql packet"); + if (flags.memory_hint != AQLPROFILE_MEMORY_HINT_DEVICE_UNCACHED) + { + CHECK_HSA(hsa_amd_memory_pool_allocate(AgentInfo::cpu_pool, size, + HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, ptr)); + CHECK_HSA(hsa_amd_memory_fill(*ptr, 0u, size / sizeof(uint32_t))); + return hsa_amd_agents_allow_access(1, &packet->hsa_agent, nullptr, *ptr); + } + else + { + CHECK_HSA(hsa_amd_memory_pool_allocate(AgentInfo::kernarg_pool, size, + HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, ptr)); + CHECK_HSA(hsa_amd_memory_fill(*ptr, 0u, size / sizeof(uint32_t))); + return hsa_amd_agents_allow_access(1, &packet->hsa_agent, nullptr, *ptr); + } +} + +void AQLPacket::Free(void* ptr, void* data) +{ + if(ptr == nullptr) return; + hsa_amd_memory_pool_free(ptr); +} + +hsa_status_t AQLPacket::Copy(void* dst, const void* src, size_t size, void* data) +{ + if(size == 0) return HSA_STATUS_SUCCESS; + return hsa_memory_copy(dst, src, size); +} + +AQLPacket::AQLPacket( + AgentInfo& agent, + const std::vector& _counters +): hsa_agent(agent.hsa_agent) { + constexpr hsa_ext_amd_aql_pm4_packet_t null_amd_aql_pm4_packet = + { + .header = 0, + .pm4_command = {0}, + .completion_signal = {.handle = 0} + }; + + packets.start_packet = null_amd_aql_pm4_packet; + packets.stop_packet = null_amd_aql_pm4_packet; + packets.read_packet = null_amd_aql_pm4_packet; + + aqlprofile_pmc_profile_t profile{}; + std::vector events; + for (auto& counter : _counters) + { + auto& event = agent.counters.at(counter).at(0); + results[event] = 0; + prev_results[event] = 0; + counter_names[event] = counter; + for (auto& ev : agent.counters.at(counter)) + events.push_back(ev); + } + + profile.agent = agent.handle; + profile.events = events.data(); + profile.event_count = static_cast(events.size()); + + CHECK_HSA(aqlprofile_pmc_create_packets(&this->handle, + &this->packets, + profile, + &AQLPacket::Alloc, + &AQLPacket::Free, + &AQLPacket::Copy, + this)); + + packets.start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + packets.stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + packets.read_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; +} + +void AQLPacket::iterate() +{ + for (auto& [key, value] : results) + { + prev_results[key] = value; + results[key] = 0; + } + CHECK_HSA(aqlprofile_pmc_iterate_data(this->handle, data_callback, this)); +} diff --git a/projects/aqlprofile/test/integration/counter.hpp b/projects/aqlprofile/test/integration/counter.hpp new file mode 100644 index 0000000000..ea938dfb2c --- /dev/null +++ b/projects/aqlprofile/test/integration/counter.hpp @@ -0,0 +1,72 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include +#include +#include +#include +#include "aql_profile_v2.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "agent.hpp" + +class AQLPacket +{ + using desc_t = aqlprofile_buffer_desc_flags_t; +public: + AQLPacket(AgentInfo& _agent, const std::vector& counters); + ~AQLPacket() { aqlprofile_pmc_delete_packets(this->handle); }; + + void iterate(); + + static hsa_status_t Alloc(void** ptr, size_t size, desc_t flags, void* data); + static void Free(void* ptr, void* data); + static hsa_status_t Copy(void* dst, const void* src, size_t size, void* data); + + std::map get() + { + std::map ret; + for (auto& [event, counter] : counter_names) + ret.emplace(counter, results.at(event)); + return ret; + } + + std::map counter_names; + std::map results; + std::map prev_results; + + aqlprofile_handle_t handle{0}; + hsa_agent_t hsa_agent; + aqlprofile_pmc_aql_packets_t packets; + + bool delta = false; +}; diff --git a/projects/aqlprofile/test/integration/exportmap b/projects/aqlprofile/test/integration/exportmap new file mode 100644 index 0000000000..52a99830fb --- /dev/null +++ b/projects/aqlprofile/test/integration/exportmap @@ -0,0 +1,4 @@ +{ + global: HSA_AMD_TOOL_PRIORITY; OnLoad; + local: *; +}; \ No newline at end of file diff --git a/projects/aqlprofile/test/integration/intercept.cpp b/projects/aqlprofile/test/integration/intercept.cpp new file mode 100644 index 0000000000..c5e6a3aa6e --- /dev/null +++ b/projects/aqlprofile/test/integration/intercept.cpp @@ -0,0 +1,168 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define CHECK_HSA(x) { auto _status = (x); if (_status != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << ':' << __LINE__ << std::endl; abort(); } } + +extern "C" const uint32_t HSA_AMD_TOOL_PRIORITY = 25; +decltype(hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled_fn = nullptr; +decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn = nullptr; +decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn = nullptr; +decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn = nullptr; +decltype(hsa_signal_store_screlease)* hsa_signal_store_screlease_fn = nullptr; +decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn = nullptr; +decltype(hsa_queue_add_write_index_relaxed)* hsa_queue_add_write_index_relaxed_fn = nullptr; + +decltype(hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info_fn = nullptr; +decltype(hsa_agent_get_info)* hsa_agent_get_info_fn = nullptr; +decltype(hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools_fn = nullptr; +decltype(hsa_queue_create)* hsa_queue_create_fn = nullptr; + +hsa_amd_memory_pool_t cpu_pool; + +hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data) +{ + hsa_amd_segment_t segment; + CHECK_HSA(hsa_amd_memory_pool_get_info_fn(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment)); + + if(HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS; + + uint32_t flag; + CHECK_HSA(hsa_amd_memory_pool_get_info_fn(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag)); + uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT; + + if (karg_st == 0) cpu_pool = pool; + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t iterate_agent_cb(hsa_agent_t agent, void* userdata) +{ + hsa_device_type_t type; + + CHECK_HSA(hsa_agent_get_info_fn(agent, HSA_AGENT_INFO_DEVICE, &type)); + if (type != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS; + + CHECK_HSA(hsa_amd_agent_iterate_memory_pools_fn(agent, FindGlobalPool, nullptr)); + return HSA_STATUS_SUCCESS; +} + +bool queue_submit(hsa_queue_t* queue, hsa_ext_amd_aql_pm4_packet_t* packet) +{ + const uint64_t write_idx = hsa_queue_add_write_index_relaxed_fn(queue, 1); + + size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t); + auto* queue_slot = reinterpret_cast(size_t(queue->base_address) + index); // NOLINT + + const auto* slot_data = reinterpret_cast(packet); + + std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t)); + auto* header = reinterpret_cast*>(queue_slot); + + header->store(slot_data[0], std::memory_order_release); + hsa_signal_store_screlease_fn(queue->doorbell_signal, write_idx); + + int loops = 0; + while(hsa_queue_load_read_index_relaxed_fn(queue) <= write_idx) + { + loops++; + usleep(1); + if(loops > 10000) + { + std::cerr << "Packet submission failed!" << std::endl; + return false; + } + } + return true; +} + +void set_profiler_active_on_queue( + hsa_agent_t hsa_agent, + hsa_queue_t* queue +) { + hsa_ext_amd_aql_pm4_packet_t packet{}; + hsa_ven_amd_aqlprofile_profile_t profile{}; + profile.agent = hsa_agent; + + // Query for cmd buffer size + CHECK_HSA(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD, nullptr)); + + // Allocate cmd buffer + const size_t mask = 0x1000 - 1; + auto size = (profile.command_buffer.size + mask) & ~mask; + + CHECK_HSA(hsa_amd_memory_pool_allocate_fn(cpu_pool, size, + HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, &profile.command_buffer.ptr)); + CHECK_HSA(hsa_amd_agents_allow_access_fn(1, &hsa_agent, nullptr, profile.command_buffer.ptr)); + + CHECK_HSA(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD, &packet)); + + queue_submit(queue, &packet); + hsa_amd_memory_pool_free_fn(profile.command_buffer.ptr); +} + + +hsa_status_t QueueCreateInterceptor(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type, + void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), + void* data, uint32_t private_segment_size, + uint32_t group_segment_size, hsa_queue_t** queue) +{ + CHECK_HSA(hsa_queue_create_fn(agent, size, type, callback, data, private_segment_size, group_segment_size, queue)); + //CHECK_HSA(hsa_amd_profiling_set_profiler_enabled_fn(*queue, true)); + set_profiler_active_on_queue(agent, *queue); + return HSA_STATUS_SUCCESS; +} + +extern "C" __attribute__((visibility("default"))) bool +OnLoad(HsaApiTable* table, uint64_t, uint64_t, const char* const*) +{ + hsa_queue_create_fn = table->core_->hsa_queue_create_fn; + // Install the Queue intercept + table->core_->hsa_queue_create_fn = QueueCreateInterceptor; + hsa_amd_profiling_set_profiler_enabled_fn = table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn; + hsa_amd_memory_pool_allocate_fn = table->amd_ext_->hsa_amd_memory_pool_allocate_fn; + hsa_amd_agents_allow_access_fn = table->amd_ext_->hsa_amd_agents_allow_access_fn; + hsa_amd_memory_pool_free_fn = table->amd_ext_->hsa_amd_memory_pool_free_fn; + hsa_signal_store_screlease_fn = table->core_->hsa_signal_store_screlease_fn; + hsa_queue_load_read_index_relaxed_fn = table->core_->hsa_queue_load_read_index_relaxed_fn; + hsa_queue_add_write_index_relaxed_fn = table->core_->hsa_queue_add_write_index_relaxed_fn; + hsa_amd_memory_pool_get_info_fn = table->amd_ext_->hsa_amd_memory_pool_get_info_fn; + hsa_amd_agent_iterate_memory_pools_fn = table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn; + hsa_agent_get_info_fn = table->core_->hsa_agent_get_info_fn; + + CHECK_HSA(table->core_->hsa_iterate_agents_fn(iterate_agent_cb, nullptr)); + + return true; +} diff --git a/projects/aqlprofile/test/integration/main.cpp b/projects/aqlprofile/test/integration/main.cpp new file mode 100644 index 0000000000..3dc903d6f5 --- /dev/null +++ b/projects/aqlprofile/test/integration/main.cpp @@ -0,0 +1,419 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "counter.hpp" +#include "workload.hpp" +#include "hip/hip_runtime.h" + +#define DATA_SIZE (64*4) + +#define HIP_API_CALL(CALL) do { if ((CALL) != hipSuccess) abort(); } while(0) + + +//#define ATTEMPT_GMI +#define DOT2_ARCH + +class hipMemory +{ +public: + hipMemory(size_t size) + { + HIP_API_CALL(hipMalloc(&ptr, size * sizeof(float))); + HIP_API_CALL(hipMemset(ptr, 0, size * sizeof(float))); + } + ~hipMemory() + { + if(ptr) HIP_API_CALL(hipFree(ptr)); + } + + hipMemory(hipMemory& other) = delete; + hipMemory& operator=(hipMemory& other) = delete; + + float* ptr = nullptr; +}; + +class Stream +{ +public: + Stream() { HIP_API_CALL(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); } + ~Stream() { HIP_API_CALL(hipStreamDestroy(stream)); } + + Stream(Stream& other) = delete; + Stream& operator=(Stream& other) = delete; + + void synchronize() { HIP_API_CALL(hipStreamSynchronize(stream)); } + + hipStream_t stream; +}; + +class HIPWorkload : public IWorkload +{ +public: + HIPWorkload(AgentInfo& agent, const std::vector& counters) + { + col = std::make_unique(agent, counters); + } + virtual ~HIPWorkload() {}; + virtual std::string_view name() = 0; + + std::map collect(Queue& queue) + { + assert(col); + return col->iterate(queue, *this); + } + + void printcounters(Queue& queue) + { + std::cout << "Name: " << name() << std::endl; + for (auto& [name, v] : collect(queue)) std::cout << " - " << name << ": " << v << std::endl; + } + + std::unique_ptr col{nullptr}; + hipMemory src{DATA_SIZE}; + hipMemory dst{DATA_SIZE}; + Stream stream{}; +}; + + +__global__ void copy_kernel(float* a, const float* b) +{ + int idx = threadIdx.x + blockIdx.x*blockDim.x; + if (idx < DATA_SIZE) + a[idx] = b[idx]; +} + +__global__ void atomic_kernel(float* a, const float* b) +{ + int idx = threadIdx.x + blockIdx.x*blockDim.x; + if (idx < DATA_SIZE) + atomicAdd(a+threadIdx.x, b[idx]); +} + +__global__ void iops_kernel_trans() +{ + // 3 F16 Trans OPS + asm volatile("v_cos_f16 v0, v0; v_cos_f16 v1, v1; v_cos_f16 v2, v2;"); + // 2 F32 Trans OPS + asm volatile("v_cos_f32 v3, v3; v_cos_f32 v4, v4"); +} + +__global__ void iops_kernel1() +{ + asm volatile("v_add_f16 v2, v1, v0"); // 1 F16 OPS + asm volatile("v_fma_f32 v3, v1, v2, v3"); // 2 F32 OPs + + asm volatile("v_add_f64 v[0:1], v[2:3], v[4:5]"); // 1 F64 OP + asm volatile("v_fma_f64 v[0:1], v[2:3], v[4:5], v[6:7]"); // 2 F64 OP + asm volatile("v_fma_f64 v[0:1], v[2:3], v[4:5], v[6:7]"); // 2 F64 OP +} + +__global__ void iops_kernel2() +{ + #if defined(__gfx940__) || defined(__gfx90a__) || defined(__gfx1030__) + // Supported architectures + asm volatile("v_dot2_f32_f16 v0, v1, v2, v3"); +#else + // Fallback or skip + asm volatile("v_add_f32 v4, v5, v6"); // 1 F32 OP + asm volatile("v_fma_f64 v[0:1], v[0:1], v[2:3], v[4:5]"); // 2 F64 OPs +#endif + +} + +class CopyWorkload : public HIPWorkload +{ +public: + CopyWorkload(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + copy_kernel<<>>(dst.ptr, src.ptr); + stream.synchronize(); + } + virtual std::string_view name() override { return "CopyWorkload"; }; +}; + +class AtomicWorkload : public HIPWorkload +{ +public: + AtomicWorkload(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + atomic_kernel<<>>(dst.ptr, src.ptr); + stream.synchronize(); + } + virtual std::string_view name() override { return "AtomicWorkload"; }; +}; + +class IOPSWorkload1 : public HIPWorkload +{ +public: + IOPSWorkload1(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + iops_kernel1<<>>(); + stream.synchronize(); + } + virtual std::string_view name() override { return "IOPSWorkload1"; }; +}; + +class IOPSWorkload2 : public HIPWorkload +{ +public: + IOPSWorkload2(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + iops_kernel2<<>>(); + stream.synchronize(); + } + virtual std::string_view name() override { return "IOPSWorkload2"; }; +}; + +class IOPSWorkload3 : public HIPWorkload +{ +public: + IOPSWorkload3(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + iops_kernel_trans<<>>(); + stream.synchronize(); + } + virtual std::string_view name() override { return "Trans IOPSWorkload"; }; +}; + +class GMIWorkload : public HIPWorkload +{ +public: + GMIWorkload(AgentInfo& agent, const std::vector& counters): HIPWorkload(agent, counters) {} + virtual void run() override + { + auto policies = std::vector{hipHostMallocDefault, hipHostMallocCoherent, hipHostMallocNonCoherent}; + for (auto& policy : policies) + { + float* srchost; + float* dsthost; + HIP_API_CALL(hipHostMalloc(&srchost, DATA_SIZE * sizeof(float), policy)); + HIP_API_CALL(hipHostMalloc(&dsthost, DATA_SIZE * sizeof(float), policy)); + + for (size_t i=0; i>>(dsthost, srchost); + stream.synchronize(); + atomic_kernel<<>>(srchost, dsthost); + stream.synchronize(); + copy_kernel<<>>(dst.ptr, src.ptr); + stream.synchronize(); + + HIP_API_CALL(hipHostFree(srchost)); + HIP_API_CALL(hipHostFree(dsthost)); + } + } + virtual std::string_view name() override { return "GMIWorkload"; }; +}; + +auto tcp1_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("TCP_CACHE_ACCESS"); + counters.push_back("TCP_CACHE_MISS"); + counters.push_back("TCP_READ"); + counters.push_back("TCP_WRITE"); + counters.push_back("TCC_EA0_WRREQ_DRAM"); + counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM"); + counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM_32B"); + counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM"); + } + else if (gfxip.find("gfx94") == 0) + { + counters.push_back("TCP_READ"); + counters.push_back("TCP_WRITE"); + counters.push_back("TCC_REQ"); + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_ATOMIC"); + counters.push_back("TCC_EA0_ATOMIC"); + } + return counters; +} + +auto tcp2_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("TCP_CACHE_MISS_TG0"); + counters.push_back("TCP_CACHE_MISS_TG1"); + counters.push_back("TCP_CACHE_MISS_TG2"); + counters.push_back("TCP_CACHE_MISS_TG3"); + } + else if (gfxip.find("gfx94") == 0) + { + counters.push_back("TCP_READ"); + counters.push_back("TCP_WRITE"); + counters.push_back("TCC_REQ"); + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_ATOMIC"); + counters.push_back("TCC_EA0_ATOMIC"); + } + return counters; +} + +auto atomic_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM"); + counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM_32B"); + counters.push_back("TCC_EA0_WRREQ_DRAM"); + counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM"); + } + else if (gfxip.find("gfx94") == 0) + { + counters.push_back("TCP_READ"); + counters.push_back("TCP_WRITE"); + counters.push_back("TCC_REQ"); + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_ATOMIC"); + counters.push_back("TCC_EA0_ATOMIC"); + } + return counters; +} + +auto iops_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("SQ_INSTS_VALU_FLOPS_FP16"); + counters.push_back("SQ_INSTS_VALU_FLOPS_FP32"); + counters.push_back("SQ_INSTS_VALU_FLOPS_FP64"); + counters.push_back("SQ_INSTS_VALU_FLOPS_FP16_TRANS"); + counters.push_back("SQ_INSTS_VALU_FLOPS_FP32_TRANS"); + counters.push_back("SQ_INSTS_VALU_FLOPS_FP64_TRANS"); + } + return counters; +} + +auto gmi_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_EA0_RDREQ_GMI_32B"); + counters.push_back("TCC_EA0_WRREQ_GMI_32B"); + counters.push_back("TCC_EA0_ATOMIC_GMI_32B"); + } + else if (gfxip.find("gfx94") == 0) + { + counters.push_back("TCC_EA0_WRREQ_CREDIT_STALL"); + counters.push_back("TCC_EA0_WRREQ_IO_CREDIT_STALL"); + counters.push_back("TCC_EA0_WRREQ_GMI_CREDIT_STALL"); + counters.push_back("TCC_EA0_WRREQ_DRAM_CREDIT_STALL"); + } + return counters; +} + +auto io_counters(std::string_view gfxip) +{ + std::vector counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"}; + + if (gfxip.find("gfx95") == 0) + { + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_EA0_RDREQ_IO_32B"); + counters.push_back("TCC_EA0_WRREQ_IO_32B"); + counters.push_back("TCC_EA0_ATOMIC_IO_32B"); + } + else if (gfxip.find("gfx94") == 0) + { + counters.push_back("TCC_EA0_RDREQ"); + counters.push_back("TCC_EA0_RDREQ_IO_CREDIT_STALL"); + counters.push_back("TCC_EA0_RDREQ_GMI_CREDIT_STALL"); + counters.push_back("TCC_EA0_RDREQ_DRAM_CREDIT_STALL"); + } + return counters; +} + +void printcounters(const std::map& map) +{ + for (auto& [name, v] : map) std::cout << " - " << name << ": " << v << std::endl; +} + +int main() +{ + CHECK_HSA(hsa_init()); + AgentInfo::iterate_agents(); + + auto agent = AgentInfo::gpu_agents.at(0); + + { + Queue queue(agent); + CopyWorkload tcp1(*agent, tcp1_counters(agent->gfxip)); + CopyWorkload tcp2(*agent, tcp2_counters(agent->gfxip)); + AtomicWorkload atomic(*agent, atomic_counters(agent->gfxip)); + IOPSWorkload1 iops1(*agent, iops_counters(agent->gfxip)); + IOPSWorkload2 iops2(*agent, iops_counters(agent->gfxip)); + IOPSWorkload3 iops3(*agent, iops_counters(agent->gfxip)); + + // warmup + tcp1.run(); + tcp2.run(); + atomic.run(); + + // Test + tcp1.printcounters(queue); + tcp2.printcounters(queue); + atomic.printcounters(queue); + iops1.printcounters(queue); + iops2.printcounters(queue); + iops3.printcounters(queue); + +#ifdef ATTEMPT_GMI + GMIWorkload(*agent, gmi_counters(agent->gfxip)).printcounters(queue); + GMIWorkload(*agent, io_counters(agent->gfxip)).printcounters(queue); +#endif + } + + CHECK_HSA(hsa_shut_down()); + return 0; +} diff --git a/projects/aqlprofile/test/integration/workload.cpp b/projects/aqlprofile/test/integration/workload.cpp new file mode 100644 index 0000000000..209c45a660 --- /dev/null +++ b/projects/aqlprofile/test/integration/workload.cpp @@ -0,0 +1,68 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include +#include "counter.hpp" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "workload.hpp" +#include "counter.hpp" + +Collection::Collection(AgentInfo& agent, const std::vector& counters) + : packet(std::make_unique(agent, counters)) {} + +Collection::~Collection() {} + +std::map Collection::iterate(Queue& queue, IWorkload& load) +{ + start(queue); + load.run(); + stop(queue); + return packet->get(); +} + +void Collection::start(Queue& queue) +{ + assert(packet); + queue.flush(); + queue.Submit(&packet->packets.start_packet); +} + +void Collection::stop(Queue& queue) +{ + assert(packet); + queue.flush(); + + queue.Submit(&packet->packets.read_packet); + queue.Submit(&packet->packets.stop_packet); + packet->iterate(); +} diff --git a/projects/aqlprofile/test/integration/workload.hpp b/projects/aqlprofile/test/integration/workload.hpp new file mode 100644 index 0000000000..0397aee541 --- /dev/null +++ b/projects/aqlprofile/test/integration/workload.hpp @@ -0,0 +1,61 @@ +// MIT License +// +// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include +#include +#include +#include "counter.hpp" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "agent.hpp" + +class IWorkload +{ +public: + IWorkload() {} + virtual ~IWorkload() {}; + virtual void run() = 0; +}; + +class Collection +{ +public: + Collection(AgentInfo& agent, const std::vector& counters); + virtual ~Collection(); + + std::map iterate(Queue& queue, IWorkload& load); + +private: + void start(Queue& queue); + void stop(Queue& queue); + std::unique_ptr packet{nullptr}; +};