diff --git a/rocrtst/common/base_rocr_utils.cc b/rocrtst/common/base_rocr_utils.cc
index 14a87a168a..298b646862 100755
--- a/rocrtst/common/base_rocr_utils.cc
+++ b/rocrtst/common/base_rocr_utils.cc
@@ -438,6 +438,35 @@ hsa_kernel_dispatch_packet_t * WriteAQLToQueue(BaseRocR* test, uint64_t *ind) {
   return queue_aql_packet;
 }
 
+void
+WriteAQLToQueueLoc(hsa_queue_t *queue, uint64_t indx,
+                                      hsa_kernel_dispatch_packet_t *aql_pkt) {
+  assert(queue);
+  assert(aql_pkt);
+
+  void *queue_base = queue->base_address;
+  const uint32_t queue_mask = queue->size - 1;
+  hsa_kernel_dispatch_packet_t* queue_aql_packet;
+
+  queue_aql_packet =
+       &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(queue_base))
+                                                        [indx & queue_mask];
+
+  queue_aql_packet->workgroup_size_x = aql_pkt->workgroup_size_x;
+  queue_aql_packet->workgroup_size_y = aql_pkt->workgroup_size_y;
+  queue_aql_packet->workgroup_size_z = aql_pkt->workgroup_size_z;
+  queue_aql_packet->grid_size_x = aql_pkt->grid_size_x;
+  queue_aql_packet->grid_size_y = aql_pkt->grid_size_y;
+  queue_aql_packet->grid_size_z = aql_pkt->grid_size_z;
+  queue_aql_packet->private_segment_size =
+                                     aql_pkt->private_segment_size;
+  queue_aql_packet->group_segment_size =
+                                       aql_pkt->group_segment_size;
+  queue_aql_packet->kernel_object = aql_pkt->kernel_object;
+  queue_aql_packet->kernarg_address = aql_pkt->kernarg_address;
+  queue_aql_packet->completion_signal = aql_pkt->completion_signal;
+}
+
 // Allocate a buffer in the kern_arg_pool for the kernel arguments and write
 // the arguments to buffer
 hsa_status_t AllocAndSetKernArgs(BaseRocR* test, void* args, size_t arg_size) {
diff --git a/rocrtst/common/base_rocr_utils.h b/rocrtst/common/base_rocr_utils.h
index 55cd7051f5..c86d6f42d2 100755
--- a/rocrtst/common/base_rocr_utils.h
+++ b/rocrtst/common/base_rocr_utils.h
@@ -98,6 +98,8 @@ hsa_status_t InitializeAQLPacket(const BaseRocR* test,
 /// \returns Pointer to dispatch packet in queue that was written to
 hsa_kernel_dispatch_packet_t* WriteAQLToQueue(BaseRocR* test, uint64_t *ind);
 
+void WriteAQLToQueueLoc(hsa_queue_t *queue, uint64_t indx,
+                                      hsa_kernel_dispatch_packet_t *aql_pkt);
 /// This function writes the first 32 bits of an aql packet to the provided
 /// aql packet. This function is meant to be called immediately before
 /// ringing door_bell signal.
diff --git a/rocrtst/suites/functional/debug_basic.cc b/rocrtst/suites/functional/debug_basic.cc
index 5bd009b538..13474022d3 100755
--- a/rocrtst/suites/functional/debug_basic.cc
+++ b/rocrtst/suites/functional/debug_basic.cc
@@ -70,15 +70,6 @@ typedef struct test_debug_data_t {
 
 static void TestDebugTrap(hsa_status_t status, hsa_queue_t *source, void *data);
 
-// This wrapper atomically writes the provided header and setup to the
-// provided AQL packet. The provided AQL packet address should be in the
-// queue memory space.
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
 #define RET_IF_HSA_ERR(err) { \
   if ((err) != HSA_STATUS_SUCCESS) { \
     const char* msg = 0; \
@@ -166,31 +157,6 @@ static void PrintDebugSubtestHeader(const char *header) {
   std::cout << "  *** Debug Basic Subtest: " << header << " ***" << std::endl;
 }
 
-void WriteAQLToQueue(hsa_kernel_dispatch_packet_t const* in_aql,
-                     hsa_queue_t* q) {
-  void* queue_base = q->base_address;
-  const uint32_t queue_mask = q->size - 1;
-  uint64_t que_idx = hsa_queue_add_write_index_relaxed(q, 1);
-
-  hsa_kernel_dispatch_packet_t* queue_aql_packet;
-
-  queue_aql_packet =
-    &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(queue_base))
-    [que_idx & queue_mask];
-
-  queue_aql_packet->workgroup_size_x = in_aql->workgroup_size_x;
-  queue_aql_packet->workgroup_size_y = in_aql->workgroup_size_y;
-  queue_aql_packet->workgroup_size_z = in_aql->workgroup_size_z;
-  queue_aql_packet->grid_size_x = in_aql->grid_size_x;
-  queue_aql_packet->grid_size_y = in_aql->grid_size_y;
-  queue_aql_packet->grid_size_z = in_aql->grid_size_z;
-  queue_aql_packet->private_segment_size = in_aql->private_segment_size;
-  queue_aql_packet->group_segment_size = in_aql->group_segment_size;
-  queue_aql_packet->kernel_object = in_aql->kernel_object;
-  queue_aql_packet->kernarg_address = in_aql->kernarg_address;
-  queue_aql_packet->completion_signal = in_aql->completion_signal;
-}
-
 void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent,
                                             hsa_agent_t gpuAgent) {
   hsa_status_t err;
@@ -323,9 +289,9 @@ void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent,
   // write to command queue
   uint64_t index = hsa_queue_load_write_index_relaxed(queue);
 
-  // This function simply copies the data we've collected so far into our
-  // local AQL packet, except the the setup and header fields.
-  WriteAQLToQueue(&aql, queue);
+  hsa_queue_store_write_index_relaxed(queue, index + 1);
+
+  rocrtst::WriteAQLToQueueLoc(queue, index, &aql);
 
   uint32_t aql_header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
   aql_header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -334,12 +300,10 @@ void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent,
                 HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
 
   void* q_base = queue->base_address;
-  AtomicSetPacketHeader(aql_header, aql.setup,
+  rocrtst::AtomicSetPacketHeader(aql_header, aql.setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
-  hsa_queue_store_write_index_relaxed(queue, index + 1);
-
   // ringdoor bell
   hsa_signal_store_relaxed(queue->doorbell_signal, index);
 
diff --git a/rocrtst/suites/functional/ipc.cc b/rocrtst/suites/functional/ipc.cc
index 8ed1a06924..c88bcfe886 100755
--- a/rocrtst/suites/functional/ipc.cc
+++ b/rocrtst/suites/functional/ipc.cc
@@ -243,15 +243,6 @@ void IPCTest::SetUp(void) {
   return;
 }
 
-// This wrapper atomically writes the provided header and setup to the
-// provided AQL packet. The provided AQL packet address should be in the
-// queue memory space.
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
 // Do a few extra iterations as we toss out some of the inital and final
 // iterations when calculating statistics
 uint32_t IPCTest::RealIterationNum(void) {
diff --git a/rocrtst/suites/functional/memory_access.cc b/rocrtst/suites/functional/memory_access.cc
index ce26127958..69869d9543 100755
--- a/rocrtst/suites/functional/memory_access.cc
+++ b/rocrtst/suites/functional/memory_access.cc
@@ -281,8 +281,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent,
     err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
     ASSERT_EQ(err, HSA_STATUS_SUCCESS);
 
-
-
     // Fill the dispatch packet with
     // workgroup_size, grid_size, kernelArgs and completion signal
     // Put it on the queue and launch the kernel by ringing the doorbell
@@ -296,12 +294,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent,
     memset(&aql, 0, sizeof(aql));
 
     // initialize aql packet
-    aql.header =
-      (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
-      (1 << HSA_PACKET_HEADER_BARRIER) |
-      (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
-      (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
-    aql.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
     aql.workgroup_size_x = 256;
     aql.workgroup_size_y = 1;
     aql.workgroup_size_z = 1;
@@ -319,10 +311,21 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent,
 
     // write to command queue
     uint64_t index = hsa_queue_load_write_index_relaxed(queue);
-    reinterpret_cast<hsa_kernel_dispatch_packet_t*>
-                              (queue->base_address)[index & queue_mask] = aql;
     hsa_queue_store_write_index_relaxed(queue, index + 1);
 
+    rocrtst::WriteAQLToQueueLoc(queue, index, &aql);
+
+    hsa_kernel_dispatch_packet_t *q_base_addr =
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>(queue->base_address);
+    rocrtst::AtomicSetPacketHeader(
+        (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+           (1 << HSA_PACKET_HEADER_BARRIER) |
+          (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+           (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE),
+                  (1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                          (&q_base_addr[index & queue_mask]));
+
     // ringdoor bell
     hsa_signal_store_relaxed(queue->doorbell_signal, index);
     // wait for the signal and reset it for future use
@@ -367,9 +370,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent,
   }
 }
 
-
-
-
 // Test to check cpu can read & write to GPU memory
 void MemoryAccessTest::CPUAccessToGPUMemoryTest(hsa_agent_t cpuAgent,
                                                  hsa_agent_t gpuAgent,
diff --git a/rocrtst/suites/functional/memory_allocation.cc b/rocrtst/suites/functional/memory_allocation.cc
index 098593c000..ae52fa6580 100755
--- a/rocrtst/suites/functional/memory_allocation.cc
+++ b/rocrtst/suites/functional/memory_allocation.cc
@@ -64,17 +64,6 @@ static const uint32_t kNumBufferElements = 256;
 static const int kValue = 5;
 
 
-// This wrapper atomically writes the provided header and setup to the
-// provided AQL packet. The provided AQL packet address should be in the
-// queue memory space.
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
-
-
 MemoryAllocationTest::MemoryAllocationTest(bool launch_GroupMemory,
                                            bool launch_BasicAllocateFree) : TestBase() {
   set_num_iteration(10);  // Number of iterations to execute of the main test;
@@ -179,9 +168,6 @@ static void PrintMemorySubtestHeader(const char *header) {
 
 static const int kMemoryAllocSize = 1024;
 
-
-
-
 void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent,
                                                    hsa_agent_t gpuAgent) {
   hsa_status_t err;
@@ -312,11 +298,11 @@ void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent,
 
     // Load index for writing header later to command queue at same index
     uint64_t index = hsa_queue_load_write_index_relaxed(queue);
+    hsa_queue_store_write_index_relaxed(queue, index + 1);
 
     // This function simply copies the data we've collected so far into our
     // local AQL packet, except the the setup and header fields.
-    WriteAQLPktToQueue(queue);
-
+    rocrtst::WriteAQLToQueueLoc(queue, index, &aql());
 
     aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
     aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -326,11 +312,10 @@ void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent,
 
     void* q_base = queue->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(aql().header, aql().setup,
+    rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
-
     // ringdoor bell
     hsa_signal_store_relaxed(queue->doorbell_signal, index);
 
diff --git a/rocrtst/suites/functional/memory_atomics.cc b/rocrtst/suites/functional/memory_atomics.cc
index 0fc498f182..25e1745fb0 100755
--- a/rocrtst/suites/functional/memory_atomics.cc
+++ b/rocrtst/suites/functional/memory_atomics.cc
@@ -63,18 +63,6 @@
 static const uint32_t kNumBufferElements = 256;
 static const int kValue = 5;
 
-
-// This wrapper atomically writes the provided header and setup to the
-// provided AQL packet. The provided AQL packet address should be in the
-// queue memory space.
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
-
-
 MemoryAtomic::MemoryAtomic(AtomicTest testtype) :
     TestBase() {
   set_num_iteration(10);  // Number of iterations to execute of the main test;
@@ -186,19 +174,6 @@ void MemoryAtomic::Close() {
   TestBase::Close();
 }
 
-void MemoryAtomic::WriteAQLPktToQueue(hsa_queue_t* q) {
-  void* queue_base = q->base_address;
-  const uint32_t queue_mask = q->size - 1;
-  uint64_t index = hsa_queue_add_write_index_relaxed(q, 1);
-
-      reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                                     queue_base)[index & queue_mask] = aql();
-}
-
-
-
-
-
 typedef struct  __attribute__ ((aligned(16)))  args_t {
   int *a;
   int *b;
@@ -207,8 +182,6 @@ typedef struct  __attribute__ ((aligned(16)))  args_t {
   int n;
   } args;
 
-
-
 static const char kSubTestSeparator[] = "  **************************";
 
 
@@ -473,11 +446,9 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent,
 
   // Load index for writing header later to command queue at same index
   uint64_t index = hsa_queue_load_write_index_relaxed(queue);
+  hsa_queue_store_write_index_relaxed(queue, index + 1);
 
-  // This function simply copies the data we've collected so far into our
-  // local AQL packet, except the the setup and header fields.
-  WriteAQLPktToQueue(queue);
-
+  rocrtst::WriteAQLToQueueLoc(queue, index, &aql());
 
   aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
   aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -487,7 +458,7 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent,
 
   void* q_base = queue->base_address;
   // Set the Aql packet header
-  AtomicSetPacketHeader(aql().header, aql().setup,
+  rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                       &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                           (q_base))[index & queue_mask]);
 
@@ -554,8 +525,6 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent,
   }
 }
 
-
-
 void MemoryAtomic::MemoryAtomicTest(void) {
   hsa_status_t err;
   // find all cpu agents
diff --git a/rocrtst/suites/functional/signal_kernel.cc b/rocrtst/suites/functional/signal_kernel.cc
old mode 100644
new mode 100755
index 97d41be03e..bbf991be03
--- a/rocrtst/suites/functional/signal_kernel.cc
+++ b/rocrtst/suites/functional/signal_kernel.cc
@@ -66,12 +66,6 @@ static unsigned int NumOfKernels = 1;
   } \
 }
 
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
 SignalKernelTest::SignalKernelTest(SignalKernelType type_) : TestBase() {
   set_num_iteration(10);  // Number of iterations to execute of the main test;
             // This is a default value which can be overridden
@@ -263,10 +257,11 @@ void SignalKernelTest::KernelSetFunction(SignalKernelType type_) {
 
     // write to command queue
     uint64_t index = hsa_queue_load_write_index_relaxed(queue);
-    reinterpret_cast<hsa_kernel_dispatch_packet_t*>
-        (queue->base_address)[index & queue_mask] = dispatch_packet;
     hsa_queue_store_write_index_relaxed(queue, index + 1);
 
+    rocrtst::WriteAQLToQueueLoc(queue, index, &dispatch_packet);
+
+
     dispatch_packet.header |= HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
     dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
     dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
@@ -275,7 +270,7 @@ void SignalKernelTest::KernelSetFunction(SignalKernelType type_) {
 
     void* q_base = queue->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup,
+    rocrtst::AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
@@ -446,10 +441,11 @@ void SignalKernelTest::TestSignalKernelMultiWait(void) {
     const uint32_t queue_mask = queue->size - 1;
     // write to command queue
     uint64_t index = hsa_queue_load_write_index_relaxed(queue);
-    reinterpret_cast<hsa_kernel_dispatch_packet_t*>
-        (queue->base_address)[index & queue_mask] = dispatch_packet;
     hsa_queue_store_write_index_relaxed(queue, index + 1);
 
+    rocrtst::WriteAQLToQueueLoc(queue, index, &dispatch_packet);
+
+
     dispatch_packet.header |= HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
     dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
     dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
@@ -458,7 +454,7 @@ void SignalKernelTest::TestSignalKernelMultiWait(void) {
 
     void* q_base = queue->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup,
+    rocrtst::AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
diff --git a/rocrtst/suites/negative/queue_validation.cc b/rocrtst/suites/negative/queue_validation.cc
index 896a2cf5c6..24e96f18e7 100755
--- a/rocrtst/suites/negative/queue_validation.cc
+++ b/rocrtst/suites/negative/queue_validation.cc
@@ -70,16 +70,6 @@ typedef struct test_validation_data_t {
 
 static void CallbackQueueErrorHandling(hsa_status_t status, hsa_queue_t *source, void *data);
 
-// This wrapper atomically writes the provided header and setup to the
-// provided AQL packet. The provided AQL packet address should be in the
-// queue memory space.
-static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
-                                  hsa_kernel_dispatch_packet_t* queue_packet) {
-  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
-                   header | (setup << 16), __ATOMIC_RELEASE);
-}
-
-
 QueueValidation::QueueValidation(bool launch_InvalidDimension,
                                  bool launch_InvalidGroupMemory,
                                  bool launch_InvalidKernelObject,
@@ -182,16 +172,6 @@ static void PrintDebugSubtestHeader(const char *header) {
   std::cout << "  *** QueueValidation Subtest: " << header << " ***" << std::endl;
 }
 
-void QueueValidation::WriteAQLPktToQueue(hsa_queue_t* q) {
-  void* queue_base = q->base_address;
-  const uint32_t queue_mask = q->size - 1;
-  uint64_t index = hsa_queue_add_write_index_relaxed(q, 1);
-
-      reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                                     queue_base)[index & queue_mask] = aql();
-}
-
-
 void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
                                             hsa_agent_t gpuAgent) {
   hsa_status_t err;
@@ -234,11 +214,9 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
 
     // Load index for writing header later to command queue at same index
     uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]);
+    hsa_queue_store_write_index_relaxed(queue[ii], index + 1);
 
-    // This function simply copies the data we've collected so far into our
-    // local AQL packet, except the the setup and header fields.
-    WriteAQLPktToQueue(queue[ii]);
-
+    rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql());
 
     aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
     aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -248,7 +226,7 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
 
     void* q_base = queue[ii]->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(aql().header, aql().setup,
+    rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
@@ -321,11 +299,9 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent,
 
     // Load index for writing header later to command queue at same index
     uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]);
+    hsa_queue_store_write_index_relaxed(queue[ii], index + 1);
 
-    // This function simply copies the data we've collected so far into our
-    // local AQL packet, except the the setup and header fields.
-    WriteAQLPktToQueue(queue[ii]);
-
+    rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql());
 
     aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
     aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -335,7 +311,7 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent,
 
     void* q_base = queue[ii]->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(aql().header, aql().setup,
+    rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
@@ -406,11 +382,9 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent
 
     // Load index for writing header later to command queue at same index
     uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]);
+    hsa_queue_store_write_index_relaxed(queue[ii], index + 1);
 
-    // This function simply copies the data we've collected so far into our
-    // local AQL packet, except the the setup and header fields.
-    WriteAQLPktToQueue(queue[ii]);
-
+    rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql());
 
     aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
     aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
@@ -420,7 +394,7 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent
 
     void* q_base = queue[ii]->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(aql().header, aql().setup,
+    rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
@@ -488,11 +462,9 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent,
 
     // Load index for writing header later to command queue at same index
     uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]);
+    hsa_queue_store_write_index_relaxed(queue[ii], index + 1);
 
-    // This function simply copies the data we've collected so far into our
-    // local AQL packet, except the the setup and header fields.
-    WriteAQLPktToQueue(queue[ii]);
-
+    rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql());
     // setting the invalid packet type
     aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
     aql().header |=  -1 << HSA_PACKET_HEADER_TYPE;
@@ -500,7 +472,7 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent,
 
     void* q_base = queue[ii]->base_address;
     // Set the Aql packet header
-    AtomicSetPacketHeader(aql().header, aql().setup,
+    rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                         &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                             (q_base))[index & queue_mask]);
 
@@ -577,11 +549,9 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen
 
       // Load index for writing header later to command queue at same index
       uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]);
+      hsa_queue_store_write_index_relaxed(queue[ii], index + 1);
 
-      // This function simply copies the data we've collected so far into our
-      // local AQL packet, except the the setup and header fields.
-      WriteAQLPktToQueue(queue[ii]);
-
+      rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql());
       aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
       aql().header |= HSA_FENCE_SCOPE_SYSTEM <<
                     HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
@@ -590,7 +560,7 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen
 
       void* q_base = queue[ii]->base_address;
       // Set the Aql packet header
-      AtomicSetPacketHeader(aql().header, aql().setup,
+      rocrtst::AtomicSetPacketHeader(aql().header, aql().setup,
                           &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
                           (q_base))[index & queue_mask]);
 
diff --git a/rocrtst/suites/negative/queue_validation.h b/rocrtst/suites/negative/queue_validation.h
index c72531888e..5ce9f8cf9e 100755
--- a/rocrtst/suites/negative/queue_validation.h
+++ b/rocrtst/suites/negative/queue_validation.h
@@ -98,7 +98,6 @@ class QueueValidation : public TestBase {
   void QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent);
   void QueueValidationForInvalidPacket(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent);
   void QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent);
-  void WriteAQLPktToQueue(hsa_queue_t* q);
 };
 
 #endif  // ROCRTST_SUITES_NEGATIVE_QUEUE_VALIDATION_H_
