GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-04-21 17:19:36 +08:00
Code Issues Packages Projects Releases Wiki Activity

bug #1526 - CUDA compilation fails on CUDA 9.x SDK when arch is set to compute_60 and/or above

Browse Source
This commit is contained in:
Jiandong Ruan 2018-09-08 12:05:33 -07:00
parent edfb7962fd
commit 6dcd2642aa
1 changed files with 8 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
Eigen/src/Core/arch/GPU/Half.h
View File

@ -272,7 +272,11 @@ namespace half_impl {
// conversion steps back and forth. // conversion steps back and forth.
EIGEN_STRONG_INLINE __device__ half operator + (const half& a, const half& b) { EIGEN_STRONG_INLINE __device__ half operator + (const half& a, const half& b) {
#if defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER >= 90000
return __hadd(::__half(a), ::__half(b));
#else
return __hadd(a, b); return __hadd(a, b);
#endif
} }
EIGEN_STRONG_INLINE __device__ half operator * (const half& a, const half& b) { EIGEN_STRONG_INLINE __device__ half operator * (const half& a, const half& b) {
return __hmul(a, b); return __hmul(a, b);
@ -281,9 +285,13 @@ EIGEN_STRONG_INLINE __device__ half operator - (const half& a, const half& b) {
return __hsub(a, b); return __hsub(a, b);
} }
EIGEN_STRONG_INLINE __device__ half operator / (const half& a, const half& b) { EIGEN_STRONG_INLINE __device__ half operator / (const half& a, const half& b) {
#if defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER >= 90000
return __hdiv(a, b);
#else
float num = __half2float(a); float num = __half2float(a);
float denom = __half2float(b); float denom = __half2float(b);
return __float2half(num / denom); return __float2half(num / denom);
#endif
} }
EIGEN_STRONG_INLINE __device__ half operator - (const half& a) { EIGEN_STRONG_INLINE __device__ half operator - (const half& a) {
return __hneg(a); return __hneg(a);