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

Mojo function

st_shared_v4_b32

def st_shared_v4_b32[dtype: DType, //](dst: Pointer[Scalar[dtype], address_space=AddressSpace.SHARED, _safe=False], elem_off: Int, packed: SIMD[DType.uint32, SIMDLength(4)])

Explicit 16 B st.shared.v4.b32 (one STS.128) to dst[elem_off].

Forces the wide vector store regardless of how packed's four words were produced. A plain .store() of a SIMD[uint32, 4] scalarizes to 4x STS.32 when the words come from a long-lived accumulator (the split-K combine's o_final): ptxas cannot fuse the non-contiguous in-place F2FP pack outputs, and the resulting 4 B stores hit only every 4th bank (4-way conflict, 4x wavefronts). The v4.b32 operand mandates a contiguous register quad, so this stays one bank-conflict-free 16 B transaction.

dtype is the shared buffer's element type -- any 1-byte (fp8), 2-byte (bf16/f16), or 4-byte (f32) output; elem_off is in dtype elements. packed is a fixed 16 B (four u32) for every dtype: pack_row folds the per-u32 element count (2 for bf16, 4 for fp8, 1 for f32) into that width, so one call stores one SWIZZLE_NONE block (fa4_ws_intracta_combine stores its f32 output this way).

Parameters:

dtype (DType): Element dtype of the shared buffer; any 1-byte (fp8), 2-byte (bf16/f16), or 4-byte (f32) output dtype (inferred).

Args:

dst (Pointer[Scalar[dtype], address_space=AddressSpace.SHARED, _safe=False]): Shared-memory pointer to the store target.
elem_off (Int): Element offset into dst in dtype-element units.
packed (SIMD[DType.uint32, SIMDLength(4)]): Four u32 words forming one 16 B SWIZZLE_NONE block to store.