diff --git a/rocrtst/suites/performance/dispatch_time.cc b/rocrtst/suites/performance/dispatch_time.cc
index 7184ae55c3..0e4ca03cad 100755
--- a/rocrtst/suites/performance/dispatch_time.cc
+++ b/rocrtst/suites/performance/dispatch_time.cc
@@ -173,31 +173,37 @@ size_t DispatchTime::RealIterationNum() {
 void DispatchTime::RunSingle() {
   std::vector<double> timer;
 
-  int it = RealIterationNum();
+  uint32_t it = RealIterationNum();
   const uint32_t queue_mask = main_queue()->size - 1;
 
   // queue should be empty
   ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()),
             hsa_queue_load_write_index_scacquire(main_queue()));
 
-  void *q_base_addr = main_queue()->base_address;
-  for (int i = 0; i < it; i++) {
+  hsa_kernel_dispatch_packet_t *q_base_addr =
+      reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                                 (main_queue()->base_address);
+
+  if (it > main_queue()->size) {
+    it = main_queue()->size;
+  }
+  for (uint32_t i = 0; i < it; i++) {
     // Obtain the current queue write index.
     uint64_t index = hsa_queue_add_write_index_relaxed(main_queue(), 1);
 
-    ASSERT_LT(index, main_queue()->size + index);
-
     // Write the aql packet at the calculated queue index address.
+    rocrtst::WriteAQLToQueueLoc(main_queue(), index, &aql());
 
-    reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                                     q_base_addr)[index & queue_mask] = aql();
     // Get timing stamp and ring the doorbell to dispatch the kernel.
     rocrtst::PerfTimer p_timer;
     int id = p_timer.CreateTimer();
     p_timer.StartTimer(id);
-    reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                        q_base_addr)[index & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+
+    rocrtst::AtomicSetPacketHeader(
+        HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE,
+        aql().setup,
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                        (&(q_base_addr)[index & queue_mask]));
 
     hsa_signal_store_screlease(main_queue()->doorbell_signal, index);
 
@@ -206,7 +212,6 @@ void DispatchTime::RunSingle() {
          HSA_SIGNAL_CONDITION_LT, 1, (uint64_t) - 1, HSA_WAIT_STATE_ACTIVE)) {
     }
 
