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

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

Browse Source
This commit is contained in:
Dave Pugmire 2020-11-02 16:28:57 -05:00
commit a7f049ac1a
36 changed files with 632 additions and 550 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
CMake/VTKmDeviceAdapters.cmake
View File

@ -167,7 +167,10 @@ if(VTKm_ENABLE_CUDA)
# 6 - volta
# - Uses: --generate-code=arch=compute_70,code=sm_70
# 7 - turing
# - Uses: --generate-code=arch=compute_75code=sm_75
# - Uses: --generate-code=arch=compute_75,code=sm_75
# 8 - ampere
# - Uses: --generate-code=arch=compute_80,code=sm_80
# - Uses: --generate-code=arch=compute_86,code=sm_86
# 8 - all
# - Uses: --generate-code=arch=compute_30,code=sm_30
# - Uses: --generate-code=arch=compute_35,code=sm_35
@ -175,12 +178,14 @@ if(VTKm_ENABLE_CUDA)
# - Uses: --generate-code=arch=compute_60,code=sm_60
# - Uses: --generate-code=arch=compute_70,code=sm_70
# - Uses: --generate-code=arch=compute_75,code=sm_75
# - Uses: --generate-code=arch=compute_80,code=sm_80
# - Uses: --generate-code=arch=compute_86,code=sm_86
# 8 - none
#
#specify the property
set(VTKm_CUDA_Architecture "native" CACHE STRING "Which GPU Architecture(s) to compile for")
set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing all none)
set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing ampere all none)
#detect what the property is set too
if(VTKm_CUDA_Architecture STREQUAL "native")
@ -234,13 +239,18 @@ if(VTKm_ENABLE_CUDA)
set(arch_flags --generate-code=arch=compute_70,code=sm_70)
elseif(VTKm_CUDA_Architecture STREQUAL "turing")
set(arch_flags --generate-code=arch=compute_75,code=sm_75)
elseif(VTKm_CUDA_Architecture STREQUAL "ampere")
set(arch_flags --generate-code=arch=compute_80,code=sm_80)
set(arch_flags --generate-code=arch=compute_86,code=sm_86)
elseif(VTKm_CUDA_Architecture STREQUAL "all")
set(arch_flags --generate-code=arch=compute_30,code=sm_30
--generate-code=arch=compute_35,code=sm_35
--generate-code=arch=compute_50,code=sm_50
--generate-code=arch=compute_60,code=sm_60
--generate-code=arch=compute_70,code=sm_70
--generate-code=arch=compute_75,code=sm_75)
--generate-code=arch=compute_75,code=sm_75
--generate-code=arch=compute_80,code=sm_80
--generate-code=arch=compute_86,code=sm_86)
endif()
string(REPLACE ";" " " arch_flags "${arch_flags}")

56
benchmarking/BenchmarkFilters.cxx
View File

@ -347,10 +347,11 @@ void BenchContour(::benchmark::State& state)
{
const vtkm::cont::DeviceAdapterId device = Config.Device;
const vtkm::Id numIsoVals = static_cast<vtkm::Id>(state.range(0));
const bool mergePoints = static_cast<bool>(state.range(1));
const bool normals = static_cast<bool>(state.range(2));
const bool fastNormals = static_cast<bool>(state.range(3));
const bool isStructured = static_cast<vtkm::Id>(state.range(0));
const vtkm::Id numIsoVals = static_cast<vtkm::Id>(state.range(1));
const bool mergePoints = static_cast<bool>(state.range(2));
const bool normals = static_cast<bool>(state.range(3));
const bool fastNormals = static_cast<bool>(state.range(4));
vtkm::filter::Contour filter;
filter.SetActiveField(PointScalarsName, vtkm::cont::Field::Association::POINTS);
@ -376,11 +377,14 @@ void BenchContour(::benchmark::State& state)
filter.SetComputeFastNormalsForUnstructured(fastNormals);
vtkm::cont::Timer timer{ device };
vtkm::cont::DataSet input = isStructured ? InputDataSet : UnstructuredInputDataSet;
for (auto _ : state)
{
(void)_;
timer.Start();
auto result = filter.Execute(InputDataSet);
auto result = filter.Execute(input);
::benchmark::DoNotOptimize(result);
timer.Stop();
@ -390,13 +394,17 @@ void BenchContour(::benchmark::State& state)
void BenchContourGenerator(::benchmark::internal::Benchmark* bm)
{
bm->ArgNames({ "NIsoVals", "MergePts", "GenNormals", "FastNormals" });
bm->ArgNames({ "IsStructuredDataSet", "NIsoVals", "MergePts", "GenNormals", "FastNormals" });
auto helper = [&](const vtkm::Id numIsoVals) {
bm->Args({ numIsoVals, 0, 0, 0 });
bm->Args({ numIsoVals, 1, 0, 0 });
bm->Args({ numIsoVals, 0, 1, 0 });
bm->Args({ numIsoVals, 0, 1, 1 });
bm->Args({ 0, numIsoVals, 0, 0, 0 });
bm->Args({ 0, numIsoVals, 1, 0, 0 });
bm->Args({ 0, numIsoVals, 0, 1, 0 });
bm->Args({ 0, numIsoVals, 0, 1, 1 });
bm->Args({ 1, numIsoVals, 0, 0, 0 });
bm->Args({ 1, numIsoVals, 1, 0, 0 });
bm->Args({ 1, numIsoVals, 0, 1, 0 });
bm->Args({ 1, numIsoVals, 0, 1, 1 });
};
helper(1);
@ -405,7 +413,7 @@ void BenchContourGenerator(::benchmark::internal::Benchmark* bm)
}
// :TODO: Disabled until SIGSEGV in Countour when passings field is resolved
//VTKM_BENCHMARK_APPLY(BenchContour, BenchContourGenerator);
VTKM_BENCHMARK_APPLY(BenchContour, BenchContourGenerator);
void BenchExternalFaces(::benchmark::State& state)
{
@ -993,26 +1001,22 @@ void InitDataSet(int& argc, char** argv)
source.SetExtent({ 0 }, { waveletDim - 1 });
InputDataSet = source.Execute();
vtkm::cont::DataSet input = vtkm::cont::testing::MakeTestDataSet().Make2DUniformDataSet2();
vtkm::filter::Triangulate triangulateFilter;
triangulateFilter.SetFieldsToPass(
vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL));
UnstructuredInputDataSet = triangulateFilter.Execute(input);
}
if (tetra)
{
std::cerr << "[InitDataSet] Tetrahedralizing dataset...\n";
vtkm::filter::Tetrahedralize tet;
tet.SetFieldsToPass(vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL));
InputDataSet = tet.Execute(InputDataSet);
}
FindFields();
CreateMissingFields();
std::cerr
<< "[InitDataSet] Create UnstructuredInputDataSet from Tetrahedralized InputDataSet...\n";
vtkm::filter::Tetrahedralize tet;
tet.SetFieldsToPass(vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL));
UnstructuredInputDataSet = tet.Execute(InputDataSet);
if (tetra)
{
InputDataSet = UnstructuredInputDataSet;
}
inputGenTimer.Stop();
std::cerr << "[InitDataSet] DataSet initialization took " << inputGenTimer.GetElapsedTime()

3
vtkm/Math.h
View File

