Shuenhoy/vtk-m2
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into cmake_refactor

Browse Source
This commit is contained in:
Matthew Letter 2017-11-28 13:36:02 -07:00
commit fac43bd812
103 changed files with 2609 additions and 2282 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
version.txt
View File

@ -1 +1 @@
1.0.0
1.1.0

2
vtkm/CMakeLists.txt
View File

@ -35,6 +35,7 @@ set(headers
CellShape.h
CellTraits.h
Hash.h
ImplicitFunction.h
ListTag.h
Math.h
Matrix.h
@ -55,6 +56,7 @@ set(headers
VectorAnalysis.h
VecTraits.h
VecVariable.h
VirtualObjectBase.h
UnaryPredicates.h
)

643
vtkm/ImplicitFunction.h Normal file
View File

@ -0,0 +1,643 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_ImplicitFunction_h
#define vtk_m_ImplicitFunction_h
#include <vtkm/Math.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/VirtualObjectBase.h>
namespace vtkm
{
//============================================================================
class VTKM_ALWAYS_EXPORT ImplicitFunction : public vtkm::VirtualObjectBase
{
public:
using Scalar = vtkm::FloatDefault;
using Vector = vtkm::Vec<Scalar, 3>;
VTKM_EXEC_CONT virtual Scalar Value(const Vector& point) const = 0;
VTKM_EXEC_CONT virtual Vector Gradient(const Vector& point) const = 0;
VTKM_EXEC_CONT Scalar Value(Scalar x, Scalar y, Scalar z) const
{
return this->Value(Vector(x, y, z));
}
VTKM_EXEC_CONT Vector Gradient(Scalar x, Scalar y, Scalar z) const
{
return this->Gradient(Vector(x, y, z));
}
};
//============================================================================
/// A helpful functor that calls the (virtual) value method of a given ImplicitFunction. Can be
/// passed to things that expect a functor instead of an ImplictFunction class (like an array
/// transform).
///
class VTKM_ALWAYS_EXPORT ImplicitFunctionValue
{
public:
using Scalar = vtkm::ImplicitFunction::Scalar;
using Vector = vtkm::ImplicitFunction::Vector;
VTKM_EXEC_CONT ImplicitFunctionValue()
: Function(nullptr)
{
}
VTKM_EXEC_CONT ImplicitFunctionValue(const ImplicitFunction* function)
: Function(function)
{
}
VTKM_EXEC_CONT Scalar operator()(const Vector& point) const
{
return this->Function->Value(point);
}
private:
const vtkm::ImplicitFunction* Function;
};
/// A helpful functor that calls the (virtual) gradient method of a given ImplicitFunction. Can be
/// passed to things that expect a functor instead of an ImplictFunction class (like an array
/// transform).
///
class VTKM_ALWAYS_EXPORT ImplicitFunctionGradient
{
public:
using Scalar = vtkm::ImplicitFunction::Scalar;
using Vector = vtkm::ImplicitFunction::Vector;
VTKM_EXEC_CONT ImplicitFunctionGradient()
: Function(nullptr)
{
}
VTKM_EXEC_CONT ImplicitFunctionGradient(const ImplicitFunction* function)
: Function(function)
{
}
VTKM_EXEC_CONT Vector operator()(const Vector& point) const
{
return this->Function->Gradient(point);
}
private:
const vtkm::ImplicitFunction* Function;
};
//============================================================================
/// \brief Implicit function for a box
class VTKM_ALWAYS_EXPORT Box : public ImplicitFunction
{
public:
VTKM_EXEC_CONT Box()
: MinPoint(Vector(Scalar(0)))
, MaxPoint(Vector(Scalar(0)))
{
}
VTKM_EXEC_CONT Box(const Vector& minPoint, const Vector& maxPoint)
: MinPoint(minPoint)
, MaxPoint(maxPoint)
{
}
VTKM_EXEC_CONT Box(Scalar xmin, Scalar xmax, Scalar ymin, Scalar ymax, Scalar zmin, Scalar zmax)
: MinPoint(xmin, ymin, zmin)
, MaxPoint(xmax, ymax, zmax)
{
}
VTKM_CONT void SetMinPoint(const Vector& point)
{
this->MinPoint = point;
this->Modified();
}
VTKM_CONT void SetMaxPoint(const Vector& point)
{
this->MaxPoint = point;
this->Modified();
}
VTKM_EXEC_CONT const Vector& GetMinPoint() const { return this->MinPoint; }
VTKM_EXEC_CONT const Vector& GetMaxPoint() const { return this->MaxPoint; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
{
Scalar minDistance = vtkm::NegativeInfinity32();
Scalar diff, t, dist;
Scalar distance = Scalar(0.0);
vtkm::IdComponent inside = 1;
for (vtkm::IdComponent d = 0; d < 3; d++)
{
diff = this->MaxPoint[d] - this->MinPoint[d];
if (diff != Scalar(0.0))
{
t = (point[d] - this->MinPoint[d]) / diff;
// Outside before the box
if (t < Scalar(0.0))
{
inside = 0;
dist = this->MinPoint[d] - point[d];
}
// Outside after the box
else if (t > Scalar(1.0))
{
inside = 0;
dist = point[d] - this->MaxPoint[d];
}
else
{
// Inside the box in lower half
if (t <= Scalar(0.5))
{
dist = MinPoint[d] - point[d];
}
// Inside the box in upper half
else
{
dist = point[d] - MaxPoint[d];
}
if (dist > minDistance)
{
minDistance = dist;
}
}
}
else
{
dist = vtkm::Abs(point[d] - MinPoint[d]);
if (dist > Scalar(0.0))
{
inside = 0;
}
}
if (dist > Scalar(0.0))
{
distance += dist * dist;
}
}
distance = vtkm::Sqrt(distance);
if (inside)
{
return minDistance;
}
else
{
return distance;
}
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
{
vtkm::IdComponent minAxis = 0;
Scalar dist = 0.0;
Scalar minDist = vtkm::Infinity32();
vtkm::Vec<vtkm::IdComponent, 3> location;
Vector normal(Scalar(0));
Vector inside(Scalar(0));
Vector outside(Scalar(0));
Vector center((this->MaxPoint + this->MinPoint) * Scalar(0.5));
// Compute the location of the point with respect to the box
// Point will lie in one of 27 separate regions around or within the box
// Gradient vector is computed differently in each of the regions.
for (vtkm::IdComponent d = 0; d < 3; d++)
{
if (point[d] < this->MinPoint[d])
{
// Outside the box low end
location[d] = 0;
outside[d] = -1.0;
}
else if (point[d] > this->MaxPoint[d])
{
// Outside the box high end
location[d] = 2;
outside[d] = 1.0;
}
else
{
location[d] = 1;
if (point[d] <= center[d])
{
// Inside the box low end
dist = point[d] - this->MinPoint[d];
inside[d] = -1.0;
}
else
{
// Inside the box high end
dist = this->MaxPoint[d] - point[d];
inside[d] = 1.0;
}
if (dist < minDist) // dist is negative
{
minDist = dist;
minAxis = d;
}
}
}
vtkm::Id indx = location[0] + 3 * location[1] + 9 * location[2];
switch (indx)
{
// verts - gradient points away from center point
case 0:
case 2:
case 6:
case 8:
case 18:
case 20:
case 24:
case 26:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
normal[d] = point[d] - center[d];
}
vtkm::Normalize(normal);
break;
// edges - gradient points out from axis of cube
case 1:
case 3:
case 5:
case 7:
case 9:
case 11:
case 15:
case 17:
case 19:
case 21:
case 23:
case 25:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
if (outside[d] != 0.0)
{
normal[d] = point[d] - center[d];
}
else
{
normal[d] = 0.0;
}
}
vtkm::Normalize(normal);
break;
// faces - gradient points perpendicular to face
case 4:
case 10:
case 12:
case 14:
case 16:
case 22:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
normal[d] = outside[d];
}
break;
// interior - gradient is perpendicular to closest face
case 13:
normal[0] = normal[1] = normal[2] = 0.0;
normal[minAxis] = inside[minAxis];
break;
default:
VTKM_ASSERT(false);
break;
}
return normal;
}
private:
Vector MinPoint;
Vector MaxPoint;
};
//============================================================================
/// \brief Implicit function for a cylinder
class VTKM_ALWAYS_EXPORT Cylinder : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Cylinder()
: Center(Scalar(0))
, Axis(Scalar(1), Scalar(0), Scalar(0))
, Radius(Scalar(0.2))
{
}
VTKM_EXEC_CONT Cylinder(const Vector& axis, Scalar radius)
: Center(Scalar(0))
, Axis(axis)
, Radius(radius)
{
}
VTKM_EXEC_CONT Cylinder(const Vector& center, const Vector& axis, Scalar radius)
: Center(center)
, Axis(axis)
, Radius(radius)
{
}
VTKM_CONT void SetCenter(const Vector& center)
{
this->Center = center;
this->Modified();
}
VTKM_CONT void SetAxis(const Vector& axis)
{
this->Axis = axis;
this->Modified();
}
VTKM_CONT void SetRadius(Scalar radius)
{
this->Radius = radius;
this->Modified();
}
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
{
Vector x2c = point - this->Center;
FloatDefault proj = vtkm::dot(this->Axis, x2c);
return vtkm::dot(x2c, x2c) - (proj * proj) - (this->Radius * this->Radius);
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
{
Vector x2c = point - this->Center;
FloatDefault t = this->Axis[0] * x2c[0] + this->Axis[1] * x2c[1] + this->Axis[2] * x2c[2];
vtkm::Vec<FloatDefault, 3> closestPoint = this->Center + (this->Axis * t);
return (point - closestPoint) * FloatDefault(2);
}
private:
Vector Center;
Vector Axis;
Scalar Radius;
};
//============================================================================
/// \brief Implicit function for a frustum
class VTKM_ALWAYS_EXPORT Frustum : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Frustum() = default;
VTKM_EXEC_CONT Frustum(const Vector points[6], const Vector normals[6])
{
this->SetPlanes(points, normals);
}
VTKM_EXEC_CONT explicit Frustum(const Vector points[8]) { this->CreateFromPoints(points); }
VTKM_EXEC void SetPlanes(const Vector points[6], const Vector normals[6])
{
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
this->Points[index] = points[index];
}
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
this->Normals[index] = normals[index];
}
this->Modified();
}
VTKM_EXEC void SetPlane(int idx, const Vector& point, const Vector& normal)
{
VTKM_ASSERT((idx >= 0) && (idx < 6));
this->Points[idx] = point;
this->Normals[idx] = normal;
this->Modified();
}
VTKM_EXEC_CONT void GetPlanes(Vector points[6], Vector normals[6]) const
{
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
points[index] = this->Points[index];
}
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
normals[index] = this->Normals[index];
}
}
VTKM_EXEC_CONT const Vector* GetPoints() const { return this->Points; }
VTKM_EXEC_CONT const Vector* GetNormals() const { return this->Normals; }
// The points should be specified in the order of hex-cell vertices
VTKM_EXEC_CONT void CreateFromPoints(const Vector points[8])
{
// XXX(clang-format-3.9): 3.8 is silly. 3.9 makes it look like this.
// clang-format off
int planes[6][3] = {
{ 3, 2, 0 }, { 4, 5, 7 }, { 0, 1, 4 }, { 1, 2, 5 }, { 2, 3, 6 }, { 3, 0, 7 }
};
// clang-format on
for (int i = 0; i < 6; ++i)
{
const Vector& v0 = points[planes[i][0]];
const Vector& v1 = points[planes[i][1]];
const Vector& v2 = points[planes[i][2]];
this->Points[i] = v0;
this->Normals[i] = vtkm::Normal(vtkm::Cross(v2 - v0, v1 - v0));
}
this->Modified();
}
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
{
Scalar maxVal = vtkm::NegativeInfinity<Scalar>();
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
const Vector& p = this->Points[index];
const Vector& n = this->Normals[index];
const Scalar val = vtkm::dot(point - p, n);
maxVal = vtkm::Max(maxVal, val);
}
return maxVal;
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
{
Scalar maxVal = vtkm::NegativeInfinity<Scalar>();
vtkm::Id maxValIdx = 0;
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
{
const Vector& p = this->Points[index];
const Vector& n = this->Normals[index];
Scalar val = vtkm::dot(point - p, n);
if (val > maxVal)
{
maxVal = val;
maxValIdx = index;
}
}
return this->Normals[maxValIdx];
}
private:
Vector Points[6];
Vector Normals[6];
};
//============================================================================
/// \brief Implicit function for a plane
class VTKM_ALWAYS_EXPORT Plane : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Plane()
: Origin(Scalar(0))
, Normal(Scalar(0), Scalar(0), Scalar(1))
{
}
VTKM_EXEC_CONT explicit Plane(const Vector& normal)
: Origin(Scalar(0))
, Normal(normal)
{
}
VTKM_EXEC_CONT Plane(const Vector& origin, const Vector& normal)
: Origin(origin)
, Normal(normal)
{
}
VTKM_CONT void SetOrigin(const Vector& origin)
{
this->Origin = origin;
this->Modified();
}
VTKM_CONT void SetNormal(const Vector& normal)
{
this->Normal = normal;
this->Modified();
}
VTKM_EXEC_CONT const Vector& GetOrigin() const { return this->Origin; }
VTKM_EXEC_CONT const Vector& GetNormal() const { return this->Normal; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
{
return vtkm::dot(point - this->Origin, this->Normal);
}
VTKM_EXEC_CONT Vector Gradient(const Vector&) const override { return this->Normal; }
private:
Vector Origin;
Vector Normal;
};
//============================================================================
/// \brief Implicit function for a sphere
class VTKM_ALWAYS_EXPORT Sphere : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Sphere()
: Radius(Scalar(0.2))
, Center(Scalar(0))
{
}
VTKM_EXEC_CONT explicit Sphere(Scalar radius)
: Radius(radius)
, Center(Scalar(0))
{
}
VTKM_EXEC_CONT Sphere(Vector center, Scalar radius)
: Radius(radius)
, Center(center)
{
}
VTKM_CONT void SetRadius(Scalar radius)
{
this->Radius = radius;
this->Modified();
}
VTKM_CONT void SetCenter(const Vector& center)
{
this->Center = center;
this->Modified();
}
VTKM_EXEC_CONT Scalar GetRadius() const { return this->Radius; }
VTKM_EXEC_CONT const Vector& GetCenter() const { return this->Center; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
{
return vtkm::MagnitudeSquared(point - this->Center) - (this->Radius * this->Radius);
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
{
return Scalar(2) * (point - this->Center);
}
private:
Scalar Radius;
Vector Center;
};
} // namespace vtkm
#ifdef VTKM_CUDA
// Cuda seems to have a bug where it expects the template class VirtualObjectTransfer
// to be instantiated in a consitent order among all the translation units of an
// executable. Failing to do so results in random crashes and incorrect results.
// We workaroud this issue by explicitly instantiating VirtualObjectTransfer for
// all the implicit functions here.
#include <vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h>
VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(vtkm::Box);
VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(vtkm::Cylinder);
VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(vtkm::Frustum);
VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(vtkm::Plane);
VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(vtkm::Sphere);
#endif
#endif //vtk_m_ImplicitFunction_h

19
vtkm/ListTag.h
View File

@ -73,7 +73,7 @@ struct ListTagJoin : detail::ListRoot
using list = typename detail::ListJoin<typename ListTag1::list, typename ListTag2::list>::type;
};
/// A tag that consits of elements that are found in both tags. This struct
/// A tag that consists of elements that are found in both tags. This struct
/// can be subclassed and still behave like a list tag.
template <typename ListTag1, typename ListTag2>
struct ListTagIntersect : detail::ListRoot
@ -85,13 +85,24 @@ struct ListTagIntersect : detail::ListRoot
/// For each typename represented by the list tag, call the functor with a
/// default instance of that type.
///
template <typename Functor, typename ListTag>
VTKM_CONT void ListForEach(Functor&& f, ListTag)
template <typename Functor, typename ListTag, typename... Args>
VTKM_CONT void ListForEach(Functor&& f, ListTag, Args&&... args)
{
VTKM_IS_LIST_TAG(ListTag);
detail::ListForEachImpl(f, typename ListTag::list());
detail::ListForEachImpl(
std::forward<Functor>(f), typename ListTag::list{}, std::forward<Args>(args)...);
}
/// Generate a tag that is the cross product of two other tags. The resulting
// a tag has the form of Tag< std::pair<A1,B1>, std::pair<A1,B2> .... >
///
template <typename ListTag1, typename ListTag2>
struct ListCrossProduct : detail::ListRoot
{
using list =
typename detail::ListCrossProductImpl<typename ListTag1::list, typename ListTag2::list>::type;
};
/// Checks to see if the given \c Type is in the list pointed to by \c ListTag.
/// There is a static boolean named \c value that is set to true if the type is
/// contained in the list and false otherwise.

82
vtkm/VirtualObjectBase.h Normal file
View File

@ -0,0 +1,82 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_VirtualObjectBase_h
#define vtk_m_VirtualObjectBase_h
#include <vtkm/Types.h>
namespace vtkm
{
/// \brief Base class for virtual objects that work in the execution environment
///
/// Any class built in VTK-m that has virtual methods and is intended to work in both the control
/// and execution environment should inherit from \c VirtualObjectBase. Hierarchies under \c
/// VirtualObjectBase can be used in conjunction with \c VirtualObjectHandle to transfer from the
/// control environment (where they are set up) to the execution environment (where they are used).
///
/// In addition to inheriting from \c VirtualObjectBase, virtual objects have to satisfy 2 other
/// conditions to work correctly. First, they have to be a plain old data type that can be copied
/// with \c memcpy (with the exception of the virtual table, which \c VirtualObjectHandle will take
/// care of). Second, if the object changes its state in the control environment, it should call
/// \c Modified on itself so the \c VirtualObjectHandle will know it update the object in the
/// execution environment.
///
class VTKM_ALWAYS_EXPORT VirtualObjectBase
{
public:
VTKM_EXEC_CONT virtual ~VirtualObjectBase() = default;
VTKM_EXEC_CONT void Modified() { this->ModifiedCount++; }
VTKM_EXEC_CONT vtkm::Id GetModifiedCount() const { return this->ModifiedCount; }
protected:
VTKM_EXEC_CONT VirtualObjectBase()
: ModifiedCount(0)
{
}
VTKM_EXEC_CONT VirtualObjectBase(const VirtualObjectBase&) = default;
VTKM_EXEC_CONT VirtualObjectBase(VirtualObjectBase&& other)
: ModifiedCount(other.ModifiedCount)
{
}
VTKM_EXEC_CONT VirtualObjectBase& operator=(const VirtualObjectBase&)
{
this->Modified();
return *this;
}
VTKM_EXEC_CONT VirtualObjectBase& operator=(VirtualObjectBase&&)
{
this->Modified();
return *this;
}
private:
vtkm::Id ModifiedCount;
};
} // namespace vtkm
#endif //vtk_m_VirtualObjectBase_h

78
vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx
View File

@ -24,7 +24,7 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/cont/Timer.h>
#include <vtkm/worklet/DispatcherMapField.h>
@ -264,7 +264,7 @@ public:
typedef void ControlSignature(FieldIn<Vec3>, FieldOut<Scalar>);
typedef void ExecutionSignature(_1, _2);
EvaluateImplicitFunction(const ImplicitFunction& function)
EvaluateImplicitFunction(const ImplicitFunction* function)
: Function(function)
{
}
@ -272,11 +272,11 @@ public:
template <typename VecType, typename ScalarType>
VTKM_EXEC void operator()(const VecType& point, ScalarType& val) const
{
val = this->Function.Value(point);
val = this->Function->Value(point);
}
private:
ImplicitFunction Function;
const ImplicitFunction* Function;
};
template <typename T1, typename T2>
@ -286,7 +286,7 @@ public:
typedef void ControlSignature(FieldIn<Vec3>, FieldOut<Scalar>);
typedef void ExecutionSignature(_1, _2);
Evaluate2ImplicitFunctions(const T1& f1, const T2& f2)
Evaluate2ImplicitFunctions(const T1* f1, const T2* f2)
: Function1(f1)
, Function2(f2)
{
@ -295,12 +295,12 @@ public:
template <typename VecType, typename ScalarType>
VTKM_EXEC void operator()(const VecType& point, ScalarType& val) const
{
val = this->Function1.Value(point) + this->Function2.Value(point);
val = this->Function1->Value(point) + this->Function2->Value(point);
}
private:
T1 Function1;
T2 Function2;
const T1* Function1;
const T2* Function2;
};
struct ValueTypes : vtkm::ListTagBase<vtkm::Float32, vtkm::Float64>
@ -680,7 +680,7 @@ private:
{
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> Points;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> Result;
vtkm::cont::Sphere Sphere1, Sphere2;
vtkm::Sphere Sphere1, Sphere2;
};
static ImplicitFunctionBenchData MakeImplicitFunctionBenchData()
@ -703,8 +703,8 @@ private:
portal.Set(i, vtkm::make_Vec(distx(rangen), disty(rangen), distz(rangen)));
}
data.Sphere1 = vtkm::cont::Sphere({ 0.22f, 0.33f, 0.44f }, 0.55f);
data.Sphere2 = vtkm::cont::Sphere({ 0.22f, 0.33f, 0.11f }, 0.77f);
data.Sphere1 = vtkm::Sphere({ 0.22f, 0.33f, 0.44f }, 0.55f);
data.Sphere2 = vtkm::Sphere({ 0.22f, 0.33f, 0.11f }, 0.77f);
return data;
}
@ -720,10 +720,13 @@ private:
VTKM_CONT
vtkm::Float64 operator()()
{
using EvalWorklet = EvaluateImplicitFunction<vtkm::cont::Sphere>;
using EvalWorklet = EvaluateImplicitFunction<vtkm::Sphere>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
EvalWorklet eval(Internal.Sphere1);
auto handle = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto function =
static_cast<const vtkm::Sphere*>(handle.PrepareForExecution(DeviceAdapterTag()));
EvalWorklet eval(function);
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
@ -734,8 +737,8 @@ private:
std::string Description() const
{
std::stringstream description;
description << "Implicit Function (vtkm::cont::Sphere) on "
<< Internal.Points.GetNumberOfValues() << " points";
description << "Implicit Function (vtkm::Sphere) on " << Internal.Points.GetNumberOfValues()
<< " points";
return description.str();
}
@ -743,9 +746,9 @@ private:
};
template <typename Value>
struct BenchDynamicImplicitFunction
struct BenchVirtualImplicitFunction
{
BenchDynamicImplicitFunction()
BenchVirtualImplicitFunction()
: Internal(MakeImplicitFunctionBenchData())
{
}
@ -753,10 +756,11 @@ private:
VTKM_CONT
vtkm::Float64 operator()()
{
using EvalWorklet = EvaluateImplicitFunction<vtkm::exec::ImplicitFunction>;
using EvalWorklet = EvaluateImplicitFunction<vtkm::ImplicitFunction>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
EvalWorklet eval(Internal.Sphere1.PrepareForExecution(DeviceAdapterTag()));
auto sphere = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
EvalWorklet eval(sphere.PrepareForExecution(DeviceAdapterTag()));
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
@ -767,7 +771,7 @@ private:
std::string Description() const
{
std::stringstream description;
description << "Implicit Function (DynamicImplicitFunction) on "
description << "Implicit Function (VirtualImplicitFunction) on "
<< Internal.Points.GetNumberOfValues() << " points";
return description.str();
}
@ -786,10 +790,14 @@ private:
VTKM_CONT
vtkm::Float64 operator()()
{
using EvalWorklet = Evaluate2ImplicitFunctions<vtkm::cont::Sphere, vtkm::cont::Sphere>;
using EvalWorklet = Evaluate2ImplicitFunctions<vtkm::Sphere, vtkm::Sphere>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
EvalWorklet eval(Internal.Sphere1, Internal.Sphere2);
auto h1 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto h2 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere2);
auto f1 = static_cast<const vtkm::Sphere*>(h1.PrepareForExecution(DeviceAdapterTag()));
auto f2 = static_cast<const vtkm::Sphere*>(h2.PrepareForExecution(DeviceAdapterTag()));
EvalWorklet eval(f1, f2);
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
@ -800,8 +808,8 @@ private:
std::string Description() const
{
std::stringstream description;
description << "Implicit Function 2x(vtkm::cont::Sphere) on "
<< Internal.Points.GetNumberOfValues() << " points";
description << "Implicit Function 2x(vtkm::Sphere) on " << Internal.Points.GetNumberOfValues()
<< " points";
return description.str();
}
@ -809,9 +817,9 @@ private:
};
template <typename Value>
struct Bench2DynamicImplicitFunctions
struct Bench2VirtualImplicitFunctions
{
Bench2DynamicImplicitFunctions()
Bench2VirtualImplicitFunctions()
: Internal(MakeImplicitFunctionBenchData())
{
}
@ -820,11 +828,13 @@ private:
vtkm::Float64 operator()()
{
using EvalWorklet =
Evaluate2ImplicitFunctions<vtkm::exec::ImplicitFunction, vtkm::exec::ImplicitFunction>;
Evaluate2ImplicitFunctions<vtkm::ImplicitFunction, vtkm::ImplicitFunction>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
EvalWorklet eval(Internal.Sphere1.PrepareForExecution(DeviceAdapterTag()),
Internal.Sphere2.PrepareForExecution(DeviceAdapterTag()));
auto s1 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto s2 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere2);
EvalWorklet eval(s1.PrepareForExecution(DeviceAdapterTag()),
s2.PrepareForExecution(DeviceAdapterTag()));
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
@ -835,7 +845,7 @@ private:
std::string Description() const
{
std::stringstream description;
description << "Implicit Function 2x(DynamicImplicitFunction) on "
description << "Implicit Function 2x(VirtualImplicitFunction) on "
<< Internal.Points.GetNumberOfValues() << " points";
return description.str();
}
@ -844,9 +854,9 @@ private:
};
VTKM_MAKE_BENCHMARK(ImplicitFunction, BenchImplicitFunction);
VTKM_MAKE_BENCHMARK(ImplicitFunctionDynamic, BenchDynamicImplicitFunction);
VTKM_MAKE_BENCHMARK(ImplicitFunctionVirtual, BenchVirtualImplicitFunction);
VTKM_MAKE_BENCHMARK(ImplicitFunction2, Bench2ImplicitFunctions);
VTKM_MAKE_BENCHMARK(ImplicitFunctionDynamic2, Bench2DynamicImplicitFunctions);
VTKM_MAKE_BENCHMARK(ImplicitFunctionVirtual2, Bench2VirtualImplicitFunctions);
public:
static VTKM_CONT int Run(int benchmarks)
@ -887,9 +897,9 @@ public:
std::cout << "\nBenchmarking Implicit Function\n";
VTKM_RUN_BENCHMARK(ImplicitFunction, FloatDefaultType());
VTKM_RUN_BENCHMARK(ImplicitFunctionDynamic, FloatDefaultType());
VTKM_RUN_BENCHMARK(ImplicitFunctionVirtual, FloatDefaultType());
VTKM_RUN_BENCHMARK(ImplicitFunction2, FloatDefaultType());
VTKM_RUN_BENCHMARK(ImplicitFunctionDynamic2, FloatDefaultType());
VTKM_RUN_BENCHMARK(ImplicitFunctionVirtual2, FloatDefaultType());
}
return 0;

7
vtkm/cont/CMakeLists.txt
View File

