Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-07-11 09:33:37 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into vtk-m-cmake_refactor

Browse Source
This commit is contained in:
Robert Maynard 2018-02-23 14:52:00 -05:00
commit e630ac5aa4
416 changed files with 8533 additions and 3402 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
CMake/FindTBB.cmake
View File

@ -174,7 +174,7 @@ endif ()
# will never adequately match the user's setup, so there is no feasible way
# to detect the "best" version to use. The user will have to manually
# select the right files. (Chances are the distributions are shipping their
# custom version of tbb, anyway, so the problem is probably nonexistant.)
# custom version of tbb, anyway, so the problem is probably nonexistent.)
if (WIN32 AND MSVC)
set(COMPILER_PREFIX "vc7.1")
if (MSVC_VERSION EQUAL 1400)

2
CMake/VTKmCheckPyexpander.cmake
View File

@ -70,7 +70,7 @@ if(${diff_result})
# If diff returned non-zero, it failed and the two files are different.
get_filename_component(filename ${SOURCE_FILE} NAME)
message(SEND_ERROR
"The source file ${filename} does not match the generated file. If you have modified this file directly, then you have messed up. Modify the ${filename}.in file instead and then copy the pyexpander result to ${filename}. If you modified ${filename}.in, then you might just need to copy the pyresult back to the source directory. If you have not modifed either, then you have likely checked out an inappropriate change. Check the git logs to see what changes were made.
"The source file ${filename} does not match the generated file. If you have modified this file directly, then you have messed up. Modify the ${filename}.in file instead and then copy the pyexpander result to ${filename}. If you modified ${filename}.in, then you might just need to copy the pyresult back to the source directory. If you have not modified either, then you have likely checked out an inappropriate change. Check the git logs to see what changes were made.
If the changes have resulted from modifying ${filename}.in, then you can finish by moving ${GENERATED_FILE}.save over ${SOURCE_FILE}")
else()
# Now that we have done the comparison, remove the generated file so there is

1
CMake/VTKmCheckSourceInBuild.cmake
View File

@ -37,6 +37,7 @@ set(FILES_TO_CHECK
)
set(EXCEPTIONS
kxsort.h
)
set(DIRECTORY_EXCEPTIONS

2
CMake/VTKmConfig.cmake.in Executable file → Normal file
View File

@ -53,7 +53,7 @@ set(VTKm_ENABLE_EGL_CONTEXT "@VTKm_ENABLE_EGL_CONTEXT@")
set(VTKm_ENABLE_MPI "@VTKm_ENABLE_MPI@")
# This is true when the package is still in the build directory (not installed)
if(CMAKE_CURRENT_LIST_DIR STREQUAL "@VTKm_BINARY_DIR@/@VTKm_INSTALL_CONFIG_DIR@")
if(CMAKE_CURRENT_LIST_DIR STREQUAL "@VTKm_BUILD_CMAKE_BASE_DIR@/@VTKm_INSTALL_CONFIG_DIR@")
set(VTKm_PACKAGE_IN_BUILD TRUE)
endif()

2
CMake/doxyfile.in
View File

@ -157,6 +157,8 @@ EXCLUDE_SYMLINKS = NO
EXCLUDE_PATTERNS = */testing/*
EXCLUDE_PATTERNS += */examples/*
EXCLUDE_PATTERNS += UnitTest*
EXCLUDE_PATTERNS += */vtkmdiy/fmt/*
EXCLUDE_PATTERNS += */vtkmdiy/thread/*
EXCLUDE_SYMBOLS = thrust
EXCLUDE_SYMBOLS += detail

59
CMakeLists.txt
View File

@ -55,9 +55,15 @@ endif()
if (NOT DEFINED VTKm_INSTALL_BIN_DIR)
set(VTKm_INSTALL_BIN_DIR "bin")
endif()
if (NOT DEFINED VTKm_INSTALL_CMAKE_MODULE_DIR)
set(VTKm_INSTALL_CMAKE_MODULE_DIR "share/vtkm-${VTKm_VERSION_MAJOR}.${VTKm_VERSION_MINOR}/cmake")
if (NOT DEFINED VTKm_INSTALL_SHARE_DIR)
set(VTKm_INSTALL_SHARE_DIR "share/vtkm-${VTKm_VERSION_MAJOR}.${VTKm_VERSION_MINOR}")
endif()
if (NOT DEFINED VTKm_INSTALL_CMAKE_MODULE_DIR)
set(VTKm_INSTALL_CMAKE_MODULE_DIR "${VTKm_INSTALL_SHARE_DIR}/cmake")
endif()
if (NOT DEFINED VTKm_BUILD_CMAKE_BASE_DIR)
set(VTKm_BUILD_CMAKE_BASE_DIR "${VTKm_BINARY_DIR}")
endif ()
set(VTKm_BINARY_INCLUDE_DIR "${VTKm_BINARY_DIR}/include")
@ -65,9 +71,6 @@ if (NOT DEFINED VTKm_EXPORT_NAME)
set(VTKm_EXPORT_NAME "VTKmTargets")
endif()
set(VTKm_EXPORT_NAME "VTKmTargets")
#-----------------------------------------------------------------------------
# vtkm_option(variable doc [initial])
# Provides an option if it is not already defined.
@ -108,14 +111,19 @@ vtkm_option(VTKm_INSTALL_ONLY_LIBRARIES "install only vtk-m libraries and no hea
# VTK does.
vtkm_option(VTKm_USE_DEFAULT_SYMBOL_VISIBILITY "Don't explicitly hide symbols from libraries." OFF)
vtkm_option(BUILD_SHARED_LIBS "Build VTK-m with shared libraries" ON)
set(VTKm_BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS})
mark_as_advanced(
VTKm_NO_ASSERT
VTKm_INSTALL_ONLY_LIBRARIES
VTKm_USE_DEFAULT_SYMBOL_VISIBILITY
)
vtkm_option(BUILD_SHARED_LIBS "Build VTK-m with shared libraries" ON)
set(VTKm_BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS})
#-----------------------------------------------------------------------------
# When using C++11 support make sure you use the standard C++ extensions rather
# than compiler-specific versions of the extensions (to preserve portability).
set(CMAKE_CXX_EXTENSIONS False)
# Setup default build types
include(VTKmBuildType)
@ -215,7 +223,7 @@ include(CMakePackageConfigHelpers)
configure_package_config_file(
${VTKm_SOURCE_DIR}/CMake/VTKmConfig.cmake.in
${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfig.cmake
${VTKm_BUILD_CMAKE_BASE_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfig.cmake
INSTALL_DESTINATION ${VTKm_INSTALL_CONFIG_DIR}
PATH_VARS
VTKm_INSTALL_INCLUDE_DIR
@ -226,25 +234,25 @@ configure_package_config_file(
)
write_basic_package_version_file(
${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfigVersion.cmake
${VTKm_BUILD_CMAKE_BASE_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfigVersion.cmake
VERSION ${VTKm_VERSION}
COMPATIBILITY ExactVersion )
if(NOT VTKm_INSTALL_ONLY_LIBRARIES)
install(
FILES
${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfig.cmake
${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfigVersion.cmake
${VTKm_BUILD_CMAKE_BASE_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfig.cmake
${VTKm_BUILD_CMAKE_BASE_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmConfigVersion.cmake
DESTINATION ${VTKm_INSTALL_CONFIG_DIR}
)
# Install the readme and license files.
install(FILES ${VTKm_SOURCE_DIR}/README.md
DESTINATION ${VTKm_INSTALL_CONFIG_DIR}
DESTINATION ${VTKm_INSTALL_SHARE_DIR}
RENAME VTKmREADME.md
)
install(FILES ${VTKm_SOURCE_DIR}/LICENSE.txt
DESTINATION ${VTKm_INSTALL_CONFIG_DIR}
DESTINATION ${VTKm_INSTALL_SHARE_DIR}
RENAME VTKmLICENSE.txt
)
@ -277,17 +285,20 @@ if(NOT VTKm_INSTALL_ONLY_LIBRARIES)
)
endif()
# Enable CPack packaging
set(CPACK_PACKAGE_DESCRIPTION_FILE ${VTKm_SOURCE_DIR}/README.md)
set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "The VTKm Toolkit")
set(CPACK_PACKAGE_NAME "VTKm")
set(CPACK_PACKAGE_VERSION_MAJOR ${VTKm_VERSION_MAJOR})
set(CPACK_PACKAGE_VERSION_MINOR ${VTKm_VERSION_MINOR})
set(CPACK_PACKAGE_VERSION_PATCH ${VTKm_VERSION_PATCH})
set(CPACK_PACKAGE_FILE_NAME "VTKm-${VTKm_VERSION}")
set(CPACK_RESOURCE_FILE_LICENSE ${VTKm_SOURCE_DIR}/LICENSE.txt)
set(CPACK_RESOURCE_FILE_README ${VTKm_SOURCE_DIR}/README.md)
include(CPack)
vtkm_option(VTKm_ENABLE_CPACK "Enable CPack packaging of VTKm" ON)
if (VTKm_ENABLE_CPACK)
# Enable CPack packaging
set(CPACK_PACKAGE_DESCRIPTION_FILE ${VTKm_SOURCE_DIR}/README.md)
set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "The VTKm Toolkit")
set(CPACK_PACKAGE_NAME "VTKm")
set(CPACK_PACKAGE_VERSION_MAJOR ${VTKm_VERSION_MAJOR})
set(CPACK_PACKAGE_VERSION_MINOR ${VTKm_VERSION_MINOR})
set(CPACK_PACKAGE_VERSION_PATCH ${VTKm_VERSION_PATCH})
set(CPACK_PACKAGE_FILE_NAME "VTKm-${VTKm_VERSION}")
set(CPACK_RESOURCE_FILE_LICENSE ${VTKm_SOURCE_DIR}/LICENSE.txt)
set(CPACK_RESOURCE_FILE_README ${VTKm_SOURCE_DIR}/README.md)
include(CPack)
endif ()
#-----------------------------------------------------------------------------
#add the benchmarking folder

2
CONTRIBUTING.md
View File

@ -436,7 +436,7 @@ authorized developers may add a comment with a single *trailing* line:
Do: merge
to ask that the change be merged into the upstream repository. By
convention, only merge if you have recieved `+1` . Do not request a merge
convention, only merge if you have received `+1` . Do not request a merge
if any `-1` review comments have not been resolved.
### Merge Success ###

6
CTestCustom.cmake.in
View File

@ -20,7 +20,7 @@
list(APPEND CTEST_CUSTOM_WARNING_EXCEPTION
".*warning: ignoring loop annotation.*"
".*diy.include.diy.*WShadow.*" # exclude `diy` shadow warnings.
".*diy.include.diy.*note: shadowed.*" # exclude `diy` shadow warnings.
".*diy.include.diy.storage.hpp.*Wunused-result.*" # this is a TODO in DIY.
".*warning: Included by graph for.*not generated, too many nodes. Consider increasing DOT_GRAPH_MAX_NODES."
".*diy.include.*diy.mpi.io.hpp.*Wconversion.*"
".*diy.include.*diy.mpi.datatypes.hpp.*Wunused-parameter.*"
)

2
LICENSE.txt
View File

@ -74,6 +74,8 @@ CMake/FindTBB.cmake
CMake/FindGLEW.cmake
Utilities
vtkm/cont/tbb/internal/parallel_sort.h
vtkm/cont/tbb/internal/parallel_radix_sort_tbb.h
vtkm/cont/tbb/internal/kxsort.h
vtkm/testing/OptionParser.h
vtkm/internal/brigand.hpp
version.txt

48
benchmarking/BenchmarkDeviceAdapter.cxx
View File

@ -169,13 +169,13 @@ const static std::string DIVIDER(40, '-');
template <class DeviceAdapterTag>
class BenchmarkDeviceAdapter
{
typedef vtkm::cont::StorageTagBasic StorageTag;
using StorageTag = vtkm::cont::StorageTagBasic;
typedef vtkm::cont::ArrayHandle<vtkm::Id, StorageTag> IdArrayHandle;
using IdArrayHandle = vtkm::cont::ArrayHandle<vtkm::Id, StorageTag>;
typedef vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag> Algorithm;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag>;
typedef vtkm::cont::Timer<DeviceAdapterTag> Timer;
using Timer = vtkm::cont::Timer<DeviceAdapterTag>;
public:
// Various kernels used by the different benchmarks to accelerate
@ -183,8 +183,8 @@ public:
template <typename Value>
struct FillTestValueKernel : vtkm::exec::FunctorBase
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
typedef typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal PortalType;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
using PortalType = typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal;
PortalType Output;
@ -200,8 +200,8 @@ public:
template <typename Value>
struct FillScaledTestValueKernel : vtkm::exec::FunctorBase
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
typedef typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal PortalType;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
using PortalType = typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal;
PortalType Output;
const vtkm::Id IdScale;
@ -219,8 +219,8 @@ public:
template <typename Value>
struct FillModuloTestValueKernel : vtkm::exec::FunctorBase
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
typedef typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal PortalType;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
using PortalType = typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal;
PortalType Output;
const vtkm::Id Modulus;
@ -238,8 +238,8 @@ public:
template <typename Value>
struct FillBinaryTestValueKernel : vtkm::exec::FunctorBase
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
typedef typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal PortalType;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
using PortalType = typename ValueArrayHandle::template ExecutionTypes<DeviceAdapterTag>::Portal;
PortalType Output;
const vtkm::Id Modulus;
@ -261,7 +261,7 @@ private:
template <typename Value>
struct BenchCopy
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
ValueArrayHandle ValueHandle_src;
ValueArrayHandle ValueHandle_dst;
@ -302,7 +302,7 @@ private:
template <typename Value>
struct BenchCopyIf
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
const vtkm::Id PERCENT_VALID;
const vtkm::Id N_VALID;
@ -359,7 +359,7 @@ private:
template <typename Value>
struct BenchLowerBounds
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
const vtkm::Id N_VALS;
const vtkm::Id PERCENT_VALUES;
@ -415,10 +415,10 @@ private:
template <typename Value>
struct BenchReduce
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
ValueArrayHandle InputHandle;
// We don't actually use this, but we need it to prevent sufficently
// We don't actually use this, but we need it to prevent sufficiently
// smart compilers from optimizing the Reduce call out.
Value Result;
@ -461,7 +461,7 @@ private:
template <typename Value>
struct BenchReduceByKey
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
const vtkm::Id N_KEYS;
const vtkm::Id PERCENT_KEYS;
@ -518,7 +518,7 @@ private:
template <typename Value>
struct BenchScanInclusive
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
ValueArrayHandle ValueHandle, OutHandle;
VTKM_CONT
@ -553,7 +553,7 @@ private:
template <typename Value>
struct BenchScanExclusive
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
ValueArrayHandle ValueHandle, OutHandle;
@ -589,7 +589,7 @@ private:
template <typename Value>
struct BenchSort
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
ValueArrayHandle ValueHandle;
std::mt19937 Rng;
@ -631,7 +631,7 @@ private:
template <typename Value>
struct BenchSortByKey
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
std::mt19937 Rng;
vtkm::Id N_KEYS;
@ -797,7 +797,7 @@ private:
template <typename Value>
struct BenchUnique
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
const vtkm::Id N_VALID;
const vtkm::Id PERCENT_VALID;
@ -853,7 +853,7 @@ private:
template <typename Value>
struct BenchUpperBounds
{
typedef vtkm::cont::ArrayHandle<Value, StorageTag> ValueArrayHandle;
using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;
const vtkm::Id N_VALS;
const vtkm::Id PERCENT_VALS;

10
benchmarking/BenchmarkFieldAlgorithms.cxx
View File

@ -205,7 +205,7 @@ class GenerateEdges : public vtkm::worklet::WorkletMapPointToCell
public:
typedef void ControlSignature(CellSetIn cellset, WholeArrayOut<> edgeIds);
typedef void ExecutionSignature(PointIndices, ThreadIndices, _2);
typedef _1 InputDomain;
using InputDomain = _1;
template <typename ConnectivityInVec, typename ThreadIndicesType, typename IdPairTableType>
VTKM_EXEC void operator()(const ConnectivityInVec& connectivity,
@ -234,7 +234,7 @@ public:
WholeArrayIn<> inputField,
FieldOut<> output);
typedef void ExecutionSignature(_1, _2, _3, _4);
typedef _1 InputDomain;
using InputDomain = _1;
template <typename WeightType, typename T, typename S, typename D>
VTKM_EXEC void operator()(const vtkm::Id2& low_high,
@ -321,11 +321,11 @@ using StorageListTag = ::vtkm::cont::StorageListTagBasic;
template <class DeviceAdapterTag>
class BenchmarkFieldAlgorithms
{
typedef vtkm::cont::StorageTagBasic StorageTag;
using StorageTag = vtkm::cont::StorageTagBasic;
typedef vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag> Algorithm;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag>;
typedef vtkm::cont::Timer<DeviceAdapterTag> Timer;
using Timer = vtkm::cont::Timer<DeviceAdapterTag>;
using ValueDynamicHandle = vtkm::cont::DynamicArrayHandleBase<ValueTypes, StorageListTag>;
using InterpDynamicHandle = vtkm::cont::DynamicArrayHandleBase<InterpValueTypes, StorageListTag>;

14
benchmarking/BenchmarkTopologyAlgorithms.cxx
View File

@ -59,7 +59,7 @@ public:
CellSetIn cellset,
FieldOutCell<> outCells);
typedef void ExecutionSignature(_1, PointCount, _3);
typedef _2 InputDomain;
using InputDomain = _2;
template <typename PointValueVecType, typename OutType>
VTKM_EXEC void operator()(const PointValueVecType& pointValues,
@ -81,7 +81,7 @@ class AverageCellToPoint : public vtkm::worklet::WorkletMapCellToPoint
public:
typedef void ControlSignature(FieldInCell<> inCells, CellSetIn topology, FieldOut<> outPoints);
typedef void ExecutionSignature(_1, _3, CellCount);
typedef _2 InputDomain;
using InputDomain = _2;
template <typename CellVecType, typename OutType>
VTKM_EXEC void operator()(const CellVecType& cellValues,
@ -110,7 +110,7 @@ public:
CellSetIn cellset,
FieldOutCell<IdComponentType> outCaseId);
typedef void ExecutionSignature(_1, _3);
typedef _2 InputDomain;
using InputDomain = _2;
T IsoValue;
@ -123,7 +123,7 @@ public:
template <typename FieldInType>
VTKM_EXEC void operator()(const FieldInType& fieldIn, vtkm::IdComponent& caseNumber) const
{
typedef typename vtkm::VecTraits<FieldInType>::ComponentType FieldType;
using FieldType = typename vtkm::VecTraits<FieldInType>::ComponentType;
const FieldType iso = static_cast<FieldType>(this->IsoValue);
caseNumber = ((fieldIn[0] > iso) | (fieldIn[1] > iso) << 1 | (fieldIn[2] > iso) << 2 |
@ -143,11 +143,11 @@ using StorageListTag = ::vtkm::cont::StorageListTagBasic;
template <class DeviceAdapterTag>
class BenchmarkTopologyAlgorithms
{
typedef vtkm::cont::StorageTagBasic StorageTag;
using StorageTag = vtkm::cont::StorageTagBasic;
typedef vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag> Algorithm;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag>;
typedef vtkm::cont::Timer<DeviceAdapterTag> Timer;
using Timer = vtkm::cont::Timer<DeviceAdapterTag>;
using ValueDynamicHandle = vtkm::cont::DynamicArrayHandleBase<ValueTypes, StorageListTag>;

2
examples/isosurface/IsosurfaceUniformGrid.cxx
View File

@ -245,7 +245,7 @@ int main(int argc, char* argv[])
vtkm::filter::MarchingCubes filter;
filter.SetGenerateNormals(true);
filter.SetMergeDuplicatePoints(false);
filter.SetIsoValue(0, 0.5);
filter.SetIsoValue(0, 0.1);
vtkm::filter::Result result = filter.Execute(dataSet, dataSet.GetField("nodevar"));
filter.MapFieldOntoOutput(result, dataSet.GetField("nodevar"));

16
examples/tau_timing/TauTiming.cxx
View File

@ -64,14 +64,14 @@ VTKM_EXEC_CONT vtkm::Vec<T, 3> Normalize(vtkm::Vec<T, 3> v)
return one / magnitude * v;
}
typedef VTKM_DEFAULT_DEVICE_ADAPTER_TAG DeviceAdapter;
using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
class TangleField : public vtkm::worklet::WorkletMapField
{
public:
typedef void ControlSignature(FieldIn<IdType> vertexId, FieldOut<Scalar> v);
typedef void ExecutionSignature(_1, _2);
typedef _1 InputDomain;
using InputDomain = _1;
const vtkm::Id xdim, ydim, zdim;
const vtkm::FloatDefault xmin, ymin, zmin, xmax, ymax, zmax;
@ -343,9 +343,9 @@ void RenderRTTest(const vtkm::cont::DataSet& ds, int N)
if (printProgress && i % 10 == 0)
std::cout << " " << i << " of " << N << std::endl;
typedef vtkm::rendering::MapperRayTracer M;
typedef vtkm::rendering::CanvasRayTracer C;
typedef vtkm::rendering::View3D V3;
using M = vtkm::rendering::MapperRayTracer;
using C = vtkm::rendering::CanvasRayTracer;
using V3 = vtkm::rendering::View3D;
//std::cout<<"Render: "<<i<<std::endl;
vtkm::rendering::ColorTable colorTable("thermal");
@ -361,9 +361,9 @@ void RenderVolTest(const vtkm::cont::DataSet& ds, int N)
if (printProgress && i % 10 == 0)
std::cout << " " << i << " of " << N << std::endl;
typedef vtkm::rendering::MapperVolume M;
typedef vtkm::rendering::CanvasRayTracer C;
typedef vtkm::rendering::View3D V3;
using M = vtkm::rendering::MapperVolume;
using C = vtkm::rendering::CanvasRayTracer;
using V3 = vtkm::rendering::View3D;
//std::cout<<"Render: "<<i<<std::endl;
vtkm::rendering::ColorTable colorTable("thermal");

2
vtkm/BinaryPredicates.h
View File

@ -102,4 +102,4 @@ struct LogicalOr
} // namespace vtkm
#endif //vtk_m_BinaryPredicates_h
#endif //vtk_m_BinaryPredicates_h

2
vtkm/Bounds.h
View File

@ -140,7 +140,7 @@ struct Bounds
/// \b Expand bounds to include other bounds.
///
/// This version of \c Include expands these bounds just enough to include
/// that of another bounds. Esentially it is the union of the two bounds.
/// that of another bounds. Essentially it is the union of the two bounds.
///
VTKM_EXEC_CONT
void Include(const vtkm::Bounds& bounds)

18
vtkm/CellTraits.h
View File

@ -66,7 +66,7 @@ struct CellTraits
/// a convenient way to overload a function based on topological dimensions
/// (which is usually more efficient than conditionals).
///
typedef vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS> TopologicalDimensionsTag;
using TopologicalDimensionsTag = vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS>;
/// \brief A tag specifying whether the number of points is fixed.
///
@ -74,7 +74,7 @@ struct CellTraits
/// \c CellTraitsTagSizeVariable, then the number of points is not known at
/// compile time.
///
typedef vtkm::CellTraitsTagSizeFixed IsSizeFixed;
using IsSizeFixed = vtkm::CellTraitsTagSizeFixed;
/// \brief Number of points in the cell.
///
@ -95,19 +95,19 @@ struct CellTraits
template <> \
struct CellTraits<vtkm::CellShapeTag##name> \
{ \
const static vtkm::IdComponent TOPOLOGICAL_DIMENSIONS = dimensions; \
typedef vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS> TopologicalDimensionsTag; \
typedef vtkm::CellTraitsTagSizeFixed IsSizeFixed; \
static const vtkm::IdComponent NUM_POINTS = numPoints; \
static constexpr vtkm::IdComponent TOPOLOGICAL_DIMENSIONS = dimensions; \
using TopologicalDimensionsTag = vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS>; \
using IsSizeFixed = vtkm::CellTraitsTagSizeFixed; \
static constexpr vtkm::IdComponent NUM_POINTS = numPoints; \
}
#define VTKM_DEFINE_CELL_TRAITS_VARIABLE(name, dimensions) \
template <> \
struct CellTraits<vtkm::CellShapeTag##name> \
{ \
const static vtkm::IdComponent TOPOLOGICAL_DIMENSIONS = dimensions; \
typedef vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS> TopologicalDimensionsTag; \
typedef vtkm::CellTraitsTagSizeVariable IsSizeFixed; \
static constexpr vtkm::IdComponent TOPOLOGICAL_DIMENSIONS = dimensions; \
using TopologicalDimensionsTag = vtkm::CellTopologicalDimensionsTag<TOPOLOGICAL_DIMENSIONS>; \
using IsSizeFixed = vtkm::CellTraitsTagSizeVariable; \
}
VTKM_DEFINE_CELL_TRAITS(Empty, 0, 0);

2
vtkm/Matrix.h
View File

@ -329,7 +329,7 @@ VTKM_EXEC_CONT void MatrixLUPFactorFindUpperTriangleElements(vtkm::Matrix<T, Siz
/// LU-factorization takes a matrix A and decomposes it into a lower triangular
/// matrix L and upper triangular matrix U such that A = LU. The
/// LUP-factorization also allows permutation of A, which makes the
/// decomposition always posible so long as A is not singular. In addition to
/// decomposition always possible so long as A is not singular. In addition to
/// matrices L and U, LUP also finds permutation matrix P containing all zeros
/// except one 1 per row and column such that PA = LU.
///

4
vtkm/NewtonsMethod.h
View File

@ -57,8 +57,8 @@ VTKM_EXEC_CONT NewtonsMethodResult<ScalarType, Size> NewtonsMethod(
ScalarType convergeDifference = ScalarType(1e-3),
vtkm::IdComponent maxIterations = 10)
{
typedef vtkm::Vec<ScalarType, Size> VectorType;
typedef vtkm::Matrix<ScalarType, Size, Size> MatrixType;
using VectorType = vtkm::Vec<ScalarType, Size>;
using MatrixType = vtkm::Matrix<ScalarType, Size, Size>;
VectorType x = initialGuess;

9
vtkm/Range.h
View File

@ -137,13 +137,16 @@ struct Range
/// \b Expand range to include other range.
///
/// This version of \c Include expands this range just enough to include that
/// of another range. Esentially it is the union of the two ranges.
/// of another range. Essentially it is the union of the two ranges.
///
VTKM_EXEC_CONT
void Include(const vtkm::Range& range)
{
this->Include(range.Min);
this->Include(range.Max);
if (range.IsNonEmpty())
{
this->Include(range.Min);
this->Include(range.Max);
}
}
/// \b Return the union of this and another range.

2
vtkm/RangeId.h
View File

@ -102,7 +102,7 @@ struct RangeId
/// \b Expand range to include other range.
///
/// This version of \c Include expands this range just enough to include that
/// of another range. Esentially it is the union of the two ranges.
/// of another range. Essentially it is the union of the two ranges.
///
VTKM_EXEC_CONT
void Include(const vtkm::RangeId& range)

2
vtkm/RangeId3.h
View File

@ -130,7 +130,7 @@ struct RangeId3
/// \b Expand range to include other range.
///
/// This version of \c Include expands the range just enough to include
/// the other range. Esentially it is the union of the two ranges.
/// the other range. Essentially it is the union of the two ranges.
///
VTKM_EXEC_CONT
void Include(const vtkm::RangeId3& range)

2
vtkm/Transform3D.h
View File

@ -58,7 +58,7 @@ VTKM_EXEC_CONT vtkm::Vec<T, 3> Transform3DPoint(const vtkm::Matrix<T, 4, 4>& mat
/// transformed by the given matrix in homogeneous coordinates.
///
/// Unlike Transform3DPoint, this method honors the fourth component of the
/// transformed homogeneous coordiante. This makes it applicable for perspective
/// transformed homogeneous coordinate. This makes it applicable for perspective
/// transformations, but requires some more computations.
///
template <typename T>

4
vtkm/TypeTraits.h
View File

@ -63,7 +63,7 @@ struct TypeTraitsVectorTag
};
/// The TypeTraits class provides helpful compile-time information about the
/// basic types used in VTKm (and a few others for convienience). The majority
/// basic types used in VTKm (and a few others for convenience). The majority
/// of TypeTraits contents are typedefs to tags that can be used to easily
/// override behavior of called functions.
///
@ -71,7 +71,7 @@ template <typename T>
class TypeTraits
{
public:
/// \brief A tag to determing whether the type is integer or real.
/// \brief A tag to determine whether the type is integer or real.
///
/// This tag is either TypeTraitsRealTag or TypeTraitsIntegerTag.
using NumericTag = vtkm::TypeTraitsUnknownTag;

4
vtkm/Types.h
View File

@ -48,7 +48,7 @@
* \brief Transportation controls for Control Environment Objects.
*
* vtkm::cont::arg includes the classes that allows the vtkm::worklet::Dispatchers
* to request Control Environment Objects to be transfered to the Execution Environment.
* to request Control Environment Objects to be transferred to the Execution Environment.
*
* \namespace vtkm::cont::cuda
* \brief CUDA implementation for Control Environment.
@ -749,7 +749,7 @@ protected:
/// The \c Vec class is most often used to represent vectors in the
/// mathematical sense as a quantity with a magnitude and direction. Vectors
/// are, of course, used extensively in computational geometry as well as
/// phyiscal simulations. The \c Vec class can be (and is) repurposed for more
/// physical simulations. The \c Vec class can be (and is) repurposed for more
/// general usage of holding a fixed-length sequence of objects.
///
/// There is no real limit to the size of the sequence (other than the largest

2
vtkm/UnaryPredicates.h
View File

@ -63,4 +63,4 @@ struct LogicalNot
} // namespace vtkm
#endif //vtk_m_UnaryPredicates_h
#endif //vtk_m_UnaryPredicates_h

14
vtkm/VecAxisAlignedPointCoordinates.h
View File

@ -75,7 +75,7 @@ template <vtkm::IdComponent NumDimensions>
class VecAxisAlignedPointCoordinates
{
public:
typedef vtkm::Vec<vtkm::FloatDefault, 3> ComponentType;
using ComponentType = vtkm::Vec<vtkm::FloatDefault, 3>;
static const vtkm::IdComponent NUM_COMPONENTS =
detail::VecAxisAlignedPointCoordinatesNumComponents<NumDimensions>::NUM_COMPONENTS;
@ -132,8 +132,8 @@ private:
template <vtkm::IdComponent NumDimensions>
struct TypeTraits<vtkm::VecAxisAlignedPointCoordinates<NumDimensions>>
{
typedef vtkm::TypeTraitsRealTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
using NumericTag = vtkm::TypeTraitsRealTag;
using DimensionalityTag = TypeTraitsVectorTag;
VTKM_EXEC_CONT
static vtkm::VecAxisAlignedPointCoordinates<NumDimensions> ZeroInitialization()
@ -146,11 +146,11 @@ struct TypeTraits<vtkm::VecAxisAlignedPointCoordinates<NumDimensions>>
template <vtkm::IdComponent NumDimensions>
struct VecTraits<vtkm::VecAxisAlignedPointCoordinates<NumDimensions>>
{
typedef vtkm::VecAxisAlignedPointCoordinates<NumDimensions> VecType;
using VecType = vtkm::VecAxisAlignedPointCoordinates<NumDimensions>;
typedef vtkm::Vec<vtkm::FloatDefault, 3> ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeStatic IsSizeStatic;
using ComponentType = vtkm::Vec<vtkm::FloatDefault, 3>;
using HasMultipleComponents = vtkm::VecTraitsTagMultipleComponents;
using IsSizeStatic = vtkm::VecTraitsTagSizeStatic;
static const vtkm::IdComponent NUM_COMPONENTS = VecType::NUM_COMPONENTS;

14
vtkm/VecFromPortal.h
View File

@ -89,11 +89,11 @@ template <typename PortalType>
struct TypeTraits<vtkm::VecFromPortal<PortalType>>
{
private:
typedef typename PortalType::ValueType ComponentType;
using ComponentType = typename PortalType::ValueType;
public:
typedef typename vtkm::TypeTraits<ComponentType>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
using NumericTag = typename vtkm::TypeTraits<ComponentType>::NumericTag;
using DimensionalityTag = TypeTraitsVectorTag;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
@ -106,11 +106,11 @@ public:
template <typename PortalType>
struct VecTraits<vtkm::VecFromPortal<PortalType>>
{
typedef vtkm::VecFromPortal<PortalType> VecType;
using VecType = vtkm::VecFromPortal<PortalType>;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
using ComponentType = typename VecType::ComponentType;
using HasMultipleComponents = vtkm::VecTraitsTagMultipleComponents;
using IsSizeStatic = vtkm::VecTraitsTagSizeVariable;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT

24
vtkm/VecFromPortalPermute.h
View File

@ -83,12 +83,12 @@ template <typename IndexVecType, typename PortalType>
struct TypeTraits<vtkm::VecFromPortalPermute<IndexVecType, PortalType>>
{
private:
typedef vtkm::VecFromPortalPermute<IndexVecType, PortalType> VecType;
typedef typename PortalType::ValueType ComponentType;
using VecType = vtkm::VecFromPortalPermute<IndexVecType, PortalType>;
using ComponentType = typename PortalType::ValueType;
public:
typedef typename vtkm::TypeTraits<ComponentType>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
using NumericTag = typename vtkm::TypeTraits<ComponentType>::NumericTag;
using DimensionalityTag = TypeTraitsVectorTag;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
@ -98,11 +98,11 @@ public:
template <typename IndexVecType, typename PortalType>
struct VecTraits<vtkm::VecFromPortalPermute<IndexVecType, PortalType>>
{
typedef vtkm::VecFromPortalPermute<IndexVecType, PortalType> VecType;
using VecType = vtkm::VecFromPortalPermute<IndexVecType, PortalType>;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
using ComponentType = typename VecType::ComponentType;
using HasMultipleComponents = vtkm::VecTraitsTagMultipleComponents;
using IsSizeStatic = vtkm::VecTraitsTagSizeVariable;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
@ -126,6 +126,14 @@ struct VecTraits<vtkm::VecFromPortalPermute<IndexVecType, PortalType>>
}
};
template <typename IndexVecType, typename PortalType>
inline VTKM_EXEC VecFromPortalPermute<IndexVecType, PortalType> make_VecFromPortalPermute(
const IndexVecType* index,
const PortalType& portal)
{
return VecFromPortalPermute<IndexVecType, PortalType>(index, portal);
}
} // namespace vtkm
#endif //vtk_m_VecFromPortalPermute_h

2
vtkm/VecTraits.h
View File

@ -367,7 +367,7 @@ struct VecTraitsBasic
/// \brief VecTraits for Pair types
///
/// Although a pair woudl seem better as a size-2 vector, we treat it as a
/// Although a pair would seem better as a size-2 vector, we treat it as a
/// scalar. This is because a \c Vec is assumed to have the same type for
/// every component, and a pair in general has a different type for each
/// component. Thus we treat a pair as a "scalar" unit.

29
vtkm/VecVariable.h
View File

@ -40,7 +40,7 @@ template <typename T, vtkm::IdComponent MaxSize>
class VecVariable
{
public:
typedef T ComponentType;
using ComponentType = T;
VTKM_EXEC_CONT
VecVariable()
@ -48,8 +48,8 @@ public:
{
}
template <vtkm::IdComponent SrcSize>
VTKM_EXEC_CONT VecVariable(const vtkm::VecVariable<ComponentType, SrcSize>& src)
template <typename SrcVecType>
VTKM_EXEC_CONT VecVariable(const SrcVecType& src)
: NumComponents(src.GetNumberOfComponents())
{
VTKM_ASSERT(this->NumComponents <= MaxSize);
@ -59,17 +59,6 @@ public:
}
}
template <vtkm::IdComponent SrcSize>
VTKM_EXEC_CONT VecVariable(const vtkm::Vec<ComponentType, SrcSize>& src)
: NumComponents(SrcSize)
{
VTKM_ASSERT(this->NumComponents <= MaxSize);
for (vtkm::IdComponent index = 0; index < this->NumComponents; index++)
{
this->Data[index] = src[index];
}
}
VTKM_EXEC_CONT
vtkm::IdComponent GetNumberOfComponents() const { return this->NumComponents; }
@ -105,8 +94,8 @@ private:
template <typename T, vtkm::IdComponent MaxSize>
struct TypeTraits<vtkm::VecVariable<T, MaxSize>>
{
typedef typename vtkm::TypeTraits<T>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
using NumericTag = typename vtkm::TypeTraits<T>::NumericTag;
using DimensionalityTag = TypeTraitsVectorTag;
VTKM_EXEC_CONT
static vtkm::VecVariable<T, MaxSize> ZeroInitialization()
@ -118,11 +107,11 @@ struct TypeTraits<vtkm::VecVariable<T, MaxSize>>
template <typename T, vtkm::IdComponent MaxSize>
struct VecTraits<vtkm::VecVariable<T, MaxSize>>
{
typedef vtkm::VecVariable<T, MaxSize> VecType;
using VecType = vtkm::VecVariable<T, MaxSize>;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
using ComponentType = typename VecType::ComponentType;
using HasMultipleComponents = vtkm::VecTraitsTagMultipleComponents;
using IsSizeStatic = vtkm::VecTraitsTagSizeVariable;
VTKM_EXEC_CONT
static vtkm::IdComponent GetNumberOfComponents(const VecType& vector)

2
vtkm/VectorAnalysis.h
View File

@ -22,7 +22,7 @@
#ifndef vtk_m_VectorAnalysis_h
#define vtk_m_VectorAnalysis_h
// This header file defines math functions that deal with linear albegra funcitons
// This header file defines math functions that deal with linear albegra functions
#include <vtkm/Math.h>
#include <vtkm/TypeTraits.h>

39
vtkm/cont/ArrayHandle.cxx
View File

@ -21,49 +21,10 @@
#define vtkm_cont_ArrayHandle_cxx
#include <vtkm/cont/ArrayHandle.h>
#ifdef VTKM_MSVC
#define _VTKM_SHARED_PTR_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT \
std::shared_ptr<vtkm::cont::ArrayHandle<Type, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 2>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 3>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 4>, vtkm::cont::StorageTagBasic>::InternalStruct>;
_VTKM_SHARED_PTR_INSTANTIATE(char)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int8)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt8)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int16)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt16)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int64)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt64)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Float32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Float64)
#undef _VTKM_SHARED_PTR_INSTANTIATE
#endif // VTKM_MSVC
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasicBase::ExecutionArrayInterfaceBasicBase(StorageBasicBase& storage)
: ControlStorage(storage)
{
}
ExecutionArrayInterfaceBasicBase::~ExecutionArrayInterfaceBasicBase()
{
}
} // end namespace internal
#define _VTKM_ARRAYHANDLE_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT ArrayHandle<Type, StorageTagBasic>; \

35
vtkm/cont/ArrayHandle.h
View File

@ -485,12 +485,13 @@ public:
struct VTKM_ALWAYS_EXPORT InternalStruct
{
StorageType ControlArray;
bool ControlArrayValid;
mutable StorageType ControlArray;
mutable bool ControlArrayValid;
std::unique_ptr<vtkm::cont::internal::ArrayHandleExecutionManagerBase<ValueType, StorageTag>>
mutable std::unique_ptr<
vtkm::cont::internal::ArrayHandleExecutionManagerBase<ValueType, StorageTag>>
ExecutionArray;
bool ExecutionArrayValid;
mutable bool ExecutionArrayValid;
};
VTKM_CONT
@ -661,32 +662,6 @@ VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle(
#ifndef vtkm_cont_ArrayHandle_cxx
#ifdef VTKM_MSVC
#define _VTKM_SHARED_PTR_EXPORT(Type) \
extern template class VTKM_CONT_TEMPLATE_EXPORT \
std::shared_ptr<vtkm::cont::ArrayHandle<Type, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 2>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 3>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 4>, vtkm::cont::StorageTagBasic>::InternalStruct>;
_VTKM_SHARED_PTR_EXPORT(char)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int8)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt8)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int16)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt16)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int32)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt32)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int64)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt64)
_VTKM_SHARED_PTR_EXPORT(vtkm::Float32)
_VTKM_SHARED_PTR_EXPORT(vtkm::Float64)
#undef _VTKM_SHARED_PTR_EXPORT
#endif // VTKM_MSVC
namespace vtkm
{
namespace cont

19
vtkm/cont/ArrayHandle.hxx
View File

@ -299,9 +299,8 @@ void ArrayHandle<T, S>::PrepareForDevice(DeviceAdapterTag) const
this->SyncControlArray();
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
internals->ExecutionArray.reset();
internals->ExecutionArrayValid = false;
this->Internals->ExecutionArray.reset();
this->Internals->ExecutionArrayValid = false;
}
}
@ -309,10 +308,9 @@ void ArrayHandle<T, S>::PrepareForDevice(DeviceAdapterTag) const
VTKM_ASSERT(!this->Internals->ExecutionArrayValid);
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
internals->ExecutionArray.reset(
this->Internals->ExecutionArray.reset(
new vtkm::cont::internal::ArrayHandleExecutionManager<T, StorageTag, DeviceAdapterTag>(
&internals->ControlArray));
&this->Internals->ControlArray));
}
template <typename T, typename S>
@ -322,19 +320,18 @@ void ArrayHandle<T, S>::SyncControlArray() const
{
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
if (this->Internals->ExecutionArrayValid)
{
internals->ExecutionArray->RetrieveOutputData(&internals->ControlArray);
internals->ControlArrayValid = true;
this->Internals->ExecutionArray->RetrieveOutputData(&this->Internals->ControlArray);
this->Internals->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
internals->ControlArray.Allocate(0);
internals->ControlArrayValid = true;
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
}
}

2
vtkm/cont/ArrayHandleTransform.h
View File

@ -436,7 +436,7 @@ public:
{
throw vtkm::cont::ErrorInternal(
"ArrayHandleTransform read only. "
"There should be no occurance of the ArrayHandle trying to pull "
"There should be no occurrence of the ArrayHandle trying to pull "
"data from the execution environment.");
}

4
vtkm/cont/ArrayHandleVirtualCoordinates.h
View File

@ -441,7 +441,7 @@ public:
}
template <typename ArrayHandleType>
bool IsSameType() const
bool IsType() const
{
return this->GetArrayHandleWrapper<ArrayHandleType>() != nullptr;
}
@ -488,7 +488,7 @@ void CastAndCall(const vtkm::cont::ArrayHandleVirtualCoordinates& coords,
Functor&& f,
Args&&... args)
{
if (coords.IsSameType<vtkm::cont::ArrayHandleUniformPointCoordinates>())
if (coords.IsType<vtkm::cont::ArrayHandleUniformPointCoordinates>())
{
f(coords.Cast<vtkm::cont::ArrayHandleUniformPointCoordinates>(), std::forward<Args>(args)...);
}

4
vtkm/cont/ArrayPortalToIterators.h
View File

@ -75,7 +75,7 @@ private:
PortalType Portal;
};
/// Convienience function for converting an ArrayPortal to a begin iterator.
/// Convenience function for converting an ArrayPortal to a begin iterator.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename PortalType>
@ -86,7 +86,7 @@ ArrayPortalToIteratorBegin(const PortalType& portal)
return iterators.GetBegin();
}
/// Convienience function for converting an ArrayPortal to an end iterator.
/// Convenience function for converting an ArrayPortal to an end iterator.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename PortalType>

5
vtkm/cont/ArrayRangeCompute.cxx
View File

@ -25,6 +25,11 @@ namespace vtkm
namespace cont
{
void ThrowArrayRangeComputeFailed()
{
throw vtkm::cont::ErrorExecution("Failed to run ArrayRangeComputation on any device.");
}
#define VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(T, Storage) \
VTKM_CONT \
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( \

2
vtkm/cont/ArrayRangeCompute.h
View File

@ -159,6 +159,8 @@ VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
return result;
}
VTKM_CONT_EXPORT void ThrowArrayRangeComputeFailed();
}
} // namespace vtkm::cont

105
vtkm/cont/ArrayRangeCompute.hxx
View File

@ -25,8 +25,9 @@
#include <vtkm/BinaryOperators.h>
#include <vtkm/VecTraits.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/cont/Algorithm.h>
#include <limits>
namespace vtkm
{
@ -36,73 +37,68 @@ namespace cont
namespace detail
{
template <typename ArrayHandleType>
struct ArrayRangeComputeFunctor
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
ArrayHandleType InputArray;
vtkm::cont::ArrayHandle<vtkm::Range> RangeArray;
VTKM_CONT
ArrayRangeComputeFunctor(const ArrayHandleType& input)
: InputArray(input)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device)
template <typename Device, typename T, typename S>
VTKM_CONT bool operator()(Device,
const vtkm::cont::ArrayHandle<T, S>& handle,
const vtkm::Vec<T, 2>& initialValue,
vtkm::Vec<T, 2>& result) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
using ValueType = typename ArrayHandleType::ValueType;
using VecTraits = vtkm::VecTraits<ValueType>;
const vtkm::IdComponent NumberOfComponents = VecTraits::NUM_COMPONENTS;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<Device>;
this->RangeArray.Allocate(NumberOfComponents);
if (this->InputArray.GetNumberOfValues() < 1)
{
for (vtkm::IdComponent i = 0; i < NumberOfComponents; ++i)
{
this->RangeArray.GetPortalControl().Set(i, vtkm::Range());
}
return true;
}
vtkm::Vec<ValueType, 2> initial(this->InputArray.GetPortalConstControl().Get(0));
vtkm::Vec<ValueType, 2> result =
Algorithm::Reduce(this->InputArray, initial, vtkm::MinAndMax<ValueType>());
for (vtkm::IdComponent i = 0; i < NumberOfComponents; ++i)
{
this->RangeArray.GetPortalControl().Set(
i,
vtkm::Range(VecTraits::GetComponent(result[0], i), VecTraits::GetComponent(result[1], i)));
}
result = Algorithm::Reduce(handle, initialValue, vtkm::MinAndMax<T>());
return true;
}
};
template <typename ArrayHandleType>
template <typename T, typename S>
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
const ArrayHandleType& input,
vtkm::cont::RuntimeDeviceTracker tracker)
const vtkm::cont::ArrayHandle<T, S>& input,
vtkm::cont::RuntimeDeviceTracker& tracker)
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
using VecTraits = vtkm::VecTraits<T>;
using CT = typename VecTraits::ComponentType;
//We want to minimize the amount of code that we do in try execute as
//it is repeated for each
vtkm::cont::ArrayHandle<vtkm::Range> range;
range.Allocate(VecTraits::NUM_COMPONENTS);
detail::ArrayRangeComputeFunctor<ArrayHandleType> functor(input);
if (!vtkm::cont::TryExecute(functor, tracker))
if (input.GetNumberOfValues() < 1)
{
throw vtkm::cont::ErrorExecution("Failed to run ArrayRangeComputation on any device.");
auto portal = range.GetPortalControl();
for (vtkm::IdComponent i = 0; i < VecTraits::NUM_COMPONENTS; ++i)
{
portal.Set(i, vtkm::Range());
}
}
else
{
//We used the limits, so that we don't need to sync the array handle
//
vtkm::Vec<T, 2> result;
vtkm::Vec<T, 2> initial;
initial[0] = T(std::numeric_limits<CT>::max());
initial[1] = T(std::numeric_limits<CT>::lowest());
return functor.RangeArray;
const bool rangeComputed =
vtkm::cont::TryExecute(detail::ArrayRangeComputeFunctor{}, tracker, input, initial, result);
if (!rangeComputed)
{
ThrowArrayRangeComputeFailed();
}
else
{
auto portal = range.GetPortalControl();
for (vtkm::IdComponent i = 0; i < VecTraits::NUM_COMPONENTS; ++i)
{
portal.Set(i,
vtkm::Range(VecTraits::GetComponent(result[0], i),
VecTraits::GetComponent(result[1], i)));
}
}
}
return range;
}
} // namespace detail
@ -113,7 +109,6 @@ inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
vtkm::cont::RuntimeDeviceTracker tracker)
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
return detail::ArrayRangeComputeImpl(input, tracker);
}
}

1
vtkm/cont/CMakeLists.txt
View File

@ -107,6 +107,7 @@ set(sources
Field.cxx
internal/SimplePolymorphicContainer.cxx
internal/ArrayManagerExecutionShareWithControl.cxx
internal/ArrayHandleBasicImpl.cxx
StorageBasic.cxx
)

4
vtkm/cont/CellLocator.h
View File

@ -38,7 +38,7 @@ private:
VTKM_CONT static bool IsUniformGrid(const vtkm::cont::DynamicCellSet& cellset,
const vtkm::cont::CoordinateSystem& coordinates)
{
return coordinates.GetData().IsSameType<vtkm::cont::ArrayHandleUniformPointCoordinates>() &&
return coordinates.GetData().IsType<vtkm::cont::ArrayHandleUniformPointCoordinates>() &&
(cellset.IsType<vtkm::cont::CellSetStructured<1>>() ||
cellset.IsType<vtkm::cont::CellSetStructured<2>>() ||
cellset.IsType<vtkm::cont::CellSetStructured<3>>());
@ -97,7 +97,7 @@ public:
"expecting ArrayHandleUniformPointCoordinates for coordinates");
auto cellId3 = static_cast<vtkm::Id3>((point - coords.GetOrigin()) / coords.GetSpacing());
auto cellDim = coords.GetDimensions() - vtkm::Id3(1);
auto cellDim = vtkm::Max(vtkm::Id3(1), coords.GetDimensions() - vtkm::Id3(1));
if (cellId3[0] < 0 || cellId3[0] >= cellDim[0] || cellId3[1] < 0 ||
cellId3[1] >= cellDim[1] || cellId3[2] < 0 || cellId3[2] >= cellDim[2])
{

22
vtkm/cont/DeviceAdapterAlgorithm.h
View File

@ -439,6 +439,26 @@ struct DeviceAdapterAlgorithm
///
VTKM_CONT static void Synchronize();
/// \brief Apply a given binary operation function element-wise to input arrays.
///
/// Apply the give binary operation to pairs of elements from the two input array
/// \c input1 and \c input2. The number of elements in the input arrays do not
/// have to be the same, in this case, only the smaller of the two numbers of elements
/// will be applied.
/// Outputs of the binary operation is stored in \c output.
///
template <typename T,
typename U,
typename V,
typename StorageT,
typename StorageU,
typename StorageV,
typename BinaryFunctor>
VTKM_CONT static void Transform(const vtkm::cont::ArrayHandle<T, StorageT>& input1,
const vtkm::cont::ArrayHandle<U, StorageU>& input2,
vtkm::cont::ArrayHandle<V, StorageV>& output,
BinaryFunctor binaryFunctor);
/// \brief Reduce an array to only the unique values it contains
///
/// Removes all duplicate values in \c values that are adjacent to each
@ -611,7 +631,7 @@ class DeviceAdapterAtomicArrayImplementation;
/// \brief Class providing a device-specific support for selecting the optimal
/// Task type for a given worklet.
///
/// When worklets are launched inside the execution enviornment we need to
/// When worklets are launched inside the execution environment we need to
/// ask the device adapter what is the preferred execution style, be it
/// a tiled iteration pattern, or strided. This class
///

4
vtkm/cont/DynamicArrayHandle.h
View File

@ -67,7 +67,7 @@ struct VTKM_CONT_EXPORT PolymorphicArrayHandleContainerBase
///
/// The \c PolymorphicArrayHandleContainer is similar to the
/// \c SimplePolymorphicContainer in that it can contain an object of an
/// unkown type. However, this class specifically holds ArrayHandle objects
/// unknown type. However, this class specifically holds ArrayHandle objects
/// (with different template parameters) so that it can polymorphically answer
/// simple questions about the object.
///
@ -289,7 +289,7 @@ public:
return this->CastToTypeStorage<ValueType, StorageTag>();
}
/// Given a refernce to an ArrayHandle object, casts this array to the
/// Given a references to an ArrayHandle object, casts this array to the
/// ArrayHandle's type and sets the given ArrayHandle to this array. Throws
/// \c ErrorBadType if the cast does not work. Use \c
/// ArrayHandleType to check if the cast can happen.

2
vtkm/cont/Field.h
View File

@ -308,7 +308,7 @@ void CastAndCall(const vtkm::cont::Field& field, Functor&& f, Args&&... args)
}
//@{
/// Convinience functions to build fields from C style arrays and std::vector
/// Convenience functions to build fields from C style arrays and std::vector
template <typename T>
vtkm::cont::Field make_Field(std::string name,
Field::AssociationEnum association,

10
vtkm/cont/MultiBlock.h
View File

@ -36,16 +36,16 @@ namespace cont
class VTKM_CONT_EXPORT MultiBlock
{
public:
/// creat a new MultiBlcok containng a single DataSet "ds"
/// create a new MultiBlock containng a single DataSet "ds"
VTKM_CONT
MultiBlock(const vtkm::cont::DataSet& ds);
/// creat a new MultiBlcok with the exisiting one "src"
/// create a new MultiBlock with the existing one "src"
VTKM_CONT
MultiBlock(const vtkm::cont::MultiBlock& src);
/// creat a new MultiBlcok with a DataSet vector "mblocks"
/// create a new MultiBlock with a DataSet vector "mblocks"
VTKM_CONT
MultiBlock(const std::vector<vtkm::cont::DataSet>& mblocks);
/// creat a new MultiBlcok with the capacity set to be "size"
/// create a new MultiBlock with the capacity set to be "size"
VTKM_CONT
MultiBlock(vtkm::Id size);
@ -57,7 +57,7 @@ public:
VTKM_CONT
~MultiBlock();
/// get the field "field_name" from blcok "block_index"
/// get the field "field_name" from block "block_index"
VTKM_CONT
vtkm::cont::Field GetField(const std::string& field_name, const int& block_index);

2
vtkm/cont/PointLocatorUniformGrid.h
View File

@ -175,7 +175,7 @@ public:
template <typename DeviceAdapter>
void Build(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords, DeviceAdapter)
{
typedef vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> Algorithm;
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);

4
vtkm/cont/RuntimeDeviceTracker.h
View File

@ -131,7 +131,7 @@ public:
///
/// If you want a \c RuntimeDeviceTracker with independent state, just create
/// one independently. If you want to start with the state of a source
/// \c RuntimeDeviceTracker but update the state indepenently, you can use
/// \c RuntimeDeviceTracker but update the state independently, you can use
/// \c DeepCopy method to get the initial state. Further changes will
/// not be shared.
///
@ -152,7 +152,7 @@ public:
///
/// If you want a \c RuntimeDeviceTracker with independent state, just create
/// one independently. If you want to start with the state of a source
/// \c RuntimeDeviceTracker but update the state indepenently, you can use
/// \c RuntimeDeviceTracker but update the state independently, you can use
/// \c DeepCopy method to get the initial state. Further changes will
/// not be shared.
///

2
vtkm/cont/Storage.h
View File

@ -132,7 +132,7 @@ public:
VTKM_CONT
PortalConstType GetPortalConst() const;
/// Retuns the number of entries allocated in the array.
/// Returns the number of entries allocated in the array.
VTKM_CONT
vtkm::Id GetNumberOfValues() const;

216
vtkm/cont/StorageBasic.cxx
View File

@ -20,6 +20,7 @@
#define vtkm_cont_StorageBasic_cxx
#include <vtkm/cont/StorageBasic.h>
#include <vtkm/internal/Configure.h>
#if defined(VTKM_POSIX)
#define VTKM_MEMALIGN_POSIX
@ -51,11 +52,7 @@ namespace cont
namespace internal
{
StorageBasicBase::~StorageBasicBase()
{
}
void* alloc_aligned(size_t size, size_t align)
void* StorageBasicAllocator::allocate(size_t size, size_t align)
{
#if defined(VTKM_MEMALIGN_POSIX)
void* mem = nullptr;
@ -70,14 +67,10 @@ void* alloc_aligned(size_t size, size_t align)
#else
void* mem = malloc(size);
#endif
if (mem == nullptr)
{
throw std::bad_alloc();
}
return mem;
}
void free_aligned(void* mem)
void StorageBasicAllocator::free_memory(void* mem)
{
#if defined(VTKM_MEMALIGN_POSIX)
free(mem);
@ -90,33 +83,188 @@ void free_aligned(void* mem)
#endif
}
template class VTKM_CONT_EXPORT Storage<char, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int8, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt8, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int16, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt16, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Float32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Float64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int64, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int32, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 2>, StorageTagBasic>;
StorageBasicBase::StorageBasicBase()
: Array(nullptr)
, AllocatedByteSize(0)
, NumberOfValues(0)
, DeallocateOnRelease(true)
{
}
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int64, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int32, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 3>, StorageTagBasic>;
StorageBasicBase::StorageBasicBase(const void* array,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue)
: Array(const_cast<void*>(array))
, AllocatedByteSize(static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue)
, NumberOfValues(numberOfValues)
, DeallocateOnRelease(array == nullptr ? true : false)
{
}
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<char, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Int8, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<UInt8, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 4>, StorageTagBasic>;
StorageBasicBase::~StorageBasicBase()
{
this->ReleaseResources();
}
StorageBasicBase::StorageBasicBase(const StorageBasicBase& src)
: Array(src.Array)
, AllocatedByteSize(src.AllocatedByteSize)
, NumberOfValues(src.NumberOfValues)
, DeallocateOnRelease(src.DeallocateOnRelease)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
}
StorageBasicBase StorageBasicBase::operator=(const StorageBasicBase& src)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
this->ReleaseResources();
this->Array = src.Array;
this->AllocatedByteSize = src.AllocatedByteSize;
this->NumberOfValues = src.NumberOfValues;
this->DeallocateOnRelease = src.DeallocateOnRelease;
return *this;
}
void StorageBasicBase::ReleaseResources()
{
if (this->AllocatedByteSize > 0)
{
VTKM_ASSERT(this->Array != nullptr);
if (this->DeallocateOnRelease)
{
AllocatorType{}.free_memory(this->Array);
}
this->Array = nullptr;
this->AllocatedByteSize = 0;
this->NumberOfValues = 0;
}
else
{
VTKM_ASSERT(this->Array == nullptr);
}
}
void StorageBasicBase::AllocateValues(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue)
{
if (numberOfValues < 0)
{
throw vtkm::cont::ErrorBadAllocation("Cannot allocate an array with negative size.");
}
// Check that the number of bytes won't be more than a size_t can hold.
const size_t maxNumValues = std::numeric_limits<size_t>::max() / sizeOfValue;
if (static_cast<vtkm::UInt64>(numberOfValues) > maxNumValues)
{
throw ErrorBadAllocation("Requested allocation exceeds size_t capacity.");
}
// If we are allocating less data, just shrink the array.
// (If allocation empty, drop down so we can deallocate memory.)
vtkm::UInt64 allocsize = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue;
if ((allocsize <= this->AllocatedByteSize) && (numberOfValues > 0))
{
this->NumberOfValues = numberOfValues;
return;
}
if (!this->DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue("User allocated arrays cannot be reallocated.");
}
this->ReleaseResources();
if (numberOfValues > 0)
{
this->Array = AllocatorType{}.allocate(allocsize, VTKM_ALLOCATION_ALIGNMENT);
this->AllocatedByteSize = allocsize;
this->NumberOfValues = numberOfValues;
if (this->Array == nullptr)
{
// Make sureour state is OK.
this->AllocatedByteSize = 0;
this->NumberOfValues = 0;
throw vtkm::cont::ErrorBadAllocation("Could not allocate basic control array.");
}
}
else
{
// ReleaseResources should have already set NumberOfValues to 0.
VTKM_ASSERT(this->NumberOfValues == 0);
VTKM_ASSERT(this->AllocatedByteSize == 0);
}
this->DeallocateOnRelease = true;
}
void StorageBasicBase::Shrink(vtkm::Id numberOfValues)
{
if (numberOfValues > this->NumberOfValues)
{
throw vtkm::cont::ErrorBadValue("Shrink method cannot be used to grow array.");
}
this->NumberOfValues = numberOfValues;
}
void* StorageBasicBase::GetBasePointer() const
{
return this->Array;
}
void* StorageBasicBase::GetEndPointer(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue) const
{
VTKM_ASSERT(this->NumberOfValues == numberOfValues);
if (!this->Array)
{
return nullptr;
}
auto p = static_cast<vtkm::UInt8*>(this->Array);
auto offset = static_cast<vtkm::UInt64>(this->NumberOfValues) * sizeOfValue;
return static_cast<void*>(p + offset);
}
void* StorageBasicBase::GetCapacityPointer() const
{
if (!this->Array)
{
return nullptr;
}
auto v = static_cast<vtkm::UInt8*>(this->Array) + AllocatedByteSize;
return static_cast<void*>(v);
}
#define _VTKM_STORAGE_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT Storage<Type, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 2>, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 3>, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 4>, StorageTagBasic>;
_VTKM_STORAGE_INSTANTIATE(char)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int8)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt8)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int16)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt16)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int32)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt32)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int64)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt64)
_VTKM_STORAGE_INSTANTIATE(vtkm::Float32)
_VTKM_STORAGE_INSTANTIATE(vtkm::Float64)
}
}
} // namespace vtkm::cont::internal

283
vtkm/cont/StorageBasic.h
View File

@ -28,11 +28,6 @@
#include <vtkm/cont/internal/ArrayPortalFromIterators.h>
// Defines the cache line size in bytes to align allocations to
#ifndef VTKM_CACHE_LINE_SIZE
#define VTKM_CACHE_LINE_SIZE 64
#endif
namespace vtkm
{
namespace cont
@ -46,126 +41,87 @@ struct VTKM_ALWAYS_EXPORT StorageTagBasic
namespace internal
{
VTKM_CONT_EXPORT
void* alloc_aligned(size_t size, size_t align);
VTKM_CONT_EXPORT
void free_aligned(void* mem);
/// \brief an aligned allocator
/// A simple aligned allocator type that will align allocations to `Alignment` bytes
/// TODO: Once C++11 std::allocator_traits is better used by STL and we want to drop
/// support for pre-C++11 we can drop a lot of the typedefs and functions here.
template <typename T, size_t Alignment>
struct AlignedAllocator
/// Class that does all of VTK-m allocations
/// for storage basic. This is exists so that
/// stolen arrays can call the correct free
/// function ( _aligned_malloc ) on windows
struct VTKM_CONT_EXPORT StorageBasicAllocator
{
using value_type = T;
using reference = T&;
using const_reference = const T&;
using pointer = T*;
using const_pointer = const T*;
using void_pointer = void*;
using const_void_pointer = const void*;
using difference_type = std::ptrdiff_t;
using size_type = std::size_t;
void* allocate(size_t size, size_t align);
void free_memory(void* p);
template <typename U>
struct rebind
template <typename T>
void deallocate(T* p)
{
using other = AlignedAllocator<U, Alignment>;
};
AlignedAllocator() {}
template <typename Tb>
AlignedAllocator(const AlignedAllocator<Tb, Alignment>&)
{
}
pointer allocate(size_t n)
{
return static_cast<pointer>(alloc_aligned(n * sizeof(T), Alignment));
}
void deallocate(pointer p, size_t) { free_aligned(static_cast<void*>(p)); }
pointer address(reference r) { return &r; }
const_pointer address(const_reference r) { return &r; }
size_type max_size() const { return (std::numeric_limits<size_type>::max)() / sizeof(T); }
void construct(pointer p, const T& t)
{
(void)p;
new (p) T(t);
}
void destroy(pointer p)
{
(void)p;
p->~T();
this->free_memory(static_cast<void*>(p));
}
};
template <typename T, typename U, size_t AlignA, size_t AlignB>
bool operator==(const AlignedAllocator<T, AlignA>&, const AlignedAllocator<U, AlignB>&)
{
return AlignA == AlignB;
}
template <typename T, typename U, size_t AlignA, size_t AlignB>
bool operator!=(const AlignedAllocator<T, AlignA>&, const AlignedAllocator<U, AlignB>&)
{
return AlignA != AlignB;
}
/// Base class for basic storage classes. This is currently only used by
/// Basic storage to provide a type-agnostic API for allocations, etc.
/// Base class for basic storage classes. This allow us to implement
/// vtkm::cont::Storage<T, StorageTagBasic > for any T type with no overhead
/// as all heavy logic is provide by a type-agnostic API including allocations, etc.
class VTKM_CONT_EXPORT StorageBasicBase
{
public:
StorageBasicBase() {}
virtual ~StorageBasicBase();
using AllocatorType = StorageBasicAllocator;
VTKM_CONT StorageBasicBase();
VTKM_CONT StorageBasicBase(const void* array, vtkm::Id size, vtkm::UInt64 sizeOfValue);
VTKM_CONT ~StorageBasicBase();
/// \brief Return the number of bytes allocated for this storage object.
VTKM_CONT
virtual vtkm::UInt64 GetNumberOfBytes() const = 0;
VTKM_CONT StorageBasicBase(const StorageBasicBase& src);
VTKM_CONT StorageBasicBase operator=(const StorageBasicBase& src);
/// \brief Allocates an array with the specified size in bytes.
/// \brief Return the number of bytes allocated for this storage object(Capacity).
///
///
VTKM_CONT vtkm::UInt64 GetNumberOfBytes() const { return this->AllocatedByteSize; }
/// \brief Return the number of 'T' values allocated by this storage
VTKM_CONT vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
/// \brief Allocates an array with the specified number of elements.
///
/// The allocation may be done on an already existing array, but can wipe out
/// any data already in the array. This method can throw
/// ErrorBadAllocation if the array cannot be allocated or
/// ErrorBadValue if the allocation is not feasible (for example, the
/// array storage is read-only).
VTKM_CONT
virtual void AllocateBytes(vtkm::UInt64 numberOfBytes) = 0;
VTKM_CONT void AllocateValues(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue);
/// \brief Reduces the size of the array without changing its values.
///
/// This method allows you to resize the array without reallocating it. The
/// size of the array is changed to \c numberOfBytes bytes. The data
/// in the reallocated array stays the same, but \c numberOfBytes must be
/// size of the array is changed so that it can hold \c numberOfValues values.
/// The data in the reallocated array stays the same, but \c numberOfValues must be
/// equal or less than the preexisting size. That is, this method can only be
/// used to shorten the array, not lengthen.
VTKM_CONT
virtual void ShrinkBytes(vtkm::UInt64 numberOfBytes) = 0;
VTKM_CONT void Shrink(vtkm::Id numberOfValues);
/// \brief Frees any resources (i.e. memory) stored in this array.
///
/// After calling this method GetNumberOfBytes() will return 0. The
/// resources should also be released when the Storage class is
/// destroyed.
VTKM_CONT
virtual void ReleaseResources() = 0;
VTKM_CONT void ReleaseResources();
/// \brief Returns if vtkm will deallocate this memory. VTK-m StorageBasic
/// is designed that VTK-m will not deallocate user passed memory, or
/// instances that have been stolen (\c StealArray)
VTKM_CONT bool WillDeallocate() const { return this->DeallocateOnRelease; }
/// Return the memory location of the first element of the array data.
VTKM_CONT
virtual void* GetBasePointer() const = 0;
VTKM_CONT void* GetBasePointer() const;
/// Return the memory location of the first element past the end of the array
/// data.
VTKM_CONT
virtual void* GetEndPointer() const = 0;
VTKM_CONT void* GetEndPointer(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue) const;
/// Return the memory location of the first element past the end of the
/// array's allocated memory buffer.
VTKM_CONT
virtual void* GetCapacityPointer() const = 0;
VTKM_CONT void* GetCapacityPointer() const;
protected:
void* Array;
vtkm::UInt64 AllocatedByteSize;
vtkm::Id NumberOfValues;
bool DeallocateOnRelease;
};
/// A basic implementation of an Storage object.
@ -180,70 +136,23 @@ template <typename ValueT>
class VTKM_ALWAYS_EXPORT Storage<ValueT, vtkm::cont::StorageTagBasic> : public StorageBasicBase
{
public:
using AllocatorType = vtkm::cont::internal::StorageBasicAllocator;
using ValueType = ValueT;
using PortalType = vtkm::cont::internal::ArrayPortalFromIterators<ValueType*>;
using PortalConstType = vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*>;
/// The original design of this class provided an allocator as a template
/// parameters. That messed things up, though, because other templated
/// classes assume that the \c Storage has one template parameter. There are
/// other ways to allow you to specify the allocator, but it is uncertain
/// whether that would ever be useful. So, instead of jumping through hoops
/// implementing them, just fix the allocator for now.
///
using AllocatorType = AlignedAllocator<ValueType, VTKM_CACHE_LINE_SIZE>;
public:
/// \brief construct storage that VTK-m is responsible for
VTKM_CONT
Storage();
VTKM_CONT Storage();
/// \brief construct storage that VTK-m is not responsible for
VTKM_CONT
Storage(const ValueType* array, vtkm::Id numberOfValues = 0);
VTKM_CONT Storage(const ValueType* array, vtkm::Id numberOfValues = 0);
VTKM_CONT
~Storage();
VTKM_CONT void Allocate(vtkm::Id numberOfValues);
VTKM_CONT
Storage(const Storage<ValueType, StorageTagBasic>& src);
VTKM_CONT PortalType GetPortal();
VTKM_CONT
Storage& operator=(const Storage<ValueType, StorageTagBasic>& src);
VTKM_CONT
void ReleaseResources() final;
VTKM_CONT
void Allocate(vtkm::Id numberOfValues);
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
VTKM_CONT
vtkm::UInt64 GetNumberOfBytes() const final
{
return static_cast<vtkm::UInt64>(this->NumberOfValues) *
static_cast<vtkm::UInt64>(sizeof(ValueT));
}
VTKM_CONT
void Shrink(vtkm::Id numberOfValues);
VTKM_CONT
void AllocateBytes(vtkm::UInt64) final;
VTKM_CONT
void ShrinkBytes(vtkm::UInt64) final;
VTKM_CONT
PortalType GetPortal() { return PortalType(this->Array, this->Array + this->NumberOfValues); }
VTKM_CONT
PortalConstType GetPortalConst() const
{
return PortalConstType(this->Array, this->Array + this->NumberOfValues);
}
VTKM_CONT PortalConstType GetPortalConst() const;
/// \brief Get a pointer to the underlying data structure.
///
@ -251,10 +160,9 @@ public:
/// memory associated with this array still belongs to the Storage (i.e.
/// Storage will eventually deallocate the array).
///
VTKM_CONT
ValueType* GetArray() { return this->Array; }
VTKM_CONT
const ValueType* GetArray() const { return this->Array; }
VTKM_CONT ValueType* GetArray();
VTKM_CONT const ValueType* GetArray() const;
/// \brief Take the reference away from this object.
///
@ -265,36 +173,7 @@ public:
/// VTK-m and not having to keep a VTK-m object around. Obviously the caller
/// becomes responsible for destroying the memory.
///
VTKM_CONT
ValueType* StealArray();
/// \brief Returns if vtkm will deallocate this memory. VTK-m StorageBasic
/// is designed that VTK-m will not deallocate user passed memory, or
/// instances that have been stolen (\c StealArray)
VTKM_CONT
bool WillDeallocate() const { return this->DeallocateOnRelease; }
VTKM_CONT
void* GetBasePointer() const final { return static_cast<void*>(this->Array); }
VTKM_CONT
void* GetEndPointer() const final
{
return static_cast<void*>(this->Array + this->NumberOfValues);
}
VTKM_CONT
void* GetCapacityPointer() const final
{
return static_cast<void*>(this->Array + this->AllocatedSize);
}
private:
ValueType* Array;
vtkm::Id NumberOfValues;
vtkm::Id AllocatedSize;
bool DeallocateOnRelease;
VTKM_CONT ValueType* StealArray();
};
} // namespace internal
@ -311,39 +190,23 @@ namespace internal
/// \cond
/// Make doxygen ignore this section
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<char, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int8, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt8, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int16, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt16, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int64, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt64, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Float32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Float64, StorageTagBasic>;
#define _VTKM_STORAGE_EXPORT(Type) \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<Type, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 2>, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 3>, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int64, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int32, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int64, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int32, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<char, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Int8, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<UInt8, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 4>, StorageTagBasic>;
_VTKM_STORAGE_EXPORT(char)
_VTKM_STORAGE_EXPORT(vtkm::Int8)
_VTKM_STORAGE_EXPORT(vtkm::UInt8)
_VTKM_STORAGE_EXPORT(vtkm::Int16)
_VTKM_STORAGE_EXPORT(vtkm::UInt16)
_VTKM_STORAGE_EXPORT(vtkm::Int32)
_VTKM_STORAGE_EXPORT(vtkm::UInt32)
_VTKM_STORAGE_EXPORT(vtkm::Int64)
_VTKM_STORAGE_EXPORT(vtkm::UInt64)
_VTKM_STORAGE_EXPORT(vtkm::Float32)
_VTKM_STORAGE_EXPORT(vtkm::Float64)
/// \endcond
}
}

156
vtkm/cont/StorageBasic.hxx
View File

@ -31,173 +31,55 @@ namespace internal
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage()
: Array(nullptr)
, NumberOfValues(0)
, AllocatedSize(0)
, DeallocateOnRelease(true)
: StorageBasicBase()
{
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage(const T* array, vtkm::Id numberOfValues)
: Array(const_cast<T*>(array))
, NumberOfValues(numberOfValues)
, AllocatedSize(numberOfValues)
, DeallocateOnRelease(array == nullptr ? true : false)
: StorageBasicBase(const_cast<T*>(array), numberOfValues, sizeof(T))
{
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::~Storage()
{
this->ReleaseResources();
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage(const Storage<T, StorageTagBasic>& src)
: Array(src.Array)
, NumberOfValues(src.NumberOfValues)
, AllocatedSize(src.AllocatedSize)
, DeallocateOnRelease(src.DeallocateOnRelease)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>& Storage<T, vtkm::cont::StorageTagBasic>::operator=(
const Storage<T, StorageTagBasic>& src)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
this->ReleaseResources();
this->Array = src.Array;
this->NumberOfValues = src.NumberOfValues;
this->AllocatedSize = src.AllocatedSize;
this->DeallocateOnRelease = src.DeallocateOnRelease;
return *this;
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::ReleaseResources()
{
if (this->NumberOfValues > 0)
{
VTKM_ASSERT(this->Array != nullptr);
if (this->DeallocateOnRelease)
{
AllocatorType allocator;
allocator.deallocate(this->Array, static_cast<std::size_t>(this->AllocatedSize));
}
this->Array = nullptr;
this->NumberOfValues = 0;
this->AllocatedSize = 0;
}
else
{
VTKM_ASSERT(this->Array == nullptr);
}
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::Allocate(vtkm::Id numberOfValues)
{
if (numberOfValues < 0)
{
throw vtkm::cont::ErrorBadAllocation("Cannot allocate an array with negative size.");
}
// Check that the number of bytes won't be more than a size_t can hold.
const size_t maxNumValues = std::numeric_limits<size_t>::max() / sizeof(T);
if (static_cast<vtkm::UInt64>(numberOfValues) > static_cast<vtkm::UInt64>(maxNumValues))
{
throw ErrorBadAllocation("Requested allocation exceeds size_t capacity.");
}
this->AllocateBytes(static_cast<vtkm::UInt64>(numberOfValues) *
static_cast<vtkm::UInt64>(sizeof(T)));
this->AllocateValues(numberOfValues, sizeof(T));
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::AllocateBytes(vtkm::UInt64 numberOfBytes)
typename Storage<T, vtkm::cont::StorageTagBasic>::PortalType
Storage<T, vtkm::cont::StorageTagBasic>::GetPortal()
{
const vtkm::Id numberOfValues =
static_cast<vtkm::Id>(numberOfBytes / static_cast<vtkm::UInt64>(sizeof(T)));
// If we are allocating less data, just shrink the array.
// (If allocation empty, drop down so we can deallocate memory.)
if ((numberOfValues <= this->AllocatedSize) && (numberOfValues > 0))
{
this->NumberOfValues = numberOfValues;
return;
}
if (!this->DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue("User allocated arrays cannot be reallocated.");
}
this->ReleaseResources();
try
{
if (numberOfValues > 0)
{
AllocatorType allocator;
this->Array = allocator.allocate(static_cast<std::size_t>(numberOfValues));
this->AllocatedSize = numberOfValues;
this->NumberOfValues = numberOfValues;
}
else
{
// ReleaseResources should have already set AllocatedSize to 0.
VTKM_ASSERT(this->AllocatedSize == 0);
}
}
catch (std::bad_alloc&)
{
// Make sureour state is OK.
this->Array = nullptr;
this->NumberOfValues = 0;
this->AllocatedSize = 0;
throw vtkm::cont::ErrorBadAllocation("Could not allocate basic control array.");
}
this->DeallocateOnRelease = true;
auto v = static_cast<T*>(this->Array);
return PortalType(v, v + this->NumberOfValues);
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::Shrink(vtkm::Id numberOfValues)
typename Storage<T, vtkm::cont::StorageTagBasic>::PortalConstType
Storage<T, vtkm::cont::StorageTagBasic>::GetPortalConst() const
{
this->ShrinkBytes(static_cast<vtkm::UInt64>(numberOfValues) *
static_cast<vtkm::UInt64>(sizeof(T)));
auto v = static_cast<T*>(this->Array);
return PortalConstType(v, v + this->NumberOfValues);
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::ShrinkBytes(vtkm::UInt64 numberOfBytes)
T* Storage<T, vtkm::cont::StorageTagBasic>::GetArray()
{
if (numberOfBytes > this->GetNumberOfBytes())
{
throw vtkm::cont::ErrorBadValue("Shrink method cannot be used to grow array.");
}
return static_cast<T*>(this->Array);
}
this->NumberOfValues =
static_cast<vtkm::Id>(numberOfBytes / static_cast<vtkm::UInt64>(sizeof(T)));
template <typename T>
const T* Storage<T, vtkm::cont::StorageTagBasic>::GetArray() const
{
return static_cast<T*>(this->Array);
}
template <typename T>
T* Storage<T, vtkm::cont::StorageTagBasic>::StealArray()
{
this->DeallocateOnRelease = false;
return this->Array;
return static_cast<T*>(this->Array);
}
} // namespace internal

2
vtkm/cont/Timer.h
View File

@ -27,7 +27,7 @@ namespace vtkm
namespace cont
{
/// A class that can be used to time operations in VTK-m that might be occuring
/// A class that can be used to time operations in VTK-m that might be occurring
/// in parallel. You should make sure that the device adapter for the timer
/// matches that being used to execute algorithms to ensure that the thread
/// synchronization is correct.

171
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu
View File

@ -21,182 +21,11 @@
#define vtk_m_cont_cuda_internal_ArrayManagerExecutionCuda_cu
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
using vtkm::cont::cuda::internal::CudaAllocator;
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::ExecutionArrayInterfaceBasic(
StorageBasicBase& storage)
: Superclass(storage)
{
}
DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_CUDA;
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Allocate(TypelessExecutionArray& execArray,
vtkm::UInt64 numBytes) const
{
// Detect if we can reuse a device-accessible pointer from the control env:
if (CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
const vtkm::UInt64 managedCapacity =
static_cast<vtkm::UInt64>(static_cast<const char*>(execArray.ArrayControlCapacity) -
static_cast<const char*>(execArray.ArrayControl));
if (managedCapacity >= numBytes)
{
if (execArray.Array && execArray.Array != execArray.ArrayControl)
{
this->Free(execArray);
}
execArray.Array = const_cast<void*>(execArray.ArrayControl);
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
execArray.ArrayCapacity = const_cast<void*>(execArray.ArrayControlCapacity);
return;
}
}
if (execArray.Array != nullptr)
{
const vtkm::UInt64 cap = static_cast<vtkm::UInt64>(static_cast<char*>(execArray.ArrayCapacity) -
static_cast<char*>(execArray.Array));
if (cap < numBytes)
{ // Current allocation too small -- free & realloc
this->Free(execArray);
}
else
{ // Reuse buffer if possible:
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
return;
}
}
VTKM_ASSERT(execArray.Array == nullptr);
// Attempt to allocate:
try
{
// Cast to char* so that the pointer math below will work.
char* tmp = static_cast<char*>(CudaAllocator::Allocate(static_cast<size_t>(numBytes)));
execArray.Array = tmp;
execArray.ArrayEnd = tmp + numBytes;
execArray.ArrayCapacity = tmp + numBytes;
}
catch (const std::exception& error)
{
std::ostringstream err;
err << "Failed to allocate " << numBytes << " bytes on device: " << error.what();
throw vtkm::cont::ErrorBadAllocation(err.str());
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Free(
TypelessExecutionArray& execArray) const
{
// If we're sharing a device-accessible pointer between control/exec, don't
// actually free it -- just null the pointers here:
if (execArray.Array == execArray.ArrayControl &&
CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
return;
}
if (execArray.Array != nullptr)
{
CudaAllocator::Free(execArray.Array);
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a device-accessible pointer between control and
// execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
CudaAllocator::PrepareForInput(executionPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(executionPtr,
controlPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyHostToDevice,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a cuda managed pointer between control and execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
// If we're trying to copy a shared, non-managed device pointer back to
// control throw an exception -- the pointer cannot be read from control,
// so this operation is invalid.
if (!CudaAllocator::IsManagedPointer(controlPtr))
{
throw vtkm::cont::ErrorBadValue(
"Control pointer is a CUDA device pointer that does not supported managed access.");
}
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInPlace(executionPtr, static_cast<size_t>(numBytes));
}
} // end namespace internal
VTKM_INSTANTIATE_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagCuda)
}
} // end vtkm::cont

41
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
View File

@ -26,6 +26,10 @@
#include <vtkm/cont/internal/ArrayExportMacros.h>
#include <vtkm/cont/internal/ArrayManagerExecution.h>
//This is in a separate header so that ArrayHandleBasicImpl can include
//the interface without getting any CUDA headers
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
@ -127,49 +131,16 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagCuda>
VTKM_CONT
static PortalType CreatePortal(ValueType* start, ValueType* end)
{
using ThrustPointerT = thrust::system::cuda::pointer<ValueType>;
ThrustPointerT startThrust(start);
ThrustPointerT endThrust(end);
return PortalType(startThrust, endThrust);
return PortalType(start, end);
}
VTKM_CONT
static PortalConstType CreatePortalConst(const ValueType* start, const ValueType* end)
{
using ThrustPointerT = thrust::system::cuda::pointer<const ValueType>;
ThrustPointerT startThrust(start);
ThrustPointerT endThrust(end);
return PortalConstType(startThrust, endThrust);
return PortalConstType(start, end);
}
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
: public ExecutionArrayInterfaceBasicBase
{
using Superclass = ExecutionArrayInterfaceBasicBase;
VTKM_CONT ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT DeviceAdapterId GetDeviceId() const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
} // namespace internal
#ifndef vtk_m_cont_cuda_internal_ArrayManagerExecutionCuda_cu

66
vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h
View File

@ -24,21 +24,16 @@
#include <vtkm/cont/ErrorBadAllocation.h>
#include <vtkm/cont/Storage.h>
// Disable warnings we check vtkm for but Thrust does not.
VTKM_THIRDPARTY_PRE_INCLUDE
#include <thrust/copy.h>
#include <thrust/device_malloc_allocator.h>
#include <thrust/system/cuda/vector.h>
#include <thrust/system/cuda/execution_policy.h>
VTKM_THIRDPARTY_POST_INCLUDE
#include <vtkm/cont/cuda/ErrorCuda.h>
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
#include <vtkm/cont/cuda/internal/ThrustExceptionHandler.h>
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <thrust/copy.h>
#include <thrust/device_ptr.h>
VTKM_THIRDPARTY_POST_INCLUDE
#include <limits>
namespace vtkm
@ -53,15 +48,13 @@ namespace internal
/// \c ArrayManagerExecutionThrustDevice provides an implementation for a \c
/// ArrayManagerExecution class for a thrust device adapter that is designed
/// for the cuda backend which has separate memory spaces for host and device.
/// This implementation contains a thrust::system::cuda::pointer to contain the
/// data.
template <typename T, class StorageTag>
class ArrayManagerExecutionThrustDevice
{
public:
using ValueType = T;
using PointerType = typename thrust::system::cuda::pointer<ValueType>;
using difference_type = typename PointerType::difference_type;
using PointerType = T*;
using difference_type = std::ptrdiff_t;
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
@ -71,9 +64,9 @@ public:
VTKM_CONT
ArrayManagerExecutionThrustDevice(StorageType* storage)
: Storage(storage)
, Begin(static_cast<ValueType*>(nullptr))
, End(static_cast<ValueType*>(nullptr))
, Capacity(static_cast<ValueType*>(nullptr))
, Begin(nullptr)
, End(nullptr)
, Capacity(nullptr)
{
}
@ -83,10 +76,7 @@ public:
/// Returns the size of the array.
///
VTKM_CONT
vtkm::Id GetNumberOfValues() const
{
return static_cast<vtkm::Id>(this->End.get() - this->Begin.get());
}
vtkm::Id GetNumberOfValues() const { return static_cast<vtkm::Id>(this->End - this->Begin); }
/// Allocates the appropriate size of the array and copies the given data
/// into the array.
@ -144,14 +134,13 @@ public:
PortalType PrepareForOutput(vtkm::Id numberOfValues)
{
// Can we reuse the existing buffer?
vtkm::Id curCapacity = this->Begin.get() != nullptr
? static_cast<vtkm::Id>(this->Capacity.get() - this->Begin.get())
: 0;
vtkm::Id curCapacity =
this->Begin != nullptr ? static_cast<vtkm::Id>(this->Capacity - this->Begin) : 0;
// Just mark a new end if we don't need to increase the allocation:
if (curCapacity >= numberOfValues)
{
this->End = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
this->End = this->Begin + static_cast<difference_type>(numberOfValues);
return PortalType(this->Begin, this->End);
}
@ -173,9 +162,8 @@ public:
// Attempt to allocate:
try
{
ValueType* tmp =
this->Begin =
static_cast<ValueType*>(vtkm::cont::cuda::internal::CudaAllocator::Allocate(bufferSize));
this->Begin = PointerType(tmp);
}
catch (const std::exception& error)
{
@ -184,7 +172,7 @@ public:
throw vtkm::cont::ErrorBadAllocation(err.str());
}
this->Capacity = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
this->Capacity = this->Begin + static_cast<difference_type>(numberOfValues);
this->End = this->Capacity;
return PortalType(this->Begin, this->End);
@ -206,8 +194,9 @@ public:
storage->Allocate(this->GetNumberOfValues());
try
{
::thrust::copy(
this->Begin, this->End, vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
::thrust::copy(thrust::cuda::pointer<ValueType>(this->Begin),
thrust::cuda::pointer<ValueType>(this->End),
vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
}
catch (...)
{
@ -221,22 +210,21 @@ public:
{
// The operation will succeed even if this assertion fails, but this
// is still supposed to be a precondition to Shrink.
VTKM_ASSERT(this->Begin.get() != nullptr &&
this->Begin.get() + numberOfValues <= this->End.get());
VTKM_ASSERT(this->Begin != nullptr && this->Begin + numberOfValues <= this->End);
this->End = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
this->End = this->Begin + static_cast<difference_type>(numberOfValues);
}
/// Frees all memory.
///
VTKM_CONT void ReleaseResources()
{
if (this->Begin.get() != nullptr)
if (this->Begin != nullptr)
{
vtkm::cont::cuda::internal::CudaAllocator::Free(this->Begin.get());
this->Begin = PointerType(static_cast<ValueType*>(nullptr));
this->End = PointerType(static_cast<ValueType*>(nullptr));
this->Capacity = PointerType(static_cast<ValueType*>(nullptr));
vtkm::cont::cuda::internal::CudaAllocator::Free(this->Begin);
this->Begin = nullptr;
this->End = nullptr;
this->Capacity = nullptr;
}
}
@ -258,7 +246,7 @@ private:
this->PrepareForOutput(this->Storage->GetNumberOfValues());
::thrust::copy(vtkm::cont::ArrayPortalToIteratorBegin(this->Storage->GetPortalConst()),
vtkm::cont::ArrayPortalToIteratorEnd(this->Storage->GetPortalConst()),
this->Begin);
thrust::cuda::pointer<ValueType>(this->Begin));
}
catch (...)
{

3
vtkm/cont/cuda/internal/CMakeLists.txt
View File

@ -25,11 +25,13 @@ set(headers
DeviceAdapterAlgorithmCuda.h
DeviceAdapterAlgorithmThrust.h
DeviceAdapterTagCuda.h
ExecutionArrayInterfaceBasicCuda.h
MakeThrustIterator.h
TaskTuner.h
ThrustExceptionHandler.h
VirtualObjectTransferCuda.h
)
vtkm_declare_headers(CUDA ${headers} TESTABLE ${VTKm_ENABLE_CUDA})
#-----------------------------------------------------------------------------
@ -39,5 +41,6 @@ endif()
target_sources(vtkm_cont PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/ArrayManagerExecutionCuda.cu
${CMAKE_CURRENT_SOURCE_DIR}/ExecutionArrayInterfaceBasicCuda.cu
${CMAKE_CURRENT_SOURCE_DIR}/CudaAllocator.cu
)

2
vtkm/cont/cuda/internal/CudaAllocator.h
View File

@ -42,7 +42,7 @@ struct VTKM_CONT_EXPORT CudaAllocator
/// that can be accessed concurrently by the CPU and GPUs.
static VTKM_CONT bool UsingManagedMemory();
/// Returns true if the pointer is accessable from a CUDA device.
/// Returns true if the pointer is accessible from a CUDA device.
static VTKM_CONT bool IsDevicePointer(const void* ptr);
/// Returns true if the pointer is a CUDA pointer allocated with

15
vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h
View File

@ -53,14 +53,13 @@ VTKM_THIRDPARTY_PRE_INCLUDE
#include <thrust/binary_search.h>
#include <thrust/copy.h>
#include <thrust/count.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/scan.h>
#include <thrust/sort.h>
#include <thrust/system/cpp/memory.h>
#include <thrust/system/cuda/vector.h>
#include <thrust/unique.h>
#include <vtkm/exec/cuda/internal/ExecutionPolicy.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/system/cuda/execution_policy.h>
VTKM_THIRDPARTY_POST_INCLUDE
#include <atomic>
@ -218,8 +217,7 @@ private:
OutputPortal output,
UnaryPredicate unary_predicate)
{
using IteratorType = typename detail::IteratorTraits<OutputPortal>::IteratorType;
IteratorType outputBegin = IteratorBegin(output);
auto outputBegin = IteratorBegin(output);
using ValueType = typename StencilPortal::ValueType;
@ -356,7 +354,7 @@ private:
BinaryFunctor binary_functor,
std::false_type)
{
//The portal type and the initial value ARENT the same type so we have
//The portal type and the initial value AREN'T the same type so we have
//to a slower approach, where we wrap the input portal inside a cast
//portal
using CastFunctor = vtkm::cont::internal::Cast<typename InputPortal::ValueType, T>;
@ -497,12 +495,15 @@ private:
try
{
::thrust::system::cuda::vector<ValueType> result(1);
auto end = ::thrust::inclusive_scan(ThrustCudaPolicyPerThread,
IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(output),
bop);
return *(end - 1);
::thrust::copy_n(ThrustCudaPolicyPerThread, end - 1, 1, result.begin());
return result[0];
}
catch (...)
{

199
vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.cu Normal file
View File

@ -0,0 +1,199 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
using vtkm::cont::cuda::internal::CudaAllocator;
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::ExecutionArrayInterfaceBasic(
StorageBasicBase& storage)
: Superclass(storage)
{
}
DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_CUDA;
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const
{
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue;
// Detect if we can reuse a device-accessible pointer from the control env:
if (CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
const vtkm::UInt64 managedCapacity =
static_cast<vtkm::UInt64>(static_cast<const char*>(execArray.ArrayControlCapacity) -
static_cast<const char*>(execArray.ArrayControl));
if (managedCapacity >= numBytes)
{
if (execArray.Array && execArray.Array != execArray.ArrayControl)
{
this->Free(execArray);
}
execArray.Array = const_cast<void*>(execArray.ArrayControl);
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
execArray.ArrayCapacity = const_cast<void*>(execArray.ArrayControlCapacity);
return;
}
}
if (execArray.Array != nullptr)
{
const vtkm::UInt64 cap = static_cast<vtkm::UInt64>(static_cast<char*>(execArray.ArrayCapacity) -
static_cast<char*>(execArray.Array));
if (cap < numBytes)
{ // Current allocation too small -- free & realloc
this->Free(execArray);
}
else
{ // Reuse buffer if possible:
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
return;
}
}
VTKM_ASSERT(execArray.Array == nullptr);
// Attempt to allocate:
try
{
// Cast to char* so that the pointer math below will work.
char* tmp = static_cast<char*>(CudaAllocator::Allocate(static_cast<size_t>(numBytes)));
execArray.Array = tmp;
execArray.ArrayEnd = tmp + numBytes;
execArray.ArrayCapacity = tmp + numBytes;
}
catch (const std::exception& error)
{
std::ostringstream err;
err << "Failed to allocate " << numBytes << " bytes on device: " << error.what();
throw vtkm::cont::ErrorBadAllocation(err.str());
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Free(
TypelessExecutionArray& execArray) const
{
// If we're sharing a device-accessible pointer between control/exec, don't
// actually free it -- just null the pointers here:
if (execArray.Array == execArray.ArrayControl &&
CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
return;
}
if (execArray.Array != nullptr)
{
CudaAllocator::Free(execArray.Array);
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a device-accessible pointer between control and
// execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
CudaAllocator::PrepareForInput(executionPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(executionPtr,
controlPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyHostToDevice,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a cuda managed pointer between control and execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
// If we're trying to copy a shared, non-managed device pointer back to
// control throw an exception -- the pointer cannot be read from control,
// so this operation is invalid.
if (!CudaAllocator::IsManagedPointer(controlPtr))
{
throw vtkm::cont::ErrorBadValue(
"Control pointer is a CUDA device pointer that does not supported managed access.");
}
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInPlace(executionPtr, static_cast<size_t>(numBytes));
}
} // end namespace internal
}
} // end vtkm::cont

64
vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h Normal file
View File

@ -0,0 +1,64 @@
//============================================================================
// 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 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 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_cuda_internal_ExecutionArrayInterfaceCuda_h
#define vtk_m_cont_cuda_internal_ExecutionArrayInterfaceCuda_h
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
: public ExecutionArrayInterfaceBasicBase
{
using Superclass = ExecutionArrayInterfaceBasicBase;
VTKM_CONT ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT DeviceAdapterId GetDeviceId() const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_cuda_internal_ExecutionArrayInterfaceCuda_h

128
vtkm/cont/cuda/internal/MakeThrustIterator.h
View File

@ -20,21 +20,8 @@
#ifndef vtk_m_cont_cuda_internal_MakeThrustIterator_h
#define vtk_m_cont_cuda_internal_MakeThrustIterator_h
#include <vtkm/Pair.h>
#include <vtkm/Types.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/internal/ExportMacros.h>
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
#include <vtkm/exec/cuda/internal/WrappedOperators.h>
// Disable warnings we check vtkm for but Thrust does not.
VTKM_THIRDPARTY_PRE_INCLUDE
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/system/cuda/memory.h>
VTKM_THIRDPARTY_POST_INCLUDE
#include <vtkm/exec/cuda/internal/IteratorFromArrayPortal.h>
namespace vtkm
{
@ -44,107 +31,46 @@ namespace cuda
{
namespace internal
{
namespace detail
{
// Tags to specify what type of thrust iterator to use.
struct ThrustIteratorFromArrayPortalTag
{
};
struct ThrustIteratorDevicePtrTag
{
};
// Traits to help classify what thrust iterators will be used.
template <typename IteratorType>
struct ThrustIteratorTag
{
using Type = ThrustIteratorFromArrayPortalTag;
};
template <typename T>
struct ThrustIteratorTag<thrust::system::cuda::pointer<T>>
{
using Type = ThrustIteratorDevicePtrTag;
};
template <typename T>
struct ThrustIteratorTag<thrust::system::cuda::pointer<const T>>
{
using Type = ThrustIteratorDevicePtrTag;
};
template <typename PortalType, typename Tag>
struct IteratorChooser;
template <typename PortalType>
struct IteratorChooser<PortalType, detail::ThrustIteratorFromArrayPortalTag>
{
using Type = vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>;
};
template <typename PortalType>
struct IteratorChooser<PortalType, detail::ThrustIteratorDevicePtrTag>
{
using PortalToIteratorType = vtkm::cont::ArrayPortalToIterators<PortalType>;
using Type = typename PortalToIteratorType::IteratorType;
};
template <typename PortalType>
struct IteratorTraits
{
using PortalToIteratorType = vtkm::cont::ArrayPortalToIterators<PortalType>;
using Tag = typename detail::ThrustIteratorTag<typename PortalToIteratorType::IteratorType>::Type;
using IteratorType = typename IteratorChooser<PortalType, Tag>::Type;
};
template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
PortalType portal,
detail::ThrustIteratorFromArrayPortalTag)
{
return vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>(portal);
}
template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
PortalType portal,
detail::ThrustIteratorDevicePtrTag)
{
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
return iterators.GetBegin();
}
template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
PortalType portal,
detail::ThrustIteratorFromArrayPortalTag)
inline vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> IteratorBegin(
const PortalType& portal)
{
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
::thrust::advance(iterator, static_cast<std::size_t>(portal.GetNumberOfValues()));
return iterator;
}
template <typename PortalType>
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
PortalType portal,
detail::ThrustIteratorDevicePtrTag)
inline vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> IteratorEnd(
const PortalType& portal)
{
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
return iterators.GetEnd();
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
iterator += static_cast<std::size_t>(portal.GetNumberOfValues());
return iterator;
}
} // namespace detail
template <typename PortalType>
VTKM_CONT typename detail::IteratorTraits<PortalType>::IteratorType IteratorBegin(PortalType portal)
template <typename T>
inline T* IteratorBegin(const vtkm::exec::cuda::internal::ArrayPortalFromThrust<T>& portal)
{
using IteratorTag = typename detail::IteratorTraits<PortalType>::Tag;
return detail::MakeIteratorBegin(portal, IteratorTag());
return portal.GetIteratorBegin();
}
template <typename PortalType>
VTKM_CONT typename detail::IteratorTraits<PortalType>::IteratorType IteratorEnd(PortalType portal)
template <typename T>
inline T* IteratorEnd(const vtkm::exec::cuda::internal::ArrayPortalFromThrust<T>& portal)
{
using IteratorTag = typename detail::IteratorTraits<PortalType>::Tag;
return detail::MakeIteratorEnd(portal, IteratorTag());
return portal.GetIteratorEnd();
}
template <typename T>
inline const T* IteratorBegin(
const vtkm::exec::cuda::internal::ConstArrayPortalFromThrust<T>& portal)
{
return portal.GetIteratorBegin();
}
template <typename T>
inline const T* IteratorEnd(const vtkm::exec::cuda::internal::ConstArrayPortalFromThrust<T>& portal)
{
return portal.GetIteratorEnd();
}
}
}

28
vtkm/cont/cuda/testing/UnitTestCudaShareUserProvidedManagedMemory.cu
View File

@ -78,8 +78,8 @@ void TestPrepareForInput(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInput(vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForInput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForInput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -102,8 +102,8 @@ void TestPrepareForInPlace(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInPlace(vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForInPlace.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForInPlace.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -127,8 +127,8 @@ void TestPrepareForOutput(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForOutput(32, vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForOutput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForOutput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -151,12 +151,12 @@ void TestReleaseResourcesExecution(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInput(vtkm::cont::DeviceAdapterTagCuda());
ValueType* origArray = handle.Internals->ExecutionArray;
void* origArray = handle.Internals->ExecutionArray;
handle.ReleaseResourcesExecution();
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after ReleaseResourcesExecution.");
VTKM_TEST_ASSERT(execArray == nullptr,
@ -178,10 +178,10 @@ void TestRoundTrip(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForOutput(32, vtkm::cont::DeviceAdapterTagCuda());
ValueType* origContArray = handle.Internals->ControlArray.GetArray();
void* origContArray = handle.Internals->ControlArray->GetBasePointer();
{
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForOutput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForOutput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -222,8 +222,8 @@ void TestRoundTrip(bool managed)
}
{
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after GetPortalConst.");
VTKM_TEST_ASSERT(execArray == nullptr, "Execution array not cleared after GetPortalConst.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(contArray),

337
vtkm/cont/internal/ArrayHandleBasicImpl.cxx Normal file
View File

@ -0,0 +1,337 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#define vtkm_cont_internal_ArrayHandleImpl_cxx
#include <vtkm/cont/internal/ArrayHandleBasicImpl.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
#ifdef VTKM_ENABLE_TBB
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
#endif
#ifdef VTKM_ENABLE_CUDA
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
#endif
namespace vtkm
{
namespace cont
{
namespace internal
{
TypelessExecutionArray::TypelessExecutionArray(const ArrayHandleImpl* data)
: Array(data->ExecutionArray)
, ArrayEnd(data->ExecutionArrayEnd)
, ArrayCapacity(data->ExecutionArrayCapacity)
, ArrayControl(data->ControlArray->GetBasePointer())
, ArrayControlCapacity(data->ControlArray->GetCapacityPointer())
{
}
ExecutionArrayInterfaceBasicBase::ExecutionArrayInterfaceBasicBase(StorageBasicBase& storage)
: ControlStorage(storage)
{
}
ExecutionArrayInterfaceBasicBase::~ExecutionArrayInterfaceBasicBase()
{
}
ArrayHandleImpl::~ArrayHandleImpl()
{
if (this->ExecutionArrayValid && this->ExecutionInterface != nullptr &&
this->ExecutionArray != nullptr)
{
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
}
delete this->ControlArray;
delete this->ExecutionInterface;
}
void ArrayHandleImpl::CheckControlArrayValid()
{
if (!this->ControlArrayValid)
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
}
vtkm::Id ArrayHandleImpl::GetNumberOfValues(vtkm::UInt64 sizeOfT) const
{
if (this->ControlArrayValid)
{
return this->ControlArray->GetNumberOfValues();
}
else if (this->ExecutionArrayValid)
{
auto numBytes =
static_cast<char*>(this->ExecutionArrayEnd) - static_cast<char*>(this->ExecutionArray);
return static_cast<vtkm::Id>(numBytes) / static_cast<vtkm::Id>(sizeOfT);
}
else
{
return 0;
}
}
void ArrayHandleImpl::Allocate(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT)
{
this->ReleaseResourcesExecutionInternal();
this->ControlArray->AllocateValues(numberOfValues, sizeOfT);
this->ControlArrayValid = true;
}
void ArrayHandleImpl::Shrink(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT)
{
VTKM_ASSERT(numberOfValues >= 0);
if (numberOfValues > 0)
{
vtkm::Id originalNumberOfValues = this->GetNumberOfValues(sizeOfT);
if (numberOfValues < originalNumberOfValues)
{
if (this->ControlArrayValid)
{
this->ControlArray->Shrink(numberOfValues);
}
if (this->ExecutionArrayValid)
{
auto offset = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfT;
this->ExecutionArrayEnd = static_cast<char*>(this->ExecutionArray) + offset;
}
}
else if (numberOfValues == originalNumberOfValues)
{
// Nothing to do.
}
else // numberOfValues > originalNumberOfValues
{
throw vtkm::cont::ErrorBadValue("ArrayHandle::Shrink cannot be used to grow array.");
}
VTKM_ASSERT(this->GetNumberOfValues(sizeOfT) == numberOfValues);
}
else // numberOfValues == 0
{
// If we are shrinking to 0, there is nothing to save and we might as well
// free up memory. Plus, some storage classes expect that data will be
// deallocated when the size goes to zero.
this->Allocate(0, sizeOfT);
}
}
void ArrayHandleImpl::ReleaseResources()
{
this->ReleaseResourcesExecutionInternal();
if (this->ControlArrayValid)
{
this->ControlArray->ReleaseResources();
this->ControlArrayValid = false;
}
}
void ArrayHandleImpl::PrepareForInput(vtkm::UInt64 sizeOfT) const
{
const vtkm::Id numVals = this->GetNumberOfValues(sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
if (!this->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->ControlArrayValid)
{
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
this->ExecutionInterface->CopyFromControl(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
this->ExecutionInterface->UsingForRead(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
}
void ArrayHandleImpl::PrepareForOutput(vtkm::Id numVals, vtkm::UInt64 sizeOfT)
{
// Invalidate control arrays since we expect the execution data to be
// overwritten. Don't free control resources in case they're shared with
// the execution environment.
this->ControlArrayValid = false;
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
this->ExecutionInterface->UsingForWrite(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
void ArrayHandleImpl::PrepareForInPlace(vtkm::UInt64 sizeOfT)
{
const vtkm::Id numVals = this->GetNumberOfValues(sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
if (!this->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->ControlArrayValid)
{
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
this->ExecutionInterface->CopyFromControl(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
this->ExecutionInterface->UsingForReadWrite(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
// Invalidate the control array, since we expect the values to be modified:
this->ControlArrayValid = false;
}
bool ArrayHandleImpl::PrepareForDevice(DeviceAdapterId devId, vtkm::UInt64 sizeOfT) const
{
// Check if the current device matches the last one and sync through
// the control environment if the device changes.
if (this->ExecutionInterface)
{
if (this->ExecutionInterface->GetDeviceId() == devId)
{
// All set, nothing to do.
return false;
}
else
{
// Update the device allocator:
this->SyncControlArray(sizeOfT);
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
delete this->ExecutionInterface;
this->ExecutionInterface = nullptr;
this->ExecutionArrayValid = false;
}
}
VTKM_ASSERT(this->ExecutionInterface == nullptr);
VTKM_ASSERT(!this->ExecutionArrayValid);
switch (devId)
{
#ifdef VTKM_ENABLE_TBB
case VTKM_DEVICE_ADAPTER_TBB:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>(*this->ControlArray);
break;
#endif
//this doesn't need to be guarded as a .cu file as it is calling host methods
//and not cuda code directly
#ifdef VTKM_ENABLE_CUDA
case VTKM_DEVICE_ADAPTER_CUDA:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>(*this->ControlArray);
break;
#endif
default:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>(*this->ControlArray);
break;
}
return true;
}
DeviceAdapterId ArrayHandleImpl::GetDeviceAdapterId() const
{
return this->ExecutionArrayValid ? this->ExecutionInterface->GetDeviceId()
: VTKM_DEVICE_ADAPTER_UNDEFINED;
}
void ArrayHandleImpl::SyncControlArray(vtkm::UInt64 sizeOfT) const
{
if (!this->ControlArrayValid)
{
// Need to change some state that does not change the logical state from
// an external point of view.
if (this->ExecutionArrayValid)
{
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(
static_cast<char*>(this->ExecutionArrayEnd) - static_cast<char*>(this->ExecutionArray));
const vtkm::Id numVals = static_cast<vtkm::Id>(numBytes / sizeOfT);
this->ControlArray->AllocateValues(numVals, sizeOfT);
this->ExecutionInterface->CopyToControl(
this->ExecutionArray, this->ControlArray->GetBasePointer(), numBytes);
this->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
}
}
void ArrayHandleImpl::ReleaseResourcesExecutionInternal()
{
if (this->ExecutionArrayValid)
{
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
this->ExecutionArrayValid = false;
}
}
} // end namespace internal
}
} // end vtkm::cont
#ifdef VTKM_MSVC
//Export this when being used with std::shared_ptr
template class VTKM_CONT_EXPORT std::shared_ptr<vtkm::cont::internal::ArrayHandleImpl>;
#endif

157
vtkm/cont/internal/ArrayHandleBasicImpl.h
View File

@ -35,22 +35,13 @@ namespace cont
namespace internal
{
struct ArrayHandleImpl;
/// Type-agnostic container for an execution memory buffer.
struct VTKM_ALWAYS_EXPORT TypelessExecutionArray
struct VTKM_CONT_EXPORT TypelessExecutionArray
{
VTKM_CONT
TypelessExecutionArray(void*& array,
void*& arrayEnd,
void*& arrayCapacity,
const void* arrayControl,
const void* arrayControlCapacity)
: Array(array)
, ArrayEnd(arrayEnd)
, ArrayCapacity(arrayCapacity)
, ArrayControl(arrayControl)
, ArrayControlCapacity(arrayControlCapacity)
{
}
TypelessExecutionArray(const ArrayHandleImpl* data);
void*& Array;
void*& ArrayEnd;
@ -94,7 +85,9 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicBase
/// If (capacity - base) < @a numBytes, the buffer will be freed and
/// reallocated. If (capacity - base) >= numBytes, a new end is marked.
VTKM_CONT
virtual void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const = 0;
virtual void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const = 0;
/// Release the buffer held by @a execArray and reset all pointer to null.
VTKM_CONT
@ -134,13 +127,81 @@ protected:
template <typename DeviceTag>
struct ExecutionArrayInterfaceBasic;
struct VTKM_CONT_EXPORT ArrayHandleImpl
{
VTKM_CONT
template <typename T>
explicit ArrayHandleImpl(T)
: ControlArrayValid(false)
, ControlArray(new vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>())
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
VTKM_CONT
template <typename T>
explicit ArrayHandleImpl(
const vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>& storage)
: ControlArrayValid(true)
, ControlArray(new vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>(storage))
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
VTKM_CONT ~ArrayHandleImpl();
VTKM_CONT ArrayHandleImpl(const ArrayHandleImpl&) = delete;
VTKM_CONT void operator=(const ArrayHandleImpl&) = delete;
//Throws ErrorInternal if ControlArrayValid == false
VTKM_CONT void CheckControlArrayValid() noexcept(false);
VTKM_CONT vtkm::Id GetNumberOfValues(vtkm::UInt64 sizeOfT) const;
VTKM_CONT void Allocate(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT);
VTKM_CONT void Shrink(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT);
VTKM_CONT void SyncControlArray(vtkm::UInt64 sizeofT) const;
VTKM_CONT void ReleaseResources();
VTKM_CONT void ReleaseResourcesExecutionInternal();
VTKM_CONT void PrepareForInput(vtkm::UInt64 sizeofT) const;
VTKM_CONT void PrepareForOutput(vtkm::Id numVals, vtkm::UInt64 sizeofT);
VTKM_CONT void PrepareForInPlace(vtkm::UInt64 sizeofT);
// Check if the current device matches the last one. If they don't match
// this moves all data back from execution environment and deletes the
// ExecutionInterface instance.
// Returns true when the caller needs to reallocate ExecutionInterface
VTKM_CONT bool PrepareForDevice(DeviceAdapterId devId, vtkm::UInt64 sizeofT) const;
VTKM_CONT DeviceAdapterId GetDeviceAdapterId() const;
mutable bool ControlArrayValid;
StorageBasicBase* ControlArray;
mutable ExecutionArrayInterfaceBasicBase* ExecutionInterface;
mutable bool ExecutionArrayValid;
mutable void* ExecutionArray;
mutable void* ExecutionArrayEnd;
mutable void* ExecutionArrayCapacity;
};
} // end namespace internal
/// Specialization of ArrayHandle for Basic storage. The goal here is to reduce
/// the amount of codegen for the common case of Basic storage when we build
/// the common arrays into libvtkm_cont.
template <typename T>
class ArrayHandle<T, ::vtkm::cont::StorageTagBasic> : public ::vtkm::cont::internal::ArrayHandleBase
class VTKM_ALWAYS_EXPORT ArrayHandle<T, ::vtkm::cont::StorageTagBasic>
: public ::vtkm::cont::internal::ArrayHandleBase
{
private:
using Thisclass = ArrayHandle<T, ::vtkm::cont::StorageTagBasic>;
@ -154,7 +215,6 @@ public:
using ValueType = T;
using PortalControl = typename StorageType::PortalType;
using PortalConstControl = typename StorageType::PortalConstType;
struct InternalStruct;
template <typename DeviceTag>
struct ExecutionTypes
@ -167,7 +227,6 @@ public:
VTKM_CONT ArrayHandle(const Thisclass& src);
VTKM_CONT ArrayHandle(const Thisclass&& src);
VTKM_CONT ArrayHandle(const StorageType& storage);
VTKM_CONT ArrayHandle(const std::shared_ptr<InternalStruct>& i);
VTKM_CONT ~ArrayHandle();
@ -214,65 +273,19 @@ public:
VTKM_CONT void SyncControlArray() const;
VTKM_CONT void ReleaseResourcesExecutionInternal();
struct VTKM_ALWAYS_EXPORT InternalStruct
{
InternalStruct()
: ControlArrayValid(false)
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
InternalStruct(const StorageType& storage)
: ControlArrayValid(true)
, ControlArray(storage)
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
~InternalStruct()
{
if (this->ExecutionArrayValid && this->ExecutionInterface != nullptr &&
this->ExecutionArray != nullptr)
{
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->ExecutionArray),
reinterpret_cast<void*&>(this->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->ExecutionArrayCapacity),
this->ControlArray.GetBasePointer(),
this->ControlArray.GetCapacityPointer());
this->ExecutionInterface->Free(execArray);
}
delete this->ExecutionInterface;
}
InternalStruct(const InternalStruct&) = delete;
void operator=(const InternalStruct&) = delete;
bool ControlArrayValid;
StorageType ControlArray;
internal::ExecutionArrayInterfaceBasicBase* ExecutionInterface;
bool ExecutionArrayValid;
ValueType* ExecutionArray;
ValueType* ExecutionArrayEnd;
ValueType* ExecutionArrayCapacity;
};
std::shared_ptr<InternalStruct> Internals;
std::shared_ptr<internal::ArrayHandleImpl> Internals;
};
} // end namespace cont
} // end namespace vtkm
#ifndef vtkm_cont_internal_ArrayHandleImpl_cxx
#ifdef VTKM_MSVC
extern template class VTKM_CONT_TEMPLATE_EXPORT
std::shared_ptr<vtkm::cont::internal::ArrayHandleImpl>;
#endif
#endif
#include <vtkm/cont/internal/ArrayHandleBasicImpl.hxx>
#endif // vtk_m_cont_internal_ArrayHandleBasicImpl_h

325
vtkm/cont/internal/ArrayHandleBasicImpl.hxx
View File

@ -27,10 +27,9 @@ namespace vtkm
{
namespace cont
{
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle()
: Internals(new InternalStruct)
: Internals(new internal::ArrayHandleImpl(T{}))
{
}
@ -48,13 +47,7 @@ ArrayHandle<T, StorageTagBasic>::ArrayHandle(const Thisclass&& src)
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle(const StorageType& storage)
: Internals(new InternalStruct(storage))
{
}
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle(const std::shared_ptr<InternalStruct>& i)
: Internals(i)
: Internals(new internal::ArrayHandleImpl(storage))
{
}
@ -107,15 +100,11 @@ template <typename T>
typename ArrayHandle<T, StorageTagBasic>::StorageType& ArrayHandle<T, StorageTagBasic>::GetStorage()
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray;
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
return *(static_cast<StorageType*>(this->Internals->ControlArray));
}
template <typename T>
@ -123,15 +112,11 @@ const typename ArrayHandle<T, StorageTagBasic>::StorageType&
ArrayHandle<T, StorageTagBasic>::GetStorage() const
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray;
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
return *(static_cast<const StorageType*>(this->Internals->ControlArray));
}
template <typename T>
@ -139,19 +124,17 @@ typename ArrayHandle<T, StorageTagBasic>::PortalControl
ArrayHandle<T, StorageTagBasic>::GetPortalControl()
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
// If the user writes into the iterator we return, then the execution
// array will become invalid. Play it safe and release the execution
// resources. (Use the const version to preserve the execution array.)
this->ReleaseResourcesExecutionInternal();
return this->Internals->ControlArray.GetPortal();
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
// If the user writes into the iterator we return, then the execution
// array will become invalid. Play it safe and release the execution
// resources. (Use the const version to preserve the execution array.)
this->ReleaseResourcesExecutionInternal();
StorageType* privStorage = static_cast<StorageType*>(this->Internals->ControlArray);
return privStorage->GetPortal();
}
@ -160,81 +143,30 @@ typename ArrayHandle<T, StorageTagBasic>::PortalConstControl
ArrayHandle<T, StorageTagBasic>::GetPortalConstControl() const
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray.GetPortalConst();
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
StorageType* privStorage = static_cast<StorageType*>(this->Internals->ControlArray);
return privStorage->GetPortalConst();
}
template <typename T>
vtkm::Id ArrayHandle<T, StorageTagBasic>::GetNumberOfValues() const
{
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray.GetNumberOfValues();
}
else if (this->Internals->ExecutionArrayValid)
{
return static_cast<vtkm::Id>(this->Internals->ExecutionArrayEnd -
this->Internals->ExecutionArray);
}
else
{
return 0;
}
return this->Internals->GetNumberOfValues(sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::Allocate(vtkm::Id numberOfValues)
{
this->ReleaseResourcesExecutionInternal();
this->Internals->ControlArray.Allocate(numberOfValues);
this->Internals->ControlArrayValid = true;
this->Internals->Allocate(numberOfValues, sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::Shrink(vtkm::Id numberOfValues)
{
VTKM_ASSERT(numberOfValues >= 0);
if (numberOfValues > 0)
{
vtkm::Id originalNumberOfValues = this->GetNumberOfValues();
if (numberOfValues < originalNumberOfValues)
{
if (this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Shrink(numberOfValues);
}
if (this->Internals->ExecutionArrayValid)
{
this->Internals->ExecutionArrayEnd = this->Internals->ExecutionArray + numberOfValues;
}
}
else if (numberOfValues == originalNumberOfValues)
{
// Nothing to do.
}
else // numberOfValues > originalNumberOfValues
{
throw vtkm::cont::ErrorBadValue("ArrayHandle::Shrink cannot be used to grow array.");
}
VTKM_ASSERT(this->GetNumberOfValues() == numberOfValues);
}
else // numberOfValues == 0
{
// If we are shrinking to 0, there is nothing to save and we might as well
// free up memory. Plus, some storage classes expect that data will be
// deallocated when the size goes to zero.
this->Allocate(0);
}
this->Internals->Shrink(numberOfValues, sizeof(T));
}
template <typename T>
@ -243,19 +175,13 @@ void ArrayHandle<T, StorageTagBasic>::ReleaseResourcesExecution()
// Save any data in the execution environment by making sure it is synced
// with the control environment.
this->SyncControlArray();
this->ReleaseResourcesExecutionInternal();
this->Internals->ReleaseResourcesExecutionInternal();
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::ReleaseResources()
{
this->ReleaseResourcesExecutionInternal();
if (this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.ReleaseResources();
this->Internals->ControlArrayValid = false;
}
this->Internals->ReleaseResources();
}
template <typename T>
@ -264,41 +190,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForInput(DeviceAdapterTag device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
}
this->Internals->ExecutionInterface->UsingForRead(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
return PortalFactory<DeviceAdapterTag>::CreatePortalConst(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForInput(sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortalConst(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -307,34 +204,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForOutput(vtkm::Id numVals, DeviceAdapterTag device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
// Invalidate control arrays since we expect the execution data to be
// overwritten. Don't free control resources in case they're shared with
// the execution environment.
this->Internals->ControlArrayValid = false;
internal::TypelessExecutionArray execArray(reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
const vtkm::UInt64 numBytes =
static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numVals);
this->Internals->ExecutionInterface->Allocate(execArray, numBytes);
this->Internals->ExecutionInterface->UsingForWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
return PortalFactory<DeviceAdapterTag>::CreatePortal(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForOutput(numVals, sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortal(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -343,45 +218,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForInPlace(DeviceAdapterTag device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetStorage().GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->Internals->ExecutionArray),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
}
priv->ExecutionInterface->UsingForReadWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
// Invalidate the control array, since we expect the values to be modified:
this->Internals->ControlArrayValid = false;
return PortalFactory<DeviceAdapterTag>::CreatePortal(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForInPlace(sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortal(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -389,92 +231,25 @@ template <typename DeviceAdapterTag>
void ArrayHandle<T, StorageTagBasic>::PrepareForDevice(DeviceAdapterTag) const
{
DeviceAdapterId devId = DeviceAdapterTraits<DeviceAdapterTag>::GetId();
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
// Check if the current device matches the last one and sync through
// the control environment if the device changes.
if (this->Internals->ExecutionInterface)
{
if (this->Internals->ExecutionInterface->GetDeviceId() == devId)
{
// All set, nothing to do.
return;
}
else
{
// Update the device allocator:
this->SyncControlArray();
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Free(execArray);
delete priv->ExecutionInterface;
priv->ExecutionInterface = nullptr;
priv->ExecutionArrayValid = false;
}
}
VTKM_ASSERT(priv->ExecutionInterface == nullptr);
VTKM_ASSERT(!priv->ExecutionArrayValid);
priv->ExecutionInterface =
new internal::ExecutionArrayInterfaceBasic<DeviceAdapterTag>(this->Internals->ControlArray);
this->Internals->PrepareForDevice(devId, sizeof(T));
}
template <typename T>
DeviceAdapterId ArrayHandle<T, StorageTagBasic>::GetDeviceAdapterId() const
{
return this->Internals->ExecutionArrayValid ? this->Internals->ExecutionInterface->GetDeviceId()
: VTKM_DEVICE_ADAPTER_UNDEFINED;
return this->Internals->GetDeviceAdapterId();
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::SyncControlArray() const
{
if (!this->Internals->ControlArrayValid)
{
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
if (this->Internals->ExecutionArrayValid)
{
const vtkm::Id numValues =
static_cast<vtkm::Id>(this->Internals->ExecutionArrayEnd - this->Internals->ExecutionArray);
const vtkm::UInt64 numBytes =
static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numValues);
priv->ControlArray.Allocate(numValues);
priv->ExecutionInterface->CopyToControl(
priv->ExecutionArray, priv->ControlArray.GetArray(), numBytes);
priv->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
priv->ControlArray.Allocate(0);
priv->ControlArrayValid = true;
}
}
this->Internals->SyncControlArray(sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::ReleaseResourcesExecutionInternal()
{
if (this->Internals->ExecutionArrayValid)
{
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->Internals->ExecutionArray),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
this->Internals->ExecutionInterface->Free(execArray);
this->Internals->ExecutionArrayValid = false;
}
this->Internals->ReleaseResourcesExecutionInternal();
}
}
} // end namespace vtkm::cont

2
vtkm/cont/internal/ArrayManagerExecution.h
View File

@ -40,7 +40,7 @@ namespace internal
/// any resources in its destructor.
///
/// This class typically takes on one of two forms. If the control and
/// execution environments have seperate memory spaces, then this class
/// execution environments have separate memory spaces, then this class
/// behaves how you would expect. It allocates/deallocates arrays and copies
/// data. However, if the control and execution environments share the same
/// memory space, this class should delegate all its operations to the

8
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx
View File

@ -34,12 +34,12 @@ ExecutionArrayInterfaceBasicShareWithControl::ExecutionArrayInterfaceBasicShareW
}
void ExecutionArrayInterfaceBasicShareWithControl::Allocate(TypelessExecutionArray& execArray,
vtkm::UInt64 numBytes) const
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const
{
this->ControlStorage.AllocateBytes(numBytes);
this->ControlStorage.AllocateValues(numberOfValues, sizeOfValue);
execArray.Array = this->ControlStorage.GetBasePointer();
execArray.ArrayEnd = this->ControlStorage.GetEndPointer();
execArray.ArrayEnd = this->ControlStorage.GetEndPointer(numberOfValues, sizeOfValue);
execArray.ArrayCapacity = this->ControlStorage.GetCapacityPointer();
}

4
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
View File

@ -139,7 +139,9 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicShareWithControl
VTKM_CONT ExecutionArrayInterfaceBasicShareWithControl(StorageBasicBase& storage);
VTKM_CONT void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* src, void* dst, vtkm::UInt64 bytes) const final;

33
vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
View File

@ -763,6 +763,37 @@ public:
DerivedAlgorithm::Sort(zipHandle, internal::KeyCompare<T, U, BinaryCompare>(binary_compare));
}
template <typename T,
typename U,
typename V,
typename StorageT,
typename StorageU,
typename StorageV,
typename BinaryFunctor>
VTKM_CONT static void Transform(const vtkm::cont::ArrayHandle<T, StorageT>& input1,
const vtkm::cont::ArrayHandle<U, StorageU>& input2,
vtkm::cont::ArrayHandle<V, StorageV>& output,
BinaryFunctor binaryFunctor)
{
vtkm::Id numValues = vtkm::Min(input1.GetNumberOfValues(), input2.GetNumberOfValues());
if (numValues <= 0)
{
return;
}
auto input1Portal = input1.PrepareForInput(DeviceAdapterTag());
auto input2Portal = input2.PrepareForInput(DeviceAdapterTag());
auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag());
BinaryTransformKernel<decltype(input1Portal),
decltype(input2Portal),
decltype(outputPortal),
BinaryFunctor>
binaryKernel(input1Portal, input2Portal, outputPortal, binaryFunctor);
DerivedAlgorithm::Schedule(binaryKernel, numValues);
}
//};
//--------------------------------------------------------------------------
// Unique
template <typename T, class Storage>
@ -966,7 +997,7 @@ private:
/// \brief Class providing a device-specific support for selecting the optimal
/// Task type for a given worklet.
///
/// When worklets are launched inside the execution enviornment we need to
/// When worklets are launched inside the execution environment we need to
/// ask the device adapter what is the preferred execution style, be it
/// a tiled iteration pattern, or strided. This class
///

32
vtkm/cont/internal/FunctorsGeneral.h
View File

@ -881,6 +881,38 @@ struct ScanKernel : vtkm::exec::FunctorBase
}
}
};
template <typename InPortalType1,
typename InPortalType2,
typename OutPortalType,
typename BinaryFunctor>
struct BinaryTransformKernel : vtkm::exec::FunctorBase
{
InPortalType1 InPortal1;
InPortalType2 InPortal2;
OutPortalType OutPortal;
BinaryFunctor BinaryOperator;
VTKM_CONT
BinaryTransformKernel(const InPortalType1& inPortal1,
const InPortalType2& inPortal2,
const OutPortalType& outPortal,
BinaryFunctor binaryOperator)
: InPortal1(inPortal1)
, InPortal2(inPortal2)
, OutPortal(outPortal)
, BinaryOperator(binaryOperator)
{
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC
void operator()(vtkm::Id index) const
{
this->OutPortal.Set(
index, this->BinaryOperator(this->InPortal1.Get(index), this->InPortal2.Get(index)));
}
};
}
}
} // namespace vtkm::cont::internal

2
vtkm/cont/internal/testing/UnitTestArrayPortalFromIterators.cxx
View File

@ -119,7 +119,7 @@ struct TemplatedTests
VTKM_TEST_ASSERT(const_portal.GetNumberOfValues() == ARRAY_SIZE,
"Const portal array size wrong.");
std::cout << " Check inital value." << std::endl;
std::cout << " Check initial value." << std::endl;
VTKM_TEST_ASSERT(CheckPortal(portal, ORIGINAL_VALUE), "Portal iterator has bad value.");
VTKM_TEST_ASSERT(CheckPortal(const_portal, ORIGINAL_VALUE),
"Const portal iterator has bad value.");

14
vtkm/cont/serial/internal/ArrayManagerExecutionSerial.h
View File

@ -24,6 +24,7 @@
#include <vtkm/cont/internal/ArrayManagerExecution.h>
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
namespace vtkm
{
@ -62,19 +63,6 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagSerial>
using Superclass::CreatePortalConst;
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return VTKM_DEVICE_ADAPTER_SERIAL; }
};
} // namespace internal
#ifndef vtk_m_cont_serial_internal_ArrayManagerExecutionSerial_cxx

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

@ -22,6 +22,7 @@ set(headers
ArrayManagerExecutionSerial.h
DeviceAdapterAlgorithmSerial.h
DeviceAdapterTagSerial.h
ExecutionArrayInterfaceBasicSerial.h
VirtualObjectTransferSerial.h
)
vtkm_declare_headers(${headers})
@ -29,4 +30,5 @@ vtkm_declare_headers(${headers})
target_sources(vtkm_cont PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/ArrayManagerExecutionSerial.cxx
${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterAlgorithmSerial.cxx
${CMAKE_CURRENT_SOURCE_DIR}/ExecutionArrayInterfaceBasicSerial.cxx
)

29
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h
View File

@ -281,28 +281,6 @@ public:
values_output.Shrink(writePos + 1);
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
vtkm::Id numberOfValues = input.GetNumberOfValues();
auto inputPortal = input.PrepareForInput(Device());
auto outputPortal = output.PrepareForOutput(numberOfValues, Device());
if (numberOfValues <= 0)
{
return vtkm::TypeTraits<T>::ZeroInitialization();
}
std::partial_sum(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal),
vtkm::cont::ArrayPortalToIteratorEnd(inputPortal),
vtkm::cont::ArrayPortalToIteratorBegin(outputPortal));
// Return the value at the last index in the array, which is the full sum.
return outputPortal.Get(numberOfValues - 1);
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,
@ -329,6 +307,13 @@ public:
return outputPortal.Get(numberOfValues - 1);
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
return ScanInclusive(input, output, vtkm::Sum());
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,

36
vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.cxx Normal file
View File

@ -0,0 +1,36 @@
//============================================================================
// 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 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
vtkm::cont::DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>::GetDeviceId()
const
{
return VTKM_DEVICE_ADAPTER_SERIAL;
}
} // namespace internal
}
} // namespace vtkm::cont

50
vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h Normal file
View File

@ -0,0 +1,50 @@
//============================================================================
// 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 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 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_serial_internal_ExecutionArrayInterfaceBasicSerial_h
#define vtk_m_cont_serial_internal_ExecutionArrayInterfaceBasicSerial_h
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
DeviceAdapterId GetDeviceId() const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_serial_internal_ExecutionArrayInterfaceBasicSerial_h

18
vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h
View File

@ -20,12 +20,13 @@
#ifndef vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_h
#define vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_h
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/internal/ArrayExportMacros.h>
#include <vtkm/cont/internal/ArrayManagerExecution.h>
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
@ -86,19 +87,6 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagTBB>
using Superclass::CreatePortalConst;
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return VTKM_DEVICE_ADAPTER_TBB; }
};
} // namespace internal
#ifndef vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_cxx
VTKM_EXPORT_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagTBB)

8
vtkm/cont/tbb/internal/CMakeLists.txt
View File

@ -22,7 +22,9 @@ set(headers
ArrayManagerExecutionTBB.h
DeviceAdapterAlgorithmTBB.h
DeviceAdapterTagTBB.h
ExecutionArrayInterfaceBasicTBB.h
FunctorsTBB.h
ParallelSortTBB.h
VirtualObjectTransferTBB.h
)
@ -34,7 +36,9 @@ if (VTKm_ENABLE_TBB)
endif()
endif()
vtkm_declare_headers(parallel_sort.h TESTABLE OFF)
vtkm_declare_headers(ParallelSortTBB.hxx
TESTABLE OFF)
vtkm_declare_headers(${headers} TESTABLE ${VTKm_ENABLE_TBB})
#-----------------------------------------------------------------------------
@ -45,4 +49,6 @@ endif()
target_sources(vtkm_cont PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/ArrayManagerExecutionTBB.cxx
${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterAlgorithmTBB.cxx
${CMAKE_CURRENT_SOURCE_DIR}/ExecutionArrayInterfaceBasicTBB.cxx
${CMAKE_CURRENT_SOURCE_DIR}/ParallelSortTBB.cxx
)

57
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
View File

@ -30,6 +30,7 @@
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/tbb/internal/FunctorsTBB.h>
#include <vtkm/cont/tbb/internal/ParallelSortTBB.h>
#include <vtkm/exec/tbb/internal/TaskTiling.h>
@ -242,75 +243,37 @@ public:
ScheduleTask(kernel, rangeMax);
}
//1. We need functions for each of the following
template <typename T, class Container>
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Container>& values)
{
//this is required to get sort to work with zip handles
std::less<T> lessOp;
Sort(values, lessOp);
vtkm::cont::tbb::internal::parallel_sort(values, lessOp);
}
template <typename T, class Container, class BinaryCompare>
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Container>& values,
BinaryCompare binary_compare)
{
using PortalType = typename vtkm::cont::ArrayHandle<T, Container>::template ExecutionTypes<
vtkm::cont::DeviceAdapterTagTBB>::Portal;
PortalType arrayPortal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagTBB());
using IteratorsType = vtkm::cont::ArrayPortalToIterators<PortalType>;
IteratorsType iterators(arrayPortal);
internal::WrappedBinaryOperator<bool, BinaryCompare> wrappedCompare(binary_compare);
::tbb::parallel_sort(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare);
vtkm::cont::tbb::internal::parallel_sort(values, binary_compare);
}
template <typename T, typename U, class StorageT, class StorageU>
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values)
{
SortByKey(keys, values, std::less<T>());
vtkm::cont::tbb::internal::parallel_sort_bykey(keys, values, std::less<T>());
}
template <typename T, typename U, class StorageT, class StorageU, class Compare>
template <typename T, typename U, class StorageT, class StorageU, class BinaryCompare>
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
Compare comp)
BinaryCompare binary_compare)
{
using KeyType = vtkm::cont::ArrayHandle<T, StorageT>;
VTKM_CONSTEXPR bool larger_than_64bits = sizeof(U) > sizeof(vtkm::Int64);
if (larger_than_64bits)
{
/// More efficient sort:
/// Move value indexes when sorting and reorder the value array at last
using ValueType = vtkm::cont::ArrayHandle<U, StorageU>;
using IndexType = vtkm::cont::ArrayHandle<vtkm::Id>;
using ZipHandleType = vtkm::cont::ArrayHandleZip<KeyType, IndexType>;
IndexType indexArray;
ValueType valuesScattered;
const vtkm::Id size = values.GetNumberOfValues();
Copy(ArrayHandleIndex(keys.GetNumberOfValues()), indexArray);
ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, indexArray);
Sort(zipHandle, vtkm::cont::internal::KeyCompare<T, vtkm::Id, Compare>(comp));
tbb::ScatterPortal(values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB()));
Copy(valuesScattered, values);
}
else
{
using ValueType = vtkm::cont::ArrayHandle<U, StorageU>;
using ZipHandleType = vtkm::cont::ArrayHandleZip<KeyType, ValueType>;
ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values);
Sort(zipHandle, vtkm::cont::internal::KeyCompare<T, U, Compare>(comp));
}
vtkm::cont::tbb::internal::parallel_sort_bykey(keys, values, binary_compare);
}
template <typename T, class Storage>

35
vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.cxx Normal file
View File

@ -0,0 +1,35 @@
//============================================================================
// 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 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
vtkm::cont::DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_TBB;
}
} // namespace internal
}
} // namespace vtkm::cont

53
vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h Normal file
View File

@ -0,0 +1,53 @@
//============================================================================
// 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 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 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_tbb_internal_ExecutionArrayInterfaceBasicTBB_h
#define vtk_m_cont_tbb_internal_ExecutionArrayInterfaceBasicTBB_h
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
DeviceAdapterId GetDeviceId() const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_tbb_internal_ExecutionArrayInterfaceBasicTBB_h

958
vtkm/cont/tbb/internal/ParallelSortTBB.cxx Normal file
View File

@ -0,0 +1,958 @@
//=============================================================================
//
// 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 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2018 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.
//
//=============================================================================
// Copyright 2010, Takuya Akiba
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are
// met:
//
// * Redistributions of source code must retain the above copyright
// notice, this list of conditions and the following disclaimer.
// * Redistributions in binary form must reproduce the above
// copyright notice, this list of conditions and the following disclaimer
// in the documentation and/or other materials provided with the
// distribution.
// * Neither the name of Takuya Akiba nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
// "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
// LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
// A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
// OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
// SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
// LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
// DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
// THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
// (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
// Modifications of Takuya Akiba's original GitHub source code for inclusion
// in VTK-m:
//
// - Changed parallel threading from OpenMP to TBB tasks
// - Added minimum threshold for parallel, will instead invoke serial radix sort (kxsort)
// - Added std::greater<T> and std::less<T> to interface for descending order sorts
// - Added linear scaling of threads used by the algorithm for more stable performance
// on machines with lots of available threads (KNL and Haswell)
//
// This file contains an implementation of Satish parallel radix sort
// as documented in the following citation:
//
// Fast sort on CPUs and GPUs: a case for bandwidth oblivious SIMD sort.
// N. Satish, C. Kim, J. Chhugani, A. D. Nguyen, V. W. Lee, D. Kim, and P. Dubey.
// In Proc. SIGMOD, pages 351–362, 2010
#include <algorithm>
#include <cassert>
#include <climits>
#include <cmath>
#include <cstring>
#include <functional>
#include <stdint.h>
#include <utility>
#include <vtkm/Types.h>
#include <vtkm/cont/tbb/internal/ParallelSortTBB.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <vtkm/cont/tbb/internal/kxsort.h>
#if defined(VTKM_MSVC)
// TBB's header include a #pragma comment(lib,"tbb.lib") line to make all
// consuming libraries link to tbb, this is bad behavior in a header
// based project
#pragma push_macro("__TBB_NO_IMPLICITLINKAGE")
#define __TBB_NO_IMPLICIT_LINKAGE 1
#endif // defined(VTKM_MSVC)
// TBB includes windows.h, so instead we want to include windows.h with the
// correct settings so that we don't clobber any existing function
#include <vtkm/internal/Windows.h>
#include <tbb/task.h>
#include <tbb/tbb_thread.h>
#if defined(VTKM_MSVC)
#pragma pop_macro("__TBB_NO_IMPLICITLINKAGE")
#endif
namespace vtkm
{
namespace cont
{
namespace tbb
{
namespace internal
{
const size_t MIN_BYTES_FOR_PARALLEL = 400000;
const size_t BYTES_FOR_MAX_PARALLELISM = 4000000;
const size_t MAX_CORES = ::tbb::tbb_thread::hardware_concurrency();
const double CORES_PER_BYTE =
double(MAX_CORES - 1) / double(BYTES_FOR_MAX_PARALLELISM - MIN_BYTES_FOR_PARALLEL);
const double Y_INTERCEPT = 1.0 - CORES_PER_BYTE * MIN_BYTES_FOR_PARALLEL;
namespace utility
{
// Return the number of threads that would be executed in parallel regions
inline size_t GetMaxThreads(size_t num_bytes)
{
size_t num_cores = (size_t)(CORES_PER_BYTE * double(num_bytes) + Y_INTERCEPT);
if (num_cores < 1)
{
return 1;
}
if (num_cores > MAX_CORES)
{
return MAX_CORES;
}
return num_cores;
}
} // namespace utility
namespace internal
{
// Size of the software managed buffer
const size_t kOutBufferSize = 32;
// Ascending order radix sort is a no-op
template <typename PlainType,
typename UnsignedType,
typename CompareType,
typename ValueManager,
unsigned int Base>
struct ParallelRadixCompareInternal
{
inline static void reverse(UnsignedType& t) { (void)t; }
};
// Handle descending order radix sort
template <typename PlainType, typename UnsignedType, typename ValueManager, unsigned int Base>
struct ParallelRadixCompareInternal<PlainType,
UnsignedType,
std::greater<PlainType>,
ValueManager,
Base>
{
inline static void reverse(UnsignedType& t) { t = ((1 << Base) - 1) - t; }
};
// The algorithm is implemented in this internal class
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
class ParallelRadixSortInternal
{
public:
using CompareInternal =
ParallelRadixCompareInternal<PlainType, UnsignedType, CompareType, ValueManager, Base>;
ParallelRadixSortInternal();
~ParallelRadixSortInternal();
void Init(PlainType* data, size_t num_elems);
PlainType* Sort(PlainType* data, ValueManager* value_manager);
static void InitAndSort(PlainType* data, size_t num_elems, ValueManager* value_manager);
private:
CompareInternal compare_internal_;
size_t num_elems_;
size_t num_threads_;
UnsignedType* tmp_;
size_t** histo_;
UnsignedType*** out_buf_;
size_t** out_buf_n_;
size_t *pos_bgn_, *pos_end_;
ValueManager* value_manager_;
void DeleteAll();
UnsignedType* SortInternal(UnsignedType* data, ValueManager* value_manager);
// Compute |pos_bgn_| and |pos_end_| (associated ranges for each threads)
void ComputeRanges();
// First step of each iteration of sorting
// Compute the histogram of |src| using bits in [b, b + Base)
void ComputeHistogram(unsigned int b, UnsignedType* src);
// Second step of each iteration of sorting
// Scatter elements of |src| to |dst| using the histogram
void Scatter(unsigned int b, UnsignedType* src, UnsignedType* dst);
};
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
ParallelRadixSortInternal()
: num_elems_(0)
, num_threads_(0)
, tmp_(NULL)
, histo_(NULL)
, out_buf_(NULL)
, out_buf_n_(NULL)
, pos_bgn_(NULL)
, pos_end_(NULL)
{
assert(sizeof(PlainType) == sizeof(UnsignedType));
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
~ParallelRadixSortInternal()
{
DeleteAll();
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
DeleteAll()
{
delete[] tmp_;
tmp_ = NULL;
for (size_t i = 0; i < num_threads_; ++i)
delete[] histo_[i];
delete[] histo_;
histo_ = NULL;
for (size_t i = 0; i < num_threads_; ++i)
{
for (size_t j = 0; j < 1 << Base; ++j)
{
delete[] out_buf_[i][j];
}
delete[] out_buf_n_[i];
delete[] out_buf_[i];
}
delete[] out_buf_;
delete[] out_buf_n_;
out_buf_ = NULL;
out_buf_n_ = NULL;
delete[] pos_bgn_;
delete[] pos_end_;
pos_bgn_ = pos_end_ = NULL;
num_elems_ = 0;
num_threads_ = 0;
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
Init(PlainType* data, size_t num_elems)
{
(void)data;
DeleteAll();
num_elems_ = num_elems;
num_threads_ = utility::GetMaxThreads(num_elems_ * sizeof(PlainType));
tmp_ = new UnsignedType[num_elems_];
histo_ = new size_t*[num_threads_];
for (size_t i = 0; i < num_threads_; ++i)
{
histo_[i] = new size_t[1 << Base];
}
out_buf_ = new UnsignedType**[num_threads_];
out_buf_n_ = new size_t*[num_threads_];
for (size_t i = 0; i < num_threads_; ++i)
{
out_buf_[i] = new UnsignedType*[1 << Base];
out_buf_n_[i] = new size_t[1 << Base];
for (size_t j = 0; j < 1 << Base; ++j)
{
out_buf_[i][j] = new UnsignedType[kOutBufferSize];
}
}
pos_bgn_ = new size_t[num_threads_];
pos_end_ = new size_t[num_threads_];
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
PlainType*
ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::Sort(
PlainType* data,
ValueManager* value_manager)
{
UnsignedType* src = reinterpret_cast<UnsignedType*>(data);
UnsignedType* res = SortInternal(src, value_manager);
return reinterpret_cast<PlainType*>(res);
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
InitAndSort(PlainType* data, size_t num_elems, ValueManager* value_manager)
{
ParallelRadixSortInternal prs;
prs.Init(data, num_elems);
const PlainType* res = prs.Sort(data, value_manager);
if (res != data)
{
for (size_t i = 0; i < num_elems; ++i)
data[i] = res[i];
}
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
UnsignedType*
ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
SortInternal(UnsignedType* data, ValueManager* value_manager)
{
value_manager_ = value_manager;
// Compute |pos_bgn_| and |pos_end_|
ComputeRanges();
// Iterate from lower bits to higher bits
const size_t bits = CHAR_BIT * sizeof(UnsignedType);
UnsignedType *src = data, *dst = tmp_;
for (unsigned int b = 0; b < bits; b += Base)
{
ComputeHistogram(b, src);
Scatter(b, src, dst);
std::swap(src, dst);
value_manager->Next();
}
return src;
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
ComputeRanges()
{
pos_bgn_[0] = 0;
for (size_t i = 0; i < num_threads_ - 1; ++i)
{
const size_t t = (num_elems_ - pos_bgn_[i]) / (num_threads_ - i);
pos_bgn_[i + 1] = pos_end_[i] = pos_bgn_[i] + t;
}
pos_end_[num_threads_ - 1] = num_elems_;
}
template <typename PlainType,
typename UnsignedType,
typename Encoder,
unsigned int Base,
typename Function>
class RunTask : public ::tbb::task
{
public:
RunTask(size_t binary_tree_height,
size_t binary_tree_position,
Function f,
size_t num_elems,
size_t num_threads)
: binary_tree_height_(binary_tree_height)
, binary_tree_position_(binary_tree_position)
, f_(f)
, num_elems_(num_elems)
, num_threads_(num_threads)
{
}
::tbb::task* execute()
{
size_t num_nodes_at_current_height = (size_t)pow(2, (double)binary_tree_height_);
if (num_threads_ <= num_nodes_at_current_height)
{
const size_t my_id = binary_tree_position_ - num_nodes_at_current_height;
if (my_id < num_threads_)
{
f_(my_id);
}
return NULL;
}
else
{
::tbb::empty_task& p = *new (task::allocate_continuation())::tbb::empty_task();
RunTask& left = *new (p.allocate_child()) RunTask(
binary_tree_height_ + 1, 2 * binary_tree_position_, f_, num_elems_, num_threads_);
RunTask& right = *new (p.allocate_child()) RunTask(
binary_tree_height_ + 1, 2 * binary_tree_position_ + 1, f_, num_elems_, num_threads_);
p.set_ref_count(2);
task::spawn(left);
task::spawn(right);
return NULL;
}
}
private:
size_t binary_tree_height_;
size_t binary_tree_position_;
Function f_;
size_t num_elems_;
size_t num_threads_;
};
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
ComputeHistogram(unsigned int b, UnsignedType* src)
{
// Compute local histogram
auto lambda = [=](const size_t my_id) {
const size_t my_bgn = pos_bgn_[my_id];
const size_t my_end = pos_end_[my_id];
size_t* my_histo = histo_[my_id];
memset(my_histo, 0, sizeof(size_t) * (1 << Base));
for (size_t i = my_bgn; i < my_end; ++i)
{
const UnsignedType s = Encoder::encode(src[i]);
UnsignedType t = (s >> b) & ((1 << Base) - 1);
compare_internal_.reverse(t);
++my_histo[t];
}
};
typedef RunTask<PlainType, UnsignedType, Encoder, Base, std::function<void(size_t)>> RunTaskType;
RunTaskType& root =
*new (::tbb::task::allocate_root()) RunTaskType(0, 1, lambda, num_elems_, num_threads_);
::tbb::task::spawn_root_and_wait(root);
// Compute global histogram
size_t s = 0;
for (size_t i = 0; i < 1 << Base; ++i)
{
for (size_t j = 0; j < num_threads_; ++j)
{
const size_t t = s + histo_[j][i];
histo_[j][i] = s;
s = t;
}
}
}
template <typename PlainType,
typename CompareType,
typename UnsignedType,
typename Encoder,
typename ValueManager,
unsigned int Base>
void ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>::
Scatter(unsigned int b, UnsignedType* src, UnsignedType* dst)
{
auto lambda = [=](const size_t my_id) {
const size_t my_bgn = pos_bgn_[my_id];
const size_t my_end = pos_end_[my_id];
size_t* my_histo = histo_[my_id];
UnsignedType** my_buf = out_buf_[my_id];
size_t* my_buf_n = out_buf_n_[my_id];
memset(my_buf_n, 0, sizeof(size_t) * (1 << Base));
for (size_t i = my_bgn; i < my_end; ++i)
{
const UnsignedType s = Encoder::encode(src[i]);
UnsignedType t = (s >> b) & ((1 << Base) - 1);
compare_internal_.reverse(t);
my_buf[t][my_buf_n[t]] = src[i];
value_manager_->Push(my_id, t, my_buf_n[t], i);
++my_buf_n[t];
if (my_buf_n[t] == kOutBufferSize)
{
size_t p = my_histo[t];
for (size_t j = 0; j < kOutBufferSize; ++j)
{
size_t tp = p++;
dst[tp] = my_buf[t][j];
}
value_manager_->Flush(my_id, t, kOutBufferSize, my_histo[t]);
my_histo[t] += kOutBufferSize;
my_buf_n[t] = 0;
}
}
// Flush everything
for (size_t i = 0; i < 1 << Base; ++i)
{
size_t p = my_histo[i];
for (size_t j = 0; j < my_buf_n[i]; ++j)
{
size_t tp = p++;
dst[tp] = my_buf[i][j];
}
value_manager_->Flush(my_id, i, my_buf_n[i], my_histo[i]);
}
};
typedef RunTask<PlainType, UnsignedType, Encoder, Base, std::function<void(size_t)>> RunTaskType;
RunTaskType& root =
*new (::tbb::task::allocate_root()) RunTaskType(0, 1, lambda, num_elems_, num_threads_);
::tbb::task::spawn_root_and_wait(root);
}
} // namespace internal
// Encoders encode signed/unsigned integers and floating point numbers
// to correctly ordered unsigned integers
namespace encoder
{
class EncoderDummy
{
};
class EncoderUnsigned
{
public:
template <typename UnsignedType>
inline static UnsignedType encode(UnsignedType x)
{
return x;
}
};
class EncoderSigned
{
public:
template <typename UnsignedType>
inline static UnsignedType encode(UnsignedType x)
{
return x ^ (UnsignedType(1) << (CHAR_BIT * sizeof(UnsignedType) - 1));
}
};
class EncoderDecimal
{
public:
template <typename UnsignedType>
inline static UnsignedType encode(UnsignedType x)
{
static const size_t bits = CHAR_BIT * sizeof(UnsignedType);
const UnsignedType a = x >> (bits - 1);
const UnsignedType b = (-static_cast<int>(a)) | (UnsignedType(1) << (bits - 1));
return x ^ b;
}
};
} // namespace encoder
// Value managers are used to generalize the sorting algorithm
// to sorting of keys and sorting of pairs
namespace value_manager
{
class DummyValueManager
{
public:
inline void Push(int thread, size_t bucket, size_t num, size_t from_pos)
{
(void)thread;
(void)bucket;
(void)num;
(void)from_pos;
}
inline void Flush(int thread, size_t bucket, size_t num, size_t to_pos)
{
(void)thread;
(void)bucket;
(void)num;
(void)to_pos;
}
void Next() {}
};
template <typename PlainType, typename ValueType, int Base>
class PairValueManager
{
public:
PairValueManager()
: max_elems_(0)
, max_threads_(0)
, original_(NULL)
, tmp_(NULL)
, src_(NULL)
, dst_(NULL)
, out_buf_(NULL)
{
}
~PairValueManager() { DeleteAll(); }
void Init(size_t max_elems);
void Start(ValueType* original, size_t num_elems)
{
assert(num_elems <= max_elems_);
src_ = original_ = original;
dst_ = tmp_;
}
inline void Push(int thread, size_t bucket, size_t num, size_t from_pos)
{
out_buf_[thread][bucket][num] = src_[from_pos];
}
inline void Flush(int thread, size_t bucket, size_t num, size_t to_pos)
{
for (size_t i = 0; i < num; ++i)
{
dst_[to_pos++] = out_buf_[thread][bucket][i];
}
}
void Next() { std::swap(src_, dst_); }
ValueType* GetResult() { return src_; }
private:
size_t max_elems_;
int max_threads_;
static const size_t kOutBufferSize = internal::kOutBufferSize;
ValueType *original_, *tmp_;
ValueType *src_, *dst_;
ValueType*** out_buf_;
void DeleteAll();
};
template <typename PlainType, typename ValueType, int Base>
void PairValueManager<PlainType, ValueType, Base>::Init(size_t max_elems)
{
DeleteAll();
max_elems_ = max_elems;
max_threads_ = utility::GetMaxThreads(max_elems_ * sizeof(PlainType));
tmp_ = new ValueType[max_elems];
out_buf_ = new ValueType**[max_threads_];
for (int i = 0; i < max_threads_; ++i)
{
out_buf_[i] = new ValueType*[1 << Base];
for (size_t j = 0; j < 1 << Base; ++j)
{
out_buf_[i][j] = new ValueType[kOutBufferSize];
}
}
}
template <typename PlainType, typename ValueType, int Base>
void PairValueManager<PlainType, ValueType, Base>::DeleteAll()
{
delete[] tmp_;
tmp_ = NULL;
for (int i = 0; i < max_threads_; ++i)
{
for (size_t j = 0; j < 1 << Base; ++j)
{
delete[] out_buf_[i][j];
}
delete[] out_buf_[i];
}
delete[] out_buf_;
out_buf_ = NULL;
max_elems_ = 0;
max_threads_ = 0;
}
} // namespace value_manager
// Frontend class for sorting keys
template <typename PlainType,
typename CompareType,
typename UnsignedType = PlainType,
typename Encoder = encoder::EncoderDummy,
unsigned int Base = 8>
class KeySort
{
using DummyValueManager = value_manager::DummyValueManager;
using Internal = internal::ParallelRadixSortInternal<PlainType,
CompareType,
UnsignedType,
Encoder,
DummyValueManager,
Base>;
public:
void InitAndSort(PlainType* data, size_t num_elems, const CompareType& comp)
{
(void)comp;
DummyValueManager dvm;
Internal::InitAndSort(data, num_elems, &dvm);
}
};
// Frontend class for sorting pairs
template <typename PlainType,
typename ValueType,
typename CompareType,
typename UnsignedType = PlainType,
typename Encoder = encoder::EncoderDummy,
int Base = 8>
class PairSort
{
using ValueManager = value_manager::PairValueManager<PlainType, ValueType, Base>;
using Internal = internal::
ParallelRadixSortInternal<PlainType, CompareType, UnsignedType, Encoder, ValueManager, Base>;
public:
void InitAndSort(PlainType* keys, ValueType* vals, size_t num_elems, const CompareType& comp)
{
(void)comp;
ValueManager vm;
vm.Init(num_elems);
vm.Start(vals, num_elems);
Internal::InitAndSort(keys, num_elems, &vm);
ValueType* res_vals = vm.GetResult();
if (res_vals != vals)
{
for (size_t i = 0; i < num_elems; ++i)
{
vals[i] = res_vals[i];
}
}
}
private:
};
#define KEY_SORT_CASE(plain_type, compare_type, unsigned_type, encoder_type) \
template <> \
class KeySort<plain_type, compare_type> \
: public KeySort<plain_type, compare_type, unsigned_type, encoder::Encoder##encoder_type> \
{ \
}; \
template <typename V> \
class PairSort<plain_type, V, compare_type> \
: public PairSort<plain_type, V, compare_type, unsigned_type, encoder::Encoder##encoder_type> \
{ \
};
// Unsigned integers
KEY_SORT_CASE(unsigned int, std::less<unsigned int>, unsigned int, Unsigned);
KEY_SORT_CASE(unsigned int, std::greater<unsigned int>, unsigned int, Unsigned);
KEY_SORT_CASE(unsigned short int, std::less<unsigned short int>, unsigned short int, Unsigned);
KEY_SORT_CASE(unsigned short int, std::greater<unsigned short int>, unsigned short int, Unsigned);
KEY_SORT_CASE(unsigned long int, std::less<unsigned long int>, unsigned long int, Unsigned);
KEY_SORT_CASE(unsigned long int, std::greater<unsigned long int>, unsigned long int, Unsigned);
KEY_SORT_CASE(unsigned long long int,
std::less<unsigned long long int>,
unsigned long long int,
Unsigned);
KEY_SORT_CASE(unsigned long long int,
std::greater<unsigned long long int>,
unsigned long long int,
Unsigned);
// Unsigned char
KEY_SORT_CASE(unsigned char, std::less<unsigned char>, unsigned char, Unsigned);
KEY_SORT_CASE(unsigned char, std::greater<unsigned char>, unsigned char, Unsigned);
KEY_SORT_CASE(char16_t, std::less<char16_t>, uint16_t, Unsigned);
KEY_SORT_CASE(char16_t, std::greater<char16_t>, uint16_t, Unsigned);
KEY_SORT_CASE(char32_t, std::less<char32_t>, uint32_t, Unsigned);
KEY_SORT_CASE(char32_t, std::greater<char32_t>, uint32_t, Unsigned);
KEY_SORT_CASE(wchar_t, std::less<wchar_t>, uint32_t, Unsigned);
KEY_SORT_CASE(wchar_t, std::greater<wchar_t>, uint32_t, Unsigned);
// Signed integers
KEY_SORT_CASE(char, std::less<char>, unsigned char, Signed);
KEY_SORT_CASE(char, std::greater<char>, unsigned char, Signed);
KEY_SORT_CASE(short, std::less<short>, unsigned short, Signed);
KEY_SORT_CASE(short, std::greater<short>, unsigned short, Signed);
KEY_SORT_CASE(int, std::less<int>, unsigned int, Signed);
KEY_SORT_CASE(int, std::greater<int>, unsigned int, Signed);
KEY_SORT_CASE(long, std::less<long>, unsigned long, Signed);
KEY_SORT_CASE(long, std::greater<long>, unsigned long, Signed);
KEY_SORT_CASE(long long, std::less<long long>, unsigned long long, Signed);
KEY_SORT_CASE(long long, std::greater<long long>, unsigned long long, Signed);
// |signed char| and |char| are treated as different types
KEY_SORT_CASE(signed char, std::less<signed char>, unsigned char, Signed);
KEY_SORT_CASE(signed char, std::greater<signed char>, unsigned char, Signed);
// Floating point numbers
KEY_SORT_CASE(float, std::less<float>, uint32_t, Decimal);
KEY_SORT_CASE(float, std::greater<float>, uint32_t, Decimal);
KEY_SORT_CASE(double, std::less<double>, uint64_t, Decimal);
KEY_SORT_CASE(double, std::greater<double>, uint64_t, Decimal);
#undef KEY_SORT_CASE
template <typename T, typename CompareType>
struct run_kx_radix_sort_keys
{
static void run(T* data, size_t num_elems, const CompareType& comp)
{
std::sort(data, data + num_elems, comp);
}
};
#define KX_SORT_KEYS(key_type) \
template <> \
struct run_kx_radix_sort_keys<key_type, std::less<key_type>> \
{ \
static void run(key_type* data, size_t num_elems, const std::less<key_type>& comp) \
{ \
(void)comp; \
kx::radix_sort(data, data + num_elems); \
} \
};
KX_SORT_KEYS(unsigned short int);
KX_SORT_KEYS(int);
KX_SORT_KEYS(unsigned int);
KX_SORT_KEYS(long int);
KX_SORT_KEYS(unsigned long int);
KX_SORT_KEYS(long long int);
KX_SORT_KEYS(unsigned long long int);
KX_SORT_KEYS(unsigned char);
#undef KX_SORT_KEYS
template <typename T, typename CompareType>
bool use_serial_sort_keys(T* data, size_t num_elems, const CompareType& comp)
{
size_t total_bytes = (num_elems) * sizeof(T);
if (total_bytes < MIN_BYTES_FOR_PARALLEL)
{
run_kx_radix_sort_keys<T, CompareType>::run(data, num_elems, comp);
return true;
}
return false;
}
// Generate radix sort interfaces for key and key value sorts.
#define VTKM_TBB_SORT_EXPORT(key_type) \
void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::greater<key_type>& comp) \
{ \
PairSort<key_type, vtkm::Id, std::greater<key_type>> ps; \
ps.InitAndSort(keys, vals, num_elems, comp); \
} \
void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::less<key_type>& comp) \
{ \
PairSort<key_type, vtkm::Id, std::less<key_type>> ps; \
ps.InitAndSort(keys, vals, num_elems, comp); \
} \
void parallel_radix_sort(key_type* data, size_t num_elems, const std::greater<key_type>& comp) \
{ \
if (!use_serial_sort_keys(data, num_elems, comp)) \
{ \
KeySort<key_type, std::greater<key_type>> ks; \
ks.InitAndSort(data, num_elems, comp); \
} \
} \
void parallel_radix_sort(key_type* data, size_t num_elems, const std::less<key_type>& comp) \
{ \
if (!use_serial_sort_keys(data, num_elems, comp)) \
{ \
KeySort<key_type, std::less<key_type>> ks; \
ks.InitAndSort(data, num_elems, comp); \
} \
}
VTKM_TBB_SORT_EXPORT(short int);
VTKM_TBB_SORT_EXPORT(unsigned short int);
VTKM_TBB_SORT_EXPORT(int);
VTKM_TBB_SORT_EXPORT(unsigned int);
VTKM_TBB_SORT_EXPORT(long int);
VTKM_TBB_SORT_EXPORT(unsigned long int);
VTKM_TBB_SORT_EXPORT(long long int);
VTKM_TBB_SORT_EXPORT(unsigned long long int);
VTKM_TBB_SORT_EXPORT(unsigned char);
VTKM_TBB_SORT_EXPORT(signed char);
VTKM_TBB_SORT_EXPORT(char);
VTKM_TBB_SORT_EXPORT(char16_t);
VTKM_TBB_SORT_EXPORT(char32_t);
VTKM_TBB_SORT_EXPORT(wchar_t);
VTKM_TBB_SORT_EXPORT(float);
VTKM_TBB_SORT_EXPORT(double);
#undef VTKM_TBB_SORT_EXPORT
VTKM_THIRDPARTY_POST_INCLUDE
}
}
}
}

309
vtkm/cont/tbb/internal/ParallelSortTBB.h Normal file
View File

@ -0,0 +1,309 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_tbb_internal_ParallelSort_h
#define vtk_m_cont_tbb_internal_ParallelSort_h
#include <vtkm/BinaryPredicates.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleZip.h>
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/tbb/internal/FunctorsTBB.h>
#include <vtkm/cont/tbb/internal/ParallelSortTBB.hxx>
#include <type_traits>
namespace vtkm
{
namespace cont
{
namespace tbb
{
namespace internal
{
struct RadixSortTag
{
};
struct PSortTag
{
};
template <typename T>
struct is_valid_compare_type : std::integral_constant<bool, false>
{
};
template <typename T>
struct is_valid_compare_type<std::less<T>> : std::integral_constant<bool, true>
{
};
template <typename T>
struct is_valid_compare_type<std::greater<T>> : std::integral_constant<bool, true>
{
};
template <>
struct is_valid_compare_type<vtkm::SortLess> : std::integral_constant<bool, true>
{
};
template <>
struct is_valid_compare_type<vtkm::SortGreater> : std::integral_constant<bool, true>
{
};
template <typename BComp, typename T>
BComp&& get_std_compare(BComp&& b, T&&)
{
return std::forward<BComp>(b);
}
template <typename T>
std::less<T> get_std_compare(vtkm::SortLess, T&&)
{
return std::less<T>{};
}
template <typename T>
std::greater<T> get_std_compare(vtkm::SortGreater, T&&)
{
return std::greater<T>{};
}
template <typename T, typename StorageTag, typename BinaryCompare>
struct sort_tag_type
{
using type = PSortTag;
};
template <typename T, typename BinaryCompare>
struct sort_tag_type<T, vtkm::cont::StorageTagBasic, BinaryCompare>
{
using PrimT = std::is_arithmetic<T>;
using LongDT = std::is_same<T, long double>;
using BComp = is_valid_compare_type<BinaryCompare>;
using type = typename std::conditional<PrimT::value && BComp::value && !LongDT::value,
RadixSortTag,
PSortTag>::type;
};
template <typename T, typename U, typename StorageTagT, typename StorageTagU, class BinaryCompare>
struct sortbykey_tag_type
{
using type = PSortTag;
};
template <typename T, typename U, typename BinaryCompare>
struct sortbykey_tag_type<T,
U,
vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagBasic,
BinaryCompare>
{
using PrimT = std::is_arithmetic<T>;
using PrimU = std::is_arithmetic<U>;
using LongDT = std::is_same<T, long double>;
using BComp = is_valid_compare_type<BinaryCompare>;
using type =
typename std::conditional<PrimT::value && PrimU::value && BComp::value && !LongDT::value,
RadixSortTag,
PSortTag>::type;
};
#define VTKM_TBB_SORT_EXPORT(key_type) \
VTKM_CONT_EXPORT void parallel_radix_sort( \
key_type* data, size_t num_elems, const std::greater<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort( \
key_type* data, size_t num_elems, const std::less<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::greater<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::less<key_type>& comp);
// Generate radix sort interfaces for key and key value sorts.
VTKM_TBB_SORT_EXPORT(short int);
VTKM_TBB_SORT_EXPORT(unsigned short int);
VTKM_TBB_SORT_EXPORT(int);
VTKM_TBB_SORT_EXPORT(unsigned int);
VTKM_TBB_SORT_EXPORT(long int);
VTKM_TBB_SORT_EXPORT(unsigned long int);
VTKM_TBB_SORT_EXPORT(long long int);
VTKM_TBB_SORT_EXPORT(unsigned long long int);
VTKM_TBB_SORT_EXPORT(unsigned char);
VTKM_TBB_SORT_EXPORT(signed char);
VTKM_TBB_SORT_EXPORT(char);
VTKM_TBB_SORT_EXPORT(char16_t);
VTKM_TBB_SORT_EXPORT(char32_t);
VTKM_TBB_SORT_EXPORT(wchar_t);
VTKM_TBB_SORT_EXPORT(float);
VTKM_TBB_SORT_EXPORT(double);
#undef VTKM_TBB_SORT_EXPORT
template <typename T, typename Container, class BinaryCompare>
void parallel_sort(vtkm::cont::ArrayHandle<T, Container>& values, BinaryCompare binary_compare)
{
using SortAlgorithmTag = typename sort_tag_type<T, Container, BinaryCompare>::type;
parallel_sort(values, binary_compare, SortAlgorithmTag{});
}
template <typename HandleType, class BinaryCompare>
void parallel_sort(HandleType& values, BinaryCompare binary_compare, PSortTag)
{
auto arrayPortal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagTBB());
using IteratorsType = vtkm::cont::ArrayPortalToIterators<decltype(arrayPortal)>;
IteratorsType iterators(arrayPortal);
internal::WrappedBinaryOperator<bool, BinaryCompare> wrappedCompare(binary_compare);
::tbb::parallel_sort(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare);
}
template <typename T, typename StorageT, class BinaryCompare>
void parallel_sort(vtkm::cont::ArrayHandle<T, StorageT>& values,
BinaryCompare binary_compare,
RadixSortTag)
{
auto c = get_std_compare(binary_compare, T{});
parallel_radix_sort(
values.GetStorage().GetArray(), static_cast<std::size_t>(values.GetNumberOfValues()), c);
}
template <typename T, typename StorageT, typename U, typename StorageU, class BinaryCompare>
void parallel_sort_bykey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare)
{
using SortAlgorithmTag =
typename sortbykey_tag_type<T, U, StorageT, StorageU, BinaryCompare>::type;
parallel_sort_bykey(keys, values, binary_compare, SortAlgorithmTag{});
}
template <typename T, typename StorageT, typename U, typename StorageU, class BinaryCompare>
void parallel_sort_bykey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare,
PSortTag)
{
using KeyType = vtkm::cont::ArrayHandle<T, StorageT>;
VTKM_CONSTEXPR bool larger_than_64bits = sizeof(U) > sizeof(vtkm::Int64);
if (larger_than_64bits)
{
/// More efficient sort:
/// Move value indexes when sorting and reorder the value array at last
using ValueType = vtkm::cont::ArrayHandle<U, StorageU>;
using IndexType = vtkm::cont::ArrayHandle<vtkm::Id>;
using ZipHandleType = vtkm::cont::ArrayHandleZip<KeyType, IndexType>;
IndexType indexArray;
ValueType valuesScattered;
const vtkm::Id size = values.GetNumberOfValues();
{
auto handle = ArrayHandleIndex(keys.GetNumberOfValues());
auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB());
auto outputPortal =
indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB());
tbb::CopyPortals(inputPortal, outputPortal, 0, 0, keys.GetNumberOfValues());
}
ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, indexArray);
parallel_sort(zipHandle,
vtkm::cont::internal::KeyCompare<T, vtkm::Id, BinaryCompare>(binary_compare),
PSortTag());
tbb::ScatterPortal(values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB()));
{
auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB());
auto outputPortal =
values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB());
tbb::CopyPortals(inputPortal, outputPortal, 0, 0, valuesScattered.GetNumberOfValues());
}
}
else
{
using ValueType = vtkm::cont::ArrayHandle<U, StorageU>;
using ZipHandleType = vtkm::cont::ArrayHandleZip<KeyType, ValueType>;
ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values);
parallel_sort(
zipHandle, vtkm::cont::internal::KeyCompare<T, U, BinaryCompare>(binary_compare), PSortTag{});
}
}
template <typename T, typename StorageT, typename StorageU, class BinaryCompare>
void parallel_sort_bykey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<vtkm::Id, StorageU>& values,
BinaryCompare binary_compare,
RadixSortTag)
{
auto c = get_std_compare(binary_compare, T{});
parallel_radix_sort_key_values(keys.GetStorage().GetArray(),
values.GetStorage().GetArray(),
static_cast<std::size_t>(keys.GetNumberOfValues()),
c);
}
template <typename T, typename StorageT, typename U, typename StorageU, class BinaryCompare>
void parallel_sort_bykey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare,
RadixSortTag)
{
using KeyType = vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>;
using ValueType = vtkm::cont::ArrayHandle<U, vtkm::cont::StorageTagBasic>;
using IndexType = vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic>;
using ZipHandleType = vtkm::cont::ArrayHandleZip<KeyType, IndexType>;
IndexType indexArray;
ValueType valuesScattered;
const vtkm::Id size = values.GetNumberOfValues();
{
auto handle = ArrayHandleIndex(keys.GetNumberOfValues());
auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB());
auto outputPortal =
indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB());
tbb::CopyPortals(inputPortal, outputPortal, 0, 0, keys.GetNumberOfValues());
}
if (static_cast<vtkm::Id>(sizeof(T)) * keys.GetNumberOfValues() > 400000)
{
parallel_sort_bykey(keys, indexArray, binary_compare);
}
else
{
ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, indexArray);
parallel_sort(zipHandle,
vtkm::cont::internal::KeyCompare<T, vtkm::Id, BinaryCompare>(binary_compare),
PSortTag{});
}
tbb::ScatterPortal(values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()),
valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB()));
{
auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB());
auto outputPortal =
values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB());
tbb::CopyPortals(inputPortal, outputPortal, 0, 0, valuesScattered.GetNumberOfValues());
}
}
}
}
}
}
#endif // vtk_m_cont_tbb_internal_ParallelSort_h

0
vtkm/cont/tbb/internal/parallel_sort.h → vtkm/cont/tbb/internal/ParallelSortTBB.hxx
View File

176
vtkm/cont/tbb/internal/kxsort.h Normal file
View File

@ -0,0 +1,176 @@
/* The MIT License
Copyright (c) 2016 Dinghua Li <voutcn@gmail.com>
Permission is hereby granted, free of charge, to any person obtaining
a copy of this software and associated documentation files (the
"Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be
included in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS
BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN
ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.
*/
#ifndef KXSORT_H__
#define KXSORT_H__
#include <algorithm>
#include <iterator>
namespace kx
{
static const int kRadixBits = 8;
static const size_t kInsertSortThreshold = 64;
static const int kRadixMask = (1 << kRadixBits) - 1;
static const int kRadixBin = 1 << kRadixBits;
//================= HELPING FUNCTIONS ====================
template <class T>
struct RadixTraitsUnsigned
{
static const int nBytes = sizeof(T);
int kth_byte(const T& x, int k) { return (x >> (kRadixBits * k)) & kRadixMask; }
bool compare(const T& x, const T& y) { return x < y; }
};
template <class T>
struct RadixTraitsSigned
{
static const int nBytes = sizeof(T);
static const T kMSB = T(0x80) << ((sizeof(T) - 1) * 8);
int kth_byte(const T& x, int k) { return ((x ^ kMSB) >> (kRadixBits * k)) & kRadixMask; }
bool compare(const T& x, const T& y) { return x < y; }
};
template <class RandomIt, class ValueType, class RadixTraits>
inline void insert_sort_core_(RandomIt s, RandomIt e, RadixTraits radix_traits)
{
for (RandomIt i = s + 1; i < e; ++i)
{
if (radix_traits.compare(*i, *(i - 1)))
{
RandomIt j;
ValueType tmp = *i;
*i = *(i - 1);
for (j = i - 1; j > s && radix_traits.compare(tmp, *(j - 1)); --j)
{
*j = *(j - 1);
}
*j = tmp;
}
}
}
template <class RandomIt, class ValueType, class RadixTraits, int kWhichByte>
inline void radix_sort_core_(RandomIt s, RandomIt e, RadixTraits radix_traits)
{
RandomIt last_[kRadixBin + 1];
RandomIt* last = last_ + 1;
size_t count[kRadixBin] = { 0 };
for (RandomIt i = s; i < e; ++i)
{
++count[radix_traits.kth_byte(*i, kWhichByte)];
}
last_[0] = last_[1] = s;
for (int i = 1; i < kRadixBin; ++i)
{
last[i] = last[i - 1] + count[i - 1];
}
for (int i = 0; i < kRadixBin; ++i)
{
RandomIt end = last[i - 1] + count[i];
if (end == e)
{
last[i] = e;
break;
}
while (last[i] != end)
{
ValueType swapper = *last[i];
int tag = radix_traits.kth_byte(swapper, kWhichByte);
if (tag != i)
{
do
{
std::swap(swapper, *last[tag]++);
} while ((tag = radix_traits.kth_byte(swapper, kWhichByte)) != i);
*last[i] = swapper;
}
++last[i];
}
}
if (kWhichByte > 0)
{
for (int i = 0; i < kRadixBin; ++i)
{
if (count[i] > kInsertSortThreshold)
{
radix_sort_core_<RandomIt, ValueType, RadixTraits, (kWhichByte > 0 ? (kWhichByte - 1) : 0)>(
last[i - 1], last[i], radix_traits);
}
else if (count[i] > 1)
{
insert_sort_core_<RandomIt, ValueType, RadixTraits>(last[i - 1], last[i], radix_traits);
}
}
}
}
template <class RandomIt, class ValueType, class RadixTraits>
inline void radix_sort_entry_(RandomIt s, RandomIt e, ValueType*, RadixTraits radix_traits)
{
if (e - s <= (int)kInsertSortThreshold)
insert_sort_core_<RandomIt, ValueType, RadixTraits>(s, e, radix_traits);
else
radix_sort_core_<RandomIt, ValueType, RadixTraits, RadixTraits::nBytes - 1>(s, e, radix_traits);
}
template <class RandomIt, class ValueType>
inline void radix_sort_entry_(RandomIt s, RandomIt e, ValueType*)
{
if (ValueType(-1) > ValueType(0))
{
radix_sort_entry_(s, e, (ValueType*)(0), RadixTraitsUnsigned<ValueType>());
}
else
{
radix_sort_entry_(s, e, (ValueType*)(0), RadixTraitsSigned<ValueType>());
}
}
//================= INTERFACES ====================
template <class RandomIt, class RadixTraits>
inline void radix_sort(RandomIt s, RandomIt e, RadixTraits radix_traits)
{
typename std::iterator_traits<RandomIt>::value_type* dummy(0);
radix_sort_entry_(s, e, dummy, radix_traits);
}
template <class RandomIt>
inline void radix_sort(RandomIt s, RandomIt e)
{
typename std::iterator_traits<RandomIt>::value_type* dummy(0);
radix_sort_entry_(s, e, dummy);
}
}
#endif

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

@ -31,4 +31,5 @@ set(unit_tests
UnitTestTBBPointLocatorUniformGrid.cxx
UnitTestTBBVirtualObjectHandle.cxx
)
vtkm_unit_tests(SOURCES ${unit_tests} BACKEND TBB)

2
vtkm/cont/testing/TestingCellLocatorTwoLevelUniformGrid.h
View File

@ -128,7 +128,7 @@ vtkm::cont::DataSet MakeTestDataSet(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dims,
VTKM_ASSERT(false);
}
// It is posible that the warping will result in invalid cells. So use a
// It is possible that the warping will result in invalid cells. So use a
// local random generator with a known seed that does not create invalid cells.
std::default_random_engine rgen;

1
vtkm/cont/testing/TestingDeviceAdapter.h
View File

@ -1070,7 +1070,6 @@ private:
std::cout << " Reduce with 0 values." << std::endl;
array.Shrink(0);
vtkm::Id reduce_sum_no_values = Algorithm::Reduce(array, vtkm::Id(0));
VTKM_TEST_ASSERT(reduce_sum == OFFSET * ARRAY_SIZE, "Got bad sum from Reduce");
VTKM_TEST_ASSERT(reduce_sum_with_intial_value == reduce_sum + ARRAY_SIZE,
"Got bad sum from Reduce with initial value");

2
vtkm/cont/testing/TestingPointLocatorUniformGrid.h
View File

@ -79,7 +79,7 @@ template <typename DeviceAdapter>
class TestingPointLocatorUniformGrid
{
public:
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG> Algorithm;
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>;
void TestTest() const
{
vtkm::Int32 nTrainingPoints = 1000;

2
vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
View File

@ -126,7 +126,7 @@ struct ExecutionArrayInterfaceBasic<DeviceAdapterTagTestAlgorithmGeneral>
}
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return -3; }
DeviceAdapterId GetDeviceId() const final { return -3; }
};
}
}

