Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-08 13:23:51 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Add Mask capabilities to worklets

Browse Source 
Mask objects allow you to specify which output values should be
generated when a worklet is run. That is, the Mask allows you to skip
the invocation of a worklet for any number of outputs.
This commit is contained in:
Kenneth Moreland 2019-01-20 19:14:55 -07:00
parent c27a33669b
commit 191d6e5580
38 changed files with 1608 additions and 365 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

104
docs/changelog/mask-worklets.md Normal file
View File

@ -0,0 +1,104 @@
# Allow masking of worklet invocations
There have recently been use cases where it would be helpful to mask out
some of the invocations of a worklet. The idea is that when invoking a
worklet with a mask array on the input domain, you might implement your
worklet more-or-less like the following.
```cpp
VTKM_EXEC void operator()(bool mask, /* other parameters */)
{
if (mask)
{
// Do interesting stuff
}
}
```
This works, but what if your mask has mostly false values? In that case,
you are spending tons of time loading data to and from memory where fields
are stored for no reason.
You could potentially get around this problem by adding a scatter to the
worklet. However, that will compress the output arrays to only values that
are active in the mask. That is problematic if you want the masked output
in the appropriate place in the original arrays. You will have to do some
complex (and annoying and possibly expensive) permutations of the output
arrays.
Thus, we would like a new feature similar to scatter that instead masks out
invocations so that the worklet is simply not run on those outputs.
## New Interface
The new "Mask" feature that is similar (and orthogonal) to the existing
"Scatter" feature. Worklet objects now define a `MaskType` that provides on
object that manages the selections of which invocations are skipped. The
following Mask objects are defined.
* `MaskNone` - This removes any mask of the output. All outputs are
generated. This is the default if no `MaskType` is explicitly defined.
* `MaskSelect` - The user to provides an array that specifies whether
each output is created with a 1 to mean that the output should be
created an 0 the mean that it should not.
* `MaskIndices` - The user provides an array with a list of indices for
all outputs that should be created.
It will be straightforward to implement other versions of masks. (For
example, you could make a mask class that selectes every Nth entry.) Those
could be made on an as-needed basis.
## Implementation
The implementation follows the same basic idea of how scatters are
implemented.
### Mask Classes
The mask class is required to implement the following items.
* `ThreadToOutputType` - A type for an array that maps a thread index (an
index in the array) to an output index. A reasonable type for this
could be `vtkm::cont::ArrayHandle<vtkm::Id>`.
* `GetThreadToOutputMap` - Given the range for the output (e.g. the
number of items in the output domain), returns an array of type
`ThreadToOutputType` that is the actual map.
* `GetThreadRange` - Given a range for the output (e.g. the number of
items in the output domain), returns the range for the threads (e.g.
the number of times the worklet will be invoked).
### Dispatching
The `vtkm::worklet::internal::DispatcherBase` manages a mask class in
the same way it manages the scatter class. It gets the `MaskType` from
the worklet it is templated on. It requires a `MaskType` object during
its construction.
Previously the dispatcher (and downstream) had to manage the range and
indices of inputs and threads. They now have to also manage a separate
output range/index as now all three may be different.
The `vtkm::Invocation` is changed to hold the ThreadToOutputMap array from
the mask. It likewises has a templated `ChangeThreadToOutputMap` method
added (similar to those already existing for the arrays from a scatter).
This method is used in `DispatcherBase::InvokeTransportParameters` to add
the mask's array to the invocation before calling `InvokeSchedule`.
### Thread Indices
With the addition of masks, the `ThreadIndices` classes are changed to
manage the actual output index. Previously, the output index was always the
same as the thread index. However, now these two can be different. The
`GetThreadIndices` methods of the worklet base classes have an argument
added that is the portal to the ThreadToOutputMap.
The worklet `GetThreadIndices` is called from the `Task` classes. These
classes are changed to pass in this additional argument. Since the `Task`
classes get an `Invocation` object from the dispatcher, which contains the
`ThreadToOutputMap`, this change is trivial.
## Interaction Between Mask and Scatter
Although it seems weird, it should work fine to mix scatters and masks. The
scatter will first be applied to the input to generate a (potential) list
of output elements. The mask will then be applied to these output elements.

18
vtkm/exec/arg/ThreadIndicesBasic.h
View File

@ -49,14 +49,25 @@ public:
ThreadIndicesBasic(vtkm::Id threadIndex,
vtkm::Id inIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outIndex,
vtkm::Id globalThreadIndexOffset = 0)
: InputIndex(inIndex)
, OutputIndex(threadIndex)
: ThreadIndex(threadIndex)
, InputIndex(inIndex)
, OutputIndex(outIndex)
, VisitIndex(visitIndex)
, GlobalThreadIndexOffset(globalThreadIndexOffset)
{
}
/// \brief The index of the thread or work invocation.
///
/// This index refers to which instance of the worklet is being invoked. Every invocation of the
/// worklet has a unique thread index. This is also called the work index depending on the
/// context.
///
VTKM_EXEC
vtkm::Id GetThreadIndex() const { return this->ThreadIndex; }
/// \brief The index into the input domain.
///
/// This index refers to the input element (array value, cell, etc.) that
@ -98,9 +109,10 @@ public:
///
/// Global index (for streaming)
VTKM_EXEC
vtkm::Id GetGlobalIndex() const { return (this->GlobalThreadIndexOffset + this->OutputIndex); }
vtkm::Id GetGlobalIndex() const { return (this->GlobalThreadIndexOffset + this->ThreadIndex); }
private:
vtkm::Id ThreadIndex;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;

49
vtkm/exec/arg/ThreadIndicesPointNeighborhood.h
View File

