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

[Joseph Huber via Openmp-commits]

jhuber6 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);
> ```

It is, I should probably update it as well. I think I originally copied part of that from the old plugins somewhere.

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