AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-468021: Improving profiler performance

Browse Source 
Change-Id: I91921160a19aaea1fedfa40d11b83114d6a6fc65


[ROCm/rocprofiler commit: 4d0fee6637]
This commit is contained in:
Giovanni LB
2024-02-29 18:28:31 -03:00
parent 437d39de9e
commit 4a09c57f1a
11 changed files with 512 additions and 414 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler/src/core/counters/metrics/derived_counters.xml
+39 -39
View File
@@ -353,8 +353,8 @@
<metric name="KERNEL_DURATION" expr=1 descr="The duration of the kernel dispatch"></metric>
<metric name="WriteUnitStalled" expr=100*TCC_WRREQ_STALL_max/GRBM_GUI_ACTIVE descr="The percentage of GPUTime the Write unit is stalled. Value range: 0% to 100% (bad)."></metric>
<metric name="TCP_TCP_TA_DATA_STALL_CYCLES_sum" expr=sum(TCP_TCP_TA_DATA_STALL_CYCLES,16) descr="Total number of TCP stalls TA data interface."></metric>
<metric name="TCP_TCP_TA_DATA_STALL_CYCLES_max" expr=max(TCP_TCP_TA_DATA_STALL_CYCLES,16) descr="Maximum number of TCP stalls TA data interface."></metric>
<metric name="TCP_TCP_TA_DATA_STALL_CYCLES_sum" expr=sum(TCP_TCP_TA_DATA_STALL_CYCLES,10) descr="Total number of TCP stalls TA data interface."></metric>
<metric name="TCP_TCP_TA_DATA_STALL_CYCLES_max" expr=max(TCP_TCP_TA_DATA_STALL_CYCLES,10) descr="Maximum number of TCP stalls TA data interface."></metric>
<metric name="MeanOccupancyPerCU" expr=SQ_LEVEL_WAVES*0+SQ_ACCUM_PREV_HIRES*XCC_NUM/GRBM_GUI_ACTIVE/CU_NUM descr="Mean occupancy per compute unit."></metric>
<metric name="MeanOccupancyPerActiveCU" expr=SQ_LEVEL_WAVES*0+SQ_ACCUM_PREV_HIRES*4*XCC_NUM/SQ_BUSY_CYCLES/CU_NUM descr="Mean occupancy per active compute unit."></metric>
@@ -393,43 +393,43 @@
<metric name="TD_STORE_WAVEFRONT_sum" expr=sum(TD_STORE_WAVEFRONT,16) descr="Count the wavefronts with opcode = store. Sum over TD instances."></metric>
<metric name="TD_COALESCABLE_WAVEFRONT_sum" expr=sum(TD_COALESCABLE_WAVEFRONT,16) descr="Count wavefronts that TA finds coalescable. Sum over TD instances."></metric>
<metric name="TD_SPI_STALL_sum" expr=sum(TD_SPI_STALL,16) descr="TD is stalled SPI vinit, sum of TCP instances"></metric>
<metric name="TCP_GATE_EN1_sum" expr=sum(TCP_GATE_EN1,16) descr="TCP interface clocks are turned on. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_GATE_EN2_sum" expr=sum(TCP_GATE_EN2,16) descr="TCP core clocks are turned on. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_TD_TCP_STALL_CYCLES_sum" expr=sum(TCP_TD_TCP_STALL_CYCLES,16) descr="TD stalls TCP. Sum over TCP instances."></metric>
<metric name="TCP_TCR_TCP_STALL_CYCLES_sum" expr=sum(TCP_TCR_TCP_STALL_CYCLES,16) descr="TCR stalls TCP_TCR_req interface. Sum over TCP instances."></metric>
<metric name="TCP_READ_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_READ_TAGCONFLICT_STALL_CYCLES,16) descr="Tagram conflict stall on a read. Sum over TCP instances."></metric>
<metric name="TCP_WRITE_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_WRITE_TAGCONFLICT_STALL_CYCLES,16) descr="Tagram conflict stall on a write. Sum over TCP instances."></metric>
<metric name="TCP_ATOMIC_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_ATOMIC_TAGCONFLICT_STALL_CYCLES,16) descr="Tagram conflict stall on an atomic. Sum over TCP instances."></metric>
<metric name="TCP_VOLATILE_sum" expr=sum(TCP_VOLATILE,16) descr="Total number of L1 volatile pixels/buffers from TA. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ACCESSES_sum" expr=sum(TCP_TOTAL_ACCESSES,16) descr="Total number of pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_READ+TCP_PERF_SEL_TOTAL_NONREAD. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_READ_sum" expr=sum(TCP_TOTAL_READ,16) descr="Total number of read pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_HIT_LRU_READ + TCP_PERF_SEL_TOTAL_MISS_LRU_READ + TCP_PERF_SEL_TOTAL_MISS_EVICT_READ. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_WRITE_sum" expr=sum(TCP_TOTAL_WRITE,16) descr="Total number of local write pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_MISS_LRU_WRITE+ TCP_PERF_SEL_TOTAL_MISS_EVICT_WRITE. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ATOMIC_WITH_RET_sum" expr=sum(TCP_TOTAL_ATOMIC_WITH_RET,16) descr="Total number of atomic with return pixels/buffers from TA. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ATOMIC_WITHOUT_RET_sum" expr=sum(TCP_TOTAL_ATOMIC_WITHOUT_RET,16) descr="Total number of atomic without return pixels/buffers from TA Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_WRITEBACK_INVALIDATES_sum" expr=sum(TCP_TOTAL_WRITEBACK_INVALIDATES,16) descr="Total number of cache invalidates. Equals TCP_PERF_SEL_TOTAL_WBINVL1+ TCP_PERF_SEL_TOTAL_WBINVL1_VOL+ TCP_PERF_SEL_CP_TCP_INVALIDATE+ TCP_PERF_SEL_SQ_TCP_INVALIDATE_VOL. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_REQUEST_sum" expr=sum(TCP_UTCL1_REQUEST,16) descr="Total CLIENT_UTCL1 NORMAL requests Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_TRANSLATION_MISS_sum" expr=sum(TCP_UTCL1_TRANSLATION_MISS,16) descr="Total utcl1 translation misses Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_TRANSLATION_HIT_sum" expr=sum(TCP_UTCL1_TRANSLATION_HIT,16) descr="Total utcl1 translation hits Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_PERMISSION_MISS_sum" expr=sum(TCP_UTCL1_PERMISSION_MISS,16) descr="Total utcl1 permission misses Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_CACHE_ACCESSES_sum" expr=sum(TCP_TOTAL_CACHE_ACCESSES,16) descr="Count of total cache line (tag) accesses (includes hits and misses). Sum over TCP instances."></metric>
<metric name="TCP_TA_TCP_STATE_READ_sum" expr=sum(TCP_TA_TCP_STATE_READ,16) descr="Number of state reads Sum over TCP instances."></metric>
<metric name="TCP_TCC_READ_REQ_sum" expr=sum(TCP_TCC_READ_REQ,16) descr="Total read requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_WRITE_REQ_sum" expr=sum(TCP_TCC_WRITE_REQ,16) descr="Total write requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_ATOMIC_WITH_RET_REQ_sum" expr=sum(TCP_TCC_ATOMIC_WITH_RET_REQ,16) descr="Total atomic with return requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_ATOMIC_WITHOUT_RET_REQ_sum" expr=sum(TCP_TCC_ATOMIC_WITHOUT_RET_REQ,16) descr="Total atomic without return requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_READ_REQ_sum" expr=sum(TCP_TCC_NC_READ_REQ,16) descr="Total read requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_WRITE_REQ_sum" expr=sum(TCP_TCC_NC_WRITE_REQ,16) descr="Total write requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_NC_ATOMIC_REQ,16) descr="Total atomic requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_READ_REQ_sum" expr=sum(TCP_TCC_UC_READ_REQ,16) descr="Total read requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_WRITE_REQ_sum" expr=sum(TCP_TCC_UC_WRITE_REQ,16) descr="Total write requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_UC_ATOMIC_REQ,16) descr="Total atomic requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_READ_REQ_sum" expr=sum(TCP_TCC_CC_READ_REQ,16) descr="Total write requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_WRITE_REQ_sum" expr=sum(TCP_TCC_CC_WRITE_REQ,16) descr="Total write requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_CC_ATOMIC_REQ,16) descr="Total atomic requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_READ_REQ_sum" expr=sum(TCP_TCC_RW_READ_REQ,16) descr="Total write requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_WRITE_REQ_sum" expr=sum(TCP_TCC_RW_WRITE_REQ,16) descr="Total write requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_ATOMIC_REQ_sum" expr=sum(TCP_TCC_RW_ATOMIC_REQ,16) descr="Total atomic requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_PENDING_STALL_CYCLES_sum" expr=sum(TCP_PENDING_STALL_CYCLES,16) descr="Stall due to data pending from L2. Sum over TCP instances."></metric>
<metric name="TCP_GATE_EN1_sum" expr=sum(TCP_GATE_EN1,10) descr="TCP interface clocks are turned on. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_GATE_EN2_sum" expr=sum(TCP_GATE_EN2,10) descr="TCP core clocks are turned on. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_TD_TCP_STALL_CYCLES_sum" expr=sum(TCP_TD_TCP_STALL_CYCLES,10) descr="TD stalls TCP. Sum over TCP instances."></metric>
<metric name="TCP_TCR_TCP_STALL_CYCLES_sum" expr=sum(TCP_TCR_TCP_STALL_CYCLES,10) descr="TCR stalls TCP_TCR_req interface. Sum over TCP instances."></metric>
<metric name="TCP_READ_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_READ_TAGCONFLICT_STALL_CYCLES,10) descr="Tagram conflict stall on a read. Sum over TCP instances."></metric>
<metric name="TCP_WRITE_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_WRITE_TAGCONFLICT_STALL_CYCLES,10) descr="Tagram conflict stall on a write. Sum over TCP instances."></metric>
<metric name="TCP_ATOMIC_TAGCONFLICT_STALL_CYCLES_sum" expr=sum(TCP_ATOMIC_TAGCONFLICT_STALL_CYCLES,10) descr="Tagram conflict stall on an atomic. Sum over TCP instances."></metric>
<metric name="TCP_VOLATILE_sum" expr=sum(TCP_VOLATILE,10) descr="Total number of L1 volatile pixels/buffers from TA. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ACCESSES_sum" expr=sum(TCP_TOTAL_ACCESSES,10) descr="Total number of pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_READ+TCP_PERF_SEL_TOTAL_NONREAD. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_READ_sum" expr=sum(TCP_TOTAL_READ,10) descr="Total number of read pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_HIT_LRU_READ + TCP_PERF_SEL_TOTAL_MISS_LRU_READ + TCP_PERF_SEL_TOTAL_MISS_EVICT_READ. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_WRITE_sum" expr=sum(TCP_TOTAL_WRITE,10) descr="Total number of local write pixels/buffers from TA. Equals TCP_PERF_SEL_TOTAL_MISS_LRU_WRITE+ TCP_PERF_SEL_TOTAL_MISS_EVICT_WRITE. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ATOMIC_WITH_RET_sum" expr=sum(TCP_TOTAL_ATOMIC_WITH_RET,10) descr="Total number of atomic with return pixels/buffers from TA. Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_ATOMIC_WITHOUT_RET_sum" expr=sum(TCP_TOTAL_ATOMIC_WITHOUT_RET,10) descr="Total number of atomic without return pixels/buffers from TA Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_WRITEBACK_INVALIDATES_sum" expr=sum(TCP_TOTAL_WRITEBACK_INVALIDATES,10) descr="Total number of cache invalidates. Equals TCP_PERF_SEL_TOTAL_WBINVL1+ TCP_PERF_SEL_TOTAL_WBINVL1_VOL+ TCP_PERF_SEL_CP_TCP_INVALIDATE+ TCP_PERF_SEL_SQ_TCP_INVALIDATE_VOL. Not Windowed. Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_REQUEST_sum" expr=sum(TCP_UTCL1_REQUEST,10) descr="Total CLIENT_UTCL1 NORMAL requests Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_TRANSLATION_MISS_sum" expr=sum(TCP_UTCL1_TRANSLATION_MISS,10) descr="Total utcl1 translation misses Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_TRANSLATION_HIT_sum" expr=sum(TCP_UTCL1_TRANSLATION_HIT,10) descr="Total utcl1 translation hits Sum over TCP instances."></metric>
<metric name="TCP_UTCL1_PERMISSION_MISS_sum" expr=sum(TCP_UTCL1_PERMISSION_MISS,10) descr="Total utcl1 permission misses Sum over TCP instances."></metric>
<metric name="TCP_TOTAL_CACHE_ACCESSES_sum" expr=sum(TCP_TOTAL_CACHE_ACCESSES,10) descr="Count of total cache line (tag) accesses (includes hits and misses). Sum over TCP instances."></metric>
<metric name="TCP_TA_TCP_STATE_READ_sum" expr=sum(TCP_TA_TCP_STATE_READ,10) descr="Number of state reads Sum over TCP instances."></metric>
<metric name="TCP_TCC_READ_REQ_sum" expr=sum(TCP_TCC_READ_REQ,10) descr="Total read requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_WRITE_REQ_sum" expr=sum(TCP_TCC_WRITE_REQ,10) descr="Total write requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_ATOMIC_WITH_RET_REQ_sum" expr=sum(TCP_TCC_ATOMIC_WITH_RET_REQ,10) descr="Total atomic with return requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_ATOMIC_WITHOUT_RET_REQ_sum" expr=sum(TCP_TCC_ATOMIC_WITHOUT_RET_REQ,10) descr="Total atomic without return requests from TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_READ_REQ_sum" expr=sum(TCP_TCC_NC_READ_REQ,10) descr="Total read requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_WRITE_REQ_sum" expr=sum(TCP_TCC_NC_WRITE_REQ,10) descr="Total write requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_NC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_NC_ATOMIC_REQ,10) descr="Total atomic requests with NC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_READ_REQ_sum" expr=sum(TCP_TCC_UC_READ_REQ,10) descr="Total read requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_WRITE_REQ_sum" expr=sum(TCP_TCC_UC_WRITE_REQ,10) descr="Total write requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_UC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_UC_ATOMIC_REQ,10) descr="Total atomic requests with UC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_READ_REQ_sum" expr=sum(TCP_TCC_CC_READ_REQ,10) descr="Total write requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_WRITE_REQ_sum" expr=sum(TCP_TCC_CC_WRITE_REQ,10) descr="Total write requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_CC_ATOMIC_REQ_sum" expr=sum(TCP_TCC_CC_ATOMIC_REQ,10) descr="Total atomic requests with CC mtype from this TCP to all TCCs Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_READ_REQ_sum" expr=sum(TCP_TCC_RW_READ_REQ,10) descr="Total write requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_WRITE_REQ_sum" expr=sum(TCP_TCC_RW_WRITE_REQ,10) descr="Total write requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_TCC_RW_ATOMIC_REQ_sum" expr=sum(TCP_TCC_RW_ATOMIC_REQ,10) descr="Total atomic requests with RW mtype from this TCP to all TCCs. Sum over TCP instances."></metric>
<metric name="TCP_PENDING_STALL_CYCLES_sum" expr=sum(TCP_PENDING_STALL_CYCLES,10) descr="Stall due to data pending from L2. Sum over TCP instances."></metric>
<metric name="TCP_TCR_TCP_STALL_CYCLES_PERCENT" expr=100*TCP_TCR_TCP_STALL_CYCLES_sum/TCP_GATE_EN1_sum descr="Percentage of time TCP is stalled by TCR."></metric>
<metric name="TCA_CYCLE_sum" expr=sum(TCA_CYCLE,16) descr="Number of cycles. Sum over all TCA instances "></metric>
<metric name="TCA_BUSY_sum" expr=sum(TCA_BUSY,16) descr="Number of cycles we have a request pending. Sum over all TCA instances."></metric>
projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp
+31 -25
View File
@@ -27,40 +27,43 @@ struct block_status_t {
uint32_t group_index;
};
struct aqlprofile_event_t : public hsa_ven_amd_aqlprofile_event_t
{
bool operator==(const aqlprofile_event_t& other) const {
return this->block_name == other.block_name &&
this->block_index == other.block_index &&
this->counter_id ==other.counter_id;
}
};
template <>
struct std::hash<aqlprofile_event_t>
{
std::size_t operator()(const aqlprofile_event_t& k) const {
return (int(k.block_name)<<20) ^ (int(k.counter_id)<<10) ^ int(k.block_index);
}
};
typedef struct {
std::vector<results_t*>* results;
std::unordered_map<aqlprofile_event_t, results_t*> results;
size_t index;
} callback_data_t;
static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1,
const hsa_ven_amd_aqlprofile_event_t& event2) {
return (event1.block_name == event2.block_name) && (event1.block_index == event2.block_index) &&
(event1.counter_id == event2.counter_id);
}
hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
if (info_type != HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA) return HSA_STATUS_SUCCESS;
callback_data_t* passed_data = reinterpret_cast<callback_data_t*>(data);
try {
for (auto data_it = passed_data->results->begin(); data_it != passed_data->results->end();
++data_it) {
if (info_type != HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA) continue;
if (!IsEventMatch(info_data->pmc_data.event, (*data_it)->event)) continue;
// stores event result from each event separately
(*data_it)->xcc_vals.push_back(info_data->pmc_data.result);
// stores accumulated event result from all xccs
(*data_it)->val_double += info_data->pmc_data.result;
}
} catch (std::exception& e) {
std::cout << "caught an exception in eval_metrics.cpp:pmcCallback(): " << e.what() << std::endl;
}
auto it = passed_data->results.find(aqlprofile_event_t{info_data->pmc_data.event});
if (it == passed_data->results.end())
return HSA_STATUS_ERROR;
auto* res = it->second;
res->xcc_vals.push_back(info_data->pmc_data.result);
res->val_double += info_data->pmc_data.result;
passed_data->index += 1;
return status;
return HSA_STATUS_SUCCESS;
}
@@ -187,10 +190,13 @@ bool metrics::ExtractMetricEvents(
bool metrics::GetCounterData(hsa_ven_amd_aqlprofile_profile_t* profile, hsa_agent_t gpu_agent,
std::vector<results_t*>& results_list) {
size_t gpu_xcc_count = HSASupport_Singleton::GetInstance().GetHSAAgentInfo(gpu_agent.handle).GetDeviceInfo().getXccCount();
callback_data_t callback_data{&results_list, 0};
callback_data_t callback_data{};
for (auto* res : results_list)
callback_data.results[aqlprofile_event_t{res->event}] = res;
hsa_status_t status = hsa_ven_amd_aqlprofile_iterate_data(profile, pmcCallback, &callback_data);
for (auto& data : *(callback_data.results))
for (auto& data : results_list)
{
size_t xcc_count = (data->event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_UMC) ? gpu_xcc_count : 1;
std::vector<double> xcc_results = std::move(data->xcc_vals);
projects/rocprofiler/src/core/hsa/hsa_support.cpp
+3 -3
View File
@@ -844,8 +844,8 @@ void HSASupport_Singleton::InitKsymbols() {
ksymbols_flag.exchange(false, std::memory_order_release);
}
{
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
kernel_names = new std::map<std::string, std::vector<uint64_t>>();
std::unique_lock<std::shared_mutex> lock(kernel_names_map_lock);
kernel_names = new std::unordered_map<uint64_t, std::string>();
kernel_names_flag.exchange(false, std::memory_order_release);
}
}
@@ -858,7 +858,7 @@ void HSASupport_Singleton::FinitKsymbols() {
ksymbols_flag.exchange(true, std::memory_order_release);
}
if (!kernel_names_flag.load(std::memory_order_relaxed)) {
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
std::unique_lock<std::shared_mutex> lock(kernel_names_map_lock);
kernel_names->clear();
delete kernel_names;
kernel_names_flag.exchange(true, std::memory_order_release);
projects/rocprofiler/src/core/hsa/hsa_support.h
+3 -3
View File
@@ -33,7 +33,7 @@
#include <unordered_map>
#include <memory>
#include <optional>
#include <shared_mutex>
#include "rocprofiler.h"
#include "src/core/hardware/hsa_info.h"
@@ -152,8 +152,8 @@ class HSASupport_Singleton {
HSAAgentInfo& GetHSAAgentInfo(uint64_t agent_handle);
HSAAgentInfo& GetHSAAgentInfo(Agent::DeviceInfo device_info);
Agent::DeviceInfo& GetDeviceInfo(HSAAgentInfo* agent_info);
std::mutex kernel_names_map_lock;
std::map<std::string, std::vector<uint64_t>>* kernel_names;
std::shared_mutex kernel_names_map_lock;
std::unordered_map<uint64_t, std::string>* kernel_names;
std::mutex ksymbol_map_lock;
std::map<uint64_t, std::string>* ksymbols;
std::mutex signals_timestamps_map_lock;
projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp
+163 -105
View File
@@ -43,6 +43,7 @@
#include "src/core/counters/metrics/metrics.h"
#include "src/core/hardware/hsa_info.h"
#include "src/core/hsa/packets/packets_generator.h"
#define ASSERTM(exp, msg) assert(((void)msg, exp))
@@ -160,16 +161,20 @@ void CheckPacketReqiurements() {
// Initialize the PM4 commands with having the CPU&GPU agents, the counters,
// counters count to output three packets which are start, stop and read
// packets
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names, rocprofiler_session_id_t session_id,
bool is_spm) {
std::unique_ptr<AQLPacketProfile> InitializeAqlPackets(
hsa_agent_t cpu_agent,
hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names,
rocprofiler_session_id_t session_id,
bool is_spm
) {
hsa_status_t status = HSA_STATUS_SUCCESS;
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton =
rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
if (!counters_added.load(std::memory_order_acquire)) {
if (!counters_added.load(std::memory_order_acquire))
{
for (auto& name : counter_names) {
if (rocprofiler_singleton.HasActiveSession()) {
rocprofiler_singleton.GetSession(session_id)->GetProfiler()->AddCounterName(name);
@@ -205,7 +210,10 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
}
// do {
rocprofiler::profiling_context_t* context = new rocprofiler::profiling_context_t();
auto prof_context = std::make_unique<AQLPacketProfile>(hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn);
auto* context = prof_context->context.get();
auto* profile = prof_context->profile.get();
context->gpu_agent = gpu_agent;
auto result = results_list.begin();
std::map<std::pair<uint32_t, uint32_t>, uint32_t> block_max_events_count;
@@ -292,120 +300,93 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
context->results_map = results_map;
context->metrics_dict = metricsDict[gpu_agent.handle];
hsa_ven_amd_aqlprofile_parameter_t* params = {};
packet_t* start_packet = new packet_t();
packet_t* stop_packet = new packet_t();
packet_t* read_packet = new packet_t();
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles = std::vector<
std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>();
// if (context->events_list.size() <= 0) {
// std::cerr << "Error: No events to profile" << std::endl;
// abort();
// }
// Preparing the profile structure to get the packets
hsa_ven_amd_aqlprofile_event_type_t profile_type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
if (is_spm) profile_type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE;
hsa_ven_amd_aqlprofile_profile_t* profile =
new hsa_ven_amd_aqlprofile_profile_t{gpu_agent,
profile_type,
&(context->events_list[0]),
static_cast<uint32_t>(context->events_list.size()),
params,
0,
0,
0};
profile->agent = gpu_agent;
profile->type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
profile->events = &(context->events_list[0]);
profile->event_count = static_cast<uint32_t>(context->events_list.size());
size_t ag_list_count = 1; // rocprofiler::hsa_support::GetCPUAgentList().size();
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
if (context->events_list.size() > 0) {
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
// CHECK_HSA_STATUS("Error: Getting Buffers Size", status);
if (context->events_list.size() == 0)
return prof_context;
if (profile->command_buffer.size > 0 && profile->output_buffer.size > 0) {
status = HSA_STATUS_ERROR;
size_t size = profile->command_buffer.size;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
if (size <= 0) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: Command buffer given size is " << size << std::endl;
abort();
}
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool_, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
if (status != HSA_STATUS_SUCCESS) {
profile->command_buffer.ptr = malloc(size);
/*numa_alloc_onnode(
size,
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());*/
if (profile->command_buffer.ptr == NULL) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: allocating memory for command buffer using NUMA" << std::endl;
abort();
}
} else {
// Both the CPU and GPU can access the memory
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->command_buffer.ptr);
CHECK_HSA_STATUS("Error: Allowing access to Command Buffer", status);
}
if (!is_spm) {
status = HSA_STATUS_ERROR;
size_t size = profile->output_buffer.size;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
if (size <= 0) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: Output buffer given size is " << size << std::endl;
abort();
}
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool_, size, 0,
reinterpret_cast<void**>(&profile->output_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
profile->output_buffer.ptr = malloc(size);
/*numa_alloc_onnode(
size,
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle)
.getNumaNode());*/
if (profile->output_buffer.ptr == NULL) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: allocating memory for output buffer using NUMA" << std::endl;
abort();
}
} else {
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->output_buffer.ptr);
CHECK_HSA_STATUS("Error: GPU Agent can't have output buffer access", status);
memset(profile->output_buffer.ptr, 0x0, profile->output_buffer.size);
}
} else {
profile->output_buffer.size = 0;
}
status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
// CHECK_HSA_STATUS("Error: Creating Start Packet\n", status);
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
// CHECK_HSA_STATUS("Error: Creating Stop Packet\n", status);
status = hsa_ven_amd_aqlprofile_read(profile, read_packet);
// CHECK_HSA_STATUS("Error: Creating Read Packet\n", status);
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
CHECK_HSA_STATUS("Error: Getting Buffers Size", status);
context->start_packet = start_packet;
context->stop_packet = stop_packet;
context->read_packet = read_packet;
}
if (profile->command_buffer.size == 0 || profile->output_buffer.size == 0)
{
std::cerr << __FILE__ << ":" << __LINE__ << "Error: Did not return buffer size" << std::endl;
abort();
}
// add profiles
profiles.emplace_back(std::make_pair(context, profile));
// } while (events_list.size() > 0);
return profiles;
status = HSA_STATUS_ERROR;
size_t size = (profile->command_buffer.size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool_, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
if (status != HSA_STATUS_SUCCESS) {
profile->command_buffer.ptr = malloc(size);
/*numa_alloc_onnode(
size,
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());*/
if (profile->command_buffer.ptr == NULL) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: allocating memory for command buffer using NUMA" << std::endl;
abort();
}
} else {
// Both the CPU and GPU can access the memory
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->command_buffer.ptr);
CHECK_HSA_STATUS("Error: Allowing access to Command Buffer", status);
}
if (!is_spm) {
status = HSA_STATUS_ERROR;
size_t size = (profile->output_buffer.size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool_, size, 0, &profile->output_buffer.ptr
);
if (status != HSA_STATUS_SUCCESS) {
profile->output_buffer.ptr = malloc(size);
if (profile->output_buffer.ptr == NULL) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
<< "Error: allocating memory for output buffer using NUMA" << std::endl;
abort();
}
} else {
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->output_buffer.ptr);
CHECK_HSA_STATUS("Error: GPU Agent can't have output buffer access", status);
hsa_amd_memory_fill(profile->output_buffer.ptr, 0, profile->output_buffer.size/sizeof(uint32_t));
}
} else {
profile->output_buffer.size = 0;
}
status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
CHECK_HSA_STATUS("Error: Creating Start Packet\n", status);
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
CHECK_HSA_STATUS("Error: Creating Stop Packet\n", status);
status = hsa_ven_amd_aqlprofile_read(profile, read_packet);
CHECK_HSA_STATUS("Error: Creating Read Packet\n", status);
context->start_packet = start_packet;
context->stop_packet = stop_packet;
context->read_packet = read_packet;
return prof_context;
}
// Initialize the PM4 commands with having the CPU&GPU agents, the counters,
@@ -665,10 +646,7 @@ void CreateBarrierPacket(std::vector<packet_t>* transformed_packets,
const hsa_signal_t* packet_completion_signal
) {
hsa_barrier_and_packet_t barrier{0};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE |
(1 << HSA_PACKET_HEADER_BARRIER) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE | (1 << HSA_PACKET_HEADER_BARRIER);
if (packet_completion_signal != nullptr) barrier.completion_signal = *packet_completion_signal;
if (packet_dependency_signal != nullptr) barrier.dep_signal[0] = *packet_dependency_signal;
void* barrier_ptr = &barrier;
@@ -714,4 +692,84 @@ std::vector<const hsa_kernel_dispatch_packet_s*> ExtractDispatchPackets(
return ret;
}
std::atomic<int> AQLPacketProfile::valid_profiles{0};
std::condition_variable_any AQLPacketProfile::delete_cv;
std::shared_mutex AQLPacketProfile::deleter_mutex;
bool AQLPacketProfile::IsDeletingBegin = false;
std::unordered_map<uint64_t, std::vector<std::unique_ptr<AQLPacketProfile>>> _cache;
std::mutex cache_mutex;
void AQLPacketProfile::WaitForProfileDeletion()
{
{
std::unique_lock<std::mutex> lk(cache_mutex);
_cache.clear();
}
std::unique_lock<std::shared_mutex> lk(deleter_mutex);
if (valid_profiles.load() == 0)
{
IsDeletingBegin = true;
return;
}
delete_cv.wait_for(lk, std::chrono::seconds(2), [] () {
return ::Packet::AQLPacketProfile::valid_profiles.load() == 0;
});
IsDeletingBegin = true;
}
std::unique_ptr<AQLPacketProfile> AQLPacketProfile::MoveFromCache(hsa_agent_t gpu_agent)
{
std::lock_guard<std::mutex> lk(cache_mutex);
auto agent_it = _cache.find(gpu_agent.handle);
if (agent_it == _cache.end()) return nullptr;
auto& profile_set = agent_it->second;
if (!profile_set.size()) return nullptr;
auto moved = std::move(profile_set.back());
profile_set.resize(profile_set.size()-1);
return moved;
}
void AQLPacketProfile::MoveToCache(hsa_agent_t gpu_agent, std::unique_ptr<AQLPacketProfile>&& packet)
{
if (!packet.get()) return;
auto& output_buffer = packet->profile->output_buffer;
std::lock_guard<std::mutex> lk(cache_mutex);
auto agent_it = _cache.find(gpu_agent.handle);
if (agent_it == _cache.end())
agent_it = _cache.emplace(gpu_agent.handle, std::vector<std::unique_ptr<AQLPacketProfile>>{}).first;
auto& profile_set = agent_it->second;
profile_set.emplace_back(std::move(packet));
}
AQLPacketProfile::~AQLPacketProfile()
{
std::shared_lock<std::shared_mutex> deleter_lock(deleter_mutex);
if (IsDeletingBegin) return;
int old_valid = valid_profiles.fetch_sub(1);
if (profile->output_buffer.ptr)
free_fn(profile->output_buffer.ptr);
profile->output_buffer.ptr = nullptr;
profile->output_buffer.size = 0;
if (profile->command_buffer.ptr)
free_fn(profile->command_buffer.ptr);
profile->command_buffer.ptr = nullptr;
profile->command_buffer.size = 0;
if (old_valid <= 1)
delete_cv.notify_all();
}
} // namespace Packet
projects/rocprofiler/src/core/hsa/packets/packets_generator.h
+38 -4
View File
@@ -32,6 +32,7 @@
#include <string>
#include <vector>
#include <mutex>
#include <shared_mutex>
#include "src/core/counters/metrics/eval_metrics.h"
@@ -39,10 +40,43 @@ namespace Packet {
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names, rocprofiler_session_id_t session_id,
bool is_spm = false);
class AQLPacketProfile
{
public:
AQLPacketProfile(decltype(hsa_amd_memory_pool_free) _free_fn)
{
profile = std::make_unique<hsa_ven_amd_aqlprofile_profile_t>();
context = std::make_unique<rocprofiler::profiling_context_t>();
this->free_fn = _free_fn;
valid_profiles.fetch_add(1);
};
~AQLPacketProfile();
std::unique_ptr<hsa_ven_amd_aqlprofile_profile_t> profile;
std::unique_ptr<rocprofiler::profiling_context_t> context;
static std::unique_ptr<AQLPacketProfile> MoveFromCache(hsa_agent_t gpu_agent);
static void MoveToCache(hsa_agent_t gpu_agent, std::unique_ptr<AQLPacketProfile>&& packet);
static void WaitForProfileDeletion();
static std::atomic<int> valid_profiles;
static std::condition_variable_any delete_cv;
static std::shared_mutex deleter_mutex;
static bool IsDeletingBegin;
decltype(hsa_amd_memory_pool_free)* free_fn;
};
std::unique_ptr<AQLPacketProfile> InitializeAqlPackets(
hsa_agent_t cpu_agent,
hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names,
rocprofiler_session_id_t session_id,
bool is_spm = false
);
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
void GetCommandBufferMap(std::map<size_t, uint8_t*>);
void GetOutputBufferMap(std::map<size_t, uint8_t*>);
projects/rocprofiler/src/core/hsa/queues/queue.cpp
+226 -225
View File
@@ -60,12 +60,6 @@ std::mutex sessions_pending_signal_lock;
namespace rocprofiler {
// std::atomic<uint32_t> ACTIVE_INTERRUPT_SIGNAL_COUNT{0};
// uint32_t GetCurrentActiveInterruptSignalsCount() {
// return ACTIVE_INTERRUPT_SIGNAL_COUNT.load(std::memory_order_relaxed);
// }
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> pmc_callback_data_t;
static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1,
@@ -98,19 +92,16 @@ std::string GetKernelNameFromKsymbols(uint64_t handle) {
void AddKernelNameWithDispatchID(std::string name, uint64_t id) {
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.kernel_names_map_lock);
if (hsasupport_singleton.kernel_names->find(name) == hsasupport_singleton.kernel_names->end())
hsasupport_singleton.kernel_names->emplace(name, std::vector<uint64_t>());
hsasupport_singleton.kernel_names->at(name).push_back(id);
std::unique_lock<std::shared_mutex> lock(hsasupport_singleton.kernel_names_map_lock);
(*hsasupport_singleton.kernel_names)[id] = name;
}
std::string GetKernelNameUsingDispatchID(uint64_t given_id) {
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.kernel_names_map_lock);
for (auto kernel_name : (*hsasupport_singleton.kernel_names)) {
for (auto dispatch_id : kernel_name.second) {
if (dispatch_id == given_id) return kernel_name.first;
}
}
std::shared_lock<std::shared_mutex> lock(hsasupport_singleton.kernel_names_map_lock);
auto it = hsasupport_singleton.kernel_names->find(given_id);
if (it != hsasupport_singleton.kernel_names->end())
return it->second;
return "Unknown Kernel!";
}
@@ -273,21 +264,23 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
return status;
}
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) {
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending)
{
auto* context = pending->profile->context.get();
record->counters_count =
rocprofiler_record_counters_instances_count_t{pending->context->metrics_list.size()};
rocprofiler_record_counters_instances_count_t{context->metrics_list.size()};
size_t counters_list_size =
record->counters_count.value * sizeof(rocprofiler_record_counter_instance_t);
rocprofiler_record_counter_instance_t* counters =
static_cast<rocprofiler_record_counter_instance_t*>(malloc(counters_list_size));
for (size_t i = 0; i < pending->context->metrics_list.size(); i++) {
const rocprofiler::Metric* metric = pending->context->metrics_list[i];
for (size_t i = 0; i < context->metrics_list.size(); i++) {
const rocprofiler::Metric* metric = context->metrics_list[i];
double value = 0;
std::string metric_name = metric->GetName();
auto it = pending->context->results_map.find(metric_name);
if (it != pending->context->results_map.end()) {
auto it = context->results_map.find(metric_name);
if (it != context->results_map.end())
value = it->second->val_double;
}
counters[i] = (rocprofiler_record_counter_instance_t{
// TODO(aelwazir): Moving to span once C++20 is adopted, strdup can be
// removed after that
@@ -309,6 +302,12 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
static_cast<const rocprofiler_record_counter_instance_t*>(data);
});
}
// Reset counters
for (auto& [key, value] : context->results_map)
value->val_double = 0;
for (auto* res : context->results_list)
res->val_double = 0;
}
/*
@@ -384,14 +383,15 @@ void SignalAsyncReadyHandler(const hsa_signal_t& signal, void* data) {
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalReadyHandler, data);
if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed");
}
bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data)
{
auto queue_info_session = static_cast<queue_info_session_t*>(data);
if (!queue_info_session) return true;
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton =
rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton::GetInstance();
rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id);
if (!session) return true;
@@ -402,112 +402,98 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
auto pending_signals = profiler->MovePendingSignals(queue_info_session->writer_id);
for (auto& pending : pending_signals) {
if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal))
return true;
hsa_amd_profiling_dispatch_time_t time;
hsasupport_singleton.GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
queue_info_session->agent, pending->new_signal, &time);
{
std::lock_guard<std::mutex> lock(hsasupport_singleton.signals_timestamps_map_lock);
hsasupport_singleton.signals_timestamps[pending->original_signal.handle].time =
std::make_optional(time);
}
uint32_t record_count = 1;
bool is_individual_xcc_mode = false;
uint32_t xcc_count = queue_info_session->xcc_count;
if (xcc_count > 1 && pending->counters_count > 0) { // for MI300
const char* str = getenv("ROCPROFILER_INDIVIDUAL_XCC_MODE");
if (str != NULL) is_individual_xcc_mode = (atol(str) > 0);
// for individual xcc mode, there will be xcc_count records for each dispatch
// for accumulation mode, there will be only one record for a dispatch
if (is_individual_xcc_mode) record_count = xcc_count;
}
for (uint32_t xcc_id = 0; xcc_id < record_count; xcc_id++) {
rocprofiler_record_profiler_t record{};
// TODO: (sauverma) gpu-id will need to support xcc like so- 1.1, 1.2, 1.3 ... 1.5 for
// different xcc
record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index};
record.kernel_properties = pending->kernel_properties;
record.thread_id = rocprofiler_thread_id_t{pending->thread_id};
record.queue_idx = rocprofiler_queue_index_t{pending->queue_index};
record.timestamps = rocprofiler_record_header_timestamp_t{time.start, time.end};
record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id};
record.xcc_index = xcc_id;
// Kernel Descriptor is the right record id generated in the WriteInterceptor function and
// will be used to handle the kernel name of that dispatch
record.header = rocprofiler_record_header_t{
ROCPROFILER_PROFILER_RECORD, rocprofiler_record_id_t{pending->kernel_descriptor}};
record.kernel_id = rocprofiler_kernel_id_t{pending->kernel_descriptor};
record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id};
if (pending->session_id.handle == 0) {
pending->session_id = rocprofiler_singleton.GetCurrentSessionId();
for (auto& pending : pending_signals)
{
if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal))
return true;
hsa_amd_profiling_dispatch_time_t time;
hsasupport_singleton.GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
queue_info_session->agent, pending->new_signal, &time);
{
std::lock_guard<std::mutex> lock(hsasupport_singleton.signals_timestamps_map_lock);
hsasupport_singleton.signals_timestamps[pending->original_signal.handle].time =
std::make_optional(time);
}
if (pending->counters_count > 0) {
if (xcc_id == 0 && pending->context && pending->context->metrics_list.size() > 0 &&
pending->profile) // call to GetCounterData() is required only once for a dispatch
rocprofiler::metrics::GetCounterData(pending->profile, queue_info_session->agent,
pending->context->results_list);
if (is_individual_xcc_mode)
rocprofiler::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, pending->context->results_list, pending->context->results_map,
pending->context->metrics_list, time.end - time.start);
else
rocprofiler::metrics::GetMetricsData(
pending->context->results_map, pending->context->metrics_list, time.end - time.start);
AddRecordCounters(&record, pending.get());
} else {
if (session->FindBuffer(pending->buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id);
buffer->AddRecord(record);
//hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
uint32_t record_count = 1;
uint32_t xcc_count = queue_info_session->xcc_count;
static thread_local bool is_individual_xcc_mode = [xcc_count]() {
if (xcc_count < 2) return false;
const char* str = getenv("ROCPROFILER_INDIVIDUAL_XCC_MODE");
if (str != NULL) return (atol(str) > 0);
return false;
}();
if (is_individual_xcc_mode) record_count = xcc_count;
for (uint32_t xcc_id = 0; xcc_id < record_count; xcc_id++) {
rocprofiler_record_profiler_t record{};
// TODO: (sauverma) gpu-id will need to support xcc like so- 1.1, 1.2, 1.3 ... 1.5 for
// different xcc
record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index};
record.kernel_properties = pending->kernel_properties;
record.thread_id = rocprofiler_thread_id_t{pending->thread_id};
record.queue_idx = rocprofiler_queue_index_t{pending->queue_index};
record.timestamps = rocprofiler_record_header_timestamp_t{time.start, time.end};
record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id};
record.xcc_index = xcc_id;
// Kernel Descriptor is the right record id generated in the WriteInterceptor function and
// will be used to handle the kernel name of that dispatch
record.header = rocprofiler_record_header_t{
ROCPROFILER_PROFILER_RECORD, rocprofiler_record_id_t{pending->kernel_descriptor}};
record.kernel_id = rocprofiler_kernel_id_t{pending->kernel_descriptor};
record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id};
if (pending->session_id.handle == 0) {
pending->session_id = rocprofiler_singleton.GetCurrentSessionId();
}
if (pending->counters_count > 0)
{
auto* context = pending->profile->context.get();
auto* profile = pending->profile->profile.get();
if (xcc_id == 0 && context && context->metrics_list.size() > 0 && profile)
rocprofiler::metrics::GetCounterData(profile, queue_info_session->agent,
context->results_list);
if (is_individual_xcc_mode)
rocprofiler::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, context->results_list, context->results_map,
context->metrics_list, time.end - time.start);
else
rocprofiler::metrics::GetMetricsData(context->results_map,
context->metrics_list,
time.end - time.start);
AddRecordCounters(&record, pending.get());
} else {
if (session->FindBuffer(pending->buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id);
buffer->AddRecord(record);
}
}
}
}
if (pending->counters_count > 0 && pending->profile && pending->profile->events) {
// TODO(aelwazir): we need a better way of distributing events and free them
// if (pending->profile->output_buffer.ptr)
// numa_free(pending->profile->output_buffer.ptr, pending->profile->output_buffer.size);
hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending->profile->output_buffer.ptr));
CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status);
// if (pending->profile->command_buffer.ptr)
// numa_free(pending->profile->command_buffer.ptr, pending->profile->command_buffer.size);
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending->profile->command_buffer.ptr));
CHECK_HSA_STATUS("Error: Couldn't free command buffer memory", status);
delete pending->profile;
for (auto& it : pending->context->results_map) {
delete it.second;
}
delete pending->context;
/*
Check if the dispatch ready is empty, If so, there is no more
dispatches to be launched and we return. Else, dispatch the
kernel of the queue in the front of the dispatch_ready.
*/
auto* profile = pending->profile ? pending->profile->profile.get() : nullptr;
if (pending->counters_count > 0 && profile && profile->events)
{
Packet::AQLPacketProfile::MoveToCache(queue_info_session->agent, std::move(pending->profile));
profiler_serializer_t& serializer =
profiler_serializer_t& serializer =
rocprofiler::ROCProfiler_Singleton::GetInstance().GetSerializer();
std::lock_guard<std::mutex> serializer_lock(serializer.serializer_mutex);
assert(serializer.dispatch_queue != nullptr);
hsasupport_singleton.GetCoreApiTable().hsa_signal_store_screlease_fn(
queue_info_session->block_signal, 1);
serializer.dispatch_queue = nullptr;
if (serializer.dispatch_ready.empty()) return false;
Queue* queue = serializer.dispatch_ready.front();
serializer.dispatch_ready.erase(serializer.dispatch_ready.begin());
enable_dispatch(queue);
}
std::lock_guard<std::mutex> serializer_lock(serializer.serializer_mutex);
assert(serializer.dispatch_queue != nullptr);
hsasupport_singleton.GetCoreApiTable().hsa_signal_store_screlease_fn(
queue_info_session->block_signal, 1);
serializer.dispatch_queue = nullptr;
if (serializer.dispatch_ready.empty()) return false;
Queue* queue = serializer.dispatch_ready.front();
serializer.dispatch_ready.erase(serializer.dispatch_ready.begin());
enable_dispatch(queue);
if (pending->new_signal.handle)
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
if (queue_info_session->interrupt_signal.handle)
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(
queue_info_session->interrupt_signal);
}
if (pending->new_signal.handle)
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
if (queue_info_session->interrupt_signal.handle)
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(
queue_info_session->interrupt_signal);
}
delete queue_info_session;
// ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_sub(1, std::memory_order_relaxed);
return false;
}
@@ -541,23 +527,27 @@ uint32_t replay_mode_count = 0;
rocprofiler::Session* session = nullptr;
void Queue::ResetSessionID(rocprofiler_session_id_t id) {
void Queue::ResetSessionID(rocprofiler_session_id_t id)
{
std::unique_lock<std::shared_mutex> session_id_lock(session_id_mutex);
session_id = id;
if (session_id.handle != 0)
session = rocprofiler::ROCProfiler_Singleton::GetInstance().GetSession(session_id);
else
session = nullptr;
}
void Queue::CheckNeededProfileConfigs() {
rocprofiler_session_id_t internal_session_id;
bool Queue::CheckNeededProfileConfigs()
{
std::unique_lock<std::shared_mutex> session_id_lock(session_id_mutex);
// Getting Session ID
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton =
rocprofiler::ROCProfiler_Singleton::GetInstance();
internal_session_id = rocprofiler_singleton.GetCurrentSessionId();
if (session_id.handle > 0 && internal_session_id.handle == session_id.handle) return;
if (internal_session_id.handle == 0) return;
session_id = internal_session_id;
session_id = rocprofiler_singleton.GetCurrentSessionId();
// Getting Counters count from the Session
if (session_id.handle == 0) return false;
session = rocprofiler_singleton.GetSession(session_id);
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
@@ -568,7 +558,8 @@ void Queue::CheckNeededProfileConfigs() {
is_counter_collection_mode = true;
session_data_count = session_data.size();
buffer_id = filter->GetBufferId();
} else if (session && session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) {
} else if (session &&
session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) {
is_timestamp_collection_mode = true;
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION);
@@ -583,15 +574,18 @@ void Queue::CheckNeededProfileConfigs() {
att_tracer->SetParameters(filter->GetAttParametersData());
is_att_collection_mode = true;
buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))
->GetBufferId();
->GetBufferId();
att_tracer->SetCountersNames(filter->GetCounterData());
att_tracer->SetKernelsNames(
std::get<std::vector<std::string>>(filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES)));
att_tracer->SetDispatchIds(std::get<std::vector<std::pair<uint64_t, uint64_t>>>(
filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS)));
att_tracer->SetKernelsNames(std::get<std::vector<std::string>>(
filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES)
));
att_tracer->SetDispatchIds(std::get<std::vector<std::pair<uint64_t,uint64_t>>>(
filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS)
));
} else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
is_pc_sampling_collection_mode = true;
}
return true;
}
std::atomic<uint32_t> WRITER_ID{0};
@@ -604,79 +598,85 @@ std::atomic<uint32_t> WRITER_ID{0};
* interceptor by invoking the writer function.
*/
void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer) {
std::shared_lock<std::shared_mutex> session_id_lock(session_id_mutex);
void* data, hsa_amd_queue_intercept_packet_writer writer)
{
const Packet::packet_t* packets_arr = reinterpret_cast<const Packet::packet_t*>(packets);
std::vector<Packet::packet_t> transformed_packets;
CheckNeededProfileConfigs();
rocprofiler_session_id_t session_id_snapshot = session_id;
std::shared_lock<std::shared_mutex> session_id_lock(session_id_mutex);
if (session_id.handle == 0 || session_id.handle != rocprofiler::ROCProfiler_Singleton::GetInstance().GetCurrentSessionId().handle)
{
session_id_lock.unlock();
CheckNeededProfileConfigs();
session_id_lock.lock();
}
auto& queue_info = *reinterpret_cast<Queue*>(data);
std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
if (session_id_snapshot.handle > 0 && pkt_count > 0 &&
if (session_id.handle > 0 && pkt_count > 0 &&
(is_counter_collection_mode || is_timestamp_collection_mode ||
is_pc_sampling_collection_mode) &&
session) {
// hsa_ven_amd_aqlprofile_profile_t* profile;
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles;
// Searching accross all the packets given during this write
for (size_t i = 0; i < pkt_count; ++i) {
auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
// +Skip kernel dispatch IDs not wanted
// Skip packets other than kernel dispatch packets.
if (session_id_snapshot.handle == 0 || !Packet::IsDispatchPacket(original_packet)) {
if (session_id.handle == 0 || !Packet::IsDispatchPacket(original_packet)) {
transformed_packets.emplace_back(packets_arr[i]);
continue;
}
std::unique_ptr<Packet::AQLPacketProfile> profile_packet;
// If counters found in the session
if (session_data_count > 0 && is_counter_collection_mode) {
// Get the PM4 Packets using packets_generator
profiles = Packet::InitializeAqlPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
session_data, session_id_snapshot);
replay_mode_count = profiles.size();
profile_packet = Packet::AQLPacketProfile::MoveFromCache(queue_info.GetGPUAgent());
if (!profile_packet)
profile_packet = Packet::InitializeAqlPackets(
queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), session_data, session_id);
}
uint32_t profile_id = 0;
std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*> profile;
if (session_data_count > 0 && is_counter_collection_mode) {
if (profiles.size() > 0 && replay_mode_count > 0) {
profile = profiles.at(profile_id);
hsa_signal_t ready_signal = queue_info.GetReadySignal();
hsa_signal_t block_signal = queue_info.GetBlockSignal();
/*
Creates a barrier packet with its completion signal as the
queue's ready signal.
*/
Packet::CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal);
/*
Creates a barrier packet with queue's blocksignal as its input and
completion signal.This will ensure it is no longer 0 so a later barrier
packet waiting on it to be 0 will be blocked
*/
Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal);
}
if (profile_packet.get())
{
hsa_signal_t ready_signal = queue_info.GetReadySignal();
hsa_signal_t block_signal = queue_info.GetBlockSignal();
/*
Creates a barrier packet with its completion signal as the
queue's ready signal.
*/
Packet::CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal);
/*
Creates a barrier packet with queue's blocksignal as its input and
completion signal.This will ensure it is no longer 0 so a later barrier
packet waiting on it to be 0 will be blocked
*/
Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal);
}
uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 &&
replay_mode_count > 0 && profile.first && profile.first->start_packet) {
if (session_data_count > 0 && is_counter_collection_mode && profile_packet.get())
{
auto* start_packet = profile_packet->context->start_packet;
// Adding start packet and its barrier with a dummy signal
hsa_signal_t dummy_signal{};
dummy_signal.handle = 0;
profile.first->start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets,
dummy_signal);
start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(start_packet, &transformed_packets, dummy_signal);
Packet::CreateBarrierPacket(&transformed_packets,
&profile.first->start_packet->completion_signal, nullptr);
Packet::CreateBarrierPacket(
&transformed_packets,
&start_packet->completion_signal,
nullptr
);
}
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
@@ -689,19 +689,24 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
// list to be processed by the signal interrupt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
auto* context_backup = profile_packet.get() ? profile_packet->context.get() : nullptr;
if (session) {
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (session_data_count > 0 && profile.second) {
if (session_data_count > 0 && profile_packet.get())
{
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, original_packet.completion_signal, packet.completion_signal,
session_id, buffer_id, profile.first, session_data_count, profile.second,
session_id, buffer_id, session_data_count, std::move(profile_packet),
kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id);
} else {
}
else
{
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, original_packet.completion_signal, packet.completion_signal,
session_id, buffer_id, nullptr, session_data_count, nullptr, kernel_properties,
session_id, buffer_id, session_data_count, nullptr, kernel_properties,
(uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id);
}
}
@@ -710,7 +715,7 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
// packet and create a new signal for it to get timestamps
if (original_packet.completion_signal.handle) {
hsa_barrier_and_packet_t barrier{};
barrier.header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE);
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt).completion_signal =
@@ -730,56 +735,47 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
CreateSignal(0, &interrupt_signal);
// Adding Stop and Read PM4 Packets
if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 &&
profile.first && profile.first->read_packet) {
if (session_data_count > 0 && is_counter_collection_mode && context_backup)
{
hsa_signal_t dummy_signal{};
profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets,
dummy_signal);
profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets,
interrupt_signal);
context_backup->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(context_backup->read_packet, &transformed_packets, dummy_signal);
context_backup->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(context_backup->stop_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr);
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
SignalAsyncHandler(
interrupt_signal,
new queue_info_session_t{
queue_info.GetGPUAgent(), session_id_snapshot, queue_info.GetQueueID(), writer_id,
interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(),
agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()});
} else {
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
Packet::CreateBarrierPacket(&transformed_packets, nullptr, &interrupt_signal);
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
SignalAsyncHandler(
interrupt_signal,
new queue_info_session_t{
queue_info.GetGPUAgent(), session_id_snapshot, queue_info.GetQueueID(), writer_id,
interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(),
agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()});
Packet::CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr);
}
else
Packet::CreateBarrierPacket(&transformed_packets, nullptr, &interrupt_signal);
// ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_add(1, std::memory_order_relaxed);
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
SignalAsyncHandler(
interrupt_signal,
new queue_info_session_t{
queue_info.GetGPUAgent(), session_id, queue_info.GetQueueID(), writer_id,
interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(),
agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()});
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
} else if (!is_att_collection_mode ||
!session->GetAttTracer()->ATTWriteInterceptor(packets, pkt_count, user_pkt_index,
*static_cast<Queue*>(data), writer,
buffer_id)) {
} else if (
!is_att_collection_mode||
!session ||
!session->GetAttTracer()||
!session->GetAttTracer()->ATTWriteInterceptor(
packets,
pkt_count,
user_pkt_index,
*static_cast<Queue*>(data),
writer,
buffer_id
)
) {
/* Write the original packets to the hardware queue if no profiling session is active */
writer(packets, pkt_count);
}
@@ -830,11 +826,16 @@ hsa_agent_t Queue::GetCPUAgent() { return cpu_agent_; }
uint64_t Queue::GetQueueID() { return intercept_queue_->id; }
void CheckPacketReqiurements() { Packet::CheckPacketReqiurements(); }
void CheckPacketReqiurements() {
Packet::CheckPacketReqiurements();
}
hsa_signal_t Queue::GetReadySignal() { return ready_signal_; }
hsa_signal_t Queue::GetBlockSignal() { return block_signal_; }
} // namespace queue
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/queues/queue.h
+1 -1
View File
@@ -89,7 +89,7 @@ class Queue {
hsa_signal_t GetBlockSignal();
static void ResetSessionID(rocprofiler_session_id_t id = rocprofiler_session_id_t{0});
static void CheckNeededProfileConfigs();
static bool CheckNeededProfileConfigs();
private:
static std::shared_mutex session_id_mutex;
static rocprofiler_session_id_t session_id;
projects/rocprofiler/src/core/session/profiler/profiler.cpp
+3 -3
View File
@@ -119,8 +119,8 @@ bool Profiler::HasActivePass() {
void Profiler::AddPendingSignals(
uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal,
const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id, rocprofiler::profiling_context_t* context,
uint64_t session_data_count, hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_buffer_id_t buffer_id,
uint64_t session_data_count, std::unique_ptr<Packet::AQLPacketProfile>&& profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index,
uint64_t correlation_id)
{
@@ -134,7 +134,7 @@ void Profiler::AddPendingSignals(
sessions_pending_signals_.at(writer_id).emplace_back(
new pending_signal_t{
kernel_object, original_completion_signal, new_completion_signal,
session_id_, buffer_id, context, session_data_count, profile,
session_id_, buffer_id, session_data_count, std::move(profile),
kernel_properties, thread_id, queue_index, correlation_id
}
);
projects/rocprofiler/src/core/session/profiler/profiler.h
+4 -5
View File
@@ -30,7 +30,7 @@
#include <optional>
#include <string>
#include <vector>
#include "src/core/hsa/packets/packets_generator.h"
#include "rocprofiler.h"
#include "src/core/counters/basic/basic_counter.h"
#include "src/core/counters/metrics/eval_metrics.h"
@@ -48,9 +48,8 @@ typedef struct {
hsa_signal_t new_signal;
rocprofiler_session_id_t session_id;
rocprofiler_buffer_id_t buffer_id;
rocprofiler::profiling_context_t* context;
uint64_t counters_count;
hsa_ven_amd_aqlprofile_profile_t* profile;
std::unique_ptr<::Packet::AQLPacketProfile> profile;
rocprofiler_kernel_properties_t kernel_properties;
uint32_t thread_id;
uint64_t queue_index;
@@ -73,8 +72,8 @@ class Profiler {
const hsa_signal_t& original_completion_signal,
const hsa_signal_t& new_completion_signal,
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile,
uint64_t session_data_count,
std::unique_ptr<Packet::AQLPacketProfile>&& profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
uint64_t queue_index, uint64_t correlation_id);
projects/rocprofiler/src/core/session/session.cpp
+1 -1
View File
@@ -216,7 +216,7 @@ void Session::Terminate()
GetProfiler()->WaitForPendingAndDestroy();
if (GetAttTracer())
GetAttTracer()->WaitForPendingAndDestroy();
Packet::AQLPacketProfile::WaitForProfileDeletion();
std::lock_guard<std::mutex> lock(session_lock_);
if (FindFilterWithKind(ROCPROFILER_SPM_COLLECTION)) {