@ -65,19 +65,14 @@ class ThreadIndicesPointNeighborhood
{
public:
template <typename OutToInArrayType, typename VisitArrayType, vtkm::IdComponent Dimension>
template <vtkm::IdComponent Dimension>
VTKM_EXEC ThreadIndicesPointNeighborhood(
const vtkm::Id3& outIndex,
const OutToInArrayType&,
const VisitArrayType&,
const vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint,
Dimension>& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: State(outIndex, detail::To3D(connectivity.GetPointDimensions()))
, InputIndex(0)
, OutputIndex(0)
, VisitIndex(0)
, GlobalThreadIndexOffset(globalThreadIndexOffset)
{
using ConnectivityType = vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagCell,
@ -85,41 +80,27 @@ public:
Dimension>;
using ConnRangeType = typename ConnectivityType::SchedulingRangeType;
const ConnRangeType index = detail::Deflate(outIndex, ConnRangeType());
this->InputIndex = connectivity.LogicalToFlatToIndex(index);
this->OutputIndex = this->InputIndex;
}
template <typename OutToInArrayType, typename VisitArrayType, vtkm::IdComponent Dimension>
VTKM_EXEC ThreadIndicesPointNeighborhood(
const vtkm::Id& outIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint,
Dimension>& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: State(detail::To3D(connectivity.FlatToLogicalToIndex(outToIn.Get(outIndex))),
detail::To3D(connectivity.GetPointDimensions()))
, InputIndex(outToIn.Get(outIndex))
, OutputIndex(outIndex)
, VisitIndex(static_cast<vtkm::IdComponent>(visit.Get(outIndex)))
, GlobalThreadIndexOffset(globalThreadIndexOffset)
{
this->ThreadIndex = connectivity.LogicalToFlatToIndex(index);
this->InputIndex = this->ThreadIndex;
this->VisitIndex = 0;
this->OutputIndex = this->ThreadIndex;
}
template <vtkm::IdComponent Dimension>
VTKM_EXEC ThreadIndicesPointNeighborhood(
const vtkm::Id& outIndex,
const vtkm::Id& inIndex,
const vtkm::IdComponent& visitIndex,
vtkm::Id threadIndex,
vtkm::Id inputIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outputIndex,
const vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint,
Dimension>& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: State(detail::To3D(connectivity.FlatToLogicalToIndex(inIndex)),
: State(detail::To3D(connectivity.FlatToLogicalToIndex(inputIndex)),
detail::To3D(connectivity.GetPointDimensions()))
, InputIndex(inIndex)
, OutputIndex(outIndex)
, ThreadIndex(threadIndex)
, InputIndex(inputIndex)
, OutputIndex(outputIndex)
, VisitIndex(visitIndex)
, GlobalThreadIndexOffset(globalThreadIndexOffset)
{
@ -128,6 +109,9 @@ public:
VTKM_EXEC
const vtkm::exec::BoundaryState& GetBoundaryState() const { return this->State; }
VTKM_EXEC
vtkm::Id GetThreadIndex() const { return this->ThreadIndex; }
VTKM_EXEC
vtkm::Id GetInputIndex() const { return this->InputIndex; }
@ -148,6 +132,7 @@ public:
private:
vtkm::exec::BoundaryState State;
vtkm::Id ThreadIndex;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;

3
vtkm/exec/arg/ThreadIndicesReduceByKey.h
View File

@ -48,9 +48,10 @@ public:
vtkm::Id threadIndex,
vtkm::Id inIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outIndex,
const vtkm::exec::internal::ReduceByKeyLookup<P1, P2, P3>& keyLookup,
vtkm::Id globalThreadIndexOffset = 0)
: Superclass(threadIndex, inIndex, visitIndex, globalThreadIndexOffset)
: Superclass(threadIndex, inIndex, visitIndex, outIndex, globalThreadIndexOffset)
, ValueOffset(keyLookup.Offsets.Get(inIndex))
, NumberOfValues(keyLookup.Counts.Get(inIndex))
{

122
vtkm/exec/arg/ThreadIndicesTopologyMap.h
View File

@ -100,41 +100,20 @@ public:
using CellShapeTag = typename ConnectivityType::CellShapeTag;
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename OutToInArrayType, typename VisitArrayType>
VTKM_EXEC ThreadIndicesTopologyMap(vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
vtkm::Id inputIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outputIndex,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: Superclass(threadIndex,
outToIn.Get(threadIndex),
visit.Get(threadIndex),
globalThreadIndexOffset)
: Superclass(threadIndex, inputIndex, visitIndex, outputIndex, globalThreadIndexOffset)
// The connectivity is stored in the invocation parameter at the given
// input domain index. If this class is being used correctly, the type
// of the domain will match the connectivity type used here. If there is
// a compile error here about a type mismatch, chances are a worklet has
// set its input domain incorrectly.
, IndicesFrom(connectivity.GetIndices(outToIn.Get(threadIndex)))
, CellShape(connectivity.GetCellShape(outToIn.Get(threadIndex)))
{
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC
ThreadIndicesTopologyMap(vtkm::Id threadIndex,
vtkm::Id inIndex,
vtkm::IdComponent visitIndex,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: Superclass(threadIndex, inIndex, visitIndex, globalThreadIndexOffset)
// The connectivity is stored in the invocation parameter at the given
// input domain index. If this class is being used correctly, the type
// of the domain will match the connectivity type used here. If there is
// a compile error here about a type mismatch, chances are a worklet has
// set its input domain incorrectly.
, IndicesFrom(connectivity.GetIndices(inIndex))
, CellShape(connectivity.GetCellShape(inIndex))
, IndicesFrom(connectivity.GetIndices(inputIndex))
, CellShape(connectivity.GetCellShape(inputIndex))
{
}
@ -187,61 +166,51 @@ public:
using CellShapeTag = typename ConnectivityType::CellShapeTag;
using LogicalIndexType = typename ConnectivityType::SchedulingRangeType;
template <typename OutToInArrayType, typename VisitArrayType>
VTKM_EXEC ThreadIndicesTopologyMap(vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
vtkm::Id inIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outIndex,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
{
this->InputIndex = outToIn.Get(threadIndex);
this->OutputIndex = threadIndex;
this->VisitIndex = visit.Get(threadIndex);
this->ThreadIndex = threadIndex;
this->InputIndex = inIndex;
this->VisitIndex = visitIndex;
this->OutputIndex = outIndex;
this->LogicalIndex = connectivity.FlatToLogicalToIndex(this->InputIndex);
this->IndicesFrom = connectivity.GetIndices(this->LogicalIndex);
this->CellShape = connectivity.GetCellShape(this->InputIndex);
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
template <typename OutToInArrayType, typename VisitArrayType>
VTKM_EXEC ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex,
const OutToInArrayType&,
const VisitArrayType& visit,
const ConnectivityType& connectivity,
const vtkm::Id globalThreadIndexOffset = 0)
{
// We currently only support multidimensional indices on one-to-one input-
// to-output mappings. (We don't have a use case otherwise.)
// that is why the OutToInArrayType is ignored
// That is why we treat teh threadIndex as also the inputIndex and outputIndex
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
this->ThreadIndex = index;
this->InputIndex = index;
this->OutputIndex = index;
this->VisitIndex = visit.Get(index);
this->VisitIndex = 0;
this->LogicalIndex = logicalIndex;
this->IndicesFrom = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(index);
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
VTKM_SUPPRESS_EXEC_WARNINGS
/// \brief The index of the thread or work invocation.
///
/// This index refers to which instance of the worklet is being invoked. Every invocation of the
/// worklet has a unique thread index. This is also called the work index depending on the
/// context.
///
VTKM_EXEC
ThreadIndicesTopologyMap(vtkm::Id threadIndex,
vtkm::Id vtkmNotUsed(inIndex),
vtkm::IdComponent visitIndex,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
{
this->InputIndex = threadIndex;
this->OutputIndex = threadIndex;
this->VisitIndex = visitIndex;
this->LogicalIndex = connectivity.FlatToLogicalToIndex(this->InputIndex);
this->IndicesFrom = connectivity.GetIndices(this->LogicalIndex);
this->CellShape = connectivity.GetCellShape(this->InputIndex);
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
vtkm::Id GetThreadIndex() const { return this->ThreadIndex; }
/// \brief The logical index into the input domain.
///
@ -321,9 +290,10 @@ public:
CellShapeTag GetCellShape() const { return this->CellShape; }
private:
vtkm::Id ThreadIndex;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
vtkm::Id OutputIndex;
LogicalIndexType LogicalIndex;
IndicesFromType IndicesFrom;
CellShapeTag CellShape;
@ -351,35 +321,17 @@ public:
using CellShapeTag = typename ConnectivityType::CellShapeTag;
using LogicalIndexType = typename ConnectivityType::SchedulingRangeType;
template <typename OutToInArrayType, typename VisitArrayType>
VTKM_EXEC ThreadIndicesTopologyMap(vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
vtkm::Id inputIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outputIndex,
const PermutedConnectivityType& permutation,
vtkm::Id globalThreadIndexOffset = 0)
{
this->InputIndex = outToIn.Get(threadIndex);
this->OutputIndex = threadIndex;
this->VisitIndex = visit.Get(threadIndex);
const vtkm::Id permutedIndex = permutation.Portal.Get(this->InputIndex);
this->LogicalIndex = permutation.Connectivity.FlatToLogicalToIndex(permutedIndex);
this->IndicesFrom = permutation.Connectivity.GetIndices(this->LogicalIndex);
this->CellShape = permutation.Connectivity.GetCellShape(permutedIndex);
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC
ThreadIndicesTopologyMap(vtkm::Id threadIndex,
vtkm::Id vtkmNotUsed(inIndex),
vtkm::IdComponent visitIndex,
const PermutedConnectivityType& permutation,
vtkm::Id globalThreadIndexOffset = 0)
{
this->InputIndex = threadIndex;
this->OutputIndex = threadIndex;
this->ThreadIndex = threadIndex;
this->InputIndex = inputIndex;
this->VisitIndex = visitIndex;
this->OutputIndex = outputIndex;
const vtkm::Id permutedIndex = permutation.Portal.Get(this->InputIndex);
this->LogicalIndex = permutation.Connectivity.FlatToLogicalToIndex(permutedIndex);
@ -388,6 +340,15 @@ public:
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
/// \brief The index of the thread or work invocation.
///
/// This index refers to which instance of the worklet is being invoked. Every invocation of the
/// worklet has a unique thread index. This is also called the work index depending on the
/// context.
///
VTKM_EXEC
vtkm::Id GetThreadIndex() const { return this->ThreadIndex; }
/// \brief The logical index into the input domain.
///
/// This is similar to \c GetIndex3D except the Vec size matches the actual
@ -466,9 +427,10 @@ public:
CellShapeTag GetCellShape() const { return this->CellShape; }
private:
vtkm::Id ThreadIndex;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
vtkm::Id OutputIndex;
LogicalIndexType LogicalIndex;
IndicesFromType IndicesFrom;
CellShapeTag CellShape;

14
vtkm/exec/arg/testing/UnitTestFetchArrayNeighborhoodIn.cxx
View File

@ -45,14 +45,6 @@ struct TestPortal
}
};
struct TestIndexPortal
{
using ValueType = vtkm::Id;
VTKM_EXEC_CONT
ValueType Get(vtkm::Id index) const { return index; }
};
template <typename NeighborhoodType, typename T>
void verify_neighbors(NeighborhoodType neighbors, vtkm::Id index, vtkm::Id3 index3d, T)
{
@ -127,8 +119,7 @@ struct FetchArrayNeighborhoodInTests
for (vtkm::Id i = 0; i < POINT_DIMS[0]; i++, index++)
{
index3d[0] = i;
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(
index3d, vtkm::internal::NullType(), vtkm::internal::NullType(), connectivity);
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(index3d, connectivity);
auto neighbors = fetch.Load(indices, execObject);
@ -149,8 +140,7 @@ struct FetchArrayNeighborhoodInTests
//Verify that 1D scheduling works with neighborhoods
for (vtkm::Id index = 0; index < (POINT_DIMS[0] * POINT_DIMS[1] * POINT_DIMS[2]); index++)
{
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(
index, TestIndexPortal(), TestIndexPortal(), connectivity);
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(index, index, 0, index, connectivity);
auto neighbors = fetch.Load(indices, execObject);

51
vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx
View File

@ -82,8 +82,12 @@ struct FetchArrayTopologyMapInTests
FetchType fetch;
const vtkm::Id threadIndex = 0;
const vtkm::Id outputIndex = invocation.ThreadToOutputMap.Get(threadIndex);
const vtkm::Id inputIndex = invocation.OutputToInputMap.Get(outputIndex);
const vtkm::IdComponent visitIndex = invocation.VisitArray.Get(outputIndex);
ThreadIndicesType indices(
0, invocation.OutputToInputMap, invocation.VisitArray, invocation.GetInputDomain());
threadIndex, inputIndex, visitIndex, outputIndex, invocation.GetInputDomain());
typename FetchType::ValueType value =
fetch.Load(indices, invocation.Parameters.template GetParameter<ParamIndex>());
@ -125,7 +129,8 @@ struct FetchArrayTopologyMapInTests
BaseFunctionInterface(),
BaseFunctionInterface(),
TestIndexPortal(),
TestZeroPortal()));
TestZeroPortal(),
TestIndexPortal()));
}
};
@ -157,20 +162,32 @@ void TryStructuredPointCoordinatesInvocation(const Invocation& invocation)
vtkm::Vec<vtkm::FloatDefault, 3> origin = TestValue(0, vtkm::Vec<vtkm::FloatDefault, 3>());
vtkm::Vec<vtkm::FloatDefault, 3> spacing = TestValue(1, vtkm::Vec<vtkm::FloatDefault, 3>());
vtkm::VecAxisAlignedPointCoordinates<NumDimensions> value = fetch.Load(
ThreadIndicesType(
0, invocation.OutputToInputMap, invocation.VisitArray, invocation.GetInputDomain()),
invocation.Parameters.template GetParameter<ParamIndex>());
VTKM_TEST_ASSERT(test_equal(value.GetOrigin(), origin), "Bad origin.");
VTKM_TEST_ASSERT(test_equal(value.GetSpacing(), spacing), "Bad spacing.");
{
const vtkm::Id threadIndex = 0;
const vtkm::Id outputIndex = invocation.ThreadToOutputMap.Get(threadIndex);
const vtkm::Id inputIndex = invocation.OutputToInputMap.Get(outputIndex);
const vtkm::IdComponent visitIndex = invocation.VisitArray.Get(outputIndex);
vtkm::VecAxisAlignedPointCoordinates<NumDimensions> value =
fetch.Load(ThreadIndicesType(
threadIndex, inputIndex, visitIndex, outputIndex, invocation.GetInputDomain()),
invocation.Parameters.template GetParameter<ParamIndex>());
VTKM_TEST_ASSERT(test_equal(value.GetOrigin(), origin), "Bad origin.");
VTKM_TEST_ASSERT(test_equal(value.GetSpacing(), spacing), "Bad spacing.");
}
origin[0] += spacing[0];
value = fetch.Load(
ThreadIndicesType(
1, invocation.OutputToInputMap, invocation.VisitArray, invocation.GetInputDomain()),
invocation.Parameters.template GetParameter<ParamIndex>());
VTKM_TEST_ASSERT(test_equal(value.GetOrigin(), origin), "Bad origin.");
VTKM_TEST_ASSERT(test_equal(value.GetSpacing(), spacing), "Bad spacing.");
{
const vtkm::Id threadIndex = 1;
const vtkm::Id outputIndex = invocation.ThreadToOutputMap.Get(threadIndex);
const vtkm::Id inputIndex = invocation.OutputToInputMap.Get(outputIndex);
const vtkm::IdComponent visitIndex = invocation.VisitArray.Get(outputIndex);
vtkm::VecAxisAlignedPointCoordinates<NumDimensions> value =
fetch.Load(ThreadIndicesType(
threadIndex, inputIndex, visitIndex, outputIndex, invocation.GetInputDomain()),
invocation.Parameters.template GetParameter<ParamIndex>());
VTKM_TEST_ASSERT(test_equal(value.GetOrigin(), origin), "Bad origin.");
VTKM_TEST_ASSERT(test_equal(value.GetSpacing(), spacing), "Bad spacing.");
}
}
template <vtkm::IdComponent NumDimensions>
@ -192,14 +209,16 @@ void TryStructuredPointCoordinates(
BaseFunctionInterface(),
BaseFunctionInterface(),
TestIndexPortal(),
TestZeroPortal()));
TestZeroPortal(),
TestIndexPortal()));
// Try again with topology in argument 3 and point coordinates in argument 1
TryStructuredPointCoordinatesInvocation<NumDimensions, 1>(vtkm::internal::make_Invocation<3>(
BaseFunctionInterface().Replace<3>(connectivity).template Replace<1>(coordinates),
BaseFunctionInterface(),
BaseFunctionInterface(),
TestIndexPortal(),
TestZeroPortal()));
TestZeroPortal(),
TestIndexPortal()));
}
void TryStructuredPointCoordinates()

2
vtkm/exec/cuda/internal/TaskStrided.h
View File

@ -87,6 +87,7 @@ public:
this->Worklet.GetThreadIndices(index,
this->Invocation.OutputToInputMap,
this->Invocation.VisitArray,
this->Invocation.ThreadToOutputMap,
this->Invocation.GetInputDomain(),
this->GlobalIndexOffset));
}
@ -157,6 +158,7 @@ public:
this->Worklet.GetThreadIndices(index,
this->Invocation.OutputToInputMap,
this->Invocation.VisitArray,
this->Invocation.ThreadToOutputMap,
this->Invocation.GetInputDomain(),
this->GlobalIndexOffset));
}

