Updates will be applied April 15th at 12pm EDT (UTC-0400). GitLab could be a little slow between 12 - 12:45pm EDT.

Commit 9519737b authored by Robert Maynard's avatar Robert Maynard

Adding Reduce to the DeviceAdapterAlgorithm.

parent fbee20fc
......@@ -101,12 +101,27 @@ struct DeviceAdapterAlgorithm
const vtkm::cont::ArrayHandle<vtkm::Id,CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id,COut>& values_output);
/// \brief Compute a accumulated sum operation on the input ArrayHandle
///
/// Computes an accumulated sum (or any user binary operation) on the
/// \c input ArrayHandle, returning the total sum. Reduce is
/// similar to the stl accumulate sum function, exception that Reduce
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be commutative, or you will get
/// inconsistent results.
///
/// \return The total sum.
template<typename T, class CIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input,
T initialValue);
/// \brief Compute an inclusive prefix sum operation on the input ArrayHandle.
///
/// Computes an inclusive prefix sum operation on the \c input ArrayHandle,
/// storing the results in the \c output ArrayHandle. InclusiveScan is
/// similiar to the stl partial sum function, exception that InclusiveScan
/// doesn't do a serial sumnation. This means that if you have defined a
/// similar to the stl partial sum function, exception that InclusiveScan
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be associative, or you will get
/// inconsistent results. When the input and output ArrayHandles are the same
/// ArrayHandle the operation will be done inplace.
......@@ -122,8 +137,8 @@ struct DeviceAdapterAlgorithm
///
/// Computes an exclusive prefix sum operation on the \c input ArrayHandle,
/// storing the results in the \c output ArrayHandle. ExclusiveScan is
/// similiar to the stl partial sum function, exception that ExclusiveScan
/// doesn't do a serial sumnation. This means that if you have defined a
/// similar to the stl partial sum function, exception that ExclusiveScan
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be associative, or you will get
/// inconsistent results. When the input and output ArrayHandles are the same
/// ArrayHandle the operation will be done inplace.
......
......@@ -290,6 +290,16 @@ private:
IteratorBegin(values_output));
}
template<class InputPortal>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ReducePortal(const InputPortal &input,
typename InputPortal::ValueType initialValue)
{
return ::thrust::reduce(IteratorBegin(input),
IteratorEnd(input),
initialValue);
}
template<class InputPortal, class OutputPortal>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ScanExclusivePortal(const InputPortal &input,
......@@ -524,6 +534,18 @@ public:
LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTag()),
values_output.PrepareForInPlace(DeviceAdapterTag()));
}
template<typename T, class SIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,SIn> &input, T initialValue)
{
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
return initialValue;
}
return ReducePortal(input.PrepareForInput( DeviceAdapterTag() ),
initialValue);
}
template<typename T, class SIn, class SOut>
VTKM_CONT_EXPORT static T ScanExclusive(
......
......@@ -22,8 +22,9 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/cont/ArrayHandleZip.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/StorageBasic.h>
#include <vtkm/exec/FunctorBase.h>
......@@ -316,6 +317,107 @@ public:
values_output);
}
//--------------------------------------------------------------------------
// Reduce
private:
template<int ReduceWidth, typename T, typename ArrayType, typename BinaryOperation >
struct ReduceKernel : vtkm::exec::FunctorBase
{
typedef typename ArrayType::template ExecutionTypes<
DeviceAdapterTag> ExecutionTypes;
typedef typename ExecutionTypes::PortalConst PortalConst;
PortalConst Portal;
BinaryOperation BinaryOperator;
vtkm::Id ArrayLength;
VTKM_CONT_EXPORT
ReduceKernel()
: Portal(),
BinaryOperator(),
ArrayLength(0)
{
}
VTKM_CONT_EXPORT
ReduceKernel(const ArrayType &array, BinaryOperation op)
: Portal(array.PrepareForInput( DeviceAdapterTag() ) ),
BinaryOperator(op),
ArrayLength( array.GetNumberOfValues() )
{ }
VTKM_EXEC_EXPORT
T operator()(vtkm::Id index) const
{
const vtkm::Id offset = index * ReduceWidth;
//at least the first value access to the portal will be valid
//only the rest could be invalid
T partialSum = this->Portal.Get( offset );
if( offset + ReduceWidth >= this->ArrayLength )
{
vtkm::Id currentIndex = offset + 1;
while( currentIndex < this->ArrayLength)
{
partialSum = BinaryOperator(partialSum, this->Portal.Get(currentIndex));
++currentIndex;
}
}
else
{
//optimize the usecase where all values are valid and we don't
//need to check that we might go out of bounds
for(int i=1; i < ReduceWidth; ++i)
{
partialSum = BinaryOperator(partialSum,
this->Portal.Get( offset + i )
);
}
}
return partialSum;
}
};
public:
template<typename T, class CIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input, T initialValue)
{
//Crazy Idea:
//We create a implicit array handle that wraps the input
//array handle. The implicit functor is passed the input array handle, and
//the number of elements it needs to sum. This way the implicit handle
//acts as the first level reduction. Say for example reducing 16 values
//at a time.
//
//Now that we have an implicit array that is 1/16 the length of full array
//we can use scan inclusive to compute the final sum
typedef ReduceKernel<
16,
T,
vtkm::cont::ArrayHandle<T,CIn>,
vtkm::internal::Add
> ReduceKernelType;
typedef vtkm::cont::ArrayHandleImplicit<
T,
ReduceKernelType > ReduceHandleType;
typedef vtkm::cont::ArrayHandle<
T,
vtkm::cont::StorageTagBasic> TempArrayType;
ReduceKernelType kernel(input, vtkm::internal::Add());
vtkm::Id length = (input.GetNumberOfValues() / 16);
length += (input.GetNumberOfValues() % 16 == 0) ? 0 : 1;
ReduceHandleType reduced = vtkm::cont::make_ArrayHandleImplicit<T>(kernel,
length);
TempArrayType inclusiveScan;
return initialValue + DerivedAlgorithm::ScanInclusive(reduced,
inclusiveScan);
}
//--------------------------------------------------------------------------
// Scan Exclusive
private:
......
......@@ -1053,6 +1053,33 @@ private:
VTKM_TEST_ASSERT(value == OFFSET, "Got bad unique value");
}
static VTKM_CONT_EXPORT void TestReduce()
{
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing Reduce" << std::endl;
//construct the index array
IdArrayHandle array;
Algorithm::Schedule(
ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE,
DeviceAdapterTag())),
ARRAY_SIZE);
//the output of reduce and scan inclusive should be the same
vtkm::Id reduce_sum = Algorithm::Reduce(array, vtkm::Id(0));
vtkm::Id reduce_sum_with_intial_value = Algorithm::Reduce(array,
vtkm::Id(ARRAY_SIZE));
vtkm::Id inclusive_sum = Algorithm::ScanInclusive(array, array);
VTKM_TEST_ASSERT(reduce_sum == OFFSET * ARRAY_SIZE,
"Got bad sum from Reduce");
VTKM_TEST_ASSERT(reduce_sum_with_intial_value == reduce_sum + ARRAY_SIZE,
"Got bad sum from Reduce with initial value");
VTKM_TEST_ASSERT(reduce_sum == inclusive_sum,
"Got different sums from Reduce and ScanInclusive");
}
static VTKM_CONT_EXPORT void TestScanInclusive()
{
std::cout << "-------------------------------------------" << std::endl;
......@@ -1340,14 +1367,22 @@ private:
TestAlgorithmSchedule();
TestErrorExecution();
TestReduce();
TestScanInclusive();
TestScanExclusive();
TestSort();
TestSortWithComparisonObject();
TestSortByKey();
TestLowerBoundsWithComparisonObject();
TestUpperBoundsWithComparisonObject();
TestUniqueWithComparisonObject();
TestOrderedUniqueValues(); //tests Copy, LowerBounds, Sort, Unique
// TestDispatcher();
TestStreamCompactWithStencil();
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment