GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-04-24 02:29:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

Merged in mehdi_goli/upstr_benoit/TensorSYCLImageVolumePatchFixed (pull request PR-14)

Browse Source 
Applying Benoit's comment for Fixing ImageVolumePatch.

* Applying Benoit's comment for Fixing ImageVolumePatch. Fixing conflict on cmake file.

* Fixing dealocation of the memory in ImagePatch test for SYCL.

* Fixing the automerge issue.
This commit is contained in:
Benoit Steiner 2017-07-06 05:08:13 +00:00
parent c92faf9d84
commit 62b4634ebe
5 changed files with 86 additions and 39 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

34
unsupported/Eigen/CXX11/src/Tensor/TensorImagePatch.h
View File

@ -27,6 +27,7 @@ namespace Eigen {
* patch_cols, and 1 for all the additional dimensions.
*/
namespace internal {
template<DenseIndex Rows, DenseIndex Cols, typename XprType>
struct traits<TensorImagePatchOp<Rows, Cols, XprType> > : public traits<XprType>
{
@ -93,6 +94,24 @@ class TensorImagePatchOp : public TensorBase<TensorImagePatchOp<Rows, Cols, XprT
m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(PADDING_VALID), m_padding_value(padding_value) {}
#ifdef EIGEN_USE_SYCL // this is work around for sycl as Eigen could not use c++11 deligate constructor feature
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorImagePatchOp(const XprType& expr, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex row_strides, DenseIndex col_strides,
DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
bool padding_explicit, DenseIndex padding_top, DenseIndex padding_bottom,
DenseIndex padding_left, DenseIndex padding_right, PaddingType padding_type,
Scalar padding_value)
: m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_row_strides(row_strides), m_col_strides(col_strides),
m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(padding_explicit), m_padding_top(padding_top), m_padding_bottom(padding_bottom),
m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(padding_type), m_padding_value(padding_value) {}
#endif
EIGEN_DEVICE_FUNC
DenseIndex patch_rows() const { return m_patch_rows; }
EIGEN_DEVICE_FUNC
@ -172,7 +191,11 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
RawAccess = false
};
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
#ifdef __SYCL_DEVICE_ONLY__
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator( const XprType op, const Device& device)
#else
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator( const XprType& op, const Device& device)
#endif
: m_impl(op.expression(), device)
#ifdef EIGEN_USE_SYCL
, m_op(op)
@ -429,6 +452,7 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
#ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression tree on the device
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; }
#endif
@ -511,10 +535,10 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
Scalar m_paddingValue;
TensorEvaluator<ArgType, Device> m_impl;
#ifdef EIGEN_USE_SYCL
const XprType& m_op;
#endif
#ifdef EIGEN_USE_SYCL
// Required for SYCL in order to construct the expression tree on the device
XprType m_op;
#endif
};

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

@ -448,8 +448,8 @@ struct ExprConstructor<CVQual TensorImagePatchOp<Rows, Cols, OrigXprType>, CVQu
template <typename FuncDetector>\
ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_row_strides, funcD.m_col_strides,\
funcD.m_in_row_strides, funcD.m_in_col_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, \
funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_value, funcD.m_padding_type, funcD.m_padding_explicit){}\
funcD.m_in_row_strides, funcD.m_in_col_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, funcD.m_padding_explicit, \
funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_type, funcD.m_padding_value){}\
};
SYCLTENSORIMAGEPATCHOPEXPR(const)
@ -468,8 +468,8 @@ struct ExprConstructor<CVQual TensorVolumePatchOp<Planes, Rows, Cols, OrigXprTy
ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_planes, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_plane_strides, funcD.m_row_strides, funcD.m_col_strides,\
funcD.m_in_plane_strides, funcD.m_in_row_strides, funcD.m_in_col_strides,funcD.m_plane_inflate_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, \
funcD.m_padding_top_z, funcD.m_padding_bottom_z, funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_value,\
funcD.m_padding_type, funcD.m_padding_explicit){\
funcD.m_padding_explicit, funcD.m_padding_top_z, funcD.m_padding_bottom_z, funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, \
funcD.m_padding_type, funcD.m_padding_value ){\
}\
};

23
unsupported/Eigen/CXX11/src/Tensor/TensorVolumePatch.h
View File

