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

Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m

Browse Source
This commit is contained in:
T.J. Corona 2015-11-20 13:00:31 -05:00 committed by jimmy
commit 5a4c357c6e
84 changed files with 3953 additions and 1810 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

48
CMake/UseVTKmBase.cmake Normal file
View File

@ -0,0 +1,48 @@
##============================================================================
## 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 2015 Sandia Corporation.
## Copyright 2015 UT-Battelle, LLC.
## Copyright 2015 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## 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.
##============================================================================
# This is the base configuration for using the VTK-m library. All other
# device configurations rely on this.
if (VTKm_Base_initialize_complete)
return()
endif (VTKm_Base_initialize_complete)
# Find the Boost library.
if(NOT Boost_FOUND)
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION})
endif()
if (NOT Boost_FOUND)
message(STATUS "Boost not found")
set(VTKm_Base_FOUND)
else()
set(VTKm_Base_FOUND TRUE)
endif ()
# Set up all these dependent packages (if they were all found).
if (VTKm_Base_FOUND)
set(VTKm_INCLUDE_DIRS
${VTKm_INCLUDE_DIRS}
${Boost_INCLUDE_DIRS}
)
set(VTKm_Base_initialize_complete TRUE)
endif (VTKm_Base_FOUND)

78
CMake/UseVTKmCUDA.cmake Normal file
View File

@ -0,0 +1,78 @@
##============================================================================
## 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 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## 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.
##============================================================================
if (VTKm_CUDA_initialize_complete)
return()
endif (VTKm_CUDA_initialize_complete)
vtkm_configure_device(Base)
if (VTKm_Base_FOUND)
set(VTKm_CUDA_FOUND ${VTKm_ENABLE_CUDA})
if (NOT VTKm_CUDA_FOUND)
message(STATUS "This build of VTK-m does not include CUDA.")
endif ()
#---------------------------------------------------------------------------
# Find CUDA library.
#---------------------------------------------------------------------------
if (VTKm_CUDA_FOUND)
find_package(CUDA)
mark_as_advanced(CUDA_BUILD_CUBIN
CUDA_BUILD_EMULATION
CUDA_HOST_COMPILER
CUDA_SDK_ROOT_DIR
CUDA_SEPARABLE_COMPILATION
CUDA_TOOLKIT_ROOT_DIR
CUDA_VERBOSE_BUILD
)
if (NOT CUDA_FOUND)
message(STATUS "CUDA not found")
set(VTKm_CUDA_FOUND)
endif ()
endif ()
#---------------------------------------------------------------------------
# Find Thrust library.
#---------------------------------------------------------------------------
if (VTKm_CUDA_FOUND)
find_package(Thrust)
if (NOT THRUST_FOUND)
message(STATUS "Thrust not found")
set(VTKm_CUDA_FOUND)
endif ()
endif ()
endif () # VTKm_Base_FOUND
#-----------------------------------------------------------------------------
# Set up all these dependent packages (if they were all found).
#-----------------------------------------------------------------------------
if (VTKm_CUDA_FOUND)
set(VTKm_INCLUDE_DIRS
${VTKm_INCLUDE_DIRS}
${THRUST_INCLUDE_DIRS}
)
set(VTKm_CUDA_initialize_complete TRUE)
endif (VTKm_CUDA_FOUND)

84
CMake/UseVTKmCuda.cmake
View File

@ -1,84 +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 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## 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.
##============================================================================
if (VTKm_Cuda_initialize_complete)
return()
endif (VTKm_Cuda_initialize_complete)
set(VTKm_Cuda_FOUND ${VTKm_ENABLE_CUDA})
if (NOT VTKm_Cuda_FOUND)
message(STATUS "This build of VTKm does not include Cuda.")
endif (NOT VTKm_Cuda_FOUND)
# Find the Boost library.
if (VTKm_Cuda_FOUND)
if(NOT Boost_FOUND)
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION})
endif()
if (NOT Boost_FOUND)
message(STATUS "Boost not found")
set(VTKm_Cuda_FOUND)
endif (NOT Boost_FOUND)
endif (VTKm_Cuda_FOUND)
#-----------------------------------------------------------------------------
# Find CUDA library.
#-----------------------------------------------------------------------------
if (VTKm_Cuda_FOUND)
find_package(CUDA)
mark_as_advanced(CUDA_BUILD_CUBIN
CUDA_BUILD_EMULATION
CUDA_HOST_COMPILER
CUDA_SDK_ROOT_DIR
CUDA_SEPARABLE_COMPILATION
CUDA_TOOLKIT_ROOT_DIR
CUDA_VERBOSE_BUILD
)
if (NOT CUDA_FOUND)
message(STATUS "CUDA not found")
set(VTKm_Cuda_FOUND)
endif (NOT CUDA_FOUND)
endif ()
#-----------------------------------------------------------------------------
# Find Thrust library.
#-----------------------------------------------------------------------------
if (VTKm_Cuda_FOUND)
find_package(Thrust)
if (NOT THRUST_FOUND)
message(STATUS "Thrust not found")
set(VTKm_Cuda_FOUND)
endif (NOT THRUST_FOUND)
endif ()
#-----------------------------------------------------------------------------
# Set up all these dependent packages (if they were all found).
#-----------------------------------------------------------------------------
if (VTKm_Cuda_FOUND)
cuda_include_directories(
${Boost_INCLUDE_DIRS}
${THRUST_INCLUDE_DIRS}
${VTKm_INCLUDE_DIRS}
)
set(VTKm_Cuda_initialize_complete TRUE)
endif (VTKm_Cuda_FOUND)

27
CMake/UseVTKmSerial.cmake
View File

@ -22,26 +22,15 @@ if (VTKm_Serial_initialize_complete)
return()
endif (VTKm_Serial_initialize_complete)
# Find the Boost library.
if (NOT VTKm_Serial_FOUND)
if(NOT Boost_FOUND)
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION})
endif()
vtkm_configure_device(Base)
if (NOT Boost_FOUND)
message(STATUS "Boost not found")
set(VTKm_Serial_FOUND FALSE)
else(NOT Boost_FOUND)
set(VTKm_Serial_FOUND TRUE)
endif (NOT Boost_FOUND)
endif (NOT VTKm_Serial_FOUND)
if (VTKm_Base_FOUND)
# Serial only relies on base configuration
set(VTKm_Serial_FOUND TRUE)
else () # !VTKm_Base_FOUND
set(VTKm_Serial_FOUND)
endif ()
# Set up all these dependent packages (if they were all found).
if (VTKm_Serial_FOUND)
include_directories(
${Boost_INCLUDE_DIRS}
${VTKm_INCLUDE_DIRS}
)
set(VTKm_Serial_initialize_complete TRUE)
endif (VTKm_Serial_FOUND)
endif ()

43
CMake/UseVTKmTBB.cmake
View File

@ -22,36 +22,37 @@ if (VTKm_TBB_initialize_complete)
return()
endif (VTKm_TBB_initialize_complete)
#-----------------------------------------------------------------------------
# Find TBB.
#-----------------------------------------------------------------------------
if (NOT VTKm_TBB_FOUND)
find_package(TBB REQUIRED)
set (VTKm_TBB_FOUND ${TBB_FOUND})
endif()
vtkm_configure_device(Base)
#-----------------------------------------------------------------------------
# Find the Boost library.
#-----------------------------------------------------------------------------
if (VTKm_TBB_FOUND)
if(NOT Boost_FOUND)
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION})
if (VTKm_Base_FOUND)
set(VTKm_TBB_FOUND ${VTKm_ENABLE_TBB})
if (NOT VTKm_TBB_FOUND)
message(STATUS "This build of VTK-m does not include TBB.")
endif ()
#---------------------------------------------------------------------------
# Find TBB.
#---------------------------------------------------------------------------
if (VTKm_TBB_FOUND)
find_package(TBB)
if (NOT TBB_FOUND)
message(STATUS "TBB not found")
set(VTKm_TBB_FOUND)
endif ()
endif()
if (NOT Boost_FOUND)
message(STATUS "Boost not found")
set(VTKm_TBB_FOUND FALSE)
endif()
endif()
endif ()
#-----------------------------------------------------------------------------
# Set up all these dependent packages (if they were all found).
#-----------------------------------------------------------------------------
if (VTKm_TBB_FOUND)
include_directories(
${Boost_INCLUDE_DIRS}
set(VTKm_INCLUDE_DIRS
${VTKm_INCLUDE_DIRS}
${TBB_INCLUDE_DIRS}
)
)
set(VTKm_LIBRARIES ${TBB_LIBRARIES})
set(VTKm_TBB_initialize_complete TRUE)
endif()

27
CMake/VTKmConfig.cmake.in
View File

@ -18,6 +18,12 @@
## this software.
##============================================================================
# When this file is run by CMake (usually through the find_package command),
# The following variables will be defined:
#
# VTKm_INCLUDE_DIRS - Directories containing VTK-m and dependent headers
# VTKm_LIBRARIES - The libraries required when using VTK-m
# This file should be installed in the include directory.
# Find the root directory.
get_filename_component(_dir "${CMAKE_CURRENT_LIST_FILE}" PATH)
@ -35,6 +41,27 @@ set(VTKm_REQUIRED_BOOST_VERSION "@VTKm_REQUIRED_BOOST_VERSION@")
set(VTKm_CMAKE_MODULE_PATH "@VTKm_CMAKE_MODULE_PATH_CONFIG@")
set(VTKm_ENABLE_CUDA "@VTKm_ENABLE_CUDA@")
set(VTKm_ENABLE_TBB "@VTKm_ENABLE_TBB@")
# VTKm requires some CMake Find modules not included with CMake, so
# include the CMake modules distributed with VTKm.
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${VTKm_CMAKE_MODULE_PATH})
include(VTKmMacros)
if(VTKm_FIND_REQUIRED)
set(vtkm_is_required REQUIRED)
else()
set(vtkm_is_required)
endif()
vtkm_configure_device(Serial ${vtkm_is_required})
if(VTKm_ENABLE_CUDA)
vtkm_configure_device(CUDA)
endif()
if(VTKm_ENABLE_TBB)
# If configured with TBB, the VTK-m header files might choose TBB as the
# default device. In that case, the TBB device is required.
vtkm_configure_device(TBB ${vtkm_is_required})
endif()

20
CMake/VTKmMacros.cmake
View File

@ -223,7 +223,7 @@ function(vtkm_unit_tests)
)
#set up what we possibly need to link too.
list(APPEND VTKm_UT_LIBRARIES ${TBB_LIBRARIES})
list(APPEND VTKm_UT_LIBRARIES ${VTKm_LIBRARIES})
#set up storage for the include dirs
set(VTKm_UT_INCLUDE_DIRS )
@ -374,9 +374,7 @@ function(vtkm_worklet_unit_tests device_adapter)
set(CUDA_NVCC_FLAGS ${old_nvcc_flags} )
else()
add_executable(${test_prog} ${unit_test_drivers} ${unit_test_srcs})
if("${device_adapter}" STREQUAL "VTKM_DEVICE_ADAPTER_TBB")
target_link_libraries(${test_prog} ${TBB_LIBRARIES})
endif()
target_link_libraries(${test_prog} ${VTKm_LIBRARIES})
endif()
#add a test for each worklet test file. We will inject the device
@ -421,7 +419,7 @@ function(vtkm_save_benchmarks)
#create the benchmarks driver when we are called, since
#the driver expects the files to be in the same
#directory as the test driver
#TODO: This is probably ok to use for benchmarks as well
#TODO: This is probably ok to use for benchmarks as well
create_test_sourcelist(bench_sources BenchmarkDriver.cxx ${ARGN})
#store the absolute path for the driver and all the test
@ -501,9 +499,7 @@ function(vtkm_benchmarks device_adapter)
set(CUDA_NVCC_FLAGS ${old_nvcc_flags})
else()
add_executable(${benchmark_prog} ${benchmark_drivers} ${benchmark_srcs})
if("${device_adapter}" STREQUAL "VTKM_DEVICE_ADAPTER_TBB")
target_link_libraries(${benchmark_prog} ${TBB_LIBRARIES})
endif()
target_link_libraries(${benchmark_prog} ${VTKm_LIBRARIES})
endif()
if(MSVC)
@ -553,7 +549,11 @@ macro(vtkm_configure_device device)
set(VTKm_ENABLE_${device_uppercase} ON)
include("UseVTKm${device}")
if(NOT VTKm_${device}_FOUND)
message(SEND_ERROR "Could not configure for using VTKm with ${device}")
endif(NOT VTKm_${device}_FOUND)
if ("${ARGV1}" STREQUAL "REQUIRED")
message(SEND_ERROR "Could not configure for using VTKm with ${device}")
else()
message(STATUS "Could not configure for using VTKm with ${device}")
endif()
endif()
endmacro(vtkm_configure_device)

10
CMakeLists.txt
View File

@ -112,12 +112,14 @@ endif()
#-----------------------------------------------------------------------------
# Set up devices selected.
vtkm_configure_device(Serial)
if (VTKm_ENABLE_CUDA)
vtkm_configure_device(Cuda)
endif (VTKm_ENABLE_CUDA)
if (VTKm_ENABLE_TBB)
vtkm_configure_device(TBB)
vtkm_configure_device(TBB REQUIRED)
endif (VTKm_ENABLE_TBB)
if (VTKm_ENABLE_CUDA)
vtkm_configure_device(CUDA REQUIRED)
cuda_include_directories(${VTKm_INCLUDE_DIRS})
endif (VTKm_ENABLE_CUDA)
include_directories(${VTKm_INCLUDE_DIRS})
#-----------------------------------------------------------------------------

4
examples/clipping/CMakeLists.txt
View File

@ -42,8 +42,8 @@ endif()
if(VTKm_ENABLE_TBB)
add_executable(Clipping_TBB Clipping.cxx)
target_include_directories(Clipping_TBB PRIVATE ${TBB_INCLUDE_DIRS})
target_link_libraries(Clipping_TBB ${TBB_LIBRARIES})
target_include_directories(Clipping_TBB PRIVATE ${VTKm_INCLUDE_DIRS})
target_link_libraries(Clipping_TBB ${VTKm_LIBRARIES})
set_property(
TARGET Clipping_TBB
APPEND

257
examples/clipping/Clipping.cxx
View File

@ -21,6 +21,8 @@
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/Timer.h>
#include <vtkm/io/reader/VTKDataSetReader.h>
#include <vtkm/io/writer/VTKDataSetWriter.h>
#include <vtkm/worklet/Clip.h>
@ -36,223 +38,14 @@ VTKM_THIRDPARTY_POST_INCLUDE
typedef vtkm::Vec<vtkm::Float32, 3> FloatVec3;
template<typename T>
inline void flipEndianness(T* buffer, vtkm::Id size)
{
for (vtkm::Id i = 0; i < size; ++i)
{
T val = buffer[i];
vtkm::UInt8 *bytes = reinterpret_cast<vtkm::UInt8*>(&val);
std::reverse(bytes, bytes + sizeof(T));
buffer[i] = val;
}
}
template<typename T, typename DeviceAdapter>
inline vtkm::cont::ArrayHandle<vtkm::Float32> LoadBinaryPointDataImpl(
std::ifstream &fstream, vtkm::Id numPoints, T, DeviceAdapter)
{
vtkm::cont::ArrayHandle<vtkm::Float32> result;
std::vector<T> buffer(static_cast<size_t>(numPoints));
fstream.read(reinterpret_cast<char*>(&buffer[0]),
numPoints * static_cast<vtkm::Id>(sizeof(T)));
flipEndianness(&buffer[0], numPoints);
vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>::Copy(
vtkm::cont::make_ArrayHandleCast(vtkm::cont::make_ArrayHandle(buffer), vtkm::Float32()),
result);
return result;
}
template<typename DeviceAdapter>
inline vtkm::cont::ArrayHandle<vtkm::Float32> LoadBinaryPointData(
std::ifstream &fstream, vtkm::Id numPoints, std::string type, DeviceAdapter)
{
if (type == "short")
{
return LoadBinaryPointDataImpl(fstream, numPoints, short(), DeviceAdapter());
}
else if (type == "float")
{
return LoadBinaryPointDataImpl(fstream, numPoints, vtkm::Float32(), DeviceAdapter());
}
else
{
throw std::runtime_error("only short and float types supported");
}
return vtkm::cont::ArrayHandle<vtkm::Float32>();
}
template <typename DeviceAdapter>
vtkm::cont::DataSet LoadVtkLegacyStructuredPoints(const char *fname, DeviceAdapter)
{
vtkm::Id3 dim;
FloatVec3 spacing, origin;
vtkm::cont::ArrayHandle<vtkm::Float32> scalars;
std::ifstream vtkfile;
vtkfile.open(fname);
std::string tag;
std::getline(vtkfile, tag); // version comment
std::getline(vtkfile, tag); // datset name
vtkfile >> tag;
if (tag != "BINARY")
{
throw std::runtime_error("only binary format supported");
}
for (;;)
{
vtkfile >> tag;
if (tag == "DATASET")
{
std::string dataset;
vtkfile >> dataset;
if (dataset != "STRUCTURED_POINTS")
{
throw std::runtime_error("expecting structured dataset");
}
}
else if (tag == "DIMENSIONS")
{
vtkfile >> dim[0] >> dim[1] >> dim[2];
}
else if (tag == "SPACING")
{
vtkfile >> spacing[0] >> spacing[1] >> spacing[2];
}
else if (tag == "ORIGIN")
{
vtkfile >> origin[0] >> origin[1] >> origin[2];
}
else if (tag == "POINT_DATA")
{
vtkm::Id numPoints;
std::string type;
vtkfile >> numPoints;
vtkfile >> tag;
if (tag != "SCALARS")
{
throw std::runtime_error("only scalars supported for point data");
}
vtkfile >> tag >> type >> std::ws;
std::getline(vtkfile, tag); // LOOKUP_TABLE default
scalars = LoadBinaryPointData(vtkfile, numPoints, type, DeviceAdapter());
break;
}
}
vtkfile.close();
vtkm::cont::CellSetStructured<3> cs("cells");
cs.SetPointDimensions(vtkm::make_Vec(dim[0], dim[1], dim[2]));
vtkm::cont::DataSet ds;
ds.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", 1, dim, origin,
spacing));
ds.AddField(vtkm::cont::Field("scalars", 1, vtkm::cont::Field::ASSOC_POINTS,
scalars));
ds.AddCellSet(cs);
return ds;
}
void WriteVtkLegacyUnstructuredGrid(const char *fname, const vtkm::cont::DataSet &ds)
{
std::ofstream vtkfile;
vtkfile.open(fname);
vtkfile << "# vtk DataFile Version 3.0" << std::endl;
vtkfile << "vtkm_clip_output" << std::endl;
vtkfile << "BINARY" << std::endl << "DATASET UNSTRUCTURED_GRID" << std::endl;
vtkm::cont::CoordinateSystem coords = ds.GetCoordinateSystem();
vtkm::Id numPoints = coords.GetData().GetNumberOfValues();
vtkfile << "POINTS " << numPoints << " float" << std::endl;
{
vtkm::cont::ArrayHandle<FloatVec3> coordinates =
coords.GetData().CastToArrayHandle(FloatVec3(), VTKM_DEFAULT_STORAGE_TAG());
std::vector<FloatVec3> buffer(static_cast<size_t>(numPoints));
std::copy(vtkm::cont::ArrayPortalToIteratorBegin(coordinates.GetPortalConstControl()),
vtkm::cont::ArrayPortalToIteratorEnd(coordinates.GetPortalConstControl()),
buffer.begin());
flipEndianness(reinterpret_cast<vtkm::Float32*>(&buffer[0]), numPoints * 3);
vtkfile.write(reinterpret_cast<const char*>(&buffer[0]),
numPoints * static_cast<vtkm::Id>(sizeof(FloatVec3)));
}
vtkfile << std::endl;
vtkm::cont::CellSetExplicit<> cse =
ds.GetCellSet().CastTo<vtkm::cont::CellSetExplicit<> >();
vtkm::Id numCells = cse.GetNumberOfCells();
{
std::vector<int> idxBuffer, shapeBuffer;
idxBuffer.reserve(static_cast<size_t>(numCells * 4));
shapeBuffer.reserve(static_cast<size_t>(numCells));
vtkm::Vec<vtkm::Id, 8> pointIndices;
for (vtkm::Id i = 0; i < numCells; ++i)
{
vtkm::IdComponent numCellPoints = cse.GetNumberOfPointsInCell(i);
idxBuffer.push_back(static_cast<int>(numCellPoints));
cse.GetIndices(i, pointIndices);
for (vtkm::IdComponent j = 0; j < numCellPoints; ++j)
{
idxBuffer.push_back(static_cast<int>(pointIndices[j]));
}
shapeBuffer.push_back(static_cast<int>(cse.GetCellShape(i)));
}
vtkm::Id numIndices = static_cast<vtkm::Id>(idxBuffer.size());
vtkfile << "CELLS " << numCells << " " << numIndices << std::endl;
flipEndianness(&idxBuffer[0], numIndices);
vtkfile.write(reinterpret_cast<const char*>(&idxBuffer[0]),
numIndices * static_cast<vtkm::Id>(sizeof(idxBuffer[0])));
vtkfile << std::endl;
vtkfile << "CELL_TYPES " << numCells << std::endl;
flipEndianness(&shapeBuffer[0], numCells);
vtkfile.write(reinterpret_cast<const char*>(&shapeBuffer[0]),
numCells * static_cast<vtkm::Id>(sizeof(shapeBuffer[0])));
}
vtkfile << std::endl;
vtkm::cont::Field field = ds.GetField(0);
vtkfile << "POINT_DATA " << numPoints << std::endl
<< "SCALARS " << field.GetName() << " float" << std::endl
<< "LOOKUP_TABLE default" << std::endl;
{
vtkm::cont::ArrayHandle<vtkm::Float32> scalars =
field.GetData().CastToArrayHandle(vtkm::Float32(), VTKM_DEFAULT_STORAGE_TAG());
std::vector<vtkm::Float32> buffer(static_cast<size_t>(numPoints));
std::copy(vtkm::cont::ArrayPortalToIteratorBegin(scalars.GetPortalConstControl()),
vtkm::cont::ArrayPortalToIteratorEnd(scalars.GetPortalConstControl()),
buffer.begin());
flipEndianness(&buffer[0], numPoints);
vtkfile.write(reinterpret_cast<const char*>(&buffer[0]),
numPoints * static_cast<vtkm::Id>(sizeof(vtkm::Float32)));
}
vtkfile << std::endl;
vtkfile.close();
}
int main(int argc, char *argv[])
{
if (argc != 4)
if (argc < 4)
{
std::cout << "Usage: " << std::endl
<< "$ " << argv[0]
<< " <vtk_structure_points> <isoval> <vtk_unstructured_grid>"
<< " <input_vtk_file> [fieldName] <isoval> <output_vtk_file>"
<< std::endl;
return 1;
}
@ -262,24 +55,46 @@ int main(int argc, char *argv[])
<< vtkm::cont::internal::DeviceAdapterTraits<DeviceAdapter>::GetId()
<< std::endl;
vtkm::cont::DataSet input = LoadVtkLegacyStructuredPoints(argv[1], DeviceAdapter());
vtkm::io::reader::VTKDataSetReader reader(argv[1]);
vtkm::cont::DataSet input = reader.ReadDataSet();
vtkm::Float32 clipValue = boost::lexical_cast<vtkm::Float32>(argv[2]);
vtkm::cont::Field scalarField = (argc == 5) ?
input.GetField(argv[2]) :
input.GetField(0);
vtkm::Float32 clipValue = boost::lexical_cast<vtkm::Float32>(argv[argc - 2]);
vtkm::worklet::Clip<DeviceAdapter> clip;
vtkm::cont::Timer<DeviceAdapter> total;
vtkm::cont::Timer<DeviceAdapter> timer;
vtkm::cont::CellSetExplicit<> outputCellSet =
clip.Run(input.GetCellSet(0), input.GetField(0).GetData(), clipValue);
clip.Run(input.GetCellSet(0),
scalarField.GetData().ResetTypeList(vtkm::TypeListTagScalarAll()),
clipValue);
vtkm::Float64 clipTime = timer.GetElapsedTime();
vtkm::cont::DataSet output;
output.AddCellSet(outputCellSet);
timer.Reset();
vtkm::cont::DynamicArrayHandle coords =
clip.ProcessField(input.GetCoordinateSystem(0).GetData());
vtkm::Float64 processCoordinatesTime = timer.GetElapsedTime();
output.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", 1, coords));
timer.Reset();
vtkm::cont::DynamicArrayHandle scalars =
clip.ProcessField(input.GetField(0).GetData());
for (vtkm::Id i = 0; i < input.GetNumberOfFields(); ++i)
{
vtkm::cont::Field inField = input.GetField(i);
if (inField.GetAssociation() != vtkm::cont::Field::ASSOC_POINTS)
{
continue; // clip only supports point fields for now.
}
vtkm::cont::DynamicArrayHandle data =
clip.ProcessField(inField.GetData().ResetTypeList(vtkm::TypeListTagAll()));
output.AddField(vtkm::cont::Field(inField.GetName(), 1,
vtkm::cont::Field::ASSOC_POINTS, data));
}
vtkm::Float64 processScalarsTime = timer.GetElapsedTime();
vtkm::Float64 totalTime = total.GetElapsedTime();
@ -290,13 +105,9 @@ int main(int argc, char *argv[])
<< "process scalars: " << processScalarsTime << std::endl
<< "Total: " << totalTime << std::endl;
vtkm::cont::DataSet output;
output.AddField(vtkm::cont::Field("scalars", 1, vtkm::cont::Field::ASSOC_POINTS,
scalars));
output.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", 1, coords));
output.AddCellSet(outputCellSet);
WriteVtkLegacyUnstructuredGrid(argv[3], output);
std::ofstream outFile(argv[argc - 1]);
vtkm::io::writer::VTKDataSetWriter::Write(outFile, output);
outFile.close();
return 0;
}

6
examples/hello_world/CMakeLists.txt
View File

