Testing/Performance:
* Switch from using CUDA events for timing to using an OS-based steady clock.
* Change default parameter sweep for `bench.cu`.
Bug
Bug
git-commit
git-author Bryce Adelstein Lelbach aka wash
VDVS:
Jobs:
[git-p4: depot-paths = "//sw/gpgpu/thrust/": change = 24223401]
MODIFY
internal/benchmark/bench.cu
internal/benchmark/bench.cu
#include <iostream>
+#include <cassert>
#include <cstdlib> // For `atoi`.
#include <climits> // For CHAR_BIT.
#include <cmath> // For `sqrt` and `abs`.
#endif
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
- #include <cuda_runtime.h> // For cudaSetDevice.
+ #include <thrust/system_error.h> // For `thrust::system_error`
+ #include <thrust/system/cuda/error.h> // For `thrust::cuda_category`
#endif
// We don't use THRUST_PP_STRINGIZE and THRUST_PP_CAT because they are new, and
@@ -391,6 +393,7 @@ struct experiment_driver
);
#endif
+/*
stl_average_walltime = round_to_precision(
stl_average_walltime, stl_walltime_precision
);
@@ -414,6 +417,7 @@ struct experiment_driver
tbb_walltime_uncertainty, tbb_walltime_precision
);
#endif
+*/
// Round the average throughput and throughput uncertainty to the
// significant figure of the throughput uncertainty.
@@ -432,6 +436,7 @@ struct experiment_driver
);
#endif
+/*
stl_average_throughput = round_to_precision(
stl_average_throughput, stl_throughput_precision
);
@@ -455,6 +460,7 @@ struct experiment_driver
tbb_throughput_uncertainty, tbb_throughput_precision
);
#endif
+*/
std::cout << THRUST_VERSION // Thrust Version.
<< "," << test_name // Algorithm.
@@ -522,7 +528,7 @@ private:
// Generate random input for next trial.
trial.setup(elements);
- timer e;
+ steady_timer e;
// Benchmark.
e.start();
@@ -743,6 +749,11 @@ struct sort_tester
void operator()()
{
thrust::sort(this->input.begin(), this->input.end());
+ #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+ cudaError_t err = cudaDeviceSynchronize();
+ if (err != cudaSuccess)
+ throw thrust::error_code(err, thrust::cuda_category());
+ #endif
}
};
@@ -782,6 +793,11 @@ struct transform_inplace_tester
this->input.begin(), this->input.end(), this->input.begin()
, thrust::negate<T>()
);
+ #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+ cudaError_t err = cudaDeviceSynchronize();
+ if (err != cudaSuccess)
+ throw thrust::error_code(err, thrust::cuda_category());
+ #endif
}
};
@@ -818,6 +834,11 @@ struct inclusive_scan_inplace_tester
thrust::inclusive_scan(
this->input.begin(), this->input.end(), this->input.begin()
);
+ #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+ cudaError_t err = cudaDeviceSynchronize();
+ if (err != cudaSuccess)
+ throw thrust::error_code(err, thrust::cuda_category());
+ #endif
}
};
@@ -850,6 +871,11 @@ struct copy_tester
void operator()()
{
thrust::copy(this->input.begin(), this->input.end(), this->input.begin());
+ #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+ cudaError_t err = cudaDeviceSynchronize();
+ if (err != cudaSuccess)
+ throw thrust::error_code(err, thrust::cuda_category());
+ #endif
}
};
@@ -1177,12 +1203,12 @@ int main(int argc, char** argv)
//run_core_primitives_experiments< 1LLU << 21LLU , 4 , 16 >();
//run_core_primitives_experiments< 1LLU << 22LLU , 4 , 16 >();
//run_core_primitives_experiments< 1LLU << 23LLU , 4 , 16 >();
- run_core_primitives_experiments< 1LLU << 24LLU , 3 , 8 >();
- run_core_primitives_experiments< 1LLU << 25LLU , 3 , 8 >();
-//run_core_primitives_experiments< 1LLU << 26LLU , 3 , 8 >();
-//run_core_primitives_experiments< 1LLU << 27LLU , 3 , 8 >();
-//run_core_primitives_experiments< 1LLU << 28LLU , 3 , 8 >();
-//run_core_primitives_experiments< 1LLU << 29LLU , 3 , 8 >();
+//run_core_primitives_experiments< 1LLU << 24LLU , 4 , 16 >();
+//run_core_primitives_experiments< 1LLU << 25LLU , 4 , 16 >();
+ run_core_primitives_experiments< 1LLU << 26LLU , 4 , 16 >();
+ run_core_primitives_experiments< 1LLU << 27LLU , 4 , 16 >();
+//run_core_primitives_experiments< 1LLU << 28LLU , 4 , 16 >();
+//run_core_primitives_experiments< 1LLU << 29LLU , 4 , 16 >();
return 0;
}

