Layout Mappings

A layout mapping describes how the elements of a ct::tensor_span will be arranged in memory.

cuda::tiles::layout_right

struct layout_right

A layout policy representing the row-major element order.

The last multi-dimensional index varies the fastest with the increasing elements in memory.

Example

The following code creates \(2 \times 3\) row-major ct::tensor_span over the elements \((0, 1, \ldots, 5)\):

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

int x[] = {0, 1, 2, 3, 4, 5};

// Syntax 1
ct::tensor_span t1{&x[0], ct::extents{2_ic, 3_ic}, ct::layout_right{}};

// Syntax 2
ct::layout_right_mapping m{ct::extents{2_ic, 3_ic}};
ct::tensor_span t2{&x[0], m};
\[\begin{split}\begin{pmatrix} 0 & 1 & 2 \\ 3 & 4 & 5 \end{pmatrix}\end{split}\]
template<ct::extents_like E>
using mapping = ct::layout_right_mapping<E>

Alias template yielding the associated layout mapping type for representing the row-major arrangement of an array whose shape is an instance of \(E\).

cuda::tiles::layout_right_mapping

template<ct::extents_like E>
struct layout_right_mapping

Layout mapping representing a row-major arrangement of an array whose shape is an instance of \(E\). See ct::layout_right for an example.

A specialization \(T\) of ct::layout_right_mapping models the requirements of layout mapping. Additionally, \(T\) satisfies the following constraints if the corresponding constraint is satisfied by the underlying extents like type \(E\):

  • std::is_nothrow_move_constructible_v<T> 1

  • std::is_nothrow_move_assignable_v<T> 1

  • std::is_nothrow_swappable_v<T> 1

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

Exposition Only Members

extents_type __extents

\(T\) contains an exposition only member variable __extents denoting this object’s layout mapping shape. This member exists for the purposes of defining the behavior of the implicitly generated special member functions.

Member Aliases

using extents_type = E
using rank_type = typename extents_type::rank_type
using index_type = typename extents_type::index_type
using layout_type = ct::layout_right

Constructor

__tile__ __host__ __device__ explicit layout_right_mapping(extents_type const &other) noexcept

Constructs an instance of this layout mapping by direct-list-initializing __extents with other.

is_always_strided

__tile__ __host__ __device__ static constexpr bool is_always_strided() noexcept

Yields true.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept

Returns the static stride at dimension \(i\) for a row-major layout of shape __extents.

Let \(E\) denote the extents_type and let \(N\) be its rank. The behavior is undefined if \(i\) is not in the range \([0, N)\).

If the value extents_type::static_extent(k) is ct::dynamic_extent for some value \(i < k < N\), the result is ct::dynamic_extent. Otherwise, the result is

\[\prod_{k=i + 1}^{N - 1} E_k\]

and the behavior is undefined if this value is not representable in size_t or if the value equals ct::dynamic_extent.

stride

__tile__ __host__ __device__ index_type stride(rank_type i) const noexcept

Returns the stride at dimension \(i\) for a row major layout of shape __extents.

Let \(e\) denote __extents and let \(N\) be its rank. The behavior is undefined if \(i\) is not in the range \([0, N)\).

The result is

\[\prod_{k=i + 1}^{N - 1} e_k\]

The behavior is undefined if this value is not representable in index_type.

extents

__tile__ __host__ __device__ extents_type const &extents() const noexcept

Yields the glvalue __extents.

cuda::tiles::layout_left

struct layout_left

Layout policy representing a column-major arrangement of elements. The first multi-dimensional index varies the fastest with the increasing elements in memory.

Example

The following code creates \(2 \times 3\) column-major ct::tensor_span over the elements \((0, 1, \ldots, 5)\):

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

int x[] = {0, 1, 2, 3, 4, 5};

// Syntax 1
ct::tensor_span p1{&x[0], ct::extents{2_ic, 3_ic}, ct::layout_left{}};

// Syntax 2
ct::layout_left_mapping m{ct::extents{2_ic, 3_ic}};
ct::tensor_span p2{&x[0], m};
\[\begin{split}\begin{pmatrix} 0 & 2 & 4 \\ 1 & 3 & 5 \end{pmatrix}\end{split}\]
template<ct::extents_like E>
using mapping = ct::layout_left_mapping<E>

