aboutsummaryrefslogtreecommitdiffhomepage
path: root/Eigen/src/Core/arch/GPU
diff options
context:
space:
mode:
authorGravatar Deven Desai <deven.desai.amd@gmail.com>2018-11-19 18:13:59 +0000
committerGravatar Deven Desai <deven.desai.amd@gmail.com>2018-11-19 18:13:59 +0000
commite7e6809e6b38a5928efc0b5ca9520258e4d1fb3a (patch)
tree500d4209bdf3236178e5d9e9c7a6051d9c2d02d2 /Eigen/src/Core/arch/GPU
parent6a510fe69c3d8ec0cdfa3e0f54a68c07ede68620 (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.h8
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;