AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Adding tests for code coverage (#82)

Bläddra i källkod 
* Adding tests for code coverage

* integration tests

* Fixing gtest dependencies

* Removing gmock linkage

* build fixes

* Fixing Analyze job

* CMake updates

* Fixing PSDB

* Addressing feedback

* Adding licenses

* Addressing feedback

* updating dashboard

* more API tests

* Addressing feedback

* correcting Macro

* preloading libintercept
This commit is contained in:
Bhardwaj, Gopesh
2025-06-12 19:13:44 +05:30
committad av GitHub
förälder ed7a2104ef
incheckning eca6739e13
26 ändrade filer med 2086 tillägg och 5 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
test/CMakeLists.txt
+5
Visa fil
@@ -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()
test/integration/CMakeLists.txt
+67
Visa fil
@@ -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}")
test/integration/agent.cpp
+235
Visa fil
@@ -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 <cstring>
#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); }
std::vector<std::shared_ptr<AgentInfo>> 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<aqlprofile_pmc_event_t> cnt{};
for (int i=0; i<block_cnt; i++)
{
block.block_index = i;
cnt.push_back(block);
}
counters[counter] = std::move(cnt);
}
hsa_status_t AgentInfo::get_agent_handle_cb(hsa_agent_t agent, void* userdata)
{
hsa_device_type_t type;
CHECK_HSA(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type));
if (type == HSA_DEVICE_TYPE_CPU)
{
cpu_agent = agent;
return HSA_STATUS_SUCCESS;
}
std::shared_ptr<AgentInfo> info = std::make_shared<AgentInfo>();
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_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_XCC), &info->info.xcc_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES), &info->info.se_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT), &info->info.cu_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(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<uint32_t*>(size_t(queue->base_address) + index); // NOLINT
const auto* slot_data = reinterpret_cast<const uint32_t*>(packet);
std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t));
auto* header = reinterpret_cast<std::atomic<uint32_t>*>(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<AgentInfo>& _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<<HSA_PACKET_HEADER_BARRIER);
barrier.header |= HSA_FENCE_SCOPE_SYSTEM<<HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
barrier.header |= HSA_FENCE_SCOPE_SYSTEM<<HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
Submit((hsa_ext_amd_aql_pm4_packet_t*)&barrier);
}
test/integration/agent.hpp
+88
Visa fil
@@ -0,0 +1,88 @@
// 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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
//#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<std::string, std::vector<aqlprofile_pmc_event_t>> counters;
static void iterate_agents();
static std::vector<std::shared_ptr<AgentInfo>> 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<AgentInfo>& _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<AgentInfo> agent;
};
test/integration/counter.cpp
+127
Visa fil
@@ -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 <assert.h>
#include "counter.hpp"
#include <cstring>
#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<AQLPacket*>(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<AQLPacket*>(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<std::string>& _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<aqlprofile_pmc_event_t> 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<uint32_t>(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));
}
test/integration/counter.hpp
+72
Visa fil
@@ -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 <assert.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include "aql_profile_v2.h"
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include "agent.hpp"
class AQLPacket
{
using desc_t = aqlprofile_buffer_desc_flags_t;
public:
AQLPacket(AgentInfo& _agent, const std::vector<std::string>& 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<std::string, int64_t> get()
{
std::map<std::string, int64_t> ret;
for (auto& [event, counter] : counter_names)
ret.emplace(counter, results.at(event));
return ret;
}
std::map<aqlprofile_pmc_event_t, std::string> counter_names;
std::map<aqlprofile_pmc_event_t, int64_t> results;
std::map<aqlprofile_pmc_event_t, int64_t> prev_results;
aqlprofile_handle_t handle{0};
hsa_agent_t hsa_agent;
aqlprofile_pmc_aql_packets_t packets;
bool delta = false;
};
test/integration/exportmap
+4
Visa fil
@@ -0,0 +1,4 @@
{
global: HSA_AMD_TOOL_PRIORITY; OnLoad;
local: *;
};
test/integration/intercept.cpp
+168
Visa fil
@@ -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 <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <unistd.h>
#include <vector>
#include <atomic>
#include <cstring>
#include <hsa/hsa.h>
#include <hsa_ext_amd.h>
#include <hsa_api_trace.h>
#include <hsa_ven_amd_aqlprofile.h>
#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<uint32_t*>(size_t(queue->base_address) + index); // NOLINT
const auto* slot_data = reinterpret_cast<const uint32_t*>(packet);
std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t));
auto* header = reinterpret_cast<std::atomic<uint32_t>*>(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;
}
test/integration/main.cpp
+419
Visa fil
@@ -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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#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<std::string>& counters)
{
col = std::make_unique<Collection>(agent, counters);
}
virtual ~HIPWorkload() {};
virtual std::string_view name() = 0;
std::map<std::string, int64_t> 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<Collection> 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<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(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<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
atomic_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(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<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel1<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "IOPSWorkload1"; };
};
class IOPSWorkload2 : public HIPWorkload
{
public:
IOPSWorkload2(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel2<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "IOPSWorkload2"; };
};
class IOPSWorkload3 : public HIPWorkload
{
public:
IOPSWorkload3(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel_trans<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "Trans IOPSWorkload"; };
};
class GMIWorkload : public HIPWorkload
{
public:
GMIWorkload(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
auto policies = std::vector<unsigned>{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<DATA_SIZE; i++)
srchost[i] = float(i);
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(dsthost, srchost);
stream.synchronize();
atomic_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(srchost, dsthost);
stream.synchronize();
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(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<std::string> 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<std::string> 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<std::string> 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<std::string> 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<std::string> 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<std::string> 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<std::string, int64_t>& 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;
}
test/integration/workload.cpp
+68
Visa fil
@@ -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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include "counter.hpp"
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include "workload.hpp"
#include "counter.hpp"
Collection::Collection(AgentInfo& agent, const std::vector<std::string>& counters)
: packet(std::make_unique<AQLPacket>(agent, counters)) {}
Collection::~Collection() {}
std::map<std::string, int64_t> 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();
}
test/integration/workload.hpp
+61
Visa fil
@@ -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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include "counter.hpp"
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include "agent.hpp"
class IWorkload
{
public:
IWorkload() {}
virtual ~IWorkload() {};
virtual void run() = 0;
};
class Collection
{
public:
Collection(AgentInfo& agent, const std::vector<std::string>& counters);
virtual ~Collection();
std::map<std::string, int64_t> iterate(Queue& queue, IWorkload& load);
private:
void start(Queue& queue);
void stop(Queue& queue);
std::unique_ptr<AQLPacket> packet{nullptr};
};