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

@@ -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;

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