Add internal benchmark to build
MODIFY
CMakeLists.txt
CMakeLists.txt
@@ -37,6 +37,7 @@ include(cmake/VerifyCompiler.cmake)
option(DISABLE_WERROR "Disable building with Werror" ON)
option(BUILD_TEST "Build tests" OFF)
option(BUILD_EXAMPLES "Build examples" OFF)
+option(BUILD_BENCHMARKS "Build benchmarks" OFF)
option(DOWNLOAD_ROCPRIM "Download rocPRIM and do not search for rocPRIM package" OFF)
# Set CXX flags
@@ -72,12 +73,16 @@ if(BUILD_TEST)
add_subdirectory(testing)
endif()
-# Tests
+# Examples
if(BUILD_EXAMPLES)
- message (STATUS " BUILD_EXAMPLES : ON")
add_subdirectory(examples)
endif()
+# Benchmarks
+if(BUILD_BENCHMARKS)
+ add_subdirectory(internal/benchmark)
+endif()
+
# Package
set(CPACK_DEBIAN_ARCHIVE_TYPE "gnutar")

ADD
internal/benchmark/CMakeLists.txt
+# ########################################################################
+# Copyright 2020 Advanced Micro Devices, Inc.
+# ########################################################################
+
+function(add_thrust_benchmark BENCHMARK)
+ set(BENCHMARK_SOURCE "${BENCHMARK}.cu")
+ set(BENCHMARK_TARGET "benchmark_thrust_${BENCHMARK}")
+ set_source_files_properties(${BENCHMARK_SOURCE} PROPERTIES LANGUAGE CXX COMPILE_FLAGS "-x c++")
+ add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE})
+
+ target_link_libraries(${BENCHMARK_TARGET}
+ PRIVATE
+ rocthrust
+ )
+ foreach(amdgpu_target ${AMDGPU_TARGETS})
+ target_link_libraries(${BENCHMARK_TARGET}
+ PRIVATE
+ --amdgpu-target=${amdgpu_target}
+ )
+ endforeach()
+ set_target_properties(${BENCHMARK_TARGET}
+ PROPERTIES
+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/benchmarks/"
+ )
+endfunction()
+
+# ****************************************************************************
+# Benchmarks
+# ****************************************************************************
+message (STATUS "Building benchmarks")
+
+add_thrust_benchmark("bench")

MODIFY
internal/benchmark/bench.cu
internal/benchmark/bench.cu
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/scan.h>
+#include <thrust/version.h>
#include <algorithm>
#include <numeric>
#include "tbb_algos.h"
#endif
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC // HCC compiler
+#ifdef __HIP_DEVICE_COMPILE__
+ using ::abs;
+ using ::sqrt;
+#else
+ using std::abs;
+ using std::sqrt;
+#endif
+#else // Not HCC devce compiler
+ using std::abs;
+ using std::sqrt;
+#endif
+
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
+ #include <thrust/system_error.h> // For `thrust::system_error`
+ #include <thrust/system/hip/error.h> // For `thrust::hip_category`
+#endif
+
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
#include <thrust/system_error.h> // For `thrust::system_error`
#include <thrust/system/cuda/error.h> // For `thrust::cuda_category`
@@ -173,7 +192,7 @@ T sample_standard_deviation(InputIt first, InputIt last, T average)
value_and_count<T> vc
= thrust::transform_reduce(first, last, transform, init_vc, reduce_vc);
- return std::sqrt(vc.value / T(vc.count - 1));
+ return sqrt(vc.value / T(vc.count - 1));
}
///////////////////////////////////////////////////////////////////////////////
@@ -200,8 +219,8 @@ T uncertainty_multiplicative(
, 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));
+ return abs(f)
+ * sqrt((A_unc / A) * (A_unc / A) + (B_unc / B) * (B_unc / B));
}
// Compute the propagated uncertainty from addition of two uncertain values,
@@ -217,7 +236,7 @@ T uncertainty_additive(
, T const& d, T const& B_unc
)
{
- return std::sqrt((c * c * A_unc * A_unc) + (d * d * B_unc * B_unc));
+ return sqrt((c * c * A_unc * A_unc) + (d * d * B_unc * B_unc));
}
///////////////////////////////////////////////////////////////////////////////

MODIFY
internal/benchmark/timer.h
internal/benchmark/timer.h
#pragma once
#include <cassert>
+#include <thrust/detail/config.h>
+
+#if (THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC)
+
+#define HIP_CHECK(condition) \
+ { \
+ hipError_t error = condition; \
+ if(error != hipSuccess){ \
+ std::cout << "HIP error: " << error << " line: " << __LINE__ << std::endl; \
+ exit(error); \
+ } \
+ }
+
+class hip_timer
+{
+ hipEvent_t start_;
+ hipEvent_t stop_;
+
+ public:
+ hip_timer()
+ {
+ HIP_CHECK(hipEventCreate(&start_));
+ HIP_CHECK(hipEventCreate(&stop_));
+ }
+
+ ~hip_timer()
+ {
+ HIP_CHECK(hipEventDestroy(start_));
+ HIP_CHECK(hipEventDestroy(stop_));
+ }
+
+ void start()
+ {
+ HIP_CHECK(hipEventRecord(start_, 0));
+ }
+
+ void stop()
+ {
+ HIP_CHECK(hipEventRecord(stop_, 0));
+ HIP_CHECK(hipEventSynchronize(stop_));
+ }
+
+ double milliseconds_elapsed()
+ {
+ float elapsed_time;
+ HIP_CHECK(hipEventElapsedTime(&elapsed_time, start_, stop_));
+ return elapsed_time;
+ }
+
+ double seconds_elapsed()
+ {
+ return milliseconds_elapsed() / 1000.0;
+ }
+};
+
+#endif
+
+#if (THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC)
# define CUDA_SAFE_CALL_NO_SYNC( call) do { \
cudaError err = call; \
@@ -61,6 +119,8 @@ class cuda_timer
}
};
+#endif
+
#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
#include <windows.h>
@@ -125,5 +185,3 @@ class steady_timer
}
};
#endif
-
-

Add internal benchmark to build