Alias template yielding the associated layout mapping type for representing the column-major arrangement of an array whose shape is an instance of \(E\).

cuda::tiles::layout_left_mapping

template<ct::extents_like E>
struct layout_left_mapping

Layout mapping representing a column-major arrangement of an array whose shape is an instance of \(E\). For an example, see ct::layout_left.

A specialization \(T\) of ct::layout_left_mapping models the requirements of layout mapping. Additionally, \(T\) satisfies the following constraints if the corresponding constraint is satisfied by the underlying extents like type \(E\):

  • std::is_nothrow_move_constructible_v<T> 1

  • std::is_nothrow_move_assignable_v<T> 1

  • std::is_nothrow_swappable_v<T> 1

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

Exposition Only Members

extents_type __extents

\(T\) contains an exposition only member variable __extents denoting this object’s layout mapping shape. This member exists for the purposes of defining the behavior of the implicitly generated special member functions.

Types

using extents_type = E
using rank_type = typename extents_type::rank_type
using index_type = typename extents_type::index_type
using layout_type = ct::layout_left

Constructor

__tile__ __host__ __device__ explicit layout_left_mapping(extents_type const &other) noexcept

Constructs an instance of this layout mapping by direct-list-initializing __extents with other.

is_always_strided

__tile__ __host__ __device__ static constexpr bool is_always_strided() noexcept

Yields true.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept

Returns the static stride at dimension \(i\) for a column major layout of shape __extents.

Let \(E\) denote extents_type and let \(N\) be its rank. The behavior is undefined if \(i\) is not in the range \([0, N)\).

If extents_type::static_extent(i) is ct::dynamic_extent for some value \(0 \leq i < dim\), the result is ct::dynamic_extent. Otherwise, the result is

\[\prod_{k=0}^{i - 1} E_k\]

and the behavior is undefined if this value is not representable in size_t or equals ct::dynamic_extent.

stride

__tile__ __host__ __device__ index_type stride(rank_type i) const noexcept

Returns the stride at dimension \(i\) for a column major layout of shape __extents.

Let \(e\) denote __extents and let \(N\) be its rank. The behavior is undefined if \(i\) is not in the range \([0, N)\).

The returned value is

\[\prod_{k=0}^{i - 1} e_k\]

The behavior is undefined if this value is not representable in index_type.

extents

__tile__ __host__ __device__ extents_type const &extents() const noexcept

Yields the glvalue __extents.

cuda::tiles::layout_right_padded

template<size_t A>
struct layout_right_padded

Layout policy representing a row-major element order whose last dimension is padded to ensure alignment of consecutive rows. The template parameter \(A\) specifies the desired alignment of each row. If \(A\) is ct::dynamic_extent, the alignment is instead supplied as a runtime argument when constructing an instance of layout_right_padded_mapping.

In a padded row-major layout, the last multi-dimensional index varies the fastest with the increasing elements in memory. Additionally, the position of each element anchoring a row is divisible by the desired alignment.

Example

Creates a \(2 \times 3\) ct::tensor_span object over the elements \(0, 1, \ldots, 7\) where each row is aligned to \(4\).

In the second case, the specified alignment of \(2\) is less then the row length of \(3\). The alignment is rounded up to the smallest multiple that is greater than the row length.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[] = {0, 1, 2, 3, 4, 5, 6, 7};

// Syntax 1
ct::layout_right_padded_mapping m1{ct::extents{2_ic, 3_ic}, 4_ic};
ct::tensor_span p1{&x[0], m1};

// Syntax 2
ct::layout_right_padded_mapping m2{ct::extents{2_ic, 3_ic}, 2_ic};
ct::tensor_span p2{&x[0], m2};
\[\begin{split}\begin{pmatrix} 0 & 1 & 2 \\ 4 & 5 & 6 \end{pmatrix}\end{split}\]
template<ct::extents_like E>
using mapping = ct::layout_right_padded_mapping<E, Alignment>

Alias template yielding the associated layout mapping type for representing a padded row-major arrangement of an array whose shape is an instance of \(E\).

cuda::tiles::layout_right_padded_mapping

template<ct::extents_like E, size_t A>
struct layout_right_padded_mapping

Layout mapping representing a row-major element order whose shape is an instance of \(E\). The last dimension is padded to ensure alignment of consecutive rows. For an example, see ct::layout_right_padded.

