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

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

Browse Source
This commit is contained in:
ayenpure 2018-06-30 11:56:33 -06:00
commit e2dccee099
221 changed files with 10610 additions and 1697 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

551
CMake/FindOpenMP.cmake Normal file
View File

@ -0,0 +1,551 @@
##=============================================================================
##
## 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.
##
##=============================================================================
#.rst:
# FindOpenMP
# ----------
#
# Finds OpenMP support
#
# This module can be used to detect OpenMP support in a compiler. If
# the compiler supports OpenMP, the flags required to compile with
# OpenMP support are returned in variables for the different languages.
# The variables may be empty if the compiler does not need a special
# flag to support OpenMP.
#
# Variables
# ^^^^^^^^^
#
# The module exposes the components ``C``, ``CXX``, and ``Fortran``.
# Each of these controls the various languages to search OpenMP support for.
#
# Depending on the enabled components the following variables will be set:
#
# ``OpenMP_FOUND``
# Variable indicating that OpenMP flags for all requested languages have been found.
# If no components are specified, this is true if OpenMP settings for all enabled languages
# were detected.
# ``OpenMP_VERSION``
# Minimal version of the OpenMP standard detected among the requested languages,
# or all enabled languages if no components were specified.
#
# This module will set the following variables per language in your
# project, where ``<lang>`` is one of C, CXX, or Fortran:
#
# ``OpenMP_<lang>_FOUND``
# Variable indicating if OpenMP support for ``<lang>`` was detected.
# ``OpenMP_<lang>_FLAGS``
# OpenMP compiler flags for ``<lang>``, separated by spaces.
#
# For linking with OpenMP code written in ``<lang>``, the following
# variables are provided:
#
# ``OpenMP_<lang>_LIB_NAMES``
# :ref:`;-list <CMake Language Lists>` of libraries for OpenMP programs for ``<lang>``.
# ``OpenMP_<libname>_LIBRARY``
# Location of the individual libraries needed for OpenMP support in ``<lang>``.
# ``OpenMP_<lang>_LIBRARIES``
# A list of libraries needed to link with OpenMP code written in ``<lang>``.
#
# Additionally, the module provides :prop_tgt:`IMPORTED` targets:
#
# ``OpenMP::OpenMP_<lang>``
# Target for using OpenMP from ``<lang>``.
#
# Specifically for Fortran, the module sets the following variables:
#
# ``OpenMP_Fortran_HAVE_OMPLIB_HEADER``
# Boolean indicating if OpenMP is accessible through ``omp_lib.h``.
# ``OpenMP_Fortran_HAVE_OMPLIB_MODULE``
# Boolean indicating if OpenMP is accessible through the ``omp_lib`` Fortran module.
#
# The module will also try to provide the OpenMP version variables:
#
# ``OpenMP_<lang>_SPEC_DATE``
# Date of the OpenMP specification implemented by the ``<lang>`` compiler.
# ``OpenMP_<lang>_VERSION_MAJOR``
# Major version of OpenMP implemented by the ``<lang>`` compiler.
# ``OpenMP_<lang>_VERSION_MINOR``
# Minor version of OpenMP implemented by the ``<lang>`` compiler.
# ``OpenMP_<lang>_VERSION``
# OpenMP version implemented by the ``<lang>`` compiler.
#
# The specification date is formatted as given in the OpenMP standard:
# ``yyyymm`` where ``yyyy`` and ``mm`` represents the year and month of
# the OpenMP specification implemented by the ``<lang>`` compiler.
cmake_policy(PUSH)
cmake_policy(SET CMP0012 NEW) # if() recognizes numbers and booleans
cmake_policy(SET CMP0054 NEW) # if() quoted variables not dereferenced
cmake_policy(SET CMP0057 NEW) # if IN_LIST
function(_OPENMP_FLAG_CANDIDATES LANG)
if(NOT OpenMP_${LANG}_FLAG)
unset(OpenMP_FLAG_CANDIDATES)
set(OMP_FLAG_GNU "-fopenmp")
set(OMP_FLAG_Clang "-fopenmp=libomp" "-fopenmp=libiomp5" "-fopenmp")
set(OMP_FLAG_AppleClang "-Xclang -fopenmp")
set(OMP_FLAG_HP "+Oopenmp")
if(WIN32)
set(OMP_FLAG_Intel "-Qopenmp")
elseif(CMAKE_${LANG}_COMPILER_ID STREQUAL "Intel" AND
"${CMAKE_${LANG}_COMPILER_VERSION}" VERSION_LESS "15.0.0.20140528")
set(OMP_FLAG_Intel "-openmp")
else()
set(OMP_FLAG_Intel "-qopenmp")
endif()
set(OMP_FLAG_MIPSpro "-mp")
set(OMP_FLAG_MSVC "-openmp")
set(OMP_FLAG_PathScale "-openmp")
set(OMP_FLAG_NAG "-openmp")
set(OMP_FLAG_Absoft "-openmp")
set(OMP_FLAG_PGI "-mp")
set(OMP_FLAG_Flang "-fopenmp")
set(OMP_FLAG_SunPro "-xopenmp")
set(OMP_FLAG_XL "-qsmp=omp")
# Cray compiler activate OpenMP with -h omp, which is enabled by default.
set(OMP_FLAG_Cray " " "-h omp")
# If we know the correct flags, use those
if(DEFINED OMP_FLAG_${CMAKE_${LANG}_COMPILER_ID})
set(OpenMP_FLAG_CANDIDATES "${OMP_FLAG_${CMAKE_${LANG}_COMPILER_ID}}")
# Fall back to reasonable default tries otherwise
else()
set(OpenMP_FLAG_CANDIDATES "-openmp" "-fopenmp" "-mp" " ")
endif()
set(OpenMP_${LANG}_FLAG_CANDIDATES "${OpenMP_FLAG_CANDIDATES}" PARENT_SCOPE)
else()
set(OpenMP_${LANG}_FLAG_CANDIDATES "${OpenMP_${LANG}_FLAG}" PARENT_SCOPE)
endif()
endfunction()
# sample openmp source code to test
set(OpenMP_C_CXX_TEST_SOURCE
"
#include <omp.h>
int main() {
#ifdef _OPENMP
int n = omp_get_max_threads();
return 0;
#else
breaks_on_purpose
#endif
}
")
# in Fortran, an implementation may provide an omp_lib.h header
# or omp_lib module, or both (OpenMP standard, section 3.1)
# Furthmore !$ is the Fortran equivalent of #ifdef _OPENMP (OpenMP standard, 2.2.2)
# Without the conditional compilation, some compilers (e.g. PGI) might compile OpenMP code
# while not actually enabling OpenMP, building code sequentially
set(OpenMP_Fortran_TEST_SOURCE
"
program test
@OpenMP_Fortran_INCLUDE_LINE@
!$ integer :: n
n = omp_get_num_threads()
end program test
"
)
function(_OPENMP_WRITE_SOURCE_FILE LANG SRC_FILE_CONTENT_VAR SRC_FILE_NAME SRC_FILE_FULLPATH)
set(WORK_DIR ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/FindOpenMP)
if("${LANG}" STREQUAL "C")
set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.c")
file(WRITE "${SRC_FILE}" "${OpenMP_C_CXX_${SRC_FILE_CONTENT_VAR}}")
elseif("${LANG}" STREQUAL "CXX")
set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.cpp")
file(WRITE "${SRC_FILE}" "${OpenMP_C_CXX_${SRC_FILE_CONTENT_VAR}}")
elseif("${LANG}" STREQUAL "Fortran")
set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.f90")
file(WRITE "${SRC_FILE}_in" "${OpenMP_Fortran_${SRC_FILE_CONTENT_VAR}}")
configure_file("${SRC_FILE}_in" "${SRC_FILE}" @ONLY)
endif()
set(${SRC_FILE_FULLPATH} "${SRC_FILE}" PARENT_SCOPE)
endfunction()
include(CMakeParseImplicitLinkInfo)
function(_OPENMP_GET_FLAGS LANG FLAG_MODE OPENMP_FLAG_VAR OPENMP_LIB_NAMES_VAR)
_OPENMP_FLAG_CANDIDATES("${LANG}")
_OPENMP_WRITE_SOURCE_FILE("${LANG}" "TEST_SOURCE" OpenMPTryFlag _OPENMP_TEST_SRC)
unset(OpenMP_VERBOSE_COMPILE_OPTIONS)
separate_arguments(OpenMP_VERBOSE_OPTIONS NATIVE_COMMAND "${CMAKE_${LANG}_VERBOSE_FLAG}")
foreach(_VERBOSE_OPTION IN LISTS OpenMP_VERBOSE_OPTIONS)
if(NOT _VERBOSE_OPTION MATCHES "^-Wl,")
list(APPEND OpenMP_VERBOSE_COMPILE_OPTIONS ${_VERBOSE_OPTION})
endif()
endforeach()
foreach(OPENMP_FLAG IN LISTS OpenMP_${LANG}_FLAG_CANDIDATES)
set(OPENMP_FLAGS_TEST "${OPENMP_FLAG}")
if(OpenMP_VERBOSE_COMPILE_OPTIONS)
string(APPEND OPENMP_FLAGS_TEST " ${OpenMP_VERBOSE_COMPILE_OPTIONS}")
endif()
string(REGEX REPLACE "[-/=+]" "" OPENMP_PLAIN_FLAG "${OPENMP_FLAG}")
try_compile( OpenMP_COMPILE_RESULT_${FLAG_MODE}_${OPENMP_PLAIN_FLAG} ${CMAKE_BINARY_DIR} ${_OPENMP_TEST_SRC}
CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OPENMP_FLAGS_TEST}"
LINK_LIBRARIES ${CMAKE_${LANG}_VERBOSE_FLAG}
OUTPUT_VARIABLE OpenMP_TRY_COMPILE_OUTPUT
)
if(OpenMP_COMPILE_RESULT_${FLAG_MODE}_${OPENMP_PLAIN_FLAG})
set("${OPENMP_FLAG_VAR}" "${OPENMP_FLAG}" PARENT_SCOPE)
if(CMAKE_${LANG}_VERBOSE_FLAG)
unset(OpenMP_${LANG}_IMPLICIT_LIBRARIES)
unset(OpenMP_${LANG}_IMPLICIT_LINK_DIRS)
unset(OpenMP_${LANG}_IMPLICIT_FWK_DIRS)
unset(OpenMP_${LANG}_LOG_VAR)
file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
"Detecting ${LANG} OpenMP compiler ABI info compiled with the following output:\n${OpenMP_TRY_COMPILE_OUTPUT}\n\n")
cmake_parse_implicit_link_info("${OpenMP_TRY_COMPILE_OUTPUT}"
OpenMP_${LANG}_IMPLICIT_LIBRARIES
OpenMP_${LANG}_IMPLICIT_LINK_DIRS
OpenMP_${LANG}_IMPLICIT_FWK_DIRS
OpenMP_${LANG}_LOG_VAR
"${CMAKE_${LANG}_IMPLICIT_OBJECT_REGEX}"
)
file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
"Parsed ${LANG} OpenMP implicit link information from above output:\n${OpenMP_${LANG}_LOG_VAR}\n\n")
unset(_OPENMP_LIB_NAMES)
foreach(_OPENMP_IMPLICIT_LIB IN LISTS OpenMP_${LANG}_IMPLICIT_LIBRARIES)
get_filename_component(_OPENMP_IMPLICIT_LIB_DIR "${_OPENMP_IMPLICIT_LIB}" DIRECTORY)
get_filename_component(_OPENMP_IMPLICIT_LIB_NAME "${_OPENMP_IMPLICIT_LIB}" NAME)
get_filename_component(_OPENMP_IMPLICIT_LIB_PLAIN "${_OPENMP_IMPLICIT_LIB}" NAME_WE)
string(REGEX REPLACE "([][+.*?()^$])" "\\\\\\1" _OPENMP_IMPLICIT_LIB_PLAIN_ESC "${_OPENMP_IMPLICIT_LIB_PLAIN}")
string(REGEX REPLACE "([][+.*?()^$])" "\\\\\\1" _OPENMP_IMPLICIT_LIB_PATH_ESC "${_OPENMP_IMPLICIT_LIB}")
if(NOT ( "${_OPENMP_IMPLICIT_LIB}" IN_LIST CMAKE_${LANG}_IMPLICIT_LINK_LIBRARIES
OR "${CMAKE_${LANG}_STANDARD_LIBRARIES}" MATCHES "(^| )(-Wl,)?(-l)?(${_OPENMP_IMPLICIT_LIB_PLAIN_ESC}|${_OPENMP_IMPLICIT_LIB_PATH_ESC})( |$)"
OR "${CMAKE_${LANG}_LINK_EXECUTABLE}" MATCHES "(^| )(-Wl,)?(-l)?(${_OPENMP_IMPLICIT_LIB_PLAIN_ESC}|${_OPENMP_IMPLICIT_LIB_PATH_ESC})( |$)" ) )
if(_OPENMP_IMPLICIT_LIB_DIR)
set(OpenMP_${_OPENMP_IMPLICIT_LIB_PLAIN}_LIBRARY "${_OPENMP_IMPLICIT_LIB}" CACHE FILEPATH
"Path to the ${_OPENMP_IMPLICIT_LIB_PLAIN} library for OpenMP")
else()
find_library(OpenMP_${_OPENMP_IMPLICIT_LIB_PLAIN}_LIBRARY
NAMES "${_OPENMP_IMPLICIT_LIB_NAME}"
DOC "Path to the ${_OPENMP_IMPLICIT_LIB_PLAIN} library for OpenMP"
HINTS ${OpenMP_${LANG}_IMPLICIT_LINK_DIRS}
CMAKE_FIND_ROOT_PATH_BOTH
NO_DEFAULT_PATH
)
endif()
mark_as_advanced(OpenMP_${_OPENMP_IMPLICIT_LIB_PLAIN}_LIBRARY)
list(APPEND _OPENMP_LIB_NAMES ${_OPENMP_IMPLICIT_LIB_PLAIN})
endif()
endforeach()
set("${OPENMP_LIB_NAMES_VAR}" "${_OPENMP_LIB_NAMES}" PARENT_SCOPE)
else()
# We do not know how to extract implicit OpenMP libraries for this compiler.
# Assume that it handles them automatically, e.g. the Intel Compiler on
# Windows should put the dependency in its object files.
set("${OPENMP_LIB_NAMES_VAR}" "" PARENT_SCOPE)
endif()
break()
elseif(CMAKE_${LANG}_COMPILER_ID STREQUAL "AppleClang"
AND CMAKE_${LANG}_COMPILER_VERSION VERSION_GREATER_EQUAL "7.0")
# Check for separate OpenMP library on AppleClang 7+
find_library(OpenMP_libomp_LIBRARY
NAMES omp gomp iomp5
HINTS ${CMAKE_${LANG}_IMPLICIT_LINK_DIRECTORIES}
)
mark_as_advanced(OpenMP_libomp_LIBRARY)
if(OpenMP_libomp_LIBRARY)
try_compile( OpenMP_COMPILE_RESULT_${FLAG_MODE}_${OPENMP_PLAIN_FLAG} ${CMAKE_BINARY_DIR} ${_OPENMP_TEST_SRC}
CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OPENMP_FLAGS_TEST}"
LINK_LIBRARIES ${CMAKE_${LANG}_VERBOSE_FLAG} ${OpenMP_libomp_LIBRARY}
OUTPUT_VARIABLE OpenMP_TRY_COMPILE_OUTPUT
)
if(OpenMP_COMPILE_RESULT_${FLAG_MODE}_${OPENMP_PLAIN_FLAG})
set("${OPENMP_FLAG_VAR}" "${OPENMP_FLAG}" PARENT_SCOPE)
set("${OPENMP_LIB_NAMES_VAR}" "libomp" PARENT_SCOPE)
break()
endif()
endif()
endif()
set("${OPENMP_LIB_NAMES_VAR}" "NOTFOUND" PARENT_SCOPE)
set("${OPENMP_FLAG_VAR}" "NOTFOUND" PARENT_SCOPE)
endforeach()
unset(OpenMP_VERBOSE_COMPILE_OPTIONS)
endfunction()
set(OpenMP_C_CXX_CHECK_VERSION_SOURCE
"
#include <stdio.h>
#include <omp.h>
const char ompver_str[] = { 'I', 'N', 'F', 'O', ':', 'O', 'p', 'e', 'n', 'M',
'P', '-', 'd', 'a', 't', 'e', '[',
('0' + ((_OPENMP/100000)%10)),
('0' + ((_OPENMP/10000)%10)),
('0' + ((_OPENMP/1000)%10)),
('0' + ((_OPENMP/100)%10)),
('0' + ((_OPENMP/10)%10)),
('0' + ((_OPENMP/1)%10)),
']', '\\0' };
int main()
{
puts(ompver_str);
return 0;
}
")
set(OpenMP_Fortran_CHECK_VERSION_SOURCE
"
program omp_ver
@OpenMP_Fortran_INCLUDE_LINE@
integer, parameter :: zero = ichar('0')
integer, parameter :: ompv = openmp_version
character, dimension(24), parameter :: ompver_str =&
(/ 'I', 'N', 'F', 'O', ':', 'O', 'p', 'e', 'n', 'M', 'P', '-',&
'd', 'a', 't', 'e', '[',&
char(zero + mod(ompv/100000, 10)),&
char(zero + mod(ompv/10000, 10)),&
char(zero + mod(ompv/1000, 10)),&
char(zero + mod(ompv/100, 10)),&
char(zero + mod(ompv/10, 10)),&
char(zero + mod(ompv/1, 10)), ']' /)
print *, ompver_str
end program omp_ver
")
function(_OPENMP_GET_SPEC_DATE LANG SPEC_DATE)
_OPENMP_WRITE_SOURCE_FILE("${LANG}" "CHECK_VERSION_SOURCE" OpenMPCheckVersion _OPENMP_TEST_SRC)
set(BIN_FILE "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/FindOpenMP/ompver_${LANG}.bin")
string(REGEX REPLACE "[-/=+]" "" OPENMP_PLAIN_FLAG "${OPENMP_FLAG}")
try_compile(OpenMP_SPECTEST_${LANG}_${OPENMP_PLAIN_FLAG} "${CMAKE_BINARY_DIR}" "${_OPENMP_TEST_SRC}"
CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OpenMP_${LANG}_FLAGS}"
COPY_FILE ${BIN_FILE})
if(${OpenMP_SPECTEST_${LANG}_${OPENMP_PLAIN_FLAG}})
file(STRINGS ${BIN_FILE} specstr LIMIT_COUNT 1 REGEX "INFO:OpenMP-date")
set(regex_spec_date ".*INFO:OpenMP-date\\[0*([^]]*)\\].*")
if("${specstr}" MATCHES "${regex_spec_date}")
set(${SPEC_DATE} "${CMAKE_MATCH_1}" PARENT_SCOPE)
endif()
endif()
endfunction()
macro(_OPENMP_SET_VERSION_BY_SPEC_DATE LANG)
set(OpenMP_SPEC_DATE_MAP
# Preview versions
"201611=5.0" # OpenMP 5.0 preview 1
# Combined versions, 2.5 onwards
"201511=4.5"
"201307=4.0"
"201107=3.1"
"200805=3.0"
"200505=2.5"
# C/C++ version 2.0
"200203=2.0"
# Fortran version 2.0
"200011=2.0"
# Fortran version 1.1
"199911=1.1"
# C/C++ version 1.0 (there's no 1.1 for C/C++)
"199810=1.0"
# Fortran version 1.0
"199710=1.0"
)
if(OpenMP_${LANG}_SPEC_DATE)
string(REGEX MATCHALL "${OpenMP_${LANG}_SPEC_DATE}=([0-9]+)\\.([0-9]+)" _version_match "${OpenMP_SPEC_DATE_MAP}")
else()
set(_version_match "")
endif()
if(NOT _version_match STREQUAL "")
set(OpenMP_${LANG}_VERSION_MAJOR ${CMAKE_MATCH_1})
set(OpenMP_${LANG}_VERSION_MINOR ${CMAKE_MATCH_2})
set(OpenMP_${LANG}_VERSION "${OpenMP_${LANG}_VERSION_MAJOR}.${OpenMP_${LANG}_VERSION_MINOR}")
else()
unset(OpenMP_${LANG}_VERSION_MAJOR)
unset(OpenMP_${LANG}_VERSION_MINOR)
unset(OpenMP_${LANG}_VERSION)
endif()
unset(_version_match)
unset(OpenMP_SPEC_DATE_MAP)
endmacro()
foreach(LANG IN ITEMS C CXX)
if(CMAKE_${LANG}_COMPILER_LOADED)
if(NOT DEFINED OpenMP_${LANG}_FLAGS OR "${OpenMP_${LANG}_FLAGS}" STREQUAL "NOTFOUND"
OR NOT DEFINED OpenMP_${LANG}_LIB_NAMES OR "${OpenMP_${LANG}_LIB_NAMES}" STREQUAL "NOTFOUND")
_OPENMP_GET_FLAGS("${LANG}" "${LANG}" OpenMP_${LANG}_FLAGS_WORK OpenMP_${LANG}_LIB_NAMES_WORK)
endif()
set(OpenMP_${LANG}_FLAGS "${OpenMP_${LANG}_FLAGS_WORK}"
CACHE STRING "${LANG} compiler flags for OpenMP parallelization")
set(OpenMP_${LANG}_LIB_NAMES "${OpenMP_${LANG}_LIB_NAMES_WORK}"
CACHE STRING "${LANG} compiler libraries for OpenMP parallelization")
mark_as_advanced(OpenMP_${LANG}_FLAGS OpenMP_${LANG}_LIB_NAMES)
endif()
endforeach()
if(CMAKE_Fortran_COMPILER_LOADED)
if(NOT DEFINED OpenMP_Fortran_FLAGS OR "${OpenMP_Fortran_FLAGS}" STREQUAL "NOTFOUND"
OR NOT DEFINED OpenMP_Fortran_LIB_NAMES OR "${OpenMP_Fortran_LIB_NAMES}" STREQUAL "NOTFOUND"
OR NOT DEFINED OpenMP_Fortran_HAVE_OMPLIB_MODULE)
set(OpenMP_Fortran_INCLUDE_LINE "use omp_lib\n implicit none")
_OPENMP_GET_FLAGS("Fortran" "FortranHeader" OpenMP_Fortran_FLAGS_WORK OpenMP_Fortran_LIB_NAMES_WORK)
if(OpenMP_Fortran_FLAGS_WORK)
set(OpenMP_Fortran_HAVE_OMPLIB_MODULE TRUE CACHE BOOL INTERNAL "")
endif()
set(OpenMP_Fortran_FLAGS "${OpenMP_Fortran_FLAGS_WORK}"
CACHE STRING "Fortran compiler flags for OpenMP parallelization")
set(OpenMP_Fortran_LIB_NAMES "${OpenMP_Fortran_LIB_NAMES_WORK}"
CACHE STRING "Fortran compiler libraries for OpenMP parallelization")
mark_as_advanced(OpenMP_Fortran_FLAGS OpenMP_Fortran_LIB_NAMES)
endif()
if(NOT DEFINED OpenMP_Fortran_FLAGS OR "${OpenMP_Fortran_FLAGS}" STREQUAL "NOTFOUND"
OR NOT DEFINED OpenMP_Fortran_LIB_NAMES OR "${OpenMP_Fortran_LIB_NAMES}" STREQUAL "NOTFOUND"
OR NOT DEFINED OpenMP_Fortran_HAVE_OMPLIB_HEADER)
set(OpenMP_Fortran_INCLUDE_LINE "implicit none\n include 'omp_lib.h'")
_OPENMP_GET_FLAGS("Fortran" "FortranModule" OpenMP_Fortran_FLAGS_WORK OpenMP_Fortran_LIB_NAMES_WORK)
if(OpenMP_Fortran_FLAGS_WORK)
set(OpenMP_Fortran_HAVE_OMPLIB_HEADER TRUE CACHE BOOL INTERNAL "")
endif()
set(OpenMP_Fortran_FLAGS "${OpenMP_Fortran_FLAGS_WORK}"
CACHE STRING "Fortran compiler flags for OpenMP parallelization")
set(OpenMP_Fortran_LIB_NAMES "${OpenMP_Fortran_LIB_NAMES}"
CACHE STRING "Fortran compiler libraries for OpenMP parallelization")
endif()
if(OpenMP_Fortran_HAVE_OMPLIB_MODULE)
set(OpenMP_Fortran_INCLUDE_LINE "use omp_lib\n implicit none")
else()
set(OpenMP_Fortran_INCLUDE_LINE "implicit none\n include 'omp_lib.h'")
endif()
endif()
if(NOT OpenMP_FIND_COMPONENTS)
set(OpenMP_FINDLIST C CXX Fortran)
else()
set(OpenMP_FINDLIST ${OpenMP_FIND_COMPONENTS})
endif()
unset(_OpenMP_MIN_VERSION)
include(FindPackageHandleStandardArgs)
foreach(LANG IN LISTS OpenMP_FINDLIST)
if(CMAKE_${LANG}_COMPILER_LOADED)
if (NOT OpenMP_${LANG}_SPEC_DATE AND OpenMP_${LANG}_FLAGS)
_OPENMP_GET_SPEC_DATE("${LANG}" OpenMP_${LANG}_SPEC_DATE_INTERNAL)
set(OpenMP_${LANG}_SPEC_DATE "${OpenMP_${LANG}_SPEC_DATE_INTERNAL}" CACHE
INTERNAL "${LANG} compiler's OpenMP specification date")
_OPENMP_SET_VERSION_BY_SPEC_DATE("${LANG}")
endif()
set(OpenMP_${LANG}_FIND_QUIETLY ${OpenMP_FIND_QUIETLY})
set(OpenMP_${LANG}_FIND_REQUIRED ${OpenMP_FIND_REQUIRED})
set(OpenMP_${LANG}_FIND_VERSION ${OpenMP_FIND_VERSION})
set(OpenMP_${LANG}_FIND_VERSION_EXACT ${OpenMP_FIND_VERSION_EXACT})
set(_OPENMP_${LANG}_REQUIRED_VARS OpenMP_${LANG}_FLAGS)
if("${OpenMP_${LANG}_LIB_NAMES}" STREQUAL "NOTFOUND")
set(_OPENMP_${LANG}_REQUIRED_LIB_VARS OpenMP_${LANG}_LIB_NAMES)
else()
foreach(_OPENMP_IMPLICIT_LIB IN LISTS OpenMP_${LANG}_LIB_NAMES)
list(APPEND _OPENMP_${LANG}_REQUIRED_LIB_VARS OpenMP_${_OPENMP_IMPLICIT_LIB}_LIBRARY)
endforeach()
endif()
find_package_handle_standard_args(OpenMP_${LANG}
REQUIRED_VARS OpenMP_${LANG}_FLAGS ${_OPENMP_${LANG}_REQUIRED_LIB_VARS}
VERSION_VAR OpenMP_${LANG}_VERSION
)
if(OpenMP_${LANG}_FOUND)
if(DEFINED OpenMP_${LANG}_VERSION)
if(NOT _OpenMP_MIN_VERSION OR _OpenMP_MIN_VERSION VERSION_GREATER OpenMP_${LANG}_VERSION)
set(_OpenMP_MIN_VERSION OpenMP_${LANG}_VERSION)
endif()
endif()
set(OpenMP_${LANG}_LIBRARIES "")
foreach(_OPENMP_IMPLICIT_LIB IN LISTS OpenMP_${LANG}_LIB_NAMES)
list(APPEND OpenMP_${LANG}_LIBRARIES "${OpenMP_${_OPENMP_IMPLICIT_LIB}_LIBRARY}")
endforeach()
if(NOT TARGET OpenMP::OpenMP_${LANG})
add_library(OpenMP::OpenMP_${LANG} INTERFACE IMPORTED)
endif()
if(OpenMP_${LANG}_FLAGS)
separate_arguments(_OpenMP_${LANG}_OPTIONS NATIVE_COMMAND "${OpenMP_${LANG}_FLAGS}")
set_property(TARGET OpenMP::OpenMP_${LANG} PROPERTY
INTERFACE_COMPILE_OPTIONS "$<$<COMPILE_LANGUAGE:${LANG}>:${_OpenMP_${LANG}_OPTIONS}>")
unset(_OpenMP_${LANG}_OPTIONS)
endif()
if(OpenMP_${LANG}_LIBRARIES)
set_property(TARGET OpenMP::OpenMP_${LANG} PROPERTY
INTERFACE_LINK_LIBRARIES "${OpenMP_${LANG}_LIBRARIES}")
endif()
endif()
endif()
endforeach()
unset(_OpenMP_REQ_VARS)
foreach(LANG IN ITEMS C CXX Fortran)
if((NOT OpenMP_FIND_COMPONENTS AND CMAKE_${LANG}_COMPILER_LOADED) OR LANG IN_LIST OpenMP_FIND_COMPONENTS)
list(APPEND _OpenMP_REQ_VARS "OpenMP_${LANG}_FOUND")
endif()
endforeach()
find_package_handle_standard_args(OpenMP
REQUIRED_VARS ${_OpenMP_REQ_VARS}
VERSION_VAR ${_OpenMP_MIN_VERSION}
HANDLE_COMPONENTS)
set(OPENMP_FOUND ${OpenMP_FOUND})
if(CMAKE_Fortran_COMPILER_LOADED AND OpenMP_Fortran_FOUND)
if(NOT DEFINED OpenMP_Fortran_HAVE_OMPLIB_MODULE)
set(OpenMP_Fortran_HAVE_OMPLIB_MODULE FALSE CACHE BOOL INTERNAL "")
endif()
if(NOT DEFINED OpenMP_Fortran_HAVE_OMPLIB_HEADER)
set(OpenMP_Fortran_HAVE_OMPLIB_HEADER FALSE CACHE BOOL INTERNAL "")
endif()
endif()
if(NOT ( CMAKE_C_COMPILER_LOADED OR CMAKE_CXX_COMPILER_LOADED OR CMAKE_Fortran_COMPILER_LOADED ))
message(SEND_ERROR "FindOpenMP requires the C, CXX or Fortran languages to be enabled")
endif()
unset(OpenMP_C_CXX_TEST_SOURCE)
unset(OpenMP_Fortran_TEST_SOURCE)
unset(OpenMP_C_CXX_CHECK_VERSION_SOURCE)
unset(OpenMP_Fortran_CHECK_VERSION_SOURCE)
unset(OpenMP_Fortran_INCLUDE_LINE)
cmake_policy(POP)

2
CMake/VTKmCPUVectorization.cmake
View File

@ -68,7 +68,7 @@
#
#
# guard agaisnt building vectorization_flags more than once
# guard against building vectorization_flags more than once
if(TARGET vtkm_vectorization_flags)
return()
endif()

13
CMake/VTKmConfig.cmake.in
View File

@ -27,6 +27,9 @@
# vtkm::tbb Target that contains tbb related link information
# implicitly linked to by `vtkm_cont` if tbb is enabled
#
# vtkm::openmp Target that contains openmp related link information
# implicitly linked to by `vtkm_cont` if openmp is enabled
#
# vtkm::cuda Target that contains cuda related link information
# implicitly linked to by `vtkm_cont` if cuda is enabled
#
@ -41,6 +44,7 @@
# VTKm_BUILD_SHARED_LIBS Will be enabled if VTK-m was built shared/dynamic
# VTKm_ENABLE_CUDA Will be enabled if VTK-m was built with CUDA support
# VTKm_ENABLE_TBB Will be enabled if VTK-m was built with TBB support
# VTKm_ENABLE_OPENMP Will be enabled if VTK-m was built with OpenMP support
# VTKm_ENABLE_MPI Will be enabled if VTK-m was built with MPI support
# VTKm_ENABLE_RENDERING Will be enabled if VTK-m was built with rendering support
# VTKm_ENABLE_GL_CONTEXT Will be enabled if VTK-m rendering was built with a GL context
@ -49,7 +53,13 @@
#
#
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
if (CMAKE_VERSION VERSION_LESS "3.3")
message(FATAL_ERROR "VTK-m requires CMake 3.3+")
endif()
if("${CMAKE_GENERATOR}" MATCHES "Visual Studio" AND
CMAKE_VERSION VERSION_LESS "3.11")
message(FATAL_ERROR "VTK-m requires CMake 3.11+ when using the Visual Studio Generators")
endif()
@PACKAGE_INIT@
@ -63,6 +73,7 @@ set(VTKm_VERSION "@VTKm_VERSION@")
set(VTKm_BUILD_SHARED_LIBS "@VTKm_BUILD_SHARED_LIBS@")
set(VTKm_ENABLE_CUDA "@VTKm_ENABLE_CUDA@")
set(VTKm_ENABLE_TBB "@VTKm_ENABLE_TBB@")
set(VTKm_ENABLE_OPENMP "@VTKm_ENABLE_OPENMP@")
set(VTKm_ENABLE_RENDERING "@VTKm_ENABLE_RENDERING@")
set(VTKm_ENABLE_GL_CONTEXT "@VTKm_ENABLE_GL_CONTEXT@")
set(VTKm_ENABLE_OSMESA_CONTEXT "@VTKm_ENABLE_OSMESA_CONTEXT@")

103
CMake/VTKmDeviceAdapters.cmake
View File

@ -18,10 +18,43 @@
## this software.
##============================================================================
#
function(vtkm_extract_real_library library real_library)
if(NOT UNIX)
set(${real_library} "${library}" PARENT_SCOPE)
return()
endif()
#Read in the first 4 bytes and see if they are the ELF magic number
set(_elf_magic "7f454c46")
file(READ ${library} _hex_data OFFSET 0 LIMIT 4 HEX)
if(_hex_data STREQUAL _elf_magic)
#we have opened a elf binary so this is what
#we should link too
set(${real_library} "${library}" PARENT_SCOPE)
return()
endif()
file(READ ${library} _data OFFSET 0 LIMIT 1024)
if("${_data}" MATCHES "INPUT \\(([^(]+)\\)")
#extract out the so name from REGEX MATCh command
set(_proper_so_name "${CMAKE_MATCH_1}")
#construct path to the real .so which is presumed to be in the same directory
#as the input file
get_filename_component(_so_dir "${library}" DIRECTORY)
set(${real_library} "${_so_dir}/${_proper_so_name}" PARENT_SCOPE)
else()
#unable to determine what this library is so just hope everything works
#add pass it unmodified.
set(${real_library} "${library}" PARENT_SCOPE)
endif()
endfunction()
if(VTKm_ENABLE_TBB AND NOT TARGET vtkm::tbb)
find_package(TBB REQUIRED)
# Workaround a bug in older versions of cmake prevents linking with UNKOWN IMPORTED libraries
# Workaround a bug in older versions of cmake prevents linking with UNKNOWN IMPORTED libraries
# refer to CMake issue #17245
if (CMAKE_VERSION VERSION_LESS 3.10)
add_library(vtkm::tbb SHARED IMPORTED GLOBAL)
@ -32,24 +65,49 @@ if(VTKm_ENABLE_TBB AND NOT TARGET vtkm::tbb)
set_target_properties(vtkm::tbb PROPERTIES
INTERFACE_INCLUDE_DIRECTORIES "${TBB_INCLUDE_DIRS}")
if(TBB_LIBRARY_RELEASE)
set_property(TARGET vtkm::tbb APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
set_target_properties(vtkm::tbb PROPERTIES IMPORTED_LOCATION_RELEASE "${TBB_LIBRARY_RELEASE}")
endif()
if(EXISTS "${TBB_LIBRARY_RELEASE}")
vtkm_extract_real_library("${TBB_LIBRARY_RELEASE}" 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(TBB_LIBRARY_DEBUG)
set_property(TARGET vtkm::tbb APPEND PROPERTY IMPORTED_CONFIGURATIONS DEBUG)
set_target_properties(vtkm::tbb PROPERTIES IMPORTED_LOCATION_DEBUG "${TBB_LIBRARY_DEBUG}")
endif()
if(NOT TBB_LIBRARY_RELEASE AND NOT TBB_LIBRARY_DEBUG)
set_property(TARGET vtkm::tbb APPEND PROPERTY IMPORTED_LOCATION "${TBB_LIBRARY}")
endif()
if(EXISTS "${TBB_LIBRARY_DEBUG}")
vtkm_extract_real_library("${TBB_LIBRARY_DEBUG}" real_path)
set_property(TARGET vtkm::tbb APPEND PROPERTY IMPORTED_CONFIGURATIONS DEBUG)
set_target_properties(vtkm::tbb PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
IMPORTED_LOCATION_DEBUG "${real_path}"
)
endif()
endif()
if(VTKm_ENABLE_OPENMP AND NOT TARGET vtkm::openmp)
cmake_minimum_required(VERSION 3.9...3.12 FATAL_ERROR)
find_package(OpenMP 4.0 REQUIRED COMPONENTS CXX QUIET)
add_library(vtkm::openmp INTERFACE IMPORTED GLOBAL)
if(OpenMP_CXX_FLAGS)
set_property(TARGET vtkm::openmp
APPEND PROPERTY INTERFACE_COMPILE_OPTIONS $<$<COMPILE_LANGUAGE:CXX>:${OpenMP_CXX_FLAGS}>)
if(VTKm_ENABLE_CUDA)
string(REPLACE ";" "," openmp_cuda_flags "-Xcompiler=${OpenMP_CXX_FLAGS}")
set_property(TARGET vtkm::openmp
APPEND PROPERTY INTERFACE_COMPILE_OPTIONS $<$<COMPILE_LANGUAGE:CUDA>:${openmp_cuda_flags}>)
endif()
endif()
if(OpenMP_CXX_LIBRARIES)
set_target_properties(vtkm::openmp PROPERTIES
INTERFACE_LINK_LIBRARIES "${OpenMP_CXX_LIBRARIES}")
endif()
endif()
if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
cmake_minimum_required(VERSION 3.9 FATAL_ERROR)
cmake_minimum_required(VERSION 3.9...3.12 FATAL_ERROR)
enable_language(CUDA)
#To work around https://gitlab.kitware.com/cmake/cmake/issues/17512
@ -58,7 +116,7 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
list(APPEND CMAKE_CUDA_IMPLICIT_INCLUDE_DIRECTORIES "${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES}")
endif()
# Workaround a bug in older versions of cmake prevents linking with UNKOWN IMPORTED libraries
# Workaround a bug in older versions of cmake prevents linking with UNKNOWN IMPORTED libraries
# refer to CMake issue #17245
if (CMAKE_VERSION VERSION_LESS 3.10)
add_library(vtkm::cuda STATIC IMPORTED GLOBAL)
@ -75,24 +133,27 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
if("x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
get_filename_component(VTKM_CUDA_BIN_DIR "${CMAKE_CUDA_COMPILER}" DIRECTORY)
set_property(TARGET vtkm::cuda APPEND PROPERTY IMPORTED_LOCATION "${VTKM_CUDA_BIN_DIR}/../lib/x64/cudadevrt.lib")
set_target_properties(vtkm::cuda PROPERTIES INTERFACE_INCLUDE_DIRECTORIES "${VTKM_CUDA_BIN_DIR}/../include/")
set_target_properties(vtkm::cuda PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
IMPORTED_LOCATION "${VTKM_CUDA_BIN_DIR}/../lib/x64/cudadevrt.lib"
INTERFACE_INCLUDE_DIRECTORIES "${VTKM_CUDA_BIN_DIR}/../include/"
)
else()
list(GET CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES 0 VTKM_CUDA_LIBRARY)
if(IS_ABSOLUTE "${VTKM_CUDA_LIBRARY}")
set_property(TARGET vtkm::cuda APPEND PROPERTY IMPORTED_LOCATION "${VTKM_CUDA_LIBRARY}")
set_target_properties(vtkm::cuda PROPERTIES IMPORTED_LOCATION "${VTKM_CUDA_LIBRARY}")
else()
find_library(cuda_lib
NAME ${VTKM_CUDA_LIBRARY}
PATHS ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}
)
set(VTKM_CUDA_LIBRARY ${cuda_lib})
set_property(TARGET vtkm::cuda APPEND PROPERTY IMPORTED_LOCATION "${VTKM_CUDA_LIBRARY}")
set_target_properties(vtkm::cuda PROPERTIES IMPORTED_LOCATION "${VTKM_CUDA_LIBRARY}")
unset(cuda_lib CACHE)
endif()
set_target_properties(vtkm::cuda PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES "CXX"
INTERFACE_INCLUDE_DIRECTORIES "${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}")
endif()
@ -140,7 +201,7 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
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)
#detect what the propery is set too
#detect what the property is set too
if(VTKm_CUDA_Architecture STREQUAL "native")
if(VTKM_CUDA_NATIVE_EXE_PROCESS_RAN_OUTPUT)

2
CMake/VTKmWrappers.cmake
View File

@ -61,7 +61,7 @@ endfunction(vtkm_pyexpander_generated_file)
#-----------------------------------------------------------------------------
function(vtkm_compile_as_cuda output)
# We cant use set_source_files_properties(<> PROPERTIES LANGUAGE "CUDA")
# We can't use set_source_files_properties(<> PROPERTIES LANGUAGE "CUDA")
# for the following reasons:
#
# 1. As of CMake 3.10 MSBuild cuda language support has a bug where files

19
CMakeLists.txt
View File

@ -25,11 +25,11 @@
# If you want CUDA support, you will need to have CMake 3.9 on Linux/OSX.
# We require CMake 3.11 with the MSVC generator as the $<COMPILE_LANGUAGE:>
# generator expression is not supported on older versions.
cmake_minimum_required(VERSION 3.3)
cmake_minimum_required(VERSION 3.3...3.12)
project (VTKm)
if(${CMAKE_GENERATOR} MATCHES "Visual Studio")
cmake_minimum_required(VERSION 3.11 FATAL_ERROR)
cmake_minimum_required(VERSION 3.11...3.12 FATAL_ERROR)
endif()
# Update module path
@ -86,6 +86,7 @@ endmacro ()
# Configurable Options
vtkm_option(VTKm_ENABLE_CUDA "Enable Cuda support" OFF)
vtkm_option(VTKm_ENABLE_TBB "Enable TBB support" OFF)
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)
@ -120,7 +121,7 @@ set(VTKm_BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS})
# This flag can be used to prevent VTK-m from exporting its warning flags in its
# build interface. This is useful when building VTK-m as a thirdparty library
# and the warnings are too strict for the parent project.
vtkm_option(VTKm_ENABLE_DEVELOPER_FLAGS "Enable compiler flags that are usefull while developing VTK-m" ON)
vtkm_option(VTKm_ENABLE_DEVELOPER_FLAGS "Enable compiler flags that are useful while developing VTK-m" ON)
mark_as_advanced(
VTKm_NO_ASSERT
@ -165,7 +166,16 @@ endif()
#-----------------------------------------------------------------------------
if (VTKm_ENABLE_TESTING)
enable_testing()
include(CTest)
# Only include CTest if it has not been included by a superproject. The
# variable DEFAULT_CTEST_CONFIGURATION_TYPE is a non-cached variable set by
# CTest.cmake, so we'll use that to determine if it's already included.
if(NOT DEFINED DEFAULT_CTEST_CONFIGURATION_TYPE)
include(CTest)
# Mark this as advanced to avoid confusion, since we actually rely on
# VTKm_ENABLE_TESTING.
mark_as_advanced(BUILD_TESTING)
endif()
configure_file(${VTKm_SOURCE_DIR}/CTestCustom.cmake.in
${VTKm_BINARY_DIR}/CTestCustom.cmake @ONLY)
@ -266,6 +276,7 @@ if(NOT VTKm_INSTALL_ONLY_LIBRARIES)
FILES
${VTKm_SOURCE_DIR}/CMake/FindTBB.cmake
${VTKm_SOURCE_DIR}/CMake/FindOpenGL.cmake
${VTKm_SOURCE_DIR}/CMake/FindOpenMP.cmake
DESTINATION ${VTKm_INSTALL_CMAKE_MODULE_DIR}
)