@ -69,7 +69,7 @@ set(headers
ErrorExecution.h
ErrorInternal.h
Field.h
ImplicitFunction.h
ImplicitFunctionHandle.h
MultiBlock.h
PointLocatorUniformGrid.h
RuntimeDeviceInformation.h
@ -80,14 +80,13 @@ set(headers
StorageListTag.h
Timer.h
TryExecute.h
VirtualObjectCache.h
VirtualObjectHandle.h
)
set(template_sources
ArrayHandle.hxx
CellSetExplicit.hxx
CellSetStructured.hxx
ImplicitFunction.hxx
StorageBasic.hxx
)
@ -99,7 +98,6 @@ set(sources
CoordinateSystem.cxx
DynamicArrayHandle.cxx
Field.cxx
ImplicitFunction.cxx
internal/SimplePolymorphicContainer.cxx
MultiBlock.cxx
internal/ArrayManagerExecutionShareWithControl.cxx
@ -112,6 +110,7 @@ set(device_sources
ArrayRangeCompute.cxx
CellSetExplicit.cxx
RuntimeDeviceTracker.cxx
TryExecute.cxx
)
#-----------------------------------------------------------------------------

10
vtkm/cont/CellSetExplicit.h
View File

@ -106,11 +106,11 @@ public:
virtual ~CellSetExplicit();
vtkm::Id GetNumberOfCells() const VTKM_OVERRIDE;
vtkm::Id GetNumberOfPoints() const VTKM_OVERRIDE;
vtkm::Id GetNumberOfFaces() const VTKM_OVERRIDE;
vtkm::Id GetNumberOfEdges() const VTKM_OVERRIDE;
void PrintSummary(std::ostream& out) const VTKM_OVERRIDE;
vtkm::Id GetNumberOfCells() const override;
vtkm::Id GetNumberOfPoints() const override;
vtkm::Id GetNumberOfFaces() const override;
vtkm::Id GetNumberOfEdges() const override;
void PrintSummary(std::ostream& out) const override;
VTKM_CONT vtkm::Id GetSchedulingRange(vtkm::TopologyElementTagCell) const;
VTKM_CONT vtkm::Id GetSchedulingRange(vtkm::TopologyElementTagPoint) const;

16
vtkm/cont/CoordinateSystem.h
View File

@ -72,7 +72,17 @@ struct StorageListTagCoordinateSystemDefault
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>::StorageTag>
{
};
}
}
namespace vtkm
{
template struct ListCrossProduct<::vtkm::TypeListTagFieldVec3,
::vtkm::cont::StorageListTagCoordinateSystemDefault>;
namespace cont
{
using DynamicArrayHandleCoordinateSystem =
vtkm::cont::DynamicArrayHandleBase<VTKM_DEFAULT_COORDINATE_SYSTEM_TYPE_LIST_TAG,
VTKM_DEFAULT_COORDINATE_SYSTEM_STORAGE_LIST_TAG>;
@ -214,10 +224,10 @@ public:
virtual void PrintSummary(std::ostream& out) const;
};
template <typename Functor>
void CastAndCall(const vtkm::cont::CoordinateSystem& coords, const Functor& f)
template <typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::CoordinateSystem& coords, const Functor& f, Args&&... args)
{
coords.GetData().CastAndCall(f);
coords.GetData().CastAndCall(f, std::forward<Args>(args)...);
}
namespace internal

41
vtkm/cont/DeviceAdapterListTag.h
View File

@ -40,6 +40,47 @@ struct DeviceAdapterListTagCommon : vtkm::ListTagBase<vtkm::cont::DeviceAdapterT
vtkm::cont::DeviceAdapterTagSerial>
{
};
namespace detail
{
template <typename FunctorType>
class ExecuteIfValidDeviceTag
{
private:
template <typename DeviceAdapter>
using EnableIfValid = std::enable_if<vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
template <typename DeviceAdapter>
using EnableIfInvalid = std::enable_if<!vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
public:
explicit ExecuteIfValidDeviceTag(const FunctorType& functor)
: Functor(functor)
{
}
template <typename DeviceAdapter>
typename EnableIfValid<DeviceAdapter>::type operator()(DeviceAdapter) const
{
this->Functor(DeviceAdapter());
}
template <typename DeviceAdapter>
typename EnableIfInvalid<DeviceAdapter>::type operator()(DeviceAdapter) const
{
}
private:
FunctorType Functor;
};
} // detail
template <typename DeviceList, typename Functor>
VTKM_CONT void ForEachValidDevice(DeviceList devices, const Functor& functor)
{
vtkm::ListForEach(detail::ExecuteIfValidDeviceTag<Functor>(functor), devices);
}
}
} // namespace vtkm::cont

14
vtkm/cont/DynamicArrayHandle.cxx
View File

@ -18,6 +18,8 @@
// this software.
//============================================================================
#include <sstream>
#include <typeindex>
#include <vtkm/cont/DynamicArrayHandle.h>
namespace vtkm
@ -34,6 +36,18 @@ PolymorphicArrayHandleContainerBase::PolymorphicArrayHandleContainerBase()
PolymorphicArrayHandleContainerBase::~PolymorphicArrayHandleContainerBase()
{
}
void ThrowCastAndCallException(PolymorphicArrayHandleContainerBase* ptr,
const std::type_info* type,
const std::type_info* storage)
{
std::ostringstream out;
out << "Could not find appropriate cast for array in CastAndCall1.\n"
"Array: ";
ptr->PrintSummary(out);
out << "TypeList: " << type->name() << "\nStorageList: " << storage->name() << "\n";
throw vtkm::cont::ErrorBadValue(out.str());
}
}
}
} // namespace vtkm::cont::detail

158
vtkm/cont/DynamicArrayHandle.h
View File

@ -31,10 +31,11 @@
#include <vtkm/cont/internal/DynamicTransform.h>
#include <sstream>
namespace vtkm
{
template struct ListCrossProduct<VTKM_DEFAULT_TYPE_LIST_TAG, VTKM_DEFAULT_STORAGE_LIST_TAG>;
namespace cont
{
@ -131,9 +132,8 @@ template <typename Type, typename Storage>
VTKM_CONT vtkm::cont::ArrayHandle<Type, Storage>* DynamicArrayHandleTryCast(
vtkm::cont::detail::PolymorphicArrayHandleContainerBase* arrayContainer)
{
vtkm::cont::detail::PolymorphicArrayHandleContainer<Type, Storage>* downcastContainer =
dynamic_cast<vtkm::cont::detail::PolymorphicArrayHandleContainer<Type, Storage>*>(
arrayContainer);
vtkm::cont::detail::PolymorphicArrayHandleContainer<Type, Storage>* downcastContainer = nullptr;
downcastContainer = dynamic_cast<decltype(downcastContainer)>(arrayContainer);
if (downcastContainer != nullptr)
{
return &downcastContainer->Array;
@ -357,8 +357,8 @@ public:
/// two lists to VTKM_DEFAULT_TYPE_LIST_TAG and VTK_DEFAULT_STORAGE_LIST_TAG,
/// respectively.
///
template <typename Functor>
VTKM_CONT void CastAndCall(const Functor& f) const;
template <typename Functor, typename... Args>
VTKM_CONT void CastAndCall(const Functor& f, Args&&...) const;
/// \brief Create a new array of the same type as this array.
///
@ -406,135 +406,66 @@ using DynamicArrayHandle =
namespace detail
{
template <typename Functor, typename Type>
struct DynamicArrayHandleTryStorage
struct DynamicArrayHandleTry
{
const DynamicArrayHandle* const Array;
const Functor& Function;
bool FoundCast;
VTKM_CONT
DynamicArrayHandleTryStorage(const DynamicArrayHandle& array, const Functor& f)
: Array(&array)
, Function(f)
, FoundCast(false)
DynamicArrayHandleTry(const PolymorphicArrayHandleContainerBase* const c)
: Container(c)
{
}
template <typename Storage>
VTKM_CONT void operator()(Storage)
template <typename T, typename U, typename... Args>
void operator()(std::pair<T, U>&& p, Args&&... args) const
{
this->DoCast(Storage(),
typename vtkm::cont::internal::IsValidArrayHandle<Type, Storage>::type());
using storage = vtkm::cont::internal::Storage<T, U>;
using invalid = typename std::is_base_of<vtkm::cont::internal::UndefinedStorage, storage>::type;
this->run(std::forward<decltype(p)>(p), invalid{}, args...);
}
private:
template <typename Storage>
void DoCast(Storage, std::true_type)
template <typename T, typename U, typename Functor, typename... Args>
void run(std::pair<T, U>&&, std::false_type, Functor&& f, bool& called, Args&&... args) const
{
if (!this->FoundCast && this->Array->template IsTypeAndStorage<Type, Storage>())
if (!called)
{
this->Function(this->Array->template CastToTypeStorage<Type, Storage>());
this->FoundCast = true;
using downcastType = const vtkm::cont::detail::PolymorphicArrayHandleContainer<T, U>* const;
downcastType downcastContainer = dynamic_cast<downcastType>(this->Container);
if (downcastContainer)
{
f(downcastContainer->Array, std::forward<Args>(args)...);
called = true;
}
}
}
template <typename Storage>
void DoCast(Storage, std::false_type)
template <typename T, typename U, typename... Args>
void run(std::pair<T, U>&&, std::true_type, Args&&...) const
{
// This type of array handle cannot exist, so do nothing.
}
void operator=(const DynamicArrayHandleTryStorage<Functor, Type>&) = delete;
};
template <typename Functor, typename StorageList>
struct DynamicArrayHandleTryType
{
const DynamicArrayHandle* const Array;
const Functor& Function;
bool FoundCast;
VTKM_CONT
DynamicArrayHandleTryType(const DynamicArrayHandle& array, const Functor& f)
: Array(&array)
, Function(f)
, FoundCast(false)
{
}
template <typename Type>
VTKM_CONT void operator()(Type)
{
if (this->FoundCast)
{
return;
}
using TryStorageType = DynamicArrayHandleTryStorage<Functor, Type>;
TryStorageType tryStorage = TryStorageType(*this->Array, this->Function);
vtkm::ListForEach(tryStorage, StorageList());
if (tryStorage.FoundCast)
{
this->FoundCast = true;
}
}
private:
void operator=(const DynamicArrayHandleTryType<Functor, StorageList>&) = delete;
const PolymorphicArrayHandleContainerBase* const Container;
};
VTKM_CONT_EXPORT void ThrowCastAndCallException(PolymorphicArrayHandleContainerBase*,
const std::type_info*,
const std::type_info*);
} // namespace detail
template <typename TypeList, typename StorageList>
template <typename Functor>
VTKM_CONT void DynamicArrayHandleBase<TypeList, StorageList>::CastAndCall(const Functor& f) const
template <typename Functor, typename... Args>
VTKM_CONT void DynamicArrayHandleBase<TypeList, StorageList>::CastAndCall(const Functor& f,
Args&&... args) const
{
VTKM_IS_LIST_TAG(TypeList);
VTKM_IS_LIST_TAG(StorageList);
using TryTypeType = detail::DynamicArrayHandleTryType<Functor, StorageList>;
//For optimizations we should compile once the cross product for the default types
//and make it extern
using crossProduct = typename vtkm::ListCrossProduct<TypeList, StorageList>;
// We cast this to a DynamicArrayHandle because at this point we are ignoring
// the type/storage lists in it. There is no sense in adding more unnecessary
// template cases.
// The downside to this approach is that a copy is created, causing an
// atomic increment, which affects both performance and library size.
// For these reasons we have a specialization of this method to remove
// the copy when the type/storage lists are the default
DynamicArrayHandle t(*this);
TryTypeType tryType = TryTypeType(t, f);
vtkm::ListForEach(tryType, TypeList());
if (!tryType.FoundCast)
bool called = false;
auto* ptr = this->ArrayContainer.get();
vtkm::ListForEach(
detail::DynamicArrayHandleTry(ptr), crossProduct{}, f, called, std::forward<Args>(args)...);
if (!called)
{
std::ostringstream out;
out << "Could not find appropriate cast for array in CastAndCall1.\n"
"Array: ";
this->PrintSummary(out);
out << "TypeList: " << typeid(TypeList).name()
<< "\nStorageList: " << typeid(StorageList).name() << "\n";
throw vtkm::cont::ErrorBadValue(out.str());
}
}
template <>
template <typename Functor>
VTKM_CONT void
DynamicArrayHandleBase<VTKM_DEFAULT_TYPE_LIST_TAG, VTKM_DEFAULT_STORAGE_LIST_TAG>::CastAndCall(
const Functor& f) const
{
using TryTypeType = detail::DynamicArrayHandleTryType<Functor, VTKM_DEFAULT_STORAGE_LIST_TAG>;
// We can remove the copy, as the current DynamicArrayHandle is already
// the default one, and no reason to do an atomic increment and increase
// library size, and reduce performance
TryTypeType tryType = TryTypeType(*this, f);
vtkm::ListForEach(tryType, VTKM_DEFAULT_TYPE_LIST_TAG());
if (!tryType.FoundCast)
{
throw vtkm::cont::ErrorBadValue("Could not find appropriate cast for array in CastAndCall2.");
// throw an exception
detail::ThrowCastAndCallException(ptr, &typeid(TypeList), &typeid(StorageList));
}
}
@ -549,6 +480,7 @@ struct DynamicTransformTraits<vtkm::cont::DynamicArrayHandleBase<TypeList, Stora
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_DynamicArrayHandle_h

51
vtkm/cont/DynamicCellSet.h
View File

@ -227,8 +227,8 @@ public:
/// DynamicCellSet). You can use \c ResetCellSetList to get different
/// behavior from \c CastAndCall.
///
template <typename Functor>
VTKM_CONT void CastAndCall(const Functor& f) const;
template <typename Functor, typename... Args>
VTKM_CONT void CastAndCall(const Functor& f, Args&&...) const;
/// \brief Create a new cell set of the same type as this cell set.
///
@ -272,51 +272,42 @@ private:
namespace detail
{
template <typename Functor>
struct DynamicCellSetTryCellSet
struct DynamicCellSetTry
{
vtkm::cont::internal::SimplePolymorphicContainerBase* CellSetContainer;
const Functor& Function;
bool FoundCast;
VTKM_CONT
DynamicCellSetTryCellSet(vtkm::cont::internal::SimplePolymorphicContainerBase* cellSetContainer,
const Functor& f)
: CellSetContainer(cellSetContainer)
, Function(f)
, FoundCast(false)
DynamicCellSetTry(
const vtkm::cont::internal::SimplePolymorphicContainerBase* const cellSetContainer)
: Container(cellSetContainer)
{
}
template <typename CellSetType>
VTKM_CONT void operator()(CellSetType)
template <typename CellSetType, typename Functor, typename... Args>
void operator()(CellSetType, Functor&& f, bool& called, Args&&... args) const
{
if (!this->FoundCast)
using downcastType = const vtkm::cont::internal::SimplePolymorphicContainer<CellSetType>* const;
if (!called)
{
CellSetType* cellSet = detail::DynamicCellSetTryCast<CellSetType>(this->CellSetContainer);
if (cellSet != nullptr)
downcastType downcastContainer = dynamic_cast<downcastType>(this->Container);
if (downcastContainer)
{
this->Function(*cellSet);
this->FoundCast = true;
f(downcastContainer->Item, std::forward<Args>(args)...);
called = true;
}
}
}
private:
void operator=(const DynamicCellSetTryCellSet<Functor>&) = delete;
const vtkm::cont::internal::SimplePolymorphicContainerBase* const Container;
};
} // namespace detail
template <typename CellSetList>
template <typename Functor>
VTKM_CONT void DynamicCellSetBase<CellSetList>::CastAndCall(const Functor& f) const
template <typename Functor, typename... Args>
VTKM_CONT void DynamicCellSetBase<CellSetList>::CastAndCall(const Functor& f, Args&&... args) const
{
using TryCellSetType = detail::DynamicCellSetTryCellSet<Functor>;
TryCellSetType tryCellSet = TryCellSetType(this->CellSetContainer.get(), f);
vtkm::ListForEach(tryCellSet, CellSetList());
if (!tryCellSet.FoundCast)
bool called = false;
detail::DynamicCellSetTry tryCellSet(this->CellSetContainer.get());
vtkm::ListForEach(tryCellSet, CellSetList{}, f, called, std::forward<Args>(args)...);
if (!called)
{
throw vtkm::cont::ErrorBadValue("Could not find appropriate cast for cell set.");
}

4
vtkm/cont/Error.h
View File

@ -27,7 +27,7 @@
#include <exception>
#include <string>
#include <vtkm/internal/ExportMacros.h> // For VTKM_OVERRIDE
#include <vtkm/internal/ExportMacros.h> // For VTKM_NOEXCEPT
namespace vtkm
{
@ -55,7 +55,7 @@ public:
#endif
// For std::exception compatibility:
const char* what() const VTKM_NOEXCEPT VTKM_OVERRIDE { return this->Message.c_str(); }
const char* what() const VTKM_NOEXCEPT override { return this->Message.c_str(); }
protected:
Error() {}

6
vtkm/cont/Field.h
View File

@ -401,10 +401,10 @@ private:
}
};
template <typename Functor>
void CastAndCall(const vtkm::cont::Field& field, const Functor& f)
template <typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::Field& field, const Functor& f, Args&&... args)
{
field.GetData().CastAndCall(f);
field.GetData().CastAndCall(f, std::forward<Args>(args)...);
}
namespace internal

29
vtkm/cont/ImplicitFunction.cxx
View File

@ -1,29 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/ImplicitFunction.h>
namespace vtkm
{
namespace cont
{
ImplicitFunction::~ImplicitFunction() = default;
}
} // vtkm::cont

286
vtkm/cont/ImplicitFunction.h
View File

@ -1,286 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_ImplicitFunction_h
#define vtk_m_cont_ImplicitFunction_h
#include <vtkm/internal/Configure.h>
#include <vtkm/cont/VirtualObjectCache.h>
#include <vtkm/exec/ImplicitFunction.h>
#include <memory>
namespace vtkm
{
namespace cont
{
class VTKM_CONT_EXPORT ImplicitFunction
{
public:
virtual ~ImplicitFunction();
template <typename DeviceAdapter>
vtkm::exec::ImplicitFunction PrepareForExecution(DeviceAdapter device) const
{
if (!this->Cache->GetValid())
{
this->SetDefaultDevices();
}
return this->Cache->GetVirtualObject(device);
}
void Modified() { this->Cache->SetRefreshFlag(true); }
protected:
using CacheType = vtkm::cont::VirtualObjectCache<vtkm::exec::ImplicitFunction>;
ImplicitFunction()
: Cache(new CacheType)
{
}
ImplicitFunction(ImplicitFunction&& other)
: Cache(std::move(other.Cache))
{
}
ImplicitFunction& operator=(ImplicitFunction&& other)
{
if (this != &other)
{
this->Cache = std::move(other.Cache);
}
return *this;
}
virtual void SetDefaultDevices() const = 0;
std::unique_ptr<CacheType> Cache;
};
template <typename Derived>
class VTKM_ALWAYS_EXPORT ImplicitFunctionImpl : public ImplicitFunction
{
public:
template <typename DeviceAdapterList>
void ResetDevices(DeviceAdapterList devices)
{
this->Cache->Bind(static_cast<const Derived*>(this), devices);
}
protected:
ImplicitFunctionImpl() = default;
ImplicitFunctionImpl(const ImplicitFunctionImpl&)
: ImplicitFunction()
{
}
// Cannot default due to a bug in VS2013
ImplicitFunctionImpl(ImplicitFunctionImpl&& other)
: ImplicitFunction(std::move(other))
{
}
ImplicitFunctionImpl& operator=(const ImplicitFunctionImpl&) { return *this; }
// Cannot default due to a bug in VS2013
ImplicitFunctionImpl& operator=(ImplicitFunctionImpl&& other)
{
ImplicitFunction::operator=(std::move(other));
return *this;
}
void SetDefaultDevices() const override { this->Cache->Bind(static_cast<const Derived*>(this)); }
};
//============================================================================
// ImplicitFunctions:
//============================================================================
/// \brief Implicit function for a box
class VTKM_ALWAYS_EXPORT Box : public ImplicitFunctionImpl<Box>
{
public:
Box();
Box(vtkm::Vec<FloatDefault, 3> minPoint, vtkm::Vec<FloatDefault, 3> maxPoint);
Box(FloatDefault xmin,
FloatDefault xmax,
FloatDefault ymin,
FloatDefault ymax,
FloatDefault zmin,
FloatDefault zmax);
void SetMinPoint(const vtkm::Vec<FloatDefault, 3>& point);
void SetMaxPoint(const vtkm::Vec<FloatDefault, 3>& point);
const vtkm::Vec<FloatDefault, 3>& GetMinPoint() const;
const vtkm::Vec<FloatDefault, 3>& GetMaxPoint() const;
VTKM_EXEC_CONT
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const;
private:
vtkm::Vec<FloatDefault, 3> MinPoint;
vtkm::Vec<FloatDefault, 3> MaxPoint;
};
//============================================================================
/// \brief Implicit function for a cylinder
class VTKM_ALWAYS_EXPORT Cylinder : public ImplicitFunctionImpl<Cylinder>
{
public:
Cylinder();
Cylinder(const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault radius);
Cylinder(const vtkm::Vec<FloatDefault, 3>& center,
const vtkm::Vec<FloatDefault, 3>& axis,
FloatDefault radius);
void SetCenter(const vtkm::Vec<FloatDefault, 3>& center);
void SetAxis(const vtkm::Vec<FloatDefault, 3>& axis);
void SetRadius(FloatDefault radius);
const vtkm::Vec<FloatDefault, 3>& GetCenter() const;
const vtkm::Vec<FloatDefault, 3>& GetAxis() const;
FloatDefault GetRadius() const;
VTKM_EXEC_CONT
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const;
private:
vtkm::Vec<FloatDefault, 3> Center;
vtkm::Vec<FloatDefault, 3> Axis;
FloatDefault Radius;
};
//============================================================================
/// \brief Implicit function for a frustum
class VTKM_ALWAYS_EXPORT Frustum : public ImplicitFunctionImpl<Frustum>
{
public:
Frustum();
Frustum(const vtkm::Vec<FloatDefault, 3> points[6], const vtkm::Vec<FloatDefault, 3> normals[6]);
explicit Frustum(const vtkm::Vec<FloatDefault, 3> points[8]);
void SetPlanes(const vtkm::Vec<FloatDefault, 3> points[6],
const vtkm::Vec<FloatDefault, 3> normals[6]);
void SetPlane(int idx, vtkm::Vec<FloatDefault, 3>& point, vtkm::Vec<FloatDefault, 3>& normal);
void GetPlanes(vtkm::Vec<FloatDefault, 3> points[6], vtkm::Vec<FloatDefault, 3> normals[6]) const;
const vtkm::Vec<FloatDefault, 3>* GetPoints() const;
const vtkm::Vec<FloatDefault, 3>* GetNormals() const;
// The points should be specified in the order of hex-cell vertices
void CreateFromPoints(const vtkm::Vec<FloatDefault, 3> points[8]);
VTKM_EXEC_CONT
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault, FloatDefault, FloatDefault) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>&) const;
private:
vtkm::Vec<FloatDefault, 3> Points[6];
vtkm::Vec<FloatDefault, 3> Normals[6];
};
//============================================================================
/// \brief Implicit function for a plane
class VTKM_ALWAYS_EXPORT Plane : public ImplicitFunctionImpl<Plane>
{
public:
Plane();
explicit Plane(const vtkm::Vec<FloatDefault, 3>& normal);
Plane(const vtkm::Vec<FloatDefault, 3>& origin, const vtkm::Vec<FloatDefault, 3>& normal);
void SetOrigin(const vtkm::Vec<FloatDefault, 3>& origin);
void SetNormal(const vtkm::Vec<FloatDefault, 3>& normal);
const vtkm::Vec<FloatDefault, 3>& GetOrigin() const;
const vtkm::Vec<FloatDefault, 3>& GetNormal() const;
VTKM_EXEC_CONT
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault, FloatDefault, FloatDefault) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>&) const;
private:
vtkm::Vec<FloatDefault, 3> Origin;
vtkm::Vec<FloatDefault, 3> Normal;
};
//============================================================================
/// \brief Implicit function for a sphere
class VTKM_ALWAYS_EXPORT Sphere : public ImplicitFunctionImpl<Sphere>
{
public:
Sphere();
explicit Sphere(FloatDefault radius);
Sphere(vtkm::Vec<FloatDefault, 3> center, FloatDefault radius);
void SetRadius(FloatDefault radius);
void SetCenter(const vtkm::Vec<FloatDefault, 3>& center);
FloatDefault GetRadius() const;
const vtkm::Vec<FloatDefault, 3>& GetCenter() const;
VTKM_EXEC_CONT
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const;
VTKM_EXEC_CONT
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>& x) const;
private:
FloatDefault Radius;
vtkm::Vec<FloatDefault, 3> Center;
};
}
} // vtkm::cont
#include <vtkm/cont/ImplicitFunction.hxx>
#endif // vtk_m_cont_ImplicitFunction_h

640
vtkm/cont/ImplicitFunction.hxx
View File

