Skip to main content
Log in

Mojo function

producer_main_loop

producer_main_loop[a_type: DType, b_type: DType, a_tile_layout: Layout, b_tile_layout: Layout, a_smem_layout: Layout, b_smem_layout: Layout, a_desc_layout: Layout, b_desc_layout: Layout, pipeline_stages: Int, /, *, block_tile_shape: Index[3], cluster_shape: StaticTuple[SIMD[int32, 1], 3] = StaticTuple(__init__[__mlir_type.!pop.int_literal](1), __init__[__mlir_type.!pop.int_literal](1), __init__[__mlir_type.!pop.int_literal](1)), partitioned_multicast: Bool = False](a_tma_op: TMATensorTile[a_type, a_tile_layout, a_desc_layout], b_tma_op: TMATensorTile[b_type, b_tile_layout, b_desc_layout], a_smem_iter: LayoutTensorIter[a_type, a_smem_layout, MutableAnyOrigin, address_space=AddressSpace(3), alignment=128, circular=circular, axis=axis, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked], b_smem_iter: LayoutTensorIter[b_type, b_smem_layout, MutableAnyOrigin, address_space=AddressSpace(3), alignment=128, circular=circular, axis=axis, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked], num_k_iters: Int, m_coord: UInt, n_coord: UInt, rank_n: UInt, rank_m: UInt, mut write_pipeline_states: PipelineState[pipeline_stages], empty_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3), alignment=8], full_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3), alignment=8])