From cc44c489520a752d426d10e19f03ffbaeef54b9b 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> [ROCm/rocprofiler-sdk commit: 2225153c23b8325e676779758b84a5a72e03819e] --- .../source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp | 2 +- .../rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 0dfcdaca40..9e30fcfbfa 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/projects/rocprofiler-sdk/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/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp index 2d4b6bca9f..d8d0139e22 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/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); }