Merge topic 'cell-locators-not-virtual'

047d79672 Fix CUDA compilation error with Lagrangian filter
aec9890e9 Remove check for CellLocator.h in SourceInInstall test
2399741a9 Change Probe filter to use CellLocatorChooser
e61c54f87 Add CellLocatorChooser
23c823d4b Fix compile errors and warnings related to new CellLocator structure
47429a316 Fix export issues with CellLocatorBase
8922f600e Use GNU attributes for deprecated
0a5f5d55c Remove virtual methods from cell locators

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Nickolas Davis <nadavi@sandia.gov>
Merge-request: !2271
This commit is contained in:
Kenneth Moreland 2021-02-17 00:15:53 +00:00 committed by Kitware Robot
commit 3aa2c7521a
43 changed files with 1089 additions and 430 deletions

@ -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

@ -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

@ -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

@ -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()

@ -14,6 +14,8 @@ namespace vtkm
{
namespace cont
{
VTKM_DEPRECATED_SUPPRESS_BEGIN
CellLocator::~CellLocator() = default;
VTKM_DEPRECATED_SUPPRESS_END
}
}

@ -12,6 +12,7 @@
#include <vtkm/cont/vtkm_cont_export.h>
#include <vtkm/Deprecated.h>
#include <vtkm/Types.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DynamicCellSet.h>
@ -19,13 +20,21 @@
#include <vtkm/exec/CellLocator.h>
#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

