Gather Scatter View CUDA 13.4

A ct::gather_scatter_view is a wrapper around a tensor span like type that supports accessing non-adjacent chunks of an array along a gather-scatter dimension.

When indexing into the view, the gather-scatter dimension accepts a tile of indices indicating the sections along that dimension which should be accessed. The remaining dimensions accept scalar indices that specify the starting position of each chunk. A load produces a tile which is the concatenation of the loaded chunks in the order specified by the gather-scatter indices.

The ct::gather_scatter_view type was introduced in CUDA 13.4.

Example

In the following example, non-adjacent sections along dimension $0$ of an $8 \times 10$ matrix are loaded and concatenated. The ct::shape{4_ic, 4_ic} argument of ct::gather_scatter_view specifies the shape of the resulting load operation. Each chunk starts at column $3$ as specified by the second argument to the load API. Row $5$ is loaded twice and its data is duplicated in the result tile.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
using i32x4 = ct::tile<int, ct::shape<4>>;

int x[8][10] = {
  {0,  1,  2,  3,  4,  5,  6,  7,  8,  9},
  {10, 11, 12, 13, 14, 15, 16, 17, 18, 19},
  {20, 21, 22, 23, 24, 25, 26, 27, 28, 29},
  {30, 31, 32, 33, 34, 35, 36, 37, 38, 39},
  {40, 41, 42, 43, 44, 45, 46, 47, 48, 49},
  {50, 51, 52, 53, 54, 55, 56, 57, 58, 59},
  {60, 61, 62, 63, 64, 65, 66, 67, 68, 69},
  {70, 71, 72, 73, 74, 75, 76, 77, 78, 79}
};

ct::tensor_span t{&x[0][0], ct::extents{8_ic, 10_ic}};
ct::gather_scatter_view v{t, /*result shape*/ct::shape{4_ic, 4_ic}, /*gather dim*/0_ic};

int idxData[4] = {5, 0, 3, 5};
auto idx = ct::load(&idxData[0] + ct::iota<i32x4>());

auto r = v.load(idx, 3);

\[ \begin{align}\begin{aligned}\begin{pmatrix}\\\begin{split}\begin{matrix} 0\hphantom{0} & 1\hphantom{0} & 2\hphantom{0} \end{matrix} & \quad & \fbox{$\begin{matrix} 3\hphantom{0} & 4\hphantom{0} & 5\hphantom{0} & 6\hphantom{0} \end{matrix}$} & \quad & \begin{matrix} 7\hphantom{0} & 8\hphantom{0} & 9\hphantom{0} \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 10 & 11 & 12 \end{matrix} & \quad & \begin{matrix} 13 & 14 & 15 & 16 \end{matrix} & \quad & \begin{matrix} 17 & 18 & 19 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 20 & 21 & 22 \end{matrix} & \quad & \begin{matrix} 23 & 24 & 25 & 26 \end{matrix} & \quad & \begin{matrix} 27 & 28 & 29 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 30 & 31 & 32 \end{matrix} & \quad & \fbox{$\begin{matrix} 33 & 34 & 35 & 36 \end{matrix}$} & \quad & \begin{matrix} 37 & 38 & 39 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 40 & 41 & 42 \end{matrix} & \quad & \begin{matrix} 43 & 44 & 45 & 46 \end{matrix} & \quad & \begin{matrix} 47 & 48 & 49 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 50 & 51 & 52 \end{matrix} & \quad & \fbox{$\fbox{$\begin{matrix} 53 & 54 & 55 & 56 \end{matrix}$}$} & \quad & \begin{matrix} 57 & 58 & 59 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 60 & 61 & 62 \end{matrix} & \quad & \begin{matrix} 63 & 64 & 65 & 66 \end{matrix} & \quad & \begin{matrix} 67 & 68 & 69 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 70 & 71 & 72 \end{matrix} & \quad & \begin{matrix} 73 & 74 & 75 & 76 \end{matrix} & \quad & \begin{matrix} 77 & 78 & 79 \end{matrix} \\\end{split}\\\end{pmatrix}\end{aligned}\end{align} \]