62
vtkm/cont/testing/UnitTestStorageBasic.cxx
View File

@ -26,16 +26,6 @@
#include <vtkm/VecTraits.h>
#include <vtkm/cont/testing/Testing.h>
// We use these to check if the aligned allocator provided by
// StorageBasic can be used with all STL containers
#include <deque>
#include <list>
#include <map>
#include <queue>
#include <set>
#include <stack>
#include <vector>
namespace
{
@ -72,53 +62,6 @@ struct TemplatedTests
typename vtkm::VecTraits<ValueType>::ComponentType STOLEN_ARRAY_VALUE() { return 29; }
void TestAlignedAllocatorSTL()
{
using Allocator = typename StorageType::AllocatorType;
std::vector<ValueType, Allocator> vec(ARRAY_SIZE, ValueType());
StorageType store(&vec[0], ARRAY_SIZE);
}
// This test checks that we can compile and use the allocator with all
// STL containers
void CompileSTLAllocator()
{
using Allocator = typename StorageType::AllocatorType;
using PairAllocator =
typename StorageType::AllocatorType::template rebind<std::pair<ValueType, ValueType>>::other;
std::vector<ValueType, Allocator> v;
ValueType value = vtkm::TypeTraits<ValueType>::ZeroInitialization();
v.push_back(value);
std::deque<ValueType, Allocator> d;
d.push_front(value);
std::list<ValueType, Allocator> l;
l.push_front(value);
std::set<ValueType, std::less<ValueType>, Allocator> set;
set.insert(value);
std::map<ValueType, ValueType, std::less<ValueType>, PairAllocator> m;
m[value] = value;
std::multiset<ValueType, std::less<ValueType>, Allocator> ms;
ms.insert(value);
std::multimap<ValueType, ValueType, std::less<ValueType>, PairAllocator> mm;
mm.insert(std::pair<ValueType, ValueType>(value, value));
std::stack<ValueType, std::deque<ValueType, Allocator>> stack;
stack.push(value);
std::queue<ValueType, std::deque<ValueType, Allocator>> queue;
queue.push(value);
std::priority_queue<ValueType, std::vector<ValueType, Allocator>> pqueue;
pqueue.push(value);
}
/// Returned value should later be passed to StealArray2. It is best to
/// put as much between the two test parts to maximize the chance of a
/// deallocated array being overridden (and thus detected).
@ -153,7 +96,7 @@ struct TemplatedTests
"Stolen array did not retain values.");
}
typename StorageType::AllocatorType allocator;
allocator.deallocate(stolenArray, ARRAY_SIZE);
allocator.deallocate(stolenArray);
}
void BasicAllocation()
@ -199,9 +142,6 @@ struct TemplatedTests
BasicAllocation();
StealArray2(stolenArray);
TestAlignedAllocatorSTL();
CompileSTLAllocator();
}
};

