From 61fc78bbda22f425a97761be57c219928d929ddc Mon Sep 17 00:00:00 2001
From: Rasmus Munk Larsen <rmlarsen@google.com>
Date: Tue, 13 Oct 2020 23:53:11 +0000
Subject: Get rid of nested template specialization in TensorReductionGpu.h,
 which was broken by c6953f799b01d36f4236b64f351cc1446e0abe17.

---
 unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

(limited to 'unsupported/Eigen/CXX11/src')

diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
index 02a514c0f..db4e8d866 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
@@ -80,8 +80,8 @@ __device__ inline double atomicExchCustom(double* address, double val) {
 }
 
 #ifdef EIGEN_HAS_GPU_FP16
-template <template <typename T> class R>
-__device__ inline void atomicReduce(half2* output, half2 accum, R<half>& reducer) {
+template <typename R>
+__device__ inline void atomicReduce(half2* output, half2 accum, R& reducer) {
   unsigned int oldval = *reinterpret_cast<unsigned int*>(output);
   unsigned int newval = oldval;
   reducer.reducePacket(accum, reinterpret_cast<half2*>(&newval));
@@ -99,9 +99,8 @@ __device__ inline void atomicReduce(half2* output, half2 accum, R<half>& reducer
   }
 }
 // reduction should be associative since reduction is not atomic in wide vector but atomic in half2 operations
-template <template <typename T> class R>
-__device__ inline void atomicReduce(Packet4h2* output, Packet4h2 accum,
-                                    R<half>& reducer) {
+template <typename R>
+__device__ inline void atomicReduce(Packet4h2* output, Packet4h2 accum, R& reducer) {
   half2* houtput=reinterpret_cast<half2*>(output);
   half2* haccum=reinterpret_cast<half2*>(&accum);
   for(int i=0;i<4;++i){
-- 
cgit v1.2.3