44 #ifndef KOKKOS_CRS_HPP
45 #define KOKKOS_CRS_HPP
79 template<
class DataType,
81 class Arg2Type = void,
82 typename SizeType =
typename ViewTraits<DataType*, Arg1Type, Arg2Type, void >::size_type>
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;
106 KOKKOS_FUNCTION
Crs() =
default;
107 KOKKOS_FUNCTION
Crs(
Crs const &) =
default;
108 KOKKOS_FUNCTION
Crs(
Crs &&) =
default;
109 KOKKOS_FUNCTION
Crs& operator=(
Crs const &) =
default;
110 KOKKOS_FUNCTION
Crs& operator=(
Crs &&) =
default;
111 KOKKOS_FUNCTION ~
Crs() =
default;
117 template<
class EntriesType,
class RowMapType>
118 KOKKOS_INLINE_FUNCTION
119 Crs(
const RowMapType& row_map_,
const EntriesType& entries_)
120 : row_map(row_map_), entries(entries_)
126 KOKKOS_INLINE_FUNCTION
128 return (row_map.
extent(0) != 0) ?
129 row_map.
extent(0) -
static_cast<size_type
> (1) :
130 static_cast<size_type> (0);
136 template<
class OutCounts,
141 void get_crs_transpose_counts(
143 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in,
144 std::string
const& name =
"transpose_counts");
146 template<
class OutCounts,
148 typename OutCounts::value_type get_crs_row_map_from_counts(
151 std::string
const& name =
"row_map");
153 template<
class DataType,
158 Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
159 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in);
170 template <
class InCrs,
class OutCounts>
171 class GetCrsTransposeCounts {
173 using execution_space =
typename InCrs::execution_space;
174 using self_type = GetCrsTransposeCounts<InCrs, OutCounts>;
175 using index_type =
typename InCrs::size_type;
180 KOKKOS_INLINE_FUNCTION
181 void operator()(index_type i)
const {
182 atomic_increment( &out[in.entries(i)] );
184 GetCrsTransposeCounts(InCrs
const& arg_in, OutCounts
const& arg_out):
185 in(arg_in),out(arg_out) {
186 using policy_type = RangePolicy<index_type, execution_space>;
188 const closure_type closure(*
this, policy_type(0, index_type(in.entries.size())));
190 execution_space::fence();
194 template <
class InCounts,
class OutRowMap>
195 class CrsRowMapFromCounts {
197 using execution_space =
typename InCounts::execution_space;
198 using value_type =
typename OutRowMap::value_type;
199 using index_type =
typename InCounts::size_type;
204 last_value_type m_last_value;
206 KOKKOS_INLINE_FUNCTION
207 void operator()(index_type i, value_type& update,
bool final_pass)
const {
208 if (i < m_in.size()) {
210 if (final_pass) m_out(i + 1) = update;
211 }
else if (final_pass) {
213 m_last_value() = update;
216 KOKKOS_INLINE_FUNCTION
217 void init(value_type& update)
const { update = 0; }
218 KOKKOS_INLINE_FUNCTION
219 void join(
volatile value_type& update,
const volatile value_type& input)
const {
222 using self_type = CrsRowMapFromCounts<InCounts, OutRowMap>;
223 CrsRowMapFromCounts(InCounts
const& arg_in, OutRowMap
const& arg_out):
224 m_in(arg_in), m_out(arg_out), m_last_value(
"last_value") {
226 value_type execute() {
227 using policy_type = RangePolicy<index_type, execution_space>;
229 closure_type closure(*
this, policy_type(0, m_in.size() + 1));
231 auto last_value = Kokkos::create_mirror_view(m_last_value);
237 template <
class InCrs,
class OutCrs>
238 class FillCrsTransposeEntries {
240 using execution_space =
typename InCrs::execution_space;
241 using memory_space =
typename InCrs::memory_space;
242 using value_type =
typename OutCrs::entries_type::value_type;
243 using index_type =
typename InCrs::size_type;
245 using counters_type = View<index_type*, memory_space>;
248 counters_type counters;
250 KOKKOS_INLINE_FUNCTION
251 void operator()(index_type i)
const {
252 auto begin = in.row_map(i);
253 auto end = in.row_map(i + 1);
254 for (
auto j = begin; j < end; ++j) {
255 auto ti = in.entries(j);
256 auto tbegin = out.row_map(ti);
257 auto tj = atomic_fetch_add( &counters(ti), 1 );
258 out.entries( tbegin + tj ) = i;
261 using self_type = FillCrsTransposeEntries<InCrs, OutCrs>;
262 FillCrsTransposeEntries(InCrs
const& arg_in, OutCrs
const& arg_out):
263 in(arg_in),out(arg_out),
264 counters(
"counters", arg_out.numRows()) {
265 using policy_type = RangePolicy<index_type, execution_space>;
267 const closure_type closure(*
this, policy_type(0, index_type(in.numRows())));
269 execution_space::fence();
281 template<
class OutCounts,
286 void get_crs_transpose_counts(
288 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in,
289 std::string
const& name) {
290 using InCrs = Crs<DataType, Arg1Type, Arg2Type, SizeType>;
291 out = OutCounts(name, in.numRows());
292 Kokkos::Impl::GetCrsTransposeCounts<InCrs, OutCounts> functor(in, out);
295 template<
class OutRowMap,
297 typename OutRowMap::value_type get_crs_row_map_from_counts(
300 std::string
const& name) {
301 out = OutRowMap(ViewAllocateWithoutInitializing(name), in.size() + 1);
302 Kokkos::Impl::CrsRowMapFromCounts<InCounts, OutRowMap> functor(in, out);
303 return functor.execute();
306 template<
class DataType,
311 Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
312 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in)
314 typedef Crs<DataType, Arg1Type, Arg2Type, SizeType> crs_type ;
315 typedef typename crs_type::memory_space memory_space ;
316 typedef View<SizeType*, memory_space> counts_type ;
319 Kokkos::get_crs_transpose_counts(counts, in);
320 Kokkos::get_crs_row_map_from_counts(out.row_map, counts,
323 out.entries = decltype(out.entries)(
"transpose_entries", in.entries.size());
325 FillCrsTransposeEntries<crs_type, crs_type> entries_functor(in, out);
328 template<
class CrsType,
330 class ExecutionSpace =
typename CrsType::execution_space>
331 struct CountAndFillBase;
333 template<
class CrsType,
335 class ExecutionSpace>
336 struct CountAndFillBase {
337 using data_type =
typename CrsType::size_type;
338 using size_type =
typename CrsType::size_type;
339 using row_map_type =
typename CrsType::row_map_type;
340 using counts_type = row_map_type;
343 counts_type m_counts;
345 inline void operator()(Count, size_type i)
const {
346 m_counts(i) = m_functor(i,
nullptr);
349 inline void operator()(Fill, size_type i)
const {
350 auto j = m_crs.row_map(i);
357 (j ==
static_cast<decltype(j)
>(m_crs.entries.extent(0))) ?
358 nullptr : (&(m_crs.entries(j)));
361 CountAndFillBase(CrsType& crs, Functor
const& f):
367 #if defined( KOKKOS_ENABLE_CUDA )
368 template<
class CrsType,
370 struct CountAndFillBase<CrsType, Functor, Kokkos::Cuda> {
371 using data_type =
typename CrsType::size_type;
372 using size_type =
typename CrsType::size_type;
373 using row_map_type =
typename CrsType::row_map_type;
374 using counts_type = row_map_type;
377 counts_type m_counts;
379 __device__
inline void operator()(Count, size_type i)
const {
380 m_counts(i) = m_functor(i,
nullptr);
383 __device__
inline void operator()(Fill, size_type i)
const {
384 auto j = m_crs.row_map(i);
391 (j ==
static_cast<decltype(j)
>(m_crs.entries.extent(0))) ?
392 nullptr : (&(m_crs.entries(j)));
395 CountAndFillBase(CrsType& crs, Functor
const& f):
402 template<
class CrsType,
404 struct CountAndFill :
public CountAndFillBase<CrsType, Functor> {
405 using base_type = CountAndFillBase<CrsType, Functor>;
406 using typename base_type::data_type;
407 using typename base_type::size_type;
408 using typename base_type::counts_type;
409 using typename base_type::Count;
410 using typename base_type::Fill;
411 using entries_type =
typename CrsType::entries_type;
412 using self_type = CountAndFill<CrsType, Functor>;
413 CountAndFill(CrsType& crs, size_type nrows, Functor
const& f):
416 using execution_space =
typename CrsType::execution_space;
417 this->m_counts = counts_type(
"counts", nrows);
419 using count_policy_type = RangePolicy<size_type, execution_space, Count>;
420 using count_closure_type =
422 const count_closure_type closure(*
this, count_policy_type(0, nrows));
425 auto nentries = Kokkos::
426 get_crs_row_map_from_counts(this->m_crs.row_map, this->m_counts);
427 this->m_counts = counts_type();
428 this->m_crs.entries = entries_type(
"entries", nentries);
430 using fill_policy_type = RangePolicy<size_type, execution_space, Fill>;
431 using fill_closure_type =
433 const fill_closure_type closure(*
this, fill_policy_type(0, nrows));
440 template<
class CrsType,
442 void count_and_fill_crs(
444 typename CrsType::size_type nrows,
446 Kokkos::CountAndFill<CrsType, Functor>(crs, nrows, f);
Implementation detail of parallel_scan.
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.
Compressed row storage array.
KOKKOS_INLINE_FUNCTION size_type numRows() const
Return number of rows in the graph.
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...
Implementation of the ParallelFor operator that has a partial specialization for the device...
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 noexcept
rank() to be implemented