Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable > Struct Template Reference

GeneralScalarContraction is a template class that provides the scalar value of Tensor -Tensor contraction operation, when all the dimensions are contracting dimensions. This Kernel reduces two tensors to an scalar. More...

#include <TensorContractionSycl.h>

Public Types

typedef cl::sycl::accessor< OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local > Scratch
 

Public Member Functions

EIGEN_DEVICE_FUNC GeneralScalarContraction (Scratch scratch_, const LhsMapper lhs_, const RhsMapper rhs_, OutAccessor out_res_, const StorageIndex rng_)
 
EIGEN_DEVICE_FUNC void operator() (cl::sycl::nd_item< 1 > itemID) const
 

Public Attributes

Scratch scratch
 
const LhsMapper lhs
 
const RhsMapper rhs
 
OutAccessor out_res
 
const StorageIndex rng
 

Detailed Description

template<typename OutScalar, typename LhsScalar, typename RhsScalar, typename OutAccessor, typename LhsMapper, typename RhsMapper, typename StorageIndex, bool Vectorizable>
struct Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >

GeneralScalarContraction is a template class that provides the scalar value of Tensor -Tensor contraction operation, when all the dimensions are contracting dimensions. This Kernel reduces two tensors to an scalar.

Template Parameters
OutScalardetermines the output scalar type
LhsScalardetermines the left-hand-side scalar type
RhsScalardetermines the right-hand-side scalar type
OutAccessordetermines the sycl accessor type for out put (please see the sycl-1.2.1 specification (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for accessor definition)
LhsMapperdetermines the tensor contraction mapper type for left-hand-side matrix
RhsMapperdetermines the tensor contraction mapper type for right-hand-side matrix
StorageIndexdetermines the StorageIndex Type
Vectorizabledetermines whether or not the vectorization is enabled for the Eigen expression.
Parameters
scratchlocal memory containing tiles of LHS and RHS tensors for each work-group
lhsdetermines the left-hand-side flattened tensor (tensor mapper)
rhsdetermines the right-hand-side flattened tensor (tensor mapper)
out_resdetermines the output tensor containing the contraction result
rngdetermines the total input data size

Member Typedef Documentation

◆ Scratch

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::Scratch

Constructor & Destructor Documentation

◆ GeneralScalarContraction()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
EIGEN_DEVICE_FUNC Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::GeneralScalarContraction ( Scratch  scratch_,
const LhsMapper  lhs_,
const RhsMapper  rhs_,
OutAccessor  out_res_,
const StorageIndex  rng_ 
)
inline
1245  : scratch(scratch_), lhs(lhs_), rhs(rhs_), out_res(out_res_), rng(rng_) {}
OutAccessor out_res
Definition: TensorContractionSycl.h:1240
Scratch scratch
Definition: TensorContractionSycl.h:1237
const RhsMapper rhs
Definition: TensorContractionSycl.h:1239
const StorageIndex rng
Definition: TensorContractionSycl.h:1241
const LhsMapper lhs
Definition: TensorContractionSycl.h:1238

Member Function Documentation

◆ operator()()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
EIGEN_DEVICE_FUNC void Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::operator() ( cl::sycl::nd_item< 1 >  itemID) const
inline
1247  {
1248  auto out_ptr = out_res;
1249  OutScalar *scratch_ptr = scratch.get_pointer();
1250 
1251  StorageIndex globalid = itemID.get_global_id(0);
1252  StorageIndex localid = itemID.get_local_id(0);
1253  OutScalar accumulator = OutScalar(0);
1254  for (StorageIndex i = globalid; i < rng; i += itemID.get_global_range(0)) {
1255  accumulator = Eigen::internal::pmadd(lhs(0, i), rhs(i, 0), accumulator);
1256  }
1257  auto out_scratch_ptr = scratch_ptr + localid;
1258  *out_scratch_ptr = accumulator;
1259  for (StorageIndex offset = itemID.get_local_range(0) >> 1; offset > 0; offset >>= 1) {
1260  itemID.barrier(cl::sycl::access::fence_space::local_space);
1261  if (localid < offset) {
1262  *out_scratch_ptr = (accumulator += out_scratch_ptr[offset]);
1263  }
1264  }
1265  if (localid == 0) {
1266  out_ptr[itemID.get_group(0)] = accumulator;
1267  }
1268  }
int i
Definition: BiCGSTAB_step_by_step.cpp:9
EIGEN_STRONG_INLINE Packet4f pmadd(const Packet4f &a, const Packet4f &b, const Packet4f &c)
Definition: AltiVec/PacketMath.h:1218

References i, Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::lhs, Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::out_res, Eigen::internal::pmadd(), Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::rhs, Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::rng, and Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::scratch.

Member Data Documentation

◆ lhs

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
const LhsMapper Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::lhs

◆ out_res

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
OutAccessor Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::out_res

◆ rhs

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
const RhsMapper Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::rhs

◆ rng

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
const StorageIndex Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::rng

◆ scratch

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , bool Vectorizable>
Scratch Eigen::TensorSycl::internal::GeneralScalarContraction< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Vectorizable >::scratch

The documentation for this struct was generated from the following file: