graph
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
gpu_csr-impl.hpp
Go to the documentation of this file.
1 #pragma once
2 
3 #include <set>
4 
5 #include "gpu_csr.hpp"
6 #include "pangolin/logger.hpp"
7 
8 #ifdef __CUDACC__
9 #define PANGOLIN_CUDA_MEMBER __host__ __device__
10 #else
11 #define PANGOLIN_CUDA_MEMBER
12 #endif
13 
15 
16 template<typename Index>
17 GPUCSR<Index>::GPUCSR() : maxCol_(0) {
18 }
19 
20 template<typename Index>
22  if (rowOffset_.size() == 0) {
23  return 0;
24  } else {
25  return rowOffset_.size() - 1;
26  }
27 }
28 
29 template<typename Index>
30 uint64_t GPUCSR<Index>::num_nodes() const {
31  std::set<Index> nodes;
32  // add all dsts
33  for (Index ci = 0; ci < col_.size(); ++ci) {
34  nodes.insert(col_[ci]);
35  }
36  // add non-zero sources
37  for (Index i = 0; i < rowOffset_.size() - 1; ++i) {
38  Index row_start = rowOffset_[i];
39  Index row_end = rowOffset_[i+1];
40  if (row_start != row_end) {
41  nodes.insert(i);
42  }
43  }
44  return nodes.size();
45 }
46 
47 template<typename Index>
48 GPUCSR<Index> GPUCSR<Index>::from_edgelist(const EdgeList &es, bool (*edgeFilter)(const Edge &)) {
49  GPUCSR<Index> csr;
50 
51 
52  if (es.size() == 0) {
53  LOG(warn, "constructing from empty edge list");
54  return csr;
55  }
56 
57 
58  for (const auto &edge : es) {
59 
60  const Index src = static_cast<Index>(edge.first);
61  const Index dst = static_cast<Index>(edge.second);
62 
63  // edge has a new src and should be in a new row
64  // even if the edge is filtered out, we need to add empty rows
65  while (csr.rowOffset_.size() != size_t(src + 1))
66  {
67  // expecting inputs to be sorted by src, so it should be at least
68  // as big as the current largest row we have recored
69  assert(src >= csr.rowOffset_.size());
70  // SPDLOG_TRACE(logger::console, "node {} edges start at {}", edge.src_, csr.edgeSrc_.size());
71  csr.rowOffset_.push_back(csr.col_.size());
72  }
73 
74  // filter or add the edge
75  if (nullptr != edgeFilter && edgeFilter(edge)) {
76  continue;
77  } else {
78  csr.col_.push_back(dst);
79  csr.maxCol_ = std::max(dst, csr.maxCol_);
80  }
81  }
82 
83  // add the final length of the non-zeros to the offset array
84  csr.rowOffset_.push_back(csr.col_.size());
85 
86  return csr;
87 }
88 
89 template<typename Index>
91  GPUCSRView<Index> view;
92  view.nnz_ = nnz();
93  view.num_rows_ = num_rows();
94  view.rowOffset_ = rowOffset_.data();
95  view.col_ = col_.data();
96  return view;
97 }
98 
99 
101 
102 #undef PANGOLIN_CUDA_MEMBER
A CSR matrix backed by CUDA Unified Memory.
Definition: gpu_csr.hpp:16
static GPUCSR< Index > from_edgelist(const EdgeList &es, bool(*edgeFilter)(const Edge &)=nullptr)
Definition: gpu_csr-impl.hpp:48
a read-only view of a GPUCSRm suitable for passing to a GPU kernel by value.
Definition: gpu_csr.hpp:23
Vector< Index > col_
non-zero column indices
Definition: gpu_csr.hpp:55
PANGOLIN_CUDA_MEMBER uint64_t num_rows() const
number of matrix rows
Definition: gpu_csr-impl.hpp:21
PANGOLIN_BEGIN_NAMESPACE()
std::vector< Edge > EdgeList
Definition: edge_list.hpp:9
#define LOG(level,...)
Definition: logger.hpp:25
Vector< Index > rowOffset_
offset in col_ that each row starts at
Definition: gpu_csr.hpp:54
GPUCSRView< Index > view() const
create a GPUCSRView for this GPUCSR
Definition: gpu_csr-impl.hpp:90
Index maxCol_
Definition: gpu_csr.hpp:51
uint64_t nnz_
number of non-zeros
Definition: gpu_csr.hpp:26
uint64_t num_rows_
length of rowOffset - 1
Definition: gpu_csr.hpp:27
const Index * rowOffset_
offset in col_ that each row starts at
Definition: gpu_csr.hpp:31
const Index * col_
non-zero column indices
Definition: gpu_csr.hpp:32
uint64_t num_nodes() const
number of unique row/col indices
Definition: gpu_csr-impl.hpp:30
#define PANGOLIN_CUDA_MEMBER
Definition: gpu_csr-impl.hpp:11
#define PANGOLIN_END_NAMESPACE()
std::pair< Uint, Uint > Edge
Definition: edge.hpp:14