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 ofelect()or anelect()-derivedelect_mask); the PTX instruction is skipped when this is 0.
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!