Modernize the performance tests a bit.

Fixes NVIDIA#144

performance/SConscript

@@ -20,21 +20,21 @@ cu_builder = Builder(action = cu_build_function,
 my_env.Append(BUILDERS = {'CUFile' : cu_builder})
 
 # define a rule to build a report from an executable
-report_builder = Builder(action = os.path.join('"' + str(my_env.Dir('.')), '$SOURCE" > $TARGET'),
-                         suffix = '.xml',
-                         src_suffix = my_env['PROGSUFFIX'])
-my_env.Append(BUILDERS = {'Report' : report_builder})
+xml_builder = Builder(action = os.path.join('"' + str(my_env.Dir('.')), '$SOURCE" > $TARGET'),
+                      suffix = '.xml',
+                      src_suffix = my_env['PROGSUFFIX'])
+my_env.Append(BUILDERS = {'XMLFile' : xml_builder})
 
 my_env.Append(CPPPATH = [Dir('.').srcnode(), Dir('#/testing')])
 
 cu_list = []
 program_list = []
-report_list = []
+xml_list = []
 
 build_files = [os.path.join('build', f) for f in ['perftest.py', 'test_function_template.cxx']]
 
 # describe dependency graph:
-# report -> program -> .cu -> .test
+# xml -> program -> .cu -> .test
 for test in my_env.Glob('*.test'):
   cu = my_env.CUFile(test)
   my_env.Depends(cu, build_files)
@@ -43,21 +43,21 @@ for test in my_env.Glob('*.test'):
   prog = my_env.Program(cu)
   program_list.append(prog)
 
-  report = my_env.Report(prog)
-  report_list.append(report)
+  xml = my_env.XMLFile(prog)
+  xml_list.append(xml)
 
 # make aliases for groups of targets
-reports = my_env.Alias("reports", report_list)
-programs = my_env.Alias("programs", program_list)
+run_performance_tests_alias = my_env.Alias("run_performance_tests", xml_list)
+performance_tests_alias     = my_env.Alias("performance_tests", program_list)
 
 # when no build target is specified, by default we build the programs
-my_env.Default(programs)
+my_env.Default(performance_tests_alias)
 
 # output a help message
 my_env.Help("""
 Type: 'scons' to build all performance test programs.
-Type: 'scons reports' to run all performance tests and output reports.
+Type: 'scons run_performance_tests' to run all performance tests and output reports.
 Type: 'scons <test name>' to build a single performance test program of interest.
-Type: 'scons <test name>.xml' to run a single performance test of interest and output a report.
+Type: 'scons <test name>.xml' to run a single performance test of interest and output a report in an XML file.
 """)
 

performance/copy_if.test

@@ -1,18 +1,15 @@
 PREAMBLE = \
     """
-    #include <thrust/detail/backend/cuda/copy_if.h>
     #include <thrust/copy.h>
     #include <thrust/device_vector.h>
     #include <thrust/host_vector.h>
     #include <unittest/unittest.h>
     #include <thrust/sequence.h>
 
-    using namespace thrust;
-
     struct pred
     {
-        __host__ __device__
-            bool operator()(int x) { return bool(x); }
+      __host__ __device__
+      bool operator()(int x) { return bool(x); }
     };
 
     """
@@ -28,19 +25,16 @@ INITIALIZE = \
     thrust::device_vector<int> d_stencil = h_stencil;
     thrust::device_vector<int> d_output  = h_output;
 
-    typedef thrust::device_vector<int>::iterator OutputIterator;
-
-
     size_t h_count = thrust::copy_if(h_input.begin(), h_input.end(), h_stencil.begin(), h_output.begin(), pred()) - h_output.begin();
-    size_t d_count =       $Function(d_input.begin(), d_input.end(), d_stencil.begin(), d_output.begin(), pred()) - d_output.begin();
+    size_t d_count = thrust::copy_if(d_input.begin(), d_input.end(), d_stencil.begin(), d_output.begin(), pred()) - d_output.begin();
 
     ASSERT_EQUAL(h_output, d_output);
     ASSERT_EQUAL(h_count, d_count);
     """
 
 TIME = \
     """
-    $Function(d_input.begin(), d_input.end(), d_stencil.begin(), d_output.begin(), pred());
+    thrust::copy_if(d_input.begin(), d_input.end(), d_stencil.begin(), d_output.begin(), pred());
     """
 
 FINALIZE = \
@@ -50,8 +44,7 @@ FINALIZE = \
     RECORD_BANDWIDTH((2*sizeof(int) + 2*sizeof(float)) *  double($InputSize));
     """
 
-Functions  = ['thrust::copy_if','thrust::detail::backend::cuda::copy_if']
 InputSizes = [2**N for N in range(20, 27)]
 
-TestVariables = [('Function',Functions), ('InputSize', InputSizes)]
+TestVariables = [('InputSize', InputSizes)]
 

performance/indirect_sort.test

@@ -15,8 +15,7 @@ PREAMBLE = \
         __host__ __device__
         bool operator()(IndexType a, IndexType b)
         {
-            return comp(thrust::detail::backend::dereference(first, a),
-                        thrust::detail::backend::dereference(first, b));
+            return comp(thrust::raw_reference_cast(first[a]), thrust::raw_reference_cast(first[b]));
         }    
     };
 