3
CTestCustom.cmake.in
View File

@ -28,6 +28,9 @@ list(APPEND CTEST_CUSTOM_WARNING_EXCEPTION
# disable doxygen warning about not generating graph
".*warning: Included by graph for"
# Doxygen warns when creating output directory:
"Notice: Output directory.*does not exist. I have created it for you."
# disable doxygen warnings from CONTRIBUTING.md, CodingConventions.md.
# these files are really intended for Gitlab, hence we don't want to use
# doxygen tags in them.

6
README.md
View File

@ -64,15 +64,17 @@ VTK-m Requires:
+ MSVC 2015+
+ [CMake](http://www.cmake.org/download/)
+ CMake 3.3+ (for any build)
+ CMake 3.9+ (for CUDA build)
+ CMake 3.9+ (for CUDA build or OpenMP build)
+ CMake 3.11+ (for Visual Studio generator)
Optional dependencies are:
+ CUDA Device Adapter
+ [Cuda Toolkit 7+](https://developer.nvidia.com/cuda-toolkit)
+ [Cuda Toolkit 7.5+](https://developer.nvidia.com/cuda-toolkit)
+ TBB Device Adapter
+ [TBB](https://www.threadingbuildingblocks.org/)
+ OpenMP Device Adapter
+ Requires a compiler that supports OpenMP >= 4.0.
+ OpenGL Rendering
+ The rendering module contains multiple rendering implementations
including standalone rendering code. The rendering module also

19
benchmarking/BenchmarkDeviceAdapter.cxx
View File

@ -46,7 +46,9 @@
#if VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
#include <tbb/task_scheduler_init.h>
#endif // TBB
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
#include <omp.h>
#endif
// This benchmark has a number of commandline options to customize its behavior.
// See The BenchDevAlgoConfig documentations for details.
@ -748,6 +750,7 @@ private:
const vtkm::Id N_VALID;
const vtkm::Id PERCENT_VALID;
ValueArrayHandle ValueHandle;
IndexArrayHandle IndexHandle;
VTKM_CONT
BenchStableSortIndicesUnique(vtkm::Id percent_valid)
@ -759,12 +762,14 @@ private:
FillModuloTestValueKernel<Value>(
N_VALID, this->ValueHandle.PrepareForOutput(arraySize, DeviceAdapterTag())),
arraySize);
this->IndexHandle = SSI::Sort(this->ValueHandle);
}
VTKM_CONT
vtkm::Float64 operator()()
{
IndexArrayHandle indices = SSI::Sort(this->ValueHandle);
IndexArrayHandle indices;
Algorithm::Copy(this->IndexHandle, indices);
Timer timer;
SSI::Unique(this->ValueHandle, indices);
return timer.GetElapsedTime();
@ -1190,6 +1195,8 @@ int main(int argc, char* argv[])
{
#if VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
int numThreads = tbb::task_scheduler_init::automatic;
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
int numThreads = omp_get_max_threads();
#endif // TBB
vtkm::benchmarking::BenchDevAlgoConfig& config = vtkm::benchmarking::Config;
@ -1320,8 +1327,12 @@ int main(int argc, char* argv[])
std::istringstream parse(argv[i]);
parse >> numThreads;
std::cout << "Selected " << numThreads << " TBB threads." << std::endl;
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
std::istringstream parse(argv[i]);
parse >> numThreads;
std::cout << "Selected " << numThreads << " OpenMP threads." << std::endl;
#else
std::cerr << "NumThreads valid only on TBB. Ignoring." << std::endl;
std::cerr << "NumThreads not valid on this device. Ignoring." << std::endl;
#endif // TBB
}
else
@ -1334,6 +1345,8 @@ int main(int argc, char* argv[])
#if VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
// Must not be destroyed as long as benchmarks are running:
tbb::task_scheduler_init init(numThreads);
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
omp_set_num_threads(numThreads);
#endif // TBB
if (config.BenchmarkFlags == 0)

9
benchmarking/CMakeLists.txt
View File

@ -30,6 +30,12 @@ function(add_benchmark name files)
target_compile_definitions(${name}_TBB PRIVATE "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_TBB")
endif()
if (TARGET vtkm::openmp)
add_executable(${name}_OPENMP ${files})
list(APPEND benchmarks ${name}_OPENMP)
target_compile_definitions(${name}_OPENMP PRIVATE "VTKM_DEVICE_ADAPTER=VTKM_DEVICE_ADAPTER_OPENMP")
endif()
if (TARGET vtkm::cuda)
get_filename_component(fname "${name}" NAME_WE)
get_filename_component(fullpath "${name}.cxx" ABSOLUTE)
@ -70,6 +76,9 @@ if(TARGET vtkm_rendering)
if(TARGET BenchmarkRayTracing_TBB)
target_link_libraries(BenchmarkRayTracing_TBB PRIVATE vtkm_rendering)
endif()
if(TARGET BenchmarkRayTracing_OPENMP)
target_link_libraries(BenchmarkRayTracing_OPENMP PRIVATE vtkm_rendering)
endif()
if(TARGET BenchmarkRayTracing_CUDA)
target_link_libraries(BenchmarkRayTracing_CUDA PRIVATE vtkm_rendering)
endif()

4
docs/changelog/Allow-custom-range-in-histogram-filter.md Normal file
View File

@ -0,0 +1,4 @@
#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.

2
docs/changelog/add-float-version-functions-pi.md
View File

@ -2,4 +2,4 @@
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 serveral conversion warnings.
template ones and fix several conversion warnings.

9
docs/changelog/allow-CoordinateSystem-to-unload-execution-resources.md Normal file
View File

@ -0,0 +1,9 @@
# 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.

2
docs/changelog/cell-measure.md
View File

@ -45,7 +45,7 @@ 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.
Finaly, note that cell measures may return invalid (NaN) or infinite (Inf, -Inf)
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.

75
docs/changelog/geometry.md Normal file
View File

@ -0,0 +1,75 @@
# 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.

5
docs/changelog/openmp-backend.md Normal file
View File

@ -0,0 +1,5 @@
# 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.

7
docs/changelog/swap.md Normal file
View File

@ -0,0 +1,7 @@
# 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.

94
docs/changelog/vec-initializer-lists.md Normal file
View File

@ -0,0 +1,94 @@
# Support initializer lists for Vec
Add constructors to the `vtkm::Vec` classes that accept `std::initializer_list`. The
main advantage of this addition is that it makes it much easier to initialize `Vec`s
of arbitrary length.
Although previously some `Vec` classes could be constructed with values listed in
their parameters, that only worked for initializes up to size 4.
``` cpp
vtkm::Vec<vtkm::Float64, 3> vec1{1.1, 2.2, 3.3}; // New better initializer
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(1.1, 2.2, 3.3); // Old style that still works but
// probably should be deprecated.
```
Nested initializer lists work to initialize `Vec`s of `Vec`s.
``` cpp
vtkm::Vec<vtkm::Vec<vtkm::Float64, 2>, 3> vec{ {1.1, 2.2}, {3.3, 4.4}, {5.5, 6.6} };
```
The nice thing about the `std::initializer_list` implementation is that it works for any
size `Vec`. That keeps us from jumping through hoops for larger `Vec`s.
``` cpp
vtkm::Vec<vtkm::Float64, 5> vec1{1.1, 2.2, 3.3, 4.4, 5.5}; // Works fine.
vtkm::Vec<vtkm::Float64, 5> vec2(1.1, 2.2, 3.3, 4.4, 5.5); // ERROR! This constructor
// not implemented!
```
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`s 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.
``` 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.
There is also a version of `vtkm::make_Vec` that accepts an `std::initializer_list`,
but it is not very useful because you have to separately specify the length of the
`Vec` (due to some limitations of `std::initializer_list` in C++11).
``` cpp
// Creates a vtkm::Vec<vtkm::Float64, 3>
auto vec1 = vtkm::make_Vec<3>({1.1, 2.2, 3.3});
// Creates exactly the same thing
auto vec2 = vtkm::Vec<vtkm::Float64, 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};
```

2
docs/changelog/vec-vec-constructors.md
View File

@ -1,7 +1,7 @@
# 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 initalize
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>;

2
examples/CMakeLists.txt
View File

@ -34,7 +34,7 @@ add_subdirectory(histogram)
add_subdirectory(isosurface)
add_subdirectory(multi_backend)
add_subdirectory(particle_advection)
add_subdirectory(temporaladvection)
add_subdirectory(temporal_advection)
add_subdirectory(redistribute_points)
add_subdirectory(rendering)
add_subdirectory(streamline)

2
examples/clipping/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(Clipping CXX)
#Find the VTK-m package

2
examples/contour_tree/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(ContourTree CXX)
#Find the VTK-m package

2
examples/cosmotools/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(CosmoTools CXX)
#Find the VTK-m package

2
examples/demo/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(VTKmDemo CXX)
#Find the VTK-m package

2
examples/game_of_life/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(GameOfLife CXX)
#Find the VTK-m package

2
examples/hello_world/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(HelloWorld CXX)
#Find the VTK-m package

2
examples/histogram/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(Histogram CXX)
#Find the VTK-m package

2
examples/isosurface/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(IsoSurface CXX)
#Find the VTK-m package

2
examples/multi_backend/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(MultiBackend CXX)
#Find the VTK-m package

24
examples/multi_backend/MultiDeviceGradient.hxx
View File

@ -27,15 +27,15 @@
namespace
{
vtkm::Id deterine_cuda_gpu_count()
int determine_cuda_gpu_count()
{
vtkm::Id count = 0;
int count = 0;
#if defined(VTKM_ENABLE_CUDA)
int numberOfDevices = 0;
auto res = cudaGetDeviceCount(&numberOfDevices);
if (res == cudaSuccess)
{
count = static_cast<vtkm::Id>(numberOfDevices);
count = numberOfDevices;
}
#endif
return count;
@ -115,7 +115,7 @@ VTKM_CONT MultiDeviceGradient::MultiDeviceGradient()
const bool runOnCuda = tracker.CanRunOn(vtkm::cont::DeviceAdapterTagCuda{});
//Note currently the virtual implementation has some issues
//In a multi-threaded enviornment only cuda can be used or
//In a multi-threaded environment only cuda can be used or
//all SMP backends ( Serial, TBB, OpenMP ).
//Once this issue is resolved we can enable CUDA + TBB in
//this example
@ -126,16 +126,16 @@ VTKM_CONT MultiDeviceGradient::MultiDeviceGradient()
if (runOnCuda)
{
std::cout << "adding cuda workers" << std::endl;
const vtkm::Id gpu_count = deterine_cuda_gpu_count();
for (vtkm::Id i = 0; i < gpu_count; ++i)
const int gpu_count = determine_cuda_gpu_count();
for (int i = 0; i < gpu_count; ++i)
{
//The number of workers per GPU is purely arbitrary currently,
//but in general we want multiple of them so we can overlap compute
//and transfer
this->Workers.emplace_back(process_block_cuda, std::ref(this->Queue), i);
this->Workers.emplace_back(process_block_cuda, std::ref(this->Queue), i);
this->Workers.emplace_back(process_block_cuda, std::ref(this->Queue), i);
this->Workers.emplace_back(process_block_cuda, std::ref(this->Queue), i);
this->Workers.emplace_back(std::bind(process_block_cuda, std::ref(this->Queue), i));
this->Workers.emplace_back(std::bind(process_block_cuda, std::ref(this->Queue), i));
this->Workers.emplace_back(std::bind(process_block_cuda, std::ref(this->Queue), i));
this->Workers.emplace_back(std::bind(process_block_cuda, std::ref(this->Queue), i));
}
}
//Step 3. Launch a worker that will use tbb (if enabled).
@ -144,7 +144,7 @@ VTKM_CONT MultiDeviceGradient::MultiDeviceGradient()
else if (runOnTbb)
{
std::cout << "adding a tbb worker" << std::endl;
this->Workers.emplace_back(process_block_tbb, std::ref(this->Queue));
this->Workers.emplace_back(std::bind(process_block_tbb, std::ref(this->Queue)));
}
}
@ -174,7 +174,7 @@ inline VTKM_CONT vtkm::cont::MultiBlock MultiDeviceGradient::PrepareForExecution
//Step 2. Construct the multi-block we are going to fill. The size signature
//to MultiBlock just reserves size
vtkm::cont::MultiBlock output;
output.AddBlocks(std::vector<vtkm::cont::DataSet>(mb.GetNumberOfBlocks()));
output.AddBlocks(std::vector<vtkm::cont::DataSet>(static_cast<size_t>(mb.GetNumberOfBlocks())));
vtkm::cont::MultiBlock* outPtr = &output;

2
examples/particle_advection/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(ParticleAdvection CXX)
#Find the VTK-m package

1
examples/particle_advection/ParticleAdvectionTBB.cxx
View File

@ -23,4 +23,3 @@
#define __BUILDING_TBB_VERSION__
#include "ParticleAdvection.cxx"
#include <tbb/task_scheduler_init.h>

2
examples/redistribute_points/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(RedistributePoints CXX)
#Find the VTK-m package

2
examples/rendering/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(RenderingExample CXX)
#Find the VTK-m package

2
examples/streamline/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(Streamline CXX)
#Find the VTK-m package

2
examples/tau_timing/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(TauTiming CXX)
#Find the VTK-m package

16
examples/temporaladvection/CMakeLists.txt → examples/temporal_advection/CMakeLists.txt
View File

@ -21,21 +21,21 @@
##=============================================================================
#Find the VTK-m package
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
project(ParticleAdvection CXX)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(TemporalAdvection CXX)
#Find the VTK-m package
find_package(VTKm REQUIRED QUIET)
add_executable(exeserial ParticleAdvection.cxx)
target_link_libraries(exeserial PRIVATE vtkm_cont)
add_executable(Temporal_Advection_SERIAL TemporalAdvection.cxx)
target_link_libraries(Temporal_Advection_SERIAL PRIVATE vtkm_cont)
if(TARGET vtkm::tbb)
add_executable(exetbb ParticleAdvectionTBB.cxx)
target_link_libraries(exetbb PRIVATE vtkm_cont)
add_executable(Temporal_Advection_TBB TemporalAdvectionTBB.cxx)
target_link_libraries(Temporal_Advection_TBB PRIVATE vtkm_cont)
endif()
if(TARGET vtkm::cuda)
add_executable(execuda ParticleAdvection.cu)
target_link_libraries(execuda PRIVATE vtkm_cont)
add_executable(Temporal_Advection_CUDA TemporalAdvection.cu)
target_link_libraries(Temporal_Advection_CUDA PRIVATE vtkm_cont)
endif()

2
examples/temporaladvection/ParticleAdvection.cu → examples/temporal_advection/TemporalAdvection.cu
View File

@ -20,4 +20,4 @@
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_CUDA
#include "ParticleAdvection.cxx"
#include "TemporalAdvection.cxx"

9
examples/temporaladvection/ParticleAdvection.cxx → examples/temporal_advection/TemporalAdvection.cxx
View File

@ -38,11 +38,11 @@
// The way to verify if the code produces correct streamlines
// is to do a visual test by using VisIt/ParaView to visualize
// the file written by this method.
int renderAndWriteDataSet(vtkm::cont::DataSet& dataset)
int renderAndWriteDataSet(const vtkm::cont::DataSet& dataset)
{
std::cout << "Trying to render the dataset" << std::endl;
vtkm::io::writer::VTKDataSetWriter writer("pathlines.vtk");
writer.WriteDataSet(dataset, static_cast<vtkm::Id>(0));
writer.WriteDataSet(dataset);
return 0;
}
@ -62,7 +62,6 @@ void RunTest(vtkm::Id numSteps, vtkm::Float32 stepSize, vtkm::Id advectType)
// as VTKm evolves.
vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3>> fieldArray1;
vtkm::cont::ArrayHandle<vtkm::Vec<FieldType, 3>> fieldArray2;
vtkm::Id numValues;
vtkm::io::reader::BOVDataSetReader reader1("slice1.bov");
vtkm::cont::DataSet ds1 = reader1.ReadDataSet();
@ -94,7 +93,7 @@ void RunTest(vtkm::Id numSteps, vtkm::Float32 stepSize, vtkm::Id advectType)
// This example does not work on scale, works on basic 11 particles
// so the results are more tractible
std::vector<vtkm::Vec<FieldType, 3>> seeds;
vtkm::Id x = 0, y = 5, z = 0;
FieldType x = 0, y = 5, z = 0;
for (int i = 0; i <= 11; i++)
{
vtkm::Vec<FieldType, 3> point;
@ -138,7 +137,7 @@ int main(int argc, char** argv)
}
numSteps = atoi(argv[1]);
stepSize = atof(argv[2]);
stepSize = static_cast<vtkm::Float32>(atof(argv[2]));
advectionType = atoi(argv[3]);
RunTest(numSteps, stepSize, advectionType);

5
examples/temporaladvection/ParticleAdvectionTBB.cxx → examples/temporal_advection/TemporalAdvectionTBB.cxx
View File

@ -20,7 +20,4 @@
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_TBB
#define __BUILDING_TBB_VERSION__
#include "ParticleAdvection.cxx"
#include <tbb/task_scheduler_init.h>
#include "TemporalAdvection.cxx"

2
examples/tetrahedra/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(Tetrahedra CXX)
#Find the VTK-m package

2
examples/unified_memory/CMakeLists.txt
View File

@ -19,7 +19,7 @@
## this software.
##
##=============================================================================
cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
cmake_minimum_required(VERSION 3.3...3.12 FATAL_ERROR)
project(UnifiedMemory CXX)
#Find the VTK-m package

15
vtkm/CMakeLists.txt
View File

@ -35,6 +35,7 @@ set(headers
CellShape.h
CellTraits.h
Flags.h
Geometry.h
Hash.h
ImplicitFunction.h
ListTag.h
@ -46,6 +47,7 @@ set(headers
RangeId.h
RangeId3.h
StaticAssert.h
Swap.h
TopologyElementTag.h
Transform3D.h
TypeListTag.h
@ -59,11 +61,20 @@ set(headers
VecVariable.h
VirtualObjectBase.h
UnaryPredicates.h
)
)
set(template_sources
Geometry.hxx
)
vtkm_pyexpander_generated_file(Math.h)
vtkm_declare_headers(${headers})
vtkm_declare_headers(
${headers}
${template_sources}
EXCLUDE_FROM_TESTING
${template_sources}
)
#-----------------------------------------------------------------------------
#first add all the components vtkm that are shared between control and exec

425
vtkm/Geometry.h Normal file
View File

@ -0,0 +1,425 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_Geometry_h
#define vtk_m_Geometry_h
#include <vtkm/VectorAnalysis.h>
namespace vtkm
{
// Forward declarations of geometric types:
template <typename CoordType, int Dim, bool IsTwoSided>
struct Ray;
template <typename CoordType, int Dim>
struct LineSegment;
template <typename CoordType>
struct Plane;
template <typename CoordType, int Dim>
struct Sphere;
/// Represent an infinite or semi-infinite line segment with a point and a direction
///
/// The \a IsTwoSided template parameter indicates whether the class represents
/// an infinite line extending in both directions from the base point or a
/// semi-infinite ray extending only in the positive direction from the base points.
template <typename CoordType = vtkm::FloatDefault, int Dim = 3, bool IsTwoSided = false>
struct Ray
{
static constexpr int Dimension = Dim;
using Vector = vtkm::Vec<CoordType, Dim>;
static constexpr bool TwoSided = IsTwoSided;
Vector Origin;
Vector Direction; // always stored as a unit length.
/// Construct a default 2-D ray, from (0,0) pointing along the +x axis.
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 2, int>::type = 0>
VTKM_EXEC_CONT Ray();
/// Construct a default 3-D ray from (0,0,0) pointing along the +x axis.
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 3, int>::type = 0>
VTKM_EXEC_CONT Ray();
/// Construct a ray from a point and direction.
VTKM_EXEC_CONT
Ray(const Vector& point, const Vector& direction);
/// Construct a ray from a line segment.
VTKM_EXEC_CONT
Ray(const LineSegment<CoordType, Dim>& segment);
/// Return whether the ray is valid or not.
///
/// It is possible for an invalid direction (zero length)
/// to be passed to the constructor. When this happens,
/// the constructor divides by zero, leaving Inf in all
/// components.
VTKM_EXEC_CONT
bool IsValid() const;
/// Compute a point along the line. \a param values > 0 lie on the ray.
VTKM_EXEC_CONT
Vector Evaluate(CoordType param) const;
/// Return the minmum distance from \a point to this line/ray.
///
/// Note that when the direction has zero length, this simplifies
/// the distance between \a point and the ray's origin.
/// Otherwise, the distance returned is either the perpendicular
/// distance from \a point to the line or the distance
/// to the ray's origin.
VTKM_EXEC_CONT
CoordType DistanceTo(const Vector& point) const;
/// Return the minimum distance between the ray/line and \a point.
VTKM_EXEC_CONT
CoordType DistanceTo(const Vector& point, CoordType& param, Vector& projectedPoint) const;
/// Compute the non-degenerate point where two 2-D rays intersect, or return false.
///
/// If true is returned, then the rays intersect in a unique point and
/// \a point is set to that location.
///
/// If false is returned, then either
/// (1) the rays are parallel and may be
/// coincident (intersect everywhere) or offset (intersect nowhere); or
/// (2) the lines intersect but not the rays (because the the intersection
/// occurs in the negative parameter space of one or both rays).
/// In the latter case (2), the \a point is still set to the location of the
/// intersection.
///
/// The tolerance \a tol is the minimum acceptable denominator used to
/// compute the intersection point coordinates and thus dictates the
/// maximum distance from the segments at which intersections will be
/// reported as valid.
template <bool OtherTwoSided, int Dim_ = Dim, typename std::enable_if<Dim_ == 2, int>::type = 0>
VTKM_EXEC_CONT bool Intersect(const Ray<CoordType, Dim, OtherTwoSided>& other,
Vector& point,
CoordType tol = 0.f);
};
/// Represent a finite line segment with a pair of points
template <typename CoordType = vtkm::FloatDefault, int Dim = 3>
struct LineSegment
{
static constexpr int Dimension = Dim;
using Vector = vtkm::Vec<CoordType, Dim>;
Vector Endpoints[2];
/// Construct a default segment from (0,0) to (1,0).
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 2, int>::type = 0>
VTKM_EXEC_CONT LineSegment();
/// Construct a default segment from (0,0,0) to (1,0,0).
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 3, int>::type = 0>
VTKM_EXEC_CONT LineSegment();
/// Construct a segment spanning points \a p0 and \a p1.
VTKM_EXEC_CONT
LineSegment(const Vector& p0, const Vector& p1);
/// Return whether this line segment has an infinitesimal extent (i.e., whether the endpoints are coincident)
VTKM_EXEC_CONT
bool IsSingular(CoordType tol2 = static_cast<CoordType>(1.0e-6f)) const;
/// Construct a plane bisecting this line segment (only when Dimension is 3).
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 3, int>::type = 0>
VTKM_EXEC_CONT Plane<CoordType> PerpendicularBisector() const;
/// Construct a perpendicular bisector to this line segment (only when Dimension is 2).
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 2, int>::type = 0>
VTKM_EXEC_CONT Ray<CoordType, Dim, true> PerpendicularBisector() const;
/// Return the midpoint of the line segment
VTKM_EXEC_CONT
Vector Center() const { return this->Evaluate(0.5f); }
/// Return the vector pointing to endpoint 1 from endpoint 0.
/// This vector is not of unit length and in the case of
/// degenerate lines, may have zero length.
/// Call vtkm::Normal() on the return value if you want a normalized result.
VTKM_EXEC_CONT
Vector Direction() const { return this->Endpoints[1] - this->Endpoints[0]; }
/// Compute a point along the line. \a param values in [0,1] lie on the line segment.
VTKM_EXEC_CONT
Vector Evaluate(CoordType param) const;
/// Return the minmum distance from \a point to this line segment.
///
/// Note that when the endpoints are coincident, this simplifies
/// the distance between \a point and either endpoint.
/// Otherwise, the distance returned is either the perpendicular
/// distance from \a point to the line segment or the distance
/// to the nearest endpoint (whichever is smaller).
VTKM_EXEC_CONT
CoordType DistanceTo(const Vector& point) const;
/// Return the minimum distance between the line segment and \a point.
VTKM_EXEC_CONT
CoordType DistanceTo(const Vector& point, CoordType& param, Vector& projectedPoint) const;
/// Compute the non-degenerate point where two (infinite) 2-D line segments intersect, or return false.
///
/// If true is returned, then the lines intersect in a unique point and
/// \a point is set to that location.
///
/// If false is returned, then the lines are parallel and either they are
/// coincident (intersect everywhere) or offset (intersect nowhere).
///
/// The tolerance \a tol is the minimum acceptable denominator used to
/// compute the intersection point coordinates and thus dictates the
/// maximum distance from the segments at which intersections will be
/// reported as valid.
template <int Dim_ = Dim, typename std::enable_if<Dim_ == 2, int>::type = 0>
VTKM_EXEC_CONT bool IntersectInfinite(const LineSegment<CoordType, Dim>& other,
Vector& point,
CoordType tol = 0.f);
};
/// Represent a plane with a base point (origin) and normal vector.
template <typename CoordType = vtkm::FloatDefault>
struct Plane
{
using Vector = vtkm::Vec<CoordType, 3>;
Vector Origin;
Vector Normal;
/// Construct a default plane whose base point is the origin and whose normal is (0,0,1)
VTKM_EXEC_CONT
Plane();
/// Construct a plane with the given \a origin and \a normal.
VTKM_EXEC_CONT
Plane(const Vector& origin, const Vector& normal, CoordType tol2 = static_cast<CoordType>(1e-8f));
/// Return true if the plane's normal is well-defined to within the given tolerance.
VTKM_EXEC_CONT
bool IsValid() const { return !vtkm::IsInf(this->Normal[0]); }
/// Return the **signed** distance from the plane to the point.
VTKM_EXEC_CONT
CoordType DistanceTo(const Vector& point) const;
/// Return the closest point in the plane to the given point.
VTKM_EXEC_CONT
Vector ClosestPoint(const Vector& point) const;
/// Intersect this plane with the ray (or line if the ray is two-sided).
///
/// Returns true if there is a non-degenrate intersection (i.e., an isolated point of intersection).
/// Returns false if there is no intersection *or* if the intersection is degenerate (i.e., the
/// entire ray/line lies in the plane).
/// In the latter case, \a lineInPlane will be true upon exit.
///
/// If this method returns true, then \a parameter will be set to a number indicating
/// where along the ray/line the plane hits and \a point will be set to that location.
/// If the input is a ray, the \a parameter will be non-negative.
template <bool IsTwoSided>
VTKM_EXEC_CONT bool Intersect(const Ray<CoordType, 3, IsTwoSided>& ray,
CoordType& parameter,
Vector& point,
bool& lineInPlane,
CoordType tol = CoordType(1e-6f)) const;
/// Intersect this plane with the line \a segment.
///
/// Returns true if there is a non-degenrate intersection (i.e., an isolated point of intersection).
/// Returns false if there is no intersection *or* if the intersection is degenerate (i.e., the
/// entire line segment lies in the plane).
/// In the latter case, \a lineInPlane will be true upon exit.
///
/// If this method returns true, then \a parameter will be set to a number in [0,1] indicating
/// where along the line segment the plane hits.
VTKM_EXEC_CONT
bool Intersect(const LineSegment<CoordType>& segment,
CoordType& parameter,
bool& lineInPlane) const;
/// Intersect this plane with the line \a segment.
///
/// Returns true if there is a non-degenrate intersection (i.e., an isolated point of intersection).
/// Returns false if there is no intersection *or* if the intersection is degenerate (i.e., the
/// entire line segment lines in the plane).
/// In the latter case, \a lineInPlane will be true upon exit.
///
/// If this method returns true, then \a parameter will be set to a number in [0,1] indicating
/// where along the line segment the plane hits and \a point will be set to that location.
VTKM_EXEC_CONT
bool Intersect(const LineSegment<CoordType>& segment,
CoordType& parameter,
Vector& point,
bool& lineInPlane) const;
/// Intersect this plane with another plane.
///
/// Returns true if there is a non-degenrate intersection (i.e., a line of intersection).
/// Returns false if there is no intersection *or* if the intersection is degenerate
/// (i.e., the planes are coincident).
/// In the latter case, \a coincident will be true upon exit and \a segment will
/// unmodified.
///
/// If this method returns true, then the resulting \a segment will have its
/// base point on the line of intersection and its second point will be a unit
/// length away in the direction of the cross produce of the input plane normals
/// (this plane crossed with the \a other).
///
/// The tolerance \a tol is the minimum squared length of the cross-product
/// of the two plane normals. It is also compared to the squared distance of
/// the base point of \a other away from \a this plane when considering whether
/// the planes are coincident.
VTKM_EXEC_CONT
bool Intersect(const Plane<CoordType>& other,
Ray<CoordType, 3, true>& ray,
bool& coincident,
CoordType tol2 = static_cast<CoordType>(1e-6f)) const;
};
/// Represent a sphere of the given \a Dimension.
/// If a constructor is given an invalid specification, then
/// the Radius of the resulting sphere will be -1.
template <typename CoordType = vtkm::FloatDefault, int Dim = 3>
struct Sphere
{
static constexpr int Dimension = Dim;
using Vector = vtkm::Vec<CoordType, Dim>;
Vector Center;
CoordType Radius;
/// Construct a default sphere (unit radius at the origin).
VTKM_EXEC_CONT Sphere();
/// Construct a sphere from a center point and radius.
VTKM_EXEC_CONT Sphere(const Vector& center, CoordType radius);
/// Return true if the sphere is valid (i.e., has a strictly positive radius).
VTKM_EXEC_CONT
bool IsValid() const { return this->Radius > 0.f; }
/// Return whether the point lies strictly inside the sphere.
VTKM_EXEC_CONT
bool Contains(const Vector& point, CoordType tol2 = 0.f) const;
/// Classify a point as inside (-1), on (0), or outside (+1) of the sphere.
///
/// The tolerance \a tol2 is the maximum allowable difference in squared
/// magnitude between the squared radius and the squared distance between
/// the \a point and Center.
VTKM_EXEC_CONT
int Classify(const Vector& point, CoordType tol2 = 0.f) const;
};
// -----------------------------------------------------------------------------
// Synonyms
//
// These "using" statements aim to make it easier to use the templated
// structs above when working with a particular dimension and/or the
// default floating-point type.
/// Lines are two-sided rays:
template <typename CoordType, int Dim = 3>
using Line = Ray<CoordType, Dim, true>;
// Shortcuts for 2D and 3D rays, lines, and line segments:
template <typename CoordType>
using Ray2 = Ray<CoordType, 2>;
template <typename CoordType>
using Ray3 = Ray<CoordType, 3>;
template <typename CoordType>
using Line2 = Line<CoordType, 2>;
template <typename CoordType>
using Line3 = Line<CoordType, 3>;
template <typename CoordType>
using LineSegment2 = LineSegment<CoordType, 2>;
template <typename CoordType>
using LineSegment3 = LineSegment<CoordType, 3>;
/// Circle is an alias for a 2-Dimensional sphere.
template <typename T>
using Circle = Sphere<T, 2>;
// Aliases for d-dimensional spheres.
template <typename T>
using Sphere2 = Sphere<T, 2>;
template <typename T>
using Sphere3 = Sphere<T, 3>;
// Shortcuts for default floating-point types
using Ray2d = Ray2<vtkm::FloatDefault>;
using Ray3d = Ray3<vtkm::FloatDefault>;
using Line2d = Line2<vtkm::FloatDefault>;
using Line3d = Line3<vtkm::FloatDefault>;
using LineSegment2d = LineSegment2<vtkm::FloatDefault>;
using LineSegment3d = LineSegment3<vtkm::FloatDefault>;
using Plane3d = Plane<vtkm::FloatDefault>;
using Circle2d = Circle<vtkm::FloatDefault>;
using Sphere2d = Sphere2<vtkm::FloatDefault>;
using Sphere3d = Sphere3<vtkm::FloatDefault>;
// -----------------------------------------------------------------------------
// Construction techniques
//
// These are free functions that create instances of geometric structs by taking
// in data that is not identical to the state of the struct and converting it
// into state for the struct.
/// Construct a plane from a point plus one of: a line, a ray, or a line segment.
///
/// The plane returned will contain the point and the line/ray/segment.
/// The plane normal will be the cross product of the line/ray/segment's direction
/// and the vector from the line/ray/segment's origin to the given \a point.
/// If the \a point is collinear with the line/ray/line-segment, an invalid
/// plane will be returned.
template <typename CoordType, bool IsTwoSided>
VTKM_EXEC_CONT vtkm::Plane<CoordType> make_PlaneFromPointAndLine(
const vtkm::Vec<CoordType, 3>& point,
const vtkm::Ray<CoordType, 3, IsTwoSided>& ray,
CoordType tol2 = static_cast<CoordType>(1e-8f));
template <typename CoordType>
VTKM_EXEC_CONT vtkm::Plane<CoordType> make_PlaneFromPointAndLineSegment(
const vtkm::Vec<CoordType, 3>& point,
const vtkm::LineSegment3<CoordType>& segment,
CoordType tol2 = static_cast<CoordType>(1e-8f));
/// Construct a circle from 3 points.
template <typename CoordType>
VTKM_EXEC_CONT vtkm::Circle<CoordType> make_CircleFrom3Points(
const typename vtkm::Vec<CoordType, 2>& p0,
const typename vtkm::Vec<CoordType, 2>& p1,
const typename vtkm::Vec<CoordType, 2>& p2,
CoordType tol = static_cast<CoordType>(1e-6f));
/// Construct a sphere from 4 points.
template <typename CoordType>
VTKM_EXEC_CONT vtkm::Sphere<CoordType, 3> make_SphereFrom4Points(
const vtkm::Vec<CoordType, 3>& a0,
const vtkm::Vec<CoordType, 3>& a1,
const vtkm::Vec<CoordType, 3>& a2,
const vtkm::Vec<CoordType, 3>& a3,
CoordType tol = static_cast<CoordType>(1e-6f));
} // namespace vtkm
#include "vtkm/Geometry.hxx"
#endif // vtk_m_Geometry_h

647
vtkm/Geometry.hxx Normal file
View File

@ -0,0 +1,647 @@
//============================================================================
// 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.
//============================================================================
namespace vtkm
{
// -----------------------------------------------------------------------------
// Ray
template <typename CoordType, int Dim, bool IsTwoSided>
template <int Dim_, typename std::enable_if<Dim_ == 2, int>::type>
Ray<CoordType, Dim, IsTwoSided>::Ray()
: Origin{ 0.f }
, Direction{ 1.f, 0.f }
{
}
template <typename CoordType, int Dim, bool IsTwoSided>
template <int Dim_, typename std::enable_if<Dim_ == 3, int>::type>
Ray<CoordType, Dim, IsTwoSided>::Ray()
: Origin{ 0.f }
, Direction{ 1.f, 0.f, 0.f }
{
}
template <typename CoordType, int Dim, bool IsTwoSided>
Ray<CoordType, Dim, IsTwoSided>::Ray(const LineSegment<CoordType, Dim>& segment)
: Origin(segment.Endpoints[0])
, Direction(vtkm::Normal(segment.Direction()))
{
}
template <typename CoordType, int Dim, bool IsTwoSided>
Ray<CoordType, Dim, IsTwoSided>::Ray(const Vector& point, const Vector& direction)
: Origin(point)
, Direction(vtkm::Normal(direction))
{
}
template <typename CoordType, int Dim, bool IsTwoSided>
typename Ray<CoordType, Dim, IsTwoSided>::Vector Ray<CoordType, Dim, IsTwoSided>::Evaluate(
CoordType param) const
{
auto pointOnLine = this->Origin + this->Direction * param;
return pointOnLine;
}
template <typename CoordType, int Dim, bool IsTwoSided>
bool Ray<CoordType, Dim, IsTwoSided>::IsValid() const
{
// At least on Ubuntu 17.10, cuda 9.1 will fail with an internal
// compiler error when calling vtkm::IsInf() here. But the fix
// below works. The fix should be removed as soon as our dashboards
// allow it.
#if __CUDACC_VER_MAJOR__ == 9 && __CUDACC_VER_MINOR__ == 1
return !isinf(this->Direction[0]);
#else
return !vtkm::IsInf(this->Direction[0]);
#endif
}
template <typename CoordType, int Dim, bool IsTwoSided>
CoordType Ray<CoordType, Dim, IsTwoSided>::DistanceTo(const Vector& point) const
{
Vector closest;
CoordType param;
return this->DistanceTo(point, param, closest);
}
template <typename CoordType, int Dim, bool IsTwoSided>
CoordType Ray<CoordType, Dim, IsTwoSided>::DistanceTo(const Vector& point,
CoordType& param,
Vector& projectedPoint) const
{
const auto& dir = this->Direction;
auto mag2 = vtkm::MagnitudeSquared(dir);
if (mag2 <= static_cast<CoordType>(0.0f))
{
// We have a point, not a line segment.
projectedPoint = this->Origin;
param = static_cast<CoordType>(0.0f);
return vtkm::Magnitude(point - this->Origin);
}
// Find the closest point on the line, then clamp to the ray if the
// parameter value is negative.
param = vtkm::Dot(point - this->Origin, dir) / mag2;
if (!TwoSided)
{
param = vtkm::Max(param, static_cast<CoordType>(0.0f));
}
// Compute the distance between the closest point and the input point.
projectedPoint = this->Evaluate(param);
auto dist = vtkm::Magnitude(point - projectedPoint);
return dist;
}
template <typename CoordType, int Dim, bool IsTwoSided>
template <bool OtherTwoSided, int Dim_, typename std::enable_if<Dim_ == 2, int>::type>
bool Ray<CoordType, Dim, IsTwoSided>::Intersect(const Ray<CoordType, Dim, OtherTwoSided>& other,
Vector& point,
CoordType tol)
{
auto d1 = this->Direction;
auto d2 = other.Direction;
auto denom = d1[0] * d2[1] - d1[1] * d2[0];
if (vtkm::Abs(denom) < tol)
{ // The lines are coincident or at least parallel.
return false;
}
const auto& a = this->Origin;
const auto& b = other.Origin;
CoordType numerU = a[1] * d2[0] + d2[1] * b[0] - b[1] * d2[0] - d2[1] * a[0];
CoordType uParam = numerU / denom;
point = a + uParam * d1;
if (TwoSided && OtherTwoSided)
{
return true;
}
else
{
CoordType numerV = d1[0] * (a[1] - b[1]) - d1[1] * (a[0] - b[0]);
CoordType vParam = numerV / denom;
return (TwoSided || (uParam + tol) > 0) && (OtherTwoSided || (vParam + tol) > 0);
}
}
// -----------------------------------------------------------------------------
// LineSegment
template <typename CoordType, int Dim>
template <int Dim_, typename std::enable_if<Dim_ == 2, int>::type>
LineSegment<CoordType, Dim>::LineSegment()
: Endpoints{ { 0.f }, { 1.f, 0.f } }
{
}
template <typename CoordType, int Dim>
template <int Dim_, typename std::enable_if<Dim_ == 3, int>::type>
LineSegment<CoordType, Dim>::LineSegment()
: Endpoints{ { 0.f }, { 1.f, 0.f, 0.f } }
{
}
template <typename CoordType, int Dim>
LineSegment<CoordType, Dim>::LineSegment(const Vector& p0, const Vector& p1)
: Endpoints{ p0, p1 }
{
}
template <typename CoordType, int Dim>
bool LineSegment<CoordType, Dim>::IsSingular(CoordType tol2) const
{
return vtkm::MagnitudeSquared(this->Direction()) < tol2;
}
template <typename CoordType, int Dim>
template <int Dim_, typename std::enable_if<Dim_ == 2, int>::type>
Ray<CoordType, Dim, true> LineSegment<CoordType, Dim>::PerpendicularBisector() const
{
const Vector dir = this->Direction();
const Vector perp(-dir[1], dir[0]);
const Vector mid = this->Center();
return Ray<CoordType, Dim, true>(mid, perp);
}
template <typename CoordType, int Dim>
template <int Dim_, typename std::enable_if<Dim_ == 3, int>::type>
Plane<CoordType> LineSegment<CoordType, Dim>::PerpendicularBisector() const
{
return Plane<CoordType>(this->Center(), this->Direction());
}
template <typename CoordType, int Dim>
typename LineSegment<CoordType, Dim>::Vector LineSegment<CoordType, Dim>::Evaluate(
CoordType param) const
{
auto pointOnLine = this->Endpoints[0] * (1.0f - param) + this->Endpoints[1] * param;
return pointOnLine;
}
template <typename CoordType, int Dim>
CoordType LineSegment<CoordType, Dim>::DistanceTo(const Vector& point) const
{
Vector closest;
CoordType param;
return this->DistanceTo(point, param, closest);
}
template <typename CoordType, int Dim>
CoordType LineSegment<CoordType, Dim>::DistanceTo(const Vector& point,
CoordType& param,
Vector& projectedPoint) const
{
auto dir = this->Endpoints[1] - this->Endpoints[0];
auto mag2 = vtkm::MagnitudeSquared(dir);
if (mag2 <= static_cast<CoordType>(0.0f))
{
// We have a point, not a line segment.
projectedPoint = this->Endpoints[0];
param = static_cast<CoordType>(0.0f);
return vtkm::Magnitude(point - this->Endpoints[0]);
}
// Find the closest point on the line, then clamp to the line segment
param = vtkm::Clamp(vtkm::Dot(point - this->Endpoints[0], dir) / mag2,
static_cast<CoordType>(0.0f),
static_cast<CoordType>(1.0f));
// Compute the distance between the closest point and the input point.
projectedPoint = this->Evaluate(param);
auto dist = vtkm::Magnitude(point - projectedPoint);
return dist;
}
template <typename CoordType, int Dim>
template <int Dim_, typename std::enable_if<Dim_ == 2, int>::type>
bool LineSegment<CoordType, Dim>::IntersectInfinite(const LineSegment<CoordType, Dim>& other,
Vector& point,
CoordType tol)
{
auto d1 = this->Direction();
auto d2 = other.Direction();
auto denom = d1[0] * d2[1] - d1[1] * d2[0];
if (vtkm::Abs(denom) < tol)
{ // The lines are coincident or at least parallel.
return false;
}
const auto& a = this->Endpoints;
const auto& b = other.Endpoints;
CoordType numerX = (a[0][0] * a[1][1] - a[0][1] * a[1][0]) * -d2[0] -
(b[0][0] * b[1][1] - b[0][1] * b[1][0]) * -d1[0];
CoordType numerY = (a[0][0] * a[1][1] - a[0][1] * a[1][0]) * -d2[1] -
(b[0][0] * b[1][1] - b[0][1] * b[1][0]) * -d1[1];
point = Vector(numerX / denom, numerY / denom);
return true;
}
// -----------------------------------------------------------------------------
// Plane
template <typename CoordType>
Plane<CoordType>::Plane()
: Origin{ 0.f, 0.f, 0.f }
, Normal{ 0.f, 0.f, 1.f }
{
}
template <typename CoordType>
Plane<CoordType>::Plane(const Vector& origin, const Vector& normal, CoordType tol2)
: Origin(origin)
, Normal(vtkm::Normal(normal))
{
if (tol2 > 0.0f && vtkm::MagnitudeSquared(normal) < tol2)
{
auto inf = vtkm::Infinity<CoordType>();
Normal = vtkm::Vec<CoordType, 3>(inf, inf, inf);
}
}
template <typename CoordType>
CoordType Plane<CoordType>::DistanceTo(const Vector& point) const
{
auto dist = vtkm::Dot(point - this->Origin, this->Normal);
return dist;
}
template <typename CoordType>
typename Plane<CoordType>::Vector Plane<CoordType>::ClosestPoint(const Vector& point) const
{
auto vop = vtkm::Project(point - this->Origin, this->Normal);
auto closest = point - vop;
return closest;
}
template <typename CoordType>
template <bool IsTwoSided>
bool Plane<CoordType>::Intersect(const Ray<CoordType, 3, IsTwoSided>& ray,
CoordType& parameter,
Vector& point,
bool& lineInPlane,
CoordType tol) const
{
CoordType d0 = this->DistanceTo(ray.Origin);
CoordType dirDot = vtkm::Dot(this->Normal, ray.Direction);
// If the ray/line lies parallel to the plane, the intersection is degenerate:
if (vtkm::Abs(dirDot) < tol)
{
lineInPlane = (vtkm::Abs(d0) < tol);
return false;
}
lineInPlane = false;
parameter = -d0 / dirDot;
// If we have a ray (not a line) and it points away from the
// side of the plane where its origin lies, then there is no
// intersection.
if (!IsTwoSided)
{
if (parameter < 0.0f)
{
return false;
}
}
// Check whether an endpoint lies in the plane:
if (vtkm::Abs(d0) < tol)
{
parameter = static_cast<CoordType>(0.0f);
point = ray.Origin;
return true;
}
// The perpendicular distance of the origin to the plane
// forms one side of a triangle whose hypotenuse is the
// parameter value (because ray.Direction has unit length).
// The dot product of the plane normal and ray direction
// is the cosine of the angle between the hypotenuse and
// the shortest path to the plane, so....
point = ray.Origin + parameter * ray.Direction;
return true;
}
template <typename CoordType>
bool Plane<CoordType>::Intersect(const LineSegment<CoordType>& segment,
CoordType& parameter,
bool& lineInPlane) const
{
Vector point;
return this->Intersect(segment, parameter, point, lineInPlane);
}
template <typename CoordType>
bool Plane<CoordType>::Intersect(const LineSegment<CoordType>& segment,
CoordType& parameter,
Vector& point,
bool& lineInPlane) const
{
CoordType d0 = this->DistanceTo(segment.Endpoints[0]);
CoordType d1 = this->DistanceTo(segment.Endpoints[1]);
if (d0 == 0 && d1 == 0)
{
// The entire segment lies in the plane.
lineInPlane = true;
return true;
}
lineInPlane = false;
// Check whether an endpoint lies in the plane:
if (d0 == 0)
{
parameter = static_cast<CoordType>(0.0f);
point = segment.Endpoints[0];
return true;
}
if (d1 == 0)
{
parameter = static_cast<CoordType>(1.0f);
point = segment.Endpoints[1];
return true;
}
// See whether endpoints lie on opposite sides of the plane:
//
// Note that the vtkm::SignBit comments below cause an internal compiler
// error on cuda 9.1, ubuntu 17.10 or they would be used:
bool c0 = d0 < 0; // !!vtkm::SignBit(d0);
bool c1 = d1 < 0; // !!vtkm::SignBit(d1);
CoordType a0 = vtkm::Abs(d0);
CoordType a1 = vtkm::Abs(d1);
if (c0 == c1)
{
// Both endpoints lie to the same side of the plane, so
// there is no intersection. Report the closest endpoint.
parameter = static_cast<CoordType>(a0 < a1 ? 0.0f : 1.0f);
point = segment.Endpoints[a0 < a1 ? 0 : 1];
return false;
}
// Endpoint distances have the opposite sign; there must be
// an intersection. It must occur at distance 0, and distance
// varies linearly from d0 to d1, so...
parameter = a0 / (a0 + a1);
point = segment.Endpoints[0] * (1.0f - parameter) + segment.Endpoints[1] * parameter;
return true;
}
template <typename CoordType>
bool Plane<CoordType>::Intersect(const Plane<CoordType>& other,
Ray<CoordType, 3, true>& ray,
bool& coincident,
CoordType tol2) const
{
auto dir = vtkm::Cross(this->Normal, other.Normal);
auto mag2 = vtkm::MagnitudeSquared(dir);
if (mag2 < tol2)
{ // The planes are parallel.
auto dist = this->DistanceTo(other.Origin);
coincident = dist * dist < tol2;
return false;
}
// The planes intersect. We want to find a point on the new plane
// and we want it to be near the other plane base points to avoid
// precision issues in the future.
// So, project each plane origin to the other plane along a line
// perpendicular to the plane and the output line. Both of these
// points are on the output line. Average the two points. The result
// will still be on the line and will be closer to the two base points.
auto nn = vtkm::Normal(dir);
auto moveDir01 = vtkm::Cross(this->Normal, nn);
auto moveDir02 = vtkm::Cross(other.Normal, nn);
Ray<CoordType, 3, true> bra(this->Origin, moveDir01);
Ray<CoordType, 3, true> brb(other.Origin, moveDir02);
vtkm::Vec<CoordType, 3> p0a;
vtkm::Vec<CoordType, 3> p0b;
CoordType pDummyA, pDummyB;
bool bDummyA, bDummyB;
auto tol = vtkm::Sqrt(tol2);
this->Intersect(brb, pDummyA, p0a, bDummyA, tol);
other.Intersect(bra, pDummyB, p0b, bDummyB, tol);
ray = vtkm::Ray<CoordType, 3, true>((p0a + p0b) * 0.5f, nn);
return true;
}
// -----------------------------------------------------------------------------
// Sphere
template <typename CoordType, int Dim>
Sphere<CoordType, Dim>::Sphere()
: Center{ 0.f }
, Radius(static_cast<CoordType>(1.f))
{
}
template <typename CoordType, int Dim>
Sphere<CoordType, Dim>::Sphere(const Vector& center, CoordType radius)
: Center(center)
, Radius(radius <= 0.f ? static_cast<CoordType>(-1.0f) : radius)
{
}
template <typename CoordType, int Dim>
bool Sphere<CoordType, Dim>::Contains(const Vector& point, CoordType tol2) const
{
return this->Classify(point, tol2) < 0;
}
template <typename CoordType, int Dim>
int Sphere<CoordType, Dim>::Classify(const Vector& point, CoordType tol2) const
{
if (!this->IsValid())
{
return 1; // All points are outside invalid spheres.
}
auto d2 = vtkm::MagnitudeSquared(point - this->Center);
auto r2 = this->Radius * this->Radius;
return d2 < r2 - tol2 ? -1 : (d2 > r2 + tol2 ? +1 : 0);
}
// -----------------------------------------------------------------------------
// Construction techniques
template <typename CoordType, bool IsTwoSided>
vtkm::Plane<CoordType> make_PlaneFromPointAndLine(const vtkm::Vec<CoordType, 3>& point,
const vtkm::Ray<CoordType, 3, IsTwoSided>& ray,
CoordType tol2)
{
auto tmpDir = point - ray.Origin;
return vtkm::Plane<CoordType>(point, vtkm::Cross(ray.Direction, tmpDir), tol2);
}
template <typename CoordType>
vtkm::Plane<CoordType> make_PlaneFromPointAndLineSegment(
const vtkm::Vec<CoordType, 3>& point,
const vtkm::LineSegment3<CoordType>& segment,
CoordType tol2)
{
auto tmpDir = point - segment.Endpoints[0];
return vtkm::Plane<CoordType>(point, vtkm::Cross(segment.Direction(), tmpDir), tol2);
}
template <typename CoordType>
vtkm::Circle<CoordType> make_CircleFrom3Points(const typename vtkm::Vec<CoordType, 2>& p0,
const typename vtkm::Vec<CoordType, 2>& p1,
const typename vtkm::Vec<CoordType, 2>& p2,
CoordType tol)
{
constexpr int Dim = 2;
using Vector = typename vtkm::Circle<CoordType>::Vector;
vtkm::LineSegment<CoordType, Dim> l01(p0, p1);
vtkm::LineSegment<CoordType, Dim> l02(p0, p2);
auto pb01 = l01.PerpendicularBisector();
auto pb02 = l02.PerpendicularBisector();
Vector center;
CoordType radius;
if (!pb01.IsValid() || !pb02.IsValid())
{
center = Vector(0.f, 0.f);
radius = -1.f;
return vtkm::Circle<CoordType>(center, radius);
}
auto isValid = pb01.Intersect(pb02, center, tol);
radius = isValid ? vtkm::Magnitude(center - p0) : static_cast<CoordType>(-1.0f);
if (!isValid)
{ // Initialize center to something.
center = Vector(vtkm::Nan<CoordType>(), vtkm::Nan<CoordType>());
}
return vtkm::Circle<CoordType>(center, radius);
}
template <typename CoordType>
vtkm::Sphere<CoordType, 3> make_SphereFrom4Points(const vtkm::Vec<CoordType, 3>& a0,
const vtkm::Vec<CoordType, 3>& a1,
const vtkm::Vec<CoordType, 3>& a2,
const vtkm::Vec<CoordType, 3>& a3,
CoordType tol)
{
// Choose p3 such that the min(p3 - p[012]) is larger than any other choice of p3.
// From: http://steve.hollasch.net/cgindex/geometry/sphere4pts.html,
// retrieved 2018-06-12 and paraphrased in terms of local variable names:
//
// If circlePointInPlaneOfP3-p3 is much smaller than
// circlePointInPlaneOfP3-circleCenterWorld, then the
// sphere center will be very close to circleCenterWorld
// and subject to error.
// It's best to choose p3 so that the least of p0-p3, p1-p3, and
// p2-p3 is larger than for any other.
CoordType d0 = vtkm::MagnitudeSquared(a1 - a0);
CoordType d1 = vtkm::MagnitudeSquared(a2 - a0);
CoordType d2 = vtkm::MagnitudeSquared(a3 - a0);
CoordType d3 = vtkm::MagnitudeSquared(a2 - a1);
CoordType d4 = vtkm::MagnitudeSquared(a3 - a1);
CoordType d5 = vtkm::MagnitudeSquared(a3 - a2);
CoordType sel0 = vtkm::Min(d0, vtkm::Min(d1, d2));
CoordType sel1 = vtkm::Min(d0, vtkm::Min(d3, d4));
CoordType sel2 = vtkm::Min(d1, vtkm::Min(d3, d5));
CoordType sel3 = vtkm::Min(d2, vtkm::Min(d4, d5));
CoordType selm = vtkm::Max(vtkm::Max(sel0, sel1), vtkm::Max(sel2, sel3));
vtkm::Vec<CoordType, 3> p0 = a0;
vtkm::Vec<CoordType, 3> p1 = a1;
vtkm::Vec<CoordType, 3> p2 = a2;
vtkm::Vec<CoordType, 3> p3 = a3;
if (sel0 == selm)
{
p3 = a0;
p0 = a3;
}
else if (sel1 == selm)
{
p3 = a1;
p1 = a3;
}
else if (sel2 == selm)
{
p3 = a2;
p2 = a3;
}
else // sel3 == selm
{
// do nothing.
}
vtkm::Vec<vtkm::Vec<CoordType, 3>, 3> axes;
axes[1] = p1 - p0;
axes[2] = p2 - p0;
axes[0] = vtkm::Cross(axes[1], axes[2]);
vtkm::Vec<vtkm::Vec<CoordType, 3>, 3> basis;
int rank = vtkm::Orthonormalize(axes, basis, tol);
if (rank < 3)
{
vtkm::Sphere<CoordType, 3> invalid;
invalid.Radius = -1.0f;
return invalid;
}
// Project points to plane and fit a circle.
auto p0_p = vtkm::Vec<CoordType, 2>{ 0.f }; // This is p0's new coordinate...
auto p1_p = vtkm::Vec<CoordType, 2>(vtkm::ProjectedDistance(axes[1], basis[1]),
vtkm::ProjectedDistance(axes[1], basis[2]));
auto p2_p = vtkm::Vec<CoordType, 2>(vtkm::ProjectedDistance(axes[2], basis[1]),
vtkm::ProjectedDistance(axes[2], basis[2]));
auto circle = make_CircleFrom3Points(p0_p, p1_p, p2_p);
if (!circle.IsValid())
{
vtkm::Sphere<CoordType, 3> invalid;
invalid.Radius = -1.0f;
return invalid;
}
vtkm::Vec<CoordType, 3> circleCenterWorld =
p0 + circle.Center[0] * basis[1] + circle.Center[1] * basis[2];
vtkm::Line3<CoordType> centerRay(circleCenterWorld, basis[0]);
// If our remaining unused point (p3) lines on centerRay,
// use one of the other points to locate the sphere's center:
vtkm::Vec<CoordType, 3> circlePointInPlaneOfP3;
if (vtkm::Abs(centerRay.DistanceTo(p3)) < tol)
{
circlePointInPlaneOfP3 = p0;
}
else
{
Plane<CoordType> pp3(circleCenterWorld, basis[0]);
circlePointInPlaneOfP3 =
circleCenterWorld + vtkm::Normal(pp3.ClosestPoint(p3) - circleCenterWorld) * circle.Radius;
}
auto bisectorPlane =
vtkm::LineSegment3<CoordType>(circlePointInPlaneOfP3, p3).PerpendicularBisector();
vtkm::Vec<CoordType, 3> sphereCenter;
CoordType param;
bool lineInPlane;
if (!bisectorPlane.Intersect(centerRay, param, sphereCenter, lineInPlane, tol))
{
vtkm::Sphere<CoordType, 3> invalid;
invalid.Radius = -1.0f;
return invalid;
}
auto sphereRadius = vtkm::Magnitude(sphereCenter - p3);
vtkm::Sphere<CoordType, 3> result(sphereCenter, sphereRadius);
;
return result;
}
} // namespace vtkm

30
vtkm/ImplicitFunction.h
View File

@ -146,7 +146,7 @@ public:
VTKM_EXEC_CONT const Vector& GetMaxPoint() const { return this->MaxPoint; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
VTKM_EXEC_CONT Scalar Value(const Vector& point) const final
{
Scalar minDistance = vtkm::NegativeInfinity32();
Scalar diff, t, dist;
@ -214,7 +214,7 @@ public:
}
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const final
{
vtkm::IdComponent minAxis = 0;
Scalar dist = 0.0;
@ -343,7 +343,7 @@ private:
//============================================================================
/// \brief Implicit function for a cylinder
class VTKM_ALWAYS_EXPORT Cylinder : public vtkm::ImplicitFunction
class VTKM_ALWAYS_EXPORT Cylinder final : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Cylinder()
@ -385,14 +385,14 @@ public:
this->Modified();
}
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
VTKM_EXEC_CONT Scalar Value(const Vector& point) const final
{
Vector x2c = point - this->Center;
FloatDefault proj = vtkm::Dot(this->Axis, x2c);
return vtkm::Dot(x2c, x2c) - (proj * proj) - (this->Radius * this->Radius);
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const final
{
Vector x2c = point - this->Center;
FloatDefault t = this->Axis[0] * x2c[0] + this->Axis[1] * x2c[1] + this->Axis[2] * x2c[2];
@ -409,7 +409,7 @@ private:
//============================================================================
/// \brief Implicit function for a frustum
class VTKM_ALWAYS_EXPORT Frustum : public vtkm::ImplicitFunction
class VTKM_ALWAYS_EXPORT Frustum final : public vtkm::ImplicitFunction
{
public:
Frustum() = default;
@ -480,7 +480,7 @@ public:
this->Modified();
}
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
VTKM_EXEC_CONT Scalar Value(const Vector& point) const final
{
Scalar maxVal = vtkm::NegativeInfinity<Scalar>();
for (vtkm::Id index : { 0, 1, 2, 3, 4, 5 })
@ -493,7 +493,7 @@ public:
return maxVal;
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const final
{
Scalar maxVal = vtkm::NegativeInfinity<Scalar>();
vtkm::Id maxValIdx = 0;
@ -518,7 +518,7 @@ private:
//============================================================================
/// \brief Implicit function for a plane
class VTKM_ALWAYS_EXPORT Plane : public vtkm::ImplicitFunction
class VTKM_ALWAYS_EXPORT Plane final : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Plane()
@ -554,12 +554,12 @@ public:
VTKM_EXEC_CONT const Vector& GetOrigin() const { return this->Origin; }
VTKM_EXEC_CONT const Vector& GetNormal() const { return this->Normal; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
VTKM_EXEC_CONT Scalar Value(const Vector& point) const final
{
return vtkm::Dot(point - this->Origin, this->Normal);
}
VTKM_EXEC_CONT Vector Gradient(const Vector&) const override { return this->Normal; }
VTKM_EXEC_CONT Vector Gradient(const Vector&) const final { return this->Normal; }
private:
Vector Origin;
@ -568,7 +568,7 @@ private:
//============================================================================
/// \brief Implicit function for a sphere
class VTKM_ALWAYS_EXPORT Sphere : public vtkm::ImplicitFunction
class VTKM_ALWAYS_EXPORT Sphere final : public vtkm::ImplicitFunction
{
public:
VTKM_EXEC_CONT Sphere()
@ -605,12 +605,12 @@ public:
VTKM_EXEC_CONT const Vector& GetCenter() const { return this->Center; }
VTKM_EXEC_CONT Scalar Value(const Vector& point) const override
VTKM_EXEC_CONT Scalar Value(const Vector& point) const final
{
return vtkm::MagnitudeSquared(point - this->Center) - (this->Radius * this->Radius);
}
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const override
VTKM_EXEC_CONT Vector Gradient(const Vector& point) const final
{
return Scalar(2) * (point - this->Center);
}
@ -625,7 +625,7 @@ private:
#ifdef VTKM_CUDA
// Cuda seems to have a bug where it expects the template class VirtualObjectTransfer
// to be instantiated in a consitent order among all the translation units of an
// to be instantiated in a consistent order among all the translation units of an
// executable. Failing to do so results in random crashes and incorrect results.
// We workaroud this issue by explicitly instantiating VirtualObjectTransfer for
// all the implicit functions here.

16
vtkm/ListTag.h
View File

@ -73,6 +73,22 @@ struct ListTagJoin : detail::ListRoot
using list = typename detail::ListJoin<typename ListTag1::list, typename ListTag2::list>::type;
};
/// A tag that is constructed by appending \c Type to \c ListTag.
template <typename ListTag, typename Type>
struct ListTagAppend : detail::ListRoot
{
using list = typename detail::ListJoin<typename ListTag::list, detail::ListBase<Type>>::type;
};
/// Append \c Type to \c ListTag only if \c ListTag does not already contain \c Type.
/// No checks are performed to see if \c ListTag itself has only unqiue elements.
template <typename ListTag, typename Type>
struct ListTagAppendUnique : detail::ListRoot
{
using list = typename detail::ListAppendUniqueImpl<typename ListTag::list, Type>::type;
};
/// A tag that consists of elements that are found in both tags. This struct
/// can be subclassed and still behave like a list tag.
template <typename ListTag1, typename ListTag2>

12
vtkm/Math.h
View File

@ -1815,6 +1815,18 @@ static inline VTKM_EXEC_CONT T Min(const T& x, const T& y)
return detail::Min(x, y, typename vtkm::TypeTraits<T>::DimensionalityTag());
}
/// Clamp \p x to the given range.
///
inline VTKM_EXEC_CONT vtkm::Float32 Clamp(vtkm::Float32 x, vtkm::Float32 lo, vtkm::Float32 hi)
{
return x > lo ? (x < hi ? x : hi) : lo;
}
inline VTKM_EXEC_CONT vtkm::Float64 Clamp(vtkm::Float64 x, vtkm::Float64 lo, vtkm::Float64 hi)
{
return x > lo ? (x < hi ? x : hi) : lo;
}
//-----------------------------------------------------------------------------
//#ifdef VTKM_CUDA

12
vtkm/Math.h.in
View File

@ -637,6 +637,18 @@ static inline VTKM_EXEC_CONT T Min(const T& x, const T& y)
return detail::Min(x, y, typename vtkm::TypeTraits<T>::DimensionalityTag());
}
/// Clamp \p x to the given range.
///
inline VTKM_EXEC_CONT vtkm::Float32 Clamp(vtkm::Float32 x, vtkm::Float32 lo, vtkm::Float32 hi)
{
return x > lo ? (x < hi ? x : hi) : lo;
}
inline VTKM_EXEC_CONT vtkm::Float64 Clamp(vtkm::Float64 x, vtkm::Float64 lo, vtkm::Float64 hi)
{
return x > lo ? (x < hi ? x : hi) : lo;
}
//-----------------------------------------------------------------------------
//#ifdef VTKM_CUDA

54
vtkm/Swap.h Normal file
View File

@ -0,0 +1,54 @@
//============================================================================
// 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.
//============================================================================
#ifndef vtk_m_Swap_h
#define vtk_m_Swap_h
#include <vtkm/internal/ExportMacros.h>
#ifdef __CUDACC__
#include <thrust/swap.h>
#else
#include <algorithm>
#endif
namespace vtkm
{
/// Performs a swap operation. Safe to call from cuda code.
#ifdef __CUDACC__
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)
{
using namespace thrust;
swap(a, b);
}
#else
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)
{
using namespace std;
swap(a, b);
}
#endif
} // end namespace vtkm
#endif //vtk_m_Swap_h

13
vtkm/TypeListTag.h
View File

@ -92,6 +92,19 @@ struct TypeListTagFieldVec4
{
};
/// A list containing common types for floating-point vectors. Specifically contains
/// 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>>
{
};
/// A list containing common types for values in fields. Specifically contains
/// floating point scalars and vectors of size 2, 3, and 4 with floating point
/// components.

105
vtkm/Types.h
View File

@ -613,11 +613,12 @@ public:
using ComponentType = T;
static constexpr vtkm::IdComponent NUM_COMPONENTS = Size;
protected:
VecBase() = default;
VTKM_EXEC_CONT
explicit VecBase(const ComponentType& value)
// The enable_if predicate will disable this constructor for Size=1 so that
// the variadic constructor constexpr VecBase(T, Ts&&...) is called instead.
template <vtkm::IdComponent Size2 = Size, typename std::enable_if<Size2 != 1, int>::type = 0>
VTKM_EXEC_CONT explicit VecBase(const ComponentType& value)
{
for (vtkm::IdComponent i = 0; i < Size; ++i)
{
@ -625,6 +626,38 @@ protected:
}
}
template <typename... Ts>
VTKM_EXEC_CONT constexpr VecBase(ComponentType value0, Ts&&... values)
: Components{ value0, values... }
{
VTKM_STATIC_ASSERT(sizeof...(Ts) + 1 == Size);
}
VTKM_EXEC_CONT
VecBase(std::initializer_list<ComponentType> values)
{
ComponentType* dest = this->Components;
auto src = values.begin();
if (values.size() == 1)
{
for (vtkm::IdComponent i = 0; i < Size; ++i)
{
this->Components[i] = *src;
++dest;
}
}
else
{
VTKM_ASSERT((values.size() == NUM_COMPONENTS) &&
"Vec object initialized wrong number of components.");
for (; src != values.end(); ++src)
{
*dest = *src;
++dest;
}
}
}
#if (!(defined(VTKM_CUDA) && (__CUDACC_VER_MAJOR__ < 8)))
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
#pragma GCC diagnostic push
@ -787,17 +820,15 @@ public:
static constexpr vtkm::IdComponent NUM_COMPONENTS = Size;
#endif
using Superclass::Superclass;
Vec() = default;
VTKM_EXEC_CONT explicit Vec(const T& value)
: Superclass(value)
{
}
template <typename OtherType>
VTKM_EXEC_CONT explicit Vec(const Vec<OtherType, Size>& src)
: Superclass(src)
#if defined(_MSC_VER) && _MSC_VER < 1910
template <typename... Ts>
constexpr Vec(T value, Ts&&... values)
: Superclass(value, std::forward<Ts>(values)...)
{
}
#endif
inline VTKM_EXEC_CONT void CopyInto(Vec<T, Size>& dest) const { dest = *this; }
};
@ -847,7 +878,7 @@ class VTKM_ALWAYS_EXPORT Vec<T, 1> : public detail::VecBase<T, 1, Vec<T, 1>>
public:
Vec() = default;
VTKM_EXEC_CONT explicit Vec(const T& value)
VTKM_EXEC_CONT constexpr Vec(const T& value)
: Superclass(value)
{
}
@ -869,7 +900,7 @@ class VTKM_ALWAYS_EXPORT Vec<T, 2> : public detail::VecBase<T, 2, Vec<T, 2>>
public:
Vec() = default;
VTKM_EXEC_CONT explicit Vec(const T& value)
VTKM_EXEC_CONT Vec(const T& value)
: Superclass(value)
{
}
@ -881,10 +912,9 @@ public:
}
VTKM_EXEC_CONT
Vec(const T& x, const T& y)
constexpr Vec(const T& x, const T& y)
: Superclass(x, y)
{
this->Components[0] = x;
this->Components[1] = y;
}
};
@ -898,7 +928,7 @@ class VTKM_ALWAYS_EXPORT Vec<T, 3> : public detail::VecBase<T, 3, Vec<T, 3>>
public:
Vec() = default;
VTKM_EXEC_CONT explicit Vec(const T& value)
VTKM_EXEC_CONT Vec(const T& value)
: Superclass(value)
{
}
@ -910,11 +940,9 @@ public:
}
VTKM_EXEC_CONT
Vec(const T& x, const T& y, const T& z)
constexpr Vec(const T& x, const T& y, const T& z)
: Superclass(x, y, z)
{
this->Components[0] = x;
this->Components[1] = y;
this->Components[2] = z;
}
};
@ -929,7 +957,7 @@ class VTKM_ALWAYS_EXPORT Vec<T, 4> : public detail::VecBase<T, 4, Vec<T, 4>>
public:
Vec() = default;
VTKM_EXEC_CONT explicit Vec(const T& value)
VTKM_EXEC_CONT Vec(const T& value)
: Superclass(value)
{
}
@ -941,37 +969,20 @@ public:
}
VTKM_EXEC_CONT
Vec(const T& x, const T& y, const T& z, const T& w)
constexpr Vec(const T& x, const T& y, const T& z, const T& w)
: Superclass(x, y, z, w)
{
this->Components[0] = x;
this->Components[1] = y;
this->Components[2] = z;
this->Components[3] = w;
}
};
/// Initializes and returns a Vec of length 2.
/// Initializes and returns a Vec containing all the arguments. The arguments should all be the
/// same type or compile issues will occur.
///
template <typename T>
VTKM_EXEC_CONT vtkm::Vec<T, 2> make_Vec(const T& x, const T& y)
template <typename T, typename... Ts>
VTKM_EXEC_CONT constexpr vtkm::Vec<T, vtkm::IdComponent(sizeof...(Ts) + 1)> make_Vec(T value0,
Ts&&... args)
{
return vtkm::Vec<T, 2>(x, y);
}
/// Initializes and returns a Vec of length 3.
///
template <typename T>
VTKM_EXEC_CONT vtkm::Vec<T, 3> make_Vec(const T& x, const T& y, const T& z)
{
return vtkm::Vec<T, 3>(x, y, z);
}
/// Initializes and returns a Vec of length 4.
///
template <typename T>
VTKM_EXEC_CONT vtkm::Vec<T, 4> make_Vec(const T& x, const T& y, const T& z, const T& w)
{
return vtkm::Vec<T, 4>(x, y, z, w);
return vtkm::Vec<T, vtkm::IdComponent(sizeof...(Ts) + 1)>(value0, T(args)...);
}
/// \brief A Vec-like representation for short arrays.

78
vtkm/VectorAnalysis.h
View File

@ -208,6 +208,84 @@ TriangleNormal(const vtkm::Vec<T, 3>& a, const vtkm::Vec<T, 3>& b, const vtkm::V
return vtkm::Cross(b - a, c - a);
}
//-----------------------------------------------------------------------------
/// \brief Project a vector onto another vector.
///
/// This method computes the orthogonal projection of the vector v onto u;
/// that is, it projects its first argument onto its second.
///
/// Note that if the vector \a u has zero length, the output
/// vector will have all its entries equal to NaN.
template <typename T, int N>
VTKM_EXEC_CONT vtkm::Vec<T, N> Project(const vtkm::Vec<T, N>& v, const vtkm::Vec<T, N>& u)
{
T uu = vtkm::Dot(u, u);
T uv = vtkm::Dot(u, v);
T factor = uv / uu;
vtkm::Vec<T, N> result = factor * u;
return result;
}
//-----------------------------------------------------------------------------
/// \brief Project a vector onto another vector, returning only the projected distance.
///
/// This method computes the orthogonal projection of the vector v onto u;
/// that is, it projects its first argument onto its second.
///
/// Note that if the vector \a u has zero length, the output will be NaN.
template <typename T, int N>
VTKM_EXEC_CONT T ProjectedDistance(const vtkm::Vec<T, N>& v, const vtkm::Vec<T, N>& u)
{
T uu = vtkm::Dot(u, u);
T uv = vtkm::Dot(u, v);
T factor = uv / uu;
return factor;
}
//-----------------------------------------------------------------------------
/// \brief Perform Gram-Schmidt orthonormalization for 3-D vectors.
///
/// See https://en.wikipedia.org/wiki/Gram%E2%80%93Schmidt_process for details.
/// The first output vector will always be parallel to the first input vector.
/// The remaining output vectors will be orthogonal and unit length and have
/// the same handedness as their corresponding input vectors.
///
/// This method is geometric.
/// It does not require a matrix solver.
/// However, unlike the algebraic eigensolver techniques which do use matrix
/// inversion, this method may return zero-length output vectors if some input
/// vectors are collinear. The number of non-zero (to within the specified
/// tolerance, \a tol ) output vectors is the return value.
template <typename T, int N>
VTKM_EXEC_CONT int Orthonormalize(const vtkm::Vec<vtkm::Vec<T, N>, N>& inputs,
vtkm::Vec<vtkm::Vec<T, N>, N>& outputs,
T tol = static_cast<T>(1e-6))
{
int j = 0; // j is the number of non-zero-length, non-collinear inputs encountered.
vtkm::Vec<vtkm::Vec<T, N>, N> u;
for (int i = 0; i < N; ++i)
{
u[j] = inputs[i];
for (int k = 0; k < j; ++k)
{
u[j] -= vtkm::Project(inputs[i], u[k]);
}
T rmag = vtkm::RMagnitude(u[j]);
if (rmag * tol > 1.0)
{
// skip this vector, it is zero-length or collinear with others.
continue;
}
outputs[j] = rmag * u[j];
++j;
}
for (int i = j; i < N; ++i)
{
outputs[j] = Vec<T, N>{ 0. };
}
return j;
}
} // namespace vtkm
#endif //vtk_m_VectorAnalysis_h