@ -21,7 +21,7 @@
#include <vtkm/cont/ArrayHandleTransform.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/ErrorBadDevice.h>
#include <vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h>
#include <vtkm/exec/CellLocatorBoundingIntervalHierarchy.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/worklet/WorkletMapField.h>
@ -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 <typename CellSetType>
void operator()(
const CellSetType& cellset,
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token,
vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator>& bihExec,
const vtkm::cont::ArrayHandle<vtkm::exec::CellLocatorBoundingIntervalHierarchyNode>& nodes,
const vtkm::cont::ArrayHandle<vtkm::Id>& 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<CellSetType>;
ExecutionType* execObject =
new ExecutionType(nodes, processedCellIds, cellset, coords, device, token);
bihExec.Reset(execObject);
execObject = vtkm::exec::CellLocatorBoundingIntervalHierarchy<CellSetType>(
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

@ -16,8 +16,11 @@
#include <vtkm/Types.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleTransform.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/cont/internal/CellLocatorBase.h>
#include <vtkm/exec/CellLocatorBoundingIntervalHierarchy.h>
#include <vtkm/exec/CellLocatorMultiplexer.h>
#include <vtkm/worklet/spatialstructure/BoundingIntervalHierarchy.h>
@ -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<CellLocatorBoundingIntervalHierarchy>
{
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorBoundingIntervalHierarchy>;
public:
using SupportedCellSets = VTKM_DEFAULT_CELL_SET_LIST;
using CellLocatorExecList =
vtkm::ListTransform<SupportedCellSets, vtkm::exec::CellLocatorBoundingIntervalHierarchy>;
using ExecObjType = vtkm::ListApply<CellLocatorExecList, vtkm::exec::CellLocatorMultiplexer>;
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<vtkm::exec::CellLocatorBoundingIntervalHierarchyNode> Nodes;
vtkm::cont::ArrayHandle<vtkm::Id> ProcessedCellIds;
mutable vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator> ExecutionObjectHandle;
friend Superclass;
VTKM_CONT void Build();
struct MakeExecObject;
};
} // namespace cont

@ -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 <vtkm/cont/CastAndCall.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>
namespace vtkm
{
namespace cont
{
namespace detail
{
template <typename CellSetType, typename CoordinateSystemArrayType>
struct CellLocatorChooserImpl
{
using type = vtkm::cont::CellLocatorTwoLevel;
};
using UniformArray = vtkm::cont::ArrayHandleUniformPointCoordinates;
template <>
struct CellLocatorChooserImpl<vtkm::cont::CellSetStructured<3>, UniformArray>
{
using type = vtkm::cont::CellLocatorUniformGrid;
};
using RectilinearArray =
vtkm::cont::ArrayHandleCartesianProduct<vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>;
template <>
struct CellLocatorChooserImpl<vtkm::cont::CellSetStructured<3>, 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 <typename CellSetType, typename CoordinateSystemArrayType>
using CellLocatorChooser =
typename detail::CellLocatorChooserImpl<CellSetType, CoordinateSystemArrayType>::type;
namespace detail
{
struct CastAndCallCellLocatorChooserFunctor
{
template <typename CellLocatorType, typename Functor, typename... Args>
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>(args)...);
}
template <typename CellSetType, typename Functor, typename... Args>
void operator()(const CellSetType& cellSet,
const vtkm::cont::CoordinateSystem& coordinateSystem,
Functor&& functor,
Args&&... args) const
{
this->CallFunctorWithLocator<vtkm::cont::CellLocatorTwoLevel>(
cellSet, coordinateSystem, std::forward<Functor>(functor), std::forward<Args>(args)...);
}
template <typename Functor, typename... Args>
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<detail::UniformArray>())
{
this->CallFunctorWithLocator<vtkm::cont::CellLocatorUniformGrid>(
cellSet, coordinateSystem, std::forward<Functor>(functor), std::forward<Args>(args)...);
}
else if (coordArray.IsType<detail::RectilinearArray>())
{
this->CallFunctorWithLocator<vtkm::cont::CellLocatorRectilinearGrid>(
cellSet, coordinateSystem, std::forward<Functor>(functor), std::forward<Args>(args)...);
}
else
{
this->CallFunctorWithLocator<vtkm::cont::CellLocatorTwoLevel>(
cellSet, coordinateSystem, std::forward<Functor>(functor), std::forward<Args>(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 <typename CellSetList, typename Functor, typename... Args>
VTKM_CONT void CastAndCallCellLocatorChooser(
const vtkm::cont::DynamicCellSetBase<CellSetList>& cellSet,
const vtkm::cont::CoordinateSystem& coordinateSystem,
Functor&& functor,
Args&&... args)
{
vtkm::cont::CastAndCall(cellSet,
detail::CastAndCallCellLocatorChooserFunctor{},
coordinateSystem,
std::forward<Functor>(functor),
std::forward<Args>(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 <typename Functor, typename... Args>
VTKM_CONT void CastAndCallCellLocatorChooser(const vtkm::cont::DataSet& dataSet,
Functor&& functor,
Args&&... args)
{
CastAndCallCellLocatorChooser(dataSet.GetCellSet(),
dataSet.GetCoordinateSystem(),
std::forward<Functor>(functor),
std::forward<Args>(args)...);
}
}
} // namespace vtkm::cont
#endif //vtk_m_cont_CellLocatorChooser_h

@ -20,39 +20,19 @@
namespace
{
VTKM_CONT
void DefaultConfigurator(std::unique_ptr<vtkm::cont::CellLocator>& locator,
const vtkm::cont::DynamicCellSet& cellSet,
const vtkm::cont::CoordinateSystem& coords)
template <typename LocatorImplType, typename LocatorVariantType>
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::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>;
if (cellSet.IsType<StructuredCellSet>() && coords.GetData().IsType<UniformCoordinates>())
constexpr vtkm::IdComponent LOCATOR_INDEX =
LocatorVariantType::template GetIndexOf<LocatorImplType>();
if (locatorVariant.GetIndex() != LOCATOR_INDEX)
{
if (!dynamic_cast<vtkm::cont::CellLocatorUniformGrid*>(locator.get()))
{
locator.reset(new vtkm::cont::CellLocatorUniformGrid);
}
locatorVariant = LocatorImplType{};
}
else if (cellSet.IsType<StructuredCellSet>() && coords.GetData().IsType<RectilinearCoordinates>())
{
if (!dynamic_cast<vtkm::cont::CellLocatorRectilinearGrid*>(locator.get()))
{
locator.reset(new vtkm::cont::CellLocatorRectilinearGrid);
}
}
else if (!dynamic_cast<vtkm::cont::CellLocatorTwoLevel*>(locator.get()))
{
locator.reset(new vtkm::cont::CellLocatorTwoLevel);
}
locator->SetCellSet(cellSet);
locator->SetCoordinates(coords);
LocatorImplType& locatorImpl = locatorVariant.template Get<LOCATOR_INDEX>();
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::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>;
vtkm::cont::DynamicCellSet cellSet = this->GetCellSet();
vtkm::cont::CoordinateSystem coords = this->GetCoordinates();
if (cellSet.IsType<StructuredCellSet>() && coords.GetData().IsType<UniformCoordinates>())
{
BuildForType<vtkm::cont::CellLocatorUniformGrid>(*this, this->LocatorImpl);
}
else if (cellSet.IsType<StructuredCellSet>() && coords.GetData().IsType<RectilinearCoordinates>())
{
BuildForType<vtkm::cont::CellLocatorRectilinearGrid>(*this, this->LocatorImpl);
}
else
{
BuildForType<vtkm::cont::CellLocatorTwoLevel>(*this, this->LocatorImpl);
}
}
VTKM_CONT CellLocatorGeneral::~CellLocatorGeneral() = default;
struct CellLocatorGeneral::PrepareFunctor
{
template <typename LocatorType>
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

@ -10,7 +10,13 @@
#ifndef vtk_m_cont_CellLocatorGeneral_h
#define vtk_m_cont_CellLocatorGeneral_h
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/exec/CellLocatorMultiplexer.h>
#include <vtkm/cont/internal/Variant.h>
#include <functional>
#include <memory>
@ -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<CellLocatorGeneral>
{
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorGeneral>;
public:
VTKM_CONT CellLocatorGeneral();
using ContLocatorList = vtkm::List<vtkm::cont::CellLocatorUniformGrid,
vtkm::cont::CellLocatorRectilinearGrid,
vtkm::cont::CellLocatorTwoLevel>;
VTKM_CONT ~CellLocatorGeneral() override;
using ExecLocatorList =
vtkm::List<vtkm::cont::internal::ExecutionObjectType<vtkm::cont::CellLocatorUniformGrid>,
vtkm::cont::internal::ExecutionObjectType<vtkm::cont::CellLocatorRectilinearGrid>,
vtkm::cont::internal::ExecutionObjectType<vtkm::cont::CellLocatorTwoLevel>>;
/// Get the current underlying cell locator
///
VTKM_CONT const vtkm::cont::CellLocator* GetCurrentLocator() const { return this->Locator.get(); }
using ExecObjType = vtkm::ListApply<ExecLocatorList, vtkm::exec::CellLocatorMultiplexer>;
/// 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<vtkm::cont::CellLocator>&,
const vtkm::cont::DynamicCellSet&,
const vtkm::cont::CoordinateSystem&);
VTKM_CONT void SetConfigurator(const std::function<ConfiguratorSignature>& configurator)
{
this->Configurator = configurator;
}
VTKM_CONT const std::function<ConfiguratorSignature>& 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<vtkm::cont::CellLocator> Locator;
std::function<ConfiguratorSignature> Configurator;
vtkm::cont::internal::ListAsVariant<ContLocatorList> LocatorImpl;
friend Superclass;
VTKM_CONT void Build();
struct PrepareFunctor;
};
}
} // vtkm::cont

@ -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<vtkm::FloatDefault>;
using RectilinearType = vtkm::cont::ArrayHandleCartesianProduct<AxisHandle, AxisHandle, AxisHandle>;
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<Structured3DType>(),
this->GetCoordinates().GetData().AsArrayHandle<RectilinearType>(),
device,
token);
this->ExecutionObjectHandle.Reset(execObject);
return ExecObjType(this->PlaneSize,
this->RowSize,
this->GetCellSet().template Cast<Structured3DType>(),
this->GetCoordinates().GetData().template AsArrayHandle<RectilinearType>(),
device,
token);
}
else
{
using ExecutionType = vtkm::exec::CellLocatorRectilinearGrid<2>;
ExecutionType* execObject =
new ExecutionType(this->PlaneSize,
this->RowSize,
this->GetCellSet().Cast<Structured2DType>(),
this->GetCoordinates().GetData().AsArrayHandle<RectilinearType>(),
device,
token);
this->ExecutionObjectHandle.Reset(execObject);
return ExecObjType(this->PlaneSize,
this->RowSize,
this->GetCellSet().template Cast<Structured2DType>(),
this->GetCoordinates().GetData().template AsArrayHandle<RectilinearType>(),
device,
token);
}
return this->ExecutionObjectHandle.PrepareForExecution(device, token);
}
} //namespace cont
} //namespace vtkm

