Strided View CUDA 13.4

A ct::strided_view is a wrapper around a tensor span like type that divides it into statically sized chunks which may be loaded and stored from memory. The spacing between chunks is determined by a compile time striding factor. Depending on the strides, there may be overlap or gaps between chunks.

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

Example

The following code subdivides the provided $5 \times 10$ tensor span into chunks of size $2 \times 2$. Chunks are spaced $3$ elements apart along the rows and $4$ elements apart along the columns. The chunk at index $(0, 1)$ is loaded.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[4][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},
};

ct::tensor_span t{&x[0][0], ct::extents{5_ic, 10_ic}};
ct::strided_view p{t, ct::shape{2_ic, 2_ic}, ct::strides{3_ic, 4_ic}};
auto r = p.load(0, 1);

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

Example

The following code subdivides the provided $4 \times 4$ tensor span into chunks of size $2 \times 2$. Chunks are spaced $1$ element apart in each dimension causing overlap. The chunk at index $(1, 1)$ is loaded.

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

ct::tensor_span t{&x[0][0], ct::extents{4_ic, 4_ic}};
ct::strided_view p{t, ct::shape{2_ic, 2_ic}, ct::strides{1_ic, 1_ic}};
auto r = p.load(1, 1);

\[\begin{split}\begin{array}{c|c|c|c|c} 0 & 1 & 2 & 3 \\ \hline 4 & 5 & 6 & 7 \\ \hline 8 & 9 & 10 & 11 \\ \hline 12 & 13 & 14 & 15 \end{array} \rightarrow \begin{pmatrix} 5 & 6 \\ 9 & 10 \end{pmatrix}\end{split}\]

cuda::tiles::strided_view

template<ct::tensor_span_like Span, ct::tile_shape Shape, ct::strides_like Strides> requires /* atomic constraint */ struct strided_view

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

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

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 ranks of Span, Shape, and Strides are identical.
The stride type of Strides is the same as the index type of the shape of Span.
All stride values of Strides are static and greater than $0$.

The program is ill-formed if any of Span, Shape, or Strides 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 Members

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.

Strides __strides: Exposition only member representing the strides between chunks of the view.

Exposition Only Definitions

strided view mapping

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

Let let $S$ denote Strides and let $N$ be its rank.

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 view_shape_type.

The mapping is a new index

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

Note

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

strided view index space

The index space of the strided view is the set of chunk indices which correspond to in bounds or partially out of bound chunks.

Let $S$ denote Strides, let $e$ denote the object __span.extents() and let $N$ be their rank. The index space is the set of non-negative indices $I = (i_0, i_1, \ldots, i_{N-1})$ satisfying

\[i_k \cdot S_k < e_k \quad 0 \leq k < N\]

Constructor

__tile__ __host__ __device__ strided_view(span_type span, view_shape_type, Strides strides) noexcept

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

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 a tile from a chunk specified by the indices Idx.

Example

In the following example, a $4 \times 11$ tensor span is partitioned into $2 \times 4$ chunks with a stride length of $1 \times 1$. Chunk index $(1, 8)$ is loaded using the masked load API. The part of the chunk which is out of bounds is filled with NaN values.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
float x[4][11] = {
  {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}
};

ct::tensor_span t{&x[0][0], ct::extents{4_ic, 11_ic}};
ct::strided_view p{t, ct::shape{2_ic, 4_ic}, ct::strides{1_ic, 1_ic}};

auto r = p.load_masked(ct::view_padding_nan_t{}, 1, 8);

