o پi/,@sddlmZddlmZddlmZddlmZmZm Z ddl m Z m Z ddl mZmZmZmZddl mZmZmZmZddl mZGd d d eZd ed efd dZeddGdddeZGdddeZeddGdddeZdS))Optional) dataclassN)BooleanInt32 const_expr) if_generateand_) MbarrierArrayCooperativeGroup PipelineOppipeline_init_wait) PipelineAsyncPipelineTmaAsync PipelineStatePipelineUserType)PipelineTmaUmmac@s"eZdZdefddZddZdS)PipelineStateWAdvancenum_iterationscCsH|jt|7_|jt|}||j}|j|N_||j|_dSN)_countr_indexstages_phase)selfr new_index num_crossingsrB/home/ubuntu/.local/lib/python3.10/site-packages/quack/pipeline.py advance_iterss  z#PipelineStateWAdvance.advance_iterscCs(t|jt|dt|dt|dS)Nr)rrrrvaluesrrr__new_from_mlir_values__s"z.PipelineStateWAdvance.__new_from_mlir_values__N)__name__ __module__ __qualname__rrr#rrrrrs rtypercCsP|tjurt|tdtdtdS|tjur$t|tdtdtdSJd)zz Creates a pipeline state. Producers are assumed to start with an empty buffer and have a flipped phase bit of 1. rrFzBError: invalid PipelineUserType specified for make_pipeline_state.)rProducerrrConsumer)r'rrrrmake_pipeline_states  r*T)frozenc@eZdZdZedddddededededejd e ej d e e fd d Z  dde de ede efddZde fddZdS)PipelineTmaCpAsynczZ PipelineTmaCpAsync is used for CpAsync + TMA producers and AsyncThread consumers N)barrier_storagecta_layout_vmnktidx num_stagesproducer_groupconsumer_grouptx_countr.r/r0cCst|tjstdt|tj}tj}||f} ||f} t |j dd|| |} t |j dd||| } |durBtj \}} } |durKt d}t||\}}|dus^t|dkrad}n|}d}t|t| | ||||S)a This helper function computes any necessary attributes and returns an instance of PipelineTmaAsync. :param barrier_storage: Pointer to the smem address for this pipeline's mbarriers :type barrier_storage: cute.Pointer :param num_stages: Number of buffer stages for this pipeline :type num_stages: Int32 :param producer_group: CooperativeGroup for the producer agent :type producer_group: CooperativeGroup :param consumer_group: CooperativeGroup for the consumer agent :type consumer_group: CooperativeGroup :param tx_count: Number of bytes expected to be written to the transaction barrier for one stage :type tx_count: int :param cta_layout_vmnk: Layout of the cluster shape :type cta_layout_vmnk: cute.Layout | None :param tidx: thread index to consumer async threads :type tidx: Int32 | None 7Expected barrier_storage to be a cute.Pointer, but got  min_alignN)rrrrr) isinstancecutePointer ValueErrorr'r TmaLoad AsyncThreadr _make_sync_objectalignarch thread_idx make_layoutr init_empty_barrier_arrive_signalsizer r-)r1r2r3r4r.r/r0 producer_type consumer_typeproducerconsumersync_object_fullsync_object_empty_dst_rankis_signalling_thread producer_maskrrrcreate:sF    zPipelineTmaCpAsync.createTstatetry_acquire_token is_tma_warpcs8t|dup|dkfddt|fdddS)zk TMA producer commit conditionally waits on buffer empty and sets the transaction barrier. NrcjjjSrrKwaitindexphaserrrQrrz5PipelineTmaCpAsync.producer_acquire..cjjjSrrJarriverWrOrrYrrrZr[)rrrQrRrSrrYrproducer_acquires   z#PipelineTmaCpAsync.producer_acquirecCtj||dSzJ We need the mbarrier to track the completion of cp.async Nr:rAcp_async_mbarrier_arrive_noincproducer_get_barrierrYrrrproducer_cpasync_commitz*PipelineTmaCpAsync.producer_cpasync_commitNTr$r%r&__doc__ staticmethodintr r:r;rLayoutrrPrrr`rfrrrrr-4s> I r-c @sReZdZ  ddejdedeeefdede e ddf d d Z d d Z d dZ dS)MbarrierArrayWDropCountrNr.r1agentr4 drop_countreturncCs||_||_||_|\|_|_|jj|_||_|jdkr td|jdkr)td|jt j ur8|jdkr8tdt |durD|j||_|j|_ | dS)Nrz3Error: Mbarrier stage count must be greater than 0.z4Error: Mbarrier arrive count must be greater than 0.z=Error: Mbarrier tx count must not be less than 0 for TMA ops.)r.r4r1op_typecgrE arrive_countrpr<r r=r mbarrier_base mbarrier_init)rr.r1ror4rprrr__init__s        z MbarrierArrayWDropCount.__init__cCs |j|jgSr)r.rp)rrrr__extract_mlir_values__s z/MbarrierArrayWDropCount.__extract_mlir_values__cCs$t|d|j|j|jf|j|dS)Nrr)rnr1rrrsr4r!rrrr#sz0MbarrierArrayWDropCount.__new_from_mlir_values__)rN)r$r%r&r:r;rltupler r rrrwrxr#rrrrrns$   rnc@r,)PipelineTmaCpAsyncUmmazr PipelineTmaCpAsync is used for CpAsync + TMA producers and UMMA consumers (e.g. Blackwell mainloops) N)r.r/producer_drop_countr1r2r3r4r.r/r{cCst|tjstdt|tj}tj}||f} ||f} t|j dd|| ||d} t |j dd||| } |dusBt |dkrGd} d}n t |} t |}|dus_tj |dgd dkretjjjjntjjjj}| }t|t| | || |||S) aB This helper function computes any necessary attributes and returns an instance of PipelineTmaUmma. :param barrier_storage: Pointer to the smem address for this pipeline's mbarriers :type barrier_storage: cute.Pointer :param num_stages: Number of buffer stages for this pipeline :type num_stages: Int32 :param producer_group: `CooperativeGroup` for the producer agent :type producer_group: CooperativeGroup :param consumer_group: `CooperativeGroup` for the consumer agent :type consumer_group: CooperativeGroup :param tx_count: Number of bytes expected to be written to the transaction barrier for one stage :type tx_count: int :param cta_layout_vmnk: Layout of the cluster shape :type cta_layout_vmnk: cute.Layout | None r5r6r7)rpNrTr)mode)r9r:r;r<r'r r= TCGen05Mmarnr@r r?rEr_compute_mcast_arrival_mask_compute_is_leader_ctanvgputcgen05CtaGroupONETWOr rz)r1r2r3r4r.r/r{rFrGrHrIrJrKrO is_leader_cta cta_group consumer_maskrrrrPsL      zPipelineTmaCpAsyncUmma.createTrQrRrScs@t|dup|dkfddttj|fdddS)z TMA producer commit conditionally waits on buffer empty and sets the transaction barrier for leader threadblocks. NrcrTrrUrrYrrrZ%r[z9PipelineTmaCpAsyncUmma.producer_acquire..cr\rr]rrYrrrZ+r[)rrrr_rrYrr`s    z'PipelineTmaCpAsyncUmma.producer_acquirecCrarbrcrYrrrrf.rgz.PipelineTmaCpAsyncUmma.producer_cpasync_commitrhrirrrrrzs> O rz)typingr dataclassesr cutlass.cuter:cutlassrrrcutlass.cutlass_dslrrcutlass.pipeliner r r r r rrrrrrlr*r-rnrzrrrrs    g)