@ -10,27 +10,35 @@
#ifndef vtkm_cont_CellLocatorRectilinearGrid_h
#define vtkm_cont_CellLocatorRectilinearGrid_h
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/cont/internal/CellLocatorBase.h>
#include <vtkm/exec/CellLocatorRectilinearGrid.h>
namespace vtkm
{
namespace cont
{
class VTKM_CONT_EXPORT CellLocatorRectilinearGrid : public vtkm::cont::CellLocator
class VTKM_CONT_EXPORT CellLocatorRectilinearGrid
: public vtkm::cont::internal::CellLocatorBase<CellLocatorRectilinearGrid>
{
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorRectilinearGrid>;
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<vtkm::FloatDefault>;
using RectilinearType =
vtkm::cont::ArrayHandleCartesianProduct<AxisHandle, AxisHandle, AxisHandle>;
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<vtkm::exec::CellLocator> ExecutionObjectHandle;
friend Superclass;
VTKM_CONT void Build();
};
} //namespace cont

@ -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<CellSetType>(self.TopLevel,
self.LeafDimensions,
self.LeafStartIndex,
self.CellStartIndex,
self.CellCount,
self.CellIds,
cellSet,
self.GetCoordinates(),
device,
token);
self.ExecutionObjectHandle.Reset(execObject);
using CellStructuredType = CellSetContToExec<CellSetType>;
execObject = vtkm::exec::CellLocatorTwoLevel<CellStructuredType>(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;
}
//----------------------------------------------------------------------------

@ -11,9 +11,11 @@
#define vtk_m_cont_CellLocatorTwoLevel_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/cont/CellSetList.h>
#include <vtkm/cont/internal/CellLocatorBase.h>
#include <vtkm/exec/CellLocatorMultiplexer.h>
#include <vtkm/exec/CellLocatorTwoLevel.h>
@ -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<CellLocatorTwoLevel>
{
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorTwoLevel>;
template <typename CellSetCont>
using CellSetContToExec =
typename CellSetCont::template ExecConnectivityType<vtkm::TopologyElementTagCell,
vtkm::TopologyElementTagPoint>;
public:
using SupportedCellSets = VTKM_DEFAULT_CELL_SET_LIST;
using CellExecObjectList = vtkm::ListTransform<SupportedCellSets, CellSetContToExec>;
using CellLocatorExecList =
vtkm::ListTransform<CellExecObjectList, vtkm::exec::CellLocatorTwoLevel>;
using ExecObjType = vtkm::ListApply<CellLocatorExecList, vtkm::exec::CellLocatorMultiplexer>;
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<vtkm::Id> CellCount;
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
mutable vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator> ExecutionObjectHandle;
struct MakeExecObject;
};
}
} // vtkm::cont

