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

__ldg is only available with cuda architectures >= 3.5

Browse Source
This commit is contained in:
Benoit Steiner 2016-05-11 15:22:13 -07:00
parent 02f76dae2d
commit 9091351dbe
1 changed files with 2 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
Eigen/src/Core/arch/CUDA/PacketMathHalf.h
View File

@ -72,7 +72,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<half>(half* to, co
template<> template<>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Aligned>(const half* from) { EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Aligned>(const half* from) {
#if __CUDA_ARCH__ >= 320 #if __CUDA_ARCH__ >= 350
return __ldg((const half2*)from); return __ldg((const half2*)from);
#else #else
return __halves2half2(*(from+0), *(from+1)); return __halves2half2(*(from+0), *(from+1));
@ -81,7 +81,7 @@ template<>
template<> template<>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Unaligned>(const half* from) { EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Unaligned>(const half* from) {
#if __CUDA_ARCH__ >= 320 #if __CUDA_ARCH__ >= 350
return __halves2half2(__ldg(from+0), __ldg(from+1)); return __halves2half2(__ldg(from+0), __ldg(from+1));
#else #else
return __halves2half2(*(from+0), *(from+1)); return __halves2half2(*(from+0), *(from+1));