[libc-commits] [libc] [openmp] [libomptarget][nextgen-plugin] Use SCRELEASE/SCACQUIRE in packet headers (PR #85678)
Gheorghe-Teodor Bercea via libc-commits
libc-commits at lists.llvm.org
Tue Mar 19 12:51:56 PDT 2024
https://github.com/doru1004 updated https://github.com/llvm/llvm-project/pull/85678
>From 14252f55620ac01d6f6d1d06fcbfcd9ddb9ddae6 Mon Sep 17 00:00:00 2001
From: Doru Bercea <doru.bercea at amd.com>
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/src/rtl.cpp | 23 ++++++++++---------
2 files changed, 14 insertions(+), 12 deletions(-)
diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp
index e3911eda2bd82a..7fd45acfd47e97 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<uint32_t *>(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/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
index fce7454bf2800d..c147cefe58e9b4 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<uint32_t *>(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<uint32_t *>(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<uint32_t *>(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<uint32_t *>(Packet), HeaderWord,
+ __ATOMIC_RELEASE);
// Signal the doorbell about the published packet.
hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId);
More information about the libc-commits
mailing list