From 5406f2469706a2cce32d85213fd714c0e3be8f64 Mon Sep 17 00:00:00 2001 From: neoblizz Date: Thu, 21 Jul 2022 17:17:37 -0700 Subject: [PATCH 1/8] Scripts + quick and easy spmm. --- examples/CMakeLists.txt | 1 + examples/spmm/CMakeLists.txt | 16 ++++ examples/spmm/helpers.hxx | 80 ++++++++++++++++ examples/spmm/thread_mapped.cu | 44 +++++++++ .../loops/algorithms/spmm/thread_mapped.cuh | 94 +++++++++++++++++++ include/loops/container/matrix.cuh | 53 +++++++++++ scripts/spmv.py | 23 +++++ 7 files changed, 311 insertions(+) create mode 100644 examples/spmm/CMakeLists.txt create mode 100644 examples/spmm/helpers.hxx create mode 100644 examples/spmm/thread_mapped.cu create mode 100644 include/loops/algorithms/spmm/thread_mapped.cuh create mode 100644 include/loops/container/matrix.cuh create mode 100644 scripts/spmv.py diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 74a425f..c83ab69 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -2,4 +2,5 @@ add_subdirectory(range) add_subdirectory(saxpy) add_subdirectory(spmv) +add_subdirectory(spmm) # end /* Add examples' subdirectories */ \ No newline at end of file diff --git a/examples/spmm/CMakeLists.txt b/examples/spmm/CMakeLists.txt new file mode 100644 index 0000000..4607ffd --- /dev/null +++ b/examples/spmm/CMakeLists.txt @@ -0,0 +1,16 @@ +# begin /* Add application */ +set(SOURCES + thread_mapped.cu +) + +foreach(SOURCE IN LISTS SOURCES) + get_filename_component(TEST_NAME ${SOURCE} NAME_WLE) + add_executable(loops.spmm.${TEST_NAME} ${SOURCE}) + target_link_libraries(loops.spmm.${TEST_NAME} PRIVATE loops) + set_target_properties(loops.spmm.${TEST_NAME} + PROPERTIES + CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES} + ) + message(STATUS "Example Added: loops.spmm.${TEST_NAME}") +endforeach() +# end /* Add application */ \ No newline at end of file diff --git a/examples/spmm/helpers.hxx b/examples/spmm/helpers.hxx new file mode 100644 index 0000000..0eed062 --- /dev/null +++ b/examples/spmm/helpers.hxx @@ -0,0 +1,80 @@ +/** + * @file helpers.hxx + * @author Muhammad Osama (mosama@ucdavis.edu) + * @brief Header file for SpMM. + * @version 0.1 + * @date 2022-02-03 + * + * @copyright Copyright (c) 2022 + * + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +struct parameters_t { + std::string filename; + bool validate; + bool verbose; + cxxopts::Options options; + + /** + * @brief Construct a new parameters object and parse command line arguments. + * + * @param argc Number of command line arguments. + * @param argv Command line arguments. + */ + parameters_t(int argc, char** argv) + : options(argv[0], "Sparse Matrix-Matrix Multiplication") { + // Add command line options + options.add_options()("h,help", "Print help") // help + ("m,market", "Matrix file", cxxopts::value()) // mtx + ("validate", "CPU validation") // validate + ("v,verbose", "Verbose output"); // verbose + + // Parse command line arguments + auto result = options.parse(argc, argv); + + if (result.count("help") || (result.count("market") == 0)) { + std::cout << options.help({""}) << std::endl; + std::exit(0); + } + + if (result.count("market") == 1) { + filename = result["market"].as(); + if (loops::is_market(filename)) { + } else { + std::cout << options.help({""}) << std::endl; + std::exit(0); + } + } else { + std::cout << options.help({""}) << std::endl; + std::exit(0); + } + + if (result.count("validate") == 1) { + validate = true; + } else { + validate = false; + } + + if (result.count("verbose") == 1) { + verbose = true; + } else { + verbose = false; + } + } +}; diff --git a/examples/spmm/thread_mapped.cu b/examples/spmm/thread_mapped.cu new file mode 100644 index 0000000..5944a0e --- /dev/null +++ b/examples/spmm/thread_mapped.cu @@ -0,0 +1,44 @@ +/** + * @file thread_mapped.cu + * @author Muhammad Osama (mosama@ucdavis.edu) + * @brief Sparse Matrix-Matrix Multiplication example. + * @version 0.1 + * @date 2022-02-03 + * + * @copyright Copyright (c) 2022 + * + */ + +#include "helpers.hxx" +#include + +using namespace loops; + +int main(int argc, char** argv) { + using index_t = int; + using offset_t = int; + using type_t = float; + + // ... I/O parameters, mtx, etc. + parameters_t parameters(argc, argv); + + matrix_market_t mtx; + csr_t csr(mtx.load(parameters.filename)); + + // Input and output matrices. + std::size_t n = 10; + matrix_t B(csr.cols, n); + matrix_t C(csr.rows, n); + + // Generate random numbers between [0, 10]. + generate::random::uniform_distribution(B.m_data.begin(), B.m_data.end(), 1, + 10); + + // Run the benchmark. + util::timer_t timer; + timer.start(); + algorithms::spmm::thread_mapped(csr, B, C); + timer.stop(); + + std::cout << "Elapsed (ms):\t" << timer.milliseconds() << std::endl; +} \ No newline at end of file diff --git a/include/loops/algorithms/spmm/thread_mapped.cuh b/include/loops/algorithms/spmm/thread_mapped.cuh new file mode 100644 index 0000000..49b54eb --- /dev/null +++ b/include/loops/algorithms/spmm/thread_mapped.cuh @@ -0,0 +1,94 @@ +/** + * @file thread_mapped.cuh + * @author Muhammad Osama (mosama@ucdavis.edu) + * @brief Sparse Matrix-Matrix Multiplication kernels. + * @version 0.1 + * @date 2022-02-03 + * + * @copyright Copyright (c) 2022 + * + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace loops { +namespace algorithms { +namespace spmm { + +template +__global__ void __thread_mapped(setup_t config, + const std::size_t a_rows, + const std::size_t a_cols, + const std::size_t a_nnz, + const offset_t* offsets, + const index_t* indices, + const type_t* values, + const matrix_t B, + matrix_t C) { + for (auto row : config.tiles()) { + for (auto col : + custom_stride_range(std::size_t(0), B.cols, std::size_t(1))) { + type_t sum = 0; + for (auto nz : config.atoms(row)) { + sum += values[nz] * B(nz, col); + } + + // Output + C(row, col) = sum; + } + } +} + +/** + * @brief Sparse-Matrix Matrix Multiplication API. + * + * @tparam index_t Type of column indices. + * @tparam offset_t Type of row offsets. + * @tparam type_t Type of values. + * @param csr CSR matrix (GPU). + * @param n Number of columns in the B-matrix. + * @param B Input matrix B (GPU). + * @param C Output matrix C (GPU). + * @param stream CUDA stream. + */ +template +void thread_mapped(csr_t& csr, + matrix_t& B, + matrix_t& C, + cudaStream_t stream = 0) { + // Create a schedule. + constexpr std::size_t block_size = 128; + + /// Set-up kernel launch parameters and run the kernel. + + // Create a schedule. + using setup_t = schedule::setup; + setup_t config(csr.offsets.data().get(), csr.rows, csr.nnzs); + + std::size_t grid_size = (csr.rows + block_size - 1) / block_size; + launch::non_cooperative( + stream, __thread_mapped, grid_size, + block_size, config, csr.rows, csr.cols, csr.nnzs, + csr.offsets.data().get(), csr.indices.data().get(), + csr.values.data().get(), B, C); + + cudaStreamSynchronize(stream); +} + +} // namespace spmm +} // namespace algorithms +} // namespace loops \ No newline at end of file diff --git a/include/loops/container/matrix.cuh b/include/loops/container/matrix.cuh new file mode 100644 index 0000000..f520668 --- /dev/null +++ b/include/loops/container/matrix.cuh @@ -0,0 +1,53 @@ +#pragma once + +#include +#include + +namespace loops { + +template +struct matrix_t { + std::size_t rows; + std::size_t cols; + + vector_t m_data; + value_t* m_data_ptr; + + matrix_t() : rows(0), cols(0), m_data(), m_data_ptr(nullptr) {} + + matrix_t(std::size_t r, std::size_t c) + : rows(r), + cols(c), + m_data(r * c), + m_data_ptr(memory::raw_pointer_cast(m_data.data())) {} + + __host__ __device__ matrix_t(const matrix_t& other) + : rows(other.rows), cols(other.cols), m_data_ptr(other.m_data_ptr) {} + + __host__ __device__ __forceinline__ value_t operator()(int r, int c) const { + std::size_t idx = (cols * r) + c; + return m_data_ptr[idx]; + } + + __host__ __device__ __forceinline__ value_t& operator()(int r, int c) { + std::size_t idx = (cols * r) + c; + return m_data_ptr[idx]; + } + + __host__ __device__ __forceinline__ value_t + operator[](std::size_t index) const { + std::size_t r = index / cols; + std::size_t c = index % cols; + std::size_t idx = (cols * r) + c; + return m_data_ptr[idx]; + } + + __host__ __device__ __forceinline__ value_t& operator[](std::size_t index) { + std::size_t r = index / cols; + std::size_t c = index % cols; + std::size_t idx = (cols * r) + c; + return m_data_ptr[idx]; + } +}; + +} // namespace loops \ No newline at end of file diff --git a/scripts/spmv.py b/scripts/spmv.py new file mode 100644 index 0000000..9bc1f40 --- /dev/null +++ b/scripts/spmv.py @@ -0,0 +1,23 @@ +import subprocess +from pathlib import Path + + +spmv_variants = ["work_oriented", "group_mapped", "original", "thread_mapped"] + +# Modify the following lines to change the parameters of the experiment. +executable_path = "~/loops/build/bin" +dataset_path = "~/essentials/datasets" +datasets = ["chesapeake", "hollywood-2009"] + +# Build commands for each variand and dataset. +commands = [] +for variant in spmv_variants: + for dataset in datasets: + commands.append(str(executable_path) + "/loops.spmv." + str(variant) + " -m " + str(dataset_path) + "/" + str(dataset) + "/" + str(dataset) + ".mtx" + ' -v --validate') + +# Run all commands. +for command in commands: + popen = subprocess.Popen(command, stdout=subprocess.PIPE, shell=True) + popen.wait() + output = popen.stdout.read() + print(output) \ No newline at end of file From 485f7d3f41cbfcc9fa23fc8c97428568720bff3c Mon Sep 17 00:00:00 2001 From: neoblizz Date: Fri, 22 Jul 2022 11:14:18 -0700 Subject: [PATCH 2/8] Debugging loaders. --- include/loops/container/market.hxx | 73 +++++++++++++++++------------- 1 file changed, 42 insertions(+), 31 deletions(-) diff --git a/include/loops/container/market.hxx b/include/loops/container/market.hxx index e0592e3..1f37f25 100644 --- a/include/loops/container/market.hxx +++ b/include/loops/container/market.hxx @@ -81,6 +81,9 @@ struct matrix_market_t { matrix_market_data_t data; // Data type matrix_market_storage_scheme_t scheme; // Storage scheme + // mtx are generally written as coordinate format + coo_t coo; + matrix_market_t() {} ~matrix_market_t() {} @@ -116,10 +119,13 @@ struct matrix_market_t { exit(1); } - // mtx are generally written as coordinate formaat - coo_t coo( - (std::size_t)num_rows, (std::size_t)num_columns, - (std::size_t)num_nonzeros); + // Allocate memory for the matrix. + coo.rows = (std::size_t)num_rows; + coo.cols = (std::size_t)num_columns; + coo.nnzs = (std::size_t)num_nonzeros; + coo.row_indices.resize(num_nonzeros); + coo.col_indices.resize(num_nonzeros); + coo.values.resize(num_nonzeros); if (mm_is_coordinate(code)) format = matrix_market_format_t::coordinate; @@ -131,11 +137,22 @@ struct matrix_market_t { // pattern matrix defines sparsity pattern, but not values for (vertex_t i = 0; i < num_nonzeros; ++i) { - assert(fscanf(file, " %d %d \n", &(coo.row_indices[i]), - &(coo.col_indices[i])) == 2); - coo.row_indices[i]--; // adjust from 1-based to 0-based indexing - coo.col_indices[i]--; - coo.values[i] = (weight_t)1.0; // use value 1.0 for all nonzero entries + vertex_t I = 0; + vertex_t J = 0; + assert(fscanf(file, " %d %d \n", &I, &J) == 2); + + if (i < 40) + std::cout << "(" << I << ", " << J << ") " << std::endl; + // adjust from 1-based to 0-based indexing + coo.row_indices[i] = (vertex_t)I - 1; + coo.col_indices[i] = (vertex_t)J - 1; + + // use value 1.0 for all nonzero entries + coo.values[i] = (weight_t)1.0; + + if (i < 40) + std::cout << "(" << coo.row_indices[i] << ", " << coo.col_indices[i] + << ") = " << coo.values[i] << std::endl; } } else if (mm_is_real(code) || mm_is_integer(code)) { if (mm_is_real(code)) @@ -167,38 +184,32 @@ struct matrix_market_t { ++off_diagonals; } - vertex_t _nonzeros = 2 * off_diagonals + (coo.nnzs - off_diagonals); - - vector_t new_I(_nonzeros); - vector_t new_J(_nonzeros); - vector_t new_V(_nonzeros); - - vertex_t* _I = new_I.data(); - vertex_t* _J = new_J.data(); - weight_t* _V = new_V.data(); + // Duplicate off-diagonal entries for symmetric matrix. + std::size_t _nonzeros = 2 * off_diagonals + (coo.nnzs - off_diagonals); + coo_t temp(coo.rows, coo.cols, + _nonzeros); vertex_t ptr = 0; for (vertex_t i = 0; i < coo.nnzs; ++i) { if (coo.row_indices[i] != coo.col_indices[i]) { - _I[ptr] = coo.row_indices[i]; - _J[ptr] = coo.col_indices[i]; - _V[ptr] = coo.values[i]; + temp.row_indices[ptr] = coo.row_indices[i]; + temp.col_indices[ptr] = coo.col_indices[i]; + temp.values[ptr] = coo.values[i]; ++ptr; - _J[ptr] = coo.row_indices[i]; - _I[ptr] = coo.col_indices[i]; - _V[ptr] = coo.values[i]; + temp.col_indices[ptr] = coo.row_indices[i]; + temp.row_indices[ptr] = coo.col_indices[i]; + temp.values[ptr] = coo.values[i]; ++ptr; } else { - _I[ptr] = coo.row_indices[i]; - _J[ptr] = coo.col_indices[i]; - _V[ptr] = coo.values[i]; + temp.row_indices[ptr] = coo.row_indices[i]; + temp.col_indices[ptr] = coo.col_indices[i]; + temp.values[ptr] = coo.values[i]; ++ptr; } } - coo.row_indices = new_I; - coo.col_indices = new_J; - coo.values = new_V; - coo.nnzs = _nonzeros; + + // Move data to the original COO matrix. + coo = temp; } // end symmetric case fclose(file); From 9eac64f2f61472fbc8f528ef69381af78b347fec Mon Sep 17 00:00:00 2001 From: neoblizz Date: Fri, 22 Jul 2022 12:26:05 -0700 Subject: [PATCH 3/8] Known issue: nvbench messes up mtx.When benchmarking is enabled, nvbench is somehow causing the matrix-market loader to report nothing on the reads of I, J, K coordinate values read from the file. Resulting in a COO that is all (-1, -1) = 1 as entries. --- CMakeLists.txt | 2 +- examples/spmv/CMakeLists.txt | 17 ++++++++++++----- include/loops/container/market.hxx | 26 ++++++++++++++------------ 3 files changed, 27 insertions(+), 18 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fb452fa..9668aef 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -198,7 +198,7 @@ endif(LOOPS_BUILD_TESTS) #################################################### option(LOOPS_BUILD_BENCHMARKS "If on, builds loops with benchmarking support." - ON) + OFF) # Subdirectories for examples, testing and documentation if(LOOPS_BUILD_BENCHMARKS) diff --git a/examples/spmv/CMakeLists.txt b/examples/spmv/CMakeLists.txt index 5a52cda..ae56709 100644 --- a/examples/spmv/CMakeLists.txt +++ b/examples/spmv/CMakeLists.txt @@ -8,12 +8,19 @@ set(SOURCES foreach(SOURCE IN LISTS SOURCES) get_filename_component(TEST_NAME ${SOURCE} NAME_WLE) - add_executable(loops.spmv.${TEST_NAME} ${SOURCE}) - target_link_libraries(loops.spmv.${TEST_NAME} PRIVATE loops) - set_target_properties(loops.spmv.${TEST_NAME} - PROPERTIES + string(PREPEND TEST_NAME "loops.spmv.") + add_executable(${TEST_NAME} ${SOURCE}) + target_link_libraries(${TEST_NAME} + PRIVATE loops + ) + + set_target_properties(${TEST_NAME} + PROPERTIES + CXX_STANDARD 17 + CUDA_STANDARD 17 CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES} ) - message(STATUS "Example Added: loops.spmv.${TEST_NAME}") + + message(STATUS "Example Added: ${TEST_NAME}") endforeach() # end /* Add application */ \ No newline at end of file diff --git a/include/loops/container/market.hxx b/include/loops/container/market.hxx index 1f37f25..fb7da1f 100644 --- a/include/loops/container/market.hxx +++ b/include/loops/container/market.hxx @@ -95,31 +95,33 @@ struct matrix_market_t { * @param _filename input file name (.mtx) * @return coordinate sparse format */ - auto load(std::string _filename) { + coo_t& load(std::string _filename) { filename = _filename; dataset = extract_dataset(extract_filename(filename)); file_t file; - // Load MTX information + /// Load MTX information from the file. if ((file = fopen(filename.c_str(), "r")) == NULL) { std::cerr << "File could not be opened: " << filename << std::endl; exit(1); } + /// TODO: Add support for mtx with no banners. if (mm_read_banner(file, &code) != 0) { std::cerr << "Could not process Matrix Market banner" << std::endl; exit(1); } - int num_rows, num_columns, num_nonzeros; // XXX: requires all ints intially + /// TODO: Update C-interface to support unsigned ints instead. + int num_rows, num_columns, num_nonzeros; if ((mm_read_mtx_crd_size(file, &num_rows, &num_columns, &num_nonzeros)) != 0) { std::cerr << "Could not read file info (M, N, NNZ)" << std::endl; exit(1); } - // Allocate memory for the matrix. + /// Allocate memory for the matrix. coo.rows = (std::size_t)num_rows; coo.cols = (std::size_t)num_columns; coo.nnzs = (std::size_t)num_nonzeros; @@ -132,6 +134,7 @@ struct matrix_market_t { else format = matrix_market_format_t::array; + /// Pattern matrices do not have nonzero values. if (mm_is_pattern(code)) { data = matrix_market_data_t::pattern; @@ -141,20 +144,17 @@ struct matrix_market_t { vertex_t J = 0; assert(fscanf(file, " %d %d \n", &I, &J) == 2); - if (i < 40) - std::cout << "(" << I << ", " << J << ") " << std::endl; // adjust from 1-based to 0-based indexing coo.row_indices[i] = (vertex_t)I - 1; coo.col_indices[i] = (vertex_t)J - 1; // use value 1.0 for all nonzero entries coo.values[i] = (weight_t)1.0; - - if (i < 40) - std::cout << "(" << coo.row_indices[i] << ", " << coo.col_indices[i] - << ") = " << coo.values[i] << std::endl; } - } else if (mm_is_real(code) || mm_is_integer(code)) { + } + + /// Real or Integer matrices have real or integer values. + else if (mm_is_real(code) || mm_is_integer(code)) { if (mm_is_real(code)) data = matrix_market_data_t::real; else @@ -176,8 +176,10 @@ struct matrix_market_t { exit(1); } - if (mm_is_symmetric(code)) { // duplicate off diagonal entries + /// Symmetric matrices have symmetric halves. + if (mm_is_symmetric(code)) { scheme = matrix_market_storage_scheme_t::symmetric; + vertex_t off_diagonals = 0; for (vertex_t i = 0; i < coo.nnzs; ++i) { if (coo.row_indices[i] != coo.col_indices[i]) From fbf112a0f15d3bfcd581a41d3e7e2c1f284b3610 Mon Sep 17 00:00:00 2001 From: neoblizz Date: Sun, 24 Jul 2022 15:12:48 -0700 Subject: [PATCH 4/8] Changes to support nvbench (assert causes issues). --- include/loops/container/detail/mmio.cpp | 98 +++++++++-------- include/loops/container/detail/mmio.hxx | 55 ++++++---- include/loops/container/market.hxx | 139 ++++++++++++------------ include/loops/error.hxx | 49 +++++++++ 4 files changed, 206 insertions(+), 135 deletions(-) create mode 100644 include/loops/error.hxx diff --git a/include/loops/container/detail/mmio.cpp b/include/loops/container/detail/mmio.cpp index 8ff672a..b4b8646 100644 --- a/include/loops/container/detail/mmio.cpp +++ b/include/loops/container/detail/mmio.cpp @@ -10,22 +10,24 @@ #include #include #include +#include #include +#include int mm_read_unsymmetric_sparse(const char* fname, - int* M_, - int* N_, - int* nz_, + std::size_t* M_, + std::size_t* N_, + std::size_t* nz_, double** val_, - int** I_, - int** J_) { + std::size_t** I_, + std::size_t** J_) { FILE* f; MM_typecode matcode; - int M, N, nz; - int i; + std::size_t M, N, nz; + std::size_t i; double* val; - int *I, *J; + std::size_t *I, *J; if ((f = fopen(fname, "r")) == NULL) return -1; @@ -57,8 +59,8 @@ int mm_read_unsymmetric_sparse(const char* fname, /* reseve memory for matrices */ - I = (int*)malloc(nz * sizeof(int)); - J = (int*)malloc(nz * sizeof(int)); + I = (std::size_t*)malloc(nz * sizeof(std::size_t)); + J = (std::size_t*)malloc(nz * sizeof(std::size_t)); val = (double*)malloc(nz * sizeof(double)); *val_ = val; @@ -70,8 +72,11 @@ int mm_read_unsymmetric_sparse(const char* fname, /* (ANSI C X3.159-1989, Sec. 4.9.6.2, p. 136 lines 13-15) */ for (i = 0; i < nz; ++i) { - if (fscanf(f, "%d %d %lg\n", &I[i], &J[i], &val[i])) + if (fscanf(f, "%zu %zu %lg\n", &I[i], &J[i], &val[i])) ; + + loops::error::throw_if_exception(I[i] == 0, "Market file is zero-indexed"); + loops::error::throw_if_exception(J[i] == 0, "Market file is zero-indexed"); I[i]--; /* adjust from 1-based to 0-based */ J[i]--; } @@ -175,7 +180,10 @@ int mm_write_mtx_crd_size(FILE* f, int M, int N, int nz) { return 0; } -int mm_read_mtx_crd_size(FILE* f, int* M, int* N, int* nz) { +int mm_read_mtx_crd_size(FILE* f, + std::size_t* M, + std::size_t* N, + std::size_t* nz) { char line[MM_MAX_LINE_LENGTH]; int num_items_read; @@ -189,12 +197,12 @@ int mm_read_mtx_crd_size(FILE* f, int* M, int* N, int* nz) { } while (line[0] == '%'); /* line[] is either blank or has M,N, nz */ - if (sscanf(line, "%d %d %d", M, N, nz) == 3) + if (sscanf(line, "%zu %zu %zu", M, N, nz) == 3) return 0; else do { - num_items_read = fscanf(f, "%d %d %d", M, N, nz); + num_items_read = fscanf(f, "%zu %zu %zu", M, N, nz); if (num_items_read == EOF) return MM_PREMATURE_EOF; } while (num_items_read != 3); @@ -242,29 +250,29 @@ int mm_write_mtx_array_size(FILE* f, int M, int N) { /******************************************************************/ int mm_read_mtx_crd_data(FILE* f, - int M, - int N, - int nz, - int I[], - int J[], + std::size_t M, + std::size_t N, + std::size_t nz, + std::size_t I[], + std::size_t J[], double val[], MM_typecode matcode) { - int i; + std::size_t i; if (mm_is_complex(matcode)) { for (i = 0; i < nz; i++) - if (fscanf(f, "%d %d %lg %lg", &I[i], &J[i], &val[2 * i], + if (fscanf(f, "%zu %zu %lg %lg", &I[i], &J[i], &val[2 * i], &val[2 * i + 1]) != 4) return MM_PREMATURE_EOF; } else if (mm_is_real(matcode)) { for (i = 0; i < nz; i++) { - if (fscanf(f, "%d %d %lg\n", &I[i], &J[i], &val[i]) != 3) + if (fscanf(f, "%zu %zu %lg\n", &I[i], &J[i], &val[i]) != 3) return MM_PREMATURE_EOF; } } else if (mm_is_pattern(matcode)) { for (i = 0; i < nz; i++) - if (fscanf(f, "%d %d", &I[i], &J[i]) != 2) + if (fscanf(f, "%zu %zu", &I[i], &J[i]) != 2) return MM_PREMATURE_EOF; } else return MM_UNSUPPORTED_TYPE; @@ -273,22 +281,22 @@ int mm_read_mtx_crd_data(FILE* f, } int mm_read_mtx_crd_entry(FILE* f, - int* I, - int* J, + std::size_t* I, + std::size_t* J, double* real, double* imag, MM_typecode matcode) { if (mm_is_complex(matcode)) { - if (fscanf(f, "%d %d %lg %lg", I, J, real, imag) != 4) + if (fscanf(f, "%zu %zu %lg %lg", I, J, real, imag) != 4) return MM_PREMATURE_EOF; } else if (mm_is_real(matcode)) { - if (fscanf(f, "%d %d %lg\n", I, J, real) != 3) + if (fscanf(f, "%zu %zu %lg\n", I, J, real) != 3) return MM_PREMATURE_EOF; } else if (mm_is_pattern(matcode)) { - if (fscanf(f, "%d %d", I, J) != 2) + if (fscanf(f, "%zu %zu", I, J) != 2) return MM_PREMATURE_EOF; } else return MM_UNSUPPORTED_TYPE; @@ -305,11 +313,11 @@ int mm_read_mtx_crd_entry(FILE* f, ************************************************************************/ int mm_read_mtx_crd(char* fname, - int* M, - int* N, - int* nz, - int** I, - int** J, + std::size_t* M, + std::size_t* N, + std::size_t* nz, + std::size_t** I, + std::size_t** J, double** val, MM_typecode* matcode) { int ret_code; @@ -330,8 +338,8 @@ int mm_read_mtx_crd(char* fname, if ((ret_code = mm_read_mtx_crd_size(f, M, N, nz)) != 0) return ret_code; - *I = (int*)malloc(*nz * sizeof(int)); - *J = (int*)malloc(*nz * sizeof(int)); + *I = (std::size_t*)malloc(*nz * sizeof(std::size_t)); + *J = (std::size_t*)malloc(*nz * sizeof(std::size_t)); *val = NULL; if (mm_is_complex(*matcode)) { @@ -370,15 +378,15 @@ int mm_write_banner(FILE* f, MM_typecode matcode) { } int mm_write_mtx_crd(char fname[], - int M, - int N, - int nz, - int I[], - int J[], + std::size_t M, + std::size_t N, + std::size_t nz, + std::size_t I[], + std::size_t J[], double val[], MM_typecode matcode) { FILE* f; - int i; + std::size_t i; if (strcmp(fname, "stdout") == 0) f = stdout; @@ -390,18 +398,18 @@ int mm_write_mtx_crd(char fname[], fprintf(f, "%s\n", mm_typecode_to_str(matcode)); /* print matrix sizes and nonzeros */ - fprintf(f, "%d %d %d\n", M, N, nz); + fprintf(f, "%zu %zu %zu\n", M, N, nz); /* print values */ if (mm_is_pattern(matcode)) for (i = 0; i < nz; i++) - fprintf(f, "%d %d\n", I[i], J[i]); + fprintf(f, "%zu %zu\n", I[i], J[i]); else if (mm_is_real(matcode)) for (i = 0; i < nz; i++) - fprintf(f, "%d %d %20.16g\n", I[i], J[i], val[i]); + fprintf(f, "%zu %zu %20.16g\n", I[i], J[i], val[i]); else if (mm_is_complex(matcode)) for (i = 0; i < nz; i++) - fprintf(f, "%d %d %20.16g %20.16g\n", I[i], J[i], val[2 * i], + fprintf(f, "%zu %zu %20.16g %20.16g\n", I[i], J[i], val[2 * i], val[2 * i + 1]); else { if (f != stdout) diff --git a/include/loops/container/detail/mmio.hxx b/include/loops/container/detail/mmio.hxx index 73e933a..ba76168 100644 --- a/include/loops/container/detail/mmio.hxx +++ b/include/loops/container/detail/mmio.hxx @@ -4,12 +4,14 @@ * @brief Matrix-Market file format header file, see mmio.cpp for implementation * details. * @version 0.1 - * @date 2022-02-03 + * @date 2020-10-12 * - * @copyright Copyright (c) 2022 + * @copyright Copyright (c) 2020 * */ +#pragma once + #if defined(__cplusplus) extern "C" { #endif @@ -34,12 +36,19 @@ typedef char MM_typecode[4]; char* mm_typecode_to_str(MM_typecode matcode); int mm_read_banner(FILE* f, MM_typecode* matcode); -int mm_read_mtx_crd_size(FILE* f, int* M, int* N, int* nz); -int mm_read_mtx_array_size(FILE* f, int* M, int* N); +int mm_read_mtx_crd_size(FILE* f, + std::size_t* M, + std::size_t* N, + std::size_t* nz); + +int mm_read_mtx_array_size(FILE* f, std::size_t* M, std::size_t* N); int mm_write_banner(FILE* f, MM_typecode matcode); -int mm_write_mtx_crd_size(FILE* f, int M, int N, int nz); -int mm_write_mtx_array_size(FILE* f, int M, int N); +int mm_write_mtx_crd_size(FILE* f, + std::size_t M, + std::size_t N, + std::size_t nz); +int mm_write_mtx_array_size(FILE* f, std::size_t M, std::size_t N); /********************* MM_typecode query fucntions ***************************/ @@ -126,37 +135,37 @@ int mm_is_valid(MM_typecode matcode); /* too complex for a macro */ /* high level routines */ int mm_write_mtx_crd(char fname[], - int M, - int N, - int nz, - int I[], - int J[], + std::size_t M, + std::size_t N, + std::size_t nz, + std::size_t I[], + std::size_t J[], double val[], MM_typecode matcode); int mm_read_mtx_crd_data(FILE* f, - int M, - int N, - int nz, - int I[], - int J[], + std::size_t M, + std::size_t N, + std::size_t nz, + std::size_t I[], + std::size_t J[], double val[], MM_typecode matcode); int mm_read_mtx_crd_entry(FILE* f, - int* I, - int* J, + std::size_t* I, + std::size_t* J, double* real, double* img, MM_typecode matcode); int mm_read_unsymmetric_sparse(const char* fname, - int* M_, - int* N_, - int* nz_, + std::size_t* M_, + std::size_t* N_, + std::size_t* nz_, double** val_, - int** I_, - int** J_); + std::size_t** I_, + std::size_t** J_); #endif diff --git a/include/loops/container/market.hxx b/include/loops/container/market.hxx index fb7da1f..4f2d083 100644 --- a/include/loops/container/market.hxx +++ b/include/loops/container/market.hxx @@ -3,20 +3,23 @@ * @author Muhammad Osama (mosama@ucdavis.edu) * @brief * @version 0.1 - * @date 2022-02-03 + * @date 2020-10-09 * - * @copyright Copyright (c) 2022 + * @copyright Copyright (c) 2020 * */ #pragma once #include +#include #include + #include #include #include +#include namespace loops { @@ -64,7 +67,7 @@ enum matrix_market_storage_scheme_t { general, hermitian, symmetric, skew }; * * Indices are 1-based i.2. A(1,1) is the first element. */ -template +template struct matrix_market_t { // typedef FILE* file_t; // typedef MM_typecode matrix_market_code_t; @@ -81,9 +84,6 @@ struct matrix_market_t { matrix_market_data_t data; // Data type matrix_market_storage_scheme_t scheme; // Storage scheme - // mtx are generally written as coordinate format - coo_t coo; - matrix_market_t() {} ~matrix_market_t() {} @@ -95,129 +95,134 @@ struct matrix_market_t { * @param _filename input file name (.mtx) * @return coordinate sparse format */ - coo_t& load(std::string _filename) { + auto load(std::string _filename) { filename = _filename; dataset = extract_dataset(extract_filename(filename)); file_t file; - /// Load MTX information from the file. + // Load MTX information if ((file = fopen(filename.c_str(), "r")) == NULL) { std::cerr << "File could not be opened: " << filename << std::endl; exit(1); } - /// TODO: Add support for mtx with no banners. if (mm_read_banner(file, &code) != 0) { std::cerr << "Could not process Matrix Market banner" << std::endl; exit(1); } - /// TODO: Update C-interface to support unsigned ints instead. - int num_rows, num_columns, num_nonzeros; + std::size_t num_rows, num_columns, num_nonzeros; if ((mm_read_mtx_crd_size(file, &num_rows, &num_columns, &num_nonzeros)) != 0) { std::cerr << "Could not read file info (M, N, NNZ)" << std::endl; exit(1); } - /// Allocate memory for the matrix. - coo.rows = (std::size_t)num_rows; - coo.cols = (std::size_t)num_columns; - coo.nnzs = (std::size_t)num_nonzeros; - coo.row_indices.resize(num_nonzeros); - coo.col_indices.resize(num_nonzeros); - coo.values.resize(num_nonzeros); + error::throw_if_exception( + num_rows >= std::numeric_limits::max() || + num_columns >= std::numeric_limits::max(), + "index_t overflow"); + error::throw_if_exception( + num_nonzeros >= std::numeric_limits::max(), + "offset_t overflow"); + + // mtx are generally written as coordinate format + coo_t coo( + (index_t)num_rows, (index_t)num_columns, (offset_t)num_nonzeros); if (mm_is_coordinate(code)) format = matrix_market_format_t::coordinate; else format = matrix_market_format_t::array; - /// Pattern matrices do not have nonzero values. if (mm_is_pattern(code)) { data = matrix_market_data_t::pattern; // pattern matrix defines sparsity pattern, but not values - for (vertex_t i = 0; i < num_nonzeros; ++i) { - vertex_t I = 0; - vertex_t J = 0; - assert(fscanf(file, " %d %d \n", &I, &J) == 2); - - // adjust from 1-based to 0-based indexing - coo.row_indices[i] = (vertex_t)I - 1; - coo.col_indices[i] = (vertex_t)J - 1; - - // use value 1.0 for all nonzero entries - coo.values[i] = (weight_t)1.0; + for (index_t i = 0; i < num_nonzeros; ++i) { + std::size_t row_index{0}, col_index{0}; + auto num_assigned = fscanf(file, " %zu %zu \n", &row_index, &col_index); + error::throw_if_exception(num_assigned != 2, + "Could not read edge from market file"); + error::throw_if_exception(row_index == 0, + "Market file is zero-indexed"); + error::throw_if_exception(col_index == 0, + "Market file is zero-indexed"); + // set and adjust from 1-based to 0-based indexing + coo.row_indices[i] = (index_t)row_index - 1; + coo.col_indices[i] = (index_t)col_index - 1; + coo.values[i] = (type_t)1.0; // use value 1.0 for all nonzero entries } - } - - /// Real or Integer matrices have real or integer values. - else if (mm_is_real(code) || mm_is_integer(code)) { + } else if (mm_is_real(code) || mm_is_integer(code)) { if (mm_is_real(code)) data = matrix_market_data_t::real; else data = matrix_market_data_t::integer; - for (vertex_t i = 0; i < coo.nnzs; ++i) { - vertex_t I = 0; - vertex_t J = 0; - double V = 0.0f; + for (index_t i = 0; i < coo.nnzs; ++i) { + std::size_t row_index{0}, col_index{0}; + double weight{0.0}; - assert(fscanf(file, " %d %d %lf \n", &I, &J, &V) == 3); + auto num_assigned = + fscanf(file, " %zu %zu %lf \n", &row_index, &col_index, &weight); - coo.row_indices[i] = (vertex_t)I - 1; - coo.col_indices[i] = (vertex_t)J - 1; - coo.values[i] = (weight_t)V; + error::throw_if_exception( + num_assigned != 3, "Could not read weighted edge from market file"); + error::throw_if_exception(row_index == 0, + "Market file is zero-indexed"); + error::throw_if_exception(col_index == 0, + "Market file is zero-indexed"); + + coo.row_indices[i] = (index_t)row_index - 1; + coo.col_indices[i] = (index_t)col_index - 1; + coo.values[i] = (type_t)weight; } } else { std::cerr << "Unrecognized matrix market format type" << std::endl; exit(1); } - /// Symmetric matrices have symmetric halves. - if (mm_is_symmetric(code)) { + if (mm_is_symmetric(code)) { // duplicate off diagonal entries scheme = matrix_market_storage_scheme_t::symmetric; - - vertex_t off_diagonals = 0; - for (vertex_t i = 0; i < coo.nnzs; ++i) { + index_t off_diagonals = 0; + for (index_t i = 0; i < coo.nnzs; ++i) { if (coo.row_indices[i] != coo.col_indices[i]) ++off_diagonals; } - // Duplicate off-diagonal entries for symmetric matrix. - std::size_t _nonzeros = 2 * off_diagonals + (coo.nnzs - off_diagonals); - coo_t temp(coo.rows, coo.cols, - _nonzeros); + index_t _nonzeros = 2 * off_diagonals + (coo.nnzs - off_diagonals); + + vector_t _I(_nonzeros); + vector_t _J(_nonzeros); + vector_t _V(_nonzeros); - vertex_t ptr = 0; - for (vertex_t i = 0; i < coo.nnzs; ++i) { + index_t ptr = 0; + for (index_t i = 0; i < coo.nnzs; ++i) { if (coo.row_indices[i] != coo.col_indices[i]) { - temp.row_indices[ptr] = coo.row_indices[i]; - temp.col_indices[ptr] = coo.col_indices[i]; - temp.values[ptr] = coo.values[i]; + _I[ptr] = coo.row_indices[i]; + _J[ptr] = coo.col_indices[i]; + _V[ptr] = coo.values[i]; ++ptr; - temp.col_indices[ptr] = coo.row_indices[i]; - temp.row_indices[ptr] = coo.col_indices[i]; - temp.values[ptr] = coo.values[i]; + _J[ptr] = coo.row_indices[i]; + _I[ptr] = coo.col_indices[i]; + _V[ptr] = coo.values[i]; ++ptr; } else { - temp.row_indices[ptr] = coo.row_indices[i]; - temp.col_indices[ptr] = coo.col_indices[i]; - temp.values[ptr] = coo.values[i]; + _I[ptr] = coo.row_indices[i]; + _J[ptr] = coo.col_indices[i]; + _V[ptr] = coo.values[i]; ++ptr; } } - - // Move data to the original COO matrix. - coo = temp; + coo.row_indices = _I; + coo.col_indices = _J; + coo.values = _V; + coo.nnzs = _nonzeros; } // end symmetric case fclose(file); - return coo; } }; - } // namespace loops \ No newline at end of file diff --git a/include/loops/error.hxx b/include/loops/error.hxx new file mode 100644 index 0000000..b6cd291 --- /dev/null +++ b/include/loops/error.hxx @@ -0,0 +1,49 @@ +#pragma once + +#include +#include +#include + +namespace loops { + +/** + * @namespace error + * Error utilities for exception handling within device and host code. + */ +namespace error { + +typedef cudaError_t error_t; + +/** + * @brief Exception class for errors in device code. + * + */ +struct exception_t : std::exception { + std::string report; + + exception_t(error_t _status, std::string _message = "") { + report = cudaGetErrorString(_status) + std::string("\t: ") + _message; + } + + exception_t(std::string _message = "") { report = _message; } + virtual const char* what() const noexcept { return report.c_str(); } +}; + +/** + * @brief Throw an exception if the given error code is not cudaSuccess. + * + * @param status error_t error code (equivalent to cudaError_t). + * @param message custom message to be appended to the error message. + */ +inline void throw_if_exception(error_t status, std::string message = "") { + if (status != cudaSuccess) + throw exception_t(status, message); +} + +inline void throw_if_exception(bool is_exception, std::string message = "") { + if (is_exception) + throw exception_t(message); +} + +} // namespace error +} // namespace loops \ No newline at end of file From bdab92426044cc86d01990cab8b08f4653ac74b8 Mon Sep 17 00:00:00 2001 From: neoblizz Date: Mon, 25 Jul 2022 09:24:14 -0700 Subject: [PATCH 5/8] Colors are always nice. --- cmake/FetchColors.cmake | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) create mode 100644 cmake/FetchColors.cmake diff --git a/cmake/FetchColors.cmake b/cmake/FetchColors.cmake new file mode 100644 index 0000000..a171487 --- /dev/null +++ b/cmake/FetchColors.cmake @@ -0,0 +1,19 @@ +if(NOT WIN32) + string(ASCII 27 Esc) + set(ColourReset "${Esc}[m") + set(ColourBold "${Esc}[1m") + set(Red "${Esc}[31m") + set(Green "${Esc}[32m") + set(Yellow "${Esc}[33m") + set(Blue "${Esc}[34m") + set(Magenta "${Esc}[35m") + set(Cyan "${Esc}[36m") + set(White "${Esc}[37m") + set(BoldRed "${Esc}[1;31m") + set(BoldGreen "${Esc}[1;32m") + set(BoldYellow "${Esc}[1;33m") + set(BoldBlue "${Esc}[1;34m") + set(BoldMagenta "${Esc}[1;35m") + set(BoldCyan "${Esc}[1;36m") + set(BoldWhite "${Esc}[1;37m") +endif() \ No newline at end of file From 9420fcd12a93baec3f64dd58f7b7e998568c8109 Mon Sep 17 00:00:00 2001 From: neoblizz Date: Mon, 25 Jul 2022 09:24:27 -0700 Subject: [PATCH 6/8] NVBench issue resolved, enabled. --- CMakeLists.txt | 19 ++++++++++++++----- include/loops/container/market.hxx | 3 ++- 2 files changed, 16 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 9668aef..7e74edf 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -34,6 +34,7 @@ set(PROJECT_DEPS_DIR externals) # end /* Dependencies directory */ # begin /* Include cmake modules */ +include(${PROJECT_SOURCE_DIR}/cmake/FetchColors.cmake) include(${PROJECT_SOURCE_DIR}/cmake/FetchThrustCUB.cmake) include(${PROJECT_SOURCE_DIR}/cmake/FetchModernGPU.cmake) include(${PROJECT_SOURCE_DIR}/cmake/FetchCXXOpts.cmake) @@ -58,13 +59,21 @@ set(CMAKE_VERBOSE_MAKEFILE OFF) ############ ADD LIBRARY: LOOPS (HEADER-ONLY) ############ add_library(loops INTERFACE) + #################################################### -############### SET TARGET PROPERTIES ############## +############### SET SM ARCHITECTURE ################ #################################################### -if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) - set(CMAKE_CUDA_ARCHITECTURES 75) -endif() +## Note: This applies to NVBench as well. +## Can be used for applications by extracting the +## CUDA_ARCHITECTURES property from loops project. +## see: get_target_properties() +message(STATUS "${Magenta}Set SM Architecture using -DCMAKE_CUDA_ARCHITECTURES=75,80${ColourReset}") +set(CMAKE_CUDA_ARCHITECTURES 75) + +#################################################### +############### SET TARGET PROPERTIES ############## +#################################################### set_target_properties(loops PROPERTIES CXX_STANDARD 17 @@ -198,7 +207,7 @@ endif(LOOPS_BUILD_TESTS) #################################################### option(LOOPS_BUILD_BENCHMARKS "If on, builds loops with benchmarking support." - OFF) + ON) # Subdirectories for examples, testing and documentation if(LOOPS_BUILD_BENCHMARKS) diff --git a/include/loops/container/market.hxx b/include/loops/container/market.hxx index 4f2d083..500975c 100644 --- a/include/loops/container/market.hxx +++ b/include/loops/container/market.hxx @@ -1,7 +1,8 @@ /** * @file matrix_market.hxx * @author Muhammad Osama (mosama@ucdavis.edu) - * @brief + * @brief Matrix Market format reader. + * @see http://math.nist.gov/MatrixMarket/ * @version 0.1 * @date 2020-10-09 * From f343eb96c25b7122657c1e206bfd2d98a35a4a68 Mon Sep 17 00:00:00 2001 From: neoblizz Date: Mon, 25 Jul 2022 10:37:59 -0700 Subject: [PATCH 7/8] Clean-up: header files and junk --- CMakeLists.txt | 1 - benchmarks/spmv/CMakeLists.txt | 2 +- benchmarks/spmv/parameters.hxx | 2 ++ benchmarks/spmv/test_benchmarks.sh | 52 --------------------------- benchmarks/spmv/work_oriented.cu | 42 +++++++++++++++------- include/loops/container/container.hxx | 15 ++++++++ include/loops/util/gitsha1.c.in | 2 -- include/loops/util/gitsha1.hxx | 1 - include/loops/util/gitsha1make.c | 5 --- 9 files changed, 48 insertions(+), 74 deletions(-) delete mode 100755 benchmarks/spmv/test_benchmarks.sh delete mode 100644 include/loops/util/gitsha1.c.in delete mode 100644 include/loops/util/gitsha1.hxx delete mode 100644 include/loops/util/gitsha1make.c diff --git a/CMakeLists.txt b/CMakeLists.txt index 7e74edf..b839b7f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -133,7 +133,6 @@ target_link_libraries(loops ################# TARGET SOURCES ################### #################################################### target_sources(loops - INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/util/gitsha1make.c" INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/container/detail/mmio.cpp" ) diff --git a/benchmarks/spmv/CMakeLists.txt b/benchmarks/spmv/CMakeLists.txt index 8f154ca..b1a9dd3 100644 --- a/benchmarks/spmv/CMakeLists.txt +++ b/benchmarks/spmv/CMakeLists.txt @@ -8,7 +8,7 @@ foreach(SOURCE IN LISTS BENCHMARK_SOURCES) add_executable(${BENCHMARK_NAME} ${SOURCE}) target_link_libraries(${BENCHMARK_NAME} PRIVATE loops - PRIVATE nvbench::main + # PRIVATE nvbench::main PRIVATE nvbench::nvbench ) diff --git a/benchmarks/spmv/parameters.hxx b/benchmarks/spmv/parameters.hxx index 9ef9b27..6b09907 100644 --- a/benchmarks/spmv/parameters.hxx +++ b/benchmarks/spmv/parameters.hxx @@ -17,6 +17,8 @@ #include +std::string filename; + struct parameters_t { std::string filename; bool help = false; diff --git a/benchmarks/spmv/test_benchmarks.sh b/benchmarks/spmv/test_benchmarks.sh deleted file mode 100755 index 93d7e86..0000000 --- a/benchmarks/spmv/test_benchmarks.sh +++ /dev/null @@ -1,52 +0,0 @@ -# ------------------------------------------------------------------------ -# Algorithm benchmarking tests -# Run this from build directory -# If error CUPTI_ERROR_INSUFFICIENT_PRIVILEGES: run with sudo -# Make sure to pass -DESSENTIALS_BUILD_BENCHMARKS=ON -DNVBench_ENABLE_CUPTI=ON to CMake -# ------------------------------------------------------------------------ - -#!/bin/bash - -DATASET_DIR="../datasets" -BIN_DIR="./bin" - -# Where to store output -JSON_DIR="json" - -# Used for all algorithms except SPGEMM -MATRIX_FILE="${DATASET_DIR}/chesapeake/chesapeake.mtx" - -# Used for Geo -COORDINATES_FILE="${DATASET_DIR}/geolocation/sample.labels" - -# Used for SPGEMM -A_MATRIX="${DATASET_DIR}/spgemm/a.mtx" -B_MATRIX="${DATASET_DIR}/spgemm/b.mtx" - -make bc_bench -make bfs_bench -make color_bench -make geo_bench -make hits_bench -make kcore_bench -make mst_bench -make ppr_bench -make pr_bench -make spgemm_bench -make spmv_bench -make sssp_bench -make tc_bench - -${BIN_DIR}/bc_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/bc.json -${BIN_DIR}/bfs_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/bfs.json -${BIN_DIR}/color_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/color.json -${BIN_DIR}/geo_bench -m ${MATRIX_FILE} -c ${COORDINATES_FILE} --json ${JSON_DIR}/geo.json -${BIN_DIR}/hits_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/hits.json -${BIN_DIR}/kcore_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/kcore.json -${BIN_DIR}/mst_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/mst.json -${BIN_DIR}/ppr_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/ppr.json -${BIN_DIR}/pr_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/pr.json -${BIN_DIR}/spgemm_bench -a ${A_MATRIX} -b ${B_MATRIX} --json ${JSON_DIR}/spgemm.json -${BIN_DIR}/spmv_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/spmv.json -${BIN_DIR}/sssp_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/sssp.json -${BIN_DIR}/tc_bench -m ${MATRIX_FILE} --json ${JSON_DIR}/tc.json diff --git a/benchmarks/spmv/work_oriented.cu b/benchmarks/spmv/work_oriented.cu index 2f2419b..12d77a6 100644 --- a/benchmarks/spmv/work_oriented.cu +++ b/benchmarks/spmv/work_oriented.cu @@ -11,11 +11,8 @@ #include "parameters.hxx" -#include -#include -#include -#include #include +#include #include #include @@ -30,17 +27,14 @@ void work_oriented_bench(nvbench::state& state, nvbench::type_list) { using offset_t = int; using type_t = value_t; - /// Get sample CSR matrix and create x, y vectors. - std::size_t rows = 1 << 20; - std::size_t cols = 1 << 20; - - csr_t csr; - generate::random::csr(rows, cols, 0.0000001, csr); + matrix_market_t mtx; + csr_t csr(mtx.load(filename)); vector_t x(csr.rows); vector_t y(csr.rows); - generate::random::uniform_distribution(x.begin(), x.end(), 1, 10); + generate::random::uniform_distribution(x.begin(), x.end(), type_t(1.0), + type_t(10.0)); #if LOOPS_CUPTI_SUPPORTED /// Add CUPTI metrics to collect for the state. @@ -59,4 +53,28 @@ void work_oriented_bench(nvbench::state& state, nvbench::type_list) { // Define a type_list to use for the type axis: using value_types = nvbench::type_list; -NVBENCH_BENCH_TYPES(work_oriented_bench, NVBENCH_TYPE_AXES(value_types)); \ No newline at end of file +NVBENCH_BENCH_TYPES(work_oriented_bench, NVBENCH_TYPE_AXES(value_types)); + +int main(int argc, char** argv) { + parameters_t params(argc, argv); + filename = params.filename; + + if (params.help) { + // Print NVBench help. + const char* args[1] = {"-h"}; + NVBENCH_MAIN_BODY(1, args); + } else { + // Create a new argument array without matrix filename to pass to NVBench. + char* args[argc - 2]; + int j = 0; + for (int i = 0; i < argc; i++) { + if (strcmp(argv[i], "--market") == 0 || strcmp(argv[i], "-m") == 0) { + i++; + continue; + } + args[j] = argv[i]; + j++; + } + NVBENCH_MAIN_BODY(argc - 2, args); + } +} \ No newline at end of file diff --git a/include/loops/container/container.hxx b/include/loops/container/container.hxx index e69de29..e55595f 100644 --- a/include/loops/container/container.hxx +++ b/include/loops/container/container.hxx @@ -0,0 +1,15 @@ +/** + * @file container.hxx + * @author Muhammad Osama (mosama@ucdavis.edu) + * @brief Includes all the container headers. + * @version 0.1 + * @date 2022-07-25 + * + * @copyright Copyright (c) 2022 + * + */ + +#include +#include +#include +#include \ No newline at end of file diff --git a/include/loops/util/gitsha1.c.in b/include/loops/util/gitsha1.c.in deleted file mode 100644 index 18461a3..0000000 --- a/include/loops/util/gitsha1.c.in +++ /dev/null @@ -1,2 +0,0 @@ -#define GIT_SHA1 "@GIT_SHA1@" -const char g_GIT_SHA1[] = GIT_SHA1; \ No newline at end of file diff --git a/include/loops/util/gitsha1.hxx b/include/loops/util/gitsha1.hxx deleted file mode 100644 index e6e0842..0000000 --- a/include/loops/util/gitsha1.hxx +++ /dev/null @@ -1 +0,0 @@ -extern const char g_GIT_SHA1[]; \ No newline at end of file diff --git a/include/loops/util/gitsha1make.c b/include/loops/util/gitsha1make.c deleted file mode 100644 index 96dc5be..0000000 --- a/include/loops/util/gitsha1make.c +++ /dev/null @@ -1,5 +0,0 @@ -#ifndef GIT_SHA1 -#define GIT_SHA1 "GIT-SHA-NOT-FOUND" -#endif - -const char g_GIT_SHA1[] = GIT_SHA1; From ae25951569f8231ad0be78197b00e76faf8ff37b Mon Sep 17 00:00:00 2001 From: neoblizz Date: Mon, 25 Jul 2022 12:29:15 -0700 Subject: [PATCH 8/8] Simplified main-code for benchmark. --- benchmarks/spmv/parameters.hxx | 59 +++++++++++++++++++++++--------- benchmarks/spmv/work_oriented.cu | 23 +------------ 2 files changed, 43 insertions(+), 39 deletions(-) diff --git a/benchmarks/spmv/parameters.hxx b/benchmarks/spmv/parameters.hxx index 6b09907..0d19f00 100644 --- a/benchmarks/spmv/parameters.hxx +++ b/benchmarks/spmv/parameters.hxx @@ -16,48 +16,73 @@ #include #include +#include std::string filename; struct parameters_t { - std::string filename; - bool help = false; - cxxopts::Options options; - /** * @brief Construct a new parameters object and parse command line arguments. * * @param argc Number of command line arguments. * @param argv Command line arguments. */ - parameters_t(int argc, char** argv) : options(argv[0], "SPMV Benchmarking") { - options.allow_unrecognised_options(); + parameters_t(int argc, char** argv) + : m_options(argv[0], "SPMV Benchmarking"), m_argc(argc) { + m_options.allow_unrecognised_options(); // Add command line options - options.add_options()("h,help", "Print help") // help + m_options.add_options()("h,help", "Print help") // help ("m,market", "Matrix file", cxxopts::value()); // mtx - // Parse command line arguments - auto result = options.parse(argc, argv); + // Parse command line arguments. + auto result = m_options.parse(argc, argv); + // Print help if requested if (result.count("help")) { - help = true; - std::cout << options.help({""}); + m_help = true; + std::cout << m_options.help({""}); std::cout << " [optional nvbench args]" << std::endl << std::endl; - // Do not exit so we also print NVBench help. - } else { + const char* argh[1] = {"-h"}; + NVBENCH_MAIN_BODY(1, argh); + } + + // Get matrix market file or error if not specified. + else { if (result.count("market") == 1) { - filename = result["market"].as(); - if (!loops::is_market(filename)) { - std::cout << options.help({""}); + this->m_filename = result["market"].as(); + filename = m_filename; + if (!loops::is_market(m_filename)) { + std::cout << m_options.help({""}); std::cout << " [optional nvbench args]" << std::endl << std::endl; std::exit(0); } + + // Remove loops parameters and pass the rest to nvbench. + for (int i = 0; i < argc; i++) { + if (strcmp(argv[i], "--market") == 0 || strcmp(argv[i], "-m") == 0) { + i++; + continue; + } + m_args.push_back(argv[i]); + } + } else { - std::cout << options.help({""}); + std::cout << m_options.help({""}); std::cout << " [optional nvbench args]" << std::endl << std::endl; std::exit(0); } } } + + /// Helpers for NVBENCH_MAIN_BODY call. + int nvbench_argc() { return m_argc - 2; } + auto nvbench_argv() { return m_args.data(); } + + private: + std::string m_filename; + cxxopts::Options m_options; + std::vector m_args; + bool m_help = false; + int m_argc; }; \ No newline at end of file diff --git a/benchmarks/spmv/work_oriented.cu b/benchmarks/spmv/work_oriented.cu index 12d77a6..efe396a 100644 --- a/benchmarks/spmv/work_oriented.cu +++ b/benchmarks/spmv/work_oriented.cu @@ -15,8 +15,6 @@ #include #include -#include - #define LOOPS_CUPTI_SUPPORTED 0 using namespace loops; @@ -57,24 +55,5 @@ NVBENCH_BENCH_TYPES(work_oriented_bench, NVBENCH_TYPE_AXES(value_types)); int main(int argc, char** argv) { parameters_t params(argc, argv); - filename = params.filename; - - if (params.help) { - // Print NVBench help. - const char* args[1] = {"-h"}; - NVBENCH_MAIN_BODY(1, args); - } else { - // Create a new argument array without matrix filename to pass to NVBench. - char* args[argc - 2]; - int j = 0; - for (int i = 0; i < argc; i++) { - if (strcmp(argv[i], "--market") == 0 || strcmp(argv[i], "-m") == 0) { - i++; - continue; - } - args[j] = argv[i]; - j++; - } - NVBENCH_MAIN_BODY(argc - 2, args); - } + NVBENCH_MAIN_BODY(params.nvbench_argc(), params.nvbench_argv()); } \ No newline at end of file