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

2024-09-08 13:23:51 +00:00 · 2018-12-03 07:44:31 -08:00 · 2018-12-03 07:44:31 -08:00 · c062f2e26c
commit c062f2e26c
parent 0fa7335294 bc13764228
616 changed files with 50277 additions and 14046 deletions
--- a/CMake/VTKmCompilerDynamicAnalysisFlags.cmake
+++ b/CMake/VTKmCompilerDynamicAnalysisFlags.cmake
@ -0,0 +1,42 @@
+##============================================================================
+##  Copyright (c) Kitware, Inc.
+##  All rights reserved.
+##  See LICENSE.txt for details.
+##  This software is distributed WITHOUT ANY WARRANTY; without even
+##  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+##  PURPOSE.  See the above copyright notice for more information.
+##
+##  Copyright 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
+##  Copyright 2014 UT-Battelle, LLC.
+##  Copyright 2014 Los Alamos National Security.
+##
+##  Under the terms of Contract DE-NA0003525 with NTESS,
+##  the U.S. Government retains certain rights in this software.
+##
+##  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+##  Laboratory (LANL), the U.S. Government retains certain rights in
+##  this software.
+##============================================================================
+
+#-----------------------------------------------------------------------------
+# check if this is a sanitizer build. If so, set up the environment.
+
+function(vtkm_check_sanitizer_build)
+  string (FIND "${CTEST_MEMORYCHECK_TYPE}" "Sanitizer" SANITIZER_BUILD)
+  if (${SANITIZER_BUILD} GREATER -1)
+    # This is a sanitizer build.
+    # Configure the sanitizer blacklist file
+    set (SANITIZER_BLACKLIST "${VTKm_BINARY_DIR}/sanitizer_blacklist.txt")
+    configure_file (
+      "${VTKm_SOURCE_DIR}/Utilities/DynamicAnalysis/sanitizer_blacklist.txt.in"
+      ${SANITIZER_BLACKLIST}
+      @ONLY
+      )
+
+    # Add the compiler flags for blacklist
+    set (FSANITIZE_BLACKLIST "\"-fsanitize-blacklist=${SANITIZER_BLACKLIST}\"")
+    foreach (entity C CXX SHARED_LINKER EXE_LINKER MODULE_LINKER)
+      set (CMAKE_${entity}_FLAGS "${CMAKE_${entity}_FLAGS} ${FSANITIZE_BLACKLIST}")
+    endforeach ()
+  endif ()
+endfunction()
--- a/CMake/VTKmCompilerFlags.cmake
+++ b/CMake/VTKmCompilerFlags.cmake
@ -51,8 +51,9 @@ add_library(vtkm_compiler_flags INTERFACE)
 # When building libraries/tests that are part of the VTK-m repository
 # inherit the properties from vtkm_developer_flags and vtkm_vectorization_flags.
 # The flags are intended only for VTK-m itself and are not needed by consumers.
-# We will export vtkm_vectorization_flags in general so consumer can enable
-# vectorization if they so desire
+# We will export vtkm_developer_flags, and vtkm_vectorization_flags in general
+# so consumer can either enable vectorization or use VTK-m's build flags if
+# they so desire
 if (VTKm_ENABLE_DEVELOPER_FLAGS)
  target_link_libraries(vtkm_compiler_flags
    INTERFACE $<BUILD_INTERFACE:vtkm_developer_flags>)
@ -92,26 +93,29 @@ add_library(vtkm_developer_flags INTERFACE)
 # about failures to vectorize.
 if (CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND
    CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 3.4)
-  target_compile_options(vtkm_developer_flags INTERFACE $<BUILD_INTERFACE:$<$<COMPILE_LANGUAGE:CXX>:-Wno-pass-failed>>)
+  target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CXX>:-Wno-pass-failed>)
 elseif(CMAKE_CXX_COMPILER_ID STREQUAL "AppleClang" AND
       CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 6.99)
-  target_compile_options(vtkm_developer_flags INTERFACE $<BUILD_INTERFACE:$<$<COMPILE_LANGUAGE:CXX>:-Wno-pass-failed>>)
+  target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CXX>:-Wno-pass-failed>)
 endif()

 if(VTKM_COMPILER_IS_MSVC)
  target_compile_definitions(vtkm_developer_flags INTERFACE "_SCL_SECURE_NO_WARNINGS"
                                                            "_CRT_SECURE_NO_WARNINGS")

-  #CMake COMPILE_LANGUAGE doesn't work with MSVC, ans since we want these flags
-  #only for C++ compilation we have to resort to setting CMAKE_CXX_FLAGS :(
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4702 /wd4505")
-  set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=\"/wd4702 /wd4505\" -Xcudafe=\"--diag_suppress=1394 --diag_suppress=766 --display_error_number\"")
+  #Setup MSVC warnings with CUDA and CXX
+  target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CXX>:-wd4702 -wd4505>)
+  if(TARGET vtkm::cuda)
+    target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-wd4702,-wd4505  -Xcudafe=--display_error_number,--diag_suppress=1394,--diag_suppress=766>)
+  endif()

  if(MSVC_VERSION LESS 1900)
    # In VS2013 the C4127 warning has a bug in the implementation and
    # generates false positive warnings for lots of template code
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -wd4127")
-    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=\"/wd4127\"")
+    target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CXX>:-wd4127>)
+    if(TARGET vtkm::cuda)
+      target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-wd4127>)
+    endif()
  endif()

 elseif(VTKM_COMPILER_IS_ICC)
@ -137,13 +141,9 @@ elseif(VTKM_COMPILER_IS_GNU OR VTKM_COMPILER_IS_CLANG)
    (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 6.99) )
    list(APPEND cxx_flags -Wno-strict-overflow)
  endif()
-  target_compile_options(vtkm_developer_flags
-    INTERFACE $<BUILD_INTERFACE:$<$<COMPILE_LANGUAGE:CXX>:${cxx_flags}>>
-    )
+  target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CXX>:${cxx_flags}>)
  if(TARGET vtkm::cuda)
-    target_compile_options(vtkm_developer_flags
-      INTERFACE $<BUILD_INTERFACE:$<$<COMPILE_LANGUAGE:CUDA>:${cuda_flags}>>
-      )
+    target_compile_options(vtkm_developer_flags INTERFACE $<$<COMPILE_LANGUAGE:CUDA>:${cuda_flags}>)
  endif()
 endif()

--- a/CMake/VTKmDeviceAdapters.cmake
+++ b/CMake/VTKmDeviceAdapters.cmake
@ -72,6 +72,16 @@ if(VTKm_ENABLE_TBB AND NOT TARGET vtkm::tbb)
      IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
      IMPORTED_LOCATION_RELEASE "${real_path}"
      )
+  elseif(EXISTS "${TBB_LIBRARY}")
+    #When VTK-m is mixed with OSPray we could use the OSPray FindTBB file
+    #which doesn't define TBB_LIBRARY_RELEASE but instead defined only
+    #TBB_LIBRARY
+    vtkm_extract_real_library("${TBB_LIBRARY}" real_path)
+    set_property(TARGET vtkm::tbb APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
+    set_target_properties(vtkm::tbb PROPERTIES
+      IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
+      IMPORTED_LOCATION_RELEASE "${real_path}"
+      )
  endif()

  if(EXISTS "${TBB_LIBRARY_DEBUG}")
@ -188,18 +198,21 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
  #   - Uses: --generate-code=arch=compute_60,code=sm_60
  # 6 - volta
  #   - Uses: --generate-code=arch=compute_70,code=sm_70
-  # 7 - all
+  # 7 - turing
+  #   - Uses: --generate-code=arch=compute_75code=sm_75
+  # 8 - all
  #   - Uses: --generate-code=arch=compute_30,code=sm_30
  #   - Uses: --generate-code=arch=compute_35,code=sm_35
  #   - Uses: --generate-code=arch=compute_50,code=sm_50
  #   - Uses: --generate-code=arch=compute_60,code=sm_60
  #   - Uses: --generate-code=arch=compute_70,code=sm_70
+  #   - Uses: --generate-code=arch=compute_75,code=sm_75
  # 8 - none
  #

  #specify the property
  set(VTKm_CUDA_Architecture "native" CACHE STRING "Which GPU Architecture(s) to compile for")
-  set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta all none)
+  set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing all none)

  #detect what the property is set too
  if(VTKm_CUDA_Architecture STREQUAL "native")
@ -253,12 +266,15 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
    set(arch_flags --generate-code=arch=compute_60,code=sm_60)
  elseif(VTKm_CUDA_Architecture STREQUAL "volta")
    set(arch_flags --generate-code=arch=compute_70,code=sm_70)
+  elseif(VTKm_CUDA_Architecture STREQUAL "turing")
+    set(arch_flags --generate-code=arch=compute_75,code=sm_75)
  elseif(VTKm_CUDA_Architecture STREQUAL "all")
    set(arch_flags --generate-code=arch=compute_30,code=sm_30
                   --generate-code=arch=compute_35,code=sm_35
                   --generate-code=arch=compute_50,code=sm_50
                   --generate-code=arch=compute_60,code=sm_60
-                   --generate-code=arch=compute_70,code=sm_70)
+                   --generate-code=arch=compute_70,code=sm_70
+                   --generate-code=arch=compute_75,code=sm_75)
  endif()

  string(REPLACE ";" " " arch_flags "${arch_flags}")
--- a/CMake/VTKmWrappers.cmake
+++ b/CMake/VTKmWrappers.cmake
@ -95,6 +95,13 @@ endfunction()
 function(vtkm_add_header_build_test name dir_prefix use_cuda)
  set(hfiles ${ARGN})

+  #only attempt to add a test build executable if we have any headers to
+  #test. this might not happen when everything depends on thrust.
+  list(LENGTH hfiles num_srcs)
+  if (${num_srcs} EQUAL 0)
+    return()
+  endif()
+
  set(ext "cxx")
  if(use_cuda)
    set(ext "cu")
@ -104,10 +111,9 @@ function(vtkm_add_header_build_test name dir_prefix use_cuda)
  foreach (header ${hfiles})
    get_source_file_property(cant_be_tested ${header} VTKm_CANT_BE_HEADER_TESTED)
    if( NOT cant_be_tested )
-      get_filename_component(headername ${header} NAME_WE)
-      get_filename_component(headerextension ${header} EXT)
-      string(SUBSTRING ${headerextension} 1 -1 headerextension)
-      set(src ${CMAKE_CURRENT_BINARY_DIR}/TB_${headername}_${headerextension}.${ext})
+      string(REPLACE "/" "_" headername "${header}")
+      string(REPLACE "." "_" headername "${headername}")
+      set(src ${CMAKE_CURRENT_BINARY_DIR}/TB_${headername}.${ext})

      #By using file generate we will not trigger CMake execution when
      #a header gets touched
@ -118,9 +124,10 @@ function(vtkm_add_header_build_test name dir_prefix use_cuda)
 //This is used by headers that include thrust to properly define a proper
 //device backend / system
 #define VTKM_TEST_HEADER_BUILD
-#include <${dir_prefix}/${headername}.${headerextension}>
+#include <${dir_prefix}/${header}>
 int ${headername}_${headerextension}_testbuild_symbol;"
        )
+
      list(APPEND srcs ${src})
    endif()
  endforeach()
@ -129,13 +136,6 @@ int ${headername}_${headerextension}_testbuild_symbol;"
    PROPERTIES HEADER_FILE_ONLY TRUE
    )

-  #only attempt to add a test build executable if we have any headers to
-  #test. this might not happen when everything depends on thrust.
-  list(LENGTH srcs num_srcs)
-  if (${num_srcs} EQUAL 0)
-    return()
-  endif()
-
  if(TARGET TestBuild_${name})
    #If the target already exists just add more sources to it
    target_sources(TestBuild_${name} PRIVATE ${srcs})
@ -206,6 +206,7 @@ function(vtkm_generate_export_header lib_name)

 endfunction(vtkm_generate_export_header)

+#-----------------------------------------------------------------------------
 function(vtkm_install_headers dir_prefix)
  if(NOT VTKm_INSTALL_ONLY_LIBRARIES)
    set(hfiles ${ARGN})
@ -345,6 +346,7 @@ endfunction(vtkm_library)
 #   LIBRARIES <dependent_library_list>
 #   TEST_ARGS <argument_list>
 #   MPI
+#   ALL_BACKENDS
 #   <options>
 #   )
 #
@ -362,6 +364,8 @@ endfunction(vtkm_library)
 #
 # [MPI]       : when specified, the tests should be run in parallel if
 #               MPI is enabled.
+# [ALL_BACKENDS] : when specified, the tests would test against all enabled
+#                  backends. BACKEND argument would be ignored.
 #
 function(vtkm_unit_tests)
  if (NOT VTKm_ENABLE_TESTING)
@ -369,7 +373,7 @@ function(vtkm_unit_tests)
  endif()

  set(options)
-  set(global_options ${options} MPI)
+  set(global_options ${options} MPI ALL_BACKENDS)
  set(oneValueArgs BACKEND NAME)
  set(multiValueArgs SOURCES LIBRARIES TEST_ARGS)
  cmake_parse_arguments(VTKm_UT
@ -378,17 +382,33 @@ function(vtkm_unit_tests)
    )
  vtkm_parse_test_options(VTKm_UT_SOURCES "${options}" ${VTKm_UT_SOURCES})

-  set(test_prog )
+  set(test_prog)
  set(backend ${VTKm_UT_BACKEND})

+  set(enable_all_backends ${VTKm_UT_ALL_BACKENDS})
+  set(all_backends SERIAL)
+  if (VTKm_ENABLE_CUDA)
+    list(APPEND all_backends CUDA)
+  endif()
+  if (VTKm_ENABLE_TBB)
+    list(APPEND all_backends TBB)
+  endif()
+  if (VTKm_ENABLE_OPENMP)
+    list(APPEND all_backends OPENMP)
+  endif()
+
  if(VTKm_UT_NAME)
    set(test_prog "${VTKm_UT_NAME}")
  else()
    vtkm_get_kit_name(kit)
    set(test_prog "UnitTests_${kit}")
  endif()
+
  if(backend)
    set(test_prog "${test_prog}_${backend}")
+    set(all_backends ${backend})
+  elseif(NOT enable_all_backends)
+    set (all_backends "NO_BACKEND")
  endif()

  if(VTKm_UT_MPI)
@ -403,7 +423,8 @@ function(vtkm_unit_tests)
  #the creation of the test source list needs to occur before the labeling as
  #cuda. This is so that we get the correctly named entry points generated
  create_test_sourcelist(test_sources ${test_prog}.cxx ${VTKm_UT_SOURCES} ${extraArgs})
-  if(backend STREQUAL "CUDA")
+  #if all backends are enabled, we can use cuda compiler to handle all possible backends.
+  if(backend STREQUAL "CUDA" OR (enable_all_backends AND VTKm_ENABLE_CUDA))
    vtkm_compile_as_cuda(cu_srcs ${VTKm_UT_SOURCES})
    set(VTKm_UT_SOURCES ${cu_srcs})
  endif()
@ -414,43 +435,50 @@ function(vtkm_unit_tests)
  set_property(TARGET ${test_prog} PROPERTY RUNTIME_OUTPUT_DIRECTORY ${VTKm_EXECUTABLE_OUTPUT_PATH})

  target_link_libraries(${test_prog} PRIVATE vtkm_cont ${VTKm_UT_LIBRARIES})
-  if(backend)
-    target_compile_definitions(${test_prog} PRIVATE "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_${backend}")
-  endif()

-  #determine the timeout for all the tests based on the backend. CUDA tests
-  #generally require more time because of kernel generation.
-  set(timeout 180)
-  set(run_serial False)
-  if(backend STREQUAL "CUDA")
-    set(timeout 1500)
-  elseif(backend STREQUAL "OPENMP")
-    #We need to have all OpenMP tests run serially as they
-    #will uses all the system cores, and we will cause a N*N thread
-    #explosion which causes the tests to run slower than when run
-    #serially
-    set(run_serial True)
-  endif()
-
-
-  foreach (test ${VTKm_UT_SOURCES})
-    get_filename_component(tname ${test} NAME_WE)
-    if(VTKm_UT_MPI AND VTKm_ENABLE_MPI)
-      add_test(NAME ${tname}${backend}
-        COMMAND ${MPIEXEC} ${MPIEXEC_NUMPROC_FLAG} 3 ${MPIEXEC_PREFLAGS}
-                $<TARGET_FILE:${test_prog}> ${tname} ${VTKm_UT_TEST_ARGS}
-                ${MPIEXEC_POSTFLAGS}
-        )
-    else()
-      add_test(NAME ${tname}${backend}
-        COMMAND ${test_prog} ${tname} ${VTKm_UT_TEST_ARGS}
-        )
+  foreach(current_backend ${all_backends})
+    set (device_command_line_argument --device=${current_backend})
+    if (current_backend STREQUAL "NO_BACKEND")
+      set (current_backend "")
+      set(device_command_line_argument "")
    endif()
-    set_tests_properties("${tname}${backend}" PROPERTIES
-      TIMEOUT ${timeout}
-      RUN_SERIAL ${run_serial}
-    )
-  endforeach (test)
+    foreach (test ${VTKm_UT_SOURCES})
+      get_filename_component(tname ${test} NAME_WE)
+      if(VTKm_UT_MPI AND VTKm_ENABLE_MPI)
+        add_test(NAME ${tname}${current_backend}
+          COMMAND ${MPIEXEC} ${MPIEXEC_NUMPROC_FLAG} 3 ${MPIEXEC_PREFLAGS}
+                  $<TARGET_FILE:${test_prog}> ${tname} ${device_command_line_argument} ${VTKm_UT_TEST_ARGS}
+                  ${MPIEXEC_POSTFLAGS}
+          )
+      else()
+        add_test(NAME ${tname}${current_backend}
+          COMMAND ${test_prog} ${tname} ${device_command_line_argument} ${VTKm_UT_TEST_ARGS}
+          )
+      endif()
+
+      #determine the timeout for all the tests based on the backend. CUDA tests
+      #generally require more time because of kernel generation.
+      if (current_backend STREQUAL "CUDA")
+        set(timeout 1500)
+      else()
+        set(timeout 180)
+      endif()
+      if(current_backend STREQUAL "OPENMP")
+        #We need to have all OpenMP tests run serially as they
+        #will uses all the system cores, and we will cause a N*N thread
+        #explosion which causes the tests to run slower than when run
+        #serially
+        set(run_serial True)
+      else()
+        set(run_serial False)
+      endif()
+
+      set_tests_properties("${tname}${current_backend}" PROPERTIES
+        TIMEOUT ${timeout}
+        RUN_SERIAL ${run_serial}
+      )
+    endforeach (test)
+  endforeach(current_backend)

 endfunction(vtkm_unit_tests)

--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@ -97,6 +97,7 @@ vtkm_option(VTKm_ENABLE_OPENMP "Enable OpenMP support" OFF)
 vtkm_option(VTKm_ENABLE_RENDERING "Enable rendering library" ON)
 vtkm_option(VTKm_ENABLE_TESTING "Enable VTKm Testing" ON)
 vtkm_option(VTKm_ENABLE_BENCHMARKS "Enable VTKm Benchmarking" OFF)
+vtkm_option(VTKm_ENABLE_LOGGING "Enable VTKm Logging" OFF)
 vtkm_option(VTKm_ENABLE_MPI "Enable MPI support" OFF)

 vtkm_option(VTKm_ENABLE_DOCUMENTATION "Build Doxygen documentation" OFF)
@ -199,6 +200,11 @@ if (VTKm_ENABLE_TESTING)
  # Find Pyexpander in case somebody wants to update the auto generated
  # faux variadic template code
  find_package(Pyexpander QUIET)
+
+  #-----------------------------------------------------------------------------
+  # Setup compiler flags for dynamic analysis if needed
+  include(VTKmCompilerDynamicAnalysisFlags)
+  vtkm_check_sanitizer_build()
 endif (VTKm_ENABLE_TESTING)

 #-----------------------------------------------------------------------------
