Shuenhoy/vtk-m2
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into add-external-faces

Browse Source
This commit is contained in:
Robert Maynard 2015-08-24 16:22:35 -04:00
commit a23125575a
64 changed files with 3096 additions and 418 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
CMake/FindBoostHeaders.cmake
View File

@ -148,6 +148,16 @@ endfunction()
#-------------------------------------------------------------------------------
# main.
#-------------------------------------------------------------------------------
#
# Act like the real boost find package by setting up the Boost version variables
# We want to keep the original Boost.cmake names so it is easier to patch this
# file, when Boost.cmake gets updated.
#
set(Boost_FIND_VERSION ${BoostHeaders_FIND_VERSION})
set(Boost_FIND_VERSION_MAJOR ${BoostHeaders_FIND_VERSION_MAJOR})
set(Boost_FIND_VERSION_MINOR ${BoostHeaders_FIND_VERSION_MINOR})
set(Boost_FIND_VERSION_PATCH ${BoostHeaders_FIND_VERSION_PATCH})
set(Boost_DEBUG ${BoostHeaders_DEBUG})
# Check the version of Boost against the requested version.
if(Boost_FIND_VERSION AND NOT Boost_FIND_VERSION_MINOR)

78
CMake/FindGLEW.cmake Normal file
View File

@ -0,0 +1,78 @@
# - Find the OpenGL Extension Wrangler Library (GLEW)
# This module defines the following variables:
# GLEW_INCLUDE_DIRS - include directories for GLEW
# GLEW_LIBRARIES - libraries to link against GLEW
# GLEW_FOUND - true if GLEW has been found and can be used
#=============================================================================
# Copyright 2012 Benjamin Eikel
#
# CMake - Cross Platform Makefile Generator
# Copyright 2000-2011 Kitware, Inc., Insight Software Consortium
# All rights reserved.
#
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions
# are met:
#
# * Redistributions of source code must retain the above copyright
# notice, this list of conditions and the following disclaimer.
#
# * Redistributions in binary form must reproduce the above copyright
# notice, this list of conditions and the following disclaimer in the
# documentation and/or other materials provided with the distribution.
#
# * Neither the names of Kitware, Inc., the Insight Software Consortium,
# nor the names of their contributors may be used to endorse or promote
# products derived from this software without specific prior written
# permission.
#
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
# "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
# LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
# A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
# HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
# SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
# LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
# DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
# THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#
# ------------------------------------------------------------------------------
#
# The above copyright and license notice applies to distributions of
# CMake in source and binary form. Some source files contain additional
# notices of original copyright by their contributors; see each source
# for details. Third-party software packages supplied with CMake under
# compatible licenses provide their own copyright notices documented in
# corresponding subdirectories.
#
# ------------------------------------------------------------------------------
#
# CMake was initially developed by Kitware with the following sponsorship:
#
# * National Library of Medicine at the National Institutes of Health
# as part of the Insight Segmentation and Registration Toolkit (ITK).
#
# * US National Labs (Los Alamos, Livermore, Sandia) ASC Parallel
# Visualization Initiative.
#
# * National Alliance for Medical Image Computing (NAMIC) is funded by the
# National Institutes of Health through the NIH Roadmap for Medical Research,
# Grant U54 EB005149.
#
# * Kitware, Inc.
find_path(GLEW_INCLUDE_DIR GL/glew.h)
find_library(GLEW_LIBRARY NAMES GLEW glew32 glew glew32s PATH_SUFFIXES lib64)
set(GLEW_INCLUDE_DIRS ${GLEW_INCLUDE_DIR})
set(GLEW_LIBRARIES ${GLEW_LIBRARY})
include(FindPackageHandleStandardArgs)
find_package_handle_standard_args(GLEW
REQUIRED_VARS GLEW_INCLUDE_DIR GLEW_LIBRARY)
mark_as_advanced(GLEW_INCLUDE_DIR GLEW_LIBRARY)

44
CMakeLists.txt
View File

