For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /get-started.md).

Mojo function

store_global_pred

def store_global_pred[dtype: DType, address_space: AddressSpace, //](ptr: Pointer[Scalar[dtype], address_space=address_space], value: Scalar[dtype], pred: Int32)

Issue a global store predicated on pred != 0.

Equivalent to:

if pred != 0:
    ptr[] = value

but folds the guard into a PTX @%p predicate on the store, so ptxas emits a single predicated STG instead of the BSSY/BRA/NOP/BSYNC reconvergence quartet an if costs.

That is a real trade, not a free win. The if form gives ptxas a basic block boundary, which caps the scheduling region; this form does not, so ptxas may hoist the producers of value across neighbouring code and spill. Measure spill counts, not just the instruction count, when switching a site to this helper.

There is no ~{memory} clobber, matching expect_bytes_pred and the st.shared.v4.b32 helper above: adding one would force LLVM to reload every value it has cached from memory at each call site. Callers must therefore not read back what they store here without an explicit fence.

Parameters:

dtype (DType): Element dtype of the store; must be 4 or 8 bytes (inferred).
address_space (AddressSpace): Address space of ptr (inferred).

Args:

ptr (Pointer[Scalar[dtype], address_space=address_space]): Destination address. Must point into global memory -- the instruction is st.global, so a generic pointer into shared or local memory is undefined behaviour rather than a compile error.
value (Scalar[dtype]): The value to store when pred is nonzero.
pred (Int32): Runtime predicate; the store is skipped when this is 0.