diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h index b18200166..afa1a2697 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h @@ -193,16 +193,18 @@ static __global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self __syncthreads(); if (gridDim.x == 1 && first_index == 0) { - reducer.reduce(__low2half(*scratch), output); - reducer.reduce(__high2half(*scratch), output); + half tmp = __low2half(*scratch); + reducer.reduce(__high2half(*scratch), &tmp); + *output = tmp; } } template __global__ void ReductionCleanupKernelHalfFloat(Op& reducer, half* output, half2* scratch) { eigen_assert(threadIdx.x == 1); - reducer.reduce(__low2half(*scratch), output); - reducer.reduce(__high2half(*scratch), output); + half tmp = __low2half(*scratch); + reducer.reduce(__high2half(*scratch), &tmp); + *output = tmp; } #endif