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

pack_row

def pack_row[n: Int, //, output_type: DType, w: Int, start: Int = Int(0)](o_vals: Array[Float32, n]) -> SIMD[DType.uint32, SIMDLength(4)]

Cast the w f32 O lanes o_vals[start : start + w] to output_type and pack them into one 16 B SWIZZLE_NONE store register (exactly four u32).

per_u32 = 4 // size_of[output_type]() output elements pack into each u32 (2 for a 2-byte bf16/f16 output, 4 for a 1-byte fp8 output), so a full 16 B block is w == 4 * per_u32 f32 lanes (8 for bf16/f16, 16 for fp8). The return width is a fixed 4 so the wide-store helper st_shared_v4_b32 takes it without a symbolic-width unification. Scale-free sibling of scale_pack_o_row, used by the split-K combine where o_final is already normalized.

o_vals is a tcgen05_ld / accumulator result; start/w window it. Each u32 is built from an SIMD[f32, per_u32] chunk (f32x2 for bf16 -- wider SIMD scalarizes; f32x4 for fp8, mirroring the MLA fp8 store path); only the packed u32 store register is built wide.

Parameters:

n (Int): Total number of f32 O lanes in o_vals (inferred).
output_type (DType): Target dtype to cast the f32 lanes to; must be a 1-byte (fp8) or 2-byte (bf16/f16) dtype.
w (Int): Number of f32 lanes to pack; must equal 4 * per_u32 (8 for bf16/f16, 16 for fp8) to fill one 16 B block.
start (Int): Starting index into o_vals for the window (defaults to 0).

Args:

o_vals (Array[Float32, n]): tcgen05_ld or accumulator result holding the f32 O lanes; the window [start, start + w) is packed.

Returns:

SIMD[DType.uint32, SIMDLength(4)]