Add set_difference and unit tests.

performance/set_difference.test

@@ -0,0 +1,45 @@
+PREAMBLE = \
+    """
+    #include <thrust/set_operations.h>
+    #include <thrust/sort.h>
+    #include <algorithm>
+    """
+
+INITIALIZE = \
+    """
+    thrust::host_vector<$InputType> h_a = unittest::random_integers<$InputType>($InputSize);
+    thrust::host_vector<$InputType> h_b = unittest::random_integers<$InputType>($InputSize);
+    thrust::sort(h_a.begin(), h_a.end());
+    thrust::sort(h_b.begin(), h_b.end());
+
+    thrust::host_vector<$InputType> h_result(h_a.size());
+    thrust::host_vector<$InputType>::iterator new_end = 
+      thrust::set_difference(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), h_result.begin());
+    h_result.resize(new_end - h_result.begin());
+
+    thrust::device_vector<$InputType> d_a = h_a, d_b = h_b;
+
+    thrust::device_vector<$InputType> d_result(h_result.size());
+    thrust::set_difference(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), d_result.begin());
+
+    ASSERT_EQUAL(h_result, d_result);
+    """
+
+TIME = \
+    """
+    thrust::set_difference(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), d_result.begin());
+    """
+
+FINALIZE = \
+    """
+    RECORD_TIME();
+    RECORD_BANDWIDTH((2 * double($InputSize) + d_result.size()) * sizeof($InputType));
+    RECORD_SORTING_RATE(2 * double($InputSize))
+    """
+
+
+InputTypes = ['char', 'short', 'int', 'long', 'float', 'double']
+InputSizes = [2**N for N in range(10, 25)]
+
+TestVariables = [('InputType', InputTypes), ('InputSize', InputSizes)]
+

testing/set_difference.cu

