Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-07-30 18:50:59 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Convert execution preparation to use tokens

Browse Source 
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.
This commit is contained in:
Kenneth Moreland 2020-01-21 13:18:03 -07:00
parent 642b59f4fd
commit ad0a53af71
125 changed files with 1698 additions and 1087 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
benchmarking/BenchmarkFieldAlgorithms.cxx
View File

@ -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<const vtkm::Sphere*>(handle.PrepareForExecution(device));
auto function = static_cast<const vtkm::Sphere*>(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<const vtkm::Sphere*>(h1.PrepareForExecution(device));
auto f2 = static_cast<const vtkm::Sphere*>(h2.PrepareForExecution(device));
auto f1 = static_cast<const vtkm::Sphere*>(h1.PrepareForExecution(device, token));
auto f2 = static_cast<const vtkm::Sphere*>(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 };

3
benchmarking/BenchmarkFilters.cxx
View File

@ -549,9 +549,10 @@ struct PrepareForInput
CellSetT& mcellSet = const_cast<CellSetT&>(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();

7
vtkm/cont/Algorithm.h
View File

@ -27,12 +27,11 @@ namespace cont
namespace detail
{
template <typename Device, typename T>
inline auto DoPrepareArgForExec(T&& object, vtkm::cont::Token& token, std::true_type)
-> decltype(std::declval<T>().PrepareForExecution(Device()))
inline auto DoPrepareArgForExec(T&& object, vtkm::cont::Token& token, std::true_type) -> decltype(
vtkm::cont::internal::CallPrepareForExecution(std::forward<T>(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<T>(object), Device{}, token);
}
template <typename Device, typename T>

13
vtkm/cont/ArrayHandle.h
View File

@ -13,6 +13,7 @@
#include <vtkm/cont/vtkm_cont_export.h>
#include <vtkm/Assert.h>
#include <vtkm/Deprecated.h>
#include <vtkm/Flags.h>
#include <vtkm/Types.h>
@ -495,24 +496,24 @@ public:
DeviceAdapterTag,
vtkm::cont::Token& token);
// TODO: Deprecate these
template <typename DeviceAdapterTag>
VTKM_CONT
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInput now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::PortalConst PrepareForInput(DeviceAdapterTag) const
{
vtkm::cont::Token token;
return this->PrepareForInput(DeviceAdapterTag{}, token);
}
template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForOutput(
vtkm::Id numberOfValues,
DeviceAdapterTag)
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForOutput now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::Portal
PrepareForOutput(vtkm::Id numberOfValues, DeviceAdapterTag)
{
vtkm::cont::Token token;
return this->PrepareForOutput(numberOfValues, DeviceAdapterTag{}, token);
}
template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForInPlace(DeviceAdapterTag)
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInPlace now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForInPlace(DeviceAdapterTag)
{
vtkm::cont::Token token;
return this->PrepareForInPlace(DeviceAdapterTag{}, token);

8
vtkm/cont/ArrayHandleCartesianProduct.h
View File

@ -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

11
vtkm/cont/CellLocator.h
View File

@ -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; }

21
vtkm/cont/CellLocatorBoundingIntervalHierarchy.cxx
View File

@ -453,6 +453,7 @@ struct CellLocatorBIHPrepareForExecutionFunctor
bool operator()(
DeviceAdapter,
const CellSetType& cellset,
vtkm::cont::Token& token,
vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator>& bihExec,
const vtkm::cont::ArrayHandle<vtkm::exec::CellLocatorBoundingIntervalHierarchyNode>& nodes,
const vtkm::cont::ArrayHandle<vtkm::Id>& processedCellIds,
@ -461,7 +462,7 @@ struct CellLocatorBIHPrepareForExecutionFunctor
using ExecutionType =
vtkm::exec::CellLocatorBoundingIntervalHierarchyExec<DeviceAdapter, CellSetType>;
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 <typename CellSet, typename... Args>
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>(args)...);
const bool success = vtkm::cont::TryExecuteOnDevice(device,
CellLocatorBIHPrepareForExecutionFunctor(),
cellset,
token,
std::forward<Args>(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);
;
}

4
vtkm/cont/CellLocatorBoundingIntervalHierarchy.h
View File

@ -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

5
vtkm/cont/CellLocatorGeneral.cxx
View File

@ -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;
}

3
vtkm/cont/CellLocatorGeneral.h
View File

@ -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;

11
vtkm/cont/CellLocatorRectilinearGrid.cxx
View File

@ -68,10 +68,12 @@ struct CellLocatorRectilinearGridPrepareForExecutionFunctor
template <typename DeviceAdapter, typename... Args>
VTKM_CONT bool operator()(DeviceAdapter,
vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator>& execLocator,
vtkm::cont::Token& token,
Args&&... args) const
{
using ExecutionType = vtkm::exec::CellLocatorRectilinearGrid<DeviceAdapter, dimensions>;
ExecutionType* execObject = new ExecutionType(std::forward<Args>(args)..., DeviceAdapter());
ExecutionType* execObject =
new ExecutionType(std::forward<Args>(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<Structured3DType>(),
@ -99,6 +103,7 @@ const vtkm::exec::CellLocator* CellLocatorRectilinearGrid::PrepareForExecution(
device,
CellLocatorRectilinearGridPrepareForExecutionFunctor<2>(),
this->ExecutionObjectHandle,
token,
this->PlaneSize,
this->RowSize,
this->GetCellSet().template Cast<Structured2DType>(),
@ -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

3
vtkm/cont/CellLocatorRectilinearGrid.h
View File

@ -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;

17
vtkm/cont/CellLocatorUniformBins.cxx
View File

@ -457,6 +457,7 @@ struct CellLocatorUniformBins::MakeExecObject
template <typename CellSetType, typename DeviceAdapter>
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 <typename DeviceAdapter>
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);
}
//----------------------------------------------------------------------------

30
vtkm/cont/CellLocatorUniformBins.h
View File

@ -94,13 +94,15 @@ private:
using ArrayPortalConst =
typename vtkm::cont::ArrayHandle<T>::template ExecutionTypes<DeviceAdapter>::PortalConst;
using CoordsPortalType =
decltype(vtkm::cont::ArrayHandleVirtualCoordinates{}.PrepareForInput(DeviceAdapter{}));
using CoordsPortalType = decltype(vtkm::cont::ArrayHandleVirtualCoordinates{}.PrepareForInput(
DeviceAdapter{},
std::declval<vtkm::cont::Token&>()));
using CellSetP2CExecType =
decltype(std::declval<CellSetType>().PrepareForInput(DeviceAdapter{},
vtkm::TopologyElementTagCell{},
vtkm::TopologyElementTagPoint{}));
vtkm::TopologyElementTagPoint{},
std::declval<vtkm::cont::Token&>()));
// 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<vtkm::Id>& cellCount,
const vtkm::cont::ArrayHandle<vtkm::Id>& 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;

11
vtkm/cont/CellLocatorUniformGrid.cxx
View File

@ -80,10 +80,12 @@ struct CellLocatorUniformGridPrepareForExecutionFunctor
template <typename DeviceAdapter, typename... Args>
VTKM_CONT bool operator()(DeviceAdapter,
vtkm::cont::VirtualObjectHandle<vtkm::exec::CellLocator>& execLocator,
vtkm::cont::Token& token,
Args&&... args) const
{
using ExecutionType = vtkm::exec::CellLocatorUniformGrid<DeviceAdapter, dimensions>;
ExecutionType* execObject = new ExecutionType(std::forward<Args>(args)..., DeviceAdapter());
ExecutionType* execObject =
new ExecutionType(std::forward<Args>(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

3
vtkm/cont/CellLocatorUniformGrid.h
View File

@ -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;

4
vtkm/cont/CellSetExtrude.cxx
View File

@ -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)
{

4
vtkm/cont/ColorTable.h
View File

@ -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;

32
vtkm/cont/ColorTable.hxx
View File

@ -48,13 +48,15 @@ struct transfer_color_table_to_device
{
template <typename DeviceAdapter>
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 <typename DeviceAdapter, typename ColorTable, typename... Args>
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>(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);
}
}
}

71
vtkm/cont/ExecutionObjectBase.h
View File

@ -52,9 +52,11 @@ struct CheckPrepareForExecution
struct CheckPrepareForExecutionDeprecated
{
VTKM_DEPRECATED_SUPPRESS_BEGIN
template <typename T>
static auto check(T* p)
-> decltype(p->PrepareForExecution(vtkm::cont::DeviceAdapterTagSerial{}), std::true_type());
VTKM_DEPRECATED_SUPPRESS_END
template <typename T>
static auto check(...) -> std::false_type;
@ -64,7 +66,7 @@ struct CheckPrepareForExecutionDeprecated
template <typename T>
using IsExecutionObjectBase =
std::is_base_of<vtkm::cont::ExecutionObjectBase, typename std::decay<T>::type>;
typename std::is_base_of<vtkm::cont::ExecutionObjectBase, typename std::decay<T>::type>::type;
template <typename T>
struct HasPrepareForExecution
@ -88,37 +90,6 @@ struct HasPrepareForExecutionDeprecated
::vtkm::cont::internal::HasPrepareForExecutionDeprecated<execObject>::value, \
"Provided type does not have requisite PrepareForExecution method.")
namespace detail
{
template <typename T, typename Device>
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 <typename T, typename Device>
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 <typename T, typename Device>
VTKM_CONT auto CallPrepareForExecution(T&& execObject, Device device, vtkm::cont::Token& token)
-> decltype(detail::CallPrepareForExecutionImpl(std::forward<T>(execObject),
device,
token,
HasPrepareForExecution<T>{},
HasPrepareForExecutionDeprecated<T>{}))
-> decltype(execObject.PrepareForExecution(device, token))
{
VTKM_IS_EXECUTION_OBJECT(T);
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
return detail::CallPrepareForExecutionImpl(std::forward<T>(execObject),
device,
token,
HasPrepareForExecution<T>{},
HasPrepareForExecutionDeprecated<T>{});
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 <typename T, typename Device, typename TokenType>
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<vtkm::cont::Token, typename std::decay<TokenType>::type>::value));
return execObject.PrepareForExecution(device);
}
/// \brief Gets the type of the execution-side object for an ExecutionObject.

5
vtkm/cont/ImplicitFunctionHandle.h
View File

@ -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

19
vtkm/cont/PointLocator.h
View File

@ -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::exec::PointLocator>;
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;

18
vtkm/cont/PointLocatorUniformGrid.cxx
View File

@ -101,7 +101,8 @@ struct PointLocatorUniformGrid::PrepareExecutionObjectFunctor
template <typename DeviceAdapter>
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<vtkm::FloatDefault>(self.Range[0].Min),
static_cast<vtkm::FloatDefault>(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);

3
vtkm/cont/PointLocatorUniformGrid.h
View File

@ -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
{

3
vtkm/cont/StorageExtrude.h
View File

@ -466,7 +466,8 @@ template <typename T, typename Device>
class VTKM_ALWAYS_EXPORT ArrayTransfer<T, internal::StorageTagExtrude, Device>
{
using BaseT = typename VecTraits<T>::BaseComponentType;
using TPortalType = decltype(vtkm::cont::ArrayHandle<BaseT>{}.PrepareForInput(Device{}));
using TPortalType = decltype(
vtkm::cont::ArrayHandle<BaseT>{}.PrepareForInput(Device{}, std::declval<vtkm::cont::Token&>()));
public:
using ValueType = T;

9
vtkm/cont/StorageVirtual.hxx
View File

@ -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<decltype(portal)>;
vtkm::cont::detail::TransferToDevice<DerivedPortal> 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<decltype(portal)>;
vtkm::cont::detail::TransferToDevice<DerivedPortal> 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<decltype(portal)>;
vtkm::cont::detail::TransferToDevice<DerivedPortal> transfer;
return transfer(device, payload, portal);

14
vtkm/cont/VirtualObjectHandle.h
View File

@ -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<const VirtualBaseType*>(this->Internals->PrepareForExecution(deviceId));
return static_cast<const VirtualBaseType*>(
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

3
vtkm/cont/arg/TransportTagArrayIn.h
View File

@ -37,7 +37,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayIn, ContObjectType, Device>
{
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInput(Device()));
using ExecObjectType = decltype(
std::declval<ContObjectType>().PrepareForInput(Device(), std::declval<vtkm::cont::Token&>()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object,

3
vtkm/cont/arg/TransportTagArrayInOut.h
View File

@ -40,7 +40,8 @@ struct Transport<vtkm::cont::arg::TransportTagArrayInOut, ContObjectType, Device
// is not an array handle as an argument that is expected to be one.
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInPlace(Device()));
using ExecObjectType = decltype(
std::declval<ContObjectType>().PrepareForInPlace(Device(), std::declval<vtkm::cont::Token&>()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType& object,

4
vtkm/cont/arg/TransportTagArrayOut.h
View File

@ -40,7 +40,9 @@ struct Transport<vtkm::cont::arg::TransportTagArrayOut, ContObjectType, Device>
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType =
decltype(std::declval<ContObjectType>().PrepareForOutput(vtkm::Id{}, Device()));
decltype(std::declval<ContObjectType>().PrepareForOutput(vtkm::Id{},
Device{},
std::declval<vtkm::cont::Token&>()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(ContObjectType& object,

7
vtkm/cont/arg/TransportTagCellSetIn.h
View File

@ -43,8 +43,11 @@ struct Transport<vtkm::cont::arg::TransportTagCellSetIn<VisitTopology, IncidentT
{
VTKM_IS_CELL_SET(ContObjectType);
using ExecObjectType = decltype(
std::declval<ContObjectType>().PrepareForInput(Device(), VisitTopology(), IncidentTopology()));
using ExecObjectType =
decltype(std::declval<ContObjectType>().PrepareForInput(Device(),
VisitTopology(),
IncidentTopology(),
std::declval<vtkm::cont::Token&>()));
template <typename InputDomainType>
VTKM_CONT ExecObjectType operator()(const ContObjectType& object,

3
vtkm/cont/arg/TransportTagTopologyFieldIn.h
View File

@ -78,7 +78,8 @@ struct Transport<vtkm::cont::arg::TransportTagTopologyFieldIn<TopologyElementTag
VTKM_IS_ARRAY_HANDLE(ContObjectType);
using ExecObjectType = decltype(std::declval<ContObjectType>().PrepareForInput(Device()));
using ExecObjectType = decltype(
std::declval<ContObjectType>().PrepareForInput(Device(), std::declval<vtkm::cont::Token&>()));
VTKM_CONT
ExecObjectType operator()(const ContObjectType& object,

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

@ -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<vtkm::UInt64>(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 <typename T, typename U, class SIn, class SStencil, class SOut>
@ -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 <typename T, class SIn, class SVal, class SOut, class BinaryCompare>
@ -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 <typename T, typename U, class SIn>
@ -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 <typename T, typename U, class SIn, class BinaryFunctor>
@ -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 <typename T,
@ -1335,12 +1347,13 @@ public:
{
return;
}
vtkm::Id reduced_size =
ReduceByKeyPortal(keys.PrepareForInput(DeviceAdapterTagCuda()),
values.PrepareForInput(DeviceAdapterTagCuda()),
keys_output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()),
values_output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda()),
binary_functor);
vtkm::cont::Token token;
vtkm::Id reduced_size = ReduceByKeyPortal(
keys.PrepareForInput(DeviceAdapterTagCuda(), token),
values.PrepareForInput(DeviceAdapterTagCuda(), token),
keys_output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token),
values_output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token),
binary_functor);
keys_output.Shrink(reduced_size);
values_output.Shrink(reduced_size);
@ -1355,7 +1368,7 @@ public:
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTagCuda());
output.Allocate(0);
return vtkm::TypeTraits<T>::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 <typename T, class SIn, class SOut, class BinaryFunctor>
@ -1379,7 +1393,7 @@ public:
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTagCuda());
output.Allocate(0);
return vtkm::TypeTraits<T>::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 <typename T, class SIn, class SOut>
@ -1403,7 +1419,7 @@ public:
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTagCuda());
output.Allocate(0);
return vtkm::TypeTraits<T>::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 <typename T, class SIn, class SOut, class BinaryFunctor>
@ -1426,7 +1443,7 @@ public:
const vtkm::Id numberOfValues = input.GetNumberOfValues();
if (numberOfValues <= 0)
{
output.PrepareForOutput(0, DeviceAdapterTagCuda());
output.Allocate(0);
return vtkm::TypeTraits<T>::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 <typename T, typename U, typename KIn, typename VIn, typename VOut>
@ -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 <typename T,
@ -1478,18 +1502,20 @@ 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()),
output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token),
::thrust::equal_to<T>(),
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<T>::ZeroInitialization(),
::thrust::equal_to<T>(),
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<T>(),
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 <typename T, class Storage, class BinaryCompare>
@ -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 <typename T, typename U, class StorageT, class StorageU>
@ -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 <typename T, typename U, class StorageT, class StorageU, class BinaryCompare>
@ -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 <typename T, class SIn, class SVal, class SOut, class BinaryCompare>
@ -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()

16
vtkm/cont/internal/ArrayHandleBasicImpl.h
View File

@ -403,25 +403,25 @@ public:
DeviceAdapterTag device,
vtkm::cont::Token& token);
// TODO: Deprecate these
template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::PortalConst PrepareForInput(
DeviceAdapterTag device) const
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInput now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::PortalConst
PrepareForInput(DeviceAdapterTag device) const
{
vtkm::cont::Token token;
return this->PrepareForInput(device, token);
}
template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForOutput(
vtkm::Id numVals,
DeviceAdapterTag device)
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForOutput now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::Portal
PrepareForOutput(vtkm::Id numVals, DeviceAdapterTag device)
{
vtkm::cont::Token token;
return this->PrepareForOutput(numVals, device, token);
}
template <typename DeviceAdapterTag>
VTKM_CONT typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForInPlace(
DeviceAdapterTag device)
VTKM_CONT VTKM_DEPRECATED(1.6, "PrepareForInPlace now requires a vtkm::cont::Token object.")
typename ExecutionTypes<DeviceAdapterTag>::Portal PrepareForInPlace(DeviceAdapterTag device)
{
vtkm::cont::Token token;
return this->PrepareForInPlace(device, token);

13
vtkm/cont/internal/ConnectivityExplicitInternals.h
View File

@ -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<decltype(offInPortal)> cellIdCalc{ offInPortal };
PassThrough idxCalc{};
ConnIdxToCellIdCalc<decltype(offInPortal)> 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<vtkm::UInt8>(CELL_SHAPE_VERTEX), numberOfPoints);

195
vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
View File

@ -101,9 +101,11 @@ private:
{
using OutputArrayType = vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>;
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<decltype(inputPortal), decltype(outputPortal)> 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<vtkm::UInt64> 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<vtkm::Id>(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<decltype(inputPortal), decltype(outputPortal)> kernel(inputPortal, outputPortal);
DerivedAlgorithm::Schedule(kernel, inSize);
@ -175,26 +183,36 @@ public:
using IndexArrayType = vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic>;
IndexArrayType indices;
auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag());
auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag());
{
vtkm::cont::Token token;
StencilToIndexFlagKernel<decltype(stencilPortal), decltype(indexPortal), UnaryPredicate>
indexKernel(stencilPortal, indexPortal, unary_predicate);
auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag(), token);
auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag(), token);
DerivedAlgorithm::Schedule(indexKernel, arrayLength);
StencilToIndexFlagKernel<decltype(stencilPortal), decltype(indexPortal), UnaryPredicate>
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<decltype(inputPortal),
decltype(stencilPortal),
decltype(indexPortal),
decltype(outputPortal),
UnaryPredicate>
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<decltype(inputPortal),
decltype(stencilPortal),
decltype(indexPortal),
decltype(outputPortal),
UnaryPredicate>
copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate);
DerivedAlgorithm::Schedule(copyKernel, arrayLength);
}
}
template <typename T, typename U, class CIn, class CStencil, class COut>
@ -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<decltype(inputPortal), decltype(outputPortal)> 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<vtkm::UInt64> 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<DeviceAdapterTag>::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<DeviceAdapterTag>::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<decltype(portal)> 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<decltype(portal)> 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<decltype(inputPortal), decltype(valuesPortal), decltype(outputPortal)> 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<decltype(inputPortal),
decltype(valuesPortal),
@ -504,6 +542,8 @@ public:
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::Token token;
//Crazy Idea:
//We create a implicit array handle that wraps the input
//array handle. The implicit functor is passed the input array handle, and
@ -513,7 +553,7 @@ public:
//
//Now that we have an implicit array that is 1/16 the length of full array
//we can use scan inclusive to compute the final sum
auto inputPortal = input.PrepareForInput(DeviceAdapterTag());
auto inputPortal = input.PrepareForInput(DeviceAdapterTag(), token);
ReduceKernel<decltype(inputPortal), U, BinaryFunctor> kernel(
inputPortal, initialValue, binary_functor);
@ -605,8 +645,9 @@ public:
vtkm::cont::ArrayHandle<ReduceKeySeriesStates> 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<decltype(inputPortal), decltype(keyStatePortal)> kernel(
inputPortal, keyStatePortal);
DerivedAlgorithm::Schedule(kernel, numberOfKeys);
@ -668,8 +709,10 @@ public:
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> 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<decltype(inputPortal), decltype(outputPortal), BinaryFunctor>
inclusiveToExclusive(inputPortal, outputPortal, binaryFunctor, initialValue);
@ -710,8 +753,10 @@ public:
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> 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<decltype(inputPortal), decltype(outputPortal), BinaryFunctor>
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<ReduceKeySeriesStates> 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<decltype(inputPortal), decltype(keyStatePortal)> 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<ValueT, vtkm::cont::StorageTagBasic> 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<ValueT,
decltype(inputPortal),
@ -891,7 +938,9 @@ public:
return vtkm::TypeTraits<T>::ZeroInitialization();
}
auto portal = output.PrepareForInPlace(DeviceAdapterTag());
vtkm::cont::Token token;
auto portal = output.PrepareForInPlace(DeviceAdapterTag(), token);
using ScanKernelType = ScanKernel<decltype(portal), BinaryFunctor>;
@ -945,8 +994,9 @@ public:
vtkm::cont::ArrayHandle<ReduceKeySeriesStates> 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<decltype(inputPortal), decltype(keyStatePortal)> 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<decltype(portal), BinaryCompare>;
using CrossoverKernel = BitonicSortCrossoverKernel<decltype(portal), BinaryCompare>;
@ -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<decltype(input1Portal),
decltype(input2Portal),
@ -1101,12 +1155,17 @@ public:
using WrappedBOpType = internal::WrappedBinaryOperator<bool, BinaryCompare>;
WrappedBOpType wrappedCompare(binary_compare);
auto valuesPortal = values.PrepareForInput(DeviceAdapterTag());
auto stencilPortal = stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag());
ClassifyUniqueComparisonKernel<decltype(valuesPortal), decltype(stencilPortal), WrappedBOpType>
classifyKernel(valuesPortal, stencilPortal, wrappedCompare);
{
vtkm::cont::Token token;
auto valuesPortal = values.PrepareForInput(DeviceAdapterTag(), token);
auto stencilPortal = stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag(), token);
ClassifyUniqueComparisonKernel<decltype(valuesPortal),
decltype(stencilPortal),
WrappedBOpType>
classifyKernel(valuesPortal, stencilPortal, wrappedCompare);
DerivedAlgorithm::Schedule(classifyKernel, inputSize);
DerivedAlgorithm::Schedule(classifyKernel, inputSize);
}
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> 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<decltype(inputPortal), decltype(valuesPortal), decltype(outputPortal)> 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<decltype(inputPortal),
decltype(valuesPortal),

