From 876f392c396318f33454168db36ed54308e54e0d Mon Sep 17 00:00:00 2001
From: Deven Desai <deven.desai.amd@gmail.com>
Date: Wed, 11 Jul 2018 10:39:54 -0400
Subject: 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.
---
 Eigen/Core | 17 ++++++++++++-----
 1 file changed, 12 insertions(+), 5 deletions(-)

(limited to 'Eigen/Core')

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"
-- 
cgit v1.2.3