@ -24,18 +24,18 @@ if(OPENGL_FOUND AND GLUT_FOUND)
include_directories(${CMAKE_CURRENT_SOURCE_DIR})
add_executable(HelloWorld_SERIAL HelloWorld.cxx)
target_include_directories(HelloWorld_SERIAL PRIVATE ${OPENGL_INCLUDE_DIR} ${GLEW_INCLUDE_DIR} ${GLUT_INCLUDE_DIR})
target_link_libraries(HelloWorld_SERIAL ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(HelloWorld_SERIAL ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
if(VTKm_Cuda_FOUND)
vtkm_disable_troublesome_thrust_warnings()
cuda_add_executable(HelloWorld_CUDA HelloWorld.cu)
target_link_libraries(HelloWorld_CUDA ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(HelloWorld_CUDA ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
if(VTKm_TBB_FOUND)
add_executable(HelloWorld_TBB HelloWorld.cxx)
target_include_directories(HelloWorld_TBB PRIVATE ${OPENGL_INCLUDE_DIR} ${GLEW_INCLUDE_DIR} ${GLUT_INCLUDE_DIR})
target_link_libraries(HelloWorld_TBB ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_link_libraries(HelloWorld_TBB ${OPENGL_LIBRARIES} ${GLEW_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
endif()

17
examples/hello_world/HelloWorld.cxx
View File

@ -56,9 +56,10 @@ struct HelloVTKMInterop
vtkm::Vec< vtkm::Int32, 2 > Dims;
GLuint ProgramId;
GLuint VBOId;
GLuint VAOId;
GLuint ColorId;
vtkm::opengl::BufferState VBOState;
vtkm::opengl::BufferState ColorState;
vtkm::cont::Timer<DeviceAdapter> Timer;
@ -70,9 +71,9 @@ struct HelloVTKMInterop
HelloVTKMInterop(vtkm::Int32 width, vtkm::Int32 height):
Dims(256,256),
ProgramId(),
VBOId(),
VAOId(),
ColorId(),
VBOState(),
ColorState(),
Timer(),
InputData(),
InHandle(),
@ -119,11 +120,11 @@ struct HelloVTKMInterop
glUniformMatrix4fv( unifLoc, 1, GL_FALSE, mvp );
glEnableVertexAttribArray(0);
glBindBuffer(GL_ARRAY_BUFFER, this->VBOId);
glBindBuffer(GL_ARRAY_BUFFER, *this->VBOState.GetHandle());
glVertexAttribPointer( 0, 3, GL_FLOAT, GL_FALSE, 0, 0 );
glEnableClientState(GL_COLOR_ARRAY);
glBindBuffer(GL_ARRAY_BUFFER, this->ColorId);
glBindBuffer(GL_ARRAY_BUFFER, *this->ColorState.GetHandle());
glColorPointer(4, GL_UNSIGNED_BYTE, 0, 0 );
glDrawArrays( GL_POINTS, 0, arraySize );
@ -165,8 +166,8 @@ struct HelloVTKMInterop
GenerateSurfaceWorklet worklet( t );
DispatcherType(worklet).Invoke( this->InHandle, this->OutCoords, this->OutColors );
vtkm::opengl::TransferToOpenGL( this->OutCoords, this->VBOId, DeviceAdapter() );
vtkm::opengl::TransferToOpenGL( this->OutColors, this->ColorId, DeviceAdapter() );
vtkm::opengl::TransferToOpenGL( this->OutCoords, this->VBOState, DeviceAdapter() );
vtkm::opengl::TransferToOpenGL( this->OutColors, this->ColorState, DeviceAdapter() );
this->render();
if(t > 10)

8
examples/isosurface/CMakeLists.txt
View File

@ -23,18 +23,18 @@
if(OPENGL_FOUND AND GLUT_FOUND)
add_executable(IsosurfaceUniformGrid_SERIAL IsosurfaceUniformGrid.cxx)
target_include_directories(IsosurfaceUniformGrid_SERIAL PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(IsosurfaceUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(IsosurfaceUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
if(VTKm_Cuda_FOUND)
vtkm_disable_troublesome_thrust_warnings()
cuda_add_executable(IsosurfaceUniformGrid_CUDA IsosurfaceUniformGrid.cu)
target_link_libraries(IsosurfaceUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(IsosurfaceUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
if(VTKm_ENABLE_TBB)
add_executable(IsosurfaceUniformGrid_TBB IsosurfaceUniformGridTBB.cxx)
target_include_directories(IsosurfaceUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR} ${TBB_INCLUDE_DIRS})
target_link_libraries(IsosurfaceUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_include_directories(IsosurfaceUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(IsosurfaceUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
endif()

41
examples/isosurface/IsosurfaceUniformGrid.cxx
View File

@ -69,11 +69,11 @@ public:
typedef _1 InputDomain;
const vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::FloatDefault xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::Id cellsPerLayer;
VTKM_CONT_EXPORT
TangleField(const vtkm::Id3 dims, const float mins[3], const float maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
TangleField(const vtkm::Id3 dims, const vtkm::FloatDefault mins[3], const vtkm::FloatDefault maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { };
VTKM_EXEC_EXPORT
@ -83,9 +83,9 @@ public:
const vtkm::Id y = (vertexId / (xdim)) % (ydim);
const vtkm::Id z = vertexId / cellsPerLayer;
const float fx = static_cast<float>(x) / static_cast<float>(xdim-1);
const float fy = static_cast<float>(y) / static_cast<float>(xdim-1);
const float fz = static_cast<float>(z) / static_cast<float>(xdim-1);
const vtkm::FloatDefault fx = static_cast<vtkm::FloatDefault>(x) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fy = static_cast<vtkm::FloatDefault>(y) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fz = static_cast<vtkm::FloatDefault>(z) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::Float32 xx = 3.0f*(xmin+(xmax-xmin)*(fx));
const vtkm::Float32 yy = 3.0f*(ymin+(ymax-ymin)*(fy));
@ -103,15 +103,22 @@ vtkm::cont::DataSet MakeIsosurfaceTestDataSet(vtkm::Id3 dims)
const vtkm::Id3 vdims(dims[0] + 1, dims[1] + 1, dims[2] + 1);
float mins[3] = {-1.0f, -1.0f, -1.0f};
float maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::FloatDefault mins[3] = {-1.0f, -1.0f, -1.0f};
vtkm::FloatDefault maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray;
vtkm::cont::ArrayHandleCounting<vtkm::Id> vertexCountImplicitArray(0, 1, vdims[0]*vdims[1]*vdims[2]);
vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher(TangleField(vdims, mins, maxs));
tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray);
vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims);
vtkm::Vec<vtkm::FloatDefault,3> origin(0.0f, 0.0f, 0.0f);
vtkm::Vec<vtkm::FloatDefault,3> spacing(
1.0f/static_cast<vtkm::FloatDefault>(dims[0]),
1.0f/static_cast<vtkm::FloatDefault>(dims[2]),
1.0f/static_cast<vtkm::FloatDefault>(dims[1]));
vtkm::cont::ArrayHandleUniformPointCoordinates
coordinates(vdims, origin, spacing);
dataSet.AddCoordinateSystem(
vtkm::cont::CoordinateSystem("coordinates", 1, coordinates));
@ -135,9 +142,9 @@ void initializeGL()
glEnable(GL_DEPTH_TEST);
glShadeModel(GL_SMOOTH);
float white[] = { 0.8f, 0.8f, 0.8f, 1.0f };
float black[] = { 0.0f, 0.0f, 0.0f, 1.0f };
float lightPos[] = { 10.0f, 10.0f, 10.5f, 1.0f };
vtkm::FloatDefault white[] = { 0.8f, 0.8f, 0.8f, 1.0f };
vtkm::FloatDefault black[] = { 0.0f, 0.0f, 0.0f, 1.0f };
vtkm::FloatDefault lightPos[] = { 10.0f, 10.0f, 10.5f, 1.0f };
glLightfv(GL_LIGHT0, GL_AMBIENT, white);
glLightfv(GL_LIGHT0, GL_DIFFUSE, white);
@ -199,7 +206,7 @@ void mouseMove(int x, int y)
if (mouse_state == 0)
{
vtkm::Float32 pideg = static_cast<float>(vtkm::Pi()/180.0);
vtkm::Float32 pideg = static_cast<vtkm::Float32>(vtkm::Pi_2());
Quaternion newRotX;
newRotX.setEulerAngles(-0.2f*dx*pideg/180.0f, 0.0f, 0.0f);
qrot.mul(newRotX);
@ -243,6 +250,16 @@ int main(int argc, char* argv[])
std::cout << "Number of output vertices: " << verticesArray.GetNumberOfValues() << std::endl;
std::cout << "vertices: ";
vtkm::cont::printSummary_ArrayHandle(verticesArray, std::cout);
std::cout << std::endl;
std::cout << "normals: ";
vtkm::cont::printSummary_ArrayHandle(normalsArray, std::cout);
std::cout << std::endl;
std::cout << "scalars: ";
vtkm::cont::printSummary_ArrayHandle(scalarsArray, std::cout);
std::cout << std::endl;
lastx = lasty = 0;
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE | GLUT_DEPTH);

2
examples/multi_backend/CMakeLists.txt
View File

@ -27,5 +27,5 @@ else()
endif()
if(VTKm_TBB_FOUND)
target_link_libraries(MultiBackend ${TBB_LIBRARIES})
target_link_libraries(MultiBackend ${VTKm_LIBRARIES})
endif()

32
examples/tetrahedra/CMakeLists.txt
View File

@ -23,50 +23,50 @@
if(OPENGL_FOUND AND GLUT_FOUND)
add_executable(TetrahedralizeExplicitGrid_SERIAL TetrahedralizeExplicitGrid.cxx)
target_include_directories(TetrahedralizeExplicitGrid_SERIAL PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TetrahedralizeExplicitGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TetrahedralizeExplicitGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TriangulateExplicitGrid_SERIAL TriangulateExplicitGrid.cxx)
target_include_directories(TriangulateExplicitGrid_SERIAL PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TriangulateExplicitGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TriangulateExplicitGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TetrahedralizeUniformGrid_SERIAL TetrahedralizeUniformGrid.cxx)
target_include_directories(TetrahedralizeUniformGrid_SERIAL PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TetrahedralizeUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TetrahedralizeUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TriangulateUniformGrid_SERIAL TriangulateUniformGrid.cxx)
target_include_directories(TriangulateUniformGrid_SERIAL PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TriangulateUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TriangulateUniformGrid_SERIAL ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
if(VTKm_Cuda_FOUND)
cuda_add_executable(TetrahedralizeExplicitGrid_CUDA TetrahedralizeExplicitGrid.cu)
target_link_libraries(TetrahedralizeExplicitGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TetrahedralizeExplicitGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
cuda_add_executable(TriangulateExplicitGrid_CUDA TriangulateExplicitGrid.cu)
target_link_libraries(TriangulateExplicitGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TriangulateExplicitGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
cuda_add_executable(TetrahedralizeUniformGrid_CUDA TetrahedralizeUniformGrid.cu)
target_link_libraries(TetrahedralizeUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TetrahedralizeUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
cuda_add_executable(TriangulateUniformGrid_CUDA TriangulateUniformGrid.cu)
target_link_libraries(TriangulateUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES})
target_link_libraries(TriangulateUniformGrid_CUDA ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
if(VTKm_ENABLE_TBB)
add_executable(TetrahedralizeExplicitGrid_TBB TetrahedralizeExplicitGridTBB.cxx)
target_include_directories(TetrahedralizeExplicitGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR} ${TBB_INCLUDE_DIRS})
target_link_libraries(TetrahedralizeExplicitGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_include_directories(TetrahedralizeExplicitGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TetrahedralizeExplicitGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TriangulateExplicitGrid_TBB TriangulateExplicitGridTBB.cxx)
target_include_directories(TriangulateExplicitGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR} ${TBB_INCLUDE_DIRS})
target_link_libraries(TriangulateExplicitGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_include_directories(TriangulateExplicitGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TriangulateExplicitGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TetrahedralizeUniformGrid_TBB TetrahedralizeUniformGridTBB.cxx)
target_include_directories(TetrahedralizeUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR} ${TBB_INCLUDE_DIRS})
target_link_libraries(TetrahedralizeUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_include_directories(TetrahedralizeUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TetrahedralizeUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
add_executable(TriangulateUniformGrid_TBB TriangulateUniformGridTBB.cxx)
target_include_directories(TriangulateUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR} ${TBB_INCLUDE_DIRS})
target_link_libraries(TriangulateUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${TBB_LIBRARIES})
target_include_directories(TriangulateUniformGrid_TBB PRIVATE ${GLUT_INCLUDE_DIR} ${OPENGL_INCLUDE_DIR})
target_link_libraries(TriangulateUniformGrid_TBB ${OPENGL_LIBRARIES} ${GLUT_LIBRARIES} ${VTKm_LIBRARIES})
endif()
endif()

6
vtkm/Types.h
View File

@ -175,6 +175,12 @@ namespace internal {
//-----------------------------------------------------------------------------
/// Placeholder class for when a type is not applicable.
///
struct NullType { };
//-----------------------------------------------------------------------------
template<vtkm::IdComponent Size>
struct VecEquals
{

17
vtkm/cont/ArrayHandleCast.h
View File

@ -53,18 +53,19 @@ class ArrayHandleCast :
public vtkm::cont::ArrayHandleTransform<
T,
ArrayHandleType,
internal::Cast<typename ArrayHandleType::ValueType, T> >
internal::Cast<typename ArrayHandleType::ValueType, T>,
internal::Cast<T, typename ArrayHandleType::ValueType> >
{
public:
VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleCast,
(ArrayHandleCast<T, ArrayHandleType>),
(vtkm::cont::ArrayHandleTransform<
T,
ArrayHandleType,
internal::Cast<typename ArrayHandleType::ValueType, T> >));
ArrayHandleCast,
(ArrayHandleCast<T, ArrayHandleType>),
(vtkm::cont::ArrayHandleTransform<
T,
ArrayHandleType,
internal::Cast<typename ArrayHandleType::ValueType, T>,
internal::Cast<T, typename ArrayHandleType::ValueType> >));
VTKM_CONT_EXPORT
ArrayHandleCast(const ArrayHandleType &handle)
: Superclass(handle)
{ }

14
vtkm/cont/ArrayHandleGroupVec.h
View File

@ -345,6 +345,20 @@ public:
: Superclass(StorageType(sourceArray)) { }
};
/// \c make_ArrayHandleGroupVec is convenience function to generate an
/// ArrayHandleGroupVec. It takes in an ArrayHandle and the number of components
/// (as a specified template parameter), and returns an array handle with
/// consecutive entries grouped in a Vec.
///
template<vtkm::IdComponent NUM_COMPONENTS,
typename ArrayHandleType>
VTKM_CONT_EXPORT
vtkm::cont::ArrayHandleGroupVec<ArrayHandleType, NUM_COMPONENTS>
make_ArrayHandleGroupVec(const ArrayHandleType &array)
{
return vtkm::cont::ArrayHandleGroupVec<ArrayHandleType,NUM_COMPONENTS>(array);
}
}
} // namespace vtkm::cont

327
vtkm/cont/ArrayHandleTransform.h
View File

@ -27,20 +27,38 @@
#include <vtkm/cont/ErrorControlBadType.h>
#include <vtkm/cont/ErrorControlInternal.h>
namespace vtkm {
namespace cont {
namespace internal {
/// Tag used in place of an inverse functor.
struct NullFunctorType {};
}
}
} // namespace vtkm::cont::internal
namespace vtkm {
namespace exec {
namespace internal {
typedef vtkm::cont::internal::NullFunctorType NullFunctorType;
/// \brief An array portal that transforms a value from another portal.
///
template<typename ValueType_, typename PortalType_, typename FunctorType_,
typename InverseFunctorType_=NullFunctorType>
class ArrayPortalExecTransform;
template<typename ValueType_, typename PortalType_, typename FunctorType_>
class ArrayPortalExecTransform
class ArrayPortalExecTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType>
{
public:
typedef PortalType_ PortalType;
typedef ValueType_ ValueType;
typedef FunctorType_ FunctorType;
VTKM_CONT_EXPORT
VTKM_EXEC_CONT_EXPORT
ArrayPortalExecTransform(const PortalType &portal = PortalType(),
const FunctorType &functor = FunctorType())
: Portal(portal), Functor(functor)
@ -51,7 +69,7 @@ public:
/// type casting that the iterators do (like the non-const to const cast).
///
template<class OtherV, class OtherP, class OtherF>
VTKM_CONT_EXPORT
VTKM_EXEC_CONT_EXPORT
ArrayPortalExecTransform(const ArrayPortalExecTransform<OtherV,OtherP,OtherF> &src)
: Portal(src.GetPortal()),
Functor(src.GetFunctor())
@ -73,11 +91,49 @@ public:
VTKM_EXEC_CONT_EXPORT
const FunctorType &GetFunctor() const { return this->Functor; }
private:
protected:
PortalType Portal;
FunctorType Functor;
};
template<typename ValueType_, typename PortalType_,
typename FunctorType_, typename InverseFunctorType_>
class ArrayPortalExecTransform : public ArrayPortalExecTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType>
{
public:
typedef ArrayPortalExecTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType> Superclass;
typedef PortalType_ PortalType;
typedef ValueType_ ValueType;
typedef FunctorType_ FunctorType;
typedef InverseFunctorType_ InverseFunctorType;
VTKM_EXEC_CONT_EXPORT
ArrayPortalExecTransform(const PortalType &portal = PortalType(),
const FunctorType &functor = FunctorType(),
const InverseFunctorType& inverseFunctor = InverseFunctorType())
: Superclass(portal,functor), InverseFunctor(inverseFunctor)
{ }
template<class OtherV, class OtherP, class OtherF, class OtherInvF>
VTKM_EXEC_CONT_EXPORT
ArrayPortalExecTransform(const ArrayPortalExecTransform<OtherV,OtherP,OtherF,OtherInvF> &src)
: Superclass(src), InverseFunctor(src.GetInverseFunctor())
{ }
VTKM_EXEC_EXPORT
void Set(vtkm::Id index, const ValueType& value) const {
return this->Portal.Set(index,this->InverseFunctor(value));
}
VTKM_EXEC_CONT_EXPORT
const InverseFunctorType &GetInverseFunctor() const {
return this->InverseFunctor; }
private:
InverseFunctorType InverseFunctor;
};
}
}
} // namespace vtkm::exec::internal
@ -90,8 +146,12 @@ namespace internal {
/// \brief An array portal that transforms a value from another portal.
///
template<typename ValueType_, typename PortalType_, typename FunctorType_,
typename InverseFunctorType=NullFunctorType>
class ArrayPortalContTransform;
template<typename ValueType_, typename PortalType_, typename FunctorType_>
class ArrayPortalContTransform
class ArrayPortalContTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType>
{
public:
typedef PortalType_ PortalType;
@ -131,16 +191,53 @@ public:
VTKM_CONT_EXPORT
const FunctorType &GetFunctor() const { return this->Functor; }
private:
protected:
PortalType Portal;
FunctorType Functor;
};
template<typename ValueType, typename ArrayHandleType, typename FunctorType>
struct StorageTagTransform { };
template<typename ValueType_, typename PortalType_, typename FunctorType_,
typename InverseFunctorType_>
class ArrayPortalContTransform : public ArrayPortalContTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType>
{
public:
typedef ArrayPortalContTransform<ValueType_,PortalType_,FunctorType_,NullFunctorType> Superclass;
typedef PortalType_ PortalType;
typedef ValueType_ ValueType;
typedef FunctorType_ FunctorType;
typedef InverseFunctorType_ InverseFunctorType;
VTKM_CONT_EXPORT
ArrayPortalContTransform(const PortalType &portal = PortalType(),
const FunctorType &functor = FunctorType(),
const InverseFunctorType &inverseFunctor = InverseFunctorType())
: Superclass(portal,functor), InverseFunctor(inverseFunctor)
{ }
template<class OtherV, class OtherP, class OtherF, class OtherInvF>
VTKM_CONT_EXPORT
ArrayPortalContTransform(const ArrayPortalContTransform<OtherV,OtherP,OtherF,OtherInvF> &src)
: Superclass(src), InverseFunctor(src.GetInverseFunctor())
{ }
VTKM_CONT_EXPORT
void Set(vtkm::Id index, const ValueType& value) const {
this->Portal.Set(index,this->InverseFunctor(value));
}
VTKM_CONT_EXPORT
const InverseFunctorType &GetInverseFunctor() const { return this->InverseFunctor; }
private:
InverseFunctorType InverseFunctor;
};
template<typename ValueType, typename ArrayHandleType, typename FunctorType,
typename InverseFunctorType=NullFunctorType>
struct StorageTagTransform {};
template<typename T, typename ArrayHandleType, typename FunctorType>
class Storage<T, StorageTagTransform<T, ArrayHandleType, FunctorType > >
class Storage<T, StorageTagTransform<T, ArrayHandleType, FunctorType, NullFunctorType > >
{
public:
typedef T ValueType;
@ -221,6 +318,92 @@ private:
bool Valid;
};
template<typename T, typename ArrayHandleType, typename FunctorType,
typename InverseFunctorType>
class Storage<T,
StorageTagTransform<T, ArrayHandleType, FunctorType, InverseFunctorType> >
{
public:
typedef T ValueType;
typedef ArrayPortalContTransform<ValueType,
typename ArrayHandleType::PortalControl, FunctorType, InverseFunctorType>
PortalType;
typedef ArrayPortalContTransform<ValueType,
typename ArrayHandleType::PortalConstControl,FunctorType,InverseFunctorType>
PortalConstType;
VTKM_CONT_EXPORT
Storage() : Valid(false) { }
VTKM_CONT_EXPORT
Storage(const ArrayHandleType &array,
const FunctorType &functor,
const InverseFunctorType &inverseFunctor)
: Array(array), Functor(functor), InverseFunctor(inverseFunctor), Valid(true) { }
VTKM_CONT_EXPORT
PortalType GetPortal() {
VTKM_ASSERT_CONT(this->Valid);
return PortalType(this->Array.GetPortalControl(),
this->Functor,
this->InverseFunctor);
}
VTKM_CONT_EXPORT
PortalConstType GetPortalConst() const {
VTKM_ASSERT_CONT(this->Valid);
return PortalConstType(this->Array.GetPortalConstControl(),
this->Functor,
this->InverseFunctor);
}
VTKM_CONT_EXPORT
vtkm::Id GetNumberOfValues() const {
VTKM_ASSERT_CONT(this->Valid);
return this->Array.GetNumberOfValues();
}
VTKM_CONT_EXPORT
void Allocate(vtkm::Id numberOfValues) {
this->Array.Allocate(numberOfValues);
this->Valid = true;
}
VTKM_CONT_EXPORT
void Shrink(vtkm::Id numberOfValues) {
this->Array.Shrink(numberOfValues);
}
VTKM_CONT_EXPORT
void ReleaseResources() {
this->Array.ReleaseResources();
this->Valid = false;
}
VTKM_CONT_EXPORT
const ArrayHandleType &GetArray() const {
VTKM_ASSERT_CONT(this->Valid);
return this->Array;
}
VTKM_CONT_EXPORT
const FunctorType &GetFunctor() const {
return this->Functor;
}
VTKM_CONT_EXPORT
const InverseFunctorType &GetInverseFunctor() const {
return this->InverseFunctor;
}
private:
ArrayHandleType Array;
FunctorType Functor;
InverseFunctorType InverseFunctor;
bool Valid;
};
template<typename T,
typename ArrayHandleType,
typename FunctorType,
@ -295,6 +478,83 @@ private:
FunctorType Functor;
};
template<typename T,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType,
typename Device>
class ArrayTransfer<
T, StorageTagTransform<T,ArrayHandleType,FunctorType,InverseFunctorType>,
Device>
{
typedef StorageTagTransform<T,ArrayHandleType,
FunctorType,InverseFunctorType> StorageTag;
typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType;
public:
typedef T ValueType;
typedef typename StorageType::PortalType PortalControl;
typedef typename StorageType::PortalConstType PortalConstControl;
typedef vtkm::exec::internal::ArrayPortalExecTransform<
ValueType,
typename ArrayHandleType::template ExecutionTypes<Device>::Portal,
FunctorType, InverseFunctorType> PortalExecution;
typedef vtkm::exec::internal::ArrayPortalExecTransform<
ValueType,
typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst,
FunctorType, InverseFunctorType> PortalConstExecution;
VTKM_CONT_EXPORT
ArrayTransfer(StorageType *storage)
: Array(storage->GetArray()),
Functor(storage->GetFunctor()),
InverseFunctor(storage->GetInverseFunctor()) { }
VTKM_CONT_EXPORT
vtkm::Id GetNumberOfValues() const {
return this->Array.GetNumberOfValues();
}
VTKM_CONT_EXPORT
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) {
return PortalConstExecution(this->Array.PrepareForInput(Device()),this->Functor,this->InverseFunctor);
}
VTKM_CONT_EXPORT
PortalExecution PrepareForInPlace(bool &vtkmNotUsed(updateData)) {
return PortalExecution(this->Array.PrepareForInPlace(Device()),this->Functor,this->InverseFunctor);
}
VTKM_CONT_EXPORT
PortalExecution PrepareForOutput(vtkm::Id numberOfValues) {
return PortalExecution(this->Array.PrepareForOutput(numberOfValues,
Device()),this->Functor,this->InverseFunctor);
}
VTKM_CONT_EXPORT
void RetrieveOutputData(StorageType *vtkmNotUsed(storage)) const {
// Implementation of this method should be unnecessary. The internal
// array handle should automatically retrieve the output data as necessary.
}
VTKM_CONT_EXPORT
void Shrink(vtkm::Id numberOfValues) {
this->Array.Shrink(numberOfValues);
}
VTKM_CONT_EXPORT
void ReleaseResources() {
this->Array.ReleaseResourcesExecution();
}
private:
ArrayHandleType Array;
FunctorType Functor;
InverseFunctorType InverseFunctor;
};
} // namespace internal
/// \brief Implicitly transform values of one array to another with a functor.
@ -308,10 +568,16 @@ private:
/// the functor operator should work in both the control and execution
/// environments.
///
template <typename T,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType=internal::NullFunctorType>
class ArrayHandleTransform;
template <typename T,
typename ArrayHandleType,
typename FunctorType>
class ArrayHandleTransform
class ArrayHandleTransform<T,ArrayHandleType,FunctorType,internal::NullFunctorType>
: public vtkm::cont::ArrayHandle<
T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType> >
{
@ -327,7 +593,7 @@ public:
T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType> >));
private:
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType;
public:
VTKM_CONT_EXPORT
@ -348,6 +614,45 @@ make_ArrayHandleTransform(HandleType handle, FunctorType functor)
return ArrayHandleTransform<T,HandleType,FunctorType>(handle,functor);
}
// ArrayHandleTransform with inverse functors enabled (no need to subclass from
// ArrayHandleTransform without inverse functors: nothing to inherit).
template <typename T,
typename ArrayHandleType,
typename FunctorType,
typename InverseFunctorType>
class ArrayHandleTransform
: public vtkm::cont::ArrayHandle<
T,
internal::StorageTagTransform<T, ArrayHandleType, FunctorType,
InverseFunctorType> >
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
public:
VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleTransform,
(ArrayHandleTransform<T,ArrayHandleType,FunctorType,InverseFunctorType>),
(vtkm::cont::ArrayHandle<
T, internal::StorageTagTransform<T, ArrayHandleType, FunctorType,
InverseFunctorType> >));
private:
typedef vtkm::cont::internal::Storage<T, StorageTag> StorageType;
public:
ArrayHandleTransform(const ArrayHandleType &handle,
const FunctorType &functor = FunctorType(),
const InverseFunctorType &inverseFunctor = InverseFunctorType())
: Superclass(StorageType(handle, functor, inverseFunctor)) { }
};
template <typename T, typename HandleType, typename FunctorType, typename InverseFunctorType>
VTKM_CONT_EXPORT
vtkm::cont::ArrayHandleTransform<T, HandleType, FunctorType, InverseFunctorType>
make_ArrayHandleTransform(HandleType handle, FunctorType functor, InverseFunctorType inverseFunctor)
{
return ArrayHandleTransform<T,HandleType,FunctorType,InverseFunctorType>(handle,functor,inverseFunctor);
}
}
} // namespace vtkm::cont

5
vtkm/cont/cuda/DeviceAdapterCuda.h
View File

@ -23,6 +23,11 @@
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
#ifdef VTKM_CUDA
//This is required to be first so that we get patches for thrust included
//in the correct order
#include <vtkm/exec/cuda/internal/ThrustPatches.h>
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
#include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h>
#endif

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

@ -74,7 +74,7 @@ public:
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< const T > PortalConstType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< T > PortalConstType;
VTKM_CONT_EXPORT
ArrayManagerExecutionThrustDevice(StorageType *storage)

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

@ -40,7 +40,7 @@ void buildIndexOffsets(const NumIndicesArrayType& numIndices,
//We first need to make sure that NumIndices and IndexOffsetArrayType
//have the same type so we can call scane exclusive
typedef vtkm::cont::ArrayHandleCast< vtkm::Id,
NumIndicesArrayType > CastedNumIndicesType;
NumIndicesArrayType > CastedNumIndicesType;
// Although technically we are making changes to this object, the changes
// are logically consistent with the previous state, so we consider it

3
vtkm/cont/testing/TestingFancyArrayHandles.h
View File

@ -421,7 +421,8 @@ private:
typedef vtkm::cont::ArrayHandleIndex InputArrayType;
InputArrayType input(ARRAY_SIZE);
vtkm::cont::ArrayHandleCast<CastToType, InputArrayType> castArray =
vtkm::cont::ArrayHandleCast<CastToType,
InputArrayType> castArray =
vtkm::cont::make_ArrayHandleCast(input, CastToType());
vtkm::cont::ArrayHandle<CastToType> result;

18
vtkm/exec/CellDerivative.h
View File

@ -181,15 +181,15 @@ PermutePyramidToHex(const FieldVecType &field)
//
#define VTKM_ACCUM_JACOBIAN_3D(pointIndex, weight0, weight1, weight2) \
jacobian(0,0) += wCoords[pointIndex][0] * (weight0); \
jacobian(1,0) += wCoords[pointIndex][1] * (weight0); \
jacobian(2,0) += wCoords[pointIndex][2] * (weight0); \
jacobian(0,1) += wCoords[pointIndex][0] * (weight1); \
jacobian(1,1) += wCoords[pointIndex][1] * (weight1); \
jacobian(2,1) += wCoords[pointIndex][2] * (weight1); \
jacobian(0,2) += wCoords[pointIndex][0] * (weight2); \
jacobian(1,2) += wCoords[pointIndex][1] * (weight2); \
jacobian(2,2) += wCoords[pointIndex][2] * (weight2)
jacobian(0,0) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight0)); \
jacobian(1,0) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight0)); \
jacobian(2,0) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight0)); \
jacobian(0,1) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight1)); \
jacobian(1,1) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight1)); \
jacobian(2,1) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight1)); \
jacobian(0,2) += static_cast<JacobianType>(wCoords[pointIndex][0] * (weight2)); \
jacobian(1,2) += static_cast<JacobianType>(wCoords[pointIndex][1] * (weight2)); \
jacobian(2,2) += static_cast<JacobianType>(wCoords[pointIndex][2] * (weight2))
template<typename WorldCoordType,
typename ParametricCoordType,

