warp.tile#

warp.tile(
x: Any,
preserve_type: bool = False,
) → Tile[Any, tuple]#
  • Kernel
  • Differentiable

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 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 CPU Tile Semantics.

Parameters:

  • x – A per-thread local value, e.g. scalar, vector, or matrix.

  • 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.

Example

This example shows how to create a linear sequence from thread variables.

@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())

[ 0  2  4  6  8 10 12 14 16 18 20 22 24 26 28 30]