@ -1,640 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/VectorAnalysis.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <algorithm>
#include <limits>
namespace vtkm
{
namespace cont
{
//============================================================================
inline Box::Box()
: MinPoint(vtkm::Vec<FloatDefault, 3>(FloatDefault(0)))
, MaxPoint(vtkm::Vec<FloatDefault, 3>(FloatDefault(1)))
{
}
inline Box::Box(vtkm::Vec<FloatDefault, 3> minPoint, vtkm::Vec<FloatDefault, 3> maxPoint)
: MinPoint(minPoint)
, MaxPoint(maxPoint)
{
}
inline Box::Box(FloatDefault xmin,
FloatDefault xmax,
FloatDefault ymin,
FloatDefault ymax,
FloatDefault zmin,
FloatDefault zmax)
{
MinPoint[0] = xmin;
MaxPoint[0] = xmax;
MinPoint[1] = ymin;
MaxPoint[1] = ymax;
MinPoint[2] = zmin;
MaxPoint[2] = zmax;
}
inline void Box::SetMinPoint(const vtkm::Vec<FloatDefault, 3>& point)
{
this->MinPoint = point;
this->Modified();
}
inline void Box::SetMaxPoint(const vtkm::Vec<FloatDefault, 3>& point)
{
this->MaxPoint = point;
this->Modified();
}
inline const vtkm::Vec<FloatDefault, 3>& Box::GetMinPoint() const
{
return this->MinPoint;
}
inline const vtkm::Vec<FloatDefault, 3>& Box::GetMaxPoint() const
{
return this->MaxPoint;
}
VTKM_EXEC_CONT
inline FloatDefault Box::Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->Value(vtkm::Vec<vtkm::FloatDefault, 3>(x, y, z));
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Box::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const
{
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));
}
VTKM_EXEC_CONT
inline FloatDefault Box::Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
FloatDefault minDistance = vtkm::NegativeInfinity32();
FloatDefault diff, t, dist;
FloatDefault distance = FloatDefault(0.0);
vtkm::IdComponent inside = 1;
for (vtkm::IdComponent d = 0; d < 3; d++)
{
diff = this->MaxPoint[d] - this->MinPoint[d];
if (diff != FloatDefault(0.0))
{
t = (x[d] - this->MinPoint[d]) / diff;
// Outside before the box
if (t < FloatDefault(0.0))
{
inside = 0;
dist = this->MinPoint[d] - x[d];
}
// Outside after the box
else if (t > FloatDefault(1.0))
{
inside = 0;
dist = x[d] - this->MaxPoint[d];
}
else
{
// Inside the box in lower half
if (t <= FloatDefault(0.5))
{
dist = MinPoint[d] - x[d];
}
// Inside the box in upper half
else
{
dist = x[d] - MaxPoint[d];
}
if (dist > minDistance)
{
minDistance = dist;
}
}
}
else
{
dist = vtkm::Abs(x[d] - MinPoint[d]);
if (dist > FloatDefault(0.0))
{
inside = 0;
}
}
if (dist > FloatDefault(0.0))
{
distance += dist * dist;
}
}
distance = vtkm::Sqrt(distance);
if (inside)
{
return minDistance;
}
else
{
return distance;
}
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Box::Gradient(const vtkm::Vec<FloatDefault, 3>& x) const
{
vtkm::IdComponent minAxis = 0;
FloatDefault dist = 0.0;
FloatDefault minDist = vtkm::Infinity32();
vtkm::Vec<vtkm::IdComponent, 3> location;
vtkm::Vec<FloatDefault, 3> normal(FloatDefault(0));
vtkm::Vec<FloatDefault, 3> inside(FloatDefault(0));
vtkm::Vec<FloatDefault, 3> outside(FloatDefault(0));
vtkm::Vec<FloatDefault, 3> center((this->MaxPoint + this->MinPoint) * FloatDefault(0.5));
// Compute the location of the point with respect to the box
// Point will lie in one of 27 separate regions around or within the box
// Gradient vector is computed differently in each of the regions.
for (vtkm::IdComponent d = 0; d < 3; d++)
{
if (x[d] < this->MinPoint[d])
{
// Outside the box low end
location[d] = 0;
outside[d] = -1.0;
}
else if (x[d] > this->MaxPoint[d])
{
// Outside the box high end
location[d] = 2;
outside[d] = 1.0;
}
else
{
location[d] = 1;
if (x[d] <= center[d])
{
// Inside the box low end
dist = x[d] - this->MinPoint[d];
inside[d] = -1.0;
}
else
{
// Inside the box high end
dist = this->MaxPoint[d] - x[d];
inside[d] = 1.0;
}
if (dist < minDist) // dist is negative
{
minDist = dist;
minAxis = d;
}
}
}
vtkm::Id indx = location[0] + 3 * location[1] + 9 * location[2];
switch (indx)
{
// verts - gradient points away from center point
case 0:
case 2:
case 6:
case 8:
case 18:
case 20:
case 24:
case 26:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
normal[d] = x[d] - center[d];
}
vtkm::Normalize(normal);
break;
// edges - gradient points out from axis of cube
case 1:
case 3:
case 5:
case 7:
case 9:
case 11:
case 15:
case 17:
case 19:
case 21:
case 23:
case 25:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
if (outside[d] != 0.0)
{
normal[d] = x[d] - center[d];
}
else
{
normal[d] = 0.0;
}
}
vtkm::Normalize(normal);
break;
// faces - gradient points perpendicular to face
case 4:
case 10:
case 12:
case 14:
case 16:
case 22:
for (vtkm::IdComponent d = 0; d < 3; d++)
{
normal[d] = outside[d];
}
break;
// interior - gradient is perpendicular to closest face
case 13:
normal[0] = normal[1] = normal[2] = 0.0;
normal[minAxis] = inside[minAxis];
break;
default:
VTKM_ASSERT(false);
break;
}
return normal;
}
//============================================================================
inline Cylinder::Cylinder()
: Center(FloatDefault(0))
, Axis(vtkm::make_Vec(FloatDefault(1), FloatDefault(0), FloatDefault(0)))
, Radius(FloatDefault(0.2))
{
}
inline Cylinder::Cylinder(const vtkm::Vec<FloatDefault, 3>& axis, FloatDefault radius)
: Center(FloatDefault(0))
, Axis(vtkm::Normal(axis))
, Radius(radius)
{
}
inline Cylinder::Cylinder(const vtkm::Vec<FloatDefault, 3>& center,
const vtkm::Vec<FloatDefault, 3>& axis,
FloatDefault radius)
: Center(center)
, Axis(vtkm::Normal(axis))
, Radius(radius)
{
}
inline void Cylinder::SetCenter(const vtkm::Vec<FloatDefault, 3>& center)
{
this->Center = center;
this->Modified();
}
inline void Cylinder::SetAxis(const vtkm::Vec<FloatDefault, 3>& axis)
{
this->Axis = vtkm::Normal(axis);
this->Modified();
}
inline void Cylinder::SetRadius(FloatDefault radius)
{
this->Radius = radius;
this->Modified();
}
inline const vtkm::Vec<FloatDefault, 3>& Cylinder::GetCenter() const
{
return this->Center;
}
inline const vtkm::Vec<FloatDefault, 3>& Cylinder::GetAxis() const
{
return this->Axis;
}
inline FloatDefault Cylinder::GetRadius() const
{
return this->Radius;
}
VTKM_EXEC_CONT
inline FloatDefault Cylinder::Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
vtkm::Vec<FloatDefault, 3> x2c = x - this->Center;
FloatDefault proj = vtkm::dot(this->Axis, x2c);
return vtkm::dot(x2c, x2c) - (proj * proj) - (this->Radius * this->Radius);
}
VTKM_EXEC_CONT
inline FloatDefault Cylinder::Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->Value(vtkm::Vec<vtkm::FloatDefault, 3>(x, y, z));
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Cylinder::Gradient(const vtkm::Vec<FloatDefault, 3>& x) const
{
vtkm::Vec<FloatDefault, 3> x2c = x - this->Center;
FloatDefault t = this->Axis[0] * x2c[0] + this->Axis[1] * x2c[1] + this->Axis[2] * x2c[2];
vtkm::Vec<FloatDefault, 3> closestPoint = this->Center + (this->Axis * t);
return (x - closestPoint) * FloatDefault(2);
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Cylinder::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const
{
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));
}
//============================================================================
inline Frustum::Frustum()
{
std::fill(this->Points, this->Points + 6, vtkm::Vec<FloatDefault, 3>{});
std::fill(this->Normals, this->Normals + 6, vtkm::Vec<FloatDefault, 3>{});
}
inline Frustum::Frustum(const vtkm::Vec<FloatDefault, 3> points[6],
const vtkm::Vec<FloatDefault, 3> normals[6])
{
std::copy(points, points + 6, this->Points);
std::copy(normals, normals + 6, this->Normals);
}
inline Frustum::Frustum(const vtkm::Vec<FloatDefault, 3> points[8])
{
this->CreateFromPoints(points);
}
inline void Frustum::SetPlanes(const vtkm::Vec<FloatDefault, 3> points[6],
const vtkm::Vec<FloatDefault, 3> normals[6])
{
std::copy(points, points + 6, this->Points);
std::copy(normals, normals + 6, this->Normals);
this->Modified();
}
inline void Frustum::GetPlanes(vtkm::Vec<FloatDefault, 3> points[6],
vtkm::Vec<FloatDefault, 3> normals[6]) const
{
std::copy(this->Points, this->Points + 6, points);
std::copy(this->Normals, this->Normals + 6, normals);
}
inline const vtkm::Vec<FloatDefault, 3>* Frustum::GetPoints() const
{
return this->Points;
}
inline const vtkm::Vec<FloatDefault, 3>* Frustum::GetNormals() const
{
return this->Normals;
}
inline void Frustum::SetPlane(int idx,
vtkm::Vec<FloatDefault, 3>& point,
vtkm::Vec<FloatDefault, 3>& normal)
{
if (idx < 0 || idx >= 6)
{
std::string msg = "Plane idx ";
msg += std::to_string(idx) + " is out of range [0, 5]";
throw vtkm::cont::ErrorBadValue(msg);
}
this->Points[idx] = point;
this->Normals[idx] = normal;
this->Modified();
}
inline void Frustum::CreateFromPoints(const vtkm::Vec<FloatDefault, 3> points[8])
{
// XXX(clang-format-3.9): 3.8 is silly. 3.9 makes it look like this.
// clang-format off
int planes[6][3] = {
{ 3, 2, 0 }, { 4, 5, 7 }, { 0, 1, 4 }, { 1, 2, 5 }, { 2, 3, 6 }, { 3, 0, 7 }
};
// clang-format on
for (int i = 0; i < 6; ++i)
{
auto& v0 = points[planes[i][0]];
auto& v1 = points[planes[i][1]];
auto& v2 = points[planes[i][2]];
this->Points[i] = v0;
this->Normals[i] = vtkm::Normal(vtkm::Cross(v2 - v0, v1 - v0));
this->Modified();
}
}
VTKM_EXEC_CONT
inline FloatDefault Frustum::Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
FloatDefault maxVal = -std::numeric_limits<FloatDefault>::max();
for (int i = 0; i < 6; ++i)
{
auto& p = this->Points[i];
auto& n = this->Normals[i];
FloatDefault val = ((x - p[0]) * n[0]) + ((y - p[1]) * n[1]) + ((z - p[2]) * n[2]);
maxVal = vtkm::Max(maxVal, val);
}
return maxVal;
}
VTKM_EXEC_CONT
inline FloatDefault Frustum::Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->Value(x[0], x[1], x[2]);
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Frustum::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const
{
FloatDefault maxVal = -std::numeric_limits<FloatDefault>::max();
int maxValIdx = 0;
for (int i = 0; i < 6; ++i)
{
auto& p = this->Points[i];
auto& n = this->Normals[i];
FloatDefault val = ((x - p[0]) * n[0]) + ((y - p[1]) * n[1]) + ((z - p[2]) * n[2]);
if (val > maxVal)
{
maxVal = val;
maxValIdx = i;
}
}
return this->Normals[maxValIdx];
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Frustum::Gradient(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->Gradient(x[0], x[1], x[2]);
}
//============================================================================
inline Plane::Plane()
: Origin(FloatDefault(0))
, Normal(FloatDefault(0), FloatDefault(0), FloatDefault(1))
{
}
inline Plane::Plane(const vtkm::Vec<FloatDefault, 3>& normal)
: Origin(FloatDefault(0))
, Normal(normal)
{
}
inline Plane::Plane(const vtkm::Vec<FloatDefault, 3>& origin,
const vtkm::Vec<FloatDefault, 3>& normal)
: Origin(origin)
, Normal(normal)
{
}
inline void Plane::SetOrigin(const vtkm::Vec<FloatDefault, 3>& origin)
{
this->Origin = origin;
this->Modified();
}
inline void Plane::SetNormal(const vtkm::Vec<FloatDefault, 3>& normal)
{
this->Normal = normal;
this->Modified();
}
inline const vtkm::Vec<FloatDefault, 3>& Plane::GetOrigin() const
{
return this->Origin;
}
inline const vtkm::Vec<FloatDefault, 3>& Plane::GetNormal() const
{
return this->Normal;
}
VTKM_EXEC_CONT
inline FloatDefault Plane::Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return ((x - this->Origin[0]) * this->Normal[0]) + ((y - this->Origin[1]) * this->Normal[1]) +
((z - this->Origin[2]) * this->Normal[2]);
}
VTKM_EXEC_CONT
inline FloatDefault Plane::Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->Value(x[0], x[1], x[2]);
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Plane::Gradient(FloatDefault, FloatDefault, FloatDefault) const
{
return this->Normal;
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Plane::Gradient(const vtkm::Vec<FloatDefault, 3>&) const
{
return this->Normal;
}
//============================================================================
inline Sphere::Sphere()
: Radius(FloatDefault(0.2))
, Center(FloatDefault(0))
{
}
inline Sphere::Sphere(FloatDefault radius)
: Radius(radius)
, Center(FloatDefault(0))
{
}
inline Sphere::Sphere(vtkm::Vec<FloatDefault, 3> center, FloatDefault radius)
: Radius(radius)
, Center(center)
{
}
inline void Sphere::SetRadius(FloatDefault radius)
{
this->Radius = radius;
this->Modified();
}
inline void Sphere::SetCenter(const vtkm::Vec<FloatDefault, 3>& center)
{
this->Center = center;
this->Modified();
}
inline FloatDefault Sphere::GetRadius() const
{
return this->Radius;
}
inline const vtkm::Vec<FloatDefault, 3>& Sphere::GetCenter() const
{
return this->Center;
}
VTKM_EXEC_CONT
inline FloatDefault Sphere::Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return ((x - this->Center[0]) * (x - this->Center[0]) +
(y - this->Center[1]) * (y - this->Center[1]) +
(z - this->Center[2]) * (z - this->Center[2])) -
(this->Radius * this->Radius);
}
VTKM_EXEC_CONT
inline FloatDefault Sphere::Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->Value(x[0], x[1], x[2]);
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Sphere::Gradient(FloatDefault x,
FloatDefault y,
FloatDefault z) const
{
return this->Gradient(vtkm::Vec<FloatDefault, 3>(x, y, z));
}
VTKM_EXEC_CONT
inline vtkm::Vec<FloatDefault, 3> Sphere::Gradient(const vtkm::Vec<FloatDefault, 3>& x) const
{
return FloatDefault(2) * (x - this->Center);
}
}
} // vtkm::cont

78
vtkm/cont/ImplicitFunctionHandle.h Normal file
View File

@ -0,0 +1,78 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_ImplicitFunctionHandle_h
#define vtk_m_cont_ImplicitFunctionHandle_h
#include <vtkm/ImplicitFunction.h>
#include <vtkm/cont/VirtualObjectHandle.h>
namespace vtkm
{
namespace cont
{
class VTKM_ALWAYS_EXPORT ImplicitFunctionHandle
: public vtkm::cont::VirtualObjectHandle<vtkm::ImplicitFunction>
{
private:
using Superclass = vtkm::cont::VirtualObjectHandle<vtkm::ImplicitFunction>;
public:
ImplicitFunctionHandle() = default;
template <typename ImplicitFunctionType,
typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
explicit ImplicitFunctionHandle(ImplicitFunctionType* function,
bool acquireOwnership = true,
DeviceAdapterList devices = DeviceAdapterList())
: Superclass(function, acquireOwnership, devices)
{
}
};
template <typename ImplicitFunctionType,
typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
VTKM_CONT ImplicitFunctionHandle
make_ImplicitFunctionHandle(ImplicitFunctionType&& func,
DeviceAdapterList devices = DeviceAdapterList())
{
using IFType = typename std::remove_reference<ImplicitFunctionType>::type;
return ImplicitFunctionHandle(
new IFType(std::forward<ImplicitFunctionType>(func)), true, devices);
}
template <typename ImplicitFunctionType, typename... Args>
VTKM_CONT ImplicitFunctionHandle make_ImplicitFunctionHandle(Args&&... args)
{
return ImplicitFunctionHandle(new ImplicitFunctionType(std::forward<Args>(args)...),
true,
VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
}
template <typename ImplicitFunctionType, typename DeviceAdapterList, typename... Args>
VTKM_CONT ImplicitFunctionHandle make_ImplicitFunctionHandle(Args&&... args)
{
return ImplicitFunctionHandle(
new ImplicitFunctionType(std::forward<Args>(args)...), true, DeviceAdapterList());
}
}
} // vtkm::cont
#endif // vtk_m_cont_ImplicitFunctionHandle_h

21
vtkm/cont/RuntimeDeviceTracker.cxx
View File

@ -29,7 +29,10 @@
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
#include <algorithm>
#include <map>
#include <mutex>
#include <sstream>
#include <thread>
#define VTKM_MAX_DEVICE_ADAPTER_ID 8
@ -155,8 +158,22 @@ void RuntimeDeviceTracker::ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId,
VTKM_CONT
vtkm::cont::RuntimeDeviceTracker GetGlobalRuntimeDeviceTracker()
{
static vtkm::cont::RuntimeDeviceTracker globalTracker;
return globalTracker;
static std::mutex mtx;
static std::map<std::thread::id, vtkm::cont::RuntimeDeviceTracker> globalTrackers;
std::thread::id this_id = std::this_thread::get_id();
std::unique_lock<std::mutex> lock(mtx);
auto iter = globalTrackers.find(this_id);
if (iter != globalTrackers.end())
{
return iter->second;
}
else
{
vtkm::cont::RuntimeDeviceTracker tracker;
globalTrackers[this_id] = tracker;
return tracker;
}
}
}
} // namespace vtkm::cont

17
vtkm/cont/RuntimeDeviceTracker.h
View File

@ -74,6 +74,17 @@ public:
this->SetDeviceState(Traits::GetId(), Traits::GetName(), false);
}
/// Report a failure to allocate memory on a device, this will flag the
/// device as being unusable for all future invocations of the instance of
/// the filter.
///
VTKM_CONT void ReportAllocationFailure(vtkm::Int8 deviceId,
const std::string& name,
const vtkm::cont::ErrorBadAllocation&)
{
this->SetDeviceState(deviceId, name, false);
}
/// Reset the tracker for the given device. This will discard any updates
/// caused by reported failures
///
@ -196,13 +207,13 @@ private:
bool runtimeExists);
};
/// \brief Get the global \c RuntimeDeviceTracker.
/// \brief Get the \c RuntimeDeviceTracker for the current thread.
///
/// Many features in VTK-m will attempt to run algorithms on the "best
/// available device." This often is determined at runtime as failures in
/// one device are recorded and that device is disabled. To prevent having
/// to check over and over again, VTK-m features generally use the global
/// device adapter so that these choices are marked and shared.
/// to check over and over again, VTK-m uses per thread runtime device tracker
/// so that these choices are marked and shared.
///
VTKM_CONT_EXPORT
VTKM_CONT

77
vtkm/cont/TryExecute.cxx Normal file
View File

@ -0,0 +1,77 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2016 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2016 UT-Battelle, LLC.
// Copyright 2016 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/ErrorBadAllocation.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/TryExecute.h>
namespace vtkm
{
namespace cont
{
namespace detail
{
void HandleTryExecuteException(vtkm::Int8 deviceId,
const std::string& name,
vtkm::cont::RuntimeDeviceTracker& tracker)
{
try
{
//re-throw the last exception
throw;
}
catch (vtkm::cont::ErrorBadAllocation& e)
{
std::cerr << "caught ErrorBadAllocation " << e.GetMessage() << std::endl;
//currently we only consider OOM errors worth disabling a device for
//than we fallback to another device
tracker.ReportAllocationFailure(deviceId, name, e);
}
catch (vtkm::cont::ErrorBadType& e)
{
//should bad type errors should stop the execution, instead of
//deferring to another device adapter?
std::cerr << "caught ErrorBadType : " << e.GetMessage() << std::endl;
}
catch (vtkm::cont::ErrorBadValue& e)
{
//should bad value errors should stop the filter, instead of deferring
//to another device adapter?
std::cerr << "caught ErrorBadValue : " << e.GetMessage() << std::endl;
}
catch (vtkm::cont::Error& e)
{
//general errors should be caught and let us try the next device adapter.
std::cerr << "exception is: " << e.GetMessage() << std::endl;
}
catch (std::exception& e)
{
std::cerr << "caught standard exception: " << e.what() << std::endl;
}
catch (...)
{
std::cerr << "unknown exception caught" << std::endl;
}
}
}
}
}

333
vtkm/cont/TryExecute.h
View File

@ -21,10 +21,6 @@
#define vtk_m_cont_TryExecute_h
#include <vtkm/cont/DeviceAdapterListTag.h>
#include <vtkm/cont/ErrorBadAllocation.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
namespace vtkm
@ -35,123 +31,171 @@ namespace cont
namespace detail
{
template <typename Functor, typename Device, bool DeviceAdapterValid>
struct TryExecuteRunIfValid;
VTKM_CONT_EXPORT void HandleTryExecuteException(vtkm::Int8,
const std::string&,
vtkm::cont::RuntimeDeviceTracker&);
template <typename Functor, typename Device>
struct TryExecuteRunIfValid<Functor, Device, false>
template <typename DeviceTag, typename Functor, typename... Args>
bool TryExecuteIfValid(std::true_type,
DeviceTag tag,
Functor&& f,
vtkm::cont::RuntimeDeviceTracker& tracker,
Args&&... args)
{
VTKM_CONT
static bool Run(Functor&, const vtkm::cont::RuntimeDeviceTracker&) { return false; }
};
template <typename Functor, typename Device>
struct TryExecuteRunIfValid<Functor, Device, true>
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
VTKM_CONT
static bool Run(Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker)
if (tracker.CanRunOn(tag))
{
if (tracker.CanRunOn(Device()))
try
{
try
{
return functor(Device());
}
catch (vtkm::cont::ErrorBadAllocation& e)
{
std::cerr << "caught ErrorBadAllocation " << e.GetMessage() << std::endl;
//currently we only consider OOM errors worth disabling a device for
//than we fallback to another device
tracker.ReportAllocationFailure(Device(), e);
}
catch (vtkm::cont::ErrorBadType& e)
{
//should bad type errors should stop the execution, instead of
//deferring to another device adapter?
std::cerr << "caught ErrorBadType : " << e.GetMessage() << std::endl;
}
catch (vtkm::cont::ErrorBadValue& e)
{
//should bad value errors should stop the filter, instead of deferring
//to another device adapter?
std::cerr << "caught ErrorBadValue : " << e.GetMessage() << std::endl;
}
catch (vtkm::cont::Error& e)
{
//general errors should be caught and let us try the next device adapter.
std::cerr << "exception is: " << e.GetMessage() << std::endl;
}
catch (std::exception& e)
{
std::cerr << "caught standard exception: " << e.what() << std::endl;
}
catch (...)
{
std::cerr << "unknown exception caught" << std::endl;
}
return f(tag, std::forward<Args>(args)...);
}
// If we are here, then the functor was either never run or failed.
return false;
}
};
template <typename FunctorType>
struct TryExecuteImpl
{
// Warning, these are a references. Make sure referenced objects do not go
// out of scope.
FunctorType& Functor;
vtkm::cont::RuntimeDeviceTracker Tracker;
bool Success;
VTKM_CONT
TryExecuteImpl(
FunctorType& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
: Functor(functor)
, Tracker(tracker)
, Success(false)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device)
{
if (!this->Success)
catch (...)
{
using DeviceTraits = vtkm::cont::DeviceAdapterTraits<Device>;
this->Success = detail::TryExecuteRunIfValid<FunctorType, Device, DeviceTraits::Valid>::Run(
this->Functor, this->Tracker);
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceTag>;
HandleTryExecuteException(Traits::GetId(), Traits::GetName(), tracker);
}
return this->Success;
}
private:
void operator=(const TryExecuteImpl<FunctorType>&) = delete;
// If we are here, then the functor was either never run or failed.
return false;
}
template <typename DeviceTag, typename Functor, typename... Args>
bool TryExecuteIfValid(std::false_type,
DeviceTag,
Functor&&,
vtkm::cont::RuntimeDeviceTracker&,
Args&&...)
{
return false;
}
struct TryExecuteWrapper
{
template <typename DeviceTag, typename Functor, typename... Args>
void operator()(DeviceTag tag,
Functor&& f,
vtkm::cont::RuntimeDeviceTracker& tracker,
bool& ran,
Args&&... args) const
{
if (!ran)
{
using DeviceTraits = vtkm::cont::DeviceAdapterTraits<DeviceTag>;
ran = TryExecuteIfValid(std::integral_constant<bool, DeviceTraits::Valid>(),
tag,
std::forward<Functor>(f),
std::forward<decltype(tracker)>(tracker),
std::forward<Args>(args)...);
}
}
};
template <typename Functor, typename DeviceList, typename... Args>
VTKM_CONT bool TryExecuteImpl(Functor&& functor,
std::true_type,
std::true_type,
vtkm::cont::RuntimeDeviceTracker& tracker,
DeviceList list,
Args&&... args)
{
bool success = false;
detail::TryExecuteWrapper task;
vtkm::ListForEach(
task, list, std::forward<Functor>(functor), tracker, success, std::forward<Args>(args)...);
return success;
}
template <typename Functor, typename... Args>
VTKM_CONT bool TryExecuteImpl(Functor&& functor, std::false_type, std::false_type, Args&&... args)
{
bool success = false;
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
detail::TryExecuteWrapper task;
vtkm::ListForEach(task,
VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG(),
std::forward<Functor>(functor),
tracker,
success,
std::forward<Args>(args)...);
return success;
}
template <typename Functor, typename Arg1, typename... Args>
VTKM_CONT bool TryExecuteImpl(Functor&& functor,
std::true_type t,
std::false_type,
Arg1&& arg1,
Args&&... args)
{
return TryExecuteImpl(std::forward<Functor>(functor),
t,
t,
std::forward<Arg1>(arg1),
VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG(),
std::forward<Args>(args)...);
}
template <typename Functor, typename Arg1, typename... Args>
VTKM_CONT bool TryExecuteImpl(Functor&& functor,
std::false_type,
std::true_type t,
Arg1&& arg1,
Args&&... args)
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
return TryExecuteImpl(std::forward<Functor>(functor),
t,
t,
tracker,
std::forward<Arg1>(arg1),
std::forward<Args>(args)...);
}
} // namespace detail
///@{
/// \brief Try to execute a functor on a list of devices until one succeeds.
///
/// This function takes a functor and a list of devices. It then tries to run
/// the functor for each device (in the order given in the list) until the
/// execution succeeds.
/// This function takes a functor and optionally a list of devices and \c RuntimeDeviceTracker.
/// It then tries to run the functor for each device (in the order given in the list) until the
/// execution succeeds. The optional \c RuntimeDeviceTracker allows for monitoring for certain
/// failures across calls to TryExecute and skip trying devices with a history of failure.
///
/// The functor parentheses operator should take exactly one argument, which is
/// the \c DeviceAdapterTag to use. The functor should return a \c bool that is
/// \c true if the execution succeeds, \c false if it fails. If an exception is
/// thrown from the functor, then the execution is assumed to have failed.
/// The TryExecute is also able to perfectly forward arbitrary arguments onto the functor.
/// These arguments must be placed after the optional \c RuntimeDeviceTracker, and device adapter
/// list and will passed to the functor in the same order as listed.
///
/// This function also optionally takes a \c RuntimeDeviceTracker, which will
/// monitor for certain failures across calls to TryExecute and skip trying
/// devices with a history of failure.
/// The functor must implement the function call operator ( \c operator() ) with a return type of
/// \c bool and that is \c true if the execution succeeds, \c false if it fails. If an exception
/// is thrown from the functor, then the execution is assumed to have failed. The functor call
/// operator must also take at least one argument being the required \c DeviceAdapterTag to use.
///
/// \code{.cpp}
/// struct TryCallExample
/// {
/// template<typename DeviceList>
/// bool operator()(DeviceList tags, int) const
/// {
/// return true;
/// }
/// };
///
///
/// // Executing without a runtime tracker or device list
/// vtkm::cont::TryExecute(TryCallExample(), int{42});
///
/// // Executing with a runtime tracker
/// auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
/// vtkm::cont::TryExecute(TryCallExample(), tracker, int{42});
///
/// // Executing with a device list
/// using DeviceList = vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial>;
/// vtkm::cont::TryExecute(TryCallExample(), DeviceList(), int{42});
///
/// // Executing with a runtime tracker and device list
/// vtkm::cont::TryExecute(EdgeCaseFunctor(), tracker, DeviceList(), int{42});
///
/// \endcode
///
/// This function returns \c true if the functor succeeded on a device,
/// \c false otherwise.
@ -159,46 +203,57 @@ private:
/// If no device list is specified, then \c VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG
/// is used.
///
template <typename Functor, typename DeviceList>
VTKM_CONT bool TryExecute(const Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker,
DeviceList)
{
detail::TryExecuteImpl<const Functor> internals(functor, tracker);
vtkm::ListForEach(internals, DeviceList());
return internals.Success;
}
template <typename Functor, typename DeviceList>
VTKM_CONT bool TryExecute(Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker, DeviceList)
{
detail::TryExecuteImpl<Functor> internals(functor, tracker);
vtkm::ListForEach(internals, DeviceList());
return internals.Success;
}
template <typename Functor, typename DeviceList>
VTKM_CONT bool TryExecute(const Functor& functor, DeviceList)
{
return vtkm::cont::TryExecute(functor, vtkm::cont::GetGlobalRuntimeDeviceTracker(), DeviceList());
}
template <typename Functor, typename DeviceList>
VTKM_CONT bool TryExecute(Functor& functor, DeviceList)
{
return vtkm::cont::TryExecute(functor, vtkm::cont::GetGlobalRuntimeDeviceTracker(), DeviceList());
}
/// If no \c RuntimeDeviceTracker specified, then \c GetGlobalRuntimeDeviceTracker()
/// is used.
template <typename Functor>
VTKM_CONT bool TryExecute(
const Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
VTKM_CONT bool TryExecute(Functor&& functor)
{
return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
//we haven't been passed either a runtime tracker or a device list
return detail::TryExecuteImpl(functor, std::false_type{}, std::false_type{});
}
template <typename Functor>
VTKM_CONT bool TryExecute(
Functor& functor,
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker())
template <typename Functor, typename Arg1>
VTKM_CONT bool TryExecute(Functor&& functor, Arg1&& arg1)
{
return vtkm::cont::TryExecute(functor, tracker, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
//determine if we are being passed a device adapter or runtime tracker as our argument
using is_deviceAdapter = typename std::is_base_of<vtkm::detail::ListRoot, Arg1>::type;
using is_tracker = typename std::is_base_of<vtkm::cont::RuntimeDeviceTracker,
typename std::remove_reference<Arg1>::type>::type;
return detail::TryExecuteImpl(
functor, is_tracker{}, is_deviceAdapter{}, std::forward<Arg1>(arg1));
}
template <typename Functor, typename Arg1, typename Arg2, typename... Args>
VTKM_CONT bool TryExecute(Functor&& functor, Arg1&& arg1, Arg2&& arg2, Args&&... args)
{
//So arg1 can be runtime or device adapter
//if arg1 is runtime, we need to see if arg2 is device adapter
using is_arg1_tracker =
typename std::is_base_of<vtkm::cont::RuntimeDeviceTracker,
typename std::remove_reference<Arg1>::type>::type;
using is_arg2_devicelist = typename std::is_base_of<vtkm::detail::ListRoot, Arg2>::type;
//We now know what of three states we are currently at
using has_runtime_and_deviceAdapter =
brigand::bool_<is_arg1_tracker::value && is_arg2_devicelist::value>;
using has_just_runtime = brigand::bool_<is_arg1_tracker::value && !is_arg2_devicelist::value>;
using has_just_devicelist = typename std::is_base_of<vtkm::detail::ListRoot, Arg1>::type;
//With this information we can now compute if we have a runtime tracker and/or
//the device adapter and enable the correct flags
using first_true =
brigand::bool_<has_runtime_and_deviceAdapter::value || has_just_runtime::value>;
using second_true =
brigand::bool_<has_runtime_and_deviceAdapter::value || has_just_devicelist::value>;
return detail::TryExecuteImpl(functor,
first_true{},
second_true{},
std::forward<Arg1>(arg1),
std::forward<Arg2>(arg2),
std::forward<Args>(args)...);
}
//@} //block doxygen all TryExecute functions
}
} // namespace vtkm::cont