2
vtkm/cont/Algorithm.h
View File

@ -271,7 +271,7 @@ struct UpperBoundsFunctor
return true;
}
};
} // annonymous namespace
} // anonymous namespace
struct Algorithm
{

9
vtkm/cont/ArrayHandle.h
View File

@ -82,21 +82,18 @@ struct IsInValidArrayHandle
{
};
/// Checks to see if the ArrayHandle for the given DeviceAdatper allows
/// Checks to see if the ArrayHandle allows
/// writing, as some ArrayHandles (Implicit) don't support writing.
/// This check is compatible with the C++11 type_traits.
/// It contains a typedef named type that is either
/// std::true_type or std::false_type.
/// Both of these have a typedef named value with the respective boolean value.
///
template <typename ArrayHandle, typename DeviceAdapterTag>
template <typename ArrayHandle>
struct IsWriteableArrayHandle
{
private:
template <typename T>
using ExecutionTypes = typename ArrayHandle::template ExecutionTypes<T>;
using ValueType = typename ExecutionTypes<DeviceAdapterTag>::Portal::ValueType;
using ValueType = typename ArrayHandle::PortalControl::ValueType;
//All ArrayHandles that use ImplicitStorage as the final writable location
//will have a value type of void*, which is what we are trying to detect

9
vtkm/cont/ArrayHandleSwizzle.h
View File

@ -386,6 +386,15 @@ VTKM_CONT ArrayHandleSwizzle<ArrayHandleType, OutSize> make_ArrayHandleSwizzle(
{
return ArrayHandleSwizzle<ArrayHandleType, OutSize>(array, map);
}
template <typename ArrayHandleType, typename... SwizzleIndexTypes>
VTKM_CONT ArrayHandleSwizzle<ArrayHandleType, vtkm::IdComponent(sizeof...(SwizzleIndexTypes) + 1)>
make_ArrayHandleSwizzle(const ArrayHandleType& array,
vtkm::IdComponent swizzleIndex0,
SwizzleIndexTypes... swizzleIndices)
{
return make_ArrayHandleSwizzle(array, vtkm::make_Vec(swizzleIndex0, swizzleIndices...));
}
}
} // namespace vtkm::cont

9
vtkm/cont/ArrayHandleVirtualCoordinates.h
View File

@ -181,6 +181,7 @@ public:
VTKM_CONT virtual void Allocate(vtkm::Id numberOfValues) = 0;
VTKM_CONT virtual void Shrink(vtkm::Id numberOfValues) = 0;
VTKM_CONT virtual void ReleaseResources() = 0;
VTKM_CONT virtual void ReleaseResourcesExecution() = 0;
VTKM_CONT virtual PortalConst PrepareForInput(vtkm::cont::DeviceAdapterId deviceId) = 0;
VTKM_CONT virtual Portal PrepareForOutput(vtkm::Id numberOfValues,
@ -241,6 +242,8 @@ public:
VTKM_CONT void ReleaseResources() override { this->Array.ReleaseResources(); }
VTKM_CONT void ReleaseResourcesExecution() override { this->Array.ReleaseResourcesExecution(); }
VTKM_CONT PortalConst PrepareForInput(vtkm::cont::DeviceAdapterId deviceId) override
{
PortalConst portal;
@ -406,8 +409,10 @@ public:
VTKM_CONT
void Shrink(vtkm::Id numberOfValues) { this->Array->Shrink(numberOfValues); }
// ArrayTransfer should only be capable of releasing resources in the execution
// environment
VTKM_CONT
void ReleaseResources() { this->Array->ReleaseResources(); }
void ReleaseResources() { this->Array->ReleaseResourcesExecution(); }
private:
CoordinatesArrayHandleBase* Array;
@ -513,7 +518,7 @@ void CastAndCall(const typename vtkm::cont::ArrayHandleVirtualCoordinates::Super
#ifdef VTKM_CUDA
// Cuda seems to have a bug where it expects the template class VirtualObjectTransfer
// to be instantiated in a consitent order among all the translation units of an
// to be instantiated in a consistent order among all the translation units of an
// executable. Failing to do so results in random crashes and incorrect results.
// We workaroud this issue by explicitly instantiating VirtualObjectTransfer for
// all the portal types here.

7
vtkm/cont/CMakeLists.txt
View File

@ -119,6 +119,9 @@ set(sources
BoundsGlobalCompute.cxx
CellSet.cxx
CellSetStructured.cxx
DataSetBuilderExplicit.cxx
DataSetBuilderRectilinear.cxx
DataSetBuilderUniform.cxx
DynamicArrayHandle.cxx
EnvironmentTracker.cxx
Field.cxx
@ -156,6 +159,7 @@ add_subdirectory(arg)
add_subdirectory(diy)
add_subdirectory(serial)
add_subdirectory(tbb)
add_subdirectory(openmp)
add_subdirectory(cuda)
set(backends )
@ -165,6 +169,9 @@ endif()
if(TARGET vtkm::cuda)
list(APPEND backends vtkm::cuda)
endif()
if(TARGET vtkm::openmp)
list(APPEND backends vtkm::openmp)
endif()
target_link_libraries(vtkm_cont PUBLIC vtkm_compiler_flags ${backends})
if(TARGET vtkm_diy)
# This will become a required dependency eventually.

2
vtkm/cont/CellSet.h
View File

@ -70,6 +70,8 @@ public:
virtual void PrintSummary(std::ostream&) const = 0;
virtual void ReleaseResourcesExecution() = 0;
protected:
std::string Name;
};

1
vtkm/cont/CellSetExplicit.h
View File

@ -111,6 +111,7 @@ public:
vtkm::Id GetNumberOfFaces() const override;
vtkm::Id GetNumberOfEdges() const override;
void PrintSummary(std::ostream& out) const override;
void ReleaseResourcesExecution() override;
VTKM_CONT vtkm::Id GetSchedulingRange(vtkm::TopologyElementTagCell) const;
VTKM_CONT vtkm::Id GetSchedulingRange(vtkm::TopologyElementTagPoint) const;

13
vtkm/cont/CellSetExplicit.hxx
View File

@ -101,6 +101,19 @@ void CellSetExplicit<ShapeStorageTag,
this->CellToPoint.PrintSummary(out);
}
template <typename ShapeStorageTag,
typename NumIndicesStorageTag,
typename ConnectivityStorageTag,
typename OffsetStorageTag>
void CellSetExplicit<ShapeStorageTag,
NumIndicesStorageTag,
ConnectivityStorageTag,
OffsetStorageTag>::ReleaseResourcesExecution()
{
this->PointToCell.ReleaseResourcesExecution();
this->CellToPoint.ReleaseResourcesExecution();
}
//----------------------------------------------------------------------------
template <typename ShapeStorageTag,

9
vtkm/cont/CellSetPermutation.h
View File

@ -205,6 +205,15 @@ public:
VTKM_CONT
vtkm::Id GetNumberOfEdges() const override { return -1; }
VTKM_CONT
void ReleaseResourcesExecution() override
{
this->ValidCellIds.ReleaseResourcesExecution();
this->FullCellSet.ReleaseResourcesExecution();
this->CellToPoint.ReleaseResourcesExecution();
}
VTKM_CONT
vtkm::IdComponent GetNumberOfPointsInCell(vtkm::Id cellIndex) const
{

15
vtkm/cont/CellSetStructured.h
View File

@ -34,7 +34,7 @@ namespace cont
{
template <vtkm::IdComponent DIMENSION>
class VTKM_ALWAYS_EXPORT CellSetStructured : public CellSet
class VTKM_ALWAYS_EXPORT CellSetStructured final : public CellSet
{
private:
using Thisclass = vtkm::cont::CellSetStructured<DIMENSION>;
@ -55,13 +55,16 @@ public:
Thisclass& operator=(const Thisclass& src);
virtual vtkm::Id GetNumberOfCells() const { return this->Structure.GetNumberOfCells(); }
vtkm::Id GetNumberOfCells() const override { return this->Structure.GetNumberOfCells(); }
virtual vtkm::Id GetNumberOfPoints() const { return this->Structure.GetNumberOfPoints(); }
vtkm::Id GetNumberOfPoints() const override { return this->Structure.GetNumberOfPoints(); }
virtual vtkm::Id GetNumberOfFaces() const { return -1; }
vtkm::Id GetNumberOfFaces() const override { return -1; }
virtual vtkm::Id GetNumberOfEdges() const { return -1; }
vtkm::Id GetNumberOfEdges() const override { return -1; }
// Since the entire topology is defined by by three integers, nothing to do here.
void ReleaseResourcesExecution() override {}
void SetPointDimensions(SchedulingRangeType dimensions)
{
@ -95,7 +98,7 @@ public:
typename ExecutionTypes<DeviceAdapter, FromTopology, ToTopology>::ExecObjectType
PrepareForInput(DeviceAdapter, FromTopology, ToTopology) const;
virtual void PrintSummary(std::ostream& out) const;
void PrintSummary(std::ostream& out) const override;
private:
InternalsType Structure;

2
vtkm/cont/ColorTable.cxx
View File

@ -197,7 +197,7 @@ void ColorTable::SetColorSpace(ColorSpace space)
break;
}
default:
throw vtkm::cont::ErrorBadType("unkown vtkm::cont::ColorType requested");
throw vtkm::cont::ErrorBadType("unknown vtkm::cont::ColorType requested");
}
}
}

8
vtkm/cont/ColorTable.h
View File

@ -605,7 +605,7 @@ public:
/// Will use the current range of the color table to generate evenly spaced
/// values using either vtkm::Float32 or vtkm::Float64 space.
/// Will use vtkm::Float32 space when the difference between the float and double
/// values when the range is withing float space and the following are within a tolerance:
/// values when the range is within float space and the following are within a tolerance:
///
/// - (max-min) / numSamples
/// - ((max-min) / numSamples) * numSamples
@ -620,7 +620,7 @@ public:
/// Will use the current range of the color table to generate evenly spaced
/// values using either vtkm::Float32 or vtkm::Float64 space.
/// Will use vtkm::Float32 space when the difference between the float and double
/// values when the range is withing float space and the following are within a tolerance:
/// values when the range is within float space and the following are within a tolerance:
///
/// - (max-min) / numSamples
/// - ((max-min) / numSamples) * numSamples
@ -635,7 +635,7 @@ public:
/// Will use the current range of the color table to generate evenly spaced
/// values using either vtkm::Float32 or vtkm::Float64 space.
/// Will use vtkm::Float32 space when the difference between the float and double
/// values when the range is withing float space and the following are within a tolerance:
/// values when the range is within float space and the following are within a tolerance:
///
/// - (max-min) / numSamples
/// - ((max-min) / numSamples) * numSamples
@ -650,7 +650,7 @@ public:
/// Will use the current range of the color table to generate evenly spaced
/// values using either vtkm::Float32 or vtkm::Float64 space.
/// Will use vtkm::Float32 space when the difference between the float and double
/// values when the range is withing float space and the following are within a tolerance:
/// values when the range is within float space and the following are within a tolerance:
///
/// - (max-min) / numSamples
/// - ((max-min) / numSamples) * numSamples

10
vtkm/cont/CoordinateSystem.h
View File

@ -81,7 +81,15 @@ public:
VTKM_CONT
vtkm::Bounds GetBounds() const;
virtual void PrintSummary(std::ostream& out) const;
virtual void PrintSummary(std::ostream& out) const override;
/// Releases any resources being used in the execution environment (that are
/// not being shared by the control environment).
VTKM_CONT void ReleaseResourcesExecution() override
{
this->Superclass::ReleaseResourcesExecution();
this->GetData().ReleaseResourcesExecution();
}
};
template <typename Functor, typename... Args>

110
vtkm/cont/DataSetBuilderExplicit.cxx Normal file
View File

@ -0,0 +1,110 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/DataSetBuilderExplicit.h>
namespace vtkm
{
namespace cont
{
VTKM_CONT
DataSetBuilderExplicitIterative::DataSetBuilderExplicitIterative()
{
}
VTKM_CONT
void DataSetBuilderExplicitIterative::Begin(const std::string& coordName,
const std::string& cellName)
{
this->coordNm = coordName;
this->cellNm = cellName;
this->points.resize(0);
this->shapes.resize(0);
this->numIdx.resize(0);
this->connectivity.resize(0);
}
//Define points.
VTKM_CONT
vtkm::cont::DataSet DataSetBuilderExplicitIterative::Create()
{
DataSetBuilderExplicit dsb;
return dsb.Create(points, shapes, numIdx, connectivity, coordNm, cellNm);
}
VTKM_CONT
vtkm::Id DataSetBuilderExplicitIterative::AddPoint(const vtkm::Vec<vtkm::Float32, 3>& pt)
{
points.push_back(pt);
vtkm::Id id = static_cast<vtkm::Id>(points.size());
return id;
}
VTKM_CONT
vtkm::Id DataSetBuilderExplicitIterative::AddPoint(const vtkm::Float32& x,
const vtkm::Float32& y,
const vtkm::Float32& z)
{
points.push_back(vtkm::make_Vec(x, y, z));
vtkm::Id id = static_cast<vtkm::Id>(points.size());
return id;
}
//Define cells.
VTKM_CONT
void DataSetBuilderExplicitIterative::AddCell(vtkm::UInt8 shape)
{
this->shapes.push_back(shape);
this->numIdx.push_back(0);
}
VTKM_CONT
void DataSetBuilderExplicitIterative::AddCell(const vtkm::UInt8& shape,
const std::vector<vtkm::Id>& conn)
{
this->shapes.push_back(shape);
this->numIdx.push_back(static_cast<vtkm::IdComponent>(conn.size()));
connectivity.insert(connectivity.end(), conn.begin(), conn.end());
}
VTKM_CONT
void DataSetBuilderExplicitIterative::AddCell(const vtkm::UInt8& shape,
const vtkm::Id* conn,
const vtkm::IdComponent& n)
{
this->shapes.push_back(shape);
this->numIdx.push_back(n);
for (int i = 0; i < n; i++)
{
connectivity.push_back(conn[i]);
}
}
VTKM_CONT
void DataSetBuilderExplicitIterative::AddCellPoint(vtkm::Id pointIndex)
{
VTKM_ASSERT(this->numIdx.size() > 0);
this->connectivity.push_back(pointIndex);
this->numIdx.back() += 1;
}
}
} // end namespace vtkm::cont

66
vtkm/cont/DataSetBuilderExplicit.h
View File

@ -33,7 +33,7 @@ namespace cont
//Coordinates builder??
//Need a singlecellset handler.
class DataSetBuilderExplicit
class VTKM_CONT_EXPORT DataSetBuilderExplicit
{
template <typename T>
VTKM_CONT static void CopyInto(const std::vector<T>& input, vtkm::cont::ArrayHandle<T>& output)
@ -313,42 +313,24 @@ inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicit::BuildDataSet(
return dataSet;
}
class DataSetBuilderExplicitIterative
class VTKM_CONT_EXPORT DataSetBuilderExplicitIterative
{
public:
VTKM_CONT
DataSetBuilderExplicitIterative() {}
DataSetBuilderExplicitIterative();
VTKM_CONT
void Begin(const std::string& coordName = "coords", const std::string& cellName = "cells")
{
this->coordNm = coordName;
this->cellNm = cellName;
this->points.resize(0);
this->shapes.resize(0);
this->numIdx.resize(0);
this->connectivity.resize(0);
}
void Begin(const std::string& coordName = "coords", const std::string& cellName = "cells");
//Define points.
VTKM_CONT
vtkm::cont::DataSet Create();
VTKM_CONT
vtkm::Id AddPoint(const vtkm::Vec<vtkm::Float32, 3>& pt)
{
points.push_back(pt);
vtkm::Id id = static_cast<vtkm::Id>(points.size());
return id;
}
vtkm::Id AddPoint(const vtkm::Vec<vtkm::Float32, 3>& pt);
VTKM_CONT
vtkm::Id AddPoint(const vtkm::Float32& x, const vtkm::Float32& y, const vtkm::Float32& z = 0)
{
points.push_back(vtkm::make_Vec(x, y, z));
vtkm::Id id = static_cast<vtkm::Id>(points.size());
return id;
}
vtkm::Id AddPoint(const vtkm::Float32& x, const vtkm::Float32& y, const vtkm::Float32& z = 0);
template <typename T>
VTKM_CONT vtkm::Id AddPoint(const T& x, const T& y, const T& z = 0)
@ -365,38 +347,16 @@ public:
//Define cells.
VTKM_CONT
void AddCell(vtkm::UInt8 shape)
{
this->shapes.push_back(shape);
this->numIdx.push_back(0);
}
void AddCell(vtkm::UInt8 shape);
VTKM_CONT
void AddCell(const vtkm::UInt8& shape, const std::vector<vtkm::Id>& conn)
{
this->shapes.push_back(shape);
this->numIdx.push_back(static_cast<vtkm::IdComponent>(conn.size()));
connectivity.insert(connectivity.end(), conn.begin(), conn.end());
}
void AddCell(const vtkm::UInt8& shape, const std::vector<vtkm::Id>& conn);
VTKM_CONT
void AddCell(const vtkm::UInt8& shape, const vtkm::Id* conn, const vtkm::IdComponent& n)
{
this->shapes.push_back(shape);
this->numIdx.push_back(n);
for (int i = 0; i < n; i++)
{
connectivity.push_back(conn[i]);
}
}
void AddCell(const vtkm::UInt8& shape, const vtkm::Id* conn, const vtkm::IdComponent& n);
VTKM_CONT
void AddCellPoint(vtkm::Id pointIndex)
{
VTKM_ASSERT(this->numIdx.size() > 0);
this->connectivity.push_back(pointIndex);
this->numIdx.back() += 1;
}
void AddCellPoint(vtkm::Id pointIndex);
private:
std::string coordNm, cellNm;
@ -406,12 +366,6 @@ private:
std::vector<vtkm::IdComponent> numIdx;
std::vector<vtkm::Id> connectivity;
};
inline VTKM_CONT vtkm::cont::DataSet DataSetBuilderExplicitIterative::Create()
{
DataSetBuilderExplicit dsb;
return dsb.Create(points, shapes, numIdx, connectivity, coordNm, cellNm);
}
}
}

33
vtkm/cont/DataSetBuilderRectilinear.cxx Normal file
View File

@ -0,0 +1,33 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/DataSetBuilderRectilinear.h>
namespace vtkm
{
namespace cont
{
VTKM_CONT
DataSetBuilderRectilinear::DataSetBuilderRectilinear()
{
}
}
} // end namespace vtkm::cont

6
vtkm/cont/DataSetBuilderRectilinear.h
View File

@ -31,7 +31,7 @@ namespace vtkm
namespace cont
{
class DataSetBuilderRectilinear
class VTKM_CONT_EXPORT DataSetBuilderRectilinear
{
template <typename T, typename U>
VTKM_CONT static void CopyInto(const std::vector<T>& input, vtkm::cont::ArrayHandle<U>& output)
@ -55,7 +55,7 @@ class DataSetBuilderRectilinear
public:
VTKM_CONT
DataSetBuilderRectilinear() {}
DataSetBuilderRectilinear();
//1D grids.
template <typename T>
@ -220,7 +220,7 @@ private:
vtkm::cont::CoordinateSystem cs(coordNm, coords);
dataSet.AddCoordinateSystem(cs);
// compute the dimensions of the cellset by counting the number of axises
// compute the dimensions of the cellset by counting the number of axes
// with >1 dimension
int ndims = 0;
vtkm::Id dims[3];

115
vtkm/cont/DataSetBuilderUniform.cxx Normal file
View File

@ -0,0 +1,115 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/DataSetBuilderUniform.h>
namespace vtkm
{
namespace cont
{
VTKM_CONT
DataSetBuilderUniform::DataSetBuilderUniform()
{
}
VTKM_CONT
vtkm::cont::DataSet DataSetBuilderUniform::Create(const vtkm::Id& dimension,
std::string coordNm,
std::string cellNm)
{
return CreateDataSet(vtkm::Id3(dimension, 1, 1), VecType(0), VecType(1), coordNm, cellNm);
}
VTKM_CONT
vtkm::cont::DataSet DataSetBuilderUniform::Create(const vtkm::Id2& dimensions,
std::string coordNm,
std::string cellNm)
{
return CreateDataSet(
vtkm::Id3(dimensions[0], dimensions[1], 1), VecType(0), VecType(1), coordNm, cellNm);
}
VTKM_CONT
vtkm::cont::DataSet DataSetBuilderUniform::Create(const vtkm::Id3& dimensions,
std::string coordNm,
std::string cellNm)
{
return CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], dimensions[2]),
VecType(0),
VecType(1),
coordNm,
cellNm);
}
VTKM_CONT
vtkm::cont::DataSet DataSetBuilderUniform::CreateDataSet(
const vtkm::Id3& dimensions,
const vtkm::Vec<vtkm::FloatDefault, 3>& origin,
const vtkm::Vec<vtkm::FloatDefault, 3>& spacing,
std::string coordNm,
std::string cellNm)
{
vtkm::Id dims[3] = { 1, 1, 1 };
int ndims = 0;
for (int i = 0; i < 3; ++i)
{
if (dimensions[i] > 1)
{
if (spacing[i] <= 0.0f)
{
throw vtkm::cont::ErrorBadValue("spacing must be > 0.0");
}
dims[ndims++] = dimensions[i];
}
}
vtkm::cont::DataSet dataSet;
vtkm::cont::ArrayHandleUniformPointCoordinates coords(dimensions, origin, spacing);
vtkm::cont::CoordinateSystem cs(coordNm, coords);
dataSet.AddCoordinateSystem(cs);
if (ndims == 1)
{
vtkm::cont::CellSetStructured<1> cellSet(cellNm);
cellSet.SetPointDimensions(dims[0]);
dataSet.AddCellSet(cellSet);
}
else if (ndims == 2)
{
vtkm::cont::CellSetStructured<2> cellSet(cellNm);
cellSet.SetPointDimensions(vtkm::Id2(dims[0], dims[1]));
dataSet.AddCellSet(cellSet);
}
else if (ndims == 3)
{
vtkm::cont::CellSetStructured<3> cellSet(cellNm);
cellSet.SetPointDimensions(vtkm::Id3(dims[0], dims[1], dims[2]));
dataSet.AddCellSet(cellSet);
}
else
{
throw vtkm::cont::ErrorBadValue("Invalid cell set dimension");
}
return dataSet;
}
}
} // end namespace vtkm::cont

