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/arch/GPU | |
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/arch/GPU')
-rw-r--r-- | Eigen/src/Core/arch/GPU/Half.h | 8 |
1 files changed, 6 insertions, 2 deletions
diff --git a/Eigen/src/Core/arch/GPU/Half.h b/Eigen/src/Core/arch/GPU/Half.h index f87d8a18c..7873f8ec0 100644 --- a/Eigen/src/Core/arch/GPU/Half.h +++ b/Eigen/src/Core/arch/GPU/Half.h @@ -52,7 +52,9 @@ namespace half_impl { #if !defined(EIGEN_HAS_GPU_FP16) // Make our own __half_raw definition that is similar to CUDA's. struct __half_raw { - EIGEN_DEVICE_FUNC __half_raw() : x(0) {} + // The default constructor cannot initialize its member, otherwise the + // derived class Eigen::Half cannot be used as __shared__ variable in HIPCC. + EIGEN_DEVICE_FUNC __half_raw() {} explicit EIGEN_DEVICE_FUNC __half_raw(unsigned short raw) : x(raw) {} unsigned short x; }; @@ -70,7 +72,9 @@ struct __half_raw { // so we need to implicitly convert "__half_raw" to "__half" to avoid having to explicitly make // that conversiion in each call to a "__h*" routine...that is why we have "operator __half" routine struct __half_raw { - EIGEN_DEVICE_FUNC __half_raw() : x(0) {} + // The default constructor cannot initialize its member, otherwise the + // derived class Eigen::Half cannot be used as __shared__ variable in HIPCC. + EIGEN_DEVICE_FUNC __half_raw() {} explicit EIGEN_DEVICE_FUNC __half_raw(unsigned short raw) : x(raw) {} union { unsigned short x; |