2
vtkm/exec/arg/CMakeLists.txt
View File

@ -32,8 +32,10 @@ set(headers
FetchTagTopologyIn.h
FromCount.h
FromIndices.h
ThreadIndices.h
ThreadIndicesBasic.h
ThreadIndicesTopologyMap.h
VisitIndex.h
WorkIndex.h
)

2
vtkm/exec/arg/FetchTagArrayDirectIn.h
View File

@ -50,7 +50,7 @@ struct Fetch<
ValueType Load(const ThreadIndicesType &indices,
const ExecObjectType &arrayPortal) const
{
return arrayPortal.Get(indices.GetIndex());
return arrayPortal.Get(indices.GetInputIndex());
}
VTKM_EXEC_EXPORT

13
vtkm/exec/arg/FetchTagArrayDirectInOut.h
View File

@ -34,6 +34,15 @@ namespace arg {
/// indexing, so the thread index given to \c Store is used as the index into
/// the array.
///
/// When using \c FetchTagArrayDirectInOut with a worklet invocation with a
/// scatter, it is a bit undefined how the in/out array should be indexed.
/// Should it be the size of the input arrays and written back there, or
/// should it be the size of the output arrays and pre-filled with the output.
/// The implementation indexes based on the output because it is safer. The
/// output will have a unique index for each worklet instance, so you don't
/// have to worry about writes stomping on each other (which they would
/// inevitably do if index as input).
///
struct FetchTagArrayDirectInOut { };
@ -51,7 +60,7 @@ struct Fetch<
ValueType Load(const ThreadIndicesType &indices,
const ExecObjectType &arrayPortal) const
{
return arrayPortal.Get(indices.GetIndex());
return arrayPortal.Get(indices.GetOutputIndex());
}
VTKM_SUPPRESS_EXEC_WARNINGS
@ -60,7 +69,7 @@ struct Fetch<
const ExecObjectType &arrayPortal,
const ValueType &value) const
{
arrayPortal.Set(indices.GetIndex(), value);
arrayPortal.Set(indices.GetOutputIndex(), value);
}
};

2
vtkm/exec/arg/FetchTagArrayDirectOut.h
View File

@ -59,7 +59,7 @@ struct Fetch<
const ExecObjectType &arrayPortal,
const ValueType &value) const
{
arrayPortal.Set(indices.GetIndex(), value);
arrayPortal.Set(indices.GetOutputIndex(), value);
}
};

6
vtkm/exec/arg/FetchTagArrayTopologyMapIn.h
View File

@ -76,7 +76,11 @@ struct FetchArrayTopologyMapInImplementation
static ValueType Load(const ThreadIndicesType &indices,
const FieldExecObjectType &field)
{
return ValueType(indices.GetIndicesFrom(), field);
// It is important that we give the VecFromPortalPermute (ValueType) a
// pointer that will stay around during the time the Vec is valid. Thus, we
// should make sure that indices is a reference that goes up the stack at
// least as far as the returned VecFromPortalPermute is used.
return ValueType(indices.GetIndicesFromPointer(), field);
}
};

85
vtkm/exec/arg/ThreadIndices.h Normal file
View File

@ -0,0 +1,85 @@
//============================================================================
// 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 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_arg_ThreadIndices_h
#define vtk_m_exec_arg_ThreadIndices_h
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Aspect tag to use for getting the thread indices.
///
/// The \c AspectTagThreadIndices aspect tag causes the \c Fetch class to
/// ignore whatever data is in the associated execution object and return the
/// thread indices.
///
struct AspectTagThreadIndices { };
/// \brief The \c ExecutionSignature tag to use to get the thread indices
///
/// When a worklet is dispatched, it broken into pieces defined by the input
/// domain and scheduled on independent threads. During this process multiple
/// indices associated with the input and output can be generated. This tag in
/// the \c ExecutionSignature passes the index for this work. \c WorkletBase
/// contains a typedef that points to this class.
///
struct ThreadIndices : vtkm::exec::arg::ExecutionSignatureTagBase
{
// The index does not really matter because the fetch is going to ignore it.
// However, it still has to point to a valid parameter in the
// ControlSignature because the templating is going to grab a fetch tag
// whether we use it or not. 1 should be guaranteed to be valid since you
// need at least one argument for the input domain.
static const vtkm::IdComponent INDEX = 1;
typedef vtkm::exec::arg::AspectTagThreadIndices AspectTag;
};
template<typename FetchTag, typename ThreadIndicesType, typename ExecObjectType>
struct Fetch<FetchTag,
vtkm::exec::arg::AspectTagThreadIndices,
ThreadIndicesType,
ExecObjectType>
{
typedef const ThreadIndicesType &ValueType;
VTKM_EXEC_EXPORT
const ThreadIndicesType &
Load(const ThreadIndicesType &indices, const ExecObjectType &) const
{
return indices;
}
VTKM_EXEC_EXPORT
void Store(const ThreadIndicesType &,
const ExecObjectType &,
const ThreadIndicesType &) const
{
// Store is a no-op.
}
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_ThreadIndices_h

36
vtkm/exec/arg/ThreadIndicesBasic.h
View File

@ -44,8 +44,10 @@ class ThreadIndicesBasic
public:
template<typename Invocation>
VTKM_EXEC_EXPORT
ThreadIndicesBasic(vtkm::Id threadIndex, const Invocation &)
: Index(threadIndex)
ThreadIndicesBasic(vtkm::Id threadIndex, const Invocation &invocation)
: InputIndex(invocation.OutputToInputMap.Get(threadIndex)),
OutputIndex(threadIndex),
VisitIndex(invocation.VisitArray.Get(threadIndex))
{
}
@ -53,10 +55,10 @@ public:
///
/// This index refers to the input element (array value, cell, etc.) that
/// this thread is being invoked for. This is the typical index used during
/// fetches.
/// Fetch::Load.
///
VTKM_EXEC_EXPORT
vtkm::Id GetIndex() const { return this->Index; }
vtkm::Id GetInputIndex() const { return this->InputIndex; }
/// \brief The 3D index into the input domain.
///
@ -67,10 +69,32 @@ public:
/// first component with the remaining components set to 0.
///
VTKM_EXEC_EXPORT
vtkm::Id3 GetIndex3D() const { return vtkm::Id3(this->GetIndex(), 0, 0); }
vtkm::Id3 GetInputIndex3D() const
{
return vtkm::Id3(this->GetInputIndex(), 0, 0);
}
/// \brief The index into the output domain.
///
/// This index refers to the output element (array value, cell, etc.) that
/// this thread is creating. This is the typical index used during
/// Fetch::Store.
///
VTKM_EXEC_EXPORT
vtkm::Id GetOutputIndex() const { return this->OutputIndex; }
/// \brief The visit index.
///
/// When multiple output indices have the same input index, they are
/// distinguished using the visit index.
///
VTKM_EXEC_EXPORT
vtkm::IdComponent GetVisitIndex() const { return this->VisitIndex; }
private:
vtkm::Id Index;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
};
}

92
vtkm/exec/arg/ThreadIndicesTopologyMap.h
View File

@ -85,8 +85,8 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
this->IndicesFrom = connectivity.GetIndices(this->GetIndex());
this->CellShape = connectivity.GetCellShape(this->GetIndex());
this->IndicesFrom = connectivity.GetIndices(this->GetInputIndex());
this->CellShape = connectivity.GetCellShape(this->GetInputIndex());
}
/// \brief The input indices of the "from" elements.
@ -97,7 +97,22 @@ public:
/// containing the indices to the "from" elements.
///
VTKM_EXEC_EXPORT
IndicesFromType GetIndicesFrom() const { return this->IndicesFrom; }
const IndicesFromType &GetIndicesFrom() const { return this->IndicesFrom; }
/// \brief The input indices of the "from" elements in pointer form.
///
/// Returns the same object as GetIndicesFrom except that it returns a
/// pointer to the internally held object rather than a reference or copy.
/// Since the from indices can be a sizeable Vec (8 entries is common), it is
/// best not to have a bunch a copies. Thus, you can pass around a pointer
/// instead. However, care should be taken to make sure that this object does
/// not go out of scope, at which time the returned pointer becomes invalid.
///
VTKM_EXEC_EXPORT
const IndicesFromType *GetIndicesFromPointer() const
{
return &this->IndicesFrom;
}
/// \brief The shape of the input cell.
///
@ -179,12 +194,12 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
const LogicalIndexType logicalIndex = connectivity.FlatToLogicalToIndex(threadIndex);
this->Index = threadIndex;
this->LogicalIndex = logicalIndex;
this->IndicesFrom = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(threadIndex);
this->InputIndex = invocation.OutputToInputMap.Get(threadIndex);
this->OutputIndex = threadIndex;
this->VisitIndex = invocation.VisitArray.Get(threadIndex);
this->LogicalIndex = connectivity.FlatToLogicalToIndex(this->InputIndex);
this->IndicesFrom = connectivity.GetIndices(this->LogicalIndex);
this->CellShape = connectivity.GetCellShape(this->InputIndex);
}
template<typename Invocation>
@ -198,11 +213,14 @@ public:
// set its input domain incorrectly.
const ConnectivityType &connectivity = invocation.GetInputDomain();
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const LogicalIndexType logicalIndex =
detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
this->Index = index;
// We currently only support multidimensional indices on one-to-one input-
// to-output mappings. (We don't have a use case otherwise.)
this->InputIndex = this->OutputIndex = index;
this->VisitIndex = invocation.VisitArray.Get(index);
this->LogicalIndex = logicalIndex;
this->IndicesFrom = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(index);
@ -226,9 +244,9 @@ public:
/// fetches.
///
VTKM_EXEC_EXPORT
vtkm::Id GetIndex() const
vtkm::Id GetInputIndex() const
{
return this->Index;
return this->InputIndex;
}
/// \brief The 3D index into the input domain.
@ -237,11 +255,34 @@ public:
/// for the input.
///
VTKM_EXEC_EXPORT
vtkm::Id3 GetIndex3D() const
vtkm::Id3 GetInputIndex3D() const
{
return detail::InflateTo3D(this->GetIndexLogical());
}
/// \brief The index into the output domain.
///
/// This index refers to the output element (array value, cell, etc.) that
/// this thread is creating. This is the typical index used during
/// Fetch::Store.
///
VTKM_EXEC_EXPORT
vtkm::Id GetOutputIndex() const
{
return this->OutputIndex;
}
/// \brief The visit index.
///
/// When multiple output indices have the same input index, they are
/// distinguished using the visit index.
///
VTKM_EXEC_EXPORT
vtkm::IdComponent GetVisitIndex() const
{
return this->VisitIndex;
}
/// \brief The input indices of the "from" elements.
///
/// A topology map has "from" and "to" elements (for example from points to
@ -250,7 +291,22 @@ public:
/// containing the indices to the "from" elements.
///
VTKM_EXEC_EXPORT
IndicesFromType GetIndicesFrom() const { return this->IndicesFrom; }
const IndicesFromType &GetIndicesFrom() const { return this->IndicesFrom; }
/// \brief The input indices of the "from" elements in pointer form.
///
/// Returns the same object as GetIndicesFrom except that it returns a
/// pointer to the internally held object rather than a reference or copy.
/// Since the from indices can be a sizeable Vec (8 entries is common), it is
/// best not to have a bunch a copies. Thus, you can pass around a pointer
/// instead. However, care should be taken to make sure that this object does
/// not go out of scope, at which time the returned pointer becomes invalid.
///
VTKM_EXEC_EXPORT
const IndicesFromType *GetIndicesFromPointer() const
{
return &this->IndicesFrom;
}
/// \brief The shape of the input cell.
///
@ -263,7 +319,9 @@ public:
CellShapeTag GetCellShape() const { return this->CellShape; }
private:
vtkm::Id Index;
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
LogicalIndexType LogicalIndex;
IndicesFromType IndicesFrom;
CellShapeTag CellShape;

87
vtkm/exec/arg/VisitIndex.h Normal file
View File

@ -0,0 +1,87 @@
//============================================================================
// 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 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_arg_VisitIndex_h
#define vtk_m_exec_arg_VisitIndex_h
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Aspect tag to use for getting the work index.
///
/// The \c AspectTagVisitIndex aspect tag causes the \c Fetch class to ignore
/// whatever data is in the associated execution object and return the visit
/// index.
///
struct AspectTagVisitIndex { };
/// \brief The \c ExecutionSignature tag to use to get the visit index
///
/// When a worklet is dispatched, there is a scatter operation defined that
/// optionally allows each input to go to multiple output entries. When one
/// input is assigned to multiple outputs, there needs to be a mechanism to
/// uniquely identify which output is which. The visit index is a value between
/// 0 and the number of outputs a particular input goes to. This tag in the \c
/// ExecutionSignature passes the visit index for this work. \c WorkletBase
/// contains a typedef that points to this class.
///
struct VisitIndex : vtkm::exec::arg::ExecutionSignatureTagBase
{
// The index does not really matter because the fetch is going to ignore it.
// However, it still has to point to a valid parameter in the
// ControlSignature because the templating is going to grab a fetch tag
// whether we use it or not. 1 should be guaranteed to be valid since you
// need at least one argument for the input domain.
static const vtkm::IdComponent INDEX = 1;
typedef vtkm::exec::arg::AspectTagVisitIndex AspectTag;
};
template<typename FetchTag, typename ThreadIndicesType, typename ExecObjectType>
struct Fetch<FetchTag,
vtkm::exec::arg::AspectTagVisitIndex,
ThreadIndicesType,
ExecObjectType>
{
typedef vtkm::IdComponent ValueType;
VTKM_EXEC_EXPORT
vtkm::IdComponent Load(const ThreadIndicesType &indices,
const ExecObjectType &) const
{
return indices.GetVisitIndex();
}
VTKM_EXEC_EXPORT
void Store(const ThreadIndicesType &,
const ExecObjectType &,
const ValueType &) const
{
// Store is a no-op.
}
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_VisitIndex_h

2
vtkm/exec/arg/WorkIndex.h
View File

@ -63,7 +63,7 @@ struct Fetch<FetchTag,
VTKM_EXEC_EXPORT
vtkm::Id Load(const ThreadIndicesType &indices, const ExecObjectType &) const
{
return indices.GetIndex();
return indices.GetOutputIndex();
}
VTKM_EXEC_EXPORT

5
vtkm/exec/arg/testing/CMakeLists.txt
View File

@ -17,6 +17,11 @@
## Laboratory (LANL), the U.S. Government retains certain rights in
## this software.
##============================================================================
set(headers
ThreadIndicesTesting.h
)
vtkm_declare_headers(${headers})
set(unit_tests
UnitTestExecutionSignatureTag.cxx

70
vtkm/exec/arg/testing/ThreadIndicesTesting.h Normal file
View File

@ -0,0 +1,70 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_arg_testing_ThreadIndicesTesting_h
#define vtk_m_exec_arg_testing_ThreadIndicesTesting_h
#include <vtkm/Types.h>
namespace vtkm {
namespace exec {
namespace arg {
/// \brief Simplified version of ThreadIndices for unit testing purposes
///
class ThreadIndicesTesting
{
public:
VTKM_EXEC_CONT_EXPORT
ThreadIndicesTesting(vtkm::Id index)
: InputIndex(index), OutputIndex(index), VisitIndex(0) { }
VTKM_EXEC_CONT_EXPORT
ThreadIndicesTesting(vtkm::Id inputIndex,
vtkm::Id outputIndex,
vtkm::IdComponent visitIndex)
: InputIndex(inputIndex), OutputIndex(outputIndex), VisitIndex(visitIndex)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetInputIndex() const { return this->InputIndex; }
VTKM_EXEC_CONT_EXPORT
vtkm::Id3 GetInputIndex3D() const
{
return vtkm::Id3(this->GetInputIndex(), 0, 0);
}
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetOutputIndex() const { return this->OutputIndex; }
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetVisitIndex() const { return this->VisitIndex; }
private:
vtkm::Id InputIndex;
vtkm::Id OutputIndex;
vtkm::IdComponent VisitIndex;
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_testing_ThreadIndicesTesting_h

50
vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectIn.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -47,59 +44,36 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectInTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectIn,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
VTKM_TEST_ASSERT(test_equal(value, TestValue(index, T())),
"Got invalid value from Load.");
value = T(T(2)*value);
// This should be a no-op, but we should be able to call it.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
}
void operator()() const
{
std::cout << "Trying ArrayDirectIn fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -107,11 +81,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectInTests<1,T>()();
FetchArrayDirectInTests<2,T>()();
FetchArrayDirectInTests<3,T>()();
FetchArrayDirectInTests<4,T>()();
FetchArrayDirectInTests<5,T>()();
FetchArrayDirectInTests<T>()();
}
};

49
vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectInOut.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -59,19 +56,18 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectInTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectInOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
@ -80,19 +76,15 @@ struct FetchArrayDirectInTests
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
VTKM_TEST_ASSERT(test_equal(value, TestValue(index, T())),
"Got invalid value from Load.");
value = T(T(2)*value);
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
VTKM_TEST_ASSERT(g_NumSets == ARRAY_SIZE,
@ -100,23 +92,6 @@ struct FetchArrayDirectInTests
"Store method must be wrong.");
}
void operator()() const
{
std::cout << "Trying ArrayDirectInOut fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -124,11 +99,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectInTests<1,T>()();
FetchArrayDirectInTests<2,T>()();
FetchArrayDirectInTests<3,T>()();
FetchArrayDirectInTests<4,T>()();
FetchArrayDirectInTests<5,T>()();
FetchArrayDirectInTests<T>()();
}
};

49
vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagArrayDirectOut.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
@ -51,19 +48,18 @@ struct TestPortal
}
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename T>
template<typename T>
struct FetchArrayDirectOutTests
{
template<typename Invocation>
void TryInvocation(const Invocation &invocation) const
void operator()()
{
TestPortal<T> execObject;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestPortal<T> > FetchType;
FetchType fetch;
@ -72,19 +68,15 @@ struct FetchArrayDirectOutTests
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
// This is a no-op, but should be callable.
T value = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
T value = fetch.Load(indices, execObject);
value = TestValue(index, T());
// The portal will check to make sure we are setting a good value.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
value);
fetch.Store(indices, execObject, value);
}
VTKM_TEST_ASSERT(g_NumSets == ARRAY_SIZE,
@ -92,23 +84,6 @@ struct FetchArrayDirectOutTests
"Store method must be wrong.");
}
void operator()() const
{
std::cout << "Trying ArrayDirectOut fetch on parameter " << ParamIndex
<< " with type " << vtkm::testing::TypeName<T>::Name()
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
this->TryInvocation(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestPortal<T>()),
NullParam(),
NullParam()));
}
};
struct TryType
@ -116,11 +91,7 @@ struct TryType
template<typename T>
void operator()(T) const
{
FetchArrayDirectOutTests<1,T>()();
FetchArrayDirectOutTests<2,T>()();
FetchArrayDirectOutTests<3,T>()();
FetchArrayDirectOutTests<4,T>()();
FetchArrayDirectOutTests<5,T>()();
FetchArrayDirectOutTests<T>()();
}
};

46
vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx
View File

@ -47,7 +47,21 @@ struct TestPortal
}
};
struct NullParam { };
struct TestIndexPortal
{
typedef vtkm::Id ValueType;
VTKM_EXEC_CONT_EXPORT
ValueType Get(vtkm::Id index) const { return index; }
};
struct TestZeroPortal
{
typedef vtkm::IdComponent ValueType;
VTKM_EXEC_CONT_EXPORT
ValueType Get(vtkm::Id) const { return 0; }
};
template<vtkm::IdComponent InputDomainIndex,
vtkm::IdComponent ParamIndex,
@ -102,7 +116,11 @@ struct FetchArrayTopologyMapInTests
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
void(vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType)>
BaseFunctionInterface;
vtkm::internal::ConnectivityStructuredInternals<3> connectivityInternals;
@ -115,8 +133,10 @@ struct FetchArrayTopologyMapInTests
BaseFunctionInterface()
.Replace<InputDomainIndex>(connectivity)
.template Replace<ParamIndex>(TestPortal<T>()),
NullParam(),
NullParam()));
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal()));
}
};
@ -175,7 +195,11 @@ void TryStructuredPointCoordinates(
const vtkm::internal::ArrayPortalUniformPointCoordinates &coordinates)
{
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
void(vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType,
vtkm::internal::NullType)>
BaseFunctionInterface;
// Try with topology in argument 1 and point coordinates in argument 2
@ -184,8 +208,10 @@ void TryStructuredPointCoordinates(
BaseFunctionInterface()
.Replace<1>(connectivity)
.template Replace<2>(coordinates),
NullParam(),
NullParam())
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal())
);
// Try again with topology in argument 3 and point coordinates in argument 1
TryStructuredPointCoordinatesInvocation<NumDimensions,1>(
@ -193,8 +219,10 @@ void TryStructuredPointCoordinates(
BaseFunctionInterface()
.Replace<3>(connectivity)
.template Replace<1>(coordinates),
NullParam(),
NullParam())
vtkm::internal::NullType(),
vtkm::internal::NullType(),
TestIndexPortal(),
TestZeroPortal())
);
}

50
vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx
View File