13
vtkm/cont/internal/ReverseConnectivityBuilder.h
View File

@ -146,7 +146,8 @@ public:
{
using Algo = vtkm::cont::DeviceAdapterAlgorithm<Device>;
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<vtkm::IdComponent> atomicCounter{ rNumIndices };
auto ac = atomicCounter.PrepareForExecution(Device());
vtkm::cont::Token token;
auto ac = atomicCounter.PrepareForExecution(Device(), token);
using BuildHisto =
rcb::BuildHistogram<decltype(ac), decltype(connPortal), RConnToConnIdxCalc>;
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<vtkm::IdComponent> 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<decltype(ac),
decltype(connPortal),

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

@ -14,6 +14,7 @@
#include <vtkm/VirtualObjectBase.h>
#include <vtkm/cont/DeviceAdapterTag.h>
#include <vtkm/cont/Token.h>
#include <array>
#include <memory>
@ -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<std::size_t>(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; }

5
vtkm/cont/internal/VirtualObjectTransferShareWithControl.h
View File

@ -13,6 +13,8 @@
#include <vtkm/StaticAssert.h>
#include <vtkm/VirtualObjectBase.h>
#include <vtkm/cont/Token.h>
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;
}

44
vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h
View File

@ -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<typename std::decay<U>::type> fastPath;
return ReduceHelper::Execute(portal, initialValue, binary_functor, fastPath);
@ -238,13 +243,15 @@ public:
return vtkm::TypeTraits<T>::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<InPortalT, OutPortalT, BinaryFunctor>;
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<InPortalT, OutPortalT, BinaryFunctor>;
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<decltype(iter)>::type;
@ -347,6 +356,7 @@ public:
Uniqifier uniquifier(iter, portal.GetNumberOfValues(), binary_compare);
vtkm::Id outSize = uniquifier.Execute();
token.DetachFromAll();
values.Shrink(outSize);
}