@ -2562,7 +2562,8 @@ static inline VTKM_EXEC_CONT vtkm::Vec<T, N> CopySign(const vtkm::Vec<T, N>& x,
inline VTKM_EXEC_CONT vtkm::Float32 Frexp(vtkm::Float32 x, vtkm::Int32 *exponent)
{
#ifdef VTKM_CUDA
// See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
#if defined(VTKM_CUDA) || defined(VTKM_HIP)
return VTKM_CUDA_MATH_FUNCTION_32(frexp)(x, exponent);
#else
return std::frexp(x, exponent);

3
vtkm/Math.h.in
View File

@ -1164,7 +1164,8 @@ static inline VTKM_EXEC_CONT vtkm::Vec<T, N> CopySign(const vtkm::Vec<T, N>& x,
inline VTKM_EXEC_CONT vtkm::Float32 Frexp(vtkm::Float32 x, vtkm::Int32 *exponent)
{
#ifdef VTKM_CUDA
// See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
#if defined(VTKM_CUDA) || defined(VTKM_HIP)
return VTKM_CUDA_MATH_FUNCTION_32(frexp)(x, exponent);
#else
return std::frexp(x, exponent);

18
vtkm/Swap.h
View File

@ -13,7 +13,7 @@
#include <vtkm/internal/ExportMacros.h>
#ifdef __CUDACC__
#ifdef VTKM_CUDA
#include <thrust/swap.h>
#else
#include <algorithm>
@ -23,13 +23,27 @@ namespace vtkm
{
/// Performs a swap operation. Safe to call from cuda code.
#ifdef __CUDACC__
#if defined(VTKM_CUDA)
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)
{
using namespace thrust;
swap(a, b);
}
#elif defined(VTKM_HIP)
template <typename T>
__host__ void Swap(T& a, T& b)
{
using namespace std;
swap(a, b);
}
template <typename T>
__device__ void Swap(T& a, T& b)
{
T temp = a;
a = b;
b = temp;
}
#else
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)

18
vtkm/cont/ArrayHandleConstant.h
View File

@ -47,20 +47,13 @@ using StorageTagConstantSuperclass =
template <typename T>
struct Storage<T, vtkm::cont::StorageTagConstant> : Storage<T, StorageTagConstantSuperclass<T>>
{
using Superclass = Storage<T, StorageTagConstantSuperclass<T>>;
using Superclass::Superclass;
};
template <typename T, typename Device>
struct ArrayTransfer<T, vtkm::cont::StorageTagConstant, Device>
: ArrayTransfer<T, StorageTagConstantSuperclass<T>, Device>
{
using Superclass = ArrayTransfer<T, StorageTagConstantSuperclass<T>, Device>;
using Superclass::Superclass;
};
} // namespace internal
template <typename T>
VTKM_ARRAY_HANDLE_NEW_STYLE(T, vtkm::cont::StorageTagConstant);
/// \brief An array handle with a constant value.
///
/// ArrayHandleConstant is an implicit array handle with a constant value. A
@ -79,9 +72,8 @@ public:
VTKM_CONT
ArrayHandleConstant(T value, vtkm::Id numberOfValues = 0)
: Superclass(typename internal::Storage<T, StorageTag>::PortalConstType(
internal::ConstantFunctor<T>(value),
numberOfValues))
: Superclass(internal::FunctorToArrayHandleImplicitBuffers(internal::ConstantFunctor<T>(value),
numberOfValues))
{
}
};

37
vtkm/cont/ArrayHandleCounting.h
View File

@ -10,8 +10,7 @@
#ifndef vtk_m_cont_ArrayHandleCounting_h
#define vtk_m_cont_ArrayHandleCounting_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageImplicit.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/VecTraits.h>
@ -52,24 +51,6 @@ public:
{
}
template <typename OtherValueType>
VTKM_EXEC_CONT ArrayPortalCounting(const ArrayPortalCounting<OtherValueType>& src)
: Start(src.Start)
, Step(src.Step)
, NumberOfValues(src.NumberOfValues)
{
}
template <typename OtherValueType>
VTKM_EXEC_CONT ArrayPortalCounting<ValueType>& operator=(
const ArrayPortalCounting<OtherValueType>& src)
{
this->Start = src.Start;
this->Step = src.Step;
this->NumberOfValues = src.NumberOfValues;
return *this;
}
VTKM_EXEC_CONT
ValueType GetStart() const { return this->Start; }
@ -98,20 +79,13 @@ using StorageTagCountingSuperclass =
template <typename T>
struct Storage<T, vtkm::cont::StorageTagCounting> : Storage<T, StorageTagCountingSuperclass<T>>
{
using Superclass = Storage<T, StorageTagCountingSuperclass<T>>;
using Superclass::Superclass;
};
template <typename T, typename Device>
struct ArrayTransfer<T, vtkm::cont::StorageTagCounting, Device>
: ArrayTransfer<T, StorageTagCountingSuperclass<T>, Device>
{
using Superclass = ArrayTransfer<T, StorageTagCountingSuperclass<T>, Device>;
using Superclass::Superclass;
};
} // namespace internal
template <typename T>
VTKM_ARRAY_HANDLE_NEW_STYLE(T, vtkm::cont::StorageTagCounting);
/// ArrayHandleCounting is a specialization of ArrayHandle. By default it
/// contains a increment value, that is increment for each step between zero
/// and the passed in length
@ -126,7 +100,8 @@ public:
VTKM_CONT
ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::Id length)
: Superclass(internal::ArrayPortalCounting<CountingValueType>(start, step, length))
: Superclass(internal::PortalToArrayHandleImplicitBuffers(
internal::ArrayPortalCounting<CountingValueType>(start, step, length)))
{
}
};

75
vtkm/cont/ArrayHandleImplicit.cxx Normal file
View File

@ -0,0 +1,75 @@
//============================================================================
// 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.
//============================================================================
#include <vtkm/cont/ArrayHandleImplicit.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
BufferMetaDataImplicit::BufferMetaDataImplicit(const BufferMetaDataImplicit& src)
: Deleter(src.Deleter)
, Copier(src.Copier)
{
if (src.Portal)
{
VTKM_ASSERT(src.Deleter);
VTKM_ASSERT(src.Copier);
this->Portal = src.Copier(src.Portal);
}
else
{
this->Portal = nullptr;
}
}
BufferMetaDataImplicit::~BufferMetaDataImplicit()
{
if (this->Portal)
{
VTKM_ASSERT(this->Deleter);
this->Deleter(this->Portal);
this->Portal = nullptr;
}
}
std::unique_ptr<vtkm::cont::internal::BufferMetaData> BufferMetaDataImplicit::DeepCopy() const
{
return std::unique_ptr<vtkm::cont::internal::BufferMetaData>(new BufferMetaDataImplicit(*this));
}
namespace detail
{
vtkm::cont::internal::BufferMetaDataImplicit* GetImplicitMetaData(
const vtkm::cont::internal::Buffer& buffer)
{
vtkm::cont::internal::BufferMetaDataImplicit* metadata =
dynamic_cast<vtkm::cont::internal::BufferMetaDataImplicit*>(buffer.GetMetaData());
VTKM_ASSERT(metadata && "Buffer for implicit array does not have correct metadata.");
return metadata;
}
} // namespace detail
}
}
} // namespace vtkm::cont::internal
namespace vtkm
{
namespace cont
{
}
} // namespace vtkm::cont::detail

192
vtkm/cont/ArrayHandleImplicit.h
View File

