Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-10-05 01:49:02 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Rename CellLocatorUniformBins to CellLocatorTwoLevel

Browse Source 
The new name is more descriptive and prevents confusion with
CellLocatorUniformGrid.
This commit is contained in:
Kenneth Moreland 2020-09-21 15:35:30 -06:00
parent c9b763aca9
commit b012c42beb
22 changed files with 390 additions and 328 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
CTestCustom.cmake.in
View File

@ -55,7 +55,7 @@ list(APPEND CTEST_CUSTOM_WARNING_EXCEPTION
"nvlink warning : .*ArrayPortalVirtual.* has address taken but no possible call to it"
"nvlink warning : .*CellLocatorBoundingIntervalHierarchyExec.* has address taken but no possible call to it"
"nvlink warning : .*CellLocatorRectilinearGrid.* has address taken but no possible call to it"
"nvlink warning : .*CellLocatorUniformBins.* has address taken but no possible call to it"
"nvlink warning : .*CellLocatorTwoLevel.* has address taken but no possible call to it"
"nvlink warning : .*CellLocatorUniformGrid.* has address taken but no possible call to it"
)

3
vtkm/cont/CMakeLists.txt
View File

@ -55,6 +55,7 @@ set(headers
CellLocatorBoundingIntervalHierarchy.h
CellLocatorGeneral.h
CellLocatorRectilinearGrid.h
CellLocatorTwoLevel.h
CellLocatorUniformBins.h
CellLocatorUniformGrid.h
CellSet.h
@ -166,7 +167,7 @@ set(sources
CellLocatorBoundingIntervalHierarchy.cxx
CellLocatorGeneral.cxx
CellLocatorRectilinearGrid.cxx
CellLocatorUniformBins.cxx
CellLocatorTwoLevel.cxx
CellLocatorUniformGrid.cxx
CellSet.cxx
CellSetExplicit.cxx

6
vtkm/cont/CellLocatorGeneral.cxx
View File

@ -13,7 +13,7 @@
#include <vtkm/cont/ArrayHandleCartesianProduct.h>
#include <vtkm/cont/ArrayHandleUniformPointCoordinates.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorUniformBins.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/CellSetStructured.h>
@ -46,9 +46,9 @@ void DefaultConfigurator(std::unique_ptr<vtkm::cont::CellLocator>& locator,
locator.reset(new vtkm::cont::CellLocatorRectilinearGrid);
}
}
else if (!dynamic_cast<vtkm::cont::CellLocatorUniformBins*>(locator.get()))
else if (!dynamic_cast<vtkm::cont::CellLocatorTwoLevel*>(locator.get()))
{
locator.reset(new vtkm::cont::CellLocatorUniformBins);
locator.reset(new vtkm::cont::CellLocatorTwoLevel);
}
locator->SetCellSet(cellSet);

36
vtkm/cont/CellLocatorUniformBins.cxx → vtkm/cont/CellLocatorTwoLevel.cxx
View File

@ -7,7 +7,7 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/CellLocatorUniformBins.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/Algorithm.h>
#include <vtkm/cont/ArrayCopy.h>
@ -345,7 +345,7 @@ namespace cont
//----------------------------------------------------------------------------
/// Builds the cell locator lookup structure
///
VTKM_CONT void CellLocatorUniformBins::Build()
VTKM_CONT void CellLocatorTwoLevel::Build()
{
vtkm::cont::Invoker invoke;
@ -453,53 +453,53 @@ VTKM_CONT void CellLocatorUniformBins::Build()
}
//----------------------------------------------------------------------------
struct CellLocatorUniformBins::MakeExecObject
struct CellLocatorTwoLevel::MakeExecObject
{
template <typename CellSetType, typename DeviceAdapter>
VTKM_CONT void operator()(const CellSetType& cellSet,
DeviceAdapter,
vtkm::cont::Token& token,
const CellLocatorUniformBins& self) const
const CellLocatorTwoLevel& self) const
{
auto execObject =
new vtkm::exec::CellLocatorUniformBins<CellSetType, DeviceAdapter>(self.TopLevel,
self.LeafDimensions,
self.LeafStartIndex,
self.CellStartIndex,
self.CellCount,
self.CellIds,
cellSet,
self.GetCoordinates(),
token);
new vtkm::exec::CellLocatorTwoLevel<CellSetType, DeviceAdapter>(self.TopLevel,
self.LeafDimensions,
self.LeafStartIndex,
self.CellStartIndex,
self.CellCount,
self.CellIds,
cellSet,
self.GetCoordinates(),
token);
self.ExecutionObjectHandle.Reset(execObject);
}
};
struct CellLocatorUniformBins::PrepareForExecutionFunctor
struct CellLocatorTwoLevel::PrepareForExecutionFunctor
{
template <typename DeviceAdapter>
VTKM_CONT bool operator()(DeviceAdapter,
vtkm::cont::Token& token,
const CellLocatorUniformBins& self) const
const CellLocatorTwoLevel& self) const
{
self.GetCellSet().CastAndCall(MakeExecObject{}, DeviceAdapter{}, token, self);
return true;
}
};
VTKM_CONT const vtkm::exec::CellLocator* CellLocatorUniformBins::PrepareForExecution(
VTKM_CONT const vtkm::exec::CellLocator* CellLocatorTwoLevel::PrepareForExecution(
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const
{
if (!vtkm::cont::TryExecuteOnDevice(device, PrepareForExecutionFunctor(), token, *this))
{
throwFailedRuntimeDeviceTransfer("CellLocatorUniformBins", device);
throwFailedRuntimeDeviceTransfer("CellLocatorTwoLevel", device);
}
return this->ExecutionObjectHandle.PrepareForExecution(device, token);
}
//----------------------------------------------------------------------------
void CellLocatorUniformBins::PrintSummary(std::ostream& out) const
void CellLocatorTwoLevel::PrintSummary(std::ostream& out) const
{
out << "DensityL1: " << this->DensityL1 << "\n";
out << "DensityL2: " << this->DensityL2 << "\n";

93
vtkm/cont/CellLocatorTwoLevel.h Normal file
View File

@ -0,0 +1,93 @@
//============================================================================
// 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_CellLocatorTwoLevel_h
#define vtk_m_cont_CellLocatorTwoLevel_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/exec/CellLocatorTwoLevel.h>
namespace vtkm
{
namespace cont
{
/// \brief A locator that uses 2 nested levels of grids.
///
/// `CellLocatorTwoLevel` creates a cell search structure using two levels of structured
/// grids. The first level is a coarse grid that covers the entire region of the data.
/// It is expected that the distributions of dataset cells in this coarse grid will be
/// very uneven. Within each bin of the coarse grid, a second level grid is defined within
/// the spatial bounds of the coarse bin. The size of this second level grid is proportional
/// to the number of cells in the first level. In this way, the second level grids adapt
/// to the distribution of cells being located. The adaption is not perfect, but it is
/// has very good space efficiency and is fast to generate and use.
///
/// The algorithm used in `CellLocatorTwoLevel` is described in the following publication:
///
/// 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
{
public:
CellLocatorTwoLevel()
: DensityL1(32.0f)
, DensityL2(2.0f)
{
}
/// Get/Set the desired approximate number of cells per level 1 bin
///
void SetDensityL1(vtkm::FloatDefault val)
{
this->DensityL1 = val;
this->SetModified();
}
vtkm::FloatDefault GetDensityL1() const { return this->DensityL1; }
/// Get/Set the desired approximate number of cells per level 1 bin
///
void SetDensityL2(vtkm::FloatDefault val)
{
this->DensityL2 = val;
this->SetModified();
}
vtkm::FloatDefault GetDensityL2() const { return this->DensityL2; }
void PrintSummary(std::ostream& out) const;
const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const override;
private:
VTKM_CONT void Build() override;
vtkm::FloatDefault DensityL1, DensityL2;
vtkm::internal::cl_uniform_bins::Grid TopLevel;
vtkm::cont::ArrayHandle<vtkm::internal::cl_uniform_bins::DimVec3> LeafDimensions;
vtkm::cont::ArrayHandle<vtkm::Id> LeafStartIndex;
vtkm::cont::ArrayHandle<vtkm::Id> CellStartIndex;
vtkm::cont::ArrayHandle<vtkm::Id> CellCount;
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
mutable vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator> ExecutionObjectHandle;
struct MakeExecObject;
struct PrepareForExecutionFunctor;
};
}
} // vtkm::cont
#endif // vtk_m_cont_CellLocatorTwoLevel_h

278
vtkm/cont/CellLocatorUniformBins.h
View File

@ -10,283 +10,21 @@
#ifndef vtk_m_cont_CellLocatorUniformBins_h
#define vtk_m_cont_CellLocatorUniformBins_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/VirtualObjectHandle.h>
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/Math.h>
#include <vtkm/Types.h>
#include <vtkm/VecFromPortalPermute.h>
#include <vtkm/VecTraits.h>
namespace vtkm
{
namespace internal
{
namespace cl_uniform_bins
{
using DimensionType = vtkm::Int16;
using DimVec3 = vtkm::Vec<DimensionType, 3>;
using FloatVec3 = vtkm::Vec3f;
struct Grid
{
DimVec3 Dimensions;
FloatVec3 Origin;
FloatVec3 BinSize;
};
struct Bounds
{
FloatVec3 Min;
FloatVec3 Max;
};
VTKM_EXEC inline vtkm::Id ComputeFlatIndex(const DimVec3& idx, const DimVec3 dim)
{
return idx[0] + (dim[0] * (idx[1] + (dim[1] * idx[2])));
}
VTKM_EXEC inline Grid ComputeLeafGrid(const DimVec3& idx, const DimVec3& dim, const Grid& l1Grid)
{
return { dim,
l1Grid.Origin + (static_cast<FloatVec3>(idx) * l1Grid.BinSize),
l1Grid.BinSize / static_cast<FloatVec3>(dim) };
}
template <typename PointsVecType>
VTKM_EXEC inline Bounds ComputeCellBounds(const PointsVecType& points)
{
using CoordsType = typename vtkm::VecTraits<PointsVecType>::ComponentType;
auto numPoints = vtkm::VecTraits<PointsVecType>::GetNumberOfComponents(points);
CoordsType minp = points[0], maxp = points[0];
for (vtkm::IdComponent i = 1; i < numPoints; ++i)
{
minp = vtkm::Min(minp, points[i]);
maxp = vtkm::Max(maxp, points[i]);
}
return { FloatVec3(minp), FloatVec3(maxp) };
}
}
}
} // vtkm::internal::cl_uniform_bins
namespace vtkm
{
namespace exec
{
//--------------------------------------------------------------------
template <typename CellSetType, typename DeviceAdapter>
class VTKM_ALWAYS_EXPORT CellLocatorUniformBins : public vtkm::exec::CellLocator
{
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>::template ExecutionTypes<DeviceAdapter>::PortalConst;
using CoordsPortalType =
typename vtkm::cont::CoordinateSystem::MultiplexerArrayType::ExecutionTypes<
DeviceAdapter>::PortalConst;
using CellSetP2CExecType =
decltype(std::declval<CellSetType>().PrepareForInput(DeviceAdapter{},
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.
template <typename CellShapeTag, typename CoordsType>
VTKM_EXEC static vtkm::ErrorCode PointInsideCell(FloatVec3 point,
CellShapeTag cellShape,
CoordsType cellPoints,
FloatVec3& parametricCoordinates,
bool& inside)
{
auto bounds = vtkm::internal::cl_uniform_bins::ComputeCellBounds(cellPoints);
if (point[0] >= bounds.Min[0] && point[0] <= bounds.Max[0] && point[1] >= bounds.Min[1] &&
point[1] <= bounds.Max[1] && point[2] >= bounds.Min[2] && point[2] <= bounds.Max[2])
{
VTKM_RETURN_ON_ERROR(vtkm::exec::WorldCoordinatesToParametricCoordinates(
cellPoints, point, cellShape, parametricCoordinates));
inside = vtkm::exec::CellInside(parametricCoordinates, cellShape);
}
else
{
inside = false;
}
// Return success error code even point is not inside this cell
return vtkm::ErrorCode::Success;
}
public:
VTKM_CONT CellLocatorUniformBins(const vtkm::internal::cl_uniform_bins::Grid& topLevelGrid,
const vtkm::cont::ArrayHandle<DimVec3>& leafDimensions,
const vtkm::cont::ArrayHandle<vtkm::Id>& leafStartIndex,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellStartIndex,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellCount,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellIds,
const CellSetType& cellSet,
const vtkm::cont::CoordinateSystem& coords,
vtkm::cont::Token& token)
: TopLevel(topLevelGrid)
, LeafDimensions(leafDimensions.PrepareForInput(DeviceAdapter{}, token))
, LeafStartIndex(leafStartIndex.PrepareForInput(DeviceAdapter{}, token))
, CellStartIndex(cellStartIndex.PrepareForInput(DeviceAdapter{}, token))
, CellCount(cellCount.PrepareForInput(DeviceAdapter{}, token))
, CellIds(cellIds.PrepareForInput(DeviceAdapter{}, token))
, CellSet(cellSet.PrepareForInput(DeviceAdapter{},
vtkm::TopologyElementTagCell{},
vtkm::TopologyElementTagPoint{},
token))
, Coords(coords.GetDataAsMultiplexer().PrepareForInput(DeviceAdapter{}, token))
{
}
VTKM_EXEC_CONT virtual ~CellLocatorUniformBins() 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
{
using namespace vtkm::internal::cl_uniform_bins;
cellId = -1;
DimVec3 binId3 = static_cast<DimVec3>((point - this->TopLevel.Origin) / this->TopLevel.BinSize);
if (binId3[0] >= 0 && binId3[0] < this->TopLevel.Dimensions[0] && binId3[1] >= 0 &&
binId3[1] < this->TopLevel.Dimensions[1] && binId3[2] >= 0 &&
binId3[2] < this->TopLevel.Dimensions[2])
{
vtkm::Id binId = ComputeFlatIndex(binId3, this->TopLevel.Dimensions);
auto ldim = this->LeafDimensions.Get(binId);
if (!ldim[0] || !ldim[1] || !ldim[2])
{
return vtkm::ErrorCode::CellNotFound;
}
auto leafGrid = ComputeLeafGrid(binId3, ldim, this->TopLevel);
DimVec3 leafId3 = static_cast<DimVec3>((point - leafGrid.Origin) / leafGrid.BinSize);
// precision issues may cause leafId3 to be out of range so clamp it
leafId3 = vtkm::Max(DimVec3(0), vtkm::Min(ldim - DimVec3(1), leafId3));
vtkm::Id leafStart = this->LeafStartIndex.Get(binId);
vtkm::Id leafId = leafStart + ComputeFlatIndex(leafId3, leafGrid.Dimensions);
vtkm::Id start = this->CellStartIndex.Get(leafId);
vtkm::Id end = start + this->CellCount.Get(leafId);
for (vtkm::Id i = start; i < end; ++i)
{
vtkm::Id cid = this->CellIds.Get(i);
auto indices = this->CellSet.GetIndices(cid);
auto pts = vtkm::make_VecFromPortalPermute(&indices, this->Coords);
FloatVec3 pc;
bool inside;
VTKM_RETURN_ON_ERROR(
PointInsideCell(point, this->CellSet.GetCellShape(cid), pts, pc, inside));
if (inside)
{
cellId = cid;
parametric = pc;
return vtkm::ErrorCode::Success;
}
}
}
return vtkm::ErrorCode::CellNotFound;
}
private:
vtkm::internal::cl_uniform_bins::Grid TopLevel;
ArrayPortalConst<DimVec3> LeafDimensions;
ArrayPortalConst<vtkm::Id> LeafStartIndex;
ArrayPortalConst<vtkm::Id> CellStartIndex;
ArrayPortalConst<vtkm::Id> CellCount;
ArrayPortalConst<vtkm::Id> CellIds;
CellSetP2CExecType CellSet;
CoordsPortalType Coords;
};
}
} // vtkm::exec
#include <vtkm/Deprecated.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
namespace vtkm
{
namespace cont
{
//----------------------------------------------------------------------------
class VTKM_CONT_EXPORT CellLocatorUniformBins : public vtkm::cont::CellLocator
struct VTKM_ALWAYS_EXPORT VTKM_DEPRECATED(1.6, "Replaced with CellLocatorTwoLevel.")
CellLocatorUniformBins : vtkm::cont::CellLocatorTwoLevel
{
public:
CellLocatorUniformBins()
: DensityL1(32.0f)
, DensityL2(2.0f)
{
}
/// Get/Set the desired approximate number of cells per level 1 bin
///
void SetDensityL1(vtkm::FloatDefault val)
{
this->DensityL1 = val;
this->SetModified();
}
vtkm::FloatDefault GetDensityL1() const { return this->DensityL1; }
/// Get/Set the desired approximate number of cells per level 1 bin
///
void SetDensityL2(vtkm::FloatDefault val)
{
this->DensityL2 = val;
this->SetModified();
}
vtkm::FloatDefault GetDensityL2() const { return this->DensityL2; }
void PrintSummary(std::ostream& out) const;
const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const override;
private:
VTKM_CONT void Build() override;
vtkm::FloatDefault DensityL1, DensityL2;
vtkm::internal::cl_uniform_bins::Grid TopLevel;
vtkm::cont::ArrayHandle<vtkm::internal::cl_uniform_bins::DimVec3> LeafDimensions;
vtkm::cont::ArrayHandle<vtkm::Id> LeafStartIndex;
vtkm::cont::ArrayHandle<vtkm::Id> CellStartIndex;
vtkm::cont::ArrayHandle<vtkm::Id> CellCount;
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
mutable vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator> ExecutionObjectHandle;
struct MakeExecObject;
struct PrepareForExecutionFunctor;
};
}
} // vtkm::cont
#endif // vtk_m_cont_CellLocatorUniformBins_h
}
} // namespace vtkm::cont
#endif //vtk_m_cont_CellLocatorUniformBins_h

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

@ -14,7 +14,7 @@ set(unit_tests
UnitTestCudaArrayHandleMultiplexer.cu
UnitTestCudaBitField.cu
UnitTestCudaCellLocatorRectilinearGrid.cu
UnitTestCudaCellLocatorUniformBins.cu
UnitTestCudaCellLocatorTwoLevel.cu
UnitTestCudaCellLocatorUniformGrid.cu
UnitTestCudaComputeRange.cu
UnitTestCudaColorTable.cu

6
vtkm/cont/cuda/testing/UnitTestCudaCellLocatorUniformBins.cu → vtkm/cont/cuda/testing/UnitTestCudaCellLocatorTwoLevel.cu
View File

@ -8,12 +8,12 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/testing/TestingCellLocatorUniformBins.h>
#include <vtkm/cont/testing/TestingCellLocatorTwoLevel.h>
int UnitTestCudaCellLocatorUniformBins(int argc, char* argv[])
int UnitTestCudaCellLocatorTwoLevel(int argc, char* argv[])
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
return vtkm::cont::testing::Testing::Run(
TestingCellLocatorUniformBins<vtkm::cont::DeviceAdapterTagCuda>, argc, argv);
TestingCellLocatorTwoLevel<vtkm::cont::DeviceAdapterTagCuda>, argc, argv);
}

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

@ -14,7 +14,7 @@ set(unit_tests
UnitTestKokkosArrayHandleMultiplexer.cxx
UnitTestKokkosBitField.cxx
UnitTestKokkosCellLocatorRectilinearGrid.cxx
UnitTestKokkosCellLocatorUniformBins.cxx
UnitTestKokkosCellLocatorTwoLevel.cxx
UnitTestKokkosCellLocatorUniformGrid.cxx
UnitTestKokkosComputeRange.cxx
UnitTestKokkosColorTable.cxx

6
vtkm/cont/kokkos/testing/UnitTestKokkosCellLocatorUniformBins.cxx → vtkm/cont/kokkos/testing/UnitTestKokkosCellLocatorTwoLevel.cxx
View File

@ -8,12 +8,12 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/testing/TestingCellLocatorUniformBins.h>
#include <vtkm/cont/testing/TestingCellLocatorTwoLevel.h>
int UnitTestKokkosCellLocatorUniformBins(int argc, char* argv[])
int UnitTestKokkosCellLocatorTwoLevel(int argc, char* argv[])
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagKokkos{});
return vtkm::cont::testing::Testing::Run(
TestingCellLocatorUniformBins<vtkm::cont::DeviceAdapterTagKokkos>, argc, argv);
TestingCellLocatorTwoLevel<vtkm::cont::DeviceAdapterTagKokkos>, argc, argv);
}

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