MODIFY
internal/benchmark/combine_benchmark_results.py
internal/benchmark/combine_benchmark_results.py
@@ -724,13 +724,13 @@ class record_aggregator(object):
# uncertainty and insert the combined values into the results.
sigdig = find_significant_digit(combined_sample_standard_deviation)
- combined_arithmetic_mean = round_with_int_conversion(
- combined_arithmetic_mean, sigdig
- )
+# combined_arithmetic_mean = round_with_int_conversion(
+# combined_arithmetic_mean, sigdig
+# )
- combined_sample_standard_deviation = round_with_int_conversion(
- combined_sample_standard_deviation, sigdig
- )
+# combined_sample_standard_deviation = round_with_int_conversion(
+# combined_sample_standard_deviation, sigdig
+# )
combined_dependent_values[quantity] = combined_arithmetic_mean
combined_dependent_values[uncertainty] = combined_sample_standard_deviation

MODIFY
internal/benchmark/compare_benchmark_results.py
internal/benchmark/compare_benchmark_results.py
@@ -581,7 +581,7 @@ def process_program_arguments():
ap.add_argument(
"-t", "--change-threshold",
help = ("Treat relative changes less than this amount (a percentage) as "
- "statistically insignificant. The default is 5%."),
+ "statistically insignificant. The default is 5%%."),
action = "store", type = float, default = 5,
metavar = "PERCENTAGE"
)
@@ -998,13 +998,13 @@ class record_aggregator(object):
# uncertainty and insert the combined values into the results.
sigdig = find_significant_digit(combined_sample_standard_deviation)
- combined_arithmetic_mean = round_with_int_conversion(
- combined_arithmetic_mean, sigdig
- )
+# combined_arithmetic_mean = round_with_int_conversion(
+# combined_arithmetic_mean, sigdig
+# )
- combined_sample_standard_deviation = round_with_int_conversion(
- combined_sample_standard_deviation, sigdig
- )
+# combined_sample_standard_deviation = round_with_int_conversion(
+# combined_sample_standard_deviation, sigdig
+# )
combined_dependent_values[quantity] = combined_arithmetic_mean
combined_dependent_values[uncertainty] = combined_sample_standard_deviation
@@ -1243,12 +1243,12 @@ with io_manager(args.baseline_input_file,
find_significant_digit(abs_change_unc),
)
- abs_change = round_with_int_conversion(
- abs_change, abs_change_sigdig
- )
- abs_change_unc = round_with_int_conversion(
- abs_change_unc, abs_change_sigdig
- )
+# abs_change = round_with_int_conversion(
+# abs_change, abs_change_sigdig
+# )
+# abs_change_unc = round_with_int_conversion(
+# abs_change_unc, abs_change_sigdig
+# )
except:
# Any value errors should be due to NaNs returned by
# `percent_change_uncertainty` because quantities or change in
@@ -1261,12 +1261,12 @@ with io_manager(args.baseline_input_file,
find_significant_digit(per_change_unc)
)
- per_change = round_with_int_conversion(
- per_change, per_change_sigdig
- )
- per_change_unc = round_with_int_conversion(
- per_change_unc, per_change_sigdig
- )
+# per_change = round_with_int_conversion(
+# per_change, per_change_sigdig
+# )
+# per_change_unc = round_with_int_conversion(
+# per_change_unc, per_change_sigdig
+# )
except:
# Any value errors should be due to NaNs returned by
# `percent_change_uncertainty` because quantities or change in

