diff options
11 files changed, 251 insertions, 27 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index b2ddea2ba..7f0f16de3 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -123,9 +123,45 @@ struct SyclDevice { // some runtime conditions that can be applied here EIGEN_STRONG_INLINE bool isDeviceSuitable() const { return true; } + template <typename T> EIGEN_STRONG_INLINE std::map<const void *, std::shared_ptr<void>>::iterator find_nearest(const T* ptr) const { + auto it1 = buffer_map.find(ptr); + if (it1 != buffer_map.end()){ + return it1; + } + else{ + for(std::map<const void *, std::shared_ptr<void>>::iterator it=buffer_map.begin(); it!=buffer_map.end(); ++it){ + auto size = ((cl::sycl::buffer<T, 1>*)it->second.get())->get_size(); + if((static_cast<const T*>(it->first) < ptr) && (ptr < (static_cast<const T*>(it->first)) + size)) return it; + } + } + return buffer_map.end(); + } + /// the memcpy function - EIGEN_STRONG_INLINE void memcpy(void *dst, const void *src, size_t n) const { - ::memcpy(dst, src, n); + template<typename T> EIGEN_STRONG_INLINE void memcpy(void *dst, const T *src, size_t n) const { + auto it1 = find_nearest(src); + auto it2 = find_nearest(static_cast<T*>(dst)); + if ((it1 != buffer_map.end()) && (it2!=buffer_map.end())) { + auto offset= (src - (static_cast<const T*>(it1->first))); + auto i= ((static_cast<T*>(dst)) - const_cast<T*>((static_cast<const T*>(it2->first)))); + size_t rng, GRange, tileSize; + parallel_for_setup(n/sizeof(T), tileSize, rng, GRange); + m_queue.submit([&](cl::sycl::handler &cgh) { + auto src_acc =((cl::sycl::buffer<T, 1>*)it1->second.get())-> template get_access<cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer>(cgh); + auto dst_acc =((cl::sycl::buffer<T, 1>*)it2->second.get())-> template get_access<cl::sycl::access::mode::discard_write, cl::sycl::access::target::global_buffer>(cgh); + typedef decltype(src_acc) DevToDev; + cgh.parallel_for<DevToDev>( cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), [=](cl::sycl::nd_item<1> itemID) { + auto globalid=itemID.get_global_linear_id(); + if (globalid< rng) { + dst_acc[globalid+i ]=src_acc[globalid+offset]; + } + }); + }); + m_queue.throw_asynchronous(); + } else{ + eigen_assert("no source or destination device memory found."); + } + //::memcpy(dst, src, n); } /// The memcpyHostToDevice is used to copy the device only pointer to a host pointer. Using the device @@ -136,7 +172,7 @@ struct SyclDevice { template<typename T> EIGEN_STRONG_INLINE void memcpyHostToDevice(T *dst, const T *src, size_t n) const { auto host_acc= get_sycl_buffer(n, dst)-> template get_access<cl::sycl::access::mode::discard_write, cl::sycl::access::target::host_buffer>(); - memcpy(host_acc.get_pointer(), src, n); + ::memcpy(host_acc.get_pointer(), src, n); } /// The memcpyDeviceToHost is used to copy the data from host to device. Here, in order to avoid double copying the data. We create a sycl /// buffer with map_allocator for the destination pointer with a discard_write accessor on it. The lifespan of the buffer is bound to the @@ -145,21 +181,22 @@ struct SyclDevice { /// would be available on the dst pointer using fast copy technique (map_allocator). In this case we can make sure that we copy the data back /// to the cpu only once per function call. template<typename T> EIGEN_STRONG_INLINE void memcpyDeviceToHost(T *dst, const T *src, size_t n) const { - auto it = buffer_map.find(src); + auto it = find_nearest(src); + auto offset = src- (static_cast<const T*>(it->first)); if (it != buffer_map.end()) { size_t rng, GRange, tileSize; parallel_for_setup(n/sizeof(T), tileSize, rng, GRange); - + // Assuming that the dst is the start of the destination pointer auto dest_buf = cl::sycl::buffer<T, 1, cl::sycl::map_allocator<T>>(dst, cl::sycl::range<1>(rng)); typedef decltype(dest_buf) SYCLDTOH; m_queue.submit([&](cl::sycl::handler &cgh) { auto src_acc= (static_cast<cl::sycl::buffer<T, 1>*>(it->second.get()))-> template get_access<cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer>(cgh); auto dst_acc =dest_buf.template get_access<cl::sycl::access::mode::discard_write, cl::sycl::access::target::global_buffer>(cgh); cgh.parallel_for<SYCLDTOH>( cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), [=](cl::sycl::nd_item<1> itemID) { - auto globalid=itemID.get_global_linear_id(); - if (globalid< dst_acc.get_size()) { - dst_acc[globalid] = src_acc[globalid]; - } + auto globalid=itemID.get_global_linear_id(); + if (globalid< dst_acc.get_size()) { + dst_acc[globalid] = src_acc[globalid + offset]; + } }); }); m_queue.throw_asynchronous(); @@ -176,12 +213,12 @@ struct SyclDevice { m_queue.submit([&](cl::sycl::handler &cgh) { auto buf_acc =get_sycl_buffer(n, buff)-> template get_access<cl::sycl::access::mode::discard_write, cl::sycl::access::target::global_buffer>(cgh); cgh.parallel_for<SyclDevice>( cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), [=](cl::sycl::nd_item<1> itemID) { - auto globalid=itemID.get_global_linear_id(); - auto buf_ptr= reinterpret_cast<typename cl::sycl::global_ptr<unsigned char>::pointer_t>((&(*buf_acc.get_pointer()))); - if (globalid< buf_acc.get_size()) { - for(size_t i=0; i<sizeof(T); i++) - buf_ptr[globalid*sizeof(T) + i] = c; - } + auto globalid=itemID.get_global_linear_id(); + auto buf_ptr= reinterpret_cast<typename cl::sycl::global_ptr<unsigned char>::pointer_t>((&(*buf_acc.get_pointer()))); + if (globalid< buf_acc.get_size()) { + for(size_t i=0; i<sizeof(T); i++) + buf_ptr[globalid*sizeof(T) + i] = c; + } }); }); m_queue.throw_asynchronous(); diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h index d34f1e328..2b56340bd 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h @@ -299,6 +299,16 @@ template <typename Index> struct MemcpyTriggerForSlicing<Index, GpuDevice> { EIGEN_DEVICE_FUNC bool operator ()(Index val) const { return val > 4*1024*1024; } }; #endif + +// It is very expensive to start the memcpy kernel on GPU: we therefore only +// use it for large copies. +#ifdef EIGEN_USE_SYCL +template <typename Index> struct MemcpyTriggerForSlicing<Index, const Eigen::SyclDevice> { + EIGEN_DEVICE_FUNC MemcpyTriggerForSlicing(const SyclDevice&) { } + EIGEN_DEVICE_FUNC bool operator ()(Index val) const { return val > 4*1024*1024; } +}; +#endif + } // Eval as rvalue @@ -493,7 +503,14 @@ struct TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, ArgType>, Devi } return NULL; } - + /// used by stcl + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const{ + return m_impl; + } + /// used by stcl + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const StartIndices& startIndices() const{ + return m_offsets; + } protected: EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const { diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h index 8729c86ee..bb847afad 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h @@ -48,9 +48,9 @@ struct DeviceConvertor{ /// specialisation of the \ref ConvertToDeviceExpression struct when the node /// type is TensorMap #define TENSORMAPCONVERT(CVQual)\ -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_>\ -struct ConvertToDeviceExpression<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_> > {\ - typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\ +template <typename T, int Options2_, template <class> class MakePointer_>\ +struct ConvertToDeviceExpression<CVQual TensorMap<T, Options2_, MakePointer_> > {\ + typedef CVQual TensorMap<T, Options2_, MakeGlobalPointer> Type;\ }; TENSORMAPCONVERT(const) @@ -114,6 +114,16 @@ KERNELBROKERCONVERTREDUCTION(const) KERNELBROKERCONVERTREDUCTION() #undef KERNELBROKERCONVERTREDUCTION +#define KERNELBROKERCONVERTSLICEOP(CVQual)\ +template<typename StartIndices, typename Sizes, typename XprType>\ +struct ConvertToDeviceExpression<CVQual TensorSlicingOp <StartIndices, Sizes, XprType> >{\ + typedef CVQual TensorSlicingOp<StartIndices, Sizes, typename ConvertToDeviceExpression<XprType>::Type> Type;\ +}; + +KERNELBROKERCONVERTSLICEOP(const) +KERNELBROKERCONVERTSLICEOP() +#undef KERNELBROKERCONVERTSLICEOP + } // namespace internal } // namespace TensorSycl } // namespace Eigen diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h index 7ed3a3a56..c3152513c 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h @@ -45,17 +45,18 @@ struct ExprConstructor; /// specialisation of the \ref ExprConstructor struct when the node type is /// TensorMap #define TENSORMAP(CVQual)\ -template <typename Scalar_, int Options_, int Options2_, int Options3_, int NumIndices_, typename IndexType_,\ +template <typename T, int Options2_, int Options3_,\ template <class> class MakePointer_, size_t N, typename... Params>\ -struct ExprConstructor< CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer>,\ -CVQual PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_, MakePointer_>, N>, Params...>{\ - typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\ +struct ExprConstructor< CVQual TensorMap<T, Options2_, MakeGlobalPointer>,\ +CVQual PlaceHolder<CVQual TensorMap<T, Options3_, MakePointer_>, N>, Params...>{\ + typedef CVQual TensorMap<T, Options2_, MakeGlobalPointer> Type;\ Type expr;\ template <typename FuncDetector>\ ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)\ : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), fd.dimensions())) {}\ }; + TENSORMAP(const) TENSORMAP() #undef TENSORMAP @@ -224,6 +225,25 @@ SYCLREDUCTIONEXPR(const) SYCLREDUCTIONEXPR() #undef SYCLREDUCTIONEXPR + + +#define SYCLSLICEOPEXPR(CVQual)\ +template<typename StartIndices, typename Sizes, typename OrigXprType, typename XprType, typename... Params>\ +struct ExprConstructor<CVQual TensorSlicingOp <StartIndices, Sizes, OrigXprType> , CVQual TensorSlicingOp<StartIndices, Sizes, XprType>, Params... >{\ + typedef ExprConstructor<OrigXprType, XprType, Params...> my_xpr_type;\ + typedef CVQual TensorSlicingOp<StartIndices, Sizes, typename my_xpr_type::Type> Type ;\ + my_xpr_type xprExpr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.startIndices(), funcD.dimensions()) {}\ +}; + +SYCLSLICEOPEXPR(const) +SYCLSLICEOPEXPR() +#undef SYCLSLICEOPEXPR + + /// template deduction for \ref ExprConstructor struct template <typename OrigExpr, typename IndexExpr, typename FuncD, typename... Params> auto createDeviceExpression(FuncD &funcD, const utility::tuple::Tuple<Params...> &t) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h index b1da6858e..461aef128 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h @@ -191,6 +191,20 @@ template <typename OP, typename Dim, typename Expr, typename Dev> struct ExtractAccessor<TensorEvaluator<TensorReductionOp<OP, Dim, Expr>, Dev> > : ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> >{}; + +/// specialisation of the \ref ExtractAccessor struct when the node type is +/// const TensorSlicingOp. This is a special case where there is no OP +template <typename StartIndices, typename Sizes, typename XprType, typename Dev> +struct ExtractAccessor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){ + return AccessorConstructor::getTuple(cgh, eval.impl()); + } +}; + +template <typename StartIndices, typename Sizes, typename XprType, typename Dev> +struct ExtractAccessor<TensorEvaluator<TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > +:ExtractAccessor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> >{}; /// template deduction for \ref ExtractAccessor template <typename Evaluator> auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h index 427125343..ef56391ff 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h @@ -165,6 +165,23 @@ struct FunctorExtractor<TensorEvaluator<const TensorReductionOp<Op, Dims, ArgTyp template<typename Op, typename Dims, typename ArgType, template <class> class MakePointer_, typename Device> struct FunctorExtractor<TensorEvaluator<TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device>> : FunctorExtractor<TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device>>{}; + +/// specialisation of the \ref FunctorExtractor struct when the node type is +/// const TensorSlicingOp. This is an specialisation without OP so it has to be separated. +template <typename StartIndices, typename Sizes, typename XprType, typename Dev> +struct FunctorExtractor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > { + FunctorExtractor<TensorEvaluator<XprType, Dev> > xprExpr; + const StartIndices m_offsets; + const Sizes m_dimensions; + FunctorExtractor(const TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev>& expr) + : xprExpr(expr.impl()), m_offsets(expr.startIndices()), m_dimensions(expr.dimensions()) {} + EIGEN_STRONG_INLINE const StartIndices& startIndices() const {return m_offsets;} + EIGEN_STRONG_INLINE const Sizes& dimensions() const {return m_dimensions;} +}; + +template <typename StartIndices, typename Sizes, typename XprType, typename Dev> +struct FunctorExtractor<TensorEvaluator<TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > +:FunctorExtractor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > {}; /// template deduction function for FunctorExtractor template <typename Evaluator> auto inline extractFunctors(const Evaluator& evaluator)-> FunctorExtractor<Evaluator> { diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h index 25d1fac9b..54d2a8bdd 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h @@ -103,6 +103,15 @@ struct LeafCount<const TensorReductionOp<OP, Dim, Expr> > { template <typename OP, typename Dim, typename Expr> struct LeafCount<TensorReductionOp<OP, Dim, Expr> >: LeafCount<const TensorReductionOp<OP, Dim, Expr> >{}; +/// specialisation of the \ref LeafCount struct when the node type is const TensorSlicingOp +template <typename StartIndices, typename Sizes, typename XprType> +struct LeafCount<const TensorSlicingOp<StartIndices, Sizes, XprType> >:CategoryCount<XprType>{}; + +/// specialisation of the \ref LeafCount struct when the node type is TensorSlicingOp +template <typename StartIndices, typename Sizes, typename XprType> +struct LeafCount<TensorSlicingOp<StartIndices, Sizes, XprType> > +: LeafCount<const TensorSlicingOp<StartIndices, Sizes, XprType> >{}; + /// specialisation of the \ref LeafCount struct when the node type is TensorEvalToOp template <typename Expr> struct LeafCount<TensorEvalToOp<Expr> >: LeafCount<const TensorEvalToOp<Expr> >{}; diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h index d4c250c6d..0340b777f 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h @@ -122,9 +122,9 @@ ASSIGNEXPR() /// specialisation of the \ref PlaceHolderExpression when the node is /// TensorMap #define TENSORMAPEXPR(CVQual)\ -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_, size_t N>\ -struct PlaceHolderExpression< CVQual TensorMap< Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> {\ - typedef CVQual PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> Type;\ +template <typename T, int Options2_, template <class> class MakePointer_, size_t N>\ +struct PlaceHolderExpression< CVQual TensorMap< T, Options2_, MakePointer_>, N> {\ + typedef CVQual PlaceHolder<CVQual TensorMap<T, Options2_, MakePointer_>, N> Type;\ }; TENSORMAPEXPR(const) @@ -167,6 +167,20 @@ SYCLREDUCTION(const) SYCLREDUCTION() #undef SYCLREDUCTION + +/// specialisation of the \ref PlaceHolderExpression when the node is +/// TensorCwiseSelectOp +#define SLICEOPEXPR(CVQual)\ +template <typename StartIndices, typename Sizes, typename XprType, size_t N>\ +struct PlaceHolderExpression<CVQual TensorSlicingOp<StartIndices, Sizes, XprType>, N> {\ + typedef CVQual TensorSlicingOp<StartIndices, Sizes, typename CalculateIndex<N, XprType>::ArgType> Type;\ +}; + +SLICEOPEXPR(const) +SLICEOPEXPR() +#undef SLICEOPEXPR + + /// template deduction for \ref PlaceHolderExpression struct template <typename Expr> struct createPlaceHolderExpression { diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index b5fa1c845..f988cb465 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt @@ -146,6 +146,7 @@ if(EIGEN_TEST_CXX11) ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11") ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11") ei_add_test_sycl(cxx11_tensor_reduction_sycl "-std=c++11") + ei_add_test_sycl(cxx11_tensor_morphing_sycl "-std=c++11") endif(EIGEN_TEST_SYCL) # It should be safe to always run these tests as there is some fallback code for # older compiler that don't support cxx11. diff --git a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp index 7201bfe37..11b608ea4 100644 --- a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp +++ b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp @@ -47,7 +47,8 @@ static void test_broadcast_sycl(const Eigen::SyclDevice &sycl_device){ float * gpu_in_data = static_cast<float*>(sycl_device.allocate(input.dimensions().TotalSize()*sizeof(float))); float * gpu_out_data = static_cast<float*>(sycl_device.allocate(out.dimensions().TotalSize()*sizeof(float))); - TensorMap<Tensor<float, 4>> gpu_in(gpu_in_data, in_range); + TensorMap<TensorFixedSize<float, Sizes<2, 3, 5, 7>>> gpu_in(gpu_in_data, in_range); + //TensorMap<Tensor<float, 4>> gpu_in(gpu_in_data, in_range); TensorMap<Tensor<float, 4>> gpu_out(gpu_out_data, out_range); sycl_device.memcpyHostToDevice(gpu_in_data, input.data(),(input.dimensions().TotalSize())*sizeof(float)); gpu_out.device(sycl_device) = gpu_in.broadcast(broadcasts); diff --git a/unsupported/test/cxx11_tensor_morphing_sycl.cpp b/unsupported/test/cxx11_tensor_morphing_sycl.cpp new file mode 100644 index 000000000..8a03b826e --- /dev/null +++ b/unsupported/test/cxx11_tensor_morphing_sycl.cpp @@ -0,0 +1,84 @@ +// This file is part of Eigen, a lightweight C++ template library +// for linear algebra. +// +// Copyright (C) 2016 +// Mehdi Goli Codeplay Software Ltd. +// Ralph Potter Codeplay Software Ltd. +// Luke Iwanski Codeplay Software Ltd. +// Contact: <eigen@codeplay.com> +// Benoit Steiner <benoit.steiner.goog@gmail.com> +// +// This Source Code Form is subject to the terms of the Mozilla +// Public License v. 2.0. If a copy of the MPL was not distributed +// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. + + +#define EIGEN_TEST_NO_LONGDOUBLE +#define EIGEN_TEST_NO_COMPLEX +#define EIGEN_TEST_FUNC cxx11_tensor_morphing_sycl +#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int +#define EIGEN_USE_SYCL + + +#include "main.h" +#include <unsupported/Eigen/CXX11/Tensor> + +using Eigen::array; +using Eigen::SyclDevice; +using Eigen::Tensor; +using Eigen::TensorMap; + + +static void test_simple_slice(const Eigen::SyclDevice &sycl_device) +{ + int sizeDim1 = 2; + int sizeDim2 = 3; + int sizeDim3 = 5; + int sizeDim4 = 7; + int sizeDim5 = 11; + array<int, 5> tensorRange = {{sizeDim1, sizeDim2, sizeDim3, sizeDim4, sizeDim5}}; + Tensor<float, 5> tensor(tensorRange); + tensor.setRandom(); + array<int, 5> slice1_range ={{1, 1, 1, 1, 1}}; + Tensor<float, 5> slice1(slice1_range); + + float* gpu_data1 = static_cast<float*>(sycl_device.allocate(tensor.size()*sizeof(float))); + float* gpu_data2 = static_cast<float*>(sycl_device.allocate(slice1.size()*sizeof(float))); + TensorMap<Tensor<float, 5>> gpu1(gpu_data1, tensorRange); + TensorMap<Tensor<float, 5>> gpu2(gpu_data2, slice1_range); + Eigen::DSizes<ptrdiff_t, 5> indices(1,2,3,4,5); + Eigen::DSizes<ptrdiff_t, 5> sizes(1,1,1,1,1); + sycl_device.memcpyHostToDevice(gpu_data1, tensor.data(),(tensor.size())*sizeof(float)); + gpu2.device(sycl_device)=gpu1.slice(indices, sizes); + sycl_device.memcpyDeviceToHost(slice1.data(), gpu_data2,(slice1.size())*sizeof(float)); + VERIFY_IS_EQUAL(slice1(0,0,0,0,0), tensor(1,2,3,4,5)); + + + array<int, 5> slice2_range ={{1,1,2,2,3}}; + Tensor<float, 5> slice2(slice2_range); + float* gpu_data3 = static_cast<float*>(sycl_device.allocate(slice2.size()*sizeof(float))); + TensorMap<Tensor<float, 5>> gpu3(gpu_data3, slice2_range); + Eigen::DSizes<ptrdiff_t, 5> indices2(1,1,3,4,5); + Eigen::DSizes<ptrdiff_t, 5> sizes2(1,1,2,2,3); + gpu3.device(sycl_device)=gpu1.slice(indices2, sizes2); + sycl_device.memcpyDeviceToHost(slice2.data(), gpu_data3,(slice2.size())*sizeof(float)); + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 2; ++j) { + for (int k = 0; k < 3; ++k) { + VERIFY_IS_EQUAL(slice2(0,0,i,j,k), tensor(1,1,3+i,4+j,5+k)); + } + } + } + sycl_device.deallocate(gpu_data1); + sycl_device.deallocate(gpu_data2); + sycl_device.deallocate(gpu_data3); +} + +void test_cxx11_tensor_morphing_sycl() +{ + /// Currentlly it only works on cpu. Adding GPU cause LLVM ERROR in cunstructing OpenCL Kernel at runtime. + cl::sycl::cpu_selector s; + Eigen::SyclDevice sycl_device(s); + CALL_SUBTEST(test_simple_slice(sycl_device)); + +} |