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

Merge branch 'scatter-worklets' into 'master'

Browse Source 
Scatter in worklets

Add the functionality to perform a scatter operation from input to output in a worklet invocation. This allows you to, for example, specify a variable amount of outputs generated for each input.

See merge request !221
This commit is contained in:
Kenneth Moreland 2015-11-11 13:09:47 -05:00
commit 1a538ca196
50 changed files with 2643 additions and 1217 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

41
examples/isosurface/IsosurfaceUniformGrid.cxx
View File

@ -69,11 +69,11 @@ public:
typedef _1 InputDomain;
const vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::FloatDefault xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::Id cellsPerLayer;
VTKM_CONT_EXPORT
TangleField(const vtkm::Id3 dims, const float mins[3], const float maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
TangleField(const vtkm::Id3 dims, const vtkm::FloatDefault mins[3], const vtkm::FloatDefault maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { };
VTKM_EXEC_EXPORT
@ -83,9 +83,9 @@ public:
const vtkm::Id y = (vertexId / (xdim)) % (ydim);
const vtkm::Id z = vertexId / cellsPerLayer;
const float fx = static_cast<float>(x) / static_cast<float>(xdim-1);
const float fy = static_cast<float>(y) / static_cast<float>(xdim-1);
const float fz = static_cast<float>(z) / static_cast<float>(xdim-1);
const vtkm::FloatDefault fx = static_cast<vtkm::FloatDefault>(x) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fy = static_cast<vtkm::FloatDefault>(y) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fz = static_cast<vtkm::FloatDefault>(z) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::Float32 xx = 3.0f*(xmin+(xmax-xmin)*(fx));
const vtkm::Float32 yy = 3.0f*(ymin+(ymax-ymin)*(fy));
@ -103,15 +103,22 @@ vtkm::cont::DataSet MakeIsosurfaceTestDataSet(vtkm::Id3 dims)
const vtkm::Id3 vdims(dims[0] + 1, dims[1] + 1, dims[2] + 1);
float mins[3] = {-1.0f, -1.0f, -1.0f};
float maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::FloatDefault mins[3] = {-1.0f, -1.0f, -1.0f};
vtkm::FloatDefault maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray;
vtkm::cont::ArrayHandleCounting<vtkm::Id> vertexCountImplicitArray(0, 1, vdims[0]*vdims[1]*vdims[2]);
vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher(TangleField(vdims, mins, maxs));
tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray);
vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims);
vtkm::Vec<vtkm::FloatDefault,3> origin(0.0f, 0.0f, 0.0f);
vtkm::Vec<vtkm::FloatDefault,3> spacing(
1.0f/static_cast<vtkm::FloatDefault>(dims[0]),
1.0f/static_cast<vtkm::FloatDefault>(dims[2]),
1.0f/static_cast<vtkm::FloatDefault>(dims[1]));
vtkm::cont::ArrayHandleUniformPointCoordinates
coordinates(vdims, origin, spacing);
dataSet.AddCoordinateSystem(
vtkm::cont::CoordinateSystem("coordinates", 1, coordinates));
@ -135,9 +142,9 @@ void initializeGL()
glEnable(GL_DEPTH_TEST);
glShadeModel(GL_SMOOTH);
float white[] = { 0.8f, 0.8f, 0.8f, 1.0f };
float black[] = { 0.0f, 0.0f, 0.0f, 1.0f };
float lightPos[] = { 10.0f, 10.0f, 10.5f, 1.0f };
vtkm::FloatDefault white[] = { 0.8f, 0.8f, 0.8f, 1.0f };
vtkm::FloatDefault black[] = { 0.0f, 0.0f, 0.0f, 1.0f };
vtkm::FloatDefault lightPos[] = { 10.0f, 10.0f, 10.5f, 1.0f };
glLightfv(GL_LIGHT0, GL_AMBIENT, white);
glLightfv(GL_LIGHT0, GL_DIFFUSE, white);
@ -199,7 +206,7 @@ void mouseMove(int x, int y)
if (mouse_state == 0)
{
vtkm::Float32 pideg = static_cast<float>(vtkm::Pi()/180.0);
vtkm::Float32 pideg = static_cast<vtkm::Float32>(vtkm::Pi_2());
Quaternion newRotX;
newRotX.setEulerAngles(-0.2f*dx*pideg/180.0f, 0.0f, 0.0f);
qrot.mul(newRotX);
@ -243,6 +250,16 @@ int main(int argc, char* argv[])
std::cout << "Number of output vertices: " << verticesArray.GetNumberOfValues() << std::endl;
std::cout << "vertices: ";
vtkm::cont::printSummary_ArrayHandle(verticesArray, std::cout);
std::cout << std::endl;
std::cout << "normals: ";
vtkm::cont::printSummary_ArrayHandle(normalsArray, std::cout);
std::cout << std::endl;
std::cout << "scalars: ";
vtkm::cont::printSummary_ArrayHandle(scalarsArray, std::cout);
std::cout << std::endl;
lastx = lasty = 0;
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE | GLUT_DEPTH);

6
vtkm/Types.h
View File

@ -175,6 +175,12 @@ namespace internal {
//-----------------------------------------------------------------------------
/// Placeholder class for when a type is not applicable.
///
struct NullType { };
//-----------------------------------------------------------------------------
template<vtkm::IdComponent Size>
struct VecEquals
{

14
vtkm/cont/ArrayHandleGroupVec.h
View File

@ -345,6 +345,20 @@ public:
: Superclass(StorageType(sourceArray)) { }
};
/// \c make_ArrayHandleGroupVec is convenience function to generate an
/// ArrayHandleGroupVec. It takes in an ArrayHandle and the number of components
/// (as a specified template parameter), and returns an array handle with
/// consecutive entries grouped in a Vec.
///
template<vtkm::IdComponent NUM_COMPONENTS,
typename ArrayHandleType>
VTKM_CONT_EXPORT
vtkm::cont::ArrayHandleGroupVec<ArrayHandleType, NUM_COMPONENTS>
make_ArrayHandleGroupVec(const ArrayHandleType &array)
{
return vtkm::cont::ArrayHandleGroupVec<ArrayHandleType,NUM_COMPONENTS>(array);
}
}
} // namespace vtkm::cont

2
vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h
View File

@ -74,7 +74,7 @@ public:
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< const T > PortalConstType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< T > PortalConstType;
VTKM_CONT_EXPORT
ArrayManagerExecutionThrustDevice(StorageType *storage)

18
vtkm/exec/CellDerivative.h
View File

@ -181,15 +181,15 @@ PermutePyramidToHex(const FieldVecType &field)
//
#define VTKM_ACCUM_JACOBIAN_3D(pointIndex, weight0, weight1, weight2) \
jacobian(0,0) += wCoords[pointIndex][0] * (weight0); \
jacobian(1,0) += wCoords[pointIndex][1] * (weight0); \
jacobian(2,0) += wCoords[pointIndex][2] * (weight0); \
jacobian(0,1) += wCoords[pointIndex][0] * (weight1); \
jacobian(1,1) += wCoords[pointIndex][1] * (weight1); \
jacobian(2,1) += wCoords[pointIndex][2] * (weight1); \
jacobian(0,2) += wCoords[pointIndex][0] * (weight2); \
jacobian(1,2) += wCoords[pointIndex][1] * (weight2); \
jacobian(2,2) += wCoords[pointIndex][2] * (weight2)
jacobian(0,0) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight0)); \
jacobian(1,0) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight0)); \
jacobian(2,0) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight0)); \
jacobian(0,1) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight1)); \
jacobian(1,1) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight1)); \
jacobian(2,1) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight1)); \
jacobian(0,2) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight2)); \
jacobian(1,2) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight2)); \
jacobian(2,2) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight2))
template<typename WorldCoordType,
typename ParametricCoordType,

2
vtkm/exec/arg/CMakeLists.txt
View File

@ -32,8 +32,10 @@ set(headers
FetchTagTopologyIn.h
FromCount.h
FromIndices.h
ThreadIndices.h
ThreadIndicesBasic.h
ThreadIndicesTopologyMap.h
VisitIndex.h
WorkIndex.h
)

2
vtkm/exec/arg/FetchTagArrayDirectIn.h
View File

@ -50,7 +50,7 @@ struct Fetch<
ValueType Load(const ThreadIndicesType &indices,
const ExecObjectType &arrayPortal) const
{
return arrayPortal.Get(indices.GetIndex());
return arrayPortal.Get(indices.GetInputIndex());
}
VTKM_EXEC_EXPORT

13
vtkm/exec/arg/FetchTagArrayDirectInOut.h
View File

@ -34,6 +34,15 @@ namespace arg {
/// indexing, so the thread index given to \c Store is used as the index into
/// the array.
///
/// When using \c FetchTagArrayDirectInOut with a worklet invocation with a
/// scatter, it is a bit undefined how the in/out array should be indexed.
/// Should it be the size of the input arrays and written back there, or
/// should it be the size of the output arrays and pre-filled with the output.
/// The implementation indexes based on the output because it is safer. The
/// output will have a unique index for each worklet instance, so you don't
/// have to worry about writes stomping on each other (which they would
/// inevitably do if index as input).
///
struct FetchTagArrayDirectInOut { };
@ -51,7 +60,7 @@ struct Fetch<
ValueType Load(const ThreadIndicesType &indices,
const ExecObjectType &arrayPortal) const
{
return arrayPortal.Get(indices.GetIndex());
return arrayPortal.Get(indices.GetOutputIndex());
}
VTKM_SUPPRESS_EXEC_WARNINGS
@ -60,7 +69,7 @@ struct Fetch<
const ExecObjectType &arrayPortal,
const ValueType &value) const
{
arrayPortal.Set(indices.GetIndex(), value);
arrayPortal.Set(indices.GetOutputIndex(), value);
}
};

2
vtkm/exec/arg/FetchTagArrayDirectOut.h
View File

@ -59,7 +59,7 @@ struct Fetch<
const ExecObjectType &arrayPortal,
const ValueType &value) const
{
arrayPortal.Set(indices.GetIndex(), value);
arrayPortal.Set(indices.GetOutputIndex(), value);
}
};

6
vtkm/exec/arg/FetchTagArrayTopologyMapIn.h
View File

@ -76,7 +76,11 @@ struct FetchArrayTopologyMapInImplementation
static ValueType Load(const ThreadIndicesType &indices,
const FieldExecObjectType &field)
{
return ValueType(indices.GetIndicesFrom(), field);
// It is important that we give the VecFromPortalPermute (ValueType) a
// pointer that will stay around during the time the Vec is valid. Thus, we
// should make sure that indices is a reference that goes up the stack at
// least as far as the returned VecFromPortalPermute is used.
return ValueType(indices.GetIndicesFromPointer(), field);
}
};

85
vtkm/exec/arg/ThreadIndices.h Normal file
View File

@ -0,0 +1,85 @@
//============================================================================
// 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 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_exec_arg_ThreadIndices_h
#define vtk_m_exec_arg_ThreadIndices_h
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Aspect tag to use for getting the thread indices.
///
/// The \c AspectTagThreadIndices aspect tag causes the \c Fetch class to
/// ignore whatever data is in the associated execution object and return the
/// thread indices.
///
struct AspectTagThreadIndices { };
/// \brief The \c ExecutionSignature tag to use to get the thread indices
///
/// When a worklet is dispatched, it broken into pieces defined by the input
/// domain and scheduled on independent threads. During this process multiple
/// indices associated with the input and output can be generated. This tag in
/// the \c ExecutionSignature passes the index for this work. \c WorkletBase
/// contains a typedef that points to this class.
///
struct ThreadIndices : vtkm::exec::arg::ExecutionSignatureTagBase
{
// The index does not really matter because the fetch is going to ignore it.
// However, it still has to point to a valid parameter in the
// ControlSignature because the templating is going to grab a fetch tag
// whether we use it or not. 1 should be guaranteed to be valid since you
// need at least one argument for the input domain.
static const vtkm::IdComponent INDEX = 1;
typedef vtkm::exec::arg::AspectTagThreadIndices AspectTag;
};
template<typename FetchTag, typename ThreadIndicesType, typename ExecObjectType>
struct Fetch<FetchTag,
vtkm::exec::arg::AspectTagThreadIndices,
ThreadIndicesType,
ExecObjectType>
{
typedef const ThreadIndicesType &ValueType;
VTKM_EXEC_EXPORT
const ThreadIndicesType &
Load(const ThreadIndicesType &indices, const ExecObjectType &) const
{
return indices;
}
VTKM_EXEC_EXPORT
void Store(const ThreadIndicesType &,
const ExecObjectType &,
const ThreadIndicesType &) const
{
// Store is a no-op.
}
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_ThreadIndices_h

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

@ -44,8 +44,10 @@ class ThreadIndicesBasic
public:
template<typename Invocation>
VTKM_EXEC_EXPORT
ThreadIndicesBasic(vtkm::Id threadIndex, const Invocation &)
: Index(threadIndex)
ThreadIndicesBasic(vtkm::Id threadIndex, const Invocation &invocation)
: InputIndex(invocation.OutputToInputMap.Get(threadIndex)),
OutputIndex(threadIndex),
VisitIndex(invocation.VisitArray.Get(threadIndex))
{
}
@ -53,10 +55,10 @@ public:
///
/// This index refers to the input element (array value, cell, etc.) that
/// this thread is being invoked for. This is the typical index used during
/// fetches.
/// Fetch::Load.
///
VTKM_EXEC_EXPORT
vtkm::Id GetIndex() const { return this->Index; }
vtkm::Id GetInputIndex() const { return this->InputIndex; }
/// \brief The 3D index into the input domain.
///
@ -67,10 +69,32 @@ public:
/// first component with the remaining components set to 0.
///
VTKM_EXEC_EXPORT
vtkm::Id3 GetIndex3D() const { return vtkm::Id3(this->GetIndex(), 0, 0); }
vtkm::Id3 GetInputIndex3D() const
{
return vtkm::Id3(this->GetInputIndex(), 0, 0);
}
/// \brief The index into the output domain.
///
/// This index refers to the output element (array value, cell, etc.) that
/// this thread is creating. This is the typical index used during
/// Fetch::Store.
///
VTKM_EXEC_EXPORT
vtkm::Id GetOutputIndex() const { return this->OutputIndex; }
/// \brief The visit index.
///
/// When multiple output indices have the same input index, they are
/// distinguished using the visit index.
///
VTKM_EXEC_EXPORT
vtkm::IdComponent GetVisitIndex() const { return this->VisitIndex; }
private:
vtkm::Id Index;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
};
}

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