@ -11,30 +11,15 @@
#define vtk_m_cont_ArrayHandleImplicit_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageImplicit.h>
#include <vtkmstd/is_trivial.h>
namespace vtkm
{
namespace cont
{
namespace detail
namespace internal
{
template <class FunctorType_>
class VTKM_ALWAYS_EXPORT ArrayPortalImplicit;
/// A convenience class that provides a typedef to the appropriate tag for
/// a implicit array container.
template <typename FunctorType>
struct ArrayHandleImplicitTraits
{
using ValueType = decltype(FunctorType{}(vtkm::Id{}));
using StorageTag = vtkm::cont::StorageTagImplicit<ArrayPortalImplicit<FunctorType>>;
using Superclass = vtkm::cont::ArrayHandle<ValueType, StorageTag>;
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
};
/// \brief An array portal that returns the result of a functor
///
/// This array portal is similar to an implicit array i.e an array that is
@ -48,8 +33,8 @@ template <class FunctorType_>
class VTKM_ALWAYS_EXPORT ArrayPortalImplicit
{
public:
using ValueType = typename ArrayHandleImplicitTraits<FunctorType_>::ValueType;
using FunctorType = FunctorType_;
using ValueType = decltype(FunctorType{}(vtkm::Id{}));
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
@ -82,8 +67,170 @@ private:
vtkm::Id NumberOfValues;
};
} // namespace internal
namespace cont
{
/// \brief An implementation for read-only implicit arrays.
///
/// It is sometimes the case that you want VTK-m to operate on an array of
/// implicit values. That is, rather than store the data in an actual array, it
/// is gerenated on the fly by a function. This is handled in VTK-m by creating
/// an ArrayHandle in VTK-m with a StorageTagImplicit type of \c Storage. This
/// tag itself is templated to specify an ArrayPortal that generates the
/// desired values. An ArrayHandle created with this tag will raise an error on
/// any operation that tries to modify it.
///
template <class ArrayPortalType>
struct VTKM_ALWAYS_EXPORT StorageTagImplicit
{
using PortalType = ArrayPortalType;
};
namespace internal
{
struct VTKM_CONT_EXPORT BufferMetaDataImplicit : vtkm::cont::internal::BufferMetaData
{
void* Portal;
using DeleterType = void(void*);
DeleterType* Deleter;
using CopierType = void*(void*);
CopierType* Copier;
template <typename PortalType>
BufferMetaDataImplicit(const PortalType& portal)
: Portal(new PortalType(portal))
, Deleter([](void* p) { delete reinterpret_cast<PortalType*>(p); })
, Copier([](void* p) -> void* { return new PortalType(*reinterpret_cast<PortalType*>(p)); })
{
}
VTKM_CONT BufferMetaDataImplicit(const BufferMetaDataImplicit& src);
BufferMetaDataImplicit& operator=(const BufferMetaDataImplicit&) = delete;
VTKM_CONT ~BufferMetaDataImplicit() override;
VTKM_CONT std::unique_ptr<vtkm::cont::internal::BufferMetaData> DeepCopy() const override;
};
namespace detail
{
VTKM_CONT_EXPORT vtkm::cont::internal::BufferMetaDataImplicit* GetImplicitMetaData(
const vtkm::cont::internal::Buffer& buffer);
} // namespace detail
template <class ArrayPortalType>
struct VTKM_ALWAYS_EXPORT
Storage<typename ArrayPortalType::ValueType, StorageTagImplicit<ArrayPortalType>>
{
VTKM_IS_TRIVIALLY_COPYABLE(ArrayPortalType);
using ReadPortalType = ArrayPortalType;
// Note that this portal is almost certainly read-only, so you will probably get
// an error if you try to write to it.
using WritePortalType = ArrayPortalType;
// Implicit array has one buffer that should be empty (NumberOfBytes = 0), but holds
// the metadata for the array.
VTKM_CONT static vtkm::IdComponent GetNumberOfBuffers() { return 1; }
VTKM_CONT static vtkm::Id GetNumberOfValues(const vtkm::cont::internal::Buffer* buffers)
{
vtkm::cont::internal::BufferMetaDataImplicit* metadata =
detail::GetImplicitMetaData(buffers[0]);
VTKM_ASSERT(metadata->Portal);
return reinterpret_cast<ArrayPortalType*>(metadata->Portal)->GetNumberOfValues();
}
VTKM_CONT static void ResizeBuffers(vtkm::Id numValues,
vtkm::cont::internal::Buffer* buffers,
vtkm::CopyFlag,
vtkm::cont::Token&)
{
if (numValues == GetNumberOfValues(buffers))
{
// In general, we don't allow resizing of the array, but if it was "allocated" to the
// correct size, we will allow that.
}
else
{
throw vtkm::cont::ErrorBadAllocation("Cannot allocate/resize implicit arrays.");
}
}
VTKM_CONT static ReadPortalType CreateReadPortal(const vtkm::cont::internal::Buffer* buffers,
vtkm::cont::DeviceAdapterId,
vtkm::cont::Token&)
{
vtkm::cont::internal::BufferMetaDataImplicit* metadata =
detail::GetImplicitMetaData(buffers[0]);
VTKM_ASSERT(metadata->Portal);
return *reinterpret_cast<ReadPortalType*>(metadata->Portal);
}
VTKM_CONT static WritePortalType CreateWritePortal(const vtkm::cont::internal::Buffer*,
vtkm::cont::DeviceAdapterId,
vtkm::cont::Token&)
{
throw vtkm::cont::ErrorBadAllocation("Cannot write to implicit arrays.");
}
};
/// Given an array portal, returns the buffers for the `ArrayHandle` with a storage that
/// is (or is compatible with) a storage tag of `StorageTagImplicit<PortalType>`.
template <typename PortalType>
VTKM_CONT inline std::vector<vtkm::cont::internal::Buffer> PortalToArrayHandleImplicitBuffers(
const PortalType& portal)
{
std::vector<vtkm::cont::internal::Buffer> buffers(1);
buffers[0].SetMetaData(std::unique_ptr<vtkm::cont::internal::BufferMetaData>(
new vtkm::cont::internal::BufferMetaDataImplicit(portal)));
return buffers;
}
/// Given a functor and the number of values, returns the buffers for the `ArrayHandleImplicit`
/// for the given functor.
template <typename FunctorType>
VTKM_CONT inline std::vector<vtkm::cont::internal::Buffer> FunctorToArrayHandleImplicitBuffers(
const FunctorType& functor,
vtkm::Id numValues)
{
return PortalToArrayHandleImplicitBuffers(
vtkm::internal::ArrayPortalImplicit<FunctorType>(functor, numValues));
}
} // namespace internal
namespace detail
{
/// A convenience class that provides a typedef to the appropriate tag for
/// a implicit array container.
template <typename FunctorType>
struct ArrayHandleImplicitTraits
{
using ValueType = decltype(FunctorType{}(vtkm::Id{}));
using PortalType = vtkm::internal::ArrayPortalImplicit<FunctorType>;
using StorageTag = vtkm::cont::StorageTagImplicit<PortalType>;
using Superclass = vtkm::cont::ArrayHandle<ValueType, StorageTag>;
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
};
} // namespace detail
// This can go away once ArrayHandle is replaced with ArrayHandleNewStyle
template <typename PortalType>
VTKM_ARRAY_HANDLE_NEW_STYLE(typename PortalType::ValueType,
vtkm::cont::StorageTagImplicit<PortalType>);
/// \brief An \c ArrayHandle that computes values on the fly.
///
/// \c ArrayHandleImplicit is a specialization of ArrayHandle.
@ -97,6 +244,7 @@ class VTKM_ALWAYS_EXPORT ArrayHandleImplicit
{
private:
using ArrayTraits = typename detail::ArrayHandleImplicitTraits<FunctorType>;
using PortalType = typename ArrayTraits::PortalType;
public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleImplicit,
@ -105,7 +253,7 @@ public:
VTKM_CONT
ArrayHandleImplicit(FunctorType functor, vtkm::Id length)
: Superclass(typename ArrayTraits::StorageType::PortalType(functor, length))
: Superclass(internal::PortalToArrayHandleImplicitBuffers(PortalType(functor, length)))
{
}
};
@ -144,7 +292,7 @@ struct SerializableTypeString<vtkm::cont::ArrayHandleImplicit<Functor>>
template <typename Functor>
struct SerializableTypeString<vtkm::cont::ArrayHandle<
typename vtkm::cont::detail::ArrayHandleImplicitTraits<Functor>::ValueType,
vtkm::cont::StorageTagImplicit<vtkm::cont::detail::ArrayPortalImplicit<Functor>>>>
vtkm::cont::StorageTagImplicit<vtkm::internal::ArrayPortalImplicit<Functor>>>>
: SerializableTypeString<vtkm::cont::ArrayHandleImplicit<Functor>>
{
};
@ -183,7 +331,7 @@ public:
template <typename Functor>
struct Serialization<vtkm::cont::ArrayHandle<
typename vtkm::cont::detail::ArrayHandleImplicitTraits<Functor>::ValueType,
vtkm::cont::StorageTagImplicit<vtkm::cont::detail::ArrayPortalImplicit<Functor>>>>
vtkm::cont::StorageTagImplicit<vtkm::internal::ArrayPortalImplicit<Functor>>>>
: Serialization<vtkm::cont::ArrayHandleImplicit<Functor>>
{
};

35
vtkm/cont/ArrayHandleIndex.h
View File

@ -14,6 +14,17 @@
namespace vtkm
{
namespace internal
{
struct VTKM_ALWAYS_EXPORT IndexFunctor
{
VTKM_EXEC_CONT vtkm::Id operator()(vtkm::Id index) const { return index; }
};
} // namespace internal
namespace cont
{
@ -24,32 +35,19 @@ struct VTKM_ALWAYS_EXPORT StorageTagIndex
namespace internal
{
struct VTKM_ALWAYS_EXPORT IndexFunctor
{
VTKM_EXEC_CONT
vtkm::Id operator()(vtkm::Id index) const { return index; }
};
using StorageTagIndexSuperclass =
typename vtkm::cont::ArrayHandleImplicit<IndexFunctor>::StorageTag;
typename vtkm::cont::ArrayHandleImplicit<vtkm::internal::IndexFunctor>::StorageTag;
template <>
struct Storage<vtkm::Id, vtkm::cont::StorageTagIndex> : Storage<vtkm::Id, StorageTagIndexSuperclass>
{
using Superclass = Storage<vtkm::Id, StorageTagIndexSuperclass>;
using Superclass::Superclass;
};
template <typename Device>
struct ArrayTransfer<vtkm::Id, vtkm::cont::StorageTagIndex, Device>
: ArrayTransfer<vtkm::Id, StorageTagIndexSuperclass, Device>
{
using Superclass = ArrayTransfer<vtkm::Id, StorageTagIndexSuperclass, Device>;
using Superclass::Superclass;
};
} // namespace internal
template <>
VTKM_ARRAY_HANDLE_NEW_STYLE(vtkm::Id, vtkm::cont::StorageTagIndex);
/// \brief An implicit array handle containing the its own indices.
///
/// \c ArrayHandleIndex is an implicit array handle containing the values
@ -65,8 +63,7 @@ public:
VTKM_CONT
ArrayHandleIndex(vtkm::Id length)
: Superclass(
typename internal::Storage<vtkm::Id, StorageTagIndex>::PortalType(internal::IndexFunctor{},
length))
internal::FunctorToArrayHandleImplicitBuffers(vtkm::internal::IndexFunctor{}, length))
{
}
};

20
vtkm/cont/ArrayHandleUniformPointCoordinates.h
View File

@ -10,8 +10,7 @@
#ifndef vtk_m_cont_ArrayHandleUniformPointCoordinates_h
#define vtk_m_cont_ArrayHandleUniformPointCoordinates_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageImplicit.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/internal/ArrayPortalUniformPointCoordinates.h>
namespace vtkm
@ -33,22 +32,13 @@ template <>
struct Storage<vtkm::Vec3f, vtkm::cont::StorageTagUniformPoints>
: Storage<vtkm::Vec3f, StorageTagUniformPointsSuperclass>
{
using Superclass = Storage<vtkm::Vec3f, StorageTagUniformPointsSuperclass>;
using Superclass::Superclass;
};
template <typename Device>
struct ArrayTransfer<vtkm::Vec3f, vtkm::cont::StorageTagUniformPoints, Device>
: ArrayTransfer<vtkm::Vec3f, StorageTagUniformPointsSuperclass, Device>
{
using Superclass = ArrayTransfer<vtkm::Vec3f, StorageTagUniformPointsSuperclass, Device>;
using Superclass::Superclass;
};
} // namespace internal
template <>
VTKM_ARRAY_HANDLE_NEW_STYLE(vtkm::Vec3f, vtkm::cont::StorageTagUniformPoints);
/// ArrayHandleUniformPointCoordinates is a specialization of ArrayHandle. It
/// contains the information necessary to compute the point coordinates in a
/// uniform orthogonal grid (extent, origin, and spacing) and implicitly
@ -70,7 +60,7 @@ public:
ArrayHandleUniformPointCoordinates(vtkm::Id3 dimensions,
ValueType origin = ValueType(0.0f, 0.0f, 0.0f),
ValueType spacing = ValueType(1.0f, 1.0f, 1.0f))
: Superclass(StorageType(
: Superclass(internal::PortalToArrayHandleImplicitBuffers(
vtkm::internal::ArrayPortalUniformPointCoordinates(dimensions, origin, spacing)))
{
}

20
vtkm/cont/BitField.cxx
View File

@ -41,7 +41,7 @@ namespace cont
namespace detail
{
vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData(
vtkm::cont::internal::BufferMetaDataBitField* GetBitFieldMetaData(
const vtkm::cont::internal::Buffer& buffer)
{
vtkm::cont::internal::BufferMetaData* generalMetaData = buffer.GetMetaData();
@ -49,22 +49,22 @@ vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData(
{
VTKM_LOG_F(vtkm::cont::LogLevel::Warn, "BitField has buffer with no metadata. Setting.");
const_cast<vtkm::cont::internal::Buffer&>(buffer).SetMetaData(
vtkm::cont::internal::BitFieldMetaData{});
vtkm::cont::internal::BufferMetaDataBitField{});
generalMetaData = buffer.GetMetaData();
VTKM_ASSERT(generalMetaData != nullptr);
}
vtkm::cont::internal::BitFieldMetaData* metadata =
dynamic_cast<vtkm::cont::internal::BitFieldMetaData*>(generalMetaData);
vtkm::cont::internal::BufferMetaDataBitField* metadata =
dynamic_cast<vtkm::cont::internal::BufferMetaDataBitField*>(generalMetaData);
if (metadata == nullptr)
{
VTKM_LOG_F(vtkm::cont::LogLevel::Error,
"BitField has a buffer with metadata of the wrong type. "
"Replacing, but this will likely cause problems.");
const_cast<vtkm::cont::internal::Buffer&>(buffer).SetMetaData(
vtkm::cont::internal::BitFieldMetaData{});
vtkm::cont::internal::BufferMetaDataBitField{});
generalMetaData = buffer.GetMetaData();
metadata = dynamic_cast<vtkm::cont::internal::BitFieldMetaData*>(generalMetaData);
metadata = dynamic_cast<vtkm::cont::internal::BufferMetaDataBitField*>(generalMetaData);
VTKM_ASSERT(metadata != nullptr);
}
@ -79,18 +79,18 @@ vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData(
namespace internal
{
BitFieldMetaData::~BitFieldMetaData() {}
BufferMetaDataBitField::~BufferMetaDataBitField() {}
std::unique_ptr<vtkm::cont::internal::BufferMetaData> BitFieldMetaData::DeepCopy() const
std::unique_ptr<vtkm::cont::internal::BufferMetaData> BufferMetaDataBitField::DeepCopy() const
{
return std::unique_ptr<vtkm::cont::internal::BufferMetaData>(new BitFieldMetaData(*this));
return std::unique_ptr<vtkm::cont::internal::BufferMetaData>(new BufferMetaDataBitField(*this));
}
} // namespace internal
BitField::BitField()
{
this->Buffer.SetMetaData(internal::BitFieldMetaData{});
this->Buffer.SetMetaData(internal::BufferMetaDataBitField{});
}
vtkm::Id BitField::GetNumberOfBits() const

6
vtkm/cont/BitField.h
View File

@ -36,11 +36,11 @@ namespace internal
struct StorageTagBitField;
struct VTKM_CONT_EXPORT BitFieldMetaData : vtkm::cont::internal::BufferMetaData
struct VTKM_CONT_EXPORT BufferMetaDataBitField : vtkm::cont::internal::BufferMetaData
{
vtkm::Id NumberOfBits = 0;
VTKM_CONT ~BitFieldMetaData() override;
VTKM_CONT ~BufferMetaDataBitField() override;
VTKM_CONT std::unique_ptr<vtkm::cont::internal::BufferMetaData> DeepCopy() const override;
};
@ -99,7 +99,7 @@ struct BitCoordinate
vtkm::Int32 BitOffset; // [0, bitsInWord)
};
VTKM_CONT_EXPORT vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData(
VTKM_CONT_EXPORT vtkm::cont::internal::BufferMetaDataBitField* GetBitFieldMetaData(
const vtkm::cont::internal::Buffer& buffer);
/// Portal for performing bit or word operations on a BitField.

2
vtkm/cont/CMakeLists.txt
View File

@ -109,7 +109,6 @@ set(headers
Serialization.h
Storage.h
StorageExtrude.h
StorageImplicit.h
StorageList.h
StorageListTag.h
Timer.h
@ -137,6 +136,7 @@ set(template_sources
set(sources
ArrayHandle.cxx
ArrayHandleBasic.cxx
ArrayHandleImplicit.cxx
ArrayHandleSOA.cxx
BitField.cxx
ColorTablePresets.cxx

177
vtkm/cont/StorageImplicit.h
View File

@ -1,177 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_cont_StorageImplicit
#define vtk_m_cont_StorageImplicit
#include <vtkm/Types.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/Storage.h>
#include <vtkm/cont/Token.h>
#include <vtkm/cont/internal/ArrayTransfer.h>
namespace vtkm
{
namespace cont
{
/// \brief An implementation for read-only implicit arrays.
///
/// It is sometimes the case that you want VTK-m to operate on an array of
/// implicit values. That is, rather than store the data in an actual array, it
/// is gerenated on the fly by a function. This is handled in VTK-m by creating
/// an ArrayHandle in VTK-m with a StorageTagImplicit type of \c Storage. This
/// tag itself is templated to specify an ArrayPortal that generates the
/// desired values. An ArrayHandle created with this tag will raise an error on
/// any operation that tries to modify it.
///
template <class ArrayPortalType>
struct VTKM_ALWAYS_EXPORT StorageTagImplicit
{
using PortalType = ArrayPortalType;
};
namespace internal
{
template <class ArrayPortalType>
class VTKM_ALWAYS_EXPORT
Storage<typename ArrayPortalType::ValueType, StorageTagImplicit<ArrayPortalType>>
{
using ClassType =
Storage<typename ArrayPortalType::ValueType, StorageTagImplicit<ArrayPortalType>>;
public:
using ValueType = typename ArrayPortalType::ValueType;
using PortalConstType = ArrayPortalType;
// Note that this portal is likely to be read-only, so you will probably get an error
// if you try to write to it.
using PortalType = ArrayPortalType;
VTKM_CONT
Storage(const PortalConstType& portal = PortalConstType())
: Portal(portal)
, NumberOfValues(portal.GetNumberOfValues())
{
}
VTKM_CONT Storage(const ClassType&) = default;
VTKM_CONT Storage(ClassType&&) = default;
VTKM_CONT ClassType& operator=(const ClassType&) = default;
VTKM_CONT ClassType& operator=(ClassType&&) = default;
// All these methods do nothing but raise errors.
VTKM_CONT
PortalType GetPortal() { throw vtkm::cont::ErrorBadValue("Implicit arrays are read-only."); }
VTKM_CONT
PortalConstType GetPortalConst() const { return this->Portal; }
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
VTKM_CONT
void Allocate(vtkm::Id numberOfValues)
{
VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues());
this->NumberOfValues = numberOfValues;
}
VTKM_CONT
void Shrink(vtkm::Id numberOfValues)
{
VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues());
this->NumberOfValues = numberOfValues;
}
VTKM_CONT
void ReleaseResources() {}
private:
PortalConstType Portal;
vtkm::Id NumberOfValues;
};
template <typename T, class ArrayPortalType, class DeviceAdapterTag>
class ArrayTransfer<T, StorageTagImplicit<ArrayPortalType>, DeviceAdapterTag>
{
public:
using StorageTag = StorageTagImplicit<ArrayPortalType>;
using StorageType = vtkm::cont::internal::Storage<T, StorageTag>;
using ValueType = T;
using PortalControl = typename StorageType::PortalType;
using PortalConstControl = typename StorageType::PortalConstType;
using PortalExecution = PortalControl;
using PortalConstExecution = PortalConstControl;
VTKM_CONT
ArrayTransfer(StorageType* storage)
: Storage(storage)
{
}
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->Storage->GetNumberOfValues(); }
VTKM_CONT
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData), vtkm::cont::Token&)
{
return this->Storage->GetPortalConst();
}
#if defined(VTKM_GCC) && defined(VTKM_ENABLE_OPENMP) && (__GNUC__ == 6 && __GNUC_MINOR__ == 1)
// When using GCC 6.1 with OpenMP enabled we cause a compiler ICE that is
// an identified compiler regression (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=71210)
// The easiest way to work around this is to make sure we aren't building with >= O2
#define NO_OPTIMIZE_FUNC_ATTRIBUTE __attribute__((optimize(1)))
#else // gcc 6.1 openmp compiler ICE workaround
#define NO_OPTIMIZE_FUNC_ATTRIBUTE
#endif
VTKM_CONT
NO_OPTIMIZE_FUNC_ATTRIBUTE
PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData), vtkm::cont::Token&)
{
throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output or in place.");
}
VTKM_CONT
NO_OPTIMIZE_FUNC_ATTRIBUTE
PortalExecution PrepareForOutput(vtkm::Id vtkmNotUsed(numberOfValues), vtkm::cont::Token&)
{
throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output.");
}
#undef NO_OPTIMIZE_FUNC_ATTRIBUTE
VTKM_CONT
void RetrieveOutputData(StorageType* vtkmNotUsed(controlArray)) const
{
throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output.");
}
VTKM_CONT
void Shrink(vtkm::Id vtkmNotUsed(numberOfValues))
{
throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be resized.");
}
VTKM_CONT
void ReleaseResources() {}
private:
StorageType* Storage;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_StorageImplicit

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

@ -1639,6 +1639,8 @@ public:
static void ScheduleTask(vtkm::exec::cuda::internal::TaskStrided1D<WType, IType>& functor,
vtkm::Id numInstances)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
VTKM_ASSERT(numInstances >= 0);
if (numInstances < 1)
{
@ -1671,6 +1673,8 @@ public:
static void ScheduleTask(vtkm::exec::cuda::internal::TaskStrided3D<WType, IType>& functor,
vtkm::Id3 rangeMax)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
VTKM_ASSERT((rangeMax[0] >= 0) && (rangeMax[1] >= 0) && (rangeMax[2] >= 0));
if ((rangeMax[0] < 1) || (rangeMax[1] < 1) || (rangeMax[2] < 1))
{

4
vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h
View File

@ -165,6 +165,8 @@ public:
vtkm::exec::kokkos::internal::TaskBasic1D<WType, IType>& functor,
vtkm::Id numInstances)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
if (numInstances < 1)
{
// No instances means nothing to run. Just return.
@ -187,6 +189,8 @@ public:
vtkm::exec::kokkos::internal::TaskBasic3D<WType, IType>& functor,
vtkm::Id3 rangeMax)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
if ((rangeMax[0] < 1) || (rangeMax[1] < 1) || (rangeMax[2] < 1))
{
// No instances means nothing to run. Just return.

4
vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx
View File

@ -24,6 +24,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagOpenMP>::ScheduleTask(
vtkm::exec::openmp::internal::TaskTiling1D& functor,
vtkm::Id size)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
static constexpr vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE];
errorString[0] = '\0';
@ -66,6 +68,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagOpenMP>::ScheduleTask(
vtkm::exec::openmp::internal::TaskTiling3D& functor,
vtkm::Id3 size)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
static constexpr vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE];
errorString[0] = '\0';

4
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx
View File

@ -19,6 +19,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
vtkm::exec::tbb::internal::TaskTiling1D& functor,
vtkm::Id size)
{
VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, "Schedule Task TBB 1D");
const vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE];
errorString[0] = '\0';
@ -40,6 +42,8 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
vtkm::exec::tbb::internal::TaskTiling3D& functor,
vtkm::Id3 size)
{
VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, "Schedule Task TBB 3D");
static constexpr vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 };
const vtkm::Id MESSAGE_SIZE = 1024;
char errorString[MESSAGE_SIZE];

10
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
View File

@ -286,6 +286,8 @@ public:
template <typename T, class Container>
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Container>& values)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
//this is required to get sort to work with zip handles
std::less<T> lessOp;
vtkm::cont::tbb::sort::parallel_sort(values, lessOp);
@ -295,6 +297,8 @@ public:
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Container>& values,
BinaryCompare binary_compare)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::tbb::sort::parallel_sort(values, binary_compare);
}
@ -302,6 +306,8 @@ public:
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::tbb::sort::parallel_sort_bykey(keys, values, std::less<T>());
}
@ -310,6 +316,8 @@ public:
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::tbb::sort::parallel_sort_bykey(keys, values, binary_compare);
}
@ -323,6 +331,8 @@ public:
VTKM_CONT static void Unique(vtkm::cont::ArrayHandle<T, Storage>& values,
BinaryCompare binary_compare)
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::Id outputSize;
{
vtkm::cont::Token token;

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

@ -81,7 +81,6 @@ set(unit_tests
UnitTestRuntimeDeviceInformation.cxx
UnitTestRuntimeDeviceNames.cxx
UnitTestScopedRuntimeDeviceTracker.cxx
UnitTestStorageImplicit.cxx
UnitTestStorageList.cxx
UnitTestStorageListTag.cxx
UnitTestTimer.cxx

55
vtkm/cont/testing/UnitTestArrayHandleCounting.cxx
View File

@ -18,57 +18,8 @@ namespace UnitTestArrayHandleCountingNamespace
const vtkm::Id ARRAY_SIZE = 10;
// An unusual data type that represents a number with a string of a
// particular length. This makes sure that the ArrayHandleCounting
// works correctly with type casts.
class StringInt
{
public:
StringInt() {}
StringInt(vtkm::Id v)
{
VTKM_ASSERT(v >= 0);
for (vtkm::Id i = 0; i < v; i++)
{
++(*this);
}
}
operator vtkm::Id() const { return vtkm::Id(this->Value.size()); }
StringInt operator+(const StringInt& rhs) const { return StringInt(this->Value + rhs.Value); }
StringInt operator*(const StringInt& rhs) const
{
StringInt result;
for (vtkm::Id i = 0; i < rhs; i++)
{
result = result + *this;
}
return result;
}
bool operator==(const StringInt& other) const { return this->Value.size() == other.Value.size(); }
StringInt& operator++()
{
this->Value.append(".");
return *this;
}
private:
StringInt(const std::string& v)
: Value(v)
{
}
std::string Value;
};
} // namespace UnitTestArrayHandleCountingNamespace
VTKM_BASIC_TYPE_VECTOR(UnitTestArrayHandleCountingNamespace::StringInt)
namespace UnitTestArrayHandleCountingNamespace
{
@ -81,7 +32,7 @@ struct TemplatedTests
using PortalType =
typename vtkm::cont::internal::Storage<ValueType,
typename ArrayHandleType::StorageTag>::PortalConstType;
typename ArrayHandleType::StorageTag>::ReadPortalType;
void operator()(const ValueType& startingValue, const ValueType& step)
{
@ -90,7 +41,7 @@ struct TemplatedTests
ArrayHandleType arrayMake =
vtkm::cont::make_ArrayHandleCounting(startingValue, step, ARRAY_SIZE);
ArrayHandleType2 arrayHandle = ArrayHandleType2(PortalType(startingValue, step, ARRAY_SIZE));
ArrayHandleType2 arrayHandle = ArrayHandleType(startingValue, step, ARRAY_SIZE);
VTKM_TEST_ASSERT(arrayConst.GetNumberOfValues() == ARRAY_SIZE,
"Counting array using constructor has wrong size.");
@ -127,8 +78,6 @@ void TestArrayHandleCounting()
TemplatedTests<vtkm::Float32>()(3.0f, -0.5f);
TemplatedTests<vtkm::Float64>()(0.0, 1.0);
TemplatedTests<vtkm::Float64>()(-3.0, 2.0);
TemplatedTests<StringInt>()(StringInt(0), StringInt(1));
TemplatedTests<StringInt>()(StringInt(10), StringInt(2));
}

134
vtkm/cont/testing/UnitTestStorageImplicit.cxx
View File

@ -1,134 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/Types.h>
#include <vtkm/VecTraits.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/StorageImplicit.h>
#include <vtkm/cont/internal/IteratorFromArrayPortal.h>
#include <vtkm/cont/testing/Testing.h>
#if defined(VTKM_STORAGE)
#undef VTKM_STORAGE
#endif
#define VTKM_STORAGE VTKM_STORAGE_ERROR
namespace
{
const vtkm::Id ARRAY_SIZE = 10;
template <typename T>
struct TestImplicitStorage
{
using ValueType = T;
ValueType Temp;
VTKM_EXEC_CONT
TestImplicitStorage()
: Temp(1)
{
}
VTKM_EXEC_CONT
vtkm::Id GetNumberOfValues() const { return ARRAY_SIZE; }
VTKM_EXEC_CONT
ValueType Get(vtkm::Id vtkmNotUsed(index)) const { return Temp; }
};
template <typename T>
struct TemplatedTests
{
using StorageTagType = vtkm::cont::StorageTagImplicit<TestImplicitStorage<T>>;
using StorageType = vtkm::cont::internal::Storage<T, StorageTagType>;
using ValueType = typename StorageType::ValueType;
using PortalType = typename StorageType::PortalType;
void BasicAllocation()
{
StorageType arrayStorage;
// The implicit portal defined for this test always returns ARRAY_SIZE for the
// number of values. We should get that.
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE,
"Implicit Storage GetNumberOfValues returned wrong size.");
// Make sure you can allocate and shrink to any value <= the reported portal size.
arrayStorage.Allocate(ARRAY_SIZE / 2);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE / 2,
"Cannot re-Allocate array to half size.");
arrayStorage.Allocate(0);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == 0, "Cannot re-Allocate array to zero.");
arrayStorage.Allocate(ARRAY_SIZE);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE,
"Cannot re-Allocate array to original size.");
arrayStorage.Shrink(ARRAY_SIZE / 2);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE / 2,
"Cannot Shrink array to half size.");
arrayStorage.Shrink(0);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == 0, "Cannot Shrink array to zero.");
arrayStorage.Shrink(ARRAY_SIZE);
VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE,
"Cannot Shrink array to original size.");
//verify that calling ReleaseResources doesn't throw an exception
arrayStorage.ReleaseResources();
//verify that you can allocate after releasing resources.
arrayStorage.Allocate(ARRAY_SIZE);
}
void BasicAccess()
{
TestImplicitStorage<T> portal;
vtkm::cont::ArrayHandle<T, StorageTagType> implictHandle(portal);
VTKM_TEST_ASSERT(implictHandle.GetNumberOfValues() == ARRAY_SIZE, "handle has wrong size");
VTKM_TEST_ASSERT(implictHandle.ReadPortal().Get(0) == T(1), "portals first values should be 1");
}
void operator()()
{
BasicAllocation();
BasicAccess();
}
};
struct TestFunctor
{
template <typename T>
void operator()(T) const
{
TemplatedTests<T> tests;
tests();
}
};
void TestStorageBasic()
{
vtkm::testing::Testing::TryTypes(TestFunctor());
}
} // Anonymous namespace
int UnitTestStorageImplicit(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestStorageBasic, argc, argv);
}

