warp.tile
=========

.. currentmodule:: warp._src.lang

.. function:: tile(x: Any, preserve_type: bool = False) -> Tile[Any, tuple]

   .. wp-function-tags::
      :kernel: true
      :python: false
      :differentiable: true

   Construct a new tile from per-thread kernel values.
   
   This function converts values computed using scalar kernel code to a tile
   representation for input into collective operations. Each thread of the block
   contributes one value, so the tile's trailing dimension is always ``block_dim``:
   
   * If the input value is a scalar, then the resulting tile has ``shape=(block_dim,)``
   * If the input value is a vector, then the resulting tile has
     ``shape=(length(vector), block_dim)``
   * If the input value is a vector, and ``preserve_type=True``, then the resulting
     tile has ``dtype=vector`` and ``shape=(block_dim,)``
   * If the input value is a matrix, then the resulting tile has
     ``shape=(rows, cols, block_dim)``
   * If the input value is a matrix, and ``preserve_type=True``, then the resulting
     tile has ``dtype=matrix`` and ``shape=(block_dim,)``
   
   Quaternion values are supported with ``preserve_type=True``. Use
   :func:`~warp.untile` to convert the tile back to per-thread values.
   
   Every thread of the block must reach this call. On CPU the effective block width is
   ``1``, so the tile has a single element regardless of the requested ``block_dim`` -
   see :ref:`CPU Tile Semantics <cpu_tile_semantics>`.
   
   :param x: A per-thread local value, e.g. scalar, vector, or matrix.
   :param preserve_type: If true, the tile will have the same data type as the input value.
                         Must be a compile-time constant.
   
   :returns: If ``preserve_type=True``, a tile of type ``x.type`` of length ``block_dim``.
             Otherwise, an N-dimensional tile such that the first N-1 dimensions match the
             shape of ``x`` and the final dimension is of size ``block_dim``.
   
   .. rubric:: Example
   
   This example shows how to create a linear sequence from thread variables.
   
   .. testcode::
       :skipif: wp.get_device() == "cpu" or wp.get_cuda_device_count() == 0
   
       @wp.kernel
       def store_doubled_thread_indices(out: wp.array[int]):
           i = wp.tid()
           t = wp.tile(i * 2)
           wp.tile_store(out, t)
   
       out = wp.zeros(16, dtype=int)
   
       wp.launch(store_doubled_thread_indices, dim=16, outputs=[out], block_dim=16)
   
       print(out.numpy())
   
   .. testoutput::
   
       [ 0  2  4  6  8 10 12 14 16 18 20 22 24 26 28 30]