diff --git a/Accelerators/Vtkm/DataModel/vtkmlib/CellSetConverters.cxx b/Accelerators/Vtkm/DataModel/vtkmlib/CellSetConverters.cxx
index 06a69d8cb5cb0774093ac3ee2766e49d5532e8e0..e5065c854dbef317cc925c929b5254e4bf1408b0 100644
--- a/Accelerators/Vtkm/DataModel/vtkmlib/CellSetConverters.cxx
+++ b/Accelerators/Vtkm/DataModel/vtkmlib/CellSetConverters.cxx
@@ -33,6 +33,7 @@
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/BinaryPredicates.h>
+#include <vtkm/Swap.h>
#include "vtkCellArray.h"
#include "vtkCellType.h"
@@ -40,12 +41,6 @@
#include "vtkNew.h"
#include "vtkUnsignedCharArray.h"
-#if defined(VTKM_CUDA)
-#define FUNC_SCOPE __device__
-#else
-#define FUNC_SCOPE
-#endif
-
namespace tovtkm
{
VTK_ABI_NAMESPACE_BEGIN
@@ -57,16 +52,10 @@ struct ReorderHex : vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldInOut);
- FUNC_SCOPE void operator()(vtkm::Vec<vtkm::Id, 8>& indices) const
+ VTKM_EXEC void operator()(vtkm::Vec<vtkm::Id, 8>& indices) const
{
- auto doSwap = [&](vtkm::IdComponent id1, vtkm::IdComponent id2) {
- const auto t = indices[id1];
- indices[id1] = indices[id2];
- indices[id2] = t;
- };
-
- doSwap(2, 3);
- doSwap(6, 7);
+ vtkm::Swap(indices[2], indices[3]);
+ vtkm::Swap(indices[6], indices[7]);
}
};
@@ -111,14 +100,7 @@ struct BuildSingleTypeVoxelCellSetVisitor
vtkm::cont::ArrayCopy(
vtkm::cont::make_ArrayHandle(origData, numIds, vtkm::CopyFlag::Off), connHandle);
- // reorder cells from voxel->hex: which only can run on
- // devices that have shared memory / vtable with the CPU
- vtkm::cont::ScopedRuntimeDeviceTracker tracker(
- vtkm::cont::DeviceAdapterTagAny{}, vtkm::cont::RuntimeDeviceTrackerMode::Disable);
- tracker.ResetDevice(vtkm::cont::DeviceAdapterTagTBB{});
- tracker.ResetDevice(vtkm::cont::DeviceAdapterTagOpenMP{});
- tracker.ResetDevice(vtkm::cont::DeviceAdapterTagSerial{});
-
+ // reorder cells from voxel->hex
vtkm::cont::Invoker invoke;
invoke(ReorderHex{}, vtkm::cont::make_ArrayHandleGroupVec<8>(connHandle));
}
diff --git a/Accelerators/Vtkm/Filters/Testing/Cxx/CMakeLists.txt b/Accelerators/Vtkm/Filters/Testing/Cxx/CMakeLists.txt
index 3b3f89d6e55629df9451d8cc5ca5fd23de968981..ed37753a4884649140e3c8fc8ecd4697bb0ac024 100644
--- a/Accelerators/Vtkm/Filters/Testing/Cxx/CMakeLists.txt
+++ b/Accelerators/Vtkm/Filters/Testing/Cxx/CMakeLists.txt
@@ -1,5 +1,6 @@
vtk_add_test_cxx(vtkAcceleratorsVTKmFiltersCxxTests tests
+ TestVTKMAbort.cxx,NO_VALID
TestVTKMCleanGrid.cxx
TestVTKMCoordinateSystemTransform.cxx,NO_VALID
TestVTKMClip.cxx
diff --git a/Accelerators/Vtkm/Filters/Testing/Cxx/TestVTKMAbort.cxx b/Accelerators/Vtkm/Filters/Testing/Cxx/TestVTKMAbort.cxx
new file mode 100644
index 0000000000000000000000000000000000000000..86ea0f9c2599386f384db96dbab8833c227f7872
--- /dev/null
+++ b/Accelerators/Vtkm/Filters/Testing/Cxx/TestVTKMAbort.cxx
@@ -0,0 +1,109 @@
+/*=========================================================================
+
+ Program: Visualization Toolkit
+ Module: TestAbortExecute.cxx
+
+ Copyright (c) Ken Martin, Will Schroeder, Bill Lorensen
+ All rights reserved.
+ See Copyright.txt or http://www.kitware.com/Copyright.htm for details.
+
+ This software is distributed WITHOUT ANY WARRANTY; without even
+ the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+ PURPOSE. See the above copyright notice for more information.
+
+=========================================================================*/
+
+#include "vtkInformation.h"
+#include "vtkLogger.h"
+#include "vtkPlane.h"
+#include "vtkRTAnalyticSource.h"
+#include "vtkShrinkFilter.h"
+#include "vtkUnstructuredGrid.h"
+#include "vtkmClip.h"
+#include "vtkmContour.h"
+
+#include <vtkm/cont/Initialize.h>
+
+int TestVTKMAbort(int, char*[])
+{
+ vtkNew<vtkRTAnalyticSource> wavelet;
+ vtkNew<vtkShrinkFilter> shrink;
+ vtkNew<vtkmContour> contour;
+ vtkNew<vtkmClip> clip;
+
+ wavelet->SetWholeExtent(0, 10, 0, 10, 0, 10);
+
+ shrink->SetInputConnection(wavelet->GetOutputPort());
+
+ contour->SetInputConnection(shrink->GetOutputPort());
+ contour->GenerateValues(5, -6, 250);
+
+ vtkNew<vtkPlane> clipPlane;
+ clipPlane->SetNormal(1, 0, 0);
+ clipPlane->SetOrigin(0, 0, 0);
+
+ clip->SetInputConnection(contour->GetOutputPort());
+ clip->SetClipFunction(clipPlane);
+
+ //--------------------------------------------------------------------------
+ std::cout << "Run 1 with abort on contour\n";
+
+ contour->SetAbortExecuteAndUpdateTime();
+ clip->Update();
+
+ if (!contour->GetAbortExecute())
+ {
+ vtkLog(ERROR, "Contour AbortExecute flag is not set.");
+ return 1;
+ }
+
+ if (shrink->GetAbortExecute() || wavelet->GetAbortExecute() || clip->GetAbortExecute())
+ {
+ vtkLog(ERROR, "Shrink, Wavelet, or Clip AbortExecute flag is set.");
+ return 1;
+ }
+
+ if (!contour->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()) ||
+ !clip->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()))
+ {
+ vtkLog(ERROR, "Contour, or Clip ABORTED flag is not set.");
+ return 1;
+ }
+
+ if (clip->GetOutput()->GetNumberOfPoints())
+ {
+ vtkLog(ERROR, "Found output data.");
+ return 1;
+ }
+
+ //--------------------------------------------------------------------------
+ std::cout << "Run 2 with no aborts\n";
+ contour->SetAbortExecute(0);
+ clip->Update();
+
+ if (wavelet->GetAbortExecute() || shrink->GetAbortExecute() || contour->GetAbortExecute() ||
+ clip->GetAbortExecute())
+ {
+ vtkLog(ERROR, "Wavelet, Shrink, Contour, or Clip AbortExecute flag is set.");
+ return 1;
+ }
+
+ if (wavelet->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()) ||
+ shrink->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()) ||
+ contour->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()) ||
+ clip->GetOutputInformation(0)->Get(vtkAlgorithm::ABORTED()))
+ {
+ vtkLog(ERROR, "Wavelet, Shrink, Contour, or Clip ABORTED flag is set.");
+ return 1;
+ }
+
+ if (!clip->GetOutput()->GetNumberOfPoints())
+ {
+ vtkLog(ERROR, "No output data.");
+ return 1;
+ }
+
+ std::cout << "Tests successful\n";
+
+ return 0;
+}
diff --git a/Accelerators/Vtkm/Filters/vtkmContour.cxx b/Accelerators/Vtkm/Filters/vtkmContour.cxx
index 07a2642ab2147551308613cc49d6673135a03235..585c6266855de720d508e592f583ebbe7150fcae 100644
--- a/Accelerators/Vtkm/Filters/vtkmContour.cxx
+++ b/Accelerators/Vtkm/Filters/vtkmContour.cxx
@@ -33,6 +33,7 @@
#include "vtkmlib/PolyDataConverter.h"
#include <vtkm/cont/ErrorFilterExecution.h>
+#include <vtkm/cont/ErrorUserAbort.h>
#include <vtkm/filter/contour/Contour.h>
#include <vtkm/worklet/WorkletMapField.h>
@@ -185,6 +186,8 @@ int vtkmContour::RequestData(
try
{
+ vtkm::cont::ScopedRuntimeDeviceTracker rtdt([&]() { return this->CheckAbort(); });
+
if (!this->CanProcessInput(input))
{
throw vtkm::cont::ErrorFilterExecution(
@@ -236,6 +239,11 @@ int vtkmContour::RequestData(
filter.GetNormalArrayName().c_str(), vtkDataSetAttributes::NORMALS);
}
}
+ catch (const vtkm::cont::ErrorUserAbort&)
+ {
+ // vtkm detected an abort request, clear the output
+ output->Initialize();
+ }
catch (const vtkm::cont::Error& e)
{
vtkWarningMacro(<< "VTK-m failed with message: " << e.GetMessage() << "\n"
diff --git a/Documentation/release/dev/use-vtkm-abort.md b/Documentation/release/dev/use-vtkm-abort.md
new file mode 100644
index 0000000000000000000000000000000000000000..ef27b44bb58750cfbe8f6bff1963116a21f101c2
--- /dev/null
+++ b/Documentation/release/dev/use-vtkm-abort.md
@@ -0,0 +1,4 @@
+# Add abort functionality to the vtkmContour filter in VTK::AcceleratorsVTKmFilters
+
+VTk-m recently added an abort feature that is compatible with VTK's abort implementation. To
+demonstrate this, we have added abort capability to the `vtkmContour` filter.