warp.tile_stack_pop

warp.tile_stack_pop(s: Any) → tuple[Any, int]

Kernel

Pop a value from a tile stack (cooperative).

All threads in the block must call this function. Each calling thread races for a slot.

Parameters:

s – The tile stack

Returns:

A tuple (value, slot) where value is the popped element (or the default value if the stack was empty) and slot is the index of the popped element (the slot it previously occupied), or -1 if the stack was empty. When non-negative, slot lies in [0, capacity-1]. Consistent with tile_stack_push() which also uses -1 to indicate failure.

Example

CAP = wp.constant(8)

@wp.kernel
def pop_kernel(out: wp.array[int]):
    _i, j = wp.tid()
    s = wp.tile_stack(capacity=CAP, dtype=int)
    wp.tile_stack_push(s, j * 10, j < 4)

    val, slot = wp.tile_stack_pop(s)
    if slot != -1:
        out[slot] = val

out = wp.full(8, -1, dtype=int)
wp.launch_tiled(pop_kernel, dim=[1], inputs=[out], block_dim=8)

vals = out.numpy()
print(sorted(vals[vals >= 0].tolist()))

[0, 10, 20, 30]