Respect VTK-m convention of parameters all or nothing on a line

clang-format BinPack settings have been disabled to make sure that the
VTK-m style guideline is obeyed.
This commit is contained in:
Robert Maynard 2017-05-26 13:53:28 -04:00
parent 3297dbab41
commit 5dd346007b
416 changed files with 30907 additions and 8604 deletions

@ -6,6 +6,8 @@ AlignOperands: false
AlwaysBreakAfterReturnType: None AlwaysBreakAfterReturnType: None
AlwaysBreakAfterDefinitionReturnType: None AlwaysBreakAfterDefinitionReturnType: None
BreakBeforeBraces: Allman BreakBeforeBraces: Allman
BinPackArguments: false
BinPackParameters: false
ColumnLimit: 100 ColumnLimit: 100
Standard: Cpp11 Standard: Cpp11
# This requires clang-format 4.0 (at least). # This requires clang-format 4.0 (at least).

@ -59,8 +59,10 @@ int main(int argc, char* argv[])
vtkm::cont::Timer<DeviceAdapter> total; vtkm::cont::Timer<DeviceAdapter> total;
vtkm::cont::Timer<DeviceAdapter> timer; vtkm::cont::Timer<DeviceAdapter> timer;
vtkm::cont::CellSetExplicit<> outputCellSet = vtkm::cont::CellSetExplicit<> outputCellSet =
clip.Run(input.GetCellSet(0), scalarField.GetData().ResetTypeList(vtkm::TypeListTagScalarAll()), clip.Run(input.GetCellSet(0),
clipValue, DeviceAdapter()); scalarField.GetData().ResetTypeList(vtkm::TypeListTagScalarAll()),
clipValue,
DeviceAdapter());
vtkm::Float64 clipTime = timer.GetElapsedTime(); vtkm::Float64 clipTime = timer.GetElapsedTime();
vtkm::cont::DataSet output; vtkm::cont::DataSet output;

@ -34,11 +34,15 @@
#include <iostream> #include <iostream>
void makeScene(const vtkm::cont::DataSet& inputData, const vtkm::rendering::ColorTable& colorTable, void makeScene(const vtkm::cont::DataSet& inputData,
const std::string& fieldName, vtkm::rendering::Scene& scene) const vtkm::rendering::ColorTable& colorTable,
const std::string& fieldName,
vtkm::rendering::Scene& scene)
{ {
scene.AddActor(vtkm::rendering::Actor(inputData.GetCellSet(), inputData.GetCoordinateSystem(), scene.AddActor(vtkm::rendering::Actor(inputData.GetCellSet(),
inputData.GetField(fieldName), colorTable)); inputData.GetCoordinateSystem(),
inputData.GetField(fieldName),
colorTable));
} }
// This example reads an input vtk file specified on the command-line (or generates a default // This example reads an input vtk file specified on the command-line (or generates a default