71
vtkm/cont/DataSetBuilderUniform.h
View File

@ -28,13 +28,13 @@ namespace vtkm
namespace cont
{
class DataSetBuilderUniform
class VTKM_CONT_EXPORT DataSetBuilderUniform
{
using VecType = vtkm::Vec<vtkm::FloatDefault, 3>;
public:
VTKM_CONT
DataSetBuilderUniform() {}
DataSetBuilderUniform();
//1D uniform grid
template <typename T>
@ -55,10 +55,7 @@ public:
VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id& dimension,
std::string coordNm = "coords",
std::string cellNm = "cells")
{
return CreateDataSet(vtkm::Id3(dimension, 1, 1), VecType(0), VecType(1), coordNm, cellNm);
}
std::string cellNm = "cells");
//2D uniform grids.
template <typename T>
@ -82,11 +79,7 @@ public:
VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id2& dimensions,
std::string coordNm = "coords",
std::string cellNm = "cells")
{
return CreateDataSet(
vtkm::Id3(dimensions[0], dimensions[1], 1), VecType(0), VecType(1), coordNm, cellNm);
}
std::string cellNm = "cells");
//3D uniform grids.
template <typename T>
@ -111,14 +104,7 @@ public:
VTKM_CONT
static vtkm::cont::DataSet Create(const vtkm::Id3& dimensions,
std::string coordNm = "coords",
std::string cellNm = "cells")
{
return CreateDataSet(vtkm::Id3(dimensions[0], dimensions[1], dimensions[2]),
VecType(0),
VecType(1),
coordNm,
cellNm);
}
std::string cellNm = "cells");
private:
VTKM_CONT
@ -126,52 +112,7 @@ private:
const vtkm::Vec<vtkm::FloatDefault, 3>& origin,
const vtkm::Vec<vtkm::FloatDefault, 3>& spacing,
std::string coordNm,
std::string cellNm)
{
vtkm::Id dims[3] = { 1, 1, 1 };
int ndims = 0;
for (int i = 0; i < 3; ++i)
{
if (dimensions[i] > 1)
{
if (spacing[i] <= 0.0f)
{
throw vtkm::cont::ErrorBadValue("spacing must be > 0.0");
}
dims[ndims++] = dimensions[i];
}
}
vtkm::cont::DataSet dataSet;
vtkm::cont::ArrayHandleUniformPointCoordinates coords(dimensions, origin, spacing);
vtkm::cont::CoordinateSystem cs(coordNm, coords);
dataSet.AddCoordinateSystem(cs);
if (ndims == 1)
{
vtkm::cont::CellSetStructured<1> cellSet(cellNm);
cellSet.SetPointDimensions(dims[0]);
dataSet.AddCellSet(cellSet);
}
else if (ndims == 2)
{
vtkm::cont::CellSetStructured<2> cellSet(cellNm);
cellSet.SetPointDimensions(vtkm::Id2(dims[0], dims[1]));
dataSet.AddCellSet(cellSet);
}
else if (ndims == 3)
{
vtkm::cont::CellSetStructured<3> cellSet(cellNm);
cellSet.SetPointDimensions(vtkm::Id3(dims[0], dims[1], dims[2]));
dataSet.AddCellSet(cellSet);
}
else
{
throw vtkm::cont::ErrorBadValue("Invalid cell set dimension");
}
return dataSet;
}
std::string cellNm);
};
} // namespace cont

