mirror/eigen
2
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2024-12-15 07:10:37 +08:00
Code Issues Packages Projects Releases Wiki Activity

Only enable optimized reductions of fp16 if the reduction functor supports them

Browse Source
This commit is contained in:
Benoit Steiner 2016-05-31 10:33:40 -07:00
parent b6e306f189
commit 5aeb3687c4
1 changed files with 2 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h
View File

@ -291,7 +291,7 @@ struct FullReducer<Self, Op, GpuDevice, Vectorizable> {
#ifdef EIGEN_HAS_CUDA_FP16
static const bool HasOptimizedImplementation = !Op::IsStateful &&
(internal::is_same<typename Self::CoeffReturnType, float>::value ||
internal::is_same<typename Self::CoeffReturnType, Eigen::half>::value);
(internal::is_same<typename Self::CoeffReturnType, Eigen::half>::value && Op::PacketAccess));
#else
static const bool HasOptimizedImplementation = !Op::IsStateful &&
internal::is_same<typename Self::CoeffReturnType, float>::value;
@ -475,12 +475,6 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
template <typename Self, typename Op>
struct InnerReductionLauncher {
// Unfortunately nvidia doesn't support well exotic types such as complex,
// so reduce the scope of the optimized version of the code to the simple case
// of floats.
static const bool HasOptimizedImplementation = !Op::IsStateful &&
internal::is_same<typename Self::CoeffReturnType, float>::value;
template <typename OutputType>
static EIGEN_DEVICE_FUNC bool run(const Self&, Op&, const GpuDevice&, OutputType*, typename Self::Index, typename Self::Index) {
assert(false && "Should only be called to reduce floats and half floats on a gpu device");
@ -561,7 +555,7 @@ struct InnerReducer<Self, Op, GpuDevice> {
#ifdef EIGEN_HAS_CUDA_FP16
static const bool HasOptimizedImplementation = !Op::IsStateful &&
(internal::is_same<typename Self::CoeffReturnType, float>::value ||
internal::is_same<typename Self::CoeffReturnType, Eigen::half>::value);
(internal::is_same<typename Self::CoeffReturnType, Eigen::half>::value && Op::PacketAccess));
#else
static const bool HasOptimizedImplementation = !Op::IsStateful &&
internal::is_same<typename Self::CoeffReturnType, float>::value;