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