Commit 2382c4d6 authored by Kenneth Moreland's avatar Kenneth Moreland Committed by Kitware Robot
Browse files

Merge topic 'pointlocator-general-interface'

4459ab91 Merge branch 'master' into 'pointlocator-general-interface'
51fd4a11 Fix warning about __host__/__device__ on default constructor
6f75cd00 Fix crash in CUDA compiler
439beaae Make point locator tests have consistent devices
33f1f2dd Make sure all source files are listed in CMake
367ca3e2 Correct error of grabbing reference of stack variable
693c8ea8 Update PointLocatorUniformGrid.h
d8ff2ba7

 Update PointLocator.h to add a newline at the end of file.
...
Acked-by: Kitware Robot's avatarKitware Robot <kwrobot@kitware.com>
Acked-by: Li-Ta Lo's avatarLi-Ta Lo <ollie@lanl.gov>
Merge-request: !1292
parents 9a1b44cc 4459ab91
......@@ -84,6 +84,7 @@ set(headers
FieldRangeGlobalCompute.h
ImplicitFunctionHandle.h
MultiBlock.h
PointLocator.h
PointLocatorUniformGrid.h
RuntimeDeviceInformation.h
RuntimeDeviceTracker.h
......
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2018 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_PointLocator_h
#define vtk_m_cont_PointLocator_h
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/ExecutionObjectBase.h>
#include <vtkm/exec/PointLocator.h>
namespace vtkm
{
namespace cont
{
class PointLocator : public vtkm::cont::ExecutionObjectBase
{
public:
PointLocator()
: dirty(true)
{
}
vtkm::cont::CoordinateSystem GetCoords() const { return coordinates; }
void SetCoords(const vtkm::cont::CoordinateSystem& coords)
{
coordinates = coords;
dirty = true;
}
virtual void Build() = 0;
void Update()
{
if (dirty)
Build();
dirty = false;
}
template <typename DeviceAdapter>
VTKM_CONT const vtkm::exec::PointLocator* PrepareForExecution(DeviceAdapter) const
{
vtkm::cont::DeviceAdapterId deviceId = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetId();
return PrepareForExecutionImp(deviceId).PrepareForExecution(DeviceAdapter());
}
//VTKM_CONT virtual const vtkm::exec::PointLocator*
using HandleType = vtkm::cont::VirtualObjectHandle<vtkm::exec::PointLocator>;
VTKM_CONT virtual const HandleType PrepareForExecutionImp(
vtkm::cont::DeviceAdapterId deviceId) const = 0;
private:
vtkm::cont::CoordinateSystem coordinates;
bool dirty;
};
}
}
#endif // vtk_m_cont_PointLocator_h
......@@ -23,21 +23,25 @@
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/PointLocator.h>
#include <vtkm/exec/PointLocatorUniformGrid.h>
namespace vtkm
{
namespace worklet
namespace cont
{
template <typename T>
class PointLocatorUniformGrid
class PointLocatorUniformGrid : public vtkm::cont::PointLocator
{
public:
PointLocatorUniformGrid(const vtkm::Vec<T, 3>& _min,
const vtkm::Vec<T, 3>& _max,
PointLocatorUniformGrid(const vtkm::Vec<vtkm::FloatDefault, 3>& _min,
const vtkm::Vec<vtkm::FloatDefault, 3>& _max,
const vtkm::Vec<vtkm::Id, 3>& _dims)
: Min(_min)
: PointLocator()
, Min(_min)
, Max(_max)
, Dims(_dims)
{
......@@ -51,7 +55,9 @@ public:
using ExecutionSignature = void(_1, _2);
VTKM_CONT
BinPointsWorklet(vtkm::Vec<T, 3> _min, vtkm::Vec<T, 3> _max, vtkm::Vec<vtkm::Id, 3> _dims)
BinPointsWorklet(vtkm::Vec<vtkm::FloatDefault, 3> _min,
vtkm::Vec<vtkm::FloatDefault, 3> _max,
vtkm::Vec<vtkm::Id, 3> _dims)
: Min(_min)
, Dims(_dims)
, Dxdydz((_max - Min) / Dims)
......@@ -66,175 +72,126 @@ public:
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<vtkm::FloatDefault, 3> Min;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<T, 3> Dxdydz;
vtkm::Vec<vtkm::FloatDefault, 3> Dxdydz;
};
class UniformGridSearch : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(FieldIn<> query,
WholeArrayIn<> coordIn,
WholeArrayIn<IdType> pointId,
WholeArrayIn<IdType> cellLower,
WholeArrayIn<IdType> cellUpper,
FieldOut<IdType> neighborId,
FieldOut<> distance);
using ExecutionSignature = void(_1, _2, _3, _4, _5, _6, _7);
/// \brief Construct a 3D uniform grid for nearest neighbor search.
///
/// \param coords An ArrayHandle of x, y, z coordinates of input points.
/// \param device Tag for selecting device adapter
VTKM_CONT
UniformGridSearch(const vtkm::Vec<T, 3>& _min,
const vtkm::Vec<T, 3>& _max,
const vtkm::Vec<vtkm::Id, 3>& _dims)
: Min(_min)
, Dims(_dims)
, Dxdydz((_max - _min) / _dims)
struct BuildFunctor
{
BuildFunctor(vtkm::cont::PointLocatorUniformGrid* self)
: Self(self)
{
}
template <typename CoordiVecType,
typename IdPortalType,
typename CoordiPortalType,
typename IdType,
typename CoordiType>
VTKM_EXEC void operator()(const CoordiVecType& queryCoord,
const CoordiPortalType& coordi_Handle,
const IdPortalType& pointId,
const IdPortalType& cellLower,
const IdPortalType& cellUpper,
IdType& nnId,
CoordiType& nnDis) const
template <typename Device>
bool operator()(Device)
{
auto nlayers = vtkm::Max(vtkm::Max(Dims[0], Dims[1]), Dims[2]);
vtkm::Vec<vtkm::Id, 3> xyz = (queryCoord - Min) / Dxdydz;
float min_distance = std::numeric_limits<float>::max();
vtkm::Id neareast = -1;
for (vtkm::Id layer = 0; layer < nlayers; layer++)
{
vtkm::Id minx = vtkm::Max(vtkm::Id(), xyz[0] - layer);
vtkm::Id maxx = vtkm::Min(Dims[0] - 1, xyz[0] + layer);
vtkm::Id miny = vtkm::Max(vtkm::Id(), xyz[1] - layer);
vtkm::Id maxy = vtkm::Min(Dims[1] - 1, xyz[1] + layer);
vtkm::Id minz = vtkm::Max(vtkm::Id(), xyz[2] - layer);
vtkm::Id maxz = vtkm::Min(Dims[2] - 1, xyz[2] + layer);
for (auto i = minx; i <= maxx; i++)
{
for (auto j = miny; j <= maxy; j++)
{
for (auto k = minz; k <= maxz; k++)
{
if (i == (xyz[0] + layer) || i == (xyz[0] - layer) || j == (xyz[1] + layer) ||
j == (xyz[1] - layer) || k == (xyz[2] + layer) || k == (xyz[2] - layer))
{
auto cellid = i + j * Dims[0] + k * Dims[0] * Dims[1];
auto lower = cellLower.Get(cellid);
auto upper = cellUpper.Get(cellid);
for (auto index = lower; index < upper; index++)
{
auto pointid = pointId.Get(index);
auto point = coordi_Handle.Get(pointid);
auto dx = point[0] - queryCoord[0];
auto dy = point[1] - queryCoord[1];
auto dz = point[2] - queryCoord[2];
auto distance2 = dx * dx + dy * dy + dz * dz;
if (distance2 < min_distance)
{
neareast = pointid;
min_distance = distance2;
nlayers = layer + 2;
}
}
}
}
}
}
}
nnId = neareast;
nnDis = vtkm::Sqrt(min_distance);
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<Device>;
// Save training data points.
Algorithm::Copy(this->Self->GetCoords().GetData(), this->Self->coords);
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(
0, 1, this->Self->coords.GetNumberOfValues());
Algorithm::Copy(pointCounting, this->Self->pointIds);
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(this->Self->Min, this->Self->Max, this->Self->Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet, Device> dispatchCellId(cellIdWorklet);
dispatchCellId.Invoke(this->Self->coords, this->Self->cellIds);
// Group points of the same cell together by sorting them according to the cell ids
Algorithm::SortByKey(this->Self->cellIds, this->Self->pointIds);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(
0, 1, this->Self->Dims[0] * this->Self->Dims[1] * this->Self->Dims[2]);
Algorithm::UpperBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellUpper);
Algorithm::LowerBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellLower);
return true;
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<T, 3> Dxdydz;
vtkm::cont::PointLocatorUniformGrid* Self;
};
/// \brief Construct a 3D uniform grid for nearest neighbor search.
///
/// \param coords An ArrayHandle of x, y, z coordinates of input points.
/// \param device Tag for selecting device adapter
template <typename DeviceAdapter>
void Build(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, DeviceAdapter device)
void Build() override
{
(void)device;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(0, 1, coords.GetNumberOfValues());
Algorithm::Copy(pointCounting, PointIds);
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(Min, Max, Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet, DeviceAdapter> dispatchCellId(
cellIdWorklet);
dispatchCellId.Invoke(coords, CellIds);
// Group points of the same cell together by sorting them according to the cell ids
Algorithm::SortByKey(CellIds, PointIds);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(0, 1, Dims[0] * Dims[1] * Dims[2]);
Algorithm::UpperBounds(CellIds, cell_ids_counting, CellUpper);
Algorithm::LowerBounds(CellIds, cell_ids_counting, CellLower);
}
BuildFunctor functor(this);
bool success = vtkm::cont::TryExecute(functor);
if (!success)
{
throw vtkm::cont::ErrorExecution("Could not build point locator structure");
}
};
/// \brief Nearest neighbor search using a Uniform Grid
///
/// Parallel search of nearesat neighbor for each point in the \c queryPoints in the set of
/// \c coords. Returns neareast neighbot in \c nearestNeighborIds and distances to nearest
/// neighbor in \c distances.
///
/// \param coords Point coordinates for training dataset.
/// \param queryPoints Point coordinates to query for nearest neighbors.
/// \param nearestNeighborIds Neareast neighbor in the training dataset for each points in
/// the test set
/// \param distances Distance between query points and their nearest neighbors.
/// \param device Tag for selecting device adapter.
template <typename DeviceAdapter>
void FindNearestPoint(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& queryPoints,
vtkm::cont::ArrayHandle<vtkm::Id>& nearestNeighborIds,
vtkm::cont::ArrayHandle<T>& distances,
DeviceAdapter device)
{
(void)device;
UniformGridSearch uniformGridSearch(Min, Max, Dims);
using HandleType = vtkm::cont::VirtualObjectHandle<vtkm::exec::PointLocator>;
vtkm::worklet::DispatcherMapField<UniformGridSearch, DeviceAdapter> searchDispatcher(
uniformGridSearch);
searchDispatcher.Invoke(
queryPoints, coords, PointIds, CellLower, CellUpper, nearestNeighborIds, distances);
struct PrepareForExecutionFunctor
{
template <typename DeviceAdapter>
VTKM_CONT void operator()(DeviceAdapter,
const vtkm::cont::PointLocatorUniformGrid& self,
HandleType& handle) const
{
//vtkm::exec::PointLocatorUniformGrid* locator =
vtkm::exec::PointLocatorUniformGrid<DeviceAdapter>* h =
new vtkm::exec::PointLocatorUniformGrid<DeviceAdapter>(
self.Min,
self.Max,
self.Dims,
self.coords.PrepareForInput(DeviceAdapter()),
self.pointIds.PrepareForInput(DeviceAdapter()),
self.cellLower.PrepareForInput(DeviceAdapter()),
self.cellUpper.PrepareForInput(DeviceAdapter()));
handle.Reset(h);
//return handle.PrepareForExecution(DeviceAdapter());
}
};
VTKM_CONT
//const vtkm::exec::PointLocator *
const HandleType PrepareForExecutionImp(vtkm::cont::DeviceAdapterId deviceId) const override
{
// TODO: call VirtualObjectHandle::PrepareForExecution() and return vtkm::exec::PointLocator
// TODO: how to convert deviceId back to DeviceAdapter tag?
//using DeviceList = vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda,
// vtkm::cont::DeviceAdapterTagTBB,
// vtkm::cont::DeviceAdapterTagSerial>;
using DeviceList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG;
//HandleType ExecHandle; // = new HandleType(locator, false);
vtkm::cont::internal::FindDeviceAdapterTagAndCall(
deviceId, DeviceList(), PrepareForExecutionFunctor(), *this, ExecHandle);
return ExecHandle;
//return ExecHandle.PrepareForExecution(DeviceAdapter());
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<T, 3> Max;
vtkm::Vec<vtkm::FloatDefault, 3> Min;
vtkm::Vec<vtkm::FloatDefault, 3> Max;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::cont::ArrayHandle<vtkm::Id> PointIds;
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
vtkm::cont::ArrayHandle<vtkm::Id> CellLower;
vtkm::cont::ArrayHandle<vtkm::Id> CellUpper;
// TODO: how to convert CoordinateSystem to ArrayHandle<Vec<Float, 3>>?
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> coords;
vtkm::cont::ArrayHandle<vtkm::Id> pointIds;
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<vtkm::Id> cellLower;
vtkm::cont::ArrayHandle<vtkm::Id> cellUpper;
// TODO: std::unique_ptr/std::shared_ptr?
mutable HandleType ExecHandle;
};
}
}
......
......@@ -27,8 +27,6 @@
int UnitTestCudaPointLocatorUniformGrid(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
return vtkm::cont::testing::Testing::Run(
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagSerial>());
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagCuda>());
}
......@@ -25,9 +25,6 @@
int UnitTestOpenMPPointLocatorUniformGrid(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagOpenMP{});
return vtkm::cont::testing::Testing::Run(
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagOpenMP>());
}
......@@ -27,8 +27,6 @@
int UnitTestSerialPointLocatorUniformGrid(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagSerial{});
return vtkm::cont::testing::Testing::Run(
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagSerial>());
}
......@@ -27,8 +27,6 @@
int UnitTestTBBPointLocatorUniformGrid(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagTBB{});
return vtkm::cont::testing::Testing::Run(
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagSerial>());
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagTBB>());
}
......@@ -21,17 +21,20 @@
#ifndef vtk_m_cont_testing_TestingPointLocatorUniformGrid_h
#define vtk_m_cont_testing_TestingPointLocatorUniformGrid_h
//#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_SERIAL
#include <random>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/cont/PointLocatorUniformGrid.h>
#include <vtkm/exec/PointLocatorUniformGrid.h>
////brute force method /////
template <typename CoordiVecT, typename CoordiPortalT, typename CoordiT>
VTKM_EXEC_CONT vtkm::Id NNSVerify3D(CoordiVecT qc, CoordiPortalT coordiPortal, CoordiT& dis)
VTKM_EXEC_CONT vtkm::Id NNSVerify3D(CoordiVecT qc, CoordiPortalT coordiPortal, CoordiT& dis2)
{
dis = std::numeric_limits<CoordiT>::max();
dis2 = std::numeric_limits<CoordiT>::max();
vtkm::Id nnpIdx = -1;
for (vtkm::Int32 i = 0; i < coordiPortal.GetNumberOfValues(); i++)
......@@ -39,12 +42,11 @@ VTKM_EXEC_CONT vtkm::Id NNSVerify3D(CoordiVecT qc, CoordiPortalT coordiPortal, C
CoordiT splitX = coordiPortal.Get(i)[0];
CoordiT splitY = coordiPortal.Get(i)[1];
CoordiT splitZ = coordiPortal.Get(i)[2];
CoordiT _dis =
vtkm::Sqrt((splitX - qc[0]) * (splitX - qc[0]) + (splitY - qc[1]) * (splitY - qc[1]) +
(splitZ - qc[2]) * (splitZ - qc[2]));
if (_dis < dis)
CoordiT _dis2 = (splitX - qc[0]) * (splitX - qc[0]) + (splitY - qc[1]) * (splitY - qc[1]) +
(splitZ - qc[2]) * (splitZ - qc[2]);
if (_dis2 < dis2)
{
dis = _dis;
dis2 = _dis2;
nnpIdx = i;
}
}
......@@ -75,6 +77,30 @@ public:
}
};
class PointLocatorUniformGridWorklet : public vtkm::worklet::WorkletMapField
{
public:
typedef void ControlSignature(FieldIn<> qcIn,
ExecObject locator,
FieldOut<> nnIdOut,
FieldOut<> nnDistOut);
typedef void ExecutionSignature(_1, _2, _3, _4);
VTKM_CONT
PointLocatorUniformGridWorklet() {}
// TODO: change IdType, it is used for other purpose.
template <typename CoordiVecType, typename Locator, typename IdType, typename CoordiType>
VTKM_EXEC void operator()(const CoordiVecType& qc,
const Locator& locator,
IdType& nnIdOut,
CoordiType& nnDis) const
{
locator->FindNearestNeighbor(qc, nnIdOut, nnDis);
}
};
template <typename DeviceAdapter>
class TestingPointLocatorUniformGrid
{
......@@ -82,8 +108,8 @@ public:
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
void TestTest() const
{
vtkm::Int32 nTrainingPoints = 1000;
vtkm::Int32 nTestingPoint = 1000;
vtkm::Int32 nTrainingPoints = 5;
vtkm::Int32 nTestingPoint = 1;
std::vector<vtkm::Vec<vtkm::Float32, 3>> coordi;
......@@ -97,10 +123,17 @@ public:
}
auto coordi_Handle = vtkm::cont::make_ArrayHandle(coordi);
vtkm::worklet::PointLocatorUniformGrid<vtkm::Float32> uniformGrid(
vtkm::cont::CoordinateSystem coord("points", coordi_Handle);
// TODO: locator needs to be a pointer to have runtime polymorphism.
//vtkm::cont::PointLocator * locator = new vtkm::cont::PointLocatorUniformGrid(
// { 0.0f, 0.0f, 0.0f }, { 10.0f, 10.0f, 10.0f }, { 5, 5, 5 });
vtkm::cont::PointLocatorUniformGrid locator(
{ 0.0f, 0.0f, 0.0f }, { 10.0f, 10.0f, 10.0f }, { 5, 5, 5 });
uniformGrid.Build(coordi_Handle, DeviceAdapter());
locator.SetCoords(coord);
locator.Build();
///// randomly generate testing points/////
std::vector<vtkm::Vec<vtkm::Float32, 3>> qcVec;
for (vtkm::Int32 i = 0; i < nTestingPoint; i++)
{
......@@ -111,9 +144,12 @@ public:
vtkm::cont::ArrayHandle<vtkm::Id> nnId_Handle;
vtkm::cont::ArrayHandle<vtkm::Float32> nnDis_Handle;
uniformGrid.FindNearestPoint(
coordi_Handle, qc_Handle, nnId_Handle, nnDis_Handle, DeviceAdapter());
PointLocatorUniformGridWorklet pointLocatorUniformGridWorklet;
vtkm::worklet::DispatcherMapField<PointLocatorUniformGridWorklet, DeviceAdapter>
locatorDispatcher(pointLocatorUniformGridWorklet);
locatorDispatcher.Invoke(qc_Handle, locator, nnId_Handle, nnDis_Handle);
// brute force
vtkm::cont::ArrayHandle<vtkm::Id> bfnnId_Handle;
vtkm::cont::ArrayHandle<vtkm::Float32> bfnnDis_Handle;
NearestNeighborSearchBruteForce3DWorklet nnsbf3dWorklet;
......@@ -122,7 +158,7 @@ public:
nnsbf3DDispatcher.Invoke(
qc_Handle, vtkm::cont::make_ArrayHandle(coordi), bfnnId_Handle, bfnnDis_Handle);
///// verfity search result /////
///// verify search result /////
bool passTest = true;
for (vtkm::Int32 i = 0; i < nTestingPoint; i++)
{
......@@ -138,11 +174,14 @@ public:
passTest = false;
}
}
VTKM_TEST_ASSERT(passTest, "Uniform Grid NN search result incorrect.");
}
void operator()() const { this->TestTest(); }
void operator()() const
{
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(DeviceAdapter());
this->TestTest();
}
};
#endif // vtk_m_cont_testing_TestingPointLocatorUniformGrid_h
......@@ -34,6 +34,8 @@ set(headers
FunctorBase.h
Jacobian.h
ParametricCoordinates.h
PointLocator.h
PointLocatorUniformGrid.h