performance/merge_sort.test

@@ -1,7 +1,16 @@
 PREAMBLE = \
     """
     #include <thrust/sort.h>
-    #include <thrust/detail/backend/cuda/detail/stable_merge_sort.h>
+
+    template<typename T>
+      struct my_less
+    {
+      __host__ __device__
+      bool operator()(const T &x, const T &y) const
+      {
+        return x < y;
+      }
+    };
     """
 
 INITIALIZE = \
@@ -12,15 +21,15 @@ INITIALIZE = \
 
     // test sort
     thrust::stable_sort(h_keys.begin(), h_keys.end());
-    thrust::detail::backend::cuda::detail::stable_merge_sort(d_keys.begin(), d_keys.end(), thrust::less<$KeyType>());
+    thrust::stable_sort(d_keys.begin(), d_keys.end(), my_less<$KeyType>());
 
     ASSERT_EQUAL(d_keys, h_keys);
     """
 
 TIME = \
     """
     thrust::copy(d_keys_copy.begin(), d_keys_copy.end(), d_keys.begin());
-    thrust::detail::backend::cuda::detail::stable_merge_sort(d_keys.begin(), d_keys.end(), thrust::less<$KeyType>());
+    thrust::stable_sort(d_keys.begin(), d_keys.end(), my_less<$KeyType>());
     """
 
 FINALIZE = \

performance/radix_sort.test

@@ -1,7 +1,6 @@
 PREAMBLE = \
     """
     #include <thrust/sort.h>
-    #include <thrust/detail/backend/cuda/detail/stable_radix_sort.h>
     """
 
 INITIALIZE = \
@@ -12,15 +11,15 @@ INITIALIZE = \
 
     // test sort
     thrust::stable_sort(h_keys.begin(), h_keys.end());
-    thrust::detail::backend::cuda::detail::stable_radix_sort(d_keys.begin(), d_keys.end());
+    thrust::stable_sort(d_keys.begin(), d_keys.end());
 
     ASSERT_EQUAL(d_keys, h_keys);
     """
 
 TIME = \
     """
     thrust::copy(d_keys_copy.begin(), d_keys_copy.end(), d_keys.begin());
-    thrust::detail::backend::cuda::detail::stable_radix_sort(d_keys.begin(), d_keys.end());
+    thrust::stable_sort(d_keys.begin(), d_keys.end());
     """
 
 FINALIZE = \

performance/radix_sort_bits.test