2
vtkm/cont/DeviceAdapterListTag.h
View File

@ -27,6 +27,7 @@
#include <vtkm/ListTag.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/openmp/DeviceAdapterOpenMP.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
@ -37,6 +38,7 @@ namespace cont
struct DeviceAdapterListTagCommon : vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda,
vtkm::cont::DeviceAdapterTagTBB,
vtkm::cont::DeviceAdapterTagOpenMP,
vtkm::cont::DeviceAdapterTagSerial>
{
};

8
vtkm/cont/Field.h
View File

@ -276,6 +276,14 @@ public:
VTKM_CONT
virtual void PrintSummary(std::ostream& out) const;
VTKM_CONT
virtual void ReleaseResourcesExecution()
{
// TODO: Call ReleaseResourcesExecution on the data when
// the DynamicArrayHandle class is able to do so.
this->Range.ReleaseResourcesExecution();
}
private:
std::string Name; ///< name of field

49
vtkm/cont/PointLocator.h
View File

@ -7,8 +7,8 @@
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
// Copyright 2014 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.
@ -20,9 +20,7 @@
#ifndef vtk_m_cont_PointLocator_h
#define vtk_m_cont_PointLocator_h
#include <vtkm/Types.h>
#include <vtkm/cont/CoordinateSystem.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ExecutionObjectBase.h>
#include <vtkm/exec/PointLocator.h>
@ -33,49 +31,46 @@ namespace cont
class PointLocator : public vtkm::cont::ExecutionObjectBase
{
public:
PointLocator()
: Dirty(true)
: dirty(true)
{
}
vtkm::cont::CoordinateSystem GetCoordinates() const { return Coords; }
vtkm::cont::CoordinateSystem GetCoords() const { return coordinates; }
void SetCoordinates(const vtkm::cont::CoordinateSystem& coords)
void SetCoords(const vtkm::cont::CoordinateSystem& coords)
{
Coords = coords;
Dirty = true;
coordinates = coords;
dirty = true;
}
virtual void Build() = 0;
void Update()
{
if (Dirty)
if (dirty)
Build();
Dirty = false;
dirty = false;
}
template <typename DeviceAdapter>
VTKM_CONT const vtkm::exec::PointLocator* PrepareForExecution(DeviceAdapter)
VTKM_CONT const vtkm::exec::PointLocator* PrepareForExecution(DeviceAdapter) const
{
vtkm::cont::DeviceAdapterId deviceId = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetId();
return PrepareForExecution(deviceId);
return PrepareForExecutionImp(deviceId).PrepareForExecution(DeviceAdapter());
}
protected:
void SetDirty() { Dirty = true; }
VTKM_CONT virtual void Build() = 0;
VTKM_CONT virtual const vtkm::exec::PointLocator* PrepareForExecutionImpl(
const vtkm::Int8 device) = 0;
//VTKM_CONT virtual const vtkm::exec::PointLocator*
using HandleType = vtkm::cont::VirtualObjectHandle<vtkm::exec::PointLocator>;
VTKM_CONT virtual const HandleType PrepareForExecutionImp(
vtkm::cont::DeviceAdapterId deviceId) const = 0;
private:
vtkm::cont::CoordinateSystem Coords;
bool Dirty;
vtkm::cont::CoordinateSystem coordinates;
bool dirty;
};
} // namespace cont
} // namespace vtkm
}
}
#endif // vtk_m_cont_PointLocator_h

