From 0dcc1ba5a7c6f7ee5f45381f325db7c0a4a982c0 Mon Sep 17 00:00:00 2001 From: Doru Bercea Date: Mon, 18 Mar 2024 13:52:32 -0400 Subject: [PATCH] Use SCRELEASE/SCACQUIRE in packet header --- libc/utils/gpu/loader/amdgpu/Loader.cpp | 3 ++- .../plugins-nextgen/amdgpu/dynamic_hsa/hsa.h | 4 ++-- .../plugins-nextgen/amdgpu/src/rtl.cpp | 23 ++++++++++--------- 3 files changed, 16 insertions(+), 14 deletions(-) diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp index e3911eda2bd82..7fd45acfd47e9 100644 --- a/libc/utils/gpu/loader/amdgpu/Loader.cpp +++ b/libc/utils/gpu/loader/amdgpu/Loader.cpp @@ -276,7 +276,8 @@ hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable, (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); uint32_t header_word = header | (setup << 16u); - __atomic_store_n((uint32_t *)&packet->header, header_word, __ATOMIC_RELEASE); + __atomic_store_n(reinterpret_cast(packet), header_word, + __ATOMIC_RELEASE); hsa_signal_store_relaxed(queue->doorbell_signal, packet_id); // Wait until the kernel has completed execution on the device. Periodically diff --git a/openmp/libomptarget/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/openmp/libomptarget/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h index 64a1d3308aed0..0ac7661d7e007 100644 --- a/openmp/libomptarget/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h +++ b/openmp/libomptarget/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h @@ -216,8 +216,8 @@ typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t; typedef enum { HSA_PACKET_HEADER_TYPE = 0, - HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9, - HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11 + HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE = 9, + HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE = 11 } hsa_packet_header_t; typedef enum { diff --git a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp index fce7454bf2800..c147cefe58e9b 100644 --- a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp @@ -826,15 +826,15 @@ struct AMDGPUQueueTy { /// Assumes the queue lock is acquired. void publishKernelPacket(uint64_t PacketId, uint16_t Setup, hsa_kernel_dispatch_packet_t *Packet) { - uint32_t *PacketPtr = reinterpret_cast(Packet); - - uint16_t Header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; - Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; - Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + uint16_t Header = + (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); // Publish the packet. Do not modify the package after this point. uint32_t HeaderWord = Header | (Setup << 16u); - __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE); + __atomic_store_n(reinterpret_cast(Packet), HeaderWord, + __ATOMIC_RELEASE); // Signal the doorbell about the published packet. hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId); @@ -845,15 +845,16 @@ struct AMDGPUQueueTy { /// barrier dependencies (signals) are satisfied. Assumes the queue is locked void publishBarrierPacket(uint64_t PacketId, hsa_barrier_and_packet_t *Packet) { - uint32_t *PacketPtr = reinterpret_cast(Packet); uint16_t Setup = 0; - uint16_t Header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; - Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; - Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + uint16_t Header = + (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); // Publish the packet. Do not modify the package after this point. uint32_t HeaderWord = Header | (Setup << 16u); - __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE); + __atomic_store_n(reinterpret_cast(Packet), HeaderWord, + __ATOMIC_RELEASE); // Signal the doorbell about the published packet. hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId);