o c۷i)@sddlmZddlmZddlmZddlmZmZm Z ddl m Z m Z m Z ddlmZmZmZddlmZmZmZddlmZdd lmZmZGd d d eZd ed efddZeddGdddeZGdddeZeddGdddeZdS))Optional) dataclassN)BooleanInt32 const_expr) if_generateand_ dsl_user_op) MbarrierArrayCooperativeGroup PipelineOp)PipelineTmaAsync PipelineStatePipelineUserType)PipelineTmaUmma)Agent agent_syncc@s.eZdZeddddefddZddZdS)PipelineStateWAdvanceNlocipnum_iterationscCsH|jt|7_|jt|}||j}|j|N_||j|_dSN)_countr_indexstages_phase)selfrrr new_index num_crossingsr D/home/ubuntu/vllm_env/lib/python3.10/site-packages/quack/pipeline.py advance_iterss  z#PipelineStateWAdvance.advance_iterscCs(t|jt|dt|dt|dS)Nr)rrrrvaluesr r r!__new_from_mlir_values__s"z.PipelineStateWAdvance.__new_from_mlir_values__)__name__ __module__ __qualname__r rr"r'r r r r!rs 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. rr#FzBError: invalid PipelineUserType specified for make_pipeline_state.)rProducerrrConsumer)r+rr r r!make_pipeline_state s  r.T)frozenc @sfeZdZdZeddZe  dddddedee d ee fd d Z eddddefd d Z dS)PipelineTmaCpAsynczZ PipelineTmaCpAsync is used for CpAsync + TMA producers and AsyncThread consumers cOs"tj|i|}t|dt|S)N __class__)r createobject __setattr__r0)argskwargsobjr r r!r2<szPipelineTmaCpAsync.createNTrstatetry_acquire_token is_tma_warpcsLt|dup|dkfdddt|fddddS)zk TMA producer commit conditionally waits on buffer empty and sets the transaction barrier. NrcjjjjdSNrsync_object_emptywaitindexphaser rrrr8r r!Sz5PipelineTmaCpAsync.producer_acquire..rcjjjjdSr<sync_object_fullarriver@ producer_maskr rBr r!rC[rD)rrr8r9r:rrr rBr!producer_acquireDs  z#PipelineTmaCpAsync.producer_acquirecC"tjj|j|||d||ddSzJ We need the mbarrier to track the completion of cp.async rNcutearchcp_async_mbarrier_arrive_noincproducer_get_barrierrr8rrr r r!producer_cpasync_commit` z*PipelineTmaCpAsync.producer_cpasync_commitNT) r(r)r*__doc__ staticmethodr2r rrrrKrTr r r r!r06s&  r0c@s^eZdZe  dddddejdedeee fdede e d df d d Z d d Z ddZdS)MbarrierArrayWDropCountrNrbarrier_storage num_stagesagenttx_count drop_countreturncCs||_||_||_|\|_|_|jj|_||_|jdkr td|jdkr)td|jt j ur8|jdkr8tdt |durD|j||_|j|_ |j ||ddS)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)rZr]r[op_typecgsize arrive_countr^ ValueErrorr TmaLoadr mbarrier_base mbarrier_init)rrZr[r\r]r^rrr r r!__init__ks        z MbarrierArrayWDropCount.__init__cCs |j|jgSr)rZr^)rr r r!__extract_mlir_values__s z/MbarrierArrayWDropCount.__extract_mlir_values__cCs$t|d|j|j|jf|j|dS)Nrr#)rYr[r`rar]r%r r r!r'sz0MbarrierArrayWDropCount.__new_from_mlir_values__)rN)r(r)r*r rOPointerinttupler r rrrhrir'r r r r!rYjs,  " rYc@seZdZdZeedddddddddededed ed ej d e ej d e eefd e de efddZe  dddddede ede efddZeddddefddZdS)PipelineTmaCpAsyncUmmazr PipelineTmaCpAsync is used for CpAsync + TMA producers and UMMA consumers (e.g. Blackwell mainloops) N)r#r#F)rZcta_layout_vmnk mcast_mode_mn defer_syncproducer_drop_countrrr[producer_groupconsumer_groupr]rZrnrorprqc  CsTt|tjstdt|tj} tj} | |f} | |f}t|j dd|| ||| | d}t j |j dd|||| | d}|dusJtj || | ddkrOd}d}nt j ||| | d}t j|| | d}|dusptj |d g| | d dkrvtjjjjntjjjj}|}|stj|dustj || | ddkrttjnttjdd t|||||||S) aCreates and initializes a new PipelineTmaUmma instance. :param num_stages: Number of buffer stages for this pipeline :type num_stages: int :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 barrier_storage: Pointer to the shared memory address for this pipeline's mbarriers :type barrier_storage: cute.Pointer, optional :param cta_layout_vmnk: Layout of the cluster shape :type cta_layout_vmnk: cute.Layout, optional :param mcast_mode_mn: Tuple specifying multicast modes for m and n dimensions (each 0 or 1) :type mcast_mode_mn: tuple[int, int], optional :raises ValueError: If barrier_storage is not a cute.Pointer instance :return: A new PipelineTmaUmma instance configured with the provided parameters :rtype: PipelineTmaUmma z7Expected barrier_storage to be a cute.Pointer, but got ) min_align)r^rrrNr#Tr)moderr) is_relaxed) isinstancerOrj TypeErrorr+r re TCGen05MmarYalignr_make_sync_objectrb_compute_mcast_arrival_mask_compute_is_leader_ctanvgputcgen05CtaGroupONETWOrPmbarrier_init_fencerr ThreadBlockThreadBlockClusterrm)r[rrrsr]rZrnrorprqrr producer_type consumer_typeproducerconsumerrGr>rI is_leader_cta cta_group consumer_maskr r r!r2sd $       zPipelineTmaCpAsyncUmma.createTrr8r9r:csTt|dup|dkfdddttj|fddddS)z TMA producer commit conditionally waits on buffer empty and sets the transaction barrier for leader threadblocks. Nrcr;r<r=r rBr r!rCrDz9PipelineTmaCpAsyncUmma.producer_acquire..rcrEr<rFr rBr r!rCrD)rrrrJr rBr!rKs  z'PipelineTmaCpAsyncUmma.producer_acquirecCrLrMrNrSr r r!rTrUz.PipelineTmaCpAsyncUmma.producer_cpasync_commitrV)r(r)r*rWr rXrkr rOrjrLayoutrlboolrr2rrrKrTr r r r!rmsZ   b rm)typingr dataclassesr cutlass.cuterOcutlassrrrcutlass.cutlass_dslrrr cutlass.pipeliner r r r rrrrrrrkr.r0rYrmr r r r!s    3-