@@ -17,15 +17,15 @@ INITIALIZE = \
 
     // test sort
     thrust::stable_sort(h_keys.begin(), h_keys.end());
-    thrust::detail::backend::cuda::detail::stable_radix_sort(d_keys.begin(), d_keys.end());
+    thrust::stable_sort(d_keys.begin(), d_keys.end());
 
     ASSERT_EQUAL(d_keys, h_keys);
     """
 
 TIME = \
     """
     thrust::copy(d_keys_copy.begin(), d_keys_copy.end(), d_keys.begin());
-    thrust::detail::backend::cuda::detail::stable_radix_sort(d_keys.begin(), d_keys.end());
+    thrust::stable_sort(d_keys.begin(), d_keys.end());
     """
 
 FINALIZE = \

performance/radix_sort_by_key.test

@@ -18,7 +18,7 @@ INITIALIZE = \
 
     // test sort
     thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
-    thrust::detail::backend::cuda::detail::stable_radix_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
 
     ASSERT_EQUAL(d_keys,   h_keys);
     ASSERT_EQUAL(d_values, h_values);
@@ -27,7 +27,7 @@ INITIALIZE = \
 TIME = \
     """
     thrust::copy(d_keys_copy.begin(), d_keys_copy.end(), d_keys.begin());
-    thrust::detail::backend::cuda::detail::stable_radix_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
     """
 
 FINALIZE = \

performance/report.py

@@ -3,8 +3,8 @@
 #valid formats are png, pdf, ps, eps and svg
 #if format=None the plot will be displayed
 format = 'png'
-#output = print_results
-output = plot_results
+output = print_results
+#output = plot_results
 
 for function in ['fill', 'reduce', 'inner_product', 'gather', 'merge']:
     output(function + '.xml', 'InputType', 'InputSize', 'Bandwidth', format=format)
@@ -18,6 +18,9 @@
 for method in ['sort', 'comparison_sort', 'radix_sort']:
     output(method + '.xml',    'KeyType', 'InputSize', 'Sorting', title='thrust::' + method, format=format)
     output(method + '_by_key.xml', 'KeyType', 'InputSize', 'Sorting', title='thrust::' + method + '_by_key', format=format)
+
+for method in ['set_difference', 'set_intersection', 'set_symmetric_difference', 'set_union']:
+  output(method + '.xml', 'InputType', 'InputSize', 'Sorting', title='thrust::' + method, format=format)
 
 output('stl_sort.xml', 'KeyType', 'InputSize', 'Sorting', title='std::sort', format=format)
 

performance/set_symmetric_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_symmetric_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_symmetric_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_symmetric_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)]
+

performance/sort_by_key.test

@@ -2,6 +2,17 @@ PREAMBLE = \
     """
     #include <thrust/sort.h>
     #include <thrust/sequence.h>
+
+    template<typename T>
+    struct my_less
+    {
+      __host__ __device__
+      bool operator()(const T &x, const T& y) const
+      {
+        return x < y;
+      }
+    };
+
     """
 
 INITIALIZE = \
@@ -17,8 +28,8 @@ INITIALIZE = \
     thrust::device_vector<$KeyType> d_keys_copy = d_keys;
 
     // test sort
-    thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
-    thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+    thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin(), my_less<$KeyType>());
 
     ASSERT_EQUAL(d_keys,   h_keys);
     ASSERT_EQUAL(d_values, h_values);
@@ -27,7 +38,7 @@ INITIALIZE = \
 TIME = \
     """
     thrust::copy(d_keys_copy.begin(), d_keys_copy.end(), d_keys.begin());
-    thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin(), my_less<$KeyType>());
     """
 
 FINALIZE = \

performance/sort_large.test

@@ -5,6 +5,7 @@ PREAMBLE = \
     template <typename T>
     struct my_less : public thrust::binary_function<T,T,bool>
     {
+        __host__ __device__
         bool operator()(const T& a, const T& b) const
         {
             return a < b;