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

Fixing TensorReductionSycl for min and max.

Browse Source
This commit is contained in:
Mehdi Goli
2017-02-06 18:05:23 +00:00
parent bc128f9f3b
commit 42bd5c4e7b
4 changed files with 33 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h
View File

@@ -26,10 +26,10 @@ namespace Eigen {
namespace internal {
template<typename CoeffReturnType> struct syclGenericBufferReducer{
template<typename BufferTOut, typename BufferTIn>
static void run(BufferTOut& bufOut, BufferTIn& bufI, const Eigen::SyclDevice& dev, size_t length, size_t local){
template<typename OP, typename BufferTOut, typename BufferTIn>
static void run(OP op, BufferTOut& bufOut, BufferTIn& bufI, const Eigen::SyclDevice& dev, size_t length, size_t local){
do {
auto f = [length, local, bufOut, &bufI](cl::sycl::handler& h) mutable {
auto f = [length, local, op, &bufOut, &bufI](cl::sycl::handler& h) mutable {
cl::sycl::nd_range<1> r{cl::sycl::range<1>{std::max(length, local)},
cl::sycl::range<1>{std::min(length, local)}};
/* Two accessors are used: one to the buffer that is being reduced,
@@ -43,7 +43,7 @@ static void run(BufferTOut& bufOut, BufferTIn& bufI, const Eigen::SyclDevice& de
/* The parallel_for invocation chosen is the variant with an nd_item
* parameter, since the code requires barriers for correctness. */
h.parallel_for(r, TensorSycl::internal::GenericKernelReducer< CoeffReturnType, OutputAccessor, InputAccessor, LocalAccessor>(aOut, aI, scratch, length, local));
h.parallel_for(r, TensorSycl::internal::GenericKernelReducer<CoeffReturnType, OP, OutputAccessor, InputAccessor, LocalAccessor>(op, aOut, aI, scratch, length, local));
};
dev.sycl_queue().submit(f);
dev.asynchronousExec();
@@ -123,7 +123,7 @@ struct FullReducer<Self, Op, const Eigen::SyclDevice, Vectorizable> {
// getting final out buffer at the moment the created buffer is true because there is no need for assign
auto out_buffer =dev.get_sycl_buffer(output);
/// This is used to recursively reduce the tmp value to an element of 1;
syclGenericBufferReducer<CoeffReturnType>::run(out_buffer, temp_global_buffer,dev, GRange, outTileSize);
syclGenericBufferReducer<CoeffReturnType>::run(reducer, out_buffer, temp_global_buffer,dev, GRange, outTileSize);
}
};

2
unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h
View File

@@ -35,7 +35,7 @@ namespace Eigen {
namespace TensorSycl {
namespace internal {
template<typename CoeffReturnType, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer;
template<typename CoeffReturnType, typename OP, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer;
/// This struct is used for special expression nodes with no operations (for example assign and selectOP).

30
unsupported/Eigen/CXX11/src/Tensor/TensorSyclFunctors.h
View File

@@ -18,13 +18,14 @@ namespace Eigen {
namespace TensorSycl {
namespace internal {
template<typename CoeffReturnType, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer{
template<typename CoeffReturnType, typename OP, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer{
OP op;
OutputAccessor aOut;
InputAccessor aI;
LocalAccessor scratch;
size_t length, local;
GenericKernelReducer(OutputAccessor aOut_, InputAccessor aI_, LocalAccessor scratch_, size_t length_, size_t local_)
: aOut(aOut_), aI(aI_), scratch(scratch_), length(length_), local(local_){}
GenericKernelReducer(OP op_, OutputAccessor aOut_, InputAccessor aI_, LocalAccessor scratch_, size_t length_, size_t local_)
: op(op_), aOut(aOut_), aI(aI_), scratch(scratch_), length(length_), local(local_){}
void operator()(cl::sycl::nd_item<1> itemID) {
size_t globalid = itemID.get_global(0);
size_t localid = itemID.get_local(0);
@@ -44,7 +45,12 @@ namespace internal {
auto min = (length < local) ? length : local;
for (size_t offset = min / 2; offset > 0; offset /= 2) {
if (localid < offset) {
scratch[localid] += scratch[localid + offset];
auto accum = op.initialize();
op.reduce(scratch[localid], &accum);
op.reduce(scratch[localid + offset], &accum);
op.finalize(accum);
scratch[localid]=accum;
//scratch[localid] += scratch[localid + offset];
}
itemID.barrier(cl::sycl::access::fence_space::local_space);
}
@@ -131,11 +137,21 @@ public:
if(globalid<rng)
tmp_global_accessor.get_pointer()[globalid]=Eigen::internal::InnerMostDimReducer<decltype(device_self_evaluator), Op, false>::reduce(device_self_evaluator, static_cast<typename DevExpr::Index>(red_factor*globalid), red_factor, const_cast<Op&>(op));
else
tmp_global_accessor.get_pointer()[globalid]=static_cast<CoeffReturnType>(0);
tmp_global_accessor.get_pointer()[globalid]=static_cast<CoeffReturnType>(op.initialize());
if(remaining!=0 && globalid==0 )
if(remaining!=0 && globalid==0 ){
// this will add the rest of input buffer when the input size is not devidable to red_factor.
tmp_global_accessor.get_pointer()[0]+=Eigen::internal::InnerMostDimReducer<decltype(device_self_evaluator), Op, false>::reduce(device_self_evaluator, static_cast<typename DevExpr::Index>(red_factor*(rng)), static_cast<typename DevExpr::Index>(remaining), const_cast<Op&>(op));
// tmp_global_accessor.get_pointer()[0]+=
auto remaining_reduce =Eigen::internal::InnerMostDimReducer<decltype(device_self_evaluator), Op, false>::
reduce(device_self_evaluator, static_cast<typename DevExpr::Index>(red_factor*(rng)), static_cast<typename DevExpr::Index>(remaining), const_cast<Op&>(op));
auto accum = op.initialize();
op.reduce(tmp_global_accessor.get_pointer()[0], &accum);
op.reduce(remaining_reduce, &accum);
op.finalize(accum);
tmp_global_accessor.get_pointer()[0]=accum;
}
}
};