Kokkos Core Kernels Package  Version of the Day
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups Pages
Kokkos_Crs.hpp
1 /*
2 //@HEADER
3 // ************************************************************************
4 //
5 // Kokkos v. 3.0
6 // Copyright (2020) National Technology & Engineering
7 // Solutions of Sandia, LLC (NTESS).
8 //
9 // Under the terms of Contract DE-NA0003525 with NTESS,
10 // the U.S. Government retains certain rights in this software.
11 //
12 // Redistribution and use in source and binary forms, with or without
13 // modification, are permitted provided that the following conditions are
14 // met:
15 //
16 // 1. Redistributions of source code must retain the above copyright
17 // notice, this list of conditions and the following disclaimer.
18 //
19 // 2. Redistributions in binary form must reproduce the above copyright
20 // notice, this list of conditions and the following disclaimer in the
21 // documentation and/or other materials provided with the distribution.
22 //
23 // 3. Neither the name of the Corporation nor the names of the
24 // contributors may be used to endorse or promote products derived from
25 // this software without specific prior written permission.
26 //
27 // THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
28 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
29 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
30 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
31 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
32 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
33 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
34 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
35 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
36 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
37 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
38 //
39 // Questions? Contact Christian R. Trott (crtrott@sandia.gov)
40 //
41 // ************************************************************************
42 //@HEADER
43 */
44 
45 #ifndef KOKKOS_CRS_HPP
46 #define KOKKOS_CRS_HPP
47 
48 namespace Kokkos {
49 
80 template <class DataType, class Arg1Type, class Arg2Type = void,
81  typename SizeType = typename ViewTraits<DataType*, Arg1Type, Arg2Type,
82  void>::size_type>
83 class Crs {
84  protected:
86 
87  public:
88  typedef DataType data_type;
89  typedef typename traits::array_layout array_layout;
90  typedef typename traits::execution_space execution_space;
91  typedef typename traits::memory_space memory_space;
92  typedef typename traits::device_type device_type;
93  typedef SizeType size_type;
94 
96  typedef Crs<DataType, array_layout, typename traits::host_mirror_space,
97  SizeType>
98  HostMirror;
101 
102  row_map_type row_map;
103  entries_type entries;
104 
105  /*
106  * Default Constructors, operators and destructor
107  */
108  KOKKOS_DEFAULTED_FUNCTION Crs() = default;
109  KOKKOS_DEFAULTED_FUNCTION Crs(Crs const&) = default;
110  KOKKOS_DEFAULTED_FUNCTION Crs(Crs&&) = default;
111  KOKKOS_DEFAULTED_FUNCTION Crs& operator=(Crs const&) = default;
112  KOKKOS_DEFAULTED_FUNCTION Crs& operator=(Crs&&) = default;
113  KOKKOS_DEFAULTED_FUNCTION ~Crs() = default;
114 
119  template <class EntriesType, class RowMapType>
120  KOKKOS_INLINE_FUNCTION Crs(const RowMapType& row_map_,
121  const EntriesType& entries_)
122  : row_map(row_map_), entries(entries_) {}
123 
126  KOKKOS_INLINE_FUNCTION
127  size_type numRows() const {
128  return (row_map.extent(0) != 0)
129  ? row_map.extent(0) - static_cast<size_type>(1)
130  : static_cast<size_type>(0);
131  }
132 };
133 
134 /*--------------------------------------------------------------------------*/
135 
136 template <class OutCounts, class DataType, class Arg1Type, class Arg2Type,
137  class SizeType>
138 void get_crs_transpose_counts(
139  OutCounts& out, Crs<DataType, Arg1Type, Arg2Type, SizeType> const& in,
140  std::string const& name = "transpose_counts");
141 
142 template <class OutCounts, class InCrs>
143 typename OutCounts::value_type get_crs_row_map_from_counts(
144  OutCounts& out, InCrs const& in, std::string const& name = "row_map");
145 
146 template <class DataType, class Arg1Type, class Arg2Type, class SizeType>
147 void transpose_crs(Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
148  Crs<DataType, Arg1Type, Arg2Type, SizeType> const& in);
149 
150 } // namespace Kokkos
151 
152 /*--------------------------------------------------------------------------*/
153 
154 /*--------------------------------------------------------------------------*/
155 
156 namespace Kokkos {
157 namespace Impl {
158 
159 template <class InCrs, class OutCounts>
160 class GetCrsTransposeCounts {
161  public:
162  using execution_space = typename InCrs::execution_space;
163  using self_type = GetCrsTransposeCounts<InCrs, OutCounts>;
164  using index_type = typename InCrs::size_type;
165 
166  private:
167  InCrs in;
168  OutCounts out;
169 
170  public:
171  KOKKOS_INLINE_FUNCTION
172  void operator()(index_type i) const { atomic_increment(&out[in.entries(i)]); }
173  GetCrsTransposeCounts(InCrs const& arg_in, OutCounts const& arg_out)
174  : in(arg_in), out(arg_out) {
175  using policy_type = RangePolicy<index_type, execution_space>;
177  const closure_type closure(*this,
178  policy_type(0, index_type(in.entries.size())));
179  closure.execute();
180  execution_space().fence();
181  }
182 };
183 
184 template <class InCounts, class OutRowMap>
185 class CrsRowMapFromCounts {
186  public:
187  using execution_space = typename InCounts::execution_space;
188  using value_type = typename OutRowMap::value_type;
189  using index_type = typename InCounts::size_type;
190  using last_value_type = Kokkos::View<value_type, execution_space>;
191 
192  private:
193  InCounts m_in;
194  OutRowMap m_out;
195  last_value_type m_last_value;
196 
197  public:
198  KOKKOS_INLINE_FUNCTION
199  void operator()(index_type i, value_type& update, bool final_pass) const {
200  if (i < m_in.size()) {
201  update += m_in(i);
202  if (final_pass) m_out(i + 1) = update;
203  } else if (final_pass) {
204  m_out(0) = 0;
205  m_last_value() = update;
206  }
207  }
208  KOKKOS_INLINE_FUNCTION
209  void init(value_type& update) const { update = 0; }
210  KOKKOS_INLINE_FUNCTION
211  void join(volatile value_type& update,
212  const volatile value_type& input) const {
213  update += input;
214  }
215  using self_type = CrsRowMapFromCounts<InCounts, OutRowMap>;
216  CrsRowMapFromCounts(InCounts const& arg_in, OutRowMap const& arg_out)
217  : m_in(arg_in), m_out(arg_out), m_last_value("last_value") {}
218  value_type execute() {
219  using policy_type = RangePolicy<index_type, execution_space>;
221  closure_type closure(*this, policy_type(0, m_in.size() + 1));
222  closure.execute();
223  auto last_value = Kokkos::create_mirror_view(m_last_value);
224  Kokkos::deep_copy(last_value, m_last_value);
225  return last_value();
226  }
227 };
228 
229 template <class InCrs, class OutCrs>
230 class FillCrsTransposeEntries {
231  public:
232  using execution_space = typename InCrs::execution_space;
233  using memory_space = typename InCrs::memory_space;
234  using value_type = typename OutCrs::entries_type::value_type;
235  using index_type = typename InCrs::size_type;
236 
237  private:
238  using counters_type = View<index_type*, memory_space>;
239  InCrs in;
240  OutCrs out;
241  counters_type counters;
242 
243  public:
244  KOKKOS_INLINE_FUNCTION
245  void operator()(index_type i) const {
246  auto begin = in.row_map(i);
247  auto end = in.row_map(i + 1);
248  for (auto j = begin; j < end; ++j) {
249  auto ti = in.entries(j);
250  auto tbegin = out.row_map(ti);
251  auto tj = atomic_fetch_add(&counters(ti), 1);
252  out.entries(tbegin + tj) = i;
253  }
254  }
255  using self_type = FillCrsTransposeEntries<InCrs, OutCrs>;
256  FillCrsTransposeEntries(InCrs const& arg_in, OutCrs const& arg_out)
257  : in(arg_in), out(arg_out), counters("counters", arg_out.numRows()) {
258  using policy_type = RangePolicy<index_type, execution_space>;
260  const closure_type closure(*this, policy_type(0, index_type(in.numRows())));
261  closure.execute();
262  execution_space().fence();
263  }
264 };
265 
266 } // namespace Impl
267 } // namespace Kokkos
268 
269 /*--------------------------------------------------------------------------*/
270 
271 /*--------------------------------------------------------------------------*/
272 
273 namespace Kokkos {
274 
275 template <class OutCounts, class DataType, class Arg1Type, class Arg2Type,
276  class SizeType>
277 void get_crs_transpose_counts(
278  OutCounts& out, Crs<DataType, Arg1Type, Arg2Type, SizeType> const& in,
279  std::string const& name) {
280  using InCrs = Crs<DataType, Arg1Type, Arg2Type, SizeType>;
281  out = OutCounts(name, in.numRows());
282  Kokkos::Impl::GetCrsTransposeCounts<InCrs, OutCounts> functor(in, out);
283 }
284 
285 template <class OutRowMap, class InCounts>
286 typename OutRowMap::value_type get_crs_row_map_from_counts(
287  OutRowMap& out, InCounts const& in, std::string const& name) {
288  out = OutRowMap(ViewAllocateWithoutInitializing(name), in.size() + 1);
289  Kokkos::Impl::CrsRowMapFromCounts<InCounts, OutRowMap> functor(in, out);
290  return functor.execute();
291 }
292 
293 template <class DataType, class Arg1Type, class Arg2Type, class SizeType>
294 void transpose_crs(Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
295  Crs<DataType, Arg1Type, Arg2Type, SizeType> const& in) {
296  typedef Crs<DataType, Arg1Type, Arg2Type, SizeType> crs_type;
297  typedef typename crs_type::memory_space memory_space;
298  typedef View<SizeType*, memory_space> counts_type;
299  {
300  counts_type counts;
301  Kokkos::get_crs_transpose_counts(counts, in);
302  Kokkos::get_crs_row_map_from_counts(out.row_map, counts,
303  "tranpose_row_map");
304  }
305  out.entries = decltype(out.entries)("transpose_entries", in.entries.size());
306  Kokkos::Impl::FillCrsTransposeEntries<crs_type, crs_type> entries_functor(
307  in, out);
308 }
309 
310 template <class CrsType, class Functor,
311  class ExecutionSpace = typename CrsType::execution_space>
312 struct CountAndFillBase;
313 
314 template <class CrsType, class Functor, class ExecutionSpace>
315 struct CountAndFillBase {
316  using data_type = typename CrsType::data_type;
317  using size_type = typename CrsType::size_type;
318  using row_map_type = typename CrsType::row_map_type;
319  using counts_type = row_map_type;
320  CrsType m_crs;
321  Functor m_functor;
322  counts_type m_counts;
323  struct Count {};
324  inline void operator()(Count, size_type i) const {
325  m_counts(i) = m_functor(i, nullptr);
326  }
327  struct Fill {};
328  inline void operator()(Fill, size_type i) const {
329  auto j = m_crs.row_map(i);
330  /* we don't want to access entries(entries.size()), even if its just to get
331  its address and never use it. this can happen when row (i) is empty and
332  all rows after it are also empty. we could compare to row_map(i + 1), but
333  that is a read from global memory, whereas dimension_0() should be part
334  of the View in registers (or constant memory) */
335  data_type* fill = (j == static_cast<decltype(j)>(m_crs.entries.extent(0)))
336  ? nullptr
337  : (&(m_crs.entries(j)));
338  m_functor(i, fill);
339  }
340  CountAndFillBase(CrsType& crs, Functor const& f) : m_crs(crs), m_functor(f) {}
341 };
342 
343 #if defined(KOKKOS_ENABLE_CUDA)
344 template <class CrsType, class Functor>
345 struct CountAndFillBase<CrsType, Functor, Kokkos::Cuda> {
346  using data_type = typename CrsType::data_type;
347  using size_type = typename CrsType::size_type;
348  using row_map_type = typename CrsType::row_map_type;
349  using counts_type = row_map_type;
350  CrsType m_crs;
351  Functor m_functor;
352  counts_type m_counts;
353  struct Count {};
354  __device__ inline void operator()(Count, size_type i) const {
355  m_counts(i) = m_functor(i, nullptr);
356  }
357  struct Fill {};
358  __device__ inline void operator()(Fill, size_type i) const {
359  auto j = m_crs.row_map(i);
360  /* we don't want to access entries(entries.size()), even if its just to get
361  its address and never use it. this can happen when row (i) is empty and
362  all rows after it are also empty. we could compare to row_map(i + 1), but
363  that is a read from global memory, whereas dimension_0() should be part
364  of the View in registers (or constant memory) */
365  data_type* fill = (j == static_cast<decltype(j)>(m_crs.entries.extent(0)))
366  ? nullptr
367  : (&(m_crs.entries(j)));
368  m_functor(i, fill);
369  }
370  CountAndFillBase(CrsType& crs, Functor const& f) : m_crs(crs), m_functor(f) {}
371 };
372 #endif
373 
374 template <class CrsType, class Functor>
375 struct CountAndFill : public CountAndFillBase<CrsType, Functor> {
376  using base_type = CountAndFillBase<CrsType, Functor>;
377  using typename base_type::Count;
378  using typename base_type::counts_type;
379  using typename base_type::data_type;
380  using typename base_type::Fill;
381  using typename base_type::size_type;
382  using entries_type = typename CrsType::entries_type;
383  using self_type = CountAndFill<CrsType, Functor>;
384  CountAndFill(CrsType& crs, size_type nrows, Functor const& f)
385  : base_type(crs, f) {
386  using execution_space = typename CrsType::execution_space;
387  this->m_counts = counts_type("counts", nrows);
388  {
389  using count_policy_type = RangePolicy<size_type, execution_space, Count>;
390  using count_closure_type =
392  const count_closure_type closure(*this, count_policy_type(0, nrows));
393  closure.execute();
394  }
395  auto nentries = Kokkos::get_crs_row_map_from_counts(this->m_crs.row_map,
396  this->m_counts);
397  this->m_counts = counts_type();
398  this->m_crs.entries = entries_type("entries", nentries);
399  {
400  using fill_policy_type = RangePolicy<size_type, execution_space, Fill>;
401  using fill_closure_type =
403  const fill_closure_type closure(*this, fill_policy_type(0, nrows));
404  closure.execute();
405  }
406  crs = this->m_crs;
407  }
408 };
409 
410 template <class CrsType, class Functor>
411 void count_and_fill_crs(CrsType& crs, typename CrsType::size_type nrows,
412  Functor const& f) {
413  Kokkos::CountAndFill<CrsType, Functor>(crs, nrows, f);
414 }
415 
416 } // namespace Kokkos
417 
418 #endif /* #define KOKKOS_CRS_HPP */
void deep_copy(const View< DT, DP...> &dst, typename ViewTraits< DT, DP...>::const_value_type &value, typename std::enable_if< std::is_same< typename ViewTraits< DT, DP...>::specialize, void >::value >::type *=nullptr)
Deep copy a value from Host memory into a view.
KOKKOS_INLINE_FUNCTION constexpr std::enable_if< std::is_integral< iType >::value, size_t >::type extent(const iType &r) const noexcept
rank() to be implemented
Implementation detail of parallel_scan.
Compressed row storage array.
Definition: Kokkos_Crs.hpp:83
KOKKOS_INLINE_FUNCTION size_type numRows() const
Return number of rows in the graph.
Definition: Kokkos_Crs.hpp:127
KOKKOS_INLINE_FUNCTION Crs(const RowMapType &row_map_, const EntriesType &entries_)
Assign to a view of the rhs array. If the old view is the last view then allocated memory is dealloca...
Definition: Kokkos_Crs.hpp:120
Implementation of the ParallelFor operator that has a partial specialization for the device...
Traits class for accessing attributes of a View.