310
vtkm/cont/VirtualObjectCache.h
View File

@ -1,310 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_VirtualObjectCache_h
#define vtk_m_cont_VirtualObjectCache_h
#include <vtkm/cont/DeviceAdapterListTag.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/internal/DeviceAdapterTag.h>
#include <vtkm/cont/internal/VirtualObjectTransfer.h>
#include <array>
#include <type_traits>
#define VTKM_MAX_DEVICE_ADAPTER_ID 8
namespace vtkm
{
namespace cont
{
/// \brief Implements VTK-m's execution side <em> Virtual Methods </em>
/// functionality.
///
/// The template parameter \c VirtualObject is the class that acts as the
/// interface. Following is a method for implementing such classes:
/// 1. Create a <tt>const void*</tt> member variable that will hold a
/// reference to the target class.
/// 2. For each virtual-like method:
/// a. Create a typedef for a function with the same signature,
/// except for an extra <tt>const void*</tt> argument.
/// b. Create a function pointer member variable with the type of the
/// associated typedef from 2a.
/// c. Create an implementation of the method, that calls the associated
/// function pointer with the void pointer to the target class as one of
/// the arguments.
/// 3. Create the following template function:
/// \code{.cpp}
/// template<typename TargetClass>
/// VTKM_EXEC void Bind(const TargetClass *deviceTarget);
/// \endcode
/// This function should set the void pointer from 1 to \c deviceTarget,
/// and assign the function pointers from 2b to functions that cast their
/// first argument to <tt>const TargetClass*</tt> and call the corresponding
/// member function on it.
///
/// Both \c VirtualObject and target class objects should be bitwise copyable.
///
/// \sa vtkm::exec::ImplicitFunction, vtkm::cont::ImplicitFunction
///
template <typename VirtualObject>
class VirtualObjectCache
{
public:
VirtualObjectCache()
: Target(nullptr)
, CurrentDevice(VTKM_DEVICE_ADAPTER_UNDEFINED)
, DeviceState(nullptr)
, RefreshFlag(false)
{
}
~VirtualObjectCache() { this->Reset(); }
VirtualObjectCache(const VirtualObjectCache& other)
: Target(other.Target)
, Transfers(other.Transfers)
, CurrentDevice(VTKM_DEVICE_ADAPTER_UNDEFINED)
, DeviceState(nullptr)
, RefreshFlag(false)
{
}
VirtualObjectCache& operator=(const VirtualObjectCache& other)
{
if (this != &other)
{
this->Target = other.Target;
this->Transfers = other.Transfers;
this->CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED;
this->DeviceState = nullptr;
this->RefreshFlag = false;
this->Object = VirtualObject();
}
return *this;
}
VirtualObjectCache(VirtualObjectCache&& other)
: Target(other.Target)
, Transfers(other.Transfers)
, CurrentDevice(other.CurrentDevice)
, DeviceState(other.DeviceState)
, RefreshFlag(other.RefreshFlag)
, Object(other.Object)
{
other.CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED;
other.DeviceState = nullptr;
}
VirtualObjectCache& operator=(VirtualObjectCache&& other)
{
if (this != &other)
{
this->Target = other.Target;
this->Transfers = std::move(other.Transfers);
this->CurrentDevice = other.CurrentDevice;
this->DeviceState = other.DeviceState;
this->RefreshFlag = other.RefreshFlag;
this->Object = std::move(other.Object);
other.CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED;
other.DeviceState = nullptr;
}
return *this;
}
/// Reset to the default constructed state
void Reset()
{
this->ReleaseResources();
this->Target = nullptr;
this->Transfers.fill(TransferInterface());
}
void ReleaseResources()
{
if (this->CurrentDevice > 0)
{
this->GetCurrentTransfer().Cleanup(this->DeviceState);
this->CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED;
this->DeviceState = nullptr;
this->RefreshFlag = false;
this->Object = VirtualObject();
}
}
/// Get if in a valid state (a target is bound)
bool GetValid() const { return this->Target != nullptr; }
// Set/Get if the cached virtual object should be refreshed to the current
// state of the target
void SetRefreshFlag(bool value) { this->RefreshFlag = value; }
bool GetRefreshFlag() const { return this->RefreshFlag; }
/// Bind to \c target. The lifetime of target is expected to be managed
/// externally, and should be valid for as long as it is bound.
/// Also accepts a list-tag of device adapters where the virtual
/// object may be used (default = \c VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG).
///
template <typename TargetClass, typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
void Bind(const TargetClass* target, DeviceAdapterList devices = DeviceAdapterList())
{
this->Reset();
this->Target = target;
vtkm::ListForEach(CreateTransferInterface<TargetClass>(this->Transfers.data()), devices);
}
/// Get a \c VirtualObject for \c DeviceAdapter.
/// VirtualObjectCache and VirtualObject are analogous to ArrayHandle and Portal
/// The returned VirtualObject will be invalidated if:
/// 1. A new VirtualObject is requested for a different DeviceAdapter
/// 2. VirtualObjectCache is destroyed
/// 3. Reset or ReleaseResources is called
///
template <typename DeviceAdapter>
VirtualObject GetVirtualObject(DeviceAdapter)
{
using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
if (!this->GetValid())
{
throw vtkm::cont::ErrorBadValue("No target object bound");
}
vtkm::cont::DeviceAdapterId deviceId = DeviceInfo::GetId();
if (deviceId < 0 || deviceId >= VTKM_MAX_DEVICE_ADAPTER_ID)
{
std::string msg = "Device '" + DeviceInfo::GetName() + "' has invalid ID of " +
std::to_string(deviceId) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " +
std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")";
throw vtkm::cont::ErrorBadType(msg);
}
if (this->CurrentDevice != deviceId)
{
this->ReleaseResources();
std::size_t idx = static_cast<std::size_t>(deviceId);
TransferInterface& transfer = this->Transfers[idx];
if (!TransferInterfaceValid(transfer))
{
std::string msg = DeviceInfo::GetName() + " was not in the list specified in Bind";
throw vtkm::cont::ErrorBadType(msg);
}
this->CurrentDevice = deviceId;
this->DeviceState = transfer.Create(this->Object, this->Target);
}
else if (this->RefreshFlag)
{
this->GetCurrentTransfer().Update(this->DeviceState, this->Target);
}
this->RefreshFlag = false;
return this->Object;
}
private:
struct TransferInterface
{
using CreateSig = void*(VirtualObject&, const void*);
using UpdateSig = void(void*, const void*);
using CleanupSig = void(void*);
CreateSig* Create = nullptr;
UpdateSig* Update = nullptr;
CleanupSig* Cleanup = nullptr;
};
static bool TransferInterfaceValid(const TransferInterface& t) { return t.Create != nullptr; }
TransferInterface& GetCurrentTransfer()
{
return this->Transfers[static_cast<std::size_t>(this->CurrentDevice)];
}
template <typename TargetClass>
class CreateTransferInterface
{
private:
template <typename DeviceAdapter>
using EnableIfValid = std::enable_if<vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
template <typename DeviceAdapter>
using EnableIfInvalid = std::enable_if<!vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
public:
CreateTransferInterface(TransferInterface* transfers)
: Transfers(transfers)
{
}
// Use SFINAE to create entries for valid device adapters only
template <typename DeviceAdapter>
typename EnableIfValid<DeviceAdapter>::type operator()(DeviceAdapter) const
{
using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
if (DeviceInfo::GetId() >= 0 && DeviceInfo::GetId() < VTKM_MAX_DEVICE_ADAPTER_ID)
{
using TransferImpl =
internal::VirtualObjectTransfer<VirtualObject, TargetClass, DeviceAdapter>;
std::size_t id = static_cast<std::size_t>(DeviceInfo::GetId());
TransferInterface& transfer = this->Transfers[id];
transfer.Create = &TransferImpl::Create;
transfer.Update = &TransferImpl::Update;
transfer.Cleanup = &TransferImpl::Cleanup;
}
else
{
std::string msg = "Device '" + DeviceInfo::GetName() + "' has invalid ID of " +
std::to_string(DeviceInfo::GetId()) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " +
std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")";
throw vtkm::cont::ErrorBadType(msg);
}
}
template <typename DeviceAdapter>
typename EnableIfInvalid<DeviceAdapter>::type operator()(DeviceAdapter) const
{
}
private:
TransferInterface* Transfers;
};
const void* Target;
std::array<TransferInterface, VTKM_MAX_DEVICE_ADAPTER_ID> Transfers;
vtkm::cont::DeviceAdapterId CurrentDevice;
void* DeviceState;
bool RefreshFlag;
VirtualObject Object;
};
}
} // vtkm::cont
#undef VTKM_MAX_DEVICE_ADAPTER_ID
#endif // vtk_m_cont_VirtualObjectCache_h

277
vtkm/cont/VirtualObjectHandle.h Normal file
View File

@ -0,0 +1,277 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_VirtualObjectHandle_h
#define vtk_m_cont_VirtualObjectHandle_h
#include <vtkm/cont/DeviceAdapterListTag.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/internal/DeviceAdapterTag.h>
#include <vtkm/cont/internal/VirtualObjectTransfer.h>
#include <array>
#include <type_traits>
#define VTKM_MAX_DEVICE_ADAPTER_ID 8
namespace vtkm
{
namespace cont
{
/// \brief Implements VTK-m's execution side <em> Virtual Methods </em> functionality.
///
/// The template parameter \c VirtualBaseType is the base class that acts as the
/// interface. This base clase must inherit from \c vtkm::VirtualObjectBase. See the
/// documentation of that class to see the other requirements.
///
/// A derived object can be bound to the handle either during construction or using the \c Reset
/// function on previously constructed handles. These function accept a control side pointer to
/// the derived class object, a boolean flag stating if the handle should acquire ownership of
/// the object (i.e. manage the lifetime), and a type-list of device adapter tags where the object
/// is expected to be used.
///
/// To get an execution side pointer call the \c PrepareForExecution function. The device adapter
/// passed to this function should be one of the device in the list passed during the set up of
/// the handle.
///
///
/// \sa vtkm::VirtualObjectBase
///
template <typename VirtualBaseType>
class VTKM_ALWAYS_EXPORT VirtualObjectHandle
{
VTKM_STATIC_ASSERT_MSG((std::is_base_of<vtkm::VirtualObjectBase, VirtualBaseType>::value),
"All virtual objects must be subclass of vtkm::VirtualObjectBase.");
public:
VTKM_CONT VirtualObjectHandle()
: Internals(new InternalStruct)
{
}
template <typename VirtualDerivedType,
typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
VTKM_CONT explicit VirtualObjectHandle(VirtualDerivedType* derived,
bool acquireOwnership = true,
DeviceAdapterList devices = DeviceAdapterList())
: Internals(new InternalStruct)
{
this->Reset(derived, acquireOwnership, devices);
}
/// Get if in a valid state (a target is bound)
VTKM_CONT bool GetValid() const { return this->Internals->VirtualObject != nullptr; }
/// Get if this handle owns the control side target object
VTKM_CONT bool OwnsObject() const { return this->Internals->Owner; }
/// Get the control side pointer to the virtual object
VTKM_CONT VirtualBaseType* Get() const { return this->Internals->VirtualObject; }
/// Reset the underlying derived type object
template <typename VirtualDerivedType,
typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
VTKM_CONT void Reset(VirtualDerivedType* derived,
bool acquireOwnership = true,
DeviceAdapterList devices = DeviceAdapterList())
{
this->Reset();
if (derived)
{
VTKM_STATIC_ASSERT_MSG((std::is_base_of<VirtualBaseType, VirtualDerivedType>::value),
"Tried to bind a type that is not a subclass of the base class.");
this->Internals->VirtualObject = derived;
this->Internals->Owner = acquireOwnership;
vtkm::cont::ForEachValidDevice(
devices,
CreateTransferInterface<VirtualDerivedType>(this->Internals->Transfers.data(), derived));
}
}
void Reset() { this->Internals->Reset(); }
/// Release all the execution side resources
VTKM_CONT void ReleaseExecutionResources() { this->Internals->ReleaseExecutionResources(); }
/// Get a valid \c VirtualBaseType* with the current control side state for \c DeviceAdapter.
/// VirtualObjectHandle and the returned pointer are analogous to ArrayHandle and Portal
/// The returned pointer will be invalidated if:
/// 1. A new pointer is requested for a different DeviceAdapter
/// 2. VirtualObjectHandle is destroyed
/// 3. Reset or ReleaseResources is called
///
template <typename DeviceAdapter>
VTKM_CONT const VirtualBaseType* PrepareForExecution(DeviceAdapter) const
{
using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
if (!this->GetValid())
{
throw vtkm::cont::ErrorBadValue("No target object bound");
}
vtkm::cont::DeviceAdapterId deviceId = DeviceInfo::GetId();
if (deviceId < 0 || deviceId >= VTKM_MAX_DEVICE_ADAPTER_ID)
{
std::string msg = "Device '" + DeviceInfo::GetName() + "' has invalid ID of " +
std::to_string(deviceId) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " +
std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")";
throw vtkm::cont::ErrorBadType(msg);
}
if (!this->Internals->Current || this->Internals->Current->GetDeviceId() != deviceId)
{
if (!this->Internals->Transfers[static_cast<std::size_t>(deviceId)])
{
std::string msg = DeviceInfo::GetName() + " was not in the list specified in Bind";
throw vtkm::cont::ErrorBadType(msg);
}
if (this->Internals->Current)
{
this->Internals->Current->ReleaseResources();
}
this->Internals->Current =
this->Internals->Transfers[static_cast<std::size_t>(deviceId)].get();
}
return this->Internals->Current->PrepareForExecution();
}
private:
class TransferInterface
{
public:
VTKM_CONT virtual ~TransferInterface() = default;
VTKM_CONT virtual vtkm::cont::DeviceAdapterId GetDeviceId() const = 0;
VTKM_CONT virtual const VirtualBaseType* PrepareForExecution() = 0;
VTKM_CONT virtual void ReleaseResources() = 0;
};
template <typename VirtualDerivedType, typename DeviceAdapter>
class TransferInterfaceImpl : public TransferInterface
{
public:
VTKM_CONT TransferInterfaceImpl(const VirtualDerivedType* virtualObject)
: LastModifiedCount(-1)
, VirtualObject(virtualObject)
, Transfer(virtualObject)
{
}
VTKM_CONT vtkm::cont::DeviceAdapterId GetDeviceId() const override
{
return vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetId();
}
VTKM_CONT const VirtualBaseType* PrepareForExecution() override
{
vtkm::Id modifiedCount = this->VirtualObject->GetModifiedCount();
bool updateData = (this->LastModifiedCount != modifiedCount);
const VirtualBaseType* executionObject = this->Transfer.PrepareForExecution(updateData);
this->LastModifiedCount = modifiedCount;
return executionObject;
}
VTKM_CONT void ReleaseResources() override { this->Transfer.ReleaseResources(); }
private:
vtkm::Id LastModifiedCount;
const VirtualDerivedType* VirtualObject;
vtkm::cont::internal::VirtualObjectTransfer<VirtualDerivedType, DeviceAdapter> Transfer;
};
template <typename VirtualDerivedType>
class CreateTransferInterface
{
public:
CreateTransferInterface(std::unique_ptr<TransferInterface>* transfers,
const VirtualDerivedType* virtualObject)
: Transfers(transfers)
, VirtualObject(virtualObject)
{
}
template <typename DeviceAdapter>
void operator()(DeviceAdapter) const
{
using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
if (DeviceInfo::GetId() < 0 || DeviceInfo::GetId() >= VTKM_MAX_DEVICE_ADAPTER_ID)
{
std::string msg = "Device '" + DeviceInfo::GetName() + "' has invalid ID of " +
std::to_string(DeviceInfo::GetId()) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " +
std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")";
throw vtkm::cont::ErrorBadType(msg);
}
using TransferImpl = TransferInterfaceImpl<VirtualDerivedType, DeviceAdapter>;
this->Transfers[DeviceInfo::GetId()].reset(new TransferImpl(this->VirtualObject));
}
private:
std::unique_ptr<TransferInterface>* Transfers;
const VirtualDerivedType* VirtualObject;
};
struct InternalStruct
{
VirtualBaseType* VirtualObject = nullptr;
bool Owner = false;
std::array<std::unique_ptr<TransferInterface>, VTKM_MAX_DEVICE_ADAPTER_ID> Transfers;
TransferInterface* Current = nullptr;
VTKM_CONT void ReleaseExecutionResources()
{
if (this->Current)
{
this->Current->ReleaseResources();
this->Current = nullptr;
}
}
VTKM_CONT void Reset()
{
this->ReleaseExecutionResources();
for (auto& transfer : this->Transfers)
{
transfer.reset(nullptr);
}
if (this->Owner)
{
delete this->VirtualObject;
}
this->VirtualObject = nullptr;
this->Owner = false;
}
VTKM_CONT ~InternalStruct() { this->Reset(); }
};
std::shared_ptr<InternalStruct> Internals;
};
}
} // vtkm::cont
#undef VTKM_MAX_DEVICE_ADAPTER_ID
#endif // vtk_m_cont_VirtualObjectHandle_h

25
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu
View File

@ -133,6 +133,7 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
// execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
CudaAllocator::PrepareForInput(executionPtr, numBytes);
return;
}
@ -160,6 +161,7 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const voi
}
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
@ -170,6 +172,29 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const voi
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInPlace(executionPtr, static_cast<size_t>(numBytes));
}
} // end namespace internal
VTKM_INSTANTIATE_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagCuda)

10
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
View File

@ -159,6 +159,16 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
VTKM_CONT void CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
} // namespace internal

40
vtkm/cont/cuda/internal/CudaAllocator.cu
View File

@ -52,6 +52,7 @@ bool CudaAllocator::UsingManagedMemory()
bool CudaAllocator::IsDevicePointer(const void* ptr)
{
CudaAllocator::Initialize();
if (!ptr)
{
return false;
@ -72,7 +73,7 @@ bool CudaAllocator::IsDevicePointer(const void* ptr)
bool CudaAllocator::IsManagedPointer(const void* ptr)
{
if (!ptr)
if (!ptr || !ManagedMemorySupported)
{
return false;
}
@ -109,23 +110,17 @@ void* CudaAllocator::Allocate(std::size_t numBytes)
void CudaAllocator::Free(void* ptr)
{
CudaAllocator::Initialize();
VTKM_CUDA_CALL(cudaFree(ptr));
}
void CudaAllocator::PrepareForControl(const void* ptr, std::size_t numBytes)
{
CudaAllocator::Initialize();
if (ManagedMemorySupported)
if (IsManagedPointer(ptr))
{
#if CUDART_VERSION >= 8000
// TODO these hints need to be benchmarked and adjusted once we start
// sharing the pointers between cont/exec
VTKM_CUDA_CALL(
cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, cudaCpuDeviceId));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseUnsetReadMostly, cudaCpuDeviceId));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId));
VTKM_CUDA_CALL(cudaMemPrefetchAsync(ptr, numBytes, cudaCpuDeviceId, cudaStreamPerThread));
#endif // CUDA >= 8.0
}
@ -133,15 +128,14 @@ void CudaAllocator::PrepareForControl(const void* ptr, std::size_t numBytes)
void CudaAllocator::PrepareForInput(const void* ptr, std::size_t numBytes)
{
CudaAllocator::Initialize();
if (ManagedMemorySupported)
if (IsManagedPointer(ptr))
{
#if CUDART_VERSION >= 8000
int dev;
VTKM_CUDA_CALL(cudaGetDevice(&dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetReadMostly, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetReadMostly, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetAccessedBy, dev));
VTKM_CUDA_CALL(cudaMemPrefetchAsync(ptr, numBytes, dev, cudaStreamPerThread));
#endif // CUDA >= 8.0
}
@ -149,15 +143,14 @@ void CudaAllocator::PrepareForInput(const void* ptr, std::size_t numBytes)
void CudaAllocator::PrepareForOutput(const void* ptr, std::size_t numBytes)
{
CudaAllocator::Initialize();
if (ManagedMemorySupported)
if (IsManagedPointer(ptr))
{
#if CUDART_VERSION >= 8000
int dev;
VTKM_CUDA_CALL(cudaGetDevice(&dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseUnsetReadMostly, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseUnsetReadMostly, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetAccessedBy, dev));
VTKM_CUDA_CALL(cudaMemPrefetchAsync(ptr, numBytes, dev, cudaStreamPerThread));
#endif // CUDA >= 8.0
}
@ -165,15 +158,14 @@ void CudaAllocator::PrepareForOutput(const void* ptr, std::size_t numBytes)
void CudaAllocator::PrepareForInPlace(const void* ptr, std::size_t numBytes)
{
CudaAllocator::Initialize();
if (ManagedMemorySupported)
if (IsManagedPointer(ptr))
{
#if CUDART_VERSION >= 8000
int dev;
VTKM_CUDA_CALL(cudaGetDevice(&dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseUnsetReadMostly, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetPreferredLocation, dev));
// VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseUnsetReadMostly, dev));
VTKM_CUDA_CALL(cudaMemAdvise(ptr, numBytes, cudaMemAdviseSetAccessedBy, dev));
VTKM_CUDA_CALL(cudaMemPrefetchAsync(ptr, numBytes, dev, cudaStreamPerThread));
#endif // CUDA >= 8.0
}

112
vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h
View File

@ -20,6 +20,7 @@
#ifndef vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h
#define vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/cuda/ErrorCuda.h>
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
#include <vtkm/cont/internal/VirtualObjectTransfer.h>
@ -34,45 +35,108 @@ namespace internal
namespace detail
{
template <typename VirtualObject, typename TargetClass>
__global__ void CreateKernel(VirtualObject* object, const TargetClass* target)
template <typename VirtualDerivedType>
__global__ void ConstructVirtualObjectKernel(VirtualDerivedType* deviceObject,
const VirtualDerivedType* targetObject)
{
object->Bind(target);
// Use the "placement new" syntax to construct an object in pre-allocated memory
new (deviceObject) VirtualDerivedType(*targetObject);
}
template <typename VirtualDerivedType>
__global__ void UpdateVirtualObjectKernel(VirtualDerivedType* deviceObject,
const VirtualDerivedType* targetObject)
{
*deviceObject = *targetObject;
}
template <typename VirtualDerivedType>
__global__ void DeleteVirtualObjectKernel(VirtualDerivedType* deviceObject)
{
deviceObject->~VirtualDerivedType();
}
} // detail
template <typename VirtualObject, typename TargetClass>
struct VirtualObjectTransfer<VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagCuda>
template <typename VirtualDerivedType>
struct VirtualObjectTransfer<VirtualDerivedType, vtkm::cont::DeviceAdapterTagCuda>
{
static void* Create(VirtualObject& object, const void* target)
VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject)
: ControlObject(virtualObject)
, ExecutionObject(nullptr)
{
TargetClass* cutarget;
VTKM_CUDA_CALL(cudaMalloc(&cutarget, sizeof(TargetClass)));
VTKM_CUDA_CALL(cudaMemcpyAsync(
cutarget, target, sizeof(TargetClass), cudaMemcpyHostToDevice, cudaStreamPerThread));
VirtualObject* cuobject;
VTKM_CUDA_CALL(cudaMalloc(&cuobject, sizeof(VirtualObject)));
detail::CreateKernel<<<1, 1, 0, cudaStreamPerThread>>>(cuobject, cutarget);
VTKM_CUDA_CHECK_ASYNCHRONOUS_ERROR();
VTKM_CUDA_CALL(cudaMemcpyAsync(
&object, cuobject, sizeof(VirtualObject), cudaMemcpyDeviceToHost, cudaStreamPerThread));
VTKM_CUDA_CALL(cudaFree(cuobject));
return cutarget;
}
static void Update(void* deviceState, const void* target)
VTKM_CONT ~VirtualObjectTransfer() { this->ReleaseResources(); }
VirtualObjectTransfer(const VirtualObjectTransfer&) = delete;
void operator=(const VirtualObjectTransfer&) = delete;
VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool updateData)
{
VTKM_CUDA_CALL(cudaMemcpyAsync(
deviceState, target, sizeof(TargetClass), cudaMemcpyHostToDevice, cudaStreamPerThread));
if (this->ExecutionObject == nullptr)
{
// deviceTarget will hold a byte copy of the host object on the device. The virtual table
// will be wrong.
VirtualDerivedType* deviceTarget;
VTKM_CUDA_CALL(cudaMalloc(&deviceTarget, sizeof(VirtualDerivedType)));
VTKM_CUDA_CALL(cudaMemcpy(
deviceTarget, this->ControlObject, sizeof(VirtualDerivedType), cudaMemcpyHostToDevice));
// Allocate memory for the object that will eventually be a correct copy on the device.
VTKM_CUDA_CALL(cudaMalloc(&this->ExecutionObject, sizeof(VirtualDerivedType)));
// Initialize the device object
detail::ConstructVirtualObjectKernel<<<1, 1>>>(this->ExecutionObject, deviceTarget);
VTKM_CUDA_CHECK_ASYNCHRONOUS_ERROR();
// Clean up intermediate copy
VTKM_CUDA_CALL(cudaFree(deviceTarget));
}
else if (updateData)
{
// deviceTarget will hold a byte copy of the host object on the device. The virtual table
// will be wrong.
VirtualDerivedType* deviceTarget;
VTKM_CUDA_CALL(cudaMalloc(&deviceTarget, sizeof(VirtualDerivedType)));
VTKM_CUDA_CALL(cudaMemcpy(
deviceTarget, this->ControlObject, sizeof(VirtualDerivedType), cudaMemcpyHostToDevice));
// Initialize the device object
detail::UpdateVirtualObjectKernel<<<1, 1>>>(this->ExecutionObject, deviceTarget);
VTKM_CUDA_CHECK_ASYNCHRONOUS_ERROR();
// Clean up intermediate copy
VTKM_CUDA_CALL(cudaFree(deviceTarget));
}
else
{
// Nothing to do. The device object is already up to date.
}
return this->ExecutionObject;
}
static void Cleanup(void* deviceState) { VTKM_CUDA_CALL(cudaFree(deviceState)); }
VTKM_CONT void ReleaseResources()
{
if (this->ExecutionObject != nullptr)
{
detail::DeleteVirtualObjectKernel<<<1, 1>>>(this->ExecutionObject);
VTKM_CUDA_CALL(cudaFree(this->ExecutionObject));
this->ExecutionObject = nullptr;
}
}
private:
const VirtualDerivedType* ControlObject;
VirtualDerivedType* ExecutionObject;
};
}
}
} // vtkm::cont::internal
#define VTKM_EXPLICITLY_INSTANTIATE_TRANSFER(DerivedType) \
template class vtkm::cont::internal::VirtualObjectTransfer<DerivedType, \
vtkm::cont::DeviceAdapterTagCuda>
#endif // vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h