62
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h
View File

@ -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<U, BinaryFunctor> 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<decltype(arrayPortal)> iterators(arrayPortal);
internal::WrappedBinaryOperator<bool, BinaryCompare> 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<decltype(arrayPortal)> iterators(arrayPortal);
internal::WrappedBinaryOperator<bool, BinaryCompare> wrappedCompare(binary_compare);
auto end = std::unique(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare);
values.Shrink(static_cast<vtkm::Id>(end - iterators.GetBegin()));
vtkm::Id newSize = static_cast<vtkm::Id>(end - iterators.GetBegin());
token.DetachFromAll();
values.Shrink(newSize);
}
VTKM_CONT static void Synchronize()

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

@ -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 <typename T,
@ -172,14 +180,17 @@ public:
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::Token token;
vtkm::Id inputSize = keys.GetNumberOfValues();
VTKM_ASSERT(inputSize == values.GetNumberOfValues());
vtkm::Id outputSize =
tbb::ReduceByKeyPortals(keys.PrepareForInput(DeviceAdapterTagTBB()),
values.PrepareForInput(DeviceAdapterTagTBB()),
keys_output.PrepareForOutput(inputSize, DeviceAdapterTagTBB()),
values_output.PrepareForOutput(inputSize, DeviceAdapterTagTBB()),
binary_functor);
vtkm::Id outputSize = tbb::ReduceByKeyPortals(
keys.PrepareForInput(DeviceAdapterTagTBB(), token),
values.PrepareForInput(DeviceAdapterTagTBB(), token),
keys_output.PrepareForOutput(inputSize, DeviceAdapterTagTBB(), token),
values_output.PrepareForOutput(inputSize, DeviceAdapterTagTBB(), token),
binary_functor);
token.DetachFromAll();
keys_output.Shrink(outputSize);
values_output.Shrink(outputSize);
}
@ -190,9 +201,10 @@ public:
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::Token token;
return tbb::ScanInclusivePortals(
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),
vtkm::Add());
}
@ -203,9 +215,10 @@ public:
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);
vtkm::cont::Token token;
return tbb::ScanInclusivePortals(
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);
}
@ -215,9 +228,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),
vtkm::Add(),
vtkm::TypeTraits<T>::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<T, Storage>& 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);
}