2
vtkm/filter/CMakeLists.txt
View File

@ -81,6 +81,7 @@ set(extra_headers
MeshQuality.h
NDEntropy.h
NDHistogram.h
ParticleDensityNearestGridPoint.h
ParticleAdvection.h
Pathline.h
PointAverage.h
@ -129,6 +130,7 @@ set(extra_header_template_sources
MeshQuality.hxx
NDEntropy.hxx
NDHistogram.hxx
ParticleDensityNearestGridPoint.hxx
ParticleAdvection.hxx
Pathline.hxx
PointAverage.hxx

12
vtkm/filter/Filter.hxx
View File

@ -249,9 +249,6 @@ inline VTKM_CONT Filter<Derived>::~Filter()
template <typename Derived>
inline VTKM_CONT vtkm::cont::DataSet Filter<Derived>::Execute(const vtkm::cont::DataSet& input)
{
VTKM_LOG_SCOPE(
vtkm::cont::LogLevel::Perf, "Filter: '%s'", vtkm::cont::TypeToString<Derived>().c_str());
Derived* self = static_cast<Derived*>(this);
vtkm::cont::PartitionedDataSet output = self->Execute(vtkm::cont::PartitionedDataSet(input));
if (output.GetNumberOfPartitions() > 1)
@ -267,7 +264,8 @@ inline VTKM_CONT vtkm::cont::PartitionedDataSet Filter<Derived>::Execute(
const vtkm::cont::PartitionedDataSet& input)
{
VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf,
"Filter (PartitionedDataSet): '%s'",
"Filter (%d partitions): '%s'",
(int)input.GetNumberOfPartitions(),
vtkm::cont::TypeToString<Derived>().c_str());
Derived* self = static_cast<Derived*>(this);
@ -293,9 +291,6 @@ VTKM_CONT vtkm::cont::DataSet Filter<Derived>::Execute(
const vtkm::cont::DataSet& input,
vtkm::filter::PolicyBase<DerivedPolicy> policy)
{
VTKM_LOG_SCOPE(
vtkm::cont::LogLevel::Perf, "Filter: '%s'", vtkm::cont::TypeToString<Derived>().c_str());
Derived* self = static_cast<Derived*>(this);
VTKM_DEPRECATED_SUPPRESS_BEGIN
vtkm::cont::PartitionedDataSet output =
@ -316,7 +311,8 @@ VTKM_CONT vtkm::cont::PartitionedDataSet Filter<Derived>::Execute(
vtkm::filter::PolicyBase<DerivedPolicy> policy)
{
VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf,
"Filter (PartitionedDataSet): '%s'",
"Filter (%d partitions): '%s'",
(int)input.GetNumberOfPartitions(),
vtkm::cont::TypeToString<Derived>().c_str());
Derived* self = static_cast<Derived*>(this);

53
vtkm/filter/ParticleDensityNearestGridPoint.h Normal file
View File

@ -0,0 +1,53 @@
//============================================================================
// 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.
//============================================================================
#ifndef vtk_m_filter_particle_density_ngp_h
#define vtk_m_filter_particle_density_ngp_h
#include <vtkm/filter/FilterField.h>
namespace vtkm
{
namespace filter
{
/// \brief Estimate the density of particles using the Nearest Grid Point method
// We only need the CoordinateSystem of the input dataset thus a FilterField
class ParticleDensityNearestGridPoint
: public vtkm::filter::FilterField<ParticleDensityNearestGridPoint>
{
public:
// ParticleDensity only support turning 2D/3D particle positions into density
using SupportedTypes = vtkm::TypeListFieldVec3;
//
ParticleDensityNearestGridPoint(const vtkm::Id3& dimension,
const vtkm::Vec3f& origin,
const vtkm::Vec3f& spacing);
ParticleDensityNearestGridPoint(const vtkm::Id3& dimension, const vtkm::Bounds& bounds);
template <typename T, typename StorageType, typename Policy>
VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input,
const vtkm::cont::ArrayHandle<T, StorageType>& field,
const vtkm::filter::FieldMetadata& fieldMeta,
vtkm::filter::PolicyBase<Policy> policy);
private:
vtkm::Id3 Dimension; // Cell dimension
vtkm::Vec3f Origin;
vtkm::Vec3f Spacing;
};
}
}
#include <vtkm/filter/ParticleDensityNearestGridPoint.hxx>
#endif //vtk_m_filter_particle_density_ngp_h

119
vtkm/filter/ParticleDensityNearestGridPoint.hxx Normal file
View File

@ -0,0 +1,119 @@
//============================================================================
// 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.
//============================================================================
#ifndef vtk_m_filter_particle_density_ngp_hxx
#define vtk_m_filter_particle_density_ngp_hxx
#include "ParticleDensityNearestGridPoint.h"
#include <vtkm/cont/CellLocatorUniformGrid.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/filter/PolicyBase.h>
#include <vtkm/worklet/WorkletMapField.h>
namespace vtkm
{
namespace worklet
{
class NGPWorklet : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(FieldIn coords, ExecObject locator, AtomicArrayInOut density);
using ExecutionSignature = void(_1, _2, _3);
template <typename Point, typename CellLocatorExecObj, typename AtomicArray>
VTKM_EXEC void operator()(const Point& point,
const CellLocatorExecObj& locator,
AtomicArray& density) const
{
vtkm::Id cellId{};
vtkm::Vec3f parametric;
// Find the cell containing the point
if (locator->FindCell(point, cellId, parametric) == vtkm::ErrorCode::Success)
{
// increment density
density.Add(cellId, 1);
}
// FIXME: what does mean when it is not found?
// We simply ignore that particular particle.
}
}; //NGPWorklet
} //worklet
} //vtkm
namespace vtkm
{
namespace filter
{
inline VTKM_CONT ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoint(
const vtkm::Id3& dimension,
const vtkm::Vec3f& origin,
const vtkm::Vec3f& spacing)
: Dimension(dimension)
, Origin(origin)
, Spacing(spacing)
{
}
ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoint(const Id3& dimension,
const vtkm::Bounds& bounds)
: Dimension(dimension)
, Origin({ static_cast<vtkm::FloatDefault>(bounds.X.Min),
static_cast<vtkm::FloatDefault>(bounds.Y.Min),
static_cast<vtkm::FloatDefault>(bounds.Z.Min) })
, Spacing(vtkm::Vec3f{ static_cast<vtkm::FloatDefault>(bounds.X.Length()),
static_cast<vtkm::FloatDefault>(bounds.Y.Length()),
static_cast<vtkm::FloatDefault>(bounds.Z.Length()) } /
Dimension)
{
}
template <typename T, typename StorageType, typename Policy>
inline VTKM_CONT vtkm::cont::DataSet ParticleDensityNearestGridPoint::DoExecute(
const vtkm::cont::DataSet&,
const vtkm::cont::ArrayHandle<T, StorageType>& field,
const vtkm::filter::FieldMetadata&,
vtkm::filter::PolicyBase<Policy>)
{
// TODO: it really doesn't need to be a UniformGrid, any CellSet with CellLocator will work.
// Make it another input rather an output generated.
// We stores density as CellField which conforms to physicists' idea of particle density
// better. However, VTK/VTKm's idea of "Image" Dataset and the ImageConnectivity filter
// expect a PointField. For better separation of concerns, we create a uniform dataset
// that has the cell dimension as expected and later convert the dataset to its dual.
auto uniform = vtkm::cont::DataSetBuilderUniform::Create(
this->Dimension + vtkm::Id3{ 1, 1, 1 }, this->Origin, this->Spacing);
// Create a CellLocator
vtkm::cont::CellLocatorUniformGrid locator;
locator.SetCellSet(uniform.GetCellSet());
locator.SetCoordinates(uniform.GetCoordinateSystem());
locator.Update();
// We create an ArrayHandle and pass it to the Worklet as AtomicArrayInOut.
// However the ArrayHandle needs to be allocated and initialized first. The
// easily way to do it is to copy from an ArrayHandleConstant
vtkm::cont::ArrayHandle<vtkm::Id> density;
vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleConstant<vtkm::Id>(0, uniform.GetNumberOfCells()),
density);
this->Invoke(vtkm::worklet::NGPWorklet{}, field, locator, density);
uniform.AddField(vtkm::cont::make_FieldCell("density", density));
return uniform;
}
}
}
#endif //vtk_m_filter_particle_density_ngp_hxx

2
vtkm/filter/testing/CMakeLists.txt
View File

@ -46,6 +46,8 @@ set(unit_tests
UnitTestMeshQualityFilter.cxx
UnitTestNDEntropyFilter.cxx
UnitTestNDHistogramFilter.cxx
UnitTestParticleAdvectionFilter.cxx
UnitTestParticleDensity.cxx
UnitTestPartitionedDataSetFilters.cxx
UnitTestPartitionedDataSetHistogramFilter.cxx
UnitTestPointAverageFilter.cxx

57
vtkm/filter/testing/UnitTestParticleDensity.cxx Normal file
View File

@ -0,0 +1,57 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/ArrayHandleRandomUniformReal.h>
#include <vtkm/cont/DataSetBuilderExplicit.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/filter/ParticleDensityNearestGridPoint.h>
#include <vtkm/worklet/DescriptiveStatistics.h>
void TestNGP()
{
const vtkm::Id N = 1000000;
auto composite = vtkm::cont::make_ArrayHandleCompositeVector(
vtkm::cont::ArrayHandleRandomUniformReal<vtkm::Float32>(N, 0xceed),
vtkm::cont::ArrayHandleRandomUniformReal<vtkm::Float32>(N, 0xdeed),
vtkm::cont::ArrayHandleRandomUniformReal<vtkm::Float32>(N, 0xabba));
vtkm::cont::ArrayHandle<vtkm::Vec3f> positions;
vtkm::cont::ArrayCopy(composite, positions);
vtkm::cont::ArrayHandle<vtkm::Id> connectivity;
vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleIndex(N), connectivity);
auto dataSet = vtkm::cont::DataSetBuilderExplicit::Create(
positions, vtkm::CellShapeTagVertex{}, 1, connectivity);
auto cellDims = vtkm::Id3{ 3, 3, 3 };
vtkm::filter::ParticleDensityNearestGridPoint filter{
cellDims, { 0.f, 0.f, 0.f }, vtkm::Vec3f{ 1.f / 3.f, 1.f / 3.f, 1.f / 3.f }
};
filter.SetUseCoordinateSystemAsField(true);
auto density = filter.Execute(dataSet);
vtkm::cont::ArrayHandle<vtkm::Id> field;
density.GetCellField("density").GetData().AsArrayHandle<vtkm::Id>(field);
auto field_f = vtkm::cont::make_ArrayHandleCast<vtkm::Float32>(field);
auto result = vtkm::worklet::DescriptiveStatistics::Run(field_f);
VTKM_TEST_ASSERT(test_equal(result.Sum(), N));
VTKM_TEST_ASSERT(test_equal(result.Mean(), N / density.GetNumberOfCells()));
}
void TestParticleDensity()
{
TestNGP();
}
int UnitTestParticleDensity(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestParticleDensity, argc, argv);
}

