Half-Precision Intrinsics

[iyaja]

This adds support for some half-precision intrinsics as mentioned in #391.

The implementations so far:

• Alternative definitions for some functions like expm1 which did not have a fallback implementation for Float16 arguments in Base. Some Base math functions also rely on libm, which does not work for the CUDA definitions.
• shfl_recurse for Float16 via reinterpret, as in done for Float64.
• atomic_add! for Float16. PTX 6.3, which is the current supported version, only has an f16 atomic add instruction. Other atomics can be implemented in via atomic_cas!, but it seems like like CUDA headers (i.e. cuda_fp16.h) only provides a definition for __half atomicAdd(__half *address, __half val) and none of the other atomics.

The device/intrinsics test suite has also been updated to include the above definitions.

[codecov]

Codecov Report

Merging #871 (92a661f) into master (57577a1) will increase coverage by 0.06%.
The diff coverage is n/a.

@@            Coverage Diff             @@
##           master     #871      +/-   ##
==========================================
+ Coverage   77.06%   77.13%   +0.06%     
==========================================
  Files         121      121              
  Lines        7530     7517      -13     
==========================================
- Hits         5803     5798       -5     
+ Misses       1727     1719       -8     

Impacted Files	Coverage Δ
lib/cudadrv/memory.jl	85.87% <0.00%> (+0.48%)	⬆️
src/pointer.jl	77.94% <0.00%> (+5.57%)	⬆️

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 3b871a0...92a661f. Read the comment docs.

[maleadt]

All green :-) Let's merge this.