General Operations

Thread Functions

cuda::tiles::bid

__tile__ uint3 bid() noexcept;

Yields the block index of the actively executing tile block. The block indices are organized into \(x\), \(y\), and \(z\) and are determined by the launch parameters of the tile kernel.

Example

The following code copies the elements from array in to out using \(16\) blocks arranged in a \(4 \times 4\) grid. Each block handles \(8\) elements of data at a time. The ct::bid() function is used to select the starting offset of the data that will be handled by the block.

#include "cuda_tile.h"

namespace ct = ::cuda::tiles;
using i32x8 = ct::tile<int, ct::shape<8>>;

__tile_global__ void blockCopy(int* in, int* out) {
  int offset = ct::bid().x * ct::num_blocks().y * 8 + ct::bid().y * 8;

  auto data = ct::load(in + offset + ct::iota<i32x8>());
  ct::store(out + offset + ct::iota<i32x8>(), data);
}

void callBlockCopy() {
  int* pOut, *pIn;
  cudaMalloc(&pIn, 16 * 8 * sizeof(int));
  cudaMemset(pIn, 0, 16 * 8 * sizeof(int));
  cudaMalloc(&pOut, 16 * 8 * sizeof(int));
  blockCopy<<<dim3(4, 4, 1), 1>>>(pIn, pOut);
}

cuda::tiles::num_blocks

__tile__ dim3 num_blocks() noexcept;: Yields the number of blocks in the tile grid. The number of blocks is determined by the kernel launch parameters.

Type Classification

cuda::tiles::integral

template<typename T> concept integral = /* atomic constraint */;: Determines if T is a possibly cv-qualified integral type.

cuda::tiles::tile_like

template<typename T> concept tile_like = /* atomic constraint */;: Indicates whether T is a tile like type.

cuda::tiles::pointer_tile

template<typename T> concept pointer_tile = ct::tile_like<T> && /* atomic constraint */: Indicates whether T is a pointer tile type.

cuda::tiles::numeric_tile

template<typename T> concept numeric_tile = ct::tile_like<T> && /* atomic constraint */;: Indicates whether T is a numeric tile type.

cuda::tiles::arithmetic_tile

template<typename T> concept arithmetic_tile = ct::numeric_tile<T> && /* atomic constraint */;: Indicates whether T is an arithmetic tile type.

cuda::tiles::floating_point_tile

template<typename T> concept floating_point_tile = ct::numeric_tile<T> && /* atomic constraint */;: Indicates whether T is a floating point tile type.

cuda::tiles::basic_floating_point_tile

template<typename T> concept basic_floating_point_tile = ct::arithmetic_tile<T> && ct::floating_point_tile<T> && /* atomic constraint */;: Indicates whether T is a basic floating point tile type.

cuda::tiles::restricted_floating_point_tile

template<typename T> concept restricted_floating_point_tile = ct::floating_point_tile<T> && /* atomic constraint */;: Indicates whether T is a restricted floating point tile type.

cuda::tiles::integral_tile

template<typename T> concept integral_tile = ct::arithmetic_tile<T> && /* atomic constraint */;: Indicates whether T is an integral tile type.

cuda::tiles::scalar

template<typename T> concept scalar = ct::tile_like<T> && /* atomic constraint */;: Indicates whether T is a scalar type.

cuda::tiles::pointer_scalar

template<typename T> concept pointer_scalar = ct::pointer_tile<T> && ct::scalar<T>;: Indicates whether T is a pointer scalar type.

cuda::tiles::numeric_scalar

template<typename T> concept numeric_scalar = ct::numeric_tile<T> && ct::scalar<T>;: Indicates whether T is a numeric scalar type.

cuda::tiles::arithmetic_scalar

template<typename T> concept arithmetic_scalar = ct::arithmetic_tile<T> && ct::scalar<T>;: Indicates whether T is an arithmetic scalar type.

Note

ct::arithmetic_scalar subsumes ct::numeric_tile and ct::numeric_scalar.

cuda::tiles::floating_point_scalar

template<typename T> concept floating_point_scalar = ct::floating_point_tile<T> && ct::scalar<T>;: Indicates whether T is a floating point scalar type.

Note

ct::floating_point_scalar subsumes ct::numeric_tile and ct::numeric_scalar.

cuda::tiles::basic_floating_point_scalar

template<typename T> concept basic_floating_point_scalar = ct::basic_floating_point_tile<T> && ct::scalar<T>;

Indicates whether T is a basic floating point scalar type.

Note

ct::basic_floating_point_scalar subsumes all of the following:

