devtools/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2026-04-10 11:34:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

Replace instances of EIGEN_CONSTEXPR macro

Browse Source
This commit is contained in:
Tyler Veness
2024-11-15 21:11:01 -08:00
parent 33f5f59614
commit d6689a15d7
78 changed files with 445 additions and 501 deletions
Download Patch File Download Diff File Expand all files Collapse all files

107
unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h
View File

@@ -45,19 +45,19 @@ namespace internal {
template <typename Scalar, typename StorageIndex, StorageIndex NCWindow, StorageIndex CFactor, StorageIndex NCFactor>
struct TVPanelSize {
// LocalThreadSizeC: determines total number of thread per workgroup for the contracting dimension
static EIGEN_CONSTEXPR StorageIndex LocalThreadSizeC = EIGEN_SYCL_LOCAL_THREAD_DIM0;
static constexpr StorageIndex LocalThreadSizeC = EIGEN_SYCL_LOCAL_THREAD_DIM0;
// LocalThreadSizeNC: determines total number of thread per workgroup for the non-contracting dimension
static EIGEN_CONSTEXPR StorageIndex LocalThreadSizeNC = EIGEN_SYCL_LOCAL_THREAD_DIM1;
static constxpr StorageIndex LocalThreadSizeNC = EIGEN_SYCL_LOCAL_THREAD_DIM1;
// TileSizeDimNC: determines the tile size for the non-contracting dimension
static EIGEN_CONSTEXPR StorageIndex TileSizeDimNC = NCWindow / NCFactor;
static constexpr StorageIndex TileSizeDimNC = NCWindow / NCFactor;
// TileSizeDimC: determines the tile size for the contracting dimension
static EIGEN_CONSTEXPR StorageIndex TileSizeDimC = CFactor * LocalThreadSizeNC * LocalThreadSizeC;
static constexpr StorageIndex TileSizeDimC = CFactor * LocalThreadSizeNC * LocalThreadSizeC;
// WorkLoadPerThreadNC : determines workload per thread for loading the non-contracting dimension
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadNC = TileSizeDimNC / LocalThreadSizeNC;
static constexpr StorageIndex WorkLoadPerThreadNC = TileSizeDimNC / LocalThreadSizeNC;
// WorkLoadPerThreadC: determines workload per thread for loading the non-contracting dimension
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadC = TileSizeDimC / LocalThreadSizeC;
static constexpr StorageIndex WorkLoadPerThreadC = TileSizeDimC / LocalThreadSizeC;
// BC : determines if supporting bank conflict is required
static EIGEN_CONSTEXPR bool BC = false;
static constexpr bool BC = false;
};
#endif
@@ -81,40 +81,40 @@ struct TVPanelSize {
template <typename Scalar, typename StorageIndex, StorageIndex REG_SIZE_M, StorageIndex REG_SIZE_N, StorageIndex TSDK>
struct TTPanelSize {
// TileSizeDimK: determines Tile size for dimension K. The packet size is assumed to be considered
static EIGEN_CONSTEXPR StorageIndex TileSizeDimK = TSDK;
static constexpr StorageIndex TileSizeDimK = TSDK;
// WorkLoadPerThreadM : determines workload per thread for loading the M dimension This can be varied based on the
// available register on a chosen device(can be controlled by EIGEN_SYCL_REG_M macro//
#ifndef EIGEN_SYCL_REG_M
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadM = REG_SIZE_M;
static constexpr StorageIndex WorkLoadPerThreadM = REG_SIZE_M;
#else
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadM = EIGEN_SYCL_REG_M;
static constexpr StorageIndex WorkLoadPerThreadM = EIGEN_SYCL_REG_M;
#endif
// WorkLoadPerThreadN : determines workload per thread for loading the N dimension This can be varied based on the
// available register on a chosen device(can be controlled by EIGEN_SYCL_REG_N macro
#ifndef EIGEN_SYCL_REG_N
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadN = REG_SIZE_N;
static constexpr StorageIndex WorkLoadPerThreadN = REG_SIZE_N;
#else
static EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadN = EIGEN_SYCL_REG_N;
static constexpr StorageIndex WorkLoadPerThreadN = EIGEN_SYCL_REG_N;
#endif
// LocalThreadSizeM: determines total number of thread per workgroup for the m dimension
static EIGEN_CONSTEXPR StorageIndex LocalThreadSizeM = EIGEN_SYCL_LOCAL_THREAD_DIM0;
static constexpr StorageIndex LocalThreadSizeM = EIGEN_SYCL_LOCAL_THREAD_DIM0;
// LocalThreadSizeN: determines total number of thread per workgroup for the n dimension
static EIGEN_CONSTEXPR StorageIndex LocalThreadSizeN = EIGEN_SYCL_LOCAL_THREAD_DIM1;
static constexpr StorageIndex LocalThreadSizeN = EIGEN_SYCL_LOCAL_THREAD_DIM1;
// TileSizeDimM: determines the tile size for the m dimension
static EIGEN_CONSTEXPR StorageIndex TileSizeDimM = LocalThreadSizeM * WorkLoadPerThreadM;
static constexpr StorageIndex TileSizeDimM = LocalThreadSizeM * WorkLoadPerThreadM;
// TileSizeDimN: determines the tile size for the n dimension
static EIGEN_CONSTEXPR StorageIndex TileSizeDimN = LocalThreadSizeN * WorkLoadPerThreadN;
static constexpr StorageIndex TileSizeDimN = LocalThreadSizeN * WorkLoadPerThreadN;
// LoadPerThreadLhs: determines workload per thread for loading Lhs Tensor. This must be divisible by packetsize
static EIGEN_CONSTEXPR StorageIndex LoadPerThreadLhs =
static constexpr StorageIndex LoadPerThreadLhs =
((TileSizeDimK * WorkLoadPerThreadM * WorkLoadPerThreadN) / (TileSizeDimN));
// LoadPerThreadRhs: determines workload per thread for loading Rhs Tensor. This must be divisible by packetsize
static EIGEN_CONSTEXPR StorageIndex LoadPerThreadRhs =
static constexpr StorageIndex LoadPerThreadRhs =
((TileSizeDimK * WorkLoadPerThreadM * WorkLoadPerThreadN) / (TileSizeDimM));
// BC : determines if supporting bank conflict is required
static EIGEN_CONSTEXPR bool BC = true;
static constexpr bool BC = true;
// DoubleBuffer: determines if double buffering technique should be used (This can be disabled by
// EIGEN_SYCL_DISABLE_DOUBLE_BUFFER macro when the device does not have sufficient local memory)
static EIGEN_CONSTEXPR bool DoubleBuffer =
static constexpr bool DoubleBuffer =
#ifdef EIGEN_SYCL_DISABLE_DOUBLE_BUFFER
false;
#else
@@ -220,7 +220,7 @@ static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!PacketLoad, Packe
template <typename StorageIndex, StorageIndex ld, data_source dt, typename PacketType, typename DataScalar>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<dt != data_source::global_mem, void> write(
PacketType &packet_data, DataScalar ptr) {
EIGEN_CONSTEXPR int PacketSize = Eigen::internal::unpacket_traits<PacketType>::size;
constexpr int PacketSize = Eigen::internal::unpacket_traits<PacketType>::size;
EIGEN_UNROLL_LOOP
for (int i = 0; i < PacketSize; i++) {
*ptr = PacketWrapper<PacketType, PacketSize>::scalarize(i, packet_data);
@@ -320,14 +320,14 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool check_boundary<false>(bool cond) {
*/
template <bool is_transposed, bool is_rhs_, bool packet_load_, typename PacketType>
struct BlockProperties {
static EIGEN_CONSTEXPR bool packet_load = packet_load_;
static constexpr bool packet_load = packet_load_;
typedef typename Eigen::internal::unpacket_traits<PacketType>::type OutScalar;
static EIGEN_CONSTEXPR bool is_rhs = is_rhs_;
static constexpr bool is_rhs = is_rhs_;
typedef std::conditional_t<packet_load, PacketType, OutScalar> OutType;
static EIGEN_CONSTEXPR int elements_per_access = Eigen::internal::unpacket_traits<OutType>::size;
static EIGEN_CONSTEXPR bool is_coalesced_layout = !(is_transposed ^ is_rhs);
static EIGEN_CONSTEXPR int nc_stride = (is_coalesced_layout ? elements_per_access : 1);
static EIGEN_CONSTEXPR int c_stride = (is_coalesced_layout ? 1 : elements_per_access);
static constexpr int elements_per_access = Eigen::internal::unpacket_traits<OutType>::size;
static constexpr bool is_coalesced_layout = !(is_transposed ^ is_rhs);
static constexpr int nc_stride = (is_coalesced_layout ? elements_per_access : 1);
static constexpr int c_stride = (is_coalesced_layout ? 1 : elements_per_access);
};
/*!
@@ -458,11 +458,11 @@ class TensorContractionKernel {
public:
typedef typename Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketReturnType
PacketReturnType;
static EIGEN_CONSTEXPR int PacketSize =
static constexpr int PacketSize =
Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketSize;
static EIGEN_CONSTEXPR bool is_lhs_transposed =
static constexpr bool is_lhs_transposed =
!::Eigen::internal::TensorContractionInputMapperTrait<LhsMapper>::inner_dim_contiguous;
static EIGEN_CONSTEXPR bool is_rhs_transposed =
static constexpr bool is_rhs_transposed =
!::Eigen::internal::TensorContractionInputMapperTrait<RhsMapper>::inner_dim_contiguous;
typedef BlockProperties<is_lhs_transposed, false, input_mapper_properties::is_lhs_matrix && Vectorizable,
@@ -473,20 +473,20 @@ class TensorContractionKernel {
PacketReturnType>
RHSBlockProperties;
static EIGEN_CONSTEXPR StorageIndex NStride =
static constexpr StorageIndex NStride =
contraction_tp == contraction_type::local ? Properties::WorkLoadPerThreadN : RHSBlockProperties::nc_stride;
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Scratch;
typedef cl::sycl::multi_ptr<OutScalar, cl::sycl::access::address_space::local_space> local_ptr;
typedef OutScalar * /*cl::sycl::multi_ptr<OutScalar, cl::sycl::access::address_space::private_space>*/ private_ptr;
typedef std::conditional_t<contraction_tp == contraction_type::local, local_ptr, private_ptr> tile_ptr;
static EIGEN_CONSTEXPR StorageIndex LSDL = contraction_tp == contraction_type::local
? Properties::TileSizeDimM + Properties::BC
: Properties::WorkLoadPerThreadM;
static EIGEN_CONSTEXPR StorageIndex LSDR = contraction_tp == contraction_type::local
? Properties::TileSizeDimN + Properties::BC
: Properties::WorkLoadPerThreadN;
static EIGEN_CONSTEXPR StorageIndex LocalOffset = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN;
static constexpr StorageIndex LSDL = contraction_tp == contraction_type::local
? Properties::TileSizeDimM + Properties::BC
: Properties::WorkLoadPerThreadM;
static constexpr StorageIndex LSDR = contraction_tp == contraction_type::local
? Properties::TileSizeDimN + Properties::BC
: Properties::WorkLoadPerThreadN;
static constexpr StorageIndex LocalOffset = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN;
/**
* \brief MemHolder this is a place holder struct for creating memory hierarchy in SYCL. Inside SYCL kernel it is not
@@ -638,7 +638,7 @@ class TensorContractionKernel {
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void compute_block_per_tile(OutScalar *lhs_block_ptr, OutScalar *rhs_block_ptr,
PacketReturnType *privateRes) const {
StorageIndex idx = 0;
EIGEN_CONSTEXPR StorageIndex lhs_stride =
constexpr StorageIndex lhs_stride =
contraction_tp == contraction_type::local ? (PacketSize * Properties::LocalThreadSizeM) : 1;
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTN = 0; wLPTN < Properties::WorkLoadPerThreadN; wLPTN++) {
@@ -668,8 +668,7 @@ class TensorContractionKernel {
// when local memory is not used M and N are both accessed in a coalesced way. However, when local memory is
// available the k*N is transposed in the local to N*K therefore, each blocks operates on blockId*
// WorkLoadPerThreadN slice of N
EIGEN_CONSTEXPR StorageIndex GlobalNStride =
contraction_tp == contraction_type::local ? 1 : Properties::LocalThreadSizeN;
constexpr StorageIndex GlobalNStride = contraction_tp == contraction_type::local ? 1 : Properties::LocalThreadSizeN;
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTN = 0; wLPTN < Properties::WorkLoadPerThreadN / PrivateNStride; wLPTN++) {
// output leading dimension
@@ -713,9 +712,9 @@ class TensorContractionKernel {
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<contract_tp == contraction_type::no_local> extract_block(
const Input &inpt, PrivateReg private_ptr, const std::pair<StorageIndex, StorageIndex> &,
const StorageIndex &ncOffset, const StorageIndex cOffset) const {
EIGEN_CONSTEXPR StorageIndex LocalThreadSizeNC =
constexpr StorageIndex LocalThreadSizeNC =
InputBlockProperties::is_rhs ? Properties::LocalThreadSizeN : Properties::LocalThreadSizeM;
EIGEN_CONSTEXPR StorageIndex WorkLoadPerThreadNC =
constexpr StorageIndex WorkLoadPerThreadNC =
InputBlockProperties::is_rhs ? Properties::WorkLoadPerThreadN : Properties::WorkLoadPerThreadM;
const StorageIndex &NC = InputBlockProperties::is_rhs ? triple_dim.N : triple_dim.M;
@@ -891,11 +890,11 @@ class TensorContractionKernel {
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<contract_tp == contraction_type::local> extract_block(
const Input &inpt, Local local_ptr, const std::pair<StorageIndex, StorageIndex> &local_index,
const StorageIndex &ncOffset, const StorageIndex cOffset) const {
EIGEN_CONSTEXPR StorageIndex TileSizeDimNC =
constexpr StorageIndex TileSizeDimNC =
InputBlockProperties::is_rhs ? Properties::TileSizeDimN : Properties::TileSizeDimM;
EIGEN_CONSTEXPR StorageIndex LoadPerThread =
constexpr StorageIndex LoadPerThread =
InputBlockProperties::is_rhs ? Properties::LoadPerThreadRhs : Properties::LoadPerThreadLhs;
EIGEN_CONSTEXPR StorageIndex LSD = InputBlockProperties::is_rhs ? LSDR : LSDL;
constexpr StorageIndex LSD = InputBlockProperties::is_rhs ? LSDR : LSDL;
static_assert(((LocalOffset % (TileSizeDimNC / InputBlockProperties::nc_stride) == 0) &&
(LocalOffset % (Properties::TileSizeDimK / InputBlockProperties::c_stride) == 0)),
" LocalOffset must be divisible by stride");
@@ -995,11 +994,11 @@ template <typename OutScalar, typename OutAccessor, typename VectorMapper, typen
struct GeneralVectorTensor {
typedef typename Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketReturnType
PacketReturnType;
static EIGEN_CONSTEXPR int PacketSize =
static constexpr int PacketSize =
Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketSize;
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Scratch;
static EIGEN_CONSTEXPR StorageIndex OutScratchOffset =
static constexpr StorageIndex OutScratchOffset =
KFactor * Properties::LocalThreadSizeC * Properties::LocalThreadSizeNC;
// Since the access layout for a vector can always be coalesced, when LHS is a vector, we pass false and false to make
@@ -1328,8 +1327,8 @@ struct TensorEvaluator<const TensorContractionOp<Indices, LeftArgType, RightArgT
template <bool lhs_inner_dim_contiguous, bool rhs_inner_dim_contiguous, bool rhs_inner_dim_reordered>
struct input_mapper_propertis {
static EIGEN_CONSTEXPR bool is_lhs_matrix = (LDims == 2 && ContractDims == 1) || lhs_inner_dim_contiguous;
static EIGEN_CONSTEXPR bool is_rhs_matrix =
static constexpr bool is_lhs_matrix = (LDims == 2 && ContractDims == 1) || lhs_inner_dim_contiguous;
static constexpr bool is_rhs_matrix =
(RDims == 2 && ContractDims == 1) || (rhs_inner_dim_contiguous && !rhs_inner_dim_reordered);
};
@@ -1537,9 +1536,9 @@ struct TensorEvaluator<const TensorContractionOp<Indices, LeftArgType, RightArgT
void EIGEN_ALWAYS_INLINE LaunchVT(EvaluatorPointerType buffer, const VectorMapper &vec, const TensorMapper &mat,
StorageIndex NC, StorageIndex C) const {
const StorageIndex nonContractDim = NC;
EIGEN_CONSTEXPR StorageIndex NCFactor = 1;
EIGEN_CONSTEXPR StorageIndex CFactor = 1;
EIGEN_CONSTEXPR StorageIndex NCWindow = 16;
constexpr StorageIndex NCFactor = 1;
constexpr StorageIndex CFactor = 1;
constexpr StorageIndex NCWindow = 16;
typedef Eigen::TensorSycl::internal::TVPanelSize<CoeffReturnType, StorageIndex, NCWindow, CFactor, NCFactor>
Properties;
const StorageIndex roundUpC = Eigen::TensorSycl::internal::roundUp(C, Properties::TileSizeDimC);
@@ -1601,7 +1600,7 @@ struct TensorEvaluator<const TensorContractionOp<Indices, LeftArgType, RightArgT
(EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1 - 1)),
"The Local thread size must be a power of 2 for the reduction "
"operation");
EIGEN_CONSTEXPR StorageIndex local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1;
constexpr StorageIndex local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1;
// Here we force the code not to be more than 2-step reduction: Our empirical research shows that if each thread
// reduces at least 512 elementss individually, we get better performance.

26
unsupported/Eigen/CXX11/src/Tensor/TensorIndexList.h
View File

@@ -42,10 +42,10 @@ struct NumTraits<type2index<n>> {
typedef Index Real;
enum { IsComplex = 0, RequireInitialization = false, ReadCost = 1, AddCost = 1, MulCost = 1 };
EIGEN_DEVICE_FUNC static EIGEN_CONSTEXPR EIGEN_STRONG_INLINE Real epsilon() { return 0; }
EIGEN_DEVICE_FUNC static EIGEN_CONSTEXPR EIGEN_STRONG_INLINE Real dummy_precision() { return 0; }
EIGEN_DEVICE_FUNC static EIGEN_CONSTEXPR EIGEN_STRONG_INLINE Real highest() { return n; }
EIGEN_DEVICE_FUNC static EIGEN_CONSTEXPR EIGEN_STRONG_INLINE Real lowest() { return n; }
EIGEN_DEVICE_FUNC static constexpr EIGEN_STRONG_INLINE Real epsilon() { return 0; }
EIGEN_DEVICE_FUNC static constexpr EIGEN_STRONG_INLINE Real dummy_precision() { return 0; }
EIGEN_DEVICE_FUNC static constexpr EIGEN_STRONG_INLINE Real highest() { return n; }
EIGEN_DEVICE_FUNC static constexpr EIGEN_STRONG_INLINE Real lowest() { return n; }
};
namespace internal {
@@ -569,47 +569,47 @@ struct index_pair_second_statically_eq_impl<const IndexPairList<FirstType, Other
namespace Eigen {
namespace internal {
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_known_statically(Index i) {
static EIGEN_DEVICE_FUNC constexpr bool index_known_statically(Index i) {
return index_known_statically_impl<T>::run(i);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool all_indices_known_statically() {
static EIGEN_DEVICE_FUNC constexpr bool all_indices_known_statically() {
return all_indices_known_statically_impl<T>::run();
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool indices_statically_known_to_increase() {
static EIGEN_DEVICE_FUNC constexpr bool indices_statically_known_to_increase() {
return indices_statically_known_to_increase_impl<T>::run();
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_statically_eq(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_statically_eq(Index i, Index value) {
return index_statically_eq_impl<T>::run(i, value);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_statically_ne(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_statically_ne(Index i, Index value) {
return index_statically_ne_impl<T>::run(i, value);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_statically_gt(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_statically_gt(Index i, Index value) {
return index_statically_gt_impl<T>::run(i, value);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_statically_lt(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_statically_lt(Index i, Index value) {
return index_statically_lt_impl<T>::run(i, value);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_pair_first_statically_eq(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_pair_first_statically_eq(Index i, Index value) {
return index_pair_first_statically_eq_impl<T>::run(i, value);
}
template <typename T>
static EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bool index_pair_second_statically_eq(Index i, Index value) {
static EIGEN_DEVICE_FUNC constexpr bool index_pair_second_statically_eq(Index i, Index value) {
return index_pair_second_statically_eq_impl<T>::run(i, value);
}

19
unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h
View File

@@ -212,9 +212,9 @@ struct Pair {
typedef U first_type;
typedef V second_type;
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Pair() : first(), second() {}
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Pair() : first(), second() {}
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Pair(const U& f, const V& s) : first(f), second(s) {}
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Pair(const U& f, const V& s) : first(f), second(s) {}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void swap(Pair& rhs) {
using numext::swap;
@@ -224,20 +224,20 @@ struct Pair {
};
template <typename U, typename V>
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator==(const Pair<U, V>& x, const Pair<U, V>& y) {
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator==(const Pair<U, V>& x, const Pair<U, V>& y) {
return (x.first == y.first && x.second == y.second);
}
template <typename U, typename V>
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator!=(const Pair<U, V>& x, const Pair<U, V>& y) {
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator!=(const Pair<U, V>& x, const Pair<U, V>& y) {
return !(x == y);
}
// Can't use std::pairs on cuda devices
template <typename Idx>
struct IndexPair {
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE IndexPair() : first(0), second(0) {}
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE IndexPair(Idx f, Idx s) : first(f), second(s) {}
constexpr EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE IndexPair() : first(0), second(0) {}
constexpr EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE IndexPair(Idx f, Idx s) : first(f), second(s) {}
EIGEN_DEVICE_FUNC void set(IndexPair<Idx> val) {
first = val.first;
@@ -251,19 +251,18 @@ struct IndexPair {
namespace internal {
template <typename IndexType, typename Index, Index First, Index... Is>
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, 1 + sizeof...(Is)> customIndices2Array(
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, 1 + sizeof...(Is)> customIndices2Array(
IndexType& idx, numeric_list<Index, First, Is...>) {
return {static_cast<Index>(idx[First]), static_cast<Index>(idx[Is])...};
}
template <typename IndexType, typename Index>
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, 0> customIndices2Array(IndexType&,
numeric_list<Index>) {
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, 0> customIndices2Array(IndexType&, numeric_list<Index>) {
return array<Index, 0>();
}
/** Make an array (for index/dimensions) out of a custom index */
template <typename Index, std::size_t NumIndices, typename IndexType>
EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, NumIndices> customIndices2Array(IndexType& idx) {
constexpr EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE array<Index, NumIndices> customIndices2Array(IndexType& idx) {
return customIndices2Array(idx, typename gen_numeric_list<Index, NumIndices>::type{});
}

22
unsupported/Eigen/CXX11/src/Tensor/TensorReductionSycl.h
View File

@@ -402,9 +402,9 @@ struct SecondStepPartialReduction {
template <typename Index, Index LTP, Index LTR, bool BC_>
struct ReductionPannel {
static EIGEN_CONSTEXPR Index LocalThreadSizeP = LTP;
static EIGEN_CONSTEXPR Index LocalThreadSizeR = LTR;
static EIGEN_CONSTEXPR bool BC = BC_;
static constexpr Index LocalThreadSizeP = LTP;
static constexpr Index LocalThreadSizeR = LTR;
static constexpr bool BC = BC_;
};
template <typename Self, typename Op, TensorSycl::internal::reduction_dim rt>
@@ -430,7 +430,7 @@ struct PartialReducerLauncher {
"The Local thread size must be a power of 2 for the reduction "
"operation");
EIGEN_CONSTEXPR Index localRange = PannelParameters::LocalThreadSizeP * PannelParameters::LocalThreadSizeR;
constexpr Index localRange = PannelParameters::LocalThreadSizeP * PannelParameters::LocalThreadSizeR;
// In this step, we force the code not to be more than 2-step reduction:
// Our empirical research shows that if each thread reduces at least 64
// elements individually, we get better performance. However, this can change
@@ -445,7 +445,7 @@ struct PartialReducerLauncher {
const Index rNumGroups = num_coeffs_to_reduce > reductionPerThread * localRange ? std::min(rGroups, localRange) : 1;
const Index globalRange = pNumGroups * rNumGroups * localRange;
EIGEN_CONSTEXPR Index scratchSize =
constexpr Index scratchSize =
PannelParameters::LocalThreadSizeR * (PannelParameters::LocalThreadSizeP + PannelParameters::BC);
auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(globalRange), cl::sycl::range<1>(localRange));
if (rNumGroups > 1) {
@@ -482,15 +482,15 @@ template <typename Self, typename Op, bool Vectorizable>
struct FullReducer<Self, Op, Eigen::SyclDevice, Vectorizable> {
typedef typename Self::CoeffReturnType CoeffReturnType;
typedef typename Self::EvaluatorPointerType EvaluatorPointerType;
static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true;
static EIGEN_CONSTEXPR int PacketSize = Self::PacketAccess ? Self::PacketSize : 1;
static constexpr bool HasOptimizedImplementation = true;
static constexpr int PacketSize = Self::PacketAccess ? Self::PacketSize : 1;
static void run(const Self &self, Op &reducer, const Eigen::SyclDevice &dev, EvaluatorPointerType data) {
typedef std::conditional_t<Self::PacketAccess, typename Self::PacketReturnType, CoeffReturnType> OutType;
static_assert(!((EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1) &
(EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1 - 1)),
"The Local thread size must be a power of 2 for the reduction "
"operation");
EIGEN_CONSTEXPR Index local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1;
constexpr Index local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1;
typename Self::Index inputSize = self.impl().dimensions().TotalSize();
// In this step we force the code not to be more than 2-step reduction:
@@ -535,7 +535,7 @@ struct FullReducer<Self, Op, Eigen::SyclDevice, Vectorizable> {
// col reduction
template <typename Self, typename Op>
struct OuterReducer<Self, Op, Eigen::SyclDevice> {
static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true;
static constexpr bool HasOptimizedImplementation = true;
static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev,
typename Self::EvaluatorPointerType output, typename Self::Index num_coeffs_to_reduce,
@@ -549,7 +549,7 @@ struct OuterReducer<Self, Op, Eigen::SyclDevice> {
// row reduction
template <typename Self, typename Op>
struct InnerReducer<Self, Op, Eigen::SyclDevice> {
static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true;
static constexpr bool HasOptimizedImplementation = true;
static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev,
typename Self::EvaluatorPointerType output, typename Self::Index num_coeffs_to_reduce,
@@ -566,7 +566,7 @@ struct InnerReducer<Self, Op, Eigen::SyclDevice> {
// generic partial reduction
template <typename Self, typename Op>
struct GenericReducer<Self, Op, Eigen::SyclDevice> {
static EIGEN_CONSTEXPR bool HasOptimizedImplementation = false;
static constexpr bool HasOptimizedImplementation = false;
static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev,
typename Self::EvaluatorPointerType output, typename Self::Index num_values_to_reduce,
typename Self::Index num_coeffs_to_preserve) {

2
unsupported/Eigen/CXX11/src/Tensor/TensorScan.h
View File

@@ -197,7 +197,7 @@ struct ScanLauncher {
// threads. Currently set to twice the cache line size on Intel and ARM
// processors.
EIGEN_STRONG_INLINE Index AdjustBlockSize(Index item_size, Index block_size) {
EIGEN_CONSTEXPR Index kBlockAlignment = 128;
constexpr Index kBlockAlignment = 128;
const Index items_per_cacheline = numext::maxi<Index>(1, kBlockAlignment / item_size);
return items_per_cacheline * numext::div_ceil(block_size, items_per_cacheline);
}

6
unsupported/Eigen/CXX11/src/Tensor/TensorScanSycl.h
View File

@@ -51,7 +51,7 @@ namespace internal {
template <typename index_t>
struct ScanParameters {
// must be power of 2
static EIGEN_CONSTEXPR index_t ScanPerThread = 8;
static constexpr index_t ScanPerThread = 8;
const index_t total_size;
const index_t non_scan_size;
const index_t scan_size;
@@ -86,7 +86,7 @@ template <typename Evaluator, typename CoeffReturnType, typename OutAccessor, ty
struct ScanKernelFunctor {
typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local>
LocalAccessor;
static EIGEN_CONSTEXPR int PacketSize = ScanParameters<Index>::ScanPerThread / 2;
static constexpr int PacketSize = ScanParameters<Index>::ScanPerThread / 2;
LocalAccessor scratch;
Evaluator dev_eval;
@@ -288,7 +288,7 @@ template <typename CoeffReturnType, typename InAccessor, typename OutAccessor, t
struct ScanAdjustmentKernelFunctor {
typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local>
LocalAccessor;
static EIGEN_CONSTEXPR int PacketSize = ScanParameters<Index>::ScanPerThread / 2;
static constexpr int PacketSize = ScanParameters<Index>::ScanPerThread / 2;
InAccessor in_ptr;
OutAccessor out_ptr;
const ScanParameters<Index> scanParameters;

2
unsupported/Eigen/CXX11/src/TensorSymmetry/DynamicSymmetry.h
View File

@@ -106,7 +106,7 @@ class DynamicSGroup {
int one;
int two;
int flags;
constexpr inline Generator(int one_, int two_, int flags_) : one(one_), two(two_), flags(flags_) {}
constexpr Generator(int one_, int two_, int flags_) : one(one_), two(two_), flags(flags_) {}
};
std::size_t m_numIndices;

20
unsupported/Eigen/CXX11/src/TensorSymmetry/StaticSymmetry.h
View File

@@ -25,7 +25,7 @@ struct tensor_static_symgroup_permutate<numeric_list<int, nn...>> {
constexpr static std::size_t N = sizeof...(nn);
template <typename T>
constexpr static inline std::array<T, N> run(const std::array<T, N>& indices) {
constexpr static std::array<T, N> run(const std::array<T, N>& indices) {
return {{indices[nn]...}};
}
};
@@ -51,7 +51,7 @@ struct tensor_static_symgroup_identity_ctor {
template <typename iib>
struct tensor_static_symgroup_multiply_helper {
template <int... iia>
constexpr static inline numeric_list<int, get<iia, iib>::value...> helper(numeric_list<int, iia...>) {
constexpr static numeric_list<int, get<iia, iib>::value...> helper(numeric_list<int, iia...>) {
return numeric_list<int, get<iia, iib>::value...>();
}
};
@@ -107,9 +107,9 @@ struct tensor_static_symgroup {
};
template <typename Index, std::size_t N, int... ii, int... jj>
constexpr static inline std::array<Index, N> tensor_static_symgroup_index_permute(std::array<Index, N> idx,
internal::numeric_list<int, ii...>,
internal::numeric_list<int, jj...>) {
constexpr static std::array<Index, N> tensor_static_symgroup_index_permute(std::array<Index, N> idx,
internal::numeric_list<int, ii...>,
internal::numeric_list<int, jj...>) {
return {{idx[ii]..., idx[jj]...}};
}
@@ -179,9 +179,9 @@ class StaticSGroup {
typedef typename group_elements::type ge;
public:
constexpr inline StaticSGroup() {}
constexpr inline StaticSGroup(const StaticSGroup<Gen...>&) {}
constexpr inline StaticSGroup(StaticSGroup<Gen...>&&) {}
constexpr StaticSGroup() {}
constexpr StaticSGroup(const StaticSGroup<Gen...>&) {}
constexpr StaticSGroup(StaticSGroup<Gen...>&&) {}
template <typename Op, typename RV, typename Index, std::size_t N, typename... Args>
static inline RV apply(const std::array<Index, N>& idx, RV initial, Args&&... args) {
@@ -196,8 +196,8 @@ class StaticSGroup {
constexpr static std::size_t static_size = ge::count;
constexpr static inline std::size_t size() { return ge::count; }
constexpr static inline int globalFlags() { return group_elements::global_flags; }
constexpr static std::size_t size() { return ge::count; }
constexpr static int globalFlags() { return group_elements::global_flags; }
template <typename Tensor_, typename... IndexTypes>
inline internal::tensor_symmetry_value_setter<Tensor_, StaticSGroup<Gen...>> operator()(

6
unsupported/Eigen/CXX11/src/util/CXX11Workarounds.h
View File

@@ -36,15 +36,15 @@ namespace internal {
*/
template <std::size_t I_, class T>
constexpr inline T& array_get(std::vector<T>& a) {
constexpr T& array_get(std::vector<T>& a) {
return a[I_];
}
template <std::size_t I_, class T>
constexpr inline T&& array_get(std::vector<T>&& a) {
constexpr T&& array_get(std::vector<T>&& a) {
return a[I_];
}
template <std::size_t I_, class T>
constexpr inline T const& array_get(std::vector<T> const& a) {
constexpr T const& array_get(std::vector<T> const& a) {
return a[I_];
}