Commit Graph

1236 Commits

Author SHA1 Message Date
Gael Guennebaud
f2dc048df9 complete implementation of Packet16h (AVX512) 2018-07-06 17:43:11 +02:00
Gael Guennebaud
f4d623ffa7 Complete Packet8h implementation and test it in packetmath unit test 2018-07-06 17:13:36 +02:00
Deven Desai
b6cc0961b1 updates based on PR feedback
There are two major changes (and a few minor ones which are not listed here...see PR discussion for details)

1. Eigen::half implementations for HIP and CUDA have been merged.
This means that
- `CUDA/Half.h` and `HIP/hcc/Half.h` got merged to a new file `GPU/Half.h`
- `CUDA/PacketMathHalf.h` and `HIP/hcc/PacketMathHalf.h` got merged to a new file `GPU/PacketMathHalf.h`
- `CUDA/TypeCasting.h` and `HIP/hcc/TypeCasting.h` got merged to a new file `GPU/TypeCasting.h`

After this change the `HIP/hcc` directory only contains one file `math_constants.h`. That will go away too once that file becomes a part of the HIP install.

2. new macros EIGEN_GPUCC, EIGEN_GPU_COMPILE_PHASE and EIGEN_HAS_GPU_FP16 have been added and the code has been updated to use them where appropriate.
- `EIGEN_GPUCC` is the same as `(EIGEN_CUDACC || EIGEN_HIPCC)`
- `EIGEN_GPU_DEVICE_COMPILE` is the same as `(EIGEN_CUDA_ARCH || EIGEN_HIP_DEVICE_COMPILE)`
- `EIGEN_HAS_GPU_FP16` is the same as `(EIGEN_HAS_CUDA_FP16 or EIGEN_HAS_HIP_FP16)`
2018-06-14 10:21:54 -04:00
Deven Desai
ba972fb6b4 moving Half headers from CUDA dir to GPU dir, removing the HIP versions 2018-06-13 12:26:18 -04:00
Deven Desai
d1d22ef0f4 syncing this fork with upstream 2018-06-13 12:09:52 -04:00
Rasmus Munk Larsen
bda71ad394 Fix typo in pbend for AltiVec. 2018-06-22 15:04:35 -07:00
Michael Figurnov
30fa3d0454 Merge from eigen/eigen 2018-06-07 17:57:56 +01:00
Michael Figurnov
4bd158fa37 Derivative of the incomplete Gamma function and the sample of a Gamma random variable.
In addition to igamma(a, x), this code implements:
* igamma_der_a(a, x) = d igamma(a, x) / da -- derivative of igamma with respect to the parameter
* gamma_sample_der_alpha(alpha, sample) -- reparameterization derivative of a Gamma(alpha, 1) random variable sample with respect to the alpha parameter

The derivatives are computed by forward mode differentiation of the igamma(a, x) code. Although gamma_sample_der_alpha can be implemented via igamma_der_a, a separate function is more accurate and efficient due to analytical cancellation of some terms. All three functions are implemented by a method parameterized with "mode" that always computes the derivatives, but does not return them unless required by the mode. The compiler is expected to (and, based on benchmarks, does) skip the unnecessary computations depending on the mode.
2018-06-06 18:49:26 +01:00
Deven Desai
8fbd47052b Adding support for using Eigen in HIP kernels.
This commit enables the use of Eigen on HIP kernels / AMD GPUs. Support has been added along the same lines as what already exists for using Eigen in CUDA kernels / NVidia GPUs.

Application code needs to explicitly define EIGEN_USE_HIP when using Eigen in HIP kernels. This is because some of the CUDA headers get picked up by default during Eigen compile (irrespective of whether or not the underlying compiler is CUDACC/NVCC, for e.g. Eigen/src/Core/arch/CUDA/Half.h). In order to maintain this behavior, the EIGEN_USE_HIP macro is used to switch to using the HIP version of those header files (see Eigen/Core and unsupported/Eigen/CXX11/Tensor)


Use the "-DEIGEN_TEST_HIP" cmake option to enable the HIP specific unit tests.
2018-06-06 10:12:58 -04:00
Michael Figurnov
f216854453 Exponentially scaled modified Bessel functions of order zero and one.
The functions are conventionally called i0e and i1e. The exponentially scaled version is more numerically stable. The standard Bessel functions can be obtained as i0(x) = exp(|x|) i0e(x)

