diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 02a9e70d..e7cf89e7 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -4,4 +4,5 @@ add_subdirectory(bfs) add_subdirectory(color) add_subdirectory(geo) add_subdirectory(pr) +add_subdirectory(mtx2bin) # end /* Add examples' subdirectories */ diff --git a/examples/bfs/bfs.cu b/examples/bfs/bfs.cu index 8f2f61da..4bf20ce4 100644 --- a/examples/bfs/bfs.cu +++ b/examples/bfs/bfs.cu @@ -1,4 +1,5 @@ #include +#include "bfs_cpu.hxx" // Reference implementation using namespace gunrock; using namespace memory; @@ -22,8 +23,12 @@ void test_bfs(int num_arguments, char** argument_array) { std::string filename = argument_array[1]; io::matrix_market_t mm; - format::csr_t csr; + + using csr_t = + format::csr_t; + csr_t csr; csr.from_coo(mm.load(filename)); + thrust::device_vector row_indices(csr.number_of_nonzeros); thrust::device_vector column_offsets(csr.number_of_columns + 1); @@ -54,17 +59,42 @@ void test_bfs(int num_arguments, char** argument_array) { // -- // Run problem - float elapsed = gunrock::bfs::run(G, single_source, distances.data().get(), + float gpu_elapsed = gunrock::bfs::run(G, single_source, distances.data().get(), predecessors.data().get()); + // -- + // CPU Run + + thrust::host_vector h_distances(n_vertices); + thrust::host_vector h_predecessors(n_vertices); + + float cpu_elapsed = bfs_cpu::run( + csr, single_source, h_distances.data(), h_predecessors.data()); + + int n_errors = bfs_cpu::compute_error(distances, h_distances); + // -- // Log - std::cout << "Distances (output) = "; - thrust::copy(distances.begin(), distances.end(), + std::cout << "GPU Distances (output) = "; + thrust::copy(distances.begin(), + (distances.size() < 40) ? distances.begin() + distances.size() + : distances.begin() + 40, std::ostream_iterator(std::cout, " ")); std::cout << std::endl; - std::cout << "BFS Elapsed Time: " << elapsed << " (ms)" << std::endl; + + std::cout << "CPU Distances (output) = "; + thrust::copy(h_distances.begin(), + (h_distances.size() < 40) + ? h_distances.begin() + h_distances.size() + : h_distances.begin() + 40, + std::ostream_iterator(std::cout, " ")); + std::cout << std::endl; + + std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl; + std::cout << "CPU Elapsed Time : " << cpu_elapsed << " (ms)" << std::endl; + std::cout << "Number of errors : " << n_errors << std::endl; + } int main(int argc, char** argv) { diff --git a/examples/bfs/bfs_cpu.hxx b/examples/bfs/bfs_cpu.hxx new file mode 100644 index 00000000..40a0aef5 --- /dev/null +++ b/examples/bfs/bfs_cpu.hxx @@ -0,0 +1,80 @@ +#pragma once + +#include +#include +#include + +namespace bfs_cpu { + +using namespace std; +using namespace std::chrono; + +template +class prioritize { + public: + bool operator()(pair& p1, pair& p2) { + return p1.second > p2.second; + } +}; + +template +float run(csr_t& csr, + vertex_t& single_source, + vertex_t* distances, + vertex_t* predecessors) { + thrust::host_vector row_offsets(csr.row_offsets); // Copy data to CPU + thrust::host_vector column_indices(csr.column_indices); + + for (vertex_t i = 0; i < csr.number_of_rows; i++) + distances[i] = std::numeric_limits::max(); + + auto t_start = high_resolution_clock::now(); + + distances[single_source] = 0; + + priority_queue, + std::vector>, + prioritize> pq; + + pq.push(make_pair(single_source, 0.0)); + + while (!pq.empty()) { + pair curr = pq.top(); + pq.pop(); + + vertex_t curr_node = curr.first; + vertex_t curr_dist = curr.second; + + vertex_t start = row_offsets[curr_node]; + vertex_t end = row_offsets[curr_node + 1]; + + for (vertex_t offset = start; offset < end; offset++) { + vertex_t neib = column_indices[offset]; + vertex_t new_dist = curr_dist + 1; + if (new_dist < distances[neib]) { + distances[neib] = new_dist; + pq.push(make_pair(neib, new_dist)); + } + } + } + + auto t_stop = high_resolution_clock::now(); + auto elapsed = duration_cast(t_stop - t_start).count(); + return (float)elapsed / 1000; +} + +template +int compute_error(thrust::device_vector _gpu_result, + thrust::host_vector cpu_result) { + thrust::host_vector gpu_result(_gpu_result); + + int n_errors = 0; + for (int i = 0; i < cpu_result.size(); i++) { + if (gpu_result[i] != cpu_result[i]) { + n_errors++; + } + } + return n_errors; +} + +} // namespace bfs_cpu \ No newline at end of file diff --git a/examples/color/color.cu b/examples/color/color.cu index ad6dcd35..ab938096 100644 --- a/examples/color/color.cu +++ b/examples/color/color.cu @@ -48,16 +48,15 @@ void test_color(int num_arguments, char** argument_array) { // -- // Run problem - float elapsed = gunrock::color::run(G, colors.data().get()); + float gpu_elapsed = gunrock::color::run(G, colors.data().get()); // -- // Log - std::cout << "Colors (output) = "; - thrust::copy(colors.begin(), colors.end(), - std::ostream_iterator(std::cout, " ")); + std::cout << "colors[:40] = "; + gunrock::print::head(colors, 40); std::cout << std::endl; - std::cout << "color Elapsed Time: " << elapsed << " (ms)" << std::endl; + std::cout << "GPU Elapsed Time: " << gpu_elapsed << " (ms)" << std::endl; } int main(int argc, char** argv) { diff --git a/examples/mtx2bin/CMakeLists.txt b/examples/mtx2bin/CMakeLists.txt new file mode 100644 index 00000000..fdc08cd4 --- /dev/null +++ b/examples/mtx2bin/CMakeLists.txt @@ -0,0 +1,21 @@ +# begin /* Set the application name. */ +set(APPLICATION_NAME mtx2bin) +# end /* Set the application name. */ + +# begin /* Add CUDA executables */ +add_executable(${APPLICATION_NAME}) + +set(SOURCE_LIST + ${APPLICATION_NAME}.cu +) + +target_sources(${APPLICATION_NAME} PRIVATE ${SOURCE_LIST}) +target_link_libraries(${APPLICATION_NAME} PRIVATE essentials) +get_target_property(ESSENTIALS_ARCHITECTURES essentials CUDA_ARCHITECTURES) +set_target_properties(${APPLICATION_NAME} + PROPERTIES + CUDA_ARCHITECTURES ${ESSENTIALS_ARCHITECTURES} +) # XXX: Find a better way to inherit essentials properties. + +message("-- Example Added: ${APPLICATION_NAME}") +# end /* Add CUDA executables */ \ No newline at end of file diff --git a/examples/mtx2bin/mtx2bin.cu b/examples/mtx2bin/mtx2bin.cu new file mode 100644 index 00000000..3f3b0dd3 --- /dev/null +++ b/examples/mtx2bin/mtx2bin.cu @@ -0,0 +1,44 @@ +#include // EXIT_SUCCESS + +#include + +using namespace gunrock; +using namespace memory; + +void test_mtx2bin(int num_arguments, char** argument_array) { + if (num_arguments != 2) { + std::cerr << "usage: ./bin/mtx2bin " << std::endl; + exit(1); + } + + // -- + // Define types + + using vertex_t = int; + using edge_t = int; + using weight_t = float; + + // -- + // IO + + std::string inpath = argument_array[1]; + std::string outpath = inpath + ".csr"; + + io::matrix_market_t mm; + + using csr_t = format::csr_t; + csr_t csr; + csr.from_coo(mm.load(inpath)); + + std::cout << "csr.number_of_rows = " << csr.number_of_rows << std::endl; + std::cout << "csr.number_of_columns = " << csr.number_of_columns << std::endl; + std::cout << "csr.number_of_nonzeros = " << csr.number_of_nonzeros << std::endl; + std::cout << "writing to = " << outpath << std::endl; + + csr.write_binary(outpath); +} + +int main(int argc, char** argv) { + test_mtx2bin(argc, argv); + return EXIT_SUCCESS; +} \ No newline at end of file diff --git a/examples/pr/pr.cu b/examples/pr/pr.cu index 32569a47..b2d63bc1 100644 --- a/examples/pr/pr.cu +++ b/examples/pr/pr.cu @@ -16,18 +16,24 @@ void test_sssp(int num_arguments, char** argument_array) { using edge_t = int; using weight_t = float; + using csr_t = format::csr_t; + csr_t csr; + // -- // IO std::string filename = argument_array[1]; - - io::matrix_market_t mm; - - using csr_t = - format::csr_t; - csr_t csr; - csr.from_coo(mm.load(filename)); - + + if(util::is_market(filename)) { + io::matrix_market_t mm; + csr.from_coo(mm.load(filename)); + } else if(util::is_binary_csr(filename)) { + csr.read_binary(filename); + } else { + std::cerr << "Unknown file format: " << filename << std::endl; + exit(1); + } + // -- // Build graph @@ -54,19 +60,13 @@ void test_sssp(int num_arguments, char** argument_array) { // -- // GPU Run - std::cout << "gunrock::pr::run -- starting" << std::endl; float gpu_elapsed = gunrock::pr::run(G, alpha, tol, p.data().get()); - std::cout << "gunrock::pr::run -- complete" << std::endl; // -- // Log + Validate - std::cout << "GPU p (output) = "; - thrust::copy(p.begin(), - (p.size() < 40) ? p.begin() + p.size() - : p.begin() + 40, - std::ostream_iterator(std::cout, " ")); - std::cout << std::endl; + std::cout << "GPU p[:40] = "; + gunrock::print::head(p, 40); std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl; } diff --git a/examples/sssp/sssp.cu b/examples/sssp/sssp.cu index 9961a6dd..728a4503 100644 --- a/examples/sssp/sssp.cu +++ b/examples/sssp/sssp.cu @@ -71,20 +71,11 @@ void test_sssp(int num_arguments, char** argument_array) { // -- // Log + Validate - std::cout << "GPU Distances (output) = "; - thrust::copy(distances.begin(), - (distances.size() < 40) ? distances.begin() + distances.size() - : distances.begin() + 40, - std::ostream_iterator(std::cout, " ")); - std::cout << std::endl; + std::cout << "GPU distances[:40] = "; + gunrock::print::head(distances, 40); std::cout << "CPU Distances (output) = "; - thrust::copy(h_distances.begin(), - (h_distances.size() < 40) - ? h_distances.begin() + h_distances.size() - : h_distances.begin() + 40, - std::ostream_iterator(std::cout, " ")); - std::cout << std::endl; + gunrock::print::head(h_distances, 40); std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl; std::cout << "CPU Elapsed Time : " << cpu_elapsed << " (ms)" << std::endl; diff --git a/include/gunrock/applications/application.hxx b/include/gunrock/applications/application.hxx index e04a1d46..07570b7b 100644 --- a/include/gunrock/applications/application.hxx +++ b/include/gunrock/applications/application.hxx @@ -19,6 +19,7 @@ // Utility includes #include +#include // Format includes #include diff --git a/include/gunrock/formats/csr.hxx b/include/gunrock/formats/csr.hxx index b3c0c7d1..03d20da5 100644 --- a/include/gunrock/formats/csr.hxx +++ b/include/gunrock/formats/csr.hxx @@ -144,6 +144,73 @@ struct csr_t { return *this; // CSR representation (with possible duplicates) } + void read_binary(std::string filename) { + FILE* file = fopen(filename.c_str(), "rb"); + + // Read metadata + fread(&number_of_rows, sizeof(index_t), 1, file); + fread(&number_of_columns, sizeof(index_t), 1, file); + fread(&number_of_nonzeros, sizeof(offset_t), 1, file); + + row_offsets.resize(number_of_rows + 1); + column_indices.resize(number_of_nonzeros); + nonzero_values.resize(number_of_nonzeros); + + if(space == memory_space_t::device) { + assert(space == memory_space_t::device); + + thrust::host_vector h_row_offsets(number_of_rows + 1); + thrust::host_vector h_column_indices(number_of_nonzeros); + thrust::host_vector h_nonzero_values(number_of_nonzeros); + + fread(memory::raw_pointer_cast(h_row_offsets.data()), sizeof(offset_t), number_of_rows + 1, file); + fread(memory::raw_pointer_cast(h_column_indices.data()), sizeof(index_t), number_of_nonzeros, file); + fread(memory::raw_pointer_cast(h_nonzero_values.data()), sizeof(value_t), number_of_nonzeros, file); + + // Copy data from host to device + row_offsets = h_row_offsets; + column_indices = h_column_indices; + nonzero_values = h_nonzero_values; + + } else { + assert(space == memory_space_t::host); + + fread(memory::raw_pointer_cast(row_offsets.data()), sizeof(offset_t), number_of_rows + 1, file); + fread(memory::raw_pointer_cast(column_indices.data()), sizeof(index_t), number_of_nonzeros, file); + fread(memory::raw_pointer_cast(nonzero_values.data()), sizeof(value_t), number_of_nonzeros, file); + } + } + + void write_binary(std::string filename) { + FILE* file = fopen(filename.c_str(), "wb"); + + // Write metadata + fwrite(&number_of_rows, sizeof(index_t), 1, file); + fwrite(&number_of_columns, sizeof(index_t), 1, file); + fwrite(&number_of_nonzeros, sizeof(offset_t), 1, file); + + // Write data + if(space == memory_space_t::device) { + assert(space == memory_space_t::device); + + thrust::host_vector h_row_offsets(row_offsets); + thrust::host_vector h_column_indices(column_indices); + thrust::host_vector h_nonzero_values(nonzero_values); + + fwrite(memory::raw_pointer_cast(h_row_offsets.data()), sizeof(offset_t), number_of_rows + 1, file); + fwrite(memory::raw_pointer_cast(h_column_indices.data()), sizeof(index_t), number_of_nonzeros, file); + fwrite(memory::raw_pointer_cast(h_nonzero_values.data()), sizeof(value_t), number_of_nonzeros, file); + } else { + assert(space == memory_space_t::host); + + fwrite(memory::raw_pointer_cast(row_offsets.data()), sizeof(offset_t), number_of_rows + 1, file); + fwrite(memory::raw_pointer_cast(column_indices.data()), sizeof(index_t), number_of_nonzeros, file); + fwrite(memory::raw_pointer_cast(nonzero_values.data()), sizeof(value_t), number_of_nonzeros, file); + } + + fclose(file); + } + }; // struct csr_t } // namespace format diff --git a/include/gunrock/util/filepath.hxx b/include/gunrock/util/filepath.hxx index 294a955e..1e1d4e71 100644 --- a/include/gunrock/util/filepath.hxx +++ b/include/gunrock/util/filepath.hxx @@ -15,5 +15,16 @@ std::string extract_dataset(std::string filename) { return filename.substr(0, lastindex); } +bool is_market(std::string filename) { + return ( + (filename.substr(filename.size() - 4) == ".mtx" ) || + (filename.substr(filename.size() - 5) == ".mmio") + ); +} + +bool is_binary_csr(std::string filename) { + return filename.substr(filename.size() - 4) == ".csr"; +} + } // namespace util } // namespace gunrock \ No newline at end of file diff --git a/include/gunrock/util/print.hxx b/include/gunrock/util/print.hxx new file mode 100644 index 00000000..6cc79c92 --- /dev/null +++ b/include/gunrock/util/print.hxx @@ -0,0 +1,30 @@ +/** + * @file print.hxx + * + * @brief + */ + +#pragma once + +namespace gunrock { + +/** + * @namespace print + * Print utilities. + */ +namespace print { + +// Print the first k elements of a `thrust` vector +template +void head(vec_t& x, int k) { + thrust::copy( + x.begin(), + (x.size() < k) ? x.begin() + x.size() : x.begin() + k, + std::ostream_iterator(std::cout, " ") + ); + std::cout << std::endl; +} + + +} // namespace print +} // namespace gunrock \ No newline at end of file