From 1296abdf82875d847c2a24afd09b301b38c0dd48 Mon Sep 17 00:00:00 2001 From: Antonio Sanchez Date: Mon, 8 Mar 2021 07:32:54 -0800 Subject: 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. ``` --- Eigen/src/Core/arch/Default/Half.h | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) (limited to 'Eigen') 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) {} -- cgit v1.2.3