19
vtkm/internal/ArrayPortalUniformPointCoordinates.h
View File

@ -43,25 +43,6 @@ public:
{
}
VTKM_EXEC_CONT
ArrayPortalUniformPointCoordinates(const ArrayPortalUniformPointCoordinates& src)
: Dimensions(src.Dimensions)
, NumberOfValues(src.NumberOfValues)
, Origin(src.Origin)
, Spacing(src.Spacing)
{
}
VTKM_EXEC_CONT
ArrayPortalUniformPointCoordinates& operator=(const ArrayPortalUniformPointCoordinates& src)
{
this->Dimensions = src.Dimensions;
this->NumberOfValues = src.NumberOfValues;
this->Origin = src.Origin;
this->Spacing = src.Spacing;
return *this;
}
VTKM_EXEC_CONT
vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }

6
vtkm/testing/CMakeLists.txt
View File

@ -55,11 +55,9 @@ set(unit_tests_device
)
#suppress gcc note:
#variable tracking size limit exceeded with -fvar-tracking-assignments, retrying without
if (VTKM_COMPILER_IS_GNU)
set_source_files_properties(UnitTestBounds.cxx PROPERTIES COMPILE_OPTIONS "-fno-var-tracking-assignments")
endif()
set_source_files_properties(UnitTestBounds.cxx PROPERTIES
COMPILE_OPTIONS "$<$<AND:$<COMPILE_LANGUAGE:CXX>,$<CXX_COMPILER_ID:GNU>>:-fno-var-tracking-assignments>")
vtkm_unit_tests(SOURCES ${unit_tests})

