
    uki&              	       t   d Z ddlmZ ddl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 ee	j                  e	j                   f   Zd
edee	j$                     fdZd
edee	j$                     fdZd
edee	j$                     fdZd
edee	j$                     fdZd
edee	j$                     fdZd
edee	j$                     fdZd
edefdZd
edefdZd
edefdZd
edefdZd
edefdZd
edefdZd
edefdZ d
edefdZ!d
edefdZ"d
edefdZ#d
edefdZ$d
edefdZ%d
edefdZ&de'd
ed e(de	j$                  d!z  fd"Z)d
ed#e	jT                  de'de	j$                  d!z  fd$Z+ ee+d%&      Z, ee+d'&      Z-d
e	j                   defd(Z.d
e	j                   defd)Z/d
e	j                   defd*Z0d+e	jT                  defd,Z1d-e	jT                  d.e'de	j$                  d!z  fd/Z2d-e	jT                  de	j$                  d!z  fd0Z3d-e	jT                  de	j$                  d!z  fd1Z4d2ejj                  defd3Z6y!)4z/Layout & transform inference convenience utils.    )Sequence)partial)castUnion)ir   )fragmented_array)tcgen05)utilsopreturnc                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the in_layouts attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an in_layouts attribute.
  
in_layoutsz' does not have an in_layouts attribute.
attributes
ValueErrorr   s    f/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/mosaic/gpu/inference_utils.pyr   r      s2     &
tBC
DD	|	$$    c                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the out_layouts attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an out_layouts attribute.
  out_layoutsz( does not have an out_layouts attribute.r   r   s    r   r   r   (   s2     "--'
tCD
EE	}	%%r   c                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the in_transforms attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an in_transforms attribute.
  in_transformsz* does not have an in_transforms attribute.r   r   s    r   r   r   3   s2     BMM)
tEF
GG		''r   c                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the out_transforms attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an out_transforms attribute.
  out_transformsz+ does not have an out_transforms attribute.r   r   s    r   r   r   >   s3     R]]*
tFG
HH	'	((r   c                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the in_tmem_layouts attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an in_tmem_layouts attribute.
  in_tmem_layoutsz, does not have an in_tmem_layouts attribute.r   r   s    r   r   r   I   s3     bmm+
tGH
II	(	))r   c                 X    d| j                   vrt        |  d      | j                   d   S )zReturns the out_tmem_layouts attribute of the given operation.

  Raises:
    ValueError: If the operation does not have an out_tmem_layouts attribute.
  out_tmem_layoutsz- does not have an out_tmem_layouts attribute.r   r   s    r   r   r   T   s3     r}},
tHI
JJ	)	**r   c                 :    t        d | j                  D              S )zJReturns 'true' if the operation operands should be assigned a TMEM layout.c              3      K   | ]D  }t         j                  j                  |j                        xr t	        j
                  |       F y wNr   
MemRefType
isinstancetyper   is_tmem_ref.0vs     r   	<genexpr>z-should_have_in_tmem_layout.<locals>.<genexpr>a   >      
 	mmqvv&?5+<+<Q+??   A
Aanyoperandsr   s    r   should_have_in_tmem_layoutr1   _   s     	 {{ 
 r   c                 :    t        d | j                  D              S )zIReturns 'true' if the operation results should be assigned a TMEM layout.c              3      K   | ]D  }t         j                  j                  |j                        xr t	        j
                  |       F y wr"   r#   r(   s     r   r+   z.should_have_out_tmem_layout.<locals>.<genexpr>i   r,   r-   r/   resultsr   s    r   should_have_out_tmem_layoutr6   g   s     	 zz 
 r   c                 2    t        |       xs t        |       S )zAReturns 'true' if the operation should be assigned a TMEM layout.)r1   r6   r   s    r   should_have_tmem_layoutr8   o   s    	#B	'	J+Fr+JJr   c                     d| j                   v S )Nr   r   r   s    r   has_in_tmem_layouts_setr;   t   s    	bmm	++r   c                     d| j                   v S )Nr   r:   r   s    r   has_out_tmem_layouts_setr=   x   s    	r}}	,,r   c                 :    t        d | j                  D              S )zEReturns 'true' if the operation operands should be assigned a layout.c              3   n   K   | ]-  }t         j                  j                  |j                         / y wr"   r   
VectorTyper%   r&   r(   s     r   r+   z(should_have_in_layout.<locals>.<genexpr>~   s#     C!R]]%%aff-C   35r.   r   s    r   should_have_in_layoutrC   |   s    	Cr{{C	CCr   c                 :    t        d | j                  D              S )zDReturns 'true' if the operation results should be assigned a layout.c              3   n   K   | ]-  }t         j                  j                  |j                         / y wr"   r@   r(   s     r   r+   z)should_have_out_layout.<locals>.<genexpr>   s#     B!R]]%%aff-BrB   r4   r   s    r   should_have_out_layoutrF      s    	BrzzB	BBr   c                 2    t        |       xs t        |       S )z<Returns 'true' if the operation should be assigned a layout.)rC   rF   r   s    r   should_have_layoutrH      s    	r	"	@&<R&@@r   c                     d| j                   v S )Nr   r:   r   s    r   has_in_layouts_setrJ      s    		&&r   c                     d| j                   v S )Nr   r:   r   s    r   has_out_layouts_setrL      s    	"--	''r   c                 2    t        |       xs t        |       S r"   )rJ   rL   r   s    r   has_any_layout_setrN      s    	B		:#6r#::r   c                     d| j                   v S )Nr   r:   r   s    r   has_in_transforms_setrP      s    	BMM	))r   c                     d| j                   v S )Nr   r:   r   s    r   has_out_transforms_setrR      s    	R]]	**r   	attr_nameindexNc                 h    | |j                   vry|j                   |    }|sy|j                   |    |   S )zReturns `op.attributes[attr_name][index]` if it exists, otherwise None.

  If `op.attributes[attr_name]` exists, then `index` must be a valid index into
  the attribute array.
  Nr:   )rS   r   rT   attrs       r   attr_elementrW      s;     bmm#	y	!$		y	!%	((r   operandc                     |dk(  rd }n|dk(  rt         }nt        d|       | j                  D cg c]  } ||      s| c}j                  |      }t	        || |      S c c}w )Nr   c                 T    t         j                  j                  | j                        S r"   r@   r*   s    r   <lambda>z&_in_attr_for_operand.<locals>.<lambda>       "--22166: r   r   Unknown attribute: )is_transformable_smem_memrefr   r0   rT   rW   )r   rX   rS   	predicateooperand_numbers         r   _in_attr_for_operandrc      sj    
 ,:IO#,I
*9+6
77!{{;!ilA;AA'J.	i^	44 <s   A# A#r   )rS   r   c                 H    t        t        t        | j                              S )zAReturns 'True' if the operation should be assigned in transforms.)r/   mapr_   r0   r   s    r   should_have_in_transformsrf      s    	S-r{{;	<<r   c                 H    t        t        t        | j                              S )zBReturns 'True' if the operation should be assigned out transforms.)r/   re   r_   r5   r   s    r   should_have_out_transformsrh      s    	S-rzz:	;;r   c                 2    t        |       xs t        |       S )zEReturns 'True' if the operation should be assigned in/out transforms.)rf   rh   r   s    r   should_have_transformsrj      s    	"2	&	H*DR*HHr   r*   c                     t         j                  j                  d      }t         j                  j	                  | j
                        xr0 | j
                  j                  |k7  xr t        j                  |       S )zLWhether the value is a memref in SMEM on which transforms should be applied.z!mosaic_gpu.barrier)	r   Typeparser$   r%   r&   element_typer   is_smem_ref)r*   
barrier_tys     r   r_   r_      s\    ww}}23*mmqvv& 
&&


+ 

A
	r   value	attr_typec                    |dk(  rd }n|dk(  rt         }nt        d|       d|z   }d|z   }| j                  }t        |t        j
                        rS||j                  vry |j                  D cg c]  } ||      s| c}j                  |       }|j                  |   |   S t        |t        j                        ry|j                  }t        t        j                  |      }	||j                  vry |	j                  D cg c]  } ||      s| c}j                  |       }
|j                  |   |
   S t        | d      c c}w c c}w )Nlayoutsc                 T    t         j                  j                  | j                        S r"   r@   r[   s    r   r\   z_value_attr.<locals>.<lambda>   r]   r   
transformsr^   in_out_z* is not a function block nor an operation.)r_   r   ownerr%   r   	Operationr   r5   rT   Blockr   	argumentsNotImplementedError)rq   rr   r`   in_attr_typeout_attr_typery   rvalue_result_numberowner_opblockvalue_arg_numbers              r   _value_attrr      sH   ):IL ,I
*9+6
77",9$-
++%r||$E,,,&+mmDy|1DJJ M*+>?? rxx {{H5!E8...#(??CailCII%P|,-=>>	9:	  E Ds   1E?EEEc                     t         j                  j                  | j                        st	        |  d      t        | d      S )zxReturns the layout for a given value as defined by its owner.

  Raises:
    ValueError: If `result` is not a Vector.
  z is not a vector.rt   )r   rA   r%   r&   r   r   rq   s    r   value_layoutr     s:     
	!	!%**	-
w/0
11	UI	&&r   c                     t         j                  j                  | j                        st	        |  d      t        | d      S )z|Returns the transforms for a given value as defined by its owner.

  Raises:
    ValueError: If `result` is not a memref.
  z is not a memref.rv   )r   r$   r%   r&   r   r   r   s    r   value_transformsr     s:     
	!	!%**	-
w/0
11	UL	))r   layoutc           	         t        | t        j                        sy| t        j                  t        j                  t        j
                  t        j                  t        j                  t        j                  t        j                  t        j                  hv ryt        | j                  j                  d         dk7  ry| j                  j                  d   d   }|dz  dk(  xr | t        j                  |      k(  S )NFTr      r      )r%   faTiledLayoutWGMMA_LAYOUTWGMMA_LAYOUT_ACC_32BITWGMMA_LAYOUT_UPCAST_2XWGMMA_LAYOUT_UPCAST_4XWGMMA_TRANSPOSED_LAYOUTWGMMA_LAYOUT_8BITTCGEN05_LAYOUTTCGEN05_TRANSPOSED_LAYOUTlentilingtilesr
   fa_m64_collective_layout)r   columnss     r   is_mma_layoutr     s    	FBNN	+oo  ""	 	 		Q	 A%MM"1%'	2	 
0099r   )7__doc__collections.abcr   	functoolsr   typingr   r   jax._src.lib.mlirr    r	   r   r
   r   rz   OpViewMlirOperation	Attributer   r   r   r   r   r   boolr1   r6   r8   r;   r=   rC   rF   rH   rJ   rL   rN   rP   rR   strintrW   Valuerc   in_layout_for_operandin_transforms_for_operandrf   rh   rj   r_   r   r   r   FragmentedLayoutr    r   r   <module>r      sI   6 $     $  bllBII-.%= %Xbll%; %&M &hr||&< &(m ((> ()} )",,)? )* *(2<<*@ *+ +8BLL+A += T M d K K$ K
, ,$ ,- -4 -Dm D D
C} C C
A= AT A
'= 'T '(M (d (;= ;T ;*m * *+} + +))%).1)\\D) 55XX5 5 \\D	5"  L  $O 
=")) = =
<299 < <
Iryy IT I
BHH   rxx  C  BLL44G  F	' 	'R\\D%8 	'	*BHH 	*)< 	*"-- $ r   