diff --git a/CMakeLists.txt b/CMakeLists.txt index dfa42689..808a9113 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,7 +12,8 @@ else() endif() # Thrust project -project(rocthrust LANGUAGES CXX) +# Note: C is required here for dependencies +project(rocthrust LANGUAGES CXX C) #Adding CMAKE_PREFIX_PATH list( APPEND CMAKE_PREFIX_PATH /opt/rocm/llvm /opt/rocm ${ROCM_PATH} ) diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 218f6122..5720f2da 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -62,4 +62,50 @@ if(BUILD_TEST) ) find_package(GTest REQUIRED CONFIG PATHS ${GTEST_ROOT}) endif() + + # SQlite (for run-to-run bitwise-reproducibility tests) + # Note: SQLite 3.36.0 enabled the backup API by default, which we need + # for cache serialization. We also want to use a static SQLite, + # and distro static libraries aren't typically built + # position-independent. + include( FetchContent ) + + if(DEFINED ENV{SQLITE_3_43_2_SRC_URL}) + set(SQLITE_3_43_2_SRC_URL_INIT $ENV{SQLITE_3_43_2_SRC_URL}) + else() + set(SQLITE_3_43_2_SRC_URL_INIT https://www.sqlite.org/2023/sqlite-amalgamation-3430200.zip) + endif() + set(SQLITE_3_43_2_SRC_URL ${SQLITE_3_43_2_SRC_URL_INIT} CACHE STRING "Location of SQLite source code") + set(SQLITE_SRC_3_43_2_SHA3_256 af02b88cc922e7506c6659737560c0756deee24e4e7741d4b315af341edd8b40 CACHE STRING "SHA3-256 hash of SQLite source code") + + # embed SQLite + if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.24) + # use extract timestamp for fetched files instead of timestamps in the archive + cmake_policy(SET CMP0135 NEW) + endif() + + message("Downloading SQLite.") + FetchContent_Declare(sqlite_local + URL ${SQLITE_3_43_2_SRC_URL} + URL_HASH SHA3_256=${SQLITE_SRC_3_43_2_SHA3_256} + ) + FetchContent_MakeAvailable(sqlite_local) + + add_library(sqlite3 OBJECT ${sqlite_local_SOURCE_DIR}/sqlite3.c) + target_include_directories(sqlite3 PUBLIC ${sqlite_local_SOURCE_DIR}) + set_target_properties( sqlite3 PROPERTIES + C_VISIBILITY_PRESET "hidden" + VISIBILITY_INLINES_HIDDEN ON + POSITION_INDEPENDENT_CODE ON + LINKER_LANGUAGE CXX + ) + + # We don't need extensions, and omitting them from SQLite removes the + # need for dlopen/dlclose from within rocThrust. + # We also don't need the shared cache, and omitting it yields some performance improvements. + target_compile_options( + sqlite3 + PRIVATE -DSQLITE_OMIT_LOAD_EXTENSION + PRIVATE -DSQLITE_OMIT_SHARED_CACHE + ) endif() diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 483d5c97..5c2f9f8b 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -54,11 +54,14 @@ function(add_rocthrust_test TEST) target_include_directories(${TEST_TARGET} SYSTEM BEFORE PUBLIC $ + ${sqlite_local_SOURCE_DIR} ) target_link_libraries(${TEST_TARGET} PRIVATE rocthrust roc::rocprim_hip + PUBLIC + sqlite3 ) if (TARGET GTest::GTest) target_link_libraries(${TEST_TARGET} diff --git a/test/bitwise_repro/bwr_db.hpp b/test/bitwise_repro/bwr_db.hpp new file mode 100644 index 00000000..448bc221 --- /dev/null +++ b/test/bitwise_repro/bwr_db.hpp @@ -0,0 +1,320 @@ +// Copyright (C) 2024 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#ifndef BWR_DB_HPP +#define BWR_DB_HPP + +#include +#include + +#include "bwr_utils.hpp" + +/*! \brief Database that can be used to store function call information between runs. +* This allows us to test whether the input has changed between runs. +*/ +class BitwiseReproDB +{ +public: + /*! \brief Enum class used to control the mode the database operates in. + * There are two reasons that a row might not be found in the database: + * 1. There is a run-to-run reproducibility error. + * 2. The database entries for this architecture/rocm version/rocThrust version + * haven't been generated yet. + * + * This enum class allows us to distinguish between these two cases. + * * + * In test mode, if an entry is not found, it is not inserted. This allows + * errors to be detected (case 1 above). + * + * In generate_mode, if an entry is not found, it is inserted (case 2 above). + * No run-to-run errors will be detected while in generate mode (the match functions + * will always report that a match was found). + */ + enum class Mode + { + test_mode, + generate_mode + }; + + /*! \brief Database constructor. Will create the SQLite database file if it doesn't already exist (and db_path is not null). + * \param db_path Path to the database file (eg. "./repro.db"). If null, database won't be created. + * \param mode The database mode (generate or test). See the enum class above for details. + */ + BitwiseReproDB(const char* db_path, const BitwiseReproDB::Mode mode) : + m_db_conn(nullptr), + m_insert_stmt(nullptr), + m_match_stmt(nullptr), + m_mode(mode) + { + if (!db_path) + throw std::runtime_error("No database path given (ROCTHRUST_REPRO_DB_PATH environment variable is not set)."); + + int ret = sqlite3_open(db_path, &m_db_conn); + if (ret != SQLITE_OK) + throw std::runtime_error("Cannot open run-to-run bitwise reproducibility database: " + std::string(db_path)); + + // Access to a database file may occur in parallel. + // Increase default sqlite timeout, so diferent process + // can wait for one another. + sqlite3_busy_timeout(m_db_conn, 30000); + + // Set sqlite3 engine to WAL mode to avoid potential deadlocks with multiple + // concurrent processes (if a deadlock occurs, the busy timeout is not honored). + ret = sqlite3_exec(m_db_conn, "PRAGMA journal_mode = WAL", nullptr, nullptr, nullptr); + if(ret != SQLITE_OK) + throw std::runtime_error("Error setting WAL mode: " + std::string(sqlite3_errmsg(m_db_conn))); + + // Create the rocthrust_test_run table if it doesn't already exist. + ret = sqlite3_exec(m_db_conn, + BitwiseReproDB::get_create_table_sql().c_str(), + nullptr, + nullptr, + nullptr); + if(ret != SQLITE_OK) + throw std::runtime_error("Error creating table: " + + std::string(sqlite3_errmsg(m_db_conn))); + + // Initialize prepared statements. + prepare_match_stmt(); + prepare_insert_stmt(); + } + + /*! \brief Destructor - cleans up and closes the database connection. + */ + ~BitwiseReproDB() + { + sqlite3_finalize(m_insert_stmt); + sqlite3_finalize(m_match_stmt); + sqlite3_close(m_db_conn); + } + + /*! \brief Given a pair of input and output "tokens" (which uniquely identify a function call), + * looks for a match in the database. If the DB is in generate mode and a match is not found, + * the (input, output) token pair will be inserted into the database. In test mode, no insertion + * is performed. + * + * \param input_token String that uniquely identifies a function call's inputs. See bwr_utils.hpp for details. + * \param output_token String that uniquely identifies a function call's outputs. See bwr_utils.hpp for details. + * \param match_found [out] In test mode, set to true if the (input, output) token pair is already in the database. + * In generate mode, rows are inserted if they don't already exist, and this is always set to true. + * \param inserted [out] Set to true if a row was inserted. + */ + void match( + const std::string& input_token, + const std::string& output_token, + bool& match_found, + bool& inserted) + { + match_found = false; + inserted = false; + + // A test_run is a convenience struct that encapsulates all the information in a single row of the database table. + // Create one using the given input/output pair. + const rocthrust_test_run test_run(input_token, output_token); + + // Do a select to check for a matching existing row. + const int match_count = select(test_run); + // Note: Because of our database constraints (the unique index), + // we know that match_count will either be 0 or 1 here. + match_found = (match_count == 1); + + // Only insert if we are in generate mode and an entry does + // not already exist. + if (m_mode == Mode::generate_mode && !match_found) + { + try + { + inserted = insert(test_run); + // If the insertion was successful, set match to true, since a matching row now exists. + match_found = inserted; + } + catch(const std::exception& e) + { + std::cerr << e.what() << '\n'; + std::cerr << "input_token: " << input_token << std::endl; + std::cerr << "output_token: " << output_token << std::endl; + std::cerr << "match_found: " << int(match_found) << std::endl; + std::cerr << "inserted: " << int(inserted) << std::endl; + } + } + } + + /*! \brief See above. This overload exists for convenience - you can call it when you don't need to check if anything was inserted. + * \return A bool indicating whether or not a match was found. + */ + bool match( + const std::string& input_token, + const std::string& output_token + ) + { + bool match_found; + bool inserted; + match(input_token, output_token, match_found, inserted); + + return match_found; + } + +private: + struct rocthrust_test_run + { + rocthrust_test_run( + const std::string& input_token, + const std::string& output_token, + const std::string& rocm_version, + const std::string& rocthrust_version, + const std::string& gpu_arch + ) : input_token(input_token), + output_token(output_token), + rocm_version(rocm_version), + rocthrust_version(rocthrust_version), + gpu_arch(gpu_arch) + { + } + + rocthrust_test_run( + const std::string& input_token, + const std::string& output_token + ) : input_token(input_token), + output_token(output_token), + rocm_version(bwr_utils::get_rocm_version()), + rocthrust_version(bwr_utils::get_rocthrust_version()), + gpu_arch(bwr_utils::get_gpu_arch()) + { + } + + std::string input_token; + std::string output_token; + std::string rocm_version; + std::string rocthrust_version; + std::string gpu_arch; + }; + + static const std::string get_create_table_sql() + { + return "CREATE TABLE IF NOT EXISTS rocthrust_test_run(" + "input_token TEXT NOT NULL, " + "output_token TEXT NOT NULL, " + "rocm_version TEXT NOT NULL, " + "rocthrust_version TEXT NOT NULL, " + "gpu_arch TEXT NOT NULL);" + "CREATE UNIQUE INDEX IF NOT EXISTS id_index_unique_run ON rocthrust_test_run(" + "input_token, rocm_version, rocthrust_version, gpu_arch);"; + } + + static const std::string get_insert_sql() + { + return "INSERT INTO rocthrust_test_run(" + "input_token, output_token, rocm_version, rocthrust_version, gpu_arch) " + "VALUES (?, ?, ?, ?, ?);"; + } + + static const std::string get_match_sql() + { + return "SELECT COUNT(*) FROM rocthrust_test_run WHERE " + "input_token = ? AND output_token = ? AND rocm_version = ? AND rocthrust_version = ? AND gpu_arch = ?;"; + } + + void prepare_match_stmt() + { + static const std::string match_sql = get_match_sql(); + + const int ret = sqlite3_prepare_v2(m_db_conn, match_sql.c_str(), -1, &m_match_stmt, nullptr); + if (ret != SQLITE_OK) + throw std::runtime_error("Cannot prepare match statement: " + + std::string(sqlite3_errmsg(m_db_conn))); + } + + void prepare_insert_stmt() + { + static const std::string insert_sql = get_insert_sql(); + + const int ret = sqlite3_prepare_v2(m_db_conn, insert_sql.c_str(), -1, &m_insert_stmt, nullptr); + if (ret != SQLITE_OK) + throw std::runtime_error("Cannot prepare insert statement: " + + std::string(sqlite3_errmsg(m_db_conn))); + } + + int select(const rocthrust_test_run& test_run) + { + int count = 0; + bind_match_stmt(test_run); + const int ret = sqlite3_step(m_match_stmt); + if (ret != SQLITE_ROW) + { + throw std::runtime_error(std::string("Error executing select statement: ") + + std::string(sqlite3_errmsg(m_db_conn))); + } + + // Note: select indices start at 0 + count = sqlite3_column_int(m_match_stmt, 0); + + return count; + } + + bool insert(const rocthrust_test_run& test_run) + { + bind_insert_stmt(test_run); + const int ret = sqlite3_step(m_insert_stmt); + const bool inserted = (ret == SQLITE_DONE); + if (!inserted) + throw std::runtime_error(std::string("Error executing insert statement: ") + + std::string(sqlite3_errmsg(m_db_conn))); + + return inserted; + } + + void bind_insert_stmt(const rocthrust_test_run& test_run) + { + // Note: bind indices start at 1 + sqlite3_reset(m_insert_stmt); + bind_text(m_insert_stmt, test_run.input_token, 1); + bind_text(m_insert_stmt, test_run.output_token, 2); + bind_text(m_insert_stmt, test_run.rocm_version, 3); + bind_text(m_insert_stmt, test_run.rocthrust_version, 4); + bind_text(m_insert_stmt, test_run.gpu_arch, 5); + } + + void bind_match_stmt(const rocthrust_test_run& test_run) + { + // Note: bind indices start at 1 + sqlite3_reset(m_match_stmt); + bind_text(m_match_stmt, test_run.input_token, 1); + bind_text(m_match_stmt, test_run.output_token, 2); + bind_text(m_match_stmt, test_run.rocm_version, 3); + bind_text(m_match_stmt, test_run.rocthrust_version, 4); + bind_text(m_match_stmt, test_run.gpu_arch, 5); + } + + void bind_text(sqlite3_stmt* stmt, const std::string& text, const int index) + { + const int ret = sqlite3_bind_text(stmt, index, text.c_str(), -1, SQLITE_TRANSIENT); + if (ret != SQLITE_OK) + throw std::runtime_error(std::string("Error binding text field in insert statement:\n" + "index: " + std::to_string(index) + "\n" + "value: " + text)); + } + + sqlite3* m_db_conn; + sqlite3_stmt* m_insert_stmt; + sqlite3_stmt* m_match_stmt; + Mode m_mode; +}; + +#endif // BWR_DB_HPP \ No newline at end of file diff --git a/test/bitwise_repro/bwr_utils.hpp b/test/bitwise_repro/bwr_utils.hpp new file mode 100644 index 00000000..ffd001e0 --- /dev/null +++ b/test/bitwise_repro/bwr_utils.hpp @@ -0,0 +1,503 @@ +// Copyright (C) 2024 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#ifndef BWR_UTILS_HPP +#define BWR_UTILS_HPP + +#include "../thrust/include/rocthrust_version.hpp" +#include +#include + +// These macros are be used to get a stringified version of a type name (see get_typename_str below) . +// If data_type matches key_type, then we return a stringified version of key_type. +#define STRING(s) #s +#define IF_TYPE_TEST(data_type, key_type) if (std::is_same::value) return STRING(key_type); +#define ELSE_IF_TYPE_TEST(data_type, key_type) else IF_TYPE_TEST(data_type, key_type) + +namespace bwr_utils +{ + +// These members don't need to be visible outside this file. +namespace +{ + // The separator character used in the ROCm and rocThrust versions. + static const std::string ver_sep = std::string("."); + + /*! \brief This function returns a string version of the template parameter typename it's passed. + * \note Because C++ has no reflection/introspection, we must manually define strings for each + * type we want to be able to do this for. + * + * Currently, only types used in tests in test/test_reproducibility.cpp are supported here. + * Input types passed passed to TokenHelper (below) also need to be defined here. + * + * \tparam The type to find a string name for. + * \return String version of the type, or empty string if not found. + */ + template + std::string get_typename_str() + { + IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::host_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, thrust::device_vector) + ELSE_IF_TYPE_TEST(T, short) + ELSE_IF_TYPE_TEST(T, int) + ELSE_IF_TYPE_TEST(T, long long) + ELSE_IF_TYPE_TEST(T, unsigned short) + ELSE_IF_TYPE_TEST(T, unsigned int) + ELSE_IF_TYPE_TEST(T, unsigned long long) + ELSE_IF_TYPE_TEST(T, float) + ELSE_IF_TYPE_TEST(T, double) + else + throw std::runtime_error("Unable to lookup type name in get_typename_str()"); + + return ""; + } + + /*! \brief Builds a table that's used to cache values for the CRC algorithm. + * \note: We're assuming the system is little-endian here. + * \param table Pointer to the buffer to use to store the table (must be of size 256). + */ + void make_table(uint32_t* table) + { + table[0] = 0; + uint32_t crc = 1; + uint8_t i = 128; + do + { + // Note: The constant below is the polynomial for the + // CRC-32/ISO-HDLC version of the 32-bit algorithm. + if (crc & 1) + crc = (crc >> 1) ^ 0xedb88320; + else + crc >>= 1; + + for (uint32_t j = 0; j < 256; j += 2 * i) + { + table[i + j] = crc ^ table[j]; + } + i >>= 1; + } while (i > 0); + } + + /*! \brief Performs a 32-bit cyclic redundancy check on the data it's passed. + * \note For more information on this algorithm and the optimizations used here, + * see: https://en.wikipedia.org/wiki/Computation_of_cyclic_redundancy_checks. + * \param data Pointer to the data to compute the check for, as a byte-array. + * \param len Number of bytes in the data buffer. + */ + uint32_t crc(const uint8_t* data, size_t len) + { + // Precompute values we know we'll use frequently and store them in a + // lookup table. + static uint32_t table[256]; + static bool table_exists = false; + if (!table_exists) + { + make_table(table); + table_exists = true; + } + + // Start with a negated version of (unsigned) 0 to handle the + // case where data contains zeros. + uint32_t crc = 0xffffffff; + + for (size_t i = 0; i < len; i++) + { + uint32_t index = (crc ^ data[i]) & 0xff; + crc = (crc >> 8) ^ table[index]; + } + + // Under the negation we performed at the start. + crc = crc ^ 0xffffffff; + + return crc; + } + + /*! \brief Uses a cyclic redundancy check (CRC) to hash the contents of the buffer passed in. + * \note The result does not include type or length information, which is necessary to uniquely + * identify a buffer (eg. two vectors of zeros that are of different lengths will both + * produce the same CRC hash; two vectors of zeros that are the same length but different integral + * types may also produce the same CRC hash.) + * + * \tparam T Buffer element type + * \param buffer Pointer to the buffer to hash + * \param size Number of elements in the range to hash + * \return String hash value as indicated above + */ + template + std::string hash_buffer_crc(T* buffer, const size_t size) + { + const uint8_t* bytes = reinterpret_cast(buffer); + const size_t num_bytes = size * sizeof(T); + const uint32_t hash_val = crc(bytes, num_bytes); + + std::stringstream sstream; + sstream << hash_val; + + return sstream.str(); + } + + /*! \brief Uses a cyclic redundancy check (CRC) to hash the contents of the buffer passed in. + * \note Includes length, but not type information. The length information is important because + * two zero-containing vectors of different lengths will both hash to zero. + * + * \param begin Iterator pointing to the begining of the range to hash + * \param end Iterator pointing to one-past-the-end of the range to hash + * \return String of the form "hash,length" + */ + template + std::string hash_vector(const T begin, const T end) + { + const size_t size = end - begin; + return hash_buffer_crc(thrust::raw_pointer_cast(&(*begin)), size) + "," + std::to_string(size); + } + + /*! \brief Builds a compound string hash value from existing string hashes. + * + * \param begin Iterator pointing to the beginning of the vector of string hashes to combine + * \param end Iterator pointing to the end of the vector of string hashes to combine + * \return String hash value of the form "(,)" + */ + std::string build_compound_token(const std::vector::const_iterator begin, const std::vector::const_iterator end) + { + std::string token = "("; + if (begin != end) + { + for (auto it = begin; it != end; it++) + { + token += *it; + if (it + 1 != end) + token += ","; + } + } + token += ")"; + + return token; + } + + /*! \brief Builds a string "token" from a function call's inputs. This token is unique to a particular function call. + * + * \param list Vector of strings where the first element is the function name, and the remainder are hashes representing inputs to the function call. + * \return String hash value of the form "fcn_name(,,...)" + */ + std::string build_input_token(const std::vector& list) + { + return *(list.begin()) + build_compound_token(list.begin() + 1, list.end()); + } + + /*! \brief Builds a string "token" from a function call's outputs. This token is unique to a particular return value. + * + * \param list Vector of string hashes representing outputs to the function call. + * \return String hash value of the form "(,,...)" + */ + std::string build_output_token(const std::vector& list) + { + return build_compound_token(list.begin(), list.end()); + } +} + +/*! \brief Builds and returns a string containing the rocThrust version number. + * + * \return String of the form "..". + */ +std::string get_rocthrust_version() +{ + static const std::string rocthrust_ver = std::to_string(ROCTHRUST_VERSION_MAJOR) + ver_sep + + std::to_string(ROCTHRUST_VERSION_MINOR) + ver_sep + + std::to_string(ROCTHRUST_VERSION_PATCH); + return rocthrust_ver; +} + +/*! \brief Builds and returns a string containing the GPU architecture. + * + * \return String of the form "gfx". + */ +std::string get_gpu_arch() +{ + hipDeviceProp_t device_prop; + if(hipGetDeviceProperties(&device_prop, 0) != hipSuccess) + throw std::runtime_error("hipGetDeviceProperties failure"); + + static const std::string gpu_arch(device_prop.gcnArchName); + + return gpu_arch; +} + +/*! \brief Builds and returns a string containing the ROCm version number. + * \note We intentionally omit the "HIP_PATCH_VERSION" here, since we don't + * anticipate results to change at that granularity, and storing results at + * that granularity would require frequent database updates. + * + * \return String of the form ".". + */ +std::string get_rocm_version() +{ + static const std::string runtime_ver = std::to_string(HIP_VERSION_MAJOR) + ver_sep + + std::to_string(HIP_VERSION_MINOR); + + return runtime_ver; +} + +/*! \brief Returns a "token" string the uniquely identifies a scalar value. + * + * \tparam T Type of the scalar + * \param val Scalar value to hash + * \return String of the form "scalar(value)" + */ +template +std::string get_scalar_token(const T& val) +{ + return "scalar<" + get_typename_str() + ">(" + std::to_string(val) + ")"; +} + +/*! \brief Returns a "token" string the uniquely identifies a vector. + * \note This version accepts an existing hash string and type string + * (obtained from hash_vector and get_typename_str, respectively). This + * allows callers to make multiple calls to this function without invoking + * hash_vector multiple times, since it may be expensive if the vector is + * large. + * + * \param vec_hash Hash string for the vector, obtained from hash_vector. + * \param data_type Data type name string, obtained from get_typename_str. + * \return String of the form "vector(hash)". + */ +std::string get_vector_token(const std::string& vec_hash, const std::string& data_type) +{ + return "vector<" + data_type + ">(" + vec_hash + ")"; +} + +/*! \brief Returns a "token" string the uniquely identifies a vector iterator. + * \note This version accepts an existing hash string and type string + * (obtained from hash_vector and get_typename_str, respectively). This + * allows callers to make multiple calls to this function without invoking + * hash_vector multiple times, since it may be expensive if the vector is + * large. + * + * \param vec_hash Hash string for the vector, obtained from hash_vector. + * \param data_type Data type name string, obtained from get_typename_str. + * \param offset The iterator's current offset from the beginning of the vector. + * \return String of the form "iter(vector(vec_hash),offset)". + */ +std::string get_iterator_token(const std::string& vec_hash, const std::string& data_type, const size_t offset) +{ + return "iter(" + get_vector_token(vec_hash, data_type) + "," + std::to_string(offset) + ")"; +} + +/*! \brief Returns a "token" string that uniquely identifies a functor (callable object). + * + * \tparam T The data type that the functor operates on. + * \param functor_type String representing the functor type, without the datatype (eg. "thrust::plus" for functor thrust::plus). + * \return Strin gof the form "function(function_type)" + */ +template +std::string get_functor_token(const std::string& functor_type) +{ + return "functor<" + get_typename_str() + ">(" + functor_type + ")"; +} + +/*! \brief This class helps create input and output tokens representing a function call. + * You can call build_input_token and build_output_token, passing them string and iterators, + * and the class will combine them together into a single token. + * Eg. + * // Suppose we want to record this call: + * std::vector input = {...}; + * thrust::inclusive_scan(policy, d_input.begin(), d_input.end(), d_output.begin(), thrust::plus); + * + * TokenHelper helper; + * token_helper.build_input_token( + * "thrust::inclusive_scan", + * d_input.begin(), + * d_input.end(), + * {bwr_utils::get_functor_token("thrust::plus")} + * ); + * + * token_helper.build_output_token(d_output.begin(), d_output.size()); + * + * // Can access the input/output tokens using: + * std::string input_token = helper.get_input_token(); + * std::string output_Token = helper.get_output_token(); + */ +class TokenHelper +{ +public: + TokenHelper() = default; + ~TokenHelper() = default; + + /*! \brief Returns a "token" string that uniquely identifies a function call's inputs. + * \note It's assumed that you'll pass in all the function call's inputs in the order they appear in the call from left to right. + * + * \tparam Iter The input vector iterator type + * \param fcn_name The name of the function being called + * \param input_begin Iterator to the beginning of the input data vector + * \param input_end Iterator to the end of the input data vector + * \param extra_inputs A vector of (ordered) extra input tokens obtained from the bwr_utils::get_*_token functions. + */ + template + void build_input_token(const std::string& fcn_name, const Iter input_begin, const Iter input_end, std::vector&& extra_inputs = {}) + { + save_input_token(fcn_name, input_begin, input_end, {}, std::forward&&>(extra_inputs)); + } + + /*! \brief Returns a "token" string that uniquely identifies a function call's inputs. This version accepts both values and keys. + * \note It's assumed that you'll pass in all the function call's inputs in the order they appear in the call from left to right. + * + * \tparam KeyIter The key input vector iterator type + * \tparam ValueIter The value input vector iterator type + * \param fcn_name The name of the function being called + * \param key_input_begin Iterator to the beginning of the key input data vector + * \param key_input_end Iterator to the end of the key input data vector + * \param key_input_begin Iterator to the beginning of the value input data vector + * \param extra_inputs A vector of (ordered) extra input tokens obtained from the bwr_utils::get_*_token functions. + */ + template + void build_input_token(const std::string& fcn_name, const KeyIter key_input_begin, const KeyIter key_input_end, const ValueIter value_input_begin, std::vector&& extra_inputs = {}) + { + save_input_token(fcn_name, key_input_begin, key_input_end, std::vector({value_input_begin}), std::forward&&>(extra_inputs)); + } + + /*! \brief Returns a "token" string that uniquely identifies a function call's output. + * + * \tparam Iter The output vector iterator type + * \param output_begin Iterator to the beginning of the output data vector + * \param size Number of elements in the output vector. + */ + template + void build_output_token(const Iter output_begin, const size_t size) + { + save_output_token(output_begin, {}, size); + } + + /*! \brief Returns a "token" string that uniquely identifies a function call's outputs. This version accepts both values and keys. + * + * \tparam KeyIter The key output vector iterator type + * \tparam ValueIter The value output vector iterator type + * \param key_output_begin Iterator to the beginning of the key output data vector + * \param value_output_begin Iterator to the beginning of the value output data vector + */ + template + void build_output_token(const KeyIter key_output_begin, const ValueIter value_output_begin, const size_t size) + { + save_output_token(key_output_begin, std::vector({value_output_begin}), size); + } + + /*! \brief Retrieves the input token that was generated by the last call to build_input_token(). + * + * \return The input token as described above, or the empty string if build_input_token hasn't been called yet. + */ + std::string get_input_token() const + { + return m_input_token; + } + + /*! \brief Retrieves the output token that was generated by the last call to build_output_token(). + * + * \return The output token as described above, or the empty string if build_output_token hasn't been called yet. + */ + std::string get_output_token() const + { + return m_output_token; + } + +private: + template + void save_input_token(std::string fcn_name, KeyIter key_input_begin, KeyIter key_input_end, std::vector value_begins, std::vector&& extra_inputs = {}) + { + using KeyDataType = typename std::iterator_traits::value_type; + using ValueDataType = typename std::iterator_traits::value_type; + const std::string key_data_type = get_typename_str(); + const size_t size = key_input_end - key_input_begin; + const std::string key_input_data_hash = bwr_utils::hash_vector(key_input_begin, key_input_end); + + // Build a vector of tokens to pass to bwr_utils::build_input_token. + // We will always have the key iterators: + std::vector subtokens = {fcn_name, + bwr_utils::get_iterator_token(key_input_data_hash, key_data_type, 0), + bwr_utils::get_iterator_token(key_input_data_hash, key_data_type, size) + }; + + // But we may or may not have a value iterator: + for (auto value_input_begin : value_begins) + { + const std::string value_data_type = get_typename_str(); + const std::string value_input_data_hash = bwr_utils::hash_vector(value_input_begin, value_input_begin + size); + subtokens.push_back( + bwr_utils::get_iterator_token(value_input_data_hash, value_data_type, 0) + ); + } + + subtokens.insert( + subtokens.end(), + std::make_move_iterator(extra_inputs.begin()), + std::make_move_iterator(extra_inputs.end()) + ); + + // Save the resulting compound token so that it can be retrieved later + m_input_token = bwr_utils::build_input_token(subtokens); + } + + template + void save_output_token(const KeyIter key_output_begin, std::vector value_begins, const size_t size) + { + using KeyDataType = typename std::iterator_traits::value_type; + using ValueDataType = typename std::iterator_traits::value_type; + + const std::string key_output_data_hash = bwr_utils::hash_vector(key_output_begin, key_output_begin + size); + const std::string key_data_type = bwr_utils::get_typename_str(); + + // Build a vector of tokens to pass to bwr_utils::build_output_token. + // We will always have the key iterators: + std::vector subtokens = { + bwr_utils::get_vector_token(key_output_data_hash, key_data_type), + }; + + // But we may or may not have a value iterator: + for (auto value_output_begin : value_begins) + { + const std::string value_data_type = bwr_utils::get_typename_str(); + const std::string value_output_data_hash = bwr_utils::hash_vector(value_output_begin, value_output_begin + size); + subtokens.push_back( + bwr_utils::get_vector_token(value_output_data_hash, value_data_type) + ); + } + + // Save the resulting compound token so that it can be retrieved later + m_output_token = bwr_utils::build_output_token(subtokens); + } + + std::string m_input_token; + std::string m_output_token; +}; + +} // end namespace bwr_utils + +#endif // BRW_UTILS_HPP \ No newline at end of file diff --git a/test/test_header.hpp b/test/test_header.hpp index 1d33f36e..144ca7a5 100644 --- a/test/test_header.hpp +++ b/test/test_header.hpp @@ -16,6 +16,7 @@ */ #include "../testing/unittest/random.h" +#include "bitwise_repro/bwr_db.hpp" #include @@ -115,6 +116,54 @@ inline int set_device_from_ctest() } } +// If enabled, set up the database for inter-run bitwise reproducibility testing. +// Inter-run testing is enabled through the following environment variables: +// ROCTHRUST_BWR_PATH - path to the database (or where it should be created) +// ROCTHRUST_BWR_GENERATE - if set to 1, info about any function calls not +// found in the database will be inserted. No errors will be reported in this mode. +namespace inter_run_bwr +{ + // Disable this testing by default. + bool enabled = false; + + // This code doesn't need to be visible outside this file. + namespace + { + const static std::string path_env = "ROCTHRUST_BWR_PATH"; + const static std::string generate_env = "ROCTHRUST_BWR_GENERATE"; + + // Check the environment variables to see if the database should be + // instantiated, and if so, what mode it should be in. + std::unique_ptr create_db() + { + // Get the path to the database from an environment variable. + const char* db_path = std::getenv(path_env.c_str()); + const char* db_mode = std::getenv(generate_env.c_str()); + if (db_path) + { + // Check if we are allowed to insert rows into the database if + // we encounter calls that aren't already recorded. + BitwiseReproDB::Mode mode = BitwiseReproDB::Mode::test_mode; + if (db_mode && std::stoi(db_mode) > 0) + mode = BitwiseReproDB::Mode::generate_mode; + + enabled = true; + return std::make_unique(db_path, mode); + } + else if (db_mode) + { + throw std::runtime_error("ROCTHRUST_BWR_GENERATE is defined, but no database path was given.\n" + "Please set ROCTHRUST_BWR_PATH to the database path."); + } + + return nullptr; + } + } + + // Create/open the run-to-run bitwise reproducibility database. + std::unique_ptr db = create_db(); +} + // Input type parameter template > struct Params diff --git a/test/test_reproducibility.cpp b/test/test_reproducibility.cpp index 57450fff..9ee62818 100644 --- a/test/test_reproducibility.cpp +++ b/test/test_reproducibility.cpp @@ -22,6 +22,7 @@ #include #include "test_header.hpp" +#include "bitwise_repro/bwr_utils.hpp" typedef ::testing::Types< Params, std::decay_t>, @@ -70,6 +71,12 @@ void assert_reproducible(const thrust::device_vector& d_a, const thrust::devi ASSERT_NO_FATAL_FAILURE(assert_bit_eq(h_a, h_b)); } +void check_bwr_match(const bwr_utils::TokenHelper& token_helper) +{ + if (inter_run_bwr::enabled && inter_run_bwr::db) + ASSERT_TRUE(inter_run_bwr::db->match(token_helper.get_input_token(), token_helper.get_output_token())); +} + TYPED_TEST(ReproducibilityTests, Scan) { using Vector = typename TestFixture::input_type; @@ -77,6 +84,8 @@ TYPED_TEST(ReproducibilityTests, Scan) using T = typename Vector::value_type; using ScanOp = eepy_scan_op>; + bwr_utils::TokenHelper token_helper; + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); Policy policy; @@ -100,17 +109,51 @@ TYPED_TEST(ReproducibilityTests, Scan) // inclusive thrust::inclusive_scan( policy, d_input.begin(), d_input.end(), d_output_0.begin(), ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::inclusive_scan", + d_input.begin(), + d_input.end(), + {bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::inclusive_scan( policy, d_input.begin(), d_input.end(), d_output_1.begin(), ScanOp(true)); + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); // exclusive thrust::exclusive_scan( policy, d_input.begin(), d_input.end(), d_output_0.begin(), T {42}, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::exclusive_scan", + d_input.begin(), + d_input.end(), + {bwr_utils::get_scalar_token(T {42}), bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::exclusive_scan( policy, d_input.begin(), d_input.end(), d_output_1.begin(), T {42}, ScanOp(true)); + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); } } @@ -123,6 +166,8 @@ TYPED_TEST(ReproducibilityTests, ScanByKey) using T = typename Vector::value_type; using ScanOp = eepy_scan_op>; + bwr_utils::TokenHelper token_helper; + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); Policy policy; @@ -167,6 +212,19 @@ TYPED_TEST(ReproducibilityTests, ScanByKey) d_output_0.begin(), thrust::equal_to {}, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::inclusive_scan_by_key", + d_keys.begin(), + d_keys.end(), + d_input.begin(), + {bwr_utils::get_functor_token("thrust::equal_to"), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::inclusive_scan_by_key(policy, d_keys.begin(), d_keys.end(), @@ -175,6 +233,12 @@ TYPED_TEST(ReproducibilityTests, ScanByKey) thrust::equal_to {}, ScanOp(true)); + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); // exclusive @@ -185,6 +249,19 @@ TYPED_TEST(ReproducibilityTests, ScanByKey) d_output_0.begin(), T {123}, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::exclusive_scan_by_key", + d_keys.begin(), + d_keys.end(), + d_input.begin(), + {bwr_utils::get_scalar_token(T {123}), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::exclusive_scan_by_key(policy, d_keys.begin(), d_keys.end(), @@ -193,6 +270,12 @@ TYPED_TEST(ReproducibilityTests, ScanByKey) T {123}, ScanOp(true)); + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); } } @@ -205,6 +288,8 @@ TYPED_TEST(ReproducibilityTests, ReduceByKey) using T = typename Vector::value_type; using ScanOp = eepy_scan_op>; + bwr_utils::TokenHelper token_helper; + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); Policy policy; @@ -250,6 +335,19 @@ TYPED_TEST(ReproducibilityTests, ReduceByKey) d_vals_output_0.begin(), thrust::equal_to {}, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::reduce_by_key", + d_keys.begin(), + d_keys.end(), + d_vals.begin(), + {bwr_utils::get_functor_token("thrust::equal_to"), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::reduce_by_key(policy, d_keys.begin(), d_keys.end(), @@ -259,6 +357,12 @@ TYPED_TEST(ReproducibilityTests, ReduceByKey) thrust::equal_to {}, ScanOp(true)); + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_keys_output_1.begin(), d_vals_output_1.begin(), d_keys.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_keys_output_0, d_keys_output_1); assert_reproducible(d_vals_output_0, d_vals_output_1); } @@ -272,6 +376,8 @@ TYPED_TEST(ReproducibilityTests, TransformScan) using T = typename Vector::value_type; using ScanOp = eepy_scan_op>; + bwr_utils::TokenHelper token_helper; + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); Policy policy; @@ -297,12 +403,31 @@ TYPED_TEST(ReproducibilityTests, TransformScan) d_output_0.begin(), thrust::negate(), ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::transform_inclusive_scan", + d_input.begin(), + d_input.end(), + {bwr_utils::get_functor_token("thrust::negate"), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::transform_inclusive_scan(policy, d_input.begin(), d_input.end(), d_output_1.begin(), thrust::negate(), ScanOp(true)); + + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); thrust::transform_exclusive_scan(policy, @@ -312,6 +437,19 @@ TYPED_TEST(ReproducibilityTests, TransformScan) thrust::negate(), (T)11, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::transform_exclusive_scan", + d_input.begin(), + d_input.end(), + {bwr_utils::get_functor_token("thrust::negate"), + bwr_utils::get_scalar_token((T) 11), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::transform_exclusive_scan(policy, d_input.begin(), d_input.end(), @@ -319,6 +457,13 @@ TYPED_TEST(ReproducibilityTests, TransformScan) thrust::negate(), (T)11, ScanOp(true)); + + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); // in-place scans @@ -330,12 +475,31 @@ TYPED_TEST(ReproducibilityTests, TransformScan) d_output_0.begin(), thrust::negate(), ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::transform_inclusive_scan", + d_output_1.begin(), + d_output_1.end(), + {bwr_utils::get_functor_token("thrust::negate"), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::transform_inclusive_scan(policy, d_output_1.begin(), d_output_1.end(), d_output_1.begin(), thrust::negate(), ScanOp(true)); + + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); d_output_0 = d_input; @@ -347,6 +511,19 @@ TYPED_TEST(ReproducibilityTests, TransformScan) thrust::negate(), (T)11, ScanOp(false)); + + if (inter_run_bwr::enabled) + { + token_helper.build_input_token( + "thrust::transform_exclusive_scan", + d_output_1.begin(), + d_output_1.end(), + {bwr_utils::get_functor_token("thrust::negate"), + bwr_utils::get_scalar_token((T) 11), + bwr_utils::get_functor_token("thrust::plus")} + ); + } + thrust::transform_exclusive_scan(policy, d_output_1.begin(), d_output_1.end(), @@ -354,6 +531,13 @@ TYPED_TEST(ReproducibilityTests, TransformScan) thrust::negate(), (T)11, ScanOp(true)); + + if (inter_run_bwr::enabled) + { + token_helper.build_output_token(d_output_1.begin(), d_output_1.size()); + check_bwr_match(token_helper); + } + assert_reproducible(d_output_0, d_output_1); } }