43
vtkm/cont/tbb/internal/ParallelSortTBB.h
View File

@ -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<decltype(arrayPortal)>;
IteratorsType iterators(arrayPortal);
@ -106,10 +107,11 @@ void parallel_sort_bykey(vtkm::cont::ArrayHandle<T, StorageT>& 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<T, StorageT>& keys,
vtkm::cont::internal::KeyCompare<T, vtkm::Id, BinaryCompare>(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<T, StorageT>& 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<T, StorageT>& 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());
}
}

48
vtkm/cont/testing/TestingArrayHandles.h
View File

@ -141,6 +141,7 @@ private:
arrayHandle = vtkm::cont::ArrayHandle<T>();
VTKM_TEST_ASSERT(arrayHandle.GetPortalConstControl().GetNumberOfValues() == 0,
"Uninitialized array does not give portal with zero values.");
vtkm::cont::Token token;
arrayHandle = vtkm::cont::ArrayHandle<T>();
arrayHandle.Shrink(0);
arrayHandle = vtkm::cont::ArrayHandle<T>();
@ -148,11 +149,11 @@ private:
arrayHandle = vtkm::cont::ArrayHandle<T>();
arrayHandle.ReleaseResources();
arrayHandle = vtkm::cont::make_ArrayHandle(std::vector<T>());
arrayHandle.PrepareForInput(DeviceAdapterTag());
arrayHandle.PrepareForInput(DeviceAdapterTag(), token);
arrayHandle = vtkm::cont::ArrayHandle<T>();
arrayHandle.PrepareForInPlace(DeviceAdapterTag());
arrayHandle.PrepareForInPlace(DeviceAdapterTag(), token);
arrayHandle = vtkm::cont::ArrayHandle<T>();
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<T>::template ExecutionTypes<DeviceAdapterTag>::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<T> result;
@ -189,9 +192,11 @@ private:
std::cout << "Check out inplace." << std::endl;
{ //as inplace
vtkm::cont::Token token;
typename vtkm::cont::ArrayHandle<T>::template ExecutionTypes<DeviceAdapterTag>::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<T> 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<T>::template ExecutionTypes<DeviceAdapterTag>::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<T>::template ExecutionTypes<DeviceAdapterTag>::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<T> result;
@ -268,9 +278,11 @@ private:
std::cout << "Check out inplace." << std::endl;
{ //as inplace
vtkm::cont::Token token;
typename vtkm::cont::ArrayHandle<T>::template ExecutionTypes<DeviceAdapterTag>::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<T> 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<T>::template ExecutionTypes<DeviceAdapterTag>::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<T>::template ExecutionTypes<DeviceAdapterTag>::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<T>::template ExecutionTypes<DeviceAdapterTag>::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.");
}

14
vtkm/cont/testing/TestingBitField.h
View File

@ -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

10
vtkm/cont/testing/TestingCellLocatorRectilinearGrid.h
View File

@ -35,11 +35,14 @@ public:
using RectilinearPortalType =
typename RectilinearType::template ExecutionTypes<DeviceAdapter>::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<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> parametric;
vtkm::cont::ArrayHandle<bool> match;
vtkm::cont::Token token;
LocatorWorklet<DeviceAdapter> worklet(
bounds, dims, coords.GetData().template Cast<RectilinearType>());
bounds, dims, coords.GetData().template Cast<RectilinearType>(), token);
vtkm::worklet::DispatcherMapField<LocatorWorklet<DeviceAdapter>> dispatcher(worklet);
dispatcher.SetDevice(DeviceAdapter());

11
vtkm/cont/testing/TestingColorTable.h
View File

@ -476,10 +476,13 @@ public:
auto samples = vtkm::cont::make_ArrayHandle(data, nvals);
vtkm::cont::ArrayHandle<vtkm::Vec4ui_8> colors;
TransferFunction transfer(table.PrepareForExecution(DeviceAdapterTag{}));
vtkm::worklet::DispatcherMapField<TransferFunction> dispatcher(transfer);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(samples, colors);
{
vtkm::cont::Token token;
TransferFunction transfer(table.PrepareForExecution(DeviceAdapterTag{}, token));
vtkm::worklet::DispatcherMapField<TransferFunction> 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 },

243
vtkm/cont/testing/TestingDeviceAdapter.h
View File

