 /*
  * sort
  */
+
+template <typename T0, typename T1>
+class tuple_less {
+public:
+    __device__ bool operator()(tuple<T0, T1> i, tuple<T0, T1> j) {
+        T0 i0 = get<0>(i), j0 = get<0>(j);
+        T1 i1 = get<1>(i), j1 = get<1>(j);
+        return i0 < j0 || i0 == j0 && i1 < j1;
+    }
+};
 
 template <typename T>
 void cupy::thrust::_sort(void *data_start, size_t *keys_start, const std::vector<ptrdiff_t>& shape) {

#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/sort.h>
#include "cupy_common.h"

                  dp_keys_first,
                  divides<size_t>());

        stable_sort(
            make_zip_iterator(make_tuple(dp_keys_first, dp_data_first)),
            make_zip_iterator(make_tuple(dp_keys_last, dp_data_last)),
            tuple_less<size_t, T>());
    }
}

