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

Merged from trunk

Browse Source
This commit is contained in:
Benoit Steiner
2014-10-30 21:59:22 -07:00
parent fcecafde3a d62bfe73a9
commit 67fcf47ecb
5 changed files with 38 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h
View File

@@ -149,26 +149,26 @@ class TensorExecutor<Expression, ThreadPoolDevice, Vectorizable>
// GPU: the evaluation of the expression is offloaded to a GPU.
#if defined(EIGEN_USE_GPU) && defined(__CUDACC__)
template <typename Evaluator>
template <typename Evaluator, typename Index>
__global__ void
__launch_bounds__(1024)
EigenMetaKernel(Evaluator eval, unsigned int size) {
EigenMetaKernel(Evaluator eval, Index size) {
const int first_index = blockIdx.x * blockDim.x + threadIdx.x;
const int step_size = blockDim.x * gridDim.x;
const Index first_index = blockIdx.x * blockDim.x + threadIdx.x;
const Index step_size = blockDim.x * gridDim.x;
if (!Evaluator::PacketAccess || !Evaluator::IsAligned) {
// Use the scalar path
for (int i = first_index; i < size; i += step_size) {
for (Index i = first_index; i < size; i += step_size) {
eval.evalScalar(i);
}
}
else {
// Use the vector path
const int PacketSize = unpacket_traits<typename Evaluator::PacketReturnType>::size;
const int vectorized_step_size = step_size * PacketSize;
const int vectorized_size = (size / PacketSize) * PacketSize;
int i = first_index * PacketSize;
const Index PacketSize = unpacket_traits<typename Evaluator::PacketReturnType>::size;
const Index vectorized_step_size = step_size * PacketSize;
const Index vectorized_size = (size / PacketSize) * PacketSize;
Index i = first_index * PacketSize;
for ( ; i < vectorized_size; i += vectorized_step_size) {
eval.evalPacket(i);
}
@@ -193,7 +193,7 @@ class TensorExecutor<Expression, GpuDevice, Vectorizable>
const int block_size = maxCudaThreadsPerBlock();
const Index size = array_prod(evaluator.dimensions());
EigenMetaKernel<TensorEvaluator<Expression, GpuDevice> > <<<num_blocks, block_size, 0, device.stream()>>>(evaluator, size);
EigenMetaKernel<TensorEvaluator<Expression, GpuDevice>, Index><<<num_blocks, block_size, 0, device.stream()>>>(evaluator, size);
assert(cudaGetLastError() == cudaSuccess);
}
evaluator.cleanup();

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

@@ -459,7 +459,7 @@ struct TensorEvaluator<TensorSlicingOp<StartIndices, Sizes, ArgType>, Device>
this->m_impl.template writePacket<StoreMode>(inputIndices[0], x);
}
else {
CoeffReturnType values[packetSize];
EIGEN_ALIGN_DEFAULT CoeffReturnType values[packetSize];
internal::pstore<CoeffReturnType, PacketReturnType>(values, x);
this->m_impl.coeffRef(inputIndices[0]) = values[0];
this->m_impl.coeffRef(inputIndices[1]) = values[packetSize-1];

18
unsupported/Eigen/CXX11/src/Tensor/TensorPadding.h
View File

@@ -98,7 +98,6 @@ struct TensorEvaluator<const TensorPaddingOp<PaddingDimensions, ArgType>, Device
for (int i = 0; i < NumDims; ++i) {
m_dimensions[i] += m_padding[i].first + m_padding[i].second;
}
const typename TensorEvaluator<ArgType, Device>::Dimensions& input_dims = m_impl.dimensions();
m_inputStrides[0] = 1;
m_outputStrides[0] = 1;
@@ -125,6 +124,7 @@ struct TensorEvaluator<const TensorPaddingOp<PaddingDimensions, ArgType>, Device
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
{
eigen_assert(index < dimensions().TotalSize());
Index inputIndex = 0;
for (int i = NumDims - 1; i > 0; --i) {
const Index idx = index / m_outputStrides[i];
@@ -151,11 +151,11 @@ struct TensorEvaluator<const TensorPaddingOp<PaddingDimensions, ArgType>, Device
const Index initialIndex = index;
Index inputIndex = 0;
for (int i = NumDims - 1; i > 0; --i) {
const int first = index;
const int last = index + packetSize - 1;
const int lastPaddedLeft = m_padding[i].first * m_outputStrides[i];
const int firstPaddedRight = (m_dimensions[i] - m_padding[i].second) * m_outputStrides[i];
const int lastPaddedRight = m_outputStrides[i+1];
const Index first = index;
const Index last = index + packetSize - 1;
const Index lastPaddedLeft = m_padding[i].first * m_outputStrides[i];
const Index firstPaddedRight = (m_dimensions[i] - m_padding[i].second) * m_outputStrides[i];
const Index lastPaddedRight = m_outputStrides[i+1];
if (last < lastPaddedLeft) {
// all the coefficient are in the padding zone.
@@ -179,9 +179,9 @@ struct TensorEvaluator<const TensorPaddingOp<PaddingDimensions, ArgType>, Device
const Index last = index + packetSize - 1;
const Index first = index;
const int lastPaddedLeft = m_padding[0].first;
const int firstPaddedRight = (m_dimensions[0] - m_padding[0].second);
const int lastPaddedRight = m_outputStrides[1];
const Index lastPaddedLeft = m_padding[0].first;
const Index firstPaddedRight = (m_dimensions[0] - m_padding[0].second);
const Index lastPaddedRight = m_outputStrides[1];
if (last < lastPaddedLeft) {
// all the coefficient are in the padding zone.