2
vtkm/cont/cuda/testing/CMakeLists.txt
View File

@ -29,6 +29,6 @@ set(unit_tests
UnitTestCudaMath.cu
UnitTestCudaShareUserProvidedManagedMemory.cu
UnitTestCudaPointLocatorUniformGrid.cxx
UnitTestCudaVirtualObjectCache.cu
UnitTestCudaVirtualObjectHandle.cu
)
vtkm_unit_tests(SOURCES ${unit_tests})

10
vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu → vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectHandle.cu
View File

@ -17,22 +17,22 @@
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/testing/TestingVirtualObjectCache.h>
#include <vtkm/cont/testing/TestingVirtualObjectHandle.h>
namespace
{
void TestVirtualObjectCache()
void TestVirtualObjectHandle()
{
using DeviceAdapterList =
vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial, vtkm::cont::DeviceAdapterTagCuda>;
vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
vtkm::cont::testing::TestingVirtualObjectHandle<DeviceAdapterList>::Run();
}
} // anonymous namespace
int UnitTestCudaVirtualObjectCache(int, char* [])
int UnitTestCudaVirtualObjectHandle(int, char* [])
{
return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
return vtkm::cont::testing::Testing::Run(TestVirtualObjectHandle);
}

11
vtkm/cont/internal/ArrayHandleBasicImpl.h
View File

@ -112,6 +112,17 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicBase
void* controlPtr,
vtkm::UInt64 numBytes) const = 0;
VTKM_CONT virtual void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const = 0;
VTKM_CONT virtual void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const = 0;
VTKM_CONT virtual void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const = 0;
protected:
StorageBasicBase& ControlStorage;
};

27
vtkm/cont/internal/ArrayHandleBasicImpl.hxx
View File

@ -267,6 +267,9 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInput(DeviceAdapterTag device) const
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
@ -284,9 +287,6 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInput(DeviceAdapterTag device) const
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
@ -294,6 +294,8 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInput(DeviceAdapterTag device) const
this->Internals->ExecutionArrayValid = true;
}
this->Internals->ExecutionInterface->UsingForRead(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
return PortalFactory<DeviceAdapterTag>::CreatePortalConst(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
@ -320,11 +322,17 @@ ArrayHandle<T, StorageTagBasic>::PrepareForOutput(vtkm::Id numVals, DeviceAdapte
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
this->Internals->ExecutionInterface->Allocate(
execArray, static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numVals));
const vtkm::UInt64 numBytes =
static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numVals);
this->Internals->ExecutionInterface->Allocate(execArray, numBytes);
this->Internals->ExecutionInterface->UsingForWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
return PortalFactory<DeviceAdapterTag>::CreatePortal(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
}
@ -339,6 +347,9 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInPlace(DeviceAdapterTag device)
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
// Initialize an empty array if needed:
@ -355,9 +366,6 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInPlace(DeviceAdapterTag device)
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
@ -366,6 +374,9 @@ ArrayHandle<T, StorageTagBasic>::PrepareForInPlace(DeviceAdapterTag device)
this->Internals->ExecutionArrayValid = true;
}
priv->ExecutionInterface->UsingForReadWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
// Invalidate the control array, since we expect the values to be modified:
this->Internals->ControlArrayValid = false;

16
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx
View File

@ -77,6 +77,22 @@ void ExecutionArrayInterfaceBasicShareWithControl::CopyToControl(const void* src
std::copy(srcBegin, srcEnd, dstBegin);
}
}
void ExecutionArrayInterfaceBasicShareWithControl::UsingForRead(const void*,
const void*,
vtkm::UInt64) const
{
}
void ExecutionArrayInterfaceBasicShareWithControl::UsingForWrite(const void*,
const void*,
vtkm::UInt64) const
{
}
void ExecutionArrayInterfaceBasicShareWithControl::UsingForReadWrite(const void*,
const void*,
vtkm::UInt64) const
{
}
}
}
} // end namespace vtkm::cont::internal

11
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
View File

@ -141,8 +141,19 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicShareWithControl
VTKM_CONT void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* src, void* dst, vtkm::UInt64 bytes) const final;
VTKM_CONT void CopyToControl(const void* src, void* dst, vtkm::UInt64 bytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
}
}

60
vtkm/cont/internal/DynamicTransform.h
View File

@ -22,6 +22,8 @@
#include "vtkm/internal/IndexTag.h"
#include <utility>
namespace vtkm
{
namespace cont
@ -29,6 +31,10 @@ namespace cont
template <typename T, typename S>
class ArrayHandle;
class CoordinateSystem;
class Field;
template <vtkm::IdComponent>
class CellSetStructured;
template <typename T>
@ -41,57 +47,75 @@ class CellSetPermutation;
/// A Generic interface to CastAndCall. The default implementation simply calls
/// DynamicObject's CastAndCall, but specializations of this function exist for
/// other classes (e.g. Field, CoordinateSystem, ArrayHandle).
template <typename DynamicObject, typename Functor>
void CastAndCall(const DynamicObject& dynamicObject, const Functor& f)
template <typename DynamicObject, typename Functor, typename... Args>
void CastAndCall(const DynamicObject& dynamicObject, const Functor& f, Args&&... args)
{
dynamicObject.CastAndCall(f);
dynamicObject.CastAndCall(f, std::forward<Args>(args)...);
}
/// A specialization of CastAndCall for basic CoordinateSystem to make
/// it be treated just like any other dynamic object
// actually implemented in vtkm/cont/CoordinateSystem
template <typename Functor, typename... Args>
void CastAndCall(const CoordinateSystem& coords, const Functor& f, Args&&... args);
/// A specialization of CastAndCall for basic Field to make
/// it be treated just like any other dynamic object
// actually implemented in vtkm/cont/Field
template <typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::Field& field, const Functor& f, Args&&... args);
/// A specialization of CastAndCall for basic ArrayHandle types,
/// Since the type is already known no deduction is needed.
/// This specialization is used to simplify numerous worklet algorithms
template <typename T, typename U, typename Functor>
void CastAndCall(const vtkm::cont::ArrayHandle<T, U>& handle, const Functor& f)
template <typename T, typename U, typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::ArrayHandle<T, U>& handle, const Functor& f, Args&&... args)
{
f(handle);
f(handle, std::forward<Args>(args)...);
}
/// A specialization of CastAndCall for basic CellSetStructured types,
/// Since the type is already known no deduction is needed.
/// This specialization is used to simplify numerous worklet algorithms
template <vtkm::IdComponent Dim, typename Functor>
void CastAndCall(const vtkm::cont::CellSetStructured<Dim>& cellset, const Functor& f)
template <vtkm::IdComponent Dim, typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::CellSetStructured<Dim>& cellset,
const Functor& f,
Args&&... args)
{
f(cellset);
f(cellset, std::forward<Args>(args)...);
}
/// A specialization of CastAndCall for basic CellSetSingleType types,
/// Since the type is already known no deduction is needed.
/// This specialization is used to simplify numerous worklet algorithms
template <typename ConnectivityStorageTag, typename Functor>
template <typename ConnectivityStorageTag, typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::CellSetSingleType<ConnectivityStorageTag>& cellset,
const Functor& f)
const Functor& f,
Args&&... args)
{
f(cellset);
f(cellset, std::forward<Args>(args)...);
}
/// A specialization of CastAndCall for basic CellSetExplicit types,
/// Since the type is already known no deduction is needed.
/// This specialization is used to simplify numerous worklet algorithms
template <typename T, typename S, typename U, typename V, typename Functor>
void CastAndCall(const vtkm::cont::CellSetExplicit<T, S, U, V>& cellset, const Functor& f)
template <typename T, typename S, typename U, typename V, typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::CellSetExplicit<T, S, U, V>& cellset,
const Functor& f,
Args&&... args)
{
f(cellset);
f(cellset, std::forward<Args>(args)...);
}
/// A specialization of CastAndCall for basic CellSetPermutation types,
/// Since the type is already known no deduction is needed.
/// This specialization is used to simplify numerous worklet algorithms
template <typename PermutationType, typename CellSetType, typename Functor>
template <typename PermutationType, typename CellSetType, typename Functor, typename... Args>
void CastAndCall(const vtkm::cont::CellSetPermutation<PermutationType, CellSetType>& cellset,
const Functor& f)
const Functor& f,
Args&&... args)
{
f(cellset);
f(cellset, std::forward<Args>(args)...);
}
namespace internal

30
vtkm/cont/internal/VirtualObjectTransfer.h
View File

@ -20,6 +20,8 @@
#ifndef vtk_m_cont_internal_VirtualObjectTransfer_h
#define vtk_m_cont_internal_VirtualObjectTransfer_h
#include <vtkm/VirtualObjectBase.h>
namespace vtkm
{
namespace cont
@ -27,23 +29,31 @@ namespace cont
namespace internal
{
template <typename VirtualObject, typename TargetClass, typename DeviceAdapter>
template <typename VirtualDerivedType, typename DeviceAdapter>
struct VirtualObjectTransfer
#ifdef VTKM_DOXYGEN_ONLY
{
/// Takes a void* to host copy of the target object, transfers it to the
/// device, binds it to the VirtualObject, and returns a void* to an internal
/// state structure.
/// A VirtualObjectTransfer is constructed with a pointer to the derived type that (eventually)
/// gets transferred to the execution environment of the given DeviceAdapter.
///
static void* Create(VirtualObject& object, const void* hostTarget);
VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject);
/// Performs cleanup of the device state used to track the VirtualObject on
/// the device.
/// \brief Transfers the virtual object to the execution environment.
///
static void Cleanup(void* deviceState);
/// This method transfers the virtual object to the execution environment and returns a pointer
/// to the object that can be used in the execution environment (but not necessarily the control
/// environment). If the \c updateData flag is true, then the data is always copied to the
/// execution environment (such as if the data were updated since the last call to \c
/// PrepareForExecution). If the \c updateData flag is false and the object was already
/// transferred previously, the previously created object is returned.
///
VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool updateData);
/// Update the device state with the state of target
static void Update(void* deviceState, const void* target);
/// \brief Frees up any resources in the execution environment.
///
/// Any previously returned virtual object from \c PrepareForExecution becomes invalid.
///
VTKM_CONT void ReleaseResources();
}
#endif
;

23
vtkm/cont/internal/VirtualObjectTransferShareWithControl.h
View File

@ -20,6 +20,11 @@
#ifndef vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h
#define vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h
#include <vtkm/StaticAssert.h>
#include <vtkm/VirtualObjectBase.h>
#include <typeinfo>
namespace vtkm
{
namespace cont
@ -27,19 +32,23 @@ namespace cont
namespace internal
{
template <typename VirtualObject, typename TargetClass>
template <typename VirtualDerivedType>
struct VirtualObjectTransferShareWithControl
{
static void* Create(VirtualObject& object, const void* target)
VTKM_CONT VirtualObjectTransferShareWithControl(const VirtualDerivedType* virtualObject)
: VirtualObject(virtualObject)
{
static int dummyState = 0;
object.Bind(static_cast<const TargetClass*>(target));
return &dummyState;
}
static void Update(void*, const void*) {}
VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool vtkmNotUsed(updateData))
{
return this->VirtualObject;
}
static void Cleanup(void*) {}
VTKM_CONT void ReleaseResources() {}
private:
const VirtualDerivedType* VirtualObject;
};
}
}

5
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h
View File

@ -83,6 +83,11 @@ public:
auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial());
auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTagSerial());
if (inSize <= 0)
{
return;
}
using InputType = decltype(inputPortal.Get(0));
using OutputType = decltype(outputPortal.Get(0));

10
vtkm/cont/serial/internal/VirtualObjectTransferSerial.h
View File

@ -31,10 +31,14 @@ namespace cont
namespace internal
{
template <typename VirtualObject, typename TargetClass>
struct VirtualObjectTransfer<VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagSerial>
: public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
template <typename VirtualDerivedType>
struct VirtualObjectTransfer<VirtualDerivedType, vtkm::cont::DeviceAdapterTagSerial>
: VirtualObjectTransferShareWithControl<VirtualDerivedType>
{
VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject)
: VirtualObjectTransferShareWithControl<VirtualDerivedType>(virtualObject)
{
}
};
}
}

2
vtkm/cont/serial/testing/CMakeLists.txt
View File

@ -28,6 +28,6 @@ set(unit_tests
UnitTestSerialDeviceAdapter.cxx
UnitTestSerialImplicitFunction.cxx
UnitTestSerialPointLocatorUniformGrid.cxx
UnitTestSerialVirtualObjectCache.cxx
UnitTestSerialVirtualObjectHandle.cxx
)
vtkm_unit_tests(TBB SOURCES ${unit_tests})

10
vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx → vtkm/cont/serial/testing/UnitTestSerialVirtualObjectHandle.cxx
View File

@ -17,21 +17,21 @@
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/testing/TestingVirtualObjectCache.h>
#include <vtkm/cont/testing/TestingVirtualObjectHandle.h>
namespace
{
void TestVirtualObjectCache()
void TestVirtualObjectHandle()
{
using DeviceAdapterList = vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial>;
vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
vtkm::cont::testing::TestingVirtualObjectHandle<DeviceAdapterList>::Run();
}
} // anonymous namespace
int UnitTestSerialVirtualObjectCache(int, char* [])
int UnitTestSerialVirtualObjectHandle(int, char* [])
{
return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
return vtkm::cont::testing::Testing::Run(TestVirtualObjectHandle);
}

10
vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h
View File

@ -31,10 +31,14 @@ namespace cont
namespace internal
{
template <typename VirtualObject, typename TargetClass>
struct VirtualObjectTransfer<VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagTBB>
: public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
template <typename VirtualDerivedType>
struct VirtualObjectTransfer<VirtualDerivedType, vtkm::cont::DeviceAdapterTagTBB>
: VirtualObjectTransferShareWithControl<VirtualDerivedType>
{
VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject)
: VirtualObjectTransferShareWithControl<VirtualDerivedType>(virtualObject)
{
}
};
}
}

2
vtkm/cont/tbb/testing/CMakeLists.txt
View File

@ -28,6 +28,6 @@ set(unit_tests
UnitTestTBBDeviceAdapter.cxx
UnitTestTBBImplicitFunction.cxx
UnitTestTBBPointLocatorUniformGrid.cxx
UnitTestTBBVirtualObjectCache.cxx
UnitTestTBBVirtualObjectHandle.cxx
)
vtkm_unit_tests(TBB SOURCES ${unit_tests})

10
vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx → vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectHandle.cxx
View File

@ -17,22 +17,22 @@
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/testing/TestingVirtualObjectCache.h>
#include <vtkm/cont/testing/TestingVirtualObjectHandle.h>
namespace
{
void TestVirtualObjectCache()
void TestVirtualObjectHandle()
{
using DeviceAdapterList =
vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial, vtkm::cont::DeviceAdapterTagTBB>;
vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
vtkm::cont::testing::TestingVirtualObjectHandle<DeviceAdapterList>::Run();
}
} // anonymous namespace
int UnitTestTBBVirtualObjectCache(int, char* [])
int UnitTestTBBVirtualObjectHandle(int, char* [])
{
return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
return vtkm::cont::testing::Testing::Run(TestVirtualObjectHandle);
}

2
vtkm/cont/testing/CMakeLists.txt
View File

@ -31,7 +31,7 @@ set(headers
TestingFancyArrayHandles.h
TestingImplicitFunction.h
TestingPointLocatorUniformGrid.h
TestingVirtualObjectCache.h
TestingVirtualObjectHandle.h
)
vtkm_declare_headers(${headers})

70
vtkm/cont/testing/TestingDeviceAdapter.h
View File

@ -316,44 +316,31 @@ public:
class VirtualObjectTransferKernel
{
public:
struct Interface
struct Interface : public vtkm::VirtualObjectBase
{
using FooSig = vtkm::Id(const void*);
template <typename T>
VTKM_EXEC void Bind(const T* target)
{
this->Target = target;
this->FooPtr = [](const void* t) { return static_cast<const T*>(t)->Foo(); };
}
VTKM_EXEC
vtkm::Id Foo() const { return this->FooPtr(this->Target); }
const void* Target;
FooSig* FooPtr;
VTKM_EXEC virtual vtkm::Id Foo() const = 0;
};
struct Concrete
struct Concrete : public Interface
{
vtkm::Id Foo() const { return this->Value; }
VTKM_EXEC vtkm::Id Foo() const override { return this->Value; }
vtkm::Id Value = 0;
};
VirtualObjectTransferKernel(const Interface& vo, IdArrayHandle& result)
VirtualObjectTransferKernel(const Interface* vo, IdArrayHandle& result)
: Virtual(vo)
, Result(result.PrepareForInPlace(DeviceAdapterTag()))
{
}
VTKM_EXEC
void operator()(vtkm::Id) const { this->Result.Set(0, this->Virtual.Foo()); }
void operator()(vtkm::Id) const { this->Result.Set(0, this->Virtual->Foo()); }
VTKM_CONT void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {}
private:
Interface Virtual;
const Interface* Virtual;
IdPortalType Result;
};
@ -512,10 +499,9 @@ private:
VTKM_CONT
static void TestVirtualObjectTransfer()
{
using VirtualObject = typename VirtualObjectTransferKernel::Interface;
using BaseType = typename VirtualObjectTransferKernel::Interface;
using TargetType = typename VirtualObjectTransferKernel::Concrete;
using Transfer =
vtkm::cont::internal::VirtualObjectTransfer<VirtualObject, TargetType, DeviceAdapterTag>;
using Transfer = vtkm::cont::internal::VirtualObjectTransfer<TargetType, DeviceAdapterTag>;
IdArrayHandle result;
result.Allocate(1);
@ -524,21 +510,21 @@ private:
TargetType target;
target.Value = 5;
VirtualObject vo;
void* state = Transfer::Create(vo, &target);
Transfer transfer(&target);
const BaseType* base = transfer.PrepareForExecution(false);
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing VirtualObjectTransfer" << std::endl;
Algorithm::Schedule(VirtualObjectTransferKernel(vo, result), 1);
Algorithm::Schedule(VirtualObjectTransferKernel(base, result), 1);
VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 5, "Did not get expected result");
target.Value = 10;
Transfer::Update(state, &target);
Algorithm::Schedule(VirtualObjectTransferKernel(vo, result), 1);
base = transfer.PrepareForExecution(true);
Algorithm::Schedule(VirtualObjectTransferKernel(base, result), 1);
VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 10, "Did not get expected result");
Transfer::Cleanup(state);
transfer.ReleaseResources();
}
static VTKM_CONT void TestAlgorithmSchedule()
@ -1839,6 +1825,15 @@ private:
}
};
template <typename T, typename U>
struct TestCopy<vtkm::Pair<T, U>>
{
static vtkm::Pair<T, U> get(vtkm::Id i)
{
return vtkm::make_Pair(TestCopy<T>::get(i), TestCopy<U>::get(i));
}
};
template <typename T>
static VTKM_CONT void TestCopyArrays()
{
@ -1862,9 +1857,10 @@ private:
VTKM_TEST_ASSERT(temp.GetNumberOfValues() == COPY_ARRAY_SIZE, "Copy Needs to Resize Array");
typename std::vector<T>::const_iterator c = testData.begin();
const auto& portal = temp.GetPortalConstControl();
for (vtkm::Id i = 0; i < COPY_ARRAY_SIZE; i += 50, c += 50)
{
T value = temp.GetPortalConstControl().Get(i);
T value = portal.Get(i);
VTKM_TEST_ASSERT(value == *c, "Got bad value (Copy)");
}
}
@ -2015,19 +2011,15 @@ private:
{
std::cout << "-------------------------------------------------" << std::endl;
std::cout << "Testing Copy to same array type" << std::endl;
TestCopyArrays<vtkm::Vec<vtkm::Float32, 4>>();
TestCopyArrays<vtkm::Vec<vtkm::Float64, 4>>();
TestCopyArrays<vtkm::Vec<vtkm::Float32, 3>>();
TestCopyArrays<vtkm::Vec<vtkm::UInt8, 4>>();
//
TestCopyArrays<vtkm::Vec<vtkm::UInt8, 2>>();
TestCopyArrays<vtkm::Vec<vtkm::UInt16, 2>>();
TestCopyArrays<vtkm::Vec<vtkm::UInt32, 2>>();
TestCopyArrays<vtkm::Vec<vtkm::UInt64, 2>>();
TestCopyArrays<vtkm::Pair<vtkm::Id, vtkm::Float32>>();
TestCopyArrays<vtkm::Pair<vtkm::Id, vtkm::Vec<vtkm::Float32, 3>>>();
//
TestCopyArrays<vtkm::Float32>();
TestCopyArrays<vtkm::Float64>();
//
TestCopyArrays<vtkm::Int8>();
TestCopyArrays<vtkm::Int16>();
TestCopyArrays<vtkm::Int32>();
TestCopyArrays<vtkm::Int64>();
//
@ -2035,8 +2027,6 @@ private:
TestCopyArrays<vtkm::UInt16>();
TestCopyArrays<vtkm::UInt32>();
TestCopyArrays<vtkm::UInt64>();
//
TestCopyArrays<vtkm::Id>();
}
static VTKM_CONT void TestCopyArraysInDiffTypes()

59
vtkm/cont/testing/TestingImplicitFunction.h
View File

