
    uki:#                        d Z ddlZddlmZ ddlmZ ddlmZ	 ddlm
Z
  ej                  d      Zd	e	j                  d
ej                  fdZdej                  d
e	j                  fdZdej                  d
efdZ ej                  d      Zd	e	j(                  d
ej                  fdZdej                  d
e	j(                  fdZdej                  d
efdZ ej                  d      Zd	e	j2                  d
ej                  fdZ ej                  d      Z ej                  d      Z ej                  d      Zdej                  d
e	j2                  fdZdej                  d
efdZd	e	j@                  d
ej                  fdZ!dej                  d
e	j@                  fdZ"de	j                  de	j2                  d
efdZ# ej                  d      Z$dej                  d
efd Z% ej                  d!      Z&dej                  d
efd"Z' ej                  d#      Z(dej                  d
efd$Z)d%e
jT                  ejV                  z  d
ej                  fd&Z,d%ej                  d
e
jT                  ejV                  z  fd'Z-y)(zLayout utilities.    N)mosaic_gpu_dialect)ir   )fragmented_array)launch_contextz2^#mosaic_gpu.WGSplatFragLayout<\[(?P<shape>.*)\]>$layoutreturnc                 n    t         j                  j                  dt        | j                         d      S )zNConstructs a #mosaic_gpu.WGSplatFragLayout attribute from a WGSplatFragLayout.z#mosaic_gpu.WGSplatFragLayout<>)r   	Attributeparselistshaper   s    ^/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/mosaic/gpu/layouts.pyto_splat_fragmented_layout_attrr      s/    			&tFLL'9&:!<
     attrc                     t         j                  t        |             }|st        d|        t	        j
                  t        d |j                  d      j                  d      D                    S )zConstructs a WGSplatFragLayout from a #mosaic_gpu.WGSplatFragLayout attribute.

  Raises:
    ValueError: If the attribute is not a #mosaic_gpu.WGSplatFragLayout
      attribute.
  z8Expected a #mosaic_gpu.WGSplatFragLayout attribute, got c              3   2   K   | ]  }t        |        y wNint.0ss     r   	<genexpr>z4from_splat_fragmented_layout_attr.<locals>.<genexpr>4        BQ#a&B   r   ,)r   )	%_splat_fragmented_layout_attr_pattern	fullmatchstr
ValueErrorfaWGSplatFragLayouttuplegroupsplitr   matchs     r   !from_splat_fragmented_layout_attrr,   &   si     0
9
9#d)
D%	

B4&I  
		B%++g"6"<"<S"ABB
 r   c                 P    t        t        j                  t        |                   S r   )boolr!   searchr#   r   s    r   is_splat_fragmented_layoutr1   8   s    	3::3t9E	FFr   zJ^#mosaic_gpu.WGStridedFragLayout<\[(?P<shape>.*)\], (?P<vector_size>\d+)>$c                     t         j                  j                  dt        | j                         d| j
                   d      S )zRConstructs a #mosaic_gpu.WGStridedFragLayout attribute from a WGStridedFragLayout.z #mosaic_gpu.WGStridedFragLayout<, r   )r   r   r   r   r   vec_sizer   s    r   !to_strided_fragmented_layout_attrr5   A   sB     
		(fll);(< =

//	!
 r   c                    t         j                  t        |             }|st        d|        t	        j
                  t        d |j                  d      j                  d      D              t        |j                  d                  S )zConstructs a WGStridedFragLayout from a #mosaic_gpu.WGStridedFragLayout attribute.

  Raises:
    ValueError: If the attribute is not a #mosaic_gpu.WGStridedFragLayout
      attribute.
  z:Expected a #mosaic_gpu.WGStridedFragLayout attribute, got c              3   2   K   | ]  }t        |        y wr   r   r   s     r   r   z6from_strided_fragmented_layout_attr.<locals>.<genexpr>[   r   r   r   r    vector_size)r   r4   )
'_strided_fragmented_layout_attr_patternr"   r#   r$   r%   WGStridedFragLayoutr'   r(   r)   r   r*   s     r   #from_strided_fragmented_layout_attrr;   K   sz     2
;
;CI
F%	