@ -31,13 +31,21 @@
struct ExampleFieldWorklet : public vtkm::worklet::WorkletMapField struct ExampleFieldWorklet : public vtkm::worklet::WorkletMapField
{ {
typedef void ControlSignature(FieldIn<>, FieldIn<>, FieldIn<>, FieldOut<>, FieldOut<>, typedef void ControlSignature(FieldIn<>,
FieldIn<>,
FieldIn<>,
FieldOut<>,
FieldOut<>,
FieldOut<>); FieldOut<>);
typedef void ExecutionSignature(_1, _2, _3, _4, _5, _6); typedef void ExecutionSignature(_1, _2, _3, _4, _5, _6);
template <typename T, typename U, typename V> template <typename T, typename U, typename V>
VTKM_EXEC void operator()(const vtkm::Vec<T, 3>& vec, const U& scalar1, const V& scalar2, VTKM_EXEC void operator()(const vtkm::Vec<T, 3>& vec,
vtkm::Vec<T, 3>& out_vec, U& out_scalar1, V& out_scalar2) const const U& scalar1,
const V& scalar2,
vtkm::Vec<T, 3>& out_vec,
U& out_scalar1,
V& out_scalar2) const
{ {
out_vec = vec * scalar1; out_vec = vec * scalar1;
out_scalar1 = static_cast<U>(scalar1 + scalar2); out_scalar1 = static_cast<U>(scalar1 + scalar2);

@ -86,8 +86,8 @@ struct HelloVTKMInterop
{ {
for (int j = 0; j < dim; ++j) for (int j = 0; j < dim; ++j)
{ {
this->InputData.push_back(vtkm::Vec<T, 3>(2.f * static_cast<T>(i / dim) - 1.f, 0.f, this->InputData.push_back(vtkm::Vec<T, 3>(
2.f * static_cast<T>(j / dim) - 1.f)); 2.f * static_cast<T>(i / dim) - 1.f, 0.f, 2.f * static_cast<T>(j / dim) - 1.f));
} }
} }
@ -143,7 +143,8 @@ struct HelloVTKMInterop
typedef void ExecutionSignature(_1, _2, _3); typedef void ExecutionSignature(_1, _2, _3);
VTKM_EXEC VTKM_EXEC
void operator()(const vtkm::Vec<T, 3>& input, vtkm::Vec<T, 3>& output, void operator()(const vtkm::Vec<T, 3>& input,
vtkm::Vec<T, 3>& output,
vtkm::Vec<vtkm::UInt8, 4>& color) const vtkm::Vec<vtkm::UInt8, 4>& color) const
{ {
output[0] = input[0]; output[0] = input[0];

@ -116,8 +116,8 @@ vtkm::cont::DataSet MakeIsosurfaceTestDataSet(vtkm::Id3 dims)
vtkm::Float32 maxs[3] = { 1.0f, 1.0f, 1.0f }; vtkm::Float32 maxs[3] = { 1.0f, 1.0f, 1.0f };
vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray; vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray;
vtkm::cont::ArrayHandleCounting<vtkm::Id> vertexCountImplicitArray(0, 1, vdims[0] * vdims[1] * vtkm::cont::ArrayHandleCounting<vtkm::Id> vertexCountImplicitArray(
vdims[2]); 0, 1, vdims[0] * vdims[1] * vdims[2]);
vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher( vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher(
TangleField(vdims, mins, maxs)); TangleField(vdims, mins, maxs));
tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray); tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray);

@ -46,7 +46,8 @@ struct GenerateSurfaceWorklet : public vtkm::worklet::WorkletMapField
typedef void ExecutionSignature(_1, _2, _3); typedef void ExecutionSignature(_1, _2, _3);
template <typename T> template <typename T>
VTKM_EXEC void operator()(const vtkm::Vec<T, 3>& input, vtkm::Vec<T, 3>& output, VTKM_EXEC void operator()(const vtkm::Vec<T, 3>& input,
vtkm::Vec<T, 3>& output,
vtkm::Vec<vtkm::UInt8, 4>& color) const vtkm::Vec<vtkm::UInt8, 4>& color) const
{ {
output[0] = input[0]; output[0] = input[0];
@ -93,8 +94,8 @@ std::vector<vtkm::Vec<T, 3>> make_testData(int size)
{ {
for (int j = 0; j < size; ++j) for (int j = 0; j < size; ++j)
{ {
data.push_back(vtkm::Vec<T, 3>(2.f * static_cast<T>(i / size) - 1.f, 0.f, data.push_back(vtkm::Vec<T, 3>(
2.f * static_cast<T>(j / size) - 1.f)); 2.f * static_cast<T>(i / size) - 1.f, 0.f, 2.f * static_cast<T>(j / size) - 1.f));
} }
} }
return data; return data;
@ -102,9 +103,9 @@ std::vector<vtkm::Vec<T, 3>> make_testData(int size)
//This is the list of devices to compile in support for. The order of the //This is the list of devices to compile in support for. The order of the
//devices determines the runtime preference. //devices determines the runtime preference.
struct DevicesToTry struct DevicesToTry : vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda,
: vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda, vtkm::cont::DeviceAdapterTagTBB, vtkm::cont::DeviceAdapterTagTBB,
vtkm::cont::DeviceAdapterTagSerial> vtkm::cont::DeviceAdapterTagSerial>
{ {
}; };

@ -138,7 +138,8 @@ int main(int argc, char* argv[])
vtkm::rendering::MapperGL mapper; vtkm::rendering::MapperGL mapper;
vtkm::rendering::Scene scene; vtkm::rendering::Scene scene;
scene.AddActor(vtkm::rendering::Actor(ds.GetCellSet(), ds.GetCoordinateSystem(), scene.AddActor(vtkm::rendering::Actor(ds.GetCellSet(),
ds.GetCoordinateSystem(),
ds.GetField("pointvar"), ds.GetField("pointvar"),
vtkm::rendering::ColorTable("thermal"))); vtkm::rendering::ColorTable("thermal")));

@ -70,9 +70,10 @@ vtkm::cont::DataSet MakeTetrahedralizeTestDataSet(vtkm::Id3 dim)
// Place uniform grid on a set physical space so OpenGL drawing is easier // Place uniform grid on a set physical space so OpenGL drawing is easier
const vtkm::Id3 vdims(dim[0] + 1, dim[1] + 1, dim[2] + 1); const vtkm::Id3 vdims(dim[0] + 1, dim[1] + 1, dim[2] + 1);
const vtkm::Vec<vtkm::Float32, 3> origin = vtkm::make_Vec(0.0f, 0.0f, 0.0f); const vtkm::Vec<vtkm::Float32, 3> origin = vtkm::make_Vec(0.0f, 0.0f, 0.0f);
const vtkm::Vec<vtkm::Float32, 3> spacing = vtkm::make_Vec( const vtkm::Vec<vtkm::Float32, 3> spacing =
1.0f / static_cast<vtkm::Float32>(dim[0]), 1.0f / static_cast<vtkm::Float32>(dim[1]), vtkm::make_Vec(1.0f / static_cast<vtkm::Float32>(dim[0]),
1.0f / static_cast<vtkm::Float32>(dim[2])); 1.0f / static_cast<vtkm::Float32>(dim[1]),
1.0f / static_cast<vtkm::Float32>(dim[2]));
// Generate coordinate system // Generate coordinate system
vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims, origin, spacing); vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims, origin, spacing);

@ -54,8 +54,12 @@ struct Bounds
} }
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
VTKM_EXEC_CONT Bounds(const T1& minX, const T2& maxX, const T3& minY, const T4& maxY, VTKM_EXEC_CONT Bounds(const T1& minX,
const T5& minZ, const T6& maxZ) const T2& maxX,
const T3& minY,
const T4& maxY,
const T5& minZ,
const T6& maxZ)
: X(vtkm::Range(minX, maxX)) : X(vtkm::Range(minX, maxX))
, Y(vtkm::Range(minY, maxY)) , Y(vtkm::Range(minY, maxY))
, Z(vtkm::Range(minZ, maxZ)) , Z(vtkm::Range(minZ, maxZ))

@ -1738,8 +1738,8 @@ static inline VTKM_EXEC_CONT T Max(const T& x, const T& y, vtkm::TypeTraitsVecto
T result; T result;
for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++) for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
{ {
Traits::SetComponent(result, index, Traits::SetComponent(
vtkm::Max(Traits::GetComponent(x, index), Traits::GetComponent(y, index))); result, index, vtkm::Max(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
} }
return result; return result;
} }
@ -1757,8 +1757,8 @@ static inline VTKM_EXEC_CONT T Min(const T& x, const T& y, vtkm::TypeTraitsVecto
T result; T result;
for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++) for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
{ {
Traits::SetComponent(result, index, Traits::SetComponent(
vtkm::Min(Traits::GetComponent(x, index), Traits::GetComponent(y, index))); result, index, vtkm::Min(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
} }
return result; return result;
} }

@ -115,7 +115,8 @@ private:
/// ///
template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol> template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol>
VTKM_EXEC_CONT const vtkm::Vec<T, NumCol>& MatrixGetRow( VTKM_EXEC_CONT const vtkm::Vec<T, NumCol>& MatrixGetRow(
const vtkm::Matrix<T, NumRow, NumCol>& matrix, vtkm::IdComponent rowIndex) const vtkm::Matrix<T, NumRow, NumCol>& matrix,
vtkm::IdComponent rowIndex)
{ {
return matrix[rowIndex]; return matrix[rowIndex];
} }
@ -139,7 +140,8 @@ VTKM_EXEC_CONT vtkm::Vec<T, NumRow> MatrixGetColumn(const vtkm::Matrix<T, NumRow
/// ///
template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol> template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol>
VTKM_EXEC_CONT void MatrixSetRow(vtkm::Matrix<T, NumRow, NumCol>& matrix, VTKM_EXEC_CONT void MatrixSetRow(vtkm::Matrix<T, NumRow, NumCol>& matrix,
vtkm::IdComponent rowIndex, const vtkm::Vec<T, NumCol>& rowValues) vtkm::IdComponent rowIndex,
const vtkm::Vec<T, NumCol>& rowValues)
{ {
matrix[rowIndex] = rowValues; matrix[rowIndex] = rowValues;
} }
@ -159,7 +161,9 @@ VTKM_EXEC_CONT void MatrixSetColumn(vtkm::Matrix<T, NumRow, NumCol>& matrix,
/// Standard matrix multiplication. /// Standard matrix multiplication.
/// ///
template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol, template <typename T,
vtkm::IdComponent NumRow,
vtkm::IdComponent NumCol,
vtkm::IdComponent NumInternal> vtkm::IdComponent NumInternal>
VTKM_EXEC_CONT vtkm::Matrix<T, NumRow, NumCol> MatrixMultiply( VTKM_EXEC_CONT vtkm::Matrix<T, NumRow, NumCol> MatrixMultiply(
const vtkm::Matrix<T, NumRow, NumInternal>& leftFactor, const vtkm::Matrix<T, NumRow, NumInternal>& leftFactor,
@ -185,7 +189,8 @@ VTKM_EXEC_CONT vtkm::Matrix<T, NumRow, NumCol> MatrixMultiply(
/// ///
template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol> template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol>
VTKM_EXEC_CONT vtkm::Vec<T, NumRow> MatrixMultiply( VTKM_EXEC_CONT vtkm::Vec<T, NumRow> MatrixMultiply(
const vtkm::Matrix<T, NumRow, NumCol>& leftFactor, const vtkm::Vec<T, NumCol>& rightFactor) const vtkm::Matrix<T, NumRow, NumCol>& leftFactor,
const vtkm::Vec<T, NumCol>& rightFactor)
{ {
vtkm::Vec<T, NumRow> product; vtkm::Vec<T, NumRow> product;
for (vtkm::IdComponent rowIndex = 0; rowIndex < NumRow; rowIndex++) for (vtkm::IdComponent rowIndex = 0; rowIndex < NumRow; rowIndex++)
@ -199,7 +204,8 @@ VTKM_EXEC_CONT vtkm::Vec<T, NumRow> MatrixMultiply(
/// ///
template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol> template <typename T, vtkm::IdComponent NumRow, vtkm::IdComponent NumCol>
VTKM_EXEC_CONT vtkm::Vec<T, NumCol> MatrixMultiply( VTKM_EXEC_CONT vtkm::Vec<T, NumCol> MatrixMultiply(
const vtkm::Vec<T, NumRow>& leftFactor, const vtkm::Matrix<T, NumRow, NumCol>& rightFactor) const vtkm::Vec<T, NumRow>& leftFactor,
const vtkm::Matrix<T, NumRow, NumCol>& rightFactor)
{ {
vtkm::Vec<T, NumCol> product; vtkm::Vec<T, NumCol> product;
for (vtkm::IdComponent colIndex = 0; colIndex < NumCol; colIndex++) for (vtkm::IdComponent colIndex = 0; colIndex < NumCol; colIndex++)
@ -260,7 +266,8 @@ namespace detail
template <typename T, vtkm::IdComponent Size> template <typename T, vtkm::IdComponent Size>
VTKM_EXEC_CONT void MatrixLUPFactorFindPivot(vtkm::Matrix<T, Size, Size>& A, VTKM_EXEC_CONT void MatrixLUPFactorFindPivot(vtkm::Matrix<T, Size, Size>& A,
vtkm::Vec<vtkm::IdComponent, Size>& permutation, vtkm::Vec<vtkm::IdComponent, Size>& permutation,
vtkm::IdComponent topCornerIndex, T& inversionParity, vtkm::IdComponent topCornerIndex,
T& inversionParity,
bool& valid) bool& valid)
{ {
vtkm::IdComponent maxRowIndex = topCornerIndex; vtkm::IdComponent maxRowIndex = topCornerIndex;
@ -351,7 +358,8 @@ VTKM_EXEC_CONT void MatrixLUPFactorFindUpperTriangleElements(vtkm::Matrix<T, Siz
template <typename T, vtkm::IdComponent Size> template <typename T, vtkm::IdComponent Size>
VTKM_EXEC_CONT void MatrixLUPFactor(vtkm::Matrix<T, Size, Size>& A, VTKM_EXEC_CONT void MatrixLUPFactor(vtkm::Matrix<T, Size, Size>& A,
vtkm::Vec<vtkm::IdComponent, Size>& permutation, vtkm::Vec<vtkm::IdComponent, Size>& permutation,
T& inversionParity, bool& valid) T& inversionParity,
bool& valid)
{ {
// Initialize permutation. // Initialize permutation.
for (vtkm::IdComponent index = 0; index < Size; index++) for (vtkm::IdComponent index = 0; index < Size; index++)
@ -374,7 +382,8 @@ VTKM_EXEC_CONT void MatrixLUPFactor(vtkm::Matrix<T, Size, Size>& A,
/// ///
template <typename T, vtkm::IdComponent Size> template <typename T, vtkm::IdComponent Size>
VTKM_EXEC_CONT vtkm::Vec<T, Size> MatrixLUPSolve( VTKM_EXEC_CONT vtkm::Vec<T, Size> MatrixLUPSolve(
const vtkm::Matrix<T, Size, Size>& LU, const vtkm::Vec<vtkm::IdComponent, Size>& permutation, const vtkm::Matrix<T, Size, Size>& LU,
const vtkm::Vec<vtkm::IdComponent, Size>& permutation,
const vtkm::Vec<T, Size>& b) const vtkm::Vec<T, Size>& b)
{ {
// The LUP-factorization gives us PA = LU or equivalently A = inv(P)LU. // The LUP-factorization gives us PA = LU or equivalently A = inv(P)LU.
@ -418,7 +427,8 @@ VTKM_EXEC_CONT vtkm::Vec<T, Size> MatrixLUPSolve(
/// ///
template <typename T, vtkm::IdComponent Size> template <typename T, vtkm::IdComponent Size>
VTKM_EXEC_CONT vtkm::Vec<T, Size> SolveLinearSystem(const vtkm::Matrix<T, Size, Size>& A, VTKM_EXEC_CONT vtkm::Vec<T, Size> SolveLinearSystem(const vtkm::Matrix<T, Size, Size>& A,
const vtkm::Vec<T, Size>& b, bool& valid) const vtkm::Vec<T, Size>& b,
bool& valid)
{ {
// First, we will make an LUP-factorization to help us. // First, we will make an LUP-factorization to help us.
vtkm::Matrix<T, Size, Size> LU = A; vtkm::Matrix<T, Size, Size> LU = A;

@ -37,13 +37,17 @@ namespace vtkm
/// returned. /// returned.
/// ///
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ScalarType, vtkm::IdComponent Size, typename JacobianFunctor, template <typename ScalarType,
vtkm::IdComponent Size,
typename JacobianFunctor,
typename FunctionFunctor> typename FunctionFunctor>
VTKM_EXEC_CONT vtkm::Vec<ScalarType, Size> NewtonsMethod( VTKM_EXEC_CONT vtkm::Vec<ScalarType, Size> NewtonsMethod(
JacobianFunctor jacobianEvaluator, FunctionFunctor functionEvaluator, JacobianFunctor jacobianEvaluator,
FunctionFunctor functionEvaluator,
vtkm::Vec<ScalarType, Size> desiredFunctionOutput, vtkm::Vec<ScalarType, Size> desiredFunctionOutput,
vtkm::Vec<ScalarType, Size> initialGuess = vtkm::Vec<ScalarType, Size>(ScalarType(0)), vtkm::Vec<ScalarType, Size> initialGuess = vtkm::Vec<ScalarType, Size>(ScalarType(0)),
ScalarType convergeDifference = ScalarType(1e-3), vtkm::IdComponent maxIterations = 10) ScalarType convergeDifference = ScalarType(1e-3),
vtkm::IdComponent maxIterations = 10)
{ {
typedef vtkm::Vec<ScalarType, Size> VectorType; typedef vtkm::Vec<ScalarType, Size> VectorType;
typedef vtkm::Matrix<ScalarType, Size, Size> MatrixType; typedef vtkm::Matrix<ScalarType, Size, Size> MatrixType;

@ -93,7 +93,8 @@ VTKM_EXEC_CONT vtkm::Vec<T, 3> Transform3DVector(const vtkm::Matrix<T, 4, 4>& ma
/// transformation matrix for those scales. /// transformation matrix for those scales.
/// ///
template <typename T> template <typename T>
VTKM_EXEC_CONT vtkm::Matrix<T, 4, 4> Transform3DScale(const T& scaleX, const T& scaleY, VTKM_EXEC_CONT vtkm::Matrix<T, 4, 4> Transform3DScale(const T& scaleX,
const T& scaleY,
const T& scaleZ) const T& scaleZ)
{ {
vtkm::Matrix<T, 4, 4> scaleMatrix(T(0)); vtkm::Matrix<T, 4, 4> scaleMatrix(T(0));

@ -27,132 +27,148 @@
#include <vtkm/ListTag.h> #include <vtkm/ListTag.h>
#include <vtkm/Types.h> #include <vtkm/Types.h>
namespace vtkm { namespace vtkm
{
/// A list containing the type vtkm::Id. /// A list containing the type vtkm::Id.
/// ///
struct TypeListTagId : vtkm::ListTagBase<vtkm::Id> { }; struct TypeListTagId : vtkm::ListTagBase<vtkm::Id>
{
};
/// A list containing the type vtkm::Id2. /// A list containing the type vtkm::Id2.
/// ///
struct TypeListTagId2 : vtkm::ListTagBase<vtkm::Id2> { }; struct TypeListTagId2 : vtkm::ListTagBase<vtkm::Id2>
{
};
/// A list containing the type vtkm::Id3. /// A list containing the type vtkm::Id3.
/// ///
struct TypeListTagId3 : vtkm::ListTagBase<vtkm::Id3> { }; struct TypeListTagId3 : vtkm::ListTagBase<vtkm::Id3>
{
};
/// A list containing the type vtkm::IdComponent /// A list containing the type vtkm::IdComponent
/// ///
struct TypeListTagIdComponent : vtkm::ListTagBase<vtkm::IdComponent> { }; struct TypeListTagIdComponent : vtkm::ListTagBase<vtkm::IdComponent>
{
};
/// A list containing types used to index arrays. Contains vtkm::Id, vtkm::Id2, /// A list containing types used to index arrays. Contains vtkm::Id, vtkm::Id2,
/// and vtkm::Id3. /// and vtkm::Id3.
/// ///
struct TypeListTagIndex struct TypeListTagIndex : vtkm::ListTagBase<vtkm::Id, vtkm::Id2, vtkm::Id3>
: vtkm::ListTagBase<vtkm::Id,vtkm::Id2,vtkm::Id3> { }; {
};
/// A list containing types used for scalar fields. Specifically, contains /// A list containing types used for scalar fields. Specifically, contains
/// floating point numbers of different widths (i.e. vtkm::Float32 and /// floating point numbers of different widths (i.e. vtkm::Float32 and
/// vtkm::Float64). /// vtkm::Float64).
struct TypeListTagFieldScalar : vtkm::ListTagBase<vtkm::Float32,vtkm::Float64> { }; struct TypeListTagFieldScalar : vtkm::ListTagBase<vtkm::Float32, vtkm::Float64>
{
};
/// A list containing types for values for fields with two dimensional /// A list containing types for values for fields with two dimensional
/// vectors. /// vectors.
/// ///
struct TypeListTagFieldVec2 struct TypeListTagFieldVec2
: vtkm::ListTagBase<vtkm::Vec<vtkm::Float32,2>, : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 2>, vtkm::Vec<vtkm::Float64, 2>>
vtkm::Vec<vtkm::Float64,2> > { }; {
};
/// A list containing types for values for fields with three dimensional /// A list containing types for values for fields with three dimensional
/// vectors. /// vectors.
/// ///
struct TypeListTagFieldVec3 struct TypeListTagFieldVec3
: vtkm::ListTagBase<vtkm::Vec<vtkm::Float32,3>, : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 3>, vtkm::Vec<vtkm::Float64, 3>>
vtkm::Vec<vtkm::Float64,3> > { }; {
};
/// A list containing types for values for fields with four dimensional /// A list containing types for values for fields with four dimensional
/// vectors. /// vectors.
/// ///
struct TypeListTagFieldVec4 struct TypeListTagFieldVec4
: vtkm::ListTagBase<vtkm::Vec<vtkm::Float32,4>, : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 4>, vtkm::Vec<vtkm::Float64, 4>>
vtkm::Vec<vtkm::Float64,4> > { }; {
};
/// A list containing common types for values in fields. Specifically contains /// A list containing common types for values in fields. Specifically contains
/// floating point scalars and vectors of size 2, 3, and 4 with floating point /// floating point scalars and vectors of size 2, 3, and 4 with floating point
/// components. /// components.
/// ///
struct TypeListTagField struct TypeListTagField : vtkm::ListTagBase<vtkm::Float32,
: vtkm::ListTagBase<vtkm::Float32, vtkm::Float64,
vtkm::Float64, vtkm::Vec<vtkm::Float32, 2>,
vtkm::Vec<vtkm::Float32,2>, vtkm::Vec<vtkm::Float64, 2>,
vtkm::Vec<vtkm::Float64,2>, vtkm::Vec<vtkm::Float32, 3>,
vtkm::Vec<vtkm::Float32,3>, vtkm::Vec<vtkm::Float64, 3>,
vtkm::Vec<vtkm::Float64,3>, vtkm::Vec<vtkm::Float32, 4>,
vtkm::Vec<vtkm::Float32,4>, vtkm::Vec<vtkm::Float64, 4>>
vtkm::Vec<vtkm::Float64,4> > {
{ }; };
/// A list of all scalars defined in vtkm/Types.h. A scalar is a type that /// A list of all scalars defined in vtkm/Types.h. A scalar is a type that
/// holds a single number. /// holds a single number.
/// ///
struct TypeListTagScalarAll struct TypeListTagScalarAll : vtkm::ListTagBase<vtkm::Int8,
: vtkm::ListTagBase<vtkm::Int8, vtkm::UInt8,
vtkm::UInt8, vtkm::Int16,
vtkm::Int16, vtkm::UInt16,
vtkm::UInt16, vtkm::Int32,
vtkm::Int32, vtkm::UInt32,
vtkm::UInt32, vtkm::Int64,
vtkm::Int64, vtkm::UInt64,
vtkm::UInt64, vtkm::Float32,
vtkm::Float32, vtkm::Float64>
vtkm::Float64> {
{ }; };
/// A list of the most commonly use Vec classes. Specifically, these are /// A list of the most commonly use Vec classes. Specifically, these are
/// vectors of size 2, 3, or 4 containing either unsigned bytes, signed /// vectors of size 2, 3, or 4 containing either unsigned bytes, signed
/// integers of 32 or 64 bits, or floating point values of 32 or 64 bits. /// integers of 32 or 64 bits, or floating point values of 32 or 64 bits.
/// ///
struct TypeListTagVecCommon struct TypeListTagVecCommon : vtkm::ListTagBase<vtkm::Vec<vtkm::UInt8, 2>,
: vtkm::ListTagBase<vtkm::Vec<vtkm::UInt8,2>, vtkm::Vec<vtkm::Int32, 2>,
vtkm::Vec<vtkm::Int32,2>, vtkm::Vec<vtkm::Int64, 2>,
vtkm::Vec<vtkm::Int64,2>, vtkm::Vec<vtkm::Float32, 2>,
vtkm::Vec<vtkm::Float32,2>, vtkm::Vec<vtkm::Float64, 2>,
vtkm::Vec<vtkm::Float64,2>, vtkm::Vec<vtkm::UInt8, 3>,
vtkm::Vec<vtkm::UInt8,3>, vtkm::Vec<vtkm::Int32, 3>,
vtkm::Vec<vtkm::Int32,3>, vtkm::Vec<vtkm::Int64, 3>,
vtkm::Vec<vtkm::Int64,3>, vtkm::Vec<vtkm::Float32, 3>,
vtkm::Vec<vtkm::Float32,3>, vtkm::Vec<vtkm::Float64, 3>,
vtkm::Vec<vtkm::Float64,3>, vtkm::Vec<vtkm::UInt8, 4>,
vtkm::Vec<vtkm::UInt8,4>, vtkm::Vec<vtkm::Int32, 4>,
vtkm::Vec<vtkm::Int32,4>, vtkm::Vec<vtkm::Int64, 4>,
vtkm::Vec<vtkm::Int64,4>, vtkm::Vec<vtkm::Float32, 4>,
vtkm::Vec<vtkm::Float32,4>, vtkm::Vec<vtkm::Float64, 4>>
vtkm::Vec<vtkm::Float64,4> > {
{ }; };
namespace internal { namespace internal
{
/// A list of uncommon Vec classes with length up to 4. This is not much /// A list of uncommon Vec classes with length up to 4. This is not much
/// use in general, but is used when joined with \c TypeListTagVecCommon /// use in general, but is used when joined with \c TypeListTagVecCommon
/// to get a list of all vectors up to size 4. /// to get a list of all vectors up to size 4.
/// ///
struct TypeListTagVecUncommon struct TypeListTagVecUncommon : vtkm::ListTagBase<vtkm::Vec<vtkm::Int8, 2>,
: vtkm::ListTagBase<vtkm::Vec<vtkm::Int8,2>, vtkm::Vec<vtkm::Int16, 2>,
vtkm::Vec<vtkm::Int16,2>, vtkm::Vec<vtkm::UInt16, 2>,
vtkm::Vec<vtkm::UInt16,2>, vtkm::Vec<vtkm::UInt32, 2>,
vtkm::Vec<vtkm::UInt32,2>, vtkm::Vec<vtkm::UInt64, 2>,
vtkm::Vec<vtkm::UInt64,2>, vtkm::Vec<vtkm::Int8, 3>,
vtkm::Vec<vtkm::Int8,3>, vtkm::Vec<vtkm::Int16, 3>,
vtkm::Vec<vtkm::Int16,3>, vtkm::Vec<vtkm::UInt16, 3>,
vtkm::Vec<vtkm::UInt16,3>, vtkm::Vec<vtkm::UInt32, 3>,
vtkm::Vec<vtkm::UInt32,3>, vtkm::Vec<vtkm::UInt64, 3>,
vtkm::Vec<vtkm::UInt64,3>, vtkm::Vec<vtkm::Int8, 4>,
vtkm::Vec<vtkm::Int8,4>, vtkm::Vec<vtkm::Int16, 4>,
vtkm::Vec<vtkm::Int16,4>, vtkm::Vec<vtkm::UInt16, 4>,
vtkm::Vec<vtkm::UInt16,4>, vtkm::Vec<vtkm::UInt32, 4>,
vtkm::Vec<vtkm::UInt32,4>, vtkm::Vec<vtkm::UInt64, 4>>
vtkm::Vec<vtkm::UInt64,4> > {
{ }; };
} // namespace internal } // namespace internal
@ -160,35 +176,35 @@ struct TypeListTagVecUncommon
/// lengths between 2 and 4. /// lengths between 2 and 4.
/// ///
struct TypeListTagVecAll struct TypeListTagVecAll
: vtkm::ListTagJoin< : vtkm::ListTagJoin<vtkm::TypeListTagVecCommon, vtkm::internal::TypeListTagVecUncommon>
vtkm::TypeListTagVecCommon, vtkm::internal::TypeListTagVecUncommon> {
{ }; };
/// A list of all basic types listed in vtkm/Types.h. Does not include all /// A list of all basic types listed in vtkm/Types.h. Does not include all
/// possible VTK-m types like arbitrarily typed and sized Vecs (only up to /// possible VTK-m types like arbitrarily typed and sized Vecs (only up to
/// length 4) or math types like matrices. /// length 4) or math types like matrices.
/// ///
struct TypeListTagAll struct TypeListTagAll : vtkm::ListTagJoin<vtkm::TypeListTagScalarAll, vtkm::TypeListTagVecAll>
: vtkm::ListTagJoin<vtkm::TypeListTagScalarAll, vtkm::TypeListTagVecAll> {
{ }; };
/// A list of the most commonly used types across multiple domains. Includes /// A list of the most commonly used types across multiple domains. Includes
/// integers, floating points, and 3 dimensional vectors of floating points. /// integers, floating points, and 3 dimensional vectors of floating points.
/// ///
struct TypeListTagCommon struct TypeListTagCommon : vtkm::ListTagBase<vtkm::Int32,
: vtkm::ListTagBase<vtkm::Int32, vtkm::Int64,
vtkm::Int64, vtkm::Float32,
vtkm::Float32, vtkm::Float64,
vtkm::Float64, vtkm::Vec<vtkm::Float32, 3>,
vtkm::Vec<vtkm::Float32,3>, vtkm::Vec<vtkm::Float64, 3>>
vtkm::Vec<vtkm::Float64,3> > {
{ }; };
// Special implementation of ListContains for TypeListTagAll to always be // Special implementation of ListContains for TypeListTagAll to always be
// true. Although TypeListTagAll is necessarily finite, the point is to // true. Although TypeListTagAll is necessarily finite, the point is to
// be all inclusive. Besides, this should speed up the compilation when // be all inclusive. Besides, this should speed up the compilation when
// checking a list that should contain everything. // checking a list that should contain everything.
template<typename Type> template <typename Type>
struct ListContains<vtkm::TypeListTagAll, Type> struct ListContains<vtkm::TypeListTagAll, Type>
{ {
static const bool value = true; static const bool value = true;

@ -109,13 +109,16 @@ struct VecTraits
/// Returns the value in a given component of the vector. /// Returns the value in a given component of the vector.
/// ///
VTKM_EXEC_CONT static const ComponentType& GetComponent( VTKM_EXEC_CONT static const ComponentType& GetComponent(
const typename std::remove_const<VecType>::type& vector, vtkm::IdComponent component); const typename std::remove_const<VecType>::type& vector,
vtkm::IdComponent component);
VTKM_EXEC_CONT static ComponentType& GetComponent( VTKM_EXEC_CONT static ComponentType& GetComponent(
typename std::remove_const<VecType>::type& vector, vtkm::IdComponent component); typename std::remove_const<VecType>::type& vector,
vtkm::IdComponent component);
/// Changes the value in a given component of the vector. /// Changes the value in a given component of the vector.
/// ///
VTKM_EXEC_CONT static void SetComponent(VecType& vector, vtkm::IdComponent component, VTKM_EXEC_CONT static void SetComponent(VecType& vector,
vtkm::IdComponent component,
ComponentType value); ComponentType value);
/// Copies the components in the given vector into a given Vec object. /// Copies the components in the given vector into a given Vec object.
@ -183,7 +186,8 @@ struct VecTraits<vtkm::Vec<T, Size>>
/// Changes the value in a given component of the vector. /// Changes the value in a given component of the vector.
/// ///
VTKM_EXEC_CONT static void SetComponent(VecType& vector, vtkm::IdComponent component, VTKM_EXEC_CONT static void SetComponent(VecType& vector,
vtkm::IdComponent component,
ComponentType value) ComponentType value)
{ {
vector[component] = value; vector[component] = value;
@ -346,7 +350,8 @@ struct VecTraitsBasic
VTKM_EXEC_CONT VTKM_EXEC_CONT
static ComponentType& GetComponent(ScalarType& vector, vtkm::IdComponent) { return vector; } static ComponentType& GetComponent(ScalarType& vector, vtkm::IdComponent) { return vector; }
VTKM_EXEC_CONT static void SetComponent(ScalarType& vector, vtkm::IdComponent, VTKM_EXEC_CONT static void SetComponent(ScalarType& vector,
vtkm::IdComponent,
ComponentType value) ComponentType value)
{ {
vector = value; vector = value;

@ -142,7 +142,8 @@ struct VecTraits<vtkm::VecVariable<T, MaxSize>>
} }
VTKM_EXEC_CONT VTKM_EXEC_CONT
static void SetComponent(VecType& vector, vtkm::IdComponent componentIndex, static void SetComponent(VecType& vector,
vtkm::IdComponent componentIndex,
const ComponentType& value) const ComponentType& value)
{ {
vector[componentIndex] = value; vector[componentIndex] = value;

@ -41,7 +41,8 @@ namespace vtkm
/// extrapolates. If w=0 => v0 is returned if w=1 => v1 is returned. /// extrapolates. If w=0 => v0 is returned if w=1 => v1 is returned.
/// ///
template <typename ValueType, typename WeightType> template <typename ValueType, typename WeightType>
VTKM_EXEC_CONT ValueType Lerp(const ValueType& value0, const ValueType& value1, VTKM_EXEC_CONT ValueType Lerp(const ValueType& value0,
const ValueType& value1,
const WeightType& weight) const WeightType& weight)
{ {
return static_cast<ValueType>((WeightType(1) - weight) * value0 + weight * value1); return static_cast<ValueType>((WeightType(1) - weight) * value0 + weight * value1);
@ -80,14 +81,16 @@ namespace detail
{ {
template <typename T> template <typename T>
VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type MagnitudeTemplate( VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type MagnitudeTemplate(
T x, vtkm::TypeTraitsScalarTag) T x,
vtkm::TypeTraitsScalarTag)
{ {
return static_cast<typename detail::FloatingPointReturnType<T>::Type>(vtkm::Abs(x)); return static_cast<typename detail::FloatingPointReturnType<T>::Type>(vtkm::Abs(x));
} }
template <typename T> template <typename T>
VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type MagnitudeTemplate( VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type MagnitudeTemplate(
const T& x, vtkm::TypeTraitsVectorTag) const T& x,
vtkm::TypeTraitsVectorTag)
{ {
return vtkm::Sqrt(vtkm::MagnitudeSquared(x)); return vtkm::Sqrt(vtkm::MagnitudeSquared(x));
} }
@ -113,14 +116,16 @@ namespace detail
{ {
template <typename T> template <typename T>
VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type RMagnitudeTemplate( VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type RMagnitudeTemplate(
T x, vtkm::TypeTraitsScalarTag) T x,
vtkm::TypeTraitsScalarTag)
{ {
return T(1) / vtkm::Abs(x); return T(1) / vtkm::Abs(x);
} }
template <typename T> template <typename T>
VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type RMagnitudeTemplate( VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type RMagnitudeTemplate(
const T& x, vtkm::TypeTraitsVectorTag) const T& x,
vtkm::TypeTraitsVectorTag)
{ {
return vtkm::RSqrt(vtkm::MagnitudeSquared(x)); return vtkm::RSqrt(vtkm::MagnitudeSquared(x));
} }
@ -179,7 +184,8 @@ VTKM_EXEC_CONT void Normalize(T& x)
/// ///
template <typename T> template <typename T>
VTKM_EXEC_CONT vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Cross( VTKM_EXEC_CONT vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Cross(
const vtkm::Vec<T, 3>& x, const vtkm::Vec<T, 3>& y) const vtkm::Vec<T, 3>& x,
const vtkm::Vec<T, 3>& y)
{ {
return vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>( return vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
x[1] * y[2] - x[2] * y[1], x[2] * y[0] - x[0] * y[2], x[0] * y[1] - x[1] * y[0]); x[1] * y[2] - x[2] * y[1], x[2] * y[0] - x[0] * y[2], x[0] * y[1] - x[1] * y[0]);
@ -193,8 +199,8 @@ VTKM_EXEC_CONT vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> C
/// to that triangle/plane. /// to that triangle/plane.
/// ///
template <typename T> template <typename T>
VTKM_EXEC_CONT vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> TriangleNormal( VTKM_EXEC_CONT vtkm::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>
const vtkm::Vec<T, 3>& a, const vtkm::Vec<T, 3>& b, const vtkm::Vec<T, 3>& c) TriangleNormal(const vtkm::Vec<T, 3>& a, const vtkm::Vec<T, 3>& b, const vtkm::Vec<T, 3>& c)
{ {
return vtkm::Cross(b - a, c - a); return vtkm::Cross(b - a, c - a);
} }

@ -604,10 +604,16 @@ private:
VTKM_MAKE_BENCHMARK(UpperBounds30, BenchUpperBounds, 30); VTKM_MAKE_BENCHMARK(UpperBounds30, BenchUpperBounds, 30);
public: public:
struct ValueTypes struct ValueTypes : vtkm::ListTagBase<vtkm::UInt8,
: vtkm::ListTagBase<vtkm::UInt8, vtkm::UInt32, vtkm::Int32, vtkm::Int64, vtkm::UInt32,
vtkm::Vec<vtkm::Int32, 2>, vtkm::Vec<vtkm::UInt8, 4>, vtkm::Float32, vtkm::Int32,
vtkm::Float64, vtkm::Vec<vtkm::Float64, 3>, vtkm::Vec<vtkm::Float32, 4>> vtkm::Int64,
vtkm::Vec<vtkm::Int32, 2>,
vtkm::Vec<vtkm::UInt8, 4>,
vtkm::Float32,
vtkm::Float64,
vtkm::Vec<vtkm::Float64, 3>,
vtkm::Vec<vtkm::Float32, 4>>
{ {
}; };

@ -64,7 +64,10 @@ class BlackScholes : public vtkm::worklet::WorkletMapField
T Volatility; T Volatility;
public: public:
typedef void ControlSignature(FieldIn<Scalar>, FieldIn<Scalar>, FieldIn<Scalar>, FieldOut<Scalar>, typedef void ControlSignature(FieldIn<Scalar>,
FieldIn<Scalar>,
FieldIn<Scalar>,
FieldOut<Scalar>,
FieldOut<Scalar>); FieldOut<Scalar>);
typedef void ExecutionSignature(_1, _2, _3, _4, _5); typedef void ExecutionSignature(_1, _2, _3, _4, _5);
@ -99,8 +102,8 @@ public:
} }
template <typename U, typename V, typename W> template <typename U, typename V, typename W>
VTKM_EXEC void operator()(const U& sp, const V& os, const W& oy, T& callResult, VTKM_EXEC void operator()(const U& sp, const V& os, const W& oy, T& callResult, T& putResult)
T& putResult) const const
{ {
const T stockPrice = static_cast<T>(sp); const T stockPrice = static_cast<T>(sp);
const T optionStrike = static_cast<T>(os); const T optionStrike = static_cast<T>(os);
@ -226,13 +229,15 @@ class InterpolateField : public vtkm::worklet::WorkletMapField
{ {
public: public:
typedef void ControlSignature(FieldIn<Id2Type> interpolation_ids, typedef void ControlSignature(FieldIn<Id2Type> interpolation_ids,
FieldIn<Scalar> interpolation_weights, WholeArrayIn<> inputField, FieldIn<Scalar> interpolation_weights,
WholeArrayIn<> inputField,
FieldOut<> output); FieldOut<> output);
typedef void ExecutionSignature(_1, _2, _3, _4); typedef void ExecutionSignature(_1, _2, _3, _4);
typedef _1 InputDomain; typedef _1 InputDomain;
template <typename WeightType, typename T, typename S, typename D> template <typename WeightType, typename T, typename S, typename D>
VTKM_EXEC void operator()(const vtkm::Id2& low_high, const WeightType& weight, VTKM_EXEC void operator()(const vtkm::Id2& low_high,
const WeightType& weight,
const vtkm::exec::ExecutionWholeArrayConst<T, S, D>& inPortal, const vtkm::exec::ExecutionWholeArrayConst<T, S, D>& inPortal,
T& result) const T& result) const
{ {
@ -241,8 +246,10 @@ public:
} }
template <typename WeightType, typename T, typename S, typename D, typename U> template <typename WeightType, typename T, typename S, typename D, typename U>
VTKM_EXEC void operator()(const vtkm::Id2&, const WeightType&, VTKM_EXEC void operator()(const vtkm::Id2&,
const vtkm::exec::ExecutionWholeArrayConst<T, S, D>&, U&) const const WeightType&,
const vtkm::exec::ExecutionWholeArrayConst<T, S, D>&,
U&) const
{ {
//the inPortal and result need to be the same type so this version only //the inPortal and result need to be the same type so this version only
//exists to generate code when using dynamic arrays //exists to generate code when using dynamic arrays
@ -300,9 +307,10 @@ struct ValueTypes : vtkm::ListTagBase<vtkm::Float32, vtkm::Float64>
{ {
}; };
struct InterpValueTypes struct InterpValueTypes : vtkm::ListTagBase<vtkm::Float32,
: vtkm::ListTagBase<vtkm::Float32, vtkm::Float64, vtkm::Vec<vtkm::Float32, 3>, vtkm::Float64,
vtkm::Vec<vtkm::Float64, 3>> vtkm::Vec<vtkm::Float32, 3>,
vtkm::Vec<vtkm::Float64, 3>>
{ {
}; };
using StorageListTag = ::vtkm::cont::StorageListTagBasic; using StorageListTag = ::vtkm::cont::StorageListTagBasic;
@ -371,8 +379,8 @@ private:
BlackScholes<Value> worklet(RISKFREE, VOLATILITY); BlackScholes<Value> worklet(RISKFREE, VOLATILITY);
vtkm::worklet::DispatcherMapField<BlackScholes<Value>> dispatcher(worklet); vtkm::worklet::DispatcherMapField<BlackScholes<Value>> dispatcher(worklet);
dispatcher.Invoke(this->StockPrice, this->OptionStrike, this->OptionYears, callResultHandle, dispatcher.Invoke(
putResultHandle); this->StockPrice, this->OptionStrike, this->OptionYears, callResultHandle, putResultHandle);
return timer.GetElapsedTime(); return timer.GetElapsedTime();
} }

@ -54,14 +54,16 @@ enum BenchmarkName
class AveragePointToCell : public vtkm::worklet::WorkletMapPointToCell class AveragePointToCell : public vtkm::worklet::WorkletMapPointToCell
{ {
public: public:
typedef void ControlSignature(FieldInPoint<> inPoints, CellSetIn cellset, typedef void ControlSignature(FieldInPoint<> inPoints,
CellSetIn cellset,
FieldOutCell<> outCells); FieldOutCell<> outCells);
typedef void ExecutionSignature(_1, PointCount, _3); typedef void ExecutionSignature(_1, PointCount, _3);
typedef _2 InputDomain; typedef _2 InputDomain;
template <typename PointValueVecType, typename OutType> template <typename PointValueVecType, typename OutType>
VTKM_EXEC void operator()(const PointValueVecType& pointValues, VTKM_EXEC void operator()(const PointValueVecType& pointValues,
const vtkm::IdComponent& numPoints, OutType& average) const const vtkm::IdComponent& numPoints,
OutType& average) const
{ {
OutType sum = static_cast<OutType>(pointValues[0]); OutType sum = static_cast<OutType>(pointValues[0]);
for (vtkm::IdComponent pointIndex = 1; pointIndex < numPoints; ++pointIndex) for (vtkm::IdComponent pointIndex = 1; pointIndex < numPoints; ++pointIndex)
@ -81,7 +83,8 @@ public:
typedef _2 InputDomain; typedef _2 InputDomain;
template <typename CellVecType, typename OutType> template <typename CellVecType, typename OutType>
VTKM_EXEC void operator()(const CellVecType& cellValues, OutType& avgVal, VTKM_EXEC void operator()(const CellVecType& cellValues,
OutType& avgVal,
const vtkm::IdComponent& numCellIDs) const const vtkm::IdComponent& numCellIDs) const
{ {
//simple functor that returns the average cell Value. //simple functor that returns the average cell Value.
@ -99,7 +102,8 @@ template <typename T>
class Classification : public vtkm::worklet::WorkletMapPointToCell class Classification : public vtkm::worklet::WorkletMapPointToCell
{ {
public: public:
typedef void ControlSignature(FieldInPoint<> inNodes, CellSetIn cellset, typedef void ControlSignature(FieldInPoint<> inNodes,
CellSetIn cellset,
FieldOutCell<IdComponentType> outCaseId); FieldOutCell<IdComponentType> outCaseId);
typedef void ExecutionSignature(_1, _3); typedef void ExecutionSignature(_1, _3);
typedef _2 InputDomain; typedef _2 InputDomain;

@ -420,7 +420,8 @@ public:
/// ///
template <typename DeviceAdapterTag> template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForOutput( VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForOutput(
vtkm::Id numberOfValues, DeviceAdapterTag); vtkm::Id numberOfValues,
DeviceAdapterTag);
/// Prepares this array to be used in an in-place operation (both as input /// Prepares this array to be used in an in-place operation (both as input
/// and output) in the execution environment. If necessary, copies data to /// and output) in the execution environment. If necessary, copies data to
@ -508,15 +509,17 @@ namespace detail
{ {
template <typename T> template <typename T>
VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value( VTKM_NEVER_EXPORT VTKM_CONT inline void
const T& value, std::ostream& out, vtkm::VecTraitsTagSingleComponent) printSummary_ArrayHandle_Value(const T& value, std::ostream& out, vtkm::VecTraitsTagSingleComponent)
{ {
out << value; out << value;
} }
template <typename T> template <typename T>
VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value( VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value(
const T& value, std::ostream& out, vtkm::VecTraitsTagMultipleComponents) const T& value,
std::ostream& out,
vtkm::VecTraitsTagMultipleComponents)
{ {
using Traits = vtkm::VecTraits<T>; using Traits = vtkm::VecTraits<T>;
using ComponentType = typename Traits::ComponentType; using ComponentType = typename Traits::ComponentType;
@ -534,7 +537,8 @@ VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value(
VTKM_NEVER_EXPORT VTKM_NEVER_EXPORT
VTKM_CONT VTKM_CONT
inline void printSummary_ArrayHandle_Value(UInt8 value, std::ostream& out, inline void printSummary_ArrayHandle_Value(UInt8 value,
std::ostream& out,
vtkm::VecTraitsTagSingleComponent) vtkm::VecTraitsTagSingleComponent)
{ {
out << static_cast<int>(value); out << static_cast<int>(value);
@ -542,7 +546,8 @@ inline void printSummary_ArrayHandle_Value(UInt8 value, std::ostream& out,
VTKM_NEVER_EXPORT VTKM_NEVER_EXPORT
VTKM_CONT VTKM_CONT
inline void printSummary_ArrayHandle_Value(Int8 value, std::ostream& out, inline void printSummary_ArrayHandle_Value(Int8 value,
std::ostream& out,
vtkm::VecTraitsTagSingleComponent) vtkm::VecTraitsTagSingleComponent)
{ {
out << static_cast<int>(value); out << static_cast<int>(value);
@ -550,14 +555,16 @@ inline void printSummary_ArrayHandle_Value(Int8 value, std::ostream& out,
template <typename T1, typename T2> template <typename T1, typename T2>
VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value( VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value(
const vtkm::Pair<T1, T2>& value, std::ostream& out, vtkm::VecTraitsTagSingleComponent) const vtkm::Pair<T1, T2>& value,
std::ostream& out,
vtkm::VecTraitsTagSingleComponent)
{ {
out << "{"; out << "{";
printSummary_ArrayHandle_Value(value.first, out, printSummary_ArrayHandle_Value(
typename vtkm::VecTraits<T1>::HasMultipleComponents()); value.first, out, typename vtkm::VecTraits<T1>::HasMultipleComponents());
out << ","; out << ",";
printSummary_ArrayHandle_Value(value.second, out, printSummary_ArrayHandle_Value(
typename vtkm::VecTraits<T2>::HasMultipleComponents()); value.second, out, typename vtkm::VecTraits<T2>::HasMultipleComponents());
out << "}"; out << "}";
} }
@ -565,7 +572,8 @@ VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle_Value(
template <typename T, typename StorageT> template <typename T, typename StorageT>
VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle( VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle(
const vtkm::cont::ArrayHandle<T, StorageT>& array, std::ostream& out) const vtkm::cont::ArrayHandle<T, StorageT>& array,
std::ostream& out)
{ {
using ArrayType = vtkm::cont::ArrayHandle<T, StorageT>; using ArrayType = vtkm::cont::ArrayHandle<T, StorageT>;
using PortalType = typename ArrayType::PortalConstControl; using PortalType = typename ArrayType::PortalConstControl;

@ -186,7 +186,8 @@ void ArrayHandle<T, S>::CopyInto(IteratorType dest, DeviceAdapterTag) const
{ {
PortalConstControl portal = this->GetPortalConstControl(); PortalConstControl portal = this->GetPortalConstControl();
std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal), std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal),
vtkm::cont::ArrayPortalToIteratorEnd(portal), dest); vtkm::cont::ArrayPortalToIteratorEnd(portal),
dest);
} }
} }

@ -34,7 +34,9 @@ namespace internal
/// \brief An array portal that acts as a 3D cartesian product of 3 arrays. /// \brief An array portal that acts as a 3D cartesian product of 3 arrays.
/// ///
template <typename ValueType_, typename PortalTypeFirst_, typename PortalTypeSecond_, template <typename ValueType_,
typename PortalTypeFirst_,
typename PortalTypeSecond_,
typename PortalTypeThird_> typename PortalTypeThird_>
class VTKM_ALWAYS_EXPORT ArrayPortalCartesianProduct class VTKM_ALWAYS_EXPORT ArrayPortalCartesianProduct
{ {
@ -101,8 +103,8 @@ public:
vtkm::Id i2 = idx12 / dim1; vtkm::Id i2 = idx12 / dim1;
vtkm::Id i3 = index / dim12; vtkm::Id i3 = index / dim12;
return vtkm::make_Vec(this->PortalFirst.Get(i1), this->PortalSecond.Get(i2), return vtkm::make_Vec(
this->PortalThird.Get(i3)); this->PortalFirst.Get(i1), this->PortalSecond.Get(i2), this->PortalThird.Get(i3));
} }
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
@ -190,12 +192,16 @@ public:
typedef T ValueType; typedef T ValueType;
typedef vtkm::exec::internal::ArrayPortalCartesianProduct< typedef vtkm::exec::internal::ArrayPortalCartesianProduct<
ValueType, typename FirstHandleType::PortalControl, typename SecondHandleType::PortalControl, ValueType,
typename FirstHandleType::PortalControl,
typename SecondHandleType::PortalControl,
typename ThirdHandleType::PortalControl> typename ThirdHandleType::PortalControl>
PortalType; PortalType;
typedef vtkm::exec::internal::ArrayPortalCartesianProduct< typedef vtkm::exec::internal::ArrayPortalCartesianProduct<
ValueType, typename FirstHandleType::PortalConstControl, ValueType,
typename SecondHandleType::PortalConstControl, typename ThirdHandleType::PortalConstControl> typename FirstHandleType::PortalConstControl,
typename SecondHandleType::PortalConstControl,
typename ThirdHandleType::PortalConstControl>
PortalConstType; PortalConstType;
VTKM_CONT VTKM_CONT
@ -207,7 +213,8 @@ public:
} }
VTKM_CONT VTKM_CONT
Storage(const FirstHandleType& array1, const SecondHandleType& array2, Storage(const FirstHandleType& array1,
const SecondHandleType& array2,
const ThirdHandleType& array3) const ThirdHandleType& array3)
: FirstArray(array1) : FirstArray(array1)
, SecondArray(array2) , SecondArray(array2)
@ -218,7 +225,8 @@ public:
VTKM_CONT VTKM_CONT
PortalType GetPortal() PortalType GetPortal()
{ {
return PortalType(this->FirstArray.GetPortalControl(), this->SecondArray.GetPortalControl(), return PortalType(this->FirstArray.GetPortalControl(),
this->SecondArray.GetPortalControl(),
this->ThirdArray.GetPortalControl()); this->ThirdArray.GetPortalControl());
} }
@ -271,10 +279,14 @@ private:
ThirdHandleType ThirdArray; ThirdHandleType ThirdArray;
}; };
template <typename T, typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType, template <typename T,
typename FirstHandleType,
typename SecondHandleType,
typename ThirdHandleType,
typename Device> typename Device>
class ArrayTransfer< class ArrayTransfer<T,
T, StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>, Device> StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>,
Device>
{ {
typedef StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType> StorageTag; typedef StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType> StorageTag;
typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType; typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType;
@ -286,13 +298,15 @@ public:
typedef typename StorageType::PortalConstType PortalConstControl; typedef typename StorageType::PortalConstType PortalConstControl;
typedef vtkm::exec::internal::ArrayPortalCartesianProduct< typedef vtkm::exec::internal::ArrayPortalCartesianProduct<
ValueType, typename FirstHandleType::template ExecutionTypes<Device>::Portal, ValueType,
typename FirstHandleType::template ExecutionTypes<Device>::Portal,
typename SecondHandleType::template ExecutionTypes<Device>::Portal, typename SecondHandleType::template ExecutionTypes<Device>::Portal,
typename ThirdHandleType::template ExecutionTypes<Device>::Portal> typename ThirdHandleType::template ExecutionTypes<Device>::Portal>
PortalExecution; PortalExecution;
typedef vtkm::exec::internal::ArrayPortalCartesianProduct< typedef vtkm::exec::internal::ArrayPortalCartesianProduct<
ValueType, typename FirstHandleType::template ExecutionTypes<Device>::PortalConst, ValueType,
typename FirstHandleType::template ExecutionTypes<Device>::PortalConst,
typename SecondHandleType::template ExecutionTypes<Device>::PortalConst, typename SecondHandleType::template ExecutionTypes<Device>::PortalConst,
typename ThirdHandleType::template ExecutionTypes<Device>::PortalConst> typename ThirdHandleType::template ExecutionTypes<Device>::PortalConst>
PortalConstExecution; PortalConstExecution;
@ -371,7 +385,8 @@ private:
/// ///
template <typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType> template <typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType>
class ArrayHandleCartesianProduct class ArrayHandleCartesianProduct
: public internal::ArrayHandleCartesianProductTraits<FirstHandleType, SecondHandleType, : public internal::ArrayHandleCartesianProductTraits<FirstHandleType,
SecondHandleType,
ThirdHandleType>::Superclass ThirdHandleType>::Superclass
{ {
// If the following line gives a compile error, then the FirstHandleType // If the following line gives a compile error, then the FirstHandleType
@ -384,7 +399,8 @@ public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleCartesianProduct, ArrayHandleCartesianProduct,
(ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>), (ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>),
(typename internal::ArrayHandleCartesianProductTraits<FirstHandleType, SecondHandleType, (typename internal::ArrayHandleCartesianProductTraits<FirstHandleType,
SecondHandleType,
ThirdHandleType>::Superclass)); ThirdHandleType>::Superclass));
private: private:
@ -406,7 +422,8 @@ public:
template <typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType> template <typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType>
VTKM_CONT VTKM_CONT
vtkm::cont::ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType> vtkm::cont::ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>
make_ArrayHandleCartesianProduct(const FirstHandleType& first, const SecondHandleType& second, make_ArrayHandleCartesianProduct(const FirstHandleType& first,
const SecondHandleType& second,
const ThirdHandleType& third) const ThirdHandleType& third)
{ {
return ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>( return ArrayHandleCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>(

@ -49,14 +49,17 @@ struct VTKM_ALWAYS_EXPORT Cast
/// ///
template <typename T, typename ArrayHandleType> template <typename T, typename ArrayHandleType>
class ArrayHandleCast class ArrayHandleCast
: public vtkm::cont::ArrayHandleTransform<T, ArrayHandleType, : public vtkm::cont::ArrayHandleTransform<T,
ArrayHandleType,
internal::Cast<typename ArrayHandleType::ValueType, T>, internal::Cast<typename ArrayHandleType::ValueType, T>,
internal::Cast<T, typename ArrayHandleType::ValueType>> internal::Cast<T, typename ArrayHandleType::ValueType>>
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleCast, (ArrayHandleCast<T, ArrayHandleType>), ArrayHandleCast,
(vtkm::cont::ArrayHandleTransform<T, ArrayHandleType, (ArrayHandleCast<T, ArrayHandleType>),
(vtkm::cont::ArrayHandleTransform<T,
ArrayHandleType,
internal::Cast<typename ArrayHandleType::ValueType, T>, internal::Cast<typename ArrayHandleType::ValueType, T>,
internal::Cast<T, typename ArrayHandleType::ValueType>>)); internal::Cast<T, typename ArrayHandleType::ValueType>>));

@ -123,7 +123,8 @@ struct CompositeVectorArrayToPortalCont
template <typename ArrayHandleType, vtkm::IdComponent Index> template <typename ArrayHandleType, vtkm::IdComponent Index>
VTKM_CONT typename ReturnType<ArrayHandleType, Index>::type operator()( VTKM_CONT typename ReturnType<ArrayHandleType, Index>::type operator()(
const ArrayHandleType& array, vtkm::internal::IndexTag<Index>) const const ArrayHandleType& array,
vtkm::internal::IndexTag<Index>) const
{ {
return array.GetPortalConstControl(); return array.GetPortalConstControl();
} }
@ -140,7 +141,8 @@ struct CompositeVectorArrayToPortalExec
template <typename ArrayHandleType, vtkm::IdComponent Index> template <typename ArrayHandleType, vtkm::IdComponent Index>
VTKM_CONT typename ReturnType<ArrayHandleType, Index>::type operator()( VTKM_CONT typename ReturnType<ArrayHandleType, Index>::type operator()(
const ArrayHandleType& array, vtkm::internal::IndexTag<Index>) const const ArrayHandleType& array,
vtkm::internal::IndexTag<Index>) const
{ {
return array.PrepareForInput(DeviceAdapterTag()); return array.PrepareForInput(DeviceAdapterTag());
} }
@ -464,7 +466,8 @@ class ArrayHandleCompositeVector
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleCompositeVector, (ArrayHandleCompositeVector<Signature>), ArrayHandleCompositeVector,
(ArrayHandleCompositeVector<Signature>),
(typename internal::ArrayHandleCompositeVectorTraits<Signature>::Superclass)); (typename internal::ArrayHandleCompositeVectorTraits<Signature>::Superclass));
VTKM_CONT VTKM_CONT
@ -506,13 +509,18 @@ public:
ComponentMapType(sourceComponent1, sourceComponent2, sourceComponent3))) ComponentMapType(sourceComponent1, sourceComponent2, sourceComponent3)))
{ {
} }
template <typename ArrayHandleType1, typename ArrayHandleType2, typename ArrayHandleType3, template <typename ArrayHandleType1,
typename ArrayHandleType2,
typename ArrayHandleType3,
typename ArrayHandleType4> typename ArrayHandleType4>
VTKM_CONT ArrayHandleCompositeVector( VTKM_CONT ArrayHandleCompositeVector(const ArrayHandleType1& array1,
const ArrayHandleType1& array1, vtkm::IdComponent sourceComponent1, vtkm::IdComponent sourceComponent1,
const ArrayHandleType2& array2, vtkm::IdComponent sourceComponent2, const ArrayHandleType2& array2,
const ArrayHandleType3& array3, vtkm::IdComponent sourceComponent3, vtkm::IdComponent sourceComponent2,
const ArrayHandleType4& array4, vtkm::IdComponent sourceComponent4) const ArrayHandleType3& array3,
vtkm::IdComponent sourceComponent3,
const ArrayHandleType4& array4,
vtkm::IdComponent sourceComponent4)
: Superclass(StorageType( : Superclass(StorageType(
vtkm::internal::make_FunctionInterface<ValueType>(array1, array2, array3, array4), vtkm::internal::make_FunctionInterface<ValueType>(array1, array2, array3, array4),
ComponentMapType(sourceComponent1, sourceComponent2, sourceComponent3, sourceComponent4))) ComponentMapType(sourceComponent1, sourceComponent2, sourceComponent3, sourceComponent4)))
@ -533,8 +541,10 @@ public:
/// OutArrayType outArray = vtkm::cont::make_ArrayHandleCompositeVector(a1,a2); /// OutArrayType outArray = vtkm::cont::make_ArrayHandleCompositeVector(a1,a2);
/// \endcode /// \endcode
/// ///
template <typename ArrayHandleType1, typename ArrayHandleType2 = void, template <typename ArrayHandleType1,
typename ArrayHandleType3 = void, typename ArrayHandleType4 = void> typename ArrayHandleType2 = void,
typename ArrayHandleType3 = void,
typename ArrayHandleType4 = void>
struct ArrayHandleCompositeVectorType struct ArrayHandleCompositeVectorType
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType1); VTKM_IS_ARRAY_HANDLE(ArrayHandleType1);
@ -545,8 +555,10 @@ struct ArrayHandleCompositeVectorType
private: private:
typedef typedef
typename vtkm::VecTraits<typename ArrayHandleType1::ValueType>::ComponentType ComponentType; typename vtkm::VecTraits<typename ArrayHandleType1::ValueType>::ComponentType ComponentType;
typedef vtkm::Vec<ComponentType, 4> Signature(ArrayHandleType1, ArrayHandleType2, typedef vtkm::Vec<ComponentType, 4> Signature(ArrayHandleType1,
ArrayHandleType3, ArrayHandleType4); ArrayHandleType2,
ArrayHandleType3,
ArrayHandleType4);
public: public:
typedef vtkm::cont::ArrayHandleCompositeVector<Signature> type; typedef vtkm::cont::ArrayHandleCompositeVector<Signature> type;
@ -562,7 +574,8 @@ struct ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, ArrayH
private: private:
typedef typedef
typename vtkm::VecTraits<typename ArrayHandleType1::ValueType>::ComponentType ComponentType; typename vtkm::VecTraits<typename ArrayHandleType1::ValueType>::ComponentType ComponentType;
typedef vtkm::Vec<ComponentType, 3> Signature(ArrayHandleType1, ArrayHandleType2, typedef vtkm::Vec<ComponentType, 3> Signature(ArrayHandleType1,
ArrayHandleType2,
ArrayHandleType3); ArrayHandleType3);
public: public:
@ -622,8 +635,10 @@ make_ArrayHandleCompositeVector(const ArrayHandleType1& array1, vtkm::IdComponen
} }
template <typename ArrayHandleType1, typename ArrayHandleType2> template <typename ArrayHandleType1, typename ArrayHandleType2>
VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2>::type VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2>::type
make_ArrayHandleCompositeVector(const ArrayHandleType1& array1, vtkm::IdComponent sourceComponent1, make_ArrayHandleCompositeVector(const ArrayHandleType1& array1,
const ArrayHandleType2& array2, vtkm::IdComponent sourceComponent2) vtkm::IdComponent sourceComponent1,
const ArrayHandleType2& array2,
vtkm::IdComponent sourceComponent2)
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType1); VTKM_IS_ARRAY_HANDLE(ArrayHandleType1);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType2); VTKM_IS_ARRAY_HANDLE(ArrayHandleType2);
@ -631,39 +646,55 @@ make_ArrayHandleCompositeVector(const ArrayHandleType1& array1, vtkm::IdComponen
array1, sourceComponent1, array2, sourceComponent2); array1, sourceComponent1, array2, sourceComponent2);
} }
template <typename ArrayHandleType1, typename ArrayHandleType2, typename ArrayHandleType3> template <typename ArrayHandleType1, typename ArrayHandleType2, typename ArrayHandleType3>
VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1,
ArrayHandleType2,
ArrayHandleType3>::type ArrayHandleType3>::type
make_ArrayHandleCompositeVector(const ArrayHandleType1& array1, vtkm::IdComponent sourceComponent1, make_ArrayHandleCompositeVector(const ArrayHandleType1& array1,
const ArrayHandleType2& array2, vtkm::IdComponent sourceComponent2, vtkm::IdComponent sourceComponent1,
const ArrayHandleType3& array3, vtkm::IdComponent sourceComponent3) const ArrayHandleType2& array2,
vtkm::IdComponent sourceComponent2,
const ArrayHandleType3& array3,
vtkm::IdComponent sourceComponent3)
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType1); VTKM_IS_ARRAY_HANDLE(ArrayHandleType1);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType2); VTKM_IS_ARRAY_HANDLE(ArrayHandleType2);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType3); VTKM_IS_ARRAY_HANDLE(ArrayHandleType3);
return typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, return
ArrayHandleType3>::type(array1, sourceComponent1, typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, ArrayHandleType3>::
array2, sourceComponent2, type(array1, sourceComponent1, array2, sourceComponent2, array3, sourceComponent3);
array3, sourceComponent3);
} }
template <typename ArrayHandleType1, typename ArrayHandleType2, typename ArrayHandleType3, template <typename ArrayHandleType1,
typename ArrayHandleType2,
typename ArrayHandleType3,
typename ArrayHandleType4> typename ArrayHandleType4>
VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, VTKM_CONT typename ArrayHandleCompositeVectorType<ArrayHandleType1,
ArrayHandleType3, ArrayHandleType4>::type ArrayHandleType2,
make_ArrayHandleCompositeVector(const ArrayHandleType1& array1, vtkm::IdComponent sourceComponent1, ArrayHandleType3,
const ArrayHandleType2& array2, vtkm::IdComponent sourceComponent2, ArrayHandleType4>::type
const ArrayHandleType3& array3, vtkm::IdComponent sourceComponent3, make_ArrayHandleCompositeVector(const ArrayHandleType1& array1,
const ArrayHandleType4& array4, vtkm::IdComponent sourceComponent4) vtkm::IdComponent sourceComponent1,
const ArrayHandleType2& array2,
vtkm::IdComponent sourceComponent2,
const ArrayHandleType3& array3,
vtkm::IdComponent sourceComponent3,
const ArrayHandleType4& array4,
vtkm::IdComponent sourceComponent4)
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType1); VTKM_IS_ARRAY_HANDLE(ArrayHandleType1);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType2); VTKM_IS_ARRAY_HANDLE(ArrayHandleType2);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType3); VTKM_IS_ARRAY_HANDLE(ArrayHandleType3);
VTKM_IS_ARRAY_HANDLE(ArrayHandleType4); VTKM_IS_ARRAY_HANDLE(ArrayHandleType4);
return return typename ArrayHandleCompositeVectorType<ArrayHandleType1,
typename ArrayHandleCompositeVectorType<ArrayHandleType1, ArrayHandleType2, ArrayHandleType3, ArrayHandleType2,
ArrayHandleType4>::type(array1, sourceComponent1, ArrayHandleType3,
array2, sourceComponent2, ArrayHandleType4>::type(array1,
array3, sourceComponent3, sourceComponent1,
array4, sourceComponent4); array2,
sourceComponent2,
array3,
sourceComponent3,
array4,
sourceComponent4);
} }
} }
} // namespace vtkm::cont } // namespace vtkm::cont

@ -200,7 +200,8 @@ private:
template <typename ArrayHandleType1, typename ArrayHandleType2, typename Device> template <typename ArrayHandleType1, typename ArrayHandleType2, typename Device>
class ArrayTransfer<typename ArrayHandleType1::ValueType, class ArrayTransfer<typename ArrayHandleType1::ValueType,
StorageTagConcatenate<ArrayHandleType1, ArrayHandleType2>, Device> StorageTagConcatenate<ArrayHandleType1, ArrayHandleType2>,
Device>
{ {
public: public:
typedef typename ArrayHandleType1::ValueType ValueType; typedef typename ArrayHandleType1::ValueType ValueType;
@ -299,7 +300,8 @@ class ArrayHandleConcatenate
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleConcatenate, (ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2>), ArrayHandleConcatenate,
(ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2>),
(vtkm::cont::ArrayHandle<typename ArrayHandleType1::ValueType, (vtkm::cont::ArrayHandle<typename ArrayHandleType1::ValueType,
StorageTagConcatenate<ArrayHandleType1, ArrayHandleType2>>)); StorageTagConcatenate<ArrayHandleType1, ArrayHandleType2>>));
@ -316,7 +318,8 @@ public:
template <typename ArrayHandleType1, typename ArrayHandleType2> template <typename ArrayHandleType1, typename ArrayHandleType2>
VTKM_CONT ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2> make_ArrayHandleConcatenate( VTKM_CONT ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2> make_ArrayHandleConcatenate(
const ArrayHandleType1& array1, const ArrayHandleType2& array2) const ArrayHandleType1& array1,
const ArrayHandleType2& array2)
{ {
return ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2>(array1, array2); return ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2>(array1, array2);
} }

@ -62,7 +62,8 @@ template <typename T>
class ArrayHandleConstant : public vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T>> class ArrayHandleConstant : public vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T>>
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleConstant, (ArrayHandleConstant<T>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleConstant,
(ArrayHandleConstant<T>),
(vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T>>)); (vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T>>));
VTKM_CONT VTKM_CONT

@ -113,15 +113,17 @@ struct ArrayHandleCountingTraits
/// contains a increment value, that is increment for each step between zero /// contains a increment value, that is increment for each step between zero
/// and the passed in length /// and the passed in length
template <typename CountingValueType> template <typename CountingValueType>
class ArrayHandleCounting class ArrayHandleCounting : public vtkm::cont::ArrayHandle<
: public vtkm::cont::ArrayHandle< CountingValueType,
CountingValueType, typename internal::ArrayHandleCountingTraits<CountingValueType>::Tag> typename internal::ArrayHandleCountingTraits<CountingValueType>::Tag>
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleCounting, (ArrayHandleCounting<CountingValueType>), ArrayHandleCounting,
(ArrayHandleCounting<CountingValueType>),
(vtkm::cont::ArrayHandle< (vtkm::cont::ArrayHandle<
CountingValueType, typename internal::ArrayHandleCountingTraits<CountingValueType>::Tag>)); CountingValueType,
typename internal::ArrayHandleCountingTraits<CountingValueType>::Tag>));
VTKM_CONT VTKM_CONT
ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::Id length) ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::Id length)
@ -133,8 +135,8 @@ public:
/// A convenience function for creating an ArrayHandleCounting. It takes the /// A convenience function for creating an ArrayHandleCounting. It takes the
/// value to start counting from and and the number of times to increment. /// value to start counting from and and the number of times to increment.
template <typename CountingValueType> template <typename CountingValueType>
VTKM_CONT vtkm::cont::ArrayHandleCounting<CountingValueType> make_ArrayHandleCounting( VTKM_CONT vtkm::cont::ArrayHandleCounting<CountingValueType>
CountingValueType start, CountingValueType step, vtkm::Id length) make_ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::Id length)
{ {
return vtkm::cont::ArrayHandleCounting<CountingValueType>(start, step, length); return vtkm::cont::ArrayHandleCounting<CountingValueType>(start, step, length);
} }

@ -221,7 +221,8 @@ template <typename ValueType_>
class ArrayHandleDiscard : public internal::ArrayHandleDiscardTraits<ValueType_>::Superclass class ArrayHandleDiscard : public internal::ArrayHandleDiscardTraits<ValueType_>::Superclass
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleDiscard, (ArrayHandleDiscard<ValueType_>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleDiscard,
(ArrayHandleDiscard<ValueType_>),
(typename internal::ArrayHandleDiscardTraits<ValueType_>::Superclass)); (typename internal::ArrayHandleDiscardTraits<ValueType_>::Superclass));
}; };

@ -136,9 +136,9 @@ public:
typedef vtkm::exec::internal::ArrayPortalGroupVec<typename SourceArrayHandleType::PortalControl, typedef vtkm::exec::internal::ArrayPortalGroupVec<typename SourceArrayHandleType::PortalControl,
NUM_COMPONENTS> NUM_COMPONENTS>
PortalType; PortalType;
typedef vtkm::exec::internal::ArrayPortalGroupVec< typedef vtkm::exec::internal::
typename SourceArrayHandleType::PortalConstControl, NUM_COMPONENTS> ArrayPortalGroupVec<typename SourceArrayHandleType::PortalConstControl, NUM_COMPONENTS>
PortalConstType; PortalConstType;
VTKM_CONT VTKM_CONT
Storage() Storage()
@ -235,10 +235,12 @@ public:
typedef typename StorageType::PortalConstType PortalConstControl; typedef typename StorageType::PortalConstType PortalConstControl;
typedef vtkm::exec::internal::ArrayPortalGroupVec< typedef vtkm::exec::internal::ArrayPortalGroupVec<
typename SourceArrayHandleType::template ExecutionTypes<Device>::Portal, NUM_COMPONENTS> typename SourceArrayHandleType::template ExecutionTypes<Device>::Portal,
NUM_COMPONENTS>
PortalExecution; PortalExecution;
typedef vtkm::exec::internal::ArrayPortalGroupVec< typedef vtkm::exec::internal::ArrayPortalGroupVec<
typename SourceArrayHandleType::template ExecutionTypes<Device>::PortalConst, NUM_COMPONENTS> typename SourceArrayHandleType::template ExecutionTypes<Device>::PortalConst,
NUM_COMPONENTS>
PortalConstExecution; PortalConstExecution;
VTKM_CONT VTKM_CONT
@ -333,7 +335,8 @@ class ArrayHandleGroupVec
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleGroupVec, (ArrayHandleGroupVec<SourceArrayHandleType, NUM_COMPONENTS>), ArrayHandleGroupVec,
(ArrayHandleGroupVec<SourceArrayHandleType, NUM_COMPONENTS>),
(vtkm::cont::ArrayHandle< (vtkm::cont::ArrayHandle<
vtkm::Vec<typename SourceArrayHandleType::ValueType, NUM_COMPONENTS>, vtkm::Vec<typename SourceArrayHandleType::ValueType, NUM_COMPONENTS>,
vtkm::cont::internal::StorageTagGroupVec<SourceArrayHandleType, NUM_COMPONENTS>>)); vtkm::cont::internal::StorageTagGroupVec<SourceArrayHandleType, NUM_COMPONENTS>>));

@ -96,7 +96,8 @@ public:
} }
return ValueType(this->SourcePortal, return ValueType(this->SourcePortal,
static_cast<vtkm::IdComponent>(nextOffsetIndex - offsetIndex), offsetIndex); static_cast<vtkm::IdComponent>(nextOffsetIndex - offsetIndex),
offsetIndex);
} }
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
@ -135,7 +136,8 @@ namespace arg
// can also ignore the Store because the data is already set in the array at // can also ignore the Store because the data is already set in the array at
// that point. // that point.
template <typename ThreadIndicesType, typename SourcePortalType, typename OffsetsPortalType> template <typename ThreadIndicesType, typename SourcePortalType, typename OffsetsPortalType>
struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut, vtkm::exec::arg::AspectTagDefault, struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
ThreadIndicesType, ThreadIndicesType,
vtkm::exec::internal::ArrayPortalGroupVecVariable<SourcePortalType, OffsetsPortalType>> vtkm::exec::internal::ArrayPortalGroupVecVariable<SourcePortalType, OffsetsPortalType>>
{ {
@ -187,7 +189,8 @@ public:
using ValueType = vtkm::VecFromPortal<typename SourceArrayHandleType::PortalControl>; using ValueType = vtkm::VecFromPortal<typename SourceArrayHandleType::PortalControl>;
using PortalType = vtkm::exec::internal::ArrayPortalGroupVecVariable< using PortalType = vtkm::exec::internal::ArrayPortalGroupVecVariable<
typename SourceArrayHandleType::PortalControl, typename OffsetsArrayHandleType::PortalControl>; typename SourceArrayHandleType::PortalControl,
typename OffsetsArrayHandleType::PortalControl>;
using PortalConstType = vtkm::exec::internal::ArrayPortalGroupVecVariable< using PortalConstType = vtkm::exec::internal::ArrayPortalGroupVecVariable<
typename SourceArrayHandleType::PortalConstControl, typename SourceArrayHandleType::PortalConstControl,
typename OffsetsArrayHandleType::PortalConstControl>; typename OffsetsArrayHandleType::PortalConstControl>;
@ -402,9 +405,10 @@ public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleGroupVecVariable, ArrayHandleGroupVecVariable,
(ArrayHandleGroupVecVariable<SourceArrayHandleType, OffsetsArrayHandleType>), (ArrayHandleGroupVecVariable<SourceArrayHandleType, OffsetsArrayHandleType>),
(vtkm::cont::ArrayHandle<vtkm::VecFromPortal<typename SourceArrayHandleType::PortalControl>, (vtkm::cont::ArrayHandle<
vtkm::cont::internal::StorageTagGroupVecVariable< vtkm::VecFromPortal<typename SourceArrayHandleType::PortalControl>,
SourceArrayHandleType, OffsetsArrayHandleType>>)); vtkm::cont::internal::StorageTagGroupVecVariable<SourceArrayHandleType,
OffsetsArrayHandleType>>));
using ComponentType = typename SourceArrayHandleType::ValueType; using ComponentType = typename SourceArrayHandleType::ValueType;
@ -496,7 +500,8 @@ VTKM_CONT void DoConvertNumComponentsToOffsets(const NumComponentsArrayType& num
template <typename NumComponentsArrayType, typename OffsetsStorage> template <typename NumComponentsArrayType, typename OffsetsStorage>
VTKM_CONT void ConvertNumComponentsToOffsets( VTKM_CONT void ConvertNumComponentsToOffsets(
const NumComponentsArrayType& numComponentsArray, const NumComponentsArrayType& numComponentsArray,
vtkm::cont::ArrayHandle<vtkm::Id, OffsetsStorage>& offsetsArray, vtkm::Id& sourceArraySize) vtkm::cont::ArrayHandle<vtkm::Id, OffsetsStorage>& offsetsArray,
vtkm::Id& sourceArraySize)
{ {
VTKM_IS_ARRAY_HANDLE(NumComponentsArrayType); VTKM_IS_ARRAY_HANDLE(NumComponentsArrayType);
@ -513,7 +518,8 @@ VTKM_CONT void ConvertNumComponentsToOffsets(
} }
template <typename NumComponentsArrayType> template <typename NumComponentsArrayType>
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Id> ConvertNumComponentsToOffsets( VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Id> ConvertNumComponentsToOffsets(
const NumComponentsArrayType& numComponentsArray, vtkm::Id& sourceArraySize) const NumComponentsArrayType& numComponentsArray,
vtkm::Id& sourceArraySize)
{ {
VTKM_IS_ARRAY_HANDLE(NumComponentsArrayType); VTKM_IS_ARRAY_HANDLE(NumComponentsArrayType);

@ -115,7 +115,8 @@ private:
typedef typename detail::ArrayHandleImplicitTraits<T, FunctorType> ArrayTraits; typedef typename detail::ArrayHandleImplicitTraits<T, FunctorType> ArrayTraits;
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleImplicit, (ArrayHandleImplicit<T, FunctorType>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleImplicit,
(ArrayHandleImplicit<T, FunctorType>),
(typename ArrayTraits::Superclass)); (typename ArrayTraits::Superclass));
VTKM_CONT VTKM_CONT
@ -131,7 +132,8 @@ public:
template <typename T, typename FunctorType> template <typename T, typename FunctorType>
VTKM_CONT vtkm::cont::ArrayHandleImplicit<T, FunctorType> make_ArrayHandleImplicit( VTKM_CONT vtkm::cont::ArrayHandleImplicit<T, FunctorType> make_ArrayHandleImplicit(
FunctorType functor, vtkm::Id length) FunctorType functor,
vtkm::Id length)
{ {
return ArrayHandleImplicit<T, FunctorType>(functor, length); return ArrayHandleImplicit<T, FunctorType>(functor, length);
} }

@ -205,7 +205,8 @@ private:
template <typename IndexArrayType, typename ValueArrayType, typename Device> template <typename IndexArrayType, typename ValueArrayType, typename Device>
class ArrayTransfer<typename ValueArrayType::ValueType, class ArrayTransfer<typename ValueArrayType::ValueType,
StorageTagPermutation<IndexArrayType, ValueArrayType>, Device> StorageTagPermutation<IndexArrayType, ValueArrayType>,
Device>
{ {
public: public:
typedef typename ValueArrayType::ValueType ValueType; typedef typename ValueArrayType::ValueType ValueType;
@ -342,7 +343,8 @@ class ArrayHandlePermutation
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandlePermutation, (ArrayHandlePermutation<IndexArrayHandleType, ValueArrayHandleType>), ArrayHandlePermutation,
(ArrayHandlePermutation<IndexArrayHandleType, ValueArrayHandleType>),
(vtkm::cont::ArrayHandle< (vtkm::cont::ArrayHandle<
typename ValueArrayHandleType::ValueType, typename ValueArrayHandleType::ValueType,
internal::StorageTagPermutation<IndexArrayHandleType, ValueArrayHandleType>>)); internal::StorageTagPermutation<IndexArrayHandleType, ValueArrayHandleType>>));

@ -238,7 +238,8 @@ class ArrayHandleReverse : public vtkm::cont::ArrayHandle<typename ArrayHandleTy
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleReverse, (ArrayHandleReverse<ArrayHandleType>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleReverse,
(ArrayHandleReverse<ArrayHandleType>),
(vtkm::cont::ArrayHandle<typename ArrayHandleType::ValueType, (vtkm::cont::ArrayHandle<typename ArrayHandleType::ValueType,
StorageTagReverse<ArrayHandleType>>)); StorageTagReverse<ArrayHandleType>>));

@ -47,7 +47,9 @@ public:
} }
VTKM_CONT VTKM_CONT
ArrayPortalStreaming(const PortalType& inputPortal, vtkm::Id blockIndex, vtkm::Id blockSize, ArrayPortalStreaming(const PortalType& inputPortal,
vtkm::Id blockIndex,
vtkm::Id blockSize,
vtkm::Id curBlockSize) vtkm::Id curBlockSize)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, BlockIndex(blockIndex) , BlockIndex(blockIndex)
@ -137,7 +139,9 @@ public:
} }
VTKM_CONT VTKM_CONT
Storage(const ArrayHandleInputType inputArray, vtkm::Id blockSize, vtkm::Id blockIndex, Storage(const ArrayHandleInputType inputArray,
vtkm::Id blockSize,
vtkm::Id blockIndex,
vtkm::Id curBlockSize) vtkm::Id curBlockSize)
: InputArray(inputArray) : InputArray(inputArray)
, BlockSize(blockSize) , BlockSize(blockSize)
@ -158,8 +162,8 @@ public:
PortalConstType GetPortalConst() const PortalConstType GetPortalConst() const
{ {
VTKM_ASSERT(this->Valid); VTKM_ASSERT(this->Valid);
return PortalConstType(this->InputArray.GetPortalConstControl(), BlockSize, BlockIndex, return PortalConstType(
CurBlockSize); this->InputArray.GetPortalConstControl(), BlockSize, BlockIndex, CurBlockSize);
} }
VTKM_CONT VTKM_CONT
@ -226,7 +230,8 @@ class ArrayHandleStreaming
StorageTagStreaming<ArrayHandleInputType>> StorageTagStreaming<ArrayHandleInputType>>
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleStreaming, (ArrayHandleStreaming<ArrayHandleInputType>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleStreaming,
(ArrayHandleStreaming<ArrayHandleInputType>),
(vtkm::cont::ArrayHandle<typename ArrayHandleInputType::ValueType, (vtkm::cont::ArrayHandle<typename ArrayHandleInputType::ValueType,
StorageTagStreaming<ArrayHandleInputType>>)); StorageTagStreaming<ArrayHandleInputType>>));
@ -235,8 +240,10 @@ private:
public: public:
VTKM_CONT VTKM_CONT
ArrayHandleStreaming(const ArrayHandleInputType& inputArray, const vtkm::Id blockIndex, ArrayHandleStreaming(const ArrayHandleInputType& inputArray,
const vtkm::Id blockSize, const vtkm::Id curBlockSize) const vtkm::Id blockIndex,
const vtkm::Id blockSize,
const vtkm::Id curBlockSize)
: Superclass(StorageType(inputArray, blockIndex, blockSize, curBlockSize)) : Superclass(StorageType(inputArray, blockIndex, blockSize, curBlockSize))
{ {
this->GetPortalConstControl().SetBlockIndex(blockIndex); this->GetPortalConstControl().SetBlockIndex(blockIndex);

@ -52,7 +52,9 @@ typedef vtkm::cont::internal::NullFunctorType NullFunctorType;
/// \brief An array portal that transforms a value from another portal. /// \brief An array portal that transforms a value from another portal.
/// ///
template <typename ValueType_, typename PortalType_, typename FunctorType_, template <typename ValueType_,
typename PortalType_,
typename FunctorType_,
typename InverseFunctorType_ = NullFunctorType> typename InverseFunctorType_ = NullFunctorType>
class VTKM_ALWAYS_EXPORT ArrayPortalTransform; class VTKM_ALWAYS_EXPORT ArrayPortalTransform;
@ -117,7 +119,9 @@ protected:
FunctorType Functor; FunctorType Functor;
}; };
template <typename ValueType_, typename PortalType_, typename FunctorType_, template <typename ValueType_,
typename PortalType_,
typename FunctorType_,
typename InverseFunctorType_> typename InverseFunctorType_>
class VTKM_ALWAYS_EXPORT ArrayPortalTransform class VTKM_ALWAYS_EXPORT ArrayPortalTransform
: public ArrayPortalTransform<ValueType_, PortalType_, FunctorType_, NullFunctorType> : public ArrayPortalTransform<ValueType_, PortalType_, FunctorType_, NullFunctorType>
@ -174,7 +178,9 @@ namespace cont
namespace internal namespace internal
{ {
template <typename ValueType, typename ArrayHandleType, typename FunctorType, template <typename ValueType,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType = NullFunctorType> typename InverseFunctorType = NullFunctorType>
struct VTKM_ALWAYS_EXPORT StorageTagTransform struct VTKM_ALWAYS_EXPORT StorageTagTransform
{ {
@ -194,8 +200,9 @@ public:
typedef void* IteratorType; typedef void* IteratorType;
}; };
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<ValueType,
ValueType, typename ArrayHandleType::PortalConstControl, FunctorType> typename ArrayHandleType::PortalConstControl,
FunctorType>
PortalConstType; PortalConstType;
VTKM_CONT VTKM_CONT
@ -275,11 +282,15 @@ class Storage<T, StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFun
public: public:
typedef T ValueType; typedef T ValueType;
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<ValueType,
ValueType, typename ArrayHandleType::PortalControl, FunctorType, InverseFunctorType> typename ArrayHandleType::PortalControl,
FunctorType,
InverseFunctorType>
PortalType; PortalType;
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<ValueType,
ValueType, typename ArrayHandleType::PortalConstControl, FunctorType, InverseFunctorType> typename ArrayHandleType::PortalConstControl,
FunctorType,
InverseFunctorType>
PortalConstType; PortalConstType;
VTKM_CONT VTKM_CONT
@ -289,7 +300,8 @@ public:
} }
VTKM_CONT VTKM_CONT
Storage(const ArrayHandleType& array, const FunctorType& functor, Storage(const ArrayHandleType& array,
const FunctorType& functor,
const InverseFunctorType& inverseFunctor) const InverseFunctorType& inverseFunctor)
: Array(array) : Array(array)
, Functor(functor) , Functor(functor)
@ -309,8 +321,8 @@ public:
PortalConstType GetPortalConst() const PortalConstType GetPortalConst() const
{ {
VTKM_ASSERT(this->Valid); VTKM_ASSERT(this->Valid);
return PortalConstType(this->Array.GetPortalConstControl(), this->Functor, return PortalConstType(
this->InverseFunctor); this->Array.GetPortalConstControl(), this->Functor, this->InverseFunctor);
} }
VTKM_CONT VTKM_CONT
@ -372,7 +384,9 @@ public:
//meant to be an invalid writeable execution portal //meant to be an invalid writeable execution portal
typedef typename StorageType::PortalType PortalExecution; typedef typename StorageType::PortalType PortalExecution;
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<
ValueType, typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst, FunctorType> ValueType,
typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst,
FunctorType>
PortalConstExecution; PortalConstExecution;
VTKM_CONT VTKM_CONT
@ -427,9 +441,13 @@ private:
FunctorType Functor; FunctorType Functor;
}; };
template <typename T, typename ArrayHandleType, typename FunctorType, typename InverseFunctorType, template <typename T,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType,
typename Device> typename Device>
class ArrayTransfer<T, StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>, class ArrayTransfer<T,
StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>,
Device> Device>
{ {
typedef StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType> StorageTag; typedef StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType> StorageTag;
@ -442,11 +460,15 @@ public:
typedef typename StorageType::PortalConstType PortalConstControl; typedef typename StorageType::PortalConstType PortalConstControl;
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<
ValueType, typename ArrayHandleType::template ExecutionTypes<Device>::Portal, FunctorType, ValueType,
typename ArrayHandleType::template ExecutionTypes<Device>::Portal,
FunctorType,
InverseFunctorType> InverseFunctorType>
PortalExecution; PortalExecution;
typedef vtkm::exec::internal::ArrayPortalTransform< typedef vtkm::exec::internal::ArrayPortalTransform<
ValueType, typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst, FunctorType, ValueType,
typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst,
FunctorType,
InverseFunctorType> InverseFunctorType>
PortalConstExecution; PortalConstExecution;
@ -464,22 +486,22 @@ public:
VTKM_CONT VTKM_CONT
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData))
{ {
return PortalConstExecution(this->Array.PrepareForInput(Device()), this->Functor, return PortalConstExecution(
this->InverseFunctor); this->Array.PrepareForInput(Device()), this->Functor, this->InverseFunctor);
} }
VTKM_CONT VTKM_CONT
PortalExecution PrepareForInPlace(bool& vtkmNotUsed(updateData)) PortalExecution PrepareForInPlace(bool& vtkmNotUsed(updateData))
{ {
return PortalExecution(this->Array.PrepareForInPlace(Device()), this->Functor, return PortalExecution(
this->InverseFunctor); this->Array.PrepareForInPlace(Device()), this->Functor, this->InverseFunctor);
} }
VTKM_CONT VTKM_CONT
PortalExecution PrepareForOutput(vtkm::Id numberOfValues) PortalExecution PrepareForOutput(vtkm::Id numberOfValues)
{ {
return PortalExecution(this->Array.PrepareForOutput(numberOfValues, Device()), this->Functor, return PortalExecution(
this->InverseFunctor); this->Array.PrepareForOutput(numberOfValues, Device()), this->Functor, this->InverseFunctor);
} }
VTKM_CONT VTKM_CONT
@ -514,7 +536,9 @@ private:
/// the functor operator should work in both the control and execution /// the functor operator should work in both the control and execution
/// environments. /// environments.
/// ///
template <typename T, typename ArrayHandleType, typename FunctorType, template <typename T,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType = internal::NullFunctorType> typename InverseFunctorType = internal::NullFunctorType>
class ArrayHandleTransform; class ArrayHandleTransform;
@ -529,7 +553,8 @@ class ArrayHandleTransform<T, ArrayHandleType, FunctorType, internal::NullFuncto
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleTransform, (ArrayHandleTransform<T, ArrayHandleType, FunctorType>), ArrayHandleTransform,
(ArrayHandleTransform<T, ArrayHandleType, FunctorType>),
(vtkm::cont::ArrayHandle<T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType>>)); (vtkm::cont::ArrayHandle<T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType>>));
private: private:
@ -549,7 +574,8 @@ public:
template <typename T, typename HandleType, typename FunctorType> template <typename T, typename HandleType, typename FunctorType>
VTKM_CONT vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType> make_ArrayHandleTransform( VTKM_CONT vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType> make_ArrayHandleTransform(
HandleType handle, FunctorType functor) HandleType handle,
FunctorType functor)
{ {
return ArrayHandleTransform<T, HandleType, FunctorType>(handle, functor); return ArrayHandleTransform<T, HandleType, FunctorType>(handle, functor);
} }
@ -559,7 +585,8 @@ VTKM_CONT vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType> make_Arra
template <typename T, typename ArrayHandleType, typename FunctorType, typename InverseFunctorType> template <typename T, typename ArrayHandleType, typename FunctorType, typename InverseFunctorType>
class ArrayHandleTransform class ArrayHandleTransform
: public vtkm::cont::ArrayHandle< : public vtkm::cont::ArrayHandle<
T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>> T,
internal::StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>>
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType); VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
@ -568,13 +595,15 @@ public:
ArrayHandleTransform, ArrayHandleTransform,
(ArrayHandleTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>), (ArrayHandleTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>),
(vtkm::cont::ArrayHandle< (vtkm::cont::ArrayHandle<
T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>>)); T,
internal::StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType>>));
private: private:
typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType; typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType;
public: public:
ArrayHandleTransform(const ArrayHandleType& handle, const FunctorType& functor = FunctorType(), ArrayHandleTransform(const ArrayHandleType& handle,
const FunctorType& functor = FunctorType(),
const InverseFunctorType& inverseFunctor = InverseFunctorType()) const InverseFunctorType& inverseFunctor = InverseFunctorType())
: Superclass(StorageType(handle, functor, inverseFunctor)) : Superclass(StorageType(handle, functor, inverseFunctor))
{ {
@ -585,8 +614,8 @@ template <typename T, typename HandleType, typename FunctorType, typename Invers
VTKM_CONT vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType, InverseFunctorType> VTKM_CONT vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType, InverseFunctorType>
make_ArrayHandleTransform(HandleType handle, FunctorType functor, InverseFunctorType inverseFunctor) make_ArrayHandleTransform(HandleType handle, FunctorType functor, InverseFunctorType inverseFunctor)
{ {
return ArrayHandleTransform<T, HandleType, FunctorType, InverseFunctorType>(handle, functor, return ArrayHandleTransform<T, HandleType, FunctorType, InverseFunctorType>(
inverseFunctor); handle, functor, inverseFunctor);
} }
} }
} // namespace vtkm::cont } // namespace vtkm::cont

@ -139,7 +139,8 @@ class Storage<T, StorageTagZip<FirstHandleType, SecondHandleType>>
public: public:
typedef T ValueType; typedef T ValueType;
typedef vtkm::exec::internal::ArrayPortalZip<ValueType, typename FirstHandleType::PortalControl, typedef vtkm::exec::internal::ArrayPortalZip<ValueType,
typename FirstHandleType::PortalControl,
typename SecondHandleType::PortalControl> typename SecondHandleType::PortalControl>
PortalType; PortalType;
typedef vtkm::exec::internal::ArrayPortalZip<ValueType, typedef vtkm::exec::internal::ArrayPortalZip<ValueType,
@ -226,12 +227,14 @@ public:
typedef typename StorageType::PortalConstType PortalConstControl; typedef typename StorageType::PortalConstType PortalConstControl;
typedef vtkm::exec::internal::ArrayPortalZip< typedef vtkm::exec::internal::ArrayPortalZip<
ValueType, typename FirstHandleType::template ExecutionTypes<Device>::Portal, ValueType,
typename FirstHandleType::template ExecutionTypes<Device>::Portal,
typename SecondHandleType::template ExecutionTypes<Device>::Portal> typename SecondHandleType::template ExecutionTypes<Device>::Portal>
PortalExecution; PortalExecution;
typedef vtkm::exec::internal::ArrayPortalZip< typedef vtkm::exec::internal::ArrayPortalZip<
ValueType, typename FirstHandleType::template ExecutionTypes<Device>::PortalConst, ValueType,
typename FirstHandleType::template ExecutionTypes<Device>::PortalConst,
typename SecondHandleType::template ExecutionTypes<Device>::PortalConst> typename SecondHandleType::template ExecutionTypes<Device>::PortalConst>
PortalConstExecution; PortalConstExecution;
@ -316,7 +319,8 @@ class ArrayHandleZip
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS( VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleZip, (ArrayHandleZip<FirstHandleType, SecondHandleType>), ArrayHandleZip,
(ArrayHandleZip<FirstHandleType, SecondHandleType>),
(typename internal::ArrayHandleZipTraits<FirstHandleType, SecondHandleType>::Superclass)); (typename internal::ArrayHandleZipTraits<FirstHandleType, SecondHandleType>::Superclass));
private: private:
@ -335,7 +339,8 @@ public:
/// ///
template <typename FirstHandleType, typename SecondHandleType> template <typename FirstHandleType, typename SecondHandleType>
VTKM_CONT vtkm::cont::ArrayHandleZip<FirstHandleType, SecondHandleType> make_ArrayHandleZip( VTKM_CONT vtkm::cont::ArrayHandleZip<FirstHandleType, SecondHandleType> make_ArrayHandleZip(
const FirstHandleType& first, const SecondHandleType& second) const FirstHandleType& first,
const SecondHandleType& second)
{ {
return ArrayHandleZip<FirstHandleType, SecondHandleType>(first, second); return ArrayHandleZip<FirstHandleType, SecondHandleType>(first, second);
} }

@ -104,9 +104,10 @@ VTKM_CONT
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle< const vtkm::cont::ArrayHandle<
vtkm::Vec<vtkm::Float32, 3>, vtkm::Vec<vtkm::Float32, 3>,
typename vtkm::cont::ArrayHandleCompositeVector<vtkm::Vec<vtkm::Float32, 3>( typename vtkm::cont::ArrayHandleCompositeVector<
vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::Vec<vtkm::Float32, 3>(vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>)>::StorageTag>& input, vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>)>::StorageTag>& input,
vtkm::cont::RuntimeDeviceTracker tracker) vtkm::cont::RuntimeDeviceTracker tracker)
{ {
return detail::ArrayRangeComputeImpl(input, tracker); return detail::ArrayRangeComputeImpl(input, tracker);
@ -116,9 +117,10 @@ VTKM_CONT
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle< const vtkm::cont::ArrayHandle<
vtkm::Vec<vtkm::Float64, 3>, vtkm::Vec<vtkm::Float64, 3>,
typename vtkm::cont::ArrayHandleCompositeVector<vtkm::Vec<vtkm::Float64, 3>( typename vtkm::cont::ArrayHandleCompositeVector<
vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::Vec<vtkm::Float64, 3>(vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>)>::StorageTag>& input, vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>)>::StorageTag>& input,
vtkm::cont::RuntimeDeviceTracker tracker) vtkm::cont::RuntimeDeviceTracker tracker)
{ {
return detail::ArrayRangeComputeImpl(input, tracker); return detail::ArrayRangeComputeImpl(input, tracker);

@ -93,7 +93,8 @@ VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::UInt8, 4, vtkm::cont::StorageTagBasic)
VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float32, 4, vtkm::cont::StorageTagBasic); VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float32, 4, vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float64, 4, vtkm::cont::StorageTagBasic); VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float64, 4, vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::FloatDefault, 3, VTKM_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::FloatDefault,
3,
vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag); vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag);
#undef VTKM_ARRAY_RANGE_COMPUTE_EXPORT_T #undef VTKM_ARRAY_RANGE_COMPUTE_EXPORT_T
@ -105,9 +106,10 @@ VTKM_CONT
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle< const vtkm::cont::ArrayHandle<
vtkm::Vec<vtkm::Float32, 3>, vtkm::Vec<vtkm::Float32, 3>,
typename vtkm::cont::ArrayHandleCompositeVector<vtkm::Vec<vtkm::Float32, 3>( typename vtkm::cont::ArrayHandleCompositeVector<
vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::Vec<vtkm::Float32, 3>(vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>)>::StorageTag>& input, vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>)>::StorageTag>& input,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()); vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker());
VTKM_CONT_EXPORT VTKM_CONT_EXPORT
@ -115,16 +117,18 @@ VTKM_CONT
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle< const vtkm::cont::ArrayHandle<
vtkm::Vec<vtkm::Float64, 3>, vtkm::Vec<vtkm::Float64, 3>,
typename vtkm::cont::ArrayHandleCompositeVector<vtkm::Vec<vtkm::Float64, 3>( typename vtkm::cont::ArrayHandleCompositeVector<
vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::Vec<vtkm::Float64, 3>(vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>)>::StorageTag>& input, vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>)>::StorageTag>& input,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()); vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker());
// Implementation of cartesian products // Implementation of cartesian products
template <typename T, typename ArrayType1, typename ArrayType2, typename ArrayType3> template <typename T, typename ArrayType1, typename ArrayType2, typename ArrayType3>
VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle< const vtkm::cont::ArrayHandle<
T, vtkm::cont::internal::StorageTagCartesianProduct<ArrayType1, ArrayType2, ArrayType3>>& input, T,
vtkm::cont::internal::StorageTagCartesianProduct<ArrayType1, ArrayType2, ArrayType3>>& input,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()) vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
{ {
vtkm::cont::ArrayHandle<vtkm::Range> result; vtkm::cont::ArrayHandle<vtkm::Range> result;

@ -90,7 +90,8 @@ struct ArrayRangeComputeFunctor
template <typename ArrayHandleType> template <typename ArrayHandleType>
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl( inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
const ArrayHandleType& input, vtkm::cont::RuntimeDeviceTracker tracker) const ArrayHandleType& input,
vtkm::cont::RuntimeDeviceTracker tracker)
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType); VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
@ -108,7 +109,8 @@ inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
template <typename ArrayHandleType> template <typename ArrayHandleType>
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const ArrayHandleType& input, vtkm::cont::RuntimeDeviceTracker tracker) const ArrayHandleType& input,
vtkm::cont::RuntimeDeviceTracker tracker)
{ {
VTKM_IS_ARRAY_HANDLE(ArrayHandleType); VTKM_IS_ARRAY_HANDLE(ArrayHandleType);

@ -73,15 +73,18 @@ template <typename ShapeStorageTag = VTKM_DEFAULT_SHAPE_STORAGE_TAG,
typename OffsetsStorageTag = VTKM_DEFAULT_OFFSETS_STORAGE_TAG> typename OffsetsStorageTag = VTKM_DEFAULT_OFFSETS_STORAGE_TAG>
class VTKM_ALWAYS_EXPORT CellSetExplicit : public CellSet class VTKM_ALWAYS_EXPORT CellSetExplicit : public CellSet
{ {
typedef CellSetExplicit<ShapeStorageTag, NumIndicesStorageTag, ConnectivityStorageTag, typedef CellSetExplicit<ShapeStorageTag,
NumIndicesStorageTag,
ConnectivityStorageTag,
OffsetsStorageTag> OffsetsStorageTag>
Thisclass; Thisclass;
template <typename FromTopology, typename ToTopology> template <typename FromTopology, typename ToTopology>
struct ConnectivityChooser struct ConnectivityChooser
{ {
typedef typename detail::CellSetExplicitConnectivityChooser< typedef
Thisclass, FromTopology, ToTopology>::ConnectivityType ConnectivityType; typename detail::CellSetExplicitConnectivityChooser<Thisclass, FromTopology, ToTopology>::
ConnectivityType ConnectivityType;
typedef typename ConnectivityType::ShapeArrayType ShapeArrayType; typedef typename ConnectivityType::ShapeArrayType ShapeArrayType;
typedef typename ConnectivityType::NumIndicesArrayType NumIndicesArrayType; typedef typename ConnectivityType::NumIndicesArrayType NumIndicesArrayType;
@ -296,14 +299,16 @@ public:
typedef typename ConnectivityTypes::IndexOffsetArrayType::template ExecutionTypes< typedef typename ConnectivityTypes::IndexOffsetArrayType::template ExecutionTypes<
DeviceAdapter>::PortalConst IndexOffsetPortalType; DeviceAdapter>::PortalConst IndexOffsetPortalType;
typedef vtkm::exec::ConnectivityExplicit<ShapePortalType, IndicePortalType, typedef vtkm::exec::ConnectivityExplicit<ShapePortalType,
ConnectivityPortalType, IndexOffsetPortalType> IndicePortalType,
ConnectivityPortalType,
IndexOffsetPortalType>
ExecObjectType; ExecObjectType;
}; };
template <typename Device, typename FromTopology, typename ToTopology> template <typename Device, typename FromTopology, typename ToTopology>
typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType PrepareForInput( typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType
Device, FromTopology, ToTopology) const PrepareForInput(Device, FromTopology, ToTopology) const
{ {
this->BuildConnectivity(Device(), FromTopology(), ToTopology()); this->BuildConnectivity(Device(), FromTopology(), ToTopology());
@ -322,7 +327,9 @@ public:
template <typename Device, typename FromTopology, typename ToTopology> template <typename Device, typename FromTopology, typename ToTopology>
VTKM_CONT void BuildConnectivity(Device, FromTopology, ToTopology) const VTKM_CONT void BuildConnectivity(Device, FromTopology, ToTopology) const
{ {
typedef CellSetExplicit<ShapeStorageTag, NumIndicesStorageTag, ConnectivityStorageTag, typedef CellSetExplicit<ShapeStorageTag,
NumIndicesStorageTag,
ConnectivityStorageTag,
OffsetsStorageTag> OffsetsStorageTag>
CSE; CSE;
@ -334,7 +341,8 @@ public:
} }
template <typename Device> template <typename Device>
VTKM_CONT void CreateConnectivity(Device, vtkm::TopologyElementTagPoint, VTKM_CONT void CreateConnectivity(Device,
vtkm::TopologyElementTagPoint,
vtkm::TopologyElementTagCell) vtkm::TopologyElementTagCell)
{ {
// nothing to do // nothing to do
@ -344,7 +352,9 @@ public:
class ExpandIndices : public vtkm::worklet::WorkletMapField class ExpandIndices : public vtkm::worklet::WorkletMapField
{ {
public: public:
typedef void ControlSignature(FieldIn<> cellIndex, FieldIn<> offset, FieldIn<> numIndices, typedef void ControlSignature(FieldIn<> cellIndex,
FieldIn<> offset,
FieldIn<> numIndices,
WholeArrayOut<> cellIndices); WholeArrayOut<> cellIndices);
typedef void ExecutionSignature(_1, _2, _3, _4); typedef void ExecutionSignature(_1, _2, _3, _4);
typedef _1 InputDomain; typedef _1 InputDomain;
@ -353,8 +363,10 @@ public:
ExpandIndices() {} ExpandIndices() {}
template <typename PortalType> template <typename PortalType>
VTKM_EXEC void operator()(const vtkm::Id& cellIndex, const vtkm::Id& offset, VTKM_EXEC void operator()(const vtkm::Id& cellIndex,
const vtkm::Id& numIndices, const PortalType& cellIndices) const const vtkm::Id& offset,
const vtkm::Id& numIndices,
const PortalType& cellIndices) const
{ {
VTKM_ASSERT(cellIndices.GetNumberOfValues() >= offset + numIndices); VTKM_ASSERT(cellIndices.GetNumberOfValues() >= offset + numIndices);
vtkm::Id startIndex = offset; vtkm::Id startIndex = offset;
@ -366,7 +378,8 @@ public:
}; };
template <typename Device> template <typename Device>
VTKM_CONT void CreateConnectivity(Device, vtkm::TopologyElementTagCell, VTKM_CONT void CreateConnectivity(Device,
vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint) vtkm::TopologyElementTagPoint)
{ {
// PointToCell connectivity array (point indices) will be // PointToCell connectivity array (point indices) will be
@ -397,7 +410,9 @@ public:
this->PointToCell.BuildIndexOffsets(Device()); this->PointToCell.BuildIndexOffsets(Device());
vtkm::worklet::DispatcherMapField<ExpandIndices, Device> expandDispatcher; vtkm::worklet::DispatcherMapField<ExpandIndices, Device> expandDispatcher;
expandDispatcher.Invoke(index, this->PointToCell.IndexOffsets, this->PointToCell.NumIndices, expandDispatcher.Invoke(index,
this->PointToCell.IndexOffsets,
this->PointToCell.NumIndices,
this->CellToPoint.Connectivity); this->CellToPoint.Connectivity);
// SortByKey where key is PointToCell connectivity and value is the expanded cellIndex // SortByKey where key is PointToCell connectivity and value is the expanded cellIndex
@ -490,9 +505,11 @@ private:
return this->Ivar; \ return this->Ivar; \
} }
VTKM_GET_CONNECTIVITY_METHOD(vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, VTKM_GET_CONNECTIVITY_METHOD(vtkm::TopologyElementTagPoint,
vtkm::TopologyElementTagCell,
PointToCell) PointToCell)
VTKM_GET_CONNECTIVITY_METHOD(vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, VTKM_GET_CONNECTIVITY_METHOD(vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint,
CellToPoint) CellToPoint)
#undef VTKM_GET_CONNECTIVITY_METHOD #undef VTKM_GET_CONNECTIVITY_METHOD
@ -512,15 +529,17 @@ namespace detail
template <typename Storage1, typename Storage2, typename Storage3, typename Storage4> template <typename Storage1, typename Storage2, typename Storage3, typename Storage4>
struct CellSetExplicitConnectivityChooser< struct CellSetExplicitConnectivityChooser<
vtkm::cont::CellSetExplicit<Storage1, Storage2, Storage3, Storage4>, vtkm::cont::CellSetExplicit<Storage1, Storage2, Storage3, Storage4>,
vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell> vtkm::TopologyElementTagPoint,
vtkm::TopologyElementTagCell>
{ {
typedef vtkm::cont::internal::ConnectivityExplicitInternals<Storage1, Storage2, Storage3, typedef vtkm::cont::internal::
Storage4> ConnectivityExplicitInternals<Storage1, Storage2, Storage3, Storage4>
ConnectivityType; ConnectivityType;
}; };
template <typename CellSetType> template <typename CellSetType>
struct CellSetExplicitConnectivityChooser<CellSetType, vtkm::TopologyElementTagCell, struct CellSetExplicitConnectivityChooser<CellSetType,
vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint> vtkm::TopologyElementTagPoint>
{ {
//only specify the shape type as it will be constant as everything //only specify the shape type as it will be constant as everything

@ -55,8 +55,10 @@ template <typename ShapeStorageTag = VTKM_DEFAULT_SHAPE_STORAGE_TAG,
typename ConnectivityStorageTag = VTKM_DEFAULT_CONNECTIVITY_STORAGE_TAG, typename ConnectivityStorageTag = VTKM_DEFAULT_CONNECTIVITY_STORAGE_TAG,
typename OffsetsStorageTag = VTKM_DEFAULT_OFFSETS_STORAGE_TAG> typename OffsetsStorageTag = VTKM_DEFAULT_OFFSETS_STORAGE_TAG>
struct VTKM_ALWAYS_EXPORT CellSetListTagExplicit struct VTKM_ALWAYS_EXPORT CellSetListTagExplicit
: vtkm::ListTagBase<vtkm::cont::CellSetExplicit<ShapeStorageTag, NumIndicesStorageTag, : vtkm::ListTagBase<vtkm::cont::CellSetExplicit<ShapeStorageTag,
ConnectivityStorageTag, OffsetsStorageTag>> NumIndicesStorageTag,
ConnectivityStorageTag,
OffsetsStorageTag>>
{ {
}; };
@ -64,9 +66,10 @@ struct VTKM_ALWAYS_EXPORT CellSetListTagExplicitDefault : CellSetListTagExplicit
{ {
}; };
struct VTKM_ALWAYS_EXPORT CellSetListTagCommon struct VTKM_ALWAYS_EXPORT CellSetListTagCommon : vtkm::ListTagBase<vtkm::cont::CellSetStructured<2>,
: vtkm::ListTagBase<vtkm::cont::CellSetStructured<2>, vtkm::cont::CellSetStructured<3>, vtkm::cont::CellSetStructured<3>,
vtkm::cont::CellSetExplicit<>, vtkm::cont::CellSetSingleType<>> vtkm::cont::CellSetExplicit<>,
vtkm::cont::CellSetSingleType<>>
{ {
}; };

@ -44,7 +44,8 @@ class CellSetGeneralPermutation : public CellSet
public: public:
VTKM_CONT VTKM_CONT
CellSetGeneralPermutation(const PermutationArrayHandleType& validCellIds, CellSetGeneralPermutation(const PermutationArrayHandleType& validCellIds,
const OriginalCellSet& cellset, const std::string& name) const OriginalCellSet& cellset,
const std::string& name)
: CellSet(name) : CellSet(name)
, ValidCellIds(validCellIds) , ValidCellIds(validCellIds)
, FullCellSet(cellset) , FullCellSet(cellset)
@ -94,15 +95,15 @@ public:
typedef typename PermutationArrayHandleType::template ExecutionTypes<Device>::PortalConst typedef typename PermutationArrayHandleType::template ExecutionTypes<Device>::PortalConst
ExecPortalType; ExecPortalType;
typedef typename OriginalCellSet::template ExecutionTypes< typedef typename OriginalCellSet::template ExecutionTypes<Device, FromTopology, ToTopology>::
Device, FromTopology, ToTopology>::ExecObjectType OrigExecObjectType; ExecObjectType OrigExecObjectType;
typedef vtkm::exec::ConnectivityPermuted<ExecPortalType, OrigExecObjectType> ExecObjectType; typedef vtkm::exec::ConnectivityPermuted<ExecPortalType, OrigExecObjectType> ExecObjectType;
}; };
template <typename Device, typename FromTopology, typename ToTopology> template <typename Device, typename FromTopology, typename ToTopology>
typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType PrepareForInput( typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType
Device d, FromTopology f, ToTopology t) const PrepareForInput(Device d, FromTopology f, ToTopology t) const
{ {
typedef typedef
typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType ConnectivityType; typename ExecutionTypes<Device, FromTopology, ToTopology>::ExecObjectType ConnectivityType;
@ -143,7 +144,8 @@ class CellSetPermutation
public: public:
VTKM_CONT VTKM_CONT
CellSetPermutation(const PermutationArrayHandleType& validCellIds, const OriginalCellSet& cellset, CellSetPermutation(const PermutationArrayHandleType& validCellIds,
const OriginalCellSet& cellset,
const std::string& name = std::string()) const std::string& name = std::string())
: ParentType(validCellIds, cellset, name) : ParentType(validCellIds, cellset, name)
{ {
@ -166,19 +168,21 @@ public:
template <typename OriginalCellSet, typename PermutationArrayHandleType> template <typename OriginalCellSet, typename PermutationArrayHandleType>
vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType> make_CellSetPermutation( vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType> make_CellSetPermutation(
const PermutationArrayHandleType& cellIndexMap, const OriginalCellSet& cellSet, const PermutationArrayHandleType& cellIndexMap,
const OriginalCellSet& cellSet,
const std::string& name) const std::string& name)
{ {
VTKM_IS_CELL_SET(OriginalCellSet); VTKM_IS_CELL_SET(OriginalCellSet);
VTKM_IS_ARRAY_HANDLE(PermutationArrayHandleType); VTKM_IS_ARRAY_HANDLE(PermutationArrayHandleType);
return vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType>(cellIndexMap, return vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType>(
cellSet, name); cellIndexMap, cellSet, name);
} }
template <typename OriginalCellSet, typename PermutationArrayHandleType> template <typename OriginalCellSet, typename PermutationArrayHandleType>
vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType> make_CellSetPermutation( vtkm::cont::CellSetPermutation<OriginalCellSet, PermutationArrayHandleType> make_CellSetPermutation(
const PermutationArrayHandleType& cellIndexMap, const OriginalCellSet& cellSet) const PermutationArrayHandleType& cellIndexMap,
const OriginalCellSet& cellSet)
{ {
VTKM_IS_CELL_SET(OriginalCellSet); VTKM_IS_CELL_SET(OriginalCellSet);
VTKM_IS_ARRAY_HANDLE(PermutationArrayHandleType); VTKM_IS_ARRAY_HANDLE(PermutationArrayHandleType);

@ -50,7 +50,8 @@ class VTKM_ALWAYS_EXPORT CellSetSingleType
typedef vtkm::cont::CellSetSingleType<ConnectivityStorageTag> Thisclass; typedef vtkm::cont::CellSetSingleType<ConnectivityStorageTag> Thisclass;
typedef vtkm::cont::CellSetExplicit< typedef vtkm::cont::CellSetExplicit<
typename vtkm::cont::ArrayHandleConstant<vtkm::UInt8>::StorageTag, typename vtkm::cont::ArrayHandleConstant<vtkm::UInt8>::StorageTag,
typename vtkm::cont::ArrayHandleConstant<vtkm::IdComponent>::StorageTag, ConnectivityStorageTag, typename vtkm::cont::ArrayHandleConstant<vtkm::IdComponent>::StorageTag,
ConnectivityStorageTag,
typename vtkm::cont::ArrayHandleCounting<vtkm::Id>::StorageTag> typename vtkm::cont::ArrayHandleCounting<vtkm::Id>::StorageTag>
Superclass; Superclass;
@ -178,7 +179,9 @@ public:
//This is the way you can fill the memory from another system without copying //This is the way you can fill the memory from another system without copying
VTKM_CONT VTKM_CONT
void Fill(vtkm::Id numPoints, vtkm::UInt8 shapeId, vtkm::IdComponent numberOfPointsPerCell, void Fill(vtkm::Id numPoints,
vtkm::UInt8 shapeId,
vtkm::IdComponent numberOfPointsPerCell,
const vtkm::cont::ArrayHandle<vtkm::Id, ConnectivityStorageTag>& connectivity) const vtkm::cont::ArrayHandle<vtkm::Id, ConnectivityStorageTag>& connectivity)
{ {
this->NumberOfPoints = numPoints; this->NumberOfPoints = numPoints;
@ -218,7 +221,8 @@ public:
private: private:
template <typename CellShapeTag> template <typename CellShapeTag>
void CheckNumberOfPointsPerCell(CellShapeTag, vtkm::CellTraitsTagSizeFixed, void CheckNumberOfPointsPerCell(CellShapeTag,
vtkm::CellTraitsTagSizeFixed,
vtkm::IdComponent numVertices) const vtkm::IdComponent numVertices) const
{ {
if (numVertices != vtkm::CellTraits<CellShapeTag>::NUM_POINTS) if (numVertices != vtkm::CellTraits<CellShapeTag>::NUM_POINTS)
@ -228,7 +232,8 @@ private:
} }
template <typename CellShapeTag> template <typename CellShapeTag>
void CheckNumberOfPointsPerCell(CellShapeTag, vtkm::CellTraitsTagSizeVariable, void CheckNumberOfPointsPerCell(CellShapeTag,
vtkm::CellTraitsTagSizeVariable,
vtkm::IdComponent vtkmNotUsed(numVertices)) const vtkm::IdComponent vtkmNotUsed(numVertices)) const
{ {
// Technically, a shape with a variable number of points probably has a // Technically, a shape with a variable number of points probably has a

@ -92,8 +92,8 @@ public:
}; };
template <typename DeviceAdapter, typename FromTopology, typename ToTopology> template <typename DeviceAdapter, typename FromTopology, typename ToTopology>
typename ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType PrepareForInput( typename ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType
DeviceAdapter, FromTopology, ToTopology) const; PrepareForInput(DeviceAdapter, FromTopology, ToTopology) const;
virtual void PrintSummary(std::ostream& out) const; virtual void PrintSummary(std::ostream& out) const;

@ -50,8 +50,8 @@ typename CellSetStructured<DIMENSION>::SchedulingRangeType
template <vtkm::IdComponent DIMENSION> template <vtkm::IdComponent DIMENSION>
template <typename DeviceAdapter, typename FromTopology, typename ToTopology> template <typename DeviceAdapter, typename FromTopology, typename ToTopology>
typename CellSetStructured<DIMENSION>::template ExecutionTypes<DeviceAdapter, FromTopology, typename CellSetStructured<
ToTopology>::ExecObjectType DIMENSION>::template ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType
CellSetStructured<DIMENSION>::PrepareForInput(DeviceAdapter, FromTopology, ToTopology) const CellSetStructured<DIMENSION>::PrepareForInput(DeviceAdapter, FromTopology, ToTopology) const
{ {
typedef typename ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType typedef typename ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType

@ -35,7 +35,8 @@ void CoordinateSystem::PrintSummary(std::ostream& out) const
VTKM_CONT VTKM_CONT
void CoordinateSystem::GetRange(vtkm::Range* range) const void CoordinateSystem::GetRange(vtkm::Range* range) const
{ {
this->Superclass::GetRange(range, VTKM_DEFAULT_COORDINATE_SYSTEM_TYPE_LIST_TAG(), this->Superclass::GetRange(range,
VTKM_DEFAULT_COORDINATE_SYSTEM_TYPE_LIST_TAG(),
VTKM_DEFAULT_COORDINATE_SYSTEM_STORAGE_LIST_TAG()); VTKM_DEFAULT_COORDINATE_SYSTEM_STORAGE_LIST_TAG());
} }

@ -44,13 +44,15 @@ namespace cont
namespace detail namespace detail
{ {
typedef vtkm::cont::ArrayHandleCompositeVectorType< typedef vtkm::cont::ArrayHandleCompositeVectorType<vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::cont::ArrayHandle<vtkm::Float32>, vtkm::cont::ArrayHandle<vtkm::Float32>,
vtkm::cont::ArrayHandle<vtkm::Float32>>::type ArrayHandleCompositeVectorFloat32_3Default; vtkm::cont::ArrayHandle<vtkm::Float32>>::type
ArrayHandleCompositeVectorFloat32_3Default;
typedef vtkm::cont::ArrayHandleCompositeVectorType< typedef vtkm::cont::ArrayHandleCompositeVectorType<vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::cont::ArrayHandle<vtkm::Float64>, vtkm::cont::ArrayHandle<vtkm::Float64>,
vtkm::cont::ArrayHandle<vtkm::Float64>>::type ArrayHandleCompositeVectorFloat64_3Default; vtkm::cont::ArrayHandle<vtkm::Float64>>::type
ArrayHandleCompositeVectorFloat64_3Default;
} // namespace detail } // namespace detail
@ -60,13 +62,14 @@ typedef vtkm::cont::ArrayHandleCompositeVectorType<
/// by default (unless it is defined before including VTK-m headers. /// by default (unless it is defined before including VTK-m headers.
/// ///
struct StorageListTagCoordinateSystemDefault struct StorageListTagCoordinateSystemDefault
: vtkm::ListTagBase< : vtkm::ListTagBase<vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagBasic, vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag, vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag,
detail::ArrayHandleCompositeVectorFloat32_3Default::StorageTag, detail::ArrayHandleCompositeVectorFloat32_3Default::StorageTag,
detail::ArrayHandleCompositeVectorFloat64_3Default::StorageTag, detail::ArrayHandleCompositeVectorFloat64_3Default::StorageTag,
vtkm::cont::ArrayHandleCartesianProduct< vtkm::cont::ArrayHandleCartesianProduct<
vtkm::cont::ArrayHandle<vtkm::FloatDefault>, vtkm::cont::ArrayHandle<vtkm::FloatDefault>, vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>::StorageTag> vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>::StorageTag>
{ {
}; };
@ -113,10 +116,12 @@ public:
/// ///
VTKM_CONT VTKM_CONT
CoordinateSystem( CoordinateSystem(
std::string name, vtkm::Id3 dimensions, std::string name,
vtkm::Id3 dimensions,
vtkm::Vec<vtkm::FloatDefault, 3> origin = vtkm::Vec<vtkm::FloatDefault, 3>(0.0f, 0.0f, 0.0f), vtkm::Vec<vtkm::FloatDefault, 3> origin = vtkm::Vec<vtkm::FloatDefault, 3>(0.0f, 0.0f, 0.0f),
vtkm::Vec<vtkm::FloatDefault, 3> spacing = vtkm::Vec<vtkm::FloatDefault, 3>(1.0f, 1.0f, 1.0f)) vtkm::Vec<vtkm::FloatDefault, 3> spacing = vtkm::Vec<vtkm::FloatDefault, 3>(1.0f, 1.0f, 1.0f))
: Superclass(name, ASSOC_POINTS, : Superclass(name,
ASSOC_POINTS,
vtkm::cont::DynamicArrayHandle( vtkm::cont::DynamicArrayHandle(
vtkm::cont::ArrayHandleUniformPointCoordinates(dimensions, origin, spacing))) vtkm::cont::ArrayHandleUniformPointCoordinates(dimensions, origin, spacing)))
{ {
@ -145,8 +150,8 @@ public:
{ {
VTKM_IS_LIST_TAG(TypeList); VTKM_IS_LIST_TAG(TypeList);
this->Superclass::GetRange(range, TypeList(), this->Superclass::GetRange(
VTKM_DEFAULT_COORDINATE_SYSTEM_STORAGE_LIST_TAG()); range, TypeList(), VTKM_DEFAULT_COORDINATE_SYSTEM_STORAGE_LIST_TAG());
} }
template <typename TypeList, typename StorageList> template <typename TypeList, typename StorageList>

@ -67,8 +67,9 @@ public:
} }
VTKM_CONT VTKM_CONT
vtkm::Id GetFieldIndex(const std::string& name, vtkm::cont::Field::AssociationEnum assoc = vtkm::Id GetFieldIndex(
vtkm::cont::Field::ASSOC_ANY) const const std::string& name,
vtkm::cont::Field::AssociationEnum assoc = vtkm::cont::Field::ASSOC_ANY) const
{ {
bool found; bool found;
vtkm::Id index = this->FindFieldIndex(name, assoc, found); vtkm::Id index = this->FindFieldIndex(name, assoc, found);
@ -234,7 +235,8 @@ private:
std::vector<vtkm::cont::DynamicCellSet> CellSets; std::vector<vtkm::cont::DynamicCellSet> CellSets;
VTKM_CONT VTKM_CONT
vtkm::Id FindFieldIndex(const std::string& name, vtkm::cont::Field::AssociationEnum association, vtkm::Id FindFieldIndex(const std::string& name,
vtkm::cont::Field::AssociationEnum association,
bool& found) const bool& found) const
{ {
for (std::size_t index = 0; index < this->Fields.size(); ++index) for (std::size_t index = 0; index < this->Fields.size(); ++index)

@ -59,8 +59,8 @@ public:
const std::string& cellNm = "cells") const std::string& cellNm = "cells")
{ {
std::vector<T> yVals(xVals.size(), 0), zVals(xVals.size(), 0); std::vector<T> yVals(xVals.size(), 0), zVals(xVals.size(), 0);
return DataSetBuilderExplicit::Create(xVals, yVals, zVals, shapes, numIndices, connectivity, return DataSetBuilderExplicit::Create(
coordsNm, cellNm); xVals, yVals, zVals, shapes, numIndices, connectivity, coordsNm, cellNm);
} }
template <typename T> template <typename T>
@ -73,27 +73,33 @@ public:
const std::string& cellNm = "cells") const std::string& cellNm = "cells")
{ {
std::vector<T> zVals(xVals.size(), 0); std::vector<T> zVals(xVals.size(), 0);
return DataSetBuilderExplicit::Create(xVals, yVals, zVals, shapes, numIndices, connectivity, return DataSetBuilderExplicit::Create(
coordsNm, cellNm); xVals, yVals, zVals, shapes, numIndices, connectivity, coordsNm, cellNm);
} }
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create( VTKM_CONT static vtkm::cont::DataSet Create(const std::vector<T>& xVals,
const std::vector<T>& xVals, const std::vector<T>& yVals, const std::vector<T>& zVals, const std::vector<T>& yVals,
const std::vector<vtkm::UInt8>& shapes, const std::vector<vtkm::IdComponent>& numIndices, const std::vector<T>& zVals,
const std::vector<vtkm::Id>& connectivity, const std::string& coordsNm = "coords", const std::vector<vtkm::UInt8>& shapes,
const std::string& cellNm = "cells"); const std::vector<vtkm::IdComponent>& numIndices,
const std::vector<vtkm::Id>& connectivity,
const std::string& coordsNm = "coords",
const std::string& cellNm = "cells");
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create( VTKM_CONT static vtkm::cont::DataSet Create(
const vtkm::cont::ArrayHandle<T>& xVals, const vtkm::cont::ArrayHandle<T>& yVals, const vtkm::cont::ArrayHandle<T>& xVals,
const vtkm::cont::ArrayHandle<T>& zVals, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<T>& yVals,
const vtkm::cont::ArrayHandle<T>& zVals,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm = "coords", const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm = "coords",
const std::string& cellNm = "cells") const std::string& cellNm = "cells")
{ {
return DataSetBuilderExplicit::BuildDataSet(xVals, yVals, zVals, shapes, numIndices, return DataSetBuilderExplicit::BuildDataSet(
connectivity, coordsNm, cellNm); xVals, yVals, zVals, shapes, numIndices, connectivity, coordsNm, cellNm);
} }
template <typename T> template <typename T>
@ -109,11 +115,12 @@ public:
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm = "coords", const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm = "coords",
const std::string& cellNm = "cells") const std::string& cellNm = "cells")
{ {
return DataSetBuilderExplicit::BuildDataSet(coords, shapes, numIndices, connectivity, coordsNm, return DataSetBuilderExplicit::BuildDataSet(
cellNm); coords, shapes, numIndices, connectivity, coordsNm, cellNm);
} }
template <typename T, typename CellShapeTag> template <typename T, typename CellShapeTag>
@ -126,21 +133,27 @@ public:
template <typename T, typename CellShapeTag> template <typename T, typename CellShapeTag>
VTKM_CONT static vtkm::cont::DataSet Create( VTKM_CONT static vtkm::cont::DataSet Create(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, CellShapeTag tag, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
vtkm::IdComponent numberOfPointsPerCell, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, CellShapeTag tag,
const std::string& coordsNm = "coords", const std::string& cellNm = "cells") vtkm::IdComponent numberOfPointsPerCell,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm = "coords",
const std::string& cellNm = "cells")
{ {
return DataSetBuilderExplicit::BuildDataSet(coords, tag, numberOfPointsPerCell, connectivity, return DataSetBuilderExplicit::BuildDataSet(
coordsNm, cellNm); coords, tag, numberOfPointsPerCell, connectivity, coordsNm, cellNm);
} }
private: private:
template <typename T> template <typename T>
static vtkm::cont::DataSet BuildDataSet( static vtkm::cont::DataSet BuildDataSet(
const vtkm::cont::ArrayHandle<T>& X, const vtkm::cont::ArrayHandle<T>& Y, const vtkm::cont::ArrayHandle<T>& X,
const vtkm::cont::ArrayHandle<T>& Z, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<T>& Y,
const vtkm::cont::ArrayHandle<T>& Z,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm); const std::string& cellNm);
template <typename T> template <typename T>
@ -148,21 +161,30 @@ private:
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm); const std::string& cellNm);
template <typename T, typename CellShapeTag> template <typename T, typename CellShapeTag>
VTKM_CONT static vtkm::cont::DataSet BuildDataSet( VTKM_CONT static vtkm::cont::DataSet BuildDataSet(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, CellShapeTag tag, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
vtkm::IdComponent numberOfPointsPerCell, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, CellShapeTag tag,
const std::string& coordsNm, const std::string& cellNm); vtkm::IdComponent numberOfPointsPerCell,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm);
}; };
template <typename T> template <typename T>
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create( inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create(
const std::vector<T>& xVals, const std::vector<T>& yVals, const std::vector<T>& zVals, const std::vector<T>& xVals,
const std::vector<vtkm::UInt8>& shapes, const std::vector<vtkm::IdComponent>& numIndices, const std::vector<T>& yVals,
const std::vector<vtkm::Id>& connectivity, const std::string& coordsNm, const std::string& cellNm) const std::vector<T>& zVals,
const std::vector<vtkm::UInt8>& shapes,
const std::vector<vtkm::IdComponent>& numIndices,
const std::vector<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm)
{ {
VTKM_ASSERT(xVals.size() == yVals.size() && yVals.size() == zVals.size() && xVals.size() > 0); VTKM_ASSERT(xVals.size() == yVals.size() && yVals.size() == zVals.size() && xVals.size() > 0);
@ -183,10 +205,13 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create(
template <typename T> template <typename T>
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet( inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
const vtkm::cont::ArrayHandle<T>& X, const vtkm::cont::ArrayHandle<T>& Y, const vtkm::cont::ArrayHandle<T>& X,
const vtkm::cont::ArrayHandle<T>& Z, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<T>& Y,
const vtkm::cont::ArrayHandle<T>& Z,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm) const std::string& cellNm)
{ {
VTKM_ASSERT(X.GetNumberOfValues() == Y.GetNumberOfValues() && VTKM_ASSERT(X.GetNumberOfValues() == Y.GetNumberOfValues() &&
@ -207,9 +232,12 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
template <typename T> template <typename T>
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create( inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create(
const std::vector<vtkm::Vec<T, 3>>& coords, const std::vector<vtkm::UInt8>& shapes, const std::vector<vtkm::Vec<T, 3>>& coords,
const std::vector<vtkm::IdComponent>& numIndices, const std::vector<vtkm::Id>& connectivity, const std::vector<vtkm::UInt8>& shapes,
const std::string& coordsNm, const std::string& cellNm) const std::vector<vtkm::IdComponent>& numIndices,
const std::vector<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm)
{ {
vtkm::cont::ArrayHandle<Vec<T, 3>> coordsArray; vtkm::cont::ArrayHandle<Vec<T, 3>> coordsArray;
DataSetBuilderExplicit::CopyInto(coords, coordsArray); DataSetBuilderExplicit::CopyInto(coords, coordsArray);
@ -229,7 +257,8 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes, const vtkm::cont::ArrayHandle<vtkm::UInt8>& shapes,
const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices, const vtkm::cont::ArrayHandle<vtkm::IdComponent>& numIndices,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, const std::string& coordsNm, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm) const std::string& cellNm)
{ {
vtkm::cont::DataSet dataSet; vtkm::cont::DataSet dataSet;
@ -246,9 +275,12 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
template <typename T, typename CellShapeTag> template <typename T, typename CellShapeTag>
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create( inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create(
const std::vector<vtkm::Vec<T, 3>>& coords, CellShapeTag tag, const std::vector<vtkm::Vec<T, 3>>& coords,
vtkm::IdComponent numberOfPointsPerCell, const std::vector<vtkm::Id>& connectivity, CellShapeTag tag,
const std::string& coordsNm, const std::string& cellNm) vtkm::IdComponent numberOfPointsPerCell,
const std::vector<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm)
{ {
vtkm::cont::ArrayHandle<Vec<T, 3>> coordsArray; vtkm::cont::ArrayHandle<Vec<T, 3>> coordsArray;
DataSetBuilderExplicit::CopyInto(coords, coordsArray); DataSetBuilderExplicit::CopyInto(coords, coordsArray);
@ -256,15 +288,18 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::Create(
vtkm::cont::ArrayHandle<vtkm::Id> Cc; vtkm::cont::ArrayHandle<vtkm::Id> Cc;
DataSetBuilderExplicit::CopyInto(connectivity, Cc); DataSetBuilderExplicit::CopyInto(connectivity, Cc);
return DataSetBuilderExplicit::Create(coordsArray, tag, numberOfPointsPerCell, Cc, coordsNm, return DataSetBuilderExplicit::Create(
cellNm); coordsArray, tag, numberOfPointsPerCell, Cc, coordsNm, cellNm);
} }
template <typename T, typename CellShapeTag> template <typename T, typename CellShapeTag>
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet( inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, CellShapeTag tag, const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
vtkm::IdComponent numberOfPointsPerCell, const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity, CellShapeTag tag,
const std::string& coordsNm, const std::string& cellNm) vtkm::IdComponent numberOfPointsPerCell,
const vtkm::cont::ArrayHandle<vtkm::Id>& connectivity,
const std::string& coordsNm,
const std::string& cellNm)
{ {
vtkm::cont::DataSet dataSet; vtkm::cont::DataSet dataSet;
@ -317,8 +352,8 @@ public:
template <typename T> template <typename T>
VTKM_CONT vtkm::Id AddPoint(const T& x, const T& y, const T& z = 0) VTKM_CONT vtkm::Id AddPoint(const T& x, const T& y, const T& z = 0)
{ {
return AddPoint(static_cast<vtkm::Float32>(x), static_cast<vtkm::Float32>(y), return AddPoint(
static_cast<vtkm::Float32>(z)); static_cast<vtkm::Float32>(x), static_cast<vtkm::Float32>(y), static_cast<vtkm::Float32>(z));
} }
template <typename T> template <typename T>

@ -68,12 +68,14 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx, T* xvals, std::string coordNm = "coords", VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx,
T* xvals,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
T yvals = 0, zvals = 0; T yvals = 0, zvals = 0;
return DataSetBuilderRectilinear::BuildDataSet(nx, 1, 1, xvals, &yvals, &zvals, coordNm, return DataSetBuilderRectilinear::BuildDataSet(
cellNm); nx, 1, 1, xvals, &yvals, &zvals, coordNm, cellNm);
} }
template <typename T> template <typename T>
@ -101,13 +103,16 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx, vtkm::Id ny, T* xvals, T* yvals, VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx,
vtkm::Id ny,
T* xvals,
T* yvals,
std::string coordNm = "coords", std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
T zvals = 0; T zvals = 0;
return DataSetBuilderRectilinear::BuildDataSet(nx, ny, 1, xvals, yvals, &zvals, coordNm, return DataSetBuilderRectilinear::BuildDataSet(
cellNm); nx, ny, 1, xvals, yvals, &zvals, coordNm, cellNm);
} }
template <typename T> template <typename T>
@ -124,12 +129,17 @@ public:
//3D grids. //3D grids.
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx, vtkm::Id ny, vtkm::Id nz, T* xvals, VTKM_CONT static vtkm::cont::DataSet Create(vtkm::Id nx,
T* yvals, T* zvals, std::string coordNm = "coords", vtkm::Id ny,
vtkm::Id nz,
T* xvals,
T* yvals,
T* zvals,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return DataSetBuilderRectilinear::BuildDataSet(nx, ny, nz, xvals, yvals, zvals, coordNm, return DataSetBuilderRectilinear::BuildDataSet(
cellNm); nx, ny, nz, xvals, yvals, zvals, coordNm, cellNm);
} }
template <typename T> template <typename T>
@ -157,7 +167,8 @@ private:
VTKM_CONT static vtkm::cont::DataSet BuildDataSet(const std::vector<T>& xvals, VTKM_CONT static vtkm::cont::DataSet BuildDataSet(const std::vector<T>& xvals,
const std::vector<T>& yvals, const std::vector<T>& yvals,
const std::vector<T>& zvals, const std::vector<T>& zvals,
std::string coordNm, std::string cellNm) std::string coordNm,
std::string cellNm)
{ {
vtkm::cont::ArrayHandle<vtkm::FloatDefault> Xc, Yc, Zc; vtkm::cont::ArrayHandle<vtkm::FloatDefault> Xc, Yc, Zc;
DataSetBuilderRectilinear::CopyInto(xvals, Xc); DataSetBuilderRectilinear::CopyInto(xvals, Xc);
@ -168,9 +179,14 @@ private:
} }
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet BuildDataSet(vtkm::Id nx, vtkm::Id ny, vtkm::Id nz, VTKM_CONT static vtkm::cont::DataSet BuildDataSet(vtkm::Id nx,
const T* xvals, const T* yvals, const T* zvals, vtkm::Id ny,
std::string coordNm, std::string cellNm) vtkm::Id nz,
const T* xvals,
const T* yvals,
const T* zvals,
std::string coordNm,
std::string cellNm)
{ {
vtkm::cont::ArrayHandle<vtkm::FloatDefault> Xc, Yc, Zc; vtkm::cont::ArrayHandle<vtkm::FloatDefault> Xc, Yc, Zc;
DataSetBuilderRectilinear::CopyInto(xvals, nx, Xc); DataSetBuilderRectilinear::CopyInto(xvals, nx, Xc);
@ -184,7 +200,8 @@ private:
VTKM_CONT static vtkm::cont::DataSet BuildDataSet(const vtkm::cont::ArrayHandle<T>& X, VTKM_CONT static vtkm::cont::DataSet BuildDataSet(const vtkm::cont::ArrayHandle<T>& X,
const vtkm::cont::ArrayHandle<T>& Y, const vtkm::cont::ArrayHandle<T>& Y,
const vtkm::cont::ArrayHandle<T>& Z, const vtkm::cont::ArrayHandle<T>& Z,
std::string coordNm, std::string cellNm) std::string coordNm,
std::string cellNm)
{ {
vtkm::cont::DataSet dataSet; vtkm::cont::DataSet dataSet;

@ -38,17 +38,23 @@ public:
//1D uniform grid //1D uniform grid
template <typename T> template <typename T>
VTKM_CONT static vtkm::cont::DataSet Create(const vtkm::Id& dimension, const T& origin, VTKM_CONT static vtkm::cont::DataSet Create(const vtkm::Id& dimension,
const T& spacing, std::string coordNm = "coords", const T& origin,
const T& spacing,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return DataSetBuilderUniform::CreateDataSet( return DataSetBuilderUniform::CreateDataSet(
vtkm::Id3(dimension, 1, 1), VecType(static_cast<vtkm::FloatDefault>(origin), 0, 0), vtkm::Id3(dimension, 1, 1),
VecType(static_cast<vtkm::FloatDefault>(spacing), 1, 1), coordNm, cellNm); VecType(static_cast<vtkm::FloatDefault>(origin), 0, 0),
VecType(static_cast<vtkm::FloatDefault>(spacing), 1, 1),
coordNm,
cellNm);
} }
VTKM_CONT VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id& dimension, std::string coordNm = "coords", static vtkm::cont::DataSet Create(const vtkm::Id& dimension,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return CreateDataSet(vtkm::Id3(dimension, 1, 1), VecType(0), VecType(1), coordNm, cellNm); return CreateDataSet(vtkm::Id3(dimension, 1, 1), VecType(0), VecType(1), coordNm, cellNm);
@ -62,21 +68,24 @@ public:
std::string coordNm = "coords", std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return DataSetBuilderUniform::CreateDataSet( return DataSetBuilderUniform::CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], 1),
vtkm::Id3(dimensions[0], dimensions[1], 1), VecType(static_cast<vtkm::FloatDefault>(origin[0]),
VecType(static_cast<vtkm::FloatDefault>(origin[0]), static_cast<vtkm::FloatDefault>(origin[1]),
static_cast<vtkm::FloatDefault>(origin[1]), 0), 0),
VecType(static_cast<vtkm::FloatDefault>(spacing[0]), VecType(static_cast<vtkm::FloatDefault>(spacing[0]),
static_cast<vtkm::FloatDefault>(spacing[1]), 1), static_cast<vtkm::FloatDefault>(spacing[1]),
coordNm, cellNm); 1),
coordNm,
cellNm);
} }
VTKM_CONT VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id2& dimensions, std::string coordNm = "coords", static vtkm::cont::DataSet Create(const vtkm::Id2& dimensions,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], 1), VecType(0), VecType(1), return CreateDataSet(
coordNm, cellNm); vtkm::Id3(dimensions[0], dimensions[1], 1), VecType(0), VecType(1), coordNm, cellNm);
} }
//3D uniform grids. //3D uniform grids.
@ -95,15 +104,20 @@ public:
VecType(static_cast<vtkm::FloatDefault>(spacing[0]), VecType(static_cast<vtkm::FloatDefault>(spacing[0]),
static_cast<vtkm::FloatDefault>(spacing[1]), static_cast<vtkm::FloatDefault>(spacing[1]),
static_cast<vtkm::FloatDefault>(spacing[2])), static_cast<vtkm::FloatDefault>(spacing[2])),
coordNm, cellNm); coordNm,
cellNm);
} }
VTKM_CONT VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id3& dimensions, std::string coordNm = "coords", static vtkm::cont::DataSet Create(const vtkm::Id3& dimensions,
std::string coordNm = "coords",
std::string cellNm = "cells") std::string cellNm = "cells")
{ {
return CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], dimensions[2]), VecType(0), return CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], dimensions[2]),
VecType(1), coordNm, cellNm); VecType(0),
VecType(1),
coordNm,
cellNm);
} }
private: private:
@ -111,7 +125,8 @@ private:
static vtkm::cont::DataSet CreateDataSet(const vtkm::Id3& dimensions, static vtkm::cont::DataSet CreateDataSet(const vtkm::Id3& dimensions,
const vtkm::Vec<vtkm::FloatDefault, 3>& origin, const vtkm::Vec<vtkm::FloatDefault, 3>& origin,
const vtkm::Vec<vtkm::FloatDefault, 3>& spacing, const vtkm::Vec<vtkm::FloatDefault, 3>& spacing,
std::string coordNm, std::string cellNm) std::string coordNm,
std::string cellNm)
{ {
vtkm::Id dims[3]; vtkm::Id dims[3];
int ndims = 0; int ndims = 0;

@ -37,36 +37,42 @@ public:
//Point centered fields. //Point centered fields.
VTKM_CONT VTKM_CONT
static void AddPointField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, static void AddPointField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const vtkm::cont::DynamicArrayHandle& field) const vtkm::cont::DynamicArrayHandle& field)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field));
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const vtkm::cont::ArrayHandle<T, Storage>& field) const vtkm::cont::ArrayHandle<T, Storage>& field)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field));
} }
template <typename T> template <typename T>
VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const std::vector<T>& field) const std::vector<T>& field)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field));
} }
template <typename T> template <typename T>
VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddPointField(vtkm::cont::DataSet& dataSet,
const T* field, const vtkm::Id& n) const std::string& fieldName,
const T* field,
const vtkm::Id& n)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field, n)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_POINTS, field, n));
} }
//Cell centered field //Cell centered field
VTKM_CONT VTKM_CONT
static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, static void AddCellField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const vtkm::cont::DynamicArrayHandle& field, const vtkm::cont::DynamicArrayHandle& field,
const std::string& cellSetName) const std::string& cellSetName)
{ {
@ -74,7 +80,8 @@ public:
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const vtkm::cont::ArrayHandle<T, Storage>& field, const vtkm::cont::ArrayHandle<T, Storage>& field,
const std::string& cellSetName) const std::string& cellSetName)
{ {
@ -82,29 +89,36 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const std::vector<T>& field, const std::string& cellSetName) const std::string& fieldName,
const std::vector<T>& field,
const std::string& cellSetName)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_CELL_SET, cellSetName, field)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_CELL_SET, cellSetName, field));
} }
template <typename T> template <typename T>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const T* field, const vtkm::Id& n, const std::string& fieldName,
const T* field,
const vtkm::Id& n,
const std::string& cellSetName) const std::string& cellSetName)
{ {
dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_CELL_SET, cellSetName, field, n)); dataSet.AddField(Field(fieldName, vtkm::cont::Field::ASSOC_CELL_SET, cellSetName, field, n));
} }
VTKM_CONT VTKM_CONT
static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, static void AddCellField(vtkm::cont::DataSet& dataSet,
const vtkm::cont::DynamicArrayHandle& field, vtkm::Id cellSetIndex = 0) const std::string& fieldName,
const vtkm::cont::DynamicArrayHandle& field,
vtkm::Id cellSetIndex = 0)
{ {
std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName(); std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName();
DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName); DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName);
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const std::string& fieldName,
const vtkm::cont::ArrayHandle<T, Storage>& field, const vtkm::cont::ArrayHandle<T, Storage>& field,
vtkm::Id cellSetIndex = 0) vtkm::Id cellSetIndex = 0)
{ {
@ -112,16 +126,21 @@ public:
DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName); DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName);
} }
template <typename T> template <typename T>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const std::vector<T>& field, vtkm::Id cellSetIndex = 0) const std::string& fieldName,
const std::vector<T>& field,
vtkm::Id cellSetIndex = 0)
{ {
std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName(); std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName();
DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName); DataSetFieldAdd::AddCellField(dataSet, fieldName, field, cellSetName);
} }
template <typename T> template <typename T>
VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet, const std::string& fieldName, VTKM_CONT static void AddCellField(vtkm::cont::DataSet& dataSet,
const T* field, const vtkm::Id& n, vtkm::Id cellSetIndex = 0) const std::string& fieldName,
const T* field,
const vtkm::Id& n,
vtkm::Id cellSetIndex = 0)
{ {
std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName(); std::string cellSetName = dataSet.GetCellSet(cellSetIndex).GetName();
DataSetFieldAdd::AddCellField(dataSet, fieldName, field, n, cellSetName); DataSetFieldAdd::AddCellField(dataSet, fieldName, field, n, cellSetName);

@ -112,7 +112,8 @@ struct DeviceAdapterAlgorithm
/// ///
template <typename T, typename U, class CIn, class COut> template <typename T, typename U, class CIn, class COut>
VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::Id inputStartIndex, vtkm::Id numberOfElementsToCopy, vtkm::Id inputStartIndex,
vtkm::Id numberOfElementsToCopy,
vtkm::cont::ArrayHandle<U, COut>& output, vtkm::cont::ArrayHandle<U, COut>& output,
vtkm::Id outputIndex = 0); vtkm::Id outputIndex = 0);
@ -180,7 +181,8 @@ struct DeviceAdapterAlgorithm
/// ///
/// \return The total sum. /// \return The total sum.
template <typename T, typename U, class CIn, class BinaryFunctor> template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue, VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor); BinaryFunctor binary_functor);
/// \brief Compute a accumulated sum operation on the input key value pairs /// \brief Compute a accumulated sum operation on the input key value pairs
@ -191,7 +193,12 @@ struct DeviceAdapterAlgorithm
/// values inside that range. Once finished a single key and value is created /// values inside that range. Once finished a single key and value is created
/// for each segment. /// for each segment.
/// ///
template <typename T, typename U, class CKeyIn, class CValIn, class CKeyOut, class CValOut, template <typename T,
typename U,
class CKeyIn,
class CValIn,
class CKeyOut,
class CValOut,
class BinaryFunctor> class BinaryFunctor>
VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, CKeyIn>& keys, VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, CKeyIn>& keys,
const vtkm::cont::ArrayHandle<U, CValIn>& values, const vtkm::cont::ArrayHandle<U, CValIn>& values,
@ -251,7 +258,11 @@ struct DeviceAdapterAlgorithm
/// applied to all values inside that range. Once finished the result is /// applied to all values inside that range. Once finished the result is
/// stored in \c values_output ArrayHandle. /// stored in \c values_output ArrayHandle.
/// ///
template <typename T, typename U, typename KIn, typename VIn, typename VOut, template <typename T,
typename U,
typename KIn,
typename VIn,
typename VOut,
typename BinaryFunctor> typename BinaryFunctor>
VTKM_CONT static void ScanInclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ScanInclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -310,7 +321,8 @@ struct DeviceAdapterAlgorithm
VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output, vtkm::cont::ArrayHandle<U, VOut>& output,
const U& initialValue, BinaryFunctor binaryFunctor); const U& initialValue,
BinaryFunctor binaryFunctor);
/// \brief Compute a segmented exclusive prefix sum operation on the input key value pairs. /// \brief Compute a segmented exclusive prefix sum operation on the input key value pairs.
/// ///

