diff options
author | Mehdi Goli <mehdi.goli@codeplay.com> | 2019-11-28 10:08:54 +0000 |
---|---|---|
committer | Mehdi Goli <mehdi.goli@codeplay.com> | 2019-11-28 10:08:54 +0000 |
commit | 00f32752f7d0b193c6788691c3cf0b76457a044d (patch) | |
tree | 792e46110f0751ea8802fa9d403d1472d5977ac3 /unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h | |
parent | ea51a9eace7e4f0ea839e61eb2df85ccfb94aee8 (diff) |
[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
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h')
-rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h | 120 |
1 files changed, 0 insertions, 120 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h deleted file mode 100644 index 7b8bd2df7..000000000 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h +++ /dev/null @@ -1,120 +0,0 @@ -// This file is part of Eigen, a lightweight C++ template library -// for linear algebra. -// -// 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/. - -// General include header of SYCL target for Tensor Module -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H - -#ifdef EIGEN_USE_SYCL - -// global pointer to set different attribute state for a class -template <class T> -struct MakeGlobalPointer { - typedef typename cl::sycl::global_ptr<T>::pointer_t Type; - typedef typename cl::sycl::global_ptr<T>::reference_t RefType; -}; - -// global pointer to set different attribute state for a class -template <class T> -struct MakeLocalPointer { - typedef typename cl::sycl::local_ptr<T>::pointer_t Type; - typedef typename cl::sycl::local_ptr<T>::reference_t RefType; -}; - - -namespace Eigen { - template<typename StrideDims, typename XprType> class TensorTupleReducerDeviceOp; - template<typename StrideDims, typename ArgType> struct TensorEvaluator<const TensorTupleReducerDeviceOp<StrideDims, ArgType>, SyclKernelDevice>; -namespace internal { - -#ifdef __SYCL_DEVICE_ONLY__ -template<typename A, typename B> struct TypeConversion { - template<typename T> - static typename MakeGlobalPointer<A>::Type get_address_space_pointer(typename MakeGlobalPointer<T>::Type p); - template<typename T> - static typename MakeLocalPointer<A>::Type get_address_space_pointer(typename MakeLocalPointer<T>::Type p); - - template<typename T> - static A* get_address_space_pointer(T* p); - typedef decltype(get_address_space_pointer(B())) type; -}; - -#endif -} -namespace TensorSycl { -namespace internal { - - template<typename CoeffReturnType, typename OP, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer; -/// This struct is used for special expression nodes with no operations (for example assign and selectOP). - struct NoOP; - -template<bool IsConst, typename T> struct GetType{ - typedef const T Type; -}; -template<typename T> struct GetType<false, T>{ - typedef T Type; -}; - -template <bool Conds, size_t X , size_t Y > struct ValueCondition { - static constexpr size_t Res =X; -}; -template<size_t X, size_t Y> struct ValueCondition<false, X, Y> { - static constexpr size_t Res =Y; -}; - -} -} -} - -// tuple construction -#include "TensorSyclTuple.h" - -// counting number of leaf at compile time -#include "TensorSyclLeafCount.h" - -// The index PlaceHolder takes the actual expression and replaces the actual -// data on it with the place holder. It uses the same pre-order expression tree -// traverse as the leaf count in order to give the right access number to each -// node in the expression -#include "TensorSyclPlaceHolderExpr.h" - -// creation of an accessor tuple from a tuple of SYCL buffers -#include "TensorSyclExtractAccessor.h" - -// this is used to change the address space type in tensor map for GPU -#include "TensorSyclConvertToDeviceExpression.h" - -// this is used to extract the functors -#include "TensorSyclExtractFunctors.h" - -// this is used to create tensormap on the device -// this is used to construct the expression on the device -#include "TensorSyclExprConstructor.h" - -/// this is used for extracting tensor reduction -#include "TensorReductionSycl.h" - -// TensorArgMaxSycl.h -#include "TensorArgMaxSycl.h" - -/// this is used for extracting tensor convolution -#include "TensorConvolutionSycl.h" - -// kernel execution using fusion -#include "TensorSyclRun.h" -//sycl functors -#include "TensorSyclFunctors.h" - -#include "TensorContractionSycl.h" - -#endif // end of EIGEN_USE_SYCL -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H |