Extents

A ct::extents is an extents like type suitable for describing the dimensions of a ct::tile or ct::tensor_span.

Example

In the following code, the extents object x and y represents the shape \(4 \times 7\). where the first dimension is known at compile time and the second is determined at runtime.

The extents object z represents the shape \(33 \times 0\) where both dimensions are known at compile time.

In all cases, the index type used for index computations is uint32_t.

namespace ct = ::cuda::tiles;
using namespace ct::literals;

ct::extents                                   x{4_ic, 7};
ct::extents<uint32_t, 4, ct::dynamic_extent>  y{7};
ct::shape<33, 0>                              z;

cuda::tiles::extents

template<ct::integral I, size_t... Exts>
struct extents;

An extents like type with static or dynamic extents specified by Exts and index type specified by I.

A specialization \(T\) of extents models extents like and satisfies the following requirements:

• std::regular 1

• std::is_nothrow_move_constructible_v 2

• std::is_nothrow_move_assignable_v 2

• std::is_nothrow_swappable_v 2

\(T\) is trivially copyable 3.

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

The shape described by an extents object is determined by the static_extent() and extent() member functions.

Exposition Only Member

index_type __dynamic_extents[rank_dynamic()];

For the purposes of defining the object size, alignment, and behavior of implicitly generated member functions, \(T\) has an exposition only member variable __dynamic_extents.

The member variable determines the value of the dynamic extents of an object of type \(T\).

Member Aliases

using index_type = I;

The index type used for storing the dynamic extent values.

using rank_type = size_t;

The rank type which is suitable for selecting a particular extent in the static_extent() and extent() APIs.

rank

__tile__ __host__ __device__ static constexpr rank_type rank() noexcept;

Yields the size of the Exts parameter pack.

rank_dynamic

__tile__ __host__ __device__ static constexpr rank_type rank_dynamic() noexcept;

Yields the number of ct::dynamic_extent values in the Exts parameter pack.

static_extent

__tile__ __host__ __device__ static constexpr size_t static_extent(rank_type i) noexcept;

Yields the value of the \(i^{th}\) element of the Exts parameter pack. The behavior is undefined if \(i\) is not less than the length of the Exts pack.

extent

__tile__ __host__ __device__ constexpr index_type extent(rank_type i) const noexcept;

Yields the value of the \(i^{th}\) element of the Exts pack or the runtime value for the \(i^{th}\) dimension if \(i\) corresponds to a ct::dynamic_extent value in the Exts pack.

When \(i^{th}\) references a dynamic extent, the result is the \(N^{th}\) element of __dynamic_extents where \(N\) is the number of dynamic extents whose pack index in Exts is less than \(i\).

The behavior is undefined if \(i\) is not less than rank() or if \(i\) is a static extent whose value is not representable in index_type.

Default Construction

constexpr extents() noexcept = default;

Constructs a ct::extents object where each element of __dynamic_extents is initialized to \(0\).

Construction From Dynamic Extents

template<typename ...OtherExtents>
requires /* atomic constraint */
__tile__ __host__ __device__ explicit constexpr extents(OtherExtents... other) noexcept;

Constructs a ct::extents from a pack specifying the dynamic extent values. The __dynamic_extents member is initialized as if by __dynamic_extents{index_type(other)...}.

The atomic constraint validates that:

1. Each of OtherExtents is convertible to index_type.

2. The size of the OtherExtents pack equals rank_dynamic() and is not \(0\).

Example

The following code constructs an extents object modeling the shape \(8 \times 42 \times 3\).

namespace ct = ::cuda::tiles;
ct::extents<int32_t, 8, ct::dynamic_extent, 3> e{42};

Construction From Extents

template<typename ...OtherExtents>
requires /*  atomic constraint */
__tile__ __host__ __device__ explicit constexpr extents(OtherExtents... other) noexcept;

Constructs a ct::extents from a pack specifying all the extent values. The __dynamic_extents member is initialized as by considering each element of the other pack whose index \(N\) corresponds to a dynamic extent value in Exts.

Each such element, after conversion to index_type, is placed at index \(M\) of __dynamic_extents where \(M\) is the number of dynamic extent values whose pack index in Exts is less than \(N\).

The behavior is undefined if any element of other corresponding to a static extent does not equal the static extent value after conversion to index_type.

The atomic constraint validates that:

1. Each of OtherExtents is convertible to index_type.

2. The size of the OtherExtents pack equals rank(), does not equal rank_dynamic() and is not \(0\).

Example

The following code constructs an extents object modeling the shape \(8 \times 42 \times 3\).

namespace ct = ::cuda::tiles;
ct::extents<int32_t, 8, ct::dynamic_extent, 3> e{8, 42, 3};

Deduction Guide

template<typename ...Ts>
extents(Ts...) -> extents<uint32_t, extent-constant-or-dynamic<Ts>...>;

Deduction guide enabling CTAD from a collection of ct::integral_constant or integral arguments.

Example

The following code produces an extents object of type ct::extents<uint32_t, 4, ct::dynamic_extent>.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
ct::extents x{4_ic, 7};

cuda::tiles::dynamic_extent

enum /* unspecified identifier */ : size_t

enumerator dynamic_extent = /* see below */

Yields the sentinel value for encoding a dynamic extent in an extents like type. The value is the maximum value which can be represented by size_t.

Note

The program is ill-formed when attempting to form the address of ct::dynamic_extent.

cuda::tiles::shape

template<ct::size_t... Ts>
requires ct::shape_like<ct::extents<uint32_t, Ts...>>
using shape;

Convenience alias for producing a shape like extents whose index type is uint32_t.

Equality Comparison

template<ct::extents_like Lhs, ct::extents_like Rhs>
__tile__ __host__ __device__ constexpr bool operator==(Lhs const &lhs, Rhs const &rhs) noexcept;

Indicates whether lhs is extent equivalent to rhs.

Note

A corresponding overload for operator!= is available via rewritten operator candidates 4.

extent-constant-or-dynamic

template<typename T>
inline constexpr size_t extent-constant-or-dynamic<T>

Exposition only variable template defined as follows:

• If \(T\) is a specialization of ct::integral_constant, yields the result of converting T::value to size_t.

• Otherwise, yields ct::dynamic_extent.

Footnotes

1

See § 18.6 [concepts.object] of ISO/IEC 14882:2024

2(1,2,3)

See § 21.3.3 [meta.type.synop] of ISO/IEC 14882:2024

3

See § 6.8.1 [basic.types.general] of ISO/IEC 14882:2024

4

See § 12.2.2.3 [over.match.oper] of ISO/IEC 14882:2024