@ -22,7 +22,7 @@
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DeviceAdapterListTag.h>
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/cont/internal/DeviceAdapterTag.h>
#include <vtkm/cont/testing/MakeTestDataSet.h>
#include <vtkm/cont/testing/Testing.h>
@ -50,7 +50,7 @@ public:
typedef void ControlSignature(FieldIn<Vec3>, FieldOut<Scalar>);
typedef void ExecutionSignature(_1, _2);
EvaluateImplicitFunction(const vtkm::exec::ImplicitFunction& function)
EvaluateImplicitFunction(const vtkm::ImplicitFunction* function)
: Function(function)
{
}
@ -58,16 +58,16 @@ public:
template <typename VecType, typename ScalarType>
VTKM_EXEC void operator()(const VecType& point, ScalarType& val) const
{
val = this->Function.Value(point);
val = this->Function->Value(point);
}
private:
vtkm::exec::ImplicitFunction Function;
const vtkm::ImplicitFunction* Function;
};
template <typename DeviceAdapter>
void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points,
const vtkm::cont::ImplicitFunction& function,
const vtkm::cont::ImplicitFunctionHandle& function,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>& values,
DeviceAdapter device)
{
@ -141,13 +141,16 @@ private:
template <typename DeviceAdapter>
void TestBox(DeviceAdapter device)
{
std::cout << "Testing vtkm::cont::Box on "
std::cout << "Testing vtkm::Box on "
<< vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName() << "\n";
vtkm::cont::Box box({ 0.0f, -0.5f, -0.5f }, { 1.5f, 1.5f, 0.5f });
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), box, values, device);
this->Input.GetCoordinateSystem(0),
vtkm::cont::make_ImplicitFunctionHandle(
vtkm::Box({ 0.0f, -0.5f, -0.5f }, { 1.5f, 1.5f, 0.5f })),
values,
device);
std::array<vtkm::FloatDefault, 8> expected = {
{ 0.0f, -0.5f, 0.5f, 0.5f, 0.0f, -0.5f, 0.5f, 0.5f }
@ -159,17 +162,20 @@ private:
template <typename DeviceAdapter>
void TestCylinder(DeviceAdapter device)
{
std::cout << "Testing vtkm::cont::Cylinder on "
std::cout << "Testing vtkm::Cylinder on "
<< vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName() << "\n";
vtkm::cont::Cylinder cylinder;
vtkm::Cylinder cylinder;
cylinder.SetCenter({ 0.0f, 0.0f, 1.0f });
cylinder.SetAxis({ 0.0f, 1.0f, 0.0f });
cylinder.SetRadius(1.0f);
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), cylinder, values, device);
this->Input.GetCoordinateSystem(0),
vtkm::cont::ImplicitFunctionHandle(&cylinder, false),
values,
device);
std::array<vtkm::FloatDefault, 8> expected = {
{ 0.0f, 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f, -1.0f }
@ -181,7 +187,7 @@ private:
template <typename DeviceAdapter>
void TestFrustum(DeviceAdapter device)
{
std::cout << "Testing vtkm::cont::Frustum on "
std::cout << "Testing vtkm::Frustum on "
<< vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName() << "\n";
vtkm::Vec<vtkm::FloatDefault, 3> points[8] = {
@ -194,12 +200,15 @@ private:
{ 1.5f, 1.5f, 1.5f }, // 6
{ 0.5f, 1.5f, 1.5f } // 7
};
vtkm::cont::Frustum frustum;
vtkm::Frustum frustum;
frustum.CreateFromPoints(points);
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), frustum, values, device);
this->Input.GetCoordinateSystem(0),
vtkm::cont::make_ImplicitFunctionHandle(frustum),
values,
device);
std::array<vtkm::FloatDefault, 8> expected = {
{ 0.0f, 0.0f, 0.0f, 0.0f, 0.316228f, 0.316228f, -0.316228f, 0.316228f }
@ -211,23 +220,25 @@ private:
template <typename DeviceAdapter>
void TestPlane(DeviceAdapter device)
{
std::cout << "Testing vtkm::cont::Plane on "
std::cout << "Testing vtkm::Plane on "
<< vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName() << "\n";
vtkm::cont::Plane plane({ 0.5f, 0.5f, 0.5f }, { 1.0f, 0.0f, 1.0f });
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
auto planeHandle = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Plane>(
vtkm::make_Vec(0.5f, 0.5f, 0.5f), vtkm::make_Vec(1.0f, 0.0f, 1.0f));
auto plane = static_cast<vtkm::Plane*>(planeHandle.Get());
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), plane, values, device);
this->Input.GetCoordinateSystem(0), planeHandle, values, device);
std::array<vtkm::FloatDefault, 8> expected1 = {
{ -1.0f, 0.0f, 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f }
};
VTKM_TEST_ASSERT(implicit_function_detail::TestArrayEqual(values, expected1),
"Result does not match expected values");
plane.SetNormal({ -1.0f, 0.0f, -1.0f });
plane->SetNormal({ -1.0f, 0.0f, -1.0f });
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), plane, values, device);
this->Input.GetCoordinateSystem(0), planeHandle, values, device);
std::array<vtkm::FloatDefault, 8> expected2 = {
{ 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f, -1.0f, 0.0f }
};
@ -238,13 +249,15 @@ private:
template <typename DeviceAdapter>
void TestSphere(DeviceAdapter device)
{
std::cout << "Testing vtkm::cont::Sphere on "
std::cout << "Testing vtkm::Sphere on "
<< vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName() << "\n";
vtkm::cont::Sphere sphere({ 0.0f, 0.0f, 0.0f }, 1.0f);
vtkm::cont::ArrayHandle<vtkm::FloatDefault> values;
implicit_function_detail::EvaluateOnCoordinates(
this->Input.GetCoordinateSystem(0), sphere, values, device);
this->Input.GetCoordinateSystem(0),
vtkm::cont::make_ImplicitFunctionHandle<vtkm::Sphere>(vtkm::make_Vec(0.0f, 0.0f, 0.0f), 1.0f),
values,
device);
std::array<vtkm::FloatDefault, 8> expected = {
{ -1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 2.0f, 1.0f }

110
vtkm/cont/testing/TestingVirtualObjectCache.h → vtkm/cont/testing/TestingVirtualObjectHandle.h
View File

@ -17,14 +17,14 @@
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_testing_TestingVirtualObjectCache_h
#define vtk_m_cont_testing_TestingVirtualObjectCache_h
#ifndef vtk_m_cont_testing_TestingVirtualObjectHandle_h
#define vtk_m_cont_testing_TestingVirtualObjectHandle_h
#include <vtkm/Types.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleTransform.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/VirtualObjectCache.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/cont/testing/Testing.h>
#define ARRAY_LEN 8
@ -39,62 +39,76 @@ namespace testing
namespace virtual_object_detail
{
class Transformer
class Transformer : public vtkm::VirtualObjectBase
{
public:
template <typename T>
VTKM_EXEC void Bind(const T* target)
VTKM_EXEC
virtual vtkm::FloatDefault Eval(vtkm::FloatDefault val) const = 0;
};
class Square : public Transformer
{
public:
VTKM_EXEC
vtkm::FloatDefault Eval(vtkm::FloatDefault val) const override { return val * val; }
};
class Multiply : public Transformer
{
public:
VTKM_CONT
void SetMultiplicand(vtkm::FloatDefault val)
{
this->Concrete = target;
this->Caller = [](const void* concrete, vtkm::FloatDefault val) {
return static_cast<const T*>(concrete)->operator()(val);
};
this->Multiplicand = val;
this->Modified();
}
VTKM_CONT
vtkm::FloatDefault GetMultiplicand() const { return this->Multiplicand; }
VTKM_EXEC
vtkm::FloatDefault operator()(vtkm::FloatDefault val) const
vtkm::FloatDefault Eval(vtkm::FloatDefault val) const override
{
return this->Caller(this->Concrete, val);
return val * this->Multiplicand;
}
private:
using Signature = vtkm::FloatDefault(const void*, vtkm::FloatDefault);
const void* Concrete;
Signature* Caller;
vtkm::FloatDefault Multiplicand = 0.0f;
};
struct Square
class TransformerFunctor
{
VTKM_EXEC
vtkm::FloatDefault operator()(vtkm::FloatDefault val) const { return val * val; }
};
public:
TransformerFunctor() = default;
explicit TransformerFunctor(const Transformer* impl)
: Impl(impl)
{
}
struct Multiply
{
VTKM_EXEC
vtkm::FloatDefault operator()(vtkm::FloatDefault val) const { return val * this->Multiplicand; }
vtkm::FloatDefault operator()(vtkm::FloatDefault val) const { return this->Impl->Eval(val); }
vtkm::FloatDefault Multiplicand;
private:
const Transformer* Impl;
};
} // virtual_object_detail
template <typename DeviceAdapterList>
class TestingVirtualObjectCache
class TestingVirtualObjectHandle
{
private:
using FloatArrayHandle = vtkm::cont::ArrayHandle<vtkm::FloatDefault>;
using ArrayTransform =
vtkm::cont::ArrayHandleTransform<FloatArrayHandle, virtual_object_detail::Transformer>;
using TransformerCache = vtkm::cont::VirtualObjectCache<virtual_object_detail::Transformer>;
vtkm::cont::ArrayHandleTransform<FloatArrayHandle, virtual_object_detail::TransformerFunctor>;
using TransformerHandle = vtkm::cont::VirtualObjectHandle<virtual_object_detail::Transformer>;
class TestStage1
{
public:
TestStage1(const FloatArrayHandle& input, TransformerCache& manager)
TestStage1(const FloatArrayHandle& input, TransformerHandle& handle)
: Input(&input)
, Manager(&manager)
, Handle(&handle)
{
}
@ -107,7 +121,9 @@ private:
for (int n = 0; n < 2; ++n)
{
ArrayTransform transformed(*this->Input, this->Manager->GetVirtualObject(device));
virtual_object_detail::TransformerFunctor tfnctr(this->Handle->PrepareForExecution(device));
ArrayTransform transformed(*this->Input, tfnctr);
FloatArrayHandle output;
Algorithm::Copy(transformed, output);
auto portal = output.GetPortalConstControl();
@ -120,14 +136,14 @@ private:
if (n == 0)
{
std::cout << "\tReleaseResources and test again..." << std::endl;
this->Manager->ReleaseResources();
this->Handle->ReleaseExecutionResources();
}
}
}
private:
const FloatArrayHandle* Input;
TransformerCache* Manager;
TransformerHandle* Handle;
};
class TestStage2
@ -135,10 +151,10 @@ private:
public:
TestStage2(const FloatArrayHandle& input,
virtual_object_detail::Multiply& mul,
TransformerCache& manager)
TransformerHandle& handle)
: Input(&input)
, Mul(&mul)
, Manager(&manager)
, Handle(&handle)
{
}
@ -149,17 +165,18 @@ private:
std::cout << "\tDeviceAdapter: " << vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName()
<< std::endl;
this->Mul->Multiplicand = 2;
this->Manager->SetRefreshFlag(true);
this->Mul->SetMultiplicand(2);
for (int n = 0; n < 2; ++n)
{
ArrayTransform transformed(*this->Input, this->Manager->GetVirtualObject(device));
virtual_object_detail::TransformerFunctor tfnctr(this->Handle->PrepareForExecution(device));
ArrayTransform transformed(*this->Input, tfnctr);
FloatArrayHandle output;
Algorithm::Copy(transformed, output);
auto portal = output.GetPortalConstControl();
for (vtkm::Id i = 0; i < ARRAY_LEN; ++i)
{
VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i) * this->Mul->Multiplicand,
VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i) * this->Mul->GetMultiplicand(),
"\tIncorrect result");
}
std::cout << "\tSuccess." << std::endl;
@ -167,8 +184,7 @@ private:
if (n == 0)
{
std::cout << "\tUpdate and test again..." << std::endl;
this->Mul->Multiplicand = 3;
this->Manager->SetRefreshFlag(true);
this->Mul->SetMultiplicand(3);
}
}
}
@ -176,7 +192,7 @@ private:
private:
const FloatArrayHandle* Input;
virtual_object_detail::Multiply* Mul;
TransformerCache* Manager;
TransformerHandle* Handle;
};
public:
@ -190,20 +206,20 @@ public:
portal.Set(i, vtkm::FloatDefault(i));
}
TransformerCache manager;
TransformerHandle handle;
std::cout << "Testing with concrete type 1 (Square)..." << std::endl;
virtual_object_detail::Square sqr;
manager.Bind(&sqr, DeviceAdapterList());
vtkm::ListForEach(TestStage1(input, manager), DeviceAdapterList());
handle.Reset(&sqr, false, DeviceAdapterList());
vtkm::ListForEach(TestStage1(input, handle), DeviceAdapterList());
std::cout << "Reset..." << std::endl;
manager.Reset();
handle.Reset();
std::cout << "Testing with concrete type 2 (Multiply)..." << std::endl;
virtual_object_detail::Multiply mul;
manager.Bind(&mul, DeviceAdapterList());
vtkm::ListForEach(TestStage2(input, mul, manager), DeviceAdapterList());
handle.Reset(&mul, false, DeviceAdapterList());
vtkm::ListForEach(TestStage2(input, mul, handle), DeviceAdapterList());
}
};
}

12
vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
View File

@ -90,12 +90,14 @@ public:
}
};
template <typename VirtualObject, typename TargetClass>
struct VirtualObjectTransfer<VirtualObject,
TargetClass,
vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral>
: public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
template <typename TargetClass>
struct VirtualObjectTransfer<TargetClass, vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral>
: public VirtualObjectTransferShareWithControl<TargetClass>
{
VirtualObjectTransfer(const TargetClass* target)
: VirtualObjectTransferShareWithControl<TargetClass>(target)
{
}
};
template <typename T>

27
vtkm/cont/testing/UnitTestDynamicArrayHandle.cxx
View File

@ -106,14 +106,12 @@ struct TestValueFunctor
T operator()(vtkm::Id index) const { return TestValue(index, T()); }
};
bool CheckCalled;
struct CheckFunctor
{
template <typename T, typename Storage>
void operator()(vtkm::cont::ArrayHandle<T, Storage> array) const
void operator()(vtkm::cont::ArrayHandle<T, Storage> array, bool& called) const
{
CheckCalled = true;
called = true;
std::cout << " Checking for type: " << typeid(T).name() << std::endl;
VTKM_TEST_ASSERT(array.GetNumberOfValues() == ARRAY_SIZE, "Unexpected array size.");
@ -138,10 +136,11 @@ void CheckDynamicArray(vtkm::cont::DynamicArrayHandle array, vtkm::IdComponent n
{
BasicDynamicArrayChecks(array, numComponents);
array.CastAndCall(CheckFunctor());
bool called = false;
array.CastAndCall(CheckFunctor(), called);
VTKM_TEST_ASSERT(
CheckCalled, "The functor was never called (and apparently a bad value exception not thrown).");
called, "The functor was never called (and apparently a bad value exception not thrown).");
}
template <typename TypeList, typename StorageList>
@ -150,10 +149,11 @@ void CheckDynamicArray(vtkm::cont::DynamicArrayHandleBase<TypeList, StorageList>
{
BasicDynamicArrayChecks(array, numComponents);
CastAndCall(array, CheckFunctor());
bool called = false;
CastAndCall(array, CheckFunctor(), called);
VTKM_TEST_ASSERT(
CheckCalled, "The functor was never called (and apparently a bad value exception not thrown).");
called, "The functor was never called (and apparently a bad value exception not thrown).");
}
template <typename T>
@ -223,8 +223,6 @@ void TryNewInstance(T, DynamicArrayType originalArray)
template <typename T>
void TryDefaultType(T)
{
CheckCalled = false;
vtkm::cont::DynamicArrayHandle array = CreateDynamicArray(T());
CheckDynamicArray(array, vtkm::VecTraits<T>::NUM_COMPONENTS);
@ -237,8 +235,6 @@ struct TryBasicVTKmType
template <typename T>
void operator()(T) const
{
CheckCalled = false;
vtkm::cont::DynamicArrayHandle array = CreateDynamicArray(T());
CheckDynamicArray(array.ResetTypeList(vtkm::TypeListTagAll()),
@ -263,10 +259,7 @@ void TryUnusualType()
std::cout << " Caught exception for unrecognized type." << std::endl;
}
CheckCalled = false;
CheckDynamicArray(array.ResetTypeList(TypeListTagString()), 1);
VTKM_TEST_ASSERT(
CheckCalled, "The functor was never called (and apparently a bad value exception not thrown).");
std::cout << " Found type when type list was reset." << std::endl;
}
@ -284,7 +277,6 @@ void TryUnusualStorage()
std::cout << " Caught exception for unrecognized storage." << std::endl;
}
CheckCalled = false;
CheckDynamicArray(array.ResetStorageList(StorageListTagUnusual()), 1);
std::cout << " Found instance when storage list was reset." << std::endl;
}
@ -334,17 +326,14 @@ void TryUnusualTypeAndStorage()
VTKM_TEST_FAIL("ResetTypeAndStorageLists should have handled the custom type/storage.");
}
CheckCalled = false;
CheckDynamicArray(
array.ResetTypeList(TypeListTagString()).ResetStorageList(StorageListTagUnusual()), 1);
std::cout << " Found instance when type and storage lists were reset." << std::endl;
CheckCalled = false;
CheckDynamicArray(
array.ResetStorageList(StorageListTagUnusual()).ResetTypeList(TypeListTagString()), 1);
std::cout << " Found instance when storage and type lists were reset." << std::endl;
CheckCalled = false;
CheckDynamicArray(array.ResetTypeAndStorageLists(TypeListTagString(), StorageListTagUnusual()),
1);
std::cout << " Found instance when storage and type lists were reset." << std::endl;

19
vtkm/cont/testing/UnitTestDynamicCellSet.cxx
View File

@ -34,17 +34,16 @@ struct NonDefaultCellSetList
{
};
bool CheckCalled;
template <typename ExpectedCellType>
struct CheckFunctor
{
void operator()(const ExpectedCellType&) const { CheckCalled = true; }
void operator()(const ExpectedCellType&, bool& called) const { called = true; }
template <typename UnexpectedType>
void operator()(const UnexpectedType&) const
void operator()(const UnexpectedType&, bool& called) const
{
VTKM_TEST_FAIL("CastAndCall functor called with wrong type.");
called = false;
}
};
@ -60,17 +59,17 @@ void CheckDynamicCellSet(const CellSetType& cellSet,
dynamicCellSet.template Cast<CellSetType>();
CheckCalled = false;
dynamicCellSet.CastAndCall(CheckFunctor<CellSetType>());
bool called = false;
dynamicCellSet.CastAndCall(CheckFunctor<CellSetType>(), called);
VTKM_TEST_ASSERT(
CheckCalled, "The functor was never called (and apparently a bad value exception not thrown).");
called, "The functor was never called (and apparently a bad value exception not thrown).");
CheckCalled = false;
CastAndCall(dynamicCellSet, CheckFunctor<CellSetType>());
called = false;
CastAndCall(dynamicCellSet, CheckFunctor<CellSetType>(), called);
VTKM_TEST_ASSERT(
CheckCalled, "The functor was never called (and apparently a bad value exception not thrown).");
called, "The functor was never called (and apparently a bad value exception not thrown).");
}
template <typename CellSetType, typename CellSetList>

121
vtkm/cont/testing/UnitTestTryExecute.cxx
View File

@ -21,6 +21,7 @@
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_ERROR
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
@ -36,30 +37,27 @@ static const vtkm::Id ARRAY_SIZE = 10;
struct TryExecuteTestFunctor
{
vtkm::IdComponent NumCalls;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> InArray;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> OutArray;
VTKM_CONT
TryExecuteTestFunctor(vtkm::cont::ArrayHandle<vtkm::FloatDefault> inArray,
vtkm::cont::ArrayHandle<vtkm::FloatDefault> outArray)
TryExecuteTestFunctor()
: NumCalls(0)
, InArray(inArray)
, OutArray(outArray)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device)
VTKM_CONT bool operator()(Device,
const vtkm::cont::ArrayHandle<vtkm::FloatDefault>& in,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>& out)
{
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<Device>;
Algorithm::Copy(this->InArray, this->OutArray);
Algorithm::Copy(in, out);
this->NumCalls++;
return true;
}
};
template <typename DeviceList>
void TryExecuteWithList(DeviceList, bool expectSuccess)
void TryExecuteWithDevice(DeviceList, bool expectSuccess)
{
vtkm::cont::ArrayHandle<vtkm::FloatDefault> inArray;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> outArray;
@ -67,9 +65,9 @@ void TryExecuteWithList(DeviceList, bool expectSuccess)
inArray.Allocate(ARRAY_SIZE);
SetPortal(inArray.GetPortalControl());
TryExecuteTestFunctor functor(inArray, outArray);
TryExecuteTestFunctor functor;
bool result = vtkm::cont::TryExecute(functor, DeviceList());
bool result = vtkm::cont::TryExecute(functor, DeviceList(), inArray, outArray);
if (expectSuccess)
{
@ -81,6 +79,96 @@ void TryExecuteWithList(DeviceList, bool expectSuccess)
{
VTKM_TEST_ASSERT(!result, "Call returned true when expected failure.");
}
//verify the ability to pass rvalue functors
vtkm::cont::ArrayHandle<vtkm::FloatDefault> outArray2;
result = vtkm::cont::TryExecute(TryExecuteTestFunctor(), DeviceList(), inArray, outArray2);
if (expectSuccess)
{
VTKM_TEST_ASSERT(result, "Call returned failure when expected success.");
CheckPortal(outArray2.GetPortalConstControl());
}
else
{
VTKM_TEST_ASSERT(!result, "Call returned true when expected failure.");
}
}
template <typename DeviceList>
void TryExecuteAllExplicit(DeviceList, bool expectSuccess)
{
vtkm::cont::RuntimeDeviceTracker tracker;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> inArray;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> outArray;
inArray.Allocate(ARRAY_SIZE);
SetPortal(inArray.GetPortalControl());
bool result =
vtkm::cont::TryExecute(TryExecuteTestFunctor(), tracker, DeviceList(), inArray, outArray);
if (expectSuccess)
{
VTKM_TEST_ASSERT(result, "Call returned failure when expected success.");
CheckPortal(outArray.GetPortalConstControl());
}
else
{
VTKM_TEST_ASSERT(!result, "Call returned true when expected failure.");
}
}
struct EdgeCaseFunctor
{
template <typename DeviceList>
bool operator()(DeviceList, int, float, bool) const
{
return true;
}
template <typename DeviceList>
bool operator()(DeviceList) const
{
return true;
}
};
void TryExecuteAllEdgeCases()
{
using ValidDevice = vtkm::cont::DeviceAdapterTagSerial;
using SingleValidList = vtkm::ListTagBase<ValidDevice>;
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
std::cout << "TryExecute no Runtime, no Device, no parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor());
std::cout << "TryExecute no Runtime, no Device, with parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor(), int{ 42 }, float{ 3.14f }, bool{ true });
std::cout << "TryExecute with Runtime, no Device, no parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor(), tracker);
std::cout << "TryExecute with Runtime, no Device, with parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor(), tracker, int{ 42 }, float{ 3.14f }, bool{ true });
std::cout << "TryExecute no Runtime, with Device, no parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor(), SingleValidList());
std::cout << "TryExecute no Runtime, with Device, with parameters." << std::endl;
vtkm::cont::TryExecute(
EdgeCaseFunctor(), SingleValidList(), int{ 42 }, float{ 3.14f }, bool{ true });
std::cout << "TryExecute with Runtime, with Device, no parameters." << std::endl;
vtkm::cont::TryExecute(EdgeCaseFunctor(), tracker, SingleValidList());
std::cout << "TryExecute with Runtime, with Device, with parameters." << std::endl;
vtkm::cont::TryExecute(
EdgeCaseFunctor(), tracker, SingleValidList(), int{ 42 }, float{ 3.14f }, bool{ true });
}
template <typename DeviceList>
void TryExecuteTests(DeviceList list, bool expectSuccess)
{
TryExecuteAllExplicit(list, expectSuccess);
TryExecuteWithDevice(list, expectSuccess);
}
static void Run()
@ -88,21 +176,24 @@ static void Run()
using ValidDevice = vtkm::cont::DeviceAdapterTagSerial;
using InvalidDevice = vtkm::cont::DeviceAdapterTagError;
TryExecuteAllEdgeCases();
std::cout << "Try a list with a single entry." << std::endl;
using SingleValidList = vtkm::ListTagBase<ValidDevice>;
TryExecuteWithList(SingleValidList(), true);
TryExecuteTests(SingleValidList(), true);
std::cout << "Try a list with two valid devices." << std::endl;
using DoubleValidList = vtkm::ListTagBase<ValidDevice, ValidDevice>;
TryExecuteWithList(DoubleValidList(), true);
TryExecuteTests(DoubleValidList(), true);
std::cout << "Try a list with only invalid device." << std::endl;
using SingleInvalidList = vtkm::ListTagBase<InvalidDevice>;
TryExecuteWithList(SingleInvalidList(), false);
TryExecuteTests(SingleInvalidList(), false);
std::cout << "Try a list with an invalid and valid device." << std::endl;
using InvalidAndValidList = vtkm::ListTagBase<InvalidDevice, ValidDevice>;
TryExecuteWithList(InvalidAndValidList(), true);
TryExecuteTests(InvalidAndValidList(), true);
}
} // anonymous namespace

22
vtkm/filter/ClipWithImplicitFunction.h
View File

@ -21,12 +21,10 @@
#ifndef vtk_m_filter_ClipWithImplicitFunction_h
#define vtk_m_filter_ClipWithImplicitFunction_h
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/filter/FilterDataSet.h>
#include <vtkm/worklet/Clip.h>
#include <memory>
namespace vtkm
{
namespace filter
@ -34,26 +32,18 @@ namespace filter
/// \brief Clip a dataset using an implicit function
///
/// Clip a dataset using a given implicit function value, such as vtkm::cont::Sphere
/// or vtkm::cont::Frustum.
/// Clip a dataset using a given implicit function value, such as vtkm::Sphere
/// or vtkm::Frustum.
/// The resulting geometry will not be water tight.
class ClipWithImplicitFunction : public vtkm::filter::FilterDataSet<ClipWithImplicitFunction>
{
public:
template <typename ImplicitFunctionType, typename DerivedPolicy>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>& policy);
template <typename ImplicitFunctionType>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func)
void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func)
{
this->Function = func;
}
std::shared_ptr<vtkm::cont::ImplicitFunction> GetImplicitFunction() const
{
return this->Function;
}
const vtkm::cont::ImplicitFunctionHandle& GetImplicitFunction() const { return this->Function; }
template <typename DerivedPolicy, typename DeviceAdapter>
vtkm::filter::Result DoExecute(const vtkm::cont::DataSet& input,
@ -70,7 +60,7 @@ public:
const DeviceAdapter& tag);
private:
std::shared_ptr<vtkm::cont::ImplicitFunction> Function;
vtkm::cont::ImplicitFunctionHandle Function;
vtkm::worklet::Clip Worklet;
};
}

12
vtkm/filter/ClipWithImplicitFunction.hxx
View File

@ -53,16 +53,6 @@ struct PointMapHelper
} // end namespace clipwithimplicitfunction
//-----------------------------------------------------------------------------
template <typename ImplicitFunctionType, typename DerivedPolicy>
inline void ClipWithImplicitFunction::SetImplicitFunction(
const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>&)
{
func->ResetDevices(DerivedPolicy::DeviceAdapterList);
this->Function = func;
}
//-----------------------------------------------------------------------------
template <typename DerivedPolicy, typename DeviceAdapter>
inline vtkm::filter::Result ClipWithImplicitFunction::DoExecute(
@ -79,7 +69,7 @@ inline vtkm::filter::Result ClipWithImplicitFunction::DoExecute(
input.GetCoordinateSystem(this->GetActiveCoordinateSystemIndex());
vtkm::cont::CellSetExplicit<> outputCellSet = this->Worklet.Run(
vtkm::filter::ApplyPolicy(cells, policy), *this->Function, inputCoords, device);
vtkm::filter::ApplyPolicy(cells, policy), this->Function, inputCoords, device);
// compute output coordinates
vtkm::cont::DynamicArrayHandle outputCoordsArray;

16
vtkm/filter/ExtractGeometry.h
View File

@ -21,7 +21,7 @@
#ifndef vtk_m_filter_ExtractGeometry_h
#define vtk_m_filter_ExtractGeometry_h
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/filter/FilterDataSet.h>
#include <vtkm/worklet/ExtractGeometry.h>
@ -52,20 +52,12 @@ public:
ExtractGeometry();
// Set the volume of interest to extract
template <typename ImplicitFunctionType, typename DerivedPolicy>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>& policy);
template <typename ImplicitFunctionType>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func)
void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func)
{
this->Function = func;
}
std::shared_ptr<vtkm::cont::ImplicitFunction> GetImplicitFunction() const
{
return this->Function;
}
const vtkm::cont::ImplicitFunctionHandle& GetImplicitFunction() const { return this->Function; }
VTKM_CONT
bool GetExtractInside() { return this->ExtractInside; }
@ -111,7 +103,7 @@ private:
bool ExtractInside;
bool ExtractBoundaryCells;
bool ExtractOnlyBoundaryCells;
std::shared_ptr<vtkm::cont::ImplicitFunction> Function;
vtkm::cont::ImplicitFunctionHandle Function;
vtkm::worklet::ExtractGeometry Worklet;
};

15
vtkm/filter/ExtractGeometry.hxx
View File

@ -34,7 +34,7 @@ struct CallWorker
vtkm::cont::DynamicCellSet& Output;
vtkm::worklet::ExtractGeometry& Worklet;
const vtkm::cont::CoordinateSystem& Coords;
const vtkm::cont::ImplicitFunction& Function;
const vtkm::cont::ImplicitFunctionHandle& Function;
bool ExtractInside;
bool ExtractBoundaryCells;
bool ExtractOnlyBoundaryCells;
@ -42,7 +42,7 @@ struct CallWorker
CallWorker(vtkm::cont::DynamicCellSet& output,
vtkm::worklet::ExtractGeometry& worklet,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::cont::ImplicitFunction& function,
const vtkm::cont::ImplicitFunctionHandle& function,
bool extractInside,
bool extractBoundaryCells,
bool extractOnlyBoundaryCells)
@ -76,15 +76,6 @@ namespace vtkm
namespace filter
{
//-----------------------------------------------------------------------------
template <typename ImplicitFunctionType, typename DerivedPolicy>
inline void ExtractGeometry::SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>&)
{
func->ResetDevices(DerivedPolicy::DeviceAdapterList);
this->Function = func;
}
//-----------------------------------------------------------------------------
inline VTKM_CONT ExtractGeometry::ExtractGeometry()
: vtkm::filter::FilterDataSet<ExtractGeometry>()
@ -110,7 +101,7 @@ inline VTKM_CONT vtkm::filter::Result ExtractGeometry::DoExecute(
CallWorker<DeviceAdapter> worker(outCells,
this->Worklet,
coords,
*this->Function,
this->Function,
this->ExtractInside,
this->ExtractBoundaryCells,
this->ExtractOnlyBoundaryCells);

17
vtkm/filter/ExtractPoints.h
View File

@ -21,7 +21,7 @@
#ifndef vtk_m_filter_ExtractPoints_h
#define vtk_m_filter_ExtractPoints_h
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/filter/CleanGrid.h>
#include <vtkm/filter/FilterDataSet.h>
#include <vtkm/worklet/ExtractPoints.h>
@ -53,21 +53,12 @@ public:
void SetCompactPoints(bool value) { this->CompactPoints = value; }
/// Set the volume of interest to extract
template <typename ImplicitFunctionType, typename DerivedPolicy>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>& policy);
/// Set the volume of interest to extract
template <typename ImplicitFunctionType>
void SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func)
void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func)
{
this->Function = func;
}
std::shared_ptr<vtkm::cont::ImplicitFunction> GetImplicitFunction() const
{
return this->Function;
}
const vtkm::cont::ImplicitFunctionHandle& GetImplicitFunction() const { return this->Function; }
VTKM_CONT
bool GetExtractInside() { return this->ExtractInside; }
@ -93,7 +84,7 @@ public:
private:
bool ExtractInside;
std::shared_ptr<vtkm::cont::ImplicitFunction> Function;
vtkm::cont::ImplicitFunctionHandle Function;
bool CompactPoints;
vtkm::filter::CleanGrid Compactor;

