tilize

template<bool is_fp32_dest_acc_en = DST_ACCUM_MODE>
void ckernel::tilize_init(uint32_t icb, uint32_t block, uint32_t ocb, uint32_t call_line = __builtin_LINE())

Initializes the tilize operation. Should be called once at the beginning of a kernel.

Return value: None

Param Type Name

Function icb

Function block

Function ocb

template<bool neginf_srcA = true, bool zero_srcA_reduce = false, bool is_fp32_dest_acc_en = DST_ACCUM_MODE>
void ckernel::tilizeA_B_reduce_init(uint32_t icb0, uint32_t icb1_scaler, uint32_t block, uint32_t call_line = __builtin_LINE())

Short initializes the tilize operation with reduction.

Param Type Name

Template neginf_srcA

Template zero_srcA_reduce

Function icb0

Function icb1_scaler

Function block

Unpack face geometry for operand A comes from circular-buffer metadata (JIT unpack_tile_* arrays), e.g. set_unpack_face_geometry / set_tile_dims on the host.

template<bool is_fp32_dest_acc_en = DST_ACCUM_MODE>
void ckernel::tilize_init_short_with_dt(uint32_t old_icb, uint32_t new_icb, uint32_t block, uint32_t ocb)

Re-initializes the tilize operation and reconfigures the unpacker with CB data type.

Return value: None

Param Type Name

Function old_icb

Function new_icb

Function block

Function ocb

template<bool is_fp32_dest_acc_en = DST_ACCUM_MODE>
void ckernel::tilize_block(uint32_t icb, uint32_t block, uint32_t ocb, uint32_t input_tile_index = 0, uint32_t output_tile_index = 0)

Performs the tilize operation on a block.

Return value: None

Param Type Name

Function icb

Function block

Function ocb

Function input_tile_index

Function output_tile_index

template<bool neginf_srcA = true, std::uint32_t reload_srcB = true, bool zero_srcA = false, bool zero_srcA_reduce = false>
void ckernel::unpack_tilizeA_B_block(uint32_t icb0, uint32_t icb1, uint32_t block, uint32_t tile_idx_b)

Unpacks and tilizes a block from two input CBs.

Return value: None

Param Type Name

Template neginf_srcA

Template reload_srcB

Template zero_srcA

Template zero_srcA_reduce

Function icb0

Function icb1

Function block

Function tile_idx_b

Operand A face geometry is read from circular-buffer unpack metadata.

void ckernel::tilize_uninit(uint32_t icb, uint32_t ocb)

Uninitializes the tilize operation before re-initializing for another operation.

NOTE: This function is not in line with our programming model, and will be removed by the end of 2025 as a part of tt-metal#22904. NOTE: Does nothing on Quasar because there is no persistent tilize unpack/pack state to undo.

Return value: None

Param Type Name

Function icb

Function ocb

template<bool is_fp32_dest_acc_en = DST_ACCUM_MODE>
void ckernel::tilize_uninit_with_dt(uint32_t old_icb, uint32_t new_icb, uint32_t ocb)

Uninitializes the tilize operation and reconfigures the unpacker with CB data types.

NOTE: This function is not in line with our programming model, and will be removed by the end of 2025 as a part of tt-metal#22904.

Return value: None

Param Type Name

Function old_icb

Function new_icb

Function ocb