@ -85,8 +85,8 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
this->IndicesFrom = connectivity.GetIndices(this->GetIndex());
this->CellShape = connectivity.GetCellShape(this->GetIndex());
this->IndicesFrom = connectivity.GetIndices(this->GetInputIndex());
this->CellShape = connectivity.GetCellShape(this->GetInputIndex());
}
/// \brief The input indices of the "from" elements.
@ -97,7 +97,22 @@ public:
/// containing the indices to the "from" elements.
///
VTKM_EXEC_EXPORT
IndicesFromType GetIndicesFrom() const { return this->IndicesFrom; }
const IndicesFromType &GetIndicesFrom() const { return this->IndicesFrom; }
/// \brief The input indices of the "from" elements in pointer form.
///
/// Returns the same object as GetIndicesFrom except that it returns a
/// pointer to the internally held object rather than a reference or copy.
/// Since the from indices can be a sizeable Vec (8 entries is common), it is
/// best not to have a bunch a copies. Thus, you can pass around a pointer
/// instead. However, care should be taken to make sure that this object does
/// not go out of scope, at which time the returned pointer becomes invalid.
///
VTKM_EXEC_EXPORT
const IndicesFromType *GetIndicesFromPointer() const
{
return &this->IndicesFrom;
}
/// \brief The shape of the input cell.
///
@ -179,12 +194,12 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
const LogicalIndexType logicalIndex = connectivity.FlatToLogicalToIndex(threadIndex);
this->Index = threadIndex;
this->LogicalIndex = logicalIndex;
this->IndicesFrom = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(threadIndex);
this->InputIndex = invocation.OutputToInputMap.Get(threadIndex);
this->OutputIndex = threadIndex;
this->VisitIndex = invocation.VisitArray.Get(threadIndex);
this->LogicalIndex = connectivity.FlatToLogicalToIndex(this->InputIndex);
this->IndicesFrom = connectivity.GetIndices(this->LogicalIndex);
this->CellShape = connectivity.GetCellShape(this->InputIndex);
}
template<typename Invocation>
@ -198,11 +213,14 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const LogicalIndexType logicalIndex =
detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
this->Index = index;
// We currently only support multidimensional indices on one-to-one input-
// to-output mappings. (We don't have a use case otherwise.)
this->InputIndex = this->OutputIndex = index;
this->VisitIndex = invocation.VisitArray.Get(index);
this->LogicalIndex = logicalIndex;
this->IndicesFrom = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(index);
@ -226,9 +244,9 @@ public:
/// fetches.
///
VTKM_EXEC_EXPORT
vtkm::Id GetIndex() const
vtkm::Id GetInputIndex() const
{
return this->Index;
return this->InputIndex;
}
/// \brief The 3D index into the input domain.
@ -237,11 +255,34 @@ public:
/// for the input.
///
VTKM_EXEC_EXPORT
vtkm::Id3 GetIndex3D() const
vtkm::Id3 GetInputIndex3D() const
{
return detail::InflateTo3D(this->GetIndexLogical());
}
/// \brief The index into the output domain.
///
/// This index refers to the output element (array value, cell, etc.) that
/// this thread is creating. This is the typical index used during
/// Fetch::Store.
///
VTKM_EXEC_EXPORT
vtkm::Id GetOutputIndex() const
{
return this->OutputIndex;
}
/// \brief The visit index.
///
/// When multiple output indices have the same input index, they are
/// distinguished using the visit index.
///
VTKM_EXEC_EXPORT
vtkm::IdComponent GetVisitIndex() const
{
return this->VisitIndex;
}
/// \brief The input indices of the "from" elements.
///
/// A topology map has "from" and "to" elements (for example from points to
@ -250,7 +291,22 @@ public:
/// containing the indices to the "from" elements.
///
VTKM_EXEC_EXPORT
IndicesFromType GetIndicesFrom() const { return this->IndicesFrom; }
const IndicesFromType &GetIndicesFrom() const { return this->IndicesFrom; }
/// \brief The input indices of the "from" elements in pointer form.
///
/// Returns the same object as GetIndicesFrom except that it returns a
/// pointer to the internally held object rather than a reference or copy.
/// Since the from indices can be a sizeable Vec (8 entries is common), it is
/// best not to have a bunch a copies. Thus, you can pass around a pointer
/// instead. However, care should be taken to make sure that this object does
/// not go out of scope, at which time the returned pointer becomes invalid.
///
VTKM_EXEC_EXPORT
const IndicesFromType *GetIndicesFromPointer() const
{
return &this->IndicesFrom;
}
/// \brief The shape of the input cell.
///
@ -263,7 +319,9 @@ public:
CellShapeTag GetCellShape() const { return this->CellShape; }
private:
vtkm::Id Index;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
LogicalIndexType LogicalIndex;
IndicesFromType IndicesFrom;
CellShapeTag CellShape;

87
vtkm/exec/arg/VisitIndex.h Normal file
View File

@ -0,0 +1,87 @@
//============================================================================
// 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 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_exec_arg_VisitIndex_h
#define vtk_m_exec_arg_VisitIndex_h
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Aspect tag to use for getting the work index.
///
/// The \c AspectTagVisitIndex aspect tag causes the \c Fetch class to ignore
/// whatever data is in the associated execution object and return the visit
/// index.
///
struct AspectTagVisitIndex { };
/// \brief The \c ExecutionSignature tag to use to get the visit index
///
/// When a worklet is dispatched, there is a scatter operation defined that
/// optionally allows each input to go to multiple output entries. When one
/// input is assigned to multiple outputs, there needs to be a mechanism to
/// uniquely identify which output is which. The visit index is a value between
/// 0 and the number of outputs a particular input goes to. This tag in the \c
/// ExecutionSignature passes the visit index for this work. \c WorkletBase
/// contains a typedef that points to this class.
///
struct VisitIndex : vtkm::exec::arg::ExecutionSignatureTagBase
{
// The index does not really matter because the fetch is going to ignore it.
// However, it still has to point to a valid parameter in the
// ControlSignature because the templating is going to grab a fetch tag
// whether we use it or not. 1 should be guaranteed to be valid since you
// need at least one argument for the input domain.
static const vtkm::IdComponent INDEX = 1;
typedef vtkm::exec::arg::AspectTagVisitIndex AspectTag;
};
template<typename FetchTag, typename ThreadIndicesType, typename ExecObjectType>
struct Fetch<FetchTag,
vtkm::exec::arg::AspectTagVisitIndex,
ThreadIndicesType,
ExecObjectType>
{
typedef vtkm::IdComponent ValueType;
VTKM_EXEC_EXPORT
vtkm::IdComponent Load(const ThreadIndicesType &indices,
const ExecObjectType &) const
{
return indices.GetVisitIndex();
}
VTKM_EXEC_EXPORT
void Store(const ThreadIndicesType &,
const ExecObjectType &,
const ValueType &) const
{
// Store is a no-op.
}
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_VisitIndex_h

2
vtkm/exec/arg/WorkIndex.h
View File

@ -63,7 +63,7 @@ struct Fetch<FetchTag,
VTKM_EXEC_EXPORT
vtkm::Id Load(const ThreadIndicesType &indices, const ExecObjectType &) const
{
return indices.GetIndex();
return indices.GetOutputIndex();
}
VTKM_EXEC_EXPORT

5
vtkm/exec/arg/testing/CMakeLists.txt
View File

@ -17,6 +17,11 @@
## Laboratory (LANL), the U.S. Government retains certain rights in
## this software.
##============================================================================
set(headers
ThreadIndicesTesting.h
)
vtkm_declare_headers(${headers})
set(unit_tests
UnitTestExecutionSignatureTag.cxx

70
vtkm/exec/arg/testing/ThreadIndicesTesting.h Normal file
View File

@ -0,0 +1,70 @@
//============================================================================
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_exec_arg_testing_ThreadIndicesTesting_h
#define vtk_m_exec_arg_testing_ThreadIndicesTesting_h
#include <vtkm/Types.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Simplified version of ThreadIndices for unit testing purposes
///
class ThreadIndicesTesting
{
public:
VTKM_EXEC_CONT_EXPORT
ThreadIndicesTesting(vtkm::Id index)
: InputIndex(index), OutputIndex(index), VisitIndex(0) { }
VTKM_EXEC_CONT_EXPORT
ThreadIndicesTesting(vtkm::Id inputIndex,
vtkm::Id outputIndex,
vtkm::IdComponent visitIndex)
: InputIndex(inputIndex), OutputIndex(outputIndex), VisitIndex(visitIndex)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetInputIndex() const { return this->InputIndex; }
VTKM_EXEC_CONT_EXPORT
vtkm::Id3 GetInputIndex3D() const
{
return vtkm::Id3(this->GetInputIndex(), 0, 0);
}
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetOutputIndex() const { return this->OutputIndex; }
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetVisitIndex() const { return this->VisitIndex; }
private:
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_testing_ThreadIndicesTesting_h

50
vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectIn.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -47,59 +44,36 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectInTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectIn,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
VTKM_TEST_ASSERT(test_equal(value, TestValue(index, T())),
"Got invalid value from Load.");
value = T(T(2)*value);
// This should be a no-op, but we should be able to call it.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
}
void operator()() const
{
std::cout << "Trying ArrayDirectIn fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -107,11 +81,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectInTests<1,T>()();
FetchArrayDirectInTests<2,T>()();
FetchArrayDirectInTests<3,T>()();
FetchArrayDirectInTests<4,T>()();
FetchArrayDirectInTests<5,T>()();
FetchArrayDirectInTests<T>()();
}
};

49
vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectInOut.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -59,19 +56,18 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectInTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectInOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
@ -80,19 +76,15 @@ struct FetchArrayDirectInTests
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
VTKM_TEST_ASSERT(test_equal(value, TestValue(index, T())),
"Got invalid value from Load.");
value = T(T(2)*value);
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
VTKM_TEST_ASSERT(g_NumSets == ARRAY_SIZE,
@ -100,23 +92,6 @@ struct FetchArrayDirectInTests
"Store method must be wrong.");
}
void operator()() const
{
std::cout << "Trying ArrayDirectInOut fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -124,11 +99,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectInTests<1,T>()();
FetchArrayDirectInTests<2,T>()();
FetchArrayDirectInTests<3,T>()();
FetchArrayDirectInTests<4,T>()();
FetchArrayDirectInTests<5,T>()();
FetchArrayDirectInTests<T>()();
}
};

49
vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectOut.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -51,19 +48,18 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectOutTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
@ -72,19 +68,15 @@ struct FetchArrayDirectOutTests
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
// This is a no-op, but should be callable.
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
value = TestValue(index, T());
// The portal will check to make sure we are setting a good value.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
VTKM_TEST_ASSERT(g_NumSets == ARRAY_SIZE,
@ -92,23 +84,6 @@ struct FetchArrayDirectOutTests
"Store method must be wrong.");
}
void operator()() const
{
std::cout << "Trying ArrayDirectOut fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -116,11 +91,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectOutTests<1,T>()();
FetchArrayDirectOutTests<2,T>()();
FetchArrayDirectOutTests<3,T>()();
FetchArrayDirectOutTests<4,T>()();
FetchArrayDirectOutTests<5,T>()();
FetchArrayDirectOutTests<T>()();
}
};

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

@ -47,7 +47,21 @@ struct TestPortal
}
};
struct NullParam { };
struct TestIndexPortal
{
typedef vtkm::Id ValueType;
VTKM_EXEC_CONT_EXPORT
ValueType Get(vtkm::Id index) const { return index; }
};
struct TestZeroPortal
{
typedef vtkm::IdComponent ValueType;
VTKM_EXEC_CONT_EXPORT
ValueType Get(vtkm::Id) const { return 0; }
};
template<vtkm::IdComponent InputDomainIndex,
vtkm::IdComponent ParamIndex,
@ -102,7 +116,11 @@ struct FetchArrayTopologyMapInTests
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
void(vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType)>
BaseFunctionInterface;
vtkm::internal::ConnectivityStructuredInternals<3> connectivityInternals;
@ -115,8 +133,10 @@ struct FetchArrayTopologyMapInTests
BaseFunctionInterface()
.Replace<InputDomainIndex>(connectivity)
.template Replace<ParamIndex>(TestPortal<T>()),
NullParam(),
NullParam()));
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal()));
}
};
@ -175,7 +195,11 @@ void TryStructuredPointCoordinates(
const vtkm::internal::ArrayPortalUniformPointCoordinates &coordinates)
{
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
void(vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType)>
BaseFunctionInterface;
// Try with topology in argument 1 and point coordinates in argument 2
@ -184,8 +208,10 @@ void TryStructuredPointCoordinates(
BaseFunctionInterface()
.Replace<1>(connectivity)
.template Replace<2>(coordinates),
NullParam(),
NullParam())
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal())
);
// Try again with topology in argument 3 and point coordinates in argument 1
TryStructuredPointCoordinatesInvocation<NumDimensions,1>(
@ -193,8 +219,10 @@ void TryStructuredPointCoordinates(
BaseFunctionInterface()
.Replace<3>(connectivity)
.template Replace<1>(coordinates),
NullParam(),
NullParam())
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal())
);
}

50
vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagExecObject.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/exec/ExecutionObjectBase.h>
@ -40,64 +37,37 @@ struct TestExecutionObject : public vtkm::exec::ExecutionObjectBase
vtkm::Int32 Number;
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename Invocation>
void TryInvocation(const Invocation &invocation)
void TryInvocation()
{
TestExecutionObject execObjectStore(EXPECTED_NUMBER);
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagExecObject,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestExecutionObject> FetchType;
FetchType fetch;
vtkm::exec::arg::ThreadIndicesBasic indices(0, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(0);
TestExecutionObject execObject = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
TestExecutionObject execObject = fetch.Load(indices, execObjectStore);
VTKM_TEST_ASSERT(execObject.Number == EXPECTED_NUMBER,
"Did not load object correctly.");
execObject.Number = -1;
// This should be a no-op.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
execObject);
fetch.Store(indices, execObjectStore, execObject);
// Data in Invocation should not have changed.
execObject = invocation.Parameters.template GetParameter<ParamIndex>();
VTKM_TEST_ASSERT(execObject.Number == EXPECTED_NUMBER,
VTKM_TEST_ASSERT(execObjectStore.Number == EXPECTED_NUMBER,
"Fetch changed read-only execution object.");
}
template<vtkm::IdComponent ParamIndex>
void TryParamIndex()
{
std::cout << "Trying ExecObject fetch on parameter " << ParamIndex
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
TryInvocation<ParamIndex>(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestExecutionObject(EXPECTED_NUMBER)),
NullParam(),
NullParam()));
}
void TestExecObjectFetch()
{
TryParamIndex<1>();
TryParamIndex<2>();
TryParamIndex<3>();
TryParamIndex<4>();
TryParamIndex<5>();
TryInvocation();
}
} // anonymous namespace

35
vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx
View File

@ -22,56 +22,39 @@
#include <vtkm/exec/arg/FetchTagArrayDirectIn.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
namespace {
struct NullParam { };
template<typename Invocation>
void TryInvocation(const Invocation &invocation)
void TestWorkIndexFetch()
{
std::cout << "Trying WorkIndex fetch." << std::endl;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectIn, // Not used but probably common.
vtkm::exec::arg::AspectTagWorkIndex,
vtkm::exec::arg::ThreadIndicesBasic,
NullParam> FetchType;
vtkm::exec::arg::ThreadIndicesTesting,
vtkm::internal::NullType> FetchType;
FetchType fetch;
for (vtkm::Id index = 0; index < 10; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
vtkm::Id value = fetch.Load(indices, NullParam());
vtkm::Id value = fetch.Load(indices, vtkm::internal::NullType());
VTKM_TEST_ASSERT(value == index,
"Fetch did not give correct work index.");
value++;
// This should be a no-op.
fetch.Store(indices, NullParam(), value);
fetch.Store(indices, vtkm::internal::NullType(), value);
}
}
void TestWorkIndexFetch()
{
std::cout << "Trying WorkIndex fetch." << std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
TryInvocation(vtkm::internal::make_Invocation<1>(BaseFunctionInterface(),
NullParam(),
NullParam()));
}
} // anonymous namespace
int UnitTestFetchWorkIndex(int, char *[])

