From c6f7b338343ead9617558857c91fd3e03e347c3f Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Wed, 18 Jan 2017 10:45:28 +0000 Subject: [PATCH] Applying Benoit's comment. Embedding synchronisation inside device memcpy so there is no need to externally call synchronise() for device memcopy. --- .../Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 6 +++--- .../Eigen/CXX11/src/Tensor/TensorForcedEval.h | 6 +++--- .../Eigen/CXX11/src/Tensor/TensorMorphing.h | 14 ++++---------- .../Eigen/CXX11/src/Tensor/TensorStriding.h | 14 +++----------- 4 files changed, 13 insertions(+), 27 deletions(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index 268d9d148..8f8d1caad 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -234,7 +234,7 @@ struct SyclDevice { auto dst_acc =it2->second.template get_access(cgh); cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), MemCopyFunctor(src_acc, dst_acc, rng, i, offset)); }); - asynchronousExec(); + synchronize(); } /// The memcpyHostToDevice is used to copy the device only pointer to a host pointer. Using the device @@ -265,7 +265,7 @@ struct SyclDevice { auto dst_acc =dest_buf.template get_access(cgh); cgh.parallel_for( cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), MemCopyFunctor(src_acc, dst_acc, rng, 0, offset)); }); - asynchronousExec(); + synchronize(); } /// returning the sycl queue EIGEN_STRONG_INLINE cl::sycl::queue& sycl_queue() const { return m_queue_stream->m_queue;} @@ -308,7 +308,7 @@ struct SyclDevice { } EIGEN_STRONG_INLINE void asynchronousExec() const { - ///FIXEDME:: currently there is a race condition regarding the asynch scheduler. + ///FIXEDME:: currently there is a race condition regarding the asynch scheduler. //sycl_queue().throw_asynchronous();// does not pass. Temporarily disabled sycl_queue().wait_and_throw(); //pass diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h b/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h index 822e22c2d..abe85c860 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h @@ -143,12 +143,12 @@ struct TensorEvaluator, Device> return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize); } - CoeffReturnType* data() const { return m_buffer; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { return m_buffer; } /// required by sycl in order to extract the sycl accessor - const TensorEvaluator& impl() { return m_impl; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator& impl() { return m_impl; } /// used by sycl in order to build the sycl buffer - const Device& device() const{return m_device;} + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const{return m_device;} private: TensorEvaluator m_impl; const ArgType m_op; diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h index dbe11c7af..6ddd2ca18 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h @@ -736,22 +736,12 @@ struct TensorEvaluator::value; ++i) { eigen_assert(m_strides[i] != 0 && "0 stride is invalid"); if(m_strides[i]>0){ - #ifndef __SYCL_DEVICE_ONLY__ startIndicesClamped[i] = clamp(op.startIndices()[i], 0, m_impl.dimensions()[i]); stopIndicesClamped[i] = clamp(op.stopIndices()[i], 0, m_impl.dimensions()[i]); - #else - startIndicesClamped[i] = cl::sycl::clamp(static_cast(op.startIndices()[i]), static_cast(0), static_cast(m_impl.dimensions()[i])); - stopIndicesClamped[i] = cl::sycl::clamp(static_cast(op.stopIndices()[i]), static_cast(0), static_cast(m_impl.dimensions()[i])); - #endif }else{ /* implies m_strides[i]<0 by assert */ - #ifndef __SYCL_DEVICE_ONLY__ startIndicesClamped[i] = clamp(op.startIndices()[i], -1, m_impl.dimensions()[i] - 1); stopIndicesClamped[i] = clamp(op.stopIndices()[i], -1, m_impl.dimensions()[i] - 1); - #else - startIndicesClamped[i] = cl::sycl::clamp(static_cast(op.startIndices()[i]), static_cast(-1), static_cast(m_impl.dimensions()[i] - 1)); - stopIndicesClamped[i] = cl::sycl::clamp(static_cast(op.stopIndices()[i]), static_cast(-1), static_cast(m_impl.dimensions()[i] - 1)); - #endif } m_startIndices[i] = startIndicesClamped[i]; } @@ -867,7 +857,11 @@ struct TensorEvaluator m_outputStrides; diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorStriding.h b/unsupported/Eigen/CXX11/src/Tensor/TensorStriding.h index 93615e5c2..2237140e7 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorStriding.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorStriding.h @@ -121,11 +121,7 @@ struct TensorEvaluator, Device> { m_dimensions = m_impl.dimensions(); for (int i = 0; i < NumDims; ++i) { -#ifndef __SYCL_DEVICE_ONLY__ - m_dimensions[i] = ceilf(static_cast(m_dimensions[i]) / op.strides()[i]); -#else - m_dimensions[i] = cl::sycl::ceil(static_cast(m_dimensions[i]) / op.strides()[i]); -#endif + m_dimensions[i] =Eigen::numext::ceil(static_cast(m_dimensions[i]) / op.strides()[i]); } const typename TensorEvaluator::Dimensions& input_dims = m_impl.dimensions(); @@ -233,8 +229,6 @@ struct TensorEvaluator, Device> /// required by sycl in order to extract the accessor Strides functor() const { return m_strides; } - - protected: EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const { @@ -264,7 +258,6 @@ struct TensorEvaluator, Device> const Strides m_strides; }; - // Eval as lvalue template struct TensorEvaluator, Device> @@ -299,10 +292,9 @@ struct TensorEvaluator, Device> } /// required by sycl in order to extract the accessor - const TensorEvaluator& impl() const { return this->m_impl; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator& impl() const { return this->m_impl; } /// required by sycl in order to extract the accessor - Strides functor() const { return this->m_strides; } - + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Strides functor() const { return this->m_strides; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void writePacket(Index index, const PacketReturnType& x)