DTFK  
		B%++g"6"<"<S"ABB5;;}-.
 r   c                 P    t        t        j                  t        |                   S r   )r.   r9   r/   r#   r0   s    r   is_strided_fragmented_layoutr=   `   s    	5<<SYG	HHr   z^#mosaic_gpu.TiledLayout<\[(?P<tiling>.*)\], warp_dims\s*=\s*\[(?P<warp_dims>.*)\], lane_dims\s*=\s*\[(?P<lane_dims>.*)\], vector_dim\s*=\s*(?P<vector_dim>[-\d]+)>$c                    dt         t        j                  z  dt        fdd ddj	                  fd| j
                  j                  D              z   dz   }dd	j	                  fd
| j                  D              z   dz   }dd	j	                  fd| j                  D              z   dz   }t        j                  j                  d| d| d| d| j                   d	      S )zBConstructs a #mosaic_gpu.TiledLayout attribute from a TiledLayout.dr	   c                 l    t        | t        j                        rd| j                   dS t	        |       S )Nz#mosaic_gpu.Replicated<times=r   )
isinstancer%   
Replicatedtimesr#   )r?   s    r   _int_or_replicatedz0to_tiled_layout_attr.<locals>._int_or_replicatedq   s,    !R]]#,QWWIQ77q6Mr   c                 >    ddj                  d | D              z   dz   S )N[r3   c              3   2   K   | ]  }t        |        y wr   )r#   )r   r?   s     r   r   z9to_tiled_layout_attr.<locals>.<lambda>.<locals>.<genexpr>v   s     )?Q#a&)?r   ])join)tiles    r   <lambda>z&to_tiled_layout_attr.<locals>.<lambda>v   s    #		)?$)? ??#E r   rF   r3   c              3   .   K   | ]  } |        y wr    )r   rJ   tile_strs     r   r   z'to_tiled_layout_attr.<locals>.<genexpr>w   s     Jd8D>J   rH   r    c              3   .   K   | ]  } |        y wr   rM   r   r?   rD   s     r   r   z'to_tiled_layout_attr.<locals>.<genexpr>y        Eq'*ErO   c              3   .   K   | ]  } |        y wr   rM   rQ   s     r   r   z'to_tiled_layout_attr.<locals>.<genexpr>|   rR   rO   z#mosaic_gpu.TiledLayout<z, warp_dims=z, lane_dims=z, vector_dim=r   )r   r%   rB   r#   rI   tilingtiles	warp_dims	lane_dimsr   r   r   
vector_dim)r   rT   rV   rW   rD   rN   s       @@r   to_tiled_layout_attrrY   l   s    
C"--/ C 
 F(Jfmm6I6IJJJSP&	CHHEF4D4DEEEK  
CHHEF4D4DEEEK  
		 YK @;mF,=,=+>aA
 r   z\]\s*,\s*\[z^(?P<num>[-\d]+)(\s*:\s*\w+)?$z=^#mosaic_gpu.Replicated<\s*times\s*=\s*(?P<times>\d+)\s*>\s*$c           
         t         j                  t        |             }|st        d|        dt        dt        t
        j                  z  fd|j                  d      }g }t        |      dkD  rt        j                  |dd       }t        d	 |D              }t        j                  t        j                  |      t        fd
|j                  d      j                  d      D              t        fd|j                  d      j                  d      D              t	        |j                  d                  S )zConstructs a TiledLayout from a #mosaic_gpu.TiledLayout attribute.

  Raises:
    ValueError: If the attribute is not a #mosaic_gpu.TiledLayout
      attribute.
  z2Expected a #mosaic_gpu.TiledLayout attribute, got replicated_dimr	   c                    t         j                  |       }|r-t        j                  t	        |j                  d                  S t        j                  |       }|rt	        |j                  d            S t        d|        )NrC   numz%Unexpected format for replicated dim )_replicated_patternr"   r%   rB   r   r(   _int_patternr$   )r[   r+   s     r   rD   z2from_tiled_layout_attr.<locals>._int_or_replicated   sm    )).9E]]3u{{73455"">2EU#$$
<^<LM
NNr   rT      r   c           	   3   l   K   | ],  }t        t        t        |j                  d                    . yw)r    N)r'   mapr   r)   )r   tss     r   r   z)from_tiled_layout_attr.<locals>.<genexpr>   s$     F2c#rxx}-.Fs   24c              3   J   K   | ]  } |j                                 y wr   stripr   r   rD   s     r   r   z)from_tiled_layout_attr.<locals>.<genexpr>   $       QWWY
'    #rV   r    c              3   J   K   | ]  } |j                                 y wr   rf   rh   s     r   r   z)from_tiled_layout_attr.<locals>.<genexpr>   ri   rj   rW   rX   )rT   rV   rW   rX   )_tiled_layout_attr_patternr"   r#   r$   r   r%   rB   r(   len_list_of_lists_delimiterr)   r'   TiledLayoutTiling)r   r+   
tiling_strtile_stringsrU   rD   s        @r   from_tiled_layout_attrrs      s    %
.
.s4y
9%	

