Ë
    ÄÓ^jZ  ã                   óœ  — d dl mZ d dlmZ d dlmZ d dlmZmZm	Z	 d dl
mZmZmZ d dlmZmZmZ d dlmZmZ d dlmZmZ d d	lmZ d d
lmZ d dlmZ d dlmZ d dlmZ d dlm Z! d dlm"Z# d„ Z$dededefd„Z%d„ Z& G d„ de«      Z'dede(fd„Z) G d„ d«      Z* ed¬«       G d„ d e«      «       Z e$ee«      e_+         ed¬«       G d!„ d"e*e«      «       Z ed¬«       G d#„ d$e*e«      «       Z ed¬«       G d%„ d&e*e«      «       Z e$ee«      e_+         ed¬«       G d'„ d(e*e«      «       Z e$ee«      e_+         ed¬«       G d)„ d*e*e!«      «       Z  e$e!e «      e _+         ed¬«       G d+„ d,e*e#«      «       Z" e$e#e"«      e"_+         ed¬«       G d-„ d.e*e«      «       Z, e$ee,«      e,_+         G d/„ d0e«      Z- ed¬«       G d1„ d2e«      «       Z.y)3é    )ÚOptional)Ú	dataclassN)ÚBooleanÚInt32Ú
const_expr)Úif_generateÚand_Údsl_user_op)ÚMbarrierArrayÚCooperativeGroupÚ
PipelineOp)ÚPipelineStateÚPipelineUserType)ÚAgentÚ
agent_sync)ÚNamedBarrier)ÚPipelineAsync)ÚPipelineCpAsync)ÚPipelineTmaAsync)ÚPipelineTmaUmma)ÚPipelineUmmaAsync)ÚPipelineAsyncUmmac                 ó(   ‡ ‡— t         ˆˆ fd„«       }|S )zPCreate a static factory that constructs parent_cls then re-classes to child_cls.c                  óZ   •—  ‰j                   | i |¤Ž}t        j                  |d‰«       |S )NÚ	__class__)ÚcreateÚobjectÚ__setattr__)ÚargsÚkwargsÚobjÚ	child_clsÚ
parent_clss      €€úg/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_vendor/quack/pipeline.pyr   z _override_create.<locals>.create   s1   ø€ àˆj×Ñ Ð0¨Ñ0ˆä×Ñ˜3 ¨YÔ7Øˆ
ó    )Ústaticmethod)r#   r"   r   s   `` r$   Ú_override_creater'      s   ù€ ô ôó ðð €Mr%   ÚindexÚphaseÚreturnc                 ó2   — t        dt        d«      | |¬«      S )zPConstruct a PipelineState from index and phase (count/stages unused by callers).r   )ÚstagesÚcountr(   r)   )r   r   )r(   r)   s     r$   Ú_make_stater.   %   s   € ä ¬¨q«¸ÀeÔLÐLr%   c                 ó  — t        |«      r]t        |«      rt        j                  j                  «        t        j                  j	                  «       5   | ||||¬«       ddd«       y | ||||¬«       y# 1 sw Y   yxY w)zGOptionally wrap a parent pipeline method call in sync_warp + elect_one.©ÚlocÚipN)r   ÚcuteÚarchÚ	sync_warpÚ	elect_one)Úparent_methodÚselfÚstater6   Úsyncwarpr1   r2   s          r$   Ú_call_with_elect_oner;   *   sj   € ä�)ÔÜ�hÔÜ�I‰I×ÑÔ!Ü�Y‰Y× Ñ Ó"ñ 	7Ù˜$ ¨3°2Õ6÷	7ð 	7ñ 	�d˜E s¨rÖ2÷	7ð 	7ús   ÁA6Á6A?c                   ó0   — e Zd Zedddœdefd„«       Zd„ Zy)ÚPipelineStateWAdvanceNr0   Únum_iterationsc                óà   — | xj                   t        |«      z  c_         | j                  t        |«      z   }|| j                  z  }| xj                  |z  c_        || j                  z  | _        y ©N)Ú_countr   Ú_indexr,   Ú_phase)r8   r>   r1   r2   Ú	new_indexÚnum_crossingss         r$   Úadvance_itersz#PipelineStateWAdvance.advance_iters9   sT   € à�Š”u˜^Ó,Ñ,�Ø—K‘K¤%¨Ó"7Ñ7ˆ	à! T§[¡[Ñ0ˆØ�Š�}Ñ$�Ø $§+¡+Ñ-ˆ�r%   c           	      óz   — t        | j                  t        |d   «      t        |d   «      t        |d   «      «      S )Nr   é   é   )r=   r,   r   ©r8   Úvaluess     r$   Ú__new_from_mlir_values__z.PipelineStateWAdvance.__new_from_mlir_values__C   s7   € Ü$Ø�K‰Kœ˜v a™yÓ)¬5°¸±Ó+;¼UÀ6È!Á9Ó=Mó
ð 	
r%   )Ú__name__Ú
__module__Ú__qualname__r
   r   rF   rL   © r%   r$   r=   r=   8   s%   „ ØØ:>À4ò .¨Eò .ó ð.ó
