CUDA mutexes for complicated reductions

[ianthomas23]

Currently datashader only has CUDA support for relatively simple Reduction classes. This is primarily due to the parallel nature of CUDA hardware which means we have to support multiple GPU threads writing to the same element (pixel) of an aggregation, so we only support reductions which have a corresponding CUDA atomic operation such as numba.cuda.atomic.max. For more complicated reductions such as the planned max_n reduction we will need to implement some form of mutex ourselves, to limit one thread at a time accessing the relevant shared data. The CUDA name for this is "atomicCAS" or "atomic compare and swap".

This will need some experimentation. The initial implementation could try to limit access to a single element (pixel) of one aggregation to a single thread at a time. Benchmarking will be important here to identify a workable and efficient solution.

Useful reference: https://towardsdatascience.com/cuda-by-numba-examples-c583474124b0