268
vtkm/cont/PointLocatorUniformGrid.h
View File

@ -23,21 +23,25 @@
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/PointLocator.h>
#include <vtkm/exec/PointLocatorUniformGrid.h>
namespace vtkm
{
namespace worklet
namespace cont
{
template <typename T>
class PointLocatorUniformGrid
class PointLocatorUniformGrid : public vtkm::cont::PointLocator
{
public:
PointLocatorUniformGrid(const vtkm::Vec<T, 3>& _min,
const vtkm::Vec<T, 3>& _max,
PointLocatorUniformGrid(const vtkm::Vec<vtkm::FloatDefault, 3>& _min,
const vtkm::Vec<vtkm::FloatDefault, 3>& _max,
const vtkm::Vec<vtkm::Id, 3>& _dims)
: Min(_min)
: PointLocator()
, Min(_min)
, Max(_max)
, Dims(_dims)
{
@ -51,7 +55,9 @@ public:
using ExecutionSignature = void(_1, _2);
VTKM_CONT
BinPointsWorklet(vtkm::Vec<T, 3> _min, vtkm::Vec<T, 3> _max, vtkm::Vec<vtkm::Id, 3> _dims)
BinPointsWorklet(vtkm::Vec<vtkm::FloatDefault, 3> _min,
vtkm::Vec<vtkm::FloatDefault, 3> _max,
vtkm::Vec<vtkm::Id, 3> _dims)
: Min(_min)
, Dims(_dims)
, Dxdydz((_max - Min) / Dims)
@ -66,172 +72,126 @@ public:
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<vtkm::FloatDefault, 3> Min;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<T, 3> Dxdydz;
};
class UniformGridSearch : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(FieldIn<> query,
WholeArrayIn<> coordIn,
WholeArrayIn<IdType> pointId,
WholeArrayIn<IdType> cellLower,
WholeArrayIn<IdType> cellUpper,
FieldOut<IdType> neighborId,
FieldOut<> distance);
using ExecutionSignature = void(_1, _2, _3, _4, _5, _6, _7);
VTKM_CONT
UniformGridSearch(const vtkm::Vec<T, 3>& _min,
const vtkm::Vec<T, 3>& _max,
const vtkm::Vec<vtkm::Id, 3>& _dims)
: Min(_min)
, Dims(_dims)
, Dxdydz((_max - _min) / _dims)
{
}
template <typename CoordiVecType,
typename IdPortalType,
typename CoordiPortalType,
typename IdType,
typename CoordiType>
VTKM_EXEC void operator()(const CoordiVecType& queryCoord,
const CoordiPortalType& coordi_Handle,
const IdPortalType& pointId,
const IdPortalType& cellLower,
const IdPortalType& cellUpper,
IdType& nnId,
CoordiType& nnDis) const
{
auto nlayers = vtkm::Max(vtkm::Max(Dims[0], Dims[1]), Dims[2]);
vtkm::Vec<vtkm::Id, 3> xyz = (queryCoord - Min) / Dxdydz;
float min_distance = std::numeric_limits<float>::max();
vtkm::Id neareast = -1;
for (vtkm::Id layer = 0; layer < nlayers; layer++)
{
vtkm::Id minx = vtkm::Max(vtkm::Id(), xyz[0] - layer);
vtkm::Id maxx = vtkm::Min(Dims[0] - 1, xyz[0] + layer);
vtkm::Id miny = vtkm::Max(vtkm::Id(), xyz[1] - layer);
vtkm::Id maxy = vtkm::Min(Dims[1] - 1, xyz[1] + layer);
vtkm::Id minz = vtkm::Max(vtkm::Id(), xyz[2] - layer);
vtkm::Id maxz = vtkm::Min(Dims[2] - 1, xyz[2] + layer);
for (auto i = minx; i <= maxx; i++)
{
for (auto j = miny; j <= maxy; j++)
{
for (auto k = minz; k <= maxz; k++)
{
if (i == (xyz[0] + layer) || i == (xyz[0] - layer) || j == (xyz[1] + layer) ||
j == (xyz[1] - layer) || k == (xyz[2] + layer) || k == (xyz[2] - layer))
{
auto cellid = i + j * Dims[0] + k * Dims[0] * Dims[1];
auto lower = cellLower.Get(cellid);
auto upper = cellUpper.Get(cellid);
for (auto index = lower; index < upper; index++)
{
auto pointid = pointId.Get(index);
auto point = coordi_Handle.Get(pointid);
auto dx = point[0] - queryCoord[0];
auto dy = point[1] - queryCoord[1];
auto dz = point[2] - queryCoord[2];
auto distance2 = dx * dx + dy * dy + dz * dz;
if (distance2 < min_distance)
{
neareast = pointid;
min_distance = distance2;
nlayers = layer + 2;
}
}
}
}
}
}
}
nnId = neareast;
nnDis = vtkm::Sqrt(min_distance);
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<T, 3> Dxdydz;
vtkm::Vec<vtkm::FloatDefault, 3> Dxdydz;
};
/// \brief Construct a 3D uniform grid for nearest neighbor search.
///
/// \param coords An ArrayHandle of x, y, z coordinates of input points.
/// \param device Tag for selecting device adapter
template <typename DeviceAdapter>
void Build(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
DeviceAdapter vtkmNotUsed(device))
struct BuildFunctor
{
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(0, 1, coords.GetNumberOfValues());
Algorithm::Copy(pointCounting, PointIds);
BuildFunctor(vtkm::cont::PointLocatorUniformGrid* self)
: Self(self)
{
}
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(Min, Max, Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet, DeviceAdapter> dispatchCellId(
cellIdWorklet);
dispatchCellId.Invoke(coords, CellIds);
template <typename Device>
bool operator()(Device)
{
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<Device>;
// Group points of the same cell together by sorting them according to the cell ids
Algorithm::SortByKey(CellIds, PointIds);
// Save training data points.
Algorithm::Copy(this->Self->GetCoords().GetData(), this->Self->coords);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(0, 1, Dims[0] * Dims[1] * Dims[2]);
Algorithm::UpperBounds(CellIds, cell_ids_counting, CellUpper);
Algorithm::LowerBounds(CellIds, cell_ids_counting, CellLower);
}
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(
0, 1, this->Self->coords.GetNumberOfValues());
Algorithm::Copy(pointCounting, this->Self->pointIds);
/// \brief Nearest neighbor search using a Uniform Grid
///
/// Parallel search of nearesat neighbor for each point in the \c queryPoints in the set of
/// \c coords. Returns neareast neighbot in \c nearestNeighborIds and distances to nearest
/// neighbor in \c distances.
///
/// \param coords Point coordinates for training dataset.
/// \param queryPoints Point coordinates to query for nearest neighbors.
/// \param nearestNeighborIds Neareast neighbor in the training dataset for each points in
/// the test set
/// \param distances Distance between query points and their nearest neighbors.
/// \param device Tag for selecting device adapter.
template <typename DeviceAdapter>
void FindNearestPoint(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& coords,
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>>& queryPoints,
vtkm::cont::ArrayHandle<vtkm::Id>& nearestNeighborIds,
vtkm::cont::ArrayHandle<T>& distances,
DeviceAdapter)
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(this->Self->Min, this->Self->Max, this->Self->Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet, Device> dispatchCellId(cellIdWorklet);
dispatchCellId.Invoke(this->Self->coords, this->Self->cellIds);
// Group points of the same cell together by sorting them according to the cell ids
Algorithm::SortByKey(this->Self->cellIds, this->Self->pointIds);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(
0, 1, this->Self->Dims[0] * this->Self->Dims[1] * this->Self->Dims[2]);
Algorithm::UpperBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellUpper);
Algorithm::LowerBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellLower);
return true;
}
private:
vtkm::cont::PointLocatorUniformGrid* Self;
};
void Build() override
{
UniformGridSearch uniformGridSearch(Min, Max, Dims);
BuildFunctor functor(this);
vtkm::worklet::DispatcherMapField<UniformGridSearch, DeviceAdapter> searchDispatcher(
uniformGridSearch);
searchDispatcher.Invoke(
queryPoints, coords, PointIds, CellLower, CellUpper, nearestNeighborIds, distances);
bool success = vtkm::cont::TryExecute(functor);
if (!success)
{
throw vtkm::cont::ErrorExecution("Could not build point locator structure");
}
};
using HandleType = vtkm::cont::VirtualObjectHandle<vtkm::exec::PointLocator>;
struct PrepareForExecutionFunctor
{
template <typename DeviceAdapter>
VTKM_CONT void operator()(DeviceAdapter,
const vtkm::cont::PointLocatorUniformGrid& self,
HandleType& handle) const
{
//vtkm::exec::PointLocatorUniformGrid* locator =
vtkm::exec::PointLocatorUniformGrid<DeviceAdapter>* h =
new vtkm::exec::PointLocatorUniformGrid<DeviceAdapter>(
self.Min,
self.Max,
self.Dims,
self.coords.PrepareForInput(DeviceAdapter()),
self.pointIds.PrepareForInput(DeviceAdapter()),
self.cellLower.PrepareForInput(DeviceAdapter()),
self.cellUpper.PrepareForInput(DeviceAdapter()));
handle.Reset(h);
//return handle.PrepareForExecution(DeviceAdapter());
}
};
VTKM_CONT
//const vtkm::exec::PointLocator *
const HandleType PrepareForExecutionImp(vtkm::cont::DeviceAdapterId deviceId) const override
{
// TODO: call VirtualObjectHandle::PrepareForExecution() and return vtkm::exec::PointLocator
// TODO: how to convert deviceId back to DeviceAdapter tag?
//using DeviceList = vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagCuda,
// vtkm::cont::DeviceAdapterTagTBB,
// vtkm::cont::DeviceAdapterTagSerial>;
using DeviceList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG;
//HandleType ExecHandle; // = new HandleType(locator, false);
vtkm::cont::internal::FindDeviceAdapterTagAndCall(
deviceId, DeviceList(), PrepareForExecutionFunctor(), *this, ExecHandle);
return ExecHandle;
//return ExecHandle.PrepareForExecution(DeviceAdapter());
}
private:
vtkm::Vec<T, 3> Min;
vtkm::Vec<T, 3> Max;
vtkm::Vec<vtkm::FloatDefault, 3> Min;
vtkm::Vec<vtkm::FloatDefault, 3> Max;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::cont::ArrayHandle<vtkm::Id> PointIds;
vtkm::cont::ArrayHandle<vtkm::Id> CellIds;
vtkm::cont::ArrayHandle<vtkm::Id> CellLower;
vtkm::cont::ArrayHandle<vtkm::Id> CellUpper;
// TODO: how to convert CoordinateSystem to ArrayHandle<Vec<Float, 3>>?
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> coords;
vtkm::cont::ArrayHandle<vtkm::Id> pointIds;
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<vtkm::Id> cellLower;
vtkm::cont::ArrayHandle<vtkm::Id> cellUpper;
// TODO: std::unique_ptr/std::shared_ptr?
mutable HandleType ExecHandle;
};
}
}