@ -20,10 +20,7 @@
#include <vtkm/exec/arg/FetchTagExecObject.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/exec/ExecutionObjectBase.h>
@ -40,64 +37,37 @@ struct TestExecutionObject : public vtkm::exec::ExecutionObjectBase
vtkm::Int32 Number;
};
struct NullParam { };
template<vtkm::IdComponent ParamIndex, typename Invocation>
void TryInvocation(const Invocation &invocation)
void TryInvocation()
{
TestExecutionObject execObjectStore(EXPECTED_NUMBER);
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagExecObject,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic,
vtkm::exec::arg::ThreadIndicesTesting,
TestExecutionObject> FetchType;
FetchType fetch;
vtkm::exec::arg::ThreadIndicesBasic indices(0, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(0);
TestExecutionObject execObject = fetch.Load(
indices, invocation.Parameters.template GetParameter<ParamIndex>());
TestExecutionObject execObject = fetch.Load(indices, execObjectStore);
VTKM_TEST_ASSERT(execObject.Number == EXPECTED_NUMBER,
"Did not load object correctly.");
execObject.Number = -1;
// This should be a no-op.
fetch.Store(
indices,
invocation.Parameters.template GetParameter<ParamIndex>(),
execObject);
fetch.Store(indices, execObjectStore, execObject);
// Data in Invocation should not have changed.
execObject = invocation.Parameters.template GetParameter<ParamIndex>();
VTKM_TEST_ASSERT(execObject.Number == EXPECTED_NUMBER,
VTKM_TEST_ASSERT(execObjectStore.Number == EXPECTED_NUMBER,
"Fetch changed read-only execution object.");
}
template<vtkm::IdComponent ParamIndex>
void TryParamIndex()
{
std::cout << "Trying ExecObject fetch on parameter " << ParamIndex
<< std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
TryInvocation<ParamIndex>(vtkm::internal::make_Invocation<1>(
BaseFunctionInterface().Replace<ParamIndex>(
TestExecutionObject(EXPECTED_NUMBER)),
NullParam(),
NullParam()));
}
void TestExecObjectFetch()
{
TryParamIndex<1>();
TryParamIndex<2>();
TryParamIndex<3>();
TryParamIndex<4>();
TryParamIndex<5>();
TryInvocation();
}
} // anonymous namespace

35
vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx
View File

@ -22,56 +22,39 @@
#include <vtkm/exec/arg/FetchTagArrayDirectIn.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/internal/FunctionInterface.h>
#include <vtkm/internal/Invocation.h>
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
#include <vtkm/testing/Testing.h>
namespace {
struct NullParam { };
template<typename Invocation>
void TryInvocation(const Invocation &invocation)
void TestWorkIndexFetch()
{
std::cout << "Trying WorkIndex fetch." << std::endl;
typedef vtkm::exec::arg::Fetch<
vtkm::exec::arg::FetchTagArrayDirectIn, // Not used but probably common.
vtkm::exec::arg::AspectTagWorkIndex,
vtkm::exec::arg::ThreadIndicesBasic,
NullParam> FetchType;
vtkm::exec::arg::ThreadIndicesTesting,
vtkm::internal::NullType> FetchType;
FetchType fetch;
for (vtkm::Id index = 0; index < 10; index++)
{
vtkm::exec::arg::ThreadIndicesBasic indices(index, invocation);
vtkm::exec::arg::ThreadIndicesTesting indices(index);
vtkm::Id value = fetch.Load(indices, NullParam());
vtkm::Id value = fetch.Load(indices, vtkm::internal::NullType());
VTKM_TEST_ASSERT(value == index,
"Fetch did not give correct work index.");
value++;
// This should be a no-op.
fetch.Store(indices, NullParam(), value);
fetch.Store(indices, vtkm::internal::NullType(), value);
}
}
void TestWorkIndexFetch()
{
std::cout << "Trying WorkIndex fetch." << std::endl;
typedef vtkm::internal::FunctionInterface<
void(NullParam,NullParam,NullParam,NullParam,NullParam)>
BaseFunctionInterface;
TryInvocation(vtkm::internal::make_Invocation<1>(BaseFunctionInterface(),
NullParam(),
NullParam()));
}
} // anonymous namespace
int UnitTestFetchWorkIndex(int, char *[])

14
vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
View File

@ -96,7 +96,7 @@ template<typename T, typename Enable = void>
struct load_through_texture
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
return *(data.get());
}
@ -109,7 +109,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTextureLoad<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseScalarTextureLoad");
@ -125,7 +125,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextureLoads<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseVecTextureLoads");
@ -193,7 +193,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
typedef typename boost::remove_const<T>::type NonConstT;
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseMultipleScalarTextureLoads");
@ -204,7 +204,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
}
__device__
static T getAs(const thrust::system::cuda::pointer<T>& data)
static T getAs(const thrust::system::cuda::pointer<const T>& data)
{
//we need to fetch each component individually
const vtkm::IdComponent NUM_COMPONENTS= T::NUM_COMPONENTS;
@ -291,8 +291,8 @@ public:
VTKM_EXEC_CONT_EXPORT ConstArrayPortalFromThrust() { }
VTKM_CONT_EXPORT
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< T > begin,
const thrust::system::cuda::pointer< T > end)
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< const T > begin,
const thrust::system::cuda::pointer< const T > end)
: BeginIterator( begin ),
EndIterator( end )
{

85
vtkm/exec/cuda/internal/ThrustPatches.h Normal file
View File

@ -0,0 +1,85 @@
//============================================================================
// 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 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_cuda_internal_ThrustPatches_h
#define vtk_m_exec_cuda_internal_ThrustPatches_h
//Forward declare of WrappedBinaryOperator
namespace vtkm { namespace exec { namespace cuda { namespace internal {
template<typename T, typename F>
class WrappedBinaryOperator;
} } } } //namespace vtkm::exec::cuda::internal
namespace thrust { namespace system { namespace cuda { namespace detail {
namespace bulk_ { namespace detail { namespace accumulate_detail {
//So for thrust 1.8.0 - 1.8.2 the inclusive_scan has a bug when accumulating
//values when the binary operators states it is not commutative.
//For more complex value types, we patch thrust/bulk with fix that is found
//in issue: https://github.com/thrust/thrust/issues/692
//
//This specialization needs to be included before ANY thrust includes otherwise
//other device code inside thrust that calls it will not see it
template<typename ConcurrentGroup, typename RandomAccessIterator, typename Size, typename T, typename F>
__device__ T destructive_accumulate_n(ConcurrentGroup &g, RandomAccessIterator first, Size n, T init,
vtkm::exec::cuda::internal::WrappedBinaryOperator<T, F> binary_op)
{
typedef typename ConcurrentGroup::size_type size_type;
size_type tid = g.this_exec.index();
T x = init;
if(tid < n)
{
x = first[tid];
}
g.wait();
for(size_type offset = 1; offset < g.size(); offset += offset)
{
if(tid >= offset && tid - offset < n)
{
x = binary_op(first[tid - offset], x);
}
g.wait();
if(tid < n)
{
first[tid] = x;
}
g.wait();
}
T result = binary_op(init, first[n - 1]);
g.wait();
return result;
}
} } } //namespace bulk_::detail::accumulate_detail
} } } } //namespace thrust::system::cuda::detail
#endif //vtk_m_exec_cuda_internal_ThrustPatches_h

18
vtkm/exec/cuda/internal/WrappedOperators.h
View File

@ -240,19 +240,17 @@ struct WrappedBinaryPredicate
}
} //namespace vtkm::exec::cuda::internal
namespace thrust
{
namespace detail
{
//So for thrust 1.8.0 - 1.8.2 the inclusive_scan has a bug when accumulating
//values when the binary operators states it is not commutative. At the
//same time the is_commutative condition is used to perform faster paths. So
//We state that all WrappedBinaryOperator are commutative.
namespace thrust { namespace detail {
//
// We tell Thrust that our WrappedBinaryOperator is commutative so that we
// activate numerous fast paths inside thrust which are only available when
// the binary functor is commutative and the T type is is_arithmetic
//
//
template< typename T, typename F>
struct is_commutative< vtkm::exec::cuda::internal::WrappedBinaryOperator<T, F> > :
public thrust::detail::is_arithmetic<T> { };
}
}
} } //namespace thrust::detail
#endif //vtk_m_exec_cuda_internal_WrappedOperators_h

10
vtkm/exec/internal/VecFromPortalPermute.h
View File

@ -45,12 +45,12 @@ public:
VecFromPortalPermute() { }
VTKM_EXEC_EXPORT
VecFromPortalPermute(const IndexVecType &indices, const PortalType &portal)
VecFromPortalPermute(const IndexVecType *indices, const PortalType &portal)
: Indices(indices), Portal(portal) { }
VTKM_EXEC_EXPORT
vtkm::IdComponent GetNumberOfComponents() const {
return this->Indices.GetNumberOfComponents();
return this->Indices->GetNumberOfComponents();
}
template<vtkm::IdComponent DestSize>
@ -61,18 +61,18 @@ public:
vtkm::Min(DestSize, this->GetNumberOfComponents());
for (vtkm::IdComponent index = 0; index < numComponents; index++)
{
dest[index] = this->Portal.Get(this->Indices[index]);
dest[index] = (*this)[index];
}
}
VTKM_EXEC_EXPORT
ComponentType operator[](vtkm::IdComponent index) const
{
return this->Portal.Get(this->Indices[index]);
return this->Portal.Get((*this->Indices)[index]);
}
private:
IndexVecType Indices;
const IndexVecType *Indices;
PortalType Portal;
};

200
vtkm/exec/internal/WorkletInvokeFunctorDetail.h
View File

@ -84,6 +84,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1>
@ -94,14 +96,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -136,6 +142,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1>
VTKM_EXEC_EXPORT
@ -145,14 +153,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -180,6 +192,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -191,14 +205,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -242,6 +260,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2>
@ -252,14 +272,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -296,6 +320,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -308,14 +334,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -368,6 +398,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -379,14 +411,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -432,6 +468,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -445,14 +483,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -514,6 +556,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -526,14 +570,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -588,6 +636,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -602,14 +652,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -680,6 +734,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -693,14 +749,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -764,6 +824,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -779,14 +841,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -866,6 +932,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -880,14 +948,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -960,6 +1032,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -976,14 +1050,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1072,6 +1150,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1087,14 +1167,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1176,6 +1260,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1193,14 +1279,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1298,6 +1388,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1314,14 +1406,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1412,6 +1508,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1430,14 +1528,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1544,6 +1646,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1561,14 +1665,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1668,6 +1776,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename R,
typename P1,
@ -1687,14 +1797,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<R(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;
@ -1810,6 +1924,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
typename P1,
typename P2,
@ -1828,14 +1944,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<void(P1,P2,P3,P4,P5,P6,P7,P8,P9,P10)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
typedef InvocationToFetch<ThreadIndicesType,Invocation,1> FetchInfo1;
typedef typename FetchInfo1::type FetchType1;

20
vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in
View File

@ -146,6 +146,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
$template_params(num_params)>
VTKM_EXEC_EXPORT
@ -155,14 +157,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params)>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params)>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
$for(param_index in range(1, num_params+1))\
typedef InvocationToFetch<ThreadIndicesType,Invocation,$(param_index)> FetchInfo$(param_index);
@ -201,6 +207,8 @@ template<typename WorkletType,
typename ParameterInterface,
typename ControlInterface,
vtkm::IdComponent InputDomainIndex,
typename OutputToInputMapType,
typename VisitArrayType,
typename ThreadIndicesType,
$template_params(num_params, start=1)>
VTKM_EXEC_EXPORT
@ -210,14 +218,18 @@ void DoWorkletInvokeFunctor(
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params, return_type='void')>,
InputDomainIndex> &invocation,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> &invocation,
const ThreadIndicesType &threadIndices)
{
typedef vtkm::internal::Invocation<
ParameterInterface,
ControlInterface,
vtkm::internal::FunctionInterface<$signature(num_params, return_type='void')>,
InputDomainIndex> Invocation;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> Invocation;
$for(param_index in range(1, num_params+1))\
typedef InvocationToFetch<ThreadIndicesType,Invocation,$(param_index)> FetchInfo$(param_index);

2
vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx
View File

@ -87,7 +87,7 @@ struct VecFromPortalPermuteTestFunctor
indices.Append(offset + 2*index);
}
VecType vec(indices, portal);
VecType vec(&indices, portal);
VTKM_TEST_ASSERT(vec.GetNumberOfComponents() == length,
"Wrong length.");

33
vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx
View File

@ -47,6 +47,19 @@ struct TestExecObject
vtkm::Id *Value;
};
struct MyOutputToInputMapPortal
{
typedef vtkm::Id ValueType;
VTKM_EXEC_CONT_EXPORT
vtkm::Id Get(vtkm::Id index) const { return index; }
};
struct MyVisitArrayPortal
{
typedef vtkm::IdComponent ValueType;
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
};
struct TestFetchTagInput { };
struct TestFetchTagOutput { };
@ -78,7 +91,7 @@ struct Fetch<
VTKM_EXEC_EXPORT
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject) const {
return *execObject.Value + 10*indices.GetIndex();
return *execObject.Value + 10*indices.GetInputIndex();
}
VTKM_EXEC_EXPORT
@ -109,7 +122,7 @@ struct Fetch<
void Store(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject,
ValueType value) const {
*execObject.Value = value + 20*indices.GetIndex();
*execObject.Value = value + 20*indices.GetOutputIndex();
}
};
@ -141,13 +154,17 @@ typedef vtkm::internal::Invocation<
ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface1,
1> InvocationType1;
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal> InvocationType1;
typedef vtkm::internal::Invocation<
ExecutionParameterInterface,
TestControlInterface,
TestExecutionInterface2,
1> InvocationType2;
1,
MyOutputToInputMapPortal,
MyVisitArrayPortal> InvocationType2;
// Not a full worklet, but provides operators that we expect in a worklet.
struct TestWorkletProxy : vtkm::exec::FunctorBase
@ -232,7 +249,9 @@ void TestDoWorkletInvoke()
CallDoWorkletInvokeFunctor(
vtkm::internal::make_Invocation<1>(execObjects,
TestControlInterface(),
TestExecutionInterface1()),
TestExecutionInterface1(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
1);
VTKM_TEST_ASSERT(inputTestValue == 5, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 100 + 30,
@ -244,7 +263,9 @@ void TestDoWorkletInvoke()
CallDoWorkletInvokeFunctor(
vtkm::internal::make_Invocation<1>(execObjects,
TestControlInterface(),
TestExecutionInterface2()),
TestExecutionInterface2(),
MyOutputToInputMapPortal(),
MyVisitArrayPortal()),
2);
VTKM_TEST_ASSERT(inputTestValue == 6, "Input value changed.");
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 200 + 30*2,

2
vtkm/exec/testing/UnitTestParametricCoordinates.cxx
View File

@ -85,7 +85,7 @@ static void CompareCoordinates(const PointWCoordsType &pointWCoords,
shape,
workletProxy);
VTKM_TEST_ASSERT(!errorMessage.IsErrorRaised(), messageBuffer);
VTKM_TEST_ASSERT(test_equal(computedWCoords, trueWCoords),
VTKM_TEST_ASSERT(test_equal(computedWCoords, trueWCoords, 0.01),
"Computed wrong world coords from parametric coords.");
Vector3 computedPCoords

132
vtkm/internal/Invocation.h
View File

@ -34,7 +34,9 @@ namespace internal {
template<typename _ParameterInterface,
typename _ControlInterface,
typename _ExecutionInterface,
vtkm::IdComponent _InputDomainIndex>
vtkm::IdComponent _InputDomainIndex,
typename _OutputToInputMapType = vtkm::internal::NullType,
typename _VisitArrayType = vtkm::internal::NullType>
struct Invocation
{
/// \brief The types of the parameters
@ -68,8 +70,32 @@ struct Invocation
///
static const vtkm::IdComponent InputDomainIndex = _InputDomainIndex;
/// \brief An array representing the output to input map.
///
/// When a worklet is invoked, there is an optional scatter operation that
/// allows you to vary the number of outputs each input affects. This is
/// represented with a map where each output points to an input that creates
/// it.
///
typedef _OutputToInputMapType OutputToInputMapType;
/// \brief An array containing visit indices.
///
/// When a worklet is invoked, there is an optinonal scatter operation that
/// allows you to vary the number of outputs each input affects. Thus,
/// multiple outputs may point to the same input. The visit index uniquely
/// identifies which instance each is.
///
typedef _VisitArrayType VisitArrayType;
VTKM_CONT_EXPORT
Invocation(ParameterInterface parameters) : Parameters(parameters) { }
Invocation(ParameterInterface parameters,
OutputToInputMapType outputToInputMap = OutputToInputMapType(),
VisitArrayType visitArray = VisitArrayType())
: Parameters(parameters),
OutputToInputMap(outputToInputMap),
VisitArray(visitArray)
{ }
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c Parameters replaced.
@ -79,7 +105,9 @@ struct Invocation
typedef Invocation<NewParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex> type;
InputDomainIndex,
OutputToInputMapType,
VisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
@ -90,7 +118,7 @@ struct Invocation
typename ChangeParametersType<NewParameterInterface>::type
ChangeParameters(NewParameterInterface newParameters) const {
return typename ChangeParametersType<NewParameterInterface>::type(
newParameters);
newParameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -111,7 +139,7 @@ struct Invocation
typename ChangeControlInterfaceType<NewControlInterface>::type
ChangeControlInterface(NewControlInterface) const {
return typename ChangeControlInterfaceType<NewControlInterface>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -132,7 +160,7 @@ struct Invocation
typename ChangeExecutionInterfaceType<NewExecutionInterface>::type
ChangeExecutionInterface(NewExecutionInterface) const {
return typename ChangeExecutionInterfaceType<NewExecutionInterface>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
@ -154,7 +182,55 @@ struct Invocation
typename ChangeInputDomainIndexType<NewInputDomainIndex>::type
ChangeInputDomainIndex() const {
return typename ChangeInputDomainIndexType<NewInputDomainIndex>::type(
this->Parameters);
this->Parameters, this->OutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c OutputToInputMapType replaced.
///
template<typename NewOutputToInputMapType>
struct ChangeOutputToInputMapType {
typedef Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
NewOutputToInputMapType,
VisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
/// \c OutputToInputMap is replaced with that provided.
///
template<typename NewOutputToInputMapType>
VTKM_CONT_EXPORT
typename ChangeOutputToInputMapType<NewOutputToInputMapType>::type
ChangeOutputToInputMap(NewOutputToInputMapType newOutputToInputMap) const {
return typename ChangeOutputToInputMapType<NewOutputToInputMapType>::type(
this->Parameters, newOutputToInputMap, this->VisitArray);
}
/// Defines a new \c Invocation type that is the same as this type except
/// with the \c VisitArrayType replaced.
///
template<typename NewVisitArrayType>
struct ChangeVisitArrayType {
typedef Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
NewVisitArrayType> type;
};
/// Returns a new \c Invocation that is the same as this one except that the
/// \c VisitArray is replaced with that provided.
///
template<typename NewVisitArrayType>
VTKM_CONT_EXPORT
typename ChangeVisitArrayType<NewVisitArrayType>::type
ChangeVisitArray(NewVisitArrayType newVisitArray) const {
return typename ChangeVisitArrayType<NewVisitArrayType>::type(
this->Parameters, this->OutputToInputMap, newVisitArray);
}
/// A convenience typedef for the input domain type.
@ -172,13 +248,43 @@ struct Invocation
}
/// The state of an \c Invocation object holds the parameters of the
/// invocation.
/// invocation. As well as the output to input map and the visit array.
///
ParameterInterface Parameters;
OutputToInputMapType OutputToInputMap;
VisitArrayType VisitArray;
};
/// Convenience function for creating an Invocation object.
///
template<vtkm::IdComponent InputDomainIndex,
typename ControlInterface,
typename ExecutionInterface,
typename ParameterInterface,
typename OutputToInputMapType,
typename VisitArrayType>
VTKM_CONT_EXPORT
vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>
make_Invocation(const ParameterInterface &params,
ControlInterface,
ExecutionInterface,
OutputToInputMapType outputToInputMap,
VisitArrayType visitArray)
{
return vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex,
OutputToInputMapType,
VisitArrayType>(params,
outputToInputMap,
visitArray);
}
template<vtkm::IdComponent InputDomainIndex,
typename ControlInterface,
typename ExecutionInterface,
@ -192,10 +298,12 @@ make_Invocation(const ParameterInterface &params,
ControlInterface = ControlInterface(),
ExecutionInterface = ExecutionInterface())
{
return vtkm::internal::Invocation<ParameterInterface,
ControlInterface,
ExecutionInterface,
InputDomainIndex>(params);
return vtkm::internal::make_Invocation<InputDomainIndex>(
params,
ControlInterface(),
ExecutionInterface(),
vtkm::internal::NullType(),
vtkm::internal::NullType());
}
}

32
vtkm/internal/ListTagDetail.h
View File

@ -46,8 +46,6 @@ struct ListRoot { };
template<typename signature>
struct ListBase { };
struct ListParamNull { };
//-----------------------------------------------------------------------------
template<typename Functor>
@ -1151,21 +1149,21 @@ struct ListContainsImpl<ListBase<void(T1,T2,T3,T4,T5,T6,T7,T8,T9,T10,T11,T12,T13
/// A basic tag for a list of typenames. This struct can be subclassed
/// and still behave like a list tag.
template<typename T1 = vtkm::detail::ListParamNull,
typename T2 = vtkm::detail::ListParamNull,
typename T3 = vtkm::detail::ListParamNull,
typename T4 = vtkm::detail::ListParamNull,
typename T5 = vtkm::detail::ListParamNull,
typename T6 = vtkm::detail::ListParamNull,
typename T7 = vtkm::detail::ListParamNull,
typename T8 = vtkm::detail::ListParamNull,
typename T9 = vtkm::detail::ListParamNull,
typename T10 = vtkm::detail::ListParamNull,
typename T11 = vtkm::detail::ListParamNull,
typename T12 = vtkm::detail::ListParamNull,
typename T13 = vtkm::detail::ListParamNull,
typename T14 = vtkm::detail::ListParamNull,
typename T15 = vtkm::detail::ListParamNull>
template<typename T1 = vtkm::internal::NullType,
typename T2 = vtkm::internal::NullType,
typename T3 = vtkm::internal::NullType,
typename T4 = vtkm::internal::NullType,
typename T5 = vtkm::internal::NullType,
typename T6 = vtkm::internal::NullType,
typename T7 = vtkm::internal::NullType,
typename T8 = vtkm::internal::NullType,
typename T9 = vtkm::internal::NullType,
typename T10 = vtkm::internal::NullType,
typename T11 = vtkm::internal::NullType,
typename T12 = vtkm::internal::NullType,
typename T13 = vtkm::internal::NullType,
typename T14 = vtkm::internal::NullType,
typename T15 = vtkm::internal::NullType>
struct ListTagBase : detail::ListRoot
{
typedef detail::ListBase<void(T1,T2,T3,T4,T5,T6,T7,T8,T9,T10,T11,T12,T13,T14,T15)> List;

4
vtkm/internal/ListTagDetail.h.in
View File

@ -94,8 +94,6 @@ struct ListRoot { };
template<typename signature>
struct ListBase { };
struct ListParamNull { };
//-----------------------------------------------------------------------------
template<typename Functor>
@ -172,7 +170,7 @@ $endfor\
/// A basic tag for a list of typenames. This struct can be subclassed
/// and still behave like a list tag.
template<$template_params(max_base_list, default=' = vtkm::detail::ListParamNull')>
template<$template_params(max_base_list, default=' = vtkm::internal::NullType')>
struct ListTagBase : detail::ListRoot
{
typedef detail::ListBase<void($param_list(max_base_list))> List;

120
vtkm/io/reader/VTKDataSetReaderBase.h
View File

@ -33,11 +33,13 @@
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/smart_ptr/scoped_ptr.hpp>
#include <boost/lexical_cast.hpp>
#include <boost/type_traits/is_same.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
#include <algorithm>
#include <fstream>
#include <string>
#include <typeinfo>
#include <vector>
@ -92,34 +94,92 @@ inline void parseAssert(bool condition)
}
}
struct DummyFixed8Type
{
vtkm::UInt8 data;
};
template <typename T> struct StreamIOType { typedef T Type; };
template <> struct StreamIOType<vtkm::Int8> { typedef vtkm::Int16 Type; };
template <> struct StreamIOType<vtkm::UInt8> { typedef vtkm::UInt16 Type; };
// Since Fields and DataSets store data in the default DynamicArrayHandle, convert
// the data to the closest type supported by default. The following will
// need to be updated if DynamicArrayHandle or TypeListTagCommon changes.
template <typename T> struct ClosestCommonType { typedef T Type; };
template <> struct ClosestCommonType<vtkm::Int8> { typedef vtkm::Int32 Type; };
template <> struct ClosestCommonType<vtkm::UInt8> { typedef vtkm::Int32 Type; };
template <> struct ClosestCommonType<vtkm::Int16> { typedef vtkm::Int32 Type; };
template <> struct ClosestCommonType<vtkm::UInt16> { typedef vtkm::Int32 Type; };
template <> struct ClosestCommonType<vtkm::UInt32> { typedef vtkm::Int64 Type; };
template <> struct ClosestCommonType<vtkm::UInt64> { typedef vtkm::Int64 Type; };
template <typename T> struct ClosestFloat { typedef T Type; };
template <> struct ClosestFloat<vtkm::Int8> { typedef vtkm::Float32 Type; };
template <> struct ClosestFloat<vtkm::UInt8> { typedef vtkm::Float32 Type; };
template <> struct ClosestFloat<vtkm::Int16> { typedef vtkm::Float32 Type; };
template <> struct ClosestFloat<vtkm::UInt16> { typedef vtkm::Float32 Type; };
template <> struct ClosestFloat<vtkm::Int32> { typedef vtkm::Float64 Type; };
template <> struct ClosestFloat<vtkm::UInt32> { typedef vtkm::Float64 Type; };
template <> struct ClosestFloat<vtkm::Int64> { typedef vtkm::Float64 Type; };
template <> struct ClosestFloat<vtkm::UInt64> { typedef vtkm::Float64 Type; };
template <typename T>
struct TypeTraits
vtkm::cont::DynamicArrayHandle CreateDynamicArrayHandle(const std::vector<T> &vec)
{
typedef T AsciiReadType;
};
switch (vtkm::VecTraits<T>::NUM_COMPONENTS)
{
case 1:
{
typedef typename ClosestCommonType<T>::Type CommonType;
if (!boost::is_same<T, CommonType>::value)
{
std::cerr << "Type " << typeid(T).name() << " is currently unsupported. "
<< "Converting to " << typeid(CommonType).name() << "." << std::endl;
}
template <>
struct TypeTraits<vtkm::Int8>
{
typedef vtkm::Int16 AsciiReadType;
};
vtkm::cont::ArrayHandle<CommonType> output;
output.Allocate(static_cast<vtkm::Id>(vec.size()));
for (vtkm::Id i = 0; i < output.GetNumberOfValues(); ++i)
{
output.GetPortalControl().Set(i,
static_cast<CommonType>(vec[static_cast<std::size_t>(i)]));
}
template <>
struct TypeTraits<vtkm::UInt8>
{
typedef vtkm::UInt16 AsciiReadType;
};
return vtkm::cont::DynamicArrayHandle(output);
}
case 2:
case 3:
{
typedef typename vtkm::VecTraits<T>::ComponentType InComponentType;
typedef typename ClosestFloat<InComponentType>::Type OutComponentType;
typedef vtkm::Vec<OutComponentType, 3> CommonType;
if (!boost::is_same<T, CommonType>::value)
{
std::cerr << "Type " << typeid(InComponentType).name()
<< "[" << vtkm::VecTraits<T>::NUM_COMPONENTS << "] "
<< "is currently unsupported. Converting to "
<< typeid(OutComponentType).name() << "[3]." << std::endl;
}
template <>
struct TypeTraits<DummyFixed8Type>
{
typedef vtkm::Float32 AsciiReadType;
};
vtkm::cont::ArrayHandle<CommonType> output;
output.Allocate(static_cast<vtkm::Id>(vec.size()));
for (vtkm::Id i = 0; i < output.GetNumberOfValues(); ++i)
{
CommonType outval = CommonType();
for (vtkm::IdComponent j = 0; j < vtkm::VecTraits<T>::NUM_COMPONENTS; ++j)
{
outval[j] = static_cast<OutComponentType>(
vtkm::VecTraits<T>::GetComponent(vec[static_cast<std::size_t>(i)], j));
}
output.GetPortalControl().Set(i, outval);
}
return vtkm::cont::DynamicArrayHandle(output);
}
default:
{
std::cerr << "Only 1, 2, or 3 components supported. Skipping." << std::endl;
return vtkm::cont::DynamicArrayHandle(vtkm::cont::ArrayHandle<vtkm::Float32>());
}
}
}
} // namespace internal
@ -443,7 +503,7 @@ private:
std::size_t numValues;
this->DataFile->Stream >> dataName >> numValues >> std::ws;
this->SkipArray(numElements * numValues, internal::DummyFixed8Type());
this->SkipArray(numElements * numValues, internal::ColorChannel8());
}
void ReadLookupTable(std::string &dataName)
@ -453,7 +513,7 @@ private:
std::size_t numEntries;
this->DataFile->Stream >> dataName >> numEntries >> std::ws;
this->SkipArray(numEntries, vtkm::Vec<internal::DummyFixed8Type, 4>());
this->SkipArray(numEntries, vtkm::Vec<internal::ColorChannel8, 4>());
}
void ReadTextureCoordinates(std::size_t numElements, std::string &dataName,
@ -544,11 +604,7 @@ private:
std::vector<T> buffer(this->NumElements);
this->Reader->ReadArray(buffer);
vtkm::cont::ArrayHandle<T> data;
data.Allocate(static_cast<vtkm::Id>(buffer.size()));
std::copy(buffer.begin(), buffer.end(),
vtkm::cont::ArrayPortalToIteratorBegin(data.GetPortalControl()));
*this->Data = vtkm::cont::DynamicArrayHandle(data);
*this->Data = internal::CreateDynamicArrayHandle(buffer);
}
template <typename T>
@ -602,7 +658,7 @@ private:
{
for (vtkm::IdComponent j = 0; j < numComponents; ++j)
{
typename internal::TypeTraits<ComponentType>::AsciiReadType val;
typename internal::StreamIOType<ComponentType>::Type val;
this->DataFile->Stream >> val;
vtkm::VecTraits<T>::SetComponent(buffer[i], j,
static_cast<ComponentType>(val));
@ -647,7 +703,7 @@ private:
{
for (vtkm::IdComponent j = 0; j < numComponents; ++j)
{
typename internal::TypeTraits<ComponentType>::AsciiReadType val;
typename internal::StreamIOType<ComponentType>::Type val;
this->DataFile->Stream >> val;
}
}
@ -694,7 +750,7 @@ private:
}
} // vtkm::io::reader
VTKM_BASIC_TYPE_VECTOR(io::reader::internal::DummyFixed8Type)
VTKM_BASIC_TYPE_VECTOR(io::reader::internal::ColorChannel8)
VTKM_BASIC_TYPE_VECTOR(io::reader::internal::DummyBitType)
#endif // vtk_m_io_reader_VTKDataSetReaderBase_h

61
vtkm/io/reader/internal/TypeInfo.h
View File

@ -22,6 +22,7 @@
#include <vtkm/Types.h>
#include <algorithm>
#include <cassert>
#include <string>
@ -118,6 +119,7 @@ inline DataType DataTypeId(const std::string &str)
return type;
}
struct DummyBitType
{
// Needs to work with streams' << operator
@ -127,6 +129,43 @@ struct DummyBitType
}
};
class ColorChannel8
{
public:
ColorChannel8() : Data()
{ }
ColorChannel8(vtkm::UInt8 val) : Data(val)
{ }
ColorChannel8(vtkm::Float32 val)
: Data(static_cast<vtkm::UInt8>(std::min(std::max(val, 1.0f), 0.0f) * 255))
{ }
operator vtkm::Float32() const
{
return static_cast<vtkm::Float32>(this->Data)/255.0f;
}
operator vtkm::UInt8() const
{
return this->Data;
}
private:
vtkm::UInt8 Data;
};
inline std::ostream& operator<<(std::ostream& out, const ColorChannel8 &val)
{
return out << static_cast<vtkm::Float32>(val);
}
inline std::istream& operator>>(std::istream& in, ColorChannel8 &val)
{
vtkm::Float32 fval;
in >> fval;
val = ColorChannel8(fval);
return in;
}
template <typename T, typename Functor>
inline void SelectVecTypeAndCall(T, vtkm::IdComponent numComponents, const Functor &functor)
{
@ -163,28 +202,28 @@ inline void SelectTypeAndCall(DataType dtype, vtkm::IdComponent numComponents,
SelectVecTypeAndCall(DummyBitType(), numComponents, functor);
break;
case DTYPE_UNSIGNED_CHAR:
SelectVecTypeAndCall(vtkm::Int8(), numComponents, functor);
break;
case DTYPE_CHAR:
SelectVecTypeAndCall(vtkm::UInt8(), numComponents, functor);
break;
case DTYPE_UNSIGNED_SHORT:
SelectVecTypeAndCall(vtkm::Int16(), numComponents, functor);
case DTYPE_CHAR:
SelectVecTypeAndCall(vtkm::Int8(), numComponents, functor);
break;
case DTYPE_SHORT:
case DTYPE_UNSIGNED_SHORT:
SelectVecTypeAndCall(vtkm::UInt16(), numComponents, functor);
break;
case DTYPE_UNSIGNED_INT:
SelectVecTypeAndCall(vtkm::Int32(), numComponents, functor);
case DTYPE_SHORT:
SelectVecTypeAndCall(vtkm::Int16(), numComponents, functor);
break;
case DTYPE_INT:
case DTYPE_UNSIGNED_INT:
SelectVecTypeAndCall(vtkm::UInt32(), numComponents, functor);
break;
case DTYPE_INT:
SelectVecTypeAndCall(vtkm::Int32(), numComponents, functor);
break;
case DTYPE_UNSIGNED_LONG:
SelectVecTypeAndCall(vtkm::Int64(), numComponents, functor);
SelectVecTypeAndCall(vtkm::UInt64(), numComponents, functor);
break;
case DTYPE_LONG:
SelectVecTypeAndCall(vtkm::UInt64(), numComponents, functor);
SelectVecTypeAndCall(vtkm::Int64(), numComponents, functor);
break;
case DTYPE_FLOAT:
SelectVecTypeAndCall(vtkm::Float32(), numComponents, functor);

78
vtkm/io/writer/VTKDataSetWriter.h
View File

@ -156,6 +156,69 @@ public:
}
};
template <typename T> struct DataTypeName
{
static const char* Name() { return "unknown"; }
};
template <> struct DataTypeName<vtkm::Int8>
{
static const char* Name() { return "char"; }
};
template <> struct DataTypeName<vtkm::UInt8>
{
static const char* Name() { return "unsigned_char"; }
};
template <> struct DataTypeName<vtkm::Int16>
{
static const char* Name() { return "short"; }
};
template <> struct DataTypeName<vtkm::UInt16>
{
static const char* Name() { return "unsigned_short"; }
};
template <> struct DataTypeName<vtkm::Int32>
{
static const char* Name() { return "int"; }
};
template <> struct DataTypeName<vtkm::UInt32>
{
static const char* Name() { return "unsigned_int"; }
};
template <> struct DataTypeName<vtkm::Int64>
{
static const char* Name() { return "long"; }
};
template <> struct DataTypeName<vtkm::UInt64>
{
static const char* Name() { return "unsigned_long"; }
};
template <> struct DataTypeName<vtkm::Float32>
{
static const char* Name() { return "float"; }
};
template <> struct DataTypeName<vtkm::Float64>
{
static const char* Name() { return "double"; }
};
class GetDataTypeName
{
public:
GetDataTypeName(std::string &name)
: Name(&name)
{ }
template <typename ArrayHandleType>
void operator()(const ArrayHandleType &) const
{
typedef typename vtkm::VecTraits<typename ArrayHandleType::ValueType>::ComponentType
DataType;
*this->Name = DataTypeName<DataType>::Name();
}
private:
std::string *Name;
};
}
namespace vtkm
@ -179,7 +242,10 @@ private:
vtkm::Id npoints = cdata.GetNumberOfValues();
out << "POINTS " << npoints << " float" << std::endl;
std::string typeName;
cdata.CastAndCall(GetDataTypeName(typeName));
out << "POINTS " << npoints << " " << typeName << " " << std::endl;
cdata.CastAndCall(OutputPointsFunctor(out));
}
@ -251,7 +317,10 @@ private:
out << "POINT_DATA " << npoints << std::endl;
wrote_header = true;
out << "SCALARS " << field.GetName() << " float "<< ncomps << std::endl;
std::string typeName;
field.GetData().CastAndCall(GetDataTypeName(typeName));
out << "SCALARS " << field.GetName() << " " << typeName << " " << ncomps << std::endl;
out << "LOOKUP_TABLE default" << std::endl;
field.GetData().CastAndCall(OutputFieldFunctor(out));
@ -281,7 +350,10 @@ private:
out << "CELL_DATA " << ncells << std::endl;
wrote_header = true;
out << "SCALARS " << field.GetName() << " float "<< ncomps << std::endl;
std::string typeName;
field.GetData().CastAndCall(GetDataTypeName(typeName));
out << "SCALARS " << field.GetName() << " " << typeName << " " << ncomps << std::endl;
out << "LOOKUP_TABLE default" << std::endl;
field.GetData().CastAndCall(OutputFieldFunctor(out));