@ -318,8 +318,8 @@ public:
struct AtomicKernel
{
VTKM_CONT
AtomicKernel(const vtkm::cont::AtomicArray<T>& array)
: AArray(array.PrepareForExecution(DeviceAdapterTag()))
AtomicKernel(const vtkm::cont::AtomicArray<T>& array, vtkm::cont::Token& token)
: AArray(array.PrepareForExecution(DeviceAdapterTag(), token))
{
}
@ -338,8 +338,8 @@ public:
struct AtomicCASKernel
{
VTKM_CONT
AtomicCASKernel(const vtkm::cont::AtomicArray<T>& array)
: AArray(array.PrepareForExecution(DeviceAdapterTag()))
AtomicCASKernel(const vtkm::cont::AtomicArray<T>& 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<vtkm::Vec4f_32, StorageTagBasic> 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<const BaseType*>(transfer.PrepareForExecution(false));
vtkm::cont::Token transferToken;
const BaseType* base =
static_cast<const BaseType*>(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<const BaseType*>(transfer.PrepareForExecution(true));
Algorithm::Schedule(VirtualObjectTransferKernel(base, result), 1);
{
vtkm::cont::Token token;
target.Value = 10;
base = static_cast<const BaseType*>(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<vtkm::Id> 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<vtkm::Id> 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<BoolPortal>(
tracker.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()), false),
ARRAY_SIZE);
Algorithm::Schedule(GenericClearArrayKernel<BoolPortal>(
valid.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag()), false),
ARRAY_SIZE);
{
vtkm::cont::Token token;
Algorithm::Schedule(
GenericClearArrayKernel<BoolPortal>(
tracker.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token), false),
ARRAY_SIZE);
Algorithm::Schedule(GenericClearArrayKernel<BoolPortal>(
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<BoolPortal>(
tracker.PrepareForOutput(numElems, DeviceAdapterTag()), dims, false),
numElems);
Algorithm::Schedule(GenericClearArrayKernel<BoolPortal>(
valid.PrepareForOutput(numElems, DeviceAdapterTag()), dims, false),
numElems);
{
vtkm::cont::Token token;
Algorithm::Schedule(
GenericClearArrayKernel<BoolPortal>(
tracker.PrepareForOutput(numElems, DeviceAdapterTag(), token), dims, false),
numElems);
Algorithm::Schedule(
GenericClearArrayKernel<BoolPortal>(
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<IdArrayHandle, IdArrayHandle> 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<vtkm::Int32> atomic(atomicElement);
Algorithm::Schedule(AtomicKernel<vtkm::Int32>(atomic), SHORT_ARRAY_SIZE);
{
vtkm::cont::Token token;
Algorithm::Schedule(AtomicKernel<vtkm::Int32>(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<vtkm::Int64> atomic(atomicElement);
Algorithm::Schedule(AtomicKernel<vtkm::Int64>(atomic), SHORT_ARRAY_SIZE);
{
vtkm::cont::Token token;
Algorithm::Schedule(AtomicKernel<vtkm::Int64>(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<vtkm::Int32> atomic(atomicElement);
Algorithm::Schedule(AtomicCASKernel<vtkm::Int32>(atomic), SHORT_ARRAY_SIZE);
{
vtkm::cont::Token token;
Algorithm::Schedule(AtomicCASKernel<vtkm::Int32>(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<vtkm::Int64> atomic(atomicElement);
Algorithm::Schedule(AtomicCASKernel<vtkm::Int64>(atomic), SHORT_ARRAY_SIZE);
{
vtkm::cont::Token token;
Algorithm::Schedule(AtomicCASKernel<vtkm::Int64>(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");

5
vtkm/cont/testing/TestingFancyArrayHandles.h
View File

@ -188,9 +188,10 @@ struct TransformExecObject : public vtkm::cont::ExecutionAndControlObjectBase
};
template <typename DeviceAdapterTag>
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

3
vtkm/cont/testing/TestingImplicitFunction.h
View File

@ -65,7 +65,8 @@ void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points,
{
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvaluateImplicitFunction>;
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);

12
vtkm/cont/testing/TestingVirtualObjectHandle.h
View File

@ -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;

2
vtkm/cont/testing/UnitTestAlgorithm.cxx
View File

@ -147,7 +147,7 @@ struct CompFunctor
struct CompExecObject : vtkm::cont::ExecutionObjectBase
{
template <typename Device>
VTKM_CONT CompFunctor PrepareForExecution(Device)
VTKM_CONT CompFunctor PrepareForExecution(Device, vtkm::cont::Token&)
{
return CompFunctor();
}

10
vtkm/cont/testing/UnitTestArrayHandleDiscard.cxx
View File

@ -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()()

12
vtkm/cont/testing/UnitTestArrayHandleExtractComponent.cxx
View File

@ -158,11 +158,15 @@ struct ExtractComponentTests
using RefPortal =
typename ReferenceCompositeArray::template ExecutionTypes<DeviceTag>::PortalConst;
WriteTestFunctor<Portal, RefPortal> functor(extract.PrepareForInPlace(DeviceTag()),
this->RefComposite.PrepareForInput(DeviceTag()),
component);
{
vtkm::cont::Token token;
WriteTestFunctor<Portal, RefPortal> 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);
}
}

3
vtkm/cont/testing/UnitTestArrayHandleImplicit.cxx
View File

@ -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<Device>::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);

29
vtkm/cont/testing/UnitTestArrayHandlePermutation.cxx
View File

@ -58,12 +58,14 @@ struct CheckPermutationFunctor : vtkm::exec::FunctorBase
template <typename PermutedArrayHandleType, typename Device>
VTKM_CONT CheckPermutationFunctor<
typename PermutedArrayHandleType::template ExecutionTypes<Device>::PortalConst>
make_CheckPermutationFunctor(const PermutedArrayHandleType& permutedArray, Device)
make_CheckPermutationFunctor(const PermutedArrayHandleType& permutedArray,
Device,
vtkm::cont::Token& token)
{
using PermutedPortalType =
typename PermutedArrayHandleType::template ExecutionTypes<Device>::PortalConst;
CheckPermutationFunctor<PermutedPortalType> functor;
functor.PermutedPortal = permutedArray.PrepareForInput(Device());
functor.PermutedPortal = permutedArray.PrepareForInput(Device(), token);
return functor;
}
@ -87,12 +89,14 @@ struct InPlacePermutationFunctor : vtkm::exec::FunctorBase
template <typename PermutedArrayHandleType, typename Device>
VTKM_CONT InPlacePermutationFunctor<
typename PermutedArrayHandleType::template ExecutionTypes<Device>::Portal>
make_InPlacePermutationFunctor(PermutedArrayHandleType& permutedArray, Device)
make_InPlacePermutationFunctor(PermutedArrayHandleType& permutedArray,
Device,
vtkm::cont::Token& token)
{
using PermutedPortalType =
typename PermutedArrayHandleType::template ExecutionTypes<Device>::Portal;
InPlacePermutationFunctor<PermutedPortalType> functor;
functor.PermutedPortal = permutedArray.PrepareForInPlace(Device());
functor.PermutedPortal = permutedArray.PrepareForInPlace(Device(), token);
return functor;
}
@ -136,12 +140,14 @@ struct OutputPermutationFunctor : vtkm::exec::FunctorBase
template <typename PermutedArrayHandleType, typename Device>
VTKM_CONT OutputPermutationFunctor<
typename PermutedArrayHandleType::template ExecutionTypes<Device>::Portal>
make_OutputPermutationFunctor(PermutedArrayHandleType& permutedArray, Device)
make_OutputPermutationFunctor(PermutedArrayHandleType& permutedArray,
Device,
vtkm::cont::Token& token)
{
using PermutedPortalType =
typename PermutedArrayHandleType::template ExecutionTypes<Device>::Portal;
OutputPermutationFunctor<PermutedPortalType> 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());
}

3
vtkm/cont/testing/UnitTestArrayHandleSwizzle.cxx
View File

@ -161,8 +161,9 @@ struct SwizzleTests
template <typename DeviceTag, typename SwizzleHandleType>
bool operator()(DeviceTag, SwizzleHandleType& swizzle) const
{
vtkm::cont::Token token;
using Portal = typename SwizzleHandleType::template ExecutionTypes<DeviceTag>::Portal;
WriteTestFunctor<Portal> functor(swizzle.PrepareForInPlace(DeviceTag()));
WriteTestFunctor<Portal> functor(swizzle.PrepareForInPlace(DeviceTag(), token));
Algo::Schedule(functor, swizzle.GetNumberOfValues());
return true;
}

52
vtkm/cont/testing/UnitTestArrayHandleTransform.cxx
View File

@ -59,15 +59,16 @@ VTKM_CONT CheckTransformFunctor<
typename TransformedArrayHandleType::template ExecutionTypes<Device>::PortalConst>
make_CheckTransformFunctor(const OriginalArrayHandleType& originalArray,
const TransformedArrayHandleType& transformedArray,
Device)
Device,
vtkm::cont::Token& token)
{
using OriginalPortalType =
typename OriginalArrayHandleType::template ExecutionTypes<Device>::PortalConst;
using TransformedPortalType =
typename TransformedArrayHandleType::template ExecutionTypes<Device>::PortalConst;
CheckTransformFunctor<OriginalPortalType, TransformedPortalType> 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<InputValueType> 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<InputValueType> 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<InputValueType>::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);
}
}
};

7
vtkm/cont/testing/UnitTestArrayHandleVirtual.cxx
View File

@ -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&)
{

3
vtkm/cont/testing/UnitTestCellSetExplicit.cxx
View File

@ -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.");

11
vtkm/cont/testing/UnitTestDataSetPermutation.cxx
View File

@ -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<vtkm::Float32> 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<vtkm::Float32> 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<vtkm::Float32> result;

15
vtkm/cont/testing/UnitTestDataSetRectilinear.cxx
View File

@ -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<vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, 2>
pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(),
vtkm::TopologyElementTagCell(),
vtkm::TopologyElementTagPoint());
vtkm::TopologyElementTagPoint(),
token);
vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, 2>
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<vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, 3>
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<vtkm::Id, 8> retrievedPointIds = pointToCell.GetIndices(vtkm::Id3(0));
for (vtkm::IdComponent localPointIndex = 0; localPointIndex < 8; localPointIndex++)
@ -176,7 +182,8 @@ static void ThreeDimRectilinearTest()
vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, 3>
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<vtkm::Id, 6> expectedCellIds = cellToPoint.GetIndices(vtkm::Id3(0));
VTKM_TEST_ASSERT(expectedCellIds.GetNumberOfComponents() <= 6,

15
vtkm/cont/testing/UnitTestDataSetUniform.cxx
View File

@ -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<vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, 2>
pointToCell = cellSet.PrepareForInput(vtkm::cont::DeviceAdapterTagSerial(),
vtkm::TopologyElementTagCell(),
vtkm::TopologyElementTagPoint());
vtkm::TopologyElementTagPoint(),
token);
vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, 2>
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<vtkm::TopologyElementTagCell, vtkm::TopologyElementTagPoint, 3>
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<vtkm::Id, 8> retrievedPointIds = pointToCell.GetIndices(vtkm::Id3(0));
for (vtkm::IdComponent localPointIndex = 0; localPointIndex < 8; localPointIndex++)
@ -186,7 +192,8 @@ static void ThreeDimUniformTest()
vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint, vtkm::TopologyElementTagCell, 3>
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<vtkm::Id, 6> expectedCellIds = cellToPoint.GetIndices(vtkm::Id3(0));
VTKM_TEST_ASSERT(expectedCellIds.GetNumberOfComponents() <= 6,

16
vtkm/cont/testing/UnitTestMoveConstructors.cxx
View File

@ -81,16 +81,16 @@ struct IsNoExceptHandle
is_noexcept_movable<VirtualType>();
//verify the input portals of the handle
is_noexcept_movable<decltype(
std::declval<HandleType>().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}))>();
is_noexcept_movable<decltype(
std::declval<VirtualType>().PrepareForInput(vtkm::cont::DeviceAdapterTagSerial{}))>();
is_noexcept_movable<decltype(std::declval<HandleType>().PrepareForInput(
vtkm::cont::DeviceAdapterTagSerial{}, std::declval<vtkm::cont::Token&>()))>();
is_noexcept_movable<decltype(std::declval<VirtualType>().PrepareForInput(
vtkm::cont::DeviceAdapterTagSerial{}, std::declval<vtkm::cont::Token&>()))>();
//verify the output portals of the handle
is_noexcept_movable<decltype(
std::declval<HandleType>().PrepareForOutput(2, vtkm::cont::DeviceAdapterTagSerial{}))>();
is_noexcept_movable<decltype(
std::declval<VirtualType>().PrepareForOutput(2, vtkm::cont::DeviceAdapterTagSerial{}))>();
is_noexcept_movable<decltype(std::declval<HandleType>().PrepareForOutput(
2, vtkm::cont::DeviceAdapterTagSerial{}, std::declval<vtkm::cont::Token&>()))>();
is_noexcept_movable<decltype(std::declval<VirtualType>().PrepareForOutput(
2, vtkm::cont::DeviceAdapterTagSerial{}, std::declval<vtkm::cont::Token&>()))>();
}
};

11
vtkm/exec/CellLocatorBoundingIntervalHierarchyExec.h
View File

@ -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))
{
}

7
vtkm/exec/CellLocatorRectilinearGrid.h
View File

@ -49,11 +49,12 @@ public:
const vtkm::Id rowSize,
const vtkm::cont::CellSetStructured<dimensions>& 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();

5
vtkm/exec/CellLocatorUniformGrid.h
View File

@ -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))
{
}

11
vtkm/rendering/Texture2D.h
View File

@ -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 <typename Device>
VTKM_CONT Texture2DSamplerExecutionObject<Device> PrepareForExecution(Device) const
VTKM_CONT Texture2DSamplerExecutionObject<Device> PrepareForExecution(
Device,
vtkm::cont::Token& token) const
{
return Texture2DSamplerExecutionObject<Device>(
this->Width, this->Height, this->Data, this->FilterMode, this->WrapMode);
this->Width, this->Height, this->Data, this->FilterMode, this->WrapMode, token);
}
private:

52
vtkm/rendering/Wireframer.h
View File

@ -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<vtkm::Id>(Canvas->GetWidth());
vtkm::Id height = static_cast<vtkm::Id>(Canvas->GetHeight());
vtkm::Id pixelCount = width * height;
FrameBuffer.PrepareForOutput(pixelCount, DeviceTag());
if (ShowInternalZones && !IsOverlay)
if (this->ShowInternalZones && !this->IsOverlay)
{
vtkm::cont::ArrayHandleConstant<vtkm::Int64> 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<CopyIntoFrameBuffer>(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<DeviceTag> plotter(WorldToProjection,
width,
height,
subsetWidth,
subsetHeight,
xOffset,
yOffset,
isAssocPoints,
ScalarFieldRange,
ColorMap,
FrameBuffer,
Camera.GetClippingRange());
vtkm::worklet::DispatcherMapField<EdgePlotter<DeviceTag>> plotterDispatcher(plotter);
plotterDispatcher.SetDevice(DeviceTag());
plotterDispatcher.Invoke(
PointIndices, Coordinates, ScalarField.GetData().ResetTypes(vtkm::TypeListFieldScalar()));
{
vtkm::cont::Token token;
EdgePlotter<DeviceTag> plotter(WorldToProjection,
width,
height,
subsetWidth,
subsetHeight,
xOffset,
yOffset,
isAssocPoints,
ScalarFieldRange,
ColorMap,
FrameBuffer,
Camera.GetClippingRange(),
token);
vtkm::worklet::DispatcherMapField<EdgePlotter<DeviceTag>> plotterDispatcher(plotter);
plotterDispatcher.SetDevice(DeviceTag());
plotterDispatcher.Invoke(
PointIndices, Coordinates, ScalarField.GetData().ResetTypes(vtkm::TypeListFieldScalar()));
}
BufferConverter converter;
vtkm::worklet::DispatcherMapField<BufferConverter> converterDispatcher(converter);

66
vtkm/rendering/raytracing/ChannelBuffer.cxx
View File

@ -9,8 +9,10 @@
//============================================================================
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/rendering/vtkm_rendering_export.h>
@ -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 <typename T, typename BufferPortalType>
VTKM_EXEC void operator()(T& outValue,
const BufferPortalType& inBuffer,
@ -169,34 +171,6 @@ public:
}
}; //class Extract Channel
template <typename Precision>
struct ExtractChannelFunctor
{
ChannelBuffer<Precision>* Self;
vtkm::cont::ArrayHandle<Precision> Output;
vtkm::Int32 Channel;
ExtractChannelFunctor(ChannelBuffer<Precision>* self,
vtkm::cont::ArrayHandle<Precision> output,
const vtkm::Int32 channel)
: Self(self)
, Output(output)
, Channel(channel)
{
}
template <typename Device>
bool operator()(Device device)
{
Output.PrepareForOutput(Self->GetSize(), device);
vtkm::worklet::DispatcherMapField<ExtractChannel> dispatcher(
ExtractChannel(Self->GetNumChannels(), Channel));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Output, Self->Buffer);
return true;
}
};
template <typename Precision>
ChannelBuffer<Precision> ChannelBuffer<Precision>::GetChannel(const vtkm::Int32 channel)
{
@ -208,9 +182,13 @@ ChannelBuffer<Precision> ChannelBuffer<Precision>::GetChannel(const vtkm::Int32
{
return output;
}
ExtractChannelFunctor<Precision> 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 <typename Device>
bool operator()(Device device)
{
vtkm::Id totalSize = OutputLength * static_cast<vtkm::Id>(NumChannels);
Output->Buffer.PrepareForOutput(totalSize, device);
ChannelBufferOperations::InitChannels(*Output, Signature, device);
vtkm::cont::Token token;
vtkm::worklet::DispatcherMapField<Expand> dispatcher((Expand(NumChannels)));
vtkm::Id totalSize = this->OutputLength * static_cast<vtkm::Id>(this->NumChannels);
this->Output->Buffer.PrepareForOutput(totalSize, device, token);
ChannelBufferOperations::InitChannels(*this->Output, this->Signature, device);
vtkm::worklet::DispatcherMapField<Expand> 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 <typename Device>
bool operator()(Device device)
{
vtkm::Id totalSize = OutputLength * static_cast<vtkm::Id>(NumChannels);
Output->Buffer.PrepareForOutput(totalSize, device);
ChannelBufferOperations::InitConst(*Output, InitVal, device);
vtkm::cont::Token token;
vtkm::worklet::DispatcherMapField<Expand> dispatcher((Expand(NumChannels)));
vtkm::Id totalSize = this->OutputLength * static_cast<vtkm::Id>(this->NumChannels);
this->Output->Buffer.PrepareForOutput(totalSize, device, token);
ChannelBufferOperations::InitConst(*this->Output, this->InitVal, device);
vtkm::worklet::DispatcherMapField<Expand> dispatcher((Expand(this->NumChannels)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Input, SparseIndexes, Output->Buffer);
dispatcher.Invoke(this->Input, this->SparseIndexes, this->Output->Buffer);
return true;
}

9
vtkm/rendering/raytracing/ChannelBuffer.h
View File

@ -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 <typename Device>
@ -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<vtkm::Id>(NumChannels), device);
vtkm::cont::Token token;
this->Buffer.PrepareForOutput(this->Size * static_cast<vtkm::Id>(NumChannels), device, token);
}
};
}

13
vtkm/rendering/raytracing/CylinderIntersector.cxx
View File

@ -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 <typename Device>
VTKM_CONT CylinderLeafIntersector<Device> PrepareForExecution(Device) const
VTKM_CONT CylinderLeafIntersector<Device> PrepareForExecution(Device,
vtkm::cont::Token& token) const
{
return CylinderLeafIntersector<Device>(CylIds, Radii);
return CylinderLeafIntersector<Device>(this->CylIds, this->Radii, token);
}
};

