*CuTe composed layout: A ∘ offset ∘ B*


A ``ComposedLayout`` models a CuTe composed layout: a three-part composition
``A o offset o B`` where a coordinate ``c`` maps to ``A(offset + B(c))``.
Same as layouts, a composed layout is also a function that maps a coordinate to an index.

Parts:

- ``A``: either a layout or a swizzle (the outer mapping or swizzle function)
- ``offset``: an integer tuple applied between B and A
- ``B``: the base layout (defines the domain)

**Examples**

1. ``!cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">``

   +----------------+----------------+----------------+-------------------------+-------------------------+
   | Integral coord | Hier Coord (c) | B(c) = B index | offset + B(c) = A index | A o offset o B (Result) |
   +================+================+================+=========================+=========================+
   | 0              | (0,0)          | 0              | 0 + 2 = 2               | A(2) = 2                |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 1              | (1,0)          | 1              | 1 + 2 = 3               | A(3) = 3                |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 2              | (0,1)          | 2              | 2 + 2 = 4               | A(4) = 4                |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 3              | (1,1)          | 3              | 3 + 2 = 5               | A(5) = 5                |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 4              | (0,2)          | 4              | 4 + 2 = 6               | A(6) = 6                |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 5              | (1,2)          | 5              | 5 + 2 = 7               | A(7) = 7                |
   +----------------+----------------+----------------+-------------------------+-------------------------+

2. ``!cute.composed_layout<"(6,2):(1,3) o (2,1) o (2,3):(1@1,2@0)">``

   B = ``(2,3):(1@1,2@0)`` uses scaled-basis strides — mode 0 of the shape
   contributes its coord to basis dim 1, mode 1 of the shape contributes
   ``2 * coord`` to basis dim 0. So ``B(c)`` returns a rank-2 tuple:
   ``(basis-0-contribution, basis-1-contribution)``. The tuple offset ``(2,1)``
   is added per-component (matching basis-0 / basis-1 slots).
   A = ``(6,2):(1,3)`` then applies its strides to the resulting ``(n0, n1)``
   tuple coord: ``A(n0, n1) = n0 * 1 + n1 * 3``.

   +----------------+----------------+----------------+-------------------------+-------------------------+
   | Integral coord | Hier Coord (c) | B(c) = B index | offset + B(c) = A index | A o offset o B (Result) |
   +================+================+================+=========================+=========================+
   | 0              | (0,0)          | (0,0)          | (0,0) + (2,1) = (2,1)   | A(2,1) =  5             |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 1              | (1,0)          | (0,1)          | (0,1) + (2,1) = (2,2)   | A(2,2) =  8             |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 2              | (0,1)          | (2,0)          | (2,0) + (2,1) = (4,1)   | A(4,1) =  7             |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 3              | (1,1)          | (2,1)          | (2,1) + (2,1) = (4,2)   | A(4,2) = 10             |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 4              | (0,2)          | (4,0)          | (4,0) + (2,1) = (6,1)   | A(6,1) =  9             |
   +----------------+----------------+----------------+-------------------------+-------------------------+
   | 5              | (1,2)          | (4,1)          | (4,1) + (2,1) = (6,2)   | A(6,2) = 12             |
   +----------------+----------------+----------------+-------------------------+-------------------------+

Constraints:

The ``offset`` parameter is an ``int_tuple`` whose value is added — via
``cute.tuple_add`` semantics — to the result of
``cute.layout_eval(input_crd, B)``. The sum then flows as the indexing
argument into ``A``. Because of this, the output profile of
``cute.layout_eval(input_crd, B)`` and ``A``'s input profile must be
compatible.

1. Profile compatibility with ``cute.layout_eval(input_crd, B)``. The
   sum ``offset + cute.layout_eval(input_crd, B)`` must produce a
   well-formed ``int_tuple`` that ``A`` can consume. In practice this
   means:

   - A scalar offset (``0``, any static integer, ``?``) is always safe
     when ``B`` has plain integer strides: ``cute.layout_eval`` returns a
     scalar integer and scalar + scalar is well-formed.

   - A tuple offset (e.g. ``(0,0)``, ``(1,2)``) is required when ``B``'s
     strides are scaled-basis in a way that makes ``cute.layout_eval``
     return a tuple.

2. ``A`` is a swizzle. Swizzles consume only scalar indices, so
   the sum ``offset + cute.layout_eval(input_crd, B)`` must be a scalar
   (static or dynamic integer). Any surviving scaled-basis
   structure in the sum is rejected.

3. ``A`` is a layout (affine). The sum either is a scalar (and
   ``A`` indexes it directly) or must match ``A``'s shape profile so
   each leaf of the sum dispatches to the corresponding mode of
   ``A``. Profile mismatches are rejected.

4. Positive shape modes. All static shape leaves in ``B`` (and
   in ``A`` when ``A`` is a layout) must be positive.

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the composed layout.

**Examples**

.. code-block::

   !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">          // layout A
   !cute.composed_layout<"(4,5):(1,4) o (2,1) o (2,3):(1@1,2@0)">  // layout A + tuple offset
   !cute.composed_layout<"S<3,5,4> o 2 o (2,3):(1,2)">             // swizzle A
   !cute.composed_layout<"S<0,4,3> o 0 o (8,4):(1,8)">             // identity swizzle



.. _type-coord:

``!cute.coord``
~~~~~~~~~~~~~~~

*Scalar integer, underscore wildcard, or recursive tuple of coordinates*


A ``Coord`` models a CuTe coordinate: either a single leaf value (depth 0)
or a tuple of ``Coord`` elements (depth ≥ 1), arbitrarily nested.

Leaf values:

- static integer — compile-time constant
- underscore (``_``) — the CuTe slice wildcard
- dynamic integer — runtime value (``?``), default i32 width;
  ``?{i64}`` for an i64-wide dynamic

A coordinate is considered static when no leaf is dynamic (no ``?``); the
``_`` wildcard is a compile-time slice marker and counts as static.

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the coordinate value.

**Examples**

.. code-block::

   !cute.coord<"1">           // depth-0, scalar static 1
   !cute.coord<"_">           // depth-0, underscore wildcard
   !cute.coord<"?">           // depth-0, scalar dynamic
   !cute.coord<"(1,2)">       // depth-1, rank-2 static
   !cute.coord<"(_,?)">       // depth-1, underscore + dynamic
   !cute.coord<"(_,_,?)">     // depth-1, two underscores + dynamic
   !cute.coord<"(1,(2,3))">   // depth-2, nested static



.. _type-int_tuple:

``!cute.int_tuple``
~~~~~~~~~~~~~~~~~~~

*Scalar integer or recursive tuple of integers*


An ``IntTuple`` models a CuTe integer tuple: either a single integer (depth 0)
or a tuple of ``IntTuple`` elements (depth ≥ 1), arbitrarily nested.

Leaf values:

- static integer — compile-time constant
- dynamic integer — runtime value (``?``), default i32 width; ``?{i64}``
  for an i64-wide dynamic.

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the integer tuple.

**Examples**

.. code-block::

   !cute.int_tuple<"1">                   // depth-0, scalar static 1
   !cute.int_tuple<"?">                   // depth-0, scalar dynamic (i32)
   !cute.int_tuple<"?{i64}">              // depth-0, scalar dynamic (i64)
   !cute.int_tuple<"(1,2)">               // depth-1, rank-2 static tuple
   !cute.int_tuple<"(1,(2,3))">           // depth-2, nested static tuple
   !cute.int_tuple<"(?,32)">              // depth-1, mixed: dynamic + static
   !cute.int_tuple<"(?,?{i64})">          // depth-1, mixed widths



.. _type-layout:

``!cute.layout``
~~~~~~~~~~~~~~~~

*CuTe layout: a shape/stride pair*


A ``Layout`` models a CuTe layout: a shape/stride pair that maps a
coordinate space to an index space.  The shape and stride may be scalar
(depth 0) or arbitrarily nested tuples (depth ≥ 1).

Leaf values in the shape:

- static integer — compile-time constant extent
- dynamic integer — runtime extent (``?``), default i32 width;
  ``?{i64}`` for an i64-wide dynamic

Leaf values in the stride:

- static integer — compile-time constant stride
- dynamic integer — runtime stride (``?``), default i32 width;
  ``?{i64}`` for i64

- scaled basis (``V@Mn@...@M0``) — value scaled by a basis vector; a single
  index ``V@M`` addresses a top-level mode, chained indices ``V@Mn@...@M0``
  address a nested mode (corresponds to ``E<M0,...,Mn>``, see !cute.stride)

A layout is considered static when no leaf is dynamic (no ``?``).

**Examples**

1. ``A = !cute.layout<"(2,3):(1,2)">`` — affine rank-2 layout.
   The integral coord enumerates ``size(A) = 2*3 = 6`` positions;
   ``idx2crd`` maps it to the hier coord with column-major order,
   and ``A(c) = c0*1 + c1*2``. This layout is the identity (A(c) = i).

   +--------------------+----------------+------+
   | Integral coord (i) | Hier coord (c) | A(c) |
   +====================+================+======+
   | 0                  | (0, 0)         | 0    |
   +--------------------+----------------+------+
   | 1                  | (1, 0)         | 1    |
   +--------------------+----------------+------+
   | 2                  | (0, 1)         | 2    |
   +--------------------+----------------+------+
   | 3                  | (1, 1)         | 3    |
   +--------------------+----------------+------+
   | 4                  | (0, 2)         | 4    |
   +--------------------+----------------+------+
   | 5                  | (1, 2)         | 5    |
   +--------------------+----------------+------+

2. ``B = !cute.layout<"((2,3),4):((1@0@0,1@1@0),1@1)">`` — hierarchical
   basis-stride layout. Shape ``((2,3),4)`` has ``size = 6*4 = 24``. The
   basis strides decompose the linear index into a *nested* result
   int_tuple whose shape mirrors the coord: each ``1@Mn@...@M0``
   deposits its scalar contribution at the slot identified by the
   chained mode indices.

   - ``1@0@0`` (E<0,0>) → mode-0 inner-0 of result
   - ``1@1@0`` (E<0,1>) → mode-0 inner-1 of result
   - ``1@1``   (E<1>)   → mode-1 of result

   So for coord ``c = ((c00, c01), c1)``, ``B(c) = ((c00, c01), c1)`` —
   a structural identity that preserves the hierarchy. Spot checks
   (verified via ``cute.layout_eval``):

   +----------------+-------------+
   | Hier coord (c) | B(c)        |
   +================+=============+
   | ((0, 0), 0)    | ((0, 0), 0) |
   +----------------+-------------+
   | ((1, 0), 0)    | ((1, 0), 0) |
   +----------------+-------------+
   | ((0, 1), 0)    | ((0, 1), 0) |
   +----------------+-------------+
   | ((0, 0), 1)    | ((0, 0), 1) |
   +----------------+-------------+
   | ((1, 2), 0)    | ((1, 2), 0) |
   +----------------+-------------+
   | ((1, 2), 3)    | ((1, 2), 3) |
   +----------------+-------------+

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the layout.

**Examples**

.. code-block::

   !cute.layout<"(2,3):(1,2)">                   // static rank-2 layout
   !cute.layout<"(2,3):(?,2)">                   // dynamic stride element
   !cute.layout<"(?,3):(1,2)">                   // dynamic shape element
   !cute.layout<"((2,3),(4,5)):((1,2),(3,4))">   // nested shape and stride
   !cute.layout<"(2,3):(1@0,1@1)">               // scaled-basis strides
   !cute.layout<"((2,3),4):((1@0@0,1@1@0),1@1)"> // hierarchical basis strides (mode-0 rank-2 with depth-2 hierarchical; mode-1 rank-1 single-level)



.. _type-shape:

``!cute.shape``
~~~~~~~~~~~~~~~

*Scalar integer or recursive tuple of shape extents*


A ``Shape`` models a CuTe shape (extent tuple): either a single integer
(depth 0) or a tuple of ``Shape`` elements (depth ≥ 1), arbitrarily nested.

Leaf values:

