o پi4'@sddlmZddlmZddlZddlmZddlmZmZm Z ddl m Z ddl m Z mZmZmZddl mZmZddl mZdd l mZdd eejfd d Zd efddZGdddZdedefddZeddGdddeZeddGdddeZdS))Optional) dataclassN)BooleanInt32 const_expr) if_generate) PipelineAsync PipelineStateAgentCooperativeGroup)PipelineUserType PipelineOp)PipelineTmaAsync)PipelineTmaUmmacta_layout_vmnkcCs<tj|dust|dkrttjdSttjdS)zG Fences the mbarrier init and syncs the threadblock or cluster N)cutearchmbarrier_init_fencesize_syncr ThreadBlockThreadBlockCluster)rrS/home/ubuntu/.local/lib/python3.10/site-packages/flash_attn_origin/cute/pipeline.pypipeline_init_waits rgroupcCsT|tjur td|tjurtjdS|tjur&tjtj dSJd)z, Syncs all threads within an agent. zError: Not supported.Fz[Error: No explicit sync instruction exists. Please use barriers (named / mbarrier) instead.N) r ThreadNotImplementedErrorrrr sync_threadsrcluster_arrive_relaxed cluster_wait)rrrrr s    rc@szeZdZdZdedefddZdddZedefd d Z edefd d Z edefd dZ ddZ ddZ ddZdS)PipelineStateSimplea# Pipeline state contains an index and phase bit corresponding to the current position in the circular buffer. Use a single Int32 to store both the index and phase bit, then we use divmod to get the index and phase. If stages is a power of 2, divmod turns into bit twiddling. stages phase_indexcCs||_||_dSN)_stages _phase_index)selfr#r$rrr__init__8s zPipelineStateSimple.__init__returncCst|j|jSr%)r"r#r'r(rrrclone?szPipelineStateSimple.clonecCs|jSr%)r&r+rrrr#BszPipelineStateSimple.stagescCs"t|jdkr tdS|j|jS)Nrr)rr&rr'r+rrrindexGs zPipelineStateSimple.indexcCs t|jdkr |jS|j|jSNrrr&r'r+rrrphasePs zPipelineStateSimple.phasecCs2t|jdkr|jdN_dS|jd7_dSr.r/r+rrradvance\szPipelineStateSimple.advancecCs|j}|gSr%)r'ir_value)r(r$rrr__extract_mlir_values__qs z+PipelineStateSimple.__extract_mlir_values__cCst|jt|dSNr)r"r#r)r(valuesrrr__new_from_mlir_values__usz,PipelineStateSimple.__new_from_mlir_values__N)r*r")__name__ __module__ __qualname____doc__intrr)r,propertyr#r-r0r1r3r6rrrrr"1s   r"typer#cCs8|tjur t|t|S|tjurt|tdSJd)zz Creates a pipeline state. Producers are assumed to start with an empty buffer and have a flipped phase bit of 1. rFzBError: invalid PipelineUserType specified for make_pipeline_state.)r Producerr"rConsumer)r=r#rrrmake_pipeline_stateys  r@T)frozenc@s<eZdZdZeddZ  d dedeede fd d Z dS) rzH Override producer_acquire to take in extra_tx_count parameter. cOs"tj|i|}t|dt|S)N __class__)PipelineTmaAsyncOgcreateobject __setattr__r)argskwargsobjrrrrDszPipelineTmaAsync.createNrstatetry_acquire_tokenextra_tx_countcsbt|dup|dkfddt|dkr!jjjdSjj|}jj|dS) TMA producer commit conditionally waits on buffer empty and sets the transaction barrier for leader threadblocks. NrcjjjSr%sync_object_emptywaitr-r0rr(rJrrz3PipelineTmaAsync.producer_acquire..)rrsync_object_fullarriver- producer_masktx_countarrive_and_expect_tx)r(rJrKrLrXrrRrproducer_acquires    z!PipelineTmaAsync.producer_acquirer4) r7r8r9r: staticmethodrDr rrr;rZrrrrrs  rc@s|eZdZeddddddedededed ejd eej d e eefd e j e fd dZ  ddedeedefddZdS)rN)rrT)barrier_storager mcast_mode_mn init_wait num_stagesproducer_groupconsumer_grouprXr\rr]r^cCst|tjstdt|tj}tj} ||f} | |f} t |j dd|| |} t |j dd||| } |dusAt |dkrFd}d}n t ||}t |}|dus_tj |dgddkretjjjjntjjjj}|}t|rut|t | | |||||S) a  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 :param mcast_mode_mn: Tuple of two integers, specifying whether mcast is enabled for the m and n modes. At least one of the two integers must be 1. :type mcast_mode_mn: tuple[int, int] z7Expected barrier_storage to be a cute.Pointer, but got ) min_alignNrTr)mode) isinstancerPointer ValueErrorr=r TmaLoad TCGen05Mmar_make_sync_objectalignrr_compute_mcast_arrival_mask_compute_is_leader_ctanvgputcgen05CtaGroupONETWOrr)r_r`rarXr\rr]r^ producer_type consumer_typeproducerconsumerrUrPrW is_leader_cta cta_group consumer_maskrrrrDsJ    zPipelineTmaUmma.createrrJrKrLcsnt|dup|dkfddt|dkr#tjfdddSjj|tjfdddS)rMNrcrNr%rOrrRrrrSrTz2PipelineTmaUmma.producer_acquire..csjjjSr%)rUrVr-rWrrRrrrS rTcsjjSr%)rUrYr-rr(rJrXrrrSs)rrrwrUrX)r(rJrKrLrrzrrZs     z PipelineTmaUmma.producer_acquirer4)r7r8r9r[r;r rrfrLayouttuplecutlass ConstexprboolrDr rrZrrrrrs@  Qrr%)typingr dataclassesrr} cutlass.cuterrrrcutlass.cutlass_dslrcutlass.pipelinerr r r r r rrCrPipelineTmaUmmaOgr{rrr"r;r@rrrrs$      H !