@@ -0,0 +1,188 @@
+#include <unittest/unittest.h>
+#include <thrust/set_operations.h>
+#include <thrust/functional.h>
+#include <thrust/sort.h>
+
+template<typename Vector>
+void TestSetDifferenceSimple(void)
+{
+  typedef typename Vector::iterator Iterator;
+
+  Vector a(4), b(5);
+
+  a[0] = 0; a[1] = 2; a[2] = 4; a[3] = 5;
+  b[0] = 0; b[1] = 3; b[2] = 3; b[3] = 4; b[4] = 6;
+
+  Vector ref(2);
+  ref[0] = 2; ref[1] = 5;
+
+  Vector result(2);
+
+  Iterator end = thrust::set_difference(a.begin(), a.end(),
+                                        b.begin(), b.end(),
+                                        result.begin());
+
+  ASSERT_EQUAL_QUIET(result.end(), end);
+  ASSERT_EQUAL(ref, result);
+}
+DECLARE_VECTOR_UNITTEST(TestSetDifferenceSimple);
+
+
+template<typename T>
+void TestSetDifference(const size_t n)
+{
+  thrust::host_vector<T> temp = unittest::random_integers<T>(2 * n);
+  thrust::host_vector<T> h_a(temp.begin(), temp.begin() + n);
+  thrust::host_vector<T> h_b(temp.begin() + n, temp.end());
+
+  thrust::sort(h_a.begin(), h_a.end());
+  thrust::sort(h_b.begin(), h_b.end());
+
+  thrust::device_vector<T> d_a = h_a;
+  thrust::device_vector<T> d_b = h_b;
+
+  thrust::host_vector<T> h_result(n);
+  thrust::device_vector<T> d_result(n);
+
+  typename thrust::host_vector<T>::iterator h_end;
+  typename thrust::device_vector<T>::iterator d_end;
+
+  h_end = thrust::set_difference(h_a.begin(), h_a.end(),
+                                 h_b.begin(), h_b.end(),
+                                 h_result.begin());
+  h_result.resize(h_end - h_result.begin());
+
+  d_end = thrust::set_difference(d_a.begin(), d_a.end(),
+                                 d_b.begin(), d_b.end(),
+                                 d_result.begin());
+
+  d_result.resize(d_end - d_result.begin());
+
+  ASSERT_EQUAL(h_result, d_result);
+}
+DECLARE_VARIABLE_UNITTEST(TestSetDifference);
+
+
+template<typename T>
+void TestSetDifferenceEquivalentRanges(const size_t n)
+{
+  thrust::host_vector<T> temp = unittest::random_integers<T>(n);
+  thrust::host_vector<T> h_a = temp; thrust::sort(h_a.begin(), h_a.end());
+  thrust::host_vector<T> h_b = h_a;
+
+  thrust::device_vector<T> d_a = h_a;
+  thrust::device_vector<T> d_b = h_b;
+
+  thrust::host_vector<T>   h_result(n);
+  thrust::device_vector<T> d_result(n);
+
+  typename thrust::host_vector<T>::iterator   h_end;
+  typename thrust::device_vector<T>::iterator d_end;
+
+  h_end = thrust::set_difference(h_a.begin(), h_a.end(),
+                                 h_b.begin(), h_b.end(),
+                                 h_result.begin());
+  h_result.resize(h_end - h_result.begin());
+
+  d_end = thrust::set_difference(d_a.begin(), d_a.end(),
+                                 d_b.begin(), d_b.end(),
+                                 d_result.begin());
+
+  d_result.resize(d_end - d_result.begin());
+
+  ASSERT_EQUAL(h_result, d_result);
+}
+DECLARE_VARIABLE_UNITTEST(TestSetDifferenceEquivalentRanges);
+
+
+template<typename T>
+void TestSetDifferenceMultiset(const size_t n)
+{
+  thrust::host_vector<T> temp = unittest::random_integers<T>(2 * n);
+
+  // restrict elements to [min,13)
+  for(typename thrust::host_vector<T>::iterator i = temp.begin();
+      i != temp.end();
+      ++i)
+  {
+    int temp = static_cast<int>(*i);
+    temp %= 13;
+    *i = temp;
+  }
+
+  thrust::host_vector<T> h_a(temp.begin(), temp.begin() + n);
+  thrust::host_vector<T> h_b(temp.begin() + n, temp.end());
+
+  thrust::sort(h_a.begin(), h_a.end());
+  thrust::sort(h_b.begin(), h_b.end());
+
+  thrust::device_vector<T> d_a = h_a;
+  thrust::device_vector<T> d_b = h_b;
+
+  thrust::host_vector<T> h_result(n);
+  thrust::device_vector<T> d_result(n);
+
+  typename thrust::host_vector<T>::iterator h_end;
+  typename thrust::device_vector<T>::iterator d_end;
+
+  h_end = thrust::set_difference(h_a.begin(), h_a.end(),
+                                 h_b.begin(), h_b.end(),
+                                 h_result.begin());
+  h_result.resize(h_end - h_result.begin());
+
+  d_end = thrust::set_difference(d_a.begin(), d_a.end(),
+                                 d_b.begin(), d_b.end(),
+                                 d_result.begin());
+
+  d_result.resize(d_end - d_result.begin());
+
+  ASSERT_EQUAL(h_result, d_result);
+}
+DECLARE_VARIABLE_UNITTEST(TestSetDifferenceMultiset);
+
+
+template<typename U>
+  void TestSetDifferenceKeyValue(size_t n)
+{
+  typedef key_value<U,U> T;
+
+  thrust::host_vector<U> h_keys_a   = unittest::random_integers<U>(n);
+  thrust::host_vector<U> h_values_a = unittest::random_integers<U>(n);
+
+  thrust::host_vector<U> h_keys_b   = unittest::random_integers<U>(n);
+  thrust::host_vector<U> h_values_b = unittest::random_integers<U>(n);
+
+  thrust::host_vector<T> h_a(n), h_b(n);
+  for(size_t i = 0; i < n; ++i)
+  {
+    h_a[i] = T(h_keys_a[i], h_values_a[i]);
+    h_b[i] = T(h_keys_b[i], h_values_b[i]);
+  }
+
+  thrust::stable_sort(h_a.begin(), h_a.end());
+  thrust::stable_sort(h_b.begin(), h_b.end());
+
+  thrust::device_vector<T> d_a = h_a;
+  thrust::device_vector<T> d_b = h_b;
+
+  thrust::host_vector<T> h_result(n);
+  thrust::device_vector<T> d_result(n);
+
+  typename thrust::host_vector<T>::iterator h_end;
+  typename thrust::device_vector<T>::iterator d_end;
+
+  h_end = thrust::set_difference(h_a.begin(), h_a.end(),
+                                 h_b.begin(), h_b.end(),
+                                 h_result.begin());
+  h_result.resize(h_end - h_result.begin());
+
+  d_end = thrust::set_difference(d_a.begin(), d_a.end(),
+                                 d_b.begin(), d_b.end(),
+                                 d_result.begin());
+
+  d_result.resize(d_end - d_result.begin());
+
+  ASSERT_EQUAL_QUIET(h_result, d_result);
+}
+DECLARE_VARIABLE_UNITTEST(TestSetDifferenceKeyValue);
+