@ -35,9 +35,9 @@ namespace vtkm
namespace cont namespace cont
{ {
struct DeviceAdapterListTagCommon struct DeviceAdapterListTagCommon : vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda,
: vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda, vtkm::cont::DeviceAdapterTagTBB, vtkm::cont::DeviceAdapterTagTBB,
vtkm::cont::DeviceAdapterTagSerial> vtkm::cont::DeviceAdapterTagSerial>
{ {
}; };
} }

@ -342,7 +342,8 @@ public:
/// ///
template <typename NewTypeList, typename NewStorageList> template <typename NewTypeList, typename NewStorageList>
VTKM_CONT DynamicArrayHandleBase<NewTypeList, NewStorageList> ResetTypeAndStorageLists( VTKM_CONT DynamicArrayHandleBase<NewTypeList, NewStorageList> ResetTypeAndStorageLists(
NewTypeList = NewTypeList(), NewStorageList = NewStorageList()) const NewTypeList = NewTypeList(),
NewStorageList = NewStorageList()) const
{ {
VTKM_IS_LIST_TAG(NewTypeList); VTKM_IS_LIST_TAG(NewTypeList);
VTKM_IS_LIST_TAG(NewStorageList); VTKM_IS_LIST_TAG(NewStorageList);
@ -519,8 +520,9 @@ VTKM_CONT void DynamicArrayHandleBase<TypeList, StorageList>::CastAndCall(const
template <> template <>
template <typename Functor> template <typename Functor>
VTKM_CONT void DynamicArrayHandleBase< VTKM_CONT void
VTKM_DEFAULT_TYPE_LIST_TAG, VTKM_DEFAULT_STORAGE_LIST_TAG>::CastAndCall(const Functor& f) const DynamicArrayHandleBase<VTKM_DEFAULT_TYPE_LIST_TAG, VTKM_DEFAULT_STORAGE_LIST_TAG>::CastAndCall(
const Functor& f) const
{ {
typedef detail::DynamicArrayHandleTryType<Functor, VTKM_DEFAULT_STORAGE_LIST_TAG> TryTypeType; typedef detail::DynamicArrayHandleTryType<Functor, VTKM_DEFAULT_STORAGE_LIST_TAG> TryTypeType;

@ -88,7 +88,8 @@ public:
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT Field(std::string name, AssociationEnum association, VTKM_CONT Field(std::string name,
AssociationEnum association,
const ArrayHandle<T, Storage>& data) const ArrayHandle<T, Storage>& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -129,7 +130,9 @@ public:
/// constructors for cell set associations /// constructors for cell set associations
VTKM_CONT VTKM_CONT
Field(std::string name, AssociationEnum association, const std::string& cellSetName, Field(std::string name,
AssociationEnum association,
const std::string& cellSetName,
const vtkm::cont::DynamicArrayHandle& data) const vtkm::cont::DynamicArrayHandle& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -143,7 +146,9 @@ public:
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT Field(std::string name, AssociationEnum association, const std::string& cellSetName, VTKM_CONT Field(std::string name,
AssociationEnum association,
const std::string& cellSetName,
const vtkm::cont::ArrayHandle<T, Storage>& data) const vtkm::cont::ArrayHandle<T, Storage>& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -157,7 +162,9 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT Field(std::string name, AssociationEnum association, const std::string& cellSetName, VTKM_CONT Field(std::string name,
AssociationEnum association,
const std::string& cellSetName,
const std::vector<T>& data) const std::vector<T>& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -171,8 +178,11 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT Field(std::string name, AssociationEnum association, const std::string& cellSetName, VTKM_CONT Field(std::string name,
const T* data, vtkm::Id nvals) AssociationEnum association,
const std::string& cellSetName,
const T* data,
vtkm::Id nvals)
: Name(name) : Name(name)
, Association(association) , Association(association)
, AssocCellSetName(cellSetName) , AssocCellSetName(cellSetName)
@ -186,7 +196,9 @@ public:
/// constructors for logical dimension associations /// constructors for logical dimension associations
VTKM_CONT VTKM_CONT
Field(std::string name, AssociationEnum association, vtkm::IdComponent logicalDim, Field(std::string name,
AssociationEnum association,
vtkm::IdComponent logicalDim,
const vtkm::cont::DynamicArrayHandle& data) const vtkm::cont::DynamicArrayHandle& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -200,7 +212,9 @@ public:
} }
template <typename T, typename Storage> template <typename T, typename Storage>
VTKM_CONT Field(std::string name, AssociationEnum association, vtkm::IdComponent logicalDim, VTKM_CONT Field(std::string name,
AssociationEnum association,
vtkm::IdComponent logicalDim,
const vtkm::cont::ArrayHandle<T, Storage>& data) const vtkm::cont::ArrayHandle<T, Storage>& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -213,7 +227,9 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT Field(std::string name, AssociationEnum association, vtkm::IdComponent logicalDim, VTKM_CONT Field(std::string name,
AssociationEnum association,
vtkm::IdComponent logicalDim,
const std::vector<T>& data) const std::vector<T>& data)
: Name(name) : Name(name)
, Association(association) , Association(association)
@ -226,8 +242,11 @@ public:
} }
template <typename T> template <typename T>
VTKM_CONT Field(std::string name, AssociationEnum association, vtkm::IdComponent logicalDim, VTKM_CONT Field(std::string name,
const T* data, vtkm::Id nvals) AssociationEnum association,
vtkm::IdComponent logicalDim,
const T* data,
vtkm::Id nvals)
: Name(name) : Name(name)
, Association(association) , Association(association)
, AssocLogicalDim(logicalDim) , AssocLogicalDim(logicalDim)
@ -342,7 +361,8 @@ public:
tmp.Allocate(nvals); tmp.Allocate(nvals);
//copy into the memory owned by the array handle //copy into the memory owned by the array handle
std::copy(ptr, ptr + static_cast<std::size_t>(nvals), std::copy(ptr,
ptr + static_cast<std::size_t>(nvals),
vtkm::cont::ArrayPortalToIteratorBegin(tmp.GetPortalControl())); vtkm::cont::ArrayPortalToIteratorBegin(tmp.GetPortalControl()));
//assign to the dynamic array handle //assign to the dynamic array handle

@ -121,7 +121,11 @@ class VTKM_ALWAYS_EXPORT Box : public ImplicitFunctionImpl<Box>
public: public:
Box(); Box();
Box(vtkm::Vec<FloatDefault, 3> minPoint, vtkm::Vec<FloatDefault, 3> maxPoint); Box(vtkm::Vec<FloatDefault, 3> minPoint, vtkm::Vec<FloatDefault, 3> maxPoint);
Box(FloatDefault xmin, FloatDefault xmax, FloatDefault ymin, FloatDefault ymax, FloatDefault zmin, Box(FloatDefault xmin,
FloatDefault xmax,
FloatDefault ymin,
FloatDefault ymax,
FloatDefault zmin,
FloatDefault zmax); FloatDefault zmax);
void SetMinPoint(const vtkm::Vec<FloatDefault, 3>& point); void SetMinPoint(const vtkm::Vec<FloatDefault, 3>& point);
@ -152,7 +156,8 @@ class VTKM_ALWAYS_EXPORT Cylinder : public ImplicitFunctionImpl<Cylinder>
public: public:
Cylinder(); Cylinder();
Cylinder(const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault radius); Cylinder(const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault radius);
Cylinder(const vtkm::Vec<FloatDefault, 3>& center, const vtkm::Vec<FloatDefault, 3>& axis, Cylinder(const vtkm::Vec<FloatDefault, 3>& center,
const vtkm::Vec<FloatDefault, 3>& axis,
FloatDefault radius); FloatDefault radius);
void SetCenter(const vtkm::Vec<FloatDefault, 3>& center); void SetCenter(const vtkm::Vec<FloatDefault, 3>& center);

@ -41,8 +41,12 @@ inline Box::Box(vtkm::Vec<FloatDefault, 3> minPoint, vtkm::Vec<FloatDefault, 3>
{ {
} }
inline Box::Box(FloatDefault xmin, FloatDefault xmax, FloatDefault ymin, FloatDefault ymax, inline Box::Box(FloatDefault xmin,
FloatDefault zmin, FloatDefault zmax) FloatDefault xmax,
FloatDefault ymin,
FloatDefault ymax,
FloatDefault zmin,
FloatDefault zmax)
{ {
MinPoint[0] = xmin; MinPoint[0] = xmin;
MaxPoint[0] = xmax; MaxPoint[0] = xmax;
@ -81,7 +85,8 @@ inline FloatDefault Box::Value(FloatDefault x, FloatDefault y, FloatDefault z) c
} }
VTKM_EXEC_CONT VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Box::Gradient(FloatDefault x, FloatDefault y, inline vtkm::Vec<FloatDefault, 3> Box::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const FloatDefault z) const
{ {
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z)); return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));
@ -295,7 +300,8 @@ inline Cylinder::Cylinder(const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault r
} }
inline Cylinder::Cylinder(const vtkm::Vec<FloatDefault, 3>& center, inline Cylinder::Cylinder(const vtkm::Vec<FloatDefault, 3>& center,
const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault radius) const vtkm::Vec<FloatDefault, 3>& axis,
FloatDefault radius)
: Center(center) : Center(center)
, Axis(vtkm::Normal(axis)) , Axis(vtkm::Normal(axis))
, Radius(radius) , Radius(radius)
@ -359,7 +365,8 @@ inline vtkm::Vec<FloatDefault, 3> Cylinder::Gradient(const vtkm::Vec<FloatDefaul
} }
VTKM_EXEC_CONT VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Cylinder::Gradient(FloatDefault x, FloatDefault y, inline vtkm::Vec<FloatDefault, 3> Cylinder::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const FloatDefault z) const
{ {
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z)); return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));
@ -409,7 +416,8 @@ inline const vtkm::Vec<FloatDefault, 3>* Frustum::GetNormals() const
return this->Normals; return this->Normals;
} }
inline void Frustum::SetPlane(int idx, vtkm::Vec<FloatDefault, 3>& point, inline void Frustum::SetPlane(int idx,
vtkm::Vec<FloatDefault, 3>& point,
vtkm::Vec<FloatDefault, 3>& normal) vtkm::Vec<FloatDefault, 3>& normal)
{ {
if (idx < 0 || idx >= 6) if (idx < 0 || idx >= 6)
@ -466,7 +474,8 @@ inline FloatDefault Frustum::Value(const vtkm::Vec<FloatDefault, 3>& x) const
} }
VTKM_EXEC_CONT VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Frustum::Gradient(FloatDefault x, FloatDefault y, inline vtkm::Vec<FloatDefault, 3> Frustum::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const FloatDefault z) const
{ {
FloatDefault maxVal = -std::numeric_limits<FloatDefault>::max(); FloatDefault maxVal = -std::numeric_limits<FloatDefault>::max();
@ -615,7 +624,8 @@ inline FloatDefault Sphere::Value(const vtkm::Vec<FloatDefault, 3>& x) const
} }
VTKM_EXEC_CONT VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Sphere::Gradient(FloatDefault x, FloatDefault y, inline vtkm::Vec<FloatDefault, 3> Sphere::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const FloatDefault z) const
{ {
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z)); return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));

