GeezerVison

The World Wide Web, restored to 800 pixels of glory.

GeezerVison is showing docs.nvidia.com — Partition View — CUDA Tile C++ API Reference 13.5 documentation

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:

std::is_nothrow_move_constructible_v<T>
std::is_nothrow_move_assignable_v<T>
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::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;

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:

a value of the Idx pack is not representable in index_type
\(I\) is outside the partition 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
When Pad is present, it must be compatible with element_type.
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::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;

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:

a value of the Idx pack is not representable in index_type
\(I\) is outside the partition view’s index space
The tensor span function is not injective

An invocation generates a write memory operation 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.
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 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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic AND on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static int x[4][3] = {
  { 1, 2, 0b10101010 },
  { 3, 4, 0b11001100 },
  { 5, 6, 0b11110000 },
  { 7, 8, 0b00110011 }
};

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

auto a =  ct::full<ct::tile<int, ct::shape<2, 2>>>(0b00001111);

// Using the unmasked atomic_and here would be UB (see the Undefined Behavior Annex).
p.atomic_and_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 0b10101010 & \text{x} \\ 3 & 4 & 0b11001100 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 0b00001010 & \text{x} \\ 3 & 4 & 0b00001100 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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

This API was introduced in CUDA 13.4.

Note

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

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic OR on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static int x[4][3] = {
  { 1, 2, 0b10101010 },
  { 3, 4, 0b11001100 },
  { 5, 6, 0b11110000 },
  { 7, 8, 0b00110011 }
};

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

auto a =  ct::full<ct::tile<int, ct::shape<2, 2>>>(0b00001111);

// Using the unmasked atomic_or here would be UB (see the Undefined Behavior Annex).
p.atomic_or_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 0b10101010 & \text{x} \\ 3 & 4 & 0b11001100 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 0b10101111 & \text{x} \\ 3 & 4 & 0b11001111 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic XOR on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static int x[4][3] = {
  { 1, 2, 0b10101010 },
  { 3, 4, 0b11001100 },
  { 5, 6, 0b11110000 },
  { 7, 8, 0b00110011 }
};

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

auto a =  ct::full<ct::tile<int, ct::shape<2, 2>>>(0b11111111);

// Using the unmasked atomic_xor here would be UB (see the Undefined Behavior Annex).
p.atomic_xor_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 0b10101010 & \text{x} \\ 3 & 4 & 0b11001100 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 0b01010101 & \text{x} \\ 3 & 4 & 0b00110011 & \text{x} \\ \hline 5 & 6 & 0b11110000 & \text{x} \\ 7 & 8 & 0b00110011 & \text{x} \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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic maximum on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static unsigned int x[4][3] = {
  { 1, 2, 3 },
  { 3, 4, 6 },
  { 5, 6, 2 },
  { 7, 8, 8 }
};

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

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2, 2>>>(4);

// Using the unmasked atomic_max here would be UB (see the Undefined Behavior Annex).
p.atomic_max_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 3 & \text{x} \\ 3 & 4 & 6 & \text{x} \\ \hline 5 & 6 & 2 & \text{x} \\ 7 & 8 & 8 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 4 & \text{x} \\ 3 & 4 & 6 & \text{x} \\ \hline 5 & 6 & 2 & \text{x} \\ 7 & 8 & 8 & \text{x} \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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic minimum on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static unsigned int x[4][3] = {
  { 1, 2, 3 },
  { 3, 4, 6 },
  { 5, 6, 2 },
  { 7, 8, 8 }
};

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

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2, 2>>>(4);

// Using the unmasked atomic_min here would be UB (see the Undefined Behavior Annex).
p.atomic_min_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 3 & \text{x} \\ 3 & 4 & 6 & \text{x} \\ \hline 5 & 6 & 2 & \text{x} \\ 7 & 8 & 8 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 3 & \text{x} \\ 3 & 4 & 4 & \text{x} \\ \hline 5 & 6 & 2 & \text{x} \\ 7 & 8 & 8 & \text{x} \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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic addition on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static unsigned int x[4][3] = {
  { 1, 2, 1 },
  { 3, 4, 2 },
  { 5, 6, 3 },
  { 7, 8, 4 }
};

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

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2, 2>>>(10);

// Using the unmasked atomic_add here would be UB (see the Undefined Behavior Annex).
p.atomic_add_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 1 & \text{x} \\ 3 & 4 & 2 & \text{x} \\ \hline 5 & 6 & 3 & \text{x} \\ 7 & 8 & 4 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 11 & \text{x} \\ 3 & 4 & 12 & \text{x} \\ \hline 5 & 6 & 3 & \text{x} \\ 7 & 8 & 4 & \text{x} \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

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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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.

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;

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_masked(Value a, ct::memory_order_constant<Order>, ct::thread_scope_constant<Scope>, Idx... idx) const noexcept; CUDA 13.5

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}\]

Example

In the following example, a \(4 \times 3\) tensor span is partitioned into \(2 \times 2\) partitions. Partitions \((0, 1)\) and \((1, 1)\) are partially out of bounds. A masked atomic subtraction on partition \((0, 1)\) updates the in-bounds elements and leaves the out of bounds elements (denoted \(\text{x}\)) unmodified.

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

static unsigned int x[4][3] = {
  { 1, 2, 5 },
  { 3, 4, 8 },
  { 5, 6, 6 },
  { 7, 8, 9 }
};

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

auto a =  ct::full<ct::tile<unsigned int, ct::shape<2, 2>>>(3);

// Using the unmasked atomic_sub here would be UB (see the Undefined Behavior Annex).
p.atomic_sub_masked(a, ct::memory_order_relaxed_t{}, ct::thread_scope_device_t{}, 0, 1);

\[\begin{split}\left(\begin{array}{cc|cc} 1 & 2 & 5 & \text{x} \\ 3 & 4 & 8 & \text{x} \\ \hline 5 & 6 & 6 & \text{x} \\ 7 & 8 & 9 & \text{x} \end{array}\right) \rightarrow \left(\begin{array}{cc|cc} 1 & 2 & 2 & \text{x} \\ 3 & 4 & 5 & \text{x} \\ \hline 5 & 6 & 6 & \text{x} \\ 7 & 8 & 9 & \text{x} \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 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:

for the non-masked overload, 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 partition view’s index space
The tensor span function is not injective

CUDA 13.5 For the masked overload, no store operation is performed to the out of bounds elements.

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

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.

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.

You are visitor 01402503 — GeezerVison does not store the pages you look up.
Best viewed with Netscape Navigator 4.7 or Internet Explorer 6 at 800×600. No cookies required.