diff --git a/CMake/testing/VTKmCheckSourceInInstall.cmake b/CMake/testing/VTKmCheckSourceInInstall.cmake index 28b2b0e05..b38b3432a 100644 --- a/CMake/testing/VTKmCheckSourceInInstall.cmake +++ b/CMake/testing/VTKmCheckSourceInInstall.cmake @@ -120,8 +120,10 @@ function(do_verify root_dir prefix) cont/ArrayHandleVirtual.h cont/ArrayHandleVirtual.hxx cont/ArrayHandleVirtualCoordinates.h + cont/CellLocator.h cont/StorageVirtual.h cont/StorageVirtual.hxx + exec/CellLocator.h ) #by default every header in a testing directory doesn't need to be installed diff --git a/examples/lagrangian/CMakeLists.txt b/examples/lagrangian/CMakeLists.txt index c87652ea4..72e3326df 100644 --- a/examples/lagrangian/CMakeLists.txt +++ b/examples/lagrangian/CMakeLists.txt @@ -12,6 +12,19 @@ cmake_minimum_required(VERSION 3.12...3.15 FATAL_ERROR) #Find the VTK-m package find_package(VTKm REQUIRED QUIET) +if ((TARGET vtkm::cuda) OR (TARGET vtkm::kokkos_cuda)) + # CUDA architecture has a limited amount of memory available for constants. The CUDA + # compiler uses this space to hold constants for some optimizations. However, for large + # kernels, the number of constants needed might be larger than the constant space + # available. For these conditions, you have to disable this form of optimization with + # the -Xptxas --disable-optimizer-constants flags. + # TODO: Find a more elegant way to do this. Either figure out a way around this problem + # or add more general flags to vtkm_library/vtkm_unit_tests for sources with "large" kernels. + set_source_files_properties(lagrangian.cxx PROPERTIES + COMPILE_OPTIONS "-Xptxas;--disable-optimizer-constants" + ) +endif() + add_executable(Lagrangian lagrangian.cxx ABCfield.h) target_link_libraries(Lagrangian PRIVATE vtkm_filter) vtkm_add_target_information(Lagrangian diff --git a/vtkm/Deprecated.h b/vtkm/Deprecated.h index 656f71e8a..68baf1b61 100644 --- a/vtkm/Deprecated.h +++ b/vtkm/Deprecated.h @@ -133,7 +133,14 @@ // Only actually use the [[deprecated]] attribute if the compiler supports it AND // we know how to suppress deprecations when necessary. #if defined(VTK_M_DEPRECATED_ATTRIBUTE_SUPPORTED) && defined(VTKM_DEPRECATED_SUPPRESS_SUPPORTED) +#ifdef VTKM_MSVC #define VTKM_DEPRECATED(...) [[deprecated(VTK_M_DEPRECATED_MAKE_MESSAGE(__VA_ARGS__))]] +#else // !MSVC +// GCC and other compilers support the C++14 attribute [[deprecated]], but there appears to be a +// bug (or other undesirable behavior) where if you mix [[deprecated]] with __attribute__(()) you +// get compile errors. To get around this, use __attribute((deprecated)) where supported. +#define VTKM_DEPRECATED(...) __attribute__((deprecated(VTK_M_DEPRECATED_MAKE_MESSAGE(__VA_ARGS__)))) +#endif // !MSVC #else #define VTKM_DEPRECATED(...) #endif diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index f253789ae..eff13bd38 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -54,8 +54,8 @@ set(headers BoundsCompute.h BoundsGlobalCompute.h CastAndCall.h - CellLocator.h CellLocatorBoundingIntervalHierarchy.h + CellLocatorChooser.h CellLocatorGeneral.h CellLocatorRectilinearGrid.h CellLocatorTwoLevel.h @@ -164,7 +164,6 @@ set(device_sources AssignerPartitionedDataSet.cxx BoundsCompute.cxx BoundsGlobalCompute.cxx - CellLocator.cxx CellLocatorBoundingIntervalHierarchy.cxx CellLocatorGeneral.cxx CellLocatorRectilinearGrid.cxx @@ -195,6 +194,7 @@ if (NOT VTKm_NO_DEPRECATED_VIRTUAL) set(headers ${headers} ArrayHandleVirtual.h ArrayHandleVirtualCoordinates.h + CellLocator.h StorageVirtual.h ) @@ -205,6 +205,7 @@ if (NOT VTKm_NO_DEPRECATED_VIRTUAL) set(device_sources ${device_sources} ArrayHandleVirtual.cxx + CellLocator.cxx StorageVirtual.cxx ) endif() diff --git a/vtkm/cont/CellLocator.cxx b/vtkm/cont/CellLocator.cxx index 31c8c79d7..146290e45 100644 --- a/vtkm/cont/CellLocator.cxx +++ b/vtkm/cont/CellLocator.cxx @@ -14,6 +14,8 @@ namespace vtkm { namespace cont { +VTKM_DEPRECATED_SUPPRESS_BEGIN CellLocator::~CellLocator() = default; +VTKM_DEPRECATED_SUPPRESS_END } } diff --git a/vtkm/cont/CellLocator.h b/vtkm/cont/CellLocator.h index 746ca3a82..1a96564a1 100644 --- a/vtkm/cont/CellLocator.h +++ b/vtkm/cont/CellLocator.h @@ -12,6 +12,7 @@ #include +#include #include #include #include @@ -19,13 +20,21 @@ #include +#ifdef VTKM_NO_DEPRECATED_VIRTUAL +#error "CellLocator with virtual methods is removed. Do not include CellLocator.h" +#endif + namespace vtkm { namespace cont { -class VTKM_CONT_EXPORT CellLocator : public vtkm::cont::ExecutionObjectBase +class VTKM_CONT_EXPORT VTKM_DEPRECATED(1.6, + "CellLocator with virtual methods no longer supported. Use " + "CellLocatorGeneral or CellLocatorChooser.") CellLocator + : public vtkm::cont::ExecutionObjectBase { + VTKM_DEPRECATED_SUPPRESS_BEGIN public: virtual ~CellLocator(); @@ -79,6 +88,7 @@ private: vtkm::cont::CoordinateSystem Coords; bool Modified = true; }; +VTKM_DEPRECATED_SUPPRESS_END } // namespace cont } // namespace vtkm diff --git a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx index b19067c80..cb343059b 100644 --- a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx +++ b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx @@ -21,7 +21,7 @@ #include #include #include -#include +#include #include #include @@ -224,8 +224,6 @@ IdArrayHandle CalculateSplitScatterIndices(const IdArrayHandle& cellIds, } // anonymous namespace -CellLocatorBoundingIntervalHierarchy::~CellLocatorBoundingIntervalHierarchy() = default; - void CellLocatorBoundingIntervalHierarchy::Build() { @@ -446,44 +444,32 @@ void CellLocatorBoundingIntervalHierarchy::Build() //std::cout << "Total time: " << totalTimer.GetElapsedTime() << "\n"; } -namespace -{ - -struct BIHCellSetCaster +struct CellLocatorBoundingIntervalHierarchy::MakeExecObject { template - void operator()( - const CellSetType& cellset, - vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token, - vtkm::cont::VirtualObjectHandle& bihExec, - const vtkm::cont::ArrayHandle& nodes, - const vtkm::cont::ArrayHandle& processedCellIds, - const vtkm::cont::CoordinateSystem::MultiplexerArrayType& coords) const + VTKM_CONT void operator()(const CellSetType& cellSet, + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token, + const CellLocatorBoundingIntervalHierarchy& self, + ExecObjType& execObject) const { - using ExecutionType = vtkm::exec::CellLocatorBoundingIntervalHierarchyExec; - ExecutionType* execObject = - new ExecutionType(nodes, processedCellIds, cellset, coords, device, token); - bihExec.Reset(execObject); + execObject = vtkm::exec::CellLocatorBoundingIntervalHierarchy( + self.Nodes, + self.ProcessedCellIds, + cellSet, + self.GetCoordinates().GetDataAsMultiplexer(), + device, + token); } }; -} // anonymous namespace - - -const vtkm::exec::CellLocator* CellLocatorBoundingIntervalHierarchy::PrepareForExecution( - vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const +CellLocatorBoundingIntervalHierarchy::ExecObjType +CellLocatorBoundingIntervalHierarchy::PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { - this->GetCellSet().CastAndCall(BIHCellSetCaster{}, - device, - token, - this->ExecutionObjectHandle, - this->Nodes, - this->ProcessedCellIds, - this->GetCoordinates().GetDataAsMultiplexer()); - return this->ExecutionObjectHandle.PrepareForExecution(device, token); - ; + ExecObjType execObject; + this->GetCellSet().CastAndCall(MakeExecObject{}, device, token, *this, execObject); + return execObject; } } //namespace cont diff --git a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h index 188a35491..d7fccd4ea 100644 --- a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h +++ b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h @@ -16,8 +16,11 @@ #include #include #include -#include -#include + +#include + +#include +#include #include @@ -26,9 +29,19 @@ namespace vtkm namespace cont { -class VTKM_CONT_EXPORT CellLocatorBoundingIntervalHierarchy : public vtkm::cont::CellLocator +class VTKM_CONT_EXPORT CellLocatorBoundingIntervalHierarchy + : public vtkm::cont::internal::CellLocatorBase { + using Superclass = vtkm::cont::internal::CellLocatorBase; + public: + using SupportedCellSets = VTKM_DEFAULT_CELL_SET_LIST; + + using CellLocatorExecList = + vtkm::ListTransform; + + using ExecObjType = vtkm::ListApply; + VTKM_CONT CellLocatorBoundingIntervalHierarchy(vtkm::IdComponent numPlanes = 4, vtkm::IdComponent maxLeafSize = 5) @@ -39,8 +52,6 @@ public: { } - VTKM_CONT ~CellLocatorBoundingIntervalHierarchy() override; - VTKM_CONT void SetNumberOfSplittingPlanes(vtkm::IdComponent numPlanes) { @@ -61,13 +72,8 @@ public: VTKM_CONT vtkm::Id GetMaxLeafSize() { return this->MaxLeafSize; } - VTKM_CONT - const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const override; - -protected: - VTKM_CONT - void Build() override; + VTKM_CONT ExecObjType PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const; private: vtkm::IdComponent NumPlanes; @@ -75,7 +81,10 @@ private: vtkm::cont::ArrayHandle Nodes; vtkm::cont::ArrayHandle ProcessedCellIds; - mutable vtkm::cont::VirtualObjectHandle ExecutionObjectHandle; + friend Superclass; + VTKM_CONT void Build(); + + struct MakeExecObject; }; } // namespace cont diff --git a/vtkm/cont/CellLocatorChooser.h b/vtkm/cont/CellLocatorChooser.h new file mode 100644 index 000000000..91759b800 --- /dev/null +++ b/vtkm/cont/CellLocatorChooser.h @@ -0,0 +1,164 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_CellLocatorChooser_h +#define vtk_m_cont_CellLocatorChooser_h + +#include +#include +#include +#include +#include +#include + +namespace vtkm +{ +namespace cont +{ + +namespace detail +{ + +template +struct CellLocatorChooserImpl +{ + using type = vtkm::cont::CellLocatorTwoLevel; +}; + +using UniformArray = vtkm::cont::ArrayHandleUniformPointCoordinates; + +template <> +struct CellLocatorChooserImpl, UniformArray> +{ + using type = vtkm::cont::CellLocatorUniformGrid; +}; + +using RectilinearArray = + vtkm::cont::ArrayHandleCartesianProduct, + vtkm::cont::ArrayHandle, + vtkm::cont::ArrayHandle>; + +template <> +struct CellLocatorChooserImpl, RectilinearArray> +{ + using type = vtkm::cont::CellLocatorRectilinearGrid; +}; + +} // namespace detail + +/// \brief A template to select an appropriate CellLocator based on CellSet type. +/// +/// Given a concrete type for a `CellSet` subclass and a type of `ArrayHandle` for the +/// coordinate system, `CellLocatorChooser` picks an appropriate `CellLocator` for that +/// type of grid. It is a convenient class to use when you can resolve your templates +/// to discover the type of data set being used for location. +/// +template +using CellLocatorChooser = + typename detail::CellLocatorChooserImpl::type; + +namespace detail +{ + +struct CastAndCallCellLocatorChooserFunctor +{ + template + void CallFunctorWithLocator(const vtkm::cont::DynamicCellSet& cellSet, + const vtkm::cont::CoordinateSystem& coordinateSystem, + Functor&& functor, + Args&&... args) const + { + CellLocatorType locator; + locator.SetCellSet(cellSet); + locator.SetCoordinates(coordinateSystem); + + functor(locator, std::forward(args)...); + } + + template + void operator()(const CellSetType& cellSet, + const vtkm::cont::CoordinateSystem& coordinateSystem, + Functor&& functor, + Args&&... args) const + { + this->CallFunctorWithLocator( + cellSet, coordinateSystem, std::forward(functor), std::forward(args)...); + } + + template + void operator()(const vtkm::cont::CellSetStructured<3>& cellSet, + const vtkm::cont::CoordinateSystem& coordinateSystem, + Functor&& functor, + Args&&... args) const + { + auto coordArray = coordinateSystem.GetData(); + if (coordArray.IsType()) + { + this->CallFunctorWithLocator( + cellSet, coordinateSystem, std::forward(functor), std::forward(args)...); + } + else if (coordArray.IsType()) + { + this->CallFunctorWithLocator( + cellSet, coordinateSystem, std::forward(functor), std::forward(args)...); + } + else + { + this->CallFunctorWithLocator( + cellSet, coordinateSystem, std::forward(functor), std::forward(args)...); + } + } +}; + +} // namespace detail + +/// \brief Calls a functor with the appropriate type of `CellLocator`. +/// +/// Given a cell set and a coordinate system of unknown types, calls a functor with an appropriate +/// CellLocator of the given type. The CellLocator is populated with the provided cell set and +/// coordinate system. +/// +/// Any additional args are passed to the functor. +/// +template +VTKM_CONT void CastAndCallCellLocatorChooser( + const vtkm::cont::DynamicCellSetBase& cellSet, + const vtkm::cont::CoordinateSystem& coordinateSystem, + Functor&& functor, + Args&&... args) +{ + vtkm::cont::CastAndCall(cellSet, + detail::CastAndCallCellLocatorChooserFunctor{}, + coordinateSystem, + std::forward(functor), + std::forward(args)...); +} + +/// \brief Calls a functor with the appropriate type of `CellLocator`. +/// +/// Given a `DataSet`, calls a functor with an appropriate CellLocator of the given type. +/// The CellLocator is populated with the provided cell set and coordinate system. +/// +/// Any additional args are passed to the functor. +/// +template +VTKM_CONT void CastAndCallCellLocatorChooser(const vtkm::cont::DataSet& dataSet, + Functor&& functor, + Args&&... args) +{ + CastAndCallCellLocatorChooser(dataSet.GetCellSet(), + dataSet.GetCoordinateSystem(), + std::forward(functor), + std::forward(args)...); +} + +} +} // namespace vtkm::cont + +#endif //vtk_m_cont_CellLocatorChooser_h diff --git a/vtkm/cont/CellLocatorGeneral.cxx b/vtkm/cont/CellLocatorGeneral.cxx index 475aeda9b..272094f1c 100644 --- a/vtkm/cont/CellLocatorGeneral.cxx +++ b/vtkm/cont/CellLocatorGeneral.cxx @@ -20,39 +20,19 @@ namespace { -VTKM_CONT -void DefaultConfigurator(std::unique_ptr& locator, - const vtkm::cont::DynamicCellSet& cellSet, - const vtkm::cont::CoordinateSystem& coords) +template +void BuildForType(vtkm::cont::CellLocatorGeneral& locator, LocatorVariantType& locatorVariant) { - using StructuredCellSet = vtkm::cont::CellSetStructured<3>; - using UniformCoordinates = vtkm::cont::ArrayHandleUniformPointCoordinates; - using RectilinearCoordinates = - vtkm::cont::ArrayHandleCartesianProduct, - vtkm::cont::ArrayHandle, - vtkm::cont::ArrayHandle>; - - if (cellSet.IsType() && coords.GetData().IsType()) + constexpr vtkm::IdComponent LOCATOR_INDEX = + LocatorVariantType::template GetIndexOf(); + if (locatorVariant.GetIndex() != LOCATOR_INDEX) { - if (!dynamic_cast(locator.get())) - { - locator.reset(new vtkm::cont::CellLocatorUniformGrid); - } + locatorVariant = LocatorImplType{}; } - else if (cellSet.IsType() && coords.GetData().IsType()) - { - if (!dynamic_cast(locator.get())) - { - locator.reset(new vtkm::cont::CellLocatorRectilinearGrid); - } - } - else if (!dynamic_cast(locator.get())) - { - locator.reset(new vtkm::cont::CellLocatorTwoLevel); - } - - locator->SetCellSet(cellSet); - locator->SetCoordinates(coords); + LocatorImplType& locatorImpl = locatorVariant.template Get(); + locatorImpl.SetCellSet(locator.GetCellSet()); + locatorImpl.SetCoordinates(locator.GetCoordinates()); + locatorImpl.Update(); } } // anonymous namespace @@ -62,33 +42,50 @@ namespace vtkm namespace cont { -VTKM_CONT CellLocatorGeneral::CellLocatorGeneral() - : Configurator(DefaultConfigurator) +VTKM_CONT void CellLocatorGeneral::Build() { + using StructuredCellSet = vtkm::cont::CellSetStructured<3>; + using UniformCoordinates = vtkm::cont::ArrayHandleUniformPointCoordinates; + using RectilinearCoordinates = + vtkm::cont::ArrayHandleCartesianProduct, + vtkm::cont::ArrayHandle, + vtkm::cont::ArrayHandle>; + + vtkm::cont::DynamicCellSet cellSet = this->GetCellSet(); + vtkm::cont::CoordinateSystem coords = this->GetCoordinates(); + + if (cellSet.IsType() && coords.GetData().IsType()) + { + BuildForType(*this, this->LocatorImpl); + } + else if (cellSet.IsType() && coords.GetData().IsType()) + { + BuildForType(*this, this->LocatorImpl); + } + else + { + BuildForType(*this, this->LocatorImpl); + } } -VTKM_CONT CellLocatorGeneral::~CellLocatorGeneral() = default; +struct CellLocatorGeneral::PrepareFunctor +{ + template + ExecObjType operator()(LocatorType&& locator, + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const + { + return locator.PrepareForExecution(device, token); + } +}; -VTKM_CONT const vtkm::exec::CellLocator* CellLocatorGeneral::PrepareForExecution( +CellLocatorGeneral::ExecObjType CellLocatorGeneral::PrepareForExecution( vtkm::cont::DeviceAdapterId device, vtkm::cont::Token& token) const { - if (this->Locator) - { - return this->Locator->PrepareForExecution(device, token); - } - return nullptr; + this->Update(); + return this->LocatorImpl.CastAndCall(PrepareFunctor{}, device, token); } -VTKM_CONT void CellLocatorGeneral::Build() -{ - this->Configurator(this->Locator, this->GetCellSet(), this->GetCoordinates()); - this->Locator->Update(); -} - -VTKM_CONT void CellLocatorGeneral::ResetToDefaultConfigurator() -{ - this->SetConfigurator(DefaultConfigurator); -} } } // vtkm::cont diff --git a/vtkm/cont/CellLocatorGeneral.h b/vtkm/cont/CellLocatorGeneral.h index ce15763d9..a5d34af36 100644 --- a/vtkm/cont/CellLocatorGeneral.h +++ b/vtkm/cont/CellLocatorGeneral.h @@ -10,7 +10,13 @@ #ifndef vtk_m_cont_CellLocatorGeneral_h #define vtk_m_cont_CellLocatorGeneral_h -#include +#include +#include +#include + +#include + +#include #include #include @@ -20,48 +26,45 @@ namespace vtkm namespace cont { -class VTKM_CONT_EXPORT CellLocatorGeneral : public vtkm::cont::CellLocator +/// \brief A CellLocator that works generally well for any supported cell set. +/// +/// `CellLocatorGeneral` creates a `CellLocator` that acts like a multiplexer to +/// switch at runtime to any supported cell set. It is a convenient class to use +/// when the type of `CellSet` cannot be determined at runtime. +/// +/// Note that `CellLocatorGeneral` only supports a finite amount of `CellSet` types. +/// Thus, it is possible to give it a cell set type that is not supported. +/// +/// Also note that `CellLocatorGeneral` can add a significant amount of code inside +/// of worklet that uses it, and this might cause some issues with some compilers. +/// +class VTKM_CONT_EXPORT CellLocatorGeneral + : public vtkm::cont::internal::CellLocatorBase { + using Superclass = vtkm::cont::internal::CellLocatorBase; + public: - VTKM_CONT CellLocatorGeneral(); + using ContLocatorList = vtkm::List; - VTKM_CONT ~CellLocatorGeneral() override; + using ExecLocatorList = + vtkm::List, + vtkm::cont::internal::ExecutionObjectType, + vtkm::cont::internal::ExecutionObjectType>; - /// Get the current underlying cell locator - /// - VTKM_CONT const vtkm::cont::CellLocator* GetCurrentLocator() const { return this->Locator.get(); } + using ExecObjType = vtkm::ListApply; - /// A configurator can be provided to select an appropriate - /// cell locator implementation and configure its parameters, based on the - /// input cell set and cooridinates. - /// If unset, a resonable default is used. - /// - using ConfiguratorSignature = void(std::unique_ptr&, - const vtkm::cont::DynamicCellSet&, - const vtkm::cont::CoordinateSystem&); - - VTKM_CONT void SetConfigurator(const std::function& configurator) - { - this->Configurator = configurator; - } - - VTKM_CONT const std::function& GetConfigurator() const - { - return this->Configurator; - } - - VTKM_CONT void ResetToDefaultConfigurator(); - - VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const override; - -protected: - VTKM_CONT void Build() override; + VTKM_CONT ExecObjType PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const; private: - std::unique_ptr Locator; - std::function Configurator; + vtkm::cont::internal::ListAsVariant LocatorImpl; + + friend Superclass; + VTKM_CONT void Build(); + + struct PrepareFunctor; }; } } // vtkm::cont diff --git a/vtkm/cont/CellLocatorRectilinearGrid.cxx b/vtkm/cont/CellLocatorRectilinearGrid.cxx index f41f9a935..d86964cdf 100644 --- a/vtkm/cont/CellLocatorRectilinearGrid.cxx +++ b/vtkm/cont/CellLocatorRectilinearGrid.cxx @@ -20,15 +20,6 @@ namespace vtkm namespace cont { -CellLocatorRectilinearGrid::CellLocatorRectilinearGrid() = default; - -CellLocatorRectilinearGrid::~CellLocatorRectilinearGrid() = default; - -using Structured2DType = vtkm::cont::CellSetStructured<2>; -using Structured3DType = vtkm::cont::CellSetStructured<3>; -using AxisHandle = vtkm::cont::ArrayHandle; -using RectilinearType = vtkm::cont::ArrayHandleCartesianProduct; - void CellLocatorRectilinearGrid::Build() { vtkm::cont::CoordinateSystem coords = this->GetCoordinates(); @@ -59,35 +50,33 @@ void CellLocatorRectilinearGrid::Build() } } -const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution( +vtkm::exec::CellLocatorRectilinearGrid CellLocatorRectilinearGrid::PrepareForExecution( vtkm::cont::DeviceAdapterId device, vtkm::cont::Token& token) const { + this->Update(); + + using ExecObjType = vtkm::exec::CellLocatorRectilinearGrid; + if (this->Is3D) { - using ExecutionType = vtkm::exec::CellLocatorRectilinearGrid<3>; - ExecutionType* execObject = - new ExecutionType(this->PlaneSize, - this->RowSize, - this->GetCellSet().Cast(), - this->GetCoordinates().GetData().AsArrayHandle(), - device, - token); - this->ExecutionObjectHandle.Reset(execObject); + return ExecObjType(this->PlaneSize, + this->RowSize, + this->GetCellSet().template Cast(), + this->GetCoordinates().GetData().template AsArrayHandle(), + device, + token); } else { - using ExecutionType = vtkm::exec::CellLocatorRectilinearGrid<2>; - ExecutionType* execObject = - new ExecutionType(this->PlaneSize, - this->RowSize, - this->GetCellSet().Cast(), - this->GetCoordinates().GetData().AsArrayHandle(), - device, - token); - this->ExecutionObjectHandle.Reset(execObject); + return ExecObjType(this->PlaneSize, + this->RowSize, + this->GetCellSet().template Cast(), + this->GetCoordinates().GetData().template AsArrayHandle(), + device, + token); } - return this->ExecutionObjectHandle.PrepareForExecution(device, token); } + } //namespace cont } //namespace vtkm diff --git a/vtkm/cont/CellLocatorRectilinearGrid.h b/vtkm/cont/CellLocatorRectilinearGrid.h index 6e7ed95b9..63dc85a92 100644 --- a/vtkm/cont/CellLocatorRectilinearGrid.h +++ b/vtkm/cont/CellLocatorRectilinearGrid.h @@ -10,27 +10,35 @@ #ifndef vtkm_cont_CellLocatorRectilinearGrid_h #define vtkm_cont_CellLocatorRectilinearGrid_h -#include -#include +#include + +#include namespace vtkm { namespace cont { -class VTKM_CONT_EXPORT CellLocatorRectilinearGrid : public vtkm::cont::CellLocator +class VTKM_CONT_EXPORT CellLocatorRectilinearGrid + : public vtkm::cont::internal::CellLocatorBase { + using Superclass = vtkm::cont::internal::CellLocatorBase; + + using Structured2DType = vtkm::cont::CellSetStructured<2>; + using Structured3DType = vtkm::cont::CellSetStructured<3>; + // Might want to handle cartesian product of both Float32 and Float64. + using AxisHandle = vtkm::cont::ArrayHandle; + using RectilinearType = + vtkm::cont::ArrayHandleCartesianProduct; + public: - VTKM_CONT CellLocatorRectilinearGrid(); + CellLocatorRectilinearGrid() = default; - VTKM_CONT ~CellLocatorRectilinearGrid() override; + ~CellLocatorRectilinearGrid() = default; - VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( + VTKM_CONT vtkm::exec::CellLocatorRectilinearGrid PrepareForExecution( vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const override; - -protected: - VTKM_CONT void Build() override; + vtkm::cont::Token& token) const; private: vtkm::Bounds Bounds; @@ -38,7 +46,8 @@ private: vtkm::Id RowSize; bool Is3D = true; - mutable vtkm::cont::VirtualObjectHandle ExecutionObjectHandle; + friend Superclass; + VTKM_CONT void Build(); }; } //namespace cont diff --git a/vtkm/cont/CellLocatorTwoLevel.cxx b/vtkm/cont/CellLocatorTwoLevel.cxx index d17697bdc..e1294628f 100644 --- a/vtkm/cont/CellLocatorTwoLevel.cxx +++ b/vtkm/cont/CellLocatorTwoLevel.cxx @@ -461,28 +461,31 @@ struct CellLocatorTwoLevel::MakeExecObject VTKM_CONT void operator()(const CellSetType& cellSet, vtkm::cont::DeviceAdapterId device, vtkm::cont::Token& token, - const CellLocatorTwoLevel& self) const + const CellLocatorTwoLevel& self, + ExecObjType& execObject) const { - auto execObject = new vtkm::exec::CellLocatorTwoLevel(self.TopLevel, - self.LeafDimensions, - self.LeafStartIndex, - self.CellStartIndex, - self.CellCount, - self.CellIds, - cellSet, - self.GetCoordinates(), - device, - token); - self.ExecutionObjectHandle.Reset(execObject); + using CellStructuredType = CellSetContToExec; + execObject = vtkm::exec::CellLocatorTwoLevel(self.TopLevel, + self.LeafDimensions, + self.LeafStartIndex, + self.CellStartIndex, + self.CellCount, + self.CellIds, + cellSet, + self.GetCoordinates(), + device, + token); } }; -VTKM_CONT const vtkm::exec::CellLocator* CellLocatorTwoLevel::PrepareForExecution( +CellLocatorTwoLevel::ExecObjType CellLocatorTwoLevel::PrepareForExecution( vtkm::cont::DeviceAdapterId device, vtkm::cont::Token& token) const { - this->GetCellSet().CastAndCall(MakeExecObject{}, device, token, *this); - return this->ExecutionObjectHandle.PrepareForExecution(device, token); + this->Update(); + ExecObjType execObject; + this->GetCellSet().CastAndCall(MakeExecObject{}, device, token, *this, execObject); + return execObject; } //---------------------------------------------------------------------------- diff --git a/vtkm/cont/CellLocatorTwoLevel.h b/vtkm/cont/CellLocatorTwoLevel.h index c79162152..5f9ff0ea3 100644 --- a/vtkm/cont/CellLocatorTwoLevel.h +++ b/vtkm/cont/CellLocatorTwoLevel.h @@ -11,9 +11,11 @@ #define vtk_m_cont_CellLocatorTwoLevel_h #include -#include -#include +#include +#include + +#include #include @@ -38,9 +40,25 @@ namespace cont /// Javor Kalojanov, Markus Billeter, and Philipp Slusallek. "Two-Level Grids for Ray Tracing /// on GPUs." _Computer Graphics Forum_, 2011, pages 307-314. DOI 10.1111/j.1467-8659.2011.01862.x /// -class VTKM_CONT_EXPORT CellLocatorTwoLevel : public vtkm::cont::CellLocator +class VTKM_CONT_EXPORT CellLocatorTwoLevel + : public vtkm::cont::internal::CellLocatorBase { + using Superclass = vtkm::cont::internal::CellLocatorBase; + + template + using CellSetContToExec = + typename CellSetCont::template ExecConnectivityType; + public: + using SupportedCellSets = VTKM_DEFAULT_CELL_SET_LIST; + + using CellExecObjectList = vtkm::ListTransform; + using CellLocatorExecList = + vtkm::ListTransform; + + using ExecObjType = vtkm::ListApply; + CellLocatorTwoLevel() : DensityL1(32.0f) , DensityL2(2.0f) @@ -67,11 +85,13 @@ public: void PrintSummary(std::ostream& out) const; - const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const override; +public: + ExecObjType PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const; private: - VTKM_CONT void Build() override; + friend Superclass; + VTKM_CONT void Build(); vtkm::FloatDefault DensityL1, DensityL2; @@ -82,10 +102,9 @@ private: vtkm::cont::ArrayHandle CellCount; vtkm::cont::ArrayHandle CellIds; - mutable vtkm::cont::VirtualObjectHandle ExecutionObjectHandle; - struct MakeExecObject; }; + } } // vtkm::cont diff --git a/vtkm/cont/CellLocatorUniformGrid.cxx b/vtkm/cont/CellLocatorUniformGrid.cxx index 612d77ca4..c2c11bce1 100644 --- a/vtkm/cont/CellLocatorUniformGrid.cxx +++ b/vtkm/cont/CellLocatorUniformGrid.cxx @@ -13,15 +13,10 @@ #include #include -#include - namespace vtkm { namespace cont { -CellLocatorUniformGrid::CellLocatorUniformGrid() = default; - -CellLocatorUniformGrid::~CellLocatorUniformGrid() = default; using UniformType = vtkm::cont::ArrayHandleUniformPointCoordinates; using Structured2DType = vtkm::cont::CellSetStructured<2>; @@ -73,25 +68,13 @@ void CellLocatorUniformGrid::Build() this->CellDims[2] = this->PointDims[2] - 1; } -const vtkm::exec::CellLocator* CellLocatorUniformGrid::PrepareForExecution( - vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const +vtkm::exec::CellLocatorUniformGrid CellLocatorUniformGrid::PrepareForExecution( + vtkm::cont::DeviceAdapterId vtkmNotUsed(device), + vtkm::cont::Token& vtkmNotUsed(token)) const { - if (this->Is3D) - { - using ExecutionType = vtkm::exec::CellLocatorUniformGrid<3>; - ExecutionType* execObject = new ExecutionType( - this->CellDims, this->PointDims, this->Origin, this->InvSpacing, this->MaxPoint); - this->ExecutionObjectHandle.Reset(execObject); - } - else - { - using ExecutionType = vtkm::exec::CellLocatorUniformGrid<2>; - ExecutionType* execObject = new ExecutionType( - this->CellDims, this->PointDims, this->Origin, this->InvSpacing, this->MaxPoint); - this->ExecutionObjectHandle.Reset(execObject); - } - return this->ExecutionObjectHandle.PrepareForExecution(device, token); + this->Update(); + return vtkm::exec::CellLocatorUniformGrid( + this->CellDims, this->Origin, this->InvSpacing, this->MaxPoint); } } //namespace cont diff --git a/vtkm/cont/CellLocatorUniformGrid.h b/vtkm/cont/CellLocatorUniformGrid.h index ae81c43d2..0de6f9927 100644 --- a/vtkm/cont/CellLocatorUniformGrid.h +++ b/vtkm/cont/CellLocatorUniformGrid.h @@ -10,27 +10,24 @@ #ifndef vtkm_cont_CellLocatorUniformGrid_h #define vtkm_cont_CellLocatorUniformGrid_h -#include -#include +#include + +#include namespace vtkm { namespace cont { -class VTKM_CONT_EXPORT CellLocatorUniformGrid : public vtkm::cont::CellLocator +class VTKM_CONT_EXPORT CellLocatorUniformGrid + : public vtkm::cont::internal::CellLocatorBase { + using Superclass = vtkm::cont::internal::CellLocatorBase; + public: - VTKM_CONT CellLocatorUniformGrid(); - - VTKM_CONT ~CellLocatorUniformGrid() override; - - VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( + VTKM_CONT vtkm::exec::CellLocatorUniformGrid PrepareForExecution( vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) const override; - -protected: - VTKM_CONT void Build() override; + vtkm::cont::Token& token) const; private: vtkm::Id3 CellDims; @@ -40,7 +37,8 @@ private: vtkm::Vec3f MaxPoint; bool Is3D = true; - mutable vtkm::cont::VirtualObjectHandle ExecutionObjectHandle; + friend Superclass; + VTKM_CONT void Build(); }; } } // vtkm::cont diff --git a/vtkm/cont/ExecutionObjectBase.h b/vtkm/cont/ExecutionObjectBase.h index f3aa29d4b..52d248bc4 100644 --- a/vtkm/cont/ExecutionObjectBase.h +++ b/vtkm/cont/ExecutionObjectBase.h @@ -151,9 +151,9 @@ VTKM_CONT VTKM_DEPRECATED( /// environment for a particular device. This templated type gives the type for the class used /// in the execution environment for a given ExecutionObject and device. /// -template +template using ExecutionObjectType = decltype(CallPrepareForExecution(std::declval(), - Device{}, + std::declval(), std::declval())); } // namespace internal diff --git a/vtkm/cont/VirtualObjectHandle.h b/vtkm/cont/VirtualObjectHandle.h index 7255c9e7a..f9904e7b2 100644 --- a/vtkm/cont/VirtualObjectHandle.h +++ b/vtkm/cont/VirtualObjectHandle.h @@ -100,6 +100,7 @@ public: bool acquireOwnership = true, DeviceAdapterList devices = DeviceAdapterList()) { + VTKM_DEPRECATED_SUPPRESS_BEGIN VTKM_STATIC_ASSERT_MSG((std::is_base_of::value), "Tried to bind a type that is not a subclass of the base class."); @@ -118,6 +119,7 @@ public: vtkm::cont::internal::ForEachValidDevice( devices, internal::CreateTransferInterface(), this->Internals.get(), derived); } + VTKM_DEPRECATED_SUPPRESS_END } /// Release all host and execution side resources diff --git a/vtkm/cont/internal/CMakeLists.txt b/vtkm/cont/internal/CMakeLists.txt index 006e45f94..7c0c4d166 100644 --- a/vtkm/cont/internal/CMakeLists.txt +++ b/vtkm/cont/internal/CMakeLists.txt @@ -17,6 +17,7 @@ set(headers AtomicInterfaceExecution.h Buffer.h CastInvalidValue.h + CellLocatorBase.h ConnectivityExplicitInternals.h DeviceAdapterAlgorithmGeneral.h DeviceAdapterMemoryManager.h diff --git a/vtkm/cont/internal/CellLocatorBase.h b/vtkm/cont/internal/CellLocatorBase.h new file mode 100644 index 000000000..c2a0386ac --- /dev/null +++ b/vtkm/cont/internal/CellLocatorBase.h @@ -0,0 +1,197 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_internal_CellLocatorBase_h +#define vtk_m_cont_internal_CellLocatorBase_h + +#include + +#include +#include +#include +#include + +#ifndef VTKM_NO_DEPRECATED_VIRTUAL +// To support deprecated implementation +#include +#include +#include +#endif //!VTKM_NO_DEPRECATED_VIRTUAL + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +namespace detail +{ + +#ifndef VTKM_NO_DEPRECATED_VIRTUAL +VTKM_DEPRECATED_SUPPRESS_BEGIN + +// Wrong namespace, but it's only for deprecated code. +template +class VTKM_ALWAYS_EXPORT CellLocatorBaseExecWrapper : public vtkm::exec::CellLocator +{ + LocatorType Locator; + +public: + VTKM_CONT CellLocatorBaseExecWrapper(const LocatorType& locator) + : Locator(locator) + { + } + + VTKM_EXEC_CONT virtual ~CellLocatorBaseExecWrapper() noexcept override + { + // This must not be defaulted, since defaulted virtual destructors are + // troublesome with CUDA __host__ __device__ markup. + } + + VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point, + vtkm::Id& cellId, + vtkm::Vec3f& parametric) const override + { + return this->Locator.FindCell(point, cellId, parametric); + } +}; + +template +struct CellLocatorBaseWrapperPrepareForExecutionFunctor +{ + template + VTKM_CONT bool operator()(Device device, + vtkm::cont::VirtualObjectHandle& execHandle, + const LocatorType& locator, + vtkm::cont::Token& token) + { + auto execObject = locator.PrepareForExecution(device, token); + using WrapType = CellLocatorBaseExecWrapper; + execHandle.Reset(new WrapType(execObject)); + return true; + } +}; + +template +class VTKM_ALWAYS_EXPORT CellLocatorBaseWrapper : public vtkm::cont::CellLocator +{ + Derived Locator; + mutable vtkm::cont::VirtualObjectHandle ExecutionObjectHandle; + +public: + CellLocatorBaseWrapper() = default; + + CellLocatorBaseWrapper(const Derived& locator) + : Locator(locator) + { + this->SetCellSet(locator.GetCellSet()); + this->SetCoordinates(locator.GetCoordinates()); + } + + VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override + { + bool success = + vtkm::cont::TryExecuteOnDevice(device, + CellLocatorBaseWrapperPrepareForExecutionFunctor{}, + this->ExecutionObjectHandle, + this->Locator, + token); + if (!success) + { + throwFailedRuntimeDeviceTransfer("CellLocatorUniformGrid", device); + } + return this->ExecutionObjectHandle.PrepareForExecution(device, token); + } + +private: + void Build() override + { + this->Locator.SetCellSet(this->GetCellSet()); + this->Locator.SetCoordinates(this->GetCoordinates()); + this->Locator.Update(); + } +}; + +VTKM_DEPRECATED_SUPPRESS_END +#endif //!VTKM_NO_DEPRECATED_VIRTUAL + + +} // namespace detail + +/// \brief Base class for all `CellLocator` classes. +/// +/// `CellLocatorBase` uses the curiously recurring template pattern (CRTP). Subclasses +/// must provide their own type for the template parameter. Subclasses must implement +/// `Update` and `PrepareForExecution` methods. +/// +template +class VTKM_ALWAYS_EXPORT CellLocatorBase : public vtkm::cont::ExecutionObjectBase +{ + vtkm::cont::DynamicCellSet CellSet; + vtkm::cont::CoordinateSystem Coords; + mutable bool Modified = true; + +public: +#ifndef VTKM_NO_DEPRECATED_VIRTUAL + VTKM_DEPRECATED_SUPPRESS_BEGIN + // Support deprecated classes + operator detail::CellLocatorBaseWrapper() const + { + return detail::CellLocatorBaseWrapper(reinterpret_cast(*this)); + } + VTKM_DEPRECATED_SUPPRESS_END +#endif //!VTKM_NO_DEPRECATED_VIRTUAL + + const vtkm::cont::DynamicCellSet& GetCellSet() const { return this->CellSet; } + + void SetCellSet(const vtkm::cont::DynamicCellSet& cellSet) + { + this->CellSet = cellSet; + this->SetModified(); + } + + const vtkm::cont::CoordinateSystem& GetCoordinates() const { return this->Coords; } + + void SetCoordinates(const vtkm::cont::CoordinateSystem& coords) + { + this->Coords = coords; + this->SetModified(); + } + + void Update() const + { + if (this->Modified) + { + static_cast(const_cast(this))->Build(); + this->Modified = false; + } + } + + template + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object.") + vtkm::cont::internal::ExecutionObjectType PrepareForExecution( + Device device) const + { + vtkm::cont::Token token; + return this->PrepareForExecution(device, token); + } + +protected: + void SetModified() { this->Modified = true; } + bool GetModified() const { return this->Modified; } +}; + +} +} +} // vtkm::cont::internal + +#endif //vtk_m_cont_internal_CellLocatorBase_h diff --git a/vtkm/cont/testing/CMakeLists.txt b/vtkm/cont/testing/CMakeLists.txt index eb0bb6532..411e25441 100644 --- a/vtkm/cont/testing/CMakeLists.txt +++ b/vtkm/cont/testing/CMakeLists.txt @@ -55,6 +55,7 @@ set(unit_tests UnitTestArrayHandleVirtual.cxx UnitTestArrayHandleXGCCoordinates.cxx UnitTestArrayPortalToIterators.cxx + UnitTestCellLocatorChooser.cxx UnitTestCellLocatorGeneral.cxx UnitTestCellSet.cxx UnitTestCellSetExplicit.cxx diff --git a/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h b/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h index 8925992b5..b3da8c9cf 100644 --- a/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h +++ b/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h @@ -103,7 +103,7 @@ public: bool& match) const { vtkm::Id calculated = CalculateCellId(pointIn); - vtkm::ErrorCode status = locator->FindCell(pointIn, cellId, parametric); + vtkm::ErrorCode status = locator.FindCell(pointIn, cellId, parametric); if (status != vtkm::ErrorCode::Success) { this->RaiseError(vtkm::ErrorString(status)); diff --git a/vtkm/cont/testing/TestingCellLocatorTwoLevel.h b/vtkm/cont/testing/TestingCellLocatorTwoLevel.h index 8aa9bdc1a..3e7648cf8 100644 --- a/vtkm/cont/testing/TestingCellLocatorTwoLevel.h +++ b/vtkm/cont/testing/TestingCellLocatorTwoLevel.h @@ -173,7 +173,7 @@ public: vtkm::Id& cellId, vtkm::Vec3f& pcoords) const { - vtkm::ErrorCode status = locator->FindCell(point, cellId, pcoords); + vtkm::ErrorCode status = locator.FindCell(point, cellId, pcoords); if (status != vtkm::ErrorCode::Success) { this->RaiseError(vtkm::ErrorString(status)); diff --git a/vtkm/cont/testing/TestingCellLocatorUniformGrid.h b/vtkm/cont/testing/TestingCellLocatorUniformGrid.h index 185499b7b..1ba2ae699 100644 --- a/vtkm/cont/testing/TestingCellLocatorUniformGrid.h +++ b/vtkm/cont/testing/TestingCellLocatorUniformGrid.h @@ -69,7 +69,7 @@ public: bool& match) const { vtkm::Id calculated = CalculateCellId(pointIn); - vtkm::ErrorCode status = locator->FindCell(pointIn, cellId, parametric); + vtkm::ErrorCode status = locator.FindCell(pointIn, cellId, parametric); if ((status != vtkm::ErrorCode::Success) && (status != vtkm::ErrorCode::CellNotFound)) { this->RaiseError(vtkm::ErrorString(status)); diff --git a/vtkm/cont/testing/UnitTestCellLocatorChooser.cxx b/vtkm/cont/testing/UnitTestCellLocatorChooser.cxx new file mode 100644 index 000000000..869ca0166 --- /dev/null +++ b/vtkm/cont/testing/UnitTestCellLocatorChooser.cxx @@ -0,0 +1,232 @@ +//============================================================================ +// 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. +//============================================================================ +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace +{ + +std::default_random_engine RandomGenerator; + +using PointType = vtkm::Vec3f; + +//----------------------------------------------------------------------------- +vtkm::cont::DataSet MakeTestDataSetUniform() +{ + return vtkm::cont::DataSetBuilderUniform::Create( + vtkm::Id3{ 64 }, PointType{ -32.0f }, PointType{ 1.0f / 64.0f }); +} + +vtkm::cont::DataSet MakeTestDataSetRectilinear() +{ + std::uniform_real_distribution coordGen(1.0f / 128.0f, 1.0f / 32.0f); + + vtkm::cont::ArrayHandle coords[3]; + for (int i = 0; i < 3; ++i) + { + coords[i].Allocate(64); + auto portal = coords[i].WritePortal(); + + vtkm::FloatDefault cur = 0.0f; + for (vtkm::Id j = 0; j < portal.GetNumberOfValues(); ++j) + { + cur += coordGen(RandomGenerator); + portal.Set(j, cur); + } + } + + return vtkm::cont::DataSetBuilderRectilinear::Create(coords[0], coords[1], coords[2]); +} + +vtkm::cont::DataSet MakeTestDataSetCurvilinear() +{ + auto recti = MakeTestDataSetRectilinear(); + auto coords = recti.GetCoordinateSystem().GetDataAsMultiplexer(); + + vtkm::cont::ArrayHandle sheared; + sheared.Allocate(coords.GetNumberOfValues()); + + auto inPortal = coords.ReadPortal(); + auto outPortal = sheared.WritePortal(); + for (vtkm::Id i = 0; i < inPortal.GetNumberOfValues(); ++i) + { + auto val = inPortal.Get(i); + outPortal.Set(i, val + vtkm::make_Vec(val[1], val[2], val[0])); + } + + vtkm::cont::DataSet curvi; + curvi.SetCellSet(recti.GetCellSet()); + curvi.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coords", sheared)); + + return curvi; +} + +//----------------------------------------------------------------------------- +class ParametricToWorldCoordinates : public vtkm::worklet::WorkletVisitCellsWithPoints +{ +public: + using ControlSignature = void(CellSetIn cellset, + FieldInPoint coords, + FieldInOutCell pcs, + FieldOutCell wcs); + using ExecutionSignature = void(CellShape, _2, _3, _4); + + using ScatterType = vtkm::worklet::ScatterPermutation<>; + + VTKM_CONT + static ScatterType MakeScatter(const vtkm::cont::ArrayHandle& cellIds) + { + return ScatterType(cellIds); + } + + template + VTKM_EXEC void operator()(CellShapeTagType cellShape, + PointsVecType points, + const PointType& pc, + PointType& wc) const + { + auto status = vtkm::exec::CellInterpolate(points, pc, cellShape, wc); + if (status != vtkm::ErrorCode::Success) + { + this->RaiseError(vtkm::ErrorString(status)); + } + } +}; + +void GenerateRandomInput(const vtkm::cont::DataSet& ds, + vtkm::Id count, + vtkm::cont::ArrayHandle& cellIds, + vtkm::cont::ArrayHandle& pcoords, + vtkm::cont::ArrayHandle& wcoords) +{ + vtkm::Id numberOfCells = ds.GetNumberOfCells(); + + std::uniform_int_distribution cellIdGen(0, numberOfCells - 1); + std::uniform_real_distribution pcoordGen(0.0f, 1.0f); + + cellIds.Allocate(count); + pcoords.Allocate(count); + wcoords.Allocate(count); + + for (vtkm::Id i = 0; i < count; ++i) + { + cellIds.WritePortal().Set(i, cellIdGen(RandomGenerator)); + + PointType pc{ pcoordGen(RandomGenerator), + pcoordGen(RandomGenerator), + pcoordGen(RandomGenerator) }; + pcoords.WritePortal().Set(i, pc); + } + + vtkm::worklet::DispatcherMapTopology dispatcher( + ParametricToWorldCoordinates::MakeScatter(cellIds)); + dispatcher.Invoke( + ds.GetCellSet(), ds.GetCoordinateSystem().GetDataAsMultiplexer(), pcoords, wcoords); +} + +//----------------------------------------------------------------------------- +class FindCellWorklet : public vtkm::worklet::WorkletMapField +{ +public: + using ControlSignature = void(FieldIn points, + ExecObject locator, + FieldOut cellIds, + FieldOut pcoords); + using ExecutionSignature = void(_1, _2, _3, _4); + + template + VTKM_EXEC void operator()(const vtkm::Vec3f& point, + const LocatorType& locator, + vtkm::Id& cellId, + vtkm::Vec3f& pcoords) const + { + vtkm::ErrorCode status = locator.FindCell(point, cellId, pcoords); + if (status != vtkm::ErrorCode::Success) + { + this->RaiseError(vtkm::ErrorString(status)); + } + } +}; + +template +void TestWithDataSet(const vtkm::cont::DataSet& dataset) +{ + VTKM_TEST_ASSERT(dataset.GetCellSet().IsType()); + VTKM_TEST_ASSERT(dataset.GetCoordinateSystem().GetData().IsType()); + + vtkm::cont::CellLocatorChooser locator; + locator.SetCellSet(dataset.GetCellSet()); + locator.SetCoordinates(dataset.GetCoordinateSystem()); + locator.Update(); + + vtkm::cont::ArrayHandle expCellIds; + vtkm::cont::ArrayHandle expPCoords; + vtkm::cont::ArrayHandle points; + GenerateRandomInput(dataset, 128, expCellIds, expPCoords, points); + + vtkm::cont::ArrayHandle cellIds; + vtkm::cont::ArrayHandle pcoords; + + vtkm::worklet::DispatcherMapField dispatcher; + dispatcher.Invoke(points, locator, cellIds, pcoords); + + auto cellIdPortal = cellIds.ReadPortal(); + auto expCellIdsPortal = expCellIds.ReadPortal(); + auto pcoordsPortal = pcoords.ReadPortal(); + auto expPCoordsPortal = expPCoords.ReadPortal(); + for (vtkm::Id i = 0; i < 128; ++i) + { + VTKM_TEST_ASSERT(cellIdPortal.Get(i) == expCellIdsPortal.Get(i), "Incorrect cell ids"); + VTKM_TEST_ASSERT(test_equal(pcoordsPortal.Get(i), expPCoordsPortal.Get(i), 1e-3), + "Incorrect parameteric coordinates"); + } + + std::cout << "Passed.\n"; +} + +void TestCellLocatorChooser() +{ + std::cout << "Test UniformGrid dataset\n"; + TestWithDataSet, vtkm::cont::ArrayHandleUniformPointCoordinates>( + MakeTestDataSetUniform()); + + std::cout << "Test Rectilinear dataset\n"; + TestWithDataSet< + vtkm::cont::CellSetStructured<3>, + vtkm::cont::ArrayHandleCartesianProduct, + vtkm::cont::ArrayHandle, + vtkm::cont::ArrayHandle>>( + MakeTestDataSetRectilinear()); + + std::cout << "Test Curvilinear dataset\n"; + TestWithDataSet, vtkm::cont::ArrayHandle>( + MakeTestDataSetCurvilinear()); +} + +} // anonymous namespace + +int UnitTestCellLocatorChooser(int argc, char* argv[]) +{ + vtkm::cont::GetRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial{}); + return vtkm::cont::testing::Testing::Run(TestCellLocatorChooser, argc, argv); +} diff --git a/vtkm/cont/testing/UnitTestCellLocatorGeneral.cxx b/vtkm/cont/testing/UnitTestCellLocatorGeneral.cxx index 1f283c933..fc557a824 100644 --- a/vtkm/cont/testing/UnitTestCellLocatorGeneral.cxx +++ b/vtkm/cont/testing/UnitTestCellLocatorGeneral.cxx @@ -160,7 +160,7 @@ public: vtkm::Id& cellId, vtkm::Vec3f& pcoords) const { - vtkm::ErrorCode status = locator->FindCell(point, cellId, pcoords); + vtkm::ErrorCode status = locator.FindCell(point, cellId, pcoords); if (status != vtkm::ErrorCode::Success) { this->RaiseError(vtkm::ErrorString(status)); @@ -174,9 +174,6 @@ void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont:: locator.SetCoordinates(dataset.GetCoordinateSystem()); locator.Update(); - const vtkm::cont::CellLocator& curLoc = *locator.GetCurrentLocator(); - std::cout << "using locator: " << typeid(curLoc).name() << "\n"; - vtkm::cont::ArrayHandle expCellIds; vtkm::cont::ArrayHandle expPCoords; vtkm::cont::ArrayHandle points; @@ -186,8 +183,7 @@ void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont:: vtkm::cont::ArrayHandle pcoords; vtkm::worklet::DispatcherMapField dispatcher; - // CellLocatorGeneral is non-copyable. Pass it via a pointer. - dispatcher.Invoke(points, &locator, cellIds, pcoords); + dispatcher.Invoke(points, locator, cellIds, pcoords); auto cellIdPortal = cellIds.ReadPortal(); auto expCellIdsPortal = expCellIds.ReadPortal(); diff --git a/vtkm/exec/CMakeLists.txt b/vtkm/exec/CMakeLists.txt index ecf68f1fc..e56dec4fa 100644 --- a/vtkm/exec/CMakeLists.txt +++ b/vtkm/exec/CMakeLists.txt @@ -16,8 +16,8 @@ set(headers CellFace.h CellInside.h CellInterpolate.h - CellLocator.h - CellLocatorBoundingIntervalHierarchyExec.h + CellLocatorBoundingIntervalHierarchy.h + CellLocatorMultiplexer.h CellLocatorRectilinearGrid.h CellLocatorTwoLevel.h CellLocatorUniformGrid.h @@ -40,6 +40,12 @@ set(header_impls ColorTable.hxx ) +if (NOT VTKm_NO_DEPRECATED_VIRTUAL) + set(headers ${headers} + CellLocator.h + ) +endif() + #----------------------------------------------------------------------------- add_subdirectory(internal) add_subdirectory(arg) diff --git a/vtkm/exec/CellLocator.h b/vtkm/exec/CellLocator.h index f81601d5b..21f52c72a 100644 --- a/vtkm/exec/CellLocator.h +++ b/vtkm/exec/CellLocator.h @@ -16,13 +16,21 @@ #include #include +#ifdef VTKM_NO_DEPRECATED_VIRTUAL +#error "CellLocator with virtual methods is removed. Do not include CellLocator.h" +#endif + namespace vtkm { namespace exec { -class VTKM_ALWAYS_EXPORT CellLocator : public vtkm::VirtualObjectBase +class VTKM_DEPRECATED( + 1.6, + "CellLocator with virtual methods no longer supported. Use CellLocatorGeneral.") + VTKM_ALWAYS_EXPORT CellLocator : public vtkm::VirtualObjectBase { + VTKM_DEPRECATED_SUPPRESS_BEGIN public: VTKM_EXEC_CONT virtual ~CellLocator() noexcept { @@ -48,6 +56,7 @@ public: worklet.RaiseError(vtkm::ErrorString(status)); } } + VTKM_DEPRECATED_SUPPRESS_END }; } // namespace exec diff --git a/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h b/vtkm/exec/CellLocatorBoundingIntervalHierarchy.h similarity index 91% rename from vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h rename to vtkm/exec/CellLocatorBoundingIntervalHierarchy.h index d4e47cf82..d3ed4acc5 100644 --- a/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h +++ b/vtkm/exec/CellLocatorBoundingIntervalHierarchy.h @@ -7,15 +7,14 @@ // the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR // PURPOSE. See the above copyright notice for more information. //============================================================================ -#ifndef vtk_m_cont_CellLocatorBoundingIntervalHierarchyExec_h -#define vtk_m_cont_CellLocatorBoundingIntervalHierarchyExec_h +#ifndef vtk_m_exec_CellLocatorBoundingIntervalHierarchy_h +#define vtk_m_exec_CellLocatorBoundingIntervalHierarchy_h #include #include #include #include #include -#include #include namespace vtkm @@ -62,8 +61,7 @@ struct CellLocatorBoundingIntervalHierarchyNode }; // struct CellLocatorBoundingIntervalHierarchyNode template -class VTKM_ALWAYS_EXPORT CellLocatorBoundingIntervalHierarchyExec final - : public vtkm::exec::CellLocator +class VTKM_ALWAYS_EXPORT CellLocatorBoundingIntervalHierarchy { using NodeArrayHandle = vtkm::cont::ArrayHandle; @@ -71,10 +69,7 @@ class VTKM_ALWAYS_EXPORT CellLocatorBoundingIntervalHierarchyExec final public: VTKM_CONT - CellLocatorBoundingIntervalHierarchyExec() {} - - VTKM_CONT - CellLocatorBoundingIntervalHierarchyExec( + CellLocatorBoundingIntervalHierarchy( const NodeArrayHandle& nodes, const CellIdArrayHandle& cellIds, const CellSetType& cellSet, @@ -89,16 +84,10 @@ public: } - VTKM_EXEC_CONT virtual ~CellLocatorBoundingIntervalHierarchyExec() noexcept override - { - // This must not be defaulted, since defaulted virtual destructors are - // troublesome with CUDA __host__ __device__ markup. - } - VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point, vtkm::Id& cellId, - vtkm::Vec3f& parametric) const override + vtkm::Vec3f& parametric) const { cellId = -1; vtkm::Id nodeIndex = 0; @@ -133,6 +122,11 @@ public: } } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC CellLocatorBoundingIntervalHierarchy* operator->() { return this; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC const CellLocatorBoundingIntervalHierarchy* operator->() const { return this; } + private: enum struct FindCellState { @@ -282,10 +276,10 @@ private: CellIdPortal CellIds; CellSetPortal CellSet; CoordsPortal Coords; -}; // class CellLocatorBoundingIntervalHierarchyExec +}; // class CellLocatorBoundingIntervalHierarchy } // namespace exec } // namespace vtkm -#endif //vtk_m_cont_CellLocatorBoundingIntervalHierarchyExec_h +#endif //vtk_m_exec_CellLocatorBoundingIntervalHierarchy_h diff --git a/vtkm/exec/CellLocatorMultiplexer.h b/vtkm/exec/CellLocatorMultiplexer.h new file mode 100644 index 000000000..88a5f4b55 --- /dev/null +++ b/vtkm/exec/CellLocatorMultiplexer.h @@ -0,0 +1,70 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +//============================================================================ +#ifndef vtk_m_exec_CellLocatorMultiplexer_h +#define vtk_m_exec_CellLocatorMultiplexer_h + +#include +#include + +#include + +namespace vtkm +{ +namespace exec +{ + +namespace detail +{ + +struct FindCellFunctor +{ + template + VTKM_EXEC vtkm::ErrorCode operator()(Locator&& locator, + const vtkm::Vec3f& point, + vtkm::Id& cellId, + vtkm::Vec3f& parametric) const + { + return locator.FindCell(point, cellId, parametric); + } +}; + +} // namespace detail + +template +class VTKM_ALWAYS_EXPORT CellLocatorMultiplexer +{ + vtkm::exec::internal::Variant Locators; + +public: + CellLocatorMultiplexer() = default; + + template + VTKM_CONT CellLocatorMultiplexer(const Locator& locator) + : Locators(locator) + { + } + + VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point, + vtkm::Id& cellId, + vtkm::Vec3f& parametric) const + { + return this->Locators.CastAndCall(detail::FindCellFunctor{}, point, cellId, parametric); + } + + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC CellLocatorMultiplexer* operator->() { return this; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC const CellLocatorMultiplexer* operator->() const { return this; } +}; + +} +} // namespace vtkm::exec + +#endif //vtk_m_exec_CellLocatorMultiplexer_h diff --git a/vtkm/exec/CellLocatorRectilinearGrid.h b/vtkm/exec/CellLocatorRectilinearGrid.h index bf9f7db0c..66a1ae349 100644 --- a/vtkm/exec/CellLocatorRectilinearGrid.h +++ b/vtkm/exec/CellLocatorRectilinearGrid.h @@ -19,7 +19,6 @@ #include #include -#include #include #include @@ -29,55 +28,49 @@ namespace vtkm namespace exec { -template -class VTKM_ALWAYS_EXPORT CellLocatorRectilinearGrid final : public vtkm::exec::CellLocator +class VTKM_ALWAYS_EXPORT CellLocatorRectilinearGrid { private: - using VisitType = vtkm::TopologyElementTagCell; - using IncidentType = vtkm::TopologyElementTagPoint; - using CellSetPortal = vtkm::exec::ConnectivityStructured; - using AxisHandle = vtkm::cont::ArrayHandle; using RectilinearType = vtkm::cont::ArrayHandleCartesianProduct; using AxisPortalType = typename AxisHandle::ReadPortalType; using RectilinearPortalType = typename RectilinearType::ReadPortalType; + VTKM_CONT static vtkm::Id3&& ToId3(vtkm::Id3&& src) { return std::move(src); } + VTKM_CONT static vtkm::Id3 ToId3(vtkm::Id2&& src) { return vtkm::Id3(src[0], src[1], 1); } + VTKM_CONT static vtkm::Id3 ToId3(vtkm::Id&& src) { return vtkm::Id3(src, 1, 1); } + public: - VTKM_CONT - CellLocatorRectilinearGrid(const vtkm::Id planeSize, - const vtkm::Id rowSize, - const vtkm::cont::CellSetStructured& cellSet, - const RectilinearType& coords, - vtkm::cont::DeviceAdapterId device, - vtkm::cont::Token& token) + template + VTKM_CONT CellLocatorRectilinearGrid(const vtkm::Id planeSize, + const vtkm::Id rowSize, + const vtkm::cont::CellSetStructured& cellSet, + const RectilinearType& coords, + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) : PlaneSize(planeSize) , RowSize(rowSize) - , CellSet(cellSet.PrepareForInput(device, VisitType(), IncidentType(), token)) - , Coords(coords.PrepareForInput(device, token)) - , PointDimensions(cellSet.GetPointDimensions()) + , PointDimensions(ToId3(cellSet.GetPointDimensions())) + , Dimensions(dimensions) { - this->AxisPortals[0] = this->Coords.GetFirstPortal(); - this->MinPoint[0] = coords.ReadPortal().GetFirstPortal().Get(0); - this->MaxPoint[0] = coords.ReadPortal().GetFirstPortal().Get(this->PointDimensions[0] - 1); + auto coordsContPortal = coords.ReadPortal(); + RectilinearPortalType coordsExecPortal = coords.PrepareForInput(device, token); + this->AxisPortals[0] = coordsExecPortal.GetFirstPortal(); + this->MinPoint[0] = coordsContPortal.GetFirstPortal().Get(0); + this->MaxPoint[0] = coordsContPortal.GetFirstPortal().Get(this->PointDimensions[0] - 1); - this->AxisPortals[1] = this->Coords.GetSecondPortal(); - this->MinPoint[1] = coords.ReadPortal().GetSecondPortal().Get(0); - this->MaxPoint[1] = coords.ReadPortal().GetSecondPortal().Get(this->PointDimensions[1] - 1); + this->AxisPortals[1] = coordsExecPortal.GetSecondPortal(); + this->MinPoint[1] = coordsContPortal.GetSecondPortal().Get(0); + this->MaxPoint[1] = coordsContPortal.GetSecondPortal().Get(this->PointDimensions[1] - 1); if (dimensions == 3) { - this->AxisPortals[2] = this->Coords.GetThirdPortal(); - this->MinPoint[2] = coords.ReadPortal().GetThirdPortal().Get(0); - this->MaxPoint[2] = coords.ReadPortal().GetThirdPortal().Get(this->PointDimensions[2] - 1); + this->AxisPortals[2] = coordsExecPortal.GetThirdPortal(); + this->MinPoint[2] = coordsContPortal.GetThirdPortal().Get(0); + this->MaxPoint[2] = coordsContPortal.GetThirdPortal().Get(this->PointDimensions[2] - 1); } } - VTKM_EXEC_CONT virtual ~CellLocatorRectilinearGrid() noexcept override - { - // This must not be defaulted, since defaulted virtual destructors are - // troublesome with CUDA __host__ __device__ markup. - } - VTKM_EXEC inline bool IsInside(const vtkm::Vec3f& point) const { @@ -86,7 +79,7 @@ public: inside = false; if (point[1] < this->MinPoint[1] || point[1] > this->MaxPoint[1]) inside = false; - if (dimensions == 3) + if (this->Dimensions == 3) { if (point[2] < this->MinPoint[2] || point[2] > this->MaxPoint[2]) inside = false; @@ -97,7 +90,7 @@ public: VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point, vtkm::Id& cellId, - vtkm::Vec3f& parametric) const override + vtkm::Vec3f& parametric) const { if (!this->IsInside(point)) { @@ -107,7 +100,7 @@ public: // Get the Cell Id from the point. vtkm::Id3 logicalCell(0, 0, 0); - for (vtkm::Int32 dim = 0; dim < dimensions; ++dim) + for (vtkm::Int32 dim = 0; dim < this->Dimensions; ++dim) { // // When searching for points, we consider the max value of the cell @@ -151,16 +144,20 @@ public: return vtkm::ErrorCode::Success; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC CellLocatorRectilinearGrid* operator->() { return this; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC const CellLocatorRectilinearGrid* operator->() const { return this; } + private: vtkm::Id PlaneSize; vtkm::Id RowSize; - CellSetPortal CellSet; - RectilinearPortalType Coords; AxisPortalType AxisPortals[3]; - vtkm::Vec PointDimensions; + vtkm::Id3 PointDimensions; vtkm::Vec3f MinPoint; vtkm::Vec3f MaxPoint; + vtkm::Id Dimensions; }; } //namespace exec } //namespace vtkm diff --git a/vtkm/exec/CellLocatorTwoLevel.h b/vtkm/exec/CellLocatorTwoLevel.h index 77ff9876f..67f6bf955 100644 --- a/vtkm/exec/CellLocatorTwoLevel.h +++ b/vtkm/exec/CellLocatorTwoLevel.h @@ -15,9 +15,9 @@ #include #include -#include #include +#include #include #include #include @@ -83,25 +83,19 @@ namespace exec { //-------------------------------------------------------------------- -template -class VTKM_ALWAYS_EXPORT CellLocatorTwoLevel : public vtkm::exec::CellLocator +template +class VTKM_ALWAYS_EXPORT CellLocatorTwoLevel { private: using DimVec3 = vtkm::internal::cl_uniform_bins::DimVec3; using FloatVec3 = vtkm::internal::cl_uniform_bins::FloatVec3; template - using ArrayPortalConst = typename vtkm::cont::ArrayHandle::ReadPortalType; + using ReadPortal = typename vtkm::cont::ArrayHandle::ReadPortalType; using CoordsPortalType = typename vtkm::cont::CoordinateSystem::MultiplexerArrayType::ReadPortalType; - using CellSetP2CExecType = decltype( - std::declval().PrepareForInput(std::declval(), - vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{}, - std::declval())); - // TODO: This function may return false positives for non 3D cells as the // tests are done on the projection of the point on the cell. Extra checks // should be added to test if the point actually falls on the cell. @@ -129,6 +123,7 @@ private: } public: + template VTKM_CONT CellLocatorTwoLevel(const vtkm::internal::cl_uniform_bins::Grid& topLevelGrid, const vtkm::cont::ArrayHandle& leafDimensions, const vtkm::cont::ArrayHandle& leafStartIndex, @@ -153,16 +148,8 @@ public: { } - VTKM_EXEC_CONT virtual ~CellLocatorTwoLevel() noexcept override - { - // This must not be defaulted, since defaulted virtual destructors are - // troublesome with CUDA __host__ __device__ markup. - } - VTKM_EXEC - vtkm::ErrorCode FindCell(const FloatVec3& point, - vtkm::Id& cellId, - FloatVec3& parametric) const override + vtkm::ErrorCode FindCell(const FloatVec3& point, vtkm::Id& cellId, FloatVec3& parametric) const { using namespace vtkm::internal::cl_uniform_bins; @@ -213,17 +200,22 @@ public: return vtkm::ErrorCode::CellNotFound; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC CellLocatorTwoLevel* operator->() { return this; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC const CellLocatorTwoLevel* operator->() const { return this; } + private: vtkm::internal::cl_uniform_bins::Grid TopLevel; - ArrayPortalConst LeafDimensions; - ArrayPortalConst LeafStartIndex; + ReadPortal LeafDimensions; + ReadPortal LeafStartIndex; - ArrayPortalConst CellStartIndex; - ArrayPortalConst CellCount; - ArrayPortalConst CellIds; + ReadPortal CellStartIndex; + ReadPortal CellCount; + ReadPortal CellIds; - CellSetP2CExecType CellSet; + CellStructureType CellSet; CoordsPortalType Coords; }; } diff --git a/vtkm/exec/CellLocatorUniformGrid.h b/vtkm/exec/CellLocatorUniformGrid.h index 84f8610b0..502a68d30 100644 --- a/vtkm/exec/CellLocatorUniformGrid.h +++ b/vtkm/exec/CellLocatorUniformGrid.h @@ -11,6 +11,7 @@ #define vtkm_exec_celllocatoruniformgrid_h #include +#include #include #include #include @@ -18,7 +19,6 @@ #include #include -#include #include namespace vtkm @@ -27,35 +27,22 @@ namespace vtkm namespace exec { -template -class VTKM_ALWAYS_EXPORT CellLocatorUniformGrid final : public vtkm::exec::CellLocator +class VTKM_ALWAYS_EXPORT CellLocatorUniformGrid { -private: - using VisitType = vtkm::TopologyElementTagCell; - using IncidentType = vtkm::TopologyElementTagPoint; - using CellSetPortal = vtkm::exec::ConnectivityStructured; - public: VTKM_CONT CellLocatorUniformGrid(const vtkm::Id3 cellDims, - const vtkm::Id3 pointDims, const vtkm::Vec3f origin, const vtkm::Vec3f invSpacing, const vtkm::Vec3f maxPoint) : CellDims(cellDims) - , PointDims(pointDims) + , MaxCellIds(vtkm::Max(cellDims - vtkm::Id3(1), vtkm::Id3(0))) , Origin(origin) , InvSpacing(invSpacing) , MaxPoint(maxPoint) { } - VTKM_EXEC_CONT virtual ~CellLocatorUniformGrid() noexcept override - { - // This must not be defaulted, since defaulted virtual destructors are - // troublesome with CUDA __host__ __device__ markup. - } - VTKM_EXEC inline bool IsInside(const vtkm::Vec3f& point) const { bool inside = true; @@ -71,7 +58,7 @@ public: VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point, vtkm::Id& cellId, - vtkm::Vec3f& parametric) const override + vtkm::Vec3f& parametric) const { if (!this->IsInside(point)) { @@ -86,21 +73,8 @@ public: temp = temp * this->InvSpacing; //make sure that if we border the upper edge, we sample the correct cell - logicalCell = temp; - if (logicalCell[0] == this->CellDims[0]) - { - logicalCell[0]--; - } - if (logicalCell[1] == this->CellDims[1]) - { - logicalCell[1]--; - } - if (logicalCell[2] == this->CellDims[2]) - { - logicalCell[2]--; - } - if (dimensions == 2) - logicalCell[2] = 0; + logicalCell = vtkm::Min(vtkm::Id3(temp), this->MaxCellIds); + cellId = (logicalCell[2] * this->CellDims[1] + logicalCell[1]) * this->CellDims[0] + logicalCell[0]; parametric = temp - logicalCell; @@ -108,9 +82,14 @@ public: return vtkm::ErrorCode::Success; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC CellLocatorUniformGrid* operator->() { return this; } + VTKM_DEPRECATED(1.6, "Locators no longer pointers. Use . operator.") + VTKM_EXEC const CellLocatorUniformGrid* operator->() const { return this; } + private: vtkm::Id3 CellDims; - vtkm::Id3 PointDims; + vtkm::Id3 MaxCellIds; vtkm::Vec3f Origin; vtkm::Vec3f InvSpacing; vtkm::Vec3f MaxPoint; diff --git a/vtkm/filter/ParticleDensityNearestGridPoint.hxx b/vtkm/filter/ParticleDensityNearestGridPoint.hxx index 3de3f54db..4c9d20b95 100644 --- a/vtkm/filter/ParticleDensityNearestGridPoint.hxx +++ b/vtkm/filter/ParticleDensityNearestGridPoint.hxx @@ -37,7 +37,7 @@ public: vtkm::Vec3f parametric; // Find the cell containing the point - if (locator->FindCell(point, cellId, parametric) == vtkm::ErrorCode::Success) + if (locator.FindCell(point, cellId, parametric) == vtkm::ErrorCode::Success) { // increment density density.Add(cellId, 1); diff --git a/vtkm/filter/testing/CMakeLists.txt b/vtkm/filter/testing/CMakeLists.txt index bdb9763b1..538297194 100644 --- a/vtkm/filter/testing/CMakeLists.txt +++ b/vtkm/filter/testing/CMakeLists.txt @@ -94,6 +94,25 @@ if (VTKm_ENABLE_RENDERING) ) endif() +if ((TARGET vtkm::cuda) OR (TARGET vtkm::kokkos_cuda)) + # CUDA architecture has a limited amount of memory available for constants. The CUDA + # compiler uses this space to hold constants for some optimizations. However, for large + # kernels, the number of constants needed might be larger than the constant space + # available. For these conditions, you have to disable this form of optimization with + # the -Xptxas --disable-optimizer-constants flags. + # TODO: Find a more elegant way to do this. Either figure out a way around this problem + # or add more general flags to vtkm_library/vtkm_unit_tests for sources with "large" kernels. + set(large_kernel_sources + RegressionTestStreamline.cxx + UnitTestLagrangianFilter.cxx + UnitTestStreamlineFilter.cxx + UnitTestStreamSurfaceFilter.cxx + ) + set_source_files_properties(${large_kernel_sources} PROPERTIES + COMPILE_OPTIONS "-Xptxas;--disable-optimizer-constants" + ) +endif() + vtkm_unit_tests( SOURCES ${unit_tests} LIBRARIES ${libraries} diff --git a/vtkm/internal/ExportMacros.h b/vtkm/internal/ExportMacros.h index 7a5e55776..5b1456ab9 100644 --- a/vtkm/internal/ExportMacros.h +++ b/vtkm/internal/ExportMacros.h @@ -89,8 +89,8 @@ #define VTKM_ALWAYS_EXPORT #define VTKM_NEVER_EXPORT #else -#define VTKM_ALWAYS_EXPORT [[gnu::visibility("default")]] -#define VTKM_NEVER_EXPORT [[gnu::visibility("hidden")]] +#define VTKM_ALWAYS_EXPORT __attribute__((visibility("default"))) +#define VTKM_NEVER_EXPORT __attribute__((visibility("hidden"))) #endif // cuda 7.5 doesn't support static const or static constexpr variables diff --git a/vtkm/rendering/raytracing/ConnectivityTracer.cxx b/vtkm/rendering/raytracing/ConnectivityTracer.cxx index 30eb5ea18..59f3a74f3 100644 --- a/vtkm/rendering/raytracing/ConnectivityTracer.cxx +++ b/vtkm/rendering/raytracing/ConnectivityTracer.cxx @@ -528,7 +528,7 @@ public: query_distance += bumpDistance; vtkm::Vec location = origin + rdir * (query_distance); vtkm::Vec pcoords; - locator->FindCell(location, cellId, pcoords); + locator.FindCell(location, cellId, pcoords); } currentCell = cellId; diff --git a/vtkm/worklet/Probe.h b/vtkm/worklet/Probe.h index 324b8755b..0fc63d86f 100644 --- a/vtkm/worklet/Probe.h +++ b/vtkm/worklet/Probe.h @@ -12,13 +12,12 @@ #include #include -#include +#include +#include #include #include #include -#include -#include #include #include @@ -48,11 +47,21 @@ public: vtkm::Id& cellId, vtkm::Vec3f& pcoords) const { - locator->FindCell(point, cellId, pcoords); + locator.FindCell(point, cellId, pcoords); } }; private: + struct RunSelectLocator + { + template + void operator()(const LocatorType& locator, Probe& worklet, const PointsType& points) const + { + worklet.Invoke( + FindCellWorklet{}, points, locator, worklet.CellIds, worklet.ParametricCoordinates); + } + }; + template void RunImpl(const CellSetType& cells, const vtkm::cont::CoordinateSystem& coords, @@ -60,14 +69,7 @@ private: { this->InputCellSet = vtkm::cont::DynamicCellSet(cells); - vtkm::cont::CellLocatorGeneral locator; - locator.SetCellSet(this->InputCellSet); - locator.SetCoordinates(coords); - locator.Update(); - - vtkm::worklet::DispatcherMapField dispatcher; - // CellLocatorGeneral is non-copyable. Pass it via a pointer. - dispatcher.Invoke(points, &locator, this->CellIds, this->ParametricCoordinates); + vtkm::cont::CastAndCallCellLocatorChooser(cells, coords, RunSelectLocator{}, *this, points); } //============================================================================ @@ -151,8 +153,8 @@ private: this->CellIds); this->ParametricCoordinates.Allocate(points.GetNumberOfValues()); - vtkm::worklet::DispatcherMapTopology dispatcher; - dispatcher.Invoke(cells, coords, points, this->CellIds, this->ParametricCoordinates); + this->Invoke( + ProbeUniformPoints{}, cells, coords, points, this->CellIds, this->ParametricCoordinates); } //============================================================================ @@ -298,8 +300,7 @@ public: vtkm::cont::ArrayHandle GetHiddenPointsField() const { vtkm::cont::ArrayHandle field; - vtkm::worklet::DispatcherMapField dispatcher; - dispatcher.Invoke(this->CellIds, field); + this->Invoke(HiddenPointsWorklet{}, this->CellIds, field); return field; } @@ -331,8 +332,7 @@ public: vtkm::cont::ArrayHandle GetHiddenCellsField(CellSetType cellset) const { vtkm::cont::ArrayHandle field; - vtkm::worklet::DispatcherMapTopology dispatcher; - dispatcher.Invoke(cellset, this->CellIds, field); + this->Invoke(HiddenCellsWorklet{}, cellset, this->CellIds, field); return field; } @@ -343,6 +343,8 @@ private: vtkm::cont::ArrayHandle CellIds; vtkm::cont::ArrayHandle ParametricCoordinates; vtkm::cont::DynamicCellSet InputCellSet; + + vtkm::cont::Invoker Invoke; }; } } // vtkm::worklet diff --git a/vtkm/worklet/particleadvection/GridEvaluatorStatus.h b/vtkm/worklet/particleadvection/GridEvaluatorStatus.h index 76621c0c7..c3b238944 100644 --- a/vtkm/worklet/particleadvection/GridEvaluatorStatus.h +++ b/vtkm/worklet/particleadvection/GridEvaluatorStatus.h @@ -15,7 +15,6 @@ #include #include #include -#include #include #include #include diff --git a/vtkm/worklet/particleadvection/GridEvaluators.h b/vtkm/worklet/particleadvection/GridEvaluators.h index 708998a63..177c95758 100644 --- a/vtkm/worklet/particleadvection/GridEvaluators.h +++ b/vtkm/worklet/particleadvection/GridEvaluators.h @@ -16,7 +16,7 @@ #include #include #include -#include +#include #include #include #include @@ -45,7 +45,7 @@ public: ExecutionGridEvaluator() = default; VTKM_CONT - ExecutionGridEvaluator(std::shared_ptr locator, + ExecutionGridEvaluator(const vtkm::cont::CellLocatorGeneral& locator, std::shared_ptr interpolationHelper, const vtkm::Bounds& bounds, const FieldType& field, @@ -57,7 +57,7 @@ public: , GhostCells(ghostCells.PrepareForInput(device, token)) , HaveGhostCells(ghostCells.GetNumberOfValues() > 0) , InterpolationHelper(interpolationHelper->PrepareForExecution(device, token)) - , Locator(locator->PrepareForExecution(device, token)) + , Locator(locator.PrepareForExecution(device, token)) { } @@ -67,7 +67,7 @@ public: vtkm::Id cellId; Point parametric; - Locator->FindCell(point, cellId, parametric); + this->Locator.FindCell(point, cellId, parametric); if (cellId == -1) return false; @@ -94,7 +94,7 @@ public: const vtkm::FloatDefault& time, vtkm::VecVariable& out) const { - vtkm::Id cellId; + vtkm::Id cellId = -1; Point parametric; GridEvaluatorStatus status; @@ -105,7 +105,7 @@ public: status.SetTemporalBounds(); } - Locator->FindCell(point, cellId, parametric); + this->Locator.FindCell(point, cellId, parametric); if (cellId == -1) { status.SetFail(); @@ -150,7 +150,7 @@ private: GhostCellPortal GhostCells; bool HaveGhostCells; const vtkm::exec::CellInterpolationHelper* InterpolationHelper; - const vtkm::exec::CellLocator* Locator; + typename vtkm::cont::CellLocatorGeneral::ExecObjType Locator; }; template @@ -214,55 +214,23 @@ private: VTKM_CONT void InitializeLocator(const vtkm::cont::CoordinateSystem& coordinates, const vtkm::cont::DynamicCellSet& cellset) { + this->Locator.SetCoordinates(coordinates); + this->Locator.SetCellSet(cellset); + this->Locator.Update(); if (cellset.IsSameType(Structured2DType()) || cellset.IsSameType(Structured3DType())) { - if (coordinates.GetData().IsType()) - { - vtkm::cont::CellLocatorUniformGrid locator; - locator.SetCoordinates(coordinates); - locator.SetCellSet(cellset); - locator.Update(); - this->Locator = std::make_shared(locator); - } - else if (coordinates.GetData().IsType()) - { - vtkm::cont::CellLocatorRectilinearGrid locator; - locator.SetCoordinates(coordinates); - locator.SetCellSet(cellset); - locator.Update(); - this->Locator = std::make_shared(locator); - } - else - { - // Default to using an locator for explicit meshes. - vtkm::cont::CellLocatorTwoLevel locator; - locator.SetCoordinates(coordinates); - locator.SetCellSet(cellset); - locator.Update(); - this->Locator = std::make_shared(locator); - } vtkm::cont::StructuredCellInterpolationHelper interpolationHelper(cellset); this->InterpolationHelper = std::make_shared(interpolationHelper); } else if (cellset.IsSameType(vtkm::cont::CellSetSingleType<>())) { - vtkm::cont::CellLocatorTwoLevel locator; - locator.SetCoordinates(coordinates); - locator.SetCellSet(cellset); - locator.Update(); - this->Locator = std::make_shared(locator); vtkm::cont::SingleCellTypeInterpolationHelper interpolationHelper(cellset); this->InterpolationHelper = std::make_shared(interpolationHelper); } else if (cellset.IsSameType(vtkm::cont::CellSetExplicit<>())) { - vtkm::cont::CellLocatorTwoLevel locator; - locator.SetCoordinates(coordinates); - locator.SetCellSet(cellset); - locator.Update(); - this->Locator = std::make_shared(locator); vtkm::cont::ExplicitCellInterpolationHelper interpolationHelper(cellset); this->InterpolationHelper = std::make_shared(interpolationHelper); @@ -275,7 +243,7 @@ private: FieldType Field; GhostCellArrayType GhostCellArray; std::shared_ptr InterpolationHelper; - std::shared_ptr Locator; + vtkm::cont::CellLocatorGeneral Locator; }; } //namespace particleadvection diff --git a/vtkm/worklet/spatialstructure/BoundingIntervalHierarchy.h b/vtkm/worklet/spatialstructure/BoundingIntervalHierarchy.h index f2c5702be..d859b0594 100644 --- a/vtkm/worklet/spatialstructure/BoundingIntervalHierarchy.h +++ b/vtkm/worklet/spatialstructure/BoundingIntervalHierarchy.h @@ -24,7 +24,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/vtkm/worklet/testing/UnitTestBoundingIntervalHierarchy.cxx b/vtkm/worklet/testing/UnitTestBoundingIntervalHierarchy.cxx index 267074c77..18dec9f8f 100644 --- a/vtkm/worklet/testing/UnitTestBoundingIntervalHierarchy.cxx +++ b/vtkm/worklet/testing/UnitTestBoundingIntervalHierarchy.cxx @@ -50,7 +50,7 @@ struct BoundingIntervalHierarchyTester : public vtkm::worklet::WorkletMapField { vtkm::Vec3f parametric; vtkm::Id cellId; - bih->FindCell(point, cellId, parametric); + bih.FindCell(point, cellId, parametric); return (1 - static_cast(expectedId == cellId)); } }; // struct BoundingIntervalHierarchyTester