@ -13,15 +13,10 @@
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/exec/CellLocatorUniformGrid.h>
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

@ -10,27 +10,24 @@
#ifndef vtkm_cont_CellLocatorUniformGrid_h
#define vtkm_cont_CellLocatorUniformGrid_h
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/cont/internal/CellLocatorBase.h>
#include <vtkm/exec/CellLocatorUniformGrid.h>
namespace vtkm
{
namespace cont
{
class VTKM_CONT_EXPORT CellLocatorUniformGrid : public vtkm::cont::CellLocator
class VTKM_CONT_EXPORT CellLocatorUniformGrid
: public vtkm::cont::internal::CellLocatorBase<CellLocatorUniformGrid>
{
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorUniformGrid>;
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<vtkm::exec::CellLocator> ExecutionObjectHandle;
friend Superclass;
VTKM_CONT void Build();
};
}
} // vtkm::cont

@ -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 <typename ExecutionObject, typename Device>
template <typename ExecutionObject, typename Device = vtkm::cont::DeviceAdapterId>
using ExecutionObjectType = decltype(CallPrepareForExecution(std::declval<ExecutionObject>(),
Device{},
std::declval<Device>(),
std::declval<vtkm::cont::Token&>()));
} // namespace internal

@ -100,6 +100,7 @@ public:
bool acquireOwnership = true,
DeviceAdapterList devices = DeviceAdapterList())
{
VTKM_DEPRECATED_SUPPRESS_BEGIN
VTKM_STATIC_ASSERT_MSG((std::is_base_of<VirtualBaseType, VirtualDerivedType>::value),
"Tried to bind a type that is not a subclass of the base class.");
@ -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

@ -17,6 +17,7 @@ set(headers
AtomicInterfaceExecution.h
Buffer.h
CastInvalidValue.h
CellLocatorBase.h
ConnectivityExplicitInternals.h
DeviceAdapterAlgorithmGeneral.h
DeviceAdapterMemoryManager.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 <vtkm/cont/vtkm_cont_export.h>
#include <vtkm/Types.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DynamicCellSet.h>
#include <vtkm/cont/ExecutionObjectBase.h>
#ifndef VTKM_NO_DEPRECATED_VIRTUAL
// To support deprecated implementation
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/exec/CellLocator.h>
#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 <typename LocatorType>
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 <typename LocatorType>
struct CellLocatorBaseWrapperPrepareForExecutionFunctor
{
template <typename Device>
VTKM_CONT bool operator()(Device device,
vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator>& execHandle,
const LocatorType& locator,
vtkm::cont::Token& token)
{
auto execObject = locator.PrepareForExecution(device, token);
using WrapType = CellLocatorBaseExecWrapper<decltype(execObject)>;
execHandle.Reset(new WrapType(execObject));
return true;
}
};
template <typename Derived>
class VTKM_ALWAYS_EXPORT CellLocatorBaseWrapper : public vtkm::cont::CellLocator
{
Derived Locator;
mutable vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator> 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<Derived>{},
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 <typename Derived>
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<Derived>() const
{
return detail::CellLocatorBaseWrapper<Derived>(reinterpret_cast<const Derived&>(*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<Derived*>(const_cast<CellLocatorBase*>(this))->Build();
this->Modified = false;
}
}
template <typename Device>
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object.")
vtkm::cont::internal::ExecutionObjectType<Derived, Device> 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

@ -55,6 +55,7 @@ set(unit_tests
UnitTestArrayHandleVirtual.cxx
UnitTestArrayHandleXGCCoordinates.cxx
UnitTestArrayPortalToIterators.cxx
UnitTestCellLocatorChooser.cxx
UnitTestCellLocatorGeneral.cxx
UnitTestCellSet.cxx
UnitTestCellSetExplicit.cxx

@ -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));

@ -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));

@ -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));

