Skip to main content

Mojo function

expect_bytes_pred

expect_bytes_pred(mbar_ptr: UnsafePointer[address_space=AddressSpace.SHARED], bytes: Int32, pred: Int32)

Issue mbarrier.arrive.expect_tx.shared::cta.b64 predicated on pred != 0.

Equivalent to:

if pred != 0:
    mbar_ptr[].expect_bytes(bytes)

but folds the runtime branch into a single PTX @%p predicate on the mbarrier.arrive.expect_tx instruction β€” no Mojo-level if, no SASS branch, no warp divergence.

Args:

  • ​mbar_ptr (UnsafePointer[address_space=AddressSpace.SHARED]): Pointer to the shared-memory mbarrier (8-byte slot).
  • ​bytes (Int32): Expected transaction byte count for this barrier.
  • ​pred (Int32): Runtime predicate (typically the result of elect() or an elect()-derived elect_mask); the PTX instruction is skipped when this is 0.