2
vtkm/thirdparty/diy/update.sh vendored
View File

@ -8,7 +8,7 @@ readonly name="diy"
readonly ownership="Diy Upstream <kwrobot@kitware.com>"
readonly subtree="vtkm/thirdparty/$name/vtkm$name"
readonly repo="https://gitlab.kitware.com/third-party/diy2.git"
readonly tag="for/vtk-m-20200826-master"
readonly tag="for/vtk-m-20201026-master"
readonly paths="
cmake
include

4
vtkm/thirdparty/diy/vtkmdiy/include/vtkmdiy/master.hpp vendored
View File

@ -166,7 +166,7 @@ namespace diy
ExternalStorage* storage = 0, //!< storage object (path, method, etc.) for storing temporary blocks being shuffled in/out of core
SaveBlock save = 0, //!< block save function; master manages saving if save != 0
LoadBlock load_ = 0, //!< block load function; master manages loading if load != 0
QueuePolicy* q_policy = new QueueSizePolicy(4096)); //!< policy for managing message queues specifies maximum size of message queues to keep in memory
QueuePolicy* q_policy = nullptr); //!< policy for managing message queues specifies maximum size of message queues to keep in memory
inline ~Master();
inline void clear();
@ -370,7 +370,7 @@ Master(mpi::communicator comm,
LoadBlock load_,
QueuePolicy* q_policy):
blocks_(create_, destroy_, storage, save, load_),
queue_policy_(q_policy),
queue_policy_( (q_policy==nullptr) ? new QueueSizePolicy(4096): q_policy),
limit_(limit__),
#if !defined(VTKMDIY_NO_THREADS)
threads_(threads__ == -1 ? static_cast<int>(thread::hardware_concurrency()) : threads__),

