From 98a0a20feb6535d7b7f53f438b514d885425fac7 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Wed, 5 Sep 2018 10:07:23 -0600 Subject: [PATCH 1/3] Allow ArrayHandleTransform to work with ExecObject This change allows you to set a subclass of vtkm::cont::ExecutionObjectBase as a functor used in ArrayHandleTransform. This latter class will then detect that the functor is an ExecObject and will call PrepareForExecution with the appropriate device to get the actual Functor object. This change allows you to use virtual objects and other device dependent objects as functors for ArrayHandleTransform without knowing a priori what device the portal will be used on. --- .../array-handle-transform-exec-object.md | 97 +++++++ vtkm/cont/ArrayHandleTransform.h | 163 +++++++++--- vtkm/cont/RuntimeDeviceTracker.h | 23 ++ .../internal/DeviceAdapterAlgorithmCuda.h | 25 +- vtkm/cont/testing/TestingFancyArrayHandles.h | 241 +++++++++++++++++- 5 files changed, 511 insertions(+), 38 deletions(-) create mode 100644 docs/changelog/array-handle-transform-exec-object.md diff --git a/docs/changelog/array-handle-transform-exec-object.md b/docs/changelog/array-handle-transform-exec-object.md new file mode 100644 index 000000000..08a4bf9bd --- /dev/null +++ b/docs/changelog/array-handle-transform-exec-object.md @@ -0,0 +1,97 @@ +# 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::ExecutionObjectBase` that +contains a `PrepareForExecution` 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 +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 an +execution object to set this up for a device. + +``` cpp +template +struct FindValueExecutionObject : vtkm::cont::ExecutionObjectBase +{ + VTKM_IS_ARRAY_HANDLE(ArrayHandleType); + + ArrayHandleType SortedArray; + + FindValueExecutionObject() = default; + + VTKM_CONT FindValueExecutionObject(const ArrayHandleType& sortedArray) + : SortedArray(sortedArray) + { } + + template + VTKM_CONT + FindValueFunctor()(Device()))> + PrepareForExecution(Device device) + { + using FunctorType = + FindValueFunctor()(Device()))> + + return FunctorType(this->SortedArray.PrepareForInput(device)); + } +} +``` + +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(sortedArray)); +``` diff --git a/vtkm/cont/ArrayHandleTransform.h b/vtkm/cont/ArrayHandleTransform.h index 7cafbc0c9..a2fc2e022 100644 --- a/vtkm/cont/ArrayHandleTransform.h +++ b/vtkm/cont/ArrayHandleTransform.h @@ -25,6 +25,9 @@ #include #include #include +#include +#include +#include namespace vtkm { @@ -178,18 +181,94 @@ namespace cont namespace internal { +template +struct TransformFunctorManagerImpl; + +template +struct TransformFunctorManagerImpl +{ + using FunctorType = ProvidedFunctorType; + ProvidedFunctorType Functor; + + TransformFunctorManagerImpl() = default; + + VTKM_CONT + TransformFunctorManagerImpl(const ProvidedFunctorType& functor) + : Functor(functor) + { + } + + VTKM_CONT + FunctorType GetFunctor() const { return this->Functor; } +}; + +template +struct TransformFunctorManagerImpl +{ + using FunctorType = decltype(std::declval().PrepareForExecution(Device())); + ProvidedFunctorType Functor; + + TransformFunctorManagerImpl() = default; + + VTKM_CONT + TransformFunctorManagerImpl(const ProvidedFunctorType& functor) + : Functor(functor) + { + } + + VTKM_CONT + FunctorType GetFunctor() const { return this->Functor.PrepareForExecution(Device()); } +}; + +template +struct TransformFunctorManager + : TransformFunctorManagerImpl< + ProvidedFunctorType, + Device, + typename std::is_base_of::type> +{ + using Superclass = TransformFunctorManagerImpl< + ProvidedFunctorType, + Device, + typename std::is_base_of::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 + VTKM_CONT TransformFunctorManager( + const TransformFunctorManager& other) + : Superclass(other.Functor) + { + } + + template + using TransformedValueType = decltype(std::declval()(ValueType{})); +}; + template struct VTKM_ALWAYS_EXPORT StorageTagTransform { - using ValueType = decltype(FunctorType{}(typename ArrayHandleType::ValueType{})); + using FunctorManager = TransformFunctorManager; + using ValueType = + typename FunctorManager::template TransformedValueType; }; template class Storage::ValueType, StorageTagTransform> { + using FunctorManager = TransformFunctorManager; + public: using ValueType = typename StorageTagTransform::ValueType; @@ -204,7 +283,7 @@ public: using PortalConstType = vtkm::exec::internal::ArrayPortalTransform; + typename FunctorManager::FunctorType>; VTKM_CONT Storage() @@ -223,15 +302,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.GetFunctor()); } VTKM_CONT @@ -269,11 +350,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 +363,9 @@ class Storage< typename StorageTagTransform::ValueType, StorageTagTransform> { + using FunctorManager = TransformFunctorManager; + using InverseFunctorManager = TransformFunctorManager; + public: using ValueType = typename StorageTagTransform::ValueType; @@ -289,13 +373,13 @@ public: using PortalType = vtkm::exec::internal::ArrayPortalTransform; + typename FunctorManager::FunctorType, + typename InverseFunctorManager::FunctorType>; using PortalConstType = vtkm::exec::internal::ArrayPortalTransform; + typename FunctorManager::FunctorType, + typename InverseFunctorManager::FunctorType>; VTKM_CONT Storage() @@ -318,15 +402,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.GetFunctor(), + this->InverseFunctor.GetFunctor()); } 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.GetFunctor(), + this->InverseFunctor.GetFunctor()); } VTKM_CONT @@ -361,15 +452,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 +470,7 @@ class ArrayTransfer:: Device> { using StorageTag = StorageTagTransform; + using FunctorManager = TransformFunctorManager; public: using ValueType = typename StorageTagTransform::ValueType; @@ -392,7 +484,7 @@ public: using PortalConstExecution = vtkm::exec::internal::ArrayPortalTransform< ValueType, typename ArrayHandleType::template ExecutionTypes::PortalConst, - FunctorType>; + typename FunctorManager::FunctorType>; VTKM_CONT ArrayTransfer(StorageType* storage) @@ -407,7 +499,7 @@ public: VTKM_CONT PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) { - return PortalConstExecution(this->Array.PrepareForInput(Device()), this->Functor); + return PortalConstExecution(this->Array.PrepareForInput(Device()), this->Functor.GetFunctor()); } VTKM_CONT @@ -443,7 +535,7 @@ public: private: ArrayHandleType Array; - FunctorType Functor; + FunctorManager Functor; }; template { using StorageTag = StorageTagTransform; + using FunctorManager = TransformFunctorManager; + using InverseFunctorManager = TransformFunctorManager; public: using ValueType = typename StorageTagTransform::ValueType; @@ -467,13 +561,13 @@ public: using PortalExecution = vtkm::exec::internal::ArrayPortalTransform< ValueType, typename ArrayHandleType::template ExecutionTypes::Portal, - FunctorType, - InverseFunctorType>; + typename FunctorManager::FunctorType, + typename InverseFunctorManager::FunctorType>; using PortalConstExecution = vtkm::exec::internal::ArrayPortalTransform< ValueType, typename ArrayHandleType::template ExecutionTypes::PortalConst, - FunctorType, - InverseFunctorType>; + typename FunctorManager::FunctorType, + typename InverseFunctorManager::FunctorType>; VTKM_CONT ArrayTransfer(StorageType* storage) @@ -489,22 +583,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.GetFunctor(), + this->InverseFunctor.GetFunctor()); } 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.GetFunctor(), + this->InverseFunctor.GetFunctor()); } 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.GetFunctor(), + this->InverseFunctor.GetFunctor()); } VTKM_CONT @@ -522,8 +619,8 @@ public: private: ArrayHandleType Array; - FunctorType Functor; - InverseFunctorType InverseFunctor; + FunctorManager Functor; + InverseFunctorManager InverseFunctor; }; } // namespace internal diff --git a/vtkm/cont/RuntimeDeviceTracker.h b/vtkm/cont/RuntimeDeviceTracker.h index 59c2c22f4..8c2654c71 100644 --- a/vtkm/cont/RuntimeDeviceTracker.h +++ b/vtkm/cont/RuntimeDeviceTracker.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 diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h index f0f17ac84..0b2d1e18f 100644 --- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h +++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h @@ -26,7 +26,6 @@ #include #include -#include #include #include #include @@ -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 +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(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; - - vtkm::exec::internal::ArrayPortalTransform castPortal(input); + vtkm::cont::cuda::internal::CastPortal castPortal(input); vtkm::exec::cuda::internal::WrappedBinaryOperator bop(binary_functor); diff --git a/vtkm/cont/testing/TestingFancyArrayHandles.h b/vtkm/cont/testing/TestingFancyArrayHandles.h index 6cd8b905a..ce6253a62 100644 --- a/vtkm/cont/testing/TestingFancyArrayHandles.h +++ b/vtkm/cont/testing/TestingFancyArrayHandles.h @@ -38,6 +38,7 @@ #include #include #include +#include #include #include @@ -73,7 +74,7 @@ struct ValueSquared struct ValueScale { ValueScale() - : Factor() + : Factor(1.0) { } @@ -102,6 +103,107 @@ struct ValueScale private: vtkm::Float64 Factor; }; + +struct InverseValueScale +{ + InverseValueScale() + : InverseFactor(1.0) + { + } + + InverseValueScale(vtkm::Float64 factor) + : InverseFactor(1.0 / factor) + { + } + + template + VTKM_EXEC_CONT ValueType operator()(const ValueType& v) const + { + using Traits = vtkm::VecTraits; + using TTraits = vtkm::TypeTraits; + using ComponentType = typename Traits::ComponentType; + + ValueType result = TTraits::ZeroInitialization(); + for (vtkm::IdComponent i = 0; i < Traits::GetNumberOfComponents(v); ++i) + { + vtkm::Float64 vi = static_cast(Traits::GetComponent(v, i)); + vtkm::Float64 ri = vi * this->InverseFactor; + Traits::SetComponent(result, i, static_cast(ri)); + } + return result; + } + +private: + vtkm::Float64 InverseFactor; +}; + +template +struct VirtualTransformFunctorBase : public vtkm::VirtualObjectBase +{ + VirtualTransformFunctorBase() = default; + + VTKM_EXEC_CONT + virtual ValueType operator()(const ValueType& v) const = 0; +}; + +template +struct VirtualTransformFunctor : VirtualTransformFunctorBase +{ + 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 +struct TransformExecObject : public vtkm::cont::ExecutionObjectBase +{ + vtkm::cont::VirtualObjectHandle> VirtualFunctor; + + VTKM_CONT TransformExecObject() = default; + + template + 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(functor)); + } + + template + struct ExecutionObject + { + const VirtualTransformFunctorBase* FunctorPointer; + + ExecutionObject() = default; + + VTKM_CONT + ExecutionObject(const VirtualTransformFunctorBase* functorPointer) + : FunctorPointer(functorPointer) + { + } + + template + VTKM_EXEC ValueType operator()(const InValueType& value) const + { + return (*this->FunctorPointer)(value); + } + }; + + template + VTKM_CONT ExecutionObject PrepareForExecution(DeviceAdapterTag device) const + { + return ExecutionObject(this->VirtualFunctor.PrepareForExecution(device)); + } +}; } namespace vtkm @@ -476,6 +578,44 @@ private: } }; + struct TestTransformVirtualAsInput + { + template + VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const + { + using FunctorType = fancy_array_detail::ValueScale; + using VirtualFunctorType = fancy_array_detail::TransformExecObject; + + const vtkm::Id length = ARRAY_SIZE; + FunctorType functor(2.0); + VirtualFunctorType virtualFunctor(functor); + + vtkm::cont::ArrayHandle 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 result; + + vtkm::worklet::DispatcherMapField 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 @@ -917,6 +1057,88 @@ private: } }; + struct TestTransformAsOutput + { + template + 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 input; + input.Allocate(length); + SetPortal(input.GetPortalControl()); + + vtkm::cont::ArrayHandle output; + auto transformed = vtkm::cont::make_ArrayHandleTransform(output, functor, inverseFunctor); + + vtkm::worklet::DispatcherMapField 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 + 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; + + const vtkm::Id length = ARRAY_SIZE; + FunctorType functor(2.0); + InverseFunctorType inverseFunctor(2.0); + + VirtualFunctorType virtualFunctor(functor); + VirtualFunctorType virtualInverseFunctor(inverseFunctor); + + vtkm::cont::ArrayHandle input; + input.Allocate(length); + SetPortal(input.GetPortalControl()); + + vtkm::cont::ArrayHandle output; + auto transformed = + vtkm::cont::make_ArrayHandleTransform(output, virtualFunctor, virtualInverseFunctor); + + vtkm::worklet::DispatcherMapField 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 @@ -1047,6 +1269,12 @@ private: vtkm::testing::Testing::TryTypes( TestingFancyArrayHandles::TestTransformAsInput(), HandleTypesToTest()); + std::cout << "-------------------------------------------" << std::endl; + std::cout << "Testing ArrayHandleTransform with virtual as Input" << std::endl; + vtkm::testing::Testing::TryTypes( + TestingFancyArrayHandles::TestTransformVirtualAsInput(), + HandleTypesToTest()); + std::cout << "-------------------------------------------" << std::endl; std::cout << "Testing ArrayHandleTransform with Counting as Input" << std::endl; vtkm::testing::Testing::TryTypes( @@ -1105,6 +1333,17 @@ private: vtkm::testing::Testing::TryTypes( TestingFancyArrayHandles::TestViewAsOutput(), HandleTypesToTest()); + std::cout << "-------------------------------------------" << std::endl; + std::cout << "Testing ArrayHandleTransform as Output" << std::endl; + vtkm::testing::Testing::TryTypes( + TestingFancyArrayHandles::TestTransformAsOutput(), HandleTypesToTest()); + + std::cout << "-------------------------------------------" << std::endl; + std::cout << "Testing ArrayHandleTransform with virtual as Output" << std::endl; + vtkm::testing::Testing::TryTypes( + TestingFancyArrayHandles::TestTransformVirtualAsOutput(), + HandleTypesToTest()); + std::cout << "-------------------------------------------" << std::endl; std::cout << "Testing ArrayHandleDiscard as Output" << std::endl; vtkm::testing::Testing::TryTypes( From 62d3b9f4fbbcca311f78c6e77c7be27a8945acd1 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Wed, 5 Sep 2018 15:28:26 -0600 Subject: [PATCH 2/3] Correct warning about potential divide by zero I'm not sure why this warning suddenly popped up, but this was an easy fix. --- .../raytracing/ConnectivityTracer.hxx | 29 ++++++++++--------- 1 file changed, 16 insertions(+), 13 deletions(-) diff --git a/vtkm/rendering/raytracing/ConnectivityTracer.hxx b/vtkm/rendering/raytracing/ConnectivityTracer.hxx index 9e4bfeda1..499016f1f 100644 --- a/vtkm/rendering/raytracing/ConnectivityTracer.hxx +++ b/vtkm/rendering/raytracing/ConnectivityTracer.hxx @@ -328,21 +328,24 @@ public: vtkm::Vec 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 point = vtkm::Vec(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 point = vtkm::Vec(vertices.Get(cellConn[i])); + centroid = centroid + point; + xpoints[i] = point[0]; + ypoints[i] = point[1]; + zpoints[i] = point[2]; + } - FloatType invNumIndices = static_cast(1.) / static_cast(numIndices); - centroid[0] = centroid[0] * invNumIndices; - centroid[1] = centroid[1] * invNumIndices; - centroid[2] = centroid[2] * invNumIndices; + FloatType invNumIndices = static_cast(1.) / static_cast(numIndices); + centroid[0] = centroid[0] * invNumIndices; + centroid[1] = centroid[1] * invNumIndices; + centroid[2] = centroid[2] * invNumIndices; + } vtkm::Vec toCentroid = centroid - origin; vtkm::Normalize(toCentroid); From 2b05487398a1973e099c7bf95b6e99a1cf79ec0a Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Sat, 8 Sep 2018 11:07:06 -0600 Subject: [PATCH 3/3] Add ExecutionAndControlObjectBase This is a subclass of ExecutionObject and a superset of its functionality. In addition to having a PrepareForExecution method, it also has a PrepareForControl method that gets an object appropriate for the control environment. This is helpful for situations where you need code to work in both environments, such as the functor in an ArrayHandleTransform. Also added several runtime checks for execution objects and execution and cotnrol objects. --- .../array-handle-transform-exec-object.md | 23 +++-- .../execution-and-control-objects.md | 16 ++++ vtkm/cont/Algorithm.h | 15 ++-- vtkm/cont/ArrayHandleTransform.h | 89 +++++++++++-------- vtkm/cont/CMakeLists.txt | 1 + vtkm/cont/ExecutionAndControlObjectBase.h | 82 +++++++++++++++++ vtkm/cont/ExecutionObjectBase.h | 64 +++++++++++-- vtkm/cont/arg/TransportTagExecObject.h | 6 +- vtkm/cont/arg/TypeCheckTagExecObject.h | 2 +- .../testing/UnitTestTransportExecObject.cxx | 23 +++++ vtkm/cont/testing/TestingFancyArrayHandles.h | 18 ++-- .../arg/testing/UnitTestFetchExecObject.cxx | 4 +- 12 files changed, 273 insertions(+), 70 deletions(-) create mode 100644 docs/changelog/execution-and-control-objects.md create mode 100644 vtkm/cont/ExecutionAndControlObjectBase.h diff --git a/docs/changelog/array-handle-transform-exec-object.md b/docs/changelog/array-handle-transform-exec-object.md index 08a4bf9bd..d2a665dcc 100644 --- a/docs/changelog/array-handle-transform-exec-object.md +++ b/docs/changelog/array-handle-transform-exec-object.md @@ -12,8 +12,9 @@ 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::ExecutionObjectBase` that -contains a `PrepareForExecution` method. +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 @@ -56,12 +57,14 @@ struct FindValueFunctor 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 an -execution object to set this up for a device. +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 -struct FindValueExecutionObject : vtkm::cont::ExecutionObjectBase +struct FindValueExecutionObject : vtkm::cont::ExecutionAndControlObjectBase { VTKM_IS_ARRAY_HANDLE(ArrayHandleType); @@ -83,6 +86,16 @@ struct FindValueExecutionObject : vtkm::cont::ExecutionObjectBase return FunctorType(this->SortedArray.PrepareForInput(device)); } + + VTKM_CONT + FundValueFunctor + PrepareForControl() + { + using FunctorType = + FindValueFunctor + + return FunctorType(this->SortedArray.GetPortalConstControl()); + } } ``` diff --git a/docs/changelog/execution-and-control-objects.md b/docs/changelog/execution-and-control-objects.md new file mode 100644 index 000000000..ebb55801c --- /dev/null +++ b/docs/changelog/execution-and-control-objects.md @@ -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`. diff --git a/vtkm/cont/Algorithm.h b/vtkm/cont/Algorithm.h index 07df19ec7..798e688b0 100644 --- a/vtkm/cont/Algorithm.h +++ b/vtkm/cont/Algorithm.h @@ -38,24 +38,25 @@ template inline auto DoPrepareArgForExec(T&& object, std::true_type) -> decltype(std::declval().PrepareForExecution(Device())) { + VTKM_IS_EXECUTION_OBJECT(T); return object.PrepareForExecution(Device{}); } template inline T&& DoPrepareArgForExec(T&& object, std::false_type) { + static_assert(!vtkm::cont::internal::IsExecutionObjectBase::value, + "Internal error: failed to detect execution object."); return std::forward(object); } template -auto PrepareArgForExec(T&& object) -> decltype(DoPrepareArgForExec( - std::forward(object), - typename std::is_base_of::type>::type{})) +auto PrepareArgForExec(T&& object) + -> decltype(DoPrepareArgForExec(std::forward(object), + vtkm::cont::internal::IsExecutionObjectBase{})) { - return DoPrepareArgForExec( - std::forward(object), - typename std::is_base_of::type>::type{}); + return DoPrepareArgForExec(std::forward(object), + vtkm::cont::internal::IsExecutionObjectBase{}); } struct CopyFunctor diff --git a/vtkm/cont/ArrayHandleTransform.h b/vtkm/cont/ArrayHandleTransform.h index a2fc2e022..c272f3150 100644 --- a/vtkm/cont/ArrayHandleTransform.h +++ b/vtkm/cont/ArrayHandleTransform.h @@ -25,9 +25,8 @@ #include #include #include -#include +#include #include -#include namespace vtkm { @@ -181,14 +180,17 @@ namespace cont namespace internal { -template +template struct TransformFunctorManagerImpl; -template -struct TransformFunctorManagerImpl +template +struct TransformFunctorManagerImpl { + VTKM_STATIC_ASSERT_MSG(!vtkm::cont::internal::IsExecutionObjectBase::value, + "Must use an ExecutionAndControlObject instead of an ExecutionObject."); + + ProvidedFunctorType Functor; using FunctorType = ProvidedFunctorType; - ProvidedFunctorType Functor; TransformFunctorManagerImpl() = default; @@ -199,14 +201,23 @@ struct TransformFunctorManagerImpl } VTKM_CONT - FunctorType GetFunctor() const { return this->Functor; } + ProvidedFunctorType PrepareForControl() const { return this->Functor; } + + template + VTKM_CONT ProvidedFunctorType PrepareForExecution(Device) const + { + return this->Functor; + } }; -template -struct TransformFunctorManagerImpl +template +struct TransformFunctorManagerImpl { - using FunctorType = decltype(std::declval().PrepareForExecution(Device())); + VTKM_IS_EXECUTION_AND_CONTROL_OBJECT(ProvidedFunctorType); + ProvidedFunctorType Functor; + // using FunctorType = decltype(std::declval().PrepareForControl()); + using FunctorType = decltype(Functor.PrepareForControl()); TransformFunctorManagerImpl() = default; @@ -217,20 +228,28 @@ struct TransformFunctorManagerImpl } VTKM_CONT - FunctorType GetFunctor() const { return this->Functor.PrepareForExecution(Device()); } + auto PrepareForControl() const -> decltype(this->Functor.PrepareForControl()) + { + return this->Functor.PrepareForControl(); + } + + template + VTKM_CONT auto PrepareForExecution(Device device) const + -> decltype(this->Functor.PrepareForExecution(device)) + { + return this->Functor.PrepareForExecution(device); + } }; -template +template struct TransformFunctorManager : TransformFunctorManagerImpl< ProvidedFunctorType, - Device, - typename std::is_base_of::type> + typename vtkm::cont::internal::IsExecutionAndControlObjectBase::type> { using Superclass = TransformFunctorManagerImpl< ProvidedFunctorType, - Device, - typename std::is_base_of::type>; + typename vtkm::cont::internal::IsExecutionAndControlObjectBase::type>; using FunctorType = typename Superclass::FunctorType; VTKM_CONT TransformFunctorManager() = default; @@ -242,13 +261,6 @@ struct TransformFunctorManager { } - template - VTKM_CONT TransformFunctorManager( - const TransformFunctorManager& other) - : Superclass(other.Functor) - { - } - template using TransformedValueType = decltype(std::declval()(ValueType{})); }; @@ -312,7 +324,7 @@ public: VTKM_ASSERT(this->Valid); vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope; vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial()); - return PortalConstType(this->Array.GetPortalConstControl(), this->Functor.GetFunctor()); + return PortalConstType(this->Array.GetPortalConstControl(), this->Functor.PrepareForControl()); } VTKM_CONT @@ -405,8 +417,8 @@ public: vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope; vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial()); return PortalType(this->Array.GetPortalControl(), - this->Functor.GetFunctor(), - this->InverseFunctor.GetFunctor()); + this->Functor.PrepareForControl(), + this->InverseFunctor.PrepareForControl()); } VTKM_CONT @@ -416,8 +428,8 @@ public: vtkm::cont::ScopedGlobalRuntimeDeviceTracker trackerScope; vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(vtkm::cont::DeviceAdapterTagSerial()); return PortalConstType(this->Array.GetPortalConstControl(), - this->Functor.GetFunctor(), - this->InverseFunctor.GetFunctor()); + this->Functor.PrepareForControl(), + this->InverseFunctor.PrepareForControl()); } VTKM_CONT @@ -470,7 +482,7 @@ class ArrayTransfer:: Device> { using StorageTag = StorageTagTransform; - using FunctorManager = TransformFunctorManager; + using FunctorManager = TransformFunctorManager; public: using ValueType = typename StorageTagTransform::ValueType; @@ -499,7 +511,8 @@ public: VTKM_CONT PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) { - return PortalConstExecution(this->Array.PrepareForInput(Device()), this->Functor.GetFunctor()); + return PortalConstExecution(this->Array.PrepareForInput(Device()), + this->Functor.PrepareForExecution(Device())); } VTKM_CONT @@ -548,8 +561,8 @@ class ArrayTransfer< Device> { using StorageTag = StorageTagTransform; - using FunctorManager = TransformFunctorManager; - using InverseFunctorManager = TransformFunctorManager; + using FunctorManager = TransformFunctorManager; + using InverseFunctorManager = TransformFunctorManager; public: using ValueType = typename StorageTagTransform::ValueType; @@ -584,24 +597,24 @@ public: PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) { return PortalConstExecution(this->Array.PrepareForInput(Device()), - this->Functor.GetFunctor(), - this->InverseFunctor.GetFunctor()); + this->Functor.PrepareForExecution(Device()), + this->InverseFunctor.PrepareForExecution(Device())); } VTKM_CONT PortalExecution PrepareForInPlace(bool& vtkmNotUsed(updateData)) { return PortalExecution(this->Array.PrepareForInPlace(Device()), - this->Functor.GetFunctor(), - this->InverseFunctor.GetFunctor()); + this->Functor.PrepareForExecution(Device()), + this->InverseFunctor.PrepareForExecution(Device())); } VTKM_CONT PortalExecution PrepareForOutput(vtkm::Id numberOfValues) { return PortalExecution(this->Array.PrepareForOutput(numberOfValues, Device()), - this->Functor.GetFunctor(), - this->InverseFunctor.GetFunctor()); + this->Functor.PrepareForExecution(Device()), + this->InverseFunctor.PrepareForExecution(Device())); } VTKM_CONT diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 2583ea4a2..c04f2bd02 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -83,6 +83,7 @@ set(headers ErrorFilterExecution.h ErrorExecution.h ErrorInternal.h + ExecutionAndControlObjectBase.h ExecutionObjectBase.h Field.h FieldRangeCompute.h diff --git a/vtkm/cont/ExecutionAndControlObjectBase.h b/vtkm/cont/ExecutionAndControlObjectBase.h new file mode 100644 index 000000000..26f340dca --- /dev/null +++ b/vtkm/cont/ExecutionAndControlObjectBase.h @@ -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 + +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 + static auto check(T* p) -> decltype(p->PrepareForControl(), std::true_type()); + + template + static auto check(...) -> std::false_type; +}; + +} // namespace detail + +template +using IsExecutionAndControlObjectBase = + std::is_base_of::type>; + +template +struct HasPrepareForControl + : decltype(detail::CheckPrepareForControl::check::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::value, \ + "Provided type is not a subclass of vtkm::cont::ExecutionAndControlObjectBase."); \ + static_assert(::vtkm::cont::internal::HasPrepareForExecution::value, \ + "Provided type does not have requisite PrepareForExecution method."); \ + static_assert(::vtkm::cont::internal::HasPrepareForControl::value, \ + "Provided type does not have requisite PrepareForControl method.") + +#endif //vtk_m_cont_ExecutionAndControlObjectBase_h diff --git a/vtkm/cont/ExecutionObjectBase.h b/vtkm/cont/ExecutionObjectBase.h index d0a30e1f7..b33b123d4 100644 --- a/vtkm/cont/ExecutionObjectBase.h +++ b/vtkm/cont/ExecutionObjectBase.h @@ -19,21 +19,75 @@ //============================================================================ #ifndef vtk_m_cont_ExecutionObjectBase_h #define vtk_m_cont_ExecutionObjectBase_h + #include + +#include + +#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 +#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 + static auto check(T* p) + -> decltype(p->PrepareForExecution(VTK_M_DEVICE_ADAPTER_TO_TEST_EXEC_OBJECT()), + std::true_type()); + + template + static auto check(...) -> std::false_type; +}; + +} // namespace detail + +template +using IsExecutionObjectBase = + std::is_base_of::type>; + +template +struct HasPrepareForExecution + : decltype(detail::CheckPrepareForExecution::check::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::value, \ + "Provided type is not a subclass of vtkm::cont::ExecutionObjectBase."); \ + static_assert(::vtkm::cont::internal::HasPrepareForExecution::value, \ + "Provided type does not have requisite PrepareForExecution method.") + #endif //vtk_m_cont_ExecutionObjectBase_h diff --git a/vtkm/cont/arg/TransportTagExecObject.h b/vtkm/cont/arg/TransportTagExecObject.h index b6b5f0cbb..bb5cf6a1d 100644 --- a/vtkm/cont/arg/TransportTagExecObject.h +++ b/vtkm/cont/arg/TransportTagExecObject.h @@ -48,10 +48,8 @@ struct Transport::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().PrepareForExecution(Device())); template diff --git a/vtkm/cont/arg/TypeCheckTagExecObject.h b/vtkm/cont/arg/TypeCheckTagExecObject.h index 67e4649b1..f45bb2f13 100644 --- a/vtkm/cont/arg/TypeCheckTagExecObject.h +++ b/vtkm/cont/arg/TypeCheckTagExecObject.h @@ -46,7 +46,7 @@ struct TypeCheckTagExecObject template struct TypeCheck { - static constexpr bool value = std::is_base_of::value; + static constexpr bool value = vtkm::cont::internal::IsExecutionObjectBase::value; }; } } diff --git a/vtkm/cont/arg/testing/UnitTestTransportExecObject.cxx b/vtkm/cont/arg/testing/UnitTestTransportExecObject.cxx index ec0987d79..67548c31d 100644 --- a/vtkm/cont/arg/testing/UnitTestTransportExecObject.cxx +++ b/vtkm/cont/arg/testing/UnitTestTransportExecObject.cxx @@ -33,6 +33,13 @@ namespace { +struct NotAnExecutionObject +{ +}; +struct InvalidExecutionObject : vtkm::cont::ExecutionObjectBase +{ +}; + template struct ExecutionObject { @@ -51,6 +58,7 @@ struct TestExecutionObject : public vtkm::cont::ExecutionObjectBase return object; } }; + template 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::value, + "Bad query"); + VTKM_TEST_ASSERT(vtkm::cont::internal::IsExecutionObjectBase::value, + "Bad query"); + VTKM_TEST_ASSERT(vtkm::cont::internal::IsExecutionObjectBase::value, + "Bad query"); + + VTKM_TEST_ASSERT(!vtkm::cont::internal::HasPrepareForExecution::value, + "Bad query"); + VTKM_TEST_ASSERT(!vtkm::cont::internal::HasPrepareForExecution::value, + "Bad query"); + VTKM_TEST_ASSERT(vtkm::cont::internal::HasPrepareForExecution::value, + "Bad query"); + std::cout << "Trying ExecObject transport with serial device." << std::endl; TryExecObjectTransport(vtkm::cont::DeviceAdapterTagSerial()); } diff --git a/vtkm/cont/testing/TestingFancyArrayHandles.h b/vtkm/cont/testing/TestingFancyArrayHandles.h index ce6253a62..16f2d8ea9 100644 --- a/vtkm/cont/testing/TestingFancyArrayHandles.h +++ b/vtkm/cont/testing/TestingFancyArrayHandles.h @@ -162,7 +162,7 @@ struct VirtualTransformFunctor : VirtualTransformFunctorBase }; template -struct TransformExecObject : public vtkm::cont::ExecutionObjectBase +struct TransformExecObject : public vtkm::cont::ExecutionAndControlObjectBase { vtkm::cont::VirtualObjectHandle> VirtualFunctor; @@ -178,15 +178,14 @@ struct TransformExecObject : public vtkm::cont::ExecutionObjectBase this->VirtualFunctor.Reset(new VirtualTransformFunctor(functor)); } - template - struct ExecutionObject + struct FunctorWrapper { const VirtualTransformFunctorBase* FunctorPointer; - ExecutionObject() = default; + FunctorWrapper() = default; VTKM_CONT - ExecutionObject(const VirtualTransformFunctorBase* functorPointer) + FunctorWrapper(const VirtualTransformFunctorBase* functorPointer) : FunctorPointer(functorPointer) { } @@ -199,9 +198,14 @@ struct TransformExecObject : public vtkm::cont::ExecutionObjectBase }; template - VTKM_CONT ExecutionObject PrepareForExecution(DeviceAdapterTag device) const + VTKM_CONT FunctorWrapper PrepareForExecution(DeviceAdapterTag device) const { - return ExecutionObject(this->VirtualFunctor.PrepareForExecution(device)); + return FunctorWrapper(this->VirtualFunctor.PrepareForExecution(device)); + } + + VTKM_CONT FunctorWrapper PrepareForControl() const + { + return FunctorWrapper(this->VirtualFunctor.Get()); } }; } diff --git a/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx b/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx index 4a1c8eb68..d23d8379d 100644 --- a/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx +++ b/vtkm/exec/arg/testing/UnitTestFetchExecObject.cxx @@ -22,8 +22,6 @@ #include -#include - #include #define EXPECTED_NUMBER 67 @@ -31,7 +29,7 @@ namespace { -struct TestExecutionObject : public vtkm::cont::ExecutionObjectBase +struct TestExecutionObject { TestExecutionObject() : Number(static_cast(0xDEADDEAD))