Implement atomic_ref<Integral>

[wmaxey]

Changelog Summary

Builds upon previous atomic refactor to create implementations of cuda::atomic_ref and cuda::std::atomic_ref

Related Work

#179 - Refactor included back-end pathways for atomic_ref types.

Details

This implementation conforms to the C++20 specification of atomic_ref.

Deviations

• cuda::std::atomic_ref is backported to C++11.
• cuda::atomic_ref has fetch_min/fetch_max overloads

Motivation

atomic_ref provides a modern abstraction for visibility and ordering of reads and writes to memory according to the C++ memory model. This interface can be used to replace uses of CUDA specific atomicOperation(_Scope) functions in device code and provides interoperability with host code.

Previously:

__global__ atomic_kernel(int *out) {
  // Calculate some work per thread
  int bucket = threadIdx.x%16; 
  // Threads maybe have a different bucket, e.g. when computing a histogram
  atomicAdd(out+bucket, result); // Can't use `atomic<int>` on raw pointers :(
}

With atomic_ref:

#include <cuda/atomic>

__global__ atomic_kernel(int *out) {
  // Calculate some work per thread
 cuda::atomic_ref<int, cuda::thread_scope_block> bucket(out[threadIdx.x%16]);

  // Each thread may have a different bucket, e.g. when computing a histogram
 bucket += result;
}

Design

Builds upon existing atomic pathways to create the atomic_ref front end. There were several changes required to properly cast into the correct lower level interface.

Tests for atomic_ref were replicated out of other atomic baseline tests, these provide coverage for both pointer and integral operations.

Host only libcxx tests were also added to provide coverage from a host standard library view and to ensure that the implementation is sound.

Testing

Virtuals:

• AUSDVS: http://ausdvs.nvidia.com/Build_Results?virtualId=1000987553&which_page=current_build
• SCDVS: http://scdvs.nvidia.com/Build_Results?virtualId=1000075758&which_page=current_build

New tests have been added that exercise all interfaces of atomic_ref

[jrhemstad]

This implementation conforms to the C++20 specification of atomic_ref.

What are we doing about padding bits? std::atomic_ref explicitly states it only compares the value representations.

[griwes]

This looks good to me; it's pretty much what I'd expect this to look like after the earlier refactor.

This probably needs a number of more tests, for I don't think all of them are crucial. What is definitely crucial is ref-to-pointer tests in the libc++ layer (similar to what you have for integral_ref already). Other than that, ship it once it's nicely squashed and tested!