AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Pull from Github

Bladeren bron 
Squashed commit of the following:

commit f029195705a15700380c6f832ba5d15d46fd6de7
Author: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
Date:   Thu Jul 13 14:38:56 2023 -0500

    Formatting workflows for source (clang-format) and cmake (cmake-format) (#4)

    * Add .cmake-format.yaml file

    * Add formatting workflow

    * provide base input for creating PR

    * Update scheme for extracting branch name

    - disable running formatting on push to amd-staging branch

    * patch .cmake-format.yaml for find_package signature

    - apparently cmake-format doesn't format the full signature of find_package

    * run formatting (clang-format v11) (#7)

    Co-authored-by: jrmadsen <jrmadsen@users.noreply.github.com>

    * run cmake formatting (cmake-format) (#6)

    Co-authored-by: jrmadsen <jrmadsen@users.noreply.github.com>

    ---------

    Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

commit bc4d135fdd8a1a9e51235f18a5d575fd2b3735e6
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Thu Jul 13 12:55:17 2023 -0500

    Removing Build cache for potential issues with auto-generated header files (#5)

    Change-Id: I9e2319f4335e2f88585ffa6fac2bd88a1c952e6e

commit ce86dea6a311d44d880fa684eb78f3329295e2a4
Author: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
Date:   Thu Jul 13 11:08:58 2023 -0500

    Fix decltype(<hsa-function>) function pointer usage (#3)

    - the following is done in several places:
        decltype(hsa_memory_allocate)* hsa_memory_allocate
    - above can cause compiler errors
    - replace decltype(<hsa-function>) with decltype(::<hsa-function>)
      - this ensures that the type within the decltype is recognized as the global scope HSA function, not the variable
    - in many places, the variable has a "_fn" suffix to prevent this issue but added '::' anyway for consistency

commit ac49fdd92a72e9c99394253a02da413a6c2e3b3a
Merge: a07946a 03a0855
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Wed Jul 12 11:36:24 2023 -0500

    Merge pull request #2 from ROCm-Developer-Tools/gerrit-amd-staging

    Pull from gerrit

commit 03a085588cffe863e8f466de67be1cfb205b675a
Merge: c26b32b a07946a
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Wed Jul 12 10:57:30 2023 -0500

    Merge branch 'amd-staging' into gerrit-amd-staging

commit a07946a5cd4c670c83c27ad1a076a9d4567ce6d7
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 15:46:04 2023 +0000

    Enabling Cached Builds

commit 525e494a7f13941077a8fd4ad6840904db4d27d4
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 04:53:54 2023 +0000

    Updating missed GPU Targets

commit 42c75862f628c9bee7cfb7dc04dff2619430efbc
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 04:43:02 2023 +0000

    Adding V1 Testing

commit 9d72fd4aee85e4b0c12e717060d2730fa5b73be1
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 03:34:31 2023 +0000

    Fixing Artifacts directory path

commit f4000cc558b3b2e4676f7994f7ce8c8e6f94518e
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 03:27:26 2023 +0000

    Fixing CMake for test build job

commit 2ce8115d4c33948c3c8f957f545a95a04e1d6cd2
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 03:16:18 2023 +0000

    Fixing Ubuntu CMake for ubuntu test build

commit 6d0ed439191be900748d0c025157f9d689a73ec7
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 01:28:41 2023 +0000

    Removing Navi21

commit e349a7642e5ae5eb03ab9fcd0a0f74f09f78cab5
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 01:14:14 2023 +0000

    Removing Navi21

commit fefd02fe68d2a4bca7ec2e381960ad004ee9fc5b
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 00:42:48 2023 +0000

    Fixing CMake Job

commit 2ea46abf7bf92643efa8c549fa70346ffbd79d65
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 00:35:13 2023 +0000

    Fixing CMake Job

commit d99d681ed1999c5fcf291dc678b11a77205fb0f3
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Wed Jul 12 00:32:13 2023 +0000

    Fixing Pull Latest Dockers and CMake Jobs

commit dfc4498072d13b4a1df3a63047d34c682c3d9a29
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Tue Jul 11 23:54:21 2023 +0000

    Fixing CMake job

commit 919efe04de707f7c702031be15c3e2c5f8442cbb
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Tue Jul 11 23:52:13 2023 +0000

    Adding Pull Last dockers job

commit be1b1256e8b0e05308e8f7e7e69bee3acca55281
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Tue Jul 11 18:25:40 2023 -0500

    Update cmake.yml

commit 212299fa4355ae6ec18f9aaacbb79c51ea6c6f97
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Tue Jul 11 18:23:35 2023 -0500

    Update cmake.yml

commit 7c2c1327086a61466cc6cac39f70865c051a8bc7
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Tue Jul 11 18:18:53 2023 -0500

    Update cmake.yml

commit 191b5ce007e612e814c1d7a3afb4ad398f3852e1
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Tue Jul 11 16:03:22 2023 -0500

    Update cmake.yml

commit 8824113d95f3e13c7ce4d0af8e0d9d8f522a6c4a
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Tue Jul 11 16:28:09 2023 +0000

    Fixing Pull from Gerrit job name

    Change-Id: I9e7ed9a27a13ca49d62c93bdadb30f0057e4d385

commit cc3d5e4b02ffb439e8cc2b3efa53527c376f9982
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Tue Jul 11 16:21:43 2023 +0000

    Adding Staging sync job

    Change-Id: I0551f43878b0678ce4b3e74e27d62357cf95ad95

commit b9be2eee71380a2e6dd34d520e92d0c4209277a0
Author: Ammar ELWazir <Ammar.ELWazir@amd.com>
Date:   Tue Jul 11 15:57:11 2023 +0000

    Fixing build.sh

    Change-Id: Ia987b0244f0875370d5fe69907b3f5e9cea914de

commit 9eee33a95a1abd656a7ac5ca10a9f245e9825431
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 21:39:46 2023 -0500

    Update cmake.yml

commit 7093b85a78497140e8b52632ca2a002bdaeacd62
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 21:33:29 2023 -0500

    Update cmake.yml

commit f54697172c72a67740f9fdfa0c217b6ea6931576
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 21:01:26 2023 -0500

    Update cmake.yml

commit 1b6620e16f8940386b0f4f04e69e2410d21c0e26
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 20:21:02 2023 -0500

    Update cmake.yml

commit a94bec740c6b42c4b79c87bca20fa87b99bf060d
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:46:35 2023 -0500

    Update cmake.yml

commit 85d6b29d4375a69d575c18ece8542c50f2ddfcc3
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:34:39 2023 -0500

    Update cmake.yml

commit 8c004887cf1435f1a6214c3d2455299a8a27bd4c
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:31:17 2023 -0500

    Update cmake.yml

commit a14a9168e17d9348a53c6e9c9a47ba1edb4c4509
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:25:46 2023 -0500

    Update cmake.yml

commit 000f2f40b84e6a2f7d4becdbf5aed01436ca4c83
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:08:18 2023 -0500

    Update cmake.yml

commit a28a53d56731cad848fa9133d1c4dbaa8fc7afa7
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 19:03:39 2023 -0500

    Update cmake.yml

commit a6a2db01027f0b01fdfbb5997ddb772c7f51b649
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 18:21:53 2023 -0500

    Update cmake.yml

commit 118ef2a88b2d44e3207c31c343da3e5e5ec6f176
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 17:55:57 2023 -0500

    Update cmake.yml

commit 03c4c232396440cd0be6d2dd7baf4ceea1c2589d
Author: Ammar ELWazir <aelwazir@amd.com>
Date:   Mon Jul 10 17:48:49 2023 -0500

    Create cmake.yml

Change-Id: I77992f15694e77cbae49c56f9ff02f4f9079235d


[ROCm/rocprofiler commit: d4a33cf33a]
This commit is contained in:
Ammar ELWazir
2023-07-13 19:48:38 +00:00
gecommit door Ammar Elwazir
bovenliggende e599708211
commit 6eb06cf201
144 gewijzigde bestanden met toevoegingen van 154836 en 166916 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/rocprofiler/src/core/activity.cpp
+20 -21
Bestand weergeven
@@ -61,11 +61,11 @@ void check_status(hsa_status_t status) {
namespace activity_prim {
// PC sampling callback data
struct pcsmp_callback_data_t {
const char* kernel_name; // sampled kernel name
void* data_buffer; // host buffer for tracing data
uint64_t id; // sample id
uint64_t cycle; // sample cycle
uint64_t pc; // sample PC
const char* kernel_name; // sampled kernel name
void* data_buffer; // host buffer for tracing data
uint64_t id; // sample id
uint64_t cycle; // sample cycle
uint64_t pc; // sample PC
};
uint32_t activity_op = UINT32_MAX;
@@ -74,9 +74,8 @@ std::atomic<activity_async_callback_t> activity_callback{NULL};
rocprofiler_t* context = NULL;
hsa_status_t trace_data_cb(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* data) {
const pcsmp_callback_data_t* pcsmp_data = (pcsmp_callback_data_t*) data;
hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) {
const pcsmp_callback_data_t* pcsmp_data = (pcsmp_callback_data_t*)data;
activity_record_t record{};
record.op = activity_op;
@@ -96,11 +95,13 @@ bool context_handler(rocprofiler_group_t group, void* arg) {
hsa_agent_t agent{};
hsa_status_t status = rocprofiler_get_agent(group.context, &agent);
check_status(status);
const rocprofiler::util::AgentInfo* agent_info = rocprofiler::util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
const rocprofiler::util::AgentInfo* agent_info =
rocprofiler::util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
pcsmp_callback_data_t pcsmp_data{};
pcsmp_data.kernel_name = (const char*)arg;
pcsmp_data.data_buffer = rocprofiler::util::HsaRsrcFactory::Instance().AllocateSysMemory(agent_info, rocprofiler::TraceProfile::GetSize());
pcsmp_data.data_buffer = rocprofiler::util::HsaRsrcFactory::Instance().AllocateSysMemory(
agent_info, rocprofiler::TraceProfile::GetSize());
status = rocprofiler_iterate_trace_data(group.context, trace_data_cb, &pcsmp_data);
check_status(status);
return false;
@@ -110,8 +111,8 @@ bool context_handler(rocprofiler_group_t group, void* arg) {
hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data, void* user_data,
rocprofiler_group_t* group) {
// context features
const rocprofiler_feature_kind_t trace_kind =
(rocprofiler_feature_kind_t)(ROCPROFILER_FEATURE_KIND_TRACE | ROCPROFILER_FEATURE_KIND_PCSMP_MOD);
const rocprofiler_feature_kind_t trace_kind = (rocprofiler_feature_kind_t)(
ROCPROFILER_FEATURE_KIND_TRACE | ROCPROFILER_FEATURE_KIND_PCSMP_MOD);
const uint32_t feature_count = 1;
const uint32_t parameter_count = 1;
rocprofiler_feature_t* features = new rocprofiler_feature_t[feature_count];
@@ -131,8 +132,8 @@ hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data,
properties.handler_arg = (void*)strdup(callback_data->kernel_name);
// Open profiling context
hsa_status_t status = rocprofiler_open(callback_data->agent, features, feature_count,
&context, 0 /*ROCPROFILER_MODE_SINGLEGROUP*/, &properties);
hsa_status_t status = rocprofiler_open(callback_data->agent, features, feature_count, &context,
0 /*ROCPROFILER_MODE_SINGLEGROUP*/, &properties);
check_status(status);
// Get group[0]
@@ -141,7 +142,7 @@ hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data,
return status;
}
} // namespace activity_prim
} // namespace activity_prim
extern "C" {
PUBLIC_API const char* GetOpName(uint32_t op) { return strdup("PCSAMPLE"); }
@@ -152,7 +153,8 @@ PUBLIC_API bool RemoveApiCallback(uint32_t op) { return true; }
PUBLIC_API bool InitActivityCallback(void* callback, void* arg) {
activity_prim::activity_arg = arg;
activity_prim::activity_callback.store((activity_async_callback_t)callback, std::memory_order_release);
activity_prim::activity_callback.store((activity_async_callback_t)callback,
std::memory_order_release);
rocprofiler_queue_callbacks_t queue_callbacks{};
queue_callbacks.dispatch = activity_prim::dispatch_callback;
@@ -191,11 +193,8 @@ struct evt_cb_entry_t {
};
evt_cb_entry_t evt_cb_table[HSA_EVT_ID_NUMBER];
hsa_status_t codeobj_evt_callback(
rocprofiler_hsa_cb_id_t id,
const rocprofiler_hsa_callback_data_t* cb_data,
void* arg)
{
hsa_status_t codeobj_evt_callback(rocprofiler_hsa_cb_id_t id,
const rocprofiler_hsa_callback_data_t* cb_data, void* arg) {
const auto evt = evt_cb_table[id].get();
activity_rtapi_callback_t evt_callback = (activity_rtapi_callback_t)evt.first;
if (evt_callback != NULL) evt_callback(ACTIVITY_DOMAIN_HSA_EVT, id, cb_data, evt.second);
projects/rocprofiler/src/core/activity.h
+1 -1
Bestand weergeven
@@ -19,4 +19,4 @@ enum hsa_evt_id_t {
// HSA EVT callback data type
typedef rocprofiler_hsa_callback_data_t hsa_evt_data_t;
#endif // _SRC_CORE_ACTIVITY_H
#endif // _SRC_CORE_ACTIVITY_H
projects/rocprofiler/src/core/context.h
+65 -67
Bestand weergeven
@@ -27,7 +27,7 @@ THE SOFTWARE.
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <unistd.h> // usleep
#include <unistd.h> // usleep
#include <atomic>
#include <list>
#include <map>
@@ -91,8 +91,7 @@ class Group {
barrier_signal_{},
dispatch_signal_{},
orig_signal_{},
record_{}
{}
record_{} {}
void Insert(const profile_info_t& info) {
const rocprofiler_feature_kind_t kind = info.rinfo->kind;
@@ -110,11 +109,10 @@ class Group {
}
hsa_status_t Finalize(const bool is_concurrent = false) {
hsa_status_t status = pmc_profile_.Finalize(start_vector_, stop_vector_,
read_vector_, is_concurrent);
hsa_status_t status =
pmc_profile_.Finalize(start_vector_, stop_vector_, read_vector_, is_concurrent);
if (status == HSA_STATUS_SUCCESS) {
status = trace_profile_.Finalize(start_vector_, stop_vector_,
read_vector_, is_concurrent);
status = trace_profile_.Finalize(start_vector_, stop_vector_, read_vector_, is_concurrent);
}
if (status == HSA_STATUS_SUCCESS) {
if (!pmc_profile_.Empty()) ++n_profiles_;
@@ -137,32 +135,20 @@ class Group {
Context* GetContext() { return context_; }
uint32_t GetIndex() const { return index_; }
void SetBarrierSignal(const hsa_signal_t &signal) {
barrier_signal_ = signal;
}
hsa_signal_t& GetBarrierSignal() {
return barrier_signal_;
}
void SetDispatchSignal(const hsa_signal_t &signal) {
dispatch_signal_ = signal;
}
hsa_signal_t& GetDispatchSignal() {
return dispatch_signal_;
}
void SetOrigSignal(const hsa_signal_t &signal) {
orig_signal_ = signal;
}
const hsa_signal_t& GetOrigSignal() const {
return orig_signal_;
}
rocprofiler_dispatch_record_t* GetRecord() {
return &record_;
}
void SetBarrierSignal(const hsa_signal_t& signal) { barrier_signal_ = signal; }
hsa_signal_t& GetBarrierSignal() { return barrier_signal_; }
void SetDispatchSignal(const hsa_signal_t& signal) { dispatch_signal_ = signal; }
hsa_signal_t& GetDispatchSignal() { return dispatch_signal_; }
void SetOrigSignal(const hsa_signal_t& signal) { orig_signal_ = signal; }
const hsa_signal_t& GetOrigSignal() const { return orig_signal_; }
rocprofiler_dispatch_record_t* GetRecord() { return &record_; }
atomic_refs_t* AtomicRefsCount() { return reinterpret_cast<atomic_refs_t*>(&refs_); }
void ResetRefsCount() { AtomicRefsCount()->store(n_profiles_, std::memory_order_release); }
void IncrRefsCount() { AtomicRefsCount()->fetch_add(1, std::memory_order_acq_rel); }
uint32_t FetchDecrRefsCount() { return AtomicRefsCount()->fetch_sub(1, std::memory_order_acq_rel); }
uint32_t FetchDecrRefsCount() {
return AtomicRefsCount()->fetch_sub(1, std::memory_order_acq_rel);
}
private:
PmcProfile pmc_profile_;
@@ -188,23 +174,23 @@ class Context {
public:
typedef std::map<std::string, rocprofiler_feature_t*> info_map_t;
static void Create(Context* obj, const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info,
const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg)
{
static void Create(Context* obj, const util::AgentInfo* agent_info, Queue* queue,
rocprofiler_feature_t* info, const uint32_t info_count,
rocprofiler_handler_t handler, void* handler_arg) {
new (obj) Context(agent_info, queue, info, info_count, handler, handler_arg);
obj->Construct(agent_info, queue, info, info_count, handler, handler_arg);
}
static void Release(Context* obj) { obj->Destruct(); }
static Context* Create(const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info,
const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg)
{
static Context* Create(const util::AgentInfo* agent_info, Queue* queue,
rocprofiler_feature_t* info, const uint32_t info_count,
rocprofiler_handler_t handler, void* handler_arg) {
Context* obj = new Context(agent_info, queue, info, info_count, handler, handler_arg);
if (obj == NULL) EXC_RAISING(HSA_STATUS_ERROR, "allocation error");
try {
obj->Construct(agent_info, queue, info, info_count, handler, handler_arg);
} catch(...) {
} catch (...) {
delete obj;
obj = NULL;
std::cerr << "Error: Context Create failed" << std::endl;
@@ -213,7 +199,9 @@ class Context {
return obj;
}
static void Destroy(Context* obj) { if (obj != NULL) delete obj; }
static void Destroy(Context* obj) {
if (obj != NULL) delete obj;
}
void Reset(const uint32_t& group_index) { set_[group_index].ResetRefsCount(); }
@@ -293,8 +281,10 @@ class Context {
hsa_rsrc_->SignalWaitRestore(tuple.completion_signal, 1);
// Restore other signals
RestoreSignals(tuple);
for (rocprofiler_feature_t* rinfo : *(tuple.info_vector)) rinfo->data.kind = ROCPROFILER_DATA_KIND_UNINIT;
callback_data_t callback_data{tuple.profile, tuple.info_vector, tuple.info_vector->size(), NULL};
for (rocprofiler_feature_t* rinfo : *(tuple.info_vector))
rinfo->data.kind = ROCPROFILER_DATA_KIND_UNINIT;
callback_data_t callback_data{tuple.profile, tuple.info_vector, tuple.info_vector->size(),
NULL};
const hsa_status_t status =
api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, DataCallback, &callback_data);
if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "context iterate data failed");
@@ -310,7 +300,8 @@ class Context {
if (expr) {
auto it = info_map_.find(name);
if (it == info_map_.end())
EXC_RAISING(HSA_STATUS_ERROR, "metric '" << name << "', rocprofiler info is not found " << this);
EXC_RAISING(HSA_STATUS_ERROR,
"metric '" << name << "', rocprofiler info is not found " << this);
rocprofiler_feature_t* info = it->second;
info->data.result_double = expr->Eval(args);
info->data.kind = ROCPROFILER_DATA_KIND_DOUBLE;
@@ -324,7 +315,7 @@ class Context {
for (auto& tuple : profile_vector) {
if (pcsmp_mode_) const_cast<profile_t*>(tuple.profile)->event_count = UINT32_MAX;
const hsa_status_t status =
api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, callback, data);
api_->hsa_ven_amd_aqlprofile_iterate_data(tuple.profile, callback, data);
if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "context iterate data failed");
}
}
@@ -342,7 +333,10 @@ class Context {
hsa_agent_t GetAgent() const { return agent_; }
Group* GetGroup(const uint32_t& index) { return &set_[index]; }
rocprofiler_handler_t GetHandler(void** arg) const { *arg = handler_arg_; return handler_; }
rocprofiler_handler_t GetHandler(void** arg) const {
*arg = handler_arg_;
return handler_;
}
// Concurrent profiling mode
static bool k_concurrent_;
@@ -358,8 +352,7 @@ class Context {
metrics_(NULL),
handler_(handler),
handler_arg_(handler_arg),
pcsmp_mode_(false)
{}
pcsmp_mode_(false) {}
~Context() { Destruct(); }
@@ -375,8 +368,7 @@ class Context {
}
void Construct(const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info,
const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg)
{
const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg) {
if (info_count == 0) {
set_.push_back(Group(agent_info_, this, 0));
return;
@@ -386,9 +378,11 @@ class Context {
if (metrics_ == NULL) EXC_RAISING(HSA_STATUS_ERROR, "MetricsDict create failed");
if (Initialize(info, info_count) == false) {
fprintf(stdout, "\nInput metrics out of HW limit. Proposed metrics group set:\n"); fflush(stdout);
fprintf(stdout, "\nInput metrics out of HW limit. Proposed metrics group set:\n");
fflush(stdout);
MetricsGroupSet(agent_info, info, info_count).Print(stdout);
fprintf(stdout, "\n"); fflush(stdout);
fprintf(stdout, "\n");
fflush(stdout);
EXC_RAISING(HSA_STATUS_ERROR, "Metrics list exceeds HW limits");
}
Finalize();
@@ -420,8 +414,8 @@ class Context {
info_map_[name] = info;
auto ret = metrics_map_.insert({name, NULL});
if (!ret.second)
EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name
<< "' is registered more then once");
EXC_RAISING(HSA_STATUS_ERROR,
"input metric '" << name << "' is registered more then once");
}
}
@@ -437,8 +431,9 @@ class Context {
if (kind == ROCPROFILER_FEATURE_KIND_METRIC) { // Processing metrics features
const Metric* metric = metrics_->Get(name);
if (metric == NULL)
EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not supported on this hardware: "
<< agent_info_->name);
EXC_RAISING(HSA_STATUS_ERROR,
"input metric '"
<< name << "' is not supported on this hardware: " << agent_info_->name);
#if 0
std::cout << " " << name << (metric->GetExpr() ? " = " + metric->GetExpr()->String() : " counter") << std::endl;
#endif
@@ -493,9 +488,9 @@ class Context {
info->kind = ROCPROFILER_FEATURE_KIND_TRACE;
const event_t* event = NULL;
if (kind & ROCPROFILER_FEATURE_KIND_PCSMP_MOD) { // PC sampling
if (kind & ROCPROFILER_FEATURE_KIND_PCSMP_MOD) { // PC sampling
pcsmp_mode_ = true;
} else if (kind & ROCPROFILER_FEATURE_KIND_SPM_MOD) { // SPM trace
} else if (kind & ROCPROFILER_FEATURE_KIND_SPM_MOD) { // SPM trace
const Metric* metric = metrics_->Get(name);
if (metric == NULL)
EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not found");
@@ -559,14 +554,14 @@ class Context {
const bool trace_local = TraceProfile::IsLocal();
util::HsaRsrcFactory* hsa_rsrc = &util::HsaRsrcFactory::Instance();
if (sample_id == 0) {
const uint32_t output_buffer_size = profile->output_buffer.size;
const uint32_t output_buffer_size64 = profile->output_buffer.size / sizeof(uint64_t);
const util::AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(profile->agent);
void* ptr = (trace_local) ? hsa_rsrc->AllocateSysMemory(agent_info, output_buffer_size) :
calloc(output_buffer_size64, sizeof(uint64_t));
rinfo->data.result_bytes.size = output_buffer_size;
rinfo->data.result_bytes.ptr = ptr;
callback_data->ptr = reinterpret_cast<char*>(ptr);
const uint32_t output_buffer_size = profile->output_buffer.size;
const uint32_t output_buffer_size64 = profile->output_buffer.size / sizeof(uint64_t);
const util::AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(profile->agent);
void* ptr = (trace_local) ? hsa_rsrc->AllocateSysMemory(agent_info, output_buffer_size)
: calloc(output_buffer_size64, sizeof(uint64_t));
rinfo->data.result_bytes.size = output_buffer_size;
rinfo->data.result_bytes.ptr = ptr;
callback_data->ptr = reinterpret_cast<char*>(ptr);
}
char* result_bytes_ptr = reinterpret_cast<char*>(rinfo->data.result_bytes.ptr);
const char* end = result_bytes_ptr + rinfo->data.result_bytes.size;
@@ -577,8 +572,10 @@ class Context {
char* dest = ptr + sizeof(*header);
if ((dest + size) >= end) {
if (dest < end) size = end - dest;
else EXC_RAISING(HSA_STATUS_ERROR, "Trace data out of output buffer");
if (dest < end)
size = end - dest;
else
EXC_RAISING(HSA_STATUS_ERROR, "Trace data out of output buffer");
}
bool suc = true;
@@ -593,7 +590,9 @@ class Context {
rinfo->data.result_bytes.instance_count = sample_id + 1;
rinfo->data.kind = ROCPROFILER_DATA_KIND_BYTES;
} else
EXC_RAISING(HSA_STATUS_ERROR, "Agent Memcpy failed, dst(" << (void*)dest << ") src(" << (void*)src << ") size(" << size << ")");
EXC_RAISING(HSA_STATUS_ERROR,
"Agent Memcpy failed, dst(" << (void*)dest << ") src(" << (void*)src
<< ") size(" << size << ")");
} else {
if (sample_id == 0) {
rinfo->data.result_bytes.ptr = profile->output_buffer.ptr;
@@ -647,8 +646,7 @@ class Context {
bool pcsmp_mode_;
};
#define CONTEXT_INSTANTIATE() \
bool rocprofiler::Context::k_concurrent_ = false;
#define CONTEXT_INSTANTIATE() bool rocprofiler::Context::k_concurrent_ = false;
} // namespace rocprofiler
projects/rocprofiler/src/core/context_pool.h
+37 -45
Bestand weergeven
@@ -31,7 +31,7 @@ THE SOFTWARE.
namespace rocprofiler {
class ContextPool {
public:
public:
typedef uint64_t index_t;
typedef std::mutex mutex_t;
@@ -41,16 +41,12 @@ class ContextPool {
std::atomic<bool> completed;
};
static ContextPool* Create(
uint32_t num_entries,
uint32_t payload_bytes,
const util::AgentInfo* agent_info,
rocprofiler_feature_t* info,
const uint32_t info_count,
rocprofiler_pool_handler_t handler,
void* handler_arg)
{
ContextPool* obj = new ContextPool(num_entries, payload_bytes, agent_info, info, info_count, handler, handler_arg);
static ContextPool* Create(uint32_t num_entries, uint32_t payload_bytes,
const util::AgentInfo* agent_info, rocprofiler_feature_t* info,
const uint32_t info_count, rocprofiler_pool_handler_t handler,
void* handler_arg) {
ContextPool* obj = new ContextPool(num_entries, payload_bytes, agent_info, info, info_count,
handler, handler_arg);
if (obj == NULL) EXC_RAISING(HSA_STATUS_ERROR, "allocation error");
return obj;
}
@@ -61,18 +57,18 @@ class ContextPool {
if (constructed_ == false) {
Construct(agent_info_, info_, info_count_);
}
const index_t write_index = write_index_.fetch_add(entry_size_bytes_, std::memory_order_relaxed);
const index_t write_index =
write_index_.fetch_add(entry_size_bytes_, std::memory_order_relaxed);
while (write_index >= (read_index_.load(std::memory_order_acquire) + array_size_bytes_)) {
check_completed();
std::this_thread::yield();
}
entry_t* entry = GetPoolEntry(write_index, pool_entry);
if (entry->completed.load(std::memory_order_relaxed) != false) EXC_RAISING(HSA_STATUS_ERROR, "Corrupted pool entry");
if (entry->completed.load(std::memory_order_relaxed) != false)
EXC_RAISING(HSA_STATUS_ERROR, "Corrupted pool entry");
}
void Flush() {
check_completed();
}
void Flush() { check_completed(); }
#if 0
template <class F>
F for_each(const F& f_p) {
@@ -95,7 +91,7 @@ class ContextPool {
return f;
}
#endif
private:
private:
static unsigned aligned64(const unsigned& size) { return (size + 0x3f) & ~0x3fu; }
static bool context_handler(rocprofiler_group_t group, void* arg) {
@@ -105,45 +101,41 @@ class ContextPool {
return true;
}
ContextPool(
uint32_t num_entries,
uint32_t payload_bytes,
const util::AgentInfo* agent_info,
rocprofiler_feature_t* info,
const uint32_t info_count,
rocprofiler_pool_handler_t pool_handler,
void* pool_handler_arg
) :
payload_off_(aligned64(sizeof(entry_t))),
entry_size_bytes_(payload_off_ + aligned64(payload_bytes)),
array_size_bytes_(entry_size_bytes_ * num_entries),
array_(NULL),
read_index_(0),
write_index_(0),
sync_flag_(false),
ContextPool(uint32_t num_entries, uint32_t payload_bytes, const util::AgentInfo* agent_info,
rocprofiler_feature_t* info, const uint32_t info_count,
rocprofiler_pool_handler_t pool_handler, void* pool_handler_arg)
: payload_off_(aligned64(sizeof(entry_t))),
entry_size_bytes_(payload_off_ + aligned64(payload_bytes)),
array_size_bytes_(entry_size_bytes_ * num_entries),
array_(NULL),
read_index_(0),
write_index_(0),
sync_flag_(false),
agent_info_(agent_info),
info_(info),
info_count_(info_count),
pool_handler_(pool_handler),
pool_handler_arg_(pool_handler_arg),
constructed_(false)
{}
agent_info_(agent_info),
info_(info),
info_count_(info_count),
pool_handler_(pool_handler),
pool_handler_arg_(pool_handler_arg),
constructed_(false) {}
void Construct(const util::AgentInfo* agent_info, rocprofiler_feature_t* info, const uint32_t info_count) {
void Construct(const util::AgentInfo* agent_info, rocprofiler_feature_t* info,
const uint32_t info_count) {
std::lock_guard<mutex_t> lck(mutex_);
if (constructed_ == false) {
array_data_ = (char*) malloc(array_size_bytes_ + 0x3f);
array_data_ = (char*)malloc(array_size_bytes_ + 0x3f);
array_ = reinterpret_cast<char*>(((intptr_t)array_data_ + 0x3f) >> 6 << 6);
if (((intptr_t)array_ & 0x3f) != 0) EXC_RAISING(HSA_STATUS_ERROR, "Pool array is not aligned");
if (((intptr_t)array_ & 0x3f) != 0)
EXC_RAISING(HSA_STATUS_ERROR, "Pool array is not aligned");
memset(array_, 0, array_size_bytes_);
const char* end = array_ + array_size_bytes_;
for (char* ptr = array_; ptr < end; ptr += entry_size_bytes_) {
entry_t* entry = reinterpret_cast<entry_t*>(ptr);
entry->pool = this;
entry->context = Context::Create(agent_info, NULL, info, info_count, ContextPool::context_handler, ptr);
entry->context =
Context::Create(agent_info, NULL, info, info_count, ContextPool::context_handler, ptr);
}
constructed_ = true;
@@ -175,7 +167,7 @@ class ContextPool {
if (sync_flag_.test_and_set(std::memory_order_acquire) == false) {
index_t read_index = read_index_.load(std::memory_order_relaxed);
const index_t write_index = write_index_.load(std::memory_order_relaxed);
while(read_index < write_index) {
while (read_index < write_index) {
rocprofiler_pool_entry_t pool_entry{};
entry_t* entry = GetPoolEntry(read_index, &pool_entry);
if (entry->completed.load(std::memory_order_acquire) == true) {
projects/rocprofiler/src/core/core_timer.h
+11 -12
Bestand weergeven
@@ -1,8 +1,7 @@
#ifndef _CORE_TIMER_H_
#define _CORE_TIMER_H_
template <int Size>
class CoreTimer {
template <int Size> class CoreTimer {
CoreTimer() {
index_ = 0;
freq_in_100mhz_ = MeasureTSCFreqHz();
@@ -20,15 +19,15 @@ class CoreTimer {
// AMD Linux timing
unsigned int unused;
n = __rdtscp(&unused);
data_[index_] = 10 * n / freq_in_100mhz_; // unit is ns
data_[index_] = 10 * n / freq_in_100mhz_; // unit is ns
index_ += 1;
}
double Print()
double Print()
private:
// timer data
double data_[Size];
private :
// timer data
double data_[Size];
// data index
uint32_t index_;
// frequency
@@ -40,20 +39,20 @@ class CoreTimer {
clock_gettime(CLOCK_MONOTONIC_RAW, &ts);
return uint64_t(ts.tv_sec) * 1000000 + ts.tv_nsec / 1000;
}
uint64_t CoreTimer::MeasureTSCFreqHz() {
// Make a coarse interval measurement of TSC ticks for 1 gigacycles.
unsigned int unused;
uint64_t tscTicksEnd;
uint64_t coarseBeginUs = CoarseTimestampUs();
uint64_t tscTicksBegin = __rdtscp(&unused);
do {
tscTicksEnd = __rdtscp(&unused);
} while (tscTicksEnd - tscTicksBegin < 1000000000);
uint64_t coarseEndUs = CoarseTimestampUs();
// Compute the TSC frequency and round to nearest 100MHz.
uint64_t coarseIntervalNs = (coarseEndUs - coarseBeginUs) * 1000;
uint64_t tscIntervalTicks = tscTicksEnd - tscTicksBegin;
@@ -61,4 +60,4 @@ class CoreTimer {
}
};
#endif // _CORE_TIMER_H_
#endif // _CORE_TIMER_H_
projects/rocprofiler/src/core/counters/derived/derived_counter.cpp.in
+5 -9
Bestand weergeven
@@ -27,8 +27,7 @@ namespace Counter {
static std::atomic<uint64_t> COUNTER_COUNTER{0};
DerivedCounter::DerivedCounter(std::string name, std::string description,
std::string gpu_name)
DerivedCounter::DerivedCounter(std::string name, std::string description, std::string gpu_name)
: Counter(name, description, gpu_name) {
metric_id_ = COUNTER_COUNTER.fetch_add(1, std::memory_order_release);
addCounterToCounterMap();
@@ -41,20 +40,17 @@ DerivedCounter::~DerivedCounter() {
uint64_t DerivedCounter::getMetricId() { return metric_id_; }
std::map<uint64_t, BasicCounter*> *DerivedCounter::getAllCounters() {
return &counters_;
}
std::map<uint64_t, BasicCounter*>* DerivedCounter::getAllCounters() { return &counters_; }
BasicCounter *DerivedCounter::getBasicCounterFromDerived(uint64_t counter_id) {
BasicCounter* DerivedCounter::getBasicCounterFromDerived(uint64_t counter_id) {
return counters_[counter_id];
}
void DerivedCounter::addBasicCounter(uint64_t counter_id,
BasicCounter *counter) {
void DerivedCounter::addBasicCounter(uint64_t counter_id, BasicCounter* counter) {
counters_.emplace(counter_id, counter);
}
@DERIVED_XML_PARSE_RESULT@
@DERIVED_XML_PARSE_RESULT @
} // namespace Counter
projects/rocprofiler/src/core/counters/derived/derived_counter.h
+1 -2
Bestand weergeven
@@ -39,8 +39,7 @@ namespace Counter {
class DerivedCounter : Counter {
public:
std::function<uint64_t()> evaluate_metric;
DerivedCounter(std::string name, std::string description,
std::string gpu_name);
DerivedCounter(std::string name, std::string description, std::string gpu_name);
~DerivedCounter();
uint64_t getMetricId();
projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp
+4 -3
Bestand weergeven
@@ -108,7 +108,7 @@ bool metrics::ExtractMetricEvents(
// adding result object for derived metric
std::lock_guard<std::mutex> lock(extract_metric_events_lock);
if(metric_names[i].compare("KERNEL_DURATION")==0) {
if (metric_names[i].compare("KERNEL_DURATION") == 0) {
if (results_map.find(metric_names[i]) == results_map.end()) {
results_map[metric_names[i]] = new results_t(metric_names[i], {}, xcc_count);
}
@@ -192,7 +192,7 @@ bool metrics::GetMetricsData(std::map<std::string, results_t*>& results_map,
auto it = results_map.find(metric->GetName());
if (it == results_map.end()) rocprofiler::fatal("metric results not found ");
results_t* res = it->second;
if(metric->GetName().compare("KERNEL_DURATION") == 0) {
if (metric->GetName().compare("KERNEL_DURATION") == 0) {
res->val_double = kernel_duration;
continue;
}
@@ -206,7 +206,8 @@ bool metrics::GetMetricsData(std::map<std::string, results_t*>& results_map,
void metrics::GetCountersAndMetricResultsByXcc(uint32_t xcc_index,
std::vector<results_t*>& results_list,
std::map<std::string, results_t*>& results_map,
std::vector<const Metric*>& metrics_list, uint64_t kernel_duration) {
std::vector<const Metric*>& metrics_list,
uint64_t kernel_duration) {
for (auto it = results_list.begin(); it != results_list.end(); it++) {
(*it)->val_double =
(*it)->xcc_vals[xcc_index]; // set val_double to hold value for specific xcc
projects/rocprofiler/src/core/counters/metrics/eval_metrics.h
+7 -6
Bestand weergeven
@@ -35,10 +35,10 @@ namespace rocprofiler {
typedef std::vector<double> xcc_results_t;
class results_t{
public:
results_t(std::string in_name, event_t in_event, uint32_t xcc_count):
name(in_name), val_double(0), event(in_event) {
class results_t {
public:
results_t(std::string in_name, event_t in_event, uint32_t xcc_count)
: name(in_name), val_double(0), event(in_event) {
xcc_vals.resize(xcc_count);
std::fill(xcc_vals.begin(), xcc_vals.end(), 0);
}
@@ -78,8 +78,9 @@ bool GetMetricsData(std::map<std::string, results_t*>& results_map,
std::vector<const Metric*>& metrics_list, uint64_t kernel_duration = 0);
void GetCountersAndMetricResultsByXcc(uint32_t xcc_index, std::vector<results_t*>& results_list,
std::map<std::string, results_t*>& results_map,
std::vector<const Metric*>& metrics_list, uint64_t kernel_duration = 0);
std::map<std::string, results_t*>& results_map,
std::vector<const Metric*>& metrics_list,
uint64_t kernel_duration = 0);
} // namespace metrics
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/metrics/exception.h
+1 -1
Bestand weergeven
@@ -45,7 +45,7 @@ THE SOFTWARE.
do { \
std::ostringstream oss; \
oss << __FUNCTION__ << "(), " << stream; \
throw rocprofiler::util::exception(error, oss.str()); \
throw rocprofiler::util::exception(error, oss.str()); \
} while (0)
#define AQL_EXC_RAISING(error, stream) \
projects/rocprofiler/src/core/counters/metrics/metrics.h
Executable → Regular
+3 -6
Bestand weergeven
@@ -221,14 +221,11 @@ class MetricsDict {
agent_name_ = agent_name_.substr(0, agent_name_.find(':'));
std::unordered_set<std::string> supported_agent_names = {
"gfx906",
"gfx908",
"gfx906", "gfx908",
"gfx90a", // Vega
"gfx940",
"gfx941",
"gfx940", "gfx941",
"gfx942", // Mi300
"gfx1030",
"gfx1031",
"gfx1030", "gfx1031",
"gfx1032", // Navi2x
"gfx1100",
"gfx1101" // Navi3x
projects/rocprofiler/src/core/counters/mmio/df_counters_mi200.h
+2 -3
Bestand weergeven
@@ -17,8 +17,8 @@ class DFPerfMonMI200 : public PerfMon {
DFPerfMonMI200(const Agent::AgentInfo& info);
~DFPerfMonMI200();
void Start() override;
void Stop() {};
void Read(std::vector<rocprofiler_counters_sampler_counter_output_t>& values) {};
void Stop(){};
void Read(std::vector<rocprofiler_counters_sampler_counter_output_t>& values){};
void SetCounterNames(std::vector<std::string>& counter_names);
mmio::mmap_type_t Type() override { return mmio::mmap_type_t::DF_PERFMON; }
@@ -31,7 +31,6 @@ class DFPerfMonMI200 : public PerfMon {
uint64_t GetFicaNodeOutboundBw(uint32_t ficaa_val);
private:
mmio::DFPerfmonMMIO* mmio_;
static std::mutex mutex_; // should be an MMIO member
projects/rocprofiler/src/core/counters/mmio/pcie_counters_mi200.cpp
+17 -25
Bestand weergeven
@@ -13,12 +13,12 @@ PciePerfMonMI200::~PciePerfMonMI200() {
mmio::MMIOManager::DestroyMMIOInstance(dynamic_cast<mmio::MMIO*>(mmio_));
}
void PciePerfMonMI200::writeRegister(uint32_t reg_offset, uint32_t value){
void PciePerfMonMI200::writeRegister(uint32_t reg_offset, uint32_t value) {
// mmio or ioctl approaches
mmio_->RegisterWriteAPI(reg_offset, value);
mmio_->RegisterWriteAPI(reg_offset, value);
}
void PciePerfMonMI200::readRegister(uint32_t reg_offset, uint32_t& value){
void PciePerfMonMI200::readRegister(uint32_t reg_offset, uint32_t& value) {
// mmio or ioctl approaches
mmio_->RegisterReadAPI(reg_offset, value);
}
@@ -35,44 +35,40 @@ void PciePerfMonMI200::SetCounterNames(std::vector<std::string>& counter_names)
}
}
void PciePerfMonMI200::Start(){
void PciePerfMonMI200::Start() {
// TODO: make sure values stored in table
// in registers header are dec and not hex
Start_RX_TILE_SCLK(event_id_);
}
void PciePerfMonMI200::Stop(){
void PciePerfMonMI200::Stop() {
// TODO: revisit correct value to stop
writeRegister(PCIE_MI200::PCIE_PERF_COUNT_CNTL, 0x2); // stop
writeRegister(PCIE_MI200::PCIE_PERF_COUNT_CNTL, 0x2); // stop
}
void PciePerfMonMI200::Read(std::vector<rocprofiler_counters_sampler_counter_output_t>& values){
uint64_t val=0;
void PciePerfMonMI200::Read(std::vector<rocprofiler_counters_sampler_counter_output_t>& values) {
uint64_t val = 0;
Read_RX_TILE_SCLK(val);
rocprofiler_counters_sampler_counter_output_t value = {
ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS,
static_cast<double>(val)
};
rocprofiler_counters_sampler_counter_output_t value = {ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS,
static_cast<double>(val)};
values.push_back(value);
}
void PciePerfMonMI200::Start_RX_TILE_TXCLK(uint32_t event){
void PciePerfMonMI200::Start_RX_TILE_TXCLK(uint32_t event) {
// Step 1: PORT SEL update
writeRegister(PCIE_MI200::PCIE_PERF_CNTL_EVENT_CI_PORT_SEL, 0x0);
// Step 2: EVENT SEL update
uint32_t value = event; // last 8 bits for event
uint32_t value = event; // last 8 bits for event
writeRegister(PCIE_MI200::PCIE_PERF_CNTL_TXCLK3, value);
// Steps 3 & 4: Performance counters initialization, enable:
// TODO: revisit. Just a single write with 0x3 might be enough (check with pcie team)
writeRegister(PCIE_MI200::PCIE_PERF_COUNT_CNTL, 0x5);
}
void PciePerfMonMI200::Read_RX_TILE_TXCLK(uint64_t& result){
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);
@@ -84,22 +80,20 @@ void PciePerfMonMI200::Read_RX_TILE_TXCLK(uint64_t& result){
result = val | lo_val;
}
void PciePerfMonMI200::Start_RX_TILE_SCLK(uint32_t event){
void PciePerfMonMI200::Start_RX_TILE_SCLK(uint32_t event) {
// Step 1: PORT SEL update
writeRegister(PCIE_MI200::PCIE_PERF_CNTL_EVENT_CI_PORT_SEL, 0x0);
// Step 2: EVENT SEL update
uint32_t value = event; // last 8 bits for event
uint32_t value = event; // last 8 bits for event
writeRegister(PCIE_MI200::PCIE_PERF_CNTL_LCLK1, value);
// Steps 3 & 4: Performance counters initialization, enable:
// TODO: revisit. Just a single write with 0x3 might be enough (check with pcie team)
writeRegister(PCIE_MI200::PCIE_PERF_COUNT_CNTL, 0x5);
}
void PciePerfMonMI200::Read_RX_TILE_SCLK(uint64_t& result){
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);
@@ -111,6 +105,4 @@ void PciePerfMonMI200::Read_RX_TILE_SCLK(uint64_t& result){
result = val | lo_val;
}
} // namespace rocprofiler
} // namespace rocprofiler
projects/rocprofiler/src/core/counters/mmio/pcie_counters_mi200.h
+1 -1
Bestand weergeven
@@ -22,7 +22,7 @@ class PciePerfMonMI200 : public PerfMon {
mmio::mmap_type_t Type() override { return mmio::mmap_type_t::PCIE_PERFMON; }
private:
// TODO : check google coding std
// TODO : check google coding std
void writeRegister(uint32_t reg_offset, uint32_t value);
void readRegister(uint32_t reg_offset, uint32_t& value);
projects/rocprofiler/src/core/counters/mmio/pcie_perfmon_registers_mi200.h
+236 -237
Bestand weergeven
@@ -4,70 +4,70 @@
#include <stdint.h>
namespace PCIE_MI200 {
// -------- RX Tile TXCLK Start --------
// Step 1: PORT SEL update
const static uint32_t PCIE_PERF_CNTL_EVENT_CI_PORT_SEL = 0x11180250;
// Step 2: EVENT SEL update
const static uint32_t PCIE_PERF_CNTL_TXCLK1 = 0x11180204;
const static uint32_t PCIE_PERF_CNTL_TXCLK2 = 0x11180210;
const static uint32_t PCIE_PERF_CNTL_TXCLK3 = 0x1118021C; //#
const static uint32_t PCIE_PERF_CNTL_TXCLK4 = 0x11180228; //#
const static uint32_t PCIE_PERF_CNTL_TXCLK5 = 0x11180258;
const static uint32_t PCIE_PERF_CNTL_TXCLK6 = 0x11180264;
const static uint32_t PCIE_PERF_CNTL_TXCLK7 = 0x11180888;
const static uint32_t PCIE_PERF_CNTL_TXCLK8 = 0x11180894;
const static uint32_t PCIE_PERF_CNTL_TXCLK9 = 0x111808A0;
const static uint32_t PCIE_PERF_CNTL_TXCLK10 = 0x111808AC;
const static uint32_t PCIE_PERF_CNTL_TXCLK1 = 0x11180204;
const static uint32_t PCIE_PERF_CNTL_TXCLK2 = 0x11180210;
const static uint32_t PCIE_PERF_CNTL_TXCLK3 = 0x1118021C; //#
const static uint32_t PCIE_PERF_CNTL_TXCLK4 = 0x11180228; //#
const static uint32_t PCIE_PERF_CNTL_TXCLK5 = 0x11180258;
const static uint32_t PCIE_PERF_CNTL_TXCLK6 = 0x11180264;
const static uint32_t PCIE_PERF_CNTL_TXCLK7 = 0x11180888;
const static uint32_t PCIE_PERF_CNTL_TXCLK8 = 0x11180894;
const static uint32_t PCIE_PERF_CNTL_TXCLK9 = 0x111808A0;
const static uint32_t PCIE_PERF_CNTL_TXCLK10 = 0x111808AC;
// Steps 3 & 4: Performance counters initialization, enable:
const static uint32_t PCIE_PERF_COUNT_CNTL = 0x11180200;
// Step 5: Performance counters read:
const static uint32_t PCIE_PERF_COUNT0_TXCLK1 = 0x11180208;
const static uint32_t PCIE_PERF_COUNT0_TXCLK2 = 0x11180214;
const static uint32_t PCIE_PERF_COUNT0_TXCLK3 = 0x11180220; //#
const static uint32_t PCIE_PERF_COUNT0_TXCLK4 = 0x1118022C; //#
const static uint32_t PCIE_PERF_COUNT0_TXCLK5 = 0x1118025C;
const static uint32_t PCIE_PERF_COUNT0_TXCLK6 = 0x11180268;
const static uint32_t PCIE_PERF_COUNT0_TXCLK7 = 0x1118088C;
const static uint32_t PCIE_PERF_COUNT0_TXCLK8 = 0x11180898;
const static uint32_t PCIE_PERF_COUNT0_TXCLK9 = 0x111808A4;
const static uint32_t PCIE_PERF_COUNT0_TXCLK10 = 0x111808B0;
const static uint32_t PCIE_PERF_COUNT0_TXCLK1 = 0x11180208;
const static uint32_t PCIE_PERF_COUNT0_TXCLK2 = 0x11180214;
const static uint32_t PCIE_PERF_COUNT0_TXCLK3 = 0x11180220; //#
const static uint32_t PCIE_PERF_COUNT0_TXCLK4 = 0x1118022C; //#
const static uint32_t PCIE_PERF_COUNT0_TXCLK5 = 0x1118025C;
const static uint32_t PCIE_PERF_COUNT0_TXCLK6 = 0x11180268;
const static uint32_t PCIE_PERF_COUNT0_TXCLK7 = 0x1118088C;
const static uint32_t PCIE_PERF_COUNT0_TXCLK8 = 0x11180898;
const static uint32_t PCIE_PERF_COUNT0_TXCLK9 = 0x111808A4;
const static uint32_t PCIE_PERF_COUNT0_TXCLK10 = 0x111808B0;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK1 = 0x111808E8;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK2 = 0x111808F0;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK3 = 0x111808F8; //#
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK4 = 0x11180900; //#
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK5 = 0x11180908;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK6 = 0x11180910;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK7 = 0x11180918;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK8 = 0x11180920;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK9 = 0x11180928;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK1 = 0x111808E8;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK2 = 0x111808F0;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK3 = 0x111808F8; //#
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK4 = 0x11180900; //#
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK5 = 0x11180908;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK6 = 0x11180910;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK7 = 0x11180918;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK8 = 0x11180920;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK9 = 0x11180928;
const static uint32_t PCIE_PERF_COUNT0_UPVAL_TXCLK10 = 0x11180930;
const static uint32_t PCIE_PERF_COUNT1_TXCLK1 = 0x1118020C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK2 = 0x11180218;
const static uint32_t PCIE_PERF_COUNT1_TXCLK3 = 0x11180224; //#
const static uint32_t PCIE_PERF_COUNT1_TXCLK4 = 0x11180230; //#
const static uint32_t PCIE_PERF_COUNT1_TXCLK5 = 0x11180260;
const static uint32_t PCIE_PERF_COUNT1_TXCLK6 = 0x1118026C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK7 = 0x11180890;
const static uint32_t PCIE_PERF_COUNT1_TXCLK8 = 0x1118089C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK9 = 0x111808A8;
const static uint32_t PCIE_PERF_COUNT1_TXCLK1 = 0x1118020C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK2 = 0x11180218;
const static uint32_t PCIE_PERF_COUNT1_TXCLK3 = 0x11180224; //#
const static uint32_t PCIE_PERF_COUNT1_TXCLK4 = 0x11180230; //#
const static uint32_t PCIE_PERF_COUNT1_TXCLK5 = 0x11180260;
const static uint32_t PCIE_PERF_COUNT1_TXCLK6 = 0x1118026C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK7 = 0x11180890;
const static uint32_t PCIE_PERF_COUNT1_TXCLK8 = 0x1118089C;
const static uint32_t PCIE_PERF_COUNT1_TXCLK9 = 0x111808A8;
const static uint32_t PCIE_PERF_COUNT1_TXCLK10 = 0x111808B4;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK1 = 0x111808EC;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK2 = 0x111808F4;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK3 = 0x111808FC; //#
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK4 = 0x11180904; //#
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK5 = 0x1118090C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK6 = 0x11180914;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK7 = 0x1118091C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK8 = 0x11180924;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK9 = 0x1118092C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK1 = 0x111808EC;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK2 = 0x111808F4;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK3 = 0x111808FC; //#
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK4 = 0x11180904; //#
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK5 = 0x1118090C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK6 = 0x11180914;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK7 = 0x1118091C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK8 = 0x11180924;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK9 = 0x1118092C;
const static uint32_t PCIE_PERF_COUNT1_UPVAL_TXCLK10 = 0x11180934;
@@ -127,201 +127,200 @@ const static uint32_t PCIE_PERF_COUNT1_UPVAL_LCLK8 = 0x11180974;
// -------- RX Tile SCLK End ----------
typedef enum{
TX_TILE_TXCLK = 0,
TX_TILE_SCLK = 1,
RX_TILE_TXCLK = 2,
RX_TILE_SCLK = 3,
LC_TILE_TXCLK = 4
}pcie_event_category_t;
typedef enum {
TX_TILE_TXCLK = 0,
TX_TILE_SCLK = 1,
RX_TILE_TXCLK = 2,
RX_TILE_SCLK = 3,
LC_TILE_TXCLK = 4
} pcie_event_category_t;
struct pcie_event_t{
pcie_event_t(int id, pcie_event_category_t cat): event_id(id), event_category(cat){}
int event_id;
pcie_event_category_t event_category;
struct pcie_event_t {
pcie_event_t(int id, pcie_event_category_t cat) : event_id(id), event_category(cat) {}
int event_id;
pcie_event_category_t event_category;
};
const static std::map<std::string, pcie_event_t> pcie_events_table = {
{"RX_PERF_RXP_RX_TailEdb_A[0]", {2, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[1]", {3, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[2]", {4, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[3]", {5, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[0]", {6, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[1]", {7, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[2]", {8, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[3]", {9, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[0]", {10, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[1]", {11, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[2]", {12, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[3]", {13, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[0]", {14, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[1]", {15, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[2]", {16, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[3]", {17, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_nullified_tlp_A", {18, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_valid_crc_A", {19, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_invalid_crc_A", {20, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_vendor_type1_A", {21, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_vendor_type0_A", {22, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_set_slot_power_limit_A", {23, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_unlock_A", {24, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_fatal_A", {25, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_nonfatal_A", {26, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_corr_A", {27, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pme_to_ack_A", {28, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pme_turn_off_A", {29, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pm_pme_A", {30, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pm_active_state_nak_A", {31, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intd_A", {32, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intc_A", {33, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intb_A", {34, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_inta_A", {35, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intd_A", {36, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intc_A", {37, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intb_A", {38, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_inta_A", {39, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_valid_A", {40, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_unsupported_A", {41, RX_TILE_TXCLK}},
{"RX_PERF_RCB_unexpected_cpl_A", {42, RX_TILE_TXCLK}},
{"RX_PERF_RCB_timeout_cpl_A", {43, RX_TILE_TXCLK}},
{"RX_PERF_HDS_tlphdrvalid_A", {44, RX_TILE_TXCLK}},
{"RX_PERF_HDS_tlpdatavalid_A", {45, RX_TILE_TXCLK}},
{"RX_PERF_GAN_bad_tlp_A", {46, RX_TILE_TXCLK}},
{"RX_PERF_GAN_nak_A", {47, RX_TILE_TXCLK}},
{"RX_PERF_GAN_ack_A", {48, RX_TILE_TXCLK}},
{"RX_PERF_FE_unsupported_req_A", {49, RX_TILE_TXCLK}},
{"RX_PERF_FE_unsupported_cpl_A", {50, RX_TILE_TXCLK}},
{"RX_PERF_FE_unexpected_cpl_A", {51, RX_TILE_TXCLK}},
{"RX_PERF_FE_poisoned_tlp_A", {52, RX_TILE_TXCLK}},
{"RX_PERF_FE_poisoned_cpl_A", {53, RX_TILE_TXCLK}},
{"RX_PERF_FE_malformed_tlp_A", {54, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_abort_A", {55, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MSG_A", {56, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_CFG_WR_A", {57, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_CFG_RD_A", {58, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_IO_WR_A", {59, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_IO_RD_A", {60, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MEM_WR_A", {61, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MEM_RD_A", {62, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_gt16_A", {63, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_9to16_A", {64, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_5to8_A", {65, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_2to4_A", {66, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_1_A", {67, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_gt32_A", {68, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_17to32_A", {69, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_9to16_A", {70, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_5to8_A", {71, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_2to4_A", {72, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_1_A", {73, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_CA_A", {74, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_CRS_A", {75, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_UR_A", {76, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_SC_A", {77, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_active_state_request_l1_A", {78, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_request_ack_A", {79, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_enter_l23_A", {80, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_enter_l1_A", {81, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_error_A", {82, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_crc_err_A", {83, RX_TILE_TXCLK}},
{"SB_PERF_FCC_npd_0", {84, RX_TILE_TXCLK}},
{"SB_PERF_FCC_pd_0", {85, RX_TILE_TXCLK}},
{"SB_PERF_FCC_nph_0", {86, RX_TILE_TXCLK}},
{"SB_PERF_FCC_ph_0", {87, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_rd_hdr_0", {88, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_rd_hdr_0", {89, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_wr_hdr_0", {90, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_wr_hdr_0", {91, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_data_0", {92, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_data_0", {93, RX_TILE_TXCLK}},
{"SB_PERF_invalid_crc_0", {94, RX_TILE_TXCLK}},
{"SB_PERF_valid_crc_0", {95, RX_TILE_TXCLK}},
{"SB_PERF_rd_hdr_WEN_0", {96, RX_TILE_TXCLK}},
{"SB_PERF_wr_hdr_WEN_0", {97, RX_TILE_TXCLK}},
{"SB_PERF_data_WEN_0", {98, RX_TILE_TXCLK}},
{"SB_PERF_non_post_rd_from_FE", {99, RX_TILE_TXCLK}},
{"SB_PERF_non_post_wr_from_FE", {100, RX_TILE_TXCLK}},
{"SB_PERF_post_req_from_FE", {101, RX_TILE_TXCLK}},
{"SB_PERF_non_post_rd_from_FE_0", {102, RX_TILE_TXCLK}},
{"SB_PERF_non_post_wr_from_FE_0", {103, RX_TILE_TXCLK}},
{"SB_PERF_post_req_from_FE_0", {104, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_nak_A", {111, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_ack_A", {112, RX_TILE_TXCLK}},
{"RX_PERF_allErrors_A", {113, RX_TILE_TXCLK}},
{"perf_PG_COUNT", {175, RX_TILE_TXCLK}},
{"perf_NOT_POWER_GATED", {176, RX_TILE_TXCLK}},
{"perf_POWER_GATED", {177, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[0]", {2, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[1]", {3, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[2]", {4, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEdb_A[3]", {5, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[0]", {6, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[1]", {7, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[2]", {8, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_TailEnd_A[3]", {9, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[0]", {10, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[1]", {11, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[2]", {12, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadSdp_A[3]", {13, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[0]", {14, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[1]", {15, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[2]", {16, RX_TILE_TXCLK}},
{"RX_PERF_RXP_RX_HeadStp_A[3]", {17, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_nullified_tlp_A", {18, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_valid_crc_A", {19, RX_TILE_TXCLK}},
{"RX_PERF_RXCRC_invalid_crc_A", {20, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_vendor_type1_A", {21, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_vendor_type0_A", {22, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_set_slot_power_limit_A", {23, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_unlock_A", {24, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_fatal_A", {25, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_nonfatal_A", {26, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_err_corr_A", {27, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pme_to_ack_A", {28, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pme_turn_off_A", {29, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pm_pme_A", {30, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_pm_active_state_nak_A", {31, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intd_A", {32, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intc_A", {33, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_intb_A", {34, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_deassert_inta_A", {35, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intd_A", {36, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intc_A", {37, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_intb_A", {38, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_assert_inta_A", {39, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_valid_A", {40, RX_TILE_TXCLK}},
{"RX_PERF_RMSG_unsupported_A", {41, RX_TILE_TXCLK}},
{"RX_PERF_RCB_unexpected_cpl_A", {42, RX_TILE_TXCLK}},
{"RX_PERF_RCB_timeout_cpl_A", {43, RX_TILE_TXCLK}},
{"RX_PERF_HDS_tlphdrvalid_A", {44, RX_TILE_TXCLK}},
{"RX_PERF_HDS_tlpdatavalid_A", {45, RX_TILE_TXCLK}},
{"RX_PERF_GAN_bad_tlp_A", {46, RX_TILE_TXCLK}},
{"RX_PERF_GAN_nak_A", {47, RX_TILE_TXCLK}},
{"RX_PERF_GAN_ack_A", {48, RX_TILE_TXCLK}},
{"RX_PERF_FE_unsupported_req_A", {49, RX_TILE_TXCLK}},
{"RX_PERF_FE_unsupported_cpl_A", {50, RX_TILE_TXCLK}},
{"RX_PERF_FE_unexpected_cpl_A", {51, RX_TILE_TXCLK}},
{"RX_PERF_FE_poisoned_tlp_A", {52, RX_TILE_TXCLK}},
{"RX_PERF_FE_poisoned_cpl_A", {53, RX_TILE_TXCLK}},
{"RX_PERF_FE_malformed_tlp_A", {54, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_abort_A", {55, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MSG_A", {56, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_CFG_WR_A", {57, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_CFG_RD_A", {58, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_IO_WR_A", {59, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_IO_RD_A", {60, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MEM_WR_A", {61, RX_TILE_TXCLK}},
{"RX_PERF_FE_request_MEM_RD_A", {62, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_gt16_A", {63, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_9to16_A", {64, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_5to8_A", {65, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_2to4_A", {66, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_MST_1_A", {67, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_gt32_A", {68, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_17to32_A", {69, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_9to16_A", {70, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_5to8_A", {71, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_2to4_A", {72, RX_TILE_TXCLK}},
{"RX_PERF_FE_length_SLV_1_A", {73, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_CA_A", {74, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_CRS_A", {75, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_UR_A", {76, RX_TILE_TXCLK}},
{"RX_PERF_FE_cpl_status_SC_A", {77, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_active_state_request_l1_A", {78, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_request_ack_A", {79, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_enter_l23_A", {80, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_pm_enter_l1_A", {81, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_error_A", {82, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_crc_err_A", {83, RX_TILE_TXCLK}},
{"SB_PERF_FCC_npd_0", {84, RX_TILE_TXCLK}},
{"SB_PERF_FCC_pd_0", {85, RX_TILE_TXCLK}},
{"SB_PERF_FCC_nph_0", {86, RX_TILE_TXCLK}},
{"SB_PERF_FCC_ph_0", {87, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_rd_hdr_0", {88, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_rd_hdr_0", {89, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_wr_hdr_0", {90, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_wr_hdr_0", {91, RX_TILE_TXCLK}},
{"SB_PERF_fail_crc_data_0", {92, RX_TILE_TXCLK}},
{"SB_PERF_pass_crc_data_0", {93, RX_TILE_TXCLK}},
{"SB_PERF_invalid_crc_0", {94, RX_TILE_TXCLK}},
{"SB_PERF_valid_crc_0", {95, RX_TILE_TXCLK}},
{"SB_PERF_rd_hdr_WEN_0", {96, RX_TILE_TXCLK}},
{"SB_PERF_wr_hdr_WEN_0", {97, RX_TILE_TXCLK}},
{"SB_PERF_data_WEN_0", {98, RX_TILE_TXCLK}},
{"SB_PERF_non_post_rd_from_FE", {99, RX_TILE_TXCLK}},
{"SB_PERF_non_post_wr_from_FE", {100, RX_TILE_TXCLK}},
{"SB_PERF_post_req_from_FE", {101, RX_TILE_TXCLK}},
{"SB_PERF_non_post_rd_from_FE_0", {102, RX_TILE_TXCLK}},
{"SB_PERF_non_post_wr_from_FE_0", {103, RX_TILE_TXCLK}},
{"SB_PERF_post_req_from_FE_0", {104, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_nak_A", {111, RX_TILE_TXCLK}},
{"RX_PERF_DLLP_ack_A", {112, RX_TILE_TXCLK}},
{"RX_PERF_allErrors_A", {113, RX_TILE_TXCLK}},
{"perf_PG_COUNT", {175, RX_TILE_TXCLK}},
{"perf_NOT_POWER_GATED", {176, RX_TILE_TXCLK}},
{"perf_POWER_GATED", {177, RX_TILE_TXCLK}},
{"SB_PERF_non_post_rd_to_HI", {2, RX_TILE_SCLK}},
{"SB_PERF_non_post_wr_to_HI", {3, RX_TILE_SCLK}},
{"SB_PERF_post_req_to_HI", {4, RX_TILE_SCLK}},
{"SB_PERF_non_post_rd_to_HI_0", {5, RX_TILE_SCLK}},
{"SB_PERF_non_post_wr_to_HI_0", {6, RX_TILE_SCLK}},
{"SB_PERF_post_req_to_HI_0", {7, RX_TILE_SCLK}},
{"SB_PERF_rd_hdr_REN_0", {8, RX_TILE_SCLK}},
{"SB_PERF_wr_hdr_REN_0", {9, RX_TILE_SCLK}},
{"SB_PERF_data_REN_0", {10, RX_TILE_SCLK}},
{"SB_PERF_rd_hdr_empty_0", {11, RX_TILE_SCLK}},
{"SB_PERF_wr_hdr_empty_0", {12, RX_TILE_SCLK}},
{"SB_PERF_data_empty_0", {13, RX_TILE_SCLK}},
{"CI_PERF_slv_total128BRdCpl", {29, RX_TILE_SCLK}},
{"CI_PERF_slv_total32BMemRdTx", {30, RX_TILE_SCLK}},
{"CI_PERF_slv_total64BMemRdTx", {31, RX_TILE_SCLK}},
{"CI_PERF_slv_total16BMemWrTx", {32, RX_TILE_SCLK}},
{"CI_PERF_slv_total32BMemWrTx", {33, RX_TILE_SCLK}},
{"CI_PERF_slv_total64BMemWrTx", {34, RX_TILE_SCLK}},
{"CI_PERF_slv_totalTx", {35, RX_TILE_SCLK}},
{"CI_PERF_slv_stallGrantGen", {36, RX_TILE_SCLK}},
{"CI_PERF_slv_totalGrant", {37, RX_TILE_SCLK}},
{"CI_PERF_slv_txPending", {38, RX_TILE_SCLK}},
{"CI_PERF_slv_numMemRdLT32B", {39, RX_TILE_SCLK}},
{"CI_PERF_slv_numMemRdLT16B", {40, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemTx", {41, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemRdTx", {42, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemWrTx", {43, RX_TILE_SCLK}},
{"CI_PERF_slv_numGrant0", {44, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow_ns0", {45, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow_ns0", {46, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow_s0", {47, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow_s0", {48, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow0", {49, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow0", {50, RX_TILE_SCLK}},
{"CI_PERF_slv_npNotAccepted_ns0", {51, RX_TILE_SCLK}},
{"CI_PERF_slv_npNotAccepted_s0", {52, RX_TILE_SCLK}},
{"CI_PERF_slv_num128BRdCpl0", {53, RX_TILE_SCLK}},
{"CI_PERF_slv_num32BMemRdTx0", {54, RX_TILE_SCLK}},
{"CI_PERF_slv_num64BMemRdTx0", {55, RX_TILE_SCLK}},
{"CI_PERF_slv_num16BMemWrTx0", {56, RX_TILE_SCLK}},
{"CI_PERF_slv_num32BMemWrTx0", {57, RX_TILE_SCLK}},
{"CI_PERF_slv_num64BMemWrTx0", {58, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_Bandwidth0", {59, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_Bandwidth0", {60, RX_TILE_SCLK}},
{"TX_PERF_S_RCLK_s_tag_buf_empty", {61, RX_TILE_SCLK}},
{"P_request_latency_500ns_or_more", {62, RX_TILE_SCLK}},
{"P_request_latency_250_to_500ns", {63, RX_TILE_SCLK}},
{"P_request_latency_100_to_250ns", {64, RX_TILE_SCLK}},
{"P_request_latency_100ns_or_less", {65, RX_TILE_SCLK}},
{"NP_request_latency_500ns_or_more", {66, RX_TILE_SCLK}},
{"NP_request_latency_250_to_500ns", {67, RX_TILE_SCLK}},
{"NP_request_latency_100_to_250ns", {68, RX_TILE_SCLK}},
{"NP_request_latency_100ns_or_less", {69, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_cpl_slot[0]", {70, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_tag[0]", {71, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_d_credit[0]", {72, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_h_credit[0]", {73, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_tag[0]", {74, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_d_credit[0]", {75, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_h_credit[0]", {76, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_no_tags_q", {77, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_data_credits_q", {78, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_req_credits_q", {79, RX_TILE_SCLK}},
{"CISLV_PERF_no_cpl_slots_q[0]", {80, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_no_tags_q", {81, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_data_credits_q", {82, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_req_credits_q", {83, RX_TILE_SCLK}}
};
{"SB_PERF_non_post_rd_to_HI", {2, RX_TILE_SCLK}},
{"SB_PERF_non_post_wr_to_HI", {3, RX_TILE_SCLK}},
{"SB_PERF_post_req_to_HI", {4, RX_TILE_SCLK}},
{"SB_PERF_non_post_rd_to_HI_0", {5, RX_TILE_SCLK}},
{"SB_PERF_non_post_wr_to_HI_0", {6, RX_TILE_SCLK}},
{"SB_PERF_post_req_to_HI_0", {7, RX_TILE_SCLK}},
{"SB_PERF_rd_hdr_REN_0", {8, RX_TILE_SCLK}},
{"SB_PERF_wr_hdr_REN_0", {9, RX_TILE_SCLK}},
{"SB_PERF_data_REN_0", {10, RX_TILE_SCLK}},
{"SB_PERF_rd_hdr_empty_0", {11, RX_TILE_SCLK}},
{"SB_PERF_wr_hdr_empty_0", {12, RX_TILE_SCLK}},
{"SB_PERF_data_empty_0", {13, RX_TILE_SCLK}},
{"CI_PERF_slv_total128BRdCpl", {29, RX_TILE_SCLK}},
{"CI_PERF_slv_total32BMemRdTx", {30, RX_TILE_SCLK}},
{"CI_PERF_slv_total64BMemRdTx", {31, RX_TILE_SCLK}},
{"CI_PERF_slv_total16BMemWrTx", {32, RX_TILE_SCLK}},
{"CI_PERF_slv_total32BMemWrTx", {33, RX_TILE_SCLK}},
{"CI_PERF_slv_total64BMemWrTx", {34, RX_TILE_SCLK}},
{"CI_PERF_slv_totalTx", {35, RX_TILE_SCLK}},
{"CI_PERF_slv_stallGrantGen", {36, RX_TILE_SCLK}},
{"CI_PERF_slv_totalGrant", {37, RX_TILE_SCLK}},
{"CI_PERF_slv_txPending", {38, RX_TILE_SCLK}},
{"CI_PERF_slv_numMemRdLT32B", {39, RX_TILE_SCLK}},
{"CI_PERF_slv_numMemRdLT16B", {40, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemTx", {41, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemRdTx", {42, RX_TILE_SCLK}},
{"CI_PERF_slv_totalMemWrTx", {43, RX_TILE_SCLK}},
{"CI_PERF_slv_numGrant0", {44, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow_ns0", {45, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow_ns0", {46, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow_s0", {47, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow_s0", {48, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntOverFlow0", {49, RX_TILE_SCLK}},
{"CI_PERF_slv_portCntUnderFlow0", {50, RX_TILE_SCLK}},
{"CI_PERF_slv_npNotAccepted_ns0", {51, RX_TILE_SCLK}},
{"CI_PERF_slv_npNotAccepted_s0", {52, RX_TILE_SCLK}},
{"CI_PERF_slv_num128BRdCpl0", {53, RX_TILE_SCLK}},
{"CI_PERF_slv_num32BMemRdTx0", {54, RX_TILE_SCLK}},
{"CI_PERF_slv_num64BMemRdTx0", {55, RX_TILE_SCLK}},
{"CI_PERF_slv_num16BMemWrTx0", {56, RX_TILE_SCLK}},
{"CI_PERF_slv_num32BMemWrTx0", {57, RX_TILE_SCLK}},
{"CI_PERF_slv_num64BMemWrTx0", {58, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_Bandwidth0", {59, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_Bandwidth0", {60, RX_TILE_SCLK}},
{"TX_PERF_S_RCLK_s_tag_buf_empty", {61, RX_TILE_SCLK}},
{"P_request_latency_500ns_or_more", {62, RX_TILE_SCLK}},
{"P_request_latency_250_to_500ns", {63, RX_TILE_SCLK}},
{"P_request_latency_100_to_250ns", {64, RX_TILE_SCLK}},
{"P_request_latency_100ns_or_less", {65, RX_TILE_SCLK}},
{"NP_request_latency_500ns_or_more", {66, RX_TILE_SCLK}},
{"NP_request_latency_250_to_500ns", {67, RX_TILE_SCLK}},
{"NP_request_latency_100_to_250ns", {68, RX_TILE_SCLK}},
{"NP_request_latency_100ns_or_less", {69, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_cpl_slot[0]", {70, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_tag[0]", {71, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_d_credit[0]", {72, RX_TILE_SCLK}},
{"CI_PERF_slv_MemRd_wait_for_h_credit[0]", {73, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_tag[0]", {74, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_d_credit[0]", {75, RX_TILE_SCLK}},
{"CI_PERF_slv_MemWr_wait_for_h_credit[0]", {76, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_no_tags_q", {77, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_data_credits_q", {78, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC1_req_credits_q", {79, RX_TILE_SCLK}},
{"CISLV_PERF_no_cpl_slots_q[0]", {80, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_no_tags_q", {81, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_data_credits_q", {82, RX_TILE_SCLK}},
{"CISLV_PERF_no_VC0_req_credits_q", {83, RX_TILE_SCLK}}};
}
} // namespace PCIE_MI200
#endif
projects/rocprofiler/src/core/counters/mmio/perfmon.h
+1 -1
Bestand weergeven
@@ -42,6 +42,6 @@ class PerfMon {
std::vector<std::string> counter_names_;
};
} // namespace rocprofiler
} // namespace rocprofiler
#endif
projects/rocprofiler/src/core/gpu_command.cpp
+19 -17
Bestand weergeven
@@ -31,10 +31,8 @@ THE SOFTWARE.
#include "util/hsa_rsrc_factory.h"
namespace rocprofiler {
size_t CreateGpuCommand(gpu_cmd_op_t op,
const rocprofiler::util::AgentInfo* agent_info,
packet_t* command,
const size_t& slot_count) {
size_t CreateGpuCommand(gpu_cmd_op_t op, const rocprofiler::util::AgentInfo* agent_info,
packet_t* command, const size_t& slot_count) {
if (op >= NUMBER_GPU_CMD_OP) EXC_RAISING(HSA_STATUS_ERROR, "bad op value (" << op << ")");
const bool is_legacy = (strncmp(agent_info->name, "gfx8", 4) == 0);
@@ -49,14 +47,15 @@ size_t CreateGpuCommand(gpu_cmd_op_t op,
profile.agent = agent_info->dev_id;
// Query for cmd buffer size
hsa_ven_amd_aqlprofile_info_type_t info_type =
(hsa_ven_amd_aqlprofile_info_type_t)((int)HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD + (int)op);
hsa_status_t status = hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, NULL);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "get_info(ENABLE_CMD ).size exc, op(" << int(op) << ")");
(hsa_ven_amd_aqlprofile_info_type_t)((int)HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD + (int)op);
hsa_status_t status =
hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, NULL);
if (status != HSA_STATUS_SUCCESS)
EXC_RAISING(status, "get_info(ENABLE_CMD ).size exc, op(" << int(op) << ")");
if (profile.command_buffer.size == 0) EXC_RAISING(status, "get_info(ENABLE_CMD).size == 0");
// Allocate cmd buffer
const size_t aligment_mask = 0x100 - 1;
profile.command_buffer.ptr =
hsa_rsrc->AllocateSysMemory(agent_info, profile.command_buffer.size);
profile.command_buffer.ptr = hsa_rsrc->AllocateSysMemory(agent_info, profile.command_buffer.size);
if ((reinterpret_cast<uintptr_t>(profile.command_buffer.ptr) & aligment_mask) != 0) {
EXC_RAISING(status, "profile.command_buffer.ptr bad alignment");
}
@@ -66,15 +65,18 @@ size_t CreateGpuCommand(gpu_cmd_op_t op,
packet_t packet{};
// Query for cmd buffer data
status = hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, &packet);
status =
hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, &packet);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "get_info(ENABLE_CMD).data exc");
// Check for legacy GFXIP
status = hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_legacy_get_pm4(&packet, command);
if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
if (status != HSA_STATUS_SUCCESS)
AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
} else {
// Query for cmd buffer data
status = hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, command);
status =
hsa_rsrc->AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(&profile, info_type, command);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "get_info(ENABLE_CMD).data exc");
}
@@ -91,15 +93,14 @@ struct gpu_cmd_key_t {
uint32_t node_id;
};
struct gpu_cmd_fncomp_t {
bool operator() (const gpu_cmd_key_t& a, const gpu_cmd_key_t& b) const {
bool operator()(const gpu_cmd_key_t& a, const gpu_cmd_key_t& b) const {
return (a.op < b.op) || ((a.op == b.op) && (a.node_id < b.node_id));
}
};
typedef std::map<gpu_cmd_key_t, gpu_cmd_entry_t, gpu_cmd_fncomp_t> gpu_cmd_map_t;
size_t GetGpuCommand(gpu_cmd_op_t op,
const rocprofiler::util::AgentInfo* agent_info,
packet_t** command_out) {
size_t GetGpuCommand(gpu_cmd_op_t op, const rocprofiler::util::AgentInfo* agent_info,
packet_t** command_out) {
thread_local gpu_cmd_map_t map;
// Getting NUMA node id
@@ -112,7 +113,8 @@ size_t GetGpuCommand(gpu_cmd_op_t op,
auto ret = map.insert({gpu_cmd_key_t{op, node_id}, gpu_cmd_entry_t{}});
gpu_cmd_map_t::iterator it = ret.first;
if (ret.second) {
it->second.size = CreateGpuCommand(op, agent_info, it->second.command, Profile::LEGACY_SLOT_SIZE_PKT);
it->second.size =
CreateGpuCommand(op, agent_info, it->second.command, Profile::LEGACY_SLOT_SIZE_PKT);
}
*command_out = it->second.command;
projects/rocprofiler/src/core/gpu_command.h
+3 -6
Bestand weergeven
@@ -37,9 +37,8 @@ enum gpu_cmd_op_t {
NUMBER_GPU_CMD_OP
};
size_t GetGpuCommand(gpu_cmd_op_t op,
const rocprofiler::util::AgentInfo* agent_info,
packet_t** command_out);
size_t GetGpuCommand(gpu_cmd_op_t op, const rocprofiler::util::AgentInfo* agent_info,
packet_t** command_out);
static inline size_t IssueGpuCommand(gpu_cmd_op_t op,
const rocprofiler::util::AgentInfo* agent_info,
@@ -55,9 +54,7 @@ static inline size_t IssueGpuCommand(gpu_cmd_op_t op,
return HSA_STATUS_SUCCESS;
}
static inline size_t IssueGpuCommand(gpu_cmd_op_t op,
hsa_agent_t agent,
hsa_queue_t* queue) {
static inline size_t IssueGpuCommand(gpu_cmd_op_t op, hsa_agent_t agent, hsa_queue_t* queue) {
rocprofiler::util::HsaRsrcFactory* hsa_rsrc = &rocprofiler::util::HsaRsrcFactory::Instance();
const rocprofiler::util::AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(agent);
return IssueGpuCommand(op, agent_info, queue);
projects/rocprofiler/src/core/group_set.h
+25 -26
Bestand weergeven
@@ -55,31 +55,30 @@ struct block_status_t {
// Metrics set class
class MetricsGroup {
public:
public:
// Info map type
typedef std::map<std::string, const Metric*> info_map_t;
// Blocks map type
typedef std::map<block_des_t, block_status_t, lt_block_des> blocks_map_t;
MetricsGroup(const util::AgentInfo* agent_info) :
agent_info_(agent_info)
{
MetricsGroup(const util::AgentInfo* agent_info) : agent_info_(agent_info) {
metrics_ = MetricsDict::Create(agent_info);
if (metrics_ == NULL) EXC_RAISING(HSA_STATUS_ERROR, "MetricsDict create failed");
}
void Print(FILE* file) const {
for (const Metric* metric : metrics_vec_) {
fprintf(file, " %s", metric->GetName().c_str()); fflush(stdout);
fprintf(file, " %s", metric->GetName().c_str());
fflush(stdout);
}
fprintf(file, "\n"); fflush(stdout);
fprintf(file, "\n");
fflush(stdout);
}
static const Metric* GetMetric(const MetricsDict* metrics, const std::string& name) {
// Metric object
const Metric* metric = metrics->Get(name);
if (metric == NULL)
EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not found");
if (metric == NULL) EXC_RAISING(HSA_STATUS_ERROR, "input metric '" << name << "' is not found");
return metric;
}
@@ -95,9 +94,7 @@ class MetricsGroup {
}
// Add metric
bool AddMetric(const rocprofiler_feature_t* info) {
return AddMetric(GetMetric(metrics_, info));
}
bool AddMetric(const rocprofiler_feature_t* info) { return AddMetric(GetMetric(metrics_, info)); }
bool AddMetric(const Metric* metric) {
// Blocks utilization delta
@@ -125,8 +122,9 @@ class MetricsGroup {
query.events = event;
uint32_t block_counters;
hsa_status_t status = util::HsaRsrcFactory::Instance().AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(
&query, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS, &block_counters);
hsa_status_t status =
util::HsaRsrcFactory::Instance().AqlProfileApi()->hsa_ven_amd_aqlprofile_get_info(
&query, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS, &block_counters);
if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "get block_counters info");
block_status.max_counters = block_counters;
}
@@ -141,7 +139,8 @@ class MetricsGroup {
metrics_vec_.push_back(metric);
info_map_[metric->GetName()] = metric;
for (const counter_t* counter : counters_vec) {
if (info_map_.find(counter->name) == info_map_.end()) info_map_[counter->name] = NewCounterInfo(counter->name);
if (info_map_.find(counter->name) == info_map_.end())
info_map_[counter->name] = NewCounterInfo(counter->name);
}
for (const auto& entry : blocks_delta) {
blocks_map_[entry.first] = entry.second;
@@ -150,10 +149,8 @@ class MetricsGroup {
return true;
}
private:
const Metric* NewCounterInfo(const std::string& name) const {
return GetMetric(metrics_, name);
}
private:
const Metric* NewCounterInfo(const std::string& name) const { return GetMetric(metrics_, name); }
// Agent info
const util::AgentInfo* const agent_info_;
@@ -169,10 +166,10 @@ class MetricsGroup {
// Metrics groups class
class MetricsGroupSet {
public:
MetricsGroupSet(const util::AgentInfo* agent_info, const rocprofiler_feature_t* info_array, const uint32_t info_count) :
agent_info_(agent_info)
{
public:
MetricsGroupSet(const util::AgentInfo* agent_info, const rocprofiler_feature_t* info_array,
const uint32_t info_count)
: agent_info_(agent_info) {
metrics_ = MetricsDict::Create(agent_info);
if (metrics_ == NULL) EXC_RAISING(HSA_STATUS_ERROR, "MetricsDict create failed");
Initialize(info_array, info_count);
@@ -186,12 +183,13 @@ class MetricsGroupSet {
void Print(FILE* file) const {
for (const auto* group : groups_) {
fprintf(stdout, " pmc : "); fflush(stdout);
fprintf(stdout, " pmc : ");
fflush(stdout);
group->Print(file);
}
}
private:
private:
void Initialize(const rocprofiler_feature_t* info_array, const uint32_t info_count) {
std::multimap<uint32_t, const Metric*, std::greater<uint32_t> > input_metrics;
for (unsigned i = 0; i < info_count; ++i) {
@@ -202,7 +200,8 @@ class MetricsGroupSet {
input_metrics.insert({counters_num, metric});
if (MetricsGroup(agent_info_).AddMetric(metric) == false) {
AQL_EXC_RAISING(HSA_STATUS_ERROR, "Metric '" << metric->GetName() << "' doesn't fit in one group");
AQL_EXC_RAISING(HSA_STATUS_ERROR,
"Metric '" << metric->GetName() << "' doesn't fit in one group");
}
}
#if 0
@@ -239,4 +238,4 @@ class MetricsGroupSet {
} // namespace rocprofiler
#endif // SRC_CORE_GROUP_SET_H_
#endif // SRC_CORE_GROUP_SET_H_
projects/rocprofiler/src/core/hardware/hsa_info.cpp
+14 -19
Bestand weergeven
@@ -62,33 +62,28 @@ 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)
{
rocprofiler::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)
{
rocprofiler::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(saurabh, giovanni): Remove this in 5.7
if (table->hsa_agent_get_info_fn(agent,
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_XCC), &xcc_num_) != HSA_STATUS_SUCCESS) {
xcc_num_ = 1;
if (table->hsa_agent_get_info_fn(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_XCC),
&xcc_num_) != HSA_STATUS_SUCCESS) {
xcc_num_ = 1;
}
}
projects/rocprofiler/src/core/hsa/hsa_common.cpp
+5 -7
Bestand weergeven
@@ -33,8 +33,8 @@ Agent::AgentInfo& GetAgentInfo(decltype(hsa_agent_t::handle) handle) {
if (agent_info_map.find(handle) != agent_info_map.end()) {
return agent_info_map.at(handle);
} else {
std::cerr << std::string("Error: Can't find Agent with handle(") << std::to_string(handle) <<
") in this system" << std::endl;
std::cerr << std::string("Error: Can't find Agent with handle(") << std::to_string(handle)
<< ") in this system" << std::endl;
abort();
}
}
@@ -49,9 +49,7 @@ void SetAgentInfo(decltype(hsa_agent_t::handle) handle, const Agent::AgentInfo&
}
}
std::vector<hsa_agent_t>& GetCPUAgentList() {
return cpu_agents_list;
}
std::vector<hsa_agent_t>& GetCPUAgentList() { return cpu_agents_list; }
hsa_agent_t GetAgentByIndex(uint64_t agent_index) {
std::lock_guard<std::mutex> lock(agents_map_lock);
@@ -60,8 +58,8 @@ hsa_agent_t GetAgentByIndex(uint64_t agent_index) {
return hsa_agent_t{agent_info.second.getHandle()};
}
}
std::cerr << std::string("Error: Can't find Agent with Index(") << std::to_string(agent_index) <<
") in this system" << std::endl;
std::cerr << std::string("Error: Can't find Agent with Index(") << std::to_string(agent_index)
<< ") in this system" << std::endl;
abort();
}
projects/rocprofiler/src/core/hsa/hsa_support.h
+1 -1
Bestand weergeven
@@ -95,7 +95,7 @@ namespace rocprofiler {
namespace hsa_support {
void Initialize(HsaApiTable* Table);
hsa_status_t hsa_iterate_agents_cb(hsa_agent_t agent, void *data);
hsa_status_t hsa_iterate_agents_cb(hsa_agent_t agent, void* data);
void Finalize();
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp
+8 -6
Bestand weergeven
@@ -181,7 +181,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
// TODO: validate needs to be called on each events_list[i]
// Validating the events array for the specified gpu agent
if(events_list.size() > 0) {
if (events_list.size() > 0) {
bool validate_event_result;
status =
hsa_ven_amd_aqlprofile_validate_event(gpu_agent, &events_list[0], &validate_event_result);
@@ -234,9 +234,10 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
}
}
for(auto& cname : counter_names) {
if(cname.compare("KERNEL_DURATION")==0) {
rocprofiler::Metric* metric = const_cast<rocprofiler::Metric*>(metricsDict[gpu_agent.handle]->Get(cname));
for (auto& cname : counter_names) {
if (cname.compare("KERNEL_DURATION") == 0) {
rocprofiler::Metric* metric =
const_cast<rocprofiler::Metric*>(metricsDict[gpu_agent.handle]->Get(cname));
if (metric == nullptr) std::cout << cname << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
}
@@ -315,7 +316,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
if(context->events_list.size() > 0) {
if (context->events_list.size() > 0) {
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
// CHECK_HSA_STATUS("Error: Getting Buffers Size", status);
@@ -510,7 +511,8 @@ uint8_t* AllocateLocalMemory(size_t size, hsa_amd_memory_pool_t* gpu_pool) {
return ptr;
}
hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile, size_t att_buffer_size) {
hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile,
size_t att_buffer_size) {
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
profile->command_buffer.ptr =
AllocateSysMemory(gpu_agent, profile->command_buffer.size, &agentInfo.cpu_pool);
projects/rocprofiler/src/core/hsa/queues/queue.cpp
+41 -41
Bestand weergeven
@@ -435,16 +435,18 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
pending->session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
}
if (pending->counters_count > 0) {
if (xcc_id == 0 && pending->context && pending->context->metrics_list.size() > 0 && pending->profile) // call to GetCounterData() is required only once for a dispatch
if (xcc_id == 0 && pending->context && pending->context->metrics_list.size() > 0 &&
pending->profile) // call to GetCounterData() is required only once for a dispatch
rocprofiler::metrics::GetCounterData(pending->profile, queue_info_session->agent,
pending->context->results_list);
if (is_individual_xcc_mode)
rocprofiler::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, pending->context->results_list, pending->context->results_map,
pending->context->metrics_list, time.end-time.start);
pending->context->metrics_list, time.end - time.start);
else
rocprofiler::metrics::GetMetricsData(pending->context->results_map,
pending->context->metrics_list, time.end-time.start);
pending->context->metrics_list,
time.end - time.start);
AddRecordCounters(&record, pending);
} else {
if (session->FindBuffer(pending->buffer_id)) {
@@ -652,8 +654,8 @@ void CheckNeededProfileConfigs() {
att_counters_names = filter->GetCounterData();
kernel_profile_names = std::get<std::vector<std::string>>(
filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES));
kernel_profile_dispatch_ids = std::get<std::vector<uint64_t>>(
filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS));
kernel_profile_dispatch_ids =
std::get<std::vector<uint64_t>>(filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS));
} else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
is_pc_sampling_collection_mode = true;
}
@@ -685,23 +687,20 @@ std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, i
auto& kdispatch = static_cast<const hsa_kernel_dispatch_packet_s*>(packets)[i];
// If Dispatch IDs specified, profile based on dispatch ID
for (auto id : kernel_profile_dispatch_ids)
b_profile_this_object |= id == current_writer_id;
for (auto id : kernel_profile_dispatch_ids) b_profile_this_object |= id == current_writer_id;
try {
// Can throw
const std::string& kernel_name = ksymbols->at(kdispatch.kernel_object);
// If no filters specified, auto profile this kernel
if (kernel_profile_names.size() == 0 &&
kernel_profile_dispatch_ids.size() == 0 &&
if (kernel_profile_names.size() == 0 && kernel_profile_dispatch_ids.size() == 0 &&
kernel_name.find("__amd_rocclr_") == std::string::npos)
b_profile_this_object = true;
b_profile_this_object = true;
// Try to match the mangled kernel name with given matches in input.txt
// We want to initiate att profiling if a match exists
for (const std::string& kernel_matches : kernel_profile_names)
if (kernel_name.find(kernel_matches) != std::string::npos)
b_profile_this_object = true;
if (kernel_name.find(kernel_matches) != std::string::npos) b_profile_this_object = true;
} catch (...) {
printf("Warning: Unknown name for object %lu\n", kdispatch.kernel_object);
}
@@ -711,17 +710,13 @@ std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, i
can_profile_packet.push_back(b_profile_this_object);
}
// If we're going to skip all packets, need to update writer ID
if (!b_can_profile_anypacket)
WRITER_ID.store(current_writer_id, std::memory_order_release);
if (!b_can_profile_anypacket) WRITER_ID.store(current_writer_id, std::memory_order_release);
return {can_profile_packet, b_can_profile_anypacket};
}
hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(
Packet::packet_t& start_packet,
Packet::packet_t& stop_packet,
Queue& queue_info,
Agent::AgentInfo& agentInfo
) {
hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packet,
Packet::packet_t& stop_packet, Queue& queue_info,
Agent::AgentInfo& agentInfo) {
std::vector<hsa_ven_amd_aqlprofile_parameter_t> att_params;
int num_att_counters = 0;
uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE;
@@ -731,15 +726,16 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(
case ROCPROFILER_ATT_PERFCOUNTER_NAME:
break;
case ROCPROFILER_ATT_BUFFER_SIZE:
att_buffer_size = std::max(96l<<10l, std::min(int64_t(param.value)<<20l, (1l<<32l)-(3l<<20)));
break; // Clip to [96KB, 4GB)
att_buffer_size =
std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20)));
break; // Clip to [96KB, 4GB)
case ROCPROFILER_ATT_PERFCOUNTER:
num_att_counters += 1;
break;
default:
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(param.parameter_name)),
param.value});
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(param.parameter_name)),
param.value});
}
}
@@ -760,22 +756,21 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(
printf("Only events from the SQ block can be selected for ATT.");
exit(1);
}
att_params.push_back({static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(
int(ROCPROFILER_ATT_PERFCOUNTER)),
event.counter_id | (event.counter_id ? (0xF << 24) : 0)});
att_params.push_back(
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)),
event.counter_id | (event.counter_id ? (0xF << 24) : 0)});
num_att_counters += 1;
}
hsa_ven_amd_aqlprofile_parameter_t zero_perf = {
static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)),
0};
static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)), 0};
// Fill other perfcounters with 0's
for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf);
}
// Get the PM4 Packets using packets_generator
return Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
att_params, &start_packet, &stop_packet, att_buffer_size);
return Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), att_params,
&start_packet, &stop_packet, att_buffer_size);
}
/**
@@ -866,14 +861,16 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
record_id);
if (session_data_count > 0 && profile.second) {
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id,
profile.first, session_data_count, profile.second, kernel_properties,
(uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id);
writer_id, record_id, original_packet.completion_signal,
dispatch_packet.completion_signal, session_id, buffer_id, profile.first,
session_data_count, profile.second, kernel_properties, (uint32_t)syscall(__NR_gettid),
user_pkt_index, correlation_id);
} else {
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id,
nullptr, session_data_count, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid),
user_pkt_index, correlation_id);
writer_id, record_id, original_packet.completion_signal,
dispatch_packet.completion_signal, session_id, buffer_id, nullptr, session_data_count,
nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index,
correlation_id);
}
}
@@ -893,7 +890,8 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
CreateSignal(0, &interrupt_signal);
// Adding Stop and Read PM4 Packets
if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 && profile.first && profile.first->stop_packet) {
if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 &&
profile.first && profile.first->stop_packet) {
hsa_signal_t dummy_signal{};
profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
@@ -937,7 +935,8 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
bool can_profile_anypacket = false;
std::vector<bool> can_profile_packet;
std::tie(can_profile_packet, can_profile_anypacket) = GetAllowedProfilesList(packets, pkt_count);
std::tie(can_profile_packet, can_profile_anypacket) =
GetAllowedProfilesList(packets, pkt_count);
if (!can_profile_anypacket) {
/* Write the original packets to the hardware if no patch will be profiled */
@@ -964,8 +963,9 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// increment writer ID for every packet
if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_TYPE+HSA_PACKET_HEADER_WIDTH_TYPE-1) == HSA_PACKET_TYPE_KERNEL_DISPATCH)
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) ==
HSA_PACKET_TYPE_KERNEL_DISPATCH)
writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
continue;
}
projects/rocprofiler/src/core/hsa_interceptor.h
+99 -158
Bestand weergeven
@@ -37,33 +37,37 @@ SOFTWARE.
#include "util/exception.h"
#include "util/hsa_rsrc_factory.h"
#define HSA_RT(call) \
do { \
const hsa_status_t status = call; \
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, #call); \
} while(0)
#define IS_HSA_CALLBACK(ID) \
const auto __id = ID; (void)__id; \
void *__arg = arg_.load(); (void)__arg; \
rocprofiler_hsa_callback_fun_t __callback = \
(ID == ROCPROFILER_HSA_CB_ID_ALLOCATE) ? callbacks_.allocate: \
(ID == ROCPROFILER_HSA_CB_ID_DEVICE) ? callbacks_.device: \
(ID == ROCPROFILER_HSA_CB_ID_MEMCOPY) ? callbacks_.memcopy: \
(ID == ROCPROFILER_HSA_CB_ID_SUBMIT) ? callbacks_.submit: \
(ID == ROCPROFILER_HSA_CB_ID_KSYMBOL) ? callbacks_.ksymbol: \
callbacks_.codeobj; \
if ((__callback != NULL) && (recursion_ == false))
#define DO_HSA_CALLBACK \
do { \
recursion_ = true; \
__callback(__id, &data, __arg); \
recursion_ = false; \
#define HSA_RT(call) \
do { \
const hsa_status_t status = call; \
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, #call); \
} while (0)
#define ISSUE_HSA_CALLBACK(ID) \
do { IS_HSA_CALLBACK(ID) { DO_HSA_CALLBACK; } } while(0)
#define IS_HSA_CALLBACK(ID) \
const auto __id = ID; \
(void)__id; \
void* __arg = arg_.load(); \
(void)__arg; \
rocprofiler_hsa_callback_fun_t __callback = (ID == ROCPROFILER_HSA_CB_ID_ALLOCATE) \
? callbacks_.allocate \
: (ID == ROCPROFILER_HSA_CB_ID_DEVICE) ? callbacks_.device \
: (ID == ROCPROFILER_HSA_CB_ID_MEMCOPY) ? callbacks_.memcopy \
: (ID == ROCPROFILER_HSA_CB_ID_SUBMIT) ? callbacks_.submit \
: (ID == ROCPROFILER_HSA_CB_ID_KSYMBOL) ? callbacks_.ksymbol \
: callbacks_.codeobj; \
if ((__callback != NULL) && (recursion_ == false))
#define DO_HSA_CALLBACK \
do { \
recursion_ = true; \
__callback(__id, &data, __arg); \
recursion_ = false; \
} while (0)
#define ISSUE_HSA_CALLBACK(ID) \
do { \
IS_HSA_CALLBACK(ID) { DO_HSA_CALLBACK; } \
} while (0)
// Demangle C++ symbol name
static const char* cpp_demangle(const char* symname) {
@@ -74,15 +78,15 @@ static const char* cpp_demangle(const char* symname) {
}
namespace rocprofiler {
extern decltype(hsa_memory_allocate)* hsa_memory_allocate_fn;
extern decltype(hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
extern decltype(hsa_memory_copy)* hsa_memory_copy_fn;
extern decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
extern decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
extern decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
extern decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
extern decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
extern decltype(hsa_executable_destroy)* hsa_executable_destroy_fn;
extern decltype(::hsa_memory_allocate)* hsa_memory_allocate_fn;
extern decltype(::hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
extern decltype(::hsa_memory_copy)* hsa_memory_copy_fn;
extern decltype(::hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
extern decltype(::hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
extern decltype(::hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
extern decltype(::hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
extern decltype(::hsa_executable_freeze)* hsa_executable_freeze_fn;
extern decltype(::hsa_executable_destroy)* hsa_executable_destroy_fn;
class HsaInterceptor {
public:
@@ -95,10 +99,7 @@ class HsaInterceptor {
if (enable_) {
// Fetching AMD Loader HSA extension API
HSA_RT(hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER,
1,
sizeof(hsa_ven_amd_loader_1_01_pfn_t),
&LoaderApiTable));
HSA_EXTENSION_AMD_LOADER, 1, sizeof(hsa_ven_amd_loader_1_01_pfn_t), &LoaderApiTable));
// Saving original API functions
hsa_memory_allocate_fn = table->core_->hsa_memory_allocate_fn;
@@ -131,10 +132,7 @@ class HsaInterceptor {
}
private:
static hsa_status_t HSA_API MemoryAllocate(hsa_region_t region,
size_t size,
void** ptr)
{
static hsa_status_t HSA_API MemoryAllocate(hsa_region_t region, size_t size, void** ptr) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_memory_allocate_fn(region, size, ptr));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
@@ -150,11 +148,8 @@ class HsaInterceptor {
return status;
}
static hsa_status_t MemoryAssignAgent(
void *ptr,
hsa_agent_t agent,
hsa_access_permission_t access)
{
static hsa_status_t MemoryAssignAgent(void* ptr, hsa_agent_t agent,
hsa_access_permission_t access) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_memory_assign_agent_fn(ptr, agent, access));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
@@ -169,11 +164,7 @@ class HsaInterceptor {
}
// Spawn device allow access callback
static void DeviceCallback(
uint32_t num_agents,
const hsa_agent_t* agents,
const void* ptr)
{
static void DeviceCallback(uint32_t num_agents, const hsa_agent_t* agents, const void* ptr) {
for (const hsa_agent_t* agent_p = agents; agent_p < (agents + num_agents); ++agent_p) {
hsa_agent_t agent = *agent_p;
rocprofiler_hsa_callback_data_t data{};
@@ -188,17 +179,11 @@ class HsaInterceptor {
}
// Agent allow access callback 'hsa_amd_agents_allow_access'
static hsa_status_t AgentsAllowAccess(
uint32_t num_agents,
const hsa_agent_t* agents,
const uint32_t* flags,
const void* ptr)
{
static hsa_status_t AgentsAllowAccess(uint32_t num_agents, const hsa_agent_t* agents,
const uint32_t* flags, const void* ptr) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_amd_agents_allow_access_fn(num_agents, agents, flags, ptr));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
DeviceCallback(num_agents, agents, ptr);
}
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) { DeviceCallback(num_agents, agents, ptr); }
return status;
}
@@ -218,12 +203,8 @@ class HsaInterceptor {
return HSA_STATUS_SUCCESS;
}
static hsa_status_t MemoryPoolAllocate(
hsa_amd_memory_pool_t pool,
size_t size,
uint32_t flags,
void** ptr)
{
static hsa_status_t MemoryPoolAllocate(hsa_amd_memory_pool_t pool, size_t size, uint32_t flags,
void** ptr) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_amd_memory_pool_allocate_fn(pool, size, flags, ptr));
if (size != 0) {
@@ -232,8 +213,10 @@ class HsaInterceptor {
data.allocate.ptr = *ptr;
data.allocate.size = size;
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &data.allocate.segment));
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &data.allocate.global_flag));
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
&data.allocate.segment));
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
&data.allocate.global_flag));
DO_HSA_CALLBACK;
@@ -246,9 +229,7 @@ class HsaInterceptor {
}
return status;
}
static hsa_status_t MemoryPoolFree(
void* ptr)
{
static hsa_status_t MemoryPoolFree(void* ptr) {
hsa_status_t status = HSA_STATUS_SUCCESS;
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
rocprofiler_hsa_callback_data_t data{};
@@ -260,11 +241,7 @@ class HsaInterceptor {
return status;
}
static hsa_status_t MemoryCopy(
void *dst,
const void *src,
size_t size)
{
static hsa_status_t MemoryCopy(void* dst, const void* src, size_t size) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_memory_copy_fn(dst, src, size));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_MEMCOPY) {
@@ -277,17 +254,13 @@ class HsaInterceptor {
return status;
}
static hsa_status_t MemoryAsyncCopy(
void* dst, hsa_agent_t dst_agent, const void* src,
hsa_agent_t src_agent, size_t size,
uint32_t num_dep_signals,
const hsa_signal_t* dep_signals,
hsa_signal_t completion_signal)
{
static hsa_status_t MemoryAsyncCopy(void* dst, hsa_agent_t dst_agent, const void* src,
hsa_agent_t src_agent, size_t size, uint32_t num_dep_signals,
const hsa_signal_t* dep_signals,
hsa_signal_t completion_signal) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_amd_memory_async_copy_fn(
dst, dst_agent, src, src_agent, size,
num_dep_signals, dep_signals, completion_signal));
HSA_RT(hsa_amd_memory_async_copy_fn(dst, dst_agent, src, src_agent, size, num_dep_signals,
dep_signals, completion_signal));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_MEMCOPY) {
rocprofiler_hsa_callback_data_t data{};
data.memcopy.dst = dst;
@@ -298,14 +271,11 @@ class HsaInterceptor {
return status;
}
static hsa_status_t CodeObjectCallback(
hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object,
void* arg)
{
static hsa_status_t CodeObjectCallback(hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object, void* arg) {
const int free_flag = reinterpret_cast<long>(arg);
hsa_ven_amd_loader_code_object_storage_type_t storage_type =
HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE;
HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE;
int storage_fd = -1;
uint64_t memory_base = 0;
uint64_t memory_size = 0;
@@ -316,56 +286,45 @@ class HsaInterceptor {
char* uri_str = NULL;
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE,
&storage_type));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE,
&storage_type));
if (storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE) {
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE,
&storage_fd));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE,
&storage_fd));
if (storage_fd == -1) {
printf("CodeObjectCallback: fd == -1\n"); fflush(stdout);
abort();
printf("CodeObjectCallback: fd == -1\n");
fflush(stdout);
abort();
}
} else if (storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY) {
HSA_RT(LoaderApiTable.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,
&memory_base));
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE,
&memory_base));
HSA_RT(LoaderApiTable.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,
&memory_size));
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE,
&memory_size));
}
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
&load_base));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE, &load_base));
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
&load_size));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, &load_size));
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
&load_delta));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA, &load_delta));
// Getting URI
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
&uri_len));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH, &uri_len));
uri_str = (char*)calloc(uri_len + 1, sizeof(char));
if (!uri_str) EXC_ABORT(HSA_STATUS_ERROR, "URI allocation");
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI,
uri_str));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, uri_str));
if (storage_type != HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE) {
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_CODEOBJ) {
@@ -377,8 +336,8 @@ class HsaInterceptor {
data.codeobj.load_base = load_base;
data.codeobj.load_size = load_size;
data.codeobj.load_delta = load_delta;
data.codeobj.uri_length = uri_len;
data.codeobj.uri = uri_str;
data.codeobj.uri_length = uri_len;
data.codeobj.uri = uri_str;
data.codeobj.unload = free_flag;
DO_HSA_CALLBACK;
@@ -406,12 +365,8 @@ class HsaInterceptor {
uint32_t num_agents = 0;
hsa_agent_t* agents = NULL;
pointer_info.size = sizeof(hsa_amd_pointer_info_t);
HSA_RT(hsa_amd_pointer_info(
reinterpret_cast<void*>(load_base),
&pointer_info,
malloc,
&num_agents,
&agents));
HSA_RT(hsa_amd_pointer_info(reinterpret_cast<void*>(load_base), &pointer_info, malloc,
&num_agents, &agents));
DeviceCallback(num_agents, agents, reinterpret_cast<void*>(load_base));
}
@@ -420,11 +375,8 @@ class HsaInterceptor {
return HSA_STATUS_SUCCESS;
}
static hsa_status_t KernelSymbolCallback(
hsa_executable_t executable,
hsa_executable_symbol_t symbol,
void *arg)
{
static hsa_status_t KernelSymbolCallback(hsa_executable_t executable,
hsa_executable_symbol_t symbol, void* arg) {
const int free_flag = reinterpret_cast<long>(arg);
hsa_symbol_kind_t kind = (hsa_symbol_kind_t)0;
HSA_RT(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &kind));
@@ -433,9 +385,11 @@ class HsaInterceptor {
const char* name = NULL;
uint32_t len = 0;
uint64_t obj = 0;
HSA_RT(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &obj));
HSA_RT(
hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &obj));
if (free_flag == 0) {
HSA_RT(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len));
HSA_RT(
hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len));
char sym_name[len + 1];
HSA_RT(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, sym_name));
name = cpp_demangle(sym_name);
@@ -453,10 +407,7 @@ class HsaInterceptor {
return HSA_STATUS_SUCCESS;
}
static hsa_status_t ExecutableFreeze(
hsa_executable_t executable,
const char *options)
{
static hsa_status_t ExecutableFreeze(hsa_executable_t executable, const char* options) {
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_executable_freeze_fn(executable, options));
@@ -466,39 +417,29 @@ class HsaInterceptor {
{ IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) is_codeobj_cb |= 1; }
if (is_codeobj_cb) {
LoaderApiTable.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
executable,
CodeObjectCallback,
reinterpret_cast<void*>(0));
executable, CodeObjectCallback, reinterpret_cast<void*>(0));
}
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_KSYMBOL) {
HSA_RT(hsa_executable_iterate_symbols(
executable,
KernelSymbolCallback,
reinterpret_cast<void*>(0)));
HSA_RT(hsa_executable_iterate_symbols(executable, KernelSymbolCallback,
reinterpret_cast<void*>(0)));
}
return status;
}
static hsa_status_t ExecutableDestroy(
hsa_executable_t executable)
{
static hsa_status_t ExecutableDestroy(hsa_executable_t executable) {
hsa_status_t status = HSA_STATUS_SUCCESS;
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
LoaderApiTable.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
executable,
CodeObjectCallback,
reinterpret_cast<void*>(1));
executable, CodeObjectCallback, reinterpret_cast<void*>(1));
}
{
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_KSYMBOL) {
HSA_RT(hsa_executable_iterate_symbols(
executable,
KernelSymbolCallback,
reinterpret_cast<void*>(1)));
HSA_RT(hsa_executable_iterate_symbols(executable, KernelSymbolCallback,
reinterpret_cast<void*>(1)));
}
}
projects/rocprofiler/src/core/hsa_proxy_queue.h
+3 -3
Bestand weergeven
@@ -33,9 +33,9 @@ THE SOFTWARE.
#include "util/hsa_rsrc_factory.h"
namespace rocprofiler {
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
extern decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
extern decltype(::hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(::hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
extern decltype(::hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
class HsaProxyQueue : public ProxyQueue {
public:
projects/rocprofiler/src/core/intercept_queue.h
+72 -61
Bestand weergeven
@@ -40,16 +40,13 @@ THE SOFTWARE.
#include "util/hsa_rsrc_factory.h"
namespace rocprofiler {
enum {
K_CONC_OFF = 0,
K_CONC_PMC = 1,
K_CONC_TRACE = 2
};
enum { K_CONC_OFF = 0, K_CONC_PMC = 1, K_CONC_TRACE = 2 };
extern decltype(hsa_queue_create)* hsa_queue_create_fn;
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(::hsa_queue_create)* hsa_queue_create_fn;
extern decltype(::hsa_queue_destroy)* hsa_queue_destroy_fn;
static inline void print_packet(const void* in_p, const uint32_t& in_n, const uint32_t& w_n = UINT32_MAX) {
static inline void print_packet(const void* in_p, const uint32_t& in_n,
const uint32_t& w_n = UINT32_MAX) {
const uint32_t size32 = util::HsaRsrcFactory::CMD_SLOT_SIZE_B / 4;
const uint32_t* beg = (const uint32_t*)in_p;
const uint32_t* end = beg + (in_n * size32);
@@ -85,31 +82,33 @@ class InterceptQueue {
typedef std::recursive_mutex mutex_t;
typedef std::map<uint64_t, InterceptQueue*> obj_map_t;
typedef hsa_status_t (*queue_callback_t)(hsa_queue_t*, void* data);
typedef void (*queue_event_callback_t)(hsa_status_t status, hsa_queue_t *queue, void *arg);
typedef void (*queue_event_callback_t)(hsa_status_t status, hsa_queue_t* queue, void* arg);
typedef uint32_t queue_id_t;
static void HsaIntercept(HsaApiTable* table);
static hsa_status_t InterceptQueueCreate(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue,
const bool& tracker_on) {
static hsa_status_t InterceptQueueCreate(
hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue,
const bool& tracker_on) {
std::lock_guard<mutex_t> lck(mutex_);
hsa_status_t status = HSA_STATUS_ERROR;
if (in_create_call_) EXC_ABORT(status, "recursive InterceptQueueCreate()");
in_create_call_ = true;
ProxyQueue* proxy = ProxyQueue::Create(agent, size, type, queue_event_callback, data, private_segment_size,
group_segment_size, queue, &status);
ProxyQueue* proxy =
ProxyQueue::Create(agent, size, type, queue_event_callback, data, private_segment_size,
group_segment_size, queue, &status);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "ProxyQueue::Create()");
if (tracker_on || tracker_on_) {
if (tracker_ == NULL) tracker_ = &Tracker::Instance();
status = rocprofiler::util::HsaRsrcFactory::HsaApi()->hsa_amd_profiling_set_profiler_enabled(*queue, true);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_set_profiler_enabled()");
status = rocprofiler::util::HsaRsrcFactory::HsaApi()->hsa_amd_profiling_set_profiler_enabled(
*queue, true);
if (status != HSA_STATUS_SUCCESS)
EXC_ABORT(status, "hsa_amd_profiling_set_profiler_enabled()");
}
InterceptQueue* obj = new InterceptQueue(agent, *queue, proxy);
@@ -138,15 +137,17 @@ class InterceptQueue {
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue) {
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size, group_segment_size, queue, false);
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue, false);
}
static hsa_status_t QueueCreateTracked(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue) {
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size, group_segment_size, queue, true);
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue) {
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue, true);
}
static hsa_status_t QueueDestroy(hsa_queue_t* queue) {
@@ -170,8 +171,8 @@ class InterceptQueue {
return status;
}
static void OnSubmitCB_opt(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data,
hsa_amd_queue_intercept_packet_writer writer) {
static void OnSubmitCB_opt(const void* in_packets, uint64_t count, uint64_t user_que_idx,
void* data, hsa_amd_queue_intercept_packet_writer writer) {
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(in_packets);
InterceptQueue* obj = reinterpret_cast<InterceptQueue*>(data);
Queue* proxy = obj->proxy_;
@@ -195,10 +196,10 @@ class InterceptQueue {
obj->queue_id,
completion_signal,
dispatch_packet,
NULL, // kernel_name
0, // kernel_object
NULL, // kernel_code
0, // (uint32_t)syscall(__NR_gettid),
NULL, // kernel_name
0, // kernel_object
NULL, // kernel_code
0, // (uint32_t)syscall(__NR_gettid),
NULL}; // record
// Calling dispatch callback
@@ -210,7 +211,8 @@ class InterceptQueue {
if (group.feature_count != 0) {
if (tracker_ != NULL) {
Group* context_group = context->GetGroup(group.index);
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = context_group->GetDispatchSignal();
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal =
context_group->GetDispatchSignal();
Tracker::Enable_opt(context_group, completion_signal);
context_group->IncrRefsCount();
}
@@ -254,8 +256,9 @@ class InterceptQueue {
const uint32_t tid = syscall(__NR_gettid);
hsa_queue_t* qptr = obj->queue_;
const void* slot_ptr = util::HsaRsrcFactory::GetSlotPointer(qptr, user_que_idx);
printf("OnSubmitCB: %u:%u queue(%p:%lu) in(%p, %p, %lu) hdr(%u)\n",
pid, tid, qptr, user_que_idx, in_packets, slot_ptr, count, header_val); fflush(stdout);
printf("OnSubmitCB: %u:%u queue(%p:%lu) in(%p, %p, %lu) hdr(%u)\n", pid, tid, qptr,
user_que_idx, in_packets, slot_ptr, count, header_val);
fflush(stdout);
print_packet(in_packets, count);
abort();
#endif
@@ -277,8 +280,9 @@ class InterceptQueue {
if (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) {
uint64_t kernel_object = dispatch_packet->kernel_object;
const amd_kernel_code_t* kernel_code = GetKernelCode(kernel_object);
kernel_name = (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) ?
QueryKernelName(kernel_object, kernel_code) : NULL;
kernel_name = (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH)
? QueryKernelName(kernel_object, kernel_code)
: NULL;
}
// Prepareing submit callback data
@@ -311,8 +315,11 @@ class InterceptQueue {
const bool is_serial = (k_concurrent_ == K_CONC_OFF);
if (tracker_ != NULL) {
tracker_entry = tracker_->Alloc(obj->agent_info_->dev_id, dispatch_packet->completion_signal, is_serial);
if (is_serial) const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = tracker_entry->signal;
tracker_entry = tracker_->Alloc(obj->agent_info_->dev_id,
dispatch_packet->completion_signal, is_serial);
if (is_serial)
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal =
tracker_entry->signal;
}
// Prepareing dispatch callback data
@@ -339,7 +346,9 @@ class InterceptQueue {
// Injecting profiling start/stop/read packets
if ((status != HSA_STATUS_SUCCESS) || (group.context == NULL)) {
if (tracker_entry != NULL) {
if (is_serial) const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = tracker_entry->orig;
if (is_serial)
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal =
tracker_entry->orig;
tracker_->Delete(tracker_entry);
}
} else {
@@ -351,11 +360,11 @@ class InterceptQueue {
const pkt_vector_t& read_vector = context->ReadPackets(group.index);
pkt_vector_t packets;
if (is_serial) { // serial
if (is_serial) { // serial
packets = start_vector;
packets.insert(packets.end(), *packet);
packets.insert(packets.end(), stop_vector.begin(), stop_vector.end());
} else { // concurrent
} else { // concurrent
// Insert start packets once
auto inject_start = [&packets](const pkt_vector_t& starts) mutable {
packets = starts;
@@ -363,14 +372,15 @@ class InterceptQueue {
std::call_once(once_flag_, inject_start, start_vector);
// Reads at both kernel start and end (also with barriers)
assert(read_vector.size() >= 2 * start_vector.size());
auto mid = read_vector.begin() + read_vector.size()/2;
auto mid = read_vector.begin() + read_vector.size() / 2;
// Read at kernel start
packets.insert(packets.end(), read_vector.begin(), mid);
// Kernel dispatch packet
assert(tracker_entry != NULL);
// Bind dispatch and barrier signals with tracker entry
tracker_->SetHandler(tracker_entry, context->GetGroup(group.index));
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = context->GetGroup(group.index)->GetDispatchSignal();
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal =
context->GetGroup(group.index)->GetDispatchSignal();
packets.insert(packets.end(), *packet);
// Read at kernel end
packets.insert(packets.end(), mid, read_vector.end());
@@ -379,7 +389,8 @@ class InterceptQueue {
if (tracker_entry != NULL) {
Group* context_group = context->GetGroup(group.index);
context_group->IncrRefsCount();
tracker_->EnableContext(tracker_entry, Context::Handler, reinterpret_cast<void*>(context_group));
tracker_->EnableContext(tracker_entry, Context::Handler,
reinterpret_cast<void*>(context_group));
}
if (writer != NULL) {
@@ -409,8 +420,8 @@ class InterceptQueue {
}
}
static void OnSubmitCB_ctrace(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data,
hsa_amd_queue_intercept_packet_writer writer) {
static void OnSubmitCB_ctrace(const void* in_packets, uint64_t count, uint64_t user_que_idx,
void* data, hsa_amd_queue_intercept_packet_writer writer) {
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(in_packets);
InterceptQueue* obj = reinterpret_cast<InterceptQueue*>(data);
Queue* proxy = obj->proxy_;
@@ -431,8 +442,9 @@ class InterceptQueue {
if (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) {
uint64_t kernel_object = dispatch_packet->kernel_object;
const amd_kernel_code_t* kernel_code = GetKernelCode(kernel_object);
kernel_name = (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) ?
QueryKernelName(kernel_object, kernel_code) : NULL;
kernel_name = (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH)
? QueryKernelName(kernel_object, kernel_code)
: NULL;
}
// Prepareing submit callback data
@@ -529,7 +541,9 @@ class InterceptQueue {
Stop();
}
static inline void Start() { dispatch_callback_.store(callbacks_.dispatch, std::memory_order_release); }
static inline void Start() {
dispatch_callback_.store(callbacks_.dispatch, std::memory_order_release);
}
static inline void Stop() { dispatch_callback_.store(NULL, std::memory_order_relaxed); }
static void SetSubmitCallback(rocprofiler_hsa_callback_fun_t fun, void* arg) {
@@ -545,7 +559,7 @@ class InterceptQueue {
static uint32_t k_concurrent_;
private:
static void queue_event_callback(hsa_status_t status, hsa_queue_t *queue, void *arg) {
static void queue_event_callback(hsa_status_t status, hsa_queue_t* queue, void* arg) {
if (status != HSA_STATUS_SUCCESS) {
uint32_t* read_ptr32 = (uint32_t*)util::HsaRsrcFactory::GetReadPointer(queue);
print_packet(read_ptr32, 1);
@@ -582,12 +596,13 @@ class InterceptQueue {
const uint16_t kernel_object_flag = *((uint64_t*)kernel_code + 1);
if (kernel_object_flag == 0) {
if (!util::HsaRsrcFactory::IsExecutableTracking()) {
EXC_ABORT(HSA_STATUS_ERROR, "Error: V3 code object detected - code objects tracking should be enabled\n");
EXC_ABORT(HSA_STATUS_ERROR,
"Error: V3 code object detected - code objects tracking should be enabled\n");
}
}
const char* kernel_symname = (util::HsaRsrcFactory::IsExecutableTracking()) ?
util::HsaRsrcFactory::GetKernelNameRef(kernel_object) :
GetKernelName(kernel_code->runtime_loader_kernel_symbol);
const char* kernel_symname = (util::HsaRsrcFactory::IsExecutableTracking())
? util::HsaRsrcFactory::GetKernelNameRef(kernel_object)
: GetKernelName(kernel_code->runtime_loader_kernel_symbol);
return kernel_symname;
}
@@ -618,17 +633,13 @@ class InterceptQueue {
return status;
}
InterceptQueue(const hsa_agent_t& agent, hsa_queue_t* const queue, ProxyQueue* proxy) :
queue_(queue),
proxy_(proxy)
{
InterceptQueue(const hsa_agent_t& agent, hsa_queue_t* const queue, ProxyQueue* proxy)
: queue_(queue), proxy_(proxy) {
agent_info_ = util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
queue_event_callback_ = NULL;
}
~InterceptQueue() {
ProxyQueue::Destroy(proxy_);
}
~InterceptQueue() { ProxyQueue::Destroy(proxy_); }
static const packet_word_t header_type_mask = (1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1;
projects/rocprofiler/src/core/metrics.cpp
+1 -1
Bestand weergeven
@@ -25,4 +25,4 @@ THE SOFTWARE.
namespace rocprofiler {
MetricsDict::map_t* MetricsDict::map_ = NULL;
MetricsDict::mutex_t MetricsDict::mutex_;
}
} // namespace rocprofiler
projects/rocprofiler/src/core/metrics.h
Executable → Regular
+5 -5
Bestand weergeven
@@ -202,15 +202,15 @@ class MetricsDict {
xml_->AddConst("top.const.metric", "SE_NUM", agent_info->se_num);
ImportMetrics(agent_info, "const");
agent_name_ = agent_info->name;
if (agent_name_.find(':') != std::string::npos) // Remove compiler flags from the agent_name
if (agent_name_.find(':') != std::string::npos) // Remove compiler flags from the agent_name
agent_name_ = agent_name_.substr(0, agent_name_.find(':'));
std::unordered_set<std::string> supported_agent_names = {
"gfx906", "gfx908", "gfx90a", // Vega
"gfx940", "gfx941", "gfx942", // Mi300
"gfx906", "gfx908", "gfx90a", // Vega
"gfx940", "gfx941", "gfx942", // Mi300
"gfx1030", "gfx1031", "gfx1032", // Navi2x
"gfx1100", "gfx1101" // Navi3x
"gfx1100", "gfx1101" // Navi3x
};
if (supported_agent_names.find(agent_name_) != supported_agent_names.end()) {
ImportMetrics(agent_info, agent_name_);
projects/rocprofiler/src/core/profile.h
+25 -21
Bestand weergeven
@@ -140,7 +140,7 @@ class Profile {
static void SetConcurrent(profile_t* profile) {
// Check whether conconcurrent has been set
for (const parameter_t* p = profile->parameters;
p < (profile->parameters + profile->parameter_count); ++p) {
p < (profile->parameters + profile->parameter_count); ++p) {
// If yes, stop here
if (p->parameter_name == HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_K_CONCURRENT) {
return;
@@ -148,7 +148,7 @@ class Profile {
}
// Otherwise, try to set
parameter_t* parameters = new parameter_t[profile->parameter_count+1];
parameter_t* parameters = new parameter_t[profile->parameter_count + 1];
for (unsigned i = 0; i < profile->parameter_count; ++i) {
parameters[i].parameter_name = profile->parameters[i].parameter_name;
parameters[i].value = profile->parameters[i].value;
@@ -162,15 +162,16 @@ class Profile {
}
void BarrierPacket(packet_t* packet, const hsa_signal_t& prior_signal) {
hsa_barrier_and_packet_t* barrier =
reinterpret_cast<hsa_barrier_and_packet_t*>(packet);
hsa_barrier_and_packet_t* barrier = reinterpret_cast<hsa_barrier_and_packet_t*>(packet);
barrier->header = HSA_PACKET_TYPE_BARRIER_AND;
if (prior_signal.handle) barrier->dep_signal[0] = prior_signal; // set packet dependency
else barrier->header |= 1 << HSA_PACKET_HEADER_BARRIER; // set barrier bit
if (prior_signal.handle)
barrier->dep_signal[0] = prior_signal; // set packet dependency
else
barrier->header |= 1 << HSA_PACKET_HEADER_BARRIER; // set barrier bit
}
hsa_status_t Finalize(pkt_vector_t& start_vector, pkt_vector_t& stop_vector,
pkt_vector_t& read_vector, bool is_concurrent = false) {
pkt_vector_t& read_vector, bool is_concurrent = false) {
if (is_concurrent) SetConcurrent(&profile_);
hsa_status_t status = HSA_STATUS_SUCCESS;
@@ -180,8 +181,8 @@ class Profile {
const pfn_t* api = rsrc->AqlProfileApi();
packet_t start{};
packet_t stop{};
packet_t read{}; // read at kernel start
packet_t read2{}; // read at kernel end
packet_t read{}; // read at kernel start
packet_t read2{}; // read at kernel end
// Check the profile buffer sizes
status = api->hsa_ven_amd_aqlprofile_start(&profile_, NULL);
@@ -200,12 +201,12 @@ class Profile {
#ifdef AQLPROF_NEW_API
if (profile_.type == HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC) {
rd_status = api->hsa_ven_amd_aqlprofile_read(&profile_, &read);
if (is_concurrent){ // concurrent: one more read
if (is_concurrent) { // concurrent: one more read
if (rd_status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "aqlprofile_read");
rd_status = api->hsa_ven_amd_aqlprofile_read(&profile_, &read2);
}
}
#if 0 // Read API returns error if disabled
#if 0 // Read API returns error if disabled
if (rd_status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "aqlprofile_read");
#endif
#endif
@@ -220,7 +221,8 @@ class Profile {
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
if (is_concurrent) {
status = hsa_signal_create(1, 0, NULL, &read_signal_);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
if (status != HSA_STATUS_SUCCESS)
EXC_RAISING(status, "signal_create " << std::hex << status);
read.completion_signal = read_signal_;
read2.completion_signal = completion_signal_;
} else {
@@ -239,7 +241,8 @@ class Profile {
BarrierPacket(&barrier_rd, read.completion_signal);
BarrierPacket(&barrier_rd2, dispatch_signal_);
status = hsa_signal_create(1, 0, NULL, &(barrier_signal_));
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
if (status != HSA_STATUS_SUCCESS)
EXC_RAISING(status, "signal_create " << std::hex << status);
barrier_rd2.completion_signal = barrier_signal_;
}
@@ -297,8 +300,8 @@ class Profile {
void GetProfiles(profile_vector_t& vec) {
if (!info_vector_.empty()) {
vec.push_back(profile_tuple_t{&profile_, &info_vector_, completion_signal_,
dispatch_signal_, barrier_signal_, read_signal_});
vec.push_back(profile_tuple_t{&profile_, &info_vector_, completion_signal_, dispatch_signal_,
barrier_signal_, read_signal_});
}
}
@@ -330,11 +333,12 @@ class PmcProfile : public Profile {
hsa_status_t Allocate(util::HsaRsrcFactory* rsrc) {
profile_.command_buffer.ptr =
rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
// Allocate profile output buffer from kernarg memory pool since kernarg
// memory buffer is uncached. So when GPU copies performance counter values
// to this buffer they are guaranteed to be visible to CPU.
profile_.output_buffer.ptr = rsrc->AllocateKernArgMemory(agent_info_, profile_.output_buffer.size);
profile_.output_buffer.ptr =
rsrc->AllocateKernArgMemory(agent_info_, profile_.output_buffer.size);
return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS
: HSA_STATUS_ERROR;
}
@@ -366,11 +370,11 @@ class TraceProfile : public Profile {
hsa_status_t Allocate(util::HsaRsrcFactory* rsrc) {
profile_.command_buffer.ptr =
rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
rsrc->AllocateSysMemory(agent_info_, profile_.command_buffer.size);
profile_.output_buffer.size = output_buffer_size_;
profile_.output_buffer.ptr = (output_buffer_local_) ?
rsrc->AllocateLocalMemory(agent_info_, profile_.output_buffer.size) :
rsrc->AllocateSysMemory(agent_info_, profile_.output_buffer.size);
profile_.output_buffer.ptr = (output_buffer_local_)
? rsrc->AllocateLocalMemory(agent_info_, profile_.output_buffer.size)
: rsrc->AllocateSysMemory(agent_info_, profile_.output_buffer.size);
return (profile_.command_buffer.ptr && profile_.output_buffer.ptr) ? HSA_STATUS_SUCCESS
: HSA_STATUS_ERROR;
}
projects/rocprofiler/src/core/proxy_queue.cpp
+2 -2
Bestand weergeven
@@ -38,10 +38,10 @@ ProxyQueue* ProxyQueue::Create(hsa_agent_t agent, uint32_t size, hsa_queue_type3
hsa_status_t* status) {
hsa_status_t suc = HSA_STATUS_ERROR;
ProxyQueue* instance =
(rocp_type_) ? (ProxyQueue*) new SimpleProxyQueue() : (ProxyQueue*) new HsaProxyQueue();
(rocp_type_) ? (ProxyQueue*)new SimpleProxyQueue() : (ProxyQueue*)new HsaProxyQueue();
if (instance != NULL) {
suc = instance->Init(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue);
group_segment_size, queue);
if (suc != HSA_STATUS_SUCCESS) {
delete instance;
instance = NULL;
projects/rocprofiler/src/core/rocprofiler.cpp
+80 -80
Bestand weergeven
@@ -75,34 +75,34 @@ hsa_status_t CreateQueuePro(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t
void* data, uint32_t private_segment_size, uint32_t group_segment_size,
hsa_queue_t** queue);
decltype(hsa_queue_create)* hsa_queue_create_fn;
decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
decltype(::hsa_queue_create)* hsa_queue_create_fn;
decltype(::hsa_queue_destroy)* hsa_queue_destroy_fn;
decltype(hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
decltype(hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
decltype(::hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
decltype(::hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
decltype(hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
decltype(hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
decltype(hsa_queue_add_write_index_scacq_screl)* hsa_queue_add_write_index_scacq_screl_fn;
decltype(::hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
decltype(::hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
decltype(::hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
decltype(::hsa_queue_add_write_index_scacq_screl)* hsa_queue_add_write_index_scacq_screl_fn;
decltype(hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
decltype(hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
decltype(::hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
decltype(::hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
decltype(::hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
decltype(::hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
decltype(::hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
decltype(hsa_memory_allocate)* hsa_memory_allocate_fn;
decltype(hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
decltype(hsa_memory_copy)* hsa_memory_copy_fn;
decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn;
decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
decltype(hsa_executable_destroy)* hsa_executable_destroy_fn;
decltype(::hsa_memory_allocate)* hsa_memory_allocate_fn;
decltype(::hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
decltype(::hsa_memory_copy)* hsa_memory_copy_fn;
decltype(::hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
decltype(::hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
decltype(::hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
decltype(::hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
decltype(::hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn;
decltype(::hsa_executable_freeze)* hsa_executable_freeze_fn;
decltype(::hsa_executable_destroy)* hsa_executable_destroy_fn;
::HsaApiTable* kHsaApiTable;
@@ -393,80 +393,80 @@ ROCPROFILER_EXPORT extern const uint32_t HSA_AMD_TOOL_PRIORITY = 25;
PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names) {
ONLOAD_TRACE_BEG();
rocprofiler::SaveHsaApi(table);
rocprofiler::ProxyQueue::InitFactory();
rocprofiler::SaveHsaApi(table);
rocprofiler::ProxyQueue::InitFactory();
// Checking environment to enable intercept mode
const char* intercept_env = getenv("ROCP_HSA_INTERCEPT");
// Checking environment to enable intercept mode
const char* intercept_env = getenv("ROCP_HSA_INTERCEPT");
int intercept_env_value = 0;
if (intercept_env != NULL) {
intercept_env_value = atoi(intercept_env);
int intercept_env_value = 0;
if (intercept_env != NULL) {
intercept_env_value = atoi(intercept_env);
switch (intercept_env_value) {
case 0:
case 1:
// 0: Intercepting disabled
// 1: Intercepting enabled without timestamping
rocprofiler::InterceptQueue::TrackerOn(false);
break;
case 2:
// Intercepting enabled with timestamping
rocprofiler::InterceptQueue::TrackerOn(true);
break;
default:
ERR_LOGGING("Bad ROCP_HSA_INTERCEPT env var value ("
<< intercept_env << "): "
<< "valid values are 0 (standalone), 1 (intercepting without timestamp), 2 "
"(intercepting with timestamp)");
return false;
}
switch (intercept_env_value) {
case 0:
case 1:
// 0: Intercepting disabled
// 1: Intercepting enabled without timestamping
rocprofiler::InterceptQueue::TrackerOn(false);
break;
case 2:
// Intercepting enabled with timestamping
rocprofiler::InterceptQueue::TrackerOn(true);
break;
default:
ERR_LOGGING("Bad ROCP_HSA_INTERCEPT env var value ("
<< intercept_env << "): "
<< "valid values are 0 (standalone), 1 (intercepting without timestamp), 2 "
"(intercepting with timestamp)");
return false;
}
}
// always enable excutable tracking
rocprofiler::util::HsaRsrcFactory::EnableExecutableTracking(table);
// always enable excutable tracking
rocprofiler::util::HsaRsrcFactory::EnableExecutableTracking(table);
// Loading a tool lib and setting of intercept mode
const uint32_t intercept_mode_mask = rocprofiler::LoadTool();
// Loading a tool lib and setting of intercept mode
const uint32_t intercept_mode_mask = rocprofiler::LoadTool();
if (intercept_mode_mask & rocprofiler::MEMCOPY_INTERCEPT_MODE) {
hsa_status_t status = hsa_amd_profiling_async_copy_enable(true);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_async_copy_enable");
rocprofiler::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
rocprofiler::hsa_amd_memory_async_copy_rect_fn =
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn;
table->amd_ext_->hsa_amd_memory_async_copy_fn =
rocprofiler::hsa_amd_memory_async_copy_interceptor;
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn =
rocprofiler::hsa_amd_memory_async_copy_rect_interceptor;
}
if (intercept_mode_mask & rocprofiler::HSA_INTERCEPT_MODE) {
if (intercept_mode_mask & rocprofiler::MEMCOPY_INTERCEPT_MODE) {
hsa_status_t status = hsa_amd_profiling_async_copy_enable(true);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_async_copy_enable");
rocprofiler::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
rocprofiler::hsa_amd_memory_async_copy_rect_fn =
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn;
table->amd_ext_->hsa_amd_memory_async_copy_fn =
rocprofiler::hsa_amd_memory_async_copy_interceptor;
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn =
rocprofiler::hsa_amd_memory_async_copy_rect_interceptor;
}
if (intercept_mode_mask & rocprofiler::HSA_INTERCEPT_MODE) {
if (intercept_mode_mask & rocprofiler::MEMCOPY_INTERCEPT_MODE) {
EXC_ABORT(HSA_STATUS_ERROR, "HSA_INTERCEPT and MEMCOPY_INTERCEPT conflict");
}
rocprofiler::HsaInterceptor::Enable(true);
rocprofiler::HsaInterceptor::HsaIntercept(table);
EXC_ABORT(HSA_STATUS_ERROR, "HSA_INTERCEPT and MEMCOPY_INTERCEPT conflict");
}
rocprofiler::HsaInterceptor::Enable(true);
rocprofiler::HsaInterceptor::HsaIntercept(table);
}
// HSA intercepting
if (intercept_env_value != 0) {
rocprofiler::ProxyQueue::HsaIntercept(table);
rocprofiler::InterceptQueue::HsaIntercept(table);
} else {
rocprofiler::StandaloneIntercept();
}
// HSA intercepting
if (intercept_env_value != 0) {
rocprofiler::ProxyQueue::HsaIntercept(table);
rocprofiler::InterceptQueue::HsaIntercept(table);
} else {
rocprofiler::StandaloneIntercept();
}
ONLOAD_TRACE("end intercept_mode(" << std::hex << intercept_env_value << ")"
<< " intercept_mode_mask(" << std::hex << intercept_mode_mask
<< ")" << std::dec);
ONLOAD_TRACE("end intercept_mode(" << std::hex << intercept_env_value << ")"
<< " intercept_mode_mask(" << std::hex << intercept_mode_mask
<< ")" << std::dec);
return true;
}
// HSA-runtime tool on-unload method
PUBLIC_API void OnUnload() {
ONLOAD_TRACE_BEG();
rocprofiler::UnloadTool();
rocprofiler::RestoreHsaApi();
rocprofiler::UnloadTool();
rocprofiler::RestoreHsaApi();
ONLOAD_TRACE_END();
}
projects/rocprofiler/src/core/session/att/att.cpp
+9 -11
Bestand weergeven
@@ -27,22 +27,20 @@ namespace rocprofiler {
namespace att {
AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
rocprofiler_session_id_t session_id)
rocprofiler_session_id_t session_id)
: buffer_id_(buffer_id), filter_id_(filter_id), session_id_(session_id) {}
void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal,
rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id,
hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties,
uint32_t thread_id, uint64_t queue_index) {
void AttTracer::AddPendingSignals(
uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal,
const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index) {
std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end())
sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>());
sessions_pending_signals_.at(writer_id).emplace_back(
att_pending_signal_t{kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, profile,
kernel_properties, thread_id, queue_index});
sessions_pending_signals_.at(writer_id).emplace_back(att_pending_signal_t{
kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id,
profile, kernel_properties, thread_id, queue_index});
std::atomic_thread_fence(std::memory_order_release);
}
projects/rocprofiler/src/core/session/filter.cpp
+5 -7
Bestand weergeven
@@ -40,7 +40,7 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind
}
break;
}
case ROCPROFILER_PC_SAMPLING_COLLECTION:{
case ROCPROFILER_PC_SAMPLING_COLLECTION: {
break;
}
case ROCPROFILER_ATT_TRACE_COLLECTION: {
@@ -62,8 +62,8 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind
}
case ROCPROFILER_API_TRACE: {
tracer_apis_.clear();
for (uint32_t j = 0; j < data_count; j++){
tracer_apis_.emplace_back(filter_data.trace_apis[j]);
for (uint32_t j = 0; j < data_count; j++) {
tracer_apis_.emplace_back(filter_data.trace_apis[j]);
}
break;
}
@@ -195,7 +195,7 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) {
case ROCPROFILER_FILTER_DISPATCH_IDS:
dispatch_id_filter_.clear();
for (uint32_t j = 0; j < property.data_count; j++)
dispatch_id_filter_.emplace_back(property.dispatch_ids[j]);
dispatch_id_filter_.emplace_back(property.dispatch_ids[j]);
break;
default:
break;
@@ -249,9 +249,7 @@ void Filter::SetCallback(rocprofiler_sync_callback_t& callback) {
bool Filter::HasCallback() { return has_sync_callback_; }
rocprofiler_sync_callback_t& Filter::GetCallback() {
return callback_;
}
rocprofiler_sync_callback_t& Filter::GetCallback() { return callback_; }
size_t Filter::GetPropertiesCount(rocprofiler_filter_property_kind_t kind) {
switch (kind) {
projects/rocprofiler/src/core/session/filter.h
+6 -8
Bestand weergeven
@@ -53,11 +53,8 @@ class Filter {
bool HasCallback();
void SetProperty(rocprofiler_filter_property_t property);
std::variant<
std::vector<std::string>,
uint32_t*,
std::vector<uint64_t>
> GetProperty(rocprofiler_filter_property_kind_t kind);
std::variant<std::vector<std::string>, uint32_t*, std::vector<uint64_t> > GetProperty(
rocprofiler_filter_property_kind_t kind);
size_t GetPropertiesCount(rocprofiler_filter_property_kind_t kind);
rocprofiler_spm_parameter_t* GetSpmParameterData();
@@ -74,11 +71,12 @@ class Filter {
std::vector<std::string> kernel_names_; // HIP/HSA API Functions
uint32_t dispatch_range_[2]; // Kernel Dispatches OR API Range
std::vector<std::string> profiler_counter_names_; // Counter Names to collect
std::vector<std::string> profiler_counter_names_; // Counter Names to collect
std::vector<rocprofiler_tracer_activity_domain_t> tracer_apis_; // ROCTX/HIP/HSA API
rocprofiler_spm_parameter_t* spm_parameter_; // spm parameter
std::vector<rocprofiler_att_parameter_t> att_parameters_; // ATT Parameters
rocprofiler_counters_sampler_parameters_t counters_sampler_parameters_; // sampled counters parameters
std::vector<rocprofiler_att_parameter_t> att_parameters_; // ATT Parameters
rocprofiler_counters_sampler_parameters_t
counters_sampler_parameters_; // sampled counters parameters
std::vector<uint64_t> dispatch_id_filter_;
bool has_sync_callback_{false};
projects/rocprofiler/src/core/session/profiler/profiler.cpp
+10 -8
Bestand weergeven
@@ -125,17 +125,19 @@ bool Profiler::HasActivePass() {
}
void Profiler::AddPendingSignals(
uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal,
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile, rocprofiler_kernel_properties_t kernel_properties,
uint32_t thread_id, uint64_t queue_index, uint64_t correlation_id) {
uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal,
const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id, rocprofiler::profiling_context_t* context,
uint64_t session_data_count, hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_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_);
if (sessions_pending_signals_->find(writer_id) == sessions_pending_signals_->end())
sessions_pending_signals_->emplace(writer_id, std::vector<pending_signal_t*>());
sessions_pending_signals_->at(writer_id).emplace_back(new pending_signal_t{
kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, context, session_data_count,
profile, kernel_properties, thread_id, queue_index, correlation_id});
sessions_pending_signals_->at(writer_id).emplace_back(
new pending_signal_t{kernel_object, original_completion_signal, new_completion_signal,
session_id_, buffer_id, context, session_data_count, profile,
kernel_properties, thread_id, queue_index, correlation_id});
}
const std::vector<pending_signal_t*>& Profiler::GetPendingSignals(uint32_t writer_id) {
projects/rocprofiler/src/core/session/profiler/profiler.h
+12 -9
Bestand weergeven
@@ -36,7 +36,7 @@
#include "src/core/counters/metrics/eval_metrics.h"
typedef void (*rocprofiler_add_profiler_record_t)(rocprofiler_record_profiler_t&& record,
rocprofiler_session_id_t session_id);
rocprofiler_session_id_t session_id);
typedef rocprofiler_timestamp_t (*rocprofiler_get_timestamp_t)();
@@ -68,12 +68,13 @@ class Profiler {
~Profiler();
void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
rocprofiler_buffer_id_t buffer_id,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
uint64_t queue_index, uint64_t correlation_id);
const hsa_signal_t& original_completion_signal,
const hsa_signal_t& new_completion_signal,
rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
rocprofiler::profiling_context_t* context, uint64_t session_data_count,
hsa_ven_amd_aqlprofile_profile_t* profile,
rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
uint64_t queue_index, uint64_t correlation_id);
const std::vector<pending_signal_t*>& GetPendingSignals(uint32_t writer_id);
bool CheckPendingSignalsIsEmpty();
@@ -83,8 +84,10 @@ class Profiler {
std::string& GetCounterName(rocprofiler_counter_id_t handler);
bool FindCounter(rocprofiler_counter_id_t counter_id);
size_t GetCounterInfoSize(rocprofiler_counter_info_kind_t kind, rocprofiler_counter_id_t counter_id);
const char* GetCounterInfo(rocprofiler_counter_info_kind_t kind, rocprofiler_counter_id_t counter_id);
size_t GetCounterInfoSize(rocprofiler_counter_info_kind_t kind,
rocprofiler_counter_id_t counter_id);
const char* GetCounterInfo(rocprofiler_counter_info_kind_t kind,
rocprofiler_counter_id_t counter_id);
void StartReplayPass(rocprofiler_session_id_t session_id);
void EndReplayPass();
projects/rocprofiler/src/core/session/session.h
+3 -3
Bestand weergeven
@@ -67,8 +67,8 @@ class Session {
// Filter
rocprofiler_filter_id_t CreateFilter(rocprofiler_filter_kind_t filter_kind,
rocprofiler_filter_data_t filter_data, uint64_t data_count,
rocprofiler_filter_property_t property);
rocprofiler_filter_data_t filter_data, uint64_t data_count,
rocprofiler_filter_property_t property);
bool FindFilter(rocprofiler_filter_id_t filter_id);
void DestroyFilter(rocprofiler_filter_id_t filter_id);
Filter* GetFilter(rocprofiler_filter_id_t filter_id);
@@ -83,7 +83,7 @@ class Session {
// Buffer
rocprofiler_buffer_id_t CreateBuffer(rocprofiler_buffer_callback_t buffer_callback,
size_t buffer_size);
size_t buffer_size);
bool FindBuffer(rocprofiler_buffer_id_t buffer_id);
void DestroyBuffer(rocprofiler_buffer_id_t buffer_id);
Memory::GenericBuffer* GetBuffer(rocprofiler_buffer_id_t buffer_id);
projects/rocprofiler/src/core/session/tracer/src/roctracer.cpp
+13 -24
Bestand weergeven
@@ -112,8 +112,7 @@ const char* roctracer_op_string(uint32_t domain, uint32_t op) {
case ACTIVITY_DOMAIN_EXT_API:
return "EXT_API";
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -178,8 +177,7 @@ constexpr uint32_t get_op_begin(activity_domain_t domain) {
case ACTIVITY_DOMAIN_EXT_API:
return 0;
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -200,8 +198,7 @@ constexpr uint32_t get_op_end(activity_domain_t domain) {
case ACTIVITY_DOMAIN_EXT_API:
return get_op_begin(ACTIVITY_DOMAIN_EXT_API);
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -476,11 +473,10 @@ int TracerCallback(activity_domain_t domain, uint32_t operation_id, void* data)
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
->GetBuffer((*pool)->buffer_id)
->AddRecord(
rocprofiler_record, record->kernel_name, kernel_name_size,
[](auto& rocprofiler_record, const void* data) {
rocprofiler_record.name = static_cast<const char*>(data);
});
->AddRecord(rocprofiler_record, record->kernel_name, kernel_name_size,
[](auto& rocprofiler_record, const void* data) {
rocprofiler_record.name = static_cast<const char*>(data);
});
} else {
rocprofiler::GetROCProfilerSingleton()
->GetSession((*pool)->session_id)
@@ -584,8 +580,7 @@ static void roctracer_enable_op_callback(activity_domain_t domain, uint32_t oper
user_data);
break;
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -623,8 +618,7 @@ void roctracer_disable_op_callback(activity_domain_t domain, uint32_t operation_
ROCTX_registration_group.Unregister(roctx_api_callback_table, operation_id);
break;
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -667,8 +661,7 @@ void roctracer_enable_op_activity(activity_domain_t domain, uint32_t op,
case ACTIVITY_DOMAIN_ROCTX:
break;
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -710,8 +703,7 @@ void roctracer_disable_activity(activity_domain_t domain, uint32_t op) {
case ACTIVITY_DOMAIN_ROCTX:
break;
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -774,8 +766,7 @@ void roctracer_set_properties(activity_domain_t domain, void* properties) {
break;
}
default:
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID,
"Invalid domain ID");
throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_INVALID_DOMAIN_ID, "Invalid domain ID");
}
}
@@ -791,9 +782,7 @@ static std::string getKernelNameMultiKernelMultiDevice(hipLaunchParams* launchPa
return name_str.str();
}
template <typename... Ts> struct Overloaded : Ts... {
using Ts::operator()...;
};
template <typename... Ts> struct Overloaded : Ts... { using Ts::operator()...; };
template <class... Ts> Overloaded(Ts...) -> Overloaded<Ts...>;
std::optional<std::string> GetHipKernelName(uint32_t cid, hip_api_data_t* data) {
projects/rocprofiler/src/core/simple_proxy_queue.cpp
+12 -6
Bestand weergeven
@@ -27,13 +27,19 @@ void SimpleProxyQueue::HsaIntercept(HsaApiTable* table) {
table->core_->hsa_signal_store_relaxed_fn = rocprofiler::SimpleProxyQueue::SignalStore;
table->core_->hsa_signal_store_screlease_fn = rocprofiler::SimpleProxyQueue::SignalStore;
table->core_->hsa_queue_load_write_index_relaxed_fn = rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_relaxed_fn = rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_relaxed_fn = rocprofiler::SimpleProxyQueue::GetSubmitIndex;
table->core_->hsa_queue_load_write_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::GetSubmitIndex;
table->core_->hsa_queue_load_write_index_scacquire_fn = rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_screlease_fn = rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_scacquire_fn = rocprofiler::SimpleProxyQueue::GetSubmitIndex;
table->core_->hsa_queue_load_write_index_scacquire_fn =
rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_screlease_fn =
rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_scacquire_fn =
rocprofiler::SimpleProxyQueue::GetSubmitIndex;
}
SimpleProxyQueue::queue_map_t* SimpleProxyQueue::queue_map_ = NULL;
projects/rocprofiler/src/core/simple_proxy_queue.h
+16 -16
Bestand weergeven
@@ -33,23 +33,23 @@ THE SOFTWARE.
#include "util/hsa_rsrc_factory.h"
#ifndef ROCP_PROXY_LOCK
# define ROCP_PROXY_LOCK 1
#define ROCP_PROXY_LOCK 1
#endif
namespace rocprofiler {
extern decltype(hsa_queue_create)* hsa_queue_create_fn;
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(::hsa_queue_create)* hsa_queue_create_fn;
extern decltype(::hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
extern decltype(hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
extern decltype(::hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
extern decltype(::hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
extern decltype(hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
extern decltype(hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
extern decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
extern decltype(::hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
extern decltype(::hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
extern decltype(::hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
extern decltype(hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
extern decltype(hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
extern decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
extern decltype(::hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
extern decltype(::hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
extern decltype(::hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
typedef decltype(hsa_signal_t::handle) signal_handle_t;
@@ -128,7 +128,8 @@ class SimpleProxyQueue : public ProxyQueue {
const uint64_t que_idx = hsa_queue_load_write_index_relaxed_fn(queue_);
// Waiting untill there is a free space in the queue
while (que_idx >= (hsa_queue_load_read_index_relaxed_fn(queue_) + size_));
while (que_idx >= (hsa_queue_load_read_index_relaxed_fn(queue_) + size_))
;
// Increment the write index
hsa_queue_store_write_index_relaxed_fn(queue_, que_idx + 1);
@@ -163,8 +164,7 @@ class SimpleProxyQueue : public ProxyQueue {
queue_mask_(0),
submit_index_(0),
on_submit_cb_(NULL),
on_submit_cb_data_(NULL)
{
on_submit_cb_data_(NULL) {
printf("ROCProfiler: SimpleProxyQueue is enabled\n");
fflush(stdout);
}
@@ -203,8 +203,8 @@ class SimpleProxyQueue : public ProxyQueue {
if (queue_map_ == NULL) queue_map_ = new queue_map_t;
(*queue_map_)[queue_->doorbell_signal.handle] = this;
}
else abort();
} else
abort();
}
}
if (status != HSA_STATUS_SUCCESS) abort();
projects/rocprofiler/src/core/tracker.h
+40 -28
Bestand weergeven
@@ -40,7 +40,7 @@ THE SOFTWARE.
namespace rocprofiler {
class Tracker {
public:
public:
typedef std::mutex mutex_t;
typedef util::HsaRsrcFactory::timestamp_t timestamp_t;
typedef rocprofiler_dispatch_record_t record_t;
@@ -89,7 +89,7 @@ class Tracker {
}
// Add tracker entry
entry_t* Alloc(const hsa_agent_t& agent, const hsa_signal_t& orig, bool proxy=true) {
entry_t* Alloc(const hsa_agent_t& agent, const hsa_signal_t& orig, bool proxy = true) {
hsa_status_t status = HSA_STATUS_ERROR;
// Creating a new tracker entry
@@ -108,10 +108,12 @@ class Tracker {
// Creating a proxy signal
if (proxy) {
entry->is_proxy = true;
const hsa_signal_value_t signal_value = (orig.handle) ? hsa_api_.hsa_signal_load_relaxed(orig) : 1;
const hsa_signal_value_t signal_value =
(orig.handle) ? hsa_api_.hsa_signal_load_relaxed(orig) : 1;
status = hsa_api_.hsa_signal_create(signal_value, 0, NULL, &(entry->signal));
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_signal_create");
status = hsa_api_.hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT, signal_value, Handler, entry);
status = hsa_api_.hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT,
signal_value, Handler, entry);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
}
@@ -128,7 +130,8 @@ class Tracker {
hsa_signal_t& dispatch_signal = group->GetDispatchSignal();
hsa_signal_t& handler_signal = group->GetBarrierSignal();
entry->signal = dispatch_signal;
hsa_status_t status = hsa_api_.hsa_amd_signal_async_handler(handler_signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
hsa_status_t status = hsa_api_.hsa_amd_signal_async_handler(
handler_signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
}
@@ -150,7 +153,8 @@ class Tracker {
// Debug trace
if (trace_on_) {
auto outstanding = outstanding_.fetch_add(1);
fprintf(stdout, "Tracker::Enable: entry %p, record %p, outst %lu\n", entry, entry->record, outstanding);
fprintf(stdout, "Tracker::Enable: entry %p, record %p, outst %lu\n", entry, entry->record,
outstanding);
fflush(stdout);
}
}
@@ -173,12 +177,14 @@ class Tracker {
group->GetRecord()->dispatch = util::HsaRsrcFactory::Instance().TimestampNs();
// Creating a proxy signal
const hsa_signal_value_t signal_value = (orig_signal.handle) ?
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_load_relaxed(orig_signal) : 1;
const hsa_signal_value_t signal_value = (orig_signal.handle)
? util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_load_relaxed(orig_signal)
: 1;
hsa_signal_t& dispatch_signal = group->GetDispatchSignal();
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(dispatch_signal, signal_value);
hsa_status_t status =
util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_signal_async_handler(dispatch_signal, HSA_SIGNAL_CONDITION_LT, signal_value, Handler_opt, group);
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(dispatch_signal,
signal_value);
hsa_status_t status = util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_signal_async_handler(
dispatch_signal, HSA_SIGNAL_CONDITION_LT, signal_value, Handler_opt, group);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
}
@@ -190,7 +196,8 @@ class Tracker {
record_t* record = group->GetRecord();
hsa_amd_profiling_dispatch_time_t dispatch_time{};
hsa_status_t status =
util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_profiling_get_dispatch_time(context->GetAgent(), dispatch_signal, &dispatch_time);
util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_profiling_get_dispatch_time(
context->GetAgent(), dispatch_signal, &dispatch_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
record->begin = util::HsaRsrcFactory::Instance().SysclockToNs(dispatch_time.start);
record->end = util::HsaRsrcFactory::Instance().SysclockToNs(dispatch_time.end);
@@ -203,22 +210,23 @@ class Tracker {
amd_signal_t* prof_signal_ptr = reinterpret_cast<amd_signal_t*>(dispatch_signal.handle);
orig_signal_ptr->start_ts = prof_signal_ptr->start_ts;
orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(orig_signal, signal_value);
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(orig_signal,
signal_value);
}
return Context::Handler(signal_value, arg);
}
private:
Tracker() :
outstanding_(0),
hsa_rsrc_(&(util::HsaRsrcFactory::Instance())),
hsa_api_(*(hsa_rsrc_->HsaApi()))
{}
private:
Tracker()
: outstanding_(0),
hsa_rsrc_(&(util::HsaRsrcFactory::Instance())),
hsa_api_(*(hsa_rsrc_->HsaApi())) {}
~Tracker() {
if (trace_on_) {
fprintf(stdout, "Tracker::DESTR: sig list %d, outst %lu\n", (int)(sig_list_.size()), outstanding_.load());
fprintf(stdout, "Tracker::DESTR: sig list %d, outst %lu\n", (int)(sig_list_.size()),
outstanding_.load());
fflush(stdout);
}
@@ -226,8 +234,8 @@ class Tracker {
auto end = sig_list_.end();
while (it != end) {
auto cur = it++;
// The wait should be optiona as there possible some inter kernel dependencies and it possible to wait for
// the kernels will never be lunched as the application was finished by some reason.
// The wait should be optiona as there possible some inter kernel dependencies and it possible to
// wait for the kernels will never be lunched as the application was finished by some reason.
#if 0
// FIXME: currently the signal value for tracking signals are taken from original application signal
hsa_rsrc_->SignalWait((*cur)->signal, 1);
@@ -246,20 +254,24 @@ class Tracker {
// Debug trace
if (trace_on_) {
auto outstanding = outstanding_.fetch_sub(1);
fprintf(stdout, "Tracker::Complete: entry %p, record %p, outst %lu\n", entry, entry->record, outstanding);
fprintf(stdout, "Tracker::Complete: entry %p, record %p, outst %lu\n", entry, entry->record,
outstanding);
fflush(stdout);
}
// Query begin/end and complete timestamps
if (entry->is_memcopy) {
hsa_amd_profiling_async_copy_time_t async_copy_time{};
hsa_status_t status = hsa_api_.hsa_amd_profiling_get_async_copy_time(entry->signal, &async_copy_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_async_copy_time");
hsa_status_t status =
hsa_api_.hsa_amd_profiling_get_async_copy_time(entry->signal, &async_copy_time);
if (status != HSA_STATUS_SUCCESS)
EXC_RAISING(status, "hsa_amd_profiling_get_async_copy_time");
record->begin = hsa_rsrc_->SysclockToNs(async_copy_time.start);
record->end = hsa_rsrc_->SysclockToNs(async_copy_time.end);
} else {
hsa_amd_profiling_dispatch_time_t dispatch_time{};
hsa_status_t status = hsa_api_.hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
hsa_status_t status =
hsa_api_.hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
record->begin = hsa_rsrc_->SysclockToNs(dispatch_time.start);
record->end = hsa_rsrc_->SysclockToNs(dispatch_time.end);
@@ -349,6 +361,6 @@ class Tracker {
static const bool trace_on_ = false;
};
} // namespace rocprofiler
} // namespace rocprofiler
#endif // SRC_CORE_TRACKER_H_
#endif // SRC_CORE_TRACKER_H_
projects/rocprofiler/src/core/types.h
+4 -3
Bestand weergeven
@@ -36,11 +36,12 @@ typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
typedef uint32_t packet_word_t;
typedef uint64_t timestamp_t;
inline std::ostream& operator<< (std::ostream& out, const event_t& event) {
out << "[block_name(" << event.block_name << "). block_index(" << event.block_index << "). counter_id(" << event.counter_id << ")]";
inline std::ostream& operator<<(std::ostream& out, const event_t& event) {
out << "[block_name(" << event.block_name << "). block_index(" << event.block_index
<< "). counter_id(" << event.counter_id << ")]";
return out;
}
inline std::ostream& operator<< (std::ostream& out, const parameter_t& parameter) {
inline std::ostream& operator<<(std::ostream& out, const parameter_t& parameter) {
out << "[parameter_name(" << parameter.parameter_name << "). value(" << parameter.value << ")]";
return out;
}