diff --git a/version.txt b/version.txt index 3eefcb9dd..9084fa2f7 100644 --- a/version.txt +++ b/version.txt @@ -1 +1 @@ -1.0.0 +1.1.0 diff --git a/vtkm/CMakeLists.txt b/vtkm/CMakeLists.txt index f053f6c3d..443db458f 100644 --- a/vtkm/CMakeLists.txt +++ b/vtkm/CMakeLists.txt @@ -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 ) diff --git a/vtkm/ImplicitFunction.h b/vtkm/ImplicitFunction.h new file mode 100644 index 000000000..5553933d9 --- /dev/null +++ b/vtkm/ImplicitFunction.h @@ -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 +#include +#include + +namespace vtkm +{ + +//============================================================================ +class VTKM_ALWAYS_EXPORT ImplicitFunction : public vtkm::VirtualObjectBase +{ +public: + using Scalar = vtkm::FloatDefault; + using Vector = vtkm::Vec; + + 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 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 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(); + 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(); + 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_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 diff --git a/vtkm/ListTag.h b/vtkm/ListTag.h index cdcfa3f8e..afcd6a5ac 100644 --- a/vtkm/ListTag.h +++ b/vtkm/ListTag.h @@ -73,7 +73,7 @@ struct ListTagJoin : detail::ListRoot using list = typename detail::ListJoin::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 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 -VTKM_CONT void ListForEach(Functor&& f, ListTag) +template +VTKM_CONT void ListForEach(Functor&& f, ListTag, Args&&... args) { VTKM_IS_LIST_TAG(ListTag); - detail::ListForEachImpl(f, typename ListTag::list()); + detail::ListForEachImpl( + std::forward(f), typename ListTag::list{}, std::forward(args)...); } +/// Generate a tag that is the cross product of two other tags. The resulting +// a tag has the form of Tag< std::pair, std::pair .... > +/// +template +struct ListCrossProduct : detail::ListRoot +{ + using list = + typename detail::ListCrossProductImpl::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. diff --git a/vtkm/VirtualObjectBase.h b/vtkm/VirtualObjectBase.h new file mode 100644 index 000000000..e1e94fe72 --- /dev/null +++ b/vtkm/VirtualObjectBase.h @@ -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 + +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 diff --git a/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx b/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx index 9de006251..163eaf8db 100644 --- a/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx +++ b/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx @@ -24,7 +24,7 @@ #include #include #include -#include +#include #include #include @@ -264,7 +264,7 @@ public: typedef void ControlSignature(FieldIn, FieldOut); typedef void ExecutionSignature(_1, _2); - EvaluateImplicitFunction(const ImplicitFunction& function) + EvaluateImplicitFunction(const ImplicitFunction* function) : Function(function) { } @@ -272,11 +272,11 @@ public: template 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 @@ -286,7 +286,7 @@ public: typedef void ControlSignature(FieldIn, FieldOut); 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 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 @@ -680,7 +680,7 @@ private: { vtkm::cont::ArrayHandle> Points; vtkm::cont::ArrayHandle 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; + using EvalWorklet = EvaluateImplicitFunction; using EvalDispatcher = vtkm::worklet::DispatcherMapField; - EvalWorklet eval(Internal.Sphere1); + auto handle = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1); + auto function = + static_cast(handle.PrepareForExecution(DeviceAdapterTag())); + EvalWorklet eval(function); vtkm::cont::Timer 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 - struct BenchDynamicImplicitFunction + struct BenchVirtualImplicitFunction { - BenchDynamicImplicitFunction() + BenchVirtualImplicitFunction() : Internal(MakeImplicitFunctionBenchData()) { } @@ -753,10 +756,11 @@ private: VTKM_CONT vtkm::Float64 operator()() { - using EvalWorklet = EvaluateImplicitFunction; + using EvalWorklet = EvaluateImplicitFunction; using EvalDispatcher = vtkm::worklet::DispatcherMapField; - EvalWorklet eval(Internal.Sphere1.PrepareForExecution(DeviceAdapterTag())); + auto sphere = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1); + EvalWorklet eval(sphere.PrepareForExecution(DeviceAdapterTag())); vtkm::cont::Timer 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; + using EvalWorklet = Evaluate2ImplicitFunctions; using EvalDispatcher = vtkm::worklet::DispatcherMapField; - 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(h1.PrepareForExecution(DeviceAdapterTag())); + auto f2 = static_cast(h2.PrepareForExecution(DeviceAdapterTag())); + EvalWorklet eval(f1, f2); vtkm::cont::Timer 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 - struct Bench2DynamicImplicitFunctions + struct Bench2VirtualImplicitFunctions { - Bench2DynamicImplicitFunctions() + Bench2VirtualImplicitFunctions() : Internal(MakeImplicitFunctionBenchData()) { } @@ -820,11 +828,13 @@ private: vtkm::Float64 operator()() { using EvalWorklet = - Evaluate2ImplicitFunctions; + Evaluate2ImplicitFunctions; using EvalDispatcher = vtkm::worklet::DispatcherMapField; - 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 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; diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index fcfa7be94..76cb68e74 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -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 ) #----------------------------------------------------------------------------- diff --git a/vtkm/cont/CellSetExplicit.h b/vtkm/cont/CellSetExplicit.h index 47dbed10d..d302f26f5 100644 --- a/vtkm/cont/CellSetExplicit.h +++ b/vtkm/cont/CellSetExplicit.h @@ -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; diff --git a/vtkm/cont/CoordinateSystem.h b/vtkm/cont/CoordinateSystem.h index 28f8a4ee0..1504a83bc 100644 --- a/vtkm/cont/CoordinateSystem.h +++ b/vtkm/cont/CoordinateSystem.h @@ -72,7 +72,17 @@ struct StorageListTagCoordinateSystemDefault vtkm::cont::ArrayHandle>::StorageTag> { }; +} +} +namespace vtkm +{ + +template struct ListCrossProduct<::vtkm::TypeListTagFieldVec3, + ::vtkm::cont::StorageListTagCoordinateSystemDefault>; + +namespace cont +{ using DynamicArrayHandleCoordinateSystem = vtkm::cont::DynamicArrayHandleBase; @@ -214,10 +224,10 @@ public: virtual void PrintSummary(std::ostream& out) const; }; -template -void CastAndCall(const vtkm::cont::CoordinateSystem& coords, const Functor& f) +template +void CastAndCall(const vtkm::cont::CoordinateSystem& coords, const Functor& f, Args&&... args) { - coords.GetData().CastAndCall(f); + coords.GetData().CastAndCall(f, std::forward(args)...); } namespace internal diff --git a/vtkm/cont/DeviceAdapterListTag.h b/vtkm/cont/DeviceAdapterListTag.h index a41d57c24..d38bb09ca 100644 --- a/vtkm/cont/DeviceAdapterListTag.h +++ b/vtkm/cont/DeviceAdapterListTag.h @@ -40,6 +40,47 @@ struct DeviceAdapterListTagCommon : vtkm::ListTagBase { }; + +namespace detail +{ + +template +class ExecuteIfValidDeviceTag +{ +private: + template + using EnableIfValid = std::enable_if::Valid>; + + template + using EnableIfInvalid = std::enable_if::Valid>; + +public: + explicit ExecuteIfValidDeviceTag(const FunctorType& functor) + : Functor(functor) + { + } + + template + typename EnableIfValid::type operator()(DeviceAdapter) const + { + this->Functor(DeviceAdapter()); + } + + template + typename EnableIfInvalid::type operator()(DeviceAdapter) const + { + } + +private: + FunctorType Functor; +}; +} // detail + +template +VTKM_CONT void ForEachValidDevice(DeviceList devices, const Functor& functor) +{ + vtkm::ListForEach(detail::ExecuteIfValidDeviceTag(functor), devices); +} } } // namespace vtkm::cont diff --git a/vtkm/cont/DynamicArrayHandle.cxx b/vtkm/cont/DynamicArrayHandle.cxx index e53ad69ff..0d7ce827d 100644 --- a/vtkm/cont/DynamicArrayHandle.cxx +++ b/vtkm/cont/DynamicArrayHandle.cxx @@ -18,6 +18,8 @@ // this software. //============================================================================ +#include +#include #include 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 diff --git a/vtkm/cont/DynamicArrayHandle.h b/vtkm/cont/DynamicArrayHandle.h index d9a29a08d..bc600e648 100644 --- a/vtkm/cont/DynamicArrayHandle.h +++ b/vtkm/cont/DynamicArrayHandle.h @@ -31,10 +31,11 @@ #include -#include - namespace vtkm { + +template struct ListCrossProduct; + namespace cont { @@ -131,9 +132,8 @@ template VTKM_CONT vtkm::cont::ArrayHandle* DynamicArrayHandleTryCast( vtkm::cont::detail::PolymorphicArrayHandleContainerBase* arrayContainer) { - vtkm::cont::detail::PolymorphicArrayHandleContainer* downcastContainer = - dynamic_cast*>( - arrayContainer); + vtkm::cont::detail::PolymorphicArrayHandleContainer* downcastContainer = nullptr; + downcastContainer = dynamic_cast(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 - VTKM_CONT void CastAndCall(const Functor& f) const; + template + 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 -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 - VTKM_CONT void operator()(Storage) + template + void operator()(std::pair&& p, Args&&... args) const { - this->DoCast(Storage(), - typename vtkm::cont::internal::IsValidArrayHandle::type()); + using storage = vtkm::cont::internal::Storage; + using invalid = typename std::is_base_of::type; + this->run(std::forward(p), invalid{}, args...); } -private: - template - void DoCast(Storage, std::true_type) + template + void run(std::pair&&, std::false_type, Functor&& f, bool& called, Args&&... args) const { - if (!this->FoundCast && this->Array->template IsTypeAndStorage()) + if (!called) { - this->Function(this->Array->template CastToTypeStorage()); - this->FoundCast = true; + using downcastType = const vtkm::cont::detail::PolymorphicArrayHandleContainer* const; + downcastType downcastContainer = dynamic_cast(this->Container); + if (downcastContainer) + { + f(downcastContainer->Array, std::forward(args)...); + called = true; + } } } - template - void DoCast(Storage, std::false_type) + template + void run(std::pair&&, std::true_type, Args&&...) const { - // This type of array handle cannot exist, so do nothing. } - void operator=(const DynamicArrayHandleTryStorage&) = delete; -}; - -template -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 - VTKM_CONT void operator()(Type) - { - if (this->FoundCast) - { - return; - } - using TryStorageType = DynamicArrayHandleTryStorage; - TryStorageType tryStorage = TryStorageType(*this->Array, this->Function); - - vtkm::ListForEach(tryStorage, StorageList()); - if (tryStorage.FoundCast) - { - this->FoundCast = true; - } - } - -private: - void operator=(const DynamicArrayHandleTryType&) = delete; + const PolymorphicArrayHandleContainerBase* const Container; }; +VTKM_CONT_EXPORT void ThrowCastAndCallException(PolymorphicArrayHandleContainerBase*, + const std::type_info*, + const std::type_info*); } // namespace detail template -template -VTKM_CONT void DynamicArrayHandleBase::CastAndCall(const Functor& f) const +template +VTKM_CONT void DynamicArrayHandleBase::CastAndCall(const Functor& f, + Args&&... args) const { - VTKM_IS_LIST_TAG(TypeList); - VTKM_IS_LIST_TAG(StorageList); - using TryTypeType = detail::DynamicArrayHandleTryType; + //For optimizations we should compile once the cross product for the default types + //and make it extern + using crossProduct = typename vtkm::ListCrossProduct; - // 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)...); + 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 -VTKM_CONT void -DynamicArrayHandleBase::CastAndCall( - const Functor& f) const -{ - - using TryTypeType = detail::DynamicArrayHandleTryType; - - // 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 void CastAndCall(const Functor& f) const; + template + 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 -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 - VTKM_CONT void operator()(CellSetType) + template + void operator()(CellSetType, Functor&& f, bool& called, Args&&... args) const { - if (!this->FoundCast) + using downcastType = const vtkm::cont::internal::SimplePolymorphicContainer* const; + if (!called) { - CellSetType* cellSet = detail::DynamicCellSetTryCast(this->CellSetContainer); - if (cellSet != nullptr) + downcastType downcastContainer = dynamic_cast(this->Container); + if (downcastContainer) { - this->Function(*cellSet); - this->FoundCast = true; + f(downcastContainer->Item, std::forward(args)...); + called = true; } } } -private: - void operator=(const DynamicCellSetTryCellSet&) = delete; + const vtkm::cont::internal::SimplePolymorphicContainerBase* const Container; }; } // namespace detail template -template -VTKM_CONT void DynamicCellSetBase::CastAndCall(const Functor& f) const +template +VTKM_CONT void DynamicCellSetBase::CastAndCall(const Functor& f, Args&&... args) const { - using TryCellSetType = detail::DynamicCellSetTryCellSet; - 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)...); + if (!called) { throw vtkm::cont::ErrorBadValue("Could not find appropriate cast for cell set."); } diff --git a/vtkm/cont/Error.h b/vtkm/cont/Error.h index b2144b38c..4a036a76e 100644 --- a/vtkm/cont/Error.h +++ b/vtkm/cont/Error.h @@ -27,7 +27,7 @@ #include #include -#include // For VTKM_OVERRIDE +#include // 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() {} diff --git a/vtkm/cont/Field.h b/vtkm/cont/Field.h index 8efced0f6..79e1386ca 100644 --- a/vtkm/cont/Field.h +++ b/vtkm/cont/Field.h @@ -401,10 +401,10 @@ private: } }; -template -void CastAndCall(const vtkm::cont::Field& field, const Functor& f) +template +void CastAndCall(const vtkm::cont::Field& field, const Functor& f, Args&&... args) { - field.GetData().CastAndCall(f); + field.GetData().CastAndCall(f, std::forward(args)...); } namespace internal diff --git a/vtkm/cont/ImplicitFunction.cxx b/vtkm/cont/ImplicitFunction.cxx deleted file mode 100644 index ae0201390..000000000 --- a/vtkm/cont/ImplicitFunction.cxx +++ /dev/null @@ -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 - -namespace vtkm -{ -namespace cont -{ - -ImplicitFunction::~ImplicitFunction() = default; -} -} // vtkm::cont diff --git a/vtkm/cont/ImplicitFunction.h b/vtkm/cont/ImplicitFunction.h deleted file mode 100644 index 48128b219..000000000 --- a/vtkm/cont/ImplicitFunction.h +++ /dev/null @@ -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 - -#include -#include - -#include - -namespace vtkm -{ -namespace cont -{ - -class VTKM_CONT_EXPORT ImplicitFunction -{ -public: - virtual ~ImplicitFunction(); - - template - 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; - - 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 Cache; -}; - -template -class VTKM_ALWAYS_EXPORT ImplicitFunctionImpl : public ImplicitFunction -{ -public: - template - void ResetDevices(DeviceAdapterList devices) - { - this->Cache->Bind(static_cast(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(this)); } -}; - -//============================================================================ -// ImplicitFunctions: - -//============================================================================ -/// \brief Implicit function for a box -class VTKM_ALWAYS_EXPORT Box : public ImplicitFunctionImpl -{ -public: - Box(); - Box(vtkm::Vec minPoint, vtkm::Vec maxPoint); - Box(FloatDefault xmin, - FloatDefault xmax, - FloatDefault ymin, - FloatDefault ymax, - FloatDefault zmin, - FloatDefault zmax); - - void SetMinPoint(const vtkm::Vec& point); - void SetMaxPoint(const vtkm::Vec& point); - - const vtkm::Vec& GetMinPoint() const; - const vtkm::Vec& GetMaxPoint() const; - - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec& x) const; - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const; - - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec& x) const; - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const; - -private: - vtkm::Vec MinPoint; - vtkm::Vec MaxPoint; -}; - -//============================================================================ -/// \brief Implicit function for a cylinder -class VTKM_ALWAYS_EXPORT Cylinder : public ImplicitFunctionImpl -{ -public: - Cylinder(); - Cylinder(const vtkm::Vec& axis, FloatDefault radius); - Cylinder(const vtkm::Vec& center, - const vtkm::Vec& axis, - FloatDefault radius); - - void SetCenter(const vtkm::Vec& center); - void SetAxis(const vtkm::Vec& axis); - void SetRadius(FloatDefault radius); - - const vtkm::Vec& GetCenter() const; - const vtkm::Vec& GetAxis() const; - FloatDefault GetRadius() const; - - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec& x) const; - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const; - - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec& x) const; - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const; - -private: - vtkm::Vec Center; - vtkm::Vec Axis; - FloatDefault Radius; -}; - -//============================================================================ -/// \brief Implicit function for a frustum -class VTKM_ALWAYS_EXPORT Frustum : public ImplicitFunctionImpl -{ -public: - Frustum(); - Frustum(const vtkm::Vec points[6], const vtkm::Vec normals[6]); - explicit Frustum(const vtkm::Vec points[8]); - - void SetPlanes(const vtkm::Vec points[6], - const vtkm::Vec normals[6]); - void SetPlane(int idx, vtkm::Vec& point, vtkm::Vec& normal); - - void GetPlanes(vtkm::Vec points[6], vtkm::Vec normals[6]) const; - const vtkm::Vec* GetPoints() const; - const vtkm::Vec* GetNormals() const; - - // The points should be specified in the order of hex-cell vertices - void CreateFromPoints(const vtkm::Vec points[8]); - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const; - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec& x) const; - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault, FloatDefault, FloatDefault) const; - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec&) const; - -private: - vtkm::Vec Points[6]; - vtkm::Vec Normals[6]; -}; - -//============================================================================ -/// \brief Implicit function for a plane -class VTKM_ALWAYS_EXPORT Plane : public ImplicitFunctionImpl -{ -public: - Plane(); - explicit Plane(const vtkm::Vec& normal); - Plane(const vtkm::Vec& origin, const vtkm::Vec& normal); - - void SetOrigin(const vtkm::Vec& origin); - void SetNormal(const vtkm::Vec& normal); - - const vtkm::Vec& GetOrigin() const; - const vtkm::Vec& GetNormal() const; - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const; - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec& x) const; - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault, FloatDefault, FloatDefault) const; - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec&) const; - -private: - vtkm::Vec Origin; - vtkm::Vec Normal; -}; - -//============================================================================ -/// \brief Implicit function for a sphere -class VTKM_ALWAYS_EXPORT Sphere : public ImplicitFunctionImpl -{ -public: - Sphere(); - explicit Sphere(FloatDefault radius); - Sphere(vtkm::Vec center, FloatDefault radius); - - void SetRadius(FloatDefault radius); - void SetCenter(const vtkm::Vec& center); - - FloatDefault GetRadius() const; - const vtkm::Vec& GetCenter() const; - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const; - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec& x) const; - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const; - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec& x) const; - -private: - FloatDefault Radius; - vtkm::Vec Center; -}; -} -} // vtkm::cont - -#include - -#endif // vtk_m_cont_ImplicitFunction_h diff --git a/vtkm/cont/ImplicitFunction.hxx b/vtkm/cont/ImplicitFunction.hxx deleted file mode 100644 index 4cc3464d1..000000000 --- a/vtkm/cont/ImplicitFunction.hxx +++ /dev/null @@ -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 -#include - -#include -#include - -namespace vtkm -{ -namespace cont -{ - -//============================================================================ -inline Box::Box() - : MinPoint(vtkm::Vec(FloatDefault(0))) - , MaxPoint(vtkm::Vec(FloatDefault(1))) -{ -} - -inline Box::Box(vtkm::Vec minPoint, vtkm::Vec 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& point) -{ - this->MinPoint = point; - this->Modified(); -} - -inline void Box::SetMaxPoint(const vtkm::Vec& point) -{ - this->MaxPoint = point; - this->Modified(); -} - -inline const vtkm::Vec& Box::GetMinPoint() const -{ - return this->MinPoint; -} - -inline const vtkm::Vec& 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(x, y, z)); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Box::Gradient(FloatDefault x, - FloatDefault y, - FloatDefault z) const -{ - return this->Gradient(vtkm::Vec(x, y, z)); -} - -VTKM_EXEC_CONT -inline FloatDefault Box::Value(const vtkm::Vec& 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 Box::Gradient(const vtkm::Vec& x) const -{ - vtkm::IdComponent minAxis = 0; - FloatDefault dist = 0.0; - FloatDefault minDist = vtkm::Infinity32(); - vtkm::Vec location; - vtkm::Vec normal(FloatDefault(0)); - vtkm::Vec inside(FloatDefault(0)); - vtkm::Vec outside(FloatDefault(0)); - vtkm::Vec 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& axis, FloatDefault radius) - : Center(FloatDefault(0)) - , Axis(vtkm::Normal(axis)) - , Radius(radius) -{ -} - -inline Cylinder::Cylinder(const vtkm::Vec& center, - const vtkm::Vec& axis, - FloatDefault radius) - : Center(center) - , Axis(vtkm::Normal(axis)) - , Radius(radius) -{ -} - -inline void Cylinder::SetCenter(const vtkm::Vec& center) -{ - this->Center = center; - this->Modified(); -} - -inline void Cylinder::SetAxis(const vtkm::Vec& axis) -{ - this->Axis = vtkm::Normal(axis); - this->Modified(); -} - -inline void Cylinder::SetRadius(FloatDefault radius) -{ - this->Radius = radius; - this->Modified(); -} - -inline const vtkm::Vec& Cylinder::GetCenter() const -{ - return this->Center; -} - -inline const vtkm::Vec& Cylinder::GetAxis() const -{ - return this->Axis; -} - -inline FloatDefault Cylinder::GetRadius() const -{ - return this->Radius; -} - -VTKM_EXEC_CONT -inline FloatDefault Cylinder::Value(const vtkm::Vec& x) const -{ - vtkm::Vec 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(x, y, z)); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Cylinder::Gradient(const vtkm::Vec& x) const -{ - vtkm::Vec x2c = x - this->Center; - FloatDefault t = this->Axis[0] * x2c[0] + this->Axis[1] * x2c[1] + this->Axis[2] * x2c[2]; - vtkm::Vec closestPoint = this->Center + (this->Axis * t); - return (x - closestPoint) * FloatDefault(2); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Cylinder::Gradient(FloatDefault x, - FloatDefault y, - FloatDefault z) const -{ - return this->Gradient(vtkm::Vec(x, y, z)); -} - -//============================================================================ -inline Frustum::Frustum() -{ - std::fill(this->Points, this->Points + 6, vtkm::Vec{}); - std::fill(this->Normals, this->Normals + 6, vtkm::Vec{}); -} - -inline Frustum::Frustum(const vtkm::Vec points[6], - const vtkm::Vec normals[6]) -{ - std::copy(points, points + 6, this->Points); - std::copy(normals, normals + 6, this->Normals); -} - -inline Frustum::Frustum(const vtkm::Vec points[8]) -{ - this->CreateFromPoints(points); -} - -inline void Frustum::SetPlanes(const vtkm::Vec points[6], - const vtkm::Vec normals[6]) -{ - std::copy(points, points + 6, this->Points); - std::copy(normals, normals + 6, this->Normals); - this->Modified(); -} - -inline void Frustum::GetPlanes(vtkm::Vec points[6], - vtkm::Vec normals[6]) const -{ - std::copy(this->Points, this->Points + 6, points); - std::copy(this->Normals, this->Normals + 6, normals); -} - -inline const vtkm::Vec* Frustum::GetPoints() const -{ - return this->Points; -} - -inline const vtkm::Vec* Frustum::GetNormals() const -{ - return this->Normals; -} - -inline void Frustum::SetPlane(int idx, - vtkm::Vec& point, - vtkm::Vec& 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 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::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& x) const -{ - return this->Value(x[0], x[1], x[2]); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Frustum::Gradient(FloatDefault x, - FloatDefault y, - FloatDefault z) const -{ - FloatDefault maxVal = -std::numeric_limits::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 Frustum::Gradient(const vtkm::Vec& 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& normal) - : Origin(FloatDefault(0)) - , Normal(normal) -{ -} - -inline Plane::Plane(const vtkm::Vec& origin, - const vtkm::Vec& normal) - : Origin(origin) - , Normal(normal) -{ -} - -inline void Plane::SetOrigin(const vtkm::Vec& origin) -{ - this->Origin = origin; - this->Modified(); -} - -inline void Plane::SetNormal(const vtkm::Vec& normal) -{ - this->Normal = normal; - this->Modified(); -} - -inline const vtkm::Vec& Plane::GetOrigin() const -{ - return this->Origin; -} - -inline const vtkm::Vec& 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& x) const -{ - return this->Value(x[0], x[1], x[2]); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Plane::Gradient(FloatDefault, FloatDefault, FloatDefault) const -{ - return this->Normal; -} - -VTKM_EXEC_CONT -inline vtkm::Vec Plane::Gradient(const vtkm::Vec&) 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 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& center) -{ - this->Center = center; - this->Modified(); -} - -inline FloatDefault Sphere::GetRadius() const -{ - return this->Radius; -} - -inline const vtkm::Vec& 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& x) const -{ - return this->Value(x[0], x[1], x[2]); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Sphere::Gradient(FloatDefault x, - FloatDefault y, - FloatDefault z) const -{ - return this->Gradient(vtkm::Vec(x, y, z)); -} - -VTKM_EXEC_CONT -inline vtkm::Vec Sphere::Gradient(const vtkm::Vec& x) const -{ - return FloatDefault(2) * (x - this->Center); -} -} -} // vtkm::cont diff --git a/vtkm/cont/ImplicitFunctionHandle.h b/vtkm/cont/ImplicitFunctionHandle.h new file mode 100644 index 000000000..035176d0d --- /dev/null +++ b/vtkm/cont/ImplicitFunctionHandle.h @@ -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 +#include + +namespace vtkm +{ +namespace cont +{ + +class VTKM_ALWAYS_EXPORT ImplicitFunctionHandle + : public vtkm::cont::VirtualObjectHandle +{ +private: + using Superclass = vtkm::cont::VirtualObjectHandle; + +public: + ImplicitFunctionHandle() = default; + + template + explicit ImplicitFunctionHandle(ImplicitFunctionType* function, + bool acquireOwnership = true, + DeviceAdapterList devices = DeviceAdapterList()) + : Superclass(function, acquireOwnership, devices) + { + } +}; + +template +VTKM_CONT ImplicitFunctionHandle +make_ImplicitFunctionHandle(ImplicitFunctionType&& func, + DeviceAdapterList devices = DeviceAdapterList()) +{ + using IFType = typename std::remove_reference::type; + return ImplicitFunctionHandle( + new IFType(std::forward(func)), true, devices); +} + +template +VTKM_CONT ImplicitFunctionHandle make_ImplicitFunctionHandle(Args&&... args) +{ + return ImplicitFunctionHandle(new ImplicitFunctionType(std::forward(args)...), + true, + VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG()); +} + +template +VTKM_CONT ImplicitFunctionHandle make_ImplicitFunctionHandle(Args&&... args) +{ + return ImplicitFunctionHandle( + new ImplicitFunctionType(std::forward(args)...), true, DeviceAdapterList()); +} +} +} // vtkm::cont + +#endif // vtk_m_cont_ImplicitFunctionHandle_h diff --git a/vtkm/cont/RuntimeDeviceTracker.cxx b/vtkm/cont/RuntimeDeviceTracker.cxx index 33ef3b43c..d0b31714e 100644 --- a/vtkm/cont/RuntimeDeviceTracker.cxx +++ b/vtkm/cont/RuntimeDeviceTracker.cxx @@ -29,7 +29,10 @@ #include #include +#include +#include #include +#include #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 globalTrackers; + std::thread::id this_id = std::this_thread::get_id(); + + std::unique_lock 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 diff --git a/vtkm/cont/RuntimeDeviceTracker.h b/vtkm/cont/RuntimeDeviceTracker.h index 3b355b756..ebdfe5e72 100644 --- a/vtkm/cont/RuntimeDeviceTracker.h +++ b/vtkm/cont/RuntimeDeviceTracker.h @@ -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 diff --git a/vtkm/cont/TryExecute.cxx b/vtkm/cont/TryExecute.cxx new file mode 100644 index 000000000..a2b1fe66c --- /dev/null +++ b/vtkm/cont/TryExecute.cxx @@ -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 +#include +#include +#include + +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; + } +} +} +} +} diff --git a/vtkm/cont/TryExecute.h b/vtkm/cont/TryExecute.h index 702a7086a..b5b253401 100644 --- a/vtkm/cont/TryExecute.h +++ b/vtkm/cont/TryExecute.h @@ -21,10 +21,6 @@ #define vtk_m_cont_TryExecute_h #include -#include -#include -#include - #include namespace vtkm @@ -35,123 +31,171 @@ namespace cont namespace detail { -template -struct TryExecuteRunIfValid; +VTKM_CONT_EXPORT void HandleTryExecuteException(vtkm::Int8, + const std::string&, + vtkm::cont::RuntimeDeviceTracker&); -template -struct TryExecuteRunIfValid +template +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 -struct TryExecuteRunIfValid -{ - 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)...); } - - // If we are here, then the functor was either never run or failed. - return false; - } -}; - -template -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 - VTKM_CONT bool operator()(Device) - { - if (!this->Success) + catch (...) { - using DeviceTraits = vtkm::cont::DeviceAdapterTraits; - - this->Success = detail::TryExecuteRunIfValid::Run( - this->Functor, this->Tracker); + using Traits = vtkm::cont::DeviceAdapterTraits; + HandleTryExecuteException(Traits::GetId(), Traits::GetName(), tracker); } - - return this->Success; } -private: - void operator=(const TryExecuteImpl&) = delete; + // If we are here, then the functor was either never run or failed. + return false; +} + +template +bool TryExecuteIfValid(std::false_type, + DeviceTag, + Functor&&, + vtkm::cont::RuntimeDeviceTracker&, + Args&&...) +{ + return false; +} + +struct TryExecuteWrapper +{ + template + void operator()(DeviceTag tag, + Functor&& f, + vtkm::cont::RuntimeDeviceTracker& tracker, + bool& ran, + Args&&... args) const + { + if (!ran) + { + using DeviceTraits = vtkm::cont::DeviceAdapterTraits; + ran = TryExecuteIfValid(std::integral_constant(), + tag, + std::forward(f), + std::forward(tracker), + std::forward(args)...); + } + } }; +template +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), tracker, success, std::forward(args)...); + return success; +} + +template +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), + tracker, + success, + std::forward(args)...); + return success; +} + +template +VTKM_CONT bool TryExecuteImpl(Functor&& functor, + std::true_type t, + std::false_type, + Arg1&& arg1, + Args&&... args) +{ + return TryExecuteImpl(std::forward(functor), + t, + t, + std::forward(arg1), + VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG(), + std::forward(args)...); +} + +template +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), + t, + t, + tracker, + std::forward(arg1), + std::forward(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 +/// 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::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 -VTKM_CONT bool TryExecute(const Functor& functor, - vtkm::cont::RuntimeDeviceTracker tracker, - DeviceList) -{ - detail::TryExecuteImpl internals(functor, tracker); - vtkm::ListForEach(internals, DeviceList()); - return internals.Success; -} -template -VTKM_CONT bool TryExecute(Functor& functor, vtkm::cont::RuntimeDeviceTracker tracker, DeviceList) -{ - detail::TryExecuteImpl internals(functor, tracker); - vtkm::ListForEach(internals, DeviceList()); - return internals.Success; -} -template -VTKM_CONT bool TryExecute(const Functor& functor, DeviceList) -{ - return vtkm::cont::TryExecute(functor, vtkm::cont::GetGlobalRuntimeDeviceTracker(), DeviceList()); -} -template -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 -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 -VTKM_CONT bool TryExecute( - Functor& functor, - vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()) +template +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::type; + using is_tracker = typename std::is_base_of::type>::type; + + return detail::TryExecuteImpl( + functor, is_tracker{}, is_deviceAdapter{}, std::forward(arg1)); } +template +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::type>::type; + using is_arg2_devicelist = typename std::is_base_of::type; + + //We now know what of three states we are currently at + using has_runtime_and_deviceAdapter = + brigand::bool_; + using has_just_runtime = brigand::bool_; + using has_just_devicelist = typename std::is_base_of::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_; + using second_true = + brigand::bool_; + + return detail::TryExecuteImpl(functor, + first_true{}, + second_true{}, + std::forward(arg1), + std::forward(arg2), + std::forward(args)...); +} + +//@} //block doxygen all TryExecute functions } } // namespace vtkm::cont diff --git a/vtkm/cont/VirtualObjectCache.h b/vtkm/cont/VirtualObjectCache.h deleted file mode 100644 index b82738497..000000000 --- a/vtkm/cont/VirtualObjectCache.h +++ /dev/null @@ -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 -#include -#include -#include -#include - -#include -#include - -#define VTKM_MAX_DEVICE_ADAPTER_ID 8 - -namespace vtkm -{ -namespace cont -{ - -/// \brief Implements VTK-m's execution side Virtual Methods -/// 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 const void* 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 const void* 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 -/// 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 const TargetClass* 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 -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 - void Bind(const TargetClass* target, DeviceAdapterList devices = DeviceAdapterList()) - { - this->Reset(); - - this->Target = target; - vtkm::ListForEach(CreateTransferInterface(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 - VirtualObject GetVirtualObject(DeviceAdapter) - { - using DeviceInfo = vtkm::cont::DeviceAdapterTraits; - - 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(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(this->CurrentDevice)]; - } - - template - class CreateTransferInterface - { - private: - template - using EnableIfValid = std::enable_if::Valid>; - - template - using EnableIfInvalid = std::enable_if::Valid>; - - public: - CreateTransferInterface(TransferInterface* transfers) - : Transfers(transfers) - { - } - - // Use SFINAE to create entries for valid device adapters only - template - typename EnableIfValid::type operator()(DeviceAdapter) const - { - using DeviceInfo = vtkm::cont::DeviceAdapterTraits; - - if (DeviceInfo::GetId() >= 0 && DeviceInfo::GetId() < VTKM_MAX_DEVICE_ADAPTER_ID) - { - using TransferImpl = - internal::VirtualObjectTransfer; - - std::size_t id = static_cast(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 EnableIfInvalid::type operator()(DeviceAdapter) const - { - } - - private: - TransferInterface* Transfers; - }; - - const void* Target; - std::array 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 diff --git a/vtkm/cont/VirtualObjectHandle.h b/vtkm/cont/VirtualObjectHandle.h new file mode 100644 index 000000000..c372d3095 --- /dev/null +++ b/vtkm/cont/VirtualObjectHandle.h @@ -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 +#include +#include +#include +#include + +#include +#include + +#define VTKM_MAX_DEVICE_ADAPTER_ID 8 + +namespace vtkm +{ +namespace cont +{ + +/// \brief Implements VTK-m's execution side Virtual Methods 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 +class VTKM_ALWAYS_EXPORT VirtualObjectHandle +{ + VTKM_STATIC_ASSERT_MSG((std::is_base_of::value), + "All virtual objects must be subclass of vtkm::VirtualObjectBase."); + +public: + VTKM_CONT VirtualObjectHandle() + : Internals(new InternalStruct) + { + } + + template + 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 + VTKM_CONT void Reset(VirtualDerivedType* derived, + bool acquireOwnership = true, + DeviceAdapterList devices = DeviceAdapterList()) + { + this->Reset(); + if (derived) + { + VTKM_STATIC_ASSERT_MSG((std::is_base_of::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(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 + VTKM_CONT const VirtualBaseType* PrepareForExecution(DeviceAdapter) const + { + using DeviceInfo = vtkm::cont::DeviceAdapterTraits; + + 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(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(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 + 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::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 Transfer; + }; + + template + class CreateTransferInterface + { + public: + CreateTransferInterface(std::unique_ptr* transfers, + const VirtualDerivedType* virtualObject) + : Transfers(transfers) + , VirtualObject(virtualObject) + { + } + + template + void operator()(DeviceAdapter) const + { + using DeviceInfo = vtkm::cont::DeviceAdapterTraits; + + 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; + this->Transfers[DeviceInfo::GetId()].reset(new TransferImpl(this->VirtualObject)); + } + + private: + std::unique_ptr* Transfers; + const VirtualDerivedType* VirtualObject; + }; + + struct InternalStruct + { + VirtualBaseType* VirtualObject = nullptr; + bool Owner = false; + std::array, 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 Internals; +}; +} +} // vtkm::cont + +#undef VTKM_MAX_DEVICE_ADAPTER_ID + +#endif // vtk_m_cont_VirtualObjectHandle_h diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu index 136c6e048..5118ee02f 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu @@ -133,6 +133,7 @@ void ExecutionArrayInterfaceBasic::CopyFromControl( // execution: if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr)) { + CudaAllocator::PrepareForInput(executionPtr, numBytes); return; } @@ -160,6 +161,7 @@ void ExecutionArrayInterfaceBasic::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::CopyToControl(const voi cudaStreamPerThread)); } +void ExecutionArrayInterfaceBasic::UsingForRead(const void* controlPtr, + const void* executionPtr, + vtkm::UInt64 numBytes) const +{ + CudaAllocator::PrepareForInput(executionPtr, static_cast(numBytes)); +} + +void ExecutionArrayInterfaceBasic::UsingForWrite(const void* controlPtr, + const void* executionPtr, + vtkm::UInt64 numBytes) const +{ + CudaAllocator::PrepareForOutput(executionPtr, static_cast(numBytes)); +} + +void ExecutionArrayInterfaceBasic::UsingForReadWrite( + const void* controlPtr, + const void* executionPtr, + vtkm::UInt64 numBytes) const +{ + CudaAllocator::PrepareForInPlace(executionPtr, static_cast(numBytes)); +} + + } // end namespace internal VTKM_INSTANTIATE_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagCuda) diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h index 404b8cc8b..1f77176e4 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h @@ -159,6 +159,16 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic 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 diff --git a/vtkm/cont/cuda/internal/CudaAllocator.cu b/vtkm/cont/cuda/internal/CudaAllocator.cu index 8db04ba8e..aa7c35ea2 100644 --- a/vtkm/cont/cuda/internal/CudaAllocator.cu +++ b/vtkm/cont/cuda/internal/CudaAllocator.cu @@ -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 } diff --git a/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h b/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h index 8900bafe5..80430b19b 100644 --- a/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h +++ b/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h @@ -20,6 +20,7 @@ #ifndef vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h #define vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h +#include #include #include #include @@ -34,45 +35,108 @@ namespace internal namespace detail { -template -__global__ void CreateKernel(VirtualObject* object, const TargetClass* target) +template +__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 +__global__ void UpdateVirtualObjectKernel(VirtualDerivedType* deviceObject, + const VirtualDerivedType* targetObject) +{ + *deviceObject = *targetObject; +} + +template +__global__ void DeleteVirtualObjectKernel(VirtualDerivedType* deviceObject) +{ + deviceObject->~VirtualDerivedType(); } } // detail -template -struct VirtualObjectTransfer +template +struct VirtualObjectTransfer { - 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 + #endif // vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h diff --git a/vtkm/cont/cuda/testing/CMakeLists.txt b/vtkm/cont/cuda/testing/CMakeLists.txt index e0db2179b..4ceeb0a28 100644 --- a/vtkm/cont/cuda/testing/CMakeLists.txt +++ b/vtkm/cont/cuda/testing/CMakeLists.txt @@ -29,6 +29,6 @@ set(unit_tests UnitTestCudaMath.cu UnitTestCudaShareUserProvidedManagedMemory.cu UnitTestCudaPointLocatorUniformGrid.cxx - UnitTestCudaVirtualObjectCache.cu + UnitTestCudaVirtualObjectHandle.cu ) vtkm_unit_tests(SOURCES ${unit_tests}) diff --git a/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu b/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectHandle.cu similarity index 79% rename from vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu rename to vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectHandle.cu index b0b00e461..193b7184a 100644 --- a/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu +++ b/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectHandle.cu @@ -17,22 +17,22 @@ // Laboratory (LANL), the U.S. Government retains certain rights in // this software. //============================================================================ -#include +#include namespace { -void TestVirtualObjectCache() +void TestVirtualObjectHandle() { using DeviceAdapterList = vtkm::ListTagBase; - vtkm::cont::testing::TestingVirtualObjectCache::Run(); + vtkm::cont::testing::TestingVirtualObjectHandle::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); } diff --git a/vtkm/cont/internal/ArrayHandleBasicImpl.h b/vtkm/cont/internal/ArrayHandleBasicImpl.h index 737d8814a..465bfb222 100644 --- a/vtkm/cont/internal/ArrayHandleBasicImpl.h +++ b/vtkm/cont/internal/ArrayHandleBasicImpl.h @@ -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; }; diff --git a/vtkm/cont/internal/ArrayHandleBasicImpl.hxx b/vtkm/cont/internal/ArrayHandleBasicImpl.hxx index 272e03fab..be03fe2a4 100644 --- a/vtkm/cont/internal/ArrayHandleBasicImpl.hxx +++ b/vtkm/cont/internal/ArrayHandleBasicImpl.hxx @@ -267,6 +267,9 @@ ArrayHandle::PrepareForInput(DeviceAdapterTag device) const InternalStruct* priv = const_cast(this->Internals.get()); this->PrepareForDevice(device); + const vtkm::UInt64 numBytes = static_cast(sizeof(ValueType)) * + static_cast(this->GetStorage().GetNumberOfValues()); + if (!this->Internals->ExecutionArrayValid) { @@ -284,9 +287,6 @@ ArrayHandle::PrepareForInput(DeviceAdapterTag device) const this->Internals->ControlArray.GetBasePointer(), this->Internals->ControlArray.GetCapacityPointer()); - const vtkm::UInt64 numBytes = static_cast(sizeof(ValueType)) * - static_cast(this->GetStorage().GetNumberOfValues()); - priv->ExecutionInterface->Allocate(execArray, numBytes); priv->ExecutionInterface->CopyFromControl( @@ -294,6 +294,8 @@ ArrayHandle::PrepareForInput(DeviceAdapterTag device) const this->Internals->ExecutionArrayValid = true; } + this->Internals->ExecutionInterface->UsingForRead( + priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes); return PortalFactory::CreatePortalConst(this->Internals->ExecutionArray, this->Internals->ExecutionArrayEnd); @@ -320,11 +322,17 @@ ArrayHandle::PrepareForOutput(vtkm::Id numVals, DeviceAdapte this->Internals->ControlArray.GetBasePointer(), this->Internals->ControlArray.GetCapacityPointer()); - this->Internals->ExecutionInterface->Allocate( - execArray, static_cast(sizeof(ValueType)) * static_cast(numVals)); + const vtkm::UInt64 numBytes = + static_cast(sizeof(ValueType)) * static_cast(numVals); + + this->Internals->ExecutionInterface->Allocate(execArray, numBytes); + + this->Internals->ExecutionInterface->UsingForWrite( + priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes); this->Internals->ExecutionArrayValid = true; + return PortalFactory::CreatePortal(this->Internals->ExecutionArray, this->Internals->ExecutionArrayEnd); } @@ -339,6 +347,9 @@ ArrayHandle::PrepareForInPlace(DeviceAdapterTag device) this->PrepareForDevice(device); + const vtkm::UInt64 numBytes = static_cast(sizeof(ValueType)) * + static_cast(this->GetStorage().GetNumberOfValues()); + if (!this->Internals->ExecutionArrayValid) { // Initialize an empty array if needed: @@ -355,9 +366,6 @@ ArrayHandle::PrepareForInPlace(DeviceAdapterTag device) this->Internals->ControlArray.GetBasePointer(), this->Internals->ControlArray.GetCapacityPointer()); - vtkm::UInt64 numBytes = static_cast(sizeof(ValueType)) * - static_cast(this->GetStorage().GetNumberOfValues()); - priv->ExecutionInterface->Allocate(execArray, numBytes); priv->ExecutionInterface->CopyFromControl( @@ -366,6 +374,9 @@ ArrayHandle::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; diff --git a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx index 1424af77c..d9347dfdd 100644 --- a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx +++ b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx @@ -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 diff --git a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h index faf0934a3..0589c1481 100644 --- a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h +++ b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h @@ -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; }; } } diff --git a/vtkm/cont/internal/DynamicTransform.h b/vtkm/cont/internal/DynamicTransform.h index 444070944..1195723c0 100644 --- a/vtkm/cont/internal/DynamicTransform.h +++ b/vtkm/cont/internal/DynamicTransform.h @@ -22,6 +22,8 @@ #include "vtkm/internal/IndexTag.h" +#include + namespace vtkm { namespace cont @@ -29,6 +31,10 @@ namespace cont template class ArrayHandle; + +class CoordinateSystem; +class Field; + template class CellSetStructured; template @@ -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 -void CastAndCall(const DynamicObject& dynamicObject, const Functor& f) +template +void CastAndCall(const DynamicObject& dynamicObject, const Functor& f, Args&&... args) { - dynamicObject.CastAndCall(f); + dynamicObject.CastAndCall(f, std::forward(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 +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 +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 -void CastAndCall(const vtkm::cont::ArrayHandle& handle, const Functor& f) +template +void CastAndCall(const vtkm::cont::ArrayHandle& handle, const Functor& f, Args&&... args) { - f(handle); + f(handle, std::forward(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 -void CastAndCall(const vtkm::cont::CellSetStructured& cellset, const Functor& f) +template +void CastAndCall(const vtkm::cont::CellSetStructured& cellset, + const Functor& f, + Args&&... args) { - f(cellset); + f(cellset, std::forward(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 +template void CastAndCall(const vtkm::cont::CellSetSingleType& cellset, - const Functor& f) + const Functor& f, + Args&&... args) { - f(cellset); + f(cellset, std::forward(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 -void CastAndCall(const vtkm::cont::CellSetExplicit& cellset, const Functor& f) +template +void CastAndCall(const vtkm::cont::CellSetExplicit& cellset, + const Functor& f, + Args&&... args) { - f(cellset); + f(cellset, std::forward(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 +template void CastAndCall(const vtkm::cont::CellSetPermutation& cellset, - const Functor& f) + const Functor& f, + Args&&... args) { - f(cellset); + f(cellset, std::forward(args)...); } namespace internal diff --git a/vtkm/cont/internal/VirtualObjectTransfer.h b/vtkm/cont/internal/VirtualObjectTransfer.h index cc6edea2e..43a000e0d 100644 --- a/vtkm/cont/internal/VirtualObjectTransfer.h +++ b/vtkm/cont/internal/VirtualObjectTransfer.h @@ -20,6 +20,8 @@ #ifndef vtk_m_cont_internal_VirtualObjectTransfer_h #define vtk_m_cont_internal_VirtualObjectTransfer_h +#include + namespace vtkm { namespace cont @@ -27,23 +29,31 @@ namespace cont namespace internal { -template +template 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 ; diff --git a/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h index a3bd32daa..52139a37d 100644 --- a/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h +++ b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h @@ -20,6 +20,11 @@ #ifndef vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h #define vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h +#include +#include + +#include + namespace vtkm { namespace cont @@ -27,19 +32,23 @@ namespace cont namespace internal { -template +template 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(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; }; } } diff --git a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h index d4bd8d06b..cd110a105 100644 --- a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h +++ b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h @@ -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)); diff --git a/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h b/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h index 47ec90663..435759925 100644 --- a/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h +++ b/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h @@ -31,10 +31,14 @@ namespace cont namespace internal { -template -struct VirtualObjectTransfer - : public VirtualObjectTransferShareWithControl +template +struct VirtualObjectTransfer + : VirtualObjectTransferShareWithControl { + VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject) + : VirtualObjectTransferShareWithControl(virtualObject) + { + } }; } } diff --git a/vtkm/cont/serial/testing/CMakeLists.txt b/vtkm/cont/serial/testing/CMakeLists.txt index 983ccbf80..a06d7b2d1 100644 --- a/vtkm/cont/serial/testing/CMakeLists.txt +++ b/vtkm/cont/serial/testing/CMakeLists.txt @@ -28,6 +28,6 @@ set(unit_tests UnitTestSerialDeviceAdapter.cxx UnitTestSerialImplicitFunction.cxx UnitTestSerialPointLocatorUniformGrid.cxx - UnitTestSerialVirtualObjectCache.cxx + UnitTestSerialVirtualObjectHandle.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx b/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectHandle.cxx similarity index 78% rename from vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx rename to vtkm/cont/serial/testing/UnitTestSerialVirtualObjectHandle.cxx index ac33c0318..79c5e6b21 100644 --- a/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx +++ b/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectHandle.cxx @@ -17,21 +17,21 @@ // Laboratory (LANL), the U.S. Government retains certain rights in // this software. //============================================================================ -#include +#include namespace { -void TestVirtualObjectCache() +void TestVirtualObjectHandle() { using DeviceAdapterList = vtkm::ListTagBase; - vtkm::cont::testing::TestingVirtualObjectCache::Run(); + vtkm::cont::testing::TestingVirtualObjectHandle::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); } diff --git a/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h b/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h index 31d804718..9c598afc6 100644 --- a/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h +++ b/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h @@ -31,10 +31,14 @@ namespace cont namespace internal { -template -struct VirtualObjectTransfer - : public VirtualObjectTransferShareWithControl +template +struct VirtualObjectTransfer + : VirtualObjectTransferShareWithControl { + VTKM_CONT VirtualObjectTransfer(const VirtualDerivedType* virtualObject) + : VirtualObjectTransferShareWithControl(virtualObject) + { + } }; } } diff --git a/vtkm/cont/tbb/testing/CMakeLists.txt b/vtkm/cont/tbb/testing/CMakeLists.txt index 38b793c8b..7e0063174 100644 --- a/vtkm/cont/tbb/testing/CMakeLists.txt +++ b/vtkm/cont/tbb/testing/CMakeLists.txt @@ -28,6 +28,6 @@ set(unit_tests UnitTestTBBDeviceAdapter.cxx UnitTestTBBImplicitFunction.cxx UnitTestTBBPointLocatorUniformGrid.cxx - UnitTestTBBVirtualObjectCache.cxx + UnitTestTBBVirtualObjectHandle.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx b/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectHandle.cxx similarity index 79% rename from vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx rename to vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectHandle.cxx index 908e5fce7..073c13d34 100644 --- a/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx +++ b/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectHandle.cxx @@ -17,22 +17,22 @@ // Laboratory (LANL), the U.S. Government retains certain rights in // this software. //============================================================================ -#include +#include namespace { -void TestVirtualObjectCache() +void TestVirtualObjectHandle() { using DeviceAdapterList = vtkm::ListTagBase; - vtkm::cont::testing::TestingVirtualObjectCache::Run(); + vtkm::cont::testing::TestingVirtualObjectHandle::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); } diff --git a/vtkm/cont/testing/CMakeLists.txt b/vtkm/cont/testing/CMakeLists.txt index a5efd683c..8af665b81 100644 --- a/vtkm/cont/testing/CMakeLists.txt +++ b/vtkm/cont/testing/CMakeLists.txt @@ -31,7 +31,7 @@ set(headers TestingFancyArrayHandles.h TestingImplicitFunction.h TestingPointLocatorUniformGrid.h - TestingVirtualObjectCache.h + TestingVirtualObjectHandle.h ) vtkm_declare_headers(${headers}) diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h index 1e71fb779..dffdb96f1 100644 --- a/vtkm/cont/testing/TestingDeviceAdapter.h +++ b/vtkm/cont/testing/TestingDeviceAdapter.h @@ -316,44 +316,31 @@ public: class VirtualObjectTransferKernel { public: - struct Interface + struct Interface : public vtkm::VirtualObjectBase { - using FooSig = vtkm::Id(const void*); - - template - VTKM_EXEC void Bind(const T* target) - { - this->Target = target; - this->FooPtr = [](const void* t) { return static_cast(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; + using Transfer = vtkm::cont::internal::VirtualObjectTransfer; 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 + struct TestCopy> + { + static vtkm::Pair get(vtkm::Id i) + { + return vtkm::make_Pair(TestCopy::get(i), TestCopy::get(i)); + } + }; + template 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::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>(); - TestCopyArrays>(); + TestCopyArrays>(); + TestCopyArrays>(); // - TestCopyArrays>(); - TestCopyArrays>(); - TestCopyArrays>(); - TestCopyArrays>(); + TestCopyArrays>(); + TestCopyArrays>>(); // TestCopyArrays(); TestCopyArrays(); // - TestCopyArrays(); - TestCopyArrays(); TestCopyArrays(); TestCopyArrays(); // @@ -2035,8 +2027,6 @@ private: TestCopyArrays(); TestCopyArrays(); TestCopyArrays(); - // - TestCopyArrays(); } static VTKM_CONT void TestCopyArraysInDiffTypes() diff --git a/vtkm/cont/testing/TestingImplicitFunction.h b/vtkm/cont/testing/TestingImplicitFunction.h index 8cc801da4..fae11ac9a 100644 --- a/vtkm/cont/testing/TestingImplicitFunction.h +++ b/vtkm/cont/testing/TestingImplicitFunction.h @@ -22,7 +22,7 @@ #include #include -#include +#include #include #include #include @@ -50,7 +50,7 @@ public: typedef void ControlSignature(FieldIn, FieldOut); typedef void ExecutionSignature(_1, _2); - EvaluateImplicitFunction(const vtkm::exec::ImplicitFunction& function) + EvaluateImplicitFunction(const vtkm::ImplicitFunction* function) : Function(function) { } @@ -58,16 +58,16 @@ public: template 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 void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points, - const vtkm::cont::ImplicitFunction& function, + const vtkm::cont::ImplicitFunctionHandle& function, vtkm::cont::ArrayHandle& values, DeviceAdapter device) { @@ -141,13 +141,16 @@ private: template void TestBox(DeviceAdapter device) { - std::cout << "Testing vtkm::cont::Box on " + std::cout << "Testing vtkm::Box on " << vtkm::cont::DeviceAdapterTraits::GetName() << "\n"; - vtkm::cont::Box box({ 0.0f, -0.5f, -0.5f }, { 1.5f, 1.5f, 0.5f }); vtkm::cont::ArrayHandle 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 expected = { { 0.0f, -0.5f, 0.5f, 0.5f, 0.0f, -0.5f, 0.5f, 0.5f } @@ -159,17 +162,20 @@ private: template void TestCylinder(DeviceAdapter device) { - std::cout << "Testing vtkm::cont::Cylinder on " + std::cout << "Testing vtkm::Cylinder on " << vtkm::cont::DeviceAdapterTraits::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 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 expected = { { 0.0f, 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f, -1.0f } @@ -181,7 +187,7 @@ private: template void TestFrustum(DeviceAdapter device) { - std::cout << "Testing vtkm::cont::Frustum on " + std::cout << "Testing vtkm::Frustum on " << vtkm::cont::DeviceAdapterTraits::GetName() << "\n"; vtkm::Vec 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 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 expected = { { 0.0f, 0.0f, 0.0f, 0.0f, 0.316228f, 0.316228f, -0.316228f, 0.316228f } @@ -211,23 +220,25 @@ private: template void TestPlane(DeviceAdapter device) { - std::cout << "Testing vtkm::cont::Plane on " + std::cout << "Testing vtkm::Plane on " << vtkm::cont::DeviceAdapterTraits::GetName() << "\n"; - vtkm::cont::Plane plane({ 0.5f, 0.5f, 0.5f }, { 1.0f, 0.0f, 1.0f }); - vtkm::cont::ArrayHandle values; + auto planeHandle = vtkm::cont::make_ImplicitFunctionHandle( + vtkm::make_Vec(0.5f, 0.5f, 0.5f), vtkm::make_Vec(1.0f, 0.0f, 1.0f)); + auto plane = static_cast(planeHandle.Get()); + vtkm::cont::ArrayHandle values; implicit_function_detail::EvaluateOnCoordinates( - this->Input.GetCoordinateSystem(0), plane, values, device); + this->Input.GetCoordinateSystem(0), planeHandle, values, device); std::array 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 expected2 = { { 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f, -1.0f, 0.0f } }; @@ -238,13 +249,15 @@ private: template void TestSphere(DeviceAdapter device) { - std::cout << "Testing vtkm::cont::Sphere on " + std::cout << "Testing vtkm::Sphere on " << vtkm::cont::DeviceAdapterTraits::GetName() << "\n"; - vtkm::cont::Sphere sphere({ 0.0f, 0.0f, 0.0f }, 1.0f); vtkm::cont::ArrayHandle values; implicit_function_detail::EvaluateOnCoordinates( - this->Input.GetCoordinateSystem(0), sphere, values, device); + this->Input.GetCoordinateSystem(0), + vtkm::cont::make_ImplicitFunctionHandle(vtkm::make_Vec(0.0f, 0.0f, 0.0f), 1.0f), + values, + device); std::array expected = { { -1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 2.0f, 1.0f } diff --git a/vtkm/cont/testing/TestingVirtualObjectCache.h b/vtkm/cont/testing/TestingVirtualObjectHandle.h similarity index 65% rename from vtkm/cont/testing/TestingVirtualObjectCache.h rename to vtkm/cont/testing/TestingVirtualObjectHandle.h index acc65b4ab..f0fec6516 100644 --- a/vtkm/cont/testing/TestingVirtualObjectCache.h +++ b/vtkm/cont/testing/TestingVirtualObjectHandle.h @@ -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 #include #include #include -#include +#include #include #define ARRAY_LEN 8 @@ -39,62 +39,76 @@ namespace testing namespace virtual_object_detail { -class Transformer +class Transformer : public vtkm::VirtualObjectBase { public: - template - 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(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 -class TestingVirtualObjectCache +class TestingVirtualObjectHandle { private: using FloatArrayHandle = vtkm::cont::ArrayHandle; using ArrayTransform = - vtkm::cont::ArrayHandleTransform; - using TransformerCache = vtkm::cont::VirtualObjectCache; + vtkm::cont::ArrayHandleTransform; + using TransformerHandle = vtkm::cont::VirtualObjectHandle; 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::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()); } }; } diff --git a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx index 09d34247d..8aac3de04 100644 --- a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx +++ b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx @@ -90,12 +90,14 @@ public: } }; -template -struct VirtualObjectTransfer - : public VirtualObjectTransferShareWithControl +template +struct VirtualObjectTransfer + : public VirtualObjectTransferShareWithControl { + VirtualObjectTransfer(const TargetClass* target) + : VirtualObjectTransferShareWithControl(target) + { + } }; template diff --git a/vtkm/cont/testing/UnitTestDynamicArrayHandle.cxx b/vtkm/cont/testing/UnitTestDynamicArrayHandle.cxx index 5517e5af2..99868f9a1 100644 --- a/vtkm/cont/testing/UnitTestDynamicArrayHandle.cxx +++ b/vtkm/cont/testing/UnitTestDynamicArrayHandle.cxx @@ -106,14 +106,12 @@ struct TestValueFunctor T operator()(vtkm::Id index) const { return TestValue(index, T()); } }; -bool CheckCalled; - struct CheckFunctor { template - void operator()(vtkm::cont::ArrayHandle array) const + void operator()(vtkm::cont::ArrayHandle 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 @@ -150,10 +149,11 @@ void CheckDynamicArray(vtkm::cont::DynamicArrayHandleBase { 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 @@ -223,8 +223,6 @@ void TryNewInstance(T, DynamicArrayType originalArray) template void TryDefaultType(T) { - CheckCalled = false; - vtkm::cont::DynamicArrayHandle array = CreateDynamicArray(T()); CheckDynamicArray(array, vtkm::VecTraits::NUM_COMPONENTS); @@ -237,8 +235,6 @@ struct TryBasicVTKmType template 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; diff --git a/vtkm/cont/testing/UnitTestDynamicCellSet.cxx b/vtkm/cont/testing/UnitTestDynamicCellSet.cxx index a0f4f051b..27c53d6cd 100644 --- a/vtkm/cont/testing/UnitTestDynamicCellSet.cxx +++ b/vtkm/cont/testing/UnitTestDynamicCellSet.cxx @@ -34,17 +34,16 @@ struct NonDefaultCellSetList { }; -bool CheckCalled; - template struct CheckFunctor { - void operator()(const ExpectedCellType&) const { CheckCalled = true; } + void operator()(const ExpectedCellType&, bool& called) const { called = true; } template - 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(); - CheckCalled = false; - dynamicCellSet.CastAndCall(CheckFunctor()); + bool called = false; + dynamicCellSet.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)."); - CheckCalled = false; - CastAndCall(dynamicCellSet, CheckFunctor()); + called = false; + CastAndCall(dynamicCellSet, 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 diff --git a/vtkm/cont/testing/UnitTestTryExecute.cxx b/vtkm/cont/testing/UnitTestTryExecute.cxx index c00788db4..f4afe2009 100644 --- a/vtkm/cont/testing/UnitTestTryExecute.cxx +++ b/vtkm/cont/testing/UnitTestTryExecute.cxx @@ -21,6 +21,7 @@ #define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_ERROR #include +#include #include #include @@ -36,30 +37,27 @@ static const vtkm::Id ARRAY_SIZE = 10; struct TryExecuteTestFunctor { vtkm::IdComponent NumCalls; - vtkm::cont::ArrayHandle InArray; - vtkm::cont::ArrayHandle OutArray; VTKM_CONT - TryExecuteTestFunctor(vtkm::cont::ArrayHandle inArray, - vtkm::cont::ArrayHandle outArray) + TryExecuteTestFunctor() : NumCalls(0) - , InArray(inArray) - , OutArray(outArray) { } template - VTKM_CONT bool operator()(Device) + VTKM_CONT bool operator()(Device, + const vtkm::cont::ArrayHandle& in, + vtkm::cont::ArrayHandle& out) { using Algorithm = vtkm::cont::DeviceAdapterAlgorithm; - Algorithm::Copy(this->InArray, this->OutArray); + Algorithm::Copy(in, out); this->NumCalls++; return true; } }; template -void TryExecuteWithList(DeviceList, bool expectSuccess) +void TryExecuteWithDevice(DeviceList, bool expectSuccess) { vtkm::cont::ArrayHandle inArray; vtkm::cont::ArrayHandle 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 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 +void TryExecuteAllExplicit(DeviceList, bool expectSuccess) +{ + vtkm::cont::RuntimeDeviceTracker tracker; + vtkm::cont::ArrayHandle inArray; + vtkm::cont::ArrayHandle 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 + bool operator()(DeviceList, int, float, bool) const + { + return true; + } + template + bool operator()(DeviceList) const + { + return true; + } +}; + +void TryExecuteAllEdgeCases() +{ + using ValidDevice = vtkm::cont::DeviceAdapterTagSerial; + using SingleValidList = vtkm::ListTagBase; + 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 +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; - TryExecuteWithList(SingleValidList(), true); + TryExecuteTests(SingleValidList(), true); std::cout << "Try a list with two valid devices." << std::endl; using DoubleValidList = vtkm::ListTagBase; - TryExecuteWithList(DoubleValidList(), true); + TryExecuteTests(DoubleValidList(), true); std::cout << "Try a list with only invalid device." << std::endl; using SingleInvalidList = vtkm::ListTagBase; - TryExecuteWithList(SingleInvalidList(), false); + TryExecuteTests(SingleInvalidList(), false); std::cout << "Try a list with an invalid and valid device." << std::endl; using InvalidAndValidList = vtkm::ListTagBase; - TryExecuteWithList(InvalidAndValidList(), true); + TryExecuteTests(InvalidAndValidList(), true); } } // anonymous namespace diff --git a/vtkm/filter/ClipWithImplicitFunction.h b/vtkm/filter/ClipWithImplicitFunction.h index 98f32e918..2aef2ccc4 100644 --- a/vtkm/filter/ClipWithImplicitFunction.h +++ b/vtkm/filter/ClipWithImplicitFunction.h @@ -21,12 +21,10 @@ #ifndef vtk_m_filter_ClipWithImplicitFunction_h #define vtk_m_filter_ClipWithImplicitFunction_h -#include +#include #include #include -#include - 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 { public: - template - void SetImplicitFunction(const std::shared_ptr& func, - const vtkm::filter::PolicyBase& policy); - - template - void SetImplicitFunction(const std::shared_ptr& func) + void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func) { this->Function = func; } - std::shared_ptr GetImplicitFunction() const - { - return this->Function; - } + const vtkm::cont::ImplicitFunctionHandle& GetImplicitFunction() const { return this->Function; } template vtkm::filter::Result DoExecute(const vtkm::cont::DataSet& input, @@ -70,7 +60,7 @@ public: const DeviceAdapter& tag); private: - std::shared_ptr Function; + vtkm::cont::ImplicitFunctionHandle Function; vtkm::worklet::Clip Worklet; }; } diff --git a/vtkm/filter/ClipWithImplicitFunction.hxx b/vtkm/filter/ClipWithImplicitFunction.hxx index 5294ef55f..fdbd02b14 100644 --- a/vtkm/filter/ClipWithImplicitFunction.hxx +++ b/vtkm/filter/ClipWithImplicitFunction.hxx @@ -53,16 +53,6 @@ struct PointMapHelper } // end namespace clipwithimplicitfunction -//----------------------------------------------------------------------------- -template -inline void ClipWithImplicitFunction::SetImplicitFunction( - const std::shared_ptr& func, - const vtkm::filter::PolicyBase&) -{ - func->ResetDevices(DerivedPolicy::DeviceAdapterList); - this->Function = func; -} - //----------------------------------------------------------------------------- template 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; diff --git a/vtkm/filter/ExtractGeometry.h b/vtkm/filter/ExtractGeometry.h index 98a8551e6..2f36d53ab 100644 --- a/vtkm/filter/ExtractGeometry.h +++ b/vtkm/filter/ExtractGeometry.h @@ -21,7 +21,7 @@ #ifndef vtk_m_filter_ExtractGeometry_h #define vtk_m_filter_ExtractGeometry_h -#include +#include #include #include @@ -52,20 +52,12 @@ public: ExtractGeometry(); // Set the volume of interest to extract - template - void SetImplicitFunction(const std::shared_ptr& func, - const vtkm::filter::PolicyBase& policy); - - template - void SetImplicitFunction(const std::shared_ptr& func) + void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func) { this->Function = func; } - std::shared_ptr 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 Function; + vtkm::cont::ImplicitFunctionHandle Function; vtkm::worklet::ExtractGeometry Worklet; }; diff --git a/vtkm/filter/ExtractGeometry.hxx b/vtkm/filter/ExtractGeometry.hxx index 863e42214..837cff238 100644 --- a/vtkm/filter/ExtractGeometry.hxx +++ b/vtkm/filter/ExtractGeometry.hxx @@ -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 -inline void ExtractGeometry::SetImplicitFunction(const std::shared_ptr& func, - const vtkm::filter::PolicyBase&) -{ - func->ResetDevices(DerivedPolicy::DeviceAdapterList); - this->Function = func; -} - //----------------------------------------------------------------------------- inline VTKM_CONT ExtractGeometry::ExtractGeometry() : vtkm::filter::FilterDataSet() @@ -110,7 +101,7 @@ inline VTKM_CONT vtkm::filter::Result ExtractGeometry::DoExecute( CallWorker worker(outCells, this->Worklet, coords, - *this->Function, + this->Function, this->ExtractInside, this->ExtractBoundaryCells, this->ExtractOnlyBoundaryCells); diff --git a/vtkm/filter/ExtractPoints.h b/vtkm/filter/ExtractPoints.h index 87227239c..1d15907e9 100644 --- a/vtkm/filter/ExtractPoints.h +++ b/vtkm/filter/ExtractPoints.h @@ -21,7 +21,7 @@ #ifndef vtk_m_filter_ExtractPoints_h #define vtk_m_filter_ExtractPoints_h -#include +#include #include #include #include @@ -53,21 +53,12 @@ public: void SetCompactPoints(bool value) { this->CompactPoints = value; } /// Set the volume of interest to extract - template - void SetImplicitFunction(const std::shared_ptr& func, - const vtkm::filter::PolicyBase& policy); - - /// Set the volume of interest to extract - template - void SetImplicitFunction(const std::shared_ptr& func) + void SetImplicitFunction(const vtkm::cont::ImplicitFunctionHandle& func) { this->Function = func; } - std::shared_ptr 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 Function; + vtkm::cont::ImplicitFunctionHandle Function; bool CompactPoints; vtkm::filter::CleanGrid Compactor; diff --git a/vtkm/filter/ExtractPoints.hxx b/vtkm/filter/ExtractPoints.hxx index a55c403a2..62f5be33d 100644 --- a/vtkm/filter/ExtractPoints.hxx +++ b/vtkm/filter/ExtractPoints.hxx @@ -44,15 +44,6 @@ namespace vtkm namespace filter { -//----------------------------------------------------------------------------- -template -inline void ExtractPoints::SetImplicitFunction(const std::shared_ptr& func, - const vtkm::filter::PolicyBase&) -{ - func->ResetDevices(DerivedPolicy::DeviceAdapterList); - this->Function = func; -} - //----------------------------------------------------------------------------- inline VTKM_CONT ExtractPoints::ExtractPoints() : vtkm::filter::FilterDataSet() @@ -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); diff --git a/vtkm/filter/FilterDataSet.hxx b/vtkm/filter/FilterDataSet.hxx index f56e2f684..c735f7009 100644 --- a/vtkm/filter/FilterDataSet.hxx +++ b/vtkm/filter/FilterDataSet.hxx @@ -139,11 +139,11 @@ inline VTKM_CONT bool FilterDataSet::MapFieldOntoOutput( { vtkm::filter::FieldMetadata metaData(field); typedef internal::ResolveFieldTypeAndMap FunctorType; - FunctorType functor( - static_cast(this), result, metaData, policy, this->Tracker, valid); + FunctorType functor(static_cast(this), result, metaData, policy, valid); typedef vtkm::filter::FilterTraits 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 diff --git a/vtkm/filter/FilterDataSetWithField.hxx b/vtkm/filter/FilterDataSetWithField.hxx index e03c5dea0..1b3981886 100644 --- a/vtkm/filter/FilterDataSetWithField.hxx +++ b/vtkm/filter/FilterDataSetWithField.hxx @@ -162,10 +162,11 @@ inline VTKM_CONT Result FilterDataSetWithField::PrepareForExecution( Result result; typedef internal::ResolveFieldTypeAndExecute FunctorType; - FunctorType functor(static_cast(this), input, metaData, policy, this->Tracker, result); + FunctorType functor(static_cast(this), input, metaData, policy, result); typedef vtkm::filter::FilterTraits 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::PrepareForExecution( //determine the field type first Result result; typedef internal::ResolveFieldTypeAndExecute FunctorType; - FunctorType functor(static_cast(this), input, metaData, policy, this->Tracker, result); + FunctorType functor(static_cast(this), input, metaData, policy, result); typedef vtkm::filter::FilterTraits 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::MapFieldOntoOutput( { vtkm::filter::FieldMetadata metaData(field); typedef internal::ResolveFieldTypeAndMap FunctorType; - FunctorType functor( - static_cast(this), result, metaData, policy, this->Tracker, valid); + FunctorType functor(static_cast(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 diff --git a/vtkm/filter/FilterField.hxx b/vtkm/filter/FilterField.hxx index 9cf06885d..7f1d5bf29 100644 --- a/vtkm/filter/FilterField.hxx +++ b/vtkm/filter/FilterField.hxx @@ -159,10 +159,11 @@ FilterField::PrepareForExecution(const vtkm::cont::DataSet& input, Result result; typedef internal::ResolveFieldTypeAndExecute FunctorType; - FunctorType functor(static_cast(this), input, metaData, policy, this->Tracker, result); + FunctorType functor(static_cast(this), input, metaData, policy, result); typedef vtkm::filter::FilterTraits 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::PrepareForExecution(const vtkm::cont::DataSet& input, Result result; typedef internal::ResolveFieldTypeAndExecute FunctorType; - FunctorType functor(static_cast(this), input, metaData, policy, this->Tracker, result); + FunctorType functor(static_cast(this), input, metaData, policy, result); typedef vtkm::filter::FilterTraits 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; } } diff --git a/vtkm/filter/internal/ResolveFieldTypeAndExecute.h b/vtkm/filter/internal/ResolveFieldTypeAndExecute.h index 61d9bdd4f..98bf1847a 100644 --- a/vtkm/filter/internal/ResolveFieldTypeAndExecute.h +++ b/vtkm/filter/internal/ResolveFieldTypeAndExecute.h @@ -35,6 +35,18 @@ namespace filter namespace internal { +struct ResolveFieldTypeAndExecuteForDevice +{ + template + 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 struct ResolveFieldTypeAndExecute { @@ -44,56 +56,28 @@ struct ResolveFieldTypeAndExecute const vtkm::cont::DataSet& InputData; const vtkm::filter::FieldMetadata& Metadata; const vtkm::filter::PolicyBase& Policy; - vtkm::cont::RuntimeDeviceTracker Tracker; ResultType& Result; ResolveFieldTypeAndExecute(Derived* derivedClass, const vtkm::cont::DataSet& inputData, const vtkm::filter::FieldMetadata& fieldMeta, const vtkm::filter::PolicyBase& policy, - const vtkm::cont::RuntimeDeviceTracker& tracker, ResultType& result) : DerivedClass(derivedClass) , InputData(inputData) , Metadata(fieldMeta) , Policy(policy) - , Tracker(tracker) , Result(result) { } -private: template - struct ResolveFieldTypeAndExecuteForDevice + void operator()(const vtkm::cont::ArrayHandle& field, + vtkm::cont::RuntimeDeviceTracker& tracker) const { - typedef vtkm::cont::ArrayHandle FieldArrayHandle; - ResolveFieldTypeAndExecuteForDevice(const Self& instance, const FieldArrayHandle& field) - : Instance(instance) - , Field(field) - { - } - - const Self& Instance; - const vtkm::cont::ArrayHandle& Field; - - template - 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&) = delete; - }; - -public: - template - void operator()(const vtkm::cont::ArrayHandle& field) const - { - ResolveFieldTypeAndExecuteForDevice 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: diff --git a/vtkm/filter/internal/ResolveFieldTypeAndMap.h b/vtkm/filter/internal/ResolveFieldTypeAndMap.h index 632519c5d..379aef2ae 100644 --- a/vtkm/filter/internal/ResolveFieldTypeAndMap.h +++ b/vtkm/filter/internal/ResolveFieldTypeAndMap.h @@ -43,6 +43,15 @@ namespace filter { namespace internal { +struct ResolveFieldTypeAndMapForDevice +{ + template + bool operator()(DeviceAdapterTag tag, InstanceType&& instance, FieldType&& field) const + { + return instance.DerivedClass->DoMapField( + instance.InputResult, field, instance.Metadata, instance.Policy, tag); + } +}; template struct ResolveFieldTypeAndMap @@ -53,63 +62,28 @@ struct ResolveFieldTypeAndMap vtkm::filter::Result& InputResult; const vtkm::filter::FieldMetadata& Metadata; const vtkm::filter::PolicyBase& Policy; - vtkm::cont::RuntimeDeviceTracker Tracker; bool& RanProperly; ResolveFieldTypeAndMap(Derived* derivedClass, vtkm::filter::Result& inResult, const vtkm::filter::FieldMetadata& fieldMeta, const vtkm::filter::PolicyBase& policy, - const vtkm::cont::RuntimeDeviceTracker& tracker, bool& ran) : DerivedClass(derivedClass) , InputResult(inResult) , Metadata(fieldMeta) , Policy(policy) - , Tracker(tracker) , RanProperly(ran) { } -private: template - struct ResolveFieldTypeAndMapForDevice + void operator()(const vtkm::cont::ArrayHandle& field, + vtkm::cont::RuntimeDeviceTracker& tracker) const { - typedef vtkm::cont::ArrayHandle FieldArrayHandle; - ResolveFieldTypeAndMapForDevice(const Self& instance, const FieldArrayHandle& field) - : Instance(instance) - , Field(field) - , Valid(false) - { - } - - const Self& Instance; - const vtkm::cont::ArrayHandle& Field; - mutable bool Valid; - - template - 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&) = delete; - }; - -public: - template - void operator()(const vtkm::cont::ArrayHandle& field) const - { - ResolveFieldTypeAndMapForDevice 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: diff --git a/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx b/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx index 0ce3bb440..1a7c0fc3d 100644 --- a/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx +++ b/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx @@ -60,11 +60,10 @@ void TestClipStructured() vtkm::Vec center(1, 1, 0); vtkm::FloatDefault radius(0.5); - auto sphere = std::make_shared(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")); diff --git a/vtkm/filter/testing/UnitTestExtractGeometryFilter.cxx b/vtkm/filter/testing/UnitTestExtractGeometryFilter.cxx index a6c77e52a..33621da5c 100644 --- a/vtkm/filter/testing/UnitTestExtractGeometryFilter.cxx +++ b/vtkm/filter/testing/UnitTestExtractGeometryFilter.cxx @@ -40,7 +40,7 @@ public: // Implicit function vtkm::Vec minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(0.5f, 0.5f, 0.5f); vtkm::Vec maxPoint(3.5f, 3.5f, 3.5f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(0.5f, 0.5f, 0.5f); vtkm::Vec maxPoint(3.5f, 3.5f, 3.5f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint); // Setup and run filter to extract by volume of interest vtkm::filter::ExtractGeometry extractGeometry; diff --git a/vtkm/filter/testing/UnitTestExtractPointsFilter.cxx b/vtkm/filter/testing/UnitTestExtractPointsFilter.cxx index 7966f60d8..8d7dbd3cb 100644 --- a/vtkm/filter/testing/UnitTestExtractPointsFilter.cxx +++ b/vtkm/filter/testing/UnitTestExtractPointsFilter.cxx @@ -40,7 +40,7 @@ public: // Implicit function vtkm::Vec minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 center(2.f, 2.f, 2.f); vtkm::FloatDefault radius(1.8f); - auto sphere = std::make_shared(center, radius); + auto sphere = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(0.f, 0.f, 0.f); vtkm::Vec maxPoint(1.f, 1.f, 1.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(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 minPoint(0.f, 0.f, 0.f); vtkm::Vec maxPoint(1.f, 1.f, 1.f); - auto box = std::make_shared(minPoint, maxPoint); + auto box = vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint); // Setup and run filter to extract by volume of interest vtkm::filter::ExtractPoints extractPoints; diff --git a/vtkm/internal/ExportMacros.h b/vtkm/internal/ExportMacros.h index 43b0e6837..218160427 100644 --- a/vtkm/internal/ExportMacros.h +++ b/vtkm/internal/ExportMacros.h @@ -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 diff --git a/vtkm/internal/ListTagDetail.h b/vtkm/internal/ListTagDetail.h index c0fbce138..0d812f609 100644 --- a/vtkm/internal/ListTagDetail.h +++ b/vtkm/internal/ListTagDetail.h @@ -165,30 +165,30 @@ struct ListIntersect using type = SameListTag; }; -template -VTKM_CONT void ListForEachImpl(Functor&&, brigand::empty_sequence) +template +VTKM_CONT void ListForEachImpl(Functor&&, brigand::list<>, Args&&...) { } -template -VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list) +template +VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list, Args&&... args) { - f(T1{}); + f(T1{}, std::forward(args)...); } -template -VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list) +template +VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list, Args&&... args) { - f(T1{}); - f(T2{}); + f(T1{}, std::forward(args)...); + f(T2{}, std::forward(args)...); } -template -VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list) +template +VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list, Args&&... args) { - f(T1{}); - f(T2{}); - f(T3{}); + f(T1{}, std::forward(args)...); + f(T2{}, std::forward(args)...); + f(T3{}, std::forward(args)...); } template -VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list) + typename... ArgTypes, + typename... Args> +VTKM_CONT void ListForEachImpl(Functor&& f, + brigand::list&&, + Args&&... args) { - f(T1{}); - f(T2{}); - f(T3{}); - f(T4{}); - ListForEachImpl(f, brigand::list()); + f(T1{}, std::forward(args)...); + f(T2{}, std::forward(args)...); + f(T3{}, std::forward(args)...); + f(T4{}, std::forward(args)...); + ListForEachImpl( + std::forward(f), brigand::list{}, std::forward(args)...); } + +template +struct ListCrossProductAppend +{ + using type = brigand::push_back>; +}; + +template +struct ListCrossProductImplUnrollR2 +{ + using P = + brigand::fold, + ListCrossProductAppend>>; + + using type = brigand::append; +}; + +template +struct ListCrossProductImpl +{ + using type = brigand::fold< + R2, + brigand::list<>, + ListCrossProductImplUnrollR2>>; +}; + + + } // namespace detail //----------------------------------------------------------------------------- diff --git a/vtkm/rendering/AxisAnnotation2D.h b/vtkm/rendering/AxisAnnotation2D.h index 1c7da50d4..1ae48e751 100644 --- a/vtkm/rendering/AxisAnnotation2D.h +++ b/vtkm/rendering/AxisAnnotation2D.h @@ -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 diff --git a/vtkm/rendering/AxisAnnotation3D.cxx b/vtkm/rendering/AxisAnnotation3D.cxx index 38322e1a1..94a0b3913 100644 --- a/vtkm/rendering/AxisAnnotation3D.cxx +++ b/vtkm/rendering/AxisAnnotation3D.cxx @@ -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; } } diff --git a/vtkm/rendering/AxisAnnotation3D.h b/vtkm/rendering/AxisAnnotation3D.h index f3870541d..6fa953d01 100644 --- a/vtkm/rendering/AxisAnnotation3D.h +++ b/vtkm/rendering/AxisAnnotation3D.h @@ -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 diff --git a/vtkm/rendering/CanvasEGL.h b/vtkm/rendering/CanvasEGL.h index ae163cb3e..d55cf0d67 100644 --- a/vtkm/rendering/CanvasEGL.h +++ b/vtkm/rendering/CanvasEGL.h @@ -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 Internals; diff --git a/vtkm/rendering/CanvasGL.h b/vtkm/rendering/CanvasGL.h index d9a1a368e..f94deb4db 100644 --- a/vtkm/rendering/CanvasGL.h +++ b/vtkm/rendering/CanvasGL.h @@ -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& point0, const vtkm::Vec& 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& point0, const vtkm::Vec& point1, const vtkm::Vec& point2, const vtkm::Vec& point3, - const vtkm::rendering::Color& color) const VTKM_OVERRIDE; + const vtkm::rendering::Color& color) const override; void AddText(const vtkm::Vec& position, vtkm::Float32 scale, @@ -82,7 +82,7 @@ protected: vtkm::Float32 windowAspect, const vtkm::Vec& anchor, const vtkm::rendering::Color& color, - const std::string& text) const VTKM_OVERRIDE; + const std::string& text) const override; private: vtkm::rendering::BitmapFont Font; diff --git a/vtkm/rendering/CanvasOSMesa.h b/vtkm/rendering/CanvasOSMesa.h index a72697876..208659f35 100644 --- a/vtkm/rendering/CanvasOSMesa.h +++ b/vtkm/rendering/CanvasOSMesa.h @@ -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 Internals; diff --git a/vtkm/rendering/CanvasRayTracer.h b/vtkm/rendering/CanvasRayTracer.h index 42604dfa9..fd8c44ba1 100644 --- a/vtkm/rendering/CanvasRayTracer.h +++ b/vtkm/rendering/CanvasRayTracer.h @@ -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& rays, const vtkm::cont::ArrayHandle& colors, diff --git a/vtkm/rendering/MapperConnectivity.h b/vtkm/rendering/MapperConnectivity.h index 5b3310f8d..c2c8e1cb6 100644 --- a/vtkm/rendering/MapperConnectivity.h +++ b/vtkm/rendering/MapperConnectivity.h @@ -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: diff --git a/vtkm/rendering/MapperGL.h b/vtkm/rendering/MapperGL.h index dd48c60f4..7253ddbc9 100644 --- a/vtkm/rendering/MapperGL.h +++ b/vtkm/rendering/MapperGL.h @@ -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; diff --git a/vtkm/rendering/MapperRayTracer.h b/vtkm/rendering/MapperRayTracer.h index 55317a2bc..5d26cdb27 100644 --- a/vtkm/rendering/MapperRayTracer.h +++ b/vtkm/rendering/MapperRayTracer.h @@ -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; diff --git a/vtkm/rendering/MapperVolume.h b/vtkm/rendering/MapperVolume.h index 2e11f1039..70ae39ea9 100644 --- a/vtkm/rendering/MapperVolume.h +++ b/vtkm/rendering/MapperVolume.h @@ -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); diff --git a/vtkm/rendering/MapperWireframer.h b/vtkm/rendering/MapperWireframer.h index fcbfc3d9a..32d0afec6 100644 --- a/vtkm/rendering/MapperWireframer.h +++ b/vtkm/rendering/MapperWireframer.h @@ -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; diff --git a/vtkm/rendering/TextAnnotationBillboard.h b/vtkm/rendering/TextAnnotationBillboard.h index 57b23ac59..687bb57e8 100644 --- a/vtkm/rendering/TextAnnotationBillboard.h +++ b/vtkm/rendering/TextAnnotationBillboard.h @@ -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 diff --git a/vtkm/rendering/TextAnnotationScreen.h b/vtkm/rendering/TextAnnotationScreen.h index f9f406b30..04cebd99c 100644 --- a/vtkm/rendering/TextAnnotationScreen.h +++ b/vtkm/rendering/TextAnnotationScreen.h @@ -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 diff --git a/vtkm/rendering/View1D.h b/vtkm/rendering/View1D.h index c877af173..7923f05e6 100644 --- a/vtkm/rendering/View1D.h +++ b/vtkm/rendering/View1D.h @@ -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(); diff --git a/vtkm/rendering/View2D.h b/vtkm/rendering/View2D.h index f50dd1912..669dd1107 100644 --- a/vtkm/rendering/View2D.h +++ b/vtkm/rendering/View2D.h @@ -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(); diff --git a/vtkm/rendering/View3D.h b/vtkm/rendering/View3D.h index ed0e7cbb5..37644c83c 100644 --- a/vtkm/rendering/View3D.h +++ b/vtkm/rendering/View3D.h @@ -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 diff --git a/vtkm/rendering/WorldAnnotatorGL.h b/vtkm/rendering/WorldAnnotatorGL.h index f7a9973f5..6dd5b5ec8 100644 --- a/vtkm/rendering/WorldAnnotatorGL.h +++ b/vtkm/rendering/WorldAnnotatorGL.h @@ -42,7 +42,7 @@ public: const vtkm::Vec& point1, vtkm::Float32 lineWidth, const vtkm::rendering::Color& color, - bool inFront) const VTKM_OVERRIDE; + bool inFront) const override; void AddText(const vtkm::Vec& origin, const vtkm::Vec& right, @@ -50,7 +50,7 @@ public: vtkm::Float32 scale, const vtkm::Vec& anchor, const vtkm::rendering::Color& color, - const std::string& text) const VTKM_OVERRIDE; + const std::string& text) const override; private: BitmapFont Font; diff --git a/vtkm/rendering/internal/RunTriangulator.cxx b/vtkm/rendering/internal/RunTriangulator.cxx index e36325d24..ec7de2063 100644 --- a/vtkm/rendering/internal/RunTriangulator.cxx +++ b/vtkm/rendering/internal/RunTriangulator.cxx @@ -59,7 +59,7 @@ struct TriangulatorFunctor void RunTriangulator(const vtkm::cont::DynamicCellSet& cellSet, vtkm::cont::ArrayHandle>& 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? diff --git a/vtkm/rendering/internal/RunTriangulator.h b/vtkm/rendering/internal/RunTriangulator.h index f10994e9d..a8bf4ea57 100644 --- a/vtkm/rendering/internal/RunTriangulator.h +++ b/vtkm/rendering/internal/RunTriangulator.h @@ -43,7 +43,7 @@ void RunTriangulator( const vtkm::cont::DynamicCellSet& cellSet, vtkm::cont::ArrayHandle>& indices, vtkm::Id& numberOfTriangles, - const vtkm::cont::RuntimeDeviceTracker& tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()); + vtkm::cont::RuntimeDeviceTracker tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker()); } } } // namespace vtkm::rendering::internal diff --git a/vtkm/rendering/raytracing/Camera.cxx b/vtkm/rendering/raytracing/Camera.cxx index ec32bfc8f..05fd9d379 100644 --- a/vtkm/rendering/raytracing/Camera.cxx +++ b/vtkm/rendering/raytracing/Camera.cxx @@ -741,6 +741,9 @@ VTKM_CONT void Camera::CreateRaysOnDevice(Ray& rays, vtkm::worklet::DispatcherMapField, Device>(MemSet(0.f)) .Invoke(rays.MinDistance); + vtkm::worklet::DispatcherMapField, Device>(MemSet(0.f)) + .Invoke(rays.Distance); + //Reset the Rays Hit Index to -2 vtkm::worklet::DispatcherMapField, Device>(MemSet(-2)) .Invoke(rays.HitIdx); diff --git a/vtkm/rendering/raytracing/ChannelBuffer.cxx b/vtkm/rendering/raytracing/ChannelBuffer.cxx index 2d5b7a647..b34cab808 100644 --- a/vtkm/rendering/raytracing/ChannelBuffer.cxx +++ b/vtkm/rendering/raytracing/ChannelBuffer.cxx @@ -100,10 +100,10 @@ ChannelBuffer::ChannelBuffer() template ChannelBuffer::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 ChannelBuffer::GetChannel(const vtkm::Int32 throw vtkm::cont::ErrorBadValue("ChannelBuffer: invalid channel to extract"); ChannelBuffer output(1, this->Size); output.SetName(this->Name); - + if (this->Size == 0) + { + return output; + } ExtractChannelFunctor functor(this, output.Buffer, channel); vtkm::cont::TryExecute(functor); @@ -373,7 +376,6 @@ public: { InvDeltaScalar = 1.f / (maxScalar - minScalar); } - //std::cout<<"Min scalar "<); typedef void ExecutionSignature(_1); diff --git a/vtkm/rendering/raytracing/ConnectivityTracer.h b/vtkm/rendering/raytracing/ConnectivityTracer.h index 4a0a8064a..510839ea4 100644 --- a/vtkm/rendering/raytracing/ConnectivityTracer.h +++ b/vtkm/rendering/raytracing/ConnectivityTracer.h @@ -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 + class AddPathLengths : public vtkm::worklet::WorkletMapField + { + public: + VTKM_CONT + AddPathLengths() {} + + typedef void ControlSignature(FieldIn, // ray status + FieldIn, // cell enter distance + FieldIn, // cell exit distance + FieldInOut); // 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 class Integrate : public vtkm::worklet::WorkletMapField { @@ -1138,6 +1170,17 @@ public: this->IntersectTime += timer.GetElapsedTime(); } + template + void AccumulatePathLengths(Ray& rays, detail::RayTracking& tracker, Device) + { + vtkm::worklet::DispatcherMapField, Device>( + AddPathLengths()) + .Invoke(rays.Status, + *(tracker.EnterDist), + *(tracker.ExitDist), + rays.GetBuffer("path_lengths").Buffer); + } + template void FindLostRays(Ray& rays, detail::RayTracking& tracker, Device) { @@ -1210,7 +1253,6 @@ public: bool divideEmisByAbsorp = false; vtkm::cont::ArrayHandle absorp = rays.Buffers.at(0).Buffer; vtkm::cont::ArrayHandle emission = rays.GetBuffer("emission").Buffer; - vtkm::worklet::DispatcherMapField, Device>( IntegrateEmission(rays.Buffers.at(0).GetNumChannels(), divideEmisByAbsorp)) .Invoke(rays.Status, @@ -1242,13 +1284,27 @@ public: template void PrintDebugRay(Ray& 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& rays, Device) { vtkm::worklet::DispatcherMapField, Device>( - AdvanceRay(FloatType(0.0000001))) + AdvanceRay(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 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 diff --git a/vtkm/rendering/raytracing/MortonCodes.h b/vtkm/rendering/raytracing/MortonCodes.h index 392d1dd80..9d259e13e 100644 --- a/vtkm/rendering/raytracing/MortonCodes.h +++ b/vtkm/rendering/raytracing/MortonCodes.h @@ -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); diff --git a/vtkm/rendering/raytracing/Ray.h b/vtkm/rendering/raytracing/Ray.h index ae7c7d6c6..dff325e95 100644 --- a/vtkm/rendering/raytracing/Ray.h +++ b/vtkm/rendering/raytracing/Ray.h @@ -181,9 +181,28 @@ public: VTKM_CONT Ray(const vtkm::Int32 size, Device, bool enableIntersectionData = false) { NumRays = size; - IntersectionDataEnabled = enableIntersectionData; + ChannelBuffer 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 + 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 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(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) { diff --git a/vtkm/testing/UnitTestListTag.cxx b/vtkm/testing/UnitTestListTag.cxx index b1ad73191..3e1a68d82 100644 --- a/vtkm/testing/UnitTestListTag.cxx +++ b/vtkm/testing/UnitTestListTag.cxx @@ -33,10 +33,6 @@ namespace template struct TestClass { - enum - { - NUMBER = N - }; }; struct TestListTag1 : vtkm::ListTagBase> @@ -63,36 +59,50 @@ struct TestListTagIntersect : vtkm::ListTagIntersect +{ +}; + struct TestListTagUniversal : vtkm::ListTagUniversal { }; +template +std::pair test_number(std::pair, TestClass>) +{ + return std::make_pair(N, M); +} + +template +int test_number(TestClass) +{ + return N; +} + +template struct MutableFunctor { - std::vector FoundTypes; + std::vector FoundTypes; - template - VTKM_CONT void operator()(T) + template + VTKM_CONT void operator()(U u) { - this->FoundTypes.push_back(T::NUMBER); + this->FoundTypes.push_back(test_number(u)); } }; -std::vector g_FoundType; - +template struct ConstantFunctor { - ConstantFunctor() { g_FoundType.erase(g_FoundType.begin(), g_FoundType.end()); } - - template - VTKM_CONT void operator()(T) const + template + VTKM_CONT void operator()(U u, VectorType& vector) const { - g_FoundType.push_back(T::NUMBER); + vector.push_back(test_number(u)); } }; -template -void CheckSame(const vtkm::Vec& expected, const std::vector& found) +template +void CheckSame(const vtkm::Vec& expected, const std::vector& found) { VTKM_TEST_ASSERT(static_cast(found.size()) == N, "Got wrong number of items."); @@ -137,13 +147,15 @@ void TryList(const vtkm::Vec& expected, ListTag) VTKM_IS_LIST_TAG(ListTag); std::cout << " Try mutable for each" << std::endl; - MutableFunctor functor; + MutableFunctor 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 foundTypes; + ConstantFunctor 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& expected, ListTag) CheckContains(TestClass<43>(), ListTag(), functor.FoundTypes); CheckContains(TestClass<44>(), ListTag(), functor.FoundTypes); } +template +void TryList(const vtkm::Vec, N>& expected, ListTag) +{ + VTKM_IS_LIST_TAG(ListTag); + + std::cout << " Try mutable for each" << std::endl; + MutableFunctor> functor; + vtkm::ListForEach(functor, ListTag()); + CheckSame(expected, functor.FoundTypes); + + std::cout << " Try constant for each" << std::endl; + std::vector> foundTypes; + ConstantFunctor> cfunc; + vtkm::ListForEach(cfunc, ListTag(), foundTypes); + CheckSame(expected, foundTypes); +} template void TryList(const vtkm::Vec&, TestListTagUniversal tag) @@ -207,6 +235,12 @@ void TestLists() std::cout << "ListTagIntersect" << std::endl; TryList(vtkm::Vec(31, 32, 33), TestListTagIntersect()); + std::cout << "ListTagCrossProduct" << std::endl; + TryList(vtkm::Vec, 3>({ 31, 11 }, { 32, 11 }, { 33, 11 }), + TestListTagCrossProduct()); + + + std::cout << "ListTagUniversal" << std::endl; TryList(vtkm::Vec(1, 2, 3, 4), TestListTagUniversal()); } diff --git a/vtkm/worklet/Clip.h b/vtkm/worklet/Clip.h index 3bea0a996..d43a7531f 100644 --- a/vtkm/worklet/Clip.h +++ b/vtkm/worklet/Clip.h @@ -30,7 +30,7 @@ #include #include #include -#include +#include #include #include @@ -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 - clipScalars(handle, this->Function); + vtkm::cont::ArrayHandleTransform 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 vtkm::cont::CellSetExplicit<> Run(const vtkm::cont::DynamicCellSetBase& cellSet, - const vtkm::cont::ImplicitFunction& clipFunction, + const vtkm::cont::ImplicitFunctionHandle& clipFunction, const vtkm::cont::CoordinateSystem& coords, DeviceAdapter device) { diff --git a/vtkm/worklet/ExtractGeometry.h b/vtkm/worklet/ExtractGeometry.h index b5da5d916..c744d61a8 100644 --- a/vtkm/worklet/ExtractGeometry.h +++ b/vtkm/worklet/ExtractGeometry.h @@ -27,7 +27,7 @@ #include #include #include -#include +#include 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(indx)]; vtkm::Vec 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 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, diff --git a/vtkm/worklet/ExtractPoints.h b/vtkm/worklet/ExtractPoints.h index 0e9015e0c..fb0a7f54e 100644 --- a/vtkm/worklet/ExtractPoints.h +++ b/vtkm/worklet/ExtractPoints.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include 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& 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 vtkm::cont::CellSetSingleType<> Run(const CellSetType& cellSet, const CoordinateType& coordinates, - const vtkm::cont::ImplicitFunction& implicitFunction, + const vtkm::cont::ImplicitFunctionHandle& implicitFunction, bool extractInside, DeviceAdapter device) { diff --git a/vtkm/worklet/MarchingCubes.h b/vtkm/worklet/MarchingCubes.h index 9dc669bda..02d056775 100644 --- a/vtkm/worklet/MarchingCubes.h +++ b/vtkm/worklet/MarchingCubes.h @@ -686,11 +686,11 @@ struct GenerateNormalsDeduced // NormalsWorkletPass1 pass1(*edges); vtkm::worklet::DispatcherMapTopology(pass1).Invoke( - *cellset, *cellset, coordinates, *field, *normals); + *cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *normals); NormalsWorkletPass2 pass2(*edges); vtkm::worklet::DispatcherMapTopology(pass2).Invoke( - *cellset, *cellset, coordinates, *field, *weights, *normals); + *cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *weights, *normals); } }; diff --git a/vtkm/worklet/VertexClustering.h b/vtkm/worklet/VertexClustering.h index 38474ba0e..92d11711a 100644 --- a/vtkm/worklet/VertexClustering.h +++ b/vtkm/worklet/VertexClustering.h @@ -78,28 +78,20 @@ struct SelectRepresentativePoint : public vtkm::worklet::WorkletReduceByKey return pointsIn[pointsIn.GetNumberOfComponents() / 2]; } - template struct RunTrampoline { - const vtkm::worklet::Keys& Keys; - vtkm::cont::DynamicArrayHandle& OutputPoints; - - VTKM_CONT - RunTrampoline(const vtkm::worklet::Keys& keys, vtkm::cont::DynamicArrayHandle& output) - : Keys(keys) - , OutputPoints(output) + template + VTKM_CONT void operator()(const InputPointsArrayType& points, + const vtkm::worklet::Keys& keys, + vtkm::cont::DynamicArrayHandle& output, + DeviceAdapterTag) const { - } - template - VTKM_CONT void operator()(const InputPointsArrayType& points) const - { vtkm::cont::ArrayHandle out; - vtkm::worklet::DispatcherReduceByKey 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& keys, const InputDynamicPointsArrayType& inputPoints, - DeviceAdapterTag) + DeviceAdapterTag tag) { vtkm::cont::DynamicArrayHandle output; - RunTrampoline trampoline(keys, output); - vtkm::cont::CastAndCall(inputPoints, trampoline); + RunTrampoline trampoline; + vtkm::cont::CastAndCall(inputPoints, trampoline, keys, output, tag); return output; } }; diff --git a/vtkm/worklet/testing/UnitTestClipping.cxx b/vtkm/worklet/testing/UnitTestClipping.cxx index 484782c5e..993442131 100644 --- a/vtkm/worklet/testing/UnitTestClipping.cxx +++ b/vtkm/worklet/testing/UnitTestClipping.cxx @@ -30,7 +30,7 @@ #include #include #include -#include +#include #include #include @@ -243,13 +243,15 @@ void TestClippingWithImplicitFunction() { vtkm::Vec center(1, 1, 0); vtkm::FloatDefault radius(0.5); - vtkm::cont::Sphere sphere(center, radius); vtkm::cont::DataSet ds = MakeTestDatasetStructured(); vtkm::worklet::Clip clip; vtkm::cont::CellSetExplicit<> outputCellSet = - clip.Run(ds.GetCellSet(0), sphere, ds.GetCoordinateSystem("coords"), DeviceAdapter()); + clip.Run(ds.GetCellSet(0), + vtkm::cont::make_ImplicitFunctionHandle(center, radius), + ds.GetCoordinateSystem("coords"), + DeviceAdapter()); vtkm::cont::ArrayHandleUniformPointCoordinates coordsIn; ds.GetCoordinateSystem("coords").GetData().CopyTo(coordsIn); diff --git a/vtkm/worklet/testing/UnitTestExtractGeometry.cxx b/vtkm/worklet/testing/UnitTestExtractGeometry.cxx index a646e99af..c1af77958 100644 --- a/vtkm/worklet/testing/UnitTestExtractGeometry.cxx +++ b/vtkm/worklet/testing/UnitTestExtractGeometry.cxx @@ -77,7 +77,6 @@ public: // Implicit function vtkm::Vec minPoint(0.5f, 0.0f, 0.0f); vtkm::Vec maxPoint(2.0f, 2.0f, 2.0f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = true; bool extractBoundaryCells = false; @@ -88,7 +87,7 @@ public: vtkm::cont::DynamicCellSet outCellSet = extractGeometry.Run(cellSet, dataset.GetCoordinateSystem("coordinates"), - box, + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), extractInside, extractBoundaryCells, extractOnlyBoundaryCells, @@ -187,7 +186,6 @@ public: // Implicit function vtkm::Vec minPoint(1.0f, 1.0f, 1.0f); vtkm::Vec maxPoint(3.0f, 3.0f, 3.0f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = true; bool extractBoundaryCells = false; @@ -198,7 +196,7 @@ public: vtkm::cont::DynamicCellSet outCellSet = extractGeometry.Run(cellSet, dataset.GetCoordinateSystem("coords"), - box, + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), extractInside, extractBoundaryCells, extractOnlyBoundaryCells, @@ -229,7 +227,6 @@ public: // Implicit function vtkm::Vec center(2.f, 2.f, 2.f); vtkm::FloatDefault radius(1.8f); - vtkm::cont::Sphere sphere(center, radius); bool extractInside = true; bool extractBoundaryCells = false; @@ -240,7 +237,7 @@ public: vtkm::cont::DynamicCellSet outCellSet = extractGeometry.Run(cellSet, dataset.GetCoordinateSystem("coords"), - sphere, + vtkm::cont::make_ImplicitFunctionHandle(center, radius), extractInside, extractBoundaryCells, extractOnlyBoundaryCells, diff --git a/vtkm/worklet/testing/UnitTestExtractPoints.cxx b/vtkm/worklet/testing/UnitTestExtractPoints.cxx index 2be56b703..3227aad52 100644 --- a/vtkm/worklet/testing/UnitTestExtractPoints.cxx +++ b/vtkm/worklet/testing/UnitTestExtractPoints.cxx @@ -68,7 +68,6 @@ public: // Implicit function vtkm::Vec minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = true; // Output dataset contains input coordinate system and point data @@ -77,11 +76,12 @@ public: // Output data set with cell set containing extracted points vtkm::worklet::ExtractPoints extractPoints; - OutCellSetType outCellSet = extractPoints.Run(dataset.GetCellSet(0), - dataset.GetCoordinateSystem("coords"), - box, - extractInside, - DeviceAdapter()); + OutCellSetType outCellSet = + extractPoints.Run(dataset.GetCellSet(0), + dataset.GetCoordinateSystem("coords"), + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), + extractInside, + DeviceAdapter()); outDataSet.AddCellSet(outCellSet); VTKM_TEST_ASSERT(test_equal(outCellSet.GetNumberOfCells(), 27), @@ -100,7 +100,6 @@ public: // Implicit function vtkm::Vec minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = false; // Output dataset contains input coordinate system and point data @@ -109,11 +108,12 @@ public: // Output data set with cell set containing extracted points vtkm::worklet::ExtractPoints extractPoints; - OutCellSetType outCellSet = extractPoints.Run(dataset.GetCellSet(0), - dataset.GetCoordinateSystem("coords"), - box, - extractInside, - DeviceAdapter()); + OutCellSetType outCellSet = + extractPoints.Run(dataset.GetCellSet(0), + dataset.GetCoordinateSystem("coords"), + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), + extractInside, + DeviceAdapter()); outDataSet.AddCellSet(outCellSet); VTKM_TEST_ASSERT(test_equal(outCellSet.GetNumberOfCells(), 98), @@ -132,7 +132,6 @@ public: // Implicit function vtkm::Vec center(2.f, 2.f, 2.f); vtkm::FloatDefault radius(1.8f); - vtkm::cont::Sphere sphere(center, radius); bool extractInside = true; // Output dataset contains input coordinate system and point data @@ -141,11 +140,12 @@ public: // Output data set with cell set containing extracted points vtkm::worklet::ExtractPoints extractPoints; - OutCellSetType outCellSet = extractPoints.Run(dataset.GetCellSet(0), - dataset.GetCoordinateSystem("coords"), - sphere, - extractInside, - DeviceAdapter()); + OutCellSetType outCellSet = + extractPoints.Run(dataset.GetCellSet(0), + dataset.GetCoordinateSystem("coords"), + vtkm::cont::make_ImplicitFunctionHandle(center, radius), + extractInside, + DeviceAdapter()); outDataSet.AddCellSet(outCellSet); VTKM_TEST_ASSERT(test_equal(outCellSet.GetNumberOfCells(), 27), @@ -164,7 +164,6 @@ public: // Implicit function vtkm::Vec minPoint(0.f, 0.f, 0.f); vtkm::Vec maxPoint(1.f, 1.f, 1.f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = true; // Output dataset contains input coordinate system and point data @@ -173,11 +172,12 @@ public: // Output data set with cell set containing extracted points vtkm::worklet::ExtractPoints extractPoints; - OutCellSetType outCellSet = extractPoints.Run(dataset.GetCellSet(0), - dataset.GetCoordinateSystem("coordinates"), - box, - extractInside, - DeviceAdapter()); + OutCellSetType outCellSet = + extractPoints.Run(dataset.GetCellSet(0), + dataset.GetCoordinateSystem("coordinates"), + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), + extractInside, + DeviceAdapter()); outDataSet.AddCellSet(outCellSet); VTKM_TEST_ASSERT(test_equal(outCellSet.GetNumberOfCells(), 8), @@ -196,7 +196,6 @@ public: // Implicit function vtkm::Vec minPoint(0.f, 0.f, 0.f); vtkm::Vec maxPoint(1.f, 1.f, 1.f); - vtkm::cont::Box box(minPoint, maxPoint); bool extractInside = false; // Output dataset contains input coordinate system and point data @@ -205,11 +204,12 @@ public: // Output data set with cell set containing extracted points vtkm::worklet::ExtractPoints extractPoints; - OutCellSetType outCellSet = extractPoints.Run(dataset.GetCellSet(0), - dataset.GetCoordinateSystem("coordinates"), - box, - extractInside, - DeviceAdapter()); + OutCellSetType outCellSet = + extractPoints.Run(dataset.GetCellSet(0), + dataset.GetCoordinateSystem("coordinates"), + vtkm::cont::make_ImplicitFunctionHandle(minPoint, maxPoint), + extractInside, + DeviceAdapter()); outDataSet.AddCellSet(outCellSet); VTKM_TEST_ASSERT(test_equal(outCellSet.GetNumberOfCells(), 3),