23
vtkm/exec/cuda/internal/testing/UnitTestTaskStrided.cu
View File

@ -74,6 +74,13 @@ struct MyVisitArrayPortal
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct MyThreadToOutputMapPortal
{
using ValueType = vtkm::Id;
VTKM_EXEC_CONT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct TestFetchTagInput
{
};
@ -169,14 +176,16 @@ using InvocationType1 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestExecutionInterface1,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
using InvocationType2 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface2,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
template <typename TaskType>
static __global__ void ScheduleTaskStrided(TaskType task, vtkm::Id start, vtkm::Id end)
@ -202,17 +211,20 @@ struct TestWorkletProxy : vtkm::exec::FunctorBase
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};
@ -227,17 +239,20 @@ struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};

16
vtkm/exec/internal/TaskSingular.h
View File

@ -65,13 +65,15 @@ public:
VTKM_EXEC void operator()(T index) const
{
//Todo: rename this function to DoTaskSingular
detail::DoWorkletInvokeFunctor(this->Worklet,
this->Invocation,
this->Worklet.GetThreadIndices(index,
this->Invocation.OutputToInputMap,
this->Invocation.VisitArray,
this->Invocation.GetInputDomain(),
GlobalIndexOffset));
detail::DoWorkletInvokeFunctor(
this->Worklet,
this->Invocation,
this->Worklet.GetThreadIndices(index,
this->Invocation.OutputToInputMap,
this->Invocation.VisitArray,
this->Invocation.ThreadToOutputMap,
this->Invocation.GetInputDomain(),
GlobalIndexOffset));
}
private:

140
vtkm/exec/internal/WorkletInvokeFunctorDetail.h
View File

@ -91,6 +91,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1>
@ -101,7 +102,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -147,7 +148,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1>
VTKM_EXEC void DoWorkletInvokeFunctor(
@ -157,7 +158,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -197,6 +198,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -208,7 +210,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -264,7 +266,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2>
@ -275,7 +277,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -325,6 +327,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -337,7 +340,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -403,7 +406,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -415,7 +418,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -475,6 +478,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -488,7 +492,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -564,7 +568,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -577,7 +581,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -647,6 +651,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -661,7 +666,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -747,7 +752,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -761,7 +766,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -841,6 +846,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -856,7 +862,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -952,7 +958,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -967,7 +973,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -1057,6 +1063,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1073,7 +1080,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -1179,7 +1186,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1195,7 +1202,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -1295,6 +1302,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1312,7 +1320,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -1428,7 +1436,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1445,7 +1453,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -1555,6 +1563,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1573,7 +1582,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -1699,7 +1708,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1717,7 +1726,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -1837,6 +1846,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1856,7 +1866,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -1992,7 +2002,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -2011,7 +2021,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -2141,6 +2151,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -2161,7 +2172,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -2307,7 +2318,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -2327,7 +2338,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -2467,6 +2478,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -2488,7 +2500,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -2644,7 +2656,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -2665,7 +2677,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -2815,6 +2827,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -2837,7 +2850,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -3003,7 +3016,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -3025,7 +3038,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -3185,6 +3198,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -3208,7 +3222,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -3384,7 +3398,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -3407,7 +3421,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -3577,6 +3591,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -3601,7 +3616,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -3787,7 +3802,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -3811,7 +3826,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -3991,6 +4006,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -4016,7 +4032,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -4212,7 +4228,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -4237,7 +4253,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -4427,6 +4443,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -4453,7 +4470,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -4659,7 +4676,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -4685,7 +4702,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -4885,6 +4902,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -4912,7 +4930,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -5128,7 +5146,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -5155,7 +5173,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -5365,6 +5383,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -5393,7 +5412,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18, P19)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -5619,7 +5638,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -5647,7 +5666,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18, P19)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =
@ -5867,6 +5886,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -5896,7 +5916,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<R(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18, P19, P20)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -6132,7 +6152,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -6161,7 +6181,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<void(P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11, P12, P13, P14, P15, P16, P17, P18, P19, P20)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =

7
vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in
View File

@ -153,6 +153,7 @@ template <typename WorkletType,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadToOutputMapType,
typename ThreadIndicesType,
$template_params(num_params)>
VTKM_EXEC void DoWorkletInvokeFunctor(
@ -162,7 +163,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<$signature(num_params)>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation = vtkm::internal::Invocation<ParameterInterface,
@ -212,7 +213,7 @@ template <typename WorkletType,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename VisitArrayType, typename ThreadToOutputMapType,
typename ThreadIndicesType,
$template_params(num_params, start=1)>
VTKM_EXEC void DoWorkletInvokeFunctor(
@ -222,7 +223,7 @@ VTKM_EXEC void DoWorkletInvokeFunctor(
vtkm::internal::FunctionInterface<$signature(num_params, return_type='void')>,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>& invocation,
VisitArrayType, ThreadToOutputMapType>& invocation,
const ThreadIndicesType& threadIndices)
{
using Invocation =

45
vtkm/exec/internal/testing/TestingTaskTiling.h
View File

@ -76,6 +76,13 @@ struct MyVisitArrayPortal
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct MyThreadToOutputMapPortal
{
using ValueType = vtkm::Id;
VTKM_EXEC_CONT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct TestFetchTagInput
{
};
@ -181,14 +188,16 @@ using InvocationType1 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestExecutionInterface1,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
using InvocationType2 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface2,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
// Not a full worklet, but provides operators that we expect in a worklet.
struct TestWorkletProxy : vtkm::exec::FunctorBase
@ -201,33 +210,42 @@ struct TestWorkletProxy : vtkm::exec::FunctorBase
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutputArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutputArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id3& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id index = vtkm::Dot(threadIndex, vtkm::Id3(1, 8, 64));
return vtkm::exec::arg::ThreadIndicesBasic(
index, outToIn.Get(index), visit.Get(index), globalThreadIndexOffset);
const vtkm::Id flatThreadIndex = vtkm::Dot(threadIndex, vtkm::Id3(1, 8, 64));
const vtkm::Id outIndex = threadToOut.Get(flatThreadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(flatThreadIndex,
outToIn.Get(outIndex),
visit.Get(outIndex),
outIndex,
globalThreadIndexOffset);
}
};
@ -241,33 +259,42 @@ struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutputArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id3& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutputArrayType& threadToOutput,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id index = vtkm::Dot(threadIndex, vtkm::Id3(1, 8, 64));
return vtkm::exec::arg::ThreadIndicesBasic(
index, outToIn.Get(index), visit.Get(index), globalThreadIndexOffset);
const vtkm::Id outputIndex = threadToOutput.Get(index);
return vtkm::exec::arg::ThreadIndicesBasic(index,
outToIn.Get(outputIndex),
visit.Get(outputIndex),
outputIndex,
globalThreadIndexOffset);
}
};

23
vtkm/exec/internal/testing/UnitTestTaskSingular.cxx
View File

@ -64,6 +64,13 @@ struct MyVisitArrayPortal
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct MyThreadToOutputMapPortal
{
using ValueType = vtkm::Id;
VTKM_EXEC_CONT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct TestFetchTagInput
{
};
@ -159,14 +166,16 @@ using InvocationType1 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestExecutionInterface1,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
using InvocationType2 = vtkm::internal::Invocation<ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface2,
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal>;
MyVisitArrayPortal,
MyThreadToOutputMapPortal>;
// Not a full worklet, but provides operators that we expect in a worklet.
struct TestWorkletProxy : vtkm::exec::FunctorBase
@ -180,17 +189,20 @@ struct TestWorkletProxy : vtkm::exec::FunctorBase
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};
@ -205,17 +217,20 @@ struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};

25
vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx
View File

