diff options
author | Mehdi Goli <mehdi.goli@codeplay.com> | 2016-11-10 18:45:12 +0000 |
---|---|---|
committer | Mehdi Goli <mehdi.goli@codeplay.com> | 2016-11-10 18:45:12 +0000 |
commit | 2e704d4257f235dd1f3224cd590e4fca4e3aaf96 (patch) | |
tree | e56e7d5b886830cde7cf6fa90c448880fb578665 /unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h | |
parent | 75c080b1762b8b83f6c2bb7baf95478a049b45d4 (diff) |
Adding Memset; optimising MecopyDeviceToHost by removing double copying;
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h')
-rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h | 13 |
1 files changed, 3 insertions, 10 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h index 3daecb045..db23bd7b0 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h @@ -188,15 +188,8 @@ struct InnerReducer<Self, Op, const Eigen::SyclDevice> { typedef const typename Self::ChildType HostExpr; /// this is the child of reduction typedef typename TensorSycl::internal::createPlaceHolderExpression<HostExpr>::Type PlaceHolderExpr; auto functors = TensorSycl::internal::extractFunctors(self.impl()); - - size_t tileSize =dev.m_queue.get_device(). template get_info<cl::sycl::info::device::max_work_group_size>()/2; - - size_t GRange=num_coeffs_to_preserve; - if (tileSize>GRange) tileSize=GRange; - else if(GRange>tileSize){ - size_t xMode = GRange % tileSize; - if (xMode != 0) GRange += (tileSize - xMode); - } + size_t range, GRange, tileSize; + dev.parallel_for_setup(num_coeffs_to_preserve, tileSize, range, GRange); // getting final out buffer at the moment the created buffer is true because there is no need for assign /// creating the shared memory for calculating reduction. /// This one is used to collect all the reduced value of shared memory as we dont have global barrier on GPU. Once it is saved we can @@ -223,7 +216,7 @@ struct InnerReducer<Self, Op, const Eigen::SyclDevice> { auto device_self_evaluator = Eigen::TensorEvaluator<decltype(device_self_expr), Eigen::DefaultDevice>(device_self_expr, Eigen::DefaultDevice()); /// const cast added as a naive solution to solve the qualifier drop error auto globalid=itemID.get_global_linear_id(); - if (globalid< static_cast<size_t>(num_coeffs_to_preserve)) { + if (globalid< range) { typename DeiceSelf::CoeffReturnType accum = functor.initialize(); GenericDimReducer<DeiceSelf::NumReducedDims-1, DeiceSelf, Op>::reduce(device_self_evaluator, device_self_evaluator.firstInput(globalid),const_cast<Op&>(functor), &accum); functor.finalize(accum); |