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;
104 Crs() : row_map(), entries() {}
107 Crs(
const Crs& rhs) : row_map(rhs.row_map), entries(rhs.entries)
110 template<
class EntriesType,
class RowMapType>
111 Crs(
const RowMapType& row_map_,
const EntriesType& entries_) : row_map(row_map_), entries(entries_)
119 row_map = rhs.row_map;
120 entries = rhs.entries;
131 KOKKOS_INLINE_FUNCTION
133 return (row_map.
extent(0) != 0) ?
134 row_map.
extent(0) -
static_cast<size_type
> (1) :
135 static_cast<size_type> (0);
141 template<
class OutCounts,
146 void get_crs_transpose_counts(
148 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in,
149 std::string
const& name =
"transpose_counts");
151 template<
class OutCounts,
153 typename OutCounts::value_type get_crs_row_map_from_counts(
156 std::string
const& name =
"row_map");
158 template<
class DataType,
163 Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
164 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in);
175 template <
class InCrs,
class OutCounts>
176 class GetCrsTransposeCounts {
178 using execution_space =
typename InCrs::execution_space;
179 using self_type = GetCrsTransposeCounts<InCrs, OutCounts>;
180 using index_type =
typename InCrs::size_type;
185 KOKKOS_INLINE_FUNCTION
186 void operator()(index_type i)
const {
187 atomic_increment( &out[in.entries(i)] );
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())));
195 execution_space::fence();
199 template <
class InCounts,
class OutRowMap>
200 class CrsRowMapFromCounts {
202 using execution_space =
typename InCounts::execution_space;
203 using value_type =
typename OutRowMap::value_type;
204 using index_type =
typename InCounts::size_type;
209 last_value_type m_last_value;
211 KOKKOS_INLINE_FUNCTION
212 void operator()(index_type i, value_type& update,
bool final_pass)
const {
213 if (i < m_in.size()) {
215 if (final_pass) m_out(i + 1) = update;
216 }
else if (final_pass) {
218 m_last_value() = update;
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 {
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") {
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));
236 auto last_value = Kokkos::create_mirror_view(m_last_value);
242 template <
class InCrs,
class OutCrs>
243 class FillCrsTransposeEntries {
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;
250 using counters_type = View<index_type*, memory_space>;
253 counters_type counters;
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;
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())));
274 execution_space::fence();
286 template<
class OutCounts,
291 void get_crs_transpose_counts(
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);
300 template<
class OutRowMap,
302 typename OutRowMap::value_type get_crs_row_map_from_counts(
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();
311 template<
class DataType,
316 Crs<DataType, Arg1Type, Arg2Type, SizeType>& out,
317 Crs<DataType, Arg1Type, Arg2Type, SizeType>
const& in)
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 ;
324 Kokkos::get_crs_transpose_counts(counts, in);
325 Kokkos::get_crs_row_map_from_counts(out.row_map, counts,
328 out.entries = decltype(out.entries)(
"transpose_entries", in.entries.size());
330 FillCrsTransposeEntries<crs_type, crs_type> entries_functor(in, out);
333 template<
class CrsType,
335 class ExecutionSpace =
typename CrsType::execution_space>
336 struct CountAndFillBase;
338 template<
class CrsType,
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;
348 counts_type m_counts;
350 inline void operator()(Count, size_type i)
const {
351 m_counts(i) = m_functor(i,
nullptr);
354 inline void operator()(Fill, size_type i)
const {
355 auto j = m_crs.row_map(i);
362 (j ==
static_cast<decltype(j)
>(m_crs.entries.extent(0))) ?
363 nullptr : (&(m_crs.entries(j)));
366 CountAndFillBase(CrsType& crs, Functor
const& f):
372 #if defined( KOKKOS_ENABLE_CUDA ) 373 template<
class CrsType,
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;
382 counts_type m_counts;
384 __device__
inline void operator()(Count, size_type i)
const {
385 m_counts(i) = m_functor(i,
nullptr);
388 __device__
inline void operator()(Fill, size_type i)
const {
389 auto j = m_crs.row_map(i);
396 (j ==
static_cast<decltype(j)
>(m_crs.entries.extent(0))) ?
397 nullptr : (&(m_crs.entries(j)));
400 CountAndFillBase(CrsType& crs, Functor
const& f):
407 template<
class CrsType,
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):
421 using execution_space =
typename CrsType::execution_space;
422 this->m_counts = counts_type(
"counts", nrows);
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));
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);
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));
445 template<
class CrsType,
447 void count_and_fill_crs(
449 typename CrsType::size_type nrows,
451 Kokkos::CountAndFill<CrsType, Functor>(crs, nrows, f);
KOKKOS_INLINE_FUNCTION size_type numRows() const
Return number of rows in the graph.
Crs(const Crs &rhs)
Copy constructor (shallow copy).
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...
Implementation detail of parallel_scan.
Compressed row storage array.
Implementation of the ParallelFor operator that has a partial specialization for the device...
Crs()
Construct an empty view.
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.
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