diff options
author | Antonio Sanchez <cantonios@google.com> | 2021-03-08 07:32:54 -0800 |
---|---|---|
committer | Antonio Sanchez <cantonios@google.com> | 2021-03-08 07:32:54 -0800 |
commit | 1296abdf82875d847c2a24afd09b301b38c0dd48 (patch) | |
tree | 66ab706a937bcf872050c391b1f046ff795b2976 /Eigen/src/Core/arch/Default | |
parent | 60452431412b27517207b002094176a7dc06271c (diff) |
Fix non-trivial Half constructor for CUDA.
Both CUDA and HIP require trivial default constructors for types used
in shared memory. Otherwise failing with
```
error: initialization is not supported for __shared__ variables.
```
Diffstat (limited to 'Eigen/src/Core/arch/Default')
-rw-r--r-- | Eigen/src/Core/arch/Default/Half.h | 11 |
1 files changed, 6 insertions, 5 deletions
diff --git a/Eigen/src/Core/arch/Default/Half.h b/Eigen/src/Core/arch/Default/Half.h index 12ed8f676..5fae1eb5a 100644 --- a/Eigen/src/Core/arch/Default/Half.h +++ b/Eigen/src/Core/arch/Default/Half.h @@ -85,11 +85,12 @@ namespace half_impl { #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE) // Make our own __half_raw definition that is similar to CUDA's. struct __half_raw { -#if (defined(EIGEN_HAS_HIP_FP16) && !defined(EIGEN_HIP_DEVICE_COMPILE)) - // Eigen::half can be used as the datatype for some shared memory declarations (in Eigen and TF) - // (In HIP) The element type for shared memory declaration cannot have non-trivial constructors - // and hence the following special casing (which skips the zero-initilization) - // Note that this check gets done even in the host compilation phase, and hence the need for this +#if (defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_DEVICE_COMPILE_PHASE)) + // Eigen::half can be used as the datatype for shared memory declarations (in Eigen and TF) + // The element type for shared memory cannot have non-trivial constructors + // and hence the following special casing (which skips the zero-initilization). + // Note that this check gets done even in the host compilation phase, and + // hence the need for this EIGEN_DEVICE_FUNC __half_raw() {} #else EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw() : x(0) {} |