\[\begin{split}\begin{pmatrix} \begin{matrix} 0\hphantom{0} & 1\hphantom{0} & 2\hphantom{0} & 3\hphantom{0} & 4\hphantom{0} & 5\hphantom{0} & 6\hphantom{0} & 7\hphantom{0} \end{matrix} & \begin{matrix} 8\hphantom{0} & 9\hphantom{0} & 10 & \text{x} \end{matrix} \\ \begin{matrix} 11 & 12 & 13 & 14 & 15 & 16 & 17 & 18 \\ 22 & 23 & 24 & 25 & 26 & 27 & 28 & 29 \end{matrix} & \fbox{$\begin{matrix} 19 & 20 & 21 & \text{x} \\ 30 & 31 & 32 & \text{x} \end{matrix}$} \\ \begin{matrix} 33 & 34 & 35 & 36 & 37 & 38 & 39 & 40 \end{matrix} & \begin{matrix} 41 & 42 & 43 & \text{x} \end{matrix} \end{pmatrix} \rightarrow \begin{pmatrix} 19 & 20 & 21 & \text{NaN} \\ 30 & 31 & 32 & \text{NaN} \end{pmatrix}\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type.

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

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

where $p$ is the strided 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 following hold:

a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
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.
Each element of the Idx pack 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

The indices Idx specify which chunk should be loaded. A fully out of bounds chunk always yields undefined behavior. A partially out of bounds chunk yields undefined behavior for the non-masked variants.

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 partition specified by indices Idx.

Example

In the following example, a $4 \times 8$ tensor span is split into $2 \times 2$ chunks with strides $2 \times 3$ and the bottom right chunk is updated with a new value.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[4][8] = {
  {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}
};

ct::tensor_span t{&x[0][0], ct::extents{4_ic, 8_ic}};
ct::strided_view p{t, ct::shape{2_ic, 2_ic}, ct::strides{2_ic, 3_ic}};

auto a = 60  + 10 * ct::iota<ct::tile<int, ct::shape<2, 2>>>();
p.store(a, 1, 2);

\[\begin{split}\begin{pmatrix} 60 & 70 \\ 80 & 90 \end{pmatrix} \rightarrow \begin{pmatrix} \fbox{$\begin{matrix} 0 & 1 \\ 8 & 9 \end{matrix}$} & \begin{matrix} 2 \\ 10 \end{matrix} & \fbox{$\begin{matrix} 3 & 4 \\ 11 & 12 \end{matrix}$} & \begin{matrix} 5 \\ 13 \end{matrix} & \fbox{$\begin{matrix} 6 & 7 \\ 14 & 15 \end{matrix}$} \\ \fbox{$\begin{matrix} 8 & 9 \\ 16 & 17 \end{matrix}$} & \begin{matrix} 10 \\ 18 \end{matrix} & \fbox{$\begin{matrix} 11 & 12 \\ 19 & 20 \end{matrix}$} & \begin{matrix} 13 \\ 21 \end{matrix} & \fbox{$\begin{matrix} 60 & 70 \\ 80 & 90 \end{matrix}$} \end{pmatrix}\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_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 strided 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:

a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates write memory operations for the values at the addresses $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.
Each element of the Idx pack scalar convertible to index_type
When specified, the values Order, and Scope are all enumerators of their respective types.

Atomic And

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

Performs elementwise atomic bitwise AND operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic AND operation.

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

static int x[] = { 0x1, 0x2, 0b10101010, 0b01010101 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<int, ct::shape<2>>>(0b00001111);
p.atomic_and(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 0b10101010 & 0b01010101 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 0b00001010 & 0b00000101 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

Let $M$ denote the memory location

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

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

Let $v$ denote the value $a(J)$, and $k$ denote the value in memory location $M$. The bitwise AND between $v$ and $k$ is computed as if by invoking ct::operator&(k, v) and the result is stored to memory location $M$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
the tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is a $32$ or $64$ bit integral scalar type

Note

Unlike the non-member atomic_and, this function does not return any value (void return type).

Atomic Or

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

Performs elementwise atomic bitwise OR operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic OR operation.

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

static int x[] = { 0x1, 0x2, 0b10101010, 0b01010101 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<int, ct::shape<2>>>(0b00001111);
p.atomic_or(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 0b10101010 & 0b01010101 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 0b10101111 & 0b01011111 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

Let $M$ denote the memory location

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

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

Let $v$ denote the value $a(J)$, and $k$ denote the value in memory location $M$. The bitwise OR between $v$ and $k$ is computed as if by invoking ct::operator|(k, v) and the result is stored to memory location $M$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is a $32$ or $64$ bit integral scalar type

Note

Unlike the non-member atomic_or, this function does not return any value (void return type).

Atomic Xor

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

Performs elementwise atomic bitwise XOR operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic XOR operation.

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

static int x[] = { 0x1, 0x2, 0b10101010, 0b01010101 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<int, ct::shape<2>>>(0b11111111);
p.atomic_xor(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 0b10101010 & 0b01010101 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 0b01010101 & 0b10101010 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

Let $M$ denote the memory location

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

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

Let $v$ denote the value $a(J)$, and $k$ denote the value in memory location $M$. The bitwise XOR between $v$ and $k$ is computed as if by invoking ct::operator^(k, v) and the result is stored to memory location $M$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is a $32$ or $64$ bit integral scalar type

Note

Unlike the non-member atomic_xor, this function does not return any value (void return type).

Atomic Max

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

Tparam-line-spec

Performs elementwise atomic maximum operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic maximum operation.

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

static unsigned int x[] = { 1, 2, 3, 6 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2>>>(4);
p.atomic_max(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 3 & 6 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 4 & 6 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

Let $M$ denote the memory location

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

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

Let $v$ denote the value $a(J)$, and $k$ denote the value in memory location $M$. The maximum between $v$ and $k$ is computed as if by invoking ct::max(k, v) and the result is stored to memory location $M$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is a $32$ or $64$ bit integral scalar type

Note

Unlike the non-member atomic_max, this function does not return any value (void return type).

Atomic Min

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

Performs elementwise atomic minimum operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic minimum operation.

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

static unsigned int x[] = { 1, 2, 3, 6 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2>>>(4);
p.atomic_min(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 3 & 6 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 3 & 4 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

Let $M$ denote the memory location

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

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

Let $v$ denote the value $a(J)$, and $k$ denote the value in memory location $M$. The minimum between $v$ and $k$ is computed as if by invoking ct::min(k, v) and the result is stored to memory location $M$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is a $32$ or $64$ bit integral scalar type

Note

Unlike the non-member atomic_min, this function does not return any value (void return type).

Atomic Add

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

Performs elementwise atomic addition operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic add operation.

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

static unsigned int x[] = { 1, 2, 3, 4 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2>>>(10);
p.atomic_add(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 3 & 4 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 13 & 14 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

The value $k$ in memory location $t(p(I, J))$ is loaded, where $p$ is the strided view mapping of this object and $t$ is the tensor span function of __span. Let $v$ denote the value $a(J)$.

The sum between $v$ and $k$ is computed as if by invoking

If v is a integral scalar: ct::add(k, v). If this produces undefined behavior, the behavior of the operation as a whole is undefined.
If v is a basic floating point scalar:

ct::add<ct::rounding_mode::round_ties_to_even, SubMode>(k, v)

where the value of SubMode is not specified.

The result of the computation is stored back in the memory location $t(p(I, J))$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is one of

A $32$ or $64$ bit integral scalar type

double, float, or __half

__nv_bfloat16 when targeting architecture >= sm_90.

Note

Unlike the non-member atomic_add, this function does not return any value (void return type).

Atomic Sub

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

Performs elementwise atomic subtraction operation of the tile converted operand a and the chunk specified by the indices Idx and stores the result.

Example

In the following example, a $1D$ tensor span is split into $2$ chunks and the 2nd chunk is updated with an atomic subtraction operation.

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

static unsigned int x[] = { 1, 2, 3, 4 };

ct::tensor_span t{&x[0], ct::extents{4_ic}};
ct::strided_view p{t, ct::shape{2_ic}, ct::strides{2_ic}};

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2>>>(3);
p.atomic_sub(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 1);

\[\begin{split}\begin{pmatrix} 1 & 2 & 3 & 4 \\ \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc} 1 & 2 & 0 & 1 \\ \end{array} \right)\end{split}\]

Let $I$ denote the values specified by the pack Idx and let $J$ be an index in the index space of view_tile_type. Let $a$ denote the value of a after tile conversion to view_tile_type.

The value $k$ in memory location $t(p(I, J))$ is loaded, where $p$ is the strided view mapping of this object and $t$ is the tensor span function of __span. Let $v$ denote the value $a(J)$.

The difference between $v$ and $k$ is computed as if by invoking

If v is a integral scalar: ct::add(k, -v). If this produces undefined behavior, the behavior of the operation as a whole is undefined.

Note

The behavior is undefined if $v$ is a signed integral type and overflow occurs in the unary negation even if the difference between $k$ and $v$ is representable in the target integer type.
If v is a basic floating point scalar:

ct::sub<ct::rounding_mode::round_ties_to_even, SubMode>(k, v)

where the value of SubMode is not specified.

The result of the computation is stored back in the memory location $t(p(I, J))$.

The behavior is undefined if any of the following hold:

the value $p(I, J)$ is not in index space of __span
a value of the Idx pack is not representable in index_type
$I$ is outside the strided view’s index space
The tensor span function is not injective

An invocation generates strong read-write memory operations for the values at the addresses $t(p(I, J))$ for each $J \in \mathbb{J}$s.

The atomic constraint requires that:

The size of the Idx pack matches the rank of view_shape_type
Each element of the Idx pack scalar convertible to index_type
Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block
The element type of view_tile_type is one of

A $32$ or $64$ bit integral scalar type

double, float, or __half

__nv_bfloat16 when targeting architecture >= sm_90.

Note

Unlike the non-member atomic_sub, this function does not return any value (void return type).

span

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