298 void packetStoreRelease(uint32_t *Packet, uint16_t Header, uint16_t Rest) {  in packetStoreRelease()  argument
299   __atomic_store_n(Packet, Header | (Rest << 16), __ATOMIC_RELEASE);  in packetStoreRelease()
1438     hsa_kernel_dispatch_packet_t *Packet =  in runRegionLocked()  local
1442     Packet->setup = UINT16_C(1) << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;  in runRegionLocked()
1443     Packet->workgroup_size_x = WorkgroupSize;  in runRegionLocked()
1444     Packet->workgroup_size_y = 1;  in runRegionLocked()
1445     Packet->workgroup_size_z = 1;  in runRegionLocked()
1446     Packet->reserved0 = 0;  in runRegionLocked()
1447     Packet->grid_size_x = GridSize;  in runRegionLocked()
1448     Packet->grid_size_y = 1;  in runRegionLocked()
1449     Packet->grid_size_z = 1;  in runRegionLocked()
1450     Packet->private_segment_size = KernelInfoEntry.private_segment_size;  in runRegionLocked()
1451     Packet->group_segment_size = GroupSegmentSize;  in runRegionLocked()
1452     Packet->kernel_object = KernelInfoEntry.kernel_object;  in runRegionLocked()
1453     Packet->kernarg_address = 0;     // use the block allocator  in runRegionLocked()
1454     Packet->reserved2 = 0;           // impl writes id_ here  in runRegionLocked()
1455     Packet->completion_signal = {0}; // may want a pool of signals  in runRegionLocked()
1530       Packet->kernarg_address = KernArg;  in runRegionLocked()
1538     Packet->completion_signal = S;  in runRegionLocked()
1539     hsa_signal_store_relaxed(Packet->completion_signal, 1);  in runRegionLocked()
1542     core::packetStoreRelease(reinterpret_cast<uint32_t *>(Packet),  in runRegionLocked()
1543                              core::createHeader(), Packet->setup);  in runRegionLocked()