From 007b76fc4a42e5b0ce5b81acb6339916ab5532fe Mon Sep 17 00:00:00 2001 From: Nico <31079890+ntrost57@users.noreply.github.com> Date: Thu, 25 May 2023 06:59:58 +0200 Subject: [PATCH 1/2] Change rocsparse test/benchmark matrix search path (#494) (#306) Co-authored-by: Yvan Mokwinski --- clients/benchmarks/CMakeLists.txt | 3 +- .../benchmarks/rocsparse_bench_cmdlines.hpp | 16 +- .../common/rocsparse_clients_envariables.cpp | 275 ++++++++++++++++++ .../common/rocsparse_importer_rocalution.cpp | 3 +- .../common/rocsparse_importer_rocsparseio.cpp | 4 +- clients/common/rocsparse_matrix_factory.cpp | 43 ++- clients/common/rocsparse_parse_data.cpp | 18 +- .../include/rocsparse_clients_envariables.hpp | 99 +++++++ clients/tests/CMakeLists.txt | 1 + 9 files changed, 443 insertions(+), 19 deletions(-) create mode 100644 clients/common/rocsparse_clients_envariables.cpp create mode 100644 clients/include/rocsparse_clients_envariables.hpp diff --git a/clients/benchmarks/CMakeLists.txt b/clients/benchmarks/CMakeLists.txt index 0e70c719..1662a323 100644 --- a/clients/benchmarks/CMakeLists.txt +++ b/clients/benchmarks/CMakeLists.txt @@ -1,5 +1,5 @@ # ######################################################################## -# Copyright (C) 2018-2022 Advanced Micro Devices, Inc. All rights Reserved. +# Copyright (C) 2018-2023 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 @@ -59,6 +59,7 @@ set(ROCSPARSE_CLIENTS_COMMON ../common/rocsparse_importer_rocalution.cpp ../common/rocsparse_importer_rocsparseio.cpp ../common/rocsparse_importer_matrixmarket.cpp + ../common/rocsparse_clients_envariables.cpp ) diff --git a/clients/benchmarks/rocsparse_bench_cmdlines.hpp b/clients/benchmarks/rocsparse_bench_cmdlines.hpp index 6e555ec9..7702bb01 100644 --- a/clients/benchmarks/rocsparse_bench_cmdlines.hpp +++ b/clients/benchmarks/rocsparse_bench_cmdlines.hpp @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ -* Copyright (C) 2021-2022 Advanced Micro Devices, Inc. All rights Reserved. +* Copyright (C) 2021-2023 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 @@ -23,6 +23,7 @@ * ************************************************************************ */ #pragma once +#include "rocsparse_clients_envariables.hpp" #include #include #include @@ -573,6 +574,19 @@ class rocsparse_bench_cmdlines public: static void help(std::ostream& out) { + out << "" << std::endl; + out << "Specific environment variables:" << std::endl; + for(const auto v : rocsparse_clients_envariables::s_var_bool_all) + { + out << rocsparse_clients_envariables::get_name(v) << " " + << rocsparse_clients_envariables::get_description(v) << std::endl; + } + for(const auto v : rocsparse_clients_envariables::s_var_string_all) + { + out << rocsparse_clients_envariables::get_name(v) << " " + << rocsparse_clients_envariables::get_description(v) << std::endl; + } + out << "" << std::endl; out << "Benchmarks options:" << std::endl; out << "--bench-x flag to preceed the main option " << std::endl; out << "--bench-o output JSON file, (default = a.json)" << std::endl; diff --git a/clients/common/rocsparse_clients_envariables.cpp b/clients/common/rocsparse_clients_envariables.cpp new file mode 100644 index 00000000..e36b323e --- /dev/null +++ b/clients/common/rocsparse_clients_envariables.cpp @@ -0,0 +1,275 @@ +/* ************************************************************************ + * Copyright (C) 2023 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. + * + * ************************************************************************ */ + +#include "rocsparse_clients_envariables.hpp" +#include "rocsparse-types.h" +#include + +constexpr rocsparse_clients_envariables::var_bool rocsparse_clients_envariables::s_var_bool_all[]; +constexpr rocsparse_clients_envariables::var_string + rocsparse_clients_envariables::s_var_string_all[]; + +template +static inline constexpr std::size_t countof(T (&)[N]) +{ + return N; +} + +static constexpr size_t s_var_bool_size = countof(rocsparse_clients_envariables::s_var_bool_all); +static constexpr size_t s_var_string_size + = countof(rocsparse_clients_envariables::s_var_string_all); + +static constexpr const char* s_var_bool_names[s_var_bool_size] = {"ROCSPARSE_CLIENTS_VERBOSE"}; +static constexpr const char* s_var_string_names[s_var_string_size] + = {"ROCSPARSE_CLIENTS_MATRICES_DIR"}; +static constexpr const char* s_var_bool_descriptions[s_var_bool_size] = {"0: disabled, 1: enabled"}; +static constexpr const char* s_var_string_descriptions[s_var_string_size] + = {"Full path of the matrices directory"}; + +/// +/// @brief Grab an environment variable value. +/// @return true if the operation is successful, false otherwise. +/// +template +static bool rocsparse_getenv(const char* name, bool& defined, T& val); + +template <> +bool rocsparse_getenv(const char* name, bool& defined, bool& val) +{ + val = false; + const char* getenv_str = getenv(name); + defined = (getenv_str != nullptr); + if(defined) + { + auto getenv_int = atoi(getenv_str); + if((getenv_int != 0) && (getenv_int != 1)) + { + std::cerr << "rocsparse error, invalid environment variable " << name + << " must be 0 or 1." << std::endl; + val = false; + return false; + } + else + { + val = (getenv_int == 1); + return true; + } + } + else + { + return true; + } +} + +template <> +bool rocsparse_getenv(const char* name, bool& defined, std::string& val) +{ + const char* getenv_str = getenv(name); + defined = (getenv_str != nullptr); + if(defined) + { + val = getenv_str; + } + return true; +} + +struct rocsparse_clients_envariables_impl +{ + +public: + // + // \brief Return value of a Boolean variable. + // + inline bool get(rocsparse_clients_envariables::var_bool v) const + { + return this->m_var_bool[v]; + }; + + // + // \brief Return value of a string variable. + // + inline const char* get(rocsparse_clients_envariables::var_string v) const + { + return this->m_var_string[v].c_str(); + }; + + // + // \brief Is a Boolean variable defined ? + // + inline bool is_defined(rocsparse_clients_envariables::var_bool v) const + { + return this->m_var_bool_defined[v]; + }; + + // + // \brief Is a string variable defined ? + // + inline bool is_defined(rocsparse_clients_envariables::var_string v) const + { + return this->m_var_string_defined[v]; + }; + + // + // Return the unique instance. + // + static rocsparse_clients_envariables_impl& Instance(); + +private: + ~rocsparse_clients_envariables_impl() = default; + rocsparse_clients_envariables_impl(const rocsparse_clients_envariables_impl&) = delete; + rocsparse_clients_envariables_impl& operator=(const rocsparse_clients_envariables_impl&) + = delete; + + bool m_var_bool[s_var_bool_size]{}; + bool m_var_bool_defined[s_var_bool_size]{}; + + std::string m_var_string[s_var_string_size]{}; + bool m_var_string_defined[s_var_string_size]{}; + + rocsparse_clients_envariables_impl() + { + for(auto tag : rocsparse_clients_envariables::s_var_bool_all) + { + switch(tag) + { + case rocsparse_clients_envariables::VERBOSE: + { + const bool success = rocsparse_getenv( + s_var_bool_names[tag], this->m_var_bool_defined[tag], this->m_var_bool[tag]); + if(!success) + { + std::cerr << "rocsparse_getenv failed on fetching " << s_var_bool_names[tag] + << std::endl; + throw(rocsparse_status_invalid_value); + } + break; + } + } + } + + for(auto tag : rocsparse_clients_envariables::s_var_string_all) + { + switch(tag) + { + case rocsparse_clients_envariables::MATRICES_DIR: + { + const bool success = rocsparse_getenv(s_var_string_names[tag], + this->m_var_string_defined[tag], + this->m_var_string[tag]); + if(!success) + { + std::cerr << "rocsparse_getenv failed on fetching " << s_var_string_names[tag] + << std::endl; + throw(rocsparse_status_invalid_value); + } + break; + } + } + } + + if(this->m_var_bool[rocsparse_clients_envariables::VERBOSE]) + { + for(auto tag : rocsparse_clients_envariables::s_var_bool_all) + { + switch(tag) + { + case rocsparse_clients_envariables::VERBOSE: + { + const bool v = this->m_var_bool[tag]; + std::cout << "" + << "env variable " << s_var_bool_names[tag] << " : " + << ((this->m_var_bool_defined[tag]) ? ((v) ? "enabled" : "disabled") + : "") + << std::endl; + break; + } + } + } + + for(auto tag : rocsparse_clients_envariables::s_var_string_all) + { + switch(tag) + { + case rocsparse_clients_envariables::MATRICES_DIR: + { + const std::string v = this->m_var_string[tag]; + std::cout << "" + << "env variable " << s_var_string_names[tag] << " : " + << ((this->m_var_string_defined[tag]) ? this->m_var_string[tag] + : "") + << std::endl; + break; + } + } + } + } + } +}; + +rocsparse_clients_envariables_impl& rocsparse_clients_envariables_impl::Instance() +{ + static rocsparse_clients_envariables_impl instance; + return instance; +} + +bool rocsparse_clients_envariables::is_defined(rocsparse_clients_envariables::var_string v) +{ + return rocsparse_clients_envariables_impl::Instance().is_defined(v); +} + +const char* rocsparse_clients_envariables::get(rocsparse_clients_envariables::var_string v) +{ + return rocsparse_clients_envariables_impl::Instance().get(v); +} + +const char* rocsparse_clients_envariables::get_name(rocsparse_clients_envariables::var_string v) +{ + return s_var_string_names[v]; +} + +const char* + rocsparse_clients_envariables::get_description(rocsparse_clients_envariables::var_string v) +{ + return s_var_string_descriptions[v]; +} + +bool rocsparse_clients_envariables::is_defined(rocsparse_clients_envariables::var_bool v) +{ + return rocsparse_clients_envariables_impl::Instance().is_defined(v); +} + +bool rocsparse_clients_envariables::get(rocsparse_clients_envariables::var_bool v) +{ + return rocsparse_clients_envariables_impl::Instance().get(v); +} + +const char* rocsparse_clients_envariables::get_name(rocsparse_clients_envariables::var_bool v) +{ + return s_var_bool_names[v]; +} + +const char* + rocsparse_clients_envariables::get_description(rocsparse_clients_envariables::var_bool v) +{ + return s_var_bool_descriptions[v]; +} diff --git a/clients/common/rocsparse_importer_rocalution.cpp b/clients/common/rocsparse_importer_rocalution.cpp index ae65860a..63c28823 100644 --- a/clients/common/rocsparse_importer_rocalution.cpp +++ b/clients/common/rocsparse_importer_rocalution.cpp @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ - * Copyright (C) 2021-2022 Advanced Micro Devices, Inc. All rights Reserved. + * Copyright (C) 2021-2023 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 @@ -140,6 +140,7 @@ rocsparse_status rocsparse_importer_rocalution::import_sparse_csx( this->m_info_csx.in = new std::ifstream(this->m_filename, std::ios::in | std::ios::binary); if(!this->m_info_csx.in->is_open()) { + std::cerr << "cannot open file '" << this->m_filename << "'" << std::endl; return rocsparse_status_internal_error; } std::string header; diff --git a/clients/common/rocsparse_importer_rocsparseio.cpp b/clients/common/rocsparse_importer_rocsparseio.cpp index 8ddbec4b..24136d8e 100644 --- a/clients/common/rocsparse_importer_rocsparseio.cpp +++ b/clients/common/rocsparse_importer_rocsparseio.cpp @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ - * Copyright (C) 2021-2022 Advanced Micro Devices, Inc. All rights Reserved. + * Copyright (C) 2021-2023 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 @@ -160,7 +160,7 @@ rocsparse_importer_rocsparseio::rocsparse_importer_rocsparseio(const std::string istatus = rocsparseio_open(&this->m_handle, rocsparseio_rwmode_read, this->m_filename.c_str()); if(istatus != rocsparseio_status_success) { - std::cerr << "Problem with closing rocsparseio_open" << std::endl; + std::cerr << "cannot open file '" << this->m_filename << "'" << std::endl; throw rocsparse_status_internal_error; } #else diff --git a/clients/common/rocsparse_matrix_factory.cpp b/clients/common/rocsparse_matrix_factory.cpp index 0b3412f5..c6af72a1 100644 --- a/clients/common/rocsparse_matrix_factory.cpp +++ b/clients/common/rocsparse_matrix_factory.cpp @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ - * Copyright (C) 2021-2022 Advanced Micro Devices, Inc. All rights Reserved. + * Copyright (C) 2021-2023 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 @@ -23,8 +23,27 @@ * ************************************************************************ */ #include "rocsparse_matrix_factory.hpp" +#include "rocsparse_clients_envariables.hpp" #include "rocsparse_init.hpp" +static void get_matrix_full_filename(std::string& full_filename_, + const std::string& filename_, + const std::string& extension_, + const bool timing_) +{ + const bool envar_is_defined + = rocsparse_clients_envariables::is_defined(rocsparse_clients_envariables::MATRICES_DIR); + std::string path = rocsparse_exepath() + "../matrices/"; + if(envar_is_defined) + { + path = rocsparse_clients_envariables::get(rocsparse_clients_envariables::MATRICES_DIR); + path += "/"; + } + + full_filename_ = timing_ ? ((envar_is_defined) ? (path + filename_) : filename_) + : (path + filename_ + extension_); +} + // // Destructor. // @@ -96,30 +115,28 @@ rocsparse_matrix_factory::rocsparse_matrix_factory(const Arguments& case rocsparse_matrix_file_rocalution: { - std::string filename = arg.timing - ? arg.filename - : rocsparse_exepath() + "../matrices/" + arg.filename + ".csr"; + std::string full_filename; + get_matrix_full_filename(full_filename, arg.filename, ".csr", arg.timing); + this->m_instance - = new rocsparse_matrix_factory_rocalution(filename.c_str(), to_int); + = new rocsparse_matrix_factory_rocalution(full_filename.c_str(), to_int); break; } case rocsparse_matrix_file_rocsparseio: { - std::string filename = arg.timing - ? arg.filename - : rocsparse_exepath() + "../matrices/" + arg.filename + ".bin"; + std::string full_filename; + get_matrix_full_filename(full_filename, arg.filename, ".bin", arg.timing); this->m_instance - = new rocsparse_matrix_factory_rocsparseio(filename.c_str(), to_int); + = new rocsparse_matrix_factory_rocsparseio(full_filename.c_str(), to_int); break; } case rocsparse_matrix_file_mtx: { - std::string filename = arg.timing - ? arg.filename - : rocsparse_exepath() + "../matrices/" + arg.filename + ".mtx"; - this->m_instance = new rocsparse_matrix_factory_mtx(filename.c_str()); + std::string full_filename; + get_matrix_full_filename(full_filename, arg.filename, ".mtx", arg.timing); + this->m_instance = new rocsparse_matrix_factory_mtx(full_filename.c_str()); break; } diff --git a/clients/common/rocsparse_parse_data.cpp b/clients/common/rocsparse_parse_data.cpp index 686b3ddc..ee70caca 100644 --- a/clients/common/rocsparse_parse_data.cpp +++ b/clients/common/rocsparse_parse_data.cpp @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ - * Copyright (C) 2019-2022 Advanced Micro Devices, Inc. All rights Reserved. + * Copyright (C) 2019-2023 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 @@ -44,6 +44,8 @@ namespace fs = std::experimental::filesystem; #include #endif +#include "rocsparse_clients_envariables.hpp" + // Parse YAML data static std::string rocsparse_parse_yaml(const std::string& yaml) { @@ -142,6 +144,20 @@ bool rocsparse_parse_data(int& argc, char** argv, const std::string& default_fil std::cout << "\n" << argv[0] << " [ --data | --yaml ] ...\n" << std::endl; + + std::cout << "" << std::endl; + std::cout << "Specific environment variables:" << std::endl; + for(const auto v : rocsparse_clients_envariables::s_var_bool_all) + { + std::cout << rocsparse_clients_envariables::get_name(v) << " " + << rocsparse_clients_envariables::get_description(v) << std::endl; + } + for(const auto v : rocsparse_clients_envariables::s_var_string_all) + { + std::cout << rocsparse_clients_envariables::get_name(v) << " " + << rocsparse_clients_envariables::get_description(v) << std::endl; + } + std::cout << "" << std::endl; } } } diff --git a/clients/include/rocsparse_clients_envariables.hpp b/clients/include/rocsparse_clients_envariables.hpp new file mode 100644 index 00000000..a28483a0 --- /dev/null +++ b/clients/include/rocsparse_clients_envariables.hpp @@ -0,0 +1,99 @@ +/*! \file */ +/* ************************************************************************ + * Copyright (C) 2023 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. + * + * ************************************************************************ */ + +#pragma once + +/// +/// @brief Definition of a utility struct to grab environment variables +/// for the rocsparse clients. +// +/// The corresponding environment variable is the literal enum string +/// with the prefix ROCSPARSE_CLIENTS_. +/// Example: rocsparse_clients_envariables::VERBOSE will have a one to one correspondance with the environment variable +/// ROCSPARSE_CLIENTS_VERBOSE. +/// Obviously it loads environment variables at run time. +/// +struct rocsparse_clients_envariables +{ + + /// + /// @brief Enumerate Boolean environment variables. + /// + typedef enum var_bool_ : int32_t + { + VERBOSE + } var_bool; + + static constexpr var_bool s_var_bool_all[] = {VERBOSE}; + + /// + /// @brief Return value of a Boolean variable. + /// + static bool get(var_bool v); + + /// + /// @brief Is the Boolean enviromnent variable defined ? + /// + static bool is_defined(var_bool v); + + /// + /// @brief Return the name of a Boolean variable. + /// + static const char* get_name(var_bool v); + + /// + /// @brief Return the description of a Boolean variable. + /// + static const char* get_description(var_bool v); + + /// + /// @brief Enumerate string environment variables. + /// + typedef enum var_string_ : int32_t + { + MATRICES_DIR + } var_string; + + static constexpr var_string s_var_string_all[] = {MATRICES_DIR}; + + /// + /// @brief Return value of a string variable. + /// + static const char* get(var_string v); + + /// + /// @brief Return the name of a string variable. + /// + static const char* get_name(var_string v); + + /// + /// @brief Return the description of a string variable. + /// + static const char* get_description(var_string v); + + /// + /// @brief Is the string enviromnent variable defined ? + /// + static bool is_defined(var_string v); +}; diff --git a/clients/tests/CMakeLists.txt b/clients/tests/CMakeLists.txt index 966b0c32..52d5d60c 100644 --- a/clients/tests/CMakeLists.txt +++ b/clients/tests/CMakeLists.txt @@ -329,6 +329,7 @@ set(ROCSPARSE_CLIENTS_COMMON ../common/rocsparse_importer_rocalution.cpp ../common/rocsparse_importer_rocsparseio.cpp ../common/rocsparse_importer_matrixmarket.cpp + ../common/rocsparse_clients_envariables.cpp ) add_executable(rocsparse-test rocsparse_test_main.cpp ${ROCSPARSE_TEST_SOURCES} ${ROCSPARSE_CLIENTS_COMMON} ${ROCSPARSE_CLIENTS_TESTINGS}) From d5016c91f00b2cd48334d719589a86c813ad7f1d Mon Sep 17 00:00:00 2001 From: James Sandham <33790278+jsandham@users.noreply.github.com> Date: Tue, 30 May 2023 10:39:22 -0700 Subject: [PATCH 2/2] Fix xnack csritilu0 sync_split_fusion failures (#552) (#317) * Fix xnack csritilu0 sync_split_fusion failures (#552) * fix xnack failures * 2023 * formatting * fix clang formatting errors * clang formatting * clang formatting * clang formatting --------- Co-authored-by: Angelo Gonzales <48339672+Angeloo01@users.noreply.github.com> Co-authored-by: jsandham --- clients/common/rocsparse_host.cpp | 73 +++++++++---------- clients/testings/testing_bsr2csr.cpp | 1 - clients/testings/testing_prune_csr2csr.cpp | 2 +- .../testing_prune_csr2csr_by_percentage.cpp | 1 - library/include/rocsparse-types.h | 2 +- .../src/conversion/rocsparse_coo2dense.cpp | 1 - library/src/include/common.h | 2 +- .../rocsparse_csritilu0_buffer_size.cpp | 3 +- .../itilu0/rocsparse_csritilu0_compute.cpp | 8 +- .../itilu0/rocsparse_csritilu0_history.cpp | 5 +- .../itilu0/rocsparse_csritilu0_preprocess.cpp | 5 +- .../rocsparse_csritilu0x_sync_fusion.cpp | 2 + 12 files changed, 48 insertions(+), 57 deletions(-) diff --git a/clients/common/rocsparse_host.cpp b/clients/common/rocsparse_host.cpp index 4a7dc294..6f37f4b9 100644 --- a/clients/common/rocsparse_host.cpp +++ b/clients/common/rocsparse_host.cpp @@ -7410,7 +7410,7 @@ void host_gebsr_to_csr(rocsparse_direction direction, { for(rocsparse_int c = 0; c < col_block_dim; ++c) { - rocsparse_int col = col_block_dim * j + c; + rocsparse_int col = col_block_dim * j + c; rocsparse_int index = start * row_block_dim * col_block_dim + (end - start) * col_block_dim * r + (k - start) * col_block_dim + c; @@ -9957,42 +9957,41 @@ template void host_coosort_by_column(rocsparse_int M, std::vector& coo_row_ind, std::vector& coo_col_ind, std::vector& coo_val); - -#define INSTANTIATE1(TYPE) \ - template void host_bsr_to_csr(rocsparse_direction direction, \ - rocsparse_int mb, \ - rocsparse_int nb, \ - rocsparse_int nnzb, \ - const std::vector& bsr_val, \ - const std::vector& bsr_row_ptr, \ - const std::vector& bsr_col_ind, \ - rocsparse_int block_dim, \ - rocsparse_index_base bsr_base, \ - std::vector& csr_val, \ - std::vector& csr_row_ptr, \ - std::vector& csr_col_ind, \ - rocsparse_index_base csr_base); \ - template void host_csr_to_bsr(rocsparse_direction direction, \ - rocsparse_int m, \ - rocsparse_int n, \ - rocsparse_int nnz, \ - const std::vector& csr_val, \ - const std::vector& csr_row_ptr, \ - const std::vector& csr_col_ind, \ - rocsparse_int block_dim, \ - rocsparse_index_base csr_base, \ - std::vector& bsr_val, \ - std::vector& bsr_row_ptr, \ - std::vector& bsr_col_ind, \ - rocsparse_index_base bsr_base); \ - template void host_bsrpad_value(rocsparse_int m, \ - rocsparse_int mb, \ - rocsparse_int nnzb, \ - rocsparse_int block_dim, \ - TYPE value, \ - TYPE * bsr_val, \ - const rocsparse_int* bsr_row_ptr, \ - const rocsparse_int* bsr_col_ind, \ +#define INSTANTIATE1(TYPE) \ + template void host_bsr_to_csr(rocsparse_direction direction, \ + rocsparse_int mb, \ + rocsparse_int nb, \ + rocsparse_int nnzb, \ + const std::vector& bsr_val, \ + const std::vector& bsr_row_ptr, \ + const std::vector& bsr_col_ind, \ + rocsparse_int block_dim, \ + rocsparse_index_base bsr_base, \ + std::vector& csr_val, \ + std::vector& csr_row_ptr, \ + std::vector& csr_col_ind, \ + rocsparse_index_base csr_base); \ + template void host_csr_to_bsr(rocsparse_direction direction, \ + rocsparse_int m, \ + rocsparse_int n, \ + rocsparse_int nnz, \ + const std::vector& csr_val, \ + const std::vector& csr_row_ptr, \ + const std::vector& csr_col_ind, \ + rocsparse_int block_dim, \ + rocsparse_index_base csr_base, \ + std::vector& bsr_val, \ + std::vector& bsr_row_ptr, \ + std::vector& bsr_col_ind, \ + rocsparse_index_base bsr_base); \ + template void host_bsrpad_value(rocsparse_int m, \ + rocsparse_int mb, \ + rocsparse_int nnzb, \ + rocsparse_int block_dim, \ + TYPE value, \ + TYPE * bsr_val, \ + const rocsparse_int* bsr_row_ptr, \ + const rocsparse_int* bsr_col_ind, \ rocsparse_index_base bsr_base); #define INSTANTIATE2(ITYPE, TTYPE) \ diff --git a/clients/testings/testing_bsr2csr.cpp b/clients/testings/testing_bsr2csr.cpp index e0cc926f..ea579ed2 100644 --- a/clients/testings/testing_bsr2csr.cpp +++ b/clients/testings/testing_bsr2csr.cpp @@ -68,7 +68,6 @@ void testing_bsr2csr_bad_arg(const Arguments& arg) // Check block_dim == 0 block_dim = 0; EXPECT_ROCSPARSE_STATUS(rocsparse_bsr2csr(PARAMS), rocsparse_status_invalid_size); - #undef PARAMS } diff --git a/clients/testings/testing_prune_csr2csr.cpp b/clients/testings/testing_prune_csr2csr.cpp index 24b821bf..e3c0874b 100644 --- a/clients/testings/testing_prune_csr2csr.cpp +++ b/clients/testings/testing_prune_csr2csr.cpp @@ -128,7 +128,7 @@ void testing_prune_csr2csr(const Arguments& arg) device_vector d_csr_row_ptr_A(M + 1); device_vector d_csr_col_ind_A(nnz_A); device_vector d_csr_val_A(nnz_A); - device_scalar d_threshold(h_threshold); + device_scalar d_threshold(h_threshold); // Copy data from CPU to device CHECK_HIP_ERROR(hipMemcpy( diff --git a/clients/testings/testing_prune_csr2csr_by_percentage.cpp b/clients/testings/testing_prune_csr2csr_by_percentage.cpp index e92619ee..a2322cd2 100644 --- a/clients/testings/testing_prune_csr2csr_by_percentage.cpp +++ b/clients/testings/testing_prune_csr2csr_by_percentage.cpp @@ -94,7 +94,6 @@ void testing_prune_csr2csr_by_percentage_bad_arg(const Arguments& arg) rocsparse_status_not_implemented); EXPECT_ROCSPARSE_STATUS(rocsparse_prune_csr2csr_by_percentage(PARAMS), rocsparse_status_not_implemented); - CHECK_ROCSPARSE_ERROR( rocsparse_set_mat_storage_mode(csr_descr_A, rocsparse_storage_mode_sorted)); CHECK_ROCSPARSE_ERROR( diff --git a/library/include/rocsparse-types.h b/library/include/rocsparse-types.h index 41eaa287..4d1ec6c7 100644 --- a/library/include/rocsparse-types.h +++ b/library/include/rocsparse-types.h @@ -489,7 +489,7 @@ typedef enum rocsparse_itilu0_alg_ rocsparse_itilu0_alg_sync_split = 3, /**< Synchronous ITILU0 algorithm with explicit storage splitting */ rocsparse_itilu0_alg_sync_split_fusion - = 4 /**< Semi-synchronous ITILU0 algorithm with explicit storage splitting, this algorithm is having accuracy issues and is now falling back on \ref rocsparse_itilu0_alg_sync_split */ + = 4 /**< Semi-synchronous ITILU0 algorithm with explicit storage splitting */ } rocsparse_itilu0_alg; /*! \ingroup types_module diff --git a/library/src/conversion/rocsparse_coo2dense.cpp b/library/src/conversion/rocsparse_coo2dense.cpp index fff9b945..bcd6a741 100644 --- a/library/src/conversion/rocsparse_coo2dense.cpp +++ b/library/src/conversion/rocsparse_coo2dense.cpp @@ -26,7 +26,6 @@ #include "rocsparse_coo2dense.hpp" - #include "common.h" #include "coo2dense_device.h" diff --git a/library/src/include/common.h b/library/src/include/common.h index 7573f148..c58ae3ae 100644 --- a/library/src/include/common.h +++ b/library/src/include/common.h @@ -961,7 +961,7 @@ __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL template __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL -void memset2d_kernel(I m, I n, T value, T* __restrict__ data, I ld, rocsparse_order order) + void memset2d_kernel(I m, I n, T value, T* __restrict__ data, I ld, rocsparse_order order) { I gid = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x; diff --git a/library/src/precond/itilu0/rocsparse_csritilu0_buffer_size.cpp b/library/src/precond/itilu0/rocsparse_csritilu0_buffer_size.cpp index e3bc8c1e..38fdb3a9 100644 --- a/library/src/precond/itilu0/rocsparse_csritilu0_buffer_size.cpp +++ b/library/src/precond/itilu0/rocsparse_csritilu0_buffer_size.cpp @@ -49,9 +49,8 @@ static rocsparse_status buffer_size_dispatch(rocsparse_itilu0_alg alg_, P&&... p } case rocsparse_itilu0_alg_sync_split_fusion: { - // Fall back to the sync split algorithm. return rocsparse_csritilu0_driver_t< - rocsparse_itilu0_alg_sync_split>::buffer_size::run(parameters...); + rocsparse_itilu0_alg_sync_split_fusion>::buffer_size::run(parameters...); } } return rocsparse_status_invalid_value; diff --git a/library/src/precond/itilu0/rocsparse_csritilu0_compute.cpp b/library/src/precond/itilu0/rocsparse_csritilu0_compute.cpp index bf1c2c05..770260ef 100644 --- a/library/src/precond/itilu0/rocsparse_csritilu0_compute.cpp +++ b/library/src/precond/itilu0/rocsparse_csritilu0_compute.cpp @@ -48,12 +48,8 @@ static rocsparse_status compute_dispatch(rocsparse_itilu0_alg alg_, P&&... param } case rocsparse_itilu0_alg_sync_split_fusion: { - // Fall back to the sync split algorithm. - std::cerr << "// rocSPARSE.WARNING: algorithm sync_split_fusion is temporarily disabled " - "and replaced with sync_split" - << std::endl; - return rocsparse_csritilu0_driver_t::compute::run( - parameters...); + return rocsparse_csritilu0_driver_t< + rocsparse_itilu0_alg_sync_split_fusion>::compute::run(parameters...); } } diff --git a/library/src/precond/itilu0/rocsparse_csritilu0_history.cpp b/library/src/precond/itilu0/rocsparse_csritilu0_history.cpp index 0c4b7b8c..3c5b50fe 100644 --- a/library/src/precond/itilu0/rocsparse_csritilu0_history.cpp +++ b/library/src/precond/itilu0/rocsparse_csritilu0_history.cpp @@ -48,9 +48,8 @@ static rocsparse_status history_dispatch(rocsparse_itilu0_alg alg_, P&&... param } case rocsparse_itilu0_alg_sync_split_fusion: { - // Fall back to the sync split algorithm. - return rocsparse_csritilu0_driver_t< - rocsparse_itilu0_alg_sync_split>::history, J>::run(parameters...); + return rocsparse_csritilu0_driver_t:: + history, J>::run(parameters...); } } diff --git a/library/src/precond/itilu0/rocsparse_csritilu0_preprocess.cpp b/library/src/precond/itilu0/rocsparse_csritilu0_preprocess.cpp index fbff123d..61544972 100644 --- a/library/src/precond/itilu0/rocsparse_csritilu0_preprocess.cpp +++ b/library/src/precond/itilu0/rocsparse_csritilu0_preprocess.cpp @@ -49,9 +49,8 @@ static rocsparse_status preprocess_dispatch(rocsparse_itilu0_alg alg_, P&&... pa } case rocsparse_itilu0_alg_sync_split_fusion: { - // Fall back to the sync split algorithm. - return rocsparse_csritilu0_driver_t::preprocess::run( - parameters...); + return rocsparse_csritilu0_driver_t< + rocsparse_itilu0_alg_sync_split_fusion>::preprocess::run(parameters...); } } diff --git a/library/src/precond/itilu0/rocsparse_csritilu0x_sync_fusion.cpp b/library/src/precond/itilu0/rocsparse_csritilu0x_sync_fusion.cpp index 84373950..59b42985 100644 --- a/library/src/precond/itilu0/rocsparse_csritilu0x_sync_fusion.cpp +++ b/library/src/precond/itilu0/rocsparse_csritilu0x_sync_fusion.cpp @@ -1051,6 +1051,8 @@ struct rocsparse_csritilu0x_driver_t handle_->stream); } + RETURN_IF_HIP_ERROR(hipStreamSynchronize(handle_->stream)); + if(compute_nrm_residual && compute_nrm_corr) { nrm_indicator = std::max(nrm_residual, nrm_corr);