(Noticed whilst reviewing #6695)

From the docs for numba.cuda.atomic.compare_and_swap:

class compare_and_swap(ary, old, val)

Conditionally assign val to the first element of an 1D array ary if the current value matches old.
Returns the current value as if it is loaded atomically.

It seem

Just an FYI whilst I was trawling through the ROCm GitHub page:

https://rocmdocs.amd.com/en/latest/Programming_Guides/Programming-Guides.html#