14
vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
View File

@ -96,7 +96,7 @@ template<typename T, typename Enable = void>
struct load_through_texture
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
return *(data.get());
}
@ -109,7 +109,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTextureLoad<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseScalarTextureLoad");
@ -125,7 +125,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextureLoads<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseVecTextureLoads");
@ -193,7 +193,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
typedef typename boost::remove_const<T>::type NonConstT;
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseMultipleScalarTextureLoads");
@ -204,7 +204,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
}
__device__
static T getAs(const thrust::system::cuda::pointer<T>& data)
static T getAs(const thrust::system::cuda::pointer<const T>& data)
{
//we need to fetch each component individually
const vtkm::IdComponent NUM_COMPONENTS= T::NUM_COMPONENTS;
@ -291,8 +291,8 @@ public:
VTKM_EXEC_CONT_EXPORT ConstArrayPortalFromThrust() { }
VTKM_CONT_EXPORT
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< T > begin,
const thrust::system::cuda::pointer< T > end)
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< const T > begin,
const thrust::system::cuda::pointer< const T > end)
: BeginIterator( begin ),
EndIterator( end )
{

10
vtkm/exec/internal/VecFromPortalPermute.h
View File

@ -45,12 +45,12 @@ public:
VecFromPortalPermute() { }
VTKM_EXEC_EXPORT
VecFromPortalPermute(const IndexVecType &indices, const PortalType &portal)
VecFromPortalPermute(const IndexVecType *indices, const PortalType &portal)
: Indices(indices), Portal(portal) { }
VTKM_EXEC_EXPORT
vtkm::IdComponent GetNumberOfComponents() const {
return this->Indices.GetNumberOfComponents();
return this->Indices->GetNumberOfComponents();
}
template<vtkm::IdComponent DestSize>
@ -61,18 +61,18 @@ public:
vtkm::Min(DestSize, this->GetNumberOfComponents());
for (vtkm::IdComponent index = 0; index < numComponents; index++)
{
dest[index] = this->Portal.Get(this->Indices[index]);
dest[index] = (*this)[index];
}
}
VTKM_EXEC_EXPORT
ComponentType operator[](vtkm::IdComponent index) const
{
return this->Portal.Get(this->Indices[index]);
return this->Portal.Get((*this->Indices)[index]);
}
private:
IndexVecType Indices;
const IndexVecType *Indices;
PortalType Portal;
};

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

@ -84,6 +84,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1>
@ -94,14 +96,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -136,6 +142,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1>
VTKM_EXEC_EXPORT
@ -145,14 +153,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -180,6 +192,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -191,14 +205,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -242,6 +260,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2>
@ -252,14 +272,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -296,6 +320,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -308,14 +334,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -368,6 +398,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -379,14 +411,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -432,6 +468,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -445,14 +483,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -514,6 +556,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -526,14 +570,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -588,6 +636,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -602,14 +652,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -680,6 +734,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -693,14 +749,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -764,6 +824,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -779,14 +841,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -866,6 +932,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -880,14 +948,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -960,6 +1032,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -976,14 +1050,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1072,6 +1150,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1087,14 +1167,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1176,6 +1260,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1193,14 +1279,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1298,6 +1388,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1314,14 +1406,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1412,6 +1508,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1430,14 +1528,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1544,6 +1646,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1561,14 +1665,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1668,6 +1776,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1687,14 +1797,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1810,6 +1924,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1828,14 +1944,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;

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

@ -146,6 +146,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
$template_params(num_params)>
VTKM_EXEC_EXPORT
@ -155,14 +157,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
$for(param_index in range(1, num_params+1))\
typedef InvocationToFetch<ThreadIndicesType,Invocation,$(param_index)> FetchInfo$(param_index);
@ -201,6 +207,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
$template_params(num_params, start=1)>
VTKM_EXEC_EXPORT
@ -210,14 +218,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params, return_type='void')>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params, return_type='void')>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
$for(param_index in range(1, num_params+1))\
typedef InvocationToFetch<ThreadIndicesType,Invocation,$(param_index)> FetchInfo$(param_index);

2
vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx
View File

@ -87,7 +87,7 @@ struct VecFromPortalPermuteTestFunctor
indices.Append(offset + 2*index);
}
VecType vec(indices, portal);
VecType vec(&indices, portal);
VTKM_TEST_ASSERT(vec.GetNumberOfComponents() == length,
"Wrong length.");

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

