mirror/eigen
2
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2024-12-27 07:29:52 +08:00
Code Issues Packages Projects Releases Wiki Activity
1fd5ce1002
eigen/unsupported/test/cxx11_tensor_reduction_sycl.cpp

1015 lines
41 KiB
C++
Raw Normal View History

Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
// 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: <eigen@codeplay.com>
//
// 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
Get rid of EIGEN_TEST_FUNC, unit tests must now be declared with EIGEN_DECLARE_TEST(mytest) { /* code */ }. This provide several advantages: - more flexibility in designing unit tests - unit tests can be glued to speed up compilation - unit tests are compiled with same predefined macros, which is a requirement for zapcc
2018-07-17 20:46:15 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
#define EIGEN_USE_SYCL
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
#define EIGEN_HAS_CONSTEXPR 1
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
#include "main.h"
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
#include <unsupported/Eigen/CXX11/Tensor>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
template <typename DataType, int DataLayout, typename IndexType>
static void test_full_reductions_sum_sycl(
const Eigen::SyclDevice& sycl_device) {
const IndexType num_rows = 753;
const IndexType num_cols = 537;
array<IndexType, 2> tensorRange = {{num_rows, num_cols}};
array<IndexType, 2> outRange = {{1, 1}};
Tensor<DataType, 2, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> full_redux(outRange);
Tensor<DataType, 2, DataLayout, IndexType> full_redux_gpu(outRange);
in.setRandom();
auto dim = DSizes<IndexType, 2>(1, 1);
full_redux = in.sum().reshape(dim);
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) * (full_redux_gpu.dimensions().TotalSize()));
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> out_gpu(gpu_out_data,
outRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.sum().reshape(dim);
sycl_device.memcpyDeviceToHost(
full_redux_gpu.data(), gpu_out_data,
(full_redux_gpu.dimensions().TotalSize()) * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
std::cout << "SYCL FULL :" << full_redux_gpu(0, 0)
<< ", CPU FULL: " << full_redux(0, 0) << "\n";
VERIFY_IS_APPROX(full_redux_gpu(0, 0), full_redux(0, 0));
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
template <typename DataType, int DataLayout, typename IndexType>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
static void test_full_reductions_sum_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using scalar_tensor = Tensor<DataType, 0, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
const IndexType n_elems = internal::array_prod(tensor_range);
data_tensor in(tensor_range);
scalar_tensor full_redux;
scalar_tensor full_redux_gpu;
in.setRandom();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
const IndexType offset = 64;
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
full_redux = in_offset.sum();
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data =
static_cast<DataType*>(sycl_device.allocate(sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<scalar_tensor> out_gpu(gpu_out_data);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.sum();
sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data,
sizeof(DataType));
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
// 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_max_sycl(
const Eigen::SyclDevice& sycl_device) {
const IndexType num_rows = 4096;
const IndexType num_cols = 4096;
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
array<IndexType, 2> tensorRange = {{num_rows, num_cols}};
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
Tensor<DataType, 2, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 0, DataLayout, IndexType> full_redux;
Tensor<DataType, 0, DataLayout, IndexType> full_redux_gpu;
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
in.setRandom();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
full_redux = in.maximum();
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
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));
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
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.maximum();
sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data,
sizeof(DataType));
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_max_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using scalar_tensor = Tensor<DataType, 0, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
const IndexType n_elems = internal::array_prod(tensor_range);
data_tensor in(tensor_range);
scalar_tensor full_redux;
scalar_tensor full_redux_gpu;
in.setRandom();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
// Set the initial value to be the max.
// As we don't include this in the reduction the result should not be 2.
in(0) = static_cast<DataType>(2);
const IndexType offset = 64;
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
full_redux = in_offset.maximum();
VERIFY_IS_NOT_EQUAL(full_redux(), in(0));
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data =
static_cast<DataType*>(sycl_device.allocate(sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<scalar_tensor> out_gpu(gpu_out_data);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.maximum();
sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data,
sizeof(DataType));
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
// Check that the CPU and GPU reductions return the same result.
VERIFY_IS_APPROX(full_redux_gpu(), full_redux());
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
template <typename DataType, int DataLayout, typename IndexType>
static void test_full_reductions_mean_sycl(
const Eigen::SyclDevice& sycl_device) {
const IndexType num_rows = 4096;
const IndexType num_cols = 4096;
array<IndexType, 2> tensorRange = {{num_rows, num_cols}};
array<IndexType, 1> argRange = {{num_cols}};
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 0;
// red_axis[1]=1;
Tensor<DataType, 2, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> in_arg1(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> in_arg2(tensorRange);
Tensor<bool, 1, DataLayout, IndexType> out_arg_cpu(argRange);
Tensor<bool, 1, DataLayout, IndexType> out_arg_gpu(argRange);
Tensor<bool, 1, DataLayout, IndexType> out_arg_gpu_helper(argRange);
Tensor<DataType, 0, DataLayout, IndexType> full_redux;
Tensor<DataType, 0, DataLayout, IndexType> full_redux_gpu;
in.setRandom();
in_arg1.setRandom();
in_arg2.setRandom();
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_in_arg1_data = static_cast<DataType*>(sycl_device.allocate(
in_arg1.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_in_arg2_data = static_cast<DataType*>(sycl_device.allocate(
in_arg2.dimensions().TotalSize() * sizeof(DataType)));
bool* gpu_out_arg__gpu_helper_data = static_cast<bool*>(sycl_device.allocate(
out_arg_gpu.dimensions().TotalSize() * sizeof(DataType)));
bool* gpu_out_arg_data = static_cast<bool*>(sycl_device.allocate(
out_arg_gpu.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, 2, DataLayout, IndexType>> in_Arg1_gpu(
gpu_in_arg1_data, tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> in_Arg2_gpu(
gpu_in_arg2_data, tensorRange);
TensorMap<Tensor<bool, 1, DataLayout, IndexType>> out_Argout_gpu(
gpu_out_arg_data, argRange);
TensorMap<Tensor<bool, 1, DataLayout, IndexType>> out_Argout_gpu_helper(
gpu_out_arg__gpu_helper_data, argRange);
TensorMap<Tensor<DataType, 0, DataLayout, IndexType>> out_gpu(gpu_out_data);
// CPU VERSION
out_arg_cpu =
(in_arg1.argmax(1) == in_arg2.argmax(1))
.select(out_arg_cpu.constant(true), out_arg_cpu.constant(false));
full_redux = (out_arg_cpu.template cast<float>())
.reduce(red_axis, Eigen::internal::MeanReducer<DataType>());
// GPU VERSION
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
sycl_device.memcpyHostToDevice(
gpu_in_arg1_data, in_arg1.data(),
(in_arg1.dimensions().TotalSize()) * sizeof(DataType));
sycl_device.memcpyHostToDevice(
gpu_in_arg2_data, in_arg2.data(),
(in_arg2.dimensions().TotalSize()) * sizeof(DataType));
out_Argout_gpu_helper.device(sycl_device) =
(in_Arg1_gpu.argmax(1) == in_Arg2_gpu.argmax(1));
out_Argout_gpu.device(sycl_device) =
(out_Argout_gpu_helper)
.select(out_Argout_gpu.constant(true),
out_Argout_gpu.constant(false));
out_gpu.device(sycl_device) =
(out_Argout_gpu.template cast<float>())
.reduce(red_axis, Eigen::internal::MeanReducer<DataType>());
sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data,
sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
std::cout << "SYCL : " << full_redux_gpu() << " , CPU : " << full_redux()
<< '\n';
VERIFY_IS_EQUAL(full_redux_gpu(), full_redux());
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_in_arg1_data);
sycl_device.deallocate(gpu_in_arg2_data);
sycl_device.deallocate(gpu_out_arg__gpu_helper_data);
sycl_device.deallocate(gpu_out_arg_data);
sycl_device.deallocate(gpu_out_data);
}
template <typename DataType, int DataLayout, typename IndexType>
static void test_full_reductions_mean_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using scalar_tensor = Tensor<DataType, 0, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
const IndexType n_elems = internal::array_prod(tensor_range);
data_tensor in(tensor_range);
scalar_tensor full_redux;
scalar_tensor full_redux_gpu;
in.setRandom();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
const IndexType offset = 64;
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
full_redux = in_offset.mean();
VERIFY_IS_NOT_EQUAL(full_redux(), in(0));
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data =
static_cast<DataType*>(sycl_device.allocate(sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<scalar_tensor> out_gpu(gpu_out_data);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * 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);
}
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
template <typename DataType, int DataLayout, typename IndexType>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
static void test_full_reductions_mean_with_odd_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
// This is a particular case which illustrates a possible problem when the
// number of local threads in a workgroup is even, but is not a power of two.
using data_tensor = Tensor<DataType, 1, DataLayout, IndexType>;
using scalar_tensor = Tensor<DataType, 0, DataLayout, IndexType>;
// 2177 = (17 * 128) + 1 gives rise to 18 local threads.
// 8708 = 4 * 2177 = 4 * (17 * 128) + 4 uses 18 vectorised local threads.
const IndexType n_elems = 8707;
array<IndexType, 1> tensor_range = {{n_elems}};
data_tensor in(tensor_range);
DataType full_redux;
DataType full_redux_gpu;
TensorMap<scalar_tensor> red_cpu(&full_redux);
TensorMap<scalar_tensor> red_gpu(&full_redux_gpu);
const DataType const_val = static_cast<DataType>(0.6391);
in = in.constant(const_val);
Eigen::IndexList<Eigen::type2index<0>> red_axis;
red_cpu = in.reduce(red_axis, Eigen::internal::MeanReducer<DataType>());
VERIFY_IS_APPROX(const_val, red_cpu());
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data =
static_cast<DataType*>(sycl_device.allocate(sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data, tensor_range);
TensorMap<scalar_tensor> out_gpu(gpu_out_data);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
out_gpu.device(sycl_device) =
in_gpu.reduce(red_axis, Eigen::internal::MeanReducer<DataType>());
sycl_device.memcpyDeviceToHost(red_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);
}
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
template <typename DataType, int DataLayout, typename IndexType>
static void test_full_reductions_min_sycl(
const Eigen::SyclDevice& sycl_device) {
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
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();
Fixing TensorReductionSycl for min and max.
2017-02-07 02:05:23 +08:00
full_redux = in.minimum();
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
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));
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 0, DataLayout, IndexType>> out_gpu(gpu_out_data);
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
Fixing TensorReductionSycl for min and max.
2017-02-07 02:05:23 +08:00
out_gpu.device(sycl_device) = in_gpu.minimum();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyDeviceToHost(full_redux_gpu.data(), gpu_out_data,
sizeof(DataType));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
// 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);
}
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
template <typename DataType, int DataLayout, typename IndexType>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
static void test_full_reductions_min_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using scalar_tensor = Tensor<DataType, 0, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
const IndexType n_elems = internal::array_prod(tensor_range);
data_tensor in(tensor_range);
scalar_tensor full_redux;
scalar_tensor full_redux_gpu;
in.setRandom();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
// Set the initial value to be the min.
// As we don't include this in the reduction the result should not be -2.
in(0) = static_cast<DataType>(-2);
const IndexType offset = 64;
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
full_redux = in_offset.minimum();
VERIFY_IS_NOT_EQUAL(full_redux(), in(0));
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data =
static_cast<DataType*>(sycl_device.allocate(sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<scalar_tensor> out_gpu(gpu_out_data);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.minimum();
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());
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
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_max_sycl(
const Eigen::SyclDevice& sycl_device) {
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
IndexType dim_x = 145;
IndexType dim_y = 1;
IndexType dim_z = 67;
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
array<IndexType, 3> tensorRange = {{dim_x, dim_y, dim_z}};
Eigen::array<IndexType, 1> red_axis;
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
red_axis[0] = 0;
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
array<IndexType, 2> reduced_tensorRange = {{dim_y, dim_z}};
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
Tensor<DataType, 3, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux(reduced_tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux_gpu(reduced_tensorRange);
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
in.setRandom();
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
redux = in.maximum(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu.dimensions().TotalSize() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> out_gpu(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.maximum(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu.data(), gpu_out_data,
redux_gpu.dimensions().TotalSize() * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType j = 0; j < reduced_tensorRange[0]; j++)
for (IndexType k = 0; k < reduced_tensorRange[1]; k++)
VERIFY_IS_APPROX(redux_gpu(j, k), redux(j, k));
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_max_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using reduced_tensor = Tensor<DataType, 1, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
array<IndexType, 1> reduced_range = {{num_cols}};
const IndexType n_elems = internal::array_prod(tensor_range);
const IndexType n_reduced = num_cols;
data_tensor in(tensor_range);
reduced_tensor redux;
reduced_tensor redux_gpu(reduced_range);
in.setRandom();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
// Set maximum value outside of the considered range.
for (IndexType i = 0; i < n_reduced; i++) {
in(i) = static_cast<DataType>(2);
}
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 0;
const IndexType offset = 64;
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
redux = in_offset.maximum(red_axis);
for (IndexType i = 0; i < n_reduced; i++) {
VERIFY_IS_NOT_EQUAL(redux(i), in(i));
}
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(
sycl_device.allocate(n_reduced * sizeof(DataType)));
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<reduced_tensor> out_gpu(gpu_out_data, reduced_range);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.maximum(red_axis);
sycl_device.memcpyDeviceToHost(redux_gpu.data(), gpu_out_data,
n_reduced * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType i = 0; i < n_reduced; i++) {
VERIFY_IS_APPROX(redux_gpu(i), redux(i));
}
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
template <typename DataType, int DataLayout, typename IndexType>
static void test_last_dim_reductions_max_with_offset_sycl(
const Eigen::SyclDevice& sycl_device) {
using data_tensor = Tensor<DataType, 2, DataLayout, IndexType>;
using reduced_tensor = Tensor<DataType, 1, DataLayout, IndexType>;
const IndexType num_rows = 64;
const IndexType num_cols = 64;
array<IndexType, 2> tensor_range = {{num_rows, num_cols}};
array<IndexType, 1> full_reduced_range = {{num_rows}};
array<IndexType, 1> reduced_range = {{num_rows - 1}};
const IndexType n_elems = internal::array_prod(tensor_range);
const IndexType n_reduced = reduced_range[0];
data_tensor in(tensor_range);
reduced_tensor redux(full_reduced_range);
reduced_tensor redux_gpu(reduced_range);
in.setRandom();
redux.setZero();
array<IndexType, 2> tensor_offset_range(tensor_range);
tensor_offset_range[0] -= 1;
// Set maximum value outside of the considered range.
for (IndexType i = 0; i < n_reduced; i++) {
in(i) = static_cast<DataType>(2);
}
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 1;
const IndexType offset = 64;
// Introduce an offset in both the input and the output.
TensorMap<data_tensor> in_offset(in.data() + offset, tensor_offset_range);
TensorMap<reduced_tensor> red_offset(redux.data() + 1, reduced_range);
red_offset = in_offset.maximum(red_axis);
// Check that the first value hasn't been changed and that the reduced values
// are not equal to the previously set maximum in the input outside the range.
VERIFY_IS_EQUAL(redux(0), static_cast<DataType>(0));
for (IndexType i = 0; i < n_reduced; i++) {
VERIFY_IS_NOT_EQUAL(red_offset(i), in(i));
}
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(n_elems * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(
sycl_device.allocate((n_reduced + 1) * sizeof(DataType)));
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
TensorMap<data_tensor> in_gpu(gpu_in_data + offset, tensor_offset_range);
TensorMap<reduced_tensor> out_gpu(gpu_out_data + 1, reduced_range);
sycl_device.memcpyHostToDevice(gpu_in_data, in.data(),
n_elems * sizeof(DataType));
Fixing TensorReductionSycl for min and max.
2017-02-07 02:05:23 +08:00
out_gpu.device(sycl_device) = in_gpu.maximum(red_axis);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyDeviceToHost(redux_gpu.data(), out_gpu.data(),
n_reduced * sizeof(DataType));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
// Check that the CPU and GPU reductions return the same result.
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
for (IndexType i = 0; i < n_reduced; i++) {
VERIFY_IS_APPROX(redux_gpu(i), red_offset(i));
}
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_sum_sycl(
const Eigen::SyclDevice& sycl_device, IndexType dim_x, IndexType dim_y) {
array<IndexType, 2> tensorRange = {{dim_x, dim_y}};
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 0;
array<IndexType, 1> reduced_tensorRange = {{dim_y}};
Tensor<DataType, 2, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 1, DataLayout, IndexType> redux(reduced_tensorRange);
Tensor<DataType, 1, DataLayout, IndexType> redux_gpu(reduced_tensorRange);
in.setRandom();
redux = in.sum(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu.dimensions().TotalSize() * sizeof(DataType)));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 1, DataLayout, IndexType>> out_gpu(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.sum(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu.data(), gpu_out_data,
redux_gpu.dimensions().TotalSize() * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType i = 0; i < redux.size(); i++) {
VERIFY_IS_APPROX(redux_gpu.data()[i], redux.data()[i]);
}
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
template <typename DataType, int DataLayout, typename IndexType>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
static void test_first_dim_reductions_mean_sycl(
const Eigen::SyclDevice& sycl_device) {
IndexType dim_x = 145;
IndexType dim_y = 1;
IndexType dim_z = 67;
array<IndexType, 3> tensorRange = {{dim_x, dim_y, dim_z}};
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 0;
array<IndexType, 2> reduced_tensorRange = {{dim_y, dim_z}};
Tensor<DataType, 3, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux(reduced_tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux_gpu(reduced_tensorRange);
in.setRandom();
redux = in.mean(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu.dimensions().TotalSize() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> out_gpu(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.mean(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu.data(), gpu_out_data,
redux_gpu.dimensions().TotalSize() * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType j = 0; j < reduced_tensorRange[0]; j++)
for (IndexType k = 0; k < reduced_tensorRange[1]; k++)
VERIFY_IS_APPROX(redux_gpu(j, k), redux(j, k));
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
template <typename DataType, int DataLayout, typename IndexType>
static void test_last_dim_reductions_mean_sycl(
const Eigen::SyclDevice& sycl_device) {
IndexType dim_x = 64;
IndexType dim_y = 1;
IndexType dim_z = 32;
array<IndexType, 3> tensorRange = {{dim_x, dim_y, dim_z}};
Eigen::array<IndexType, 1> red_axis;
red_axis[0] = 2;
array<IndexType, 2> reduced_tensorRange = {{dim_x, dim_y}};
Tensor<DataType, 3, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux(reduced_tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux_gpu(reduced_tensorRange);
in.setRandom();
redux = in.mean(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu.dimensions().TotalSize() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> out_gpu(
gpu_out_data, reduced_tensorRange);
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
out_gpu.device(sycl_device) = in_gpu.mean(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu.data(), gpu_out_data,
redux_gpu.dimensions().TotalSize() * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType j = 0; j < reduced_tensorRange[0]; j++)
for (IndexType k = 0; k < reduced_tensorRange[1]; k++)
VERIFY_IS_APPROX(redux_gpu(j, k), redux(j, k));
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
template <typename DataType, int DataLayout, typename IndexType>
static void test_last_dim_reductions_sum_sycl(
const Eigen::SyclDevice& sycl_device) {
IndexType dim_x = 64;
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
IndexType dim_y = 1;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
IndexType dim_z = 32;
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
array<IndexType, 3> tensorRange = {{dim_x, dim_y, dim_z}};
Eigen::array<IndexType, 1> red_axis;
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
red_axis[0] = 2;
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
array<IndexType, 2> reduced_tensorRange = {{dim_x, dim_y}};
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
Reducing warnings in Sycl backend.
2017-02-01 23:29:53 +08:00
Tensor<DataType, 3, DataLayout, IndexType> in(tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux(reduced_tensorRange);
Tensor<DataType, 2, DataLayout, IndexType> redux_gpu(reduced_tensorRange);
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
in.setRandom();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
redux = in.sum(red_axis);
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu.dimensions().TotalSize() * sizeof(DataType)));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
TensorMap<Tensor<DataType, 3, DataLayout, IndexType>> in_gpu(gpu_in_data,
tensorRange);
TensorMap<Tensor<DataType, 2, DataLayout, IndexType>> out_gpu(
gpu_out_data, reduced_tensorRange);
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyHostToDevice(
gpu_in_data, in.data(), (in.dimensions().TotalSize()) * sizeof(DataType));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
out_gpu.device(sycl_device) = in_gpu.sum(red_axis);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
sycl_device.memcpyDeviceToHost(
redux_gpu.data(), gpu_out_data,
redux_gpu.dimensions().TotalSize() * sizeof(DataType));
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
// Check that the CPU and GPU reductions return the same result.
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
for (IndexType j = 0; j < reduced_tensorRange[0]; j++)
for (IndexType k = 0; k < reduced_tensorRange[1]; k++)
VERIFY_IS_APPROX(redux_gpu(j, k), redux(j, k));
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
}
Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.
2016-11-09 01:08:02 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
template <typename DataType, int DataLayout, typename IndexType>
static void test_last_reductions_sum_sycl(
const Eigen::SyclDevice& sycl_device) {
auto tensorRange = Sizes<64, 32>(64, 32);
// auto red_axis = Sizes<0,1>(0,1);
Eigen::IndexList<Eigen::type2index<1>> red_axis;
auto reduced_tensorRange = Sizes<64>(64);
TensorFixedSize<DataType, Sizes<64, 32>, DataLayout> in_fix;
TensorFixedSize<DataType, Sizes<64>, DataLayout> redux_fix;
TensorFixedSize<DataType, Sizes<64>, DataLayout> redux_gpu_fix;
in_fix.setRandom();
redux_fix = in_fix.sum(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in_fix.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu_fix.dimensions().TotalSize() * sizeof(DataType)));
TensorMap<TensorFixedSize<DataType, Sizes<64, 32>, DataLayout>> in_gpu_fix(
gpu_in_data, tensorRange);
TensorMap<TensorFixedSize<DataType, Sizes<64>, DataLayout>> out_gpu_fix(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in_fix.data(),
(in_fix.dimensions().TotalSize()) * sizeof(DataType));
out_gpu_fix.device(sycl_device) = in_gpu_fix.sum(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu_fix.data(), gpu_out_data,
redux_gpu_fix.dimensions().TotalSize() * sizeof(DataType));
// Check that the CPU and GPU reductions return the same result.
for (IndexType j = 0; j < reduced_tensorRange[0]; j++) {
VERIFY_IS_APPROX(redux_gpu_fix(j), redux_fix(j));
}
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
}
Only runs the cxx11_tensor_reduction_sycl on devices that are available.
2016-11-19 08:31:14 +08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
template <typename DataType, int DataLayout, typename IndexType>
static void test_last_reductions_mean_sycl(
const Eigen::SyclDevice& sycl_device) {
auto tensorRange = Sizes<64, 32>(64, 32);
Eigen::IndexList<Eigen::type2index<1>> red_axis;
auto reduced_tensorRange = Sizes<64>(64);
TensorFixedSize<DataType, Sizes<64, 32>, DataLayout> in_fix;
TensorFixedSize<DataType, Sizes<64>, DataLayout> redux_fix;
TensorFixedSize<DataType, Sizes<64>, DataLayout> redux_gpu_fix;
in_fix.setRandom();
redux_fix = in_fix.mean(red_axis);
DataType* gpu_in_data = static_cast<DataType*>(
sycl_device.allocate(in_fix.dimensions().TotalSize() * sizeof(DataType)));
DataType* gpu_out_data = static_cast<DataType*>(sycl_device.allocate(
redux_gpu_fix.dimensions().TotalSize() * sizeof(DataType)));
TensorMap<TensorFixedSize<DataType, Sizes<64, 32>, DataLayout>> in_gpu_fix(
gpu_in_data, tensorRange);
TensorMap<TensorFixedSize<DataType, Sizes<64>, DataLayout>> out_gpu_fix(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(
gpu_in_data, in_fix.data(),
(in_fix.dimensions().TotalSize()) * sizeof(DataType));
out_gpu_fix.device(sycl_device) = in_gpu_fix.mean(red_axis);
sycl_device.memcpyDeviceToHost(
redux_gpu_fix.data(), gpu_out_data,
redux_gpu_fix.dimensions().TotalSize() * sizeof(DataType));
sycl_device.synchronize();
// Check that the CPU and GPU reductions return the same result.
for (IndexType j = 0; j < reduced_tensorRange[0]; j++) {
VERIFY_IS_APPROX(redux_gpu_fix(j), redux_fix(j));
}
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
// SYCL supports a generic case of reduction where the accumulator is a
// different type than the input data This is an example on how to get if a
// Tensor contains nan and/or inf in one reduction
template <typename InT, typename OutT>
struct CustomReducer {
static const bool PacketAccess = false;
static const bool IsStateful = false;
static constexpr OutT InfBit = 1;
static constexpr OutT NanBit = 2;
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const InT x,
OutT* accum) const {
if (Eigen::numext::isinf(x))
*accum |= InfBit;
else if (Eigen::numext::isnan(x))
*accum |= NanBit;
}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const OutT x,
OutT* accum) const {
*accum |= x;
}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE OutT initialize() const {
return OutT(0);
}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE OutT finalize(const OutT accum) const {
return accum;
}
};
template <typename DataType, typename AccumType, int DataLayout,
typename IndexType>
static void test_full_reductions_custom_sycl(
const Eigen::SyclDevice& sycl_device) {
constexpr IndexType InSize = 64;
auto tensorRange = Sizes<InSize>(InSize);
Eigen::IndexList<Eigen::type2index<0>> dims;
auto reduced_tensorRange = Sizes<>();
TensorFixedSize<DataType, Sizes<InSize>, DataLayout> in_fix;
TensorFixedSize<AccumType, Sizes<>, DataLayout> redux_gpu_fix;
CustomReducer<DataType, AccumType> reducer;
in_fix.setRandom();
size_t in_size_bytes = in_fix.dimensions().TotalSize() * sizeof(DataType);
DataType* gpu_in_data =
static_cast<DataType*>(sycl_device.allocate(in_size_bytes));
AccumType* gpu_out_data =
static_cast<AccumType*>(sycl_device.allocate(sizeof(AccumType)));
TensorMap<TensorFixedSize<DataType, Sizes<InSize>, DataLayout>> in_gpu_fix(
gpu_in_data, tensorRange);
TensorMap<TensorFixedSize<AccumType, Sizes<>, DataLayout>> out_gpu_fix(
gpu_out_data, reduced_tensorRange);
sycl_device.memcpyHostToDevice(gpu_in_data, in_fix.data(), in_size_bytes);
out_gpu_fix.device(sycl_device) = in_gpu_fix.reduce(dims, reducer);
sycl_device.memcpyDeviceToHost(redux_gpu_fix.data(), gpu_out_data,
sizeof(AccumType));
VERIFY_IS_EQUAL(redux_gpu_fix(0), AccumType(0));
sycl_device.deallocate(gpu_in_data);
sycl_device.deallocate(gpu_out_data);
}
template <typename DataType, typename Dev>
void sycl_reduction_test_full_per_device(const Dev& sycl_device) {
test_full_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device);
test_full_reductions_sum_sycl<DataType, ColMajor, int64_t>(sycl_device);
test_full_reductions_min_sycl<DataType, ColMajor, int64_t>(sycl_device);
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
test_full_reductions_min_sycl<DataType, RowMajor, int64_t>(sycl_device);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
test_full_reductions_max_sycl<DataType, ColMajor, int64_t>(sycl_device);
test_full_reductions_max_sycl<DataType, RowMajor, int64_t>(sycl_device);
test_full_reductions_mean_sycl<DataType, ColMajor, int64_t>(sycl_device);
test_full_reductions_mean_sycl<DataType, RowMajor, int64_t>(sycl_device);
test_full_reductions_custom_sycl<DataType, int, RowMajor, int64_t>(
sycl_device);
test_full_reductions_custom_sycl<DataType, int, ColMajor, int64_t>(
sycl_device);
sycl_device.synchronize();
}
template <typename DataType, typename Dev>
void sycl_reduction_full_offset_per_device(const Dev& sycl_device) {
test_full_reductions_sum_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
test_full_reductions_sum_with_offset_sycl<DataType, ColMajor, int64_t>(
sycl_device);
test_full_reductions_min_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
test_full_reductions_min_with_offset_sycl<DataType, ColMajor, int64_t>(
sycl_device);
test_full_reductions_max_with_offset_sycl<DataType, ColMajor, int64_t>(
sycl_device);
test_full_reductions_max_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
test_full_reductions_mean_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
test_full_reductions_mean_with_offset_sycl<DataType, ColMajor, int64_t>(
sycl_device);
test_full_reductions_mean_with_odd_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
sycl_device.synchronize();
}
template <typename DataType, typename Dev>
void sycl_reduction_test_first_dim_per_device(const Dev& sycl_device) {
test_first_dim_reductions_sum_sycl<DataType, ColMajor, int64_t>(sycl_device,
4197, 4097);
test_first_dim_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device,
4197, 4097);
test_first_dim_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device,
129, 8);
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
test_first_dim_reductions_max_sycl<DataType, RowMajor, int64_t>(sycl_device);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
test_first_dim_reductions_max_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
sycl_device.synchronize();
}
template <typename DataType, typename Dev>
void sycl_reduction_test_last_dim_per_device(const Dev& sycl_device) {
Adding mean to TensorReductionSycl.h
2017-02-07 23:43:17 +08:00
test_last_dim_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
test_last_dim_reductions_max_with_offset_sycl<DataType, RowMajor, int64_t>(
sycl_device);
test_last_reductions_sum_sycl<DataType, ColMajor, int64_t>(sycl_device);
test_last_reductions_sum_sycl<DataType, RowMajor, int64_t>(sycl_device);
test_last_reductions_mean_sycl<DataType, ColMajor, int64_t>(sycl_device);
test_last_reductions_mean_sycl<DataType, RowMajor, int64_t>(sycl_device);
sycl_device.synchronize();
Modifying TensorDeviceSycl.h to always create buffer of type uint8_t and convert them to the actual type at the execution on the device; adding the queue interface class to separate the lifespan of sycl queue and buffers,created for that queue, from Eigen::SyclDevice; modifying sycl tests to support the evaluation of the results for both row major and column major data layout on all different devices that are supported by Sycl{CPU; GPU; and Host}.
2016-11-19 00:20:42 +08:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
Get rid of EIGEN_TEST_FUNC, unit tests must now be declared with EIGEN_DECLARE_TEST(mytest) { /* code */ }. This provide several advantages: - more flexibility in designing unit tests - unit tests can be glued to speed up compilation - unit tests are compiled with same predefined macros, which is a requirement for zapcc
2018-07-17 20:46:15 +08:00
EIGEN_DECLARE_TEST(cxx11_tensor_reduction_sycl) {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 18:08:54 +08:00
for (const auto& device : Eigen::get_sycl_supported_devices()) {
std::cout << "Running on "
<< device.template get_info<cl::sycl::info::device::name>()
<< std::endl;
QueueInterface queueInterface(device);
auto sycl_device = Eigen::SyclDevice(&queueInterface);
CALL_SUBTEST_1(sycl_reduction_test_full_per_device<float>(sycl_device));
CALL_SUBTEST_2(sycl_reduction_full_offset_per_device<float>(sycl_device));
CALL_SUBTEST_3(
sycl_reduction_test_first_dim_per_device<float>(sycl_device));
CALL_SUBTEST_4(sycl_reduction_test_last_dim_per_device<float>(sycl_device));
Only runs the cxx11_tensor_reduction_sycl on devices that are available.
2016-11-19 08:31:14 +08:00
}
Removed the sycl include from Eigen/Core and moved it to Unsupported/Eigen/CXX11/Tensor; added TensorReduction for sycl (full reduction and partial reduction); added TensorReduction test case for sycl (full reduction and partial reduction); fixed the tile size on TensorSyclRun.h based on the device max work group size;
2016-11-05 02:18:19 +08:00
}
Reference in New Issue Copy Permalink