diff --git a/CMake/UseVTKmBase.cmake b/CMake/UseVTKmBase.cmake new file mode 100644 index 000000000..d1241f0b3 --- /dev/null +++ b/CMake/UseVTKmBase.cmake @@ -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) diff --git a/CMake/UseVTKmCUDA.cmake b/CMake/UseVTKmCUDA.cmake new file mode 100644 index 000000000..219886150 --- /dev/null +++ b/CMake/UseVTKmCUDA.cmake @@ -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) diff --git a/CMake/UseVTKmCuda.cmake b/CMake/UseVTKmCuda.cmake deleted file mode 100644 index 04f046196..000000000 --- a/CMake/UseVTKmCuda.cmake +++ /dev/null @@ -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) diff --git a/CMake/UseVTKmSerial.cmake b/CMake/UseVTKmSerial.cmake index 76479bedc..e96bbbfd0 100644 --- a/CMake/UseVTKmSerial.cmake +++ b/CMake/UseVTKmSerial.cmake @@ -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 () diff --git a/CMake/UseVTKmTBB.cmake b/CMake/UseVTKmTBB.cmake index e008ea982..df48a35e2 100644 --- a/CMake/UseVTKmTBB.cmake +++ b/CMake/UseVTKmTBB.cmake @@ -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() diff --git a/CMake/VTKmConfig.cmake.in b/CMake/VTKmConfig.cmake.in index 627df6113..1e871e5ea 100644 --- a/CMake/VTKmConfig.cmake.in +++ b/CMake/VTKmConfig.cmake.in @@ -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() diff --git a/CMake/VTKmMacros.cmake b/CMake/VTKmMacros.cmake index 2c5d82b61..99668f6d8 100644 --- a/CMake/VTKmMacros.cmake +++ b/CMake/VTKmMacros.cmake @@ -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) diff --git a/CMakeLists.txt b/CMakeLists.txt index 2dfb8313a..4cead0a42 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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}) #----------------------------------------------------------------------------- diff --git a/examples/clipping/CMakeLists.txt b/examples/clipping/CMakeLists.txt index 4c87986dd..f8c8fb95f 100644 --- a/examples/clipping/CMakeLists.txt +++ b/examples/clipping/CMakeLists.txt @@ -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 diff --git a/examples/clipping/Clipping.cxx b/examples/clipping/Clipping.cxx index 40055c5c3..d3e4a89c1 100644 --- a/examples/clipping/Clipping.cxx +++ b/examples/clipping/Clipping.cxx @@ -21,6 +21,8 @@ #include #include #include +#include +#include #include @@ -36,223 +38,14 @@ VTKM_THIRDPARTY_POST_INCLUDE typedef vtkm::Vec FloatVec3; -template -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(&val); - std::reverse(bytes, bytes + sizeof(T)); - buffer[i] = val; - } -} - -template -inline vtkm::cont::ArrayHandle LoadBinaryPointDataImpl( - std::ifstream &fstream, vtkm::Id numPoints, T, DeviceAdapter) -{ - vtkm::cont::ArrayHandle result; - - std::vector buffer(static_cast(numPoints)); - fstream.read(reinterpret_cast(&buffer[0]), - numPoints * static_cast(sizeof(T))); - flipEndianness(&buffer[0], numPoints); - - vtkm::cont::DeviceAdapterAlgorithm::Copy( - vtkm::cont::make_ArrayHandleCast(vtkm::cont::make_ArrayHandle(buffer), vtkm::Float32()), - result); - - return result; -} - -template -inline vtkm::cont::ArrayHandle 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(); -} - -template -vtkm::cont::DataSet LoadVtkLegacyStructuredPoints(const char *fname, DeviceAdapter) -{ - vtkm::Id3 dim; - FloatVec3 spacing, origin; - vtkm::cont::ArrayHandle 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 coordinates = - coords.GetData().CastToArrayHandle(FloatVec3(), VTKM_DEFAULT_STORAGE_TAG()); - std::vector buffer(static_cast(numPoints)); - - std::copy(vtkm::cont::ArrayPortalToIteratorBegin(coordinates.GetPortalConstControl()), - vtkm::cont::ArrayPortalToIteratorEnd(coordinates.GetPortalConstControl()), - buffer.begin()); - - flipEndianness(reinterpret_cast(&buffer[0]), numPoints * 3); - vtkfile.write(reinterpret_cast(&buffer[0]), - numPoints * static_cast(sizeof(FloatVec3))); - } - vtkfile << std::endl; - - vtkm::cont::CellSetExplicit<> cse = - ds.GetCellSet().CastTo >(); - vtkm::Id numCells = cse.GetNumberOfCells(); - { - std::vector idxBuffer, shapeBuffer; - - idxBuffer.reserve(static_cast(numCells * 4)); - shapeBuffer.reserve(static_cast(numCells)); - - vtkm::Vec pointIndices; - for (vtkm::Id i = 0; i < numCells; ++i) - { - vtkm::IdComponent numCellPoints = cse.GetNumberOfPointsInCell(i); - idxBuffer.push_back(static_cast(numCellPoints)); - cse.GetIndices(i, pointIndices); - for (vtkm::IdComponent j = 0; j < numCellPoints; ++j) - { - idxBuffer.push_back(static_cast(pointIndices[j])); - } - shapeBuffer.push_back(static_cast(cse.GetCellShape(i))); - } - vtkm::Id numIndices = static_cast(idxBuffer.size()); - - vtkfile << "CELLS " << numCells << " " << numIndices << std::endl; - flipEndianness(&idxBuffer[0], numIndices); - vtkfile.write(reinterpret_cast(&idxBuffer[0]), - numIndices * static_cast(sizeof(idxBuffer[0]))); - vtkfile << std::endl; - vtkfile << "CELL_TYPES " << numCells << std::endl; - flipEndianness(&shapeBuffer[0], numCells); - vtkfile.write(reinterpret_cast(&shapeBuffer[0]), - numCells * static_cast(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 scalars = - field.GetData().CastToArrayHandle(vtkm::Float32(), VTKM_DEFAULT_STORAGE_TAG()); - std::vector buffer(static_cast(numPoints)); - - std::copy(vtkm::cont::ArrayPortalToIteratorBegin(scalars.GetPortalConstControl()), - vtkm::cont::ArrayPortalToIteratorEnd(scalars.GetPortalConstControl()), - buffer.begin()); - - flipEndianness(&buffer[0], numPoints); - vtkfile.write(reinterpret_cast(&buffer[0]), - numPoints * static_cast(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] - << " " + << " [fieldName] " << std::endl; return 1; } @@ -262,24 +55,46 @@ int main(int argc, char *argv[]) << vtkm::cont::internal::DeviceAdapterTraits::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(argv[2]); + vtkm::cont::Field scalarField = (argc == 5) ? + input.GetField(argv[2]) : + input.GetField(0); + + vtkm::Float32 clipValue = boost::lexical_cast(argv[argc - 2]); vtkm::worklet::Clip clip; vtkm::cont::Timer total; vtkm::cont::Timer 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; } diff --git a/examples/hello_world/CMakeLists.txt b/examples/hello_world/CMakeLists.txt index 83425051d..44aa8d3b1 100644 --- a/examples/hello_world/CMakeLists.txt +++ b/examples/hello_world/CMakeLists.txt @@ -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() diff --git a/examples/hello_world/HelloWorld.cxx b/examples/hello_world/HelloWorld.cxx index 753791c7b..c23e676d7 100755 --- a/examples/hello_world/HelloWorld.cxx +++ b/examples/hello_world/HelloWorld.cxx @@ -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 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) diff --git a/examples/isosurface/CMakeLists.txt b/examples/isosurface/CMakeLists.txt index 027d77ffb..7d8c69a60 100644 --- a/examples/isosurface/CMakeLists.txt +++ b/examples/isosurface/CMakeLists.txt @@ -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() diff --git a/examples/isosurface/IsosurfaceUniformGrid.cxx b/examples/isosurface/IsosurfaceUniformGrid.cxx index 1aff52916..e8b0f2a03 100644 --- a/examples/isosurface/IsosurfaceUniformGrid.cxx +++ b/examples/isosurface/IsosurfaceUniformGrid.cxx @@ -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(x) / static_cast(xdim-1); - const float fy = static_cast(y) / static_cast(xdim-1); - const float fz = static_cast(z) / static_cast(xdim-1); + const vtkm::FloatDefault fx = static_cast(x) / static_cast(xdim-1); + const vtkm::FloatDefault fy = static_cast(y) / static_cast(xdim-1); + const vtkm::FloatDefault fz = static_cast(z) / static_cast(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 fieldArray; vtkm::cont::ArrayHandleCounting vertexCountImplicitArray(0, 1, vdims[0]*vdims[1]*vdims[2]); vtkm::worklet::DispatcherMapField tangleFieldDispatcher(TangleField(vdims, mins, maxs)); tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray); - vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims); + vtkm::Vec origin(0.0f, 0.0f, 0.0f); + vtkm::Vec spacing( + 1.0f/static_cast(dims[0]), + 1.0f/static_cast(dims[2]), + 1.0f/static_cast(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(vtkm::Pi()/180.0); + vtkm::Float32 pideg = static_cast(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); diff --git a/examples/multi_backend/CMakeLists.txt b/examples/multi_backend/CMakeLists.txt index a8829ba7f..b3027c1a5 100644 --- a/examples/multi_backend/CMakeLists.txt +++ b/examples/multi_backend/CMakeLists.txt @@ -27,5 +27,5 @@ else() endif() if(VTKm_TBB_FOUND) - target_link_libraries(MultiBackend ${TBB_LIBRARIES}) + target_link_libraries(MultiBackend ${VTKm_LIBRARIES}) endif() diff --git a/examples/tetrahedra/CMakeLists.txt b/examples/tetrahedra/CMakeLists.txt index 9b4ae0ecb..399427201 100644 --- a/examples/tetrahedra/CMakeLists.txt +++ b/examples/tetrahedra/CMakeLists.txt @@ -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() diff --git a/vtkm/Types.h b/vtkm/Types.h index 2d7f20c74..fbc795850 100644 --- a/vtkm/Types.h +++ b/vtkm/Types.h @@ -175,6 +175,12 @@ namespace internal { //----------------------------------------------------------------------------- +/// Placeholder class for when a type is not applicable. +/// +struct NullType { }; + +//----------------------------------------------------------------------------- + template struct VecEquals { diff --git a/vtkm/cont/ArrayHandleCast.h b/vtkm/cont/ArrayHandleCast.h index a57e6cc80..8e1567a73 100644 --- a/vtkm/cont/ArrayHandleCast.h +++ b/vtkm/cont/ArrayHandleCast.h @@ -53,18 +53,19 @@ class ArrayHandleCast : public vtkm::cont::ArrayHandleTransform< T, ArrayHandleType, - internal::Cast > + internal::Cast, + internal::Cast > { public: VTKM_ARRAY_HANDLE_SUBCLASS( - ArrayHandleCast, - (ArrayHandleCast), - (vtkm::cont::ArrayHandleTransform< - T, - ArrayHandleType, - internal::Cast >)); + ArrayHandleCast, + (ArrayHandleCast), + (vtkm::cont::ArrayHandleTransform< + T, + ArrayHandleType, + internal::Cast, + internal::Cast >)); - VTKM_CONT_EXPORT ArrayHandleCast(const ArrayHandleType &handle) : Superclass(handle) { } diff --git a/vtkm/cont/ArrayHandleGroupVec.h b/vtkm/cont/ArrayHandleGroupVec.h index 4431322c9..07e0a6fad 100644 --- a/vtkm/cont/ArrayHandleGroupVec.h +++ b/vtkm/cont/ArrayHandleGroupVec.h @@ -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_CONT_EXPORT +vtkm::cont::ArrayHandleGroupVec +make_ArrayHandleGroupVec(const ArrayHandleType &array) +{ + return vtkm::cont::ArrayHandleGroupVec(array); +} + } } // namespace vtkm::cont diff --git a/vtkm/cont/ArrayHandleTransform.h b/vtkm/cont/ArrayHandleTransform.h index f329764ec..6d6c71504 100644 --- a/vtkm/cont/ArrayHandleTransform.h +++ b/vtkm/cont/ArrayHandleTransform.h @@ -27,20 +27,38 @@ #include #include +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 +class ArrayPortalExecTransform; + template -class ArrayPortalExecTransform +class ArrayPortalExecTransform { 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 - VTKM_CONT_EXPORT + VTKM_EXEC_CONT_EXPORT ArrayPortalExecTransform(const ArrayPortalExecTransform &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 +class ArrayPortalExecTransform : public ArrayPortalExecTransform +{ +public: + typedef ArrayPortalExecTransform 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 + VTKM_EXEC_CONT_EXPORT + ArrayPortalExecTransform(const ArrayPortalExecTransform &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 +class ArrayPortalContTransform; + template -class ArrayPortalContTransform +class ArrayPortalContTransform { 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 -struct StorageTagTransform { }; +template +class ArrayPortalContTransform : public ArrayPortalContTransform +{ +public: + typedef ArrayPortalContTransform 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 + VTKM_CONT_EXPORT + ArrayPortalContTransform(const ArrayPortalContTransform &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 +struct StorageTagTransform {}; template -class Storage > +class Storage > { public: typedef T ValueType; @@ -221,6 +318,92 @@ private: bool Valid; }; +template +class Storage > +{ +public: + typedef T ValueType; + + typedef ArrayPortalContTransform + PortalType; + typedef ArrayPortalContTransform + 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 +class ArrayTransfer< + T, StorageTagTransform, + Device> +{ + typedef StorageTagTransform StorageTag; + typedef vtkm::cont::internal::Storage 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::Portal, + FunctorType, InverseFunctorType> PortalExecution; + typedef vtkm::exec::internal::ArrayPortalExecTransform< + ValueType, + typename ArrayHandleType::template ExecutionTypes::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 +class ArrayHandleTransform; + template -class ArrayHandleTransform +class ArrayHandleTransform : public vtkm::cont::ArrayHandle< T, internal::StorageTagTransform > { @@ -327,7 +593,7 @@ public: T, internal::StorageTagTransform >)); private: - typedef vtkm::cont::internal::Storage StorageType; + typedef vtkm::cont::internal::Storage StorageType; public: VTKM_CONT_EXPORT @@ -348,6 +614,45 @@ make_ArrayHandleTransform(HandleType handle, FunctorType functor) return ArrayHandleTransform(handle,functor); } +// ArrayHandleTransform with inverse functors enabled (no need to subclass from +// ArrayHandleTransform without inverse functors: nothing to inherit). +template +class ArrayHandleTransform + : public vtkm::cont::ArrayHandle< + T, + internal::StorageTagTransform > +{ + VTKM_IS_ARRAY_HANDLE(ArrayHandleType); + +public: + VTKM_ARRAY_HANDLE_SUBCLASS( + ArrayHandleTransform, + (ArrayHandleTransform), + (vtkm::cont::ArrayHandle< + T, internal::StorageTagTransform >)); + +private: + typedef vtkm::cont::internal::Storage StorageType; + + public: + ArrayHandleTransform(const ArrayHandleType &handle, + const FunctorType &functor = FunctorType(), + const InverseFunctorType &inverseFunctor = InverseFunctorType()) + : Superclass(StorageType(handle, functor, inverseFunctor)) { } +}; + +template +VTKM_CONT_EXPORT +vtkm::cont::ArrayHandleTransform +make_ArrayHandleTransform(HandleType handle, FunctorType functor, InverseFunctorType inverseFunctor) +{ + return ArrayHandleTransform(handle,functor,inverseFunctor); +} } } // namespace vtkm::cont diff --git a/vtkm/cont/cuda/DeviceAdapterCuda.h b/vtkm/cont/cuda/DeviceAdapterCuda.h index 48f5bbbfc..4476dbb71 100644 --- a/vtkm/cont/cuda/DeviceAdapterCuda.h +++ b/vtkm/cont/cuda/DeviceAdapterCuda.h @@ -23,6 +23,11 @@ #include #ifdef VTKM_CUDA + +//This is required to be first so that we get patches for thrust included +//in the correct order +#include + #include #include #endif diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h index c57397c96..d3163292c 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h @@ -74,7 +74,7 @@ public: typedef vtkm::cont::internal::Storage 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) diff --git a/vtkm/cont/internal/ConnectivityExplicitInternals.h b/vtkm/cont/internal/ConnectivityExplicitInternals.h index 5f43b8687..580b3c467 100644 --- a/vtkm/cont/internal/ConnectivityExplicitInternals.h +++ b/vtkm/cont/internal/ConnectivityExplicitInternals.h @@ -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 diff --git a/vtkm/cont/testing/TestingFancyArrayHandles.h b/vtkm/cont/testing/TestingFancyArrayHandles.h index 2c18741a7..1958bab5d 100644 --- a/vtkm/cont/testing/TestingFancyArrayHandles.h +++ b/vtkm/cont/testing/TestingFancyArrayHandles.h @@ -421,7 +421,8 @@ private: typedef vtkm::cont::ArrayHandleIndex InputArrayType; InputArrayType input(ARRAY_SIZE); - vtkm::cont::ArrayHandleCast castArray = + vtkm::cont::ArrayHandleCast castArray = vtkm::cont::make_ArrayHandleCast(input, CastToType()); vtkm::cont::ArrayHandle result; diff --git a/vtkm/exec/CellDerivative.h b/vtkm/exec/CellDerivative.h index 0ea02a442..825833848 100644 --- a/vtkm/exec/CellDerivative.h +++ b/vtkm/exec/CellDerivative.h @@ -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(wCoords[pointIndex][0] * (weight0)); \ + jacobian(1,0) += static_cast(wCoords[pointIndex][1] * (weight0)); \ + jacobian(2,0) += static_cast(wCoords[pointIndex][2] * (weight0)); \ + jacobian(0,1) += static_cast(wCoords[pointIndex][0] * (weight1)); \ + jacobian(1,1) += static_cast(wCoords[pointIndex][1] * (weight1)); \ + jacobian(2,1) += static_cast(wCoords[pointIndex][2] * (weight1)); \ + jacobian(0,2) += static_cast(wCoords[pointIndex][0] * (weight2)); \ + jacobian(1,2) += static_cast(wCoords[pointIndex][1] * (weight2)); \ + jacobian(2,2) += static_cast(wCoords[pointIndex][2] * (weight2)) template +#include + +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 +struct Fetch +{ + 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 diff --git a/vtkm/exec/arg/ThreadIndicesBasic.h b/vtkm/exec/arg/ThreadIndicesBasic.h index f5e2daaf0..720259c3b 100644 --- a/vtkm/exec/arg/ThreadIndicesBasic.h +++ b/vtkm/exec/arg/ThreadIndicesBasic.h @@ -44,8 +44,10 @@ class ThreadIndicesBasic public: template 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; }; } diff --git a/vtkm/exec/arg/ThreadIndicesTopologyMap.h b/vtkm/exec/arg/ThreadIndicesTopologyMap.h index 7a18c0d32..94abf2b74 100644 --- a/vtkm/exec/arg/ThreadIndicesTopologyMap.h +++ b/vtkm/exec/arg/ThreadIndicesTopologyMap.h @@ -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 @@ -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; diff --git a/vtkm/exec/arg/VisitIndex.h b/vtkm/exec/arg/VisitIndex.h new file mode 100644 index 000000000..a59f63139 --- /dev/null +++ b/vtkm/exec/arg/VisitIndex.h @@ -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 +#include + +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 +struct Fetch +{ + 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 diff --git a/vtkm/exec/arg/WorkIndex.h b/vtkm/exec/arg/WorkIndex.h index f4a3cc096..092c4923a 100644 --- a/vtkm/exec/arg/WorkIndex.h +++ b/vtkm/exec/arg/WorkIndex.h @@ -63,7 +63,7 @@ struct Fetch + +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 diff --git a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx index 374219c1f..5f31e6f3b 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn.cxx @@ -20,10 +20,7 @@ #include -#include - -#include -#include +#include #include @@ -47,59 +44,36 @@ struct TestPortal } }; -struct NullParam { }; - -template +template struct FetchArrayDirectInTests { - - template - void TryInvocation(const Invocation &invocation) const + void operator()() { + TestPortal execObject; + typedef vtkm::exec::arg::Fetch< vtkm::exec::arg::FetchTagArrayDirectIn, vtkm::exec::arg::AspectTagDefault, - vtkm::exec::arg::ThreadIndicesBasic, + vtkm::exec::arg::ThreadIndicesTesting, TestPortal > 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()); + 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(), - value); + fetch.Store(indices, execObject, value); } } - void operator()() const - { - std::cout << "Trying ArrayDirectIn fetch on parameter " << ParamIndex - << " with type " << vtkm::testing::TypeName::Name() - << std::endl; - - typedef vtkm::internal::FunctionInterface< - void(NullParam,NullParam,NullParam,NullParam,NullParam)> - BaseFunctionInterface; - - this->TryInvocation(vtkm::internal::make_Invocation<1>( - BaseFunctionInterface().Replace( - TestPortal()), - NullParam(), - NullParam())); - } - }; struct TryType @@ -107,11 +81,7 @@ struct TryType template void operator()(T) const { - FetchArrayDirectInTests<1,T>()(); - FetchArrayDirectInTests<2,T>()(); - FetchArrayDirectInTests<3,T>()(); - FetchArrayDirectInTests<4,T>()(); - FetchArrayDirectInTests<5,T>()(); + FetchArrayDirectInTests()(); } }; diff --git a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx index ee9c49c91..5d7f2b5a1 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectInOut.cxx @@ -20,10 +20,7 @@ #include -#include - -#include -#include +#include #include @@ -59,19 +56,18 @@ struct TestPortal } }; -struct NullParam { }; - -template +template struct FetchArrayDirectInTests { - template - void TryInvocation(const Invocation &invocation) const + void operator()() { + TestPortal execObject; + typedef vtkm::exec::arg::Fetch< vtkm::exec::arg::FetchTagArrayDirectInOut, vtkm::exec::arg::AspectTagDefault, - vtkm::exec::arg::ThreadIndicesBasic, + vtkm::exec::arg::ThreadIndicesTesting, TestPortal > 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()); + 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(), - 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::Name() - << std::endl; - - typedef vtkm::internal::FunctionInterface< - void(NullParam,NullParam,NullParam,NullParam,NullParam)> - BaseFunctionInterface; - - this->TryInvocation(vtkm::internal::make_Invocation<1>( - BaseFunctionInterface().Replace( - TestPortal()), - NullParam(), - NullParam())); - } - }; struct TryType @@ -124,11 +99,7 @@ struct TryType template void operator()(T) const { - FetchArrayDirectInTests<1,T>()(); - FetchArrayDirectInTests<2,T>()(); - FetchArrayDirectInTests<3,T>()(); - FetchArrayDirectInTests<4,T>()(); - FetchArrayDirectInTests<5,T>()(); + FetchArrayDirectInTests()(); } }; diff --git a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx index 3d748dfe6..3ad8590e9 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchArrayDirectOut.cxx @@ -20,10 +20,7 @@ #include -#include - -#include -#include +#include #include @@ -51,19 +48,18 @@ struct TestPortal } }; -struct NullParam { }; - -template +template struct FetchArrayDirectOutTests { - template - void TryInvocation(const Invocation &invocation) const + void operator()() { + TestPortal execObject; + typedef vtkm::exec::arg::Fetch< vtkm::exec::arg::FetchTagArrayDirectOut, vtkm::exec::arg::AspectTagDefault, - vtkm::exec::arg::ThreadIndicesBasic, + vtkm::exec::arg::ThreadIndicesTesting, TestPortal > 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()); + 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(), - 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::Name() - << std::endl; - - typedef vtkm::internal::FunctionInterface< - void(NullParam,NullParam,NullParam,NullParam,NullParam)> - BaseFunctionInterface; - - this->TryInvocation(vtkm::internal::make_Invocation<1>( - BaseFunctionInterface().Replace( - TestPortal()), - NullParam(), - NullParam())); - } - }; struct TryType @@ -116,11 +91,7 @@ struct TryType template void operator()(T) const { - FetchArrayDirectOutTests<1,T>()(); - FetchArrayDirectOutTests<2,T>()(); - FetchArrayDirectOutTests<3,T>()(); - FetchArrayDirectOutTests<4,T>()(); - FetchArrayDirectOutTests<5,T>()(); + FetchArrayDirectOutTests()(); } }; diff --git a/vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx b/vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx index ce0547f76..e7e282da5 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchArrayTopologyMapIn.cxx @@ -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 + 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(connectivity) .template Replace(TestPortal()), - 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( @@ -193,8 +219,10 @@ void TryStructuredPointCoordinates( BaseFunctionInterface() .Replace<3>(connectivity) .template Replace<1>(coordinates), - NullParam(), - NullParam()) + vtkm::internal::NullType(), + vtkm::internal::NullType(), + TestIndexPortal(), + TestZeroPortal()) ); } diff --git a/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx b/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx index 5f8bcd3c6..42a1df5fb 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx @@ -20,10 +20,7 @@ #include -#include - -#include -#include +#include #include @@ -40,64 +37,37 @@ struct TestExecutionObject : public vtkm::exec::ExecutionObjectBase vtkm::Int32 Number; }; -struct NullParam { }; - -template -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()); + 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(), - execObject); + fetch.Store(indices, execObjectStore, execObject); // Data in Invocation should not have changed. - execObject = invocation.Parameters.template GetParameter(); - VTKM_TEST_ASSERT(execObject.Number == EXPECTED_NUMBER, + VTKM_TEST_ASSERT(execObjectStore.Number == EXPECTED_NUMBER, "Fetch changed read-only execution object."); } -template -void TryParamIndex() -{ - std::cout << "Trying ExecObject fetch on parameter " << ParamIndex - << std::endl; - - typedef vtkm::internal::FunctionInterface< - void(NullParam,NullParam,NullParam,NullParam,NullParam)> - BaseFunctionInterface; - - TryInvocation(vtkm::internal::make_Invocation<1>( - BaseFunctionInterface().Replace( - TestExecutionObject(EXPECTED_NUMBER)), - NullParam(), - NullParam())); -} - void TestExecObjectFetch() { - TryParamIndex<1>(); - TryParamIndex<2>(); - TryParamIndex<3>(); - TryParamIndex<4>(); - TryParamIndex<5>(); + TryInvocation(); } } // anonymous namespace diff --git a/vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx b/vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx index ae586bd60..d93ed0b6e 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchWorkIndex.cxx @@ -22,56 +22,39 @@ #include -#include - -#include -#include +#include #include namespace { -struct NullParam { }; - -template -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 *[]) diff --git a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h index eb704cefa..e63ee27a9 100644 --- a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h +++ b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h @@ -96,7 +96,7 @@ template struct load_through_texture { __device__ - static T get(const thrust::system::cuda::pointer& data) + static T get(const thrust::system::cuda::pointer& data) { return *(data.get()); } @@ -109,7 +109,7 @@ template struct load_through_texture::type >::type > { __device__ - static T get(const thrust::system::cuda::pointer& data) + static T get(const thrust::system::cuda::pointer& data) { #if __CUDA_ARCH__ >= 350 // printf("__CUDA_ARCH__ UseScalarTextureLoad"); @@ -125,7 +125,7 @@ template struct load_through_texture::type >::type > { __device__ - static T get(const thrust::system::cuda::pointer& data) + static T get(const thrust::system::cuda::pointer& data) { #if __CUDA_ARCH__ >= 350 // printf("__CUDA_ARCH__ UseVecTextureLoads"); @@ -193,7 +193,7 @@ struct load_through_texture::type NonConstT; __device__ - static T get(const thrust::system::cuda::pointer& data) + static T get(const thrust::system::cuda::pointer& data) { #if __CUDA_ARCH__ >= 350 // printf("__CUDA_ARCH__ UseMultipleScalarTextureLoads"); @@ -204,7 +204,7 @@ struct load_through_texture& data) + static T getAs(const thrust::system::cuda::pointer& 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 ) { diff --git a/vtkm/exec/cuda/internal/ThrustPatches.h b/vtkm/exec/cuda/internal/ThrustPatches.h new file mode 100644 index 000000000..df3fbb6f7 --- /dev/null +++ b/vtkm/exec/cuda/internal/ThrustPatches.h @@ -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 +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 +__device__ T destructive_accumulate_n(ConcurrentGroup &g, RandomAccessIterator first, Size n, T init, + vtkm::exec::cuda::internal::WrappedBinaryOperator 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 diff --git a/vtkm/exec/cuda/internal/WrappedOperators.h b/vtkm/exec/cuda/internal/WrappedOperators.h index 3c4e06f8d..6ca67b73e 100644 --- a/vtkm/exec/cuda/internal/WrappedOperators.h +++ b/vtkm/exec/cuda/internal/WrappedOperators.h @@ -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 > : public thrust::detail::is_arithmetic { }; -} -} +} } //namespace thrust::detail #endif //vtk_m_exec_cuda_internal_WrappedOperators_h diff --git a/vtkm/exec/internal/VecFromPortalPermute.h b/vtkm/exec/internal/VecFromPortalPermute.h index 268b00924..5dead1212 100644 --- a/vtkm/exec/internal/VecFromPortalPermute.h +++ b/vtkm/exec/internal/VecFromPortalPermute.h @@ -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 @@ -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; }; diff --git a/vtkm/exec/internal/WorkletInvokeFunctorDetail.h b/vtkm/exec/internal/WorkletInvokeFunctorDetail.h index 247638fe7..4c1719b2f 100644 --- a/vtkm/exec/internal/WorkletInvokeFunctorDetail.h +++ b/vtkm/exec/internal/WorkletInvokeFunctorDetail.h @@ -84,6 +84,8 @@ template @@ -94,14 +96,18 @@ void DoWorkletInvokeFunctor( ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -136,6 +142,8 @@ template VTKM_EXEC_EXPORT @@ -145,14 +153,18 @@ void DoWorkletInvokeFunctor( ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -180,6 +192,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -242,6 +260,8 @@ template @@ -252,14 +272,18 @@ void DoWorkletInvokeFunctor( ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -296,6 +320,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -368,6 +398,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -432,6 +468,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -514,6 +556,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -588,6 +636,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -680,6 +734,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -764,6 +824,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -866,6 +932,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -960,6 +1032,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1072,6 +1150,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1176,6 +1260,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1298,6 +1388,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1412,6 +1508,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1544,6 +1646,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1668,6 +1776,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; @@ -1810,6 +1924,8 @@ template, - InputDomainIndex> &invocation, + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> &invocation, const ThreadIndicesType &threadIndices) { typedef vtkm::internal::Invocation< ParameterInterface, ControlInterface, vtkm::internal::FunctionInterface, - InputDomainIndex> Invocation; + InputDomainIndex, + OutputToInputMapType, + VisitArrayType> Invocation; typedef InvocationToFetch FetchInfo1; typedef typename FetchInfo1::type FetchType1; diff --git a/vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in b/vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in index 73509438f..49844a135 100644 --- a/vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in +++ b/vtkm/exec/internal/WorkletInvokeFunctorDetail.h.in @@ -146,6 +146,8 @@ template 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 FetchInfo$(param_index); @@ -201,6 +207,8 @@ template 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 FetchInfo$(param_index); diff --git a/vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx b/vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx index a350d8d91..a721ac4ff 100644 --- a/vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx +++ b/vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx @@ -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."); diff --git a/vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx b/vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx index 16574a531..f31c570a5 100644 --- a/vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx +++ b/vtkm/exec/internal/testing/UnitTestWorkletInvokeFunctor.cxx @@ -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, diff --git a/vtkm/exec/testing/UnitTestParametricCoordinates.cxx b/vtkm/exec/testing/UnitTestParametricCoordinates.cxx index 99885a053..81ff68edf 100644 --- a/vtkm/exec/testing/UnitTestParametricCoordinates.cxx +++ b/vtkm/exec/testing/UnitTestParametricCoordinates.cxx @@ -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 diff --git a/vtkm/internal/Invocation.h b/vtkm/internal/Invocation.h index 7e88fb116..2a056297b 100644 --- a/vtkm/internal/Invocation.h +++ b/vtkm/internal/Invocation.h @@ -34,7 +34,9 @@ namespace internal { template + 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 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::type ChangeParameters(NewParameterInterface newParameters) const { return typename ChangeParametersType::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::type ChangeControlInterface(NewControlInterface) const { return typename ChangeControlInterfaceType::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::type ChangeExecutionInterface(NewExecutionInterface) const { return typename ChangeExecutionInterfaceType::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::type ChangeInputDomainIndex() const { return typename ChangeInputDomainIndexType::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 + struct ChangeOutputToInputMapType { + typedef Invocation type; + }; + + /// Returns a new \c Invocation that is the same as this one except that the + /// \c OutputToInputMap is replaced with that provided. + /// + template + VTKM_CONT_EXPORT + typename ChangeOutputToInputMapType::type + ChangeOutputToInputMap(NewOutputToInputMapType newOutputToInputMap) const { + return typename ChangeOutputToInputMapType::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 + struct ChangeVisitArrayType { + typedef Invocation type; + }; + + /// Returns a new \c Invocation that is the same as this one except that the + /// \c VisitArray is replaced with that provided. + /// + template + VTKM_CONT_EXPORT + typename ChangeVisitArrayType::type + ChangeVisitArray(NewVisitArrayType newVisitArray) const { + return typename ChangeVisitArrayType::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_CONT_EXPORT +vtkm::internal::Invocation +make_Invocation(const ParameterInterface ¶ms, + ControlInterface, + ExecutionInterface, + OutputToInputMapType outputToInputMap, + VisitArrayType visitArray) +{ + return vtkm::internal::Invocation(params, + outputToInputMap, + visitArray); +} template(params); + return vtkm::internal::make_Invocation( + params, + ControlInterface(), + ExecutionInterface(), + vtkm::internal::NullType(), + vtkm::internal::NullType()); } } diff --git a/vtkm/internal/ListTagDetail.h b/vtkm/internal/ListTagDetail.h index c90151847..1e4166f0a 100644 --- a/vtkm/internal/ListTagDetail.h +++ b/vtkm/internal/ListTagDetail.h @@ -46,8 +46,6 @@ struct ListRoot { }; template struct ListBase { }; -struct ListParamNull { }; - //----------------------------------------------------------------------------- template @@ -1151,21 +1149,21 @@ struct ListContainsImpl +template struct ListTagBase : detail::ListRoot { typedef detail::ListBase List; diff --git a/vtkm/internal/ListTagDetail.h.in b/vtkm/internal/ListTagDetail.h.in index 2b9b611db..047f8bc47 100644 --- a/vtkm/internal/ListTagDetail.h.in +++ b/vtkm/internal/ListTagDetail.h.in @@ -94,8 +94,6 @@ struct ListRoot { }; template struct ListBase { }; -struct ListParamNull { }; - //----------------------------------------------------------------------------- template @@ -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 List; diff --git a/vtkm/io/reader/VTKDataSetReaderBase.h b/vtkm/io/reader/VTKDataSetReaderBase.h index b4dfea5b3..344b141d8 100644 --- a/vtkm/io/reader/VTKDataSetReaderBase.h +++ b/vtkm/io/reader/VTKDataSetReaderBase.h @@ -33,11 +33,13 @@ VTKM_THIRDPARTY_PRE_INCLUDE #include #include +#include VTKM_THIRDPARTY_POST_INCLUDE #include #include #include +#include #include @@ -92,34 +94,92 @@ inline void parseAssert(bool condition) } } -struct DummyFixed8Type -{ - vtkm::UInt8 data; -}; +template struct StreamIOType { typedef T Type; }; +template <> struct StreamIOType { typedef vtkm::Int16 Type; }; +template <> struct StreamIOType { 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 struct ClosestCommonType { typedef T Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int32 Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int32 Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int32 Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int32 Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int64 Type; }; +template <> struct ClosestCommonType { typedef vtkm::Int64 Type; }; + +template struct ClosestFloat { typedef T Type; }; +template <> struct ClosestFloat { typedef vtkm::Float32 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float32 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float32 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float32 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float64 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float64 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float64 Type; }; +template <> struct ClosestFloat { typedef vtkm::Float64 Type; }; template -struct TypeTraits +vtkm::cont::DynamicArrayHandle CreateDynamicArrayHandle(const std::vector &vec) { - typedef T AsciiReadType; -}; + switch (vtkm::VecTraits::NUM_COMPONENTS) + { + case 1: + { + typedef typename ClosestCommonType::Type CommonType; + if (!boost::is_same::value) + { + std::cerr << "Type " << typeid(T).name() << " is currently unsupported. " + << "Converting to " << typeid(CommonType).name() << "." << std::endl; + } -template <> -struct TypeTraits -{ - typedef vtkm::Int16 AsciiReadType; -}; + vtkm::cont::ArrayHandle output; + output.Allocate(static_cast(vec.size())); + for (vtkm::Id i = 0; i < output.GetNumberOfValues(); ++i) + { + output.GetPortalControl().Set(i, + static_cast(vec[static_cast(i)])); + } -template <> -struct TypeTraits -{ - typedef vtkm::UInt16 AsciiReadType; -}; + return vtkm::cont::DynamicArrayHandle(output); + } + case 2: + case 3: + { + typedef typename vtkm::VecTraits::ComponentType InComponentType; + typedef typename ClosestFloat::Type OutComponentType; + typedef vtkm::Vec CommonType; + if (!boost::is_same::value) + { + std::cerr << "Type " << typeid(InComponentType).name() + << "[" << vtkm::VecTraits::NUM_COMPONENTS << "] " + << "is currently unsupported. Converting to " + << typeid(OutComponentType).name() << "[3]." << std::endl; + } -template <> -struct TypeTraits -{ - typedef vtkm::Float32 AsciiReadType; -}; + vtkm::cont::ArrayHandle output; + output.Allocate(static_cast(vec.size())); + for (vtkm::Id i = 0; i < output.GetNumberOfValues(); ++i) + { + CommonType outval = CommonType(); + for (vtkm::IdComponent j = 0; j < vtkm::VecTraits::NUM_COMPONENTS; ++j) + { + outval[j] = static_cast( + vtkm::VecTraits::GetComponent(vec[static_cast(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()); + } + } +} } // 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()); + this->SkipArray(numEntries, vtkm::Vec()); } void ReadTextureCoordinates(std::size_t numElements, std::string &dataName, @@ -544,11 +604,7 @@ private: std::vector buffer(this->NumElements); this->Reader->ReadArray(buffer); - vtkm::cont::ArrayHandle data; - data.Allocate(static_cast(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 @@ -602,7 +658,7 @@ private: { for (vtkm::IdComponent j = 0; j < numComponents; ++j) { - typename internal::TypeTraits::AsciiReadType val; + typename internal::StreamIOType::Type val; this->DataFile->Stream >> val; vtkm::VecTraits::SetComponent(buffer[i], j, static_cast(val)); @@ -647,7 +703,7 @@ private: { for (vtkm::IdComponent j = 0; j < numComponents; ++j) { - typename internal::TypeTraits::AsciiReadType val; + typename internal::StreamIOType::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 diff --git a/vtkm/io/reader/internal/TypeInfo.h b/vtkm/io/reader/internal/TypeInfo.h index 9a04fe9f6..ceff949de 100644 --- a/vtkm/io/reader/internal/TypeInfo.h +++ b/vtkm/io/reader/internal/TypeInfo.h @@ -22,6 +22,7 @@ #include +#include #include #include @@ -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(std::min(std::max(val, 1.0f), 0.0f) * 255)) + { } + operator vtkm::Float32() const + { + return static_cast(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(val); +} + +inline std::istream& operator>>(std::istream& in, ColorChannel8 &val) +{ + vtkm::Float32 fval; + in >> fval; + val = ColorChannel8(fval); + return in; +} + + template 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); diff --git a/vtkm/io/writer/VTKDataSetWriter.h b/vtkm/io/writer/VTKDataSetWriter.h index 0f142453b..7ad68cc19 100644 --- a/vtkm/io/writer/VTKDataSetWriter.h +++ b/vtkm/io/writer/VTKDataSetWriter.h @@ -156,6 +156,69 @@ public: } }; +template struct DataTypeName +{ + static const char* Name() { return "unknown"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "char"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "unsigned_char"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "short"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "unsigned_short"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "int"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "unsigned_int"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "long"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "unsigned_long"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "float"; } +}; +template <> struct DataTypeName +{ + static const char* Name() { return "double"; } +}; + +class GetDataTypeName +{ +public: + GetDataTypeName(std::string &name) + : Name(&name) + { } + + template + void operator()(const ArrayHandleType &) const + { + typedef typename vtkm::VecTraits::ComponentType + DataType; + *this->Name = DataTypeName::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)); diff --git a/vtkm/opengl/BufferState.h b/vtkm/opengl/BufferState.h new file mode 100644 index 000000000..671fb4c85 --- /dev/null +++ b/vtkm/opengl/BufferState.h @@ -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 +#include + +VTKM_THIRDPARTY_PRE_INCLUDE +#include +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 + 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 Resource; +}; + +}} + +#endif //vtk_m_opengl_BufferState_h diff --git a/vtkm/opengl/CMakeLists.txt b/vtkm/opengl/CMakeLists.txt index 96c62fed9..69281f5c0 100644 --- a/vtkm/opengl/CMakeLists.txt +++ b/vtkm/opengl/CMakeLists.txt @@ -19,6 +19,7 @@ ##============================================================================ set(headers + BufferState.h TransferToOpenGL.h ) diff --git a/vtkm/opengl/TransferToOpenGL.h b/vtkm/opengl/TransferToOpenGL.h index 6320fc594..fec93d8bd 100644 --- a/vtkm/opengl/TransferToOpenGL.h +++ b/vtkm/opengl/TransferToOpenGL.h @@ -21,37 +21,21 @@ #define vtk_m_opengl_TransferToOpenGL_h #include +#include #include 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 -VTKM_CONT_EXPORT -GLenum TransferToOpenGL(vtkm::cont::ArrayHandle handle, - GLuint& openGLHandle, - DeviceAdapterTag) -{ - vtkm::opengl::internal::TransferToOpenGL 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 handle, template VTKM_CONT_EXPORT void TransferToOpenGL(vtkm::cont::ArrayHandle handle, - GLuint& openGLHandle, - GLenum type, + BufferState& state, DeviceAdapterTag) { - vtkm::opengl::internal::TransferToOpenGL toGL(type); - toGL.Transfer(handle,openGLHandle); + vtkm::opengl::internal::TransferToOpenGL toGL(state); + return toGL.Transfer(handle); } }} diff --git a/vtkm/opengl/cuda/internal/TransferToOpenGL.h b/vtkm/opengl/cuda/internal/TransferToOpenGL.h index 2508c1801..9fb196c40 100644 --- a/vtkm/opengl/cuda/internal/TransferToOpenGL.h +++ b/vtkm/opengl/cuda/internal/TransferToOpenGL.h @@ -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 { 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 + (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 &handle, - GLuint& openGLHandle ) const + vtkm::cont::ArrayHandle &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(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(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(size); //get the device pointers typedef vtkm::cont::ArrayHandle HandleType; @@ -134,14 +215,12 @@ public: thrust::cuda::pointer(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; }; diff --git a/vtkm/opengl/internal/BufferTypePicker.h b/vtkm/opengl/internal/BufferTypePicker.h index 69b1d4ee6..2848800d2 100644 --- a/vtkm/opengl/internal/BufferTypePicker.h +++ b/vtkm/opengl/internal/BufferTypePicker.h @@ -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 diff --git a/vtkm/opengl/internal/TransferToOpenGL.h b/vtkm/opengl/internal/TransferToOpenGL.h index a647f4182..c62dd8b76 100644 --- a/vtkm/opengl/internal/TransferToOpenGL.h +++ b/vtkm/opengl/internal/TransferToOpenGL.h @@ -25,7 +25,9 @@ #include #include -#include +#include + + namespace vtkm { namespace opengl { namespace internal { @@ -37,7 +39,7 @@ template VTKM_CONT_EXPORT void CopyFromHandle( vtkm::cont::ArrayHandle& 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 VTKM_CONT_EXPORT void CopyFromHandle( vtkm::cont::ArrayHandle& 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(sizeof(ValueType)) * static_cast(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 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& handle, - GLuint& openGLHandle ) const + vtkm::cont::ArrayHandle& 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; }; } diff --git a/vtkm/opengl/testing/TestingOpenGLInterop.h b/vtkm/opengl/testing/TestingOpenGLInterop.h index c960a511a..e45c309ca 100644 --- a/vtkm/opengl/testing/TestingOpenGLInterop.h +++ b/vtkm/opengl/testing/TestingOpenGLInterop.h @@ -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) { diff --git a/vtkm/testing/Testing.h b/vtkm/testing/Testing.h index 2aa7c1c47..119b855b0 100644 --- a/vtkm/testing/Testing.h +++ b/vtkm/testing/Testing.h @@ -28,10 +28,6 @@ #include #include -VTKM_THIRDPARTY_PRE_INCLUDE -#include -VTKM_THIRDPARTY_POST_INCLUDE - #include #include #include diff --git a/vtkm/worklet/CMakeLists.txt b/vtkm/worklet/CMakeLists.txt index 7e1dd9585..bb1ecb733 100644 --- a/vtkm/worklet/CMakeLists.txt +++ b/vtkm/worklet/CMakeLists.txt @@ -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 diff --git a/vtkm/worklet/Clip.h b/vtkm/worklet/Clip.h index abbdb922f..2851faf1c 100644 --- a/vtkm/worklet/Clip.h +++ b/vtkm/worklet/Clip.h @@ -103,6 +103,20 @@ void swap(T &v1, T &v2) v2 = temp; } +template +VTKM_EXEC_CONT_EXPORT +T Scale(const T &val, vtkm::Float64 scale) +{ + return static_cast(scale * static_cast(val)); +} + +template +VTKM_EXEC_CONT_EXPORT +vtkm::Vec Scale(const vtkm::Vec &val, + vtkm::Float64 scale) +{ + return val * scale; +} template class ExecutionConnectivityExplicit : vtkm::exec::ExecutionObjectBase @@ -261,7 +275,7 @@ public: { public: typedef void ControlSignature(TopologyIn topology, - FieldInPoint scalars, + FieldInPoint scalars, FieldOutCell clipTableIdxs, FieldOutCell stats); typedef void ExecutionSignature(_2, CellShape, FromCount, _3, _4); @@ -322,7 +336,7 @@ public: { public: typedef void ControlSignature(TopologyIn topology, - FieldInPoint scalars, + FieldInPoint scalars, FieldInCell clipTableIdxs, FieldInCell 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::ComponentType w = - static_cast::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: diff --git a/vtkm/worklet/IsosurfaceUniformGrid.h b/vtkm/worklet/IsosurfaceUniformGrid.h index 787cea1ee..dc60b1454 100644 --- a/vtkm/worklet/IsosurfaceUniformGrid.h +++ b/vtkm/worklet/IsosurfaceUniformGrid.h @@ -21,20 +21,26 @@ #ifndef vtk_m_worklet_IsosurfaceUniformGrid_h #define vtk_m_worklet_IsosurfaceUniformGrid_h -#include -#include -#include -#include -#include -#include #include - -#include -#include -#include -#include #include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + + #include "MarchingCubesDataTables.h" namespace vtkm { @@ -51,197 +57,147 @@ public: public: typedef void ControlSignature(FieldInPoint inNodes, TopologyIn topology, - FieldOutCell outNumVertices); - typedef void ExecutionSignature(_1, _3); + FieldOutCell<> outNumTriangles, + ExecObject numTrianglesTable); + typedef void ExecutionSignature(_1, _3, _4); typedef _2 InputDomain; - typedef vtkm::cont::ArrayHandle IdArrayHandle; - typedef typename IdArrayHandle::ExecutionTypes::PortalConst IdPortalType; - IdPortalType VertexTable; FieldType Isovalue; VTKM_CONT_EXPORT - ClassifyCell(IdPortalType vTable, FieldType isovalue) : - VertexTable(vTable), + ClassifyCell(FieldType isovalue) : Isovalue(isovalue) { } - template + template 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 inputCellId, - FieldIn 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 pcoordIn, // Input point coordinates + FieldOutCell<> vertexOut, // Vertices for output triangles + // TODO: Have a better way to iterate over and interpolate fields + FieldInPoint 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::template ExecutionTypes::PortalConst FieldPortalType; - FieldPortalType Field, Source; - - typedef typename vtkm::cont::ArrayHandle >::template ExecutionTypes::Portal VectorPortalType; - VectorPortalType Vertices; - VectorPortalType Normals; - - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal OutputPortalType; - OutputPortalType Scalars; - - typedef typename vtkm::cont::ArrayHandle IdArrayHandle; - typedef typename IdArrayHandle::ExecutionTypes::PortalConst IdPortalType; - IdPortalType TriTable; - - const vtkm::Id cellsPerLayer, pointsPerLayer; - - template + 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 + VTKM_CONT_EXPORT + IsoSurfaceGenerate(FieldType isovalue, + const CountArrayType &countArray, + Device) + : Isovalue(isovalue), Scatter(countArray, Device()) { } + + template // 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 p[8]; - - // If we have offset and spacing, can we simplify this computation - { - vtkm::Vec offset = vtkm::make_Vec(xmin+(xmax-xmin), - ymin+(ymax-ymin), - zmin+(zmax-zmin) ); - - vtkm::Vec spacing = vtkm::make_Vec( 1.0f /((float)(xdim-1)), - 1.0f /((float)(ydim-1)), - 1.0f /((float)(zdim-1))); - - vtkm::Vec firstPoint = offset * spacing * vtkm::make_Vec( x, y, z ); - vtkm::Vec 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(cubeindex*16) + (inputIteration * 3); - for (int v = 0; v < 3; v++) + const vtkm::Id triTableOffset = + static_cast(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 edgePCoord0 = + vtkm::exec::ParametricCoordinatesPoint( + fieldIn.GetNumberOfComponents(), edgeVertex0, shape, *this); + const vtkm::Vec edgePCoord1 = + vtkm::exec::ParametricCoordinatesPoint( + fieldIn.GetNumberOfComponents(), edgeVertex1, shape, *this); + const vtkm::Vec interpPCoord = + vtkm::Lerp(edgePCoord0, edgePCoord1, interpolant); + normalsOut[triVertex] = + vtkm::Normal(vtkm::exec::CellDerivative( + fieldIn, coords, interpPCoord, shape, *this)); } - - vtkm::Vec vertex0 = this->Vertices.Get(outputVertId + 0); - vtkm::Vec vertex1 = this->Vertices.Get(outputVertId + 1); - vtkm::Vec vertex2 = this->Vertices.Get(outputVertId + 2); - - vtkm::Vec 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 > normalsArray, vtkm::cont::ArrayHandle scalarsArray) { - typedef typename vtkm::cont::DeviceAdapterAlgorithm DeviceAlgorithms; - // Set up the Marching Cubes case tables - vtkm::cont::ArrayHandle vertexTableArray = - vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numVerticesTable, + vtkm::cont::ArrayHandle numTrianglesTable = + vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::numTrianglesTable, 256); - vtkm::cont::ArrayHandle triangleTableArray = + vtkm::cont::ArrayHandle triangleTableArray = vtkm::cont::make_ArrayHandle(vtkm::worklet::internal::triTable, 256*16); + typedef vtkm::exec::ExecutionWholeArrayConst + 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 numOutputTrisPerCell; + vtkm::cont::ArrayHandle 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 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 + 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 >()), + 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) + ); } }; diff --git a/vtkm/worklet/MarchingCubesDataTables.h b/vtkm/worklet/MarchingCubesDataTables.h index 15fe72152..89bc3877a 100755 --- a/vtkm/worklet/MarchingCubesDataTables.h +++ b/vtkm/worklet/MarchingCubesDataTables.h @@ -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, diff --git a/vtkm/worklet/ScatterCounting.h b/vtkm/worklet/ScatterCounting.h new file mode 100644 index 000000000..b5b2df5d8 --- /dev/null +++ b/vtkm/worklet/ScatterCounting.h @@ -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 +#include +#include +#include +#include + +#include + +#include + +namespace vtkm { +namespace worklet { + +namespace detail { + +template +struct ReverseInputToOutputMapKernel : vtkm::exec::FunctorBase +{ + typedef typename + vtkm::cont::ArrayHandle::ExecutionTypes::PortalConst + InputMapType; + typedef typename + vtkm::cont::ArrayHandle::ExecutionTypes::Portal + OutputMapType; + typedef typename + vtkm::cont::ArrayHandle::ExecutionTypes::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 +struct SubtractToVisitIndexKernel : vtkm::exec::FunctorBase +{ + typedef typename + vtkm::cont::ArrayHandle::ExecutionTypes::PortalConst + StartsOfGroupsType; + typedef typename + vtkm::cont::ArrayHandle::ExecutionTypes::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(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 + VTKM_CONT_EXPORT + ScatterCounting(const CountArrayType &countArray, Device) + { + this->BuildArrays(countArray, Device()); + } + + typedef vtkm::cont::ArrayHandle OutputToInputMapType; + template + VTKM_CONT_EXPORT + OutputToInputMapType GetOutputToInputMap(RangeType) const + { + return this->OutputToInputMap; + } + + typedef vtkm::cont::ArrayHandle VisitArrayType; + template + 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 + 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 inputToOutputMap; + vtkm::Id outputSize = + vtkm::cont::DeviceAdapterAlgorithm::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 + VTKM_CONT_EXPORT + void BuildOutputToInputMapWithFind( + vtkm::Id outputSize, + vtkm::cont::ArrayHandle inputToOutputMap, + Device) + { + vtkm::cont::ArrayHandleIndex outputIndices(outputSize); + vtkm::cont::DeviceAdapterAlgorithm::UpperBounds( + inputToOutputMap, outputIndices, this->OutputToInputMap); + + // Do not need this anymore. + inputToOutputMap.ReleaseResources(); + + vtkm::cont::ArrayHandle startsOfGroups; + + // This find gives the index of the start of a group. + vtkm::cont::DeviceAdapterAlgorithm::LowerBounds( + this->OutputToInputMap, this->OutputToInputMap, startsOfGroups); + + detail::SubtractToVisitIndexKernel + kernel(startsOfGroups.PrepareForInput(Device()), + this->VisitArray.PrepareForOutput(outputSize, Device())); + vtkm::cont::DeviceAdapterAlgorithm::Schedule(kernel, outputSize); + } + + template + VTKM_CONT_EXPORT + void BuildOutputToInputMapWithIterate( + vtkm::Id outputSize, + vtkm::cont::ArrayHandle inputToOutputMap, + Device) + { + detail::ReverseInputToOutputMapKernel + kernel(inputToOutputMap.PrepareForInput(Device()), + this->OutputToInputMap.PrepareForOutput(outputSize, Device()), + this->VisitArray.PrepareForOutput(outputSize, Device()), + outputSize); + + vtkm::cont::DeviceAdapterAlgorithm::Schedule( + kernel, inputToOutputMap.GetNumberOfValues()); + } +}; + +} +} // namespace vtkm::worklet + +#endif //vtk_m_worklet_ScatterCounting_h diff --git a/vtkm/worklet/ScatterIdentity.h b/vtkm/worklet/ScatterIdentity.h new file mode 100644 index 000000000..49b2d488c --- /dev/null +++ b/vtkm/worklet/ScatterIdentity.h @@ -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 +#include + +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 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 + VTKM_CONT_EXPORT + RangeType GetOutputRange(RangeType inputRange) const + { + return inputRange; + } +}; + +} +} // namespace vtkm::worklet + +#endif //vtk_m_worklet_ScatterIdentity_h diff --git a/vtkm/worklet/ScatterUniform.h b/vtkm/worklet/ScatterUniform.h new file mode 100644 index 000000000..eda0e469d --- /dev/null +++ b/vtkm/worklet/ScatterUniform.h @@ -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 +#include +#include + +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(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 + OutputToInputMapType; + template + VTKM_CONT_EXPORT + OutputToInputMapType GetOutputToInputMap(RangeType inputRange) const + { + return OutputToInputMapType(detail::FunctorDiv(this->NumOutputsPerInput), + this->GetOutputRange(inputRange)); + } + + typedef vtkm::cont::ArrayHandleImplicit + VisitArrayType; + template + 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 diff --git a/vtkm/worklet/TetrahedralizeExplicitGrid.h b/vtkm/worklet/TetrahedralizeExplicitGrid.h index 32f6e2ced..8c3e5af85 100644 --- a/vtkm/worklet/TetrahedralizeExplicitGrid.h +++ b/vtkm/worklet/TetrahedralizeExplicitGrid.h @@ -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 #include -#include -#include -#include +#include #include +#include +#include +#include #include #include -#include #include +#include +#include #include -#include +#include 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 &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 &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 + VTKM_CONT_EXPORT + TriangulateCell(const CountArrayType &countArray) + : Scatter(countArray, DeviceAdapter()) + { } // Each cell produces triangles and write result at the offset - template + template VTKM_EXEC_EXPORT - void operator()(const vtkm::Id &offset, - const vtkm::Id &numIndices, - vtkm::exec::ExecutionWholeArray &connectivity, - CellShapeTag shape, - const CellNodeVecType &cellNodeIds) const + void operator()( + CellShapeTag shape, + const ConnectivityInVec &connectivityIn, + const vtkm::worklet::internal::TriangulateTablesExecutionObject &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 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 + VTKM_CONT_EXPORT + TetrahedralizeCell(const CellArrayType &cellArray) + : Scatter(cellArray, DeviceAdapter()) + { } // Each cell produces tetrahedra and write result at the offset - template + template VTKM_EXEC_EXPORT - void operator()(const vtkm::Id &offset, - vtkm::exec::ExecutionWholeArray &connectivity, - CellShapeTag shape, - const CellNodeVecType &cellNodeIds) const + void operator()(CellShapeTag shape, + const ConnectivityInVec &connectivityIn, + const vtkm::worklet::internal::TetrahedralizeTablesExecutionObject &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 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 DeviceAlgorithms; - // Cell sets belonging to input and output datasets vtkm::cont::CellSetExplicit<> &inCellSet = InDataSet.GetCellSet(0).template CastTo >(); @@ -286,56 +225,61 @@ public: vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell()); vtkm::cont::ArrayHandle inNumIndices = inCellSet.GetNumIndicesArray( vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell()); - vtkm::cont::ArrayHandle inConn = inCellSet.GetConnectivityArray( - vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell()); - // Determine the number of output cells each input cell will generate - vtkm::cont::ArrayHandle numOutCellArray; - vtkm::IdComponent verticesPerOutCell = 0; + // Output topology + vtkm::cont::ArrayHandle outConnectivity; if (dimensionality == 2) { - verticesPerOutCell = 3; - vtkm::worklet::DispatcherMapField 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 numOutCellArray; + vtkm::worklet::DispatcherMapField + triPerCellDispatcher; + triPerCellDispatcher.Invoke(inShapes, + inNumIndices, + tables.PrepareForInput(DeviceAdapter()), + numOutCellArray); + + // Build new cells + TriangulateCell triangulateWorklet(numOutCellArray); + vtkm::worklet::DispatcherMapTopology + triangulateDispatcher(triangulateWorklet); + triangulateDispatcher.Invoke( + inCellSet, + tables.PrepareForInput(DeviceAdapter()), + vtkm::cont::make_ArrayHandleGroupVec<3>(outConnectivity)); } else if (dimensionality == 3) { - verticesPerOutCell = 4; - vtkm::worklet::DispatcherMapField tetrahedraPerCellDispatcher; - tetrahedraPerCellDispatcher.Invoke(inShapes, numOutCellArray); + vtkm::worklet::internal::TetrahedralizeTables tables; + + // Determine the number of output cells each input cell will generate + vtkm::cont::ArrayHandle numOutCellArray; + vtkm::worklet::DispatcherMapField + tetPerCellDispatcher; + tetPerCellDispatcher.Invoke(inShapes, + tables.PrepareForInput(DeviceAdapter()), + numOutCellArray); + + // Build new cells + TetrahedralizeCell tetrahedralizeWorklet(numOutCellArray); + vtkm::worklet::DispatcherMapTopology + 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 cellOffset; - vtkm::Id numberOfOutCells = DeviceAlgorithms::ScanExclusive(numOutCellArray, - cellOffset); - vtkm::Id numberOfOutIndices = numberOfOutCells * verticesPerOutCell; - - // Information needed to build the output cell set - vtkm::cont::ArrayHandle connectivity; - connectivity.Allocate(numberOfOutIndices); - - // Call worklet to compute the connectivity - if (dimensionality == 2) + else { - vtkm::worklet::DispatcherMapTopology triangulateCellDispatcher; - triangulateCellDispatcher.Invoke( - cellOffset, - inNumIndices, - inCellSet, - vtkm::exec::ExecutionWholeArray(connectivity, numberOfOutIndices)); - } - else if (dimensionality == 3) - { - vtkm::worklet::DispatcherMapTopology tetrahedralizeCellDispatcher; - tetrahedralizeCellDispatcher.Invoke( - cellOffset, - inCellSet, - vtkm::exec::ExecutionWholeArray(connectivity, numberOfOutIndices)); + throw vtkm::cont::ErrorControlBadValue( + "Unsupported dimensionality for TetrahedralizeExplicitGrid."); } + // Add cells to output cellset - cellSet.Fill(connectivity); + cellSet.Fill(outConnectivity); } }; diff --git a/vtkm/worklet/TetrahedralizeUniformGrid.h b/vtkm/worklet/TetrahedralizeUniformGrid.h index 807d65eb1..1189517ef 100644 --- a/vtkm/worklet/TetrahedralizeUniformGrid.h +++ b/vtkm/worklet/TetrahedralizeUniformGrid.h @@ -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 #include -#include -#include -#include -#include +#include #include +#include +#include +#include +#include +#include #include -#include -#include - -#include +#include +#include +#include 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 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 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 VTKM_EXEC_EXPORT - void operator()(vtkm::Id &inputCellId, - vtkm::exec::ExecutionWholeArray &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 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 VTKM_EXEC_EXPORT - void operator()(vtkm::Id &inputCellId, - vtkm::exec::ExecutionWholeArray &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 >(); - // 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 >(); - 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 >(); - 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 cellIndicesArray(0, 1, numberOfInCells); - - // Output dataset depends on dimension and size vtkm::cont::ArrayHandle 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 triangulateCellDispatcher(triangulateCell); - triangulateCellDispatcher.Invoke( - cellIndicesArray, - vtkm::exec::ExecutionWholeArray(connectivity, numberOfOutIndices)); - + vtkm::cont::CellSetStructured<2> &inCellSet = + InDataSet.GetCellSet(0).template CastTo >(); + vtkm::worklet::DispatcherMapTopology 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 tetrahedralizeCellDispatcher(tetrahedralizeCell); - tetrahedralizeCellDispatcher.Invoke( - cellIndicesArray, - vtkm::exec::ExecutionWholeArray(connectivity, numberOfOutIndices)); + vtkm::cont::CellSetStructured<3> &inCellSet = + InDataSet.GetCellSet(0).template CastTo >(); + vtkm::worklet::DispatcherMapTopology dispatcher; + dispatcher.Invoke(inCellSet, + vtkm::cont::make_ArrayHandleGroupVec<4>(connectivity)); + } + else + { + throw vtkm::cont::ErrorControlBadValue( + "Unsupported dimensionality for TetrahedralizeUniformGrid."); } // Add cells to output cellset diff --git a/vtkm/worklet/WorkletMapTopology.h b/vtkm/worklet/WorkletMapTopology.h index 640391e3a..1415ac3da 100644 --- a/vtkm/worklet/WorkletMapTopology.h +++ b/vtkm/worklet/WorkletMapTopology.h @@ -188,6 +188,8 @@ public: template struct FieldInOutCell : FieldInOut { }; + + struct PointIndices : FromIndices { }; }; } diff --git a/vtkm/worklet/internal/CMakeLists.txt b/vtkm/worklet/internal/CMakeLists.txt index c839e33c6..edc2dce69 100644 --- a/vtkm/worklet/internal/CMakeLists.txt +++ b/vtkm/worklet/internal/CMakeLists.txt @@ -19,10 +19,11 @@ ##============================================================================ set(headers + ClipTables.h DispatcherBase.h DispatcherBaseDetailInvoke.h + TriangulateTables.h WorkletBase.h - ClipTables.h ) set_source_files_properties(DispatcherBaseDetailInvoke.h diff --git a/vtkm/worklet/internal/DispatcherBase.h b/vtkm/worklet/internal/DispatcherBase.h index c355890f1..c9b71b0eb 100644 --- a/vtkm/worklet/internal/DispatcherBase.h +++ b/vtkm/worklet/internal/DispatcherBase.h @@ -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 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 diff --git a/vtkm/worklet/internal/TriangulateTables.h b/vtkm/worklet/internal/TriangulateTables.h new file mode 100644 index 000000000..4cc1f5e91 --- /dev/null +++ b/vtkm/worklet/internal/TriangulateTables.h @@ -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 +#include + +#include + +#include +#include + +namespace vtkm { +namespace worklet { +namespace internal { + +typedef vtkm::cont::ArrayHandle + 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 +class TriangulateTablesExecutionObject : public vtkm::exec::ExecutionObjectBase +{ +public: + typedef typename TriangulateArrayHandle::ExecutionTypes::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 + 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 + VTKM_EXEC_EXPORT + vtkm::Vec + GetIndices(CellShape shape, vtkm::IdComponent triangleIndex) const + { + vtkm::Vec 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 + vtkm::worklet::internal::TriangulateTablesExecutionObject + PrepareForInput(Device) const + { + return vtkm::worklet::internal::TriangulateTablesExecutionObject( + 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 +class TetrahedralizeTablesExecutionObject + : public vtkm::exec::ExecutionObjectBase +{ +public: + typedef typename TriangulateArrayHandle::ExecutionTypes::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 + VTKM_EXEC_EXPORT + vtkm::IdComponent GetCount(CellShape shape) const + { + return this->Counts.Get(shape.Id); + } + + template + VTKM_EXEC_EXPORT + vtkm::Vec + GetIndices(CellShape shape, vtkm::IdComponent tetrahedronIndex) const + { + vtkm::Vec 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 + vtkm::worklet::internal::TetrahedralizeTablesExecutionObject + PrepareForInput(Device) const + { + return vtkm::worklet::internal::TetrahedralizeTablesExecutionObject( + this->Counts, this->Offsets, this->Indices); + } + +private: + TriangulateArrayHandle Counts; + TriangulateArrayHandle Offsets; + TriangulateArrayHandle Indices; +}; + +} +} +} + +#endif //vtk_m_worklet_internal_TriangulateTables_h diff --git a/vtkm/worklet/internal/WorkletBase.h b/vtkm/worklet/internal/WorkletBase.h index 01c79c1e9..4a7be7b57 100644 --- a/vtkm/worklet/internal/WorkletBase.h +++ b/vtkm/worklet/internal/WorkletBase.h @@ -25,13 +25,17 @@ #include #include #include +#include #include +#include #include #include #include #include +#include + 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 diff --git a/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx b/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx index a4f2defd8..efc0503bb 100644 --- a/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx +++ b/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx @@ -119,7 +119,7 @@ struct Fetch(x) / static_cast(xdim-1); - const float fy = static_cast(y) / static_cast(xdim-1); - const float fz = static_cast(z) / static_cast(xdim-1); + const vtkm::FloatDefault fx = static_cast(x) / static_cast(xdim-1); + const vtkm::FloatDefault fy = static_cast(y) / static_cast(xdim-1); + const vtkm::FloatDefault fz = static_cast(z) / static_cast(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 fieldArray; vtkm::cont::ArrayHandleIndex vertexCountImplicitArray(vdims[0]*vdims[1]*vdims[2]); vtkm::worklet::DispatcherMapField tangleFieldDispatcher(TangleField(vdims, mins, maxs)); tangleFieldDispatcher.Invoke(vertexCountImplicitArray, fieldArray); - vtkm::cont::ArrayHandleUniformPointCoordinates coordinates(vdims); + vtkm::Vec origin(0.0f, 0.0f, 0.0f); + vtkm::Vec spacing( + 1.0f/static_cast(dims[0]), + 1.0f/static_cast(dims[2]), + 1.0f/static_cast(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"); } diff --git a/vtkm/worklet/testing/UnitTestScatterCounting.cxx b/vtkm/worklet/testing/UnitTestScatterCounting.cxx new file mode 100644 index 000000000..60653c117 --- /dev/null +++ b/vtkm/worklet/testing/UnitTestScatterCounting.cxx @@ -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 + +#include +#include + +#include +#include + +#include + +#include + +namespace { + +struct TestScatterArrays +{ + vtkm::cont::ArrayHandle CountArray; + vtkm::cont::ArrayHandle OutputToInputMap; + vtkm::cont::ArrayHandle 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 + 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 + 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 + 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 + VTKM_CONT_EXPORT + TestScatterCountingWorklet(const CountArrayType &countArray) + : Scatter(countArray, VTKM_DEFAULT_DEVICE_ADAPTER_TAG()) { } + + template + 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 +void CompareArrays(vtkm::cont::ArrayHandle array1, + vtkm::cont::ArrayHandle array2) +{ + typedef typename vtkm::cont::ArrayHandle::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 dispatcher( + worklet); + + vtkm::Id inputSize = arrays.CountArray.GetNumberOfValues(); + vtkm::cont::ArrayHandleIndex inputIndices(inputSize); + vtkm::cont::ArrayHandle outputToInputMapCopy; + vtkm::cont::ArrayHandle visitCopy; + vtkm::cont::ArrayHandle 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); +}