188
vtkm/opengl/BufferState.h Normal file
View File

@ -0,0 +1,188 @@
//============================================================================
// 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 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_opengl_BufferState_h
#define vtk_m_opengl_BufferState_h
//gl headers needs to be buffer anything to do with buffer's
#include <vtkm/opengl/internal/OpenGLHeaders.h>
#include <vtkm/opengl/internal/BufferTypePicker.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/smart_ptr/scoped_ptr.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
namespace vtkm{
namespace opengl{
namespace internal
{
/// \brief Device backend and opengl interop resources management
///
/// \c TransferResource manages a context for a given device backend and a
/// single OpenGL buffer as efficiently as possible.
///
/// Default implementation is a no-op
class TransferResource
{
public:
virtual ~TransferResource() {}
};
}
/// \brief Manages the state for transferring an ArrayHandle to opengl.
///
/// \c BufferState holds all the relevant data information for a given ArrayHandle
/// mapping into OpenGL. Reusing the state information for all renders of an
/// ArrayHandle will allow for the most efficient interop between backends and
/// OpenGL ( especially for CUDA ).
///
///
/// The interop code in vtk-m uses a lazy buffer re-allocation.
///
class BufferState
{
public:
/// Construct a BufferState using an existing GLHandle
BufferState(GLuint& gLHandle):
OpenGLHandle(&gLHandle),
BufferType(GL_INVALID_VALUE),
SizeOfActiveSection(0),
CapacityOfBuffer(0),
DefaultGLHandle(0),
Resource(NULL)
{
}
/// Construct a BufferState using an existing GLHandle and type
BufferState(GLuint& gLHandle, GLenum type):
OpenGLHandle(&gLHandle),
BufferType(type),
SizeOfActiveSection(0),
CapacityOfBuffer(0),
DefaultGLHandle(0),
Resource(NULL)
{
}
BufferState():
OpenGLHandle(NULL),
BufferType(GL_INVALID_VALUE),
SizeOfActiveSection(0),
CapacityOfBuffer(0),
DefaultGLHandle(0),
Resource(NULL)
{
this->OpenGLHandle = &this->DefaultGLHandle;
}
~BufferState()
{
//don't delete this as it points to user memory, or stack allocated
//memory inside this object instance
this->OpenGLHandle = NULL;
}
/// \brief get the OpenGL buffer handle
///
GLuint* GetHandle() const { return this->OpenGLHandle; }
/// \brief return if this buffer has a valid OpenGL buffer type
///
bool HasType() const { return this->BufferType != GL_INVALID_VALUE; }
/// \brief return what OpenGL buffer type we are bound to
///
/// will return GL_INVALID_VALUE if we don't have a valid type set
GLenum GetType() const { return this->BufferType; }
/// \brief Set what type of OpenGL buffer type we should bind as
///
void SetType(GLenum type) { this->BufferType = type; }
/// \brief deduce the buffer type from the template value type that
/// was passed in, and set that as our type
///
/// Will be GL_ELEMENT_ARRAY_BUFFER for
/// vtkm::Int32, vtkm::UInt32, vtkm::Int64, vtkm::UInt64, vtkm::Id, and vtkm::IdComponent
/// will be GL_ARRAY_BUFFER for everything else.
template<typename T>
void DeduceAndSetType(T t)
{ this->BufferType = vtkm::opengl::internal::BufferTypePicker(t); }
/// \brief Get the size of the buffer in bytes
///
/// Get the size of the active section of the buffer
///This will always be <= the capacity of the buffer
vtkm::Int64 GetSize() const { return this->SizeOfActiveSection; }
//Set the size of buffer in bytes
//This will always needs to be <= the capacity of the buffer
//Note: This call should only be used internally by vtk-m
void SetSize(vtkm::Int64 size) { this->SizeOfActiveSection = size; }
/// \brief Get the capacity of the buffer in bytes
///
/// The buffers that vtk-m allocate in OpenGL use lazy resizing. This allows
/// vtk-m to not have to reallocate a buffer while the size stays the same
/// or shrinks. This allows allows the cuda to OpenGL to perform significantly
/// better as we than don't need to call cudaGraphicsGLRegisterBuffer as
/// often
vtkm::Int64 GetCapacity() const { return this->CapacityOfBuffer; }
// Helper function to compute when we should resize the capacity of the
// buffer
bool ShouldRealloc(vtkm::Int64 desiredSize) const
{
const bool haveNotEnoughRoom = this->GetCapacity() < desiredSize;
const bool haveTooMuchRoom = this->GetCapacity() > (desiredSize*2);
return (haveNotEnoughRoom || haveTooMuchRoom);
}
//Set the capacity of buffer in bytes
//The capacity of a buffer can be larger than the active size of buffer
//Note: This call should only be used internally by vtk-m
void SetCapacity(vtkm::Int64 capacity) { this->CapacityOfBuffer = capacity; }
//Note: This call should only be used internally by vtk-m
vtkm::opengl::internal::TransferResource* GetResource()
{ return this->Resource.get(); }
//Note: This call should only be used internally by vtk-m
void SetResource( vtkm::opengl::internal::TransferResource* resource)
{ this->Resource.reset(resource); }
private:
//explicitly state the BufferState doesn't support copy or move semantics
BufferState(const BufferState&);
void operator=(const BufferState&);
GLuint* OpenGLHandle;
GLenum BufferType;
vtkm::Int64 SizeOfActiveSection; //must be Int64 as size can be over 2billion
vtkm::Int64 CapacityOfBuffer; //must be Int64 as size can be over 2billion
GLuint DefaultGLHandle;
boost::scoped_ptr<vtkm::opengl::internal::TransferResource> Resource;
};
}}
#endif //vtk_m_opengl_BufferState_h

1
vtkm/opengl/CMakeLists.txt
View File

@ -19,6 +19,7 @@
##============================================================================
set(headers
BufferState.h
TransferToOpenGL.h
)

37
vtkm/opengl/TransferToOpenGL.h
View File

@ -21,37 +21,21 @@
#define vtk_m_opengl_TransferToOpenGL_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/opengl/BufferState.h>
#include <vtkm/opengl/internal/TransferToOpenGL.h>
namespace vtkm{
namespace opengl{
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible. Will return the type of array buffer
/// that we have bound the handle too. Will be GL_ELEMENT_ARRAY_BUFFER for
/// vtkm::Id, and GL_ARRAY_BUFFER for everything else.
///
/// This function keeps the buffer as the active buffer of the returned type.
///
/// This function will throw exceptions if the transfer wasn't possible
///
template<typename ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
GLenum TransferToOpenGL(vtkm::cont::ArrayHandle<ValueType,StorageTag> handle,
GLuint& openGLHandle,
DeviceAdapterTag)
{
vtkm::opengl::internal::TransferToOpenGL<ValueType, DeviceAdapterTag> toGL;
toGL.Transfer(handle,openGLHandle);
return toGL.GetType();
}
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible. Will use the given \p type as how
/// to bind the buffer.
/// to OpenGL as efficiently as possible. Will use the given \p state to determine
/// what buffer handle to use, and the type to bind the buffer handle too.
/// If the type of buffer hasn't been determined the transfer will use
/// deduceAndSetBufferType to do so. Lastly state also holds on to per backend resources
/// that allow for efficient updating to open gl
///
/// This function keeps the buffer as the active buffer of the input type.
///
@ -60,12 +44,11 @@ GLenum TransferToOpenGL(vtkm::cont::ArrayHandle<ValueType,StorageTag> handle,
template<typename ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void TransferToOpenGL(vtkm::cont::ArrayHandle<ValueType, StorageTag> handle,
GLuint& openGLHandle,
GLenum type,
BufferState& state,
DeviceAdapterTag)
{
vtkm::opengl::internal::TransferToOpenGL<ValueType, DeviceAdapterTag> toGL(type);
toGL.Transfer(handle,openGLHandle);
vtkm::opengl::internal::TransferToOpenGL<ValueType, DeviceAdapterTag> toGL(state);
return toGL.Transfer(handle);
}
}}

191
vtkm/opengl/cuda/internal/TransferToOpenGL.h
View File

@ -39,6 +39,94 @@ namespace vtkm {
namespace opengl {
namespace internal {
/// \brief cuda backend and opengl interop resource management
///
/// \c TransferResource manages cuda resource binding for a given buffer
///
///
class CudaTransferResource : public vtkm::opengl::internal::TransferResource
{
public:
CudaTransferResource():
vtkm::opengl::internal::TransferResource()
{
this->Registered = false;
}
~CudaTransferResource()
{
//unregister the buffer
if(this->Registered)
{
cudaGraphicsUnregisterResource(this->CudaResource);
}
}
bool IsRegistered() const { return Registered; }
void Register(GLuint* handle)
{
if(this->Registered)
{
//If you don't release the cuda resource before re-registering you
//will leak memory on the OpenGL side.
cudaGraphicsUnregisterResource(this->CudaResource);
}
this->Registered = true;
cudaError_t cError = cudaGraphicsGLRegisterBuffer(&this->CudaResource,
*handle,
cudaGraphicsMapFlagsWriteDiscard);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Could not register the OpenGL buffer handle to CUDA.");
}
}
void Map()
{
//map the resource into cuda, so we can copy it
cudaError_t cError =cudaGraphicsMapResources(1,&this->CudaResource);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorControlBadAllocation(
"Could not allocate enough memory in CUDA for OpenGL interop.");
}
}
template< typename ValueType >
ValueType* GetMappedPoiner( vtkm::Int64 desiredSize)
{
//get the mapped pointer
std::size_t cuda_size;
ValueType* pointer = NULL;
cudaError_t cError = cudaGraphicsResourceGetMappedPointer((void **)&pointer,
&cuda_size,
this->CudaResource);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Unable to get pointers to CUDA memory for OpenGL buffer.");
}
//assert that cuda_size is the same size as the buffer we created in OpenGL
VTKM_ASSERT_CONT(cuda_size >= desiredSize);
return pointer;
}
void UnMap()
{
cudaGraphicsUnmapResources(1, &this->CudaResource);
}
private:
bool Registered;
cudaGraphicsResource_t CudaResource;
};
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
@ -49,72 +137,65 @@ class TransferToOpenGL<ValueType, vtkm::cont::DeviceAdapterTagCuda>
{
typedef vtkm::cont::DeviceAdapterTagCuda DeviceAdapterTag;
public:
VTKM_CONT_EXPORT TransferToOpenGL():
Type( vtkm::opengl::internal::BufferTypePicker( ValueType() ) )
{}
VTKM_CONT_EXPORT explicit TransferToOpenGL(BufferState& state):
State(state),
Resource(NULL)
{
if( !this->State.HasType() )
{
this->State.DeduceAndSetType( ValueType() );
}
VTKM_CONT_EXPORT explicit TransferToOpenGL(GLenum type):
Type(type)
{}
this->Resource = dynamic_cast<vtkm::opengl::internal::CudaTransferResource*>
(state.GetResource());
if( !this->Resource )
{
vtkm::opengl::internal::CudaTransferResource* cudaResource =
new vtkm::opengl::internal::CudaTransferResource();
GLenum GetType() const { return this->Type; }
//reset the resource to be a cuda resource
this->State.SetResource( cudaResource );
this->Resource = cudaResource;
}
}
template< typename StorageTag >
VTKM_CONT_EXPORT
void Transfer (
vtkm::cont::ArrayHandle<ValueType, StorageTag> &handle,
GLuint& openGLHandle ) const
vtkm::cont::ArrayHandle<ValueType, StorageTag> &handle) const
{
//construct a cuda resource handle
cudaGraphicsResource_t cudaResource;
cudaError_t cError;
//make a buffer for the handle if the user has forgotten too
if(!glIsBuffer(openGLHandle))
{
glGenBuffers(1,&openGLHandle);
}
if(!glIsBuffer(*this->State.GetHandle()))
{
glGenBuffers(1, this->State.GetHandle());
}
//bind the buffer to the given buffer type
glBindBuffer(this->Type, openGLHandle);
glBindBuffer(this->State.GetType(), *this->State.GetHandle());
//Allocate the memory and set it as GL_DYNAMIC_DRAW draw
const vtkm::Id size = static_cast<vtkm::Id>(sizeof(ValueType))* handle.GetNumberOfValues();
glBufferData(this->Type, size, 0, GL_DYNAMIC_DRAW);
//Determine if we need to reallocate the buffer
const vtkm::Int64 size = static_cast<vtkm::Int64>(sizeof(ValueType))* handle.GetNumberOfValues();
this->State.SetSize(size);
const bool resize = this->State.ShouldRealloc(size);
if( resize )
{
//Allocate the memory and set it as GL_DYNAMIC_DRAW draw
glBufferData(this->State.GetType(), size, 0, GL_DYNAMIC_DRAW);
//register the buffer as being used by cuda
cError = cudaGraphicsGLRegisterBuffer(&cudaResource,
openGLHandle,
cudaGraphicsMapFlagsWriteDiscard);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Could not register the OpenGL buffer handle to CUDA.");
}
this->State.SetCapacity(size);
}
//map the resource into cuda, so we can copy it
cError =cudaGraphicsMapResources(1,&cudaResource);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorControlBadAllocation(
"Could not allocate enough memory in CUDA for OpenGL interop.");
}
if(!this->Resource->IsRegistered() || resize )
{
//register the buffer as being used by cuda. This needs to be done everytime
//we change the size of the buffer. That is why we only change the buffer
//size as infrequently as possible
this->Resource->Register(this->State.GetHandle());
}
//get the mapped pointer
std::size_t cuda_size;
ValueType* beginPointer=NULL;
cError = cudaGraphicsResourceGetMappedPointer((void **)&beginPointer,
&cuda_size,
cudaResource);
this->Resource->Map();
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Unable to get pointers to CUDA memory for OpenGL buffer.");
}
//assert that cuda_size is the same size as the buffer we created in OpenGL
VTKM_ASSERT_CONT(cuda_size == size);
ValueType* beginPointer= this->Resource->GetMappedPoiner<ValueType>(size);
//get the device pointers
typedef vtkm::cont::ArrayHandle<ValueType, StorageTag> HandleType;
@ -134,14 +215,12 @@ public:
thrust::cuda::pointer<ValueType>(beginPointer));
//unmap the resource
cudaGraphicsUnmapResources(1, &cudaResource);
//unregister the buffer
cudaGraphicsUnregisterResource(cudaResource);
this->Resource->UnMap();
}
private:
GLenum Type;
vtkm::opengl::BufferState& State;
vtkm::opengl::internal::CudaTransferResource* Resource;
};

4
vtkm/opengl/internal/BufferTypePicker.h
View File

