Skip to content

Commit

Permalink
Fix an issue with aync copy timestamps
Browse files Browse the repository at this point in the history
The timestamps coming from the HIP runtime for asynchronus memory
copies are corrupted (begin > end) because the HSA setting to record
timestamps is turned off by the tracer's HSA intercept.

The solution is to intercept hsa_amd_profiling_async_copy_enable and
remember the application/runtime's request so that it can be ORed with
IsEnabled(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY).

Change-Id: Ib687cbf36711563e86c2bb8bc934c7c51572bfde
  • Loading branch information
lmoriche committed Sep 28, 2022
1 parent b664937 commit 329c046
Showing 1 changed file with 15 additions and 5 deletions.
20 changes: 15 additions & 5 deletions src/roctracer/hsa_support.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -415,6 +415,14 @@ hsa_status_t ExecutableDestroyIntercept(hsa_executable_t executable) {
return saved_core_api.hsa_executable_destroy_fn(executable);
}

bool profiling_async_copy_enable = false;

hsa_status_t ProfilingAsyncCopyEnableIntercept(bool enable) {
hsa_status_t status = saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(enable);
if (status == HSA_STATUS_SUCCESS) profiling_async_copy_enable = enable;
return status;
}

void MemoryASyncCopyHandler(const Tracker::entry_t* entry) {
activity_record_t record{};
record.domain = ACTIVITY_DOMAIN_HSA_OPS;
Expand All @@ -433,8 +441,8 @@ hsa_status_t MemoryASyncCopyIntercept(void* dst, hsa_agent_t dst_agent, const vo
bool is_enabled = IsEnabled(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY);

// FIXME: what happens if the state changes before returning?
[[maybe_unused]] hsa_status_t status =
saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(is_enabled);
[[maybe_unused]] hsa_status_t status = saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(
profiling_async_copy_enable | is_enabled);
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");

if (!is_enabled) {
Expand Down Expand Up @@ -464,8 +472,8 @@ hsa_status_t MemoryASyncCopyRectIntercept(const hsa_pitched_ptr_t* dst,
bool is_enabled = IsEnabled(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY);

// FIXME: what happens if the state changes before returning?
[[maybe_unused]] hsa_status_t status =
saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(is_enabled);
[[maybe_unused]] hsa_status_t status = saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(
profiling_async_copy_enable | is_enabled);
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");

if (!is_enabled) {
Expand Down Expand Up @@ -557,6 +565,7 @@ void Initialize(HsaApiTable* table) {
// Install the HSA_OPS intercept
table->amd_ext_->hsa_amd_memory_async_copy_fn = MemoryASyncCopyIntercept;
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = MemoryASyncCopyRectIntercept;
table->amd_ext_->hsa_amd_profiling_async_copy_enable_fn = ProfilingAsyncCopyEnableIntercept;

// Install the HSA_EVT intercept
table->core_->hsa_memory_allocate_fn = MemoryAllocateIntercept;
Expand All @@ -575,7 +584,8 @@ void Initialize(HsaApiTable* table) {
}

void Finalize() {
if (hsa_status_t status = saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(false);
if (hsa_status_t status =
saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(profiling_async_copy_enable);
status != HSA_STATUS_SUCCESS)
assert(!"hsa_amd_profiling_async_copy_enable failed");

Expand Down

0 comments on commit 329c046

Please sign in to comment.