Author: Gheorghe-Teodor Bercea Date: 2024-03-20T11:38:03-04:00 New Revision: 5e2b2342f2eba197bb1e555c1e70710c055f7df1
URL: https://github.com/llvm/llvm-project/commit/5e2b2342f2eba197bb1e555c1e70710c055f7df1 DIFF: https://github.com/llvm/llvm-project/commit/5e2b2342f2eba197bb1e555c1e70710c055f7df1.diff LOG: Revert "[libomptarget][nextgen-plugin] Use SCRELEASE/SCACQUIRE in packet head…" This reverts commit 927308a52bc51ae786db1bd645ad5ef5889fdb2a. Added: Modified: libc/utils/gpu/loader/amdgpu/Loader.cpp openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp Removed: ################################################################################ diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp index 7fd45acfd47e97..e3911eda2bd82a 100644 --- a/libc/utils/gpu/loader/amdgpu/Loader.cpp +++ b/libc/utils/gpu/loader/amdgpu/Loader.cpp @@ -276,8 +276,7 @@ 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(reinterpret_cast<uint32_t *>(packet), header_word, - __ATOMIC_RELEASE); + __atomic_store_n((uint32_t *)&packet->header, 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 c147cefe58e9b4..fce7454bf2800d 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) { - 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); + 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; // Publish the packet. Do not modify the package after this point. uint32_t HeaderWord = Header | (Setup << 16u); - __atomic_store_n(reinterpret_cast<uint32_t *>(Packet), HeaderWord, - __ATOMIC_RELEASE); + __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE); // Signal the doorbell about the published packet. hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId); @@ -845,16 +845,15 @@ 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) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); + 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; // Publish the packet. Do not modify the package after this point. uint32_t HeaderWord = Header | (Setup << 16u); - __atomic_store_n(reinterpret_cast<uint32_t *>(Packet), HeaderWord, - __ATOMIC_RELEASE); + __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE); // Signal the doorbell about the published packet. hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId); _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits