Mojo function

handle_stmatrix_output

handle_stmatrix_output[c_type: DType, c_tma_layout: Layout, c_desc_layout: Layout, accum_type: DType, c_layout: Layout, c_tile_layout: Layout, c_reg_layout: Layout, //, wgmma_shape: IndexList[3], num_consumer: Int, use_tma_store: Bool, elementwise_lambda_fn: OptionalReg[fn[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> None], elementwise_compute_lambda_fn: OptionalReg[fn[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> SIMD[dtype, width]], BM: Int, BN: Int, num_m_mmas: Int, num_consumer_threads: Int, simd_size: Int](c_tma_op: TMATensorTile[c_type, c_tma_layout, c_desc_layout], c: LayoutTensor[c_type, c_layout, MutableAnyOrigin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], c_tile: LayoutTensor[c_type, c_tile_layout, MutableAnyOrigin, address_space=AddressSpace(3), alignment=128], c_reg_tile: LayoutTensor[accum_type, c_reg_layout, MutableAnyOrigin, address_space=AddressSpace(5), alignment=alignment], c_gmem_tile: LayoutTensor[c_type, layout, MutableAnyOrigin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], c_gmem_corner_coords: IndexList[len[IntTuple](flatten(c_layout.shape)), element_type=layout_int_type], warp_group_thread_idx: UInt, local_warp_group_idx: UInt, local_thread_idx: UInt, block_y: Int, block_x: Int)

Handle output using st.matrix instructions for optimized bf16 output.