@ -30,10 +30,10 @@ namespace internal {
/// helper function that guesses what OpenGL buffer type is the best default
/// given a primitive type. Currently GL_ELEMENT_ARRAY_BUFFER is used for integer
/// types, and GL_ARRAY_BUFFER is used for everything else
VTKM_CONT_EXPORT GLenum BufferTypePicker( int )
VTKM_CONT_EXPORT GLenum BufferTypePicker( vtkm::Int32 )
{ return GL_ELEMENT_ARRAY_BUFFER; }
VTKM_CONT_EXPORT GLenum BufferTypePicker( unsigned int )
VTKM_CONT_EXPORT GLenum BufferTypePicker( vtkm::UInt32 )
{ return GL_ELEMENT_ARRAY_BUFFER; }
#if VTKM_SIZE_LONG == 8

72
vtkm/opengl/internal/TransferToOpenGL.h
View File

@ -25,7 +25,9 @@
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/opengl/internal/OpenGLHeaders.h>
#include <vtkm/opengl/internal/BufferTypePicker.h>
#include <vtkm/opengl/BufferState.h>
namespace vtkm {
namespace opengl {
namespace internal {
@ -37,7 +39,7 @@ template<class ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void CopyFromHandle(
vtkm::cont::ArrayHandle<ValueType, StorageTag>& handle,
GLenum type,
vtkm::opengl::BufferState& state,
DeviceAdapterTag)
{
//Generic implementation that will work no matter what. We copy the data
@ -61,11 +63,19 @@ void CopyFromHandle(
//Note that the temporary ArrayHandle is no longer valid after this call
ValueType* temporaryStorage = tmpHandle.Internals->ControlArray.StealArray();
//Detach the current buffer
glBufferData(type, size, 0, GL_DYNAMIC_DRAW);
//Determine if we need to reallocate the buffer
state.SetSize(size);
const bool resize = state.ShouldRealloc(size);
if( resize )
{
//Allocate the memory and set it as GL_DYNAMIC_DRAW draw
glBufferData(state.GetType(), size, 0, GL_DYNAMIC_DRAW);
//Allocate the memory and set it as static draw and copy into opengl
glBufferSubData(type,0,size,temporaryStorage);
state.SetCapacity(size);
}
//copy into opengl buffer
glBufferSubData(state.GetType(),0,size,temporaryStorage);
delete[] temporaryStorage;
}
@ -74,7 +84,7 @@ template<class ValueType, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void CopyFromHandle(
vtkm::cont::ArrayHandle<ValueType, vtkm::cont::StorageTagBasic>& handle,
GLenum type,
vtkm::opengl::BufferState& state,
DeviceAdapterTag)
{
//Specialization given that we are use an C allocated array storage tag
@ -87,13 +97,21 @@ void CopyFromHandle(
static_cast<GLsizeiptr>(sizeof(ValueType)) *
static_cast<GLsizeiptr>(handle.GetNumberOfValues());
//Detach the current buffer
glBufferData(type, size, 0, GL_DYNAMIC_DRAW);
//Determine if we need to reallocate the buffer
state.SetSize(size);
const bool resize = state.ShouldRealloc(size);
if( resize )
{
//Allocate the memory and set it as GL_DYNAMIC_DRAW draw
glBufferData(state.GetType(), size, 0, GL_DYNAMIC_DRAW);
state.SetCapacity(size);
}
//Allocate the memory and set it as static draw and copy into opengl
const ValueType* memory = &(*vtkm::cont::ArrayPortalToIteratorBegin(
handle.PrepareForInput(DeviceAdapterTag())));
glBufferSubData(type,0,size,memory);
glBufferSubData(state.GetType(),0,size,memory);
}
@ -108,30 +126,28 @@ template<typename ValueType, class DeviceAdapterTag>
class TransferToOpenGL
{
public:
VTKM_CONT_EXPORT TransferToOpenGL():
Type( vtkm::opengl::internal::BufferTypePicker( ValueType() ) )
{}
VTKM_CONT_EXPORT explicit TransferToOpenGL(GLenum type):
Type(type)
{}
VTKM_CONT_EXPORT GLenum GetType() const { return this->Type; }
VTKM_CONT_EXPORT explicit TransferToOpenGL(BufferState& state):
State(state)
{
if( !this->State.HasType() )
{
this->State.DeduceAndSetType( ValueType() );
}
}
template< typename StorageTag >
VTKM_CONT_EXPORT
void Transfer (
vtkm::cont::ArrayHandle<ValueType, StorageTag>& handle,
GLuint& openGLHandle ) const
vtkm::cont::ArrayHandle<ValueType, StorageTag>& handle) const
{
//make a buffer for the handle if the user has forgotten too
if(!glIsBuffer(openGLHandle))
{
glGenBuffers(1,&openGLHandle);
}
if(!glIsBuffer(*this->State.GetHandle()))
{
glGenBuffers(1, this->State.GetHandle());
}
//bind the buffer to the given buffer type
glBindBuffer(this->Type, openGLHandle);
glBindBuffer(this->State.GetType(), *this->State.GetHandle());
//transfer the data.
//the primary concern that we have at this point is data locality and
@ -147,10 +163,10 @@ public:
//
//The end result is that we have CopyFromHandle which does number two
//from StorageTagBasic, and does the CopyInto for everything else
detail::CopyFromHandle(handle, this->Type, DeviceAdapterTag());
detail::CopyFromHandle(handle, this->State, DeviceAdapterTag());
}
private:
GLenum Type;
vtkm::opengl::BufferState& State;
};
}

6
vtkm/opengl/testing/TestingOpenGLInterop.h
View File

@ -75,7 +75,8 @@ private:
{
try
{
vtkm::opengl::TransferToOpenGL(array,handle, DeviceAdapterTag());
vtkm::opengl::BufferState state(handle);
vtkm::opengl::TransferToOpenGL(array, state, DeviceAdapterTag());
}
catch (vtkm::cont::ErrorControlBadAllocation error)
{
@ -97,7 +98,8 @@ private:
{
try
{
vtkm::opengl::TransferToOpenGL(array,handle,type, DeviceAdapterTag());
vtkm::opengl::BufferState state(handle, type);
vtkm::opengl::TransferToOpenGL(array, state, DeviceAdapterTag());
}
catch (vtkm::cont::ErrorControlBadAllocation error)
{

4
vtkm/testing/Testing.h
View File

@ -28,10 +28,6 @@
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/static_assert.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
#include <exception>
#include <iostream>
#include <sstream>

4
vtkm/worklet/CMakeLists.txt
View File

@ -23,6 +23,7 @@ set(headers
WorkletMapField.h
DispatcherMapTopology.h
WorkletMapTopology.h
AverageByKey.h
CellAverage.h
Clip.h
@ -32,6 +33,9 @@ set(headers
IsosurfaceUniformGrid.h
MarchingCubesDataTables.h
PointElevation.h
ScatterCounting.h
ScatterIdentity.h
ScatterUniform.h
TetrahedralizeExplicitGrid.h
TetrahedralizeUniformGrid.h
Threshold.h

23
vtkm/worklet/Clip.h
View File

@ -103,6 +103,20 @@ void swap(T &v1, T &v2)
v2 = temp;
}
template <typename T>
VTKM_EXEC_CONT_EXPORT
T Scale(const T &val, vtkm::Float64 scale)
{
return static_cast<T>(scale * static_cast<vtkm::Float64>(val));
}
template <typename T, vtkm::IdComponent NumComponents>
VTKM_EXEC_CONT_EXPORT
vtkm::Vec<T, NumComponents> Scale(const vtkm::Vec<T, NumComponents> &val,
vtkm::Float64 scale)
{
return val * scale;
}
template <typename DeviceAdapter>
class ExecutionConnectivityExplicit : vtkm::exec::ExecutionObjectBase
@ -261,7 +275,7 @@ public:
{
public:
typedef void ControlSignature(TopologyIn topology,
FieldInPoint<Scalar> scalars,
FieldInPoint<ScalarAll> scalars,
FieldOutCell<IdType> clipTableIdxs,
FieldOutCell<TypeClipStats> stats);
typedef void ExecutionSignature(_2, CellShape, FromCount, _3, _4);
@ -322,7 +336,7 @@ public:
{
public:
typedef void ControlSignature(TopologyIn topology,
FieldInPoint<Scalar> scalars,
FieldInPoint<ScalarAll> scalars,
FieldInCell<IdType> clipTableIdxs,
FieldInCell<TypeClipStats> cellSetIdxs,
ExecObject connectivityExplicit,
@ -623,9 +637,8 @@ public:
EdgeInterpolation ei = this->Interpolation.Get(idx);
T v1 = Field.Get(ei.Vertex1);
T v2 = Field.Get(ei.Vertex2);
typename VecTraits<T>::ComponentType w =
static_cast<typename VecTraits<T>::ComponentType>(ei.Weight);
Field.Set(this->NewPointsOffset + idx, (w * (v2 - v1)) + v1);
Field.Set(this->NewPointsOffset + idx,
internal::Scale(T(v2 - v1), ei.Weight) + v1);
}
private:

367
vtkm/worklet/IsosurfaceUniformGrid.h
View File

@ -21,20 +21,26 @@
#ifndef vtk_m_worklet_IsosurfaceUniformGrid_h
#define vtk_m_worklet_IsosurfaceUniformGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/Pair.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/exec/CellDerivative.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterCounting.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include "MarchingCubesDataTables.h"
namespace vtkm {
@ -51,197 +57,147 @@ public:
public:
typedef void ControlSignature(FieldInPoint<Scalar> inNodes,
TopologyIn topology,
FieldOutCell<IdType> outNumVertices);
typedef void ExecutionSignature(_1, _3);
FieldOutCell<> outNumTriangles,
ExecObject numTrianglesTable);
typedef void ExecutionSignature(_1, _3, _4);
typedef _2 InputDomain;
typedef vtkm::cont::ArrayHandle<vtkm::Id> IdArrayHandle;
typedef typename IdArrayHandle::ExecutionTypes<DeviceAdapter>::PortalConst IdPortalType;
IdPortalType VertexTable;
FieldType Isovalue;
VTKM_CONT_EXPORT
ClassifyCell(IdPortalType vTable, FieldType isovalue) :
VertexTable(vTable),
ClassifyCell(FieldType isovalue) :
Isovalue(isovalue)
{
}
template<typename InPointVecType>
template<typename InPointVecType,
typename NumTrianglesTablePortalType>
VTKM_EXEC_EXPORT
void operator()(const InPointVecType &pointValues,
vtkm::Id& numVertices) const
void operator()(const InPointVecType &fieldIn,
vtkm::IdComponent &numTriangles,
const NumTrianglesTablePortalType &numTrianglesTable) const
{
vtkm::Id caseNumber = (pointValues[0] > this->Isovalue);
caseNumber += (pointValues[1] > this->Isovalue)*2;
caseNumber += (pointValues[2] > this->Isovalue)*4;
caseNumber += (pointValues[3] > this->Isovalue)*8;
caseNumber += (pointValues[4] > this->Isovalue)*16;
caseNumber += (pointValues[5] > this->Isovalue)*32;
caseNumber += (pointValues[6] > this->Isovalue)*64;
caseNumber += (pointValues[7] > this->Isovalue)*128;
numVertices = this->VertexTable.Get(caseNumber) / 3;
vtkm::IdComponent caseNumber =
( (fieldIn[0] > this->Isovalue)
| (fieldIn[1] > this->Isovalue)<<1
| (fieldIn[2] > this->Isovalue)<<2
| (fieldIn[3] > this->Isovalue)<<3
| (fieldIn[4] > this->Isovalue)<<4
| (fieldIn[5] > this->Isovalue)<<5
| (fieldIn[6] > this->Isovalue)<<6
| (fieldIn[7] > this->Isovalue)<<7 );
numTriangles = numTrianglesTable.Get(caseNumber);
}
};
/// \brief Compute isosurface vertices and scalars
class IsoSurfaceGenerate : public vtkm::worklet::WorkletMapField
class IsoSurfaceGenerate : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
FieldIn<IdType> inputIteration);
typedef void ExecutionSignature(WorkIndex, _1, _2);
typedef _1 InputDomain;
typedef void ControlSignature(
TopologyIn topology, // Cell set
FieldInPoint<> fieldIn, // Input point field defining the contour
FieldInPoint<Vec3> pcoordIn, // Input point coordinates
FieldOutCell<> vertexOut, // Vertices for output triangles
// TODO: Have a better way to iterate over and interpolate fields
FieldInPoint<Scalar> scalarsIn, // Scalars to interpolate
FieldOutCell<> scalarsOut, // Interpolated scalars (one per tri vertex)
FieldOutCell<> normalsOut, // Estimated normals (one per tri vertex)
ExecObject TriTable // An array portal with the triangle table
);
typedef void ExecutionSignature(
CellShape, _2, _3, _4, _5, _6, _7, _8, VisitIndex);
const FieldType Isovalue;
vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
typedef typename vtkm::cont::ArrayHandle<FieldType>::template ExecutionTypes<DeviceAdapter>::PortalConst FieldPortalType;
FieldPortalType Field, Source;
typedef typename vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3> >::template ExecutionTypes<DeviceAdapter>::Portal VectorPortalType;
VectorPortalType Vertices;
VectorPortalType Normals;
typedef typename vtkm::cont::ArrayHandle<FieldType>::template ExecutionTypes<DeviceAdapter>::Portal OutputPortalType;
OutputPortalType Scalars;
typedef typename vtkm::cont::ArrayHandle<vtkm::Id> IdArrayHandle;
typedef typename IdArrayHandle::ExecutionTypes<DeviceAdapter>::PortalConst IdPortalType;
IdPortalType TriTable;
const vtkm::Id cellsPerLayer, pointsPerLayer;
template<typename U, typename W, typename X>
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
IsoSurfaceGenerate(FieldType ivalue, const vtkm::Id3 cdims, IdPortalType triTablePortal,
const U & field, const U & source, const W & vertices, const W & normals, const X & scalars) :
Isovalue(ivalue),
xdim(cdims[0]), ydim(cdims[1]), zdim(cdims[2]),
xmin(-1), ymin(-1), zmin(-1), xmax(1), ymax(1), zmax(1),
Field( field.PrepareForInput( DeviceAdapter() ) ),
Source( source.PrepareForInput( DeviceAdapter() ) ),
Vertices(vertices),
Normals(normals),
Scalars(scalars),
TriTable(triTablePortal),
cellsPerLayer(xdim * ydim),
pointsPerLayer ((xdim+1)*(ydim+1))
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
IsoSurfaceGenerate(FieldType isovalue,
const CountArrayType &countArray,
Device)
: Isovalue(isovalue), Scatter(countArray, Device()) { }
template<typename CellShapeTag,
typename FieldInType, // Vec-like, one per input point
typename CoordType, // Vec-like (one per input point) of Vec-3
typename VertexOutType, // Vec-3 of Vec-3 coordinates (for triangle)
typename ScalarInType, // Vec-like, one per input point
typename ScalarOutType, // Vec-3 (one value per tri vertex)
typename NormalOutType, // Vec-3 of Vec-3
typename TriTablePortalType> // Array portal
VTKM_EXEC_EXPORT
void operator()(vtkm::Id outputCellId, vtkm::Id inputCellId, vtkm::Id inputLowerBounds) const
void operator()(
CellShapeTag shape,
const FieldInType &fieldIn, // Input point field defining the contour
const CoordType &coords, // Input point coordinates
VertexOutType &vertexOut, // Vertices for output triangles
// TODO: Have a better way to iterate over and interpolate fields
const ScalarInType &scalarsIn, // Scalars to interpolate
ScalarOutType &scalarsOut, // Interpolated scalars (one per tri vertex)
NormalOutType &normalsOut, // Estimated normals (one per tri vertex)
const TriTablePortalType &triTable, // An array portal with the triangle table
vtkm::IdComponent visitIndex
) const
{
// Get data for this cell
const int verticesForEdge[] = { 0, 1, 1, 2, 3, 2, 0, 3,
4, 5, 5, 6, 7, 6, 4, 7,
0, 4, 1, 5, 2, 6, 3, 7 };
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
const vtkm::Id z = inputCellId / cellsPerLayer;
// Compute indices for the eight vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1) + z * pointsPerLayer;
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
const vtkm::Id i4 = i0 + pointsPerLayer;
const vtkm::Id i5 = i1 + pointsPerLayer;
const vtkm::Id i6 = i2 + pointsPerLayer;
const vtkm::Id i7 = i3 + pointsPerLayer;
// Get the field values at these eight vertices
FieldType f[8];
f[0] = this->Field.Get(i0);
f[1] = this->Field.Get(i1);
f[2] = this->Field.Get(i2);
f[3] = this->Field.Get(i3);
f[4] = this->Field.Get(i4);
f[5] = this->Field.Get(i5);
f[6] = this->Field.Get(i6);
f[7] = this->Field.Get(i7);
const vtkm::IdComponent verticesForEdge[] = { 0, 1, 1, 2, 3, 2, 0, 3,
4, 5, 5, 6, 7, 6, 4, 7,
0, 4, 1, 5, 2, 6, 3, 7 };
// Compute the Marching Cubes case number for this cell
unsigned int cubeindex = 0;
cubeindex += (f[0] > this->Isovalue);
cubeindex += (f[1] > this->Isovalue)*2;
cubeindex += (f[2] > this->Isovalue)*4;
cubeindex += (f[3] > this->Isovalue)*8;
cubeindex += (f[4] > this->Isovalue)*16;
cubeindex += (f[5] > this->Isovalue)*32;
cubeindex += (f[6] > this->Isovalue)*64;
cubeindex += (f[7] > this->Isovalue)*128;
// printf("inputCellId: %i \n",inputCellId);
// printf("x: %i, y: %i, z: %i \n",x, y, z);
// printf("i0: %i \n",i0);
// printf("f0 %F\n", f[0]);
// printf("cubeindex: %i \n",cubeindex);
// printf("numCells: %i \n",(vtkm::worklet::internal::numVerticesTable[cubeindex]/3) );
// Compute the coordinates of the uniform regular grid at each of the cell's eight vertices
vtkm::Vec<FieldType, 3> p[8];
// If we have offset and spacing, can we simplify this computation
{
vtkm::Vec<FieldType, 3> offset = vtkm::make_Vec(xmin+(xmax-xmin),
ymin+(ymax-ymin),
zmin+(zmax-zmin) );
vtkm::Vec<FieldType, 3> spacing = vtkm::make_Vec( 1.0f /((float)(xdim-1)),
1.0f /((float)(ydim-1)),
1.0f /((float)(zdim-1)));
vtkm::Vec<FieldType, 3> firstPoint = offset * spacing * vtkm::make_Vec( x, y, z );
vtkm::Vec<FieldType, 3> secondPoint = offset * spacing * vtkm::make_Vec( x+1, y+1, z+1 );
p[0] = vtkm::make_Vec( firstPoint[0], firstPoint[1], firstPoint[2]);
p[1] = vtkm::make_Vec( secondPoint[0], firstPoint[1], firstPoint[2]);
p[2] = vtkm::make_Vec( secondPoint[0], secondPoint[1], firstPoint[2]);
p[3] = vtkm::make_Vec( firstPoint[0], secondPoint[1], firstPoint[2]);
p[4] = vtkm::make_Vec( firstPoint[0], firstPoint[1], secondPoint[2]);
p[5] = vtkm::make_Vec( secondPoint[0], firstPoint[1], secondPoint[2]);
p[6] = vtkm::make_Vec( secondPoint[0], secondPoint[1], secondPoint[2]);
p[7] = vtkm::make_Vec( firstPoint[0], secondPoint[1], secondPoint[2]);
}
// Get the scalar source values at the eight vertices
FieldType s[8];
s[0] = this->Source.Get(i0);
s[1] = this->Source.Get(i1);
s[2] = this->Source.Get(i2);
s[3] = this->Source.Get(i3);
s[4] = this->Source.Get(i4);
s[5] = this->Source.Get(i5);
s[6] = this->Source.Get(i6);
s[7] = this->Source.Get(i7);
vtkm::IdComponent caseNumber =
( (fieldIn[0] > this->Isovalue)
| (fieldIn[1] > this->Isovalue)<<1
| (fieldIn[2] > this->Isovalue)<<2
| (fieldIn[3] > this->Isovalue)<<3
| (fieldIn[4] > this->Isovalue)<<4
| (fieldIn[5] > this->Isovalue)<<5
| (fieldIn[6] > this->Isovalue)<<6
| (fieldIn[7] > this->Isovalue)<<7 );
// Interpolate for vertex positions and associated scalar values
const vtkm::Id inputIteration = (outputCellId - inputLowerBounds);
const vtkm::Id outputVertId = outputCellId * 3;
const vtkm::Id cellOffset = static_cast<vtkm::Id>(cubeindex*16) + (inputIteration * 3);
for (int v = 0; v < 3; v++)
const vtkm::Id triTableOffset =
static_cast<vtkm::Id>(caseNumber*16 + visitIndex*3);
for (vtkm::IdComponent triVertex = 0;
triVertex < 3;
triVertex++)
{
const vtkm::Id edge = this->TriTable.Get(cellOffset + v);
const int v0 = verticesForEdge[2*edge];
const int v1 = verticesForEdge[2*edge + 1];
const FieldType t = (this->Isovalue - f[v0]) / (f[v1] - f[v0]);
this->Vertices.Set(outputVertId + v, vtkm::Lerp(p[v0], p[v1], t));
this->Scalars.Set(outputVertId + v, vtkm::Lerp(s[v0], s[v1], t));
const vtkm::IdComponent edgeIndex =
triTable.Get(triTableOffset + triVertex);
const vtkm::IdComponent edgeVertex0 = verticesForEdge[2*edgeIndex + 0];
const vtkm::IdComponent edgeVertex1 = verticesForEdge[2*edgeIndex + 1];
const FieldType fieldValue0 = fieldIn[edgeVertex0];
const FieldType fieldValue1 = fieldIn[edgeVertex1];
const FieldType interpolant =
(this->Isovalue - fieldValue0) / (fieldValue1 - fieldValue0);
vertexOut[triVertex] = vtkm::Lerp(coords[edgeVertex0],
coords[edgeVertex1],
interpolant);
scalarsOut[triVertex] = vtkm::Lerp(scalarsIn[edgeVertex0],
scalarsIn[edgeVertex1],
interpolant);
const vtkm::Vec<vtkm::FloatDefault,3> edgePCoord0 =
vtkm::exec::ParametricCoordinatesPoint(
fieldIn.GetNumberOfComponents(), edgeVertex0, shape, *this);
const vtkm::Vec<vtkm::FloatDefault,3> edgePCoord1 =
vtkm::exec::ParametricCoordinatesPoint(
fieldIn.GetNumberOfComponents(), edgeVertex1, shape, *this);
const vtkm::Vec<vtkm::FloatDefault,3> interpPCoord =
vtkm::Lerp(edgePCoord0, edgePCoord1, interpolant);
normalsOut[triVertex] =
vtkm::Normal(vtkm::exec::CellDerivative(
fieldIn, coords, interpPCoord, shape, *this));
}
vtkm::Vec<FieldType, 3> vertex0 = this->Vertices.Get(outputVertId + 0);
vtkm::Vec<FieldType, 3> vertex1 = this->Vertices.Get(outputVertId + 1);
vtkm::Vec<FieldType, 3> vertex2 = this->Vertices.Get(outputVertId + 2);
vtkm::Vec<FieldType, 3> curNorm = vtkm::Cross(vertex1-vertex0, vertex2-vertex0);
vtkm::Normalize(curNorm);
this->Normals.Set(outputVertId + 0, curNorm);
this->Normals.Set(outputVertId + 1, curNorm);
this->Normals.Set(outputVertId + 2, curNorm);
}
private:
const FieldType Isovalue;
ScatterType Scatter;
};
@ -276,68 +232,49 @@ public:
vtkm::cont::ArrayHandle< vtkm::Vec<CoordinateType,3> > normalsArray,
vtkm::cont::ArrayHandle<FieldType> scalarsArray)
{
typedef typename vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> DeviceAlgorithms;
// Set up the Marching Cubes case tables
vtkm::cont::ArrayHandle<vtkm::Id> vertexTableArray =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numVerticesTable,
vtkm::cont::ArrayHandle<vtkm::IdComponent> numTrianglesTable =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numTrianglesTable,
256);
vtkm::cont::ArrayHandle<vtkm::Id> triangleTableArray =
vtkm::cont::ArrayHandle<vtkm::IdComponent> triangleTableArray =
vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::triTable,
256*16);
typedef vtkm::exec::ExecutionWholeArrayConst<vtkm::IdComponent, VTKM_DEFAULT_STORAGE_TAG, DeviceAdapter>
TableArrayExecObjectType;
// Call the ClassifyCell functor to compute the Marching Cubes case numbers
// for each cell, and the number of vertices to be generated
ClassifyCell classifyCell(vertexTableArray.PrepareForInput(DeviceAdapter()),
isovalue);
ClassifyCell classifyCell(isovalue);
typedef typename vtkm::worklet::DispatcherMapTopology<
ClassifyCell,
DeviceAdapter> ClassifyCellDispatcher;
ClassifyCellDispatcher classifyCellDispatcher(classifyCell);
vtkm::cont::ArrayHandle<vtkm::Id> numOutputTrisPerCell;
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutputTrisPerCell;
classifyCellDispatcher.Invoke( field,
this->DataSet.GetCellSet(0),
numOutputTrisPerCell);
numOutputTrisPerCell,
TableArrayExecObjectType(numTrianglesTable));
// Compute the number of valid input cells and those ids
const vtkm::Id numOutputCells = DeviceAlgorithms::ScanInclusive(numOutputTrisPerCell,
numOutputTrisPerCell);
IsoSurfaceGenerate isosurface(isovalue, numOutputTrisPerCell, DeviceAdapter());
// Terminate if no cells has triangles left
if (numOutputCells == 0) return;
vtkm::cont::ArrayHandle<vtkm::Id> validCellIndicesArray, inputCellIterationNumber;
vtkm::cont::ArrayHandleIndex validCellCountImplicitArray(numOutputCells);
DeviceAlgorithms::UpperBounds(numOutputTrisPerCell,
validCellCountImplicitArray,
validCellIndicesArray);
numOutputTrisPerCell.ReleaseResources();
// Compute for each output triangle what iteration of the input cell generates it
DeviceAlgorithms::LowerBounds(validCellIndicesArray,
validCellIndicesArray,
inputCellIterationNumber);
// Generate a single triangle per cell
const vtkm::Id numTotalVertices = numOutputCells * 3;
IsoSurfaceGenerate isosurface(isovalue,
this->CDims,
triangleTableArray.PrepareForInput(DeviceAdapter()),
field,
field,
verticesArray.PrepareForOutput(numTotalVertices, DeviceAdapter()),
normalsArray.PrepareForOutput(numTotalVertices, DeviceAdapter()),
scalarsArray.PrepareForOutput(numTotalVertices, DeviceAdapter())
);
typedef typename vtkm::worklet::DispatcherMapField< IsoSurfaceGenerate,
DeviceAdapter> IsoSurfaceDispatcher;
IsoSurfaceDispatcher isosurfaceDispatcher(isosurface);
isosurfaceDispatcher.Invoke(validCellIndicesArray,
inputCellIterationNumber);
vtkm::worklet::DispatcherMapTopology<IsoSurfaceGenerate, DeviceAdapter>
isosurfaceDispatcher(isosurface);
isosurfaceDispatcher.Invoke(
// Currently forcing cell set to be structured. Eventually we should
// relax this as we support other grid types.
this->DataSet.GetCellSet(0).ResetCellSetList(
vtkm::ListTagBase<vtkm::cont::CellSetStructured<3> >()),
field,
this->DataSet.GetCoordinateSystem(0).GetData(),
vtkm::cont::make_ArrayHandleGroupVec<3>(verticesArray),
field, // This is silly. The field will interp to isovalue
vtkm::cont::make_ArrayHandleGroupVec<3>(scalarsArray),
vtkm::cont::make_ArrayHandleGroupVec<3>(normalsArray),
TableArrayExecObjectType(triangleTableArray)
);
}
};

498
vtkm/worklet/MarchingCubesDataTables.h
View File

@ -27,267 +27,267 @@ namespace worklet {
namespace internal {
const vtkm::Id numVerticesTable[256] = {
const vtkm::IdComponent numTrianglesTable[256] = {
0,
1,
1,
2,
1,
2,
2,
3,
1,
2,
2,
3,
2,
3,
3,
6,
2,
1,
2,
2,
3,
6,
6,
9,
3,
6,
6,
9,
6,
9,
9,
6,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
6,
9,
12,
12,
9,
9,
12,
12,
9,
12,
15,
15,
6,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
12,
15,
15,
12,
6,
9,
9,
12,
9,
12,
6,
9,
9,
12,
12,
15,
12,
15,
9,
6,
9,
12,
12,
9,
12,
15,
9,
6,
12,
15,
15,
12,
15,
6,
12,
2,
3,
3,
6,
6,
9,
6,
9,
9,
12,
6,
9,
9,
12,
9,
12,
12,
9,
6,
9,
9,
12,
9,
12,
12,
15,
9,
6,
12,
9,
12,
9,
15,
6,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
12,
15,
15,
12,
9,
12,
12,
9,
12,
15,
15,
12,
12,
9,
15,
6,
15,
12,
6,
3,
6,
9,
9,
12,
9,
12,
12,
15,
9,
12,
12,
15,
6,
9,
9,
6,
9,
12,
12,
15,
12,
15,
15,
6,
12,
9,
15,
12,
9,
6,
12,
3,
9,
12,
12,
15,
12,
15,
9,
12,
12,
15,
15,
6,
9,
12,
6,
3,
6,
9,
9,
6,
9,
12,
6,
3,
9,
6,
12,
3,
6,
4,
2,
3,
3,
4,
3,
4,
4,
3,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
2,
3,
4,
4,
3,
3,
4,
4,
3,
4,
5,
5,
2,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
4,
5,
5,
4,
2,
3,
3,
4,
3,
4,
2,
3,
3,
4,
4,
5,
4,
5,
3,
2,
3,
4,
4,
3,
4,
5,
3,
2,
4,
5,
5,
4,
5,
2,
4,
1,
1,
2,
2,
3,
2,
3,
3,
4,
2,
3,
3,
4,
3,
4,
4,
3,
2,
3,
3,
4,
3,
4,
4,
5,
3,
2,
4,
3,
4,
3,
5,
2,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
4,
5,
5,
4,
3,
4,
4,
3,
4,
5,
5,
4,
4,
3,
5,
2,
5,
4,
2,
1,
2,
3,
3,
4,
3,
4,
4,
5,
3,
4,
4,
5,
2,
3,
3,
2,
3,
4,
4,
5,
4,
5,
5,
2,
4,
3,
5,
4,
3,
2,
4,
1,
3,
4,
4,
5,
4,
5,
3,
4,
4,
5,
5,
2,
3,
4,
2,
1,
2,
3,
3,
2,
3,
4,
2,
1,
3,
2,
4,
1,
2,
1,
1,
0,
};
const vtkm::Id triTable[256*16] =
const vtkm::IdComponent triTable[256*16] =
{
#define X -1
X, X, X, X, X, X, X, X, X, X, X, X, X, X, X, X,

280
vtkm/worklet/ScatterCounting.h Normal file
View File

@ -0,0 +1,280 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_worklet_ScatterCounting_h
#define vtk_m_worklet_ScatterCounting_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/ErrorControlBadValue.h>
#include <vtkm/exec/FunctorBase.h>
#include <sstream>
namespace vtkm {
namespace worklet {
namespace detail {
template<typename Device>
struct ReverseInputToOutputMapKernel : vtkm::exec::FunctorBase
{
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::PortalConst
InputMapType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::Portal
OutputMapType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::IdComponent>::ExecutionTypes<Device>::Portal
VisitType;
InputMapType InputToOutputMap;
OutputMapType OutputToInputMap;
VisitType Visit;
vtkm::Id OutputSize;
VTKM_CONT_EXPORT
ReverseInputToOutputMapKernel(const InputMapType &inputToOutputMap,
const OutputMapType &outputToInputMap,
const VisitType &visit,
vtkm::Id outputSize)
: InputToOutputMap(inputToOutputMap),
OutputToInputMap(outputToInputMap),
Visit(visit),
OutputSize(outputSize)
{ }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex) const
{
vtkm::Id outputStartIndex;
if (inputIndex > 0)
{
outputStartIndex = this->InputToOutputMap.Get(inputIndex-1);
}
else
{
outputStartIndex = 0;
}
vtkm::Id outputEndIndex = this->InputToOutputMap.Get(inputIndex);
vtkm::IdComponent visitIndex = 0;
for (vtkm::Id outputIndex = outputStartIndex;
outputIndex < outputEndIndex;
outputIndex++)
{
this->OutputToInputMap.Set(outputIndex, inputIndex);
this->Visit.Set(outputIndex, visitIndex);
visitIndex++;
}
}
};
template<typename Device>
struct SubtractToVisitIndexKernel : vtkm::exec::FunctorBase
{
typedef typename
vtkm::cont::ArrayHandle<vtkm::Id>::ExecutionTypes<Device>::PortalConst
StartsOfGroupsType;
typedef typename
vtkm::cont::ArrayHandle<vtkm::IdComponent>::ExecutionTypes<Device>::Portal
VisitType;
StartsOfGroupsType StartsOfGroups;
VisitType Visit;
VTKM_CONT_EXPORT
SubtractToVisitIndexKernel(const StartsOfGroupsType &startsOfGroups,
const VisitType &visit)
: StartsOfGroups(startsOfGroups), Visit(visit)
{ }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex) const
{
vtkm::Id startOfGroup = this->StartsOfGroups.Get(inputIndex);
vtkm::IdComponent visitIndex =
static_cast<vtkm::IdComponent>(inputIndex - startOfGroup);
this->Visit.Set(inputIndex, visitIndex);
}
};
} // namespace detail
/// \brief A scatter that maps input to some numbers of output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterCounting establishes a 1 to
/// N mapping from input to output. That is, every input element generates 0 or
/// more output elements associated with it. The output elements are grouped by
/// the input associated.
///
/// A counting scatter takes an array of counts for each input. The data is
/// taken in the constructor and the index arrays are derived from that. So
/// changing the counts after the scatter is created will have no effect.
///
struct ScatterCounting
{
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
ScatterCounting(const CountArrayType &countArray, Device)
{
this->BuildArrays(countArray, Device());
}
typedef vtkm::cont::ArrayHandle<vtkm::Id> OutputToInputMapType;
template<typename RangeType>
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(RangeType) const
{
return this->OutputToInputMap;
}
typedef vtkm::cont::ArrayHandle<vtkm::IdComponent> VisitArrayType;
template<typename RangeType>
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(RangeType) const
{
return this->VisitArray;
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id inputRange) const
{
if (inputRange != this->InputRange)
{
std::stringstream msg;
msg << "ScatterCounting initialized with input domain of size "
<< this->InputRange
<< " but used with a worklet invoke of size "
<< inputRange << std::endl;
throw vtkm::cont::ErrorControlBadValue(msg.str());
}
return this->VisitArray.GetNumberOfValues();
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id3 inputRange) const
{
return this->GetOutputRange(inputRange[0]*inputRange[1]*inputRange[2]);
}
private:
vtkm::Id InputRange;
OutputToInputMapType OutputToInputMap;
VisitArrayType VisitArray;
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
void BuildArrays(const CountArrayType &count, Device)
{
VTKM_IS_ARRAY_HANDLE(CountArrayType);
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
this->InputRange = count.GetNumberOfValues();
// Currently we are treating the input to output map as a temporary
// variable. However, it is possible that this could, be useful elsewhere,
// so we may want to save this and make it available.
//
// The input to output map is actually built off by one. The first entry
// is actually for the second value. The last entry is the total number of
// output. This off-by-one is so that an upper bound find will work when
// building the output to input map.
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap;
vtkm::Id outputSize =
vtkm::cont::DeviceAdapterAlgorithm<Device>::ScanInclusive(
vtkm::cont::make_ArrayHandleCast(count, vtkm::Id()),
inputToOutputMap);
// We have implemented two different ways to compute the output to input
// map. The first way is to use a binary search on each output index into
// the input map. The second way is to schedule on each input and
// iteratively fill all the output indices for that input. The first way is
// faster for output sizes that are small relative to the input (typical in
// Marching Cubes, for example) and also tends to be well load balanced.
// The second way is faster for larger outputs (typical in triangulation,
// for example). We will use the first method for small output sizes and
// the second for large output sizes. Toying with this might be a good
// place for optimization.
if (outputSize < this->InputRange)
{
this->BuildOutputToInputMapWithFind(
outputSize, inputToOutputMap, Device());
}
else
{
this->BuildOutputToInputMapWithIterate(
outputSize, inputToOutputMap, Device());
}
}
template<typename Device>
VTKM_CONT_EXPORT
void BuildOutputToInputMapWithFind(
vtkm::Id outputSize,
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap,
Device)
{
vtkm::cont::ArrayHandleIndex outputIndices(outputSize);
vtkm::cont::DeviceAdapterAlgorithm<Device>::UpperBounds(
inputToOutputMap, outputIndices, this->OutputToInputMap);
// Do not need this anymore.
inputToOutputMap.ReleaseResources();
vtkm::cont::ArrayHandle<vtkm::Id> startsOfGroups;
// This find gives the index of the start of a group.
vtkm::cont::DeviceAdapterAlgorithm<Device>::LowerBounds(
this->OutputToInputMap, this->OutputToInputMap, startsOfGroups);
detail::SubtractToVisitIndexKernel<Device>
kernel(startsOfGroups.PrepareForInput(Device()),
this->VisitArray.PrepareForOutput(outputSize, Device()));
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(kernel, outputSize);
}
template<typename Device>
VTKM_CONT_EXPORT
void BuildOutputToInputMapWithIterate(
vtkm::Id outputSize,
vtkm::cont::ArrayHandle<vtkm::Id> inputToOutputMap,
Device)
{
detail::ReverseInputToOutputMapKernel<Device>
kernel(inputToOutputMap.PrepareForInput(Device()),
this->OutputToInputMap.PrepareForOutput(outputSize, Device()),
this->VisitArray.PrepareForOutput(outputSize, Device()),
outputSize);
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(
kernel, inputToOutputMap.GetNumberOfValues());
}
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterCounting_h

77
vtkm/worklet/ScatterIdentity.h Normal file
View File

@ -0,0 +1,77 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_worklet_ScatterIdentity_h
#define vtk_m_worklet_ScatterIdentity_h
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/cont/ArrayHandleIndex.h>
namespace vtkm {
namespace worklet {
/// \brief A scatter that maps input directly to output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterIdentity establishes a 1 to
/// 1 mapping from input to output (and vice versa). That is, every input
/// element generates one output element associated with it. This is the
/// default for basic maps.
///
struct ScatterIdentity
{
typedef vtkm::cont::ArrayHandleIndex OutputToInputMapType;
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(vtkm::Id outputRange) const
{
return OutputToInputMapType(outputRange);
}
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(vtkm::Id3 outputRange) const
{
return this->GetOutputToInputMap(
outputRange[0]*outputRange[1]*outputRange[2]);
}
typedef vtkm::cont::ArrayHandleConstant<vtkm::IdComponent> VisitArrayType;
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(vtkm::Id outputRange) const
{
return VisitArrayType(1, outputRange);
}
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(vtkm::Id3 outputRange) const
{
return this->GetVisitArray(outputRange[0]*outputRange[1]*outputRange[2]);
}
template<typename RangeType>
VTKM_CONT_EXPORT
RangeType GetOutputRange(RangeType inputRange) const
{
return inputRange;
}
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterIdentity_h

121
vtkm/worklet/ScatterUniform.h Normal file
View File

@ -0,0 +1,121 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_worklet_ScatterUniform_h
#define vtk_m_worklet_ScatterUniform_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
namespace vtkm {
namespace worklet {
namespace detail {
struct FunctorModulus
{
vtkm::IdComponent Modulus;
VTKM_EXEC_CONT_EXPORT
FunctorModulus(vtkm::IdComponent modulus = 1)
: Modulus(modulus)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent operator()(vtkm::Id index) const
{
return static_cast<vtkm::IdComponent>(index % this->Modulus);
}
};
struct FunctorDiv
{
vtkm::Id Divisor;
VTKM_EXEC_CONT_EXPORT
FunctorDiv(vtkm::Id divisor = 1)
: Divisor(divisor)
{ }
VTKM_EXEC_CONT_EXPORT
vtkm::Id operator()(vtkm::Id index) const
{
return index / this->Divisor;
}
};
}
/// \brief A scatter that maps input to some constant numbers of output.
///
/// The \c Scatter classes are responsible for defining how much output is
/// generated based on some sized input. \c ScatterUniform establishes a 1 to N
/// mapping from input to output. That is, every input element generates N
/// elements associated with it where N is the same for every input. The output
/// elements are grouped by the input associated.
///
struct ScatterUniform
{
VTKM_CONT_EXPORT
ScatterUniform(vtkm::IdComponent numOutputsPerInput)
: NumOutputsPerInput(numOutputsPerInput)
{ }
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id inputRange) const
{
return inputRange * this->NumOutputsPerInput;
}
VTKM_CONT_EXPORT
vtkm::Id GetOutputRange(vtkm::Id3 inputRange) const
{
return this->GetOutputRange(inputRange[0]*inputRange[1]*inputRange[2]);
}
typedef vtkm::cont::ArrayHandleImplicit<vtkm::Id, detail::FunctorDiv>
OutputToInputMapType;
template<typename RangeType>
VTKM_CONT_EXPORT
OutputToInputMapType GetOutputToInputMap(RangeType inputRange) const
{
return OutputToInputMapType(detail::FunctorDiv(this->NumOutputsPerInput),
this->GetOutputRange(inputRange));
}
typedef vtkm::cont::ArrayHandleImplicit<vtkm::IdComponent, detail::FunctorModulus>
VisitArrayType;
template<typename RangeType>
VTKM_CONT_EXPORT
VisitArrayType GetVisitArray(RangeType inputRange) const
{
return VisitArrayType(detail::FunctorModulus(this->NumOutputsPerInput),
this->GetOutputRange(inputRange));
}
private:
vtkm::IdComponent NumOutputsPerInput;
};
}
} // namespace vtkm::worklet
#endif //vtk_m_worklet_ScatterUniform_h

324
vtkm/worklet/TetrahedralizeExplicitGrid.h
View File

@ -6,9 +6,9 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -21,20 +21,21 @@
#ifndef vtk_m_worklet_TetrahedralizeExplicitGrid_h
#define vtk_m_worklet_TetrahedralizeExplicitGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/CellSetExplicit.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterCounting.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/worklet/internal/TriangulateTables.h>
namespace vtkm {
namespace worklet {
@ -52,26 +53,22 @@ public:
{
public:
typedef void ControlSignature(FieldIn<> shapes,
FieldIn<> numIndices,
FieldIn<> numPoints,
ExecObject tables,
FieldOut<> triangleCount);
typedef void ExecutionSignature(_1,_2,_3);
typedef _4 ExecutionSignature(_1,_2,_3);
typedef _1 InputDomain;
VTKM_CONT_EXPORT
TrianglesPerCell() {}
VTKM_EXEC_EXPORT
void operator()(const vtkm::UInt8 &shape,
const vtkm::IdComponent &numIndices,
vtkm::Id &triangleCount) const
vtkm::IdComponent operator()(
vtkm::UInt8 shape,
vtkm::IdComponent numPoints,
const vtkm::worklet::internal::TriangulateTablesExecutionObject<DeviceAdapter> &tables) const
{
if (shape == vtkm::CELL_SHAPE_TRIANGLE)
triangleCount = 1;
else if (shape == vtkm::CELL_SHAPE_QUAD)
triangleCount = 2;
else if (shape == vtkm::CELL_SHAPE_POLYGON)
triangleCount = numIndices - 2;
else triangleCount = 0;
return tables.GetCount(vtkm::CellShapeTagGeneric(shape), numPoints);
}
};
@ -82,26 +79,20 @@ public:
{
public:
typedef void ControlSignature(FieldIn<> shapes,
ExecObject tables,
FieldOut<> triangleCount);
typedef void ExecutionSignature(_1,_2);
typedef _3 ExecutionSignature(_1, _2);
typedef _1 InputDomain;
VTKM_CONT_EXPORT
TetrahedraPerCell() {}
VTKM_EXEC_EXPORT
void operator()(const vtkm::UInt8 &shape,
vtkm::Id &tetrahedraCount) const
vtkm::IdComponent operator()(
vtkm::UInt8 shape,
const vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<DeviceAdapter> &tables) const
{
if (shape == vtkm::CELL_SHAPE_TETRA)
tetrahedraCount = 1;
else if (shape == vtkm::CELL_SHAPE_HEXAHEDRON)
tetrahedraCount = 5;
else if (shape == vtkm::CELL_SHAPE_WEDGE)
tetrahedraCount = 3;
else if (shape == vtkm::CELL_SHAPE_PYRAMID)
tetrahedraCount = 2;
else tetrahedraCount = 0;
return tables.GetCount(vtkm::CellShapeTagGeneric(shape));
}
};
@ -112,53 +103,46 @@ public:
class TriangulateCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldInTo<> triangleOffset,
FieldInTo<> numIndices,
TopologyIn topology,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2,_4, CellShape, FromIndices);
typedef _3 InputDomain;
typedef void ControlSignature(TopologyIn topology,
ExecObject tables,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(CellShape, PointIndices, _2, _3, VisitIndex);
typedef _1 InputDomain;
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
TriangulateCell() {}
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CountArrayType>
VTKM_CONT_EXPORT
TriangulateCell(const CountArrayType &countArray)
: Scatter(countArray, DeviceAdapter())
{ }
// Each cell produces triangles and write result at the offset
template<typename CellShapeTag, typename CellNodeVecType>
template<typename CellShapeTag,
typename ConnectivityInVec,
typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &offset,
const vtkm::Id &numIndices,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity,
CellShapeTag shape,
const CellNodeVecType &cellNodeIds) const
void operator()(
CellShapeTag shape,
const ConnectivityInVec &connectivityIn,
const vtkm::worklet::internal::TriangulateTablesExecutionObject<DeviceAdapter> &tables,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Offset is in triangles, 3 vertices per triangle needed
vtkm::Id startIndex = offset * 3;
if (shape.Id == vtkm::CELL_SHAPE_TRIANGLE)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
}
else if (shape.Id == vtkm::CELL_SHAPE_QUAD)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
}
else if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
for (vtkm::IdComponent tri = 0; tri < numIndices-2; tri++)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[tri+1]);
connectivity.Set(startIndex++, cellNodeIds[tri+2]);
}
}
vtkm::Vec<vtkm::IdComponent,3> triIndices =
tables.GetIndices(shape, visitIndex);
connectivityOut[0] = connectivityIn[triIndices[0]];
connectivityOut[1] = connectivityIn[triIndices[1]];
connectivityOut[2] = connectivityIn[triIndices[2]];
}
private:
ScatterType Scatter;
};
//
@ -168,89 +152,46 @@ public:
class TetrahedralizeCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldInTo<> tetraOffset,
TopologyIn topology,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_3, CellShape, FromIndices);
typedef _2 InputDomain;
typedef void ControlSignature(TopologyIn topology,
ExecObject tables,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(CellShape, PointIndices, _2, _3, VisitIndex);
typedef _1 InputDomain;
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
TetrahedralizeCell() {}
ScatterType GetScatter() const
{
return this->Scatter;
}
template<typename CellArrayType>
VTKM_CONT_EXPORT
TetrahedralizeCell(const CellArrayType &cellArray)
: Scatter(cellArray, DeviceAdapter())
{ }
// Each cell produces tetrahedra and write result at the offset
template<typename CellShapeTag, typename CellNodeVecType>
template<typename CellShapeTag,
typename ConnectivityInVec,
typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &offset,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity,
CellShapeTag shape,
const CellNodeVecType &cellNodeIds) const
void operator()(CellShapeTag shape,
const ConnectivityInVec &connectivityIn,
const vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<DeviceAdapter> &tables,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Offset is in tetrahedra, 4 vertices per tetrahedron needed
vtkm::Id startIndex = offset * 4;
if (shape.Id == vtkm::CELL_SHAPE_TETRA)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
}
else if (shape.Id == vtkm::CELL_SHAPE_HEXAHEDRON)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[5]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[6]);
connectivity.Set(startIndex++, cellNodeIds[7]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
else if (shape.Id == vtkm::CELL_SHAPE_WEDGE)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[5]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
else if (shape.Id == vtkm::CELL_SHAPE_PYRAMID)
{
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[1]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[4]);
connectivity.Set(startIndex++, cellNodeIds[0]);
connectivity.Set(startIndex++, cellNodeIds[2]);
connectivity.Set(startIndex++, cellNodeIds[3]);
connectivity.Set(startIndex++, cellNodeIds[4]);
}
vtkm::Vec<vtkm::IdComponent,4> tetIndices =
tables.GetIndices(shape, visitIndex);
connectivityOut[0] = connectivityIn[tetIndices[0]];
connectivityOut[1] = connectivityIn[tetIndices[1]];
connectivityOut[2] = connectivityIn[tetIndices[2]];
connectivityOut[3] = connectivityIn[tetIndices[3]];
}
private:
ScatterType Scatter;
};
//
@ -270,8 +211,6 @@ public:
//
void Run()
{
typedef typename vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> DeviceAlgorithms;
// Cell sets belonging to input and output datasets
vtkm::cont::CellSetExplicit<> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetExplicit<> >();
@ -286,56 +225,61 @@ public:
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
vtkm::cont::ArrayHandle<vtkm::IdComponent> inNumIndices = inCellSet.GetNumIndicesArray(
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
vtkm::cont::ArrayHandle<vtkm::Id> inConn = inCellSet.GetConnectivityArray(
vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::Id> numOutCellArray;
vtkm::IdComponent verticesPerOutCell = 0;
// Output topology
vtkm::cont::ArrayHandle<vtkm::Id> outConnectivity;
if (dimensionality == 2)
{
verticesPerOutCell = 3;
vtkm::worklet::DispatcherMapField<TrianglesPerCell> trianglesPerCellDispatcher;
trianglesPerCellDispatcher.Invoke(inShapes, inNumIndices, numOutCellArray);
vtkm::worklet::internal::TriangulateTables tables;
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutCellArray;
vtkm::worklet::DispatcherMapField<TrianglesPerCell,DeviceAdapter>
triPerCellDispatcher;
triPerCellDispatcher.Invoke(inShapes,
inNumIndices,
tables.PrepareForInput(DeviceAdapter()),
numOutCellArray);
// Build new cells
TriangulateCell triangulateWorklet(numOutCellArray);
vtkm::worklet::DispatcherMapTopology<TriangulateCell,DeviceAdapter>
triangulateDispatcher(triangulateWorklet);
triangulateDispatcher.Invoke(
inCellSet,
tables.PrepareForInput(DeviceAdapter()),
vtkm::cont::make_ArrayHandleGroupVec<3>(outConnectivity));
}
else if (dimensionality == 3)
{
verticesPerOutCell = 4;
vtkm::worklet::DispatcherMapField<TetrahedraPerCell> tetrahedraPerCellDispatcher;
tetrahedraPerCellDispatcher.Invoke(inShapes, numOutCellArray);
vtkm::worklet::internal::TetrahedralizeTables tables;
// Determine the number of output cells each input cell will generate
vtkm::cont::ArrayHandle<vtkm::IdComponent> numOutCellArray;
vtkm::worklet::DispatcherMapField<TetrahedraPerCell,DeviceAdapter>
tetPerCellDispatcher;
tetPerCellDispatcher.Invoke(inShapes,
tables.PrepareForInput(DeviceAdapter()),
numOutCellArray);
// Build new cells
TetrahedralizeCell tetrahedralizeWorklet(numOutCellArray);
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell,DeviceAdapter>
tetrahedralizeDispatcher(tetrahedralizeWorklet);
tetrahedralizeDispatcher.Invoke(
inCellSet,
tables.PrepareForInput(DeviceAdapter()),
vtkm::cont::make_ArrayHandleGroupVec<4>(outConnectivity));
}
// Number of output cells and number of vertices needed
vtkm::cont::ArrayHandle<vtkm::Id> cellOffset;
vtkm::Id numberOfOutCells = DeviceAlgorithms::ScanExclusive(numOutCellArray,
cellOffset);
vtkm::Id numberOfOutIndices = numberOfOutCells * verticesPerOutCell;
// Information needed to build the output cell set
vtkm::cont::ArrayHandle<vtkm::Id> connectivity;
connectivity.Allocate(numberOfOutIndices);
// Call worklet to compute the connectivity
if (dimensionality == 2)
else
{
vtkm::worklet::DispatcherMapTopology<TriangulateCell> triangulateCellDispatcher;
triangulateCellDispatcher.Invoke(
cellOffset,
inNumIndices,
inCellSet,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
}
else if (dimensionality == 3)
{
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell> tetrahedralizeCellDispatcher;
tetrahedralizeCellDispatcher.Invoke(
cellOffset,
inCellSet,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
throw vtkm::cont::ErrorControlBadValue(
"Unsupported dimensionality for TetrahedralizeExplicitGrid.");
}
// Add cells to output cellset
cellSet.Fill(connectivity);
cellSet.Fill(outConnectivity);
}
};

296
vtkm/worklet/TetrahedralizeUniformGrid.h
View File

@ -6,9 +6,9 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -21,23 +21,51 @@
#ifndef vtk_m_worklet_TetrahedralizeUniformGrid_h
#define vtk_m_worklet_TetrahedralizeUniformGrid_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/CellSetSingleType.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DynamicArrayHandle.h>
#include <vtkm/cont/ErrorControlBadValue.h>
#include <vtkm/cont/Field.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/exec/ExecutionWholeArray.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterUniform.h>
#include <vtkm/worklet/WorkletMapTopology.h>
namespace vtkm {
namespace worklet {
namespace detail {
VTKM_EXEC_CONSTANT_EXPORT
const static vtkm::IdComponent StructuredTriangleIndices[2][3] = {
{ 0, 1, 2 },
{ 0, 2, 3 }
};
VTKM_EXEC_CONSTANT_EXPORT
const static vtkm::IdComponent StructuredTetrahedronIndices[2][5][4] = {
{
{ 0, 1, 3, 4 },
{ 1, 4, 5, 6 },
{ 1, 4, 6, 3 },
{ 1, 3, 6, 2 },
{ 3, 6, 7, 4 }
},
{
{ 2, 1, 5, 0 },
{ 0, 2, 3, 7 },
{ 2, 5, 6, 7 },
{ 0, 7, 4, 5 },
{ 0, 2, 7, 5 }
}
};
} // namespace detail
/// \brief Compute the tetrahedralize cells for a uniform grid data set
template <typename DeviceAdapter>
class TetrahedralizeFilterUniformGrid
@ -48,63 +76,35 @@ public:
// Worklet to turn quads into triangles
// Vertices remain the same and each cell is processed with needing topology
//
class TriangulateCell : public vtkm::worklet::WorkletMapField
class TriangulateCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2);
typedef void ControlSignature(TopologyIn topology,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(PointIndices, _2, VisitIndex);
typedef _1 InputDomain;
vtkm::Id xdim, ydim;
typedef vtkm::worklet::ScatterUniform ScatterType;
VTKM_CONT_EXPORT
TriangulateCell(const vtkm::Id2 &cdims) :
xdim(cdims[0]), ydim(cdims[1])
ScatterType GetScatter() const
{
return ScatterType(2);
}
// Each hexahedron cell produces five tetrahedron cells
VTKM_CONT_EXPORT
TriangulateCell()
{ }
// Each quad cell produces 2 triangle cells
template<typename ConnectivityInVec, typename ConnectivityOutVec>
VTKM_EXEC_EXPORT
void operator()(vtkm::Id &inputCellId,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity) const
void operator()(const ConnectivityInVec &connectivityIn,
ConnectivityOutVec &connectivityOut,
vtkm::IdComponent visitIndex) const
{
// Calculate the i,j indices for this input cell id
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
// Calculate the type of triangle generated because it alternates
vtkm::Id indexType = (x + y) % 2;
// Compute indices for the four vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1);
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
// Set the triangles for this cell based on vertex index and index type of cell
// 2 triangles per quad, 3 indices per triangle
vtkm::Id startIndex = inputCellId * 2 * 3;
if (indexType == 0)
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
}
else
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
}
connectivityOut[0] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][0]];
connectivityOut[1] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][1]];
connectivityOut[2] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][2]];
}
};
@ -112,105 +112,45 @@ public:
// Worklet to turn hexahedra into tetrahedra
// Vertices remain the same and each cell is processed with needing topology
//
class TetrahedralizeCell : public vtkm::worklet::WorkletMapField
class TetrahedralizeCell : public vtkm::worklet::WorkletMapPointToCell
{
public:
typedef void ControlSignature(FieldIn<IdType> inputCellId,
ExecObject connectivity);
typedef void ExecutionSignature(_1,_2);
typedef void ControlSignature(TopologyIn topology,
FieldOutCell<> connectivityOut);
typedef void ExecutionSignature(PointIndices, _2, ThreadIndices);
typedef _1 InputDomain;
vtkm::Id xdim, ydim, zdim;
const vtkm::Id cellsPerLayer, pointsPerLayer;
typedef vtkm::worklet::ScatterUniform ScatterType;
VTKM_CONT_EXPORT
TetrahedralizeCell(const vtkm::Id3 &cdims) :
xdim(cdims[0]), ydim(cdims[1]), zdim(cdims[2]),
cellsPerLayer(xdim * ydim),
pointsPerLayer((xdim+1) * (ydim+1))
ScatterType GetScatter() const
{
return ScatterType(5);
}
VTKM_CONT_EXPORT
TetrahedralizeCell()
{ }
// Each hexahedron cell produces five tetrahedron cells
template<typename ConnectivityInVec,
typename ConnectivityOutVec,
typename ThreadIndicesType>
VTKM_EXEC_EXPORT
void operator()(vtkm::Id &inputCellId,
vtkm::exec::ExecutionWholeArray<vtkm::Id> &connectivity) const
void operator()(const ConnectivityInVec &connectivityIn,
ConnectivityOutVec &connectivityOut,
const ThreadIndicesType threadIndices) const
{
// Calculate the i,j,k indices for this input cell id
const vtkm::Id x = inputCellId % xdim;
const vtkm::Id y = (inputCellId / xdim) % ydim;
const vtkm::Id z = inputCellId / cellsPerLayer;
vtkm::Id3 inputIndex = threadIndices.GetInputIndex3D();
// Calculate the type of tetrahedron generated because it alternates
vtkm::Id indexType = (x + y + z) % 2;
vtkm::Id indexType = (inputIndex[0] + inputIndex[1] + inputIndex[2]) % 2;
// Compute indices for the eight vertices of this cell
const vtkm::Id i0 = x + y*(xdim+1) + z * pointsPerLayer;
const vtkm::Id i1 = i0 + 1;
const vtkm::Id i2 = i0 + 1 + (xdim + 1); //xdim is cell dim
const vtkm::Id i3 = i0 + (xdim + 1); //xdim is cell dim
const vtkm::Id i4 = i0 + pointsPerLayer;
const vtkm::Id i5 = i1 + pointsPerLayer;
const vtkm::Id i6 = i2 + pointsPerLayer;
const vtkm::Id i7 = i3 + pointsPerLayer;
vtkm::IdComponent visitIndex = threadIndices.GetVisitIndex();
// Set the tetrahedra for this cell based on vertex index and index type of cell
// 5 tetrahedra per hexahedron, 4 indices per tetrahedron
vtkm::Id startIndex = inputCellId * 5 * 4;
if (indexType == 0)
{
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i4);
}
else
{
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i1);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i3);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i6);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i4);
connectivity.Set(startIndex++, i5);
connectivity.Set(startIndex++, i0);
connectivity.Set(startIndex++, i2);
connectivity.Set(startIndex++, i7);
connectivity.Set(startIndex++, i5);
}
connectivityOut[0] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][0]];
connectivityOut[1] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][1]];
connectivityOut[2] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][2]];
connectivityOut[3] = connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][3]];
}
};
@ -236,64 +176,28 @@ public:
vtkm::cont::CellSetSingleType<> & cellSet =
OutDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetSingleType<> >();
// Get dimensionality from the explicit cell set
vtkm::IdComponent dim = cellSet.GetDimensionality();
vtkm::Id outCellsPerInCell = 0;
vtkm::IdComponent verticesPerOutCell = 0;
vtkm::Id numberOfInCells = 0;
vtkm::Id2 cdims2;
vtkm::Id3 cdims3;
// From the uniform dimension get more information
if (dim == 2)
{
outCellsPerInCell = 2;
verticesPerOutCell = 3;
vtkm::cont::CellSetStructured<2> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<2> >();
cdims2 = inCellSet.GetSchedulingRange(vtkm::TopologyElementTagCell());
numberOfInCells = cdims2[0] * cdims2[1];
}
else if (dim == 3)
{
outCellsPerInCell = 5;
verticesPerOutCell = 4;
vtkm::cont::CellSetStructured<3> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<3> >();
cdims3 = inCellSet.GetSchedulingRange(vtkm::TopologyElementTagCell());
numberOfInCells = cdims3[0] * cdims3[1] * cdims3[2];
}
vtkm::Id numberOfOutCells = outCellsPerInCell * numberOfInCells;
vtkm::Id numberOfOutIndices = verticesPerOutCell * numberOfOutCells;
// Cell indices are just counting array
vtkm::cont::ArrayHandleCounting<vtkm::Id> cellIndicesArray(0, 1, numberOfInCells);
// Output dataset depends on dimension and size
vtkm::cont::ArrayHandle<vtkm::Id> connectivity;
connectivity.Allocate(numberOfOutIndices);
// Call the TetrahedralizeCell functor to compute tetrahedra or triangles
if (dim == 2)
if (cellSet.GetDimensionality() == 2)
{
TriangulateCell triangulateCell(cdims2);
vtkm::worklet::DispatcherMapField<TriangulateCell> triangulateCellDispatcher(triangulateCell);
triangulateCellDispatcher.Invoke(
cellIndicesArray,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
vtkm::cont::CellSetStructured<2> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<2> >();
vtkm::worklet::DispatcherMapTopology<TriangulateCell,DeviceAdapter> dispatcher;
dispatcher.Invoke(inCellSet,
vtkm::cont::make_ArrayHandleGroupVec<3>(connectivity));
}
else if (dim == 3)
else if (cellSet.GetDimensionality() == 3)
{
TetrahedralizeCell tetrahedralizeCell(cdims3);
vtkm::worklet::DispatcherMapField<TetrahedralizeCell> tetrahedralizeCellDispatcher(tetrahedralizeCell);
tetrahedralizeCellDispatcher.Invoke(
cellIndicesArray,
vtkm::exec::ExecutionWholeArray<vtkm::Id>(connectivity, numberOfOutIndices));
vtkm::cont::CellSetStructured<3> &inCellSet =
InDataSet.GetCellSet(0).template CastTo<vtkm::cont::CellSetStructured<3> >();
vtkm::worklet::DispatcherMapTopology<TetrahedralizeCell,DeviceAdapter> dispatcher;
dispatcher.Invoke(inCellSet,
vtkm::cont::make_ArrayHandleGroupVec<4>(connectivity));
}
else
{
throw vtkm::cont::ErrorControlBadValue(
"Unsupported dimensionality for TetrahedralizeUniformGrid.");
}
// Add cells to output cellset

2
vtkm/worklet/WorkletMapTopology.h
View File

@ -188,6 +188,8 @@ public:
template<typename TypeList = AllTypes >
struct FieldInOutCell : FieldInOut<TypeList> { };
struct PointIndices : FromIndices { };
};
}

