Mojo function

tcgen05_ld

tcgen05_ld[*, datapaths: Int, bits: Int, repeat: Int, dtype: DType, pack: Bool, width: Int = (div_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "_mlir_value">), 1024) + -1) if (((bits * datapaths * repeat) < 0) & ((rem_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "_mlir_value">), 1024) == 0) ^ True)) else div_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "_mlir_value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "_mlir_value">), 1024)](tmem_addr: UInt32) -> SIMD[dtype, width]

Loads data from tensor memory into registers.

Parameters:

datapaths (Int): The first dimension of the shape.
bits (Int): The second dimension of the shape.
repeat (Int): The repeat factor.
dtype (DType): The data type to load.
pack (Bool): Whether to pack two 16-bit chunks of adjacent columns into a single 32-bit register.
width (Int): The number elements in the result vector.

Args:

tmem_addr (UInt32): The address of the tensor memory to load from.

Returns:

SIMD: The SIMD register containing the loaded data.