eigen/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h

// 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