From 3df3bcf3c5a6060317aa22861e720df3f681f8bb Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Mon, 24 Jun 2024 14:28:23 -0400 Subject: [PATCH] Remove legacy benchmarks and other dvs/p4 remenants. (#1901) Removes dead code and other assorted eldritch horrors from bygone eras. --- thrust/CMakeLists.txt | 6 - thrust/internal/benchmark/CMakeLists.txt | 31 - thrust/internal/benchmark/README.txt | 30 - thrust/internal/benchmark/bench.cu | 1213 --------------- thrust/internal/benchmark/bench.mk | 20 - .../benchmark/combine_benchmark_results.py | 816 ---------- .../benchmark/compare_benchmark_results.py | 1307 ----------------- thrust/internal/benchmark/random.h | 92 -- thrust/internal/benchmark/tbb_algos.h | 230 --- thrust/internal/benchmark/timer.h | 137 -- thrust/internal/build/common_build.mk | 88 -- thrust/internal/build/common_compiler.mk | 159 -- thrust/internal/build/common_detect.mk | 14 - thrust/internal/build/generic_example.mk | 12 - thrust/internal/build/generic_test.mk | 22 - thrust/internal/build/testframework.mk | 16 - thrust/internal/build/warningstester.mk | 63 - .../warningstester_create_uber_header.py | 54 - thrust/internal/racecheck.sh | 26 - thrust/internal/rename_cub_namespace.sh | 6 - .../internal/reverse_rename_cub_namespace.sh | 6 - thrust/internal/scripts/eris_perf.py | 188 --- .../internal/scripts/refresh_from_github2.sh | 96 -- thrust/internal/scripts/tounix | 6 - thrust/internal/scripts/wiki2tex.py | 194 --- thrust/internal/test/dvstest.lst | 425 ------ thrust/internal/test/thrust_nightly.pl | 600 -------- thrust/internal/test/unittest.lst | 1267 ---------------- thrust/internal/test/unittest_omp.lst | 808 ---------- thrust/internal/test/warningstester.cu | 7 - thrust/testing/CMakeLists.txt | 1 - thrust/testing/cuda/adjacent_difference.mk | 1 - thrust/testing/cuda/binary_search.mk | 1 - thrust/testing/cuda/complex.mk | 1 - thrust/testing/cuda/copy.mk | 1 - thrust/testing/cuda/copy_if.mk | 1 - thrust/testing/cuda/count.mk | 1 - thrust/testing/cuda/cudart.mk | 1 - thrust/testing/cuda/equal.mk | 1 - thrust/testing/cuda/fill.mk | 1 - thrust/testing/cuda/find.mk | 1 - thrust/testing/cuda/for_each.mk | 1 - thrust/testing/cuda/gather.mk | 1 - thrust/testing/cuda/generate.mk | 1 - thrust/testing/cuda/inner_product.mk | 1 - thrust/testing/cuda/is_partitioned.mk | 1 - thrust/testing/cuda/is_sorted.mk | 1 - thrust/testing/cuda/is_sorted_until.mk | 1 - thrust/testing/cuda/logical.mk | 1 - thrust/testing/cuda/managed_memory_pointer.mk | 1 - thrust/testing/cuda/max_element.mk | 1 - thrust/testing/cuda/memory.mk | 1 - thrust/testing/cuda/merge.mk | 1 - thrust/testing/cuda/merge_by_key.mk | 1 - thrust/testing/cuda/merge_sort.mk | 1 - thrust/testing/cuda/min_element.mk | 1 - thrust/testing/cuda/minmax_element.mk | 1 - thrust/testing/cuda/mismatch.mk | 1 - thrust/testing/cuda/pair_sort.mk | 1 - thrust/testing/cuda/pair_sort_by_key.mk | 1 - thrust/testing/cuda/partition.mk | 1 - thrust/testing/cuda/partition_point.mk | 1 - thrust/testing/cuda/reduce.mk | 1 - thrust/testing/cuda/reduce_by_key.mk | 1 - thrust/testing/cuda/remove.mk | 1 - thrust/testing/cuda/replace.mk | 1 - thrust/testing/cuda/reverse.mk | 1 - thrust/testing/cuda/scan.mk | 1 - thrust/testing/cuda/scan_by_key.mk | 1 - thrust/testing/cuda/scatter.mk | 1 - thrust/testing/cuda/sequence.mk | 1 - thrust/testing/cuda/set_difference.mk | 1 - thrust/testing/cuda/set_difference_by_key.mk | 1 - thrust/testing/cuda/set_intersection.mk | 1 - .../testing/cuda/set_intersection_by_key.mk | 1 - .../testing/cuda/set_symmetric_difference.mk | 1 - .../cuda/set_symmetric_difference_by_key.mk | 1 - thrust/testing/cuda/set_union.mk | 1 - thrust/testing/cuda/set_union_by_key.mk | 1 - thrust/testing/cuda/sort.mk | 1 - thrust/testing/cuda/sort_by_key.mk | 1 - thrust/testing/cuda/stream_per_thread.mk | 1 - thrust/testing/cuda/swap_ranges.mk | 1 - thrust/testing/cuda/tabulate.mk | 1 - thrust/testing/cuda/transform.mk | 1 - thrust/testing/cuda/transform_reduce.mk | 1 - thrust/testing/cuda/transform_scan.mk | 1 - thrust/testing/cuda/uninitialized_copy.mk | 1 - thrust/testing/cuda/uninitialized_fill.mk | 1 - thrust/testing/cuda/unique.mk | 1 - thrust/testing/cuda/unique_by_key.mk | 1 - thrust/testing/regression/CMakeLists.txt | 20 - ...1706__thrust_optional_reference_emplace.cu | 40 - ...y_wrong_element_type_default_comparator.cu | 27 - ...vbug_2318871__zip_iterator_with_complex.cu | 38 - ..._928_nvbug_2341455__reduce_with_complex.cu | 9 - ...nvbug_1632709__reduce_large_input_sizes.cu | 17 - ...g_1940974__merge_with_constant_iterator.cu | 34 - ...ary_static_on_get_occ_device_properties.cu | 5 - ...__scan_requires_assignability_from_zero.cu | 22 - ...requires_assignability_from_zero.fixed0.cu | 26 - ...requires_assignability_from_zero.fixed1.cu | 23 - 102 files changed, 8261 deletions(-) delete mode 100644 thrust/internal/benchmark/CMakeLists.txt delete mode 100644 thrust/internal/benchmark/README.txt delete mode 100644 thrust/internal/benchmark/bench.cu delete mode 100644 thrust/internal/benchmark/bench.mk delete mode 100755 thrust/internal/benchmark/combine_benchmark_results.py delete mode 100755 thrust/internal/benchmark/compare_benchmark_results.py delete mode 100644 thrust/internal/benchmark/random.h delete mode 100644 thrust/internal/benchmark/tbb_algos.h delete mode 100644 thrust/internal/benchmark/timer.h delete mode 100644 thrust/internal/build/common_build.mk delete mode 100644 thrust/internal/build/common_compiler.mk delete mode 100644 thrust/internal/build/common_detect.mk delete mode 100644 thrust/internal/build/generic_example.mk delete mode 100644 thrust/internal/build/generic_test.mk delete mode 100644 thrust/internal/build/testframework.mk delete mode 100644 thrust/internal/build/warningstester.mk delete mode 100644 thrust/internal/build/warningstester_create_uber_header.py delete mode 100755 thrust/internal/racecheck.sh delete mode 100755 thrust/internal/rename_cub_namespace.sh delete mode 100755 thrust/internal/reverse_rename_cub_namespace.sh delete mode 100755 thrust/internal/scripts/eris_perf.py delete mode 100755 thrust/internal/scripts/refresh_from_github2.sh delete mode 100755 thrust/internal/scripts/tounix delete mode 100644 thrust/internal/scripts/wiki2tex.py delete mode 100755 thrust/internal/test/dvstest.lst delete mode 100755 thrust/internal/test/thrust_nightly.pl delete mode 100644 thrust/internal/test/unittest.lst delete mode 100644 thrust/internal/test/unittest_omp.lst delete mode 100644 thrust/internal/test/warningstester.cu delete mode 100644 thrust/testing/cuda/adjacent_difference.mk delete mode 100644 thrust/testing/cuda/binary_search.mk delete mode 100644 thrust/testing/cuda/complex.mk delete mode 100644 thrust/testing/cuda/copy.mk delete mode 100644 thrust/testing/cuda/copy_if.mk delete mode 100644 thrust/testing/cuda/count.mk delete mode 100644 thrust/testing/cuda/cudart.mk delete mode 100644 thrust/testing/cuda/equal.mk delete mode 100644 thrust/testing/cuda/fill.mk delete mode 100644 thrust/testing/cuda/find.mk delete mode 100644 thrust/testing/cuda/for_each.mk delete mode 100644 thrust/testing/cuda/gather.mk delete mode 100644 thrust/testing/cuda/generate.mk delete mode 100644 thrust/testing/cuda/inner_product.mk delete mode 100644 thrust/testing/cuda/is_partitioned.mk delete mode 100644 thrust/testing/cuda/is_sorted.mk delete mode 100644 thrust/testing/cuda/is_sorted_until.mk delete mode 100644 thrust/testing/cuda/logical.mk delete mode 100644 thrust/testing/cuda/managed_memory_pointer.mk delete mode 100644 thrust/testing/cuda/max_element.mk delete mode 100644 thrust/testing/cuda/memory.mk delete mode 100644 thrust/testing/cuda/merge.mk delete mode 100644 thrust/testing/cuda/merge_by_key.mk delete mode 100644 thrust/testing/cuda/merge_sort.mk delete mode 100644 thrust/testing/cuda/min_element.mk delete mode 100644 thrust/testing/cuda/minmax_element.mk delete mode 100644 thrust/testing/cuda/mismatch.mk delete mode 100644 thrust/testing/cuda/pair_sort.mk delete mode 100644 thrust/testing/cuda/pair_sort_by_key.mk delete mode 100644 thrust/testing/cuda/partition.mk delete mode 100644 thrust/testing/cuda/partition_point.mk delete mode 100644 thrust/testing/cuda/reduce.mk delete mode 100644 thrust/testing/cuda/reduce_by_key.mk delete mode 100644 thrust/testing/cuda/remove.mk delete mode 100644 thrust/testing/cuda/replace.mk delete mode 100644 thrust/testing/cuda/reverse.mk delete mode 100644 thrust/testing/cuda/scan.mk delete mode 100644 thrust/testing/cuda/scan_by_key.mk delete mode 100644 thrust/testing/cuda/scatter.mk delete mode 100644 thrust/testing/cuda/sequence.mk delete mode 100644 thrust/testing/cuda/set_difference.mk delete mode 100644 thrust/testing/cuda/set_difference_by_key.mk delete mode 100644 thrust/testing/cuda/set_intersection.mk delete mode 100644 thrust/testing/cuda/set_intersection_by_key.mk delete mode 100644 thrust/testing/cuda/set_symmetric_difference.mk delete mode 100644 thrust/testing/cuda/set_symmetric_difference_by_key.mk delete mode 100644 thrust/testing/cuda/set_union.mk delete mode 100644 thrust/testing/cuda/set_union_by_key.mk delete mode 100644 thrust/testing/cuda/sort.mk delete mode 100644 thrust/testing/cuda/sort_by_key.mk delete mode 100644 thrust/testing/cuda/stream_per_thread.mk delete mode 100644 thrust/testing/cuda/swap_ranges.mk delete mode 100644 thrust/testing/cuda/tabulate.mk delete mode 100644 thrust/testing/cuda/transform.mk delete mode 100644 thrust/testing/cuda/transform_reduce.mk delete mode 100644 thrust/testing/cuda/transform_scan.mk delete mode 100644 thrust/testing/cuda/uninitialized_copy.mk delete mode 100644 thrust/testing/cuda/uninitialized_fill.mk delete mode 100644 thrust/testing/cuda/unique.mk delete mode 100644 thrust/testing/cuda/unique_by_key.mk delete mode 100644 thrust/testing/regression/CMakeLists.txt delete mode 100644 thrust/testing/regression/gh_1706__thrust_optional_reference_emplace.cu delete mode 100644 thrust/testing/regression/gh_911__merge_by_key_wrong_element_type_default_comparator.cu delete mode 100644 thrust/testing/regression/gh_919_nvbug_2318871__zip_iterator_with_complex.cu delete mode 100644 thrust/testing/regression/gh_928_nvbug_2341455__reduce_with_complex.cu delete mode 100644 thrust/testing/regression/nvbug_1632709__reduce_large_input_sizes.cu delete mode 100644 thrust/testing/regression/nvbug_1940974__merge_with_constant_iterator.cu delete mode 100644 thrust/testing/regression/nvbug_1965743__unnecessary_static_on_get_occ_device_properties.cu delete mode 100644 thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.cu delete mode 100644 thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed0.cu delete mode 100644 thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed1.cu diff --git a/thrust/CMakeLists.txt b/thrust/CMakeLists.txt index a779a85917..a7da7329fb 100644 --- a/thrust/CMakeLists.txt +++ b/thrust/CMakeLists.txt @@ -60,11 +60,6 @@ option(THRUST_ENABLE_EXAMPLES "Build Thrust examples." "ON") option(THRUST_ENABLE_BENCHMARKS "Build Thrust runtime benchmarks." "${CCCL_ENABLE_BENCHMARKS}") option(THRUST_INCLUDE_CUB_CMAKE "Build CUB tests and examples. (Requires CUDA)." "OFF") -# Mark this option as advanced for now. We'll revisit this later once the new -# benchmarks are ready. For now, we just need to expose a way to compile -# bench.cu from CMake for NVIDIA's internal builds. -mark_as_advanced(THRUST_ENABLE_BENCHMARKS) - # Check if we're actually building anything before continuing. If not, no need # to search for deps, etc. This is a common approach for packagers that just # need the install rules. See GH issue NVIDIA/thrust#1211. @@ -135,7 +130,6 @@ endif() if (THRUST_ENABLE_BENCHMARKS) add_subdirectory(benchmarks) - add_subdirectory(internal/benchmark) endif() if (THRUST_INCLUDE_CUB_CMAKE AND THRUST_CUDA_FOUND) diff --git a/thrust/internal/benchmark/CMakeLists.txt b/thrust/internal/benchmark/CMakeLists.txt deleted file mode 100644 index 5fbd365da8..0000000000 --- a/thrust/internal/benchmark/CMakeLists.txt +++ /dev/null @@ -1,31 +0,0 @@ -if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC") - # MSVC builds fail at runtime. Benchmarks are linux-only for now. - message(STATUS "Thrust benchmarking is not available on MSVC.") - return() -endif() - -add_custom_target(thrust.all.bench) - -foreach(thrust_target IN LISTS THRUST_TARGETS) - thrust_get_target_property(config_host ${thrust_target} HOST) - thrust_get_target_property(config_device ${thrust_target} DEVICE) - thrust_get_target_property(config_prefix ${thrust_target} PREFIX) - - # Skip non cpp.cuda targets: - if (NOT config_host STREQUAL "CPP" OR - NOT config_device STREQUAL "CUDA") - continue() - endif() - - set(bench_target ${config_prefix}.bench) - - add_executable(${bench_target} bench.cu) - target_link_libraries(${bench_target} PRIVATE ${thrust_target}) - target_include_directories(${bench_target} PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}") - thrust_clone_target_properties(${bench_target} ${thrust_target}) - thrust_configure_cuda_target(${bench_target} RDC ${THRUST_FORCE_RDC}) - thrust_fix_clang_nvcc_build_for(${bench_target}) - - add_dependencies(thrust.all.bench ${bench_target}) - add_dependencies(${config_prefix}.all ${bench_target}) -endforeach() diff --git a/thrust/internal/benchmark/README.txt b/thrust/internal/benchmark/README.txt deleted file mode 100644 index 026da00649..0000000000 --- a/thrust/internal/benchmark/README.txt +++ /dev/null @@ -1,30 +0,0 @@ -Directions for compiling and running the benchmark with Ubuntu Linux: - -Install Intel's Threading Building Blocks library (TBB): -$ sudo apt-get install libtbb-dev - -Compile the benchmark: -$ nvcc -O3 -arch=sm_20 bench.cu -ltbb -o bench - -Run the benchmark: -$ ./bench - -Typical output (Tesla C2050): - -Benchmarking with input size 33554432 -Core Primitive Performance (elements per second) - Algorithm, STL, TBB, Thrust - reduce, 3121746688, 3739585536, 26134038528 - transform, 1869492736, 2347719424, 13804681216 - scan, 1394143744, 1439394816, 5039195648 - sort, 11070660, 34622352, 673543168 -Sorting Performance (keys per second) - Type, STL, TBB, Thrust - char, 24050078, 62987040, 2798874368 - short, 15644141, 41275164, 1428603008 - int, 11062616, 33478628, 682295744 - long, 11249874, 33972564, 219719184 - float, 9850043, 29011806, 692407232 -double, 9700181, 27153626, 224345568 - -The reported numbers are performance rates in "elements per second" (higher is better). diff --git a/thrust/internal/benchmark/bench.cu b/thrust/internal/benchmark/bench.cu deleted file mode 100644 index 4ec2391fef..0000000000 --- a/thrust/internal/benchmark/bench.cu +++ /dev/null @@ -1,1213 +0,0 @@ -#include - -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include // For CHAR_BIT. -#include // For `sqrt` and `abs`. -#include // For `atoi`. -#include -#include -#include -#include -#include -#include - -#include "random.h" -#include "timer.h" -#include // For `intN_t`. - -#if defined(HAVE_TBB) -# include "tbb_algos.h" -#endif - -#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA -# include // For `thrust::cuda_category` -# include // For `thrust::system_error` -#endif - -// We don't use THRUST_PP_STRINGIZE and THRUST_PP_CAT because they are new, and -// we want this benchmark to be backwards-compatible to older versions of Thrust. -#define PP_STRINGIZE_(expr) #expr -#define PP_STRINGIZE(expr) PP_STRINGIZE_(expr) - -#define PP_CAT(a, b) a##b -/////////////////////////////////////////////////////////////////////////////// - -template -struct squared_difference -{ -private: - T const average; - -public: - __host__ __device__ squared_difference(squared_difference const& rhs) - : average(rhs.average) - {} - - __host__ __device__ squared_difference(T average_) - : average(average_) - {} - - __host__ __device__ T operator()(T x) const - { - return (x - average) * (x - average); - } -}; - -template -struct value_and_count -{ - T value; - uint64_t count; - - __host__ __device__ value_and_count(value_and_count const& other) - : value(other.value) - , count(other.count) - {} - - __host__ __device__ value_and_count(T const& value_) - : value(value_) - , count(1) - {} - - __host__ __device__ value_and_count(T const& value_, uint64_t count_) - : value(value_) - , count(count_) - {} - - __host__ __device__ value_and_count& operator=(value_and_count const& other) - { - value = other.value; - count = other.count; - return *this; - } - - __host__ __device__ value_and_count& operator=(T const& value_) - { - value = value_; - count = 1; - return *this; - } -}; - -template -struct counting_op -{ -private: - ReduceOp reduce; - -public: - __host__ __device__ counting_op() - : reduce() - {} - - __host__ __device__ counting_op(counting_op const& other) - : reduce(other.reduce) - {} - - __host__ __device__ counting_op(ReduceOp const& reduce_) - : reduce(reduce_) - {} - - __host__ __device__ value_and_count operator()(value_and_count const& x, T const& y) const - { - return value_and_count(reduce(x.value, y), x.count + 1); - } - - __host__ __device__ value_and_count operator()(value_and_count const& x, value_and_count const& y) const - { - return value_and_count(reduce(x.value, y.value), x.count + y.count); - } -}; - -template -T arithmetic_mean(InputIt first, InputIt last, T init) -{ - value_and_count init_vc(init, 0); - - counting_op> reduce_vc; - - value_and_count vc = thrust::reduce(first, last, init_vc, reduce_vc); - - return vc.value / vc.count; -} - -template -typename thrust::iterator_traits::value_type arithmetic_mean(InputIt first, InputIt last) -{ - typedef typename thrust::iterator_traits::value_type T; - return arithmetic_mean(first, last, T()); -} - -template -T sample_standard_deviation(InputIt first, InputIt last, T average) -{ - value_and_count init_vc(T(), 0); - - counting_op> reduce_vc; - - squared_difference transform(average); - - value_and_count vc = thrust::transform_reduce(first, last, transform, init_vc, reduce_vc); - - return std::sqrt(vc.value / T(vc.count - 1)); -} - -/////////////////////////////////////////////////////////////////////////////// - -// Formulas for propagation of uncertainty from: -// -// https://en.wikipedia.org/wiki/Propagation_of_uncertainty#Example_formulas -// -// Even though it's Wikipedia, I trust it as I helped write that table. -// -// XXX Replace with a proper reference. - -// Compute the propagated uncertainty from the multiplication of two uncertain -// values, `A +/- A_unc` and `B +/- B_unc`. Given `f = AB` or `f = A/B`, where -// `A != 0` and `B != 0`, the uncertainty in `f` is approximately: -// -// f_unc = abs(f) * sqrt((A_unc / A) ^ 2 + (B_unc / B) ^ 2) -// -template -__host__ __device__ T uncertainty_multiplicative(T const& f, T const& A, T const& A_unc, T const& B, T const& B_unc) -{ - return std::abs(f) * std::sqrt((A_unc / A) * (A_unc / A) + (B_unc / B) * (B_unc / B)); -} - -// Compute the propagated uncertainty from addition of two uncertain values, -// `A +/- A_unc` and `B +/- B_unc`. Given `f = cA + dB` (where `c` and `d` are -// certain constants), the uncertainty in `f` is approximately: -// -// f_unc = sqrt(c ^ 2 * A_unc ^ 2 + d ^ 2 * B_unc ^ 2) -// -template -__host__ __device__ T uncertainty_additive(T const& c, T const& A_unc, T const& d, T const& B_unc) -{ - return std::sqrt((c * c * A_unc * A_unc) + (d * d * B_unc * B_unc)); -} - -/////////////////////////////////////////////////////////////////////////////// - -// Return the significant digit of `x`. The result is the number of digits -// after the decimal place to round to (negative numbers indicate rounding -// before the decimal place) -template -int find_significant_digit(T x) -{ - if (x == T(0)) - { - return T(0); - } - return -int(std::floor(std::log10(std::abs(x)))); -} - -// Round `x` to `ndigits` after the decimal place (Python-style). -template -T round_to_precision(T x, N ndigits) -{ - double m = (x < 0.0) ? -1.0 : 1.0; - double pwr = std::pow(T(10.0), ndigits); - return (std::floor(x * m * pwr + 0.5) / pwr) * m; -} - -/////////////////////////////////////////////////////////////////////////////// - -void print_experiment_header() -{ // {{{ - std::cout - << "Thrust Version" - << "," - << "Algorithm" - << "," - << "Element Type" - << "," - << "Element Size" - << "," - << "Elements per Trial" - << "," - << "Total Input Size" - << "," - << "STL Trials" - << "," - << "STL Average Walltime" - << "," - << "STL Walltime Uncertainty" - << "," - << "STL Average Throughput" - << "," - << "STL Throughput Uncertainty" - << "," - << "Thrust Trials" - << "," - << "Thrust Average Walltime" - << "," - << "Thrust Walltime Uncertainty" - << "," - << "Thrust Average Throughput" - << "," - << "Thrust Throughput Uncertainty" -#if defined(HAVE_TBB) - << "," - << "TBB Trials" - << "," - << "TBB Average Walltime" - << "," - << "TBB Walltime Uncertainty" - << "," - << "TBB Average Throughput" - << "," - << "TBB Throughput Uncertainty" -#endif - << std::endl; - - std::cout - << "" // Thrust Version. - << "," - << "" // Algorithm. - << "," - << "" // Element Type. - << "," - << "bits/element" // Element Size. - << "," - << "elements" // Elements per Trial. - << "," - << "MiBs" // Total Input Size. - << "," - << "trials" // STL Trials. - << "," - << "secs" // STL Average Walltime. - << "," - << "secs" // STL Walltime Uncertainty. - << "," - << "elements/sec" // STL Average Throughput. - << "," - << "elements/sec" // STL Throughput Uncertainty. - << "," - << "trials" // Thrust Trials. - << "," - << "secs" // Thrust Average Walltime. - << "," - << "secs" // Thrust Walltime Uncertainty. - << "," - << "elements/sec" // Thrust Average Throughput. - << "," - << "elements/sec" // Thrust Throughput Uncertainty. -#if defined(HAVE_TBB) - << "," - << "trials" // TBB Trials. - << "," - << "secs" // TBB Average Walltime. - << "," - << "secs" // TBB Walltime Uncertainty. - << "," - << "elements/sec" // TBB Average Throughput. - << "," - << "elements/sec" // TBB Throughput Uncertainty. -#endif - << std::endl; -} // }}} - -/////////////////////////////////////////////////////////////////////////////// - -struct experiment_results -{ - double const average_time; // Arithmetic mean of trial times in seconds. - double const stdev_time; // Sample standard deviation of trial times. - - experiment_results(double average_time_, double stdev_time_) - : average_time(average_time_) - , stdev_time(stdev_time_) - {} -}; - -/////////////////////////////////////////////////////////////////////////////// - -template