Elwazir, Ammar
fb6ec25d41
* Adding Integeration tests for later building option
* Update CMakeLists.txt
---------
Co-authored-by: Ammar ELWazir <aelwazir@amd.com>
[ROCm/aqlprofile commit: e4bb0df9d5]
87 行
2.9 KiB
C++
87 行
2.9 KiB
C++
// 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 "aqlprofile-sdk/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;
|
|
}; |