@ -47,6 +47,19 @@ struct TestExecObject
vtkm::Id *Value;
};
struct MyOutputToInputMapPortal
{
typedef vtkm::Id ValueType;
VTKM_EXEC_CONT_EXPORT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct MyVisitArrayPortal
{
typedef vtkm::IdComponent ValueType;
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct TestFetchTagInput { };
struct TestFetchTagOutput { };
@ -78,7 +91,7 @@ struct Fetch<
VTKM_EXEC_EXPORT
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject) const {
return *execObject.Value + 10*indices.GetIndex();
return *execObject.Value + 10*indices.GetInputIndex();
}
VTKM_EXEC_EXPORT
@ -109,7 +122,7 @@ struct Fetch<
void Store(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject,
ValueType value) const {
*execObject.Value = value + 20*indices.GetIndex();
*execObject.Value = value + 20*indices.GetOutputIndex();
}
};
@ -141,13 +154,17 @@ typedef vtkm::internal::Invocation<
ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface1,
1> InvocationType1;
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal> InvocationType1;
typedef vtkm::internal::Invocation<
ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface2,
1> InvocationType2;
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal> InvocationType2;
// Not a full worklet, but provides operators that we expect in a worklet.
struct TestWorkletProxy : vtkm::exec::FunctorBase
@ -232,7 +249,9 @@ void TestDoWorkletInvoke()
CallDoWorkletInvokeFunctor(
vtkm::internal::make_Invocation<1>(execObjects,
TestControlInterface(),
TestExecutionInterface1()),
TestExecutionInterface1(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
1);
VTKM_TEST_ASSERT(inputTestValue == 5, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 100 + 30,
@ -244,7 +263,9 @@ void TestDoWorkletInvoke()
CallDoWorkletInvokeFunctor(
vtkm::internal::make_Invocation<1>(execObjects,
TestControlInterface(),
TestExecutionInterface2()),
TestExecutionInterface2(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
2);
VTKM_TEST_ASSERT(inputTestValue == 6, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 200 + 30*2,

132
vtkm/internal/Invocation.h
View File

@ -34,7 +34,9 @@ namespace internal {
template<typename _ParameterInterface,
typename _ControlInterface,
typename _ExecutionInterface,
vtkm::IdComponent _InputDomainIndex>
vtkm::IdComponent _InputDomainIndex,
typename _OutputToInputMapType = vtkm::internal::NullType,
typename _VisitArrayType = vtkm::internal::NullType>
struct Invocation
{
/// \brief The types of the parameters
@ -68,8 +70,32 @@ struct Invocation
///
static const vtkm::IdComponent InputDomainIndex = _InputDomainIndex;
/// \brief An array representing the output to input map.
///
/// When a worklet is invoked, there is an optional scatter operation that
/// allows you to vary the number of outputs each input affects. This is
/// represented with a map where each output points to an input that creates
/// it.
///
typedef _OutputToInputMapType OutputToInputMapType;
/// \brief An array containing visit indices.
///
/// When a worklet is invoked, there is an optinonal scatter operation that
/// allows you to vary the number of outputs each input affects. Thus,
/// multiple outputs may point to the same input. The visit index uniquely
/// identifies which instance each is.
///
typedef _VisitArrayType VisitArrayType;
VTKM_CONT_EXPORT
Invocation(ParameterInterface parameters) : Parameters(parameters) { }
Invocation(ParameterInterface parameters,
OutputToInputMapType outputToInputMap = OutputToInputMapType(),
VisitArrayType visitArray = VisitArrayType())
: Parameters(parameters),
OutputToInputMap(outputToInputMap),
VisitArray(visitArray)
{ }
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c Parameters replaced.
@ -79,7 +105,9 @@ struct Invocation
typedef Invocation<NewParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex> type;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -90,7 +118,7 @@ struct Invocation
typename ChangeParametersType<NewParameterInterface>::type
ChangeParameters(NewParameterInterface newParameters) const {
return typename ChangeParametersType<NewParameterInterface>::type(
newParameters);
newParameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -111,7 +139,7 @@ struct Invocation
typename ChangeControlInterfaceType<NewControlInterface>::type
ChangeControlInterface(NewControlInterface) const {
return typename ChangeControlInterfaceType<NewControlInterface>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -132,7 +160,7 @@ struct Invocation
typename ChangeExecutionInterfaceType<NewExecutionInterface>::type
ChangeExecutionInterface(NewExecutionInterface) const {
return typename ChangeExecutionInterfaceType<NewExecutionInterface>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -154,7 +182,55 @@ struct Invocation
typename ChangeInputDomainIndexType<NewInputDomainIndex>::type
ChangeInputDomainIndex() const {
return typename ChangeInputDomainIndexType<NewInputDomainIndex>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c OutputToInputMapType replaced.
///
template<typename NewOutputToInputMapType>
struct ChangeOutputToInputMapType {
typedef Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
NewOutputToInputMapType,
VisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
/// \c OutputToInputMap is replaced with that provided.
///
template<typename NewOutputToInputMapType>
VTKM_CONT_EXPORT
typename ChangeOutputToInputMapType<NewOutputToInputMapType>::type
ChangeOutputToInputMap(NewOutputToInputMapType newOutputToInputMap) const {
return typename ChangeOutputToInputMapType<NewOutputToInputMapType>::type(
this->Parameters, newOutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c VisitArrayType replaced.
///
template<typename NewVisitArrayType>
struct ChangeVisitArrayType {
typedef Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
NewVisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
/// \c VisitArray is replaced with that provided.
///
template<typename NewVisitArrayType>
VTKM_CONT_EXPORT
typename ChangeVisitArrayType<NewVisitArrayType>::type
ChangeVisitArray(NewVisitArrayType newVisitArray) const {
return typename ChangeVisitArrayType<NewVisitArrayType>::type(
this->Parameters, this->OutputToInputMap, newVisitArray);
}
/// A convenience typedef for the input domain type.
@ -172,13 +248,43 @@ struct Invocation
}
/// The state of an \c Invocation object holds the parameters of the
/// invocation.
/// invocation. As well as the output to input map and the visit array.
///
ParameterInterface Parameters;
OutputToInputMapType OutputToInputMap;
VisitArrayType VisitArray;
};
/// Convenience function for creating an Invocation object.
///
template<vtkm::IdComponent InputDomainIndex,
typename ControlInterface,
typename ExecutionInterface,
typename ParameterInterface,
typename OutputToInputMapType,
typename VisitArrayType>
VTKM_CONT_EXPORT
vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>
make_Invocation(const ParameterInterface &params,
ControlInterface,
ExecutionInterface,
OutputToInputMapType outputToInputMap,
VisitArrayType visitArray)
{
return vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>(params,
outputToInputMap,
visitArray);
}
template<vtkm::IdComponent InputDomainIndex,
typename ControlInterface,
typename ExecutionInterface,
@ -192,10 +298,12 @@ make_Invocation(const ParameterInterface &params,
ControlInterface = ControlInterface(),
ExecutionInterface = ExecutionInterface())
{
return vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex>(params);
return vtkm::internal::make_Invocation<InputDomainIndex>(
params,
ControlInterface(),
ExecutionInterface(),
vtkm::internal::NullType(),
vtkm::internal::NullType());
}
}

32
vtkm/internal/ListTagDetail.h
View File

@ -46,8 +46,6 @@ struct ListRoot { };
template<typename signature>
struct ListBase { };
struct ListParamNull { };
//-----------------------------------------------------------------------------
template<typename Functor>
@ -1151,21 +1149,21 @@ struct ListContainsImpl<ListBase<void(T1,T2,T3,T4,T5,T6,T7,T8,T9,T10,T11,T12,T13
/// A basic tag for a list of typenames. This struct can be subclassed
/// and still behave like a list tag.
template<typename T1 = vtkm::detail::ListParamNull,
typename T2 = vtkm::detail::ListParamNull,
typename T3 = vtkm::detail::ListParamNull,
typename T4 = vtkm::detail::ListParamNull,
typename T5 = vtkm::detail::ListParamNull,
typename T6 = vtkm::detail::ListParamNull,
typename T7 = vtkm::detail::ListParamNull,
typename T8 = vtkm::detail::ListParamNull,
typename T9 = vtkm::detail::ListParamNull,
typename T10 = vtkm::detail::ListParamNull,
typename T11 = vtkm::detail::ListParamNull,
typename T12 = vtkm::detail::ListParamNull,
typename T13 = vtkm::detail::ListParamNull,
typename T14 = vtkm::detail::ListParamNull,
typename T15 = vtkm::detail::ListParamNull>
template<typename T1 = vtkm::internal::NullType,
typename T2 = vtkm::internal::NullType,
typename T3 = vtkm::internal::NullType,
typename T4 = vtkm::internal::NullType,
typename T5 = vtkm::internal::NullType,
typename T6 = vtkm::internal::NullType,
typename T7 = vtkm::internal::NullType,
typename T8 = vtkm::internal::NullType,
typename T9 = vtkm::internal::NullType,
typename T10 = vtkm::internal::NullType,
typename T11 = vtkm::internal::NullType,
typename T12 = vtkm::internal::NullType,
typename T13 = vtkm::internal::NullType,
typename T14 = vtkm::internal::NullType,
typename T15 = vtkm::internal::NullType>
struct ListTagBase : detail::ListRoot
{
typedef detail::ListBase<void(T1,T2,T3,T4,T5,T6,T7,T8,T9,T10,T11,T12,T13,T14,T15)> List;

4
vtkm/internal/ListTagDetail.h.in
View File

@ -94,8 +94,6 @@ struct ListRoot { };
template<typename signature>
struct ListBase { };
struct ListParamNull { };
//-----------------------------------------------------------------------------
template<typename Functor>
@ -172,7 +170,7 @@ $endfor\
/// A basic tag for a list of typenames. This struct can be subclassed
/// and still behave like a list tag.
template<$template_params(max_base_list, default=' = vtkm::detail::ListParamNull')>
template<$template_params(max_base_list, default=' = vtkm::internal::NullType')>
struct ListTagBase : detail::ListRoot
{
typedef detail::ListBase<void($param_list(max_base_list))> List;

4
vtkm/testing/Testing.h
View File

@ -28,10 +28,6 @@
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/static_assert.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
#include <exception>
#include <iostream>
#include <sstream>

4
vtkm/worklet/CMakeLists.txt
View File

@ -23,6 +23,7 @@ set(headers
WorkletMapField.h
DispatcherMapTopology.h
WorkletMapTopology.h
AverageByKey.h
CellAverage.h
Clip.h
@ -32,6 +33,9 @@ set(headers
IsosurfaceUniformGrid.h
MarchingCubesDataTables.h
PointElevation.h
ScatterCounting.h
ScatterIdentity.h
ScatterUniform.h
TetrahedralizeExplicitGrid.h
TetrahedralizeUniformGrid.h
Threshold.h

367
vtkm/worklet/IsosurfaceUniformGrid.h
View File

@ -21,20 +21,26 @@
#ifndef vtk_m_worklet_IsosurfaceUniformGrid_h
#define vtk_m_worklet_IsosurfaceUniformGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/Pair.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/exec/CellDerivative.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterCounting.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include "MarchingCubesDataTables.h"
namespace vtkm {
@ -51,197 +57,147 @@ public:
public:
typedef void ControlSignature(FieldInPoint<Scalar> inNodes,
TopologyIn topology,
FieldOutCell<IdType> outNumVertices);
typedef void ExecutionSignature(_1, _3);
FieldOutCell<> outNumTriangles,
ExecObject numTrianglesTable);
typedef void ExecutionSignature(_1, _3, _4);
typedef _2 InputDomain;
typedef vtkm::cont::ArrayHandle<vtkm::Id> IdArrayHandle;
typedef typename IdArrayHandle::ExecutionTypes<DeviceAdapter>::PortalConst IdPortalType;
IdPortalType VertexTable;
FieldType Isovalue;
VTKM_CONT_EXPORT
ClassifyCell(IdPortalType vTable, FieldType isovalue) :
VertexTable(vTable),
ClassifyCell(FieldType isovalue) :
Isovalue(isovalue)
{
}
template<typename InPointVecType>
template<typename InPointVecType,
typename NumTrianglesTablePortalType>
VTKM_EXEC_EXPORT
void operator()(const InPointVecType &pointValues,
vtkm::Id& numVertices) const
void operator()(const InPointVecType &fieldIn,
vtkm::IdComponent &numTriangles,
const NumTrianglesTablePortalType &numTrianglesTable) const
{
vtkm::Id caseNumber = (pointValues[0] > this->Isovalue);
caseNumber += (pointValues[1] > this->Isovalue)*2;
caseNumber += (pointValues[2] > this->Isovalue)*4;
caseNumber += (pointValues[3] > this->Isovalue)*8;
caseNumber += (pointValues[4] > this->Isovalue)*16;
caseNumber += (pointValues[5] > this->Isovalue)*32;
caseNumber += (pointValues[6] > this->Isovalue)*64;
caseNumber += (pointValues[7] > this->Isovalue)*128;
numVertices = this->VertexTable.Get(caseNumber) / 3;
vtkm::IdComponent caseNumber =
( (fieldIn[0] > this->Isovalue)
| (fieldIn[1] > this->Isovalue)<<1
| (fieldIn[2] > this->Isovalue)<<2
| (fieldIn[3] > this->Isovalue)<<3
| (fieldIn[4] > this->Isovalue)<<4
| (fieldIn[5] > this->Isovalue)<<5
| (fieldIn[6] > this->Isovalue)<<6
| (fieldIn[7] > this->Isovalue)<<7 );
numTriangles = numTrianglesTable.Get(caseNumber);
}
};
/// \brief Compute isosurface vertices and scalars
class IsoSurfaceGenerate : public vtkm::worklet::WorkletMapField
class IsoSurfaceGenerate : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
FieldIn<IdType> inputIteration);
typedef void ExecutionSignature(WorkIndex, _1, _2);
typedef _1 InputDomain;
typedef void ControlSignature(
TopologyIn topology, // Cell set
FieldInPoint<> fieldIn, // Input point field defining the contour
FieldInPoint<Vec3> pcoordIn, // Input point coordinates
FieldOutCell<> vertexOut, // Vertices for output triangles
// TODO: Have a better way to iterate over and interpolate fields
FieldInPoint<Scalar> scalarsIn, // Scalars to interpolate
FieldOutCell<> scalarsOut, // Interpolated scalars (one per tri vertex)
FieldOutCell<> normalsOut, // Estimated normals (one per tri vertex)
ExecObject TriTable // An array portal with the triangle table
);
typedef void ExecutionSignature(
CellShape, _2, _3, _4, _5, _6, _7, _8, VisitIndex);
const FieldType Isovalue;
vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
typedef typename vtkm::cont::ArrayHandle<FieldType>::template ExecutionTypes<DeviceAdapter>::PortalConst FieldPortalType;
FieldPortalType Field, Source;
typedef typename vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3> >::template ExecutionTypes<DeviceAdapter>::Portal VectorPortalType;
VectorPortalType Vertices;
VectorPortalType Normals;
typedef typename vtkm::cont::ArrayHandle<FieldType>::template ExecutionTypes<DeviceAdapter>::Portal OutputPortalType;
OutputPortalType Scalars;
typedef typename vtkm::cont::ArrayHandle<vtkm::Id> IdArrayHandle;
typedef typename IdArrayHandle::ExecutionTypes<DeviceAdapter>::PortalConst IdPortalType;
IdPortalType TriTable;
const vtkm::Id cellsPerLayer, pointsPerLayer;
template<typename U, typename W, typename X>
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
IsoSurfaceGenerate(FieldType ivalue, const vtkm::Id3 cdims, IdPortalType triTablePortal,
const U & field, const U & source, const W & vertices, const W & normals, const X & scalars) :
Isovalue(ivalue),
xdim(cdims[0]), ydim(cdims[1]), zdim(cdims[2]),
xmin(-1), ymin(-1), zmin(-1), xmax(1), ymax(1), zmax(1),
Field( field.PrepareForInput( DeviceAdapter() ) ),
Source( source.PrepareForInput( DeviceAdapter() ) ),
Vertices(vertices),
Normals(normals),
Scalars(scalars),
TriTable(triTablePortal),
cellsPerLayer(xdim * ydim),
pointsPerLayer ((xdim+1)*(ydim+1))
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
IsoSurfaceGenerate(FieldType isovalue,
const CountArrayType &countArray,
Device)
: Isovalue(isovalue), Scatter(countArray, Device()) { }
template<typename CellShapeTag,
typename FieldInType, // Vec-like, one per input point
typename CoordType, // Vec-like (one per input point) of Vec-3
typename VertexOutType, // Vec-3 of Vec-3 coordinates (for triangle)
typename ScalarInType, // Vec-like, one per input point
typename ScalarOutType, // Vec-3 (one value per tri vertex)
typename NormalOutType, // Vec-3 of Vec-3
typename TriTablePortalType> // Array portal
VTKM_EXEC_EXPORT
void operator()(vtkm::Id outputCellId, vtkm::Id inputCellId, vtkm::Id inputLowerBounds) const
void operator()(
CellShapeTag shape,
const FieldInType &fieldIn, // Input point field defining the contour
const CoordType &coords, // Input point coordinates
VertexOutType &vertexOut, // Vertices for output triangles
// TODO: Have a better way to iterate over and interpolate fields
const ScalarInType &scalarsIn, // Scalars to interpolate
ScalarOutType &scalarsOut, // Interpolated scalars (one per tri vertex)
NormalOutType &normalsOut, // Estimated normals (one per tri vertex)
const TriTablePortalType &triTable, // An array portal with the triangle table
vtkm::IdComponent visitIndex
) const
{
// Get data for this cell
const int verticesForEdge[] = { 0, 1, 1, 2, 3, 2, 0, 3,
4, 5, 5, 6, 7, 6, 4, 7,
0, 4, 1, 5, 2, 6, 3, 7 };
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
const vtkm::Id z = inputCellId / cellsPerLayer;
// Compute indices for the eight vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1) + z * pointsPerLayer;
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
const vtkm::Id i4 = i0 + pointsPerLayer;
const vtkm::Id i5 = i1 + pointsPerLayer;
const vtkm::Id i6 = i2 + pointsPerLayer;
const vtkm::Id i7 = i3 + pointsPerLayer;
// Get the field values at these eight vertices
FieldType f[8];
f[0] = this->Field.Get(i0);
f[1] = this->Field.Get(i1);
f[2] = this->Field.Get(i2);
f[3] = this->Field.Get(i3);
f[4] = this->Field.Get(i4);
f[5] = this->Field.Get(i5);
f[6] = this->Field.Get(i6);
f[7] = this->Field.Get(i7);
const vtkm::IdComponent verticesForEdge[] = { 0, 1, 1, 2, 3, 2, 0, 3,
4, 5, 5, 6, 7, 6, 4, 7,
0, 4, 1, 5, 2, 6, 3, 7 };
// Compute the Marching Cubes case number for this cell
unsigned int cubeindex = 0;
cubeindex += (f[0] > this->Isovalue);
cubeindex += (f[1] > this->Isovalue)*2;
cubeindex += (f[2] > this->Isovalue)*4;
cubeindex += (f[3] > this->Isovalue)*8;
cubeindex += (f[4] > this->Isovalue)*16;
cubeindex += (f[5] > this->Isovalue)*32;
cubeindex += (f[6] > this->Isovalue)*64;
cubeindex += (f[7] > this->Isovalue)*128;
// printf("inputCellId: %i \n",inputCellId);
// printf("x: %i, y: %i, z: %i \n",x, y, z);
// printf("i0: %i \n",i0);
// printf("f0 %F\n", f[0]);
// printf("cubeindex: %i \n",cubeindex);
// printf("numCells: %i \n",(vtkm::worklet::internal::numVerticesTable[cubeindex]/3) );
// Compute the coordinates of the uniform regular grid at each of the cell's eight vertices
vtkm::Vec<FieldType, 3> p[8];
// If we have offset and spacing, can we simplify this computation
{
vtkm::Vec<FieldType, 3> offset = vtkm::make_Vec(xmin+(xmax-xmin),
ymin+(ymax-ymin),
zmin+(zmax-zmin) );
vtkm::Vec<FieldType, 3> spacing = vtkm::make_Vec( 1.0f /((float)(xdim-1)),
1.0f /((float)(ydim-1)),
1.0f /((float)(zdim-1)));
vtkm::Vec<FieldType, 3> firstPoint = offset * spacing * vtkm::make_Vec( x, y, z );
vtkm::Vec<FieldType, 3> secondPoint = offset * spacing * vtkm::make_Vec( x+1, y+1, z+1 );
p[0] = vtkm::make_Vec( firstPoint[0], firstPoint[1], firstPoint[2]);
p[1] = vtkm::make_Vec( secondPoint[0], firstPoint[1], firstPoint[2]);
p[2] = vtkm::make_Vec( secondPoint[0], secondPoint[1], firstPoint[2]);
p[3] = vtkm::make_Vec( firstPoint[0], secondPoint[1], firstPoint[2]);
p[4] = vtkm::make_Vec( firstPoint[0], firstPoint[1], secondPoint[2]);
p[5] = vtkm::make_Vec( secondPoint[0], firstPoint[1], secondPoint[2]);
p[6] = vtkm::make_Vec( secondPoint[0], secondPoint[1], secondPoint[2]);
p[7] = vtkm::make_Vec( firstPoint[0], secondPoint[1], secondPoint[2]);
}
// Get the scalar source values at the eight vertices
FieldType s[8];
s[0] = this->Source.Get(i0);
s[1] = this->Source.Get(i1);
s[2] = this->Source.Get(i2);
s[3] = this->Source.Get(i3);
s[4] = this->Source.Get(i4);
s[5] = this->Source.Get(i5);
s[6] = this->Source.Get(i6);
s[7] = this->Source.Get(i7);
vtkm::IdComponent caseNumber =
( (fieldIn[0] > this->Isovalue)
| (fieldIn[1] > this->Isovalue)<<1
| (fieldIn[2] > this->Isovalue)<<2
| (fieldIn[3] > this->Isovalue)<<3
| (fieldIn[4] > this->Isovalue)<<4
| (fieldIn[5] > this->Isovalue)<<5
| (fieldIn[6] > this->Isovalue)<<6
| (fieldIn[7] > this->Isovalue)<<7 );
// Interpolate for vertex positions and associated scalar values
const vtkm::Id inputIteration = (outputCellId - inputLowerBounds);
const vtkm::Id outputVertId = outputCellId * 3;
const vtkm::Id cellOffset = static_cast<vtkm::Id>(cubeindex*16) + (inputIteration * 3);
for (int v = 0; v < 3; v++)
const vtkm::Id triTableOffset =
static_cast<vtkm::Id>(caseNumber*16 + visitIndex*3);
for (vtkm::IdComponent triVertex = 0;
triVertex < 3;
triVertex++)
{
const vtkm::Id edge = this->TriTable.Get(cellOffset + v);
const int v0 = verticesForEdge[2*edge];
const int v1 = verticesForEdge[2*edge + 1];
const FieldType t = (this->Isovalue - f[v0]) / (f[v1] - f[v0]);
this->Vertices.Set(outputVertId + v, vtkm::Lerp(p[v0], p[v1], t));
this->Scalars.Set(outputVertId + v, vtkm::Lerp(s[v0], s[v1], t));
const vtkm::IdComponent edgeIndex =
triTable.Get(triTableOffset + triVertex);
const vtkm::IdComponent edgeVertex0 = verticesForEdge[2*edgeIndex + 0];
const vtkm::IdComponent edgeVertex1 = verticesForEdge[2*edgeIndex + 1];
const FieldType fieldValue0 = fieldIn[edgeVertex0];
const FieldType fieldValue1 = fieldIn[edgeVertex1];
const FieldType interpolant =
(this->Isovalue - fieldValue0) / (fieldValue1 - fieldValue0);
vertexOut[triVertex] = vtkm::Lerp(coords[edgeVertex0],
coords[edgeVertex1],
interpolant);
scalarsOut[triVertex] = vtkm::Lerp(scalarsIn[edgeVertex0],
scalarsIn[edgeVertex1],
interpolant);
const vtkm::Vec<vtkm::FloatDefault,3> edgePCoord0 =
vtkm::exec::ParametricCoordinatesPoint(
fieldIn.GetNumberOfComponents(), edgeVertex0, shape, *this);
const vtkm::Vec<vtkm::FloatDefault,3> edgePCoord1 =
vtkm::exec::ParametricCoordinatesPoint(
fieldIn.GetNumberOfComponents(), edgeVertex1, shape, *this);
const vtkm::Vec<vtkm::FloatDefault,3> interpPCoord =
vtkm::Lerp(edgePCoord0, edgePCoord1, interpolant);
normalsOut[triVertex] =
vtkm::Normal(vtkm::exec::CellDerivative(
fieldIn, coords, interpPCoord, shape, *this));
}
vtkm::Vec<FieldType, 3> vertex0 = this->Vertices.Get(outputVertId + 0);
vtkm::Vec<FieldType, 3> vertex1 = this->Vertices.Get(outputVertId + 1);
vtkm::Vec<FieldType, 3> vertex2 = this->Vertices.Get(outputVertId + 2);
vtkm::Vec<FieldType, 3> curNorm = vtkm::Cross(vertex1-vertex0, vertex2-vertex0);
vtkm::Normalize(curNorm);
this->Normals.Set(outputVertId + 0, curNorm);
this->Normals.Set(outputVertId + 1, curNorm);
this->Normals.Set(outputVertId + 2, curNorm);
}
private:
const FieldType Isovalue;
ScatterType Scatter;
};
@ -276,68 +232,49 @@ public:
vtkm::cont::ArrayHandle< vtkm::Vec<CoordinateType,3> > normalsArray,
vtkm::cont::ArrayHandle<FieldType> scalarsArray)
{
typedef typename vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> DeviceAlgorithms;
// Set up the Marching Cubes case tables
vtkm::cont::ArrayHandle<vtkm::Id> vertexTableArray =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numVerticesTable,
vtkm::cont::ArrayHandle<vtkm::IdComponent> numTrianglesTable =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numTrianglesTable,
256);
vtkm::cont::ArrayHandle<vtkm::Id> triangleTableArray =
vtkm::cont::ArrayHandle<vtkm::IdComponent> triangleTableArray =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::triTable,
256*16);
typedef vtkm::exec::ExecutionWholeArrayConst<vtkm::IdComponent, VTKM_DEFAULT_STORAGE_TAG, DeviceAdapter>
TableArrayExecObjectType;
// Call the ClassifyCell functor to compute the Marching Cubes case numbers
// for each cell, and the number of vertices to be generated
ClassifyCell classifyCell(vertexTableArray.PrepareForInput(DeviceAdapter()),
isovalue);
ClassifyCell classifyCell(isovalue);
typedef typename vtkm::worklet::DispatcherMapTopology<
ClassifyCell,
DeviceAdapter> ClassifyCellDispatcher;
ClassifyCellDispatcher classifyCellDispatcher(classifyCell);
vtkm::cont::ArrayHandle<vtkm::Id> numOutputTrisPerCell;
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutputTrisPerCell;
classifyCellDispatcher.Invoke( field,
this->DataSet.GetCellSet(0),
numOutputTrisPerCell);
numOutputTrisPerCell,
TableArrayExecObjectType(numTrianglesTable));
// Compute the number of valid input cells and those ids
const vtkm::Id numOutputCells = DeviceAlgorithms::ScanInclusive(numOutputTrisPerCell,
numOutputTrisPerCell);
IsoSurfaceGenerate isosurface(isovalue, numOutputTrisPerCell, DeviceAdapter());
// Terminate if no cells has triangles left
if (numOutputCells == 0) return;
vtkm::cont::ArrayHandle<vtkm::Id> validCellIndicesArray, inputCellIterationNumber;
vtkm::cont::ArrayHandleIndex validCellCountImplicitArray(numOutputCells);
DeviceAlgorithms::UpperBounds(numOutputTrisPerCell,
validCellCountImplicitArray,
validCellIndicesArray);
numOutputTrisPerCell.ReleaseResources();
// Compute for each output triangle what iteration of the input cell generates it
DeviceAlgorithms::LowerBounds(validCellIndicesArray,
validCellIndicesArray,
inputCellIterationNumber);
// Generate a single triangle per cell
const vtkm::Id numTotalVertices = numOutputCells * 3;
IsoSurfaceGenerate isosurface(isovalue,
this->CDims,
triangleTableArray.PrepareForInput(DeviceAdapter()),
field,
field,
verticesArray.PrepareForOutput(numTotalVertices, DeviceAdapter()),
normalsArray.PrepareForOutput(numTotalVertices, DeviceAdapter()),
scalarsArray.PrepareForOutput(numTotalVertices, DeviceAdapter())
);
typedef typename vtkm::worklet::DispatcherMapField< IsoSurfaceGenerate,
DeviceAdapter> IsoSurfaceDispatcher;
IsoSurfaceDispatcher isosurfaceDispatcher(isosurface);
isosurfaceDispatcher.Invoke(validCellIndicesArray,
inputCellIterationNumber);
vtkm::worklet::DispatcherMapTopology<IsoSurfaceGenerate, DeviceAdapter>
isosurfaceDispatcher(isosurface);
isosurfaceDispatcher.Invoke(
// Currently forcing cell set to be structured. Eventually we should
// relax this as we support other grid types.
this->DataSet.GetCellSet(0).ResetCellSetList(
vtkm::ListTagBase<vtkm::cont::CellSetStructured<3> >()),
field,
this->DataSet.GetCoordinateSystem(0).GetData(),
vtkm::cont::make_ArrayHandleGroupVec<3>(verticesArray),
field, // This is silly. The field will interp to isovalue
vtkm::cont::make_ArrayHandleGroupVec<3>(scalarsArray),
vtkm::cont::make_ArrayHandleGroupVec<3>(normalsArray),
TableArrayExecObjectType(triangleTableArray)
);
}
};

498
vtkm/worklet/MarchingCubesDataTables.h
View File

@ -27,267 +27,267 @@ namespace worklet {
namespace internal {
const vtkm::Id numVerticesTable[256] = {
const vtkm::IdComponent numTrianglesTable[256] = {
0,
1,
1,
2,
1,
2,
2,
3,
1,
2,
2,
3,
2,
3,
3,
6,
2,
1,
2,
2,
3,
6,
6,
9,
3,
6,
6,
9,
6,
9,
9,
6,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
6,
9,
12,
12,
9,
9,
12,
12,
9,
12,
15,
15,
6,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
12,
15,
15,
12,
6,
9,
9,
12,
9,
12,
6,
9,
9,
12,
12,
15,
12,
15,
9,
6,
9,
12,
12,
9,
12,
15,
9,
6,
12,
15,
15,
12,
15,
6,
12,
2,
3,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
12,
9,
12,
12,
15,
9,
6,
12,
9,
12,
9,
15,
6,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
12,
15,
15,
12,
9,
12,
12,
9,
12,
15,
15,
12,
12,
9,
15,
6,
15,
12,
6,
3,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
6,
9,
9,
6,
9,
12,
12,
15,
12,
15,
15,
6,
12,
9,
15,
12,
9,
6,
12,
3,
9,
12,
12,
15,
12,
15,
9,
12,
12,
15,
15,
6,
9,
12,
6,
3,
6,
9,
9,
6,
9,
12,
6,
3,
9,
6,
12,
3,
6,
4,
2,
3,
3,
4,
3,
4,
4,
3,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
2,
3,
4,
4,
3,
3,
4,
4,
3,
4,
5,
5,
2,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
4,
5,
5,
4,
2,
3,
3,
4,
3,
4,
2,
3,
3,
4,
4,
5,
4,
5,
3,
2,
3,
4,
4,
3,
4,
5,
3,
2,
4,
5,
5,
4,
5,
2,
4,
1,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
4,
3,
4,
4,
5,
3,
2,
4,
3,
4,
3,
5,
2,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
4,
5,
5,
4,
3,
4,
4,
3,
4,
5,
5,
4,
4,
3,
5,
2,
5,
4,
2,
1,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
2,
3,
3,
2,
3,
4,
4,
5,
4,
5,
5,
2,
4,
3,
5,
4,
3,
2,
4,
1,
3,
4,
4,
5,
4,
5,
3,
4,
4,
5,
5,
2,
3,
4,
2,
1,
2,
3,
3,
2,
3,
4,
2,
1,
3,
2,
4,
1,
2,
1,
1,
0,
};
const vtkm::Id triTable[256*16] =
const vtkm::IdComponent triTable[256*16] =
{
#define X -1
X, X, X, X, X, X, X, X, X, X, X, X, X, X, X, X,

280
vtkm/worklet/ScatterCounting.h Normal file
View File

@ -0,0 +1,280 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_ScatterCounting_h
#define vtk_m_worklet_ScatterCounting_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/ErrorControlBadValue.h>
#include <vtkm/exec/FunctorBase.h>
#include <sstream>
namespace vtkm {
namespace worklet {
namespace detail {
template<typename Device>
struct ReverseInputToOutputMapKernel : vtkm::exec::FunctorBase
{
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::PortalConst
InputMapType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::Portal
OutputMapType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::IdComponent>::ExecutionTypes<Device>::Portal
VisitType;
InputMapType InputToOutputMap;
OutputMapType OutputToInputMap;
VisitType Visit;
vtkm::Id OutputSize;
VTKM_CONT_EXPORT
ReverseInputToOutputMapKernel(const InputMapType &inputToOutputMap,
const OutputMapType &outputToInputMap,
const VisitType &visit,
vtkm::Id outputSize)
: InputToOutputMap(inputToOutputMap),
OutputToInputMap(outputToInputMap),
Visit(visit),
OutputSize(outputSize)
{ }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex) const
{
vtkm::Id outputStartIndex;
if (inputIndex > 0)
{
outputStartIndex = this->InputToOutputMap.Get(inputIndex-1);
}
else
{
outputStartIndex = 0;
}
vtkm::Id outputEndIndex = this->InputToOutputMap.Get(inputIndex);
vtkm::IdComponent visitIndex = 0;
for (vtkm::Id outputIndex = outputStartIndex;
outputIndex < outputEndIndex;
outputIndex++)
{
this->OutputToInputMap.Set(outputIndex, inputIndex);
this->Visit.Set(outputIndex, visitIndex);
visitIndex++;
}
}
};
template<typename Device>
struct SubtractToVisitIndexKernel : vtkm::exec::FunctorBase
{
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::PortalConst
StartsOfGroupsType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::IdComponent>::ExecutionTypes<Device>::Portal
VisitType;
StartsOfGroupsType StartsOfGroups;
VisitType Visit;
VTKM_CONT_EXPORT
SubtractToVisitIndexKernel(const StartsOfGroupsType &startsOfGroups,
const VisitType &visit)
: StartsOfGroups(startsOfGroups), Visit(visit)
{ }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex) const
{
vtkm::Id startOfGroup = this->StartsOfGroups.Get(inputIndex);
vtkm::IdComponent visitIndex =
static_cast<vtkm::IdComponent>(inputIndex - startOfGroup);
this->Visit.Set(inputIndex, visitIndex);
}
};
} // namespace detail
/// \brief A scatter that maps input to some numbers of output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterCounting establishes a 1 to
/// N mapping from input to output. That is, every input element generates 0 or
/// more output elements associated with it. The output elements are grouped by
/// the input associated.
///
/// A counting scatter takes an array of counts for each input. The data is
/// taken in the constructor and the index arrays are derived from that. So
/// changing the counts after the scatter is created will have no effect.
///
struct ScatterCounting
{
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
ScatterCounting(const CountArrayType &countArray, Device)
{
this->BuildArrays(countArray, Device());
}
typedef vtkm::cont::ArrayHandle<vtkm::Id> OutputToInputMapType;
template<typename RangeType>
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(RangeType) const
{
return this->OutputToInputMap;
}
typedef vtkm::cont::ArrayHandle<vtkm::IdComponent> VisitArrayType;
template<typename RangeType>
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(RangeType) const
{
return this->VisitArray;
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id inputRange) const
{
if (inputRange != this->InputRange)
{
std::stringstream msg;
msg << "ScatterCounting initialized with input domain of size "
<< this->InputRange
<< " but used with a worklet invoke of size "
<< inputRange << std::endl;
throw vtkm::cont::ErrorControlBadValue(msg.str());
}
return this->VisitArray.GetNumberOfValues();
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id3 inputRange) const
{
return this->GetOutputRange(inputRange[0]*inputRange[1]*inputRange[2]);
}
private:
vtkm::Id InputRange;
OutputToInputMapType OutputToInputMap;
VisitArrayType VisitArray;
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
void BuildArrays(const CountArrayType &count, Device)
{
VTKM_IS_ARRAY_HANDLE(CountArrayType);
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
this->InputRange = count.GetNumberOfValues();
// Currently we are treating the input to output map as a temporary
// variable. However, it is possible that this could, be useful elsewhere,
// so we may want to save this and make it available.
//
// The input to output map is actually built off by one. The first entry
// is actually for the second value. The last entry is the total number of
// output. This off-by-one is so that an upper bound find will work when
// building the output to input map.
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap;
vtkm::Id outputSize =
vtkm::cont::DeviceAdapterAlgorithm<Device>::ScanInclusive(
vtkm::cont::make_ArrayHandleCast(count, vtkm::Id()),
inputToOutputMap);
// We have implemented two different ways to compute the output to input
// map. The first way is to use a binary search on each output index into
// the input map. The second way is to schedule on each input and
// iteratively fill all the output indices for that input. The first way is
// faster for output sizes that are small relative to the input (typical in
// Marching Cubes, for example) and also tends to be well load balanced.
// The second way is faster for larger outputs (typical in triangulation,
// for example). We will use the first method for small output sizes and
// the second for large output sizes. Toying with this might be a good
// place for optimization.
if (outputSize < this->InputRange)
{
this->BuildOutputToInputMapWithFind(
outputSize, inputToOutputMap, Device());
}
else
{
this->BuildOutputToInputMapWithIterate(
outputSize, inputToOutputMap, Device());
}
}
template<typename Device>
VTKM_CONT_EXPORT
void BuildOutputToInputMapWithFind(
vtkm::Id outputSize,
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap,
Device)
{
vtkm::cont::ArrayHandleIndex outputIndices(outputSize);
vtkm::cont::DeviceAdapterAlgorithm<Device>::UpperBounds(
inputToOutputMap, outputIndices, this->OutputToInputMap);
// Do not need this anymore.
inputToOutputMap.ReleaseResources();
vtkm::cont::ArrayHandle<vtkm::Id> startsOfGroups;
// This find gives the index of the start of a group.
vtkm::cont::DeviceAdapterAlgorithm<Device>::LowerBounds(
this->OutputToInputMap, this->OutputToInputMap, startsOfGroups);
detail::SubtractToVisitIndexKernel<Device>
kernel(startsOfGroups.PrepareForInput(Device()),
this->VisitArray.PrepareForOutput(outputSize, Device()));
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(kernel, outputSize);
}
template<typename Device>
VTKM_CONT_EXPORT
void BuildOutputToInputMapWithIterate(
vtkm::Id outputSize,
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap,
Device)
{
detail::ReverseInputToOutputMapKernel<Device>
kernel(inputToOutputMap.PrepareForInput(Device()),
this->OutputToInputMap.PrepareForOutput(outputSize, Device()),
this->VisitArray.PrepareForOutput(outputSize, Device()),
outputSize);
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(
kernel, inputToOutputMap.GetNumberOfValues());
}
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterCounting_h

77
vtkm/worklet/ScatterIdentity.h Normal file
View File

@ -0,0 +1,77 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_ScatterIdentity_h
#define vtk_m_worklet_ScatterIdentity_h
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/cont/ArrayHandleIndex.h>
namespace vtkm {
namespace worklet {
/// \brief A scatter that maps input directly to output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterIdentity establishes a 1 to
/// 1 mapping from input to output (and vice versa). That is, every input
/// element generates one output element associated with it. This is the
/// default for basic maps.
///
struct ScatterIdentity
{
typedef vtkm::cont::ArrayHandleIndex OutputToInputMapType;
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(vtkm::Id outputRange) const
{
return OutputToInputMapType(outputRange);
}
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(vtkm::Id3 outputRange) const
{
return this->GetOutputToInputMap(
outputRange[0]*outputRange[1]*outputRange[2]);
}
typedef vtkm::cont::ArrayHandleConstant<vtkm::IdComponent> VisitArrayType;
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(vtkm::Id outputRange) const
{
return VisitArrayType(1, outputRange);
}
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(vtkm::Id3 outputRange) const
{
return this->GetVisitArray(outputRange[0]*outputRange[1]*outputRange[2]);
}
template<typename RangeType>
VTKM_CONT_EXPORT
RangeType GetOutputRange(RangeType inputRange) const
{
return inputRange;
}
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterIdentity_h

121
vtkm/worklet/ScatterUniform.h Normal file
View File

@ -0,0 +1,121 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_ScatterUniform_h
#define vtk_m_worklet_ScatterUniform_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
namespace vtkm {
namespace worklet {
namespace detail {
struct FunctorModulus
{
vtkm::IdComponent Modulus;
VTKM_EXEC_CONT_EXPORT
FunctorModulus(vtkm::IdComponent modulus = 1)
: Modulus(modulus)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent operator()(vtkm::Id index) const
{
return static_cast<vtkm::IdComponent>(index % this->Modulus);
}
};
struct FunctorDiv
{
vtkm::Id Divisor;
VTKM_EXEC_CONT_EXPORT
FunctorDiv(vtkm::Id divisor = 1)
: Divisor(divisor)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::Id operator()(vtkm::Id index) const
{
return index / this->Divisor;
}
};
}
/// \brief A scatter that maps input to some constant numbers of output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterUniform establishes a 1 to N
/// mapping from input to output. That is, every input element generates N
/// elements associated with it where N is the same for every input. The output
/// elements are grouped by the input associated.
///
struct ScatterUniform
{
VTKM_CONT_EXPORT
ScatterUniform(vtkm::IdComponent numOutputsPerInput)
: NumOutputsPerInput(numOutputsPerInput)
{ }
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id inputRange) const
{
return inputRange * this->NumOutputsPerInput;
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id3 inputRange) const
{
return this->GetOutputRange(inputRange[0]*inputRange[1]*inputRange[2]);
}
typedef vtkm::cont::ArrayHandleImplicit<vtkm::Id, detail::FunctorDiv>
OutputToInputMapType;
template<typename RangeType>
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(RangeType inputRange) const
{
return OutputToInputMapType(detail::FunctorDiv(this->NumOutputsPerInput),
this->GetOutputRange(inputRange));
}
typedef vtkm::cont::ArrayHandleImplicit<vtkm::IdComponent, detail::FunctorModulus>
VisitArrayType;
template<typename RangeType>
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(RangeType inputRange) const
{
return VisitArrayType(detail::FunctorModulus(this->NumOutputsPerInput),
this->GetOutputRange(inputRange));
}
private:
vtkm::IdComponent NumOutputsPerInput;
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterUniform_h

324
vtkm/worklet/TetrahedralizeExplicitGrid.h
View File

@ -6,9 +6,9 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -21,20 +21,21 @@
#ifndef vtk_m_worklet_TetrahedralizeExplicitGrid_h
#define vtk_m_worklet_TetrahedralizeExplicitGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/CellSetExplicit.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterCounting.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/worklet/internal/TriangulateTables.h>
namespace vtkm {
namespace worklet {
@ -52,26 +53,22 @@ public:
{
public:
typedef void ControlSignature(FieldIn<> shapes,
FieldIn<> numIndices,
FieldIn<> numPoints,
ExecObject tables,
FieldOut<> triangleCount);
typedef void ExecutionSignature(_1,_2,_3);
typedef _4 ExecutionSignature(_1,_2,_3);
typedef _1 InputDomain;
VTKM_CONT_EXPORT
TrianglesPerCell() {}
VTKM_EXEC_EXPORT
void operator()(const vtkm::UInt8 &shape,
const vtkm::IdComponent &numIndices,
vtkm::Id &triangleCount) const
vtkm::IdComponent operator()(
vtkm::UInt8 shape,
vtkm::IdComponent numPoints,
const vtkm::worklet::internal::TriangulateTablesExecutionObject<DeviceAdapter> &tables) const
{
if (shape == vtkm::CELL_SHAPE_TRIANGLE)
triangleCount = 1;
else if (shape == vtkm::CELL_SHAPE_QUAD)
triangleCount = 2;
else if (shape == vtkm::CELL_SHAPE_POLYGON)
triangleCount = numIndices - 2;
else triangleCount = 0;
return tables.GetCount(vtkm::CellShapeTagGeneric(shape), numPoints);
}
};
@ -82,26 +79,20 @@ public:
{
public:
typedef void ControlSignature(FieldIn<> shapes,
ExecObject tables,
FieldOut<> triangleCount);
typedef void ExecutionSignature(_1,_2);
typedef _3 ExecutionSignature(_1, _2);
typedef _1 InputDomain;
VTKM_CONT_EXPORT
TetrahedraPerCell() {}
VTKM_EXEC_EXPORT
void operator()(const vtkm::UInt8 &shape,
vtkm::Id &tetrahedraCount) const
vtkm::IdComponent operator()(
vtkm::UInt8 shape,
const vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<DeviceAdapter> &tables) const
{
if (shape == vtkm::CELL_SHAPE_TETRA)
tetrahedraCount = 1;
else if (shape == vtkm::CELL_SHAPE_HEXAHEDRON)
tetrahedraCount = 5;
else if (shape == vtkm::CELL_SHAPE_WEDGE)
tetrahedraCount = 3;
else if (shape == vtkm::CELL_SHAPE_PYRAMID)
tetrahedraCount = 2;
else tetrahedraCount = 0;
return tables.GetCount(vtkm::CellShapeTagGeneric(shape));
}
};
@ -112,53 +103,46 @@ public:
class TriangulateCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldInTo<> triangleOffset,
FieldInTo<> numIndices,
TopologyIn topology,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2,_4, CellShape, FromIndices);
typedef _3 InputDomain;
typedef void ControlSignature(TopologyIn topology,
ExecObject tables,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(CellShape, PointIndices, _2, _3, VisitIndex);
typedef _1 InputDomain;
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
TriangulateCell() {}
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CountArrayType>
VTKM_CONT_EXPORT
TriangulateCell(const CountArrayType &countArray)
: Scatter(countArray, DeviceAdapter())
{ }
// Each cell produces triangles and write result at the offset
template<typename CellShapeTag, typename CellNodeVecType>
template<typename CellShapeTag,
typename ConnectivityInVec,
typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &offset,
const vtkm::Id &numIndices,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity,
CellShapeTag shape,
const CellNodeVecType &cellNodeIds) const
void operator()(
CellShapeTag shape,
const ConnectivityInVec &connectivityIn,
const vtkm::worklet::internal::TriangulateTablesExecutionObject<DeviceAdapter> &tables,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Offset is in triangles, 3 vertices per triangle needed
vtkm::Id startIndex = offset * 3;
if (shape.Id == vtkm::CELL_SHAPE_TRIANGLE)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
}
else if (shape.Id == vtkm::CELL_SHAPE_QUAD)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
}
else if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
for (vtkm::IdComponent tri = 0; tri < numIndices-2; tri++)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[tri+1]);
connectivity.Set(startIndex++, cellNodeIds[tri+2]);
}
}
vtkm::Vec<vtkm::IdComponent,3> triIndices =
tables.GetIndices(shape, visitIndex);
connectivityOut[0] = connectivityIn[triIndices[0]];
connectivityOut[1] = connectivityIn[triIndices[1]];
connectivityOut[2] = connectivityIn[triIndices[2]];
}
private:
ScatterType Scatter;
};
//
@ -168,89 +152,46 @@ public:
class TetrahedralizeCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldInTo<> tetraOffset,
TopologyIn topology,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_3, CellShape, FromIndices);
typedef _2 InputDomain;
typedef void ControlSignature(TopologyIn topology,
ExecObject tables,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(CellShape, PointIndices, _2, _3, VisitIndex);
typedef _1 InputDomain;
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
TetrahedralizeCell() {}
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CellArrayType>
VTKM_CONT_EXPORT
TetrahedralizeCell(const CellArrayType &cellArray)
: Scatter(cellArray, DeviceAdapter())
{ }
// Each cell produces tetrahedra and write result at the offset
template<typename CellShapeTag, typename CellNodeVecType>
template<typename CellShapeTag,
typename ConnectivityInVec,
typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &offset,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity,
CellShapeTag shape,
const CellNodeVecType &cellNodeIds) const
void operator()(CellShapeTag shape,
const ConnectivityInVec &connectivityIn,
const vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<DeviceAdapter> &tables,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Offset is in tetrahedra, 4 vertices per tetrahedron needed
vtkm::Id startIndex = offset * 4;
if (shape.Id == vtkm::CELL_SHAPE_TETRA)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
}
else if (shape.Id == vtkm::CELL_SHAPE_HEXAHEDRON)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[5]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[7]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
else if (shape.Id == vtkm::CELL_SHAPE_WEDGE)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[5]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
else if (shape.Id == vtkm::CELL_SHAPE_PYRAMID)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
vtkm::Vec<vtkm::IdComponent,4> tetIndices =
tables.GetIndices(shape, visitIndex);
connectivityOut[0] = connectivityIn[tetIndices[0]];
connectivityOut[1] = connectivityIn[tetIndices[1]];
connectivityOut[2] = connectivityIn[tetIndices[2]];
connectivityOut[3] = connectivityIn[tetIndices[3]];
}
private:
ScatterType Scatter;
};
//
@ -270,8 +211,6 @@ public:
//
void Run()
{
typedef typename vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> DeviceAlgorithms;
// Cell sets belonging to input and output datasets
vtkm::cont::CellSetExplicit<> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetExplicit<> >();
@ -286,56 +225,61 @@ public:
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
vtkm::cont::ArrayHandle<vtkm::IdComponent> inNumIndices = inCellSet.GetNumIndicesArray(
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
vtkm::cont::ArrayHandle<vtkm::Id> inConn = inCellSet.GetConnectivityArray(
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::Id> numOutCellArray;
vtkm::IdComponent verticesPerOutCell = 0;
// Output topology
vtkm::cont::ArrayHandle<vtkm::Id> outConnectivity;
if (dimensionality == 2)
{
verticesPerOutCell = 3;
vtkm::worklet::DispatcherMapField<TrianglesPerCell> trianglesPerCellDispatcher;
trianglesPerCellDispatcher.Invoke(inShapes, inNumIndices, numOutCellArray);
vtkm::worklet::internal::TriangulateTables tables;
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutCellArray;
vtkm::worklet::DispatcherMapField<TrianglesPerCell,DeviceAdapter>
triPerCellDispatcher;
triPerCellDispatcher.Invoke(inShapes,
inNumIndices,
tables.PrepareForInput(DeviceAdapter()),
numOutCellArray);
// Build new cells
TriangulateCell triangulateWorklet(numOutCellArray);
vtkm::worklet::DispatcherMapTopology<TriangulateCell,DeviceAdapter>
triangulateDispatcher(triangulateWorklet);
triangulateDispatcher.Invoke(
inCellSet,
tables.PrepareForInput(DeviceAdapter()),
vtkm::cont::make_ArrayHandleGroupVec<3>(outConnectivity));
}
else if (dimensionality == 3)
{
verticesPerOutCell = 4;
vtkm::worklet::DispatcherMapField<TetrahedraPerCell> tetrahedraPerCellDispatcher;
tetrahedraPerCellDispatcher.Invoke(inShapes, numOutCellArray);
vtkm::worklet::internal::TetrahedralizeTables tables;
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutCellArray;
vtkm::worklet::DispatcherMapField<TetrahedraPerCell,DeviceAdapter>
tetPerCellDispatcher;
tetPerCellDispatcher.Invoke(inShapes,
tables.PrepareForInput(DeviceAdapter()),
numOutCellArray);
// Build new cells
TetrahedralizeCell tetrahedralizeWorklet(numOutCellArray);
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell,DeviceAdapter>
tetrahedralizeDispatcher(tetrahedralizeWorklet);
tetrahedralizeDispatcher.Invoke(
inCellSet,
tables.PrepareForInput(DeviceAdapter()),
vtkm::cont::make_ArrayHandleGroupVec<4>(outConnectivity));
}
// Number of output cells and number of vertices needed
vtkm::cont::ArrayHandle<vtkm::Id> cellOffset;
vtkm::Id numberOfOutCells = DeviceAlgorithms::ScanExclusive(numOutCellArray,
cellOffset);
vtkm::Id numberOfOutIndices = numberOfOutCells * verticesPerOutCell;
// Information needed to build the output cell set
vtkm::cont::ArrayHandle<vtkm::Id> connectivity;
connectivity.Allocate(numberOfOutIndices);
// Call worklet to compute the connectivity
if (dimensionality == 2)
else
{
vtkm::worklet::DispatcherMapTopology<TriangulateCell> triangulateCellDispatcher;
triangulateCellDispatcher.Invoke(
cellOffset,
inNumIndices,
inCellSet,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
}
else if (dimensionality == 3)
{
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell> tetrahedralizeCellDispatcher;
tetrahedralizeCellDispatcher.Invoke(
cellOffset,
inCellSet,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
throw vtkm::cont::ErrorControlBadValue(
"Unsupported dimensionality for TetrahedralizeExplicitGrid.");
}
// Add cells to output cellset
cellSet.Fill(connectivity);
cellSet.Fill(outConnectivity);
}
};

296
vtkm/worklet/TetrahedralizeUniformGrid.h
View File

@ -6,9 +6,9 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -21,23 +21,51 @@
#ifndef vtk_m_worklet_TetrahedralizeUniformGrid_h
#define vtk_m_worklet_TetrahedralizeUniformGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/CellSetSingleType.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/ErrorControlBadValue.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterUniform.h>
#include <vtkm/worklet/WorkletMapTopology.h>
namespace vtkm {
namespace worklet {
namespace detail {
VTKM_EXEC_CONSTANT_EXPORT
const static vtkm::IdComponent StructuredTriangleIndices[2][3] = {
{ 0, 1, 2 },
{ 0, 2, 3 }
};
VTKM_EXEC_CONSTANT_EXPORT
const static vtkm::IdComponent StructuredTetrahedronIndices[2][5][4] = {
{
{ 0, 1, 3, 4 },
{ 1, 4, 5, 6 },
{ 1, 4, 6, 3 },
{ 1, 3, 6, 2 },
{ 3, 6, 7, 4 }
},
{
{ 2, 1, 5, 0 },
{ 0, 2, 3, 7 },
{ 2, 5, 6, 7 },
{ 0, 7, 4, 5 },
{ 0, 2, 7, 5 }
}
};
} // namespace detail
/// \brief Compute the tetrahedralize cells for a uniform grid data set
template <typename DeviceAdapter>
class TetrahedralizeFilterUniformGrid
@ -48,63 +76,35 @@ public:
// Worklet to turn quads into triangles
// Vertices remain the same and each cell is processed with needing topology
//
class TriangulateCell : public vtkm::worklet::WorkletMapField
class TriangulateCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2);
typedef void ControlSignature(TopologyIn topology,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(PointIndices, _2, VisitIndex);
typedef _1 InputDomain;
vtkm::Id xdim, ydim;
typedef vtkm::worklet::ScatterUniform ScatterType;
VTKM_CONT_EXPORT
TriangulateCell(const vtkm::Id2 &cdims) :
xdim(cdims[0]), ydim(cdims[1])
ScatterType GetScatter() const
{
return ScatterType(2);
}
// Each hexahedron cell produces five tetrahedron cells
VTKM_CONT_EXPORT
TriangulateCell()
{ }
// Each quad cell produces 2 triangle cells
template<typename ConnectivityInVec, typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(vtkm::Id &inputCellId,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity) const
void operator()(const ConnectivityInVec &connectivityIn,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Calculate the i,j indices for this input cell id
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
// Calculate the type of triangle generated because it alternates
vtkm::Id indexType = (x + y) % 2;
// Compute indices for the four vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1);
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
// Set the triangles for this cell based on vertex index and index type of cell
// 2 triangles per quad, 3 indices per triangle
vtkm::Id startIndex = inputCellId * 2 * 3;
if (indexType == 0)
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
}
else
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
}
connectivityOut[0] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][0]];
connectivityOut[1] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][1]];
connectivityOut[2] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][2]];
}
};
@ -112,105 +112,45 @@ public:
// Worklet to turn hexahedra into tetrahedra
// Vertices remain the same and each cell is processed with needing topology
//
class TetrahedralizeCell : public vtkm::worklet::WorkletMapField
class TetrahedralizeCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2);
typedef void ControlSignature(TopologyIn topology,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(PointIndices, _2, ThreadIndices);
typedef _1 InputDomain;
vtkm::Id xdim, ydim, zdim;
const vtkm::Id cellsPerLayer, pointsPerLayer;
typedef vtkm::worklet::ScatterUniform ScatterType;
VTKM_CONT_EXPORT
TetrahedralizeCell(const vtkm::Id3 &cdims) :
xdim(cdims[0]), ydim(cdims[1]), zdim(cdims[2]),
cellsPerLayer(xdim * ydim),
pointsPerLayer((xdim+1) * (ydim+1))
ScatterType GetScatter() const
{
return ScatterType(5);
}
VTKM_CONT_EXPORT
TetrahedralizeCell()
{ }
// Each hexahedron cell produces five tetrahedron cells
template<typename ConnectivityInVec,
typename ConnectivityOutVec,
typename ThreadIndicesType>
VTKM_EXEC_EXPORT
void operator()(vtkm::Id &inputCellId,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity) const
void operator()(const ConnectivityInVec &connectivityIn,
ConnectivityOutVec &connectivityOut,
const ThreadIndicesType threadIndices) const
{
// Calculate the i,j,k indices for this input cell id
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
const vtkm::Id z = inputCellId / cellsPerLayer;
vtkm::Id3 inputIndex = threadIndices.GetInputIndex3D();
// Calculate the type of tetrahedron generated because it alternates
vtkm::Id indexType = (x + y + z) % 2;
vtkm::Id indexType = (inputIndex[0] + inputIndex[1] + inputIndex[2]) % 2;
// Compute indices for the eight vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1) + z * pointsPerLayer;
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
const vtkm::Id i4 = i0 + pointsPerLayer;
const vtkm::Id i5 = i1 + pointsPerLayer;
const vtkm::Id i6 = i2 + pointsPerLayer;
const vtkm::Id i7 = i3 + pointsPerLayer;
vtkm::IdComponent visitIndex = threadIndices.GetVisitIndex();
// Set the tetrahedra for this cell based on vertex index and index type of cell
// 5 tetrahedra per hexahedron, 4 indices per tetrahedron
vtkm::Id startIndex = inputCellId * 5 * 4;
if (indexType == 0)
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i4);
}
else
{
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i5);
}
connectivityOut[0] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][0]];
connectivityOut[1] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][1]];
connectivityOut[2] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][2]];
connectivityOut[3] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][3]];
}
};
@ -236,64 +176,28 @@ public:
vtkm::cont::CellSetSingleType<> & cellSet =
OutDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetSingleType<> >();
// Get dimensionality from the explicit cell set
vtkm::IdComponent dim = cellSet.GetDimensionality();
vtkm::Id outCellsPerInCell = 0;
vtkm::IdComponent verticesPerOutCell = 0;
vtkm::Id numberOfInCells = 0;
vtkm::Id2 cdims2;
vtkm::Id3 cdims3;
// From the uniform dimension get more information
if (dim == 2)
{
outCellsPerInCell = 2;
verticesPerOutCell = 3;
vtkm::cont::CellSetStructured<2> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<2> >();
cdims2 = inCellSet.GetSchedulingRange(vtkm::TopologyElementTagCell());
numberOfInCells = cdims2[0] * cdims2[1];
}
else if (dim == 3)
{
outCellsPerInCell = 5;
verticesPerOutCell = 4;
vtkm::cont::CellSetStructured<3> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<3> >();
cdims3 = inCellSet.GetSchedulingRange(vtkm::TopologyElementTagCell());
numberOfInCells = cdims3[0] * cdims3[1] * cdims3[2];
}
vtkm::Id numberOfOutCells = outCellsPerInCell * numberOfInCells;
vtkm::Id numberOfOutIndices = verticesPerOutCell * numberOfOutCells;
// Cell indices are just counting array
vtkm::cont::ArrayHandleCounting<vtkm::Id> cellIndicesArray(0, 1, numberOfInCells);
// Output dataset depends on dimension and size
vtkm::cont::ArrayHandle<vtkm::Id> connectivity;
connectivity.Allocate(numberOfOutIndices);
// Call the TetrahedralizeCell functor to compute tetrahedra or triangles
if (dim == 2)
if (cellSet.GetDimensionality() == 2)
{
TriangulateCell triangulateCell(cdims2);
vtkm::worklet::DispatcherMapField<TriangulateCell> triangulateCellDispatcher(triangulateCell);
triangulateCellDispatcher.Invoke(
cellIndicesArray,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
vtkm::cont::CellSetStructured<2> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<2> >();
vtkm::worklet::DispatcherMapTopology<TriangulateCell,DeviceAdapter> dispatcher;
dispatcher.Invoke(inCellSet,
vtkm::cont::make_ArrayHandleGroupVec<3>(connectivity));
}
else if (dim == 3)
else if (cellSet.GetDimensionality() == 3)
{
TetrahedralizeCell tetrahedralizeCell(cdims3);
vtkm::worklet::DispatcherMapField<TetrahedralizeCell> tetrahedralizeCellDispatcher(tetrahedralizeCell);
tetrahedralizeCellDispatcher.Invoke(
cellIndicesArray,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
vtkm::cont::CellSetStructured<3> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<3> >();
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell,DeviceAdapter> dispatcher;
dispatcher.Invoke(inCellSet,
vtkm::cont::make_ArrayHandleGroupVec<4>(connectivity));
}
else
{
throw vtkm::cont::ErrorControlBadValue(
"Unsupported dimensionality for TetrahedralizeUniformGrid.");
}
// Add cells to output cellset

2
vtkm/worklet/WorkletMapTopology.h
View File

@ -188,6 +188,8 @@ public:
template<typename TypeList = AllTypes >
struct FieldInOutCell : FieldInOut<TypeList> { };
struct PointIndices : FromIndices { };
};
}

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

