[libc-commits] [libc] 9e390a1 - [libc][Obvious] Fix missing semicolon in AMDGPU loader implementation

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2023-10-30T14:58:46-05:00
New Revision: 9e390a140857355b3524924075a261b9d06ae850

URL: https://github.com/llvm/llvm-project/commit/9e390a140857355b3524924075a261b9d06ae850
DIFF: https://github.com/llvm/llvm-project/commit/9e390a140857355b3524924075a261b9d06ae850.diff

LOG: [libc][Obvious] Fix missing semicolon in AMDGPU loader implementation

Summary:
Title

Added: 
    

Modified: 
    libc/utils/gpu/loader/amdgpu/Loader.cpp

Removed: 
    


################################################################################
diff  --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp
index c2a11fd8aab72b4..2f99076a720e2aa 100644
--- a/libc/utils/gpu/loader/amdgpu/Loader.cpp
+++ b/libc/utils/gpu/loader/amdgpu/Loader.cpp
@@ -248,9 +248,8 @@ hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
       (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 header_word =
-      header | (setup << 16u) __atomic_store_n((uint32_t *)&packet->header,
-                                               header_word, __ATOMIC_RELEASE);
+  uint32_t header_word = header | (setup << 16u);
+  __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