@ -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 <vtkm/cont/CellLocatorChooser.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DataSetBuilderRectilinear.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/exec/CellInterpolate.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterPermutation.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <ctime>
#include <random>
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<vtkm::FloatDefault> coordGen(1.0f / 128.0f, 1.0f / 32.0f);
vtkm::cont::ArrayHandle<vtkm::FloatDefault> 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<PointType> 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<vtkm::Id>& cellIds)
{
return ScatterType(cellIds);
}
template <typename CellShapeTagType, typename PointsVecType>
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<vtkm::Id>& cellIds,
vtkm::cont::ArrayHandle<PointType>& pcoords,
vtkm::cont::ArrayHandle<PointType>& wcoords)
{
vtkm::Id numberOfCells = ds.GetNumberOfCells();
std::uniform_int_distribution<vtkm::Id> cellIdGen(0, numberOfCells - 1);
std::uniform_real_distribution<vtkm::FloatDefault> 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<ParametricToWorldCoordinates> 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 <typename LocatorType>
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 <typename CellSetType, typename CoordinateSystemArrayType>
void TestWithDataSet(const vtkm::cont::DataSet& dataset)
{
VTKM_TEST_ASSERT(dataset.GetCellSet().IsType<CellSetType>());
VTKM_TEST_ASSERT(dataset.GetCoordinateSystem().GetData().IsType<CoordinateSystemArrayType>());
vtkm::cont::CellLocatorChooser<CellSetType, CoordinateSystemArrayType> locator;
locator.SetCellSet(dataset.GetCellSet());
locator.SetCoordinates(dataset.GetCoordinateSystem());
locator.Update();
vtkm::cont::ArrayHandle<vtkm::Id> expCellIds;
vtkm::cont::ArrayHandle<PointType> expPCoords;
vtkm::cont::ArrayHandle<PointType> points;
GenerateRandomInput(dataset, 128, expCellIds, expPCoords, points);
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::worklet::DispatcherMapField<FindCellWorklet> 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::CellSetStructured<3>, vtkm::cont::ArrayHandleUniformPointCoordinates>(
MakeTestDataSetUniform());
std::cout << "Test Rectilinear dataset\n";
TestWithDataSet<
vtkm::cont::CellSetStructured<3>,
vtkm::cont::ArrayHandleCartesianProduct<vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>>>(
MakeTestDataSetRectilinear());
std::cout << "Test Curvilinear dataset\n";
TestWithDataSet<vtkm::cont::CellSetStructured<3>, vtkm::cont::ArrayHandle<PointType>>(
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);
}

@ -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<vtkm::Id> expCellIds;
vtkm::cont::ArrayHandle<PointType> expPCoords;
vtkm::cont::ArrayHandle<PointType> points;
@ -186,8 +183,7 @@ void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont::
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::worklet::DispatcherMapField<FindCellWorklet> 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();

@ -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)