MODIFY
internal/benchmark/timer.h
internal/benchmark/timer.h
#pragma once
-#include <cuda.h>
+#include <cassert>
# define CUDA_SAFE_CALL_NO_SYNC( call) do { \
cudaError err = call; \
if( cudaSuccess != err) { \
- fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", \
+ fprintf(stderr, "CUDA error in file '%s' in line %i : %s.\n", \
__FILE__, __LINE__, cudaGetErrorString( err) ); \
exit(EXIT_FAILURE); \
} } while (0)
CUDA_SAFE_CALL_NO_SYNC(call); \
cudaError err = cudaDeviceSynchronize(); \
if( cudaSuccess != err) { \
- fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", \
+ fprintf(stderr, "CUDA error in file '%s' in line %i : %s.\n", \
__FILE__, __LINE__, cudaGetErrorString( err) ); \
exit(EXIT_FAILURE); \
} } while (0)
-class timer
+class cuda_timer
{
- cudaEvent_t _start;
- cudaEvent_t _end;
+ cudaEvent_t start_;
+ cudaEvent_t stop_;
public:
- timer()
+ cuda_timer()
{
- CUDA_SAFE_CALL(cudaEventCreate(&_start));
- CUDA_SAFE_CALL(cudaEventCreate(&_end));
+ CUDA_SAFE_CALL(cudaEventCreate(&start_));
+ CUDA_SAFE_CALL(cudaEventCreate(&stop_));
}
- ~timer()
+ ~cuda_timer()
{
- CUDA_SAFE_CALL(cudaEventDestroy(_start));
- CUDA_SAFE_CALL(cudaEventDestroy(_end));
+ CUDA_SAFE_CALL(cudaEventDestroy(start_));
+ CUDA_SAFE_CALL(cudaEventDestroy(stop_));
}
void start()
{
- CUDA_SAFE_CALL(cudaEventRecord(_start, 0));
+ CUDA_SAFE_CALL(cudaEventRecord(start_, 0));
}
void stop()
{
- CUDA_SAFE_CALL(cudaEventRecord(_end, 0));
- CUDA_SAFE_CALL(cudaEventSynchronize(_end));
+ CUDA_SAFE_CALL(cudaEventRecord(stop_, 0));
+ CUDA_SAFE_CALL(cudaEventSynchronize(stop_));
}
double milliseconds_elapsed()
{
float elapsed_time;
- CUDA_SAFE_CALL(cudaEventElapsedTime(&elapsed_time, _start, _end));
+ CUDA_SAFE_CALL(cudaEventElapsedTime(&elapsed_time, start_, stop_));
return elapsed_time;
}
@@ -61,4 +61,69 @@ class timer
}
};
+#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
+#include <windows.h>
+
+class steady_timer
+{
+ LARGE_INTEGER frequency_; // Cached to avoid system calls.
+ LARGE_INTEGER start_;
+ LARGE_INTEGER stop_;
+
+ public:
+ steady_timer() : start_(), stop_(), frequency_()
+ {
+ BOOL const r = QueryPerformanceFrequency(&frequency_);
+ assert(0 != r);
+ }
+
+ void start()
+ {
+ BOOL const r = QueryPerformanceCounter(&start_);
+ assert(0 != r);
+ }
+
+ void stop()
+ {
+ BOOL const r = QueryPerformanceCounter(&stop_);
+ assert(0 != r);
+ }
+
+ double seconds_elapsed()
+ {
+ return double(stop_.QuadPart - start_.QuadPart)
+ / double(frequency_.QuadPart);
+ }
+};
+#else
+#include <time.h>
+
+class steady_timer
+{
+ timespec start_;
+ timespec stop_;
+
+ public:
+ steady_timer() : start_(), stop_() {}
+
+ void start()
+ {
+ int const r = clock_gettime(CLOCK_MONOTONIC, &start_);
+ assert(0 == r);
+ }
+
+ void stop()
+ {
+ int const r = clock_gettime(CLOCK_MONOTONIC, &stop_);
+ assert(0 == r);
+ }
+
+ double seconds_elapsed()
+ {
+ return double(stop_.tv_sec - start_.tv_sec)
+ + double(stop_.tv_nsec - start_.tv_nsec) * 1.0e-9;
+ }
+};
+#endif
+

Testing/Performance:
* Switch from using CUDA events for timing to using an OS-based steady clock.
* Change default parameter sweep for `bench.cu`.
Bug 2011463
Bug 200397103
git-commit 691763c6dbd35b1b8f340909c4b390abb83f8117
git-author Bryce Adelstein Lelbach aka wash <brycelelbach@gmail.com>
VDVS: http://ausdvs.nvidia.com/Build_Results?virtualId=1000165171&which_page=current_build

Jobs: 200397103-2006 2011463-2006
[git-p4: depot-paths = "//sw/gpgpu/thrust/": change = 24223401]