- static integer — compile-time constant (typically a positive
  extent, but ``0`` is permitted because some valid uses of
  ``!cute.shape`` carry counts or paddings; positivity is
  enforced where it matters, e.g. in ``!cute.layout``'s verifier
  against the layout's shape mode).

- dynamic integer — runtime value (``?``), default i32 width;
  ``?{i64}`` for an i64-wide dynamic

A shape is considered static when no leaf is dynamic (no ``?``).

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the shape value.

**Examples**

.. code-block::

   !cute.shape<"1">           // depth-0, scalar static 1
   !cute.shape<"?">           // depth-0, scalar dynamic (i32)
   !cute.shape<"?{i64}">      // depth-0, scalar dynamic (i64)
   !cute.shape<"(1,2)">       // depth-1, rank-2 static shape
   !cute.shape<"(?,32)">      // depth-1, mixed: dynamic + static
   !cute.shape<"(1,(2,3))">   // depth-2, nested static shape



.. _type-stride:

``!cute.stride``
~~~~~~~~~~~~~~~~

*Scalar or recursive tuple of stride elements*


A ``Stride`` models a CuTe stride: either a single leaf value (depth 0)
or a tuple of ``Stride`` elements (depth ≥ 1), arbitrarily nested.

Leaf values:

- static integer — compile-time constant (e.g. ``1``, ``16``)
- dynamic integer — runtime value (``?``), default i32 width; ``?{i64}``
  for an i64-wide dynamic

- scaled basis (``V@Mn@...@M0``) — value scaled by a basis vector; a single
  index ``V@M`` addresses a top-level mode, chained indices ``V@Mn@...@M0``
  address a nested mode (corresponds to ``E<M0,...,Mn>``); ``V``
  may be static (e.g. ``4@0``) or dynamic (e.g. ``?@0``). Note the Mn@...@M0
  order, the last M0 is the outer most mode. For example, ``1@1@0`` is the stride
  for ((x, 1), ...). In other words, ``1@1@0`` is the stride for the mode-1 of the
  mode-0 of the coordinate space.

A stride is considered static when no leaf is dynamic (no ``?``);

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the stride value.

**Examples**

.. code-block::

   !cute.stride<"1">             // depth-0, scalar static stride 1
   !cute.stride<"?">             // depth-0, scalar dynamic
   !cute.stride<"4@0">           // depth-0, static scaled basis: 4 × e_0
   !cute.stride<"?@0">           // depth-0, dynamic scaled basis
   !cute.stride<"1@1@0">         // depth-0, hierarchical basis: E<0,1>
   !cute.stride<"(16,32)">       // depth-1, rank-2 static stride
   !cute.stride<"(16,?)">        // depth-1, static + dynamic
   !cute.stride<"(0,?@0)">       // depth-1, zero-int + dynamic scaled basis
   !cute.stride<"(1@1@0,2@0)">   // depth-1, hierarchical + regular basis
   !cute.stride<"(16,(32,4))">   // depth-2, nested static stride



.. _type-swizzle:

``!cute.swizzle``
~~~~~~~~~~~~~~~~~

*CuTe swizzle: S<num_bits, num_base, num_shift>*


A ``Swizzle`` models a CuTe swizzle function — a bit-manipulation
transformation parameterized by three compile-time integers:

- ``num_bits``: the number of bits involved in the XOR
- ``num_base``: the base bit position
- ``num_shift``: the shift amount

Swizzle has no dynamic parts, it is always static.

Assembly format — ``< "string" >`` where the string is the canonical
printed form ``S<num_bits,num_base,num_shift>``.

**Examples**

.. code-block::

   !cute.swizzle<"S<0,4,3>">  // canonical identity swizzle
   !cute.swizzle<"S<3,5,4>">  // non-identity swizzle
   !cute.swizzle<"S<1,0,2>">  // example from CuTe GEMM kernels



.. _type-tile:

``!cute.tile``
~~~~~~~~~~~~~~

*CuTe tile: a recursive sequence of layouts and underscores*


A ``Tile`` models a CuTe tiler: a recursive sequence of ``Layout`` elements
and underscore wildcards (depth ≥ 0).

Leaf values:

- layout: a shape/stride pair (the primary tile element)
- underscore (``_``): the CuTe slice wildcard; keeps the corresponding mode

``isStatic()`` returns ``true`` only when all layout leaves are fully static.

Assembly format — ``< "string" >`` where the string is the canonical
printed form of the tile, enclosed in brackets with elements separated by
semicolons.

**Examples**

.. code-block::

   !cute.tile<"[(2,3):(1,2)]">              // single-layout tile
   !cute.tile<"[(2,3):(1,2);(4,5):(1,4)]">  // two-layout tile
   !cute.tile<"[_]">                        // underscore tile (slice all)
   !cute.tile<"[(2,3):(1,2);_]">            // layout + underscore



Operations
----------

Constructor Operations
~~~~~~~~~~~~~~~~~~~~~~

Build cute values (shapes, strides, coords, layouts, tuples) directly from runtime operands.

.. _op-cute.make_composed_layout:

``cute.make_composed_layout``
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Construct a ComposedLayout from inner (layout or swizzle), offset, and outer*

**Assembly format**

.. code-block::

   `(` $inner `,` $offset `,` $outer `)` attr-dict `:`
       `(` type($inner) `,` type($offset) `,` type($outer) `)` `->` type($result)

**Operands**

- ``$inner``: inner layout or swizzle
- ``$offset``: offset int_tuple
- ``$outer``: outer layout

**Results**

- ``$result``: assembled composed layout

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Constructs a ``!cute.composed_layout`` from three operands. A
composed layout represents the function composition
``inner ∘ offset ∘ outer``. For a coordinate ``c``, evaluation proceeds
as ``inner(offset + outer(c))``. The result type is inferred from the
three operand types.

**Preconditions**

- ``$inner`` (the A component, applied last) is ``!cute.layout`` or
  ``!cute.swizzle``; its kind must match the A component encoded in
  the result type.

  - When ``$inner`` is ``!cute.layout``, the result A must be an affine
    layout that exactly matches the operand.

  - When ``$inner`` is ``!cute.swizzle``, the result A must be a
    swizzle whose parameters exactly match the operand.

- ``$offset`` is ``!cute.int_tuple`` and exactly matches the offset
  component encoded in the result type.

- ``$outer`` is ``!cute.layout`` (the B component, applied first) and
  exactly matches the B layout component of the result type.

**Postconditions**

- Result is the ``!cute.composed_layout`` whose A, offset, and B
  components match the three operands and the result type.

**Examples**

.. code-block::

   // Layout A, scalar offset, static outer
   %cl = cute.make_composed_layout(%a, %off, %b)
           : (!cute.layout<"(4,5):(1,4)">, !cute.int_tuple<"2">,
              !cute.layout<"(2,3):(1,2)">)
          -> !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">

   // Swizzle A, zero offset, static outer
   %cl = cute.make_composed_layout(%sw, %off, %b)
           : (!cute.swizzle<"S<3,5,4>">, !cute.int_tuple<"0">,
              !cute.layout<"(8,4):(1,8)">)
          -> !cute.composed_layout<"S<3,5,4> o 0 o (8,4):(1,8)">


**Errors**

- inner layout does not match result A

  .. code-block::

     %cl = cute.make_composed_layout(%a, %off, %b)
             : (!cute.layout<"(4,5):(2,8)">, !cute.int_tuple<"2">,
                !cute.layout<"(2,3):(1,2)">)
            -> !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">


- inner is a layout but result A is a swizzle

  .. code-block::

     %cl = cute.make_composed_layout(%a, %off, %b)
             : (!cute.layout<"(4,5):(1,4)">, !cute.int_tuple<"2">,
                !cute.layout<"(2,3):(1,2)">)
            -> !cute.composed_layout<"S<3,5,4> o 2 o (2,3):(1,2)">


- inner is a swizzle but result A is a layout

  .. code-block::

     %cl = cute.make_composed_layout(%sw, %off, %b)
             : (!cute.swizzle<"S<3,5,4>">, !cute.int_tuple<"2">,
                !cute.layout<"(2,3):(1,2)">)
            -> !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">


- swizzle params don't match result A

  .. code-block::

     %cl = cute.make_composed_layout(%sw, %off, %b)
             : (!cute.swizzle<"S<1,5,4>">, !cute.int_tuple<"2">,
                !cute.layout<"(2,3):(1,2)">)
            -> !cute.composed_layout<"S<3,5,4> o 2 o (2,3):(1,2)">


- offset operand does not match result offset

  .. code-block::

     %cl = cute.make_composed_layout(%a, %off, %b)
             : (!cute.layout<"(4,5):(1,4)">, !cute.int_tuple<"5">,
                !cute.layout<"(2,3):(1,2)">)
            -> !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">



.. _op-cute.make_coord:

``cute.make_coord``
^^^^^^^^^^^^^^^^^^^

*Construct a Coord from its runtime-valued leaves*

**Assembly format**

.. code-block::

   `(` $operands `)` attr-dict `:` functional-type($operands, $result)

**Operands**

- ``$operands`` *(variadic)*: dynamic elements

**Results**

- ``$result``: assembled coord

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Constructs a ``!cute.coord`` by supplying runtime values for its
dynamic slots. The result type specifies the full coordinate
structure; static values and underscore (``_``) placeholder slots are
encoded in the type itself and require no operand. Each ``?``
(dynamic) slot requires exactly one operand, supplied left-to-right
in depth-first traversal order. Zero operands are valid when the
result type has no ``?`` slots (fully static, or all dynamic-looking
slots are underscores).

**Preconditions**

- Each operand fills one ``?`` slot, in depth-first left-to-right order:

  - ``i32`` or ``i64`` fills a scalar ``?`` coordinate component;
    bit-width must match the slot's width annotation.

  - rank-1 ``!cute.coord`` fills a nested sub-coord slot; higher-rank
    sub-coords are rejected.

- Underscore (``_``) slots are static and consume no operand.
- The number of operands equals the number of ``?`` slots in the
  result type.

**Postconditions**

- Result is the ``!cute.coord`` whose structure matches the result
  type and whose dynamic leaves take the operand values.

**Examples**

.. code-block::

   %c = cute.make_coord() : () -> !cute.coord<"(1,2)">
   %c = cute.make_coord(%i, %j) : (i32, i32) -> !cute.coord<"(?,?)">
   %c = cute.make_coord(%i) : (i32) -> !cute.coord<"(?,_)">


**Errors**

- too few operands — underscore does not count as a dynamic leaf

  .. code-block::

     %0 = cute.make_coord() : () -> !cute.coord<"(?,_)">


- too many operands for the `?` slots

  .. code-block::

     %0 = cute.make_coord(%i, %j) : (i64, i64) -> !cute.coord<"(?,2)">


- operand width mismatch (i64 supplied for an i32-width `?` slot)

  .. code-block::

     %0 = cute.make_coord(%i) : (i64) -> !cute.coord<"?">


- non-integer-leaf operand rejected at ODS level

  .. code-block::

     %0 = cute.make_coord(%t) : (!cute.coord<"(?,?)">) -> !cute.coord<"(?,?)">



.. _op-cute.make_int_tuple:

``cute.make_int_tuple``
^^^^^^^^^^^^^^^^^^^^^^^

*Construct an IntTuple from its runtime-valued leaves*

**Assembly format**

.. code-block::

   `(` $operands `)` attr-dict `:` functional-type($operands, $result)

**Operands**

- ``$operands`` *(variadic)*: dynamic elements

**Results**

- ``$result``: assembled int_tuple

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Constructs an ``!cute.int_tuple`` by supplying runtime values for its
dynamic slots. The result type specifies the full tuple structure;
static values are encoded in the type itself and require no operand.
Each ``?`` (dynamic) slot in the result type requires exactly one
operand, supplied left-to-right in depth-first traversal order
across the tuple tree. Zero operands are valid when the result type
has no ``?`` slots (fully static).

**Preconditions**

- Each operand fills one ``?`` slot, in depth-first left-to-right order:

  - ``i32`` or ``i64`` fills a scalar ``?`` leaf; bit-width must match the
    slot's width annotation.

  - rank-1 ``!cute.int_tuple`` (holding a single integer or dynamic-
    integer element) fills a nested sub-tuple slot; higher-rank
    sub-tuples are rejected.

- The number of operands equals the number of ``?`` slots in the
  result type.

**Postconditions**

- Result is the ``!cute.int_tuple`` whose structure matches the result
  type and whose dynamic leaves take the operand values.

**Examples**

.. code-block::

   // fully static — no operands
   %t = cute.make_int_tuple() : () -> !cute.int_tuple<"4">

   // scalar dynamic — one i32 leaf
   %t = cute.make_int_tuple(%n) : (i32) -> !cute.int_tuple<"?">

   // flat tuple, two dynamic leaves
   %t = cute.make_int_tuple(%m, %n) : (i32, i32) -> !cute.int_tuple<"(?,?)">

   // nested sub-tuple as operand
   %inner = cute.make_int_tuple(%k) : (i32) -> !cute.int_tuple<"?">
   %t = cute.make_int_tuple(%m, %inner)
          : (i32, !cute.int_tuple<"?">) -> !cute.int_tuple<"(?,?)">

   // fully static tuple — no operands
   %t = cute.make_int_tuple() : () -> !cute.int_tuple<"(4,3)">


**Errors**

- too few operands for the `?` slots

  .. code-block::

     %0 = cute.make_int_tuple(%n) : (i64) -> !cute.int_tuple<"(?,?)">


- too many operands for the `?` slots

  .. code-block::

     %0 = cute.make_int_tuple(%m, %n) : (i64, i64) -> !cute.int_tuple<"(?,3)">


- operand width mismatch (i64 supplied for an i32-width `?` slot)

  .. code-block::

     %0 = cute.make_int_tuple(%n) : (i64) -> !cute.int_tuple<"?">


- depth>0 int_tuple operand rejected

  .. code-block::

     %0 = cute.make_int_tuple(%t)
            : (!cute.int_tuple<"(?,?)">) -> !cute.int_tuple<"(?,?)">



.. _op-cute.make_layout:

``cute.make_layout``
^^^^^^^^^^^^^^^^^^^^

*Construct a Layout from an explicit shape and stride*

**Assembly format**

.. code-block::

   `(` $shape `,` $stride `)` attr-dict `:`
       functional-type(operands, $layout)

**Operands**

- ``$shape``: shape
- ``$stride``: stride

**Results**

- ``$layout``: assembled layout

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Constructs a ``!cute.layout`` from a ``!cute.shape`` and a ``!cute.stride``
operand. Both operands are required and must exactly match the
shape and stride components of the result type.

**Preconditions**

- Operand ``$shape`` is a ``!cute.shape`` and exactly matches the
  result's shape.

- Operand ``$stride`` is a ``!cute.stride`` and exactly matches the
  result's stride.

**Postconditions**

- Result is the ``!cute.layout`` whose shape and stride match the
  result type, with dynamic leaves taking values from the operands.

**Examples**

.. code-block::

   %l = cute.make_layout(%s, %d)
          : (!cute.shape<"(?,2)">, !cute.stride<"(?,4)">) -> !cute.layout<"(?,2):(?,4)">


**Errors**

- shape operand does not match result shape

  .. code-block::

     %l = cute.make_layout(%s, %d)
            : (!cute.shape<"(?,2)">, !cute.stride<"(?,4)">) -> !cute.layout<"(4,2):(?,4)">


- stride operand does not match result stride

  .. code-block::

     %l = cute.make_layout(%s, %d)
            : (!cute.shape<"(?,2)">, !cute.stride<"(?,8)">) -> !cute.layout<"(?,2):(?,4)">



.. _op-cute.make_shape:

``cute.make_shape``
^^^^^^^^^^^^^^^^^^^

*Construct a Shape from its runtime-valued leaves*

**Assembly format**

.. code-block::

   `(` $operands `)` attr-dict `:` functional-type($operands, $result)

**Operands**

- ``$operands`` *(variadic)*: dynamic elements

**Results**

- ``$result``: assembled shape

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Constructs a ``!cute.shape`` by supplying runtime values for its
dynamic slots. The result type specifies the full shape structure;
static extents are encoded in the type itself and require no
operand. Each ``?`` (dynamic) extent in the result type requires
exactly one operand, supplied left-to-right in depth-first traversal
order across the shape tree. Zero operands are valid when the
result type has no ``?`` slots (fully static).

**Preconditions**

- Each operand fills one ``?`` slot, in depth-first left-to-right order:

  - ``i32`` or ``i64`` fills a scalar ``?`` extent; bit-width must match
    the slot's width annotation.

  - rank-1 ``!cute.shape`` fills a nested sub-shape slot; higher-rank
    sub-shapes are rejected.

- The number of operands equals the number of ``?`` slots in the
  result type.

**Postconditions**

- Result is the ``!cute.shape`` whose structure matches the result
  type and whose dynamic leaves take the operand values.

**Examples**

.. code-block::

   %s = cute.make_shape() : () -> !cute.shape<"(128,64)">
   %s = cute.make_shape(%m, %n) : (i32, i32) -> !cute.shape<"(?,?)">
   %s = cute.make_shape(%m) : (i32) -> !cute.shape<"(?,64)">


**Errors**

- too few operands for the `?` slots

  .. code-block::

     %0 = cute.make_shape(%n) : (i64) -> !cute.shape<"(?,?)">


- too many operands for the `?` slots

  .. code-block::

     %0 = cute.make_shape(%m, %n) : (i64, i64) -> !cute.shape<"(?,3)">


- operand width mismatch (i64 supplied for an i32-width `?` slot)

  .. code-block::

     %0 = cute.make_shape(%n) : (i64) -> !cute.shape<"?">


- non-integer-leaf operand rejected at ODS level

  .. code-block::

     %0 = cute.make_shape(%t) : (!cute.shape<"(?,?)">) -> !cute.shape<"(?,?)">



.. _op-cute.make_stride:

``cute.make_stride``
^^^^^^^^^^^^^^^^^^^^

*Construct a Stride from its runtime-valued leaves*

**Assembly format**

.. code-block::

   `(` $operands `)` attr-dict `:` functional-type($operands, $result)

**Operands**

- ``$operands`` *(variadic)*: dynamic elements

**Results**

- ``$result``: assembled stride

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Constructs a ``!cute.stride`` by supplying runtime values for its
dynamic slots. The result type specifies the full stride structure;
static strides and static scaled-basis elements (e.g. ``4@0``) are
encoded in the type itself and require no operand. Each ``?`` slot —
including the dynamic scale component of a ``?@M`` (dynamic-scale,
static-basis) element — requires exactly one operand, supplied
left-to-right in depth-first traversal order. Zero operands are
valid when the result type has no ``?`` slots (fully static).

**Preconditions**

- Each operand fills one ``?`` slot, in depth-first left-to-right order:

  - ``i32`` or ``i64`` fills a scalar ``?`` stride or the dynamic scale of
    a ``?@M`` element; bit-width must match the slot's width annotation.

  - rank-1 ``!cute.stride`` fills a nested sub-stride slot; higher-rank
    sub-strides are rejected.

- The number of operands equals the number of ``?`` slots in the
  result type.

**Postconditions**

- Result is the ``!cute.stride`` whose structure matches the result
  type and whose dynamic leaves take the operand values.

**Examples**

.. code-block::

   %d = cute.make_stride() : () -> !cute.stride<"(1,4)">
   %d = cute.make_stride(%s) : (i32) -> !cute.stride<"?">
   %d = cute.make_stride(%s) : (i32) -> !cute.stride<"?@0">
   %d = cute.make_stride() : () -> !cute.stride<"4@0">


**Errors**

- too few operands for the `?` slots

  .. code-block::

     %0 = cute.make_stride(%s) : (i64) -> !cute.stride<"(?,?)">


- too many operands for the `?` slots

  .. code-block::

     %0 = cute.make_stride(%s, %t) : (i64, i64) -> !cute.stride<"(1,?)">


- operand width mismatch (i64 supplied for an i32-width `?` slot)

  .. code-block::

     %0 = cute.make_stride(%s) : (i64) -> !cute.stride<"?">


- non-integer-leaf operand rejected at ODS level

  .. code-block::

     %0 = cute.make_stride(%t) : (!cute.stride<"(?,?)">) -> !cute.stride<"(?,?)">



.. _op-cute.make_tile:

``cute.make_tile``
^^^^^^^^^^^^^^^^^^

*Construct a Tile from its runtime-valued integer leaves*

**Assembly format**

.. code-block::

   `(` $operands `)` attr-dict `:` functional-type($operands, $result)

**Operands**

- ``$operands`` *(variadic)*: dynamic integer leaves

**Results**

- ``$result``: assembled tile

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Constructs a ``!cute.tile`` by supplying runtime values for the
dynamic integer leaves of its constituent layouts. A tile is a
fixed-length tuple of layout slots and underscore (``_``) slots,
fully described by the result type.

Operands are assigned to ``?`` leaves across all constituent layouts
in left-to-right, depth-first order: all dynamic leaves of the
first layout slot are filled first, then the second, and so on.
Underscore slots and fully-static layout slots carry no leaves and
consume no operands. Zero operands are valid when every layout
slot is fully static or underscore.

**Preconditions**

- Each operand fills one ``?`` leaf, in left-to-right depth-first
  order across all constituent layout shapes and strides:

  - ``i32`` or ``i64`` fills a scalar ``?`` leaf; bit-width must match
    the slot's width annotation.

  - rank-1 ``!cute.int_tuple`` fills a nested sub-tuple integer leaf;
    higher-rank sub-tuples are rejected.

- Underscore (``_``) slots and fully-static layout slots consume no
  operand.

- The number of operands equals the total count of ``?`` leaves
  across all constituent layout shapes and strides in the result
  type.

**Postconditions**

- Result is the ``!cute.tile`` whose layout slots, underscores, and
  dynamic-leaf values match the result type with operands filling
  the ``?`` slots in order.

**Examples**

.. code-block::

   // fully static — no operands
   %t = cute.make_tile() : () -> !cute.tile<"[(2,3):(1,2)]">

   // underscore slot — no operands
   %t = cute.make_tile() : () -> !cute.tile<"[_]">

   // static layout + underscore — no operands
   %t = cute.make_tile() : () -> !cute.tile<"[(2,3):(1,2);_]">

   // one dynamic layout with 2 leaves (one shape ?, one stride ?)
   %t = cute.make_tile(%m, %n) : (i32, i32) -> !cute.tile<"[(?,3):(1,?)]">

   // dynamic layout + underscore (underscore contributes no leaves)
   %t = cute.make_tile(%m, %n) : (i32, i32) -> !cute.tile<"[(?,3):(1,?);_]">

   // two dynamic layouts, 2 leaves each
   %t = cute.make_tile(%m, %n, %p, %q) : (i32, i32, i32, i32)
          -> !cute.tile<"[(?,3):(1,?);(?,2):(1,?)]">


**Errors**

- too many operands for a fully-static tile

  .. code-block::

     %t = cute.make_tile(%n) : (i64) -> !cute.tile<"[(2,3):(1,2)]">


- too many operands for an underscore tile

  .. code-block::

     %t = cute.make_tile(%n) : (i64) -> !cute.tile<"[_]">


- too few operands for a dynamic layout slot

  .. code-block::

     %t = cute.make_tile() : () -> !cute.tile<"[(?,3):(1,?)]">


- operand width mismatch (i64 supplied for i32-width `?` leaf)

  .. code-block::

     %t = cute.make_tile(%n) : (i64) -> !cute.tile<"[?:1]">


- depth>0 int_tuple operand rejected at ODS level

  .. code-block::

     %tile = cute.make_tile(%t)
               : (!cute.int_tuple<"(?,?)">) -> !cute.tile<"[(?,?):(1,?)]">



.. _op-cute.static:

``cute.static``
^^^^^^^^^^^^^^^

*Materialize a compile-time constant CuTe value*

**Assembly format**

.. code-block::

   attr-dict `:` type($result)

**Results**

- ``$result``: static CuTe value

**Traits**: ``AlwaysSpeculatableImplTrait``, ``ConstantLike``

**Interfaces**: ``ConditionallySpeculatable``

Description
"""""""""""


Produces a compile-time constant whose value is fully encoded in the
result type. Takes no operands. Two ``cute.static`` ops with identical
result types represent the same value.

**Preconditions**

- None — this op takes no operands.

**Postconditions**

- The result is a static CuTe type with no dynamic (``?``) slots.
- Accepts any of ``!cute.layout``, ``!cute.shape``, ``!cute.stride``,
  ``!cute.tile``, ``!cute.swizzle``, ``!cute.composed_layout``,
  ``!cute.int_tuple``, or ``!cute.coord`` (each with no ``?`` slots).

**Examples**

.. code-block::

   %0 = cute.static : !cute.layout<"(2,3):(1,2)">
   %1 = cute.static : !cute.stride<"(1,4)">
   %2 = cute.static : !cute.shape<"5">



Advanced Constructor Operations

Read individual modes, leaves, or scalar values out of cute values.

.. _op-cute.composed_get_inner:

``cute.composed_get_inner``
^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Extract the inner (A) component from a composed layout*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input) `->` type($result)

**Operands**

- ``$input``: input composed layout

**Results**

- ``$result``: inner A component

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the inner A component of the input ``!cute.composed_layout``.
A composed layout has the form ``A o offset o B``; this op extracts A.

**Preconditions**

- ``$input`` is ``!cute.composed_layout``.

**Postconditions**

- Result is ``!cute.layout`` when A is an affine layout, or
  ``!cute.swizzle`` when A is a swizzle.

- Result equals the A component of the input.

**Examples**

.. code-block::

   // Affine A: result is !cute.layout
   %a = cute.composed_get_inner(%cl)
          : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
         -> !cute.layout<"(4,5):(1,4)">

   // Swizzle A: result is !cute.swizzle
   %a = cute.composed_get_inner(%cl)
          : !cute.composed_layout<"S<3,5,4> o 0 o (8,4):(1,8)">
         -> !cute.swizzle<"S<3,5,4>">


**Errors**

- operand is not a composed_layout

  .. code-block::

     %a = cute.composed_get_inner(%l)
            : !cute.layout<"(4,5):(1,4)"> -> !cute.layout<"(4,5):(1,4)">


- declared result kind does not match the input's A

  .. code-block::

     // (input has affine A, but result is declared as a swizzle)
     %a = cute.composed_get_inner(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
           -> !cute.swizzle<"S<3,5,4>">


- declared result value does not match the input's A

  .. code-block::

     %a = cute.composed_get_inner(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
           -> !cute.layout<"(2,3):(1,2)">



.. _op-cute.composed_get_offset:

``cute.composed_get_offset``
^^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Extract the offset from a composed layout*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input) `->` qualified(type($result))

**Operands**

- ``$input``: input composed layout

**Results**

- ``$result``: offset

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the offset ``!cute.int_tuple`` from the input
``!cute.composed_layout`` ``A o offset o B``.

**Preconditions**

- ``$input`` is ``!cute.composed_layout``.

**Postconditions**

- Result is ``!cute.int_tuple`` and equals the offset component of
  the input.

**Examples**

.. code-block::

   // Scalar offset
   %off = cute.composed_get_offset(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
           -> !cute.int_tuple<"2">

   // Tuple offset with a scaled-basis outer B
   %off = cute.composed_get_offset(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o (0,0) o (2,3):(1@0,1@1)">
           -> !cute.int_tuple<"(0,0)">


**Errors**

- operand is not a composed_layout

  .. code-block::

     %off = cute.composed_get_offset(%l)
              : !cute.layout<"(4,5):(1,4)"> -> !cute.int_tuple<"2">


- declared result offset does not match the input's offset

  .. code-block::

     %off = cute.composed_get_offset(%cl)
              : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
             -> !cute.int_tuple<"5">



.. _op-cute.composed_get_outer:

``cute.composed_get_outer``
^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Extract the outer (B) layout from a composed layout*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input) `->` type($result)

**Operands**

- ``$input``: input composed layout

**Results**

- ``$result``: outer B layout

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the outer B ``!cute.layout`` from the input
``!cute.composed_layout`` ``A o offset o B``.

**Preconditions**

- ``$input`` is ``!cute.composed_layout``.

**Postconditions**

- Result is ``!cute.layout`` and equals the B (outer) component of
  the input.

**Examples**

.. code-block::

   %b = cute.composed_get_outer(%cl)
          : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
         -> !cute.layout<"(2,3):(1,2)">


**Errors**

- operand is not a composed_layout

  .. code-block::

     %b = cute.composed_get_outer(%l)
            : !cute.layout<"(2,3):(1,2)"> -> !cute.layout<"(2,3):(1,2)">


- declared result layout does not match the input's B

  .. code-block::

     %b = cute.composed_get_outer(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
           -> !cute.layout<"(4,5):(1,4)">



.. _op-cute.get_layouts_from_tile:

``cute.get_layouts_from_tile``
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Extract all layout elements from a tile, skipping underscores*

**Assembly format**

.. code-block::

   `(` $tile `)` attr-dict `:` type($tile)

**Operands**

- ``$tile``: input tile

**Results**

- ``$layouts`` *(variadic)*: extracted layout slots

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Flattens the input ``!cute.tile`` (in DFS / left-to-right order) and
returns one ``!cute.layout`` per layout leaf. Underscore elements are
silently skipped — the result count equals the number of layout
leaves in the tile.

**Preconditions**

- ``$tile`` is ``!cute.tile``.
- ``$tile`` contains at least one layout leaf. An all-underscore
  tile (no layout slots at any depth) is a use error — there is
  nothing to extract.

**Postconditions**

- The number of results equals the number of layout leaves of
  ``$tile`` after recursive flattening (underscore slots are
  skipped at every depth). The result is always non-empty.

- Each result is ``!cute.layout`` and equals the corresponding
  layout leaf of ``$tile`` in DFS left-to-right order.

- Layout leaves preserve their original values verbatim — static,
  dynamic, and scaled-basis strides all flow through unchanged.

**Examples**

.. code-block::

   // Two-layout tile: two results
   %l0, %l1 = cute.get_layouts_from_tile(%t)
                : !cute.tile<"[(4,8):(1,4);(2,3):(1,2)]">
   // %l0 : !cute.layout<"(4,8):(1,4)">, %l1 : !cute.layout<"(2,3):(1,2)">

   // Tile with an underscore: only layout slots are returned
   %l0 = cute.get_layouts_from_tile(%t) : !cute.tile<"[4:1;_]">
   // %l0 : !cute.layout<"4:1">

   // Nested tile (tile of tiles): DFS flatten yields all layout
   // leaves in left-to-right order, recursively skipping underscores.
   %l0, %l1, %l2, %l3 = cute.get_layouts_from_tile(%t)
                         : !cute.tile<"[4:1;[2:1;3:1];5:1]">
   // %l0 : !cute.layout<"4:1">, %l1 : !cute.layout<"2:1">,
   // %l2 : !cute.layout<"3:1">, %l3 : !cute.layout<"5:1">

   // Tile with a dynamic layout slot: the dynamic positions flow
   // through to the layout result.
   %l0, %l1 = cute.get_layouts_from_tile(%t)
                : !cute.tile<"[(?,4):(1,?);(2,3):(1,2)]">
   // %l0 : !cute.layout<"(?,4):(1,?)">, %l1 : !cute.layout<"(2,3):(1,2)">


**Errors**

- parser-level type constraint: operand is not a ``!cute.tile``

  .. code-block::

     %l0 = cute.get_layouts_from_tile(%l) : !cute.layout<"(4,8):(1,4)">


- result-binding count: declared result count does not match the
  layout-leaf count of the tile

  .. code-block::

     %l0 = cute.get_layouts_from_tile(%t) : !cute.tile<"[4:1;_;(2,3):(1,2)]">


- all-underscore tile has no layout leaves

  .. code-block::

     cute.get_layouts_from_tile(%t) : !cute.tile<"[_;_]">



.. _op-cute.get_leaves:

``cute.get_leaves``
^^^^^^^^^^^^^^^^^^^

*Flatten a tuple-like value and return one result per leaf*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input)

**Operands**

- ``$input``: tuple-like input

**Results**

- ``$results`` *(variadic)*: depth-1 leaf values

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Recursively flattens all nesting levels of input and returns one SSA value
per scalar leaf. Result count and per-result types are fully inferred from
the input type.

Per input kind, the leaf result type is:

- ``!cute.shape`` → one ``!cute.shape`` per scalar leaf.
- ``!cute.stride`` → one ``!cute.stride`` per scalar leaf.
- ``!cute.coord`` → one ``!cute.coord`` per scalar leaf.
- ``!cute.int_tuple`` → one ``!cute.int_tuple`` per scalar leaf.
- ``!cute.tile`` → one ``!cute.layout`` per layout slot, one
  ``!cute.tile<"_">`` per underscore slot.

**Preconditions**

- ``$input`` is one of ``!cute.shape``, ``!cute.stride``, ``!cute.coord``,
  ``!cute.int_tuple``, or ``!cute.tile``.

- The total leaf count is fully determined at compile time from the CuTe
  type encoded in ``$input``'s MLIR type attribute (e.g., `!cute.shape<"(4,(2,3))">`
  statically encodes 3 leaves).

**Postconditions**

- The number of results equals the total number of scalar leaves in ``$input``,
  produced by recursively flattening all levels of nesting.

- Each result type matches the per-input-kind rules above.

**Examples**

.. code-block::

   // shape: each leaf is !cute.shape — (4,(2,3)) flattens to three scalars
   %s0, %s1, %s2 = cute.get_leaves(%sh) : !cute.shape<"(4,(2,3))">
   // %s0 : !cute.shape<"4">, %s1 : !cute.shape<"2">, %s2 : !cute.shape<"3">

   // stride: each leaf is !cute.stride
   %t0, %t1 = cute.get_leaves(%st) : !cute.stride<"(1,4)">
   // %t0 : !cute.stride<"1">, %t1 : !cute.stride<"4">

   // coord: each leaf is !cute.coord
   %c0, %c1 = cute.get_leaves(%co) : !cute.coord<"(0,1)">
   // %c0 : !cute.coord<"0">, %c1 : !cute.coord<"1">

   // int_tuple: each leaf is !cute.int_tuple
   %i0, %i1 = cute.get_leaves(%it) : !cute.int_tuple<"(1,2)">
   // %i0 : !cute.int_tuple<"1">, %i1 : !cute.int_tuple<"2">

   // tile: layout slots → !cute.layout, underscore slots → !cute.tile<"_">
   %l0, %l1 = cute.get_leaves(%ti) : !cute.tile<"[(4,8):(1,4);(2,3):(1,2)]">
   // %l0 : !cute.layout<"(4,8):(1,4)">, %l1 : !cute.layout<"(2,3):(1,2)">

   %l0, %u, %l1 = cute.get_leaves(%ti) : !cute.tile<"[(4,8):(1,4);_;(2,3):(1,2)]">
   // %l0 : !cute.layout<"(4,8):(1,4)">, %u : !cute.tile<"_">, %l1 : !cute.layout<"(2,3):(1,2)">


**Errors**

- input is not one of the allowed CuTe tuple types (e.g. composed_layout)

  .. code-block::

     %r:1 = cute.get_leaves(%cl)
              : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">


- declared result count does not match the input's leaf count

  .. code-block::

     %r:2 = cute.get_leaves(%sh) : !cute.shape<"(4,(2,3))">



.. _op-cute.get:

``cute.get``
^^^^^^^^^^^^

*Extract a mode from a CuTe value by hierarchical index*

**Assembly format**

.. code-block::

   (`<` $mode^ `>`)? `(` $input `)` attr-dict
       `:` type($input) `->` type($result)

**Operands**

- ``$input``: input value
- ``$mode`` *(attribute, DenseI32ArrayAttr)*: hierarchical mode path

**Results**

- ``$result``: result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Extracts the element at the given hierarchical mode path. Modes
are applied sequentially: ``<[1, 0]>`` means "take element 1, then
element 0 from that". If ``mode`` is absent (or ``<[]>`` — the empty
mode array), the input value is returned unchanged (identity).
The result is always the same kind as the input.

Per input kind:

- ``!cute.shape``, ``!cute.int_tuple``, ``!cute.coord``,
  ``!cute.stride``, ``!cute.layout`` — descend into the
  nested-tuple structure; the result is the sub-element at
  the chosen path.

- ``!cute.tile`` — ``cute.get<[N]>`` selects slot N of the tile.
  The result is a tile wrapping the selected slot, not the
  bare layout / underscore. Picking a layout slot yields
  ``!cute.tile<"L">``; picking an underscore slot yields
  ``!cute.tile<"_">``. Component extraction with the slot's bare
  type can be done by composing with ``cute.get_leaves``.

- ``!cute.composed_layout`` — ``cute.get<[N]>`` slices the
  composed layout at sub-mode N of the outer B layout,
  preserving the inner A component and offset. The result is
  a new composed layout
  ``A o offset o get<[N]>(B)``, *not* a Layout. Component
  extraction (inner A / offset / outer B alone) belongs to
  ``cute.composed_get_inner`` / ``_offset`` / ``_outer``.

**Preconditions**

- ``$input`` is one of ``!cute.shape``, ``!cute.int_tuple``,
  ``!cute.coord``, ``!cute.stride``, ``!cute.tile``, ``!cute.layout``,
  or ``!cute.composed_layout``.

- Each index in ``$mode`` is in ``[0, rank)`` at its level
  (negative indices and out-of-range indices are rejected).

- For composed_layout inputs, ``$mode`` indices are interpreted
  against the outer B layout's rank at each level.

- ``$mode`` indices stop at a scalar leaf. Once the descent
  resolves to a scalar (a shape/stride/coord/int_tuple leaf,
  or a rank-1 layout / composed_layout whose shape is a
  scalar), no further indices are accepted: ``cute.get<[0, 0]>``
  on ``!cute.shape<"(4,2)">`` is a use error, not a no-op.

**Postconditions**

- Result has the same kind as ``$input`` (shape→shape,
  tile→tile, composed_layout→composed_layout, etc.). The result
  is always one of the cute kinds; to extract a raw ``i32``,
  ``i64``, or ``index`` for arithmetic, feed the ``cute.get`` result
  through ``cute.get_scalars``.

- Result equals ``$input`` when ``$mode`` is absent or empty
  (identity).

- Otherwise, result equals the element reached by descending
  through ``$mode`` indices in order — with the kind-specific
  semantics described above.

**Examples**

.. code-block::

   // Single-level: extract mode 1 from a shape — result is the nested tuple at that mode
   %r = cute.get<[1]> (%s) : !cute.shape<"(4,(2,3))"> -> !cute.shape<"(2,3)">

   // Single-level: extract mode 0 from a layout — scalar leaf
   %r = cute.get<[0]> (%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.layout<"4:1">

   // Two-level: mode 1 → then mode 0 of that — (4,(2,3)) → (2,3) → 2
   %r = cute.get<[1, 0]> (%s) : !cute.shape<"(4,(2,3))"> -> !cute.shape<"2">

   // Two-level: mode 0 → then mode 1 of that — ((4,2),(3,8)):((1,4),(8,24)) → (4,2):(1,4) → 2:4
   %r = cute.get<[0, 1]> (%l)
          : !cute.layout<"((4,2),(3,8)):((1,4),(8,24))"> -> !cute.layout<"2:4">

   // Three-level: mode 1 → mode 1 → mode 0 — (4,(2,(3,5))) → (2,(3,5)) → (3,5) → 3
   %r = cute.get<[1, 1, 0]> (%s) : !cute.shape<"(4,(2,(3,5)))"> -> !cute.shape<"3">

   // Three-level: mode 0 → mode 0 → mode 1 — deeply nested layout → scalar leaf
   %r = cute.get<[0, 0, 1]> (%l)
          : !cute.layout<"(((4,2),3),(8,16)):(((1,4),8),(24,384))"> -> !cute.layout<"2:4">

   // Identity get (no mode) — returns input unchanged
   %r = cute.get (%s) : !cute.shape<"(4,2)"> -> !cute.shape<"(4,2)">

   // Identity get (empty mode `<[]>`) — same as no-mode form
   %r = cute.get<[]> (%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.layout<"(4,8):(1,4)">

   // Tile operand: result is a single-slot tile, not a bare layout
   %r = cute.get<[0]> (%t)
          : !cute.tile<"[(4,8):(1,4);(2,3):(1,2)]"> -> !cute.tile<"(4,8):(1,4)">

   // Tile operand picking an underscore slot
   %r = cute.get<[1]> (%t)
          : !cute.tile<"[(4,8):(1,4);_]"> -> !cute.tile<"_">

   // ComposedLayout operand: mode 0 applied to outer B, inner+offset preserved
   //   outer B = (2,3):(1,2); get<[0]>(B) = 2:1
   %r = cute.get<[0]> (%cl)
          : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
         -> !cute.composed_layout<"(4,5):(1,4) o 2 o 2:1">


**Errors**

- top-level mode index out of range

  .. code-block::

     %r = cute.get<[5]>(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.layout<"(4,8):(1,4)">


- negative mode index is out of range

  .. code-block::

     %r = cute.get<[-1]>(%i) : !cute.int_tuple<"(1,2,3)"> -> !cute.int_tuple<"1">


- nested mode index out of range

  .. code-block::

     %r = cute.get<[1, 5]>(%s) : !cute.shape<"(4,(2,3))"> -> !cute.shape<"2">


- descent past a scalar leaf — after `[0]` the value is

  .. code-block::

     // the scalar `4`, so the trailing `0` has no scalar to descend into
     %r = cute.get<[0, 0]>(%s) : !cute.shape<"(4,2)"> -> !cute.shape<"4">



.. _op-cute.get_scalars:

``cute.get_scalars``
^^^^^^^^^^^^^^^^^^^^

*Extract scalar integers from a CuTe value*

**Assembly format**

.. code-block::

   (`<` `{` `only_dynamic` `}` `>` $only_dynamic^)? `(` $cute_value `)` attr-dict
       `:` qualified(type($cute_value))

**Operands**

- ``$cute_value``: CuTe value
- ``$only_dynamic`` *(attribute, UnitAttr)*: if set, return only the dynamic leaves

**Results**

- ``$scalars`` *(variadic)*: extracted scalar integers

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the scalar integers stored in a CuTe value. Result count
and per-result types are inferred from the input type and from the
``only_dynamic`` attribute.

When ``only_dynamic`` is set, only runtime-dynamic scalars are
returned. When absent, all scalars (static and dynamic) are
returned.

Per input kind:

- ``!cute.shape``, ``!cute.stride``, ``!cute.coord``, ``!cute.int_tuple``,
  ``!cute.tile`` — each integer leaf emits one ``i32`` or ``i64``
  result, with width matching the leaf's ``?`` annotation (default
  ``i32`` for ``?``). Static leaves emit ``i64`` when ``only_dynamic`` is
  absent.

- ``!cute.layout`` — scalars from shape, then scalars from stride.
- ``!cute.composed_layout`` — scalars from A, then offset, then B.
- ``!cute.swizzle`` — always three ``i32`` results: ``num_bits``,
  ``num_base``, ``num_shift``.

**Preconditions**

- ``$cute_value`` is any CuTe type.

**Postconditions**

- The number of results equals the count of scalars selected from
  ``$cute_value`` per the rules above (taking ``only_dynamic`` into
  account).

- Each result is ``i32`` or ``i64``, with width matching the source
  leaf annotation.

**Examples**

.. code-block::

   // All scalars from a mixed layout: 2 static i64 + 2 dynamic i32
   %s0, %s1, %d0, %d1 = cute.get_scalars (%l) : !cute.layout<"(4,8):(1,?)">

   // Only dynamic scalars: 1 dynamic i32
   %d = cute.get_scalars<{only_dynamic}> (%l) : !cute.layout<"(4,?):(1,4)">

   // Swizzle: always 3 x i32
   %nb, %nbase, %ns = cute.get_scalars (%sw) : !cute.swizzle<"S<3,5,4>">


**Errors**

- declared result count does not match the input's scalar count

  .. code-block::

     %s:1 = cute.get_scalars (%l) : !cute.layout<"(4,8):(1,?)">


- declared result width does not match a leaf's width annotation

  .. code-block::

     // (leaf is i32 ? but result is declared as i64)
     %d = cute.get_scalars<{only_dynamic}> (%l) : !cute.layout<"(4,?):(1,4)">
     // ... where the produced %d is bound to an i64 by the surrounding context



.. _op-cute.get_shape:

``cute.get_shape``
^^^^^^^^^^^^^^^^^^

*Extract the shape from a layout, composed layout, or tile*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input) `->` qualified(type($result))

**Operands**

- ``$input``: input layout or tile

**Results**

- ``$result``: shape

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the ``!cute.shape`` describing the domain of the input.

Per input kind:

- ``!cute.layout`` — the layout's shape (e.g. ``(4,8)`` from
  ``(4,8):(1,4)``).

- ``!cute.composed_layout`` — the shape of the outer B layout.
- ``!cute.tile`` — a shape whose i-th element is the shape of the
  i-th layout in the tile (e.g. ``[(4,8):(1,4);(2,3):(1,2)]`` →
  ``((4,8),(2,3))``).

Tile operands with underscore (``_``) slots are rejected: an
underscore slot has no shape, so it cannot contribute an
element to the extracted shape. If a tile has underscore slots,
project the layout slots out first via
``cute.get_layouts_from_tile`` and then take the shape of each
layout.

**Preconditions**

- ``$input`` is one of ``!cute.layout``, ``!cute.composed_layout``, or
  ``!cute.tile``.

- If ``$input`` is a ``!cute.tile``, it must contain no underscore
  slots.

**Postconditions**

- Result is ``!cute.shape``.
- Result equals the layout's shape (for layout / composed_layout
  inputs) or the per-slot shape tuple (for tile inputs).

**Examples**

.. code-block::

   // From a layout
   %s = cute.get_shape(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.shape<"(4,8)">

   // From a composed layout — returns B layout's shape
   %s = cute.get_shape(%cl)
          : !cute.composed_layout<"S<3,5,4> o 0 o (8,4):(1,8)">
         -> !cute.shape<"(8,4)">

   // From a tile — one shape entry per layout slot
   %s = cute.get_shape(%t)
          : !cute.tile<"[(4,8):(1,4);(2,3):(1,2)]">
         -> !cute.shape<"((4,8),(2,3))">



**Errors**

- operand is not a layout, composed_layout, or tile

  .. code-block::

     %s = cute.get_shape(%x) : !cute.shape<"(4,8)"> -> !cute.shape<"(4,8)">


- declared result shape does not match the input's shape

  .. code-block::

     %s = cute.get_shape(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.shape<"(2,3)">


- tile operand contains an underscore slot

  .. code-block::

     %s = cute.get_shape(%t)
            : !cute.tile<"[(4,8):(1,4);_]"> -> !cute.shape<"((4,8),x)">



.. _op-cute.get_stride:

``cute.get_stride``
^^^^^^^^^^^^^^^^^^^

*Extract the stride from a layout*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` type($input) `->` qualified(type($result))

**Operands**

- ``$input``: input layout

**Results**

- ``$result``: stride

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns the ``!cute.stride`` of the input ``!cute.layout``.
Scaled-basis strides are preserved verbatim.

**Preconditions**

- ``$input`` is ``!cute.layout``.

**Postconditions**

- Result is ``!cute.stride`` and equals the input layout's stride.

**Examples**

.. code-block::

   // Plain integer strides
   %s = cute.get_stride(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.stride<"(1,4)">

   // Scaled-basis strides — preserved verbatim
   %s = cute.get_stride(%l)
          : !cute.layout<"(2,3):(1@0,1@1)"> -> !cute.stride<"(1@0,1@1)">


**Errors**

- operand is not a layout

  .. code-block::

     %s = cute.get_stride(%cl)
            : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)">
           -> !cute.stride<"(1,4)">


- declared result stride does not match the input layout's stride

  .. code-block::

     %s = cute.get_stride(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.stride<"(2,3)">



.. _op-cute.select:

``cute.select``
^^^^^^^^^^^^^^^

*Select top-level modes from a CuTe value*

**Assembly format**

.. code-block::

   `<` $mode `>` `(` $input `)` attr-dict
       `:` type($input) `->` type($result)

**Operands**

- ``$input``: input value
- ``$mode`` *(attribute, DenseI32ArrayAttr)*: top-level mode indices

**Results**

- ``$result``: result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns a new value of the same kind containing only the top-level
modes at the given indices. Indices may be repeated or permuted.
The result is always wrapped in a tuple, even when a single mode
is selected. Unlike ``cute.get``, modes apply only at the top level
— there is no hierarchical descent.

**Preconditions**

- ``$input`` is one of ``!cute.shape``, ``!cute.int_tuple``,
  ``!cute.coord``, ``!cute.stride``, ``!cute.tile``, ``!cute.layout``, or
  ``!cute.composed_layout``.

- Each index in ``$mode`` is in ``[0, rank)`` of ``$input`` (negative
  indices and out-of-range indices are rejected).

**Postconditions**

- Result has the same kind as ``$input``.
- Result has rank equal to ``len($mode)``; mode ``i`` of the result
  equals top-level mode ``$mode[i]`` of ``$input``, copied as-is.

Per input kind:

- ``!cute.shape``, ``!cute.int_tuple``, ``!cute.coord``, ``!cute.stride`` —
  pick the named top-level modes and tuple them in the given
  order. Repeated indices duplicate the picked mode.

- ``!cute.layout`` — pick top-level (shape, stride) pairs; each
  picked mode retains its original shape and stride values.

- ``!cute.tile`` — pick top-level slots; layout slots and
  underscore slots are treated uniformly (an underscore at the
  picked index appears as an underscore in the result tile).

- ``!cute.composed_layout`` — ``cute.select`` operates on the
  outer B layout, slicing it to the selected modes. The
  inner A and offset are preserved.

- Empty mode ``<[]>`` — produces a rank-0 result (e.g.
  ``!cute.shape<"()">``). Contrast with ``cute.get<[]>``, which is
  identity.

**Examples**

.. code-block::

   // Keep two non-contiguous modes from a rank-3 shape
   %r = cute.select<[0, 2]> (%s)
          : !cute.shape<"(4,2,8)"> -> !cute.shape<"(4,8)">

   // Permute two modes of a layout (indices reordered)
   %r = cute.select<[2, 0]> (%l)
          : !cute.layout<"(4,2,8):(1,4,8)"> -> !cute.layout<"(8,4):(8,1)">

   // Repeat a mode — mode 1 appears twice in the output
   %r = cute.select<[1, 1]> (%i)
          : !cute.int_tuple<"(1,2,3)"> -> !cute.int_tuple<"(2,2)">

   // Single-mode select wraps result in a tuple: (4,8):(1,4) → (4):(1)
   %r = cute.select<[0]> (%l)
          : !cute.layout<"(4,8):(1,4)"> -> !cute.layout<"(4):(1)">

   // Tile: pick two top-level slots; underscore slot at index 1 is
   // copied as-is when selected.
   %r = cute.select<[1, 2]> (%t)
          : !cute.tile<"[(4,8):(1,4);_;(2,3):(1,2)]">
         -> !cute.tile<"[_;(2,3):(1,2)]">

   // ComposedLayout: select applies to outer B; inner + offset kept.
   //   outer B = (2,3,8):(1,2,6); select<[2,0]>(B) = (8,2):(6,1)
   %r = cute.select<[2, 0]> (%cl)
          : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3,8):(1,2,6)">
         -> !cute.composed_layout<"(4,5):(1,4) o 2 o (8,2):(6,1)">

   // Empty mode produces a rank-0 result.
   %r = cute.select<[]> (%s)
          : !cute.shape<"(4,8)"> -> !cute.shape<"()">


**Errors**

- mode index out of range

  .. code-block::

     %r = cute.select<[5]>(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.layout<"(4):(1)">


- mode index out of range on a shape

  .. code-block::

     %r = cute.select<[3]>(%s) : !cute.shape<"(4,8,2)"> -> !cute.shape<"(4)">


- negative mode index

  .. code-block::

     %r = cute.select<[-1]>(%i) : !cute.int_tuple<"(1,2,3)"> -> !cute.int_tuple<"(1)">



.. _op-cute.to_coord:

``cute.to_coord``
^^^^^^^^^^^^^^^^^

*Cast a !cute.int_tuple to !cute.coord*

**Assembly format**

.. code-block::

   `(` $src `)` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

**Operands**

- ``$src``: source int_tuple

**Results**

- ``$result``: coord result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Reinterprets a `!cute.int_tuple` as a `!cute.coord` with identical
structure and integer values. The result type is inferred from the
source. Because the source is a `!cute.int_tuple`, the result never
contains underscore leaves (`_`); construct those via the
`!cute.coord` constructor path instead.

**Preconditions**

- `$src` is `!cute.int_tuple`.
- `$src` must not contain error elements.

**Postconditions**

- Result is `!cute.coord` with the same hierarchical structure and
  integer values as `$src`.
- Dynamic leaves (`?`) are preserved verbatim through the cast.

**Examples**

.. code-block::

   // Static.
   %c = cute.to_coord(%i) : !cute.int_tuple<"(0,1)"> -> !cute.coord<"(0,1)">

   // Dynamic leaf preserved.
   %c = cute.to_coord(%i) : !cute.int_tuple<"(?,3)"> -> !cute.coord<"(?,3)">

   // Nested — hierarchical structure preserved.
   %c = cute.to_coord(%i)
          : !cute.int_tuple<"(0,(1,2),3)"> -> !cute.coord<"(0,(1,2),3)">



.. _op-cute.to_int_tuple:

``cute.to_int_tuple``
^^^^^^^^^^^^^^^^^^^^^

*Cast a scalar-tuple CuTe type to !cute.int_tuple*

**Assembly format**

.. code-block::

   `(` $src `)` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

**Operands**

- ``$src``: source scalar-tuple

**Results**

- ``$result``: int_tuple result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Converts a scalar-tuple CuTe type to a ``!cute.int_tuple`` with
identical integer values. The result type is inferred from the
source.

**Preconditions**

- ``$src`` is one of ``!cute.shape``, ``!cute.stride``, ``!cute.coord``,
  or ``!cute.int_tuple``.

- ``$src`` must not contain scaled-basis strides (e.g. ``1@0``).
- ``$src`` must not contain underscore (``_``) elements (reachable
  via ``!cute.coord``, which accepts underscores).

- ``$src`` must not contain error elements.

**Postconditions**

- Result is ``!cute.int_tuple`` with the same hierarchical structure
  and integer values as ``$src`` (interpreting any wrapper kind —
  shape, stride, coord — as a plain integer tuple).
- Dynamic leaves (``?``) are preserved verbatim through the cast.

- ``int_tuple`` → ``int_tuple`` is the identity case (allowed; result
  equals source).

**Examples**

.. code-block::

   // Per-kind static casts.
   %it = cute.to_int_tuple(%s)  : !cute.shape<"(4,8)">  -> !cute.int_tuple<"(4,8)">
   %it = cute.to_int_tuple(%st) : !cute.stride<"(1,4)"> -> !cute.int_tuple<"(1,4)">
   %it = cute.to_int_tuple(%c)  : !cute.coord<"(0,1)">  -> !cute.int_tuple<"(0,1)">

   // Identity: int_tuple → int_tuple.
   %it = cute.to_int_tuple(%i) : !cute.int_tuple<"(2,3)"> -> !cute.int_tuple<"(2,3)">

   // Dynamic source — `?` leaves flow through.
   %it = cute.to_int_tuple(%s) : !cute.shape<"(?,8)"> -> !cute.int_tuple<"(?,8)">

   // Nested tuple — hierarchical structure preserved.
   %it = cute.to_int_tuple(%s)
           : !cute.shape<"(4,(2,3),8)"> -> !cute.int_tuple<"(4,(2,3),8)">

   // Dynamic stride leaf preserved through the cast.
   %it = cute.to_int_tuple(%st)
           : !cute.stride<"(8,?)"> -> !cute.int_tuple<"(8,?)">


**Errors**

- scalar source has a scaled-basis stride

  .. code-block::

     %r = cute.to_int_tuple(%s) : !cute.stride<"4@0"> -> !cute.int_tuple<"4">


- tuple source has a dynamic scaled-basis component

  .. code-block::

     %r = cute.to_int_tuple(%s) : !cute.stride<"(1,?@0)"> -> !cute.int_tuple<"(1,4)">


- coord source contains an underscore element

  .. code-block::

     %r = cute.to_int_tuple(%c) : !cute.coord<"(4,_)"> -> !cute.int_tuple<"(4,8)">



.. _op-cute.to_shape:

``cute.to_shape``
^^^^^^^^^^^^^^^^^

*Cast a !cute.int_tuple to !cute.shape*

**Assembly format**

.. code-block::

   `(` $src `)` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

**Operands**

- ``$src``: source int_tuple

**Results**

- ``$result``: shape result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Reinterprets a `!cute.int_tuple` as a `!cute.shape` with identical
structure and integer values. The result type is inferred from the
source.

**Preconditions**

- `$src` is `!cute.int_tuple`.
- `$src` must not contain error elements.

**Postconditions**

- Result is `!cute.shape` with the same hierarchical structure and
  integer values as `$src`.
- Dynamic leaves (`?`) are preserved verbatim through the cast.

**Examples**

.. code-block::

   // Static.
   %s = cute.to_shape(%i) : !cute.int_tuple<"(4,8)"> -> !cute.shape<"(4,8)">

   // Dynamic leaf preserved.
   %s = cute.to_shape(%i) : !cute.int_tuple<"(?,8)"> -> !cute.shape<"(?,8)">

   // Nested — hierarchical structure preserved.
   %s = cute.to_shape(%i)
          : !cute.int_tuple<"(4,(2,3),8)"> -> !cute.shape<"(4,(2,3),8)">



.. _op-cute.to_stride:

``cute.to_stride``
^^^^^^^^^^^^^^^^^^

*Cast a !cute.int_tuple to !cute.stride*

**Assembly format**

.. code-block::

   `(` $src `)` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

**Operands**

- ``$src``: source int_tuple

**Results**

- ``$result``: stride result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Reinterprets a `!cute.int_tuple` as a `!cute.stride` with identical
structure and integer values. The result type is inferred from the
source. Because the source is a `!cute.int_tuple`, the result never
contains scaled-basis leaves (`V@M`); construct those via the
`!cute.stride` constructor path instead.

**Preconditions**

- `$src` is `!cute.int_tuple`.
- `$src` must not contain error elements.

**Postconditions**

- Result is `!cute.stride` with the same hierarchical structure and
  integer values as `$src`.
- Dynamic leaves (`?`) are preserved verbatim through the cast.

**Examples**

.. code-block::

   // Static.
   %st = cute.to_stride(%i) : !cute.int_tuple<"(1,4)"> -> !cute.stride<"(1,4)">

   // Dynamic leaf preserved.
   %st = cute.to_stride(%i) : !cute.int_tuple<"(1,?)"> -> !cute.stride<"(1,?)">

   // Nested — hierarchical structure preserved.
   %st = cute.to_stride(%i)
           : !cute.int_tuple<"(1,(4,8),32)"> -> !cute.stride<"(1,(4,8),32)">



Layout Algebra Operations

Compute sizes, products, and coord/index conversions.

.. _op-cute.append_to_rank:

``cute.append_to_rank``
^^^^^^^^^^^^^^^^^^^^^^^

*Append element to input until rank reaches N*

**Assembly format**

.. code-block::

   `<` $rank `>` `(` $input `,` $element `)` attr-dict `:`
       qualified(type($input)) `,` qualified(type($element))

**Operands**

- ``$rank`` *(attribute, I32Attr)*: target rank N (must be > 0)
- ``$input``: input value to extend to rank N
- ``$element``: element used to pad up to rank N

**Results**

- ``$result``: extended value with rank N

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Repeatedly appends ``$element`` to the right of ``$input`` until the
result has rank ``$rank``. If ``$input`` already has rank >= ``$rank``, it
is returned unchanged. Useful for broadening an int_tuple, coord,
stride, shape, or layout to a target rank using an identity-like
filler (e.g. ``1`` for a shape, ``1:0`` for a layout).

**Preconditions**

- ``$rank`` is a positive integer attribute (``> 0``).
- ``$input`` and ``$element`` types are compatible:

  - For ``!cute.int_tuple``, ``!cute.coord``, ``!cute.stride``, ``!cute.shape``:
    ``$element`` must be the same kind as ``$input``.

  - For ``!cute.layout``: ``$element`` must be ``!cute.layout``.
  - For ``!cute.composed_layout``: ``$element`` must be ``!cute.layout``.

**Postconditions**

- ``$result`` has the same kind as ``$input``.
- ``rank($result) == max(rank($input), $rank)``.
- The leading ``rank($input)`` modes of ``$result`` equal ``$input``; any
  trailing modes equal ``$element``.

**Examples**

.. code-block::

   // Shape rank 2 -> rank 4, padding with shape<"1">.
   // Result: !cute.shape<"(4,8,1,1)">.
   %r = cute.append_to_rank<4> (%in, %e)
          : !cute.shape<"(4,8)">, !cute.shape<"1">

   // Layout rank 1 -> rank 3, padding with identity layout.
   // Result: !cute.layout<"(4,1,1):(1,0,0)">.
   %r = cute.append_to_rank<3> (%in, %e)
          : !cute.layout<"4:1">, !cute.layout<"1:0">

   // Dynamic shape rank 2 -> rank 4 — dynamic extents pass through unchanged.
   // Result: !cute.shape<"(?,8,1,1)">.
   %r = cute.append_to_rank<4> (%in, %e)
          : !cute.shape<"(?,8)">, !cute.shape<"1">

   // No-op: target rank (2) <= input rank (3) — input returned unchanged.
   // Result: !cute.shape<"(4,8,2)">.
   %r = cute.append_to_rank<2> (%in, %e)
          : !cute.shape<"(4,8,2)">, !cute.shape<"1">


**Errors**

- ``$rank`` is not positive.
- ``$input`` and ``$element`` are not compatible kinds.


.. _op-cute.cosize:

``cute.cosize``
^^^^^^^^^^^^^^^

*Compute codomain size of a layout-like value*

**Assembly format**

.. code-block::

   (`<` $mode^ `>`)? `(` $input `)` attr-dict `:`
       functional-type($input, $result)

**Operands**

- ``$input``: input layout-like value
- ``$mode`` *(attribute, DenseI32ArrayAttr)*: hierarchical mode path

**Results**

- ``$result``: computed cosize as int_tuple

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Computes the codomain size of ``$input``, optionally at a nested
``$mode``. For layouts with integer strides, the cosize of a layout
``l`` is the number of distinct indices ``l`` can produce
(``cosize(l) = l(size(l) - 1) + 1``). For layouts with scaled-basis strides,
the cosize is computed as the product of the hierarchical codomain coordinate,
``size(coshape(l))``, where coshape is the smallest bounding coordinate of the
codomain.

**Preconditions**

- ``$input`` is ``!cute.layout`` or ``!cute.composed_layout``.
- When ``$mode`` is present, each index in ``$mode`` is in
  ``[0, rank)`` at its level; ``$mode`` is traversed depth-first: ``mode[0]``
  selects a top-level element, ``mode[1]`` selects within that element's subtuple, etc.
  Each element of ``$mode`` is checked against the rank of the currently selected
  sub-object.

- When ``$input`` is ``!cute.composed_layout``, its inner layout
  must be a swizzle (not affine). ``cosize`` is undefined otherwise.

- ``$mode`` indices stop at a scalar leaf. Once the descent
  resolves to a rank-1 layout with a scalar shape, no further
  indices are accepted: ``cute.cosize<[1, 0]>`` on
  ``!cute.layout<"(4,3,2):(1,4,12)">`` is a use error — after
  ``[1]`` we land on sub-layout ``3:4`` (scalar), so the trailing
  ``0`` has no scalar to descend into. Use ``<[1]>`` instead.

**Postconditions**

- Result is ``!cute.int_tuple`` (a scalar).
- With no ``$mode``: result equals ``cosize($input)``.
- With ``$mode``: result equals the cosize of the sub-object selected by
  ``cute.get<$mode>($input)``, i.e. the product of all extents within
  the selected sub-mode.
- When ``$input`` is ``!cute.composed_layout`` (a o b), the result equals ``cosize(b)``.

**Examples**

.. code-block::

   %csz = cute.cosize (%layout) : (!cute.layout<"(4,3,2):(1,4,12)">) -> !cute.int_tuple<"24">
   %csz = cute.cosize<[1]> (%layout) : (!cute.layout<"(4,3,2):(1,4,12)">) -> !cute.int_tuple<"9">


**Errors**

- top-level mode index out of range

  .. code-block::

     %csz = cute.cosize<[3]> (%src)
              : (!cute.layout<"(4,3,2):(1,4,12)">) -> !cute.int_tuple<"4">


- descent past a scalar leaf — use `<[1]>` instead

  .. code-block::

     %csz = cute.cosize<[1, 0]> (%src)
              : (!cute.layout<"(4,3,2):(1,4,12)">) -> !cute.int_tuple<"9">



.. _op-cute.crd2idx:

``cute.crd2idx``
^^^^^^^^^^^^^^^^

*Pack a coordinate into a linear index given a shape*

**Assembly format**

.. code-block::

   `(` $coord `,` $shape `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$coord``: coordinate
- ``$shape``: shape

**Results**

- ``$result``: linear index

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Packs `$coord` into a single linear index using **implicit
column-major strides** derived from `$shape`. For a shape
`(s_0, s_1, ..., s_{n-1})` and a coord
`(c_0, c_1, ..., c_{n-1})` the result is

`c_0 + c_1 * s_0 + c_2 * s_0 * s_1 + ... + c_{n-1} * s_0 * ... * s_{n-2}`.

Nested coords / shapes recurse mode-by-mode. The result is
always a scalar `!cute.int_tuple`.

This is the structural pack — no layout / stride / order
operand. Use `cute.layout_eval` instead when the strides come
from a specific layout.

**Preconditions:**

- `$coord` is `!cute.coord` and congruent with `$shape`.
- `$shape` is `!cute.shape`.

**Postconditions:**

- `$result` is a scalar `!cute.int_tuple`.
- `$result` is static iff every leaf of `$coord` and every contributing extent of `$shape` is static.

**Examples:**

.. code-block::

   // Static coord (1,2) on shape (4,8) → 1*1 + 2*4 = 9.
   %idx = cute.crd2idx(%crd, %shape)
            : (!cute.coord<"(1,2)">, !cute.shape<"(4,8)">)
           -> !cute.int_tuple<"9">

   // Dynamic coord, static shape → result is dynamic.
   %idx = cute.crd2idx(%crd, %shape)
            : (!cute.coord<"(?,?)">, !cute.shape<"(4,8)">)
           -> !cute.int_tuple<"?">

   // Static coord, dynamic shape → result is dynamic (an unknown
   // extent affects the column-major weighting of later modes).
   %idx = cute.crd2idx(%crd, %shape)
            : (!cute.coord<"(1,2)">, !cute.shape<"(?,8)">)
           -> !cute.int_tuple<"?">


**Errors:**

- Coord rank does not match shape rank:

  .. code-block::

     // error: 'cute.crd2idx' op expects coord and shape to be congruent
     %idx = cute.crd2idx(%crd, %shape)
              : (!cute.coord<"(1,2,3)">, !cute.shape<"(4,8)">)
             -> !cute.int_tuple<"?">


- Declared result type does not match the inferred index:

  .. code-block::

     // error: 'cute.crd2idx' op inferred type(s) '!cute.int_tuple<"9">'
     //   are incompatible with return type(s) of operation
     //   '!cute.int_tuple<"7">'
     %idx = cute.crd2idx(%crd, %shape)
              : (!cute.coord<"(1,2)">, !cute.shape<"(4,8)">)
             -> !cute.int_tuple<"7">



.. _op-cute.idx2crd:

``cute.idx2crd``
^^^^^^^^^^^^^^^^

*Convert an index to a coordinate with respect to a shape*

**Assembly format**

.. code-block::

   `(` $index `,` $shape `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$index``: index (scalar int_tuple)
- ``$shape``: shape

**Results**

- ``$result``: coordinate

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Computes the coordinate corresponding to ``$index`` in colexicographic
(column-major) order with respect to ``$shape``. For a 1D shape ``s``, the
result is ``idx mod s``; for higher-rank shapes, the index is unraveled
digit-by-digit using the running products of ``$shape``'s extents.

**Preconditions**

- ``$index`` is either a scalar ``!cute.int_tuple`` (rank ≤ 1)
  *or* a multi-element int_tuple whose rank matches ``$shape``'s
  rank (the latter form pre-decomposes the linear index per
  mode).

- ``$shape`` is an ``!cute.shape``.

**Postconditions**

- ``$result`` is an ``!cute.coord`` whose profile matches ``$shape``.
- ``$result`` is fully static iff ``$index`` is static and every extent of
  ``$shape`` is static; otherwise every element of the result is dynamic.

**Examples**

.. code-block::

   // Static index and shape: coord resolved at compile time.
   %crd = cute.idx2crd(%idx, %shape)
            : (!cute.int_tuple<"9">, !cute.shape<"(4,8)">)
           -> !cute.coord<"(1,2)">

   // Dynamic index, static shape: result coord is all-dynamic.
   %crd = cute.idx2crd(%idx, %shape)
            : (!cute.int_tuple<"?">, !cute.shape<"(4,8)">)
           -> !cute.coord<"(?,?)">

   // Static index, dynamic extent: result coord is all-dynamic (an unknown
   // extent affects the column-major walk for both modes).
   %crd = cute.idx2crd(%idx, %shape)
            : (!cute.int_tuple<"9">, !cute.shape<"(?,8)">)
           -> !cute.coord<"(?,?)">


**Errors**

- ``$index``'s rank is greater than 1 and does not match ``$shape``'s rank.
- ``$result``'s profile is not congruent with ``$shape``.
- ``$result`` value disagrees with the column-major unravel of ``$index``
  against ``$shape`` when both sides are static.


.. _op-cute.increment_coord:

``cute.increment_coord``
^^^^^^^^^^^^^^^^^^^^^^^^

*Increment a coordinate by one within its bounding shape*

**Assembly format**

.. code-block::

   `(` $coord `,` $shape `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$coord``: coordinate to increment
- ``$shape``: bounding shape

**Results**

- ``$result``: incremented coordinate

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Increments ``$coord`` by 1 in column-major (row-first) order within the
coordinate space bounded by ``$shape``. The coordinate is treated as a
multi-dimensional digit, so when the lowest mode reaches its extent it
wraps to 0 and carries into the next mode.

**Preconditions**

- ``$coord`` and ``$shape`` are congruent (same hierarchical rank/depth
  structure).

- ``$coord`` contains no wildcard (``_``) elements.

**Postconditions**

- ``$result`` is an ``!cute.coord`` with the same profile as ``$coord``.
- When all inputs are static, ``$result`` equals the column-major
  successor of ``$coord`` modulo ``$shape`` (wrap-around at the highest
  mode is not specified — callers are expected to keep the iteration
  within ``size($shape)``).

- Any dynamic element in ``$coord`` or ``$shape`` makes the entire
  ``$result`` dynamic.

**Examples**

.. code-block::

   // Static: (1,2) + 1 in space (4,8) -> (2,2).
   %r = cute.increment_coord(%c, %s)
          : (!cute.coord<"(1,2)">, !cute.shape<"(4,8)">)
         -> !cute.coord<"(2,2)">

   // Static wrap: (3,2) + 1 in space (4,8) -> (0,3).
   %r = cute.increment_coord(%c, %s)
          : (!cute.coord<"(3,2)">, !cute.shape<"(4,8)">)
         -> !cute.coord<"(0,3)">

   // Dynamic coord: result has the same profile with dynamic values.
   %r = cute.increment_coord(%c, %s)
          : (!cute.coord<"(?,?)">, !cute.shape<"(4,8)">)
         -> !cute.coord<"(?,?)">


**Errors**

- ``$coord`` and ``$shape`` are not congruent.
- ``$coord`` contains a wildcard (``_``) element.
- ``$result``'s profile differs from ``$coord``'s profile.
- ``$result`` value disagrees with the column-major successor of ``$coord``
  modulo ``$shape`` when both sides are static.


.. _op-cute.layout_eval:

``cute.layout_eval``
^^^^^^^^^^^^^^^^^^^^

*Convert a coordinate to an index with respect to a layout*

**Assembly format**

.. code-block::

   `(` $coord `,` $layout `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$coord``: coordinate
- ``$layout``: layout or composed layout

**Results**

- ``$result``: index

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Computes the dot product of ``$coord`` with the strides of ``$layout``,
giving the index into the layout's codomain. For a layout with shape
``s`` and stride ``d``, this evaluates to ``sum_i coord[i] * d[i]``.

Result shape depends on the layout's stride form:

- *Plain integer strides* — ``$result`` is a scalar ``!cute.int_tuple``
  containing ``sum_i coord[i] * stride[i]``.

- *Scaled-basis strides* (``v@k`` or hierarchical ``v@k@...``) — each
  basis index ``@k`` projects its contribution into a separate axis
  of the codomain. ``$result`` is an ``!cute.int_tuple`` whose
  structure is induced by the basis dimensions of the stride,
  not by the shape's nesting. For flat basis (``(1@0,1@1)``) the
  result is rank-``numBasisDims``; for hierarchical basis
  (``(1@0@0,1@1@0)``) the result groups contributions by *outer*
  basis dim with inner basis dims as sub-tuples.

**Preconditions**

- ``$coord`` is either a leaf (scalar) ``!cute.coord`` *or* a non-leaf
  coord whose top-level rank matches ``$layout``'s shape rank
  (for composed layouts, the rank source is the inner B layout
  since the coord flows into B first: ``A(offset + B(coord))``).

- ``$layout`` is an ``!cute.layout`` or ``!cute.composed_layout``.

**Postconditions**

- ``$result`` is an ``!cute.int_tuple`` whose value equals the basis-
  aware dot product ``sum_i coord[i] * stride[i]``. Plain-stride
  layouts produce a scalar; basis-strided layouts produce a tuple
  grouped by basis dim as described above.

- ``$result`` is static iff every ``coord[i]`` is static and every
  contributing ``stride[i]`` is static; otherwise it is dynamic.

**Examples**

.. code-block::

   // Static coord and layout: index resolved at compile time.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(1,2)">, !cute.layout<"(4,8):(1,4)">)
           -> !cute.int_tuple<"9">

   // Dynamic coord, static layout: result is dynamic.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(?,?)">, !cute.layout<"(4,8):(1,4)">)
           -> !cute.int_tuple<"?">

   // Static coord, dynamic stride: result is dynamic (depends on the stride).
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(1,2)">, !cute.layout<"(4,8):(1,?)">)
           -> !cute.int_tuple<"?">

   // Flat scaled-basis stride: per-basis-dim tuple result. Each
   // natural-coord component lands in its matching basis dimension.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(1,2)">, !cute.layout<"(4,8):(1@0,1@1)">)
           -> !cute.int_tuple<"(1,2)">

   // Reversed flat basis: stride order is `(1@1,1@0)` so the first
   // coord component contributes to basis dim 1, the second to
   // basis dim 0. Result is per-basis-dim — i.e. `(2,1)` — not
   // per-coord-mode.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(1,2)">, !cute.layout<"(4,8):(1@1,1@0)">)
           -> !cute.int_tuple<"(2,1)">

   // Nested shape with flat-basis strides over three basis dims.
   // The natural coord `(1,2)` decomposes: rank-0 coord `1` walks
   // the nested mode-0 `(2,3)` as flat index 1 → `(1,0)`; rank-1
   // coord `2` is the mode-1 scalar. Per-basis result: `(1,0,2)`.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"(1,2)">,
               !cute.layout<"((2,3),4):((1@0,1@1),1@2)">)
           -> !cute.int_tuple<"(1,0,2)">

   // Hierarchical scaled-basis stride: result groups by *outer*
   // basis dim, not by outer shape mode.
   %idx = cute.layout_eval(%crd, %layout)
            : (!cute.coord<"0">,
               !cute.layout<"((2,3),4):((1@1@0,1@0@0),1@1)">)
           -> !cute.int_tuple<"((0,0),0)">


**Errors**

- ``$coord``'s profile is not congruent with ``$layout``'s shape.
- ``$result``'s value or structure disagrees with the basis-aware
  dot product when both sides are static.


.. _op-cute.prepend_to_rank:

``cute.prepend_to_rank``
^^^^^^^^^^^^^^^^^^^^^^^^

*Prepend element to input until rank reaches N*

**Assembly format**

.. code-block::

   `<` $rank `>` `(` $input `,` $element `)` attr-dict `:`
       qualified(type($input)) `,` qualified(type($element))

**Operands**

- ``$rank`` *(attribute, I32Attr)*: target rank N (must be > 0)
- ``$input``: input value to extend to rank N
- ``$element``: element used to pad up to rank N

**Results**

- ``$result``: extended value with rank N

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Repeatedly prepends ``$element`` to the left of ``$input`` until the
result has rank ``$rank``. If ``$input`` already has rank >= ``$rank``, it
is returned unchanged. Useful for broadening an int_tuple, coord,
stride, shape, or layout to a target rank using an identity-like
filler (e.g. ``1`` for a shape, ``1:0`` for a layout).

**Preconditions**

- ``$rank`` is a positive integer attribute (``> 0``).
- ``$input`` and ``$element`` types are compatible:

  - For ``!cute.int_tuple``, ``!cute.coord``, ``!cute.stride``, ``!cute.shape``:
    ``$element`` must be the same kind as ``$input``.

  - For ``!cute.layout``: ``$element`` must be ``!cute.layout``.
  - For ``!cute.composed_layout``: ``$element`` must be ``!cute.layout``.

**Postconditions**

- ``$result`` has the same kind as ``$input``.
- ``rank($result) == max(rank($input), $rank)``.
- The trailing ``rank($input)`` modes of ``$result`` equal ``$input``; any
  leading modes equal ``$element``.

**Examples**

.. code-block::

   // Shape rank 2 -> rank 4, padding with shape<"1">.
   // Result: !cute.shape<"(1,1,4,8)">.
   %r = cute.prepend_to_rank<4> (%in, %e)
          : !cute.shape<"(4,8)">, !cute.shape<"1">

   // Layout rank 1 -> rank 3, padding with identity layout.
   // Result: !cute.layout<"(1,1,4):(0,0,1)">.
   %r = cute.prepend_to_rank<3> (%in, %e)
          : !cute.layout<"4:1">, !cute.layout<"1:0">

   // Dynamic shape rank 2 -> rank 4 — dynamic extents pass through unchanged.
   // Result: !cute.shape<"(1,1,?,8)">.
   %r = cute.prepend_to_rank<4> (%in, %e)
          : !cute.shape<"(?,8)">, !cute.shape<"1">

   // No-op: target rank (2) <= input rank (3) — input returned unchanged.
   // Result: !cute.shape<"(4,8,2)">.
   %r = cute.prepend_to_rank<2> (%in, %e)
          : !cute.shape<"(4,8,2)">, !cute.shape<"1">


**Errors**

- ``$rank`` is not positive.
- ``$input`` and ``$element`` are not compatible kinds.


.. _op-cute.size:

``cute.size``
^^^^^^^^^^^^^

*Compute multiplicative size of a tuple/layout-like value*

**Assembly format**

.. code-block::

   (`<` $mode^ `>`)? `(` $input `)` attr-dict `:`
       functional-type($input, $result)

**Operands**

- ``$input``: input tuple/layout-like value
- ``$mode`` *(attribute, DenseI32ArrayAttr)*: hierarchical mode path

**Results**

- ``$result``: computed size as int_tuple

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Computes the multiplicative size of ``$input``, optionally at a
nested ``$mode``. The size is the product of all extents at the
selected mode (or of the whole input when ``$mode`` is absent).

For layouts and composed layouts, "size" refers to the domain
extents (shape), not the codomain (cosize).

**Preconditions**

- ``$input`` is one of ``!cute.shape``, ``!cute.int_tuple``,
  ``!cute.layout``, or ``!cute.composed_layout``.

- When ``$mode`` is present, each index in ``$mode`` is in
  ``[0, rank)`` at its level; ``$mode`` is traversed depth-first: ``mode[0]``
  selects a top-level element, ``mode[1]`` selects within that element's subtuple, etc.
  Each element of ``$mode`` is checked against the rank of the currently selected
  sub-object.

- ``$mode`` indices stop at a scalar leaf. Once the descent
  resolves to a scalar, no further indices are accepted:
  ``cute.size<[0, 0]>`` on ``!cute.shape<"(4,2)">`` is a use
  error, not a no-op.

**Postconditions**

- Result is a scalar ``!cute.int_tuple`` (rank 0, i.e. a single integer value).
- With no ``$mode``: result equals the product of all extents in
  ``$input``. For composed layouts, this is the size of the base
  layout's shape (Size of B, if the composed layout is ``A o O o B``).

- With ``$mode``: result equals the size of the sub-object selected by
  ``cute.get<$mode>($input)``, i.e. the product of all extents within
  the selected sub-mode.

**Examples**

.. code-block::

   %sz = cute.size<[1]> (%shape) : (!cute.shape<"(1,(2,4),9)">) -> !cute.int_tuple<"8">
   %sz = cute.size (%layout) : (!cute.layout<"(4,(16,32)):(1,(4,64))">) -> !cute.int_tuple<"2048">


**Errors**

- top-level mode index out of range

  .. code-block::

     %sz = cute.size<[3]> (%src) : (!cute.shape<"(4,3)">) -> !cute.int_tuple<"4">


- nested mode index out of range

  .. code-block::

     %sz = cute.size<[1, 5]> (%src) : (!cute.shape<"(4,(3,2))">) -> !cute.int_tuple<"3">


- descent past a scalar leaf — after `[0]` the value is

  .. code-block::

     // the scalar `4`, so the trailing `0` has no scalar to descend into
     %sz = cute.size<[0, 0]> (%src) : (!cute.shape<"(4,2)">) -> !cute.int_tuple<"4">



.. _op-cute.tuple_product_each:

``cute.tuple_product_each``
^^^^^^^^^^^^^^^^^^^^^^^^^^^

*Compute per-top-level-element product of an int_tuple or shape*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` functional-type($input, $result)

**Operands**

- ``$input``: input int_tuple or shape

**Results**

- ``$result``: per-element product (same kind as input)

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Applies ``product()`` to each top-level element of ``$input``, returning a
depth-1 value of the same kind as ``$input``. A scalar input is first
wrapped into a 1-element tuple. For example,
``tuple_product_each((2,(3,4))) = (2,12)`` and
``tuple_product_each((4,3,2)) = (4,3,2)``.

**Preconditions**

- ``$input`` is an ``!cute.int_tuple`` or ``!cute.shape``.

**Postconditions**

- ``$result`` kind matches ``$input`` kind (``shape -> shape``,
  ``int_tuple -> int_tuple``); it is a depth-1 value whose rank equals
  the rank of ``$input`` (or ``1`` if ``$input`` is a scalar).

- Each output element is the product of the corresponding top-level
  sub-tuple of ``$input``. A dynamic factor anywhere in a sub-tuple makes
  that output element dynamic.

**Examples**

.. code-block::

   // Static int_tuple: per-mode product, depth flattened to 1.
   %p = cute.tuple_product_each(%t)
          : (!cute.int_tuple<"(2,(3,4))">) -> !cute.int_tuple<"(2,12)">

   // Dynamic shape: each output mode is dynamic. Result kind matches
   // input kind (shape -> shape).
   %p = cute.tuple_product_each(%t)
          : (!cute.shape<"(?,(3,?))">) -> !cute.shape<"(?,?)">


**Errors**

- ``$result`` kind does not match ``$input`` kind.
- ``$result`` rank does not match ``$input`` rank (or is not ``1`` for scalar
  input).

- ``$result`` element does not equal the product of the corresponding
  top-level sub-tuple of ``$input``.


.. _op-cute.tuple_product:

``cute.tuple_product``
^^^^^^^^^^^^^^^^^^^^^^

*Compute the product of all leaf elements in an int_tuple or shape*

**Assembly format**

.. code-block::

   `(` $input `)` attr-dict `:` functional-type($input, $result)

**Operands**

- ``$input``: input int_tuple or shape

**Results**

- ``$result``: product of all leaf elements (same kind as input)

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Folds all leaf elements of ``$input`` by multiplication. For
example, ``tuple_product((2,(3,4))) = 24``.

**Preconditions**

- ``$input`` is ``!cute.int_tuple`` or ``!cute.shape``.

**Postconditions**

- Result kind matches input kind (``shape -> shape``,
  ``int_tuple -> int_tuple``); result is a scalar (rank-1) equal
  to the product of all leaves of ``$input``.

- When ``$input`` has any dynamic leaves, the result is dynamic.

**Examples**

.. code-block::

   // Static int_tuple: full numerical product
   %p = cute.tuple_product(%t) : (!cute.int_tuple<"(2,(3,4))">) -> !cute.int_tuple<"24">

   // Dynamic shape with one static leaf: result is dynamic.
   // Result kind matches input kind (shape -> shape).
   %p = cute.tuple_product(%t) : (!cute.shape<"(?,(3,?))">) -> !cute.shape<"?">


**Errors**

- operand is not an int_tuple or shape

  .. code-block::

     %p = cute.tuple_product(%c) : (!cute.coord<"(2,3)">) -> !cute.int_tuple<"6">


- declared result kind does not match input kind

  .. code-block::

     %p = cute.tuple_product(%t)
            : (!cute.shape<"(2,(3,4))">) -> !cute.int_tuple<"24">


- declared result does not match the inferred product

  .. code-block::

     %p = cute.tuple_product(%t)
            : (!cute.int_tuple<"(2,(3,4))">) -> !cute.int_tuple<"5">



Arithmetic Operations
~~~~~~~~~~~~~~~~~~~~~

Per-leaf integer arithmetic and comparisons on tuples.

.. _op-cute.ceil_div:

``cute.ceil_div``
^^^^^^^^^^^^^^^^^

*Computes the ceiling division of an arithmetic tuple by an arithmetic or tile tiler.*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$input``: input (dividend)
- ``$tiler``: tiler (divisor)

**Results**

- ``$result``: ceil-divided result

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Computes ``ceil($input / $tiler)`` element-wise. ``$input`` is
``!cute.int_tuple`` or ``!cute.shape``. ``$tiler`` is ``!cute.int_tuple``,
``!cute.shape``, or ``!cute.tile`` (independently chosen from
``$input``). The result has the same kind as ``$input``. An
``!cute.int_tuple`` or ``!cute.shape`` tiler preserves the input profile;
an ``!cute.tile`` tiler may change it.

For an ``!cute.int_tuple`` or ``!cute.shape`` tiler, the operation has
four cases based on the top-level structure of the operands:

- (tuple, tuple) — ``rank($input) >= rank($tiler)`` is required;
  elements of ``$input`` beyond ``rank($tiler)`` are passed through
  unchanged.

- (tuple, scalar) — sequential consumption: divides the first
  element of ``$input`` by ``$tiler``, reduces ``$tiler`` by that
  amount, and repeats across ``$input``.

- (scalar, tuple) — reduces to ``ceil_div($input, product($tiler))``.
- (scalar, scalar) — standard ``ceil(input / tiler)``.

An ``!cute.tile`` tiler instead uses shape/tile semantics. It preserves
the input kind but may change the result profile.

**Preconditions**

- ``$input`` is ``!cute.int_tuple`` or ``!cute.shape``.
- ``$tiler`` is ``!cute.int_tuple``, ``!cute.shape``, or ``!cute.tile``.
- For an arithmetic (tuple, tuple) tiler:
  ``rank($input) >= rank($tiler)``.
- At every (scalar, scalar) leaf the tiler is non-zero.

**Postconditions**

- ``$result`` has the same kind as ``$input``. It has the same profile
  when ``$tiler`` is ``!cute.int_tuple`` or ``!cute.shape``; a
  ``!cute.tile`` tiler may change the profile.
- When all inputs are static, ``$result`` equals the
  ceiling-division value per the case rules above; dynamic inputs
  propagate as dynamic in the result.

**Examples**

.. code-block::

   // Scalar / scalar.
   %r = cute.ceil_div(%a, %b)
          : (!cute.int_tuple<"7">, !cute.int_tuple<"3">)
         -> !cute.int_tuple<"3">

   // Tuple / tuple — element-wise.
   %r = cute.ceil_div(%a, %b)
          : (!cute.int_tuple<"(8,10)">, !cute.int_tuple<"(3,4)">)
         -> !cute.int_tuple<"(3,3)">

   // Tuple / tuple with extra input element passing through unchanged.
   %r = cute.ceil_div(%a, %b)
          : (!cute.int_tuple<"(8,10,6)">, !cute.int_tuple<"(3,4)">)
         -> !cute.int_tuple<"(3,3,6)">

   // A tile tiler may change the result profile.
   %r = cute.ceil_div(%a, %tile)
          : (!cute.int_tuple<"(8,16)">, !cute.tile<"((4,8)):((1,4))">)
         -> !cute.int_tuple<"4">


**Errors**

- Tiler leaf is zero.
- Arithmetic (tuple, tuple) operands where
  ``rank($tiler) > rank($input)``.
- ``$result`` disagrees with the case-rule value when both sides
  are static.

.. code-block::

   // error: expects valid operands, but cannot compute ceil_div for
   //        '!cute.int_tuple<"(4,8)">' and '!cute.int_tuple<"(2,4,2)">':
   //        rank(tiler)=3 > rank(input)=2
   cute.ceil_div(%a, %b) : (!cute.int_tuple<"(4,8)">, !cute.int_tuple<"(2,4,2)">) -> ...

   // error: expects valid operands, but cannot compute ceil_div for
   //        '!cute.int_tuple<"7">' and '!cute.int_tuple<"0">': tiler is zero
   cute.ceil_div(%a, %b) : (!cute.int_tuple<"7">, !cute.int_tuple<"0">) -> ...



.. _op-cute.elem_less:

``cute.elem_less``
^^^^^^^^^^^^^^^^^^

*Element-wise less-than comparison of two CuTe tuple-like values*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$lhs``: left-hand side
- ``$rhs``: right-hand side

**Results**

- ``$result``: less-than result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns ``true`` (``i1``) if every leaf of ``$lhs`` is strictly less than
the corresponding leaf of ``$rhs``, otherwise returns ``false``. Each
operand may be ``!cute.int_tuple``, ``!cute.shape``, or ``!cute.coord``,
chosen independently.

**Preconditions**

- ``$lhs`` and ``$rhs`` are each ``!cute.int_tuple``, ``!cute.shape``, or
  ``!cute.coord`` (independently chosen).

- ``$lhs`` and ``$rhs`` are congruent — same hierarchical rank/depth
  structure (matching ranks at every level). Rank or nesting
  mismatches at any depth are rejected.

**Postconditions**

- ``$result`` is ``i1``.
- When all leaves are static, ``$result`` is fully resolved as the
  AND of ``lhs[i] < rhs[i]`` over all leaves.

- When any leaf is dynamic, ``$result`` is dynamic (an SSA ``i1``
  value whose runtime evaluates the same comparison).

**Examples**

.. code-block::

   %r = cute.elem_less(%a, %b)
          : (!cute.int_tuple<"(3,4)">, !cute.int_tuple<"(5,6)">)
         -> i1

   %r = cute.elem_less(%a, %b)
          : (!cute.shape<"(3,4)">, !cute.shape<"(5,6)">)
         -> i1

   // Mixed kinds — coord vs shape with matching profile.
   %r = cute.elem_less(%a, %b)
          : (!cute.coord<"(3,4)">, !cute.shape<"(5,6)">)
         -> i1


**Errors**

- ``$lhs`` and ``$rhs`` are not congruent (top-level rank differs, or
  nested structure differs at any depth).

.. code-block::

   // error: 'cute.elem_less' op expects congruent inputs, but got
   //        '!cute.int_tuple<"(3,4)">' and '!cute.int_tuple<"(3,4,5)">'
   cute.elem_less(%a, %b) : (!cute.int_tuple<"(3,4)">, !cute.int_tuple<"(3,4,5)">) -> i1

   // error: 'cute.elem_less' op expects congruent inputs, but got
   //        '!cute.int_tuple<"((1,2),3)">' and '!cute.int_tuple<"(4,5)">'
   cute.elem_less(%a, %b) : (!cute.int_tuple<"((1,2),3)">, !cute.int_tuple<"(4,5)">) -> i1



.. _op-cute.equal:

``cute.equal``
^^^^^^^^^^^^^^

*Equality comparison of two CuTe values*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$lhs``: left-hand side
- ``$rhs``: right-hand side

**Results**

- ``$result``: equality result

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Returns ``true`` (``i1``) if ``$lhs`` and ``$rhs`` are structurally and
value-equal, otherwise returns ``false``. Each operand may be
``!cute.int_tuple``, ``!cute.shape``, ``!cute.coord``, ``!cute.stride``,
or ``!cute.layout``, but ``$lhs`` and ``$rhs`` must be the same kind.

**Preconditions**

- ``$lhs`` and ``$rhs`` are the same CuTe type kind (e.g. both
  ``!cute.shape``, both ``!cute.layout``, etc.). Cross-kind comparison
  is rejected even when the underlying values look similar
  (e.g. ``int_tuple`` vs ``shape``).

**Postconditions**

- ``$result`` is ``i1``.
- When both operands are fully static, ``$result`` is fully
  resolved by structural and value equality.

- When either operand has dynamic components, ``$result`` is
  dynamic (an SSA ``i1`` evaluated at runtime).

**Examples**

.. code-block::

   // Equal int_tuple — true.
   %r = cute.equal(%a, %b)
          : (!cute.int_tuple<"(3,4)">, !cute.int_tuple<"(3,4)">)
         -> i1

   // Different shape values — false.
   %r = cute.equal(%a, %b)
          : (!cute.shape<"(3,4)">, !cute.shape<"(3,5)">)
         -> i1

   // Equal layout — true.
   %r = cute.equal(%a, %b)
          : (!cute.layout<"(4,8):(1,4)">, !cute.layout<"(4,8):(1,4)">)
         -> i1


**Errors**

- ``$lhs`` and ``$rhs`` have different kinds (cross-kind comparison
  is rejected, even with matching profiles).

.. code-block::

   // error: 'cute.equal' op expects lhs and rhs to be the same kind,
   //        but got '!cute.int_tuple<"(3,4)">' and '!cute.shape<"(3,4)">'
   cute.equal(%a, %b) : (!cute.int_tuple<"(3,4)">, !cute.shape<"(3,4)">) -> i1



.. _op-cute.shape_div:

``cute.shape_div``
^^^^^^^^^^^^^^^^^^

*CuTe shape division*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$a``: dividend shape
- ``$b``: divisor shape

**Results**

- ``$result``: shape quotient

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Divides shape ``$a`` by shape ``$b`` using CuTe's shape-division
rules. Both operands and the result are ``!cute.shape``. The
operation has four cases based on the top-level structure of the
operands:

- (tuple, tuple) — element-wise ``shape_div`` on corresponding
  elements. The tuple ranks must match; Traverses each element then selects
  independently among these four cases.

- (tuple, scalar) — sequential consumption: divides the first
  element of ``$a`` by ``$b``, reduces ``$b`` by that amount, and
  repeats across ``$a``. Example: ``shape_div((4,5,6), 40) -> (1,1,3)``.

- (scalar, tuple) — reduces to ``shape_div($a, product($b))``.
- (scalar, scalar) — requires ``a % b == 0`` or ``b % a == 0``;
  returns ``a / b`` rounded away from zero (result is 1 or -1 when
  ``|a| < |b|``).

**Preconditions**

- ``$a`` and ``$b`` are ``!cute.shape``.
- At each recursive (tuple, tuple) case, the tuple ranks must match.
- For (scalar, scalar) leaves the divisor must be non-zero and one
  side must divide the other (``a % b == 0`` or ``b % a == 0``).

- For (tuple, scalar) the running divisor must be fully consumed
  or remain divisible at each step.

**Postconditions**

- ``$result`` is an ``!cute.shape`` whose profile follows the case
  rules above.

- When all inputs are static, ``$result`` equals the CuTe
  shape-division value; dynamic inputs propagate as dynamic in
  the result.

**Examples**

.. code-block::

   // Scalar / scalar.
   %r = cute.shape_div(%a, %b)
          : (!cute.shape<"12">, !cute.shape<"4">)
         -> !cute.shape<"3">

   // Tuple / tuple — element-wise.
   %r = cute.shape_div(%a, %b)
          : (!cute.shape<"(8,10)">, !cute.shape<"(2,5)">)
         -> !cute.shape<"(4,2)">

   // Tuple / scalar — sequential consumption.
   %r = cute.shape_div(%a, %b)
          : (!cute.shape<"(4,5,6)">, !cute.shape<"40">)
         -> !cute.shape<"(1,1,3)">


**Errors**

- Unequal ranks at any recursive (tuple, tuple) case.
- Divisor leaf is zero.
- (scalar, scalar) leaves where neither side divides the other.
- ``$result`` shape disagrees with the case-rule value when both
  sides are static.

.. code-block::

   // error: expects valid operands, but cannot compute shape_div for
   //        '!cute.shape<"4">' and '!cute.shape<"0">': divisor is zero
   cute.shape_div(%a, %b) : (!cute.shape<"4">, !cute.shape<"0">) -> ...

   // error: expects valid operands, but cannot compute shape_div for
   //        '!cute.shape<"7">' and '!cute.shape<"3">': 7 is not divisible
   //        by 3 (neither 7 % 3 == 0 nor 3 % 7 == 0)
   cute.shape_div(%a, %b) : (!cute.shape<"7">, !cute.shape<"3">) -> ...



.. _op-cute.tuple_add:

``cute.tuple_add``
^^^^^^^^^^^^^^^^^^

*Element-wise addition of two arithmetic tuples*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$lhs``: left-hand side
- ``$rhs``: right-hand side

**Results**

- ``$result``: element-wise sum

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Adds ``$lhs`` and ``$rhs`` element-wise, recursing into nested
sub-tuples. The two operands must be the same kind
(``!cute.int_tuple`` or ``!cute.shape``); the result has the same kind.

When the two operands have different lengths, elements beyond the
common rank are copied through unchanged from whichever side
provides them — as if the shorter operand were zero-padded. The
same rule applies recursively to nested sub-tuples:

.. code-block::

   (1,2,3) + (4,5)   =  (5,7,3)        // lhs extra element passes through
   (1,2)   + (4,5,6) =  (5,7,6)        // rhs extra element passes through
   ((1,2),(3,4),(5,6)) + ((10,11),(20,21))
                      =  ((11,13),(23,25),(5,6))   // extra nested sub-tuple
   ((1,2),(3,4)) + ((10,11),(20,21,22))
                      =  ((11,13),(23,25,22))       // extra inner element


**Preconditions**

- ``$lhs`` and ``$rhs`` are both ``!cute.int_tuple`` or both ``!cute.shape``
  (matching kinds).

- Scalar-tuple identity — A zero scalar paired with a tuple is
  the additive identity (``0 + t == t + 0 == t``). A non-zero scalar
  mixed with a non-scalar tuple is invalid.

- Nesting constraint — Within the common rank, both operands
  must mirror each other structurally: at every position
  ``i < min(rank(lhs), rank(rhs))``, either both elements are leaves
  or both are sub-tuples. A leaf-vs-sub-tuple mismatch at any depth
  is invalid.

**Postconditions**

- ``$result`` has the same kind as ``$lhs`` / ``$rhs``.
- ``rank($result) == max(rank($lhs), rank($rhs))``.
- For positions in the common rank, ``$result[i]`` is the recursive
  element-wise sum of ``$lhs[i]`` and ``$rhs[i]``. Trailing positions
  pass through unchanged from the longer operand.

**Examples**

.. code-block::

   // Static int_tuple.
   %r = cute.tuple_add(%a, %b)
          : (!cute.int_tuple<"(3,4)">, !cute.int_tuple<"(1,2)">)
         -> !cute.int_tuple<"(4,6)">

   // Static shape.
   %r = cute.tuple_add(%a, %b)
          : (!cute.shape<"(6,8)">, !cute.shape<"(2,4)">)
         -> !cute.shape<"(8,12)">

   // Rank mismatch — extra lhs element passes through unchanged.
   %r = cute.tuple_add(%a, %b)
          : (!cute.int_tuple<"(1,2,3)">, !cute.int_tuple<"(4,5)">)
         -> !cute.int_tuple<"(5,7,3)">

   // Dynamic lhs + static rhs.
   %r = cute.tuple_add(%a, %b)
          : (!cute.int_tuple<"(?,?)">, !cute.int_tuple<"(1,2)">)
         -> !cute.int_tuple<"(?,?)">


**Errors**

- ``$lhs`` and ``$rhs`` have different kinds (e.g. int_tuple vs shape).
- A non-zero scalar is paired with a non-scalar tuple.
- Within the common rank, a leaf is paired with a sub-tuple at any
  depth.

- ``$result`` kind differs from ``$lhs``/``$rhs``, or its rank/values
  disagree with the element-wise sum when both sides are static.


.. _op-cute.tuple_sub:

``cute.tuple_sub``
^^^^^^^^^^^^^^^^^^

*Element-wise subtraction of two arithmetic tuples*

**Assembly format**

.. code-block::

   `(` operands `)` attr-dict `:` functional-type(operands, results)

**Operands**

- ``$lhs``: left-hand side
- ``$rhs``: right-hand side

**Results**

- ``$result``: element-wise difference

**Traits**: ``AlwaysSpeculatableImplTrait``

**Interfaces**: ``ConditionallySpeculatable``, ``InferTypeOpInterface``

Description
"""""""""""


Subtracts ``$rhs`` from ``$lhs`` element-wise (``lhs - rhs``), recursing
into nested sub-tuples. The two operands must be the same kind
(``!cute.int_tuple`` or ``!cute.shape``); the result has the same kind.

When the two operands have different lengths, elements beyond the
common rank are copied through unchanged from whichever side
provides them — as if the shorter operand were zero-padded. The
same rule applies recursively to nested sub-tuples:

.. code-block::

   (5,8,3) - (1,2)     =  (4,6,3)      // extra lhs element passes through
   (10,8)  - (3,2,5)   =  (7,6,5)      // extra rhs element passes through
   ((5,8),(3,4),(1,2)) - ((1,2),(1,1))
                        =  ((4,6),(2,3),(1,2))   // extra nested sub-tuple
   ((5,8),(3,4)) - ((1,2),(1,1,1))
                        =  ((4,6),(2,3,1))        // extra inner element


**Preconditions**

- ``$lhs`` and ``$rhs`` are both ``!cute.int_tuple`` or both ``!cute.shape``
  (matching kinds).

- Scalar-tuple identity — A zero scalar ``$rhs`` paired with a
  non-scalar tuple ``$lhs`` is the additive identity (``t - 0 == t``).
  A zero scalar ``$lhs`` paired with a non-scalar tuple ``$rhs`` is
  invalid (subtraction is not commutative). Any non-zero scalar
  mixed with a non-scalar tuple at the top level is also invalid.

- Nesting constraint — Within the common rank, both operands
  must mirror each other structurally: at every position
  ``i < min(rank(lhs), rank(rhs))``, either both elements are leaves
  or both are sub-tuples. A leaf-vs-sub-tuple mismatch at any depth
  is invalid.

**Postconditions**

- ``$result`` has the same kind as ``$lhs`` / ``$rhs``.
- ``rank($result) == max(rank($lhs), rank($rhs))``.
- For positions in the common rank, ``$result[i]`` is the recursive
  element-wise difference ``$lhs[i] - $rhs[i]``. Trailing positions
  pass through unchanged from the longer operand.

**Examples**

.. code-block::

   // Static int_tuple.
   %r = cute.tuple_sub(%a, %b)
          : (!cute.int_tuple<"(5,8)">, !cute.int_tuple<"(1,2)">)
         -> !cute.int_tuple<"(4,6)">

   // Rank mismatch — extra lhs element passes through unchanged.
   %r = cute.tuple_sub(%a, %b)
          : (!cute.int_tuple<"(5,8,3)">, !cute.int_tuple<"(1,2)">)
         -> !cute.int_tuple<"(4,6,3)">

   // Rank mismatch — extra rhs element passes through unchanged.
   %r = cute.tuple_sub(%a, %b)
          : (!cute.int_tuple<"(10,8)">, !cute.int_tuple<"(3,2,5)">)
         -> !cute.int_tuple<"(7,6,5)">

   // Dynamic lhs + static rhs.
   %r = cute.tuple_sub(%a, %b)
          : (!cute.int_tuple<"(?,?)">, !cute.int_tuple<"(1,2)">)
         -> !cute.int_tuple<"(?,?)">


**Errors**

- ``$lhs`` and ``$rhs`` have different kinds (e.g. shape vs int_tuple).
- A zero scalar ``$lhs`` is paired with a non-scalar ``$rhs``, or any
  non-zero scalar is paired with a non-scalar tuple at the top
  level.

- Within the common rank, a leaf is paired with a sub-tuple at any
  depth.

- ``$result`` kind differs from ``$lhs``/``$rhs``, or its rank/values
  disagree with the element-wise difference when both sides are
  static.


Tiling and Partitioning Products Operations

Ops that print or trace runtime values for debugging.

.. _op-cute.print:

``cute.print``
^^^^^^^^^^^^^^

*Print a CuTe value at runtime for debugging*

**Assembly format**

.. code-block::

   `(` $value `)` `:` type($value) attr-dict

**Operands**

- ``$value``: CuTe value to print

**Effects**

- ``::mlir::MemoryEffects::Write`` on ``::mlir::SideEffects::DefaultResource``

Description
"""""""""""


Prints the runtime value of a CuTe operand to stdout for debugging.
Has side effects. Produces no result.

**Preconditions**

- ``$value`` is any CuTe type (static or dynamic, scalar or tuple).

**Postconditions**

- None — this op produces no result.

**Examples**

.. code-block::

   cute.print(%x) : !cute.layout<"(2,3):(1,2)">
   cute.print(%y) : !cute.stride<"?">



Passes
------

.. _pass-cute-expand-ops:

``--cute-expand-ops``
~~~~~~~~~~~~~~~~~~~~~

*Expand selected cute ops via dialect conversion patterns*


Partial dialect conversion that rewrites high-level cute algebra ops into
sequences of constructor (``make_*``) ops + ``arith``/``scf`` ops. The base
constructor ops, ``cute.static``, ``cute.print``, and ``cute.get_scalars`` are
left legal and are lowered later by ``cute-to-base``.

Each illegal op rewrites to one of two shapes:

.. code-block::

   // Static result -> fold to cute.static
   %res = cute.algebra_op(%input) : ... -> !type_with_static_modes
   return %res : !type_with_static_modes

   // Becomes:
   %res = cute.static : !type_with_static_modes
   return %res : !type_with_static_modes

   // Dynamic result -> extract dynamic leaves and rebuild via constructors
   %res = cute.algebra_op(%input) : ... -> !type_with_dynamic_modes
   return %res : !type_with_dynamic_modes

   // Becomes:
   %scalars = cute.get_scalars(%input) <{only_dynamic}> : ...
   %t0 = arith.<...>(%scalars#0, %scalars#1) : ...
   ...
   %res = cute.make_<...>(%t0, ...) : !type_with_dynamic_modes
   return %res : !type_with_dynamic_modes



.. _pass-cute-fold-static:

``--cute-fold-static``

*Lower remaining cute ops to base dialects (LLVM/arith/scf/ub)*


Full dialect conversion from ``cute`` to LLVM/arith/scf/ub. Applied after
``cute-expand-ops`` has reduced the IR to the 10 surviving constructor /
static / print / get_scalars ops, this pass lowers each of
those ops along with all cute types.

Type conversion: every cute type maps to an LLVM struct holding the
dynamic leaves only (cutegen's ``llvm_type_sparse_flat_t``):

- ``!cute.shape``/``stride``/``int_tuple``/``coord``/``tile`` → ``LLVM struct``
  of dynamic-leaf integers.

- ``!cute.layout`` → ``LLVM struct(shape_struct, stride_struct)``.
- ``!cute.composed_layout`` →
  ``LLVM struct(inner_struct, offset_struct, outer_struct)``.

- ``!cute.swizzle`` → ``LLVM struct()`` (empty).

Lowers ``cute.static``, ``cute.get_scalars``, the 7 constructor ops
(``make_*``), and ``cute.print`` — ``cute.print`` emits ``gpu.printf``
inside a GPU module/launch and falls back to ``llvm.call @printf``
(with a generated ``llvm.mlir.global`` format string) on the host.