@ -22,6 +22,7 @@ namespace Eigen {
* dimensions.
*/
namespace internal {
template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType>
struct traits<TensorVolumePatchOp<Planes, Rows, Cols, XprType> > : public traits<XprType>
{
@ -34,6 +35,7 @@ struct traits<TensorVolumePatchOp<Planes, Rows, Cols, XprType> > : public traits
static const int NumDimensions = XprTraits::NumDimensions + 1;
static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
};
template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType>
@ -89,6 +91,24 @@ class TensorVolumePatchOp : public TensorBase<TensorVolumePatchOp<Planes, Rows,
m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(PADDING_VALID), m_padding_value(padding_value) {}
#ifdef EIGEN_USE_SYCL // this is work around for sycl as Eigen could not use c++11 deligate constructor feature
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorVolumePatchOp(const XprType& expr, DenseIndex patch_planes, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex plane_strides, DenseIndex row_strides, DenseIndex col_strides,
DenseIndex in_plane_strides, DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex plane_inflate_strides, DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
bool padding_explicit, DenseIndex padding_top_z, DenseIndex padding_bottom_z,
DenseIndex padding_top, DenseIndex padding_bottom, DenseIndex padding_left,
DenseIndex padding_right, PaddingType padding_type, Scalar padding_value)
: m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides),
m_col_inflate_strides(col_inflate_strides), m_padding_explicit(padding_explicit), m_padding_top_z(padding_top_z),
m_padding_bottom_z(padding_bottom_z), m_padding_top(padding_top), m_padding_bottom(padding_bottom), m_padding_left(padding_left),
m_padding_right(padding_right), m_padding_type(padding_type), m_padding_value(padding_value) {}
#endif
EIGEN_DEVICE_FUNC
DenseIndex patch_planes() const { return m_patch_planes; }
EIGEN_DEVICE_FUNC
@ -515,6 +535,7 @@ struct TensorEvaluator<const TensorVolumePatchOp<Planes, Rows, Cols, ArgType>, D
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
#ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression on the device
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; }
#endif
@ -614,8 +635,10 @@ struct TensorEvaluator<const TensorVolumePatchOp<Planes, Rows, Cols, ArgType>, D
TensorEvaluator<ArgType, Device> m_impl;
#ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression on the device
XprType m_op;
#endif
};

4
unsupported/test/cxx11_tensor_image_patch_sycl.cpp
View File

@ -13,7 +13,7 @@
#define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX
#define EIGEN_TEST_FUNC cxx11_tensor_image_patchOP_sycl
#define EIGEN_TEST_FUNC cxx11_tensor_image_patch_sycl
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
#define EIGEN_USE_SYCL
@ -1084,7 +1084,7 @@ test_patch_padding_same_sycl<DataType, int64_t>(sycl_device);
test_patch_no_extra_dim_sycl<DataType, int64_t>(sycl_device);
test_imagenet_patches_sycl<DataType, int64_t>(sycl_device);
}
void test_cxx11_tensor_image_patchOP_sycl()
void test_cxx11_tensor_image_patch_sycl()
{
for (const auto& device :Eigen::get_sycl_supported_devices()) {
CALL_SUBTEST(sycl_tensor_image_patch_test_per_device<float>(device));

4
unsupported/test/cxx11_tensor_volume_patch_sycl.cpp
View File

@ -13,7 +13,7 @@
#define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX
#define EIGEN_TEST_FUNC cxx11_tensor_volume_patchOP_sycl
#define EIGEN_TEST_FUNC cxx11_tensor_volume_patch_sycl
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
#define EIGEN_USE_SYCL
@ -214,7 +214,7 @@ std::cout << "Running on " << s.template get_info<cl::sycl::info::device::name>(
test_single_voxel_patch_sycl<DataType, int64_t>(sycl_device);
test_entire_volume_patch_sycl<DataType, int64_t>(sycl_device);
}
void test_cxx11_tensor_volume_patchOP_sycl()
void test_cxx11_tensor_volume_patch_sycl()
{
for (const auto& device :Eigen::get_sycl_supported_devices()) {
CALL_SUBTEST(sycl_tensor_volume_patch_test_per_device<float>(device));