@ -64,6 +64,13 @@ struct MyVisitArrayPortal
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct MyThreadToOutputMapPortal
{
using ValueType = vtkm::Id;
VTKM_EXEC_CONT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct TestFetchTagInput
{
};
@ -163,28 +170,34 @@ struct TestWorkletProxy : vtkm::exec::FunctorBase
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const G& globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};
template <typename Invocation>
void CallDoWorkletInvokeFunctor(const Invocation& invocation, vtkm::Id index)
{
const vtkm::Id outputIndex = invocation.ThreadToOutputMap.Get(index);
vtkm::exec::internal::detail::DoWorkletInvokeFunctor(
TestWorkletProxy(),
invocation,
vtkm::exec::arg::ThreadIndicesBasic(
index, invocation.OutputToInputMap.Get(index), invocation.VisitArray.Get(index)));
vtkm::exec::arg::ThreadIndicesBasic(index,
invocation.OutputToInputMap.Get(outputIndex),
invocation.VisitArray.Get(outputIndex),
outputIndex));
}
void TestDoWorkletInvoke()
@ -204,7 +217,8 @@ void TestDoWorkletInvoke()
TestControlInterface(),
TestExecutionInterface1(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
MyVisitArrayPortal(),
MyThreadToOutputMapPortal()),
1);
VTKM_TEST_ASSERT(inputTestValue == 5, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 100 + 30, "Output value not set right.");
@ -216,7 +230,8 @@ void TestDoWorkletInvoke()
TestControlInterface(),
TestExecutionInterface2(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
MyVisitArrayPortal(),
MyThreadToOutputMapPortal()),
2);
VTKM_TEST_ASSERT(inputTestValue == 6, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 200 + 30 * 2, "Output value not set right.");

2
vtkm/exec/serial/internal/TaskTiling.h
View File

@ -66,6 +66,7 @@ void VTKM_NEVER_EXPORT TaskTiling1DExecute(void* w,
worklet->GetThreadIndices(index,
invocation->OutputToInputMap,
invocation->VisitArray,
invocation->ThreadToOutputMap,
invocation->GetInputDomain(),
globalIndexOffset));
}
@ -110,6 +111,7 @@ void VTKM_NEVER_EXPORT TaskTiling3DExecute(void* w,
worklet->GetThreadIndices(index,
invocation->OutputToInputMap,
invocation->VisitArray,
invocation->ThreadToOutputMap,
invocation->GetInputDomain(),
globalIndexOffset));
}

107
vtkm/internal/Invocation.h
View File

@ -38,7 +38,8 @@ template <typename _ParameterInterface,
typename _ExecutionInterface,
vtkm::IdComponent _InputDomainIndex,
typename _OutputToInputMapType = vtkm::internal::NullType,
typename _VisitArrayType = vtkm::internal::NullType>
typename _VisitArrayType = vtkm::internal::NullType,
typename _ThreadToOutputMapType = vtkm::internal::NullType>
struct Invocation
{
/// \brief The types of the parameters
@ -90,15 +91,25 @@ struct Invocation
///
using VisitArrayType = _VisitArrayType;
/// \brief An array representing the thread to output map.
///
/// When a worklet is invoked, there is an optional mask operation that will
/// prevent the worklet to be run on masked-out elements of the output. This
/// is represented with a map where each thread points to an output it creates.
///
using ThreadToOutputMapType = _ThreadToOutputMapType;
/// \brief Default Invocation constructors that holds the given parameters
/// by reference.
VTKM_CONT
Invocation(const ParameterInterface& parameters,
OutputToInputMapType outputToInputMap = OutputToInputMapType(),
VisitArrayType visitArray = VisitArrayType())
VisitArrayType visitArray = VisitArrayType(),
ThreadToOutputMapType threadToOutputMap = ThreadToOutputMapType())
: Parameters(parameters)
, OutputToInputMap(outputToInputMap)
, VisitArray(visitArray)
, ThreadToOutputMap(threadToOutputMap)
{
}
@ -113,7 +124,8 @@ struct Invocation
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>;
VisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -124,7 +136,7 @@ struct Invocation
const NewParameterInterface& newParameters) const
{
return typename ChangeParametersType<NewParameterInterface>::type(
newParameters, this->OutputToInputMap, this->VisitArray);
newParameters, this->OutputToInputMap, this->VisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -133,8 +145,13 @@ struct Invocation
template <typename NewControlInterface>
struct ChangeControlInterfaceType
{
using type =
Invocation<ParameterInterface, NewControlInterface, ExecutionInterface, InputDomainIndex>;
using type = Invocation<ParameterInterface,
NewControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -145,7 +162,7 @@ struct Invocation
NewControlInterface) const
{
return typename ChangeControlInterfaceType<NewControlInterface>::type(
this->Parameters, this->OutputToInputMap, this->VisitArray);
this->Parameters, this->OutputToInputMap, this->VisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -154,8 +171,13 @@ struct Invocation
template <typename NewExecutionInterface>
struct ChangeExecutionInterfaceType
{
using type =
Invocation<ParameterInterface, NewExecutionInterface, ExecutionInterface, InputDomainIndex>;
using type = Invocation<ParameterInterface,
NewExecutionInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -166,7 +188,7 @@ struct Invocation
NewExecutionInterface) const
{
return typename ChangeExecutionInterfaceType<NewExecutionInterface>::type(
this->Parameters, this->OutputToInputMap, this->VisitArray);
this->Parameters, this->OutputToInputMap, this->VisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -175,8 +197,13 @@ struct Invocation
template <vtkm::IdComponent NewInputDomainIndex>
struct ChangeInputDomainIndexType
{
using type =
Invocation<ParameterInterface, ControlInterface, ExecutionInterface, NewInputDomainIndex>;
using type = Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
NewInputDomainIndex,
OutputToInputMapType,
VisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -187,7 +214,7 @@ struct Invocation
ChangeInputDomainIndex() const
{
return typename ChangeInputDomainIndexType<NewInputDomainIndex>::type(
this->Parameters, this->OutputToInputMap, this->VisitArray);
this->Parameters, this->OutputToInputMap, this->VisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -201,7 +228,8 @@ struct Invocation
ExecutionInterface,
InputDomainIndex,
NewOutputToInputMapType,
VisitArrayType>;
VisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -212,7 +240,7 @@ struct Invocation
ChangeOutputToInputMap(NewOutputToInputMapType newOutputToInputMap) const
{
return typename ChangeOutputToInputMapType<NewOutputToInputMapType>::type(
this->Parameters, newOutputToInputMap, this->VisitArray);
this->Parameters, newOutputToInputMap, this->VisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -226,7 +254,8 @@ struct Invocation
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
NewVisitArrayType>;
NewVisitArrayType,
ThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -237,7 +266,33 @@ struct Invocation
NewVisitArrayType newVisitArray) const
{
return typename ChangeVisitArrayType<NewVisitArrayType>::type(
this->Parameters, this->OutputToInputMap, newVisitArray);
this->Parameters, this->OutputToInputMap, newVisitArray, this->ThreadToOutputMap);
}
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c ThreadToOutputMapType replaced.
///
template <typename NewThreadToOutputMapType>
struct ChangeThreadToOutputMapType
{
using type = Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType,
NewThreadToOutputMapType>;
};
/// Returns a new \c Invocation that is the same as this one except that the
/// \c OutputToInputMap is replaced with that provided.
///
template <typename NewThreadToOutputMapType>
VTKM_CONT typename ChangeThreadToOutputMapType<NewThreadToOutputMapType>::type
ChangeThreadToOutputMap(NewThreadToOutputMapType newThreadToOutputMap) const
{
return typename ChangeThreadToOutputMapType<NewThreadToOutputMapType>::type(
this->Parameters, this->OutputToInputMap, this->VisitArray, newThreadToOutputMap);
}
/// A convenience alias for the input domain type.
@ -268,6 +323,7 @@ struct Invocation
ParameterInterface Parameters;
OutputToInputMapType OutputToInputMap;
VisitArrayType VisitArray;
ThreadToOutputMapType ThreadToOutputMap;
private:
// Do not allow assignment of one Invocation to another. It is too expensive.
@ -276,7 +332,8 @@ private:
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>&) = delete;
VisitArrayType,
ThreadToOutputMapType>&) = delete;
};
/// Convenience function for creating an Invocation object.
@ -286,25 +343,30 @@ template <vtkm::IdComponent InputDomainIndex,
typename ExecutionInterface,
typename ParameterInterface,
typename OutputToInputMapType,
typename VisitArrayType>
typename VisitArrayType,
typename ThreadToOutputMapType>
VTKM_CONT vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>
VisitArrayType,
ThreadToOutputMapType>
make_Invocation(const ParameterInterface& params,
ControlInterface,
ExecutionInterface,
OutputToInputMapType outputToInputMap,
VisitArrayType visitArray)
VisitArrayType visitArray,
ThreadToOutputMapType threadToOutputMap)
{
return vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>(params, outputToInputMap, visitArray);
VisitArrayType,
ThreadToOutputMapType>(
params, outputToInputMap, visitArray, threadToOutputMap);
}
template <vtkm::IdComponent InputDomainIndex,
typename ControlInterface,
@ -320,6 +382,7 @@ VTKM_CONT vtkm::internal::
ControlInterface(),
ExecutionInterface(),
vtkm::internal::NullType(),
vtkm::internal::NullType(),
vtkm::internal::NullType());
}
}

4
vtkm/worklet/CMakeLists.txt
View File