14
vtkm/cont/StorageImplicit.h
View File

@ -70,6 +70,7 @@ public:
VTKM_CONT
Storage(const PortalConstType& portal = PortalConstType())
: Portal(portal)
, NumberOfValues(portal.GetNumberOfValues())
{
}
@ -79,22 +80,25 @@ public:
VTKM_CONT
PortalConstType GetPortalConst() const { return this->Portal; }
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->Portal.GetNumberOfValues(); }
vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
VTKM_CONT
void Allocate(vtkm::Id vtkmNotUsed(numberOfValues))
void Allocate(vtkm::Id numberOfValues)
{
throw vtkm::cont::ErrorBadValue("Implicit arrays are read-only.");
VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues());
this->NumberOfValues = numberOfValues;
}
VTKM_CONT
void Shrink(vtkm::Id vtkmNotUsed(numberOfValues))
void Shrink(vtkm::Id numberOfValues)
{
throw vtkm::cont::ErrorBadValue("Implicit arrays are read-only.");
VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues());
this->NumberOfValues = numberOfValues;
}
VTKM_CONT
void ReleaseResources() {}
private:
PortalConstType Portal;
vtkm::Id NumberOfValues;
};
template <typename T, class ArrayPortalType, class DeviceAdapterTag>

2
vtkm/cont/arg/TransportTagArrayIn.h
View File