1
vtkm/exec/CMakeLists.txt
View File

@ -31,7 +31,6 @@ set(headers
ExecutionObjectBase.h
ExecutionWholeArray.h
FunctorBase.h
ImplicitFunction.h
Jacobian.h
ParametricCoordinates.h
TaskBase.h

48
vtkm/exec/CellDerivative.h
View File

@ -45,9 +45,9 @@ namespace exec
namespace
{
#define VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(pointIndex, weight0, weight1, weight2) \
parametricDerivative[0] += field[pointIndex] * weight0; \
parametricDerivative[1] += field[pointIndex] * weight1; \
parametricDerivative[2] += field[pointIndex] * weight2
parametricDerivative[0] += field[pointIndex] * (weight0); \
parametricDerivative[1] += field[pointIndex] * (weight1); \
parametricDerivative[2] += field[pointIndex] * (weight2)
// Find the derivative of a field in parametric space. That is, find the
// vector [ds/du, ds/dv, ds/dw].
@ -64,14 +64,7 @@ VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 3> ParametricDerivativ
GradientType rc = GradientType(FieldType(1)) - pc;
GradientType parametricDerivative(FieldType(0));
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(0, -rc[1] * rc[2], -rc[0] * rc[2], -rc[0] * rc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(1, rc[1] * rc[2], -pc[0] * rc[2], -pc[0] * rc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(2, pc[1] * rc[2], pc[0] * rc[2], -pc[0] * pc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(3, -pc[1] * rc[2], rc[0] * rc[2], -rc[0] * pc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(4, -rc[1] * pc[2], -rc[0] * pc[2], rc[0] * rc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(5, rc[1] * pc[2], -pc[0] * pc[2], pc[0] * rc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(6, pc[1] * pc[2], pc[0] * pc[2], pc[0] * pc[1]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D(7, -pc[1] * pc[2], rc[0] * pc[2], rc[0] * pc[1]);
VTKM_DERIVATIVE_WEIGHTS_HEXAHEDRON(pc, rc, VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D);
return parametricDerivative;
}
@ -81,22 +74,16 @@ VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 3> ParametricDerivativ
const vtkm::Vec<ParametricCoordType, 3>& pcoords,
vtkm::CellShapeTagWedge)
{
#if 0
// This is not working. Just leverage the hexahedron code that is working.
using FieldType = typename FieldVecType::ComponentType;
using GradientType = vtkm::Vec<FieldType,3>;
using GradientType = vtkm::Vec<FieldType, 3>;
GradientType pc(pcoords);
GradientType rc = GradientType(1) - pc;
GradientType rc = GradientType(FieldType(1)) - pc;
GradientType parametricDerivative(0);
GradientType parametricDerivative(FieldType(0));
VTKM_DERIVATIVE_WEIGHTS_WEDGE(pc, rc, VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D);
return parametricDerivative;
#else
return ParametricDerivative(
vtkm::exec::internal::PermuteWedgeToHex(field), pcoords, vtkm::CellShapeTagHexahedron());
#endif
}
template <typename FieldVecType, typename ParametricCoordType>
@ -105,29 +92,23 @@ VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 3> ParametricDerivativ
const vtkm::Vec<ParametricCoordType, 3>& pcoords,
vtkm::CellShapeTagPyramid)
{
#if 0
// This is not working. Just leverage the hexahedron code that is working.
using FieldType = typename FieldVecType::ComponentType;
using GradientType = vtkm::Vec<FieldType,3>;
using GradientType = vtkm::Vec<FieldType, 3>;
GradientType pc(pcoords);
GradientType rc = GradientType(1) - pc;
GradientType rc = GradientType(FieldType(1)) - pc;
GradientType parametricDerivative(0);
GradientType parametricDerivative(FieldType(0));
VTKM_DERIVATIVE_WEIGHTS_PYRAMID(pc, rc, VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D);
return parametricDerivative;
#else
return ParametricDerivative(
vtkm::exec::internal::PermutePyramidToHex(field), pcoords, vtkm::CellShapeTagHexahedron());
#endif
}
#undef VTKM_ACCUM_PARAMETRIC_DERIVATIVE_3D
#define VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D(pointIndex, weight0, weight1) \
parametricDerivative[0] += field[pointIndex] * weight0; \
parametricDerivative[1] += field[pointIndex] * weight1
parametricDerivative[0] += field[pointIndex] * (weight0); \
parametricDerivative[1] += field[pointIndex] * (weight1)
template <typename FieldVecType, typename ParametricCoordType>
VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 2> ParametricDerivative(
@ -142,10 +123,7 @@ VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 2> ParametricDerivativ
GradientType rc = GradientType(FieldType(1)) - pc;
GradientType parametricDerivative(FieldType(0));
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D(0, -rc[1], -rc[0]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D(1, rc[1], -pc[0]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D(2, pc[1], pc[0]);
VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D(3, -pc[1], rc[0]);
VTKM_DERIVATIVE_WEIGHTS_QUAD(pc, rc, VTKM_ACCUM_PARAMETRIC_DERIVATIVE_2D);
return parametricDerivative;
}

2
vtkm/exec/CellFace.h
View File

@ -265,7 +265,7 @@ static inline VTKM_EXEC vtkm::VecCConst<vtkm::IdComponent> CellFaceLocalIndices(
numPointsInFace);
}
/// \brief Returns a canonical identifer for a cell face
/// \brief Returns a canonical identifier for a cell face
///
/// Given information about a cell face and the global point indices for that cell, returns a
/// vtkm::Id3 that contains values that are unique to that face. The values for two faces will be

8
vtkm/exec/CellInterpolate.h
View File

@ -67,8 +67,8 @@ VTKM_EXEC typename WorldCoordVector::ComponentType ReverseInterpolateTriangle(
//
// In this diagram, the distance between p0 and the point marked x divided by
// the length of the edge it is on is equal, by proportionality, to the u
// parametric coordiante. (The v coordinate follows the other edge
// accordingly.) Thus, if we can find the interesection at x (or more
// parametric coordinate. (The v coordinate follows the other edge
// accordingly.) Thus, if we can find the intersection at x (or more
// specifically the distance between p0 and x), then we can find that
// parametric coordinate.
//
@ -94,13 +94,13 @@ VTKM_EXEC typename WorldCoordVector::ComponentType ReverseInterpolateTriangle(
// compute it, we are done. We can skip the part about finding the actual
// coordinates of the intersection.
//
// Solving for the interesection is as simple as substituting the line's
// Solving for the intersection is as simple as substituting the line's
// definition of p(d) into p for the plane equation. With some basic algebra
// you get:
//
// d = dot((wcoords - p0), planeNormal)/dot((p1-p0), planeNormal)
//
// From here, the u coordiante is simply d. The v coordinate follows
// From here, the u coordinate is simply d. The v coordinate follows
// similarly.
//

145
vtkm/exec/ImplicitFunction.h
View File

@ -1,145 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_exec_ImplicitFunction_h
#define vtk_m_exec_ImplicitFunction_h
#include <vtkm/Types.h>
namespace vtkm
{
namespace exec
{
class VTKM_ALWAYS_EXPORT ImplicitFunction
{
public:
ImplicitFunction()
: Function(nullptr)
, ValueCaller(nullptr)
, GradientCaller(nullptr)
{
}
VTKM_EXEC
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->ValueCaller(this->Function, x, y, z);
}
VTKM_EXEC
FloatDefault Value(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->ValueCaller(this->Function, x[0], x[1], x[2]);
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->GradientCaller(this->Function, x, y, z);
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3>& x) const
{
return this->GradientCaller(this->Function, x[0], x[1], x[2]);
}
template <typename T>
VTKM_EXEC void Bind(const T* function)
{
this->Function = function;
this->ValueCaller = [](const void* t, FloatDefault x, FloatDefault y, FloatDefault z) {
return static_cast<const T*>(t)->Value(x, y, z);
};
this->GradientCaller = [](const void* t, FloatDefault x, FloatDefault y, FloatDefault z) {
return static_cast<const T*>(t)->Gradient(x, y, z);
};
}
private:
using ValueCallerSig = FloatDefault(const void*, FloatDefault, FloatDefault, FloatDefault);
using GradientCallerSig = vtkm::Vec<FloatDefault, 3>(const void*,
FloatDefault,
FloatDefault,
FloatDefault);
const void* Function;
ValueCallerSig* ValueCaller;
GradientCallerSig* GradientCaller;
};
/// \brief A function object that evaluates the contained implicit function
class VTKM_ALWAYS_EXPORT ImplicitFunctionValue
{
public:
ImplicitFunctionValue() = default;
explicit ImplicitFunctionValue(const vtkm::exec::ImplicitFunction& func)
: Function(func)
{
}
VTKM_EXEC
FloatDefault operator()(const vtkm::Vec<FloatDefault, 3> x) const
{
return this->Function.Value(x);
}
VTKM_EXEC
FloatDefault operator()(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->Function.Value(x, y, z);
}
private:
vtkm::exec::ImplicitFunction Function;
};
/// \brief A function object that computes the gradient of the contained implicit
/// function and the specified point.
class VTKM_ALWAYS_EXPORT ImplicitFunctionGradient
{
public:
ImplicitFunctionGradient() = default;
explicit ImplicitFunctionGradient(const vtkm::exec::ImplicitFunction& func)
: Function(func)
{
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> operator()(const vtkm::Vec<FloatDefault, 3> x) const
{
return this->Function.Gradient(x);
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> operator()(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->Function.Gradient(x, y, z);
}
private:
vtkm::exec::ImplicitFunction Function;
};
}
} // vtkm::exec
#endif // vtk_m_exec_ImplicitFunction_h

137
vtkm/exec/Jacobian.h
View File

@ -68,98 +68,38 @@ struct Space2D
}
};
// Given a series of point values for a wedge, return a new series of point
// for a hexahedron that has the same interpolation within the wedge.
template <typename FieldVecType>
VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 8> PermuteWedgeToHex(
const FieldVecType& field)
{
vtkm::Vec<typename FieldVecType::ComponentType, 8> hexField;
hexField[0] = field[0];
hexField[1] = field[2];
hexField[2] = field[2] + field[1] - field[0];
hexField[3] = field[1];
hexField[4] = field[3];
hexField[5] = field[5];
hexField[6] = field[5] + field[4] - field[3];
hexField[7] = field[4];
return hexField;
}
// Given a series of point values for a pyramid, return a new series of point
// for a hexahedron that has the same interpolation within the pyramid.
template <typename FieldVecType>
VTKM_EXEC vtkm::Vec<typename FieldVecType::ComponentType, 8> PermutePyramidToHex(
const FieldVecType& field)
{
using T = typename FieldVecType::ComponentType;
vtkm::Vec<T, 8> hexField;
T baseCenter = T(0.25f) * (field[0] + field[1] + field[2] + field[3]);
hexField[0] = field[0];
hexField[1] = field[1];
hexField[2] = field[2];
hexField[3] = field[3];
hexField[4] = field[4] + (field[0] - baseCenter);
hexField[5] = field[4] + (field[1] - baseCenter);
hexField[6] = field[4] + (field[2] - baseCenter);
hexField[7] = field[4] + (field[3] - baseCenter);
return hexField;
}
} //namespace internal
#define VTKM_DERIVATIVE_WEIGHTS_HEXAHEDRON(pc, rc, call) \
call(0, -rc[1] * rc[2], -rc[0] * rc[2], -rc[0] * rc[1]); \
call(1, rc[1] * rc[2], -pc[0] * rc[2], -pc[0] * rc[1]); \
call(2, pc[1] * rc[2], pc[0] * rc[2], -pc[0] * pc[1]); \
call(3, -pc[1] * rc[2], rc[0] * rc[2], -rc[0] * pc[1]); \
call(4, -rc[1] * pc[2], -rc[0] * pc[2], rc[0] * rc[1]); \
call(5, rc[1] * pc[2], -pc[0] * pc[2], pc[0] * rc[1]); \
call(6, pc[1] * pc[2], pc[0] * pc[2], pc[0] * pc[1]); \
call(7, -pc[1] * pc[2], rc[0] * pc[2], rc[0] * pc[1])
#define VTKM_DERIVATIVE_WEIGHTS_VOXEL(pc, rc, call) \
call(0, -rc[1] * rc[2], -rc[0] * rc[2], -rc[0] * rc[1]); \
call(1, rc[1] * rc[2], -pc[0] * rc[2], -pc[0] * rc[1]); \
call(2, -pc[1] * rc[2], rc[0] * rc[2], -rc[0] * pc[1]); \
call(3, pc[1] * rc[2], pc[0] * rc[2], -pc[0] * pc[1]); \
call(4, -rc[1] * pc[2], -rc[0] * pc[2], rc[0] * rc[1]); \
call(5, rc[1] * pc[2], -pc[0] * pc[2], pc[0] * rc[1]); \
call(6, -pc[1] * pc[2], rc[0] * pc[2], rc[0] * pc[1]); \
call(7, pc[1] * pc[2], pc[0] * pc[2], pc[0] * pc[1])
call(0, (-rc[1] * rc[2]), (-rc[0] * rc[2]), (-rc[0] * rc[1])); \
call(1, (rc[1] * rc[2]), (-pc[0] * rc[2]), (-pc[0] * rc[1])); \
call(2, (pc[1] * rc[2]), (pc[0] * rc[2]), (-pc[0] * pc[1])); \
call(3, (-pc[1] * rc[2]), (rc[0] * rc[2]), (-rc[0] * pc[1])); \
call(4, (-rc[1] * pc[2]), (-rc[0] * pc[2]), (rc[0] * rc[1])); \
call(5, (rc[1] * pc[2]), (-pc[0] * pc[2]), (pc[0] * rc[1])); \
call(6, (pc[1] * pc[2]), (pc[0] * pc[2]), (pc[0] * pc[1])); \
call(7, (-pc[1] * pc[2]), (rc[0] * pc[2]), (rc[0] * pc[1]))
#define VTKM_DERIVATIVE_WEIGHTS_WEDGE(pc, rc, call) \
call(0, -rc[2], -rc[2], -1.0f + pc[0] + pc[1]); \
call(1, 0.0f, rc[2], -pc[1]); \
call(2, rc[2], 0.0f, -pc[0]); \
call(3, -pc[2], -pc[2], 1.0f - pc[0] - pc[1]); \
call(4, 0.0f, pc[2], pc[1]); \
call(5, pc[2], 0.0f, pc[0])
call(0, (-rc[2]), (-rc[2]), (pc[0] - rc[1])); \
call(1, (0.0f), (rc[2]), (-pc[1])); \
call(2, (rc[2]), (0.0f), (-pc[0])); \
call(3, (-pc[2]), (-pc[2]), (rc[0] - pc[1])); \
call(4, (0.0f), (pc[2]), (pc[1])); \
call(5, (pc[2]), (0.0f), (pc[0]))
#define VTKM_DERIVATIVE_WEIGHTS_PYRAMID(pc, rc, call) \
call(0, -rc[1] * rc[2], -rc[0] * rc[2], -rc[0] * rc[1]); \
call(1, rc[1] * rc[2], -pc[0] * rc[2], -pc[0] * rc[1]); \
call(2, pc[1] * rc[2], pc[0] * rc[2], -pc[0] * pc[1]); \
call(3, -pc[1] * rc[2], rc[0] * rc[2], -rc[0] * pc[1]); \
call(3, 0.0f, 0.0f, 1.0f)
call(0, (-rc[1] * rc[2]), (-rc[0] * rc[2]), (-rc[0] * rc[1])); \
call(1, (rc[1] * rc[2]), (-pc[0] * rc[2]), (-pc[0] * rc[1])); \
call(2, (pc[1] * rc[2]), (pc[0] * rc[2]), (-pc[0] * pc[1])); \
call(3, (-pc[1] * rc[2]), (rc[0] * rc[2]), (-rc[0] * pc[1])); \
call(4, (0.0f), (0.0f), (1.0f))
#define VTKM_DERIVATIVE_WEIGHTS_QUAD(pc, rc, call) \
call(0, -rc[1], -rc[0]); \
call(1, rc[1], -pc[0]); \
call(2, pc[1], pc[0]); \
call(3, -pc[1], rc[0])
#define VTKM_DERIVATIVE_WEIGHTS_PIXEL(pc, rc, call) \
call(0, -rc[1], -rc[0]); \
call(1, rc[1], -pc[0]); \
call(2, -pc[1], rc[0]); \
call(3, pc[1], pc[0])
call(0, (-rc[1]), (-rc[0])); \
call(1, (rc[1]), (-pc[0])); \
call(2, (pc[1]), (pc[0])); \
call(3, (-pc[1]), (rc[0]))
//-----------------------------------------------------------------------------
// This returns the Jacobian of a hexahedron's (or other 3D cell's) coordinates
@ -205,17 +145,11 @@ VTKM_EXEC void JacobianFor3DCell(const WorldCoordType& wCoords,
vtkm::Matrix<JacobianType, 3, 3>& jacobian,
vtkm::CellShapeTagWedge)
{
#if 0
// This is not working. Just leverage the hexahedron code that is working.
vtkm::Vec<JacobianType,3> pc(pcoords);
vtkm::Vec<JacobianType,3> rc = vtkm::Vec<JacobianType,3>(1) - pc;
vtkm::Vec<JacobianType, 3> pc(pcoords);
vtkm::Vec<JacobianType, 3> rc = vtkm::Vec<JacobianType, 3>(JacobianType(1)) - pc;
jacobian = vtkm::Matrix<JacobianType,3,3>(0);
jacobian = vtkm::Matrix<JacobianType, 3, 3>(JacobianType(0));
VTKM_DERIVATIVE_WEIGHTS_WEDGE(pc, rc, VTKM_ACCUM_JACOBIAN_3D);
#else
JacobianFor3DCell(
internal::PermuteWedgeToHex(wCoords), pcoords, jacobian, vtkm::CellShapeTagHexahedron());
#endif
}
template <typename WorldCoordType, typename ParametricCoordType, typename JacobianType>
@ -224,17 +158,11 @@ VTKM_EXEC void JacobianFor3DCell(const WorldCoordType& wCoords,
vtkm::Matrix<JacobianType, 3, 3>& jacobian,
vtkm::CellShapeTagPyramid)
{
#if 0
// This is not working. Just leverage the hexahedron code that is working.
vtkm::Vec<JacobianType,3> pc(pcoords);
vtkm::Vec<JacobianType,3> rc = vtkm::Vec<JacobianType,3>(1) - pc;
vtkm::Vec<JacobianType, 3> pc(pcoords);
vtkm::Vec<JacobianType, 3> rc = vtkm::Vec<JacobianType, 3>(JacobianType(1)) - pc;
jacobian = vtkm::Matrix<JacobianType,3,3>(0);
jacobian = vtkm::Matrix<JacobianType, 3, 3>(JacobianType(0));
VTKM_DERIVATIVE_WEIGHTS_PYRAMID(pc, rc, VTKM_ACCUM_JACOBIAN_3D);
#else
JacobianFor3DCell(
internal::PermutePyramidToHex(wCoords), pcoords, jacobian, vtkm::CellShapeTagHexahedron());
#endif
}
// Derivatives in quadrilaterals are computed in much the same way as
@ -316,13 +244,6 @@ void JacobianFor2DCell(const WorldCoordType &wCoords,
#undef VTKM_ACCUM_JACOBIAN_3D
#undef VTKM_ACCUM_JACOBIAN_2D
#undef VTKM_DERIVATIVE_WEIGHTS_HEXAHEDRON
#undef VTKM_DERIVATIVE_WEIGHTS_VOXEL
#undef VTKM_DERIVATIVE_WEIGHTS_WEDGE
#undef VTKM_DERIVATIVE_WEIGHTS_PYRAMID
#undef VTKM_DERIVATIVE_WEIGHTS_QUAD
#undef VTKM_DERIVATIVE_WEIGHTS_PIXEL
}
} // namespace vtkm::exec
#endif //vtk_m_exec_Jacobian_h

449
vtkm/exec/ParametricCoordinates.h
View File

@ -28,6 +28,7 @@
#include <vtkm/exec/CellInterpolate.h>
#include <vtkm/exec/FunctorBase.h>
#include <vtkm/exec/Jacobian.h>
#include <vtkm/exec/internal/FastVec.h>
#include <vtkm/internal/Assume.h>
namespace vtkm
@ -586,7 +587,77 @@ ParametricCoordinatesToWorldCoordinates(const WorldCoordVector& pointWCoords,
}
//-----------------------------------------------------------------------------
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector&,
const typename WorldCoordVector::ComponentType&,
vtkm::CellShapeTagEmpty,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
worklet.RaiseError("Attempted to find point coordinates in empty cell.");
success = false;
return typename WorldCoordVector::ComponentType();
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType&,
vtkm::CellShapeTagVertex,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
{
(void)pointWCoords; // Silence compiler warnings.
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 1);
success = true;
return typename WorldCoordVector::ComponentType(0, 0, 0);
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagLine,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 2);
success = true;
// Because this is a line, there is only one valid parametric coordinate. Let
// vec be the vector from the first point to the second point
// (pointWCoords[1] - pointWCoords[0]), which is the direction of the line.
// dot(vec,wcoords-pointWCoords[0])/mag(vec) is the orthoginal projection of
// wcoords on the line and represents the distance between the orthoginal
// projection and pointWCoords[0]. The parametric coordinate is the fraction
// of this over the length of the segment, which is mag(vec). Thus, the
// parametric coordinate is dot(vec,wcoords-pointWCoords[0])/mag(vec)^2.
using Vector3 = typename WorldCoordVector::ComponentType;
using T = typename Vector3::ComponentType;
Vector3 vec = pointWCoords[1] - pointWCoords[0];
T numerator = vtkm::dot(vec, wcoords - pointWCoords[0]);
T denominator = vtkm::MagnitudeSquared(vec);
return Vector3(numerator / denominator, 0, 0);
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagTriangle,
bool& success,
const vtkm::exec::FunctorBase&)
{
success = true;
vtkm::exec::internal::FastVec<WorldCoordVector, 3> local(pointWCoords);
return vtkm::exec::internal::ReverseInterpolateTriangle(local.Get(), wcoords);
}
//-----------------------------------------------------------------------------
namespace detail
{
@ -655,160 +726,37 @@ public:
}
};
template <typename WorldCoordVector, typename CellShapeTag>
class JacobianFunctor3DCell
{
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector3 = vtkm::Vec<T, 3>;
using Matrix3x3 = vtkm::Matrix<T, 3, 3>;
const WorldCoordVector* PointWCoords;
public:
VTKM_EXEC
JacobianFunctor3DCell(const WorldCoordVector* pointWCoords)
: PointWCoords(pointWCoords)
{
}
VTKM_EXEC
Matrix3x3 operator()(const Vector3& pcoords) const
{
Matrix3x3 jacobian;
vtkm::exec::JacobianFor3DCell(*this->PointWCoords, pcoords, jacobian, CellShapeTag());
return jacobian;
}
};
template <typename WorldCoordVector, typename CellShapeTag>
class CoordinatesFunctor3DCell
{
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector3 = vtkm::Vec<T, 3>;
const WorldCoordVector* PointWCoords;
const vtkm::exec::FunctorBase* Worklet;
public:
VTKM_EXEC
CoordinatesFunctor3DCell(const WorldCoordVector* pointWCoords,
const vtkm::exec::FunctorBase* worklet)
: PointWCoords(pointWCoords)
, Worklet(worklet)
{
}
VTKM_EXEC
Vector3 operator()(Vector3 pcoords) const
{
return vtkm::exec::ParametricCoordinatesToWorldCoordinates(
*this->PointWCoords, pcoords, CellShapeTag(), *this->Worklet);
}
};
template <typename WorldCoordVector, typename CellShapeTag>
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates3D(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
CellShapeTag,
bool& success,
const vtkm::exec::FunctorBase& worklet)
WorldCoordinatesToParametricCoordinatesQuad(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
auto result = vtkm::NewtonsMethod(
JacobianFunctor3DCell<WorldCoordVector, CellShapeTag>(&pointWCoords),
CoordinatesFunctor3DCell<WorldCoordVector, CellShapeTag>(&pointWCoords, &worklet),
wcoords,
typename WorldCoordVector::ComponentType(0.5f, 0.5f, 0.5f));
success = result.Valid;
return result.Solution;
}
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector2 = vtkm::Vec<T, 2>;
using Vector3 = vtkm::Vec<T, 3>;
// We have an underdetermined system in 3D, so create a 2D space in the
// plane that the polygon sits.
vtkm::exec::internal::Space2D<T> space(pointWCoords[0], pointWCoords[1], pointWCoords[3]);
auto result = vtkm::NewtonsMethod(
JacobianFunctorQuad<WorldCoordVector, vtkm::CellShapeTagQuad>(&pointWCoords, &space),
CoordinatesFunctorQuad<WorldCoordVector, vtkm::CellShapeTagQuad>(
&pointWCoords, &space, &worklet),
space.ConvertCoordToSpace(wcoords),
Vector2(0.5f, 0.5f));
success = result.Valid;
return Vector3(result.Solution[0], result.Solution[1], 0);
}
} // namespace detail
//-----------------------------------------------------------------------------
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagGeneric shape,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
typename WorldCoordVector::ComponentType result;
switch (shape.Id)
{
vtkmGenericCellShapeMacro(result = WorldCoordinatesToParametricCoordinates(
pointWCoords, wcoords, CellShapeTag(), success, worklet));
default:
success = false;
worklet.RaiseError("Unknown cell shape sent to world 2 parametric.");
return typename WorldCoordVector::ComponentType();
}
return result;
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector&,
const typename WorldCoordVector::ComponentType&,
vtkm::CellShapeTagEmpty,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
worklet.RaiseError("Attempted to find point coordinates in empty cell.");
success = false;
return typename WorldCoordVector::ComponentType();
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType&,
vtkm::CellShapeTagVertex,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
{
(void)pointWCoords; // Silence compiler warnings.
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 1);
success = true;
return typename WorldCoordVector::ComponentType(0, 0, 0);
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagLine,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 2);
success = true;
// Because this is a line, there is only one vaild parametric coordinate. Let
// vec be the vector from the first point to the second point
// (pointWCoords[1] - pointWCoords[0]), which is the direction of the line.
// dot(vec,wcoords-pointWCoords[0])/mag(vec) is the orthoginal projection of
// wcoords on the line and represents the distance between the orthoginal
// projection and pointWCoords[0]. The parametric coordinate is the fraction
// of this over the length of the segment, which is mag(vec). Thus, the
// parametric coordinate is dot(vec,wcoords-pointWCoords[0])/mag(vec)^2.
using Vector3 = typename WorldCoordVector::ComponentType;
using T = typename Vector3::ComponentType;
Vector3 vec = pointWCoords[1] - pointWCoords[0];
T numerator = vtkm::dot(vec, wcoords - pointWCoords[0]);
T denominator = vtkm::MagnitudeSquared(vec);
return Vector3(numerator / denominator, 0, 0);
}
static inline VTKM_EXEC vtkm::Vec<vtkm::FloatDefault, 3> WorldCoordinatesToParametricCoordinates(
const vtkm::VecAxisAlignedPointCoordinates<1>& pointWCoords,
const vtkm::VecAxisAlignedPointCoordinates<2>& pointWCoords,
const vtkm::Vec<vtkm::FloatDefault, 3>& wcoords,
vtkm::CellShapeTagLine,
vtkm::CellShapeTagQuad,
bool& success,
const FunctorBase&)
{
@ -820,14 +768,16 @@ template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagTriangle,
vtkm::CellShapeTagQuad,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
const vtkm::exec::FunctorBase& worklet)
{
success = true;
return vtkm::exec::internal::ReverseInterpolateTriangle(pointWCoords, wcoords);
vtkm::exec::internal::FastVec<WorldCoordVector, 4> local(pointWCoords);
return detail::WorldCoordinatesToParametricCoordinatesQuad(
local.Get(), wcoords, success, worklet);
}
//-----------------------------------------------------------------------------
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
@ -951,57 +901,16 @@ WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
triangleWCoords, trianglePCoords, vtkm::CellShapeTagTriangle(), worklet);
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagQuad,
bool& success,
const vtkm::exec::FunctorBase& worklet)
//-----------------------------------------------------------------------------
namespace detail
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 4);
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector2 = vtkm::Vec<T, 2>;
using Vector3 = vtkm::Vec<T, 3>;
// We have an underdetermined system in 3D, so create a 2D space in the
// plane that the polygon sits.
vtkm::exec::internal::Space2D<T> space(pointWCoords[0], pointWCoords[1], pointWCoords[3]);
auto result = vtkm::NewtonsMethod(
detail::JacobianFunctorQuad<WorldCoordVector, vtkm::CellShapeTagQuad>(&pointWCoords, &space),
detail::CoordinatesFunctorQuad<WorldCoordVector, vtkm::CellShapeTagQuad>(
&pointWCoords, &space, &worklet),
space.ConvertCoordToSpace(wcoords),
Vector2(0.5f, 0.5f));
success = result.Valid;
return Vector3(result.Solution[0], result.Solution[1], 0);
}
static inline VTKM_EXEC vtkm::Vec<vtkm::FloatDefault, 3> WorldCoordinatesToParametricCoordinates(
const vtkm::VecAxisAlignedPointCoordinates<2>& pointWCoords,
const vtkm::Vec<vtkm::FloatDefault, 3>& wcoords,
vtkm::CellShapeTagQuad,
bool& success,
const FunctorBase&)
{
success = true;
return (wcoords - pointWCoords.GetOrigin()) / pointWCoords.GetSpacing();
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagTetra,
bool& success,
const vtkm::exec::FunctorBase& vtkmNotUsed(worklet))
WorldCoordinatesToParametricCoordinatesTetra(
const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords)
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 4);
success = true;
// We solve the world to parametric coordinates problem for tetrahedra
// similarly to that for triangles. Before understanding this code, you
// should understand the triangle code (in ReverseInterpolateTriangle in
@ -1025,16 +934,13 @@ WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
// d = dot((wcoords - p0), planeNormal)/dot((p1-p0), planeNormal)
//
using Vector3 = typename WorldCoordVector::ComponentType;
const auto vec0 = pointWCoords[1] - pointWCoords[0];
const auto vec1 = pointWCoords[2] - pointWCoords[0];
const auto vec2 = pointWCoords[3] - pointWCoords[0];
const auto coordVec = wcoords - pointWCoords[0];
Vector3 pcoords;
const Vector3 vec0 = pointWCoords[1] - pointWCoords[0];
const Vector3 vec1 = pointWCoords[2] - pointWCoords[0];
const Vector3 vec2 = pointWCoords[3] - pointWCoords[0];
const Vector3 coordVec = wcoords - pointWCoords[0];
Vector3 planeNormal = vtkm::Cross(vec1, vec2);
typename WorldCoordVector::ComponentType pcoords;
auto planeNormal = vtkm::Cross(vec1, vec2);
pcoords[0] = vtkm::dot(coordVec, planeNormal) / vtkm::dot(vec0, planeNormal);
planeNormal = vtkm::Cross(vec0, vec2);
@ -1045,21 +951,94 @@ WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
return pcoords;
}
} // detail
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagHexahedron,
vtkm::CellShapeTagTetra,
bool& success,
const vtkm::exec::FunctorBase& worklet)
const vtkm::exec::FunctorBase&)
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 8);
return detail::WorldCoordinatesToParametricCoordinates3D(
pointWCoords, wcoords, vtkm::CellShapeTagHexahedron(), success, worklet);
success = true;
vtkm::exec::internal::FastVec<WorldCoordVector, 4> local(pointWCoords);
return detail::WorldCoordinatesToParametricCoordinatesTetra(local.Get(), wcoords);
}
//-----------------------------------------------------------------------------
namespace detail
{
template <typename WorldCoordVector, typename CellShapeTag>
class JacobianFunctor3DCell
{
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector3 = vtkm::Vec<T, 3>;
using Matrix3x3 = vtkm::Matrix<T, 3, 3>;
const WorldCoordVector* PointWCoords;
public:
VTKM_EXEC
JacobianFunctor3DCell(const WorldCoordVector* pointWCoords)
: PointWCoords(pointWCoords)
{
}
VTKM_EXEC
Matrix3x3 operator()(const Vector3& pcoords) const
{
Matrix3x3 jacobian;
vtkm::exec::JacobianFor3DCell(*this->PointWCoords, pcoords, jacobian, CellShapeTag());
return jacobian;
}
};
template <typename WorldCoordVector, typename CellShapeTag>
class CoordinatesFunctor3DCell
{
using T = typename WorldCoordVector::ComponentType::ComponentType;
using Vector3 = vtkm::Vec<T, 3>;
const WorldCoordVector* PointWCoords;
const vtkm::exec::FunctorBase* Worklet;
public:
VTKM_EXEC
CoordinatesFunctor3DCell(const WorldCoordVector* pointWCoords,
const vtkm::exec::FunctorBase* worklet)
: PointWCoords(pointWCoords)
, Worklet(worklet)
{
}
VTKM_EXEC
Vector3 operator()(Vector3 pcoords) const
{
return vtkm::exec::ParametricCoordinatesToWorldCoordinates(
*this->PointWCoords, pcoords, CellShapeTag(), *this->Worklet);
}
};
template <typename WorldCoordVector, typename CellShapeTag>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates3D(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
CellShapeTag,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
auto result = vtkm::NewtonsMethod(
JacobianFunctor3DCell<WorldCoordVector, CellShapeTag>(&pointWCoords),
CoordinatesFunctor3DCell<WorldCoordVector, CellShapeTag>(&pointWCoords, &worklet),
wcoords,
typename WorldCoordVector::ComponentType(0.5f, 0.5f, 0.5f));
success = result.Valid;
return result.Solution;
}
} // detail
static inline VTKM_EXEC vtkm::Vec<vtkm::FloatDefault, 3> WorldCoordinatesToParametricCoordinates(
const vtkm::VecAxisAlignedPointCoordinates<3>& pointWCoords,
const vtkm::Vec<vtkm::FloatDefault, 3>& wcoords,
@ -1071,6 +1050,19 @@ static inline VTKM_EXEC vtkm::Vec<vtkm::FloatDefault, 3> WorldCoordinatesToParam
return (wcoords - pointWCoords.GetOrigin()) / pointWCoords.GetSpacing();
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagHexahedron,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
vtkm::exec::internal::FastVec<WorldCoordVector, 8> local(pointWCoords);
return detail::WorldCoordinatesToParametricCoordinates3D(
local.Get(), wcoords, vtkm::CellShapeTagHexahedron(), success, worklet);
}
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
@ -1079,10 +1071,9 @@ WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 6);
vtkm::exec::internal::FastVec<WorldCoordVector, 6> local(pointWCoords);
return detail::WorldCoordinatesToParametricCoordinates3D(
pointWCoords, wcoords, vtkm::CellShapeTagWedge(), success, worklet);
local.Get(), wcoords, vtkm::CellShapeTagWedge(), success, worklet);
}
template <typename WorldCoordVector>
@ -1093,10 +1084,32 @@ WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
VTKM_ASSERT(pointWCoords.GetNumberOfComponents() == 5);
vtkm::exec::internal::FastVec<WorldCoordVector, 5> local(pointWCoords);
return detail::WorldCoordinatesToParametricCoordinates3D(
pointWCoords, wcoords, vtkm::CellShapeTagPyramid(), success, worklet);
local.Get(), wcoords, vtkm::CellShapeTagPyramid(), success, worklet);
}
//-----------------------------------------------------------------------------
template <typename WorldCoordVector>
static inline VTKM_EXEC typename WorldCoordVector::ComponentType
WorldCoordinatesToParametricCoordinates(const WorldCoordVector& pointWCoords,
const typename WorldCoordVector::ComponentType& wcoords,
vtkm::CellShapeTagGeneric shape,
bool& success,
const vtkm::exec::FunctorBase& worklet)
{
typename WorldCoordVector::ComponentType result;
switch (shape.Id)
{
vtkmGenericCellShapeMacro(result = WorldCoordinatesToParametricCoordinates(
pointWCoords, wcoords, CellShapeTag(), success, worklet));
default:
success = false;
worklet.RaiseError("Unknown cell shape sent to world 2 parametric.");
return typename WorldCoordVector::ComponentType();
}
return result;
}
}
} // namespace vtkm::exec

Some files were not shown because too many files have changed in this diff Show More