@ -50,7 +50,10 @@ set(headers
Magnitude.h
MarchingCubes.h
Mask.h
MaskIndices.h
MaskNone.h
MaskPoints.h
MaskSelect.h
NDimsEntropy.h
NDimsHistMarginalization.h
NDimsHistogram.h
@ -107,6 +110,7 @@ set(sources_no_device
# be compiled with a device-specific compiler (like CUDA).
set(sources_device
Keys.cxx
MaskSelect.cxx
ScatterCounting.cxx
)

18
vtkm/worklet/DispatcherMapField.h
View File

@ -42,21 +42,9 @@ class DispatcherMapField
using ScatterType = typename Superclass::ScatterType;
public:
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct
// the scatter). By convention, worklets that define a custom scatter type usually provide a
// static method named MakeScatter that constructs a scatter object.
VTKM_CONT
DispatcherMapField(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType())
: Superclass(worklet, scatter)
{
}
VTKM_CONT
DispatcherMapField(const ScatterType& scatter)
: Superclass(WorkletType(), scatter)
template <typename... T>
VTKM_CONT DispatcherMapField(T&&... args)
: Superclass(std::forward<T>(args)...)
{
}

18
vtkm/worklet/DispatcherMapTopology.h
View File

@ -45,21 +45,9 @@ class DispatcherMapTopology
using ScatterType = typename Superclass::ScatterType;
public:
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct
// the scatter). By convention, worklets that define a custom scatter type usually provide a
// static method named MakeScatter that constructs a scatter object.
VTKM_CONT
DispatcherMapTopology(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType())
: Superclass(worklet, scatter)
{
}
VTKM_CONT
DispatcherMapTopology(const ScatterType& scatter)
: Superclass(WorkletType(), scatter)
template <typename... T>
VTKM_CONT DispatcherMapTopology(T&&... args)
: Superclass(std::forward<T>(args)...)
{
}

18
vtkm/worklet/DispatcherPointNeighborhood.h
View File

@ -46,21 +46,9 @@ class DispatcherPointNeighborhood
using ScatterType = typename Superclass::ScatterType;
public:
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct
// the scatter). By convention, worklets that define a custom scatter type usually provide a
// static method named MakeScatter that constructs a scatter object.
VTKM_CONT
DispatcherPointNeighborhood(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType())
: Superclass(worklet, scatter)
{
}
VTKM_CONT
DispatcherPointNeighborhood(const ScatterType& scatter)
: Superclass(WorkletType(), scatter)
template <typename... T>
VTKM_CONT DispatcherPointNeighborhood(T&&... args)
: Superclass(std::forward<T>(args)...)
{
}

18
vtkm/worklet/DispatcherReduceByKey.h
View File

@ -45,21 +45,9 @@ class DispatcherReduceByKey
using ScatterType = typename Superclass::ScatterType;
public:
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct
// the scatter). By convention, worklets that define a custom scatter type usually provide a
// static method named MakeScatter that constructs a scatter object.
VTKM_CONT
DispatcherReduceByKey(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType())
: Superclass(worklet, scatter)
{
}
VTKM_CONT
DispatcherReduceByKey(const ScatterType& scatter)
: Superclass(WorkletType(), scatter)
template <typename... T>
VTKM_CONT DispatcherReduceByKey(T&&... args)
: Superclass(std::forward<T>(args)...)
{
}

27
vtkm/worklet/DispatcherStreamingMapField.h
View File

@ -189,24 +189,12 @@ class DispatcherStreamingMapField
WorkletType,
vtkm::worklet::WorkletMapField>;
using ScatterType = typename Superclass::ScatterType;
using MaskType = typename WorkletType::MaskType;
public:
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct
// the scatter). By convention, worklets that define a custom scatter type usually provide a
// static method named MakeScatter that constructs a scatter object.
VTKM_CONT
DispatcherStreamingMapField(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType())
: Superclass(worklet, scatter)
, NumberOfBlocks(1)
{
}
VTKM_CONT
DispatcherStreamingMapField(const ScatterType& scatter)
: Superclass(WorkletType(), scatter)
template <typename... T>
VTKM_CONT DispatcherStreamingMapField(T&&... args)
: Superclass(std::forward<T>(args)...)
, NumberOfBlocks(1)
{
}
@ -312,11 +300,16 @@ private:
this->Scatter.GetOutputToInputMap(inputRange);
typename ScatterType::VisitArrayType visitArray = this->Scatter.GetVisitArray(inputRange);
// Get the arrays used for masking output elements.
typename MaskType::ThreadToOutputMapType threadToOutputMap =
this->Mask.GetThreadToOutputMap(inputRange);
// Replace the parameters in the invocation with the execution object and
// pass to next step of Invoke. Also add the scatter information.
this->InvokeSchedule(invocation.ChangeParameters(execObjectParameters)
.ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device))
.ChangeVisitArray(visitArray.PrepareForInput(device)),
.ChangeVisitArray(visitArray.PrepareForInput(device))
.ChangeThreadToOutputMap(threadToOutputMap.PrepareForInput(device)),
outputRange,
globalIndexOffset,
device);

4
vtkm/worklet/MarchingCubes.h
View File

@ -597,7 +597,7 @@ public:
//Optimization for structured cellsets so we can call StructuredPointGradient
//and have way faster gradients
vtkm::exec::ConnectivityStructured<Cell, Point, 3> pointGeom(geometry);
vtkm::exec::arg::ThreadIndicesPointNeighborhood tpn(pointId, pointId, 0, pointGeom, 0);
vtkm::exec::arg::ThreadIndicesPointNeighborhood tpn(pointId, pointId, 0, pointId, pointGeom, 0);
const auto& boundary = tpn.GetBoundaryState();
auto pointPortal = pointCoordinates.GetPortal();
@ -681,7 +681,7 @@ public:
//Optimization for structured cellsets so we can call StructuredPointGradient
//and have way faster gradients
vtkm::exec::ConnectivityStructured<Cell, Point, 3> pointGeom(geometry);
vtkm::exec::arg::ThreadIndicesPointNeighborhood tpn(pointId, pointId, 0, pointGeom, 0);
vtkm::exec::arg::ThreadIndicesPointNeighborhood tpn(pointId, pointId, 0, pointId, pointGeom, 0);
const auto& boundary = tpn.GetBoundaryState();
auto pointPortal = pointCoordinates.GetPortal();

93
vtkm/worklet/MaskIndices.h Normal file
View File