r%   r=   Útyper,   c                 óü   — | t         j                  u r)t        |t        d«      t        d«      t        d«      «      S | t         j                  u r)t        |t        d«      t        d«      t        d«      «      S J d«       ‚)zz
    Creates a pipeline state. Producers are assumed to start with an empty buffer and have a flipped phase bit of 1.
    r   rH   zBError: invalid PipelineUserType specified for make_pipeline_state.)r   ÚProducerr=   r   ÚConsumer)rQ   r,   s     r$   Úmake_pipeline_staterU   I   sg   € ð Ô×(Ñ(Ñ(Ü$ V¬U°1«X´u¸Q³xÄÀqÃÓJÐJØ	Ô!×*Ñ*Ñ	*Ü$ V¬U°1«X´u¸Q³xÄÀqÃÓJÐJàZÐZÓZˆur%   c            
       ó¬   — e Zd ZdZe	 ddddœdededee   fd„«       Zedddœdefd„«       Z	e	 ddddœdeded	ee   fd
„«       Z
edddœdefd„«       Zy)Ú_PipelineIndexPhaseMixinz_Mixin providing _w_index_phase / _w_index methods that delegate to PipelineState-based parents.Nr0   r(   r)   Útry_acquire_tokenc                óF   — t        ||«      }| j                  ||||¬«       y ©Nr0   )r.   Úproducer_acquire)r8   r(   r)   rX   r1   r2   r9   s          r$   Úproducer_acquire_w_index_phasez7_PipelineIndexPhaseMixin.producer_acquire_w_index_phase[   s(   € ô ˜E 5Ó)ˆØ×Ñ˜eÐ%6¸CÀBÐÕGr%   c                óV   — t        |t        d«      «      }| j                  |||¬«       y ©Nr   r0   )r.   r   Úproducer_commit©r8   r(   r1   r2   r9   s        r$   Úproducer_commit_w_indexz0_PipelineIndexPhaseMixin.producer_commit_w_indexh   s'   € ä˜E¤5¨£8Ó,ˆØ×Ñ˜U¨°ÐÕ3r%   Útry_wait_tokenc                óF   — t        ||«      }| j                  ||||¬«       y rZ   )r.   Úconsumer_wait)r8   r(   r)   rb   r1   r2   r9   s          r$   Úconsumer_wait_w_index_phasez4_PipelineIndexPhaseMixin.consumer_wait_w_index_phasem   s'   € ô ˜E 5Ó)ˆØ×Ñ˜5 .°c¸bÐÕAr%   c                óV   — t        |t        d«      «      }| j                  |||¬«       y r^   )r.   r   Úconsumer_releaser`   s        r$   Úconsumer_release_w_indexz1_PipelineIndexPhaseMixin.consumer_release_w_indexz   s'   € ä˜E¤5¨£8Ó,ˆØ×Ñ˜e¨°ÐÕ4r%   r@   )rM   rN   rO   Ú__doc__r
   r   r   r   r\   ra   re   rh   rP   r%   r$   rW   rW   X   sÝ   „ Ùiàð
 04ð	
Hð Øò
Hàð
Hð ð
Hð $ GÑ,ò	
Hó ð
Hð Ø;?ÀDò 4¨Uò 4ó ð4ð ð
 -1ð	
Bð Øò
Bàð
Bð ð
Bð ! Ñ)ò	
Bó ð
Bð Ø<@ÀTò 5¨eò 5ó ñ5r%   rW   T)Úfrozenc                   ób   — e Zd Z eed«      Zedddœdeddfd„«       Zedddœdeddfd„«       Z	y)r   Nr0   r(   r*   c                óx   — t         j                  j                  | j                  |z   | j                  ||¬«       y)z”
        The aligned flavor of arrive is used when all threads in the CTA will execute the
        same instruction. See PTX documentation.
        ©Ú
barrier_idÚnumber_of_threadsr1   r2   N)r3   r4   Úbarrier_arrivern   Únum_threads©r8   r(   r1   r2   s       r$   Úarrive_w_indexzNamedBarrier.arrive_w_index‡   s6   € ô 	�	‰	× Ñ Ø—‘¨Ñ.Ø"×.Ñ.ØØð	 	!õ 	
r%   c                óx   — t         j                  j                  | j                  |z   | j                  ||¬«       y )Nrm   )r3   r4   Úbarrierrn   rq   rr   s       r$   Úarrive_and_wait_w_indexz$NamedBarrier.arrive_and_wait_w_index”   s4   € ä�	‰	×ÑØ—‘¨Ñ.Ø"×.Ñ.ØØð	 	õ 	
r%   )
rM   rN   rO   r'   ÚNamedBarrierOgr   r
   r   rs   rv   rP   r%   r$   r   r   ƒ   sZ   „ á˜n¨dÓ3€FàØ26¸4ò 

 Eð 

ÀDò 

ó ð

ð Ø;?ÀDò 
¨Uð 
ÈTò 
ó ñ
r%   r   c            
       ó´   — e Zd ZU dZdZeed<   dZeed<   dZeed<   dZ	eed<   e
dddddœd	ed
ededefd„«       Zedddœdefd„«       Zedddœdefd„«       Zy)r   aÊ  
    PipelineAsync with optional elect_one for producer_commit and consumer_release.

    When elect_one_*=True (set at create time), only one elected thread per warp
    signals the barrier arrive. This is useful when the mask count is set to 1 per warp.

    Args (to create):
        elect_one_commit: If True, only elected thread signals producer_commit.
        syncwarp_before_commit: If True (default), issue syncwarp before elect_one.
        elect_one_release: If True, only elected thread signals consumer_release.
        syncwarp_before_release: If True (default), issue syncwarp before elect_one.
            Set syncwarp to False when threads are already converged (e.g. after wgmma wait_group).
    FÚ_elect_one_commitTÚ_syncwarp_before_commitÚ_elect_one_releaseÚ_syncwarp_before_release)Úelect_one_commitÚsyncwarp_before_commitÚelect_one_releaseÚsyncwarp_before_releaser}   r~   r   r€   c                 ó  — t        j                  |i |¤Ž}t        j                  |dt        «       t        j                  |d| «       t        j                  |d|«       t        j                  |d|«       t        j                  |d|«       |S )Nr   ry   rz   r{   r|   )ÚPipelineAsyncOgr   r   r   r   )r}   r~   r   r€   r   r    r!   s          r$   r   zPipelineAsync.create¹   s€   € ô ×$Ñ$ dÐ5¨fÑ5ˆÜ×Ñ˜3 ¬]Ô;Ü×Ñ˜3Ð 3Ð5EÔFÜ×Ñ˜3Ð 9Ð;QÔRÜ×Ñ˜3Ð 4Ð6GÔHÜ×Ñ˜3Ð :Ð<SÔTØˆ
r%   Nr0   r9   c          	      ój   — t        t        j                  | || j                  | j                  ||«       y r@   )r;   r‚   r_   ry   rz   ©r8   r9   r1   r2   s       r$   r_   zPipelineAsync.producer_commitÊ   s1   € äÜ×+Ñ+ØØØ×"Ñ"Ø×(Ñ(ØØõ	
r%   c          	      ój   — t        t        j                  | || j                  | j                  ||«       y r@   )r;   r‚   rg   r{   r|   r„   s       r$   rg   zPipelineAsync.consumer_releaseÖ   s1   € äÜ×,Ñ,ØØØ×#Ñ#Ø×)Ñ)ØØõ	
r%   )rM   rN   rO   ri   ry   ÚboolÚ__annotations__rz   r{   r|   r&   r   r
   r   r_   rg   rP   r%   r$   r   r   ¤   s»   … ñð $Ð�tÓ#Ø$(Ð˜TÓ(Ø$Ð˜Ó$Ø%)Ð˜dÓ)àð "'Ø'+Ø"'Ø(,òàðð !%ðð  ð	ð
 "&òó ðð  Ø;?ÀDò 	
 ]ò 	
ó ð	
ð Ø<@ÀTò 	
 mò 	
ó ñ	
r%   r   c                   ój   — e Zd ZU dZeed<   dZeed<   edddœdedefd„«       Ze	d	d	d
œde
fd„«       Zy	)r   Fr{   Tr|   )r   r€   r   r€   c                 óÂ   — t        j                  |i |¤Ž}t        j                  |dt        «       t        j                  |d| «       t        j                  |d|«       |S )Nr   r{   r|   )ÚPipelineCpAsyncOgr   r   r   r   )r   r€   r   r    r!   s        r$   r   zPipelineCpAsync.createî   sX   € ô  ×&Ñ&¨Ð7°Ñ7ˆÜ×Ñ˜3 ¬_Ô=Ü×Ñ˜3Ð 4Ð6GÔHÜ×Ñ˜3Ð :Ð<SÔTØˆ
r%   Nr0   r9   c          	      ój   — t        t        j                  | || j                  | j                  ||«       y r@   )r;   rŠ   rg   r{   r|   r„   s       r$   rg   z PipelineCpAsync.consumer_releaseû   s1   € äÜ×.Ñ.ØØØ×#Ñ#Ø×)Ñ)ØØõ	
r%   )rM   rN   rO   r{   r†   r‡   r|   r&   r   r
   r   rg   rP   r%   r$   r   r   é   sd   … à$Ð˜Ó$Ø%)Ð˜dÓ)àð #(Ø(,ò
àð
ð "&ò
ó ð
ð Ø<@ÀTò 	
 mò 	
ó ñ	
r%   r   c            	       óB   — e Zd ZdZe	 	 ddddœdedee   defd„«       Z	y)	r   ú>Override producer_acquire to take in extra_tx_count parameter.Nr0   r9   rX   Úextra_tx_countc                óP  ‡ ‡‡‡— t        |du xs |dk(  ˆˆˆ ˆfd„‰‰¬«       t        |dk(  «      r4‰ j                  j                  ‰j                  ‰ j
                  ‰‰¬«       y‰ j                  j                  |z   }‰ j                  j                  ‰j                  |‰‰¬«       y)úƒ
        TMA producer commit conditionally waits on buffer empty and sets the transaction barrier for leader threadblocks.
        Nr   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   ©Úsync_object_emptyÚwaitr(   r)   ©r2   r1   r8   r9   s   €€€€r$   ú<lambda>z3PipelineTmaAsync.producer_acquire.<locals>.<lambda>   ó*   ø€ �D×*Ñ*×/Ñ/°·±¸U¿[¹[ÈcÐVXÐ/ÓY€ r%   r0   )r   r   Úsync_object_fullÚarriver(   Úproducer_maskÚtx_countÚarrive_and_expect_tx)r8   r9   rX   rŽ   r1   r2   r›   s   ``  `` r$   r[   z!PipelineTmaAsync.producer_acquire  s™   û€ ô 	Ø Ð%Ò?Ð):¸aÑ)?ÞYØØõ		
ô �n¨Ñ)Ô*Ø×!Ñ!×(Ñ(¨¯©°d×6HÑ6HÈcÐVXÐ(ÕYà×,Ñ,×5Ñ5¸ÑFˆHØ×!Ñ!×6Ñ6°u·{±{ÀHÐRUÐZ\Ð6Õ]r%   )Nr   ©
rM   rN   rO   ri   r
   r   r   r   Úintr[   rP   r%   r$   r   r     sS   „ áHàð 04Øð	^ð Øò^àð^ð $ GÑ,ð^ð ò	^ó ñ^r%   r   c                   óN   — e Zd ZdZe	 	 	 d	dddœdedee   dee   defd„«       Z	y)
r   r�   Nr0   r9   rX   Úis_tma_warprŽ   c                ó@  ‡ ‡‡‡‡— t        |du xs |dk(  ˆˆˆ ˆfd„‰‰¬«       t        |dk(  «      r*t        t        ‰ j                  |«      ˆˆˆ ˆfd„‰‰¬«       y‰ j                  j
                  |z   Št        t        ‰ j                  |«      ˆˆˆ ˆˆfd„‰‰¬«       y)r�   Nr   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   r’   r•   s   €€€€r$   r–   z2PipelineTmaUmma.producer_acquire.<locals>.<lambda>E  r—   r%   r0   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   ©r˜   r™   r(   rš   r•   s   €€€€r$   r–   z2PipelineTmaUmma.producer_acquire.<locals>.<lambda>N  s0   ø€ ˜×-Ñ-×4Ñ4Ø—K‘K ×!3Ñ!3¸Àð 5ó € r%   c                  óV   •— ‰j                   j                  ‰j                  ‰‰‰ ¬«      S rZ   )r˜   rœ   r(   )r2   r1   r8   r9   r›   s   €€€€€r$   r–   z2PipelineTmaUmma.producer_acquire.<locals>.<lambda>X  s+   ø€ ˜×-Ñ-×BÑBØ—K‘K ¨s°rð Có € r%   )r   r   r	   Úis_leader_ctar˜   r›   )r8   r9   rX   r    rŽ   r1   r2   r›   s   ``   ``@r$   r[   z PipelineTmaUmma.producer_acquire5  sš   ü€ ô 	Ø Ð%Ò?Ð):¸aÑ)?ÞYØØõ		
ô �n¨Ñ)Ô*ÜÜ�T×'Ñ'¨Ó5öð Øöð ×,Ñ,×5Ñ5¸ÑFˆHÜÜ�T×'Ñ'¨Ó5÷ð Øör%   )NTr   r�   rP   r%   r$   r   r   1  s_   „ áHàð 04Ø)-Øð'ð Øò'àð'ð $ GÑ,ð'ð ˜gÑ&ð	'ð
 ò'ó ñ'r%   r   c                   ó   — e Zd Zy)r   N©rM   rN   rO   rP   r%   r$   r   r   f  ó   „ àr%   r   c                   ó   — e Zd Zy)r   Nr¨   rP   r%   r$   r   r   q  r©   r%   r   c            
       óf   — e Zd ZdZe	 	 d	dddœdedee   dee   fd„«       Zedddœdefd„«       Z	y)
ÚPipelineTmaCpAsyncz¼
    PipelineTmaCpAsync is used for CpAsync + TMA producers and AsyncThread consumers.
    Compared to PipelineTmaAsync, producer_acquire gates the full-barrier arrive on is_tma_warp.
    Nr0   r9   rX   r    c                ór   ‡ ‡‡‡— t        |d u xs |dk(  ˆˆˆ ˆfd„‰‰¬«       t        |ˆˆˆ ˆfd„‰‰¬«       y )Nr   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   r’   r•   s   €€€€r$   r–   z5PipelineTmaCpAsync.producer_acquire.<locals>.<lambda>�  r—   r%   r0   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   r¤   r•   s   €€€€r$   r–   z5PipelineTmaCpAsync.producer_acquire.<locals>.<lambda>—  ó-   ø€ �D×)Ñ)×0Ñ0°·±¸d×>PÑ>PÐVYÐ^`Ð0Óa€ r%   )r   ©r8   r9   rX   r    r1   r2   s   ``  ``r$   r[   z#PipelineTmaCpAsync.producer_acquireƒ  s?   û€ ô 	Ø Ð%Ò?Ð):¸aÑ)?ÞYØØõ		
ô 	ØÞaØØö		
r%   c                ól   — t         j                  j                  | j                  |||¬«      ||¬«       y)z9We need the mbarrier to track the completion of cp.async.r0   N©r3   r4   Úcp_async_mbarrier_arrive_noincÚproducer_get_barrierr„   s       r$   Úproducer_cpasync_commitz*PipelineTmaCpAsync.producer_cpasync_commitœ  s4   € ô 	�	‰	×0Ñ0Ø×%Ñ% e°¸Ð%Ó<À#È"ð 	1õ 	
r%   ©NT)
rM   rN   rO   ri   r
   r   r   r   r[   r¶   rP   r%   r$   r¬   r¬   |  sv   „ ñð
 ð 04Ø)-ð	