@ -19,10 +19,11 @@
##============================================================================
set(headers
ClipTables.h
DispatcherBase.h
DispatcherBaseDetailInvoke.h
TriangulateTables.h
WorkletBase.h
ClipTables.h
)
set_source_files_properties(DispatcherBaseDetailInvoke.h

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

@ -441,19 +441,23 @@ protected:
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id numInstances,
DeviceAdapter tag) const
DeviceAdapter device) const
{
this->InvokeTransportParameters(invocation, numInstances, tag);
this->InvokeTransportParameters(
invocation,
this->Worklet.GetScatter().GetOutputRange(numInstances),
device);
}
template<typename Invocation, typename DeviceAdapter>
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id2 dimensions,
DeviceAdapter tag) const
DeviceAdapter device) const
{
vtkm::Id3 dim3d(dimensions[0], dimensions[1], 1);
this->InvokeTransportParameters(invocation, dim3d, tag);
this->BasicInvoke(invocation,
vtkm::Id3(dimensions[0], dimensions[1], 1),
device);
}
@ -461,9 +465,12 @@ protected:
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id3 dimensions,
DeviceAdapter tag) const
DeviceAdapter device) const
{
this->InvokeTransportParameters(invocation, dimensions, tag);
this->InvokeTransportParameters(
invocation,
this->Worklet.GetScatter().GetOutputRange(dimensions),
device);
}
WorkletType Worklet;
@ -477,7 +484,7 @@ private:
VTKM_CONT_EXPORT
void InvokeTransportParameters(const Invocation &invocation,
RangeType range,
DeviceAdapter tag) const
DeviceAdapter device) const
{
// The first step in invoking a worklet is to transport the arguments to
// the execution environment. The invocation object passed to this function
@ -499,11 +506,21 @@ private:
ExecObjectParameters execObjectParameters =
parameters.StaticTransformCont(TransportFunctorType(range));
// Get the arrays used for scattering input to output.
typename WorkletType::ScatterType::OutputToInputMapType outputToInputMap =
this->Worklet.GetScatter().GetOutputToInputMap(range);
typename WorkletType::ScatterType::VisitArrayType visitArray =
this->Worklet.GetScatter().GetVisitArray(range);
// Replace the parameters in the invocation with the execution object and
// pass to next step of Invoke.
this->InvokeSchedule(invocation.ChangeParameters(execObjectParameters),
range,
tag);
// pass to next step of Invoke. Also add the scatter information.
this->InvokeSchedule(
invocation
.ChangeParameters(execObjectParameters)
.ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device))
.ChangeVisitArray(visitArray.PrepareForInput(device)),
range,
device);
}
template<typename Invocation, typename RangeType, typename DeviceAdapter>

