Integer Range

ct::irange is a forward range representing an increasing sequence of integers separated by a step. The range is constructed with a lower bound \(L\), exclusive upper bound \(U\), and a step \(S > 0\). The range sweeps through the following values in increasing order:

\[\{ L + i \cdot S \quad | \quad i \in \mathbb{Z} \quad i \geq 0 \quad L + i \cdot S < U \}\]

A ct::irange provides the compiler with structured information about the loop iteration bounds which may be used to better optimize the code.

Example

The following code prints the values from the sequence \((5, 7, 9, 11)\). The lower bound is \(5\), upper bound is \(12\) and step is \(2\).

namespace ct = ::cuda::tiles;
for (auto idx : ct::irange(5, 12, 2)) {
  printf("%i\n", idx);
}

irange-sentinel

template<ct::integral T> struct irange-sentinel

Exposition only sentinel type corresponding to ct::irange-iterator. It models the end position of a ct::irange.

The program is ill-formed if \(T\) is cv-qualified.

Exposition Only Members

T __end: Exposition only member variable modeling the upper bound of a ct::irange. This member exists for the purposes of describing the implicitly defined special member functions.

Default Construction

constexpr irange-sentinel() noexcept = default: Constructs an instance of ct::irange-sentinel in a valid but unspecified state. The behavior is undefined if any operation other than destruction or copy/move assignment is performed on such an object.

Construction From End

__tile__ __host__ __device__ constexpr irange-sentinel(T end) noexcept: Direct-list-initializes __end from end.

cuda::tiles::irange_iterator

template<ct::integral T> struct irange_iterator

Forward iterator of a ct::irange. The integer results and step computation is performed in the integral type T.

The program is ill-formed if T is const qualified.

Exposition Only Members

The following members exist for the purposes of describing the implicitly defined special member functions.

T __current: Exposition only member variable modeling the current value of the iterator yielded by a dereference.

T __step: Exposition only member modeling the value to add to this iterator during an increment.

Member Aliases

using difference_type = int64_t
using value_type = T 

Default Construction

constexpr irange_iterator() noexcept = default: Constructs this object in a valid but unspecified state. The behavior is undefined if any operation other than destruction or copy/move assignment is performed on such an object.

Construction From Current and Step

__tile__ __host__ __device__ constexpr irange_iterator(T current, T step) noexcept: Direct-list-initializes __current from current and __step from step.

Dereference

__tile__ __host__ __device__ constexpr T operator*() const noexcept: Yields __current.

Pre-Increment

__host__ __device__ __tile__ constexpr irange_iterator &operator++() noexcept: Advances __current by __step as if by __current += __step. Yields reference to *this.

Post-Increment

__host__ __device__ __tile__ constexpr irange_iterator operator++(int) noexcept: Advances __current by __step as if by __current += __step. Yields a copy of the iterator prior to the increment operation.

Iterator Comparison

__host__ __device__ __tile__ constexpr bool operator==(irange_iterator const &other) const noexcept: Compares the __current value of *this and other for equality. The behavior is undefined if other was derived from a different ct::irange instance than *this.

Note

operator!= is available by rewritten operator candidates.

Sentinel Comparison

__host__ __device__ __tile__ constexpr bool operator==(irange-sentinel<T> const &other) const noexcept: Returns true if __current is not less than other.__end.

Note

operator!= is available by rewritten operator candidates.

cuda::tiles::irange

template<ct::integral T> struct irange

Forward range modeling an increasing sequence of integers separated by a step. For an example, see the introduction.

The program is ill-formed if T is const qualified.

Exposition Only Members

The following exposition only members exist for the purposes of describing the implicitly defined special member functions.

T __lb: Exposition only member variable modeling the inclusive lower bound of the range.

T __ub: Exposition only member variable modeling the exclusive upper bound of the range.

T __step: Exposition only member variable modeling the step size of the range.

Aliases

using iterator = irange_iterator<T>
using const_iterator = irange_iterator<T>

Construction From Bounds and Step

__host__ __device__ __tile__ constexpr irange(T lb, T ub, T step = 1) noexcept

Direct-list-initializes __ub from ub, __lb from lb and __step from step. Let \(N\) denote the largest value representable by T. The behavior is undefined if any of the following conditions hold after the initialization of members:

__step <= 0
__ub > N - __step + 1.

Example

The following iteration sweeps through the sequence \((-10, -9, -8, -7)\). The step is not explicitly specified and defaults to \(1\).

namespace ct = ::cuda::tiles;
for (auto idx : ct::irange(-10, -6)) {
  printf("%i\n", idx);
}

begin

__host__ __device__ __tile__ constexpr iterator begin() const noexcept: Yields iterator{__lb, __step}.

end

__host__ __device__ __tile__ constexpr irange-sentinel<T> end() const noexcept: Yields irange-sentinel<T>{__ub}.

empty

__host__ __device__ __tile__ constexpr bool empty() const noexcept

Determines if iteration over this range would yield zero elements.

Example

The following code yields true in both calls empty() because the lower bound is greater than or equal to the upper bound.

namespace ct = ::cuda::tiles;
ct::irange x(20, 10);
auto r0 = x.empty();

ct::irange y(20, 20);
auto r1 = y.empty();

size

__host__ __device__ __tile__ constexpr iterator::difference_type size() const noexcept: Determines the number of elements this range would produce during iteration. This value direct-initializes the return type, possibly triggering overflow.

lower_bound

__host__ __device__ __tile__ constexpr T lower_bound() const noexcept: Yields __lb.

upper_bound

__host__ __device__ __tile__ constexpr T upper_bound() const noexcept: Yields __ub.

step

__host__ __device__ __tile__ constexpr T step() const noexcept: Yields __step.

Deduction Guides

template<typename U1, typename U2> irange(U1, U2) -> ct::irange<ct::arithmetic_tile_conversion_t<U1, U2>>;
template<typename U1, typename U2, typename U3> irange(U1, U2, U3) -> ct::irange<ct::arithmetic_tile_conversion_t<ct::arithmetic_tile_conversion_t<U1, U2>, U3>>;: Deduction guides enabling CTAD from integer arguments.