ð Øò
àð
ð $ GÑ,ð
ð ˜gÑ&ò	
ó ð
ð0 ØCGÈDò 
¨]ò 
ó ñ
r%   r¬   c                   ót   — e Zd Ze	 	 ddddœdej
                  dedeee	f   dede
e   ddfd	„«       Zd
„ Zd„ Zy)ÚMbarrierArrayWDropCountNr0   Úbarrier_storageÚ
num_stagesÚagentr›   Ú
drop_countr*   c                óî  — || _         || _        || _        |\  | _        | _        | j                  j
                  | _        || _        | j                  dk  rt        d«      ‚| j                  dk  rt        d«      ‚| j                  t        j                  u r| j                  dk  rt        d«      ‚t        |d u«      r| j                  |z
  | _        | j                   | _        | j                  ||¬«       y )Nr   z3Error: 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.r0   )rº   r›   r»   Úop_typeÚcgÚsizeÚarrive_countr½   Ú
ValueErrorr   ÚTmaLoadr   Úmbarrier_baseÚmbarrier_init)r8   rº   r»   r¼   r›   r½   r1   r2   s           r$   Ú__init__z MbarrierArrayWDropCount.__init__«  sß   € ð  /ˆÔØ ˆŒØ$ˆŒØ %ÑˆŒ�d”gØ ŸG™GŸL™LˆÔØ$ˆŒà�?‰?˜aÒÜÐRÓSÐSØ×Ñ Ò!ÜÐSÓTÐTØ�<‰<œ:×-Ñ-Ñ-°$·-±-À!Ò2CÜÐ\Ó]Ð]ä�j¨Ð,Ô-Ø $× 1Ñ 1°JÑ >ˆDÔð "×1Ñ1ˆÔð 	×Ñ˜s rÐÕ*r%   c                 ó2   — | j                   | j                  gS r@   )rº   r½   )r8   s    r$   Ú__extract_mlir_values__z/MbarrierArrayWDropCount.__extract_mlir_values__Î  s   € Ø×$Ñ$ d§o¡oÐ6Ð6r%   c                 ó€   — t        |d   | j                  | j                  | j                  f| j                  |d   «      S )Nr   rH   )r¹   r»   r¿   rÀ   r›   rJ   s     r$   rL   z0MbarrierArrayWDropCount.__new_from_mlir_values__Ñ  s:   € Ü&Ø�1‰I�t—‘¨¯©°t·w±wÐ(?ÀÇÁÐPVÐWXÑPYó
ð 	
r%   )r   N)rM   rN   rO   r
   r3   ÚPointerrž   Útupler   r   r   r   rÇ   rÉ   rL   rP   r%   r$   r¹   r¹   ª  s�   „ Øð Ø&*ð +ð Øò +àŸ™ð +ð ð +ð �ZÐ!1Ð1Ñ2ð	 +ð
 ð +ð ˜U‘Oð +ð 
