diff options
| author |  Deven Desai <deven.desai.amd@gmail.com> | 2018-07-11 10:39:54 -0400 |
|---|---|---|
| committer | Deven Desai <deven.desai.amd@gmail.com> | 2018-07-11 10:39:54 -0400 |
| commit | 876f392c396318f33454168db36ed54308e54e0d (patch) | |
| tree | a727bc91873b5c0aeec05312176a0f39e2cb64d5 /Eigen/Core | |
| parent | 1fe0b749042320501c59378f2860d9322b0c6e19 (diff) | |
Updates corresponding to the latest round of PR feedback
The major changes are
1. Moving CUDA/PacketMath.h to GPU/PacketMath.h
2. Moving CUDA/MathFunctions.h to GPU/MathFunction.h
3. Moving CUDA/CudaSpecialFunctions.h to GPU/GpuSpecialFunctions.h
The above three changes effectively enable the Eigen "Packet" layer for the HIP platform
4. Merging the "hip_basic" and "cuda_basic" unit tests into one ("gpu_basic")
5. Updating the "EIGEN_DEVICE_FUNC" marking in some places
The change has been tested on the HIP and CUDA platforms.
Diffstat (limited to 'Eigen/Core')
| -rw-r--r-- | Eigen/Core | 17 |
1 files changed, 12 insertions, 5 deletions
diff --git a/Eigen/Core b/Eigen/Core index 4336de91d..647a10831 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -81,6 +81,7 @@ // clang++ always considers constexpr functions as implicitly __host__ __device__ #define EIGEN_CONSTEXPR_ARE_DEVICE_FUNC #endif + #endif #elif defined(EIGEN_HIPCC) // Do not try to vectorize on HIP @@ -92,7 +93,7 @@ // We need hip_runtime.h to ensure that that EIGEN_USING_STD_MATH macro // works properly on the device side #include <hip/hip_runtime.h> - + #if defined(__HIP_DEVICE_COMPILE__) && !defined(EIGEN_NO_HIP) // analogous to EIGEN_CUDA_ARCH, but for HIP #define EIGEN_HIP_DEVICE_COMPILE __HIP_DEVICE_COMPILE__ @@ -356,7 +357,7 @@ #endif #if defined EIGEN_CUDACC - #define EIGEN_VECTORIZE_CUDA + #define EIGEN_VECTORIZE_GPU #include <vector_types.h> #if EIGEN_CUDACC_VER >= 70500 #define EIGEN_HAS_CUDA_FP16 @@ -369,14 +370,20 @@ #endif #if defined(EIGEN_HIPCC) && defined(EIGEN_HIP_DEVICE_COMPILE) + + #define EIGEN_VECTORIZE_GPU + #include <hip/hip_vector_types.h> + #define EIGEN_HAS_HIP_FP16 #include <hip/hip_fp16.h> + #define HIP_PATCH_WITH_NEW_FP16 18215 #if (HIP_VERSION_PATCH < HIP_PATCH_WITH_NEW_FP16) #define EIGEN_HAS_OLD_HIP_FP16 // Old HIP implementation does not have a explicit typedef for "half2" typedef __half2 half2; #endif + #endif #if defined(EIGEN_HAS_CUDA_FP16) || defined(EIGEN_HAS_HIP_FP16) @@ -550,9 +557,9 @@ using std::ptrdiff_t; #include "src/Core/arch/GPU/PacketMathHalf.h" #include "src/Core/arch/GPU/TypeCasting.h" -#if defined EIGEN_VECTORIZE_CUDA - #include "src/Core/arch/CUDA/PacketMath.h" - #include "src/Core/arch/CUDA/MathFunctions.h" +#if defined EIGEN_VECTORIZE_GPU + #include "src/Core/arch/GPU/PacketMath.h" + #include "src/Core/arch/GPU/MathFunctions.h" #endif #include "src/Core/arch/Default/Settings.h" |