mirror of
https://gitlab.kitware.com/vtk/vtk-m
synced 2024-09-08 13:23:51 +00:00
Merge topic 'array-handle-transform-exec-object-2'
2b0548739 Add ExecutionAndControlObjectBase 62d3b9f4f Correct warning about potential divide by zero 98a0a20fe Allow ArrayHandleTransform to work with ExecObject Acked-by: Kitware Robot <kwrobot@kitware.com> Acked-by: Robert Maynard <robert.maynard@kitware.com> Merge-request: !1403
This commit is contained in:
Kenneth Moreland
Kitware Robot
commit
1acc610799
110
docs/changelog/array-handle-transform-exec-object.md
Normal file
110
docs/changelog/array-handle-transform-exec-object.md
Normal file
@ -0,0 +1,110 @@
|
||||
# Allow ArrayHandleTransform to work with ExecObject
|
||||
|
||||
Previously, the `ArrayHandleTransform` class only worked with plain old
|
||||
data (POD) objects as is functors. For simple transforms, this makes sense
|
||||
since all the data comes from a target `ArrayHandle` that will be sent to
|
||||
the device through a different path. However, this also requires the
|
||||
transform to be known at compile time.
|
||||
|
||||
However, there are cases where the functor cannot be a POD object and has
|
||||
to be built for a specific device. There are numerous reasons for this. One
|
||||
might be that you need some lookup tables. Another might be you want to
|
||||
support a virtual object, which has to be initialized for a particular
|
||||
device. The standard way to implement this in VTK-m is to create an
|
||||
"executive object." This actually means that we create a wrapper around
|
||||
executive objects that inherits from
|
||||
`vtkm::cont::ExecutionAndControlObjectBase` that contains a
|
||||
`PrepareForExecution` method and a `PrepareForControl` method.
|
||||
|
||||
As an example, consider the use case of a special `ArrayHandle` that takes
|
||||
the value in one array and returns the index of that value in another
|
||||
sorted array. We can do that by creating a functor that finds a value in an
|
||||
array and returns the index.
|
||||
|
||||
``` cpp
|
||||
template <typename ArrayPortalType>
|
||||
struct FindValueFunctor
|
||||
{
|
||||
ArrayPortalType SortedArrayPortal;
|
||||
|
||||
FindValueFunctor() = default;
|
||||
|
||||
VTKM_CONT FindValueFunctor(const ArrayPortalType& sortedPortal)
|
||||
: SortedArrayPortal(sortedPortal)
|
||||
{ }
|
||||
|
||||
VTKM_EXEC vtkm::Id operator()(const typename PortalType::ValueType& value)
|
||||
{
|
||||
vtkm::Id leftIndex = 0;
|
||||
vtkm::Id rightIndex = this->SortedArrayPortal.GetNubmerOfValues();
|
||||
while (leftIndex < rightIndex)
|
||||
{
|
||||
vtkm::Id middleIndex = (leftIndex + rightIndex) / 2;
|
||||
auto middleValue = this->SortedArrayPortal.Get(middleIndex);
|
||||
if (middleValue <= value)
|
||||
{
|
||||
rightIndex = middleValue;
|
||||
}
|
||||
else
|
||||
{
|
||||
leftIndex = middleValue + 1;
|
||||
}
|
||||
}
|
||||
return leftIndex;
|
||||
}
|
||||
};
|
||||
```
|
||||
|
||||
Simple enough, except that the type of `ArrayPortalType` depends on what
|
||||
device the functor runs on (not to mention its memory might need to be
|
||||
moved to different hardware). We can now solve this problem by creating a
|
||||
functor objecgt set this up for a device. `ArrayHandle`s also need to be
|
||||
able to provide portals that run in the control environment, and for that
|
||||
we need a special version of the functor for the control environment.
|
||||
|
||||
``` cpp
|
||||
template <typename ArrayHandleType>
|
||||
struct FindValueExecutionObject : vtkm::cont::ExecutionAndControlObjectBase
|
||||
{
|
||||
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
|
||||
|
||||
ArrayHandleType SortedArray;
|
||||
|
||||
FindValueExecutionObject() = default;
|
||||
|
||||
VTKM_CONT FindValueExecutionObject(const ArrayHandleType& sortedArray)
|
||||
: SortedArray(sortedArray)
|
||||
{ }
|
||||
|
||||
template <typename Device>
|
||||
VTKM_CONT
|
||||
FindValueFunctor<decltype(std::declval<FunctorType>()(Device()))>
|
||||
PrepareForExecution(Device device)
|
||||
{
|
||||
using FunctorType =
|
||||
FindValueFunctor<decltype(std::declval<FunctorType>()(Device()))>
|
||||
|
||||
return FunctorType(this->SortedArray.PrepareForInput(device));
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
FundValueFunctor<typename ArrayHandleType::PortalConstControl>
|
||||
PrepareForControl()
|
||||
{
|
||||
using FunctorType =
|
||||
FindValueFunctor<typename ArrayHandleType::PortalConstControl>
|
||||
|
||||
return FunctorType(this->SortedArray.GetPortalConstControl());
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
Now you can use this execution object in an `ArrayHandleTransform`. It will
|
||||
automatically be detected as an execution object and be converted to a
|
||||
functor in the execution environment.
|
||||
|
||||
``` cpp
|
||||
auto transformArray =
|
||||
vtkm::cont::make_ArrayHandleTransform(
|
||||
inputArray, FindValueExecutionObject<decltype(sortedArray)>(sortedArray));
|
||||
```
|
||||
16
docs/changelog/execution-and-control-objects.md
Normal file
16
docs/changelog/execution-and-control-objects.md
Normal file
@ -0,0 +1,16 @@
|
||||
# Add new execution and control objects
|
||||
|
||||
[Recent changes to execution objects](change-execution-object-creation.md)
|
||||
now have execution objects behave as factories that create an object
|
||||
specific for a particular device. Sometimes, you also need to be able to
|
||||
get an object that behaves properly in the control environment. For these
|
||||
cases, a sublcass to `vtkm::cont::ExecutionObjectBase` was created.
|
||||
|
||||
This subclass is called `vtkm::cont::ExecutionAndControlObjectBase`. In
|
||||
addition to the `PrepareForExecution` method required by its superclass,
|
||||
these objects also need to provide a `PrepareForControl` method to get an
|
||||
equivalent object that works in the control environment.
|
||||
|
||||
See [the changelog for execution objects in
|
||||
`ArrayHandleTransform`](array-handle-transform-exec-object.md) for an
|
||||
example of using a `vtkm::cont::ExecutionAndControlObjectBase`.
|
||||
@ -38,24 +38,25 @@ template <typename Device, typename T>
|
||||
inline auto DoPrepareArgForExec(T&& object, std::true_type)
|
||||
-> decltype(std::declval<T>().PrepareForExecution(Device()))
|
||||
{
|
||||
VTKM_IS_EXECUTION_OBJECT(T);
|
||||
return object.PrepareForExecution(Device{});
|
||||
}
|
||||
|
||||
template <typename Device, typename T>
|
||||
inline T&& DoPrepareArgForExec(T&& object, std::false_type)
|
||||
{
|
||||
static_assert(!vtkm::cont::internal::IsExecutionObjectBase<T>::value,
|
||||
"Internal error: failed to detect execution object.");
|
||||
return std::forward<T>(object);
|
||||
}
|
||||
|
||||
template <typename Device, typename T>
|
||||
auto PrepareArgForExec(T&& object) -> decltype(DoPrepareArgForExec<Device>(
|
||||
std::forward<T>(object),
|
||||
typename std::is_base_of<vtkm::cont::ExecutionObjectBase, typename std::decay<T>::type>::type{}))
|
||||
auto PrepareArgForExec(T&& object)
|
||||
-> decltype(DoPrepareArgForExec<Device>(std::forward<T>(object),
|
||||
vtkm::cont::internal::IsExecutionObjectBase<T>{}))
|
||||
{
|
||||
return DoPrepareArgForExec<Device>(
|
||||
std::forward<T>(object),
|
||||
typename std::is_base_of<vtkm::cont::ExecutionObjectBase,
|
||||
typename std::decay<T>::type>::type{});
|
||||
return DoPrepareArgForExec<Device>(std::forward<T>(object),
|
||||
vtkm::cont::internal::IsExecutionObjectBase<T>{});
|
||||
}
|
||||
|
||||
struct CopyFunctor
|
||||
|
||||
@ -25,6 +25,8 @@
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/ErrorBadType.h>
|
||||
#include <vtkm/cont/ErrorInternal.h>
|
||||
#include <vtkm/cont/ExecutionAndControlObjectBase.h>
|
||||
#include <vtkm/cont/RuntimeDeviceTracker.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
@ -178,18 +180,107 @@ namespace cont
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <typename ProvidedFunctorType, typename FunctorIsExecContObject>
|
||||
struct TransformFunctorManagerImpl;
|
||||
|
||||
template <typename ProvidedFunctorType>
|
||||
struct TransformFunctorManagerImpl<ProvidedFunctorType, std::false_type>
|
||||
{
|
||||
VTKM_STATIC_ASSERT_MSG(!vtkm::cont::internal::IsExecutionObjectBase<ProvidedFunctorType>::value,
|
||||
"Must use an ExecutionAndControlObject instead of an ExecutionObject.");
|
||||
|
||||
ProvidedFunctorType Functor;
|
||||
using FunctorType = ProvidedFunctorType;
|
||||
|
||||
TransformFunctorManagerImpl() = default;
|
||||
|
||||
VTKM_CONT
|
||||
TransformFunctorManagerImpl(const ProvidedFunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
ProvidedFunctorType PrepareForControl() const { return this->Functor; }
|
||||
|
||||
template <typename Device>
|
||||
VTKM_CONT ProvidedFunctorType PrepareForExecution(Device) const
|
||||
{
|
||||
return this->Functor;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename ProvidedFunctorType>
|
||||
struct TransformFunctorManagerImpl<ProvidedFunctorType, std::true_type>
|
||||
{
|
||||
VTKM_IS_EXECUTION_AND_CONTROL_OBJECT(ProvidedFunctorType);
|
||||
|
||||
ProvidedFunctorType Functor;
|
||||
// using FunctorType = decltype(std::declval<ProvidedFunctorType>().PrepareForControl());
|
||||
using FunctorType = decltype(Functor.PrepareForControl());
|
||||
|
||||
TransformFunctorManagerImpl() = default;
|
||||
|
||||
VTKM_CONT
|
||||
TransformFunctorManagerImpl(const ProvidedFunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
auto PrepareForControl() const -> decltype(this->Functor.PrepareForControl())
|
||||
{
|
||||
return this->Functor.PrepareForControl();
|
||||
}
|
||||
|
||||
template <typename Device>
|
||||
VTKM_CONT auto PrepareForExecution(Device device) const
|
||||
-> decltype(this->Functor.PrepareForExecution(device))
|
||||
{
|
||||
return this->Functor.PrepareForExecution(device);
|
||||
}
|
||||
};
|
||||
|
||||
template <typename ProvidedFunctorType>
|
||||
struct TransformFunctorManager
|
||||
: TransformFunctorManagerImpl<
|
||||
ProvidedFunctorType,
|
||||
typename vtkm::cont::internal::IsExecutionAndControlObjectBase<ProvidedFunctorType>::type>
|
||||
{
|
||||
using Superclass = TransformFunctorManagerImpl<
|
||||
ProvidedFunctorType,
|
||||
typename vtkm::cont::internal::IsExecutionAndControlObjectBase<ProvidedFunctorType>::type>;
|
||||
using FunctorType = typename Superclass::FunctorType;
|
||||
|
||||
VTKM_CONT TransformFunctorManager() = default;
|
||||
|
||||
VTKM_CONT TransformFunctorManager(const TransformFunctorManager& other) = default;
|
||||
|
||||
VTKM_CONT TransformFunctorManager(const ProvidedFunctorType& functor)
|
||||
: Superclass(functor)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename ValueType>
|
||||
using TransformedValueType = decltype(std::declval<FunctorType>()(ValueType{}));
|
||||
};
|
||||
|
||||
template <typename ArrayHandleType,
|
||||
typename FunctorType,
|
||||
typename InverseFunctorType = NullFunctorType>
|
||||
struct VTKM_ALWAYS_EXPORT StorageTagTransform
|
||||
{
|
||||
using ValueType = decltype(FunctorType{}(typename ArrayHandleType::ValueType{}));
|
||||
using FunctorManager = TransformFunctorManager<FunctorType>;
|
||||
using ValueType =
|
||||
typename FunctorManager::template TransformedValueType<typename ArrayHandleType::ValueType>;
|
||||
};
|
||||
|
||||
template <typename ArrayHandleType, typename FunctorType>
|
||||
class Storage<typename StorageTagTransform<ArrayHandleType, FunctorType>::ValueType,
|
||||
StorageTagTransform<ArrayHandleType, FunctorType>>
|
||||
{
|
||||
using FunctorManager = TransformFunctorManager<FunctorType>;
|
||||
|
||||
public:
|
||||
using ValueType = typename StorageTagTransform<ArrayHandleType, FunctorType>::ValueType;
|
||||
|
||||
@ -204,7 +295,7 @@ public:
|
||||
using PortalConstType =
|
||||
vtkm::exec::internal::ArrayPortalTransform<ValueType,
|
||||
typename ArrayHandleType::PortalConstControl,
|
||||
FunctorType>;
|
||||
typename FunctorManager::FunctorType>;
|
||||
|
||||
VTKM_CONT
|
||||
Storage()
|
||||
@ -223,15 +314,17 @@ public:
|
||||
VTKM_CONT
|
||||
PortalType GetPortal()
|
||||
{
|
||||
VTKM_ASSERT(this->Valid);
|
||||
return PortalType(this->Array.GetPortalControl(), this->Functor);
|
||||
throw vtkm::cont::ErrorBadType(
|
||||
"ArrayHandleTransform is read only. Cannot get writable portal.");
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
PortalConstType GetPortalConst() const
|
||||
{
|
||||
VTKM_ASSERT(this->Valid);
|
||||
return PortalConstType(this->Array.GetPortalConstControl(), this->Functor);
|
||||
vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope;
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial());
|
||||
return PortalConstType(this->Array.GetPortalConstControl(), this->Functor.PrepareForControl());
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
@ -269,11 +362,11 @@ public:
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
const FunctorType& GetFunctor() const { return this->Functor; }
|
||||
const FunctorManager& GetFunctor() const { return this->Functor; }
|
||||
|
||||
private:
|
||||
ArrayHandleType Array;
|
||||
FunctorType Functor;
|
||||
FunctorManager Functor;
|
||||
bool Valid;
|
||||
};
|
||||
|
||||
@ -282,6 +375,9 @@ class Storage<
|
||||
typename StorageTagTransform<ArrayHandleType, FunctorType, InverseFunctorType>::ValueType,
|
||||
StorageTagTransform<ArrayHandleType, FunctorType, InverseFunctorType>>
|
||||
{
|
||||
using FunctorManager = TransformFunctorManager<FunctorType>;
|
||||
using InverseFunctorManager = TransformFunctorManager<InverseFunctorType>;
|
||||
|
||||
public:
|
||||
using ValueType =
|
||||
typename StorageTagTransform<ArrayHandleType, FunctorType, InverseFunctorType>::ValueType;
|
||||
@ -289,13 +385,13 @@ public:
|
||||
using PortalType =
|
||||
vtkm::exec::internal::ArrayPortalTransform<ValueType,
|
||||
typename ArrayHandleType::PortalControl,
|
||||
FunctorType,
|
||||
InverseFunctorType>;
|
||||
typename FunctorManager::FunctorType,
|
||||
typename InverseFunctorManager::FunctorType>;
|
||||
using PortalConstType =
|
||||
vtkm::exec::internal::ArrayPortalTransform<ValueType,
|
||||
typename ArrayHandleType::PortalConstControl,
|
||||
FunctorType,
|
||||
InverseFunctorType>;
|
||||
typename FunctorManager::FunctorType,
|
||||
typename InverseFunctorManager::FunctorType>;
|
||||
|
||||
VTKM_CONT
|
||||
Storage()
|
||||
@ -318,15 +414,22 @@ public:
|
||||
PortalType GetPortal()
|
||||
{
|
||||
VTKM_ASSERT(this->Valid);
|
||||
return PortalType(this->Array.GetPortalControl(), this->Functor, this->InverseFunctor);
|
||||
vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope;
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial());
|
||||
return PortalType(this->Array.GetPortalControl(),
|
||||
this->Functor.PrepareForControl(),
|
||||
this->InverseFunctor.PrepareForControl());
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
PortalConstType GetPortalConst() const
|
||||
{
|
||||
VTKM_ASSERT(this->Valid);
|
||||
return PortalConstType(
|
||||
this->Array.GetPortalConstControl(), this->Functor, this->InverseFunctor);
|
||||
vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope;
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial());
|
||||
return PortalConstType(this->Array.GetPortalConstControl(),
|
||||
this->Functor.PrepareForControl(),
|
||||
this->InverseFunctor.PrepareForControl());
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
@ -361,15 +464,15 @@ public:
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
const FunctorType& GetFunctor() const { return this->Functor; }
|
||||
const FunctorManager& GetFunctor() const { return this->Functor; }
|
||||
|
||||
VTKM_CONT
|
||||
const InverseFunctorType& GetInverseFunctor() const { return this->InverseFunctor; }
|
||||
const InverseFunctorManager& GetInverseFunctor() const { return this->InverseFunctor; }
|
||||
|
||||
private:
|
||||
ArrayHandleType Array;
|
||||
FunctorType Functor;
|
||||
InverseFunctorType InverseFunctor;
|
||||
FunctorManager Functor;
|
||||
InverseFunctorManager InverseFunctor;
|
||||
bool Valid;
|
||||
};
|
||||
|
||||
@ -379,6 +482,7 @@ class ArrayTransfer<typename StorageTagTransform<ArrayHandleType, FunctorType>::
|
||||
Device>
|
||||
{
|
||||
using StorageTag = StorageTagTransform<ArrayHandleType, FunctorType>;
|
||||
using FunctorManager = TransformFunctorManager<FunctorType>;
|
||||
|
||||
public:
|
||||
using ValueType = typename StorageTagTransform<ArrayHandleType, FunctorType>::ValueType;
|
||||
@ -392,7 +496,7 @@ public:
|
||||
using PortalConstExecution = vtkm::exec::internal::ArrayPortalTransform<
|
||||
ValueType,
|
||||
typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst,
|
||||
FunctorType>;
|
||||
typename FunctorManager::FunctorType>;
|
||||
|
||||
VTKM_CONT
|
||||
ArrayTransfer(StorageType* storage)
|
||||
@ -407,7 +511,8 @@ public:
|
||||
VTKM_CONT
|
||||
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData))
|
||||
{
|
||||
return PortalConstExecution(this->Array.PrepareForInput(Device()), this->Functor);
|
||||
return PortalConstExecution(this->Array.PrepareForInput(Device()),
|
||||
this->Functor.PrepareForExecution(Device()));
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
@ -443,7 +548,7 @@ public:
|
||||
|
||||
private:
|
||||
ArrayHandleType Array;
|
||||
FunctorType Functor;
|
||||
FunctorManager Functor;
|
||||
};
|
||||
|
||||
template <typename ArrayHandleType,
|
||||
@ -456,6 +561,8 @@ class ArrayTransfer<
|
||||
Device>
|
||||
{
|
||||
using StorageTag = StorageTagTransform<ArrayHandleType, FunctorType, InverseFunctorType>;
|
||||
using FunctorManager = TransformFunctorManager<FunctorType>;
|
||||
using InverseFunctorManager = TransformFunctorManager<InverseFunctorType>;
|
||||
|
||||
public:
|
||||
using ValueType = typename StorageTagTransform<ArrayHandleType, FunctorType>::ValueType;
|
||||
@ -467,13 +574,13 @@ public:
|
||||
using PortalExecution = vtkm::exec::internal::ArrayPortalTransform<
|
||||
ValueType,
|
||||
typename ArrayHandleType::template ExecutionTypes<Device>::Portal,
|
||||
FunctorType,
|
||||
InverseFunctorType>;
|
||||
typename FunctorManager::FunctorType,
|
||||
typename InverseFunctorManager::FunctorType>;
|
||||
using PortalConstExecution = vtkm::exec::internal::ArrayPortalTransform<
|
||||
ValueType,
|
||||
typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst,
|
||||
FunctorType,
|
||||
InverseFunctorType>;
|
||||
typename FunctorManager::FunctorType,
|
||||
typename InverseFunctorManager::FunctorType>;
|
||||
|
||||
VTKM_CONT
|
||||
ArrayTransfer(StorageType* storage)
|
||||
@ -489,22 +596,25 @@ public:
|
||||
VTKM_CONT
|
||||
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData))
|
||||
{
|
||||
return PortalConstExecution(
|
||||
this->Array.PrepareForInput(Device()), this->Functor, this->InverseFunctor);
|
||||
return PortalConstExecution(this->Array.PrepareForInput(Device()),
|
||||
this->Functor.PrepareForExecution(Device()),
|
||||
this->InverseFunctor.PrepareForExecution(Device()));
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
PortalExecution PrepareForInPlace(bool& vtkmNotUsed(updateData))
|
||||
{
|
||||
return PortalExecution(
|
||||
this->Array.PrepareForInPlace(Device()), this->Functor, this->InverseFunctor);
|
||||
return PortalExecution(this->Array.PrepareForInPlace(Device()),
|
||||
this->Functor.PrepareForExecution(Device()),
|
||||
this->InverseFunctor.PrepareForExecution(Device()));
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
PortalExecution PrepareForOutput(vtkm::Id numberOfValues)
|
||||
{
|
||||
return PortalExecution(
|
||||
this->Array.PrepareForOutput(numberOfValues, Device()), this->Functor, this->InverseFunctor);
|
||||
return PortalExecution(this->Array.PrepareForOutput(numberOfValues, Device()),
|
||||
this->Functor.PrepareForExecution(Device()),
|
||||
this->InverseFunctor.PrepareForExecution(Device()));
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
@ -522,8 +632,8 @@ public:
|
||||
|
||||
private:
|
||||
ArrayHandleType Array;
|
||||
FunctorType Functor;
|
||||
InverseFunctorType InverseFunctor;
|
||||
FunctorManager Functor;
|
||||
InverseFunctorManager InverseFunctor;
|
||||
};
|
||||
|
||||
} // namespace internal
|
||||
|
||||
@ -83,6 +83,7 @@ set(headers
|
||||
ErrorFilterExecution.h
|
||||
ErrorExecution.h
|
||||
ErrorInternal.h
|
||||
ExecutionAndControlObjectBase.h
|
||||
ExecutionObjectBase.h
|
||||
Field.h
|
||||
FieldRangeCompute.h
|
||||
|
||||
82
vtkm/cont/ExecutionAndControlObjectBase.h
Normal file
82
vtkm/cont/ExecutionAndControlObjectBase.h
Normal 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 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
|
||||
// Copyright 2018 UT-Battelle, LLC.
|
||||
// Copyright 2018 Los Alamos National Security.
|
||||
//
|
||||
// Under the terms of Contract DE-NA0003525 with NTESS,
|
||||
// the U.S. Government retains certain rights in this software.
|
||||
//
|
||||
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
|
||||
// Laboratory (LANL), the U.S. Government retains certain rights in
|
||||
// this software.
|
||||
//============================================================================
|
||||
#ifndef vtk_m_cont_ExecutionAndControlObjectBase_h
|
||||
#define vtk_m_cont_ExecutionAndControlObjectBase_h
|
||||
|
||||
#include <vtkm/cont/ExecutionObjectBase.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
|
||||
/// Base \c ExecutionAndControlObjectBase class. These are objects that behave
|
||||
/// as execution objects but can also be use din the control environment.
|
||||
/// Any subclass of \c ExecutionAndControlObjectBase must implement a
|
||||
/// \c PrepareForExecution method that takes a device adapter tag and returns
|
||||
/// an object for that device as well as a \c PrepareForControl that simply
|
||||
/// returns an object that works in the control environment.
|
||||
///
|
||||
struct ExecutionAndControlObjectBase : vtkm::cont::ExecutionObjectBase
|
||||
{
|
||||
};
|
||||
|
||||
namespace internal
|
||||
{
|
||||
|
||||
namespace detail
|
||||
{
|
||||
|
||||
struct CheckPrepareForControl
|
||||
{
|
||||
template <typename T>
|
||||
static auto check(T* p) -> decltype(p->PrepareForControl(), std::true_type());
|
||||
|
||||
template <typename T>
|
||||
static auto check(...) -> std::false_type;
|
||||
};
|
||||
|
||||
} // namespace detail
|
||||
|
||||
template <typename T>
|
||||
using IsExecutionAndControlObjectBase =
|
||||
std::is_base_of<vtkm::cont::ExecutionAndControlObjectBase, typename std::decay<T>::type>;
|
||||
|
||||
template <typename T>
|
||||
struct HasPrepareForControl
|
||||
: decltype(detail::CheckPrepareForControl::check<typename std::decay<T>::type>(nullptr))
|
||||
{
|
||||
};
|
||||
|
||||
} // namespace internal
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
/// Checks that the argument is a proper execution object.
|
||||
///
|
||||
#define VTKM_IS_EXECUTION_AND_CONTROL_OBJECT(execObject) \
|
||||
static_assert(::vtkm::cont::internal::IsExecutionAndControlObjectBase<execObject>::value, \
|
||||
"Provided type is not a subclass of vtkm::cont::ExecutionAndControlObjectBase."); \
|
||||
static_assert(::vtkm::cont::internal::HasPrepareForExecution<execObject>::value, \
|
||||
"Provided type does not have requisite PrepareForExecution method."); \
|
||||
static_assert(::vtkm::cont::internal::HasPrepareForControl<execObject>::value, \
|
||||
"Provided type does not have requisite PrepareForControl method.")
|
||||
|
||||
#endif //vtk_m_cont_ExecutionAndControlObjectBase_h
|
||||
@ -19,21 +19,75 @@
|
||||
//============================================================================
|
||||
#ifndef vtk_m_cont_ExecutionObjectBase_h
|
||||
#define vtk_m_cont_ExecutionObjectBase_h
|
||||
|
||||
#include <vtkm/Types.h>
|
||||
|
||||
#include <vtkm/cont/DeviceAdapter.h>
|
||||
|
||||
#if ((VTKM_DEVICE_ADAPTER > 0) && (VTKM_DEVICE_ADAPTER < VTKM_MAX_DEVICE_ADAPTER_ID))
|
||||
// Use the default device adapter tag for testing whether execution objects are valid.
|
||||
#define VTK_M_DEVICE_ADAPTER_TO_TEST_EXEC_OBJECT VTKM_DEFAULT_DEVICE_ADAPTER_TAG
|
||||
#else
|
||||
// The default device adapter is invalid. Perhaps the error device adapter is being used.
|
||||
// In this case, try the serial device adapter instead. It should always be valid.
|
||||
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
|
||||
#define VTK_M_DEVICE_ADAPTER_TO_TEST_EXEC_OBJECT ::vtkm::cont::DeviceAdapterTagSerial
|
||||
#endif
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
|
||||
/// Base \c ExecutionObjectBase for execution objects to inherit from so that
|
||||
/// you can use an arbitrary object as a parameter in an execution environment
|
||||
/// function. Any method you want to use on the execution side must have the
|
||||
/// VTKM_EXEC modifier.
|
||||
/// \tparam Device
|
||||
|
||||
class ExecutionObjectBase
|
||||
/// function. Any subclass of \c ExecutionObjectBase must implement a
|
||||
/// \c PrepareForExecution method that takes a device adapter tag and returns
|
||||
/// an object for that device.
|
||||
///
|
||||
struct ExecutionObjectBase
|
||||
{
|
||||
};
|
||||
|
||||
namespace internal
|
||||
{
|
||||
|
||||
namespace detail
|
||||
{
|
||||
|
||||
struct CheckPrepareForExecution
|
||||
{
|
||||
template <typename T>
|
||||
static auto check(T* p)
|
||||
-> decltype(p->PrepareForExecution(VTK_M_DEVICE_ADAPTER_TO_TEST_EXEC_OBJECT()),
|
||||
std::true_type());
|
||||
|
||||
template <typename T>
|
||||
static auto check(...) -> std::false_type;
|
||||
};
|
||||
|
||||
} // namespace detail
|
||||
|
||||
template <typename T>
|
||||
using IsExecutionObjectBase =
|
||||
std::is_base_of<vtkm::cont::ExecutionObjectBase, typename std::decay<T>::type>;
|
||||
|
||||
template <typename T>
|
||||
struct HasPrepareForExecution
|
||||
: decltype(detail::CheckPrepareForExecution::check<typename std::decay<T>::type>(nullptr))
|
||||
{
|
||||
};
|
||||
|
||||
} // namespace internal
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
/// Checks that the argument is a proper execution object.
|
||||
///
|
||||
#define VTKM_IS_EXECUTION_OBJECT(execObject) \
|
||||
static_assert(::vtkm::cont::internal::IsExecutionObjectBase<execObject>::value, \
|
||||
"Provided type is not a subclass of vtkm::cont::ExecutionObjectBase."); \
|
||||
static_assert(::vtkm::cont::internal::HasPrepareForExecution<execObject>::value, \
|
||||
"Provided type does not have requisite PrepareForExecution method.")
|
||||
|
||||
#endif //vtk_m_cont_ExecutionObjectBase_h
|
||||
|
||||
@ -235,6 +235,29 @@ thread_local static vtkm::cont::RuntimeDeviceTracker runtimeDeviceTracker;
|
||||
VTKM_CONT_EXPORT
|
||||
VTKM_CONT
|
||||
vtkm::cont::RuntimeDeviceTracker GetGlobalRuntimeDeviceTracker();
|
||||
|
||||
struct ScopedGlobalRuntimeDeviceTracker
|
||||
{
|
||||
vtkm::cont::RuntimeDeviceTracker SavedTracker;
|
||||
|
||||
VTKM_CONT ScopedGlobalRuntimeDeviceTracker()
|
||||
: SavedTracker(vtkm::cont::GetGlobalRuntimeDeviceTracker().DeepCopy())
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT ScopedGlobalRuntimeDeviceTracker(vtkm::cont::RuntimeDeviceTracker tracker)
|
||||
: SavedTracker(vtkm::cont::GetGlobalRuntimeDeviceTracker().DeepCopy())
|
||||
{
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().DeepCopy(tracker);
|
||||
}
|
||||
|
||||
VTKM_CONT ~ScopedGlobalRuntimeDeviceTracker()
|
||||
{
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().DeepCopy(this->SavedTracker);
|
||||
}
|
||||
|
||||
ScopedGlobalRuntimeDeviceTracker(const ScopedGlobalRuntimeDeviceTracker&) = delete;
|
||||
};
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
||||
@ -48,10 +48,8 @@ struct Transport<vtkm::cont::arg::TransportTagExecObject, ContObjectType, Device
|
||||
{
|
||||
// If you get a compile error here, it means you tried to use an object that is not an execution
|
||||
// object as an argument that is expected to be one. All execution objects are expected to
|
||||
// inherit from vtkm::cont::ExecutionObjectBase.
|
||||
VTKM_STATIC_ASSERT_MSG(
|
||||
(std::is_base_of<vtkm::cont::ExecutionObjectBase, ContObjectType>::value),
|
||||
"All execution objects are expected to inherit from vtkm::cont::ExecutionObjectBase");
|
||||
// inherit from vtkm::cont::ExecutionObjectBase and have a PrepareForExecution method.
|
||||
VTKM_IS_EXECUTION_OBJECT(ContObjectType);
|
||||
|
||||
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForExecution(Device()));
|
||||
template <typename InputDomainType>
|
||||
|
||||
@ -46,7 +46,7 @@ struct TypeCheckTagExecObject
|
||||
template <typename Type>
|
||||
struct TypeCheck<TypeCheckTagExecObject, Type>
|
||||
{
|
||||
static constexpr bool value = std::is_base_of<vtkm::cont::ExecutionObjectBase, Type>::value;
|
||||
static constexpr bool value = vtkm::cont::internal::IsExecutionObjectBase<Type>::value;
|
||||
};
|
||||
}
|
||||
}
|
||||
|
||||
@ -33,6 +33,13 @@
|
||||
namespace
|
||||
{
|
||||
|
||||
struct NotAnExecutionObject
|
||||
{
|
||||
};
|
||||
struct InvalidExecutionObject : vtkm::cont::ExecutionObjectBase
|
||||
{
|
||||
};
|
||||
|
||||
template <typename Device>
|
||||
struct ExecutionObject
|
||||
{
|
||||
@ -51,6 +58,7 @@ struct TestExecutionObject : public vtkm::cont::ExecutionObjectBase
|
||||
return object;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename Device>
|
||||
struct TestKernel : public vtkm::exec::FunctorBase
|
||||
{
|
||||
@ -83,6 +91,21 @@ void TryExecObjectTransport(Device)
|
||||
|
||||
void TestExecObjectTransport()
|
||||
{
|
||||
std::cout << "Checking ExecObject queries." << std::endl;
|
||||
VTKM_TEST_ASSERT(!vtkm::cont::internal::IsExecutionObjectBase<NotAnExecutionObject>::value,
|
||||
"Bad query");
|
||||
VTKM_TEST_ASSERT(vtkm::cont::internal::IsExecutionObjectBase<InvalidExecutionObject>::value,
|
||||
"Bad query");
|
||||
VTKM_TEST_ASSERT(vtkm::cont::internal::IsExecutionObjectBase<TestExecutionObject>::value,
|
||||
"Bad query");
|
||||
|
||||
VTKM_TEST_ASSERT(!vtkm::cont::internal::HasPrepareForExecution<NotAnExecutionObject>::value,
|
||||
"Bad query");
|
||||
VTKM_TEST_ASSERT(!vtkm::cont::internal::HasPrepareForExecution<InvalidExecutionObject>::value,
|
||||
"Bad query");
|
||||
VTKM_TEST_ASSERT(vtkm::cont::internal::HasPrepareForExecution<TestExecutionObject>::value,
|
||||
"Bad query");
|
||||
|
||||
std::cout << "Trying ExecObject transport with serial device." << std::endl;
|
||||
TryExecObjectTransport(vtkm::cont::DeviceAdapterTagSerial());
|
||||
}
|
||||
|
||||
@ -26,7 +26,6 @@
|
||||
#include <vtkm/UnaryPredicates.h>
|
||||
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/ArrayHandleCast.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
#include <vtkm/cont/ErrorExecution.h>
|
||||
#include <vtkm/cont/vtkm_cont_export.h>
|
||||
@ -126,6 +125,26 @@ __global__ void SumExclusiveScan(T a, T b, T result, BinaryOperationType binary_
|
||||
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
|
||||
#pragma GCC diagnostic pop
|
||||
#endif
|
||||
|
||||
template <typename PortalType, typename OutValueType>
|
||||
struct CastPortal
|
||||
{
|
||||
using ValueType = OutValueType;
|
||||
|
||||
PortalType Portal;
|
||||
|
||||
VTKM_CONT
|
||||
CastPortal(const PortalType& portal)
|
||||
: Portal(portal)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC
|
||||
vtkm::Id GetNumberOfValues() const { return this->Portal.GetNumberOfValues(); }
|
||||
|
||||
VTKM_EXEC
|
||||
ValueType Get(vtkm::Id index) const { return static_cast<OutValueType>(this->Portal.Get(index)); }
|
||||
};
|
||||
}
|
||||
} // end namespace cuda::internal
|
||||
|
||||
@ -313,9 +332,7 @@ private:
|
||||
//The portal type and the initial value AREN'T the same type so we have
|
||||
//to a slower approach, where we wrap the input portal inside a cast
|
||||
//portal
|
||||
using CastFunctor = vtkm::cont::internal::Cast<typename InputPortal::ValueType, T>;
|
||||
|
||||
vtkm::exec::internal::ArrayPortalTransform<T, InputPortal, CastFunctor> castPortal(input);
|
||||
vtkm::cont::cuda::internal::CastPortal<InputPortal, T> castPortal(input);
|
||||
|
||||
vtkm::exec::cuda::internal::WrappedBinaryOperator<T, BinaryFunctor> bop(binary_functor);
|
||||
|
||||
|
||||
@ -38,6 +38,7 @@
|
||||
#include <vtkm/cont/ArrayHandleTransform.h>
|
||||
#include <vtkm/cont/ArrayHandleView.h>
|
||||
#include <vtkm/cont/ArrayHandleZip.h>
|
||||
#include <vtkm/cont/VirtualObjectHandle.h>
|
||||
|
||||
#include <vtkm/worklet/DispatcherMapField.h>
|
||||
#include <vtkm/worklet/WorkletMapField.h>
|
||||
@ -73,7 +74,7 @@ struct ValueSquared
|
||||
struct ValueScale
|
||||
{
|
||||
ValueScale()
|
||||
: Factor()
|
||||
: Factor(1.0)
|
||||
{
|
||||
}
|
||||
|
||||
@ -102,6 +103,111 @@ struct ValueScale
|
||||
private:
|
||||
vtkm::Float64 Factor;
|
||||
};
|
||||
|
||||
struct InverseValueScale
|
||||
{
|
||||
InverseValueScale()
|
||||
: InverseFactor(1.0)
|
||||
{
|
||||
}
|
||||
|
||||
InverseValueScale(vtkm::Float64 factor)
|
||||
: InverseFactor(1.0 / factor)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename ValueType>
|
||||
VTKM_EXEC_CONT ValueType operator()(const ValueType& v) const
|
||||
{
|
||||
using Traits = vtkm::VecTraits<ValueType>;
|
||||
using TTraits = vtkm::TypeTraits<ValueType>;
|
||||
using ComponentType = typename Traits::ComponentType;
|
||||
|
||||
ValueType result = TTraits::ZeroInitialization();
|
||||
for (vtkm::IdComponent i = 0; i < Traits::GetNumberOfComponents(v); ++i)
|
||||
{
|
||||
vtkm::Float64 vi = static_cast<vtkm::Float64>(Traits::GetComponent(v, i));
|
||||
vtkm::Float64 ri = vi * this->InverseFactor;
|
||||
Traits::SetComponent(result, i, static_cast<ComponentType>(ri));
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
private:
|
||||
vtkm::Float64 InverseFactor;
|
||||
};
|
||||
|
||||
template <typename ValueType>
|
||||
struct VirtualTransformFunctorBase : public vtkm::VirtualObjectBase
|
||||
{
|
||||
VirtualTransformFunctorBase() = default;
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
virtual ValueType operator()(const ValueType& v) const = 0;
|
||||
};
|
||||
|
||||
template <typename ValueType, typename FunctorType>
|
||||
struct VirtualTransformFunctor : VirtualTransformFunctorBase<ValueType>
|
||||
{
|
||||
FunctorType Functor;
|
||||
|
||||
VTKM_CONT
|
||||
VirtualTransformFunctor(const FunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
ValueType operator()(const ValueType& v) const override { return this->Functor(v); }
|
||||
};
|
||||
|
||||
template <typename ValueType>
|
||||
struct TransformExecObject : public vtkm::cont::ExecutionAndControlObjectBase
|
||||
{
|
||||
vtkm::cont::VirtualObjectHandle<VirtualTransformFunctorBase<ValueType>> VirtualFunctor;
|
||||
|
||||
VTKM_CONT TransformExecObject() = default;
|
||||
|
||||
template <typename FunctorType>
|
||||
VTKM_CONT TransformExecObject(const FunctorType& functor)
|
||||
{
|
||||
// Need to make sure the serial device is supported, since that is what is used on the
|
||||
// control side.
|
||||
vtkm::cont::ScopedGlobalRuntimeDeviceTracker scopedTracker;
|
||||
vtkm::cont::GetGlobalRuntimeDeviceTracker().ResetDevice(vtkm::cont::DeviceAdapterTagSerial());
|
||||
this->VirtualFunctor.Reset(new VirtualTransformFunctor<ValueType, FunctorType>(functor));
|
||||
}
|
||||
|
||||
struct FunctorWrapper
|
||||
{
|
||||
const VirtualTransformFunctorBase<ValueType>* FunctorPointer;
|
||||
|
||||
FunctorWrapper() = default;
|
||||
|
||||
VTKM_CONT
|
||||
FunctorWrapper(const VirtualTransformFunctorBase<ValueType>* functorPointer)
|
||||
: FunctorPointer(functorPointer)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename InValueType>
|
||||
VTKM_EXEC ValueType operator()(const InValueType& value) const
|
||||
{
|
||||
return (*this->FunctorPointer)(value);
|
||||
}
|
||||
};
|
||||
|
||||
template <typename DeviceAdapterTag>
|
||||
VTKM_CONT FunctorWrapper PrepareForExecution(DeviceAdapterTag device) const
|
||||
{
|
||||
return FunctorWrapper(this->VirtualFunctor.PrepareForExecution(device));
|
||||
}
|
||||
|
||||
VTKM_CONT FunctorWrapper PrepareForControl() const
|
||||
{
|
||||
return FunctorWrapper(this->VirtualFunctor.Get());
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
namespace vtkm
|
||||
@ -476,6 +582,44 @@ private:
|
||||
}
|
||||
};
|
||||
|
||||
struct TestTransformVirtualAsInput
|
||||
{
|
||||
template <typename ValueType>
|
||||
VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const
|
||||
{
|
||||
using FunctorType = fancy_array_detail::ValueScale;
|
||||
using VirtualFunctorType = fancy_array_detail::TransformExecObject<ValueType>;
|
||||
|
||||
const vtkm::Id length = ARRAY_SIZE;
|
||||
FunctorType functor(2.0);
|
||||
VirtualFunctorType virtualFunctor(functor);
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> input;
|
||||
auto transformed = vtkm::cont::make_ArrayHandleTransform(input, virtualFunctor);
|
||||
|
||||
input.Allocate(length);
|
||||
SetPortal(input.GetPortalControl());
|
||||
|
||||
vtkm::cont::printSummary_ArrayHandle(transformed, std::cout);
|
||||
std::cout << std::endl;
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> result;
|
||||
|
||||
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
|
||||
dispatcher.Invoke(transformed, result);
|
||||
|
||||
//verify that the control portal works
|
||||
for (vtkm::Id i = 0; i < length; ++i)
|
||||
{
|
||||
const ValueType result_v = result.GetPortalConstControl().Get(i);
|
||||
const ValueType correct_value = functor(TestValue(i, ValueType()));
|
||||
const ValueType control_value = transformed.GetPortalConstControl().Get(i);
|
||||
VTKM_TEST_ASSERT(test_equal(result_v, correct_value), "Transform Handle Failed");
|
||||
VTKM_TEST_ASSERT(test_equal(result_v, control_value), "Transform Handle Control Failed");
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct TestCountingTransformAsInput
|
||||
{
|
||||
template <typename ValueType>
|
||||
@ -917,6 +1061,88 @@ private:
|
||||
}
|
||||
};
|
||||
|
||||
struct TestTransformAsOutput
|
||||
{
|
||||
template <typename ValueType>
|
||||
VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const
|
||||
{
|
||||
using FunctorType = fancy_array_detail::ValueScale;
|
||||
using InverseFunctorType = fancy_array_detail::InverseValueScale;
|
||||
|
||||
const vtkm::Id length = ARRAY_SIZE;
|
||||
FunctorType functor(2.0);
|
||||
InverseFunctorType inverseFunctor(2.0);
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> input;
|
||||
input.Allocate(length);
|
||||
SetPortal(input.GetPortalControl());
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> output;
|
||||
auto transformed = vtkm::cont::make_ArrayHandleTransform(output, functor, inverseFunctor);
|
||||
|
||||
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
|
||||
dispatcher.Invoke(input, transformed);
|
||||
|
||||
vtkm::cont::printSummary_ArrayHandle(transformed, std::cout);
|
||||
std::cout << std::endl;
|
||||
|
||||
//verify that the control portal works
|
||||
for (vtkm::Id i = 0; i < length; ++i)
|
||||
{
|
||||
const ValueType result_v = output.GetPortalConstControl().Get(i);
|
||||
const ValueType correct_value = inverseFunctor(TestValue(i, ValueType()));
|
||||
const ValueType control_value = transformed.GetPortalConstControl().Get(i);
|
||||
VTKM_TEST_ASSERT(test_equal(result_v, correct_value), "Transform Handle Failed");
|
||||
VTKM_TEST_ASSERT(test_equal(functor(result_v), control_value),
|
||||
"Transform Handle Control Failed");
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct TestTransformVirtualAsOutput
|
||||
{
|
||||
template <typename ValueType>
|
||||
VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const
|
||||
{
|
||||
using FunctorType = fancy_array_detail::ValueScale;
|
||||
using InverseFunctorType = fancy_array_detail::InverseValueScale;
|
||||
|
||||
using VirtualFunctorType = fancy_array_detail::TransformExecObject<ValueType>;
|
||||
|
||||
const vtkm::Id length = ARRAY_SIZE;
|
||||
FunctorType functor(2.0);
|
||||
InverseFunctorType inverseFunctor(2.0);
|
||||
|
||||
VirtualFunctorType virtualFunctor(functor);
|
||||
VirtualFunctorType virtualInverseFunctor(inverseFunctor);
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> input;
|
||||
input.Allocate(length);
|
||||
SetPortal(input.GetPortalControl());
|
||||
|
||||
vtkm::cont::ArrayHandle<ValueType> output;
|
||||
auto transformed =
|
||||
vtkm::cont::make_ArrayHandleTransform(output, virtualFunctor, virtualInverseFunctor);
|
||||
|
||||
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
|
||||
dispatcher.Invoke(input, transformed);
|
||||
|
||||
vtkm::cont::printSummary_ArrayHandle(transformed, std::cout);
|
||||
std::cout << std::endl;
|
||||
|
||||
//verify that the control portal works
|
||||
for (vtkm::Id i = 0; i < length; ++i)
|
||||
{
|
||||
const ValueType result_v = output.GetPortalConstControl().Get(i);
|
||||
const ValueType correct_value = inverseFunctor(TestValue(i, ValueType()));
|
||||
const ValueType control_value = transformed.GetPortalConstControl().Get(i);
|
||||
VTKM_TEST_ASSERT(test_equal(result_v, correct_value), "Transform Handle Failed");
|
||||
VTKM_TEST_ASSERT(test_equal(functor(result_v), control_value),
|
||||
"Transform Handle Control Failed");
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct TestZipAsOutput
|
||||
{
|
||||
template <typename KeyType, typename ValueType>
|
||||
@ -1047,6 +1273,12 @@ private:
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
TestingFancyArrayHandles<DeviceAdapterTag>::TestTransformAsInput(), HandleTypesToTest());
|
||||
|
||||
std::cout << "-------------------------------------------" << std::endl;
|
||||
std::cout << "Testing ArrayHandleTransform with virtual as Input" << std::endl;
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
TestingFancyArrayHandles<DeviceAdapterTag>::TestTransformVirtualAsInput(),
|
||||
HandleTypesToTest());
|
||||
|
||||
std::cout << "-------------------------------------------" << std::endl;
|
||||
std::cout << "Testing ArrayHandleTransform with Counting as Input" << std::endl;
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
@ -1105,6 +1337,17 @@ private:
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
TestingFancyArrayHandles<DeviceAdapterTag>::TestViewAsOutput(), HandleTypesToTest());
|
||||
|
||||
std::cout << "-------------------------------------------" << std::endl;
|
||||
std::cout << "Testing ArrayHandleTransform as Output" << std::endl;
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
TestingFancyArrayHandles<DeviceAdapterTag>::TestTransformAsOutput(), HandleTypesToTest());
|
||||
|
||||
std::cout << "-------------------------------------------" << std::endl;
|
||||
std::cout << "Testing ArrayHandleTransform with virtual as Output" << std::endl;
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
TestingFancyArrayHandles<DeviceAdapterTag>::TestTransformVirtualAsOutput(),
|
||||
HandleTypesToTest());
|
||||
|
||||
std::cout << "-------------------------------------------" << std::endl;
|
||||
std::cout << "Testing ArrayHandleDiscard as Output" << std::endl;
|
||||
vtkm::testing::Testing::TryTypes(
|
||||
|
||||
@ -22,8 +22,6 @@
|
||||
|
||||
#include <vtkm/exec/arg/testing/ThreadIndicesTesting.h>
|
||||
|
||||
#include <vtkm/cont/ExecutionObjectBase.h>
|
||||
|
||||
#include <vtkm/testing/Testing.h>
|
||||
|
||||
#define EXPECTED_NUMBER 67
|
||||
@ -31,7 +29,7 @@
|
||||
namespace
|
||||
{
|
||||
|
||||
struct TestExecutionObject : public vtkm::cont::ExecutionObjectBase
|
||||
struct TestExecutionObject
|
||||
{
|
||||
TestExecutionObject()
|
||||
: Number(static_cast<vtkm::Int32>(0xDEADDEAD))
|
||||
|
||||
@ -328,21 +328,24 @@ public:
|
||||
|
||||
vtkm::Vec<FloatType, 3> centroid(0., 0., 0.);
|
||||
const vtkm::Int32 numIndices = MeshConn.GetCellIndices(cellConn, currentCell);
|
||||
//load local cell data
|
||||
for (int i = 0; i < numIndices; ++i)
|
||||
if (numIndices > 0)
|
||||
{
|
||||
BOUNDS_CHECK(vertices, cellConn[i]);
|
||||
vtkm::Vec<FloatType, 3> point = vtkm::Vec<FloatType, 3>(vertices.Get(cellConn[i]));
|
||||
centroid = centroid + point;
|
||||
xpoints[i] = point[0];
|
||||
ypoints[i] = point[1];
|
||||
zpoints[i] = point[2];
|
||||
}
|
||||
//load local cell data
|
||||
for (int i = 0; i < numIndices; ++i)
|
||||
{
|
||||
BOUNDS_CHECK(vertices, cellConn[i]);
|
||||
vtkm::Vec<FloatType, 3> point = vtkm::Vec<FloatType, 3>(vertices.Get(cellConn[i]));
|
||||
centroid = centroid + point;
|
||||
xpoints[i] = point[0];
|
||||
ypoints[i] = point[1];
|
||||
zpoints[i] = point[2];
|
||||
}
|
||||
|
||||
FloatType invNumIndices = static_cast<FloatType>(1.) / static_cast<FloatType>(numIndices);
|
||||
centroid[0] = centroid[0] * invNumIndices;
|
||||
centroid[1] = centroid[1] * invNumIndices;
|
||||
centroid[2] = centroid[2] * invNumIndices;
|
||||
FloatType invNumIndices = static_cast<FloatType>(1.) / static_cast<FloatType>(numIndices);
|
||||
centroid[0] = centroid[0] * invNumIndices;
|
||||
centroid[1] = centroid[1] * invNumIndices;
|
||||
centroid[2] = centroid[2] * invNumIndices;
|
||||
}
|
||||
|
||||
vtkm::Vec<FloatType, 3> toCentroid = centroid - origin;
|
||||
vtkm::Normalize(toCentroid);
|
||||
|
||||
Loading…
Reference in New Issue
Block a user