aboutsummaryrefslogtreecommitdiffhomepage
path: root/Eigen
diff options
context:
space:
mode:
authorGravatar Antonio Sanchez <cantonios@google.com>2021-03-08 07:32:54 -0800
committerGravatar Antonio Sanchez <cantonios@google.com>2021-03-08 07:32:54 -0800
commit1296abdf82875d847c2a24afd09b301b38c0dd48 (patch)
tree66ab706a937bcf872050c391b1f046ff795b2976 /Eigen
parent60452431412b27517207b002094176a7dc06271c (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')
-rw-r--r--Eigen/src/Core/arch/Default/Half.h11
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) {}