@ -14,7 +14,7 @@ set(unit_tests
UnitTestOpenMPArrayHandleMultiplexer.cxx
UnitTestOpenMPBitField.cxx
UnitTestOpenMPCellLocatorRectilinearGrid.cxx
UnitTestOpenMPCellLocatorUniformBins.cxx
UnitTestOpenMPCellLocatorTwoLevel.cxx
UnitTestOpenMPCellLocatorUniformGrid.cxx
UnitTestOpenMPColorTable.cxx
UnitTestOpenMPComputeRange.cxx

6
vtkm/cont/openmp/testing/UnitTestOpenMPCellLocatorUniformBins.cxx → vtkm/cont/openmp/testing/UnitTestOpenMPCellLocatorTwoLevel.cxx
View File

@ -8,12 +8,12 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/openmp/DeviceAdapterOpenMP.h>
#include <vtkm/cont/testing/TestingCellLocatorUniformBins.h>
#include <vtkm/cont/testing/TestingCellLocatorTwoLevel.h>
int UnitTestOpenMPCellLocatorUniformBins(int argc, char* argv[])
int UnitTestOpenMPCellLocatorTwoLevel(int argc, char* argv[])
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagOpenMP{});
return vtkm::cont::testing::Testing::Run(
TestingCellLocatorUniformBins<vtkm::cont::DeviceAdapterTagOpenMP>, argc, argv);
TestingCellLocatorTwoLevel<vtkm::cont::DeviceAdapterTagOpenMP>, argc, argv);
}

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