ct::floating_point_tile
ct::arithmetic_tile
ct::floating_point_scalar
arithmetic_scalar

cuda::tiles::restricted_floating_point_scalar

template<typename T> concept restricted_floating_point_scalar = ct::restricted_floating_point_tile<T> && ct::scalar<T>;: Indicates whether T is a restricted floating point scalar type.

Note

ct::restricted_floating_point_scalar subsumes ct::floating_point_tile and ct::floating_point_scalar.

cuda::tiles::integral_scalar

template<typename T> concept integral_scalar = ct::integral_tile<T> && ct::scalar<T> && ct::integral<T>;: Indicates whether T is an integral scalar type.

Note

ct::integral_scalar subsumes ct::arithmetic_tile and ct::arithmetic_scalar.

Extents Like Properties

cuda::tiles::extents_like

template<typename T> concept extents_like = /* atomic constraint */: Indicates whether T is an extents like type.

cuda::tiles::extents_equal

template<ct::extents_like T, ct::extents_like U> __tile__ __host__ __device__ constexpr bool extents_equal(T const &x, U const &y) noexcept;: Indicates whether the two extents like objects x and y are extent equivalent.

cuda::tiles::shape_like

template<typename T> concept shape_like = ct::extents_like<T> && /* atomic constraint */;: Indicates whether T is a shape like type.

cuda::tiles::shape_size_v

template<ct::shape_like T> inline constexpr size_t shape_size_v = /* see below */;: Yields the size of the shape like type T. If the size is not representable within the type size_t, the behavior is undefined.

cuda::tiles::same_shape

template<typename T, typename U> concept same_shape = ct::shape_like<T> && ct::shape_like && /* atomic constraint */;: Indicates whether shape like types T and U are shape equivalent.

Tile Like Properties

cuda::tiles::tile_element_t

template<ct::tile_like T> using tile_element_t = /* see below */;: Yields the element type of a tile like type T.

cuda::tiles::tile_shape_t

template<ct::tile_like T> using tile_shape_t = /* see below */;: Yields the shape type of a tile like type T.

cuda::tiles::tile_size_v

template<ct::tile_like T> inline constexpr size_t tile_size_v = ct::shape_size_v<ct::tile_shape_t<T>>;: Yields the tile size of tile like type T.

cuda::tiles::tile_rank_v

template<ct::tile_like T> inline constexpr size_t tile_rank_v = ct::tile_shape_t<T>::rank();: Yields the rank of tile like type T.

cuda::tiles::tile_with_element_t

template<ct::tile_like T, ct::scalar E> using tile_with_element_t = /* see below */;: Yields a tile like type whose element type is E and whose shape matches that of T. If T is a specialization of ct::tile, the result is also a specialization of tile. If T is a scalar type, the result is a Scalars type. The result type is never cv-qualified.

cuda::tiles::tile_shape

template<typename T> concept tile_shape = ct::shape_like<T> && /* atomic constraint */: Indicates whether T is a shape like type that is tile compatible.

Conversions

cuda::tiles::scalar_convertible_to

template<typename T, typename U> concept scalar_convertible_to = ct::scalar<T> && ct::scalar && /* atomic constraint */: Indicates whether there exists a scalar conversion from objects of type T to type U.

cuda::tiles::non_narrowing_scalar_convertible_to

template<typename T, typename U> concept non_narrowing_scalar_convertible_to = ct::scalar_convertible_to<T, U> && /* atomic constraint */: Indicates whether there exists a non-narrowing scalar conversion from objects of type T to type U.

cuda::tiles::tile_convertible_to

template<typename T, typename U> concept tile_convertible_to = ct::tile_like<T> && ct::tile_like && /* atomic constraint */;: Indicates whether there exists a tile conversion from an object of type T to type U.

cuda::tiles::non_narrowing_tile_convertible_to

template<typename T, typename U> concept non_narrowing_tile_convertible_to = ct::tile_convertible_to<T, U> && /* atomic constraint */;: Indicates whether there exists a non-narrowing tile conversion from an object of type T to type U.

cuda::tiles::bool_tile_convertible

template<typename T> concept bool_tile_convertible = ct::tile_like<T> && ct::tile_convertible_to<T, ct::tile_with_element_t<T, bool>>;: Indicates whether T bool tile convertible.

cuda::tiles::shape_broadcastable_to

template<typename T, typename U> concept shape_broadcastable_to = ct::shape_like<T> && ct::shape_like && /* atomic constraint */: Indicates whether shape like type T may be broadcasted to type U.

cuda::tiles::broadcastable_to

