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); }