@ -0,0 +1,93 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#ifndef vtk_m_worklet_MaskIndices_h
#define vtk_m_worklet_MaskIndices_h
#include <vtkm/cont/Algorithm.h>
namespace vtkm
{
namespace worklet
{
/// \brief Mask using a given array of indices to include in the output.
///
/// \c MaskIndices is a worklet mask object that is used to select elements in the output of a
/// worklet to include in the output. This is done by providing a mask array. This array contains
/// an entry for every output to create. Any output index not included is not generated.
///
/// It is OK to give indices that are out of order, but any index must be provided at most one
/// time. It is an error to have the same index listed twice.
///
class MaskIndices
{
public:
using ThreadToOutputMapType = vtkm::cont::ArrayHandle<vtkm::Id>;
//@{
/// \brief Construct using an index array.
///
/// When you construct a \c MaskSelect with the \c IndexArray tag, you provide an array
/// containing an index for each output to produce. It is OK to give indices that are out of
/// order, but any index must be provided at most one time. It is an error to have the same index
/// listed twice.
///
/// Note that depending on the type of the array passed in, the index may be shallow copied
/// or deep copied into the state of this mask object. Thus, it is a bad idea to alter the
/// array once given to this object.
///
explicit MaskIndices(
const vtkm::cont::ArrayHandle<vtkm::Id>& indexArray,
vtkm::cont::DeviceAdapterId vtkmNotUsed(device) = vtkm::cont::DeviceAdapterTagAny())
{
this->ThreadToOutputMap = indexArray;
}
template <typename T, typename S>
explicit MaskIndices(const vtkm::cont::ArrayHandle<T, S>& indexArray,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
{
vtkm::cont::Algorithm::Copy(device, indexArray, this->ThreadToOutputMap);
}
//@}
// TODO? Create a version that accepts a VariantArrayHandle. Is this needed?
template <typename RangeType>
vtkm::Id GetThreadRange(RangeType vtkmNotUsed(outputRange)) const
{
return this->ThreadToOutputMap.GetNumberOfValues();
}
template <typename RangeType>
ThreadToOutputMapType GetThreadToOutputMap(RangeType vtkmNotUsed(outputRange)) const
{
return this->ThreadToOutputMap;
}
private:
ThreadToOutputMapType ThreadToOutputMap;
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_MaskIndices_h

61
vtkm/worklet/MaskNone.h Normal file
View File

@ -0,0 +1,61 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#ifndef vtk_m_worklet_MaskNone_h
#define vtk_m_worklet_MaskNone_h
#include <vtkm/cont/ArrayHandleIndex.h>
namespace vtkm
{
namespace worklet
{
/// \brief Default mask object that does not suppress anything.
///
/// \c MaskNone is a worklet mask object that does not suppress any items in the output
/// domain. This is the default mask object so that the worklet is run for every possible
/// output element.
///
struct MaskNone
{
template <typename RangeType>
VTKM_CONT RangeType GetThreadRange(RangeType outputRange) const
{
return outputRange;
}
using ThreadToOutputMapType = vtkm::cont::ArrayHandleIndex;
VTKM_CONT ThreadToOutputMapType GetThreadToOutputMap(vtkm::Id outputRange) const
{
return vtkm::cont::ArrayHandleIndex(outputRange);
}
VTKM_CONT ThreadToOutputMapType GetThreadToOutputMap(const vtkm::Id3& outputRange) const
{
return this->GetThreadToOutputMap(outputRange[0] * outputRange[1] * outputRange[2]);
}
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_MaskNone_h

133
vtkm/worklet/MaskSelect.cxx Normal file
View File

@ -0,0 +1,133 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#include <vtkm/worklet/MaskSelect.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/Algorithm.h>
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleView.h>
namespace
{
struct ReverseOutputToThreadMap : vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn outputToThreadMap,
FieldIn maskArray,
WholeArrayOut threadToOutputMap);
using ExecutionSignature = void(_1, InputIndex, _2, _3);
template <typename MaskType, typename ThreadToOutputPortal>
VTKM_EXEC void operator()(vtkm::Id threadIndex,
vtkm::Id outputIndex,
MaskType mask,
ThreadToOutputPortal threadToOutput) const
{
if (mask)
{
threadToOutput.Set(threadIndex, outputIndex);
}
}
};
VTKM_CONT static vtkm::worklet::MaskSelect::ThreadToOutputMapType BuildThreadToOutputMapWithFind(
vtkm::Id numThreads,
vtkm::cont::ArrayHandle<vtkm::Id> outputToThreadMap,
vtkm::cont::DeviceAdapterId device)
{
vtkm::worklet::MaskSelect::ThreadToOutputMapType threadToOutputMap;
vtkm::Id outputSize = outputToThreadMap.GetNumberOfValues();
vtkm::cont::ArrayHandleIndex threadIndices(numThreads);
vtkm::cont::Algorithm::UpperBounds(
device,
vtkm::cont::make_ArrayHandleView(outputToThreadMap, 1, outputSize - 1),
threadIndices,
threadToOutputMap);
return threadToOutputMap;
}
template <typename MaskArrayType>
VTKM_CONT static vtkm::worklet::MaskSelect::ThreadToOutputMapType BuildThreadToOutputMapWithCopy(
vtkm::Id numThreads,
const vtkm::cont::ArrayHandle<vtkm::Id>& outputToThreadMap,
const MaskArrayType& maskArray,
vtkm::cont::DeviceAdapterId device)
{
vtkm::worklet::MaskSelect::ThreadToOutputMapType threadToOutputMap;
threadToOutputMap.Allocate(numThreads);
vtkm::worklet::DispatcherMapField<ReverseOutputToThreadMap> dispatcher;
dispatcher.SetDevice(device);
dispatcher.Invoke(outputToThreadMap, maskArray, threadToOutputMap);
return threadToOutputMap;
}
struct MaskBuilder
{
template <typename ArrayHandleType>
void operator()(const ArrayHandleType& maskArray,
vtkm::worklet::MaskSelect::ThreadToOutputMapType& threadToOutputMap,
vtkm::cont::DeviceAdapterId device)
{
vtkm::cont::ArrayHandle<vtkm::Id> outputToThreadMap;
vtkm::Id numThreads = vtkm::cont::Algorithm::ScanExclusive(
device, vtkm::cont::make_ArrayHandleCast<vtkm::Id>(maskArray), outputToThreadMap);
VTKM_ASSERT(numThreads < maskArray.GetNumberOfValues());
// We have implemented two different ways to compute the thread to output map. The first way is
// to use a binary search on each thread index into the output map. The second way is to
// schedule on each output and copy the the index to the thread map. The first way is faster
// for output sizes that are small relative to the input and also tends to be well load
// balanced. The second way is faster for larger outputs.
//
// The former is obviously faster for one thread and the latter is obviously faster when all
// outputs have a thread. We have to guess for values in the middle. I'm using if the square of
// the number of threads is less than the number of outputs because it is easy to compute.
if ((numThreads * numThreads) < maskArray.GetNumberOfValues())
{
threadToOutputMap = BuildThreadToOutputMapWithFind(numThreads, outputToThreadMap, device);
}
else
{
threadToOutputMap =
BuildThreadToOutputMapWithCopy(numThreads, outputToThreadMap, maskArray, device);
}
}
};
} // anonymous namespace
vtkm::worklet::MaskSelect::ThreadToOutputMapType vtkm::worklet::MaskSelect::Build(
const VariantArrayHandleMask& maskArray,
vtkm::cont::DeviceAdapterId device)
{
vtkm::worklet::MaskSelect::ThreadToOutputMapType threadToOutputMap;
maskArray.CastAndCall(MaskBuilder(), threadToOutputMap, device);
return threadToOutputMap;
}

94
vtkm/worklet/MaskSelect.h Normal file
View File

@ -0,0 +1,94 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#ifndef vtk_m_worklet_MaskSelect_h
#define vtk_m_worklet_MaskSelect_h
#include <vtkm/worklet/vtkm_worklet_export.h>
#include <vtkm/cont/VariantArrayHandle.h>
namespace vtkm
{
namespace worklet
{
/// \brief Mask using arrays to select specific elements to suppress.
///
/// \c MaskSelect is a worklet mask object that is used to select elementsin the output of a
/// worklet to suppress the invocation. That is, the worklet will only be invoked for elements in
/// the output that are not masked out by the given array.
///
/// \c MaskSelect is initialized with a mask array. This array should contain a 0 for any entry
/// that should be masked and a 1 for any output that should be generated. It is an error to have
/// any value that is not a 0 or 1. This method is slower than specifying an index array.
///
class VTKM_WORKLET_EXPORT MaskSelect
{
struct MaskTypes : vtkm::ListTagBase<vtkm::Int32,
vtkm::Int64,
vtkm::UInt32,
vtkm::UInt64,
vtkm::Int8,
vtkm::UInt8,
char>
{
};
using VariantArrayHandleMask = vtkm::cont::VariantArrayHandleBase<MaskTypes>;
public:
using ThreadToOutputMapType = vtkm::cont::ArrayHandle<vtkm::Id>;
MaskSelect(const VariantArrayHandleMask& maskArray,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
{
this->ThreadToOutputMap = this->Build(maskArray, device);
}
template <typename TypeList>
MaskSelect(const vtkm::cont::VariantArrayHandleBase<TypeList>& indexArray,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
{
this->ThreadToOutputMap = this->Build(VariantArrayHandleMask(indexArray), device);
}
template <typename RangeType>
vtkm::Id GetThreadRange(RangeType vtkmNotUsed(outputRange)) const
{
return this->ThreadToOutputMap.GetNumberOfValues();
}
template <typename RangeType>
ThreadToOutputMapType GetThreadToOutputMap(RangeType vtkmNotUsed(outputRange)) const
{
return this->ThreadToOutputMap;
}
private:
ThreadToOutputMapType ThreadToOutputMap;
VTKM_CONT ThreadToOutputMapType Build(const VariantArrayHandleMask& maskArray,
vtkm::cont::DeviceAdapterId device);
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_MaskSelect_h

43
vtkm/worklet/WorkletMapTopology.h
View File

@ -163,20 +163,49 @@ public:
/// Topology map worklets use topology map indices.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename T,
typename OutToInArrayType,
template <typename OutToInArrayType,
typename VisitArrayType,
typename InputDomainType,
typename G>
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType> GetThreadIndices(
const T& threadIndex,
vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType& connectivity,
const G& globalThreadIndexOffset) const
vtkm::Id globalThreadIndexOffset) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType>(threadIndex,
outToIn.Get(outIndex),
visit.Get(outIndex),
outIndex,
connectivity,
globalThreadIndexOffset);
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType> GetThreadIndices(
const vtkm::Id3& threadIndex,
const OutToInArrayType& vtkmNotUsed(outToIn),
const VisitArrayType& vtkmNotUsed(visit),
const ThreadToOutArrayType& vtkmNotUsed(threadToOut),
const InputDomainType& connectivity,
vtkm::Id globalThreadIndexOffset = 0) const
{
using ScatterCheck = std::is_same<ScatterType, vtkm::worklet::ScatterIdentity>;
VTKM_STATIC_ASSERT_MSG(ScatterCheck::value,
"Scheduling on 3D topologies only works with default ScatterIdentity.");
using MaskCheck = std::is_same<MaskType, vtkm::worklet::MaskNone>;
VTKM_STATIC_ASSERT_MSG(MaskCheck::value,
"Scheduling on 3D topologies only works with default MaskNone.");
return vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType>(
threadIndex, outToIn, visit, connectivity, globalThreadIndexOffset);
threadIndex, connectivity, globalThreadIndexOffset);
}
};

41
vtkm/worklet/WorkletPointNeighborhood.h
View File

@ -185,22 +185,51 @@ public:
/// Point neighborhood worklets use the related thread indices class.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename T,
typename IndexType,
typename OutToInArrayType,
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
vtkm::IdComponent Dimension>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesPointNeighborhood GetThreadIndices(
const IndexType& threadIndex,
vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint,
Dimension>& inputDomain, //this should be explicitly
const T& globalThreadIndexOffset = 0) const
vtkm::Id globalThreadIndexOffset = 0) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesPointNeighborhood(threadIndex,
outToIn.Get(outIndex),
visit.Get(outIndex),
outIndex,
inputDomain,
globalThreadIndexOffset);
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesPointNeighborhood GetThreadIndices(
const vtkm::Id3& threadIndex,
const OutToInArrayType& vtkmNotUsed(outToIn),
const VisitArrayType& vtkmNotUsed(visit),
const ThreadToOutArrayType& vtkmNotUsed(threadToOut),
const InputDomainType& connectivity,
vtkm::Id globalThreadIndexOffset = 0) const
{
using ScatterCheck = std::is_same<ScatterType, vtkm::worklet::ScatterIdentity>;
VTKM_STATIC_ASSERT_MSG(ScatterCheck::value,
"Scheduling on 3D topologies only works with default ScatterIdentity.");
using MaskCheck = std::is_same<MaskType, vtkm::worklet::MaskNone>;
VTKM_STATIC_ASSERT_MSG(MaskCheck::value,
"Scheduling on 3D topologies only works with default MaskNone.");
return vtkm::exec::arg::ThreadIndicesPointNeighborhood(
threadIndex, outToIn, visit, inputDomain, globalThreadIndexOffset);
threadIndex, connectivity, globalThreadIndexOffset);
}
};
}

15
vtkm/worklet/WorkletReduceByKey.h
View File

@ -178,20 +178,23 @@ public:
/// Reduce by key worklets use the related thread indices class.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename T,
typename OutToInArrayType,
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesReduceByKey GetThreadIndices(
const T& threadIndex,
vtkm::Id threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType& inputDomain,
const T& globalThreadIndexOffset = 0) const
vtkm::Id globalThreadIndexOffset = 0) const
{
const vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesReduceByKey(threadIndex,
outToIn.Get(threadIndex),
visit.Get(threadIndex),
outToIn.Get(outIndex),
visit.Get(outIndex),
outIndex,
inputDomain,
globalThreadIndexOffset);
}

2
vtkm/worklet/gradient/PointGradient.h
View File

@ -73,7 +73,7 @@ struct PointGradient : public vtkm::worklet::WorkletMapCellToPoint
for (vtkm::IdComponent i = 0; i < numCells; ++i)
{
const vtkm::Id cellId = cellIds[i];
CellThreadIndices cellIndices(cellId, cellId, 0, geometry);
CellThreadIndices cellIndices(cellId, cellId, 0, cellId, geometry);
const CellShapeTag cellShape = cellIndices.GetCellShape();

71
vtkm/worklet/internal/DispatcherBase.h
View File

@ -266,8 +266,9 @@ struct DispatcherBaseTryExecuteFunctor
Invocation& invocation,
const RangeType& dimensions)
{
auto outputRange = self->Scatter.GetOutputRange(dimensions);
self->InvokeTransportParameters(
invocation, dimensions, self->Scatter.GetOutputRange(dimensions), device);
invocation, dimensions, outputRange, self->Mask.GetThreadRange(outputRange), device);
return true;
}
};
@ -659,6 +660,7 @@ public:
//@}
using ScatterType = typename WorkletType::ScatterType;
using MaskType = typename WorkletType::MaskType;
template <typename... Args>
VTKM_CONT void Invoke(Args&&... args) const
@ -670,10 +672,58 @@ public:
}
protected:
// If you get a compile error here about there being no appropriate constructor for ScatterType
// or MapType, then that probably means that the worklet you are trying to execute has defined a
// custom ScatterType or MaskType and that you need to create one (because there is no default
// way to construct the scatter or mask).
VTKM_CONT
DispatcherBase(const WorkletType& worklet, const ScatterType& scatter)
DispatcherBase(const WorkletType& worklet = WorkletType(),
const ScatterType& scatter = ScatterType(),
const MaskType& mask = MaskType())
: Worklet(worklet)
, Scatter(scatter)
, Mask(mask)
, Device(vtkm::cont::DeviceAdapterTagAny())
{
}
// If you get a compile error here about there being no appropriate constructor for MaskType,
// then that probably means that the worklet you are trying to execute has defined a custom
// MaskType and that you need to create one (because there is no default way to construct the
// mask).
VTKM_CONT
DispatcherBase(const ScatterType& scatter, const MaskType& mask = MaskType())
: Worklet(WorkletType())
, Scatter(scatter)
, Mask(mask)
, Device(vtkm::cont::DeviceAdapterTagAny())
{
}
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct the
// scatter).
VTKM_CONT
DispatcherBase(const WorkletType& worklet,
const MaskType& mask,
const ScatterType& scatter = ScatterType())
: Worklet(worklet)
, Scatter(scatter)
, Mask(mask)
, Device(vtkm::cont::DeviceAdapterTagAny())
{
}
// If you get a compile error here about there being no appropriate constructor for ScatterType,
// then that probably means that the worklet you are trying to execute has defined a custom
// ScatterType and that you need to create one (because there is no default way to construct the
// scatter).
VTKM_CONT
DispatcherBase(const MaskType& mask, const ScatterType& scatter = ScatterType())
: Worklet(WorkletType())
, Scatter(scatter)
, Mask(mask)
, Device(vtkm::cont::DeviceAdapterTagAny())
{
}
@ -718,6 +768,7 @@ protected:
WorkletType Worklet;
ScatterType Scatter;
MaskType Mask;
private:
// Dispatchers cannot be copied
@ -729,10 +780,12 @@ private:
template <typename Invocation,
typename InputRangeType,
typename OutputRangeType,
typename ThreadRangeType,
typename DeviceAdapter>
VTKM_CONT void InvokeTransportParameters(Invocation& invocation,
const InputRangeType& inputRange,
OutputRangeType&& outputRange,
ThreadRangeType&& threadRange,
DeviceAdapter device) const
{
// The first step in invoking a worklet is to transport the arguments to
@ -758,17 +811,21 @@ private:
TransportFunctorType(invocation.GetInputDomain(), inputRange, outputRange));
// Get the arrays used for scattering input to output.
typename WorkletType::ScatterType::OutputToInputMapType outputToInputMap =
typename ScatterType::OutputToInputMapType outputToInputMap =
this->Scatter.GetOutputToInputMap(inputRange);
typename WorkletType::ScatterType::VisitArrayType visitArray =
this->Scatter.GetVisitArray(inputRange);
typename ScatterType::VisitArrayType visitArray = this->Scatter.GetVisitArray(inputRange);
// Get the arrays used for masking output elements.
typename MaskType::ThreadToOutputMapType threadToOutputMap =
this->Mask.GetThreadToOutputMap(inputRange);
// Replace the parameters in the invocation with the execution object and
// pass to next step of Invoke. Also add the scatter information.
this->InvokeSchedule(invocation.ChangeParameters(execObjectParameters)
.ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device))
.ChangeVisitArray(visitArray.PrepareForInput(device)),
outputRange,
.ChangeVisitArray(visitArray.PrepareForInput(device))
.ChangeThreadToOutputMap(threadToOutputMap.PrepareForInput(device)),
threadRange,
device);
}