template<typename T, typename U> concept broadcastable_to = ct::tile_like<T> && ct::tile_shape && ct::shape_broadcastable_to<ct::tile_shape_t<T>, U>;: Indicates whether tile like type T is broadcastable to shape U.

cuda::tiles::shape_broadcast_compatible

template<typename T, typename U> concept shape_broadcast_compatible = ct::shape_like<T> && ct::shape_like && /* atomic constraint */;: Indicates whether shape like types T and U are mutual broadcast compatible.

cuda::tiles::shape_broadcast_t

template<ct::shape_like T, ct::shape_like U> requires ct::shape_broadcast_compatible<T, U> using shape_broadcast_t = /* see below */;: Yields the mutual broadcast shape of shape like types T and U.

cuda::tiles::broadcast_compatible

template<typename T, typename U> concept broadcast_compatible = ct::tile_like<T> && ct::tile_like && /* atomic constraint */;: Indicates whether T and U are broadcast compatible.

cuda::tiles::mutual_broadcast_t

template<ct::tile_like T, ct::tile_like U, ct::scalar E> requires ct::broadcast_compatible<T, U> using mutual_broadcast_t = /* see below */;: Yields a cv-unqualified tile like type whose shape matches the mutual broadcast shape of the shapes of T and U and whose element type is E. The result type is a scalar if and only if both T and U are scalars.

cuda::tiles::arithmetic_tile_convertible

template<typename T, typename U> concept arithmetic_tile_convertible = ct::arithmetic_tile<T> && ct::arithmetic_tile && ct::broadcast_compatible<T, U> && /* atomic constraint */: Tests whether an arithmetic tile conversion exists between objects of types T and U.

cuda::tiles::arithmetic_tile_conversion_t

template<ct::arithmetic_tile T, ct::arithmetic_tile U> requires ct::arithmetic_tile_convertible<T, U> using arithmetic_tile_conversion_t = /* see below */;: Yields the common type \(C\) after running arithmetic tile conversion on two objects of types T and U.

cuda::tiles::arithmetic_tile_comparable

template<typename T, typename U> concept arithmetic_tile_comparable = ct::arithmetic_tile<T> && ct::arithmetic_tile && broadcast_compatible<T, U> && /* atomic constraint */: Tests whether an arithmetic comparison conversion exists between objects of types T and U.

cuda::tiles::arithmetic_tile_comparison_t

template<ct::arithmetic_tile T, ct::arithmetic_tile U> requires ct::arithmetic_tile_comparable<T, U> using arithmetic_tile_comparison_t = /* see below */;: Yields the result type \(C\) after running arithmetic comparison conversion on two objects of types T and U.

cuda::tiles::arithmetic_tile_promotion_t

template<ct::arithmetic_tile T> using arithmetic_tile_promotion_t = /* see below */;: Yields the result type after applying arithmetic tile promotion to a hypothetical object a of type T.

Tensor Span Like Properties

cuda::tiles::layout_mapping

template<typename T> concept layout_mapping = /* atomic constraint */: Indicates whether T is a layout mapping type.

cuda::tiles::accessor_policy

template<typename T> concept accessor_policy = /* atomic constraint */: Indicates whether T is an accessor policy type.

cuda::tiles::tensor_span_like

template<typename T> concept tensor_span_like = /* atomic constraint */: Indicates whether T is a tensor span like type.

cuda::tiles::layout_mapping_equal

template<ct::layout_mapping Lhs, ct::layout_mapping Rhs> __tile__ __host__ __device__ constexpr bool layout_mapping_equal(Lhs const &x, Rhs const &y) noexcept;: Indicates whether layout mappings x and y are layout mapping equivalent.

Exposition Only Entities

The following exposition only entities are defined below in terms of a C++23 1 standard library API. These APIs are not necessarily defined or available when including the cuda_tile.h header.

template<typename T> using remove-cv-t = std::remove_cv_t<T>;
template<typename T> using remove-pointer-t = std::remove_pointer_t<T>;
template<typename T> using make-signed-t = std::make_signed_t<T>;
using nullptr-t = decltype(nullptr);

template<typename T> inline constexpr bool is-const-v = std::is_const_v<T>;
template<typename T> inline constexpr bool is-volatile-v = std::is_volatile_v<T>;
template<typename T, typename U> inline constexpr bool is-convertible-v = std::is_convertible_v<T, U>;
template<typename T, typename U> inline constexpr bool is-nothrow-convertible-v = std::is_nothrow_convertible_v<T, U>;
template<typename T, typename ...Args> inline constexpr bool is-constructible-v = std::is_constructible_v<T, Args...>;

Footnotes

1: See ISO/IEC 14882:2024