From ad0a53af7159e7450c6ca6b7829a47d35db1abb7 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Tue, 21 Jan 2020 13:18:03 -0700 Subject: [PATCH] Convert execution preparation to use tokens Marked the old versions of PrepareFor* that do not use tokens as deprecated and moved all of the code to use the new versions that require a token. This makes the scope of the execution object more explicit so that it will be kept while in use and can potentially be reclaimed afterward. --- benchmarking/BenchmarkFieldAlgorithms.cxx | 14 +- benchmarking/BenchmarkFilters.cxx | 3 +- vtkm/cont/Algorithm.h | 7 +- vtkm/cont/ArrayHandle.h | 13 +- vtkm/cont/ArrayHandleCartesianProduct.h | 8 +- vtkm/cont/CellLocator.h | 11 +- .../CellLocatorBoundingIntervalHierarchy.cxx | 21 +- .../CellLocatorBoundingIntervalHierarchy.h | 4 +- vtkm/cont/CellLocatorGeneral.cxx | 5 +- vtkm/cont/CellLocatorGeneral.h | 3 +- vtkm/cont/CellLocatorRectilinearGrid.cxx | 11 +- vtkm/cont/CellLocatorRectilinearGrid.h | 3 +- vtkm/cont/CellLocatorUniformBins.cxx | 17 +- vtkm/cont/CellLocatorUniformBins.h | 30 ++- vtkm/cont/CellLocatorUniformGrid.cxx | 11 +- vtkm/cont/CellLocatorUniformGrid.h | 3 +- vtkm/cont/CellSetExtrude.cxx | 4 +- vtkm/cont/ColorTable.h | 4 + vtkm/cont/ColorTable.hxx | 32 ++- vtkm/cont/ExecutionObjectBase.h | 71 +++-- vtkm/cont/ImplicitFunctionHandle.h | 5 +- vtkm/cont/PointLocator.h | 19 +- vtkm/cont/PointLocatorUniformGrid.cxx | 18 +- vtkm/cont/PointLocatorUniformGrid.h | 3 +- vtkm/cont/StorageExtrude.h | 3 +- vtkm/cont/StorageVirtual.hxx | 9 +- vtkm/cont/VirtualObjectHandle.h | 14 +- vtkm/cont/arg/TransportTagArrayIn.h | 3 +- vtkm/cont/arg/TransportTagArrayInOut.h | 3 +- vtkm/cont/arg/TransportTagArrayOut.h | 4 +- vtkm/cont/arg/TransportTagCellSetIn.h | 7 +- vtkm/cont/arg/TransportTagTopologyFieldIn.h | 3 +- .../internal/DeviceAdapterAlgorithmCuda.h | 191 ++++++++------ vtkm/cont/internal/ArrayHandleBasicImpl.h | 16 +- .../internal/ConnectivityExplicitInternals.h | 13 +- .../internal/DeviceAdapterAlgorithmGeneral.h | 195 +++++++++----- .../internal/ReverseConnectivityBuilder.h | 13 +- vtkm/cont/internal/VirtualObjectTransfer.h | 19 +- .../VirtualObjectTransferShareWithControl.h | 5 +- .../internal/DeviceAdapterAlgorithmOpenMP.h | 44 ++-- .../internal/DeviceAdapterAlgorithmSerial.h | 62 +++-- .../tbb/internal/DeviceAdapterAlgorithmTBB.h | 64 +++-- vtkm/cont/tbb/internal/ParallelSortTBB.h | 43 ++-- vtkm/cont/testing/TestingArrayHandles.h | 48 +++- vtkm/cont/testing/TestingBitField.h | 14 +- .../TestingCellLocatorRectilinearGrid.h | 10 +- vtkm/cont/testing/TestingColorTable.h | 11 +- vtkm/cont/testing/TestingDeviceAdapter.h | 243 ++++++++++++------ vtkm/cont/testing/TestingFancyArrayHandles.h | 5 +- vtkm/cont/testing/TestingImplicitFunction.h | 3 +- .../cont/testing/TestingVirtualObjectHandle.h | 12 +- vtkm/cont/testing/UnitTestAlgorithm.cxx | 2 +- .../testing/UnitTestArrayHandleDiscard.cxx | 10 +- .../UnitTestArrayHandleExtractComponent.cxx | 12 +- .../testing/UnitTestArrayHandleImplicit.cxx | 3 +- .../UnitTestArrayHandlePermutation.cxx | 29 ++- .../testing/UnitTestArrayHandleSwizzle.cxx | 3 +- .../testing/UnitTestArrayHandleTransform.cxx | 52 ++-- .../testing/UnitTestArrayHandleVirtual.cxx | 7 +- vtkm/cont/testing/UnitTestCellSetExplicit.cxx | 3 +- .../testing/UnitTestDataSetPermutation.cxx | 11 +- .../testing/UnitTestDataSetRectilinear.cxx | 15 +- vtkm/cont/testing/UnitTestDataSetUniform.cxx | 15 +- .../cont/testing/UnitTestMoveConstructors.cxx | 16 +- ...CellLocatorBoundingIntervalHierarchyExec.h | 11 +- vtkm/exec/CellLocatorRectilinearGrid.h | 7 +- vtkm/exec/CellLocatorUniformGrid.h | 5 +- vtkm/rendering/Texture2D.h | 11 +- vtkm/rendering/Wireframer.h | 52 ++-- vtkm/rendering/raytracing/ChannelBuffer.cxx | 66 ++--- vtkm/rendering/raytracing/ChannelBuffer.h | 9 +- .../raytracing/CylinderIntersector.cxx | 13 +- .../raytracing/MeshConnectivityBase.h | 24 +- .../raytracing/MeshConnectivityBuilder.cxx | 3 +- .../raytracing/MeshConnectivityContainers.cxx | 62 +++-- .../raytracing/MeshConnectivityContainers.h | 15 +- vtkm/rendering/raytracing/QuadIntersector.cxx | 8 +- vtkm/rendering/raytracing/Ray.h | 62 ++--- vtkm/rendering/raytracing/RayOperations.h | 43 ++-- .../rendering/raytracing/RayTracingTypeDefs.h | 5 +- .../raytracing/SphereIntersector.cxx | 13 +- .../raytracing/TriangleIntersector.cxx | 18 +- .../raytracing/VolumeRendererStructured.cxx | 45 ++-- vtkm/testing/TestingAlgorithms.h | 30 ++- vtkm/worklet/Clip.h | 17 +- vtkm/worklet/Contour.h | 21 +- vtkm/worklet/DispatcherStreamingMapField.h | 15 +- vtkm/worklet/PointMerge.h | 3 +- vtkm/worklet/StableSortIndices.h | 8 +- vtkm/worklet/clip/ClipTables.h | 9 +- vtkm/worklet/contour/ContourTables.h | 20 +- vtkm/worklet/contourtree/EdgePeakComparator.h | 13 +- .../contourtree/VertexMergeComparator.h | 7 +- .../contourtree_augmented/ContourTreeMaker.h | 4 +- .../activegraph/EdgePeakComparator.h | 36 +-- .../activegraph/HyperArcSuperNodeComparator.h | 7 +- .../activegraph/SuperArcNodeComparator.h | 20 +- .../ContourTreeNodeComparator.h | 14 +- .../ContourTreeSuperNodeComparator.h | 15 +- ...TransferLeafChains_TransferToContourTree.h | 34 +-- .../mesh_dem/SimulatedSimplicityComperator.h | 4 +- .../mesh_dem_meshtypes/ContourTreeMesh.h | 15 +- .../Freudenthal_2D_Triangulation.h | 12 +- .../Freudenthal_3D_Triangulation.h | 22 +- .../MarchingCubes_3D_Triangulation.h | 32 ++- .../mesh_dem_meshtypes/MeshBoundary.h | 55 ++-- .../MeshStructureContourTreeMesh.h | 7 +- .../MeshStructureFreudenthal2D.h | 11 +- .../MeshStructureFreudenthal3D.h | 16 +- .../MeshStructureMarchingCubes.h | 27 +- .../contourtreemesh/ArcComparator.h | 9 +- .../contourtreemesh/CombinedVector.h | 11 +- .../CombinedVectorDifferentFromNext.h | 9 +- .../SuperArcVolumetricComparator.h | 13 +- .../SuperNodeBranchComparator.h | 18 +- vtkm/worklet/internal/DispatcherBase.h | 12 +- vtkm/worklet/internal/TriangulateTables.h | 32 ++- .../testing/UnitTestDispatcherBase.cxx | 2 +- .../CellInterpolationHelper.h | 54 ++-- .../particleadvection/GridEvaluators.h | 14 +- vtkm/worklet/particleadvection/Integrators.h | 34 ++- vtkm/worklet/particleadvection/Particles.h | 32 ++- .../TemporalGridEvaluators.h | 12 +- .../worklet/testing/UnitTestOrientNormals.cxx | 30 +-- .../UnitTestWorkletMapFieldExecArg.cxx | 2 +- 125 files changed, 1698 insertions(+), 1087 deletions(-) diff --git a/benchmarking/BenchmarkFieldAlgorithms.cxx b/benchmarking/BenchmarkFieldAlgorithms.cxx index c7b60baae..68b3a418e 100644 --- a/benchmarking/BenchmarkFieldAlgorithms.cxx +++ b/benchmarking/BenchmarkFieldAlgorithms.cxx @@ -814,8 +814,9 @@ void BenchImplicitFunction(::benchmark::State& state) state.SetLabel(desc.str()); } + vtkm::cont::Token token; auto handle = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere1); - auto function = static_cast(handle.PrepareForExecution(device)); + auto function = static_cast(handle.PrepareForExecution(device, token)); EvalWorklet eval(function); vtkm::cont::Timer timer{ device }; @@ -847,8 +848,9 @@ void BenchVirtualImplicitFunction(::benchmark::State& state) state.SetLabel(desc.str()); } + vtkm::cont::Token token; auto sphere = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere1); - EvalWorklet eval(sphere.PrepareForExecution(device)); + EvalWorklet eval(sphere.PrepareForExecution(device, token)); vtkm::cont::Timer timer{ device }; vtkm::cont::Invoker invoker{ device }; @@ -879,10 +881,11 @@ void Bench2ImplicitFunctions(::benchmark::State& state) state.SetLabel(desc.str()); } + vtkm::cont::Token token; auto h1 = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere1); auto h2 = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere2); - auto f1 = static_cast(h1.PrepareForExecution(device)); - auto f2 = static_cast(h2.PrepareForExecution(device)); + auto f1 = static_cast(h1.PrepareForExecution(device, token)); + auto f2 = static_cast(h2.PrepareForExecution(device, token)); EvalWorklet eval(f1, f2); vtkm::cont::Timer timer{ device }; @@ -914,9 +917,10 @@ void Bench2VirtualImplicitFunctions(::benchmark::State& state) state.SetLabel(desc.str()); } + vtkm::cont::Token token; auto s1 = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere1); auto s2 = vtkm::cont::make_ImplicitFunctionHandle(data.Sphere2); - EvalWorklet eval(s1.PrepareForExecution(device), s2.PrepareForExecution(device)); + EvalWorklet eval(s1.PrepareForExecution(device, token), s2.PrepareForExecution(device, token)); vtkm::cont::Timer timer{ device }; vtkm::cont::Invoker invoker{ device }; diff --git a/benchmarking/BenchmarkFilters.cxx b/benchmarking/BenchmarkFilters.cxx index 0d9b2e977..033bcaeb6 100644 --- a/benchmarking/BenchmarkFilters.cxx +++ b/benchmarking/BenchmarkFilters.cxx @@ -549,9 +549,10 @@ struct PrepareForInput CellSetT& mcellSet = const_cast(cellSet); mcellSet.ResetConnectivity(vtkm::TopologyElementTagPoint{}, vtkm::TopologyElementTagCell{}); + vtkm::cont::Token token; this->Timer.Start(); auto result = cellSet.PrepareForInput( - DeviceTag{}, vtkm::TopologyElementTagPoint{}, vtkm::TopologyElementTagCell{}); + DeviceTag{}, vtkm::TopologyElementTagPoint{}, vtkm::TopologyElementTagCell{}, token); ::benchmark::DoNotOptimize(result); this->Timer.Stop(); diff --git a/vtkm/cont/Algorithm.h b/vtkm/cont/Algorithm.h index a8ca86d0c..e492be0c8 100644 --- a/vtkm/cont/Algorithm.h +++ b/vtkm/cont/Algorithm.h @@ -27,12 +27,11 @@ namespace cont namespace detail { template -inline auto DoPrepareArgForExec(T&& object, vtkm::cont::Token& token, std::true_type) - -> decltype(std::declval().PrepareForExecution(Device())) +inline auto DoPrepareArgForExec(T&& object, vtkm::cont::Token& token, std::true_type) -> decltype( + vtkm::cont::internal::CallPrepareForExecution(std::forward(object), Device{}, token)) { VTKM_IS_EXECUTION_OBJECT(T); - return vtkm::cont::internal::CallPrepareForExecution(object, Device{}, token); - return object.PrepareForExecution(Device{}); + return vtkm::cont::internal::CallPrepareForExecution(std::forward(object), Device{}, token); } template diff --git a/vtkm/cont/ArrayHandle.h b/vtkm/cont/ArrayHandle.h index a563183a2..a5e0a02fe 100644 --- a/vtkm/cont/ArrayHandle.h +++ b/vtkm/cont/ArrayHandle.h @@ -13,6 +13,7 @@ #include #include +#include #include #include @@ -495,24 +496,24 @@ public: DeviceAdapterTag, vtkm::cont::Token& token); - // TODO: Deprecate these template - VTKM_CONT + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInput now requires a vtkm::cont::Token object.") typename ExecutionTypes::PortalConst PrepareForInput(DeviceAdapterTag) const { vtkm::cont::Token token; return this->PrepareForInput(DeviceAdapterTag{}, token); } template - VTKM_CONT typename ExecutionTypes::Portal PrepareForOutput( - vtkm::Id numberOfValues, - DeviceAdapterTag) + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForOutput now requires a vtkm::cont::Token object.") + typename ExecutionTypes::Portal + PrepareForOutput(vtkm::Id numberOfValues, DeviceAdapterTag) { vtkm::cont::Token token; return this->PrepareForOutput(numberOfValues, DeviceAdapterTag{}, token); } template - VTKM_CONT typename ExecutionTypes::Portal PrepareForInPlace(DeviceAdapterTag) + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInPlace now requires a vtkm::cont::Token object.") + typename ExecutionTypes::Portal PrepareForInPlace(DeviceAdapterTag) { vtkm::cont::Token token; return this->PrepareForInPlace(DeviceAdapterTag{}, token); diff --git a/vtkm/cont/ArrayHandleCartesianProduct.h b/vtkm/cont/ArrayHandleCartesianProduct.h index 27b9f034e..619d47eef 100644 --- a/vtkm/cont/ArrayHandleCartesianProduct.h +++ b/vtkm/cont/ArrayHandleCartesianProduct.h @@ -337,11 +337,11 @@ public: } VTKM_CONT - PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData), vtkm::cont::Token&) + PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData), vtkm::cont::Token& token) { - return PortalConstExecution(this->FirstArray.PrepareForInput(Device()), - this->SecondArray.PrepareForInput(Device()), - this->ThirdArray.PrepareForInput(Device())); + return PortalConstExecution(this->FirstArray.PrepareForInput(Device(), token), + this->SecondArray.PrepareForInput(Device(), token), + this->ThirdArray.PrepareForInput(Device(), token)); } VTKM_CONT diff --git a/vtkm/cont/CellLocator.h b/vtkm/cont/CellLocator.h index caa6e5d41..746ca3a82 100644 --- a/vtkm/cont/CellLocator.h +++ b/vtkm/cont/CellLocator.h @@ -56,7 +56,16 @@ public: } VTKM_CONT virtual const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const = 0; + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const = 0; + + VTKM_CONT + VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object.") + const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device) const + { + vtkm::cont::Token token; + return this->PrepareForExecution(device, token); + } protected: void SetModified() { this->Modified = true; } diff --git a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx index c24353aa6..381dd028b 100644 --- a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx +++ b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx @@ -453,6 +453,7 @@ struct CellLocatorBIHPrepareForExecutionFunctor bool operator()( DeviceAdapter, const CellSetType& cellset, + vtkm::cont::Token& token, vtkm::cont::VirtualObjectHandle& bihExec, const vtkm::cont::ArrayHandle& nodes, const vtkm::cont::ArrayHandle& processedCellIds, @@ -461,7 +462,7 @@ struct CellLocatorBIHPrepareForExecutionFunctor using ExecutionType = vtkm::exec::CellLocatorBoundingIntervalHierarchyExec; ExecutionType* execObject = - new ExecutionType(nodes, processedCellIds, cellset, coords, DeviceAdapter()); + new ExecutionType(nodes, processedCellIds, cellset, coords, DeviceAdapter(), token); bihExec.Reset(execObject); return true; } @@ -470,11 +471,17 @@ struct CellLocatorBIHPrepareForExecutionFunctor struct BIHCellSetCaster { template - void operator()(CellSet&& cellset, vtkm::cont::DeviceAdapterId device, Args&&... args) const + void operator()(CellSet&& cellset, + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token, + Args&&... args) const { //We need to go though CastAndCall first - const bool success = vtkm::cont::TryExecuteOnDevice( - device, CellLocatorBIHPrepareForExecutionFunctor(), cellset, std::forward(args)...); + const bool success = vtkm::cont::TryExecuteOnDevice(device, + CellLocatorBIHPrepareForExecutionFunctor(), + cellset, + token, + std::forward(args)...); if (!success) { throwFailedRuntimeDeviceTransfer("BoundingIntervalHierarchy", device); @@ -485,15 +492,17 @@ struct BIHCellSetCaster const vtkm::exec::CellLocator* CellLocatorBoundingIntervalHierarchy::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { this->GetCellSet().CastAndCall(BIHCellSetCaster{}, device, + token, this->ExecutionObjectHandle, this->Nodes, this->ProcessedCellIds, this->GetCoordinates().GetData()); - return this->ExecutionObjectHandle.PrepareForExecution(device); + return this->ExecutionObjectHandle.PrepareForExecution(device, token); ; } diff --git a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h index e5b2d5219..188a35491 100644 --- a/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h +++ b/vtkm/cont/CellLocatorBoundingIntervalHierarchy.h @@ -62,8 +62,8 @@ public: vtkm::Id GetMaxLeafSize() { return this->MaxLeafSize; } VTKM_CONT - const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const override; + const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override; protected: VTKM_CONT diff --git a/vtkm/cont/CellLocatorGeneral.cxx b/vtkm/cont/CellLocatorGeneral.cxx index 091cb4cd5..22d7de220 100644 --- a/vtkm/cont/CellLocatorGeneral.cxx +++ b/vtkm/cont/CellLocatorGeneral.cxx @@ -70,11 +70,12 @@ VTKM_CONT CellLocatorGeneral::CellLocatorGeneral() VTKM_CONT CellLocatorGeneral::~CellLocatorGeneral() = default; VTKM_CONT const vtkm::exec::CellLocator* CellLocatorGeneral::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { if (this->Locator) { - return this->Locator->PrepareForExecution(device); + return this->Locator->PrepareForExecution(device, token); } return nullptr; } diff --git a/vtkm/cont/CellLocatorGeneral.h b/vtkm/cont/CellLocatorGeneral.h index 4ec53f4bc..ce15763d9 100644 --- a/vtkm/cont/CellLocatorGeneral.h +++ b/vtkm/cont/CellLocatorGeneral.h @@ -53,7 +53,8 @@ public: VTKM_CONT void ResetToDefaultConfigurator(); VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const override; + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override; protected: VTKM_CONT void Build() override; diff --git a/vtkm/cont/CellLocatorRectilinearGrid.cxx b/vtkm/cont/CellLocatorRectilinearGrid.cxx index eadb75baa..700f9581f 100644 --- a/vtkm/cont/CellLocatorRectilinearGrid.cxx +++ b/vtkm/cont/CellLocatorRectilinearGrid.cxx @@ -68,10 +68,12 @@ struct CellLocatorRectilinearGridPrepareForExecutionFunctor template VTKM_CONT bool operator()(DeviceAdapter, vtkm::cont::VirtualObjectHandle& execLocator, + vtkm::cont::Token& token, Args&&... args) const { using ExecutionType = vtkm::exec::CellLocatorRectilinearGrid; - ExecutionType* execObject = new ExecutionType(std::forward(args)..., DeviceAdapter()); + ExecutionType* execObject = + new ExecutionType(std::forward(args)..., DeviceAdapter(), token); execLocator.Reset(execObject); return true; } @@ -79,7 +81,8 @@ struct CellLocatorRectilinearGridPrepareForExecutionFunctor } const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { bool success = false; if (this->Is3D) @@ -88,6 +91,7 @@ const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution( device, CellLocatorRectilinearGridPrepareForExecutionFunctor<3>(), this->ExecutionObjectHandle, + token, this->PlaneSize, this->RowSize, this->GetCellSet().template Cast(), @@ -99,6 +103,7 @@ const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution( device, CellLocatorRectilinearGridPrepareForExecutionFunctor<2>(), this->ExecutionObjectHandle, + token, this->PlaneSize, this->RowSize, this->GetCellSet().template Cast(), @@ -108,7 +113,7 @@ const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution( { throwFailedRuntimeDeviceTransfer("CellLocatorRectilinearGrid", device); } - return this->ExecutionObjectHandle.PrepareForExecution(device); + return this->ExecutionObjectHandle.PrepareForExecution(device, token); } } //namespace cont } //namespace vtkm diff --git a/vtkm/cont/CellLocatorRectilinearGrid.h b/vtkm/cont/CellLocatorRectilinearGrid.h index 3aa985530..6e7ed95b9 100644 --- a/vtkm/cont/CellLocatorRectilinearGrid.h +++ b/vtkm/cont/CellLocatorRectilinearGrid.h @@ -26,7 +26,8 @@ public: VTKM_CONT ~CellLocatorRectilinearGrid() override; VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const override; + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override; protected: VTKM_CONT void Build() override; diff --git a/vtkm/cont/CellLocatorUniformBins.cxx b/vtkm/cont/CellLocatorUniformBins.cxx index e0fd7d5bc..407f32ceb 100644 --- a/vtkm/cont/CellLocatorUniformBins.cxx +++ b/vtkm/cont/CellLocatorUniformBins.cxx @@ -457,6 +457,7 @@ struct CellLocatorUniformBins::MakeExecObject template VTKM_CONT void operator()(const CellSetType& cellSet, DeviceAdapter, + vtkm::cont::Token& token, const CellLocatorUniformBins& self) const { auto execObject = @@ -467,7 +468,8 @@ struct CellLocatorUniformBins::MakeExecObject self.CellCount, self.CellIds, cellSet, - self.GetCoordinates()); + self.GetCoordinates(), + token); self.ExecutionObjectHandle.Reset(execObject); } }; @@ -475,21 +477,24 @@ struct CellLocatorUniformBins::MakeExecObject struct CellLocatorUniformBins::PrepareForExecutionFunctor { template - VTKM_CONT bool operator()(DeviceAdapter, const CellLocatorUniformBins& self) const + VTKM_CONT bool operator()(DeviceAdapter, + vtkm::cont::Token& token, + const CellLocatorUniformBins& self) const { - self.GetCellSet().CastAndCall(MakeExecObject{}, DeviceAdapter{}, self); + self.GetCellSet().CastAndCall(MakeExecObject{}, DeviceAdapter{}, token, self); return true; } }; VTKM_CONT const vtkm::exec::CellLocator* CellLocatorUniformBins::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { - if (!vtkm::cont::TryExecuteOnDevice(device, PrepareForExecutionFunctor(), *this)) + if (!vtkm::cont::TryExecuteOnDevice(device, PrepareForExecutionFunctor(), token, *this)) { throwFailedRuntimeDeviceTransfer("CellLocatorUniformBins", device); } - return this->ExecutionObjectHandle.PrepareForExecution(device); + return this->ExecutionObjectHandle.PrepareForExecution(device, token); } //---------------------------------------------------------------------------- diff --git a/vtkm/cont/CellLocatorUniformBins.h b/vtkm/cont/CellLocatorUniformBins.h index 9c94938d6..23c883362 100644 --- a/vtkm/cont/CellLocatorUniformBins.h +++ b/vtkm/cont/CellLocatorUniformBins.h @@ -94,13 +94,15 @@ private: using ArrayPortalConst = typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst; - using CoordsPortalType = - decltype(vtkm::cont::ArrayHandleVirtualCoordinates{}.PrepareForInput(DeviceAdapter{})); + using CoordsPortalType = decltype(vtkm::cont::ArrayHandleVirtualCoordinates{}.PrepareForInput( + DeviceAdapter{}, + std::declval())); using CellSetP2CExecType = decltype(std::declval().PrepareForInput(DeviceAdapter{}, vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{})); + vtkm::TopologyElementTagPoint{}, + std::declval())); // TODO: This function may return false positives for non 3D cells as the // tests are done on the projection of the point on the cell. Extra checks @@ -132,17 +134,19 @@ public: const vtkm::cont::ArrayHandle& cellCount, const vtkm::cont::ArrayHandle& cellIds, const CellSetType& cellSet, - const vtkm::cont::CoordinateSystem& coords) + const vtkm::cont::CoordinateSystem& coords, + vtkm::cont::Token& token) : TopLevel(topLevelGrid) - , LeafDimensions(leafDimensions.PrepareForInput(DeviceAdapter{})) - , LeafStartIndex(leafStartIndex.PrepareForInput(DeviceAdapter{})) - , CellStartIndex(cellStartIndex.PrepareForInput(DeviceAdapter{})) - , CellCount(cellCount.PrepareForInput(DeviceAdapter{})) - , CellIds(cellIds.PrepareForInput(DeviceAdapter{})) + , LeafDimensions(leafDimensions.PrepareForInput(DeviceAdapter{}, token)) + , LeafStartIndex(leafStartIndex.PrepareForInput(DeviceAdapter{}, token)) + , CellStartIndex(cellStartIndex.PrepareForInput(DeviceAdapter{}, token)) + , CellCount(cellCount.PrepareForInput(DeviceAdapter{}, token)) + , CellIds(cellIds.PrepareForInput(DeviceAdapter{}, token)) , CellSet(cellSet.PrepareForInput(DeviceAdapter{}, vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{})) - , Coords(coords.GetData().PrepareForInput(DeviceAdapter{})) + vtkm::TopologyElementTagPoint{}, + token)) + , Coords(coords.GetData().PrepareForInput(DeviceAdapter{}, token)) { } @@ -254,8 +258,8 @@ public: void PrintSummary(std::ostream& out) const; - const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const override; + const vtkm::exec::CellLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override; private: VTKM_CONT void Build() override; diff --git a/vtkm/cont/CellLocatorUniformGrid.cxx b/vtkm/cont/CellLocatorUniformGrid.cxx index a1cf8dd88..773a34500 100644 --- a/vtkm/cont/CellLocatorUniformGrid.cxx +++ b/vtkm/cont/CellLocatorUniformGrid.cxx @@ -80,10 +80,12 @@ struct CellLocatorUniformGridPrepareForExecutionFunctor template VTKM_CONT bool operator()(DeviceAdapter, vtkm::cont::VirtualObjectHandle& execLocator, + vtkm::cont::Token& token, Args&&... args) const { using ExecutionType = vtkm::exec::CellLocatorUniformGrid; - ExecutionType* execObject = new ExecutionType(std::forward(args)..., DeviceAdapter()); + ExecutionType* execObject = + new ExecutionType(std::forward(args)..., DeviceAdapter(), token); execLocator.Reset(execObject); return true; } @@ -91,7 +93,8 @@ struct CellLocatorUniformGridPrepareForExecutionFunctor } const vtkm::exec::CellLocator* CellLocatorUniformGrid::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { bool success = true; if (this->Is3D) @@ -99,6 +102,7 @@ const vtkm::exec::CellLocator* CellLocatorUniformGrid::PrepareForExecution( success = vtkm::cont::TryExecuteOnDevice(device, CellLocatorUniformGridPrepareForExecutionFunctor<3>(), this->ExecutionObjectHandle, + token, this->CellDims, this->PointDims, this->Origin, @@ -111,6 +115,7 @@ const vtkm::exec::CellLocator* CellLocatorUniformGrid::PrepareForExecution( success = vtkm::cont::TryExecuteOnDevice(device, CellLocatorUniformGridPrepareForExecutionFunctor<2>(), this->ExecutionObjectHandle, + token, this->CellDims, this->PointDims, this->Origin, @@ -122,7 +127,7 @@ const vtkm::exec::CellLocator* CellLocatorUniformGrid::PrepareForExecution( { throwFailedRuntimeDeviceTransfer("CellLocatorUniformGrid", device); } - return this->ExecutionObjectHandle.PrepareForExecution(device); + return this->ExecutionObjectHandle.PrepareForExecution(device, token); } } //namespace cont diff --git a/vtkm/cont/CellLocatorUniformGrid.h b/vtkm/cont/CellLocatorUniformGrid.h index 78d4f49dc..ae81c43d2 100644 --- a/vtkm/cont/CellLocatorUniformGrid.h +++ b/vtkm/cont/CellLocatorUniformGrid.h @@ -26,7 +26,8 @@ public: VTKM_CONT ~CellLocatorUniformGrid() override; VTKM_CONT const vtkm::exec::CellLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const override; + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const override; protected: VTKM_CONT void Build() override; diff --git a/vtkm/cont/CellSetExtrude.cxx b/vtkm/cont/CellSetExtrude.cxx index 3ada91610..24645aa5b 100644 --- a/vtkm/cont/CellSetExtrude.cxx +++ b/vtkm/cont/CellSetExtrude.cxx @@ -168,9 +168,11 @@ vtkm::IdComponent CellSetExtrude::GetNumberOfPointsInCell(vtkm::Id) const void CellSetExtrude::GetCellPointIds(vtkm::Id id, vtkm::Id* ptids) const { + vtkm::cont::Token token; auto conn = this->PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{}); + vtkm::TopologyElementTagPoint{}, + token); auto indices = conn.GetIndices(id); for (int i = 0; i < 6; ++i) { diff --git a/vtkm/cont/ColorTable.h b/vtkm/cont/ColorTable.h index e194907b6..fd839954c 100644 --- a/vtkm/cont/ColorTable.h +++ b/vtkm/cont/ColorTable.h @@ -688,6 +688,10 @@ public: /// \brief returns a virtual object pointer of the exec color table /// /// This pointer is only valid as long as the ColorTable is unmodified + inline const vtkm::exec::ColorTableBase* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const; + + VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object") inline const vtkm::exec::ColorTableBase* PrepareForExecution( vtkm::cont::DeviceAdapterId deviceId) const; diff --git a/vtkm/cont/ColorTable.hxx b/vtkm/cont/ColorTable.hxx index 8b83e7b52..a7a3a02a2 100644 --- a/vtkm/cont/ColorTable.hxx +++ b/vtkm/cont/ColorTable.hxx @@ -48,13 +48,15 @@ struct transfer_color_table_to_device { template - inline bool operator()(DeviceAdapter device, vtkm::cont::ColorTable::TransferState&& state) const + inline bool operator()(DeviceAdapter device, + vtkm::cont::ColorTable::TransferState&& state, + vtkm::cont::Token& token) const { - auto p1 = state.ColorPosHandle.PrepareForInput(device); - auto p2 = state.ColorRGBHandle.PrepareForInput(device); - auto p3 = state.OpacityPosHandle.PrepareForInput(device); - auto p4 = state.OpacityAlphaHandle.PrepareForInput(device); - auto p5 = state.OpacityMidSharpHandle.PrepareForInput(device); + auto p1 = state.ColorPosHandle.PrepareForInput(device, token); + auto p2 = state.ColorRGBHandle.PrepareForInput(device, token); + auto p3 = state.OpacityPosHandle.PrepareForInput(device, token); + auto p4 = state.OpacityAlphaHandle.PrepareForInput(device, token); + auto p5 = state.OpacityMidSharpHandle.PrepareForInput(device, token); //The rest of the data member on portal are set when-ever the user //modifies the ColorTable instance and don't need to specified here @@ -78,7 +80,9 @@ struct map_color_table template inline bool operator()(DeviceAdapter device, ColorTable&& colors, Args&&... args) const { - vtkm::worklet::colorconversion::TransferFunction transfer(colors->PrepareForExecution(device)); + vtkm::cont::Token token; + vtkm::worklet::colorconversion::TransferFunction transfer( + colors->PrepareForExecution(device, token)); vtkm::cont::Invoker invoke(device); invoke(transfer, std::forward(args)...); return true; @@ -348,7 +352,8 @@ bool ColorTable::Sample(vtkm::Int32 numSamples, //--------------------------------------------------------------------------- const vtkm::exec::ColorTableBase* ColorTable::PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { //Build the ColorTable instance that is needed for execution if (this->NeedToCreateExecutionColorTable()) @@ -400,13 +405,20 @@ const vtkm::exec::ColorTableBase* ColorTable::PrepareForExecution( if (info.NeedsTransfer) { bool transfered = vtkm::cont::TryExecuteOnDevice( - device, detail::transfer_color_table_to_device{}, std::move(info)); + device, detail::transfer_color_table_to_device{}, std::move(info), token); if (!transfered) { throwFailedRuntimeDeviceTransfer("ColorTable", device); } } - return this->GetExecutionHandle()->PrepareForExecution(device); + return this->GetExecutionHandle()->PrepareForExecution(device, token); +} + +const vtkm::exec::ColorTableBase* ColorTable::PrepareForExecution( + vtkm::cont::DeviceAdapterId device) const +{ + vtkm::cont::Token token; + return this->PrepareForExecution(device, token); } } } diff --git a/vtkm/cont/ExecutionObjectBase.h b/vtkm/cont/ExecutionObjectBase.h index bb3c90b3c..c0a12dbd4 100644 --- a/vtkm/cont/ExecutionObjectBase.h +++ b/vtkm/cont/ExecutionObjectBase.h @@ -52,9 +52,11 @@ struct CheckPrepareForExecution struct CheckPrepareForExecutionDeprecated { + VTKM_DEPRECATED_SUPPRESS_BEGIN template static auto check(T* p) -> decltype(p->PrepareForExecution(vtkm::cont::DeviceAdapterTagSerial{}), std::true_type()); + VTKM_DEPRECATED_SUPPRESS_END template static auto check(...) -> std::false_type; @@ -64,7 +66,7 @@ struct CheckPrepareForExecutionDeprecated template using IsExecutionObjectBase = - std::is_base_of::type>; + typename std::is_base_of::type>::type; template struct HasPrepareForExecution @@ -88,37 +90,6 @@ struct HasPrepareForExecutionDeprecated ::vtkm::cont::internal::HasPrepareForExecutionDeprecated::value, \ "Provided type does not have requisite PrepareForExecution method.") -namespace detail -{ - -template -VTKM_CONT auto CallPrepareForExecutionImpl(T&& execObject, - Device device, - vtkm::cont::Token& token, - std::true_type, - std::false_type) - -> decltype(execObject.PrepareForExecution(device, token)) -{ - return execObject.PrepareForExecution(device, token); -} - -template -VTKM_DEPRECATED( - 1.6, - "ExecutionObjects now require a PrepareForExecution that takes a vtkm::cont::Token object." - "PrepareForExecution(Device) is deprecated. Implement PrepareForExecution(Device, Token).") -VTKM_CONT auto CallPrepareForExecutionImpl(T&& execObject, - Device device, - vtkm::cont::Token&, - std::false_type, - std::true_type) - -> decltype(execObject.PrepareForExecution(device)) -{ - return execObject.PrepareForExecution(device); -} - -} // namespace detail - /// \brief Gets the object to use in the execution environment from an ExecutionObject. /// /// An execution object (that is, an object inheriting from `vtkm::cont::ExecutionObjectBase`) is @@ -128,20 +99,36 @@ VTKM_CONT auto CallPrepareForExecutionImpl(T&& execObject, /// template VTKM_CONT auto CallPrepareForExecution(T&& execObject, Device device, vtkm::cont::Token& token) - -> decltype(detail::CallPrepareForExecutionImpl(std::forward(execObject), - device, - token, - HasPrepareForExecution{}, - HasPrepareForExecutionDeprecated{})) + -> decltype(execObject.PrepareForExecution(device, token)) { VTKM_IS_EXECUTION_OBJECT(T); VTKM_IS_DEVICE_ADAPTER_TAG(Device); - return detail::CallPrepareForExecutionImpl(std::forward(execObject), - device, - token, - HasPrepareForExecution{}, - HasPrepareForExecutionDeprecated{}); + return execObject.PrepareForExecution(device, token); +} + +// If you get a deprecation warning at this function, it means that an ExecutionObject is using the +// old style PrepareForExecution. Update its PrepareForExecution method to accept both a device and +// a token. +// +// Developer note: the third template argument, TokenType, is expected to be a vtkm::cont::Token +// (which is ignored). The reason why it is a template argument instead of just the type expected +// is so that ExecObjects that implement both versions of PrepareForExecution (for backward +// compatibility) will match the non-deprecated version instead of being ambiguous. +template +VTKM_CONT VTKM_DEPRECATED( + 1.6, + "ExecutionObjects now require a PrepareForExecution that takes a vtkm::cont::Token object. " + "PrepareForExecution(Device) is deprecated. Implement PrepareForExecution(Device, " + "Token).") auto CallPrepareForExecution(T&& execObject, Device device, TokenType&) + -> decltype(execObject.PrepareForExecution(device)) +{ + VTKM_IS_EXECUTION_OBJECT(T); + VTKM_IS_DEVICE_ADAPTER_TAG(Device); + VTKM_STATIC_ASSERT( + (std::is_same::type>::value)); + + return execObject.PrepareForExecution(device); } /// \brief Gets the type of the execution-side object for an ExecutionObject. diff --git a/vtkm/cont/ImplicitFunctionHandle.h b/vtkm/cont/ImplicitFunctionHandle.h index 7c2a6deb9..2f355fdd7 100644 --- a/vtkm/cont/ImplicitFunctionHandle.h +++ b/vtkm/cont/ImplicitFunctionHandle.h @@ -93,9 +93,10 @@ public: VTKM_CONT const vtkm::cont::ImplicitFunctionHandle& GetHandle() const { return this->Handle; } VTKM_CONT - vtkm::ImplicitFunctionValue PrepareForExecution(vtkm::cont::DeviceAdapterId device) const + vtkm::ImplicitFunctionValue PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { - return vtkm::ImplicitFunctionValue(this->Handle.PrepareForExecution(device)); + return vtkm::ImplicitFunctionValue(this->Handle.PrepareForExecution(device, token)); } VTKM_CONT vtkm::ImplicitFunctionValue PrepareForControl() const diff --git a/vtkm/cont/PointLocator.h b/vtkm/cont/PointLocator.h index 280df8eea..4acfcf480 100644 --- a/vtkm/cont/PointLocator.h +++ b/vtkm/cont/PointLocator.h @@ -47,11 +47,19 @@ public: } } - VTKM_CONT const vtkm::exec::PointLocator* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const + VTKM_CONT const vtkm::exec::PointLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const { - this->PrepareExecutionObject(this->ExecutionObjectHandle, device); - return this->ExecutionObjectHandle.PrepareForExecution(device); + this->PrepareExecutionObject(this->ExecutionObjectHandle, device, token); + return this->ExecutionObjectHandle.PrepareForExecution(device, token); + } + + VTKM_CONT + VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object") + const vtkm::exec::PointLocator* PrepareForExecution(vtkm::cont::DeviceAdapterId device) const + { + vtkm::cont::Token token; + return this->PrepareForExecution(device, token); } protected: @@ -64,7 +72,8 @@ protected: using ExecutionObjectHandleType = vtkm::cont::VirtualObjectHandle; VTKM_CONT virtual void PrepareExecutionObject(ExecutionObjectHandleType& execObjHandle, - vtkm::cont::DeviceAdapterId deviceId) const = 0; + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const = 0; private: vtkm::cont::CoordinateSystem Coords; diff --git a/vtkm/cont/PointLocatorUniformGrid.cxx b/vtkm/cont/PointLocatorUniformGrid.cxx index 85bb430e6..f584fdf04 100644 --- a/vtkm/cont/PointLocatorUniformGrid.cxx +++ b/vtkm/cont/PointLocatorUniformGrid.cxx @@ -101,7 +101,8 @@ struct PointLocatorUniformGrid::PrepareExecutionObjectFunctor template VTKM_CONT bool operator()(DeviceAdapter, const vtkm::cont::PointLocatorUniformGrid& self, - ExecutionObjectHandleType& handle) const + ExecutionObjectHandleType& handle, + vtkm::cont::Token& token) const { auto rmin = vtkm::make_Vec(static_cast(self.Range[0].Min), static_cast(self.Range[1].Min), @@ -114,10 +115,10 @@ struct PointLocatorUniformGrid::PrepareExecutionObjectFunctor rmin, rmax, self.Dims, - self.GetCoordinates().GetData().PrepareForInput(DeviceAdapter()), - self.PointIds.PrepareForInput(DeviceAdapter()), - self.CellLower.PrepareForInput(DeviceAdapter()), - self.CellUpper.PrepareForInput(DeviceAdapter())); + self.GetCoordinates().GetData().PrepareForInput(DeviceAdapter(), token), + self.PointIds.PrepareForInput(DeviceAdapter(), token), + self.CellLower.PrepareForInput(DeviceAdapter(), token), + self.CellUpper.PrepareForInput(DeviceAdapter(), token)); handle.Reset(h); return true; } @@ -125,10 +126,11 @@ struct PointLocatorUniformGrid::PrepareExecutionObjectFunctor VTKM_CONT void PointLocatorUniformGrid::PrepareExecutionObject( ExecutionObjectHandleType& execObjHandle, - vtkm::cont::DeviceAdapterId deviceId) const + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const { - const bool success = - vtkm::cont::TryExecuteOnDevice(deviceId, PrepareExecutionObjectFunctor(), *this, execObjHandle); + const bool success = vtkm::cont::TryExecuteOnDevice( + deviceId, PrepareExecutionObjectFunctor(), *this, execObjHandle, token); if (!success) { throwFailedRuntimeDeviceTransfer("PointLocatorUniformGrid", deviceId); diff --git a/vtkm/cont/PointLocatorUniformGrid.h b/vtkm/cont/PointLocatorUniformGrid.h index 26ccbd777..d8cc04cbc 100644 --- a/vtkm/cont/PointLocatorUniformGrid.h +++ b/vtkm/cont/PointLocatorUniformGrid.h @@ -66,7 +66,8 @@ protected: struct PrepareExecutionObjectFunctor; VTKM_CONT void PrepareExecutionObject(ExecutionObjectHandleType& execObjHandle, - vtkm::cont::DeviceAdapterId deviceId) const override; + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const override; bool IsRangeInvalid() const { diff --git a/vtkm/cont/StorageExtrude.h b/vtkm/cont/StorageExtrude.h index debd37dae..99c656e85 100644 --- a/vtkm/cont/StorageExtrude.h +++ b/vtkm/cont/StorageExtrude.h @@ -466,7 +466,8 @@ template class VTKM_ALWAYS_EXPORT ArrayTransfer { using BaseT = typename VecTraits::BaseComponentType; - using TPortalType = decltype(vtkm::cont::ArrayHandle{}.PrepareForInput(Device{})); + using TPortalType = decltype( + vtkm::cont::ArrayHandle{}.PrepareForInput(Device{}, std::declval())); public: using ValueType = T; diff --git a/vtkm/cont/StorageVirtual.hxx b/vtkm/cont/StorageVirtual.hxx index 7622496f2..1275cc33a 100644 --- a/vtkm/cont/StorageVirtual.hxx +++ b/vtkm/cont/StorageVirtual.hxx @@ -143,7 +143,8 @@ struct PortalWrapperToDevice Handle&& handle, vtkm::cont::internal::TransferInfoArray& payload) const { - auto portal = handle.PrepareForInput(device); + vtkm::cont::Token token; + auto portal = handle.PrepareForInput(device, token); using DerivedPortal = vtkm::ArrayPortalWrapper; vtkm::cont::detail::TransferToDevice transfer; return transfer(device, payload, portal); @@ -158,14 +159,16 @@ struct PortalWrapperToDevice using ACCESS_MODE = vtkm::cont::internal::detail::StorageVirtual::OutputMode; if (mode == ACCESS_MODE::WRITE) { - auto portal = handle.PrepareForOutput(numberOfValues, device); + vtkm::cont::Token token; + auto portal = handle.PrepareForOutput(numberOfValues, device, token); using DerivedPortal = vtkm::ArrayPortalWrapper; vtkm::cont::detail::TransferToDevice transfer; return transfer(device, payload, portal); } else { - auto portal = handle.PrepareForInPlace(device); + vtkm::cont::Token token; + auto portal = handle.PrepareForInPlace(device, token); using DerivedPortal = vtkm::ArrayPortalWrapper; vtkm::cont::detail::TransferToDevice transfer; return transfer(device, payload, portal); diff --git a/vtkm/cont/VirtualObjectHandle.h b/vtkm/cont/VirtualObjectHandle.h index 5147f1ef0..85d08c8f7 100644 --- a/vtkm/cont/VirtualObjectHandle.h +++ b/vtkm/cont/VirtualObjectHandle.h @@ -133,7 +133,8 @@ public: /// 2. VirtualObjectHandle is destroyed /// 3. Reset or ReleaseResources is called /// - VTKM_CONT const VirtualBaseType* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId) const + VTKM_CONT const VirtualBaseType* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const { const bool validId = this->Internals->DeviceIdIsValid(deviceId); if (!validId) @@ -142,7 +143,16 @@ public: return nullptr; } - return static_cast(this->Internals->PrepareForExecution(deviceId)); + return static_cast( + this->Internals->PrepareForExecution(deviceId, token)); + } + + VTKM_CONT + VTKM_DEPRECATED(1.6, "PrepareForExecution now requires a vtkm::cont::Token object.") + const VirtualBaseType* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId) const + { + vtkm::cont::Token token; + return this->PrepareForExecution(deviceId, token); } /// Used as part of the \c ExecutionAndControlObjectBase interface. Returns the same pointer diff --git a/vtkm/cont/arg/TransportTagArrayIn.h b/vtkm/cont/arg/TransportTagArrayIn.h index fbcdddcdd..b346cfae2 100644 --- a/vtkm/cont/arg/TransportTagArrayIn.h +++ b/vtkm/cont/arg/TransportTagArrayIn.h @@ -37,7 +37,8 @@ struct Transport { VTKM_IS_ARRAY_HANDLE(ContObjectType); - using ExecObjectType = decltype(std::declval().PrepareForInput(Device())); + using ExecObjectType = decltype( + std::declval().PrepareForInput(Device(), std::declval())); template VTKM_CONT ExecObjectType operator()(const ContObjectType& object, diff --git a/vtkm/cont/arg/TransportTagArrayInOut.h b/vtkm/cont/arg/TransportTagArrayInOut.h index eecd3731b..b70d905f2 100644 --- a/vtkm/cont/arg/TransportTagArrayInOut.h +++ b/vtkm/cont/arg/TransportTagArrayInOut.h @@ -40,7 +40,8 @@ struct Transport().PrepareForInPlace(Device())); + using ExecObjectType = decltype( + std::declval().PrepareForInPlace(Device(), std::declval())); template VTKM_CONT ExecObjectType operator()(ContObjectType& object, diff --git a/vtkm/cont/arg/TransportTagArrayOut.h b/vtkm/cont/arg/TransportTagArrayOut.h index de46691a3..0bc50a8c7 100644 --- a/vtkm/cont/arg/TransportTagArrayOut.h +++ b/vtkm/cont/arg/TransportTagArrayOut.h @@ -40,7 +40,9 @@ struct Transport VTKM_IS_ARRAY_HANDLE(ContObjectType); using ExecObjectType = - decltype(std::declval().PrepareForOutput(vtkm::Id{}, Device())); + decltype(std::declval().PrepareForOutput(vtkm::Id{}, + Device{}, + std::declval())); template VTKM_CONT ExecObjectType operator()(ContObjectType& object, diff --git a/vtkm/cont/arg/TransportTagCellSetIn.h b/vtkm/cont/arg/TransportTagCellSetIn.h index 2ebdd34b0..284e66cd1 100644 --- a/vtkm/cont/arg/TransportTagCellSetIn.h +++ b/vtkm/cont/arg/TransportTagCellSetIn.h @@ -43,8 +43,11 @@ struct Transport().PrepareForInput(Device(), VisitTopology(), IncidentTopology())); + using ExecObjectType = + decltype(std::declval().PrepareForInput(Device(), + VisitTopology(), + IncidentTopology(), + std::declval())); template VTKM_CONT ExecObjectType operator()(const ContObjectType& object, diff --git a/vtkm/cont/arg/TransportTagTopologyFieldIn.h b/vtkm/cont/arg/TransportTagTopologyFieldIn.h index 59229b901..893827fa0 100644 --- a/vtkm/cont/arg/TransportTagTopologyFieldIn.h +++ b/vtkm/cont/arg/TransportTagTopologyFieldIn.h @@ -78,7 +78,8 @@ struct Transport().PrepareForInput(Device())); + using ExecObjectType = decltype( + std::declval().PrepareForInput(Device(), std::declval())); VTKM_CONT ExecObjectType operator()(const ContObjectType& object, diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h index f5c6e5a1a..10131a6ed 100644 --- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h +++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h @@ -1112,9 +1112,11 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + vtkm::Id numBits = bits.GetNumberOfBits(); - auto bitsPortal = bits.PrepareForInput(DeviceAdapterTagCuda{}); - auto indicesPortal = indices.PrepareForOutput(numBits, DeviceAdapterTagCuda{}); + auto bitsPortal = bits.PrepareForInput(DeviceAdapterTagCuda{}, token); + auto indicesPortal = indices.PrepareForOutput(numBits, DeviceAdapterTagCuda{}, token); // Use a uint64 for accumulator, as atomicAdd does not support signed int64. numBits = BitFieldToUnorderedSetPortal(bitsPortal, indicesPortal); @@ -1135,8 +1137,9 @@ public: output.Shrink(inSize); return; } - CopyPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(inSize, DeviceAdapterTagCuda())); + vtkm::cont::Token token; + CopyPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(inSize, DeviceAdapterTagCuda(), token)); } template @@ -1153,9 +1156,11 @@ public: return; } - vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - stencil.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(size, DeviceAdapterTagCuda()), + vtkm::cont::Token token; + + vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + stencil.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(size, DeviceAdapterTagCuda(), token), ::vtkm::NotZeroInitialized()); //yes on the stencil output.Shrink(newSize); } @@ -1174,9 +1179,10 @@ public: output.Shrink(size); return; } - vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - stencil.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(size, DeviceAdapterTagCuda()), + vtk::cont::Token token; + vtkm::Id newSize = CopyIfPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + stencil.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(size, DeviceAdapterTagCuda(), token), unary_predicate); output.Shrink(newSize); } @@ -1230,10 +1236,11 @@ public: output = temp; } } - CopySubRangePortal(input.PrepareForInput(DeviceAdapterTagCuda()), + vtkm::cont::Token token; + CopySubRangePortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), inputStartIndex, numberOfElementsToCopy, - output.PrepareForInPlace(DeviceAdapterTagCuda()), + output.PrepareForInPlace(DeviceAdapterTagCuda(), token), outputIndex); return true; } @@ -1254,9 +1261,10 @@ public: VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); vtkm::Id numberOfValues = values.GetNumberOfValues(); - LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda())); + vtkm::cont::Token token; + LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token)); } template @@ -1268,9 +1276,10 @@ public: VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); vtkm::Id numberOfValues = values.GetNumberOfValues(); - LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), + vtkm::cont::Token token; + LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), binary_compare); } @@ -1280,8 +1289,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values_output.PrepareForInPlace(DeviceAdapterTagCuda())); + vtkm::cont::Token token; + LowerBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values_output.PrepareForInPlace(DeviceAdapterTagCuda(), token)); } template @@ -1294,7 +1304,8 @@ public: { return initialValue; } - return ReducePortal(input.PrepareForInput(DeviceAdapterTagCuda()), initialValue); + vtkm::cont::Token token; + return ReducePortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), initialValue); } template @@ -1309,8 +1320,9 @@ public: { return initialValue; } + vtkm::cont::Token token; return ReducePortal( - input.PrepareForInput(DeviceAdapterTagCuda()), initialValue, binary_functor); + input.PrepareForInput(DeviceAdapterTagCuda(), token), initialValue, binary_functor); } template ::ZeroInitialization(); } @@ -1363,9 +1376,10 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda()); - return ScanExclusivePortal(inputPortal, - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda())); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda(), token); + return ScanExclusivePortal( + inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token)); } template @@ -1379,7 +1393,7 @@ public: const vtkm::Id numberOfValues = input.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); return vtkm::TypeTraits::ZeroInitialization(); } @@ -1387,11 +1401,13 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda()); - return ScanExclusivePortal(inputPortal, - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), - binary_functor, - initialValue); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda(), token); + return ScanExclusivePortal( + inputPortal, + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), + binary_functor, + initialValue); } template @@ -1403,7 +1419,7 @@ public: const vtkm::Id numberOfValues = input.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); return vtkm::TypeTraits::ZeroInitialization(); } @@ -1411,9 +1427,10 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda()); - return ScanInclusivePortal(inputPortal, - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda())); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda(), token); + return ScanInclusivePortal( + inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token)); } template @@ -1426,7 +1443,7 @@ public: const vtkm::Id numberOfValues = input.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); return vtkm::TypeTraits::ZeroInitialization(); } @@ -1434,9 +1451,12 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagCuda(), token); return ScanInclusivePortal( - inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), binary_functor); + inputPortal, + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), + binary_functor); } template @@ -1449,17 +1469,21 @@ public: const vtkm::Id numberOfValues = keys.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); + return; } //We need call PrepareForInput on the input argument before invoking a //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda()); + vtkm::cont::Token token; + auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda(), token); ScanInclusiveByKeyPortal( - keysPortal, valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda())); + keysPortal, + valuesPortal, + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token)); } template (), binary_functor); } @@ -1504,7 +1530,7 @@ public: const vtkm::Id numberOfValues = keys.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); return; } @@ -1512,11 +1538,12 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda()); + vtkm::cont::Token token; + auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda(), token); ScanExclusiveByKeyPortal(keysPortal, valuesPortal, - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), vtkm::TypeTraits::ZeroInitialization(), ::thrust::equal_to(), vtkm::Add()); @@ -1539,7 +1566,7 @@ public: const vtkm::Id numberOfValues = keys.GetNumberOfValues(); if (numberOfValues <= 0) { - output.PrepareForOutput(0, DeviceAdapterTagCuda()); + output.Allocate(0); return; } @@ -1547,11 +1574,12 @@ public: //function. The order of execution of parameters of a function is undefined //so we need to make sure input is called before output, or else in-place //use case breaks. - auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda()); + vtkm::cont::Token token; + auto keysPortal = keys.PrepareForInput(DeviceAdapterTagCuda(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTagCuda(), token); ScanExclusiveByKeyPortal(keysPortal, valuesPortal, - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), initialValue, ::thrust::equal_to(), binary_functor); @@ -1691,7 +1719,8 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - SortPortal(values.PrepareForInPlace(DeviceAdapterTagCuda())); + vtkm::cont::Token token; + SortPortal(values.PrepareForInPlace(DeviceAdapterTagCuda(), token)); } template @@ -1700,7 +1729,8 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - SortPortal(values.PrepareForInPlace(DeviceAdapterTagCuda()), binary_compare); + vtkm::cont::Token token; + SortPortal(values.PrepareForInPlace(DeviceAdapterTagCuda(), token), binary_compare); } template @@ -1709,8 +1739,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTagCuda()), - values.PrepareForInPlace(DeviceAdapterTagCuda())); + vtkm::cont::Token token; + SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTagCuda(), token), + values.PrepareForInPlace(DeviceAdapterTagCuda(), token)); } template @@ -1720,8 +1751,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTagCuda()), - values.PrepareForInPlace(DeviceAdapterTagCuda()), + vtkm::cont::Token token; + SortByKeyPortal(keys.PrepareForInPlace(DeviceAdapterTagCuda(), token), + values.PrepareForInPlace(DeviceAdapterTagCuda(), token), binary_compare); } @@ -1730,7 +1762,8 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - vtkm::Id newSize = UniquePortal(values.PrepareForInPlace(DeviceAdapterTagCuda())); + vtkm::cont::Token token; + vtkm::Id newSize = UniquePortal(values.PrepareForInPlace(DeviceAdapterTagCuda(), token)); values.Shrink(newSize); } @@ -1741,8 +1774,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; vtkm::Id newSize = - UniquePortal(values.PrepareForInPlace(DeviceAdapterTagCuda()), binary_compare); + UniquePortal(values.PrepareForInPlace(DeviceAdapterTagCuda(), token), binary_compare); values.Shrink(newSize); } @@ -1755,9 +1789,10 @@ public: VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); vtkm::Id numberOfValues = values.GetNumberOfValues(); - UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda())); + vtkm::cont::Token token; + UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token)); } template @@ -1769,9 +1804,10 @@ public: VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); vtkm::Id numberOfValues = values.GetNumberOfValues(); - UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values.PrepareForInput(DeviceAdapterTagCuda()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()), + vtkm::cont::Token token; + UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values.PrepareForInput(DeviceAdapterTagCuda(), token), + output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), binary_compare); } @@ -1781,8 +1817,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda()), - values_output.PrepareForInPlace(DeviceAdapterTagCuda())); + vtkm::cont::Token token; + UpperBoundsPortal(input.PrepareForInput(DeviceAdapterTagCuda(), token), + values_output.PrepareForInPlace(DeviceAdapterTagCuda(), token)); } VTKM_CONT static void Synchronize() diff --git a/vtkm/cont/internal/ArrayHandleBasicImpl.h b/vtkm/cont/internal/ArrayHandleBasicImpl.h index 0034828d3..a69afae0f 100644 --- a/vtkm/cont/internal/ArrayHandleBasicImpl.h +++ b/vtkm/cont/internal/ArrayHandleBasicImpl.h @@ -403,25 +403,25 @@ public: DeviceAdapterTag device, vtkm::cont::Token& token); - // TODO: Deprecate these template - VTKM_CONT typename ExecutionTypes::PortalConst PrepareForInput( - DeviceAdapterTag device) const + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInput now requires a vtkm::cont::Token object.") + typename ExecutionTypes::PortalConst + PrepareForInput(DeviceAdapterTag device) const { vtkm::cont::Token token; return this->PrepareForInput(device, token); } template - VTKM_CONT typename ExecutionTypes::Portal PrepareForOutput( - vtkm::Id numVals, - DeviceAdapterTag device) + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForOutput now requires a vtkm::cont::Token object.") + typename ExecutionTypes::Portal + PrepareForOutput(vtkm::Id numVals, DeviceAdapterTag device) { vtkm::cont::Token token; return this->PrepareForOutput(numVals, device, token); } template - VTKM_CONT typename ExecutionTypes::Portal PrepareForInPlace( - DeviceAdapterTag device) + VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInPlace now requires a vtkm::cont::Token object.") + typename ExecutionTypes::Portal PrepareForInPlace(DeviceAdapterTag device) { vtkm::cont::Token token; return this->PrepareForInPlace(device, token); diff --git a/vtkm/cont/internal/ConnectivityExplicitInternals.h b/vtkm/cont/internal/ConnectivityExplicitInternals.h index bac581e6f..c0a1d5d3d 100644 --- a/vtkm/cont/internal/ConnectivityExplicitInternals.h +++ b/vtkm/cont/internal/ConnectivityExplicitInternals.h @@ -172,13 +172,16 @@ void ComputeRConnTable(RConnTableT& rConnTable, auto& rOffsets = rConnTable.Offsets; const vtkm::Id rConnSize = conn.GetNumberOfValues(); - const auto offInPortal = connTable.Offsets.PrepareForInput(Device{}); + { + vtkm::cont::Token token; + const auto offInPortal = connTable.Offsets.PrepareForInput(Device{}, token); - PassThrough idxCalc{}; - ConnIdxToCellIdCalc cellIdCalc{ offInPortal }; + PassThrough idxCalc{}; + ConnIdxToCellIdCalc cellIdCalc{ offInPortal }; - vtkm::cont::internal::ReverseConnectivityBuilder builder; - builder.Run(conn, rConn, rOffsets, idxCalc, cellIdCalc, numberOfPoints, rConnSize, Device()); + vtkm::cont::internal::ReverseConnectivityBuilder builder; + builder.Run(conn, rConn, rOffsets, idxCalc, cellIdCalc, numberOfPoints, rConnSize, Device()); + } rConnTable.Shapes = vtkm::cont::make_ArrayHandleConstant( static_cast(CELL_SHAPE_VERTEX), numberOfPoints); diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h index 9cf686b11..5d576c2eb 100644 --- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h +++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h @@ -101,9 +101,11 @@ private: { using OutputArrayType = vtkm::cont::ArrayHandle; + vtkm::cont::Token token; + OutputArrayType output; - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(1, DeviceAdapterTag()); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(1, DeviceAdapterTag(), token); CopyKernel kernel( inputPortal, outputPortal, index); @@ -125,8 +127,10 @@ public: vtkm::Id numBits = bits.GetNumberOfBits(); - auto bitsPortal = bits.PrepareForInput(DeviceAdapterTag{}); - auto indicesPortal = indices.PrepareForOutput(numBits, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto bitsPortal = bits.PrepareForInput(DeviceAdapterTag{}, token); + auto indicesPortal = indices.PrepareForOutput(numBits, DeviceAdapterTag{}, token); std::atomic popCount; popCount.store(0, std::memory_order_seq_cst); @@ -137,6 +141,8 @@ public: DerivedAlgorithm::Schedule(functor, functor.GetNumberOfInstances()); DerivedAlgorithm::Synchronize(); + token.DetachFromAll(); + numBits = static_cast(popCount.load(std::memory_order_seq_cst)); indices.Shrink(numBits); @@ -151,9 +157,11 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + const vtkm::Id inSize = input.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTag()); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTag(), token); CopyKernel kernel(inputPortal, outputPortal); DerivedAlgorithm::Schedule(kernel, inSize); @@ -175,26 +183,36 @@ public: using IndexArrayType = vtkm::cont::ArrayHandle; IndexArrayType indices; - auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag()); - auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag()); + { + vtkm::cont::Token token; - StencilToIndexFlagKernel - indexKernel(stencilPortal, indexPortal, unary_predicate); + auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag(), token); + auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag(), token); - DerivedAlgorithm::Schedule(indexKernel, arrayLength); + StencilToIndexFlagKernel + indexKernel(stencilPortal, indexPortal, unary_predicate); + + DerivedAlgorithm::Schedule(indexKernel, arrayLength); + } vtkm::Id outArrayLength = DerivedAlgorithm::ScanExclusive(indices, indices); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag()); + { + vtkm::cont::Token token; - CopyIfKernel - copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate); - DerivedAlgorithm::Schedule(copyKernel, arrayLength); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag(), token); + auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag(), token); + + CopyIfKernel + copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate); + DerivedAlgorithm::Schedule(copyKernel, arrayLength); + } } template @@ -260,8 +278,10 @@ public: } } - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForInPlace(DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForInPlace(DeviceAdapterTag(), token); CopyKernel kernel( inputPortal, outputPortal, inputStartIndex, outputIndex); @@ -275,7 +295,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - auto bitsPortal = bits.PrepareForInput(DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto bitsPortal = bits.PrepareForInput(DeviceAdapterTag{}, token); std::atomic popCount; popCount.store(0, std::memory_order_relaxed); @@ -301,7 +323,9 @@ public: return; } - auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}, token); using WordType = typename vtkm::cont::BitField::template ExecutionTypes::WordTypePreferred; @@ -325,7 +349,9 @@ public: return; } - auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}, token); using WordType = typename vtkm::cont::BitField::template ExecutionTypes::WordTypePreferred; @@ -352,7 +378,9 @@ public: return; } - auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}, token); // If less than 32 bits, repeat the word until we get a 32 bit pattern. // Using this for the pattern prevents races while writing small numbers @@ -381,7 +409,9 @@ public: return; } - auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = bits.PrepareForOutput(numBits, DeviceAdapterTag{}, token); // If less than 32 bits, repeat the word until we get a 32 bit pattern. // Using this for the pattern prevents races while writing small numbers @@ -409,7 +439,9 @@ public: return; } - auto portal = handle.PrepareForOutput(numValues, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = handle.PrepareForOutput(numValues, DeviceAdapterTag{}, token); FillArrayHandleFunctor functor{ portal, value }; DerivedAlgorithm::Schedule(functor, numValues); } @@ -428,7 +460,9 @@ public: return; } - auto portal = handle.PrepareForOutput(numValues, DeviceAdapterTag{}); + vtkm::cont::Token token; + + auto portal = handle.PrepareForOutput(numValues, DeviceAdapterTag{}, token); FillArrayHandleFunctor functor{ portal, value }; DerivedAlgorithm::Schedule(functor, numValues); } @@ -444,9 +478,11 @@ public: vtkm::Id arraySize = values.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag(), token); LowerBoundsKernel kernel( inputPortal, valuesPortal, outputPortal); @@ -464,9 +500,11 @@ public: vtkm::Id arraySize = values.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag(), token); LowerBoundsComparisonKernel kernel( inputPortal, initialValue, binary_functor); @@ -605,8 +645,9 @@ public: vtkm::cont::ArrayHandle keystate; { - auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + vtkm::cont::Token token; + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag(), token); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag(), token); ReduceStencilGeneration kernel( inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); @@ -668,8 +709,10 @@ public: vtkm::cont::ArrayHandle inclusiveScan; T result = DerivedAlgorithm::ScanInclusive(input, inclusiveScan, binaryFunctor); - auto inputPortal = inclusiveScan.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = inclusiveScan.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag(), token); InclusiveToExclusiveKernel inclusiveToExclusive(inputPortal, outputPortal, binaryFunctor, initialValue); @@ -710,8 +753,10 @@ public: vtkm::cont::ArrayHandle inclusiveScan; T result = DerivedAlgorithm::ScanInclusive(input, inclusiveScan, binaryFunctor); - auto inputPortal = inclusiveScan.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(numValues + 1, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = inclusiveScan.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(numValues + 1, DeviceAdapterTag(), token); InclusiveToExtendedKernel inclusiveToExtended(inputPortal, @@ -760,7 +805,7 @@ public: } else if (numberOfKeys == 1) { - output.PrepareForOutput(1, DeviceAdapterTag()); + output.Allocate(1); output.GetPortalControl().Set(0, initialValue); return; } @@ -772,8 +817,9 @@ public: vtkm::cont::ArrayHandle keystate; { - auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + vtkm::cont::Token token; + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag(), token); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag(), token); ReduceStencilGeneration kernel( inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); @@ -782,9 +828,10 @@ public: // 2. Shift input and initialize elements at head flags position to initValue vtkm::cont::ArrayHandle temp; { - auto inputPortal = values.PrepareForInput(DeviceAdapterTag()); - auto keyStatePortal = keystate.PrepareForInput(DeviceAdapterTag()); - auto tempPortal = temp.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + vtkm::cont::Token token; + auto inputPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto keyStatePortal = keystate.PrepareForInput(DeviceAdapterTag(), token); + auto tempPortal = temp.PrepareForOutput(numberOfKeys, DeviceAdapterTag(), token); ShiftCopyAndInit::ZeroInitialization(); } - auto portal = output.PrepareForInPlace(DeviceAdapterTag()); + vtkm::cont::Token token; + + auto portal = output.PrepareForInPlace(DeviceAdapterTag(), token); using ScanKernelType = ScanKernel; @@ -945,8 +994,9 @@ public: vtkm::cont::ArrayHandle keystate; { - auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + vtkm::cont::Token token; + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag(), token); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag(), token); ReduceStencilGeneration kernel( inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); @@ -992,7 +1042,9 @@ public: } numThreads /= 2; - auto portal = values.PrepareForInPlace(DeviceAdapterTag()); + vtkm::cont::Token token; + + auto portal = values.PrepareForInPlace(DeviceAdapterTag(), token); using MergeKernel = BitonicSortMergeKernel; using CrossoverKernel = BitonicSortCrossoverKernel; @@ -1066,9 +1118,11 @@ public: return; } - auto input1Portal = input1.PrepareForInput(DeviceAdapterTag()); - auto input2Portal = input2.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto input1Portal = input1.PrepareForInput(DeviceAdapterTag(), token); + auto input2Portal = input2.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag(), token); BinaryTransformKernel; WrappedBOpType wrappedCompare(binary_compare); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); - auto stencilPortal = stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag()); - ClassifyUniqueComparisonKernel - classifyKernel(valuesPortal, stencilPortal, wrappedCompare); + { + vtkm::cont::Token token; + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto stencilPortal = stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag(), token); + ClassifyUniqueComparisonKernel + classifyKernel(valuesPortal, stencilPortal, wrappedCompare); - DerivedAlgorithm::Schedule(classifyKernel, inputSize); + DerivedAlgorithm::Schedule(classifyKernel, inputSize); + } vtkm::cont::ArrayHandle outputArray; @@ -1127,9 +1186,11 @@ public: vtkm::Id arraySize = values.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag(), token); UpperBoundsKernel kernel( inputPortal, valuesPortal, outputPortal); @@ -1146,9 +1207,11 @@ public: vtkm::Id arraySize = values.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); - auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); - auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag(), token); UpperBoundsKernelComparisonKernel; - auto connPortal = conn.PrepareForInput(Device()); + vtkm::cont::Token connToken; + auto connPortal = conn.PrepareForInput(Device(), connToken); auto zeros = vtkm::cont::make_ArrayHandleConstant(vtkm::IdComponent{ 0 }, numberOfPoints); // Compute RConn offsets by atomically building a histogram and doing an @@ -163,7 +164,8 @@ public: { // Build histogram: vtkm::cont::AtomicArray atomicCounter{ rNumIndices }; - auto ac = atomicCounter.PrepareForExecution(Device()); + vtkm::cont::Token token; + auto ac = atomicCounter.PrepareForExecution(Device(), token); using BuildHisto = rcb::BuildHistogram; BuildHisto histoGen{ ac, connPortal, rConnToConnCalc }; @@ -194,9 +196,10 @@ public: // (out) RConn: | 0 1 2 | 0 1 | 0 | 1 2 3 | 2 3 | 3 | { vtkm::cont::AtomicArray atomicCounter{ rNumIndices }; - auto ac = atomicCounter.PrepareForExecution(Device()); - auto rOffsetPortal = rOffsets.PrepareForInput(Device()); - auto rConnPortal = rConn.PrepareForOutput(rConnSize, Device()); + vtkm::cont::Token token; + auto ac = atomicCounter.PrepareForExecution(Device(), token); + auto rOffsetPortal = rOffsets.PrepareForInput(Device(), token); + auto rConnPortal = rConn.PrepareForOutput(rConnSize, Device(), token); using GenRConnT = rcb::GenerateRConn #include +#include #include #include @@ -45,7 +46,8 @@ struct VirtualObjectTransfer /// PrepareForExecution). If the \c updateData flag is false and the object was already /// transferred previously, the previously created object is returned. /// - VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool updateData); + VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool updateData, + vtkm::cont::Token& token); /// \brief Frees up any resources in the execution environment. /// @@ -61,7 +63,9 @@ class VTKM_CONT_EXPORT TransferInterface public: VTKM_CONT virtual ~TransferInterface(); - VTKM_CONT virtual const vtkm::VirtualObjectBase* PrepareForExecution(vtkm::Id) = 0; + VTKM_CONT virtual const vtkm::VirtualObjectBase* PrepareForExecution( + vtkm::Id hostModifiedCount, + vtkm::cont::Token& token) = 0; VTKM_CONT virtual void ReleaseResources() = 0; }; @@ -75,10 +79,12 @@ public: { } - VTKM_CONT const vtkm::VirtualObjectBase* PrepareForExecution(vtkm::Id hostModifiedCount) override + VTKM_CONT const vtkm::VirtualObjectBase* PrepareForExecution(vtkm::Id hostModifiedCount, + vtkm::cont::Token& token) override { bool updateData = (this->LastModifiedCount != hostModifiedCount); - const vtkm::VirtualObjectBase* executionObject = this->Transfer.PrepareForExecution(updateData); + const vtkm::VirtualObjectBase* executionObject = + this->Transfer.PrepareForExecution(updateData, token); this->LastModifiedCount = hostModifiedCount; return executionObject; } @@ -138,12 +144,13 @@ struct VTKM_CONT_EXPORT TransferState } } - const vtkm::VirtualObjectBase* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId) const + const vtkm::VirtualObjectBase* PrepareForExecution(vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const { //make sure the device is up to date auto index = static_cast(deviceId.GetValue()); vtkm::Id count = this->HostPointer->GetModifiedCount(); - return this->DeviceTransferState[index]->PrepareForExecution(count); + return this->DeviceTransferState[index]->PrepareForExecution(count, token); } vtkm::VirtualObjectBase* HostPtr() const { return this->HostPointer; } diff --git a/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h index 49f8480db..35570cdae 100644 --- a/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h +++ b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h @@ -13,6 +13,8 @@ #include #include +#include + namespace vtkm { namespace cont @@ -28,7 +30,8 @@ struct VirtualObjectTransferShareWithControl { } - VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool vtkmNotUsed(updateData)) + VTKM_CONT const VirtualDerivedType* PrepareForExecution(bool vtkmNotUsed(updateData), + vtkm::cont::Token&) { return this->VirtualObject; } diff --git a/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h b/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h index 47804292b..8f87ea5b8 100644 --- a/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h +++ b/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h @@ -54,8 +54,9 @@ public: output.Allocate(0); return; } - auto inputPortal = input.PrepareForInput(DevTag()); - auto outputPortal = output.PrepareForOutput(inSize, DevTag()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DevTag(), token); + auto outputPortal = output.PrepareForOutput(inSize, DevTag(), token); CopyHelper(inputPortal, outputPortal, 0, 0, inSize); } @@ -86,9 +87,10 @@ public: output.Allocate(0); return; } - auto inputPortal = input.PrepareForInput(DevTag()); - auto stencilPortal = stencil.PrepareForInput(DevTag()); - auto outputPortal = output.PrepareForOutput(inSize, DevTag()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DevTag(), token); + auto stencilPortal = stencil.PrepareForInput(DevTag(), token); + auto outputPortal = output.PrepareForOutput(inSize, DevTag(), token); auto inIter = vtkm::cont::ArrayPortalToIteratorBegin(inputPortal); auto stencilIter = vtkm::cont::ArrayPortalToIteratorBegin(stencilPortal); @@ -113,6 +115,7 @@ public: } vtkm::Id numValues = helper.Reduce(outIter); + token.DetachFromAll(); output.Shrink(numValues); } @@ -168,8 +171,9 @@ public: } } - auto inputPortal = input.PrepareForInput(DevTag()); - auto outputPortal = output.PrepareForInPlace(DevTag()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DevTag(), token); + auto outputPortal = output.PrepareForInPlace(DevTag(), token); CopyHelper(inputPortal, outputPortal, inputStartIndex, outputIndex, numberOfValuesToCopy); @@ -193,7 +197,8 @@ public: using namespace vtkm::cont::openmp; - auto portal = input.PrepareForInput(DevTag()); + vtkm::cont::Token token; + auto portal = input.PrepareForInput(DevTag(), token); const OpenMPReductionSupported::type> fastPath; return ReduceHelper::Execute(portal, initialValue, binary_functor, fastPath); @@ -238,13 +243,15 @@ public: return vtkm::TypeTraits::ZeroInitialization(); } - using InPortalT = decltype(input.PrepareForInput(DevTag())); - using OutPortalT = decltype(output.PrepareForOutput(0, DevTag())); + vtkm::cont::Token token; + using InPortalT = decltype(input.PrepareForInput(DevTag(), token)); + using OutPortalT = decltype(output.PrepareForOutput(0, DevTag(), token)); using Impl = openmp::ScanInclusiveHelper; vtkm::Id numVals = input.GetNumberOfValues(); - Impl impl( - input.PrepareForInput(DevTag()), output.PrepareForOutput(numVals, DevTag()), binaryFunctor); + Impl impl(input.PrepareForInput(DevTag(), token), + output.PrepareForOutput(numVals, DevTag(), token), + binaryFunctor); return impl.Execute(vtkm::Id2(0, numVals)); } @@ -271,13 +278,14 @@ public: return initialValue; } - using InPortalT = decltype(input.PrepareForInput(DevTag())); - using OutPortalT = decltype(output.PrepareForOutput(0, DevTag())); + vtkm::cont::Token token; + using InPortalT = decltype(input.PrepareForInput(DevTag(), token)); + using OutPortalT = decltype(output.PrepareForOutput(0, DevTag(), token)); using Impl = openmp::ScanExclusiveHelper; vtkm::Id numVals = input.GetNumberOfValues(); - Impl impl(input.PrepareForInput(DevTag()), - output.PrepareForOutput(numVals, DevTag()), + Impl impl(input.PrepareForInput(DevTag(), token), + output.PrepareForOutput(numVals, DevTag(), token), binaryFunctor, initialValue); @@ -339,7 +347,8 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - auto portal = values.PrepareForInPlace(DevTag()); + vtkm::cont::Token token; + auto portal = values.PrepareForInPlace(DevTag(), token); auto iter = vtkm::cont::ArrayPortalToIteratorBegin(portal); using IterT = typename std::decay::type; @@ -347,6 +356,7 @@ public: Uniqifier uniquifier(iter, portal.GetNumberOfValues(), binary_compare); vtkm::Id outSize = uniquifier.Execute(); + token.DetachFromAll(); values.Shrink(outSize); } diff --git a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h index 3783ce4ec..e2791d2e3 100644 --- a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h +++ b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h @@ -71,9 +71,11 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + const vtkm::Id inSize = input.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial()); - auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTagSerial()); + auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial(), token); + auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTagSerial(), token); if (inSize <= 0) { @@ -111,9 +113,10 @@ public: vtkm::Id inputSize = input.GetNumberOfValues(); VTKM_ASSERT(inputSize == stencil.GetNumberOfValues()); - auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial()); - auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTagSerial()); - auto outputPortal = output.PrepareForOutput(inputSize, DeviceAdapterTagSerial()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial(), token); + auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTagSerial(), token); + auto outputPortal = output.PrepareForOutput(inputSize, DeviceAdapterTagSerial(), token); vtkm::Id readPos = 0; vtkm::Id writePos = 0; @@ -180,8 +183,9 @@ public: } } - auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial()); - auto outputPortal = output.PrepareForInPlace(DeviceAdapterTagSerial()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagSerial(), token); + auto outputPortal = output.PrepareForInPlace(DeviceAdapterTagSerial(), token); auto inIter = vtkm::cont::ArrayPortalToIteratorBegin(inputPortal); auto outIter = vtkm::cont::ArrayPortalToIteratorBegin(outputPortal); @@ -211,8 +215,10 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + internal::WrappedBinaryOperator wrappedOp(binary_functor); - auto inputPortal = input.PrepareForInput(Device()); + auto inputPortal = input.PrepareForInput(Device(), token); return std::accumulate(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal), vtkm::cont::ArrayPortalToIteratorEnd(inputPortal), initialValue, @@ -234,8 +240,10 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - auto keysPortalIn = keys.PrepareForInput(Device()); - auto valuesPortalIn = values.PrepareForInput(Device()); + vtkm::cont::Token token; + + auto keysPortalIn = keys.PrepareForInput(Device(), token); + auto valuesPortalIn = values.PrepareForInput(Device(), token); const vtkm::Id numberOfKeys = keys.GetNumberOfValues(); VTKM_ASSERT(numberOfKeys == values.GetNumberOfValues()); @@ -246,8 +254,8 @@ public: return; } - auto keysPortalOut = keys_output.PrepareForOutput(numberOfKeys, Device()); - auto valuesPortalOut = values_output.PrepareForOutput(numberOfKeys, Device()); + auto keysPortalOut = keys_output.PrepareForOutput(numberOfKeys, Device(), token); + auto valuesPortalOut = values_output.PrepareForOutput(numberOfKeys, Device(), token); vtkm::Id writePos = 0; vtkm::Id readPos = 0; @@ -295,8 +303,10 @@ public: vtkm::Id numberOfValues = input.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(Device()); - auto outputPortal = output.PrepareForOutput(numberOfValues, Device()); + vtkm::cont::Token token; + + auto inputPortal = input.PrepareForInput(Device(), token); + auto outputPortal = output.PrepareForOutput(numberOfValues, Device(), token); if (numberOfValues <= 0) { @@ -333,8 +343,9 @@ public: vtkm::Id numberOfValues = input.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(Device()); - auto outputPortal = output.PrepareForOutput(numberOfValues, Device()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(Device(), token); + auto outputPortal = output.PrepareForOutput(numberOfValues, Device(), token); if (numberOfValues <= 0) { @@ -411,9 +422,11 @@ private: const vtkm::Id n = values.GetNumberOfValues(); VTKM_ASSERT(n == index.GetNumberOfValues()); - auto valuesPortal = values.PrepareForInput(Device()); - auto indexPortal = index.PrepareForInput(Device()); - auto valuesOutPortal = values_out.PrepareForOutput(n, Device()); + vtkm::cont::Token token; + + auto valuesPortal = values.PrepareForInput(Device(), token); + auto indexPortal = index.PrepareForInput(Device(), token); + auto valuesOutPortal = values_out.PrepareForOutput(n, Device(), token); for (vtkm::Id i = 0; i < n; i++) { @@ -489,7 +502,9 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - auto arrayPortal = values.PrepareForInPlace(Device()); + vtkm::cont::Token token; + + auto arrayPortal = values.PrepareForInPlace(Device(), token); vtkm::cont::ArrayPortalToIterators iterators(arrayPortal); internal::WrappedBinaryOperator wrappedCompare(binary_compare); @@ -510,12 +525,15 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - auto arrayPortal = values.PrepareForInPlace(Device()); + vtkm::cont::Token token; + auto arrayPortal = values.PrepareForInPlace(Device(), token); vtkm::cont::ArrayPortalToIterators iterators(arrayPortal); internal::WrappedBinaryOperator wrappedCompare(binary_compare); auto end = std::unique(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare); - values.Shrink(static_cast(end - iterators.GetBegin())); + vtkm::Id newSize = static_cast(end - iterators.GetBegin()); + token.DetachFromAll(); + values.Shrink(newSize); } VTKM_CONT static void Synchronize() diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h index 18c778a43..2f609ea7c 100644 --- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h +++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h @@ -43,9 +43,11 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + const vtkm::Id inSize = input.GetNumberOfValues(); - auto inputPortal = input.PrepareForInput(DeviceAdapterTagTBB()); - auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTagTBB()); + auto inputPortal = input.PrepareForInput(DeviceAdapterTagTBB(), token); + auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTagTBB(), token); tbb::CopyPortals(inputPortal, outputPortal, 0, 0, inSize); } @@ -69,13 +71,16 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; + vtkm::Id inputSize = input.GetNumberOfValues(); VTKM_ASSERT(inputSize == stencil.GetNumberOfValues()); vtkm::Id outputSize = - tbb::CopyIfPortals(input.PrepareForInput(DeviceAdapterTagTBB()), - stencil.PrepareForInput(DeviceAdapterTagTBB()), - output.PrepareForOutput(inputSize, DeviceAdapterTagTBB()), + tbb::CopyIfPortals(input.PrepareForInput(DeviceAdapterTagTBB(), token), + stencil.PrepareForInput(DeviceAdapterTagTBB(), token), + output.PrepareForOutput(inputSize, DeviceAdapterTagTBB(), token), unary_predicate); + token.DetachFromAll(); output.Shrink(outputSize); } @@ -129,8 +134,9 @@ public: } } - auto inputPortal = input.PrepareForInput(DeviceAdapterTagTBB()); - auto outputPortal = output.PrepareForInPlace(DeviceAdapterTagTBB()); + vtkm::cont::Token token; + auto inputPortal = input.PrepareForInput(DeviceAdapterTagTBB(), token); + auto outputPortal = output.PrepareForInPlace(DeviceAdapterTagTBB(), token); tbb::CopyPortals( inputPortal, outputPortal, inputStartIndex, outputIndex, numberOfElementsToCopy); @@ -153,8 +159,10 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); - return tbb::ReducePortals( - input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), initialValue, binary_functor); + vtkm::cont::Token token; + return tbb::ReducePortals(input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + initialValue, + binary_functor); } template ::ZeroInitialization()); } @@ -230,9 +244,10 @@ public: { VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::Token token; return tbb::ScanExclusivePortals( - input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), - output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB()), + input.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + output.PrepareForOutput(input.GetNumberOfValues(), vtkm::cont::DeviceAdapterTagTBB(), token), binary_functor, initialValue); } @@ -307,8 +322,9 @@ public: VTKM_CONT static void Unique(vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { + vtkm::cont::Token token; vtkm::Id outputSize = - tbb::UniquePortals(values.PrepareForInPlace(DeviceAdapterTagTBB()), binary_compare); + tbb::UniquePortals(values.PrepareForInPlace(DeviceAdapterTagTBB(), token), binary_compare); values.Shrink(outputSize); } diff --git a/vtkm/cont/tbb/internal/ParallelSortTBB.h b/vtkm/cont/tbb/internal/ParallelSortTBB.h index 8a6302357..eec6a5df7 100644 --- a/vtkm/cont/tbb/internal/ParallelSortTBB.h +++ b/vtkm/cont/tbb/internal/ParallelSortTBB.h @@ -51,7 +51,8 @@ void parallel_sort(HandleType& values, BinaryCompare binary_compare, vtkm::cont::internal::radix::PSortTag) { - auto arrayPortal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagTBB()); + vtkm::cont::Token token; + auto arrayPortal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagTBB(), token); using IteratorsType = vtkm::cont::ArrayPortalToIterators; IteratorsType iterators(arrayPortal); @@ -106,10 +107,11 @@ void parallel_sort_bykey(vtkm::cont::ArrayHandle& keys, const vtkm::Id size = values.GetNumberOfValues(); { + vtkm::cont::Token token; auto handle = ArrayHandleIndex(keys.GetNumberOfValues()); - auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB()); + auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB(), token); auto outputPortal = - indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB()); + indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB(), token); tbb::CopyPortals(inputPortal, outputPortal, 0, 0, keys.GetNumberOfValues()); } @@ -118,14 +120,19 @@ void parallel_sort_bykey(vtkm::cont::ArrayHandle& keys, vtkm::cont::internal::KeyCompare(binary_compare), PSortTag()); - tbb::ScatterPortal(values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), - indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), - valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB())); + { + vtkm::cont::Token token; + tbb::ScatterPortal( + values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB(), token)); + } { - auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB()); + vtkm::cont::Token token; + auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB(), token); auto outputPortal = - values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB()); + values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB(), token); tbb::CopyPortals(inputPortal, outputPortal, 0, 0, valuesScattered.GetNumberOfValues()); } } @@ -172,10 +179,11 @@ void parallel_sort_bykey(vtkm::cont::ArrayHandle& keys, const vtkm::Id size = values.GetNumberOfValues(); { + vtkm::cont::Token token; auto handle = ArrayHandleIndex(keys.GetNumberOfValues()); - auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB()); + auto inputPortal = handle.PrepareForInput(DeviceAdapterTagTBB(), token); auto outputPortal = - indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB()); + indexArray.PrepareForOutput(keys.GetNumberOfValues(), DeviceAdapterTagTBB(), token); tbb::CopyPortals(inputPortal, outputPortal, 0, 0, keys.GetNumberOfValues()); } @@ -191,14 +199,19 @@ void parallel_sort_bykey(vtkm::cont::ArrayHandle& keys, vtkm::cont::internal::radix::PSortTag{}); } - tbb::ScatterPortal(values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), - indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB()), - valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB())); + { + vtkm::cont::Token token; + tbb::ScatterPortal( + values.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + indexArray.PrepareForInput(vtkm::cont::DeviceAdapterTagTBB(), token), + valuesScattered.PrepareForOutput(size, vtkm::cont::DeviceAdapterTagTBB(), token)); + } { - auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB()); + vtkm::cont::Token token; + auto inputPortal = valuesScattered.PrepareForInput(DeviceAdapterTagTBB(), token); auto outputPortal = - values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB()); + values.PrepareForOutput(valuesScattered.GetNumberOfValues(), DeviceAdapterTagTBB(), token); tbb::CopyPortals(inputPortal, outputPortal, 0, 0, valuesScattered.GetNumberOfValues()); } } diff --git a/vtkm/cont/testing/TestingArrayHandles.h b/vtkm/cont/testing/TestingArrayHandles.h index 315bc9ffb..a45ec24d1 100644 --- a/vtkm/cont/testing/TestingArrayHandles.h +++ b/vtkm/cont/testing/TestingArrayHandles.h @@ -141,6 +141,7 @@ private: arrayHandle = vtkm::cont::ArrayHandle(); VTKM_TEST_ASSERT(arrayHandle.GetPortalConstControl().GetNumberOfValues() == 0, "Uninitialized array does not give portal with zero values."); + vtkm::cont::Token token; arrayHandle = vtkm::cont::ArrayHandle(); arrayHandle.Shrink(0); arrayHandle = vtkm::cont::ArrayHandle(); @@ -148,11 +149,11 @@ private: arrayHandle = vtkm::cont::ArrayHandle(); arrayHandle.ReleaseResources(); arrayHandle = vtkm::cont::make_ArrayHandle(std::vector()); - arrayHandle.PrepareForInput(DeviceAdapterTag()); + arrayHandle.PrepareForInput(DeviceAdapterTag(), token); arrayHandle = vtkm::cont::ArrayHandle(); - arrayHandle.PrepareForInPlace(DeviceAdapterTag()); + arrayHandle.PrepareForInPlace(DeviceAdapterTag(), token); arrayHandle = vtkm::cont::ArrayHandle(); - arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()); + arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token); } }; @@ -177,9 +178,11 @@ private: std::cout << "Check out execution array behavior." << std::endl; { //as input + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst executionPortal; - executionPortal = arrayHandle.PrepareForInput(DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForInput(DeviceAdapterTag(), token); + token.DetachFromAll(); //use a worklet to verify the input transfer worked properly vtkm::cont::ArrayHandle result; @@ -189,9 +192,11 @@ private: std::cout << "Check out inplace." << std::endl; { //as inplace + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal executionPortal; - executionPortal = arrayHandle.PrepareForInPlace(DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForInPlace(DeviceAdapterTag(), token); + token.DetachFromAll(); //use a worklet to verify the inplace transfer worked properly vtkm::cont::ArrayHandle result; @@ -202,9 +207,10 @@ private: std::cout << "Check out output." << std::endl; { //as output with same length as user provided. This should work //as no new memory needs to be allocated + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal executionPortal; - executionPortal = arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token); //we can't verify output contents as those aren't fetched, we //can just make sure the allocation didn't throw an exception @@ -217,7 +223,9 @@ private: { //you should not be able to allocate a size larger than the //user provided and get the results - arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag()); + vtkm::cont::Token token; + arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag(), token); + token.DetachFromAll(); arrayHandle.GetPortalControl(); } catch (vtkm::cont::Error&) @@ -256,9 +264,11 @@ private: std::cout << "Check out execution array behavior." << std::endl; { //as input + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst executionPortal; - executionPortal = arrayHandle.PrepareForInput(DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForInput(DeviceAdapterTag(), token); + token.DetachFromAll(); //use a worklet to verify the input transfer worked properly vtkm::cont::ArrayHandle result; @@ -268,9 +278,11 @@ private: std::cout << "Check out inplace." << std::endl; { //as inplace + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal executionPortal; - executionPortal = arrayHandle.PrepareForInPlace(DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForInPlace(DeviceAdapterTag(), token); + token.DetachFromAll(); //use a worklet to verify the inplace transfer worked properly vtkm::cont::ArrayHandle result; @@ -281,9 +293,11 @@ private: std::cout << "Check out output." << std::endl; { //as output with same length as user provided. This should work //as no new memory needs to be allocated + vtkm::cont::Token token; typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal executionPortal; - executionPortal = arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()); + executionPortal = arrayHandle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token); + token.DetachFromAll(); //we can't verify output contents as those aren't fetched, we //can just make sure the allocation didn't throw an exception @@ -296,7 +310,9 @@ private: { //you should not be able to allocate a size larger than the //user provided and get the results - arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag()); + vtkm::cont::Token token; + arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag(), token); + token.DetachFromAll(); arrayHandle.GetPortalControl(); } catch (vtkm::cont::Error&) @@ -320,10 +336,11 @@ private: VTKM_TEST_ASSERT(arrayHandle.GetNumberOfValues() == 0, "ArrayHandle has wrong number of entries."); { + vtkm::cont::Token token; using ExecutionPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal; ExecutionPortalType executionPortal = - arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag()); + arrayHandle.PrepareForOutput(ARRAY_SIZE * 2, DeviceAdapterTag(), token); //we drop down to manually scheduling so that we don't need //need to bring in array handle counting @@ -349,9 +366,11 @@ private: std::cout << "Try in place operation." << std::endl; { + vtkm::cont::Token token; using ExecutionPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal; - ExecutionPortalType executionPortal = arrayHandle.PrepareForInPlace(DeviceAdapterTag()); + ExecutionPortalType executionPortal = + arrayHandle.PrepareForInPlace(DeviceAdapterTag(), token); //in place can't be done through the dispatcher //instead we have to drop down to manually scheduling @@ -396,7 +415,8 @@ private: VTKM_TEST_ASSERT(a1 == a2, "Shallow copied array not equal."); VTKM_TEST_ASSERT(!(a1 != a2), "Shallow copied array not equal."); - a1.PrepareForInPlace(DeviceAdapterTag()); + vtkm::cont::Token token; + a1.PrepareForInPlace(DeviceAdapterTag(), token); VTKM_TEST_ASSERT(a1 == a2, "Shallow copied array not equal."); VTKM_TEST_ASSERT(!(a1 != a2), "Shallow copied array not equal."); } diff --git a/vtkm/cont/testing/TestingBitField.h b/vtkm/cont/testing/TestingBitField.h index ad9588588..81cdae6a1 100644 --- a/vtkm/cont/testing/TestingBitField.h +++ b/vtkm/cont/testing/TestingBitField.h @@ -462,9 +462,10 @@ struct TestingBitField VTKM_CONT static void TestExecutionPortals() { + vtkm::cont::Token token; auto field = RandomBitField(); - auto portal = field.PrepareForInPlace(DeviceAdapterTag{}); - auto portalConst = field.PrepareForInput(DeviceAdapterTag{}); + auto portal = field.PrepareForInPlace(DeviceAdapterTag{}, token); + auto portalConst = field.PrepareForInput(DeviceAdapterTag{}, token); HelpTestPortalsExecution(portal, portalConst); } @@ -581,10 +582,13 @@ struct TestingBitField " got: ", numBits); + vtkm::cont::Token token; Algo::Schedule( - ArrayHandleBitFieldChecker{ handle.PrepareForInPlace(DeviceAdapterTag{}), false }, numBits); - Algo::Schedule(ArrayHandleBitFieldChecker{ handle.PrepareForInPlace(DeviceAdapterTag{}), true }, - numBits); + ArrayHandleBitFieldChecker{ handle.PrepareForInPlace(DeviceAdapterTag{}, token), false }, + numBits); + Algo::Schedule( + ArrayHandleBitFieldChecker{ handle.PrepareForInPlace(DeviceAdapterTag{}, token), true }, + numBits); } VTKM_CONT diff --git a/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h b/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h index e82501595..f758d943b 100644 --- a/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h +++ b/vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h @@ -35,11 +35,14 @@ public: using RectilinearPortalType = typename RectilinearType::template ExecutionTypes::PortalConst; - LocatorWorklet(vtkm::Bounds& bounds, vtkm::Id3& dims, const RectilinearType& coords) + LocatorWorklet(vtkm::Bounds& bounds, + vtkm::Id3& dims, + const RectilinearType& coords, + vtkm::cont::Token& token) : Bounds(bounds) , Dims(dims) { - RectilinearPortalType coordsPortal = coords.PrepareForInput(DeviceAdapter()); + RectilinearPortalType coordsPortal = coords.PrepareForInput(DeviceAdapter(), token); xAxis = coordsPortal.GetFirstPortal(); yAxis = coordsPortal.GetSecondPortal(); zAxis = coordsPortal.GetThirdPortal(); @@ -173,8 +176,9 @@ public: vtkm::cont::ArrayHandle cellIds; vtkm::cont::ArrayHandle parametric; vtkm::cont::ArrayHandle match; + vtkm::cont::Token token; LocatorWorklet worklet( - bounds, dims, coords.GetData().template Cast()); + bounds, dims, coords.GetData().template Cast(), token); vtkm::worklet::DispatcherMapField> dispatcher(worklet); dispatcher.SetDevice(DeviceAdapter()); diff --git a/vtkm/cont/testing/TestingColorTable.h b/vtkm/cont/testing/TestingColorTable.h index a464b5c5c..dc9d3c901 100644 --- a/vtkm/cont/testing/TestingColorTable.h +++ b/vtkm/cont/testing/TestingColorTable.h @@ -476,10 +476,13 @@ public: auto samples = vtkm::cont::make_ArrayHandle(data, nvals); vtkm::cont::ArrayHandle colors; - TransferFunction transfer(table.PrepareForExecution(DeviceAdapterTag{})); - vtkm::worklet::DispatcherMapField dispatcher(transfer); - dispatcher.SetDevice(DeviceAdapterTag()); - dispatcher.Invoke(samples, colors); + { + vtkm::cont::Token token; + TransferFunction transfer(table.PrepareForExecution(DeviceAdapterTag{}, token)); + vtkm::worklet::DispatcherMapField dispatcher(transfer); + dispatcher.SetDevice(DeviceAdapterTag()); + dispatcher.Invoke(samples, colors); + } const vtkm::Vec4ui_8 correct_sampling_points[nvals] = { { 14, 28, 31, 255 }, { 21, 150, 21, 255 }, diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h index 71c7d2801..13834f3c5 100644 --- a/vtkm/cont/testing/TestingDeviceAdapter.h +++ b/vtkm/cont/testing/TestingDeviceAdapter.h @@ -318,8 +318,8 @@ public: struct AtomicKernel { VTKM_CONT - AtomicKernel(const vtkm::cont::AtomicArray& array) - : AArray(array.PrepareForExecution(DeviceAdapterTag())) + AtomicKernel(const vtkm::cont::AtomicArray& array, vtkm::cont::Token& token) + : AArray(array.PrepareForExecution(DeviceAdapterTag(), token)) { } @@ -338,8 +338,8 @@ public: struct AtomicCASKernel { VTKM_CONT - AtomicCASKernel(const vtkm::cont::AtomicArray& array) - : AArray(array.PrepareForExecution(DeviceAdapterTag())) + AtomicCASKernel(const vtkm::cont::AtomicArray& array, vtkm::cont::Token& token) + : AArray(array.PrepareForExecution(DeviceAdapterTag(), token)) { } @@ -378,9 +378,11 @@ public: vtkm::Id Value = 0; }; - VirtualObjectTransferKernel(const Interface* vo, IdArrayHandle& result) + VirtualObjectTransferKernel(const Interface* vo, + IdArrayHandle& result, + vtkm::cont::Token& token) : Virtual(vo) - , Result(result.PrepareForInPlace(DeviceAdapterTag())) + , Result(result.PrepareForInPlace(DeviceAdapterTag(), token)) { } @@ -527,7 +529,11 @@ private: // Do an operation just so we know the values are placed in the execution // environment and they change. We are only calling on half the array // because we are about to shrink. - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{})), ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + ARRAY_SIZE); + } // Change size. handle.Shrink(ARRAY_SIZE); @@ -557,9 +563,10 @@ private: try { std::cout << "Do array allocation that should fail." << std::endl; + vtkm::cont::Token token; vtkm::cont::ArrayHandle bigArray; const vtkm::Id bigSize = 0x7FFFFFFFFFFFFFFFLL; - bigArray.PrepareForOutput(bigSize, DeviceAdapterTag{}); + bigArray.PrepareForOutput(bigSize, DeviceAdapterTag{}, token); // It does not seem reasonable to get here. The previous call should fail. VTKM_TEST_FAIL("A ridiculously sized allocation succeeded. Either there " "was a failure that was not reported but should have been " @@ -620,14 +627,22 @@ private: target.Value = 5; Transfer transfer(&target); - const BaseType* base = static_cast(transfer.PrepareForExecution(false)); + vtkm::cont::Token transferToken; + const BaseType* base = + static_cast(transfer.PrepareForExecution(false, transferToken)); - Algorithm::Schedule(VirtualObjectTransferKernel(base, result), 1); + { + vtkm::cont::Token token; + Algorithm::Schedule(VirtualObjectTransferKernel(base, result, token), 1); + } VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 5, "Did not get expected result"); - target.Value = 10; - base = static_cast(transfer.PrepareForExecution(true)); - Algorithm::Schedule(VirtualObjectTransferKernel(base, result), 1); + { + vtkm::cont::Token token; + target.Value = 10; + base = static_cast(transfer.PrepareForExecution(true, token)); + Algorithm::Schedule(VirtualObjectTransferKernel(base, result, token), 1); + } VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 10, "Did not get expected result"); transfer.ReleaseResources(); @@ -643,10 +658,17 @@ private: vtkm::cont::ArrayHandle handle; std::cout << "Running clear." << std::endl; - Algorithm::Schedule(ClearArrayKernel(handle.PrepareForOutput(1, DeviceAdapterTag{})), 1); + { + vtkm::cont::Token token; + Algorithm::Schedule(ClearArrayKernel(handle.PrepareForOutput(1, DeviceAdapterTag{}, token)), + 1); + } std::cout << "Running add." << std::endl; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{})), 1); + { + vtkm::cont::Token token; + Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), 1); + } std::cout << "Checking results." << std::endl; for (vtkm::Id index = 0; index < 1; index++) @@ -665,11 +687,19 @@ private: vtkm::cont::ArrayHandle handle; std::cout << "Running clear." << std::endl; - Algorithm::Schedule(ClearArrayKernel(handle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag{})), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(handle.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag{}, token)), + ARRAY_SIZE); + } std::cout << "Running add." << std::endl; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{})), ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + ARRAY_SIZE); + } std::cout << "Checking results." << std::endl; for (vtkm::Id index = 0; index < ARRAY_SIZE; index++) @@ -691,11 +721,18 @@ private: //size is selected to be larger than the CUDA backend can launch in a //single invocation when compiled for SM_2 support const vtkm::Id size = 8400000; - Algorithm::Schedule(ClearArrayKernel(handle.PrepareForOutput(size, DeviceAdapterTag{})), - size); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(handle.PrepareForOutput(size, DeviceAdapterTag{}, token)), size); + } std::cout << "Running add." << std::endl; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{})), size); + { + vtkm::cont::Token token; + Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + size); + } std::cout << "Checking results." << std::endl; //Rather than testing for correctness every value of a large array, @@ -721,14 +758,21 @@ private: vtkm::Id3 maxRange(DIM_SIZE); std::cout << "Running clear." << std::endl; - Algorithm::Schedule( - ClearArrayKernel( - handle.PrepareForOutput(DIM_SIZE * DIM_SIZE * DIM_SIZE, DeviceAdapterTag{}), maxRange), - maxRange); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel( + handle.PrepareForOutput(DIM_SIZE * DIM_SIZE * DIM_SIZE, DeviceAdapterTag{}, token), + maxRange), + maxRange); + } std::cout << "Running add." << std::endl; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}), maxRange), - maxRange); + { + vtkm::cont::Token token; + Algorithm::Schedule( + AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token), maxRange), maxRange); + } std::cout << "Checking results." << std::endl; const vtkm::Id maxId = DIM_SIZE * DIM_SIZE * DIM_SIZE; @@ -751,17 +795,24 @@ private: // Initialize tracker with 'false' values std::cout << "Allocating and initializing memory" << std::endl; - Algorithm::Schedule(GenericClearArrayKernel( - tracker.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()), false), - ARRAY_SIZE); - Algorithm::Schedule(GenericClearArrayKernel( - valid.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()), false), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + GenericClearArrayKernel( + tracker.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token), false), + ARRAY_SIZE); + Algorithm::Schedule(GenericClearArrayKernel( + valid.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token), false), + ARRAY_SIZE); + } std::cout << "Running Overlap kernel." << std::endl; - Algorithm::Schedule(OverlapKernel(tracker.PrepareForInPlace(DeviceAdapterTag()), - valid.PrepareForInPlace(DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(OverlapKernel(tracker.PrepareForInPlace(DeviceAdapterTag(), token), + valid.PrepareForInPlace(DeviceAdapterTag(), token)), + ARRAY_SIZE); + } std::cout << "Checking results." << std::endl; @@ -788,18 +839,26 @@ private: // Initialize tracker with 'false' values std::cout << "Allocating and initializing memory" << std::endl; - Algorithm::Schedule(GenericClearArrayKernel( - tracker.PrepareForOutput(numElems, DeviceAdapterTag()), dims, false), - numElems); - Algorithm::Schedule(GenericClearArrayKernel( - valid.PrepareForOutput(numElems, DeviceAdapterTag()), dims, false), - numElems); + { + vtkm::cont::Token token; + Algorithm::Schedule( + GenericClearArrayKernel( + tracker.PrepareForOutput(numElems, DeviceAdapterTag(), token), dims, false), + numElems); + Algorithm::Schedule( + GenericClearArrayKernel( + valid.PrepareForOutput(numElems, DeviceAdapterTag(), token), dims, false), + numElems); + } std::cout << "Running Overlap kernel." << std::endl; - Algorithm::Schedule(OverlapKernel(tracker.PrepareForInPlace(DeviceAdapterTag()), - valid.PrepareForInPlace(DeviceAdapterTag()), - dims), - dims); + { + vtkm::cont::Token token; + Algorithm::Schedule(OverlapKernel(tracker.PrepareForInPlace(DeviceAdapterTag(), token), + valid.PrepareForInPlace(DeviceAdapterTag(), token), + dims), + dims); + } std::cout << "Checking results." << std::endl; @@ -823,10 +882,15 @@ private: std::cout << " Standard call" << std::endl; //construct the index array - Algorithm::Schedule( - OffsetPlusIndexKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), ARRAY_SIZE); - Algorithm::Schedule( - MarkOddNumbersKernel(stencil.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + OffsetPlusIndexKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + Algorithm::Schedule( + MarkOddNumbersKernel(stencil.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + } Algorithm::CopyIf(array, stencil, result); VTKM_TEST_ASSERT(result.GetNumberOfValues() == array.GetNumberOfValues() / 2, @@ -1230,8 +1294,12 @@ private: //construct the index array IdArrayHandle array; - Algorithm::Schedule(ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + } //the output of reduce and scan inclusive should be the same std::cout << " Reduce with initial value of 0." << std::endl; @@ -1329,12 +1397,14 @@ private: std::cout << "-------------------------------------------" << std::endl; std::cout << "Testing Reduce with ArrayHandleZip" << std::endl; { + vtkm::cont::Token token; IdArrayHandle keys, values; - Algorithm::Schedule(ClearArrayKernel(keys.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + Algorithm::Schedule( + ClearArrayKernel(keys.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), ARRAY_SIZE); - Algorithm::Schedule(ClearArrayKernel(values.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + Algorithm::Schedule( + ClearArrayKernel(values.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); vtkm::cont::ArrayHandleZip zipped(keys, values); @@ -1767,8 +1837,12 @@ private: std::cout << " size " << ARRAY_SIZE << std::endl; //construct the index array IdArrayHandle array; - Algorithm::Schedule(ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + } //we know have an array whose sum is equal to OFFSET * ARRAY_SIZE, //let's validate that @@ -1856,11 +1930,15 @@ private: //construct the index array IdArrayHandle array; - Algorithm::Schedule(ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + Algorithm::Schedule( + AddArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), ARRAY_SIZE); + } - Algorithm::Schedule(AddArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); //we know have an array whose sum is equal to OFFSET * ARRAY_SIZE, //let's validate that IdArrayHandle result; @@ -1897,8 +1975,12 @@ private: std::cout << " size " << ARRAY_SIZE << std::endl; //construct the index array IdArrayHandle array; - Algorithm::Schedule(ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + } // we know have an array whose sum = (OFFSET * ARRAY_SIZE), // let's validate that @@ -1996,8 +2078,12 @@ private: //construct the index array IdArrayHandle array; - Algorithm::Schedule(ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag())), - ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule( + ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); + } // we now have an array whose sum = (OFFSET * ARRAY_SIZE), // let's validate that @@ -2128,9 +2214,10 @@ private: int nkernels = 0; try { + vtkm::cont::Token token; + IdArrayHandle idArray; - idArray.Allocate(ARRAY_SIZE); - auto portal = idArray.PrepareForInPlace(DeviceAdapterTag{}); + auto portal = idArray.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag{}, token); Algorithm::Schedule(OneErrorKernel(), ARRAY_SIZE); for (; nkernels < 100; ++nkernels) @@ -2447,7 +2534,10 @@ private: vtkm::cont::make_ArrayHandle(singleElement); vtkm::cont::AtomicArray atomic(atomicElement); - Algorithm::Schedule(AtomicKernel(atomic), SHORT_ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AtomicKernel(atomic, token), SHORT_ARRAY_SIZE); + } vtkm::Int32 expected = vtkm::Int32(atomicCount); vtkm::Int32 actual = atomicElement.GetPortalControl().Get(0); VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int32"); @@ -2461,7 +2551,10 @@ private: vtkm::cont::make_ArrayHandle(singleElement); vtkm::cont::AtomicArray atomic(atomicElement); - Algorithm::Schedule(AtomicKernel(atomic), SHORT_ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AtomicKernel(atomic, token), SHORT_ARRAY_SIZE); + } vtkm::Int64 expected = vtkm::Int64(atomicCount); vtkm::Int64 actual = atomicElement.GetPortalControl().Get(0); VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int64"); @@ -2475,7 +2568,10 @@ private: vtkm::cont::make_ArrayHandle(singleElement); vtkm::cont::AtomicArray atomic(atomicElement); - Algorithm::Schedule(AtomicCASKernel(atomic), SHORT_ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AtomicCASKernel(atomic, token), SHORT_ARRAY_SIZE); + } vtkm::Int32 expected = vtkm::Int32(atomicCount); vtkm::Int32 actual = atomicElement.GetPortalControl().Get(0); VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic CAS Int32"); @@ -2489,7 +2585,10 @@ private: vtkm::cont::make_ArrayHandle(singleElement); vtkm::cont::AtomicArray atomic(atomicElement); - Algorithm::Schedule(AtomicCASKernel(atomic), SHORT_ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(AtomicCASKernel(atomic, token), SHORT_ARRAY_SIZE); + } vtkm::Int64 expected = vtkm::Int64(atomicCount); vtkm::Int64 actual = atomicElement.GetPortalControl().Get(0); VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic CAS Int64"); diff --git a/vtkm/cont/testing/TestingFancyArrayHandles.h b/vtkm/cont/testing/TestingFancyArrayHandles.h index 0d061e58b..e1a94afe2 100644 --- a/vtkm/cont/testing/TestingFancyArrayHandles.h +++ b/vtkm/cont/testing/TestingFancyArrayHandles.h @@ -188,9 +188,10 @@ struct TransformExecObject : public vtkm::cont::ExecutionAndControlObjectBase }; template - VTKM_CONT FunctorWrapper PrepareForExecution(DeviceAdapterTag device) const + VTKM_CONT FunctorWrapper PrepareForExecution(DeviceAdapterTag device, + vtkm::cont::Token& token) const { - return FunctorWrapper(this->VirtualFunctor.PrepareForExecution(device)); + return FunctorWrapper(this->VirtualFunctor.PrepareForExecution(device, token)); } VTKM_CONT FunctorWrapper PrepareForControl() const diff --git a/vtkm/cont/testing/TestingImplicitFunction.h b/vtkm/cont/testing/TestingImplicitFunction.h index 675606689..c9d92d59e 100644 --- a/vtkm/cont/testing/TestingImplicitFunction.h +++ b/vtkm/cont/testing/TestingImplicitFunction.h @@ -65,7 +65,8 @@ void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points, { using EvalDispatcher = vtkm::worklet::DispatcherMapField; - EvaluateImplicitFunction eval(function.PrepareForExecution(device)); + vtkm::cont::Token token; + EvaluateImplicitFunction eval(function.PrepareForExecution(device, token)); EvalDispatcher dispatcher(eval); dispatcher.SetDevice(DeviceAdapter()); dispatcher.Invoke(points, values, gradients); diff --git a/vtkm/cont/testing/TestingVirtualObjectHandle.h b/vtkm/cont/testing/TestingVirtualObjectHandle.h index 7fd8ebc20..7930ae745 100644 --- a/vtkm/cont/testing/TestingVirtualObjectHandle.h +++ b/vtkm/cont/testing/TestingVirtualObjectHandle.h @@ -111,7 +111,9 @@ private: for (int n = 0; n < 2; ++n) { - virtual_object_detail::TransformerFunctor tfnctr(this->Handle->PrepareForExecution(device)); + vtkm::cont::Token token; + virtual_object_detail::TransformerFunctor tfnctr( + this->Handle->PrepareForExecution(device, token)); ArrayTransform transformed(*this->Input, tfnctr); FloatArrayHandle output; @@ -119,7 +121,7 @@ private: auto portal = output.GetPortalConstControl(); for (vtkm::Id i = 0; i < ARRAY_LEN; ++i) { - VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i * i), "\tIncorrect result"); + VTKM_TEST_ASSERT(test_equal(portal.Get(i), i * i), "\tIncorrect result"); } std::cout << "\tSuccess." << std::endl; @@ -158,7 +160,9 @@ private: this->Mul->SetMultiplicand(2); for (int n = 0; n < 2; ++n) { - virtual_object_detail::TransformerFunctor tfnctr(this->Handle->PrepareForExecution(device)); + vtkm::cont::Token token; + virtual_object_detail::TransformerFunctor tfnctr( + this->Handle->PrepareForExecution(device, token)); ArrayTransform transformed(*this->Input, tfnctr); FloatArrayHandle output; @@ -166,7 +170,7 @@ private: auto portal = output.GetPortalConstControl(); for (vtkm::Id i = 0; i < ARRAY_LEN; ++i) { - VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i) * this->Mul->GetMultiplicand(), + VTKM_TEST_ASSERT(test_equal(portal.Get(i), i * this->Mul->GetMultiplicand()), "\tIncorrect result"); } std::cout << "\tSuccess." << std::endl; diff --git a/vtkm/cont/testing/UnitTestAlgorithm.cxx b/vtkm/cont/testing/UnitTestAlgorithm.cxx index 2b01cdcbd..0284fc9a8 100644 --- a/vtkm/cont/testing/UnitTestAlgorithm.cxx +++ b/vtkm/cont/testing/UnitTestAlgorithm.cxx @@ -147,7 +147,7 @@ struct CompFunctor struct CompExecObject : vtkm::cont::ExecutionObjectBase { template - VTKM_CONT CompFunctor PrepareForExecution(Device) + VTKM_CONT CompFunctor PrepareForExecution(Device, vtkm::cont::Token&) { return CompFunctor(); } diff --git a/vtkm/cont/testing/UnitTestArrayHandleDiscard.cxx b/vtkm/cont/testing/UnitTestArrayHandleDiscard.cxx index 7ffd018ea..bdb6603c6 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleDiscard.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleDiscard.cxx @@ -75,18 +75,20 @@ struct Test for (vtkm::Id i = 0; i < NUM_KEYS; ++i) { - VTKM_TEST_ASSERT(outputs.Get(i) == refData[i], "Unexpected output value after ReduceByKey."); + VTKM_TEST_ASSERT(test_equal(outputs.Get(i), refData[i]), + "Unexpected output value after ReduceByKey."); } } void TestPrepareExceptions() { + vtkm::cont::Token token; DiscardHandle handle; handle.Allocate(50); try { - handle.PrepareForInput(DeviceTag()); + handle.PrepareForInput(DeviceTag(), token); } catch (vtkm::cont::ErrorBadValue&) { @@ -95,7 +97,7 @@ struct Test try { - handle.PrepareForInPlace(DeviceTag()); + handle.PrepareForInPlace(DeviceTag(), token); } catch (vtkm::cont::ErrorBadValue&) { @@ -103,7 +105,7 @@ struct Test } // Shouldn't fail: - handle.PrepareForOutput(ARRAY_SIZE, DeviceTag()); + handle.PrepareForOutput(ARRAY_SIZE, DeviceTag(), token); } void operator()() diff --git a/vtkm/cont/testing/UnitTestArrayHandleExtractComponent.cxx b/vtkm/cont/testing/UnitTestArrayHandleExtractComponent.cxx index a7b5c667d..5febe8c77 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleExtractComponent.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleExtractComponent.cxx @@ -158,11 +158,15 @@ struct ExtractComponentTests using RefPortal = typename ReferenceCompositeArray::template ExecutionTypes::PortalConst; - WriteTestFunctor functor(extract.PrepareForInPlace(DeviceTag()), - this->RefComposite.PrepareForInput(DeviceTag()), - component); + { + vtkm::cont::Token token; + WriteTestFunctor functor( + extract.PrepareForInPlace(DeviceTag(), token), + this->RefComposite.PrepareForInput(DeviceTag(), token), + component); + Algo::Schedule(functor, extract.GetNumberOfValues()); + } - Algo::Schedule(functor, extract.GetNumberOfValues()); this->ValidateWriteTestArray(composite, component); } } diff --git a/vtkm/cont/testing/UnitTestArrayHandleImplicit.cxx b/vtkm/cont/testing/UnitTestArrayHandleImplicit.cxx index 8d13ce362..0737426d7 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleImplicit.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleImplicit.cxx @@ -53,9 +53,10 @@ struct ImplicitTests } //verify that the execution portal works + vtkm::cont::Token token; using Device = vtkm::cont::DeviceAdapterTagSerial; using CEPortal = typename ImplicitHandle::template ExecutionTypes::PortalConst; - CEPortal execPortal = implict.PrepareForInput(Device()); + CEPortal execPortal = implict.PrepareForInput(Device(), token); for (int i = 0; i < ARRAY_SIZE; ++i) { const ValueType v = execPortal.Get(i); diff --git a/vtkm/cont/testing/UnitTestArrayHandlePermutation.cxx b/vtkm/cont/testing/UnitTestArrayHandlePermutation.cxx index 237133830..c22fdd9c2 100644 --- a/vtkm/cont/testing/UnitTestArrayHandlePermutation.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandlePermutation.cxx @@ -58,12 +58,14 @@ struct CheckPermutationFunctor : vtkm::exec::FunctorBase template VTKM_CONT CheckPermutationFunctor< typename PermutedArrayHandleType::template ExecutionTypes::PortalConst> -make_CheckPermutationFunctor(const PermutedArrayHandleType& permutedArray, Device) +make_CheckPermutationFunctor(const PermutedArrayHandleType& permutedArray, + Device, + vtkm::cont::Token& token) { using PermutedPortalType = typename PermutedArrayHandleType::template ExecutionTypes::PortalConst; CheckPermutationFunctor functor; - functor.PermutedPortal = permutedArray.PrepareForInput(Device()); + functor.PermutedPortal = permutedArray.PrepareForInput(Device(), token); return functor; } @@ -87,12 +89,14 @@ struct InPlacePermutationFunctor : vtkm::exec::FunctorBase template VTKM_CONT InPlacePermutationFunctor< typename PermutedArrayHandleType::template ExecutionTypes::Portal> -make_InPlacePermutationFunctor(PermutedArrayHandleType& permutedArray, Device) +make_InPlacePermutationFunctor(PermutedArrayHandleType& permutedArray, + Device, + vtkm::cont::Token& token) { using PermutedPortalType = typename PermutedArrayHandleType::template ExecutionTypes::Portal; InPlacePermutationFunctor functor; - functor.PermutedPortal = permutedArray.PrepareForInPlace(Device()); + functor.PermutedPortal = permutedArray.PrepareForInPlace(Device(), token); return functor; } @@ -136,12 +140,14 @@ struct OutputPermutationFunctor : vtkm::exec::FunctorBase template VTKM_CONT OutputPermutationFunctor< typename PermutedArrayHandleType::template ExecutionTypes::Portal> -make_OutputPermutationFunctor(PermutedArrayHandleType& permutedArray, Device) +make_OutputPermutationFunctor(PermutedArrayHandleType& permutedArray, + Device, + vtkm::cont::Token& token) { using PermutedPortalType = typename PermutedArrayHandleType::template ExecutionTypes::Portal; OutputPermutationFunctor functor; - functor.PermutedPortal = permutedArray.PrepareForOutput(ARRAY_SIZE, Device()); + functor.PermutedPortal = permutedArray.PrepareForOutput(ARRAY_SIZE, Device(), token); return functor; } @@ -217,16 +223,21 @@ struct PermutationTests VTKM_TEST_ASSERT(permutationArray.GetPortalConstControl().GetNumberOfValues() == ARRAY_SIZE, "Permutation portal wrong size."); + vtkm::cont::Token token; + std::cout << "Test initial values in execution environment" << std::endl; - Algorithm::Schedule(make_CheckPermutationFunctor(permutationArray, Device()), ARRAY_SIZE); + Algorithm::Schedule(make_CheckPermutationFunctor(permutationArray, Device(), token), + ARRAY_SIZE); std::cout << "Try in place operation" << std::endl; - Algorithm::Schedule(make_InPlacePermutationFunctor(permutationArray, Device()), ARRAY_SIZE); + Algorithm::Schedule(make_InPlacePermutationFunctor(permutationArray, Device(), token), + ARRAY_SIZE); CheckInPlaceResult(valueArray.GetPortalControl()); CheckInPlaceResult(valueArray.GetPortalConstControl()); std::cout << "Try output operation" << std::endl; - Algorithm::Schedule(make_OutputPermutationFunctor(permutationArray, Device()), ARRAY_SIZE); + Algorithm::Schedule(make_OutputPermutationFunctor(permutationArray, Device(), token), + ARRAY_SIZE); CheckOutputResult(valueArray.GetPortalConstControl()); CheckOutputResult(valueArray.GetPortalControl()); } diff --git a/vtkm/cont/testing/UnitTestArrayHandleSwizzle.cxx b/vtkm/cont/testing/UnitTestArrayHandleSwizzle.cxx index 25e6b6ab2..f4c37a0b8 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleSwizzle.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleSwizzle.cxx @@ -161,8 +161,9 @@ struct SwizzleTests template bool operator()(DeviceTag, SwizzleHandleType& swizzle) const { + vtkm::cont::Token token; using Portal = typename SwizzleHandleType::template ExecutionTypes::Portal; - WriteTestFunctor functor(swizzle.PrepareForInPlace(DeviceTag())); + WriteTestFunctor functor(swizzle.PrepareForInPlace(DeviceTag(), token)); Algo::Schedule(functor, swizzle.GetNumberOfValues()); return true; } diff --git a/vtkm/cont/testing/UnitTestArrayHandleTransform.cxx b/vtkm/cont/testing/UnitTestArrayHandleTransform.cxx index cb9234c98..c0549d7bf 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleTransform.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleTransform.cxx @@ -59,15 +59,16 @@ VTKM_CONT CheckTransformFunctor< typename TransformedArrayHandleType::template ExecutionTypes::PortalConst> make_CheckTransformFunctor(const OriginalArrayHandleType& originalArray, const TransformedArrayHandleType& transformedArray, - Device) + Device, + vtkm::cont::Token& token) { using OriginalPortalType = typename OriginalArrayHandleType::template ExecutionTypes::PortalConst; using TransformedPortalType = typename TransformedArrayHandleType::template ExecutionTypes::PortalConst; CheckTransformFunctor functor; - functor.OriginalPortal = originalArray.PrepareForInput(Device()); - functor.TransformedPortal = transformedArray.PrepareForInput(Device()); + functor.OriginalPortal = originalArray.PrepareForInput(Device(), token); + functor.TransformedPortal = transformedArray.PrepareForInput(Device(), token); return functor; } @@ -117,16 +118,20 @@ struct TransformTests MySquare functor; std::cout << "Test a transform handle with a counting handle as the values" << std::endl; - vtkm::cont::ArrayHandleCounting counting = vtkm::cont::make_ArrayHandleCounting( - InputValueType(OutputValueType(0)), InputValueType(1), ARRAY_SIZE); - CountingTransformHandle countingTransformed = - vtkm::cont::make_ArrayHandleTransform(counting, functor); + { + vtkm::cont::Token token; + vtkm::cont::ArrayHandleCounting counting = + vtkm::cont::make_ArrayHandleCounting( + InputValueType(OutputValueType(0)), InputValueType(1), ARRAY_SIZE); + CountingTransformHandle countingTransformed = + vtkm::cont::make_ArrayHandleTransform(counting, functor); - CheckControlPortals(counting, countingTransformed); + CheckControlPortals(counting, countingTransformed); - std::cout << " Verify that the execution portal works" << std::endl; - Algorithm::Schedule(make_CheckTransformFunctor(counting, countingTransformed, Device()), - ARRAY_SIZE); + std::cout << " Verify that the execution portal works" << std::endl; + Algorithm::Schedule( + make_CheckTransformFunctor(counting, countingTransformed, Device(), token), ARRAY_SIZE); + } std::cout << "Test a transform handle with a normal handle as the values" << std::endl; //we are going to connect the two handles up, and than fill @@ -136,27 +141,38 @@ struct TransformTests using Portal = typename vtkm::cont::ArrayHandle::PortalControl; input.Allocate(ARRAY_SIZE); - Portal portal = input.GetPortalControl(); - for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index) { - portal.Set(index, TestValue(index, InputValueType())); + Portal portal = input.GetPortalControl(); + for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index) + { + portal.Set(index, TestValue(index, InputValueType())); + } } CheckControlPortals(input, thandle); std::cout << " Verify that the execution portal works" << std::endl; - Algorithm::Schedule(make_CheckTransformFunctor(input, thandle, Device()), ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(make_CheckTransformFunctor(input, thandle, Device(), token), ARRAY_SIZE); + } std::cout << "Modify array handle values to ensure transform gets updated" << std::endl; - for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index) { - portal.Set(index, TestValue(index * index, InputValueType())); + Portal portal = input.GetPortalControl(); + for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index) + { + portal.Set(index, TestValue(index * index, InputValueType())); + } } CheckControlPortals(input, thandle); std::cout << " Verify that the execution portal works" << std::endl; - Algorithm::Schedule(make_CheckTransformFunctor(input, thandle, Device()), ARRAY_SIZE); + { + vtkm::cont::Token token; + Algorithm::Schedule(make_CheckTransformFunctor(input, thandle, Device(), token), ARRAY_SIZE); + } } }; diff --git a/vtkm/cont/testing/UnitTestArrayHandleVirtual.cxx b/vtkm/cont/testing/UnitTestArrayHandleVirtual.cxx index a0ca39b5e..0b3b04a89 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleVirtual.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleVirtual.cxx @@ -133,9 +133,10 @@ struct Test try { - virt.PrepareForInput(DeviceTag()); - virt.PrepareForInPlace(DeviceTag()); - virt.PrepareForOutput(ARRAY_SIZE, DeviceTag()); + vtkm::cont::Token token; + virt.PrepareForInput(DeviceTag(), token); + virt.PrepareForInPlace(DeviceTag(), token); + virt.PrepareForOutput(ARRAY_SIZE, DeviceTag(), token); } catch (vtkm::cont::ErrorBadValue&) { diff --git a/vtkm/cont/testing/UnitTestCellSetExplicit.cxx b/vtkm/cont/testing/UnitTestCellSetExplicit.cxx index 6f9c5fb86..c9cba5180 100644 --- a/vtkm/cont/testing/UnitTestCellSetExplicit.cxx +++ b/vtkm/cont/testing/UnitTestCellSetExplicit.cxx @@ -167,7 +167,8 @@ void TestCellSetExplicit() "CellToPoint table exists before PrepareForInput."); // Test a raw PrepareForInput call: - cellset.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, PointTag{}, CellTag{}); + vtkm::cont::Token token; + cellset.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, PointTag{}, CellTag{}, token); VTKM_TEST_ASSERT(VTKM_PASS_COMMAS(cellset.HasConnectivity(PointTag{}, CellTag{})), "CellToPoint table missing after PrepareForInput."); diff --git a/vtkm/cont/testing/UnitTestDataSetPermutation.cxx b/vtkm/cont/testing/UnitTestDataSetPermutation.cxx index 54c552817..ffea22391 100644 --- a/vtkm/cont/testing/UnitTestDataSetPermutation.cxx +++ b/vtkm/cont/testing/UnitTestDataSetPermutation.cxx @@ -108,10 +108,12 @@ void TestDataSet_Explicit() vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint>::ExecObjectType; + vtkm::cont::Token token; ExecObjectType execConnectivity; execConnectivity = subset.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); //run a basic for-each topology algorithm on this vtkm::cont::ArrayHandle result; @@ -153,8 +155,9 @@ void TestDataSet_Structured2D() using DeviceAdapterTag = vtkm::cont::DeviceAdapterTagSerial; //verify that PrepareForInput exists + vtkm::cont::Token token; subset.PrepareForInput( - DeviceAdapterTag(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint()); + DeviceAdapterTag(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint(), token); //run a basic for-each topology algorithm on this vtkm::cont::ArrayHandle result; @@ -192,9 +195,11 @@ void TestDataSet_Structured3D() subset.PrintSummary(std::cout); //verify that PrepareForInput exists + vtkm::cont::Token token; subset.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); //run a basic for-each topology algorithm on this vtkm::cont::ArrayHandle result; diff --git a/vtkm/cont/testing/UnitTestDataSetRectilinear.cxx b/vtkm/cont/testing/UnitTestDataSetRectilinear.cxx index f63bbfe7e..06ccf6dfc 100644 --- a/vtkm/cont/testing/UnitTestDataSetRectilinear.cxx +++ b/vtkm/cont/testing/UnitTestDataSetRectilinear.cxx @@ -76,14 +76,18 @@ static void TwoDimRectilinearTest() VTKM_TEST_ASSERT(shape == vtkm::CELL_SHAPE_QUAD, "Incorrect element type."); } + vtkm::cont::Token token; + vtkm::exec::ConnectivityStructured pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); vtkm::exec::ConnectivityStructured cellToPoint = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagPoint(), - vtkm::TopologyElementTagCell()); + vtkm::TopologyElementTagCell(), + token); vtkm::Id cells[2][4] = { { 0, 1, 4, 3 }, { 1, 2, 5, 4 } }; for (vtkm::Id cellIndex = 0; cellIndex < 2; cellIndex++) @@ -161,10 +165,12 @@ static void ThreeDimRectilinearTest() } //Test regular connectivity. + vtkm::cont::Token token; vtkm::exec::ConnectivityStructured pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); vtkm::Id expectedPointIds[8] = { 0, 1, 4, 3, 6, 7, 10, 9 }; vtkm::Vec retrievedPointIds = pointToCell.GetIndices(vtkm::Id3(0)); for (vtkm::IdComponent localPointIndex = 0; localPointIndex < 8; localPointIndex++) @@ -176,7 +182,8 @@ static void ThreeDimRectilinearTest() vtkm::exec::ConnectivityStructured cellToPoint = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagPoint(), - vtkm::TopologyElementTagCell()); + vtkm::TopologyElementTagCell(), + token); vtkm::Id retrievedCellIds[6] = { 0, -1, -1, -1, -1, -1 }; vtkm::VecVariable expectedCellIds = cellToPoint.GetIndices(vtkm::Id3(0)); VTKM_TEST_ASSERT(expectedCellIds.GetNumberOfComponents() <= 6, diff --git a/vtkm/cont/testing/UnitTestDataSetUniform.cxx b/vtkm/cont/testing/UnitTestDataSetUniform.cxx index 062242d78..894419734 100644 --- a/vtkm/cont/testing/UnitTestDataSetUniform.cxx +++ b/vtkm/cont/testing/UnitTestDataSetUniform.cxx @@ -79,14 +79,17 @@ static void TwoDimUniformTest() VTKM_TEST_ASSERT(shape == vtkm::CELL_SHAPE_QUAD, "Incorrect element type."); } + vtkm::cont::Token token; vtkm::exec::ConnectivityStructured pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); vtkm::exec::ConnectivityStructured cellToPoint = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagPoint(), - vtkm::TopologyElementTagCell()); + vtkm::TopologyElementTagCell(), + token); vtkm::Id cells[2][4] = { { 0, 1, 4, 3 }, { 1, 2, 5, 4 } }; for (vtkm::Id cellIndex = 0; cellIndex < 2; cellIndex++) @@ -170,11 +173,14 @@ static void ThreeDimUniformTest() VTKM_TEST_ASSERT(shape == vtkm::CELL_SHAPE_HEXAHEDRON, "Incorrect element type."); } + vtkm::cont::Token token; + //Test uniform connectivity. vtkm::exec::ConnectivityStructured pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint()); + vtkm::TopologyElementTagPoint(), + token); vtkm::Id expectedPointIds[8] = { 0, 1, 4, 3, 6, 7, 10, 9 }; vtkm::Vec retrievedPointIds = pointToCell.GetIndices(vtkm::Id3(0)); for (vtkm::IdComponent localPointIndex = 0; localPointIndex < 8; localPointIndex++) @@ -186,7 +192,8 @@ static void ThreeDimUniformTest() vtkm::exec::ConnectivityStructured cellToPoint = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(), vtkm::TopologyElementTagPoint(), - vtkm::TopologyElementTagCell()); + vtkm::TopologyElementTagCell(), + token); vtkm::Id retrievedCellIds[6] = { 0, -1, -1, -1, -1, -1 }; vtkm::VecVariable expectedCellIds = cellToPoint.GetIndices(vtkm::Id3(0)); VTKM_TEST_ASSERT(expectedCellIds.GetNumberOfComponents() <= 6, diff --git a/vtkm/cont/testing/UnitTestMoveConstructors.cxx b/vtkm/cont/testing/UnitTestMoveConstructors.cxx index cfdfa697c..2dd09d2ba 100644 --- a/vtkm/cont/testing/UnitTestMoveConstructors.cxx +++ b/vtkm/cont/testing/UnitTestMoveConstructors.cxx @@ -81,16 +81,16 @@ struct IsNoExceptHandle is_noexcept_movable(); //verify the input portals of the handle - is_noexcept_movable().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}))>(); - is_noexcept_movable().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}))>(); + is_noexcept_movable().PrepareForInput( + vtkm::cont::DeviceAdapterTagSerial{}, std::declval()))>(); + is_noexcept_movable().PrepareForInput( + vtkm::cont::DeviceAdapterTagSerial{}, std::declval()))>(); //verify the output portals of the handle - is_noexcept_movable().PrepareForOutput(2, vtkm::cont::DeviceAdapterTagSerial{}))>(); - is_noexcept_movable().PrepareForOutput(2, vtkm::cont::DeviceAdapterTagSerial{}))>(); + is_noexcept_movable().PrepareForOutput( + 2, vtkm::cont::DeviceAdapterTagSerial{}, std::declval()))>(); + is_noexcept_movable().PrepareForOutput( + 2, vtkm::cont::DeviceAdapterTagSerial{}, std::declval()))>(); } }; diff --git a/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h b/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h index d8206c9c4..46e60659a 100644 --- a/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h +++ b/vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h @@ -78,11 +78,12 @@ public: const CellIdArrayHandle& cellIds, const CellSetType& cellSet, const vtkm::cont::ArrayHandleVirtualCoordinates& coords, - DeviceAdapter) - : Nodes(nodes.PrepareForInput(DeviceAdapter())) - , CellIds(cellIds.PrepareForInput(DeviceAdapter())) - , CellSet(cellSet.PrepareForInput(DeviceAdapter(), VisitType(), IncidentType())) - , Coords(coords.PrepareForInput(DeviceAdapter())) + DeviceAdapter, + vtkm::cont::Token& token) + : Nodes(nodes.PrepareForInput(DeviceAdapter(), token)) + , CellIds(cellIds.PrepareForInput(DeviceAdapter(), token)) + , CellSet(cellSet.PrepareForInput(DeviceAdapter(), VisitType(), IncidentType(), token)) + , Coords(coords.PrepareForInput(DeviceAdapter(), token)) { } diff --git a/vtkm/exec/CellLocatorRectilinearGrid.h b/vtkm/exec/CellLocatorRectilinearGrid.h index 6f5c85dad..c47de9c7a 100644 --- a/vtkm/exec/CellLocatorRectilinearGrid.h +++ b/vtkm/exec/CellLocatorRectilinearGrid.h @@ -49,11 +49,12 @@ public: const vtkm::Id rowSize, const vtkm::cont::CellSetStructured& cellSet, const RectilinearType& coords, - DeviceAdapter) + DeviceAdapter, + vtkm::cont::Token& token) : PlaneSize(planeSize) , RowSize(rowSize) - , CellSet(cellSet.PrepareForInput(DeviceAdapter(), VisitType(), IncidentType())) - , Coords(coords.PrepareForInput(DeviceAdapter())) + , CellSet(cellSet.PrepareForInput(DeviceAdapter(), VisitType(), IncidentType(), token)) + , Coords(coords.PrepareForInput(DeviceAdapter(), token)) , PointDimensions(cellSet.GetPointDimensions()) { this->AxisPortals[0] = this->Coords.GetFirstPortal(); diff --git a/vtkm/exec/CellLocatorUniformGrid.h b/vtkm/exec/CellLocatorUniformGrid.h index a90c1d780..611f13dab 100644 --- a/vtkm/exec/CellLocatorUniformGrid.h +++ b/vtkm/exec/CellLocatorUniformGrid.h @@ -45,13 +45,14 @@ public: const vtkm::Vec3f invSpacing, const vtkm::Vec3f maxPoint, const vtkm::cont::ArrayHandleVirtualCoordinates& coords, - DeviceAdapter) + DeviceAdapter, + vtkm::cont::Token& token) : CellDims(cellDims) , PointDims(pointDims) , Origin(origin) , InvSpacing(invSpacing) , MaxPoint(maxPoint) - , Coords(coords.PrepareForInput(DeviceAdapter())) + , Coords(coords.PrepareForInput(DeviceAdapter(), token)) { } diff --git a/vtkm/rendering/Texture2D.h b/vtkm/rendering/Texture2D.h index 9c563fe5a..5382b5b48 100644 --- a/vtkm/rendering/Texture2D.h +++ b/vtkm/rendering/Texture2D.h @@ -107,10 +107,11 @@ public: vtkm::Id height, const TextureDataHandle& data, TextureFilterMode filterMode, - TextureWrapMode wrapMode) + TextureWrapMode wrapMode, + vtkm::cont::Token& token) : Width(width) , Height(height) - , Data(data.PrepareForInput(Device())) + , Data(data.PrepareForInput(Device(), token)) , FilterMode(filterMode) , WrapMode(wrapMode) { @@ -224,10 +225,12 @@ public: } template - VTKM_CONT Texture2DSamplerExecutionObject PrepareForExecution(Device) const + VTKM_CONT Texture2DSamplerExecutionObject PrepareForExecution( + Device, + vtkm::cont::Token& token) const { return Texture2DSamplerExecutionObject( - this->Width, this->Height, this->Data, this->FilterMode, this->WrapMode); + this->Width, this->Height, this->Data, this->FilterMode, this->WrapMode, token); } private: diff --git a/vtkm/rendering/Wireframer.h b/vtkm/rendering/Wireframer.h index e9f5a7cb7..09335d181 100644 --- a/vtkm/rendering/Wireframer.h +++ b/vtkm/rendering/Wireframer.h @@ -167,7 +167,8 @@ public: const vtkm::Range& fieldRange, const ColorMapHandle& colorMap, const AtomicPackedFrameBuffer& frameBuffer, - const vtkm::Range& clippingRange) + const vtkm::Range& clippingRange, + vtkm::cont::Token& token) : WorldToProjection(worldToProjection) , Width(width) , Height(height) @@ -176,9 +177,9 @@ public: , XOffset(xOffset) , YOffset(yOffset) , AssocPoints(assocPoints) - , ColorMap(colorMap.PrepareForInput(DeviceTag())) + , ColorMap(colorMap.PrepareForInput(DeviceTag(), token)) , ColorMapSize(vtkm::Float32(colorMap.GetNumberOfValues() - 1)) - , FrameBuffer(frameBuffer.PrepareForExecution(DeviceTag())) + , FrameBuffer(frameBuffer.PrepareForExecution(DeviceTag(), token)) , FieldMin(vtkm::Float32(fieldRange.Min)) { InverseFieldDelta = 1.0f / vtkm::Float32(fieldRange.Length()); @@ -479,19 +480,18 @@ private: vtkm::Id width = static_cast(Canvas->GetWidth()); vtkm::Id height = static_cast(Canvas->GetHeight()); vtkm::Id pixelCount = width * height; - FrameBuffer.PrepareForOutput(pixelCount, DeviceTag()); - if (ShowInternalZones && !IsOverlay) + if (this->ShowInternalZones && !this->IsOverlay) { vtkm::cont::ArrayHandleConstant clear(ClearValue, pixelCount); - vtkm::cont::Algorithm::Copy(clear, FrameBuffer); + vtkm::cont::Algorithm::Copy(clear, this->FrameBuffer); } else { - VTKM_ASSERT(SolidDepthBuffer.GetNumberOfValues() == pixelCount); + VTKM_ASSERT(this->SolidDepthBuffer.GetNumberOfValues() == pixelCount); CopyIntoFrameBuffer bufferCopy; vtkm::worklet::DispatcherMapField(bufferCopy) - .Invoke(Canvas->GetColorBuffer(), SolidDepthBuffer, FrameBuffer); + .Invoke(this->Canvas->GetColorBuffer(), this->SolidDepthBuffer, this->FrameBuffer); } // // detect a 2D camera and set the correct viewport. @@ -527,22 +527,26 @@ private: } const bool isAssocPoints = ScalarField.IsFieldPoint(); - EdgePlotter plotter(WorldToProjection, - width, - height, - subsetWidth, - subsetHeight, - xOffset, - yOffset, - isAssocPoints, - ScalarFieldRange, - ColorMap, - FrameBuffer, - Camera.GetClippingRange()); - vtkm::worklet::DispatcherMapField> plotterDispatcher(plotter); - plotterDispatcher.SetDevice(DeviceTag()); - plotterDispatcher.Invoke( - PointIndices, Coordinates, ScalarField.GetData().ResetTypes(vtkm::TypeListFieldScalar())); + { + vtkm::cont::Token token; + EdgePlotter plotter(WorldToProjection, + width, + height, + subsetWidth, + subsetHeight, + xOffset, + yOffset, + isAssocPoints, + ScalarFieldRange, + ColorMap, + FrameBuffer, + Camera.GetClippingRange(), + token); + vtkm::worklet::DispatcherMapField> plotterDispatcher(plotter); + plotterDispatcher.SetDevice(DeviceTag()); + plotterDispatcher.Invoke( + PointIndices, Coordinates, ScalarField.GetData().ResetTypes(vtkm::TypeListFieldScalar())); + } BufferConverter converter; vtkm::worklet::DispatcherMapField converterDispatcher(converter); diff --git a/vtkm/rendering/raytracing/ChannelBuffer.cxx b/vtkm/rendering/raytracing/ChannelBuffer.cxx index bc0380d3f..1ade78bd2 100644 --- a/vtkm/rendering/raytracing/ChannelBuffer.cxx +++ b/vtkm/rendering/raytracing/ChannelBuffer.cxx @@ -9,8 +9,10 @@ //============================================================================ #include +#include #include #include +#include #include #include @@ -156,8 +158,8 @@ public: , ChannelNum(channel) { } - using ControlSignature = void(FieldOut, WholeArrayIn); - using ExecutionSignature = void(_1, _2, WorkIndex); + using ControlSignature = void(FieldOut, WholeArrayIn, FieldIn); + using ExecutionSignature = void(_1, _2, _3); template VTKM_EXEC void operator()(T& outValue, const BufferPortalType& inBuffer, @@ -169,34 +171,6 @@ public: } }; //class Extract Channel -template -struct ExtractChannelFunctor -{ - ChannelBuffer* Self; - vtkm::cont::ArrayHandle Output; - vtkm::Int32 Channel; - - ExtractChannelFunctor(ChannelBuffer* self, - vtkm::cont::ArrayHandle output, - const vtkm::Int32 channel) - : Self(self) - , Output(output) - , Channel(channel) - { - } - - template - bool operator()(Device device) - { - Output.PrepareForOutput(Self->GetSize(), device); - vtkm::worklet::DispatcherMapField dispatcher( - ExtractChannel(Self->GetNumChannels(), Channel)); - dispatcher.SetDevice(Device()); - dispatcher.Invoke(Output, Self->Buffer); - return true; - } -}; - template ChannelBuffer ChannelBuffer::GetChannel(const vtkm::Int32 channel) { @@ -208,9 +182,13 @@ ChannelBuffer ChannelBuffer::GetChannel(const vtkm::Int32 { return output; } - ExtractChannelFunctor functor(this, output.Buffer, channel); - vtkm::cont::TryExecute(functor); + vtkm::cont::Invoker invoke; + invoke(ExtractChannel(this->GetNumChannels(), channel), + output.Buffer, + this->Buffer, + vtkm::cont::ArrayHandleIndex(this->GetSize())); + return output; } @@ -272,13 +250,15 @@ struct ExpandFunctorSignature template bool operator()(Device device) { - vtkm::Id totalSize = OutputLength * static_cast(NumChannels); - Output->Buffer.PrepareForOutput(totalSize, device); - ChannelBufferOperations::InitChannels(*Output, Signature, device); + vtkm::cont::Token token; - vtkm::worklet::DispatcherMapField dispatcher((Expand(NumChannels))); + vtkm::Id totalSize = this->OutputLength * static_cast(this->NumChannels); + this->Output->Buffer.PrepareForOutput(totalSize, device, token); + ChannelBufferOperations::InitChannels(*this->Output, this->Signature, device); + + vtkm::worklet::DispatcherMapField dispatcher((Expand(this->NumChannels))); dispatcher.SetDevice(Device()); - dispatcher.Invoke(Input, SparseIndexes, Output->Buffer); + dispatcher.Invoke(this->Input, this->SparseIndexes, this->Output->Buffer); return true; } @@ -313,13 +293,15 @@ struct ExpandFunctor template bool operator()(Device device) { - vtkm::Id totalSize = OutputLength * static_cast(NumChannels); - Output->Buffer.PrepareForOutput(totalSize, device); - ChannelBufferOperations::InitConst(*Output, InitVal, device); + vtkm::cont::Token token; - vtkm::worklet::DispatcherMapField dispatcher((Expand(NumChannels))); + vtkm::Id totalSize = this->OutputLength * static_cast(this->NumChannels); + this->Output->Buffer.PrepareForOutput(totalSize, device, token); + ChannelBufferOperations::InitConst(*this->Output, this->InitVal, device); + + vtkm::worklet::DispatcherMapField dispatcher((Expand(this->NumChannels))); dispatcher.SetDevice(Device()); - dispatcher.Invoke(Input, SparseIndexes, Output->Buffer); + dispatcher.Invoke(this->Input, this->SparseIndexes, this->Output->Buffer); return true; } diff --git a/vtkm/rendering/raytracing/ChannelBuffer.h b/vtkm/rendering/raytracing/ChannelBuffer.h index 25f1c6c92..9487dc75e 100644 --- a/vtkm/rendering/raytracing/ChannelBuffer.h +++ b/vtkm/rendering/raytracing/ChannelBuffer.h @@ -100,7 +100,8 @@ public: this->NumChannels = numChannels; this->Size = size; - this->Buffer.PrepareForOutput(this->Size * this->NumChannels, Device()); + vtkm::cont::Token token; + this->Buffer.PrepareForOutput(this->Size * this->NumChannels, Device(), token); } @@ -117,7 +118,8 @@ public: return; this->NumChannels = numChannels; - this->Buffer.PrepareForOutput(this->Size * this->NumChannels, Device()); + vtkm::cont::Token token; + this->Buffer.PrepareForOutput(this->Size * this->NumChannels, Device(), token); } template @@ -126,7 +128,8 @@ public: if (newSize < 0) throw vtkm::cont::ErrorBadValue("ChannelBuffer resize: Size must be greater than -1 "); this->Size = newSize; - this->Buffer.PrepareForOutput(this->Size * static_cast(NumChannels), device); + vtkm::cont::Token token; + this->Buffer.PrepareForOutput(this->Size * static_cast(NumChannels), device, token); } }; } diff --git a/vtkm/rendering/raytracing/CylinderIntersector.cxx b/vtkm/rendering/raytracing/CylinderIntersector.cxx index 835227529..6ba21c854 100644 --- a/vtkm/rendering/raytracing/CylinderIntersector.cxx +++ b/vtkm/rendering/raytracing/CylinderIntersector.cxx @@ -155,9 +155,11 @@ public: CylinderLeafIntersector() {} - CylinderLeafIntersector(const IdHandle& cylIds, const FloatHandle& radii) - : CylIds(cylIds.PrepareForInput(Device())) - , Radii(radii.PrepareForInput(Device())) + CylinderLeafIntersector(const IdHandle& cylIds, + const FloatHandle& radii, + vtkm::cont::Token& token) + : CylIds(cylIds.PrepareForInput(Device(), token)) + , Radii(radii.PrepareForInput(Device(), token)) { } @@ -309,9 +311,10 @@ public: } template - VTKM_CONT CylinderLeafIntersector PrepareForExecution(Device) const + VTKM_CONT CylinderLeafIntersector PrepareForExecution(Device, + vtkm::cont::Token& token) const { - return CylinderLeafIntersector(CylIds, Radii); + return CylinderLeafIntersector(this->CylIds, this->Radii, token); } }; diff --git a/vtkm/rendering/raytracing/MeshConnectivityBase.h b/vtkm/rendering/raytracing/MeshConnectivityBase.h index 30850689f..df4c5fc9a 100644 --- a/vtkm/rendering/raytracing/MeshConnectivityBase.h +++ b/vtkm/rendering/raytracing/MeshConnectivityBase.h @@ -178,12 +178,13 @@ public: const IdHandle& faceOffsets, const IdHandle& cellConn, const IdHandle& cellOffsets, - const UCharHandle& shapes) - : FaceConnPortal(faceConnectivity.PrepareForInput(Device())) - , FaceOffsetsPortal(faceOffsets.PrepareForInput(Device())) - , CellConnPortal(cellConn.PrepareForInput(Device())) - , CellOffsetsPortal(cellOffsets.PrepareForInput(Device())) - , ShapesPortal(shapes.PrepareForInput(Device())) + const UCharHandle& shapes, + vtkm::cont::Token& token) + : FaceConnPortal(faceConnectivity.PrepareForInput(Device(), token)) + , FaceOffsetsPortal(faceOffsets.PrepareForInput(Device(), token)) + , CellConnPortal(cellConn.PrepareForInput(Device(), token)) + , CellOffsetsPortal(cellOffsets.PrepareForInput(Device(), token)) + , ShapesPortal(shapes.PrepareForInput(Device(), token)) { } @@ -218,7 +219,7 @@ public: VTKM_EXEC vtkm::UInt8 GetCellShape(const vtkm::Id& cellId) const override { - BOUNDS_CHECK(ShapesPortal, cellId) + BOUNDS_CHECK(ShapesPortal, cellId); return ShapesPortal.Get(cellId); } @@ -253,10 +254,11 @@ public: CountingHandle& cellOffsets, vtkm::Int32 shapeId, vtkm::Int32 numIndices, - vtkm::Int32 numFaces) - : FaceConnPortal(faceConn.PrepareForInput(Device())) - , CellConnectivityPortal(cellConn.PrepareForInput(Device())) - , CellOffsetsPortal(cellOffsets.PrepareForInput(Device())) + vtkm::Int32 numFaces, + vtkm::cont::Token& token) + : FaceConnPortal(faceConn.PrepareForInput(Device(), token)) + , CellConnectivityPortal(cellConn.PrepareForInput(Device(), token)) + , CellOffsetsPortal(cellOffsets.PrepareForInput(Device(), token)) , ShapeId(shapeId) , NumIndices(numIndices) , NumFaces(numFaces) diff --git a/vtkm/rendering/raytracing/MeshConnectivityBuilder.cxx b/vtkm/rendering/raytracing/MeshConnectivityBuilder.cxx index 76f9fceae..852b0924f 100644 --- a/vtkm/rendering/raytracing/MeshConnectivityBuilder.cxx +++ b/vtkm/rendering/raytracing/MeshConnectivityBuilder.cxx @@ -792,9 +792,10 @@ struct StructuredTrianglesFunctor vtkm::cont::CellSetStructured<3>& cellSet) const { VTKM_IS_DEVICE_ADAPTER_TAG(Device); + vtkm::cont::Token token; vtkm::worklet::DispatcherMapField dispatch( StructuredExternalTriangles(cellSet.PrepareForInput( - Device(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint()))); + Device(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint(), token))); dispatch.SetDevice(Device()); dispatch.Invoke(counting, triangles); diff --git a/vtkm/rendering/raytracing/MeshConnectivityContainers.cxx b/vtkm/rendering/raytracing/MeshConnectivityContainers.cxx index 78c0dbd1d..a41b0b66b 100644 --- a/vtkm/rendering/raytracing/MeshConnectivityContainers.cxx +++ b/vtkm/rendering/raytracing/MeshConnectivityContainers.cxx @@ -38,9 +38,10 @@ VTKM_CONT void MeshConnContainer::FindEntryImpl(Ray& rays) Intersector.IntersectRays(rays, getCellIndex); } -MeshWrapper MeshConnContainer::PrepareForExecution(const vtkm::cont::DeviceAdapterId deviceId) +MeshWrapper MeshConnContainer::PrepareForExecution(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) { - return MeshWrapper(const_cast(this->Construct(deviceId))); + return MeshWrapper(const_cast(this->Construct(deviceId, token))); } void MeshConnContainer::FindEntry(Ray& rays) @@ -80,7 +81,8 @@ UnstructuredContainer::UnstructuredContainer(const vtkm::cont::CellSetExplicit<> UnstructuredContainer::~UnstructuredContainer(){}; const MeshConnectivityBase* UnstructuredContainer::Construct( - const vtkm::cont::DeviceAdapterId deviceId) + const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) { switch (deviceId.GetValue()) { @@ -92,10 +94,11 @@ const MeshConnectivityBase* UnstructuredContainer::Construct( this->FaceOffsets, this->CellConn, this->CellOffsets, - this->Shapes); + this->Shapes, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(OMP()); + return Handle.PrepareForExecution(OMP(), token); #endif #ifdef VTKM_ENABLE_TBB case VTKM_DEVICE_ADAPTER_TBB: @@ -105,10 +108,11 @@ const MeshConnectivityBase* UnstructuredContainer::Construct( this->FaceOffsets, this->CellConn, this->CellOffsets, - this->Shapes); + this->Shapes, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(TBB()); + return Handle.PrepareForExecution(TBB(), token); #endif #ifdef VTKM_ENABLE_CUDA case VTKM_DEVICE_ADAPTER_CUDA: @@ -118,10 +122,11 @@ const MeshConnectivityBase* UnstructuredContainer::Construct( this->FaceOffsets, this->CellConn, this->CellOffsets, - this->Shapes); + this->Shapes, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(CUDA()); + return Handle.PrepareForExecution(CUDA(), token); #endif case VTKM_DEVICE_ADAPTER_SERIAL: VTKM_FALLTHROUGH; @@ -132,10 +137,11 @@ const MeshConnectivityBase* UnstructuredContainer::Construct( this->FaceOffsets, this->CellConn, this->CellOffsets, - this->Shapes); + this->Shapes, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(SERIAL()); + return Handle.PrepareForExecution(SERIAL(), token); } } @@ -187,7 +193,8 @@ UnstructuredSingleContainer::UnstructuredSingleContainer( } const MeshConnectivityBase* UnstructuredSingleContainer::Construct( - const vtkm::cont::DeviceAdapterId deviceId) + const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) { switch (deviceId.GetValue()) { @@ -200,10 +207,11 @@ const MeshConnectivityBase* UnstructuredSingleContainer::Construct( this->CellOffsets, this->ShapeId, this->NumIndices, - this->NumFaces); + this->NumFaces, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(OMP()); + return Handle.PrepareForExecution(OMP(), token); #endif #ifdef VTKM_ENABLE_TBB case VTKM_DEVICE_ADAPTER_TBB: @@ -214,10 +222,11 @@ const MeshConnectivityBase* UnstructuredSingleContainer::Construct( this->CellOffsets, this->ShapeId, this->NumIndices, - this->NumFaces); + this->NumFaces, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(TBB()); + return Handle.PrepareForExecution(TBB(), token); #endif #ifdef VTKM_ENABLE_CUDA case VTKM_DEVICE_ADAPTER_CUDA: @@ -228,10 +237,11 @@ const MeshConnectivityBase* UnstructuredSingleContainer::Construct( this->CellOffsets, this->ShapeId, this->NumIndices, - this->NumFaces); + this->NumFaces, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(CUDA()); + return Handle.PrepareForExecution(CUDA(), token); #endif case VTKM_DEVICE_ADAPTER_SERIAL: VTKM_FALLTHROUGH; @@ -243,10 +253,11 @@ const MeshConnectivityBase* UnstructuredSingleContainer::Construct( this->CellOffsets, this->ShapeId, this->NumIndices, - this->NumFaces); + this->NumFaces, + token); Handle = make_MeshConnHandle(conn); } - return Handle.PrepareForExecution(SERIAL()); + return Handle.PrepareForExecution(SERIAL(), token); } } @@ -267,7 +278,8 @@ StructuredContainer::StructuredContainer(const vtkm::cont::CellSetStructured<3>& } const MeshConnectivityBase* StructuredContainer::Construct( - const vtkm::cont::DeviceAdapterId deviceId) + const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) { MeshConnStructured conn(CellDims, PointDims); @@ -277,20 +289,20 @@ const MeshConnectivityBase* StructuredContainer::Construct( { #ifdef VTKM_ENABLE_OPENMP case VTKM_DEVICE_ADAPTER_OPENMP: - return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagOpenMP()); + return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagOpenMP(), token); #endif #ifdef VTKM_ENABLE_TBB case VTKM_DEVICE_ADAPTER_TBB: - return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagTBB()); + return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagTBB(), token); #endif #ifdef VTKM_ENABLE_CUDA case VTKM_DEVICE_ADAPTER_CUDA: - return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagCuda()); + return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagCuda(), token); #endif case VTKM_DEVICE_ADAPTER_SERIAL: VTKM_FALLTHROUGH; default: - return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagSerial()); + return Handle.PrepareForExecution(vtkm::cont::DeviceAdapterTagSerial(), token); } } } diff --git a/vtkm/rendering/raytracing/MeshConnectivityContainers.h b/vtkm/rendering/raytracing/MeshConnectivityContainers.h index f0bbd63a5..2a19f34b8 100644 --- a/vtkm/rendering/raytracing/MeshConnectivityContainers.h +++ b/vtkm/rendering/raytracing/MeshConnectivityContainers.h @@ -27,9 +27,11 @@ public: MeshConnContainer(); virtual ~MeshConnContainer(); - virtual const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId) = 0; + virtual const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) = 0; - MeshWrapper PrepareForExecution(const vtkm::cont::DeviceAdapterId deviceId); + MeshWrapper PrepareForExecution(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token); template VTKM_CONT void FindEntryImpl(Ray& rays); @@ -79,7 +81,8 @@ public: virtual ~UnstructuredContainer(); - const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId); + const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) override; }; class StructuredContainer : public MeshConnContainer @@ -102,7 +105,8 @@ public: const vtkm::cont::CoordinateSystem& coords, Id4Handle& triangles); - const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId) override; + const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) override; }; //structure mesh conn @@ -138,7 +142,8 @@ public: IdHandle& faceConn, Id4Handle& externalFaces); - const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId) override; + const MeshConnectivityBase* Construct(const vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) override; }; //UnstructuredSingleContainer } diff --git a/vtkm/rendering/raytracing/QuadIntersector.cxx b/vtkm/rendering/raytracing/QuadIntersector.cxx index b43c8ccbc..2117fe583 100644 --- a/vtkm/rendering/raytracing/QuadIntersector.cxx +++ b/vtkm/rendering/raytracing/QuadIntersector.cxx @@ -109,8 +109,8 @@ public: QuadLeafIntersector() {} - QuadLeafIntersector(const IdHandle& quadIds) - : QuadIds(quadIds.PrepareForInput(Device())) + QuadLeafIntersector(const IdHandle& quadIds, vtkm::cont::Token& token) + : QuadIds(quadIds.PrepareForInput(Device(), token)) { } @@ -306,9 +306,9 @@ public: } template - VTKM_CONT QuadLeafIntersector PrepareForExecution(Device) const + VTKM_CONT QuadLeafIntersector PrepareForExecution(Device, vtkm::cont::Token& token) const { - return QuadLeafIntersector(QuadIds); + return QuadLeafIntersector(QuadIds, token); } }; diff --git a/vtkm/rendering/raytracing/Ray.h b/vtkm/rendering/raytracing/Ray.h index 6742fa380..10f869cb9 100644 --- a/vtkm/rendering/raytracing/Ray.h +++ b/vtkm/rendering/raytracing/Ray.h @@ -141,17 +141,18 @@ public: return; } + vtkm::cont::Token token; IntersectionDataEnabled = true; - IntersectionX.PrepareForOutput(NumRays, Device()); - IntersectionY.PrepareForOutput(NumRays, Device()); - IntersectionZ.PrepareForOutput(NumRays, Device()); - U.PrepareForOutput(NumRays, Device()); - V.PrepareForOutput(NumRays, Device()); - Scalar.PrepareForOutput(NumRays, Device()); + IntersectionX.PrepareForOutput(NumRays, Device(), token); + IntersectionY.PrepareForOutput(NumRays, Device(), token); + IntersectionZ.PrepareForOutput(NumRays, Device(), token); + U.PrepareForOutput(NumRays, Device(), token); + V.PrepareForOutput(NumRays, Device(), token); + Scalar.PrepareForOutput(NumRays, Device(), token); - NormalX.PrepareForOutput(NumRays, Device()); - NormalY.PrepareForOutput(NumRays, Device()); - NormalZ.PrepareForOutput(NumRays, Device()); + NormalX.PrepareForOutput(NumRays, Device(), token); + NormalY.PrepareForOutput(NumRays, Device(), token); + NormalZ.PrepareForOutput(NumRays, Device(), token); } void DisableIntersectionData() @@ -206,39 +207,40 @@ public: VTKM_CONT void Resize(const vtkm::Int32 size, Device) { NumRays = size; + vtkm::cont::Token token; if (IntersectionDataEnabled) { - IntersectionX.PrepareForOutput(NumRays, Device()); - IntersectionY.PrepareForOutput(NumRays, Device()); - IntersectionZ.PrepareForOutput(NumRays, Device()); + IntersectionX.PrepareForOutput(NumRays, Device(), token); + IntersectionY.PrepareForOutput(NumRays, Device(), token); + IntersectionZ.PrepareForOutput(NumRays, Device(), token); - U.PrepareForOutput(NumRays, Device()); - V.PrepareForOutput(NumRays, Device()); + U.PrepareForOutput(NumRays, Device(), token); + V.PrepareForOutput(NumRays, Device(), token); - Scalar.PrepareForOutput(NumRays, Device()); + Scalar.PrepareForOutput(NumRays, Device(), token); - NormalX.PrepareForOutput(NumRays, Device()); - NormalY.PrepareForOutput(NumRays, Device()); - NormalZ.PrepareForOutput(NumRays, Device()); + NormalX.PrepareForOutput(NumRays, Device(), token); + NormalY.PrepareForOutput(NumRays, Device(), token); + NormalZ.PrepareForOutput(NumRays, Device(), token); } - OriginX.PrepareForOutput(NumRays, Device()); - OriginY.PrepareForOutput(NumRays, Device()); - OriginZ.PrepareForOutput(NumRays, Device()); + OriginX.PrepareForOutput(NumRays, Device(), token); + OriginY.PrepareForOutput(NumRays, Device(), token); + OriginZ.PrepareForOutput(NumRays, Device(), token); - DirX.PrepareForOutput(NumRays, Device()); - DirY.PrepareForOutput(NumRays, Device()); - DirZ.PrepareForOutput(NumRays, Device()); + DirX.PrepareForOutput(NumRays, Device(), token); + DirY.PrepareForOutput(NumRays, Device(), token); + DirZ.PrepareForOutput(NumRays, Device(), token); - Distance.PrepareForOutput(NumRays, Device()); + Distance.PrepareForOutput(NumRays, Device(), token); - MinDistance.PrepareForOutput(NumRays, Device()); - MaxDistance.PrepareForOutput(NumRays, Device()); - Status.PrepareForOutput(NumRays, Device()); + MinDistance.PrepareForOutput(NumRays, Device(), token); + MaxDistance.PrepareForOutput(NumRays, Device(), token); + Status.PrepareForOutput(NumRays, Device(), token); - HitIdx.PrepareForOutput(NumRays, Device()); - PixelIdx.PrepareForOutput(NumRays, Device()); + HitIdx.PrepareForOutput(NumRays, Device(), token); + PixelIdx.PrepareForOutput(NumRays, Device(), token); Intersection = vtkm::cont::make_ArrayHandleCompositeVector(IntersectionX, IntersectionY, IntersectionZ); diff --git a/vtkm/rendering/raytracing/RayOperations.h b/vtkm/rendering/raytracing/RayOperations.h index 1f1a23fc3..020645a17 100644 --- a/vtkm/rendering/raytracing/RayOperations.h +++ b/vtkm/rendering/raytracing/RayOperations.h @@ -289,35 +289,36 @@ public: return; //nothing to do rays.NumRays = newSize; + vtkm::cont::Token token; if (rays.IntersectionDataEnabled) { - rays.IntersectionX.PrepareForOutput(rays.NumRays, Device()); - rays.IntersectionY.PrepareForOutput(rays.NumRays, Device()); - rays.IntersectionZ.PrepareForOutput(rays.NumRays, Device()); - rays.U.PrepareForOutput(rays.NumRays, Device()); - rays.V.PrepareForOutput(rays.NumRays, Device()); - rays.Scalar.PrepareForOutput(rays.NumRays, Device()); + rays.IntersectionX.PrepareForOutput(rays.NumRays, Device(), token); + rays.IntersectionY.PrepareForOutput(rays.NumRays, Device(), token); + rays.IntersectionZ.PrepareForOutput(rays.NumRays, Device(), token); + rays.U.PrepareForOutput(rays.NumRays, Device(), token); + rays.V.PrepareForOutput(rays.NumRays, Device(), token); + rays.Scalar.PrepareForOutput(rays.NumRays, Device(), token); - rays.NormalX.PrepareForOutput(rays.NumRays, Device()); - rays.NormalY.PrepareForOutput(rays.NumRays, Device()); - rays.NormalZ.PrepareForOutput(rays.NumRays, Device()); + rays.NormalX.PrepareForOutput(rays.NumRays, Device(), token); + rays.NormalY.PrepareForOutput(rays.NumRays, Device(), token); + rays.NormalZ.PrepareForOutput(rays.NumRays, Device(), token); } - rays.OriginX.PrepareForOutput(rays.NumRays, Device()); - rays.OriginY.PrepareForOutput(rays.NumRays, Device()); - rays.OriginZ.PrepareForOutput(rays.NumRays, Device()); + rays.OriginX.PrepareForOutput(rays.NumRays, Device(), token); + rays.OriginY.PrepareForOutput(rays.NumRays, Device(), token); + rays.OriginZ.PrepareForOutput(rays.NumRays, Device(), token); - rays.DirX.PrepareForOutput(rays.NumRays, Device()); - rays.DirY.PrepareForOutput(rays.NumRays, Device()); - rays.DirZ.PrepareForOutput(rays.NumRays, Device()); + rays.DirX.PrepareForOutput(rays.NumRays, Device(), token); + rays.DirY.PrepareForOutput(rays.NumRays, Device(), token); + rays.DirZ.PrepareForOutput(rays.NumRays, Device(), token); - rays.Distance.PrepareForOutput(rays.NumRays, Device()); - rays.MinDistance.PrepareForOutput(rays.NumRays, Device()); - rays.MaxDistance.PrepareForOutput(rays.NumRays, Device()); - rays.Status.PrepareForOutput(rays.NumRays, Device()); - rays.HitIdx.PrepareForOutput(rays.NumRays, Device()); - rays.PixelIdx.PrepareForOutput(rays.NumRays, Device()); + rays.Distance.PrepareForOutput(rays.NumRays, Device(), token); + rays.MinDistance.PrepareForOutput(rays.NumRays, Device(), token); + rays.MaxDistance.PrepareForOutput(rays.NumRays, Device(), token); + rays.Status.PrepareForOutput(rays.NumRays, Device(), token); + rays.HitIdx.PrepareForOutput(rays.NumRays, Device(), token); + rays.PixelIdx.PrepareForOutput(rays.NumRays, Device(), token); const size_t bufferCount = static_cast(rays.Buffers.size()); for (size_t i = 0; i < bufferCount; ++i) diff --git a/vtkm/rendering/raytracing/RayTracingTypeDefs.h b/vtkm/rendering/raytracing/RayTracingTypeDefs.h index 883a0897c..a3ff2edad 100644 --- a/vtkm/rendering/raytracing/RayTracingTypeDefs.h +++ b/vtkm/rendering/raytracing/RayTracingTypeDefs.h @@ -27,10 +27,7 @@ namespace rendering { // A more useful bounds check that tells you where it happened #ifndef NDEBUG -#define BOUNDS_CHECK(HANDLE, INDEX) \ - { \ - BoundsCheck((HANDLE), (INDEX), __FILE__, __LINE__); \ - } +#define BOUNDS_CHECK(HANDLE, INDEX) BoundsCheck((HANDLE), (INDEX), __FILE__, __LINE__) #else #define BOUNDS_CHECK(HANDLE, INDEX) #endif diff --git a/vtkm/rendering/raytracing/SphereIntersector.cxx b/vtkm/rendering/raytracing/SphereIntersector.cxx index 7f1653947..a9d8ffd7f 100644 --- a/vtkm/rendering/raytracing/SphereIntersector.cxx +++ b/vtkm/rendering/raytracing/SphereIntersector.cxx @@ -133,9 +133,11 @@ public: SphereLeafIntersector() {} - SphereLeafIntersector(const IdHandle& pointIds, const FloatHandle& radii) - : PointIds(pointIds.PrepareForInput(Device())) - , Radii(radii.PrepareForInput(Device())) + SphereLeafIntersector(const IdHandle& pointIds, + const FloatHandle& radii, + vtkm::cont::Token& token) + : PointIds(pointIds.PrepareForInput(Device(), token)) + , Radii(radii.PrepareForInput(Device(), token)) { } @@ -201,9 +203,10 @@ public: } template - VTKM_CONT SphereLeafIntersector PrepareForExecution(Device) const + VTKM_CONT SphereLeafIntersector PrepareForExecution(Device, + vtkm::cont::Token& token) const { - return SphereLeafIntersector(PointIds, Radii); + return SphereLeafIntersector(this->PointIds, this->Radii, token); } }; diff --git a/vtkm/rendering/raytracing/TriangleIntersector.cxx b/vtkm/rendering/raytracing/TriangleIntersector.cxx index a9746a2de..dd71f072a 100644 --- a/vtkm/rendering/raytracing/TriangleIntersector.cxx +++ b/vtkm/rendering/raytracing/TriangleIntersector.cxx @@ -40,8 +40,8 @@ public: public: WaterTightLeafIntersector() = default; - WaterTightLeafIntersector(const Id4Handle& triangles) - : Triangles(triangles.PrepareForInput(Device())) + WaterTightLeafIntersector(const Id4Handle& triangles, vtkm::cont::Token& token) + : Triangles(triangles.PrepareForInput(Device(), token)) { } @@ -93,8 +93,8 @@ public: public: MollerTriLeafIntersector() {} - MollerTriLeafIntersector(const Id4Handle& triangles) - : Triangles(triangles.PrepareForInput(Device())) + MollerTriLeafIntersector(const Id4Handle& triangles, vtkm::cont::Token& token) + : Triangles(triangles.PrepareForInput(Device(), token)) { } @@ -148,9 +148,10 @@ public: } template - VTKM_CONT MollerTriLeafIntersector PrepareForExecution(Device) const + VTKM_CONT MollerTriLeafIntersector PrepareForExecution(Device, + vtkm::cont::Token& token) const { - return MollerTriLeafIntersector(Triangles); + return MollerTriLeafIntersector(this->Triangles, token); } }; @@ -167,9 +168,10 @@ public: } template - VTKM_CONT WaterTightLeafIntersector PrepareForExecution(Device) const + VTKM_CONT WaterTightLeafIntersector PrepareForExecution(Device, + vtkm::cont::Token& token) const { - return WaterTightLeafIntersector(Triangles); + return WaterTightLeafIntersector(this->Triangles, token); } }; diff --git a/vtkm/rendering/raytracing/VolumeRendererStructured.cxx b/vtkm/rendering/raytracing/VolumeRendererStructured.cxx index 622c23f57..086ddddc9 100644 --- a/vtkm/rendering/raytracing/VolumeRendererStructured.cxx +++ b/vtkm/rendering/raytracing/VolumeRendererStructured.cxx @@ -57,11 +57,13 @@ protected: public: RectilinearLocator(const CartesianArrayHandle& coordinates, - vtkm::cont::CellSetStructured<3>& cellset) - : Coordinates(coordinates.PrepareForInput(Device())) + vtkm::cont::CellSetStructured<3>& cellset, + vtkm::cont::Token& token) + : Coordinates(coordinates.PrepareForInput(Device(), token)) , Conn(cellset.PrepareForInput(Device(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint())) + vtkm::TopologyElementTagPoint(), + token)) { CoordPortals[0] = Coordinates.GetFirstPortal(); CoordPortals[1] = Coordinates.GetSecondPortal(); @@ -199,11 +201,14 @@ protected: Conn; public: - UniformLocator(const UniformArrayHandle& coordinates, vtkm::cont::CellSetStructured<3>& cellset) - : Coordinates(coordinates.PrepareForInput(Device())) + UniformLocator(const UniformArrayHandle& coordinates, + vtkm::cont::CellSetStructured<3>& cellset, + vtkm::cont::Token& token) + : Coordinates(coordinates.PrepareForInput(Device(), token)) , Conn(cellset.PrepareForInput(Device(), vtkm::TopologyElementTagCell(), - vtkm::TopologyElementTagPoint())) + vtkm::TopologyElementTagPoint(), + token)) { Origin = Coordinates.GetOrigin(); PointDimensions = Conn.GetPointDimensions(); @@ -310,8 +315,9 @@ public: const vtkm::Float32& minScalar, const vtkm::Float32& maxScalar, const vtkm::Float32& sampleDistance, - const LocatorType& locator) - : ColorMap(colorMap.PrepareForInput(DeviceAdapterTag())) + const LocatorType& locator, + vtkm::cont::Token& token) + : ColorMap(colorMap.PrepareForInput(DeviceAdapterTag(), token)) , MinScalar(minScalar) , SampleDistance(sampleDistance) , InverseDeltaScalar(minScalar) @@ -503,8 +509,9 @@ public: const vtkm::Float32& minScalar, const vtkm::Float32& maxScalar, const vtkm::Float32& sampleDistance, - const LocatorType& locator) - : ColorMap(colorMap.PrepareForInput(DeviceAdapterTag())) + const LocatorType& locator, + vtkm::cont::Token& token) + : ColorMap(colorMap.PrepareForInput(DeviceAdapterTag(), token)) , MinScalar(minScalar) , SampleDistance(sampleDistance) , InverseDeltaScalar(minScalar) @@ -813,9 +820,10 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray

(); - UniformLocator locator(vertices, Cellset); + UniformLocator locator(vertices, Cellset, token); if (isAssocPoints) { @@ -824,7 +832,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray

(); - RectilinearLocator locator(vertices, Cellset); + RectilinearLocator locator(vertices, Cellset, token); if (isAssocPoints) { vtkm::worklet::DispatcherMapField>> @@ -862,7 +873,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray

::PortalConst, typename IdArray::ExecutionTypes::PortalConst, typename IdArray::ExecutionTypes::Portal>; - Functor functor{ needles.PrepareForInput(Device{}), - hayStack.PrepareForInput(Device{}), - results.PrepareForOutput(needles.GetNumberOfValues(), Device{}) }; + Functor functor{ needles.PrepareForInput(Device{}, token), + hayStack.PrepareForInput(Device{}, token), + results.PrepareForOutput(needles.GetNumberOfValues(), Device{}, token) }; Algo::Schedule(functor, needles.GetNumberOfValues()); + token.DetachFromAll(); + // Verify: auto needlesPortal = needles.GetPortalConstControl(); auto hayStackPortal = hayStack.GetPortalConstControl(); @@ -132,15 +136,19 @@ struct TestLowerBound IdArray hayStack = vtkm::cont::make_ArrayHandle(hayStackData); IdArray results; + vtkm::cont::Token token; + using Functor = Impl::PortalConst, typename IdArray::ExecutionTypes::PortalConst, typename IdArray::ExecutionTypes::Portal>; - Functor functor{ needles.PrepareForInput(Device{}), - hayStack.PrepareForInput(Device{}), - results.PrepareForOutput(needles.GetNumberOfValues(), Device{}) }; + Functor functor{ needles.PrepareForInput(Device{}, token), + hayStack.PrepareForInput(Device{}, token), + results.PrepareForOutput(needles.GetNumberOfValues(), Device{}, token) }; Algo::Schedule(functor, needles.GetNumberOfValues()); + token.DetachFromAll(); + // Verify: auto resultsPortal = results.GetPortalConstControl(); for (vtkm::Id i = 0; i < needles.GetNumberOfValues(); ++i) @@ -187,15 +195,19 @@ struct TestUpperBound IdArray hayStack = vtkm::cont::make_ArrayHandle(hayStackData); IdArray results; + vtkm::cont::Token token; + using Functor = Impl::PortalConst, typename IdArray::ExecutionTypes::PortalConst, typename IdArray::ExecutionTypes::Portal>; - Functor functor{ needles.PrepareForInput(Device{}), - hayStack.PrepareForInput(Device{}), - results.PrepareForOutput(needles.GetNumberOfValues(), Device{}) }; + Functor functor{ needles.PrepareForInput(Device{}, token), + hayStack.PrepareForInput(Device{}, token), + results.PrepareForOutput(needles.GetNumberOfValues(), Device{}, token) }; Algo::Schedule(functor, needles.GetNumberOfValues()); + token.DetachFromAll(); + // Verify: auto resultsPortal = results.GetPortalConstControl(); for (vtkm::Id i = 0; i < needles.GetNumberOfValues(); ++i) diff --git a/vtkm/worklet/Clip.h b/vtkm/worklet/Clip.h index 53fef3427..38a1086ad 100644 --- a/vtkm/worklet/Clip.h +++ b/vtkm/worklet/Clip.h @@ -139,11 +139,12 @@ public: vtkm::cont::ArrayHandle numberOfIndices, vtkm::cont::ArrayHandle connectivity, vtkm::cont::ArrayHandle offsets, - ClipStats stats) - : Shapes(shapes.PrepareForOutput(stats.NumberOfCells, Device())) - , NumberOfIndices(numberOfIndices.PrepareForOutput(stats.NumberOfCells, Device())) - , Connectivity(connectivity.PrepareForOutput(stats.NumberOfIndices, Device())) - , Offsets(offsets.PrepareForOutput(stats.NumberOfCells, Device())) + ClipStats stats, + vtkm::cont::Token& token) + : Shapes(shapes.PrepareForOutput(stats.NumberOfCells, Device(), token)) + , NumberOfIndices(numberOfIndices.PrepareForOutput(stats.NumberOfCells, Device(), token)) + , Connectivity(connectivity.PrepareForOutput(stats.NumberOfIndices, Device(), token)) + , Offsets(offsets.PrepareForOutput(stats.NumberOfCells, Device(), token)) { } @@ -196,10 +197,12 @@ public: } template - VTKM_CONT ExecutionConnectivityExplicit PrepareForExecution(Device) const + VTKM_CONT ExecutionConnectivityExplicit PrepareForExecution( + Device, + vtkm::cont::Token& token) const { ExecutionConnectivityExplicit execConnectivity( - this->Shapes, this->NumberOfIndices, this->Connectivity, this->Offsets, this->Stats); + this->Shapes, this->NumberOfIndices, this->Connectivity, this->Offsets, this->Stats, token); return execConnectivity; } diff --git a/vtkm/worklet/Contour.h b/vtkm/worklet/Contour.h index c1315f0f7..b357e9ae9 100644 --- a/vtkm/worklet/Contour.h +++ b/vtkm/worklet/Contour.h @@ -152,11 +152,12 @@ public: vtkm::cont::ArrayHandle& interpWeights, vtkm::cont::ArrayHandle& interpIds, vtkm::cont::ArrayHandle& interpCellIds, - vtkm::cont::ArrayHandle& interpContourId) - : InterpWeightsPortal(interpWeights.PrepareForOutput(3 * size, DeviceAdapter())) - , InterpIdPortal(interpIds.PrepareForOutput(3 * size, DeviceAdapter())) - , InterpCellIdPortal(interpCellIds.PrepareForOutput(3 * size, DeviceAdapter())) - , InterpContourPortal(interpContourId.PrepareForOutput(3 * size, DeviceAdapter())) + vtkm::cont::ArrayHandle& interpContourId, + vtkm::cont::Token& token) + : InterpWeightsPortal(interpWeights.PrepareForOutput(3 * size, DeviceAdapter(), token)) + , InterpIdPortal(interpIds.PrepareForOutput(3 * size, DeviceAdapter(), token)) + , InterpCellIdPortal(interpCellIds.PrepareForOutput(3 * size, DeviceAdapter(), token)) + , InterpContourPortal(interpContourId.PrepareForOutput(3 * size, DeviceAdapter(), token)) { // Interp needs to be 3 times longer than size as they are per point of the // output triangle @@ -182,10 +183,14 @@ public: } template - VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter) + VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token) { - return ExecObject( - this->Size, this->InterpWeights, this->InterpIds, this->InterpCellIds, this->InterpContourId); + return ExecObject(this->Size, + this->InterpWeights, + this->InterpIds, + this->InterpCellIds, + this->InterpContourId, + token); } private: diff --git a/vtkm/worklet/DispatcherStreamingMapField.h b/vtkm/worklet/DispatcherStreamingMapField.h index d7be3c318..9c41ef24d 100644 --- a/vtkm/worklet/DispatcherStreamingMapField.h +++ b/vtkm/worklet/DispatcherStreamingMapField.h @@ -300,13 +300,14 @@ private: // Replace the parameters in the invocation with the execution object and // pass to next step of Invoke. Also add the scatter information. - this->InvokeSchedule(invocation.ChangeParameters(execObjectParameters) - .ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device)) - .ChangeVisitArray(visitArray.PrepareForInput(device)) - .ChangeThreadToOutputMap(threadToOutputMap.PrepareForInput(device)), - outputRange, - globalIndexOffset, - device); + this->InvokeSchedule( + invocation.ChangeParameters(execObjectParameters) + .ChangeOutputToInputMap(outputToInputMap.PrepareForInput(device, token)) + .ChangeVisitArray(visitArray.PrepareForInput(device, token)) + .ChangeThreadToOutputMap(threadToOutputMap.PrepareForInput(device, token)), + outputRange, + globalIndexOffset, + device); } template diff --git a/vtkm/worklet/PointMerge.h b/vtkm/worklet/PointMerge.h index 1809003ab..939dde941 100644 --- a/vtkm/worklet/PointMerge.h +++ b/vtkm/worklet/PointMerge.h @@ -136,8 +136,7 @@ public: // Because this class is a POD, we can reuse it in both control and execution environments. - template - BinLocator PrepareForExecution(Device) const + BinLocator PrepareForExecution(vtkm::cont::DeviceAdapterId, vtkm::cont::Token&) const { return *this; } diff --git a/vtkm/worklet/StableSortIndices.h b/vtkm/worklet/StableSortIndices.h index 15ea0c8bf..9891188f1 100644 --- a/vtkm/worklet/StableSortIndices.h +++ b/vtkm/worklet/StableSortIndices.h @@ -77,9 +77,9 @@ struct StableSortIndices template IndirectSortPredicate::PortalConst> - PrepareForExecution(Device) const + PrepareForExecution(Device, vtkm::cont::Token& token) const { - auto keyPortal = this->KeyArray.PrepareForInput(Device()); + auto keyPortal = this->KeyArray.PrepareForInput(Device(), token); return IndirectSortPredicate(keyPortal); } }; @@ -116,9 +116,9 @@ struct StableSortIndices template IndirectUniquePredicate::PortalConst> - PrepareForExecution(Device) const + PrepareForExecution(Device, vtkm::cont::Token& token) const { - auto keyPortal = this->KeyArray.PrepareForInput(Device()); + auto keyPortal = this->KeyArray.PrepareForInput(Device(), token); return IndirectUniquePredicate(keyPortal); } }; diff --git a/vtkm/worklet/clip/ClipTables.h b/vtkm/worklet/clip/ClipTables.h index 851b77e26..ee16f7949 100644 --- a/vtkm/worklet/clip/ClipTables.h +++ b/vtkm/worklet/clip/ClipTables.h @@ -2530,12 +2530,13 @@ public: } template - DevicePortal PrepareForExecution(DeviceAdapter) + DevicePortal PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token) { DevicePortal portal; - portal.ClipTablesDataPortal = this->ClipTablesDataArray.PrepareForInput(DeviceAdapter()); - portal.ClipTablesIndicesPortal = this->ClipTablesIndicesArray.PrepareForInput(DeviceAdapter()); - portal.CellEdgesPortal = this->CellEdgesArray.PrepareForInput(DeviceAdapter()); + portal.ClipTablesDataPortal = this->ClipTablesDataArray.PrepareForInput(DeviceAdapter(), token); + portal.ClipTablesIndicesPortal = + this->ClipTablesIndicesArray.PrepareForInput(DeviceAdapter(), token); + portal.CellEdgesPortal = this->CellEdgesArray.PrepareForInput(DeviceAdapter(), token); return portal; } diff --git a/vtkm/worklet/contour/ContourTables.h b/vtkm/worklet/contour/ContourTables.h index 74f75fe51..2c7ec6875 100644 --- a/vtkm/worklet/contour/ContourTables.h +++ b/vtkm/worklet/contour/ContourTables.h @@ -517,15 +517,15 @@ public: } template - ExecObject PrepareForExecution(DeviceAdapter) + ExecObject PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token) { ExecObject execObject; execObject.NumVerticesPerCellPortal = - this->NumVerticesPerCellArray.PrepareForInput(DeviceAdapter()); + this->NumVerticesPerCellArray.PrepareForInput(DeviceAdapter(), token); execObject.NumTrianglesTableOffsetPortal = - this->NumTrianglesTableOffsetArray.PrepareForInput(DeviceAdapter()); + this->NumTrianglesTableOffsetArray.PrepareForInput(DeviceAdapter(), token); execObject.NumTrianglesTablePortal = - this->NumTrianglesTableArray.PrepareForInput(DeviceAdapter()); + this->NumTrianglesTableArray.PrepareForInput(DeviceAdapter(), token); return execObject; } @@ -589,14 +589,16 @@ public: }; template - ExecObject PrepareForExecution(DeviceAdapter) + ExecObject PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token) { ExecObject execObject; - execObject.EdgeTablePortal = this->EdgeTableArray.PrepareForInput(DeviceAdapter()); - execObject.EdgeTableOffsetPortal = this->EdgeTableOffsetArray.PrepareForInput(DeviceAdapter()); - execObject.TriangleTablePortal = this->TriangleTableArray.PrepareForInput(DeviceAdapter()); + execObject.EdgeTablePortal = this->EdgeTableArray.PrepareForInput(DeviceAdapter(), token); + execObject.EdgeTableOffsetPortal = + this->EdgeTableOffsetArray.PrepareForInput(DeviceAdapter(), token); + execObject.TriangleTablePortal = + this->TriangleTableArray.PrepareForInput(DeviceAdapter(), token); execObject.TriangleTableOffsetPortal = - this->TriangleTableOffsetArray.PrepareForInput(DeviceAdapter()); + this->TriangleTableOffsetArray.PrepareForInput(DeviceAdapter(), token); return execObject; } diff --git a/vtkm/worklet/contourtree/EdgePeakComparator.h b/vtkm/worklet/contourtree/EdgePeakComparator.h index 6812e5ab6..152a79657 100644 --- a/vtkm/worklet/contourtree/EdgePeakComparator.h +++ b/vtkm/worklet/contourtree/EdgePeakComparator.h @@ -179,13 +179,14 @@ public: }; template - VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter) const + VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter, + vtkm::cont::Token& token) const { - return ExecObject(this->Values.PrepareForInput(DeviceAdapter()), - this->ValueIndex.PrepareForInput(DeviceAdapter()), - this->EdgeFar.PrepareForInput(DeviceAdapter()), - this->EdgeNear.PrepareForInput(DeviceAdapter()), - this->ArcArray.PrepareForInput(DeviceAdapter()), + return ExecObject(this->Values.PrepareForInput(DeviceAdapter(), token), + this->ValueIndex.PrepareForInput(DeviceAdapter(), token), + this->EdgeFar.PrepareForInput(DeviceAdapter(), token), + this->EdgeNear.PrepareForInput(DeviceAdapter(), token), + this->ArcArray.PrepareForInput(DeviceAdapter(), token), this->IsJoinGraph); } }; // EdgePeakComparator diff --git a/vtkm/worklet/contourtree/VertexMergeComparator.h b/vtkm/worklet/contourtree/VertexMergeComparator.h index 56e6db092..ec5de295a 100644 --- a/vtkm/worklet/contourtree/VertexMergeComparator.h +++ b/vtkm/worklet/contourtree/VertexMergeComparator.h @@ -163,10 +163,11 @@ public: }; template - VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter) const + VTKM_CONT ExecObject PrepareForExecution(DeviceAdapter, + vtkm::cont::Token& token) const { - return ExecObject(this->Values.PrepareForInput(DeviceAdapter()), - this->Extrema.PrepareForInput(DeviceAdapter()), + return ExecObject(this->Values.PrepareForInput(DeviceAdapter(), token), + this->Extrema.PrepareForInput(DeviceAdapter(), token), this->IsJoinTree); } }; // VertexMergeComparator diff --git a/vtkm/worklet/contourtree_augmented/ContourTreeMaker.h b/vtkm/worklet/contourtree_augmented/ContourTreeMaker.h index 7c7c0c6fd..f5388a9e5 100644 --- a/vtkm/worklet/contourtree_augmented/ContourTreeMaker.h +++ b/vtkm/worklet/contourtree_augmented/ContourTreeMaker.h @@ -725,6 +725,7 @@ struct LeafChainsToContourTree template bool operator()(DeviceAdapter device, Args&&... args) const { + vtkm::cont::Token token; contourtree_maker_inc_ns::TransferLeafChains_TransferToContourTree worklet( this->NumIterations, // (input) this->IsJoin, // (input) @@ -732,7 +733,8 @@ struct LeafChainsToContourTree this->Indegree, // (input) this->Outbound, // (input) this->Inbound, // (input) - this->Inwards); // (input) + this->Inwards, // (input) + token); vtkm::worklet::DispatcherMapField dispatcher(worklet); dispatcher.SetDevice(device); dispatcher.Invoke(std::forward(args)...); diff --git a/vtkm/worklet/contourtree_augmented/activegraph/EdgePeakComparator.h b/vtkm/worklet/contourtree_augmented/activegraph/EdgePeakComparator.h index f81b62bdc..ed3f0f841 100644 --- a/vtkm/worklet/contourtree_augmented/activegraph/EdgePeakComparator.h +++ b/vtkm/worklet/contourtree_augmented/activegraph/EdgePeakComparator.h @@ -77,11 +77,14 @@ public: // constructor - takes vectors as parameters VTKM_CONT - EdgePeakComparatorImpl(const IdArrayType& edgeFar, const IdArrayType& edgeNear, bool joinGraph) + EdgePeakComparatorImpl(const IdArrayType& edgeFar, + const IdArrayType& edgeNear, + bool joinGraph, + vtkm::cont::Token& token) : IsJoinGraph(joinGraph) { // constructor - EdgeFarPortal = edgeFar.PrepareForInput(DeviceAdapter()); - EdgeNearPortal = edgeNear.PrepareForInput(DeviceAdapter()); + this->EdgeFarPortal = edgeFar.PrepareForInput(DeviceAdapter(), token); + this->EdgeNearPortal = edgeNear.PrepareForInput(DeviceAdapter(), token); } // constructor // () operator - gets called to do comparison @@ -89,40 +92,40 @@ public: bool operator()(const vtkm::Id& i, const vtkm::Id& j) const { // operator() // start by comparing the indices of the far end - vtkm::Id farIndex1 = EdgeFarPortal.Get(i); - vtkm::Id farIndex2 = EdgeFarPortal.Get(j); + vtkm::Id farIndex1 = this->EdgeFarPortal.Get(i); + vtkm::Id farIndex2 = this->EdgeFarPortal.Get(j); // first compare the far end if (farIndex1 < farIndex2) { - return true ^ IsJoinGraph; + return true ^ this->IsJoinGraph; } if (farIndex2 < farIndex1) { - return false ^ IsJoinGraph; + return false ^ this->IsJoinGraph; } // then compare the indices of the near end (which are guaranteed to be sorted!) - vtkm::Id nearIndex1 = EdgeNearPortal.Get(i); - vtkm::Id nearIndex2 = EdgeNearPortal.Get(j); + vtkm::Id nearIndex1 = this->EdgeNearPortal.Get(i); + vtkm::Id nearIndex2 = this->EdgeNearPortal.Get(j); if (nearIndex1 < nearIndex2) { - return true ^ IsJoinGraph; + return true ^ this->IsJoinGraph; } if (nearIndex2 < nearIndex1) { - return false ^ IsJoinGraph; + return false ^ this->IsJoinGraph; } // if the near indices match, compare the edge IDs if (i < j) { - return false ^ IsJoinGraph; + return false ^ this->IsJoinGraph; } if (j < i) { - return true ^ IsJoinGraph; + return true ^ this->IsJoinGraph; } // fallback can happen when multiple paths end at same extremum @@ -149,9 +152,12 @@ public: } template - VTKM_CONT EdgePeakComparatorImpl PrepareForExecution(DeviceAdapter) const + VTKM_CONT EdgePeakComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) const { - return EdgePeakComparatorImpl(this->EdgeFar, this->EdgeNear, this->JoinGraph); + return EdgePeakComparatorImpl( + this->EdgeFar, this->EdgeNear, this->JoinGraph, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/activegraph/HyperArcSuperNodeComparator.h b/vtkm/worklet/contourtree_augmented/activegraph/HyperArcSuperNodeComparator.h index f1c9b624e..7ec8dea11 100644 --- a/vtkm/worklet/contourtree_augmented/activegraph/HyperArcSuperNodeComparator.h +++ b/vtkm/worklet/contourtree_augmented/activegraph/HyperArcSuperNodeComparator.h @@ -134,11 +134,12 @@ public: template VTKM_CONT HyperArcSuperNodeComparatorImpl PrepareForExecution( - DeviceAdapter device) const + DeviceAdapter device, + vtkm::cont::Token& token) const { return HyperArcSuperNodeComparatorImpl( - this->Hyperparents.PrepareForInput(device), - this->SuperID.PrepareForInput(device), + this->Hyperparents.PrepareForInput(device, token), + this->SuperID.PrepareForInput(device, token), this->IsJoinTree); } diff --git a/vtkm/worklet/contourtree_augmented/activegraph/SuperArcNodeComparator.h b/vtkm/worklet/contourtree_augmented/activegraph/SuperArcNodeComparator.h index c48a0ec7f..a9ead909b 100644 --- a/vtkm/worklet/contourtree_augmented/activegraph/SuperArcNodeComparator.h +++ b/vtkm/worklet/contourtree_augmented/activegraph/SuperArcNodeComparator.h @@ -78,10 +78,12 @@ public: // constructor - takes vectors as parameters VTKM_CONT - SuperArcNodeComparatorImpl(const IdArrayType& superparents, bool joinSweep) + SuperArcNodeComparatorImpl(const IdArrayType& superparents, + bool joinSweep, + vtkm::cont::Token& token) : IsJoinSweep(joinSweep) { // constructor - SuperparentsPortal = superparents.PrepareForInput(DeviceAdapter()); + SuperparentsPortal = superparents.PrepareForInput(DeviceAdapter(), token); } // constructor // () operator - gets called to do comparison @@ -94,16 +96,16 @@ public: // now test on that if (superarcI < superarcJ) - return false ^ IsJoinSweep; + return false ^ this->IsJoinSweep; if (superarcJ < superarcI) - return true ^ IsJoinSweep; + return true ^ this->IsJoinSweep; // if that fails, we share the hyperarc, and sort on supernode index // since that's guaranteed to be pre-sorted if (i < j) - return false ^ IsJoinSweep; + return false ^ this->IsJoinSweep; if (j < i) - return true ^ IsJoinSweep; + return true ^ this->IsJoinSweep; // fallback just in case return false; @@ -127,9 +129,11 @@ public: } template - VTKM_CONT SuperArcNodeComparatorImpl PrepareForExecution(DeviceAdapter) const + VTKM_CONT SuperArcNodeComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) const { - return SuperArcNodeComparatorImpl(this->Superparents, this->JoinSweep); + return SuperArcNodeComparatorImpl(this->Superparents, this->JoinSweep, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeNodeComparator.h b/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeNodeComparator.h index b3ab2055c..d4a858cb0 100644 --- a/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeNodeComparator.h +++ b/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeNodeComparator.h @@ -80,10 +80,12 @@ public: // constructor VTKM_CONT - ContourTreeNodeComparatorImpl(const IdArrayType& superparents, const IdArrayType& superarcs) + ContourTreeNodeComparatorImpl(const IdArrayType& superparents, + const IdArrayType& superarcs, + vtkm::cont::Token& token) { - this->SuperparentsPortal = superparents.PrepareForInput(DeviceAdapter()); - this->SuperarcsPortal = superarcs.PrepareForInput(DeviceAdapter()); + this->SuperparentsPortal = superparents.PrepareForInput(DeviceAdapter(), token); + this->SuperarcsPortal = superarcs.PrepareForInput(DeviceAdapter(), token); } // () operator - gets called to do comparison @@ -122,9 +124,11 @@ public: } template - VTKM_CONT ContourTreeNodeComparatorImpl PrepareForExecution(DeviceAdapter) + VTKM_CONT ContourTreeNodeComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) { - return ContourTreeNodeComparatorImpl(this->Superparents, this->Superarcs); + return ContourTreeNodeComparatorImpl(this->Superparents, this->Superarcs, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeSuperNodeComparator.h b/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeSuperNodeComparator.h index b9c317af0..bab80f136 100644 --- a/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeSuperNodeComparator.h +++ b/vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeSuperNodeComparator.h @@ -83,11 +83,12 @@ public: VTKM_CONT ContourTreeSuperNodeComparatorImpl(const IdArrayType& hyperparents, const IdArrayType& supernodes, - const IdArrayType& whenTransferred) + const IdArrayType& whenTransferred, + vtkm::cont::Token& token) { - this->HyperparentsPortal = hyperparents.PrepareForInput(DeviceAdapter()); - this->SupernodesPortal = supernodes.PrepareForInput(DeviceAdapter()); - this->WhenTransferredPortal = whenTransferred.PrepareForInput(DeviceAdapter()); + this->HyperparentsPortal = hyperparents.PrepareForInput(DeviceAdapter(), token); + this->SupernodesPortal = supernodes.PrepareForInput(DeviceAdapter(), token); + this->WhenTransferredPortal = whenTransferred.PrepareForInput(DeviceAdapter(), token); } // () operator - gets called to do comparison @@ -138,10 +139,12 @@ public: } template - VTKM_CONT ContourTreeSuperNodeComparatorImpl PrepareForExecution(DeviceAdapter) + VTKM_CONT ContourTreeSuperNodeComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) { return ContourTreeSuperNodeComparatorImpl( - this->Hyperparents, this->Supernodes, this->WhenTransferred); + this->Hyperparents, this->Supernodes, this->WhenTransferred, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/contourtreemaker/TransferLeafChains_TransferToContourTree.h b/vtkm/worklet/contourtree_augmented/contourtreemaker/TransferLeafChains_TransferToContourTree.h index ce46c4f8f..959636a5a 100644 --- a/vtkm/worklet/contourtree_augmented/contourtreemaker/TransferLeafChains_TransferToContourTree.h +++ b/vtkm/worklet/contourtree_augmented/contourtreemaker/TransferLeafChains_TransferToContourTree.h @@ -108,15 +108,16 @@ public: const IdArrayType& indegree, const IdArrayType& outbound, const IdArrayType& inbound, - const IdArrayType& inwards) + const IdArrayType& inwards, + vtkm::cont::Token& token) : NumIterations(nIterations) , isJoin(IsJoin) { - OutdegreePortal = outdegree.PrepareForInput(DeviceAdapter()); - IndegreePortal = indegree.PrepareForInput(DeviceAdapter()); - OutboundPortal = outbound.PrepareForInput(DeviceAdapter()); - InboundPortal = inbound.PrepareForInput(DeviceAdapter()); - InwardsPortal = inwards.PrepareForInput(DeviceAdapter()); + this->OutdegreePortal = outdegree.PrepareForInput(DeviceAdapter(), token); + this->IndegreePortal = indegree.PrepareForInput(DeviceAdapter(), token); + this->OutboundPortal = outbound.PrepareForInput(DeviceAdapter(), token); + this->InboundPortal = inbound.PrepareForInput(DeviceAdapter(), token); + this->InwardsPortal = inwards.PrepareForInput(DeviceAdapter(), token); } @@ -128,30 +129,33 @@ public: const OutFieldPortalType& contourTreeSuperarcsPortal, const OutFieldPortalType& contourTreeWhenTransferredPortal) const { - if ((OutdegreePortal.Get(superID) == 0) && (IndegreePortal.Get(superID) == 1)) + if ((this->OutdegreePortal.Get(superID) == 0) && (this->IndegreePortal.Get(superID) == 1)) { // a leaf - contourTreeHyperparentsPortal.Set(superID, superID | (isJoin ? 0 : IS_ASCENDING)); + contourTreeHyperparentsPortal.Set(superID, superID | (this->isJoin ? 0 : IS_ASCENDING)); contourTreeHyperarcsPortal.Set( - superID, MaskedIndex(InboundPortal.Get(superID)) | (isJoin ? 0 : IS_ASCENDING)); + superID, MaskedIndex(this->InboundPortal.Get(superID)) | (this->isJoin ? 0 : IS_ASCENDING)); contourTreeSuperarcsPortal.Set( - superID, MaskedIndex(InwardsPortal.Get(superID)) | (isJoin ? 0 : IS_ASCENDING)); + superID, MaskedIndex(this->InwardsPortal.Get(superID)) | (this->isJoin ? 0 : IS_ASCENDING)); contourTreeWhenTransferredPortal.Set(superID, this->NumIterations | IS_HYPERNODE); } // a leaf else { // not a leaf // retrieve the out neighbour - vtkm::Id outNeighbour = MaskedIndex(OutboundPortal.Get(superID)); + vtkm::Id outNeighbour = MaskedIndex(this->OutboundPortal.Get(superID)); // test whether outneighbour is a leaf - if ((OutdegreePortal.Get(outNeighbour) != 0) || (IndegreePortal.Get(outNeighbour) != 1)) + if ((this->OutdegreePortal.Get(outNeighbour) != 0) || + (this->IndegreePortal.Get(outNeighbour) != 1)) { } else { // set superarc, &c. - contourTreeSuperarcsPortal.Set( - superID, MaskedIndex(InwardsPortal.Get(superID)) | (isJoin ? 0 : IS_ASCENDING)); - contourTreeHyperparentsPortal.Set(superID, outNeighbour | (isJoin ? 0 : IS_ASCENDING)); + contourTreeSuperarcsPortal.Set(superID, + MaskedIndex(this->InwardsPortal.Get(superID)) | + (this->isJoin ? 0 : IS_ASCENDING)); + contourTreeHyperparentsPortal.Set(superID, + outNeighbour | (this->isJoin ? 0 : IS_ASCENDING)); contourTreeWhenTransferredPortal.Set(superID, this->NumIterations | IS_SUPERNODE); } } // not a leaf diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem/SimulatedSimplicityComperator.h b/vtkm/worklet/contourtree_augmented/mesh_dem/SimulatedSimplicityComperator.h index c8ddbdb41..b21e8987a 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem/SimulatedSimplicityComperator.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem/SimulatedSimplicityComperator.h @@ -115,10 +115,10 @@ public: template VTKM_CONT SimulatedSimplicityIndexComparatorImpl - PrepareForExecution(DeviceAdapter device) const + PrepareForExecution(DeviceAdapter device, vtkm::cont::Token& token) const { return SimulatedSimplicityIndexComparatorImpl( - this->Values.PrepareForInput(device)); + this->Values.PrepareForInput(device, token)); } private: diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/ContourTreeMesh.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/ContourTreeMesh.h index b55e3e9e2..66bb99f5a 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/ContourTreeMesh.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/ContourTreeMesh.h @@ -123,7 +123,8 @@ public: template contourtree_mesh_inc_ns::MeshStructureContourTreeMesh PrepareForExecution( - DeviceTag) const; + DeviceTag, + vtkm::cont::Token& token) const; ContourTreeMesh() {} @@ -457,10 +458,10 @@ void ContourTreeMesh::SetPrepareForExecutionBehavior(bool getMax) template template contourtree_mesh_inc_ns::MeshStructureContourTreeMesh - ContourTreeMesh::PrepareForExecution(DeviceTag) const +ContourTreeMesh::PrepareForExecution(DeviceTag, vtkm::cont::Token& token) const { return contourtree_mesh_inc_ns::MeshStructureContourTreeMesh( - this->Neighbours, this->FirstNeighbour, this->MaxNeighbours, this->mGetMax); + this->Neighbours, this->FirstNeighbour, this->MaxNeighbours, this->mGetMax, token); } struct NotNoSuchElement @@ -478,12 +479,13 @@ void ContourTreeMesh::MergeWith(ContourTreeMesh& other) other.DebugPrint("OTHER ContourTreeMesh", __FILE__, __LINE__); #endif + vtkm::cont::Token allToken; mesh_dem_contourtree_mesh_inc::CombinedVectorExecObj allGlobalIndicesExecObj( this->GlobalMeshIndex, other.GlobalMeshIndex); - auto allGlobalIndices = allGlobalIndicesExecObj.PrepareForExecution(DeviceTag()); + auto allGlobalIndices = allGlobalIndicesExecObj.PrepareForExecution(DeviceTag(), allToken); mesh_dem_contourtree_mesh_inc::CombinedVectorExecObj allSortedValuesExecObj( this->SortedValues, other.SortedValues); - auto allSortedValues = allSortedValuesExecObj.PrepareForExecution(DeviceTag()); + auto allSortedValues = allSortedValuesExecObj.PrepareForExecution(DeviceTag(), allToken); //auto allGlobalIndices = CombinedVectorthisGlobalMeshIndex, other.GlobalMeshIndex); // Create combined sort order @@ -515,9 +517,10 @@ void ContourTreeMesh::MergeWith(ContourTreeMesh& other) IdArrayType overallSortIndex; overallSortIndex.Allocate(overallSortOrder.GetNumberOfValues()); { + vtkm::cont::Token token; // Functor return 0,1 for each element of a CombinedVector depending on whethern the current value is different from the next mesh_dem_contourtree_mesh_inc::CombinedVectorDifferentFromNext - differentFromNextFunctor(&allGlobalIndices, overallSortOrder); + differentFromNextFunctor(&allGlobalIndices, overallSortOrder, token); auto differentFromNextArr = vtkm::cont::make_ArrayHandleTransform( vtkm::cont::ArrayHandleIndex(overallSortIndex.GetNumberOfValues() - 1), differentFromNextFunctor); diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_2D_Triangulation.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_2D_Triangulation.h index 29ea10d0d..9b3e37a2d 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_2D_Triangulation.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_2D_Triangulation.h @@ -81,7 +81,8 @@ public: void SetPrepareForExecutionBehavior(bool getMax); template - MeshStructureFreudenthal2D PrepareForExecution(DeviceTag) const; + MeshStructureFreudenthal2D PrepareForExecution(DeviceTag, + vtkm::cont::Token& token) const; Mesh_DEM_Triangulation_2D_Freudenthal(vtkm::Id ncols, vtkm::Id nrows); @@ -99,7 +100,7 @@ Mesh_DEM_Triangulation_2D_Freudenthal::Mesh_DEM_Triangulation_2D : Mesh_DEM_Triangulation_2D(ncols, nrows) { - EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( + this->EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( m2d_freudenthal::EdgeBoundaryDetectionMasks, m2d_freudenthal::N_INCIDENT_EDGES); } @@ -114,7 +115,9 @@ void Mesh_DEM_Triangulation_2D_Freudenthal::SetPrepareForExecuti template template MeshStructureFreudenthal2D - Mesh_DEM_Triangulation_2D_Freudenthal::PrepareForExecution(DeviceTag) const +Mesh_DEM_Triangulation_2D_Freudenthal::PrepareForExecution( + DeviceTag, + vtkm::cont::Token& token) const { return MeshStructureFreudenthal2D(this->NumColumns, this->NumRows, @@ -122,7 +125,8 @@ MeshStructureFreudenthal2D this->UseGetMax, this->SortIndices, this->SortOrder, - EdgeBoundaryDetectionMasks); + this->EdgeBoundaryDetectionMasks, + token); } template diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_3D_Triangulation.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_3D_Triangulation.h index 132cede5e..68553dc2d 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_3D_Triangulation.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/Freudenthal_3D_Triangulation.h @@ -86,7 +86,8 @@ public: void SetPrepareForExecutionBehavior(bool getMax); template - MeshStructureFreudenthal3D PrepareForExecution(DeviceTag) const; + MeshStructureFreudenthal3D PrepareForExecution(DeviceTag, + vtkm::cont::Token& token) const; Mesh_DEM_Triangulation_3D_Freudenthal(vtkm::Id ncols, vtkm::Id nrows, vtkm::Id nslices); @@ -106,12 +107,12 @@ Mesh_DEM_Triangulation_3D_Freudenthal::Mesh_DEM_Triangulation_3D { // Initialize the case tables in vtkm - EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( + this->EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( m3d_freudenthal::EdgeBoundaryDetectionMasks, m3d_freudenthal::N_INCIDENT_EDGES); - NeighbourOffsets = vtkm::cont::make_ArrayHandleGroupVec<3>(vtkm::cont::make_ArrayHandle( + this->NeighbourOffsets = vtkm::cont::make_ArrayHandleGroupVec<3>(vtkm::cont::make_ArrayHandle( m3d_freudenthal::NeighbourOffsets, m3d_freudenthal::N_INCIDENT_EDGES * 3)); - LinkComponentCaseTable = vtkm::cont::make_ArrayHandle(m3d_freudenthal::LinkComponentCaseTable, - m3d_freudenthal::LINK_COMPONENT_CASES); + this->LinkComponentCaseTable = vtkm::cont::make_ArrayHandle( + m3d_freudenthal::LinkComponentCaseTable, m3d_freudenthal::LINK_COMPONENT_CASES); } @@ -126,7 +127,9 @@ void Mesh_DEM_Triangulation_3D_Freudenthal::SetPrepareForExecuti template template MeshStructureFreudenthal3D - Mesh_DEM_Triangulation_3D_Freudenthal::PrepareForExecution(DeviceTag) const +Mesh_DEM_Triangulation_3D_Freudenthal::PrepareForExecution( + DeviceTag, + vtkm::cont::Token& token) const { return MeshStructureFreudenthal3D(this->NumColumns, this->NumRows, @@ -135,9 +138,10 @@ MeshStructureFreudenthal3D this->UseGetMax, this->SortIndices, this->SortOrder, - EdgeBoundaryDetectionMasks, - NeighbourOffsets, - LinkComponentCaseTable); + this->EdgeBoundaryDetectionMasks, + this->NeighbourOffsets, + this->LinkComponentCaseTable, + token); } diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MarchingCubes_3D_Triangulation.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MarchingCubes_3D_Triangulation.h index c760ff6c8..6093fb9df 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MarchingCubes_3D_Triangulation.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MarchingCubes_3D_Triangulation.h @@ -90,7 +90,8 @@ public: void SetPrepareForExecutionBehavior(bool getMax); template - MeshStructureMarchingCubes PrepareForExecution(DeviceTag) const; + MeshStructureMarchingCubes PrepareForExecution(DeviceTag, + vtkm::cont::Token& token) const; Mesh_DEM_Triangulation_3D_MarchingCubes(vtkm::Id ncols, vtkm::Id nrows, vtkm::Id nslices); @@ -110,9 +111,9 @@ Mesh_DEM_Triangulation_3D_MarchingCubes::Mesh_DEM_Triangulation_ { // Initialize the case tables in vtkm - EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( + this->EdgeBoundaryDetectionMasks = vtkm::cont::make_ArrayHandle( m3d_marchingcubes::EdgeBoundaryDetectionMasks, m3d_marchingcubes::N_ALL_NEIGHBOURS); - CubeVertexPermutations = vtkm::cont::make_ArrayHandleGroupVec< + this->CubeVertexPermutations = vtkm::cont::make_ArrayHandleGroupVec< m3d_marchingcubes:: CubeVertexPermutations_PermVecLength>( // create 2D array of vectors of lenghts ...PermVecLength vtkm::cont::make_ArrayHandle( @@ -120,7 +121,7 @@ Mesh_DEM_Triangulation_3D_MarchingCubes::Mesh_DEM_Triangulation_ m3d_marchingcubes::CubeVertexPermutations_NumPermutations * m3d_marchingcubes::CubeVertexPermutations_PermVecLength // total number of elements )); - LinkVertexConnectionsSix = vtkm::cont::make_ArrayHandleGroupVec< + this->LinkVertexConnectionsSix = vtkm::cont::make_ArrayHandleGroupVec< m3d_marchingcubes:: VertexConnections_VecLength>( // create 2D array of vectors o lenght ...VecLength vtkm::cont::make_ArrayHandle( @@ -128,7 +129,7 @@ Mesh_DEM_Triangulation_3D_MarchingCubes::Mesh_DEM_Triangulation_ m3d_marchingcubes::LinkVertexConnectionsSix_NumPairs * m3d_marchingcubes::VertexConnections_VecLength // total number of elements )); - LinkVertexConnectionsEighteen = vtkm::cont::make_ArrayHandleGroupVec< + this->LinkVertexConnectionsEighteen = vtkm::cont::make_ArrayHandleGroupVec< m3d_marchingcubes:: VertexConnections_VecLength>( // create 2D array of vectors o lenght ...VecLength vtkm::cont::make_ArrayHandle( @@ -136,9 +137,9 @@ Mesh_DEM_Triangulation_3D_MarchingCubes::Mesh_DEM_Triangulation_ m3d_marchingcubes::LinkVertexConnectionsEighteen_NumPairs * m3d_marchingcubes::VertexConnections_VecLength // total number of elements )); - InCubeConnectionsSix = vtkm::cont::make_ArrayHandle( + this->InCubeConnectionsSix = vtkm::cont::make_ArrayHandle( m3d_marchingcubes::InCubeConnectionsSix, m3d_marchingcubes::InCubeConnectionsSix_NumElements); - InCubeConnectionsEighteen = + this->InCubeConnectionsEighteen = vtkm::cont::make_ArrayHandle(m3d_marchingcubes::InCubeConnectionsEighteen, m3d_marchingcubes::InCubeConnectionsEighteen_NumElements); } @@ -154,7 +155,9 @@ void Mesh_DEM_Triangulation_3D_MarchingCubes::SetPrepareForExecu template template MeshStructureMarchingCubes - Mesh_DEM_Triangulation_3D_MarchingCubes::PrepareForExecution(DeviceTag) const +Mesh_DEM_Triangulation_3D_MarchingCubes::PrepareForExecution( + DeviceTag, + vtkm::cont::Token& token) const { return MeshStructureMarchingCubes(this->NumColumns, this->NumRows, @@ -162,12 +165,13 @@ MeshStructureMarchingCubes this->UseGetMax, this->SortIndices, this->SortOrder, - EdgeBoundaryDetectionMasks, - CubeVertexPermutations, - LinkVertexConnectionsSix, - LinkVertexConnectionsEighteen, - InCubeConnectionsSix, - InCubeConnectionsEighteen); + this->EdgeBoundaryDetectionMasks, + this->CubeVertexPermutations, + this->LinkVertexConnectionsSix, + this->LinkVertexConnectionsEighteen, + this->InCubeConnectionsSix, + this->InCubeConnectionsEighteen, + token); } template diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshBoundary.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshBoundary.h index 897b24427..011fecd7e 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshBoundary.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshBoundary.h @@ -89,21 +89,24 @@ public: } VTKM_CONT - MeshBoundary2D(vtkm::Id nrows, vtkm::Id ncols, const IdArrayType& sortOrder) + MeshBoundary2D(vtkm::Id nrows, + vtkm::Id ncols, + const IdArrayType& sortOrder, + vtkm::cont::Token& token) : MeshStructure(mesh_dem::MeshStructure2D(nrows, ncols)) { - SortOrderPortal = sortOrder.PrepareForInput(DeviceTag()); + this->SortOrderPortal = sortOrder.PrepareForInput(DeviceTag(), token); } VTKM_EXEC_CONT bool liesOnBoundary(const vtkm::Id index) const { vtkm::Id meshSortOrderValue = this->SortOrderPortal.Get(index); - const vtkm::Id row = MeshStructure.VertexRow(meshSortOrderValue); - const vtkm::Id col = MeshStructure.VertexColumn(meshSortOrderValue); + const vtkm::Id row = this->MeshStructure.VertexRow(meshSortOrderValue); + const vtkm::Id col = this->MeshStructure.VertexColumn(meshSortOrderValue); - return (row == 0) || (col == 0) || (row == MeshStructure.NumRows - 1) || - (col == MeshStructure.NumColumns - 1); + return (row == 0) || (col == 0) || (row == this->MeshStructure.NumRows - 1) || + (col == this->MeshStructure.NumColumns - 1); } private: @@ -125,9 +128,9 @@ public: VTKM_CONT template - MeshBoundary2D PrepareForExecution(DeviceTag) const + MeshBoundary2D PrepareForExecution(DeviceTag, vtkm::cont::Token& token) const { - return MeshBoundary2D(this->NumRows, this->NumColumns, this->SortOrder); + return MeshBoundary2D(this->NumRows, this->NumColumns, this->SortOrder, token); } private: @@ -152,21 +155,25 @@ public: } VTKM_CONT - MeshBoundary3D(vtkm::Id nrows, vtkm::Id ncols, vtkm::Id nslices, const IdArrayType& sortOrder) + MeshBoundary3D(vtkm::Id nrows, + vtkm::Id ncols, + vtkm::Id nslices, + const IdArrayType& sortOrder, + vtkm::cont::Token& token) : MeshStructure(mesh_dem::MeshStructure3D(nrows, ncols, nslices)) { - SortOrderPortal = sortOrder.PrepareForInput(DeviceTag()); + this->SortOrderPortal = sortOrder.PrepareForInput(DeviceTag(), token); } VTKM_EXEC_CONT bool liesOnBoundary(const vtkm::Id index) const { vtkm::Id meshSortOrderValue = this->SortOrderPortal.Get(index); - const vtkm::Id row = MeshStructure.VertexRow(meshSortOrderValue); - const vtkm::Id col = MeshStructure.VertexColumn(meshSortOrderValue); - const vtkm::Id sli = MeshStructure.VertexSlice(meshSortOrderValue); - return (row == 0) || (col == 0) || (sli == 0) || (row == MeshStructure.NumRows - 1) || - (col == MeshStructure.NumColumns - 1) || (sli == MeshStructure.NumSlices - 1); + const vtkm::Id row = this->MeshStructure.VertexRow(meshSortOrderValue); + const vtkm::Id col = this->MeshStructure.VertexColumn(meshSortOrderValue); + const vtkm::Id sli = this->MeshStructure.VertexSlice(meshSortOrderValue); + return (row == 0) || (col == 0) || (sli == 0) || (row == this->MeshStructure.NumRows - 1) || + (col == this->MeshStructure.NumColumns - 1) || (sli == this->MeshStructure.NumSlices - 1); } protected: @@ -193,10 +200,10 @@ public: VTKM_CONT template - MeshBoundary3D PrepareForExecution(DeviceTag) const + MeshBoundary3D PrepareForExecution(DeviceTag, vtkm::cont::Token& token) const { return MeshBoundary3D( - this->NumRows, this->NumColumns, this->NumSlices, this->SortOrder); + this->NumRows, this->NumColumns, this->NumSlices, this->SortOrder, token); } protected: @@ -222,20 +229,21 @@ public: vtkm::Id totalNRows, vtkm::Id totalNCols, vtkm::Id3 minIdx, - vtkm::Id3 maxIdx) + vtkm::Id3 maxIdx, + vtkm::cont::Token& token) : TotalNRows(totalNRows) , TotalNCols(totalNCols) , MinIdx(minIdx) , MaxIdx(maxIdx) { - assert(TotalNRows > 0 && TotalNCols > 0); - this->GlobalMeshIndexPortal = globalMeshIndex.PrepareForInput(DeviceTag()); + assert(this->TotalNRows > 0 && this->TotalNCols > 0); + this->GlobalMeshIndexPortal = globalMeshIndex.PrepareForInput(DeviceTag(), token); } VTKM_EXEC_CONT bool liesOnBoundary(const vtkm::Id index) const { - vtkm::Id idx = GlobalMeshIndexPortal.Get(index); + vtkm::Id idx = this->GlobalMeshIndexPortal.Get(index); vtkm::Id3 rcs; rcs[0] = vtkm::Id((idx % (this->TotalNRows * this->TotalNCols)) / this->TotalNCols); rcs[1] = vtkm::Id(idx % this->TotalNCols); @@ -280,10 +288,11 @@ public: VTKM_CONT template - MeshBoundaryContourTreeMesh PrepareForExecution(DeviceTag) const + MeshBoundaryContourTreeMesh PrepareForExecution(DeviceTag, + vtkm::cont::Token& token) const { return MeshBoundaryContourTreeMesh( - GlobalMeshIndex, this->TotalNRows, this->TotalNCols, this->MinIdx, this->MaxIdx); + this->GlobalMeshIndex, this->TotalNRows, this->TotalNCols, this->MinIdx, this->MaxIdx, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureContourTreeMesh.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureContourTreeMesh.h index cbad2cdfe..d8089be15 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureContourTreeMesh.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureContourTreeMesh.h @@ -102,12 +102,13 @@ public: MeshStructureContourTreeMesh(const cpp2_ns::IdArrayType neighbours, const cpp2_ns::IdArrayType firstNeighbour, const vtkm::Id maxneighbours, - bool getmax) + bool getmax, + vtkm::cont::Token& token) : MaxNeighbours(maxneighbours) , GetMax(getmax) { - NeighboursPortal = neighbours.PrepareForInput(DeviceAdapter()); - FirstNeighbourPortal = firstNeighbour.PrepareForInput(DeviceAdapter()); + this->NeighboursPortal = neighbours.PrepareForInput(DeviceAdapter(), token); + this->FirstNeighbourPortal = firstNeighbour.PrepareForInput(DeviceAdapter(), token); } VTKM_EXEC diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal2D.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal2D.h index 977f8836a..786799056 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal2D.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal2D.h @@ -96,15 +96,16 @@ public: bool getmax, const IdArrayType& sortIndices, const IdArrayType& SortOrder, - const m2d_freudenthal::EdgeBoundaryDetectionMasksType& EdgeBoundaryDetectionMasksIn) + const m2d_freudenthal::EdgeBoundaryDetectionMasksType& EdgeBoundaryDetectionMasksIn, + vtkm::cont::Token& token) : mesh_dem::MeshStructure2D(ncols, nrows) , GetMax(getmax) , NumIncidentEdges(nincident_edges) { - SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter()); - SortOrderPortal = SortOrder.PrepareForInput(DeviceAdapter()); - EdgeBoundaryDetectionMasksPortal = - EdgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter()); + this->SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter(), token); + this->SortOrderPortal = SortOrder.PrepareForInput(DeviceAdapter(), token); + this->EdgeBoundaryDetectionMasksPortal = + EdgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter(), token); } VTKM_EXEC diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal3D.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal3D.h index 21e624008..e95579f87 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal3D.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureFreudenthal3D.h @@ -109,17 +109,19 @@ public: const IdArrayType& sortOrder, const m3d_freudenthal::EdgeBoundaryDetectionMasksType& edgeBoundaryDetectionMasksIn, const m3d_freudenthal::NeighbourOffsetsType& neighbourOffsetsIn, - const m3d_freudenthal::LinkComponentCaseTableType& linkComponentCaseTableIn) + const m3d_freudenthal::LinkComponentCaseTableType& linkComponentCaseTableIn, + vtkm::cont::Token& token) : mesh_dem::MeshStructure3D(ncols, nrows, nslices) , GetMax(getmax) , NumIncidentEdge(nincident_edges) { - SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter()); - SortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter()); - EdgeBoundaryDetectionMasksPortal = - edgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter()); - NeighbourOffsetsPortal = neighbourOffsetsIn.PrepareForInput(DeviceAdapter()); - LinkComponentCaseTablePortal = linkComponentCaseTableIn.PrepareForInput(DeviceAdapter()); + this->SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter(), token); + this->SortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter(), token); + this->EdgeBoundaryDetectionMasksPortal = + edgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter(), token); + this->NeighbourOffsetsPortal = neighbourOffsetsIn.PrepareForInput(DeviceAdapter(), token); + this->LinkComponentCaseTablePortal = + linkComponentCaseTableIn.PrepareForInput(DeviceAdapter(), token); } VTKM_EXEC diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureMarchingCubes.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureMarchingCubes.h index 120e5bcaf..acb0cc621 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureMarchingCubes.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/MeshStructureMarchingCubes.h @@ -118,20 +118,25 @@ public: const m3d_marchingcubes::LinkVertexConnectionsType& LinkVertexConnectionsSixIn, const m3d_marchingcubes::LinkVertexConnectionsType& LinkVertexConnectionsEighteenIn, const m3d_marchingcubes::InCubeConnectionsType& InCubeConnectionsSixIn, - const m3d_marchingcubes::InCubeConnectionsType& InCubeConnectionsEighteenIn) + const m3d_marchingcubes::InCubeConnectionsType& InCubeConnectionsEighteenIn, + vtkm::cont::Token& token) : mesh_dem::MeshStructure3D(ncols, nrows, nslices) , GetMax(getmax) { - SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter()); - SortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter()); - EdgeBoundaryDetectionMasksPortal = - EdgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter()); - CubeVertexPermutationsPortal = CubeVertexPermutationsIn.PrepareForInput(DeviceAdapter()); - LinkVertexConnectionsSixPortal = LinkVertexConnectionsSixIn.PrepareForInput(DeviceAdapter()); - LinkVertexConnectionsEighteenPortal = - LinkVertexConnectionsEighteenIn.PrepareForInput(DeviceAdapter()); - InCubeConnectionsSixPortal = InCubeConnectionsSixIn.PrepareForInput(DeviceAdapter()); - InCubeConnectionsEighteenPortal = InCubeConnectionsEighteenIn.PrepareForInput(DeviceAdapter()); + this->SortIndicesPortal = sortIndices.PrepareForInput(DeviceAdapter(), token); + this->SortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter(), token); + this->EdgeBoundaryDetectionMasksPortal = + EdgeBoundaryDetectionMasksIn.PrepareForInput(DeviceAdapter(), token); + this->CubeVertexPermutationsPortal = + CubeVertexPermutationsIn.PrepareForInput(DeviceAdapter(), token); + this->LinkVertexConnectionsSixPortal = + LinkVertexConnectionsSixIn.PrepareForInput(DeviceAdapter(), token); + this->LinkVertexConnectionsEighteenPortal = + LinkVertexConnectionsEighteenIn.PrepareForInput(DeviceAdapter(), token); + this->InCubeConnectionsSixPortal = + InCubeConnectionsSixIn.PrepareForInput(DeviceAdapter(), token); + this->InCubeConnectionsEighteenPortal = + InCubeConnectionsEighteenIn.PrepareForInput(DeviceAdapter(), token); } VTKM_EXEC diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/ArcComparator.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/ArcComparator.h index 90ba4a264..6798c3ee7 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/ArcComparator.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/ArcComparator.h @@ -87,9 +87,9 @@ public: // constructor - takes vectors as parameters VTKM_CONT - ArcComparatorImpl(const IdArrayType& ct_arcs) + ArcComparatorImpl(const IdArrayType& ct_arcs, vtkm::cont::Token& token) { // constructor - this->ArcsPortal = ct_arcs.PrepareForInput(DeviceAdapter()); + this->ArcsPortal = ct_arcs.PrepareForInput(DeviceAdapter(), token); } // constructor // () operator - gets called to do comparison @@ -126,9 +126,10 @@ public: } template - VTKM_CONT ArcComparatorImpl PrepareForExecution(DeviceAdapter) const + VTKM_CONT ArcComparatorImpl PrepareForExecution(DeviceAdapter, + vtkm::cont::Token& token) const { - return ArcComparatorImpl(this->Arcs); + return ArcComparatorImpl(this->Arcs, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVector.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVector.h index dfce1f7a0..8c39fb827 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVector.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVector.h @@ -85,10 +85,11 @@ public: TArrayPortalType; VTKM_CONT CombinedVector(const vtkm::cont::ArrayHandle& ThisVector, - const vtkm::cont::ArrayHandle& OtherVector) + const vtkm::cont::ArrayHandle& OtherVector, + vtkm::cont::Token& token) { - this->ThisVectorPortal = ThisVector.PrepareForInput(DeviceAdapter()); - this->OtherVectorPortal = OtherVector.PrepareForInput(DeviceAdapter()); + this->ThisVectorPortal = ThisVector.PrepareForInput(DeviceAdapter(), token); + this->OtherVectorPortal = OtherVector.PrepareForInput(DeviceAdapter(), token); } // See contourtree_augmented/Types.h for definitions of IsThis() and CV_OTHER_FLAG @@ -123,9 +124,9 @@ public: } template - CombinedVector PrepareForExecution(DeviceTag) const + CombinedVector PrepareForExecution(DeviceTag, vtkm::cont::Token& token) const { - return CombinedVector(this->ThisVector, this->OtherVector); + return CombinedVector(this->ThisVector, this->OtherVector, token); } vtkm::Id GetNumberOfValues() const diff --git a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVectorDifferentFromNext.h b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVectorDifferentFromNext.h index 965e8dc13..3093917a7 100644 --- a/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVectorDifferentFromNext.h +++ b/vtkm/worklet/contourtree_augmented/mesh_dem_meshtypes/contourtreemesh/CombinedVectorDifferentFromNext.h @@ -94,17 +94,18 @@ public: VTKM_CONT CombinedVectorDifferentFromNext(CombinedVector* inDataArray, - const IdArrayType& sortOrder) + const IdArrayType& sortOrder, + vtkm::cont::Token& token) : DataArray(inDataArray) { - OverallSortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter()); + this->OverallSortOrderPortal = sortOrder.PrepareForInput(DeviceAdapter(), token); } VTKM_EXEC_CONT vtkm::Id operator()(vtkm::Id i) const { - vtkm::Id currGlobalIdx = (*DataArray)[OverallSortOrderPortal.Get(i)]; - vtkm::Id nextGlobalIdx = (*DataArray)[OverallSortOrderPortal.Get(i + 1)]; + vtkm::Id currGlobalIdx = (*this->DataArray)[this->OverallSortOrderPortal.Get(i)]; + vtkm::Id nextGlobalIdx = (*this->DataArray)[this->OverallSortOrderPortal.Get(i + 1)]; return (currGlobalIdx != nextGlobalIdx) ? 1 : 0; } diff --git a/vtkm/worklet/contourtree_augmented/processcontourtree/SuperArcVolumetricComparator.h b/vtkm/worklet/contourtree_augmented/processcontourtree/SuperArcVolumetricComparator.h index 7f580f8b4..c576027e2 100644 --- a/vtkm/worklet/contourtree_augmented/processcontourtree/SuperArcVolumetricComparator.h +++ b/vtkm/worklet/contourtree_augmented/processcontourtree/SuperArcVolumetricComparator.h @@ -84,11 +84,12 @@ public: // constructor SuperArcVolumetricComparatorImpl(const IdArrayType& Weight, const EdgePairArray& SuperarcList, - bool PairsAtLowEnd) + bool PairsAtLowEnd, + vtkm::cont::Token& token) : pairsAtLowEnd(PairsAtLowEnd) { // constructor - weightPortal = Weight.PrepareForInput(DeviceAdapter()); - superarcListPortal = SuperarcList.PrepareForInput(DeviceAdapter()); + weightPortal = Weight.PrepareForInput(DeviceAdapter(), token); + superarcListPortal = SuperarcList.PrepareForInput(DeviceAdapter(), token); } // constructor @@ -165,10 +166,12 @@ public: } template - VTKM_CONT SuperArcVolumetricComparatorImpl PrepareForExecution(DeviceAdapter) + VTKM_CONT SuperArcVolumetricComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) { return SuperArcVolumetricComparatorImpl( - this->Weight, this->SuperArcList, this->PairsAtLowEnd); + this->Weight, this->SuperArcList, this->PairsAtLowEnd, token); } private: diff --git a/vtkm/worklet/contourtree_augmented/processcontourtree/SuperNodeBranchComparator.h b/vtkm/worklet/contourtree_augmented/processcontourtree/SuperNodeBranchComparator.h index 01b8842bb..d19c199f8 100644 --- a/vtkm/worklet/contourtree_augmented/processcontourtree/SuperNodeBranchComparator.h +++ b/vtkm/worklet/contourtree_augmented/processcontourtree/SuperNodeBranchComparator.h @@ -78,10 +78,12 @@ public: IdPortalType SupernodesPortal; // constructor - SuperNodeBranchComparatorImpl(const IdArrayType& WhichBranch, const IdArrayType& supernodes) + SuperNodeBranchComparatorImpl(const IdArrayType& whichBranch, + const IdArrayType& supernodes, + vtkm::cont::Token& token) { // constructor - WhichBranchPortal = WhichBranch.PrepareForInput(DeviceAdapter()); - SupernodesPortal = supernodes.PrepareForInput(DeviceAdapter()); + this->WhichBranchPortal = whichBranch.PrepareForInput(DeviceAdapter(), token); + this->SupernodesPortal = supernodes.PrepareForInput(DeviceAdapter(), token); } // constructor // () operator - gets called to do comparison @@ -89,8 +91,8 @@ public: bool operator()(const vtkm::Id& i, const vtkm::Id& j) const { // operator() // retrieve which branch the supernodes are on - vtkm::Id branchI = MaskedIndex(WhichBranchPortal.Get(i)); - vtkm::Id branchJ = MaskedIndex(WhichBranchPortal.Get(j)); + vtkm::Id branchI = MaskedIndex(this->WhichBranchPortal.Get(i)); + vtkm::Id branchJ = MaskedIndex(this->WhichBranchPortal.Get(j)); // and test them if (branchI < branchJ) @@ -123,9 +125,11 @@ public: } template - VTKM_CONT SuperNodeBranchComparatorImpl PrepareForExecution(DeviceAdapter) + VTKM_CONT SuperNodeBranchComparatorImpl PrepareForExecution( + DeviceAdapter, + vtkm::cont::Token& token) { - return SuperNodeBranchComparatorImpl(this->WhichBranch, this->Supernodes); + return SuperNodeBranchComparatorImpl(this->WhichBranch, this->Supernodes, token); } private: diff --git a/vtkm/worklet/internal/DispatcherBase.h b/vtkm/worklet/internal/DispatcherBase.h index 13b85484a..c6eae321a 100644 --- a/vtkm/worklet/internal/DispatcherBase.h +++ b/vtkm/worklet/internal/DispatcherBase.h @@ -757,14 +757,14 @@ private: typename Invocation::ControlInterface, typename Invocation::ExecutionInterface, Invocation::InputDomainIndex, - decltype(outputToInputMap.PrepareForInput(device)), - decltype(visitArray.PrepareForInput(device)), - decltype(threadToOutputMap.PrepareForInput(device)), + decltype(outputToInputMap.PrepareForInput(device, token)), + decltype(visitArray.PrepareForInput(device, token)), + decltype(threadToOutputMap.PrepareForInput(device, token)), DeviceAdapter> changedInvocation(execObjectParameters, - outputToInputMap.PrepareForInput(device), - visitArray.PrepareForInput(device), - threadToOutputMap.PrepareForInput(device)); + outputToInputMap.PrepareForInput(device, token), + visitArray.PrepareForInput(device, token), + threadToOutputMap.PrepareForInput(device, token)); this->InvokeSchedule(changedInvocation, threadRange, device); } diff --git a/vtkm/worklet/internal/TriangulateTables.h b/vtkm/worklet/internal/TriangulateTables.h index ba1621e98..aeb01236b 100644 --- a/vtkm/worklet/internal/TriangulateTables.h +++ b/vtkm/worklet/internal/TriangulateTables.h @@ -90,10 +90,11 @@ public: VTKM_CONT TriangulateTablesExecutionObject(const TriangulateArrayHandle& counts, const TriangulateArrayHandle& offsets, - const TriangulateArrayHandle& indices) - : Counts(counts.PrepareForInput(DeviceAdapter())) - , Offsets(offsets.PrepareForInput(DeviceAdapter())) - , Indices(indices.PrepareForInput(DeviceAdapter())) + const TriangulateArrayHandle& indices, + vtkm::cont::Token& token) + : Counts(counts.PrepareForInput(DeviceAdapter(), token)) + , Offsets(offsets.PrepareForInput(DeviceAdapter(), token)) + , Indices(indices.PrepareForInput(DeviceAdapter(), token)) { } @@ -140,13 +141,16 @@ class TriangulateTablesExecutionObjectFactory : public vtkm::cont::ExecutionObje { public: template - VTKM_CONT TriangulateTablesExecutionObject PrepareForExecution(Device) const + VTKM_CONT TriangulateTablesExecutionObject PrepareForExecution( + Device, + vtkm::cont::Token& token) const { if (BasicImpl) { return TriangulateTablesExecutionObject(); } - return TriangulateTablesExecutionObject(this->Counts, this->Offsets, this->Indices); + return TriangulateTablesExecutionObject( + this->Counts, this->Offsets, this->Indices, token); } VTKM_CONT TriangulateTablesExecutionObjectFactory() @@ -301,10 +305,11 @@ public: VTKM_CONT TetrahedralizeTablesExecutionObject(const TriangulateArrayHandle& counts, const TriangulateArrayHandle& offsets, - const TriangulateArrayHandle& indices) - : Counts(counts.PrepareForInput(DeviceAdapter())) - , Offsets(offsets.PrepareForInput(DeviceAdapter())) - , Indices(indices.PrepareForInput(DeviceAdapter())) + const TriangulateArrayHandle& indices, + vtkm::cont::Token& token) + : Counts(counts.PrepareForInput(DeviceAdapter(), token)) + , Offsets(offsets.PrepareForInput(DeviceAdapter(), token)) + , Indices(indices.PrepareForInput(DeviceAdapter(), token)) { } @@ -336,13 +341,16 @@ class TetrahedralizeTablesExecutionObjectFactory : public vtkm::cont::ExecutionO { public: template - VTKM_CONT TetrahedralizeTablesExecutionObject PrepareForExecution(Device) const + VTKM_CONT TetrahedralizeTablesExecutionObject PrepareForExecution( + Device, + vtkm::cont::Token& token) const { if (BasicImpl) { return TetrahedralizeTablesExecutionObject(); } - return TetrahedralizeTablesExecutionObject(this->Counts, this->Offsets, this->Indices); + return TetrahedralizeTablesExecutionObject( + this->Counts, this->Offsets, this->Indices, token); } VTKM_CONT diff --git a/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx b/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx index 682141fdb..373c57365 100644 --- a/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx +++ b/vtkm/worklet/internal/testing/UnitTestDispatcherBase.cxx @@ -67,7 +67,7 @@ struct TestExecObjectType : vtkm::cont::ExecutionObjectBase f(*this, std::forward(args)...); } template - VTKM_CONT ExecutionObject PrepareForExecution(Device) const + VTKM_CONT ExecutionObject PrepareForExecution(Device, vtkm::cont::Token&) const { ExecutionObject object; object.Value = this->Value; diff --git a/vtkm/worklet/particleadvection/CellInterpolationHelper.h b/vtkm/worklet/particleadvection/CellInterpolationHelper.h index d19b84358..3215679a7 100644 --- a/vtkm/worklet/particleadvection/CellInterpolationHelper.h +++ b/vtkm/worklet/particleadvection/CellInterpolationHelper.h @@ -111,10 +111,11 @@ public: VTKM_CONT SingleCellTypeInterpolationHelper(vtkm::UInt8 cellShape, vtkm::IdComponent pointsPerCell, - const ConnType& connectivity) + const ConnType& connectivity, + vtkm::cont::Token& token) : CellShape(cellShape) , PointsPerCell(pointsPerCell) - , Connectivity(connectivity.PrepareForInput(DeviceAdapter())) + , Connectivity(connectivity.PrepareForInput(DeviceAdapter(), token)) { } @@ -156,10 +157,11 @@ public: VTKM_CONT ExplicitCellInterpolationHelper(const ShapeType& shape, const OffsetType& offset, - const ConnType& connectivity) - : Shape(shape.PrepareForInput(DeviceAdapter())) - , Offset(offset.PrepareForInput(DeviceAdapter())) - , Connectivity(connectivity.PrepareForInput(DeviceAdapter())) + const ConnType& connectivity, + vtkm::cont::Token& token) + : Shape(shape.PrepareForInput(DeviceAdapter(), token)) + , Offset(offset.PrepareForInput(DeviceAdapter(), token)) + , Connectivity(connectivity.PrepareForInput(DeviceAdapter(), token)) { } @@ -199,7 +201,8 @@ public: virtual ~CellInterpolationHelper() = default; VTKM_CONT virtual const vtkm::exec::CellInterpolationHelper* PrepareForExecution( - vtkm::cont::DeviceAdapterId device) const = 0; + vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const = 0; }; class StructuredCellInterpolationHelper : public vtkm::cont::CellInterpolationHelper @@ -237,7 +240,8 @@ public: VTKM_CONT const vtkm::exec::CellInterpolationHelper* PrepareForExecution( - vtkm::cont::DeviceAdapterId deviceId) const override + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const override { auto& tracker = vtkm::cont::GetRuntimeDeviceTracker(); const bool valid = tracker.CanRunOn(deviceId); @@ -250,7 +254,7 @@ public: ExecutionType* execObject = new ExecutionType(this->CellDims, this->PointDims, this->Is3D); this->ExecHandle.Reset(execObject); - return this->ExecHandle.PrepareForExecution(deviceId); + return this->ExecHandle.PrepareForExecution(deviceId, token); } private: @@ -293,11 +297,14 @@ public: template VTKM_CONT bool operator()(DeviceAdapter, const vtkm::cont::SingleCellTypeInterpolationHelper& contInterpolator, - HandleType& execInterpolator) const + HandleType& execInterpolator, + vtkm::cont::Token& token) const { using ExecutionType = vtkm::exec::SingleCellTypeInterpolationHelper; - ExecutionType* execObject = new ExecutionType( - contInterpolator.CellShape, contInterpolator.PointsPerCell, contInterpolator.Connectivity); + ExecutionType* execObject = new ExecutionType(contInterpolator.CellShape, + contInterpolator.PointsPerCell, + contInterpolator.Connectivity, + token); execInterpolator.Reset(execObject); return true; } @@ -305,15 +312,16 @@ public: VTKM_CONT const vtkm::exec::CellInterpolationHelper* PrepareForExecution( - vtkm::cont::DeviceAdapterId deviceId) const override + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const override { - const bool success = - vtkm::cont::TryExecuteOnDevice(deviceId, SingleCellTypeFunctor(), *this, this->ExecHandle); + const bool success = vtkm::cont::TryExecuteOnDevice( + deviceId, SingleCellTypeFunctor(), *this, this->ExecHandle, token); if (!success) { throwFailedRuntimeDeviceTransfer("SingleCellTypeInterpolationHelper", deviceId); } - return this->ExecHandle.PrepareForExecution(deviceId); + return this->ExecHandle.PrepareForExecution(deviceId, token); } private: @@ -350,11 +358,12 @@ public: template VTKM_CONT bool operator()(DeviceAdapter, const vtkm::cont::ExplicitCellInterpolationHelper& contInterpolator, - HandleType& execInterpolator) const + HandleType& execInterpolator, + vtkm::cont::Token& token) const { using ExecutionType = vtkm::exec::ExplicitCellInterpolationHelper; ExecutionType* execObject = new ExecutionType( - contInterpolator.Shape, contInterpolator.Offset, contInterpolator.Connectivity); + contInterpolator.Shape, contInterpolator.Offset, contInterpolator.Connectivity, token); execInterpolator.Reset(execObject); return true; } @@ -362,15 +371,16 @@ public: VTKM_CONT const vtkm::exec::CellInterpolationHelper* PrepareForExecution( - vtkm::cont::DeviceAdapterId deviceId) const override + vtkm::cont::DeviceAdapterId deviceId, + vtkm::cont::Token& token) const override { - const bool success = - vtkm::cont::TryExecuteOnDevice(deviceId, ExplicitCellFunctor(), *this, this->ExecHandle); + const bool success = vtkm::cont::TryExecuteOnDevice( + deviceId, ExplicitCellFunctor(), *this, this->ExecHandle, token); if (!success) { throwFailedRuntimeDeviceTransfer("ExplicitCellInterpolationHelper", deviceId); } - return this->ExecHandle.PrepareForExecution(deviceId); + return this->ExecHandle.PrepareForExecution(deviceId, token); } private: diff --git a/vtkm/worklet/particleadvection/GridEvaluators.h b/vtkm/worklet/particleadvection/GridEvaluators.h index 1bfbea7cf..cc1743a63 100644 --- a/vtkm/worklet/particleadvection/GridEvaluators.h +++ b/vtkm/worklet/particleadvection/GridEvaluators.h @@ -48,12 +48,13 @@ public: ExecutionGridEvaluator(std::shared_ptr locator, std::shared_ptr interpolationHelper, const vtkm::Bounds& bounds, - const FieldArrayType& field) + const FieldArrayType& field, + vtkm::cont::Token& token) : Bounds(bounds) - , Field(field.PrepareForInput(DeviceAdapter())) + , Field(field.PrepareForInput(DeviceAdapter(), token)) { - Locator = locator->PrepareForExecution(DeviceAdapter()); - InterpolationHelper = interpolationHelper->PrepareForExecution(DeviceAdapter()); + Locator = locator->PrepareForExecution(DeviceAdapter(), token); + InterpolationHelper = interpolationHelper->PrepareForExecution(DeviceAdapter(), token); } template @@ -206,10 +207,11 @@ public: template VTKM_CONT ExecutionGridEvaluator PrepareForExecution( - DeviceAdapter) const + DeviceAdapter, + vtkm::cont::Token& token) const { return ExecutionGridEvaluator( - this->Locator, this->InterpolationHelper, this->Bounds, this->Vectors); + this->Locator, this->InterpolationHelper, this->Bounds, this->Vectors, token); } private: diff --git a/vtkm/worklet/particleadvection/Integrators.h b/vtkm/worklet/particleadvection/Integrators.h index e2cf5ac81..f7bc01f23 100644 --- a/vtkm/worklet/particleadvection/Integrators.h +++ b/vtkm/worklet/particleadvection/Integrators.h @@ -75,11 +75,13 @@ public: }; template - VTKM_CONT const ExecObject* PrepareForExecution(Device) const + VTKM_CONT const ExecObject* PrepareForExecution(Device, vtkm::cont::Token& token) const { this->PrepareForExecutionImpl( - Device(), const_cast&>(this->ExecObjectHandle)); - return this->ExecObjectHandle.PrepareForExecution(Device()); + Device(), + const_cast&>(this->ExecObjectHandle), + token); + return this->ExecObjectHandle.PrepareForExecution(Device(), token); } private: @@ -92,7 +94,8 @@ protected: VTKM_CONT virtual void PrepareForExecutionImpl( vtkm::cont::DeviceAdapterId device, - vtkm::cont::VirtualObjectHandle& execObjectHandle) const = 0; + vtkm::cont::VirtualObjectHandle& execObjectHandle, + vtkm::cont::Token& token) const = 0; template class ExecObjectBaseImpl : public ExecObject @@ -231,10 +234,11 @@ struct IntegratorPrepareForExecutionFunctor vtkm::cont::VirtualObjectHandle& execObjectHandle, const EvaluatorType& evaluator, vtkm::FloatDefault stepLength, - vtkm::FloatDefault tolerance) const + vtkm::FloatDefault tolerance, + vtkm::cont::Token& token) const { - IntegratorType* integrator = - new IntegratorType(evaluator.PrepareForExecution(Device()), stepLength, tolerance); + IntegratorType* integrator = new IntegratorType( + evaluator.PrepareForExecution(Device(), token), stepLength, tolerance); execObjectHandle.Reset(integrator); return true; } @@ -258,13 +262,13 @@ public: template class ExecObject : public Integrator::ExecObjectBaseImpl< - decltype(std::declval().PrepareForExecution(Device())), + vtkm::cont::internal::ExecutionObjectType, typename RK4Integrator::template ExecObject> { VTKM_IS_DEVICE_ADAPTER_TAG(Device); using FieldEvaluateExecType = - decltype(std::declval().PrepareForExecution(Device())); + vtkm::cont::internal::ExecutionObjectType; using Superclass = Integrator::ExecObjectBaseImpl>; @@ -320,14 +324,16 @@ private: protected: VTKM_CONT virtual void PrepareForExecutionImpl( vtkm::cont::DeviceAdapterId device, - vtkm::cont::VirtualObjectHandle& execObjectHandle) const override + vtkm::cont::VirtualObjectHandle& execObjectHandle, + vtkm::cont::Token& token) const override { vtkm::cont::TryExecuteOnDevice(device, detail::IntegratorPrepareForExecutionFunctor(), execObjectHandle, this->Evaluator, this->StepLength, - this->Tolerance); + this->Tolerance, + token); } }; @@ -383,14 +389,16 @@ private: protected: VTKM_CONT virtual void PrepareForExecutionImpl( vtkm::cont::DeviceAdapterId device, - vtkm::cont::VirtualObjectHandle& execObjectHandle) const override + vtkm::cont::VirtualObjectHandle& execObjectHandle, + vtkm::cont::Token& token) const override { vtkm::cont::TryExecuteOnDevice(device, detail::IntegratorPrepareForExecutionFunctor(), execObjectHandle, this->Evaluator, this->StepLength, - this->Tolerance); + this->Tolerance, + token); } }; //EulerIntegrator diff --git a/vtkm/worklet/particleadvection/Particles.h b/vtkm/worklet/particleadvection/Particles.h index 9e66c5644..ce2b4ab8d 100644 --- a/vtkm/worklet/particleadvection/Particles.h +++ b/vtkm/worklet/particleadvection/Particles.h @@ -13,6 +13,9 @@ #include #include +#include +#include +#include #include #include #include @@ -34,9 +37,11 @@ public: { } - ParticleExecutionObject(vtkm::cont::ArrayHandle particleArray, vtkm::Id maxSteps) + ParticleExecutionObject(vtkm::cont::ArrayHandle particleArray, + vtkm::Id maxSteps, + vtkm::cont::Token& token) { - Particles = particleArray.PrepareForInPlace(Device()); + Particles = particleArray.PrepareForInPlace(Device(), token); MaxSteps = maxSteps; } @@ -108,10 +113,11 @@ class Particles : public vtkm::cont::ExecutionObjectBase public: template VTKM_CONT vtkm::worklet::particleadvection::ParticleExecutionObject PrepareForExecution( - Device) const + Device, + vtkm::cont::Token& token) const { - return vtkm::worklet::particleadvection::ParticleExecutionObject(this->ParticleArray, - this->MaxSteps); + return vtkm::worklet::particleadvection::ParticleExecutionObject( + this->ParticleArray, this->MaxSteps, token); } VTKM_CONT @@ -147,14 +153,15 @@ public: vtkm::cont::ArrayHandle historyArray, vtkm::cont::ArrayHandle validPointArray, vtkm::cont::ArrayHandle stepCountArray, - vtkm::Id maxSteps) - : ParticleExecutionObject(pArray, maxSteps) + vtkm::Id maxSteps, + vtkm::cont::Token& token) + : ParticleExecutionObject(pArray, maxSteps, token) , Length(maxSteps + 1) { vtkm::Id numPos = pArray.GetNumberOfValues(); - History = historyArray.PrepareForOutput(numPos * Length, Device()); - ValidPoint = validPointArray.PrepareForInPlace(Device()); - StepCount = stepCountArray.PrepareForInPlace(Device()); + History = historyArray.PrepareForOutput(numPos * Length, Device(), token); + ValidPoint = validPointArray.PrepareForInPlace(Device(), token); + StepCount = stepCountArray.PrepareForInPlace(Device(), token); } VTKM_EXEC @@ -212,14 +219,15 @@ public: template VTKM_CONT vtkm::worklet::particleadvection::StateRecordingParticleExecutionObject - PrepareForExecution(Device) const + PrepareForExecution(Device, vtkm::cont::Token& token) const { return vtkm::worklet::particleadvection::StateRecordingParticleExecutionObject( this->ParticleArray, this->HistoryArray, this->ValidPointArray, this->StepCountArray, - this->MaxSteps); + this->MaxSteps, + token); } VTKM_CONT StateRecordingParticles(vtkm::cont::ArrayHandle& pArray, const vtkm::Id& maxSteps) diff --git a/vtkm/worklet/particleadvection/TemporalGridEvaluators.h b/vtkm/worklet/particleadvection/TemporalGridEvaluators.h index 269a1c0d2..c89b4f817 100644 --- a/vtkm/worklet/particleadvection/TemporalGridEvaluators.h +++ b/vtkm/worklet/particleadvection/TemporalGridEvaluators.h @@ -37,9 +37,10 @@ public: ExecutionTemporalGridEvaluator(const GridEvaluator& evaluatorOne, const vtkm::FloatDefault timeOne, const GridEvaluator& evaluatorTwo, - const vtkm::FloatDefault timeTwo) - : EvaluatorOne(evaluatorOne.PrepareForExecution(DeviceAdapter())) - , EvaluatorTwo(evaluatorTwo.PrepareForExecution(DeviceAdapter())) + const vtkm::FloatDefault timeTwo, + vtkm::cont::Token& token) + : EvaluatorOne(evaluatorOne.PrepareForExecution(DeviceAdapter(), token)) + , EvaluatorTwo(evaluatorTwo.PrepareForExecution(DeviceAdapter(), token)) , TimeOne(timeOne) , TimeTwo(timeTwo) , TimeDiff(timeTwo - timeOne) @@ -144,10 +145,11 @@ public: template VTKM_CONT ExecutionTemporalGridEvaluator PrepareForExecution( - DeviceAdapter) const + DeviceAdapter, + vtkm::cont::Token& token) const { return ExecutionTemporalGridEvaluator( - this->EvaluatorOne, this->TimeOne, this->EvaluatorTwo, this->TimeTwo); + this->EvaluatorOne, this->TimeOne, this->EvaluatorTwo, this->TimeTwo, token); } private: diff --git a/vtkm/worklet/testing/UnitTestOrientNormals.cxx b/vtkm/worklet/testing/UnitTestOrientNormals.cxx index d7fbadd72..90ae5cb6c 100644 --- a/vtkm/worklet/testing/UnitTestOrientNormals.cxx +++ b/vtkm/worklet/testing/UnitTestOrientNormals.cxx @@ -90,14 +90,6 @@ struct ValidateNormals using NormalType = vtkm::Vec; using NormalsArrayType = vtkm::cont::ArrayHandleVirtual; using NormalsPortalType = decltype(std::declval().GetPortalConstControl()); - using ConnType = - decltype(std::declval().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, - vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{})); - using RConnType = - decltype(std::declval().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, - vtkm::TopologyElementTagPoint{}, - vtkm::TopologyElementTagCell{})); using PointsType = decltype(std::declval().GetData().GetPortalConstControl()); @@ -105,8 +97,6 @@ struct ValidateNormals CellSetType Cells; PointsType Points; - ConnType Conn; - RConnType RConn; NormalsArrayType PointNormalsArray; NormalsPortalType PointNormals; @@ -169,12 +159,6 @@ struct ValidateNormals vtkm::TopologyElementTagPoint{}); this->Cells.GetConnectivityArray(vtkm::TopologyElementTagCell{}, vtkm::TopologyElementTagPoint{}); - this->Conn = this->Cells.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, - vtkm::TopologyElementTagCell{}, - vtkm::TopologyElementTagPoint{}); - this->RConn = this->Cells.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, - vtkm::TopologyElementTagPoint{}, - vtkm::TopologyElementTagCell{}); if (this->CheckPoints) { @@ -254,6 +238,16 @@ private: queue.emplace_back(startPtIdx, startRefNormal); this->VisitedPoints.SetBit(startPtIdx, true); + vtkm::cont::Token token; + auto connections = this->Cells.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, + vtkm::TopologyElementTagCell{}, + vtkm::TopologyElementTagPoint{}, + token); + auto reverseConnections = this->Cells.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}, + vtkm::TopologyElementTagPoint{}, + vtkm::TopologyElementTagCell{}, + token); + while (!queue.empty()) { const vtkm::Id curPtIdx = queue.back().first; @@ -278,7 +272,7 @@ private: } // Lookup and visit neighbor cells: - const auto neighborCells = this->RConn.GetIndices(curPtIdx); + const auto neighborCells = reverseConnections.GetIndices(curPtIdx); const auto numNeighborCells = neighborCells.GetNumberOfComponents(); for (vtkm::IdComponent nCellIdx = 0; nCellIdx < numNeighborCells; ++nCellIdx) { @@ -307,7 +301,7 @@ private: } // Lookup and visit points in this cell: - const auto neighborPoints = this->Conn.GetIndices(curCellIdx); + const auto neighborPoints = connections.GetIndices(curCellIdx); const auto numNeighborPoints = neighborPoints.GetNumberOfComponents(); for (vtkm::IdComponent nPtIdx = 0; nPtIdx < numNeighborPoints; ++nPtIdx) { diff --git a/vtkm/worklet/testing/UnitTestWorkletMapFieldExecArg.cxx b/vtkm/worklet/testing/UnitTestWorkletMapFieldExecArg.cxx index cd36f3359..2cf5f1501 100644 --- a/vtkm/worklet/testing/UnitTestWorkletMapFieldExecArg.cxx +++ b/vtkm/worklet/testing/UnitTestWorkletMapFieldExecArg.cxx @@ -20,7 +20,7 @@ struct SimpleExecObject : vtkm::cont::ExecutionObjectBase { template - Device PrepareForExecution(Device) const + Device PrepareForExecution(Device, vtkm::cont::Token&) const { return Device(); }