GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-09-23 14:53:13 +08:00
Code Issues Packages Projects Releases Wiki Activity

Fix non-trivial Half constructor for CUDA.

Browse Source 
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.
```
This commit is contained in:
Antonio Sanchez 2021-03-08 07:32:54 -08:00
parent 6045243141
commit 1296abdf82
1 changed files with 6 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
Eigen/src/Core/arch/Default/Half.h
View File

@ -85,11 +85,12 @@ namespace half_impl {
#if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE) #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE)
// Make our own __half_raw definition that is similar to CUDA's. // Make our own __half_raw definition that is similar to CUDA's.
struct __half_raw { struct __half_raw {
#if (defined(EIGEN_HAS_HIP_FP16) && !defined(EIGEN_HIP_DEVICE_COMPILE)) #if (defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_DEVICE_COMPILE_PHASE))
// Eigen::half can be used as the datatype for some shared memory declarations (in Eigen and TF) // Eigen::half can be used as the datatype for shared memory declarations (in Eigen and TF)
// (In HIP) The element type for shared memory declaration cannot have non-trivial constructors // The element type for shared memory cannot have non-trivial constructors
// and hence the following special casing (which skips the zero-initilization) // 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 // Note that this check gets done even in the host compilation phase, and
// hence the need for this
EIGEN_DEVICE_FUNC __half_raw() {} EIGEN_DEVICE_FUNC __half_raw() {}
#else #else
EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw() : x(0) {} EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw() : x(0) {}