mirror of
https://gitlab.com/libeigen/eigen.git
synced 2026-04-10 11:34:33 +08:00
Merged eigen/eigen into default
This commit is contained in:
@@ -75,16 +75,16 @@ large enough to hold all the data.
|
||||
|
||||
// Map a tensor of ints on top of stack-allocated storage.
|
||||
int storage[128]; // 2 x 4 x 2 x 8 = 128
|
||||
TensorMap<int, 4> t_4d(storage, 2, 4, 2, 8);
|
||||
TensorMap<Tensor<int, 4>> t_4d(storage, 2, 4, 2, 8);
|
||||
|
||||
// The same storage can be viewed as a different tensor.
|
||||
// You can also pass the sizes as an array.
|
||||
TensorMap<int, 2> t_2d(storage, 16, 8);
|
||||
TensorMap<Tensor<int, 2>> t_2d(storage, 16, 8);
|
||||
|
||||
// You can also map fixed-size tensors. Here we get a 1d view of
|
||||
// the 2d fixed-size tensor.
|
||||
Tensor<float, Sizes<4, 5>> t_4x3;
|
||||
TensorMap<float, 1> t_12(t_4x3, 12);
|
||||
TensorMap<Tensor<float, 1>> t_12(t_4x3, 12);
|
||||
|
||||
|
||||
#### Class TensorRef
|
||||
|
||||
@@ -50,6 +50,7 @@ template <DenseIndex DimId>
|
||||
struct DimensionId
|
||||
{
|
||||
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE DimensionId(DenseIndex dim) {
|
||||
EIGEN_UNUSED_VARIABLE(dim);
|
||||
eigen_assert(dim == DimId);
|
||||
}
|
||||
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE DenseIndex actualDim() const {
|
||||
|
||||
@@ -529,7 +529,6 @@ EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rh
|
||||
float2 rhs_shmem2[][8], const Index m_size,
|
||||
const Index n_size, const Index k_size,
|
||||
const Index base_m, const Index base_n) {
|
||||
typedef float Scalar;
|
||||
|
||||
// prefetch registers
|
||||
float4 lhs_pf0, rhs_pf0;
|
||||
@@ -540,27 +539,27 @@ EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rh
|
||||
}
|
||||
|
||||
|
||||
#define prefetch_lhs(reg, row, col) \
|
||||
if (!CHECK_LHS_BOUNDARY) { \
|
||||
if (col < k_size) { \
|
||||
reg =lhs.loadPacket<Unaligned>(row, col); \
|
||||
} \
|
||||
} else { \
|
||||
if (col < k_size) { \
|
||||
if (row + 3 < m_size) { \
|
||||
reg =lhs.loadPacket<Unaligned>(row, col); \
|
||||
} else if (row + 2 < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
reg.y =lhs(row + 1, col); \
|
||||
reg.z =lhs(row + 2, col); \
|
||||
} else if (row + 1 < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
reg.y =lhs(row + 1, col); \
|
||||
} else if (row < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
} \
|
||||
} \
|
||||
} \
|
||||
#define prefetch_lhs(reg, row, col) \
|
||||
if (!CHECK_LHS_BOUNDARY) { \
|
||||
if (col < k_size) { \
|
||||
reg =lhs.template loadPacket<Unaligned>(row, col); \
|
||||
} \
|
||||
} else { \
|
||||
if (col < k_size) { \
|
||||
if (row + 3 < m_size) { \
|
||||
reg =lhs.template loadPacket<Unaligned>(row, col); \
|
||||
} else if (row + 2 < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
reg.y =lhs(row + 1, col); \
|
||||
reg.z =lhs(row + 2, col); \
|
||||
} else if (row + 1 < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
reg.y =lhs(row + 1, col); \
|
||||
} else if (row < m_size) { \
|
||||
reg.x =lhs(row + 0, col); \
|
||||
} \
|
||||
} \
|
||||
} \
|
||||
|
||||
|
||||
Index lhs_vert = base_m+threadIdx.x*4;
|
||||
@@ -578,7 +577,7 @@ EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rh
|
||||
if (!CHECK_RHS_BOUNDARY) {
|
||||
if ((rhs_vert + 3) < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf0 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
} else if (rhs_vert + 2 < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0.x = rhs(rhs_vert, rhs_horiz0);
|
||||
@@ -593,7 +592,7 @@ EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rh
|
||||
} else {
|
||||
if (rhs_horiz0 < n_size) {
|
||||
if ((rhs_vert + 3) < k_size) {
|
||||
rhs_pf0 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf0 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
} else if ((rhs_vert + 2) < k_size) {
|
||||
rhs_pf0.x = rhs(rhs_vert, rhs_horiz0);
|
||||
rhs_pf0.y = rhs(rhs_vert + 1, rhs_horiz0);
|
||||
@@ -766,7 +765,6 @@ EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs,
|
||||
float2 rhs_shmem2[][8], const Index m_size,
|
||||
const Index n_size, const Index k_size,
|
||||
const Index base_m, const Index base_n) {
|
||||
typedef float Scalar;
|
||||
|
||||
// prefetch registers
|
||||
float4 lhs_pf0, lhs_pf1, lhs_pf2, lhs_pf3;
|
||||
@@ -790,37 +788,37 @@ EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs,
|
||||
|
||||
if (!CHECK_LHS_BOUNDARY) {
|
||||
if ((threadIdx.y/4+k+24) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf3 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+24));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf3 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+24));
|
||||
} else if ((threadIdx.y/4+k+16) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
} else if ((threadIdx.y/4+k+8) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
} else if ((threadIdx.y/4+k) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
}
|
||||
} else {
|
||||
// just CHECK_LHS_BOUNDARY
|
||||
if (lhs_vert + 3 < m_size) {
|
||||
if ((threadIdx.y/4+k+24) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf3 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+24));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf3 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+24));
|
||||
} else if ((threadIdx.y/4+k+16) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf2 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+16));
|
||||
} else if ((threadIdx.y/4+k+8) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf1 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k+8));
|
||||
} else if ((threadIdx.y/4+k) < k_size) {
|
||||
lhs_pf0 =lhs.loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
lhs_pf0 =lhs.template loadPacket<Unaligned>(lhs_vert, (threadIdx.y/4+k));
|
||||
}
|
||||
} else if (lhs_vert + 2 < m_size) {
|
||||
if ((threadIdx.y/4+k+24) < k_size) {
|
||||
@@ -909,8 +907,8 @@ EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs,
|
||||
if (!CHECK_RHS_BOUNDARY) {
|
||||
if ((rhs_vert + 3) < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf1 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz1);
|
||||
rhs_pf0 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf1 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz1);
|
||||
} else if (rhs_vert + 2 < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0.x = rhs(rhs_vert, rhs_horiz0);
|
||||
@@ -932,8 +930,8 @@ EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs,
|
||||
if (rhs_horiz1 < n_size) {
|
||||
if ((rhs_vert + 3) < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf1 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz1);
|
||||
rhs_pf0 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf1 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz1);
|
||||
} else if (rhs_vert + 2 < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0.x = rhs(rhs_vert, rhs_horiz0);
|
||||
@@ -954,7 +952,7 @@ EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs,
|
||||
} else if (rhs_horiz0 < n_size) {
|
||||
if ((rhs_vert + 3) < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0 = rhs.loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
rhs_pf0 = rhs.template loadPacket<Unaligned>(rhs_vert, rhs_horiz0);
|
||||
} else if ((rhs_vert + 2) < k_size) {
|
||||
// just CHECK_RHS_BOUNDARY
|
||||
rhs_pf0.x = rhs(rhs_vert, rhs_horiz0);
|
||||
@@ -1137,9 +1135,6 @@ EigenFloatContractionKernel(const LhsMapper lhs, const RhsMapper rhs,
|
||||
typedef float2 LHS_MEM[64][32];
|
||||
typedef float2 RHS_MEM[128][8];
|
||||
|
||||
typedef float2 LHS_MEM16x16[32][16];
|
||||
typedef float2 RHS_MEM16x16[64][8];
|
||||
|
||||
const Index m_block_idx = blockIdx.x;
|
||||
const Index n_block_idx = blockIdx.y;
|
||||
|
||||
|
||||
@@ -217,7 +217,10 @@ struct GpuDevice {
|
||||
EIGEN_UNUSED_VARIABLE(err)
|
||||
assert(err == cudaSuccess);
|
||||
#else
|
||||
eigen_assert(false && "The default device should be used instead to generate kernel code");
|
||||
EIGEN_UNUSED_VARIABLE(dst);
|
||||
EIGEN_UNUSED_VARIABLE(src);
|
||||
EIGEN_UNUSED_VARIABLE(n);
|
||||
eigen_assert(false && "The default device should be used instead to generate kernel code");
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
@@ -11,6 +11,17 @@
|
||||
#ifndef EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H
|
||||
#define EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H
|
||||
|
||||
// clang is incompatible with the CUDA syntax wrt making a kernel a class friend,
|
||||
// so we'll use a macro to make clang happy.
|
||||
#ifndef KERNEL_FRIEND
|
||||
#if defined(__clang__) && defined(__CUDA__)
|
||||
#define KERNEL_FRIEND friend __global__
|
||||
#else
|
||||
#define KERNEL_FRIEND friend
|
||||
#endif
|
||||
#endif
|
||||
|
||||
|
||||
namespace Eigen {
|
||||
|
||||
|
||||
@@ -681,15 +692,15 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType, MakePointer_>,
|
||||
template <typename S, typename O, bool V> friend struct internal::FullReducerShard;
|
||||
#endif
|
||||
#if defined(EIGEN_USE_GPU) && defined(__CUDACC__)
|
||||
template <int B, int N, typename S, typename R, typename I> friend void internal::FullReductionKernel(R, const S, I, typename S::CoeffReturnType*, unsigned int*);
|
||||
template <int B, int N, typename S, typename R, typename I> KERNEL_FRIEND void internal::FullReductionKernel(R, const S, I, typename S::CoeffReturnType*, unsigned int*);
|
||||
#ifdef EIGEN_HAS_CUDA_FP16
|
||||
template <typename S, typename R, typename I> friend void internal::ReductionInitFullReduxKernelHalfFloat(R, const S, I, half2*);
|
||||
template <int B, int N, typename S, typename R, typename I> friend void internal::FullReductionKernelHalfFloat(R, const S, I, half*, half2*);
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::InnerReductionKernelHalfFloat(R, const S, I, I, half*);
|
||||
template <typename S, typename R, typename I> KERNEL_FRIEND void internal::ReductionInitFullReduxKernelHalfFloat(R, const S, I, half2*);
|
||||
template <int B, int N, typename S, typename R, typename I> KERNEL_FRIEND void internal::FullReductionKernelHalfFloat(R, const S, I, half*, half2*);
|
||||
template <int NPT, typename S, typename R, typename I> KERNEL_FRIEND void internal::InnerReductionKernelHalfFloat(R, const S, I, I, half*);
|
||||
#endif
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::InnerReductionKernel(R, const S, I, I, typename S::CoeffReturnType*);
|
||||
template <int NPT, typename S, typename R, typename I> KERNEL_FRIEND void internal::InnerReductionKernel(R, const S, I, I, typename S::CoeffReturnType*);
|
||||
|
||||
template <int NPT, typename S, typename R, typename I> friend void internal::OuterReductionKernel(R, const S, I, I, typename S::CoeffReturnType*);
|
||||
template <int NPT, typename S, typename R, typename I> KERNEL_FRIEND void internal::OuterReductionKernel(R, const S, I, I, typename S::CoeffReturnType*);
|
||||
#endif
|
||||
|
||||
#if defined(EIGEN_USE_SYCL)
|
||||
|
||||
@@ -287,7 +287,6 @@ struct FullReductionLauncher<
|
||||
void>::type> {
|
||||
static void run(const Self& self, Op& reducer, const GpuDevice& device, OutputType* output, typename Self::Index num_coeffs) {
|
||||
typedef typename Self::Index Index;
|
||||
typedef typename Self::CoeffReturnType Scalar;
|
||||
const int block_size = 256;
|
||||
const int num_per_thread = 128;
|
||||
const int num_blocks = divup<int>(num_coeffs, block_size * num_per_thread);
|
||||
|
||||
@@ -31,12 +31,12 @@ namespace Eigen {
|
||||
*
|
||||
* \sa Tensor
|
||||
*/
|
||||
template<typename T, typename Dimensions, int Options_> class TensorStorage;
|
||||
template<typename T, typename Dimensions, int Options> class TensorStorage;
|
||||
|
||||
|
||||
// Pure fixed-size storage
|
||||
template<typename T, int Options_, typename FixedDimensions>
|
||||
class TensorStorage<T, FixedDimensions, Options_>
|
||||
template<typename T, typename FixedDimensions, int Options_>
|
||||
class TensorStorage
|
||||
{
|
||||
private:
|
||||
static const std::size_t Size = FixedDimensions::total_size;
|
||||
@@ -66,7 +66,7 @@ class TensorStorage<T, FixedDimensions, Options_>
|
||||
|
||||
|
||||
// pure dynamic
|
||||
template<typename T, int Options_, typename IndexType, int NumIndices_>
|
||||
template<typename T, typename IndexType, int NumIndices_, int Options_>
|
||||
class TensorStorage<T, DSizes<IndexType, NumIndices_>, Options_>
|
||||
{
|
||||
public:
|
||||
|
||||
@@ -23,6 +23,7 @@ struct static_val {
|
||||
|
||||
template <typename T>
|
||||
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static_val(const T& v) {
|
||||
EIGEN_UNUSED_VARIABLE(v);
|
||||
eigen_assert(v == n);
|
||||
}
|
||||
};
|
||||
|
||||
Reference in New Issue
Block a user