@ -14,7 +14,7 @@ set(unit_tests
UnitTestSerialArrayHandleMultiplexer.cxx
UnitTestSerialBitField.cxx
UnitTestSerialCellLocatorRectilinearGrid.cxx
UnitTestSerialCellLocatorUniformBins.cxx
UnitTestSerialCellLocatorTwoLevel.cxx
UnitTestSerialCellLocatorUniformGrid.cxx
UnitTestSerialComputeRange.cxx
UnitTestSerialColorTable.cxx

6
vtkm/cont/serial/testing/UnitTestSerialCellLocatorUniformBins.cxx → vtkm/cont/serial/testing/UnitTestSerialCellLocatorTwoLevel.cxx
View File

@ -9,12 +9,12 @@
//============================================================================
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/testing/TestingCellLocatorUniformBins.h>
#include <vtkm/cont/testing/TestingCellLocatorTwoLevel.h>
int UnitTestSerialCellLocatorUniformBins(int argc, char* argv[])
int UnitTestSerialCellLocatorTwoLevel(int argc, char* argv[])
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagSerial{});
return vtkm::cont::testing::Testing::Run(
TestingCellLocatorUniformBins<vtkm::cont::DeviceAdapterTagSerial>, argc, argv);
TestingCellLocatorTwoLevel<vtkm::cont::DeviceAdapterTagSerial>, argc, argv);
}

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

