Fixes issue with Numba CUDA target altering max/min semantics.

[stuartarchibald]

As title.

Closes #932

[stuartarchibald]

This is entirely untested and relies on numba/numba#5941, open question about what to do for version Numba 0.50. which has neither the old atomic min/max behaviour nor the new nanmin/nanmax functions.

[jbednar]

Thanks! I haven't tested it, but hopefully @jonmmease and/or @philippjfr can.

[jbednar]

Suggested change

	# WHAT DO YOU WANT TO DO HERE?
	raise ImportError("Datashader's CUDA support requires numba!=0.50.0")

[jbednar]

Suggested change

	cuda_atomic_max(data, (i, j), lower)
	cuda_atomic_min(data, (i, j), upper)
	cuda_atomic_nanmax(data, (i, j), lower)
	cuda_atomic_nanmin(data, (i, j), upper)

[stuartarchibald]

Thanks for the review @jbednar, once 0.51 is out I'll revisit (or someone else can, I have no particular attachment to the patch, feel free to push to my branch :)) with view to test this and make it actually work opposed to it just being a guess!

[dHannasch]

Thanks! I haven't tested it, but hopefully @jonmmease and/or @philippjfr can.

If anyone's up for testing this, I've fixed the picky little errors at #940. (See stuartarchibald#1.)
(https://github.com/dHannasch/datashader.git@fix/numba_cuda_atomic_minmax_test_with_branch also contains some monkeying with .travis.yml, but if you're going to pip install to test manually, I don't think that will get in your way.)
You can install the prerequisite branch of Numba with pip install git+https://github.com/stuartarchibald/numba.git@fix/cuda_atomic_nanminmax.

Note however that #940 has one more change that might be important. @philippjfr asked why we couldn't just import numba.cuda.atomic.nanmin and numba.cuda.atomic.nanmax directly (that is, why do we need to wrap them in @cuda.jit defs?), and I didn't have an answer, so I tried splitting the difference and using them directly in reductions.py _append_cuda, but using the @cuda.jit wrappers in masked_clip_2d_kernel. So if something goes wonky in _append_cuda but not in masked_clip_2d_kernel, blame me, not @stuartarchibald.
(If it turns out that this is fine, then this can be integrated with the version check as suggested by @philippjfr, so as to work with 0.49 as well.)

Also, if you're going to manually test, best to just test Python 3.7 for now. There are currently unrelated problems with Dask that make the test suite fail on 2.7 and 3.6.

[stuartarchibald]

xref numba/numba#6020 which is @gmarkall carrying on with some extra details on top of the original PR on the Numba side.

[AjayThorve]

@dHannasch I tested this with numba 0.51(https://github.com/numba/numba/releases/tag/0.51.0rc1), python 3.7.8.

The version check seems to work as well. Just a couple of fixes required

[philippjfr]

I'd like to go ahead and merge and then fix things up in a separate PR. Any objections?

[AjayThorve]

I'd like to go ahead and merge and then fix things up in a separate PR. Any objections?

@philippjfr that works