11
vtkm/filter/ExtractPoints.hxx
View File

@ -44,15 +44,6 @@ namespace vtkm
namespace filter
{
//-----------------------------------------------------------------------------
template <typename ImplicitFunctionType, typename DerivedPolicy>
inline void ExtractPoints::SetImplicitFunction(const std::shared_ptr<ImplicitFunctionType>& func,
const vtkm::filter::PolicyBase<DerivedPolicy>&)
{
func->ResetDevices(DerivedPolicy::DeviceAdapterList);
this->Function = func;
}
//-----------------------------------------------------------------------------
inline VTKM_CONT ExtractPoints::ExtractPoints()
: vtkm::filter::FilterDataSet<ExtractPoints>()
@ -79,7 +70,7 @@ inline vtkm::filter::Result ExtractPoints::DoExecute(
outCellSet = worklet.Run(vtkm::filter::ApplyPolicy(cells, policy),
vtkm::filter::ApplyPolicy(coords, policy),
*this->Function,
this->Function,
this->ExtractInside,
device);

6
vtkm/filter/FilterDataSet.hxx
View File

@ -139,11 +139,11 @@ inline VTKM_CONT bool FilterDataSet<Derived>::MapFieldOntoOutput(
{
vtkm::filter::FieldMetadata metaData(field);
typedef internal::ResolveFieldTypeAndMap<Derived, DerivedPolicy> FunctorType;
FunctorType functor(
static_cast<Derived*>(this), result, metaData, policy, this->Tracker, valid);
FunctorType functor(static_cast<Derived*>(this), result, metaData, policy, valid);
typedef vtkm::filter::FilterTraits<Derived> Traits;
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy, Traits()), functor);
vtkm::cont::CastAndCall(
vtkm::filter::ApplyPolicy(field, policy, Traits()), functor, this->Tracker);
}
//the bool valid will be modified by the map algorithm to hold if the

15
vtkm/filter/FilterDataSetWithField.hxx
View File

@ -162,10 +162,11 @@ inline VTKM_CONT Result FilterDataSetWithField<Derived>::PrepareForExecution(
Result result;
typedef internal::ResolveFieldTypeAndExecute<Derived, DerivedPolicy, Result> FunctorType;
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, this->Tracker, result);
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, result);
typedef vtkm::filter::FilterTraits<Derived> Traits;
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy, Traits()), functor);
vtkm::cont::CastAndCall(
vtkm::filter::ApplyPolicy(field, policy, Traits()), functor, this->Tracker);
return result;
}
@ -184,10 +185,11 @@ inline VTKM_CONT Result FilterDataSetWithField<Derived>::PrepareForExecution(
//determine the field type first
Result result;
typedef internal::ResolveFieldTypeAndExecute<Derived, DerivedPolicy, Result> FunctorType;
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, this->Tracker, result);
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, result);
typedef vtkm::filter::FilterTraits<Derived> Traits;
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy, Traits()), functor);
vtkm::cont::CastAndCall(
vtkm::filter::ApplyPolicy(field, policy, Traits()), functor, this->Tracker);
return result;
}
@ -214,10 +216,9 @@ inline VTKM_CONT bool FilterDataSetWithField<Derived>::MapFieldOntoOutput(
{
vtkm::filter::FieldMetadata metaData(field);
typedef internal::ResolveFieldTypeAndMap<Derived, DerivedPolicy> FunctorType;
FunctorType functor(
static_cast<Derived*>(this), result, metaData, policy, this->Tracker, valid);
FunctorType functor(static_cast<Derived*>(this), result, metaData, policy, valid);
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy), functor);
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy), functor, this->Tracker);
}
//the bool valid will be modified by the map algorithm to hold if the

10
vtkm/filter/FilterField.hxx
View File

@ -159,10 +159,11 @@ FilterField<Derived>::PrepareForExecution(const vtkm::cont::DataSet& input,
Result result;
typedef internal::ResolveFieldTypeAndExecute<Derived, DerivedPolicy, Result> FunctorType;
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, this->Tracker, result);
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, result);
typedef vtkm::filter::FilterTraits<Derived> Traits;
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy, Traits()), functor);
vtkm::cont::CastAndCall(
vtkm::filter::ApplyPolicy(field, policy, Traits()), functor, this->Tracker);
return result;
}
@ -181,10 +182,11 @@ FilterField<Derived>::PrepareForExecution(const vtkm::cont::DataSet& input,
Result result;
typedef internal::ResolveFieldTypeAndExecute<Derived, DerivedPolicy, Result> FunctorType;
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, this->Tracker, result);
FunctorType functor(static_cast<Derived*>(this), input, metaData, policy, result);
typedef vtkm::filter::FilterTraits<Derived> Traits;
vtkm::cont::CastAndCall(vtkm::filter::ApplyPolicy(field, policy, Traits()), functor);
vtkm::cont::CastAndCall(
vtkm::filter::ApplyPolicy(field, policy, Traits()), functor, this->Tracker);
return result;
}
}

50
vtkm/filter/internal/ResolveFieldTypeAndExecute.h
View File

@ -35,6 +35,18 @@ namespace filter
namespace internal
{
struct ResolveFieldTypeAndExecuteForDevice
{
template <typename DeviceAdapterTag, typename InstanceType, typename FieldType>
bool operator()(DeviceAdapterTag tag, InstanceType&& instance, FieldType&& field) const
{
instance.Result = instance.DerivedClass->DoExecute(
instance.InputData, field, instance.Metadata, instance.Policy, tag);
return instance.Result.IsValid();
}
};
template <typename Derived, typename DerivedPolicy, typename ResultType>
struct ResolveFieldTypeAndExecute
{
@ -44,56 +56,28 @@ struct ResolveFieldTypeAndExecute
const vtkm::cont::DataSet& InputData;
const vtkm::filter::FieldMetadata& Metadata;
const vtkm::filter::PolicyBase<DerivedPolicy>& Policy;
vtkm::cont::RuntimeDeviceTracker Tracker;
ResultType& Result;
ResolveFieldTypeAndExecute(Derived* derivedClass,
const vtkm::cont::DataSet& inputData,
const vtkm::filter::FieldMetadata& fieldMeta,
const vtkm::filter::PolicyBase<DerivedPolicy>& policy,
const vtkm::cont::RuntimeDeviceTracker& tracker,
ResultType& result)
: DerivedClass(derivedClass)
, InputData(inputData)
, Metadata(fieldMeta)
, Policy(policy)
, Tracker(tracker)
, Result(result)
{
}
private:
template <typename T, typename StorageTag>
struct ResolveFieldTypeAndExecuteForDevice
void operator()(const vtkm::cont::ArrayHandle<T, StorageTag>& field,
vtkm::cont::RuntimeDeviceTracker& tracker) const
{
typedef vtkm::cont::ArrayHandle<T, StorageTag> FieldArrayHandle;
ResolveFieldTypeAndExecuteForDevice(const Self& instance, const FieldArrayHandle& field)
: Instance(instance)
, Field(field)
{
}
const Self& Instance;
const vtkm::cont::ArrayHandle<T, StorageTag>& Field;
template <typename DeviceAdapterTag>
bool operator()(DeviceAdapterTag tag) const
{
this->Instance.Result = this->Instance.DerivedClass->DoExecute(
this->Instance.InputData, this->Field, this->Instance.Metadata, this->Instance.Policy, tag);
return this->Instance.Result.IsValid();
}
private:
void operator=(const ResolveFieldTypeAndExecuteForDevice<T, StorageTag>&) = delete;
};
public:
template <typename T, typename StorageTag>
void operator()(const vtkm::cont::ArrayHandle<T, StorageTag>& field) const
{
ResolveFieldTypeAndExecuteForDevice<T, StorageTag> doResolve(*this, field);
vtkm::cont::TryExecute(doResolve, this->Tracker, typename DerivedPolicy::DeviceAdapterList());
ResolveFieldTypeAndExecuteForDevice doResolve;
vtkm::cont::TryExecute(
doResolve, tracker, typename DerivedPolicy::DeviceAdapterList(), *this, field);
}
private:

54
vtkm/filter/internal/ResolveFieldTypeAndMap.h
View File

@ -43,6 +43,15 @@ namespace filter
{
namespace internal
{
struct ResolveFieldTypeAndMapForDevice
{
template <typename DeviceAdapterTag, typename InstanceType, typename FieldType>
bool operator()(DeviceAdapterTag tag, InstanceType&& instance, FieldType&& field) const
{
return instance.DerivedClass->DoMapField(
instance.InputResult, field, instance.Metadata, instance.Policy, tag);
}
};
template <typename Derived, typename DerivedPolicy>
struct ResolveFieldTypeAndMap
@ -53,63 +62,28 @@ struct ResolveFieldTypeAndMap
vtkm::filter::Result& InputResult;
const vtkm::filter::FieldMetadata& Metadata;
const vtkm::filter::PolicyBase<DerivedPolicy>& Policy;
vtkm::cont::RuntimeDeviceTracker Tracker;
bool& RanProperly;
ResolveFieldTypeAndMap(Derived* derivedClass,
vtkm::filter::Result& inResult,
const vtkm::filter::FieldMetadata& fieldMeta,
const vtkm::filter::PolicyBase<DerivedPolicy>& policy,
const vtkm::cont::RuntimeDeviceTracker& tracker,
bool& ran)
: DerivedClass(derivedClass)
, InputResult(inResult)
, Metadata(fieldMeta)
, Policy(policy)
, Tracker(tracker)
, RanProperly(ran)
{
}
private:
template <typename T, typename StorageTag>
struct ResolveFieldTypeAndMapForDevice
void operator()(const vtkm::cont::ArrayHandle<T, StorageTag>& field,
vtkm::cont::RuntimeDeviceTracker& tracker) const
{
typedef vtkm::cont::ArrayHandle<T, StorageTag> FieldArrayHandle;
ResolveFieldTypeAndMapForDevice(const Self& instance, const FieldArrayHandle& field)
: Instance(instance)
, Field(field)
, Valid(false)
{
}
const Self& Instance;
const vtkm::cont::ArrayHandle<T, StorageTag>& Field;
mutable bool Valid;
template <typename DeviceAdapterTag>
bool operator()(DeviceAdapterTag tag) const
{
this->Valid = this->Instance.DerivedClass->DoMapField(this->Instance.InputResult,
this->Field,
this->Instance.Metadata,
this->Instance.Policy,
tag);
return this->Valid;
}
private:
void operator=(const ResolveFieldTypeAndMapForDevice<T, StorageTag>&) = delete;
};
public:
template <typename T, typename StorageTag>
void operator()(const vtkm::cont::ArrayHandle<T, StorageTag>& field) const
{
ResolveFieldTypeAndMapForDevice<T, StorageTag> doResolve(*this, field);
vtkm::cont::TryExecute(doResolve, this->Tracker, typename DerivedPolicy::DeviceAdapterList());
this->RanProperly = doResolve.Valid;
ResolveFieldTypeAndMapForDevice doResolve;
this->RanProperly = vtkm::cont::TryExecute(
doResolve, tracker, typename DerivedPolicy::DeviceAdapterList(), *this, field);
}
private:

3
vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx
View File

@ -60,11 +60,10 @@ void TestClipStructured()
vtkm::Vec<vtkm::FloatDefault, 3> center(1, 1, 0);
vtkm::FloatDefault radius(0.5);
auto sphere = std::make_shared<vtkm::cont::Sphere>(center, radius);
vtkm::filter::Result result;
vtkm::filter::ClipWithImplicitFunction clip;
clip.SetImplicitFunction(sphere);
clip.SetImplicitFunction(vtkm::cont::make_ImplicitFunctionHandle(vtkm::Sphere(center, radius)));
result = clip.Execute(ds);
clip.MapFieldOntoOutput(result, ds.GetField("scalars"));

8
vtkm/filter/testing/UnitTestExtractGeometryFilter.cxx
View File

@ -40,7 +40,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(1.f, 1.f, 1.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.f, 3.f, 3.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractGeometry extractGeometry;
@ -76,7 +76,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(1.f, 1.f, 1.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.f, 3.f, 3.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractGeometry extractGeometry;
@ -112,7 +112,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(0.5f, 0.5f, 0.5f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.5f, 3.5f, 3.5f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractGeometry extractGeometry;
@ -147,7 +147,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(0.5f, 0.5f, 0.5f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.5f, 3.5f, 3.5f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractGeometry extractGeometry;

10
vtkm/filter/testing/UnitTestExtractPointsFilter.cxx
View File

@ -40,7 +40,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(1.f, 1.f, 1.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.f, 3.f, 3.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractPoints extractPoints;
@ -82,7 +82,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(1.f, 1.f, 1.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(3.f, 3.f, 3.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractPoints extractPoints;
@ -125,7 +125,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> center(2.f, 2.f, 2.f);
vtkm::FloatDefault radius(1.8f);
auto sphere = std::make_shared<vtkm::cont::Sphere>(center, radius);
auto sphere = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Sphere>(center, radius);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractPoints extractPoints;
@ -154,7 +154,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(0.f, 0.f, 0.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(1.f, 1.f, 1.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractPoints extractPoints;
@ -183,7 +183,7 @@ public:
// Implicit function
vtkm::Vec<vtkm::FloatDefault, 3> minPoint(0.f, 0.f, 0.f);
vtkm::Vec<vtkm::FloatDefault, 3> maxPoint(1.f, 1.f, 1.f);
auto box = std::make_shared<vtkm::cont::Box>(minPoint, maxPoint);
auto box = vtkm::cont::make_ImplicitFunctionHandle<vtkm::Box>(minPoint, maxPoint);
// Setup and run filter to extract by volume of interest
vtkm::filter::ExtractPoints extractPoints;

2
vtkm/internal/ExportMacros.h
View File

@ -89,8 +89,6 @@
#define VTKM_NOEXCEPT noexcept
#endif
#define VTKM_OVERRIDE override
// Clang will warn about weak vtables (-Wweak-vtables) on exception classes,
// but there's no good way to eliminate them in this case because MSVC (See
// http://stackoverflow.com/questions/24511376). These macros will silence the

75
vtkm/internal/ListTagDetail.h
View File

@ -165,30 +165,30 @@ struct ListIntersect<SameListTag, SameListTag>
using type = SameListTag;
};
template <typename Functor>
VTKM_CONT void ListForEachImpl(Functor&&, brigand::empty_sequence)
template <typename Functor, typename... Args>
VTKM_CONT void ListForEachImpl(Functor&&, brigand::list<>, Args&&...)
{
}
template <typename Functor, typename T1>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1>)
template <typename Functor, typename T1, typename... Args>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1>, Args&&... args)
{
f(T1{});
f(T1{}, std::forward<Args>(args)...);
}
template <typename Functor, typename T1, typename T2>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2>)
template <typename Functor, typename T1, typename T2, typename... Args>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2>, Args&&... args)
{
f(T1{});
f(T2{});
f(T1{}, std::forward<Args>(args)...);
f(T2{}, std::forward<Args>(args)...);
}
template <typename Functor, typename T1, typename T2, typename T3>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2, T3>)
template <typename Functor, typename T1, typename T2, typename T3, typename... Args>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2, T3>, Args&&... args)
{
f(T1{});
f(T2{});
f(T3{});
f(T1{}, std::forward<Args>(args)...);
f(T2{}, std::forward<Args>(args)...);
f(T3{}, std::forward<Args>(args)...);
}
template <typename Functor,
@ -196,16 +196,49 @@ template <typename Functor,
typename T2,
typename T3,
typename T4,
typename... ArgTypes>
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2, T3, T4, ArgTypes...>)
typename... ArgTypes,
typename... Args>
VTKM_CONT void ListForEachImpl(Functor&& f,
brigand::list<T1, T2, T3, T4, ArgTypes...>&&,
Args&&... args)
{
f(T1{});
f(T2{});
f(T3{});
f(T4{});
ListForEachImpl(f, brigand::list<ArgTypes...>());
f(T1{}, std::forward<Args>(args)...);
f(T2{}, std::forward<Args>(args)...);
f(T3{}, std::forward<Args>(args)...);
f(T4{}, std::forward<Args>(args)...);
ListForEachImpl(
std::forward<Functor>(f), brigand::list<ArgTypes...>{}, std::forward<Args>(args)...);
}
template <typename T, typename U, typename R>
struct ListCrossProductAppend
{
using type = brigand::push_back<T, std::pair<U, R>>;
};
template <typename T, typename U, typename R2>
struct ListCrossProductImplUnrollR2
{
using P =
brigand::fold<R2,
brigand::list<>,
ListCrossProductAppend<brigand::_state, brigand::_element, brigand::pin<U>>>;
using type = brigand::append<T, P>;
};
template <typename R1, typename R2>
struct ListCrossProductImpl
{
using type = brigand::fold<
R2,
brigand::list<>,
ListCrossProductImplUnrollR2<brigand::_state, brigand::_element, brigand::pin<R1>>>;
};
} // namespace detail
//-----------------------------------------------------------------------------

2
vtkm/rendering/AxisAnnotation2D.h
View File

@ -127,7 +127,7 @@ public:
void Render(const vtkm::rendering::Camera& camera,
const vtkm::rendering::WorldAnnotator& worldAnnotator,
vtkm::rendering::Canvas& canvas) VTKM_OVERRIDE;
vtkm::rendering::Canvas& canvas) override;
};
}
} //namespace vtkm::rendering

12
vtkm/rendering/AxisAnnotation3D.cxx
View File

@ -177,13 +177,13 @@ void AxisAnnotation3D::Render(const Camera& camera,
switch (this->Axis)
{
case 0:
tickSize[1] = this->TickMajorSize;
tickSize[1] = this->TickMinorSize;
break;
case 1:
tickSize[0] = this->TickMajorSize;
tickSize[0] = this->TickMinorSize;
break;
case 2:
tickSize[0] = this->TickMajorSize;
tickSize[0] = this->TickMinorSize;
break;
}
}
@ -192,13 +192,13 @@ void AxisAnnotation3D::Render(const Camera& camera,
switch (this->Axis)
{
case 0:
tickSize[2] = this->TickMajorSize;
tickSize[2] = this->TickMinorSize;
break;
case 1:
tickSize[2] = this->TickMajorSize;
tickSize[2] = this->TickMinorSize;
break;
case 2:
tickSize[1] = this->TickMajorSize;
tickSize[1] = this->TickMinorSize;
break;
}
}

2
vtkm/rendering/AxisAnnotation3D.h
View File

@ -123,7 +123,7 @@ public:
virtual void Render(const vtkm::rendering::Camera& camera,
const vtkm::rendering::WorldAnnotator& worldAnnotator,
vtkm::rendering::Canvas& canvas) VTKM_OVERRIDE;
vtkm::rendering::Canvas& canvas) override;
};
}
} //namespace vtkm::rendering

6
vtkm/rendering/CanvasEGL.h
View File

@ -44,11 +44,11 @@ public:
~CanvasEGL();
virtual void Initialize() VTKM_OVERRIDE;
virtual void Initialize() override;
virtual void Activate() VTKM_OVERRIDE;
virtual void Activate() override;
vtkm::rendering::Canvas* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Canvas* NewCopy() const override;
private:
std::shared_ptr<detail::CanvasEGLInternals> Internals;

30
vtkm/rendering/CanvasGL.h
View File

@ -38,43 +38,43 @@ public:
~CanvasGL();
void Initialize() VTKM_OVERRIDE;
void Initialize() override;
void Activate() VTKM_OVERRIDE;
void Activate() override;
void Clear() VTKM_OVERRIDE;
void Clear() override;
void Finish() VTKM_OVERRIDE;
void Finish() override;
vtkm::rendering::Canvas* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Canvas* NewCopy() const override;
void SetViewToWorldSpace(const vtkm::rendering::Camera& camera, bool clip) VTKM_OVERRIDE;
void SetViewToWorldSpace(const vtkm::rendering::Camera& camera, bool clip) override;
void SetViewToScreenSpace(const vtkm::rendering::Camera& camera, bool clip) VTKM_OVERRIDE;
void SetViewToScreenSpace(const vtkm::rendering::Camera& camera, bool clip) override;
void SetViewportClipping(const vtkm::rendering::Camera& camera, bool clip) VTKM_OVERRIDE;
void SetViewportClipping(const vtkm::rendering::Camera& camera, bool clip) override;
void RefreshColorBuffer() const VTKM_OVERRIDE;
void RefreshColorBuffer() const override;
virtual void RefreshDepthBuffer() const VTKM_OVERRIDE;
virtual void RefreshDepthBuffer() const override;
vtkm::rendering::WorldAnnotator* CreateWorldAnnotator() const VTKM_OVERRIDE;
vtkm::rendering::WorldAnnotator* CreateWorldAnnotator() const override;
protected:
void AddLine(const vtkm::Vec<vtkm::Float64, 2>& point0,
const vtkm::Vec<vtkm::Float64, 2>& point1,
vtkm::Float32 linewidth,
const vtkm::rendering::Color& color) const VTKM_OVERRIDE;
const vtkm::rendering::Color& color) const override;
void AddColorBar(const vtkm::Bounds& bounds,
const vtkm::rendering::ColorTable& colorTable,
bool horizontal) const VTKM_OVERRIDE;
bool horizontal) const override;
void AddColorSwatch(const vtkm::Vec<vtkm::Float64, 2>& point0,
const vtkm::Vec<vtkm::Float64, 2>& point1,
const vtkm::Vec<vtkm::Float64, 2>& point2,
const vtkm::Vec<vtkm::Float64, 2>& point3,
const vtkm::rendering::Color& color) const VTKM_OVERRIDE;
const vtkm::rendering::Color& color) const override;
void AddText(const vtkm::Vec<vtkm::Float32, 2>& position,
vtkm::Float32 scale,
@ -82,7 +82,7 @@ protected:
vtkm::Float32 windowAspect,
const vtkm::Vec<vtkm::Float32, 2>& anchor,
const vtkm::rendering::Color& color,
const std::string& text) const VTKM_OVERRIDE;
const std::string& text) const override;
private:
vtkm::rendering::BitmapFont Font;

10
vtkm/rendering/CanvasOSMesa.h
View File

@ -45,15 +45,15 @@ public:
~CanvasOSMesa();
virtual void Initialize() VTKM_OVERRIDE;
virtual void Initialize() override;
virtual void RefreshColorBuffer() const VTKM_OVERRIDE;
virtual void RefreshColorBuffer() const override;
virtual void Activate() VTKM_OVERRIDE;
virtual void Activate() override;
virtual void Finish() VTKM_OVERRIDE;
virtual void Finish() override;
vtkm::rendering::Canvas* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Canvas* NewCopy() const override;
private:
std::shared_ptr<detail::CanvasOSMesaInternals> Internals;

2
vtkm/rendering/CanvasRayTracer.h
View File

@ -37,7 +37,7 @@ public:
~CanvasRayTracer();
vtkm::rendering::Canvas* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Canvas* NewCopy() const override;
void WriteToCanvas(const vtkm::rendering::raytracing::Ray<vtkm::Float32>& rays,
const vtkm::cont::ArrayHandle<vtkm::Float32>& colors,

12
vtkm/rendering/MapperConnectivity.h
View File

@ -37,20 +37,20 @@ public:
~MapperConnectivity();
void SetSampleDistance(const vtkm::Float32&);
void SetCanvas(vtkm::rendering::Canvas* canvas) VTKM_OVERRIDE;
virtual vtkm::rendering::Canvas* GetCanvas() const VTKM_OVERRIDE;
void SetCanvas(vtkm::rendering::Canvas* canvas) override;
virtual vtkm::rendering::Canvas* GetCanvas() const override;
virtual void RenderCells(const vtkm::cont::DynamicCellSet& cellset,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::cont::Field& scalarField,
const vtkm::rendering::ColorTable&, //colorTable
const vtkm::rendering::Camera& camera,
const vtkm::Range& scalarRange) VTKM_OVERRIDE;
const vtkm::Range& scalarRange) override;
virtual void StartScene() VTKM_OVERRIDE;
virtual void EndScene() VTKM_OVERRIDE;
virtual void StartScene() override;
virtual void EndScene() override;
vtkm::rendering::Mapper* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Mapper* NewCopy() const override;
void CreateDefaultView();
protected:

12
vtkm/rendering/MapperGL.h
View File

@ -45,14 +45,14 @@ public:
const vtkm::cont::Field& scalarField,
const vtkm::rendering::ColorTable& colorTable,
const vtkm::rendering::Camera&,
const vtkm::Range& scalarRange) VTKM_OVERRIDE;
const vtkm::Range& scalarRange) override;
void StartScene() VTKM_OVERRIDE;
void EndScene() VTKM_OVERRIDE;
void SetCanvas(vtkm::rendering::Canvas* canvas) VTKM_OVERRIDE;
virtual vtkm::rendering::Canvas* GetCanvas() const VTKM_OVERRIDE;
void StartScene() override;
void EndScene() override;
void SetCanvas(vtkm::rendering::Canvas* canvas) override;
virtual vtkm::rendering::Canvas* GetCanvas() const override;
vtkm::rendering::Mapper* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Mapper* NewCopy() const override;
vtkm::rendering::CanvasGL* Canvas;
GLuint shader_programme;

12
vtkm/rendering/MapperRayTracer.h
View File

@ -38,20 +38,20 @@ public:
~MapperRayTracer();
void SetCanvas(vtkm::rendering::Canvas* canvas) VTKM_OVERRIDE;
virtual vtkm::rendering::Canvas* GetCanvas() const VTKM_OVERRIDE;
void SetCanvas(vtkm::rendering::Canvas* canvas) override;
virtual vtkm::rendering::Canvas* GetCanvas() const override;
void RenderCells(const vtkm::cont::DynamicCellSet& cellset,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::cont::Field& scalarField,
const vtkm::rendering::ColorTable& colorTable,
const vtkm::rendering::Camera& camera,
const vtkm::Range& scalarRange) VTKM_OVERRIDE;
const vtkm::Range& scalarRange) override;
virtual void StartScene() VTKM_OVERRIDE;
virtual void EndScene() VTKM_OVERRIDE;
virtual void StartScene() override;
virtual void EndScene() override;
void SetCompositeBackground(bool on);
vtkm::rendering::Mapper* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Mapper* NewCopy() const override;
private:
struct InternalsType;

12
vtkm/rendering/MapperVolume.h
View File

@ -36,20 +36,20 @@ public:
~MapperVolume();
void SetCanvas(vtkm::rendering::Canvas* canvas) VTKM_OVERRIDE;
virtual vtkm::rendering::Canvas* GetCanvas() const VTKM_OVERRIDE;
void SetCanvas(vtkm::rendering::Canvas* canvas) override;
virtual vtkm::rendering::Canvas* GetCanvas() const override;
virtual void RenderCells(const vtkm::cont::DynamicCellSet& cellset,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::cont::Field& scalarField,
const vtkm::rendering::ColorTable&, //colorTable
const vtkm::rendering::Camera& camera,
const vtkm::Range& scalarRange) VTKM_OVERRIDE;
const vtkm::Range& scalarRange) override;
virtual void StartScene() VTKM_OVERRIDE;
virtual void EndScene() VTKM_OVERRIDE;
virtual void StartScene() override;
virtual void EndScene() override;
vtkm::rendering::Mapper* NewCopy() const VTKM_OVERRIDE;
vtkm::rendering::Mapper* NewCopy() const override;
void SetSampleDistance(const vtkm::Float32 distance);
void SetCompositeBackground(const bool compositeBackground);

