Extents and Strides

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:

Each of OtherExtents is convertible to index_type.
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:

Each of OtherExtents is convertible to index_type.
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.

Strides CUDA 13.4

A ct::strides is a strides like type suitable for describing stride values in a ct::layout_strided_mapping or a ct::strided_view.

The ct::strides type and associated APIs were introduced in CUDA 13.4.

Example

The following example constructs a tensor span representing a \(256 \times 256\) shape with a stride of \(4\) between each element. A strided view over \(4 \times 4\) chunks of data is built with a stride of \(1\) in each dimension.

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

ct::extents shape{256_ic, 256_ic};
ct::strides strides{256_ic * 4_ic, 4_ic};

ct::layout_strided_mapping mapping{shape, strides};
ct::tensor_span t{ptr, mapping};

ct::strided_view v{t, ct::shape{4_ic, 4_ic}, ct::strides{1_ic, 1_ic}};

cuda::tiles::strides

template<ct::integral I, size_t... Strides> struct strides;

A strides like type with static or dynamic stride values specified by Strides and strides type specified by \(I\).

A specialization \(T\) of strides models strides 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.

Exposition Only Member

stride_type __dynamic_strides[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_strides.

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

Member Aliases

using stride_type = I;: The stride type used for storing the dynamic stride values.

using rank_type = size_t;: The rank type which is suitable for selecting a particular stride in the static_stride() and stride() APIs.

rank

__tile__ __host__ __device__ static constexpr rank_type rank() noexcept;: Yields the size of the Strides parameter pack.

rank_dynamic

__tile__ __host__ __device__ static constexpr rank_type rank_dynamic() noexcept;: Yields the number of ct::dynamic_stride values in the Strides parameter pack.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept;: Yields the value of the \(i^{th}\) element of the Strides parameter pack. The behavior is undefined if \(i\) is not less than the length of the Strides pack.

stride

__tile__ __host__ __device__ constexpr stride_type stride(rank_type i) const noexcept;

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

When \(i\) references a dynamic stride, the result is the \(N^{th}\) element of __dynamic_strides where \(N\) is the number of dynamic stride values whose pack index in Strides is less than \(i\).

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

Default Construction

constexpr strides() noexcept = default;: Constructs a ct::strides object where each element of __dynamic_strides is initialized to \(0\).

Construction From Dynamic Strides

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

Constructs a ct::strides from a pack specifying the dynamic stride values. The __dynamic_strides member is initialized as if by __dynamic_strides{stride_type(other)...}.

The atomic constraint validates that:

Each of OtherStrides is convertible to stride_type.
The size of the OtherStrides pack equals rank_dynamic() and is not \(0\).

Example

The following code constructs a strides object with \(8 \times 42 \times 3\) strides.

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

Construction From Strides

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

Constructs a ct::strides from a pack specifying all the stride values. The __dynamic_strides member is initialized by considering each element of the other pack whose index \(N\) corresponds to a dynamic stride value in Strides.

Each such element, after conversion to stride_type, is placed at index \(M\) of __dynamic_strides where \(M\) is the number of dynamic stride values whose pack index in Strides is less than \(N\).

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

The atomic constraint validates that:

Each of OtherStrides is convertible to stride_type.
The size of the OtherStrides pack equals rank(), does not equal rank_dynamic() and is not \(0\).

Example

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

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

Deduction Guide

template<typename ...Ts> strides(Ts...) -> strides<uint32_t, stride-constant-or-dynamic<Ts>...>;

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

Example

The following code produces a stride object of type ct::strides<uint32_t, 4, ct::dynamic_stride>.

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

cuda::tiles::dynamic_stride

enum __dynamic-stride-type : size_t

enumerator dynamic_stride = /* see below */: Yields the sentinel value for encoding a dynamic stride in a strides 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_stride.

Equality Comparison

template<ct::strides_like Lhs, ct::strides_like Rhs> __tile__ __host__ __device__ constexpr bool operator==(Lhs const &lhs, Rhs const &rhs) noexcept;: Indicates whether lhs is stride equivalent to rhs.

Note

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

stride-constant-or-dynamic

template<typename T> inline constexpr size_t stride-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_stride.

Footnotes

1(1,2): See § 18.6 [concepts.object] of ISO/IEC 14882:2024
2(1,2,3,4,5,6): See § 21.3.3 [meta.type.synop] of ISO/IEC 14882:2024
3(1,2): See § 6.8.1 [basic.types.general] of ISO/IEC 14882:2024
4(1,2): See § 12.2.2.3 [over.match.oper] of ISO/IEC 14882:2024