312
vtkm/worklet/internal/TriangulateTables.h Normal file
View File

@ -0,0 +1,312 @@
//============================================================================
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_internal_TriangulateTables_h
#define vtk_m_worklet_internal_TriangulateTables_h
#include <vtkm/CellShape.h>
#include <vtkm/Types.h>
#include <vtkm/exec/ExecutionObjectBase.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageBasic.h>
namespace vtkm {
namespace worklet {
namespace internal {
typedef vtkm::cont::ArrayHandle<vtkm::IdComponent, vtkm::cont::StorageTagBasic>
TriangulateArrayHandle;
static vtkm::IdComponent TriangleCountData[vtkm::NUMBER_OF_CELL_SHAPES] = {
0, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
0, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
0, // 3 = vtkm::CELL_SHAPE_LINE
0, // 4 = vtkm::CELL_SHAPE_POLY_LINE
1, // 5 = vtkm::CELL_SHAPE_TRIANGLE
0, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
0, // 8 = vtkm::CELL_SHAPE_PIXEL
2, // 9 = vtkm::CELL_SHAPE_QUAD
0, // 10 = vtkm::CELL_SHAPE_TETRA
0, // 11 = vtkm::CELL_SHAPE_VOXEL
0, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
0, // 13 = vtkm::CELL_SHAPE_WEDGE
0 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TriangleOffsetData[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
-1, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
-1, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
0, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
1, // 9 = vtkm::CELL_SHAPE_QUAD
-1, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
-1, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
-1, // 13 = vtkm::CELL_SHAPE_WEDGE
-1 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TriangleIndexData[] = {
// vtkm::CELL_SHAPE_TRIANGLE
0, 1, 2,
// vtkm::CELL_SHAPE_QUAD
0, 1, 2,
0, 2, 3
};
template<typename Device>
class TriangulateTablesExecutionObject : public vtkm::exec::ExecutionObjectBase
{
public:
typedef typename TriangulateArrayHandle::ExecutionTypes<Device>::PortalConst
PortalType;
VTKM_EXEC_CONT_EXPORT
TriangulateTablesExecutionObject()
{ }
VTKM_CONT_EXPORT
TriangulateTablesExecutionObject(const TriangulateArrayHandle &counts,
const TriangulateArrayHandle &offsets,
const TriangulateArrayHandle &indices)
: Counts(counts.PrepareForInput(Device())),
Offsets(offsets.PrepareForInput(Device())),
Indices(indices.PrepareForInput(Device()))
{ }
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::IdComponent GetCount(CellShape shape, vtkm::IdComponent numPoints) const
{
if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
return numPoints-2;
}
else
{
return this->Counts.Get(shape.Id);
}
}
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::Vec<vtkm::IdComponent, 3>
GetIndices(CellShape shape, vtkm::IdComponent triangleIndex) const
{
vtkm::Vec<vtkm::IdComponent, 3> triIndices;
if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
triIndices[0] = 0;
triIndices[1] = triangleIndex + 1;
triIndices[2] = triangleIndex + 2;
}
else
{
vtkm::IdComponent offset =
3*(this->Offsets.Get(shape.Id) + triangleIndex);
triIndices[0] = this->Indices.Get(offset + 0);
triIndices[1] = this->Indices.Get(offset + 1);
triIndices[2] = this->Indices.Get(offset + 2);
}
return triIndices;
}
private:
PortalType Counts;
PortalType Offsets;
PortalType Indices;
};
class TriangulateTables
{
public:
VTKM_CONT_EXPORT
TriangulateTables()
: Counts(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleCountData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Offsets(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleOffsetData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Indices(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleIndexData,
sizeof(vtkm::worklet::internal::TriangleIndexData)/sizeof(vtkm::IdComponent)))
{ }
template<typename Device>
vtkm::worklet::internal::TriangulateTablesExecutionObject<Device>
PrepareForInput(Device) const
{
return vtkm::worklet::internal::TriangulateTablesExecutionObject<Device>(
this->Counts, this->Offsets, this->Indices);
}
private:
TriangulateArrayHandle Counts;
TriangulateArrayHandle Offsets;
TriangulateArrayHandle Indices;
};
static vtkm::IdComponent TetrahedronCountData[vtkm::NUMBER_OF_CELL_SHAPES] = {
0, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
0, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
0, // 3 = vtkm::CELL_SHAPE_LINE
0, // 4 = vtkm::CELL_SHAPE_POLY_LINE
0, // 5 = vtkm::CELL_SHAPE_TRIANGLE
0, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
0, // 7 = vtkm::CELL_SHAPE_POLYGON
0, // 8 = vtkm::CELL_SHAPE_PIXEL
0, // 9 = vtkm::CELL_SHAPE_QUAD
1, // 10 = vtkm::CELL_SHAPE_TETRA
0, // 11 = vtkm::CELL_SHAPE_VOXEL
5, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
3, // 13 = vtkm::CELL_SHAPE_WEDGE
2 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TetrahedronOffsetData[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
-1, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
-1, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
-1, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
-1, // 9 = vtkm::CELL_SHAPE_QUAD
0, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
1, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
6, // 13 = vtkm::CELL_SHAPE_WEDGE
9 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TetrahedronIndexData[] = {
// vtkm::CELL_SHAPE_TETRA
0, 1, 2, 3,
// vtkm::CELL_SHAPE_HEXAHEDRON
0, 1, 3, 4,
1, 4, 5, 6,
1, 4, 6, 3,
1, 3, 6, 2,
3, 6, 7, 4,
// vtkm::CELL_SHAPE_WEDGE
0, 1, 2, 4,
3, 4, 5, 2,
0, 2, 3, 4,
// vtkm::CELL_SHAPE_PYRAMID
0, 1, 2, 4,
0, 2, 3, 4
};
template<typename Device>
class TetrahedralizeTablesExecutionObject
: public vtkm::exec::ExecutionObjectBase
{
public:
typedef typename TriangulateArrayHandle::ExecutionTypes<Device>::PortalConst
PortalType;
VTKM_EXEC_CONT_EXPORT
TetrahedralizeTablesExecutionObject()
{ }
VTKM_CONT_EXPORT
TetrahedralizeTablesExecutionObject(const TriangulateArrayHandle &counts,
const TriangulateArrayHandle &offsets,
const TriangulateArrayHandle &indices)
: Counts(counts.PrepareForInput(Device())),
Offsets(offsets.PrepareForInput(Device())),
Indices(indices.PrepareForInput(Device()))
{ }
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::IdComponent GetCount(CellShape shape) const
{
return this->Counts.Get(shape.Id);
}
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::Vec<vtkm::IdComponent, 4>
GetIndices(CellShape shape, vtkm::IdComponent tetrahedronIndex) const
{
vtkm::Vec<vtkm::IdComponent, 4> tetIndices;
vtkm::IdComponent offset =
4*(this->Offsets.Get(shape.Id) + tetrahedronIndex);
tetIndices[0] = this->Indices.Get(offset + 0);
tetIndices[1] = this->Indices.Get(offset + 1);
tetIndices[2] = this->Indices.Get(offset + 2);
tetIndices[3] = this->Indices.Get(offset + 3);
return tetIndices;
}
private:
PortalType Counts;
PortalType Offsets;
PortalType Indices;
};
class TetrahedralizeTables
{
public:
VTKM_CONT_EXPORT
TetrahedralizeTables()
: Counts(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronCountData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Offsets(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronOffsetData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Indices(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronIndexData,
sizeof(vtkm::worklet::internal::TetrahedronIndexData)/sizeof(vtkm::IdComponent)))
{ }
template<typename Device>
vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<Device>
PrepareForInput(Device) const
{
return vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<Device>(
this->Counts, this->Offsets, this->Indices);
}
private:
TriangulateArrayHandle Counts;
TriangulateArrayHandle Offsets;
TriangulateArrayHandle Indices;
};
}
}
}
#endif //vtk_m_worklet_internal_TriangulateTables_h

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

@ -25,13 +25,17 @@
#include <vtkm/exec/FunctorBase.h>
#include <vtkm/exec/arg/BasicArg.h>
#include <vtkm/exec/arg/FetchTagExecObject.h>
#include <vtkm/exec/arg/ThreadIndices.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/exec/arg/VisitIndex.h>
#include <vtkm/exec/arg/WorkIndex.h>
#include <vtkm/cont/arg/ControlSignatureTagBase.h>
#include <vtkm/cont/arg/TransportTagExecObject.h>
#include <vtkm/cont/arg/TypeCheckTagExecObject.h>
#include <vtkm/worklet/ScatterIdentity.h>
namespace vtkm {
namespace worklet {
namespace internal {
@ -58,8 +62,17 @@ public:
struct _9 : Arg<9> { };
/// \c ExecutionSignature tag for getting the work index.
///
typedef vtkm::exec::arg::WorkIndex WorkIndex;
/// \c ExecutionSignature tag for getting the thread indices.
///
typedef vtkm::exec::arg::ThreadIndices ThreadIndices;
/// \c ExecutionSignature tag for getting the visit index.
///
typedef vtkm::exec::arg::VisitIndex VisitIndex;
/// \c ControlSignature tag for execution object inputs.
struct ExecObject : vtkm::cont::arg::ControlSignatureTagBase {
typedef vtkm::cont::arg::TypeCheckTagExecObject TypeCheckTag;
@ -71,6 +84,17 @@ public:
/// override this by redefining this type.
typedef _1 InputDomain;
/// All worklets must define their scatter operation. The scatter defines
/// what output each input contributes to. The default scatter is the
/// identity scatter (1-to-1 input to output).
typedef vtkm::worklet::ScatterIdentity ScatterType;
/// In addition to defining the scatter type, the worklet must produce the
/// scatter. The default ScatterIdentity has no state, so just return an
/// instance.
VTKM_CONT_EXPORT
ScatterType GetScatter() const { return ScatterType(); }
/// \brief A type list containing the type vtkm::Id.
///
/// This is a convenience type to use as template arguments to \c

4
vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx
View File

@ -119,7 +119,7 @@ struct Fetch<TestFetchTagInput,
VTKM_EXEC_EXPORT
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic indices,
const TestExecObject &execObject) const {
return execObject.Array[indices.GetIndex()];
return execObject.Array[indices.GetInputIndex()];
}
VTKM_EXEC_EXPORT
@ -149,7 +149,7 @@ struct Fetch<TestFetchTagOutput,
void Store(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject,
ValueType value) const {
execObject.Array[indices.GetIndex()] = value;
execObject.Array[indices.GetOutputIndex()] = value;
}
};

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

@ -26,6 +26,7 @@ set(unit_tests
UnitTestFieldStatistics.cxx
UnitTestIsosurfaceUniformGrid.cxx
UnitTestPointElevation.cxx
UnitTestScatterCounting.cxx
UnitTestSplatKernels.cxx
UnitTestTetrahedralizeExplicitGrid.cxx
UnitTestTetrahedralizeUniformGrid.cxx

35
vtkm/worklet/testing/UnitTestIsosurfaceUniformGrid.cxx
View File

@ -38,12 +38,12 @@ public:
typedef _1 InputDomain;
const vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::FloatDefault xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::Id cellsPerLayer;
VTKM_CONT_EXPORT
TangleField(const vtkm::Id3 dims, const float mins[3], const float maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { };
TangleField(const vtkm::Id3 dims, const vtkm::FloatDefault mins[3], const vtkm::FloatDefault maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { }
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &vertexId, vtkm::Float32 &v) const
@ -52,9 +52,9 @@ public:
const vtkm::Id y = (vertexId / (xdim)) % (ydim);
const vtkm::Id z = vertexId / cellsPerLayer;
const float fx = static_cast<float>(x) / static_cast<float>(xdim-1);
const float fy = static_cast<float>(y) / static_cast<float>(xdim-1);
const float fz = static_cast<float>(z) / static_cast<float>(xdim-1);
const vtkm::FloatDefault fx = static_cast<vtkm::FloatDefault>(x) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fy = static_cast<vtkm::FloatDefault>(y) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fz = static_cast<vtkm::FloatDefault>(z) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::Float32 xx = 3.0f*(xmin+(xmax-xmin)*(fx));
const vtkm::Float32 yy = 3.0f*(ymin+(ymax-ymin)*(fy));
@ -71,15 +71,22 @@ vtkm::cont::DataSet MakeIsosurfaceTestDataSet(vtkm::Id3 dims)
const vtkm::Id3 vdims(dims[0] + 1, dims[1] + 1, dims[2] + 1);
float mins[3] = {-1.0f, -1.0f, -1.0f};
float maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::FloatDefault mins[3] = {-1.0f, -1.0f, -1.0f};
vtkm::FloatDefault maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray;
vtkm::cont::ArrayHandleIndex vertexCountImplicitArray(vdims[0]*vdims[1]*vdims[2]);
vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher(TangleField(vdims, mins, maxs));
tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray);
vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims);
vtkm::Vec<vtkm::FloatDefault,3> origin(0.0f, 0.0f, 0.0f);
vtkm::Vec<vtkm::FloatDefault,3> spacing(
1.0f/static_cast<vtkm::FloatDefault>(dims[0]),
1.0f/static_cast<vtkm::FloatDefault>(dims[2]),
1.0f/static_cast<vtkm::FloatDefault>(dims[1]));
vtkm::cont::ArrayHandleUniformPointCoordinates
coordinates(vdims, origin, spacing);
dataSet.AddCoordinateSystem(
vtkm::cont::CoordinateSystem("coordinates", 1, coordinates));
@ -116,6 +123,16 @@ void TestIsosurfaceUniformGrid()
normalsArray,
scalarsArray);
std::cout << "vertices: ";
vtkm::cont::printSummary_ArrayHandle(verticesArray, std::cout);
std::cout << std::endl;
std::cout << "normals: ";
vtkm::cont::printSummary_ArrayHandle(normalsArray, std::cout);
std::cout << std::endl;
std::cout << "scalars: ";
vtkm::cont::printSummary_ArrayHandle(scalarsArray, std::cout);
std::cout << std::endl;
VTKM_TEST_ASSERT(test_equal(verticesArray.GetNumberOfValues(), 480),
"Wrong result for Isosurface filter");
}

258
vtkm/worklet/testing/UnitTestScatterCounting.cxx Normal file
View File

@ -0,0 +1,258 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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/ScatterCounting.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 {
struct TestScatterArrays
{
vtkm::cont::ArrayHandle<vtkm::IdComponent> CountArray;
vtkm::cont::ArrayHandle<vtkm::Id> OutputToInputMap;
vtkm::cont::ArrayHandle<vtkm::IdComponent> VisitArray;
};
TestScatterArrays MakeScatterArraysShort()
{
const vtkm::Id countArraySize = 18;
const vtkm::IdComponent countArray[countArraySize] = {
1, 2, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0
};
const vtkm::Id outputSize = 6;
const vtkm::Id outputToInputMap[outputSize] = {
0, 1, 1, 4, 6, 14
};
const vtkm::IdComponent visitArray[outputSize] = {
0, 0, 1, 0, 0, 0
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(outputToInputMap, outputSize),
arrays.OutputToInputMap);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(visitArray, outputSize),
arrays.VisitArray);
return arrays;
}
TestScatterArrays MakeScatterArraysLong()
{
const vtkm::Id countArraySize = 6;
const vtkm::IdComponent countArray[countArraySize] = {
0, 1, 2, 3, 4, 5
};
const vtkm::Id outputSize = 15;
const vtkm::Id outputToInputMap[outputSize] = {
1, 2, 2, 3, 3, 3, 4, 4, 4, 4, 5, 5, 5, 5, 5
};
const vtkm::IdComponent visitArray[outputSize] = {
0, 0, 1, 0, 1, 2, 0, 1, 2, 3, 0, 1, 2, 3, 4
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(outputToInputMap, outputSize),
arrays.OutputToInputMap);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(visitArray, outputSize),
arrays.VisitArray);
return arrays;
}
TestScatterArrays MakeScatterArraysZero()
{
const vtkm::Id countArraySize = 6;
const vtkm::IdComponent countArray[countArraySize] = {
0, 0, 0, 0, 0, 0
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
arrays.OutputToInputMap.Allocate(0);
arrays.VisitArray.Allocate(0);
return arrays;
}
struct TestScatterCountingWorklet : public vtkm::worklet::WorkletMapField
{
typedef void ControlSignature(FieldIn<> inputIndices,
FieldOut<> copyIndices,
FieldOut<> recordVisit,
FieldOut<> recordWorkId);
typedef void ExecutionSignature(_1, _2 ,_3, _4, VisitIndex, WorkIndex);
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
ScatterType GetScatter() const { return this->Scatter; }
template<typename CountArrayType>
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const CountArrayType &countArray)
: Scatter(countArray, VTKM_DEFAULT_DEVICE_ADAPTER_TAG()) { }
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const CountArrayType &countArray, Device)
: Scatter(countArray, Device()) { }
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const vtkm::worklet::ScatterCounting &scatter)
: Scatter(scatter) { }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex,
vtkm::Id &indexCopy,
vtkm::IdComponent &writeVisit,
vtkm::Float32 &captureWorkId,
vtkm::IdComponent visitIndex,
vtkm::Id workId) const
{
indexCopy = inputIndex;
writeVisit = visitIndex;
captureWorkId = TestValue(workId, vtkm::Float32());
}
private:
ScatterType Scatter;
};
template<typename T>
void CompareArrays(vtkm::cont::ArrayHandle<T> array1,
vtkm::cont::ArrayHandle<T> array2)
{
typedef typename vtkm::cont::ArrayHandle<T>::PortalConstControl PortalType;
PortalType portal1 = array1.GetPortalConstControl();
PortalType portal2 = array2.GetPortalConstControl();
VTKM_TEST_ASSERT(portal1.GetNumberOfValues() == portal2.GetNumberOfValues(),
"Arrays are not the same length.");
for (vtkm::Id index = 0; index < portal1.GetNumberOfValues(); index++)
{
T value1 = portal1.Get(index);
T value2 = portal2.Get(index);
VTKM_TEST_ASSERT(value1 == value2, "Array values not equal.");
}
}
// This unit test makes sure the ScatterCounting generates the correct map
// and visit arrays.
void TestScatterArrayGeneration(const TestScatterArrays &arrays)
{
std::cout << " Testing array generation" << std::endl;
vtkm::worklet::ScatterCounting scatter(arrays.CountArray,
VTKM_DEFAULT_DEVICE_ADAPTER_TAG());
vtkm::Id inputSize = arrays.CountArray.GetNumberOfValues();
std::cout << " Checking output to input map." << std::endl;
CompareArrays(arrays.OutputToInputMap,
scatter.GetOutputToInputMap(inputSize));
std::cout << " Checking visit array." << std::endl;
CompareArrays(arrays.VisitArray,
scatter.GetVisitArray(inputSize));
}
// This is more of an integration test that makes sure the scatter works with a
// worklet invocation.
void TestScatterWorklet(const TestScatterArrays &arrays)
{
std::cout << " Testing scatter counting in a worklet." << std::endl;
vtkm::worklet::ScatterCounting scatter(arrays.CountArray,
VTKM_DEFAULT_DEVICE_ADAPTER_TAG());
TestScatterCountingWorklet worklet(scatter);
vtkm::worklet::DispatcherMapField<TestScatterCountingWorklet> dispatcher(
worklet);
vtkm::Id inputSize = arrays.CountArray.GetNumberOfValues();
vtkm::cont::ArrayHandleIndex inputIndices(inputSize);
vtkm::cont::ArrayHandle<vtkm::Id> outputToInputMapCopy;
vtkm::cont::ArrayHandle<vtkm::IdComponent> visitCopy;
vtkm::cont::ArrayHandle<vtkm::Float32> captureWorkId;
std::cout << " Invoke worklet" << std::endl;
dispatcher.Invoke(
inputIndices, outputToInputMapCopy, visitCopy, captureWorkId);
std::cout << " Check output to input map." << std::endl;
CompareArrays(outputToInputMapCopy, arrays.OutputToInputMap);
std::cout << " Check visit." << std::endl;
CompareArrays(visitCopy, arrays.VisitArray);
std::cout << " Check work id." << std::endl;
CheckPortal(captureWorkId.GetPortalConstControl());
}
void TestScatterCountingWithArrays(const TestScatterArrays &arrays)
{
TestScatterArrayGeneration(arrays);
TestScatterWorklet(arrays);
}
void TestScatterCounting()
{
std::cout << "Testing arrays with output smaller than input." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysShort());
std::cout << "Testing arrays with output larger than input." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysLong());
std::cout << "Testing arrays with zero output." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysZero());
}
} // anonymous namespace
int UnitTestScatterCounting(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestScatterCounting);
}