3
vtkm/worklet/internal/CMakeLists.txt
View File

@ -19,10 +19,11 @@
##============================================================================
set(headers
ClipTables.h
DispatcherBase.h
DispatcherBaseDetailInvoke.h
TriangulateTables.h
WorkletBase.h
ClipTables.h
)
set_source_files_properties(DispatcherBaseDetailInvoke.h

41
vtkm/worklet/internal/DispatcherBase.h
View File

@ -441,19 +441,23 @@ protected:
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id numInstances,
DeviceAdapter tag) const
DeviceAdapter device) const
{
this->InvokeTransportParameters(invocation, numInstances, tag);
this->InvokeTransportParameters(
invocation,
this->Worklet.GetScatter().GetOutputRange(numInstances),
device);
}
template<typename Invocation, typename DeviceAdapter>
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id2 dimensions,
DeviceAdapter tag) const
DeviceAdapter device) const
{
vtkm::Id3 dim3d(dimensions[0], dimensions[1], 1);
this->InvokeTransportParameters(invocation, dim3d, tag);
this->BasicInvoke(invocation,
vtkm::Id3(dimensions[0], dimensions[1], 1),
device);
}
@ -461,9 +465,12 @@ protected:
VTKM_CONT_EXPORT
void BasicInvoke(const Invocation &invocation,
vtkm::Id3 dimensions,
DeviceAdapter tag) const
DeviceAdapter device) const
{
this->InvokeTransportParameters(invocation, dimensions, tag);
this->InvokeTransportParameters(
invocation,
this->Worklet.GetScatter().GetOutputRange(dimensions),
device);
}
WorkletType Worklet;
@ -477,7 +484,7 @@ private:
VTKM_CONT_EXPORT
void InvokeTransportParameters(const Invocation &invocation,
RangeType range,
DeviceAdapter tag) const
DeviceAdapter device) const
{
// The first step in invoking a worklet is to transport the arguments to
// the execution environment. The invocation object passed to this function
@ -499,11 +506,21 @@ private:
ExecObjectParameters execObjectParameters =
parameters.StaticTransformCont(TransportFunctorType(range));
// Get the arrays used for scattering input to output.
typename WorkletType::ScatterType::OutputToInputMapType outputToInputMap =
this->Worklet.GetScatter().GetOutputToInputMap(range);
typename WorkletType::ScatterType::VisitArrayType visitArray =
this->Worklet.GetScatter().GetVisitArray(range);
// Replace the parameters in the invocation with the execution object and
// pass to next step of Invoke.
this->InvokeSchedule(invocation.ChangeParameters(execObjectParameters),
range,
tag);
// pass to next step of Invoke. Also add the scatter information.
this->InvokeSchedule(
invocation
.ChangeParameters(execObjectParameters)
.ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device))
.ChangeVisitArray(visitArray.PrepareForInput(device)),
range,
device);
}
template<typename Invocation, typename RangeType, typename DeviceAdapter>