<TFC O Or}}1D O {{8$*,_q+11*Qr2BCL
FF
F%	YYu ;;{+11#6   ;;{+11#6  U[[./
 r   c                 P    t        t        j                  t        |                   S r   )r.   rl   r/   r#   r0   s    r   is_tiled_layoutru      s    	(//D	:	;;r   c                     | xt         j                  d x\    t        |       S  xt         j                  d x\    t	        |       S  t         j
                  d x\   t        |       S  	 t        d|        )zBConstructs an MLIR attribute that corresponds to the given layout.rM   z5Unsupported layout for conversion to MLIR attribute: )r%   r&   r   r:   r5   ro   rY   NotImplementedErrorr   s    r   to_layout_attrrx      su    				,V44 
 	!			!.v66 
"		!&)) 
	
A&
J r   c                     t        |       rt        |       S t        |       rt        |       S t	        |       rt        |       S t        d|        )z+Constructs a layout from an MLIR attribute.z7Unsupported layout for conversion from MLIR attribute: )r1   r,   r=   r;   ru   rs   rw   r0   s    r   from_layout_attrrz      sR    %,T22#D).t44t!$''

A$H r   l1l2c                 j    | j                   |j                  }}t        d t        ||      D              S )Nc              3   2   K   | ]  \  }}||z  d k(    yw)r   NrM   )r   d1d2s      r   r   z1splat_is_compatible_with_tiled.<locals>.<genexpr>   s     4fb"R"W\4r   )r   base_tile_shapeallzip)r{   r|   s1s2s       r   splat_is_compatible_with_tiledr      s.     88R''b"	4B4	44r   z^#mosaic_gpu.tile<[^>]+>$c                 P    t        t        j                  t        |                   S r   )r.   _tile_transform_attr_patternr/   r#   r0   s    r   is_tile_transformr      s    	*11#d)<	==r   z^#mosaic_gpu.transpose<[^>]+>$c                 P    t        t        j                  t        |                   S r   )r.   !_transpose_transform_attr_patternr/   r#   r0   s    r   is_transpose_transformr      s    	/66s4yA	BBr   z^#mosaic_gpu.swizzle<[^>]+>$c                 P    t        t        j                  t        |                   S r   )r.   _swizzle_transform_attr_patternr/   r#   r0   s    r   is_swizzle_transformr      s    	-44SY?	@@r   	transformc                    t        | t        j                        r)t        j                  j                  | j                        S t        | t        j                        r)t        j                  j                  | j                        S t        | t        j                        rt        j                  j                  |       S t        d|        NzUnsupported transform )rA   r   TileTransformmgpuTileTransformAttrgetrT   TransposeTransformTransposeTransformAttrpermutationSwizzlingModeSwizzleTransformAttrrw   r   s    r   to_transform_attrr      s     	>778!!%%i&6&677)^>>?&&**9+@+@AA)T//0$$((33
 6ykB
CCr   c                    t        |       r2t        j                  t        j                  |       j
                        S t        |       r2t        j                  t        j                  |       j                        S t        |       r2t        j                  t        j                  |       j                        S t        d|        r   )r   r   r   r   r   rT   r   r   r   r   r   r   r   swizzlerw   r   s    r   from_transform_attrr   	  s     y!''y)00  i(,,##I.::  I&d77	BJJKK
 6ykB
CCr   ).__doc__rejax._src.libr   r   jax._src.lib.mlirr    r   r%   r   compiler!   r&   r   r   r,   r.   r1   r9   r:   r5   r;   r=   rl   ro   rY   rn   r_   r^   rs   ru   FragmentedLayoutrx   rz   r   r   r   r   r   r   r   MemRefTransformr   r   r   rM   r   r   <module>r      s    	 3   $  )3

9) %
B,@,@ R\\ BLL R=Q=Q $GR\\ Gd G +5"**+ '
""\\
,,*Ir|| I I (RZZ2 NN\\2 &2::n5 rzz;< bjjD 
(
,,(^^(V<",, <4 <2.. 2<< 2<< B,?,? 5
5"$..5	5  *rzz   
>BLL >T > %/BJJ%% !
C C$ C #-"**## Ar|| A A
D--0B0BB
D\\
DD||D##d&8&88Dr   