@ -14,7 +14,7 @@ set(unit_tests
UnitTestTBBArrayHandleMultiplexer.cxx
UnitTestTBBBitField.cxx
UnitTestTBBCellLocatorRectilinearGrid.cxx
UnitTestTBBCellLocatorUniformBins.cxx
UnitTestTBBCellLocatorTwoLevel.cxx
UnitTestTBBCellLocatorUniformGrid.cxx
UnitTestTBBColorTable.cxx
UnitTestTBBComputeRange.cxx

6
vtkm/cont/tbb/testing/UnitTestTBBCellLocatorUniformBins.cxx → vtkm/cont/tbb/testing/UnitTestTBBCellLocatorTwoLevel.cxx
View File

@ -8,12 +8,12 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/testing/TestingCellLocatorUniformBins.h>
#include <vtkm/cont/testing/TestingCellLocatorTwoLevel.h>
int UnitTestTBBCellLocatorUniformBins(int argc, char* argv[])
int UnitTestTBBCellLocatorTwoLevel(int argc, char* argv[])
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagTBB{});
return vtkm::cont::testing::Testing::Run(
TestingCellLocatorUniformBins<vtkm::cont::DeviceAdapterTagTBB>, argc, argv);
TestingCellLocatorTwoLevel<vtkm::cont::DeviceAdapterTagTBB>, argc, argv);
}

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

