I would like to use atomicSub but for floats and doubles. From what I understand of the documentation, this is not yet supported and the built in atomicSub only works with integers.

In the documentation they show how one can go about implementing atomicAdd for doubles for older GPUs with atomicCAS(), therefore I thought I would be able to easily implement my own version of atomicSub for floats by modifying it like this:

```
__device__ double atomicSubFloat(float* address, float val)
{
unsigned long long int* address_as_ull =
(unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
__double_as_longlong(val -
__longlong_as_double(assumed)));
// Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
} while (assumed != old);
return __longlong_as_double(old);
}
```

But it doesnâ€™t seem to work. Presumably â€ś__double_as_longlongâ€ť and â€ť __longlong_as_doubleâ€ť should be something else when working with floats, but I donâ€™t know what to replace them with. Does anyone know how to use atomicSub for floats and doubles?