The template parameter \(A\) specifies the desired alignment of each row. If \(A\) is ct::dynamic_extent, the alignment is instead supplied as a runtime argument in the constructor ct::layout_right_padded_mapping::layout_right_padded_mapping().

A specialization \(T\) of ct::layout_right_padded_mapping models the requirements of layout mapping. Additionally, \(T\) satisfies the following constraints if the corresponding constraint is satisfied by the underlying extents like type \(E\):

  • std::is_nothrow_move_constructible_v<T> 1

  • std::is_nothrow_move_assignable_v<T> 1

  • std::is_nothrow_swappable_v<T> 1

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

Exposition Only Members

extents_type __extents

Exposition-only member denoting this object’s layout mapping shape.

index_type __alignment;

Exposition-only member denoting the desired alignment of each row specified at runtime. This member is present only when A is ct::dynamic_extent.

using __padded_extents_t = /* see below */

Exposition-only member denoting the type of a padded extents object. Let \(E\) denote extents_type and let \(N\) be its rank.

The padded extents type is the unique ct::extents specialization \(S\) satisfying:

  1. The index type and rank of \(S\) match those of \(E\).

  2. The static extents of \(S\) match those of \(E\) with the possible exception of the dimension \(N - 1\) when \(N \neq 0\).

  3. If \(N \neq 0\), the value \(S_{N - 1}\) is

    1. ct::dynamic_extent if either \(A\) or \(E_{N - 1}\) is ct::dynamic_extent.

    2. The smallest multiple of \(A\) that is not less than \(E_{N - 1}\) otherwise. If this value is not representable in the type size_t, the behavior is undefined when constructing an instance of this object.

__tile__ __host__ __device__ constexpr __padded_extents_t __padded_extents() const noexcept

Exposition only function that yields a padded extents object based on the values of __extents and __alignment when present. Let \(e\) denote __extents and let \(N\) be its rank. Let \(a\) denote __alignment when present and A otherwise. The returned value is the unique instance \(s\) of __padded_extents_t satisfying:

  1. The extent values of \(s\) match those of \(e\) with the possible exception of of dimension \(N-1\) when \(N \neq 0\).

  2. When \(N \neq 0\), the value \(s_{N-1}\) is the smallest multiple of \(a\) that is not less than \(e_{N-1}\). If this value is not representable in the type index_type, the behavior of this function is undefined.

Member Aliases

using extents_type = E
using rank_type = typename extents_type::rank_type
using index_type = typename extents_type::index_type
using layout_type = ct::layout_right_padded<Alignment>

layout_right_padded_mapping

template<typename T>
requires is-convertible-v<T, index_type>
__tile__ __host__ __device__ layout_right_padded_mapping(extents_type const &other, T alignment) noexcept

Constructs an instance of this layout mapping by direct-list-initializing __extents with other. When __alignment is present, it is direct-initialized with alignment.

is_always_strided

__tile__ __host__ __device__ static constexpr bool is_always_strided() noexcept

Returns true.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept

Returns the static stride at dimension \(i\). Behaves as if by invoking

ct::layout_right_mapping<__padded_extents_t>::static_stride(i)

stride

__tile__ __host__ __device__ index_type stride(rank_type i) const noexcept

Returns the stride at dimension \(i\). Behaves as if by invoking

ct::layout_right_mapping{__padded_extents()}.stride(i).

extents

__tile__ __host__ __device__ extents_type const &extents() const noexcept

Yields the glvalue __extents.

Deduction Guides

template<typename E, typename A>
layout_right_padded_mapping(E const&, A) -> layout_right_padded_mapping<E, extent-constant-or-dynamic<A>>;

Deduction guide enabling class template argument deduction from an ct::integral_constant or integral alignment argument.

cuda::tiles::layout_left_padded

template<size_t A>
struct layout_left_padded

Layout policy representing a column-major element order whose first dimension is padded to ensure alignment of consecutive columns. The template parameter \(A\) specifies the desired alignment of each column. If \(A\) is ct::dynamic_extent, the alignment is instead supplied as a runtime argument when constructing an instance of layout_left_padded_mapping.

In a padded column-major layout, the first multi-dimensional index varies the fastest with the increasing elements in memory. Additionally, the position of each element that anchors a column is divisible by the desired alignment.

Example

Constructs a \(2 \times 3\) ct::tensor_span over the elements \(0, 1, \ldots, 7\) where each column is aligned to \(3\).

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};