12
vtkm/rendering/MapperWireframer.h
View File

@ -42,8 +42,8 @@ public:
MapperWireframer();
virtual ~MapperWireframer();
virtual vtkm::rendering::Canvas* GetCanvas() const VTKM_OVERRIDE;
virtual void SetCanvas(vtkm::rendering::Canvas* canvas) VTKM_OVERRIDE;
virtual vtkm::rendering::Canvas* GetCanvas() const override;
virtual void SetCanvas(vtkm::rendering::Canvas* canvas) override;
bool GetShowInternalZones() const;
void SetShowInternalZones(bool showInternalZones);
@ -51,17 +51,17 @@ public:
bool GetIsOverlay() const;
void SetIsOverlay(bool isOverlay);
virtual void StartScene() VTKM_OVERRIDE;
virtual void EndScene() VTKM_OVERRIDE;
virtual void StartScene() override;
virtual void EndScene() override;
virtual void RenderCells(const vtkm::cont::DynamicCellSet& cellset,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::cont::Field& scalarField,
const vtkm::rendering::ColorTable& colorTable,
const vtkm::rendering::Camera& camera,
const vtkm::Range& scalarRange) VTKM_OVERRIDE;
const vtkm::Range& scalarRange) override;
virtual vtkm::rendering::Mapper* NewCopy() const VTKM_OVERRIDE;
virtual vtkm::rendering::Mapper* NewCopy() const override;
private:
struct InternalsType;

2
vtkm/rendering/TextAnnotationBillboard.h
View File

@ -48,7 +48,7 @@ public:
void Render(const vtkm::rendering::Camera& camera,
const vtkm::rendering::WorldAnnotator& worldAnnotator,
vtkm::rendering::Canvas& canvas) const VTKM_OVERRIDE;
vtkm::rendering::Canvas& canvas) const override;
};
}
} // namespace vtkm::rendering

2
vtkm/rendering/TextAnnotationScreen.h
View File

@ -48,7 +48,7 @@ public:
void Render(const vtkm::rendering::Camera& camera,
const vtkm::rendering::WorldAnnotator& annotator,
vtkm::rendering::Canvas& canvas) const VTKM_OVERRIDE;
vtkm::rendering::Canvas& canvas) const override;
};
}
} // namespace vtkm::rendering

6
vtkm/rendering/View1D.h
View File

@ -45,9 +45,9 @@ public:
~View1D();
void Paint() VTKM_OVERRIDE;
void RenderScreenAnnotations() VTKM_OVERRIDE;
void RenderWorldAnnotations() VTKM_OVERRIDE;
void Paint() override;
void RenderScreenAnnotations() override;
void RenderWorldAnnotations() override;
void RenderColorLegendAnnotations();
void EnableLegend();

6
vtkm/rendering/View2D.h
View File

@ -46,11 +46,11 @@ public:
~View2D();
void Paint() VTKM_OVERRIDE;
void Paint() override;
void RenderScreenAnnotations() VTKM_OVERRIDE;
void RenderScreenAnnotations() override;
void RenderWorldAnnotations() VTKM_OVERRIDE;
void RenderWorldAnnotations() override;
private:
void UpdateCameraProperties();

6
vtkm/rendering/View3D.h
View File

@ -47,11 +47,11 @@ public:
~View3D();
void Paint() VTKM_OVERRIDE;
void Paint() override;
void RenderScreenAnnotations() VTKM_OVERRIDE;
void RenderScreenAnnotations() override;
void RenderWorldAnnotations() VTKM_OVERRIDE;
void RenderWorldAnnotations() override;
private:
// 3D-specific annotations

4
vtkm/rendering/WorldAnnotatorGL.h
View File

@ -42,7 +42,7 @@ public:
const vtkm::Vec<vtkm::Float64, 3>& point1,
vtkm::Float32 lineWidth,
const vtkm::rendering::Color& color,
bool inFront) const VTKM_OVERRIDE;
bool inFront) const override;
void AddText(const vtkm::Vec<vtkm::Float32, 3>& origin,
const vtkm::Vec<vtkm::Float32, 3>& right,
@ -50,7 +50,7 @@ public:
vtkm::Float32 scale,
const vtkm::Vec<vtkm::Float32, 2>& anchor,
const vtkm::rendering::Color& color,
const std::string& text) const VTKM_OVERRIDE;
const std::string& text) const override;
private:
BitmapFont Font;

2
vtkm/rendering/internal/RunTriangulator.cxx
View File

@ -59,7 +59,7 @@ struct TriangulatorFunctor
void RunTriangulator(const vtkm::cont::DynamicCellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Id, 4>>& indices,
vtkm::Id& numberOfTriangles,
const vtkm::cont::RuntimeDeviceTracker& tracker)
vtkm::cont::RuntimeDeviceTracker tracker)
{
// TODO: Should the rendering library support policies or some other way to
// configure with custom devices?

2
vtkm/rendering/internal/RunTriangulator.h
View File

@ -43,7 +43,7 @@ void RunTriangulator(
const vtkm::cont::DynamicCellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Id, 4>>& indices,
vtkm::Id& numberOfTriangles,
const vtkm::cont::RuntimeDeviceTracker& tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker());
vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker());
}
}
} // namespace vtkm::rendering::internal

3
vtkm/rendering/raytracing/Camera.cxx
View File

@ -741,6 +741,9 @@ VTKM_CONT void Camera::CreateRaysOnDevice(Ray<Precision>& rays,
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(0.f))
.Invoke(rays.MinDistance);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(0.f))
.Invoke(rays.Distance);
//Reset the Rays Hit Index to -2
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Id>, Device>(MemSet<vtkm::Id>(-2))
.Invoke(rays.HitIdx);

14
vtkm/rendering/raytracing/ChannelBuffer.cxx
View File

@ -100,10 +100,10 @@ ChannelBuffer<Precision>::ChannelBuffer()
template <typename Precision>
ChannelBuffer<Precision>::ChannelBuffer(const vtkm::Int32 numChannels, const vtkm::Id size)
{
if (size < 1)
throw vtkm::cont::ErrorBadValue("ChannelBuffer: Size must be greater that 0");
if (numChannels < 1)
throw vtkm::cont::ErrorBadValue("ChannelBuffer: NumChannels must be greater that 0");
if (size < 0)
throw vtkm::cont::ErrorBadValue("ChannelBuffer: Size must be greater that -1");
if (numChannels < 0)
throw vtkm::cont::ErrorBadValue("ChannelBuffer: NumChannels must be greater that -1");
this->NumChannels = numChannels;
this->Size = size;
@ -235,7 +235,10 @@ ChannelBuffer<Precision> ChannelBuffer<Precision>::GetChannel(const vtkm::Int32
throw vtkm::cont::ErrorBadValue("ChannelBuffer: invalid channel to extract");
ChannelBuffer<Precision> output(1, this->Size);
output.SetName(this->Name);
if (this->Size == 0)
{
return output;
}
ExtractChannelFunctor<Precision> functor(this, output.Buffer, channel);
vtkm::cont::TryExecute(functor);
@ -373,7 +376,6 @@ public:
{
InvDeltaScalar = 1.f / (maxScalar - minScalar);
}
//std::cout<<"Min scalar "<<minScalar<<" max "<<maxScalar<<std::endl;
}
typedef void ControlSignature(FieldInOut<>);
typedef void ExecutionSignature(_1);

85
vtkm/rendering/raytracing/ConnectivityTracer.h
View File

@ -62,8 +62,6 @@ namespace rendering
{
namespace raytracing
{
//
// Advance Ray
// After a ray leaves the mesh, we need to check to see
@ -458,6 +456,40 @@ public:
} //operator
}; //class RayBumper
template <typename FloatType>
class AddPathLengths : public vtkm::worklet::WorkletMapField
{
public:
VTKM_CONT
AddPathLengths() {}
typedef void ControlSignature(FieldIn<RayStatusType>, // ray status
FieldIn<ScalarRenderingTypes>, // cell enter distance
FieldIn<ScalarRenderingTypes>, // cell exit distance
FieldInOut<ScalarRenderingTypes>); // ray absorption data
typedef void ExecutionSignature(_1, _2, _3, _4);
VTKM_EXEC inline void operator()(const vtkm::UInt8& rayStatus,
const FloatType& enterDistance,
const FloatType& exitDistance,
FloatType& distance) const
{
if (rayStatus != RAY_ACTIVE)
{
return;
}
if (exitDistance <= enterDistance)
{
return;
}
FloatType segmentLength = exitDistance - enterDistance;
distance += segmentLength;
}
};
template <typename FloatType>
class Integrate : public vtkm::worklet::WorkletMapField
{
@ -1138,6 +1170,17 @@ public:
this->IntersectTime += timer.GetElapsedTime();
}
template <typename FloatType, typename Device>
void AccumulatePathLengths(Ray<FloatType>& rays, detail::RayTracking<FloatType>& tracker, Device)
{
vtkm::worklet::DispatcherMapField<AddPathLengths<FloatType>, Device>(
AddPathLengths<FloatType>())
.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.GetBuffer("path_lengths").Buffer);
}
template <typename FloatType, typename Device>
void FindLostRays(Ray<FloatType>& rays, detail::RayTracking<FloatType>& tracker, Device)
{
@ -1210,7 +1253,6 @@ public:
bool divideEmisByAbsorp = false;
vtkm::cont::ArrayHandle<FloatType> absorp = rays.Buffers.at(0).Buffer;
vtkm::cont::ArrayHandle<FloatType> emission = rays.GetBuffer("emission").Buffer;
vtkm::worklet::DispatcherMapField<IntegrateEmission<FloatType>, Device>(
IntegrateEmission<FloatType>(rays.Buffers.at(0).GetNumChannels(), divideEmisByAbsorp))
.Invoke(rays.Status,
@ -1242,13 +1284,27 @@ public:
template <typename FloatType>
void PrintDebugRay(Ray<FloatType>& rays, vtkm::Id rayId)
{
if (rayId < 0 || rayId > rays.Distance.GetPortalControl().GetNumberOfValues())
vtkm::Id index = -1;
for (vtkm::Id i = 0; i < rays.NumRays; ++i)
{
if (rays.PixelIdx.GetPortalControl().Get(i) == rayId)
{
index = i;
break;
}
}
if (index == -1)
{
return;
std::cout << "++++++++DEBUG RAY++++++++\n";
std::cout << "Status: " << (int)rays.Status.GetPortalControl().Get(rayId) << "\n";
std::cout << "HitIndex: " << rays.HitIdx.GetPortalControl().Get(rayId) << "\n";
std::cout << "Dist " << rays.Distance.GetPortalControl().Get(rayId) << "\n";
std::cout << "MinDist " << rays.MinDistance.GetPortalControl().Get(rayId) << "\n";
}
std::cout << "++++++++RAY " << rayId << "++++++++\n";
std::cout << "Status: " << (int)rays.Status.GetPortalControl().Get(index) << "\n";
std::cout << "HitIndex: " << rays.HitIdx.GetPortalControl().Get(index) << "\n";
std::cout << "Dist " << rays.Distance.GetPortalControl().Get(index) << "\n";
std::cout << "MinDist " << rays.MinDistance.GetPortalControl().Get(index) << "\n";
std::cout << "Origin " << rays.Origin.GetPortalConstControl().Get(index) << "\n";
std::cout << "Dir " << rays.Dir.GetPortalConstControl().Get(index) << "\n";
std::cout << "+++++++++++++++++++++++++\n";
}
@ -1256,7 +1312,7 @@ public:
void OffsetMinDistances(Ray<FloatType>& rays, Device)
{
vtkm::worklet::DispatcherMapField<AdvanceRay<FloatType>, Device>(
AdvanceRay<FloatType>(FloatType(0.0000001)))
AdvanceRay<FloatType>(FloatType(0.001)))
.Invoke(rays.Status, rays.MinDistance);
}
@ -1303,6 +1359,7 @@ public:
this->SetBoundingBox(Device());
bool hasPathLengths = rays.HasBuffer("path_lengths");
vtkm::cont::Timer<Device> timer;
this->Init(Device());
@ -1319,9 +1376,8 @@ public:
MeshConn.Construct(Device());
logger->AddLogData("active_pixels", rays.NumRays);
bool cullMissedRays = false;
bool cullMissedRays = true;
bool workRemaining = true;
if (this->CountRayStatus)
{
@ -1377,6 +1433,10 @@ public:
else
this->IntegrateCells(rays, rayTracker, Device());
if (hasPathLengths)
{
this->AccumulatePathLengths(rays, rayTracker, Device());
}
//swap enter and exit distances
rayTracker.Swap();
if (this->CountRayStatus)
@ -1405,6 +1465,7 @@ public:
}
vtkm::Float64 renderTime = renderTimer.GetElapsedTime();
this->LogTimers();
logger->AddLogData("active_pixels", rays.NumRays);
logger->CloseLogEntry(renderTime);
} //Render

6
vtkm/rendering/raytracing/MortonCodes.h
View File

@ -237,9 +237,9 @@ public:
count++;
}
//TODO: we could make this a recipical, but this is not a bottleneck.
center[0] = center[0] / vtkm::Float32(count + 1);
center[1] = center[1] / vtkm::Float32(count + 1);
center[2] = center[2] / vtkm::Float32(count + 1);
center[0] = center[0] / vtkm::Float32(count);
center[1] = center[1] / vtkm::Float32(count);
center[2] = center[2] / vtkm::Float32(count);
Normalize(center);
code = Morton3D(center[0], center[1], center[2]);
BOUNDS_CHECK(mortonCodes, offset + i);

36
vtkm/rendering/raytracing/Ray.h
View File

@ -181,9 +181,28 @@ public:
VTKM_CONT Ray(const vtkm::Int32 size, Device, bool enableIntersectionData = false)
{
NumRays = size;
IntersectionDataEnabled = enableIntersectionData;
ChannelBuffer<Precision> buffer;
this->Buffers.push_back(buffer);
DebugWidth = -1;
DebugHeight = -1;
this->Resize(size, Device());
}
VTKM_CONT void Resize(const vtkm::Int32 size)
{
this->Resize(size, vtkm::cont::DeviceAdapterTagSerial());
}
template <typename Device>
VTKM_CONT void Resize(const vtkm::Int32 size, Device)
{
NumRays = size;
if (IntersectionDataEnabled)
{
IntersectionX.PrepareForOutput(NumRays, Device());
@ -235,12 +254,11 @@ public:
DirX, inComp[0], DirY, inComp[1], DirZ, inComp[2]);
ChannelBuffer<Precision> buffer;
buffer.Resize(NumRays, Device());
this->Buffers.push_back(buffer);
DebugWidth = -1;
DebugHeight = -1;
const size_t numBuffers = this->Buffers.size();
for (size_t i = 0; i < numBuffers; ++i)
{
this->Buffers[i].Resize(NumRays, Device());
}
}
VTKM_CONT
@ -255,9 +273,9 @@ public:
VTKM_CONT
bool HasBuffer(const std::string name)
{
vtkm::Id numBuffers = static_cast<vtkm::Id>(this->Buffers.size());
size_t numBuffers = this->Buffers.size();
bool found = false;
for (vtkm::Id i = 0; i < numBuffers; ++i)
for (size_t i = 0; i < numBuffers; ++i)
{
if (this->Buffers[i].GetName() == name)
{

74
vtkm/testing/UnitTestListTag.cxx
View File

@ -33,10 +33,6 @@ namespace
template <int N>
struct TestClass
{
enum
{
NUMBER = N
};
};
struct TestListTag1 : vtkm::ListTagBase<TestClass<11>>
@ -63,36 +59,50 @@ struct TestListTagIntersect : vtkm::ListTagIntersect<TestListTag3, TestListTagJo
{
};
struct TestListTagCrossProduct : vtkm::ListCrossProduct<TestListTag3, TestListTag1>
{
};
struct TestListTagUniversal : vtkm::ListTagUniversal
{
};
template <int N, int M>
std::pair<int, int> test_number(std::pair<TestClass<N>, TestClass<M>>)
{
return std::make_pair(N, M);
}
template <int N>
int test_number(TestClass<N>)
{
return N;
}
template <typename T>
struct MutableFunctor
{
std::vector<int> FoundTypes;
std::vector<T> FoundTypes;
template <typename T>
VTKM_CONT void operator()(T)
template <typename U>
VTKM_CONT void operator()(U u)
{
this->FoundTypes.push_back(T::NUMBER);
this->FoundTypes.push_back(test_number(u));
}
};
std::vector<int> g_FoundType;
template <typename T>
struct ConstantFunctor
{
ConstantFunctor() { g_FoundType.erase(g_FoundType.begin(), g_FoundType.end()); }
template <typename T>
VTKM_CONT void operator()(T) const
template <typename U, typename VectorType>
VTKM_CONT void operator()(U u, VectorType& vector) const
{
g_FoundType.push_back(T::NUMBER);
vector.push_back(test_number(u));
}
};
template <vtkm::IdComponent N>
void CheckSame(const vtkm::Vec<int, N>& expected, const std::vector<int>& found)
template <typename T, vtkm::IdComponent N>
void CheckSame(const vtkm::Vec<T, N>& expected, const std::vector<T>& found)
{
VTKM_TEST_ASSERT(static_cast<int>(found.size()) == N, "Got wrong number of items.");
@ -137,13 +147,15 @@ void TryList(const vtkm::Vec<int, N>& expected, ListTag)
VTKM_IS_LIST_TAG(ListTag);
std::cout << " Try mutable for each" << std::endl;
MutableFunctor functor;
MutableFunctor<int> functor;
vtkm::ListForEach(functor, ListTag());
CheckSame(expected, functor.FoundTypes);
std::cout << " Try constant for each" << std::endl;
vtkm::ListForEach(ConstantFunctor(), ListTag());
CheckSame(expected, g_FoundType);
std::vector<int> foundTypes;
ConstantFunctor<int> cfunc;
vtkm::ListForEach(cfunc, ListTag(), foundTypes);
CheckSame(expected, foundTypes);
std::cout << " Try checking contents" << std::endl;
CheckContains(TestClass<11>(), ListTag(), functor.FoundTypes);
@ -157,6 +169,22 @@ void TryList(const vtkm::Vec<int, N>& expected, ListTag)
CheckContains(TestClass<43>(), ListTag(), functor.FoundTypes);
CheckContains(TestClass<44>(), ListTag(), functor.FoundTypes);
}
template <vtkm::IdComponent N, typename ListTag>
void TryList(const vtkm::Vec<std::pair<int, int>, N>& expected, ListTag)
{
VTKM_IS_LIST_TAG(ListTag);
std::cout << " Try mutable for each" << std::endl;
MutableFunctor<std::pair<int, int>> functor;
vtkm::ListForEach(functor, ListTag());
CheckSame(expected, functor.FoundTypes);
std::cout << " Try constant for each" << std::endl;
std::vector<std::pair<int, int>> foundTypes;
ConstantFunctor<std::pair<int, int>> cfunc;
vtkm::ListForEach(cfunc, ListTag(), foundTypes);
CheckSame(expected, foundTypes);
}
template <vtkm::IdComponent N>
void TryList(const vtkm::Vec<int, N>&, TestListTagUniversal tag)
@ -207,6 +235,12 @@ void TestLists()
std::cout << "ListTagIntersect" << std::endl;
TryList(vtkm::Vec<int, 3>(31, 32, 33), TestListTagIntersect());
std::cout << "ListTagCrossProduct" << std::endl;
TryList(vtkm::Vec<std::pair<int, int>, 3>({ 31, 11 }, { 32, 11 }, { 33, 11 }),
TestListTagCrossProduct());
std::cout << "ListTagUniversal" << std::endl;
TryList(vtkm::Vec<int, 4>(1, 2, 3, 4), TestListTagUniversal());
}

12
vtkm/worklet/Clip.h
View File

@ -30,7 +30,7 @@
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/DynamicCellSet.h>
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
#include <vtkm/cont/Timer.h>
#include <vtkm/exec/ExecutionWholeArray.h>
@ -522,7 +522,7 @@ public:
VTKM_CONT
ClipWithImplicitFunction(Clip* clipper,
const DynamicCellSet& cellSet,
const vtkm::exec::ImplicitFunction& function,
const vtkm::ImplicitFunction* function,
vtkm::cont::CellSetExplicit<>* result)
: Clipper(clipper)
, CellSet(&cellSet)
@ -536,8 +536,8 @@ public:
{
// Evaluate the implicit function on the input coordinates using
// ArrayHandleTransform
vtkm::cont::ArrayHandleTransform<ArrayHandleType, vtkm::exec::ImplicitFunctionValue>
clipScalars(handle, this->Function);
vtkm::cont::ArrayHandleTransform<ArrayHandleType, vtkm::ImplicitFunctionValue> clipScalars(
handle, this->Function);
// Clip at locations where the implicit function evaluates to 0
*this->Result = this->Clipper->Run(*this->CellSet, clipScalars, 0.0, DeviceAdapter());
@ -546,13 +546,13 @@ public:
private:
Clip* Clipper;
const DynamicCellSet* CellSet;
const vtkm::exec::ImplicitFunctionValue Function;
vtkm::ImplicitFunctionValue Function;
vtkm::cont::CellSetExplicit<>* Result;
};
template <typename CellSetList, typename DeviceAdapter>
vtkm::cont::CellSetExplicit<> Run(const vtkm::cont::DynamicCellSetBase<CellSetList>& cellSet,
const vtkm::cont::ImplicitFunction& clipFunction,
const vtkm::cont::ImplicitFunctionHandle& clipFunction,
const vtkm::cont::CoordinateSystem& coords,
DeviceAdapter device)
{

15
vtkm/worklet/ExtractGeometry.h
View File

@ -27,7 +27,7 @@
#include <vtkm/cont/CellSetPermutation.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
namespace vtkm
{
@ -58,7 +58,7 @@ public:
}
VTKM_CONT
ExtractCellsByVOI(const vtkm::exec::ImplicitFunction& function,
ExtractCellsByVOI(const vtkm::ImplicitFunction* function,
bool extractInside,
bool extractBoundaryCells,
bool extractOnlyBoundaryCells)
@ -77,11 +77,12 @@ public:
// Count points inside/outside volume of interest
vtkm::IdComponent inCnt = 0;
vtkm::IdComponent outCnt = 0;
for (vtkm::IdComponent indx = 0; indx < numIndices; indx++)
vtkm::Id indx;
for (indx = 0; indx < numIndices; indx++)
{
vtkm::Id ptId = connectivityIn[indx];
vtkm::Id ptId = connectivityIn[static_cast<vtkm::IdComponent>(indx)];
vtkm::Vec<FloatDefault, 3> coordinate = coordinates.Get(ptId);
vtkm::FloatDefault value = this->Function.Value(coordinate);
vtkm::FloatDefault value = this->Function->Value(coordinate);
if (value <= 0)
inCnt++;
if (value >= 0)
@ -106,7 +107,7 @@ public:
}
private:
vtkm::exec::ImplicitFunction Function;
const vtkm::ImplicitFunction* Function;
bool ExtractInside;
bool ExtractBoundaryCells;
bool ExtractOnlyBoundaryCells;
@ -155,7 +156,7 @@ public:
vtkm::cont::CellSetPermutation<CellSetType> Run(
const CellSetType& cellSet,
const vtkm::cont::CoordinateSystem& coordinates,
const vtkm::cont::ImplicitFunction& implicitFunction,
const vtkm::cont::ImplicitFunctionHandle& implicitFunction,
bool extractInside,
bool extractBoundaryCells,
bool extractOnlyBoundaryCells,

10
vtkm/worklet/ExtractPoints.h
View File

@ -26,7 +26,7 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/ImplicitFunction.h>
#include <vtkm/cont/ImplicitFunctionHandle.h>
namespace vtkm
{
@ -51,7 +51,7 @@ public:
typedef _3 ExecutionSignature(_2);
VTKM_CONT
ExtractPointsByVOI(const vtkm::exec::ImplicitFunction& function, bool extractInside)
ExtractPointsByVOI(const vtkm::ImplicitFunction* function, bool extractInside)
: Function(function)
, passValue(extractInside)
, failValue(!extractInside)
@ -62,14 +62,14 @@ public:
bool operator()(const vtkm::Vec<vtkm::Float64, 3>& coordinate) const
{
bool pass = passValue;
vtkm::Float64 value = this->Function.Value(coordinate);
vtkm::Float64 value = this->Function->Value(coordinate);
if (value > 0)
pass = failValue;
return pass;
}
private:
vtkm::exec::ImplicitFunction Function;
const vtkm::ImplicitFunction* Function;
bool passValue;
bool failValue;
};
@ -97,7 +97,7 @@ public:
template <typename CellSetType, typename CoordinateType, typename DeviceAdapter>
vtkm::cont::CellSetSingleType<> Run(const CellSetType& cellSet,
const CoordinateType& coordinates,
const vtkm::cont::ImplicitFunction& implicitFunction,
const vtkm::cont::ImplicitFunctionHandle& implicitFunction,
bool extractInside,
DeviceAdapter device)
{

4
vtkm/worklet/MarchingCubes.h
View File

@ -686,11 +686,11 @@ struct GenerateNormalsDeduced
//
NormalsWorkletPass1 pass1(*edges);
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass1>(pass1).Invoke(
*cellset, *cellset, coordinates, *field, *normals);
*cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *normals);
NormalsWorkletPass2 pass2(*edges);
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass2>(pass2).Invoke(
*cellset, *cellset, coordinates, *field, *weights, *normals);
*cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *weights, *normals);
}
};

28
vtkm/worklet/VertexClustering.h
View File

@ -78,28 +78,20 @@ struct SelectRepresentativePoint : public vtkm::worklet::WorkletReduceByKey
return pointsIn[pointsIn.GetNumberOfComponents() / 2];
}
template <typename KeyType, typename DeviceAdapterTag>
struct RunTrampoline
{
const vtkm::worklet::Keys<KeyType>& Keys;
vtkm::cont::DynamicArrayHandle& OutputPoints;
VTKM_CONT
RunTrampoline(const vtkm::worklet::Keys<KeyType>& keys, vtkm::cont::DynamicArrayHandle& output)
: Keys(keys)
, OutputPoints(output)
template <typename InputPointsArrayType, typename KeyType, typename DeviceAdapterTag>
VTKM_CONT void operator()(const InputPointsArrayType& points,
const vtkm::worklet::Keys<KeyType>& keys,
vtkm::cont::DynamicArrayHandle& output,
DeviceAdapterTag) const
{
}
template <typename InputPointsArrayType>
VTKM_CONT void operator()(const InputPointsArrayType& points) const
{
vtkm::cont::ArrayHandle<typename InputPointsArrayType::ValueType> out;
vtkm::worklet::DispatcherReduceByKey<SelectRepresentativePoint, DeviceAdapterTag> dispatcher;
dispatcher.Invoke(this->Keys, points, out);
dispatcher.Invoke(keys, points, out);
this->OutputPoints = out;
output = out;
}
};
@ -107,11 +99,11 @@ struct SelectRepresentativePoint : public vtkm::worklet::WorkletReduceByKey
VTKM_CONT static vtkm::cont::DynamicArrayHandle Run(
const vtkm::worklet::Keys<KeyType>& keys,
const InputDynamicPointsArrayType& inputPoints,
DeviceAdapterTag)
DeviceAdapterTag tag)
{
vtkm::cont::DynamicArrayHandle output;
RunTrampoline<KeyType, DeviceAdapterTag> trampoline(keys, output);
vtkm::cont::CastAndCall(inputPoints, trampoline);
RunTrampoline trampoline;
vtkm::cont::CastAndCall(inputPoints, trampoline, keys, output, tag);
return output;
}
};

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