11
vtkm/worklet/internal/WorkletBase.h
View File

@ -46,6 +46,7 @@
#include <vtkm/cont/arg/TypeCheckTagCellSet.h>
#include <vtkm/cont/arg/TypeCheckTagExecObject.h>
#include <vtkm/worklet/MaskNone.h>
#include <vtkm/worklet/ScatterIdentity.h>
namespace vtkm
@ -129,6 +130,11 @@ public:
/// identity scatter (1-to-1 input to output).
using ScatterType = vtkm::worklet::ScatterIdentity;
/// All worklets must define their mask operation. The mask defines which
/// outputs are generated. The default mask is the none mask, which generates
/// everything in the output domain.
using MaskType = vtkm::worklet::MaskNone;
/// \c ControlSignature tag for whole input arrays.
///
/// The \c WholeArrayIn control signature tag specifies an \c ArrayHandle
@ -229,16 +235,19 @@ public:
template <typename T,
typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const T& globalThreadIndexOffset = 0) const
{
vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(threadIndex), visit.Get(threadIndex), globalThreadIndexOffset);
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
};
}

3
vtkm/worklet/testing/CMakeLists.txt
View File

@ -47,7 +47,9 @@ set(unit_tests
UnitTestMagnitude.cxx
UnitTestMarchingCubes.cxx
UnitTestMask.cxx
UnitTestMaskIndices.cxx
UnitTestMaskPoints.cxx
UnitTestMaskSelect.cxx
UnitTestNormalize.cxx
UnitTestNDimsEntropy.cxx
UnitTestNDimsHistogram.cxx
@ -59,6 +61,7 @@ set(unit_tests
UnitTestProbe.cxx
UnitTestRemoveUnusedPoints.cxx
UnitTestScalarsToColors.cxx
UnitTestScatterAndMask.cxx
UnitTestScatterCounting.cxx
UnitTestScatterPermutation.cxx
UnitTestSplatKernels.cxx

140
vtkm/worklet/testing/UnitTestMaskIndices.cxx Normal file
View File

@ -0,0 +1,140 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#include <vtkm/worklet/MaskIndices.h>
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/cont/testing/MakeTestDataSet.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <ctime>
#include <random>
namespace
{
class Worklet : public vtkm::worklet::WorkletMapCellToPoint
{
public:
using ControlSignature = void(CellSetIn cellset, FieldInOutPoint outPointId);
using ExecutionSignature = void(InputIndex, _2);
using InputDomain = _1;
using MaskType = vtkm::worklet::MaskIndices;
VTKM_EXEC void operator()(vtkm::Id pointId, vtkm::Id& outPointId) const { outPointId = pointId; }
};
template <typename CellSetType>
void RunTest(const CellSetType& cellset, const vtkm::cont::ArrayHandle<vtkm::Id>& indices)
{
vtkm::Id numPoints = cellset.GetNumberOfPoints();
vtkm::cont::ArrayHandle<vtkm::Id> outPointId;
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleConstant<vtkm::Id>(-1, numPoints), outPointId);
vtkm::worklet::DispatcherMapTopology<Worklet> dispatcher(vtkm::worklet::MaskIndices{ indices });
dispatcher.Invoke(cellset, outPointId);
vtkm::cont::ArrayHandle<vtkm::Int8> stencil;
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleConstant<vtkm::Int8>(0, numPoints), stencil);
// Check that output that should be written was.
for (vtkm::Id i = 0; i < indices.GetNumberOfValues(); ++i)
{
// All unmasked indices should have been copied to the output.
vtkm::Id unmaskedIndex = indices.GetPortalConstControl().Get(i);
vtkm::Id writtenValue = outPointId.GetPortalConstControl().Get(unmaskedIndex);
VTKM_TEST_ASSERT(unmaskedIndex == writtenValue,
"Did not pass unmasked index. Expected ",
unmaskedIndex,
". Got ",
writtenValue);
// Mark index as passed.
stencil.GetPortalControl().Set(unmaskedIndex, 1);
}
// Check that output that should not be written was not.
for (vtkm::Id i = 0; i < numPoints; ++i)
{
if (stencil.GetPortalConstControl().Get(i) == 0)
{
vtkm::Id foundValue = outPointId.GetPortalConstControl().Get(i);
VTKM_TEST_ASSERT(foundValue == -1,
"Expected index ",
i,
" to be unwritten but was filled with ",
foundValue);
}
}
}
void TestMaskIndices()
{
vtkm::cont::DataSet dataset = vtkm::cont::testing::MakeTestDataSet().Make2DUniformDataSet0();
auto cellset = dataset.GetCellSet();
vtkm::Id numberOfPoints = cellset.GetNumberOfPoints();
vtkm::UInt32 seed = static_cast<vtkm::UInt32>(std::time(nullptr));
std::default_random_engine generator;
generator.seed(seed);
std::uniform_int_distribution<vtkm::Id> countDistribution(1, 2 * numberOfPoints);
std::uniform_int_distribution<vtkm::Id> ptidDistribution(0, numberOfPoints - 1);
const int iterations = 5;
std::cout << "Testing with random indices " << iterations << " times\n";
std::cout << "Seed: " << seed << std::endl;
for (int i = 1; i <= iterations; ++i)
{
std::cout << "iteration: " << i << "\n";
vtkm::Id count = countDistribution(generator);
vtkm::cont::ArrayHandle<vtkm::Id> indices;
indices.Allocate(count);
// Note that it is possible that the same index will be written twice, which is generally
// a bad idea with MaskIndices. However, the worklet will write the same value for each
// instance, so we should still get the correct result.
auto portal = indices.GetPortalControl();
std::cout << "using indices:";
for (vtkm::Id j = 0; j < count; ++j)
{
auto val = ptidDistribution(generator);
std::cout << " " << val;
portal.Set(j, val);
}
std::cout << "\n";
RunTest(cellset, indices);
}
}
} // anonymous namespace
int UnitTestMaskIndices(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestMaskIndices, argc, argv);
}

215
vtkm/worklet/testing/UnitTestMaskSelect.cxx Normal file
View File