// Syntax 1
ct::layout_left_padded_mapping m1{ct::extents{4_ic, 2_ic}, 6_ic};
ct::tensor_span p1{&x[0], m1};

// Syntax 2
ct::layout_left_padded_mapping m2{ct::extents{4_ic, 2_ic}, 3_ic};
ct::tensor_span p2{&x[0], m2};
\[\begin{split}\begin{pmatrix} 0 & 6 \\ 1 & 7 \\ 2 & 8 \\ 3 & 9 \end{pmatrix}\end{split}\]
template<ct::extents_like E>
using mapping = ct::layout_left_padded_mapping<E, Alignment>

Alias template yielding the associated layout mapping type for representing a padded column-major arrangement of an array whose shape is an instance of \(E\).

cuda::tiles::layout_left_padded_mapping

template<ct::extents_like E, size_t Alignment>
struct layout_left_padded_mapping

Layout mapping representing a column-major element order whose shape is an instance of \(E\). The first dimension is padded to ensure alignment of consecutive columns. For an example, see ct::layout_left_padded.

The template parameter \(A\) specifies the desired alignment of each column. If \(A\) is ct::dynamic_extent, the alignment is instead supplied as a runtime argument when constructing an instance of layout_left_padded_mapping.

A specialization \(T\) of ct::layout_left_padded_mapping models the requirements of layout mapping. Additionally, \(T\) satisfies the following constraints if the corresponding constraint is satisfied by the underlying extents like type \(E\):

  • std::is_nothrow_move_constructible_v<T> 1

  • std::is_nothrow_move_assignable_v<T> 1

  • std::is_nothrow_swappable_v<T> 1

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

Exposition Only Members

extents_type __extents

Exposition-only member denoting this object’s layout mapping shape.

index_type __alignment;

Exposition-only member denoting the desired alignment of each column specified at runtime. This member is present only when \(A\) is ct::dynamic_extent.

using __padded_extents_t = /* see below */

Exposition-only member denoting the type of a padded extents object. Let \(E\) denote extents_type and let \(N\) be its rank.

The padded extents type is the unique ct::extents specialization \(S\) satisfying:

  1. The index type and rank of \(S\) match those of \(E\).

  2. The static extents of \(S\) match those of \(E\) with the possible exception of the dimension \(0\) when \(N \neq 0\).

  3. If \(N \neq 0\), the value \(S_0\) is

    1. ct::dynamic_extent if either \(A\) or \(E_0\) is ct::dynamic_extent.

    2. The smallest multiple of \(A\) that is not less than \(E_0\) otherwise. If this value is not representable in the type size_t, the behavior is undefined when constructing an instance of this object.

__tile__ __host__ __device__ constexpr __padded_extents_t __padded_extents() const noexcept

Exposition only function that yields a padded extents object based on the values of __extents and __alignment when present. Let \(e\) denote __extents and let \(N\) be its rank. Let \(a\) denote __alignment when present and A otherwise. The returned value is the unique instance \(s\) of __padded_extents_t satisfying:

  1. The extent values of \(s\) match those of \(e\) with the possible exception of of dimension \(0\) when \(N \neq 0\).

  2. When \(N \neq 0\), the value \(s_0\) is the smallest multiple of \(a\) that is not less than \(e_0\). If this value is not representable in the type index_type, the behavior of this function is undefined.

Member Aliases

using extents_type = E
using rank_type = typename extents_type::rank_type
using index_type = typename extents_type::index_type
using layout_type = ct::layout_left_padded<Alignment>

Constructor

template<typename T>
requires is-convertible-v<T, index_type>
__tile__ __host__ __device__ layout_left_padded_mapping(extents_type const &other, T alignment) noexcept

Constructs an instance of this layout mapping by direct-list-initializing __extents with other. When __alignment is present, it is direct-initialized with alignment.

is_always_strided

__tile__ __host__ __device__ static constexpr bool is_always_strided() noexcept

Returns true.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept

Returns the static stride at dimension \(i\). Behaves as if by invoking

ct::layout_left_mapping<__padded_extents_t>::static_stride(i)

stride

__tile__ __host__ __device__ index_type stride(rank_type i) const noexcept

Returns the stride at dimension \(i\).

Behaves as if by invoking

ct::layout_left_mapping{__padded_extents()}.stride(i)

