triton.experimental.gluon.language.atomic_poll

triton.experimental.gluon.language.atomic_poll(pointer, expected_value, sem=None, scope=None, timeout_ns=None, _semantic=None)

Wait until the value at pointer equals expected_value.

This will spin-wait on each specified pointer until either its value equals the expected value, or the operation times out. The block waits for all polls to finish. Timed-out elements return false and acquire no results.

Parameters:

  • pointer (triton.PointerDType) – A pointer, or block of pointers, to 16-, 32-, or 64-bit integers.

  • expected_value (pointer.dtype.element_ty) – The value that ends each polling loop, broadcast to the shape of pointer.

  • sem (str, optional) – Specifies whether a successful poll has acquire semantics. Acceptable values are “acquire” (default) and “relaxed”.

  • scope (str, optional) – Defines the scope of threads that observe the synchronizing effect of the poll. Acceptable values are “gpu” (default), “cta” (cooperative thread array, thread block), and “sys” (system).

  • timeout_ns (int, optional) – Shared polling time budget for the entire operation, measured in nanoseconds by the GPU global timer. If omitted, polling has no timeout. Each element is loaded at least once, even with a zero timeout.

Returns:

A boolean with the shape of pointer, true for each element whose expected value was observed and false if its timeout expired first.

Return type:

triton.language.tensor

This function can also be called as a member function on tensor, as x.atomic_poll(...) instead of atomic_poll(x, ...).