The code is ported from Cephes and tested against SciPy.
2018-05-31 15:34:53 +01:00
Gael Guennebaud
647b724a36 Define pcast<> for SSE types even when AVX is enabled. (otherwise float are silently reinterpreted as int instead of being converted) 2018-05-29 20:46:46 +02:00
Gael Guennebaud
49262dfee6 Fix compilation and SSE support with PGI compiler 2018-05-29 15:09:31 +02:00
Gael Guennebaud
7134fa7a2e Fix compilation with MSVC by reverting to char* for _mm_prefetch except for PGI (the later being the one that has the wrong prototype). 2018-06-07 09:33:10 +02:00
Gael Guennebaud
6e7118265d Fix compilation with NEON+MSVC 2018-04-26 10:50:41 +02:00
Gael Guennebaud
e8ca5166a9 bug #1428: atempt to make NEON vectorization compilable by MSVC.
The workaround is to wrap NEON packet types to make them different c++ types.
2018-04-24 11:19:49 +02:00
Jayaram Bobba
b7b868d1c4 fix AVX512 plog 2018-04-20 13:39:18 -07:00
Gael Guennebaud
686fb57233 fix const cast in NEON 2018-04-18 18:46:34 +02:00
Gael Guennebaud
e43ca0320d bug #1520: workaround some -Wfloat-equal warnings by calling std::equal_to 2018-04-11 15:24:13 +02:00
Gael Guennebaud
13f5df9f67 Add a note on vec_min vs asm 2018-04-04 13:10:38 +02:00
Gael Guennebaud
e91e314347 bug #1494: makes pmin/pmax behave on Altivec/VSX as on x86 regading NaNs 2018-04-04 11:39:19 +02:00
Gael Guennebaud
112c899304 comment unreachable code 2018-04-03 23:16:43 +02:00
Gael Guennebaud
40b4bf3d32 AVX512: _mm512_rsqrt28_ps is available for AVX512ER only 2018-04-03 14:36:27 +02:00
Gael Guennebaud
584951ca4d Rename predux_downto4 to be more accurate on its semantic. 2018-04-03 14:28:38 +02:00
Gael Guennebaud
7b0630315f AVX512: fix psqrt and prsqrt 2018-04-03 14:12:50 +02:00
Gael Guennebaud
6719409cd9 AVX512: add missing pinsertfirst and pinsertlast, implement pblend for Packet8d, fix compilation without AVX512DQ 2018-04-03 14:11:56 +02:00
luz.paz
e3912f5e63 MIsc. source and comment typos
Found using `codespell` and `grep` from downstream FreeCAD
2018-03-11 10:01:44 -04:00
Daniel Trebbien
0c57be407d Move up the specialization of std::numeric_limits
This fixes a compilation error seen when building TensorFlow on macOS:
https://github.com/tensorflow/tensorflow/issues/17067
2018-02-18 15:35:45 -08:00
nluehr
aefd5fd5c4 Replace __float2half_rn with __float2half
The latter provides a consistent definition for CUDA 8.0 and 9.0.
2017-11-28 10:15:46 -08:00
nluehr
dd6de618c3 Fix incorrect integer cast in predux<half2>().
Bug corrupts results on Maxwell and earlier GPU architectures.
2017-11-21 10:47:00 -08:00
Christoph Hertzberg
11ddac57e5 Merged in guillaume_michel/eigen (pull request PR-334)
- Add support for NEON plog PacketMath function
2017-10-23 13:22:22 +00:00
Henry Schreiner
9bb26eb8f1 Restore __device__ 2017-10-21 00:50:38 +00:00
Henry Schreiner
4245475d22 Fixing missing inlines on device functions for newer CUDA cards 2017-10-20 03:20:13 +00:00
Konstantinos Margaritis
6c3475f110 remove debugging 2017-10-12 15:34:55 -04:00
Konstantinos Margaritis
df7644aec3 Merged eigen/eigen into default 2017-10-12 22:23:13 +03:00
Konstantinos Margaritis
c4ad358565 explicitly set conjugate mask 2017-10-11 11:05:29 -04:00
Konstantinos Margaritis
380d41fd76 added some extra debugging 2017-10-11 10:40:12 -04:00
Konstantinos Margaritis
d0b7b9d0d3 some Packet2cf pmul fixes 2017-10-11 10:17:22 -04:00
Konstantinos Margaritis
df173f5620 initial pexp() for 32-bit floats, commented out due to vec_cts() 2017-10-11 09:40:49 -04:00
Konstantinos Margaritis
3dcae2a27f initial pexp() for 32-bit floats, commented out due to vec_cts() 2017-10-11 09:40:45 -04:00
Konstantinos Margaritis
c2a2246489 fix predux_mul for z14/float 2017-10-10 13:38:32 -04:00
Konstantinos Margaritis
bc30305d29 complete z14 port 2017-10-09 16:55:10 -04:00
Gael Guennebaud
9c353dd145 Add C++11 max_digits10 for half. 2017-09-06 10:22:47 +02:00
Benoit Steiner
a4089991eb Added support for CUDA 9.0. 2017-08-31 02:49:39 +00:00
Konstantinos Margaritis
1affe3d8df Merged eigen/eigen into default 2017-08-24 12:24:01 +03:00
Gael Guennebaud
21633e585b bug #1462: remove all occurences of the deprecated __CUDACC_VER__ macro by introducing EIGEN_CUDACC_VER 2017-08-24 11:06:47 +02:00
Konstantinos Margaritis
e1e71ca4e4 initial support for z14 2017-08-06 19:53:18 -04:00
Gael Guennebaud
24fe1de9b4 merge 2017-06-15 10:17:39 +02:00
Gael Guennebaud
b240080e64 bug #1436: fix compilation of Jacobi rotations with ARM NEON, some specializations of internal::conj_helper were missing. 2017-06-15 10:16:30 +02:00
Benoit Steiner
3baef62b9a Added missing __device__ qualifier 2017-06-13 12:56:55 -07:00
Benoit Steiner
449936828c Added missing __device__ qualifier 2017-06-13 12:54:57 -07:00