[Openmp-commits] [openmp] [libomptarget][nextgen-plugin] Use SCRELEASE/SCACQUIRE in packet headers (PR #85678)

[Gheorghe-Teodor Bercea via Openmp-commits]

doru1004 wrote:

> I don't think we need to worry about the header moving honestly, that would be an incredibly ABI breaking change and it's called header for a reason. We already rely implicitly on this storing the two 16-byte values that are right next to eachother.
> 
> Right now we have "store 32-bytes at the beginning of the struct" and this turns that into "Store 32-bytes at a 16-byte pointer assuming that the out of bounds access is correct." We also do not want to make this two-16 byte stores because this should happen in a single operation. So I think it's fine the way it is.
> 
> We should probably change the name though, the documentation says it's depcreated and just makes them the same value anyway so it's NFC. From the HSA headers:
> 
> ```
> HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE = 9,
> HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
> ```

This is your code no?

```
// Initialize the packet header and set the doorbell signal to begin execution
  // by the HSA runtime.
  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 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);
```

https://github.com/llvm/llvm-project/pull/85678