24
vtkm/rendering/raytracing/MeshConnectivityBase.h
View File

@ -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)

3
vtkm/rendering/raytracing/MeshConnectivityBuilder.cxx
View File

@ -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<StructuredExternalTriangles> dispatch(
StructuredExternalTriangles(cellSet.PrepareForInput(
Device(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint())));
Device(), vtkm::TopologyElementTagCell(), vtkm::TopologyElementTagPoint(), token)));
dispatch.SetDevice(Device());
dispatch.Invoke(counting, triangles);

62
vtkm/rendering/raytracing/MeshConnectivityContainers.cxx
View File

@ -38,9 +38,10 @@ VTKM_CONT void MeshConnContainer::FindEntryImpl(Ray<T>& 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<MeshConnectivityBase*>(this->Construct(deviceId)));
return MeshWrapper(const_cast<MeshConnectivityBase*>(this->Construct(deviceId, token)));
}
void MeshConnContainer::FindEntry(Ray<vtkm::Float32>& 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);
}
}
}

15
vtkm/rendering/raytracing/MeshConnectivityContainers.h
View File

@ -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 <typename T>
VTKM_CONT void FindEntryImpl(Ray<T>& 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
}

8
vtkm/rendering/raytracing/QuadIntersector.cxx
View File

@ -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 <typename Device>
VTKM_CONT QuadLeafIntersector<Device> PrepareForExecution(Device) const
VTKM_CONT QuadLeafIntersector<Device> PrepareForExecution(Device, vtkm::cont::Token& token) const
{
return QuadLeafIntersector<Device>(QuadIds);
return QuadLeafIntersector<Device>(QuadIds, token);
}
};

