diff options
author | Deven Desai <deven.desai.amd@gmail.com> | 2018-11-19 18:13:59 +0000 |
---|---|---|
committer | Deven Desai <deven.desai.amd@gmail.com> | 2018-11-19 18:13:59 +0000 |
commit | e7e6809e6b38a5928efc0b5ca9520258e4d1fb3a (patch) | |
tree | 500d4209bdf3236178e5d9e9c7a6051d9c2d02d2 /Eigen/src/Core/util/Macros.h | |
parent | 6a510fe69c3d8ec0cdfa3e0f54a68c07ede68620 (diff) |
ROCm/HIP specfic fixes + updates
1. Eigen/src/Core/arch/GPU/Half.h
Updating the HIPCC implementation half so that it can declared as a __shared__ variable
2. Eigen/src/Core/util/Macros.h, Eigen/src/Core/util/Memory.h
introducing a EIGEN_USE_STD(func) macro that calls
- std::func be default
- ::func when eigen is being compiled with HIPCC
This change was requested in the previous HIP PR
(https://bitbucket.org/eigen/eigen/pull-requests/518/pr-with-hip-specific-fixes-for-the-eigen/diff)
3. unsupported/Eigen/CXX11/src/Tensor/TensorDeviceThreadPool.h
Removing EIGEN_DEVICE_FUNC attribute from pure virtual methods as it is not supported by HIPCC
4. unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
Disabling the template specializations of InnerMostDimReducer as they run into HIPCC link errors
Diffstat (limited to 'Eigen/src/Core/util/Macros.h')
-rw-r--r-- | Eigen/src/Core/util/Macros.h | 13 |
1 files changed, 13 insertions, 0 deletions
diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h index 9d277e26f..a7c6f50c3 100644 --- a/Eigen/src/Core/util/Macros.h +++ b/Eigen/src/Core/util/Macros.h @@ -896,6 +896,19 @@ namespace Eigen { #endif +// When compiling HIP device code with HIPCC, certain functions +// from the stdlib need to be pulled in from the global namespace +// (as opposed to from the std:: namespace). This is because HIPCC +// does not natively support all the std:: routines in device code. +// Instead it contains header files that declare the corresponding +// routines in the global namespace such they can be used in device code. +#if defined(EIGEN_HIP_DEVICE_COMPILE) + #define EIGEN_USING_STD(FUNC) using ::FUNC; +#else + #define EIGEN_USING_STD(FUNC) using std::FUNC; +#endif + + #if EIGEN_COMP_MSVC_STRICT && (EIGEN_COMP_MSVC < 1900 || EIGEN_CUDACC_VER>0) // for older MSVC versions, as well as 1900 && CUDA 8, using the base operator is sufficient (cf Bugs 1000, 1324) #define EIGEN_INHERIT_ASSIGNMENT_EQUAL_OPERATOR(Derived) \ |