@ -47,7 +47,7 @@ struct Transport<vtkm::cont::arg::TransportTagArrayIn, ContObjectType, Device>
{
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = typename ContObjectType::template ExecutionTypes<Device>::PortalConst;
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInput(Device()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object,

2
vtkm/cont/arg/TransportTagArrayInOut.h
View File

@ -50,7 +50,7 @@ struct Transport<vtkm::cont::arg::TransportTagArrayInOut, ContObjectType, Device
// is not an array handle as an argument that is expected to be one.
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = typename ContObjectType::template ExecutionTypes<Device>::Portal;
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInPlace(Device()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType object,

3
vtkm/cont/arg/TransportTagArrayOut.h
View File

@ -49,7 +49,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayOut, ContObjectType, Device>
// is not an array handle as an argument that is expected to be one.
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = typename ContObjectType::template ExecutionTypes<Device>::Portal;
using ExecObjectType =
decltype(std::declval<ContObjectType>().PrepareForOutput(vtkm::Id{}, Device()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType object,

6
vtkm/cont/arg/TransportTagCellSetIn.h
View File

@ -50,10 +50,8 @@ struct Transport<vtkm::cont::arg::TransportTagCellSetIn<FromTopology, ToTopology
{
VTKM_IS_CELL_SET(ContObjectType);
using ExecObjectType =
typename ContObjectType::template ExecutionTypes<Device,
FromTopology,
ToTopology>::ExecObjectType;
using ExecObjectType = decltype(
std::declval<ContObjectType>().PrepareForInput(Device(), FromTopology(), ToTopology()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType

3
vtkm/cont/arg/TransportTagTopologyFieldIn.h
View File

@ -87,7 +87,8 @@ struct Transport<vtkm::cont::arg::TransportTagTopologyFieldIn<TopologyElementTag
{
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = typename ContObjectType::template ExecutionTypes<Device>::PortalConst;
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInput(Device()));
VTKM_CONT
ExecObjectType operator()(const ContObjectType& object,

1
vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
View File

@ -432,6 +432,7 @@ private:
{
cuda::internal::throwAsVTKmException();
}
VTKM_CUDA_CALL(cudaStreamSynchronize(cudaStreamPerThread));
return sum[2];
}

4
vtkm/cont/cuda/internal/DeviceAdapterAtomicArrayImplementationCuda.h
View File

@ -72,16 +72,19 @@ private:
vtkm::cont::DeviceAdapterTagCuda>::Portal;
PortalType Portal;
VTKM_SUPPRESS_EXEC_WARNINGS
__device__ vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
{
return atomicAdd((unsigned long long*)address, (unsigned long long)value);
}
VTKM_SUPPRESS_EXEC_WARNINGS
__device__ vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
{
return atomicAdd(address, value);
}
VTKM_SUPPRESS_EXEC_WARNINGS
__device__ vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
const vtkm::Int32& newValue,
const vtkm::Int32& oldValue) const
@ -89,6 +92,7 @@ private:
return atomicCAS(address, oldValue, newValue);
}
VTKM_SUPPRESS_EXEC_WARNINGS
__device__ vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
const vtkm::Int64& newValue,
const vtkm::Int64& oldValue) const

19
vtkm/cont/cuda/internal/DeviceAdapterRuntimeDetectorCuda.cxx
View File

@ -65,7 +65,11 @@ DeviceAdapterRuntimeDetector<vtkm::cont::DeviceAdapterTagCuda>::DeviceAdapterRun
deviceQueryInit = true;
//first query for the number of devices
VTKM_CUDA_CALL(cudaGetDeviceCount(&numDevices));
auto res = cudaGetDeviceCount(&numDevices);
if (res != cudaSuccess)
{
numDevices = 0;
}
for (vtkm::Int32 i = 0; i < numDevices; i++)
{
@ -84,12 +88,15 @@ DeviceAdapterRuntimeDetector<vtkm::cont::DeviceAdapterTagCuda>::DeviceAdapterRun
// 3. cudaErrorNoKernelImageForDevice we built for a compute version
// greater than the device we are running on
// Most likely others that I don't even know about
vtkm::cont::cuda::internal::DetermineIfValidCudaDevice<<<1, 1, 0, cudaStreamPerThread>>>();
cudaStreamSynchronize(cudaStreamPerThread);
if (cudaSuccess != cudaGetLastError())
if (numDevices > 0)
{
numDevices = 0;
archVersion = 0;
vtkm::cont::cuda::internal::DetermineIfValidCudaDevice<<<1, 1, 0, cudaStreamPerThread>>>();
cudaStreamSynchronize(cudaStreamPerThread);
if (cudaSuccess != cudaGetLastError())
{
numDevices = 0;
archVersion = 0;
}
}
}

28
vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.cu
View File

@ -150,15 +150,6 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
static_cast<std::size_t>(numBytes),
cudaMemcpyHostToDevice,
cudaStreamPerThread));
if (CudaAllocator::IsManagedPointer(executionPtr))
{
//If we are moving memory from unmanaged host memory
//to managed host memory we have the possibility that
//the memcpy will not finish before the first usage is finished
//to work around this bug we explicitly synchronize for this
//one use case
cudaStreamSynchronize(cudaStreamPerThread);
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const void* executionPtr,
@ -179,14 +170,21 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const voi
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
else
{
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
}
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
//In all cases we have possibly multiple async calls queued up in
//our stream. We need to block on the copy back to control since
//we don't wanting it accessing memory that hasn't finished
//being used by the GPU
cudaStreamSynchronize(cudaStreamPerThread);
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(

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

@ -29,7 +29,7 @@ namespace cont
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda> final
: public ExecutionArrayInterfaceBasicBase
{
using Superclass = ExecutionArrayInterfaceBasicBase;

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

@ -28,6 +28,7 @@ set(unit_tests
UnitTestCudaDataSetExplicit.cu
UnitTestCudaDataSetSingleType.cu
UnitTestCudaDeviceAdapter.cu
UnitTestCudaGeometry.cu
UnitTestCudaImplicitFunction.cu
UnitTestCudaMath.cu
UnitTestCudaShareUserProvidedManagedMemory.cu

38
vtkm/cont/cuda/testing/UnitTestCudaGeometry.cu Normal file
View File

@ -0,0 +1,38 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 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.
//
//=============================================================================
// Make sure that the tested code is using the device adapter specified. This
// is important in the long run so we don't, for example, use the CUDA device
// for a part of an operation where the TBB device was specified.
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_ERROR
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/testing/TestingGeometry.h>
int UnitTestCudaGeometry(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
return vtkm::cont::testing::Testing::Run(
UnitTestGeometryNamespace::RunGeometryTests<vtkm::cont::DeviceAdapterTagCuda>);
}

4
vtkm/cont/cuda/testing/UnitTestCudaPointLocatorUniformGrid.cu
View File

@ -27,8 +27,6 @@
int UnitTestCudaPointLocatorUniformGrid(int, char* [])
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(vtkm::cont::DeviceAdapterTagCuda{});
return vtkm::cont::testing::Testing::Run(
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagSerial>());
TestingPointLocatorUniformGrid<vtkm::cont::DeviceAdapterTagCuda>());
}

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

@ -32,7 +32,7 @@ namespace internal
/// \brief Class that manages data in the execution environment.
///
/// This templated class must be partially specialized for each
/// DeviceAdapterTag crated, which will define the implementation for that tag.
/// DeviceAdapterTag created, which will define the implementation for that tag.
///
/// This is a class that is responsible for allocating data in the execution
/// environment and copying data back and forth between control and

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

@ -38,6 +38,9 @@ set(headers
DynamicTransform.h
FunctorsGeneral.h
IteratorFromArrayPortal.h
KXSort.h
ParallelRadixSort.h
ParallelRadixSortInterface.h
SimplePolymorphicContainer.h
StorageError.h
VirtualObjectTransfer.h

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

@ -73,8 +73,7 @@ void buildIndexOffsets(const ArrayHandleIndices& numIndices,
ArrayHandleOffsets offsets,
DeviceAdapterTag tag)
{
using IsWriteable =
vtkm::cont::internal::IsWriteableArrayHandle<ArrayHandleOffsets, DeviceAdapterTag>;
using IsWriteable = vtkm::cont::internal::IsWriteableArrayHandle<ArrayHandleOffsets>;
buildIndexOffsets(numIndices, offsets, tag, typename IsWriteable::type());
}

12
vtkm/cont/internal/DeviceAdapterDefaultSelection.h
View File

@ -36,7 +36,9 @@
// Unfortunately, VTKM_ENABLE_TBB does not guarantee that TBB is (or isn't)
// available, but there is no way to check for sure in a header library.
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_TBB
#else // !VTKM_CUDA && !VTKM_ENABLE_TBB
#elif defined(VTKM_ENABLE_OPENMP)
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_OPENMP
#else // !VTKM_CUDA && !VTKM_ENABLE_TBB && !VTKM_ENABLE_OPENMP
#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_SERIAL
#endif // !VTKM_CUDA && !VTKM_ENABLE_TBB
#endif // VTKM_DEVICE_ADAPTER
@ -68,6 +70,14 @@
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#define VTKM_DEFAULT_DEVICE_ADAPTER_TAG ::vtkm::cont::DeviceAdapterTagTBB
// OpenMP:
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
#include <vtkm/cont/openmp/internal/ArrayManagerExecutionOpenMP.h>
#include <vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h>
#include <vtkm/cont/openmp/internal/DeviceAdapterTagOpenMP.h>
#define VTKM_DEFAULT_DEVICE_ADAPTER_TAG ::vtkm::cont::DeviceAdapterTagOpenMP
// Error:
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_ERROR

1
vtkm/cont/internal/DeviceAdapterTag.h
View File

@ -32,6 +32,7 @@
#define VTKM_DEVICE_ADAPTER_SERIAL 1
#define VTKM_DEVICE_ADAPTER_CUDA 2
#define VTKM_DEVICE_ADAPTER_TBB 3
#define VTKM_DEVICE_ADAPTER_OPENMP 4
namespace vtkm
{

22
vtkm/cont/tbb/internal/kxsort.h → vtkm/cont/internal/KXSort.h
View File

@ -1,3 +1,25 @@
//=============================================================================
//
// 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.
//
//=============================================================================
/* The MIT License
Copyright (c) 2016 Dinghua Li <voutcn@gmail.com>

1069
vtkm/cont/internal/ParallelRadixSort.h Normal file
View File

File diff suppressed because it is too large Load Diff

166
vtkm/cont/internal/ParallelRadixSortInterface.h Normal file
View File

@ -0,0 +1,166 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_internal_ParallelRadixSortInterface_h
#define vtk_m_cont_internal_ParallelRadixSortInterface_h
#include <vtkm/BinaryPredicates.h>
#include <vtkm/cont/ArrayHandle.h>
#include <functional>
#include <type_traits>
namespace vtkm
{
namespace cont
{
namespace internal
{
namespace radix
{
const size_t MIN_BYTES_FOR_PARALLEL = 400000;
const size_t BYTES_FOR_MAX_PARALLELISM = 4000000;
struct RadixSortTag
{
};
struct PSortTag
{
};
// Detect supported functors for radix sort:
template <typename T>
struct is_valid_compare_type : std::integral_constant<bool, false>
{
};
template <typename T>
struct is_valid_compare_type<std::less<T>> : std::integral_constant<bool, true>
{
};
template <typename T>
struct is_valid_compare_type<std::greater<T>> : std::integral_constant<bool, true>
{
};
template <>
struct is_valid_compare_type<vtkm::SortLess> : std::integral_constant<bool, true>
{
};
template <>
struct is_valid_compare_type<vtkm::SortGreater> : std::integral_constant<bool, true>
{
};
// Convert vtkm::Sort[Less|Greater] to the std:: equivalents:
template <typename BComp, typename T>
BComp&& get_std_compare(BComp&& b, T&&)
{
return std::forward<BComp>(b);
}
template <typename T>
std::less<T> get_std_compare(vtkm::SortLess, T&&)
{
return std::less<T>{};
}
template <typename T>
std::greater<T> get_std_compare(vtkm::SortGreater, T&&)
{
return std::greater<T>{};
}
// Determine if radix sort can be used for a given ValueType, StorageType, and
// comparison functor.
template <typename T, typename StorageTag, typename BinaryCompare>
struct sort_tag_type
{
using type = PSortTag;
};
template <typename T, typename BinaryCompare>
struct sort_tag_type<T, vtkm::cont::StorageTagBasic, BinaryCompare>
{
using PrimT = std::is_arithmetic<T>;
using LongDT = std::is_same<T, long double>;
using BComp = is_valid_compare_type<BinaryCompare>;
using type = typename std::conditional<PrimT::value && BComp::value && !LongDT::value,
RadixSortTag,
PSortTag>::type;
};
template <typename KeyType,
typename ValueType,
typename KeyStorageTagType,
typename ValueStorageTagType,
class BinaryCompare>
struct sortbykey_tag_type
{
using type = PSortTag;
};
template <typename KeyType, typename ValueType, class BinaryCompare>
struct sortbykey_tag_type<KeyType,
ValueType,
vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagBasic,
BinaryCompare>
{
using PrimKey = std::is_arithmetic<KeyType>;
using PrimValue = std::is_arithmetic<ValueType>;
using LongDKey = std::is_same<KeyType, long double>;
using BComp = is_valid_compare_type<BinaryCompare>;
using type = typename std::conditional<PrimKey::value && PrimValue::value && BComp::value &&
!LongDKey::value,
RadixSortTag,
PSortTag>::type;
};
#define VTKM_INTERNAL_RADIX_SORT_DECLARE(key_type) \
VTKM_CONT_EXPORT void parallel_radix_sort( \
key_type* data, size_t num_elems, const std::greater<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort( \
key_type* data, size_t num_elems, const std::less<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::greater<key_type>& comp); \
VTKM_CONT_EXPORT void parallel_radix_sort_key_values( \
key_type* keys, vtkm::Id* vals, size_t num_elems, const std::less<key_type>& comp);
// Generate radix sort interfaces for key and key value sorts.
#define VTKM_DECLARE_RADIX_SORT() \
VTKM_INTERNAL_RADIX_SORT_DECLARE(short int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(unsigned short int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(unsigned int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(long int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(unsigned long int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(long long int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(unsigned long long int) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(unsigned char) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(signed char) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(char) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(char16_t) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(char32_t) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(wchar_t) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(float) \
VTKM_INTERNAL_RADIX_SORT_DECLARE(double)
}
}
}
} // end vtkm::cont::internal::radix
#endif // vtk_m_cont_internal_ParallelRadixSortInterface_h

32
vtkm/cont/openmp/CMakeLists.txt Normal file
View File

@ -0,0 +1,32 @@
##============================================================================
## 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.
##============================================================================
set(headers
DeviceAdapterOpenMP.h
)
add_subdirectory(internal)
vtkm_declare_headers(${headers} TESTABLE ${VTKm_ENABLE_OPENMP})
#-----------------------------------------------------------------------------
if (TARGET vtkm::openmp)
add_subdirectory(testing)
endif()

33
vtkm/cont/openmp/DeviceAdapterOpenMP.h Normal file
View File

@ -0,0 +1,33 @@
//============================================================================
// 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.
//============================================================================
#ifndef vtk_m_cont_openmp_DeviceAdapterOpenMP_h
#define vtk_m_cont_openmp_DeviceAdapterOpenMP_h
#include <vtkm/cont/openmp/internal/DeviceAdapterRuntimeDetectorOpenMP.h>
#include <vtkm/cont/openmp/internal/DeviceAdapterTagOpenMP.h>
#ifdef VTKM_ENABLE_OPENMP
#include <vtkm/cont/openmp/internal/ArrayManagerExecutionOpenMP.h>
#include <vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h>
#include <vtkm/cont/openmp/internal/VirtualObjectTransferOpenMP.h>
#endif
#endif //vtk_m_cont_openmp_DeviceAdapterOpenMP_h

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