62
vtkm/rendering/raytracing/Ray.h
View File

@ -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);

43
vtkm/rendering/raytracing/RayOperations.h
View File

@ -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<size_t>(rays.Buffers.size());
for (size_t i = 0; i < bufferCount; ++i)

5
vtkm/rendering/raytracing/RayTracingTypeDefs.h
View File

@ -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

13
vtkm/rendering/raytracing/SphereIntersector.cxx
View File

@ -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 <typename Device>
VTKM_CONT SphereLeafIntersector<Device> PrepareForExecution(Device) const
VTKM_CONT SphereLeafIntersector<Device> PrepareForExecution(Device,
vtkm::cont::Token& token) const
{
return SphereLeafIntersector<Device>(PointIds, Radii);
return SphereLeafIntersector<Device>(this->PointIds, this->Radii, token);
}
};

18
vtkm/rendering/raytracing/TriangleIntersector.cxx
View File

@ -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 <typename Device>
VTKM_CONT MollerTriLeafIntersector<Device> PrepareForExecution(Device) const
VTKM_CONT MollerTriLeafIntersector<Device> PrepareForExecution(Device,
vtkm::cont::Token& token) const
{
return MollerTriLeafIntersector<Device>(Triangles);
return MollerTriLeafIntersector<Device>(this->Triangles, token);
}
};
@ -167,9 +168,10 @@ public:
}
template <typename Device>
VTKM_CONT WaterTightLeafIntersector<Device> PrepareForExecution(Device) const
VTKM_CONT WaterTightLeafIntersector<Device> PrepareForExecution(Device,
vtkm::cont::Token& token) const
{
return WaterTightLeafIntersector<Device>(Triangles);
return WaterTightLeafIntersector<Device>(this->Triangles, token);
}
};

45
vtkm/rendering/raytracing/VolumeRendererStructured.cxx
View File

@ -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<P
if (IsUniformDataSet)
{
vtkm::cont::Token token;
vtkm::cont::ArrayHandleUniformPointCoordinates vertices;
vertices = Coordinates.Cast<vtkm::cont::ArrayHandleUniformPointCoordinates>();
UniformLocator<Device> locator(vertices, Cellset);
UniformLocator<Device> locator(vertices, Cellset, token);
if (isAssocPoints)
{
@ -824,7 +832,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator));
locator,
token));
samplerDispatcher.SetDevice(Device());
samplerDispatcher.Invoke(rays.Dir,
rays.Origin,
@ -840,7 +849,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator))
locator,
token))
.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
@ -851,9 +861,10 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
}
else
{
vtkm::cont::Token token;
CartesianArrayHandle vertices;
vertices = Coordinates.Cast<CartesianArrayHandle>();
RectilinearLocator<Device> locator(vertices, Cellset);
RectilinearLocator<Device> locator(vertices, Cellset, token);
if (isAssocPoints)
{
vtkm::worklet::DispatcherMapField<Sampler<Device, RectilinearLocator<Device>>>
@ -862,7 +873,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator));
locator,
token));
samplerDispatcher.SetDevice(Device());
samplerDispatcher.Invoke(rays.Dir,
rays.Origin,
@ -879,7 +891,8 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator));
locator,
token));
rectilinearLocatorDispatcher.SetDevice(Device());
rectilinearLocatorDispatcher.Invoke(
rays.Dir,

30
vtkm/testing/TestingAlgorithms.h
View File

@ -65,15 +65,19 @@ struct TestBinarySearch
IdArray hayStack = vtkm::cont::make_ArrayHandle(hayStackData);
IdArray results;
vtkm::cont::Token token;
using Functor = Impl<typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::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<typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::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<typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::PortalConst,
typename IdArray::ExecutionTypes<Device>::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)

17
vtkm/worklet/Clip.h
View File

@ -139,11 +139,12 @@ public:
vtkm::cont::ArrayHandle<vtkm::IdComponent> numberOfIndices,
vtkm::cont::ArrayHandle<vtkm::Id> connectivity,
vtkm::cont::ArrayHandle<vtkm::Id> 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 <typename Device>
VTKM_CONT ExecutionConnectivityExplicit<Device> PrepareForExecution(Device) const
VTKM_CONT ExecutionConnectivityExplicit<Device> PrepareForExecution(
Device,
vtkm::cont::Token& token) const
{
ExecutionConnectivityExplicit<Device> execConnectivity(
this->Shapes, this->NumberOfIndices, this->Connectivity, this->Offsets, this->Stats);
this->Shapes, this->NumberOfIndices, this->Connectivity, this->Offsets, this->Stats, token);
return execConnectivity;
}