-
     p_timer.StopTimer(id);
 
     timer.push_back(p_timer.ReadTimer(id));
@@ -237,6 +242,9 @@ void DispatchTime::RunMulti() {
   std::vector<double> timer;
   int it = RealIterationNum();
   const uint32_t queue_mask = main_queue()->size - 1;
+  hsa_kernel_dispatch_packet_t *q_base_addr =
+      reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                                 (main_queue()->base_address);
 
   // queue should be empty
   ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()),
@@ -255,35 +263,37 @@ void DispatchTime::RunMulti() {
       index[j] = hsa_queue_add_write_index_relaxed(main_queue(), 1);
 
       // Write the aql packet at the calculated queue index address.
-      (reinterpret_cast<hsa_kernel_dispatch_packet_t*>((
-                 main_queue()->base_address)))[index[j] & queue_mask] = aql();
-
-      if (j == num_batch_ - 1) {
-        (reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-            main_queue()->base_address))[index[j] & queue_mask].header |=
-                                               1 << HSA_PACKET_HEADER_BARRIER;
-      }
+      rocrtst::WriteAQLToQueueLoc(main_queue(), index[j], &aql());
     }
 
+    rocrtst::AtomicSetPacketHeader(
+        (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+        (1 << HSA_PACKET_HEADER_BARRIER),
+        aql().setup,
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                          (&q_base_addr[index[num_batch_ - 1] & queue_mask]));
+
     // Set packet header reversly; set all headers except the very first
     // one, for now.
     for (uint32_t j = num_batch_ - 1; j > 0; j--) {
-      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-         (main_queue()->base_address))[index[j] & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+      rocrtst::AtomicSetPacketHeader(
+          HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE,
+          aql().setup,
+          reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                      (&q_base_addr[index[j] & queue_mask]));
     }
 
     // Get timing stamp and ring the doorbell to dispatch the kernel.
     int id = p_timer.CreateTimer();
     p_timer.StartTimer(id);
     // Set the very first header...
-    (reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-        main_queue()->base_address))[index[0] & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+    rocrtst::AtomicSetPacketHeader(
+        HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE,
+        aql().setup,
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                     (&(q_base_addr)[index[0] & queue_mask]));
 
-    for (uint32_t j = 0; j < num_batch_; j++) {
-      hsa_signal_store_screlease(main_queue()->doorbell_signal, index[j]);
-    }
+    hsa_signal_store_screlease(main_queue()->doorbell_signal, index[num_batch_ - 1]);
 
     // Wait on the dispatch signal until the kernel is finished.
     while (hsa_signal_wait_scacquire(aql().completion_signal,
diff --git a/rocrtst/suites/performance/enqueueLatency.cc b/rocrtst/suites/performance/enqueueLatency.cc
index aef00bd7c3..282d9d9475 100755
--- a/rocrtst/suites/performance/enqueueLatency.cc
+++ b/rocrtst/suites/performance/enqueueLatency.cc
@@ -184,7 +184,9 @@ void EnqueueLatency::EnqueueSinglePacket() {
   ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()),
             hsa_queue_load_write_index_scacquire(main_queue()));
 
-  void *q_base_addr = main_queue()->base_address;
+  hsa_kernel_dispatch_packet_t *q_base_addr =
+                      reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
+                                                  main_queue()->base_address);
   rocrtst::PerfTimer p_timer;
   for (int i = 0; i < it; i++) {
     // Get timing stamp and ring the doorbell to dispatch the kernel.
@@ -196,12 +198,13 @@ void EnqueueLatency::EnqueueSinglePacket() {
     ASSERT_LT(index, main_queue()->size + index);
 
     // Write the aql packet at the calculated queue index address.
-    reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                                     q_base_addr)[index & queue_mask] = aql();
+    rocrtst::WriteAQLToQueueLoc(main_queue(), index, &aql());
 
-    reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
-                        q_base_addr)[index & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+    rocrtst::AtomicSetPacketHeader(
+        HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE,
+        aql().setup,
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                     (&(q_base_addr)[index & queue_mask]));
 
     p_timer.StopTimer(id);
 
@@ -213,8 +216,6 @@ void EnqueueLatency::EnqueueSinglePacket() {
          HSA_SIGNAL_CONDITION_LT, 1, (uint64_t) - 1, HSA_WAIT_STATE_ACTIVE)) {
     }
 
-
-
     hsa_signal_store_screlease(aql().completion_signal, 1);
 
     if (verbosity() >= VERBOSE_PROGRESS) {
@@ -238,8 +239,6 @@ void EnqueueLatency::EnqueueSinglePacket() {
   return;
 }
 
-
-
 void EnqueueLatency::EnqueueMultiPackets() {
   std::vector<double> timer;
   int it = RealIterationNum();
@@ -251,6 +250,10 @@ void EnqueueLatency::EnqueueMultiPackets() {
 
   rocrtst::PerfTimer p_timer;
 
+  hsa_kernel_dispatch_packet_t *q_base_addr =
+                      reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
+                                                  main_queue()->base_address);
+
   for (int i = 0; i < it; i++) {
     // Get timing stamp and ring the doorbell to dispatch the kernel.
     int id = p_timer.CreateTimer();
@@ -265,30 +268,28 @@ void EnqueueLatency::EnqueueMultiPackets() {
       index[j] = hsa_queue_add_write_index_relaxed(main_queue(), 1);
 
       // Write the aql packet at the calculated queue index address.
-      (reinterpret_cast<hsa_kernel_dispatch_packet_t*>((
-                 main_queue()->base_address)))[index[j] & queue_mask] = aql();
-
-      if (j == num_of_pkts_ - 1) {
-        (reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-            main_queue()->base_address))[index[j] & queue_mask].header |=
-                                               1 << HSA_PACKET_HEADER_BARRIER;
-      }
+      rocrtst::WriteAQLToQueueLoc(main_queue(), index[j], &aql());
     }
+    // Write the aql packet at the calculated queue index address.
+
+    rocrtst::AtomicSetPacketHeader(
+        (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+        (1 << HSA_PACKET_HEADER_BARRIER),
+        aql().setup,
+        reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                      (&(q_base_addr)[index[num_of_pkts_ - 1] & queue_mask]));
+
 
     // Set packet header reversly; set all headers except the very first
     // one, for now.
-    for (uint32_t j = num_of_pkts_ - 1; j > 0; j--) {
-      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-         (main_queue()->base_address))[index[j] & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+    for (int32_t j = num_of_pkts_ - 1; j >= 0; j--) {
+      rocrtst::AtomicSetPacketHeader(
+          HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE,
+          aql().setup,
+          reinterpret_cast<hsa_kernel_dispatch_packet_t *>
+                                     (&(q_base_addr)[index[j] & queue_mask]));
     }
 
-
-    // Set the very first header...
-    (reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
-        main_queue()->base_address))[index[0] & queue_mask].header |=
-                    HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
-
     p_timer.StopTimer(id);
 
     timer.push_back(p_timer.ReadTimer(id));
diff --git a/rocrtst/suites/performance/memory_async_copy.cc b/rocrtst/suites/performance/memory_async_copy.cc
index 300744ac68..c054bc41ec 100755
--- a/rocrtst/suites/performance/memory_async_copy.cc
+++ b/rocrtst/suites/performance/memory_async_copy.cc
@@ -412,9 +412,12 @@ void MemoryAsyncCopy::DisplayResults(void) const {
   }
 
   TestBase::DisplayResults();
-
+  hsa_status_t err;
   for (Transaction t : tran_) {
     DisplayBenchmark(&t);
+    err = hsa_signal_destroy(t.signal);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
     delete t.benchmark_copy_time;
     delete t.min_time;
   }
diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc
index 264eb1d45f..1314683af9 100755
--- a/rocrtst/suites/test_common/main.cc
+++ b/rocrtst/suites/test_common/main.cc
@@ -126,7 +126,7 @@ TEST(rocrtst, Test_Example) {
   RunGenericTest(&tst);
 }
 
-TEST(rocrtstFunc, DISABLED_IPC) {
+TEST(rocrtstFunc, IPC) {
   IPCTest ipc;
   RunGenericTest(&ipc);
 }
@@ -454,22 +454,22 @@ TEST(rocrtstPerf, DISABLED_Memory_Async_Copy_NUMA) {
   RunGenericTest(&numa);
 }
 
-TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Single_SpinWait) {
+TEST(rocrtstPerf, AQL_Dispatch_Time_Single_SpinWait) {
   DispatchTime dt(true, true);
   RunGenericTest(&dt);
 }
 
-TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Single_Interrupt) {
+TEST(rocrtstPerf, AQL_Dispatch_Time_Single_Interrupt) {
   DispatchTime dt(false, true);
   RunGenericTest(&dt);
 }
 
-TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Multi_SpinWait) {
+TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_SpinWait) {
   DispatchTime dt(true, false);
   RunGenericTest(&dt);
 }
 
-TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Multi_Interrupt) {
+TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_Interrupt) {
   DispatchTime dt(false, false);
   RunGenericTest(&dt);
 }