diff options
author | Mehdi Goli <mehdi.goli@codeplay.com> | 2017-02-07 15:43:17 +0000 |
---|---|---|
committer | Mehdi Goli <mehdi.goli@codeplay.com> | 2017-02-07 15:43:17 +0000 |
commit | 0ee97b60c256b31a98838324ce1909247a0133d2 (patch) | |
tree | fe2847e82a5173c4bffc21dc9da5c083a2ef7aa5 /unsupported/test/cxx11_tensor_reduction_sycl.cpp | |
parent | 42bd5c4e7b8f4b5875ae256e7ac20310161d8470 (diff) |
Adding mean to TensorReductionSycl.h
Diffstat (limited to 'unsupported/test/cxx11_tensor_reduction_sycl.cpp')
-rw-r--r-- | unsupported/test/cxx11_tensor_reduction_sycl.cpp | 53 |
1 files changed, 44 insertions, 9 deletions
diff --git a/unsupported/test/cxx11_tensor_reduction_sycl.cpp b/unsupported/test/cxx11_tensor_reduction_sycl.cpp index 251091f5b..440d48bca 100644 --- a/unsupported/test/cxx11_tensor_reduction_sycl.cpp +++ b/unsupported/test/cxx11_tensor_reduction_sycl.cpp @@ -22,7 +22,7 @@ template <typename DataType, int DataLayout, typename IndexType> -static void test_full_reductions_sycl(const Eigen::SyclDevice& sycl_device) { +static void test_full_reductions_mean_sycl(const Eigen::SyclDevice& sycl_device) { const IndexType num_rows = 452; const IndexType num_cols = 765; @@ -34,6 +34,37 @@ static void test_full_reductions_sycl(const Eigen::SyclDevice& sycl_device) { in.setRandom(); + full_redux = in.mean(); + + DataType* gpu_in_data = static_cast<DataType*>(sycl_device.allocate(in.dimensions().TotalSize()*sizeof(DataType))); + DataType* gpu_out_data =(DataType*)sycl_device.allocate(sizeof(DataType)); + + TensorMap<Tensor<DataType, 2, DataLayout, IndexType> > in_gpu(gpu_in_data, tensorRange); + TensorMap<Tensor<DataType, 0, DataLayout, IndexType> > out_gpu(gpu_out_data); + + sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),(in.dimensions().TotalSize())*sizeof(DataType)); + out_gpu.device(sycl_device) = in_gpu.mean(); + sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data, sizeof(DataType)); + // Check that the CPU and GPU reductions return the same result. + VERIFY_IS_APPROX(full_redux_gpu(), full_redux()); + sycl_device.deallocate(gpu_in_data); + sycl_device.deallocate(gpu_out_data); +} + + +template <typename DataType, int DataLayout, typename IndexType> +static void test_full_reductions_min_sycl(const Eigen::SyclDevice& sycl_device) { + + const IndexType num_rows = 876; + const IndexType num_cols = 953; + array<IndexType, 2> tensorRange = {{num_rows, num_cols}}; + + Tensor<DataType, 2, DataLayout, IndexType> in(tensorRange); + Tensor<DataType, 0, DataLayout, IndexType> full_redux; + Tensor<DataType, 0, DataLayout, IndexType> full_redux_gpu; + + in.setRandom(); + full_redux = in.minimum(); DataType* gpu_in_data = static_cast<DataType*>(sycl_device.allocate(in.dimensions().TotalSize()*sizeof(DataType))); @@ -50,8 +81,10 @@ static void test_full_reductions_sycl(const Eigen::SyclDevice& sycl_device) { sycl_device.deallocate(gpu_in_data); sycl_device.deallocate(gpu_out_data); } + + template <typename DataType, int DataLayout, typename IndexType> -static void test_first_dim_reductions_sycl(const Eigen::SyclDevice& sycl_device) { +static void test_first_dim_reductions_max_sycl(const Eigen::SyclDevice& sycl_device) { IndexType dim_x = 145; IndexType dim_y = 1; @@ -90,7 +123,7 @@ static void test_first_dim_reductions_sycl(const Eigen::SyclDevice& sycl_device) } template <typename DataType, int DataLayout, typename IndexType> -static void test_last_dim_reductions_sycl(const Eigen::SyclDevice &sycl_device) { +static void test_last_dim_reductions_sum_sycl(const Eigen::SyclDevice &sycl_device) { IndexType dim_x = 567; IndexType dim_y = 1; @@ -132,12 +165,14 @@ template<typename DataType> void sycl_reduction_test_per_device(const cl::sycl:: QueueInterface queueInterface(d); auto sycl_device = Eigen::SyclDevice(&queueInterface); - test_full_reductions_sycl<DataType, RowMajor, int64_t>(sycl_device); - test_first_dim_reductions_sycl<DataType, RowMajor, int64_t>(sycl_device); - test_last_dim_reductions_sycl<DataType, RowMajor, int64_t>(sycl_device); - test_full_reductions_sycl<DataType, ColMajor, int64_t>(sycl_device); - test_first_dim_reductions_sycl<DataType, ColMajor, int64_t>(sycl_device); - test_last_dim_reductions_sycl<DataType, ColMajor, int64_t>(sycl_device); + test_full_reductions_mean_sycl<DataType, RowMajor, int64_t>(sycl_device); + test_full_reductions_min_sycl<DataType, RowMajor, int64_t>(sycl_device); + test_first_dim_reductions_max_sycl<DataType, RowMajor, int64_t>(sycl_device); + test_last_dim_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device); + test_full_reductions_mean_sycl<DataType, ColMajor, int64_t>(sycl_device); + test_full_reductions_min_sycl<DataType, ColMajor, int64_t>(sycl_device); + test_first_dim_reductions_max_sycl<DataType, ColMajor, int64_t>(sycl_device); + test_last_dim_reductions_sum_sycl<DataType, ColMajor, int64_t>(sycl_device); } void test_cxx11_tensor_reduction_sycl() { for (const auto& device :Eigen::get_sycl_supported_devices()) { |