Merge branch 'sort_inplace_zip_arrays' into 'master'

Allow in-place sorting of ArrayHandleZip

See merge request !29

vtkm/TypeTraits.h

@@ -21,6 +21,7 @@
 #define vtk_m_TypeTraits_h
 
 #include <vtkm/Types.h>
+#include <vtkm/Pair.h>
 
 namespace vtkm {
 
@@ -120,6 +121,15 @@ struct TypeTraits<vtkm::Vec<T,Size> >
   typedef TypeTraitsVectorTag DimensionalityTag;
 };
 
+/// Traits for Pair types.
+///
+template<typename T, typename U>
+struct TypeTraits<vtkm::Pair<T,U> >
+{
+  typedef typename vtkm::TypeTraits<T>::NumericTag NumericTag;
+  typedef TypeTraitsScalarTag DimensionalityTag;
+};
+
 } // namespace vtkm
 
 #endif //vtk_m_TypeTraits_h

vtkm/cont/ArrayHandleZip.h

@@ -30,13 +30,16 @@ namespace internal {
 /// \brief An array portal that zips two portals together into a single value
 /// for the execution environment
 template<typename ValueType_,
-         typename PortalTypeFirst,
-         typename PortalTypeSecond>
+         typename PortalTypeFirst_,
+         typename PortalTypeSecond_>
 class ArrayPortalExecZip
 {
 public:
   typedef ValueType_ ValueType;
   typedef ValueType_ IteratorType;
+  typedef PortalTypeFirst_ PortalTypeFirst;
+  typedef PortalTypeSecond_ PortalTypeSecond;
+
 
   VTKM_CONT_EXPORT
   ArrayPortalExecZip(const PortalTypeFirst  &portalfirst = PortalTypeFirst(),

vtkm/cont/cuda/internal/MakeThrustIterator.h

@@ -21,6 +21,7 @@
 #define vtk_m_cont_cuda_internal_MakeThrustIterator_h
 
 #include <vtkm/Types.h>
+#include <vtkm/Pair.h>
 #include <vtkm/internal/ExportMacros.h>
 
 #include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
@@ -37,35 +38,81 @@
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
 
 #if defined(__GNUC__) || defined(____clang__)
 #pragma GCC diagnostic pop
 #endif // gcc || clang
 
+
+//needed forward declares to get zip handle to work properly
+namespace vtkm {
+namespace exec {
+namespace internal {
+
+template<typename ValueType_,
+         typename PortalTypeFirst,
+         typename PortalTypeSecond>
+class ArrayPortalExecZip;
+
+}
+}
+}
+
+//more forward declares needed to get zip handle to work properly
 namespace vtkm {
 namespace cont {
 namespace cuda {
 namespace internal {
 
+namespace detail { template<class PortalType> struct IteratorTraits; }
+
+//forward declare IteratorBegin
+template<class PortalType>
+VTKM_CONT_EXPORT
+typename detail::IteratorTraits<PortalType>::IteratorType
+IteratorBegin(PortalType portal);
+}
+}
+}
+}
+
+
+namespace vtkm {
+namespace cont {
+namespace cuda {
+namespace internal {
 namespace detail {
 
 // Tags to specify what type of thrust iterator to use.
 struct ThrustIteratorTransformTag {  };
+struct ThrustIteratorZipTag {  };
 struct ThrustIteratorDevicePtrTag {  };
 
 // Traits to help classify what thrust iterators will be used.
-template<class IteratorType>
+template<class PortalType, class IteratorType>
 struct ThrustIteratorTag {
   typedef ThrustIteratorTransformTag Type;
 };
-template<typename T>
-struct ThrustIteratorTag<T *> {
+template<typename PortalType, typename T>
+struct ThrustIteratorTag<PortalType, T *> {
   typedef ThrustIteratorDevicePtrTag Type;
 };
-template<typename T>
-struct ThrustIteratorTag<const T*> {
+template<typename PortalType, typename T>
+struct ThrustIteratorTag<PortalType, const T*> {
   typedef ThrustIteratorDevicePtrTag Type;
 };
+template<typename T, typename U, typename V>
+struct ThrustIteratorTag< vtkm::exec::internal::ArrayPortalExecZip< T, U, V >,
+                          T > {
+  //this is a real special case. ExecZip and PortalValue don't combine
+  //well together, when used with DeviceAlgorithm that has a custom operator
+  //the custom operator is actually passed the PortalValue instead of
+  //the real values, and by that point we can't fix anything since we
+  //don't know what the original operator is
+  typedef ThrustIteratorZipTag Type;
+};
+
 
 template<typename T> struct ThrustStripPointer;
 template<typename T> struct ThrustStripPointer<T *> {
@@ -80,10 +127,19 @@ template<class PortalType>
 struct PortalValue {
   typedef typename PortalType::ValueType ValueType;
 
+  VTKM_EXEC_EXPORT
+  PortalValue()
+    : Portal(),
+      Index(0) {  }
+
   VTKM_EXEC_EXPORT
   PortalValue(const PortalType &portal, vtkm::Id index)
     : Portal(portal), Index(index) {  }
 
+  VTKM_EXEC_EXPORT
+  PortalValue(const PortalValue<PortalType> &other)
+    : Portal(other.Portal), Index(other.Index) {  }
+
   VTKM_EXEC_EXPORT
   ValueType operator=(ValueType value) {
     this->Portal.Set(this->Index, value);
@@ -105,7 +161,10 @@ class LookupFunctor
                                         PortalValue<PortalType>  >
 {
   public:
-    VTKM_CONT_EXPORT LookupFunctor(PortalType portal)
+    VTKM_EXEC_EXPORT LookupFunctor()
+      : Portal() {  }
+
+    VTKM_EXEC_EXPORT LookupFunctor(PortalType portal)
       : Portal(portal) {  }
 
     VTKM_EXEC_EXPORT
@@ -127,6 +186,34 @@ struct IteratorChooser<PortalType, detail::ThrustIteratorTransformTag> {
       ::thrust::counting_iterator<vtkm::Id> > Type;
 };
 template<class PortalType>
+struct IteratorChooser<PortalType, detail::ThrustIteratorZipTag> {
+
+  //this is a real special case. ExecZip and PortalValue don't combine
+  //well together, when used with DeviceAlgorithm that has a custom operator
+  //the custom operator is actually passed the PortalValue instead of
+  //the real values, and by that point we can't fix anything since we
+  //don't know what the original operator is.
+
+  //So to fix this issue we wrap the original array portals into a thrust
+  //zip iterator and let handle everything
+  typedef typename PortalType::PortalTypeFirst PortalTypeFirst;
+  typedef typename detail::ThrustIteratorTag<
+                          PortalTypeFirst,
+                          typename PortalTypeFirst::IteratorType>::Type FirstTag;
+  typedef typename IteratorChooser<PortalTypeFirst, FirstTag>::Type FirstIterType;
+
+  typedef typename PortalType::PortalTypeSecond PortalTypeSecond;
+  typedef typename detail::ThrustIteratorTag<
+                          PortalTypeSecond,
+                          typename PortalTypeSecond::IteratorType>::Type SecondTag;
+  typedef typename IteratorChooser<PortalTypeSecond, SecondTag>::Type SecondIterType;
+
+  //Now that we have deduced the concrete types of the first and second
+  //array portals of the zip we can construct a zip iterator for those
+  typedef ::thrust::tuple<FirstIterType, SecondIterType> IteratorTuple;
+  typedef ::thrust::zip_iterator<IteratorTuple> Type;
+};
+template<class PortalType>
 struct IteratorChooser<PortalType, detail::ThrustIteratorDevicePtrTag> {
   typedef ::thrust::cuda::pointer<
       typename detail::ThrustStripPointer<
@@ -136,29 +223,48 @@ struct IteratorChooser<PortalType, detail::ThrustIteratorDevicePtrTag> {
 template<class PortalType>
 struct IteratorTraits
 {
-  typedef typename PortalType::IteratorType BaseIteratorType;
-  typedef typename detail::ThrustIteratorTag<BaseIteratorType>::Type Tag;
+  typedef typename detail::ThrustIteratorTag<
+                          PortalType,
+                          typename PortalType::IteratorType>::Type Tag;
   typedef typename IteratorChooser<PortalType, Tag>::Type IteratorType;
 };
 
 
 template<typename T>
-VTKM_CONT_EXPORT static
+VTKM_CONT_EXPORT
 ::thrust::cuda::pointer<T>
 MakeDevicePtr(T *iter)
 {
   return::thrust::cuda::pointer<T>(iter);
 }
 template<typename T>
-VTKM_CONT_EXPORT static
+VTKM_CONT_EXPORT
 ::thrust::cuda::pointer<const T>
 MakeDevicePtr(const T *iter)
 {
   return ::thrust::cuda::pointer<const T>(iter);
 }
 
+template<typename T, typename U>
+VTKM_CONT_EXPORT
+::thrust::zip_iterator<
+  ::thrust::tuple<typename IteratorTraits<T>::IteratorType,
+                  typename IteratorTraits<U>::IteratorType
+                  >
+  >
+MakeZipIterator(const T& t, const U& u)
+{
+  //todo deduce from T and U the iterator types
+  typedef typename IteratorTraits<T>::IteratorType FirstIterType;
+  typedef typename IteratorTraits<U>::IteratorType SecondIterType;
+  return ::thrust::make_zip_iterator(
+          ::thrust::make_tuple( vtkm::cont::cuda::internal::IteratorBegin(t),
+                                vtkm::cont::cuda::internal::IteratorBegin(u) )
+          );
+}
+
 template<class PortalType>
-VTKM_CONT_EXPORT static
+VTKM_CONT_EXPORT
 typename IteratorTraits<PortalType>::IteratorType
 MakeIteratorBegin(PortalType portal, detail::ThrustIteratorTransformTag)
 {
@@ -168,7 +274,17 @@ MakeIteratorBegin(PortalType portal, detail::ThrustIteratorTransformTag)
 }
 
 template<class PortalType>
-VTKM_CONT_EXPORT static
+VTKM_CONT_EXPORT
+typename IteratorTraits<PortalType>::IteratorType
+MakeIteratorBegin(PortalType portal, detail::ThrustIteratorZipTag)
+{
+    return MakeZipIterator(portal.GetFirstPortal(),
+                           portal.GetSecondPortal()
+                           );
+}
+
+template<class PortalType>
+VTKM_CONT_EXPORT
 typename IteratorTraits<PortalType>::IteratorType
 MakeIteratorBegin(PortalType portal, detail::ThrustIteratorDevicePtrTag)
 {
@@ -199,35 +315,8 @@ IteratorEnd(PortalType portal)
 }
 }
 }
+
+
 } //namespace vtkm::cont::cuda::internal
 
-namespace thrust {
-
-template< typename PortalType >
-struct less< vtkm::cont::cuda::internal::detail::PortalValue< PortalType > > :
-        public binary_function<
-          vtkm::cont::cuda::internal::detail::PortalValue< PortalType >,
-          vtkm::cont::cuda::internal::detail::PortalValue< PortalType >,
-          bool>
-{
-  typedef vtkm::cont::cuda::internal::detail::PortalValue< PortalType > T;
-  typedef typename vtkm::cont::cuda::internal::detail::PortalValue<
-                        PortalType >::ValueType ValueType;
-
-
-  /*! Function call operator. The return value is <tt>lhs < rhs</tt>.
-   */
-  __host__ __device__ bool operator()(const T &lhs, const T &rhs) const
-  {return (ValueType)lhs < (ValueType)rhs;}
-
-  /*! Function call operator. The return value is <tt>lhs < rhs</tt>.
-      specially designed to work with vtkm portal values, which can
-      be compared to their underline type
-   */
-  __host__ __device__ bool operator()(const T &lhs,
-                                      const ValueType &rhs) const
-  {return (ValueType)lhs < rhs;}
-}; // end less
-
-}
 #endif

vtkm/cont/internal/DeviceAdapterAlgorithmSerial.h

@@ -21,6 +21,7 @@
 #define vtk_m_cont_internal_DeviceAdapterAlgorithmSerial_h
 
 #include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/ArrayHandleZip.h>
 #include <vtkm/cont/ArrayPortalToIterators.h>
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
 #include <vtkm/cont/ErrorExecution.h>
@@ -395,6 +396,23 @@ public:
 
     PortalType arrayPortal = values.PrepareForInPlace(Device());
     vtkm::cont::ArrayPortalToIterators<PortalType> iterators(arrayPortal);
+
+    std::sort(iterators.GetBegin(), iterators.GetEnd());
+  }
+
+  template<typename T, typename U>
+  VTKM_CONT_EXPORT static void Sort(vtkm::cont::ArrayHandleZip<T,U>& values)
+  {
+    typedef typename vtkm::cont::ArrayHandleZip<T,U>
+        ::template ExecutionTypes<Device>::Portal PortalType;
+
+    PortalType arrayPortal = values.PrepareForInPlace(Device());
+    vtkm::cont::ArrayPortalToIterators<PortalType> iterators(arrayPortal);
+
+    //this is required to get sort to work with zip handles
+    typedef vtkm::cont::internal::ArrayHandleZipTraits< T, U > ZipTraits;
+    typedef std::less< typename ZipTraits::ValueType > LessOp;
+    internal::WrappedBinaryOperator<bool, LessOp> wrappedCompare( (LessOp()) );
     std::sort(iterators.GetBegin(), iterators.GetEnd());
   }
 
@@ -407,7 +425,10 @@ public:
 
     PortalType arrayPortal = values.PrepareForInPlace(Device());
     vtkm::cont::ArrayPortalToIterators<PortalType> iterators(arrayPortal);
-    std::sort(iterators.GetBegin(), iterators.GetEnd(), comp);
+
+
+    internal::WrappedBinaryOperator<bool,Compare> wrappedCompare(comp);
+    std::sort(iterators.GetBegin(), iterators.GetEnd(), wrappedCompare);
   }
 
   VTKM_CONT_EXPORT static void Synchronize()

vtkm/cont/internal/IteratorFromArrayPortal.h

@@ -66,6 +66,13 @@ struct IteratorFromArrayPortalValue
     return value;
   }
 
+  VTKM_CONT_EXPORT
+  bool operator<(const ValueType& value) const
+  {
+    return this->Portal.Get(this->Index) < value;
+  }
+
+
   VTKM_CONT_EXPORT
   operator ValueType(void) const
   {
@@ -100,7 +107,7 @@ public:
   operator[](std::size_t idx) const
   {
     return detail::IteratorFromArrayPortalValue<ArrayPortalType>(this->Portal,
-           static_cast<vtkm::Id>(idx) );
+           this->Index + static_cast<vtkm::Id>(idx) );
   }
 
 private:

vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h

@@ -27,6 +27,7 @@
 #include <vtkm/exec/internal/ErrorMessageBuffer.h>
 #include <vtkm/Extent.h>
 #include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/ArrayHandleZip.h>
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
 #include <vtkm/cont/ErrorExecution.h>
 #include <vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h>
@@ -506,6 +507,24 @@ public:
     ::tbb::parallel_sort(iterators.GetBegin(), iterators.GetEnd());
   }
 
+  template<typename T, typename U>
+  VTKM_CONT_EXPORT static void Sort(vtkm::cont::ArrayHandleZip<T,U>& values)
+  {
+    typedef typename vtkm::cont::ArrayHandleZip<T,U>::template
+      ExecutionTypes<vtkm::cont::DeviceAdapterTagTBB>::Portal PortalType;
+    PortalType arrayPortal = values.PrepareForInPlace(
+      vtkm::cont::DeviceAdapterTagTBB());
+
+    typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
+    IteratorsType iterators(arrayPortal);
+
+    //this is required to get sort to work with zip handles
+    typedef vtkm::cont::internal::ArrayHandleZipTraits< T, U > ZipTraits;
+    typedef std::less< typename ZipTraits::ValueType > LessOp;
+    internal::WrappedBinaryOperator<bool, LessOp> wrappedCompare( (LessOp()) );
+    std::sort(iterators.GetBegin(), iterators.GetEnd());
+  }
+
   template<typename T, class Container, class Compare>
   VTKM_CONT_EXPORT static void Sort(
       vtkm::cont::ArrayHandle<T,Container> &values, Compare comp)
@@ -518,7 +537,10 @@ public:
     typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
     IteratorsType iterators(arrayPortal);
 
-    ::tbb::parallel_sort(iterators.GetBegin(), iterators.GetEnd(), comp);
+    internal::WrappedBinaryOperator<bool,Compare> wrappedCompare(comp);
+    ::tbb::parallel_sort(iterators.GetBegin(),
+                         iterators.GetEnd(),
+                         wrappedCompare);
   }
 
 

vtkm/cont/testing/TestingDeviceAdapter.h

@@ -20,7 +20,9 @@
 #ifndef vtk_m_cont_testing_TestingDeviceAdapter_h
 #define vtk_m_cont_testing_TestingDeviceAdapter_h
 
+#include <vtkm/TypeTraits.h>
 #include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/ArrayHandleZip.h>
 #include <vtkm/cont/ArrayPortalToIterators.h>
 #include <vtkm/cont/ErrorControlOutOfMemory.h>
 #include <vtkm/cont/ErrorExecution.h>
@@ -52,21 +54,15 @@ namespace testing {
 namespace comparison {
 struct SortLess
 {
-  template<typename T>
-  VTKM_EXEC_CONT_EXPORT bool operator()(const T& a,const T& b) const
-  {
-    typedef typename vtkm::TypeTraits<T>::DimensionalityTag Dimensionality;
-    return this->compare(a,b,Dimensionality());
-  }
-  template<typename T>
-  VTKM_EXEC_CONT_EXPORT bool compare(const T& a,const T& b,
-                                     vtkm::TypeTraitsScalarTag) const
+  template<typename T, typename U>
+  VTKM_EXEC_CONT_EXPORT bool operator()(const T& a, const U& b) const
   {
     return a < b;
   }
-  template<typename T>
-  VTKM_EXEC_CONT_EXPORT bool compare(const T& a,const T& b,
-                                     vtkm::TypeTraitsVectorTag) const
+
+  template<typename T, int N>
+  VTKM_EXEC_EXPORT bool operator()(const vtkm::Vec<T,N>& a,
+                                   const vtkm::Vec<T,N>& b) const
   {
     const vtkm::IdComponent SIZE = vtkm::VecTraits<T>::NUM_COMPONENTS;
     bool valid = true;
@@ -80,21 +76,15 @@ struct SortLess
 
 struct SortGreater
 {
-  template<typename T>
-  VTKM_EXEC_CONT_EXPORT bool operator()(const T& a,const T& b) const
-  {
-    typedef typename vtkm::TypeTraits<T>::DimensionalityTag Dimensionality;
-    return this->compare(a,b,Dimensionality());
-  }
-  template<typename T>
-  VTKM_EXEC_EXPORT bool compare(const T& a,const T& b,
-                                     vtkm::TypeTraitsScalarTag) const
+  template<typename T, typename U>
+  VTKM_EXEC_CONT_EXPORT bool operator()(const T& a, const U& b) const
   {
     return a > b;
   }
-  template<typename T>
-  VTKM_EXEC_EXPORT bool compare(const T& a,const T& b,
-                                     vtkm::TypeTraitsVectorTag) const
+
+  template<typename T, int N>
+  VTKM_EXEC_EXPORT bool operator()(const vtkm::Vec<T,N>& a,
+                                   const vtkm::Vec<T,N>& b) const
   {
     const vtkm::IdComponent SIZE = vtkm::VecTraits<T>::NUM_COMPONENTS;
     bool valid = true;
@@ -119,7 +109,7 @@ struct MaxValue
 
 
 #define ERROR_MESSAGE "Got an error."
-#define ARRAY_SIZE 500
+#define ARRAY_SIZE 1000
 #define OFFSET 1000
 #define DIM_SIZE 128
 
@@ -714,7 +704,10 @@ private:
     {
       testData[i]= OFFSET+((ARRAY_SIZE-i) % 50);
     }
-    IdArrayHandle sorted = MakeArrayHandle(testData, ARRAY_SIZE);
+
+    IdArrayHandle unsorted = MakeArrayHandle(testData, ARRAY_SIZE);
+    IdArrayHandle sorted;
+    Algorithm::Copy(unsorted, sorted);
 
     //Validate the standard inplace sort is correct
     Algorithm::Sort(sorted);
@@ -725,6 +718,23 @@ private:
       vtkm::Id sorted2 = sorted.GetPortalConstControl().Get(i+1);
       VTKM_TEST_ASSERT(sorted1 <= sorted2, "Values not properly sorted.");
     }
+
+    std::cout << "-------------------------------------------------" << std::endl;
+    std::cout << "Sort of a ArrayHandleZip" << std::endl;
+
+    //verify that we can use ArrayHandleZip inplace
+    vtkm::cont::ArrayHandleZip< IdArrayHandle, IdArrayHandle> zipped(unsorted, sorted);
+
+    //verify we can use the default an custom operator sort with zip handle
+    Algorithm::Sort(zipped, comparison::SortGreater());
+    Algorithm::Sort(zipped);
+
+    for (vtkm::Id i = 0; i < ARRAY_SIZE; ++i)
+    {
+      vtkm::Pair<vtkm::Id,vtkm::Id> kv_sorted = zipped.GetPortalConstControl().Get(i);
+      VTKM_TEST_ASSERT(( OFFSET +  ( i / (ARRAY_SIZE/50)) ) == kv_sorted.first,
+                       "ArrayZipHandle improperly sorted");
+    }
   }
 
   static VTKM_CONT_EXPORT void TestSortWithComparisonObject()
@@ -787,6 +797,7 @@ private:
     Vec3ArrayHandle values = MakeArrayHandle(testValues, ARRAY_SIZE);
 
     Algorithm::SortByKey(keys,values);
+
     for(vtkm::Id i=0; i < ARRAY_SIZE; ++i)
       {
       //keys should be sorted from 1 to ARRAY_SIZE
@@ -1104,7 +1115,6 @@ private:
       VTKM_TEST_ASSERT( expectedValues[i] == v, "Incorrect reduced vale");
       }
     }
-
   }
 
   static VTKM_CONT_EXPORT void TestScanInclusive()