devtools/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2026-04-10 11:34:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

Adding TensorFixsize; adding sycl device memcpy; adding insial stage of slicing.

Browse Source
This commit is contained in:
Mehdi Goli
2016-11-14 17:51:57 +00:00
parent a5c3f15682
commit f8ca893976
11 changed files with 251 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

67
unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
View File

@@ -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();

19
unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h
View File

@@ -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
{

16
unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h
View File

@@ -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

28
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h
View File

@@ -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)

14
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h
View File

@@ -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)

17
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h
View File

@@ -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> {

9
unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h
View File

@@ -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> >{};

20
unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
View File

@@ -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 {