@ -129,8 +129,8 @@ vtkm::cont::RuntimeDeviceTracker RuntimeDeviceTracker::DeepCopy() const
VTKM_CONT VTKM_CONT
void RuntimeDeviceTracker::DeepCopy(const vtkm::cont::RuntimeDeviceTracker& src) void RuntimeDeviceTracker::DeepCopy(const vtkm::cont::RuntimeDeviceTracker& src)
{ {
std::copy_n(src.Internals->RuntimeValid, VTKM_MAX_DEVICE_ADAPTER_ID, std::copy_n(
this->Internals->RuntimeValid); src.Internals->RuntimeValid, VTKM_MAX_DEVICE_ADAPTER_ID, this->Internals->RuntimeValid);
} }
VTKM_CONT VTKM_CONT

@ -186,12 +186,14 @@ private:
VTKM_CONT_EXPORT VTKM_CONT_EXPORT
VTKM_CONT VTKM_CONT
void SetDeviceState(vtkm::cont::DeviceAdapterId deviceId, void SetDeviceState(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName, bool state); const vtkm::cont::DeviceAdapterNameType& deviceName,
bool state);
VTKM_CONT_EXPORT VTKM_CONT_EXPORT
VTKM_CONT VTKM_CONT
void ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId, void ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName, bool runtimeExists); const vtkm::cont::DeviceAdapterNameType& deviceName,
bool runtimeExists);
}; };
/// \brief Get the global \c RuntimeDeviceTracker. /// \brief Get the global \c RuntimeDeviceTracker.

