Segmented ranges #1582

upsj · 2024-04-01T10:38:24Z

This PR adds a bunch of utility abstractions for simplifying common iteration patterns in Ginkgo.

irange: Similar to the Python range([start, ] stop[, step]) function, the irange provides a range that can be used in range-for loops (and the corresponding iterators in standard library algorithms):

for (int i = 0; i < size; i++)
for (auto j = begin; j < end; j += step)
// becomes
for (auto i : irange<int>(size))
for (auto j : irange<int>(begin, end, step))
// or with C++17 CTAD
for (auto i : irange(size))
for (auto j : irange(begin, end, step))

The range-for loop has the added advantage that the iteration variable can be made const, which prevents accidentally incrementing the wrong variable, e.g. in a nested loop. The compiler generates 100% identical code for this thanks to inlining, constant folding and common subexpression elimination.

(enumerating_)indexed_iterator is pretty similar to permuting_iterator, but its intended usage is slightly different - it is mainly intended to provide strided access to an array, e.g. when stepping through a Csr matrix row with a whole warp, assigning nonzeros in a striped fashion. The enumerating_ variant returns a struct (index, value). The type is unlikely to be used directly, but used as additional functionality in the following segmented range abstraction.
segmented_((enumerating_)value_)range provides a range-of-ranges abstraction that can cover most of our typical Csr row_ptrs access patterns. The final result is intended to look as follows:

const auto begin = row_ptrs[row];
const auto end = row_ptrs[row + 1];
for (auto nz = begin; nz < end; nz++) {
    const auto col = cols[nz];
    const auto val = vals[nz];
    // ...
}
// with C++17 CTAD and structured bindings
csr_view{row_ptrs, cols, vals, num_rows}; // the actual type will probably differ
for (auto [nz, col, value] : csr_view.enumerate(row)) {
    // ...
}

Additionally, I am planning to make coalesced/striped work assignment more clear with some small helpers that turn regular ranges into strided ranges based on the warp/subgroup lane and size:

const auto begin = row_ptrs[row];
const auto end = row_ptrs[row + 1];
for (auto nz = begin + subwarp.thread_rank(); nz < end; nz += subwarp.size()) {
    const auto col = cols[nz];
    const auto val = vals[nz];
    // ...
}
// with C++17 CTAD and structured bindings
csr_view{row_ptrs, cols, vals, num_rows}; // the actual type will probably differ
for (auto [nz, col, value] : csr_view.enumerate(row).striped(subwarp)) {
    // ...
}

Similarly, we can provide a blocked(subwarp) function, which assigns a consecutive block of indices to each thread, when that is useful. The inspiration for this approach comes from cub's striped and blocked work assignment strategies

TODO:

more comprehensive tests
device tests to check compilation
strided ranges with group support

This reverts commit e473a78.

MarcelKoch

Here are some of my notes. My biggest concern is indexed|enumerated_iterator, they are too similar IMO. We could do with just enumerated_iterator. With some extra storage we could also enable it-> properly in that case.

TBH I find the iterator hierarchy difficult to follow. It would probably be much easier if the type aliases would be more direct.

MarcelKoch · 2024-04-11T08:07:01Z

core/base/integer_range.hpp

