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

Small improvement to the full reduction of fp16

Browse Source
This commit is contained in:
Benoit Steiner 2016-05-10 11:58:18 -07:00
parent 0b9e3dcd06
commit 0eb69b7552
1 changed files with 6 additions and 4 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

@ -193,16 +193,18 @@ static __global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self
__syncthreads(); __syncthreads();
if (gridDim.x == 1 && first_index == 0) { if (gridDim.x == 1 && first_index == 0) {
reducer.reduce(__low2half(*scratch), output); half tmp = __low2half(*scratch);
reducer.reduce(__high2half(*scratch), output); reducer.reduce(__high2half(*scratch), &tmp);
*output = tmp;
} }
} }
template <typename Op> template <typename Op>
__global__ void ReductionCleanupKernelHalfFloat(Op& reducer, half* output, half2* scratch) { __global__ void ReductionCleanupKernelHalfFloat(Op& reducer, half* output, half2* scratch) {
eigen_assert(threadIdx.x == 1); eigen_assert(threadIdx.x == 1);
reducer.reduce(__low2half(*scratch), output); half tmp = __low2half(*scratch);
reducer.reduce(__high2half(*scratch), output); reducer.reduce(__high2half(*scratch), &tmp);
*output = tmp;
} }
#endif #endif