@ -15,7 +15,7 @@ set(headers
TestingArrayHandles.h
TestingArrayHandleMultiplexer.h
TestingCellLocatorRectilinearGrid.h
TestingCellLocatorUniformBins.h
TestingCellLocatorTwoLevel.h
TestingCellLocatorUniformGrid.h
TestingColorTable.h
TestingComputeRange.h

12
vtkm/cont/testing/TestingCellLocatorUniformBins.h → vtkm/cont/testing/TestingCellLocatorTwoLevel.h
View File

@ -7,11 +7,11 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_cont_testing_TestingCellLocatorUniformBins_h
#define vtk_m_cont_testing_TestingCellLocatorUniformBins_h
#ifndef vtk_m_cont_testing_TestingCellLocatorTwoLevel_h
#define vtk_m_cont_testing_TestingCellLocatorTwoLevel_h
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/CellLocatorUniformBins.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/cont/testing/Testing.h>
@ -188,7 +188,7 @@ void TestCellLocator(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dim, vtkm::Id number
std::cout << "Testing " << DIMENSIONS << "D dataset with " << ds.GetNumberOfCells() << " cells\n";
vtkm::cont::CellLocatorUniformBins locator;
vtkm::cont::CellLocatorTwoLevel locator;
locator.SetDensityL1(64.0f);
locator.SetDensityL2(1.0f);
locator.SetCellSet(ds.GetCellSet());
@ -222,7 +222,7 @@ void TestCellLocator(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dim, vtkm::Id number
} // anonymous
template <typename DeviceAdapter>
void TestingCellLocatorUniformBins()
void TestingCellLocatorTwoLevel()
{
vtkm::cont::GetRuntimeDeviceTracker().ForceDevice(DeviceAdapter());
@ -234,4 +234,4 @@ void TestingCellLocatorUniformBins()
TestCellLocator(vtkm::Id2(18), 512); // 2D dataset
}
#endif // vtk_m_cont_testing_TestingCellLocatorUniformBins_h
#endif // vtk_m_cont_testing_TestingCellLocatorTwoLevel_h

1
vtkm/exec/CMakeLists.txt
View File

@ -19,6 +19,7 @@ set(headers
CellLocator.h
CellLocatorBoundingIntervalHierarchyExec.h
CellLocatorRectilinearGrid.h
CellLocatorTwoLevel.h
CellLocatorUniformGrid.h
CellMeasure.h
ColorTable.h

229
vtkm/exec/CellLocatorTwoLevel.h Normal file
View File

@ -0,0 +1,229 @@
//============================================================================
// 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_CellLocatorTwoLevel_h
#define vtk_m_exec_CellLocatorTwoLevel_h
#include <vtkm/exec/CellInside.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/Math.h>
#include <vtkm/Types.h>
#include <vtkm/VecFromPortalPermute.h>
#include <vtkm/VecTraits.h>
namespace vtkm
{
namespace internal
{
namespace cl_uniform_bins
{
using DimensionType = vtkm::Int16;
using DimVec3 = vtkm::Vec<DimensionType, 3>;
using FloatVec3 = vtkm::Vec3f;
struct Grid
{
DimVec3 Dimensions;
FloatVec3 Origin;
FloatVec3 BinSize;
};
struct Bounds
{
FloatVec3 Min;
FloatVec3 Max;
};
VTKM_EXEC inline vtkm::Id ComputeFlatIndex(const DimVec3& idx, const DimVec3 dim)
{
return idx[0] + (dim[0] * (idx[1] + (dim[1] * idx[2])));
}
VTKM_EXEC inline Grid ComputeLeafGrid(const DimVec3& idx, const DimVec3& dim, const Grid& l1Grid)
{
return { dim,
l1Grid.Origin + (static_cast<FloatVec3>(idx) * l1Grid.BinSize),
l1Grid.BinSize / static_cast<FloatVec3>(dim) };
}
template <typename PointsVecType>
VTKM_EXEC inline Bounds ComputeCellBounds(const PointsVecType& points)
{
using CoordsType = typename vtkm::VecTraits<PointsVecType>::ComponentType;
auto numPoints = vtkm::VecTraits<PointsVecType>::GetNumberOfComponents(points);
CoordsType minp = points[0], maxp = points[0];
for (vtkm::IdComponent i = 1; i < numPoints; ++i)
{
minp = vtkm::Min(minp, points[i]);
maxp = vtkm::Max(maxp, points[i]);
}
return { FloatVec3(minp), FloatVec3(maxp) };
}
}
}
} // vtkm::internal::cl_uniform_bins
namespace vtkm
{
namespace exec
{
//--------------------------------------------------------------------
template <typename CellSetType, typename DeviceAdapter>
class VTKM_ALWAYS_EXPORT CellLocatorTwoLevel : public vtkm::exec::CellLocator
{
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>::template ExecutionTypes<DeviceAdapter>::PortalConst;
using CoordsPortalType =
typename vtkm::cont::CoordinateSystem::MultiplexerArrayType::ExecutionTypes<
DeviceAdapter>::PortalConst;
using CellSetP2CExecType =
decltype(std::declval<CellSetType>().PrepareForInput(DeviceAdapter{},
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.
template <typename CellShapeTag, typename CoordsType>
VTKM_EXEC static vtkm::ErrorCode PointInsideCell(FloatVec3 point,
CellShapeTag cellShape,
CoordsType cellPoints,
FloatVec3& parametricCoordinates,
bool& inside)
{
auto bounds = vtkm::internal::cl_uniform_bins::ComputeCellBounds(cellPoints);
if (point[0] >= bounds.Min[0] && point[0] <= bounds.Max[0] && point[1] >= bounds.Min[1] &&
point[1] <= bounds.Max[1] && point[2] >= bounds.Min[2] && point[2] <= bounds.Max[2])
{
VTKM_RETURN_ON_ERROR(vtkm::exec::WorldCoordinatesToParametricCoordinates(
cellPoints, point, cellShape, parametricCoordinates));
inside = vtkm::exec::CellInside(parametricCoordinates, cellShape);
}
else
{
inside = false;
}
// Return success error code even point is not inside this cell
return vtkm::ErrorCode::Success;
}
public:
VTKM_CONT CellLocatorTwoLevel(const vtkm::internal::cl_uniform_bins::Grid& topLevelGrid,
const vtkm::cont::ArrayHandle<DimVec3>& leafDimensions,
const vtkm::cont::ArrayHandle<vtkm::Id>& leafStartIndex,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellStartIndex,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellCount,
const vtkm::cont::ArrayHandle<vtkm::Id>& cellIds,
const CellSetType& cellSet,
const vtkm::cont::CoordinateSystem& coords,
vtkm::cont::Token& token)
: TopLevel(topLevelGrid)
, LeafDimensions(leafDimensions.PrepareForInput(DeviceAdapter{}, token))
, LeafStartIndex(leafStartIndex.PrepareForInput(DeviceAdapter{}, token))
, CellStartIndex(cellStartIndex.PrepareForInput(DeviceAdapter{}, token))
, CellCount(cellCount.PrepareForInput(DeviceAdapter{}, token))
, CellIds(cellIds.PrepareForInput(DeviceAdapter{}, token))
, CellSet(cellSet.PrepareForInput(DeviceAdapter{},
vtkm::TopologyElementTagCell{},
vtkm::TopologyElementTagPoint{},
token))
, Coords(coords.GetDataAsMultiplexer().PrepareForInput(DeviceAdapter{}, token))
{
}
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
{
using namespace vtkm::internal::cl_uniform_bins;
cellId = -1;
DimVec3 binId3 = static_cast<DimVec3>((point - this->TopLevel.Origin) / this->TopLevel.BinSize);
if (binId3[0] >= 0 && binId3[0] < this->TopLevel.Dimensions[0] && binId3[1] >= 0 &&
binId3[1] < this->TopLevel.Dimensions[1] && binId3[2] >= 0 &&
binId3[2] < this->TopLevel.Dimensions[2])
{
vtkm::Id binId = ComputeFlatIndex(binId3, this->TopLevel.Dimensions);
auto ldim = this->LeafDimensions.Get(binId);
if (!ldim[0] || !ldim[1] || !ldim[2])
{
return vtkm::ErrorCode::CellNotFound;
}
auto leafGrid = ComputeLeafGrid(binId3, ldim, this->TopLevel);
DimVec3 leafId3 = static_cast<DimVec3>((point - leafGrid.Origin) / leafGrid.BinSize);
// precision issues may cause leafId3 to be out of range so clamp it
leafId3 = vtkm::Max(DimVec3(0), vtkm::Min(ldim - DimVec3(1), leafId3));
vtkm::Id leafStart = this->LeafStartIndex.Get(binId);
vtkm::Id leafId = leafStart + ComputeFlatIndex(leafId3, leafGrid.Dimensions);
vtkm::Id start = this->CellStartIndex.Get(leafId);
vtkm::Id end = start + this->CellCount.Get(leafId);
for (vtkm::Id i = start; i < end; ++i)
{
vtkm::Id cid = this->CellIds.Get(i);
auto indices = this->CellSet.GetIndices(cid);
auto pts = vtkm::make_VecFromPortalPermute(&indices, this->Coords);
FloatVec3 pc;
bool inside;
VTKM_RETURN_ON_ERROR(
PointInsideCell(point, this->CellSet.GetCellShape(cid), pts, pc, inside));
if (inside)
{
cellId = cid;
parametric = pc;
return vtkm::ErrorCode::Success;
}
}
}
return vtkm::ErrorCode::CellNotFound;
}
private:
vtkm::internal::cl_uniform_bins::Grid TopLevel;
ArrayPortalConst<DimVec3> LeafDimensions;
ArrayPortalConst<vtkm::Id> LeafStartIndex;
ArrayPortalConst<vtkm::Id> CellStartIndex;
ArrayPortalConst<vtkm::Id> CellCount;
ArrayPortalConst<vtkm::Id> CellIds;
CellSetP2CExecType CellSet;
CoordsPortalType Coords;
};
}
} // vtkm::exec
#endif //vtk_m_exec_CellLocatorTwoLevel_h

2
vtkm/worklet/particleadvection/GridEvaluatorStatus.h
View File

@ -17,7 +17,7 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorUniformBins.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>

14
vtkm/worklet/particleadvection/GridEvaluators.h
View File

@ -17,7 +17,7 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/CellLocator.h>
#include <vtkm/cont/CellLocatorRectilinearGrid.h>
#include <vtkm/cont/CellLocatorUniformBins.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>
@ -167,11 +167,11 @@ public:
else
{
// Default to using an locator for explicit meshes.
vtkm::cont::CellLocatorUniformBins locator;
vtkm::cont::CellLocatorTwoLevel locator;
locator.SetCoordinates(coordinates);
locator.SetCellSet(cellset);
locator.Update();
this->Locator = std::make_shared<vtkm::cont::CellLocatorUniformBins>(locator);
this->Locator = std::make_shared<vtkm::cont::CellLocatorTwoLevel>(locator);
}
vtkm::cont::StructuredCellInterpolationHelper interpolationHelper(cellset);
this->InterpolationHelper =
@ -179,22 +179,22 @@ public:
}
else if (cellset.IsSameType(vtkm::cont::CellSetSingleType<>()))
{
vtkm::cont::CellLocatorUniformBins locator;
vtkm::cont::CellLocatorTwoLevel locator;
locator.SetCoordinates(coordinates);
locator.SetCellSet(cellset);
locator.Update();
this->Locator = std::make_shared<vtkm::cont::CellLocatorUniformBins>(locator);
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::CellLocatorUniformBins locator;
vtkm::cont::CellLocatorTwoLevel locator;
locator.SetCoordinates(coordinates);
locator.SetCellSet(cellset);
locator.Update();
this->Locator = std::make_shared<vtkm::cont::CellLocatorUniformBins>(locator);
this->Locator = std::make_shared<vtkm::cont::CellLocatorTwoLevel>(locator);
vtkm::cont::ExplicitCellInterpolationHelper interpolationHelper(cellset);
this->InterpolationHelper =
std::make_shared<vtkm::cont::ExplicitCellInterpolationHelper>(interpolationHelper);