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

Added __ldg primitive for fp16.

Browse Source
This commit is contained in:
Benoit Steiner 2016-03-31 10:55:03 -07:00
parent b575fb1d02
commit c36ab19902
1 changed files with 7 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
Eigen/src/Core/arch/CUDA/Half.h
View File

@ -423,7 +423,14 @@ using ::ceil;
__device__ inline Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) { __device__ inline Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) {
return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width)); return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width));
} }
#endif
// ldg() has an overload for __half, but we also need one for Eigen::half.
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 320
static inline EIGEN_DEVICE_FUNC Eigen::half __ldg(const Eigen::half* ptr) {
return Eigen::internal::raw_uint16_to_half(
__ldg(reinterpret_cast<const unsigned short*>(ptr)));
}
#endif #endif