ò +ó ð +òD7ó
r%   r¹   c                   óö   — e Zd ZdZeeddddddddœdededed	ed
ej                  de
ej                     deeef   dede
e   fd„«       «       Ze	 	 ddddœdede
e   de
e   fd„«       Zedddœdefd„«       Zy)ÚPipelineTmaCpAsyncUmmazr
    PipelineTmaCpAsync is used for CpAsync + TMA producers and UMMA consumers
    (e.g. Blackwell mainloops)
    N)rH   rH   F)rº   Úcta_layout_vmnkÚmcast_mode_mnÚ
defer_syncÚproducer_drop_countr1   r2   r»   Úproducer_groupÚconsumer_groupr›   rº   rÏ   rÐ   rÑ   rÒ   c           	      óø  — t        |t        j                  «      st        dt	        |«      › �«      ‚t
        j                  }t
        j                  }||f}||f}t        |j                  d¬«      | ||||	|
¬«      }t        j                  |j                  d¬«      | z   | ||	|
¬«      }|�t        j                  ||	|
¬«      dk(  rd}d}n1t        j                  |||	|
¬«      }t        j                  ||	|
¬«      }|�t        j                  |d	g|	|
¬
«      dk(  r.t        j                  j                   j"                  j$                  n-t        j                  j                   j"                  j&                  }|}|spt        j(                  j+                  «        |�t        j                  ||	|
¬«      dk(  rt-        t.        j0                  «       nt-        t.        j2                  d¬«       t5        ||| ||||«      S )a€  Creates 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½   r1   r2   r0   NrH   Tr   )Úmoder1   r2   )Ú
is_relaxed)Ú
isinstancer3   rË   Ú	TypeErrorrQ   r   rÄ   Ú
TCGen05Mmar¹   ÚalignÚPipelineTmaUmmaOgÚ_make_sync_objectrÁ   Ú_compute_mcast_arrival_maskÚ_compute_is_leader_ctaÚnvgpuÚtcgen05ÚCtaGroupÚONEÚTWOr4   Úmbarrier_init_fencer   r   ÚThreadBlockÚThreadBlockClusterrÎ   )r»   rÓ   rÔ   r›   rº   rÏ   rÐ   rÑ   rÒ   r1   r2   Úproducer_typeÚconsumer_typeÚproducerÚconsumerr˜   r“   rš   r¦   Ú	cta_groupÚconsumer_masks                        r$   r   zPipelineTmaCpAsyncUmma.createá  sê  € ôH ˜/¬4¯<©<Ô8ÜØIÌ$ÈÓJ_ÐI`Ðaóð ô #×*Ñ*ˆÜ"×-Ñ-ˆà! >Ð2ˆØ! >Ð2ˆä2Ø×!Ñ!¨AÐ!Ó.ØØØØ*ØØô
Ðô .×?Ñ?Ø×!Ñ!¨AÐ!Ó.°Ñ;ØØØØô
Ðð Ð"¤d§i¡i°ÀSÈRÔ&PÐTUÒ&Uà ˆMà ‰Mä-×IÑIØ °C¸BôˆMô .×DÑDØ S¨RôˆMð Ð&¬$¯)©)°OÈ1È#ÐSVÐ[]Ô*^ÐbcÒ*cô �J‰J×Ñ×'Ñ'×+Ò+ä—‘×#Ñ#×,Ñ,×0Ñ0ð 	ð &ˆáÜ�I‰I×)Ñ)Ô+ØÐ&¬$¯)©)°OÈÐQSÔ*TÐXYÒ*YÜœ5×,Ñ,Õ-äœ5×3Ñ3ÀÕEä%ØØØØØØØó
ð 	
r%   r0   r9   rX   r    c                óš   ‡ ‡‡‡— t        |du xs |dk(  ˆˆˆ ˆfd„‰‰¬«       t        t        ‰ j                  |«      ˆˆˆ ˆfd„‰‰¬«       y)z‹
        TMA producer commit conditionally waits on buffer empty and sets the
        transaction barrier for leader threadblocks.
        Nr   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   r’   r•   s   €€€€r$   r–   z9PipelineTmaCpAsyncUmma.producer_acquire.<locals>.<lambda>W  r—   r%   r0   c                  ój   •— ‰j                   j                  ‰j                  ‰j                  ‰‰ ¬«      S rZ   r¤   r•   s   €€€€r$   r–   z9PipelineTmaCpAsyncUmma.producer_acquire.<locals>.<lambda>_  r°   r%   )r   r	   r¦   r±   s   ``  ``r$   r[   z'PipelineTmaCpAsyncUmma.producer_acquireG  sL   û€ ô 	Ø Ð%Ò?Ð):¸aÑ)?ÞYØØõ		
ô 	Ü�×#Ñ# [Ó1ÞaØØö		
r%   c                ól   — t         j                  j                  | j                  |||¬«      ||¬«       y)zJ
        We need the mbarrier to track the completion of cp.async
        r0   Nr³   r„   s       r$   r¶   z.PipelineTmaCpAsyncUmma.producer_cpasync_commitd  s4   € ô
 	�	‰	×0Ñ0Ø×%Ñ% e°¸Ð%Ó<À#È"ð 	1õ 	
r%   r·   )rM   rN   rO   ri   r
   r&   rž   r   r3   rË   r   ÚLayoutrÌ   r†   r   r   r   r   r[   r¶   rP   r%   r$   rÎ   rÎ   Ú  s&  „ ñð
 Øð )-Ø15Ø)/Ø Ø/3ØØòb
àðb
ð )ðb
ð )ð	b
ð
 ðb
ð Ÿ™ðb
ð " $§+¡+Ñ.ðb
ð ˜S #˜X‘ðb
ð ðb
ð & e™_òb
ó ó ðb
ðH ð 04Ø)-ð	
ð Øò
àð
ð $ GÑ,ð
ð ˜gÑ&ò	
ó ð
ð8 ØCGÈDò 
¨]ò 
ó ñ
r%   rÎ   )/Útypingr   Údataclassesr   Úcutlass.cuter3   Úcutlassr   r   r   Úcutlass.cutlass_dslr   r	   r
   Úcutlass.pipeliner   r   r   r   r   r   r   r   rw   r   r‚   r   rŠ   r   ÚPipelineTmaAsyncOgr   rÞ   r   ÚPipelineUmmaAsyncOgr   ÚPipelineAsyncUmmaOgr'   r.   r;   r=   rž   rU   rW   r   r¬   r¹   rÎ   rP   r%   r$   ú<module>rþ      sY  ðõ Ý !å ß .Ñ .ß >Ñ >ß HÑ Hß <ß .Ý ;Ý =Ý AÝ CÝ AÝ EÝ Eò
ðM�uð M Uð M¨}ó Mò
3ô
˜Mô 
ð"	[Ð.ð 	[¸ó 	[÷%5ñ %5ñV �$Ôô
�>ó 
ó ð
ñ4 ' ~°|ÓD€Ô ñ �$Ôô;
Ð,¨oó ;
ó ð;
ñH �$Ôô
Ð.Ð0Aó 
ó ð
ñF �$Ôô^Ð/Ð1Có ^ó ð^ñ: +Ð+=Ð?OÓPÐ Ô ñ �$Ôô+Ð.Ð0Aó +ó ð+ñ\ *Ð*;¸_ÓM€Ô ñ �$Ôô	Ð0Ð2Eó 	ó ð	ñ ,Ð,?ÐARÓSÐ Ô ñ �$Ôô	Ð0Ð2Eó 	ó ð	ñ ,Ð,?ÐARÓSÐ Ô ñ �$Ôô$
Ð1Ð3Eó $
ó ð$
ñN -Ð-?ÐASÓTÐ Ô ô*
˜mô *
ñ` �$ÔôP
Ð.ó P
ó ñP
r%   