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

Test fix for AQLProfile changes (#882)

Browse Source 
* Test fix for AQLProfile changes

AQL profile recently changed where the ordering
of packets now needs to be read then stop (instead
of stop and read previously). This change supports
the newer model.

* Update source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp

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

* test barrier

---------

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

[ROCm/rocprofiler-sdk commit: 2225153c23]
This commit is contained in:
Benjamin Weltonandgithub-actions[bot] authored and GitHub committed 2024-05-29 10:35:49 -05:00
1 parent d1a5031722
commit cc44c48952
2 files changed
+4 -2
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp
+1 -1
View File
@@ -148,8 +148,8 @@ queue_cb(const context::context* ctx,
}
ret_pkt->before_krn_pkt.push_back(ret_pkt->start);
ret_pkt->after_krn_pkt.push_back(ret_pkt->stop);
ret_pkt->after_krn_pkt.push_back(ret_pkt->read);
ret_pkt->after_krn_pkt.push_back(ret_pkt->stop);
for(auto& aql_pkt : ret_pkt->after_krn_pkt)
{
aql_pkt.completion_signal.handle = 0;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp
+3 -1
View File
@@ -415,7 +415,9 @@ WriteInterceptor(const void* packets,
{
get_core_table()->hsa_signal_store_screlease_fn(interrupt_signal, 0);
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
barrier.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
barrier.completion_signal = interrupt_signal;
transformed_packets.emplace_back(barrier);
}