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

Merge branch 'master' into datamodel-design

Browse Source
This commit is contained in:
Robert Maynard 2015-06-03 14:36:53 -04:00
commit 7f2ee8d050
218 changed files with 3243 additions and 476 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
CMake/FindPyexpander.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

345
CMake/FindTBB.cmake Normal file
View File

@ -0,0 +1,345 @@
# - Find ThreadingBuildingBlocks include dirs and libraries
# Use this module by invoking find_package with the form:
# find_package(TBB
# [REQUIRED] # Fail with error if TBB is not found
# ) #
# Once done, this will define
#
# TBB_FOUND - system has TBB
# TBB_INCLUDE_DIRS - the TBB include directories
# TBB_LIBRARIES - TBB libraries to be lined, doesn't include malloc or
# malloc proxy
#
# TBB_VERSION_MAJOR - Major Product Version Number
# TBB_VERSION_MINOR - Minor Product Version Number
# TBB_INTERFACE_VERSION - Engineering Focused Version Number
# TBB_COMPATIBLE_INTERFACE_VERSION - The oldest major interface version
# still supported. This uses the engineering
# focused interface version numbers.
#
# TBB_MALLOC_FOUND - system has TBB malloc library
# TBB_MALLOC_INCLUDE_DIRS - the TBB malloc include directories
# TBB_MALLOC_LIBRARIES - The TBB malloc libraries to be lined
#
# TBB_MALLOC_PROXY_FOUND - system has TBB malloc proxy library
# TBB_MALLOC_PROXY_INCLUDE_DIRS = the TBB malloc proxy include directories
# TBB_MALLOC_PROXY_LIBRARIES - The TBB malloc proxy libraries to be lined
#
#
# This module reads hints about search locations from variables:
# ENV TBB_ARCH_PLATFORM - for eg. set it to "mic" for Xeon Phi builds
# ENV TBB_ROOT or just TBB_ROOT - root directory of tbb installation
# ENV TBB_BUILD_PREFIX - specifies the build prefix for user built tbb
# libraries. Should be specified with ENV TBB_ROOT
# and optionally...
# ENV TBB_BUILD_DIR - if build directory is different than ${TBB_ROOT}/build
#
#
# Modified by Robert Maynard from the original OGRE source
#
#-------------------------------------------------------------------
# This file is part of the CMake build system for OGRE
# (Object-oriented Graphics Rendering Engine)
# For the latest info, see http://www.ogre3d.org/
#
# The contents of this file are placed in the public domain. Feel
# free to make use of it in any way you like.
#-------------------------------------------------------------------
#
#=============================================================================
# Copyright 2010-2012 Kitware, Inc.
# Copyright 2012 Rolf Eike Beer <eike@sf-mail.de>
#
# Distributed under the OSI-approved BSD License (the "License");
# see accompanying file Copyright.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 License for more information.
#=============================================================================
# (To distribute this file outside of CMake, substitute the full
# License text for the above reference.)
#=============================================================================
# FindTBB helper functions and macros
#
#===============================================
# Do the final processing for the package find.
#===============================================
macro(findpkg_finish PREFIX)
# skip if already processed during this run
if (NOT ${PREFIX}_FOUND)
if (${PREFIX}_INCLUDE_DIR AND ${PREFIX}_LIBRARY)
set(${PREFIX}_FOUND TRUE)
set (${PREFIX}_INCLUDE_DIRS ${${PREFIX}_INCLUDE_DIR})
set (${PREFIX}_LIBRARIES ${${PREFIX}_LIBRARY})
else ()
if (${PREFIX}_FIND_REQUIRED AND NOT ${PREFIX}_FIND_QUIETLY)
message(FATAL_ERROR "Required library ${PREFIX} not found.")
endif ()
endif ()
#mark the following variables as internal variables
mark_as_advanced(${PREFIX}_INCLUDE_DIR
${PREFIX}_LIBRARY
${PREFIX}_LIBRARY_DEBUG
${PREFIX}_LIBRARY_RELEASE)
endif ()
endmacro(findpkg_finish)
#===============================================
# Generate debug names from given RELEASEease names
#===============================================
macro(get_debug_names PREFIX)
foreach(i ${${PREFIX}})
set(${PREFIX}_DEBUG ${${PREFIX}_DEBUG} ${i}d ${i}D ${i}_d ${i}_D ${i}_debug ${i})
endforeach(i)
endmacro(get_debug_names)
#===============================================
# See if we have env vars to help us find tbb
#===============================================
macro(getenv_path VAR)
set(ENV_${VAR} $ENV{${VAR}})
# replace won't work if var is blank
if (ENV_${VAR})
string( REGEX REPLACE "\\\\" "/" ENV_${VAR} ${ENV_${VAR}} )
endif ()
endmacro(getenv_path)
#===============================================
# Couple a set of RELEASEease AND debug libraries
#===============================================
macro(make_library_set PREFIX)
if (${PREFIX}_RELEASE AND ${PREFIX}_DEBUG)
set(${PREFIX} optimized ${${PREFIX}_RELEASE} debug ${${PREFIX}_DEBUG})
elseif (${PREFIX}_RELEASE)
set(${PREFIX} ${${PREFIX}_RELEASE})
elseif (${PREFIX}_DEBUG)
set(${PREFIX} ${${PREFIX}_DEBUG})
endif ()
endmacro(make_library_set)
#=============================================================================
# Now to actually find TBB
#
# Get path, convert backslashes as ${ENV_${var}}
getenv_path(TBB_ROOT)
# initialize search paths
set(TBB_PREFIX_PATH ${TBB_ROOT} ${ENV_TBB_ROOT})
set(TBB_INC_SEARCH_PATH "")
set(TBB_LIB_SEARCH_PATH "")
# If user built from sources
set(TBB_BUILD_PREFIX $ENV{TBB_BUILD_PREFIX})
if (TBB_BUILD_PREFIX AND ENV_TBB_ROOT)
getenv_path(TBB_BUILD_DIR)
if (NOT ENV_TBB_BUILD_DIR)
set(ENV_TBB_BUILD_DIR ${ENV_TBB_ROOT}/build)
endif ()
# include directory under ${ENV_TBB_ROOT}/include
list(APPEND TBB_LIB_SEARCH_PATH
${ENV_TBB_BUILD_DIR}/${TBB_BUILD_PREFIX}_release
${ENV_TBB_BUILD_DIR}/${TBB_BUILD_PREFIX}_debug)
endif ()
# For Windows, let's assume that the user might be using the precompiled
# TBB packages from the main website. These use a rather awkward directory
# structure (at least for automatically finding the right files) depending
# on platform and compiler, but we'll do our best to accommodate it.
# Not adding the same effort for the precompiled linux builds, though. Those
# have different versions for CC compiler versions and linux kernels which
# will never adequately match the user's setup, so there is no feasible way
# to detect the "best" version to use. The user will have to manually
# select the right files. (Chances are the distributions are shipping their
# custom version of tbb, anyway, so the problem is probably nonexistant.)
if (WIN32 AND MSVC)
set(COMPILER_PREFIX "vc7.1")
if (MSVC_VERSION EQUAL 1400)
set(COMPILER_PREFIX "vc8")
elseif(MSVC_VERSION EQUAL 1500)
set(COMPILER_PREFIX "vc9")
elseif(MSVC_VERSION EQUAL 1600)
set(COMPILER_PREFIX "vc10")
elseif(MSVC_VERSION EQUAL 1700)
set(COMPILER_PREFIX "vc11")
elseif(MSVC_VERSION EQUAL 1800)
set(COMPILER_PREFIX "vc12")
elseif(MSVC_VERSION EQUAL 1900)
set(COMPILER_PREFIX "vc14")
endif ()
# for each prefix path, add ia32/64\${COMPILER_PREFIX}\lib to the lib search path
foreach (dir ${TBB_PREFIX_PATH})
if (CMAKE_CL_64)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/ia64/${COMPILER_PREFIX}/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/ia64/${COMPILER_PREFIX})
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/intel64/${COMPILER_PREFIX}/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/intel64/${COMPILER_PREFIX})
else ()
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/ia32/${COMPILER_PREFIX}/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/ia32/${COMPILER_PREFIX})
endif ()
endforeach ()
endif ()
# For OS X binary distribution, choose libc++ based libraries for Maverics and
# above and AppleClang
if (${CMAKE_SYSTEM_NAME} STREQUAL "Darwin" AND
NOT ${CMAKE_SYSTEM_VERSION} LESS 13.0)
set (USE_LIBCXX OFF)
cmake_policy(GET CMP0025 POLICY_VAR)
if ("${POLICY_VAR}" STREQUAL "NEW")
if (${CMAKE_CXX_COMPILER_ID} STREQUAL "AppleClang")
set (USE_LIBCXX ON)
endif ()
else ()
if (${CMAKE_CXX_COMPILER_ID} STREQUAL "Clang")
set (USE_LIBCXX ON)
endif ()
endif ()
if (${USE_LIBCXX})
foreach (dir ${TBB_PREFIX_PATH})
list (APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/libc++ ${dir}/libc++/lib)
endforeach ()
endif ()
endif ()
# check compiler ABI
if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU" AND
CMAKE_CXX_COMPILER_VERSION VERSION_LESS 4.4)
set(COMPILER_PREFIX "gcc4.1")
else () # Assume compatibility with 4.4 for other compilers
set(COMPILER_PREFIX "gcc4.4")
endif ()
# if platform architecture is explicitly specified
set(TBB_ARCH_PLATFORM $ENV{TBB_ARCH_PLATFORM})
if (TBB_ARCH_PLATFORM)
foreach (dir ${TBB_PREFIX_PATH})
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/${TBB_ARCH_PLATFORM}/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/${TBB_ARCH_PLATFORM})
endforeach ()
endif ()
foreach (dir ${TBB_PREFIX_PATH})
if (CMAKE_SIZEOF_VOID_P EQUAL 8)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/intel64)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/intel64/${COMPILER_PREFIX})
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/intel64/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/intel64/${COMPILER_PREFIX}/lib)
else ()
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/ia32)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib/ia32/${COMPILER_PREFIX})
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/ia32/lib)
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/ia32/${COMPILER_PREFIX}/lib)
endif ()
endforeach ()
# add general search paths
foreach (dir ${TBB_PREFIX_PATH})
list(APPEND TBB_LIB_SEARCH_PATH ${dir}/lib ${dir}/Lib ${dir}/lib/tbb
${dir}/Libs)
list(APPEND TBB_INC_SEARCH_PATH ${dir}/include ${dir}/Include
${dir}/include/tbb)
endforeach ()
set(TBB_LIBRARY_NAMES tbb)
get_debug_names(TBB_LIBRARY_NAMES)
find_path(TBB_INCLUDE_DIR
NAMES tbb/tbb.h
PATHS ${TBB_INC_SEARCH_PATH})
find_library(TBB_LIBRARY_RELEASE
NAMES ${TBB_LIBRARY_NAMES}
PATHS ${TBB_LIB_SEARCH_PATH})
find_library(TBB_LIBRARY_DEBUG
NAMES ${TBB_LIBRARY_NAMES_DEBUG}
PATHS ${TBB_LIB_SEARCH_PATH})
make_library_set(TBB_LIBRARY)
findpkg_finish(TBB)
#if we haven't found TBB no point on going any further
if (NOT TBB_FOUND)
return()
endif ()
#=============================================================================
# Look for TBB's malloc package
set(TBB_MALLOC_LIBRARY_NAMES tbbmalloc)
get_debug_names(TBB_MALLOC_LIBRARY_NAMES)
find_path(TBB_MALLOC_INCLUDE_DIR
NAMES tbb/tbb.h
PATHS ${TBB_INC_SEARCH_PATH})
find_library(TBB_MALLOC_LIBRARY_RELEASE
NAMES ${TBB_MALLOC_LIBRARY_NAMES}
PATHS ${TBB_LIB_SEARCH_PATH})
find_library(TBB_MALLOC_LIBRARY_DEBUG
NAMES ${TBB_MALLOC_LIBRARY_NAMES_DEBUG}
PATHS ${TBB_LIB_SEARCH_PATH})
make_library_set(TBB_MALLOC_LIBRARY)
findpkg_finish(TBB_MALLOC)
#=============================================================================
# Look for TBB's malloc proxy package
set(TBB_MALLOC_PROXY_LIBRARY_NAMES tbbmalloc_proxy)
get_debug_names(TBB_MALLOC_PROXY_LIBRARY_NAMES)
find_path(TBB_MALLOC_PROXY_INCLUDE_DIR
NAMES tbb/tbbmalloc_proxy.h
PATHS ${TBB_INC_SEARCH_PATH})
find_library(TBB_MALLOC_PROXY_LIBRARY_RELEASE
NAMES ${TBB_MALLOC_PROXY_LIBRARY_NAMES}
PATHS ${TBB_LIB_SEARCH_PATH})
find_library(TBB_MALLOC_PROXY_LIBRARY_DEBUG
NAMES ${TBB_MALLOC_PROXY_LIBRARY_NAMES_DEBUG}
PATHS ${TBB_LIB_SEARCH_PATH})
make_library_set(TBB_MALLOC_PROXY_LIBRARY)
findpkg_finish(TBB_MALLOC_PROXY)
#=============================================================================
#parse all the version numbers from tbb
if(NOT TBB_VERSION)
#only read the start of the file
file(READ
"${TBB_INCLUDE_DIR}/tbb/tbb_stddef.h"
TBB_VERSION_CONTENTS
LIMIT 2048)
string(REGEX REPLACE
".*#define TBB_VERSION_MAJOR ([0-9]+).*" "\\1"
TBB_VERSION_MAJOR "${TBB_VERSION_CONTENTS}")
string(REGEX REPLACE
".*#define TBB_VERSION_MINOR ([0-9]+).*" "\\1"
TBB_VERSION_MINOR "${TBB_VERSION_CONTENTS}")
string(REGEX REPLACE
".*#define TBB_INTERFACE_VERSION ([0-9]+).*" "\\1"
TBB_INTERFACE_VERSION "${TBB_VERSION_CONTENTS}")
string(REGEX REPLACE
".*#define TBB_COMPATIBLE_INTERFACE_VERSION ([0-9]+).*" "\\1"
TBB_COMPATIBLE_INTERFACE_VERSION "${TBB_VERSION_CONTENTS}")
endif()