312
vtkm/worklet/internal/TriangulateTables.h Normal file
View File

@ -0,0 +1,312 @@
//============================================================================
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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_worklet_internal_TriangulateTables_h
#define vtk_m_worklet_internal_TriangulateTables_h
#include <vtkm/CellShape.h>
#include <vtkm/Types.h>
#include <vtkm/exec/ExecutionObjectBase.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageBasic.h>
namespace vtkm {
namespace worklet {
namespace internal {
typedef vtkm::cont::ArrayHandle<vtkm::IdComponent, vtkm::cont::StorageTagBasic>
TriangulateArrayHandle;
static vtkm::IdComponent TriangleCountData[vtkm::NUMBER_OF_CELL_SHAPES] = {
0, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
0, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
0, // 3 = vtkm::CELL_SHAPE_LINE
0, // 4 = vtkm::CELL_SHAPE_POLY_LINE
1, // 5 = vtkm::CELL_SHAPE_TRIANGLE
0, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
0, // 8 = vtkm::CELL_SHAPE_PIXEL
2, // 9 = vtkm::CELL_SHAPE_QUAD
0, // 10 = vtkm::CELL_SHAPE_TETRA
0, // 11 = vtkm::CELL_SHAPE_VOXEL
0, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
0, // 13 = vtkm::CELL_SHAPE_WEDGE
0 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TriangleOffsetData[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
-1, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
-1, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
0, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
1, // 9 = vtkm::CELL_SHAPE_QUAD
-1, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
-1, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
-1, // 13 = vtkm::CELL_SHAPE_WEDGE
-1 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TriangleIndexData[] = {
// vtkm::CELL_SHAPE_TRIANGLE
0, 1, 2,
// vtkm::CELL_SHAPE_QUAD
0, 1, 2,
0, 2, 3
};
template<typename Device>
class TriangulateTablesExecutionObject : public vtkm::exec::ExecutionObjectBase
{
public:
typedef typename TriangulateArrayHandle::ExecutionTypes<Device>::PortalConst
PortalType;
VTKM_EXEC_CONT_EXPORT
TriangulateTablesExecutionObject()
{ }
VTKM_CONT_EXPORT
TriangulateTablesExecutionObject(const TriangulateArrayHandle &counts,
const TriangulateArrayHandle &offsets,
const TriangulateArrayHandle &indices)
: Counts(counts.PrepareForInput(Device())),
Offsets(offsets.PrepareForInput(Device())),
Indices(indices.PrepareForInput(Device()))
{ }
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::IdComponent GetCount(CellShape shape, vtkm::IdComponent numPoints) const
{
if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
return numPoints-2;
}
else
{
return this->Counts.Get(shape.Id);
}
}
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::Vec<vtkm::IdComponent, 3>
GetIndices(CellShape shape, vtkm::IdComponent triangleIndex) const
{
vtkm::Vec<vtkm::IdComponent, 3> triIndices;
if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{
triIndices[0] = 0;
triIndices[1] = triangleIndex + 1;
triIndices[2] = triangleIndex + 2;
}
else
{
vtkm::IdComponent offset =
3*(this->Offsets.Get(shape.Id) + triangleIndex);
triIndices[0] = this->Indices.Get(offset + 0);
triIndices[1] = this->Indices.Get(offset + 1);
triIndices[2] = this->Indices.Get(offset + 2);
}
return triIndices;
}
private:
PortalType Counts;
PortalType Offsets;
PortalType Indices;
};
class TriangulateTables
{
public:
VTKM_CONT_EXPORT
TriangulateTables()
: Counts(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleCountData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Offsets(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleOffsetData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Indices(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TriangleIndexData,
sizeof(vtkm::worklet::internal::TriangleIndexData)/sizeof(vtkm::IdComponent)))
{ }
template<typename Device>
vtkm::worklet::internal::TriangulateTablesExecutionObject<Device>
PrepareForInput(Device) const
{
return vtkm::worklet::internal::TriangulateTablesExecutionObject<Device>(
this->Counts, this->Offsets, this->Indices);
}
private:
TriangulateArrayHandle Counts;
TriangulateArrayHandle Offsets;
TriangulateArrayHandle Indices;
};
static vtkm::IdComponent TetrahedronCountData[vtkm::NUMBER_OF_CELL_SHAPES] = {
0, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
0, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
0, // 3 = vtkm::CELL_SHAPE_LINE
0, // 4 = vtkm::CELL_SHAPE_POLY_LINE
0, // 5 = vtkm::CELL_SHAPE_TRIANGLE
0, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
0, // 7 = vtkm::CELL_SHAPE_POLYGON
0, // 8 = vtkm::CELL_SHAPE_PIXEL
0, // 9 = vtkm::CELL_SHAPE_QUAD
1, // 10 = vtkm::CELL_SHAPE_TETRA
0, // 11 = vtkm::CELL_SHAPE_VOXEL
5, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
3, // 13 = vtkm::CELL_SHAPE_WEDGE
2 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TetrahedronOffsetData[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
-1, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
-1, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
-1, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
-1, // 9 = vtkm::CELL_SHAPE_QUAD
0, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
1, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
6, // 13 = vtkm::CELL_SHAPE_WEDGE
9 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
static vtkm::IdComponent TetrahedronIndexData[] = {
// vtkm::CELL_SHAPE_TETRA
0, 1, 2, 3,
// vtkm::CELL_SHAPE_HEXAHEDRON
0, 1, 3, 4,
1, 4, 5, 6,
1, 4, 6, 3,
1, 3, 6, 2,
3, 6, 7, 4,
// vtkm::CELL_SHAPE_WEDGE
0, 1, 2, 4,
3, 4, 5, 2,
0, 2, 3, 4,
// vtkm::CELL_SHAPE_PYRAMID
0, 1, 2, 4,
0, 2, 3, 4
};
template<typename Device>
class TetrahedralizeTablesExecutionObject
: public vtkm::exec::ExecutionObjectBase
{
public:
typedef typename TriangulateArrayHandle::ExecutionTypes<Device>::PortalConst
PortalType;
VTKM_EXEC_CONT_EXPORT
TetrahedralizeTablesExecutionObject()
{ }
VTKM_CONT_EXPORT
TetrahedralizeTablesExecutionObject(const TriangulateArrayHandle &counts,
const TriangulateArrayHandle &offsets,
const TriangulateArrayHandle &indices)
: Counts(counts.PrepareForInput(Device())),
Offsets(offsets.PrepareForInput(Device())),
Indices(indices.PrepareForInput(Device()))
{ }
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::IdComponent GetCount(CellShape shape) const
{
return this->Counts.Get(shape.Id);
}
template<typename CellShape>
VTKM_EXEC_EXPORT
vtkm::Vec<vtkm::IdComponent, 4>
GetIndices(CellShape shape, vtkm::IdComponent tetrahedronIndex) const
{
vtkm::Vec<vtkm::IdComponent, 4> tetIndices;
vtkm::IdComponent offset =
4*(this->Offsets.Get(shape.Id) + tetrahedronIndex);
tetIndices[0] = this->Indices.Get(offset + 0);
tetIndices[1] = this->Indices.Get(offset + 1);
tetIndices[2] = this->Indices.Get(offset + 2);
tetIndices[3] = this->Indices.Get(offset + 3);
return tetIndices;
}
private:
PortalType Counts;
PortalType Offsets;
PortalType Indices;
};
class TetrahedralizeTables
{
public:
VTKM_CONT_EXPORT
TetrahedralizeTables()
: Counts(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronCountData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Offsets(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronOffsetData,
vtkm::NUMBER_OF_CELL_SHAPES)),
Indices(vtkm::cont::make_ArrayHandle(
vtkm::worklet::internal::TetrahedronIndexData,
sizeof(vtkm::worklet::internal::TetrahedronIndexData)/sizeof(vtkm::IdComponent)))
{ }
template<typename Device>
vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<Device>
PrepareForInput(Device) const
{
return vtkm::worklet::internal::TetrahedralizeTablesExecutionObject<Device>(
this->Counts, this->Offsets, this->Indices);
}
private:
TriangulateArrayHandle Counts;
TriangulateArrayHandle Offsets;
TriangulateArrayHandle Indices;
};
}
}
}
#endif //vtk_m_worklet_internal_TriangulateTables_h

24
vtkm/worklet/internal/WorkletBase.h
View File

@ -25,13 +25,17 @@
#include <vtkm/exec/FunctorBase.h>
#include <vtkm/exec/arg/BasicArg.h>
#include <vtkm/exec/arg/FetchTagExecObject.h>
#include <vtkm/exec/arg/ThreadIndices.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/exec/arg/VisitIndex.h>
#include <vtkm/exec/arg/WorkIndex.h>
#include <vtkm/cont/arg/ControlSignatureTagBase.h>
#include <vtkm/cont/arg/TransportTagExecObject.h>
#include <vtkm/cont/arg/TypeCheckTagExecObject.h>
#include <vtkm/worklet/ScatterIdentity.h>
namespace vtkm {
namespace worklet {
namespace internal {
@ -58,8 +62,17 @@ public:
struct _9 : Arg<9> { };
/// \c ExecutionSignature tag for getting the work index.
///
typedef vtkm::exec::arg::WorkIndex WorkIndex;
/// \c ExecutionSignature tag for getting the thread indices.
///
typedef vtkm::exec::arg::ThreadIndices ThreadIndices;
/// \c ExecutionSignature tag for getting the visit index.
///
typedef vtkm::exec::arg::VisitIndex VisitIndex;
/// \c ControlSignature tag for execution object inputs.
struct ExecObject : vtkm::cont::arg::ControlSignatureTagBase {
typedef vtkm::cont::arg::TypeCheckTagExecObject TypeCheckTag;
@ -71,6 +84,17 @@ public:
/// override this by redefining this type.
typedef _1 InputDomain;
/// All worklets must define their scatter operation. The scatter defines
/// what output each input contributes to. The default scatter is the
/// identity scatter (1-to-1 input to output).
typedef vtkm::worklet::ScatterIdentity ScatterType;
/// In addition to defining the scatter type, the worklet must produce the
/// scatter. The default ScatterIdentity has no state, so just return an
/// instance.
VTKM_CONT_EXPORT
ScatterType GetScatter() const { return ScatterType(); }
/// \brief A type list containing the type vtkm::Id.
///
/// This is a convenience type to use as template arguments to \c

4
vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx
View File

@ -119,7 +119,7 @@ struct Fetch<TestFetchTagInput,
VTKM_EXEC_EXPORT
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic indices,
const TestExecObject &execObject) const {
return execObject.Array[indices.GetIndex()];
return execObject.Array[indices.GetInputIndex()];
}
VTKM_EXEC_EXPORT
@ -149,7 +149,7 @@ struct Fetch<TestFetchTagOutput,
void Store(const vtkm::exec::arg::ThreadIndicesBasic &indices,
const TestExecObject &execObject,
ValueType value) const {
execObject.Array[indices.GetIndex()] = value;
execObject.Array[indices.GetOutputIndex()] = value;
}
};

1
vtkm/worklet/testing/CMakeLists.txt
View File

@ -26,6 +26,7 @@ set(unit_tests
UnitTestFieldStatistics.cxx
UnitTestIsosurfaceUniformGrid.cxx
UnitTestPointElevation.cxx
UnitTestScatterCounting.cxx
UnitTestSplatKernels.cxx
UnitTestTetrahedralizeExplicitGrid.cxx
UnitTestTetrahedralizeUniformGrid.cxx

35
vtkm/worklet/testing/UnitTestIsosurfaceUniformGrid.cxx
View File

@ -38,12 +38,12 @@ public:
typedef _1 InputDomain;
const vtkm::Id xdim, ydim, zdim;
const float xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::FloatDefault xmin, ymin, zmin, xmax, ymax, zmax;
const vtkm::Id cellsPerLayer;
VTKM_CONT_EXPORT
TangleField(const vtkm::Id3 dims, const float mins[3], const float maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { };
TangleField(const vtkm::Id3 dims, const vtkm::FloatDefault mins[3], const vtkm::FloatDefault maxs[3]) : xdim(dims[0]), ydim(dims[1]), zdim(dims[2]),
xmin(mins[0]), ymin(mins[1]), zmin(mins[2]), xmax(maxs[0]), ymax(maxs[1]), zmax(maxs[2]), cellsPerLayer((xdim) * (ydim)) { }
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &vertexId, vtkm::Float32 &v) const
@ -52,9 +52,9 @@ public:
const vtkm::Id y = (vertexId / (xdim)) % (ydim);
const vtkm::Id z = vertexId / cellsPerLayer;
const float fx = static_cast<float>(x) / static_cast<float>(xdim-1);
const float fy = static_cast<float>(y) / static_cast<float>(xdim-1);
const float fz = static_cast<float>(z) / static_cast<float>(xdim-1);
const vtkm::FloatDefault fx = static_cast<vtkm::FloatDefault>(x) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fy = static_cast<vtkm::FloatDefault>(y) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::FloatDefault fz = static_cast<vtkm::FloatDefault>(z) / static_cast<vtkm::FloatDefault>(xdim-1);
const vtkm::Float32 xx = 3.0f*(xmin+(xmax-xmin)*(fx));
const vtkm::Float32 yy = 3.0f*(ymin+(ymax-ymin)*(fy));
@ -71,15 +71,22 @@ vtkm::cont::DataSet MakeIsosurfaceTestDataSet(vtkm::Id3 dims)
const vtkm::Id3 vdims(dims[0] + 1, dims[1] + 1, dims[2] + 1);
float mins[3] = {-1.0f, -1.0f, -1.0f};
float maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::FloatDefault mins[3] = {-1.0f, -1.0f, -1.0f};
vtkm::FloatDefault maxs[3] = {1.0f, 1.0f, 1.0f};
vtkm::cont::ArrayHandle<vtkm::Float32> fieldArray;
vtkm::cont::ArrayHandleIndex vertexCountImplicitArray(vdims[0]*vdims[1]*vdims[2]);
vtkm::worklet::DispatcherMapField<TangleField> tangleFieldDispatcher(TangleField(vdims, mins, maxs));
tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray);
vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims);
vtkm::Vec<vtkm::FloatDefault,3> origin(0.0f, 0.0f, 0.0f);
vtkm::Vec<vtkm::FloatDefault,3> spacing(
1.0f/static_cast<vtkm::FloatDefault>(dims[0]),
1.0f/static_cast<vtkm::FloatDefault>(dims[2]),
1.0f/static_cast<vtkm::FloatDefault>(dims[1]));
vtkm::cont::ArrayHandleUniformPointCoordinates
coordinates(vdims, origin, spacing);
dataSet.AddCoordinateSystem(
vtkm::cont::CoordinateSystem("coordinates", 1, coordinates));
@ -116,6 +123,16 @@ void TestIsosurfaceUniformGrid()
normalsArray,
scalarsArray);
std::cout << "vertices: ";
vtkm::cont::printSummary_ArrayHandle(verticesArray, std::cout);
std::cout << std::endl;
std::cout << "normals: ";
vtkm::cont::printSummary_ArrayHandle(normalsArray, std::cout);
std::cout << std::endl;
std::cout << "scalars: ";
vtkm::cont::printSummary_ArrayHandle(scalarsArray, std::cout);
std::cout << std::endl;
VTKM_TEST_ASSERT(test_equal(verticesArray.GetNumberOfValues(), 480),
"Wrong result for Isosurface filter");
}

258
vtkm/worklet/testing/UnitTestScatterCounting.cxx Normal file
View File

@ -0,0 +1,258 @@
//=============================================================================
//
// 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 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// 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/worklet/ScatterCounting.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/testing/Testing.h>
#include <vector>
namespace {
struct TestScatterArrays
{
vtkm::cont::ArrayHandle<vtkm::IdComponent> CountArray;
vtkm::cont::ArrayHandle<vtkm::Id> OutputToInputMap;
vtkm::cont::ArrayHandle<vtkm::IdComponent> VisitArray;
};
TestScatterArrays MakeScatterArraysShort()
{
const vtkm::Id countArraySize = 18;
const vtkm::IdComponent countArray[countArraySize] = {
1, 2, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0
};
const vtkm::Id outputSize = 6;
const vtkm::Id outputToInputMap[outputSize] = {
0, 1, 1, 4, 6, 14
};
const vtkm::IdComponent visitArray[outputSize] = {
0, 0, 1, 0, 0, 0
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(outputToInputMap, outputSize),
arrays.OutputToInputMap);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(visitArray, outputSize),
arrays.VisitArray);
return arrays;
}
TestScatterArrays MakeScatterArraysLong()
{
const vtkm::Id countArraySize = 6;
const vtkm::IdComponent countArray[countArraySize] = {
0, 1, 2, 3, 4, 5
};
const vtkm::Id outputSize = 15;
const vtkm::Id outputToInputMap[outputSize] = {
1, 2, 2, 3, 3, 3, 4, 4, 4, 4, 5, 5, 5, 5, 5
};
const vtkm::IdComponent visitArray[outputSize] = {
0, 0, 1, 0, 1, 2, 0, 1, 2, 3, 0, 1, 2, 3, 4
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(outputToInputMap, outputSize),
arrays.OutputToInputMap);
Algorithm::Copy(vtkm::cont::make_ArrayHandle(visitArray, outputSize),
arrays.VisitArray);
return arrays;
}
TestScatterArrays MakeScatterArraysZero()
{
const vtkm::Id countArraySize = 6;
const vtkm::IdComponent countArray[countArraySize] = {
0, 0, 0, 0, 0, 0
};
TestScatterArrays arrays;
typedef vtkm::cont::DeviceAdapterAlgorithm<VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
Algorithm;
// Need to copy arrays so that the data does not go out of scope.
Algorithm::Copy(vtkm::cont::make_ArrayHandle(countArray, countArraySize),
arrays.CountArray);
arrays.OutputToInputMap.Allocate(0);
arrays.VisitArray.Allocate(0);
return arrays;
}
struct TestScatterCountingWorklet : public vtkm::worklet::WorkletMapField
{
typedef void ControlSignature(FieldIn<> inputIndices,
FieldOut<> copyIndices,
FieldOut<> recordVisit,
FieldOut<> recordWorkId);
typedef void ExecutionSignature(_1, _2 ,_3, _4, VisitIndex, WorkIndex);
typedef vtkm::worklet::ScatterCounting ScatterType;
VTKM_CONT_EXPORT
ScatterType GetScatter() const { return this->Scatter; }
template<typename CountArrayType>
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const CountArrayType &countArray)
: Scatter(countArray, VTKM_DEFAULT_DEVICE_ADAPTER_TAG()) { }
template<typename CountArrayType, typename Device>
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const CountArrayType &countArray, Device)
: Scatter(countArray, Device()) { }
VTKM_CONT_EXPORT
TestScatterCountingWorklet(const vtkm::worklet::ScatterCounting &scatter)
: Scatter(scatter) { }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id inputIndex,
vtkm::Id &indexCopy,
vtkm::IdComponent &writeVisit,
vtkm::Float32 &captureWorkId,
vtkm::IdComponent visitIndex,
vtkm::Id workId) const
{
indexCopy = inputIndex;
writeVisit = visitIndex;
captureWorkId = TestValue(workId, vtkm::Float32());
}
private:
ScatterType Scatter;
};
template<typename T>
void CompareArrays(vtkm::cont::ArrayHandle<T> array1,
vtkm::cont::ArrayHandle<T> array2)
{
typedef typename vtkm::cont::ArrayHandle<T>::PortalConstControl PortalType;
PortalType portal1 = array1.GetPortalConstControl();
PortalType portal2 = array2.GetPortalConstControl();
VTKM_TEST_ASSERT(portal1.GetNumberOfValues() == portal2.GetNumberOfValues(),
"Arrays are not the same length.");
for (vtkm::Id index = 0; index < portal1.GetNumberOfValues(); index++)
{
T value1 = portal1.Get(index);
T value2 = portal2.Get(index);
VTKM_TEST_ASSERT(value1 == value2, "Array values not equal.");
}
}
// This unit test makes sure the ScatterCounting generates the correct map
// and visit arrays.
void TestScatterArrayGeneration(const TestScatterArrays &arrays)
{
std::cout << " Testing array generation" << std::endl;
vtkm::worklet::ScatterCounting scatter(arrays.CountArray,
VTKM_DEFAULT_DEVICE_ADAPTER_TAG());
vtkm::Id inputSize = arrays.CountArray.GetNumberOfValues();
std::cout << " Checking output to input map." << std::endl;
CompareArrays(arrays.OutputToInputMap,
scatter.GetOutputToInputMap(inputSize));
std::cout << " Checking visit array." << std::endl;
CompareArrays(arrays.VisitArray,
scatter.GetVisitArray(inputSize));
}
// This is more of an integration test that makes sure the scatter works with a
// worklet invocation.
void TestScatterWorklet(const TestScatterArrays &arrays)
{
std::cout << " Testing scatter counting in a worklet." << std::endl;
vtkm::worklet::ScatterCounting scatter(arrays.CountArray,
VTKM_DEFAULT_DEVICE_ADAPTER_TAG());
TestScatterCountingWorklet worklet(scatter);
vtkm::worklet::DispatcherMapField<TestScatterCountingWorklet> dispatcher(
worklet);
vtkm::Id inputSize = arrays.CountArray.GetNumberOfValues();
vtkm::cont::ArrayHandleIndex inputIndices(inputSize);
vtkm::cont::ArrayHandle<vtkm::Id> outputToInputMapCopy;
vtkm::cont::ArrayHandle<vtkm::IdComponent> visitCopy;
vtkm::cont::ArrayHandle<vtkm::Float32> captureWorkId;
std::cout << " Invoke worklet" << std::endl;
dispatcher.Invoke(
inputIndices, outputToInputMapCopy, visitCopy, captureWorkId);
std::cout << " Check output to input map." << std::endl;
CompareArrays(outputToInputMapCopy, arrays.OutputToInputMap);
std::cout << " Check visit." << std::endl;
CompareArrays(visitCopy, arrays.VisitArray);
std::cout << " Check work id." << std::endl;
CheckPortal(captureWorkId.GetPortalConstControl());
}
void TestScatterCountingWithArrays(const TestScatterArrays &arrays)
{
TestScatterArrayGeneration(arrays);
TestScatterWorklet(arrays);
}
void TestScatterCounting()
{
std::cout << "Testing arrays with output smaller than input." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysShort());
std::cout << "Testing arrays with output larger than input." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysLong());
std::cout << "Testing arrays with zero output." << std::endl;
TestScatterCountingWithArrays(MakeScatterArraysZero());
}
} // anonymous namespace
int UnitTestScatterCounting(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestScatterCounting);
}