@ -16,13 +16,21 @@
#include <vtkm/VirtualObjectBase.h>
#include <vtkm/exec/FunctorBase.h>
#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

@ -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 <vtkm/TopologyElementTag.h>
#include <vtkm/VecFromPortalPermute.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/CellLocator.h>
#include <vtkm/exec/ParametricCoordinates.h>
namespace vtkm
@ -62,8 +61,7 @@ struct CellLocatorBoundingIntervalHierarchyNode
}; // struct CellLocatorBoundingIntervalHierarchyNode
template <typename CellSetType>
class VTKM_ALWAYS_EXPORT CellLocatorBoundingIntervalHierarchyExec final
: public vtkm::exec::CellLocator
class VTKM_ALWAYS_EXPORT CellLocatorBoundingIntervalHierarchy
{
using NodeArrayHandle =
vtkm::cont::ArrayHandle<vtkm::exec::CellLocatorBoundingIntervalHierarchyNode>;
@ -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

@ -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 <vtkm/ErrorCode.h>
#include <vtkm/TypeList.h>
#include <vtkm/exec/internal/Variant.h>
namespace vtkm
{
namespace exec
{
namespace detail
{
struct FindCellFunctor
{
template <typename Locator>
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 <typename... LocatorTypes>
class VTKM_ALWAYS_EXPORT CellLocatorMultiplexer
{
vtkm::exec::internal::Variant<LocatorTypes...> Locators;
public:
CellLocatorMultiplexer() = default;
template <typename Locator>
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

@ -19,7 +19,6 @@
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/CellLocator.h>
#include <vtkm/exec/ConnectivityStructured.h>
#include <vtkm/exec/ParametricCoordinates.h>
@ -29,55 +28,49 @@ namespace vtkm
namespace exec
{
template <vtkm::IdComponent dimensions>
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<VisitType, IncidentType, dimensions>;
using AxisHandle = vtkm::cont::ArrayHandle<vtkm::FloatDefault>;
using RectilinearType =
vtkm::cont::ArrayHandleCartesianProduct<AxisHandle, AxisHandle, AxisHandle>;
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<dimensions>& cellSet,
const RectilinearType& coords,
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token)
template <vtkm::IdComponent dimensions>
VTKM_CONT CellLocatorRectilinearGrid(const vtkm::Id planeSize,
const vtkm::Id rowSize,
const vtkm::cont::CellSetStructured<dimensions>& 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<vtkm::Id, dimensions> PointDimensions;
vtkm::Id3 PointDimensions;
vtkm::Vec3f MinPoint;
vtkm::Vec3f MaxPoint;
vtkm::Id Dimensions;
};
} //namespace exec
} //namespace vtkm

@ -15,9 +15,9 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DeviceAdapterTag.h>
#include <vtkm/Math.h>
#include <vtkm/TopologyElementTag.h>
#include <vtkm/Types.h>
#include <vtkm/VecFromPortalPermute.h>
#include <vtkm/VecTraits.h>
@ -83,25 +83,19 @@ namespace exec
{
//--------------------------------------------------------------------
template <typename CellSetType>
class VTKM_ALWAYS_EXPORT CellLocatorTwoLevel : public vtkm::exec::CellLocator
template <typename CellStructureType>
class VTKM_ALWAYS_EXPORT CellLocatorTwoLevel
{
private:
using DimVec3 = vtkm::internal::cl_uniform_bins::DimVec3;
using FloatVec3 = vtkm::internal::cl_uniform_bins::FloatVec3;
template <typename T>
using ArrayPortalConst = typename vtkm::cont::ArrayHandle<T>::ReadPortalType;
using ReadPortal = typename vtkm::cont::ArrayHandle<T>::ReadPortalType;
using CoordsPortalType =
typename vtkm::cont::CoordinateSystem::MultiplexerArrayType::ReadPortalType;
using CellSetP2CExecType = decltype(
std::declval<CellSetType>().PrepareForInput(std::declval<vtkm::cont::DeviceAdapterId>(),
vtkm::TopologyElementTagCell{},
vtkm::TopologyElementTagPoint{},
std::declval<vtkm::cont::Token&>()));
// 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 <typename CellSetType>
VTKM_CONT CellLocatorTwoLevel(const vtkm::internal::cl_uniform_bins::Grid& topLevelGrid,
const vtkm::cont::ArrayHandle<DimVec3>& leafDimensions,
const vtkm::cont::ArrayHandle<vtkm::Id>& 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<DimVec3> LeafDimensions;
ArrayPortalConst<vtkm::Id> LeafStartIndex;
ReadPortal<DimVec3> LeafDimensions;
ReadPortal<vtkm::Id> LeafStartIndex;
ArrayPortalConst<vtkm::Id> CellStartIndex;
ArrayPortalConst<vtkm::Id> CellCount;
ArrayPortalConst<vtkm::Id> CellIds;
ReadPortal<vtkm::Id> CellStartIndex;
ReadPortal<vtkm::Id> CellCount;
ReadPortal<vtkm::Id> CellIds;
CellSetP2CExecType CellSet;
CellStructureType CellSet;
CoordsPortalType Coords;
};
}

@ -11,6 +11,7 @@
#define vtkm_exec_celllocatoruniformgrid_h
#include <vtkm/Bounds.h>
#include <vtkm/Math.h>
#include <vtkm/TopologyElementTag.h>
#include <vtkm/Types.h>
#include <vtkm/VecFromPortalPermute.h>
@ -18,7 +19,6 @@
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/CellLocator.h>
#include <vtkm/exec/ParametricCoordinates.h>
namespace vtkm
@ -27,35 +27,22 @@ namespace vtkm
namespace exec
{
template <vtkm::IdComponent dimensions>
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<VisitType, IncidentType, dimensions>;
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;

@ -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);

@ -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}

@ -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

@ -528,7 +528,7 @@ public:
query_distance += bumpDistance;
vtkm::Vec<FloatType, 3> location = origin + rdir * (query_distance);
vtkm::Vec<vtkm::FloatDefault, 3> pcoords;
locator->FindCell(location, cellId, pcoords);
locator.FindCell(location, cellId, pcoords);
}
currentCell = cellId;

@ -12,13 +12,12 @@
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocatorGeneral.h>
#include <vtkm/cont/CellLocatorChooser.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/CellInterpolate.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
@ -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 <typename LocatorType, typename PointsType>
void operator()(const LocatorType& locator, Probe& worklet, const PointsType& points) const
{
worklet.Invoke(
FindCellWorklet{}, points, locator, worklet.CellIds, worklet.ParametricCoordinates);
}
};
template <typename CellSetType, typename PointsType, typename PointsStorage>
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<FindCellWorklet> 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<ProbeUniformPoints> 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<vtkm::UInt8> GetHiddenPointsField() const
{
vtkm::cont::ArrayHandle<vtkm::UInt8> field;
vtkm::worklet::DispatcherMapField<HiddenPointsWorklet> dispatcher;
dispatcher.Invoke(this->CellIds, field);
this->Invoke(HiddenPointsWorklet{}, this->CellIds, field);
return field;
}
@ -331,8 +332,7 @@ public:
vtkm::cont::ArrayHandle<vtkm::UInt8> GetHiddenCellsField(CellSetType cellset) const
{
vtkm::cont::ArrayHandle<vtkm::UInt8> field;
vtkm::worklet::DispatcherMapTopology<HiddenCellsWorklet> dispatcher;
dispatcher.Invoke(cellset, this->CellIds, field);
this->Invoke(HiddenCellsWorklet{}, cellset, this->CellIds, field);
return field;
}
@ -343,6 +343,8 @@ private:
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
vtkm::cont::ArrayHandle<vtkm::Vec3f> ParametricCoordinates;
vtkm::cont::DynamicCellSet InputCellSet;
vtkm::cont::Invoker Invoke;
};
}
} // vtkm::worklet

