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

SWDEV-432445: ATT continuous mode

Browse Source 
Change-Id: I52732fc1dba41df3373ef6a19a428b00d729bf04
This commit is contained in:
Giovanni LB
2023-11-15 12:18:51 -03:00
parent 8d80b2969d
commit aba6a1d986
29 changed files with 1841 additions and 851 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/core/hsa/hsa_support.cpp
+88 -52
View File
@@ -417,67 +417,84 @@ struct CodeObjectCallbackArg {
bool unload;
};
hsa_status_t CodeObjectCallback(hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object, void* arg) {
hsa_evt_data_t data{};
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE,
&data.codeobj.storage_type) != HSA_STATUS_SUCCESS)
#define CheckInfo(x) if ((x) != HSA_STATUS_SUCCESS) \
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE) {
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE,
&data.codeobj.storage_file) != HSA_STATUS_SUCCESS ||
data.codeobj.storage_file == -1)
hsa_status_t CodeObjectCallback(hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object, void* arg) {
using namespace std::placeholders;
hsa_evt_data_t data{};
auto codeobj_info_func = rocprofiler::HSASupport_Singleton::GetInstance()
.GetHSALoaderApi()
.hsa_ven_amd_loader_loaded_code_object_get_info;
auto codeobj_bound = std::bind(codeobj_info_func, loaded_code_object, _1, _2);
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE,
&data.codeobj.storage_type
));
if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE)
{
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE,
&data.codeobj.storage_file
));
if (data.codeobj.storage_file == -1)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
data.codeobj.memory_base = data.codeobj.memory_size = 0;
} else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY) {
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE,
&data.codeobj.memory_base) != HSA_STATUS_SUCCESS ||
hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE,
&data.codeobj.memory_size) != HSA_STATUS_SUCCESS)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
}
else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY)
{
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE,
&data.codeobj.memory_base
));
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE,
&data.codeobj.memory_size
));
data.codeobj.storage_file = -1;
} else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE) {
return HSA_STATUS_SUCCESS; // FIXME: do we really not care about these
// code objects?
} else {
}
else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE)
{
return HSA_STATUS_SUCCESS; // FIXME: do we really not care about these code objects?
}
else
{
rocprofiler::fatal("unknown code object storage type: %d", data.codeobj.storage_type);
}
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
&data.codeobj.load_base) != HSA_STATUS_SUCCESS ||
hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
&data.codeobj.load_size) != HSA_STATUS_SUCCESS ||
hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
&data.codeobj.load_delta) != HSA_STATUS_SUCCESS)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
&data.codeobj.load_base
));
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
&data.codeobj.load_size
));
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
&data.codeobj.load_delta
));
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
&data.codeobj.uri_length) != HSA_STATUS_SUCCESS)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
&data.codeobj.uri_length
));
std::string uri_str(data.codeobj.uri_length, '\0');
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, uri_str.data()) !=
HSA_STATUS_SUCCESS)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT,
&data.codeobj.agent) != HSA_STATUS_SUCCESS)
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI,
uri_str.data()
));
CheckInfo(codeobj_bound(
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT,
&data.codeobj.agent
));
data.codeobj.uri = uri_str.c_str();
data.codeobj.unload = *static_cast<bool*>(arg) ? 1 : 0;
@@ -486,8 +503,8 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
if (data.codeobj.unload)
codeobj_capture_instance::Unload(data.codeobj.load_base);
else
codeobj_capture_instance::Load(data.codeobj.load_base, uri_str, data.codeobj.memory_base,
data.codeobj.memory_size);
codeobj_capture_instance::Load(data.codeobj.load_base,
uri_str, data.codeobj.memory_base, data.codeobj.memory_size);
hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent,
hsa_executable_iteration_callback, &(data.codeobj.unload));
@@ -783,6 +800,7 @@ const Agent::DeviceInfo& HSAAgentInfo::GetDeviceInfo() const {
return device_info_;
}
assert("Attempting to read deviceInfo for a CPU agent");
return device_info_;
}
uint64_t HSAAgentInfo::getHandle() const { return agent_.handle; }
@@ -1037,7 +1055,6 @@ void HSASupport_Singleton::HSAInitialize(HsaApiTable* table) {
// is currently doing as well as the roctracer compatibility
// code earlier in this file.
uint32_t gpu_id = 0;
hsasupport_singleton.GetCoreApiTable().hsa_agent_get_info_fn(
agent, (hsa_agent_info_t)(HSA_AMD_AGENT_INFO_DRIVER_UID), &gpu_id);
const Agent::DeviceInfo& device_info = rocprofiler_instance.GetDeviceInfo(gpu_id);
@@ -1046,6 +1063,7 @@ void HSASupport_Singleton::HSAInitialize(HsaApiTable* table) {
agent, (hsa_agent_info_t)(HSA_AMD_AGENT_INFO_NEAREST_CPU), &nearCpuAgent);
rocprofiler::HSAAgentInfo agent_info(agent, device_type);
agent_info.SetNearCpuAgent(nearCpuAgent);
agent_info.SetDeviceInfo(device_info);
Packet::InitializeGPUPool(agent, &agent_info);
hsasupport_singleton.SetHSAAgentInfo(agent, agent_info);
@@ -1121,5 +1139,23 @@ void HSASupport_Singleton::HSAFinalize() {
FinitKsymbols();
}
#define CHECK_HSA_STATUS(msg, status) \
if ((status) != HSA_STATUS_SUCCESS && (status) != HSA_STATUS_INFO_BREAK) { \
try { \
const char* emsg = nullptr; \
hsa_status_string(status, &emsg); \
if (!emsg) emsg = "<Unknown HSA Error>"; \
std::cerr << msg << std::endl; \
std::cerr << emsg << std::endl; \
} catch (std::exception & e) { \
} \
abort(); \
}
void HSASupport_Singleton::CreateSignal(uint32_t attribute, hsa_signal_t* signal) {
auto status = GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status);
}
} // namespace rocprofiler
src/core/hsa/hsa_support.h
+1
View File
@@ -169,6 +169,7 @@ class HSASupport_Singleton {
void HSAFinalize();
void InitKsymbols();
void FinitKsymbols();
void CreateSignal(uint32_t attribute, hsa_signal_t* signal);
HSASupport_Singleton(const HSASupport_Singleton&) = delete;
HSASupport_Singleton& operator=(const HSASupport_Singleton&) = delete;
};
src/core/hsa/packets/packets_generator.cpp
+109 -4
View File
@@ -18,6 +18,8 @@
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
//#define HSA_ATT_MARKER_ENABLE
#include "packets_generator.h"
#include "src/api/rocprofiler_singleton.h"
@@ -235,14 +237,12 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
rocprofiler::Metric* metric;
if (std::find(counter_names.begin(), counter_names.end(), result->name) !=
counter_names.end()) {
// std::cout << "Counter from Result List: " << result->name << std::endl;
counters_taken.insert(result->name);
metric = const_cast<rocprofiler::Metric*>(metricsDict[gpu_agent.handle]->Get(result->name));
if (metric == nullptr) std::cout << result->name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
} else {
metrics_counters_taken.insert(result->name);
// std::cout << "Counter Added: " << result->name << std::endl;
}
}
@@ -283,7 +283,6 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
}
}
if (flag) {
// std::cout << "Counter from Result Map: " << metric_name << std::endl;
counters_taken.insert(metric_name);
rocprofiler::Metric* metric =
const_cast<rocprofiler::Metric*>(metricsDict[gpu_agent.handle]->Get(metric_name));
@@ -586,10 +585,12 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wconversion-null"
// Preparing the profile structure to get the packets
auto* params = new hsa_ven_amd_aqlprofile_parameter_t[att_params.size()+1];
memcpy(params, att_params.data(), att_params.size()*sizeof(params[0]));
hsa_ven_amd_aqlprofile_profile_t* profile =
new hsa_ven_amd_aqlprofile_profile_t{gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE,
nullptr, 0,
&att_params[0], (uint32_t)att_params.size(),
params, (uint32_t)att_params.size(),
NULL, NULL};
#pragma GCC diagnostic pop
@@ -614,4 +615,108 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
return profile;
}
// Generate ATT tracer marker packets. Also generate and return
// the descriptor object which has the PM4 buffer for inserting data
hsa_ven_amd_aqlprofile_descriptor_t
GenerateATTMarkerPackets(hsa_agent_t gpu_agent, packet_t& marker_packet, uint32_t data)
{
#ifdef HSA_ATT_MARKER_ENABLE
// Preparing the profile structure to get the packets
auto pool = rocprofiler::HSASupport_Singleton::GetInstance()
.GetHSAAgentInfo(gpu_agent.handle)
.cpu_pool_;
hsa_ven_amd_aqlprofile_descriptor_t desc{AllocateSysMemory(gpu_agent, 1024, &pool), 1024};
hsa_ven_amd_aqlprofile_profile_t profile{
gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE,
nullptr, 0,
nullptr, 0,
{}, desc
};
hsa_ven_amd_aqlprofile_att_marker_channel_t channel = HSA_VEN_AMD_AQLPROFILE_ATT_CHANNEL_2;
hsa_status_t status = hsa_ven_amd_aqlprofile_att_marker(&profile, &marker_packet, data, channel);
if (status != HSA_STATUS_SUCCESS)
{
rocprofiler::warning("Could not create ATT Marker Packets.");
desc.size = 0;
desc.ptr = nullptr;
}
return desc;
#else
return {nullptr,0};
#endif
}
void AddVendorSpecificPacket(const packet_t* packet,
std::vector<packet_t>* transformed_packets,
const hsa_signal_t& packet_completion_signal) {
transformed_packets->emplace_back(*packet).completion_signal = packet_completion_signal;
}
/*
Function name: CreateBarrierPacket.
Argument : The list of transformed packets to add the
barrier packet to. Pointer to the completion signal
and the input signal of the barrier packet to be created.
Description : This packet creates the barrier packet with the given
completion signal and dependency signal. It then adds to
the transformed packets list.
*/
void CreateBarrierPacket(std::vector<packet_t>* transformed_packets,
const hsa_signal_t* packet_dependency_signal,
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);
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;
transformed_packets->emplace_back(*reinterpret_cast<packet_t*>(barrier_ptr));
}
template <typename Integral = uint64_t> constexpr Integral bit_mask(int first, int last) {
assert(last >= first && "Error: hsa_support::bit_mask -> invalid argument");
size_t num_bits = last - first + 1;
return ((num_bits >= sizeof(Integral) * 8) ? ~Integral{0}
/* num_bits exceed the size of Integral */
: ((Integral{1} << num_bits) - 1))
<< first;
}
template <typename Integral> constexpr Integral bit_extract(Integral x, int first, int last) {
return (x >> first) & bit_mask<Integral>(0, last - first);
}
bool IsDispatchPacket(const hsa_barrier_and_packet_t& packet)
{
return bit_extract(
packet.header,
HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1
) == HSA_PACKET_TYPE_KERNEL_DISPATCH;
}
// Returns a list of pointers to dispatch packets.
std::vector<const hsa_kernel_dispatch_packet_s*> ExtractDispatchPackets(
const void* packets,
int pkt_count
) {
std::vector<const hsa_kernel_dispatch_packet_s*> ret;
for (int i = 0; i < pkt_count; ++i) {
auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
// Skip packets other than kernel dispatch packets.
if (!IsDispatchPacket(original_packet)) continue;
ret.push_back(static_cast<const hsa_kernel_dispatch_packet_s*>(packets)+i);
}
return ret;
}
} // namespace Packet
src/core/hsa/packets/packets_generator.h
+17
View File
@@ -60,6 +60,8 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
std::vector<hsa_ven_amd_aqlprofile_parameter_t>& att_params, packet_t* start_packet,
packet_t* stop_packet, size_t att_buffer_size);
hsa_ven_amd_aqlprofile_descriptor_t
GenerateATTMarkerPackets(hsa_agent_t gpu_agent, packet_t& marker_packet, uint32_t data);
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
@@ -74,6 +76,21 @@ typedef struct {
att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent);
void AddVendorSpecificPacket(const packet_t* packet,
std::vector<packet_t>* transformed_packets,
const hsa_signal_t& packet_completion_signal);
void CreateBarrierPacket(std::vector<packet_t>* transformed_packets,
const hsa_signal_t* packet_dependency_signal,
const hsa_signal_t* packet_completion_signal);
bool IsDispatchPacket(const hsa_barrier_and_packet_t& packet);
// Returns a list of pointers to dispatch packets.
std::vector<const hsa_kernel_dispatch_packet_s*> ExtractDispatchPackets(
const void* packets,
int pkt_count
);
} // namespace Packet
#endif // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_
src/core/hsa/queues/queue.cpp
+67 -492
View File
@@ -29,6 +29,7 @@
#include <utility>
#include <algorithm>
#include <numa.h>
#include <unordered_map>
#include "rocprofiler.h"
#include "src/api/rocprofiler_singleton.h"
@@ -55,8 +56,6 @@
#define __NR_gettid 186
#define DEFAULT_ATT_BUFFER_SIZE 0x40000000
std::mutex sessions_pending_signal_lock;
namespace rocprofiler {
@@ -274,17 +273,6 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
return status;
}
hsa_status_t attTraceDataCallback(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;
att_trace_callback_data_t* passed_data = reinterpret_cast<att_trace_callback_data_t*>(data);
passed_data->push_back(*info_data);
// TODO: clear output buffers after copying
// either copy here or in ::AddAttRecord
return status;
}
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) {
record->counters_count =
rocprofiler_record_counters_instances_count_t{pending->context->metrics_list.size()};
@@ -323,52 +311,6 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
}
}
void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent,
att_pending_signal_t& pending) {
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
HSAAgentInfo agent_info = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle);
att_trace_callback_data_t data;
hsa_status_t status =
hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data);
if ((status & HSA_STATUS_ERROR_OUT_OF_RESOURCES) == HSA_STATUS_ERROR_OUT_OF_RESOURCES)
rocprofiler::warning("Warning: ATT buffer full!\n");
if ((status & HSA_STATUS_ERROR_EXCEPTION) == HSA_STATUS_ERROR_EXCEPTION)
rocprofiler::warning("Warning: ATT received a UTC memory error!\n");
if (status == HSA_STATUS_ERROR) fatal("Thread Trace Error!");
// Allocate memory for shader_engine_data
record->shader_engine_data = static_cast<rocprofiler_record_se_att_data_t*>(
calloc(data.size(), sizeof(rocprofiler_record_se_att_data_t)));
att_trace_callback_data_t::iterator trace_data_it;
uint32_t se_index = 0;
// iterate over the trace data collected from each shader engine
for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) {
const void* data_ptr = trace_data_it->trace_data.ptr;
const uint32_t data_size = trace_data_it->trace_data.size;
void* buffer = NULL;
if (data_size != 0) {
// Allocate buffer on CPU to copy out trace data
buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool_);
if (buffer == NULL) fatal("Trace data buffer allocation failed");
auto status =
hsasupport_singleton.GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size);
if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed");
record->shader_engine_data[se_index].buffer_ptr = buffer;
record->shader_engine_data[se_index].buffer_size = data_size;
++se_index;
// TODO: clear output buffers after copying
}
}
record->shader_engine_data_count = data.size();
}
/*
Function name: enable_dispatch
Argument : pointer to the the Queue class object
@@ -574,104 +516,6 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
return false;
}
bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
auto queue_info_session = static_cast<queue_info_session_t*>(data);
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton =
rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
if (!queue_info_session || !rocprofiler_singleton.GetSession(queue_info_session->session_id) ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetAttTracer())
return true;
rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id);
rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer();
if (!session->GetAttTracer()) return true;
std::vector<att_pending_signal_t>& pending_signals =
const_cast<std::vector<att_pending_signal_t>&>(
att_tracer->GetPendingSignals(queue_info_session->writer_id));
if (!pending_signals.empty()) {
for (auto it = pending_signals.begin(); it != pending_signals.end();
it = pending_signals.erase(it)) {
auto& pending = *it;
std::lock_guard<std::mutex> lock(session->GetSessionLock());
if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_signal))
return true;
rocprofiler_record_att_tracer_t record{};
record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
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.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id};
record.writer_id = queue_info_session->writer_id;
if (/*pending.counters_count > 0 && */ pending.profile) {
AddAttRecord(&record, queue_info_session->agent, pending);
}
// July/01/2023 -> Changed this to queue_info_session->writer_id
// so we can correlate to dispatches. kernel_id already has the descriptor.
record.header = {ROCPROFILER_ATT_TRACER_RECORD,
rocprofiler_record_id_t{pending.kernel_descriptor}};
record.intercept_list = codeobj_record::get_capture(record.header.id);
std::atomic_thread_fence(std::memory_order_release);
if (pending.session_id.handle == 0) {
pending.session_id = rocprofiler_singleton.GetCurrentSessionId();
}
if (session->FindBuffer(pending.buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
buffer->AddRecord(record);
buffer->Flush();
}
codeobj_record::free_capture(record.header.id);
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);
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;
}
}
delete queue_info_session;
std::atomic_thread_fence(std::memory_order_seq_cst);
return false;
}
/*
Function name: CreateBarrierPacket.
Argument : The list of transformed packets to add the
barrier packet to. Pointer to the completion signal
and the input signal of the barrier packet to be created.
Description : This packet creates the barrier packet with the given
completion signal and dependency signal. It then adds to
the transformed packets list.
*/
void CreateBarrierPacket(std::vector<Packet::packet_t>* transformed_packets,
const hsa_signal_t* packet_dependency_signal,
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);
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;
transformed_packets->emplace_back(*reinterpret_cast<Packet::packet_t*>(barrier_ptr));
}
void AddVendorSpecificPacket(const Packet::packet_t* packet,
std::vector<Packet::packet_t>* transformed_packets,
const hsa_signal_t& packet_completion_signal) {
transformed_packets->emplace_back(*packet).completion_signal = packet_completion_signal;
}
void SignalAsyncHandler(const hsa_signal_t& signal, void* data) {
hsa_status_t status =
@@ -680,32 +524,8 @@ void SignalAsyncHandler(const hsa_signal_t& signal, void* data) {
CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler failed", status);
}
void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) {
hsa_status_t status =
HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data);
CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler for ATT failed", status);
}
void CreateSignal(uint32_t attribute, hsa_signal_t* signal) {
hsa_status_t status =
HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_create_fn(
1, 0, nullptr, attribute, signal);
CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status);
}
template <typename Integral = uint64_t> constexpr Integral bit_mask(int first, int last) {
assert(last >= first && "Error: hsa_support::bit_mask -> invalid argument");
size_t num_bits = last - first + 1;
return ((num_bits >= sizeof(Integral) * 8) ? ~Integral{0}
/* num_bits exceed the size of Integral */
: ((Integral{1} << num_bits) - 1))
<< first;
}
/* Extract bits [last:first] from t. */
template <typename Integral> constexpr Integral bit_extract(Integral x, int first, int last) {
return (x >> first) & bit_mask<Integral>(0, last - first);
HSASupport_Singleton::GetInstance().CreateSignal(attribute, signal);
}
rocprofiler_session_id_t session_id = rocprofiler_session_id_t{0};
@@ -720,11 +540,8 @@ bool is_counter_collection_mode = false;
bool is_timestamp_collection_mode = false;
bool is_att_collection_mode = false;
bool is_pc_sampling_collection_mode = false;
std::vector<rocprofiler_att_parameter_t> att_parameters_data;
uint32_t replay_mode_count = 0;
std::vector<std::string> kernel_profile_names;
std::vector<uint64_t> kernel_profile_dispatch_ids;
std::vector<std::string> att_counters_names;
rocprofiler::Session* session = nullptr;
@@ -737,162 +554,52 @@ void CheckNeededProfileConfigs() {
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;
if (session_id.handle == 0 || internal_session_id.handle != session_id.handle) {
session_id = internal_session_id;
// Getting Counters count from the Session
if (session_id.handle > 0) {
session = rocprofiler_singleton.GetSession(session_id);
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
session_data = filter->GetCounterData();
is_counter_collection_mode = true;
session_data_count = session_data.size();
buffer_id = filter->GetBufferId();
} 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);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
buffer_id = filter->GetBufferId();
} else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
att_parameters_data = filter->GetAttParametersData();
is_att_collection_mode = true;
buffer_id =
session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))
->GetBufferId();
// Getting Counters count from the Session
session = rocprofiler_singleton.GetSession(session_id);
att_counters_names = filter->GetCounterData();
kernel_profile_names = std::get<std::vector<std::string>>(
filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES));
kernel_profile_dispatch_ids =
std::get<std::vector<uint64_t>>(filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS));
} else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
is_pc_sampling_collection_mode = true;
}
}
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
session_data = filter->GetCounterData();
is_counter_collection_mode = true;
session_data_count = session_data.size();
buffer_id = filter->GetBufferId();
} 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);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
buffer_id = filter->GetBufferId();
} else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
auto* att_tracer = session->GetAttTracer();
att_tracer->SetParameters(filter->GetAttParametersData());
is_att_collection_mode = true;
buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))
->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)
));
} else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
is_pc_sampling_collection_mode = true;
}
}
static int KernelInterceptCount = 0;
std::atomic<uint32_t> WRITER_ID{0};
std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, int pkt_count) {
std::vector<bool> can_profile_packet;
bool b_can_profile_anypacket = false;
can_profile_packet.reserve(pkt_count);
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
assert(hsasupport_singleton.ksymbols);
uint32_t current_writer_id = WRITER_ID.load(std::memory_order_relaxed);
for (int i = 0; i < pkt_count; ++i) {
auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
bool b_profile_this_object = false;
// Skip packets other than kernel dispatch packets.
if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) ==
HSA_PACKET_TYPE_KERNEL_DISPATCH) {
auto& kdispatch = static_cast<const hsa_kernel_dispatch_packet_s*>(packets)[i];
// If Dispatch IDs specified, profile based on dispatch ID
for (auto id : kernel_profile_dispatch_ids) b_profile_this_object |= id == current_writer_id;
try {
// Can throw
const std::string& kernel_name = hsasupport_singleton.ksymbols->at(kdispatch.kernel_object);
// If no filters specified, auto profile this kernel
if (kernel_profile_names.size() == 0 && kernel_profile_dispatch_ids.size() == 0 &&
kernel_name.find("__amd_rocclr_") == std::string::npos)
b_profile_this_object = true;
// Try to match the mangled kernel name with given matches in input.txt
// We want to initiate att profiling if a match exists
for (const std::string& kernel_matches : kernel_profile_names)
if (kernel_name.find(kernel_matches) != std::string::npos) b_profile_this_object = true;
} catch (...) {
rocprofiler::warning("Warning: Unknown name for object %lu\n", kdispatch.kernel_object);
}
current_writer_id += 1;
}
b_can_profile_anypacket |= b_profile_this_object;
can_profile_packet.push_back(b_profile_this_object);
}
// If we're going to skip all packets, need to update writer ID
if (!b_can_profile_anypacket) WRITER_ID.store(current_writer_id, std::memory_order_release);
return {can_profile_packet, b_can_profile_anypacket};
}
std::pair<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t> ProcessATTParams(
Packet::packet_t& start_packet, Packet::packet_t& stop_packet, Queue& queue_info,
rocprofiler::HSAAgentInfo& agentInfo) {
std::vector<hsa_ven_amd_aqlprofile_parameter_t> att_params;
int num_att_counters = 0;
uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE;
rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY;
for (rocprofiler_att_parameter_t& param : att_parameters_data) {
switch (param.parameter_name) {
case ROCPROFILER_ATT_PERFCOUNTER_NAME:
break;
case ROCPROFILER_ATT_CAPTURE_MODE:
capture_mode = static_cast<rocprofiler_codeobj_capture_mode_t>(param.value);
break;
case ROCPROFILER_ATT_BUFFER_SIZE:
att_buffer_size =
std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20)));
break; // Clip to [96KB, 4GB)
case ROCPROFILER_ATT_PERFCOUNTER:
num_att_counters += 1;
break;
default:
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(param.parameter_name)),
param.value});
}
}
if (att_counters_names.size() > 0) {
MetricsDict* metrics_dict_ = MetricsDict::Create(&agentInfo);
for (const std::string& counter_name : att_counters_names) {
const Metric* metric = metrics_dict_->Get(counter_name);
const BaseMetric* base = dynamic_cast<const BaseMetric*>(metric);
if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str());
std::vector<const counter_t*> counters;
base->GetCounters(counters);
hsa_ven_amd_aqlprofile_event_t event = counters[0]->event;
if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ)
rocprofiler::fatal("Only events from the SQ block can be selected for ATT.\n");
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)),
event.counter_id | (event.counter_id ? (0xF << 24) : 0)});
num_att_counters += 1;
}
hsa_ven_amd_aqlprofile_parameter_t zero_perf = {
static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)), 0};
// Fill other perfcounters with 0's
for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf);
}
// Get the PM4 Packets using packets_generator
return {Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), att_params,
&start_packet, &stop_packet, att_buffer_size),
capture_mode};
}
/**
* @brief This function is a queue write interceptor. It intercepts the
* packet write function. Creates an instance of packet class with the raw
@@ -902,23 +609,19 @@ std::pair<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
*/
void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer) {
static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES");
static int MAX_ATT_PROFILES = env_MAX_ATT_PROFILES ? atoi(env_MAX_ATT_PROFILES) : 1;
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;
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 &&
(is_counter_collection_mode || is_timestamp_collection_mode ||
is_pc_sampling_collection_mode) &&
session) {
// Getting Queue Data and Information
auto& queue_info = *reinterpret_cast<Queue*>(data);
std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
// hsa_ven_amd_aqlprofile_profile_t* profile;
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
@@ -929,10 +632,7 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
// +Skip kernel dispatch IDs not wanted
// Skip packets other than kernel dispatch packets.
if (session_id_snapshot.handle == 0 ||
bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) !=
HSA_PACKET_TYPE_KERNEL_DISPATCH) {
if (session_id_snapshot.handle == 0 || !Packet::IsDispatchPacket(original_packet)) {
transformed_packets.emplace_back(packets_arr[i]);
continue;
}
@@ -959,13 +659,13 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
Creates a barrier packet with its completion signal as the
queue's ready signal.
*/
CreateBarrierPacket(&transformed_packets, nullptr, &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
*/
CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal);
Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal);
}
@@ -978,10 +678,13 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
dummy_signal.handle = 0;
profile.first->start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal);
Packet::AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal);
CreateBarrierPacket(&transformed_packets, &profile.first->start_packet->completion_signal,
nullptr);
Packet::CreateBarrierPacket(
&transformed_packets,
&profile.first->start_packet->completion_signal,
nullptr
);
}
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
@@ -1043,18 +746,18 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
hsa_signal_t dummy_signal{};
profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal);
Packet::AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal);
profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, interrupt_signal);
Packet::AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr);
Packet::CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr);
}
else
CreateBarrierPacket( &transformed_packets, nullptr, &interrupt_signal);
Packet::CreateBarrierPacket( &transformed_packets, nullptr, &interrupt_signal);
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
@@ -1073,142 +776,15 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
} else if (session_id_snapshot.handle > 0 && pkt_count > 0 && is_att_collection_mode && session &&
KernelInterceptCount < MAX_ATT_PROFILES) {
// att start
// Getting Queue Data and Information
auto& queue_info = *static_cast<Queue*>(data);
std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
bool can_profile_anypacket = false;
std::vector<bool> can_profile_packet;
std::tie(can_profile_packet, can_profile_anypacket) =
GetAllowedProfilesList(packets, pkt_count);
if (!can_profile_anypacket) {
/* Write the original packets to the hardware if no patch will be profiled */
writer(packets, pkt_count);
return;
}
// Preparing att Packets
Packet::packet_t start_packet{};
Packet::packet_t stop_packet{};
hsa_ven_amd_aqlprofile_profile_t* profile = nullptr;
rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY;
if (att_parameters_data.size() > 0) {
std::tie(profile, capture_mode) =
ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo);
}
// Searching across 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];
uint32_t writer_id = 0;
// Skip all packets marked with !can_profile
if (i >= can_profile_packet.size() || can_profile_packet[i] == false) {
transformed_packets.emplace_back(packets_arr[i]);
// increment writer ID for every packet
if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) ==
HSA_PACKET_TYPE_KERNEL_DISPATCH)
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
continue;
}
KernelInterceptCount += 1;
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
if (!att_parameters_data.empty() && profile) {
// Adding start packet and its barrier with a dummy signal
hsa_signal_t dummy_signal{};
dummy_signal.handle = 0;
start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal);
CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr) ;
}
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
// Adding the dispatch packet newly created signal to the pending signals
// list to be processed by the signal interrupt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
session->GetAttTracer()->AddPendingSignals(
writer_id, record_id, original_packet.completion_signal, packet.completion_signal,
session_id_snapshot, buffer_id, profile, kernel_properties,
(uint32_t)syscall(__NR_gettid), user_pkt_index);
uint64_t off = dispatch_packet.kernel_object +
GetKernelCode(dispatch_packet.kernel_object)->kernel_code_entry_byte_offset;
codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capture_mode, off);
uint64_t IsGFX9 = HSASupport_Singleton::GetInstance()
.GetHSAAgentInfo(queue_info.GetGPUAgent().handle)
.GetDeviceInfo()
.getName()
.find("gfx9") != std::string::npos;
codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capture_mode, IsGFX9 | (off<<1));
codeobj_record::start_capture(rocprofiler_record_id_t{record_id});
codeobj_record::stop_capture(rocprofiler_record_id_t{record_id});
// Make a copy of the original packet, adding its signal to a barrier packet
if (original_packet.completion_signal.handle != 0U) {
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.dep_signal[0] = packet.completion_signal;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt).completion_signal =
original_packet.completion_signal;
}
// Adding a barrier packet with the original packet's completion signal.
hsa_signal_t interrupt_signal;
CreateSignal(0, &interrupt_signal);
// Adding Stop PM4 Packets
if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) {
stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr);
} else {
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.completion_signal = interrupt_signal;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt);
}
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
signalAsyncHandlerATT(
interrupt_signal,
new queue_info_session_t{queue_info.GetGPUAgent(), session_id_snapshot,
queue_info.GetQueueID(), writer_id, interrupt_signal});
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
// ATT end
} else {
/* Write the original packets to the hardware queue if no profiling session
* is active */
} else if (!is_att_collection_mode || !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);
}
}
@@ -1227,7 +803,6 @@ Queue::~Queue() {
{
profiler_serializer_t& serializer =
rocprofiler::ROCProfiler_Singleton::GetInstance().GetSerializer();
// std::cout << GetROCMToolObj()->
std::lock_guard<std::mutex> serializer_lock(serializer.serializer_mutex);
for (auto it = serializer.dispatch_ready.begin(); it != serializer.dispatch_ready.end();) {
if ((*it)->GetQueueID() == GetQueueID()) {
src/core/hsa/queues/queue.h
+6
View File
@@ -70,6 +70,12 @@ class Queue {
static void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer);
static bool ATTWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer);
static bool ATTSingleWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer);
static bool ATTContiguousWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer);
hsa_queue_t* GetCurrentInterceptQueue();
hsa_agent_t GetGPUAgent();
hsa_agent_t GetCPUAgent();
src/core/isa_capture/code_object_track.cpp
+14 -13
View File
@@ -48,14 +48,11 @@ std::mutex codeobj_record::mutex;
std::unordered_map<uint64_t, CodeobjPtr> codeobj_record::codeobjs{};
std::unordered_map<uint64_t, codeobj_record::RecordInstance> codeobj_record::record_id_map{};
std::unordered_set<codeobj_record*> codeobj_record::listeners;
std::atomic<uint32_t> codeobj_capture_instance::loadcount{0};
// Codeobj Record
codeobj_record::codeobj_record(rocprofiler_codeobj_capture_mode_t mode) : capture_mode(mode){};
static uint64_t timestamp() {
return rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value;
}
void codeobj_record::start_capture() {
listeners.insert(this);
for (auto& [addr, capture] : codeobjs) this->addcapture(capture);
@@ -72,10 +69,18 @@ void codeobj_record::stop_capture() {
}
// Codeobj Capture
void codeobj_capture_instance::Load(uint64_t addr, const std::string& URI, uint64_t mem_addr,
uint64_t size) {
void codeobj_capture_instance::Load(
uint64_t addr,
const std::string& URI,
uint64_t mem_addr,
uint64_t size
) {
uint32_t id = loadcount.fetch_add(1, std::memory_order_relaxed)+1;
auto time = rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value;
std::lock_guard<std::mutex> lock(codeobj_record::mutex);
auto inst = std::make_shared<codeobj_capture_instance>(addr, URI, mem_addr, size, timestamp());
auto inst = std::make_shared<codeobj_capture_instance>(addr, URI, mem_addr, size, time, id);
codeobj_record::codeobjs[addr] = inst;
for (auto* listen : codeobj_record::listeners) listen->addcapture(inst);
}
@@ -85,7 +90,8 @@ void codeobj_capture_instance::Unload(uint64_t addr) {
if (codeobj_record::codeobjs.find(addr) == codeobj_record::codeobjs.end()) return;
codeobj_record::codeobjs.at(addr)->end_time = timestamp();
auto time = rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value;
codeobj_record::codeobjs.at(addr)->end_time = time;
codeobj_record::codeobjs.erase(addr);
}
@@ -179,11 +185,6 @@ std::pair<size_t, size_t> codeobj_capture_instance::parse_uri() {
return {offset, size};
}
codeobj_capture_instance::codeobj_capture_instance(uint64_t _addr, const std::string& _uri,
uint64_t mem_addr, uint64_t mem_size,
uint64_t start_time)
: addr(_addr), start_time(start_time), URI(_uri), mem_addr(mem_addr), mem_size(mem_size) {};
void codeobj_capture_instance::setmode(rocprofiler_codeobj_capture_mode_t mode) {
// Only reset when needed & check if codeobj was not unloaded
if (static_cast<int>(mode) > capture_mode) reset(mode);
src/core/isa_capture/code_object_track.hpp
+9 -2
View File
@@ -35,22 +35,29 @@
class codeobj_capture_instance {
public:
codeobj_capture_instance(uint64_t _addr, const std::string& _uri, uint64_t mem_addr,
uint64_t mem_size, uint64_t start_time);
uint64_t mem_size, uint64_t start_time, uint32_t id)
: addr(_addr), start_time(start_time), URI(_uri),
mem_addr(mem_addr), mem_size(mem_size), load_id(id) {};
void setmode(rocprofiler_codeobj_capture_mode_t mode);
rocprofiler_intercepted_codeobj_t get() const {
const char* buf_ptr = buffer.size() ? buffer.data() : nullptr;
return {URI.c_str(), addr, mem_size, buf_ptr, buffer.size(), start_time, end_time};
return {URI.c_str(), addr, mem_size, buf_ptr, buffer.size(), start_time, end_time, load_id};
};
const uint64_t addr;
const uint64_t start_time;
const uint32_t load_id;
static void Load(uint64_t addr, const std::string& URI, uint64_t mem_addr, uint64_t mem_size);
static void Unload(uint64_t addr);
static uint32_t GetLoadCount() { return loadcount.load(std::memory_order_relaxed); }
private:
//! 32 bits ID because this is the natural channel width for ATT Markers.
//! There is no world in which 4 billions markers can be sent anyway.
static std::atomic<uint32_t> loadcount;
void reset(rocprofiler_codeobj_capture_mode_t mode);
std::pair<size_t, size_t> parse_uri();
src/core/proxy_queue.h
+2 -2
View File
@@ -29,8 +29,8 @@ THE SOFTWARE.
#include <map>
#include <mutex>
#include "core/queue.h"
#include "core/types.h"
#include "src/core/queue.h"
#include "src/core/types.h"
struct HsaApiTable;
src/core/session/att/att.cpp
+282 -9
View File
@@ -22,6 +22,10 @@
#include <cassert>
#include <atomic>
#include "src/core/hsa/packets/packets_generator.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/core/isa_capture/code_object_track.hpp"
namespace rocprofiler {
namespace att {
@@ -31,16 +35,33 @@ AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t
: buffer_id_(buffer_id), filter_id_(filter_id), session_id_(session_id) {}
void AttTracer::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, hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index) {
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,
hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties,
uint32_t thread_id, uint64_t queue_index
) {
std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end())
sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>());
sessions_pending_signals_.at(writer_id).emplace_back(att_pending_signal_t{
kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id,
profile, kernel_properties, thread_id, queue_index});
auto pending = sessions_pending_signals_.find(writer_id);
if (pending == sessions_pending_signals_.end())
pending = sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>()).first;
pending->second.emplace_back(att_pending_signal_t{
kernel_object,
original_completion_signal,
new_completion_signal,
session_id_,
buffer_id,
profile,
kernel_properties,
thread_id,
queue_index
});
}
const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t writer_id) {
@@ -50,6 +71,258 @@ const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t w
return sessions_pending_signals_.at(writer_id);
}
#define DEFAULT_ATT_BUFFER_SIZE 0x40000000
std::pair<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
AttTracer::ProcessATTParams(
hsa_ext_amd_aql_pm4_packet_t& start_packet,
hsa_ext_amd_aql_pm4_packet_t& stop_packet,
queue::Queue& queue_info,
rocprofiler::HSAAgentInfo& agentInfo
) {
std::vector<hsa_ven_amd_aqlprofile_parameter_t> att_params;
int num_att_counters = 0;
uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE;
rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY;
for (rocprofiler_att_parameter_t& param : att_parameters_data) {
switch (param.parameter_name) {
case ROCPROFILER_ATT_PERFCOUNTER_NAME:
break;
case ROCPROFILER_ATT_CAPTURE_MODE:
capture_mode = static_cast<rocprofiler_codeobj_capture_mode_t>(param.value);
break;
case ROCPROFILER_ATT_BUFFER_SIZE:
att_buffer_size =
std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20)));
break; // Clip to [96KB, 4GB)
case ROCPROFILER_ATT_PERFCOUNTER:
num_att_counters += 1;
break;
default:
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(param.parameter_name)),
param.value});
}
}
if (att_counters_names.size() > 0) {
MetricsDict* metrics_dict_ = MetricsDict::Create(&agentInfo);
for (const std::string& counter_name : att_counters_names) {
const Metric* metric = metrics_dict_->Get(counter_name);
const BaseMetric* base = dynamic_cast<const BaseMetric*>(metric);
if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str());
std::vector<const counter_t*> counters;
base->GetCounters(counters);
hsa_ven_amd_aqlprofile_event_t event = counters[0]->event;
if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ)
rocprofiler::fatal("Only events from the SQ block can be selected for ATT.\n");
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)),
event.counter_id | (event.counter_id ? (0xF << 24) : 0)});
num_att_counters += 1;
}
hsa_ven_amd_aqlprofile_parameter_t zero_perf = {
static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)), 0};
// Fill other perfcounters with 0's
for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf);
}
// Get the PM4 Packets using packets_generator
return {Packet::GenerateATTPackets(
queue_info.GetCPUAgent(),
queue_info.GetGPUAgent(),
att_params,
&start_packet,
&stop_packet,
att_buffer_size
),
capture_mode};
}
bool AttTracer::ATTWriteInterceptor(
const void* packets,
uint64_t pkt_count,
uint64_t user_pkt_index,
queue::Queue& queue,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
) {
bool IsSingleDispatchMode = kernel_profile_dispatch_ids.size() == 0;
if (session_id_.handle == 0 ||
pkt_count == 0 ||
att_parameters_data.size() == 0
) return false;
if (IsSingleDispatchMode)
return ATTSingleWriteInterceptor(packets, pkt_count, user_pkt_index, queue, writer, buffer_id);
else
return ATTContiguousWriteInterceptor(packets, pkt_count, queue, writer, buffer_id);
}
void AttTracer::signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) {
hsa_status_t status =
HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data);
if (status != HSA_STATUS_SUCCESS)
rocprofiler::fatal("Error: hsa_amd_signal_async_handler for ATT failed");
}
bool AttTracer::AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
auto queue_info_session = static_cast<queue::queue_info_session_t*>(data);
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton =
rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
if (!queue_info_session || !rocprofiler_singleton.GetSession(queue_info_session->session_id) ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetAttTracer())
return true;
rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id);
std::lock_guard<std::mutex> lock(session->GetSessionLock());
rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer();
if (!session->GetAttTracer()) return true;
std::vector<att_pending_signal_t>& pending_signals =
const_cast<std::vector<att_pending_signal_t>&>(
att_tracer->GetPendingSignals(queue_info_session->writer_id));
if (!pending_signals.empty()) {
for (auto it = pending_signals.begin(); it != pending_signals.end();
it = pending_signals.erase(it)) {
auto& pending = *it;
//if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_signal))
// return true;
rocprofiler_record_att_tracer_t record{};
record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
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.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id};
record.writer_id = queue_info_session->writer_id;
if (/*pending.counters_count > 0 && */ pending.profile) {
AddAttRecord(&record, queue_info_session->agent, pending);
}
// July/01/2023 -> Changed this to queue_info_session->writer_id
// so we can correlate to dispatches. kernel_id already has the descriptor.
record.header = {ROCPROFILER_ATT_TRACER_RECORD,
rocprofiler_record_id_t{pending.kernel_descriptor}};
record.intercept_list = codeobj_record::get_capture(record.header.id);
std::atomic_thread_fence(std::memory_order_release);
if (pending.session_id.handle == 0) {
pending.session_id = rocprofiler_singleton.GetCurrentSessionId();
}
if (session->FindBuffer(pending.buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
buffer->AddRecord(record);
buffer->Flush();
}
codeobj_record::free_capture(record.header.id);
hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->output_buffer.ptr));
if (status != HSA_STATUS_SUCCESS)
rocprofiler::warning("Error: Couldn't free output buffer memory");
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->command_buffer.ptr));
if (status != HSA_STATUS_SUCCESS)
rocprofiler::warning("Error: Couldn't free command buffer memory");
if (pending.profile->parameters)
delete[] pending.profile->parameters;
delete pending.profile;
}
}
delete queue_info_session;
std::atomic_thread_fence(std::memory_order_seq_cst);
return false;
}
void AttTracer::AddAttRecord(
rocprofiler_record_att_tracer_t* record,
hsa_agent_t gpu_agent,
att_pending_signal_t& pending
) {
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
HSAAgentInfo agent_info = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle);
std::vector<hsa_ven_amd_aqlprofile_info_data_t> data;
hsa_status_t status =
hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data);
if ((status & HSA_STATUS_ERROR_OUT_OF_RESOURCES) == HSA_STATUS_ERROR_OUT_OF_RESOURCES)
rocprofiler::warning("Warning: ATT buffer full!\n");
if ((status & HSA_STATUS_ERROR_EXCEPTION) == HSA_STATUS_ERROR_EXCEPTION)
rocprofiler::warning("Warning: ATT received a UTC memory error!\n");
if (status == HSA_STATUS_ERROR) fatal("Thread Trace Error!");
// Allocate memory for shader_engine_data
record->shader_engine_data = static_cast<rocprofiler_record_se_att_data_t*>(
calloc(data.size(), sizeof(rocprofiler_record_se_att_data_t)));
std::vector<hsa_ven_amd_aqlprofile_info_data_t>::iterator trace_data_it;
uint32_t se_index = 0;
// iterate over the trace data collected from each shader engine
for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) {
const void* data_ptr = trace_data_it->trace_data.ptr;
const uint32_t data_size = trace_data_it->trace_data.size;
void* buffer = NULL;
if (data_size != 0) {
// Allocate buffer on CPU to copy out trace data
buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool_);
if (buffer == NULL) fatal("Trace data buffer allocation failed");
auto status =
hsasupport_singleton.GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size);
if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed");
record->shader_engine_data[se_index].buffer_ptr = buffer;
record->shader_engine_data[se_index].buffer_size = data_size;
++se_index;
// TODO: clear output buffers after copying
}
}
record->shader_engine_data_count = data.size();
}
hsa_status_t AttTracer::attTraceDataCallback(
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;
auto* passed_data = reinterpret_cast<std::vector<hsa_ven_amd_aqlprofile_info_data_t>*>(data);
passed_data->push_back(*info_data);
// TODO: clear output buffers after copying
// either copy here or in ::AddAttRecord
return status;
}
std::unordered_map<uint64_t, ATTRecordSignal> AttTracer::pending_stop_packets;
std::mutex AttTracer::att_enable_disable_mutex;
} // namespace att
} // namespace rocprofiler
src/core/session/att/att.h
+111 -24
View File
@@ -27,8 +27,15 @@
#include <mutex>
#include <string>
#include <vector>
#include <unordered_map>
#include <atomic>
#include "rocprofiler.h"
#include "src/utils/helper.h"
#include "src/core/proxy_queue.h"
#include "src/core/hsa/hsa_support.h"
#include "src/core/hsa/queues/queue.h"
#include "att_header.h"
namespace rocprofiler {
@@ -44,33 +51,25 @@ typedef struct {
uint64_t queue_index;
} att_pending_signal_t;
enum rocprofiler_att_isa_dump_mode {
ISA_MODE_DUMP_ALL=0,
ISA_MODE_DUMP_OBJ,
ISA_MODE_DUMP_KERNEL,
ISA_MODE_DUMP_NONE
};
union att_header_packet_t {
struct {
uint64_t reserved : 14;
uint64_t navi : 1;
uint64_t enable : 1;
uint64_t DSIMDM : 4;
uint64_t DCU : 5;
uint64_t DSA : 1;
uint64_t SEID : 6;
uint64_t isadumpmode : 3;
};
uint64_t raw;
};
namespace att {
struct ATTRecordSignal
{
size_t record_id;
size_t writer_id;
size_t last_kernel_exec;
rocprofiler_session_id_t session_id_snapshot;
hsa_ext_amd_aql_pm4_packet_t stop_packet;
};
class AttTracer {
public:
AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
rocprofiler_session_id_t session_id);
public:
AttTracer(
rocprofiler_buffer_id_t buffer_id,
rocprofiler_filter_id_t filter_id,
rocprofiler_session_id_t session_id
);
void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
const hsa_signal_t& original_completion_signal,
@@ -82,10 +81,98 @@ class AttTracer {
const std::vector<att_pending_signal_t>& GetPendingSignals(uint32_t writer_id);
private:
bool ATTWriteInterceptor(
const void* packets,
uint64_t pkt_count,
uint64_t user_pkt_index,
queue::Queue& queue_info,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
);
void InsertMarker(
std::vector<packet_t>& transformed_packets,
hsa_agent_t agent,
uint32_t data
);
void SetParameters(const std::vector<rocprofiler_att_parameter_t>& params) {
att_parameters_data = params;
}
void SetDispatchIds(const std::vector<std::pair<uint64_t,uint64_t>>& ids) {
kernel_profile_dispatch_ids = ids;
}
void SetCountersNames(const std::vector<std::string>& names) {
att_counters_names = names;
}
void SetKernelsNames(const std::vector<std::string>& names) {
kernel_profile_names = names;
}
std::optional<std::pair<size_t, size_t>> RequiresStartPacket(size_t rstart, size_t size);
static void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data);
static bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data);
static hsa_status_t attTraceDataCallback(
hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* data
);
protected:
using packet_t = hsa_ext_amd_aql_pm4_packet_t;
static std::unordered_map<uint64_t, ATTRecordSignal> pending_stop_packets;
static std::mutex att_enable_disable_mutex;
private:
uint32_t codeobj_load_cnt = 0;
static void AddAttRecord(
rocprofiler_record_att_tracer_t* record,
hsa_agent_t gpu_agent,
att_pending_signal_t& pending
);
std::pair<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
ProcessATTParams(
hsa_ext_amd_aql_pm4_packet_t& start_packet,
hsa_ext_amd_aql_pm4_packet_t& stop_packet,
queue::Queue& queue_info,
rocprofiler::HSAAgentInfo& agentInfo
);
bool ATTSingleWriteInterceptor(
const void* packets,
uint64_t pkt_count,
uint64_t user_pkt_index,
queue::Queue& queue_info,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
);
bool ATTContiguousWriteInterceptor(
const void* packets,
uint64_t pkt_count,
queue::Queue& queue_info,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
);
static void CreateSignal(uint32_t attribute, hsa_signal_t* signal) {
HSASupport_Singleton::GetInstance().CreateSignal(attribute, signal);
}
std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, int pkt_count);
rocprofiler_buffer_id_t buffer_id_;
rocprofiler_filter_id_t filter_id_;
rocprofiler_session_id_t session_id_;
std::atomic<uint32_t> WRITER_ID{1};
std::vector<std::string> kernel_profile_names;
std::vector<std::pair<uint64_t,uint64_t>> kernel_profile_dispatch_ids;
std::vector<std::string> att_counters_names;
std::vector<rocprofiler_att_parameter_t> att_parameters_data;
std::mutex sessions_pending_signals_lock_;
std::map<uint32_t, std::vector<att_pending_signal_t>> sessions_pending_signals_;
src/core/session/att/att_header.h
+38
View File
@@ -0,0 +1,38 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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
namespace rocprofiler {
union att_header_packet_t {
struct {
uint64_t reserved : 14;
uint64_t navi : 1;
uint64_t enable : 1;
uint64_t DSIMDM : 4;
uint64_t DCU : 5;
uint64_t DSA : 1;
uint64_t SEID : 6;
};
uint64_t raw;
};
}
src/core/session/att/continuous.cpp
+195
View File
@@ -0,0 +1,195 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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 "att.h"
#include <cassert>
#include <atomic>
#include "src/core/hsa/packets/packets_generator.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/core/isa_capture/code_object_track.hpp"
#define __NR_gettid 186
namespace rocprofiler {
namespace att {
void AttTracer::InsertMarker(
std::vector<packet_t>& transformed_packets,
hsa_agent_t agent,
uint32_t data
) {
packet_t marker_packet{};
auto desc = Packet::GenerateATTMarkerPackets(agent, marker_packet, data);
if (desc.ptr && desc.size)
Packet::AddVendorSpecificPacket(&marker_packet, &transformed_packets, hsa_signal_t{.handle = 0});
}
std::optional<std::pair<size_t, size_t>> AttTracer::RequiresStartPacket(size_t rstart, size_t size)
{
for (auto& r : kernel_profile_dispatch_ids)
if (rstart <= r.first && rstart+size > r.first)
return r;
return {};
}
bool AttTracer::ATTContiguousWriteInterceptor(
const void* packets,
uint64_t pkt_count,
queue::Queue& queue_info,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
) {
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(packets);
std::vector<packet_t> transformed_packets;
// att start
// Getting Queue Data and Information
rocprofiler::HSAAgentInfo& agentInfo = rocprofiler::HSASupport_Singleton::GetInstance()
.GetHSAAgentInfo(queue_info.GetGPUAgent().handle);
auto dispatchPackets = Packet::ExtractDispatchPackets(packets, pkt_count);
if (dispatchPackets.size() == 0) return false;
size_t writer_id = WRITER_ID.fetch_add(dispatchPackets.size(), std::memory_order_relaxed);
uint32_t new_load_cnt = codeobj_capture_instance::GetLoadCount();
auto bInsertStart = RequiresStartPacket(writer_id, dispatchPackets.size());
{
std::lock_guard<std::mutex> lk(att_enable_disable_mutex);
// If att_start already exists, don't start again
auto agent_pending_packets = pending_stop_packets.find(queue_info.GetGPUAgent().handle);
if (agent_pending_packets != pending_stop_packets.end())
bInsertStart = {};
// If nothing will be added or removed, return
if (!bInsertStart && codeobj_load_cnt == new_load_cnt)
{
if (
agent_pending_packets == pending_stop_packets.end() ||
agent_pending_packets->second.last_kernel_exec > writer_id + dispatchPackets.size()
)
return false;
}
}
if (bInsertStart)
{
// Preparing att Packets
packet_t start_packet{};
packet_t stop_packet{};
hsa_ven_amd_aqlprofile_profile_t* profile = nullptr;
rocprofiler_codeobj_capture_mode_t capturem = ROCPROFILER_CAPTURE_SYMBOLS_ONLY;
std::tie(profile, capturem) = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo);
if (!profile)
{
rocprofiler::warning("Failed to create profile from queue!");
return false;
}
uint64_t IsGFX9 = HSASupport_Singleton::GetInstance()
.GetHSAAgentInfo(queue_info.GetGPUAgent().handle)
.GetDeviceInfo()
.getName()
.find("gfx9") != std::string::npos;
hsa_signal_t dummy_signal{};
dummy_signal.handle = 0;
start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal);
Packet::CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr);
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID("ATT_Contiguous", record_id);
this->AddPendingSignals(
writer_id,
record_id,
start_packet.completion_signal,
start_packet.completion_signal,
session_id_,
buffer_id,
profile,
{0},
(uint32_t)syscall(__NR_gettid),
0
);
codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capturem, IsGFX9);
codeobj_record::start_capture(rocprofiler_record_id_t{record_id});
stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
std::lock_guard<std::mutex> lk(att_enable_disable_mutex);
pending_stop_packets[queue_info.GetGPUAgent().handle]
= {record_id, writer_id, bInsertStart->second, session_id_, stop_packet};
}
if (codeobj_load_cnt != new_load_cnt)
{
codeobj_load_cnt = new_load_cnt;
InsertMarker(transformed_packets, queue_info.GetGPUAgent(), new_load_cnt);
}
// Searching across all the packets given during this write
for (size_t i = 0; i < pkt_count; ++i)
transformed_packets.emplace_back(packets_arr[i]);
{
std::lock_guard<std::mutex> lk(att_enable_disable_mutex);
auto agent_pending_packets = pending_stop_packets.find(queue_info.GetGPUAgent().handle);
if (agent_pending_packets != pending_stop_packets.end() &&
agent_pending_packets->second.last_kernel_exec <= writer_id + dispatchPackets.size()
) {
const ATTRecordSignal& rsignal = agent_pending_packets->second;
// Adding a barrier packet with the original packet's completion signal.
hsa_signal_t interrupt_signal;
CreateSignal(0, &interrupt_signal);
// Adding Stop PM4 Packets
Packet::AddVendorSpecificPacket(&rsignal.stop_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr);
// Creating Async Handler to be called every time the interrupt signal is marked complete
signalAsyncHandlerATT(interrupt_signal, new queue::queue_info_session_t{
queue_info.GetGPUAgent(),
rsignal.session_id_snapshot,
queue_info.GetQueueID(),
rsignal.writer_id,
interrupt_signal
});
codeobj_record::stop_capture(rocprofiler_record_id_t{rsignal.record_id});
pending_stop_packets.erase(queue_info.GetGPUAgent().handle);
}
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
return true;
}
} // namespace att
} // namespace rocprofiler
src/core/session/att/singlekernel.cpp
+233
View File
@@ -0,0 +1,233 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
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 "att.h"
#include <cassert>
#include <atomic>
#include "src/core/hsa/packets/packets_generator.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/core/isa_capture/code_object_track.hpp"
#define __NR_gettid 186
namespace rocprofiler {
namespace att {
std::pair<std::vector<bool>, bool>
AttTracer::GetAllowedProfilesList(const void* packets, int pkt_count) {
std::vector<bool> can_profile_packet;
bool b_can_profile_anypacket = false;
can_profile_packet.reserve(pkt_count);
rocprofiler::HSASupport_Singleton& hsasupport_singleton =
rocprofiler::HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
assert(hsasupport_singleton.ksymbols);
uint32_t current_writer_id = WRITER_ID.load(std::memory_order_relaxed);
for (int i = 0; i < pkt_count; ++i) {
auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
bool b_profile_this_object = false;
// Skip packets other than kernel dispatch packets.
if (Packet::IsDispatchPacket(original_packet)) {
auto& kdispatch = static_cast<const hsa_kernel_dispatch_packet_s*>(packets)[i];
try {
// Can throw
const std::string& kernel_name = hsasupport_singleton.ksymbols->at(kdispatch.kernel_object);
// If no filters specified, auto profile this kernel
if (kernel_profile_names.size() == 0 &&
kernel_name.find("__amd_rocclr_") == std::string::npos
) {
b_profile_this_object = true;
} else {
// Try to match the mangled kernel name with given matches in input.txt
// We want to initiate att profiling if a match exists
for (const std::string& kernel_matches : kernel_profile_names)
if (kernel_name.find(kernel_matches) != std::string::npos) b_profile_this_object = true;
}
} catch (...) {
rocprofiler::warning("Warning: Unknown name for object %lu\n", kdispatch.kernel_object);
}
current_writer_id += 1;
}
b_can_profile_anypacket |= b_profile_this_object;
can_profile_packet.push_back(b_profile_this_object);
}
// If we're going to skip all packets, need to update writer ID
if (!b_can_profile_anypacket)
WRITER_ID.store(current_writer_id, std::memory_order_release);
return {can_profile_packet, b_can_profile_anypacket};
}
bool AttTracer::ATTSingleWriteInterceptor(
const void* packets,
uint64_t pkt_count,
uint64_t user_pkt_index,
queue::Queue& queue_info,
hsa_amd_queue_intercept_packet_writer writer,
rocprofiler_buffer_id_t buffer_id
) {
static int KernelInterceptCount = 0;
static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES");
static int MAX_ATT_PROFILES = env_MAX_ATT_PROFILES ? atoi(env_MAX_ATT_PROFILES) : 1;
if (KernelInterceptCount >= MAX_ATT_PROFILES) return false;
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(packets);
std::vector<packet_t> transformed_packets;
// att start
// Getting Queue Data and Information
rocprofiler::HSAAgentInfo& agentInfo =
rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
queue_info.GetGPUAgent().handle);
bool can_profile_anypacket = false;
std::vector<bool> can_profile_packet;
std::tie(can_profile_packet, can_profile_anypacket) = GetAllowedProfilesList(packets, pkt_count);
if (!can_profile_anypacket) return false;
// Preparing att Packets
packet_t start_packet{};
packet_t stop_packet{};
hsa_ven_amd_aqlprofile_profile_t* profile = nullptr;
rocprofiler_codeobj_capture_mode_t capturem = ROCPROFILER_CAPTURE_SYMBOLS_ONLY;
std::tie(profile, capturem) = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo);
// Searching across 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];
uint32_t writer_id = 0;
// Skip all packets marked with !can_profile
if (i >= can_profile_packet.size() || can_profile_packet[i] == false) {
transformed_packets.emplace_back(packets_arr[i]);
// increment writer ID for every packet
if (Packet::IsDispatchPacket(original_packet))
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
continue;
}
KernelInterceptCount += 1;
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
if (profile) {
// Adding start packet and its barrier with a dummy signal
hsa_signal_t dummy_signal{};
dummy_signal.handle = 0;
start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal);
Packet::CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr) ;
}
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
// Adding the dispatch packet newly created signal to the pending signals
// list to be processed by the signal interrupt
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id);
this->AddPendingSignals(
writer_id,
record_id,
original_packet.completion_signal,
packet.completion_signal,
session_id_,
buffer_id,
profile,
{0},
(uint32_t)syscall(__NR_gettid),
user_pkt_index
);
uint64_t IsGFX9 = HSASupport_Singleton::GetInstance()
.GetHSAAgentInfo(queue_info.GetGPUAgent().handle)
.GetDeviceInfo()
.getName()
.find("gfx9") != std::string::npos;
codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capturem, IsGFX9);
codeobj_record::start_capture(rocprofiler_record_id_t{record_id});
codeobj_record::stop_capture(rocprofiler_record_id_t{record_id});
// Make a copy of the original packet, adding its signal to a barrier packet
if (original_packet.completion_signal.handle != 0U) {
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.dep_signal[0] = packet.completion_signal;
packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt).completion_signal =
original_packet.completion_signal;
}
// Adding a barrier packet with the original packet's completion signal.
hsa_signal_t interrupt_signal;
CreateSignal(0, &interrupt_signal);
// Adding Stop PM4 Packets
if (profile) {
stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
Packet::AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr);
} else {
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.completion_signal = interrupt_signal;
packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt);
}
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
signalAsyncHandlerATT(interrupt_signal, new queue::queue_info_session_t{
queue_info.GetGPUAgent(),
session_id_,
queue_info.GetQueueID(),
writer_id,
interrupt_signal
});
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
return true;
}
} // namespace att
} // namespace rocprofiler
src/core/session/filter.cpp
+3 -4
View File
@@ -195,7 +195,7 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) {
case ROCPROFILER_FILTER_DISPATCH_IDS:
dispatch_id_filter_.clear();
for (uint32_t j = 0; j < property.data_count; j++)
dispatch_id_filter_.emplace_back(property.dispatch_ids[j]);
dispatch_id_filter_.push_back({property.dispatch_ids[j].start, property.dispatch_ids[j].end});
break;
default:
break;
@@ -205,9 +205,8 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) {
// "profiler mode!\n");
}
}
std::variant<std::vector<std::string>, uint32_t*, std::vector<uint64_t>> Filter::GetProperty(
rocprofiler_filter_property_kind_t kind) {
std::variant<std::vector<std::string>, uint32_t*, std::vector<uint64_t>> property;
Filter::filter_property_variant_t Filter::GetProperty(rocprofiler_filter_property_kind_t kind) {
filter_property_variant_t property;
switch (kind) {
case ROCPROFILER_FILTER_GPU_NAME: {
property = agent_names_;
src/core/session/filter.h
+7 -3
View File
@@ -33,6 +33,11 @@ namespace rocprofiler {
class Filter {
public:
typedef std::variant<
std::vector<std::string>,
uint32_t*,
std::vector<std::pair<uint64_t,uint64_t>>
> filter_property_variant_t;
Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind,
rocprofiler_filter_data_t filter_data, uint64_t data_count);
~Filter();
@@ -53,8 +58,7 @@ class Filter {
bool HasCallback();
void SetProperty(rocprofiler_filter_property_t property);
std::variant<std::vector<std::string>, uint32_t*, std::vector<uint64_t> > GetProperty(
rocprofiler_filter_property_kind_t kind);
filter_property_variant_t GetProperty(rocprofiler_filter_property_kind_t kind);
size_t GetPropertiesCount(rocprofiler_filter_property_kind_t kind);
rocprofiler_spm_parameter_t* GetSpmParameterData();
@@ -77,7 +81,7 @@ class Filter {
std::vector<rocprofiler_att_parameter_t> att_parameters_; // ATT Parameters
rocprofiler_counters_sampler_parameters_t
counters_sampler_parameters_; // sampled counters parameters
std::vector<uint64_t> dispatch_id_filter_;
std::vector<std::pair<uint64_t,uint64_t>> dispatch_id_filter_;
bool has_sync_callback_{false};
rocprofiler_sync_callback_t callback_;