\[\begin{split}\begin{pmatrix} 53 & 54 & 55 & 56 \\ 3 & 4 & 5 & 6 \\ 33 & 34 & 35 & 36 \\ 53 & 54 & 55 & 56 \end{pmatrix}\end{split}\]

cuda::tiles::gather_scatter_view

template<ct::tensor_span_like Span, ct::tile_shape Shape, size_t Dim> requires /* atomic constraint */ struct gather_scatter_view

A ct::gather_scatter_view where Span indicates the tensor span type wrapped by this object, Shape indicates the shape of the tile that will be loaded or stored and Dim indicates the dimension along which non-adjacent chunks may be accessed.

All specializations $T$ of ct::gather_scatter_view model std::copyable.

$T$ satisfies the following constraints if they are satisfied by Span:

std::is_nothrow_move_constructible_v<T>
std::is_nothrow_move_assignable_v<T>
std::is_nothrow_swappable_v<T>

The atomic constraint validates that:

The rank of Span and the rank of Shape are the same.
The gather-scatter dimension Dim is less than the rank of Shape.

The program is ill-formed if Span or Shape is cv-qualified.

Aliases

using span_type = Span 
using view_shape_type = Shape 
using element_type = typename span_type::element_type
using value_type = typename span_type::value_type
using index_type = typename span_type::index_type
using view_tile_type = ct::tile<value_type, view_shape_type>

Exposition Only Member

For the purposes of defining the implicitly generated special member functions, this type contains the following member variables.

span_type __span: Exposition only member representing the wrapped tensor span object.

Exposition Only Definitions

gather scatter view mapping

The gather scatter view mapping is a function that associates an index of a chunk to a potential set of indices of the underlying tensor span which may be loaded or stored.

Let $T$ denote the view_shape_type and $N$ its rank. The chunk shape $S$ is a rank $N$ shape determined by:

\[\begin{split}S_i = \begin{cases} T_i & i \neq d \\ 1 & i = d \end{cases}\end{split}\]

where $d$ denotes the gather-scatter dimension Dim.

Let $I = (i_0, i_1, \ldots, i_{N-1})$ be a chunk index and $J = (j_0, j_1, \ldots, j_{N-1})$ be an index in the index space of $S$.

The mapping is a new index

\[p(I, J) = (i_0 + j_0, \ldots, i_{N-1} + j_{N-1})\]

Note

The result of the mapping might not be an element in the tensor span’s index space.

gather scatter view index space

The index space of the gather scatter view is the set of chunk indices which correspond to chunks that are partially or fully in-bounds or are out of bounds only along the gather-scatter dimension.

Let $e$ denote the object __span.extents(), let $N$ be its rank, and let $d$ denote the gather-scatter dimension Dim. The index space is the set of indices $I = (i_0, i_1, \ldots, i_{N-1})$ satisfying

\[\begin{split}\begin{cases} 0 \leq i_k < e_k & \quad k \neq d \\ 0 \leq i_k & \quad k = d \end{cases}\end{split}\]

Constructor

template<typename T> requires /* atomic constraint */ __tile__ __host__ __device__ gather_scatter_view(span_type span, view_shape_type, T dim) noexcept

Constructs this object by direct-list-initializing __span from span.

The behavior is undefined if dim, after conversion to size_t, does not equal Dim.

The atomic constraint validates that $T$ is convertible to size_t.

Loads

