That appears to work, although there are 2-3 other places that need similar modifications in order to work. I'll try to get you a more complete list tomorrow.
I'm paying more attention tonight to warnings coming from Eigen than I had been previously ignoring, and both with my old Macros.h based hack and the newer modifications, I'm seeing some of these: "warning: calling a __host__ function from a __host__ __device__ function is not allowed". A representative one is:
external/eigen_archive/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h(735): warning: calling a __host__ function from a __host__ __device__ function is not allowed
detected during:
instantiation of "__nv_bool Eigen::TensorEvaluator<const Eigen::TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device>::evalSubExprsIfNeeded(MakePointer_<Eigen::TensorEvaluator<const Eigen::TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device>::CoeffReturnType>::Type) [with Op=Eigen::internal::AvgPoolMeanReducer<double>, Dims=const Eigen::IndexList<Eigen::type2index<1L>, Eigen::type2index<2L>>, ArgType=const Eigen::TensorImagePatchOp<-1L, -1L, const Eigen::TensorLayoutSwapOp<const Eigen::TensorMap<Eigen::Tensor<const double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>>, MakePointer_=Eigen::MakePointer, Device=Eigen::GpuDevice]"
external/eigen_archive/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h(172): here
instantiation of "__nv_bool Eigen::TensorEvaluator<const Eigen::TensorReshapingOp<NewDimensions, ArgType>, Device>::evalSubExprsIfNeeded(Eigen::TensorEvaluator<const Eigen::TensorReshapingOp<NewDimensions, ArgType>, Device>::CoeffReturnType *) [with NewDimensions=const Eigen::DSizes<Eigen::DenseIndex, 4>, ArgType=const Eigen::TensorReductionOp<Eigen::internal::AvgPoolMeanReducer<double>, const Eigen::IndexList<Eigen::type2index<1L>, Eigen::type2index<2L>>, const Eigen::TensorImagePatchOp<-1L, -1L, const Eigen::TensorLayoutSwapOp<const Eigen::TensorMap<Eigen::Tensor<const double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>>, Eigen::MakePointer>, Device=Eigen::GpuDevice]"
external/eigen_archive/unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h(146): here
instantiation of "__nv_bool Eigen::TensorEvaluator<const Eigen::TensorAssignOp<LeftArgType, RightArgType>, Device>::evalSubExprsIfNeeded(Eigen::TensorEvaluator<const Eigen::TensorAssignOp<LeftArgType, RightArgType>, Device>::Scalar *) [with LeftArgType=Eigen::TensorLayoutSwapOp<Eigen::TensorMap<Eigen::Tensor<double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>, RightArgType=const Eigen::TensorReshapingOp<const Eigen::DSizes<Eigen::DenseIndex, 4>, const Eigen::TensorReductionOp<Eigen::internal::AvgPoolMeanReducer<double>, const Eigen::IndexList<Eigen::type2index<1L>, Eigen::type2index<2L>>, const Eigen::TensorImagePatchOp<-1L, -1L, const Eigen::TensorLayoutSwapOp<const Eigen::TensorMap<Eigen::Tensor<const double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>>, Eigen::MakePointer>>, Device=Eigen::GpuDevice]"
external/eigen_archive/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor..h(422): here
instantiation of "void Eigen::internal::TensorExecutor<_expression_, Eigen::GpuDevice, Vectorizable, Tileable>::run(const _expression_ &, const Eigen::GpuDevice &) [with _expression_=const Eigen::TensorAssignOp<Eigen::TensorLayoutSwapOp<Eigen::TensorMap<Eigen::Tensor<double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>, const Eigen::TensorReshapingOp<const Eigen::DSizes<Eigen::DenseIndex, 4>, const Eigen::TensorReductionOp<Eigen::internal::AvgPoolMeanReducer<double>, const Eigen::IndexList<Eigen::type2index<1L>, Eigen::type2index<2L>>, const Eigen::TensorImagePatchOp<-1L, -1L, const Eigen::TensorLayoutSwapOp<const Eigen::TensorMap<Eigen::Tensor<const double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>>, Eigen::MakePointer>>>, Vectorizable=false, Tileable=false]"
external/eigen_archive/unsupported/Eigen/CXX11/src/Tensor/TensorDevice.h(35): here
instantiation of "Eigen::TensorDevice<ExpressionType, DeviceType> &Eigen::TensorDevice<ExpressionType, DeviceType>::operator=(const OtherDerived &) [with ExpressionType=Eigen::TensorLayoutSwapOp<Eigen::TensorMap<Eigen::Tensor<double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>, DeviceType=tensorflow::GPUDevice, OtherDerived=Eigen::TensorReshapingOp<const Eigen::DSizes<Eigen::DenseIndex, 4>, const Eigen::TensorReductionOp<Eigen::internal::AvgPoolMeanReducer<double>, const Eigen::IndexList<Eigen::type2index<1L>, Eigen::type2index<2L>>, const Eigen::TensorImagePatchOp<-1L, -1L, const Eigen::TensorLayoutSwapOp<const Eigen::TensorMap<Eigen::Tensor<const double, 4, 1, Eigen::DenseIndex>, 16, Eigen::MakePointer>>>, Eigen::MakePointer>>]"
./tensorflow/core/kernels/avgpooling_op.h(42): here
instantiation of "void tensorflow::functor::SpatialAvgPooling<Device, T>::operator()(const Device &, tensorflow::TTypes<T, 4, Eigen::DenseIndex>::Tensor, tensorflow::TTypes<T, 4, Eigen::DenseIndex>::ConstTensor, int, int, int, int, const Eigen::PaddingType &) [with Device=tensorflow::GPUDevice, T=double]"
I'm not sure how concerned I should be about these. The build will succeed, but.... I wouldn't be at all surprised to get weird results eventually.
In this particular case, it looks like it's complaining because Eigen::GpuDevice::allocate_temp appears to be __host__ rather than __host__ __device__ (i.e. missing EIGEN_DEVICE_FUNC). I fully admit that I could be misinterpreting that or otherwise misunderstanding something basic.