2
CMake/FindThrust.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
CMake/UseVTKmCuda.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
CMake/UseVTKmSerial.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

57
CMake/UseVTKmTBB.cmake Normal file
View File

@ -0,0 +1,57 @@
##============================================================================
## Copyright (c) Kitware, Inc.
## All rights reserved.
## See LICENSE.txt for details.
## This software is distributed WITHOUT ANY WARRANTY; without even
## the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
## PURPOSE. See the above copyright notice for more information.
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.
##
## Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
## Laboratory (LANL), the U.S. Government retains certain rights in
## this software.
##============================================================================
if (VTKm_TBB_initialize_complete)
return()
endif (VTKm_TBB_initialize_complete)
#-----------------------------------------------------------------------------
# Find TBB.
#-----------------------------------------------------------------------------
if (NOT VTKm_TBB_FOUND)
find_package(TBB REQUIRED)
set (VTKm_TBB_FOUND ${TBB_FOUND})
endif()
#-----------------------------------------------------------------------------
# Find the Boost library.
#-----------------------------------------------------------------------------
if (VTKm_TBB_FOUND)
if(NOT Boost_FOUND)
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION})
endif()
if (NOT Boost_FOUND)
message(STATUS "Boost not found")
set(VTKm_TBB_FOUND FALSE)
endif()
endif()
#-----------------------------------------------------------------------------
# Set up all these dependent packages (if they were all found).
#-----------------------------------------------------------------------------
if (VTKm_TBB_FOUND)
include_directories(
${Boost_INCLUDE_DIRS}
${VTKm_INCLUDE_DIRS}
${TBB_INCLUDE_DIRS}
)
set(VTKm_TBB_initialize_complete TRUE)
endif()

2
CMake/VTKmCheckCopyright.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
CMake/VTKmCheckPyexpander.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

4
CMake/VTKmCompilerExtras.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.
@ -34,7 +34,7 @@ if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_COMPILER_IS_CLANGXX)
"${CMAKE_CXX_FLAGS_DEBUG} ${CMAKE_CXX_FLAGS_WARN}")
# Addtional warnings for GCC
set(CMAKE_CXX_FLAGS_WARN_EXTRA "-Wno-long-long -Wcast-align -Wchar-subscripts -Wextra -Wpointer-arith -Wformat -Wformat-security -Wshadow -Wunused-parameter -fno-common")
set(CMAKE_CXX_FLAGS_WARN_EXTRA "-Wno-long-long -Wcast-align -Wconversion -Wchar-subscripts -Wextra -Wpointer-arith -Wformat -Wformat-security -Wshadow -Wunused-parameter -fno-common")
if (VTKm_FORCE_ANSI)
set(CMAKE_CXX_FLAGS_WARN_EXTRA "-ansi ${CMAKE_CXX_FLAGS_WARN_EXTRA}")
endif()

2
CMake/VTKmConfig.cmake.in
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
CMake/VTKmCopyrightStatement.txt
View File

@ -8,7 +8,7 @@ PURPOSE. See the above copyright notice for more information.
Copyright 2014 Sandia Corporation.
Copyright 2014 UT-Battelle, LLC.
Copyright 2014. Los Alamos National Security
Copyright 2014 Los Alamos National Security.
Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
the U.S. Government retains certain rights in this software.

5
CMake/VTKmMacros.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.
@ -340,6 +340,9 @@ function(vtkm_worklet_unit_tests device_adapter)
cuda_add_executable(${test_prog} ${unit_test_drivers} ${unit_test_srcs})
else()
add_executable(${test_prog} ${unit_test_drivers} ${unit_test_srcs})
if("${device_adapter}" STREQUAL "VTKM_DEVICE_ADAPTER_TBB")
target_link_libraries(${test_prog} ${TBB_LIBRARIES})
endif()
endif()
#add a test for each worklet test file. We will inject the device

2
CMake/VTKmSystemInformation.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

8
CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.
@ -54,13 +54,14 @@ include(CMake/VTKmCompilerExtras.cmake)
#-----------------------------------------------------------------------------
# Configurable Options
option(VTKm_ENABLE_CUDA "Enable Cuda support" OFF)
option(VTKm_ENABLE_TBB "Enable TBB support" OFF)
option(VTKm_ENABLE_TESTING "Enable VTKm Testing" ON)
option(VTKm_USE_DOUBLE_PRECISION
"Use double precision for floating point calculations"
OFF
)
option(VTKm_USE_64BIT_IDS "Use 64-bit indices." OFF)
option(VTKm_USE_64BIT_IDS "Use 64-bit indices." ON)
if (VTKm_ENABLE_TESTING)
enable_testing()
@ -73,6 +74,9 @@ vtkm_configure_device(Serial)
if (VTKm_ENABLE_CUDA)
vtkm_configure_device(Cuda)
endif (VTKm_ENABLE_CUDA)
if (VTKm_ENABLE_TBB)
vtkm_configure_device(TBB)
endif (VTKm_ENABLE_TBB)
#-----------------------------------------------------------------------------

2
CTestConfig.cmake
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
LICENSE.txt
View File

@ -72,4 +72,6 @@ licenses.
- - - - - - - - - - - - - - - - - - - - - - - - do not remove this line
CMake/CheckCXX11Features.cmake
CMake/FindBoostHeaders.cmake
CMake/FindTBB.cmake
vtkm/cont/tbb/internal/parallel_sort.h
vtkm/testing/OptionParser.h

2
docs/CodingConventions.md
View File

@ -18,7 +18,7 @@ The statement should have the following form:
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