@ -30,13 +30,24 @@ set(VTKm_INSTALL_INCLUDE_DIR "include")
set(VTKm_INSTALL_CONFIG_DIR "include")
set(VTKm_INSTALL_CMAKE_MODULE_DIR "share/vtkm/cmake")
set(VTKm_REQUIRED_BOOST_VERSION 1.48.0)
set(VTKm_REQUIRED_BOOST_VERSION "1.48")
# include some vtkm-specific cmake code.
include(CMake/VTKmMacros.cmake)
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${VTKm_SOURCE_DIR}/CMake)
#-----------------------------------------------------------------------------
# List of Boost features used:
# * Smart Ptr
# * Meta programming language
# We check for boost first, as the device configuration code requires boost
# to be found
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION} REQUIRED)
if(NOT Boost_FOUND)
message(FATAL_ERROR "Please specify where to find boost headers (${VTKm_REQUIRED_BOOST_VERSION}+)" )
endif()
#-----------------------------------------------------------------------------
# Check for Cxx11 support.
option(VTKm_FORCE_ANSI
@ -145,11 +156,30 @@ vtkm_install_headers(
unset(VTKM_USE_DOUBLE_PRECISION)
unset(VTKM_USE_64BIT_IDS)
#-----------------------------------------------------------------------------
# List of Boost features used:
# * Smart Ptr
# * Meta programming language
find_package(BoostHeaders ${VTKm_REQUIRED_BOOST_VERSION} REQUIRED)
# Find OpenGL and GLEW, if both are found we can enable
# the OpenGL Interop support. We use
include(CMakeDependentOption)
# enable Interop only if we have OpenGL and GLEW
find_package(OpenGL)
find_package(GLEW)
find_package(GLUT)
#dependent option reads, value to set, if condition is true, otherwise
#use last value
CMAKE_DEPENDENT_OPTION(VTKm_ENABLE_OPENGL_INTEROP
"Enable OpenGL Interop will require GLEW"
ON "OPENGL_FOUND;GLEW_FOUND" OFF)
#Only enable OpenGL Interop tests if we have Interop enabled
#and we have GLUT
#dependent option reads, value to set, if condition is true, otherwise
#use last value
CMAKE_DEPENDENT_OPTION(VTKm_ENABLE_OPENGL_TESTS
"Enable OpenGL Interop Render Window Tests"
ON "VTKm_ENABLE_OPENGL_INTEROP;GLUT_FOUND" OFF)
find_package(Pyexpander)
@ -211,6 +241,9 @@ install(FILES ${VTKm_SOURCE_DIR}/LICENSE.txt
install(
FILES
${VTKm_SOURCE_DIR}/CMake/FindBoostHeaders.cmake
${VTKm_SOURCE_DIR}/CMake/FindGLEW.cmake
${VTKm_SOURCE_DIR}/CMake/FindTBB.cmake
${VTKm_SOURCE_DIR}/CMake/FindThrust.cmake
DESTINATION ${VTKm_INSTALL_CMAKE_MODULE_DIR}
)
@ -218,6 +251,7 @@ install(
install(
FILES
${VTKm_SOURCE_DIR}/CMake/UseVTKmSerial.cmake
${VTKm_SOURCE_DIR}/CMake/UseVTKmTBB.cmake
${VTKm_SOURCE_DIR}/CMake/UseVTKmCuda.cmake
DESTINATION ${VTKm_INSTALL_CMAKE_MODULE_DIR}
)

1
LICENSE.txt
View File

@ -73,5 +73,6 @@ licenses.
CMake/CheckCXX11Features.cmake
CMake/FindBoostHeaders.cmake
CMake/FindTBB.cmake
CMake/FindGLEW.cmake
vtkm/cont/tbb/internal/parallel_sort.h
vtkm/testing/OptionParser.h

5
vtkm/CMakeLists.txt
View File

@ -35,6 +35,7 @@ set(headers
TypeTraits.h
VectorAnalysis.h
VecTraits.h
VecVariable.h
UnaryPredicates.h
)
@ -47,6 +48,10 @@ vtkm_declare_headers(${headers})
add_subdirectory(testing)
add_subdirectory(internal)
if(VTKm_ENABLE_OPENGL_INTEROP)
add_subdirectory(opengl)
endif(VTKm_ENABLE_OPENGL_INTEROP)
#-----------------------------------------------------------------------------
#add the control and exec folders
add_subdirectory(cont)

6
vtkm/Matrix.h
View File

@ -557,6 +557,12 @@ public:
typedef T ComponentType;
static const vtkm::IdComponent NUM_COMPONENTS = NumRow*NumCol;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeStatic IsSizeStatic;
VTKM_EXEC_CONT_EXPORT
static vtkm::IdComponent GetNumberOfComponents(const MatrixType &) {
return NUM_COMPONENTS;
}
VTKM_EXEC_CONT_EXPORT
static const ComponentType &GetComponent(const MatrixType &matrix,

18
vtkm/Types.h
View File

@ -735,6 +735,9 @@ struct Negate
//-----------------------------------------------------------------------------
// Pre declaration
template<typename T, vtkm::IdComponent Size> class Vec;
namespace detail {
/// Base implementation of all Vec classes.
@ -766,6 +769,21 @@ protected:
}
public:
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetNumberOfComponents() { return NUM_COMPONENTS; }
template<vtkm::IdComponent OtherSize>
VTKM_EXEC_CONT_EXPORT
void CopyInto(vtkm::Vec<ComponentType,OtherSize> &dest) const
{
for (vtkm::IdComponent index = 0;
(index < NUM_COMPONENTS) && (index < OtherSize);
index++)
{
dest[index] = (*this)[index];
}
}
VTKM_EXEC_CONT_EXPORT
DerivedClass &operator=(const DerivedClass &src)
{

64
vtkm/VecTraits.h
View File

@ -38,6 +38,15 @@ struct VecTraitsTagMultipleComponents { };
///
struct VecTraitsTagSingleComponent { };
/// A tag for vectors where the number of components are known at compile time.
///
struct VecTraitsTagSizeStatic { };
/// A tag for vectors where the number of components are not determined until
/// run time.
///
struct VecTraitsTagSizeVariable { };
namespace internal {
template<vtkm::IdComponent numComponents>
@ -65,10 +74,15 @@ struct VecTraits
///
typedef typename VecType::ComponentType ComponentType;
/// Number of components in the vector.
/// Number of components in the vector. This is only defined for vectors
/// of a static size.
///
static const vtkm::IdComponent NUM_COMPONENTS = VecType::NUM_COMPONENTS;
/// Number of components in the given vector.
///
static vtkm::IdComponent GetNumberOfComponents(const VecType &vec);
/// A tag specifying whether this vector has multiple components (i.e. is a
/// "real" vector). This tag can be useful for creating specialized functions
/// when a vector is really just a scalar.
@ -76,6 +90,14 @@ struct VecTraits
typedef typename internal::VecTraitsMultipleComponentChooser<
NUM_COMPONENTS>::Type HasMultipleComponents;
/// A tag specifying whether the size of this vector is known at compile
/// time. If set to \c VecTraitsTagSizeStatic, then \c NUM_COMPONENTS is set.
/// If set to \c VecTraitsTagSizeVariable, then the number of components is
/// not known at compile time and must be queried with \c
/// GetNumberOfComponents.
///
typedef vtkm::VecTraitsTagSizeStatic IsSizeStatic;
/// Returns the value in a given component of the vector.
///
VTKM_EXEC_CONT_EXPORT static const ComponentType &GetComponent(
@ -91,11 +113,12 @@ struct VecTraits
vtkm::IdComponent component,
ComponentType value);
/// Converts whatever type this vector is into the standard VTK-m Vec.
/// Copies the components in the given vector into a given Vec object.
///
template<vktm::IdComponent destSize>
VTKM_EXEC_CONT_EXPORT
static vtkm::Vec<ComponentType,NUM_COMPONENTS>
ToVec(const VecType &vector);
static void
CopyInto(const VecType &src, vtkm::Vec<ComponentType,destSize> &dest);
};
#else // VTKM_DOXYGEN_ONLY
;
@ -121,6 +144,13 @@ struct VecTraits<vtkm::Vec<T,Size> >
///
static const vtkm::IdComponent NUM_COMPONENTS = VecType::NUM_COMPONENTS;
/// Number of components in the given vector.
///
VTKM_EXEC_CONT_EXPORT
static vtkm::IdComponent GetNumberOfComponents(const VecType &) {
return NUM_COMPONENTS;
}
/// A tag specifying whether this vector has multiple components (i.e. is a
/// "real" vector). This tag can be useful for creating specialized functions
/// when a vector is really just a scalar.
@ -128,6 +158,14 @@ struct VecTraits<vtkm::Vec<T,Size> >
typedef typename internal::VecTraitsMultipleComponentChooser<
NUM_COMPONENTS>::Type HasMultipleComponents;
/// A tag specifying whether the size of this vector is known at compile
/// time. If set to \c VecTraitsTagSizeStatic, then \c NUM_COMPONENTS is set.
/// If set to \c VecTraitsTagSizeVariable, then the number of components is
/// not known at compile time and must be queried with \c
/// GetNumberOfComponents.
///
typedef vtkm::VecTraitsTagSizeStatic IsSizeStatic;
/// Returns the value in a given component of the vector.
///
VTKM_EXEC_CONT_EXPORT
@ -151,11 +189,12 @@ struct VecTraits<vtkm::Vec<T,Size> >
/// Converts whatever type this vector is into the standard VTKm Tuple.
///
template<vtkm::IdComponent destSize>
VTKM_EXEC_CONT_EXPORT
static vtkm::Vec<ComponentType,NUM_COMPONENTS>
ToVec(const VecType &vector)
static void
CopyInto(const VecType &src, vtkm::Vec<ComponentType,destSize> &dest)
{
return vector;
src.CopyInto(dest);
}
};
@ -167,6 +206,11 @@ struct VecTraitsBasic {
typedef ScalarType ComponentType;
static const vtkm::IdComponent NUM_COMPONENTS = 1;
typedef VecTraitsTagSingleComponent HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeStatic IsSizeStatic;
static vtkm::IdComponent GetNumberOfComponents(const ScalarType &) {
return 1;
}
VTKM_EXEC_CONT_EXPORT static const ComponentType &GetComponent(
const ScalarType &vector,
@ -184,10 +228,12 @@ struct VecTraitsBasic {
vector = value;
}
template<vtkm::IdComponent destSize>
VTKM_EXEC_CONT_EXPORT
static vtkm::Vec<ScalarType,1> ToVec(const ScalarType &vector)
static void CopyInto(const ScalarType &src,
vtkm::Vec<ScalarType,destSize> &dest)
{
return vtkm::Vec<ScalarType,1>(vector);
dest[0] = src;
}
};
} // namespace internal

167
vtkm/VecVariable.h Normal file
View File

@ -0,0 +1,167 @@
//============================================================================
// 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_VecVariable_h
#define vtk_m_VecVariable_h
#include <vtkm/Math.h>
#include <vtkm/Types.h>
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
namespace vtkm {
/// \brief A short variable-length array with maximum length.
///
/// The \c VecVariable class is a Vec-like class that holds a short array of
/// some maximum length. To avoid dynamic allocations, the maximum length is
/// specified at compile time. Internally, \c VecVariable holds a \c Vec of
/// the maximum length and exposes a subsection of it.
///
template<typename T, vtkm::IdComponent MaxSize>
class VecVariable
{
public:
typedef T ComponentType;
VTKM_EXEC_CONT_EXPORT
VecVariable() : NumComponents(0) { }
template<vtkm::IdComponent SrcSize>
VTKM_EXEC_CONT_EXPORT
VecVariable(const vtkm::VecVariable<ComponentType,SrcSize> &src)
: NumComponents(src.GetNumberOfComponents())
{
for (vtkm::IdComponent index = 0; index < this->NumComponents; index++)
{
this->Data[index] = src[index];
}
}
template<vtkm::IdComponent SrcSize>
VTKM_EXEC_CONT_EXPORT
VecVariable(const vtkm::Vec<ComponentType,SrcSize> &src)
: NumComponents(SrcSize)
{
for (vtkm::IdComponent index = 0; index < this->NumComponents; index++)
{
this->Data[index] = src[index];
}
}
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetNumberOfComponents() const {
return this->NumComponents;
}
template<vtkm::IdComponent DestSize>
VTKM_EXEC_CONT_EXPORT
void CopyInto(vtkm::Vec<ComponentType,DestSize> &dest) const
{
vtkm::IdComponent numComponents = vtkm::Min(DestSize, this->NumComponents);
for (vtkm::IdComponent index = 0; index < numComponents; index++)
{
dest[index] = this->Data[index];
}
}
VTKM_EXEC_CONT_EXPORT
const ComponentType &operator[](vtkm::IdComponent index) const
{
return this->Data[index];
}
VTKM_EXEC_CONT_EXPORT
ComponentType &operator[](vtkm::IdComponent index)
{
return this->Data[index];
}
VTKM_EXEC_CONT_EXPORT
void Append(ComponentType value)
{
this->Data[this->NumComponents] = value;
this->NumComponents++;
}
private:
vtkm::Vec<T,MaxSize> Data;
vtkm::IdComponent NumComponents;
};
template<typename T, vtkm::IdComponent MaxSize>
struct TypeTraits<vtkm::VecVariable<T,MaxSize> >
{
typedef typename vtkm::TypeTraits<T>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
VTKM_EXEC_CONT_EXPORT
static vtkm::VecVariable<T,MaxSize> ZeroInitialization()
{
return vtkm::VecVariable<T,MaxSize>();
}
};
template<typename T, vtkm::IdComponent MaxSize>
struct VecTraits<vtkm::VecVariable<T,MaxSize> >
{
typedef vtkm::VecVariable<T,MaxSize> VecType;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
VTKM_EXEC_CONT_EXPORT
static vtkm::IdComponent GetNumberOfComponents(const VecType &vector) {
return vector.GetNumberOfComponents();
}
VTKM_EXEC_CONT_EXPORT
static const ComponentType &GetComponent(const VecType &vector,
vtkm::IdComponent componentIndex)
{
return vector[componentIndex];
}
VTKM_EXEC_CONT_EXPORT
static ComponentType &GetComponent(VecType &vector,
vtkm::IdComponent componentIndex)
{
return vector[componentIndex];
}
VTKM_EXEC_CONT_EXPORT
static void SetComponent(VecType &vector,
vtkm::IdComponent componentIndex,
const ComponentType &value)
{
vector[componentIndex] = value;
}
template<vtkm::IdComponent destSize>
VTKM_EXEC_CONT_EXPORT
static void CopyInto(const VecType &src,
vtkm::Vec<ComponentType,destSize> &dest)
{
src.CopyInto(dest);
}
};
} // namespace vtkm
#endif //vtk_m_VecVariable_h

5
vtkm/cont/CellSet.h
View File

@ -28,6 +28,7 @@
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/static_assert.hpp>
#include <boost/type_traits/is_base_of.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
@ -94,8 +95,8 @@ struct CellSetCheck
typedef typename boost::is_base_of<vtkm::cont::CellSet, T>::type type;
};
#define VTKM_IS_CELL_SET(type) \
BOOST_MPL_ASSERT(( ::vtkm::cont::internal::CellSetCheck<type> ))
#define VTKM_IS_CELL_SET(T) \
BOOST_STATIC_ASSERT(::vtkm::cont::internal::CellSetCheck<T>::type::value)
} // namespace internal

4
vtkm/cont/CellSetExplicit.h
View File

@ -246,11 +246,11 @@ public:
}
VTKM_CONT_EXPORT
vtkm::cont::ArrayHandle<vtkm::Id, ShapeStorageTag> &
const vtkm::cont::ArrayHandle<vtkm::Id, ShapeStorageTag> &
GetShapesArray() const { return this->PointToCell.Shapes; }
VTKM_CONT_EXPORT
vtkm::cont::ArrayHandle<vtkm::Id, NumIndicesStorageTag> &
const vtkm::cont::ArrayHandle<vtkm::Id, NumIndicesStorageTag> &
GetNumIndicesArray() const { return this->PointToCell.NumIndices; }
VTKM_CONT_EXPORT

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

@ -298,11 +298,12 @@ private:
template<class InputPortal, class ValuesPortal, class OutputPortal,
class BinaryCompare>
VTKM_CONT_EXPORT static void LowerBoundsPortal(const InputPortal &input,
const ValuesPortal &values,
const OutputPortal &output,
BinaryCompare binary_compare)
const ValuesPortal &values,
const OutputPortal &output,
BinaryCompare binary_compare)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<bool,
typedef typename InputPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType,
BinaryCompare> bop(binary_compare);
::thrust::lower_bound(thrust::cuda::par,
IteratorBegin(input),
@ -330,7 +331,8 @@ private:
typename InputPortal::ValueType initialValue,
BinaryFunctor binary_functor)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<typename InputPortal::ValueType,
typedef typename InputPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType,
BinaryFunctor> bop(binary_functor);
return ::thrust::reduce(thrust::cuda::par,
IteratorBegin(input),
@ -361,7 +363,8 @@ private:
::thrust::equal_to<typename KeysPortal::ValueType> binaryPredicate;
vtkm::exec::cuda::internal::WrappedBinaryOperator<typename ValuesPortal::ValueType,
typedef typename ValuesPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType,
BinaryFunctor> bop(binary_functor);
result_iterators = ::thrust::reduce_by_key(thrust::cuda::par,
IteratorBegin(keys),
@ -381,7 +384,7 @@ private:
typename InputPortal::ValueType ScanExclusivePortal(const InputPortal &input,
const OutputPortal &output)
{
typedef typename InputPortal::ValueType ValueType;
typedef typename OutputPortal::ValueType ValueType;
return ScanExclusivePortal(input,
output,
@ -397,7 +400,7 @@ private:
{
// Use iterator to get value so that thrust device_ptr has chance to handle
// data on device.
typedef typename InputPortal::ValueType ValueType;
typedef typename OutputPortal::ValueType ValueType;
ValueType inputEnd = *(IteratorEnd(input) - 1);
vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType,
@ -421,7 +424,7 @@ private:
typename InputPortal::ValueType ScanInclusivePortal(const InputPortal &input,
const OutputPortal &output)
{
typedef typename InputPortal::ValueType ValueType;
typedef typename OutputPortal::ValueType ValueType;
return ScanInclusivePortal(input, output, ::thrust::plus<ValueType>() );
}
@ -431,7 +434,8 @@ private:
const OutputPortal &output,
BinaryFunctor binary_functor)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<typename InputPortal::ValueType,
typedef typename OutputPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryOperator<ValueType,
BinaryFunctor> bop(binary_functor);
typedef typename detail::IteratorTraits<OutputPortal>::IteratorType
@ -458,7 +462,9 @@ private:
VTKM_CONT_EXPORT static void SortPortal(const ValuesPortal &values,
BinaryCompare binary_compare)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<bool,BinaryCompare> bop(binary_compare);
typedef typename ValuesPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType,
BinaryCompare> bop(binary_compare);
::thrust::sort(thrust::cuda::par,
IteratorBegin(values),
IteratorEnd(values),
@ -479,7 +485,9 @@ private:
const ValuesPortal &values,
BinaryCompare binary_compare)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<bool,BinaryCompare> bop(binary_compare);
typedef typename KeysPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType,
BinaryCompare> bop(binary_compare);
::thrust::sort_by_key(thrust::cuda::par,
IteratorBegin(keys),
IteratorEnd(keys),
@ -548,7 +556,10 @@ private:
{
typedef typename detail::IteratorTraits<ValuesPortal>::IteratorType
IteratorType;
vtkm::exec::cuda::internal::WrappedBinaryOperator<bool,BinaryCompare> bop(binary_compare);
typedef typename ValuesPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType,
BinaryCompare> bop(binary_compare);
IteratorType begin = IteratorBegin(values);
IteratorType newLast = ::thrust::unique(thrust::cuda::par,
begin,
@ -579,7 +590,10 @@ private:
const OutputPortal &output,
BinaryCompare binary_compare)
{
vtkm::exec::cuda::internal::WrappedBinaryOperator<bool,BinaryCompare> bop(binary_compare);
typedef typename OutputPortal::ValueType ValueType;
vtkm::exec::cuda::internal::WrappedBinaryPredicate<ValueType,
BinaryCompare> bop(binary_compare);
::thrust::upper_bound(thrust::cuda::par,
IteratorBegin(input),
IteratorEnd(input),

19
vtkm/cont/internal/IteratorFromArrayPortal.h
View File

@ -194,9 +194,26 @@ void swap( vtkm::cont::internal::detail::IteratorFromArrayPortalValue<T> a,
a.Swap(b);
}
}
}
} // namespace vtkm::cont::internal
namespace boost {
/// The boost::iterator_facade lets you redefine the reference type, which is
/// good since you cannot set an array portal from a reference in general.
/// However, the iterator_facade then checks to see if the reference type is an
/// actual reference, and if it is not it can set up some rather restrictive
/// traits that we do not want. To get around this, specialize the
/// boost::is_reference type check to declare our value class as a reference
/// type. Even though it is not a true reference type, its operators make it
/// behave like one.
///
template<typename T>
struct is_reference<
vtkm::cont::internal::detail::IteratorFromArrayPortalValue<T> >
: public boost::true_type { };
} // namespace boost
#endif //vtk_m_cont_internal_IteratorFromArrayPortal_h

17
vtkm/cont/testing/UnitTestDataSetRegular.cxx
View File

@ -89,10 +89,9 @@ TwoDimRegularTest()
vtkm::Id cells[2][4] = {{0,1,3,4}, {1,2,4,5}};
vtkm::Vec<vtkm::Id,4> pointIds;
for (vtkm::Id cellIndex = 0; cellIndex < 2; cellIndex++)
{
pointToCell.GetIndices(cellIndex, pointIds);
vtkm::Vec<vtkm::Id,4> pointIds = pointToCell.GetIndices(cellIndex);
for (vtkm::IdComponent localPointIndex = 0;
localPointIndex < 4;
localPointIndex++)
@ -112,8 +111,8 @@ TwoDimRegularTest()
for (vtkm::Id pointIndex = 0; pointIndex < 6; pointIndex++)
{
vtkm::Vec<vtkm::Id,4> retrievedCellIds;
cellToPoint.GetIndices(pointIndex, retrievedCellIds);
vtkm::Vec<vtkm::Id,4> retrievedCellIds =
cellToPoint.GetIndices(pointIndex);
for (vtkm::IdComponent cellIndex = 0; cellIndex < 4; cellIndex++)
VTKM_TEST_ASSERT(
retrievedCellIds[cellIndex] == expectedCellIds[pointIndex][cellIndex],
@ -163,8 +162,7 @@ ThreeDimRegularTest()
vtkm::TopologyElementTagPoint(),
vtkm::TopologyElementTagCell());
vtkm::Id expectedPointIds[8] = {0,1,3,4,6,7,9,10};
vtkm::Vec<vtkm::Id,8> retrievedPointIds;
pointToCell.GetIndices(0, retrievedPointIds);
vtkm::Vec<vtkm::Id,8> retrievedPointIds = pointToCell.GetIndices(0);
for (vtkm::IdComponent localPointIndex = 0;
localPointIndex < 8;
localPointIndex++)
@ -182,11 +180,10 @@ ThreeDimRegularTest()
vtkm::cont::DeviceAdapterTagSerial(),
vtkm::TopologyElementTagCell(),
vtkm::TopologyElementTagPoint());
vtkm::Vec<vtkm::Id,8> expectedCellIds;
vtkm::Id retrievedCellIds[8] = {0,-1,-1,-1,-1,-1,-1,-1};
cellToPoint.GetIndices(0, expectedCellIds);
vtkm::Id retrievedCellIds[6] = {0,-1,-1,-1,-1,-1};
vtkm::Vec<vtkm::Id,6> expectedCellIds = cellToPoint.GetIndices(0);
for (vtkm::IdComponent localPointIndex = 0;
localPointIndex < 8;
localPointIndex < 6;
localPointIndex++)
{
VTKM_TEST_ASSERT(

1
vtkm/exec/CMakeLists.txt
View File

@ -25,7 +25,6 @@ set(headers
ExecutionObjectBase.h
ExecutionWholeArray.h
FunctorBase.h
TopologyData.h
)
#-----------------------------------------------------------------------------

47
vtkm/exec/ConnectivityExplicit.h
View File

@ -20,9 +20,10 @@
#ifndef vtk_m_exec_ConnectivityExplicit_h
#define vtk_m_exec_ConnectivityExplicit_h
#include <vtkm/CellType.h>
#include <vtkm/Types.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/exec/internal/VecFromPortal.h>
namespace vtkm {
namespace exec {
@ -51,38 +52,48 @@ public:
}
VTKM_EXEC_EXPORT
vtkm::Id GetNumberOfElements()
vtkm::Id GetNumberOfElements() const
{
return Shapes.GetNumberOfValues();
return this->Shapes.GetNumberOfValues();
}
VTKM_EXEC_EXPORT
vtkm::Id GetNumberOfIndices(vtkm::Id index)
vtkm::IdComponent GetNumberOfIndices(vtkm::Id index) const
{
return NumIndices.Get(index);
// Should the NumIndices array be typed as vtkm::IdComponent instead of
// vtkm::Id? (NumIndices is really defined in
// vtkm::cont::internal::ConnectivityExplicitInternals.)
return static_cast<vtkm::IdComponent>(this->NumIndices.Get(index));
}
VTKM_EXEC_EXPORT
vtkm::Id GetCellShape(vtkm::Id index)
vtkm::CellType GetCellShape(vtkm::Id index) const
{
return Shapes.Get(index);
// Likewise, should Shapes be vtkm::Id or something smaller?
return static_cast<vtkm::CellType>(this->Shapes.Get(index));
}
template <vtkm::IdComponent ItemTupleLength>
typedef vtkm::exec::internal::VecFromPortal<ConnectivityPortalType>
IndicesType;
/// Returns a Vec-like object containing the indices for the given index.
/// The object returned is not an actual array, but rather an object that
/// loads the indices lazily out of the connectivity array. This prevents
/// us from having to know the number of indices at compile time.
///
VTKM_EXEC_EXPORT
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
IndicesType GetIndices(vtkm::Id index) const
{
vtkm::Id n = GetNumberOfIndices(index);
vtkm::Id start = IndexOffset.Get(index);
for (vtkm::IdComponent i=0; i<n && i<ItemTupleLength; i++)
ids[i] = Connectivity.Get(start+i);
vtkm::Id offset = this->IndexOffset.Get(index);
vtkm::IdComponent length = this->GetNumberOfIndices(index);
return IndicesType(this->Connectivity, length, offset);
}
private:
ShapePortalType Shapes;
NumIndicesPortalType NumIndices;
ConnectivityPortalType Connectivity;
IndexOffsetPortalType IndexOffset;
ShapePortalType Shapes;
NumIndicesPortalType NumIndices;
ConnectivityPortalType Connectivity;
IndexOffsetPortalType IndexOffset;
};
} // namespace exec

15
vtkm/exec/ConnectivityStructured.h
View File

@ -40,6 +40,8 @@ class ConnectivityStructured
typedef vtkm::internal::ConnectivityStructuredInternals<Dimension>
InternalsType;
typedef vtkm::internal::ConnectivityStructuredIndexHelper<
FromTopology,ToTopology,Dimension> Helper;
public:
typedef typename InternalsType::SchedulingRangeType SchedulingRangeType;
@ -63,9 +65,7 @@ public:
}
VTKM_EXEC_EXPORT
vtkm::Id GetNumberOfIndices(vtkm::Id index) const {
typedef vtkm::internal::ConnectivityStructuredIndexHelper<
FromTopology,ToTopology,Dimension> Helper;
vtkm::IdComponent GetNumberOfIndices(vtkm::Id index) const {
return Helper::GetNumberOfIndices(this->Internals, index);
}
// This needs some thought. What does cell shape mean when the to topology
@ -75,13 +75,12 @@ public:
return Internals.GetCellShape();
}
template <vtkm::IdComponent ItemTupleLength>
typedef typename Helper::IndicesType IndicesType;
VTKM_EXEC_EXPORT
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
IndicesType GetIndices(vtkm::Id index) const
{
typedef vtkm::internal::ConnectivityStructuredIndexHelper<
FromTopology,ToTopology,Dimension> Helper;
Helper::GetIndices(this->Internals,index,ids);
return Helper::GetIndices(this->Internals,index);
}
private:

10
vtkm/exec/ExecutionWholeArray.h
View File

@ -46,9 +46,15 @@ public:
}
VTKM_CONT_EXPORT
ExecutionWholeArray( const vtkm::cont::ArrayHandle<T,StorageTag>& handle,
ExecutionWholeArray( vtkm::cont::ArrayHandle<T,StorageTag>& handle ):
Portal( handle.PrepareForInPlace( DeviceAdapterTag()) )
{
}
VTKM_CONT_EXPORT
ExecutionWholeArray( vtkm::cont::ArrayHandle<T,StorageTag>& handle,
vtkm::Id length ):
Portal( handle.PrepareForInPlace(length, DeviceAdapterTag()) )
Portal( handle.PrepareForOutput( length, DeviceAdapterTag()) )
{
}

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

@ -22,11 +22,8 @@
#include <vtkm/exec/arg/AspectTagDefault.h>
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/TopologyData.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/type_traits.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
#include <vtkm/exec/internal/VecFromPortalPermute.h>
namespace vtkm {
namespace exec {
@ -40,6 +37,40 @@ namespace arg {
///
struct FetchTagArrayTopologyMapIn { };
namespace detail {
// This internal class defines how a TopologyMapIn fetch loads from field data
// based on the connectivity class and the object holding the field data. The
// default implementation gets a Vec of indices and an array portal for the
// field and delivers a VecFromPortalPermute. Specializations could have more
// efficient implementations. For example, if the connectivity is structured
// and the field is regular point coordinates, it is much faster to compute the
// field directly.
template<typename ConnectivityType, typename FieldExecObjectType>
struct FetchArrayTopologyMapInImplementation
{
// The connectivity classes are expected to have an IndicesType that is
// is Vec-like class that will be returned from a GetIndices method.
typedef typename ConnectivityType::IndicesType IndexVecType;
// The FieldExecObjectType is expected to behave like an ArrayPortal.
typedef FieldExecObjectType PortalType;
typedef vtkm::exec::internal::VecFromPortalPermute<
IndexVecType,PortalType> ValueType;
VTKM_EXEC_EXPORT
static ValueType Load(vtkm::Id index,
const ConnectivityType &connectivity,
const FieldExecObjectType &field)
{
return ValueType(connectivity.GetIndices(index), field);
}
};
} // namespace detail
template<typename Invocation, vtkm::IdComponent ParameterIndex>
struct Fetch<
vtkm::exec::arg::FetchTagArrayTopologyMapIn,
@ -47,45 +78,37 @@ struct Fetch<
Invocation,
ParameterIndex>
{
// The parameter for the input domain is stored in the Invocation. (It is
// also in the worklet, but it is safer to get it from the Invocation
// in case some other dispatch operation had to modify it.)
static const vtkm::IdComponent InputDomainIndex =
Invocation::InputDomainIndex;
typedef typename Invocation::ControlInterface::template
ParameterType<InputDomainIndex>::type ControlSignatureTag;
static const vtkm::IdComponent ITEM_TUPLE_LENGTH =
ControlSignatureTag::ITEM_TUPLE_LENGTH;
// Assuming that this fetch is used in a topology map, which is its
// intention, InputDomainIndex points to a connectivity object. Thus,
// ConnectivityType is one of the vtkm::exec::Connectivity* classes.
typedef typename Invocation::ParameterInterface::
template ParameterType<InputDomainIndex>::type ConnectivityType;
// The execution object associated with this parameter is expected to be
// an array portal containing the field values.
typedef typename Invocation::ParameterInterface::
template ParameterType<ParameterIndex>::type ExecObjectType;
typedef boost::remove_const<typename ExecObjectType::ValueType> NonConstType;
typedef vtkm::exec::TopologyData<typename NonConstType::type,
ITEM_TUPLE_LENGTH> ValueType;
typedef detail::FetchArrayTopologyMapInImplementation<
ConnectivityType,ExecObjectType> Implementation;
typedef typename Implementation::ValueType ValueType;
VTKM_EXEC_EXPORT
ValueType Load(vtkm::Id index, const Invocation &invocation) const
{
typedef typename Invocation::ParameterInterface ParameterInterface;
typedef typename ParameterInterface::
template ParameterType<InputDomainIndex>::type TopologyType;
TopologyType topology =
const ConnectivityType &connectivity =
invocation.Parameters.template GetParameter<InputDomainIndex>();
const ExecObjectType &field =
invocation.Parameters.template GetParameter<ParameterIndex>();
vtkm::IdComponent nids =
static_cast<vtkm::IdComponent>(topology.GetNumberOfIndices(index));
vtkm::Vec<vtkm::Id,ITEM_TUPLE_LENGTH> ids;
topology.GetIndices(index,ids);
ValueType v;
for (vtkm::IdComponent i=0; i<nids && i<ITEM_TUPLE_LENGTH; ++i)
{
v[i] = invocation.Parameters.template GetParameter<ParameterIndex>().
Get(ids[i]);
}
return v;
return Implementation::Load(index, connectivity, field);
}
VTKM_EXEC_EXPORT

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

@ -51,7 +51,7 @@ struct FromCount : vtkm::exec::arg::ExecutionSignatureTagBase
template<typename FetchTag, typename Invocation>
struct Fetch<FetchTag, vtkm::exec::arg::AspectTagFromCount, Invocation, 1>
{
typedef vtkm::Id ValueType;
typedef vtkm::IdComponent ValueType;
VTKM_EXEC_EXPORT
ValueType Load(vtkm::Id index, const Invocation &invocation) const

39
vtkm/exec/arg/FromIndices.h
View File

@ -23,8 +23,6 @@
#include <vtkm/exec/arg/Fetch.h>
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
#include <vtkm/exec/TopologyData.h>
namespace vtkm {
namespace exec {
namespace arg {
@ -50,8 +48,11 @@ struct FromIndices : vtkm::exec::arg::ExecutionSignatureTagBase
typedef vtkm::exec::arg::AspectTagFromIndices AspectTag;
};
template<typename FetchTag, typename Invocation>
struct Fetch<FetchTag, vtkm::exec::arg::AspectTagFromIndices, Invocation, 1>
template<typename FetchTag,
typename Invocation,
vtkm::IdComponent ParameterIndex>
struct Fetch<
FetchTag, vtkm::exec::arg::AspectTagFromIndices, Invocation, ParameterIndex>
{
// The parameter for the input domain is stored in the Invocation. (It is
// also in the worklet, but it is safer to get it from the Invocation
@ -59,37 +60,23 @@ struct Fetch<FetchTag, vtkm::exec::arg::AspectTagFromIndices, Invocation, 1>
static const vtkm::IdComponent InputDomainIndex =
Invocation::InputDomainIndex;
typedef typename Invocation::ControlInterface::template
ParameterType<InputDomainIndex>::type InputDomainTag;
// Assuming that this fetch is used in a topology map, which is its
// intention, InputDomainIndex points to a connectivity object. Thus,
// ConnectivityType is one of the vtkm::exec::Connectivity* classes.
typedef typename Invocation::ParameterInterface::
template ParameterType<InputDomainIndex>::type ConnectivityType;
static const vtkm::IdComponent ITEM_TUPLE_LENGTH =
InputDomainTag::ITEM_TUPLE_LENGTH;
typedef vtkm::exec::TopologyData<vtkm::Id,ITEM_TUPLE_LENGTH> ValueType;
typedef typename ConnectivityType::IndicesType ValueType;
VTKM_EXEC_EXPORT
ValueType Load(vtkm::Id index, const Invocation &invocation) const
{
// ParameterInterface (from Invocation) is a FunctionInterface type
// containing types for all objects passed to the Invoke method (with
// some dynamic casting performed so objects like DynamicArrayHandle get
// cast to ArrayHandle).
typedef typename Invocation::ParameterInterface ParameterInterface;
// This is the type for the input domain (derived from the last two things
// we got from the Invocation).
typedef typename ParameterInterface::
template ParameterType<InputDomainIndex>::type TopologyType;
// We can pull the input domain parameter (the data specifying the input
// domain) from the invocation object.
TopologyType topology =
const ConnectivityType &connectivity =
invocation.Parameters.template GetParameter<InputDomainIndex>();
ValueType v;
topology.GetIndices(index,v.vec);
return v;
return connectivity.GetIndices(index);
}
VTKM_EXEC_EXPORT

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

@ -28,6 +28,7 @@
// Disable warnings we check vtkm for but Thrust does not.
VTKM_THIRDPARTY_PRE_INCLUDE
#include <thrust/system/cuda/memory.h>
#include <boost/type_traits/remove_const.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
namespace vtkm {
@ -39,9 +40,16 @@ namespace internal {
// wrapped operator with complex value types such as
// PortalValue which happen when passed an input array that
// is implicit.
template<typename ResultType, typename Function>
struct WrappedBinaryOperator
template<typename T_, typename Function>
struct WrappedBinaryOperator
{
typedef typename boost::remove_const<T_>::type T;
//make typedefs that thust expects binary operators to have
typedef T first_argument_type;
typedef T second_argument_type;
typedef T result_type;
Function m_f;
VTKM_EXEC_EXPORT
@ -54,61 +62,128 @@ template<typename ResultType, typename Function>
: m_f(f)
{}
template<typename T, typename U>
VTKM_EXEC_EXPORT ResultType operator()(const T &x, const U &y) const
VTKM_EXEC_EXPORT T operator()(const T &x, const T &y) const
{
return m_f(x, y);
}
template<typename T, typename U>
VTKM_EXEC_EXPORT ResultType operator()(const T &x,
const PortalValue<U> &y) const
template<typename U>
VTKM_EXEC_EXPORT T operator()(const T &x,
const PortalValue<U> &y) const
{
typedef typename PortalValue<U>::ValueType ValueType;
return m_f(x, (ValueType)y);
return m_f(x, (T)y);
}
template<typename T, typename U>
VTKM_EXEC_EXPORT ResultType operator()(const PortalValue<T> &x,
const U &y) const
template<typename U>
VTKM_EXEC_EXPORT T operator()(const PortalValue<U> &x,
const T &y) const
{
typedef typename PortalValue<T>::ValueType ValueType;
return m_f((ValueType)x, y);
return m_f((T)x, y);
}
template<typename T, typename U>
VTKM_EXEC_EXPORT ResultType operator()(const PortalValue<T> &x,
const PortalValue<U> &y) const
template<typename U, typename V>
VTKM_EXEC_EXPORT T operator()(const PortalValue<U> &x,
const PortalValue<V> &y) const
{
typedef typename PortalValue<T>::ValueType ValueTypeT;
typedef typename PortalValue<U>::ValueType ValueTypeU;
return m_f((ValueTypeT)x, (ValueTypeU)y);
return m_f((T)x, (T)y);
}
template<typename T>
VTKM_EXEC_EXPORT ResultType operator()(const thrust::system::cuda::pointer<T> x,
const T* y) const
VTKM_EXEC_EXPORT T operator()(const thrust::system::cuda::pointer<T> x,
const T* y) const
{
return m_f(*x, *y);
}
template<typename T>
VTKM_EXEC_EXPORT ResultType operator()(const thrust::system::cuda::pointer<T> x,
const T& y) const
VTKM_EXEC_EXPORT T operator()(const thrust::system::cuda::pointer<T> x,
const T& y) const
{
return m_f(*x, y);
}
template<typename T>
VTKM_EXEC_EXPORT ResultType operator()(const T& x,
const thrust::system::cuda::pointer<T> y) const
VTKM_EXEC_EXPORT T operator()(const T& x,
const thrust::system::cuda::pointer<T> y) const
{
return m_f(x, *y);
}
template<typename T>
VTKM_EXEC_EXPORT ResultType operator()(const thrust::system::cuda::pointer<T> x,
const thrust::system::cuda::pointer<T> y) const
VTKM_EXEC_EXPORT T operator()(const thrust::system::cuda::pointer<T> x,
const thrust::system::cuda::pointer<T> y) const
{
return m_f(*x, *y);
}
};
template<typename T_, typename Function>
struct WrappedBinaryPredicate
{
typedef typename boost::remove_const<T_>::type T;
//make typedefs that thust expects binary operators to have
typedef T first_argument_type;
typedef T second_argument_type;
typedef T result_type;
Function m_f;
VTKM_EXEC_EXPORT
WrappedBinaryPredicate()
: m_f()
{}
VTKM_CONT_EXPORT
WrappedBinaryPredicate(const Function &f)
: m_f(f)
{}
VTKM_EXEC_EXPORT bool operator()(const T &x, const T &y) const
{
return m_f(x, y);
}
template<typename U>
VTKM_EXEC_EXPORT bool operator()(const T &x,
const PortalValue<U> &y) const
{
return m_f(x, (T)y);
}
template<typename U>
VTKM_EXEC_EXPORT bool operator()(const PortalValue<U> &x,
const T &y) const
{
return m_f((T)x, y);
}
template<typename U, typename V>
VTKM_EXEC_EXPORT bool operator()(const PortalValue<U> &x,
const PortalValue<V> &y) const
{
return m_f((T)x, (T)y);
}
VTKM_EXEC_EXPORT bool operator()(const thrust::system::cuda::pointer<T> x,
const T* y) const
{
return m_f(*x, *y);
}
VTKM_EXEC_EXPORT bool operator()(const thrust::system::cuda::pointer<T> x,
const T& y) const
{
return m_f(*x, y);
}
VTKM_EXEC_EXPORT bool operator()(const T& x,
const thrust::system::cuda::pointer<T> y) const
{
return m_f(x, *y);
}
VTKM_EXEC_EXPORT bool operator()(const thrust::system::cuda::pointer<T> x,
const thrust::system::cuda::pointer<T> y) const
{
return m_f(*x, *y);
}
@ -120,5 +195,28 @@ template<typename ResultType, typename Function>
}
} //namespace vtkm::exec::cuda::internal
#if defined(THRUST_MAJOR_VERSION) && THRUST_MAJOR_VERSION == 1 && \
THRUST_MINOR_VERSION == 8 && THRUST_SUBMINOR_VERSION < 3
//So for thrust 1.8.0 - 1.8.2 the inclusive_scan has a bug when accumulating
//values when the binary operators states it is not commutative. For now
//we can work around this issue by stating that any BinaryOperator
//from vtkm is considered to be a commutative BinaryOperator. I have
//also moved Predicates over to WrappedBinaryPredicates so that they
//don't get marked as commutative incorrectly.
//
//You can follow the status of the thrust issue at:
// https://github.com/thrust/thrust/issues/692
namespace thrust
{
namespace detail
{
template< typename T, typename F>
struct is_commutative< vtkm::exec::cuda::internal::WrappedBinaryOperator<T, F> > :
public thrust::detail::is_arithmetic<T> { };
}
}
#endif
#endif //vtk_m_exec_cuda_internal_WrappedOperators_h

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

@ -20,6 +20,8 @@
set(headers
ErrorMessageBuffer.h
VecFromPortal.h
VecFromPortalPermute.h
WorkletInvokeFunctor.h
WorkletInvokeFunctorDetail.h
)

138
vtkm/exec/internal/VecFromPortal.h Normal file
View File

@ -0,0 +1,138 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_exec_internal_VecFromPortal_h
#define vtk_m_exec_internal_VecFromPortal_h
#include <vtkm/Math.h>
#include <vtkm/Types.h>
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
namespace vtkm {
namespace exec {
namespace internal {
/// \brief A short variable-length array from a window in an ArrayPortal.
///
/// The \c VecFromPortal class is a Vec-like class that holds an array portal
/// and exposes a small window of that portal as if it were a \c Vec.
///
template<typename PortalType>
class VecFromPortal
{
public:
typedef typename PortalType::ValueType ComponentType;
VTKM_EXEC_EXPORT
VecFromPortal() : NumComponents(0), Offset(0) { }
VTKM_EXEC_EXPORT
VecFromPortal(const PortalType &portal,
vtkm::IdComponent numComponents = 0,
vtkm::Id offset = 0)
: Portal(portal), NumComponents(numComponents), Offset(offset) { }
VTKM_EXEC_EXPORT
vtkm::IdComponent GetNumberOfComponents() const {
return this->NumComponents;
}
template<vtkm::IdComponent DestSize>
VTKM_EXEC_EXPORT
void CopyInto(vtkm::Vec<ComponentType,DestSize> &dest) const
{
vtkm::IdComponent numComponents = vtkm::Min(DestSize, this->NumComponents);
for (vtkm::IdComponent index = 0; index < numComponents; index++)
{
dest[index] = this->Portal.Get(index + this->Offset);
}
}
VTKM_EXEC_EXPORT
ComponentType operator[](vtkm::IdComponent index) const
{
return this->Portal.Get(index + this->Offset);
}
private:
PortalType Portal;
vtkm::IdComponent NumComponents;
vtkm::Id Offset;
};
}
}
} // namespace vtkm::exec::internal
// Implementations of traits classes, which by definition are in the vtkm
// namespace.
namespace vtkm {
template<typename PortalType>
struct TypeTraits<vtkm::exec::internal::VecFromPortal<PortalType> >
{
private:
typedef typename PortalType::ValueType ComponentType;
public:
typedef typename vtkm::TypeTraits<ComponentType>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
VTKM_EXEC_EXPORT
static vtkm::exec::internal::VecFromPortal<PortalType> ZeroInitialization()
{
return vtkm::exec::internal::VecFromPortal<PortalType>();
}
};
template<typename PortalType>
struct VecTraits<vtkm::exec::internal::VecFromPortal<PortalType> >
{
typedef vtkm::exec::internal::VecFromPortal<PortalType> VecType;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
VTKM_EXEC_EXPORT
static vtkm::IdComponent GetNumberOfComponents(const VecType &vector) {
return vector.GetNumberOfComponents();
}
VTKM_EXEC_EXPORT
static ComponentType GetComponent(const VecType &vector,
vtkm::IdComponent componentIndex)
{
return vector[componentIndex];
}
template<vtkm::IdComponent destSize>
VTKM_EXEC_EXPORT
static void CopyInto(const VecType &src,
vtkm::Vec<ComponentType,destSize> &dest)
{
src.CopyInto(dest);
}
};
} // namespace vtkm
#endif //vtk_m_exec_internal_VecFromPortal_h

142
vtkm/exec/internal/VecFromPortalPermute.h Normal file
View File

@ -0,0 +1,142 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_exec_internal_VecFromPortalPermute_h
#define vtk_m_exec_internal_VecFromPortalPermute_h
#include <vtkm/Math.h>
#include <vtkm/Types.h>
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
namespace vtkm {
namespace exec {
namespace internal {
/// \brief A short vector from an ArrayPortal and a vector of indices.
///
/// The \c VecFromPortalPermute class is a Vec-like class that holds an array
/// portal and a second Vec-like containing indices into the array. Each value
/// of this vector is the value from the array with the respective index.
///
template<typename IndexVecType, typename PortalType>
class VecFromPortalPermute
{
public:
typedef typename PortalType::ValueType ComponentType;
VTKM_EXEC_EXPORT
VecFromPortalPermute() { }
VTKM_EXEC_EXPORT
VecFromPortalPermute(const IndexVecType &indices, const PortalType &portal)
: Indices(indices), Portal(portal) { }
VTKM_EXEC_EXPORT
vtkm::IdComponent GetNumberOfComponents() const {
return this->Indices.GetNumberOfComponents();
}
template<vtkm::IdComponent DestSize>
VTKM_EXEC_EXPORT
void CopyInto(vtkm::Vec<ComponentType,DestSize> &dest) const
{
vtkm::IdComponent numComponents =
vtkm::Min(DestSize, this->GetNumberOfComponents());
for (vtkm::IdComponent index = 0; index < numComponents; index++)
{
dest[index] = this->Portal.Get(this->Indices[index]);
}
}
VTKM_EXEC_EXPORT
ComponentType operator[](vtkm::IdComponent index) const
{
return this->Portal.Get(this->Indices[index]);
}
private:
IndexVecType Indices;
PortalType Portal;
};
}
}
} // namespace vtkm::exec::internal
// Implementations of traits classes, which by definition are in the vtkm
// namespace.
namespace vtkm {
template<typename IndexVecType, typename PortalType>
struct TypeTraits<
vtkm::exec::internal::VecFromPortalPermute<IndexVecType,PortalType> >
{
private:
typedef vtkm::exec::internal::VecFromPortalPermute<IndexVecType,PortalType>
VecType;
typedef typename PortalType::ValueType ComponentType;
public:
typedef typename vtkm::TypeTraits<ComponentType>::NumericTag NumericTag;
typedef TypeTraitsVectorTag DimensionalityTag;
VTKM_EXEC_EXPORT
static VecType ZeroInitialization()
{
return VecType();
}
};
template<typename IndexVecType, typename PortalType>
struct VecTraits<
vtkm::exec::internal::VecFromPortalPermute<IndexVecType,PortalType> >
{
typedef vtkm::exec::internal::VecFromPortalPermute<IndexVecType,PortalType>
VecType;
typedef typename VecType::ComponentType ComponentType;
typedef vtkm::VecTraitsTagMultipleComponents HasMultipleComponents;
typedef vtkm::VecTraitsTagSizeVariable IsSizeStatic;
VTKM_EXEC_EXPORT
static vtkm::IdComponent GetNumberOfComponents(const VecType &vector) {
return vector.GetNumberOfComponents();
}
VTKM_EXEC_EXPORT
static ComponentType GetComponent(const VecType &vector,
vtkm::IdComponent componentIndex)
{
return vector[componentIndex];
}
template<vtkm::IdComponent destSize>
VTKM_EXEC_EXPORT
static void CopyInto(const VecType &src,
vtkm::Vec<ComponentType,destSize> &dest)
{
src.CopyInto(dest);
}
};
} // namespace vtkm
#endif //vtk_m_exec_internal_VecFromPortalPermute_h

2
vtkm/exec/internal/testing/CMakeLists.txt
View File

@ -22,6 +22,8 @@
set(unit_tests
UnitTestErrorMessageBuffer.cxx
UnitTestVecFromPortal.cxx
UnitTestVecFromPortalPermute.cxx
UnitTestWorkletInvokeFunctor.cxx
)
vtkm_unit_tests(SOURCES ${unit_tests})

120
vtkm/exec/internal/testing/UnitTestVecFromPortal.cxx Normal file
View File

@ -0,0 +1,120 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/exec/internal/VecFromPortal.h>
#include <vtkm/testing/Testing.h>
namespace UnitTestVecFromPortalNamespace {
static const vtkm::IdComponent ARRAY_SIZE = 10;
template<typename T>
void CheckType(T, T)
{
// Check passes if this function is called correctly.
}
template<typename T>
class TestPortal
{
public:
typedef T ValueType;
VTKM_EXEC_EXPORT
vtkm::Id GetNumberOfValues() const { return ARRAY_SIZE; }
VTKM_EXEC_EXPORT
ValueType Get(vtkm::Id index) const { return TestValue(index, ValueType()); }
};
struct VecFromPortalTestFunctor
{
template<typename T>
void operator()(T) const
{
typedef TestPortal<T> PortalType;
typedef vtkm::exec::internal::VecFromPortal<PortalType> VecType;
typedef vtkm::TypeTraits<VecType> TTraits;
typedef vtkm::VecTraits<VecType> VTraits;
std::cout << "Checking VecFromPortal traits" << std::endl;
// The statements will fail to compile if the traits is not working as
// expected.
CheckType(typename TTraits::DimensionalityTag(),
vtkm::TypeTraitsVectorTag());
CheckType(typename VTraits::ComponentType(), T());
CheckType(typename VTraits::HasMultipleComponents(),
vtkm::VecTraitsTagMultipleComponents());
CheckType(typename VTraits::IsSizeStatic(),
vtkm::VecTraitsTagSizeVariable());
std::cout << "Checking VecFromPortal contents" << std::endl;
PortalType portal;
for (vtkm::Id offset = 0; offset < ARRAY_SIZE; offset++)
{
for (vtkm::IdComponent length = 0; length < ARRAY_SIZE-offset; length++)
{
VecType vec(portal, length, offset);
VTKM_TEST_ASSERT(vec.GetNumberOfComponents() == length,
"Wrong length.");
VTKM_TEST_ASSERT(VTraits::GetNumberOfComponents(vec) == length,
"Wrong length.");
vtkm::Vec<T,ARRAY_SIZE> copyDirect;
vec.CopyInto(copyDirect);
vtkm::Vec<T,ARRAY_SIZE> copyTraits;
VTraits::CopyInto(vec, copyTraits);
for (vtkm::IdComponent index = 0; index < length; index++)
{
T expected = TestValue(index+offset, T());
VTKM_TEST_ASSERT(test_equal(vec[index], expected), "Wrong value.");
VTKM_TEST_ASSERT(
test_equal(VTraits::GetComponent(vec, index), expected),
"Wrong value.");
VTKM_TEST_ASSERT(test_equal(copyDirect[index], expected),
"Wrong copied value.");
VTKM_TEST_ASSERT(test_equal(copyTraits[index], expected),
"Wrong copied value.");
}
}
}
}
};
void VecFromPortalTest()
{
vtkm::testing::Testing::TryTypes(VecFromPortalTestFunctor(),
vtkm::TypeListTagCommon());
}
} // namespace UnitTestVecFromPortalNamespace
int UnitTestVecFromPortal(int, char *[])
{
return vtkm::testing::Testing::Run(
UnitTestVecFromPortalNamespace::VecFromPortalTest);
}

132
vtkm/exec/internal/testing/UnitTestVecFromPortalPermute.cxx Normal file
View File

@ -0,0 +1,132 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/exec/internal/VecFromPortalPermute.h>
#include <vtkm/VecVariable.h>
#include <vtkm/testing/Testing.h>
namespace UnitTestVecFromPortalPermuteNamespace {
static const vtkm::IdComponent ARRAY_SIZE = 10;
template<typename T>
void CheckType(T, T)
{
// Check passes if this function is called correctly.
}
template<typename T>
class TestPortal
{
public:
typedef T ValueType;
VTKM_EXEC_EXPORT
vtkm::Id GetNumberOfValues() const { return ARRAY_SIZE; }
VTKM_EXEC_EXPORT
ValueType Get(vtkm::Id index) const { return TestValue(index, ValueType()); }
};
struct VecFromPortalPermuteTestFunctor
{
template<typename T>
void operator()(T) const
{
typedef TestPortal<T> PortalType;
typedef vtkm::VecVariable<vtkm::Id,ARRAY_SIZE> IndexVecType;
typedef vtkm::exec::internal::VecFromPortalPermute<IndexVecType,PortalType>
VecType;
typedef vtkm::TypeTraits<VecType> TTraits;
typedef vtkm::VecTraits<VecType> VTraits;
std::cout << "Checking VecFromPortal traits" << std::endl;
// The statements will fail to compile if the traits is not working as
// expected.
CheckType(typename TTraits::DimensionalityTag(),
vtkm::TypeTraitsVectorTag());
CheckType(typename VTraits::ComponentType(), T());
CheckType(typename VTraits::HasMultipleComponents(),
vtkm::VecTraitsTagMultipleComponents());
CheckType(typename VTraits::IsSizeStatic(),
vtkm::VecTraitsTagSizeVariable());
std::cout << "Checking VecFromPortal contents" << std::endl;
PortalType portal;
for (vtkm::Id offset = 0; offset < ARRAY_SIZE; offset++)
{
for (vtkm::IdComponent length = 0;
2*length + offset < ARRAY_SIZE;
length++)
{
IndexVecType indices;
for (vtkm::IdComponent index = 0; index < length; index++)
{
indices.Append(offset + 2*index);
}
VecType vec(indices, portal);
VTKM_TEST_ASSERT(vec.GetNumberOfComponents() == length,
"Wrong length.");
VTKM_TEST_ASSERT(VTraits::GetNumberOfComponents(vec) == length,
"Wrong length.");
vtkm::Vec<T,ARRAY_SIZE> copyDirect;
vec.CopyInto(copyDirect);
vtkm::Vec<T,ARRAY_SIZE> copyTraits;
VTraits::CopyInto(vec, copyTraits);
for (vtkm::IdComponent index = 0; index < length; index++)
{
T expected = TestValue(2*index+offset, T());
VTKM_TEST_ASSERT(test_equal(vec[index], expected), "Wrong value.");
VTKM_TEST_ASSERT(
test_equal(VTraits::GetComponent(vec, index), expected),
"Wrong value.");
VTKM_TEST_ASSERT(test_equal(copyDirect[index], expected),
"Wrong copied value.");
VTKM_TEST_ASSERT(test_equal(copyTraits[index], expected),
"Wrong copied value.");
}
}
}
}
};
void VecFromPortalPermuteTest()
{
vtkm::testing::Testing::TryTypes(VecFromPortalPermuteTestFunctor(),
vtkm::TypeListTagCommon());
}
} // namespace UnitTestVecFromPortalPermuteNamespace
int UnitTestVecFromPortalPermute(int, char *[])
{
return vtkm::testing::Testing::Run(
UnitTestVecFromPortalPermuteNamespace::VecFromPortalPermuteTest);
}

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

@ -27,7 +27,7 @@
#include <vtkm/testing/Testing.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <boost/mpl/assert.hpp>
#include <boost/static_assert.hpp>
#include <boost/type_traits/is_same.hpp>
VTKM_THIRDPARTY_POST_INCLUDE
@ -166,17 +166,17 @@ struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
// Check behavior of InvocationToFetch helper class.
BOOST_MPL_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType1,1>::type,
vtkm::exec::arg::Fetch<TestFetchTagInput,vtkm::exec::arg::AspectTagDefault,InvocationType1,1> > ));
BOOST_STATIC_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType1,1>::type,
vtkm::exec::arg::Fetch<TestFetchTagInput,vtkm::exec::arg::AspectTagDefault,InvocationType1,1> >::type::value ));
BOOST_MPL_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType1,2>::type,
vtkm::exec::arg::Fetch<TestFetchTagOutput,vtkm::exec::arg::AspectTagDefault,InvocationType1,2> > ));
BOOST_STATIC_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType1,2>::type,
vtkm::exec::arg::Fetch<TestFetchTagOutput,vtkm::exec::arg::AspectTagDefault,InvocationType1,2> >::type::value ));
BOOST_MPL_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType2,0>::type,
vtkm::exec::arg::Fetch<TestFetchTagOutput,vtkm::exec::arg::AspectTagDefault,InvocationType2,2> > ));
BOOST_STATIC_ASSERT(( boost::is_same<
vtkm::exec::internal::detail::InvocationToFetch<InvocationType2,0>::type,
vtkm::exec::arg::Fetch<TestFetchTagOutput,vtkm::exec::arg::AspectTagDefault,InvocationType2,2> >::type::value ));
void TestDoWorkletInvoke()
{

5
vtkm/internal/Configure.h.in
View File

@ -131,6 +131,11 @@
#define VTKM_THIRDPARTY_POST_INCLUDE
#endif
//Mark if we are building with interop enabled
#ifndef VTKm_ENABLE_OPENGL_INTEROP
#cmakedefine VTKm_ENABLE_OPENGL_INTEROP
#endif
// Determine whether we will use variadic templates (a new feature in C++11).
// Currently have VARIADIC_TEMPLATE support off.
#cmakedefine VTKM_NO_VARIADIC_TEMPLATE

184
vtkm/internal/ConnectivityStructuredInternals.h
View File

@ -73,39 +73,52 @@ public:
return this->GetNumberOfPoints();
}
static const vtkm::IdComponent NUM_POINTS_IN_CELL = 2;
static const vtkm::IdComponent MAX_CELL_TO_POINT = 2;
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetNumberOfPoints() const {return this->PointDimensions;}
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetNumberOfCells() const {return this->PointDimensions-1;}
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetNumberOfPointsInCell() const {return 2;}
vtkm::IdComponent GetNumberOfPointsInCell() const {return NUM_POINTS_IN_CELL;}
VTKM_EXEC_CONT_EXPORT
vtkm::CellType GetCellShape() const {return VTKM_LINE;}
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetPointsOfCell(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> GetPointsOfCell(vtkm::Id index) const
{
BOOST_STATIC_ASSERT(IdsLength >= 2);
ids[0] = index;
ids[1] = ids[0] + 1;
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> pointIds;
pointIds[0] = index;
pointIds[1] = pointIds[0] + 1;
return pointIds;
}
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetCellsOfPoint(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::IdComponent GetNumberOfCellsIncidentOnPoint(vtkm::Id pointIndex) const
{
BOOST_STATIC_ASSERT(IdsLength >= 2);
ids[0] = ids[1] = -1;
return
(static_cast<vtkm::IdComponent>(pointIndex > 0)
+ static_cast<vtkm::IdComponent>(pointIndex < this->PointDimensions-1));
}
VTKM_EXEC_CONT_EXPORT
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> GetCellsOfPoint(vtkm::Id index) const
{
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> cellIds;
cellIds[0] = cellIds[1] = -1;
vtkm::IdComponent idx = 0;
if (index > 0)
{
ids[idx++] = index-1;
cellIds[idx++] = index-1;
}
if (index < this->PointDimensions-1)
{
ids[idx++] = index;
cellIds[idx++] = index;
}
return cellIds;
}
VTKM_CONT_EXPORT
@ -158,57 +171,73 @@ public:
return this->GetPointDimensions();
}
static const vtkm::IdComponent NUM_POINTS_IN_CELL = 4;
static const vtkm::IdComponent MAX_CELL_TO_POINT = 4;
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetNumberOfCells() const
{
return vtkm::internal::VecProduct<2>()(this->GetCellDimensions());
}
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetNumberOfPointsInCell() const { return 4; }
vtkm::IdComponent GetNumberOfPointsInCell() const {return NUM_POINTS_IN_CELL;}
VTKM_EXEC_CONT_EXPORT
vtkm::CellType GetCellShape() const { return VTKM_PIXEL; }
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetPointsOfCell(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> GetPointsOfCell(vtkm::Id index) const
{
BOOST_STATIC_ASSERT(IdsLength >= 4);
vtkm::Id i, j;
this->CalculateLogicalPointIndices(index, i, j);
ids[0] = j*this->PointDimensions[0] + i;
ids[1] = ids[0] + 1;
ids[2] = ids[0] + this->PointDimensions[0];
ids[3] = ids[2] + 1;
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> pointIds;
pointIds[0] = j*this->PointDimensions[0] + i;
pointIds[1] = pointIds[0] + 1;
pointIds[2] = pointIds[0] + this->PointDimensions[0];
pointIds[3] = pointIds[2] + 1;
return pointIds;
}
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetCellsOfPoint(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::IdComponent GetNumberOfCellsIncidentOnPoint(vtkm::Id pointIndex) const
{
BOOST_STATIC_ASSERT(IdsLength >= 4);
vtkm::Id i, j;
this->CalculateLogicalPointIndices(pointIndex, i, j);
return
(static_cast<vtkm::IdComponent>((i > 0) && (j > 0))
+ static_cast<vtkm::IdComponent>((i < this->PointDimensions[0]-1) && (j > 0))
+ static_cast<vtkm::IdComponent>((i > 0) && (j < this->PointDimensions[1]-1))
+ static_cast<vtkm::IdComponent>(
(i < this->PointDimensions[0]-1) && (j < this->PointDimensions[1]-1)));
}
ids[0] = ids[1] = ids[2] = ids[3] = -1;
VTKM_EXEC_CONT_EXPORT
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> GetCellsOfPoint(vtkm::Id index) const
{
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> cellIds;
cellIds[0] = cellIds[1] = cellIds[2] = cellIds[3] = -1;
vtkm::Id i, j;
vtkm::IdComponent idx = 0;
CalculateLogicalPointIndices(index, i, j);
this->CalculateLogicalPointIndices(index, i, j);
if ((i > 0) && (j > 0))
{
ids[idx++] = this->CalculateCellIndex(i-1, j-1);
cellIds[idx++] = this->CalculateCellIndex(i-1, j-1);
}
if ((i < this->PointDimensions[0]-1) && (j > 0))
{
ids[idx++] = this->CalculateCellIndex(i , j-1);
cellIds[idx++] = this->CalculateCellIndex(i , j-1);
}
if ((i > 0) && (j < this->PointDimensions[1]-1))
{
ids[idx++] = this->CalculateCellIndex(i-1, j );
cellIds[idx++] = this->CalculateCellIndex(i-1, j );
}
if ((i < this->PointDimensions[0]-1) && (j < this->PointDimensions[1]-1))
{
ids[idx++] = this->CalculateCellIndex(i , j );
cellIds[idx++] = this->CalculateCellIndex(i , j );
}
return cellIds;
}
VTKM_CONT_EXPORT
@ -279,22 +308,22 @@ public:
return this->GetPointDimensions();
}
static const vtkm::IdComponent NUM_POINTS_IN_CELL = 8;
static const vtkm::IdComponent MAX_CELL_TO_POINT = 6;
VTKM_EXEC_CONT_EXPORT
vtkm::Id GetNumberOfCells() const
{
return vtkm::internal::VecProduct<3>()(this->GetCellDimensions());
}
VTKM_EXEC_CONT_EXPORT
vtkm::IdComponent GetNumberOfPointsInCell() const { return 8; }
vtkm::IdComponent GetNumberOfPointsInCell() const {return NUM_POINTS_IN_CELL;}
VTKM_EXEC_CONT_EXPORT
vtkm::CellType GetCellShape() const { return VTKM_VOXEL; }
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetPointsOfCell(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> GetPointsOfCell(vtkm::Id index) const
{
BOOST_STATIC_ASSERT(IdsLength >= 8);
vtkm::Id3 cellDimensions = this->GetCellDimensions();
vtkm::Id cellDims01 = cellDimensions[0] * cellDimensions[1];
@ -303,23 +332,25 @@ public:
vtkm::Id j = indexij / cellDimensions[0];
vtkm::Id i = indexij % cellDimensions[0];
ids[0] = (k * this->PointDimensions[1] + j) * this->PointDimensions[0] + i;
ids[1] = ids[0] + 1;
ids[2] = ids[0] + this->PointDimensions[0];
ids[3] = ids[2] + 1;
ids[4] = ids[0] + this->PointDimensions[0]*this->PointDimensions[1];
ids[5] = ids[4] + 1;
ids[6] = ids[4] + this->PointDimensions[0];
ids[7] = ids[6] + 1;
vtkm::Vec<vtkm::Id,NUM_POINTS_IN_CELL> pointIds;
pointIds[0] = (k * this->PointDimensions[1] + j) * this->PointDimensions[0] + i;
pointIds[1] = pointIds[0] + 1;
pointIds[2] = pointIds[0] + this->PointDimensions[0];
pointIds[3] = pointIds[2] + 1;
pointIds[4] = pointIds[0] + this->PointDimensions[0]*this->PointDimensions[1];
pointIds[5] = pointIds[4] + 1;
pointIds[6] = pointIds[4] + this->PointDimensions[0];
pointIds[7] = pointIds[6] + 1;
return pointIds;
}
template <vtkm::IdComponent IdsLength>
VTKM_EXEC_CONT_EXPORT
void GetCellsOfPoint(vtkm::Id index, vtkm::Vec<vtkm::Id,IdsLength> &ids) const
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> GetCellsOfPoint(vtkm::Id index) const
{
BOOST_STATIC_ASSERT(IdsLength >= 8);
vtkm::Vec<vtkm::Id,MAX_CELL_TO_POINT> cellIds;
ids[0]=ids[1]=ids[2]=ids[3]=ids[4]=ids[5]=ids[6]=ids[7]=-1;
cellIds[0]=cellIds[1]=cellIds[2]=cellIds[3]=cellIds[4]=cellIds[5]=-1;
vtkm::Id i, j, k;
vtkm::IdComponent idx=0;
@ -327,45 +358,47 @@ public:
this->CalculateLogicalPointIndices(index, i, j, k);
if ((i > 0) && (j > 0) && (k > 0))
{
ids[idx++] = this->CalculateCellIndex(i-1, j-1, k-1);
cellIds[idx++] = this->CalculateCellIndex(i-1, j-1, k-1);
}
if ((i < this->PointDimensions[0]-1) && (j > 0) && (k > 0))
{
ids[idx++] = this->CalculateCellIndex(i , j-1, k-1);
cellIds[idx++] = this->CalculateCellIndex(i , j-1, k-1);
}
if ((i > 0) && (j < this->PointDimensions[1]-1) && (k > 0))
{
ids[idx++] = this->CalculateCellIndex(i-1, j , k-1);
cellIds[idx++] = this->CalculateCellIndex(i-1, j , k-1);
}
if ((i < this->PointDimensions[0]-1) &&
(j < this->PointDimensions[1]-1) &&
(k > 0))
{
ids[idx++] = this->CalculateCellIndex(i , j , k-1);
cellIds[idx++] = this->CalculateCellIndex(i , j , k-1);
}
if ((i > 0) && (j > 0) && (k < this->PointDimensions[2]-1))
{
ids[idx++] = this->CalculateCellIndex(i-1, j-1, k);
cellIds[idx++] = this->CalculateCellIndex(i-1, j-1, k);
}
if ((i < this->PointDimensions[0]-1) &&
(j > 0) &&
(k < this->PointDimensions[2]-1))
{
ids[idx++] = this->CalculateCellIndex(i , j-1, k);
cellIds[idx++] = this->CalculateCellIndex(i , j-1, k);
}
if ((i > 0) &&
(j < this->PointDimensions[1]-1) &&
(k < this->PointDimensions[2]-1))
{
ids[idx++] = this->CalculateCellIndex(i-1, j , k);
cellIds[idx++] = this->CalculateCellIndex(i-1, j , k);
}
if ((i < this->PointDimensions[0]-1) &&
(j < this->PointDimensions[1]-1) &&
(k < this->PointDimensions[2]-1))
{
ids[idx++] = this->CalculateCellIndex(i , j , k);
cellIds[idx++] = this->CalculateCellIndex(i , j , k);
}
return cellIds;
}
VTKM_CONT_EXPORT
@ -417,19 +450,21 @@ template<vtkm::IdComponent Dimension>
struct ConnectivityStructuredIndexHelper<
vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, Dimension>
{
template <vtkm::IdComponent ItemTupleLength>
typedef vtkm::internal::ConnectivityStructuredInternals<Dimension>
ConnectivityType;
typedef vtkm::Vec<vtkm::Id,ConnectivityType::NUM_POINTS_IN_CELL> IndicesType;
VTKM_EXEC_CONT_EXPORT
static void GetIndices(
const vtkm::internal::ConnectivityStructuredInternals<Dimension> &connectivity,
vtkm::Id cellIndex,
vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
static IndicesType GetIndices(const ConnectivityType &connectivity,
vtkm::Id cellIndex)
{
connectivity.GetPointsOfCell(cellIndex, ids);
return connectivity.GetPointsOfCell(cellIndex);
}
VTKM_EXEC_CONT_EXPORT
static vtkm::IdComponent GetNumberOfIndices(
const vtkm::internal::ConnectivityStructuredInternals<Dimension> &connectivity,
const ConnectivityType &connectivity,
vtkm::Id vtkmNotUsed(cellIndex))
{
return connectivity.GetNumberOfPointsInCell();
@ -440,18 +475,27 @@ template<vtkm::IdComponent Dimension>
struct ConnectivityStructuredIndexHelper<
vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, Dimension>
{
template <vtkm::IdComponent ItemTupleLength>
typedef vtkm::internal::ConnectivityStructuredInternals<Dimension>
ConnectivityType;
// TODO: This needs to change to a Vec-like that supports a max size.
// Likewise, all the GetCellsOfPoint methods need to use it as well.
typedef vtkm::Vec<vtkm::Id,ConnectivityType::MAX_CELL_TO_POINT> IndicesType;
VTKM_EXEC_CONT_EXPORT
static void GetIndices(
const vtkm::internal::ConnectivityStructuredInternals<Dimension> &connectivity,
vtkm::Id pointIndex,
vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
static IndicesType GetIndices(const ConnectivityType &connectivity,
vtkm::Id pointIndex)
{
connectivity.GetCellsOfPoint(pointIndex,ids);
return connectivity.GetCellsOfPoint(pointIndex);
}
// TODO: Implement GetNumberOfIndices, which will rely on a
// GetNumberOfCellsOnPoint method in ConnectivityStructuredInternals
VTKM_EXEC_CONT_EXPORT
static vtkm::IdComponent GetNumberOfIndices(
const ConnectivityType &connectivity,
vtkm::Id pointIndex)
{
return connectivity.GetNumberOfCellsIncidentOnPoint(pointIndex);
}
};
}

33
vtkm/opengl/CMakeLists.txt Normal file
View File

@ -0,0 +1,33 @@
##============================================================================
## Copyright (c) Kitware, Inc.
## All rights reserved.
## See LICENSE.txt for details.
## This software is distributed WITHOUT ANY WARRANTY; without even
## the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
## PURPOSE. See the above copyright notice for more information.
##
## Copyright 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.
##============================================================================
set(headers
TransferToOpenGL.h
)
#-----------------------------------------------------------------------------
add_subdirectory(internal)
#-----------------------------------------------------------------------------
vtkm_declare_headers(${headers})
if(VTKm_ENABLE_OPENGL_TESTS)
add_subdirectory(testing)
endif()

73
vtkm/opengl/TransferToOpenGL.h Normal file
View File

@ -0,0 +1,73 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_TransferToOpenGL_h
#define vtk_m_opengl_TransferToOpenGL_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/opengl/internal/TransferToOpenGL.h>
namespace vtkm{
namespace opengl{
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible. Will return the type of array buffer
/// that we have bound the handle too. Will be GL_ELEMENT_ARRAY_BUFFER for
/// vtkm::Id, and GL_ARRAY_BUFFER for everything else.
///
/// This function keeps the buffer as the active buffer of the returned type.
///
/// This function will throw exceptions if the transfer wasn't possible
///
template<typename ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
GLenum TransferToOpenGL(vtkm::cont::ArrayHandle<ValueType,StorageTag> handle,
GLuint& openGLHandle,
DeviceAdapterTag)
{
vtkm::opengl::internal::TransferToOpenGL<ValueType, DeviceAdapterTag> toGL;
toGL.Transfer(handle,openGLHandle);
return toGL.GetType();
}
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible. Will use the given \p type as how
/// to bind the buffer.
///
/// This function keeps the buffer as the active buffer of the input type.
///
/// This function will throw exceptions if the transfer wasn't possible
///
template<typename ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void TransferToOpenGL(vtkm::cont::ArrayHandle<ValueType, StorageTag> handle,
GLuint& openGLHandle,
GLenum type,
DeviceAdapterTag)
{
vtkm::opengl::internal::TransferToOpenGL<ValueType, DeviceAdapterTag> toGL(type);
toGL.Transfer(handle,openGLHandle);
}
}}
#endif //vtk_m_opengl_TransferToOpenGL_h

34
vtkm/opengl/cuda/CMakeLists.txt Normal file
View File

@ -0,0 +1,34 @@
##============================================================================
## 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.
##============================================================================
set(headers
SetOpenGLDevice.h
)
vtkm_disable_troublesome_thrust_warnings()
#-----------------------------------------------------------------------------
vtkm_declare_headers(CUDA ${headers})
add_subdirectory(internal)
if(VTKm_ENABLE_OPENGL_TESTS)
add_subdirectory(testing)
endif()

55
vtkm/opengl/cuda/SetOpenGLDevice.h Normal file
View File

@ -0,0 +1,55 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cuda_opengl_SetOpenGLDevice_h
#define vtk_m_cuda_opengl_SetOpenGLDevice_h
#include <cuda.h>
#include <cuda_gl_interop.h>
#include <vtkm/cont/ErrorExecution.h>
namespace vtkm{
namespace opengl{
namespace cuda{
static void SetCudaGLDevice(int id)
{
//With Cuda 5.0 cudaGLSetGLDevice is deprecated and shouldn't be needed
//anymore. But it seems that macs still require you to call it or we
//segfault
#ifdef __APPLE__
cudaError_t cError = cudaGLSetGLDevice(id);
#else
cudaError_t cError = cudaSetDevice(id);
#endif
if(cError != cudaSuccess)
{
std::string cuda_error_msg("Unable to setup cuda/opengl interop. Error: ");
cuda_error_msg.append(cudaGetErrorString(cError));
throw vtkm::cont::ErrorExecution(cuda_error_msg);
}
}
}
}
} //namespace
#endif //vtk_m_cuda_opengl_SetOpenGLDevice_h

29
vtkm/opengl/cuda/internal/CMakeLists.txt Normal file
View File

@ -0,0 +1,29 @@
##============================================================================
## 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.
##============================================================================
set(headers
TransferToOpenGL.h
)
vtkm_disable_troublesome_thrust_warnings()
#-----------------------------------------------------------------------------
vtkm_declare_headers(CUDA ${headers})

149
vtkm/opengl/cuda/internal/TransferToOpenGL.h Normal file
View File

@ -0,0 +1,149 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtkm_opengl_cuda_internal_TransferToOpenGL_h
#define vtkm_opengl_cuda_internal_TransferToOpenGL_h
#include <vtkm/cont/ErrorExecution.h>
#include <vtkm/cont/ErrorControlOutOfMemory.h>
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
#include <vtkm/cont/cuda/internal/MakeThrustIterator.h>
#include <vtkm/opengl/internal/TransferToOpenGL.h>
#include <thrust/copy.h>
#include <thrust/device_ptr.h>
namespace vtkm {
namespace opengl {
namespace internal {
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible.
///
template<typename ValueType>
class TransferToOpenGL<ValueType, vtkm::cont::DeviceAdapterTagCuda>
{
typedef vtkm::cont::DeviceAdapterTagCuda DeviceAdapterTag;
public:
VTKM_CONT_EXPORT TransferToOpenGL():
Type( vtkm::opengl::internal::BufferTypePicker( ValueType() ) )
{}
VTKM_CONT_EXPORT explicit TransferToOpenGL(GLenum type):
Type(type)
{}
GLenum GetType() const { return this->Type; }
template< typename StorageTag >
VTKM_CONT_EXPORT
void Transfer (
vtkm::cont::ArrayHandle<ValueType, StorageTag> &handle,
GLuint& openGLHandle ) const
{
//construct a cuda resource handle
cudaGraphicsResource_t cudaResource;
cudaError_t cError;
//make a buffer for the handle if the user has forgotten too
if(!glIsBuffer(openGLHandle))
{
glGenBuffers(1,&openGLHandle);
}
//bind the buffer to the given buffer type
glBindBuffer(this->Type, openGLHandle);
//Allocate the memory and set it as GL_DYNAMIC_DRAW draw
const std::size_t size = sizeof(ValueType)* handle.GetNumberOfValues();
glBufferData(this->Type, size, 0, GL_DYNAMIC_DRAW);
//register the buffer as being used by cuda
cError = cudaGraphicsGLRegisterBuffer(&cudaResource,
openGLHandle,
cudaGraphicsMapFlagsWriteDiscard);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Could not register the OpenGL buffer handle to CUDA.");
}
//map the resource into cuda, so we can copy it
cError =cudaGraphicsMapResources(1,&cudaResource);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorControlOutOfMemory(
"Could not allocate enough memory in CUDA for OpenGL interop.");
}
//get the mapped pointer
std::size_t cuda_size;
ValueType* beginPointer=NULL;
cError = cudaGraphicsResourceGetMappedPointer((void **)&beginPointer,
&cuda_size,
cudaResource);
if(cError != cudaSuccess)
{
throw vtkm::cont::ErrorExecution(
"Unable to get pointers to CUDA memory for OpenGL buffer.");
}
//assert that cuda_size is the same size as the buffer we created in OpenGL
VTKM_ASSERT_CONT(cuda_size == size);
//get the device pointers
typedef vtkm::cont::ArrayHandle<ValueType, StorageTag> HandleType;
typedef typename HandleType::template
ExecutionTypes<DeviceAdapterTag>::PortalConst PortalType;
PortalType portal = handle.PrepareForInput(DeviceAdapterTag());
//Copy the data into memory that opengl owns, since we can't
//give memory from cuda to opengl
//Perhaps a direct call to thrust copy should be wrapped in a vtkm
//compatble function
::thrust::copy(vtkm::cont::cuda::internal::IteratorBegin(portal),
vtkm::cont::cuda::internal::IteratorEnd(portal),
thrust::cuda::pointer<ValueType>(beginPointer));
//unmap the resource
cudaGraphicsUnmapResources(1, &cudaResource);
//unregister the buffer
cudaGraphicsUnregisterResource(cudaResource);
}
private:
GLenum Type;
};
}
}
} //namespace vtkm::opengl::cuda::internal
#endif

27
vtkm/opengl/cuda/testing/CMakeLists.txt Normal file
View File

@ -0,0 +1,27 @@
##============================================================================
## 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.
##============================================================================
set(unit_tests
UnitTestTransferToOpenGLCuda.cu
)
vtkm_disable_troublesome_thrust_warnings()
vtkm_unit_tests(CUDA SOURCES ${unit_tests})

36
vtkm/opengl/cuda/testing/UnitTestTransferToOpenGLCuda.cu Normal file
View File

@ -0,0 +1,36 @@
//============================================================================
// 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.
//============================================================================
//silence boost threading warnings when using cuda
#define BOOST_SP_DISABLE_THREADS
//This sets up testing with the cuda device adapter
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/cuda/internal/testing/Testing.h>
#include <vtkm/opengl/testing/TestingOpenGLInterop.h>
int UnitTestTransferToOpenGLCuda(int, char *[])
{
int result = 1;
result = vtkm::opengl::testing::TestingOpenGLInterop
<vtkm::cont::cuda::DeviceAdapterTagCuda >::Run();
return vtkm::cont::cuda::internal::Testing::CheckCudaBeforeExit(result);
}

58
vtkm/opengl/internal/BufferTypePicker.h Normal file
View File

@ -0,0 +1,58 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_internal_BufferTypePicker_h
#define vtk_m_opengl_internal_BufferTypePicker_h
#include <vtkm/Types.h>
#include <vtkm/opengl/internal/OpenGLHeaders.h>
namespace vtkm {
namespace opengl {
namespace internal {
/// helper function that guesses what OpenGL buffer type is the best default
/// given a primitive type. Currently GL_ELEMENT_ARRAY_BUFFER is used for integer
/// types, and GL_ARRAY_BUFFER is used for everything else
VTKM_CONT_EXPORT GLenum BufferTypePicker( int )
{ return GL_ELEMENT_ARRAY_BUFFER; }
VTKM_CONT_EXPORT GLenum BufferTypePicker( unsigned int )
{ return GL_ELEMENT_ARRAY_BUFFER; }
#if VTKM_SIZE_LONG == 8
VTKM_CONT_EXPORT GLenum BufferTypePicker( vtkm::Int64 )
{ return GL_ELEMENT_ARRAY_BUFFER; }
VTKM_CONT_EXPORT GLenum BufferTypePicker( vtkm::UInt64 )
{ return GL_ELEMENT_ARRAY_BUFFER; }
#endif
template<typename T>
VTKM_CONT_EXPORT GLenum BufferTypePicker( T )
{ return GL_ARRAY_BUFFER; }
}
}
} //namespace vtkm::opengl::internal
#endif //vtk_m_opengl_internal_BufferTypePicker_h

33
vtkm/opengl/internal/CMakeLists.txt Normal file
View File

@ -0,0 +1,33 @@
##============================================================================
## Copyright (c) Kitware, Inc.
## All rights reserved.
## See LICENSE.txt for details.
## This software is distributed WITHOUT ANY WARRANTY; without even
## the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
## PURPOSE. See the above copyright notice for more information.
##
## Copyright 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.
##============================================================================
set(headers
BufferTypePicker.h
OpenGLHeaders.h
TransferToOpenGL.h
)
#-----------------------------------------------------------------------------
vtkm_declare_headers(${headers})
#we currently don't have to check if we have glut for these tests
if(VTKm_ENABLE_OPENGL_TESTS)
add_subdirectory(testing)
endif()

41
vtkm/opengl/internal/OpenGLHeaders.h Normal file
View File

@ -0,0 +1,41 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_internal_OpenGLHeaders_h
#define vtk_m_opengl_internal_OpenGLHeaders_h
#include <vtkm/internal/ExportMacros.h>
#if defined(__APPLE__)
# include <GL/glew.h>
# include <OpenGL/gl.h>
#else
# include <GL/glew.h>
# include <GL/gl.h>
#endif
#ifdef VTKM_CUDA
# include <cuda_runtime.h>
# include <cuda_gl_interop.h>
#endif
#endif //vtk_m_opengl_internal_OpenGLHeaders_h

165
vtkm/opengl/internal/TransferToOpenGL.h Normal file
View File

@ -0,0 +1,165 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_internal_TransferToOpenGL_h
#define vtk_m_opengl_internal_TransferToOpenGL_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageBasic.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/opengl/internal/OpenGLHeaders.h>
#include <vtkm/opengl/internal/BufferTypePicker.h>
namespace vtkm {
namespace opengl {
namespace internal {
namespace detail
{
template<class ValueType, class StorageTag, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void CopyFromHandle(
vtkm::cont::ArrayHandle<ValueType, StorageTag>& handle,
GLenum type,
DeviceAdapterTag)
{
//Generic implementation that will work no matter what. We copy the data
//in the given handle to a temporary handle using the basic storage tag.
//We then ensure the data is available in the control environment by
//synchronizing the control array. Last, we steal the array and pass the
//iterator to the rendering system
const vtkm::Id numberOfValues = handle.GetNumberOfValues();
std::size_t size = sizeof(ValueType) * numberOfValues;
//Copy the data from its specialized Storage container to a basic storage
vtkm::cont::ArrayHandle<ValueType, vtkm::cont::StorageTagBasic> tmpHandle;
vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag>::Copy(handle, tmpHandle);
//Synchronize the arrays to ensure the most current data is available in the
//control environment
tmpHandle.SyncControlArray();
//Note that the temporary ArrayHandle is no longer valid after this call
ValueType* temporaryStorage = tmpHandle.Internals->ControlArray.StealArray();
//Detach the current buffer
glBufferData(type, size, 0, GL_DYNAMIC_DRAW);
//Allocate the memory and set it as static draw and copy into opengl
glBufferSubData(type,0,size,temporaryStorage);
delete[] temporaryStorage;
}
template<class ValueType, class DeviceAdapterTag>
VTKM_CONT_EXPORT
void CopyFromHandle(
vtkm::cont::ArrayHandle<ValueType, vtkm::cont::StorageTagBasic>& handle,
GLenum type,
DeviceAdapterTag)
{
//Specialization given that we are use an C allocated array storage tag
//that allows us to directly hook into the data. We pull the data
//back to the control enviornment using PerpareForInput and give an iterator
//from the portal to OpenGL to upload to the rendering system
//This also works because we know that this class isn't used for cuda interop,
//instead we are specialized
std::size_t size = sizeof(ValueType) * handle.GetNumberOfValues();
//Detach the current buffer
glBufferData(type, size, 0, GL_DYNAMIC_DRAW);
//Allocate the memory and set it as static draw and copy into opengl
glBufferSubData(type,0,size,
vtkm::cont::ArrayPortalToIteratorBegin(
handle.PrepareForInput(DeviceAdapterTag())
));
}
} //namespace detail
/// \brief Manages transferring an ArrayHandle to opengl .
///
/// \c TransferToOpenGL manages to transfer the contents of an ArrayHandle
/// to OpenGL as efficiently as possible.
///
template<typename ValueType, class DeviceAdapterTag>
class TransferToOpenGL
{
public:
VTKM_CONT_EXPORT TransferToOpenGL():
Type( vtkm::opengl::internal::BufferTypePicker( ValueType() ) )
{}
VTKM_CONT_EXPORT explicit TransferToOpenGL(GLenum type):
Type(type)
{}
VTKM_CONT_EXPORT GLenum GetType() const { return this->Type; }
template< typename StorageTag >
VTKM_CONT_EXPORT
void Transfer (
vtkm::cont::ArrayHandle<ValueType, StorageTag>& handle,
GLuint& openGLHandle ) const
{
//make a buffer for the handle if the user has forgotten too
if(!glIsBuffer(openGLHandle))
{
glGenBuffers(1,&openGLHandle);
}
//bind the buffer to the given buffer type
glBindBuffer(this->Type, openGLHandle);
//transfer the data.
//the primary concern that we have at this point is data locality and
//the type of storage. Our options include using CopyInto and provide a
//temporary location for the values to reside before we give it to openGL
//this works for all storage types.
//
//Second option is to call PrepareForInput and get a PortalConst in the
//exection environment.
//if we are StorageTagBasic this would allow us the ability to grab
//the raw memory value and copy those, which we know are valid and remove
//a unneeded copy.
//
//The end result is that we have CopyFromHandle which does number two
//from StorageTagBasic, and does the CopyInto for everything else
detail::CopyFromHandle(handle, this->Type, DeviceAdapterTag());
}
private:
GLenum Type;
};
}
}
} //namespace vtkm::opengl::internal
//-----------------------------------------------------------------------------
// These includes are intentionally placed here after the declaration of the
// TransferToOpenGL class, so that people get the correct device adapter
/// specializations if they exist.
#if VTKM_DEVICE_ADAPTER == VTKM_DEVICE_ADAPTER_CUDA
#include <vtkm/opengl/cuda/internal/TransferToOpenGL.h>
#endif
#endif //vtk_m_opengl_internal_TransferToOpenGL_h

26
vtkm/opengl/internal/testing/CMakeLists.txt Normal file
View File

@ -0,0 +1,26 @@
##============================================================================
## 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.
##============================================================================
set(unit_tests
UnitTestBufferTypePicker.cxx
UnitTestOpenGLHeaders.cxx
)
vtkm_unit_tests(SOURCES ${unit_tests})

56
vtkm/opengl/internal/testing/UnitTestBufferTypePicker.cxx Normal file
View File

@ -0,0 +1,56 @@
//============================================================================
// 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.
//============================================================================
#include <vtkm/opengl/internal/BufferTypePicker.h>
#include <vtkm/cont/testing/Testing.h>
namespace
{
void TestBufferTypePicker()
{
//just verify that certain types match
GLenum type;
typedef unsigned int vtkmUint;
typedef vtkm::FloatDefault T;
type = vtkm::opengl::internal::BufferTypePicker(vtkm::Id());
VTKM_TEST_ASSERT(type == GL_ELEMENT_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(int());
VTKM_TEST_ASSERT(type == GL_ELEMENT_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(vtkmUint());
VTKM_TEST_ASSERT(type == GL_ELEMENT_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(vtkm::Vec<T,4>());
VTKM_TEST_ASSERT(type == GL_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(vtkm::Vec<T,3>());
VTKM_TEST_ASSERT(type == GL_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(vtkm::FloatDefault());
VTKM_TEST_ASSERT(type == GL_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(float());
VTKM_TEST_ASSERT(type == GL_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
type = vtkm::opengl::internal::BufferTypePicker(double());
VTKM_TEST_ASSERT(type == GL_ARRAY_BUFFER, "Bad OpenGL Buffer Type");
}
}
int UnitTestBufferTypePicker(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestBufferTypePicker);
}

43
vtkm/opengl/internal/testing/UnitTestOpenGLHeaders.cxx Normal file
View File

@ -0,0 +1,43 @@
//============================================================================
// 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.
//============================================================================
#include <vtkm/opengl/internal/OpenGLHeaders.h>
#include <vtkm/cont/testing/Testing.h>
namespace {
void TestOpenGLHeaders()
{
#if defined(GL_VERSION_1_3) && (GL_VERSION_1_3 == 1)
//this is pretty simple, we just verify that certain symbols exist
//and the version of openGL is high enough that our interop will work.
GLenum e = GL_ELEMENT_ARRAY_BUFFER;
GLuint u = 1;
u = u * e;
(void) u;
#else
unable_to_find_required_gl_version();
#endif
}
}
int UnitTestOpenGLHeaders(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestOpenGLHeaders);
}

33
vtkm/opengl/testing/CMakeLists.txt Normal file
View File

@ -0,0 +1,33 @@
##============================================================================
## Copyright (c) Kitware, Inc.
## All rights reserved.
## See LICENSE.txt for details.
## This software is distributed WITHOUT ANY WARRANTY; without even
## the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
## PURPOSE. See the above copyright notice for more information.
##
## Copyright 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.
##============================================================================
set(headers
TestingOpenGLInterop.h
TestingWindow.h
WindowBase.h
)
vtkm_declare_headers(${headers})
set(unit_tests
UnitTestTransferToOpenGL.cxx
)
vtkm_unit_tests(SOURCES ${unit_tests})

317
vtkm/opengl/testing/TestingOpenGLInterop.h Normal file
View File

@ -0,0 +1,317 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_testing_TestingOpenGLInterop_h
#define vtk_m_opengl_testing_TestingOpenGLInterop_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/Magnitude.h>
#include <vtkm/opengl/testing/TestingWindow.h>
#include <vtkm/opengl/TransferToOpenGL.h>
#include <vtkm/cont/testing/Testing.h>
// #include <vtkm/cont/testing/TestingGridGenerator.h>
#include <algorithm>
#include <iterator>
#include <vector>
namespace vtkm {
namespace opengl {
namespace testing {
/// This class has a single static member, Run, that tests the templated
/// DeviceAdapter for support for opengl interop.
///
template< class DeviceAdapterTag,
class StorageTag = VTKM_DEFAULT_STORAGE_TAG>
struct TestingOpenGLInterop
{
private:
//fill the array with a collection of values and return it wrapped in
//an vtkm array handle
template<typename T>
static
vtkm::cont::ArrayHandle<T,StorageTag>
FillArray(std::vector<T>& data, std::size_t length)
{
typedef typename std::vector<T>::iterator iterator;
//make sure the data array is exactly the right length
data.clear();
data.resize(length);
vtkm::Id pos = 0;
for(iterator i = data.begin(); i != data.end(); ++i, ++pos)
{ *i=T(pos); }
std::random_shuffle(data.begin(),data.end());
return vtkm::cont::make_ArrayHandle(data);
}
//Transfer the data in a vtkm ArrayHandle to open gl while making sure
//we don't throw any errors
template<typename ArrayHandleType>
static
void SafelyTransferArray(ArrayHandleType array, GLuint& handle)
{
try
{
vtkm::opengl::TransferToOpenGL(array,handle, DeviceAdapterTag());
}
catch (vtkm::cont::ErrorControlOutOfMemory error)
{
std::cout << error.GetMessage() << std::endl;
VTKM_TEST_ASSERT(true==false,
"Got an unexpected Out Of Memory error transferring to openGL");
}
catch (vtkm::cont::ErrorControlBadValue bvError)
{
std::cout << bvError.GetMessage() << std::endl;
VTKM_TEST_ASSERT(true==false,
"Got an unexpected Bad Value error transferring to openGL");
}
}
template<typename ArrayHandleType>
static
void SafelyTransferArray(ArrayHandleType array, GLuint& handle, GLenum type)
{
try
{
vtkm::opengl::TransferToOpenGL(array,handle,type, DeviceAdapterTag());
}
catch (vtkm::cont::ErrorControlOutOfMemory error)
{
std::cout << error.GetMessage() << std::endl;
VTKM_TEST_ASSERT(true==false,
"Got an unexpected Out Of Memory error transferring to openGL");
}
catch (vtkm::cont::ErrorControlBadValue bvError)
{
std::cout << bvError.GetMessage() << std::endl;
VTKM_TEST_ASSERT(true==false,
"Got an unexpected Bad Value error transferring to openGL");
}
}
//bring the data back from openGL and into a std vector. Will bind the
//passed in handle to the default buffer type for the type T
template<typename T>
static
std::vector<T> CopyGLBuffer(GLuint& handle, T t)
{
//get the type we used for this buffer.
GLenum type = vtkm::opengl::internal::BufferTypePicker(t);
//bind the buffer to the guessed buffer type, this way
//we can call CopyGLBuffer no matter what it the active buffer
glBindBuffer(type, handle);
//get the size of the buffer
int bytesInBuffer = 0;
glGetBufferParameteriv(type, GL_BUFFER_SIZE, &bytesInBuffer);
int size = ( bytesInBuffer / sizeof(T) );
//get the buffer contents and place it into a vector
std::vector<T> data;
data.resize(size);
glGetBufferSubData(type,0,bytesInBuffer,&data[0]);
return data;
}
//make a random value that we can test when loading constant values
template<typename T>
static
T MakeRandomValue(T)
{
return T(rand());
}
struct TransferFunctor
{
// std::size_t Size;
// GLuint GLHandle;
template <typename T>
void operator()(const T t) const
{
//this->Size = 10;
std::size_t Size = 10;
GLuint GLHandle;
//verify that T is able to be transfer to openGL.
//than pull down the results from the array buffer and verify
//that they match the handles contents
std::vector<T> tempData;
vtkm::cont::ArrayHandle<T,StorageTag> temp =
FillArray(tempData,Size);
//verify that the signature that doesn't have type works
SafelyTransferArray(temp,GLHandle);
bool is_buffer;
is_buffer = glIsBuffer(GLHandle);
VTKM_TEST_ASSERT(is_buffer==true,
"OpenGL buffer not filled");
std::vector<T> returnedValues = CopyGLBuffer(GLHandle, t);
//verify the results match what is in the array handle
temp.SyncControlArray();
T* expectedValues = temp.Internals->ControlArray.StealArray();
for(std::size_t i=0; i < Size; ++i)
{
VTKM_TEST_ASSERT(test_equal(*(expectedValues+i),returnedValues[i]),
"Array Handle failed to transfer properly");
}
temp.ReleaseResources();
temp = FillArray(tempData,Size*2);
GLenum type = vtkm::opengl::internal::BufferTypePicker(t);
SafelyTransferArray(temp,GLHandle,type);
is_buffer = glIsBuffer(GLHandle);
VTKM_TEST_ASSERT(is_buffer==true,
"OpenGL buffer not filled");
returnedValues = CopyGLBuffer(GLHandle, t);
//verify the results match what is in the array handle
temp.SyncControlArray();
expectedValues = temp.Internals->ControlArray.StealArray();
for(std::size_t i=0; i < Size*2; ++i)
{
VTKM_TEST_ASSERT(test_equal(*(expectedValues+i),returnedValues[i]),
"Array Handle failed to transfer properly");
}
//verify this work for a constant value array handle
T constantValue = MakeRandomValue(t);
vtkm::cont::ArrayHandleConstant<T> constant(constantValue, Size);
SafelyTransferArray(constant,GLHandle);
is_buffer = glIsBuffer(GLHandle);
VTKM_TEST_ASSERT(is_buffer==true,
"OpenGL buffer not filled");
returnedValues = CopyGLBuffer(GLHandle, constantValue);
for(std::size_t i=0; i < Size; ++i)
{
VTKM_TEST_ASSERT(test_equal(returnedValues[i],constantValue),
"Constant value array failed to transfer properly");
}
}
};
// struct TransferGridFunctor
// {
// GLuint CoordGLHandle;
// GLuint MagnitudeGLHandle;
// template <typename GridType>
// void operator()(const GridType)
// {
// //verify we are able to be transfer both coordinates and indices to openGL.
// //than pull down the results from the array buffer and verify
// //that they match the handles contents
// vtkm::cont::testing::TestGrid<GridType,
// StorageTag,
// DeviceAdapterTag> grid(64);
// vtkm::cont::ArrayHandle<vtkm::FloatDefault,
// StorageTag,
// DeviceAdapterTag> magnitudeHandle;
// vtkm::cont::DispatcherMapField< vtkm::worklet::Magnitude,
// DeviceAdapterTag> dispatcher;
// dispatcher.Invoke(grid->GetPointCoordinates(), magnitudeHandle);
// //transfer to openGL 3 handles and catch any errors
// //
// SafelyTransferArray(grid->GetPointCoordinates(),this->CoordGLHandle);
// SafelyTransferArray(magnitudeHandle,this->MagnitudeGLHandle);
// //verify all 3 handles are actually handles
// bool is_buffer = glIsBuffer(this->CoordGLHandle);
// VTKM_TEST_ASSERT(is_buffer==true,
// "Coordinates OpenGL buffer not filled");
// is_buffer = glIsBuffer(this->MagnitudeGLHandle);
// VTKM_TEST_ASSERT(is_buffer==true,
// "Magnitude OpenGL buffer not filled");
// //now that everything is openGL we have one task left.
// //transfer everything back to the host and compare it to the
// //computed values.
// std::vector<vtkm::Vec<vtkm::FloatDefault,3>> GLReturnedCoords = CopyGLBuffer(
// this->CoordGLHandle, vtkm::Vec<vtkm::FloatDefault,3>());
// std::vector<vtkm::FloatDefault> GLReturneMags = CopyGLBuffer(
// this->MagnitudeGLHandle,vtkm::FloatDefault());
// for (vtkm::Id pointIndex = 0;
// pointIndex < grid->GetNumberOfPoints();
// pointIndex++)
// {
// vtkm::Vec<vtkm::FloatDefault,3> pointCoordinateExpected = grid.GetPointCoordinates(
// pointIndex);
// vtkm::Vec<vtkm::FloatDefault,3> pointCoordinatesReturned = GLReturnedCoords[pointIndex];
// VTKM_TEST_ASSERT(test_equal(pointCoordinateExpected,
// pointCoordinatesReturned),
// "Got bad coordinate from OpenGL buffer.");
// vtkm::FloatDefault magnitudeValue = GLReturneMags[pointIndex];
// vtkm::FloatDefault magnitudeExpected =
// sqrt(vtkm::dot(pointCoordinateExpected, pointCoordinateExpected));
// VTKM_TEST_ASSERT(test_equal(magnitudeValue, magnitudeExpected),
// "Got bad magnitude from OpenGL buffer.");
// }
// }
// };
public:
VTKM_CONT_EXPORT static int Run()
{
//create a valid openGL context that we can test transfer of data
vtkm::opengl::testing::TestingWindow window;
window.Init("Testing Window", 300, 300);
//verify that we can transfer basic arrays and constant value arrays to opengl
vtkm::testing::Testing::TryAllTypes(TransferFunctor());
//verify that openGL interop works with all grid types in that we can
//transfer coordinates / verts and properties to openGL
// vtkm::cont::testing::GridTesting::TryAllGridTypes(
// TransferGridFunctor(),
// vtkm::testing::Testing::CellCheckAlwaysTrue(),
// StorageTag(),
// DeviceAdapterTag() );
return 0;
}
};
} } }
#endif //vtk_m_opengl_testing_TestingOpenGLInterop_h

82
vtkm/opengl/testing/TestingWindow.h Normal file
View File

@ -0,0 +1,82 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_opengl_testing_TestingWindow_h
#define vtk_m_opengl_testing_TestingWindow_h
#include <vtkm/internal/ExportMacros.h>
#include <vtkm/opengl/testing/WindowBase.h>
namespace vtkm{
namespace opengl{
namespace testing{
/// \brief Basic Render Window that only makes sure opengl has a valid context
///
/// Bare-bones class that fulfullis the requirements of WindowBase but
/// has no ability to interact with opengl other than to close down the window
///
///
class TestingWindow : public vtkm::opengl::testing::WindowBase<TestingWindow>
{
public:
VTKM_CONT_EXPORT TestingWindow(){};
//called after opengl is inited
VTKM_CONT_EXPORT void PostInit()
{}
VTKM_CONT_EXPORT void Display()
{}
VTKM_CONT_EXPORT void Idle()
{}
VTKM_CONT_EXPORT void ChangeSize(int vtkmNotUsed(w), int vtkmNotUsed(h) )
{}
VTKM_CONT_EXPORT void Key(unsigned char key,
int vtkmNotUsed(x), int vtkmNotUsed(y) )
{
if(key == 27) //escape pressed
{
exit(0);
}
}
VTKM_CONT_EXPORT void SpecialKey(int vtkmNotUsed(key),
int vtkmNotUsed(x), int vtkmNotUsed(y) )
{}
VTKM_CONT_EXPORT void Mouse(int vtkmNotUsed(button), int vtkmNotUsed(state),
int vtkmNotUsed(x), int vtkmNotUsed(y) )
{}
VTKM_CONT_EXPORT void MouseMove(int vtkmNotUsed(x), int vtkmNotUsed(y) )
{}
VTKM_CONT_EXPORT void PassiveMouseMove(int vtkmNotUsed(x), int vtkmNotUsed(y) )
{}
};
}
}
}
#endif //vtk_m_opengl_testing_TestingWindow_h

30
vtkm/opengl/testing/UnitTestTransferToOpenGL.cxx Normal file
View File

@ -0,0 +1,30 @@
//============================================================================
// 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.
//============================================================================
//This sets up testing with the default device adapter and array container
#include <vtkm/cont/DeviceAdapterSerial.h>
#include <vtkm/opengl/testing/TestingOpenGLInterop.h>
int UnitTestTransferToOpenGL(int, char *[])
{
return vtkm::opengl::testing::TestingOpenGLInterop<
vtkm::cont::DeviceAdapterTagSerial >::Run();
}

160
vtkm/opengl/testing/WindowBase.h Normal file
View File

@ -0,0 +1,160 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2014 Sandia Corporation.
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtkm_m_opengl_testing_WindowBase_h
#define vtkm_m_opengl_testing_WindowBase_h
//constructs a valid openGL context so that we can verify
//that vtkm to open gl bindings work
#include <string>
// OpenGL Graphics includes
//glew needs to go before glut
#include <vtkm/opengl/internal/OpenGLHeaders.h>
#if defined (__APPLE__)
# include <GLUT/glut.h>
#else
# include <GL/glut.h>
#endif
#include <vtkm/cont/ErrorControlBadValue.h>
#ifdef VTKM_CUDA
# include <vtkm/cont/cuda/ChooseCudaDevice.h>
# include <vtkm/opengl/cuda/SetOpenGLDevice.h>
#endif
#include <iostream>
namespace vtkm{
namespace opengl{
namespace testing{
namespace internal
{
template <typename T>
struct GLUTStaticCallbackHolder
{ static T* StaticGLUTResource; };
template <typename T>
T* GLUTStaticCallbackHolder<T>::StaticGLUTResource;
}
/// \brief Basic GLUT Wrapper class
///
/// This class gives the ability to wrap the glut function callbacks into
/// a single class so that you can use c++ objects. The only downside
/// is that you can only have a single window created
///
template< class Derived >
class WindowBase : private internal::GLUTStaticCallbackHolder<Derived>
{
public:
void Init(std::string title, int width, int height,
int argc, char** argv)
{
//set our selves as the static instance to call
WindowBase<Derived>::StaticGLUTResource = static_cast<Derived*>(this);
glutInit(&argc,argv);
glutInitDisplayMode( GLUT_RGBA | GLUT_DOUBLE | GLUT_DEPTH);
glutInitWindowPosition(0,0);
glutInitWindowSize(width,height);
glutCreateWindow(title.c_str());
// glewExperimental = GL_TRUE;
glewInit();
if(!glewIsSupported("GL_VERSION_2_1"))
{
std::cerr << glGetString(GL_RENDERER) << std::endl;
std::cerr << glGetString(GL_VERSION) << std::endl;
throw vtkm::cont::ErrorControlBadValue(
"Unable to create an OpenGL 2.1 Context");
}
#ifdef VTKM_CUDA
int id = vtkm::cuda::cont::FindFastestDeviceId();
vtkm::opengl::cuda::SetCudaGLDevice(id);
#endif
//attach all the glut call backs
glutDisplayFunc( WindowBase<Derived>::GLUTDisplayCallback );
glutIdleFunc( WindowBase<Derived>::GLUTIdleCallback );
glutReshapeFunc( WindowBase<Derived>::GLUTChangeSizeCallback );
glutKeyboardFunc( WindowBase<Derived>::GLUTKeyCallback );
glutSpecialFunc( WindowBase<Derived>::GLUTSpecialKeyCallback );
glutMouseFunc( WindowBase<Derived>::GLUTMouseCallback );
glutMotionFunc( WindowBase<Derived>::GLUTMouseMoveCallback );
glutPassiveMotionFunc( WindowBase<Derived>::GLUTPassiveMouseMoveCallback );
//call any custom init code you want to have
WindowBase<Derived>::StaticGLUTResource->PostInit();
}
void Init(std::string title, int width, int height)
{
int argc=0;
char** argv = 0;
Init(title,width,height,argc,argv);
}
//Init must be called before you call Start so that we have a valid
//opengl context
void Start()
{
glutMainLoop();
}
static void GLUTDisplayCallback()
{ WindowBase<Derived>::StaticGLUTResource->Display(); }
static void GLUTIdleCallback()
{ WindowBase<Derived>::StaticGLUTResource->Idle(); }
static void GLUTChangeSizeCallback(int width, int height)
{ WindowBase<Derived>::StaticGLUTResource->ChangeSize(width,height); }
static void GLUTKeyCallback(unsigned char key, int x, int y)
{ WindowBase<Derived>::StaticGLUTResource->Key(key,x,y); }
static void GLUTSpecialKeyCallback(int key, int x, int y)
{ WindowBase<Derived>::StaticGLUTResource->SpecialKey(key,x,y); }
static void GLUTMouseCallback(int button, int state ,int x, int y)
{ WindowBase<Derived>::StaticGLUTResource->Mouse(button,state,x,y); }
static void GLUTMouseMoveCallback(int x, int y)
{ WindowBase<Derived>::StaticGLUTResource->MouseMove(x,y); }
static void GLUTPassiveMouseMoveCallback(int x, int y)
{ WindowBase<Derived>::StaticGLUTResource->PassiveMouseMove(x,y); }
};
}
}
}
#endif //vtkm_m_opengl_testing_WindowBase_h

1
vtkm/testing/CMakeLists.txt
View File

@ -42,6 +42,7 @@ set(unit_tests
UnitTestUnaryPredicates.cxx
UnitTestVectorAnalysis.cxx
UnitTestVecTraits.cxx
UnitTestVecVariable.cxx
)
VTKM_unit_tests(SOURCES ${unit_tests})

9
vtkm/testing/Testing.h
View File

@ -277,10 +277,15 @@ bool test_equal(VectorType1 vector1,
{
typedef typename vtkm::VecTraits<VectorType1> Traits1;
typedef typename vtkm::VecTraits<VectorType2> Traits2;
BOOST_STATIC_ASSERT(Traits1::NUM_COMPONENTS == Traits2::NUM_COMPONENTS);
if (Traits1::GetNumberOfComponents(vector1) !=
Traits2::GetNumberOfComponents(vector2))
{
return false;
}
for (vtkm::IdComponent component = 0;
component < Traits1::NUM_COMPONENTS;
component < Traits1::GetNumberOfComponents(vector1);
component++)
{
vtkm::Float64 value1 =

2
vtkm/testing/TestingMath.h
View File

@ -160,7 +160,7 @@ struct ScalarFieldTests : public vtkm::exec::FunctorBase
T epsilon = vtkm::Epsilon<T>();
// General behavior.
VTKM_MATH_ASSERT(nan != nan, "Nan not equal itself.");
VTKM_MATH_ASSERT(nan != vtkm::Nan<T>(), "Nan not equal itself.");
VTKM_MATH_ASSERT(!(nan >= zero), "Nan not greater or less.");
VTKM_MATH_ASSERT(!(nan <= zero), "Nan not greater or less.");
VTKM_MATH_ASSERT(!(nan >= finite), "Nan not greater or less.");

30
vtkm/testing/UnitTestTypes.cxx
View File

@ -18,26 +18,6 @@
// this software.
//============================================================================
//============================================================================
// 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.
//============================================================================
#include <vtkm/Types.h>
#include <vtkm/testing/Testing.h>
@ -130,18 +110,26 @@ template<typename ComponentType, vtkm::IdComponent Size>
void GeneralVecTypeTest(const vtkm::Vec<ComponentType,Size> &)
{
typedef vtkm::Vec<ComponentType,Size> T;
VTKM_TEST_ASSERT(T::NUM_COMPONENTS == Size,
"NUM_COMPONENTS is wrong size.");
//grab the number of elements of T
T a, b, c;
ComponentType s(5);
VTKM_TEST_ASSERT(a.GetNumberOfComponents() == Size,
"GetNumberOfComponents returns wrong size.");
for(vtkm::IdComponent i=0; i < T::NUM_COMPONENTS; ++i)
{
a[i]=ComponentType((i+1)*2);
b[i]=ComponentType(i+1);
c[i]=ComponentType((i+1)*2);
}
a.CopyInto(c);
VTKM_TEST_ASSERT(test_equal(a, c), "CopyInto does not work.");
//verify prefix and postfix increment and decrement
++c[T::NUM_COMPONENTS-1];
c[T::NUM_COMPONENTS-1]++;

2
vtkm/testing/UnitTestVecTraits.cxx
View File

@ -39,7 +39,7 @@ struct TestVecTypeFunctor
{
Traits::SetComponent(vector, index, ComponentType(VecInit[index]));
}
vtkm::testing::TestVecType(vector);
vtkm::testing::TestVecType<Traits::NUM_COMPONENTS>(vector);
}
};

117
vtkm/testing/UnitTestVecVariable.cxx Normal file
View File

@ -0,0 +1,117 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
//
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/VecVariable.h>
#include <vtkm/testing/Testing.h>
namespace {
struct VecVariableTestFunctor
{
// You will get a compile fail if this does not pass
template<typename NumericTag>
void CheckNumericTag(NumericTag, NumericTag) const
{
std::cout << "NumericTag pass" << std::endl;
}
// You will get a compile fail if this does not pass
void CheckDimensionalityTag(vtkm::TypeTraitsVectorTag) const
{
std::cout << "VectorTag pass" << std::endl;
}
// You will get a compile fail if this does not pass
template<typename T>
void CheckComponentType(T, T) const
{
std::cout << "ComponentType pass" << std::endl;
}
// You will get a compile fail if this does not pass
void CheckHasMultipleComponents(vtkm::VecTraitsTagMultipleComponents) const
{
std::cout << "MultipleComponents pass" << std::endl;
}
// You will get a compile fail if this does not pass
void CheckVariableSize(vtkm::VecTraitsTagSizeVariable) const
{
std::cout << "VariableSize" << std::endl;
}
template<typename T>
void operator()(T) const
{
static const vtkm::IdComponent SIZE = 5;
typedef vtkm::Vec<T,SIZE> VecType;
typedef vtkm::VecVariable<T,SIZE> VecVariableType;
typedef vtkm::TypeTraits<VecVariableType> TTraits;
typedef vtkm::VecTraits<VecVariableType> VTraits;
std::cout << "Check NumericTag." << std::endl;
this->CheckNumericTag(typename TTraits::NumericTag(),
typename vtkm::TypeTraits<T>::NumericTag());
std::cout << "Check DimensionalityTag." << std::endl;
this->CheckDimensionalityTag(typename TTraits::DimensionalityTag());
std::cout << "Check ComponentType." << std::endl;
this->CheckComponentType(typename VTraits::ComponentType(), T());
std::cout << "Check MultipleComponents." << std::endl;
this->CheckHasMultipleComponents(typename VTraits::HasMultipleComponents());
std::cout << "Check VariableSize." << std::endl;
this->CheckVariableSize(typename VTraits::IsSizeStatic());
VecType source = TestValue(0, VecType());
VecVariableType vec1(source);
VecType vecCopy;
vec1.CopyInto(vecCopy);
VTKM_TEST_ASSERT(test_equal(vec1, vecCopy),
"Bad init or copyinto.");
vtkm::VecVariable<T,SIZE+1> vec2;
for (vtkm::IdComponent setIndex = 0; setIndex < SIZE; setIndex++)
{
VTKM_TEST_ASSERT(vec2.GetNumberOfComponents() == setIndex,
"Report wrong number of components");
vec2.Append(source[setIndex]);
}
VTKM_TEST_ASSERT(test_equal(vec2, vec1),
"Bad values from Append.");
}
};
void TestVecVariable()
{
vtkm::testing::Testing::TryTypes(VecVariableTestFunctor(),
vtkm::TypeListTagFieldScalar());
}
} // anonymous namespace
int UnitTestVecVariable(int, char *[])
{
return vtkm::testing::Testing::Run(TestVecVariable);
}

44
vtkm/testing/VecTraitsTests.h
View File

@ -46,17 +46,38 @@ inline void CompareDimensionalityTags(vtkm::TypeTraitsVectorTag,
// If we are here, everything is fine.
}
template<vtkm::IdComponent NUM_COMPONENTS, typename T>
inline void CheckIsStatic(const T &, vtkm::VecTraitsTagSizeStatic)
{
VTKM_TEST_ASSERT(vtkm::VecTraits<T>::NUM_COMPONENTS == NUM_COMPONENTS,
"Traits returns unexpected number of components");
}
template<vtkm::IdComponent NUM_COMPONENTS, typename T>
inline void CheckIsStatic(const T &, vtkm::VecTraitsTagSizeVariable)
{
// If we are here, everything is fine.
}
/// Compares some manual arithmetic through type traits to arithmetic with
/// the Tuple class.
template <class T>
template <vtkm::IdComponent NUM_COMPONENTS, typename T>
static void TestVecTypeImpl(
const typename boost::remove_const<T>::type &vector)
{
typedef typename vtkm::VecTraits<T> Traits;
typedef typename Traits::ComponentType ComponentType;
static const vtkm::IdComponent NUM_COMPONENTS = Traits::NUM_COMPONENTS;
typedef typename boost::remove_const<T>::type NonConstT;
CheckIsStatic<NUM_COMPONENTS>(vector, typename Traits::IsSizeStatic());
VTKM_TEST_ASSERT(Traits::GetNumberOfComponents(vector) == NUM_COMPONENTS,
"Traits returned wrong number of components.");
vtkm::Vec<ComponentType,NUM_COMPONENTS> vectorCopy;
Traits::CopyInto(vector, vectorCopy);
VTKM_TEST_ASSERT(test_equal(vectorCopy, vector), "CopyInto does not work.");
{
NonConstT result;
const ComponentType multiplier = 4;
@ -67,8 +88,9 @@ static void TestVecTypeImpl(
ComponentType(
multiplier*Traits::GetComponent(vector, i)));
}
VTKM_TEST_ASSERT(test_equal(Traits::ToVec(result),
multiplier*Traits::ToVec(vector)),
vtkm::Vec<ComponentType,NUM_COMPONENTS> resultCopy;
Traits::CopyInto(result, resultCopy);
VTKM_TEST_ASSERT(test_equal(resultCopy, multiplier*vectorCopy),
"Got bad result for scalar multiple");
}
@ -80,8 +102,9 @@ static void TestVecTypeImpl(
Traits::GetComponent(result, i)
= ComponentType(multiplier * Traits::GetComponent(vector, i));
}
VTKM_TEST_ASSERT(test_equal(Traits::ToVec(result),
multiplier*Traits::ToVec(vector)),
vtkm::Vec<ComponentType,NUM_COMPONENTS> resultCopy;
Traits::CopyInto(result, resultCopy);
VTKM_TEST_ASSERT(test_equal(resultCopy, multiplier*vectorCopy),
"Got bad result for scalar multiple");
}
@ -94,8 +117,7 @@ static void TestVecTypeImpl(
result = ComponentType(result + (component * component));
}
VTKM_TEST_ASSERT(
test_equal(result,
vtkm::dot(Traits::ToVec(vector), Traits::ToVec(vector))),
test_equal(result, vtkm::dot(vectorCopy, vectorCopy)),
"Got bad result for dot product");
}
@ -136,11 +158,11 @@ inline void CheckScalarComponentsTag(vtkm::VecTraitsTagSingleComponent)
/// Compares some manual arithmetic through type traits to arithmetic with
/// the Tuple class.
template <class T>
template <vtkm::IdComponent NUM_COMPONENTS, typename T>
static void TestVecType(const T &vector)
{
detail::TestVecTypeImpl<T>(vector);
detail::TestVecTypeImpl<const T>(vector);
detail::TestVecTypeImpl<NUM_COMPONENTS, T>(vector);
detail::TestVecTypeImpl<NUM_COMPONENTS, const T>(vector);
}
/// Checks to make sure that the HasMultipleComponents tag is actually for a

21
vtkm/worklet/CellAverage.h
View File

@ -23,35 +23,32 @@
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/exec/TopologyData.h>
namespace vtkm {
namespace worklet {
//simple functor that returns the average point value.
class CellAverage : public vtkm::worklet::WorkletMapTopology
{
static const int LEN_IDS = 8;
public:
typedef void ControlSignature(FieldInFrom<Scalar> inPoints,
TopologyIn<LEN_IDS> topology,
TopologyIn topology,
FieldOut<Scalar> outCells);
typedef void ExecutionSignature(_1, FromCount, _3);
typedef _2 InputDomain;
template<typename T1, typename T2>
template<typename PointValueVecType, typename OutType>
VTKM_EXEC_EXPORT
void operator()(const vtkm::exec::TopologyData<T1,LEN_IDS> &pointValues,
const vtkm::Id &count,
T2 &average) const
void operator()(const PointValueVecType &pointValues,
const vtkm::IdComponent &numPoints,
OutType &average) const
{
T1 sum = pointValues[0];
for (vtkm::IdComponent i=1; i< count; ++i)
OutType sum = static_cast<OutType>(pointValues[0]);
for (vtkm::IdComponent pointIndex = 1; pointIndex < numPoints; ++pointIndex)
{
sum += pointValues[i];
sum = sum + static_cast<OutType>(pointValues[pointIndex]);
}
average = static_cast<T2>(sum / static_cast<T1>(count));
average = sum / static_cast<OutType>(numPoints);
}
};

41
vtkm/exec/TopologyData.h → vtkm/worklet/Magnitude.h
View File

@ -6,9 +6,9 @@
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//
// Copyright 2015 Sandia Corporation.
// Copyright 2015 UT-Battelle, LLC.
// Copyright 2015 Los Alamos National Security.
// 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.
@ -17,35 +17,32 @@
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_exec_TopologyData_h
#define vtk_m_exec_TopologyData_h
#ifndef vtk_m_worklet_Magnitude_h
#define vtk_m_worklet_Magnitude_h
#include <vtkm/Types.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/VectorAnalysis.h>
namespace vtkm {
namespace exec {
namespace worklet {
template<typename T, vtkm::IdComponent ITEM_TUPLE_LENGTH = 8>
class TopologyData
class Magnitude : public vtkm::worklet::WorkletMapField
{
public:
VTKM_EXEC_EXPORT T &operator[](vtkm::IdComponent index) { return vec[index]; }
VTKM_EXEC_EXPORT const T &operator[](vtkm::IdComponent index) const { return vec[index]; }
typedef void ControlSignature(FieldIn<>, FieldOut<>);
typedef void ExecutionSignature(_1, _2);
VTKM_EXEC_EXPORT TopologyData()
template<typename T>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Vec<T,3> &inValue,
T &outValue) const
{
outValue = vtkm::Magnitude(inValue);
}
template <typename T2>
VTKM_EXEC_EXPORT TopologyData(const TopologyData<T2,ITEM_TUPLE_LENGTH> &other)
: vec(other.vec)
{
}
vtkm::Vec<T, ITEM_TUPLE_LENGTH> vec;
};
}
} // namespace vtkm::exec
} // namespace vtkm::worklet
#endif //vtk_m_exec_FunctorBase_h
#endif // vtk_m_worklet_Magnitude_h

5
vtkm/worklet/WorkletMapTopology.h
View File

@ -83,15 +83,10 @@ public:
/// \brief A control signature tag for input connectivity.
///
/// This tag takes a template argument that is a type list tag that limits
/// the possible value types in the array.
///
template< vtkm::IdComponent ItemTupleLength>
struct TopologyIn : vtkm::cont::arg::ControlSignatureTagBase {
typedef vtkm::cont::arg::TypeCheckTagTopology TypeCheckTag;
typedef vtkm::cont::arg::TransportTagTopologyIn TransportTag;
typedef vtkm::exec::arg::FetchTagTopologyIn FetchTag;
static const int ITEM_TUPLE_LENGTH = ItemTupleLength;
};
/// \brief A control signature tag for output fields.

33
vtkm/worklet/testing/UnitTestWorkletMapFieldExecArg.cxx
View File

@ -31,25 +31,27 @@
class TestExecObjectWorklet : public vtkm::worklet::WorkletMapField
{
public:
typedef void ControlSignature( FieldIn<>, ExecObject, FieldOut<>);
typedef void ExecutionSignature(_1, _2, _3);
typedef void ControlSignature( FieldIn<>, ExecObject, ExecObject, FieldOut<>);
typedef void ExecutionSignature(_1, _2, _3, _4);
template<typename T, typename StorageTag>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &index,
const vtkm::exec::ExecutionWholeArrayConst<T,StorageTag> &execArg,
const vtkm::exec::ExecutionWholeArrayConst<T,StorageTag> &execIn,
vtkm::exec::ExecutionWholeArray<T,StorageTag> &execOut,
T& out) const
{
if (!test_equal(execArg.Get(index), TestValue(index, T()) + T(100)))
if (!test_equal(execIn.Get(index), TestValue(index, T()) + T(100)))
{
this->RaiseError("Got wrong input value.");
}
out = execArg.Get(index) - T(100);
out = execIn.Get(index) - T(100);
execOut.Set(index, out);
}
template<typename T1, typename T2>
template<typename T1, typename T2, typename T3>
VTKM_EXEC_EXPORT
void operator()(const vtkm::Id &, const T1 &, const T2 &) const
void operator()(const vtkm::Id &, const T1 &, const T2 &, const T3&) const
{
this->RaiseError("Cannot call this worklet with different types.");
}
@ -78,24 +80,33 @@ struct DoTestWorklet
vtkm::cont::ArrayHandle<T> inputHandle =
vtkm::cont::make_ArrayHandle(inputArray, ARRAY_SIZE);
vtkm::cont::ArrayHandle<T> outputHandle;
vtkm::cont::ArrayHandle<T> outputFieldArray;
std::cout << "Create and run dispatcher." << std::endl;
vtkm::worklet::DispatcherMapField<WorkletType> dispatcher;
dispatcher.Invoke(counting,
vtkm::exec::ExecutionWholeArrayConst<T>(inputHandle),
outputHandle);
vtkm::exec::ExecutionWholeArray<T>(outputHandle, ARRAY_SIZE),
outputFieldArray);
std::cout << "Check result." << std::endl;
CheckPortal(outputHandle.GetPortalConstControl());
CheckPortal(outputFieldArray.GetPortalConstControl());
std::cout << "Repeat with dynamic arrays." << std::endl;
// Clear out output array.
// Clear out output arrays.
outputFieldArray = vtkm::cont::ArrayHandle<T>();
outputHandle = vtkm::cont::ArrayHandle<T>();
vtkm::cont::DynamicArrayHandle outputDynamic(outputHandle);
vtkm::cont::DynamicArrayHandle outputFieldDynamic(outputFieldArray);
dispatcher.Invoke(counting,
vtkm::exec::ExecutionWholeArrayConst<T>(inputHandle),
outputDynamic);
vtkm::exec::ExecutionWholeArray<T>(outputHandle, ARRAY_SIZE),
outputFieldDynamic);
std::cout << "Check dynamic array result." << std::endl;
CheckPortal(outputHandle.GetPortalConstControl());
CheckPortal(outputFieldArray.GetPortalConstControl());
}
};

71
vtkm/worklet/testing/UnitTestWorkletMapTopologyExplicit.cxx
View File

@ -21,6 +21,8 @@
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/Math.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/testing/Testing.h>
@ -30,11 +32,10 @@ namespace test_explicit {
class MaxPointOrCellValue : public vtkm::worklet::WorkletMapTopology
{
static const int LEN_IDS = 8;
public:
typedef void ControlSignature(FieldInTo<Scalar> inCells,
FieldInFrom<Scalar> inPoints,
TopologyIn<LEN_IDS> topology,
TopologyIn topology,
FieldOut<Scalar> outCells);
typedef void ExecutionSignature(_1, _4, _2, FromCount, CellShape, FromIndices);
typedef _3 InputDomain;
@ -42,43 +43,33 @@ public:
VTKM_CONT_EXPORT
MaxPointOrCellValue() { }
template<typename T>
template<typename InCellType,
typename OutCellType,
typename InPointVecType,
typename FromIndexType>
VTKM_EXEC_EXPORT
void operator()(const T &cellValue,
T& maxValue,
const vtkm::exec::TopologyData<T,LEN_IDS> &pointValues,
const vtkm::Id &count,
const vtkm::Id & vtkmNotUsed(type),
const vtkm::exec::TopologyData<vtkm::Id,LEN_IDS> & vtkmNotUsed(pointIDs)) const
void operator()(const InCellType &cellValue,
OutCellType &maxValue,
const InPointVecType &pointValues,
const vtkm::IdComponent &numPoints,
const vtkm::Id &vtkmNotUsed(type),
const FromIndexType &vtkmNotUsed(pointIDs)) const
{
//simple functor that returns the max of cellValue and pointValue
maxValue = cellValue;
for (vtkm::IdComponent i=0; i<count; ++i)
//simple functor that returns the max of cellValue and pointValue
maxValue = static_cast<OutCellType>(cellValue);
for (vtkm::IdComponent pointIndex = 0; pointIndex < numPoints; ++pointIndex)
{
maxValue = pointValues[i] > maxValue ? pointValues[i] : maxValue;
maxValue = vtkm::Max(maxValue,
static_cast<OutCellType>(pointValues[pointIndex]));
}
}
template<typename T1, typename T2, typename T3>
VTKM_EXEC_EXPORT
void operator()(const T1 &,
T2 &,
const vtkm::exec::TopologyData<T3,LEN_IDS> &,
const vtkm::Id &,
const vtkm::Id &,
const vtkm::exec::TopologyData<vtkm::Id,LEN_IDS> &) const
{
this->RaiseError("Cannot call this worklet with different types.");
}
};
class AveragePointToCellValue : public vtkm::worklet::WorkletMapTopology
{
static const int LEN_IDS = 8;
public:
typedef void ControlSignature(FieldInFrom<Scalar> inPoints,
TopologyIn<LEN_IDS> topology,
TopologyIn topology,
FieldOut<Scalar> outCells);
typedef void ExecutionSignature(_1, _3, FromCount);
typedef _2 InputDomain;
@ -86,28 +77,20 @@ public:
VTKM_CONT_EXPORT
AveragePointToCellValue() { }
template<typename T>
template<typename PointVecType, typename OutType>
VTKM_EXEC_EXPORT
void operator()(const vtkm::exec::TopologyData<T,LEN_IDS> &pointValues,
T& avgVal,
const vtkm::Id &count) const
void operator()(const PointVecType &pointValues,
OutType &avgVal,
const vtkm::IdComponent &numPoints) const
{
//simple functor that returns the average pointValue.
avgVal = pointValues[0];
for (vtkm::IdComponent i=1; i<count; ++i)
avgVal = static_cast<OutType>(pointValues[0]);
for (vtkm::IdComponent pointIndex = 1; pointIndex < numPoints; ++pointIndex)
{
avgVal += pointValues[i];
avgVal += static_cast<OutType>(pointValues[pointIndex]);
}
avgVal = avgVal / count;
avgVal = avgVal / static_cast<OutType>(numPoints);
}
template<typename T1, typename T2>
VTKM_EXEC_EXPORT
void operator()(const T1 &, T2 &, const vtkm::Id &) const
{
this->RaiseError("Cannot call this worklet with different types.");
}
};
}

74
vtkm/worklet/testing/UnitTestWorkletMapTopologyRegular.cxx
View File

@ -21,9 +21,9 @@
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/WorkletMapTopology.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/Math.h>
#include <vtkm/exec/TopologyData.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/cont/testing/MakeTestDataSet.h>
@ -32,56 +32,44 @@ namespace test_regular {
class MaxPointOrCellValue : public vtkm::worklet::WorkletMapTopology
{
static const int LEN_IDS = 4;
public:
typedef void ControlSignature(FieldInTo<Scalar> inCells,
FieldInFrom<Scalar> inPoints,
TopologyIn<LEN_IDS> topology,
TopologyIn topology,
FieldOut<Scalar> outCells);
//Todo: we need a way to mark what control signature item each execution signature for topology comes from
typedef void ExecutionSignature(_1, _4, _2, FromCount, CellShape, FromIndices);
typedef _3 InputDomain;
VTKM_CONT_EXPORT
MaxPointOrCellValue() { }
template<typename T>
template<typename InCellType,
typename OutCellType,
typename InPointVecType,
typename FromIndexType>
VTKM_EXEC_EXPORT
void operator()(const T &cellval,
T& max_value,
const vtkm::exec::TopologyData<T,LEN_IDS> &pointValues,
const vtkm::Id &count,
const vtkm::Id & vtkmNotUsed(type),
const vtkm::exec::TopologyData<vtkm::Id,LEN_IDS> & vtkmNotUsed(pointIDs)) const
void operator()(const InCellType &cellValue,
OutCellType &maxValue,
const InPointVecType &pointValues,
const vtkm::IdComponent &numPoints,
const vtkm::Id &vtkmNotUsed(type),
const FromIndexType &vtkmNotUsed(pointIDs)) const
{
//simple functor that returns the max of cellValue and pointValue
max_value = cellval;
for (vtkm::IdComponent i=0; i< count; ++i)
//simple functor that returns the max of cellValue and pointValue
maxValue = static_cast<OutCellType>(cellValue);
for (vtkm::IdComponent pointIndex = 0; pointIndex < numPoints; ++pointIndex)
{
max_value = pointValues[i] > max_value ? pointValues[i] : max_value;
maxValue = vtkm::Max(maxValue,
static_cast<OutCellType>(pointValues[pointIndex]));
}
}
template<typename T1, typename T2, typename T3>
VTKM_EXEC_EXPORT
void operator()(const T1 &,
T2 &,
const vtkm::exec::TopologyData<T3,LEN_IDS> &,
const vtkm::Id &,
const vtkm::Id &,
const vtkm::exec::TopologyData<vtkm::Id,LEN_IDS> &) const
{
this->RaiseError("Cannot call this worklet with different types.");
}
};
class AveragePointToCellValue : public vtkm::worklet::WorkletMapTopology
{
static const int LEN_IDS = 4;
public:
typedef void ControlSignature(FieldInFrom<Scalar> inPoints,
TopologyIn<LEN_IDS> topology,
TopologyIn topology,
FieldOut<Scalar> outCells);
typedef void ExecutionSignature(_1, _3, FromCount);
typedef _2 InputDomain;
@ -89,28 +77,20 @@ public:
VTKM_CONT_EXPORT
AveragePointToCellValue() { }
template<typename T>
template<typename PointVecType, typename OutType>
VTKM_EXEC_EXPORT
void operator()(const vtkm::exec::TopologyData<T,LEN_IDS> &pointValues,
T& avgVal,
const vtkm::Id &count) const
void operator()(const PointVecType &pointValues,
OutType &avgVal,
const vtkm::IdComponent &numPoints) const
{
//simple functor that returns the average pointValue.
avgVal = pointValues[0];
for (vtkm::IdComponent i=1; i<count; ++i)
avgVal = static_cast<OutType>(pointValues[0]);
for (vtkm::IdComponent pointIndex = 1; pointIndex < numPoints; ++pointIndex)
{
avgVal += pointValues[i];
avgVal += static_cast<OutType>(pointValues[pointIndex]);
}
avgVal = avgVal / count;
avgVal = avgVal / static_cast<OutType>(numPoints);
}
template<typename T1, typename T2>
VTKM_EXEC_EXPORT
void operator()(const T1 &, T2 &, const vtkm::Id &) const
{
this->RaiseError("Cannot call this worklet with different types.");
}
};
}