template< typename ...Idx > requires /* atomic constraint */ __tile__ view_tile_type load(Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), typename ...Idx > requires /* atomic constraint */ __tile__ view_tile_type load_masked(Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), typename ...Idx > requires /* atomic constraint */ __tile__ view_tile_type load_masked(ct::view_padding_constant<Pad>, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load(Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load(ct::memory_order_constant<Order>, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load(ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load_masked(Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load_masked(ct::view_padding_constant<Pad>, Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load_masked(ct::view_padding_constant<Pad>, ct::memory_order_constant<Order>, Idx...) const noexcept;

template< ct::view_padding Pad = ct::default_view_padding(), ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), typename ...Idx > requires ct::read_memory_order<Order> && /* atomic constraint */ __tile__ view_tile_type atomic_load_masked(ct::view_padding_constant<Pad>, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx...) const noexcept;

Loads one or more chunks specified by a tile of indices along the gather dimension Dim and the scalar indices along the remaining dimensions.

Example

In the following example, a $4 \times 4$ tile is loaded by accessing non-adjacent partial columns of the matrix. The third column is fully out of bounds while the remaining columns are partially out of bounds. The out of bounds elements receive $-\infty$ values as specified by the view padding.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
using i32x4 = ct::tile<int, ct::shape<4>>;

float x[6][10] = {
  {0,  1,  2,  3,  4,  5,  6,  7,  8,  9},
  {10, 11, 12, 13, 14, 15, 16, 17, 18, 19},
  {20, 21, 22, 23, 24, 25, 26, 27, 28, 29},
  {30, 31, 32, 33, 34, 35, 36, 37, 38, 39},
  {40, 41, 42, 43, 44, 45, 46, 47, 48, 49},
  {50, 51, 52, 53, 54, 55, 56, 57, 58, 59},
};

ct::tensor_span t{&x[0][0], ct::extents{6_ic, 10_ic}};
ct::gather_scatter_view v{t, ct::shape{4_ic, 4_ic}, 1_ic};

int idxData[4] = {5, 6, 10, 1};
auto idx = ct::load(&idxData[0] + ct::iota<i32x4>());

auto padding = ct::view_padding_negative_inf_t{};
auto r = v.load_masked(padding, 3, idx);

\[\begin{split}\begin{pmatrix} \begin{matrix} 0\hphantom{0} \\ 10 \\ 20 \end{matrix} & \begin{matrix} 1\hphantom{0} \\ 11 \\ 21 \end{matrix} & \begin{matrix} 2\hphantom{0} & 3\hphantom{0} & 4\hphantom{0} \\ 12 & 13 & 14 \\ 22 & 23 & 24 \end{matrix} & \begin{matrix} 5\hphantom{0} \\ 15 \\ 25 \end{matrix} & \begin{matrix} 6\hphantom{0} \\ 16 \\ 26 \end{matrix} & \begin{matrix} 7\hphantom{0} & 8\hphantom{0} & 9\hphantom{0} \\ 17 & 18 & 19 \\ 27 & 28 & 29 \end{matrix} & \begin{matrix} \text{X} \\ \text{X} \\ \text{X} \end{matrix} \\ \begin{matrix} 30 \\ 40 \\ 50 \\ \text{X} \end{matrix} & \fbox{$\begin{matrix} 31 \\ 41 \\ 51 \\ \text{X} \end{matrix}$} & \begin{matrix} 32 & 33 & 34 \\ 42 & 43 & 44 \\ 52 & 53 & 54 \\ \text{X} & \text{X} & \text{X} \end{matrix} & \fbox{$\begin{matrix} 35 \\ 45 \\ 55 \\ \text{X} \end{matrix}$} & \fbox{$\begin{matrix} 36 \\ 46 \\ 56 \\ \text{X} \end{matrix}$} & \begin{matrix} 37 & 38 & 39 \\ 47 & 48 & 49 \\ 57 & 58 & 59 \\ \text{X} & \text{X} & \text{X} \end{matrix} & \fbox{$\begin{matrix}\text{X} \\ \text{X} \\ \text{X} \\ \text{X}\end{matrix}$} \\ \end{pmatrix}\end{split}\]

\[\begin{split}\begin{pmatrix} 35 & 36 & -\infty & 31 \\ 45 & 46 & -\infty & 41 \\ 55 & 56 & -\infty & 51 \\ -\infty & -\infty & -\infty & -\infty \\ \end{pmatrix}\end{split}\]

Define the following terms:

Let $N$ be the rank of view_shape_type.
Let $d$ denote the gather dimension Dim.
Let $g$ denote the tile argument corresponding to element $d$ of the idx pack after elementwise scalar conversion of its elements to index_type.
Let $J = (j_0, \ldots, j_{N-1})$ be an index in the index space of the load result view_tile_type.
Given index $J$, define $I = (i_0, \ldots, i_{d-1}, g(j_d), i_{d+1}, \ldots, i_{N-1})$ to be a chunk index where $i_k$ is the $k^{th}$ element of of the idx pack after scalar conversion to index_type.

The value returned by the load is a tile object $a$ satisfying

\[a(J) = t(p(I, J))\]

where $p$ is the gather scatter view mapping of this object and $t$ is the tensor span function of __span.

If the value $p(I, J)$ is not in index space of __span, the behavior depends on the selected overload:

For the non-masked overloads, the behavior is undefined.
For masked overloads, the value of $a(J)$ is the view padding specified by Pad.

The behavior is undefined if:

Any of the tile like arguments of the idx pack has an element that is not representable in index_type.
$I$ is outside the gather scatter view’s index space for any possible $J$.
The tensor span function is not injective.

An invocation generates a read memory operation at the location $t(p(I, J))$ for each $p(I, J)$ in the index space of __span.

For the atomic overloads, the memory operations are strong and have thread scope specified by Scope and memory order specified by Order

The latency and allow_tma optimization hints may appertain to direct call expressions of the above load APIs.

The atomic constraint requires that

The size of the Idx pack matches the rank of view_shape_type.
For each element of the Idx pack:
1. If the index of this element in the pack is $d$, the type of the element must be tile convertible to a tile like object whose element type is index_type and whose shape has rank $1$ and size equal to $T_d$ where $T$ denotes view_tile_shape.
2. If the index of this element in the pack is not $d$, it must be scalar convertible to index_type
If Pad is present and its value is not zero view padding, then element_type is a basic floating point scalar.
When specified, the values Pad, Order, and Scope are all enumerators of their respective types.

Note

A single invocation of this API may generate multiple read memory operations at the same memory location if the same index is specified multiple times in the gather dimension.

Stores

template< ct::tile_like Value, typename ...Idx > requires /* atomic constraint */ && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> __tile__ void store(Value a, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void atomic_store(Value a, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint*/ __tile__ void atomic_store(Value a, ct::memory_order_constant<Order>, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void atomic_store(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx...) const noexcept;

template< ct::tile_like Value, typename ...Idx > requires ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void store_masked(Value a, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void atomic_store_masked(Value a, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void atomic_store_masked(Value a, ct::memory_order_constant<Order>, Idx...) const noexcept;

template< ct::memory_order Order, ct::thread_scope Scope = ct::default_thread_scope(), ct::tile_like Value, typename ...Idx > requires ct::write_memory_order<Order> && ct::non_narrowing_tile_convertible_to<Value, view_tile_type> && ct::storeable_tensor_span<span_type> && /* atomic constraint */ __tile__ void atomic_store_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx...) const noexcept;

Stores the tile converted operand a to the chunk specified by indices Idx.

Example

In the following example, a $2 \times 4$ tile is stored onto a $6 \times 10$ array. The first row of the tile is written to row $3$ of the array and the second row of the tile is written to row $1$.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
using i32x2 = ct::tile<int, ct::shape<2>>;

int x[6][10] = {
  {0,  1,  2,  3,  4,  5,  6,  7,  8,  9},
  {10, 11, 12, 13, 14, 15, 16, 17, 18, 19},
  {20, 21, 22, 23, 24, 25, 26, 27, 28, 29},
  {30, 31, 32, 33, 34, 35, 36, 37, 38, 39},
  {40, 41, 42, 43, 44, 45, 46, 47, 48, 49},
  {50, 51, 52, 53, 54, 55, 56, 57, 58, 59},
};

int data[2][4] = {
  {70, 71, 72, 73},
  {80, 81, 82, 83}
};

auto y = ct::load(&data[0][0] + ct::iota<ct::tile<int, ct::shape<2, 4>>>());

ct::tensor_span t{&x[0][0], ct::extents{6_ic, 10_ic}};
ct::gather_scatter_view v{t, ct::shape{2_ic, 4_ic}, 0_ic};

int idxData[2] = {3, 1};
auto idx = ct::load(&idxData[0] + ct::iota<i32x2>());

v.store(y, idx, 1);

\[\begin{split}\begin{pmatrix} 70 & 71 & 72 & 73 \\ 80 & 81 & 82 & 83 \end{pmatrix}\end{split}\]

\[ \begin{align}\begin{aligned}\begin{pmatrix}\\\begin{split}\begin{matrix} 0\hphantom{0} \end{matrix} & \quad \begin{matrix} 1\hphantom{0} & 2\hphantom{0} & 3\hphantom{0} & 4\hphantom{0} \end{matrix} & \quad \begin{matrix} 5\hphantom{0} & 6\hphantom{0} & 7\hphantom{0} & 8\hphantom{0} & 9\hphantom{0} \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 10 \end{matrix} & \quad \fbox{$\begin{matrix} 80 & 81 & 82 & 83 \end{matrix}$} & \quad \begin{matrix} 15 & 16 & 17 & 18 & 19 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 20 \end{matrix} & \quad \begin{matrix} 21 & 22 & 23 & 24 \end{matrix} & \quad \begin{matrix} 25 & 26 & 27 & 28 & 29 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 30 \end{matrix} & \quad \fbox{$\begin{matrix} 70 & 71 & 72 & 73 \end{matrix}$} & \quad \begin{matrix} 35 & 36 & 37 & 38 & 39 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 40 \end{matrix} & \quad \begin{matrix} 41 & 42 & 43 & 44 \end{matrix} & \quad \begin{matrix} 45 & 46 & 47 & 48 & 49 \end{matrix} \\\end{split}\\\begin{split}\begin{matrix} 50 \end{matrix} & \quad \begin{matrix} 51 & 52 & 53 & 54 \end{matrix} & \quad \begin{matrix} 55 & 56 & 57 & 58 & 59 \end{matrix} \\\end{split}\\\end{pmatrix}\end{aligned}\end{align} \]

Define the following terms:

Let $N$ be the rank of view_shape_type.
Let $d$ denote the scatter dimension Dim.
Let $g$ denote the tile argument corresponding to element $d$ of the idx pack after elementwise scalar conversion of its elements to index_type.
Let $J = (j_0, \ldots, j_{N-1})$ be an index in the index space of the load result view_tile_type.
Given index $J$, define $I = (i_0, \ldots, i_{d-1}, g(j_d), i_{d+1}, \ldots, i_{N-1})$ to be a chunk index where $i_k$ is the $k^{th}$ element of of the idx pack after scalar conversion to index_type.

After $a$ is tile converted to view_tile_type, the value $a(J)$ is stored to the memory location

\[t(p(I, J))\]

where $p$ is the gather scatter view mapping of this object and $t$ is the tensor span function of __span.

If the value $p(I, J)$ is not in index space of __span, the behavior depends on the selected overload:

For non-masked overloads, the behavior is undefined
For masked overloads, no store occurs at that memory location

The behavior is undefined if any of the following hold:

Any of the tile like arguments of the idx pack has an element that is not representable in index_type.
$I$ is outside the gather scatter view’s index space for any possible $J$.
The tensor span function is not injective.

An invocation generates write memory operations at the location $t(p(I, J))$ for each $p(I, J)$ in the index space of __span.

For the atomic variants, the memory operations are strong and have thread scope specified by Scope and memory order specified by Order

The latency and allow_tma optimization hints may appertain to direct call expressions of the above store APIs.

The atomic constraint requires that

The size of the Idx pack matches the rank of view_shape_type.
For each element of the Idx pack:
1. If the index of this element in the pack is $d$, the type of the element must be tile convertible to a tile like object whose element type is index_type and whose shape has rank $1$ and size equal to $T_d$ where $T$ denotes tile_view_shape.
2. If the index of this element in the pack is not $d$, it must be scalar convertible to index_type
If Pad is present and its value is not zero view padding, then element_type is a basic floating point scalar.
When specified, the values Order, and Scope are all enumerators of their respective types.

Note

A single invocation of this API may generate multiple write memory operations at the same memory location if the same index is specified multiple times in the scatter dimension. This results in undefined behavior if the memory operations are not strong.

span

span_type const &span() const noexcept;: Yields the glvalue __span.

Deduction Guide

template<typename Span, typename Shape, auto Dim> gather_scatter_view(Span, Shape, ct::integral_constant<Dim>) -> gather_scatter_view<Span, Shape, size_t(Dim)>;: Enables class template argument deduction from constructor arguments.