94
vtkm/Extent.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -281,6 +281,52 @@ ExtentPointTopologyIndexToFlatIndex(const vtkm::Vec<vtkm::Id,Dimensions> &ijk,
return flatIndex;
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentPointTopologyIndexToFlatIndex<1>(const vtkm::Vec<vtkm::Id,1> &ijk,
const vtkm::Extent<1> &extent)
{
return ijk[0] - extent.Min[0];
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentPointTopologyIndexToFlatIndex<2>(const vtkm::Vec<vtkm::Id,2> &ijk,
const vtkm::Extent<2> &extent)
{
const vtkm::Vec<vtkm::Id,2> dims = ExtentPointDimensions(extent);
const vtkm::Vec<vtkm::Id,2> deltas = ijk - extent.Min;
return (deltas[1] * dims[0] + deltas[0]);
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentPointTopologyIndexToFlatIndex<3>(const vtkm::Vec<vtkm::Id,3> &ijk,
const vtkm::Extent<3> &extent)
{
const vtkm::Vec<vtkm::Id,3> dims = ExtentPointDimensions(extent);
const vtkm::Vec<vtkm::Id,3> deltas = ijk - extent.Min;
return (deltas[2] * dims[1] + deltas[1]) * dims[0] + deltas[0];
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
@ -302,6 +348,52 @@ ExtentCellTopologyIndexToFlatIndex(const vtkm::Vec<vtkm::Id,Dimensions> &ijk,
return flatIndex;
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentCellTopologyIndexToFlatIndex<1>(const vtkm::Vec<vtkm::Id,1> &ijk,
const vtkm::Extent<1> &extent)
{
return ijk[0] - extent.Min[0];
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentCellTopologyIndexToFlatIndex<2>(const vtkm::Vec<vtkm::Id,2> &ijk,
const vtkm::Extent<2> &extent)
{
const vtkm::Vec<vtkm::Id,2> dims = ExtentCellDimensions(extent);
const vtkm::Vec<vtkm::Id,2> deltas = ijk - extent.Min;
return (deltas[1] * dims[0] + deltas[0]);
}
/// Elements in structured grids have a single index with 0 being the entry at
/// the minimum extent in every direction and then increasing first in the r
/// then the s direction and so on. This method converts topological
/// coordinates to a flat index.
///
template<>
VTKM_EXEC_CONT_EXPORT
vtkm::Id
ExtentCellTopologyIndexToFlatIndex<3>(const vtkm::Vec<vtkm::Id,3> &ijk,
const vtkm::Extent<3> &extent)
{
const vtkm::Vec<vtkm::Id,3> dims = ExtentCellDimensions(extent);
const vtkm::Vec<vtkm::Id,3> deltas = ijk - extent.Min;
return (deltas[2] * dims[1] + deltas[1]) * dims[0] + deltas[0];
}
/// Given a cell index, returns the index to the first point incident on that
/// cell.
///

2
vtkm/ListTag.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/Pair.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/TypeListTag.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/TypeTraits.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

16
vtkm/Types.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -1159,12 +1159,16 @@ T dot(const vtkm::Vec<T,4> &a, const vtkm::Vec<T,4> &b)
return (a[0]*b[0]) + (a[1]*b[1]) + (a[2]*b[2]) + (a[3]*b[3]);
}
//Integer types of a width less than an integer get implicitly casted to
//an integer when doing a multiplication.
#define VTK_M_INTEGER_PROMOTION_SCALAR_DOT(type) \
VTKM_EXEC_CONT_EXPORT type dot(type a, type b) { return static_cast<type>(a * b); }
VTK_M_INTEGER_PROMOTION_SCALAR_DOT(vtkm::Int8)
VTK_M_INTEGER_PROMOTION_SCALAR_DOT(vtkm::UInt8)
VTK_M_INTEGER_PROMOTION_SCALAR_DOT(vtkm::Int16)
VTK_M_INTEGER_PROMOTION_SCALAR_DOT(vtkm::UInt16)
#define VTK_M_SCALAR_DOT(type) \
VTKM_EXEC_CONT_EXPORT type dot(type a, type b) { return a * b; }
VTK_M_SCALAR_DOT(vtkm::Int8)
VTK_M_SCALAR_DOT(vtkm::UInt8)
VTK_M_SCALAR_DOT(vtkm::Int16)
VTK_M_SCALAR_DOT(vtkm::UInt16)
VTK_M_SCALAR_DOT(vtkm::Int32)
VTK_M_SCALAR_DOT(vtkm::UInt32)
VTK_M_SCALAR_DOT(vtkm::Int64)
@ -1178,7 +1182,7 @@ VTK_M_SCALAR_DOT(vtkm::Float64)
template<typename T>
struct not_default_constructor
{
VTKM_EXEC_CONT_EXPORT bool operator()(const T &x)
VTKM_EXEC_CONT_EXPORT bool operator()(const T &x) const
{
return (x != T());
}

2
vtkm/VecTraits.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandle.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandleCompositeVector.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

85
vtkm/cont/ArrayHandleConstant.h Normal file
View File

@ -0,0 +1,85 @@
//=============================================================================
//
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//
//=============================================================================
#ifndef vtk_m_cont_ArrayHandleConstant_h
#define vtk_m_cont_ArrayHandleConstant_h
#include <vtkm/cont/ArrayHandleImplicit.h>
namespace vtkm {
namespace cont {
namespace detail {
template<typename ValueType>
struct ConstantFunctor
{
VTKM_EXEC_CONT_EXPORT
ConstantFunctor(const ValueType &value = ValueType()) : Value(value) { }
VTKM_EXEC_CONT_EXPORT
ValueType operator()(vtkm::Id vtkmNotUsed(index)) const
{
return this->Value;
}
private:
ValueType Value;
};
} // namespace detail
/// \brief An array handle with a constant value.
///
/// ArrayHandleConstant is an implicit array handle with a constant value. A
/// constant array handle is constructed by giving a value and an array length.
/// The resulting array is of the given size with each entry the same value
/// given in the constructor. The array is defined implicitly, so there it
/// takes (almost) no memory.
///
template<typename T>
class ArrayHandleConstant
: public vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T> >
{
typedef vtkm::cont::ArrayHandleImplicit<T, detail::ConstantFunctor<T> >
Superclass;
public:
VTKM_CONT_EXPORT
ArrayHandleConstant(T value = T(), vtkm::Id numberOfValues = 0)
: Superclass(detail::ConstantFunctor<T>(value), numberOfValues) { }
};
/// make_ArrayHandleImplicit is convenience function to generate an
/// ArrayHandleImplicit. It takes a functor and the virtual length of the
/// arry.
///
template<typename T>
vtkm::cont::ArrayHandleConstant<T>
make_ArrayHandleConstant(T value, vtkm::Id numberOfValues)
{
return vtkm::cont::ArrayHandleConstant<T>(value, numberOfValues);
}
}
} // vtkm::cont
#endif //vtk_m_cont_ArrayHandleConstant_h

4
vtkm/cont/ArrayHandleCounting.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -96,7 +96,7 @@ struct ArrayHandleCountingTraits
} // namespace internal
/// ArrayHandleCountings is a specialization of ArrayHandle. By default it
/// ArrayHandleCounting is a specialization of ArrayHandle. By default it
/// contains a increment value, that is increment for each step between zero
/// and the passed in length
template <typename CountingValueType>

2
vtkm/cont/ArrayHandleImplicit.h
View File

@ -10,7 +10,7 @@
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015. Los Alamos National Security
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandlePermutation.h
View File

@ -10,7 +10,7 @@
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015. Los Alamos National Security
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandleTransform.h
View File

@ -10,7 +10,7 @@
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015. Los Alamos National Security
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandleUniformPointCoordinates.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayHandleZip.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayPortal.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ArrayPortalToIterators.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/Assert.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

6
vtkm/cont/CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.
@ -23,6 +23,7 @@ include_directories(${Boost_INCLUDE_DIRS})
set(headers
ArrayHandle.h
ArrayHandleCompositeVector.h
ArrayHandleConstant.h
ArrayHandleCounting.h
ArrayHandleImplicit.h
ArrayHandlePermutation.h
@ -73,6 +74,9 @@ vtkm_declare_headers(${impl_headers} ${headers})
if (VTKm_ENABLE_CUDA)
add_subdirectory(cuda)
endif ()
if (VTKm_ENABLE_TBB)
add_subdirectory(tbb)
endif ()
#-----------------------------------------------------------------------------
add_subdirectory(testing)

2
vtkm/cont/DeviceAdapter.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

111
vtkm/cont/DeviceAdapterAlgorithm.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -101,12 +101,61 @@ struct DeviceAdapterAlgorithm
const vtkm::cont::ArrayHandle<vtkm::Id,CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id,COut>& values_output);
/// \brief Compute a accumulated sum operation on the input ArrayHandle
///
/// Computes an accumulated sum on the \c input ArrayHandle, returning the
/// total sum. Reduce is similar to the stl accumulate sum function,
/// exception that Reduce doesn't do a serial summation. This means that if
/// you have defined a custom plus operator for T it must be commutative,
/// or you will get inconsistent results.
///
/// \return The total sum.
template<typename T, class CIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input,
T initialValue);
/// \brief Compute a accumulated sum operation on the input ArrayHandle
///
/// Computes an accumulated sum (or any user binary operation) on the
/// \c input ArrayHandle, returning the total sum. Reduce is
/// similar to the stl accumulate sum function, exception that Reduce
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be commutative, or you will get
/// inconsistent results.
///
/// \return The total sum.
template<typename T, class CIn, class BinaryOperation>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input,
T initialValue,
BinaryOperation binaryOp);
/// \brief Compute a accumulated sum operation on the input key value pairs
///
/// Computes a segmented accumulated sum (or any user binary operation) on the
/// \c keys and \c values ArrayHandle(s). Each segmented accumulated sum is
/// run on consecutive equal keys with the binary operation applied to all
/// values inside that range. Once finished a single key and value is created
/// for each segment.
///
template<typename T,
class CKeyIn, class CValIn,
class CKeyOut, class CValOut,
class BinaryOperation >
VTKM_CONT_EXPORT static void ReduceByKey(
const vtkm::cont::ArrayHandle<T,CKeyIn> &keys,
const vtkm::cont::ArrayHandle<U,CValIn> &values,
vtkm::cont::ArrayHandle<T,CKeyOut>& keys_output,
vtkm::cont::ArrayHandle<T,CValOut>& values_output,
BinaryOperation binaryOp);
/// \brief Compute an inclusive prefix sum operation on the input ArrayHandle.
///
/// Computes an inclusive prefix sum operation on the \c input ArrayHandle,
/// storing the results in the \c output ArrayHandle. InclusiveScan is
/// similiar to the stl partial sum function, exception that InclusiveScan
/// doesn't do a serial sumnation. This means that if you have defined a
/// similar to the stl partial sum function, exception that InclusiveScan
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be associative, or you will get
/// inconsistent results. When the input and output ArrayHandles are the same
/// ArrayHandle the operation will be done inplace.
@ -118,12 +167,30 @@ struct DeviceAdapterAlgorithm
const vtkm::cont::ArrayHandle<T,CIn> &input,
vtkm::cont::ArrayHandle<T,COut>& output);
/// \brief Compute an inclusive prefix sum operation on the input ArrayHandle.
///
/// Computes an inclusive prefix sum operation on the \c input ArrayHandle,
/// storing the results in the \c output ArrayHandle. InclusiveScan is
/// similar to the stl partial sum function, exception that InclusiveScan
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be associative, or you will get
/// inconsistent results. When the input and output ArrayHandles are the same
/// ArrayHandle the operation will be done inplace.
///
/// \return The total sum.
///
template<typename T, class CIn, class COut, class BinaryOperation>
VTKM_CONT_EXPORT static T ScanInclusive(
const vtkm::cont::ArrayHandle<T,CIn> &input,
vtkm::cont::ArrayHandle<T,COut>& output,
BinaryOperation binaryOp);
/// \brief Compute an exclusive prefix sum operation on the input ArrayHandle.
///
/// Computes an exclusive prefix sum operation on the \c input ArrayHandle,
/// storing the results in the \c output ArrayHandle. ExclusiveScan is
/// similiar to the stl partial sum function, exception that ExclusiveScan
/// doesn't do a serial sumnation. This means that if you have defined a
/// similar to the stl partial sum function, exception that ExclusiveScan
/// doesn't do a serial summation. This means that if you have defined a
/// custom plus operator for T it must be associative, or you will get
/// inconsistent results. When the input and output ArrayHandles are the same
/// ArrayHandle the operation will be done inplace.
@ -205,7 +272,7 @@ struct DeviceAdapterAlgorithm
/// input to remove unwanted elements. The result of the stream compaction is
/// placed in \c output. The \c input values are used as the stream
/// compaction stencil while \c input indices are used as the values to place
/// into \c ouput. The size of \c output will be modified after this call as
/// into \c output. The size of \c output will be modified after this call as
/// we can't know the number of elements that will be removed by the stream
/// compaction algorithm.
///
@ -218,10 +285,11 @@ struct DeviceAdapterAlgorithm
///
/// Calls the parallel primitive function of stream compaction on the \c
/// input to remove unwanted elements. The result of the stream compaction is
/// placed in \c output. The values in \c stencil are used as the stream
/// compaction stencil while \c input values are placed into \c ouput. The
/// size of \c output will be modified after this call as we can't know the
/// number of elements that will be removed by the stream compaction
/// placed in \c output. The values in \c stencil are used to determine which
/// \c input values are placed into \c output, with all stencil values not
/// equal to the default constructor being considered valid.
/// The size of \c output will be modified after this call as we can't know
/// the number of elements that will be removed by the stream compaction
/// algorithm.
///
template<typename T, typename U, class CIn, class CStencil, class COut>
@ -230,6 +298,25 @@ struct DeviceAdapterAlgorithm
const vtkm::cont::ArrayHandle<U,CStencil> &stencil,
vtkm::cont::ArrayHandle<T,COut> &output);
/// \brief Performs stream compaction to remove unwanted elements in the input array.
///
/// Calls the parallel primitive function of stream compaction on the \c
/// input to remove unwanted elements. The result of the stream compaction is
/// placed in \c output. The values in \c stencil are passed to the unary
/// comparison object which is used to determine which /c input values are
/// placed into \c output.
/// The size of \c output will be modified after this call as we can't know
/// the number of elements that will be removed by the stream compaction
/// algorithm.
///
template<typename T, typename U, class CIn, class CStencil,
class COut, class PredicateOperator>
VTKM_CONT_EXPORT static void StreamCompact(
const vtkm::cont::ArrayHandle<T,CIn> &input,
const vtkm::cont::ArrayHandle<U,CStencil> &stencil,
vtkm::cont::ArrayHandle<T,COut> &output,
PredicateOperator predicate);
/// \brief Completes any asynchronous operations running on the device.
///
/// Waits for any asynchronous operations running on the device to complete.
@ -399,8 +486,8 @@ public:
#include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h>
// #elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
// #include <vtkm/openmp/cont/internal/DeviceAdapterAlgorithmOpenMP.h>
// #elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
// #include <vtkm/tbb/cont/internal/DeviceAdapterAlgorithmTBB.h>
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
#include <vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h>
#endif
#endif //vtk_m_cont_DeviceAdapterAlgorithm_h

2
vtkm/cont/DeviceAdapterSerial.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/DynamicArrayHandle.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/DynamicPointCoordinates.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/Error.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControl.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControlAssert.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControlBadType.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControlBadValue.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControlInternal.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorControlOutOfMemory.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/ErrorExecution.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/PointCoordinatesArray.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/PointCoordinatesListTag.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/PointCoordinatesUniform.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/Storage.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

8
vtkm/cont/StorageBasic.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -122,7 +122,8 @@ public:
if (this->DeallocateOnRelease)
{
AllocatorType allocator;
allocator.deallocate(this->Array, this->AllocatedSize);
allocator.deallocate(this->Array,
static_cast<std::size_t>(this->AllocatedSize) );
}
this->Array = NULL;
this->NumberOfValues = 0;
@ -150,7 +151,8 @@ public:
if (numberOfValues > 0)
{
AllocatorType allocator;
this->Array = allocator.allocate(numberOfValues);
this->Array = allocator.allocate(
static_cast<std::size_t>(numberOfValues) );
this->AllocatedSize = numberOfValues;
this->NumberOfValues = numberOfValues;
}

2
vtkm/cont/StorageImplicit.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/StorageListTag.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/Timer.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestControlSignatureTag.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestTransportArrayIn.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestTransportArrayOut.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestTransportExecObject.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestTypeCheckArray.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/arg/testing/UnitTestTypeCheckExecObject.cxx
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/cuda/CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
vtkm/cont/cuda/ChooseCudaDevice.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -26,26 +26,21 @@
#include <iostream>
// Disable GCC warnings we check vtkmfor but Thrust does not.
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
// Disable warnings we check vtkm for but Thrust does not.
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic push
#endif // gcc version >= 4.6
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 2)
#pragma GCC diagnostic ignored "-Wshadow"
#pragma GCC diagnostic ignored "-Wunused-parameter"
#endif // gcc version >= 4.2
#endif // gcc && !CUDA
#pragma GCC diagnostic ignored "-Wconversion"
#endif // gcc || clang
#include <thrust/system/cuda/vector.h>
#include <thrust/device_malloc_allocator.h>
#include <thrust/copy.h>
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic pop
#endif // gcc version >= 4.6
#endif // gcc && !CUDA
#endif // gcc || clang
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
@ -81,6 +76,7 @@ class ArrayManagerExecutionThrustDevice
{
public:
typedef T ValueType;
typedef typename thrust::system::cuda::pointer<T>::difference_type difference_type;
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
@ -104,7 +100,7 @@ public:
///
VTKM_CONT_EXPORT
vtkm::Id GetNumberOfValues() const {
return this->Array.size();
return static_cast<vtkm::Id>(this->Array.size());
}
/// Allocates the appropriate size of the array and copies the given data
@ -122,8 +118,9 @@ public:
// The data in this->Array should already be valid.
}
return PortalConstType(this->Array.data(),
this->Array.data() + this->Array.size());
this->Array.data() + static_cast<difference_type>(this->Array.size()));
}
/// Allocates the appropriate size of the array and copies the given data
@ -142,7 +139,7 @@ public:
}
return PortalType(this->Array.data(),
this->Array.data() + this->Array.size());
this->Array.data() + static_cast<difference_type>(this->Array.size()));
}
/// Allocates the array to the given size.
@ -159,7 +156,7 @@ public:
try
{
this->Array.resize(numberOfValues);
this->Array.resize(static_cast<vtkm::UInt64>(numberOfValues));
}
catch (std::bad_alloc error)
{
@ -167,7 +164,7 @@ public:
}
return PortalType(this->Array.data(),
this->Array.data() + this->Array.size());
this->Array.data() + static_cast<difference_type>(this->Array.size()));
}
/// Allocates enough space in \c storage and copies the data in the
@ -176,10 +173,10 @@ public:
VTKM_CONT_EXPORT
void RetrieveOutputData(StorageType *storage) const
{
storage->Allocate(this->Array.size());
storage->Allocate(static_cast<vtkm::Id>(this->Array.size()));
::thrust::copy(
this->Array.data(),
this->Array.data() + this->Array.size(),
this->Array.data() + static_cast<difference_type>(this->Array.size()),
vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
}
@ -189,9 +186,9 @@ public:
{
// The operation will succeed even if this assertion fails, but this
// is still supposed to be a precondition to Shrink.
VTKM_ASSERT_CONT(numberOfValues <= this->Array.size());
VTKM_ASSERT_CONT(numberOfValues <= static_cast<vtkm::Id>(this->Array.size()));
this->Array.resize(numberOfValues);
this->Array.resize(static_cast<vtkm::UInt64>(numberOfValues));
}

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

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -9,7 +9,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -21,6 +21,7 @@
#ifndef vtk_m_cont_cuda_internal_DeviceAdapterThrust_h
#define vtk_m_cont_cuda_internal_DeviceAdapterThrust_h
#include <vtkm/Types.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ErrorExecution.h>
@ -30,16 +31,13 @@
#include <vtkm/exec/internal/ErrorMessageBuffer.h>
#include <vtkm/exec/internal/WorkletInvokeFunctor.h>
// Disable GCC warnings we check vtkmfor but Thrust does not.
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
// Disable warnings we check vtkm for but Thrust does not.
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic push
#endif // gcc version >= 4.6
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 2)
#pragma GCC diagnostic ignored "-Wshadow"
#pragma GCC diagnostic ignored "-Wunused-parameter"
#endif // gcc version >= 4.2
#endif // gcc && !CUDA
#pragma GCC diagnostic ignored "-Wconversion"
#endif // gcc || clang
#include <thrust/advance.h>
#include <thrust/binary_search.h>
@ -52,11 +50,9 @@
#include <thrust/iterator/counting_iterator.h>
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic pop
#endif // gcc version >= 4.6
#endif // gcc && !CUDA
#endif // gcc || clang
namespace vtkm {
namespace cont {
@ -76,28 +72,28 @@ void Schedule1DIndexKernel(FunctorType functor, vtkm::Id length)
template<class FunctorType>
__global__
void Schedule3DIndexKernel(FunctorType functor, vtkm::Id3 size)
void Schedule3DIndexKernel(FunctorType functor, dim3 size)
{
const vtkm::Id x = blockIdx.x*blockDim.x + threadIdx.x;
const vtkm::Id y = blockIdx.y*blockDim.y + threadIdx.y;
const vtkm::Id z = blockIdx.z*blockDim.z + threadIdx.z;
if (x >= size[0] || y >= size[1] || z >= size[2])
if (x >= size.x || y >= size.y || z >= size.z)
{
return;
}
//now convert back to flat memory
const int idx = x + size[0]*(y + size[1]*z);
const int idx = x + size.x*(y + size.y*z);
functor( idx );
}
inline
void compute_block_size(vtkm::Id3 rangeMax, dim3 blockSize3d, dim3& gridSize3d)
void compute_block_size(dim3 rangeMax, dim3 blockSize3d, dim3& gridSize3d)
{
gridSize3d.x = (rangeMax[0] % blockSize3d.x != 0) ? (rangeMax[0] / blockSize3d.x + 1) : (rangeMax[0] / blockSize3d.x);
gridSize3d.y = (rangeMax[1] % blockSize3d.y != 1) ? (rangeMax[1] / blockSize3d.y + 1) : (rangeMax[1] / blockSize3d.y);
gridSize3d.z = (rangeMax[2] % blockSize3d.z != 2) ? (rangeMax[2] / blockSize3d.z + 1) : (rangeMax[2] / blockSize3d.z);
gridSize3d.x = (rangeMax.x % blockSize3d.x != 0) ? (rangeMax.x / blockSize3d.x + 1) : (rangeMax.x / blockSize3d.x);
gridSize3d.y = (rangeMax.y % blockSize3d.y != 1) ? (rangeMax.y / blockSize3d.y + 1) : (rangeMax.y / blockSize3d.y);
gridSize3d.z = (rangeMax.z % blockSize3d.z != 2) ? (rangeMax.z / blockSize3d.z + 1) : (rangeMax.z / blockSize3d.z);
}
class PerfRecord
@ -121,14 +117,17 @@ public:
template<class Functor>
static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& rangeMax)
{
const dim3 ranges(static_cast<vtkm::UInt32>(rangeMax[0]),
static_cast<vtkm::UInt32>(rangeMax[1]),
static_cast<vtkm::UInt32>(rangeMax[2]) );
std::vector< PerfRecord > results;
int indexTable[16] = {1, 2, 4, 8, 12, 16, 20, 24, 28, 30, 32, 64, 128, 256, 512, 1024};
vtkm::UInt32 indexTable[16] = {1, 2, 4, 8, 12, 16, 20, 24, 28, 30, 32, 64, 128, 256, 512, 1024};
for(int i=0; i < 16; i++)
for(vtkm::UInt32 i=0; i < 16; i++)
{
for(int j=0; j < 16; j++)
for(vtkm::UInt32 j=0; j < 16; j++)
{
for(int k=0; k < 16; k++)
for(vtkm::UInt32 k=0; k < 16; k++)
{
cudaEvent_t start, stop;
cudaEventCreate(&start);
@ -150,9 +149,9 @@ static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& range
continue;
}
compute_block_size(rangeMax, blockSize3d, gridSize3d);
compute_block_size(ranges, blockSize3d, gridSize3d);
cudaEventRecord(start, 0);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, rangeMax);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, ranges);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
@ -169,12 +168,14 @@ static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& range
}
std::sort(results.begin(), results.end());
for(int i=results.size()-1; i >= 0; --i)
const vtkm::Int64 size = static_cast<vtkm::Int64>(results.size());
for(vtkm::Int64 i=1; i <= size; i++)
{
int x = results[i].blockSize.x;
int y = results[i].blockSize.y;
int z = results[i].blockSize.z;
float t = results[i].elapsedTime;
vtkm::UInt64 index = static_cast<vtkm::UInt64>(size-i);
vtkm::UInt32 x = results[index].blockSize.x;
vtkm::UInt32 y = results[index].blockSize.y;
vtkm::UInt32 z = results[index].blockSize.z;
float t = results[index].elapsedTime;
std::cout << "BlockSize of: " << x << "," << y << "," << z << " required: " << t << std::endl;
}
@ -185,9 +186,9 @@ static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& range
cudaEventCreate(&start);
cudaEventCreate(&stop);
const vtkm::Id numInstances = rangeMax[0] * rangeMax[1] * rangeMax[2];
const int blockSize = 128;
const int blocksPerGrid = (numInstances + blockSize - 1) / blockSize;
const vtkm::UInt32 numInstances = static_cast<vtkm::UInt32>(rangeMax[0] * rangeMax[1] * rangeMax[2]);
const vtkm::UInt32 blockSize = 128;
const vtkm::UInt32 blocksPerGrid = (numInstances + blockSize - 1) / blockSize;
cudaEventRecord(start, 0);
Schedule1DIndexKernel<Functor> <<<blocksPerGrid, blockSize>>> (functor, numInstances);
@ -212,9 +213,9 @@ static void compare_3d_schedule_patterns(Functor functor, const vtkm::Id3& range
dim3 blockSize3d(64,2,1);
dim3 gridSize3d;
compute_block_size(rangeMax, blockSize3d, gridSize3d);
compute_block_size(ranges, blockSize3d, gridSize3d);
cudaEventRecord(start, 0);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, rangeMax);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, ranges);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
@ -290,6 +291,62 @@ private:
IteratorBegin(values_output));
}
template<class InputPortal>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ReducePortal(const InputPortal &input,
typename InputPortal::ValueType initialValue)
{
return ::thrust::reduce(IteratorBegin(input),
IteratorEnd(input),
initialValue);
}
template<class InputPortal, class BinaryOperation>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ReducePortal(const InputPortal &input,
typename InputPortal::ValueType initialValue,
BinaryOperation binaryOP)
{
return ::thrust::reduce(IteratorBegin(input),
IteratorEnd(input),
initialValue,
binaryOP);
}
template<class KeysPortal, class ValuesPortal,
class KeysOutputPortal, class ValueOutputPortal,
class BinaryOperation>
VTKM_CONT_EXPORT static
vtkm::Id ReduceByKeyPortal(const KeysPortal &keys,
const ValuesPortal& values,
const KeysOutputPortal &keys_output,
const ValueOutputPortal &values_output,
BinaryOperation binaryOP)
{
typedef typename detail::IteratorTraits<KeysOutputPortal>::IteratorType
KeysIteratorType;
typedef typename detail::IteratorTraits<ValueOutputPortal>::IteratorType
ValuesIteratorType;
KeysIteratorType keys_out_begin = IteratorBegin(keys_output);
ValuesIteratorType values_out_begin = IteratorBegin(values_output);
::thrust::pair< KeysIteratorType, ValuesIteratorType > result_iterators;
::thrust::equal_to<typename KeysPortal::ValueType> binaryPredicate;
result_iterators = ::thrust::reduce_by_key(IteratorBegin(keys),
IteratorEnd(keys),
IteratorBegin(values),
keys_out_begin,
values_out_begin,
binaryPredicate,
binaryOP);
return static_cast<vtkm::Id>( ::thrust::distance(keys_out_begin,
result_iterators.first) );
}
template<class InputPortal, class OutputPortal>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ScanExclusivePortal(const InputPortal &input,
@ -320,6 +377,21 @@ private:
return *(IteratorEnd(output) - 1);
}
template<class InputPortal, class OutputPortal, class BinaryOperation>
VTKM_CONT_EXPORT static
typename InputPortal::ValueType ScanInclusivePortal(const InputPortal &input,
const OutputPortal &output,
BinaryOperation binaryOp)
{
::thrust::inclusive_scan(IteratorBegin(input),
IteratorEnd(input),
IteratorBegin(output),
binaryOp);
//return the value at the last index in the array, as that is the sum
return *(IteratorEnd(output) - 1);
}
template<class ValuesPortal>
VTKM_CONT_EXPORT static void SortPortal(const ValuesPortal &values)
{
@ -357,61 +429,45 @@ private:
comp);
}
template<class StencilPortal>
VTKM_CONT_EXPORT static vtkm::Id CountIfPortal(const StencilPortal &stencil)
{
typedef typename StencilPortal::ValueType ValueType;
return ::thrust::count_if(IteratorBegin(stencil),
IteratorEnd(stencil),
::vtkm::not_default_constructor<ValueType>());
}
template<class ValueIterator,
class StencilPortal,
class OutputPortal>
VTKM_CONT_EXPORT static void CopyIfPortal(ValueIterator valuesBegin,
ValueIterator valuesEnd,
const StencilPortal &stencil,
const OutputPortal &output)
{
typedef typename StencilPortal::ValueType ValueType;
::thrust::copy_if(valuesBegin,
valuesEnd,
IteratorBegin(stencil),
IteratorBegin(output),
::vtkm::not_default_constructor<ValueType>());
}
template<class ValueIterator,
class StencilArrayHandle,
class OutputArrayHandle>
VTKM_CONT_EXPORT static void RemoveIf(ValueIterator valuesBegin,
ValueIterator valuesEnd,
const StencilArrayHandle& stencil,
OutputArrayHandle& output)
{
vtkm::Id numLeft = CountIfPortal(stencil.PrepareForInput(DeviceAdapterTag()));
CopyIfPortal(valuesBegin,
valuesEnd,
stencil.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numLeft, DeviceAdapterTag()));
}
template<class InputPortal,
class StencilArrayHandle,
class OutputArrayHandle>
class OutputPortal,
class PredicateOperator>
VTKM_CONT_EXPORT static
void StreamCompactPortal(const InputPortal& inputPortal,
const StencilArrayHandle &stencil,
OutputArrayHandle& output)
vtkm::Id CopyIfPortal(ValueIterator valuesBegin,
ValueIterator valuesEnd,
StencilPortal stencil,
OutputPortal output,
PredicateOperator predicate)
{
RemoveIf(IteratorBegin(inputPortal),
IteratorEnd(inputPortal),
stencil,
output);
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType
IteratorType;
IteratorType outputBegin = IteratorBegin(output);
IteratorType newLast = ::thrust::copy_if(valuesBegin,
valuesEnd,
IteratorBegin(stencil),
outputBegin,
predicate);
return static_cast<vtkm::Id>( ::thrust::distance(outputBegin, newLast) );
}
template<class ValuePortal,
class StencilPortal,
class OutputPortal,
class PredicateOperator>
VTKM_CONT_EXPORT static
vtkm::Id CopyIfPortal(ValuePortal values,
StencilPortal stencil,
OutputPortal output,
PredicateOperator predicate)
{
return CopyIfPortal(IteratorBegin(values),
IteratorEnd(values),
stencil,
output,
predicate);
}
template<class ValuesPortal>
@ -422,7 +478,7 @@ private:
IteratorType;
IteratorType begin = IteratorBegin(values);
IteratorType newLast = ::thrust::unique(begin, IteratorEnd(values));
return ::thrust::distance(begin, newLast);
return static_cast<vtkm::Id>( ::thrust::distance(begin, newLast) );
}
template<class ValuesPortal, class Compare>
@ -433,7 +489,7 @@ private:
IteratorType;
IteratorType begin = IteratorBegin(values);
IteratorType newLast = ::thrust::unique(begin, IteratorEnd(values), comp);
return ::thrust::distance(begin, newLast);
return static_cast<vtkm::Id>( ::thrust::distance(begin, newLast) );
}
template<class InputPortal, class ValuesPortal, class OutputPortal>
@ -485,7 +541,7 @@ public:
const vtkm::cont::ArrayHandle<T,SIn> &input,
vtkm::cont::ArrayHandle<T,SOut> &output)
{
vtkm::Id numberOfValues = input.GetNumberOfValues();
const vtkm::Id numberOfValues = input.GetNumberOfValues();
CopyPortal(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()));
}
@ -525,12 +581,69 @@ public:
values_output.PrepareForInPlace(DeviceAdapterTag()));
}
template<typename T, class SIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,SIn> &input,
T initialValue)
{
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
return initialValue;
}
return ReducePortal(input.PrepareForInput( DeviceAdapterTag() ),
initialValue);
}
template<typename T, class SIn, class BinaryOperation>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,SIn> &input,
T initialValue,
BinaryOperation binaryOp)
{
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
return initialValue;
}
return ReducePortal(input.PrepareForInput( DeviceAdapterTag() ),
initialValue,
binaryOp);
}
template<typename T, typename U, class KIn, class VIn, class KOut, class VOut,
class BinaryOperation>
VTKM_CONT_EXPORT static void ReduceByKey(
const vtkm::cont::ArrayHandle<T,KIn> &keys,
const vtkm::cont::ArrayHandle<U,VIn> &values,
vtkm::cont::ArrayHandle<T,KOut> &keys_output,
vtkm::cont::ArrayHandle<U,VOut> &values_output,
BinaryOperation binaryOp)
{
//there is a concern that by default we will allocate too much
//space for the keys/values output. 1 option is to
const vtkm::Id numberOfValues = keys.GetNumberOfValues();
if (numberOfValues <= 0)
{
return;
}
vtkm::Id reduced_size =
ReduceByKeyPortal(keys.PrepareForInput( DeviceAdapterTag() ),
values.PrepareForInput( DeviceAdapterTag() ),
keys_output.PrepareForOutput( numberOfValues, DeviceAdapterTag() ),
values_output.PrepareForOutput( numberOfValues, DeviceAdapterTag() ),
binaryOp);
keys_output.Shrink( reduced_size );
values_output.Shrink( reduced_size );
}
template<typename T, class SIn, class SOut>
VTKM_CONT_EXPORT static T ScanExclusive(
const vtkm::cont::ArrayHandle<T,SIn> &input,
vtkm::cont::ArrayHandle<T,SOut>& output)
{
vtkm::Id numberOfValues = input.GetNumberOfValues();
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTag());
@ -540,12 +653,13 @@ public:
return ScanExclusivePortal(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()));
}
template<typename T, class SIn, class SOut>
VTKM_CONT_EXPORT static T ScanInclusive(
const vtkm::cont::ArrayHandle<T,SIn> &input,
vtkm::cont::ArrayHandle<T,SOut>& output)
{
vtkm::Id numberOfValues = input.GetNumberOfValues();
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTag());
@ -556,6 +670,24 @@ public:
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()));
}
template<typename T, class SIn, class SOut, class BinaryOperation>
VTKM_CONT_EXPORT static T ScanInclusive(
const vtkm::cont::ArrayHandle<T,SIn> &input,
vtkm::cont::ArrayHandle<T,SOut>& output,
BinaryOperation binaryOp)
{
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTag());
return 0;
}
return ScanInclusivePortal(input.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTag()),
binaryOp);
}
// Because of some funny code conversions in nvcc, kernels for devices have to
// be public.
#ifndef VTKM_CUDA
@ -603,8 +735,8 @@ public:
functor.SetErrorMessageBuffer(errorMessage);
const int blockSize = 128;
const int blocksPerGrid = (numInstances + blockSize - 1) / blockSize;
const vtkm::UInt32 blockSize = 128;
const vtkm::UInt32 blocksPerGrid = (static_cast<vtkm::UInt32>(numInstances) + blockSize - 1) / blockSize;
Schedule1DIndexKernel<Functor> <<<blocksPerGrid, blockSize>>> (functor, numInstances);
@ -643,6 +775,9 @@ public:
//requires the errormessage buffer be set
compare_3d_schedule_patterns(functor,rangeMax);
#endif
const dim3 ranges(static_cast<vtkm::UInt32>(rangeMax[0]),
static_cast<vtkm::UInt32>(rangeMax[1]),
static_cast<vtkm::UInt32>(rangeMax[2]) );
//currently we presume that 3d scheduling access patterns prefer accessing
//memory in the X direction. Also should be good for thin in the Z axis
@ -658,8 +793,8 @@ public:
}
dim3 gridSize3d;
compute_block_size(rangeMax, blockSize3d, gridSize3d);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, rangeMax);
compute_block_size(ranges, blockSize3d, gridSize3d);
Schedule3DIndexKernel<Functor> <<<gridSize3d, blockSize3d>>> (functor, ranges);
//sync so that we can check the results of the call.
//In the future I want move this before the schedule call, and throwing
@ -718,12 +853,13 @@ public:
const vtkm::cont::ArrayHandle<T,SStencil>& stencil,
vtkm::cont::ArrayHandle<vtkm::Id,SOut>& output)
{
vtkm::Id stencilSize = stencil.GetNumberOfValues();
RemoveIf(::thrust::make_counting_iterator<vtkm::Id>(0),
::thrust::make_counting_iterator<vtkm::Id>(stencilSize),
stencil,
output);
vtkm::Id size = stencil.GetNumberOfValues();
vtkm::Id newSize = CopyIfPortal(::thrust::make_counting_iterator<vtkm::Id>(0),
::thrust::make_counting_iterator<vtkm::Id>(size),
stencil.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(size, DeviceAdapterTag()),
::vtkm::not_default_constructor<T>());
output.Shrink(newSize);
}
template<typename T,
@ -736,7 +872,32 @@ public:
const vtkm::cont::ArrayHandle<T,SStencil>& stencil,
vtkm::cont::ArrayHandle<U,SOut>& output)
{
StreamCompactPortal(input.PrepareForInput(DeviceAdapterTag()), stencil, output);
vtkm::Id size = stencil.GetNumberOfValues();
vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTag()),
stencil.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(size, DeviceAdapterTag()),
::vtkm::not_default_constructor<T>()); //yes on the stencil
output.Shrink(newSize);
}
template<typename T,
typename U,
class SIn,
class SStencil,
class SOut,
class PredicateOperator>
VTKM_CONT_EXPORT static void StreamCompact(
const vtkm::cont::ArrayHandle<U,SIn>& input,
const vtkm::cont::ArrayHandle<T,SStencil>& stencil,
vtkm::cont::ArrayHandle<U,SOut>& output,
PredicateOperator predicate)
{
vtkm::Id size = stencil.GetNumberOfValues();
vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTag()),
stencil.PrepareForInput(DeviceAdapterTag()),
output.PrepareForOutput(size, DeviceAdapterTag()),
predicate);
output.Shrink(newSize);
}
template<typename T, class Storage>

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -25,28 +25,22 @@
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
// Disable GCC warnings we check vtkmfor but Thrust does not.
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
// Disable warnings we check vtkm for but Thrust does not.
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic push
#endif // gcc version >= 4.6
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 2)
#pragma GCC diagnostic ignored "-Wshadow"
#pragma GCC diagnostic ignored "-Wunused-parameter"
#endif // gcc version >= 4.2
#endif // gcc && !CUDA
#pragma GCC diagnostic ignored "-Wconversion"
#endif // gcc || clang
#include <thrust/system/cuda/memory.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#if defined(__GNUC__) && !defined(VTKM_CUDA)
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic pop
#endif // gcc version >= 4.6
#endif // gcc && !CUDA
#endif // gcc || clang
namespace vtkm {
namespace cont {

2
vtkm/cont/cuda/internal/testing/CMakeLists.txt
View File

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

2
vtkm/cont/cuda/internal/testing/Testing.h
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

8
vtkm/cont/cuda/testing/UnitTestCudaArrayHandle.cu
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -53,9 +53,9 @@ struct CountingTest
void operator()(const ValueType v) const
{
std::vector< ValueType > inputVector(ARRAY_SIZE);
for(int i=0; i < ARRAY_SIZE; ++i)
for(vtkm::Id i=0; i < ARRAY_SIZE; ++i)
{
inputVector[i] = v + i;
inputVector[static_cast<vtkm::UInt32>(i)] = v + i;
}
vtkm::cont::ArrayHandle< ValueType > input =
@ -65,7 +65,7 @@ struct CountingTest
dispatcher.Invoke(input, result);
//verify that the control portal works
for(int i=v; i < ARRAY_SIZE; ++i)
for(vtkm::Id i=static_cast<vtkm::Id>(v); i < ARRAY_SIZE; ++i)
{
const ValueType result_v = result.GetPortalConstControl().Get(i);
const ValueType correct_value = v + ValueType(i);

2
vtkm/cont/cuda/testing/UnitTestCudaArrayHandleFancy.cu
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

2
vtkm/cont/cuda/testing/UnitTestDeviceAdapterCuda.cu
View File

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -162,8 +162,8 @@ public:
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
// #elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_OPENMP
// #include <vtkm/openmp/cont/internal/ArrayManagerExecutionOpenMP.h>
// #elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
// #include <vtkm/tbb/cont/internal/ArrayManagerExecutionTBB.h>
#elif VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_TBB
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
#endif
#endif //vtk_m_cont_internal_ArrayManagerExecution_h

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
##
## Copyright 2014 Sandia Corporation.
## Copyright 2014 UT-Battelle, LLC.
## Copyright 2014. Los Alamos National Security
## Copyright 2014 Los Alamos National Security.
##
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
## the U.S. Government retains certain rights in this software.

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

@ -8,7 +8,7 @@
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014. Los Alamos National Security
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
@ -22,8 +22,9 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/cont/ArrayHandleZip.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/StorageBasic.h>
#include <vtkm/exec/FunctorBase.h>
@ -36,6 +37,60 @@ namespace vtkm {
namespace cont {
namespace internal {
// Binary function object wrapper which can detect and handle calling the
// wrapped operator with complex value types such as
// IteratorFromArrayPortalValue which happen when passed an input array that
// is implicit.
template<typename ResultType, typename Function>
struct WrappedBinaryOperator
{
Function m_f;
VTKM_CONT_EXPORT
WrappedBinaryOperator(const Function &f)
: m_f(f)
{}
template<typename Argument1, typename Argument2>
VTKM_CONT_EXPORT ResultType operator()(const Argument1 &x, const Argument2 &y) const
{
return m_f(x, y);
}
template<typename Argument1, typename Argument2>
VTKM_CONT_EXPORT ResultType operator()(
const detail::IteratorFromArrayPortalValue<Argument1> &x,
const detail::IteratorFromArrayPortalValue<Argument2> &y) const
{
typedef typename detail::IteratorFromArrayPortalValue<Argument1>::ValueType
ValueTypeX;
typedef typename detail::IteratorFromArrayPortalValue<Argument2>::ValueType
ValueTypeY;
return m_f( (ValueTypeX)x, (ValueTypeY)y );
}
template<typename Argument1, typename Argument2>
VTKM_CONT_EXPORT ResultType operator()(
const Argument1 &x,
const detail::IteratorFromArrayPortalValue<Argument2> &y) const
{
typedef typename detail::IteratorFromArrayPortalValue<Argument2>::ValueType
ValueTypeY;
return m_f( x, (ValueTypeY)y );
}
template<typename Argument1, typename Argument2>
VTKM_CONT_EXPORT ResultType operator()(
const detail::IteratorFromArrayPortalValue<Argument1> &x,
const Argument2 &y) const
{
typedef typename detail::IteratorFromArrayPortalValue<Argument1>::ValueType
ValueTypeX;
return m_f( (ValueTypeX)x, y );
}
};
/// \brief
///
/// This struct provides algorithms that implement "general" device adapter
@ -316,6 +371,336 @@ public:
values_output);
}
//--------------------------------------------------------------------------
// Reduce
private:
template<int ReduceWidth, typename T, typename ArrayType, typename BinaryOperation >
struct ReduceKernel : vtkm::exec::FunctorBase
{
typedef typename ArrayType::template ExecutionTypes<
DeviceAdapterTag> ExecutionTypes;
typedef typename ExecutionTypes::PortalConst PortalConst;
PortalConst Portal;
BinaryOperation BinaryOperator;
vtkm::Id ArrayLength;
VTKM_CONT_EXPORT
ReduceKernel()
: Portal(),
BinaryOperator(),
ArrayLength(0)
{
}
VTKM_CONT_EXPORT
ReduceKernel(const ArrayType &array, BinaryOperation op)
: Portal(array.PrepareForInput( DeviceAdapterTag() ) ),
BinaryOperator(op),
ArrayLength( array.GetNumberOfValues() )
{ }
VTKM_EXEC_EXPORT
T operator()(vtkm::Id index) const
{
const vtkm::Id offset = index * ReduceWidth;
//at least the first value access to the portal will be valid
//only the rest could be invalid
T partialSum = this->Portal.Get( offset );
if( offset + ReduceWidth >= this->ArrayLength )
{
vtkm::Id currentIndex = offset + 1;
while( currentIndex < this->ArrayLength)
{
partialSum = BinaryOperator(partialSum, this->Portal.Get(currentIndex));
++currentIndex;
}
}
else
{
//optimize the usecase where all values are valid and we don't
//need to check that we might go out of bounds
for(int i=1; i < ReduceWidth; ++i)
{
partialSum = BinaryOperator(partialSum,
this->Portal.Get( offset + i )
);
}
}
return partialSum;
}
};
//--------------------------------------------------------------------------
// Reduce
public:
template<typename T, class CIn>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input, T initialValue)
{
return DerivedAlgorithm::Reduce(input, initialValue, vtkm::internal::Add());
}
template<typename T, class CIn, class BinaryOperator>
VTKM_CONT_EXPORT static T Reduce(
const vtkm::cont::ArrayHandle<T,CIn> &input,
T initialValue,
BinaryOperator binaryOp)
{
//Crazy Idea:
//We create a implicit array handle that wraps the input
//array handle. The implicit functor is passed the input array handle, and
//the number of elements it needs to sum. This way the implicit handle
//acts as the first level reduction. Say for example reducing 16 values
//at a time.
//
//Now that we have an implicit array that is 1/16 the length of full array
//we can use scan inclusive to compute the final sum
typedef ReduceKernel<
16,
T,
vtkm::cont::ArrayHandle<T,CIn>,
BinaryOperator
> ReduceKernelType;
typedef vtkm::cont::ArrayHandleImplicit<
T,
ReduceKernelType > ReduceHandleType;
typedef vtkm::cont::ArrayHandle<
T,
vtkm::cont::StorageTagBasic> TempArrayType;
ReduceKernelType kernel(input, binaryOp);
vtkm::Id length = (input.GetNumberOfValues() / 16);
length += (input.GetNumberOfValues() % 16 == 0) ? 0 : 1;
ReduceHandleType reduced = vtkm::cont::make_ArrayHandleImplicit<T>(kernel,
length);
TempArrayType inclusiveScanStorage;
T scanResult = DerivedAlgorithm::ScanInclusive(reduced,
inclusiveScanStorage,
binaryOp);
return binaryOp(initialValue, scanResult);
}
//--------------------------------------------------------------------------
// Reduce By Key
private:
struct ReduceKeySeriesStates
{
//It is needed that END and START_AND_END are both odd numbers
//so that the first bit of both are 1
enum { MIDDLE=0, END=1, START=2, START_AND_END=3};
};
template<typename InputPortalType>
struct ReduceStencilGeneration : vtkm::exec::FunctorBase
{
typedef typename vtkm::cont::ArrayHandle< vtkm::UInt8 >::template ExecutionTypes<DeviceAdapterTag>
::Portal KeyStatePortalType;
InputPortalType Input;
KeyStatePortalType KeyState;
VTKM_CONT_EXPORT
ReduceStencilGeneration(const InputPortalType &input,
const KeyStatePortalType &kstate)
: Input(input),
KeyState(kstate)
{ }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id centerIndex) const
{
typedef ReduceKeySeriesStates States;
typedef typename InputPortalType::ValueType ValueType;
typedef typename KeyStatePortalType::ValueType KeyStateType;
const vtkm::Id leftIndex = centerIndex - 1;
const vtkm::Id rightIndex = centerIndex + 1;
//we need to determine which of three states this
//index is. It can be:
// 1. Middle of a set of equivalent keys.
// 2. Start of a set of equivalent keys.
// 3. End of a set of equivalent keys.
// 4. Both the start and end of a set of keys
//we don't have to worry about an array of length 1, as
//the calling code handles that use case
if(centerIndex == 0)
{
//this means we are at the start of the array
//means we are automatically START
//just need to check if we are END
const ValueType centerValue = this->Input.Get(centerIndex);
const ValueType rightValue = this->Input.Get(rightIndex);
const KeyStateType state = (rightValue == centerValue) ? States::START :
States::START_AND_END;
this->KeyState.Set(centerIndex, state);
}
else if(rightIndex == this->Input.GetNumberOfValues())
{
//this means we are at the end, so we are at least END
//just need to check if we are START
const ValueType centerValue = this->Input.Get(centerIndex);
const ValueType leftValue = this->Input.Get(leftIndex);
const KeyStateType state = (leftValue == centerValue) ? States::END :
States::START_AND_END;
this->KeyState.Set(centerIndex, state);
}
else
{
const ValueType centerValue = this->Input.Get(centerIndex);
const bool leftMatches(this->Input.Get(leftIndex) == centerValue);
const bool rightMatches(this->Input.Get(rightIndex) == centerValue);
//assume it is the middle, and check for the other use-case
KeyStateType state = States::MIDDLE;
if(!leftMatches && rightMatches)
{
state = States::START;
}
else if(leftMatches && !rightMatches)
{
state = States::END;
}
else if(!leftMatches && !rightMatches)
{
state = States::START_AND_END;
}
this->KeyState.Set(centerIndex, state);
}
}
};
struct ReduceByKeyAdd
{
template<typename T>
vtkm::Pair<T, vtkm::UInt8> operator()(const vtkm::Pair<T, vtkm::UInt8>& a,
const vtkm::Pair<T, vtkm::UInt8>& b) const
{
typedef vtkm::Pair<T, vtkm::UInt8> ReturnType;
typedef ReduceKeySeriesStates States;
//need too handle how we are going to add two numbers together
//based on the keyStates that they have
//need to optimize this logic, we can use a bit mask to determine
//the secondary value.
if(a.second == States::START && b.second == States::END)
{
return ReturnType(a.first + b.first, States::START_AND_END); //with second type as START_AND_END
}
else if((a.second == States::START || a.second == States::MIDDLE) &&
(b.second == States::MIDDLE || b.second == States::END))
{
//note that we cant have START + END as that is handled above
//as a special use case
return ReturnType(a.first + b.first, b.second); //with second type as b.second
}
else
{
return b;
}
}
};
struct ReduceByKeyUnaryStencilOp
{
bool operator()(vtkm::UInt8 keySeriesState) const
{
typedef ReduceKeySeriesStates States;
return (keySeriesState == States::END ||
keySeriesState == States::START_AND_END);
}
};
public:
template<typename T, typename U, class KIn, class VIn, class KOut, class VOut,
class BinaryOperation>
VTKM_CONT_EXPORT static void ReduceByKey(
const vtkm::cont::ArrayHandle<T,KIn> &keys,
const vtkm::cont::ArrayHandle<U,VIn> &values,
vtkm::cont::ArrayHandle<T,KOut> &keys_output,
vtkm::cont::ArrayHandle<U,VOut> &values_output,
BinaryOperation binaryOp)
{
(void) binaryOp;
VTKM_ASSERT_CONT(keys.GetNumberOfValues() == values.GetNumberOfValues());
const vtkm::Id numberOfKeys = keys.GetNumberOfValues();
if(numberOfKeys <= 1)
{ //we only have a single key/value so that is our output
DerivedAlgorithm::Copy(keys, keys_output);
DerivedAlgorithm::Copy(values, values_output);
return;
}
//we need to determine based on the keys what is the keystate for
//each key. The states are start, middle, end of a series and the special
//state start and end of a series
vtkm::cont::ArrayHandle< vtkm::UInt8 > keystate;
{
typedef typename vtkm::cont::ArrayHandle<T,KIn>::template ExecutionTypes<DeviceAdapterTag>
::PortalConst InputPortalType;
typedef typename vtkm::cont::ArrayHandle< vtkm::UInt8 >::template ExecutionTypes<DeviceAdapterTag>
::Portal KeyStatePortalType;
InputPortalType inputPortal = keys.PrepareForInput(DeviceAdapterTag());
KeyStatePortalType keyStatePortal = keystate.PrepareForOutput(numberOfKeys,
DeviceAdapterTag());
ReduceStencilGeneration<InputPortalType> kernel(inputPortal, keyStatePortal);
DerivedAlgorithm::Schedule(kernel, numberOfKeys);
}
//next step is we need to reduce the values for each key. This is done
//by running an inclusive scan over the values array using the stencil.
//
// this inclusive scan will write out two values, the first being
// the value summed currently, the second being 0 or 1, with 1 being used
// when this is a value of a key we need to write ( END or START_AND_END)
{
typedef vtkm::cont::ArrayHandle<U,VIn> ValueHandleType;
typedef vtkm::cont::ArrayHandle< vtkm::UInt8> StencilHandleType;
typedef vtkm::cont::ArrayHandleZip<ValueHandleType,
StencilHandleType> ZipHandleType;
vtkm::cont::ArrayHandle< vtkm::UInt8 > stencil;
vtkm::cont::ArrayHandle< U > reducedValues;
ZipHandleType scanInput( values, keystate);
ZipHandleType scanOutput( reducedValues, stencil);
DerivedAlgorithm::ScanInclusive(scanInput, scanOutput, ReduceByKeyAdd() );
//at this point we are done with keystate, so free the memory
keystate.ReleaseResources();
// all we need know is an efficient way of doing the write back to the
// reduced global memory. this is done by using StreamCompact with the
// stencil and values we just created with the inclusive scan
DerivedAlgorithm::StreamCompact( reducedValues,
stencil,
values_output,
ReduceByKeyUnaryStencilOp());
} //release all temporary memory
//find all the unique keys
DerivedAlgorithm::Copy(keys,keys_output);
DerivedAlgorithm::Unique(keys_output);
}
//--------------------------------------------------------------------------
// Scan Exclusive
private:
@ -386,17 +771,20 @@ public:
//--------------------------------------------------------------------------
// Scan Inclusive
private:
template<typename PortalType>
template<typename PortalType, typename BinaryOperation>
struct ScanKernel : vtkm::exec::FunctorBase
{
PortalType Portal;
BinaryOperation BinaryOperator;
vtkm::Id Stride;
vtkm::Id Offset;
vtkm::Id Distance;
VTKM_CONT_EXPORT
ScanKernel(const PortalType &portal, vtkm::Id stride, vtkm::Id offset)
ScanKernel(const PortalType &portal, BinaryOperation binaryOp,
vtkm::Id stride, vtkm::Id offset)
: Portal(portal),
BinaryOperator(binaryOp),
Stride(stride),
Offset(offset),
Distance(stride/2)
@ -414,7 +802,7 @@ private:
{
ValueType leftValue = this->Portal.Get(leftIndex);
ValueType rightValue = this->Portal.Get(rightIndex);
this->Portal.Set(rightIndex, leftValue+rightValue);
this->Portal.Set(rightIndex, BinaryOperator(leftValue,rightValue) );
}
}
};
@ -424,32 +812,45 @@ public:
VTKM_CONT_EXPORT static T ScanInclusive(
const vtkm::cont::ArrayHandle<T,CIn> &input,
vtkm::cont::ArrayHandle<T,COut>& output)
{
return DerivedAlgorithm::ScanInclusive(input,
output,
vtkm::internal::Add());
}
template<typename T, class CIn, class COut, class BinaryOperation>
VTKM_CONT_EXPORT static T ScanInclusive(
const vtkm::cont::ArrayHandle<T,CIn> &input,
vtkm::cont::ArrayHandle<T,COut>& output,
BinaryOperation binaryOp)
{
typedef typename
vtkm::cont::ArrayHandle<T,COut>
::template ExecutionTypes<DeviceAdapterTag>::Portal PortalType;
typedef ScanKernel<PortalType,BinaryOperation> ScanKernelType;
DerivedAlgorithm::Copy(input, output);
vtkm::Id numValues = output.GetNumberOfValues();
if (numValues < 1)
{
return 0;
}
{
return output.GetPortalConstControl().Get(0);
}
PortalType portal = output.PrepareForInPlace(DeviceAdapterTag());
vtkm::Id stride;
for (stride = 2; stride-1 < numValues; stride *= 2)
{
ScanKernel<PortalType> kernel(portal, stride, stride/2 - 1);
ScanKernelType kernel(portal, binaryOp, stride, stride/2 - 1);
DerivedAlgorithm::Schedule(kernel, numValues/stride);
}
// Do reverse operation on odd indices. Start at stride we were just at.
for (stride /= 2; stride > 1; stride /= 2)
{
ScanKernel<PortalType> kernel(portal, stride, stride - 1);
ScanKernelType kernel(portal, binaryOp, stride, stride - 1);
DerivedAlgorithm::Schedule(kernel, numValues/stride);
}
@ -648,24 +1049,29 @@ public:
//--------------------------------------------------------------------------
// Stream Compact
private:
template<class StencilPortalType, class OutputPortalType>
template<class StencilPortalType,
class OutputPortalType,
class PredicateOperator>
struct StencilToIndexFlagKernel
{
typedef typename StencilPortalType::ValueType StencilValueType;
StencilPortalType StencilPortal;
OutputPortalType OutputPortal;
PredicateOperator Predicate;
VTKM_CONT_EXPORT
StencilToIndexFlagKernel(StencilPortalType stencilPortal,
OutputPortalType outputPortal)
: StencilPortal(stencilPortal), OutputPortal(outputPortal) { }
OutputPortalType outputPortal,
PredicateOperator predicate)
: StencilPortal(stencilPortal),
OutputPortal(outputPortal),
Predicate(predicate) { }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id index) const
{
StencilValueType value = this->StencilPortal.Get(index);
bool flag = ::vtkm::not_default_constructor<StencilValueType>()(value);
this->OutputPortal.Set(index, flag ? 1 : 0);
this->OutputPortal.Set(index, this->Predicate(value) ? 1 : 0);
}
VTKM_CONT_EXPORT
@ -676,38 +1082,42 @@ private:
template<class InputPortalType,
class StencilPortalType,
class IndexPortalType,
class OutputPortalType>
class OutputPortalType,
class PredicateOperator>
struct CopyIfKernel
{
InputPortalType InputPortal;
StencilPortalType StencilPortal;
IndexPortalType IndexPortal;
OutputPortalType OutputPortal;
PredicateOperator Predicate;
VTKM_CONT_EXPORT
CopyIfKernel(InputPortalType inputPortal,
StencilPortalType stencilPortal,
IndexPortalType indexPortal,
OutputPortalType outputPortal)
OutputPortalType outputPortal,
PredicateOperator predicate)
: InputPortal(inputPortal),
StencilPortal(stencilPortal),
IndexPortal(indexPortal),
OutputPortal(outputPortal) { }
OutputPortal(outputPortal),
Predicate(predicate) { }
VTKM_EXEC_EXPORT
void operator()(vtkm::Id index) const
{
typedef typename StencilPortalType::ValueType StencilValueType;
StencilValueType stencilValue = this->StencilPortal.Get(index);
if (::vtkm::not_default_constructor<StencilValueType>()(stencilValue))
{
if (Predicate(stencilValue))
{
vtkm::Id outputIndex = this->IndexPortal.Get(index);
typedef typename OutputPortalType::ValueType OutputValueType;
OutputValueType value = this->InputPortal.Get(index);
this->OutputPortal.Set(outputIndex, value);
}
}
}
VTKM_CONT_EXPORT
@ -717,11 +1127,13 @@ private:
public:
template<typename T, typename U, class CIn, class CStencil, class COut>
template<typename T, typename U, class CIn, class CStencil,
class COut, class PredicateOperator>
VTKM_CONT_EXPORT static void StreamCompact(
const vtkm::cont::ArrayHandle<T,CIn>& input,
const vtkm::cont::ArrayHandle<U,CStencil>& stencil,
vtkm::cont::ArrayHandle<T,COut>& output)
vtkm::cont::ArrayHandle<T,COut>& output,
PredicateOperator predicate)
{
VTKM_ASSERT_CONT(input.GetNumberOfValues() == stencil.GetNumberOfValues());
vtkm::Id arrayLength = stencil.GetNumberOfValues();
@ -741,9 +1153,11 @@ public:
IndexPortalType indexPortal =
indices.PrepareForOutput(arrayLength, DeviceAdapterTag());
StencilToIndexFlagKernel<
StencilPortalType, IndexPortalType> indexKernel(stencilPortal,
indexPortal);
StencilToIndexFlagKernel< StencilPortalType,
IndexPortalType,
PredicateOperator> indexKernel(stencilPortal,
indexPortal,
predicate);
DerivedAlgorithm::Schedule(indexKernel, arrayLength);
@ -763,13 +1177,25 @@ public:
InputPortalType,
StencilPortalType,
IndexPortalType,
OutputPortalType>copyKernel(inputPortal,
OutputPortalType,
PredicateOperator> copyKernel(inputPortal,
stencilPortal,
indexPortal,
outputPortal);
outputPortal,
predicate);
DerivedAlgorithm::Schedule(copyKernel, arrayLength);
}
template<typename T, typename U, class CIn, class CStencil, class COut>
VTKM_CONT_EXPORT static void StreamCompact(
const vtkm::cont::ArrayHandle<T,CIn>& input,
const vtkm::cont::ArrayHandle<U,CStencil>& stencil,
vtkm::cont::ArrayHandle<T,COut>& output)
{
::vtkm::not_default_constructor<U> predicate;
DerivedAlgorithm::StreamCompact(input, stencil, output, predicate);
}
template<typename T, class CStencil, class COut>
VTKM_CONT_EXPORT static void StreamCompact(
const vtkm::cont::ArrayHandle<T,CStencil> &stencil,

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