Enable user-provided lock table for `atomic_ref<T>`

I would like to be able to use `atomic_ref` for `sizeof(T) > 8B`, i.e., `atomic_ref::is_always_lock_free == false`.

Obviously this can't rely on built-in atomic operations.Typical implementations will make use of a lock table to support this usage where you use the address of the referenced object as a lookup into a table of mutexes.

Generic support for a lock table underneath a `cuda::std::atomic_ref` would be extremely non-trivial to support on all platforms and for it to work heterogeneously.

However, a less generic solution that would still be very useful would be to allow a user to provide their own lock table.

@ogiroux's idea is to partially specialize `cuda::atomic_ref` for `is_always_lock_free == false` to contain a pointer to a lock table to be supplied by the user via the `atomic_ref` constructor at each construction. It would likely be useful to supply an opaque `cuda::atomic_lock_table` type that a user could instantiate and manage however they like.

A rough sketch of what this could look like:
```
__managed__ cuda::atomic_lock_table<1024, thread_scope_device> table;

__global__ void kernel(int4 * i){
cuda::atomic_ref ref{i[0], table};
ref.atomic_exchange( int4{1, 2, 3, 4} ); // internally locks a mutex in `table`
}
```

This would depend on https://github.com/NVIDIA/cccl/issues/949 for implementing the `atomic_lock_table`.