From 88534ba623421c956d8ffcda2d27f41d704d15ef Mon Sep 17 00:00:00 2001 From: Stanislaw Halik Date: Tue, 3 Jul 2018 07:37:12 +0200 Subject: update eigen --- .../unsupported/test/cxx11_tensor_reverse_sycl.cpp | 221 --------------------- 1 file changed, 221 deletions(-) delete mode 100644 eigen/unsupported/test/cxx11_tensor_reverse_sycl.cpp (limited to 'eigen/unsupported/test/cxx11_tensor_reverse_sycl.cpp') diff --git a/eigen/unsupported/test/cxx11_tensor_reverse_sycl.cpp b/eigen/unsupported/test/cxx11_tensor_reverse_sycl.cpp deleted file mode 100644 index 2f54844..0000000 --- a/eigen/unsupported/test/cxx11_tensor_reverse_sycl.cpp +++ /dev/null @@ -1,221 +0,0 @@ -// This file is part of Eigen, a lightweight C++ template library -// for linear algebra. -// -// Copyright (C) 2015 -// Mehdi Goli Codeplay Software Ltd. -// Ralph Potter Codeplay Software Ltd. -// Luke Iwanski Codeplay Software Ltd. -// Contact: -// -// This Source Code Form is subject to the terms of the Mozilla -// Public License v. 2.0. If a copy of the MPL was not distributed -// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. - -#define EIGEN_TEST_NO_LONGDOUBLE -#define EIGEN_TEST_NO_COMPLEX -#define EIGEN_TEST_FUNC cxx11_tensor_reverse_sycl -#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t -#define EIGEN_USE_SYCL - -#include "main.h" -#include - - -template -static void test_simple_reverse(const Eigen::SyclDevice& sycl_device) { - - IndexType dim1 = 2; - IndexType dim2 = 3; - IndexType dim3 = 5; - IndexType dim4 = 7; - - array tensorRange = {{dim1, dim2, dim3, dim4}}; - Tensor tensor(tensorRange); - Tensor reversed_tensor(tensorRange); - tensor.setRandom(); - - array dim_rev; - dim_rev[0] = false; - dim_rev[1] = true; - dim_rev[2] = true; - dim_rev[3] = false; - - DataType* gpu_in_data = static_cast(sycl_device.allocate(tensor.dimensions().TotalSize()*sizeof(DataType))); - DataType* gpu_out_data =static_cast(sycl_device.allocate(reversed_tensor.dimensions().TotalSize()*sizeof(DataType))); - - TensorMap > in_gpu(gpu_in_data, tensorRange); - TensorMap > out_gpu(gpu_out_data, tensorRange); - - sycl_device.memcpyHostToDevice(gpu_in_data, tensor.data(),(tensor.dimensions().TotalSize())*sizeof(DataType)); - out_gpu.device(sycl_device) = in_gpu.reverse(dim_rev); - sycl_device.memcpyDeviceToHost(reversed_tensor.data(), gpu_out_data, reversed_tensor.dimensions().TotalSize()*sizeof(DataType)); - // Check that the CPU and GPU reductions return the same result. - for (IndexType i = 0; i < 2; ++i) { - for (IndexType j = 0; j < 3; ++j) { - for (IndexType k = 0; k < 5; ++k) { - for (IndexType l = 0; l < 7; ++l) { - VERIFY_IS_EQUAL(tensor(i,j,k,l), reversed_tensor(i,2-j,4-k,l)); - } - } - } - } - dim_rev[0] = true; - dim_rev[1] = false; - dim_rev[2] = false; - dim_rev[3] = false; - - out_gpu.device(sycl_device) = in_gpu.reverse(dim_rev); - sycl_device.memcpyDeviceToHost(reversed_tensor.data(), gpu_out_data, reversed_tensor.dimensions().TotalSize()*sizeof(DataType)); - - for (IndexType i = 0; i < 2; ++i) { - for (IndexType j = 0; j < 3; ++j) { - for (IndexType k = 0; k < 5; ++k) { - for (IndexType l = 0; l < 7; ++l) { - VERIFY_IS_EQUAL(tensor(i,j,k,l), reversed_tensor(1-i,j,k,l)); - } - } - } - } - - dim_rev[0] = true; - dim_rev[1] = false; - dim_rev[2] = false; - dim_rev[3] = true; - out_gpu.device(sycl_device) = in_gpu.reverse(dim_rev); - sycl_device.memcpyDeviceToHost(reversed_tensor.data(), gpu_out_data, reversed_tensor.dimensions().TotalSize()*sizeof(DataType)); - - for (IndexType i = 0; i < 2; ++i) { - for (IndexType j = 0; j < 3; ++j) { - for (IndexType k = 0; k < 5; ++k) { - for (IndexType l = 0; l < 7; ++l) { - VERIFY_IS_EQUAL(tensor(i,j,k,l), reversed_tensor(1-i,j,k,6-l)); - } - } - } - } - - sycl_device.deallocate(gpu_in_data); - sycl_device.deallocate(gpu_out_data); -} - - - -template -static void test_expr_reverse(const Eigen::SyclDevice& sycl_device, bool LValue) -{ - IndexType dim1 = 2; - IndexType dim2 = 3; - IndexType dim3 = 5; - IndexType dim4 = 7; - - array tensorRange = {{dim1, dim2, dim3, dim4}}; - Tensor tensor(tensorRange); - Tensor expected(tensorRange); - Tensor result(tensorRange); - tensor.setRandom(); - - array dim_rev; - dim_rev[0] = false; - dim_rev[1] = true; - dim_rev[2] = false; - dim_rev[3] = true; - - DataType* gpu_in_data = static_cast(sycl_device.allocate(tensor.dimensions().TotalSize()*sizeof(DataType))); - DataType* gpu_out_data_expected =static_cast(sycl_device.allocate(expected.dimensions().TotalSize()*sizeof(DataType))); - DataType* gpu_out_data_result =static_cast(sycl_device.allocate(result.dimensions().TotalSize()*sizeof(DataType))); - - TensorMap > in_gpu(gpu_in_data, tensorRange); - TensorMap > out_gpu_expected(gpu_out_data_expected, tensorRange); - TensorMap > out_gpu_result(gpu_out_data_result, tensorRange); - - - sycl_device.memcpyHostToDevice(gpu_in_data, tensor.data(),(tensor.dimensions().TotalSize())*sizeof(DataType)); - - if (LValue) { - out_gpu_expected.reverse(dim_rev).device(sycl_device) = in_gpu; - } else { - out_gpu_expected.device(sycl_device) = in_gpu.reverse(dim_rev); - } - sycl_device.memcpyDeviceToHost(expected.data(), gpu_out_data_expected, expected.dimensions().TotalSize()*sizeof(DataType)); - - - array src_slice_dim; - src_slice_dim[0] = 2; - src_slice_dim[1] = 3; - src_slice_dim[2] = 1; - src_slice_dim[3] = 7; - array src_slice_start; - src_slice_start[0] = 0; - src_slice_start[1] = 0; - src_slice_start[2] = 0; - src_slice_start[3] = 0; - array dst_slice_dim = src_slice_dim; - array dst_slice_start = src_slice_start; - - for (IndexType i = 0; i < 5; ++i) { - if (LValue) { - out_gpu_result.slice(dst_slice_start, dst_slice_dim).reverse(dim_rev).device(sycl_device) = - in_gpu.slice(src_slice_start, src_slice_dim); - } else { - out_gpu_result.slice(dst_slice_start, dst_slice_dim).device(sycl_device) = - in_gpu.slice(src_slice_start, src_slice_dim).reverse(dim_rev); - } - src_slice_start[2] += 1; - dst_slice_start[2] += 1; - } - sycl_device.memcpyDeviceToHost(result.data(), gpu_out_data_result, result.dimensions().TotalSize()*sizeof(DataType)); - - for (IndexType i = 0; i < expected.dimension(0); ++i) { - for (IndexType j = 0; j < expected.dimension(1); ++j) { - for (IndexType k = 0; k < expected.dimension(2); ++k) { - for (IndexType l = 0; l < expected.dimension(3); ++l) { - VERIFY_IS_EQUAL(result(i,j,k,l), expected(i,j,k,l)); - } - } - } - } - - dst_slice_start[2] = 0; - result.setRandom(); - sycl_device.memcpyHostToDevice(gpu_out_data_result, result.data(),(result.dimensions().TotalSize())*sizeof(DataType)); - for (IndexType i = 0; i < 5; ++i) { - if (LValue) { - out_gpu_result.slice(dst_slice_start, dst_slice_dim).reverse(dim_rev).device(sycl_device) = - in_gpu.slice(dst_slice_start, dst_slice_dim); - } else { - out_gpu_result.slice(dst_slice_start, dst_slice_dim).device(sycl_device) = - in_gpu.reverse(dim_rev).slice(dst_slice_start, dst_slice_dim); - } - dst_slice_start[2] += 1; - } - sycl_device.memcpyDeviceToHost(result.data(), gpu_out_data_result, result.dimensions().TotalSize()*sizeof(DataType)); - - for (IndexType i = 0; i < expected.dimension(0); ++i) { - for (IndexType j = 0; j < expected.dimension(1); ++j) { - for (IndexType k = 0; k < expected.dimension(2); ++k) { - for (IndexType l = 0; l < expected.dimension(3); ++l) { - VERIFY_IS_EQUAL(result(i,j,k,l), expected(i,j,k,l)); - } - } - } - } -} - - - -template void sycl_reverse_test_per_device(const cl::sycl::device& d){ - std::cout << "Running on " << d.template get_info() << std::endl; - QueueInterface queueInterface(d); - auto sycl_device = Eigen::SyclDevice(&queueInterface); - test_simple_reverse(sycl_device); - test_simple_reverse(sycl_device); - test_expr_reverse(sycl_device, false); - test_expr_reverse(sycl_device, false); - test_expr_reverse(sycl_device, true); - test_expr_reverse(sycl_device, true); -} -void test_cxx11_tensor_reverse_sycl() { - for (const auto& device :Eigen::get_sycl_supported_devices()) { - CALL_SUBTEST(sycl_reverse_test_per_device(device)); - } -} -- cgit v1.2.3