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

Removing ROCMTools naming from ROCProfiler project

Browse Source 
Change-Id: Ibd6211e031157f68a1bc3dd8dd77385e695cb700


[ROCm/rocprofiler commit: 4bcef6b8de]
This commit is contained in:
Ammar ELWazir
2023-05-30 19:21:14 +00:00
committed by Ammar ELWazir
parent 3f6c4eed35
commit af5ad38656
85 changed files with 749 additions and 743 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler/include/rocprofiler/v2/rocprofiler.h
+1 -1
View File
@@ -1982,7 +1982,7 @@ typedef struct {
typedef struct {
/**
* ROCMtool General Record base header to identify the id and kind of every
* ROCProfiler General Record base header to identify the id and kind of every
* record
*/
rocprofiler_record_header_t header;
projects/rocprofiler/plugin/att/att.cpp
+1 -1
View File
@@ -79,7 +79,7 @@ class att_plugin_t {
CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME,
att_tracer_record->kernel_id, &kernel_name_c));
std::string name_demangled = rocmtools::truncate_name(rocmtools::cxx_demangle(kernel_name_c));
std::string name_demangled = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
if (name_demangled.size() > ATT_FILENAME_MAXBYTES) { // Limit filename size
name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES);
projects/rocprofiler/plugin/ctf/plugin.cpp
+4 -4
View File
@@ -138,7 +138,7 @@ class RocTxEventRecord final : public TracerEventRecord<barectf_roctx_ctx> {
: TracerEventRecord<barectf_roctx_ctx>{record, GetRecordBeginClockVal(record)},
id_{record.operation_id.id},
msg_{
rocmtools::cxx_demangle(reinterpret_cast<const char*>(record.api_data_handle.handle))} {
rocprofiler::cxx_demangle(reinterpret_cast<const char*>(record.api_data_handle.handle))} {
}
void Write(barectf_roctx_ctx& barectf_ctx) const override {
@@ -377,7 +377,7 @@ class HipOpEventRecordBegin final : public ApiOpEventRecord {
if (std::strlen(str) > 1) {
// Return demangled version.
return rocmtools::cxx_demangle(str);
return rocprofiler::cxx_demangle(str);
}
}
}
@@ -470,7 +470,7 @@ class ProfilerEventRecord : public BarectfEventRecord<barectf_profiler_ctx> {
}
// Return truncated and demangled version.
return rocmtools::truncate_name(rocmtools::cxx_demangle(kernel_name));
return rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name));
}
// Queries and returns the counter infos of the record `record` and
@@ -668,7 +668,7 @@ void Plugin::HandleTracerRecord(const rocprofiler_record_tracer_t& record,
hip_api_data_t hip_api_data =
*reinterpret_cast<const hip_api_data_t*>(record.api_data_handle.handle);
if (record.name != nullptr)
kernel_name = rocmtools::cxx_demangle(std::string(record.name));
kernel_name = rocprofiler::cxx_demangle(std::string(record.name));
else
kernel_name = "";
hip_api_tracer_.AddEventRecord(
projects/rocprofiler/plugin/file/file.cpp
+5 -5
View File
@@ -86,7 +86,7 @@ class file_plugin_t {
fs::path output_prefix(output_dir);
if (!fs::is_directory(fs::status(output_prefix))) {
if (!stream_.fail()) rocmtools::warning("Cannot open output directory '%s'", output_dir);
if (!stream_.fail()) rocprofiler::warning("Cannot open output directory '%s'", output_dir);
stream_.setstate(std::ios_base::failbit);
return;
}
@@ -150,7 +150,7 @@ class file_plugin_t {
&hsa_handles);
assert(status == HSA_STATUS_SUCCESS && "failed to iterate HSA agents");
if (hsa_handles.fail()) {
rocmtools::warning("Cannot write to '%s'", hsa_handles.name().c_str());
rocprofiler::warning("Cannot write to '%s'", hsa_handles.name().c_str());
return;
}
@@ -162,7 +162,7 @@ class file_plugin_t {
begin_ts << std::dec << app_begin_timestamp.value << std::endl;
if (begin_ts.fail()) {
rocmtools::warning("Cannot write to '%s'", begin_ts.name().c_str());
rocprofiler::warning("Cannot write to '%s'", begin_ts.name().c_str());
return;
}
@@ -207,7 +207,7 @@ class file_plugin_t {
uint64_t roctx_id;
if (tracer_record.name) {
if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API)
kernel_name = rocmtools::cxx_demangle(tracer_record.name);
kernel_name = rocprofiler::cxx_demangle(tracer_record.name);
if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) roctx_message = tracer_record.name;
}
size_t function_name_size = 0;
@@ -315,7 +315,7 @@ class file_plugin_t {
<< std::to_string(profiler_record->kernel_properties.signal_handle);
std::string kernel_name = "";
if (name_length > 1) {
kernel_name = rocmtools::truncate_name(rocmtools::cxx_demangle(kernel_name_c));
kernel_name = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
}
*output_file << "), " << std::string("obj(")
<< std::to_string(profiler_record->kernel_id.handle) << "), "
projects/rocprofiler/plugin/perfetto/perfetto.cpp
+20 -20
View File
@@ -83,7 +83,7 @@ std::string get_kernel_name(rocprofiler_record_profiler_t& profiler_record) {
CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME,
profiler_record.kernel_id, &kernel_name_c));
if (kernel_name_c && strlen(kernel_name_c) > 1)
kernel_name = rocmtools::cxx_demangle(strdup(kernel_name_c));
kernel_name = rocprofiler::cxx_demangle(strdup(kernel_name_c));
}
#pragma GCC diagnostic pop
return kernel_name;
@@ -106,7 +106,7 @@ class perfetto_plugin_t {
output_prefix_ = output_dir;
if (!fs::is_directory(fs::status(output_prefix_))) {
if (!stream_.fail()) rocmtools::warning("Cannot open output directory '%s'", output_dir);
if (!stream_.fail()) rocprofiler::warning("Cannot open output directory '%s'", output_dir);
stream_.setstate(std::ios_base::failbit);
return;
}
@@ -145,7 +145,7 @@ class perfetto_plugin_t {
output_prefix_.append(output_file_name + std::to_string(GetPid()) + "_output.pftrace");
file_descriptor_ = open(output_prefix_.string().c_str(), O_RDWR | O_CREAT | O_TRUNC, 0600);
if (file_descriptor_ == -1) rocmtools::warning("Can't open output file\n");
if (file_descriptor_ == -1) rocprofiler::warning("Can't open output file\n");
tracing_session_ = perfetto::Tracing::NewTrace();
tracing_session_->Setup(trace_cfg, file_descriptor_);
@@ -166,7 +166,7 @@ class perfetto_plugin_t {
}
}
std::string thread_track_str =
rocmtools::string_printf("Node: %s Process ID: %lu Thread ID:", hostname_, GetPid());
rocprofiler::string_printf("Node: %s Process ID: %lu Thread ID:", hostname_, GetPid());
process_track_desc.mutable_process()->set_process_name(thread_track_str);
perfetto::TrackEvent::SetTrackDescriptor(perfetto::ProcessTrack::Current(),
process_track_desc);
@@ -215,7 +215,7 @@ class perfetto_plugin_t {
rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(writing_lock);
// ToDO: rename this variable?
if (!tracing_session_) rocmtools::warning("Tracing session is deleted!\n");
if (!tracing_session_) rocprofiler::warning("Tracing session is deleted!\n");
int device_id = profiler_record.gpu_id.handle;
std::unordered_map<int, perfetto::Track>::iterator device_track_it;
@@ -230,7 +230,7 @@ class perfetto_plugin_t {
.first;
auto gpu_desc = device_track_it->second.Serialize();
gpu_desc.mutable_process()->set_pid(device_id);
std::string gpu_str = rocmtools::string_printf("Node: %s Device:", hostname_);
std::string gpu_str = rocprofiler::string_printf("Node: %s Device:", hostname_);
gpu_desc.mutable_process()->set_process_name(gpu_str);
perfetto::TrackEvent::SetTrackDescriptor(device_track_it->second, gpu_desc);
track_ids_used_.emplace_back(device_id + 1 + machine_id);
@@ -255,7 +255,7 @@ class perfetto_plugin_t {
auto queue_desc = queue_track_it->second.Serialize();
std::string queue_str =
rocmtools::string_printf("Process ID: %lu Queue %ld", GetPid(), gpu_queue_id.second);
rocprofiler::string_printf("Process ID: %lu Queue %ld", GetPid(), gpu_queue_id.second);
queue_desc.set_name(queue_str);
perfetto::TrackEvent::SetTrackDescriptor(queue_track_it->second, queue_desc);
}
@@ -268,7 +268,7 @@ class perfetto_plugin_t {
static const uint32_t lds_block_size = 128 * 4;
std::string full_kernel_name = get_kernel_name(profiler_record);
// std::string truncated_kernel_name = rocmtools::truncate_name(full_kernel_name);
// std::string truncated_kernel_name = rocprofiler::truncate_name(full_kernel_name);
// perfetto::StaticString kernel_name(truncated_kernel_name.c_str());
TRACE_EVENT_BEGIN("KERNELS", perfetto::StaticString(full_kernel_name.c_str()), queue_track,
profiler_record.timestamps.begin.value, "Full Kernel Name",
@@ -343,7 +343,7 @@ class perfetto_plugin_t {
int FlushTracerRecord(rocprofiler_record_tracer_t tracer_record,
rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(writing_lock);
if (!tracing_session_) rocmtools::warning("Tracing session is deleted!\n");
if (!tracing_session_) rocprofiler::warning("Tracing session is deleted!\n");
std::string kernel_name;
const char* function_name;
char* activity_name;
@@ -368,7 +368,7 @@ class perfetto_plugin_t {
.first;
auto gpu_desc = device_track_it->second.Serialize();
gpu_desc.mutable_process()->set_pid(device_id);
std::string gpu_str = rocmtools::string_printf("Node: %s Device:", hostname_);
std::string gpu_str = rocprofiler::string_printf("Node: %s Device:", hostname_);
gpu_desc.mutable_process()->set_process_name(gpu_str);
perfetto::TrackEvent::SetTrackDescriptor(device_track_it->second, gpu_desc);
track_ids_used_.emplace_back(1 + machine_id + device_id);
@@ -390,7 +390,7 @@ class perfetto_plugin_t {
thread_tracks_.emplace(thread_id, perfetto::ProcessTrack::Global(track_id)).first;
auto thread_track_desc = thread_track_it->second.Serialize();
std::string thread_track_str =
rocmtools::string_printf("Node: %s Process ID: %lu Thread ID:", hostname_, GetPid());
rocprofiler::string_printf("Node: %s Process ID: %lu Thread ID:", hostname_, GetPid());
thread_track_desc.mutable_process()->set_pid(thread_id);
thread_track_desc.mutable_process()->set_process_name(thread_track_str);
perfetto::TrackEvent::SetTrackDescriptor(thread_track_it->second, thread_track_desc);
@@ -418,7 +418,7 @@ class perfetto_plugin_t {
roctx_tracks_.emplace(thread_id, perfetto::Track(track_id, thread_track)).first;
auto roctx_track_desc = roctx_track_it->second.Serialize();
std::string roctx_track_str = rocmtools::string_printf("ROCTX Markers");
std::string roctx_track_str = rocprofiler::string_printf("ROCTX Markers");
roctx_track_desc.set_name(roctx_track_str);
perfetto::TrackEvent::SetTrackDescriptor(roctx_track_it->second, roctx_track_desc);
}
@@ -457,7 +457,7 @@ class perfetto_plugin_t {
hsa_track_it =
hsa_tracks_.emplace(thread_id, perfetto::Track(track_id, thread_track)).first;
auto hsa_track_desc = hsa_track_it->second.Serialize();
std::string hsa_track_str = rocmtools::string_printf("HSA API");
std::string hsa_track_str = rocprofiler::string_printf("HSA API");
hsa_track_desc.set_name(hsa_track_str);
perfetto::TrackEvent::SetTrackDescriptor(hsa_track_it->second, hsa_track_desc);
}
@@ -509,7 +509,7 @@ class perfetto_plugin_t {
hip_tracks_.emplace(thread_id, perfetto::Track(track_id, thread_track)).first;
auto hip_track_desc = hip_track_it->second.Serialize();
std::string hip_track_str = rocmtools::string_printf("HIP API");
std::string hip_track_str = rocprofiler::string_printf("HIP API");
hip_track_desc.set_name(hip_track_str);
perfetto::TrackEvent::SetTrackDescriptor(hip_track_it->second, hip_track_desc);
}
@@ -577,7 +577,7 @@ class perfetto_plugin_t {
auto stream_desc = stream_track_it->second.Serialize();
std::string stream_str =
rocmtools::string_printf("Process ID: %lu Stream %d", GetPid(), stream_id);
rocprofiler::string_printf("Process ID: %lu Stream %d", GetPid(), stream_id);
stream_desc.set_name(stream_str);
perfetto::TrackEvent::SetTrackDescriptor(stream_track_it->second, stream_desc);
track_ids_used_.emplace_back(1 + machine_id + tracer_record.agent_id.handle);
@@ -587,11 +587,11 @@ class perfetto_plugin_t {
rocprofiler_timestamp_t timestamp;
rocprofiler_get_timestamp(&timestamp);
if (tracer_record.api_data_handle.handle && tracer_record.api_data_handle.size > 1) {
kernel_name = rocmtools::cxx_demangle(
kernel_name = rocprofiler::cxx_demangle(
strdup(reinterpret_cast<const char*>(tracer_record.api_data_handle.handle)));
TRACE_EVENT_BEGIN(
"HIP_OPS",
perfetto::StaticString(strdup(rocmtools::truncate_name(kernel_name).c_str())),
perfetto::StaticString(strdup(rocprofiler::truncate_name(kernel_name).c_str())),
stream_track, tracer_record.timestamps.begin.value, "Agent ID",
tracer_record.agent_id.handle, "Process ID", GetPid(), "Kernel Name", kernel_name,
perfetto::Flow::ProcessScoped(tracer_record.correlation_id.value));
@@ -642,7 +642,7 @@ class perfetto_plugin_t {
.first;
auto queue_desc = queue_track_it->second.Serialize();
std::string queue_str = rocmtools::string_printf("Process ID: %lu Queue %ld", GetPid(),
std::string queue_str = rocprofiler::string_printf("Process ID: %lu Queue %ld", GetPid(),
gpu_queue_id.second);
queue_desc.set_name(queue_str);
perfetto::TrackEvent::SetTrackDescriptor(queue_track_it->second, queue_desc);
@@ -670,7 +670,7 @@ class perfetto_plugin_t {
break;
}
default: {
rocmtools::warning("ignored record for domain %d", tracer_record.domain);
rocprofiler::warning("ignored record for domain %d", tracer_record.domain);
break;
}
}
@@ -680,7 +680,7 @@ class perfetto_plugin_t {
int WriteBufferRecords(const rocprofiler_record_header_t* begin,
const rocprofiler_record_header_t* end,
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) {
if (!tracing_session_) rocmtools::warning("Tracing session is deleted!\n");
if (!tracing_session_) rocprofiler::warning("Tracing session is deleted!\n");
while (begin < end) {
if (!begin) return 0;
switch (begin->kind) {
projects/rocprofiler/plugin/utils.h
+1 -1
View File
@@ -38,7 +38,7 @@
// Macro to check ROCProfiler calls status
#define CHECK_ROCPROFILER(call) \
do { \
if ((call) != ROCPROFILER_STATUS_SUCCESS) rocmtools::fatal("Error: ROCProfiler API Call Error!"); \
if ((call) != ROCPROFILER_STATUS_SUCCESS) rocprofiler::fatal("Error: ROCProfiler API Call Error!"); \
} while (false)
namespace {
projects/rocprofiler/samples/common/common.h
+5 -5
View File
@@ -44,7 +44,7 @@
#define CHECK_ROCPROFILER(call) \
do { \
if ((call) != ROCPROFILER_STATUS_SUCCESS) \
rocmtools::fatal("Error: ROCProfiler API Call Error!"); \
rocprofiler::fatal("Error: ROCProfiler API Call Error!"); \
} while (false)
// Device (Kernel) functions, it must be void
@@ -114,7 +114,7 @@ void FlushTracerRecord(rocprofiler_record_tracer_t tracer_record,
if ((tracer_record.operation_id.id == 0 && tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)) {
if (tracer_record.api_data_handle.handle &&
strlen(reinterpret_cast<const char*>(tracer_record.api_data_handle.handle)) > 1)
kernel_name = rocmtools::cxx_demangle(
kernel_name = rocprofiler::cxx_demangle(
reinterpret_cast<const char*>(tracer_record.api_data_handle.handle));
}
if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_API) {
@@ -151,7 +151,7 @@ void FlushTracerRecord(rocprofiler_record_tracer_t tracer_record,
CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info(
session_id, ROCPROFILER_HIP_KERNEL_NAME, tracer_record.api_data_handle,
tracer_record.operation_id, &kernel_name_str));
if (kernel_name_str) kernel_name = rocmtools::cxx_demangle(std::string(kernel_name_str));
if (kernel_name_str) kernel_name = rocprofiler::cxx_demangle(std::string(kernel_name_str));
}
}
if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) {
@@ -166,7 +166,7 @@ void FlushTracerRecord(rocprofiler_record_tracer_t tracer_record,
session_id, ROCPROFILER_ROCTX_MESSAGE, tracer_record.api_data_handle,
tracer_record.operation_id, &roctx_message_str));
if (roctx_message_str)
roctx_message = rocmtools::cxx_demangle(std::string(strdup(roctx_message_str)));
roctx_message = rocprofiler::cxx_demangle(std::string(strdup(roctx_message_str)));
}
size_t roctx_id_size = 0;
CHECK_ROCPROFILER(rocprofiler_query_roctx_tracer_api_data_info_size(
@@ -250,7 +250,7 @@ void FlushProfilerRecord(const rocprofiler_record_profiler_t* profiler_record,
<< std::to_string(profiler_record->kernel_properties.wave_size) << "), "
<< std::string("sig(")
<< std::to_string(profiler_record->kernel_properties.signal_handle);
std::string kernel_name = rocmtools::cxx_demangle(kernel_name_c);
std::string kernel_name = rocprofiler::cxx_demangle(kernel_name_c);
output_file << "), " << std::string("obj(") << std::to_string(profiler_record->kernel_id.handle)
<< "), " << std::string("kernel-name(\"") << kernel_name << "\")"
<< std::string(", time(") << std::to_string(profiler_record->timestamps.begin.value)
projects/rocprofiler/samples/common/helper.cpp
+2 -2
View File
@@ -37,7 +37,7 @@
} \
} while (false)
namespace rocmtools {
namespace rocprofiler {
std::string string_vprintf(const char* format, va_list va) {
va_list copy;
@@ -178,4 +178,4 @@ std::string left_trim(const std::string& s) {
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/samples/common/helper.h
+2 -2
View File
@@ -30,7 +30,7 @@
#include <cxxabi.h>
// #include "exception.h"
namespace rocmtools {
namespace rocprofiler {
std::string string_vprintf(const char* format, va_list va);
@@ -69,4 +69,4 @@ bool has_counter_format(std::string const& str);
// trims the begining of the line for spaces
std::string left_trim(const std::string& s);
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/api/CMakeLists.txt
+1 -1
View File
@@ -153,7 +153,7 @@ file(GLOB ROCPROFILER_PROFILER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/
file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp)
file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp)
file(GLOB ROCPROFILER_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp)
file(GLOB ROCPROFILER_CLASS_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocmtool.cpp)
file(GLOB ROCPROFILER_CLASS_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp)
file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp)
projects/rocprofiler/src/api/rocmtool.cpp → projects/rocprofiler/src/api/rocprofiler_singleton.cpp
+39 -38
View File
@@ -18,7 +18,7 @@
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#include "rocmtool.h"
#include "rocprofiler_singleton.h"
#include <sys/syscall.h>
#include <unistd.h>
@@ -40,23 +40,23 @@ extern std::mutex sessions_pending_signal_lock;
static inline uint32_t GetTid() { return syscall(__NR_gettid); }
namespace rocmtools {
namespace rocprofiler {
rocmtool* rocmtool_obj;
ROCProfiler_Singleton* rocprofiler_singleton;
// Constructor of rocmtool
// Constructor of ROCProfiler
// Takes the buffer size, a buffer callback function and a buffer flush
// interval to allocate a buffer pool using GenericStorage Also takes the
// replay mode (application replay/kernel replay/user replay) to set the replay
// mode for the rocmtool class object
rocmtool::rocmtool() : current_session_id_(rocprofiler_session_id_t{0}) {}
// mode for the rocprofiler class object
ROCProfiler_Singleton::ROCProfiler_Singleton() : current_session_id_(rocprofiler_session_id_t{0}) {}
// Destructor of rocmtool
// Destructor of rocprofiler_singleton
// deletes the buffer pool
// Iterates over its session map and resets each session in its internal
// session map and clears them from the map. Pops labels from the range stack
// and deletes the stack.
rocmtool::~rocmtool() {
ROCProfiler_Singleton::~ROCProfiler_Singleton() {
// {
// std::lock_guard<std::mutex> lock(session_map_lock_);
// if (!sessions_.empty()) {
@@ -69,31 +69,31 @@ rocmtool::~rocmtool() {
Counter::ClearBasicCounters();
}
bool rocmtool::FindAgent(rocprofiler_agent_id_t agent_id) { return true; }
size_t rocmtool::GetAgentInfoSize(rocprofiler_agent_info_kind_t kind, rocprofiler_agent_id_t agent_id) {
bool ROCProfiler_Singleton::FindAgent(rocprofiler_agent_id_t agent_id) { return true; }
size_t ROCProfiler_Singleton::GetAgentInfoSize(rocprofiler_agent_info_kind_t kind, rocprofiler_agent_id_t agent_id) {
return 0;
}
const char* rocmtool::GetAgentInfo(rocprofiler_agent_info_kind_t kind,
const char* ROCProfiler_Singleton::GetAgentInfo(rocprofiler_agent_info_kind_t kind,
rocprofiler_agent_id_t agent_id) {
return "";
}
// TODO(aelwazir): Implement Queue Query
bool rocmtool::FindQueue(rocprofiler_queue_id_t queue_id) { return true; }
size_t rocmtool::GetQueueInfoSize(rocprofiler_queue_info_kind_t kind, rocprofiler_queue_id_t queue_id) {
bool ROCProfiler_Singleton::FindQueue(rocprofiler_queue_id_t queue_id) { return true; }
size_t ROCProfiler_Singleton::GetQueueInfoSize(rocprofiler_queue_info_kind_t kind, rocprofiler_queue_id_t queue_id) {
return 0;
}
const char* rocmtool::GetQueueInfo(rocprofiler_queue_info_kind_t kind,
const char* ROCProfiler_Singleton::GetQueueInfo(rocprofiler_queue_info_kind_t kind,
rocprofiler_queue_id_t queue_id) {
return "";
}
bool rocmtool::FindSession(rocprofiler_session_id_t session_id) {
bool ROCProfiler_Singleton::FindSession(rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(session_map_lock_);
return sessions_.find(session_id.handle) != sessions_.end();
}
rocprofiler_session_id_t rocmtool::CreateSession(rocprofiler_replay_mode_t replay_mode) {
rocprofiler_session_id_t ROCProfiler_Singleton::CreateSession(rocprofiler_replay_mode_t replay_mode) {
rocprofiler_session_id_t session_id = rocprofiler_session_id_t{GenerateUniqueSessionId()};
{
std::lock_guard<std::mutex> lock(session_map_lock_);
@@ -102,7 +102,7 @@ rocprofiler_session_id_t rocmtool::CreateSession(rocprofiler_replay_mode_t repla
return session_id;
}
void rocmtool::DestroySession(rocprofiler_session_id_t session_id) {
void ROCProfiler_Singleton::DestroySession(rocprofiler_session_id_t session_id) {
while (GetCurrentActiveInterruptSignalsCount() != 0) {
}
@@ -125,12 +125,12 @@ void rocmtool::DestroySession(rocprofiler_session_id_t session_id) {
}
}
bool rocmtool::FindDeviceProfilingSession(rocprofiler_session_id_t session_id) {
bool ROCProfiler_Singleton::FindDeviceProfilingSession(rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(device_profiling_session_map_lock_);
return dev_profiling_sessions_.find(session_id.handle) != dev_profiling_sessions_.end();
}
rocprofiler_session_id_t rocmtool::CreateDeviceProfilingSession(std::vector<std::string> counters,
rocprofiler_session_id_t ROCProfiler_Singleton::CreateDeviceProfilingSession(std::vector<std::string> counters,
int cpu_agent_index,
int gpu_agent_index) {
rocprofiler_session_id_t session_id;
@@ -149,7 +149,7 @@ rocprofiler_session_id_t rocmtool::CreateDeviceProfilingSession(std::vector<std:
return session_id;
}
void rocmtool::DestroyDeviceProfilingSession(rocprofiler_session_id_t session_id) {
void ROCProfiler_Singleton::DestroyDeviceProfilingSession(rocprofiler_session_id_t session_id) {
{
std::lock_guard<std::mutex> lock(device_profiling_session_map_lock_);
ASSERTM(dev_profiling_sessions_.find(session_id.handle) != dev_profiling_sessions_.end(),
@@ -159,15 +159,15 @@ void rocmtool::DestroyDeviceProfilingSession(rocprofiler_session_id_t session_id
}
}
DeviceProfileSession* rocmtool::GetDeviceProfilingSession(rocprofiler_session_id_t session_id) {
DeviceProfileSession* ROCProfiler_Singleton::GetDeviceProfilingSession(rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(device_profiling_session_map_lock_);
assert(dev_profiling_sessions_.find(session_id.handle) != dev_profiling_sessions_.end() &&
"Error: Can't find the session!");
return dev_profiling_sessions_.at(session_id.handle);
}
bool rocmtool::HasActiveSession() { return GetCurrentSessionId().handle > 0; }
bool rocmtool::IsActiveSession(rocprofiler_session_id_t session_id) {
bool ROCProfiler_Singleton::HasActiveSession() { return GetCurrentSessionId().handle > 0; }
bool ROCProfiler_Singleton::IsActiveSession(rocprofiler_session_id_t session_id) {
return (GetCurrentSessionId().handle == session_id.handle);
}
@@ -176,26 +176,26 @@ bool rocmtool::IsActiveSession(rocprofiler_session_id_t session_id) {
// If a given session id doesn't exist, it throws an assertion.
// If a session object exists for the given session id, the session object is
// returned.
Session* rocmtool::GetSession(rocprofiler_session_id_t session_id) {
Session* ROCProfiler_Singleton::GetSession(rocprofiler_session_id_t session_id) {
std::lock_guard<std::mutex> lock(session_map_lock_);
assert(sessions_.find(session_id.handle) != sessions_.end() && "Error: Can't find the session!");
return sessions_.at(session_id.handle);
}
// Get Current Session ID
rocprofiler_session_id_t rocmtool::GetCurrentSessionId() { return current_session_id_; }
rocprofiler_session_id_t ROCProfiler_Singleton::GetCurrentSessionId() { return current_session_id_; }
void rocmtool::SetCurrentActiveSession(rocprofiler_session_id_t session_id) {
void ROCProfiler_Singleton::SetCurrentActiveSession(rocprofiler_session_id_t session_id) {
current_session_id_ = session_id;
}
uint64_t rocmtool::GetUniqueRecordId() { return records_counter_.fetch_add(1); }
uint64_t ROCProfiler_Singleton::GetUniqueRecordId() { return records_counter_.fetch_add(1); }
uint64_t rocmtool::GetUniqueKernelDispatchId() {
uint64_t ROCProfiler_Singleton::GetUniqueKernelDispatchId() {
return kernel_dispatch_counter_.fetch_add(1, std::memory_order_relaxed);
}
size_t rocmtool::GetKernelInfoSize(rocprofiler_kernel_info_kind_t kind,
size_t ROCProfiler_Singleton::GetKernelInfoSize(rocprofiler_kernel_info_kind_t kind,
rocprofiler_kernel_id_t kernel_id) {
switch (kind) {
case ROCPROFILER_KERNEL_NAME:
@@ -205,7 +205,7 @@ size_t rocmtool::GetKernelInfoSize(rocprofiler_kernel_info_kind_t kind,
return 0;
}
}
const char* rocmtool::GetKernelInfo(rocprofiler_kernel_info_kind_t kind,
const char* ROCProfiler_Singleton::GetKernelInfo(rocprofiler_kernel_info_kind_t kind,
rocprofiler_kernel_id_t kernel_id) {
switch (kind) {
case ROCPROFILER_KERNEL_NAME:
@@ -217,19 +217,20 @@ const char* rocmtool::GetKernelInfo(rocprofiler_kernel_info_kind_t kind,
}
// TODO(aelwazir): To be implemented
bool rocmtool::CheckFilterData(rocprofiler_filter_kind_t filter_kind,
bool ROCProfiler_Singleton::CheckFilterData(rocprofiler_filter_kind_t filter_kind,
rocprofiler_filter_data_t filter_data) {
return true;
}
// End of ROCMTool Class
// End of ROCProfiler_Singleton Class
rocmtool* GetROCMToolObj() { return rocmtool_obj; }
ROCProfiler_Singleton* GetROCProfilerSingleton() { return rocprofiler_singleton; }
void InitROCMToolObj() { rocmtool_obj = new rocmtool; }
void ResetROCMToolObj() {
delete rocmtool_obj;
// if (rocmtool_obj) rocmtool_obj.reset();
void InitROCProfilerSingleton() { rocprofiler_singleton = new ROCProfiler_Singleton; }
void ResetROCProfilerSingleton() {
delete rocprofiler_singleton;
// TODO(aelwazir): We need to use std::optional or std::unique_ptr
// if (rocprofiler_singleton) rocprofiler_singleton.reset();
}
rocprofiler_timestamp_t GetCurrentTimestamp() { return hsa_support::GetCurrentTimestampNS(); }
@@ -239,4 +240,4 @@ rocprofiler_status_t IterateCounters(rocprofiler_counters_info_callback_t counte
return ROCPROFILER_STATUS_ERROR;
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/api/rocmtool.h → projects/rocprofiler/src/api/rocprofiler_singleton.h
+12 -12
View File
@@ -18,8 +18,8 @@
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#ifndef SRC_TOOLS_ROCMTOOL_H_
#define SRC_TOOLS_ROCMTOOL_H_
#ifndef SRC_TOOLS_ROCPROFILER_SINGLETON_H_
#define SRC_TOOLS_ROCPROFILER_SINGLETON_H_
#include <hsa/hsa_ven_amd_aqlprofile.h>
@@ -42,12 +42,12 @@
#include "src/core/session/session.h"
#include "src/core/session/device_profiling.h"
namespace rocmtools {
namespace rocprofiler {
class rocmtool {
class ROCProfiler_Singleton {
public:
rocmtool();
~rocmtool();
ROCProfiler_Singleton();
~ROCProfiler_Singleton();
bool FindAgent(rocprofiler_agent_id_t agent_id);
size_t GetAgentInfoSize(rocprofiler_agent_info_kind_t kind, rocprofiler_agent_id_t agent_id);
@@ -98,7 +98,7 @@ class rocmtool {
* identical if the same kernel is dispatched twice. Currently, this
* identifier is written to the `reserved2` field of the dispatch packet when
* its launch is intercepted, but this could change: a future version of
* ROCmtools may instead attempt to identify a kernel by a key with high
* ROCProfiler may instead attempt to identify a kernel by a key with high
* _probability_ of uniqueness: for example, a combination of the kernel's
* name, the queue ID to which it was dispatched, and the offset of the queue
* write pointer is likely sufficient to associate PC samples with a running
@@ -109,14 +109,14 @@ class rocmtool {
std::atomic<uint64_t> kernel_dispatch_counter_{1};
};
void InitROCMToolObj();
void ResetROCMToolObj();
rocmtool* GetROCMToolObj();
void InitROCProfilerSingleton();
void ResetROCProfilerSingleton();
ROCProfiler_Singleton* GetROCProfilerSingleton();
rocprofiler_timestamp_t GetCurrentTimestamp();
rocprofiler_status_t IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_TOOLS_ROCMTOOL_H_
#endif // SRC_TOOLS_ROCPROFILER_SINGLETON_H_
projects/rocprofiler/src/api/rocmtools.cpp → projects/rocprofiler/src/api/rocprofilerv2.cpp
+189 -189
View File
@@ -1,7 +1,7 @@
#include <atomic>
#include "src/core/hsa/hsa_support.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/utils/helper.h"
#include "rocprofiler.h"
@@ -12,7 +12,7 @@
try {
#define API_METHOD_SUFFIX \
} \
catch (rocmtools::Exception & e) { \
catch (rocprofiler::Exception & e) { \
std::cout << __FUNCTION__ << "(), " << e.what(); \
} \
return err;
@@ -20,7 +20,7 @@
#define API_INIT_CHECKER \
API_METHOD_PREFIX \
if (!api_started.load(std::memory_order_relaxed)) \
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_NOT_INITIALIZED);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_NOT_INITIALIZED);
std::atomic<bool> api_started{false};
@@ -32,9 +32,9 @@ ROCPROFILER_API uint32_t rocprofiler_version_minor() { return ROCPROFILER_VERSIO
ROCPROFILER_API const char* rocprofiler_error_str(rocprofiler_status_t status) {
switch (status) {
case ROCPROFILER_STATUS_ERROR_ALREADY_INITIALIZED:
return "ROCMTool is already initialized\n";
return "ROCProfiler is already initialized\n";
case ROCPROFILER_STATUS_ERROR_NOT_INITIALIZED:
return "ROCMTool is not initialized or already destroyed\n";
return "ROCProfiler is not initialized or already destroyed\n";
case ROCPROFILER_STATUS_ERROR_SESSION_MISSING_BUFFER:
return "Missing Buffer for a session\n";
case ROCPROFILER_STATUS_ERROR_TIMESTAMP_NOT_APPLICABLE:
@@ -111,8 +111,8 @@ ROCPROFILER_API const char* rocprofiler_error_str(rocprofiler_status_t status) {
ROCPROFILER_API rocprofiler_status_t rocprofiler_initialize() {
API_METHOD_PREFIX
if (api_started.load(std::memory_order_relaxed))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_ALREADY_INITIALIZED);
rocmtools::InitROCMToolObj();
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_ALREADY_INITIALIZED);
rocprofiler::InitROCProfilerSingleton();
api_started.exchange(true, std::memory_order_release);
API_METHOD_SUFFIX
}
@@ -120,23 +120,23 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_initialize() {
// Finalize the API
ROCPROFILER_API rocprofiler_status_t rocprofiler_finalize() {
API_INIT_CHECKER
rocmtools::ResetROCMToolObj();
rocprofiler::ResetROCProfilerSingleton();
api_started.exchange(false, std::memory_order_release);
API_METHOD_SUFFIX
}
ROCPROFILER_API rocprofiler_status_t rocprofiler_get_timestamp(rocprofiler_timestamp_t* timestamp) {
API_INIT_CHECKER
*timestamp = rocmtools::GetCurrentTimestamp();
*timestamp = rocprofiler::GetCurrentTimestamp();
if (timestamp->value <= 0)
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TIMESTAMP_NOT_APPLICABLE);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TIMESTAMP_NOT_APPLICABLE);
API_METHOD_SUFFIX
}
ROCPROFILER_API rocprofiler_status_t
rocprofiler_iterate_counters(rocprofiler_counters_info_callback_t counters_info_callback) {
API_INIT_CHECKER
return rocmtools::IterateCounters(counters_info_callback);
return rocprofiler::IterateCounters(counters_info_callback);
API_METHOD_SUFFIX
}
@@ -144,10 +144,10 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_agent_info_size(rocprofil
rocprofiler_agent_id_t agent_id,
size_t* data_size) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindAgent(agent_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND);
*data_size = rocmtools::GetROCMToolObj()->GetAgentInfoSize(kind, agent_id);
if (*data_size <= 0) throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_AGENT_INFORMATION_MISSING);
if (!rocprofiler::GetROCProfilerSingleton()->FindAgent(agent_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND);
*data_size = rocprofiler::GetROCProfilerSingleton()->GetAgentInfoSize(kind, agent_id);
if (*data_size <= 0) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_AGENT_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -155,10 +155,10 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_agent_info(rocprofiler_ag
rocprofiler_agent_id_t agent_id,
const char** data) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindAgent(agent_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND);
if (!(*data = rocmtools::GetROCMToolObj()->GetAgentInfo(kind, agent_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_AGENT_INFORMATION_MISSING);
if (!rocprofiler::GetROCProfilerSingleton()->FindAgent(agent_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND);
if (!(*data = rocprofiler::GetROCProfilerSingleton()->GetAgentInfo(kind, agent_id)))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_AGENT_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -166,10 +166,10 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_queue_info_size(rocprofil
rocprofiler_queue_id_t queue_id,
size_t* data_size) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindQueue(queue_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_NOT_FOUND);
*data_size = rocmtools::GetROCMToolObj()->GetQueueInfoSize(kind, queue_id);
if (*data_size <= 0) throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_INFORMATION_MISSING);
if (!rocprofiler::GetROCProfilerSingleton()->FindQueue(queue_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_NOT_FOUND);
*data_size = rocprofiler::GetROCProfilerSingleton()->GetQueueInfoSize(kind, queue_id);
if (*data_size <= 0) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -177,10 +177,10 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_queue_info(rocprofiler_qu
rocprofiler_queue_id_t queue_id,
const char** data) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindQueue(queue_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_NOT_FOUND);
if (!(*data = rocmtools::GetROCMToolObj()->GetQueueInfo(kind, queue_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_INFORMATION_MISSING);
if (!rocprofiler::GetROCProfilerSingleton()->FindQueue(queue_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_NOT_FOUND);
if (!(*data = rocprofiler::GetROCProfilerSingleton()->GetQueueInfo(kind, queue_id)))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_QUEUE_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -188,11 +188,11 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_kernel_info_size(rocprofi
rocprofiler_kernel_id_t kernel_id,
size_t* data_size) {
API_INIT_CHECKER
// if (!rocmtools::GetROCMToolObj()->FindKernel(kernel_id))
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_NOT_FOUND);
*data_size = rocmtools::GetROCMToolObj()->GetKernelInfoSize(kind, kernel_id);
// if (!rocprofiler::GetROCProfilerSingleton()->FindKernel(kernel_id))
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_NOT_FOUND);
*data_size = rocprofiler::GetROCProfilerSingleton()->GetKernelInfoSize(kind, kernel_id);
if (*data_size <= 0)
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -200,10 +200,10 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_kernel_info(rocprofiler_k
rocprofiler_kernel_id_t kernel_id,
const char** data) {
API_INIT_CHECKER
// if (!rocmtools::GetROCMToolObj()->FindKernel(kernel_id))
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_NOT_FOUND);
if (!(*data = rocmtools::GetROCMToolObj()->GetKernelInfo(kind, kernel_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_INFORMATION_MISSING);
// if (!rocprofiler::GetROCProfilerSingleton()->FindKernel(kernel_id))
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_NOT_FOUND);
if (!(*data = rocprofiler::GetROCProfilerSingleton()->GetKernelInfo(kind, kernel_id)))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_KERNEL_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -211,14 +211,14 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_counter_info_size(
rocprofiler_session_id_t session_id, rocprofiler_counter_info_kind_t kind,
rocprofiler_counter_id_t counter_id, size_t* data_size) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->GetProfiler()->FindCounter(counter_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND);
*data_size = rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetProfiler()->FindCounter(counter_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND);
*data_size = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetProfiler()
->GetCounterInfoSize(kind, counter_id);
if (*data_size <= 0)
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -227,13 +227,13 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_counter_info(rocprofiler_
rocprofiler_counter_id_t counter_id,
const char** data) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->GetProfiler()->FindCounter(counter_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND);
if (!(*data = rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetProfiler()->FindCounter(counter_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND);
if (!(*data = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetProfiler()
->GetCounterInfo(kind, counter_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_COUNTER_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -243,29 +243,29 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_roctx_tracer_api_data_inf
size_t* data_size) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
*data_size = rocmtools::GetROCMToolObj()
*data_size = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetROCTxApiDataInfoSize(kind, api_data_id, operation_id);
// if (*data_size <= 0)
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -275,28 +275,28 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_roctx_tracer_api_data_inf
char** data) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
if (!(*data = rocmtools::GetROCMToolObj()
if (!(*data = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetROCTxApiDataInfo(kind, api_data_id, operation_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -306,31 +306,31 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hsa_tracer_api_data_info_
size_t* data_size) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
*data_size = ((kind == ROCPROFILER_HSA_FUNCTION_NAME)
? rocmtools::tracer::GetApiCallFunctionNameSize(ACTIVITY_DOMAIN_HSA_API, operation_id)
: rocmtools::GetROCMToolObj()
? rocprofiler::tracer::GetApiCallFunctionNameSize(ACTIVITY_DOMAIN_HSA_API, operation_id)
: rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetHSAApiDataInfoSize(kind, api_data_id, operation_id));
if (*data_size <= 0)
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -340,30 +340,30 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hsa_tracer_api_data_info(
char** data) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
if (!(*data = (kind == ROCPROFILER_HSA_FUNCTION_NAME)
? rocmtools::tracer::GetApiCallFunctionName(ACTIVITY_DOMAIN_HSA_API, operation_id)
: rocmtools::GetROCMToolObj()
? rocprofiler::tracer::GetApiCallFunctionName(ACTIVITY_DOMAIN_HSA_API, operation_id)
: rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetHSAApiDataInfo(kind, api_data_id, operation_id)))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -373,31 +373,31 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hip_tracer_api_data_info_
size_t* data_size) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
*data_size = (kind == ROCPROFILER_HIP_FUNCTION_NAME)
? rocmtools::tracer::GetApiCallFunctionNameSize(ACTIVITY_DOMAIN_HIP_API, operation_id)
: rocmtools::GetROCMToolObj()
? rocprofiler::tracer::GetApiCallFunctionNameSize(ACTIVITY_DOMAIN_HIP_API, operation_id)
: rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetHIPApiDataInfoSize(kind, api_data_id, operation_id);
// if (*data_size <= 0)
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
@@ -407,45 +407,45 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hip_tracer_api_data_info(
char** data) {
API_INIT_CHECKER
// TODO(aelwazir): To be implemented
// if (!rocmtools::GetROCMToolObj()
// if (!rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHIPApiData(api_data_id)) {
// if (rocmtools::GetROCMToolObj()
// if (rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindHSAApiData(api_data_id) ||
// rocmtools::GetROCMToolObj()
// rocprofiler::GetROCProfilerSingleton()
// ->GetSession(session_id)
// ->GetTracer()
// ->FindROCTxApiData(api_data_id)) {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN);
// } else {
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND);
// }
// }
// if (!(
*data = (kind == ROCPROFILER_HIP_FUNCTION_NAME)
? rocmtools::tracer::GetApiCallFunctionName(ACTIVITY_DOMAIN_HIP_API, operation_id)
: rocmtools::GetROCMToolObj()
? rocprofiler::tracer::GetApiCallFunctionName(ACTIVITY_DOMAIN_HIP_API, operation_id)
: rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetTracer()
->GetHIPApiDataInfo(kind, api_data_id, operation_id);
// ))
// throw
// rocmtools::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
// rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING);
API_METHOD_SUFFIX
}
ROCPROFILER_API rocprofiler_status_t rocprofiler_flush_data(rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->GetBuffer(buffer_id)->Flush())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_CORRUPTED_SESSION_BUFFER);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetBuffer(buffer_id)->Flush())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_CORRUPTED_SESSION_BUFFER);
API_METHOD_SUFFIX
}
@@ -456,12 +456,12 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_next_record(const rocprofiler_r
rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!Memory::GetNextRecord(record, next))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_RECORD_CORRUPTED);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_RECORD_CORRUPTED);
API_METHOD_SUFFIX
}
@@ -469,7 +469,7 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_next_record(const rocprofiler_r
ROCPROFILER_API rocprofiler_status_t rocprofiler_create_session(rocprofiler_replay_mode_t replay_mode,
rocprofiler_session_id_t* session_id) {
API_INIT_CHECKER
*session_id = rocmtools::GetROCMToolObj()->CreateSession(replay_mode);
*session_id = rocprofiler::GetROCProfilerSingleton()->CreateSession(replay_mode);
API_METHOD_SUFFIX
}
@@ -482,14 +482,14 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_create_filter(rocprofiler_sessi
API_INIT_CHECKER
// TODO(aelwazir): CheckFilterData to be implemented
// int error_code =
// rocmtools::GetROCMToolObj()->CheckFilterData(filter_kind,
// rocprofiler::GetROCProfilerSingleton()->CheckFilterData(filter_kind,
// filter_data);
// if (error_code == -1) throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_DATA_CORRUPTED);
// if (error_code == -1) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_DATA_CORRUPTED);
// if (error_code == 0)
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_FILTER_DATA_MISMATCH);
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
*filter_id = rocmtools::GetROCMToolObj()
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_FILTER_DATA_MISMATCH);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
*filter_id = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->CreateFilter(filter_kind, filter_data, data_count, property);
API_METHOD_SUFFIX
@@ -498,11 +498,11 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_create_filter(rocprofiler_sessi
ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_filter(rocprofiler_session_id_t session_id,
rocprofiler_filter_id_t filter_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindFilter(filter_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
rocmtools::GetROCMToolObj()->GetSession(session_id)->DestroyFilter(filter_id);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindFilter(filter_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->DestroyFilter(filter_id);
API_METHOD_SUFFIX
}
@@ -510,9 +510,9 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_create_buffer(
rocprofiler_session_id_t session_id, rocprofiler_buffer_callback_t buffer_callback,
size_t buffer_size, rocprofiler_buffer_id_t* buffer_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
*buffer_id = rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
*buffer_id = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->CreateBuffer(buffer_callback, buffer_size);
API_METHOD_SUFFIX
@@ -522,11 +522,11 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_set_buffer_properties(
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
rocprofiler_buffer_property_t* buffer_properties, uint32_t buffer_properties_count) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetBuffer(buffer_id)
->SetProperties(buffer_properties, buffer_properties_count);
@@ -536,11 +536,11 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_set_buffer_properties(
ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_buffer(rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
rocmtools::GetROCMToolObj()->GetSession(session_id)->DestroyBuffer(buffer_id);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->DestroyBuffer(buffer_id);
API_METHOD_SUFFIX
}
@@ -548,17 +548,17 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_set_filter_buffer(rocprofiler_s
rocprofiler_filter_id_t filter_id,
rocprofiler_buffer_id_t buffer_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindFilter(filter_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindBuffer(buffer_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindFilter(filter_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->CheckFilterBufferSize(filter_id, buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_SIZE);
rocmtools::GetROCMToolObj()->GetSession(session_id)->GetFilter(filter_id)->SetBufferId(buffer_id);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_SIZE);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetFilter(filter_id)->SetBufferId(buffer_id);
API_METHOD_SUFFIX
}
@@ -566,14 +566,14 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_set_api_trace_sync_callback(
rocprofiler_session_id_t session_id, rocprofiler_filter_id_t filter_id,
rocprofiler_sync_callback_t callback) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->FindFilter(filter_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
if (rocmtools::GetROCMToolObj()->GetSession(session_id)->GetFilter(filter_id)->GetKind() !=
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->FindFilter(filter_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND);
if (rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetFilter(filter_id)->GetKind() !=
ROCPROFILER_API_TRACE)
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_SUPPORTED);
rocmtools::GetROCMToolObj()->GetSession(session_id)->GetFilter(filter_id)->SetCallback(callback);
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_NOT_SUPPORTED);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetFilter(filter_id)->SetCallback(callback);
API_METHOD_SUFFIX
}
@@ -585,29 +585,29 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_create_ready_session(
API_INIT_CHECKER
// TODO(aelwazir): CheckFilterData to be implemented
// int error_code =
// rocmtools::GetROCMToolObj()->CheckFilterData(filter_kind,
// rocprofiler::GetROCProfilerSingleton()->CheckFilterData(filter_kind,
// filter_data);
// if (error_code == -1) throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_FILTER_DATA_CORRUPTED);
// if (error_code == -1) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_FILTER_DATA_CORRUPTED);
// if (error_code == 0)
// throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_FILTER_DATA_MISMATCH);
*session_id = rocmtools::GetROCMToolObj()->CreateSession(replay_mode);
// throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_FILTER_DATA_MISMATCH);
*session_id = rocprofiler::GetROCProfilerSingleton()->CreateSession(replay_mode);
rocprofiler_filter_id_t filter_id =
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession(*session_id)
->CreateFilter(filter_kind, filter_data, data_count, property);
rocprofiler_buffer_id_t buffer_id = rocmtools::GetROCMToolObj()
rocprofiler_buffer_id_t buffer_id = rocprofiler::GetROCProfilerSingleton()
->GetSession(*session_id)
->CreateBuffer(buffer_callback, buffer_size);
if (filter_kind == ROCPROFILER_API_TRACE)
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession(*session_id)
->GetFilter(filter_id)
->SetCallback(callback);
if (!rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()
->GetSession(*session_id)
->CheckFilterBufferSize(filter_id, buffer_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_SIZE);
rocmtools::GetROCMToolObj()
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INCORRECT_SIZE);
rocprofiler::GetROCProfilerSingleton()
->GetSession(*session_id)
->GetFilter(filter_id)
->SetBufferId(buffer_id);
@@ -617,37 +617,37 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_create_ready_session(
// API to destroy a session by id
ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_session(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
rocmtools::GetROCMToolObj()->DestroySession(session_id);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
rocprofiler::GetROCProfilerSingleton()->DestroySession(session_id);
API_METHOD_SUFFIX
}
// API to activate a session by id
ROCPROFILER_API rocprofiler_status_t rocprofiler_start_session(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->HasFilter())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_MISSING_FILTER);
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->HasBuffer())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_MISSING_BUFFER);
if (rocmtools::GetROCMToolObj()->HasActiveSession())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_HAS_ACTIVE_SESSION);
rocmtools::GetROCMToolObj()->GetSession(session_id)->Start();
rocmtools::GetROCMToolObj()->SetCurrentActiveSession(session_id);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->HasFilter())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_MISSING_FILTER);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->HasBuffer())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_MISSING_BUFFER);
if (rocprofiler::GetROCProfilerSingleton()->HasActiveSession())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_HAS_ACTIVE_SESSION);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->Start();
rocprofiler::GetROCProfilerSingleton()->SetCurrentActiveSession(session_id);
API_METHOD_SUFFIX
}
// API to deactivate a session by id
ROCPROFILER_API rocprofiler_status_t rocprofiler_terminate_session(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocmtools::GetROCMToolObj()->IsActiveSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_ACTIVE);
rocmtools::GetROCMToolObj()->GetSession(session_id)->Terminate();
rocmtools::GetROCMToolObj()->SetCurrentActiveSession(rocprofiler_session_id_t{0});
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
if (!rocprofiler::GetROCProfilerSingleton()->IsActiveSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_ACTIVE);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->Terminate();
rocprofiler::GetROCProfilerSingleton()->SetCurrentActiveSession(rocprofiler_session_id_t{0});
API_METHOD_SUFFIX
}
@@ -656,32 +656,32 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_terminate_session(rocprofiler_s
ROCPROFILER_API rocprofiler_status_t rocprofiler_push_range(rocprofiler_session_id_t session_id,
const char* label) {
API_INIT_CHECKER
if (!label) throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_CORRUPTED_LABEL_DATA);
rocmtools::GetROCMToolObj()->GetSession(session_id)->PushRangeLabels(label);
if (!label) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_CORRUPTED_LABEL_DATA);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->PushRangeLabels(label);
API_METHOD_SUFFIX
}
// API to pop a custom label defined for a code section
ROCPROFILER_API rocprofiler_status_t rocprofiler_pop_range(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->PopRangeLabels())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_RANGE_STACK_IS_EMPTY);
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->PopRangeLabels())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_RANGE_STACK_IS_EMPTY);
API_METHOD_SUFFIX
}
ROCPROFILER_API rocprofiler_status_t rocprofiler_start_replay_pass(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->FindSession(session_id))
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
rocmtools::GetROCMToolObj()->GetSession(session_id)->GetProfiler()->StartReplayPass(session_id);
if (!rocprofiler::GetROCProfilerSingleton()->FindSession(session_id))
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetProfiler()->StartReplayPass(session_id);
API_METHOD_SUFFIX
}
ROCPROFILER_API rocprofiler_status_t rocprofiler_end_replay_pass(rocprofiler_session_id_t session_id) {
API_INIT_CHECKER
if (!rocmtools::GetROCMToolObj()->GetSession(session_id)->GetProfiler()->HasActivePass())
throw rocmtools::Exception(ROCPROFILER_STATUS_ERROR_PASS_NOT_STARTED);
rocmtools::GetROCMToolObj()->GetSession(session_id)->GetProfiler()->EndReplayPass();
if (!rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetProfiler()->HasActivePass())
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_PASS_NOT_STARTED);
rocprofiler::GetROCProfilerSingleton()->GetSession(session_id)->GetProfiler()->EndReplayPass();
API_METHOD_SUFFIX
}
@@ -691,7 +691,7 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_create
API_METHOD_PREFIX
std::vector<std::string> counters(counter_names, counter_names + num_counters);
*session_id =
rocmtools::GetROCMToolObj()->CreateDeviceProfilingSession(counters, cpu_index, gpu_index);
rocprofiler::GetROCProfilerSingleton()->CreateDeviceProfilingSession(counters, cpu_index, gpu_index);
API_METHOD_SUFFIX
}
@@ -699,7 +699,7 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_create
ROCPROFILER_API rocprofiler_status_t
rocprofiler_device_profiling_session_start(rocprofiler_session_id_t session_id) {
API_METHOD_PREFIX
rocmtools::GetROCMToolObj()->GetDeviceProfilingSession(session_id)->StartSession();
rocprofiler::GetROCProfilerSingleton()->GetDeviceProfilingSession(session_id)->StartSession();
API_METHOD_SUFFIX
}
@@ -707,7 +707,7 @@ rocprofiler_device_profiling_session_start(rocprofiler_session_id_t session_id)
ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_poll(
rocprofiler_session_id_t session_id, rocprofiler_device_profile_metric_t* data) {
API_METHOD_PREFIX
rocmtools::GetROCMToolObj()->GetDeviceProfilingSession(session_id)->PollMetrics(data);
rocprofiler::GetROCProfilerSingleton()->GetDeviceProfilingSession(session_id)->PollMetrics(data);
API_METHOD_SUFFIX
}
@@ -715,7 +715,7 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_poll(
ROCPROFILER_API rocprofiler_status_t
rocprofiler_device_profiling_session_stop(rocprofiler_session_id_t session_id) {
API_METHOD_PREFIX
rocmtools::GetROCMToolObj()->GetDeviceProfilingSession(session_id)->StopSession();
rocprofiler::GetROCProfilerSingleton()->GetDeviceProfilingSession(session_id)->StopSession();
API_METHOD_SUFFIX
}
@@ -723,7 +723,7 @@ rocprofiler_device_profiling_session_stop(rocprofiler_session_id_t session_id) {
ROCPROFILER_API rocprofiler_status_t
rocprofiler_device_profiling_session_destroy(rocprofiler_session_id_t session_id) {
API_METHOD_PREFIX
rocmtools::GetROCMToolObj()->DestroyDeviceProfilingSession(session_id);
rocprofiler::GetROCProfilerSingleton()->DestroyDeviceProfilingSession(session_id);
API_METHOD_SUFFIX
}
@@ -746,9 +746,9 @@ ROCPROFILER_EXPORT extern const uint32_t HSA_AMD_TOOL_PRIORITY = 25;
*/
ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
uint64_t failed_tool_count, const char* const* failed_tool_names) {
if (started) rocmtools::fatal("HSA Tool started already!");
if (started) rocprofiler::fatal("HSA Tool started already!");
started = true;
rocmtools::hsa_support::Initialize(table);
rocprofiler::hsa_support::Initialize(table);
return true;
}
@@ -756,8 +756,8 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
* @brief Callback function upon unloading the HSA.
*/
ROCPROFILER_EXPORT void OnUnload() {
if (!started) rocmtools::fatal("HSA Tool hasn't started yet!");
rocmtools::hsa_support::Finalize();
if (!started) rocprofiler::fatal("HSA Tool hasn't started yet!");
rocprofiler::hsa_support::Finalize();
started=false;
}
projects/rocprofiler/src/core/counters/counter.cpp
+1 -1
View File
@@ -22,7 +22,7 @@
#include <cassert>
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
namespace Counter {
projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp
+10 -9
View File
@@ -4,7 +4,7 @@
#include <set>
#include <math.h>
using namespace rocmtools;
using namespace rocprofiler;
struct block_des_t {
@@ -93,16 +93,17 @@ bool metrics::ExtractMetricEvents(
results_list holds the result objects for each event (which means, basic counters only)
*/
try {
uint32_t xcc_count = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle).getXccCount();
uint32_t xcc_count = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle).getXccCount();
for (size_t i = 0; i < metric_names.size(); i++) {
counters_vec_t counters_vec;
// TODO: saurabh
// const Metric* metric = metrics_dict->GetMetricByName(metric_names[i]);
const Metric* metric = metrics_dict->Get(metric_names[i]);
if (metric == nullptr) {
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
fatal("input metric'%s' not supported on this hardware: %s ", metric_names[i].c_str(),
agentInfo.getName().data());
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
fatal("input metric'%s' not supported on this hardware: %s ", metric_names[i].c_str(),
agentInfo.getName().data());
}
// adding result object for derived metric
@@ -113,7 +114,7 @@ bool metrics::ExtractMetricEvents(
counters_vec = metric->GetCounters();
if (counters_vec.empty())
rocmtools::fatal("bad metric '%s' is empty", metric_names[i].c_str());
rocprofiler::fatal("bad metric '%s' is empty", metric_names[i].c_str());
for (const counter_t* counter : counters_vec) {
results_t* result = nullptr;
@@ -163,7 +164,7 @@ bool metrics::ExtractMetricEvents(
// << ", Block Status Counter ID: " << block_status.counter_index << std::endl;
if (block_status.counter_index >= block_status.max_counters) {
rocmtools::fatal("Metrics specified have exceeded HW limits!");
rocprofiler::fatal("Metrics specified have exceeded HW limits!");
return false;
}
block_status.counter_index += 1;
@@ -187,7 +188,7 @@ bool metrics::ExtractMetricEvents(
bool metrics::GetCounterData(hsa_ven_amd_aqlprofile_profile_t* profile, hsa_agent_t gpu_agent,
std::vector<results_t*>& results_list) {
uint32_t xcc_count = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle).getXccCount();
uint32_t xcc_count = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle).getXccCount();
uint32_t single_xcc_buff_size = profile->output_buffer.size / (sizeof(uint64_t) * xcc_count);
callback_data_t callback_data{&results_list, 0, single_xcc_buff_size};
hsa_status_t status = hsa_ven_amd_aqlprofile_iterate_data(profile, pmcCallback, &callback_data);
@@ -201,7 +202,7 @@ bool metrics::GetMetricsData(std::map<std::string, results_t*>& results_map,
const xml::Expr* expr = metric->GetExpr();
if (expr) {
auto it = results_map.find(metric->GetName());
if (it == results_map.end()) rocmtools::fatal("metric results not found ");
if (it == results_map.end()) rocprofiler::fatal("metric results not found ");
results_t* res = it->second;
res->val_double = expr->Eval(args);
}
projects/rocprofiler/src/core/counters/metrics/eval_metrics.h
+6 -6
View File
@@ -31,7 +31,7 @@ THE SOFTWARE.
#include <atomic>
#include <set>
namespace rocmtools {
namespace rocprofiler {
typedef std::vector<double> xcc_results_t;
@@ -52,10 +52,10 @@ typedef struct {
packet_t* start_packet;
packet_t* stop_packet;
packet_t* read_packet;
rocmtools::MetricsDict* metrics_dict;
std::vector<const rocmtools::Metric*> metrics_list;
std::map<std::string, rocmtools::results_t*> results_map;
std::vector<rocmtools::results_t*> results_list;
rocprofiler::MetricsDict* metrics_dict;
std::vector<const rocprofiler::Metric*> metrics_list;
std::map<std::string, rocprofiler::results_t*> results_map;
std::vector<rocprofiler::results_t*> results_list;
std::vector<event_t> events_list;
hsa_agent_t gpu_agent;
hsa_signal_t begin_signal;
@@ -82,6 +82,6 @@ void GetCountersAndMetricResultsByXcc(uint32_t xcc_index, std::vector<results_t*
std::vector<const Metric*>& metrics_list);
} // namespace metrics
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_METRICS_EVALMETRICS_H_
projects/rocprofiler/src/core/counters/metrics/exception.h
+3 -3
View File
@@ -45,7 +45,7 @@ THE SOFTWARE.
do { \
std::ostringstream oss; \
oss << __FUNCTION__ << "(), " << stream; \
throw rocmtools::util::exception(error, oss.str()); \
throw rocprofiler::util::exception(error, oss.str()); \
} while (0)
#define AQL_EXC_RAISING(error, stream) \
@@ -55,7 +55,7 @@ THE SOFTWARE.
EXC_RAISING(error, stream << ", " << error_string); \
} while (0)
namespace rocmtools {
namespace rocprofiler {
namespace util {
class exception : public std::exception {
@@ -70,6 +70,6 @@ class exception : public std::exception {
};
} // namespace util
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_UTIL_EXCEPTION_H_
projects/rocprofiler/src/core/counters/metrics/metrics.cpp
+2 -2
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include "metrics.h"
namespace rocmtools {
namespace rocprofiler {
MetricsDict::map_t* MetricsDict::map_ = NULL;
MetricsDict::mutex_t MetricsDict::mutex_;
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/metrics/metrics.h
+4 -4
View File
@@ -41,7 +41,7 @@ THE SOFTWARE.
#include <mutex>
#include "src/core/hardware/hsa_info.h"
namespace rocmtools {
namespace rocprofiler {
struct counter_t {
std::string name;
event_t event;
@@ -105,8 +105,8 @@ class MetricsDict {
auto it = cache_->find(name);
if (it != cache_->end()) {
ret = true;
const rocmtools::ExprMetric* expr_metric =
dynamic_cast<const rocmtools::ExprMetric*>(it->second);
const rocprofiler::ExprMetric* expr_metric =
dynamic_cast<const rocprofiler::ExprMetric*>(it->second);
if (expr_metric) result = expr_metric->GetExpr()->GetStr();
}
return ret;
@@ -369,6 +369,6 @@ class MetricsDict {
static mutex_t mutex_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_COUNTERS_METRICS_METRICS_H_
projects/rocprofiler/src/core/counters/metrics/types.h
+2 -2
View File
@@ -27,7 +27,7 @@ THE SOFTWARE.
#include <hsa/hsa_ven_amd_aqlprofile.h>
namespace rocmtools {
namespace rocprofiler {
typedef hsa_ven_amd_aqlprofile_pfn_t pfn_t;
typedef hsa_ven_amd_aqlprofile_event_t event_t;
typedef hsa_ven_amd_aqlprofile_parameter_t parameter_t;
@@ -46,6 +46,6 @@ inline std::ostream& operator<<(std::ostream& out, const parameter_t& parameter)
return out;
}
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_TYPES_H_
projects/rocprofiler/src/core/counters/mmio/df_counters_mi200.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
#include "df_perfmon_registers_mi200.h"
#include "mmio.h"
namespace rocmtools {
namespace rocprofiler {
#define DF_V3_6_MAX_COUNTERS 4
@@ -73,4 +73,4 @@ void DFPerfMonMI200::Start() {
printf("CakeDlwmActiveTransferCount_2=%u\n", ficaa_out_val);
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/mmio/df_counters_mi200.h
+3 -3
View File
@@ -5,7 +5,7 @@
#include "mmio.h"
#include "perfmon.h"
namespace rocmtools {
namespace rocprofiler {
/*
One perfmon per GPU.
@@ -30,7 +30,7 @@ class DFPerfMonMI200 : public PerfMon {
void SetFicaNodeOutboundBw(uint64_t node_instance, uint32_t& ficaa_in_val);
uint64_t GetFicaNodeOutboundBw(uint32_t ficaa_val);
private:
mmio::DFPerfmonMMIO* mmio_;
@@ -39,6 +39,6 @@ class DFPerfMonMI200 : public PerfMon {
uint64_t instance_id_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif // DF_COUNTERS_MI200_H
projects/rocprofiler/src/core/counters/mmio/mmio.cpp
+2 -2
View File
@@ -23,7 +23,7 @@
#include "../../../utils/helper.h"
#include "pcie_perfmon_registers_mi200.h"
#include "df_perfmon_registers_mi200.h"
namespace rocmtools {
namespace rocprofiler {
namespace mmio {
@@ -190,4 +190,4 @@ std::map<decltype(hsa_agent_t::handle), std::vector<MMIO*>> MMIOManager::mmio_in
} // namespace mmio
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/mmio/mmio.h
+2 -2
View File
@@ -31,7 +31,7 @@
#include <unistd.h>
#include <sstream>
namespace rocmtools {
namespace rocprofiler {
namespace mmio {
@@ -126,7 +126,7 @@ class MMIOManager {
} // namespace mmio
} // namespace rocmtools
} // namespace rocprofiler
#endif
projects/rocprofiler/src/core/counters/mmio/pcie_counters_mi200.cpp
+4 -4
View File
@@ -2,7 +2,7 @@
#include "pcie_perfmon_registers_mi200.h"
#include "perfmon.h"
namespace rocmtools {
namespace rocprofiler {
PciePerfMonMI200::PciePerfMonMI200(const Agent::AgentInfo& info) : PerfMon(), mmio_(nullptr) {
mmio_ =
@@ -75,7 +75,7 @@ void PciePerfMonMI200::Start_RX_TILE_TXCLK(uint32_t event){
void PciePerfMonMI200::Read_RX_TILE_TXCLK(uint64_t& result){
// Step 5: Performance counters read:
uint32_t lo_val, hi_val;
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_TXCLK3, lo_val);
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_TXCLK3, lo_val);
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_UPVAL_TXCLK3, hi_val);
// Combine the lo and hi values and put them in result
@@ -102,7 +102,7 @@ void PciePerfMonMI200::Start_RX_TILE_SCLK(uint32_t event){
void PciePerfMonMI200::Read_RX_TILE_SCLK(uint64_t& result){
// Step 5: Performance counters read:
uint32_t lo_val, hi_val;
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_LCLK1, lo_val);
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_LCLK1, lo_val);
readRegister(PCIE_MI200::PCIE_PERF_COUNT0_UPVAL_LCLK1, hi_val);
// Combine the lo and hi values and put them in result
@@ -111,6 +111,6 @@ void PciePerfMonMI200::Read_RX_TILE_SCLK(uint64_t& result){
result = val | lo_val;
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/mmio/pcie_counters_mi200.h
+3 -3
View File
@@ -4,7 +4,7 @@
#include "mmio.h"
#include "perfmon.h"
namespace rocmtools {
namespace rocprofiler {
/*
One perfmon per GPU.
@@ -20,7 +20,7 @@ class PciePerfMonMI200 : public PerfMon {
void Stop() override;
void Read(std::vector<rocprofiler_counters_sampler_counter_output_t>& values) override;
mmio::mmap_type_t Type() override { return mmio::mmap_type_t::PCIE_PERFMON; }
private:
// TODO : check google coding std
void writeRegister(uint32_t reg_offset, uint32_t value);
@@ -38,6 +38,6 @@ class PciePerfMonMI200 : public PerfMon {
int event_id_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif
projects/rocprofiler/src/core/counters/mmio/perfmon.h
+2 -2
View File
@@ -25,7 +25,7 @@
#include "mmio.h"
#include <vector>
namespace rocmtools {
namespace rocprofiler {
class PerfMon {
public:
@@ -42,6 +42,6 @@ class PerfMon {
std::vector<std::string> counter_names_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif
projects/rocprofiler/src/core/hardware/hsa_info.cpp
+17 -12
View File
@@ -38,7 +38,7 @@ namespace Agent {
AgentInfo::AgentInfo() {}
AgentInfo::AgentInfo(const hsa_agent_t agent, ::CoreApiTable* table) : handle_(agent.handle) {
if (table->hsa_agent_get_info_fn(agent, HSA_AGENT_INFO_DEVICE, &type_) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info failed");
rocprofiler::fatal("hsa_agent_get_info failed");
table->hsa_agent_get_info_fn(agent, HSA_AGENT_INFO_NAME, name_);
@@ -62,22 +62,27 @@ AgentInfo::AgentInfo(const hsa_agent_t agent, ::CoreApiTable* table) : handle_(a
table->hsa_agent_get_info_fn(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES), &se_num_);
if (table->hsa_agent_get_info_fn(agent,
(hsa_agent_info_t)HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE,
&shader_arrays_per_se_) != HSA_STATUS_SUCCESS ||
table->hsa_agent_get_info_fn(agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU,
&waves_per_cu_) != HSA_STATUS_SUCCESS) {
rocmtools::fatal("hsa_agent_get_info for gfxip hardware configuration failed");
if (table->hsa_agent_get_info_fn(
agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE,
&shader_arrays_per_se_) != HSA_STATUS_SUCCESS ||
table->hsa_agent_get_info_fn(
agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU,
&waves_per_cu_) != HSA_STATUS_SUCCESS)
{
rocprofiler::fatal("hsa_agent_get_info for gfxip hardware configuration failed");
}
compute_units_per_sh_ = cu_num_ / (se_num_ * shader_arrays_per_se_);
wave_slots_per_simd_ = waves_per_cu_ / simds_per_cu_;
if (table->hsa_agent_get_info_fn(agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_DOMAIN,
&pci_domain_) != HSA_STATUS_SUCCESS ||
table->hsa_agent_get_info_fn(agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID,
&pci_location_id_) != HSA_STATUS_SUCCESS) {
rocmtools::fatal("hsa_agent_get_info for PCI info failed");
if (table->hsa_agent_get_info_fn(
agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_DOMAIN,
&pci_domain_) != HSA_STATUS_SUCCESS ||
table->hsa_agent_get_info_fn(
agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID,
&pci_location_id_) != HSA_STATUS_SUCCESS)
{
rocprofiler::fatal("hsa_agent_get_info for PCI info failed");
}
// TODO: (sauverma) use hsa_agent_get_info_fn(HSA_AMD_AGENT_INFO_NUM_XCC)
projects/rocprofiler/src/core/hardware/hsa_info.h
+1 -1
View File
@@ -36,7 +36,7 @@ namespace Agent {
static const uint32_t LDS_BLOCK_SIZE = 128 * 4;
// XXX TODO: This should be merged into rocmtools::hsa_support::AgentInfo and
// XXX TODO: This should be merged into rocprofiler::hsa_support::AgentInfo and
// this file should be removed entirely, as it's completely redundant
class AgentInfo {
public:
projects/rocprofiler/src/core/hsa/hsa_common.cpp
+2 -2
View File
@@ -22,7 +22,7 @@
#include "src/utils/exception.h"
namespace rocmtools {
namespace rocprofiler {
namespace hsa_support {
@@ -122,4 +122,4 @@ rocprofiler_timestamp_t GetCurrentTimestampNS() {
}
} // namespace hsa_support
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/hsa_common.h
+2 -2
View File
@@ -35,7 +35,7 @@
#define ASSERTM(exp, msg) assert(((void)msg, exp))
namespace rocmtools {
namespace rocprofiler {
namespace hsa_support {
@@ -59,6 +59,6 @@ void ResetMaps();
rocprofiler_timestamp_t GetCurrentTimestampNS();
} // namespace hsa_support
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_HSA_HSA_COMMON_H_
projects/rocprofiler/src/core/hsa/hsa_support.cpp
+91 -91
View File
@@ -48,7 +48,7 @@
#include "src/utils/helper.h"
#include "src/core/hsa/queues/queue.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include <experimental/filesystem>
namespace fs = std::experimental::filesystem;
@@ -58,11 +58,11 @@ namespace {
hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_agent_t agent,
hsa_executable_symbol_t symbol, void* args) {
hsa_symbol_kind_t type;
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &type);
if (type == HSA_SYMBOL_KIND_KERNEL) {
uint32_t name_length;
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
// TODO(aelwazir): to be removed if the HSA fixed the issue of corrupted
// names overflowing the length given
@@ -70,17 +70,17 @@ hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_
if (!(*static_cast<bool*>(args))) {
char name[name_length + 1];
uint64_t kernel_object;
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
std::string kernel_name = std::string(name).substr(0, name_length);
rocmtools::AddKernelName(kernel_object, kernel_name);
rocprofiler::AddKernelName(kernel_object, kernel_name);
} else {
uint64_t kernel_object;
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
rocmtools::RemoveKernelName(kernel_object);
rocprofiler::RemoveKernelName(kernel_object);
}
}
}
@@ -167,17 +167,17 @@ class Tracker {
entry->valid.store(ENTRY_INIT, std::memory_order_release);
// Creating a proxy signal
status = rocmtools::hsa_support::GetCoreApiTable().hsa_signal_create_fn(1, 0, NULL,
status = rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_create_fn(1, 0, NULL,
&(entry->signal));
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_signal_create failed");
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
if (status != HSA_STATUS_SUCCESS) rocprofiler::fatal("hsa_signal_create failed");
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
entry->signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_amd_signal_async_handler failed");
if (status != HSA_STATUS_SUCCESS) rocprofiler::fatal("hsa_amd_signal_async_handler failed");
}
// Delete tracker entry
inline static void Disable(entry_t* entry) {
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(entry->signal);
rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(entry->signal);
entry->valid.store(ENTRY_INV, std::memory_order_release);
}
@@ -186,19 +186,19 @@ class Tracker {
inline static void Complete(hsa_signal_value_t signal_value, entry_t* entry) {
static roctracer_timestamp_t sysclock_period = []() {
uint64_t sysclock_hz = 0;
hsa_status_t status = rocmtools::hsa_support::GetCoreApiTable().hsa_system_get_info_fn(
hsa_status_t status = rocprofiler::hsa_support::GetCoreApiTable().hsa_system_get_info_fn(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_system_get_info failed");
if (status != HSA_STATUS_SUCCESS) rocprofiler::fatal("hsa_system_get_info failed");
return (uint64_t)1000000000 / sysclock_hz;
}();
if (entry->type == COPY_ENTRY_TYPE) {
hsa_amd_profiling_async_copy_time_t async_copy_time{};
hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_get_async_copy_time_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_profiling_get_async_copy_time_fn(
entry->signal, &async_copy_time);
if (status != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_amd_profiling_get_async_copy_time failed");
rocprofiler::fatal("hsa_amd_profiling_get_async_copy_time failed");
entry->begin = async_copy_time.start * sysclock_period;
entry->end = async_copy_time.end * sysclock_period;
} else {
@@ -222,11 +222,11 @@ class Tracker {
orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
[[maybe_unused]] const hsa_signal_value_t new_value =
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(orig) - 1;
rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(orig) - 1;
assert(signal_value == new_value && "Tracker::Complete bad signal value");
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_store_screlease_fn(orig, signal_value);
rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_store_screlease_fn(orig, signal_value);
}
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(signal);
rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(signal);
delete entry;
}
@@ -244,18 +244,18 @@ class Tracker {
hsa_status_t HSA_API MemoryAllocateIntercept(hsa_region_t region, size_t size, void** ptr) {
hsa_status_t status =
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_allocate_fn(region, size, ptr);
rocprofiler::hsa_support::GetCoreApiTable().hsa_memory_allocate_fn(region, size, ptr);
if (status != HSA_STATUS_SUCCESS) return status;
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE)) {
hsa_evt_data_t data{};
data.allocate.ptr = *ptr;
data.allocate.size = size;
if (rocmtools::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
if (rocprofiler::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
region, HSA_REGION_INFO_SEGMENT, &data.allocate.segment) != HSA_STATUS_SUCCESS ||
rocmtools::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
region, HSA_REGION_INFO_GLOBAL_FLAGS, &data.allocate.global_flag) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_region_get_info failed");
rocprofiler::fatal("hsa_region_get_info failed");
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE, &data);
}
@@ -266,15 +266,15 @@ hsa_status_t HSA_API MemoryAllocateIntercept(hsa_region_t region, size_t size, v
hsa_status_t MemoryAssignAgentIntercept(void* ptr, hsa_agent_t agent,
hsa_access_permission_t access) {
hsa_status_t status =
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_assign_agent_fn(ptr, agent, access);
rocprofiler::hsa_support::GetCoreApiTable().hsa_memory_assign_agent_fn(ptr, agent, access);
if (status != HSA_STATUS_SUCCESS) return status;
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE)) {
hsa_evt_data_t data{};
data.device.ptr = ptr;
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
if (rocprofiler::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
agent, HSA_AGENT_INFO_DEVICE, &data.device.type) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info failed");
rocprofiler::fatal("hsa_agent_get_info failed");
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE, &data);
}
@@ -284,7 +284,7 @@ hsa_status_t MemoryAssignAgentIntercept(void* ptr, hsa_agent_t agent,
hsa_status_t MemoryCopyIntercept(void* dst, const void* src, size_t size) {
hsa_status_t status =
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(dst, src, size);
rocprofiler::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(dst, src, size);
if (status != HSA_STATUS_SUCCESS) return status;
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_MEMCOPY)) {
@@ -301,7 +301,7 @@ hsa_status_t MemoryCopyIntercept(void* dst, const void* src, size_t size) {
hsa_status_t MemoryPoolAllocateIntercept(hsa_amd_memory_pool_t pool, size_t size, uint32_t flags,
void** ptr) {
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
hsa_status_t status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
pool, size, flags, ptr);
if (size == 0 || status != HSA_STATUS_SUCCESS) return status;
@@ -310,12 +310,12 @@ hsa_status_t MemoryPoolAllocateIntercept(hsa_amd_memory_pool_t pool, size_t size
data.allocate.ptr = *ptr;
data.allocate.size = size;
if (rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
if (rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &data.allocate.segment) != HSA_STATUS_SUCCESS ||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &data.allocate.global_flag) !=
HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_region_get_info failed");
rocprofiler::fatal("hsa_region_get_info failed");
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE, &data);
}
@@ -326,13 +326,13 @@ hsa_status_t MemoryPoolAllocateIntercept(hsa_amd_memory_pool_t pool, size_t size
auto [pool, ptr] = *reinterpret_cast<decltype(callback_data)*>(iterate_agent_callback_data);
if (hsa_amd_memory_pool_access_t value;
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_memory_pool_get_info_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agent_memory_pool_get_info_fn(
agent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &value) != HSA_STATUS_SUCCESS ||
value != HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT)
return HSA_STATUS_SUCCESS;
auto it = agent_info_map.find(agent.handle);
if (it == agent_info_map.end()) rocmtools::fatal("agent was not found in the agent_info map");
if (it == agent_info_map.end()) rocprofiler::fatal("agent was not found in the agent_info map");
hsa_evt_data_t data{};
data.device.type = it->second.type;
@@ -343,7 +343,7 @@ hsa_status_t MemoryPoolAllocateIntercept(hsa_amd_memory_pool_t pool, size_t size
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE, &data);
return HSA_STATUS_SUCCESS;
};
rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(agent_callback, &callback_data);
rocprofiler::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(agent_callback, &callback_data);
}
return HSA_STATUS_SUCCESS;
@@ -358,7 +358,7 @@ hsa_status_t MemoryPoolFreeIntercept(void* ptr) {
}
if (ptr)
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(ptr);
return rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(ptr);
else
return HSA_STATUS_SUCCESS;
}
@@ -366,7 +366,7 @@ hsa_status_t MemoryPoolFreeIntercept(void* ptr) {
// Agent allow access callback 'hsa_amd_agents_allow_access'
hsa_status_t AgentsAllowAccessIntercept(uint32_t num_agents, const hsa_agent_t* agents,
const uint32_t* flags, const void* ptr) {
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
hsa_status_t status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
num_agents, agents, flags, ptr);
if (status != HSA_STATUS_SUCCESS) return status;
@@ -374,7 +374,7 @@ hsa_status_t AgentsAllowAccessIntercept(uint32_t num_agents, const hsa_agent_t*
while (num_agents--) {
hsa_agent_t agent = *agents++;
auto it = agent_info_map.find(agent.handle);
if (it == agent_info_map.end()) rocmtools::fatal("agent was not found in the agent_info map");
if (it == agent_info_map.end()) rocprofiler::fatal("agent was not found in the agent_info map");
hsa_evt_data_t data{};
data.device.type = it->second.type;
@@ -398,62 +398,62 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object, void* arg) {
hsa_evt_data_t data{};
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
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 (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
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 (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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 ||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
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 {
rocmtools::fatal("unknown code object storage type: %d", data.codeobj.storage_type);
rocprofiler::fatal("unknown code object storage type: %d", data.codeobj.storage_type);
}
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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 ||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
rocprofiler::hsa_support::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 ||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
std::string uri_str(data.codeobj.uri_length, '\0');
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
if (rocprofiler::hsa_support::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)
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
data.codeobj.uri = uri_str.c_str();
data.codeobj.unload = *static_cast<bool*>(arg) ? 1 : 0;
@@ -467,12 +467,12 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
hsa_status_t ExecutableFreezeIntercept(hsa_executable_t executable, const char* options) {
hsa_status_t status =
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_freeze_fn(executable, options);
rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_freeze_fn(executable, options);
if (status != HSA_STATUS_SUCCESS) return status;
// if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ)) {
bool unload = false;
rocmtools::hsa_support::GetHSALoaderApi()
rocprofiler::hsa_support::GetHSALoaderApi()
.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(executable, CodeObjectCallback,
&unload);
// }
@@ -483,19 +483,19 @@ hsa_status_t ExecutableFreezeIntercept(hsa_executable_t executable, const char*
hsa_status_t ExecutableDestroyIntercept(hsa_executable_t executable) {
// if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ)) {
bool unload = true;
rocmtools::hsa_support::GetHSALoaderApi()
rocprofiler::hsa_support::GetHSALoaderApi()
.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(executable, CodeObjectCallback,
&unload);
// }
return rocmtools::hsa_support::GetCoreApiTable().hsa_executable_destroy_fn(executable);
return rocprofiler::hsa_support::GetCoreApiTable().hsa_executable_destroy_fn(executable);
}
std::atomic<bool> profiling_async_copy_enable{false};
hsa_status_t ProfilingAsyncCopyEnableIntercept(bool enable) {
hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(enable);
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(enable);
if (status == HSA_STATUS_SUCCESS) {
profiling_async_copy_enable.exchange(enable, std::memory_order_release);
}
@@ -521,12 +521,12 @@ hsa_status_t MemoryASyncCopyIntercept(void* dst, hsa_agent_t dst_agent, const vo
// FIXME: what happens if the state changes before returning?
[[maybe_unused]] hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
profiling_async_copy_enable.load(std::memory_order_relaxed) || is_enabled);
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");
if (!is_enabled) {
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
return rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal);
}
@@ -535,7 +535,7 @@ hsa_status_t MemoryASyncCopyIntercept(void* dst, hsa_agent_t dst_agent, const vo
entry->correlation_id = CorrelationId();
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, entry->signal);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
@@ -553,12 +553,12 @@ hsa_status_t MemoryASyncCopyRectIntercept(const hsa_pitched_ptr_t* dst,
// FIXME: what happens if the state changes before returning?
[[maybe_unused]] hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
profiling_async_copy_enable.load(std::memory_order_relaxed) || is_enabled);
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");
if (!is_enabled) {
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
return rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals,
completion_signal);
}
@@ -568,7 +568,7 @@ hsa_status_t MemoryASyncCopyRectIntercept(const hsa_pitched_ptr_t* dst,
entry->correlation_id = CorrelationId();
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals,
entry->signal);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
@@ -612,7 +612,7 @@ rocprofiler_timestamp_t timestamp_ns() {
// If the HSA intercept is installed, then use the "original"
// 'hsa_system_get_info' function to avoid reporting calls for internal use
// of the HSA API by the tracer.
auto hsa_system_get_info_fn = rocmtools::hsa_support::GetCoreApiTable().hsa_system_get_info_fn;
auto hsa_system_get_info_fn = rocprofiler::hsa_support::GetCoreApiTable().hsa_system_get_info_fn;
// If the HSA intercept is not installed, use the default
// 'hsa_system_get_info'.
@@ -623,14 +623,14 @@ rocprofiler_timestamp_t timestamp_ns() {
status == HSA_STATUS_ERROR_NOT_INITIALIZED)
return rocprofiler_timestamp_t{0};
else if (status != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_system_get_info failed");
rocprofiler::fatal("hsa_system_get_info failed");
static uint64_t sysclock_period = [&]() {
uint64_t sysclock_hz = 0;
if (hsa_status_t status =
hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
status != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_system_get_info failed");
rocprofiler::fatal("hsa_system_get_info failed");
return (uint64_t)1000000000 / sysclock_hz;
}();
@@ -640,17 +640,17 @@ rocprofiler_timestamp_t timestamp_ns() {
void Initialize_roctracer(HsaApiTable* table) {
// Save the HSA core api and amd_ext api.
saved_core_api = rocmtools::hsa_support::GetCoreApiTable();
saved_amd_ext_api = rocmtools::hsa_support::GetAmdExtTable();
hsa_loader_api = rocmtools::hsa_support::GetHSALoaderApi();
saved_core_api = rocprofiler::hsa_support::GetCoreApiTable();
saved_amd_ext_api = rocprofiler::hsa_support::GetAmdExtTable();
hsa_loader_api = rocprofiler::hsa_support::GetHSALoaderApi();
// Enumerate the agents.
if (rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
if (rocprofiler::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
[](hsa_agent_t agent, void* data) {
hsa_support::AgentInfo agent_info;
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
if (rocprofiler::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
agent, HSA_AGENT_INFO_DEVICE, &agent_info.type) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info failed");
rocprofiler::fatal("hsa_agent_get_info failed");
switch (agent_info.type) {
case HSA_DEVICE_TYPE_CPU:
static int cpu_agent_count = 0;
@@ -658,10 +658,10 @@ void Initialize_roctracer(HsaApiTable* table) {
break;
case HSA_DEVICE_TYPE_GPU: {
uint32_t driver_node_id;
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
if (rocprofiler::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID),
&driver_node_id) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info failed");
rocprofiler::fatal("hsa_agent_get_info failed");
agent_info.id = driver_node_id;
} break;
@@ -674,7 +674,7 @@ void Initialize_roctracer(HsaApiTable* table) {
return HSA_STATUS_SUCCESS;
},
nullptr) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_iterate_agents failed");
rocprofiler::fatal("hsa_iterate_agents failed");
}
const char* GetApiName(uint32_t id) { return detail::GetApiName(id); }
@@ -722,7 +722,7 @@ void RegisterTracerCallback(int (*function)(activity_domain_t domain, uint32_t o
} // namespace roctracer::hsa_support
namespace rocmtools {
namespace rocprofiler {
namespace hsa_support {
@@ -845,11 +845,11 @@ void Initialize(HsaApiTable* table) {
case HSA_DEVICE_TYPE_CPU:
agent_info.setIndex(cpu_agent_count++);
cpu_agent = agent;
rocmtools::queue::InitializePools(cpu_agent, &agent_info);
rocprofiler::queue::InitializePools(cpu_agent, &agent_info);
uint32_t cpu_numa_node_id;
if (GetCoreApiTable().hsa_agent_get_info_fn(
agent, HSA_AGENT_INFO_NODE, &cpu_numa_node_id) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info(HSA_AGENT_INFO_NODE) failed");
rocprofiler::fatal("hsa_agent_get_info(HSA_AGENT_INFO_NODE) failed");
agent_info.setNumaNode(cpu_numa_node_id);
numa_node_to_cpu_agent[cpu_numa_node_id] = agent;
break;
@@ -871,11 +871,11 @@ void Initialize(HsaApiTable* table) {
uint32_t gpu_cpu_numa_node_id;
if (GetCoreApiTable().hsa_agent_get_info_fn(
agent, HSA_AGENT_INFO_NODE, &gpu_cpu_numa_node_id) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info(HSA_AGENT_INFO_NODE) failed");
rocprofiler::fatal("hsa_agent_get_info(HSA_AGENT_INFO_NODE) failed");
agent_info.setNumaNode(gpu_cpu_numa_node_id);
agent_info.setNearCpuAgent(
numa_node_to_cpu_agent[gpu_numa_nodes_near_cpu[gpu_cpu_numa_node_id]]);
rocmtools::queue::InitializeGPUPool(agent, &agent_info);
rocprofiler::queue::InitializeGPUPool(agent, &agent_info);
gpu_agents.push_back(agent);
break;
default:
@@ -886,7 +886,7 @@ void Initialize(HsaApiTable* table) {
return HSA_STATUS_SUCCESS;
},
nullptr) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_iterate_agents failed");
rocprofiler::fatal("hsa_iterate_agents failed");
for (auto& agent : gpu_agents) {
GetAgentInfo(agent.handle).cpu_pool =
@@ -895,7 +895,7 @@ void Initialize(HsaApiTable* table) {
GetAgentInfo(GetAgentInfo(agent.handle).getNearCpuAgent().handle).kernarg_pool;
}
rocmtools::queue::CheckPacketReqiurements(gpu_agents);
rocprofiler::queue::CheckPacketReqiurements(gpu_agents);
gpu_agents.clear();
numa_node_to_cpu_agent.clear();
@@ -944,22 +944,22 @@ void Finalize() {
ResetMaps();
}
static std::map<uint64_t, rocmtools::MetricsDict*> metricsDicts;
static std::map<uint64_t, rocprofiler::MetricsDict*> metricsDicts;
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback) {
if (GetCoreApiTable().hsa_iterate_agents_fn(
[](hsa_agent_t agent, void* data) {
Agent::AgentInfo agent_info{agent, &GetCoreApiTable()};
if (agent_info.getType() == HSA_DEVICE_TYPE_GPU) {
metricsDicts.emplace(agent.handle, rocmtools::MetricsDict::Create(&agent_info));
metricsDicts.emplace(agent.handle, rocprofiler::MetricsDict::Create(&agent_info));
}
return HSA_STATUS_SUCCESS;
},
nullptr) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_iterate_agents failed");
rocprofiler::fatal("hsa_iterate_agents failed");
uint32_t gpu_counter = 0;
for (auto metricsDictAgent : metricsDicts) {
rocmtools::MetricsDict* metricsDict = metricsDictAgent.second;
rocprofiler::MetricsDict* metricsDict = metricsDictAgent.second;
std::string gpu_name = metricsDict->GetAgentName();
auto nodes_vec = metricsDict->GetNodes();
@@ -1031,4 +1031,4 @@ bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback
}
} // namespace hsa_support
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/hsa_support.h
+2 -2
View File
@@ -90,7 +90,7 @@ typedef struct {
} hsa_evt_data_t;
namespace rocmtools {
namespace rocprofiler {
namespace hsa_support {
@@ -101,7 +101,7 @@ void Finalize();
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
} // namespace hsa_support
} // namespace rocmtools
} // namespace rocprofiler
#include "src/core/session/tracer/src/roctracer.h"
projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp
+41 -41
View File
@@ -19,7 +19,7 @@
THE SOFTWARE. */
#include "packets_generator.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
@@ -83,13 +83,13 @@ static hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool
if (nullptr == data) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
err = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
err = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
ASSERTM(err != HSA_STATUS_ERROR, "hsa_amd_memory_pool_get_info");
if (HSA_AMD_SEGMENT_GLOBAL != segment) {
return HSA_STATUS_SUCCESS;
}
err = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
err = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
ASSERTM(err != HSA_STATUS_ERROR, "hsa_amd_memory_pool_get_info");
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
@@ -116,11 +116,11 @@ hsa_status_t FindKernArgPool(hsa_amd_memory_pool_t pool, void* data) {
void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info) {
hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindStandardPool, &(agent_info->cpu_pool));
CHECK_HSA_STATUS("Error: Command Buffer Pool is not initialized", status);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindKernArgPool, &(agent_info->kernarg_pool));
CHECK_HSA_STATUS("Error: Output Buffer Pool is not initialized", status);
}
@@ -136,30 +136,30 @@ struct block_des_t {
uint32_t index;
};
std::map<uint32_t, rocmtools::MetricsDict*> metricsDict;
std::map<uint32_t, rocprofiler::MetricsDict*> metricsDict;
static std::atomic<bool> counters_added{false};
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
for (auto& gpu_agent : gpu_agents) {
// get the instance of MetricsDict
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
metricsDict[gpu_agent.handle] = rocmtools::MetricsDict::Create(&agentInfo);
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
metricsDict[gpu_agent.handle] = rocprofiler::MetricsDict::Create(&agentInfo);
}
}
// 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<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_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, bool is_spm) {
hsa_status_t status = HSA_STATUS_SUCCESS;
if (!counters_added.load(std::memory_order_acquire)) {
for (auto& name : counter_names) {
if (rocmtools::GetROCMToolObj()->HasActiveSession()) {
rocmtools::GetROCMToolObj()
->GetSession(rocmtools::GetROCMToolObj()->GetCurrentSessionId())
if (rocprofiler::GetROCProfilerSingleton()->HasActiveSession()) {
rocprofiler::GetROCProfilerSingleton()
->GetSession(rocprofiler::GetROCProfilerSingleton()->GetCurrentSessionId())
->GetProfiler()
->AddCounterName(name);
}
@@ -167,14 +167,14 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
counters_added.exchange(true, std::memory_order_release);
}
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
std::map<std::string, rocmtools::results_t*> results_map;
std::vector<rocmtools::event_t> events_list;
std::vector<rocmtools::results_t*> results_list;
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
std::map<std::string, rocprofiler::results_t*> results_map;
std::vector<rocprofiler::event_t> events_list;
std::vector<rocprofiler::results_t*> results_list;
std::map<std::pair<uint32_t, uint32_t>, uint64_t> event_to_max_block_count;
std::map<std::string, std::set<std::string>> metrics_counters;
if (!rocmtools::metrics::ExtractMetricEvents(
if (!rocprofiler::metrics::ExtractMetricEvents(
counter_names, gpu_agent, metricsDict[gpu_agent.handle], results_map, events_list,
results_list, event_to_max_block_count, metrics_counters)) {
std::cerr << "Error: Failed to extract metric events" << std::endl;
@@ -192,12 +192,12 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
abort();
}
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles = std::vector<
std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>();
std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>();
// do {
rocmtools::profiling_context_t* context = new rocmtools::profiling_context_t();
rocprofiler::profiling_context_t* context = new rocprofiler::profiling_context_t();
context->gpu_agent = gpu_agent;
auto result = results_list.begin();
std::map<std::pair<uint32_t, uint32_t>, uint32_t> block_max_events_count;
@@ -224,12 +224,12 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::set<std::string> metrics_counters_taken;
for (auto result : context->results_list) {
rocmtools::Metric* metric;
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<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(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 {
@@ -268,8 +268,8 @@ 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);
rocmtools::Metric* metric =
const_cast<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(metric_name));
rocprofiler::Metric* metric =
const_cast<rocprofiler::Metric*>(metricsDict[gpu_agent.handle]->Get(metric_name));
if (metric == nullptr) std::cout << metric_name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
}
@@ -302,7 +302,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
0,
0};
size_t ag_list_count = 1; // rocmtools::hsa_support::GetCPUAgentList().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;
@@ -319,13 +319,13 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
<< "Error: Command buffer given size is " << size << std::endl;
abort();
}
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
status = rocprofiler::hsa_support::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,
rocmtools::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());*/
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;
@@ -333,7 +333,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
}
} else {
// Both the CPU and GPU can access the memory
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = rocprofiler::hsa_support::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);
}
@@ -347,13 +347,13 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
<< "Error: Output buffer given size is " << size << std::endl;
abort();
}
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
status = rocprofiler::hsa_support::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,
rocmtools::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle)
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle)
.getNumaNode());*/
if (profile->output_buffer.ptr == NULL) {
std::cerr << __FILE__ << ":" << __LINE__ << " "
@@ -361,7 +361,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
abort();
}
} else {
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = rocprofiler::hsa_support::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);
@@ -409,7 +409,7 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
size_t ag_list_count = 1;
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
@@ -420,17 +420,17 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
profile->command_buffer.ptr = nullptr;
if (size <= 0) return nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->command_buffer.ptr);
CHECK_HSA_STATUS("Error: GPU Agent can't have command buffer access", status);
} else {
profile->command_buffer.ptr = numa_alloc_onnode(
profile->command_buffer.size,
rocmtools::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());
if (profile->command_buffer.ptr != nullptr) {
status = HSA_STATUS_SUCCESS;
} else {
@@ -443,12 +443,12 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
size = profile->output_buffer.size;
profile->output_buffer.ptr = nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.gpu_pool, size, 0, reinterpret_cast<void**>(&(profile->output_buffer.ptr)));
CHECK_HSA_STATUS("Error: Can't Allocate Output Buffer", status);
// Both the CPU and GPU can access the kernel arguments
if (status == HSA_STATUS_SUCCESS) {
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->output_buffer.ptr);
CHECK_HSA_STATUS("Error: Can't allow access on the Output Buffer for the GPU", status);
memset(profile->output_buffer.ptr, 0x0, profile->output_buffer.size);
@@ -473,17 +473,17 @@ bool g_output_buffer_local = true;
// Allocate system memory accessible by both CPU and GPU
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool) {
size_t ag_list_count = 1; // rocmtools::hsa_support::GetCPUAgentList().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;
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
*cpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, buffer);
}
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
@@ -501,7 +501,7 @@ uint8_t* AllocateLocalMemory(size_t size, hsa_amd_memory_pool_t* gpu_pool) {
}
hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile) {
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
profile->command_buffer.ptr =
AllocateSysMemory(gpu_agent, profile->command_buffer.size, &agentInfo.cpu_pool);
profile->output_buffer.size = g_output_buffer_size;
projects/rocprofiler/src/core/hsa/packets/packets_generator.h
+1 -1
View File
@@ -39,7 +39,7 @@ namespace Packet {
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_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, bool is_spm = false);
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
projects/rocprofiler/src/core/hsa/queues/queue.cpp
+42 -43
View File
@@ -29,7 +29,7 @@
#include <numa.h>
#include "rocprofiler.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/core/hsa/packets/packets_generator.h"
#include "src/core/hsa/hsa_support.h"
#include "src/utils/helper.h"
@@ -54,7 +54,7 @@
std::mutex sessions_pending_signal_lock;
namespace rocmtools {
namespace rocprofiler {
std::atomic<uint32_t> ACTIVE_INTERRUPT_SIGNAL_COUNT{0};
@@ -273,7 +273,7 @@ rocprofiler_kernel_properties_t set_kernel_properties(hsa_kernel_dispatch_packet
namespace queue {
using rocmtools::GetROCMToolObj;
using rocprofiler::GetROCProfilerSingleton;
hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) {
@@ -314,7 +314,7 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
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 rocmtools::Metric* metric = pending->context->metrics_list[i];
const rocprofiler::Metric* metric = pending->context->metrics_list[i];
double value = 0;
std::string metric_name = metric->GetName();
auto it = pending->context->results_map.find(metric_name);
@@ -324,11 +324,11 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
counters[i] = (rocprofiler_record_counter_instance_t{
// TODO(aelwazir): Moving to span once C++20 is adopted, strdup can be
// removed after that
rocprofiler_counter_id_t{rocmtools::profiler::GetCounterID(metric_name)},
rocprofiler_counter_id_t{rocprofiler::profiler::GetCounterID(metric_name)},
rocprofiler_record_counter_value_t{value}});
}
record->counters = counters;
rocmtools::Session* session = GetROCMToolObj()->GetSession(pending->session_id);
rocprofiler::Session* session = GetROCProfilerSingleton()->GetSession(pending->session_id);
void* initial_handle = const_cast<rocprofiler_record_counter_instance_t*>(record->counters);
if (session->FindBuffer(pending->buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id);
@@ -368,7 +368,7 @@ void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent
if (buffer == NULL) fatal("Trace data buffer allocation failed");
auto status =
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size);
rocprofiler::hsa_support::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;
@@ -383,13 +383,13 @@ void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent
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 || !GetROCMToolObj() ||
!GetROCMToolObj()->GetSession(queue_info_session->session_id) ||
!GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetProfiler())
if (!queue_info_session || !GetROCProfilerSingleton() ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id) ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id)->GetProfiler())
return true;
rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id);
rocprofiler::Session* session = GetROCProfilerSingleton()->GetSession(queue_info_session->session_id);
std::lock_guard<std::mutex> lock(session->GetSessionLock());
rocmtools::profiler::Profiler* profiler = session->GetProfiler();
rocprofiler::profiler::Profiler* profiler = session->GetProfiler();
std::vector<pending_signal_t*> pending_signals = const_cast<std::vector<pending_signal_t*>&>(
profiler->GetPendingSignals(queue_info_session->writer_id));
@@ -429,19 +429,19 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id};
if (pending->session_id.handle == 0) {
pending->session_id = GetROCMToolObj()->GetCurrentSessionId();
pending->session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
}
if (pending->counters_count > 0 && pending->context->metrics_list.size() > 0 &&
pending->profile) {
if (xcc_id == 0) // call to GetCounterData() is required only once for a dispatch
rocmtools::metrics::GetCounterData(pending->profile, queue_info_session->agent,
rocprofiler::metrics::GetCounterData(pending->profile, queue_info_session->agent,
pending->context->results_list);
if (is_individual_xcc_mode)
rocmtools::metrics::GetCountersAndMetricResultsByXcc(
rocprofiler::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, pending->context->results_list, pending->context->results_map,
pending->context->metrics_list);
else
rocmtools::metrics::GetMetricsData(pending->context->results_map,
rocprofiler::metrics::GetMetricsData(pending->context->results_map,
pending->context->metrics_list);
AddRecordCounters(&record, pending);
} else {
@@ -455,12 +455,12 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
// 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 = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
hsa_status_t status = rocprofiler::hsa_support::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 = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
status = rocprofiler::hsa_support::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;
@@ -481,16 +481,16 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
}
bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
// TODO: finish implementation to iterate trace data and add it to rocmtools record
// TODO: finish implementation to iterate trace data and add it to rocprofiler record
// and generic buffer
auto queue_info_session = static_cast<queue_info_session_t*>(data);
if (!queue_info_session || !GetROCMToolObj() ||
!GetROCMToolObj()->GetSession(queue_info_session->session_id) ||
!GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetAttTracer())
if (!queue_info_session || !GetROCProfilerSingleton() ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id) ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id)->GetAttTracer())
return true;
rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id);
rocmtools::att::AttTracer* att_tracer = session->GetAttTracer();
rocprofiler::Session* session = GetROCProfilerSingleton()->GetSession(queue_info_session->session_id);
rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer();
std::vector<att_pending_signal_t>& pending_signals =
const_cast<std::vector<att_pending_signal_t>&>(
att_tracer->GetPendingSignals(queue_info_session->writer_id));
@@ -517,16 +517,16 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
rocprofiler_record_id_t{pending.kernel_descriptor}};
if (pending.session_id.handle == 0) {
pending.session_id = GetROCMToolObj()->GetCurrentSessionId();
pending.session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
}
if (session->FindBuffer(pending.buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
buffer->AddRecord(record);
}
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
hsa_status_t status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->output_buffer.ptr));
CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
status = rocprofiler::hsa_support::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;
@@ -601,27 +601,27 @@ uint32_t replay_mode_count = 0;
std::vector<std::string> kernel_profile_names;
std::vector<std::string> att_counters_names;
rocmtools::Session* session = nullptr;
rocprofiler::Session* session = nullptr;
void ResetSessionID() { session_id = rocprofiler_session_id_t{0}; }
void CheckNeededProfileConfigs() {
rocprofiler_session_id_t internal_session_id;
if (GetROCMToolObj())
if (GetROCProfilerSingleton())
// Getting Session ID
internal_session_id = GetROCMToolObj()->GetCurrentSessionId();
internal_session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
else
internal_session_id = {0};
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 && GetROCMToolObj()) {
session = GetROCMToolObj()->GetSession(session_id);
if (session_id.handle > 0 && GetROCProfilerSingleton()) {
session = GetROCProfilerSingleton()->GetSession(session_id);
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION);
rocmtools::Filter* filter = session->GetFilter(filter_id);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
session_data = filter->GetCounterData();
is_counter_collection_mode = true;
session_data_count = session_data.size();
@@ -631,12 +631,12 @@ void CheckNeededProfileConfigs() {
is_timestamp_collection_mode = true;
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION);
rocmtools::Filter* filter = session->GetFilter(filter_id);
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);
rocmtools::Filter* filter = session->GetFilter(filter_id);
rocprofiler::Filter* filter = session->GetFilter(filter_id);
att_parameters_data = filter->GetAttParametersData();
is_att_collection_mode = true;
buffer_id =
@@ -679,7 +679,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// hsa_ven_amd_aqlprofile_profile_t* profile;
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles;
@@ -706,7 +706,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
uint32_t profile_id = 0;
// do {
std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*> profile;
std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*> profile;
if (profiles.size() > 0 && replay_mode_count > 0) profile = profiles.at(profile_id);
uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
@@ -730,10 +730,9 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
* Only PC sampling relies on this right now, so it would be better to
* only generate an ID if PC sampling is active to conserve IDs, but it's
* unlikely 64 bits' worth of identifiers will be exhausted during the
* lifetime of the ROCMToolObj.
* lifetime of the ROCProfiler_Singleton.
*/
uint64_t correlation_id = dispatch_packet.reserved2;
// dispatch_packet.reserved2 = GetROCMToolObj()->GetUniqueKernelDispatchId();
CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
// Adding the dispatch packet newly created signal to the pending signals
@@ -741,7 +740,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
if (session) {
uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId();
uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (profiles.size() > 0 && replay_mode_count > 0) {
@@ -793,7 +792,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
transformed_packets.emplace_back(*pkt);
}
Agent::AgentInfo& agentInfo =
rocmtools::hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle);
rocprofiler::hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle);
// Creating Async Handler to be called every time the interrupt signal is
// marked complete
SignalAsyncHandler(
@@ -942,7 +941,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// 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 = GetROCMToolObj()->GetUniqueRecordId();
uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (session && profile) {
@@ -1049,4 +1048,4 @@ void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
}
} // namespace queue
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/queues/queue.h
+2 -2
View File
@@ -37,7 +37,7 @@
#include "src/core/session/profiler/profiler.h"
namespace rocmtools {
namespace rocprofiler {
void InitKsymbols();
void FinitKsymbols();
@@ -95,6 +95,6 @@ void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents);
void ResetSessionID();
} // namespace queue
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_HSA_QUEUES_QUEUE_H_
projects/rocprofiler/src/core/memory/generic_buffer.cpp
+3 -3
View File
@@ -24,7 +24,7 @@
#include <atomic>
#include "rocprofiler.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
namespace Memory {
@@ -69,8 +69,8 @@ GenericBuffer::GenericBuffer(rocprofiler_session_id_t session_id, rocprofiler_bu
GenericBuffer::~GenericBuffer() {
if (is_valid_.load(std::memory_order_release)) {
std::lock_guard lock(buffer_lock_);
// if (rocmtools::GetROCMToolObj()->GetSession(session_id_))
// rocmtools::GetROCMToolObj()->GetSession(session_id_)->DisableTools(id_);
// if (rocprofiler::GetROCProfiler_Singleton()->GetSession(session_id_))
// rocprofiler::GetROCProfiler_Singleton()->GetSession(session_id_)->DisableTools(id_);
Flush();
projects/rocprofiler/src/core/session/att/att.cpp
+2 -2
View File
@@ -21,7 +21,7 @@
#include "att.h"
#include <cassert>
namespace rocmtools {
namespace rocprofiler {
namespace att {
@@ -54,4 +54,4 @@ const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t w
} // namespace att
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/att/att.h
+2 -2
View File
@@ -30,7 +30,7 @@
#include "rocprofiler.h"
namespace rocmtools {
namespace rocprofiler {
typedef struct {
uint64_t kernel_descriptor;
@@ -70,7 +70,7 @@ class AttTracer {
} // namespace att
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_SESSION_ATT_ATT_H_
projects/rocprofiler/src/core/session/counters_sampler.cpp
+9 -9
View File
@@ -20,11 +20,11 @@
#include "counters_sampler.h"
#include "src/core/hsa/hsa_support.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/core/counters/mmio/pcie_counters_mi200.h"
#include "src/core/counters/mmio/df_counters_mi200.h"
namespace rocmtools {
namespace rocprofiler {
CountersSampler::CountersSampler(rocprofiler_buffer_id_t buffer_id,
rocprofiler_filter_id_t filter_id,
@@ -35,16 +35,16 @@ CountersSampler::CountersSampler(rocprofiler_buffer_id_t buffer_id,
pci_system_initialized_(pci_system_init() == 0)
{
params_ = rocmtools::GetROCMToolObj()
params_ = rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id_)
->GetFilter(filter_id_)
->GetCountersSamplerParameterData();
std::vector<hsa_agent_t> agents;
rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
[](hsa_agent_t agent, void* arg) {
auto& agents = *reinterpret_cast<std::vector<hsa_agent_t>*>(arg);
const auto& ai = rocmtools::hsa_support::GetAgentInfo(agent.handle);
const auto& ai = rocprofiler::hsa_support::GetAgentInfo(agent.handle);
if (ai.getType() == HSA_DEVICE_TYPE_GPU) {
agents.emplace_back(agent);
}
@@ -62,7 +62,7 @@ CountersSampler::CountersSampler(rocprofiler_buffer_id_t buffer_id,
}
if (pcie_counter_names.size() > 0) {
auto agentInfo = rocmtools::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle);
auto agentInfo = rocprofiler::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle);
if (agentInfo.getName() == "gfx90a") {
PciePerfMonMI200* perfmon = new PciePerfMonMI200(agentInfo);
perfmon->SetCounterNames(pcie_counter_names);
@@ -77,7 +77,7 @@ CountersSampler::CountersSampler(rocprofiler_buffer_id_t buffer_id,
}
if (xgmi_counter_names.size() > 0) {
auto agentInfo = rocmtools::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle);
auto agentInfo = rocprofiler::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle);
if (agentInfo.getName() == "gfx90a") {
DFPerfMonMI200* perfmon = new DFPerfMonMI200(agentInfo);
perfmon->SetCounterNames(xgmi_counter_names);
@@ -132,7 +132,7 @@ void CountersSampler::Stop() {
}
void CountersSampler::AddRecord(rocprofiler_record_counters_sampler_t& record) {
const auto tool = rocmtools::GetROCMToolObj();
const auto tool = rocprofiler::GetROCProfilerSingleton();
const auto session = tool->GetSession(session_id_);
const auto buffer = session->GetBuffer(buffer_id_);
@@ -179,4 +179,4 @@ void CountersSampler::SamplerLoop() {
}
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/counters_sampler.h
+3 -3
View File
@@ -28,7 +28,7 @@
#include <mutex>
#include <vector>
namespace rocmtools {
namespace rocprofiler {
class CountersSampler {
public:
@@ -52,11 +52,11 @@ class CountersSampler {
bool pci_system_initialized_{false};
rocprofiler_counters_sampler_parameters_t params_;
std::vector<PerfMon*> perfmon_instances_;
std::atomic<bool> keep_running_{false};
std::thread sampler_thread_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif
projects/rocprofiler/src/core/session/device_profiling.cpp
+3 -3
View File
@@ -18,7 +18,7 @@
#define QUEUE_NUM_PACKETS 64
static const size_t CMD_SLOT_SIZE_B = 0x40;
using namespace rocmtools;
using namespace rocprofiler;
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> pmc_callback_data_t;
@@ -305,7 +305,7 @@ void get_hsa_agents_list(devices_t* device_list) {
}
bool rocmtools::find_hsa_agent_cpu(uint64_t index, hsa_agent_t* agent) {
bool rocprofiler::find_hsa_agent_cpu(uint64_t index, hsa_agent_t* agent) {
devices_t device_list;
get_hsa_agents_list(&device_list);
@@ -315,7 +315,7 @@ bool rocmtools::find_hsa_agent_cpu(uint64_t index, hsa_agent_t* agent) {
return true;
}
bool rocmtools::find_hsa_agent_gpu(uint64_t index, hsa_agent_t* agent) {
bool rocprofiler::find_hsa_agent_gpu(uint64_t index, hsa_agent_t* agent) {
devices_t device_list;
get_hsa_agents_list(&device_list);
projects/rocprofiler/src/core/session/device_profiling.h
+2 -2
View File
@@ -29,7 +29,7 @@
#include "src/core/counters/metrics/eval_metrics.h"
namespace rocmtools {
namespace rocprofiler {
class DeviceProfileSession {
public:
@@ -81,7 +81,7 @@ bool find_hsa_agent_cpu(uint64_t index, hsa_agent_t* agent);
bool find_hsa_agent_gpu(uint64_t index, hsa_agent_t* agent);
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_SESSION_DEVICE_PROFILING_H_
projects/rocprofiler/src/core/session/filter.cpp
+4 -4
View File
@@ -24,7 +24,7 @@
#include "src/utils/helper.h"
namespace rocmtools {
namespace rocprofiler {
Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind,
rocprofiler_filter_data_t filter_data, uint64_t data_count)
@@ -64,7 +64,7 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind
tracer_apis_.clear();
for (uint32_t j = 0; j < data_count; j++){
tracer_apis_.emplace_back(filter_data.trace_apis[j]);
}
}
break;
}
case ROCPROFILER_COUNTERS_SAMPLER: {
@@ -133,7 +133,7 @@ rocprofiler_counters_sampler_parameters_t Filter::GetCountersSamplerParameterDat
return counters_sampler_parameters_;
}
fatal(
"Error: ROCMtools filter specified is not supported for "
"Error: ROCProfiler filter specified is not supported for "
"Counters sampler mode!\n");
}
@@ -267,4 +267,4 @@ size_t Filter::GetPropertiesCount(rocprofiler_filter_property_kind_t kind) {
"kind!");
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/filter.h
+2 -2
View File
@@ -29,7 +29,7 @@
#define ASSERTM(exp, msg) assert(((void)msg, exp))
namespace rocmtools {
namespace rocprofiler {
class Filter {
public:
@@ -80,6 +80,6 @@ class Filter {
rocprofiler_sync_callback_t callback_;
};
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_SESSION_FILTER_H_
projects/rocprofiler/src/core/session/profiler/profiler.cpp
+3 -3
View File
@@ -40,7 +40,7 @@
#define ASSERTM(exp, msg) assert(((void)msg, exp))
namespace rocmtools {
namespace rocprofiler {
namespace profiler {
uint64_t GetCounterID(std::string& counter_name) {
@@ -127,7 +127,7 @@ bool Profiler::HasActivePass() {
void Profiler::AddPendingSignals(
uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& completion_signal,
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
rocmtools::profiling_context_t* context, uint64_t session_data_count,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile, rocprofiler_kernel_properties_t kernel_properties,
uint32_t thread_id, uint64_t queue_index, uint64_t correlation_id) {
std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
@@ -151,4 +151,4 @@ bool Profiler::CheckPendingSignalsIsEmpty() {
}
} // namespace profiler
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/profiler/profiler.h
+4 -4
View File
@@ -40,14 +40,14 @@ typedef void (*rocprofiler_add_profiler_record_t)(rocprofiler_record_profiler_t&
typedef rocprofiler_timestamp_t (*rocprofiler_get_timestamp_t)();
namespace rocmtools {
namespace rocprofiler {
typedef struct {
uint64_t kernel_descriptor;
hsa_signal_t signal;
rocprofiler_session_id_t session_id;
rocprofiler_buffer_id_t buffer_id;
rocmtools::profiling_context_t* context;
rocprofiler::profiling_context_t* context;
uint64_t counters_count;
hsa_ven_amd_aqlprofile_profile_t* profile;
rocprofiler_kernel_properties_t kernel_properties;
@@ -69,7 +69,7 @@ class Profiler {
void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id,
rocmtools::profiling_context_t* context, uint64_t session_data_count,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
uint64_t queue_index, uint64_t correlation_id);
@@ -102,6 +102,6 @@ class Profiler {
};
} // namespace profiler
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_TOOLS_PROFILER_PROFILER_H_
projects/rocprofiler/src/core/session/session.cpp
+4 -4
View File
@@ -38,7 +38,7 @@
#include "src/utils/helper.h"
#include "src/core/hsa/queues/queue.h"
namespace rocmtools {
namespace rocprofiler {
Session::Session(rocprofiler_replay_mode_t replay_mode, rocprofiler_session_id_t session_id)
: session_id_(session_id), is_active_(false), replay_mode_(replay_mode) {
@@ -49,7 +49,7 @@ Session::~Session() {
while (GetCurrentActiveInterruptSignalsCount() > 0) {
}
if (profiler_started_.load(std::memory_order_release)) {
rocmtools::queue::ResetSessionID();
rocprofiler::queue::ResetSessionID();
delete profiler_;
profiler_started_.exchange(false, std::memory_order_release);
}
@@ -176,7 +176,7 @@ void Session::Terminate() {
if (is_active_) {
while (GetCurrentActiveInterruptSignalsCount() > 0) {
}
rocmtools::queue::ResetSessionID();
rocprofiler::queue::ResetSessionID();
std::lock_guard<std::mutex> lock(session_lock_);
if (FindFilterWithKind(ROCPROFILER_SPM_COLLECTION)) {
{
@@ -390,4 +390,4 @@ uint64_t GenerateUniqueSessionId() {
return SESSION_COUNTER.fetch_add(1, std::memory_order_release);
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/session.h
+2 -2
View File
@@ -44,7 +44,7 @@
#define ASSERTM(exp, msg) assert(((void)msg, exp))
namespace rocmtools {
namespace rocprofiler {
class Session {
public:
@@ -135,6 +135,6 @@ class Session {
uint64_t GenerateUniqueSessionId();
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_SESSION_SESSION_H_
projects/rocprofiler/src/core/session/spm/spm.cpp
+21 -21
View File
@@ -1,7 +1,7 @@
#include "spm.h"
#include "src/core/hsa/hsa_support.h"
#include "src/utils/helper.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include <hsa/hsa.h>
@@ -58,7 +58,7 @@ std::mutex processQueueLock;
// bool* is_data_loss) {
// [[maybe_unused]] hsa_status_t status = HSA_STATUS_SUCCESS;
// #if 0
// status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_spm_set_dest_buffer_fn(
// status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_spm_set_dest_buffer_fn(
// preferred_agent, size_in_bytes, timeout, size_copied, dest, is_data_loss);
// ASSERTM(status == HSA_STATUS_SUCCESS, "ERROR: SPM set buffer failed");
// #endif
@@ -100,9 +100,9 @@ std::mutex processQueueLock;
// uint32_t pidx = preIndex.load(std::memory_order_release);
// if (spm_buffer_params[idx].len == spm_buffer_params[pidx].size) {
// std::cout << "Buffer completely filled with bytes" << spm_buffer_params[idx].len << std::endl;
// fd = fopen("SPM_rocmtool_data.txt", "wb");
// fd = fopen("SPM_rocprofiler_data.txt", "wb");
// size_t retele = fwrite(spm_buffer_params[pidx].addr, 1, spm_buffer_params[idx].len, fd);
// if (retele <= 0) rocmtools::warning("SPM Data is wrong!");
// if (retele <= 0) rocprofiler::warning("SPM Data is wrong!");
// fclose(fd);
// } else {
// std::cout << "Buffer partially filled with %d bytes" << spm_buffer_params[idx].len
@@ -156,10 +156,10 @@ std::mutex processQueueLock;
// uint64_t count = 0;
// std::vector<uint64_t> timestamp_vec;
// // Get Buffer
// rocmtools::Session* session =
// rocmtools::GetROCMToolObj()->GetSession(rocmtools::GetROCMToolObj()->GetCurrentSessionId());
// rocprofiler::Session* session =
// rocprofiler::GetROCProfilerSingleton()->GetSession(rocprofiler::GetROCProfilerSingleton()->GetCurrentSessionId());
// rocprofiler_filter_id_t filter_id = session->GetFilterIdWithKind(ROCPROFILER_SPM_COLLECTION);
// rocmtools::Filter* filter = session->GetFilter(filter_id);
// rocprofiler::Filter* filter = session->GetFilter(filter_id);
// rocprofiler_buffer_id_t buffer_id = filter->GetBufferId();
// Memory::GenericBuffer* buffer = session->GetBuffer(buffer_id);
// // Getting timestamps
@@ -188,7 +188,7 @@ std::mutex processQueueLock;
// }
// se++;
// }
// record.header.id = rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()};
// record.header.id = rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()};
// buffer->AddRecord(record);
// nSample++;
// index += 160;
@@ -197,7 +197,7 @@ std::mutex processQueueLock;
// void spmDataParse() {
// std::vector<uint16_t> lines;
// fd = fopen("SPM_rocmtool_data.txt", "rb");
// fd = fopen("SPM_rocprofiler_data.txt", "rb");
// while (!feof(fd)) {
// char bytes[2];
// size_t size = fread(&bytes, 1, 2, fd);
@@ -214,7 +214,7 @@ hsa_status_t device_cb(hsa_agent_t agent, void* data) {
hsa_device_type_t device_type;
devices_t* devices = reinterpret_cast<devices_t*>(data);
if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type) != HSA_STATUS_SUCCESS)
rocmtools::fatal("hsa_agent_get_info failed");
rocprofiler::fatal("hsa_agent_get_info failed");
switch (device_type) {
case HSA_DEVICE_TYPE_CPU:
devices->cpu_devices.push_back(agent);
@@ -233,16 +233,16 @@ void get_hsa_agents_list(devices_t* device_list) {
hsa_status_t status;
// Enumerate the agents.
status = hsa_iterate_agents(device_cb, device_list);
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_iterate_agents failed");
if (status != HSA_STATUS_SUCCESS) rocprofiler::fatal("hsa_iterate_agents failed");
}
uint64_t submitPacket(hsa_queue_t* queue, const void* packet) {
const uint32_t slot_size_b = CMD_SLOT_SIZE_B;
// advance command queue
const uint64_t write_idx =
rocmtools::hsa_support::GetCoreApiTable().hsa_queue_add_write_index_scacq_screl_fn(queue, 1);
rocprofiler::hsa_support::GetCoreApiTable().hsa_queue_add_write_index_scacq_screl_fn(queue, 1);
while ((write_idx -
rocmtools::hsa_support::GetCoreApiTable().hsa_queue_load_read_index_relaxed_fn(queue)) >=
rocprofiler::hsa_support::GetCoreApiTable().hsa_queue_load_read_index_relaxed_fn(queue)) >=
queue->size) {
sched_yield(); // TODO: remove
}
@@ -261,7 +261,7 @@ uint64_t submitPacket(hsa_queue_t* queue, const void* packet) {
header_atomic_ptr->store(slot_data[0], std::memory_order_release);
// ringdoor bell
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_store_relaxed_fn(queue->doorbell_signal,
rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_store_relaxed_fn(queue->doorbell_signal,
write_idx);
return write_idx;
@@ -271,11 +271,11 @@ uint64_t submitPacket(hsa_queue_t* queue, const void* packet) {
// // create a single-producer queue
// // TODO: check if API args are correct, especially UINT32_MAX
// hsa_status_t status;
// status = rocmtools::hsa_support::GetCoreApiTable().hsa_queue_create_fn(
// status = rocprofiler::hsa_support::GetCoreApiTable().hsa_queue_create_fn(
// gpu_agent, QUEUE_NUM_PACKETS, HSA_QUEUE_TYPE_SINGLE, nullptr, nullptr, UINT32_MAX, UINT32_MAX,
// queue);
// if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("queue creation failed");
// if (status != HSA_STATUS_SUCCESS) rocprofiler::fatal("queue creation failed");
// return (status == HSA_STATUS_SUCCESS);
// }
@@ -288,19 +288,19 @@ hsa_signal_value_t signalWait(const hsa_signal_t& signal, const hsa_signal_value
// Probably a maximum wait time should be set. We don't want application to hang because of
// unlimited wait.
// TODO2 : try 500000 assuming nanosecond granularity -- must be verified.
ret_value = rocmtools::hsa_support::GetCoreApiTable().hsa_signal_wait_scacquire_fn(
ret_value = rocprofiler::hsa_support::GetCoreApiTable().hsa_signal_wait_scacquire_fn(
signal, HSA_SIGNAL_CONDITION_LT, signal_value, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
if (ret_value == exp_value) break;
if (ret_value != signal_value)
rocmtools::fatal("Error: signalWait: signal_value(%lu), ret_value(%lu)", signal_value,
rocprofiler::fatal("Error: signalWait: signal_value(%lu), ret_value(%lu)", signal_value,
ret_value);
}
return ret_value;
}
} // namespace
namespace rocmtools {
namespace rocprofiler {
spm::SpmCounters::SpmCounters(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
@@ -417,9 +417,9 @@ rocprofiler_status_t spm::SpmCounters::stopSpm() {
status = hsa_support::GetCoreApiTable().hsa_queue_destroy_fn(queue_);
queue_ = nullptr;
}
if (status != HSA_STATUS_SUCCESS) rocmtools::warning("Queue destroy failed");
if (status != HSA_STATUS_SUCCESS) rocprofiler::warning("Queue destroy failed");
return ROCPROFILER_STATUS_SUCCESS;
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/spm/spm.h
+2 -2
View File
@@ -13,7 +13,7 @@
#include "rocprofiler.h"
namespace rocmtools {
namespace rocprofiler {
namespace spm {
@@ -52,6 +52,6 @@ class SpmCounters {
bool find_hsa_agent_cpu(uint64_t index, hsa_agent_t* agent);
bool find_hsa_agent_gpu(uint64_t index, hsa_agent_t* agent);
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_CORE_SESSION_SPM_H_
projects/rocprofiler/src/core/session/tracer/src/loader.h
+1 -1
View File
@@ -72,7 +72,7 @@ template <typename Loader> class BaseLoader {
auto function_ptr = reinterpret_cast<FunctionPtr>(::dlsym(handle_, symbol));
if (function_ptr == nullptr)
rocmtools::fatal("symbol lookup '%s' failed: %s", symbol, ::dlerror());
rocprofiler::fatal("symbol lookup '%s' failed: %s", symbol, ::dlerror());
return function_ptr;
}
projects/rocprofiler/src/core/session/tracer/src/roctracer.cpp
+28 -28
View File
@@ -41,7 +41,7 @@
#include "registration_table.h"
#include "src/core/hsa/hsa_support.h"
#include "src/utils/helper.h"
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
static inline uint32_t GetPid() {
static auto pid = syscall(__NR_getpid);
@@ -236,18 +236,18 @@ template <activity_domain_t domain> struct ApiTracer {
};
static void Exit(OperationId operation_id, TraceData* trace_data) {
if (rocmtools::GetROCMToolObj()) {
if (rocprofiler::GetROCProfilerSingleton()) {
if (auto pool = activity_table.Get(operation_id)) {
if (rocmtools::GetROCMToolObj() &&
rocmtools::GetROCMToolObj()->GetSession((*pool)->session_id) &&
rocmtools::GetROCMToolObj()
if (rocprofiler::GetROCProfilerSingleton() &&
rocprofiler::GetROCProfilerSingleton()->GetSession((*pool)->session_id) &&
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)) {
if (rocmtools::GetROCMToolObj()
if (rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->IsValid()) {
std::lock_guard<std::mutex> lock(rocmtools::GetROCMToolObj()
std::lock_guard<std::mutex> lock(rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->GetBufferLock());
@@ -255,7 +255,7 @@ template <activity_domain_t domain> struct ApiTracer {
rocprofiler_record_tracer_t record{};
record.header = rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}};
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}};
record.domain = domain;
record.operation_id = rocprofiler_tracer_operation_id_t{operation_id};
record.correlation_id =
@@ -270,7 +270,7 @@ template <activity_domain_t domain> struct ApiTracer {
rocprofiler_record_tracer_t ext_record{};
ext_record.header = rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}};
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}};
ext_record.domain = ACTIVITY_DOMAIN_EXT_API;
ext_record.operation_id =
rocprofiler_tracer_operation_id_t{ACTIVITY_EXT_OP_EXTERN_ID};
@@ -280,13 +280,13 @@ template <activity_domain_t domain> struct ApiTracer {
ext_record.phase = ROCPROFILER_PHASE_NONE;
// Write the external correlation id record directly followed by the
// activity record.
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(std::array<rocprofiler_record_tracer_t, 2>{ext_record, record});
} else {
// Write record to the buffer.
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(record);
@@ -336,7 +336,7 @@ template <activity_domain_t domain> struct ApiTracer {
if (callback_enabled) {
trace_data->phase_enter = Enter_UserCallback;
trace_data->phase_exit = [](OperationId, TraceData*) {
rocmtools::fatal("should not reach here");
rocprofiler::fatal("should not reach here");
};
}
}
@@ -377,20 +377,20 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
// If the record is for a kernel dispatch, write the kernel name in the pool's data,
// and make the record point to it. Older HIP runtimes do not provide a kernel name,
// so record.kernel_name might be null.
if (!rocmtools::GetROCMToolObj()) return 0;
if (rocmtools::GetROCMToolObj() &&
rocmtools::GetROCMToolObj()->GetSession((*pool)->session_id) &&
rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()) return 0;
if (rocprofiler::GetROCProfilerSingleton() &&
rocprofiler::GetROCProfilerSingleton()->GetSession((*pool)->session_id) &&
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)) {
std::lock_guard<std::mutex> lock(rocmtools::GetROCMToolObj()
std::lock_guard<std::mutex> lock(rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->GetBufferLock());
rocprofiler_record_tracer_t rocprofiler_record{};
rocprofiler_record.header = rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}};
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}};
rocprofiler_record.domain = domain;
rocprofiler_record.external_id = rocprofiler_tracer_external_id_t{};
rocprofiler_record.operation_id = rocprofiler_tracer_operation_id_t{record->kind};
@@ -407,7 +407,7 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
rocprofiler_record.api_data_handle.handle = strdup(record->kernel_name);
void* initial_handle = const_cast<void*>(rocprofiler_record.api_data_handle.handle);
rocprofiler_record.api_data_handle.size = (strlen(record->kernel_name) + 1);
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(rocprofiler_record, rocprofiler_record.api_data_handle.handle,
@@ -421,7 +421,7 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
static_cast<const char*>(data);
});
} else {
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(rocprofiler_record);
@@ -444,20 +444,20 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
case ACTIVITY_DOMAIN_HSA_OPS:
if (auto pool = hsa_ops_activity_table.Get(operation_id)) {
if (auto record = static_cast<activity_record_t*>(data)) {
if (!rocmtools::GetROCMToolObj()) return 0;
if (rocmtools::GetROCMToolObj() &&
rocmtools::GetROCMToolObj()->GetSession((*pool)->session_id) &&
rocmtools::GetROCMToolObj()
if (!rocprofiler::GetROCProfilerSingleton()) return 0;
if (rocprofiler::GetROCProfilerSingleton() &&
rocprofiler::GetROCProfilerSingleton()->GetSession((*pool)->session_id) &&
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)) {
std::lock_guard<std::mutex> lock(rocmtools::GetROCMToolObj()
std::lock_guard<std::mutex> lock(rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->GetBufferLock());
rocprofiler_record_tracer_t rocprofiler_record{};
rocprofiler_record.header = rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}};
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}};
rocprofiler_record.domain = domain;
rocprofiler_record.external_id = rocprofiler_tracer_external_id_t{0};
rocprofiler_record.operation_id = rocprofiler_tracer_operation_id_t{record->op};
@@ -474,7 +474,7 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
rocprofiler_record.api_data_handle.handle = strdup(record->kernel_name);
void* initial_handle = const_cast<void*>(rocprofiler_record.api_data_handle.handle);
rocprofiler_record.api_data_handle.size = strlen(record->kernel_name) + 1;
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(rocprofiler_record, rocprofiler_record.api_data_handle.handle,
@@ -488,7 +488,7 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
static_cast<const char*>(data);
});
} else {
rocmtools::GetROCMToolObj()
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(rocprofiler_record);
projects/rocprofiler/src/core/session/tracer/tracer.cpp
+11 -11
View File
@@ -11,12 +11,12 @@
#include <mutex>
#include <utility>
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/utils/helper.h"
#include "src/core/hsa/hsa_support.h"
#include "src/core/memory/generic_buffer.h"
namespace rocmtools {
namespace rocprofiler {
namespace tracer {
char* GetApiCallFunctionName(rocprofiler_tracer_activity_domain_t domain,
@@ -289,9 +289,9 @@ char* Tracer::GetHIPApiDataInfo(rocprofiler_tracer_hip_api_data_info_t kind,
void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_data, void* args) {
api_callback_data_t* args_data = reinterpret_cast<api_callback_data_t*>(args);
if (args_data && rocmtools::GetROCMToolObj() &&
rocmtools::GetROCMToolObj()->GetSession(args_data->session_id) &&
rocmtools::GetROCMToolObj()->GetSession(args_data->session_id)->GetTracer()) {
if (args_data && rocprofiler::GetROCProfilerSingleton() &&
rocprofiler::GetROCProfilerSingleton()->GetSession(args_data->session_id) &&
rocprofiler::GetROCProfilerSingleton()->GetSession(args_data->session_id)->GetTracer()) {
switch (domain) {
case ACTIVITY_DOMAIN_ROCTX: {
const roctx_api_data_t* data = reinterpret_cast<const roctx_api_data_t*>(callback_data);
@@ -299,7 +299,7 @@ void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_d
rocprofiler_record_tracer_t{
rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}},
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}},
rocprofiler_tracer_external_id_t{0}, ACTIVITY_DOMAIN_ROCTX,
rocprofiler_tracer_operation_id_t{cid},
rocprofiler_tracer_api_data_handle_t{callback_data, sizeof(*data)},
@@ -318,7 +318,7 @@ void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_d
rocprofiler_record_tracer_t{
rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}},
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}},
rocprofiler_tracer_external_id_t{0}, ACTIVITY_DOMAIN_HSA_API,
rocprofiler_tracer_operation_id_t{cid},
rocprofiler_tracer_api_data_handle_t{callback_data, sizeof(*data)},
@@ -332,7 +332,7 @@ void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_d
rocprofiler_record_tracer_t{
rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}},
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}},
rocprofiler_tracer_external_id_t{0}, ACTIVITY_DOMAIN_HSA_API,
rocprofiler_tracer_operation_id_t{cid},
rocprofiler_tracer_api_data_handle_t{callback_data, sizeof(*data)},
@@ -352,7 +352,7 @@ void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_d
rocprofiler_record_tracer_t{
rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}},
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}},
rocprofiler_tracer_external_id_t{0}, ACTIVITY_DOMAIN_HIP_API,
rocprofiler_tracer_operation_id_t{cid},
rocprofiler_tracer_api_data_handle_t{callback_data, sizeof(*data)},
@@ -389,7 +389,7 @@ void api_callback(activity_domain_t domain, uint32_t cid, const void* callback_d
rocprofiler_record_tracer_t{
rocprofiler_record_header_t{
ROCPROFILER_TRACER_RECORD,
rocprofiler_record_id_t{rocmtools::GetROCMToolObj()->GetUniqueRecordId()}},
rocprofiler_record_id_t{rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId()}},
rocprofiler_tracer_external_id_t{0}, ACTIVITY_DOMAIN_HIP_API,
rocprofiler_tracer_operation_id_t{cid},
rocprofiler_tracer_api_data_handle_t{callback_data, sizeof(*data)},
@@ -478,4 +478,4 @@ void Tracer::InitRoctracer(
}
} // namespace tracer
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/core/session/tracer/tracer.h
+2 -2
View File
@@ -37,7 +37,7 @@ typedef struct {
rocprofiler_session_id_t session_id;
} api_callback_data_t;
namespace rocmtools {
namespace rocprofiler {
namespace tracer {
char* GetApiCallFunctionName(rocprofiler_tracer_activity_domain_t domain,
@@ -108,6 +108,6 @@ class Tracer {
};
} // namespace tracer
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_TOOLS_TRACER_TRACER_H_
projects/rocprofiler/src/pcsampler/README.md
+2 -2
View File
@@ -1,10 +1,10 @@
# ROCmtools PC sampling
# ROCProfiler PC sampling
## Known issues
* Abnormal process termination with running GPU kernels while PC sampling is
active may leak resources, eventually necessitating a GPU reset or system
reboot. This is a consequence of the current userspace PC sampling
reboot. This is a consequence of the current user-space PC sampling
implementation, which must coexist with AMD GPU driver activity in the Linux
kernel.
projects/rocprofiler/src/pcsampler/core/pc_sampler.cpp
+6 -6
View File
@@ -27,13 +27,13 @@
#include <hsa/hsa.h>
#include <pciaccess.h>
#include "src/api/rocmtool.h"
#include "src/api/rocprofiler_singleton.h"
#include "src/pcsampler/session/pc_sampler.h"
#include "src/pcsampler/gfxip/gfxip.h"
#include "src/core/hsa/hsa_common.h"
#include "src/core/hsa/hsa_support.h"
namespace rocmtools::pc_sampler {
namespace rocprofiler::pc_sampler {
PCSampler::PCSampler(
rocprofiler_buffer_id_t buffer_id,
@@ -60,7 +60,7 @@ void PCSampler::Start() {
using agents_t = std::vector<hsa_agent_t>;
agents_t agents;
rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
rocprofiler::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
[](hsa_agent_t agent, void *arg){
auto &agents = *reinterpret_cast<agents_t *>(arg);
agents.emplace_back(agent);
@@ -69,7 +69,7 @@ void PCSampler::Start() {
&agents);
for (const auto &agent : agents) {
const auto &ai = rocmtools::hsa_support::GetAgentInfo(agent.handle);
const auto &ai = rocprofiler::hsa_support::GetAgentInfo(agent.handle);
if (ai.getType() != HSA_DEVICE_TYPE_GPU) {
continue;
}
@@ -88,7 +88,7 @@ void PCSampler::Stop() {
}
void PCSampler::AddRecord(rocprofiler_record_pc_sample_t &record) {
const auto tool = rocmtools::GetROCMToolObj();
const auto tool = rocprofiler::GetROCProfilerSingleton();
const auto session = tool->GetSession(session_id_);
const auto buffer = session->GetBuffer(buffer_id_);
@@ -116,4 +116,4 @@ void PCSampler::SamplerLoop() {
}
}
} // namespace rocmtools::pc_sampler
} // namespace rocprofiler::pc_sampler
projects/rocprofiler/src/pcsampler/gfxip/aldebaran_reg_offset_init.cpp
+2 -2
View File
@@ -21,7 +21,7 @@
#include "gfxip.h"
#include "aldebaran_ip_offset.h"
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
void aldebaran_reg_offset_init(device_t& dev) {
for (int i = 0; i < MAX_INSTANCE; ++i) {
@@ -30,4 +30,4 @@ void aldebaran_reg_offset_init(device_t& dev) {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/gfxip/arct_reg_offset_init.cpp
+2 -2
View File
@@ -21,7 +21,7 @@
#include "gfxip.h"
#include "arct_ip_offset.h"
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
void arct_reg_offset_init(device_t& dev) {
for (int i = 0; i < MAX_INSTANCE; ++i) {
@@ -30,4 +30,4 @@ void arct_reg_offset_init(device_t& dev) {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/gfxip/gfxip.cpp
+3 -3
View File
@@ -37,12 +37,12 @@
static const char DEBUG_DRI_PATH[] = "/sys/kernel/debug/dri/";
static const char DEV_PFX[] = "dev=";
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
namespace {
static int find_pci_instance(const std::string &pci_string) {
rocmtools::handle_t<DIR*, util::dir_closer> dir(opendir(DEBUG_DRI_PATH));
rocprofiler::handle_t<DIR*, util::dir_closer> dir(opendir(DEBUG_DRI_PATH));
if (dir.get() == nullptr) {
char *errstr = strerror(errno);
warning("Can't open debugfs dri directory: %s\n", errstr);
@@ -206,4 +206,4 @@ device_t::~device_t() {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/gfxip/gfxip.h
+5 -5
View File
@@ -43,7 +43,7 @@
#define __maybe_unused
namespace rocmtools::pc_sampler {
namespace rocprofiler::pc_sampler {
struct PCSampler;
@@ -59,7 +59,7 @@ struct fd_closer {
void operator()(int fd) { if (fd >= 0) close(fd); }
};
} // namespace rocmtools::pc_sampler::gfxip::util
} // namespace rocprofiler::pc_sampler::gfxip::util
struct amdgpu_debugfs_regs2_iocdata {
__u32 use_srbm, use_grbm, pg_lock;
@@ -114,7 +114,7 @@ struct device_t {
struct debugfs_fds {
debugfs_fds() : mmio2(-1) {}
rocmtools::handle_t<int, util::fd_closer> mmio2;
rocprofiler::handle_t<int, util::fd_closer> mmio2;
} fd_;
};
@@ -132,7 +132,7 @@ void aldebaran_reg_offset_init(device_t& dev);
void read_pc_samples_v9(const device_t& dev, PCSampler *sampler);
void read_pc_samples_v9_ioctl(const device_t& dev, PCSampler *sampler);
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
} // namespace rocmtools::pc_sampler
} // namespace rocprofiler::pc_sampler
#endif // SRC_PCSAMPLER_GFXIP_GFXIP_H_
projects/rocprofiler/src/pcsampler/gfxip/gfxip_v9.cpp
+6 -6
View File
@@ -41,7 +41,7 @@
#include "osssys/osssys_4_2_0_offset.h"
#include "osssys/osssys_4_2_0_sh_mask.h"
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
namespace {
@@ -78,7 +78,7 @@ void fill_record(
* XXX: Use of the reserved2 field in the HSA dispatch packet to uniquely
* identify kernel dispatches for PC sampling is an internal implementation
* detail which is subject to change. See the comment associated with
* rocmtool::rocmtool::kernel_dispatch_counter_.
* rocprofiler::rocprofiler::kernel_dispatch_counter_.
*/
record->pc_sample.dispatch_id =
rocprofiler_kernel_dispatch_id_t{pkt->reserved2};
@@ -107,13 +107,13 @@ void fill_record(
/*
* XXX FIXME: For consistency, this is the same method as used by
* rocmtools::queue::AsyncSignalHandler in queue.cpp to fill
* rocprofiler::queue::AsyncSignalHandler in queue.cpp to fill
* rocprofiler_record_profiler_t::gpu_id, but it should be changed; see the
* comment in rocmtools::hsa_support::Initialize about using KFD's gpu_id for
* comment in rocprofiler::hsa_support::Initialize about using KFD's gpu_id for
* more information.
*/
record->pc_sample.gpu_id = rocprofiler_agent_id_t{
(uint64_t)rocmtools::hsa_support::GetAgentInfo(hdl).getIndex()};
(uint64_t)rocprofiler::hsa_support::GetAgentInfo(hdl).getIndex()};
}
} // namespace
@@ -296,4 +296,4 @@ void read_pc_samples_v9_ioctl(const device_t& dev, PCSampler *sampler) {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/gfxip/vega10_reg_offset_init.cpp
+2 -2
View File
@@ -21,7 +21,7 @@
#include "gfxip.h"
#include "vega10_ip_offset.h"
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
void vega10_reg_offset_init(device_t& dev) {
for (int i = 0; i < MAX_INSTANCE; ++i) {
@@ -30,4 +30,4 @@ void vega10_reg_offset_init(device_t& dev) {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/gfxip/vega20_reg_offset_init.cpp
+2 -2
View File
@@ -21,7 +21,7 @@
#include "gfxip.h"
#include "vega20_ip_offset.h"
namespace rocmtools::pc_sampler::gfxip {
namespace rocprofiler::pc_sampler::gfxip {
void vega20_reg_offset_init(device_t& dev) {
for (int i = 0; i < MAX_INSTANCE; ++i) {
@@ -30,4 +30,4 @@ void vega20_reg_offset_init(device_t& dev) {
}
}
} // namespace rocmtools::pc_sampler::gfxip
} // namespace rocprofiler::pc_sampler::gfxip
projects/rocprofiler/src/pcsampler/session/pc_sampler.h
+2 -2
View File
@@ -29,7 +29,7 @@
#include "rocprofiler.h"
#include "src/pcsampler/gfxip/gfxip.h"
namespace rocmtools::pc_sampler {
namespace rocprofiler::pc_sampler {
struct PCSampler {
PCSampler(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
@@ -55,6 +55,6 @@ private:
std::thread sampler_thread_;
};
} // namespace rocmtools::pc_sampler
} // namespace rocprofiler::pc_sampler
#endif // SRC_PCSAMPLER_PC_SAMPLER_H_
projects/rocprofiler/src/tools/CMakeLists.txt
+1 -1
View File
@@ -1,5 +1,5 @@
# ############################################################################################################################################
# ROCMTOOL Tool
# ROCProfiler Tool
# ############################################################################################################################################
# Setting Default Binary output directory
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
projects/rocprofiler/src/tools/amdsys/CMakeLists.txt
+1 -1
View File
@@ -1,5 +1,5 @@
# ############################################################################################################################################
# ROCMTOOL AMDSYS Frontend
# ROCProfiler AMDSYS Frontend
# ############################################################################################################################################
set(CMAKE_BINARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
projects/rocprofiler/src/tools/tool.cpp
+1 -1
View File
@@ -65,7 +65,7 @@ namespace fs = std::experimental::filesystem;
#define CHECK_ROCPROFILER(call) \
do { \
if ((call) != ROCPROFILER_STATUS_SUCCESS) \
rocmtools::fatal("Error: ROCProfiler API Call Error!"); \
rocprofiler::fatal("Error: ROCProfiler API Call Error!"); \
} while (false)
TRACE_BUFFER_INSTANTIATE();
namespace {
projects/rocprofiler/src/utils/access_control.h
+16 -16
View File
@@ -28,16 +28,16 @@
// TODO(aelwazir): To be implemented
// type for reference counter
typedef std::atomic<uint64_t> RocmToolsRefCount;
typedef std::atomic<uint64_t> ROCProfilerRefCount;
// type for access/release control flag
typedef std::atomic<bool> RocmToolsARControl;
typedef std::atomic<bool> ROCProfilerARControl;
namespace rocmtools {
namespace rocprofiler {
// the access/release control flag can be shared or exclusive
static RocmToolsARControl kARShared = ATOMIC_VAR_INIT(0);
static RocmToolsARControl kARExclusive = ATOMIC_VAR_INIT(1);
static ROCProfilerARControl kARShared = ATOMIC_VAR_INIT(0);
static ROCProfilerARControl kARExclusive = ATOMIC_VAR_INIT(1);
/*-------------------------------------------------------------------------*/
/**
@@ -49,25 +49,25 @@ static RocmToolsARControl kARExclusive = ATOMIC_VAR_INIT(1);
class AccessControl {
public:
//!< constructor that defaults the RefCount and ARControl
AccessControl(RocmToolsRefCount rrc = ATOMIC_VAR_INIT(0),
RocmToolsARControl rac = ATOMIC_VAR_INIT(kARShared.load()));
AccessControl(ROCProfilerRefCount rrc = ATOMIC_VAR_INIT(0),
ROCProfilerARControl rac = ATOMIC_VAR_INIT(kARShared.load()));
//!< naive destructor
~AccessControl();
bool Acquire(RocmToolsARControl access_mode);
void Release(RocmToolsARControl access_mode);
bool Acquire(ROCProfilerARControl access_mode);
void Release(ROCProfilerARControl access_mode);
private:
RocmToolsRefCount rcount_; //!< refernce count, every successful access
ROCProfilerRefCount rcount_; //!< refernce count, every successful access
//! increments it
RocmToolsARControl arcontrol_; //!< shared access or exclusive?
ROCProfilerARControl arcontrol_; //!< shared access or exclusive?
std::mutex accessmutex_;
};
//<!
// AccessControl::AccessControl(RocmToolsRefCount rrc, RocmToolsARControl rac)
// AccessControl::AccessControl(ROCProfilerRefCount rrc, ROCProfilerARControl rac)
// {
// rcount_ = rrc.load();
// arcontrol_ = rac.load();
@@ -76,7 +76,7 @@ class AccessControl {
//<!
AccessControl::~AccessControl() {
// we cannot decrement reference count here, we can only hope it is 0
#ifdef ROCMTOOL_DEBUG_PRINT
#ifdef ROCPROFILER_DEBUG_PRINT
int zero = 0;
if (rcount_.compare_exchange_strong(zero&, 0, std::memory_order_acq_rel))
cout << endl
@@ -90,11 +90,11 @@ AccessControl::~AccessControl() {
#if 0
//!< three steps, 1) get scoped lock, 2) check if it is shared to shared
//!< acquire, and 3) if it is exclusive, check if refcount is 0
bool AccessControl::Acquire(RocmToolsARControl rac) {
bool AccessControl::Acquire(ROCProfilerARControl rac) {
// first, ensure we have a lock_guard
std::lock_guard<std::mutex> lg(accessmutex_);
// we should always expect the current value to be shared
RocmToolsARControl rac_expected = kARShared.load();
ROCProfilerARControl rac_expected = kARShared.load();
// if what we have is what we expect then we are done
if (rac.load(std::memory_order_relaxed) == rac_expected.load()) {
// rcount_.store(std::memory_order_acq_rel,
@@ -111,6 +111,6 @@ bool AccessControl::Acquire(RocmToolsARControl rac) {
//!< always exchange to kARShared and reduce refcount
// void AccessControl::Release(
} // namespace rocmtools
} // namespace rocprofiler
#endif // SRC_UTILS_ACCESS_CONTROL_H_
projects/rocprofiler/src/utils/exception.h
+3 -3
View File
@@ -28,8 +28,8 @@
#include "helper.h"
#include "rocprofiler.h"
// TODO(aelwazir): namespace rocmtool
namespace rocmtools {
// TODO(aelwazir): namespace rocprofiler
namespace rocprofiler {
class Exception : public std::runtime_error {
public:
@@ -46,7 +46,7 @@ class Exception : public std::runtime_error {
const rocprofiler_status_t status_;
};
}; // namespace rocmtool
}; // namespace rocprofiler
// TODO(aelwazir): throw instead of the macros
projects/rocprofiler/src/utils/handle.h
+2 -2
View File
@@ -25,7 +25,7 @@
#include <utility>
#include <cstddef>
namespace rocmtools {
namespace rocprofiler {
template <typename T, typename D>
class handle_t {
@@ -109,6 +109,6 @@ public:
}
};
} // namespace rocmtools
} // namespace rocprofiler
#endif // UTILS_HANDLE_H_
projects/rocprofiler/src/utils/helper.cpp
+4 -4
View File
@@ -53,7 +53,7 @@ struct BackTraceInfo {
void errorCallback(void* data, const char* message, int errnum) {
BackTraceInfo* info = static_cast<BackTraceInfo*>(data);
info->sstream << "ROCMTools: error: " << message << '(' << errnum << ')';
info->sstream << "ROCProfiler: error: " << message << '(' << errnum << ')';
info->error = 1;
}
@@ -95,7 +95,7 @@ int fullCallback(void* data, uintptr_t pc, const char* filename, int lineno, con
} // namespace
#endif // defined (ENABLE_BACKTRACE)
namespace rocmtools {
namespace rocprofiler {
std::string string_vprintf(const char* format, va_list va) {
va_list copy;
@@ -142,7 +142,7 @@ std::string string_printf(const char* format, ...) {
message += info.sstream.str();
#endif /* defined (ENABLE_BACKTRACE) */
std::string errmsg("ROCMTools: fatal error: " + message);
std::string errmsg("ROCProfiler: fatal error: " + message);
fputs(errmsg.c_str(), stderr);
std::cerr << errmsg << std::endl;
@@ -236,4 +236,4 @@ std::string left_trim(const std::string& s) {
}
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/utils/helper.h
+2 -2
View File
@@ -30,7 +30,7 @@
#include <cxxabi.h>
#include "exception.h"
namespace rocmtools {
namespace rocprofiler {
std::string string_vprintf(const char* format, va_list va);
@@ -69,4 +69,4 @@ bool has_counter_format(std::string const& str);
// trims the begining of the line for spaces
std::string left_trim(const std::string& s);
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/src/utils/logger.h
+11 -11
View File
@@ -40,7 +40,7 @@
#include <string>
#include <utility>
namespace rocmtools {
namespace rocprofiler {
class Logger {
public:
@@ -140,32 +140,32 @@ class Logger {
std::map<uint32_t, std::string> message_;
};
} // namespace rocmtools
} // namespace rocprofiler
#define FATAL_LOGGING(stream) \
do { \
rocmtools::Logger::Instance() << "fatal: " << rocmtools::Logger::begm \
<< stream << rocmtools::Logger::endl; \
rocprofiler::Logger::Instance() << "fatal: " << rocprofiler::Logger::begm \
<< stream << rocprofiler::Logger::endl; \
throw(ROCPROFILER_STATUS_ERROR, "Error: " << stream); \
} while (false)
#define ERR_LOGGING(stream) \
do { \
rocmtools::Logger::Instance() << "error: " << rocmtools::Logger::begm \
<< stream << rocmtools::Logger::endl; \
rocprofiler::Logger::Instance() << "error: " << rocprofiler::Logger::begm \
<< stream << rocprofiler::Logger::endl; \
} while (false)
#define INFO_LOGGING(stream) \
do { \
rocmtools::Logger::Instance() << "info: " << rocmtools::Logger::begm \
<< stream << rocmtools::Logger::endl; \
rocprofiler::Logger::Instance() << "info: " << rocprofiler::Logger::begm \
<< stream << rocprofiler::Logger::endl; \
} while (false)
#define WARN_LOGGING(stream) \
do { \
std::cerr << "ROCmTools: " << stream << std::endl; \
rocmtools::Logger::Instance() << "warning: " << rocmtools::Logger::begm \
<< stream << rocmtools::Logger::endl; \
std::cerr << "ROCProfiler: " << stream << std::endl; \
rocprofiler::Logger::Instance() << "warning: " << rocprofiler::Logger::begm \
<< stream << rocprofiler::Logger::endl; \
} while (false)
#endif // SRC_UTILS_LOGGER_H_
projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt
+1 -1
View File
@@ -3,7 +3,7 @@
844434431085362
ROCMTools: Collecting the following counters:
ROCProfiler: Collecting the following counters:
- GRBM_COUNT
Enabling Counter Collection
projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/openmp_helloworld_golden_traces.txt
+1 -1
View File
@@ -1,4 +1,4 @@
ROCMTools: Collecting the following counters:
ROCProfiler: Collecting the following counters:
- GRBM_COUNT
PASSED!
dispatch[2], gpu-id(0), kernel-name("hip_helloworld"), time(7853273641921013,7853273641924568)
projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.cpp
+4 -4
View File
@@ -885,7 +885,7 @@ TEST_F(ProfilerAPITest, WhenRunningMultipleThreadsProfilerAPIsWorkFine) {
// create as many threads as number of cores in system
std::vector<std::thread> threads(num_cpu_cores);
// inititalize profiler by creating rocmtool object
// initialize profiler by creating rocprofiler object
CheckApi(rocprofiler_initialize());
// Counter Collection with timestamps
@@ -928,7 +928,7 @@ TEST_F(ProfilerAPITest, WhenRunningMultipleThreadsProfilerAPIsWorkFine) {
// destroy session
CheckApi(rocprofiler_destroy_session(session_id));
// finalize profiler by destroying rocmtool object
// finalize profiler by destroying rocprofiler object
CheckApi(rocprofiler_finalize());
}
@@ -1061,7 +1061,7 @@ class ProfilerSPMTest : public ::testing::Test {
};
TEST_F(ProfilerSPMTest, WhenRunningSPMItCollectsSPMData) {
// inititalize rocmtools
// initialize rocprofiler
hsa_init();
CheckApi(rocprofiler_initialize());
@@ -1106,7 +1106,7 @@ TEST_F(ProfilerSPMTest, WhenRunningSPMItCollectsSPMData) {
// destroy session
CheckApi(rocprofiler_destroy_session(session_id));
// finalize spm tracing by destroying rocmtool object
// finalize spm tracing by destroying rocprofiler object
CheckApi(rocprofiler_finalize());
hsa_shut_down();
}
projects/rocprofiler/tests/featuretests/utils/csv_parser.cpp
+2 -2
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include "csv_parser.h"
namespace rocmtools {
namespace rocprofiler {
namespace tests {
namespace utility {
@@ -167,4 +167,4 @@ void CSVParser::ParseCSV(std::string buffer) {
}
} // namespace utility
} // namespace tests
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/tests/featuretests/utils/csv_parser.h
+4 -4
View File
@@ -39,7 +39,7 @@ THE SOFTWARE.
* can use this to parse different applications.
*/
/* --------------------------------------------------------------------------*/
namespace rocmtools {
namespace rocprofiler {
namespace tests {
namespace utility {
using countermap = std::map<uint32_t, std::map<uint32_t, std::string>>;
@@ -111,8 +111,8 @@ class CSVParser {
};
} // namespace utility
} // namespace tests
} // namespace rocmtools
} // namespace rocprofiler
using rocmtools::tests::utility::countermap;
using rocmtools::tests::utility::CSVParser;
using rocprofiler::tests::utility::countermap;
using rocprofiler::tests::utility::CSVParser;
#endif // TESTS_FEATURETESTS_PROFILER_UTILS_CSV_PARSER_H_
projects/rocprofiler/tests/featuretests/utils/test_utils.cpp
+2 -2
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#include "test_utils.h"
namespace rocmtools {
namespace rocprofiler {
namespace tests {
namespace utility {
@@ -84,4 +84,4 @@ bool is_installed_path() {
} // namespace utility
} // namespace tests
} // namespace rocmtools
} // namespace rocprofiler
projects/rocprofiler/tests/featuretests/utils/test_utils.h
+5 -5
View File
@@ -33,7 +33,7 @@ THE SOFTWARE.
#include <iterator>
#include <string>
namespace rocmtools {
namespace rocprofiler {
namespace tests {
namespace utility {
@@ -47,14 +47,14 @@ bool is_installed_path();
} // namespace utility
} // namespace tests
} // namespace rocmtools
} // namespace rocprofiler
// used for dl_addr to locate the running
// path for executable
int main(int argc, char** argv);
using rocmtools::tests::utility::GetNumberOfCores;
using rocmtools::tests::utility::GetRunningPath;
using rocmtools::tests::utility::is_installed_path;
using rocprofiler::tests::utility::GetNumberOfCores;
using rocprofiler::tests::utility::GetRunningPath;
using rocprofiler::tests::utility::is_installed_path;
#endif // TESTS_FEATURETESTS_PROFILER_UTILS_TEST_UTILS_H_
projects/rocprofiler/tests/memorytests/run_asan_tests.sh
+1 -1
View File
@@ -1,7 +1,7 @@
#!/bin/bash
CURRENT_DIR="$( dirname -- "$0"; )";
export PATH=$rocmtoolsdir:$PATH
export PATH=$rocprofilerdir:$PATH
echo -e "Running Memory Leaks Check From ${CURRENT_DIR}"
LD_PRELOAD=/usr/lib/x86_64-linux-gnu/libasan.so.6 ASAN_OPTIONS=detect_leaks=1 LSAN_OPTIONS=suppressions=$CURRENT_DIR/suppr.txt ${CURRENT_DIR}/../../rocprofv2 -i $CURRENT_DIR/input.txt $1 > /tmp/rocprofv2-temp 2> $2
projects/rocprofiler/tests/unittests/CMakeLists.txt
+2 -2
View File
@@ -33,7 +33,7 @@ file(GLOB ROCPROFILER_SRC_PROFILER_FILES ${PROJECT_SOURCE_DIR}/src/core/session/
file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp)
file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp)
file(GLOB ROCPROFILER_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp)
file(GLOB ROCPROFILER_SRC_CLASS_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocmtool.cpp)
file(GLOB ROCPROFILER_SRC_CLASS_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp)
file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp)
file(GLOB ROCPROFILER_SRC_API_FILES ${PROJECT_SOURCE_DIR}/src/api/*.cpp)
@@ -49,7 +49,7 @@ set(CORE_PC_SAMPLING_DIR ${PROJECT_SOURCE_DIR}/src/pcsampler)
file(GLOB CORE_PC_SAMPLING_FILES ${CORE_PC_SAMPLING_DIR}/core/*.cpp ${CORE_PC_SAMPLING_DIR}/gfxip/*.cpp ${CORE_PC_SAMPLING_DIR}/session/*.cpp)
# Compiling gtests
file(GLOB ROCPROFILER_TOOL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/rocmtools/tools/tool.cpp)
file(GLOB ROCPROFILER_TOOL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/tools/tool.cpp)
file(GLOB CORE_COUNTERS_PARENT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/counters/*.cpp)
file(GLOB CORE_COUNTERS_METRICS_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/counters/metrics/*.cpp)
projects/rocprofiler/tests/unittests/profiler_gtest.cpp
+17 -17
View File
@@ -24,7 +24,7 @@
#include <mutex>
#include <memory>
#include "api/rocmtool.h"
#include "api/rocprofiler_singleton.h"
#include "core/memory/generic_buffer.h"
#include "core/hardware/hsa_info.h"
#include "core/session/session.h"
@@ -101,8 +101,8 @@ TEST(WhenAddingARecordToBuffer, DISABLED_RecordGetsAddedSuccefully) {
class TimeStampSession : public ::testing::Test {
protected:
rocprofiler_session_id_t session_id{1234};
std::unique_ptr<rocmtools::Session> session_ptr =
std::make_unique<rocmtools::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id);
std::unique_ptr<rocprofiler::Session> session_ptr =
std::make_unique<rocprofiler::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id);
void SetUp() {
rocprofiler_filter_id_t filter_id =
@@ -159,11 +159,11 @@ TEST_F(TimeStampSession, ForANewlyCreatedSessionValidSessionIdIsReturned) {
class TestingMultipleSessions : public ::testing::Test {
protected:
std::vector<std::unique_ptr<rocmtools::Session>> session_list;
std::vector<std::unique_ptr<rocprofiler::Session>> session_list;
uint64_t number_of_sessions = 5;
void SetUp() {
for (uint64_t id = 0; id < number_of_sessions; id++) {
std::unique_ptr<rocmtools::Session> timestamp_session = std::make_unique<rocmtools::Session>(
std::unique_ptr<rocprofiler::Session> timestamp_session = std::make_unique<rocprofiler::Session>(
ROCPROFILER_NONE_REPLAY_MODE, rocprofiler_session_id_t{id});
rocprofiler_filter_id_t filter_id = timestamp_session->CreateFilter(
@@ -225,16 +225,16 @@ TEST_F(TestingMultipleSessions, DeactivatingAnActivatedSessionPasses) {
/*
* ###############################################
* ################TESTING ROCMTOOLS##############
* ################TESTING ROCProfiler##############
* ###############################################
*/
// Createing sessions with 2 different profiling mode
TEST(WhenCreatingTwoSessionsWithDiffProfilingMode, BothSessionsAreCreated) {
std::map<uint64_t, std::unique_ptr<rocmtools::Session>> sessions;
std::map<uint64_t, std::unique_ptr<rocprofiler::Session>> sessions;
sessions = std::map<uint64_t, std::unique_ptr<rocmtools::Session>>();
sessions = std::map<uint64_t, std::unique_ptr<rocprofiler::Session>>();
{
// create a counter collection session
@@ -244,7 +244,7 @@ TEST(WhenCreatingTwoSessionsWithDiffProfilingMode, BothSessionsAreCreated) {
counters.emplace_back("GRBM_COUNT");
sessions.emplace(
session_id.handle,
std::make_unique<rocmtools::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id));
std::make_unique<rocprofiler::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id));
rocprofiler_filter_id_t filter_id =
sessions.at(session_id.handle)
@@ -260,7 +260,7 @@ TEST(WhenCreatingTwoSessionsWithDiffProfilingMode, BothSessionsAreCreated) {
rocprofiler_session_id_t session_id{2};
sessions.emplace(
session_id.handle,
std::make_unique<rocmtools::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id));
std::make_unique<rocprofiler::Session>(ROCPROFILER_NONE_REPLAY_MODE, session_id));
rocprofiler_filter_id_t filter_id =
sessions.at(session_id.handle)
->CreateFilter(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION, rocprofiler_filter_data_t{},
@@ -284,7 +284,7 @@ void (*callback_fun)(const rocprofiler_record_header_t* begin,
TEST(WhenTestingCounterCollectionMode, TestSucceeds) {
rocprofiler_session_id_t session_id;
rocmtools::rocmtool toolobj;
rocprofiler::ROCProfiler_Singleton toolobj;
session_id = toolobj.CreateSession(ROCPROFILER_NONE_REPLAY_MODE);
rocprofiler_filter_id_t filter_id =
toolobj.GetSession(session_id)
@@ -295,7 +295,7 @@ TEST(WhenTestingCounterCollectionMode, TestSucceeds) {
toolobj.GetSession(session_id)->GetFilter(filter_id)->SetBufferId(buffer_id);
rocmtools::Session* session = toolobj.GetSession(session_id);
rocprofiler::Session* session = toolobj.GetSession(session_id);
EXPECT_TRUE(session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION));
toolobj.DestroySession(session_id);
}
@@ -303,7 +303,7 @@ TEST(WhenTestingCounterCollectionMode, TestSucceeds) {
TEST(WhenTestingTimeStampCollectionMode, TestSucceeds) {
rocprofiler_session_id_t session_id;
rocmtools::rocmtool toolobj;
rocprofiler::ROCProfiler_Singleton toolobj;
session_id = toolobj.CreateSession(ROCPROFILER_NONE_REPLAY_MODE);
rocprofiler_filter_id_t filter_id =
toolobj.GetSession(session_id)
@@ -314,7 +314,7 @@ TEST(WhenTestingTimeStampCollectionMode, TestSucceeds) {
toolobj.GetSession(session_id)->GetFilter(filter_id)->SetBufferId(buffer_id);
rocmtools::Session* session = toolobj.GetSession(session_id);
rocprofiler::Session* session = toolobj.GetSession(session_id);
EXPECT_TRUE(session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION));
toolobj.DestroySession(session_id);
@@ -325,7 +325,7 @@ TEST(WhenTestingApplicationReplayMode, TestSucceeds) {
counters.emplace_back("SQ_WAVES");
rocprofiler_session_id_t session_id;
rocmtools::rocmtool toolobj;
rocprofiler::ROCProfiler_Singleton toolobj;
session_id = toolobj.CreateSession(ROCPROFILER_APPLICATION_REPLAY_MODE);
rocprofiler_filter_id_t filter_id =
@@ -337,7 +337,7 @@ TEST(WhenTestingApplicationReplayMode, TestSucceeds) {
toolobj.GetSession(session_id)->CreateBuffer(callback_fun, 0x8000);
toolobj.GetSession(session_id)->GetFilter(filter_id)->SetBufferId(buffer_id);
rocmtools::Session* session = toolobj.GetSession(session_id);
rocprofiler::Session* session = toolobj.GetSession(session_id);
EXPECT_TRUE(session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION));
toolobj.DestroySession(session_id);
@@ -350,7 +350,7 @@ TEST(WhenTrucatingLongKernelNames, KernelNameGetsTruncatedProperly) {
"long, long long, long long, long long, float, float, float, float "
"const*, float*, float const*, float*, float const*) [clone .kd]";
std::string trunkated_name = rocmtools::truncate_name(long_kernel_name);
std::string trunkated_name = rocprofiler::truncate_name(long_kernel_name);
EXPECT_EQ("kernel_7r_3d_pml", trunkated_name);
}