21
vtkm/worklet/Contour.h
View File

@ -152,11 +152,12 @@ public:
vtkm::cont::ArrayHandle<vtkm::FloatDefault>& interpWeights,
vtkm::cont::ArrayHandle<vtkm::Id2>& interpIds,
vtkm::cont::ArrayHandle<vtkm::Id>& interpCellIds,
vtkm::cont::ArrayHandle<vtkm::UInt8>& 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<vtkm::UInt8>& 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 <typename DeviceAdapter>
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter)
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token)
{
return ExecObject<DeviceAdapter>(
this->Size, this->InterpWeights, this->InterpIds, this->InterpCellIds, this->InterpContourId);
return ExecObject<DeviceAdapter>(this->Size,
this->InterpWeights,
this->InterpIds,
this->InterpCellIds,
this->InterpContourId,
token);
}
private:

15
vtkm/worklet/DispatcherStreamingMapField.h
View File

@ -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 <typename Invocation, typename RangeType, typename DeviceAdapter>

3
vtkm/worklet/PointMerge.h
View File

@ -136,8 +136,7 @@ public:
// Because this class is a POD, we can reuse it in both control and execution environments.
template <typename Device>
BinLocator PrepareForExecution(Device) const
BinLocator PrepareForExecution(vtkm::cont::DeviceAdapterId, vtkm::cont::Token&) const
{
return *this;
}

8
vtkm/worklet/StableSortIndices.h
View File

@ -77,9 +77,9 @@ struct StableSortIndices
template <typename Device>
IndirectSortPredicate<typename KeyArrayType::template ExecutionTypes<Device>::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<decltype(keyPortal)>(keyPortal);
}
};
@ -116,9 +116,9 @@ struct StableSortIndices
template <typename Device>
IndirectUniquePredicate<typename KeyArrayType::template ExecutionTypes<Device>::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<decltype(keyPortal)>(keyPortal);
}
};

9
vtkm/worklet/clip/ClipTables.h
View File

@ -2530,12 +2530,13 @@ public:
}
template <typename DeviceAdapter>
DevicePortal<DeviceAdapter> PrepareForExecution(DeviceAdapter)
DevicePortal<DeviceAdapter> PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token)
{
DevicePortal<DeviceAdapter> 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;
}

20
vtkm/worklet/contour/ContourTables.h
View File

@ -517,15 +517,15 @@ public:
}
template <typename DeviceAdapter>
ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter)
ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token)
{
ExecObject<DeviceAdapter> 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 <typename DeviceAdapter>
ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter)
ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter, vtkm::cont::Token& token)
{
ExecObject<DeviceAdapter> 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;
}

13
vtkm/worklet/contourtree/EdgePeakComparator.h
View File

@ -179,13 +179,14 @@ public:
};
template <typename DeviceAdapter>
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter) const
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter,
vtkm::cont::Token& token) const
{
return ExecObject<DeviceAdapter>(this->Values.PrepareForInput(DeviceAdapter()),
this->ValueIndex.PrepareForInput(DeviceAdapter()),
this->EdgeFar.PrepareForInput(DeviceAdapter()),
this->EdgeNear.PrepareForInput(DeviceAdapter()),
this->ArcArray.PrepareForInput(DeviceAdapter()),
return ExecObject<DeviceAdapter>(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

7
vtkm/worklet/contourtree/VertexMergeComparator.h
View File

@ -163,10 +163,11 @@ public:
};
template <typename DeviceAdapter>
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter) const
VTKM_CONT ExecObject<DeviceAdapter> PrepareForExecution(DeviceAdapter,
vtkm::cont::Token& token) const
{
return ExecObject<DeviceAdapter>(this->Values.PrepareForInput(DeviceAdapter()),
this->Extrema.PrepareForInput(DeviceAdapter()),
return ExecObject<DeviceAdapter>(this->Values.PrepareForInput(DeviceAdapter(), token),
this->Extrema.PrepareForInput(DeviceAdapter(), token),
this->IsJoinTree);
}
}; // VertexMergeComparator

4
vtkm/worklet/contourtree_augmented/ContourTreeMaker.h
View File

@ -725,6 +725,7 @@ struct LeafChainsToContourTree
template <typename DeviceAdapter, typename... Args>
bool operator()(DeviceAdapter device, Args&&... args) const
{
vtkm::cont::Token token;
contourtree_maker_inc_ns::TransferLeafChains_TransferToContourTree<DeviceAdapter> 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<decltype(worklet)> dispatcher(worklet);
dispatcher.SetDevice(device);
dispatcher.Invoke(std::forward<Args>(args)...);

36
vtkm/worklet/contourtree_augmented/activegraph/EdgePeakComparator.h
View File

@ -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 <typename DeviceAdapter>
VTKM_CONT EdgePeakComparatorImpl<DeviceAdapter> PrepareForExecution(DeviceAdapter) const
VTKM_CONT EdgePeakComparatorImpl<DeviceAdapter> PrepareForExecution(
DeviceAdapter,
vtkm::cont::Token& token) const
{
return EdgePeakComparatorImpl<DeviceAdapter>(this->EdgeFar, this->EdgeNear, this->JoinGraph);
return EdgePeakComparatorImpl<DeviceAdapter>(
this->EdgeFar, this->EdgeNear, this->JoinGraph, token);
}
private:

7
vtkm/worklet/contourtree_augmented/activegraph/HyperArcSuperNodeComparator.h
View File

@ -134,11 +134,12 @@ public:
template <typename DeviceAdapter>
VTKM_CONT HyperArcSuperNodeComparatorImpl<DeviceAdapter> PrepareForExecution(
DeviceAdapter device) const
DeviceAdapter device,
vtkm::cont::Token& token) const
{
return HyperArcSuperNodeComparatorImpl<DeviceAdapter>(
this->Hyperparents.PrepareForInput(device),
this->SuperID.PrepareForInput(device),
this->Hyperparents.PrepareForInput(device, token),
this->SuperID.PrepareForInput(device, token),
this->IsJoinTree);
}

20
vtkm/worklet/contourtree_augmented/activegraph/SuperArcNodeComparator.h
View File

@ -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 <typename DeviceAdapter>
VTKM_CONT SuperArcNodeComparatorImpl<DeviceAdapter> PrepareForExecution(DeviceAdapter) const
VTKM_CONT SuperArcNodeComparatorImpl<DeviceAdapter> PrepareForExecution(
DeviceAdapter,
vtkm::cont::Token& token) const
{
return SuperArcNodeComparatorImpl<DeviceAdapter>(this->Superparents, this->JoinSweep);
return SuperArcNodeComparatorImpl<DeviceAdapter>(this->Superparents, this->JoinSweep, token);
}
private:

14
vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeNodeComparator.h
View File

@ -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 <typename DeviceAdapter>
VTKM_CONT ContourTreeNodeComparatorImpl<DeviceAdapter> PrepareForExecution(DeviceAdapter)
VTKM_CONT ContourTreeNodeComparatorImpl<DeviceAdapter> PrepareForExecution(
DeviceAdapter,
vtkm::cont::Token& token)
{
return ContourTreeNodeComparatorImpl<DeviceAdapter>(this->Superparents, this->Superarcs);
return ContourTreeNodeComparatorImpl<DeviceAdapter>(this->Superparents, this->Superarcs, token);
}
private:

15
vtkm/worklet/contourtree_augmented/contourtreemaker/ContourTreeSuperNodeComparator.h
View File

@ -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 <typename DeviceAdapter>
VTKM_CONT ContourTreeSuperNodeComparatorImpl<DeviceAdapter> PrepareForExecution(DeviceAdapter)
VTKM_CONT ContourTreeSuperNodeComparatorImpl<DeviceAdapter> PrepareForExecution(
DeviceAdapter,
vtkm::cont::Token& token)
{
return ContourTreeSuperNodeComparatorImpl<DeviceAdapter>(
this->Hyperparents, this->Supernodes, this->WhenTransferred);
this->Hyperparents, this->Supernodes, this->WhenTransferred, token);
}
private:

34
vtkm/worklet/contourtree_augmented/contourtreemaker/TransferLeafChains_TransferToContourTree.h
View File

@ -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

Some files were not shown because too many files have changed in this diff Show More