From 2225153c23b8325e676779758b84a5a72e03819e Mon Sep 17 00:00:00 2001 From: Benjamin Welton Date: Wed, 29 May 2024 08:35:49 -0700 Subject: [PATCH] Test fix for AQLProfile changes (#882) * 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> --- source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp | 2 +- source/lib/rocprofiler-sdk/hsa/queue.cpp | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 0dfcdaca40..9e30fcfbfa 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -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; diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index 2d4b6bca9f..d8d0139e22 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -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); }