Tensor Span

A ct::tensor_span is a tensor span like type representing a view to a multi-dimensional array in memory.

Example

The following example creates a \(2 \times 4\) row major view of elements of the x array. The length \(2\) is supplied at runtime while \(4\) is supplied at compile time.

namespace ct = ::cuda::tiles;
using namespace ct::literals;
int x[8] = {0, 1, 2, 3, 4, 5, 6, 7};

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

\[\begin{split}\begin{pmatrix} 0 & 1 & 2 & 3 \\ 4 & 5 & 6 & 7 \end{pmatrix}\end{split}\]

cuda::tiles::tensor_span

template< ct::scalar E, ct::extents_like Extents, typename Layout = ct::layout_right, ct::accessor_policy Accessor = ct::default_accessor<E> > requires ct::layout_mapping<Layout::mapping<Extents>> && is-same-v<E, Accessor::element_type> struct tensor_span

A ct::tensor_span is an implementation of the tensor span like abstraction. It is a handle to an in-memory multi-dimensional array along with a layout mapping object describing the array’s arrangement of elements and an accessor policy object providing customization for how elements are accessed.

A ct::tensor_span can represent arbitrary array shapes with mixed static and dynamic size and stride information.

All specializations of ct::tensor_span model tensor span like.

A specialization of ct::tensor_span may additionally satisfy the following constraints if the underlying layout mapping type Layout::mapping<Extents> and accessor policy Accessor satisfies them:

std::is_nothrow_move_constructible_v
std::is_nothrow_move_assignable_v
std::is_nothrow_swappable_v

The specialization is ill-formed if Extents, Layout, or Accessor is cv-qualified.

Exposition Only Helpers

For the purposes of defining the behavior of the implicitly defined special member functions, the following members exist.

[[no_unique_address]] mapping_type __mapping: Exposition only member variable indicating the layout of the multi-dimensional array referenced by this object.

[[no_unique_address]] accessor_type __accessor: Exposition only member describing the accessor policy for this tensor span.

[[no_unique_address]] data_handle_type __handle: Exposition only pointer to the multi-dimensional array referenced by this object.

Member Aliases

using element_type = E;
using value_type = remove-cv-t<E>
using extents_type = Extents 
using layout_type = Layout 
using mapping_type = layout_type::mapping<extents_type>
using index_type = extents_type::index_type
using rank_type = extents_type::rank_type
using accessor_type = Accessor 
using data_handle_type = accessor_type::data_handle_type
using reference = accessor_type::reference

Construction From Mapping

__tile__ __host__ __device__ constexpr tensor_span(data_handle_type handle, mapping_type const &mapping, accessor_type const &accessor = {}) noexcept;: Direct list initializes __handle from handle, __mapping from mapping, and __accessor from accessor.

Construction From Policy

__tile__ __host__ __device__ constexpr tensor_span(data_handle_type handle, extents_type const &extents, layout_type const &layout = {}, accessor_type const &accessor = {}): Direct list initializes __handle from handle, __mapping from extents, and __accessor from accessor. This function participates in overload resolution only if is-constructible-v<extents_type const&, mapping_type> is true.

Note

The argument layout is functionally ignored but may be used for the purposes of class template argument deduction.

rank

__tile__ __host__ __device__ static constexpr rank_type rank() noexcept;: Yields mapping_type::rank().

Rank Dynamic

__tile__ __host__ __device__ static constexpr rank_type rank_dynamic() noexcept: Yields mapping_type::rank_dynamic().

static_extent

__tile__ __host__ __device__ static constexpr ct::size_t static_extent(rank_type dim) noexcept: Yields extents_type::static_extent(dim).

extent

__tile__ __host__ __device__ constexpr index_type extent(rank_type dim) const noexcept: Yields __mapping.extents().extent(dim).

extents

__tile__ __host__ __device__ constexpr extents_type const &extents() const noexcept;: Yields __mapping.extents().

mapping

__tile__ __host__ __device__ constexpr mapping_type const &mapping() const noexcept;: Yields __mapping.

accessor

__tile__ __host__ __device__ constexpr accessor_type const &accessor() const noexcept;: Yields __accessor.

data_handle

__tile__ __host__ __device__ constexpr data_handle_type data_handle() const noexcept;: Yields __handle.

Deduction Guides

template<typename T, typename M, typename A = ct::default_accessor<T>> tensor_span(T*, M const&, A const& = {}) -> tensor_span<T, M::extents_type, M::layout_type, A>;
template<typename T, typename E, typename L = ct::layout_right, typename A = ct::default_accessor<T>> tensor_span(T*, E const&, L const& = {}, A const& = {}) -> tensor_span<T, E, L, A>;

Enables deduction of class template arguments from constructor arguments.

cuda::tiles::storeable_tensor_span

template<typename T> concept storeable_tensor_span = tensor_span_like<T> && /* atomic constraint */: Indicates whether \(T\) is a tensor span like type that may participate in a store operation. The atomic constraint validates that the element type of \(T\) is not cv-qualified.

cuda::tiles::default_accessor

template<ct::scalar E> struct default_accessor

A default_accessor is an accessor policy for simple contiguous memory accesses.

Aliases

using element_type = E 
using data_handle_type = E*
using reference = E&

Specializations

template<ct::scalar E> constexpr bool enable_contiguous_accessor_policy<ct::default_accessor<E>> = true: Specialization of variable template enable_contiguous_accessor_policy for ct::default_accessor.

cuda::tiles::enable_contiguous_accessor_policy

template<typename T> inline constexpr bool enable_contiguous_accessor_policy = false

Trait variable which may be specialized to mark custom types as a contiguous accessor policy.