@ -288,7 +294,7 @@ if(NOT VTKm_INSTALL_ONLY_LIBRARIES)

  # Create and install exports for external projects
  export(EXPORT ${VTKm_EXPORT_NAME}
-    FILE ${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmTargets.cmake
+    FILE ${VTKm_BUILD_CMAKE_BASE_DIR}/${VTKm_INSTALL_CONFIG_DIR}/VTKmTargets.cmake
    )
  install(EXPORT ${VTKm_EXPORT_NAME}
    DESTINATION ${VTKm_INSTALL_CONFIG_DIR}
--- a/README.md
+++ b/README.md
@ -71,6 +71,7 @@ Optional dependencies are:

  + CUDA Device Adapter
      + [Cuda Toolkit 7.5+](https://developer.nvidia.com/cuda-toolkit)
+      + Note CUDA >= 10.0 is required on Windows
  + TBB Device Adapter
      + [TBB](https://www.threadingbuildingblocks.org/)
  + OpenMP Device Adapter
@ -95,19 +96,20 @@ Optional dependencies are:
      + EGL Driver

 VTK-m has been tested on the following configurations:
-  + On Linux 
-      + GCC 4.8.5, 5.4.0, 6.4.0, Clang 3.8.0
-      + CMake 3.9.2, 3.9.3, 3.10.3
-      + CUDA 8.0.61, 9.1.85
+  + On Linux
+      + GCC 4.8.5, 5.4.0, 6.4.0, 7.3.0 Clang 3.8.0, Intel 17.0.4
+      + CMake 3.9.3, 3.10.3
+      + CUDA 8.0.61, 9.1.85, 10.0.130
      + TBB 4.4 U2, 2017 U7
  + On Windows
      + Visual Studio 2015, 2017
      + CMake 3.3, 3.11.1
-      + CUDA 9.1.85
+      + CUDA 10.0.130
      + TBB 2017 U3, 2018 U2
  + On MacOS
-      + AppleClang 6.0
-      + TBB 2017 U6
+      + AppleClang 9.1
+      + CMake 3.12.0
+      + TBB 2018


 ## Building ##
@ -133,7 +135,7 @@ Users Guide].
 ## Example##

 The VTK-m source distribution includes a number of examples. The goal of the
-VTK-m examples is to illustrate specific VTK-m concepts in a consistent and 
+VTK-m examples is to illustrate specific VTK-m concepts in a consistent and
 simple format. However, these examples only cover a small part of the
 capabilities of VTK-m.

@ -204,7 +206,7 @@ See [LICENSE.txt](LICENSE.txt) for details.
 [VTK-m download page]:      http://m.vtk.org/index.php/VTK-m_Releases
 [VTK-m git repository]:     https://gitlab.kitware.com/vtk/vtk-m/
 [VTK-m Issue Tracker]:      https://gitlab.kitware.com/vtk/vtk-m/issues
-[VTK-m Overview]:           http://m.vtk.org/images/2/29/VTKmVis2016.pptx 
+[VTK-m Overview]:           http://m.vtk.org/images/2/29/VTKmVis2016.pptx
 [VTK-m Users Guide]:        http://m.vtk.org/images/c/c8/VTKmUsersGuide.pdf
 [VTK-m users email list]:   http://vtk.org/mailman/listinfo/vtkm
 [VTK-m Wiki]:               http://m.vtk.org/
--- a/Utilities/DynamicAnalysis/lsan.supp
+++ b/Utilities/DynamicAnalysis/lsan.supp
@ -0,0 +1 @@
+leak:libX11*
--- a/Utilities/DynamicAnalysis/sanitizer_blacklist.txt.in
+++ b/Utilities/DynamicAnalysis/sanitizer_blacklist.txt.in
@ -0,0 +1,2 @@
+# Blacklist third party libraries from invoking sanitizer errors
+src:@VTKm_SOURCE_DIR@/vtkm/thirdparty/*
--- a/benchmarking/BenchmarkDeviceAdapter.cxx
+++ b/benchmarking/BenchmarkDeviceAdapter.cxx
@ -698,7 +698,7 @@ private:
  template <typename Value>
  struct BenchStableSortIndices
  {
-    using SSI = vtkm::worklet::StableSortIndices<DeviceAdapterTag>;
+    using SSI = vtkm::worklet::StableSortIndices;
    using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;

    ValueArrayHandle ValueHandle;
@ -743,7 +743,7 @@ private:
  template <typename Value>
  struct BenchStableSortIndicesUnique
  {
-    using SSI = vtkm::worklet::StableSortIndices<DeviceAdapterTag>;
+    using SSI = vtkm::worklet::StableSortIndices;
    using IndexArrayHandle = typename SSI::IndexArrayType;
    using ValueArrayHandle = vtkm::cont::ArrayHandle<Value, StorageTag>;

@ -915,6 +915,7 @@ public:
  static VTKM_CONT int Run()
  {
    std::cout << DIVIDER << "\nRunning DeviceAdapter benchmarks\n";
+    vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(DeviceAdapterTag());

    // Run fixed bytes / size tests:
    for (int sizeType = 0; sizeType < 2; ++sizeType)
--- a/benchmarking/BenchmarkFilters.cxx
+++ b/benchmarking/BenchmarkFilters.cxx
@ -32,6 +32,7 @@
 #include <vtkm/cont/CellSetStructured.h>
 #include <vtkm/cont/DataSet.h>
 #include <vtkm/cont/ErrorInternal.h>
+#include <vtkm/cont/Logging.h>
 #include <vtkm/cont/RuntimeDeviceTracker.h>
 #include <vtkm/cont/StorageBasic.h>
 #include <vtkm/cont/Timer.h>
@ -1282,6 +1283,8 @@ int BenchmarkBody(int argc, char* argv[])

 int main(int argc, char* argv[])
 {
+  vtkm::cont::InitLogging(argc, argv);
+
  int retval = 1;
  try
  {
--- a/benchmarking/BenchmarkRayTracing.cxx
+++ b/benchmarking/BenchmarkRayTracing.cxx
@ -28,12 +28,15 @@
 #include <vtkm/cont/testing/MakeTestDataSet.h>

 #include <vtkm/rendering/Camera.h>
-#include <vtkm/rendering/internal/RunTriangulator.h>
 #include <vtkm/rendering/raytracing/Ray.h>
 #include <vtkm/rendering/raytracing/RayTracer.h>
+#include <vtkm/rendering/raytracing/SphereIntersector.h>
+#include <vtkm/rendering/raytracing/TriangleExtractor.h>

 #include <vtkm/exec/FunctorBase.h>

+#include <vtkm/cont/ColorTable.hxx>
+
 #include <sstream>
 #include <string>
 #include <vector>
@ -66,18 +69,24 @@ struct BenchRayTracing
    camera.ResetToBounds(bounds);

    vtkm::cont::DynamicCellSet cellset = Data.GetCellSet();
-    vtkm::rendering::internal::RunTriangulator(cellset, Indices, NumberOfTriangles);
+
+    vtkm::rendering::raytracing::TriangleExtractor triExtractor;
+    triExtractor.ExtractCells(cellset);
+    vtkm::rendering::raytracing::TriangleIntersector* triIntersector =
+      new vtkm::rendering::raytracing::TriangleIntersector();
+    triIntersector->SetData(Coords, triExtractor.GetTriangles());
+    Tracer.AddShapeIntersector(triIntersector);

    vtkm::rendering::CanvasRayTracer canvas(1920, 1080);
    RayCamera.SetParameters(camera, canvas);
-    RayCamera.CreateRays(Rays, Coords);
+    RayCamera.CreateRays(Rays, Coords.GetBounds());

    Rays.Buffers.at(0).InitConst(0.f);

    vtkm::cont::Field field = Data.GetField("pointvar");
    vtkm::Range range = field.GetRange().GetPortalConstControl().Get(0);

-    Tracer.SetData(Coords.GetData(), Indices, field, NumberOfTriangles, range, bounds);
+    Tracer.SetField(field, range);

    vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::UInt8, 4>> temp;
    vtkm::cont::ColorTable table("cool to warm");
@ -107,7 +116,7 @@ struct BenchRayTracing
  {
    vtkm::cont::Timer<VTKM_DEFAULT_DEVICE_ADAPTER_TAG> timer;

-    RayCamera.CreateRays(Rays, Coords);
+    RayCamera.CreateRays(Rays, Coords.GetBounds());
    Tracer.Render(Rays);

    return timer.GetElapsedTime();
--- a/benchmarking/Benchmarker.h
+++ b/benchmarking/Benchmarker.h
@ -326,7 +326,7 @@ public:
    catch (std::exception& e)
    {
      std::cout << "\n"
-                << "An exception occuring during a benchmark:\n\t" << e.what() << "\n"
+                << "An exception occurring during a benchmark:\n\t" << e.what() << "\n"
                << "Attempting to continue with remaining benchmarks...\n\n";
    }
  }
--- a/docs/changelog/0-sample-topic.md
+++ b/docs/changelog/0-sample-topic.md
@ -0,0 +1,7 @@
+# 0-sample-topic
+
+This is a sample release note for the change in a topic.
+Developers should add similar notes for each topic branch
+making a noteworthy change.  Each document should be named
+and titled to match the topic name to avoid merge conflicts.
+
--- a/docs/changelog/1.3/release-notes.md
+++ b/docs/changelog/1.3/release-notes.md
--- a/docs/changelog/add-float-version-functions-pi.md
+++ b/docs/changelog/add-float-version-functions-pi.md
@ -1,5 +0,0 @@
-# Add float version operations for vtkm::Math Pi()
-
-Now PI related functions are evalulated at compile time as constexpr functions.
-It also removes the old static_cast<T>vtkm::Pi() usages with
-template ones and fix several conversion warnings.
--- a/docs/changelog/add-unload-execution-resources-to-cellsets.md
+++ b/docs/changelog/add-unload-execution-resources-to-cellsets.md
@ -1,5 +0,0 @@
-# Add a release resources API to CellSet and its derived classes
-
-We now offer the ability to unload execution memory from CellSet and its derived
-classes(CellSetExplicit, CellSetPermutation and CellSetStructured) using the ReleaseResourcesExecution.
-
--- a/docs/changelog/add_tryexecuteondevice_function.md
+++ b/docs/changelog/add_tryexecuteondevice_function.md
@ -1,31 +0,0 @@
-# TryExecuteOnDevice allows for runtime selection of which device to execute on
-
-VTK-m now offers `vtkm::cont::TryExecuteOnDevice` to allow for the user to select 
-which device to execute a function on at runtime. The original `vtkm::cont::TryExecute`
-used the first valid device, which meant users had to modify the runtime state
-through the `RuntimeTracker` which was verbose and unwieldy. 
-
-Here is an example of how you can execute a function on the device that an array handle was last executed
-on:
-```cpp
-
-struct ArrayCopyFunctor
-{
-  template <typename Device, typename InArray, typename OutArray>
-  VTKM_CONT bool operator()(Device, const InArray& src, OutArray& dest)
-  {
-    vtkm::cont::DeviceAdapterAlgorithm<Device>::Copy(src, dest);    
-    return true;
-  }
-};
-
-template<typename T, typename InStorage, typename OutStorage>
-void SmartCopy(const vtkm::cont::ArrayHandle<T, InStorage>& src, vtkm::cont::ArrayHandle<T, OutStorage>& dest)
-{
-  bool success = vtkm::cont::TryExecuteOnDevice(devId, ArrayCopyFunctor(), src, dest);
-  if (!success)
-  {
-    vtkm::cont::TryExecute(ArrayCopyFunctor(), src, dest);
-  }
-}
-```
--- a/docs/changelog/agorithms_support_runtime_device.md
+++ b/docs/changelog/agorithms_support_runtime_device.md
@ -1,26 +0,0 @@
-# vtkm::cont::Algorithm now can be told which device to use at runtime
-
-The `vtkm::cont::Algorithm` has been extended to support the user specifying
-which device to use at runtime previously Algorithm would only use the first
-enabled device, requiring users to modify the `vtkm::cont::GlobalRuntimeDeviceTracker`
-if they wanted a specific device used.
-
-To select a specific device with vtkm::cont::Algorithm pass the `vtkm::cont::DeviceAdapterId`
-as the first parameter.
-
-```cpp
-vtkm::cont::ArrayHandle<double> values;
-
-
-//call with no tag, will run on first enabled device
-auto result = vtkm::cont::Algorithm::Reduce(values, 0.0);
-
-//call with an explicit device tag, will only run on serial
-vtkm::cont::DeviceAdapterTagSerial serial;
-result = vtkm::cont::Algorithm::Reduce(serial, values, 0.0);
-
-//call with an runtime device tag, will only run on serial
-vtkm::cont::DeviceAdapterId device = serial;
-result = vtkm::cont::Algorithm::Reduce(device, values, 0.0);
-
-```
--- a/docs/changelog/allow-CoordinateSystem-to-unload-execution-resources.md
+++ b/docs/changelog/allow-CoordinateSystem-to-unload-execution-resources.md
@ -1,10 +0,0 @@
-# Add a common API for CoordinateSystem to unload execution resources
-
-We now offer the ability to unload execution memory from ArrayHandleVirtualCoordinate
-and CoordinateSystem using the ReleaseResourcesExecution method.
-
-Field now has a ReleaseResourcesExecution.
-
-This commit also fixes a bug that ArrayTransfer of ArrayHandleVirtualCoordinate
-does not release execution resources properly.
-
--- a/docs/changelog/allow-custom-range-in-histogram-filter.md
+++ b/docs/changelog/allow-custom-range-in-histogram-filter.md
@ -1,4 +0,0 @@
-#Allow histogram filter to take custom types
-
-By passing TypeList and StorageList type into FieldRangeGlobalCompute,
-upstream users(VTK) can pass custom types into the histogram filter.
--- a/docs/changelog/allow-disabling-enabling-cuda-managed-memory.md
+++ b/docs/changelog/allow-disabling-enabling-cuda-managed-memory.md
@ -1,6 +0,0 @@
-Allow disabling/enabling of CUDA managed memory through a environment variable
-
-By setting the environment variable "VTKM_MANAGEDMEMO_DISABLED" to be 1,
-users are able to disable CUDA managed memory even though the hardware is capable
-of doing so.
-
--- a/docs/changelog/array-handle-memory-ownership.md
+++ b/docs/changelog/array-handle-memory-ownership.md
@ -1,16 +0,0 @@
-# VTK-m ArrayHandle can now take ownership of a user allocated memory location
-
-Previously memory that was allocated outside of VTK-m was impossible to transfer to
-VTK-m as we didn't know how to free it. By extending the ArrayHandle constructors
-to support a Storage object that is being moved, we can clearly express that
-the ArrayHandle now owns memory it didn't allocate.
-
-Here is an example of how this is done:
-```cpp
-  T* buffer = new T[100];
-  auto user_free_function = [](void* ptr) { delete[] static_cast<T*>(ptr); };
-
-  vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>
-      storage(buffer, 100, user_free_function);
-  vtkm::cont::ArrayHandle<T> arrayHandle(std::move(storage));
-```
--- a/docs/changelog/array-handle-view.md
+++ b/docs/changelog/array-handle-view.md
@ -1,23 +0,0 @@
-# Add `ArrayHandleView` fancy array
-
-Added a new class named `ArrayHandleView` that allows you to get a subset
-of an array. You use the `ArrayHandleView` by giving it a target array, a
-starting index, and a length. Here is a simple example of usage:
-
-``` cpp
-vtkm::cont::ArrayHandle<vtkm::Id> sourceArray;
-
-vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleIndex(10), sourceArray);
-// sourceArray has [0, 1, 2, 3, 4, 5, 6, 7, 8, 9]
-
-vtkm::cont::ArrayHandleView<vtkm::cont::ArrayHandle<vtkm::Id>>
-  viewArray(sourceArray, 3, 5);
-// viewArray has [3, 4, 5, 6, 7]
-```
-
-There is also a convenience `make_ArraHandleView` function to create view
-arrays. The following makes the same view array as before.
-
-``` cpp
-auto viewArray = vtkm::cont::make_ArrayHandleView(sourceArray, 3, 5);
-```
--- a/docs/changelog/arrayhandlecompositevector-refactor.md
+++ b/docs/changelog/arrayhandlecompositevector-refactor.md
@ -1,20 +0,0 @@
-# `ArrayHandleCompositeVector` simplified and made writable.
-
-`ArrayHandleCompositeVector` is now easier to use, as its type has a more
-straightforward definition: `ArrayHandleCompositeVector<Array1, Array2, ...>`.
-Previously, a helper metaprogramming struct was needed to determine the type
-of the array handle.
-
-In addition, the new implementation supports both reading and writing, whereas
-the original version was read-only.
-
-Another notable change is that the `ArrayHandleCompositeVector` no longer
-supports component extraction from the source arrays. While the previous version
-could take a source array with a `vtkm::Vec` `ValueType` and use only a single
-component in the output, the new version requires that all input arrays have
-the same `ValueType`, which becomes the `ComponentType` of the output
-`vtkm::Vec`.
-
-When component extraction is needed, the classes `ArrayHandleSwizzle` and
-`ArrayHandleExtractComponent` have been introduced to allow the previous
-usecases to continue working efficiently.
--- a/docs/changelog/arrayhandleextractcomponent-runtime-component.md
+++ b/docs/changelog/arrayhandleextractcomponent-runtime-component.md
@ -1,7 +0,0 @@
-# `ArrayHandleExtractComponent` target component is now set at runtime.
-
-Rather than embedding the extracted component in a template parameter, the
-extract operation is now defined at runtime.
-
-This is easier to use and keeps compile times / sizes / memory requirements
-down.
--- a/docs/changelog/arrayhandleswizzle-runtime-map.md
+++ b/docs/changelog/arrayhandleswizzle-runtime-map.md
@ -1,8 +0,0 @@
-# `ArrayHandleSwizzle` component maps are now set at runtime.
-
-Rather than embedding the component map in the template parameters, the swizzle
-operation is now defined at runtime using a `vtkm::Vec<vtkm::IdComponent, N>`
-that maps the input components to the output components.
-
-This is easier to use and keeps compile times / sizes / memory requirements
-down.
--- a/docs/changelog/build-system-redesign.md
+++ b/docs/changelog/build-system-redesign.md
@ -1,61 +0,0 @@
-# Build System Redesign and new minimum CMake
-
-VTK-m CMake buildsystem was redesigned to be more declarative for consumers.
-This was done by moving away from the previous component design and instead
-to explicit targets. Additionally VTK-m now uses the native CUDA support
-introduced in CMake 3.8 and has the following minimum CMake versions:
- - Visual Studio Generator requires CMake 3.11+
- - CUDA support requires CMake 3.9+
- - Otherwise CMake 3.3+ is supported
-
-When VTK-m is found find_package it defines the following targets:
-  - `vtkm_cont`
-    - contains all common core functionality
-    - always exists
-
-  - `vtkm_rendering`  
-    - contains all the rendering code
-    - exists only when rendering is enabled
-    - rendering also provides a `vtkm_find_gl` function
-      - allows you to find the GL (EGL,MESA,Hardware), GLUT, and GLEW
-        versions that VTK-m was built with.
-
-VTK-m also provides targets that represent what device adapters it
-was built to support. The pattern for these targets are `vtkm::<device>`.
-Currently we don't provide a target for the serial device.
-
-  - `vtkm::tbb` 
-    - Target that contains tbb related link information
-       implicitly linked to by `vtkm_cont` if tbb was enabled
-
-  - `vtkm::cuda` 
-    - Target that contains cuda related link information
-       implicitly linked to by `vtkm_cont` if cuda was enabled       
-
-VTK-m can be built with specific CPU architecture vectorization/optimization flags.
-Consumers of the project can find these flags by looking at the `vtkm_vectorization_flags`
-target.
-
-So a project that wants to build an executable that uses vtk-m would look like:
-
-```cmake
-
-cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
-project(HellowWorld CXX)
-
-#Find the VTK-m package. 
-#Will automatically enable the CUDA language if needed ( and bump CMake minimum )
-
-find_package(VTKm REQUIRED)
-
-add_executable(HelloWorld HelloWorld.cxx)
-target_link_libraries(HelloWorld PRIVATE vtkm_cont)
-
-if(TARGET vtkm::cuda)
-  set_source_files_properties(HelloWorld.cxx PROPERTIES LANGUAGE CUDA)
-endif()
-
-```
-
-
-
--- a/docs/changelog/cell-measure.md
+++ b/docs/changelog/cell-measure.md
@ -1,54 +0,0 @@
-# Cell measure functions, worklet, and filter
-
-VTK-m now provides free functions, a worklet, and a filter for computing
-the integral measure of a cell (i.e., its arc length, area, or volume).
-
-The free functions are located in `vtkm/exec/CellMeasure.h` and share the
-same signature:
-
-```c++
-  template<typename OutType, typename PointVecType>
-  OutType CellMeasure(
-    const vtkm::IdComponent& numPts,
-    const PointCoordVecType& pts,
-    CellShapeTag,
-    const vtkm::exec::FunctorBase& worklet);
-```
-
-The number of points argument is provided for cell-types such as lines,
-which allow an arbitrary number of points per cell.
-See the worklet for examples of their use.
-
-The worklet is named `vtkm::worklet::CellMeasure` and takes a template
-parameter that is a tag list of measures to include.
-Cells that are not selected by the tag list return a measure of 0.
-Some convenient tag lists are predefined for you:
-
-+ `vtkm::ArcLength` will only compute the measure of cells with a 1-dimensional parameter-space.
-+ `vtkm::Area` will only compute the measure of cells with a 2-dimensional parameter-space.
-+ `vtkm::Volume` will only compute the measure of cells with a 3-dimensional parameter-space.
-+ `vtkm::AllMeasures` will compute all of the above.
-
-The filter version, named `vtkm::filter::CellMeasures` – plural since
-it produces a cell-centered array of measures — takes the same template
-parameter and tag lists as the worklet.
-By default, the output array of measure values is named "measure" but
-the filter accepts other names via the `SetCellMeasureName()` method.
-
-The only cell type that is not supported is the polygon;
-you must triangulate polygons before running this filter.
-See the unit tests for examples of how to use the worklet and filter.
-
-The cell measures are all signed: negative measures indicate that the cell is inverted.
-Simplicial cells (points, lines, triangles, tetrahedra) cannot not be inverted
-by definition and thus always return values above or equal to 0.0.
-Negative values indicate either the order in which vertices appear in its connectivity
-array is improper or the relative locations of the vertices in world coordinates
-result in a cell with a negative Jacobian somewhere in its interior.
-Finally, note that cell measures may return invalid (NaN) or infinite (Inf, -Inf)
-values if the cell is poorly defined, e.g., has coincident vertices
-or a parametric dimension larger than the space spanned by its world-coordinate
-vertices.
-
-The verdict mesh quality library was used as the source of the methods
-for approximating the cell measures.
--- a/docs/changelog/cellsetexplicit-cache-celltopoint.md
+++ b/docs/changelog/cellsetexplicit-cache-celltopoint.md
@ -1,21 +0,0 @@
-# CellSetExplicit now caches CellToPoint table when used with Invoke.
-
-Issue #268 highlighted an issue where the expensive CellToPoint table
-update was not properly cached when a CellSetExplicit was used with a
-filter. This has been corrected by ensuring that the metadata
-associated with the table survives shallow copying of the CellSet.
-
-New methods are also added to check whether the CellToPoint table
-exists, and also to reset it if needed (e.g. for benchmarking):
-
-```
-vtkm::cont::CellSetExplicit<> cellSet = ...;
-// Check if the CellToPoint table has already been computed:
-if (cellSet.HasConnectivity(vtkm::TopologyElementTagCell{},
-                            vtkm::TopologyElementTagPoint{}))
-{
-  // Reset it:
-  cellSet.ResetConnectivity(vtkm::TopologyElementTagCell{},
-                            vtkm::TopologyElementTagPoint{});
-}
-```
--- a/docs/changelog/change-execution-object-creation.md
+++ b/docs/changelog/change-execution-object-creation.md
@ -1,27 +0,0 @@
-# User defined execution objects now usable with runtime selection of device adapter
-
- Changed how Execution objects are created and passed from the cont environment to the execution environment. See chapter 13.9 on worklets in the user manual for details.
-
- Instead we will now fill out a class and call prepareForExecution() and create the execution object for the execution environment from this function. This way we do not have to template the class that extends `vtkm::cont::ExecutionObjectBase` on the device.
-
-Example of new execution object:
-```cpp
-template <typename Device>
-struct ExecutionObject
-{
-  vtkm::Int32 Number;
-};
-
-struct TestExecutionObject : public vtkm::cont::ExecutionObjectBase
-{
-  vtkm::Int32 Number;
-
-  template <typename Device>
-  VTKM_CONT ExecutionObject<Device> PrepareForExecution(Device) const
-  {
-    ExecutionObject<Device> object;
-    object.Number = this->Number;
-    return object;
-  }
-};
-```
--- a/docs/changelog/convert-enums-to-enum-structs.md
+++ b/docs/changelog/convert-enums-to-enum-structs.md
@ -1,4 +0,0 @@
-# Use the strong typed enums for vtkm::cont::Field
-
-By doing so, the compiler would not convert these enums into `int`s
-which can cause some unexpected behavior.
--- a/docs/changelog/cuda-arch-none.md
+++ b/docs/changelog/cuda-arch-none.md
@ -1,8 +0,0 @@
-# Add new option to VTKm_CUDA_Architecture
-
-A new VTKm_CUDA_Architecture option called 'none' has been added. This will
-disable all VTK-m generated cuda architecture flags, allowing the user to
-specify their own custom flags.
-
-Useful when VTK-m is used as a library in another project and the project wants
-to use its own architecture flags.
--- a/docs/changelog/cuda-asynchronous-errors.md
+++ b/docs/changelog/cuda-asynchronous-errors.md
@ -1,15 +0,0 @@
-# Worklets are now asynchronous in Cuda
-
-Worklets are now fully asynchronous in the cuda backend. This means that
-worklet errors are reported asynchonously. Existing errors are checked for
-before invocation of a new worklet and at explicit synchronization points like
-`DeviceAdapterAlgorithm<>::Synchronize()`.
-
-An important effect of this change is that functions that are synchronization
-points, like `ArrayHandle::GetPortalControl()` and
-`ArrayHandle::GetPortalConstControl()`, may now throw exception for errors from
-previously executed worklets.
-
-Worklet invocations, synchronization and error reporting happen independtly
-on different threads. Therefore, synchronization on one thread does not affect
-any other threads.
--- a/docs/changelog/cuda-deferred-free.md
+++ b/docs/changelog/cuda-deferred-free.md
@ -1,8 +0,0 @@
-# Add support for deferred freeing of cuda memory
-
-A new function, `void CudaAllocator::FreeDeferred(void* ptr, std::size_t numBytes)` has
-been added that can be used to defer the freeing of cuda memory to a later point.
-This is useful because `cudaFree` causes a global sync across all cuda streams. This function
-internally maintains a pool of to-be-freed pointers that are freed together when a
-size threshold is reached. This way a number of global syncs are collected together at
-one point.
--- a/docs/changelog/cuda-gride-stride-loops.md
+++ b/docs/changelog/cuda-gride-stride-loops.md
@ -1,49 +0,0 @@
-# VTK-m Worklets now execute on Cuda using grid stride loops
-
-Previously VTK-m Worklets used what is referred to as a monolithic kernel
-pattern for worklet execution. This assumes a single large grid of threads
-to process an entire array in a single pass. This resulted in launches that
-looked like:
-
-```cpp
-template<typename F>
-void TaskSingular(F f, vtkm::Id end)
-{
-  const vtkm::Id index = static_cast<vtkm::Id>(blockDim.x * blockIdx.x + threadIdx.x);
-  if (index < end)
-  {
-    f(index);
-  }  
-}
-
-Schedule1DIndexKernel<TaskSingular><<<totalBlocks, 128, 0, cudaStreamPerThread>>>(
-       functor, numInstances);
-```
-
-This was problematic as it had the drawbacks of:
- Not being able to reuse any infrastructure between kernel executions.
- Harder to tune performance based on the current hardware.
-
-The solution was to move to a grid stride loop strategy with a block size 
-based off the number of SM's on the executing GPU. The result is something
-that looks like:
-
-```cpp
-template<typename F>
-void TaskStrided(F f, vtkm::Id end)
-{
-  const vtkm::Id start = blockIdx.x * blockDim.x + threadIdx.x;
-  const vtkm::Id inc = blockDim.x * gridDim.x;
-  for (vtkm::Id index = start; index < end; index += inc)
-  {
-    f(index);
-  }  
-}
-Schedule1DIndexKernel<TaskStrided><<<32*numSMs, 128, 0, cudaStreamPerThread>>>(
-       functor, numInstances);
-```
-
- With a loop stride equal to grid size we maintain the optimal memory
- coalescing patterns as we had with the monolithic version. These changes
- also allow VTK-m to optimize TaskStrided so that it can reuse infrastructure
- between iterations.
--- a/docs/changelog/deviceadapterid_is_now_a_strong_type.md
+++ b/docs/changelog/deviceadapterid_is_now_a_strong_type.md
@ -1,41 +0,0 @@
-#DeviceAdapterId has becomes a real constexpr type and not an alias to vtkm::UInt8
-
-As part of the ability to support `vtkm::cont::TryExecuteOnDevice` VTK-m has made the
-DeviceAdapterId a real constexpr type instead of a vtkm::UInt8.
-
-The benefits of a real type are as follows:
-
- Easier to add functionality like range verification, which previously had
-  to be located in each user of `DeviceAdapterId`
-
- In ability to have ambiguous arguments. Previously it wasn't perfectly clear
-  what a method parameter of `vtkm::UInt8` represented. Was it actually the
-  DeviceAdapterId or something else?
-
- Ability to add subclasses that represent things such as Undefined, Error, or Any.
-
-
-The implementation of DeviceAdapterId is:
-```cpp
-struct DeviceAdapterId
-{
-  constexpr explicit DeviceAdapterId(vtkm::Int8 id)
-    : Value(id)
-  {
-  }
-
-  constexpr bool operator==(DeviceAdapterId other) const { return this->Value == other.Value; }
-  constexpr bool operator!=(DeviceAdapterId other) const { return this->Value != other.Value; }
-  constexpr bool operator<(DeviceAdapterId other) const { return this->Value < other.Value; }
-
-  constexpr bool IsValueValid() const
-  {
-    return this->Value > 0 && this->Value < VTKM_MAX_DEVICE_ADAPTER_ID;
-  }
-
-  constexpr vtkm::Int8 GetValue() const { return this->Value; }
-
-private:
-  vtkm::Int8 Value;
-};
-```
--- a/docs/changelog/deviceadaptertags_usable_for_runtime_device_selection.md
+++ b/docs/changelog/deviceadaptertags_usable_for_runtime_device_selection.md
@ -1,45 +0,0 @@
-# DeviceAdapterTags are usable for runtime device selection
-
-VTK-m DeviceAdapterTags now are both a compile time representation of which device to use, and
-also the runtime representation of that device. Previously the runtime representation was handled
-by `vtkm::cont::DeviceAdapterId`. This was done by making `DeviceAdapterTag`'s' a constexpr type that
-inherits from the constexpr `vtkm::cont::DeviceAdapterId` type.
-
-At at ten thousand foot level this change means that in general instead of using `vtkm::cont::DeviceAdapterTraits<DeviceTag>`
-you can simply use `DeviceTag`, or an instance of if `DeviceTag runtimeDeviceId;`.
-
-Previously if you wanted to get the runtime representation of a device you would do the following:
-```cpp
-template<typename DeviceTag>
-vtkm::cont::DeviceAdapterId getDeviceId()
-{
-  using Traits = vtkm::cont::DeviceAdapterTraits<DeviceTag>;
-  return Traits::GetId();
-}
-...
-vtkm::cont::DeviceAdapterId runtimeId = getDeviceId<DeviceTag>();
-```
-Now with the updates you could do the following.
-```cpp
-vtkm::cont::DeviceAdapterId runtimeId = DeviceTag();
-```
-More importantly this conversion is unnecessary as you can pass instances `DeviceAdapterTags` into methods or functions
-that want `vtkm::cont::DeviceAdapterId` as they are that type!
-
-
-Previously if you wanted to see if a DeviceAdapter was enabled you would the following:
-```cpp
-using Traits = vtkm::cont::DeviceAdapterTraits<DeviceTag>;
-constexpr auto isValid = std::integral_constant<bool, Traits::Valid>();
-```
-Now you would do:
-```cpp
-constexpr auto isValid = std::integral_constant<bool, DeviceTag::IsEnabled>();
-```
-
-So why did VTK-m make these changes? 
-
-That is a good question, and the answer for that is two fold. The VTK-m project is working better support for ArraysHandles that leverage runtime polymorphism (aka virtuals), and the ability to construct `vtkm::worklet::Dispatchers` without specifying 
-the explicit device they should run on. Both of these designs push more of the VTK-m logic to operate at runtime rather than compile time. This changes are designed to allow for consistent object usage between runtime and compile time instead of having
-to convert between compile time and runtime types.
-
--- a/docs/changelog/dispatcher-auto-device.md
+++ b/docs/changelog/dispatcher-auto-device.md
@ -1,15 +0,0 @@
-# Make DispatcherBase invoke using a TryExecute
-
-Rather than force all dispatchers to be templated on a device adapter,
-instead use a TryExecute internally within the invoke to select a device
-adapter.
-    
-Because this removes the need to declare a device when invoking a worklet,
-this commit also removes the need to declare a device in several other
-areas of the code.
-
-This changes touches quite a bit a code. The first pass of the change
-usually does the minimum amount of work, which is to change the
-compile-time specification of the device to a run-time call to `SetDevice`
-on the dispatcher. Although functionally equivalent, it might mean calling
-`TryExecute` within itself.
--- a/docs/changelog/dispatcher-invoke-supports-pointers.md
+++ b/docs/changelog/dispatcher-invoke-supports-pointers.md
@ -1,19 +0,0 @@
-# VTK-m now supports dispatcher parameters being pointers
-
-Previously it was only possible to pass values to a dispatcher when
-you wanted to invoke a VTK-m worklet. This caused problems when it came
-to designing new types that used inheritance as the types couldn't be
-past as the base type to the dispatcher. To fix this issue we now
-support invoking worklets with pointers as seen below.
-
-```cpp
-  vtkm::cont::ArrayHandle<T> input;
-  //fill input
-
-  vtkm::cont::ArrayHandle<T> output;
-  vtkm::worklet::DispatcherMapField<WorkletType> dispatcher;
-
-  dispatcher(&input, output);
-  dispatcher(input, &output); 
-  dispatcher(&input, &output); 
-```
--- a/docs/changelog/dot-product.md
+++ b/docs/changelog/dot-product.md
@ -1,6 +0,0 @@
-# Dot product function name changed.
-
-The free function `vtkm::dot()` has been renamed to `vtkm::Dot()`
-to be consistent with other vtk-m function names.
-Aliases are provided for backwards compatibility but will
-be removed in the next release.
--- a/docs/changelog/enable-developer-flags.md
+++ b/docs/changelog/enable-developer-flags.md
@ -1,11 +0,0 @@
-# Add a new cmake option: VTKm_ENABLE_DEVELOPER_FLAGS
-
-The new cmake option VTKm_ENABLE_DEVELOPER_FLAGS can be used to enable/disable
-warnings in VTK-m. It is useful to disable VTK-m's warning flags when VTK-m is
-directly embedded by a project as sub project (add_subdirectory), and the
-warnings are too strict for the project. This does not apply when using an
-installed version of VTK-m.
-
-For example, this flag is disabled in VTK.
-
-This flag is enabled by default.
--- a/docs/changelog/exec-objects-alg-args.md
+++ b/docs/changelog/exec-objects-alg-args.md
@ -1,24 +0,0 @@
-# Support ExecArg behavior in `vtkm::cont::Algorithm` methods
-
-`vtkm::cont::Algorithm` is a wrapper around `DeviceAdapterAlgorithm` that
-internally uses `TryExecute`s to select an appropriate device. The
-intention is that you can run parallel algorithms (outside of worklets)
-without having to specify a particular device.
-
-Most of the arguments given to device adapter algorithms are actually
-control-side arguments that get converted to execution objects internally
-(usually a `vtkm::cont::ArrayHandle`). However, some of the algorithms,
-take an argument that is passed directly to the execution environment, such
-as the predicate argument of `Sort`. If the argument is a plain-old-data
-(POD) type, which is common enough, then you can just pass the object
-straight through. However, if the object has any special elements that have
-to be transferred to the execution environment, such as internal arrays,
-passing this to the `vtkm::cont::Algorithm` functions becomes problematic.
-
-To cover this use case, all the `vtkm::cont::Algorithm` functions now
-support automatically transferring objects that support the `ExecObject`
-worklet convention. If any argument to any of the `vtkm::cont::Algorithm`
-functions inherits from `vtkm::cont::ExecutionObjectBase`, then the
-`PrepareForExecution` method is called with the device the algorithm is
-running on, which allows these device-specific objects to be used without
-the hassle of creating a `TryExecute`.
--- a/docs/changelog/geometry.md
+++ b/docs/changelog/geometry.md
@ -1,75 +0,0 @@
-# New geometry classes and header.
-
-There are now some additional structures available both
-the control and execution environments for representing
-geometric entities (mostly of dimensions 2 and 3).
-These new structures are now in `vtkm/Geometry.h` and
-demonstrated/tested in `vtkm/testing/TestingGeometry.h`:
-
-+ `Ray<CoordType, Dimension, IsTwoSided>`.
-  Instances of this struct represent a semi-infinite line
-  segment in a 2-D plane or in a 3-D space, depending on the
-  integer dimension specified as a template parameter.
-  Its state is the point at the start of the ray (`Origin`)
-  plus the ray's `Direction`, a unit-length vector.
-  If the third template parameter (IsTwoSided) is true, then
-  the ray serves as an infinite line. Otherwise, the ray will
-  only report intersections in its positive halfspace.
-+ `LineSegment<CoordType, Dimension>`.
-  Instances of this struct represent a finite line segment
-  in a 2-D plane or in a 3-D space, depending on the integer
-  dimension specified as a template parameter.
-  Its state is the coordinates of its `Endpoints`.
-+ `Plane<CoordType>`.
-  Instances of this struct represent a plane in 3-D.
-  Its state is the coordinates of a base point (`Origin`) and
-  a unit-length normal vector (`Normal`).
-+ `Sphere<CoordType, Dimension>`.
-  Instances of this struct represent a *d*-dimensional sphere.
-  Its state is the coordinates of its center plus a radius.
-  It is also aliased with a `using` statment to `Circle<CoordType>`
-  for the specific case of 2-D.
-
-These structures provide useful queries and generally
-interact with one another.
-For instance, it is possible to intersect lines and planes
-and compute distances.
-
-For ease of use, there are also several `using` statements
-that alias these geometric structures to names that specialize
-them for a particular dimension or other template parameter.
-As an example, `Ray<CoordType, Dimension, true>` is aliased
-to `Line<CoordType, Dimension>` and `Ray<CoordType, 3, true>`
-is aliased to `Line3<CoordType>` and `Ray<FloatDefault, 3, true>`
-is aliased to `Line3d`.
-
-## Design patterns
-
-If you plan to add a new geometric entity type,
-please adopt these conventions:
-
-+ Each geometric entity may be default-constructed.
-  The default constructor will initialize the state to some
-  valid unit-length entity, usually with some part of
-  its state at the origin of the coordinate system.
-+ Entities may always be constructed by passing in values
-  for their internal state.
-  Alternate construction methods are declared as free functions
-  such as `make_CircleFrom3Points()`
-+ Use template metaprogramming to make methods available
-  only when the template dimension gives them semantic meaning.
-  For example, a 2-D line segment's perpendicular bisector
-  is another line segment, but a 3-D line segment's perpendicular
-  line segment is a plane.
-  Note how this is accomplished and apply this pattern to
-  new geometric entities or new methods on existing entities.
-+ Some entities may have invalid state.
-  If this is possible, the entity will have an `IsValid()` method.
-  For example, a sphere may be invalid because the user or some
-  construction technique specified a zero or negative radius.
-+ When signed distance is semantically meaningful, provide it
-  in favor of or in addition to unsigned distance.
-+ Accept a tolerance parameter when appropriate,
-  but provide a sensible default value.
-  You may want to perform exact arithmetic versions of tests,
-  but please provide fast, tolerance-based versions as well.
--- a/docs/changelog/install-only-libraries.md
+++ b/docs/changelog/install-only-libraries.md
@ -1,9 +0,0 @@
-# Add a new cmake option: VTKm_INSTALL_ONLY_LIBRARIES
-
-The new cmake option VTKm_INSTALL_ONLY_LIBRARIES when enabled will cause
-VTK-m to only install libraries. This is useful for projects that are
-producing an application and don't want to ship headers or CMake infrastructure.
-
-For example, this flag is enabled by ParaView for releases.
-
-This flag is disabled by default.
--- a/docs/changelog/lagrangian-filter.md
+++ b/docs/changelog/lagrangian-filter.md
@ -1,20 +0,0 @@
-# Filter to support Lagrangian analysis capabilities.  
-
-Lagrangian analysis operates in two phases - phase one involes the extraction of flow field information. Phase two involves calculating new particle trajectories using the saved information. 
-
-The lagrangian filter can be used to extract flow field information given a time-varying vector fields. The extracted information is in the form of particle trajectories.
-
-The filter operates by first being set up with some information regarding step size, the interval at which information should be saved (write frequency), the number of seeds to be placed in the domain (specified as a reduction factor along each axis of the original dimensions).
-The step size should be equivalent to the time between vector field data input. 
-The write frequency corresponds to the number of cycles between saves. 
-
-Filter execution is called for each cycle of the simulation data. Each filter execution call requires a velocity field to advect particles forward. 
-
-The extracted particle trajectories - referred to as basis flows exist in the domain for the specified interval (write frequency). Particles are then reset along a uniform grid and new particle trajectories are calculated. 
-
-An example of using the Lagrangian filter is at vtk-m/examples/lagrangian
-The basis flows are saved into a folder named output which needs to be created in the directory in which the program is being executed.
-
-The basis flows can be interpolated using barycentric coordinate interpolation or a form of linear interpolation to calculate new particle trajectories post hoc. 
-
-An example of using basis flows generated by the Lagrangian filter is at vtk-m/examples/posthocinterpolation. The folder contains a script which specifies parameters which need to be provided to use the example.
--- a/docs/changelog/merge-worklet-testing-executables.md
+++ b/docs/changelog/merge-worklet-testing-executables.md
@ -0,0 +1,11 @@
+Merge worklet testing executables into a device dependent shared library
+
+VTK-m has been updated to replace old per device worklet testing executables with a device
+dependent shared library so that it's able to accept a device adapter
+at runtime.
+Meanwhile, it updates the testing infrastructure APIs. vtkm::cont::testing::Run
+function would call ForceDevice when needed and if users need the device
+adapter info at runtime, RunOnDevice function would pass the adapter into the functor.
+
+Optional Parser is bumped from 1.3 to 1.7.
+
--- a/docs/changelog/openmp-backend.md
+++ b/docs/changelog/openmp-backend.md
@ -1,5 +0,0 @@
-# OpenMP Device Adapter
-
-A device adapter that leverages OpenMP 4.0 for threading is now available. The
-new adapter is enabled using the CMake option `VTKm_ENABLE_OPENMP` and its
-performance is comparable to the TBB device adapter.
--- a/docs/changelog/oscillator-source.md
+++ b/docs/changelog/oscillator-source.md
@ -1,11 +0,0 @@
-# Time-varying "oscillator" source-filter and example
-
-The oscillator is a simple analytical source of time-varying data.
-It provides a function value at each point of a uniform grid that
-is computed as a sum of Gaussian kernels — each with a specified
-position, amplitude, frequency, and phase.
-
-The example (in `examples/oscillator`) generates volumetric Cinema
-datasets that can be viewed in a web browser with [ArcticViewer][].
-
-[ArcticViewer]: https://kitware.github.io/arctic-viewer/
--- a/docs/changelog/replace-random_shuffle-shuffle.md
+++ b/docs/changelog/replace-random_shuffle-shuffle.md
@ -1,5 +0,0 @@
-Replace std::random_shuffle with std::shuffle
-
-std::random_shuffle is deprecated in C++14 because it's using std::rand
-which uses a non uniform distribution and the underlying algorithm is
-unspecified. Using std::shuffle can provide a reliable result in a 64 bit version.
--- a/docs/changelog/scatter-to-dispatcher.md
+++ b/docs/changelog/scatter-to-dispatcher.md
@ -1,73 +0,0 @@
-# Scatter class moved to dispatcher
-
-Scatter classes are special objects that are associated with a worklet to
-adjust the standard 1:1 mapping of input to output in the worklet execution
-to some other mapping with multiple outputs to a single input or skipping
-over input values. A classic use case is the Marching Cubes algorithm where
-cube cases will have different numbers of output. A scatter object allows
-you to specify for each output polygon which source cube it comes from.
-
-Scatter objects have been in VTK-m for some time now (since before the 1.0
-release). The way they used to work is that the worklet completely managed
-the scatter object. It would declare the `ScatterType`, keep a copy as part
-of its state, and provide a `GetScatter` method so that the dispatcher
-could use it for scheduling.
-
-The problem with this approach is that it put control-environment-specific
-state into the worklet. The scatter object would be pushed into the
-execution environment (like a CUDA device) like the rest of the worklet
-where it could not be used. It also meant that worklets that defined their
-own scatter had to declare a bunch more code to manage the scatter.
-
-This behavior has been changed so that the dispatcher object manages the
-scatter object. The worklet still declares the type of scatter by declaring
-a `ScatterType` (defaulting to `ScatterUniform` for 1:1 mapping),
-but its responsibility ends there. When the dispatcher is constructed, it
-must be given a scatter object that matches the `ScatterType` of the
-associated worklet. (If `ScatterType` has a default constructor, then one
-can be created automatically.) A worklet may declare a static `MakeScatter`
-method for convenience, but this is not necessary.
-
-As an example, a worklet may declare a custom scatter like this.
-
-``` cpp
-  class Generate : public vtkm::worklet::WorkletMapField
-  {
-  public:
-    typedef void ControlSignature(FieldIn<Vec3> inPoints,
-                                  FieldOut<Vec3> outPoints);
-    typedef void ExecutionSignature(_1, _2);
-    using InputDomain = _1;
-
-    using ScatterType = vtkm::worklet::ScatterCounting;
-
-    template<typename CountArrayType, typename DeviceAdapterTag>
-    VTKM_CONT
-    static ScatterType MakeScatter(const CountArrayType &countArray,
-                                   DeviceAdapterTag)
-    {
-      VTKM_IS_ARRAY_HANDLE(CountArrayType);
-      return ScatterType(countArray, DeviceAdapterTag());
-    }
-```
-
-Note that the `ScatterCounting` needs to be created with the appropriate
-indexing arrays to make the scatter behave as the worklet expects, so the
-worklet provides a helpful `MakeScatter` method to make it more clear how
-to construct the scatter.
-
-This worklet can be invoked as follows.
-
-``` cpp
-    auto generateScatter =
-        ClipPoints::Generate::MakeScatter(countArray, DeviceAdapterTag());
-    vtkm::worklet::DispatcherMapField<ClipPoints::Generate, DeviceAdapterTag>
-        dispatcherGenerate(generateScatter);
-    dispatcherGenerate.Invoke(pointArray, clippedPointsArray);
-```
-
-Because the `ScatterCounting` class does not have a default constructor,
-you would get a compiler error if you failed to provide one to the
-dispatcher's constructor. The compiler error will probably not be too
-helpful the the user, but there is a detailed comment in the dispatcher's
-code where the compiler error will occur describing what the issue is.
--- a/docs/changelog/smaller-mem-footprint-cellderivative.md
+++ b/docs/changelog/smaller-mem-footprint-cellderivative.md
@ -1,5 +0,0 @@
-#  CellDerivativeFor3DCell has a better version for Vec of Vec fields.
-    
-Previously we would compute a 3x3 matrix where each element was a Vec. Using
-the jacobain of a single component is sufficient instead of computing it for
-each component. This approach saves anywhere from 2 to 3 times the memory space.
--- a/docs/changelog/spatial-search-interfaces.md
+++ b/docs/changelog/spatial-search-interfaces.md
@ -1,54 +0,0 @@
-#Interfaces for VTK-m spatial search strucutres added
-
-The objective for this feature was to add a commom interface for the VTK-m
-spatial search strucutes for ease of use for the users.
-VTK-m now distinguishes locators into two types, cell locators and point 
-locators. Cell locators can be used to query a containing cell for a point, 
-and point locators can be used to search for other points that are close to the 
-given point.
-
-All cell locators are now required to inherit from the interface
-`vtkm::cont::CellLocator`,  and all point locatos are required to inherit from 
-the interface `vtkm::cont::PointLocator` 
-
-These interfaces describe the necessary features that are required from either 
-a cell locator, or a point locator and provided an easy way to use them in the
-execution environment.
-
-By deriving new search structures from these locator interfaces, it makes it
-easier for users to build the underlying strucutres as well, abstracting away
-complicated details. After providing all the required data from a 
-`vtkm::cont::DataSet` object, the user only need to call the `Update` method 
-on the object of `vtkm::cont::CellLocator`, or `vtkm::cont::PointLocator`.
-
-For example, building the cell locator which used a Bounding Interval Hiererchy 
-tree as a search structure, provided in the class 
-`vtkm::cont::BoundingIntervalHierarchy` which inherits from 
-`vtkm::cont::CellLocator`, only requires few steps.
-
-```c++
-  // Build a bounding interval hierarchy with 5 splitting planes,
-  // and a maximum of 10 cells in the leaf node.
-  vtkm::cont::BoundingIntervalHierarchy locator(5, 10);
-  // Provide the cell set required by the search structure.
-  locator.SetCellSet(cellSet);
-  // Provide the coordinate system required by the search structure.
-  locator.SetCoordinates(coords);
-  // Cell the Update methods to finish building the underlying tree.
-  locator.Update();
-```
-Similarly, users can easily build available point locators as well.
-
-When using an object of `vtkm::cont::CellLocator`, or `vtkm::cont::PointLocator`
-in the execution environment, they need to be passed to the worklet as an 
-`ExecObject` argument. In the execution environment, users will receive a 
-pointer to an object of type `vtkm::exec::CellLocator`, or 
-`vtkm::exec::PointLocator` respectively. `vtkm::exec::CellLocator` provides a 
-method `FindCell` to use in the execution environment to query the containing 
-cell of a point. `vtkm::exec::PointLocator` provides a method 
-`FindNearestNeighbor` to query for the nearest point.
-
-As of now,  VTK-m provides only one implementation for each of the given 
-interfaces. `vtkm::cont::BoundingIntervalHierarchy` which is an implementation 
-of `vtkm::cont::CellLocator`, and `vtkm::cont::PointLocatorUniformGrid`, which 
-is an implementation of `vtkm::cont::PointLocator`.
--- a/docs/changelog/storage-free-functions.md
+++ b/docs/changelog/storage-free-functions.md
@ -1,24 +0,0 @@
-# VTK-m StorageBasic now can provide or be provided a delete function
-
-Memory that was allocated outside of VTK-m was impossible to transfer to
-VTK-m as we didn't know how to free it. This is now resolved by allowing the
-user to specify a free function to be called on release.
-
-Memory that was allocated by VTK-m and Stolen by the user needed the
-proper free function. When running on CUDA on hardware that supports
-concurrent managed access the free function of the storage could
-be cudaFree. 
-
-To properly steal memory from VTK-m you do the following:
-```cpp
-  vtkm::cont::ArrayHandle<T> arrayHandle;
-  //fill arrayHandle
-
-  //you must get the free function before calling steal array
-  auto free_function = arrayHandle.GetDeleteFunction();
-  T* ptr = arrayHandle.StealArray();
-  //use ptr
-
-
-  free_function(ptr);
-```
--- a/docs/changelog/swap.md
+++ b/docs/changelog/swap.md
@ -1,7 +0,0 @@
-# Add a CUDA-safe `vtkm::Swap` method.
-
-Added a swap implementation that is safe to call from all backends.
-
-It is not legal to call std functions from CUDA code, and the new
-`vtkm::Swap` implements a naive swap when compiled under NVCC while
-falling back to a std/ADL swap otherwise.
--- a/docs/changelog/use-std-call-once-to-construct-singletons.md
+++ b/docs/changelog/use-std-call-once-to-construct-singletons.md
@ -1,4 +0,0 @@
-# Use std::call_once to construct singeltons
-
-By using call_once from C++11, we can simplify the logic in code where we are querying same value variables from multiple threads.
-
--- a/docs/changelog/use-thread_local-in-GetGlobalRuntimeDeviceTracker.md
+++ b/docs/changelog/use-thread_local-in-GetGlobalRuntimeDeviceTracker.md
@ -1,5 +0,0 @@
-# Use thread_local in GetGlobalRuntimeDeviceTracker function if possible
-
-It will reduce the cost of getting the thread runtime device tracker,
-and will have a better runtime overhead if user constructs a lot of
-short lived threads that use VTK-m.
--- a/docs/changelog/vec-constexpr-variadic-constructor.md
+++ b/docs/changelog/vec-constexpr-variadic-constructor.md
@ -1,115 +0,0 @@
-# Support constexpr and variadic constructor for Vec
-
-Add variadic constructors to the `vtkm::Vec` classes. The main advantage of this addition
-is that it makes it much easier to initialize `Vec`s of arbitrary length.
-
-Meanwhile, `Vec` classes constructed with values listed in their parameters up to size 4
-are constructed as constant expressions at compile time to reduce runtime overhead.
-Sizes greater than 4 are not yet supported to be constructed at compile time via initializer lists
-since in C++11 constexpr does not allow for loops. Only on Windows platform with a compiler older
-than Visual Studio 2017 version 15.0, users are allowed to use initializer lists to construct a
-vec with size > 4.
-
-`vtkm::make_Vec` would always construct `Vec` at compile time if possible.
-
-``` cpp
-vtkm::Vec<vtkm::Float64, 3> vec1{1.1, 2.2, 3.3};  // New better initializer since
-                                                  // it does not allow type narrowing
-
-vtkm::Vec<vtkm::Float64, 3> vec2 = {1.1, 2.2, 3.3}; // Nice syntax also supported by
-                                                    // initializer lists.
-
-vtkm::Vec<vtkm::Float64, 3> vec3 = vtkm::make_Vec(1.1, 2.2, 3.3); // Old style that still works.
-
-vtkm::Vec<vtkm::Float64, 3> vec3(1.1, 2.2, 3.3); // Old style that still works but
-                                                 // should be deprecated. Reason listed below.
-```
-
-Nested initializer lists work to initialize `Vec` of `Vec`s. If the size is no more than 4,
-it's always constructed at compile time if possible.
-
-``` cpp
-vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 3> vec{ {1.1, 2.2}, {3.3, 4.4}, {5.5, 6.6} };
-                                                           //Constructed at compile time
-```
-
-One drawback about the `std::initializer_list` implementation is that it constructs larger
-`Vec`(size>4) of scalars or `vec`s at run time.
-
-``` cpp
-vtkm::Vec<vtkm::Float64, 5> vec1{1.1, 2.2, 3.3, 4.4, 5.5}; // Constructed at run time.
-
-vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 5> vec2{{1.1, 1.1},{2.2, 2.2},{3.3, 3.3},
-                                      {4.4, 4.4}, {5.5, 5.5}}; // Constructed at run time.
-```
-
-Parenthesis constructor would report an error if the size is larger than 4 when being used
-to construct a `Vec` of `Vec`s. If it's being used to construct a `Vec` of scalars then it's
-fine.
-
-```cpp
-vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 5> vec2({1.1, 1.1},{2.2, 2.2},{3.3, 3.3},
-                                      {4.4, 4.4}, {5.5, 5.5}); // ERROR! This type of
-                                                               // constructor not implemented!
-
-vtkm::Vec<vtkm::Float64, 5> vec1(1.1, 2.2, 3.3, 4.4, 5.5); // Constructed at compile time.
-```
-
-If a `vtkm::Vec` is initialized with a list of size one, then that one value is
-replicated for all components.
-
-``` cpp
-vtkm::Vec<vtkm::Float64, 3> vec{1.1};  // vec gets [ 1.1, 1.1, 1.1 ]
-```
-
-This "scalar" initialization also works for `Vec` of `Vec`s.
-
-``` cpp
-vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 3> vec1{ { 1.1, 2.2 } };
-// vec1 is [[1.1, 2.2], [1.1, 2.2], [1.1, 2.2]]
-
-vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 3> vec2{ { 3.3}, { 4.4 }, { 5.5 } };
-// vec2 is [[3.3, 3.3], [4.4, 4.4], [5.5, 5.5]]
-```
-
-`vtkm::make_Vec` is also updated to support an arbitrary number initial values which are
-constructed at compile time.
-
-``` cpp
-// Creates a vtkm::Vec<vtkm::Float64, 5>
-auto vec = vtkm::make_Vec(1.1, 2.2, 3.3, 4.4, 5.5);
-```
-
-This is super convenient when dealing with variadic function arguments.
-
-``` cpp
-template <typename... Ts>
-void ExampleVariadicFunction(const Ts&... params)
-{
-  auto vec = vtkm::make_Vec(params...);
-```
-
-Of course, this assumes that the type of all the parameters is the same. If not, you
-could run into compiler trouble.
-
-`vtkm::make_Vec` does not accept an `std::initializer_list`,
-
-``` cpp
-// Creates a vtkm::Vec<vtkm::Float64, 3>
-auto vec1 = vtkm::make_Vec<3>({1.1, 2.2, 3.3}); // ERROR
-
-// Creates exactly the same thing but compiles
-auto vec1 = vtkm::make_Vec<3>(1.1, 2.2, 3.3);
-```
-
-A limitation of the initializer list constructor is that the compiler has no way to
-check the length of the list or force it to a particular length. Thus it is entirely
-possible to construct a `Vec` with the wrong number of arguments. Or, more to the
-point, the compiler will let you do it, but there is an assert in the constructor to
-correct for that. (Of course, asserts are not compiled in release builds.)
-
-``` cpp
-// This will compile, but it's results are undefined when it is run.
-// In debug builds, it will fail an assert.
-vtkm::Vec<vtkm::Float64, 3> vec{1.1, 1.2};
-```
--- a/docs/changelog/vec-vec-constructors.md
+++ b/docs/changelog/vec-vec-constructors.md
@ -1,18 +0,0 @@
-# VTK-m Vec< Vec<T> > can't be constructed from Vec<U>
-
- 
-When you have a Vec<Vec<float,3>> it was possible to incorrectly initialize
-it with the contents of a Vec<double,3>. An example of this is:
-```cpp
-using Vec3d = vtkm::Vec<double, 3>;
-using Vec3f = vtkm::Vec<float, 3>;
-using Vec3x3f = vtkm::Vec<Vec3f, 3>;
-
-Vec3d x(0.0, 1.0, 2.0);
-Vec3x3f b(x); // becomes [[0,0,0],[1,1,1],[2,2,2]]
-Vec3x3f c(x, x, x); // becomes [[0,1,2],[0,1,2],[0,1,2]]
-Vec3x3f d(Vec3f(0.0f,1.0f,2.0f)) //becomes [[0,0,0],[1,1,1],[2,2,2]]
-```
-
-So the solution we have chosen is to disallow the construction of objects such
-as b. This still allows the free implicit cast to go from double to float.
--- a/docs/changelog/virtual-object-handle-non-templated-transfer.md
+++ b/docs/changelog/virtual-object-handle-non-templated-transfer.md
@ -1,21 +0,0 @@
-# VTK-m VirtualObjectHandle can transfer to a device using runtime `DeviceAdapterId` value
-
-Previously `VirtualObjectHandle` required the caller to know a compile time device adapter tag to 
-transfer data. This was problematic since in parts of VTK-m you would only have the runtime
-`vtkm::cont::DeviceAdapterId` value of the desired device. To than transfer the 
-`VirtualObjectHandle` you would have to call `FindDeviceAdapterTagAndCall`. All this extra
-work was unneeded as `VirtualObjectHandle` internally was immediately converting from
-a compile time type to a runtime value.
-
-
-Here is an example of how you can now transfer a `VirtualObjectHandle` to a device using
-a runtime value:
-```cpp
-
-template<typename BaseType>
-const BaseType* moveToDevice(VirtualObjectHandle<BaseType>& handle,
-                      vtkm::cont::vtkm::cont::DeviceAdapterId deviceId)
-{
-  return handle.PrepareForExecution(deviceId);
-}
-```
--- a/docs/changelog/vs-cmake-requirements.md
+++ b/docs/changelog/vs-cmake-requirements.md
@ -1,4 +0,0 @@
-# Visual Studio generator now requires CMake 3.11
-
-We require CMake 3.11 for the Visual Studio generators as the
-$<COMPILE_LANGUAGE:> generator expression is not supported on older versions.
--- a/docs/changelog/warp-by-vector-worklet-filter.md
+++ b/docs/changelog/warp-by-vector-worklet-filter.md
@ -1,7 +0,0 @@
-# Add a warp vector worklet and filter
-
-This commit adds a worklet that modifies point coordinates by moving points
-along point normals by the scalar amount. It's a simpified version of the
-vtkWarpScalar in VTK. Additionally the filter doesn't modify the point coordinates,
-but creates a new point coordinates that have been warped.
-Useful for showing flow profiles or mechanical deformation.
--- a/docs/changelog/warp-scalar-worklet-filter.md
+++ b/docs/changelog/warp-scalar-worklet-filter.md
@ -1,7 +0,0 @@
-Add a warpScalar worklet and filter
-
-This commit adds a worklet as well as a filter that modify point coordinates by moving points
-along point normals by the scalar amount times the scalar factor.
-It's a simpified version of the vtkWarpScalar class in VTK. Additionally the filter doesn't
-modify the point coordinates, but creates a new point coordinates that have been warped.
-
--- a/docs/changelog/wavelet_generator.md
+++ b/docs/changelog/wavelet_generator.md
@ -1,5 +0,0 @@
- # Add a `WaveletGenerator` worklet (e.g. vtkRTAnalyticSource)
-
-Add a VTK-m implementation of VTK's `vtkRTAnalyticSource`, or "Wavelet" source
-as it is known in ParaView. This is a customizable dataset with properties
-that make it useful for testing and benchmarking various algorithms.
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@ -25,6 +25,7 @@ set(CMAKE_PREFIX_PATH ${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR})

 add_subdirectory(clipping)
 add_subdirectory(contour_tree)
+add_subdirectory(contour_tree_augmented)
 add_subdirectory(cosmotools)
 add_subdirectory(demo)
 add_subdirectory(dynamic_dispatcher)
--- a/examples/clipping/Clipping.cxx
+++ b/examples/clipping/Clipping.cxx
@ -19,7 +19,6 @@
 //============================================================================
 #include <vtkm/cont/ArrayHandleCast.h>
 #include <vtkm/cont/DataSet.h>
-#include <vtkm/cont/DeviceAdapterAlgorithm.h>
 #include <vtkm/cont/Timer.h>
 #include <vtkm/io/reader/VTKDataSetReader.h>
 #include <vtkm/io/writer/VTKDataSetWriter.h>
@ -37,7 +36,6 @@ using FloatVec3 = vtkm::Vec<vtkm::Float32, 3>;
 namespace
 {

-template <typename DeviceTag>
 struct FieldMapper
 {
  vtkm::cont::DynamicArrayHandle& Output;
@ -58,11 +56,11 @@ struct FieldMapper
  {
    if (this->IsCellField)
    {
-      this->Output = this->Worklet.ProcessCellField(input, DeviceTag());
+      this->Output = this->Worklet.ProcessCellField(input);
    }
    else
    {
-      this->Output = this->Worklet.ProcessPointField(input, DeviceTag());
+      this->Output = this->Worklet.ProcessPointField(input);
    }
  }
 };
@ -79,10 +77,6 @@ int main(int argc, char* argv[])
    return 1;
  }

-  using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
-  std::cout << "Device Adapter Name: " << vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetName()
-            << std::endl;
-
  vtkm::io::reader::VTKDataSetReader reader(argv[1]);
  vtkm::cont::DataSet input = reader.ReadDataSet();

@ -91,15 +85,14 @@ int main(int argc, char* argv[])
  vtkm::Float32 clipValue = std::stof(argv[argc - 2]);
  vtkm::worklet::Clip clip;

-  vtkm::cont::Timer<DeviceAdapter> total;
-  vtkm::cont::Timer<DeviceAdapter> timer;
+  vtkm::cont::Timer<> total;
+  vtkm::cont::Timer<> timer;
  bool invertClip = false;
  vtkm::cont::CellSetExplicit<> outputCellSet =
    clip.Run(input.GetCellSet(0),
             scalarField.GetData().ResetTypeList(vtkm::TypeListTagScalarAll()),
             clipValue,
-             invertClip,
-             DeviceAdapter());
+             invertClip);
  vtkm::Float64 clipTime = timer.GetElapsedTime();

  vtkm::cont::DataSet output;
@ -108,7 +101,7 @@ int main(int argc, char* argv[])

  auto inCoords = input.GetCoordinateSystem(0).GetData();
  timer.Reset();
-  auto outCoords = clip.ProcessCellField(inCoords, DeviceAdapter());
+  auto outCoords = clip.ProcessCellField(inCoords);
  vtkm::Float64 processCoordinatesTime = timer.GetElapsedTime();
  output.AddCoordinateSystem(vtkm::cont::CoordinateSystem("coordinates", outCoords));

@ -132,7 +125,7 @@ int main(int argc, char* argv[])
    }

    vtkm::cont::DynamicArrayHandle outField;
-    FieldMapper<DeviceAdapter> fieldMapper(outField, clip, isCellField);
+    FieldMapper fieldMapper(outField, clip, isCellField);
    inField.GetData().CastAndCall(fieldMapper);
    output.AddField(vtkm::cont::Field(inField.GetName(), inField.GetAssociation(), outField));
  }
--- a/examples/contour_tree_augmented/CMakeLists.txt
+++ b/examples/contour_tree_augmented/CMakeLists.txt
@ -0,0 +1,125 @@
+##=============================================================================
+##
+##  Copyright (c) Kitware, Inc.
+##  All rights reserved.
+##  See LICENSE.txt for details.
+##
+##  This software is distributed WITHOUT ANY WARRANTY; without even
+##  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+##  PURPOSE.  See the above copyright notice for more information.
+##
+##  Copyright 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
+##  Copyright 2018 UT-Battelle, LLC.
+##  Copyright 2018 Los Alamos National Security.
+##
+##  Under the terms of Contract DE-NA0003525 with NTESS,
+##  the U.S. Government retains certain rights in this software.
+##  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+##  Laboratory (LANL), the U.S. Government retains certain rights in
+##  this software.
+##
+##=============================================================================
+##============================================================================
+## Copyright (c) 2018, The Regents of the University of California, through
+## Lawrence Berkeley National Laboratory (subject to receipt of any required approvals
+## from the U.S. Dept. of Energy).  All rights reserved.
+##
+## Redistribution and use in source and binary forms, with or without modification,
+## are permitted provided that the following conditions are met:
+##
+## (1) Redistributions of source code must retain the above copyright notice, this
+##     list of conditions and the following disclaimer.
+##
+## (2) Redistributions in binary form must reproduce the above copyright notice,
+##     this list of conditions and the following disclaimer in the documentation
+##     and/or other materials provided with the distribution.
+##
+## (3) Neither the name of the University of California, Lawrence Berkeley National
+##     Laboratory, U.S. Dept. of Energy nor the names of its contributors may be
+##     used to endorse or promote products derived from this software without
+##     specific prior written permission.
+##
+## THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+## ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+## WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
+## IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
+## INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
+## BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+## DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+## LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
+## OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
+## OF THE POSSIBILITY OF SUCH DAMAGE.
+##
+##=============================================================================
+##
+##  This code is an extension of the algorithm presented in the paper:
+##  Parallel Peak Pruning for Scalable SMP Contour Tree Computation.
+##  Hamish Carr, Gunther Weber, Christopher Sewell, and James Ahrens.
+##  Proceedings of the IEEE Symposium on Large Data Analysis and Visualization
+##  (LDAV), October 2016, Baltimore, Maryland.
+##
+##  The PPP2 algorithm and software were jointly developed by
+##  Hamish Carr (University of Leeds), Gunther H. Weber (LBNL), and
+##  Oliver Ruebel (LBNL)
+##==============================================================================
+
+#Find the VTK-m package
+find_package(VTKm REQUIRED QUIET)
+
+# MAKE DEBUG BUILD
+# set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -g3 -g -fsanitize=address")
+
+# MAKE OPTIMIZED BUILD
+# set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -O3")
+
+####################################
+# Serial builds
+####################################
+# Serial 2D / 3D / MC
+add_executable(ContourTree_PPP2_SERIAL ContourTreeApp.cxx)
+target_link_libraries(ContourTree_PPP2_SERIAL vtkm_cont)
+target_compile_definitions(ContourTree_PPP2_SERIAL PRIVATE
+                           "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_SERIAL")
+
+
+####################################
+# Serial debug builds
+####################################
+
+# Debug Serial 2D / 3D / MC
+add_executable(ContourTree_PPP2_SERIAL_DEBUG ContourTreeApp.cxx)
+target_link_libraries(ContourTree_PPP2_SERIAL_DEBUG vtkm_cont)
+target_compile_definitions(ContourTree_PPP2_SERIAL_DEBUG PRIVATE
+                           "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_SERIAL" "DEBUG_PRINT")
+
+
+if(TARGET vtkm::tbb)
+  # TBB 2D/3D/MC
+  add_executable(ContourTree_PPP2_TBB ContourTreeApp.cxx)
+  target_link_libraries(ContourTree_PPP2_TBB vtkm_cont)
+  target_compile_definitions(ContourTree_PPP2_TBB PRIVATE
+                             "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_TBB" "ENABLE_SET_NUM_THREADS")
+
+  # TBB 2D/3D/MC DEBUG
+  add_executable(ContourTree_PPP2_TBB_DEBUG ContourTreeApp.cxx)
+  target_link_libraries(ContourTree_PPP2_TBB_DEBUG vtkm_cont)
+  target_compile_definitions(ContourTree_PPP2_TBB_DEBUG PRIVATE
+                             "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_TBB" "ENABLE_SET_NUM_THREADS" "DEBUG_PRINT")
+
+endif()
+
+if(TARGET vtkm::cuda)
+  # CUDA 2D/3D/MC
+  vtkm_compile_as_cuda(cudaSource ContourTreeApp.cxx)
+  add_executable(ContourTree_PPP2_CUDA ${cudaSource})
+  target_link_libraries(ContourTree_PPP2_CUDA vtkm_cont)
+  target_compile_definitions(ContourTree_PPP2_CUDA PRIVATE
+                             "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_CUDA")
+
+  # CUDA 2D/3D/MC
+  # vtkm_compile_as_cuda(cudaSource ContourTreeApp.cxx)
+  add_executable(ContourTree_PPP2_CUDA_DEBUG ${cudaSource})
+  target_link_libraries(ContourTree_PPP2_CUDA_DEBUG vtkm_cont)
+  target_compile_definitions(ContourTree_PPP2_CUDA_DEBUG PRIVATE
+                             "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_CUDA" "DEBUG_PRINT")
+endif()
--- a/examples/contour_tree_augmented/ContourTreeApp.cxx
+++ b/examples/contour_tree_augmented/ContourTreeApp.cxx
@ -0,0 +1,435 @@
+//============================================================================
+//  Copyright (c) Kitware, Inc.
+//  All rights reserved.
+//  See LICENSE.txt for details.
+//  This software is distributed WITHOUT ANY WARRANTY; without even
+//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+//  PURPOSE.  See the above copyright notice for more information.
+//
+//  Copyright 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
+//  Copyright 2014 UT-Battelle, LLC.
+//  Copyright 2014 Los Alamos National Security.
+//
+//  Under the terms of Contract DE-NA0003525 with NTESS,
+//  the U.S. Government retains certain rights in this software.
+//
+//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+//  Laboratory (LANL), the U.S. Government retains certain rights in
+//  this software.
+//============================================================================
+// Copyright (c) 2018, The Regents of the University of California, through
+// Lawrence Berkeley National Laboratory (subject to receipt of any required approvals
+// from the U.S. Dept. of Energy).  All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// (1) Redistributions of source code must retain the above copyright notice, this
+//     list of conditions and the following disclaimer.
+//
+// (2) Redistributions in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+// (3) Neither the name of the University of California, Lawrence Berkeley National
+//     Laboratory, U.S. Dept. of Energy nor the names of its contributors may be
+//     used to endorse or promote products derived from this software without
+//     specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
+// IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
+// INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
+// BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+// DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
+// OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
+// OF THE POSSIBILITY OF SUCH DAMAGE.
+//
+//=============================================================================
+//
+//  This code is an extension of the algorithm presented in the paper:
+//  Parallel Peak Pruning for Scalable SMP Contour Tree Computation.
+//  Hamish Carr, Gunther Weber, Christopher Sewell, and James Ahrens.
+//  Proceedings of the IEEE Symposium on Large Data Analysis and Visualization
+//  (LDAV), October 2016, Baltimore, Maryland.
+//
+//  The PPP2 algorithm and software were jointly developed by
+//  Hamish Carr (University of Leeds), Gunther H. Weber (LBNL), and
+//  Oliver Ruebel (LBNL)
+//==============================================================================
+
+#define DEBUG_TIMING
+
+#ifdef ENABLE_SET_NUM_THREADS
+#include "tbb/task_scheduler_init.h"
+#endif
+
+#include <vtkm/Types.h>
+#include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/DataSet.h>
+#include <vtkm/cont/DataSetBuilderUniform.h>
+#include <vtkm/cont/DataSetFieldAdd.h>
+#include <vtkm/cont/Timer.h>
+#include <vtkm/filter/ContourTreeUniformAugmented.h>
+#include <vtkm/worklet/contourtree_augmented/PrintVectors.h>
+#include <vtkm/worklet/contourtree_augmented/ProcessContourTree.h>
+#include <vtkm/worklet/contourtree_augmented/Types.h>
+
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+#include <string>
+#include <utility>
+#include <vector>
+
+namespace cppp2_ns = vtkm::worklet::contourtree_augmented;
+
+// Simple helper class for parsing the command line options
+class ParseCL
+{
+public:
+  ParseCL() {}
+
+  void parse(std::vector<std::string>::size_type argc, char** argv)
+  {
+    mCLOptions.resize(std::vector<std::string>::size_type(argc));
+    for (std::vector<std::string>::size_type i = 1; i < argc; ++i)
+    {
+      this->mCLOptions[i] = std::string(argv[i]);
+    }
+  }
+
+  vtkm::Id findOption(const std::string& option) const
+  {
+    auto it =
+      std::find_if(this->mCLOptions.begin(),
+                   this->mCLOptions.end(),
+                   [option](const std::string& val) -> bool { return val.find(option) == 0; });
+    if (it == this->mCLOptions.end())
+    {
+      return -1;
+    }
+    else
+    {
+      return static_cast<vtkm::Id>(it - this->mCLOptions.begin());
+    }
+  }
+
+  bool hasOption(const std::string& option) const { return this->findOption(option) >= 0; }
+
+  std::string getOption(const std::string& option) const
+  {
+    vtkm::Id index = this->findOption(option);
+    if (index >= 0)
+    {
+      std::string val = this->mCLOptions[std::vector<std::string>::size_type(index)];
+      auto valPos = val.find("=");
+      if (valPos)
+      {
+        return val.substr(valPos + 1);
+      }
+    }
+    return std::string("");
+  }
+
+  const std::vector<std::string>& getOptions() const { return this->mCLOptions; }
+
+private:
+  std::vector<std::string> mCLOptions;
+};
+
+
+
+// Compute and render an isosurface for a uniform grid example
+int main(int argc, char* argv[])
+{
+  // TODO: Change timing to use logging in vtkm/cont/Logging.h
+  vtkm::cont::Timer<> totalTime;
+  vtkm::Float64 prevTime = 0;
+  vtkm::Float64 currTime = 0;
+  std::cout << "ContourTreePPP2Mesh <options> <fileName>" << std::endl;
+
+  ////////////////////////////////////////////
+  // Parse the command line options
+  ////////////////////////////////////////////
+  ParseCL parser;
+  parser.parse(std::vector<std::string>::size_type(argc), argv);
+  std::string filename = parser.getOptions().back();
+  bool computeRegularStructure = true;
+  bool useMarchingCubes = false;
+  bool computeBranchDecomposition = true;
+  bool printContourTree = false;
+  if (parser.hasOption("--augmentTree"))
+    computeRegularStructure = std::stoi(parser.getOption("--augmentTree"));
+  if (parser.hasOption("--mc"))
+    useMarchingCubes = true;
+  if (parser.hasOption("--printCT"))
+    printContourTree = true;
+  if (parser.hasOption("--branchDecomp"))
+    computeBranchDecomposition = std::stoi(parser.getOption("--branchDecomp"));
+#ifdef ENABLE_SET_NUM_THREADS
+  int numThreads = tbb::task_scheduler_init::default_num_threads();
+  if (parser.hasOption("--numThreads"))
+    numThreads = std::stoi(parser.getOption("--numThreads"));
+  tbb::task_scheduler_init schedulerInit(numThreads);
+#endif
+
+  if (argc < 2 || parser.hasOption("--help") || parser.hasOption("-h"))
+  {
+    std::cout << "Parameter is <fileName>" << std::endl;
+    std::cout << "File is expected to be ASCII with either: " << std::endl;
+    std::cout << "  - xdim ydim integers for 2D or" << std::endl;
+    std::cout << "  - xdim ydim zdim integers for 3D" << std::endl;
+    std::cout << "followed by vector data last dimension varying fastest" << std::endl;
+
+    std::cout << std::endl;
+    std::cout << "Options: (Bool options are give via int, i.e. =0 for False and =1 for True)"
+              << std::endl;
+    std::cout << "--mc              Use marching cubes interpolation for contour tree calculation. "
+                 "(Default=False)"
+              << std::endl;
+    std::cout << "--augmentTree     Compute the augmented contour tree. (Default=True)"
+              << std::endl;
+    std::cout << "--branchDecomp    Compute the volume branch decomposition for the contour tree. "
+                 "Requires --augmentTree (Default=True)"
+              << std::endl;
+    std::cout << "--printCT         Print the contour tree. (Default=False)" << std::endl;
+#ifdef ENABLE_SET_NUM_THREADS
+    std::cout << "--numThreads      Specify the number of threads to use. Available only with TBB."
+              << std::endl;
+#endif
+    return 0;
+  }
+
+  std::cout << "Settings:" << std::endl;
+  std::cout << "    filename=" << filename << std::endl;
+  std::cout << "    mc=" << useMarchingCubes << std::endl;
+  std::cout << "    augmentTree=" << computeRegularStructure << std::endl;
+  std::cout << "    branchDecomp=" << computeBranchDecomposition << std::endl;
+#ifdef ENABLE_SET_NUM_THREADS
+  std::cout << "    numThreads=" << numThreads << std::endl;
+#endif
+  std::cout << std::endl;
+
+
+  currTime = totalTime.GetElapsedTime();
+  vtkm::Float64 startUpTime = currTime - prevTime;
+  prevTime = currTime;
+
+  ///////////////////////////////////////////////
+  // Read the input data
+  ///////////////////////////////////////////////
+  std::ifstream inFile(filename);
+  if (inFile.bad())
+    return 0;
+
+  // Read the dimensions of the mesh, i.e,. number of elementes in x, y, and z
+  std::vector<vtkm::Id> dims;
+  std::string line;
+  getline(inFile, line);
+  std::istringstream linestream(line);
+  vtkm::Id dimVertices;
+  while (linestream >> dimVertices)
+  {
+    dims.push_back(dimVertices);
+  }
+
+  // Compute the number of vertices, i.e., xdim * ydim * zdim
+  auto nDims = dims.size();
+  std::vector<vtkm::Float32>::size_type nVertices = std::vector<vtkm::Float32>::size_type(
+    std::accumulate(dims.begin(), dims.end(), vtkm::Id(1), std::multiplies<vtkm::Id>()));
+
+  // Print the mesh metadata
+  std::cout << "Number of dimensions: " << nDims << std::endl;
+  std::cout << "Number of mesh vertices: " << nVertices << std::endl;
+
+  // Check the the number of dimensiosn is either 2D or 3D
+  if (nDims < 2 || nDims > 3)
+  {
+    std::cout << "The input mesh is " << nDims << "D. Input data must be either 2D or 3D."
+              << std::endl;
+    return 0;
+  }
+  if (useMarchingCubes && nDims != 3)
+  {
+    std::cout << "The input mesh is " << nDims
+              << "D. Contour tree using marching cubes only supported for 3D data." << std::endl;
+    return 0;
+  }
+
+  // read data
+  std::vector<vtkm::Float32> values(nVertices);
+  for (std::vector<vtkm::Float32>::size_type vertex = 0; vertex < nVertices; vertex++)
+  {
+    inFile >> values[vertex];
+  }
+
+  // finish reading the data
+  inFile.close();
+
+  currTime = totalTime.GetElapsedTime();
+  vtkm::Float64 dataReadTime = currTime - prevTime;
+  prevTime = currTime;
+  // build the input dataset
+  vtkm::cont::DataSetBuilderUniform dsb;
+  vtkm::cont::DataSet inDataSet;
+  // 2D data
+  if (nDims == 2)
+  {
+    vtkm::Id2 vdims;
+    vdims[0] = dims[0];
+    vdims[1] = dims[1];
+    inDataSet = dsb.Create(vdims);
+  }
+  // 3D data
+  else
+  {
+    vtkm::Id3 vdims;
+    vdims[0] = dims[0];
+    vdims[1] = dims[1];
+    vdims[2] = dims[2];
+    inDataSet = dsb.Create(vdims);
+  }
+  vtkm::cont::DataSetFieldAdd dsf;
+  dsf.AddPointField(inDataSet, "values", values);
+
+  currTime = totalTime.GetElapsedTime();
+  vtkm::Float64 buildDatasetTime = currTime - prevTime;
+  prevTime = currTime;
+
+  // Output data set is pairs of saddle and peak vertex IDs
+  vtkm::cont::DataSet result;
+
+  // Convert the mesh of values into contour tree, pairs of vertex ids
+  vtkm::filter::ContourTreePPP2 filter(useMarchingCubes, computeRegularStructure);
+  filter.SetActiveField("values");
+  result = filter.Execute(inDataSet); //, std::string("values"));
+
+  currTime = totalTime.GetElapsedTime();
+  vtkm::Float64 computeContourTreeTime = currTime - prevTime;
+  prevTime = currTime;
+
+#ifdef DEBUG_TIMING
+  std::cout << "-------------------------------------------------------------" << std::endl;
+  std::cout << "-------------------Contour Tree Timings----------------------" << std::endl;
+  // Get the timings from the contour tree computation
+  const std::vector<std::pair<std::string, vtkm::Float64>>& contourTreeTimings =
+    filter.GetTimings();
+  for (std::vector<std::pair<std::string, vtkm::Float64>>::size_type i = 0;
+       i < contourTreeTimings.size();
+       i++)
+    std::cout << std::setw(42) << std::left << contourTreeTimings[i].first << ": "
+              << contourTreeTimings[i].second << " seconds" << std::endl;
+#endif
+
+  ////////////////////////////////////////////
+  // Compute the branch decomposition
+  ////////////////////////////////////////////
+  if (computeBranchDecomposition)
+  {
+    // TODO: Change timing to use logging in vtkm/cont/Logging.h
+    vtkm::cont::Timer<> branchDecompTimer;
+    // compute the volume for each hyperarc and superarc
+    cppp2_ns::IdArrayType superarcIntrinsicWeight;
+    cppp2_ns::IdArrayType superarcDependentWeight;
+    cppp2_ns::IdArrayType supernodeTransferWeight;
+    cppp2_ns::IdArrayType hyperarcDependentWeight;
+
+    cppp2_ns::ProcessContourTree::ComputeVolumeWeights(filter.GetContourTree(),
+                                                       filter.GetNumIterations(),
+                                                       superarcIntrinsicWeight,  // (output)
+                                                       superarcDependentWeight,  // (output)
+                                                       supernodeTransferWeight,  // (output)
+                                                       hyperarcDependentWeight); // (output)
+    std::cout << std::setw(42) << std::left << "Compute Volume Weights"
+              << ": " << branchDecompTimer.GetElapsedTime() << " seconds" << std::endl;
+    branchDecompTimer.Reset();
+
+    // compute the branch decomposition by volume
+    cppp2_ns::IdArrayType whichBranch;
+    cppp2_ns::IdArrayType branchMinimum;
+    cppp2_ns::IdArrayType branchMaximum;
+    cppp2_ns::IdArrayType branchSaddle;
+    cppp2_ns::IdArrayType branchParent;
+
+    cppp2_ns::ProcessContourTree::ComputeVolumeBranchDecomposition(filter.GetContourTree(),
+                                                                   superarcDependentWeight,
+                                                                   superarcIntrinsicWeight,
+                                                                   whichBranch,   // (output)
+                                                                   branchMinimum, // (output)
+                                                                   branchMaximum, // (output)
+                                                                   branchSaddle,  // (output)
+                                                                   branchParent); // (output)
+    std::cout << std::setw(42) << std::left << "Compute Volume Branch Decomposition"
+              << ": " << branchDecompTimer.GetElapsedTime() << " seconds" << std::endl;
+  }
+  currTime = totalTime.GetElapsedTime();
+  vtkm::Float64 computeBranchDecompTime = currTime - prevTime;
+  prevTime = currTime;
+
+  //vtkm::cont::Field resultField =  result.GetField();
+  //vtkm::cont::ArrayHandle<vtkm::Pair<vtkm::Id, vtkm::Id> > saddlePeak;
+  //resultField.GetData().CopyTo(saddlePeak);
+
+  // dump out contour tree for comparison
+  if (printContourTree)
+  {
+    std::cout << "Contour Tree" << std::endl;
+    std::cout << "============" << std::endl;
+    cppp2_ns::EdgePairArray saddlePeak;
+    cppp2_ns::ProcessContourTree::CollectSortedSuperarcs(
+      filter.GetContourTree(), filter.GetSortOrder(), saddlePeak);
+    cppp2_ns::printEdgePairArray(saddlePeak);
+  }
+
+#ifdef DEBUG_TIMING
+  std::cout << "-------------------------------------------------------------" << std::endl;
+  std::cout << "--------------------------Totals-----------------------------" << std::endl;
+  std::cout << std::setw(42) << std::left << "Start-up"
+            << ": " << startUpTime << " seconds" << std::endl;
+  std::cout << std::setw(42) << std::left << "Data Read"
+            << ": " << dataReadTime << " seconds" << std::endl;
+  std::cout << std::setw(42) << std::left << "Build VTKM Dataset"
+            << ": " << buildDatasetTime << " seconds" << std::endl;
+  std::cout << std::setw(42) << std::left << "Compute Contour Tree"
+            << ": " << computeContourTreeTime << " seconds" << std::endl;
+  if (computeBranchDecomposition)
+  {
+    std::cout << std::setw(42) << std::left << "Compute Branch Decomposition"
+              << ": " << computeBranchDecompTime << " seconds" << std::endl;
+  }
+  currTime = totalTime.GetElapsedTime();
+  //vtkm::Float64 miscTime = currTime - startUpTime - dataReadTime - buildDatasetTime - computeContourTreeTime;
+  //if (computeBranchDecomposition) miscTime -= computeBranchDecompTime;
+  //std::cout<<std::setw(42)<<std::left<<"Misc. Times"<<": "<<miscTime<<" seconds"<<std::endl;
+  std::cout << std::setw(42) << std::left << "Total Time"
+            << ": " << currTime << " seconds" << std::endl;
+
+  std::cout << "-------------------------------------------------------------" << std::endl;
+  std::cout << "----------------Contour Tree Array Sizes---------------------" << std::endl;
+  const vtkm::worklet::contourtree_augmented::ContourTree& ct = filter.GetContourTree();
+  std::cout << std::setw(42) << std::left << "#Nodes"
+            << ": " << ct.nodes.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Arcs"
+            << ": " << ct.arcs.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Superparents"
+            << ": " << ct.superparents.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Superarcs"
+            << ": " << ct.superarcs.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Supernodes"
+            << ": " << ct.supernodes.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Hyperparents"
+            << ": " << ct.hyperparents.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#WhenTransferred"
+            << ": " << ct.whenTransferred.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Hypernodes"
+            << ": " << ct.hypernodes.GetNumberOfValues() << std::endl;
+  std::cout << std::setw(42) << std::left << "#Hyperarcs"
+            << ": " << ct.hyperarcs.GetNumberOfValues() << std::endl;
+
+#endif
+
+  return 0;
+}
--- a/examples/cosmotools/CMakeLists.txt
+++ b/examples/cosmotools/CMakeLists.txt
@ -25,24 +25,16 @@ project(CosmoTools CXX)
 #Find the VTK-m package
 find_package(VTKm REQUIRED QUIET)

-add_executable(CosmoCenterFinder_SERIAL CosmoCenterFinder.cxx)
-add_executable(CosmoHaloFinder_SERIAL CosmoHaloFinder.cxx)
-
-target_link_libraries(CosmoCenterFinder_SERIAL PRIVATE vtkm_cont)
-target_link_libraries(CosmoHaloFinder_SERIAL PRIVATE vtkm_cont)
-
 if(TARGET vtkm::cuda)
-  add_executable(CosmoCenterFinder_CUDA CosmoCenterFinder.cu)
-  add_executable(CosmoHaloFinder_CUDA CosmoHaloFinder.cu)
+  vtkm_compile_as_cuda(cu_srcs CosmoCenterFinder.cxx)
+  add_executable(CosmoCenterFinder ${cu_srcs})

-  target_link_libraries(CosmoCenterFinder_CUDA PRIVATE vtkm_cont)
-  target_link_libraries(CosmoHaloFinder_CUDA PRIVATE vtkm_cont)
+  vtkm_compile_as_cuda(cu_srcs CosmoHaloFinder.cxx)
+  add_executable(CosmoHaloFinder ${cu_srcs})
+else()
+  add_executable(CosmoCenterFinder CosmoCenterFinder.cxx)
+  add_executable(CosmoHaloFinder CosmoHaloFinder.cxx)
 endif()

-if(TARGET vtkm::tbb)
-  add_executable(CosmoCenterFinder_TBB CosmoCenterFinderTBB.cxx)
-  add_executable(CosmoHaloFinder_TBB CosmoHaloFinderTBB.cxx)
-
-  target_link_libraries(CosmoCenterFinder_TBB PRIVATE vtkm_cont)
-  target_link_libraries(CosmoHaloFinder_TBB PRIVATE vtkm_cont)
-endif()
+target_link_libraries(CosmoCenterFinder PRIVATE vtkm_cont)
+target_link_libraries(CosmoHaloFinder PRIVATE vtkm_cont)
--- a/examples/cosmotools/CosmoCenterFinder.cxx
+++ b/examples/cosmotools/CosmoCenterFinder.cxx
@ -19,8 +19,7 @@
 //============================================================================
 #include <vtkm/cont/ArrayHandleCast.h>
 #include <vtkm/cont/DataSet.h>
-#include <vtkm/cont/DeviceAdapterAlgorithm.h>
-#include <vtkm/cont/Timer.h>
+#include <vtkm/cont/Logging.h>
 #include <vtkm/io/reader/VTKDataSetReader.h>
 #include <vtkm/io/writer/VTKDataSetWriter.h>

@ -32,7 +31,7 @@
 #include <stdexcept>
 #include <string>

-using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
+static const vtkm::cont::LogLevel CosmoLogLevel = vtkm::cont::LogLevel(1);

 void TestCosmoCenterFinder(const char* fileName)
 {
@ -72,29 +71,27 @@ void TestCosmoCenterFinder(const char* fileName)
  vtkm::Pair<vtkm::Id, vtkm::Float32> nxnResult;
  vtkm::Pair<vtkm::Id, vtkm::Float32> mxnResult;

-  vtkm::cont::Timer<DeviceAdapter> total;
-  vtkm::cont::Timer<DeviceAdapter> timer;
-
-  // Create the worklet and run it
-  vtkm::Float32 particleMass = 1.08413e+09f;
-
+  const vtkm::Float32 particleMass = 1.08413e+09f;
  vtkm::worklet::CosmoTools cosmoTools;
-  cosmoTools.RunMBPCenterFinderNxN(
-    xLocArray, yLocArray, zLocArray, nParticles, particleMass, nxnResult, DeviceAdapter());
-  vtkm::Float64 nxnTime = timer.GetElapsedTime();

-  std::cout << "**** NxN MPB = " << nxnResult.first << "  potential = " << nxnResult.second
-            << std::endl;
-  std::cout << "**** Time for NxN: " << nxnTime << std::endl;
+  {
+    VTKM_LOG_SCOPE(CosmoLogLevel, "Executing NxN");

-  timer.Reset();
-  cosmoTools.RunMBPCenterFinderMxN(
-    xLocArray, yLocArray, zLocArray, nParticles, particleMass, mxnResult, DeviceAdapter());
-  vtkm::Float64 estTime = timer.GetElapsedTime();
+    cosmoTools.RunMBPCenterFinderNxN(
+      xLocArray, yLocArray, zLocArray, nParticles, particleMass, nxnResult);

-  std::cout << "**** MxN MPB = " << mxnResult.first << "  potential = " << mxnResult.second
-            << std::endl;
-  std::cout << "**** Time for MxN: " << estTime << std::endl;
+    VTKM_LOG_S(CosmoLogLevel,
+               "NxN MPB = " << nxnResult.first << "  potential = " << nxnResult.second);
+  }
+
+  {
+    VTKM_LOG_SCOPE(CosmoLogLevel, "Executing MxN");
+    cosmoTools.RunMBPCenterFinderMxN(
+      xLocArray, yLocArray, zLocArray, nParticles, particleMass, mxnResult);
+
+    VTKM_LOG_S(CosmoLogLevel,
+               "MxN MPB = " << mxnResult.first << "  potential = " << mxnResult.second);
+  }

  if (nxnResult.first == mxnResult.first)
    std::cout << "FOUND CORRECT PARTICLE " << mxnResult.first << " with potential "
@ -123,12 +120,20 @@ void TestCosmoCenterFinder(const char* fileName)

 int main(int argc, char* argv[])
 {
+  vtkm::cont::SetLogLevelName(CosmoLogLevel, "Cosmo");
+  vtkm::cont::SetStderrLogLevel(CosmoLogLevel);
+  vtkm::cont::InitLogging(argc, argv);
+
  if (argc < 2)
  {
    std::cout << "Usage: " << std::endl << "$ " << argv[0] << " <input_file>" << std::endl;
    return 1;
  }

+#ifndef VTKM_ENABLE_LOGGING
+  std::cout << "Warning: turn on VTKm_ENABLE_LOGGING CMake option to turn on timing." << std::endl;
+#endif
+
  TestCosmoCenterFinder(argv[1]);

  return 0;
--- a/examples/cosmotools/CosmoHaloFinder.cu
+++ b/examples/cosmotools/CosmoHaloFinder.cu
@ -1,23 +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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
-//  Copyright 2014 UT-Battelle, LLC.
-//  Copyright 2014 Los Alamos National Security.
-//
-//  Under the terms of Contract DE-NA0003525 with NTESS,
-//  the U.S. Government retains certain rights in this software.
-//
-//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
-//  Laboratory (LANL), the U.S. Government retains certain rights in
-//  this software.
-//============================================================================
-
-#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_CUDA
-
-#include "CosmoHaloFinder.cxx"
--- a/examples/cosmotools/CosmoHaloFinder.cxx
+++ b/examples/cosmotools/CosmoHaloFinder.cxx
@ -19,8 +19,7 @@
 //============================================================================
 #include <vtkm/cont/ArrayHandleCast.h>
 #include <vtkm/cont/DataSet.h>
-#include <vtkm/cont/DeviceAdapterAlgorithm.h>
-#include <vtkm/cont/Timer.h>
+#include <vtkm/cont/Logging.h>
 #include <vtkm/io/reader/VTKDataSetReader.h>
 #include <vtkm/io/writer/VTKDataSetWriter.h>

@ -32,7 +31,7 @@
 #include <stdexcept>
 #include <string>

-using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
+static const vtkm::cont::LogLevel CosmoLogLevel = vtkm::cont::LogLevel(1);

 void TestCosmoHaloFinder(const char* fileName)
 {
@ -73,28 +72,26 @@ void TestCosmoHaloFinder(const char* fileName)
  vtkm::cont::ArrayHandle<vtkm::Id> resultMBP;
  vtkm::cont::ArrayHandle<vtkm::Float32> resultPot;

-  vtkm::cont::Timer<DeviceAdapter> total;
-  vtkm::cont::Timer<DeviceAdapter> timer;
-
  // Create the worklet and run it
  vtkm::Id minHaloSize = 20;
  vtkm::Float32 linkingLength = 0.2f;
  vtkm::Float32 particleMass = 1.08413e+09f;

-  vtkm::worklet::CosmoTools cosmoTools;
-  cosmoTools.RunHaloFinder(xLocArray,
-                           yLocArray,
-                           zLocArray,
-                           nParticles,
-                           particleMass,
-                           minHaloSize,
-                           linkingLength,
-                           resultHaloId,
-                           resultMBP,
-                           resultPot,
-                           DeviceAdapter());
-  vtkm::Float64 haloTime = timer.GetElapsedTime();
-  std::cout << "**** Time for HaloFinder: " << haloTime << std::endl;
+  {
+    VTKM_LOG_SCOPE(CosmoLogLevel, "Executing HaloFinder");
+
+    vtkm::worklet::CosmoTools cosmoTools;
+    cosmoTools.RunHaloFinder(xLocArray,
+                             yLocArray,
+                             zLocArray,
+                             nParticles,
+                             particleMass,
+                             minHaloSize,
+                             linkingLength,
+                             resultHaloId,
+                             resultMBP,
+                             resultPot);
+  }

  xLocArray.ReleaseResources();
  yLocArray.ReleaseResources();
@ -117,12 +114,20 @@ void TestCosmoHaloFinder(const char* fileName)

 int main(int argc, char* argv[])
 {
+  vtkm::cont::SetLogLevelName(CosmoLogLevel, "Cosmo");
+  vtkm::cont::SetStderrLogLevel(CosmoLogLevel);
+  vtkm::cont::InitLogging(argc, argv);
+
  if (argc < 2)
  {
    std::cout << "Usage: " << std::endl << "$ " << argv[0] << " <input_file>" << std::endl;
    return 1;
  }

+#ifndef VTKM_ENABLE_LOGGING
+  std::cout << "Warning: turn on VTKm_ENABLE_LOGGING CMake option to turn on timing." << std::endl;
+#endif
+
  TestCosmoHaloFinder(argv[1]);

  return 0;
--- a/examples/cosmotools/CosmoHaloFinderTBB.cxx
+++ b/examples/cosmotools/CosmoHaloFinderTBB.cxx
@ -1,23 +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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
-//  Copyright 2014 UT-Battelle, LLC.
-//  Copyright 2014 Los Alamos National Security.
-//
-//  Under the terms of Contract DE-NA0003525 with NTESS,
-//  the U.S. Government retains certain rights in this software.
-//
-//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
-//  Laboratory (LANL), the U.S. Government retains certain rights in
-//  this software.
-//============================================================================
-
-#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_TBB
-
-#include "CosmoHaloFinder.cxx"
--- a/examples/game_of_life/GameOfLife.cxx
+++ b/examples/game_of_life/GameOfLife.cxx
@ -122,23 +122,12 @@ struct UpdateLifeState : public vtkm::worklet::WorkletPointNeighborhood3x3x3

 class GameOfLife : public vtkm::filter::FilterDataSet<GameOfLife>
 {
-  bool PrintedDeviceMsg = false;
-
 public:
-  template <typename Policy, typename Device>
+  template <typename Policy>
  VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input,
-                                          vtkm::filter::PolicyBase<Policy> policy,
-                                          Device)
+                                          vtkm::filter::PolicyBase<Policy> policy)

  {
-    if (!this->PrintedDeviceMsg)
-    {
-      using DeviceAdapterTraits = vtkm::cont::DeviceAdapterTraits<Device>;
-      std::cout << "Running GameOfLife filter on device adapter: " << DeviceAdapterTraits::GetName()
-                << std::endl;
-      this->PrintedDeviceMsg = true;
-    }
-
    using DispatcherType = vtkm::worklet::DispatcherPointNeighborhood<UpdateLifeState>;


@ -154,7 +143,6 @@ public:

    //Update the game state
    DispatcherType dispatcher;
-    dispatcher.SetDevice(Device());
    dispatcher.Invoke(vtkm::filter::ApplyPolicy(cells, policy), prevstate, state, colors);

    //save the results
@ -171,12 +159,11 @@ public:
    return output;
  }

-  template <typename T, typename StorageType, typename DerivedPolicy, typename DeviceAdapter>
+  template <typename T, typename StorageType, typename DerivedPolicy>
  VTKM_CONT bool DoMapField(vtkm::cont::DataSet&,
                            const vtkm::cont::ArrayHandle<T, StorageType>&,
                            const vtkm::filter::FieldMetadata&,
-                            const vtkm::filter::PolicyBase<DerivedPolicy>&,
-                            DeviceAdapter)
+                            vtkm::filter::PolicyBase<DerivedPolicy>)
  {
    return false;
  }
--- a/examples/multi_backend/MultiDeviceGradient.h
+++ b/examples/multi_backend/MultiDeviceGradient.h
@ -28,7 +28,7 @@

 #include <thread>

-using RuntimeTaskQueue = TaskQueue<std::function<void(const vtkm::cont::RuntimeDeviceTracker&)>>;
+using RuntimeTaskQueue = TaskQueue<std::function<void()>>;

 /// \brief Construct a MultiDeviceGradient for a given multiblock dataset
 ///
--- a/examples/multi_backend/MultiDeviceGradient.hxx
+++ b/examples/multi_backend/MultiDeviceGradient.hxx
@ -45,10 +45,10 @@ void process_block_tbb(RuntimeTaskQueue& queue)
 {
  //Step 1. Set the device adapter to this thread to TBB.
  //This makes sure that any vtkm::filters used by our
-  //task operate only on TBB
+  //task operate only on TBB. The "global" thread tracker
+  //is actually thread-local, so we can use that.
  //
-  vtkm::cont::RuntimeDeviceTracker tracker;
-  tracker.ForceDevice(vtkm::cont::DeviceAdapterTagTBB{});
+  vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagTBB{});

  while (queue.hasTasks())
  {
@ -60,7 +60,7 @@ void process_block_tbb(RuntimeTaskQueue& queue)
    //when the queue is empty and we are shutting down
    if (task != nullptr)
    {
-      task(tracker);
+      task();
    }

    //Step 4. Notify the queue that we finished processing this task
@ -73,13 +73,10 @@ void process_block_cuda(RuntimeTaskQueue& queue, int gpuId)
 {
  //Step 1. Set the device adapter to this thread to cuda.
  //This makes sure that any vtkm::filters used by our
-  //task operate only on cuda
+  //task operate only on cuda.  The "global" thread tracker
+  //is actually thread-local, so we can use that.
  //
-  vtkm::cont::RuntimeDeviceTracker tracker;
-#if defined(VTKM_ENABLE_CUDA)
-  auto error = cudaSetDevice(gpuId);
-  tracker.ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
-#endif
+  vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
  (void)gpuId;

  while (queue.hasTasks())
@ -87,12 +84,12 @@ void process_block_cuda(RuntimeTaskQueue& queue, int gpuId)
    //Step 2. Get the task to run on cuda
    auto task = queue.pop();

-    //Step 3. Run the task on TBB. We check the validity
+    //Step 3. Run the task on cuda. We check the validity
    //of the task since we could be given an empty task
    //when the queue is empty and we are shutting down
    if (task != nullptr)
    {
-      task(tracker);
+      task();
    }

    //Step 4. Notify the queue that we finished processing this task
@ -190,11 +187,8 @@ inline VTKM_CONT vtkm::cont::MultiBlock MultiDeviceGradient::PrepareForExecution
  {
    vtkm::cont::DataSet input = *block;
    this->Queue.push( //build a lambda that is the work to do
-      [=](const vtkm::cont::RuntimeDeviceTracker& tracker) {
-        //make a per thread copy of the filter
-        //and give it the device tracker
+      [=]() {
        vtkm::filter::Gradient perThreadGrad = gradient;
-        perThreadGrad.SetRuntimeDeviceTracker(tracker);

        vtkm::cont::DataSet result = perThreadGrad.Execute(input, policy);
        outPtr->ReplaceBlock(0, result);
@ -212,11 +206,8 @@ inline VTKM_CONT vtkm::cont::MultiBlock MultiDeviceGradient::PrepareForExecution
    //will allows us to have multiple works execute in a non
    //blocking manner
    this->Queue.push( //build a lambda that is the work to do
-      [=](const vtkm::cont::RuntimeDeviceTracker& tracker) {
-        //make a per thread copy of the filter
-        //and give it the device tracker
+      [=]() {
        vtkm::filter::Gradient perThreadGrad = gradient;
-        perThreadGrad.SetRuntimeDeviceTracker(tracker);

        vtkm::cont::DataSet result = perThreadGrad.Execute(input, policy);
        outPtr->ReplaceBlock(index, result);
--- a/examples/particle_advection/ParticleAdvection.cxx
+++ b/examples/particle_advection/ParticleAdvection.cxx
@ -18,10 +18,6 @@
 //  this software.
 //============================================================================

-#ifndef VTKM_DEVICE_ADAPTER
-#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_SERIAL
-#endif
-
 #include <vtkm/cont/DataSet.h>
 #include <vtkm/worklet/ParticleAdvection.h>
 #include <vtkm/worklet/particleadvection/GridEvaluators.h>
@ -66,20 +62,14 @@ void RunTest(const std::string& fname,
             vtkm::Id advectType,
             vtkm::Id seeding)
 {
-  using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
-
  using FieldType = vtkm::Float32;
  using FieldHandle = vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3>>;
-  using FieldPortalConstType =
-    typename FieldHandle::template ExecutionTypes<DeviceAdapter>::PortalConst;

  vtkm::io::reader::BOVDataSetReader rdr(fname);
  vtkm::cont::DataSet ds = rdr.ReadDataSet();

-  using RGEvalType = vtkm::worklet::particleadvection::UniformGridEvaluate<FieldPortalConstType,
-                                                                           FieldType,
-                                                                           DeviceAdapter>;
-  using RK4RGType = vtkm::worklet::particleadvection::RK4Integrator<RGEvalType, FieldType>;
+  using RGEvalType = vtkm::worklet::particleadvection::UniformGridEvaluate<FieldHandle>;
+  using RK4RGType = vtkm::worklet::particleadvection::RK4Integrator<RGEvalType>;

  vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3>> fieldArray;
  ds.GetField(0).GetData().CopyTo(fieldArray);
@ -167,12 +157,12 @@ void RunTest(const std::string& fname,
  if (advectType == 0)
  {
    vtkm::worklet::ParticleAdvection particleAdvection;
-    particleAdvection.Run(rk4, seedArray, numSteps, DeviceAdapter());
+    particleAdvection.Run(rk4, seedArray, numSteps);
  }
  else
  {
    vtkm::worklet::Streamline streamline;
-    streamline.Run(rk4, seedArray, numSteps, DeviceAdapter());
+    streamline.Run(rk4, seedArray, numSteps);
  }

  auto t1 = std::chrono::high_resolution_clock::now() - t0;
--- a/examples/redistribute_points/RedistributePoints.h
+++ b/examples/redistribute_points/RedistributePoints.h
@ -22,7 +22,7 @@
 #include <vtkm/cont/Algorithm.h>
 #include <vtkm/cont/AssignerMultiBlock.h>
 #include <vtkm/cont/BoundsGlobalCompute.h>
-#include <vtkm/cont/diy/Serialization.h>
+#include <vtkm/cont/Serialization.h>
 #include <vtkm/filter/ExtractPoints.h>
 #include <vtkm/filter/Filter.h>

@ -73,6 +73,66 @@ class Redistributor
    return extractor.Execute(input, this->Policy);
  }

+  class ConcatenateFields
+  {
+  public:
+    explicit ConcatenateFields(vtkm::Id totalSize) : TotalSize(totalSize), CurrentIdx(0) {}
+
+    void Append(const vtkm::cont::Field& field)
+    {
+      VTKM_ASSERT(this->CurrentIdx + field.GetData().GetNumberOfValues() <= this->TotalSize);
+
+      if (this->Field.GetData().GetNumberOfValues() == 0)
+      {
+        this->Field = field;
+        field.GetData().CastAndCall(Allocator{}, this->Field, this->TotalSize);
+      }
+      else
+      {
+        VTKM_ASSERT(this->Field.GetName() == field.GetName() &&
+                    this->Field.GetAssociation() == field.GetAssociation());
+      }
+
+      field.GetData().CastAndCall(Appender{}, this->Field, this->CurrentIdx);
+      this->CurrentIdx += field.GetData().GetNumberOfValues();
+    }
+
+    const vtkm::cont::Field& GetResult() const
+    {
+      return this->Field;
+    }
+
+  private:
+    struct Allocator
+    {
+      template <typename T, typename S>
+      void operator()(const vtkm::cont::ArrayHandle<T, S>&,
+                      vtkm::cont::Field& field,
+                      vtkm::Id totalSize) const
+      {
+        vtkm::cont::ArrayHandle<T> init;
+        init.Allocate(totalSize);
+        field.SetData(init);
+      }
+    };
+
+    struct Appender
+    {
+      template <typename T, typename S>
+      void operator()(const vtkm::cont::ArrayHandle<T, S>& data,
+                      vtkm::cont::Field& field,
+                      vtkm::Id currentIdx) const
+      {
+        vtkm::cont::ArrayHandle<T> farray =
+          field.GetData().template Cast<vtkm::cont::ArrayHandle<T>>();
+        vtkm::cont::Algorithm::CopySubRange(data, 0, data.GetNumberOfValues(), farray, currentIdx);
+      }
+    };
+
+    vtkm::Id TotalSize;
+    vtkm::Id CurrentIdx;
+    vtkm::cont::Field Field;
+  };

 public:
  Redistributor(const diy::RegularDecomposer<diy::ContinuousBounds>& decomposer,
@ -95,11 +155,7 @@ public:
          this->Decomposer.fill_bounds(bds, target.gid);

          auto extractedDS = this->Extract(*block, bds);
-          const auto inputAH = extractedDS.GetCoordinateSystem(0).GetData();
-          vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> outputAH;
-          outputAH.Allocate(inputAH.GetNumberOfValues());
-          vtkm::cont::Algorithm::Copy(inputAH, outputAH);
-          rp.enqueue(target, outputAH);
+          rp.enqueue(target, vtkm::filter::MakeSerializableDataSet(extractedDS, DerivedPolicy{}));
        }
        // clear our dataset.
        *block = vtkm::cont::DataSet();
@ -108,36 +164,43 @@ public:
    else
    {
      vtkm::Id numValues = 0;
-      std::vector<vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>>> received_arrays;
+      std::vector<vtkm::cont::DataSet> receives;
      for (int cc = 0; cc < rp.in_link().size(); ++cc)
      {
        auto target = rp.in_link().target(cc);
        if (rp.incoming(target.gid).size() > 0)
        {
-          vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> incomingAH;
-          rp.dequeue(target.gid, incomingAH);
-          received_arrays.push_back(incomingAH);
-          numValues += incomingAH.GetNumberOfValues();
+          auto sds = vtkm::filter::MakeSerializableDataSet(DerivedPolicy{});
+          rp.dequeue(target.gid, sds);
+          receives.push_back(sds.DataSet);
+          numValues += receives.back().GetCoordinateSystem(0).GetData().GetNumberOfValues();
        }
      }

      *block = vtkm::cont::DataSet();
-      if (received_arrays.size() == 1)
+      if (receives.size() == 1)
      {
-        auto coords = vtkm::cont::CoordinateSystem("coords", received_arrays[0]);
-        block->AddCoordinateSystem(coords);
+        *block = receives[0];
      }
-      else if (received_arrays.size() > 1)
+      else if (receives.size() > 1)
      {
-        vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> coordsAH;
-        coordsAH.Allocate(numValues);
-        vtkm::Id offset = 0;
-        for (const vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>>& receivedAH : received_arrays)
+        ConcatenateFields concatCoords(numValues);
+        for (const auto& ds: receives)
        {
-          vtkm::cont::Algorithm::CopySubRange(receivedAH, 0, receivedAH.GetNumberOfValues(), coordsAH, offset);
-          offset += receivedAH.GetNumberOfValues();
+          concatCoords.Append(ds.GetCoordinateSystem(0));
+        }
+        block->AddCoordinateSystem(vtkm::cont::CoordinateSystem(
+          concatCoords.GetResult().GetName(), concatCoords.GetResult().GetData()));
+
+        for (vtkm::IdComponent i = 0; i < receives[0].GetNumberOfFields(); ++i)
+        {
+          ConcatenateFields concatField(numValues);
+          for (const auto& ds: receives)
+          {
+            concatField.Append(ds.GetField(i));
+          }
+          block->AddField(concatField.GetResult());
        }
-        block->AddCoordinateSystem(vtkm::cont::CoordinateSystem("coords", coordsAH));
      }
    }
  }
--- a/examples/temporal_advection/TemporalAdvection.cxx
+++ b/examples/temporal_advection/TemporalAdvection.cxx
@ -18,10 +18,6 @@
 //  this software.
 //============================================================================

-#ifndef VTKM_DEVICE_ADAPTER
-#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_SERIAL
-#endif
-
 #include <vtkm/cont/DataSet.h>
 #include <vtkm/worklet/ParticleAdvection.h>
 #include <vtkm/worklet/particleadvection/Integrators.h>
@ -49,11 +45,8 @@ int renderAndWriteDataSet(const vtkm::cont::DataSet& dataset)

 void RunTest(vtkm::Id numSteps, vtkm::Float32 stepSize, vtkm::Id advectType)
 {
-  using DeviceAdapter = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
  using FieldType = vtkm::Float32;
  using FieldHandle = vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3>>;
-  using FieldPortalConstType =
-    typename FieldHandle::template ExecutionTypes<DeviceAdapter>::PortalConst;

  // These lines read two datasets, which are BOVs.
  // Currently VTKm does not support providing time series datasets
@ -73,11 +66,8 @@ void RunTest(vtkm::Id numSteps, vtkm::Float32 stepSize, vtkm::Id advectType)

  // The only change in this example and the vanilla particle advection example is
  // this example makes use of the TemporalGridEvaluator.
-  using GridEvaluator =
-    vtkm::worklet::particleadvection::TemporalGridEvaluator<FieldPortalConstType,
-                                                            FieldType,
-                                                            DeviceAdapter>;
-  using Integrator = vtkm::worklet::particleadvection::EulerIntegrator<GridEvaluator, FieldType>;
+  using GridEvaluator = vtkm::worklet::particleadvection::TemporalGridEvaluator<FieldHandle>;
+  using Integrator = vtkm::worklet::particleadvection::EulerIntegrator<GridEvaluator>;

  GridEvaluator eval(ds1.GetCoordinateSystem(),
                     ds1.GetCellSet(0),
@ -109,13 +99,12 @@ void RunTest(vtkm::Id numSteps, vtkm::Float32 stepSize, vtkm::Id advectType)
  if (advectType == 0)
  {
    vtkm::worklet::ParticleAdvection particleAdvection;
-    particleAdvection.Run(integrator, seedArray, numSteps, DeviceAdapter());
+    particleAdvection.Run(integrator, seedArray, numSteps);
  }
  else
  {
    vtkm::worklet::Streamline streamline;
-    vtkm::worklet::StreamlineResult<FieldType> res =
-      streamline.Run(integrator, seedArray, numSteps, DeviceAdapter());
+    vtkm::worklet::StreamlineResult res = streamline.Run(integrator, seedArray, numSteps);
    vtkm::cont::DataSet outData;
    vtkm::cont::CoordinateSystem outputCoords("coordinates", res.positions);
    outData.AddCellSet(res.polyLines);
--- a/version.txt
+++ b/version.txt
@ -1 +1 @@
-1.2.0
+1.3.0
--- a/vtkm/Bitset.h
+++ b/vtkm/Bitset.h
@ -0,0 +1,70 @@
+//============================================================================
+//  Copyright (c) Kitware, Inc.
+//  All rights reserved.
+//  See LICENSE.txt for details.
+//  This software is distributed WITHOUT ANY WARRANTY; without even
+//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+//  PURPOSE.  See the above copyright notice for more information.
+//
+//  Copyright 2016 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
+//  Copyright 2016 UT-Battelle, LLC.
+//  Copyright 2016 Los Alamos National Security.
+//
+//  Under the terms of Contract DE-NA0003525 with NTESS,
+//  the U.S. Government retains certain rights in this software.
+//
+//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+//  Laboratory (LANL), the U.S. Government retains certain rights in
+//  this software.
+//============================================================================
+
+#ifndef vtk_m_Bitset_h
+#define vtk_m_Bitset_h
+
+#include <assert.h>
+#include <limits>
+#include <vtkm/Types.h>
+#include <vtkm/internal/ExportMacros.h>
+
+namespace vtkm
+{
+
+/// \brief A bitmap to serve different needs.
+/// Ex. Editing particular bits in a byte(s), checkint if particular bit values
+/// are present or not. Once Cuda supports std::bitset, we should use the
+/// standard one if possible
+template <typename MaskType>
+struct Bitset
+{
+  VTKM_EXEC_CONT Bitset()
+    : Mask(0)
+  {
+  }
+
+  VTKM_EXEC_CONT void set(vtkm::Id bitIndex)
+  {
+    this->Mask = this->Mask | (static_cast<MaskType>(1) << bitIndex);
+  }
+
+  VTKM_EXEC_CONT void reset(vtkm::Id bitIndex)
+  {
+    this->Mask = this->Mask & ~(static_cast<MaskType>(1) << bitIndex);
+  }
+
+  VTKM_EXEC_CONT void toggle(vtkm::Id bitIndex)
+  {
+    this->Mask = this->Mask ^ (static_cast<MaskType>(0) << bitIndex);
+  }
+
+  VTKM_EXEC_CONT bool test(vtkm::Id bitIndex)
+  {
+    return ((this->Mask & (static_cast<MaskType>(1) << bitIndex)) != 0);
+  }
+
+private:
+  MaskType Mask;
+};
+
+} // namespace vtkm
+
+#endif //vtk_m_Bitset_h
--- a/vtkm/CMakeLists.txt
+++ b/vtkm/CMakeLists.txt
@ -31,11 +31,13 @@ set(headers
  BaseComponent.h
  BinaryPredicates.h
  BinaryOperators.h
+  Bitset.h
  Bounds.h
  CellShape.h
  CellTraits.h
  Flags.h
  Geometry.h
+  GhostCell.h
  Hash.h
  ImplicitFunction.h
  ListTag.h
@ -80,6 +82,9 @@ vtkm_declare_headers(
 #first add all the components vtkm that are shared between control and exec
 add_subdirectory(thirdparty/diy)
 add_subdirectory(thirdparty/taotuple)
+if(VTKm_ENABLE_LOGGING)
+  add_subdirectory(thirdparty/loguru)
+endif()
 add_subdirectory(testing)
 add_subdirectory(internal)

--- a/examples/cosmotools/CosmoCenterFinderTBB.cxx
+++ b/examples/cosmotools/CosmoCenterFinderTBB.cxx
@ -17,7 +17,23 @@
 //  Laboratory (LANL), the U.S. Government retains certain rights in
 //  this software.
 //============================================================================
+#ifndef vtk_m_Ghost_Cell_h
+#define vtk_m_Ghost_Cell_h

-#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_TBB
+namespace vtkm
+{

-#include "CosmoCenterFinder.cxx"
+enum struct CellClassification
+{
+  NORMAL = 0,         //Valid cell
+  DUPLICATE = 1 << 0, //Ghost cell
+  INVALID = 1 << 1,   //Cell is invalid
+  UNUSED0 = 1 << 2,
+  UNUSED1 = 1 << 3,
+  UNUSED3 = 1 << 4,
+  UNUSED4 = 1 << 5,
+  UNUSED5 = 1 << 6,
+};
+}
+
+#endif // vtk_m_Ghost_Cell_h
--- a/vtkm/Hash.h
+++ b/vtkm/Hash.h
@ -77,7 +77,7 @@ VTKM_EXEC_CONT inline vtkm::HashType HashFNV1a64(const InVecType& inVec)

 // If you get a compile error saying that there is no implementation of the class HashChooser,
 // then you have tried to make a hash from an invalid type (like a float).
-template <typename NumericTag, vtkm::Id DataSize>
+template <typename NumericTag, std::size_t DataSize>
 struct HashChooser;

 template <>
--- a/vtkm/TypeListTag.h
+++ b/vtkm/TypeListTag.h
@ -32,46 +32,46 @@ namespace vtkm

 /// A list containing the type vtkm::Id.
 ///
-struct TypeListTagId : vtkm::ListTagBase<vtkm::Id>
+struct VTKM_ALWAYS_EXPORT TypeListTagId : vtkm::ListTagBase<vtkm::Id>
 {
 };

 /// A list containing the type vtkm::Id2.
 ///
-struct TypeListTagId2 : vtkm::ListTagBase<vtkm::Id2>
+struct VTKM_ALWAYS_EXPORT TypeListTagId2 : vtkm::ListTagBase<vtkm::Id2>
 {
 };

 /// A list containing the type vtkm::Id3.
 ///
-struct TypeListTagId3 : vtkm::ListTagBase<vtkm::Id3>
+struct VTKM_ALWAYS_EXPORT TypeListTagId3 : vtkm::ListTagBase<vtkm::Id3>
 {
 };

 /// A list containing the type vtkm::IdComponent
 ///
-struct TypeListTagIdComponent : vtkm::ListTagBase<vtkm::IdComponent>
+struct VTKM_ALWAYS_EXPORT TypeListTagIdComponent : vtkm::ListTagBase<vtkm::IdComponent>
 {
 };

 /// A list containing types used to index arrays. Contains vtkm::Id, vtkm::Id2,
 /// and vtkm::Id3.
 ///
-struct TypeListTagIndex : vtkm::ListTagBase<vtkm::Id, vtkm::Id2, vtkm::Id3>
+struct VTKM_ALWAYS_EXPORT TypeListTagIndex : vtkm::ListTagBase<vtkm::Id, vtkm::Id2, vtkm::Id3>
 {
 };

 /// A list containing types used for scalar fields. Specifically, contains
 /// floating point numbers of different widths (i.e. vtkm::Float32 and
 /// vtkm::Float64).
-struct TypeListTagFieldScalar : vtkm::ListTagBase<vtkm::Float32, vtkm::Float64>
+struct VTKM_ALWAYS_EXPORT TypeListTagFieldScalar : vtkm::ListTagBase<vtkm::Float32, vtkm::Float64>
 {
 };

 /// A list containing types for values for fields with two dimensional
 /// vectors.
 ///
-struct TypeListTagFieldVec2
+struct VTKM_ALWAYS_EXPORT TypeListTagFieldVec2
  : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 2>, vtkm::Vec<vtkm::Float64, 2>>
 {
 };
@ -79,7 +79,7 @@ struct TypeListTagFieldVec2
 /// A list containing types for values for fields with three dimensional
 /// vectors.
 ///
-struct TypeListTagFieldVec3
+struct VTKM_ALWAYS_EXPORT TypeListTagFieldVec3
  : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 3>, vtkm::Vec<vtkm::Float64, 3>>
 {
 };
@ -87,7 +87,7 @@ struct TypeListTagFieldVec3
 /// A list containing types for values for fields with four dimensional
 /// vectors.
 ///
-struct TypeListTagFieldVec4
+struct VTKM_ALWAYS_EXPORT TypeListTagFieldVec4
  : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 4>, vtkm::Vec<vtkm::Float64, 4>>
 {
 };
@ -96,12 +96,13 @@ struct TypeListTagFieldVec4
 /// floating point vectors of size 2, 3, and 4 with floating point components.
 /// Scalars are not included.
 ///
-struct TypeListTagFloatVec : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 2>,
-                                               vtkm::Vec<vtkm::Float64, 2>,
-                                               vtkm::Vec<vtkm::Float32, 3>,
-                                               vtkm::Vec<vtkm::Float64, 3>,
-                                               vtkm::Vec<vtkm::Float32, 4>,
-                                               vtkm::Vec<vtkm::Float64, 4>>
+struct VTKM_ALWAYS_EXPORT TypeListTagFloatVec
+  : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 2>,
+                      vtkm::Vec<vtkm::Float64, 2>,
+                      vtkm::Vec<vtkm::Float32, 3>,
+                      vtkm::Vec<vtkm::Float64, 3>,
+                      vtkm::Vec<vtkm::Float32, 4>,
+                      vtkm::Vec<vtkm::Float64, 4>>
 {
 };

@ -109,30 +110,32 @@ struct TypeListTagFloatVec : vtkm::ListTagBase<vtkm::Vec<vtkm::Float32, 2>,
 /// floating point scalars and vectors of size 2, 3, and 4 with floating point
 /// components.
 ///
-struct TypeListTagField : vtkm::ListTagBase<vtkm::Float32,
-                                            vtkm::Float64,
-                                            vtkm::Vec<vtkm::Float32, 2>,
-                                            vtkm::Vec<vtkm::Float64, 2>,
-                                            vtkm::Vec<vtkm::Float32, 3>,
-                                            vtkm::Vec<vtkm::Float64, 3>,
-                                            vtkm::Vec<vtkm::Float32, 4>,
-                                            vtkm::Vec<vtkm::Float64, 4>>
+struct VTKM_ALWAYS_EXPORT TypeListTagField
+  : vtkm::ListTagBase<vtkm::Float32,
+                      vtkm::Float64,
+                      vtkm::Vec<vtkm::Float32, 2>,
+                      vtkm::Vec<vtkm::Float64, 2>,
+                      vtkm::Vec<vtkm::Float32, 3>,
+                      vtkm::Vec<vtkm::Float64, 3>,
+                      vtkm::Vec<vtkm::Float32, 4>,
+                      vtkm::Vec<vtkm::Float64, 4>>
 {
 };

 /// A list of all scalars defined in vtkm/Types.h. A scalar is a type that
 /// holds a single number.
 ///
-struct TypeListTagScalarAll : vtkm::ListTagBase<vtkm::Int8,
-                                                vtkm::UInt8,
-                                                vtkm::Int16,
-                                                vtkm::UInt16,
-                                                vtkm::Int32,
-                                                vtkm::UInt32,
-                                                vtkm::Int64,
-                                                vtkm::UInt64,
-                                                vtkm::Float32,
-                                                vtkm::Float64>
+struct VTKM_ALWAYS_EXPORT TypeListTagScalarAll
+  : vtkm::ListTagBase<vtkm::Int8,
+                      vtkm::UInt8,
+                      vtkm::Int16,
+                      vtkm::UInt16,
+                      vtkm::Int32,
+                      vtkm::UInt32,
+                      vtkm::Int64,
+                      vtkm::UInt64,
+                      vtkm::Float32,
+                      vtkm::Float64>
 {
 };

@ -140,21 +143,22 @@ struct TypeListTagScalarAll : vtkm::ListTagBase<vtkm::Int8,
 /// vectors of size 2, 3, or 4 containing either unsigned bytes, signed
 /// integers of 32 or 64 bits, or floating point values of 32 or 64 bits.
 ///
-struct TypeListTagVecCommon : vtkm::ListTagBase<vtkm::Vec<vtkm::UInt8, 2>,
-                                                vtkm::Vec<vtkm::Int32, 2>,
-                                                vtkm::Vec<vtkm::Int64, 2>,
-                                                vtkm::Vec<vtkm::Float32, 2>,
-                                                vtkm::Vec<vtkm::Float64, 2>,
-                                                vtkm::Vec<vtkm::UInt8, 3>,
-                                                vtkm::Vec<vtkm::Int32, 3>,
-                                                vtkm::Vec<vtkm::Int64, 3>,
-                                                vtkm::Vec<vtkm::Float32, 3>,
-                                                vtkm::Vec<vtkm::Float64, 3>,
-                                                vtkm::Vec<vtkm::UInt8, 4>,
-                                                vtkm::Vec<vtkm::Int32, 4>,
-                                                vtkm::Vec<vtkm::Int64, 4>,
-                                                vtkm::Vec<vtkm::Float32, 4>,
-                                                vtkm::Vec<vtkm::Float64, 4>>
+struct VTKM_ALWAYS_EXPORT TypeListTagVecCommon
+  : vtkm::ListTagBase<vtkm::Vec<vtkm::UInt8, 2>,
+                      vtkm::Vec<vtkm::Int32, 2>,
+                      vtkm::Vec<vtkm::Int64, 2>,
+                      vtkm::Vec<vtkm::Float32, 2>,
+                      vtkm::Vec<vtkm::Float64, 2>,
+                      vtkm::Vec<vtkm::UInt8, 3>,
+                      vtkm::Vec<vtkm::Int32, 3>,
+                      vtkm::Vec<vtkm::Int64, 3>,
+                      vtkm::Vec<vtkm::Float32, 3>,
+                      vtkm::Vec<vtkm::Float64, 3>,
+                      vtkm::Vec<vtkm::UInt8, 4>,
+                      vtkm::Vec<vtkm::Int32, 4>,
+                      vtkm::Vec<vtkm::Int64, 4>,
+                      vtkm::Vec<vtkm::Float32, 4>,
+                      vtkm::Vec<vtkm::Float64, 4>>
 {
 };

@ -165,21 +169,22 @@ namespace internal
 /// use in general, but is used when joined with \c TypeListTagVecCommon
 /// to get a list of all vectors up to size 4.
 ///
-struct TypeListTagVecUncommon : vtkm::ListTagBase<vtkm::Vec<vtkm::Int8, 2>,
-                                                  vtkm::Vec<vtkm::Int16, 2>,
-                                                  vtkm::Vec<vtkm::UInt16, 2>,
-                                                  vtkm::Vec<vtkm::UInt32, 2>,
-                                                  vtkm::Vec<vtkm::UInt64, 2>,
-                                                  vtkm::Vec<vtkm::Int8, 3>,
-                                                  vtkm::Vec<vtkm::Int16, 3>,
-                                                  vtkm::Vec<vtkm::UInt16, 3>,
-                                                  vtkm::Vec<vtkm::UInt32, 3>,
-                                                  vtkm::Vec<vtkm::UInt64, 3>,
-                                                  vtkm::Vec<vtkm::Int8, 4>,
-                                                  vtkm::Vec<vtkm::Int16, 4>,
-                                                  vtkm::Vec<vtkm::UInt16, 4>,
-                                                  vtkm::Vec<vtkm::UInt32, 4>,
-                                                  vtkm::Vec<vtkm::UInt64, 4>>
+struct VTKM_ALWAYS_EXPORT TypeListTagVecUncommon
+  : vtkm::ListTagBase<vtkm::Vec<vtkm::Int8, 2>,
+                      vtkm::Vec<vtkm::Int16, 2>,
+                      vtkm::Vec<vtkm::UInt16, 2>,
+                      vtkm::Vec<vtkm::UInt32, 2>,
+                      vtkm::Vec<vtkm::UInt64, 2>,
+                      vtkm::Vec<vtkm::Int8, 3>,
+                      vtkm::Vec<vtkm::Int16, 3>,
+                      vtkm::Vec<vtkm::UInt16, 3>,
+                      vtkm::Vec<vtkm::UInt32, 3>,
+                      vtkm::Vec<vtkm::UInt64, 3>,
+                      vtkm::Vec<vtkm::Int8, 4>,
+                      vtkm::Vec<vtkm::Int16, 4>,
+                      vtkm::Vec<vtkm::UInt16, 4>,
+                      vtkm::Vec<vtkm::UInt32, 4>,
+                      vtkm::Vec<vtkm::UInt64, 4>>
 {
 };

@ -188,7 +193,7 @@ struct TypeListTagVecUncommon : vtkm::ListTagBase<vtkm::Vec<vtkm::Int8, 2>,
 /// A list of all vector classes with standard types as components and
 /// lengths between 2 and 4.
 ///
-struct TypeListTagVecAll
+struct VTKM_ALWAYS_EXPORT TypeListTagVecAll
  : vtkm::ListTagJoin<vtkm::TypeListTagVecCommon, vtkm::internal::TypeListTagVecUncommon>
 {
 };
@ -197,19 +202,21 @@ struct TypeListTagVecAll
 /// possible VTK-m types like arbitrarily typed and sized Vecs (only up to
 /// length 4) or math types like matrices.
 ///
-struct TypeListTagAll : vtkm::ListTagJoin<vtkm::TypeListTagScalarAll, vtkm::TypeListTagVecAll>
+struct VTKM_ALWAYS_EXPORT TypeListTagAll
+  : vtkm::ListTagJoin<vtkm::TypeListTagScalarAll, vtkm::TypeListTagVecAll>
 {
 };

 /// A list of the most commonly used types across multiple domains. Includes
 /// integers, floating points, and 3 dimensional vectors of floating points.
 ///
-struct TypeListTagCommon : vtkm::ListTagBase<vtkm::Int32,
-                                             vtkm::Int64,
-                                             vtkm::Float32,
-                                             vtkm::Float64,
-                                             vtkm::Vec<vtkm::Float32, 3>,
-                                             vtkm::Vec<vtkm::Float64, 3>>
+struct VTKM_ALWAYS_EXPORT TypeListTagCommon
+  : vtkm::ListTagBase<vtkm::Int32,
+                      vtkm::Int64,
+                      vtkm::Float32,
+                      vtkm::Float64,
+                      vtkm::Vec<vtkm::Float32, 3>,
+                      vtkm::Vec<vtkm::Float64, 3>>
 {
 };

@ -217,7 +224,7 @@ struct TypeListTagCommon : vtkm::ListTagBase<vtkm::Int32,
 // true. Although TypeListTagAll is necessarily finite, the point is to
 // be all inclusive. Besides, this should speed up the compilation when
 // checking a list that should contain everything.
-template <typename Type>
+template<typename Type>
 struct ListContains<vtkm::TypeListTagAll, Type>
 {
  static constexpr bool value = true;
--- a/vtkm/cont/Algorithm.h
+++ b/vtkm/cont/Algorithm.h
@ -38,24 +38,25 @@ template <typename Device, typename T>
 inline auto DoPrepareArgForExec(T&& object, std::true_type)
  -> decltype(std::declval<T>().PrepareForExecution(Device()))
 {
+  VTKM_IS_EXECUTION_OBJECT(T);
  return object.PrepareForExecution(Device{});
 }

 template <typename Device, typename T>
 inline T&& DoPrepareArgForExec(T&& object, std::false_type)
 {
+  static_assert(!vtkm::cont::internal::IsExecutionObjectBase<T>::value,
+                "Internal error: failed to detect execution object.");
  return std::forward<T>(object);
 }

 template <typename Device, typename T>
-auto PrepareArgForExec(T&& object) -> decltype(DoPrepareArgForExec<Device>(
-  std::forward<T>(object),
-  typename std::is_base_of<vtkm::cont::ExecutionObjectBase, typename std::decay<T>::type>::type{}))
+auto PrepareArgForExec(T&& object)
+  -> decltype(DoPrepareArgForExec<Device>(std::forward<T>(object),
+                                          vtkm::cont::internal::IsExecutionObjectBase<T>{}))
 {
-  return DoPrepareArgForExec<Device>(
-    std::forward<T>(object),
-    typename std::is_base_of<vtkm::cont::ExecutionObjectBase,
-                             typename std::decay<T>::type>::type{});
+  return DoPrepareArgForExec<Device>(std::forward<T>(object),
+                                     vtkm::cont::internal::IsExecutionObjectBase<T>{});
 }

 struct CopyFunctor
@ -87,6 +88,11 @@ struct CopySubRangeFunctor
 {
  bool valid;

+  CopySubRangeFunctor()
+    : valid(false)
+  {
+  }
+
  template <typename Device, typename... Args>
  VTKM_CONT bool operator()(Device, Args&&... args)
  {
@ -181,6 +187,54 @@ struct StreamingScanExclusiveFunctor
      vtkm::cont::DeviceAdapterAlgorithm<Device>::StreamingScanExclusive(numBlocks, input, output);
    return true;
  }
+
+  template <typename Device, class CIn, class COut, class BinaryFunctor>
+  VTKM_CONT bool operator()(Device,
+                            const vtkm::Id numBlocks,
+                            const vtkm::cont::ArrayHandle<T, CIn>& input,
+                            vtkm::cont::ArrayHandle<T, COut>& output,
+                            BinaryFunctor binary_functor,
+                            const T& initialValue)
+  {
+    VTKM_IS_DEVICE_ADAPTER_TAG(Device);
+    result = vtkm::cont::DeviceAdapterAlgorithm<Device>::StreamingScanExclusive(
+      numBlocks, input, output, binary_functor, initialValue);
+    return true;
+  }
+};
+
+template <typename U>
+struct StreamingReduceFunctor
+{
+  U result;
+  StreamingReduceFunctor()
+    : result(U(0))
+  {
+  }
+  template <typename Device, typename T, class CIn>
+  VTKM_CONT bool operator()(Device,
+                            const vtkm::Id numBlocks,
+                            const vtkm::cont::ArrayHandle<T, CIn>& input,
+                            U initialValue)
+  {
+    VTKM_IS_DEVICE_ADAPTER_TAG(Device);
+    result =
+      vtkm::cont::DeviceAdapterAlgorithm<Device>::StreamingReduce(numBlocks, input, initialValue);
+    return true;
+  }
+
+  template <typename Device, typename T, class CIn, class BinaryFunctor>
+  VTKM_CONT bool operator()(Device,
+                            const vtkm::Id numBlocks,
+                            const vtkm::cont::ArrayHandle<T, CIn>& input,
+                            U InitialValue,
+                            BinaryFunctor binaryFunctor)
+  {
+    VTKM_IS_DEVICE_ADAPTER_TAG(Device);
+    result = vtkm::cont::DeviceAdapterAlgorithm<Device>::StreamingReduce(
+      numBlocks, input, InitialValue, binaryFunctor);
+    return true;
+  }
 };

 struct ScanInclusiveByKeyFunctor
@ -201,8 +255,9 @@ template <typename T>
 struct ScanExclusiveFunctor
 {
  T result;
+
  ScanExclusiveFunctor()
-    : result(T(0))
+    : result(T())
  {
  }

@ -552,6 +607,7 @@ struct Algorithm
    vtkm::cont::TryExecuteOnDevice(devId, functor, numBlocks, input, output);
    return functor.result;
  }
+
  template <typename T, class CIn, class COut>
  VTKM_CONT static T StreamingScanExclusive(const vtkm::Id numBlocks,
                                            const vtkm::cont::ArrayHandle<T, CIn>& input,
@ -560,6 +616,46 @@ struct Algorithm
    return StreamingScanExclusive(vtkm::cont::DeviceAdapterTagAny(), numBlocks, input, output);
  }

+  template <typename T, class CIn, class COut, class BinaryFunctor>
+  VTKM_CONT static T StreamingScanExclusive(const vtkm::Id numBlocks,
+                                            const vtkm::cont::ArrayHandle<T, CIn>& input,
+                                            vtkm::cont::ArrayHandle<T, COut>& output,
+                                            BinaryFunctor binary_functor,
+                                            const T& initialValue)
+  {
+    detail::StreamingScanExclusiveFunctor<T> functor;
+    vtkm::cont::TryExecuteOnDevice(vtkm::cont::DeviceAdapterTagAny(),
+                                   functor,
+                                   numBlocks,
+                                   input,
+                                   output,
+                                   binary_functor,
+                                   initialValue);
+    return functor.result;
+  }
+
+  template <typename T, typename U, class CIn>
+  VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks,
+                                     const vtkm::cont::ArrayHandle<T, CIn>& input,
+                                     U initialValue)
+  {
+    detail::StreamingReduceFunctor<U> functor;
+    vtkm::cont::TryExecuteOnDevice(
+      vtkm::cont::DeviceAdapterTagAny(), functor, numBlocks, input, initialValue);
+    return functor.result;
+  }
+
+  template <typename T, typename U, class CIn, class BinaryFunctor>
+  VTKM_CONT static U StreamingReduce(const vtkm::Id numBlocks,
+                                     const vtkm::cont::ArrayHandle<T, CIn>& input,
+                                     U initialValue,
+                                     BinaryFunctor binaryFunctor)
+  {
+    detail::StreamingReduceFunctor<U> functor;
+    vtkm::cont::TryExecuteOnDevice(
+      vtkm::cont::DeviceAdapterTagAny(), functor, numBlocks, input, initialValue, binaryFunctor);
+    return functor.result;
+  }

  template <typename T, class CIn, class COut, class BinaryFunctor>
  VTKM_CONT static T ScanInclusive(vtkm::cont::DeviceAdapterId devId,
--- a/vtkm/cont/ArrayHandle.h
+++ b/vtkm/cont/ArrayHandle.h
@ -29,6 +29,7 @@
 #include <vtkm/cont/ArrayPortalToIterators.h>
 #include <vtkm/cont/ErrorBadValue.h>
 #include <vtkm/cont/ErrorInternal.h>
+#include <vtkm/cont/Serialization.h>
 #include <vtkm/cont/Storage.h>
 #include <vtkm/cont/StorageBasic.h>

@ -293,7 +294,7 @@ public:
  /// with CUDA), then the automatically generated move constructor could be
  /// created for all devices, and it would not be valid for all devices.
  ///
-  ArrayHandle(vtkm::cont::ArrayHandle<ValueType, StorageTag>&& src);
+  ArrayHandle(vtkm::cont::ArrayHandle<ValueType, StorageTag>&& src) noexcept;

  /// Special constructor for subclass specializations that need to set the
  /// initial state of the control array. When this constructor is used, it
@ -306,7 +307,7 @@ public:
  /// initial state of the control array. When this constructor is used, it
  /// is assumed that the control array is valid.
  ///
-  ArrayHandle(StorageType&& storage);
+  ArrayHandle(StorageType&& storage) noexcept;

  /// Destructs an empty ArrayHandle.
  ///
@ -327,7 +328,7 @@ public:
  ///
  VTKM_CONT
  vtkm::cont::ArrayHandle<ValueType, StorageTag>& operator=(
-    vtkm::cont::ArrayHandle<ValueType, StorageTag>&& src);
+    vtkm::cont::ArrayHandle<ValueType, StorageTag>&& src) noexcept;

  /// Like a pointer, two \c ArrayHandles are considered equal if they point
  /// to the same location in memory.
@ -674,6 +675,50 @@ VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle(
 }
 } //namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename T>
+struct TypeString<ArrayHandle<T>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH<" + TypeString<T>::Get() + ">";
+    return name;
+  }
+};
+
+namespace internal
+{
+
+template <typename T, typename S>
+void VTKM_CONT ArrayHandleDefaultSerialization(diy::BinaryBuffer& bb,
+                                               const vtkm::cont::ArrayHandle<T, S>& obj);
+
+} // internal
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename T>
+struct Serialization<vtkm::cont::ArrayHandle<T>>
+{
+  static VTKM_CONT void save(BinaryBuffer& bb, const vtkm::cont::ArrayHandle<T>& obj)
+  {
+    vtkm::cont::internal::ArrayHandleDefaultSerialization(bb, obj);
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, vtkm::cont::ArrayHandle<T>& obj);
+};
+
+} // diy
+
 #include <vtkm/cont/ArrayHandle.hxx>
 #include <vtkm/cont/internal/ArrayHandleBasicImpl.h>
 #include <vtkm/cont/internal/ArrayExportMacros.h>
--- a/vtkm/cont/ArrayHandle.hxx
+++ b/vtkm/cont/ArrayHandle.hxx
@ -38,7 +38,7 @@ ArrayHandle<T, S>::ArrayHandle(const ArrayHandle<T, S>& src)
 }

 template <typename T, typename S>
-ArrayHandle<T, S>::ArrayHandle(ArrayHandle<T, S>&& src)
+ArrayHandle<T, S>::ArrayHandle(ArrayHandle<T, S>&& src) noexcept
  : Internals(std::move(src.Internals))
 {
 }
@ -53,7 +53,7 @@ ArrayHandle<T, S>::ArrayHandle(const typename ArrayHandle<T, S>::StorageType& st
 }

 template <typename T, typename S>
-ArrayHandle<T, S>::ArrayHandle(typename ArrayHandle<T, S>::StorageType&& storage)
+ArrayHandle<T, S>::ArrayHandle(typename ArrayHandle<T, S>::StorageType&& storage) noexcept
  : Internals(new InternalStruct)
 {
  this->Internals->ControlArray = std::move(storage);
@ -74,7 +74,7 @@ ArrayHandle<T, S>& ArrayHandle<T, S>::operator=(const ArrayHandle<T, S>& src)
 }

 template <typename T, typename S>
-ArrayHandle<T, S>& ArrayHandle<T, S>::operator=(ArrayHandle<T, S>&& src)
+ArrayHandle<T, S>& ArrayHandle<T, S>::operator=(ArrayHandle<T, S>&& src) noexcept
 {
  this->Internals = std::move(src.Internals);
  return *this;
@ -345,4 +345,49 @@ void ArrayHandle<T, S>::SyncControlArray() const
  }
 }
 }
+} // vtkm::cont
+
+namespace vtkm
+{
+namespace cont
+{
+namespace internal
+{
+
+template <typename T, typename S>
+inline void VTKM_CONT ArrayHandleDefaultSerialization(diy::BinaryBuffer& bb,
+                                                      const vtkm::cont::ArrayHandle<T, S>& obj)
+{
+  vtkm::Id count = obj.GetNumberOfValues();
+  diy::save(bb, count);
+
+  auto portal = obj.GetPortalConstControl();
+  for (vtkm::Id i = 0; i < count; ++i)
+  {
+    diy::save(bb, portal.Get(i));
+  }
 }
+}
+}
+} // vtkm::cont::internal
+
+namespace diy
+{
+
+template <typename T>
+VTKM_CONT void Serialization<vtkm::cont::ArrayHandle<T>>::load(BinaryBuffer& bb,
+                                                               vtkm::cont::ArrayHandle<T>& obj)
+{
+  vtkm::Id count = 0;
+  diy::load(bb, count);
+
+  obj.Allocate(count);
+  auto portal = obj.GetPortalControl();
+  for (vtkm::Id i = 0; i < count; ++i)
+  {
+    T val{};
+    diy::load(bb, val);
+    portal.Set(i, val);
+  }
+}
+} // diy
--- a/vtkm/cont/ArrayHandleCartesianProduct.h
+++ b/vtkm/cont/ArrayHandleCartesianProduct.h
@ -181,15 +181,16 @@ struct ArrayHandleCartesianProductTraits
  using Superclass = vtkm::cont::ArrayHandle<ValueType, Tag>;
 };

-template <typename T, typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType>
-class Storage<T, StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>>
+template <typename FirstHandleType, typename SecondHandleType, typename ThirdHandleType>
+class Storage<vtkm::Vec<typename FirstHandleType::ValueType, 3>,
+              StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>>
 {
  VTKM_IS_ARRAY_HANDLE(FirstHandleType);
  VTKM_IS_ARRAY_HANDLE(SecondHandleType);
  VTKM_IS_ARRAY_HANDLE(ThirdHandleType);

 public:
-  using ValueType = T;
+  using ValueType = vtkm::Vec<typename FirstHandleType::ValueType, 3>;

  using PortalType =
    vtkm::exec::internal::ArrayPortalCartesianProduct<ValueType,
@ -277,21 +278,22 @@ private:
  ThirdHandleType ThirdArray;
 };

-template <typename T,
-          typename FirstHandleType,
+template <typename FirstHandleType,
          typename SecondHandleType,
          typename ThirdHandleType,
          typename Device>
-class ArrayTransfer<T,
+class ArrayTransfer<vtkm::Vec<typename FirstHandleType::ValueType, 3>,
                    StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>,
                    Device>
 {
+public:
+  using ValueType = vtkm::Vec<typename FirstHandleType::ValueType, 3>;
+
+private:
  using StorageTag = StorageTagCartesianProduct<FirstHandleType, SecondHandleType, ThirdHandleType>;
-  using StorageType = vtkm::cont::internal::Storage<T, StorageTag>;
+  using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;

 public:
-  using ValueType = T;
-
  using PortalControl = typename StorageType::PortalType;
  using PortalConstControl = typename StorageType::PortalConstType;

@ -428,4 +430,74 @@ VTKM_CONT
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename AH1, typename AH2, typename AH3>
+struct TypeString<vtkm::cont::ArrayHandleCartesianProduct<AH1, AH2, AH3>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_CartesianProduct<" + TypeString<AH1>::Get() + "," +
+      TypeString<AH2>::Get() + "," + TypeString<AH3>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename AH1, typename AH2, typename AH3>
+struct TypeString<
+  vtkm::cont::ArrayHandle<vtkm::Vec<typename AH1::ValueType, 3>,
+                          vtkm::cont::internal::StorageTagCartesianProduct<AH1, AH2, AH3>>>
+  : TypeString<vtkm::cont::ArrayHandleCartesianProduct<AH1, AH2, AH3>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename AH1, typename AH2, typename AH3>
+struct Serialization<vtkm::cont::ArrayHandleCartesianProduct<AH1, AH2, AH3>>
+{
+private:
+  using Type = typename vtkm::cont::ArrayHandleCartesianProduct<AH1, AH2, AH3>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    auto storage = obj.GetStorage();
+    diy::save(bb, storage.GetFirstArray());
+    diy::save(bb, storage.GetSecondArray());
+    diy::save(bb, storage.GetThirdArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    AH1 array1;
+    AH2 array2;
+    AH3 array3;
+
+    diy::load(bb, array1);
+    diy::load(bb, array2);
+    diy::load(bb, array3);
+
+    obj = vtkm::cont::make_ArrayHandleCartesianProduct(array1, array2, array3);
+  }
+};
+
+template <typename AH1, typename AH2, typename AH3>
+struct Serialization<
+  vtkm::cont::ArrayHandle<vtkm::Vec<typename AH1::ValueType, 3>,
+                          vtkm::cont::internal::StorageTagCartesianProduct<AH1, AH2, AH3>>>
+  : Serialization<vtkm::cont::ArrayHandleCartesianProduct<AH1, AH2, AH3>>
+{
+};
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleCartesianProduct_h
--- a/vtkm/cont/ArrayHandleCast.h
+++ b/vtkm/cont/ArrayHandleCast.h
@ -103,4 +103,56 @@ VTKM_CONT
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename T1, typename T2>
+struct TypeString<vtkm::cont::internal::Cast<T1, T2>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_Cast_Functor<" + TypeString<T1>::Get() + "," + TypeString<T2>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename T, typename AH>
+struct TypeString<vtkm::cont::ArrayHandleCast<T, AH>>
+  : TypeString<
+      vtkm::cont::ArrayHandleTransform<AH,
+                                       vtkm::cont::internal::Cast<typename AH::ValueType, T>,
+                                       vtkm::cont::internal::Cast<T, typename AH::ValueType>>>
+{
+};
+}
+} // namespace vtkm::cont
+
+namespace diy
+{
+
+template <typename T1, typename T2>
+struct Serialization<vtkm::cont::internal::Cast<T1, T2>>
+{
+  static VTKM_CONT void save(BinaryBuffer&, const vtkm::cont::internal::Cast<T1, T2>&) {}
+
+  static VTKM_CONT void load(BinaryBuffer&, vtkm::cont::internal::Cast<T1, T2>&) {}
+};
+
+template <typename T, typename AH>
+struct Serialization<vtkm::cont::ArrayHandleCast<T, AH>>
+  : Serialization<
+      vtkm::cont::ArrayHandleTransform<AH,
+                                       vtkm::cont::internal::Cast<typename AH::ValueType, T>,
+                                       vtkm::cont::internal::Cast<T, typename AH::ValueType>>>
+{
+};
+
+} // diy
+
 #endif // vtk_m_cont_ArrayHandleCast_h
--- a/vtkm/cont/ArrayHandleCompositeVector.h
+++ b/vtkm/cont/ArrayHandleCompositeVector.h
@ -723,4 +723,121 @@ VTKM_CONT ArrayHandleCompositeVector<ArrayTs...> make_ArrayHandleCompositeVector
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename... AHs>
+struct TypeString<vtkm::cont::ArrayHandleCompositeVector<AHs...>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_CompositeVector<" + internal::GetVariadicTypeString(AHs{}...) + ">";
+    return name;
+  }
+};
+
+template <typename... AHs>
+struct TypeString<vtkm::cont::ArrayHandle<
+  typename vtkm::cont::internal::compvec::GetValueType<vtkmstd::tuple<AHs...>>::ValueType,
+  vtkm::cont::internal::StorageTagCompositeVector<vtkmstd::tuple<AHs...>>>>
+  : TypeString<vtkm::cont::ArrayHandleCompositeVector<AHs...>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+namespace internal
+{
+
+template <typename Functor, typename TupleType, typename... Args>
+inline void TupleForEachImpl(TupleType&& t,
+                             std::integral_constant<size_t, 0>,
+                             Functor&& f,
+                             Args&&... args)
+{
+  f(vtkmstd::get<0>(t), std::forward<Args>(args)...);
+}
+
+template <typename Functor, typename TupleType, size_t Index, typename... Args>
+inline void TupleForEachImpl(TupleType&& t,
+                             std::integral_constant<size_t, Index>,
+                             Functor&& f,
+                             Args&&... args)
+{
+  TupleForEachImpl(std::forward<TupleType>(t),
+                   std::integral_constant<size_t, Index - 1>{},
+                   std::forward<Functor>(f),
+                   std::forward<Args>(args)...);
+  f(vtkmstd::get<Index>(t), std::forward<Args>(args)...);
+}
+
+template <typename Functor, typename TupleType, typename... Args>
+inline void TupleForEach(TupleType&& t, Functor&& f, Args&&... args)
+{
+  constexpr auto size = vtkmstd::tuple_size<typename std::decay<TupleType>::type>::value;
+
+  TupleForEachImpl(std::forward<TupleType>(t),
+                   std::integral_constant<size_t, size - 1>{},
+                   std::forward<Functor>(f),
+                   std::forward<Args>(args)...);
+}
+} // internal
+
+template <typename... AHs>
+struct Serialization<vtkm::cont::ArrayHandleCompositeVector<AHs...>>
+{
+private:
+  using Type = typename vtkm::cont::ArrayHandleCompositeVector<AHs...>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+  struct SaveFunctor
+  {
+    template <typename AH>
+    void operator()(const AH& ah, BinaryBuffer& bb) const
+    {
+      diy::save(bb, ah);
+    }
+  };
+
+  struct LoadFunctor
+  {
+    template <typename AH>
+    void operator()(AH& ah, BinaryBuffer& bb) const
+    {
+      diy::load(bb, ah);
+    }
+  };
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    internal::TupleForEach(obj.GetStorage().GetArrayTuple(), SaveFunctor{}, bb);
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    vtkmstd::tuple<AHs...> arrayTuple;
+    internal::TupleForEach(arrayTuple, LoadFunctor{}, bb);
+    obj = BaseType(typename BaseType::StorageType(arrayTuple));
+  }
+};
+
+template <typename... AHs>
+struct Serialization<vtkm::cont::ArrayHandle<
+  typename vtkm::cont::internal::compvec::GetValueType<vtkmstd::tuple<AHs...>>::ValueType,
+  vtkm::cont::internal::StorageTagCompositeVector<vtkmstd::tuple<AHs...>>>>
+  : Serialization<vtkm::cont::ArrayHandleCompositeVector<AHs...>>
+{
+};
+} // diy
+
 #endif //vtk_m_ArrayHandleCompositeVector_h
--- a/vtkm/cont/ArrayHandleConcatenate.h
+++ b/vtkm/cont/ArrayHandleConcatenate.h
@ -324,4 +324,69 @@ VTKM_CONT ArrayHandleConcatenate<ArrayHandleType1, ArrayHandleType2> make_ArrayH
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename AH1, typename AH2>
+struct TypeString<vtkm::cont::ArrayHandleConcatenate<AH1, AH2>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_Concatenate<" + TypeString<AH1>::Get() + "," + TypeString<AH2>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename AH1, typename AH2>
+struct TypeString<
+  vtkm::cont::ArrayHandle<typename AH1::ValueType, vtkm::cont::StorageTagConcatenate<AH1, AH2>>>
+  : TypeString<vtkm::cont::ArrayHandleConcatenate<AH1, AH2>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename AH1, typename AH2>
+struct Serialization<vtkm::cont::ArrayHandleConcatenate<AH1, AH2>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleConcatenate<AH1, AH2>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    auto storage = obj.GetStorage();
+    diy::save(bb, storage.GetArray1());
+    diy::save(bb, storage.GetArray2());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    AH1 array1;
+    AH2 array2;
+
+    diy::load(bb, array1);
+    diy::load(bb, array2);
+
+    obj = vtkm::cont::make_ArrayHandleConcatenate(array1, array2);
+  }
+};
+
+template <typename AH1, typename AH2>
+struct Serialization<
+  vtkm::cont::ArrayHandle<typename AH1::ValueType, vtkm::cont::StorageTagConcatenate<AH1, AH2>>>
+  : Serialization<vtkm::cont::ArrayHandleConcatenate<AH1, AH2>>
+{
+};
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleConcatenate_h
--- a/vtkm/cont/ArrayHandleConstant.h
+++ b/vtkm/cont/ArrayHandleConstant.h
@ -85,4 +85,40 @@ vtkm::cont::ArrayHandleConstant<T> make_ArrayHandleConstant(T value, vtkm::Id nu
 }
 } // vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename T>
+struct TypeString<vtkm::cont::detail::ConstantFunctor<T>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_ConstantFunctor<" + TypeString<T>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename T>
+struct TypeString<vtkm::cont::ArrayHandleConstant<T>>
+  : TypeString<vtkm::cont::ArrayHandleImplicit<vtkm::cont::detail::ConstantFunctor<T>>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename T>
+struct Serialization<vtkm::cont::ArrayHandleConstant<T>>
+  : Serialization<vtkm::cont::ArrayHandleImplicit<vtkm::cont::detail::ConstantFunctor<T>>>
+{
+};
+
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleConstant_h
--- a/vtkm/cont/ArrayHandleCounting.h
+++ b/vtkm/cont/ArrayHandleCounting.h
@ -76,6 +76,12 @@ public:
    return *this;
  }

+  VTKM_EXEC_CONT
+  ValueType GetStart() const { return this->Start; }
+
+  VTKM_EXEC_CONT
+  ValueType GetStep() const { return this->Step; }
+
  VTKM_EXEC_CONT
  vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }

@ -142,4 +148,70 @@ make_ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename T>
+struct TypeString<vtkm::cont::ArrayHandleCounting<T>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_Counting<" + TypeString<T>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename T>
+struct TypeString<
+  vtkm::cont::ArrayHandle<T, typename vtkm::cont::ArrayHandleCounting<T>::StorageTag>>
+  : TypeString<vtkm::cont::ArrayHandleCounting<T>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename T>
+struct Serialization<vtkm::cont::ArrayHandleCounting<T>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleCounting<T>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    auto portal = obj.GetPortalConstControl();
+    diy::save(bb, portal.GetStart());
+    diy::save(bb, portal.GetStep());
+    diy::save(bb, portal.GetNumberOfValues());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    T start{}, step{};
+    vtkm::Id count = 0;
+
+    diy::load(bb, start);
+    diy::load(bb, step);
+    diy::load(bb, count);
+
+    obj = vtkm::cont::make_ArrayHandleCounting(start, step, count);
+  }
+};
+
+template <typename T>
+struct Serialization<
+  vtkm::cont::ArrayHandle<T, typename vtkm::cont::ArrayHandleCounting<T>::StorageTag>>
+  : Serialization<vtkm::cont::ArrayHandleCounting<T>>
+{
+};
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleCounting_h
--- a/vtkm/cont/ArrayHandleExtractComponent.h
+++ b/vtkm/cont/ArrayHandleExtractComponent.h
@ -310,4 +310,69 @@ VTKM_CONT ArrayHandleExtractComponent<ArrayHandleType> make_ArrayHandleExtractCo
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename AH>
+struct TypeString<vtkm::cont::ArrayHandleExtractComponent<AH>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_ExtractComponent<" + TypeString<AH>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename AH>
+struct TypeString<
+  vtkm::cont::ArrayHandle<typename vtkm::VecTraits<typename AH::ValueType>::ComponentType,
+                          vtkm::cont::StorageTagExtractComponent<AH>>>
+  : TypeString<vtkm::cont::ArrayHandleExtractComponent<AH>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename AH>
+struct Serialization<vtkm::cont::ArrayHandleExtractComponent<AH>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleExtractComponent<AH>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    auto storage = obj.GetStorage();
+    diy::save(bb, storage.GetComponent());
+    diy::save(bb, storage.GetArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    vtkm::IdComponent component = 0;
+    AH array;
+    diy::load(bb, component);
+    diy::load(bb, array);
+
+    obj = vtkm::cont::make_ArrayHandleExtractComponent(array, component);
+  }
+};
+
+template <typename AH>
+struct Serialization<
+  vtkm::cont::ArrayHandle<typename vtkm::VecTraits<typename AH::ValueType>::ComponentType,
+                          vtkm::cont::StorageTagExtractComponent<AH>>>
+  : Serialization<vtkm::cont::ArrayHandleExtractComponent<AH>>
+{
+};
+} // diy
+
 #endif // vtk_m_cont_ArrayHandleExtractComponent_h
--- a/vtkm/cont/ArrayHandleGroupVec.h
+++ b/vtkm/cont/ArrayHandleGroupVec.h
@ -366,4 +366,66 @@ VTKM_CONT vtkm::cont::ArrayHandleGroupVec<ArrayHandleType, NUM_COMPONENTS> make_
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename AH, vtkm::IdComponent NUM_COMPS>
+struct TypeString<vtkm::cont::ArrayHandleGroupVec<AH, NUM_COMPS>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_GroupVec<" + TypeString<AH>::Get() + "," + std::to_string(NUM_COMPS) + ">";
+    return name;
+  }
+};
+
+template <typename AH, vtkm::IdComponent NUM_COMPS>
+struct TypeString<vtkm::cont::ArrayHandle<vtkm::Vec<typename AH::ValueType, NUM_COMPS>,
+                                          vtkm::cont::internal::StorageTagGroupVec<AH, NUM_COMPS>>>
+  : TypeString<vtkm::cont::ArrayHandleGroupVec<AH, NUM_COMPS>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename AH, vtkm::IdComponent NUM_COMPS>
+struct Serialization<vtkm::cont::ArrayHandleGroupVec<AH, NUM_COMPS>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleGroupVec<AH, NUM_COMPS>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    diy::save(bb, obj.GetStorage().GetSourceArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    AH array;
+    diy::load(bb, array);
+
+    obj = vtkm::cont::make_ArrayHandleGroupVec<NUM_COMPS>(array);
+  }
+};
+
+template <typename AH, vtkm::IdComponent NUM_COMPS>
+struct Serialization<
+  vtkm::cont::ArrayHandle<vtkm::Vec<typename AH::ValueType, NUM_COMPS>,
+                          vtkm::cont::internal::StorageTagGroupVec<AH, NUM_COMPS>>>
+  : Serialization<vtkm::cont::ArrayHandleGroupVec<AH, NUM_COMPS>>
+{
+};
+
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleGroupVec_h
--- a/vtkm/cont/ArrayHandleGroupVecVariable.h
+++ b/vtkm/cont/ArrayHandleGroupVecVariable.h
@ -613,4 +613,70 @@ VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Id> ConvertNumComponentsToOffsets(
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename SAH, typename OAH>
+struct TypeString<vtkm::cont::ArrayHandleGroupVecVariable<SAH, OAH>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_GroupVecVariable<" + TypeString<SAH>::Get() + "," + TypeString<OAH>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename SAH, typename OAH>
+struct TypeString<
+  vtkm::cont::ArrayHandle<vtkm::VecFromPortal<typename SAH::PortalControl>,
+                          vtkm::cont::internal::StorageTagGroupVecVariable<SAH, OAH>>>
+  : TypeString<vtkm::cont::ArrayHandleGroupVecVariable<SAH, OAH>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename SAH, typename OAH>
+struct Serialization<vtkm::cont::ArrayHandleGroupVecVariable<SAH, OAH>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleGroupVecVariable<SAH, OAH>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    diy::save(bb, obj.GetStorage().GetSourceArray());
+    diy::save(bb, obj.GetStorage().GetOffsetsArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    SAH src;
+    OAH off;
+
+    diy::load(bb, src);
+    diy::load(bb, off);
+
+    obj = vtkm::cont::make_ArrayHandleGroupVecVariable(src, off);
+  }
+};
+
+template <typename SAH, typename OAH>
+struct Serialization<
+  vtkm::cont::ArrayHandle<vtkm::VecFromPortal<typename SAH::PortalControl>,
+                          vtkm::cont::internal::StorageTagGroupVecVariable<SAH, OAH>>>
+  : Serialization<vtkm::cont::ArrayHandleGroupVecVariable<SAH, OAH>>
+{
+};
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleGroupVecVariable_h
--- a/vtkm/cont/ArrayHandleImplicit.h
+++ b/vtkm/cont/ArrayHandleImplicit.h
@ -76,6 +76,9 @@ public:
  {
  }

+  VTKM_EXEC_CONT
+  const FunctorType& GetFunctor() const { return this->Functor; }
+
  VTKM_EXEC_CONT
  vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }

@ -142,4 +145,70 @@ VTKM_CONT vtkm::cont::ArrayHandleImplicit<FunctorType> make_ArrayHandleImplicit(
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename Functor>
+struct TypeString<vtkm::cont::ArrayHandleImplicit<Functor>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_Implicit<" + TypeString<Functor>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename Functor>
+struct TypeString<vtkm::cont::ArrayHandle<
+  typename vtkm::cont::detail::ArrayHandleImplicitTraits<Functor>::ValueType,
+  vtkm::cont::StorageTagImplicit<vtkm::cont::detail::ArrayPortalImplicit<Functor>>>>
+  : TypeString<vtkm::cont::ArrayHandleImplicit<Functor>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename Functor>
+struct Serialization<vtkm::cont::ArrayHandleImplicit<Functor>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleImplicit<Functor>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    diy::save(bb, obj.GetNumberOfValues());
+    diy::save(bb, obj.GetPortalConstControl().GetFunctor());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    vtkm::Id count = 0;
+    diy::load(bb, count);
+
+    Functor functor;
+    diy::load(bb, functor);
+
+    obj = vtkm::cont::make_ArrayHandleImplicit(functor, count);
+  }
+};
+
+template <typename Functor>
+struct Serialization<vtkm::cont::ArrayHandle<
+  typename vtkm::cont::detail::ArrayHandleImplicitTraits<Functor>::ValueType,
+  vtkm::cont::StorageTagImplicit<vtkm::cont::detail::ArrayPortalImplicit<Functor>>>>
+  : Serialization<vtkm::cont::ArrayHandleImplicit<Functor>>
+{
+};
+
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleImplicit_h
--- a/vtkm/cont/ArrayHandleIndex.h
+++ b/vtkm/cont/ArrayHandleIndex.h
@ -61,4 +61,45 @@ public:
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+
+namespace vtkm
+{
+namespace cont
+{
+
+template <>
+struct TypeString<vtkm::cont::detail::IndexFunctor>
+{
+  static VTKM_CONT const std::string Get() { return "AH_IndexFunctor"; }
+};
+
+template <>
+struct TypeString<vtkm::cont::ArrayHandleIndex>
+  : TypeString<vtkm::cont::ArrayHandleImplicit<vtkm::cont::detail::IndexFunctor>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <>
+struct Serialization<vtkm::cont::detail::IndexFunctor>
+{
+  static VTKM_CONT void save(BinaryBuffer&, const vtkm::cont::detail::IndexFunctor&) {}
+
+  static VTKM_CONT void load(BinaryBuffer&, vtkm::cont::detail::IndexFunctor&) {}
+};
+
+template <>
+struct Serialization<vtkm::cont::ArrayHandleIndex>
+  : Serialization<vtkm::cont::ArrayHandleImplicit<vtkm::cont::detail::IndexFunctor>>
+{
+};
+
+} // diy
+
 #endif //vtk_m_cont_ArrayHandleIndex_h
--- a/vtkm/cont/ArrayHandlePermutation.h
+++ b/vtkm/cont/ArrayHandlePermutation.h
@ -372,4 +372,72 @@ make_ArrayHandlePermutation(IndexArrayHandleType indexArray, ValueArrayHandleTyp
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename IdxAH, typename ValAH>
+struct TypeString<vtkm::cont::ArrayHandlePermutation<IdxAH, ValAH>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name =
+      "AH_Permutation<" + TypeString<IdxAH>::Get() + "," + TypeString<ValAH>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename IdxAH, typename ValAH>
+struct TypeString<
+  vtkm::cont::ArrayHandle<typename ValAH::ValueType,
+                          vtkm::cont::internal::StorageTagPermutation<IdxAH, ValAH>>>
+  : TypeString<vtkm::cont::ArrayHandlePermutation<IdxAH, ValAH>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename IdxAH, typename ValAH>
+struct Serialization<vtkm::cont::ArrayHandlePermutation<IdxAH, ValAH>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandlePermutation<IdxAH, ValAH>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    auto storage = obj.GetStorage();
+    diy::save(bb, storage.GetIndexArray());
+    diy::save(bb, storage.GetValueArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    IdxAH indices;
+    ValAH values;
+
+    diy::load(bb, indices);
+    diy::load(bb, values);
+
+    obj = vtkm::cont::make_ArrayHandlePermutation(indices, values);
+  }
+};
+
+template <typename IdxAH, typename ValAH>
+struct Serialization<
+  vtkm::cont::ArrayHandle<typename ValAH::ValueType,
+                          vtkm::cont::internal::StorageTagPermutation<IdxAH, ValAH>>>
+  : Serialization<vtkm::cont::ArrayHandlePermutation<IdxAH, ValAH>>
+{
+};
+
+} // diy
+
 #endif //vtk_m_cont_ArrayHandlePermutation_h
--- a/vtkm/cont/ArrayHandleReverse.h
+++ b/vtkm/cont/ArrayHandleReverse.h
@ -243,4 +243,63 @@ VTKM_CONT ArrayHandleReverse<HandleType> make_ArrayHandleReverse(const HandleTyp
 }
 } // namespace vtkm::cont

+//=============================================================================
+// Specializations of serialization related classes
+namespace vtkm
+{
+namespace cont
+{
+
+template <typename AH>
+struct TypeString<vtkm::cont::ArrayHandleReverse<AH>>
+{
+  static VTKM_CONT const std::string& Get()
+  {
+    static std::string name = "AH_Reverse<" + TypeString<AH>::Get() + ">";
+    return name;
+  }
+};
+
+template <typename AH>
+struct TypeString<
+  vtkm::cont::ArrayHandle<typename AH::ValueType, vtkm::cont::StorageTagReverse<AH>>>
+  : TypeString<vtkm::cont::ArrayHandleReverse<AH>>
+{
+};
+}
+} // vtkm::cont
+
+namespace diy
+{
+
+template <typename AH>
+struct Serialization<vtkm::cont::ArrayHandleReverse<AH>>
+{
+private:
+  using Type = vtkm::cont::ArrayHandleReverse<AH>;
+  using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
+
+public:
+  static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
+  {
+    diy::save(bb, obj.GetStorage().GetArray());
+  }
+
+  static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
+  {
+    AH array;
+    diy::load(bb, array);
+    obj = vtkm::cont::make_ArrayHandleReverse(array);
+  }
+};
+
+template <typename AH>
+struct Serialization<
+  vtkm::cont::ArrayHandle<typename AH::ValueType, vtkm::cont::StorageTagReverse<AH>>>
+  : Serialization<vtkm::cont::ArrayHandleReverse<AH>>
+{
+};
+
+} // diy
+
 #endif // vtk_m_cont_ArrayHandleReverse_h
--- a/Show More
+++ b/Show More
 @ -1 +1 @@
 .2.0
 .3.0