@ -151,7 +151,8 @@ public:
PortalType portal = this->Storage->GetPortalConst(); PortalType portal = this->Storage->GetPortalConst();
std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal), std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal),
vtkm::cont::ArrayPortalToIteratorEnd(portal), dest); vtkm::cont::ArrayPortalToIteratorEnd(portal),
dest);
} }
VTKM_CONT VTKM_CONT

@ -109,8 +109,9 @@ struct TryExecuteImpl
bool Success; bool Success;
VTKM_CONT VTKM_CONT
TryExecuteImpl(FunctorType& functor, vtkm::cont::RuntimeDeviceTracker tracker = TryExecuteImpl(
vtkm::cont::GetGlobalRuntimeDeviceTracker()) FunctorType& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
: Functor(functor) : Functor(functor)
, Tracker(tracker) , Tracker(tracker)
, Success(false) , Success(false)
@ -159,7 +160,8 @@ private:
/// is used. /// is used.
/// ///
template <typename Functor, typename DeviceList> template <typename Functor, typename DeviceList>
VTKM_CONT bool TryExecute(const Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker, VTKM_CONT bool TryExecute(const Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker,
DeviceList) DeviceList)
{ {
detail::TryExecuteImpl<const Functor> internals(functor, tracker); detail::TryExecuteImpl<const Functor> internals(functor, tracker);
@ -184,14 +186,16 @@ VTKM_CONT bool TryExecute(Functor& functor, DeviceList)
return vtkm::cont::TryExecute(functor, vtkm::cont::GetGlobalRuntimeDeviceTracker(), DeviceList()); return vtkm::cont::TryExecute(functor, vtkm::cont::GetGlobalRuntimeDeviceTracker(), DeviceList());
} }
template <typename Functor> template <typename Functor>
VTKM_CONT bool TryExecute(const Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker = VTKM_CONT bool TryExecute(
vtkm::cont::GetGlobalRuntimeDeviceTracker()) const Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
{ {
return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG()); return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
} }
template <typename Functor> template <typename Functor>
VTKM_CONT bool TryExecute(Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker = VTKM_CONT bool TryExecute(
vtkm::cont::GetGlobalRuntimeDeviceTracker()) Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
{ {
return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG()); return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
} }

@ -52,7 +52,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayIn, ContObjectType, Device>
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object, VTKM_CONT ExecObjectType operator()(const ContObjectType& object,
const InputDomainType& vtkmNotUsed(inputDomain), const InputDomainType& vtkmNotUsed(inputDomain),
vtkm::Id inputRange, vtkm::Id vtkmNotUsed(outputRange)) const vtkm::Id inputRange,
vtkm::Id vtkmNotUsed(outputRange)) const
{ {
if (object.GetNumberOfValues() != inputRange) if (object.GetNumberOfValues() != inputRange)
{ {

@ -55,7 +55,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayInOut, ContObjectType, Device
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType object, VTKM_CONT ExecObjectType operator()(ContObjectType object,
const InputDomainType& vtkmNotUsed(inputDomain), const InputDomainType& vtkmNotUsed(inputDomain),
vtkm::Id vtkmNotUsed(inputRange), vtkm::Id outputRange) const vtkm::Id vtkmNotUsed(inputRange),
vtkm::Id outputRange) const
{ {
if (object.GetNumberOfValues() != outputRange) if (object.GetNumberOfValues() != outputRange)
{ {

@ -54,7 +54,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayOut, ContObjectType, Device>
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType object, VTKM_CONT ExecObjectType operator()(ContObjectType object,
const InputDomainType& vtkmNotUsed(inputDomain), const InputDomainType& vtkmNotUsed(inputDomain),
vtkm::Id vtkmNotUsed(inputRange), vtkm::Id outputRange) const vtkm::Id vtkmNotUsed(inputRange),
vtkm::Id outputRange) const
{ {
return object.PrepareForOutput(outputRange, Device()); return object.PrepareForOutput(outputRange, Device());
} }

@ -50,13 +50,16 @@ struct TransportTagAtomicArray
template <typename T, typename Device> template <typename T, typename Device>
struct Transport<vtkm::cont::arg::TransportTagAtomicArray, struct Transport<vtkm::cont::arg::TransportTagAtomicArray,
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>, Device> vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>,
Device>
{ {
typedef vtkm::exec::AtomicArray<T, Device> ExecObjectType; typedef vtkm::exec::AtomicArray<T, Device> ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> array, VTKM_CONT ExecObjectType operator()(vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> array,
const InputDomainType&, vtkm::Id, vtkm::Id) const const InputDomainType&,
vtkm::Id,
vtkm::Id) const
{ {
// Note: we ignore the size of the domain because the randomly accessed // Note: we ignore the size of the domain because the randomly accessed
// array might not have the same size depending on how the user is using // array might not have the same size depending on how the user is using

@ -44,18 +44,20 @@ struct TransportTagCellSetIn
}; };
template <typename FromTopology, typename ToTopology, typename ContObjectType, typename Device> template <typename FromTopology, typename ToTopology, typename ContObjectType, typename Device>
struct Transport<vtkm::cont::arg::TransportTagCellSetIn<FromTopology, ToTopology>, ContObjectType, struct Transport<vtkm::cont::arg::TransportTagCellSetIn<FromTopology, ToTopology>,
ContObjectType,
Device> Device>
{ {
VTKM_IS_CELL_SET(ContObjectType); VTKM_IS_CELL_SET(ContObjectType);
typedef typedef
typename ContObjectType::template ExecutionTypes<Device, FromTopology, typename ContObjectType::template ExecutionTypes<Device,
FromTopology,
ToTopology>::ExecObjectType ExecObjectType; ToTopology>::ExecObjectType ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object, const InputDomainType&, VTKM_CONT ExecObjectType
vtkm::Id, vtkm::Id) const operator()(const ContObjectType& object, const InputDomainType&, vtkm::Id, vtkm::Id) const
{ {
return object.PrepareForInput(Device(), FromTopology(), ToTopology()); return object.PrepareForInput(Device(), FromTopology(), ToTopology());
} }

@ -56,8 +56,8 @@ struct Transport<vtkm::cont::arg::TransportTagExecObject, ContObjectType, Device
typedef ContObjectType ExecObjectType; typedef ContObjectType ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object, const InputDomainType&, VTKM_CONT ExecObjectType
vtkm::Id, vtkm::Id) const operator()(const ContObjectType& object, const InputDomainType&, vtkm::Id, vtkm::Id) const
{ {
return object; return object;
} }

@ -81,7 +81,8 @@ inline static vtkm::Id TopologyDomainSize(const vtkm::cont::CellSet& cellSet,
} // namespace detail } // namespace detail
template <typename TopologyElementTag, typename ContObjectType, typename Device> template <typename TopologyElementTag, typename ContObjectType, typename Device>
struct Transport<vtkm::cont::arg::TransportTagTopologyFieldIn<TopologyElementTag>, ContObjectType, struct Transport<vtkm::cont::arg::TransportTagTopologyFieldIn<TopologyElementTag>,
ContObjectType,
Device> Device>
{ {
VTKM_IS_ARRAY_HANDLE(ContObjectType); VTKM_IS_ARRAY_HANDLE(ContObjectType);
@ -89,8 +90,10 @@ struct Transport<vtkm::cont::arg::TransportTagTopologyFieldIn<TopologyElementTag
typedef typename ContObjectType::template ExecutionTypes<Device>::PortalConst ExecObjectType; typedef typename ContObjectType::template ExecutionTypes<Device>::PortalConst ExecObjectType;
VTKM_CONT VTKM_CONT
ExecObjectType operator()(const ContObjectType& object, const vtkm::cont::CellSet& inputDomain, ExecObjectType operator()(const ContObjectType& object,
vtkm::Id, vtkm::Id) const const vtkm::cont::CellSet& inputDomain,
vtkm::Id,
vtkm::Id) const
{ {
if (object.GetNumberOfValues() != detail::TopologyDomainSize(inputDomain, TopologyElementTag())) if (object.GetNumberOfValues() != detail::TopologyDomainSize(inputDomain, TopologyElementTag()))
{ {

@ -60,8 +60,8 @@ struct Transport<vtkm::cont::arg::TransportTagWholeArrayIn, ContObjectType, Devi
typedef vtkm::exec::ExecutionWholeArrayConst<ValueType, StorageTag, Device> ExecObjectType; typedef vtkm::exec::ExecutionWholeArrayConst<ValueType, StorageTag, Device> ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType array, const InputDomainType&, vtkm::Id, VTKM_CONT ExecObjectType
vtkm::Id) const operator()(ContObjectType array, const InputDomainType&, vtkm::Id, vtkm::Id) const
{ {
// Note: we ignore the size of the domain because the randomly accessed // Note: we ignore the size of the domain because the randomly accessed
// array might not have the same size depending on how the user is using // array might not have the same size depending on how the user is using

@ -62,8 +62,8 @@ struct Transport<vtkm::cont::arg::TransportTagWholeArrayInOut, ContObjectType, D
typedef vtkm::exec::ExecutionWholeArray<ValueType, StorageTag, Device> ExecObjectType; typedef vtkm::exec::ExecutionWholeArray<ValueType, StorageTag, Device> ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType array, const InputDomainType&, vtkm::Id, VTKM_CONT ExecObjectType
vtkm::Id) const operator()(ContObjectType array, const InputDomainType&, vtkm::Id, vtkm::Id) const
{ {
// Note: we ignore the size of the domain because the randomly accessed // Note: we ignore the size of the domain because the randomly accessed
// array might not have the same size depending on how the user is using // array might not have the same size depending on how the user is using

@ -62,8 +62,8 @@ struct Transport<vtkm::cont::arg::TransportTagWholeArrayOut, ContObjectType, Dev
typedef vtkm::exec::ExecutionWholeArray<ValueType, StorageTag, Device> ExecObjectType; typedef vtkm::exec::ExecutionWholeArray<ValueType, StorageTag, Device> ExecObjectType;
template <typename InputDomainType> template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType array, const InputDomainType&, vtkm::Id, VTKM_CONT ExecObjectType
vtkm::Id) const operator()(ContObjectType array, const InputDomainType&, vtkm::Id, vtkm::Id) const
{ {
// Note: we ignore the size of the domain because the randomly accessed // Note: we ignore the size of the domain because the randomly accessed
// array might not have the same size depending on how the user is using // array might not have the same size depending on how the user is using

@ -66,7 +66,9 @@ template <typename TypeList, typename ArrayType>
struct TypeCheck<TypeCheckTagArray<TypeList>, ArrayType> struct TypeCheck<TypeCheckTagArray<TypeList>, ArrayType>
{ {
static const bool value = detail::TypeCheckArrayValueType< static const bool value = detail::TypeCheckArrayValueType<
TypeList, ArrayType, vtkm::cont::internal::ArrayHandleCheck<ArrayType>::type::value>::value; TypeList,
ArrayType,
vtkm::cont::internal::ArrayHandleCheck<ArrayType>::type::value>::value;
}; };
} }
} }

@ -72,11 +72,12 @@ void TransportWholeCellSetIn(Device)
typedef vtkm::TopologyElementTagPoint FromType; typedef vtkm::TopologyElementTagPoint FromType;
typedef vtkm::TopologyElementTagCell ToType; typedef vtkm::TopologyElementTagCell ToType;
typedef typename vtkm::cont::CellSetExplicit<>::template ExecutionTypes< typedef typename vtkm::cont::CellSetExplicit<>::
Device, FromType, ToType>::ExecObjectType ExecObjectType; template ExecutionTypes<Device, FromType, ToType>::ExecObjectType ExecObjectType;
vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagCellSetIn<FromType, ToType>, vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagCellSetIn<FromType, ToType>,
vtkm::cont::CellSetExplicit<>, Device> vtkm::cont::CellSetExplicit<>,
Device>
transport; transport;
TestKernel<ExecObjectType> kernel; TestKernel<ExecObjectType> kernel;

@ -119,13 +119,16 @@ struct TryWholeArrayType
{ {
typedef vtkm::cont::ArrayHandle<T> ArrayHandleType; typedef vtkm::cont::ArrayHandle<T> ArrayHandleType;
typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayIn, ArrayHandleType, typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayIn,
ArrayHandleType,
Device> Device>
InTransportType; InTransportType;
typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayInOut, typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayInOut,
ArrayHandleType, Device> ArrayHandleType,
Device>
InOutTransportType; InOutTransportType;
typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayOut, ArrayHandleType, typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagWholeArrayOut,
ArrayHandleType,
Device> Device>
OutTransportType; OutTransportType;
@ -171,7 +174,8 @@ struct TryAtomicArrayType
{ {
typedef vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> ArrayHandleType; typedef vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> ArrayHandleType;
typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagAtomicArray, ArrayHandleType, typedef vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagAtomicArray,
ArrayHandleType,
Device> Device>
TransportType; TransportType;

@ -257,7 +257,8 @@ public:
VTKM_CONT VTKM_CONT
PortalConstType PrepareForInput(bool) PortalConstType PrepareForInput(bool)
{ {
return PortalConstType(this->Storage->GetDevicePointer(), this->Storage->GetDevicePointer() + return PortalConstType(this->Storage->GetDevicePointer(),
this->Storage->GetDevicePointer() +
static_cast<difference_type>(Storage->GetNumberOfValues())); static_cast<difference_type>(Storage->GetNumberOfValues()));
} }
@ -266,7 +267,8 @@ public:
VTKM_CONT VTKM_CONT
PortalType PrepareForInPlace(bool) PortalType PrepareForInPlace(bool)
{ {
return PortalType(this->Storage->GetDevicePointer(), this->Storage->GetDevicePointer() + return PortalType(this->Storage->GetDevicePointer(),
this->Storage->GetDevicePointer() +
static_cast<difference_type>(Storage->GetNumberOfValues())); static_cast<difference_type>(Storage->GetNumberOfValues()));
} }
@ -279,7 +281,8 @@ public:
this->Storage->ReleaseResources(); this->Storage->ReleaseResources();
this->Storage->Allocate(numberOfValues); this->Storage->Allocate(numberOfValues);
return PortalType(this->Storage->GetDevicePointer(), this->Storage->GetDevicePointer() + return PortalType(this->Storage->GetDevicePointer(),
this->Storage->GetDevicePointer() +
static_cast<difference_type>(Storage->GetNumberOfValues())); static_cast<difference_type>(Storage->GetNumberOfValues()));
} }
@ -326,7 +329,8 @@ template <typename T>
class ArrayHandleCuda : public vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda> class ArrayHandleCuda : public vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda>
{ {
public: public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleCuda, (ArrayHandleCuda<T>), VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleCuda,
(ArrayHandleCuda<T>),
(vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda>)); (vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda>));
VTKM_CONT VTKM_CONT
@ -340,7 +344,8 @@ public:
/// ///
template <typename T> template <typename T>
VTKM_CONT vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda> make_ArrayHandleCuda( VTKM_CONT vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda> make_ArrayHandleCuda(
T* array, vtkm::Id length) T* array,
vtkm::Id length)
{ {
typedef vtkm::cont::cuda::StorageTagCuda StorageTag; typedef vtkm::cont::cuda::StorageTagCuda StorageTag;
typedef vtkm::cont::ArrayHandle<T, StorageTag> ArrayHandleType; typedef vtkm::cont::ArrayHandle<T, StorageTag> ArrayHandleType;
@ -349,7 +354,8 @@ VTKM_CONT vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda> make_Arra
template <typename T> template <typename T>
VTKM_CONT void printSummary_ArrayHandle( VTKM_CONT void printSummary_ArrayHandle(
const vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda>& array, std::ostream& out) const vtkm::cont::ArrayHandle<T, vtkm::cont::cuda::StorageTagCuda>& array,
std::ostream& out)
{ {
vtkm::Id sz = array.GetNumberOfValues(); vtkm::Id sz = array.GetNumberOfValues();
out << "sz= " << sz << " [(on device)]"; out << "sz= " << sz << " [(on device)]";

@ -36,8 +36,8 @@
const cudaError_t vtkm_cuda_check_async_error = cudaGetLastError(); \ const cudaError_t vtkm_cuda_check_async_error = cudaGetLastError(); \
if (vtkm_cuda_check_async_error != cudaSuccess) \ if (vtkm_cuda_check_async_error != cudaSuccess) \
{ \ { \
throw ::vtkm::cont::cuda::ErrorCuda(vtkm_cuda_check_async_error, __FILE__, __LINE__, \ throw ::vtkm::cont::cuda::ErrorCuda( \
"Unchecked asynchronous error"); \ vtkm_cuda_check_async_error, __FILE__, __LINE__, "Unchecked asynchronous error"); \
} \ } \
} \ } \
VTKM_SWALLOW_SEMICOLON_POST_BLOCK VTKM_SWALLOW_SEMICOLON_POST_BLOCK
@ -79,7 +79,9 @@ public:
this->SetMessage(message.str()); this->SetMessage(message.str());
} }
ErrorCuda(cudaError_t error, const std::string& file, vtkm::Id line, ErrorCuda(cudaError_t error,
const std::string& file,
vtkm::Id line,
const std::string& description) const std::string& description)
{ {
std::stringstream message; std::stringstream message;

@ -215,8 +215,8 @@ public:
#ifdef VTKM_USE_UNIFIED_MEMORY #ifdef VTKM_USE_UNIFIED_MEMORY
cudaDeviceSynchronize(); cudaDeviceSynchronize();
#endif #endif
::thrust::copy(this->Begin, this->End, ::thrust::copy(
vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal())); this->Begin, this->End, vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
} }
catch (...) catch (...)
{ {

@ -204,7 +204,8 @@ public:
return vtkmAtomicAdd(lockedValue, value); return vtkmAtomicAdd(lockedValue, value);
} }
inline __device__ T CompareAndSwap(vtkm::Id index, const vtkm::Int64& newValue, inline __device__ T CompareAndSwap(vtkm::Id index,
const vtkm::Int64& newValue,
const vtkm::Int64& oldValue) const const vtkm::Int64& oldValue) const
{ {
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index); T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
@ -237,7 +238,8 @@ private:
const vtkm::Int64& newValue, const vtkm::Int64& newValue,
const vtkm::Int64& oldValue) const const vtkm::Int64& oldValue) const
{ {
return atomicCAS((unsigned long long int*)address, (unsigned long long int)oldValue, return atomicCAS((unsigned long long int*)address,
(unsigned long long int)oldValue,
(unsigned long long int)newValue); (unsigned long long int)newValue);
} }
}; };
@ -248,7 +250,9 @@ class DeviceTaskTypes<vtkm::cont::DeviceAdapterTagCuda>
public: public:
template <typename WorkletType, typename InvocationType> template <typename WorkletType, typename InvocationType>
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask( static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id, const WorkletType& worklet,
const InvocationType& invocation,
vtkm::Id,
vtkm::Id globalIndexOffset = 0) vtkm::Id globalIndexOffset = 0)
{ {
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>; using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
@ -257,7 +261,9 @@ public:
template <typename WorkletType, typename InvocationType> template <typename WorkletType, typename InvocationType>
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask( static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id3, const WorkletType& worklet,
const InvocationType& invocation,
vtkm::Id3,
vtkm::Id globalIndexOffset = 0) vtkm::Id globalIndexOffset = 0)
{ {
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>; using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;

@ -90,7 +90,8 @@ static __global__ void DetermineProperXGridSize(vtkm::UInt32 desired_size,
} }
template <class FunctorType> template <class FunctorType>
__global__ void Schedule1DIndexKernel(FunctorType functor, vtkm::Id numberOfKernelsInvoked, __global__ void Schedule1DIndexKernel(FunctorType functor,
vtkm::Id numberOfKernelsInvoked,
vtkm::Id length) vtkm::Id length)
{ {
//Note a cuda launch can only handle at most 2B iterations of a kernel //Note a cuda launch can only handle at most 2B iterations of a kernel
@ -152,7 +153,8 @@ public:
template <class Functor> template <class Functor>
static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& rangeMax) static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& rangeMax)
{ {
const dim3 ranges(static_cast<vtkm::UInt32>(rangeMax[0]), static_cast<vtkm::UInt32>(rangeMax[1]), const dim3 ranges(static_cast<vtkm::UInt32>(rangeMax[0]),
static_cast<vtkm::UInt32>(rangeMax[1]),
static_cast<vtkm::UInt32>(rangeMax[2])); static_cast<vtkm::UInt32>(rangeMax[2]));
std::vector<PerfRecord> results; std::vector<PerfRecord> results;
vtkm::UInt32 indexTable[16] = { 1, 2, 4, 8, 12, 16, 20, 24, 28, 30, 32, 64, 128, 256, 512, 1024 }; vtkm::UInt32 indexTable[16] = { 1, 2, 4, 8, 12, 16, 20, 24, 28, 30, 32, 64, 128, 256, 512, 1024 };
@ -286,8 +288,8 @@ private:
{ {
try try
{ {
::thrust::copy(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), ::thrust::copy(
IteratorBegin(output)); thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), IteratorBegin(output));
} }
catch (...) catch (...)
{ {
@ -296,8 +298,10 @@ private:
} }
template <class ValueIterator, class StencilPortal, class OutputPortal, class UnaryPredicate> template <class ValueIterator, class StencilPortal, class OutputPortal, class UnaryPredicate>
VTKM_CONT static vtkm::Id CopyIfPortal(ValueIterator valuesBegin, ValueIterator valuesEnd, VTKM_CONT static vtkm::Id CopyIfPortal(ValueIterator valuesBegin,
StencilPortal stencil, OutputPortal output, ValueIterator valuesEnd,
StencilPortal stencil,
OutputPortal output,
UnaryPredicate unary_predicate) UnaryPredicate unary_predicate)
{ {
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType; typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType;
@ -311,8 +315,8 @@ private:
try try
{ {
IteratorType newLast = ::thrust::copy_if(thrust::cuda::par, valuesBegin, valuesEnd, IteratorType newLast = ::thrust::copy_if(
IteratorBegin(stencil), outputBegin, up); thrust::cuda::par, valuesBegin, valuesEnd, IteratorBegin(stencil), outputBegin, up);
return static_cast<vtkm::Id>(::thrust::distance(outputBegin, newLast)); return static_cast<vtkm::Id>(::thrust::distance(outputBegin, newLast));
} }
catch (...) catch (...)
@ -323,22 +327,28 @@ private:
} }
template <class ValuePortal, class StencilPortal, class OutputPortal, class UnaryPredicate> template <class ValuePortal, class StencilPortal, class OutputPortal, class UnaryPredicate>
VTKM_CONT static vtkm::Id CopyIfPortal(ValuePortal values, StencilPortal stencil, VTKM_CONT static vtkm::Id CopyIfPortal(ValuePortal values,
OutputPortal output, UnaryPredicate unary_predicate) StencilPortal stencil,
OutputPortal output,
UnaryPredicate unary_predicate)
{ {
return CopyIfPortal(IteratorBegin(values), IteratorEnd(values), stencil, output, return CopyIfPortal(
unary_predicate); IteratorBegin(values), IteratorEnd(values), stencil, output, unary_predicate);
} }
template <class InputPortal, class OutputPortal> template <class InputPortal, class OutputPortal>
VTKM_CONT static void CopySubRangePortal(const InputPortal& input, vtkm::Id inputOffset, VTKM_CONT static void CopySubRangePortal(const InputPortal& input,
vtkm::Id size, const OutputPortal& output, vtkm::Id inputOffset,
vtkm::Id size,
const OutputPortal& output,
vtkm::Id outputOffset) vtkm::Id outputOffset)
{ {
try try
{ {
::thrust::copy_n(thrust::cuda::par, IteratorBegin(input) + inputOffset, ::thrust::copy_n(thrust::cuda::par,
static_cast<std::size_t>(size), IteratorBegin(output) + outputOffset); IteratorBegin(input) + inputOffset,
static_cast<std::size_t>(size),
IteratorBegin(output) + outputOffset);
} }
catch (...) catch (...)
{ {
@ -347,7 +357,8 @@ private:
} }
template <class InputPortal, class ValuesPortal, class OutputPortal> template <class InputPortal, class ValuesPortal, class OutputPortal>
VTKM_CONT static void LowerBoundsPortal(const InputPortal& input, const ValuesPortal& values, VTKM_CONT static void LowerBoundsPortal(const InputPortal& input,
const ValuesPortal& values,
const OutputPortal& output) const OutputPortal& output)
{ {
typedef typename ValuesPortal::ValueType ValueType; typedef typename ValuesPortal::ValueType ValueType;
@ -363,8 +374,10 @@ private:
} }
template <class InputPortal, class ValuesPortal, class OutputPortal, class BinaryCompare> template <class InputPortal, class ValuesPortal, class OutputPortal, class BinaryCompare>
VTKM_CONT static void LowerBoundsPortal(const InputPortal& input, const ValuesPortal& values, VTKM_CONT static void LowerBoundsPortal(const InputPortal& input,
const OutputPortal& output, BinaryCompare binary_compare) const ValuesPortal& values,
const OutputPortal& output,
BinaryCompare binary_compare)
{ {
typedef typename InputPortal::ValueType ValueType; typedef typename InputPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType, BinaryCompare> bop( vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType, BinaryCompare> bop(
@ -372,8 +385,13 @@ private:
try try
{ {
::thrust::lower_bound(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), ::thrust::lower_bound(thrust::cuda::par,
IteratorBegin(values), IteratorEnd(values), IteratorBegin(output), bop); IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(values),
IteratorEnd(values),
IteratorBegin(output),
bop);
} }
catch (...) catch (...)
{ {
@ -388,7 +406,8 @@ private:
} }
template <class InputPortal, typename T, class BinaryFunctor> template <class InputPortal, typename T, class BinaryFunctor>
VTKM_CONT static T ReducePortal(const InputPortal& input, T initialValue, VTKM_CONT static T ReducePortal(const InputPortal& input,
T initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
using fast_path = std::is_same<typename InputPortal::ValueType, T>; using fast_path = std::is_same<typename InputPortal::ValueType, T>;
@ -396,8 +415,10 @@ private:
} }
template <class InputPortal, typename T, class BinaryFunctor> template <class InputPortal, typename T, class BinaryFunctor>
VTKM_CONT static T ReducePortalImpl(const InputPortal& input, T initialValue, VTKM_CONT static T ReducePortalImpl(const InputPortal& input,
BinaryFunctor binary_functor, std::true_type) T initialValue,
BinaryFunctor binary_functor,
std::true_type)
{ {
//The portal type and the initial value are the same so we can use //The portal type and the initial value are the same so we can use
//the thrust reduction algorithm //the thrust reduction algorithm
@ -405,8 +426,8 @@ private:
try try
{ {
return ::thrust::reduce(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), return ::thrust::reduce(
initialValue, bop); thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), initialValue, bop);
} }
catch (...) catch (...)
{ {
@ -417,8 +438,10 @@ private:
} }
template <class InputPortal, typename T, class BinaryFunctor> template <class InputPortal, typename T, class BinaryFunctor>
VTKM_CONT static T ReducePortalImpl(const InputPortal& input, T initialValue, VTKM_CONT static T ReducePortalImpl(const InputPortal& input,
BinaryFunctor binary_functor, std::false_type) T initialValue,
BinaryFunctor binary_functor,
std::false_type)
{ {
//The portal type and the initial value ARENT the same type so we have //The portal type and the initial value ARENT the same type so we have
//to a slower approach, where we wrap the input portal inside a cast //to a slower approach, where we wrap the input portal inside a cast
@ -431,8 +454,8 @@ private:
try try
{ {
return ::thrust::reduce(thrust::cuda::par, IteratorBegin(castPortal), IteratorEnd(castPortal), return ::thrust::reduce(
initialValue, bop); thrust::cuda::par, IteratorBegin(castPortal), IteratorEnd(castPortal), initialValue, bop);
} }
catch (...) catch (...)
{ {
@ -442,9 +465,13 @@ private:
return initialValue; return initialValue;
} }
template <class KeysPortal, class ValuesPortal, class KeysOutputPortal, class ValueOutputPortal, template <class KeysPortal,
class ValuesPortal,
class KeysOutputPortal,
class ValueOutputPortal,
class BinaryFunctor> class BinaryFunctor>
VTKM_CONT static vtkm::Id ReduceByKeyPortal(const KeysPortal& keys, const ValuesPortal& values, VTKM_CONT static vtkm::Id ReduceByKeyPortal(const KeysPortal& keys,
const ValuesPortal& values,
const KeysOutputPortal& keys_output, const KeysOutputPortal& keys_output,
const ValueOutputPortal& values_output, const ValueOutputPortal& values_output,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
@ -464,9 +491,14 @@ private:
try try
{ {
result_iterators = ::thrust::reduce_by_key( result_iterators = ::thrust::reduce_by_key(vtkm_cuda_policy(),
vtkm_cuda_policy(), IteratorBegin(keys), IteratorEnd(keys), IteratorBegin(values), IteratorBegin(keys),
keys_out_begin, values_out_begin, binaryPredicate, bop); IteratorEnd(keys),
IteratorBegin(values),
keys_out_begin,
values_out_begin,
binaryPredicate,
bop);
} }
catch (...) catch (...)
{ {
@ -482,13 +514,17 @@ private:
{ {
typedef typename OutputPortal::ValueType ValueType; typedef typename OutputPortal::ValueType ValueType;
return ScanExclusivePortal(input, output, (::thrust::plus<ValueType>()), return ScanExclusivePortal(input,
output,
(::thrust::plus<ValueType>()),
vtkm::TypeTraits<ValueType>::ZeroInitialization()); vtkm::TypeTraits<ValueType>::ZeroInitialization());
} }
template <class InputPortal, class OutputPortal, class BinaryFunctor> template <class InputPortal, class OutputPortal, class BinaryFunctor>
VTKM_CONT static typename InputPortal::ValueType ScanExclusivePortal( VTKM_CONT static typename InputPortal::ValueType ScanExclusivePortal(
const InputPortal& input, const OutputPortal& output, BinaryFunctor binaryOp, const InputPortal& input,
const OutputPortal& output,
BinaryFunctor binaryOp,
typename InputPortal::ValueType initialValue) typename InputPortal::ValueType initialValue)
{ {
// Use iterator to get value so that thrust device_ptr has chance to handle // Use iterator to get value so that thrust device_ptr has chance to handle
@ -509,9 +545,12 @@ private:
vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType, BinaryFunctor> bop(binaryOp); vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType, BinaryFunctor> bop(binaryOp);
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType; typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType;
IteratorType end = IteratorType end = ::thrust::exclusive_scan(thrust::cuda::par,
::thrust::exclusive_scan(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), IteratorBegin(input),
IteratorBegin(output), initialValue, bop); IteratorEnd(input),
IteratorBegin(output),
initialValue,
bop);
//Store the new value for the end of the array. This is done because //Store the new value for the end of the array. This is done because
//with items such as the transpose array it is unsafe to pass the //with items such as the transpose array it is unsafe to pass the
@ -548,8 +587,8 @@ private:
try try
{ {
IteratorType end = ::thrust::inclusive_scan(thrust::cuda::par, IteratorBegin(input), IteratorType end = ::thrust::inclusive_scan(
IteratorEnd(input), IteratorBegin(output), bop); thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), IteratorBegin(output), bop);
return *(end - 1); return *(end - 1);
} }
catch (...) catch (...)
@ -563,19 +602,27 @@ private:
template <typename KeysPortal, typename ValuesPortal, typename OutputPortal> template <typename KeysPortal, typename ValuesPortal, typename OutputPortal>
VTKM_CONT static typename ValuesPortal::ValueType ScanInclusiveByKeyPortal( VTKM_CONT static typename ValuesPortal::ValueType ScanInclusiveByKeyPortal(
const KeysPortal& keys, const ValuesPortal& values, const OutputPortal& output) const KeysPortal& keys,
const ValuesPortal& values,
const OutputPortal& output)
{ {
using KeyType = typename KeysPortal::ValueType; using KeyType = typename KeysPortal::ValueType;
typedef typename OutputPortal::ValueType ValueType; typedef typename OutputPortal::ValueType ValueType;
return ScanInclusiveByKeyPortal(keys, values, output, ::thrust::equal_to<KeyType>(), return ScanInclusiveByKeyPortal(
::thrust::plus<ValueType>()); keys, values, output, ::thrust::equal_to<KeyType>(), ::thrust::plus<ValueType>());
} }
template <typename KeysPortal, typename ValuesPortal, typename OutputPortal, template <typename KeysPortal,
typename BinaryPredicate, typename AssociativeOperator> typename ValuesPortal,
typename OutputPortal,
typename BinaryPredicate,
typename AssociativeOperator>
VTKM_CONT static typename ValuesPortal::ValueType ScanInclusiveByKeyPortal( VTKM_CONT static typename ValuesPortal::ValueType ScanInclusiveByKeyPortal(
const KeysPortal& keys, const ValuesPortal& values, const OutputPortal& output, const KeysPortal& keys,
BinaryPredicate binary_predicate, AssociativeOperator binary_operator) const ValuesPortal& values,
const OutputPortal& output,
BinaryPredicate binary_predicate,
AssociativeOperator binary_operator)
{ {
typedef typename KeysPortal::ValueType KeyType; typedef typename KeysPortal::ValueType KeyType;
vtkm::exec::cuda::internal::WrappedBinaryOperator<KeyType, BinaryPredicate> bpred( vtkm::exec::cuda::internal::WrappedBinaryOperator<KeyType, BinaryPredicate> bpred(
@ -587,9 +634,13 @@ private:
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType; typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType;
try try
{ {
IteratorType end = IteratorType end = ::thrust::inclusive_scan_by_key(thrust::cuda::par,
::thrust::inclusive_scan_by_key(thrust::cuda::par, IteratorBegin(keys), IteratorEnd(keys), IteratorBegin(keys),
IteratorBegin(values), IteratorBegin(output), bpred, bop); IteratorEnd(keys),
IteratorBegin(values),
IteratorBegin(output),
bpred,
bop);
return *(end - 1); return *(end - 1);
} }
catch (...) catch (...)
@ -602,20 +653,30 @@ private:
} }
template <typename KeysPortal, typename ValuesPortal, typename OutputPortal> template <typename KeysPortal, typename ValuesPortal, typename OutputPortal>
VTKM_CONT static void ScanExclusiveByKeyPortal(const KeysPortal& keys, const ValuesPortal& values, VTKM_CONT static void ScanExclusiveByKeyPortal(const KeysPortal& keys,
const ValuesPortal& values,
const OutputPortal& output) const OutputPortal& output)
{ {
using KeyType = typename KeysPortal::ValueType; using KeyType = typename KeysPortal::ValueType;
typedef typename OutputPortal::ValueType ValueType; typedef typename OutputPortal::ValueType ValueType;
ScanExclusiveByKeyPortal(keys, values, output, ScanExclusiveByKeyPortal(keys,
values,
output,
vtkm::TypeTraits<ValueType>::ZeroInitialization(), vtkm::TypeTraits<ValueType>::ZeroInitialization(),
::thrust::equal_to<KeyType>(), ::thrust::plus<ValueType>()); ::thrust::equal_to<KeyType>(),
::thrust::plus<ValueType>());
} }
template <typename KeysPortal, typename ValuesPortal, typename OutputPortal, typename T, template <typename KeysPortal,
typename BinaryPredicate, typename AssociativeOperator> typename ValuesPortal,
VTKM_CONT static void ScanExclusiveByKeyPortal(const KeysPortal& keys, const ValuesPortal& values, typename OutputPortal,
const OutputPortal& output, T initValue, typename T,
typename BinaryPredicate,
typename AssociativeOperator>
VTKM_CONT static void ScanExclusiveByKeyPortal(const KeysPortal& keys,
const ValuesPortal& values,
const OutputPortal& output,
T initValue,
BinaryPredicate binary_predicate, BinaryPredicate binary_predicate,
AssociativeOperator binary_operator) AssociativeOperator binary_operator)
{ {
@ -629,9 +690,14 @@ private:
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType; typedef typename detail::IteratorTraits<OutputPortal>::IteratorType IteratorType;
try try
{ {
IteratorType end = ::thrust::exclusive_scan_by_key( IteratorType end = ::thrust::exclusive_scan_by_key(thrust::cuda::par,
thrust::cuda::par, IteratorBegin(keys), IteratorEnd(keys), IteratorBegin(values), IteratorBegin(keys),
IteratorBegin(output), initValue, bpred, bop); IteratorEnd(keys),
IteratorBegin(values),
IteratorBegin(output),
initValue,
bpred,
bop);
return; return;
} }
catch (...) catch (...)
@ -674,7 +740,8 @@ private:
} }
template <class KeysPortal, class ValuesPortal, class BinaryCompare> template <class KeysPortal, class ValuesPortal, class BinaryCompare>
VTKM_CONT static void SortByKeyPortal(const KeysPortal& keys, const ValuesPortal& values, VTKM_CONT static void SortByKeyPortal(const KeysPortal& keys,
const ValuesPortal& values,
BinaryCompare binary_compare) BinaryCompare binary_compare)
{ {
typedef typename KeysPortal::ValueType ValueType; typedef typename KeysPortal::ValueType ValueType;
@ -682,8 +749,8 @@ private:
binary_compare); binary_compare);
try try
{ {
::thrust::sort_by_key(vtkm_cuda_policy(), IteratorBegin(keys), IteratorEnd(keys), ::thrust::sort_by_key(
IteratorBegin(values), bop); vtkm_cuda_policy(), IteratorBegin(keys), IteratorEnd(keys), IteratorBegin(values), bop);
} }
catch (...) catch (...)
{ {
@ -730,13 +797,18 @@ private:
} }
template <class InputPortal, class ValuesPortal, class OutputPortal> template <class InputPortal, class ValuesPortal, class OutputPortal>
VTKM_CONT static void UpperBoundsPortal(const InputPortal& input, const ValuesPortal& values, VTKM_CONT static void UpperBoundsPortal(const InputPortal& input,
const ValuesPortal& values,
const OutputPortal& output) const OutputPortal& output)
{ {
try try
{ {
::thrust::upper_bound(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), ::thrust::upper_bound(thrust::cuda::par,
IteratorBegin(values), IteratorEnd(values), IteratorBegin(output)); IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(values),
IteratorEnd(values),
IteratorBegin(output));
} }
catch (...) catch (...)
{ {
@ -745,8 +817,10 @@ private:
} }
template <class InputPortal, class ValuesPortal, class OutputPortal, class BinaryCompare> template <class InputPortal, class ValuesPortal, class OutputPortal, class BinaryCompare>
VTKM_CONT static void UpperBoundsPortal(const InputPortal& input, const ValuesPortal& values, VTKM_CONT static void UpperBoundsPortal(const InputPortal& input,
const OutputPortal& output, BinaryCompare binary_compare) const ValuesPortal& values,
const OutputPortal& output,
BinaryCompare binary_compare)
{ {
typedef typename OutputPortal::ValueType ValueType; typedef typename OutputPortal::ValueType ValueType;
@ -754,8 +828,13 @@ private:
binary_compare); binary_compare);
try try
{ {
::thrust::upper_bound(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), ::thrust::upper_bound(thrust::cuda::par,
IteratorBegin(values), IteratorEnd(values), IteratorBegin(output), bop); IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(values),
IteratorEnd(values),
IteratorBegin(output),
bop);
} }
catch (...) catch (...)
{ {
@ -769,8 +848,11 @@ private:
{ {
try try
{ {
::thrust::upper_bound(thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), ::thrust::upper_bound(thrust::cuda::par,
IteratorBegin(values_output), IteratorEnd(values_output), IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(values_output),
IteratorEnd(values_output),
IteratorBegin(values_output)); IteratorBegin(values_output));
} }
catch (...) catch (...)
@ -811,15 +893,17 @@ public:
UnaryPredicate unary_predicate) UnaryPredicate unary_predicate)
{ {
vtkm::Id size = stencil.GetNumberOfValues(); vtkm::Id size = stencil.GetNumberOfValues();
vtkm::Id newSize = CopyIfPortal( vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTag()),
input.PrepareForInput(DeviceAdapterTag()), stencil.PrepareForInput(DeviceAdapterTag()), stencil.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(size, DeviceAdapterTag()), unary_predicate); output.PrepareForOutput(size, DeviceAdapterTag()),
unary_predicate);
output.Shrink(newSize); output.Shrink(newSize);
} }
template <typename T, typename U, class SIn, class SOut> template <typename T, typename U, class SIn, class SOut>
VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, SIn>& input, VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, SIn>& input,
vtkm::Id inputStartIndex, vtkm::Id numberOfElementsToCopy, vtkm::Id inputStartIndex,
vtkm::Id numberOfElementsToCopy,
vtkm::cont::ArrayHandle<U, SOut>& output, vtkm::cont::ArrayHandle<U, SOut>& output,
vtkm::Id outputIndex = 0) vtkm::Id outputIndex = 0)
{ {
@ -853,8 +937,10 @@ public:
output = temp; output = temp;
} }
} }
CopySubRangePortal(input.PrepareForInput(DeviceAdapterTag()), inputStartIndex, CopySubRangePortal(input.PrepareForInput(DeviceAdapterTag()),
numberOfElementsToCopy, output.PrepareForInPlace(DeviceAdapterTag()), inputStartIndex,
numberOfElementsToCopy,
output.PrepareForInPlace(DeviceAdapterTag()),
outputIndex); outputIndex);
return true; return true;
} }
@ -879,7 +965,8 @@ public:
vtkm::Id numberOfValues = values.GetNumberOfValues(); vtkm::Id numberOfValues = values.GetNumberOfValues();
LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTag()), LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), binary_compare); output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
binary_compare);
} }
template <class SIn, class SOut> template <class SIn, class SOut>
@ -902,7 +989,8 @@ public:
} }
template <typename T, typename U, class SIn, class BinaryFunctor> template <typename T, typename U, class SIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, SIn>& input, U initialValue, VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, SIn>& input,
U initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
const vtkm::Id numberOfValues = input.GetNumberOfValues(); const vtkm::Id numberOfValues = input.GetNumberOfValues();
@ -913,7 +1001,12 @@ public:
return ReducePortal(input.PrepareForInput(DeviceAdapterTag()), initialValue, binary_functor); return ReducePortal(input.PrepareForInput(DeviceAdapterTag()), initialValue, binary_functor);
} }
template <typename T, typename U, class KIn, class VIn, class KOut, class VOut, template <typename T,
typename U,
class KIn,
class VIn,
class KOut,
class VOut,
class BinaryFunctor> class BinaryFunctor>
VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -928,10 +1021,12 @@ public:
{ {
return; return;
} }
vtkm::Id reduced_size = ReduceByKeyPortal( vtkm::Id reduced_size =
keys.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()), ReduceByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()),
keys_output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
values_output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), binary_functor); keys_output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
values_output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
binary_functor);
keys_output.Shrink(reduced_size); keys_output.Shrink(reduced_size);
values_output.Shrink(reduced_size); values_output.Shrink(reduced_size);
@ -960,7 +1055,8 @@ public:
template <typename T, class SIn, class SOut, class BinaryFunctor> template <typename T, class SIn, class SOut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, SIn>& input, VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, SIn>& input,
vtkm::cont::ArrayHandle<T, SOut>& output, vtkm::cont::ArrayHandle<T, SOut>& output,
BinaryFunctor binary_functor, const T& initialValue) BinaryFunctor binary_functor,
const T& initialValue)
{ {
const vtkm::Id numberOfValues = input.GetNumberOfValues(); const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0) if (numberOfValues <= 0)
@ -976,7 +1072,8 @@ public:
input.PrepareForInput(DeviceAdapterTag()); input.PrepareForInput(DeviceAdapterTag());
return ScanExclusivePortal(input.PrepareForInput(DeviceAdapterTag()), return ScanExclusivePortal(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
binary_functor, initialValue); binary_functor,
initialValue);
} }
template <typename T, class SIn, class SOut> template <typename T, class SIn, class SOut>
@ -1044,7 +1141,11 @@ public:
output.PrepareForOutput(numberOfValues, DeviceAdapterTag())); output.PrepareForOutput(numberOfValues, DeviceAdapterTag()));
} }
template <typename T, typename U, typename KIn, typename VIn, typename VOut, template <typename T,
typename U,
typename KIn,
typename VIn,
typename VOut,
typename BinaryFunctor> typename BinaryFunctor>
VTKM_CONT static T ScanInclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static T ScanInclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -1067,7 +1168,8 @@ public:
return ScanInclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), return ScanInclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
::thrust::equal_to<T>(), binary_functor); ::thrust::equal_to<T>(),
binary_functor);
} }
template <typename T, typename U, typename KIn, typename VIn, typename VOut> template <typename T, typename U, typename KIn, typename VIn, typename VOut>
@ -1091,15 +1193,21 @@ public:
ScanExnclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), ScanExnclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
vtkm::TypeTraits<T>::ZeroInitialization(), vtkm::Add()); vtkm::TypeTraits<T>::ZeroInitialization(),
vtkm::Add());
} }
template <typename T, typename U, typename KIn, typename VIn, typename VOut, template <typename T,
typename U,
typename KIn,
typename VIn,
typename VOut,
typename BinaryFunctor> typename BinaryFunctor>
VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output, vtkm::cont::ArrayHandle<U, VOut>& output,
const U& initialValue, BinaryFunctor binary_functor) const U& initialValue,
BinaryFunctor binary_functor)
{ {
const vtkm::Id numberOfValues = keys.GetNumberOfValues(); const vtkm::Id numberOfValues = keys.GetNumberOfValues();
if (numberOfValues <= 0) if (numberOfValues <= 0)
@ -1117,7 +1225,9 @@ public:
ScanExclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), ScanExclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
initialValue, ::thrust::equal_to<T>(), binary_functor); initialValue,
::thrust::equal_to<T>(),
binary_functor);
} }
// Because of some funny code conversions in nvcc, kernels for devices have to // Because of some funny code conversions in nvcc, kernels for devices have to
// be public. // be public.
@ -1229,8 +1339,8 @@ public:
//handle datasets larger than 2B, we need to execute multiple kernels //handle datasets larger than 2B, we need to execute multiple kernels
if (totalBlocks < maxblocksPerLaunch) if (totalBlocks < maxblocksPerLaunch)
{ {
Schedule1DIndexKernel<Functor><<<totalBlocks, blockSize>>>(functor, vtkm::Id(0), Schedule1DIndexKernel<Functor><<<totalBlocks, blockSize>>>(
numInstances); functor, vtkm::Id(0), numInstances);
} }
else else
{ {
@ -1346,7 +1456,8 @@ public:
BinaryCompare binary_compare) BinaryCompare binary_compare)
{ {
SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTag()), SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTag()),
values.PrepareForInPlace(DeviceAdapterTag()), binary_compare); values.PrepareForInPlace(DeviceAdapterTag()),
binary_compare);
} }
template <typename T, class Storage> template <typename T, class Storage>
@ -1386,7 +1497,8 @@ public:
vtkm::Id numberOfValues = values.GetNumberOfValues(); vtkm::Id numberOfValues = values.GetNumberOfValues();
UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTag()), UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), binary_compare); output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
binary_compare);
} }
template <class SIn, class SOut> template <class SIn, class SOut>

