From 5842da4921d1e700bd74c57f34cdff111917cad6 Mon Sep 17 00:00:00 2001 From: Sujin Philip Date: Fri, 27 Oct 2017 10:42:22 -0400 Subject: [PATCH 1/2] Remove ArrayHandle CopyInto Fixes #170 --- vtkm/cont/ArrayHandle.h | 6 --- vtkm/cont/ArrayHandle.hxx | 39 ------------------- vtkm/cont/StorageImplicit.h | 11 ------ .../ArrayManagerExecutionThrustDevice.h | 10 ----- vtkm/cont/internal/ArrayHandleBasicImpl.h | 3 -- vtkm/cont/internal/ArrayHandleBasicImpl.hxx | 37 ------------------ .../internal/ArrayHandleExecutionManager.h | 6 --- vtkm/cont/internal/ArrayManagerExecution.h | 9 ----- .../ArrayManagerExecutionShareWithControl.h | 12 ------ vtkm/cont/internal/ArrayTransfer.h | 12 ------ vtkm/cont/testing/TestingArrayHandles.h | 36 ----------------- .../testing/UnitTestTaskSingularCuda.cu | 11 +----- vtkm/interop/internal/TransferToOpenGL.h | 9 +++-- .../testing/UnitTestRemoveUnusedPoints.cxx | 13 +++---- 14 files changed, 13 insertions(+), 201 deletions(-) diff --git a/vtkm/cont/ArrayHandle.h b/vtkm/cont/ArrayHandle.h index eb0cf4fb9..874e12175 100644 --- a/vtkm/cont/ArrayHandle.h +++ b/vtkm/cont/ArrayHandle.h @@ -360,12 +360,6 @@ public: /// VTKM_CONT vtkm::Id GetNumberOfValues() const; - /// Copies data into the given iterator for the control environment. This - /// method can skip copying into an internally managed control array. - /// - template - VTKM_CONT void CopyInto(IteratorType dest, DeviceAdapterTag) const; - /// \brief Allocates an array large enough to hold the given number of values. /// /// The allocation may be done on an already existing array, but can wipe out diff --git a/vtkm/cont/ArrayHandle.hxx b/vtkm/cont/ArrayHandle.hxx index cbca6bbf6..092ab589c 100644 --- a/vtkm/cont/ArrayHandle.hxx +++ b/vtkm/cont/ArrayHandle.hxx @@ -152,45 +152,6 @@ vtkm::Id ArrayHandle::GetNumberOfValues() const } } -template -template -void ArrayHandle::CopyInto(IteratorType dest, DeviceAdapterTag) const -{ - using pointer_type = typename std::iterator_traits::pointer; - using value_type = typename std::remove_pointer::type; - - static_assert(!std::is_const::value, "CopyInto requires a non const iterator."); - - VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag); - - if (!this->Internals->ControlArrayValid && !this->Internals->ExecutionArrayValid) - { - throw vtkm::cont::ErrorBadValue("ArrayHandle has no data to copy into Iterator."); - } - - if (!this->Internals->ControlArrayValid && - this->Internals->ExecutionArray->IsDeviceAdapter(DeviceAdapterTag())) - { - /// Dynamically cast ArrayHandleExecutionManagerBase into a concrete - /// class and call CopyInto. The dynamic conversion will be sucessful - /// becuase the check to ensure the ExecutionArray is of the type - /// DeviceAdapterTag has already passed - using ConcreteType = - vtkm::cont::internal::ArrayHandleExecutionManager; - ConcreteType* ConcreteExecutionArray = - dynamic_cast(this->Internals->ExecutionArray.get()); - - ConcreteExecutionArray->CopyInto(dest); - } - else - { - PortalConstControl portal = this->GetPortalConstControl(); - std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal), - vtkm::cont::ArrayPortalToIteratorEnd(portal), - dest); - } -} - template void ArrayHandle::Shrink(vtkm::Id numberOfValues) { diff --git a/vtkm/cont/StorageImplicit.h b/vtkm/cont/StorageImplicit.h index b02a8bafe..edf07669a 100644 --- a/vtkm/cont/StorageImplicit.h +++ b/vtkm/cont/StorageImplicit.h @@ -144,17 +144,6 @@ public: throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output."); } - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const - { - using PortalType = typename StorageType::PortalConstType; - PortalType portal = this->Storage->GetPortalConst(); - - std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal), - vtkm::cont::ArrayPortalToIteratorEnd(portal), - dest); - } - VTKM_CONT void Shrink(vtkm::Id vtkmNotUsed(numberOfValues)) { diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h index 1434ff762..a5c359f4f 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h @@ -215,16 +215,6 @@ public: } } - /// Copies the data currently in the device array into the given iterators. - /// Although the iterator is supposed to be from the control environment, - /// thrust can generally handle iterators for a device as well. - /// - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const - { - ::thrust::copy(this->Begin, this->End, dest); - } - /// Resizes the device vector. /// VTKM_CONT void Shrink(vtkm::Id numberOfValues) diff --git a/vtkm/cont/internal/ArrayHandleBasicImpl.h b/vtkm/cont/internal/ArrayHandleBasicImpl.h index 5a5455564..737d8814a 100644 --- a/vtkm/cont/internal/ArrayHandleBasicImpl.h +++ b/vtkm/cont/internal/ArrayHandleBasicImpl.h @@ -177,9 +177,6 @@ public: VTKM_CONT PortalConstControl GetPortalConstControl() const; VTKM_CONT vtkm::Id GetNumberOfValues() const; - template - VTKM_CONT void CopyInto(IteratorType dest, DeviceAdapterTag) const; - VTKM_CONT void Allocate(vtkm::Id numberOfValues); VTKM_CONT void Shrink(vtkm::Id numberOfValues); VTKM_CONT void ReleaseResourcesExecution(); diff --git a/vtkm/cont/internal/ArrayHandleBasicImpl.hxx b/vtkm/cont/internal/ArrayHandleBasicImpl.hxx index 5accbc371..272e03fab 100644 --- a/vtkm/cont/internal/ArrayHandleBasicImpl.hxx +++ b/vtkm/cont/internal/ArrayHandleBasicImpl.hxx @@ -189,43 +189,6 @@ vtkm::Id ArrayHandle::GetNumberOfValues() const } } -template -template -void ArrayHandle::CopyInto(IteratorType dest, DeviceAdapterTag) const -{ - using pointer_type = typename std::iterator_traits::pointer; - using value_type = typename std::remove_pointer::type; - - static_assert(!std::is_const::value, "CopyInto requires a non const iterator."); - - VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag); - - if (!this->Internals->ControlArrayValid && !this->Internals->ExecutionArrayValid) - { - throw vtkm::cont::ErrorBadValue("ArrayHandle has no data to copy into Iterator."); - } - - // If we can copy directly from the execution environment, do so. - DeviceAdapterId devId = DeviceAdapterTraits::GetId(); - if (!this->Internals->ControlArrayValid && std::is_pointer::value && - this->Internals->ExecutionInterface && - this->Internals->ExecutionInterface->GetDeviceId() == devId) - { - vtkm::UInt64 numBytes = static_cast(sizeof(ValueType)) * - static_cast(this->GetNumberOfValues()); - this->Internals->ExecutionInterface->CopyToControl( - this->Internals->ExecutionArray, dest, numBytes); - } - else - { - // Otherwise copy from control. - PortalConstControl portal = this->GetPortalConstControl(); - std::copy(vtkm::cont::ArrayPortalToIteratorBegin(portal), - vtkm::cont::ArrayPortalToIteratorEnd(portal), - dest); - } -} - template void ArrayHandle::Allocate(vtkm::Id numberOfValues) { diff --git a/vtkm/cont/internal/ArrayHandleExecutionManager.h b/vtkm/cont/internal/ArrayHandleExecutionManager.h index beec84853..fad83212b 100644 --- a/vtkm/cont/internal/ArrayHandleExecutionManager.h +++ b/vtkm/cont/internal/ArrayHandleExecutionManager.h @@ -217,12 +217,6 @@ public: { } - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const - { - this->Transfer.CopyInto(dest); - } - protected: VTKM_CONT vtkm::Id GetNumberOfValuesImpl() const { return this->Transfer.GetNumberOfValues(); } diff --git a/vtkm/cont/internal/ArrayManagerExecution.h b/vtkm/cont/internal/ArrayManagerExecution.h index 863757c20..4adc9e1a9 100644 --- a/vtkm/cont/internal/ArrayManagerExecution.h +++ b/vtkm/cont/internal/ArrayManagerExecution.h @@ -125,15 +125,6 @@ public: VTKM_CONT void RetrieveOutputData(vtkm::cont::internal::Storage* storage) const; - /// Similar to RetrieveOutputData except that instead of writing to the - /// controlArray itself, it writes to the given control environment - /// iterator. This allows the user to retrieve data without necessarily - /// allocating an array in the ArrayContainerControl (assuming that control - /// and exeuction have seperate memory spaces). - /// - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const; - /// \brief Reduces the size of the array without changing its values. /// /// This method allows you to resize the array without reallocating it. The diff --git a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h index 3ac5db415..faf0934a3 100644 --- a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h +++ b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h @@ -93,18 +93,6 @@ public: VTKM_ASSERT(storage == this->Storage); } - /// This methods copies data from the execution array into the given - /// iterator. - /// - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const - { - using IteratorType = typename StorageType::PortalConstType::IteratorType; - IteratorType beginIterator = this->Storage->GetPortalConst().GetIteratorBegin(); - - std::copy(beginIterator, beginIterator + this->Storage->GetNumberOfValues(), dest); - } - /// Shrinks the storage. /// VTKM_CONT diff --git a/vtkm/cont/internal/ArrayTransfer.h b/vtkm/cont/internal/ArrayTransfer.h index 42959fd02..1ef226975 100644 --- a/vtkm/cont/internal/ArrayTransfer.h +++ b/vtkm/cont/internal/ArrayTransfer.h @@ -132,18 +132,6 @@ public: this->ArrayManager.RetrieveOutputData(storage); } - /// Similar to RetrieveOutputData except that instead of writing to the - /// controlArray itself, it writes to the given control environment - /// iterator. This allows the user to retrieve data without necessarily - /// allocating an array in the ArrayContainerControl (assuming that control - /// and exeuction have seperate memory spaces). - /// - template - VTKM_CONT void CopyInto(IteratorTypeControl dest) const - { - this->ArrayManager.CopyInto(dest); - } - /// \brief Reduces the size of the array without changing its values. /// /// This method allows you to resize the array without reallocating it. The diff --git a/vtkm/cont/testing/TestingArrayHandles.h b/vtkm/cont/testing/TestingArrayHandles.h index 6074df1ab..939c500ae 100644 --- a/vtkm/cont/testing/TestingArrayHandles.h +++ b/vtkm/cont/testing/TestingArrayHandles.h @@ -220,42 +220,6 @@ private: //can just make sure the allocation didn't throw an exception } - std::cout << "Check CopyInto from control array" << std::endl; - { //Release the execution resources so that data is only - //in the control environment - arrayHandle.ReleaseResourcesExecution(); - - //Copy data from handle into iterator - T array[ARRAY_SIZE]; - arrayHandle.CopyInto(array, DeviceAdapterTag()); - array_handle_testing::CheckValues(array, array + ARRAY_SIZE, T()); - } - - std::cout << "Check CopyInto from execution array" << std::endl; - { //Copy the data to the execution environment - vtkm::cont::ArrayHandle result; - DispatcherPassThrough().Invoke(arrayHandle, result); - - //Copy data from handle into iterator - T array[ARRAY_SIZE]; - result.CopyInto(array, DeviceAdapterTag()); - array_handle_testing::CheckValues(array, array + ARRAY_SIZE, T()); - } - - if (!std::is_same::value) - { - std::cout << "Check using different device adapter" << std::endl; - //Copy the data to the execution environment - vtkm::cont::ArrayHandle result; - DispatcherPassThrough().Invoke(arrayHandle, result); - - //CopyInto allows you to copy the data even - //if you request it from a different device adapter - T array[ARRAY_SIZE]; - result.CopyInto(array, vtkm::cont::DeviceAdapterTagSerial()); - array_handle_testing::CheckValues(array, array + ARRAY_SIZE, T()); - } - { //as output with a length larger than the memory provided by the user //this should fail bool gotException = false; diff --git a/vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu b/vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu index 6d871bfbf..0ce3d8561 100644 --- a/vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu +++ b/vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu @@ -264,9 +264,6 @@ void TestNormalFunctorInvoke() std::cout << "Testing normal worklet invoke." << std::endl; vtkm::Id inputTestValues[3] = { 5, 5, 6 }; - vtkm::Id outputTestValues[3] = { static_cast(0xDEADDEAD), - static_cast(0xDEADDEAD), - static_cast(0xDEADDEAD) }; vtkm::cont::ArrayHandle input = vtkm::cont::make_ArrayHandle(inputTestValues, 3); vtkm::cont::ArrayHandle output; @@ -289,10 +286,8 @@ void TestNormalFunctorInvoke() input.SyncControlArray(); output.SyncControlArray(); - output.CopyInto(outputTestValues, DeviceAdapter()); - VTKM_TEST_ASSERT(inputTestValues[1] == 5, "Input value changed."); - VTKM_TEST_ASSERT(outputTestValues[1] == inputTestValues[1] + 100 + 30, + VTKM_TEST_ASSERT(output.GetPortalConstControl().Get(1) == inputTestValues[1] + 100 + 30, "Output value not set right."); std::cout << " Try return value." << std::endl; @@ -312,10 +307,8 @@ void TestNormalFunctorInvoke() input.SyncControlArray(); output.SyncControlArray(); - output.CopyInto(outputTestValues, DeviceAdapter()); - VTKM_TEST_ASSERT(inputTestValues[2] == 6, "Input value changed."); - VTKM_TEST_ASSERT(outputTestValues[2] == inputTestValues[2] + 200 + 30 * 2, + VTKM_TEST_ASSERT(output.GetPortalConstControl().Get(2) == inputTestValues[2] + 200 + 30 * 2, "Output value not set right."); } diff --git a/vtkm/interop/internal/TransferToOpenGL.h b/vtkm/interop/internal/TransferToOpenGL.h index 6a54b94f4..fa742cc9f 100644 --- a/vtkm/interop/internal/TransferToOpenGL.h +++ b/vtkm/interop/internal/TransferToOpenGL.h @@ -209,9 +209,9 @@ public: //transfer the data. //the primary concern that we have at this point is data locality and - //the type of storage. Our options include using CopyInto and provide a - //temporary location for the values to reside before we give it to openGL - //this works for all storage types. + //the type of storage. Our options include using DeviceAdapterAlgorithm::Copy + //and provide a temporary location for the values to reside before we give it + //to openGL this works for all storage types. // //Second option is to call PrepareForInput and get a PortalConst in the //exection environment. @@ -220,7 +220,8 @@ public: //a unneeded copy. // //The end result is that we have CopyFromHandle which does number two - //from StorageTagBasic, and does the CopyInto for everything else + //from StorageTagBasic, and does the DeviceAdapterAlgorithm::Copy for everything + //else detail::CopyFromHandle(handle, this->State, DeviceAdapterTag()); } diff --git a/vtkm/worklet/testing/UnitTestRemoveUnusedPoints.cxx b/vtkm/worklet/testing/UnitTestRemoveUnusedPoints.cxx index fe59f803d..5b4676dba 100644 --- a/vtkm/worklet/testing/UnitTestRemoveUnusedPoints.cxx +++ b/vtkm/worklet/testing/UnitTestRemoveUnusedPoints.cxx @@ -60,13 +60,12 @@ void CheckOutputCellSet(const vtkm::cont::CellSetExplicit<>& cellSet, VTKM_TEST_ASSERT(pointIds4[2] == 3, "Wrong point id for cell"); VTKM_TEST_ASSERT(pointIds4[3] == 4, "Wrong point id for cell"); - vtkm::Float32 fieldValues[5]; - field.CopyInto(fieldValues, VTKM_DEFAULT_DEVICE_ADAPTER_TAG()); - VTKM_TEST_ASSERT(test_equal(fieldValues[0], TestValue(0, vtkm::Float32())), "Bad field"); - VTKM_TEST_ASSERT(test_equal(fieldValues[1], TestValue(2, vtkm::Float32())), "Bad field"); - VTKM_TEST_ASSERT(test_equal(fieldValues[2], TestValue(4, vtkm::Float32())), "Bad field"); - VTKM_TEST_ASSERT(test_equal(fieldValues[3], TestValue(6, vtkm::Float32())), "Bad field"); - VTKM_TEST_ASSERT(test_equal(fieldValues[4], TestValue(8, vtkm::Float32())), "Bad field"); + auto fieldPortal = field.GetPortalConstControl(); + VTKM_TEST_ASSERT(test_equal(fieldPortal.Get(0), TestValue(0, vtkm::Float32())), "Bad field"); + VTKM_TEST_ASSERT(test_equal(fieldPortal.Get(1), TestValue(2, vtkm::Float32())), "Bad field"); + VTKM_TEST_ASSERT(test_equal(fieldPortal.Get(2), TestValue(4, vtkm::Float32())), "Bad field"); + VTKM_TEST_ASSERT(test_equal(fieldPortal.Get(3), TestValue(6, vtkm::Float32())), "Bad field"); + VTKM_TEST_ASSERT(test_equal(fieldPortal.Get(4), TestValue(8, vtkm::Float32())), "Bad field"); } void RunTest() From 32148cdb564ad3ab3d88ba84b8021d00f8e807ae Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 27 Oct 2017 16:42:09 -0400 Subject: [PATCH 2/2] A first pass at improving the compile time of MarchingCubes --- vtkm/worklet/MarchingCubes.h | 142 +++++++++++++++++++++++++++++++---- 1 file changed, 127 insertions(+), 15 deletions(-) diff --git a/vtkm/worklet/MarchingCubes.h b/vtkm/worklet/MarchingCubes.h index d4796ccc8..9dc669bda 100644 --- a/vtkm/worklet/MarchingCubes.h +++ b/vtkm/worklet/MarchingCubes.h @@ -662,6 +662,38 @@ private: ScatterType Scatter; }; +template +struct GenerateNormalsDeduced +{ + vtkm::cont::ArrayHandle>* normals; + const vtkm::cont::ArrayHandle* field; + const CellSet* cellset; + vtkm::cont::ArrayHandle* edges; + vtkm::cont::ArrayHandle* weights; + + template + void operator()(const CoordinateSystem& coordinates) const + { + // To save memory, the normals computation is done in two passes. In the first + // pass the gradient at the first vertex of each edge is computed and stored in + // the normals array. In the second pass the gradient at the second vertex is + // computed and the gradient of the first vertex is read from the normals array. + // The final normal is interpolated from the two gradient values and stored + // in the normals array. + // + NormalsWorkletPass1 pass1(*edges); + vtkm::worklet::DispatcherMapTopology(pass1).Invoke( + *cellset, *cellset, coordinates, *field, *normals); + + NormalsWorkletPass2 pass2(*edges); + vtkm::worklet::DispatcherMapTopology(pass2).Invoke( + *cellset, *cellset, coordinates, *field, *weights, *normals); + } +}; + template >& normals vtkm::cont::ArrayHandle& edges, vtkm::cont::ArrayHandle& weights) { - // To save memory, the normals computation is done in two passes. In the first - // pass the gradient at the first vertex of each edge is computed and stored in - // the normals array. In the second pass the gradient at the second vertex is - // computed and the gradient of the first vertex is read from the normals array. - // The final normal is interpolated from the two gradient values and stored - // in the normals array. - // - NormalsWorkletPass1 pass1(edges); - vtkm::worklet::DispatcherMapTopology(pass1).Invoke( - cellset, cellset, coordinates, field, normals); + GenerateNormalsDeduced functor; + functor.normals = &normals; + functor.field = &field; + functor.cellset = &cellset; + functor.edges = &edges; + functor.weights = &weights; - NormalsWorkletPass2 pass2(edges); - vtkm::worklet::DispatcherMapTopology(pass2).Invoke( - cellset, cellset, coordinates, field, weights, normals); + + vtkm::cont::CastAndCall(coordinates, functor); } } @@ -738,7 +765,7 @@ public: const DeviceAdapter& device) { vtkm::cont::ArrayHandle> normals; - return this->DoRun( + return this->DeduceRun( isovalues, numIsoValues, cells, coordinateSystem, input, vertices, normals, false, device); } @@ -761,7 +788,7 @@ public: vtkm::cont::ArrayHandle, StorageTagNormals> normals, const DeviceAdapter& device) { - return this->DoRun( + return this->DeduceRun( isovalues, numIsoValues, cells, coordinateSystem, input, vertices, normals, true, device); } @@ -804,6 +831,91 @@ public: void ReleaseCellMapArrays() { this->CellIdMap.ReleaseResources(); } private: + template + struct DeduceCellType + { + MarchingCubes* MC = nullptr; + const ValueType* isovalues = nullptr; + const vtkm::Id* numIsoValues = nullptr; + const CoordinateSystem* coordinateSystem = nullptr; + const vtkm::cont::ArrayHandle* inputField = nullptr; + vtkm::cont::ArrayHandle, StorageTagVertices>* vertices; + vtkm::cont::ArrayHandle, StorageTagNormals>* normals; + const bool* withNormals; + vtkm::cont::CellSetSingleType<>* result; + + template + void operator()(const CellSetType& cells) const + { + if (this->MC) + { + *this->result = this->MC->DoRun(isovalues, + *numIsoValues, + cells, + *coordinateSystem, + *inputField, + *vertices, + *normals, + *withNormals, + DeviceAdapter()); + } + } + }; + + //---------------------------------------------------------------------------- + template + vtkm::cont::CellSetSingleType<> DeduceRun( + const ValueType* isovalues, + const vtkm::Id numIsoValues, + const CellSetType& cells, + const CoordinateSystem& coordinateSystem, + const vtkm::cont::ArrayHandle& inputField, + vtkm::cont::ArrayHandle, StorageTagVertices> vertices, + vtkm::cont::ArrayHandle, StorageTagNormals> normals, + bool withNormals, + const DeviceAdapter&) + { + vtkm::cont::CellSetSingleType<> outputCells("contour"); + + DeduceCellType + functor; + functor.MC = this; + functor.isovalues = isovalues; + functor.numIsoValues = &numIsoValues; + functor.coordinateSystem = &coordinateSystem; + functor.inputField = &inputField; + functor.vertices = &vertices; + functor.normals = &normals; + functor.withNormals = &withNormals; + functor.result = &outputCells; + + vtkm::cont::CastAndCall(cells, functor); + + return outputCells; + } + //---------------------------------------------------------------------------- template