triton.language.atomic_poll

triton.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 the specified pointer until either the value equals the expected value, or the operation times out. In the event of a timeout, the operation returns false and no results may be acquired.

Parameters:

  • pointer (triton.PointerDType) – A pointer to a scalar 16-, 32-, or 64-bit integer.

  • expected_value (pointer.dtype.element_ty) – The value that ends the polling loop.

  • 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) – Maximum wall time to poll, measured in nanoseconds by the GPU global timer. If omitted, polling has no timeout. A timeout of zero still performs one load.

Returns:

True if the expected value was observed, or False if the 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, ...).