@ -97,14 +97,16 @@ struct IteratorTraits
template <typename PortalType> template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin( VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
PortalType portal, detail::ThrustIteratorFromArrayPortalTag) PortalType portal,
detail::ThrustIteratorFromArrayPortalTag)
{ {
return vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>(portal); return vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>(portal);
} }
template <typename PortalType> template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin( VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
PortalType portal, detail::ThrustIteratorDevicePtrTag) PortalType portal,
detail::ThrustIteratorDevicePtrTag)
{ {
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal); vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
return iterators.GetBegin(); return iterators.GetBegin();
@ -112,7 +114,8 @@ VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
template <typename PortalType> template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd( VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
PortalType portal, detail::ThrustIteratorFromArrayPortalTag) PortalType portal,
detail::ThrustIteratorFromArrayPortalTag)
{ {
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal); vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
::thrust::advance(iterator, static_cast<std::size_t>(portal.GetNumberOfValues())); ::thrust::advance(iterator, static_cast<std::size_t>(portal.GetNumberOfValues()));
@ -121,7 +124,8 @@ VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
template <typename PortalType> template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd( VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
PortalType portal, detail::ThrustIteratorDevicePtrTag) PortalType portal,
detail::ThrustIteratorDevicePtrTag)
{ {
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal); vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
return iterators.GetEnd(); return iterators.GetEnd();

@ -45,9 +45,9 @@ class ArrayPortalFromIterators;
/// of begin/end iterators to an ArrayPortal interface. /// of begin/end iterators to an ArrayPortal interface.
/// ///
template <class IteratorT> template <class IteratorT>
class ArrayPortalFromIterators< class ArrayPortalFromIterators<IteratorT,
IteratorT, typename std::enable_if< typename std::enable_if<!std::is_const<
!std::is_const<typename std::remove_pointer<IteratorT>::type>::value>::type> typename std::remove_pointer<IteratorT>::type>::value>::type>
{ {
public: public:
typedef typename std::iterator_traits<IteratorT>::value_type ValueType; typedef typename std::iterator_traits<IteratorT>::value_type ValueType;
@ -119,9 +119,9 @@ private:
}; };
template <class IteratorT> template <class IteratorT>
class ArrayPortalFromIterators< class ArrayPortalFromIterators<IteratorT,
IteratorT, typename std::enable_if< typename std::enable_if<std::is_const<
std::is_const<typename std::remove_pointer<IteratorT>::type>::value>::type> typename std::remove_pointer<IteratorT>::type>::value>::type>
{ {
public: public:
typedef typename std::iterator_traits<IteratorT>::value_type ValueType; typedef typename std::iterator_traits<IteratorT>::value_type ValueType;

@ -33,8 +33,10 @@ namespace internal
{ {
template <typename NumIndicesArrayType, typename IndexOffsetArrayType, typename DeviceAdapterTag> template <typename NumIndicesArrayType, typename IndexOffsetArrayType, typename DeviceAdapterTag>
void buildIndexOffsets(const NumIndicesArrayType& numIndices, IndexOffsetArrayType& offsets, void buildIndexOffsets(const NumIndicesArrayType& numIndices,
DeviceAdapterTag, std::true_type) IndexOffsetArrayType& offsets,
DeviceAdapterTag,
std::true_type)
{ {
//We first need to make sure that NumIndices and IndexOffsetArrayType //We first need to make sure that NumIndices and IndexOffsetArrayType
//have the same type so we can call scane exclusive //have the same type so we can call scane exclusive
@ -48,7 +50,9 @@ void buildIndexOffsets(const NumIndicesArrayType& numIndices, IndexOffsetArrayTy
} }
template <typename NumIndicesArrayType, typename IndexOffsetArrayType, typename DeviceAdapterTag> template <typename NumIndicesArrayType, typename IndexOffsetArrayType, typename DeviceAdapterTag>
void buildIndexOffsets(const NumIndicesArrayType&, IndexOffsetArrayType&, DeviceAdapterTag, void buildIndexOffsets(const NumIndicesArrayType&,
IndexOffsetArrayType&,
DeviceAdapterTag,
std::false_type) std::false_type)
{ {
//this is a no-op as the storage for the offsets is an implicit handle //this is a no-op as the storage for the offsets is an implicit handle
@ -59,7 +63,8 @@ void buildIndexOffsets(const NumIndicesArrayType&, IndexOffsetArrayType&, Device
} }
template <typename ArrayHandleIndices, typename ArrayHandleOffsets, typename DeviceAdapterTag> template <typename ArrayHandleIndices, typename ArrayHandleOffsets, typename DeviceAdapterTag>
void buildIndexOffsets(const ArrayHandleIndices& numIndices, ArrayHandleOffsets offsets, void buildIndexOffsets(const ArrayHandleIndices& numIndices,
ArrayHandleOffsets offsets,
DeviceAdapterTag tag) DeviceAdapterTag tag)
{ {
typedef vtkm::cont::internal::IsWriteableArrayHandle<ArrayHandleOffsets, DeviceAdapterTag> typedef vtkm::cont::internal::IsWriteableArrayHandle<ArrayHandleOffsets, DeviceAdapterTag>

@ -110,7 +110,8 @@ private:
CopyKernel<typename InputArrayType::template ExecutionTypes<DeviceAdapterTag>::PortalConst, CopyKernel<typename InputArrayType::template ExecutionTypes<DeviceAdapterTag>::PortalConst,
typename OutputArrayType::template ExecutionTypes<DeviceAdapterTag>::Portal> typename OutputArrayType::template ExecutionTypes<DeviceAdapterTag>::Portal>
kernel(input.PrepareForInput(DeviceAdapterTag()), kernel(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(1, DeviceAdapterTag()), index); output.PrepareForOutput(1, DeviceAdapterTag()),
index);
DerivedAlgorithm::Schedule(kernel, 1); DerivedAlgorithm::Schedule(kernel, 1);
@ -174,7 +175,10 @@ public:
OutputPortalType; OutputPortalType;
OutputPortalType outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag()); OutputPortalType outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag());
CopyIfKernel<InputPortalType, StencilPortalType, IndexPortalType, OutputPortalType, CopyIfKernel<InputPortalType,
StencilPortalType,
IndexPortalType,
OutputPortalType,
UnaryPredicate> UnaryPredicate>
copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate); copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate);
DerivedAlgorithm::Schedule(copyKernel, arrayLength); DerivedAlgorithm::Schedule(copyKernel, arrayLength);
@ -193,7 +197,8 @@ public:
// CopySubRange // CopySubRange
template <typename T, typename U, class CIn, class COut> template <typename T, typename U, class CIn, class COut>
VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::Id inputStartIndex, vtkm::Id numberOfElementsToCopy, vtkm::Id inputStartIndex,
vtkm::Id numberOfElementsToCopy,
vtkm::cont::ArrayHandle<U, COut>& output, vtkm::cont::ArrayHandle<U, COut>& output,
vtkm::Id outputIndex = 0) vtkm::Id outputIndex = 0)
{ {
@ -235,7 +240,9 @@ public:
} }
CopyKernel kernel(input.PrepareForInput(DeviceAdapterTag()), CopyKernel kernel(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForInPlace(DeviceAdapterTag()), inputStartIndex, outputIndex); output.PrepareForInPlace(DeviceAdapterTag()),
inputStartIndex,
outputIndex);
DerivedAlgorithm::Schedule(kernel, numberOfElementsToCopy); DerivedAlgorithm::Schedule(kernel, numberOfElementsToCopy);
return true; return true;
} }
@ -255,7 +262,8 @@ public:
DeviceAdapterTag>::PortalConst, DeviceAdapterTag>::PortalConst,
typename vtkm::cont::ArrayHandle<vtkm::Id, COut>::template ExecutionTypes< typename vtkm::cont::ArrayHandle<vtkm::Id, COut>::template ExecutionTypes<
DeviceAdapterTag>::Portal> DeviceAdapterTag>::Portal>
kernel(input.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()), kernel(input.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag())); output.PrepareForOutput(arraySize, DeviceAdapterTag()));
DerivedAlgorithm::Schedule(kernel, arraySize); DerivedAlgorithm::Schedule(kernel, arraySize);
@ -277,8 +285,10 @@ public:
typename vtkm::cont::ArrayHandle<vtkm::Id, typename vtkm::cont::ArrayHandle<vtkm::Id,
COut>::template ExecutionTypes<DeviceAdapterTag>::Portal, COut>::template ExecutionTypes<DeviceAdapterTag>::Portal,
BinaryCompare> BinaryCompare>
kernel(input.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()), kernel(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag()), binary_compare); values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag()),
binary_compare);
DerivedAlgorithm::Schedule(kernel, arraySize); DerivedAlgorithm::Schedule(kernel, arraySize);
} }
@ -300,7 +310,8 @@ public:
} }
template <typename T, typename U, class CIn, class BinaryFunctor> template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue, VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
//Crazy Idea: //Crazy Idea:
@ -320,8 +331,8 @@ public:
typedef vtkm::cont::ArrayHandleImplicit<U, ReduceKernelType> ReduceHandleType; typedef vtkm::cont::ArrayHandleImplicit<U, ReduceKernelType> ReduceHandleType;
typedef vtkm::cont::ArrayHandle<U, vtkm::cont::StorageTagBasic> TempArrayType; typedef vtkm::cont::ArrayHandle<U, vtkm::cont::StorageTagBasic> TempArrayType;
ReduceKernelType kernel(input.PrepareForInput(DeviceAdapterTag()), initialValue, ReduceKernelType kernel(
binary_functor); input.PrepareForInput(DeviceAdapterTag()), initialValue, binary_functor);
vtkm::Id length = (input.GetNumberOfValues() / 16); vtkm::Id length = (input.GetNumberOfValues() / 16);
length += (input.GetNumberOfValues() % 16 == 0) ? 0 : 1; length += (input.GetNumberOfValues() % 16 == 0) ? 0 : 1;
@ -337,14 +348,16 @@ public:
// Streaming Reduce // Streaming Reduce
template <typename T, typename U, class CIn> template <typename T, typename U, class CIn>
VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks, VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks,
const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue) const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue)
{ {
return DerivedAlgorithm::StreamingReduce(numBlocks, input, initialValue, vtkm::Add()); return DerivedAlgorithm::StreamingReduce(numBlocks, input, initialValue, vtkm::Add());
} }
template <typename T, typename U, class CIn, class BinaryFunctor> template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks, VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks,
const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue, const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
vtkm::Id fullSize = input.GetNumberOfValues(); vtkm::Id fullSize = input.GetNumberOfValues();
@ -360,8 +373,8 @@ public:
numberOfInstances = fullSize - blockSize * block; numberOfInstances = fullSize - blockSize * block;
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>> streamIn = vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>> streamIn =
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>>(input, block, blockSize, vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>>(
numberOfInstances); input, block, blockSize, numberOfInstances);
if (block == 0) if (block == 0)
lastResult = DerivedAlgorithm::Reduce(streamIn, initialValue, binary_functor); lastResult = DerivedAlgorithm::Reduce(streamIn, initialValue, binary_functor);
@ -373,7 +386,12 @@ public:
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------
// Reduce By Key // Reduce By Key
template <typename T, typename U, class KIn, class VIn, class KOut, class VOut, template <typename T,
typename U,
class KIn,
class VIn,
class KOut,
class VOut,
class BinaryFunctor> class BinaryFunctor>
VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -430,8 +448,8 @@ public:
ZipInHandleType scanInput(values, keystate); ZipInHandleType scanInput(values, keystate);
ZipOutHandleType scanOutput(reducedValues, stencil); ZipOutHandleType scanOutput(reducedValues, stencil);
DerivedAlgorithm::ScanInclusive(scanInput, scanOutput, DerivedAlgorithm::ScanInclusive(
ReduceByKeyAdd<BinaryFunctor>(binary_functor)); scanInput, scanOutput, ReduceByKeyAdd<BinaryFunctor>(binary_functor));
//at this point we are done with keystate, so free the memory //at this point we are done with keystate, so free the memory
keystate.ReleaseResources(); keystate.ReleaseResources();
@ -453,7 +471,8 @@ public:
template <typename T, class CIn, class COut, class BinaryFunctor> template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output, vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binaryFunctor, const T& initialValue) BinaryFunctor binaryFunctor,
const T& initialValue)
{ {
typedef vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> TempArrayType; typedef vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> TempArrayType;
typedef vtkm::cont::ArrayHandle<T, COut> OutputArrayType; typedef vtkm::cont::ArrayHandle<T, COut> OutputArrayType;
@ -474,7 +493,9 @@ public:
InclusiveToExclusiveKernel<SrcPortalType, DestPortalType, BinaryFunctor> inclusiveToExclusive( InclusiveToExclusiveKernel<SrcPortalType, DestPortalType, BinaryFunctor> inclusiveToExclusive(
inclusiveScan.PrepareForInput(DeviceAdapterTag()), inclusiveScan.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numValues, DeviceAdapterTag()), binaryFunctor, initialValue); output.PrepareForOutput(numValues, DeviceAdapterTag()),
binaryFunctor,
initialValue);
DerivedAlgorithm::Schedule(inclusiveToExclusive, numValues); DerivedAlgorithm::Schedule(inclusiveToExclusive, numValues);
@ -485,8 +506,8 @@ public:
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output) vtkm::cont::ArrayHandle<T, COut>& output)
{ {
return DerivedAlgorithm::ScanExclusive(input, output, vtkm::Sum(), return DerivedAlgorithm::ScanExclusive(
vtkm::TypeTraits<T>::ZeroInitialization()); input, output, vtkm::Sum(), vtkm::TypeTraits<T>::ZeroInitialization());
} }
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------
@ -495,7 +516,8 @@ public:
VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output, vtkm::cont::ArrayHandle<U, VOut>& output,
const U& initialValue, BinaryFunctor binaryFunctor) const U& initialValue,
BinaryFunctor binaryFunctor)
{ {
VTKM_ASSERT(keys.GetNumberOfValues() == values.GetNumberOfValues()); VTKM_ASSERT(keys.GetNumberOfValues() == values.GetNumberOfValues());
@ -564,8 +586,8 @@ public:
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output) vtkm::cont::ArrayHandle<U, VOut>& output)
{ {
DerivedAlgorithm::ScanExclusiveByKey(keys, values, output, DerivedAlgorithm::ScanExclusiveByKey(
vtkm::TypeTraits<U>::ZeroInitialization(), vtkm::Sum()); keys, values, output, vtkm::TypeTraits<U>::ZeroInitialization(), vtkm::Sum());
} }
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------
@ -575,15 +597,16 @@ public:
const vtkm::cont::ArrayHandle<T, CIn>& input, const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output) vtkm::cont::ArrayHandle<T, COut>& output)
{ {
return DerivedAlgorithm::StreamingScanExclusive(numBlocks, input, output, vtkm::Sum(), return DerivedAlgorithm::StreamingScanExclusive(
vtkm::TypeTraits<T>::ZeroInitialization()); numBlocks, input, output, vtkm::Sum(), vtkm::TypeTraits<T>::ZeroInitialization());
} }
template <typename T, class CIn, class COut, class BinaryFunctor> template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T StreamingScanExclusive(const vtkm::Id numBlocks, VTKM_CONT static T StreamingScanExclusive(const vtkm::Id numBlocks,
const vtkm::cont::ArrayHandle<T, CIn>& input, const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output, vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binary_functor, const T& initialValue) BinaryFunctor binary_functor,
const T& initialValue)
{ {
vtkm::Id fullSize = input.GetNumberOfValues(); vtkm::Id fullSize = input.GetNumberOfValues();
vtkm::Id blockSize = fullSize / numBlocks; vtkm::Id blockSize = fullSize / numBlocks;
@ -598,12 +621,12 @@ public:
numberOfInstances = fullSize - blockSize * block; numberOfInstances = fullSize - blockSize * block;
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>> streamIn = vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>> streamIn =
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>>(input, block, blockSize, vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, CIn>>(
numberOfInstances); input, block, blockSize, numberOfInstances);
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, COut>> streamOut = vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, COut>> streamOut =
vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, COut>>(output, block, blockSize, vtkm::cont::ArrayHandleStreaming<vtkm::cont::ArrayHandle<T, COut>>(
numberOfInstances); output, block, blockSize, numberOfInstances);
if (block == 0) if (block == 0)
{ {
@ -733,8 +756,8 @@ public:
ZipInHandleType scanInput(values, keystate); ZipInHandleType scanInput(values, keystate);
ZipOutHandleType scanOutput(reducedValues, stencil); ZipOutHandleType scanOutput(reducedValues, stencil);
DerivedAlgorithm::ScanInclusive(scanInput, scanOutput, DerivedAlgorithm::ScanInclusive(
ReduceByKeyAdd<BinaryFunctor>(binary_functor)); scanInput, scanOutput, ReduceByKeyAdd<BinaryFunctor>(binary_functor));
//at this point we are done with keystate, so free the memory //at this point we are done with keystate, so free the memory
keystate.ReleaseResources(); keystate.ReleaseResources();
DerivedAlgorithm::Copy(reducedValues, values_output); DerivedAlgorithm::Copy(reducedValues, values_output);
@ -841,11 +864,12 @@ public:
ClassifyUniqueComparisonKernel< ClassifyUniqueComparisonKernel<
typename vtkm::cont::ArrayHandle<T, Storage>::template ExecutionTypes< typename vtkm::cont::ArrayHandle<T, Storage>::template ExecutionTypes<
DeviceAdapterTag>::PortalConst, DeviceAdapterTag>::PortalConst,
typename vtkm::cont::ArrayHandle< typename vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic>::
vtkm::Id, vtkm::cont::StorageTagBasic>::template ExecutionTypes<DeviceAdapterTag>::Portal, template ExecutionTypes<DeviceAdapterTag>::Portal,
WrappedBOpType> WrappedBOpType>
classifyKernel(values.PrepareForInput(DeviceAdapterTag()), classifyKernel(values.PrepareForInput(DeviceAdapterTag()),
stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag()), wrappedCompare); stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag()),
wrappedCompare);
DerivedAlgorithm::Schedule(classifyKernel, inputSize); DerivedAlgorithm::Schedule(classifyKernel, inputSize);
@ -872,7 +896,8 @@ public:
DeviceAdapterTag>::PortalConst, DeviceAdapterTag>::PortalConst,
typename vtkm::cont::ArrayHandle<vtkm::Id, COut>::template ExecutionTypes< typename vtkm::cont::ArrayHandle<vtkm::Id, COut>::template ExecutionTypes<
DeviceAdapterTag>::Portal> DeviceAdapterTag>::Portal>
kernel(input.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()), kernel(input.PrepareForInput(DeviceAdapterTag()),
values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag())); output.PrepareForOutput(arraySize, DeviceAdapterTag()));
DerivedAlgorithm::Schedule(kernel, arraySize); DerivedAlgorithm::Schedule(kernel, arraySize);
@ -894,8 +919,10 @@ public:
typename vtkm::cont::ArrayHandle<vtkm::Id, typename vtkm::cont::ArrayHandle<vtkm::Id,
COut>::template ExecutionTypes<DeviceAdapterTag>::Portal, COut>::template ExecutionTypes<DeviceAdapterTag>::Portal,
BinaryCompare> BinaryCompare>
kernel(input.PrepareForInput(DeviceAdapterTag()), values.PrepareForInput(DeviceAdapterTag()), kernel(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag()), binary_compare); values.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(arraySize, DeviceAdapterTag()),
binary_compare);
DerivedAlgorithm::Schedule(kernel, arraySize); DerivedAlgorithm::Schedule(kernel, arraySize);
} }
@ -987,19 +1014,21 @@ private:
} }
VTKM_EXEC VTKM_EXEC
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address, const vtkm::Int32& newValue, vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
const vtkm::Int32& newValue,
const vtkm::Int32& oldValue) const const vtkm::Int32& oldValue) const
{ {
return InterlockedCompareExchange(reinterpret_cast<volatile long*>(address), newValue, return InterlockedCompareExchange(
oldValue); reinterpret_cast<volatile long*>(address), newValue, oldValue);
} }
VTKM_EXEC VTKM_EXEC
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address, const vtkm::Int64& newValue, vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
const vtkm::Int64& newValue,
const vtkm::Int64& oldValue) const const vtkm::Int64& oldValue) const
{ {
return InterlockedCompareExchange64(reinterpret_cast<volatile long long*>(address), newValue, return InterlockedCompareExchange64(
oldValue); reinterpret_cast<volatile long long*>(address), newValue, oldValue);
} }
#else //gcc built-in atomics #else //gcc built-in atomics
@ -1017,14 +1046,16 @@ private:
} }
VTKM_EXEC VTKM_EXEC
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address, const vtkm::Int32& newValue, vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
const vtkm::Int32& newValue,
const vtkm::Int32& oldValue) const const vtkm::Int32& oldValue) const
{ {
return __sync_val_compare_and_swap(address, oldValue, newValue); return __sync_val_compare_and_swap(address, oldValue, newValue);
} }
VTKM_EXEC VTKM_EXEC
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address, const vtkm::Int64& newValue, vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
const vtkm::Int64& newValue,
const vtkm::Int64& oldValue) const const vtkm::Int64& oldValue) const
{ {
return __sync_val_compare_and_swap(address, oldValue, newValue); return __sync_val_compare_and_swap(address, oldValue, newValue);
@ -1049,7 +1080,9 @@ class DeviceTaskTypes
public: public:
template <typename WorkletType, typename InvocationType> template <typename WorkletType, typename InvocationType>
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask( static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id, const WorkletType& worklet,
const InvocationType& invocation,
vtkm::Id,
vtkm::Id globalIndexOffset = 0) vtkm::Id globalIndexOffset = 0)
{ {
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>; using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
@ -1058,7 +1091,9 @@ public:
template <typename WorkletType, typename InvocationType> template <typename WorkletType, typename InvocationType>
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask( static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id3, const WorkletType& worklet,
const InvocationType& invocation,
vtkm::Id3,
vtkm::Id globalIndexOffset = 0) vtkm::Id globalIndexOffset = 0)
{ {
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>; using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;

@ -139,24 +139,28 @@ struct DynamicTransformTraits
struct DynamicTransform struct DynamicTransform
{ {
template <typename InputType, typename ContinueFunctor, vtkm::IdComponent Index> template <typename InputType, typename ContinueFunctor, vtkm::IdComponent Index>
VTKM_CONT void operator()(const InputType& input, const ContinueFunctor& continueFunc, VTKM_CONT void operator()(const InputType& input,
const ContinueFunctor& continueFunc,
vtkm::internal::IndexTag<Index>) const vtkm::internal::IndexTag<Index>) const
{ {
this->DoTransform( this->DoTransform(
input, continueFunc, input,
continueFunc,
typename vtkm::cont::internal::DynamicTransformTraits<InputType>::DynamicTag()); typename vtkm::cont::internal::DynamicTransformTraits<InputType>::DynamicTag());
} }
private: private:
template <typename InputType, typename ContinueFunctor> template <typename InputType, typename ContinueFunctor>
VTKM_CONT void DoTransform(const InputType& input, const ContinueFunctor& continueFunc, VTKM_CONT void DoTransform(const InputType& input,
const ContinueFunctor& continueFunc,
vtkm::cont::internal::DynamicTransformTagStatic) const vtkm::cont::internal::DynamicTransformTagStatic) const
{ {
continueFunc(input); continueFunc(input);
} }
template <typename InputType, typename ContinueFunctor> template <typename InputType, typename ContinueFunctor>
VTKM_CONT void DoTransform(const InputType& dynamicInput, const ContinueFunctor& continueFunc, VTKM_CONT void DoTransform(const InputType& dynamicInput,
const ContinueFunctor& continueFunc,
vtkm::cont::internal::DynamicTransformTagCastAndCall) const vtkm::cont::internal::DynamicTransformTagCastAndCall) const
{ {
CastAndCall(dynamicInput, continueFunc); CastAndCall(dynamicInput, continueFunc);

@ -68,8 +68,9 @@ struct WrappedBinaryOperator
} }
template <typename Argument1, typename Argument2> template <typename Argument1, typename Argument2>
VTKM_CONT ResultType operator()( VTKM_CONT ResultType
const Argument1& x, const vtkm::internal::ArrayPortalValueReference<Argument2>& y) const operator()(const Argument1& x,
const vtkm::internal::ArrayPortalValueReference<Argument2>& y) const
{ {
using ValueTypeY = typename vtkm::internal::ArrayPortalValueReference<Argument2>::ValueType; using ValueTypeY = typename vtkm::internal::ArrayPortalValueReference<Argument2>::ValueType;
return m_f(x, (ValueTypeY)y); return m_f(x, (ValueTypeY)y);
@ -296,7 +297,9 @@ struct ReduceByKeyUnaryStencilOp
bool operator()(ReduceKeySeriesStates keySeriesState) const { return keySeriesState.fEnd; } bool operator()(ReduceKeySeriesStates keySeriesState) const { return keySeriesState.fEnd; }
}; };
template <typename T, typename InputPortalType, typename KeyStatePortalType, template <typename T,
typename InputPortalType,
typename KeyStatePortalType,
typename OutputPortalType> typename OutputPortalType>
struct ShiftCopyAndInit : vtkm::exec::FunctorBase struct ShiftCopyAndInit : vtkm::exec::FunctorBase
{ {
@ -305,8 +308,10 @@ struct ShiftCopyAndInit : vtkm::exec::FunctorBase
OutputPortalType Output; OutputPortalType Output;
T initValue; T initValue;
ShiftCopyAndInit(const InputPortalType& _input, const KeyStatePortalType& kstate, ShiftCopyAndInit(const InputPortalType& _input,
OutputPortalType& _output, T _init) const KeyStatePortalType& kstate,
OutputPortalType& _output,
T _init)
: Input(_input) : Input(_input)
, KeyState(kstate) , KeyState(kstate)
, Output(_output) , Output(_output)
@ -336,7 +341,9 @@ struct CopyKernel
vtkm::Id OutputOffset; vtkm::Id OutputOffset;
VTKM_CONT VTKM_CONT
CopyKernel(InputPortalType inputPortal, OutputPortalType outputPortal, vtkm::Id inputOffset = 0, CopyKernel(InputPortalType inputPortal,
OutputPortalType outputPortal,
vtkm::Id inputOffset = 0,
vtkm::Id outputOffset = 0) vtkm::Id outputOffset = 0)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, OutputPortal(outputPortal) , OutputPortal(outputPortal)
@ -350,8 +357,9 @@ struct CopyKernel
void operator()(vtkm::Id index) const void operator()(vtkm::Id index) const
{ {
typedef typename OutputPortalType::ValueType ValueType; typedef typename OutputPortalType::ValueType ValueType;
this->OutputPortal.Set(index + this->OutputOffset, static_cast<ValueType>(this->InputPortal.Get( this->OutputPortal.Set(
index + this->InputOffset))); index + this->OutputOffset,
static_cast<ValueType>(this->InputPortal.Get(index + this->InputOffset)));
} }
VTKM_CONT VTKM_CONT
@ -366,7 +374,8 @@ struct LowerBoundsKernel
OutputPortalType OutputPortal; OutputPortalType OutputPortal;
VTKM_CONT VTKM_CONT
LowerBoundsKernel(InputPortalType inputPortal, ValuesPortalType valuesPortal, LowerBoundsKernel(InputPortalType inputPortal,
ValuesPortalType valuesPortal,
OutputPortalType outputPortal) OutputPortalType outputPortal)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, ValuesPortal(valuesPortal) , ValuesPortal(valuesPortal)
@ -399,7 +408,9 @@ struct LowerBoundsKernel
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {} void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {}
}; };
template <class InputPortalType, class ValuesPortalType, class OutputPortalType, template <class InputPortalType,
class ValuesPortalType,
class OutputPortalType,
class BinaryCompare> class BinaryCompare>
struct LowerBoundsComparisonKernel struct LowerBoundsComparisonKernel
{ {
@ -409,8 +420,10 @@ struct LowerBoundsComparisonKernel
BinaryCompare CompareFunctor; BinaryCompare CompareFunctor;
VTKM_CONT VTKM_CONT
LowerBoundsComparisonKernel(InputPortalType inputPortal, ValuesPortalType valuesPortal, LowerBoundsComparisonKernel(InputPortalType inputPortal,
OutputPortalType outputPortal, BinaryCompare binary_compare) ValuesPortalType valuesPortal,
OutputPortalType outputPortal,
BinaryCompare binary_compare)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, ValuesPortal(valuesPortal) , ValuesPortal(valuesPortal)
, OutputPortal(outputPortal) , OutputPortal(outputPortal)
@ -432,8 +445,10 @@ struct LowerBoundsComparisonKernel
typedef vtkm::cont::ArrayPortalToIterators<InputPortalType> InputIteratorsType; typedef vtkm::cont::ArrayPortalToIterators<InputPortalType> InputIteratorsType;
InputIteratorsType inputIterators(this->InputPortal); InputIteratorsType inputIterators(this->InputPortal);
typename InputIteratorsType::IteratorType resultPos = typename InputIteratorsType::IteratorType resultPos =
std::lower_bound(inputIterators.GetBegin(), inputIterators.GetEnd(), std::lower_bound(inputIterators.GetBegin(),
this->ValuesPortal.Get(index), this->CompareFunctor); inputIterators.GetEnd(),
this->ValuesPortal.Get(index),
this->CompareFunctor);
vtkm::Id resultIndex = vtkm::Id resultIndex =
static_cast<vtkm::Id>(std::distance(inputIterators.GetBegin(), resultPos)); static_cast<vtkm::Id>(std::distance(inputIterators.GetBegin(), resultPos));
@ -515,7 +530,8 @@ struct BitonicSortCrossoverKernel : vtkm::exec::FunctorBase
vtkm::Id GroupSize; vtkm::Id GroupSize;
VTKM_CONT VTKM_CONT
BitonicSortCrossoverKernel(const PortalType& portal, const BinaryCompare& compare, BitonicSortCrossoverKernel(const PortalType& portal,
const BinaryCompare& compare,
vtkm::Id groupSize) vtkm::Id groupSize)
: Portal(portal) : Portal(portal)
, Compare(compare) , Compare(compare)
@ -558,7 +574,8 @@ struct StencilToIndexFlagKernel
UnaryPredicate Predicate; UnaryPredicate Predicate;
VTKM_CONT VTKM_CONT
StencilToIndexFlagKernel(StencilPortalType stencilPortal, OutputPortalType outputPortal, StencilToIndexFlagKernel(StencilPortalType stencilPortal,
OutputPortalType outputPortal,
UnaryPredicate unary_predicate) UnaryPredicate unary_predicate)
: StencilPortal(stencilPortal) : StencilPortal(stencilPortal)
, OutputPortal(outputPortal) , OutputPortal(outputPortal)
@ -578,8 +595,11 @@ struct StencilToIndexFlagKernel
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {} void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {}
}; };
template <class InputPortalType, class StencilPortalType, class IndexPortalType, template <class InputPortalType,
class OutputPortalType, class PredicateOperator> class StencilPortalType,
class IndexPortalType,
class OutputPortalType,
class PredicateOperator>
struct CopyIfKernel struct CopyIfKernel
{ {
InputPortalType InputPortal; InputPortalType InputPortal;
@ -589,8 +609,10 @@ struct CopyIfKernel
PredicateOperator Predicate; PredicateOperator Predicate;
VTKM_CONT VTKM_CONT
CopyIfKernel(InputPortalType inputPortal, StencilPortalType stencilPortal, CopyIfKernel(InputPortalType inputPortal,
IndexPortalType indexPortal, OutputPortalType outputPortal, StencilPortalType stencilPortal,
IndexPortalType indexPortal,
OutputPortalType outputPortal,
PredicateOperator unary_predicate) PredicateOperator unary_predicate)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, StencilPortal(stencilPortal) , StencilPortal(stencilPortal)
@ -663,7 +685,8 @@ struct ClassifyUniqueComparisonKernel
BinaryCompare CompareFunctor; BinaryCompare CompareFunctor;
VTKM_CONT VTKM_CONT
ClassifyUniqueComparisonKernel(InputPortalType inputPortal, StencilPortalType stencilPortal, ClassifyUniqueComparisonKernel(InputPortalType inputPortal,
StencilPortalType stencilPortal,
BinaryCompare binary_compare) BinaryCompare binary_compare)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, StencilPortal(stencilPortal) , StencilPortal(stencilPortal)
@ -703,7 +726,8 @@ struct UpperBoundsKernel
OutputPortalType OutputPortal; OutputPortalType OutputPortal;
VTKM_CONT VTKM_CONT
UpperBoundsKernel(InputPortalType inputPortal, ValuesPortalType valuesPortal, UpperBoundsKernel(InputPortalType inputPortal,
ValuesPortalType valuesPortal,
OutputPortalType outputPortal) OutputPortalType outputPortal)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, ValuesPortal(valuesPortal) , ValuesPortal(valuesPortal)
@ -736,7 +760,9 @@ struct UpperBoundsKernel
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {} void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {}
}; };
template <class InputPortalType, class ValuesPortalType, class OutputPortalType, template <class InputPortalType,
class ValuesPortalType,
class OutputPortalType,
class BinaryCompare> class BinaryCompare>
struct UpperBoundsKernelComparisonKernel struct UpperBoundsKernelComparisonKernel
{ {
@ -746,8 +772,10 @@ struct UpperBoundsKernelComparisonKernel
BinaryCompare CompareFunctor; BinaryCompare CompareFunctor;
VTKM_CONT VTKM_CONT
UpperBoundsKernelComparisonKernel(InputPortalType inputPortal, ValuesPortalType valuesPortal, UpperBoundsKernelComparisonKernel(InputPortalType inputPortal,
OutputPortalType outputPortal, BinaryCompare binary_compare) ValuesPortalType valuesPortal,
OutputPortalType outputPortal,
BinaryCompare binary_compare)
: InputPortal(inputPortal) : InputPortal(inputPortal)
, ValuesPortal(valuesPortal) , ValuesPortal(valuesPortal)
, OutputPortal(outputPortal) , OutputPortal(outputPortal)
@ -769,8 +797,10 @@ struct UpperBoundsKernelComparisonKernel
typedef vtkm::cont::ArrayPortalToIterators<InputPortalType> InputIteratorsType; typedef vtkm::cont::ArrayPortalToIterators<InputPortalType> InputIteratorsType;
InputIteratorsType inputIterators(this->InputPortal); InputIteratorsType inputIterators(this->InputPortal);
typename InputIteratorsType::IteratorType resultPos = typename InputIteratorsType::IteratorType resultPos =
std::upper_bound(inputIterators.GetBegin(), inputIterators.GetEnd(), std::upper_bound(inputIterators.GetBegin(),
this->ValuesPortal.Get(index), this->CompareFunctor); inputIterators.GetEnd(),
this->ValuesPortal.Get(index),
this->CompareFunctor);
vtkm::Id resultIndex = vtkm::Id resultIndex =
static_cast<vtkm::Id>(std::distance(inputIterators.GetBegin(), resultPos)); static_cast<vtkm::Id>(std::distance(inputIterators.GetBegin(), resultPos));
@ -792,8 +822,10 @@ struct InclusiveToExclusiveKernel : vtkm::exec::FunctorBase
ValueType InitialValue; ValueType InitialValue;
VTKM_CONT VTKM_CONT
InclusiveToExclusiveKernel(const InPortalType& inPortal, const OutPortalType& outPortal, InclusiveToExclusiveKernel(const InPortalType& inPortal,
BinaryFunctor& binaryOperator, ValueType initialValue) const OutPortalType& outPortal,
BinaryFunctor& binaryOperator,
ValueType initialValue)
: InPortal(inPortal) : InPortal(inPortal)
, OutPortal(outPortal) , OutPortal(outPortal)
, BinaryOperator(binaryOperator) , BinaryOperator(binaryOperator)
@ -822,7 +854,9 @@ struct ScanKernel : vtkm::exec::FunctorBase
vtkm::Id Distance; vtkm::Id Distance;
VTKM_CONT VTKM_CONT
ScanKernel(const PortalType& portal, BinaryFunctor binary_functor, vtkm::Id stride, ScanKernel(const PortalType& portal,
BinaryFunctor binary_functor,
vtkm::Id stride,
vtkm::Id offset) vtkm::Id offset)
: Portal(portal) : Portal(portal)
, BinaryOperator(binary_functor) , BinaryOperator(binary_functor)

@ -172,18 +172,20 @@ std::ptrdiff_t operator-(vtkm::cont::internal::IteratorFromArrayPortal<PortalTyp
template <typename PortalType> template <typename PortalType>
vtkm::cont::internal::IteratorFromArrayPortal<PortalType> operator+( vtkm::cont::internal::IteratorFromArrayPortal<PortalType> operator+(
vtkm::cont::internal::IteratorFromArrayPortal<PortalType> const& iter, std::ptrdiff_t n) vtkm::cont::internal::IteratorFromArrayPortal<PortalType> const& iter,
std::ptrdiff_t n)
{ {
return vtkm::cont::internal::IteratorFromArrayPortal<PortalType>(iter.Portal, iter.Index + return vtkm::cont::internal::IteratorFromArrayPortal<PortalType>(
static_cast<vtkm::Id>(n)); iter.Portal, iter.Index + static_cast<vtkm::Id>(n));
} }
template <typename PortalType> template <typename PortalType>
vtkm::cont::internal::IteratorFromArrayPortal<PortalType> operator+( vtkm::cont::internal::IteratorFromArrayPortal<PortalType> operator+(
std::ptrdiff_t n, vtkm::cont::internal::IteratorFromArrayPortal<PortalType> const& iter) std::ptrdiff_t n,
vtkm::cont::internal::IteratorFromArrayPortal<PortalType> const& iter)
{ {
return vtkm::cont::internal::IteratorFromArrayPortal<PortalType>(iter.Portal, iter.Index + return vtkm::cont::internal::IteratorFromArrayPortal<PortalType>(
static_cast<vtkm::Id>(n)); iter.Portal, iter.Index + static_cast<vtkm::Id>(n));
} }
} }
} }

@ -80,8 +80,8 @@ struct TemplatedTests
FillIterator(array, array + ARRAY_SIZE, ORIGINAL_VALUE); FillIterator(array, array + ARRAY_SIZE, ORIGINAL_VALUE);
::vtkm::cont::internal::ArrayPortalFromIterators<ValueType*> portal(array, array + ARRAY_SIZE); ::vtkm::cont::internal::ArrayPortalFromIterators<ValueType*> portal(array, array + ARRAY_SIZE);
::vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*> const_portal(array, array + ::vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*> const_portal(
ARRAY_SIZE); array, array + ARRAY_SIZE);
std::cout << " Check that ArrayPortalToIterators is not doing indirection." << std::endl; std::cout << " Check that ArrayPortalToIterators is not doing indirection." << std::endl;
// If you get a compile error here about mismatched types, it might be // If you get a compile error here about mismatched types, it might be

@ -113,9 +113,10 @@ struct FunctionInterfaceFunctor
VTKM_TEST_FAIL("Called wrong form of functor operator."); VTKM_TEST_FAIL("Called wrong form of functor operator.");
} }
void operator()( void operator()(
const vtkm::internal::FunctionInterface< const vtkm::internal::FunctionInterface<void(vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
void(vtkm::cont::ArrayHandle<vtkm::FloatDefault>, vtkm::cont::ArrayHandle<vtkm::FloatDefault>, vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<std::string>, vtkm::cont::CellSetStructured<3>)>&) const vtkm::cont::ArrayHandle<std::string>,
vtkm::cont::CellSetStructured<3>)>&) const
{ {
std::cout << " In FunctionInterface<...> functor." << std::endl; std::cout << " In FunctionInterface<...> functor." << std::endl;
g_FunctionCalls++; g_FunctionCalls++;
@ -142,8 +143,8 @@ void TestBasicTransform()
std::cout << " Trying with unusual (string) dynamic array." << std::endl; std::cout << " Trying with unusual (string) dynamic array." << std::endl;
dynamicArray = vtkm::cont::ArrayHandle<std::string>(); dynamicArray = vtkm::cont::ArrayHandle<std::string>();
TRY_TRANSFORM(transform(dynamicArray.ResetTypeList(TypeListTagString()), TRY_TRANSFORM(transform(
ArrayHandleStringFunctor(), indexTag)); dynamicArray.ResetTypeList(TypeListTagString()), ArrayHandleStringFunctor(), indexTag));
std::cout << " Trying with structured cell set." << std::endl; std::cout << " Trying with structured cell set." << std::endl;
vtkm::cont::CellSetStructured<3> concreteCellSet; vtkm::cont::CellSetStructured<3> concreteCellSet;
@ -169,7 +170,8 @@ void TestFunctionTransform()
std::cout << " Trying dynamic cast" << std::endl; std::cout << " Trying dynamic cast" << std::endl;
TRY_TRANSFORM( TRY_TRANSFORM(
vtkm::internal::make_FunctionInterface<void>( vtkm::internal::make_FunctionInterface<void>(
scalarArray, vtkm::cont::DynamicArrayHandle(scalarArray), scalarArray,
vtkm::cont::DynamicArrayHandle(scalarArray),
vtkm::cont::DynamicArrayHandle(stringArray).ResetTypeList(TypeListTagString()), vtkm::cont::DynamicArrayHandle(stringArray).ResetTypeList(TypeListTagString()),
vtkm::cont::DynamicCellSet(structuredCellSet)) vtkm::cont::DynamicCellSet(structuredCellSet))
.DynamicTransformCont(vtkm::cont::internal::DynamicTransform(), FunctionInterfaceFunctor())); .DynamicTransformCont(vtkm::cont::internal::DynamicTransform(), FunctionInterfaceFunctor()));

@ -132,8 +132,8 @@ struct TemplatedTests
FillIterator(array, array + ARRAY_SIZE, ORIGINAL_VALUE()); FillIterator(array, array + ARRAY_SIZE, ORIGINAL_VALUE());
::vtkm::cont::internal::ArrayPortalFromIterators<ValueType*> portal(array, array + ARRAY_SIZE); ::vtkm::cont::internal::ArrayPortalFromIterators<ValueType*> portal(array, array + ARRAY_SIZE);
::vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*> const_portal(array, array + ::vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*> const_portal(
ARRAY_SIZE); array, array + ARRAY_SIZE);
std::cout << " Test read from iterator." << std::endl; std::cout << " Test read from iterator." << std::endl;
TestIteratorRead(portal); TestIteratorRead(portal);

@ -26,7 +26,8 @@ namespace cont
{ {
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask( void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
vtkm::exec::serial::internal::TaskTiling1D& functor, vtkm::Id size) vtkm::exec::serial::internal::TaskTiling1D& functor,
vtkm::Id size)
{ {
const vtkm::Id MESSAGE_SIZE = 1024; const vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE]; char errorString[MESSAGE_SIZE];
@ -50,7 +51,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
} }
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask( void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
vtkm::exec::serial::internal::TaskTiling3D& functor, vtkm::Id3 size) vtkm::exec::serial::internal::TaskTiling3D& functor,
vtkm::Id3 size)
{ {
const vtkm::Id MESSAGE_SIZE = 1024; const vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE]; char errorString[MESSAGE_SIZE];

@ -57,7 +57,8 @@ public:
} }
template <typename T, typename U, class CIn, class BinaryFunctor> template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue, VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
typedef typename vtkm::cont::ArrayHandle<T, CIn>::template ExecutionTypes<Device>::PortalConst typedef typename vtkm::cont::ArrayHandle<T, CIn>::template ExecutionTypes<Device>::PortalConst
@ -66,11 +67,17 @@ public:
internal::WrappedBinaryOperator<U, BinaryFunctor> wrappedOp(binary_functor); internal::WrappedBinaryOperator<U, BinaryFunctor> wrappedOp(binary_functor);
PortalIn inputPortal = input.PrepareForInput(Device()); PortalIn inputPortal = input.PrepareForInput(Device());
return std::accumulate(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal), return std::accumulate(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal),
vtkm::cont::ArrayPortalToIteratorEnd(inputPortal), initialValue, vtkm::cont::ArrayPortalToIteratorEnd(inputPortal),
initialValue,
wrappedOp); wrappedOp);
} }
template <typename T, typename U, class KIn, class VIn, class KOut, class VOut, template <typename T,
typename U,
class KIn,
class VIn,
class KOut,
class VOut,
class BinaryFunctor> class BinaryFunctor>
VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys, VTKM_CONT static void ReduceByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values, const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -181,7 +188,8 @@ public:
std::partial_sum(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal), std::partial_sum(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal),
vtkm::cont::ArrayPortalToIteratorEnd(inputPortal), vtkm::cont::ArrayPortalToIteratorEnd(inputPortal),
vtkm::cont::ArrayPortalToIteratorBegin(outputPortal), wrappedBinaryOp); vtkm::cont::ArrayPortalToIteratorBegin(outputPortal),
wrappedBinaryOp);
// Return the value at the last index in the array, which is the full sum. // Return the value at the last index in the array, which is the full sum.
return outputPortal.Get(numberOfValues - 1); return outputPortal.Get(numberOfValues - 1);
@ -190,7 +198,8 @@ public:
template <typename T, class CIn, class COut, class BinaryFunctor> template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output, vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binaryFunctor, const T& initialValue) BinaryFunctor binaryFunctor,
const T& initialValue)
{ {
typedef typedef
typename vtkm::cont::ArrayHandle<T, COut>::template ExecutionTypes<Device>::Portal PortalOut; typename vtkm::cont::ArrayHandle<T, COut>::template ExecutionTypes<Device>::Portal PortalOut;
@ -225,7 +234,8 @@ public:
std::partial_sum(vtkm::cont::ArrayPortalToIteratorBegin(outputPortal), std::partial_sum(vtkm::cont::ArrayPortalToIteratorBegin(outputPortal),
vtkm::cont::ArrayPortalToIteratorEnd(outputPortal), vtkm::cont::ArrayPortalToIteratorEnd(outputPortal),
vtkm::cont::ArrayPortalToIteratorBegin(outputPortal), wrappedBinaryOp); vtkm::cont::ArrayPortalToIteratorBegin(outputPortal),
wrappedBinaryOp);
return wrappedBinaryOp(outputPortal.Get(numberOfValues - 1), lastValue); return wrappedBinaryOp(outputPortal.Get(numberOfValues - 1), lastValue);
} }
@ -257,7 +267,11 @@ public:
} }
private: private:
template <typename Vin, typename I, typename Vout, class StorageVin, class StorageI, template <typename Vin,
typename I,
typename Vout,
class StorageVin,
class StorageI,
class StorageVout> class StorageVout>
VTKM_CONT static void Scatter(vtkm::cont::ArrayHandle<Vin, StorageVin>& values, VTKM_CONT static void Scatter(vtkm::cont::ArrayHandle<Vin, StorageVin>& values,
vtkm::cont::ArrayHandle<I, StorageI>& index, vtkm::cont::ArrayHandle<I, StorageI>& index,

@ -26,7 +26,8 @@ namespace cont
{ {
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask( void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
vtkm::exec::tbb::internal::TaskTiling1D& functor, vtkm::Id size) vtkm::exec::tbb::internal::TaskTiling1D& functor,
vtkm::Id size)
{ {
const vtkm::Id MESSAGE_SIZE = 1024; const vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE]; char errorString[MESSAGE_SIZE];
@ -46,7 +47,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
} }
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask( void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
vtkm::exec::tbb::internal::TaskTiling3D& functor, vtkm::Id3 size) vtkm::exec::tbb::internal::TaskTiling3D& functor,
vtkm::Id3 size)
{ {
static const vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 }; static const vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 };
const vtkm::Id MESSAGE_SIZE = 1024; const vtkm::Id MESSAGE_SIZE = 1024;
@ -57,8 +59,15 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
//memory is generally setup in a way that iterating the first range //memory is generally setup in a way that iterating the first range
//in the tightest loop has the best cache coherence. //in the tightest loop has the best cache coherence.
::tbb::blocked_range3d<vtkm::Id> range(0, size[2], TBB_GRAIN_SIZE_3D[0], 0, size[1], ::tbb::blocked_range3d<vtkm::Id> range(0,
TBB_GRAIN_SIZE_3D[1], 0, size[0], TBB_GRAIN_SIZE_3D[2]); size[2],
TBB_GRAIN_SIZE_3D[0],
0,
size[1],
TBB_GRAIN_SIZE_3D[1],
0,
size[0],
TBB_GRAIN_SIZE_3D[2]);
::tbb::parallel_for(range, [&](const ::tbb::blocked_range3d<vtkm::Id>& r) { ::tbb::parallel_for(range, [&](const ::tbb::blocked_range3d<vtkm::Id>& r) {
for (vtkm::Id k = r.pages().begin(); k != r.pages().end(); ++k) for (vtkm::Id k = r.pages().begin(); k != r.pages().end(); ++k)
{ {

@ -41,7 +41,8 @@ namespace cont
template <> template <>
struct DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB> struct DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>
: vtkm::cont::internal::DeviceAdapterAlgorithmGeneral< : vtkm::cont::internal::DeviceAdapterAlgorithmGeneral<
DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>, vtkm::cont::DeviceAdapterTagTBB> DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>,
vtkm::cont::DeviceAdapterTagTBB>
{ {
public: public:
template <typename T, typename U, class CIn> template <typename T, typename U, class CIn>
@ -51,11 +52,12 @@ public:
} }
template <typename T, typename U, class CIn, class BinaryFunctor> template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue, VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor) BinaryFunctor binary_functor)
{ {
return tbb::ReducePortals(input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), return tbb::ReducePortals(
initialValue, binary_functor); input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), initialValue, binary_functor);
} }
template <typename T, class CIn, class COut> template <typename T, class CIn, class COut>
@ -86,18 +88,21 @@ public:
return tbb::ScanExclusivePortals( return tbb::ScanExclusivePortals(
input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB()), output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB()),
vtkm::Add(), vtkm::TypeTraits<T>::ZeroInitialization()); vtkm::Add(),
vtkm::TypeTraits<T>::ZeroInitialization());
} }
template <typename T, class CIn, class COut, class BinaryFunctor> template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input, VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output, vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binary_functor, const T& initialValue) BinaryFunctor binary_functor,
const T& initialValue)
{ {
return tbb::ScanExclusivePortals( return tbb::ScanExclusivePortals(
input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB()), output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB()),
binary_functor, initialValue); binary_functor,
initialValue);
} }
VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::tbb::internal::TaskTiling1D& functor, VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::tbb::internal::TaskTiling1D& functor,
@ -151,7 +156,8 @@ public:
template <typename T, typename U, class StorageT, class StorageU, class Compare> template <typename T, typename U, class StorageT, class StorageU, class Compare>
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys, VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values, Compare comp) vtkm::cont::ArrayHandle<U, StorageU>& values,
Compare comp)
{ {
typedef vtkm::cont::ArrayHandle<T, StorageT> KeyType; typedef vtkm::cont::ArrayHandle<T, StorageT> KeyType;
if (sizeof(U) > sizeof(vtkm::Id)) if (sizeof(U) > sizeof(vtkm::Id))

@ -95,7 +95,8 @@ struct ReduceBody
BinaryOperationType BinaryOperation; BinaryOperationType BinaryOperation;
VTKM_CONT VTKM_CONT
ReduceBody(const InputPortalType& inputPortal, T initialValue, ReduceBody(const InputPortalType& inputPortal,
T initialValue,
BinaryOperationType binaryOperation) BinaryOperationType binaryOperation)
: Sum(vtkm::TypeTraits<T>::ZeroInitialization()) : Sum(vtkm::TypeTraits<T>::ZeroInitialization())
, InitialValue(initialValue) , InitialValue(initialValue)
@ -166,7 +167,8 @@ struct ReduceBody
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
template <class InputPortalType, typename T, class BinaryOperationType> template <class InputPortalType, typename T, class BinaryOperationType>
VTKM_CONT static T ReducePortals(InputPortalType inputPortal, T initialValue, VTKM_CONT static T ReducePortals(InputPortalType inputPortal,
T initialValue,
BinaryOperationType binaryOperation) BinaryOperationType binaryOperation)
{ {
typedef internal::WrappedBinaryOperator<T, BinaryOperationType> WrappedBinaryOp; typedef internal::WrappedBinaryOperator<T, BinaryOperationType> WrappedBinaryOp;
@ -204,7 +206,8 @@ struct ScanInclusiveBody
BinaryOperationType BinaryOperation; BinaryOperationType BinaryOperation;
VTKM_CONT VTKM_CONT
ScanInclusiveBody(const InputPortalType& inputPortal, const OutputPortalType& outputPortal, ScanInclusiveBody(const InputPortalType& inputPortal,
const OutputPortalType& outputPortal,
BinaryOperationType binaryOperation) BinaryOperationType binaryOperation)
: Sum(vtkm::TypeTraits<ValueType>::ZeroInitialization()) : Sum(vtkm::TypeTraits<ValueType>::ZeroInitialization())
, FirstCall(true) , FirstCall(true)
@ -292,8 +295,10 @@ struct ScanExclusiveBody
BinaryOperationType BinaryOperation; BinaryOperationType BinaryOperation;
VTKM_CONT VTKM_CONT
ScanExclusiveBody(const InputPortalType& inputPortal, const OutputPortalType& outputPortal, ScanExclusiveBody(const InputPortalType& inputPortal,
BinaryOperationType binaryOperation, const ValueType& initialValue) const OutputPortalType& outputPortal,
BinaryOperationType binaryOperation,
const ValueType& initialValue)
: Sum(initialValue) : Sum(initialValue)
, FirstCall(true) , FirstCall(true)
, InputPortal(inputPortal) , InputPortal(inputPortal)
@ -388,7 +393,8 @@ struct ScanExclusiveBody
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
template <class InputPortalType, class OutputPortalType, class BinaryOperationType> template <class InputPortalType, class OutputPortalType, class BinaryOperationType>
VTKM_CONT static typename std::remove_reference<typename OutputPortalType::ValueType>::type VTKM_CONT static typename std::remove_reference<typename OutputPortalType::ValueType>::type
ScanInclusivePortals(InputPortalType inputPortal, OutputPortalType outputPortal, ScanInclusivePortals(InputPortalType inputPortal,
OutputPortalType outputPortal,
BinaryOperationType binaryOperation) BinaryOperationType binaryOperation)
{ {
using ValueType = typename std::remove_reference<typename OutputPortalType::ValueType>::type; using ValueType = typename std::remove_reference<typename OutputPortalType::ValueType>::type;
@ -409,7 +415,9 @@ VTKM_SUPPRESS_EXEC_WARNINGS
template <class InputPortalType, class OutputPortalType, class BinaryOperationType> template <class InputPortalType, class OutputPortalType, class BinaryOperationType>
VTKM_CONT static typename std::remove_reference<typename OutputPortalType::ValueType>::type VTKM_CONT static typename std::remove_reference<typename OutputPortalType::ValueType>::type
ScanExclusivePortals( ScanExclusivePortals(
InputPortalType inputPortal, OutputPortalType outputPortal, BinaryOperationType binaryOperation, InputPortalType inputPortal,
OutputPortalType outputPortal,
BinaryOperationType binaryOperation,
typename std::remove_reference<typename OutputPortalType::ValueType>::type initialValue) typename std::remove_reference<typename OutputPortalType::ValueType>::type initialValue)
{ {
using ValueType = typename std::remove_reference<typename OutputPortalType::ValueType>::type; using ValueType = typename std::remove_reference<typename OutputPortalType::ValueType>::type;
@ -433,7 +441,8 @@ template <typename InputPortalType, typename IndexPortalType, typename OutputPor
class ScatterKernel class ScatterKernel
{ {
public: public:
VTKM_CONT ScatterKernel(InputPortalType inputPortal, IndexPortalType indexPortal, VTKM_CONT ScatterKernel(InputPortalType inputPortal,
IndexPortalType indexPortal,
OutputPortalType outputPortal) OutputPortalType outputPortal)
: ValuesPortal(inputPortal) : ValuesPortal(inputPortal)
, IndexPortal(indexPortal) , IndexPortal(indexPortal)
@ -478,7 +487,8 @@ private:
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_SUPPRESS_EXEC_WARNINGS
template <typename InputPortalType, typename IndexPortalType, typename OutputPortalType> template <typename InputPortalType, typename IndexPortalType, typename OutputPortalType>
VTKM_CONT static void ScatterPortal(InputPortalType inputPortal, IndexPortalType indexPortal, VTKM_CONT static void ScatterPortal(InputPortalType inputPortal,
IndexPortalType indexPortal,
OutputPortalType outputPortal) OutputPortalType outputPortal)
{ {
const vtkm::Id size = inputPortal.GetNumberOfValues(); const vtkm::Id size = inputPortal.GetNumberOfValues();

@ -51,7 +51,9 @@ template <typename RandomAccessIterator, typename Compare>
class quick_sort_range : private no_assign class quick_sort_range : private no_assign
{ {
inline size_t median_of_three(const RandomAccessIterator& array, size_t l, size_t m, inline size_t median_of_three(const RandomAccessIterator& array,
size_t l,
size_t m,
size_t r) const size_t r) const
{ {
return comp(array[l], array[m]) return comp(array[l], array[m])
@ -63,7 +65,8 @@ class quick_sort_range : private no_assign
const quick_sort_range& range) const const quick_sort_range& range) const
{ {
size_t offset = range.size / 8u; size_t offset = range.size / 8u;
return median_of_three(array, median_of_three(array, 0, offset, offset * 2), return median_of_three(array,
median_of_three(array, 0, offset, offset * 2),
median_of_three(array, offset * 3, offset * 4, offset * 5), median_of_three(array, offset * 3, offset * 4, offset * 5),
median_of_three(array, offset * 6, offset * 7, range.size - 1)); median_of_three(array, offset * 6, offset * 7, range.size - 1));
} }
@ -198,14 +201,16 @@ void parallel_quick_sort(RandomAccessIterator begin, RandomAccessIterator end, c
} }
parallel_for(blocked_range<RandomAccessIterator>(k + 1, end), parallel_for(blocked_range<RandomAccessIterator>(k + 1, end),
quick_sort_pretest_body<RandomAccessIterator, Compare>(comp), auto_partitioner(), quick_sort_pretest_body<RandomAccessIterator, Compare>(comp),
auto_partitioner(),
my_context); my_context);
if (my_context.is_group_execution_cancelled()) if (my_context.is_group_execution_cancelled())
do_parallel_quick_sort: do_parallel_quick_sort:
#endif /* __TBB_TASK_GROUP_CONTEXT */ #endif /* __TBB_TASK_GROUP_CONTEXT */
parallel_for(quick_sort_range<RandomAccessIterator, Compare>(begin, end - begin, comp), parallel_for(quick_sort_range<RandomAccessIterator, Compare>(begin, end - begin, comp),
quick_sort_body<RandomAccessIterator, Compare>(), auto_partitioner()); quick_sort_body<RandomAccessIterator, Compare>(),
auto_partitioner());
} }
} // namespace internal } // namespace internal
@ -248,8 +253,8 @@ void parallel_sort(RandomAccessIterator begin, RandomAccessIterator end, const C
template <typename RandomAccessIterator> template <typename RandomAccessIterator>
inline void parallel_sort(RandomAccessIterator begin, RandomAccessIterator end) inline void parallel_sort(RandomAccessIterator begin, RandomAccessIterator end)
{ {
parallel_sort(begin, end, parallel_sort(
std::less<typename std::iterator_traits<RandomAccessIterator>::value_type>()); begin, end, std::less<typename std::iterator_traits<RandomAccessIterator>::value_type>());
} }
//! Sorts the data in the range \c [begin,end) with a default comparator \c std::less<T> //! Sorts the data in the range \c [begin,end) with a default comparator \c std::less<T>

@ -528,8 +528,11 @@ inline vtkm::cont::DataSet MakeTestDataSet::Make3DExplicitDataSet1()
coords[2] = CoordType(1, 1, 0); coords[2] = CoordType(1, 1, 0);
coords[3] = CoordType(2, 1, 0); coords[3] = CoordType(2, 1, 0);
coords[4] = CoordType(2, 2, 0); coords[4] = CoordType(2, 2, 0);
CoordType coordinates[nVerts] = { CoordType(0, 0, 0), CoordType(1, 0, 0), CoordType(1, 1, 0), CoordType coordinates[nVerts] = { CoordType(0, 0, 0),
CoordType(2, 1, 0), CoordType(2, 2, 0) }; CoordType(1, 0, 0),
CoordType(1, 1, 0),
CoordType(2, 1, 0),
CoordType(2, 2, 0) };
vtkm::Float32 vars[nVerts] = { 10.1f, 20.1f, 30.2f, 40.2f, 50.3f }; vtkm::Float32 vars[nVerts] = { 10.1f, 20.1f, 30.2f, 40.2f, 50.3f };
dataSet.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", coordinates, nVerts)); dataSet.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", coordinates, nVerts));
@ -666,8 +669,9 @@ inline vtkm::cont::DataSet MakeTestDataSet::Make3DExplicitDataSet3()
const int nVerts = 4; const int nVerts = 4;
typedef vtkm::Vec<vtkm::Float32, 3> CoordType; typedef vtkm::Vec<vtkm::Float32, 3> CoordType;
CoordType coordinates[nVerts] = { CoordType(0, 0, 0), CoordType(1, 0, 0), CoordType(1, 0, 1), CoordType coordinates[nVerts] = {
CoordType(0, 1, 0) }; CoordType(0, 0, 0), CoordType(1, 0, 0), CoordType(1, 0, 1), CoordType(0, 1, 0)
};
vtkm::Float32 vars[nVerts] = { 10.1f, 10.1f, 10.2f, 30.2f }; vtkm::Float32 vars[nVerts] = { 10.1f, 10.1f, 10.2f, 30.2f };
dataSet.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", coordinates, nVerts)); dataSet.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", coordinates, nVerts));

@ -257,8 +257,9 @@ private:
{ {
gotException = true; gotException = true;
} }
VTKM_TEST_ASSERT(gotException, "PrepareForOutput should fail when asked to " VTKM_TEST_ASSERT(gotException,
"re-allocate user provided memory."); "PrepareForOutput should fail when asked to "
"re-allocate user provided memory.");
} }
} }
}; };

Some files were not shown because too many files have changed in this diff Show More