Skip to content

[REQ] Add support for tile slicing syntax, including fancy indexing #1176

Description

@daedalus5

Feature Request: NumPy-Style Slicing for Warp Tiles

Problem

Warp tiles currently lack slicing syntax support. While tile_view() exists, it requires explicit offset and shape parameters rather than the familiar NumPy slicing notation that users expect.

Current approach with tile_view():

@wp.kernel
def process_tile(data: wp.array2d(dtype=float)):
    t = wp.tile_load(data, shape=(64, 64))
    
    # Extract rows 10-20
    subset = wp.tile_view(t, offset=(10, 0), shape=(10, 64))
    
    # Extract a single row — requires knowing the full shape
    row = wp.tile_view(t, offset=(5, 0), shape=(1, 64))

What we want:

@wp.kernel
def process_tile(data: wp.array2d(dtype=float)):
    t = wp.tile_load(data, shape=(64, 64))
    
    # NumPy-style slicing
    subset = t[10:20, :]
    row = t[5, :]
    cols = t[:, 0:32]
    strided = t[::2, ::2]

    # Slice assignment
    t[0:8, :] = wp.tile_ones(shape=(8, 64), dtype=float)

Additionally, there's no support for index array slicing (fancy indexing), which is essential for gathering non-contiguous elements:

@wp.kernel
def gather_elements(data: wp.array2d(dtype=float)):
    t = wp.tile_load(data, shape=(64, 64))
    
    # Gather specific rows by index
    indices = wp.tile_arange(0, 8, dtype=int) * 8  # [0, 8, 16, 24, 32, 40, 48, 56]
    selected_rows = t[indices, :]  # Not currently possible

Implementation Notes

Relationship to Existing APIs

  • tile_view(t, offset, shape) — Should remain as the explicit API; basic slice syntax (e.g., t[0:8, :]) compiles down to tile_view calls
  • tile_slice_indexed (new primitive) — Index array slicing (e.g., t[indices, :]) would require a new gather operation
  • tile_extract(t, i, j, ...) — Unchanged; extracts scalar elements
  • t[i] — Existing integer indexing unchanged; reduces dimensionality

Compile-Time vs Runtime

Slice bounds should be compile-time constants for optimal code generation, consistent with existing tile shape requirements. The compiler can verify bounds at compile time when constants are used.

Activity

  1. added this to the 1.12.0 milestone on Jan 14, 2026
  2. self-assigned this
    on Jan 14, 2026
  3. modified the milestones: 1.12.0, 1.13.0 on Jan 29, 2026
  4. modified the milestones: 1.13.0, 1.14.0 on Apr 7, 2026
  5. modified the milestones: 1.14.0, 1.15.0 on May 5, 2026
  6. modified the milestones: 1.15.0, 1.16.0 on May 22, 2026
  7. 4 remaining items

  8. shi-eric commented on Jul 23, 2026

    @shi-eric
    Contributor

    This is addressed by c07653a.

    The change adds NumPy-style tile indexing in kernel code. Basic, strided, and reversed slices and dimension-collapsing integer indices now lower to tile views; slice assignment is supported; and 1D integer index tiles gather through the new differentiable wp.tile_slice_indexed() primitive. Existing explicit wp.tile_view() and scalar wp.tile_extract() behavior remains available.

    Public API behavior:

    Kernel scope:

    @wp.kernel
    def process_tile(data: wp.array2d[float]):
        t = wp.tile_load(data, shape=(64, 64))
    
        block = t[2:6, :]           # Select rows 2 through 5
        columns = t[:, ::2]         # Select every other column
        reversed_rows = t[::-1, :]  # Reverse the row order
        row = t[5, :]               # Select one row and collapse that dimension
    
        # Gather non-contiguous rows using a 1D integer index tile
        indices = wp.tile_arange(0, 8, dtype=int) * 8
        selected_rows = t[indices, :]
    
        # Assign a tile to a view in place
        t[0:8, :] = wp.tile_ones(shape=(8, 64), dtype=float)

    Slice bounds must be compile-time constants. Runtime integer indices are supported, including negative indices, but out-of-range runtime values are checked only in debug builds.

    Advanced-index assignment, compound assignment to views, scalar broadcasting into slices, and zero-length tile views remain unsupported.

    Test coverage:

    • warp/tests/tile/test_tile_view.py adds 39 tests covering basic, strided, reversed, dimension-collapsing, and assigned slices, including test_tile_slice_rows, test_tile_slice_strided, test_tile_slice_reverse, and test_tile_slice_assign.
    • Advanced indexing and autograd coverage includes test_tile_advanced_index, test_tile_advanced_index_duplicate_grad, test_tile_slice_assign_grad, and test_tile_gather_neg_index.
    • Runtime and negative indices, chained expressions, explicit slice offsets, overlapping assignment, bounds validation, and unsupported inputs are also covered.

    An independent probe matching the original 64×64 examples passed on CPU and CUDA. No remaining requirement from this issue was identified.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

Labels

feature requestRequest for something to be added

Projects

No projects

    Milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions