
    ^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II!YY  " 	7$326	7 	7 	dEsr2	7 	7s   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^,,KK%"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    $KKvay)5+;U6!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     ((($VU1XuQxqJJ	!**	*$VU1XuQxqJJZZZur%   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(     E5)e%6CB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'    E5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'     E5)5.cb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'    E5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 4U 4 4 
 -1	
B 
B
B 
B !)	
B 
B <@T 5e 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    nd3F264 

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5f53];3 35EF3 9;QR3 46G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 46G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[[cVX/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)*!!((d6H6HcVX(Y,,55FH!!66u{{HRUZ\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KK!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KKsr 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FH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>PVY^`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 CGD 
] 
 
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      / $ %dg GGLL$??aRSS!STT<<:---$--!2C\]]j,- $ 1 1J >D "11 	sr*r%   c                 2    | j                   | j                  gS r@   )r   r   )r8   s    r$   __extract_mlir_values__z/MbarrierArrayWDropCount.__extract_mlir_values__  s    $$doo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:    &1Ittww(?PVWXPY
 	
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	 +
  + UO + 
 +  +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!.;
 "diiSR&PTU&U M M-IICBM .DDSRM &$))O1#SV[]*^bc*c JJ''++##,,00 	 &II))+&$))OQS*TXY*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 CGD 
] 
 
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 %5V $
> 
 
4 '~|D  $;
,o ;
 ;
H $
.0A 
 
F $^/1C ^ ^: ++=?OP   $+.0A + +\ **;_M  $	02E 	 	 ,,?ARS   $	02E 	 	 ,,?ARS   $$
13E $
 $
N --?AST  *
m *
` $P
. P
 P
r%   