From e095831199ca30a3a1376c8d99bb7e3a1d87de54 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 11 Aug 2017 09:40:09 -0400 Subject: [PATCH 1/2] ArrayHandleReverse now supports shrink/allocate under all conditions. Previously ArrayHandleReverse would only work if it was provided an explicit users array to map too. But this doesn't need to be so, if a user wants to start by constructing an ArrayHandleReverse we should allow that. The side effect of this, is that some very tricky code in the DeviceAdapters can be removed, that explicitly was added to allow output to ArrayHandleReverse --- vtkm/cont/ArrayHandleReverse.h | 61 ++++++------------- .../internal/DeviceAdapterAlgorithmGeneral.h | 3 +- 2 files changed, 21 insertions(+), 43 deletions(-) diff --git a/vtkm/cont/ArrayHandleReverse.h b/vtkm/cont/ArrayHandleReverse.h index b23d7cf1d..0a64495df 100644 --- a/vtkm/cont/ArrayHandleReverse.h +++ b/vtkm/cont/ArrayHandleReverse.h @@ -97,47 +97,34 @@ public: VTKM_CONT Storage() - : valid(false) + : Array() { } VTKM_CONT Storage(const ArrayHandleType& a) - : array(a) - , valid(true){}; + : Array(a) + { + } + VTKM_CONT PortalConstType GetPortalConst() const { - VTKM_ASSERT(this->valid); - return PortalConstType(this->array.GetPortalConstControl()); + return PortalConstType(this->Array.GetPortalConstControl()); } VTKM_CONT - PortalType GetPortal() - { - VTKM_ASSERT(this->valid); - return PortalType(this->array.GetPortalControl()); - } + PortalType GetPortal() { return PortalType(this->Array.GetPortalControl()); } VTKM_CONT - vtkm::Id GetNumberOfValues() const - { - VTKM_ASSERT(this->valid); - return this->array.GetNumberOfValues(); - } + vtkm::Id GetNumberOfValues() const { return this->Array.GetNumberOfValues(); } VTKM_CONT - void Allocate(vtkm::Id vtkmNotUsed(numberOfValues)) - { - throw vtkm::cont::ErrorInternal("ArrayHandleReverse should not be allocated explicitly. "); - } + void Allocate(vtkm::Id numberOfValues) { return this->Array.Allocate(numberOfValues); } VTKM_CONT - void Shrink(vtkm::Id vtkmNotUsed(numberOfValues)) - { - throw vtkm::cont::ErrorBadType("ArrayHandleReverse cannot shrink."); - } + void Shrink(vtkm::Id numberOfValues) { return this->Array.Shrink(numberOfValues); } VTKM_CONT void ReleaseResources() @@ -148,15 +135,10 @@ public: } VTKM_CONT - const ArrayHandleType& GetArray() const - { - VTKM_ASSERT(this->valid); - return this->array; - } + const ArrayHandleType& GetArray() const { return this->Array; } private: - ArrayHandleType array; - bool valid; + ArrayHandleType Array; }; // class storage template @@ -180,29 +162,29 @@ public: VTKM_CONT ArrayTransfer(StorageType* storage) - : array(storage->GetArray()) + : Array(storage->GetArray()) { } VTKM_CONT - vtkm::Id GetNumberOfValues() const { return this->array.GetNumberOfValues(); } + vtkm::Id GetNumberOfValues() const { return this->Array.GetNumberOfValues(); } VTKM_CONT PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) { - return PortalConstExecution(this->array.PrepareForInput(Device())); + return PortalConstExecution(this->Array.PrepareForInput(Device())); } VTKM_CONT PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData)) { - return PortalExecution(this->array.PrepareForInPlace(Device())); + return PortalExecution(this->Array.PrepareForInPlace(Device())); } VTKM_CONT PortalExecution PrepareForOutput(vtkm::Id numberOfValues) { - return PortalExecution(this->array.PrepareForOutput(numberOfValues, Device())); + return PortalExecution(this->Array.PrepareForOutput(numberOfValues, Device())); } VTKM_CONT @@ -212,16 +194,13 @@ public: } VTKM_CONT - void Shrink(vtkm::Id vtkmNotUsed(numberOfValues)) - { - throw vtkm::cont::ErrorBadType("ArrayHandleReverse cannot shrink."); - } + void Shrink(vtkm::Id numberOfValues) { this->Array.Shrink(numberOfValues); } VTKM_CONT - void ReleaseResources() { this->array.ReleaseResourcesExecution(); } + void ReleaseResources() { this->Array.ReleaseResourcesExecution(); } private: - ArrayHandleType array; + ArrayHandleType Array; }; } // namespace internal diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h index 003c110ba..0b5fbe871 100644 --- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h +++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h @@ -749,8 +749,7 @@ public: StencilHandleType stencil; - vtkm::cont::ArrayHandle tempArray; - ValueOutHandleType reducedValues(tempArray); + ValueOutHandleType reducedValues; ZipInHandleType scanInput(values, keystate); ZipOutHandleType scanOutput(reducedValues, stencil); From 89f439999a3c4c774ee39d6e8a1c6fe5720fada0 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 10 Aug 2017 16:13:55 -0400 Subject: [PATCH 2/2] Reduce the amount of typedef statements in DeviceAdapters By using the auto keyword and decltype we can reduce the number of complex typedefs that exist when writing device adapter algorithms. The goal being that it is easier for developers to see the actual algorithms being implemented, by reducing the amount of template 'noise'. --- .../internal/DeviceAdapterAlgorithmThrust.h | 136 ++++--- .../internal/DeviceAdapterAlgorithmGeneral.h | 333 ++++++------------ vtkm/cont/internal/FunctorsGeneral.h | 22 +- .../internal/DeviceAdapterAlgorithmSerial.h | 86 +---- 4 files changed, 195 insertions(+), 382 deletions(-) diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h index 53c29b9a5..f340312ca 100644 --- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h +++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h @@ -216,18 +216,16 @@ private: OutputPortal output, UnaryPredicate unary_predicate) { - typedef typename detail::IteratorTraits::IteratorType IteratorType; + using ValueType = typename StencilPortal::ValueType; - IteratorType outputBegin = IteratorBegin(output); - - typedef typename StencilPortal::ValueType ValueType; + auto outputBegin = IteratorBegin(output); vtkm::exec::cuda::internal::WrappedUnaryPredicate up( unary_predicate); try { - IteratorType newLast = ::thrust::copy_if( + auto newLast = ::thrust::copy_if( thrust::cuda::par, valuesBegin, valuesEnd, IteratorBegin(stencil), outputBegin, up); return static_cast(::thrust::distance(outputBegin, newLast)); } @@ -273,7 +271,7 @@ private: const ValuesPortal& values, const OutputPortal& output) { - typedef typename ValuesPortal::ValueType ValueType; + using ValueType = typename ValuesPortal::ValueType; LowerBoundsPortal(input, values, output, ::thrust::less()); } @@ -281,7 +279,7 @@ private: VTKM_CONT static void LowerBoundsPortal(const InputPortal& input, const OutputPortal& values_output) { - typedef typename InputPortal::ValueType ValueType; + using ValueType = typename InputPortal::ValueType; LowerBoundsPortal(input, values_output, values_output, ::thrust::less()); } @@ -291,7 +289,7 @@ private: const OutputPortal& output, BinaryCompare binary_compare) { - typedef typename InputPortal::ValueType ValueType; + using ValueType = typename InputPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryPredicate bop( binary_compare); @@ -388,17 +386,14 @@ private: const ValueOutputPortal& values_output, BinaryFunctor binary_functor) { - typedef typename detail::IteratorTraits::IteratorType KeysIteratorType; - typedef typename detail::IteratorTraits::IteratorType ValuesIteratorType; + auto keys_out_begin = IteratorBegin(keys_output); + auto values_out_begin = IteratorBegin(values_output); - KeysIteratorType keys_out_begin = IteratorBegin(keys_output); - ValuesIteratorType values_out_begin = IteratorBegin(values_output); - - ::thrust::pair result_iterators; + ::thrust::pair result_iterators; ::thrust::equal_to binaryPredicate; - typedef typename ValuesPortal::ValueType ValueType; + using ValueType = typename ValuesPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bop(binary_functor); try @@ -424,7 +419,7 @@ private: VTKM_CONT static typename InputPortal::ValueType ScanExclusivePortal(const InputPortal& input, const OutputPortal& output) { - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; return ScanExclusivePortal(input, output, @@ -441,7 +436,7 @@ private: { // Use iterator to get value so that thrust device_ptr has chance to handle // data on device. - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; //we have size three so that we can store the origin end value, the //new end value, and the sum of those two @@ -456,13 +451,12 @@ private: vtkm::exec::cuda::internal::WrappedBinaryOperator bop(binaryOp); - typedef typename detail::IteratorTraits::IteratorType IteratorType; - IteratorType end = ::thrust::exclusive_scan(thrust::cuda::par, - IteratorBegin(input), - IteratorEnd(input), - IteratorBegin(output), - initialValue, - bop); + auto end = ::thrust::exclusive_scan(thrust::cuda::par, + IteratorBegin(input), + IteratorEnd(input), + IteratorBegin(output), + initialValue, + bop); //Store the new value for the end of the array. This is done because //with items such as the transpose array it is unsafe to pass the @@ -483,7 +477,7 @@ private: VTKM_CONT static typename InputPortal::ValueType ScanInclusivePortal(const InputPortal& input, const OutputPortal& output) { - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; return ScanInclusivePortal(input, output, ::thrust::plus()); } @@ -492,14 +486,12 @@ private: const OutputPortal& output, BinaryFunctor binary_functor) { - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bop(binary_functor); - typedef typename detail::IteratorTraits::IteratorType IteratorType; - try { - IteratorType end = ::thrust::inclusive_scan( + auto end = ::thrust::inclusive_scan( thrust::cuda::par, IteratorBegin(input), IteratorEnd(input), IteratorBegin(output), bop); return *(end - 1); } @@ -518,7 +510,7 @@ private: const OutputPortal& output) { using KeyType = typename KeysPortal::ValueType; - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; ScanInclusiveByKeyPortal( keys, values, output, ::thrust::equal_to(), ::thrust::plus()); } @@ -534,14 +526,13 @@ private: BinaryPredicate binary_predicate, AssociativeOperator binary_operator) { - typedef typename KeysPortal::ValueType KeyType; + using KeyType = typename KeysPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bpred( binary_predicate); - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bop( binary_operator); - typedef typename detail::IteratorTraits::IteratorType IteratorType; try { ::thrust::inclusive_scan_by_key(thrust::cuda::par, @@ -564,7 +555,7 @@ private: const OutputPortal& output) { using KeyType = typename KeysPortal::ValueType; - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; ScanExclusiveByKeyPortal(keys, values, output, @@ -586,14 +577,12 @@ private: BinaryPredicate binary_predicate, AssociativeOperator binary_operator) { - typedef typename KeysPortal::ValueType KeyType; + using KeyType = typename KeysPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bpred( binary_predicate); - typedef typename OutputPortal::ValueType ValueType; + using ValueType = typename OutputPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryOperator bop( binary_operator); - - typedef typename detail::IteratorTraits::IteratorType IteratorType; try { ::thrust::exclusive_scan_by_key(thrust::cuda::par, @@ -614,14 +603,14 @@ private: template VTKM_CONT static void SortPortal(const ValuesPortal& values) { - typedef typename ValuesPortal::ValueType ValueType; + using ValueType = typename ValuesPortal::ValueType; SortPortal(values, ::thrust::less()); } template VTKM_CONT static void SortPortal(const ValuesPortal& values, BinaryCompare binary_compare) { - typedef typename ValuesPortal::ValueType ValueType; + using ValueType = typename ValuesPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryPredicate bop( binary_compare); try @@ -646,7 +635,7 @@ private: const ValuesPortal& values, BinaryCompare binary_compare) { - typedef typename KeysPortal::ValueType ValueType; + using ValueType = typename KeysPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryPredicate bop( binary_compare); try @@ -663,11 +652,10 @@ private: template VTKM_CONT static vtkm::Id UniquePortal(const ValuesPortal values) { - typedef typename detail::IteratorTraits::IteratorType IteratorType; try { - IteratorType begin = IteratorBegin(values); - IteratorType newLast = ::thrust::unique(thrust::cuda::par, begin, IteratorEnd(values)); + auto begin = IteratorBegin(values); + auto newLast = ::thrust::unique(thrust::cuda::par, begin, IteratorEnd(values)); return static_cast(::thrust::distance(begin, newLast)); } catch (...) @@ -680,15 +668,13 @@ private: template VTKM_CONT static vtkm::Id UniquePortal(const ValuesPortal values, BinaryCompare binary_compare) { - typedef typename detail::IteratorTraits::IteratorType IteratorType; - typedef typename ValuesPortal::ValueType ValueType; - + using ValueType = typename ValuesPortal::ValueType; vtkm::exec::cuda::internal::WrappedBinaryPredicate bop( binary_compare); try { - IteratorType begin = IteratorBegin(values); - IteratorType newLast = ::thrust::unique(thrust::cuda::par, begin, IteratorEnd(values), bop); + auto begin = IteratorBegin(values); + auto newLast = ::thrust::unique(thrust::cuda::par, begin, IteratorEnd(values), bop); return static_cast(::thrust::distance(begin, newLast)); } catch (...) @@ -949,8 +935,8 @@ 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. - input.PrepareForInput(DeviceAdapterTag()); - return ScanExclusivePortal(input.PrepareForInput(DeviceAdapterTag()), + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + return ScanExclusivePortal(inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag())); } @@ -971,8 +957,8 @@ 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. - input.PrepareForInput(DeviceAdapterTag()); - return ScanExclusivePortal(input.PrepareForInput(DeviceAdapterTag()), + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + return ScanExclusivePortal(inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), binary_functor, initialValue); @@ -993,8 +979,8 @@ 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. - input.PrepareForInput(DeviceAdapterTag()); - return ScanInclusivePortal(input.PrepareForInput(DeviceAdapterTag()), + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + return ScanInclusivePortal(inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag())); } @@ -1014,10 +1000,9 @@ 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. - input.PrepareForInput(DeviceAdapterTag()); - return ScanInclusivePortal(input.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), - binary_functor); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + return ScanInclusivePortal( + inputPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), binary_functor); } template @@ -1035,11 +1020,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. - keys.PrepareForInput(DeviceAdapterTag()); - values.PrepareForInput(DeviceAdapterTag()); - ScanInclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(numberOfValues, DeviceAdapterTag())); + auto keysPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + ScanInclusiveByKeyPortal( + keysPortal, valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag())); } template (), binary_functor); @@ -1088,10 +1072,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. - keys.PrepareForInput(DeviceAdapterTag()); - values.PrepareForInput(DeviceAdapterTag()); - ScanExnclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), + auto keysPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + ScanExnclusiveByKeyPortal(keysPortal, + valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), vtkm::TypeTraits::ZeroInitialization(), vtkm::Add()); @@ -1120,10 +1104,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. - keys.PrepareForInput(DeviceAdapterTag()); - values.PrepareForInput(DeviceAdapterTag()); - ScanExclusiveByKeyPortal(keys.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), + auto keysPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + ScanExclusiveByKeyPortal(keysPortal, + valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTag()), initialValue, ::thrust::equal_to(), diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h index 0b5fbe871..dccd801f7 100644 --- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h +++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h @@ -102,16 +102,14 @@ private: template VTKM_CONT static T GetExecutionValue(const vtkm::cont::ArrayHandle& input, vtkm::Id index) { - using InputArrayType = vtkm::cont::ArrayHandle; using OutputArrayType = vtkm::cont::ArrayHandle; OutputArrayType output; + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(1, DeviceAdapterTag()); - CopyKernel::PortalConst, - typename OutputArrayType::template ExecutionTypes::Portal> - kernel(input.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(1, DeviceAdapterTag()), - index); + CopyKernel kernel( + inputPortal, outputPortal, index); DerivedAlgorithm::Schedule(kernel, 1); @@ -125,14 +123,11 @@ public: VTKM_CONT static void Copy(const vtkm::cont::ArrayHandle& input, vtkm::cont::ArrayHandle& output) { - using CopyKernelType = CopyKernel< - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal>; const vtkm::Id inSize = input.GetNumberOfValues(); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(inSize, DeviceAdapterTag()); - CopyKernelType kernel(input.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(inSize, DeviceAdapterTag())); + CopyKernel kernel(inputPortal, outputPortal); DerivedAlgorithm::Schedule(kernel, inSize); } @@ -150,35 +145,23 @@ public: using IndexArrayType = vtkm::cont::ArrayHandle; IndexArrayType indices; - using StencilPortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst; - StencilPortalType stencilPortal = stencil.PrepareForInput(DeviceAdapterTag()); + auto stencilPortal = stencil.PrepareForInput(DeviceAdapterTag()); + auto indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag()); - using IndexPortalType = - typename IndexArrayType::template ExecutionTypes::Portal; - IndexPortalType indexPortal = indices.PrepareForOutput(arrayLength, DeviceAdapterTag()); - - StencilToIndexFlagKernel indexKernel( - stencilPortal, indexPortal, unary_predicate); + StencilToIndexFlagKernel + indexKernel(stencilPortal, indexPortal, unary_predicate); DerivedAlgorithm::Schedule(indexKernel, arrayLength); vtkm::Id outArrayLength = DerivedAlgorithm::ScanExclusive(indices, indices); - using InputPortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst; - InputPortalType inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag()); - using OutputPortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal; - OutputPortalType outputPortal = output.PrepareForOutput(outArrayLength, DeviceAdapterTag()); - - CopyIfKernel copyKernel(inputPortal, stencilPortal, indexPortal, outputPortal, unary_predicate); DerivedAlgorithm::Schedule(copyKernel, arrayLength); @@ -202,11 +185,6 @@ public: vtkm::cont::ArrayHandle& output, vtkm::Id outputIndex = 0) { - using CopyKernel = CopyKernel< - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal>; - const vtkm::Id inSize = input.GetNumberOfValues(); if (inputStartIndex < 0 || numberOfElementsToCopy < 0 || outputIndex < 0 || inputStartIndex >= inSize) @@ -238,10 +216,11 @@ public: } } - CopyKernel kernel(input.PrepareForInput(DeviceAdapterTag()), - output.PrepareForInPlace(DeviceAdapterTag()), - inputStartIndex, - outputIndex); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForInPlace(DeviceAdapterTag()); + + CopyKernel kernel( + inputPortal, outputPortal, inputStartIndex, outputIndex); DerivedAlgorithm::Schedule(kernel, numberOfElementsToCopy); return true; } @@ -255,15 +234,12 @@ public: { vtkm::Id arraySize = values.GetNumberOfValues(); - LowerBoundsKernel::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::Portal> - kernel(input.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(arraySize, DeviceAdapterTag())); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + + LowerBoundsKernel kernel( + inputPortal, valuesPortal, outputPortal); DerivedAlgorithm::Schedule(kernel, arraySize); } @@ -276,18 +252,15 @@ public: { vtkm::Id arraySize = values.GetNumberOfValues(); - LowerBoundsComparisonKernel< - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal, - BinaryCompare> - kernel(input.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(arraySize, DeviceAdapterTag()), - binary_compare); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + + LowerBoundsComparisonKernel + kernel(inputPortal, valuesPortal, outputPortal, binary_compare); DerivedAlgorithm::Schedule(kernel, arraySize); } @@ -322,23 +295,15 @@ 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 - using InputPortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst; - - using ReduceKernelType = ReduceKernel; - - using ReduceHandleType = vtkm::cont::ArrayHandleImplicit; - using TempArrayType = vtkm::cont::ArrayHandle; - - ReduceKernelType kernel( - input.PrepareForInput(DeviceAdapterTag()), initialValue, binary_functor); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + ReduceKernel kernel( + inputPortal, initialValue, binary_functor); vtkm::Id length = (input.GetNumberOfValues() / 16); length += (input.GetNumberOfValues() % 16 == 0) ? 0 : 1; - ReduceHandleType reduced = vtkm::cont::make_ArrayHandleImplicit(kernel, length); + auto reduced = vtkm::cont::make_ArrayHandleImplicit(kernel, length); - TempArrayType inclusiveScanStorage; + vtkm::cont::ArrayHandle inclusiveScanStorage; const U scanResult = DerivedAlgorithm::ScanInclusive(reduced, inclusiveScanStorage, binary_functor); return scanResult; @@ -372,9 +337,8 @@ public: if (block == numBlocks - 1) numberOfInstances = fullSize - blockSize * block; - vtkm::cont::ArrayHandleStreaming> streamIn = - vtkm::cont::ArrayHandleStreaming>( - input, block, blockSize, numberOfInstances); + vtkm::cont::ArrayHandleStreaming> streamIn( + input, block, blockSize, numberOfInstances); if (block == 0) lastResult = DerivedAlgorithm::Reduce(streamIn, initialValue, binary_functor); @@ -415,17 +379,10 @@ public: vtkm::cont::ArrayHandle keystate; { - using InputPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst; - - using KeyStatePortalType = typename vtkm::cont::ArrayHandle< - ReduceKeySeriesStates>::template ExecutionTypes::Portal; - - InputPortalType inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - KeyStatePortalType keyStatePortal = - keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); - ReduceStencilGeneration kernel(inputPortal, - keyStatePortal); + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + ReduceStencilGeneration kernel( + inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); } @@ -436,17 +393,11 @@ public: // the value summed currently, the second being 0 or 1, with 1 being used // when this is a value of a key we need to write ( END or START_AND_END) { - using ValueInHandleType = vtkm::cont::ArrayHandle; - using ValueOutHandleType = vtkm::cont::ArrayHandle; - using StencilHandleType = vtkm::cont::ArrayHandle; - using ZipInHandleType = vtkm::cont::ArrayHandleZip; - using ZipOutHandleType = vtkm::cont::ArrayHandleZip; + vtkm::cont::ArrayHandle stencil; + vtkm::cont::ArrayHandle reducedValues; - StencilHandleType stencil; - ValueOutHandleType reducedValues; - - ZipInHandleType scanInput(values, keystate); - ZipOutHandleType scanOutput(reducedValues, stencil); + auto scanInput = vtkm::cont::make_ArrayHandleZip(values, keystate); + auto scanOutput = vtkm::cont::make_ArrayHandleZip(reducedValues, stencil); DerivedAlgorithm::ScanInclusive( scanInput, scanOutput, ReduceByKeyAdd(binary_functor)); @@ -474,28 +425,20 @@ public: BinaryFunctor binaryFunctor, const T& initialValue) { - using TempArrayType = vtkm::cont::ArrayHandle; - using OutputArrayType = vtkm::cont::ArrayHandle; - - using SrcPortalType = - typename TempArrayType::template ExecutionTypes::PortalConst; - using DestPortalType = - typename OutputArrayType::template ExecutionTypes::Portal; - vtkm::Id numValues = input.GetNumberOfValues(); if (numValues <= 0) { return initialValue; } - TempArrayType inclusiveScan; + vtkm::cont::ArrayHandle inclusiveScan; T result = DerivedAlgorithm::ScanInclusive(input, inclusiveScan, binaryFunctor); - InclusiveToExclusiveKernel inclusiveToExclusive( - inclusiveScan.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(numValues, DeviceAdapterTag()), - binaryFunctor, - initialValue); + auto inputPortal = inclusiveScan.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(numValues, DeviceAdapterTag()); + + InclusiveToExclusiveKernel + inclusiveToExclusive(inputPortal, outputPortal, binaryFunctor, initialValue); DerivedAlgorithm::Schedule(inclusiveToExclusive, numValues); @@ -542,39 +485,22 @@ public: vtkm::cont::ArrayHandle keystate; { - using InputPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst; - - using KeyStatePortalType = typename vtkm::cont::ArrayHandle< - ReduceKeySeriesStates>::template ExecutionTypes::Portal; - - InputPortalType inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - KeyStatePortalType keyStatePortal = - keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); - ReduceStencilGeneration kernel(inputPortal, - keyStatePortal); + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + ReduceStencilGeneration kernel( + inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); } // 2. Shift input and initialize elements at head flags position to initValue - using TempArrayType = typename vtkm::cont::ArrayHandle; - using TempPortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::Portal; - TempArrayType temp; + vtkm::cont::ArrayHandle temp; { - using InputPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst; + auto inputPortal = values.PrepareForInput(DeviceAdapterTag()); + auto keyStatePortal = keystate.PrepareForInput(DeviceAdapterTag()); + auto tempPortal = temp.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); - using KeyStatePortalType = typename vtkm::cont::ArrayHandle< - ReduceKeySeriesStates>::template ExecutionTypes::PortalConst; - - InputPortalType inputPortal = values.PrepareForInput(DeviceAdapterTag()); - KeyStatePortalType keyStatePortal = keystate.PrepareForInput(DeviceAdapterTag()); - TempPortalType tempPortal = temp.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); - - ShiftCopyAndInit kernel( - inputPortal, keyStatePortal, tempPortal, initialValue); + ShiftCopyAndInit + kernel(inputPortal, keyStatePortal, tempPortal, initialValue); DerivedAlgorithm::Schedule(kernel, numberOfKeys); } // 3. Perform a ScanInclusiveByKey @@ -620,13 +546,11 @@ public: if (block == numBlocks - 1) numberOfInstances = fullSize - blockSize * block; - vtkm::cont::ArrayHandleStreaming> streamIn = - vtkm::cont::ArrayHandleStreaming>( - input, block, blockSize, numberOfInstances); + vtkm::cont::ArrayHandleStreaming> streamIn( + input, block, blockSize, numberOfInstances); - vtkm::cont::ArrayHandleStreaming> streamOut = - vtkm::cont::ArrayHandleStreaming>( - output, block, blockSize, numberOfInstances); + vtkm::cont::ArrayHandleStreaming> streamOut( + output, block, blockSize, numberOfInstances); if (block == 0) { @@ -659,11 +583,6 @@ public: vtkm::cont::ArrayHandle& output, BinaryFunctor binary_functor) { - using PortalType = - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal; - - using ScanKernelType = ScanKernel; - DerivedAlgorithm::Copy(input, output); vtkm::Id numValues = output.GetNumberOfValues(); @@ -672,7 +591,9 @@ public: return vtkm::TypeTraits::ZeroInitialization(); } - PortalType portal = output.PrepareForInPlace(DeviceAdapterTag()); + auto portal = output.PrepareForInPlace(DeviceAdapterTag()); + using ScanKernelType = ScanKernel; + vtkm::Id stride; for (stride = 2; stride - 1 < numValues; stride *= 2) @@ -720,17 +641,10 @@ public: vtkm::cont::ArrayHandle keystate; { - using InputPortalType = typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst; - - using KeyStatePortalType = typename vtkm::cont::ArrayHandle< - ReduceKeySeriesStates>::template ExecutionTypes::Portal; - - InputPortalType inputPortal = keys.PrepareForInput(DeviceAdapterTag()); - KeyStatePortalType keyStatePortal = - keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); - ReduceStencilGeneration kernel(inputPortal, - keyStatePortal); + auto inputPortal = keys.PrepareForInput(DeviceAdapterTag()); + auto keyStatePortal = keystate.PrepareForOutput(numberOfKeys, DeviceAdapterTag()); + ReduceStencilGeneration kernel( + inputPortal, keyStatePortal); DerivedAlgorithm::Schedule(kernel, numberOfKeys); } @@ -741,18 +655,10 @@ public: // the value summed currently, the second being 0 or 1, with 1 being used // when this is a value of a key we need to write ( END or START_AND_END) { - using ValueInHandleType = vtkm::cont::ArrayHandle; - using ValueOutHandleType = vtkm::cont::ArrayHandle; - using StencilHandleType = vtkm::cont::ArrayHandle; - using ZipInHandleType = vtkm::cont::ArrayHandleZip; - using ZipOutHandleType = vtkm::cont::ArrayHandleZip; - - StencilHandleType stencil; - - ValueOutHandleType reducedValues; - - ZipInHandleType scanInput(values, keystate); - ZipOutHandleType scanOutput(reducedValues, stencil); + vtkm::cont::ArrayHandle reducedValues; + vtkm::cont::ArrayHandle stencil; + auto scanInput = vtkm::cont::make_ArrayHandleZip(values, keystate); + auto scanOutput = vtkm::cont::make_ArrayHandleZip(reducedValues, stencil); DerivedAlgorithm::ScanInclusive( scanInput, scanOutput, ReduceByKeyAdd(binary_functor)); @@ -768,17 +674,11 @@ public: VTKM_CONT static void Sort(vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { - using ArrayType = typename vtkm::cont::ArrayHandle; - using PortalType = typename ArrayType::template ExecutionTypes::Portal; - vtkm::Id numValues = values.GetNumberOfValues(); if (numValues < 2) { return; } - - PortalType portal = values.PrepareForInPlace(DeviceAdapterTag()); - vtkm::Id numThreads = 1; while (numThreads < numValues) { @@ -786,8 +686,9 @@ public: } numThreads /= 2; - using MergeKernel = BitonicSortMergeKernel; - using CrossoverKernel = BitonicSortCrossoverKernel; + auto portal = values.PrepareForInPlace(DeviceAdapterTag()); + using MergeKernel = BitonicSortMergeKernel; + using CrossoverKernel = BitonicSortCrossoverKernel; for (vtkm::Id crossoverSize = 1; crossoverSize < numValues; crossoverSize *= 2) { @@ -816,12 +717,7 @@ public: //combine the keys and values into a ZipArrayHandle //we than need to specify a custom compare function wrapper //that only checks for key side of the pair, using a custom compare functor. - using KeyType = vtkm::cont::ArrayHandle; - ; - using ValueType = vtkm::cont::ArrayHandle; - using ZipHandleType = vtkm::cont::ArrayHandleZip; - - ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); + auto zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); DerivedAlgorithm::Sort(zipHandle, internal::KeyCompare()); } @@ -834,12 +730,7 @@ public: //we than need to specify a custom compare function wrapper //that only checks for key side of the pair, using the custom compare //functor that the user passed in - using KeyType = vtkm::cont::ArrayHandle; - ; - using ValueType = vtkm::cont::ArrayHandle; - using ZipHandleType = vtkm::cont::ArrayHandleZip; - - ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); + auto zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); DerivedAlgorithm::Sort(zipHandle, internal::KeyCompare(binary_compare)); } @@ -861,15 +752,10 @@ public: using WrappedBOpType = internal::WrappedBinaryOperator; WrappedBOpType wrappedCompare(binary_compare); - ClassifyUniqueComparisonKernel< - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle:: - template ExecutionTypes::Portal, - WrappedBOpType> - classifyKernel(values.PrepareForInput(DeviceAdapterTag()), - stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag()), - wrappedCompare); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + auto stencilPortal = stencilArray.PrepareForOutput(inputSize, DeviceAdapterTag()); + ClassifyUniqueComparisonKernel + classifyKernel(valuesPortal, stencilPortal, wrappedCompare); DerivedAlgorithm::Schedule(classifyKernel, inputSize); @@ -890,16 +776,12 @@ public: { vtkm::Id arraySize = values.GetNumberOfValues(); - UpperBoundsKernel::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::Portal> - kernel(input.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(arraySize, DeviceAdapterTag())); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + UpperBoundsKernel kernel( + inputPortal, valuesPortal, outputPortal); DerivedAlgorithm::Schedule(kernel, arraySize); } @@ -911,18 +793,15 @@ public: { vtkm::Id arraySize = values.GetNumberOfValues(); - UpperBoundsKernelComparisonKernel< - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes< - DeviceAdapterTag>::PortalConst, - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal, - BinaryCompare> - kernel(input.PrepareForInput(DeviceAdapterTag()), - values.PrepareForInput(DeviceAdapterTag()), - output.PrepareForOutput(arraySize, DeviceAdapterTag()), - binary_compare); + auto inputPortal = input.PrepareForInput(DeviceAdapterTag()); + auto valuesPortal = values.PrepareForInput(DeviceAdapterTag()); + auto outputPortal = output.PrepareForOutput(arraySize, DeviceAdapterTag()); + + UpperBoundsKernelComparisonKernel + kernel(inputPortal, valuesPortal, outputPortal, binary_compare); DerivedAlgorithm::Schedule(kernel, arraySize); } diff --git a/vtkm/cont/internal/FunctorsGeneral.h b/vtkm/cont/internal/FunctorsGeneral.h index 0ffe52050..fa596b4eb 100644 --- a/vtkm/cont/internal/FunctorsGeneral.h +++ b/vtkm/cont/internal/FunctorsGeneral.h @@ -396,7 +396,7 @@ struct LowerBoundsKernel using InputIteratorsType = vtkm::cont::ArrayPortalToIterators; InputIteratorsType inputIterators(this->InputPortal); - typename InputIteratorsType::IteratorType resultPos = std::lower_bound( + auto resultPos = std::lower_bound( inputIterators.GetBegin(), inputIterators.GetEnd(), this->ValuesPortal.Get(index)); vtkm::Id resultIndex = @@ -444,11 +444,10 @@ struct LowerBoundsComparisonKernel using InputIteratorsType = vtkm::cont::ArrayPortalToIterators; InputIteratorsType inputIterators(this->InputPortal); - typename InputIteratorsType::IteratorType resultPos = - std::lower_bound(inputIterators.GetBegin(), - inputIterators.GetEnd(), - this->ValuesPortal.Get(index), - this->CompareFunctor); + auto resultPos = std::lower_bound(inputIterators.GetBegin(), + inputIterators.GetEnd(), + this->ValuesPortal.Get(index), + this->CompareFunctor); vtkm::Id resultIndex = static_cast(std::distance(inputIterators.GetBegin(), resultPos)); @@ -748,7 +747,7 @@ struct UpperBoundsKernel using InputIteratorsType = vtkm::cont::ArrayPortalToIterators; InputIteratorsType inputIterators(this->InputPortal); - typename InputIteratorsType::IteratorType resultPos = std::upper_bound( + auto resultPos = std::upper_bound( inputIterators.GetBegin(), inputIterators.GetEnd(), this->ValuesPortal.Get(index)); vtkm::Id resultIndex = @@ -796,11 +795,10 @@ struct UpperBoundsKernelComparisonKernel using InputIteratorsType = vtkm::cont::ArrayPortalToIterators; InputIteratorsType inputIterators(this->InputPortal); - typename InputIteratorsType::IteratorType resultPos = - std::upper_bound(inputIterators.GetBegin(), - inputIterators.GetEnd(), - this->ValuesPortal.Get(index), - this->CompareFunctor); + auto resultPos = std::upper_bound(inputIterators.GetBegin(), + inputIterators.GetEnd(), + this->ValuesPortal.Get(index), + this->CompareFunctor); vtkm::Id resultIndex = static_cast(std::distance(inputIterators.GetBegin(), resultPos)); diff --git a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h index 42a4eb25a..8d5e14b96 100644 --- a/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h +++ b/vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h @@ -61,11 +61,8 @@ public: U initialValue, BinaryFunctor binary_functor) { - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalIn; - internal::WrappedBinaryOperator wrappedOp(binary_functor); - PortalIn inputPortal = input.PrepareForInput(Device()); + auto inputPortal = input.PrepareForInput(Device()); return std::accumulate(vtkm::cont::ArrayPortalToIteratorBegin(inputPortal), vtkm::cont::ArrayPortalToIteratorEnd(inputPortal), initialValue, @@ -85,22 +82,12 @@ public: vtkm::cont::ArrayHandle& values_output, BinaryFunctor binary_functor) { - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalKIn; - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalVIn; - - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal PortalKOut; - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal PortalVOut; - - PortalKIn keysPortalIn = keys.PrepareForInput(Device()); - PortalVIn valuesPortalIn = values.PrepareForInput(Device()); + auto keysPortalIn = keys.PrepareForInput(Device()); + auto valuesPortalIn = values.PrepareForInput(Device()); const vtkm::Id numberOfKeys = keys.GetNumberOfValues(); - PortalKOut keysPortalOut = keys_output.PrepareForOutput(numberOfKeys, Device()); - PortalVOut valuesPortalOut = values_output.PrepareForOutput(numberOfKeys, Device()); + auto keysPortalOut = keys_output.PrepareForOutput(numberOfKeys, Device()); + auto valuesPortalOut = values_output.PrepareForOutput(numberOfKeys, Device()); vtkm::Id writePos = 0; vtkm::Id readPos = 0; @@ -141,15 +128,10 @@ public: VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle& input, vtkm::cont::ArrayHandle& output) { - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal PortalOut; - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalIn; - vtkm::Id numberOfValues = input.GetNumberOfValues(); - PortalIn inputPortal = input.PrepareForInput(Device()); - PortalOut outputPortal = output.PrepareForOutput(numberOfValues, Device()); + auto inputPortal = input.PrepareForInput(Device()); + auto outputPortal = output.PrepareForOutput(numberOfValues, Device()); if (numberOfValues <= 0) { @@ -169,17 +151,12 @@ public: vtkm::cont::ArrayHandle& output, BinaryFunctor binary_functor) { - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal PortalOut; - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalIn; - internal::WrappedBinaryOperator wrappedBinaryOp(binary_functor); vtkm::Id numberOfValues = input.GetNumberOfValues(); - PortalIn inputPortal = input.PrepareForInput(Device()); - PortalOut outputPortal = output.PrepareForOutput(numberOfValues, Device()); + auto inputPortal = input.PrepareForInput(Device()); + auto outputPortal = output.PrepareForOutput(numberOfValues, Device()); if (numberOfValues <= 0) { @@ -201,17 +178,12 @@ public: BinaryFunctor binaryFunctor, const T& initialValue) { - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal PortalOut; - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalIn; - internal::WrappedBinaryOperator wrappedBinaryOp(binaryFunctor); vtkm::Id numberOfValues = input.GetNumberOfValues(); - PortalIn inputPortal = input.PrepareForInput(Device()); - PortalOut outputPortal = output.PrepareForOutput(numberOfValues, Device()); + auto inputPortal = input.PrepareForInput(Device()); + auto outputPortal = output.PrepareForOutput(numberOfValues, Device()); if (numberOfValues <= 0) { @@ -277,21 +249,12 @@ private: vtkm::cont::ArrayHandle& index, vtkm::cont::ArrayHandle& values_out) { - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes< - Device>::PortalConst PortalVIn; - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::PortalConst - PortalI; - typedef - typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal - PortalVout; - const vtkm::Id n = values.GetNumberOfValues(); VTKM_ASSERT(n == index.GetNumberOfValues()); - PortalVIn valuesPortal = values.PrepareForInput(Device()); - PortalI indexPortal = index.PrepareForInput(Device()); - PortalVout valuesOutPortal = values_out.PrepareForOutput(n, Device()); + auto valuesPortal = values.PrepareForInput(Device()); + auto indexPortal = index.PrepareForInput(Device()); + auto valuesOutPortal = values_out.PrepareForOutput(n, Device()); for (vtkm::Id i = 0; i < n; i++) { @@ -310,12 +273,7 @@ private: //we than need to specify a custom compare function wrapper //that only checks for key side of the pair, using the custom compare //functor that the user passed in - using KeyType = vtkm::cont::ArrayHandle; - ; - using ValueType = vtkm::cont::ArrayHandle; - typedef vtkm::cont::ArrayHandleZip ZipHandleType; - - ZipHandleType zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); + auto zipHandle = vtkm::cont::make_ArrayHandleZip(keys, values); Sort(zipHandle, internal::KeyCompare(binary_compare)); } @@ -337,11 +295,8 @@ public: { /// More efficient sort: /// Move value indexes when sorting and reorder the value array at last - using ValueType = vtkm::cont::ArrayHandle; - using IndexType = vtkm::cont::ArrayHandle; - - IndexType indexArray; - ValueType valuesScattered; + vtkm::cont::ArrayHandle indexArray; + vtkm::cont::ArrayHandle valuesScattered; Copy(ArrayHandleIndex(keys.GetNumberOfValues()), indexArray); SortByKeyDirect(keys, indexArray, wrappedCompare); @@ -364,11 +319,8 @@ public: VTKM_CONT static void Sort(vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { - typedef typename vtkm::cont::ArrayHandle::template ExecutionTypes::Portal - PortalType; - - PortalType arrayPortal = values.PrepareForInPlace(Device()); - vtkm::cont::ArrayPortalToIterators iterators(arrayPortal); + auto arrayPortal = values.PrepareForInPlace(Device()); + vtkm::cont::ArrayPortalToIterators iterators(arrayPortal); internal::WrappedBinaryOperator wrappedCompare(binary_compare); std::sort(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare);