Partition View

A ct::partition_view is a wrapper around a tensor span like type that divides it into statically sized tiles which may be loaded and stored from memory.

Example

The following code subdivides the provided \(4 \times 8\) tensor span into partitions of size \(2 \times 2\). The partition identified by index \((1, 2)\) is loaded.

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::partition_view p{t, ct::shape{2_ic, 2_ic}};
auto r = p.load(1, 2);
\[\begin{split}\left( \begin{array}{cc|cc|cc|cc} 0 & 1 & 2 & 3 & 4 & 5 & 6 & 7 \\ 8 & 9 & 10 & 11 & 12 & 13 & 14 & 15 \\ \hline 16 & 17 & 18 & 19 & 20 & 21 & 22 & 23 \\ 24 & 25 & 26 & 27 & 28 & 29 & 30 & 31 \end{array} \right) \rightarrow \begin{pmatrix} 20 & 21 \\ 28 & 29 \end{pmatrix}\end{split}\]

cuda::tiles::partition_view

template<ct::tensor_span_like Span, ct::tile_shape Shape>
requires (Shape::rank() == Span::extents_type::rank())
struct partition_view

All specializations \(T\) of ct::partition_view model std::copyable.

\(T\) satisfies the following constraints if they are satisfied by Span:

  1. std::is_nothrow_move_constructible_v<T>

  2. std::is_nothrow_move_assignable_v<T>

  3. std::is_nothrow_swappable_v<T>

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

Aliases

using span_type = Span

The type of the wrapped tensor span.

using view_shape_type = Shape

The shape of the tile that will be loaded or stored.

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

span_type __span

Exposition only member containing the wrapped tensor span object. This member exists for the purpose of defining the behavior of the implicitly defined special member functions.

Exposition Only Definitions

partition view mapping

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

Let \(S\) denote view_shape_type and let \(N\) be its rank.

Let \(I = (i_0, i_1, \ldots, i_{N-1})\) be a partition 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 \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.

partition view index space

The index space of the partition view is the set of partition indices which correspond to in bounds or partially out of bound partitions.

Let \(S\) denote view_shape_type, 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__ partition_view(span_type span, view_shape_type) noexcept

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

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

Example

In the following example, a \(4 \times 11\) tensor span is partitioned into \(2 \times 4\) chunks. A partially out of bounds chunk is loaded using NaN padding 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::partition_view p{t, ct::shape{2_ic, 4_ic}};

auto r = p.load_masked(ct::view_padding_nan_t{}, 0, 2);
\[\begin{split}\left( \begin{array}{cccc|cccc|ccc} 0 & 1 & 2 & 3 & 4 & 5 & 6 & 7 & 8 & 9 & 10 \\ 11 & 12 & 13 & 14 & 15 & 16 & 17 & 18 & 19 & 20 & 21 \\ \hline 22 & 23 & 24 & 25 & 26 & 27 & 28 & 29 & 30 & 31 & 32 \\ 33 & 34 & 35 & 36 & 37 & 38 & 39 & 40 & 41 & 42 & 43 \end{array} \right) \rightarrow \begin{pmatrix} 8 & 9 & 10 & \text{NaN} \\ 19 & 20 & 21 & \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 partition 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:

  1. a value of the Idx pack is not representable in index_type

  2. \(I\) is outside the partition view’s index space

  3. The tensor span function is not injective

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

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

  1. The size of the Idx pack matches the rank of view_shape_type.

  2. Each element of the Idx pack scalar convertible to index_type

  3. If Pad is present and its value is not zero view padding, then element_type is a basic floating point scalar.

  4. When specified, the values Pad, Order, and Scope are all enumerators of their respective types.

Note

The indices Idx specify which partition should be loaded. A fully out of bounds partition always yields undefined behavior. A partially out of bounds partition 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 partitioned into \(2 \times 2\) chunks and the bottom right partition 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::partition_view p{t, ct::shape{2_ic, 2_ic}};

auto a = 100 * ct::iota<ct::tile<int, ct::shape<2, 2>>>();
p.store(a, 1, 3);
\[\begin{split}\begin{pmatrix} 0 & 100 \\ 200 & 300 \end{pmatrix} \rightarrow \left( \begin{array}{cc|cc|cc|cc} 0 & 1 & 2 & 3 & 4 & 5 & 6 & 7 \\ 8 & 9 & 10 & 11 & 12 & 13 & 14 & 15 \\ \hline 16 & 17 & 18 & 19 & 20 & 21 & 0 & 100 \\ 24 & 25 & 26 & 27 & 28 & 29 & 200 & 300 \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.

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 partition 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:

  1. a value of the Idx pack is not representable in index_type

  2. \(I\) is outside the partition view’s index space

  3. 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 \(J \in \mathbb{J}\) which is not masked.

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

  1. The size of the Idx pack matches the rank of view_shape_type.

  2. Each element of the Idx pack scalar convertible to index_type

  3. When specified, the values Order, and Scope are all enumerators of their respective types.

Atomic And CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is a \(32\) or \(64\) bit integral scalar type

This API was introduced in CUDA 13.4.

Note

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

Atomic Or CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is a \(32\) or \(64\) bit integral scalar type

This API was introduced in CUDA 13.4.

Note

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

Atomic Xor CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is a \(32\) or \(64\) bit integral scalar type

This API was introduced in CUDA 13.4.

Note

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

Atomic Max CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is a \(32\) or \(64\) bit integral scalar type

This API was introduced in CUDA 13.4.

Note

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

Atomic Min CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is a \(32\) or \(64\) bit integral scalar type

This API was introduced in CUDA 13.4.

Note

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

Atomic Add CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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

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

  2. 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is one of

  1. A \(32\) or \(64\) bit integral scalar type

  2. double, float, or __half

  3. __nv_bfloat16 when targeting architecture >= sm_90.

This API was introduced in CUDA 13.4.

Note

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

Atomic Sub CUDA 13.4

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 partition specified by the indices Idx and stores the result.

Example

In the following example, a \(1D\) tensor span is partitioned into \(2\) chunks and the 2nd partition 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::partition_view p{t, ct::shape{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 partition 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

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

  2. 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:

  1. the value \(p(I, J)\) is not in index space of __span

  2. a value of the Idx pack is not representable in index_type

  3. \(I\) is outside the partition view’s index space

  4. 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:

  1. The size of the Idx pack matches the rank of view_shape_type

  2. Each element of the Idx pack scalar convertible to index_type

  3. Order is ct::memory_order::relaxed and Scope is ct::thread_scope::device or ct::thread_scope::block

  4. The element type of view_tile_type is one of

  1. A \(32\) or \(64\) bit integral scalar type

  2. double, float, or __half

  3. __nv_bfloat16 when targeting architecture >= sm_90.

This API was introduced in CUDA 13.4.

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.