Format improvements for CSR, CSC, COO.

benchmarks/spmv/work_oriented.cu

@@ -58,5 +58,5 @@ void work_oriented_bench(nvbench::state& state, nvbench::type_list<value_t>) {
 }
 
 // Define a type_list to use for the type axis:
-using value_types = nvbench::type_list<int, float>;
+using value_types = nvbench::type_list<int, float, double>;
 NVBENCH_BENCH_TYPES(work_oriented_bench, NVBENCH_TYPE_AXES(value_types));

include/loops/container/coo.hxx

@@ -1,12 +1,26 @@
+/**
+ * @file coo.hxx
+ * @author Muhammad Osama (mosama@ucdavis.edu)
+ * @brief Interface for Coordinate format.
+ * @version 0.1
+ * @date 2022-07-21
+ *
+ * @copyright Copyright (c) 2022
+ *
+ */
+
 #pragma once
 
+#include <loops/container/formats.hxx>
+#include <loops/container/detail/convert.hxx>
 #include <loops/container/vector.hxx>
 #include <loops/memory.hxx>
 
-#include <thrust/execution_policy.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/sort.h>
 #include <thrust/tuple.h>
+#include <thrust/unique.h>
+#include <thrust/execution_policy.h>
+#include <thrust/iterator/zip_iterator.h>
 
 namespace loops {
 
@@ -30,8 +44,19 @@ struct coo_t {
   vector_t<index_t, space> col_indices;  /// J
   vector_t<value_t, space> values;       /// V
 
+  /**
+   * @brief Construct a new coo object with everything initialized to zero.
+   *
+   */
   coo_t() : rows(0), cols(0), nnzs(0), row_indices(), col_indices(), values() {}
 
+  /**
+   * @brief Construct a new coo object with the given dimensions.
+   *
+   * @param r Number of rows.
+   * @param c Number of columns.
+   * @param nnz Number of non-zero elements.
+   */
   coo_t(std::size_t r, std::size_t c, std::size_t nnz)
       : rows(r),
         cols(c),
@@ -40,6 +65,11 @@ struct coo_t {
         col_indices(nnz),
         values(nnz) {}
 
+  /**
+   * @brief Construct a new coo from another coo object on host/device.
+   *
+   * @param rhs coo_t<index_t, value_t, rhs_space>
+   */
   template <auto rhs_space>
   coo_t(const coo_t<index_t, value_t, rhs_space>& rhs)
       : rows(rhs.rows),
@@ -49,6 +79,28 @@ struct coo_t {
         col_indices(rhs.col_indices),
         values(rhs.values) {}
 
+  /**
+   * @brief Construct a new coo object from compressed sparse format (CSR).
+   *
+   * @param csr csr_t<index_t, offset_t, value_t, rhs_space>
+   */
+  template <auto rhs_space, typename offset_t>
+  coo_t(const csr_t<index_t, offset_t, value_t, rhs_space>& csr)
+      : rows(csr.rows),
+        cols(csr.cols),
+        nnzs(csr.nnzs),
+        row_indices(csr.nnzs),
+        col_indices(csr.col_indices),
+        values(csr.values) {
+    /// TODO: Do not need this copy for all cases.
+    vector_t<offset_t, space> _row_offsets = csr.offsets;
+    detail::offsets_to_indices(_row_offsets, row_indices);
+  }
+
+  /**
+   * @brief Sorts the coordinate matrix by row indices.
+   *
+   */
   void sort_by_row() {
     auto begin = thrust::make_zip_iterator(
         thrust::make_tuple(row_indices.begin(), col_indices.begin()));
@@ -57,6 +109,10 @@ struct coo_t {
     sort(begin, end);
   }
 
+  /**
+   * @brief Sorts the coordinate matrix by column indices.
+   *
+   */
   void sort_by_column() {
     auto begin = thrust::make_zip_iterator(
         thrust::make_tuple(col_indices.begin(), row_indices.begin()));
@@ -65,7 +121,44 @@ struct coo_t {
     sort(begin, end);
   }
 
+  /**
+   * @brief Sorts, removes I,J pairs.
+   *
+   */
+  void remove_duplicates() {
+    // Sort by row indices.
+    this->sort_by_row();
+    auto begin = thrust::make_zip_iterator(
+        thrust::make_tuple(row_indices.begin(), col_indices.begin()));
+    auto end = thrust::make_zip_iterator(
+        thrust::make_tuple(row_indices.end(), col_indices.end()));
+
+    // Remove duplicates.
+    auto new_it = thrust::unique_by_key(begin, end, values.begin());
+    auto first_it = thrust::get<1>(new_it);
+    nnzs = thrust::distance(values.begin(), first_it);
+
+    // Resize vectors to new size.
+    row_indices.resize(nnzs);
+    col_indices.resize(nnzs);
+    values.resize(nnzs);
+  }
+
  private:
+  /**
+   * @brief Sorting helper, uses zip iterator as begin and end.
+   *
+   * @par Example Zip iterator:
+   *  auto begin = thrust::make_zip_iterator(
+   *    thrust::make_tuple(col_indices.begin(), row_indices.begin()));
+   *  auto end = thrust::make_zip_iterator(
+   *    thrust::make_tuple(col_indices.end(), row_indices.end()));
+   *
+   * @tparam begin_it_t Begin iterator type.
+   * @tparam end_it_t End iterator type.
+   * @param begin Begin iterator (zip iterator of row and col indices).
+   * @param end End iterator (zip iterator of row and col indices).
+   */
   template <typename begin_it_t, typename end_it_t>
   void sort(begin_it_t& begin, end_it_t& end) {
     thrust::sort_by_key(begin, end, values.begin());

include/loops/container/csc.hxx

@@ -1,3 +1,14 @@
+/**
+ * @file csc.hxx
+ * @author Muhammad Osama (mosama@ucdavis.edu)
+ * @brief Interface for Compressed Sparse-Column format.
+ * @version 0.1
+ * @date 2022-07-21
+ *
+ * @copyright Copyright (c) 2022
+ *
+ */
+
 #pragma once
 
 #include <loops/container/vector.hxx>
@@ -27,17 +38,56 @@ struct csc_t {
   vector_t<index_t, space> indices;   /// Ap
   vector_t<value_t, space> values;    /// Ax
 
+  /**
+   * @brief Construct a new csc object with everything initialized to zero.
+   *
+   */
   csc_t() : rows(0), cols(0), nnzs(0), offsets(), indices(), values() {}
 
+  /**
+   * @brief Construct a new csc object with the given dimensions.
+   *
+   * @param r Number of rows.
+   * @param c Number of columns.
+   * @param nnz Number of non-zero elements.
+   */
   csc_t(std::size_t r, std::size_t c, std::size_t nnz)
       : rows(r),
         cols(c),
         nnzs(nnz),
-        offsets(c + 1),
+        offsets(r + 1),
         indices(nnz),
         values(nnz) {}
 
-  ~csc_t() {}
+  /**
+   * @brief Construct a new csc from another csc object on host/device.
+   *
+   * @param rhs csc_t<index_t, offset_t, value_t, rhs_space>
+   */
+  template <auto rhs_space>
+  csc_t(const csc_t<index_t, offset_t, value_t, rhs_space>& rhs)
+      : rows(rhs.rows),
+        cols(rhs.cols),
+        nnzs(rhs.nnzs),
+        offsets(rhs.offsets),
+        indices(rhs.indices),
+        values(rhs.values) {}
+
+  /**
+   * @brief Construct a new csc object from coordinate format (COO).
+   * @note This constructor creates a copy of the input COO matrix.
+   *
+   * @param coo coo_t<index_t, value_t, auto>
+   */
+  template <auto rhs_space>
+  csc_t(const coo_t<index_t, value_t, rhs_space>& coo)
+      : rows(coo.rows), cols(coo.cols), nnzs(coo.nnzs), offsets(coo.cols + 1) {
+    coo_t<index_t, value_t, space> __(coo);
+    __.sort_by_column();
+    indices = std::move(__.row_indices);
+    values = std::move(__.values);
+    detail::indices_to_offsets(__.col_indices, offsets);
+  }
 
 };  // struct csc_t
 

include/loops/container/csr.hxx

@@ -1,3 +1,14 @@
+/**
+ * @file csr.hxx
+ * @author Muhammad Osama (mosama@ucdavis.edu)
+ * @brief Interface for Compressed Sparse-Row format.
+ * @version 0.1
+ * @date 2022-07-21
+ *
+ * @copyright Copyright (c) 2022
+ *
+ */
+
 #pragma once
 
 #include <loops/container/vector.hxx>
@@ -44,7 +55,7 @@ struct csr_t {
    * @param c Number of columns.
    * @param nnz Number of non-zero elements.
    */
-  csr_t(const std::size_t& r, const std::size_t& c, const std::size_t& nnz)
+  csr_t(std::size_t r, std::size_t c, std::size_t nnz)
       : rows(r),
         cols(c),
         nnzs(nnz),
@@ -53,7 +64,7 @@ struct csr_t {
         values(nnz) {}
 
   /**
-   * @brief Construct a new csr from another csr object on host.
+   * @brief Construct a new csr from another csr object on host/device.
    *
    * @param rhs csr_t<index_t, offset_t, value_t, rhs_space>
    */
@@ -68,20 +79,18 @@ struct csr_t {
 
   /**
    * @brief Construct a new csr object from coordinate format (COO).
+   * @note This constructor creates a copy of the input COO matrix.
    *
    * @param coo coo_t<index_t, value_t, auto>
    */
   template <auto rhs_space>
   csr_t(const coo_t<index_t, value_t, rhs_space>& coo)
-      : rows(coo.rows),
-        cols(coo.cols),
-        nnzs(coo.nnzs),
-        offsets(coo.rows + 1),
-        indices(coo.col_indices),
-        values(coo.values) {
-    /// TODO: Do not need this copy for all cases.
-    vector_t<index_t, space> _row_indices = coo.row_indices;
-    detail::indices_to_offsets(_row_indices, offsets);
+      : rows(coo.rows), cols(coo.cols), nnzs(coo.nnzs), offsets(coo.rows + 1) {
+    coo_t<index_t, value_t, space> __(coo);
+    __.sort_by_row();
+    indices = std::move(__.col_indices);
+    values = std::move(__.values);
+    detail::indices_to_offsets(__.row_indices, offsets);
   }
 };  // struct csr_t
 

include/loops/util/generate.hxx

@@ -13,12 +13,9 @@
 #include <chrono>
 #include <thrust/random.h>
 #include <thrust/distance.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 #include <thrust/execution_policy.h>
-#include <thrust/sort.h>
-#include <thrust/unique.h>
+#include <thrust/iterator/counting_iterator.h>
 
 #include <loops/memory.hxx>
 #include <loops/container/formats.hxx>
@@ -110,26 +107,8 @@ void csr(std::size_t rows,
   uniform_distribution(coo.values.begin(), coo.values.end(), value_t(0.0),
                        value_t(1.0));
 
-  // Sort COO matrix.
-  coo.sort_by_row();
-
-  // Zip Iterators.
-  auto begin = thrust::make_zip_iterator(
-      thrust::make_tuple(coo.row_indices.begin(), coo.col_indices.begin()));
-  auto end = thrust::make_zip_iterator(
-      thrust::make_tuple(coo.row_indices.end(), coo.col_indices.end()));
-
   // Remove duplicates.
-  auto new_it = thrust::unique_by_key(begin, end, coo.values.begin());
-  auto first_it = thrust::get<1>(new_it);
-  auto new_distance = thrust::distance(coo.values.begin(), first_it);
-
-  // Resize the COO matrix.
-  coo.row_indices.resize(new_distance);
-  coo.col_indices.resize(new_distance);
-  coo.values.resize(new_distance);
-  coo.nnzs = new_distance;
-
+  coo.remove_duplicates();
   matrix = coo;
 }