testing/set_difference_descending.cu

@@ -0,0 +1,68 @@
+#include <unittest/unittest.h>
+#include <thrust/set_operations.h>
+#include <thrust/functional.h>
+#include <thrust/sort.h>
+
+template<typename Vector>
+void TestSetDifferenceAscendingSimple(void)
+{
+  typedef typename Vector::value_type T;
+  typedef typename Vector::iterator Iterator;
+
+  Vector a(4), b(5);
+
+  a[0] = 5; a[1] = 4; a[2] = 2; a[3] = 0;
+  b[0] = 6; b[1] = 4; b[2] = 3; b[3] = 3; b[4] = 0;
+
+  Vector ref(2);
+  ref[0] = 5; ref[1] = 2;
+
+  Vector result(2);
+
+  Iterator end = thrust::set_difference(a.begin(), a.end(),
+                                        b.begin(), b.end(),
+                                        result.begin(),
+                                        thrust::greater<T>());
+
+  ASSERT_EQUAL_QUIET(result.end(), end);
+  ASSERT_EQUAL(ref, result);
+}
+DECLARE_VECTOR_UNITTEST(TestSetDifferenceAscendingSimple);
+
+
+template<typename T>
+void TestSetDifferenceAscending(const size_t n)
+{
+  thrust::host_vector<T> temp = unittest::random_integers<T>(2 * n);
+  thrust::host_vector<T> h_a(temp.begin(), temp.begin() + n);
+  thrust::host_vector<T> h_b(temp.begin() + n, temp.end());
+
+  thrust::sort(h_a.begin(), h_a.end(), thrust::greater<T>());
+  thrust::sort(h_b.begin(), h_b.end(), thrust::greater<T>());
+
+  thrust::device_vector<T> d_a = h_a;
+  thrust::device_vector<T> d_b = h_b;
+
+  thrust::host_vector<T> h_result(n);
+  thrust::device_vector<T> d_result(n);
+
+  typename thrust::host_vector<T>::iterator h_end;
+  typename thrust::device_vector<T>::iterator d_end;
+
+  h_end = thrust::set_difference(h_a.begin(), h_a.end(),
+                                 h_b.begin(), h_b.end(),
+                                 h_result.begin(),
+                                 thrust::greater<T>());
+  h_result.resize(h_end - h_result.begin());
+
+  d_end = thrust::set_difference(d_a.begin(), d_a.end(),
+                                 d_b.begin(), d_b.end(),
+                                 d_result.begin(),
+                                 thrust::greater<T>());
+
+  d_result.resize(d_end - d_result.begin());
+
+  ASSERT_EQUAL(h_result, d_result);
+}
+DECLARE_VARIABLE_UNITTEST(TestSetDifferenceAscending);
+

thrust/detail/device/cuda/block/set_difference.h

@@ -0,0 +1,52 @@
+/*
+ *  Copyright 2008-2010 NVIDIA Corporation
+ *
+ *  Licensed under the Apache License, Version 2.0 (the "License");
+ *  you may not use this file except in compliance with the License.
+ *  You may obtain a copy of the License at
+ *
+ *      http://www.apache.org/licenses/LICENSE-2.0
+ *
+ *  Unless required by applicable law or agreed to in writing, software
+ *  distributed under the License is distributed on an "AS IS" BASIS,
+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ *  See the License for the specific language governing permissions and
+ *  limitations under the License.
+ */
+
+#pragma once
+
+namespace thrust
+{
+namespace detail
+{
+namespace device
+{
+namespace cuda
+{
+namespace block
+{
+
+template<typename RandomAccessIterator1,
+         typename RandomAccessIterator2,
+         typename RandomAccessIterator3,
+         typename RandomAccessIterator4,
+         typename StrictWeakOrdering>
+__device__ __forceinline__
+  RandomAccessIterator4 set_difference(RandomAccessIterator1 first1,
+                                       RandomAccessIterator1 last1,
+                                       RandomAccessIterator2 first2,
+                                       RandomAccessIterator2 last2,
+                                       RandomAccessIterator3 temporary,
+                                       RandomAccessIterator4 result,
+                                       StrictWeakOrdering comp);
+
+} // end block
+} // end cuda
+} // end device
+} // end detail
+} // end thrust
+
+#include <thrust/detail/device/cuda/block/set_difference.inl>
+
+