extents

__tile__ __host__ __device__ extents_type const &extents() const noexcept

Yields the glvalue __extents.

Deduction Guides

template<typename E, typename A>
layout_left_padded_mapping(E const&, A) -> layout_left_padded_mapping<E, extent-constant-or-dynamic<A>>;

Enables class template argument deduction from an ct::integral_constant or integral alignment argument.

cuda::tiles::layout_strided

template<ct::extents_like S>
struct layout_strided

Layout policy representing a strided arrangement of elements. The stride values are given by an instance of \(S\) when constructing the corresponding layout_strided_mapping and may include statically known or dynamically known stride values.

Example

The code below creates a \(2 \times 3\) tensor span over a sparse set of elements in the underlying array. Increment the first dimension moves \(6\) elements in memory while incrementing the second dimension moves \(2\) elements in memory.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};

ct::layout_strided_mapping m{ct::extents{2_ic, 3_ic}, ct::extents{6_ic, 2_ic}};
ct::tensor_span p{&x[0], m};
\[\begin{split}\begin{pmatrix} 0 & 2 & 4 \\ 6 & 8 & 10 \end{pmatrix}\end{split}\]
template<ct::extents_like E>
using mapping = ct::layout_strided_mapping<E, Strides>

Alias template yielding the associated layout mapping type for representing a strided layout for an array whose shape is an instance of \(E\).

cuda::tiles::layout_strided_mapping

template<ct::extents_like E, ct::extents_like S>
requires /* atomic constraint */
struct layout_strided_mapping

Layout mapping representing a layout with arbitrary strides for a shape that is an instance of \(E\). The strides for each dimension may be known either statically or at runtime and are determined by an instance of \(S\) provided at construction.

A specialization \(T\) of ct::layout_strided_mapping models the requirements of layout mapping. Additionally, \(T\) satisfies the following constraints if the corresponding constraint is satisfied by the underlying extents like type \(E\):

  • std::is_nothrow_move_constructible_v<T> 1

  • std::is_nothrow_move_assignable_v<T> 1

  • std::is_nothrow_swappable_v<T> 1

The program is ill-formed if \(E\) or \(S\) is cv-qualified.

The atomic constraint validates that \(E\) and \(S\) have the same rank and index type.

Exposition Only Members

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

extents_type __extents

Exposition-only member denoting this object’s layout mapping shape.

S __strides

Exposition-only member denoting the strides of this layout mapping.

Types

using extents_type = E
using rank_type = typename extents_type::rank_type
using index_type = typename extents_type::index_type
using layout_type = ct::layout_strided<S>

layout_strided_mapping

__tile__ __host__ __device__ layout_strided_mapping(extents_type const &other, S const &strides) noexcept

Constructs an instance of this layout mapping by performing direct-list-initializing of __extents with other and direct-list-initialization of __strides with strides.

is_always_strided

__tile__ __host__ __device__ static constexpr bool is_always_strided() noexcept

Returns true.

static_stride

__tile__ __host__ __device__ static constexpr size_t static_stride(rank_type i) noexcept

Returns the static stride at dimension \(i\) as if by invoking S::static_extent(i). The behavior is undefined if \(i\) is not in the range \([0, N)\) where \(N\) is the rank of E.

stride

__tile__ __host__ __device__ index_type stride(rank_type i) const noexcept

Returns the stride at dimension \(i\) as if by invoking __strides.extent(i). The behavior is undefined if \(i\) is not in the range \([0, N)\) where \(N\) is the rank of E.

extents

__tile__ __host__ __device__ extents_type const &extents() const noexcept

Yields the glvalue __extents.

Layout Comparison

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

Indicates whether lhs and rhs are layout mapping equivalent.

Note

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

cuda::tiles::layout_mapping_static_stride

template<typename T>
struct layout_mapping_static_stride

Traits class which may be specialized to provide the static stride values when implementing a custom layout mapping type. The primary template’s implementation delegates to a T::static_stride member function if available.

__tile__ __host__ __device__ constexpr size_t operator()(typename T::rank_type idx) noexcept;

Yields the value T::static_stride(idx). This function participates in overload resolution if the expression T::static_stride(idx) is well formed and yields a prvalue of type size_t.

Footnotes

1(1,2,3,4,5,6,7,8,9,10,11,12,13,14,15)

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

2

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