@ -15,7 +15,6 @@
#include <vtkm/Types.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>

@ -16,7 +16,7 @@
#include <vtkm/Types.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/CellLocatorGeneral.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
@ -45,7 +45,7 @@ public:
ExecutionGridEvaluator() = default;
VTKM_CONT
ExecutionGridEvaluator(std::shared_ptr<vtkm::cont::CellLocator> locator,
ExecutionGridEvaluator(const vtkm::cont::CellLocatorGeneral& locator,
std::shared_ptr<vtkm::cont::CellInterpolationHelper> 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<Point, 2>& 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 <typename FieldType>
@ -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<UniformType>())
{
vtkm::cont::CellLocatorUniformGrid locator;
locator.SetCoordinates(coordinates);
locator.SetCellSet(cellset);
locator.Update();
this->Locator = std::make_shared<vtkm::cont::CellLocatorUniformGrid>(locator);
}
else if (coordinates.GetData().IsType<RectilinearType>())
{
vtkm::cont::CellLocatorRectilinearGrid locator;
locator.SetCoordinates(coordinates);
locator.SetCellSet(cellset);
locator.Update();
this->Locator = std::make_shared<vtkm::cont::CellLocatorRectilinearGrid>(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<vtkm::cont::CellLocatorTwoLevel>(locator);
}
vtkm::cont::StructuredCellInterpolationHelper interpolationHelper(cellset);
this->InterpolationHelper =
std::make_shared<vtkm::cont::StructuredCellInterpolationHelper>(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<vtkm::cont::CellLocatorTwoLevel>(locator);
vtkm::cont::SingleCellTypeInterpolationHelper interpolationHelper(cellset);
this->InterpolationHelper =
std::make_shared<vtkm::cont::SingleCellTypeInterpolationHelper>(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<vtkm::cont::CellLocatorTwoLevel>(locator);
vtkm::cont::ExplicitCellInterpolationHelper interpolationHelper(cellset);
this->InterpolationHelper =
std::make_shared<vtkm::cont::ExplicitCellInterpolationHelper>(interpolationHelper);
@ -275,7 +243,7 @@ private:
FieldType Field;
GhostCellArrayType GhostCellArray;
std::shared_ptr<vtkm::cont::CellInterpolationHelper> InterpolationHelper;
std::shared_ptr<vtkm::cont::CellLocator> Locator;
vtkm::cont::CellLocatorGeneral Locator;
};
} //namespace particleadvection

@ -24,7 +24,7 @@
#include <vtkm/cont/ArrayHandleTransform.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h>
#include <vtkm/exec/CellLocatorBoundingIntervalHierarchy.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>

@ -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<vtkm::IdComponent>(expectedId == cellId));
}
}; // struct BoundingIntervalHierarchyTester