@ -0,0 +1,215 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#include <vtkm/worklet/MaskSelect.h>
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/testing/Testing.h>
#include <vector>
namespace
{
constexpr vtkm::Id nullValue = -2;
struct TestMaskArrays
{
vtkm::cont::ArrayHandle<vtkm::IdComponent> SelectArray;
vtkm::cont::ArrayHandle<vtkm::Id> ThreadToOutputMap;
};
TestMaskArrays MakeMaskArraysShort()
{
const vtkm::Id selectArraySize = 18;
const vtkm::IdComponent selectArray[selectArraySize] = { 1, 1, 0, 0, 0, 0, 1, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 1 };
const vtkm::Id threadRange = 4;
const vtkm::Id threadToOutputMap[threadRange] = { 0, 1, 6, 17 };
TestMaskArrays arrays;
// Need to copy arrays so that the data does not go out of scope.
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandle(selectArray, selectArraySize),
arrays.SelectArray);
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandle(threadToOutputMap, threadRange),
arrays.ThreadToOutputMap);
return arrays;
}
TestMaskArrays MakeMaskArraysLong()
{
const vtkm::Id selectArraySize = 16;
const vtkm::IdComponent selectArray[selectArraySize] = {
1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1,
};
const vtkm::Id threadRange = 15;
const vtkm::Id threadToOutputMap[threadRange] = {
0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 11, 12, 13, 14, 15
};
TestMaskArrays arrays;
// Need to copy arrays so that the data does not go out of scope.
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandle(selectArray, selectArraySize),
arrays.SelectArray);
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandle(threadToOutputMap, threadRange),
arrays.ThreadToOutputMap);
return arrays;
}
TestMaskArrays MakeMaskArraysZero()
{
const vtkm::Id selectArraySize = 6;
const vtkm::IdComponent selectArray[selectArraySize] = { 0, 0, 0, 0, 0, 0 };
TestMaskArrays arrays;
// Need to copy arrays so that the data does not go out of scope.
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandle(selectArray, selectArraySize),
arrays.SelectArray);
arrays.ThreadToOutputMap.Allocate(0);
return arrays;
}
struct TestMaskSelectWorklet : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn inputIndices, FieldInOut copyIndices);
using ExecutionSignature = void(_1, _2);
using MaskType = vtkm::worklet::MaskSelect;
VTKM_EXEC
void operator()(vtkm::Id inputIndex, vtkm::Id& indexCopy) const { indexCopy = inputIndex; }
};
template <typename T, typename SelectArrayType>
void CompareArrays(vtkm::cont::ArrayHandle<T> array1,
vtkm::cont::ArrayHandle<T> array2,
const SelectArrayType& selectArray)
{
VTKM_IS_ARRAY_HANDLE(SelectArrayType);
auto portal1 = array1.GetPortalConstControl();
auto portal2 = array2.GetPortalConstControl();
auto selectPortal = selectArray.GetPortalConstControl();
VTKM_TEST_ASSERT(portal1.GetNumberOfValues() == portal2.GetNumberOfValues());
VTKM_TEST_ASSERT(portal1.GetNumberOfValues() == selectArray.GetNumberOfValues());
for (vtkm::Id index = 0; index < portal1.GetNumberOfValues(); index++)
{
if (selectPortal.Get(index))
{
T value1 = portal1.Get(index);
T value2 = portal2.Get(index);
VTKM_TEST_ASSERT(
value1 == value2, "Array values not equal (", index, ": ", value1, " ", value2, ")");
}
else
{
T value = portal2.Get(index);
VTKM_TEST_ASSERT(value == nullValue, "Expected null value, got ", value);
}
}
}
template <typename T>
void CompareArrays(vtkm::cont::ArrayHandle<T> array1, vtkm::cont::ArrayHandle<T> array2)
{
CompareArrays(
array1, array2, vtkm::cont::make_ArrayHandleConstant<bool>(true, array1.GetNumberOfValues()));
}
// This unit test makes sure the ScatterCounting generates the correct map
// and visit arrays.
void TestMaskArrayGeneration(const TestMaskArrays& arrays)
{
std::cout << " Testing array generation" << std::endl;
vtkm::worklet::MaskSelect mask(arrays.SelectArray, vtkm::cont::DeviceAdapterTagAny());
vtkm::Id inputSize = arrays.SelectArray.GetNumberOfValues();
std::cout << " Checking thread to output map ";
vtkm::cont::printSummary_ArrayHandle(mask.GetThreadToOutputMap(inputSize), std::cout);
CompareArrays(arrays.ThreadToOutputMap, mask.GetThreadToOutputMap(inputSize));
}
// This is more of an integration test that makes sure the scatter works with a
// worklet invocation.
void TestMaskWorklet(const TestMaskArrays& arrays)
{
std::cout << " Testing mask select in a worklet." << std::endl;
vtkm::worklet::DispatcherMapField<TestMaskSelectWorklet> dispatcher(
vtkm::worklet::MaskSelect(arrays.SelectArray));
vtkm::Id inputSize = arrays.SelectArray.GetNumberOfValues();
vtkm::cont::ArrayHandle<vtkm::Id> inputIndices;
vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleIndex(inputSize), inputIndices);
vtkm::cont::ArrayHandle<vtkm::Id> selectIndexCopy;
vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleConstant<vtkm::Id>(nullValue, inputSize),
selectIndexCopy);
std::cout << " Invoke worklet" << std::endl;
dispatcher.Invoke(inputIndices, selectIndexCopy);
std::cout << " Check copied indices." << std::endl;
CompareArrays(inputIndices, selectIndexCopy, arrays.SelectArray);
}
void TestMaskSelectWithArrays(const TestMaskArrays& arrays)
{
TestMaskArrayGeneration(arrays);
TestMaskWorklet(arrays);
}
void TestMaskSelect()
{
std::cout << "Testing arrays with output smaller than input." << std::endl;
TestMaskSelectWithArrays(MakeMaskArraysShort());
std::cout << "Testing arrays with output larger than input." << std::endl;
TestMaskSelectWithArrays(MakeMaskArraysLong());
std::cout << "Testing arrays with zero output." << std::endl;
TestMaskSelectWithArrays(MakeMaskArraysZero());
}
} // anonymous namespace
int UnitTestMaskSelect(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestMaskSelect, argc, argv);
}

193
vtkm/worklet/testing/UnitTestScatterAndMask.cxx Normal file
View File

@ -0,0 +1,193 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2019 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2019 UT-Battelle, LLC.
// Copyright 2019 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/DispatcherPointNeighborhood.h>
#include <vtkm/worklet/MaskIndices.h>
#include <vtkm/worklet/ScatterUniform.h>
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/Math.h>
#include <vtkm/cont/testing/MakeTestDataSet.h>
#include <vtkm/cont/testing/Testing.h>
namespace
{
using FieldType = vtkm::Float32;
#define FieldNull vtkm::Nan32()
constexpr vtkm::IdComponent IdNull = -2;
struct FieldWorklet : vtkm::worklet::WorkletMapField
{
using ControlSignature = void(WholeCellSetIn<>, // Placeholder for interface consistency
FieldIn inputField,
FieldInOut fieldCopy,
FieldInOut visitCopy);
using ExecutionSignature = void(_2, VisitIndex, _3, _4);
using InputDomain = _2;
using ScatterType = vtkm::worklet::ScatterUniform<2>;
using MaskType = vtkm::worklet::MaskIndices;
VTKM_EXEC void operator()(FieldType inField,
vtkm::IdComponent visitIndex,
FieldType& fieldCopy,
vtkm::IdComponent& visitCopy) const
{
fieldCopy = inField;
visitCopy = visitIndex;
}
};
struct TopologyWorklet : vtkm::worklet::WorkletMapCellToPoint
{
using ControlSignature = void(CellSetIn,
FieldInPoint inputField,
FieldInOutPoint fieldCopy,
FieldInOutPoint visitCopy);
using ExecutionSignature = void(_2, VisitIndex, _3, _4);
using InputDomain = _1;
using ScatterType = vtkm::worklet::ScatterUniform<2>;
using MaskType = vtkm::worklet::MaskIndices;
VTKM_EXEC void operator()(FieldType inField,
vtkm::IdComponent visitIndex,
FieldType& fieldCopy,
vtkm::IdComponent& visitCopy) const
{
fieldCopy = inField;
visitCopy = visitIndex;
}
};
struct NeighborhoodWorklet : vtkm::worklet::WorkletPointNeighborhood
{
using ControlSignature = void(CellSetIn,
FieldIn inputField,
FieldInOut fieldCopy,
FieldInOut visitCopy);
using ExecutionSignature = void(_2, VisitIndex, _3, _4);
using InputDomain = _1;
using ScatterType = vtkm::worklet::ScatterUniform<2>;
using MaskType = vtkm::worklet::MaskIndices;
VTKM_EXEC void operator()(FieldType inField,
vtkm::IdComponent visitIndex,
FieldType& fieldCopy,
vtkm::IdComponent& visitCopy) const
{
fieldCopy = inField;
visitCopy = visitIndex;
}
};
template <typename DispatcherType>
void TestMapWorklet()
{
vtkm::cont::testing::MakeTestDataSet builder;
vtkm::cont::DataSet data = builder.Make3DUniformDataSet1();
vtkm::cont::CellSetStructured<3> cellSet =
data.GetCellSet().Cast<vtkm::cont::CellSetStructured<3>>();
vtkm::Id numPoints = cellSet.GetNumberOfPoints();
vtkm::cont::ArrayHandle<FieldType> inField;
inField.Allocate(numPoints);
SetPortal(inField.GetPortalControl());
vtkm::cont::ArrayHandle<FieldType> fieldCopy;
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleConstant(FieldNull, numPoints * 2), fieldCopy);
vtkm::cont::ArrayHandle<vtkm::IdComponent> visitCopy;
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleConstant(IdNull, numPoints * 2), visitCopy);
// The scatter is hardcoded to create 2 outputs for every input.
// Set up the mask to select a range of values in the middle.
vtkm::Id maskStart = numPoints / 2;
vtkm::Id maskEnd = (numPoints * 2) / 3;
vtkm::worklet::MaskIndices mask(
vtkm::cont::make_ArrayHandleCounting(maskStart, vtkm::Id(1), maskEnd - maskStart));
DispatcherType dispatcher(mask);
dispatcher.Invoke(cellSet, inField, fieldCopy, visitCopy);
// Check outputs
auto fieldCopyPortal = fieldCopy.GetPortalConstControl();
auto visitCopyPortal = visitCopy.GetPortalConstControl();
for (vtkm::Id outputIndex = 0; outputIndex < numPoints * 2; ++outputIndex)
{
FieldType fieldValue = fieldCopyPortal.Get(outputIndex);
vtkm::IdComponent visitValue = visitCopyPortal.Get(outputIndex);
if ((outputIndex >= maskStart) && (outputIndex < maskEnd))
{
vtkm::Id inputIndex = outputIndex / 2;
FieldType expectedField = TestValue(inputIndex, FieldType());
VTKM_TEST_ASSERT(fieldValue == expectedField,
outputIndex,
": expected ",
expectedField,
", got ",
fieldValue);
vtkm::IdComponent expectedVisit = static_cast<vtkm::IdComponent>(outputIndex % 2);
VTKM_TEST_ASSERT(visitValue == expectedVisit,
outputIndex,
": expected ",
expectedVisit,
", got ",
visitValue);
}
else
{
VTKM_TEST_ASSERT(vtkm::IsNan(fieldValue), outputIndex, ": expected NaN, got ", fieldValue);
VTKM_TEST_ASSERT(
visitValue == IdNull, outputIndex, ": expected ", IdNull, ", got ", visitValue);
}
}
}
void Test()
{
std::cout << "Try on WorkletMapField" << std::endl;
TestMapWorklet<vtkm::worklet::DispatcherMapField<FieldWorklet>>();
std::cout << "Try on WorkletMapCellToPoint" << std::endl;
TestMapWorklet<vtkm::worklet::DispatcherMapTopology<TopologyWorklet>>();
std::cout << "Try on WorkletPointNeighborhood" << std::endl;
TestMapWorklet<vtkm::worklet::DispatcherPointNeighborhood<NeighborhoodWorklet>>();
}
} // anonymous namespace
int UnitTestScatterAndMask(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(Test, argc, argv);
}