+
+
+template <typename IndexType>
+class integer_iterator {


isn't this usually called a counting iterator?

Is this even an iterator? Usually, iterators give access to something. This class only handles an integer and the multiplication with a stride. It doesn't refer to anything (see your own reference type definition).

integer_iterator<int> iterator(12, 2); for (auto&& iter : iterator) { *iter = 3; // This isn't doing anything and should result in a compiler warning }

An alternative name might be: index_generator, index_generator_iterator, or maybe even index_iterator (even though I still don't think it is a real iterator, but it has most of the properties of one).
You could even put this into the detail namespace, as I don't expect it to be used directly.

I think the name iterator is fine. There are different classes of iterators, the one you described would be an output iterator, but there are also input iterators, which you can only read from. Other libraries, e.g. thrust, also call this an iterator, so it's not uncommon.

Counting iterator implies unit stride, this iterator can have arbitrary strides, so I think it's not 100% accurate to call it that. But I agree that index_ would also be appropriate

core/base/integer_range.hpp

MarcelKoch · 2024-04-11T08:38:58Z

core/base/segmented_range.hpp

+template <typename ValueIterator,
+          typename IndexIterator = integer_iterator<
+              typename std::iterator_traits<ValueIterator>::difference_type>>
+class enumerating_indexed_iterator


TBH I don't understand why there is both indexed_iterator and enumerating_indexed_iterator. For me enumerated and indexed are (mostly) equivalent. (One difference might be that enumerated always uses the indices [0, n), while indexed could use any indices.)

From the implementation, the only difference is how you access the value and the index. I'm not sure if this warrants two separate implementations for two nearly identical types.

IMO the enumerating case is the exception, the indexed iterator is the usual case. Take a Csr SpMV, that can be implemented using an indexed_iterator<zip_iterator<IndexType*, ValueType*>>:

for (auto [column, value]: range) { sum += in[column] * value; }

I can only think of a small number of cases where we actually need the nonzero location.

I still think the names indexed and enumerated are too similar. For me they both say that I can iterate over a range and I get both the value and the corresponding index.

And your example, isn't that more of an zip operation than either indexed or enumerated? You are iterating both over the columns array and the values array.

MarcelKoch · 2024-04-11T08:52:39Z

core/base/segmented_range.hpp

+
+
+template <typename ValueIterator,
+          typename IndexIterator = integer_iterator<


Do maybe want to restrict it to always use integer_iterator? Otherwise it might be a bit more difficult to understand how it will work.

You could use this to implement a permuting_iterator, but that's probably beyond the scope of this PR.

MarcelKoch · 2024-04-11T08:55:51Z

core/base/segmented_range.hpp

+    using base_value_type = typename base::value_type;
+    using index_type = typename base::value_type;
+
+    struct enumerated {


I would suggest storing a value of this in the iterator to enable it->value(), it->index(). But in that case, I would also suggest to not store the index and value directly, but pointers, and use accessor functions.

The reason I put this into a struct is that then we can use it in structured bindings

Why did you add a special struct for something that could be a std::pair? The operators == and != are defined by std::pair as well.
Are you worried about the other operators (<, >, ...) defined by std::pair? They compare the first argument and only consider the second if the first are identical.
Or do you simply want access to index and value by those names?

I think std::pair would also not work on the device, some functions there are not marked constexpr.

Without C++17 support, the default way to access the iterator entry is like

for (auto entry : range) { std::get<0>(entry); std::get<1>(entry); entry.index; entry.value; }

the latter is much more clear to me, and less error-prone

TBH, maybe we should push for c++17 with this PR? This is not related to anything that openCARP needs, so we could use it.

core/base/segmented_range.hpp

MarcelKoch · 2024-04-11T11:29:02Z

core/test/base/integer_range.cpp

+{
+    std::vector<int> v;
+
+    for (auto i : gko::irange<int>(1, 4)) {


nit: might also provide make_irange(...) so that the template parameter doesn't need to be defined, while waiting for c++17

MarcelKoch · 2024-04-11T11:49:15Z

core/test/base/integer_range.cpp

+}
+
+
+TEST(IRangeStrided, KnowsItsProperties)


Also, size might be ambiguous. For this case it could be either 2 (how often can you increment before reaching the end), or 6 (the end of the stored half-interval).

MarcelKoch · 2024-04-11T11:54:05Z

core/test/utils/death_test_helpers.hpp

+}
+
+
+#define EXPECT_ASSERT_FAILURE(_expression, ...) \


EXPECT_EXIT is also used elsewhere, do those also need to change?

thoasm

I'm not the biggest fan of the _iterator suffix for most of the classes because they don't really reference any underlying data.
The _range suffix makes sense (classes with begin() and end(), somewhat similar to python range).
The enumerating prefix isn't descriptive enough. To me, those are pair ranges (instead of returning an index or value iterator, they return a pair of both).
Additionally, I would like to have documentation for all new classes (and maybe examples for the more complicated classes like segmented_range).

thoasm · 2024-04-17T12:36:19Z

core/base/iterator_boilerplate.hpp

+/**
+ * Implements all `random_access_iterator` operations for _iterator in terms of
+ * the already implemented advance `operator +=(difference_type)`, the
+ * difference operator `operator-(_iterator, _iterator)` and the deference


nit: typo

Suggested change

* difference operator `operator-(_iterator, _iterator)` and the deference

* difference operator `operator-(_iterator, _iterator)` and the dereference

core/base/integer_range.hpp

thoasm · 2024-04-17T15:44:56Z

core/base/integer_range.hpp

+
+
+template <typename IndexType>
+class integer_iterator {


Is this even an iterator? Usually, iterators give access to something. This class only handles an integer and the multiplication with a stride. It doesn't refer to anything (see your own reference type definition).

integer_iterator<int> iterator(12, 2); for (auto&& iter : iterator) { *iter = 3; // This isn't doing anything and should result in a compiler warning }

An alternative name might be: index_generator, index_generator_iterator, or maybe even index_iterator (even though I still don't think it is a real iterator, but it has most of the properties of one).
You could even put this into the detail namespace, as I don't expect it to be used directly.

thoasm · 2024-04-17T16:18:38Z

core/base/integer_range.hpp

+    template <typename Group>
+    constexpr irange_strided<index_type> striped(Group g) const
+    {
+        return striped(detail::group_traits<Group>::get_local_id(g),
+                       detail::group_traits<Group>::get_size(g));
+    }
+
+    constexpr irange_strided<index_type> striped(index_type local_index,
+                                                 index_type group_size) const
+    {
+        assert(local_index >= 0);
+        assert(local_index < group_size);
+        return irange_strided<index_type>{begin_index() + local_index,
+                                          end_index(), group_size};
+    }


I assume these 2 functions are the reason why you separated irange from irange_strided. Are they targeting SYCL?

They are targeting SIMD/coalescing memory accesses in general. My aim is to turn the first loop into the second:

for (auto i = begin + warp.thread_rank(); i < end; i += warp.size()); for (auto i : irange{begin, end}.striped(warp))

I will remove the templated version though, since that will be part of a separate PR.

thoasm · 2024-04-18T13:19:41Z

core/base/segmented_range.hpp

+    using base_value_type = typename base::value_type;
+    using index_type = typename base::value_type;
+
+    struct enumerated {


Why did you add a special struct for something that could be a std::pair? The operators == and != are defined by std::pair as well.
Are you worried about the other operators (<, >, ...) defined by std::pair? They compare the first argument and only consider the second if the first are identical.
Or do you simply want access to index and value by those names?

thoasm · 2024-04-18T13:37:36Z

core/base/segmented_range.hpp

+    using index_iterator = IndexIterator;
+
+private:
+    using value_traits = std::iterator_traits<index_iterator>;


copy-paste error:

Suggested change

using value_traits = std::iterator_traits<index_iterator>;

using value_traits = std::iterator_traits<value_iterator>;

upsj · 2024-07-15T09:44:29Z

Closing in favor of #1601

upsj added 6 commits March 26, 2024 23:35

add irange

fec10d9

wip: state with C++17 sentinel

e473a78

Revert "wip: state with C++17 sentinel"

6bf41f5

This reverts commit e473a78.

improve irange test

8f7be57

factor out boilerplate code

c00a440

add segmented range abstraction

390114b

upsj self-assigned this Apr 1, 2024

ginkgo-bot added reg:build This is related to the build system. reg:testing This is related to testing. mod:core This is related to the core module. labels Apr 1, 2024

upsj mentioned this pull request Apr 3, 2024

Add segmented array type #1545

Merged

6 tasks

MarcelKoch requested review from MarcelKoch, thoasm and yhmtsai April 5, 2024 09:11

MarcelKoch modified the milestone: Ginkgo 1.8.0 Apr 5, 2024

MarcelKoch reviewed Apr 11, 2024

View reviewed changes

thoasm reviewed Apr 18, 2024

View reviewed changes

This was referenced Apr 29, 2024

Simple segmented ranges #1601

Merged

Add integer range #1602

Merged

MarcelKoch removed this from the Ginkgo 1.8.0 milestone May 3, 2024

tcojean added this to the Ginkgo 1.9.0 milestone May 3, 2024

upsj closed this Jul 15, 2024

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Segmented ranges #1582

Segmented ranges #1582

upsj commented Apr 1, 2024

MarcelKoch left a comment

MarcelKoch Apr 11, 2024

thoasm Apr 17, 2024

MarcelKoch Apr 18, 2024

upsj Apr 19, 2024

MarcelKoch Apr 11, 2024

upsj Apr 12, 2024

MarcelKoch Apr 12, 2024

MarcelKoch Apr 11, 2024

upsj Apr 12, 2024

MarcelKoch Apr 11, 2024

upsj Apr 12, 2024

thoasm Apr 18, 2024

MarcelKoch Apr 18, 2024

upsj Apr 18, 2024

MarcelKoch Apr 18, 2024

MarcelKoch Apr 11, 2024

MarcelKoch Apr 11, 2024

MarcelKoch Apr 11, 2024

MarcelKoch Apr 11, 2024

thoasm left a comment

thoasm Apr 17, 2024

thoasm Apr 17, 2024

thoasm Apr 17, 2024

upsj Apr 19, 2024

thoasm Apr 18, 2024

thoasm Apr 18, 2024

upsj commented Jul 15, 2024



		template <typename ValueIterator,
		typename IndexIterator = integer_iterator<

	* difference operator `operator-(_iterator, _iterator)` and the deference
	* difference operator `operator-(_iterator, _iterator)` and the dereference

	using value_traits = std::iterator_traits<index_iterator>;
	using value_traits = std::iterator_traits<value_iterator>;

		}


		TEST(IRangeStrided, KnowsItsProperties)

Segmented ranges #1582

Segmented ranges #1582

Conversation

upsj commented Apr 1, 2024

MarcelKoch left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

thoasm left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

upsj commented Jul 15, 2024