5
vtkm/thirdparty/lcl/vtkmlcl/lcl/internal/Config.h vendored
View File

@ -13,7 +13,10 @@
#include <cstdint>
#include <type_traits>
#ifdef __CUDACC__
#if defined(__CUDACC__)
# define LCL_EXEC __device__ __host__
#elif defined(__HIP__)
#include "hip/hip_runtime.h" //required for __device__ __host__
# define LCL_EXEC __device__ __host__
#else
# define LCL_EXEC

5
vtkm/worklet/contour/FlyingEdgesPass1.h
View File

@ -141,6 +141,7 @@ struct ComputePass1 : public vtkm::worklet::WorkletVisitPointsWithCells
}
}
}
write_edge(device, startPos + (offset * end), edges, FlyingEdges3D::Below);
}
};
@ -150,7 +151,7 @@ struct launchComputePass1
VTKM_CONT bool operator()(DeviceAdapterTag device,
const ComputePass1<T>& worklet,
const vtkm::cont::ArrayHandle<T, StorageTagField>& inputField,
vtkm::cont::ArrayHandle<vtkm::UInt8> edgeCases,
vtkm::cont::ArrayHandle<vtkm::UInt8>& edgeCases,
vtkm::cont::CellSetStructured<2>& metaDataMesh2D,
Args&&... args) const
{
@ -165,7 +166,7 @@ struct launchComputePass1
VTKM_CONT bool operator()(vtkm::cont::DeviceAdapterTagCuda device,
const ComputePass1<T>& worklet,
const vtkm::cont::ArrayHandle<T, StorageTagField>& inputField,
vtkm::cont::ArrayHandle<vtkm::UInt8> edgeCases,
vtkm::cont::ArrayHandle<vtkm::UInt8>& edgeCases,
vtkm::cont::CellSetStructured<2>& metaDataMesh2D,
Args&&... args) const
{

2
vtkm/worklet/contour/FlyingEdgesPass4Common.h
View File

@ -171,7 +171,7 @@ struct Pass4TrimState
}
if (ijk[AxisToSum::xindex] >= (pdims[AxisToSum::xindex] - 2))
{
boundaryStatus[AxisToSum::yindex] += FlyingEdges3D::MaxBoundary;
boundaryStatus[AxisToSum::xindex] += FlyingEdges3D::MaxBoundary;
}
if (ijk[AxisToSum::yindex] < 1)
{