Skip to main content
Log in

Mojo function

tcgen05_ld

tcgen05_ld[*, datapaths: Int, bits: Int, repeat: Int, type: DType, pack: Bool, width: Int = (div_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "value">), 1024) + -1) if (((bits * datapaths * repeat) < 0) & ((rem_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "value">), 1024) == 0) ^ True)) else div_s(mul(#lit.struct.extract<:@stdlib::@builtin::@int::@Int bits, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int datapaths, "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int repeat, "value">), 1024)](tmem_addr: SIMD[uint32, 1]) -> SIMD[type, 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.
  • type (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 nubmer elements in the result vector.

Args:

  • tmem_addr (SIMD[uint32, 1]): The address of the tensor memory to load from.

Returns:

The SIMD register containing the loaded data.

Was this page helpful?