
    Vpf                      d Z ddlmZ ddlmZmZ ddlZddlZddlZddl	Z	ddl
mZm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 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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dl*m+Z+ ddl,m-Z- ddl,m.Z. ddl,m/Z/ ddl,m0Z0 ddl,m1Z1 ddl,m2Z2 ddl3m4Z4 dd	l3mZ5 ddl3m6Z6 dd l3m7Z8 dd	l9mZ: ddl9m6Z; dd!l<m=Z> dd"l<m?Z? ddl<m6Z@ dd#lAmBZB dd$lAmCZC dd%lAmDZD dd&lAmEZE dd'lFmGZG ddlHmIZJ dd(lKmLZL ddlIZMe?jN        ZNe:jO        ZOe5jP        eOz  ZPe:jO        jQ        ZQe:jO        jR        ZR eMjS        d)          ZTd*ZUejV        ZVeBeWcZWZXeCeYcZYZZ eMjS        d+           eMjS        d,           eMjS        d-           eMjS        d.           eMjS        d/           eMjS        d)           eMjS        d0           eMjS        d1          iZ[ej\         G d2 d3                      Z]ej\         G d4 d5                      Z^ej\         G d6 d7                      Z_dd<Z`	 dddCZa	 	 	 dddDZbddEZci Zd ee            ZfdF ZgddIZh ej\        d=J           G dK dL                      Ziej\         G dM dN                      Zj	 	 ddd]ZkdddZlddeZmddiZn G dj dkeo          ZpddpZqddtZrdu Zsdv Ztd dwZueuede@jv        <   efw                    e@jv                   d dxZxexede@jy        <   efw                    e@jy                   dy Zzdz Z{d!dZ|d"dZ}d#dZ~d Z ej\        d           G d d                      Zd dZd$dZeede6j        <   efw                    e6j                   d%dZd&dZd dZeede6jy        <   efw                    e6jy                   d dZeede6j        <   d ZeJj        e2j        j        eJj        e2j        j        eJj        e2j        j        iZeJj         ed          eJj         eMj        eMj                  j        iZ eeJj        ee          Zeedej        <   eJj        e2j        j        eJj        e2j        j        eJj        e2j        j        iZeJj         ed          eJj         eMj        eMj                  j        iZ eeJj        ee          Zeedej        <   eJj        e2j        j        eJj        e2j        j        eJj        e2j        j        iZeJj        deJj        diZ eeJj        ee          Zeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d Zd dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d Zd dZeedej        <   efw                    ej                   eedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   d dZeedej        <   efw                    ej                   d ZАd dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   efw                    ej                   d dZeedej        <   d dZeedej        <   efw                    ej                   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   d dZeedej        <   ej        dej        dej        dej        dej        dej        diZej        dej        dej        dej        dej        dej        diZd dńZ ejV        eej                  edej        <    ejV        eej                  edej        <    ejV        eej                  edej        <    ejV        eej                  edej        <    ejV        eej                  edej        <    ejV        eej                  edej        <   d dƄZeedej        <   efw                    ej                   d dǄZeedej        <   efw                    ej                   d dȄZeedej        <   d dɄZeedej        <   dʄ Zd d˄Zeedej        <   d d̄Zeede&j         <   d'dӄZd(dԄZd)d݄Zeedej        <   efw                    ej                   d dބZd d߄Zeedej        <   d dZeedej	        <   d dZ
e
edej        <   d*dZeedej        <   d dZeedej        <   d+dZeede6j        <   d+dZeede6j        <   d dZeede;j        <   d dZeede;j        <   d dZeedej        <   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZeedej        <   efw                    ej                   d dZ e edej!        <   efw                    ej!                   d Z"d dZ#e#edej$        <   d dZ%e%ede;j&        <   d dZ'e'edej(        <   d,dZ)d dZ*e*ede6j+        <   d-dZ,d dZ-e-ede;j.        <   d-dZ/e/ede;j0        <   d dZ1e1ede;j2        <   d-dZ3e3ede;j4        <   d-dZ5e5ede;j6        <   d dZ7e7ede;j8        <   d.dZ9e9edej:        <   d dZ;e;ede;j<        <   d/dZ=e=ede;j>        <   d0dZ?e?ede6j@        <   d d	ZAeAede;jB        <   d d
ZCeCede;jD        <   d ZEeEedejF        <   d ZGeGedejH        <   d ZIeIedejJ        <   d ZKeKedejL        <   d ZMeMedejN        <   d ZOddlPmQZQ   e>jR        eQjS                  eO           dS (1  z;Module for lowering JAX to Mosaic-compatible MLIR dialects.    )annotations)CallableSequenceN)AnyHashable)lax)	tree_util)ad_util)core)custom_derivatives)	debugging)dtypes)linear_util)mesh)pjit)prng)source_info_util)state)mlir)partial_eval)for_loop)version)ir)arith)func)math)memref)scf)vector)pallas_call)
primitives)utils)	discharge)indexing)safe_map)safe_zip)
split_list)unzip2)tpu)Moduleint32l         uint8int8uint16int16uint32uint64int64c                  .    e Zd ZU ded<   ded<   ded<   dS )MeshContexttuple[int, ...]
mesh_shapeztuple[str, ...]
axis_namesmesh_stridesN__name__
__module____qualname____annotations__     _/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/mosaic/lowering.pyr4   r4   ^   s6         r?   r4   c                      e Zd ZU ded<   ded<   ded<   ded<   ded	<   d
ed<   ded<   ded<   ej        Zded<   ded<   ed             Zej	        d             Z
dS )LoweringContext
ir.Context
ir_contextr5   
grid_sizestuple[Hashable, ...] | None
grid_namesmapped_dimszSequence[ir.Value] | Noneuser_grid_indicesz&list[tuple[int | pl_core.Mapped, ...]]block_shapeszsource_info_util.NameStack
name_stackzMeshContext | Nonemesh_contextzmlir.TracebackCachestraceback_cachesboolfor_verificationc                *    t          | j                  S N)lenrE   selfs    r@   	grid_rankzLoweringContext.grid_ranks   s    tr?   c              #     K    j         sd V  d S  j         }t           fdt           j                  D                       }t	          ||          }t          j        |          5  d V  d d d            d S # 1 swxY w Y   d S )Nc              3  4   K   | ]\  }}|j         v|V  d S rQ   rH   ).0idrT   s      r@   	<genexpr>z4LoweringContext.grid_name_context.<locals>.<genexpr>~   s=        aAT=M4M4M4M4M4M4M r?   )rG   tuple	enumeraterE   zipjax_coreextend_axis_env_nd)rT   rG   valid_grid_sizesgrid_envs   `   r@   grid_name_contextz!LoweringContext.grid_name_contextw   s       ? eeefJ    00     :/00H		$X	.	.  eee                 s   )A;;A?A?N)r:   r;   r<   r=   dataclassesreplacepropertyrU   
contextlibcontextmanagerrd   r>   r?   r@   rB   rB   e   s         ))))....6666((((""""'((((    8      r?   rB   c                  F    e Zd ZU ded<   ded<   ded<   ded<   ej        ZdS )	LoweringRuleContextrB   lowering_contextz Sequence[jax_core.AbstractValue]avals_in	avals_outz1Sequence[tuple[int | pl_core.Mapped, ...] | None]rJ   N)r:   r;   r<   r=   re   rf   r>   r?   r@   rk   rk      sH         ####,,,,----AAAA'''r?   rk   memory_spaceMemorySpace | Nonereturnir.Attributec                    | t           } n9| t          j        j        k    rt          } n| t          j        j        k    rt          } t          j                            d|  d          S )N#tpu.memory_space<>)	VMEMpl_coreMemorySpaceERRORSMEMINDEXr   	Attributeparsero   s    r@   _memory_space_to_tpu_memspacer      s_    LLw*000LLw*000L			@@@@	A	AAr?   Fdtype	jnp.dtypeis_kernel_boundaryrN   ir.Typec                   t          j        | t          j                  rt          j        | t          j                  rt
          j                            d          S t          j        | t          j                  rt
          j                            d          S t          j        | t          j	                  rt
          j                            d          S t          |r.t          j        | t          j        d                    rt          } t          j        |           }t          |t
          j                  r$t
          j                            |j                  S |S )N!tpu.dma_semaphore!tpu.semaphorerN   )jnp
issubdtypetpu_coresemaphore_dtypedma_semaphorer   Typer}   	semaphorebarrier_semaphoreNotImplementedErrorr   BOOL_MEMREF_TYPEr   dtype_to_ir_type
isinstanceIntegerTypeget_signlesswidth)r   r   types      r@   _dtype_to_ir_typer      s   ^E8344  
~eX344  W]]/000	x1	2	2  W]]+,,,	x9	:	:  W]]+,,, CN5#)F2C2CDD E 

&
&$bn%% >&&tz222Kr?   c                v   t          | t          j                  r| j        t          j        j        u r t          j                            d          }n| j        t          j        j	        u r t          j                            d          }nP| j        t          j        j
        u r t          j                            d          }nt          d| j         d          t          t          j                  }t          j                            d||          S t#          j        | j        t"          j                  r| j        j        j        }|t          j        }|t          j        k    rt          d|           t          |          }t          j                            |t1          t3          j        t2          j                            |          S t          | t6          j                  rN|| j        }t          |          }t          j                            |t1          | j        d	          |          S t          | t<          j                  rU|| j        }|st1          | j        |	          S t          j                             |t1          | j        |	                    S tC          |           )
Nr   r   Cannot allocate .r>   r~   z&PRNG keys must be stored in SMEM. Got Tr   )"r   r   AbstractSemaphoresem_typeSemaphoreTypeDMAr   r   r}   REGULARBARRIER
ValueErrorr   TPUMemorySpace	SEMAPHORE
MemRefTypegetr   r   r   prng_key_impl	key_shaperz   r   npr0   r   AbstractRefshaper`   ShapedArray
VectorTyper   )avalr   ro   r   r   memspaces         r@   aval_to_ir_typer      sq    h011 
B}.222344hh	(08	8	8/00hh	(08	8	8/00hh:$-:::;;;,^-EFFH=RAAAtz6?33 4J&E#(l~***NNNOOO,\::H=U$5bhry6I6I$J$J*2  4 4 4e'(( }je,\::H=U
t<<<     h*++ N}je =
*);= = = ==$*9KLLLN N N 	D!!!r?   c           	        t          | d          sit          | t                    r t          j        | t          j                  } n4t          | t                    rt          j        | t          j                  } |st          | j	                  }t          | t                    s*| j	        t          j        t          j
        t          j        fv rEt          j        |t          j                            |t          |                               j        S t          | t                    s| j	        t          j        k    rEt          j        |t          j                            |t          |                               j        S | j	        t&          j        k    rEt          j        |t          j                            |t          |                               j        S | j	        t&          j        k    rDt          j        |t          j                            t/          |                               j        S t1          | j	                  )Nr   )hasattrr   intr   arrayr+   floatfloat32r   r   r0   r-   r   
ConstantOpr   IntegerAttrr   result	FloatAttrr   bfloat16bool_BoolAttrrN   r   )x	mlir_types     r@   ir_constantr      s   	G		 "!S "
(1bh

aa	Au		 "
(1bj
!
!a	 +!!'**I3 17rxBG&DDDIr~'9'9)SVV'L'L  $%!U qw"*442<##IuQxx88  	w#,2<##IuQxx88  	w#)2;??477++  	AG$$$r?   c                l    t          j        |           j        t          | j                  d          }|S rQ   )r`   physical_avalr   rR   )r   dtype_physical_shapes     r@   _get_aval_physical_dtype_shaper      s4    !/55;	$*oo 
r?   block_mappingpl_core.BlockMapping | Nonec                \   d }t          | t          j                  r| j        }|t          j        }t          | t          j                  rt          |           d fS |t          | |          | j	        fS t          d |j        D                       }t          | ||          |j        fS )Nr~   c              3  :   K   | ]}|t           j        u rd n|V  dS    Nrw   mappedrY   bs     r@   r\   z _get_arg_type.<locals>.<genexpr>  s2      SSAQ'.((aSSSSSSr?   )r   ro   )r   rw   AbstractMemoryRefro   r   rv   r   r   r   r   r]   block_shape)r   r   ro   r   s       r@   _get_arg_typer      s     ,g/00 )$L#(lh011 '4  $&&4l;;;TZGG
SS9RSSS
S
S%d%lCCC
 r?   )initc                      e Zd ZU ded<   ded<   ded<   ded<   d	ed
<   ded<   ded<   ded<   ded<   ded<   ded<   ded<   ded<   ded<   d)dZd*dZd  Zej        d+d#            Z	d,d%Z
d-d'Zd(S ).MosaicGridMappingztuple[int, ...] | NonegridrF   rG   jax_core.Jaxprjaxprz'tuple[pl_core.BlockMapping | None, ...]block_mappingsr5   rH   ztuple[ir.Type, ...]scalar_prefetch_typesoperand_typesscratch_types
grid_typesztuple[tuple[int, ...], ...]scalar_prefetch_block_shapesoperand_block_shapesscratch_block_shapeszMeshInfo | None	mesh_infozCallable | Noneget_grid_indicesgrid_mappingpl_core.GridMappingdimension_semanticstuple[str, ...] | Noner   mesh_lib.Mesh | Nonec                    |j          _         |j         _        | _        |j         _        |j         _        |j        }|j        }t           j        j	                  |z
  |z
  }t           fdt           j                   D                       }|dt          |          z  }t          |          t          |          k    rt          d          |t           j                  k    rt          d          t           j                  t          |          z   t           j                   k    sJ d j        d|d j                     t          |          t           fdt          t           j                             D                        _        d	  j        j	        D             }	t!          |	||g          \  }
}}t#          d
 |
D                       \   _        }t          d |
D                        _        t#          d t)          | j                  D                       \   _         _        t#          d |D                       \   _        }t          d |D                        _        t#          d t          t           j                             D                       \   _        }                     |           d }| _        d S )Nc              3  4   K   | ]\  }}|j         v|V  d S rQ   rX   )rY   rZ   grT   s      r@   r\   z-MosaicGridMapping.__init__.<locals>.<genexpr>5  s=        aat7G.G.G.G.G.G.G r?   )	arbitraryz=Must have dimension semantics for each dimension of the grid.z*Must have block mappings for each operand.z%Misconfigured grid: self.mapped_dims=z, dimension_semantics=z, self.grid=c              3  L   K   | ]}|j         vrt                    nd V  dS )parallelN)rH   next)rY   rZ   rT   semantics_iters     r@   r\   z-MosaicGridMapping.__init__.<locals>.<genexpr>I  sQ       & & !")9 9 9^z& & & & & &r?   c                    g | ]	}|j         
S r>   r   )rY   invars     r@   
<listcomp>z.MosaicGridMapping.__init__.<locals>.<listcomp>N  s    :::u
:::r?   c                .    g | ]}t          |d           S rQ   r   rY   r   s     r@   r   z.MosaicGridMapping.__init__.<locals>.<listcomp>R  s2     ,+ ,+ ,+ 	dD!!,+ ,+ ,+r?   c              3  $   K   | ]}|j         V  d S rQ   r   r   s     r@   r\   z-MosaicGridMapping.__init__.<locals>.<genexpr>U  s5       .6 .6
.6 .6 .6 .6 .6 .6r?   c                4    g | ]\  }}t          ||          S r>   r   )rY   r   r   s      r@   r   z.MosaicGridMapping.__init__.<locals>.<listcomp>W  s<     <L <L <LD- 	dM**<L <L <Lr?   c                .    g | ]}t          |d           S rQ   r   r   s     r@   r   z.MosaicGridMapping.__init__.<locals>.<listcomp>Z  s/     $= $= $=&*dD!!$= $= $=r?   c              3  \   K   | ]'}t          |t          j                  s|j        nd V  (d S rQ   )r   r   r   r   r   s     r@   r\   z-MosaicGridMapping.__init__.<locals>.<genexpr>\  sN       & & %T8+EFFP

D& & & & & &r?   c                B    g | ]}t          t          j        d           S rQ   )r   rw   index_map_grid_aval)rY   _s     r@   r   z.MosaicGridMapping.__init__.<locals>.<listcomp>`  s5     ! ! ! 	g1488! ! !r?   c                    | S rQ   r>   )indicess    r@   _get_grid_indicesz5MosaicGridMapping.__init__.<locals>._get_grid_indicese  s    nr?   )r   rG   r   r   vmapped_dimsrH   num_index_operandsnum_scratch_operandsrR   invarsr]   r^   r   iterrange_dimension_semanticsr'   r(   r   r   r_   r   r   r   r   r   _prepare_mesh_infor   )rT   r   r   r   r   num_scalar_prefetchnum_scratchnum_operands	user_gridin_avalsscalar_prefetch_avalsoperand_avalsscratch_avalsr   r  r   s   `              @r@   __init__zMosaicGridMapping.__init__$  s    !DI"-DODJ&5D#0D&93K 	DJ
	
	 
     	**    I "*S^^;
9~~01111
I   s4.////CDDDt  3':#;#;;s	@ @   	t/ 	 	4G 	 	I	 	   -..N % & & & & &s49~~&&& & & ! !D
 ;:
(9:::H:D&5; ;7=- %+ ,+ ,+),+ ,+ ,+ %, %,!D ). .6 .64.6 .6 .6 )6 )6D%4: <L <L#&}d6I#J#J<L <L <L 5M 5M1D1 # $= $=.;$= $= $= > >D % & &!& & & ! !D   ! !s49~~&&! ! !  DOQ 	D!!!  -Dr?   c                v     j         s	d  _        d S t          d          j        } j        :t           fd|D                       rt          dj         d j                   t          j        t          fd|D                                 }t          j
        j        ||           _        d S )Nz:Cannot use communication in pallas_call without shard_map.c              3  *   K   | ]}|j         v V  d S rQ   )rG   )rY   arT   s     r@   r\   z7MosaicGridMapping._prepare_mesh_info.<locals>.<genexpr>s  s*      66aQ$/!666666r?   zECannot shadow axis mesh axis names with grid names. mesh axis names: z, grid names: c              3  2   K   | ]}j         |         V  d S rQ   r   )rY   r  r   s     r@   r\   z7MosaicGridMapping._prepare_mesh_info.<locals>.<genexpr>z  s:       9 9
19 9 9 9 9 9r?   )has_communicationr   r   r7   rG   anypallas_utilsstrides_from_shaper]   MeshInfo
device_idsr   )rT   r   r7   r8   s   ``  r@   r	  z$MosaicGridMapping._prepare_mesh_infoi  s	   ! dnf|
F   J"	6666:666	6	6 
HH H6:oH H
 
 	
  25 9 9 9 9)9 9 9 4 4  L do3ZNNDNNNr?   c                    d S rQ   r>   rS   s    r@   maybe_compress_gridz%MosaicGridMapping.maybe_compress_grid  s	     	Dr?   rq   rN   c                     t                      }d fd}|                     | j                              j        D ]'}|#|                     ||j                             (t          |          S )Nr   r   c                *    fd| j         D             S )Nc                |    h | ]8}t          |t          j                  rj        r|j        j        v1|j        9S r>   )r   r`   NamedAxisEffectrG   name)rY   erT   s     r@   	<setcomp>zXMosaicGridMapping.has_communication.<locals>._get_nonlocal_axis_names.<locals>.<setcomp>  sZ       8344 ?	 '(fDO&C&C & 'D&C&Cr?   )effects)r   rT   s    r@   _get_nonlocal_axis_nameszEMosaicGridMapping.has_communication.<locals>._get_nonlocal_axis_names  s0       =   r?   )r   r   )setupdater   r   index_map_jaxprrN   )rT   nonlocal_axis_namesr'  bms   `   r@   r  z#MosaicGridMapping.has_communication  s    %%      77
CCDDD! Q Q	""#;#;B<N#O#OPPP#$$$r?   tuple[Any, ...]c                    dS Nr>   r>   rS   s    r@   get_extra_argsz MosaicGridMapping.get_extra_args  s    2r?   ir.ArrayAttrc           	         dd}t           j                            t          t           j        j        t          || j                                      S )Ns
str | Nonerq   strc                    | dS d|  dS )Nz##tpu.dimension_semantics<arbitrary>z#tpu.dimension_semantics<ru   r>   r3  s    r@   _get_semanticszAMosaicGridMapping.get_dimension_semantics.<locals>._get_semantics  s    	
44-----r?   )r3  r4  rq   r5  )r   	ArrayAttrr   mapr|   r}   r  )rT   r8  s     r@   get_dimension_semanticsz)MosaicGridMapping.get_dimension_semantics  sW    . . . .
 <L 9::	
 	
  r?   N)r   r   r   r   r   r   r   r   )r   r   )rq   rN   )rq   r-  )rq   r1  )r:   r;   r<   r=   r  r	  r  	functoolscached_propertyr  r0  r;  r>   r?   r@   r   r     sJ        ))))9999,,,,$$$$$$$$!!!!;;;;33333333####C. C. C. C.JO O O O,	 	 	 % % % %        r?   r   c                  .    e Zd ZU ded<   ded<   ded<   dS )r  r5   r6   	list[str]r7   r8   Nr9   r>   r?   r@   r  r    s6         r?   r  rl   mlir.LoweringRuleContextctxrC   r   r   r   r   r   tuple[str | None, ...] | Noner   r   rO   tuple[Module, tuple[Any, ...]]c           	     6   |j         D ]vfd}t          j                  dk     r5j        j        t
          j        j        k    r                                s	 Uj        j        t
          j        j
        k    r.                                st          d |            z             d j        dd          D             \  }}	j        j        dd          \  }
}|                                 st          dk     rH|	d	z  d
k    s|	|k    o|d	k     o|dz  d
k    p||
k    o|
dk     }|st          d |            z             <|	|k    s	|	d	z  d
k    o||
k    p|dz  d
k    }|st          d |            z             xt          ||||          }|                                 t"          j                                        }t#          j        |j                  }t-          |||d|          }|j                            |           |                    |           g }|j        }|rt7          |j                   D ]\  }d| }j        j        t
          j        j
        k    r2|                    t"          j                                                   \t=          |j        j         j        |||          }|!                                s
J |            d j        D             }|tE          tG          j        j$                            z  }t"          j%                            |          }tM          |t"          j'                            |                    }tQ          j)        tT          j+                  ryj)        j,        d
gt          j                  z  x}}n)t[          tD          t]          j)        j,                   \  }}t"          j/        0                    d| d| d          |d<   |                    t"          j                            |                     |j                            |           |                    |            t"          j1                            |          |j2        d<   d |D             }t"          j%                            |          |j2        d<   t"          j3                            t"          j4        5                    d          t          |j6                            |j2        d<   t"          j3                            t"          j4        5                    d          t          |j7                            |j2        d<   |8                                |j2        d<   ||9                                fS )Nc                     d j          d j         d j        j         d j        j        j         d j        j         dS )NzBlock spec for z has block shape z, array shape z, and index_map returning z, in memory space r   )	originr   array_shape_dtyper   r*  r   outvars
block_avalro   )r,  s   r@   err_detailsz*lower_jaxpr_to_module.<locals>.err_details  sn    <	 < << </1/C/I< < *,);)A)I< < !m8	< < < =r?      FzThe Pallas TPU lowering currently supports only blocks of rank >= 2 for blocks, except those in the SMEM memory space having the same block shape as the array shape and a trivial index_map (returning all 0s).zThe Pallas TPU lowering currently supports in memory space ANY only blocks having the same block shape as the array shape and a trivial index_map (returning all 0s).c              3  :   K   | ]}|t           j        u rd n|V  dS r   r   )rY   bss     r@   r\   z(lower_jaxpr_to_module.<locals>.<genexpr>  sF       / / 7>))r / / / / / /r?   )r            r      a0  The Pallas TPU lowering currently requires that the last two dimensions of your block shape are divisible by 8 and 128 respectively, if the respective dimensions of the overall array are larger than the respective factors. If array dimensions are smaller, the block should span the full array dimension. zThe Pallas TPU lowering currently requires that the last two dimensions of your block shape are divisible by 8 and 128 respectively, or be equal to the respective dimensions of the overall array. main)mosaic_grid_mappingr#  rO   
transform_)r#  rT  rO   c                2    g | ]}|t           j        u rd n|S r   r   r   s     r@   r   z)lower_jaxpr_to_module.<locals>.<listcomp>  s5       ./qGN""!!  r?   )window_boundstransform_indicesz#tpu.element_window<,ru   window_kindwindow_paramsc                <    g | ]}|t           j        u rt          n|S r>   )rw   dynamic_grid_dimMLIR_DYNAMICr   s     r@   r   z)lower_jaxpr_to_module.<locals>.<listcomp>'  s6       ABW5551  r?   iteration_bounds@   scalar_prefetchscratch_operandsr   ):r   rR   r   rI  ro   r   r   rz   has_trivial_windowr   ANYrG  r   is_forward_compatjaxlib_versionr   r  r   r*   createSymbolTable	operationlower_jaxpr_to_funcbodyappendinsertr   r^   DictAttrr   lower_jaxpr_to_transform_funcr*  r   verifylistr   
inner_avalDenseI64ArrayAttrdictFlatSymbolRefAttrr   indexing_moderw   	Unblockedpaddingr:  r_   r|   r}   r9  
attributesr   r   r   r   r   r;  r0  )rl   rA  r   r   r   r   rO   rJ  bs0bs1as0as1evenly_divisiblerT  msym_tabfunc_opr\  r   rZ   	func_name	mlir_funcr   window_shapeblock_paramspad_lowpad_highstatic_gridr,  s                               @r@   lower_jaxpr_to_moduler    s     ' ; ;b= = = = = 2>Qm(H,C,HHH##%% I	G 
"h&=&AAA!!## 	B8:E+--HI I I
/ /.-/ / /HC#)"##.HC))++ ~
/J/J 9>7cSj6S3Y 57a<3C3J237   H
 kmm  	 #:'sa '#:%qA     	  	 *\.6 6))+++i!N1;'''	5&9$4  ' &--	..-		!$	 +S<677 %  % 2"q""i		#x'>'B	B	BR[__..////



"
-1+  i ***** 35>  k
 T89QRRSSSk)--k::l$044Y??  l 
B$g&7	8	8 
#+ !sS%8%88
8'HH!$R-=-E(FGG
'8&(l&8&88788X888'
 '
]# 2;??<88999fmmInnY*,,*:*:=*I*IG' FJ  K .0-A-E-Ek-R-RG)**,.*<*<n!!"%%s+>+T'U'U+W +W'&'+->+=+=n!!"%%s+>+L'M'M,O ,O''( 1133 
*+ 
..00	00r?   r   jax_core.AbstractValuer#  r5  rT  func.FuncOpc               P   
 t          j                  
g j        j        } 
fd}||_         t	          j        j        |d|i|          }	 |j                                         n+# t          $ r}	t          d|j         d          |	d }	~	ww xY w|j        S )Nc                    t          | g          \  }}                    |          }g dgt          |          z  j        }j        }|!t          |j        |j        |j                  }nd }t          
j
        j        j        d |t          j                    |t          j                    
  
        }t#          |g||R  }t%          	t&          j                  s
J 	            |t+          dt-          t/          j        d                              gt          t3          	j                            z  z  }|S )Nr>   rL   rM   rO   r   r+   r   )r'   r   rR   r   r   r4   r6   r7   r8   rB   r   rG   rH   r   	NameStackr   TracebackCachesjaxpr_subcompr   r   r   r   r   r   r   r   rs  )argsgrid_indicesrb  jaxpr_indicesarg_block_shapesr   rL   rl   outr   rA  rO   r   rT  num_grids            r@   	body_funcz0lower_jaxpr_to_transform_func.<locals>.body_funcD  s}   $.thZ$@$@!L/'88FFM
M""	"		9
 $-I 

	 4i6L ll l& &'"$$!-//)   (% *- *(* * *CdE-..44444 A!239W3E3E!F!FGGG*4?;;<<= =C Jr?   r#  Body failed to verify: q.
This is an internal error. Please report a bug at: https://github.com/google/jax/issues/new?assignees=sharadmv.)rR   r   r   r:   r   FuncOpfrom_py_funcr  rq  	ExceptionLoweringException)rA  r   r   r#  rT  rO   	arg_typesr  rl  r$  r  s   ``` ``    @r@   rp  rp  6  s    $/00(%0)# # # # # # # # # #J )	8	!9	84	8	8	C	C$L	   
	H$, 	H 	H 	H  	 
s   A6 6
B BBc                  	
 t          j                  
t          j                  g j        j        j        j        }g j        j        j        		 
fd}||_         t          j
        j        |d|i|          }	 |j                                         n+# t          $ r}t          d|j         d          |d }~ww xY w|j        S )Nc                    t          | g          \  }}}                    |          }t          fdt          |          D                       }j        }|!t          |j        |j        |j                  }nd }t          	j
        j        j        |t          j                    |t          j                    

  
        }t#          |g||R  S )Nc              3  4   K   | ]\  }}|j         v|V  d S rQ   rX   )rY   rZ   idxrT  s      r@   r\   z9lower_jaxpr_to_func.<locals>.body_func.<locals>.<genexpr>  sE       G G&!S&9&EEE EEEEG Gr?   r  )r'   r   r]   r^   r   r4   r6   r7   r8   rB   r   rG   rH   r   r  r   r  r  )r  r  rb  operands_and_scratchr  r   rL   rl   r  rA  rO   r   rT  r  r
  s           r@   r  z&lower_jaxpr_to_func.<locals>.body_func  s/   :Dx,-;/ ;/7L/#7&77EEL G G G GIl,C,C G G G G GM#-I 

	 4i6L ll l& &'"$$!-//)   %"14H   r?   r#  r  r  )rR   r   r   r   r   r   r   r   r:   r   r  r  r  rq  r  r  )rA  r   rT  r#  rO   r  r  rl  r$  r  r  r
  s   ``` `    @@@r@   rk  rk  v  sj    $/00(/EFF%0 ( (	)7/ /
          8 )	8	!9	84	8	8	C	C$L	   
	H$, 	H 	H 	H  	 
s   B1 1
C;CCfunr   multiple_resultsc                    d fd}|S )NrA  rk   c                2   r
n
fd}t          j        ||          }t          j        || j                  \  }}}\   |rt
          t          j        |          }| j                            | j	                  }t          ||g||R  }	s|	d         S |	S )Nc                      | i |fS rQ   r>   )r  kwr  s     r@   <lambda>z.lower_fun.<locals>.f_lowered.<locals>.<lambda>  s    ##t:Jr:J:J9L r?   rJ   r   )lu	wrap_initpetrace_to_jaxpr_dynamicrm   r   convert_constvars_jaxprrl   rf   rJ   r  )rA  r  paramsfwrapped_funr   r   constsrl   r  r  r  s             r@   	f_loweredzlower_fun.<locals>.f_lowered  s    L%L%L%L%LA,q&))K4[#,OOE1fb  &u--E+33% 4 ' '
(%
@&
@4
@
@
@C VmJr?   rA  rk   r>   )r  r  r  s   `` r@   	lower_funr    s/           
r?   c                      e Zd ZdS )r  N)r:   r;   r<   r>   r?   r@   r  r    s        $r?   r  old_name_stackr?  new_name_stacktuple[list[str], list[str]]c                    d}t          t          | |                    D ]\  }\  }}||k    r|dz   } | |d         ||d         fS )ab  Computes the popped/pushed items to the name stack after an update.

  Args:
    old_name_stack: The name stack prior to the update.
    new_name_stack: The name stack after the update.

  Returns:
    popped: A list of names popped from the name stack as part of the update.
    pushed: A list of names pushed to the name stack as part of the update.
  r   r   N)r^   
unsafe_zip)r  r  common_prefix_idxrZ   oldnews         r@   _compute_name_stack_updatesr    so      NN!K!KLL  ma#s
czzA#	)**	+^<M<N<N-O	OOr?   r  ir.ValueSequence[ir.Value]c                   |j         rJ i i dfd}dfd}dfd}t          |j        | j                  D ]
\  }}||<   t	          ||j        |           d | j        j        D             }g }	|	                    |           |j        D ]j}
t	          ||
j                  }|
j	        
                    | j        |
j	        j        z   	          }t          j        | |
j        |
j        |          }t          j        |
j	        j                  5  |5  |
j        t$          v rN|
j        t&          vrd
 t          ||
j                  D             }t	          ||
j                  }t)          | d |
j        D             d |
j        D             |          }d |j        j        D             }t-          |	|          \  }}|}	|D ]}t/          j                     |D ]}t/          j        |d           	 t%          |
j                 |g|R i |
j        }ny# t4          $ r  t6          $ rE}t5          d|
 d| dt	          d |           dt	          d |           d| d|           |d }~ww xY wt9          d|
j        j         d          |
j        j        rt	          ||
j        |           n ||
j        d         |           d d d            n# 1 swxY w Y   d d d            n# 1 swxY w Y   lt-          |	|          \  }}|D ]}t/          j                     t?          |          dk    sJ t	          ||j                  }d t          ||j                  D             }|S )Natomjax_core.Atomc                h    t          | t          j                  rd S                     | d           S rQ   )r   r`   Literalr   )r  block_shape_envs    r@   read_block_shapez'jaxpr_subcomp.<locals>.read_block_shape  s3    $()) TtT***r?   c                V    t          | t          j                  r| j        n|          S rQ   )r   r`   r  val)r  envs    r@   read_envzjaxpr_subcomp.<locals>.read_env  s%    !$(899H488s4yHr?   varjax_core.Varc                    t          |t          j        t          f          }|sJ t	          |                      || <   d S rQ   )r   r   ValueKeyScalarBundler   )r  r  is_valid_typer  s      r@   	write_envz jaxpr_subcomp.<locals>.write_env  s@    sRX$?@@M##$s))###CHHHr?   c                    g | ]	}|j         
S r>   r#  rY   scopes     r@   r   z!jaxpr_subcomp.<locals>.<listcomp>  s    EEEu
EEEr?   )rK   c                >    g | ]\  }}t          ||j                  S r>   )_ensure_mlir_valuer   )rY   r   vs      r@   r   z!jaxpr_subcomp.<locals>.<listcomp>  s8     9 9 91 'q!&11 9 9 9r?   c                    g | ]	}|j         
S r>   r   rY   r  s     r@   r   z!jaxpr_subcomp.<locals>.<listcomp>  s    (((QV(((r?   c                    g | ]	}|j         
S r>   r   r  s     r@   r   z!jaxpr_subcomp.<locals>.<listcomp>  s    )))QV)))r?   c                    g | ]	}|j         
S r>   r  r  s     r@   r   z!jaxpr_subcomp.<locals>.<listcomp>  s    KKKUejKKKr?   
   )messagelevelz Exception while lowering eqn:
  z
With context:
  z
With inval shapes=c                $    t          | dd           S )Nr   getattrts    r@   r  zjaxpr_subcomp.<locals>.<lambda>&  s    wq'4'@'@ r?   z
With inval types=c                $    t          | dd           S Nr   r  r  s    r@   r  zjaxpr_subcomp.<locals>.<lambda>(  s    ga&>&> r?   z
In jaxpr:
z
Exception: z0Unimplemented primitive in Pallas TPU lowering: z?. Please file an issue on https://github.com/google/jax/issues.r   c                j    g | ]0\  }}t          |t          j                  rt          |          n|1S r>   )r   r`   r  r   )rY   r   r  s      r@   r   z!jaxpr_subcomp.<locals>.<listcomp>>  sH       
!S #3(899@k!nnnq  r?   )r  r  )r  r  ) 	constvarsr_   r  rJ   r:  rK   stackextendeqnssource_inforf   r   _source_info_to_location	primitiver  r   user_context	tracebacklowering_rulesskip_mlir_conversionsrk   rH  r  r)   TraceStopOpTraceStartOpr  r  r   r#  r  rR   )rA  r   r  r  r  r  r   rM  initial_name_stackcurrent_name_stackeqninvalsr  locrJ   rule_contextrK   poppedpushedr   r#  ansr$  outvalsr  r  s                           @@r@   r  r    sI    _
#/+ + + + + +
I I I I I I     
 u|S%566    ieROEit$$$EE0DEEE"$.///Z 7' 7'c3:&&F/))>CO$>> *  K 
'S]CJ C 
	&s'@	A	A /' /'3 /' /'	.	(	(= 5559 9 #FCJ 7 79 9 9&+SZ88*((SZ((())S[)))	
 
 LKk.D.JKKK
4
, ,' 	 	A
/



 	3 	3D

4r
2
2
2
2
2	s}-#  '*z ## ! 	 	 	
 		 		 		!"# " "" "@@&II" " >>GG	" "
  " "  " "  		 "L}!L L LM M 	M 
	' 'Is{C((((	#+a.#&&&_/' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /' /'d /,. ..&&  aO	V				%-((' //  ' 
.s]   K0C K9 HKI3	.A I.	.I3	3AKK0KK0 K!K00K4	7K4	c                N   t          | t          j                  r| S t          | t                    r| S t          | t          j        t          j        t          t          f          r"t          | t          |j                            S t          dt          |                      )Nz1Unsupported argument to a JAX primitive of type: )r   r   r  r  r   genericndarrayr   r   r   r   r   RuntimeErrorr   )r  r   s     r@   r  r  E  s    RX J_%% J#
BJU;<< s-dj99:::
GDIIGG  r?   c                D  
 t          t          ||                    
t          
fdt          | j        |          D                       }t	          |          \  }}|rd |D             \  }nd}t          || j        |          }t          || j        |          }	||	fS )Nc              3     K   | ]B\  }}|rt                    n)t          j        d |          t          j        d |          fV  CdS )r   N)r   r!   Slice)rY   r3  indexednon_slice_idx_iters      r@   r\   z4_convert_flat_indexing_to_indexer.<locals>.<genexpr>U  sr       ! ! !W 
<dQ""J$4Q$:$:;! ! ! ! ! !r?   c                P    h | ]#}t          |t          j                  |j        $S r>   r   r!   r  r   )rY   idx_avals     r@   r%  z4_convert_flat_indexing_to_indexer.<locals>.<setcomp>]  sB     K K Kx#-h
8H#I#IKHN K K Kr?   r>   )r  r_   r]   r   r(   	NDIndexer)ref_avalnon_slice_idxnon_slice_idx_avalsindexed_dimssplatted_idx_idx_avalssplatted_idxsplatted_idx_avalsint_indexer_shape
nd_indexernd_indexer_avalsr  s             @r@   !_convert_flat_indexing_to_indexerr  R  s    C/BCCDD  ! ! ! ! HN<88	! ! !   &,,B%C%C"," K K;M K K K x~7HII*18>02 2	%	%%r?   c               z   t          j        ||          }t          j        || j        dd                    }| j        ^}}t          j        ||d d f          \  }}	t          j        ||d d f          }
|                     |
| j        d         gd gt          |
          dz
  z            } t          | g|R d|	iS )Nr   r   rm   rJ   	args_tree)	r	   tree_unflattenrm   tree_flattentree_leavesrf   rJ   rR   _load_lowering_rule)rA  reftreer  indexersindexers_avalsr  r   	args_flatr   
avals_flats              r@   _get_lowering_ruler+  g  s     %dC00(+D#,qrr2BCC.,(Q"/hd0KLL)Y$hd%KLL*$Q'I4&C
OOa4G*HI 	 	 	# 
S	B9	B	B	B		B	BBr?   c               |   t          j        ||          }t          j        || j        dd                    }| j        ^}}}	t          j        |||d f          \  }
}t          j        |||d f          }|                     || j        d         gd gt          |          dz
  z            } t          | g|
R d|iS )NrK  r   r   r  r   )	r	   r!  rm   r"  r#  rf   rJ   rR   _masked_swap_lowering_rule)rA  r%  r  r&  r  r'  r(  r  val_avalr   r)  r   r*  s                r@   _swap_lowering_ruler/  {  s     %dC00(+D#,qrr2BCC.<(Hq"/hT0JKK)Y$40 * 	$Q'I4&C
OOa4G*HI 	 	 	# 
$C	I)	I	I	Iy	I	IIr?   c                Z   t          | t          t          j        f          r,t	          | t
          j                                                  S | j        t
          j                                        k    r| S t          j
        t
          j                                        |           j        S rQ   )r   r   r   r  r   r   	IndexTyper   r   r   IndexCastOpr   r7  s    r@   _make_indexr3    s}    C$%% .q",**,,---Vr|!!!!H		2<++--q	1	1	88r?   c                Z    | rt          |          S t          |t          j                  S )Nr   )r3  r  rw   r   )cast_to_indexr   s     r@   _maybe_cast_to_indexr6    s,     q>>	AG$?	@	@	@@r?   r  +tuple[indexing.Slice | int | ir.Value, ...]r5  *tuple[ir.Value, int | ir.Value, int, bool]c                   t          | t                    rJ t          | t          j                  r&t	          || j                  }| j        }| j        }d}nht          | t                    rt	          ||           }d}d}d}n<t          j
        |           rt          d|            t	          ||           }d}d}d}||||fS )NFr   Tz+Can only use ()-shaped and slice indexing: )r   slicer$   r  r6  startsizestrider   r   r   r   )r  r5  r;  r<  r=  squeezes         r@   _index_to_start_size_strider?    s     U#####X^$$  	::E8DZFGG#s  44EDFGG	x}} LJSJJKKK 44EDFG	fg	%%r?   indexerr  ref_block_shape tuple[int | pl_core.Mapped, ...]|tuple[tuple[ir.Value, ...], tuple[int | ir.Value, ...], tuple[int, ...], tuple[bool, ...], tuple[int | pl_core.Mapped, ...]]c                  t          | j                  }g g g g f\  }}}}|D ]}|t          j        u rt	          |d          dddfnt          t          |          |          \  }	}
}}|                    |	           |                    |
           |                    |           |                    |           t          |d           }|J | j        |f            t          d t          ||          D                       }t          |          t          |          t          |          t          |          |fS )Nr   r   Tc              3  $   K   | ]\  }}||V  d S rQ   r>   )rY   r3  r>  s      r@   r\   z0_indexer_to_start_size_stride.<locals>.<genexpr>  s>       . .JAw%,.a . . . . . .r?   )
r  r   rw   r   r6  r?  r   rm  r]   r_   )r@  rA  r5  indices_iterstartssizesstridessqueeze_dimsr3  r;  r<  r=  squeeze_dim
next_indexnew_ref_block_shapes                  r@   _indexer_to_start_size_striderN    sp    go&&,)+RR&&%, % %a  !22		
 	
 )l););]KK %E4 MM%	LLNN6$$$$L$''*			go?			 . .#e\2J2J . . . . . FmmEllGnnL
 r?   r%  r  state.AbstractRefStuple[ir.Value, tuple[int | pl_core.Mapped, ...], tuple[int | pl_core.Mapped, ...]]c                   |J |                                 }t          ||d          \  }}}}}t          d |D                       st          d          t	          d |D                       }	t
          j                                        t	          fd|D                       }
t
          j        	                    |
t          |j                  | j        j                  }t          j        || ||	          j        }t#          |          rst	          fd|D                       }
t
          j        	                    |
t          |j                  | j        j                  }t          j        ||          j        }||fS )	NFr5  c              3  *   K   | ]}|d u p|dk    V  d S Nr   r>   rY   r3  s     r@   r\   z _slice_memref.<locals>.<genexpr>  s/      44qa4i!16444444r?   z-Strided slices of references are unsupported.c              3  N   K   | ] }t          |t          j                  |V  !d S rQ   r   r   r  rU  s     r@   r\   z _slice_memref.<locals>.<genexpr>  s3      DDaJq"(,C,CDDDDDDDr?   c              3  T   K   | ]"}t          |t          j                  s|nV  #d S rQ   rW  rY   r3  ir_dynamic_sizes     r@   r\   z _slice_memref.<locals>.<genexpr>  sO       < <01 !+1bh 7 7 ,qq+< < < < < <r?   r~   c              3  T   K   | ]"}t          |t          j                  s|nV  #d S rQ   rW  rY  s     r@   r\   z _slice_memref.<locals>.<genexpr>  sW       E E23 #-Q"9"9 .-E E E E E Er?   )get_indexer_shaperN  allr   r]   r   
ShapedTypeget_dynamic_sizer   r   r   r   r   ro   r)   MemRefSliceOpr   r  MemRefSqueezeOp)r%  r  r@  rA  target_shaperG  rH  rI  rJ  dynamic_sizesstatic_sizestarget_ref_tyr  squeezed_ref_tyrZ  s                 @r@   _slice_memrefrg    s   
 
	$	$	$**,,,#

   8&%, 
44G444	4	4 O
M
N
NNDD5DDDDD-M2244/ < < < <5:< < < < <,-##%hn558( $ * *- 	-fmDDK# ; E E E E7CE E E E ELm'''77X* ( , ,O 
os
3
3
:C	o	r?   c                >    |D ]}t          | |||          \  } }| |fS rQ   )rg  )r%  r  rA  r'  r@  s        r@   
_index_refri    s;     : :g(h)8: :C	o	r?   T)frozenc                  (    e Zd ZU dZded<   ded<   dS )r  a  A container class for PRNG key data.

  We pass around keys as a KeyScalarBundle in the lowering pass rather than
  as a vector, since we want the key data to live in scalar registers rather
  than vector registers. This special dataclass exists so we can return
  multiple scalar values from load_op, because the load_op primitive does
  not allow multiple results.

  Attributes:
    scalars: A list of OpResults representing scalar key data during the
      lowering pass.
  r5   r   zlist[ir.OpResult]scalarsN)r:   r;   r<   __doc__r=   r>   r?   r@   r  r    s6           r?   r  c                  |                     |          \  }}}}|                     | j                  \  }}}}|^ }	}
|^ }}|t          | j        ^}}t	          ||||	          \  }}t          j        |j                  }t          |j	                  dk    }| j        ^}}| j
        \  }t          |j        t          j                  r!|st          d          t!          | g|R d|iS |s|st          d          t#          d |j        D                       rt          d          t'          |
|d          \  }}}}}t)          d	 |D                        }|rI| j
        d
         j        rt          d          t-          |t/          j        ||          j                  S t5          j        ||j                  }|r,t9          j        t=          |d          |||          j        }n*t?          j        t=          |d          ||          j        }t-          ||          }||k    r|S t
          j         !                    |j        tE          |j                            }t?          j#        ||          j        S )N#tpu.memory_space<smem>zPRNG keys must be loaded from SMEM. Did you set the memory space to TPUMemorySpace.SMEM in the BlockSpec for the PRNG key input?r   7Indexing into a ()-shaped Ref not yet supported on TPU.c              3  Z   K   | ]&}t          |t          j                   o|j        V  'd S rQ   r  rY   r  s     r@   r\   z&_load_lowering_rule.<locals>.<genexpr>8  M       	 	
 a)***6qw	 	 	 	 	 	r?   Cannot do int indexing on TPUTrR  c              3  *   K   | ]}|d u p|dk    V  d S rT  r>   rU  s     r@   r\   z&_load_lowering_rule.<locals>.<genexpr>B  /      ??!d,a1f??????r?   r   zCan only load scalars from SMEMr   r   )$	unflattenrm   r   rJ   ri  r   r   r   r5  ro   rn   r   r   r   KeyTyr   _prng_key_load_lowering_ruler  r   rN  r]  r   _maybe_cast_load_to_boolr   LoadOpr   r`   r   r)   StridedLoadOpr   r   r   r   r   ShapeCastOp)rA  r   r)  r   r%  r'  maskr  r(  slice_indexersr  r  rA  ref_typeis_smem_loadaval_outrG  rH  rI  need_stride	load_avalload_valvec_types                          r@   r$  r$    s   $..y99#xq#,#6#6s|#D#D (NAq#NC!.Ax	
(/A#	8_n6 6#]38$$(X*++/HH,,(Q+8
++ N < ; < < < (MiMMM9MMM	 Co C
AC C C 	 		 	 	   6 4
5
55!>	" " "&%!Q
 ??w??????+ 5
}Q :8999#&-V,,35 5 5"5???) Q iD9993  H }	d;;;S&J JJP %h99((O]x~0@@B B(		Hh	/	/	66r?   c               t   |                     |          \  }}}}| j        \  }t          |j        t          j                  sJ |j        j        j        }t          |          dk    rt          d          t          |          dk    rt          d|           g }t          |d                   D ]g}t          d|f|t                                }	t          |	|d          \  }
}}}}|                    t          j        ||
          j                   ht%          |t          |          	          S )
a;  Lowering rule for loading PRNG keys from SMEM.

  PRNG key loads are currently lowered as a list of scalar loads from SMEM,
  rather than a single vector load.
  We store these scalars in a bundle type called KeyScalarBundle, which has
  special case handling for functions that consume the key such as set_seed.
  rK  Seed key_data must be 2D.r   r   5Seed key_data of shape != (1, 1) not supported. Got: r   )r   r   r  TrR  rl  r   )rx  rn   r   r   r   ry  r   r   rR   r   r]   r  r  rN  rm  r   r|  r   r  )rA  r   r)  r%  r   r  rA  load_opsrZ   r  rG  s              r@   rz  rz  W  sS    $$Y//,#q!Q+8	HNDJ	/	////N(2/Q
9
:
::
?v%%
OoOOQ Q Q (#$$ 7 7a
QF/&+gg/ / /C6  FAq!Q
 OOFM#v..56666	U?5K5K	L	L	LLr?   r  tuple[ir.Value, jnp.dtype]c                   | j         t          j        k    r|S t          t                    }| j        st          t          j                 }t          j
                            t          j                            d          |          }t          j
                            |d          }t          j        ||          }t          j        |||          j        S t%          d          )a  Casts a memref load value to bool if the requested value is a bool.

  Mosaic does not support boolean-type memrefs, since booleans
  typically live in mask registers. We instead load booleans as integers from
  memrefs and move them to mask registers on load using this function.

  Args:
    out_aval: The output aval of the load.
    val: The input value.

  Returns:
    The loaded value, and the JAX dtype of the input value.
  ra  r   z'Boolean vector loads are not supported.)r   r   r   r   r   r   _cmpi_lowering_typesr   ne_pr   r   r   r   r   r   r   CmpIOpr   r   )out_avalr  load_scalar_typepred	predicate
const_zeros         r@   r{  r{  z  s     ^sy  J&'788	 I)D""2>#>#>r#B#BDIII##$4a88J!"2J??J<	3
33::
G
H
HHr?   c                    | j         t          j        k    r|S t          | d          }t	          j        ||          j        S )zACasts a boolean value back to an integer for storing in a memref.Tr   )r   r   r   r   r   ExtUIOpr   )expected_avalr  int_out_types      r@    _maybe_cast_store_to_memref_typer    s@     CI%%J 4HHH,	|S	)	)	00r?   c                  |                     |          \  }}}}|                     | j                  \  }}	}
}|^ }}|	^ }}|t          | j        ^}}t	          ||||          \  }}t          j        |j                  }t          |j	                  dk    }| j
        \  }t          |t
          j                  s#t          |t          |
j                            }t!          d |j        D                       rt%          d          |s|st          d          t'          ||d          \  }}}}}t)          d |D                        }|rh|
j        rt%          d	          t-          j        ||          j        }t3          |
|          }t5          |
|          }t-          j        |||           |S t9          |j                  }t;          |j                  D ]5\  }}t          |t<          j                  s|                     |d
           6tC          |          fd|D             }|"                    tG          |                    }t
          j$        %                    |j        t          |j        d                    }|rtM          j'        ||||          j        }ntQ          j        |||          j        }t5          |
|          }||k    rt
          j$        %                    |j        t          |j        d                    }tQ          j)        ||          j        }t
          j$        %                    |j        t          |j        d                    }tQ          j)        ||          j        }t3          |
|          }|rtM          j*        ||||           ntQ          j        |||           |S )Nro  r  c              3  Z   K   | ]&}t          |t          j                   o|j        V  'd S rQ   r  rr  s     r@   r\   z-_masked_swap_lowering_rule.<locals>.<genexpr>  rs  r?   rt  rp  TrR  c              3  *   K   | ]}|d u p|dk    V  d S rT  r>   rU  s     r@   r\   z-_masked_swap_lowering_rule.<locals>.<genexpr>  rv  r?   zCan only store scalars to SMEMr   c                N    g | ]!}|t           j        u rd nt                    "S rW  )rw   r   r   )rY   r   mem_slice_shape_iters     r@   r   z._masked_swap_lowering_rule.<locals>.<listcomp>  sB       
 aaD)=$>$>  r?   r   r   )+rx  rm   r   rJ   ri  r   r   r   r5  ro   rn   r   r  r   r   r   r  r   r   rN  r]  r   r   r|  r   r{  r  StoreOprr  r^   r!   r  rn  r  r)  r]   r   r   r)   r}  r   r~  StridedStoreOp)rA  r   r)  r   r%  r'  r  r  r  r(  r.  r  r  r  rA  r  is_smem_storer  rG  rI  r  r   mem_slice_shaperZ   r  mem_avalmem_aval_vec_typeresult_vec_typeval_vec_typer  s                                @r@   r-  r-    s    '00;;#xd*3*=*=cl*K*K'(NHa#NC!.Ax	
(/A#	8_n6 6# ]38$$(h+,,0II-+8	C	"	" H
c%6x~%F%F
G
G
GC 	 		 	 	   6 4
5
55	 C C
AC C C ;	  &!Wa
 ??w??????+ ~ 97888]3''.F%h77F
*8S
9
9C
N3V$$$M((/()) # #daa)** #Q"""o..     / __5#9#9_::(m''hn>>>@ @ B0#vwGGNFF],c6::AF(377#m''4@@@B BO88?F=$$X^4@@@B BL

\3
/
/
6C#Hf55& %sC1111
N3V$$$	-r?   c               <    ~ |D ]}t          j        ||          }|S rQ   )r)   assume_multiple)rA  r  valuesmultiples       r@   _multiple_of_lowering_ruler    s-    	 - -h

c8
,
,CC	*r?   c                     d fd}|S )NrA  rk   c          	        | j         \  }| j        d         j        s$fd}t          |d          } || ||          S t	          j        |j        t          j                  rat          j                 }t          j                 }t          j	        
                    t          j        
                                |          }n~t	          j        |j        t          j                  rt          d          t	          j        |j        t          j                  rt          d          t          d|j         d          t          | j        d                   }t          j                            ||          }	t%          j        ||	          }
t)          j        |||
t          j        
                    d	 |D                                 }|j        S )
Nr   c                   | t           j        df         } d |D             } | |d          } t          j        |           S )N.c                    g | ]}|d z   S rW  r>   )rY   axiss     r@   r   zTreduce_lowering_rule.<locals>._lowering_rule.<locals>._proxy_fun.<locals>.<listcomp>  s    ***Tq***r?   T)r  keepdims)r   newaxisr>  )r  axes	reduce_fns     r@   
_proxy_funz@reduce_lowering_rule.<locals>._lowering_rule.<locals>._proxy_fun  sO    #+s"#**T***i$666 {3r?   Fr  r  z)Reductions over integers not implemented.zReductions over z not implemented.c                    g | ]?}t           j                            t           j                            d           |          @S )ra  )r   r   r   r   r   rr  s     r@   r   z@reduce_lowering_rule.<locals>._lowering_rule.<locals>.<listcomp>   s:    RRRR^ ; ;B ? ?CCRRRr?   )rm   rn   r   r  r   r   r   floatingr   r   r   F32Typesignedintegerr   unsignedintegerr   DenseElementsAttr	get_splatr   r   r   MultiDimReductionOpr9  r   )rA  r   r  x_avalr  proxy_loweringkindr  out_typeidentityaccopr  type_to_identitytype_to_kinds               r@   _lowering_rulez,reduce_lowering_rule.<locals>._lowering_rule  s   IV=! /
          !
u. . .n^C....
~flCL11 
>#,'dS\*cLRZ^^--s33cc	c&7	8	8 > KLLL	c&9	:	: > KLLL
<V\
<
<
<> > >s}Q/00H#--h<<H

8X
.
.C		#	
RRTRRR	
 	
	
 
B 9r?   r  r>   )r  r  r  r  s   ``` r@   reduce_lowering_ruler    s6    ( ( ( ( ( ( ( (R 
r?   z-infinf        c               F    d }t          |d          } || ||          S )Nc               b    t          j        | dd          }t          j        ||          dk    S N      ?r  r  )r   whereminargr  	float_args      r@   _proxy_reducez0_reduce_and_lowering_rule.<locals>._proxy_reduceR  1    
 	#sC((I794(((3..r?   Fr  r  r  rA  r   r  r  r  s        r@   _reduce_and_lowering_ruler  Q  C    / / / e- - -.	QT	*	*	**r?   c               F    d }t          |d          } || ||          S )Nc               b    t          j        | dd          }t          j        ||          dk    S r  )r   r  maxr  s      r@   r  z/_reduce_or_lowering_rule.<locals>._proxy_reducea  r  r?   Fr  r  r  r  s        r@   _reduce_or_lowering_ruler  `  r  r?   c                  | j         \  }| j        \  }t          j        |j        t          j                  r#d }t          |d          } || |||          S |rdgt          |          z  }t          ||j	                  D ]
\  }	}
|
||	<   t          |          }t          j                            |t          |j                            }t          j        ||          j        }||j	        k    r|S t          j                            |j	        t          |j                            }t          j        ||          j        S )Nc               |    t          j        | dd          }t          j                            |||          }|dk    S )Nr   r   )r   r  jaxr   broadcast_in_dim)r  r   broadcast_dimensionsint_val	bcast_vals        r@   r  z3_broadcast_in_dim_lowering_rule.<locals>._proxy_funz  s9    	#q!$$g'**7E;OPPi!^r?   Fr  )r   r  r   )rm   rn   r   r   r   r   r  rR   r_   r   r]   r   r   r   r   r   r~  r   BroadcastOp)rA  r  r   r  aval_inr  r  r  out_shape_listrZ   r3  	out_shaper  s                r@   _broadcast_in_dim_lowering_ruler  o  sg    |*7+8^GM39-- J
   U, , ,N>S4HJ J J J  
S3u::%N('-88  1nQn%%I}  $X^44 H 
Xs
+
+
2CHN""j]n'77 ( 
	Hc	*	*	11r?   c                	   |\  \  }}}| j         \  }t          |          }	|	j        t          fdt          j        t          j        t          j        t          j        fD                       r!t          j	        
                    d          }
n_t          j                                      r!t          j        
                    d          }
nt          | j         d         j                  t          d | j        D                       rt          d| j                   | j        \  }}|dk    r;|dk    r4| j        d         j        d         dk    r| j        d         j        | j        d         j        k    rt%          j        | j        d         j        | j         d         j                  }t          j        
                    t+          |          t-          | j         d         j                            }| j        d         j        |k    rt/          j        ||          }| j        d         j        |k    rt/          j        ||          }t          |                    |j        d         f                    }t5          j        |t          j                            ||
                    }t/          j        t          j                             d	          t5          j!        ||          |t          j"        
                    t          j        
                    t          j        #                    d
          d          g                    }t/          j$        |	|          j%        S |dk    rd}n|dk    rd}nt          |dk    rd}n|dk    rd}nt          |)|d         |d         k    rt          d          |d         }||tL          j'        j(        k    rd }nG|tL          j'        j)        k    r t          j                             d          }nt          d|           t5          j        |	t          j                            |	|
                    }tU          j+        |	||||||          }|j%        S )Nc              3  B   K   | ]}|                               V  d S rQ   )r   )rY   clsval_types     r@   r\   z-_dot_general_lowering_rule.<locals>.<genexpr>  sC       	 	
 
nnX	 	 	 	 	 	r?   r  r   c              3  F   K   | ]}t          |j                  d k    V  dS )rK  N)rR   r   rr  s     r@   r\   z-_dot_general_lowering_rule.<locals>.<genexpr>  s/      11qQW	111111r?   z,Only 2D tensors supported in dot; received: rW  r   r   z#vector.kind<add>ra  F)r   Tz%Per-operand dot precision unsupportedz#tpu.contract_precision<fp32>zUnsupported dot precision: )transpose_lhstranspose_rhs	precision),rn   r   element_typer  r   BF16Typer  Float8E5M2TypeFloat8E4M3FNTyper   r   r   r   r   r   r   rm   r   r   broadcast_shapesr   rr  r   r   r  r)  r   r   r  r  r  r|   r}   MulFOpr9  r   r~  r   r   	PrecisionDEFAULTHIGHESTr)   MatmulOp)rA  r   ydimension_numbersr  r   lhs_dimsrhs_dimsr  r  r  lhs_avalbcast_shapered_typer  redr  r  precision_attrout_tiler  r  s                        @r@   _dot_general_lowering_ruler    sQ    .8X+8X&&("( 	 	 	 	 +
*



		 	 	   6 ,

8S
)
)CC	~  ** 6
.

Xq
)
)CC
cmA.4
5
5511CL11111 
Es|EE   +(A(d**s|A/DQ/G1/L/L
|AQ 555(
,q/
q!1!7 k M%%
{

.s}Q/?/EFF k 
a	+	-	-{A..	a	+	-	-{A..xhnQ6G5IJJKKH

"&003?? C 
$
.//Q
^ ; ;B ? ?CCD	
 	
	 C h,,33MM4MM
MM4MM
|y|## GHHH!I)s}'<<<NNCM)))\''' NN GIGG
H
HH$..x== ( |1h!  "
 
r?   c                  | j         }t          j        |t          j         d                    r0|                     t          j                  } t          | |          S t          j        |t          j                  r|j        dk     r|                     t          j                  } t          j        |t          j                  r*|j        dk     r|                     t          j	                  } |                     |          S t          j        |t          j                  rHt          j        |t          j                  r|j        dk     r|                     t          j	                  } |j        dk     rzt          j
        |          j        t          j
        |          j        }}t          j        | ||          } |                     t          j                                      |          S |                     |          S t          j        |t          j         d                    r|                     t          j                  } |                     t          j	                  S t          d| d|           )NrN   to_dtyperO  zUnsupported cast: z -> )r   r   r   astyper+   _convert_helperr  itemsizer  r   iinfor  r  clipr   r   )r   r  
from_dtypeminvalmaxvals        r@   r  r    s   w*^J	& 1 122 1	A1x0000^J 122 Q
((39

a
~h--  (2Ca2G2G
((3;

a88H^J-- !
~h 122 
		q	 	 HHS[!!		Q		8,,0#)H2E2E2IHQ''xx	""))(333XXh	"(6"2"2	3	3 
((39

a88CK   KKKKKLLLr?   c                  ~~| j         d         }| j        d         j        }t          |          }t	          j        |t          j                  rt          |         }||k    r|S t          j        |t          j	                  rt          j        |t          j	                  rl|j
        |j
        k     r%|j
        dk    rt          j        ||          j        S |j
        |j
        k    r%|j
        dk    rt          j        ||          j        S nt          j        |t          j                  rt          j        |t          j                  rl|j
        |j
        k     r%|j
        dk    rt          j        ||          j        S |j
        |j
        k    r%|j
        dk    rt          j        ||          j        S n t          j        |t          j	                  rSt          j        |t          j                  r4|j
        |j
        cxk    rdk    rn nt          j        ||          j        S t          j        |t          j                  rSt          j        |t          j	                  r4|j
        |j
        cxk    rdk    rn nt          j        ||          j        S |t          j        k    r?t          j        |t          j                  r |j
        dk    rt          j        ||          S t          j        |t          j                  r|t          j        k    r|j
        dk    rt.          t0          j                 }t4          j                            t4          j                            d          |          }	t?          |          }
t4          j                            |
d          }t          j         |
|          }t          j!        |	||          j        S  tE          tG          j$        tJ          |          d          | |          S )Nr   rO  ra  r  Fr  )&rn   rm   r   r   r   r   r   r  UNSIGNED_TO_SIGNEDr  r  r   ExtFOpr   TruncFOpr  ExtSIOpTruncIOpFPToSIOpSIToFPOpr   integerextuir  r   r  r   r   r   r   r   r   r   r  r  r<  partialr  )rA  r   	new_dtype	weak_typeshardingr  	old_dtyper  r  r  
const_typer  s               r@   #_convert_element_type_lowering_ruler(    s    ]1(l1o#)X&&( ]9c122 ."9-I)H^Is|,, &92 2 &9 I...93E3J3J\(A&&--		i0	0	0Y5G15L5L^Ha((//
~i!233 9"9 9 9 I...93E3J3J]8Q''..		i0	0	0Y5G15L5L^Ha((//
~i.. 93>"4 4 9i05555A55555>(A&&--
~i!233 99 9 9i05555A55555>(A&&--39
.CK
0
0 


!
!;x###	nY,,
9
sy
 
 


!
!)D""2>#>#>r#B#BDIII"9--J##J22J!*j99J<	1j1188
+9$_yIII$)
+ 
+ 
++.
3 
3 3r?   c                @   |t           t          d |D                       rt           | j        d         j        s2t	          j        t          | j        d                   |          j        S t	          j	        t          | j        d                   |          j        S )Nc              3     K   | ]}|d u V  	d S rQ   r>   )rY   r[   s     r@   r\   z)_reshape_lowering_rule.<locals>.<genexpr>L  s&      &&qd&&&&&&r?   r   )
r   r  rm   r   r   r  r   rn   r   r~  )rA  r   	new_sizes
dimensionss       r@   _reshape_lowering_ruler-  I  s    
&&I&&&&& 
	a	 KocmA.>??CCJJ		OCM!,<==q	A	A	HHr?   c                    ~| j         \  }| j        \  }|j        s1t          j        |g dgt          |j                  z            j        S t          j        t          | j        d                   |          j        S Nr   )	rm   rn   r   r   	ExtractOprR   r   r~  r   )rA  r   r,  r  r  s        r@   _squeeze_lowering_ruler1  V  sp    |*7+8	 DArA3W]););#;<<CC		OCM!,<==q	A	A	HHr?   c               j    t          j        t          | j        d                   ||          j        S Nr   )	dimension)r)   ConcatenateOpr   rn   r   )rA  r4  xss      r@   _concatenate_lowering_ruler7  b  s5    		cmA&''y
 
 

r?   c                l    t          | j        d                   }t          j        ||          j        S r3  )r   rn   r)   IotaOpr   )rA  r   r   r4  r  s        r@   _iota_lowering_ruler:  k  s.    S]1-..(	H		2	2	2	99r?   c                   |dk    rt           t          | j        d                   }t          j        |||          j        S )N)r   r   r   )r   r   rn   r   TransposeOpr   )rA  r   permutationr  s       r@   _transpose_lowering_ruler>  s  s?    F
S]1-..(		Ha	5	5	<<r?   c                   |j         }|j         }|j        r|j         }n|j        r|j         }t          | t          j        t          j        t          t          f          rXt          |dd           t          j
                                        k    r|j        }nt          |          }t          | |          } t          |t          j        t          j        t          t          f          rXt          | dd           t          j
                                        k    r| j        }nt          |          }t          ||          }t          |j                  }|j        |j        k    rBt          j                            |t          |                    }	t%          j        |	|           } |j        |j        k    rBt          j                            |t          |                    }
t%          j        |
|          }| |fS r  )r   r$  r   r   r  numberr   r   r  r   r1  r   r   r   r   rr  r   r   r   r  )r   r  r  y_avalr  x_dtypey_dtyper   r  x_tyy_tys              r@   _bcastrF  }  s   L'L' lGG lGBJ	3677 "q&$2<#3#3#5#555&ii#G,,iAy!!ABJ	3677 "q&$2<#3#3#5#555&ii#G,,iAy!!A8>"")\X^##=Y(9'(B(BCCD4##A\X^##=Y(9'(B(BCCD4##A	
A+r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          |j                  Nr   r   )rF  rm   rn   r   r   r   r   r   AddIOpr   r  AddFOpr   rA  r   r  r  s       r@   _add_lowering_rulerL        	1cl1os|Aa8H	I	I$!Q+8^HNCK00 %<1$$^HNCL11 %<1$$HN+++r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          j        |j        t          j                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r  r   MaxSIOpr   r  MaxUIOpr  
MaximumFOpr   rK  s       r@   _max_lowering_rulerR        	1cl1os|Aa8H	I	I$!Q+8^HNC$566 )=A%%
~hnc&9:: )=A%%
~hncl33 )Aq!!((HN+++r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          j        |j        t          j                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r  r   MinSIOpr   r  MinUIOpr  
MinimumFOpr   rK  s       r@   _min_lowering_rulerX    rS  r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r   r   SubIOpr   r  SubFOpr   rK  s       r@   _sub_lowering_ruler\    rM  r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r   r   MulIOpr   r  r  r   rK  s       r@   _mul_lowering_ruler_    rM  r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          j        |j        t          j                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r   r   DivSIOpr   r  DivUIOpr  DivFOpr   rK  s       r@   _div_lowering_rulerd        	1cl1os|Aa8H	I	I$!Q+8^HNCK00 &=A%%^HNC$788 %=A%%
~hncl33 %<1$$HN+++r?   c                   t          ||| j        d         | j        d         | j        d                   \  }}| j        \  }t          j        |j        t          j                  rt          j        ||          j	        S t          j        |j        t          j
                  rt          j        ||          j	        S t          j        |j        t          j                  rt          j        ||          j	        S t          |j                  rH  )rF  rm   rn   r   r   r   r   r   RemSIOpr   r  RemUIOpr  RemFOpr   rK  s       r@   _rem_lowering_rulerj    re  r?   c                0   | j         \  }t          j        |j        t          j                  rt          j        |          j        S t          j        |j        t          j                  rt          j	        |          j        S t          |j                  rQ   )rn   r   r   r   r   r   AbsIOpr   r  AbsFOpr   )rA  r   r  s      r@   _abs_lowering_rulern    sm    +8^HNCK00 !;q>>  ^HNCL11 !;q>>  HN+++r?   c                    | j         \  }|                     t          j        d|j                  |fdg| j        R           }t          |t          j        d|j                  |          S )Nr>   r  r   rw  )	rm   rf   r`   r   r   rJ   r\  r   r   )rA  r   r  new_ctxs       r@   _neg_lowering_rulerq    sp    l)6KK$R66?*)**   ' 
GRXav|%D%D%Da	H	HHr?   c                |   t          j        | j        t           j                  r| dk                        | j                  S t          j        | j        t           j                  r=| dk                        | j                  | dk                         | j                  z
  S t          j        | j        t           j                  ro| dk                        | j                  | dk                         | j                  z
  }t          j        t          j        |           t           j	        |          S t          )Nr   r  )r   r   r   r  r  r   r  r  isnannanr   )r   r  s     r@   _sign_lowering_helperru    s    ^AGS011 $F??17###^AGS[)) =E>>!'""a!e^^AG%<%<<<^AGS\** 1r6//!'
"
"a"f__QW%=%=
=C9SYq\\37C000r?   c                B     t          t          d          | |          S NFr  )r  ru  rA  r   s     r@   _sign_lowering_rulery  -  s"    	A(5	A	A	A#q	I	IIr?   c                4    t          j        |          j        S rQ   )r   RsqrtOpr   rx  s     r@   _rsqrt_lowering_ruler|  4      	a	r?   c                4    t          j        |          j        S rQ   )r   SqrtOpr   rx  s     r@   _sqrt_lowering_ruler  ;      	Q	r?   c                4    t          j        |          j        S rQ   )r   ExpOpr   rx  s     r@   _exp_lowering_ruler  B      	A	r?   c                   t          |t          j                  s|dk    rt          j        |          j        S t          ||| j        d         | j        d         | j        d                   \  }}t          j	        ||          j        S )Ng       @r   r   )
r   r   r  r   Exp2Opr   rF  rm   rn   PowFOprA  r   r  s      r@   _pow_lowering_ruler  I  sp    	Arx	 	  !Q"WW;q>>  	1cl1os|Aa8H	I	I$!Q	Q			!!r?   c               P     t          t          j        d          | ||          S )NFr  )r  )r  lax_internal_integer_powr  s      r@   _integer_pow_lowering_ruler  T  s2    	E<,u	E	E	E	1
 
 
 r?   c                :     t          d d          | |          S )Nc                T    t          j        t          j        d          | z            S NrK  )r   expr   logr   s    r@   r  z%_exp2_lowering_rule.<locals>.<lambda>_  s    SWRVAYY]33 r?   Fr  r  rx  s     r@   _exp2_lowering_ruler  \  s/     
M33e	L	L	L	1
 
 r?   c                   t          j        |          j        }t          j        |          j        }| j        d         }t          |          }|j        dk    rt          d|          }n"t          j
        |t          d                    }t          j        ||          j        }t          j        ||          j        S )Nr   r>   r  r  )r   NegFOpr   r   r  rn   r   r   r   r   r  rJ  rc  )rA  r   neg_x	exp_neg_xr  r  onedenoms           r@   _logistic_lowering_ruler  g  s    
,q//
 %j&)]1(X&&(^r
cX
.
.
.CC

X{3'7'7
8
8C
,sI
&
&
-%	c5	!	!	((r?   c                4    t          j        |          j        S rQ   )r   SinOpr   rx  s     r@   _sin_lowering_ruler  w  r  r?   c                4    t          j        |          j        S rQ   )r   TanhOpr   rx  s     r@   _tanh_lowering_ruler  ~  r  r?   c                4    t          j        |          j        S rQ   )r   LogOpr   rx  s     r@   _log_lowering_ruler    r  r?   c                4    t          j        |          j        S rQ   )r   Log1pOpr   rx  s     r@   _log1p_lowering_ruler    r}  r?   c                   |dk    rt          j        |          j        S |dk    rt          j        |          j        S t	          d|           )Nr   r   zUnsupported rounding method: )r   RoundOpr   RoundEvenOpr   )rA  r   rounding_methods      r@   _round_lowering_ruler    sT    <??!!!A%%
OoOO
P
PPr?   r   rK     rO        c                   t          |||j        d         |j        d         |j        d                   \  }}|j        \  }}|j        |j        f}t	          d |D                       ret
          |          }t          j                            t          j	        
                    d          |          }t          j        |||          j        S t	          d |D                       ret          |          }t          j                            t          j	        
                    d          |          }t          j        |||          j        S t!          d          )Nr   r   c              3  T   K   | ]#}t          j        |t           j                  V  $d S rQ   )r   r   r   rY   r   s     r@   r\   z%_cmp_lowering_rule.<locals>.<genexpr>  s0      @@s{	+	+@@@@@@r?   ra  c              3  T   K   | ]#}t          j        |t           j                  V  $d S rQ   )r   r   r  r  s     r@   r\   z%_cmp_lowering_rule.<locals>.<genexpr>  s0      
C
C53>%..
C
C
C
C
C
Cr?   zMixed dtype operands in cmp)rF  rm   rn   r   r]  r  r   r   r   r   r   r   r  r   _cmpf_lowering_typesCmpFOpr   )	primrA  r   r  r  rA  r   r  r  s	            r@   _cmp_lowering_ruler    s$   	1cl1os|Aa8H	I	I$!Q<.&&<%&@@@@@@@ 0%D""2>#>#>r#B#BDIII<	1a((//

C
CF
C
C
CCC 0%D""2>#>#>r#B#BDIII<	1a((//9:::r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   AndIOpr   r  s      r@   _and_lowering_ruler    ;    	1	4s|	4cm	4	4	4$!Q	a			""r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   OrIOpr   r  s      r@   _or_lowering_ruler    s;    	1	4s|	4cm	4	4	4$!Q	Q			!!r?   c                l   | j         d         }t          |j                  }|j        st	          d|          }nbt          |          }t          j                            |d          }t          j
        |t          j                            ||                    }t          j        ||          j        S )Nr   )rn   r   r   r   r   r   r   r   r   r   r   r  r  XOrIOpr   )rA  r   r  out_scalar_type	minus_oner  scalar_minus_ones          r@   _not_lowering_ruler    s     ]1(%hn55/	 	B00II x((H~))/2>> "&00;KLL I 
a	#	#	**r?   c                   t          |          dk    rt          d          | j        d d         \  }}|j        t	          j        t          j                  k    rUt          | j        |g|                    t          j                  gd g          } t          d d          ||          }|s|S |\  }t          j        |||          j        S )	Nr   z+select_n only supported with <= 2 argumentsrK  rw  )rm   rn   rJ   c                    | dk    S r/  r>   r  s    r@   r  z)_select_n_lowering_rule.<locals>.<lambda>  s
    qAv r?   Fr  )rR   r   rm   r   r   r   rk   rl   r)  r  r   SelectOpr   )rA  r  r   r  	pred_avalr  	lower_ctxr  s           r@   _select_n_lowering_ruler    s    YY]]
K
L
LLl2A2&)V_****###"(#334V	  I ?9%%>>>y$OOD	 H"!	a	#	#	**r?   c                V    t          j        ||           }t          j        ||          S rQ   )r   maximumminimum)r  operandr  ress       r@   _clampr    s%    GS!!#	S#		r?   c                F     t          t          d          | |||          S )z0Compute minimum_p(maximum_p(min, operand), max).Fr  )r  r  )rA  r  r  r  s       r@   _clamp_lowering_ruler    s%    	26E	2	2	23Wc	J	JJr?   c               v   d | j         D             }t          j        |ddg|          \  }\   t          |          D ]x}|r||z
  dz
  }t	          |          }| j                            dg| j                  }	t          |	||g|R  }
t          |
          fdt          ||          D             }y|S )Nc                D    g | ]}t          |t          j                   S r>   )r   r   r   r   s     r@   r   z&_for_lowering_rule.<locals>.<listcomp>#  s5       26*T5,
-
--  r?   r>   F)should_discharger   r  c                <    g | ]\  }}|rt                    n|S r>   )r   )rY   r  r3  non_ref_args_iters      r@   r   z&_for_lowering_rule.<locals>.<listcomp>2  s@       Aq $%+!  r?   )rm   state_dischargedischarge_stater  r   rl   rf   rJ   r  r  r_   )rA  r   nstepsreverseunrollwhich_linearr  r  rZ   rl   non_ref_argsr  s              @r@   _for_lowering_ruler    s    :=,   -R5"<+;"<  )% ==  a 
1*q.aAA+33,3+, 4   !!15!CdCCCL\**   .//  DD 
+r?   r;  int | ir.Value	num_stepshas_loop_indexr  r   c               R     fd}t          |t          j                  srt          |t          j                  sX||k    rRt          |||z             D ]<}	 |t	          |	t          t          j        d                              |          }=|S |dk    rt          d|d|d          t          |t          j                  }
t          j        |
t          |t          j                            }t	          dt          t          j        d                              }t          j        |
|||          }t          j        |j                  5  |j        }|j        } |||          }t          j        |           d d d            n# 1 swxY w Y   |j        S )Nc                >   r3j                             j                  }t          |g| |R  }nf~ j                             j        d t	                             j        t	                    dz   d          z             }t          |g|R  }|S )Nr  r   )rl   rf   rJ   r  rR   )rZ   r  rl   r  rA  r  r   s      r@   	_run_bodyz+_lower_jaxpr_to_for_loop.<locals>._run_bodyA  s     
D-55' 6 ) )+UFVFQFFFFdd
-55'V5S[[1_--./ 6   +UCVCdCCCdKr?   r+   r  r   zOnly unroll=num_steps=z$ and unroll=1 supported. Got unroll=r   )r   r   r  r  r   r   r   r   r   r  rw   r   r   addir   ForOpInsertionPointrl  induction_variableinner_iter_argsYieldOpresults)rA  r   r;  r  r  r  r  r  r  rZ   lbdubdstepfor_opiv
inner_args	inner_outs   ``  ``           r@   _lower_jaxpr_to_for_loopr  <  s   
        UBH
%
%BH-- v

 5%)+,,  Y
a#4SYw5G5G#H#H
I
I
I
 dd Kq[[
JyJJJJJL L L5'"=>>#
3*9g6QRRSS#	Q"3CIg4F4F"G"G	H	H	H$9S#tT**&	%%  		"B'J	"j))IK		              
 
s   /FFFc                  t          |||z             D ]}|ra| j                            | j                  }t	          ||g|t          |t          t          j        d                              |R  }e| j                            | j        d t          |                   | j        t          |          dz   d          z             }t	          ||g||R  }|S )Nr  r+   r  r   )
r  rl   rf   rJ   r  r   r   r   r   rR   )	rA  r   r;  r  r  r  r  rZ   rl   s	            r@   !_lower_jaxpr_to_unrolled_for_loopr  j  s!    	)** D Da D-55' 6 ) )
E$*
a#4SYw5G5G#H#H
I
I
I   dd
 -55'V5S[[1_--./ 6   +UCVCdCCCdd	+r?   lineartuple[bool, ...]lengthr  
bool | int
num_consts	num_carry_split_transposec               *   ~t          |	          |z
  |z
  }
|
rt          |rt          ~~
~|j        |j        }}|rt          ~t	          j        |||          \  }}t          |	|g          \  }}	t          | j        |g          \  }}|r|	^}}	|dd          }nd}t          t          ||          }t          t          |	|          }	t          | ||||g|	R ||d}|r3t          |t          t          j        d                              g|}|S )Nr   r   r  r  r+   r  )rR   r   r   r  r  pattern_match_scan_to_fori_loopr'   rm   r:  r  r  r   r   r   r   )rA  r   r  r  r  r  r  r  r  r  num_extensivejaxpr_constsr  r  consts_avals
args_avalsloop_index_startr  s                     r@   _scan_lowering_ruler    su    d))j(94----'''mWU\%,,,&FZ % D:,//,&$'zlCC,
 "tABBJJ!6<88&	z	2	2$ 	5"F		 	$2	 	 	#  v!239W3E3E!F!FH H H C 
*r?   c          	     (   t          |||g          \  }}}	|	d d         |	dd          c\  }
}}t          |                     | j        d |dz            | j        |dz   d          z             ||
t	          j        ||
          |g|R ddd}||g|S )NrK  r   r  Tr  )r'   r  rf   rJ   r   subi)rA  
fori_jaxprcond_nconsts
cond_jaxprbody_nconsts
body_jaxprr  r   body_constscarrylbubfor_outs                r@   _lower_while_via_forir    s     %TL,+GHH![%!9eABBi.(2rD$	kk'(:,*:(:;\A-//01    jR     ' b	7	r?   c          	        t          j        ||||          \  }}|t          | g|R |||||dS t          |||g          \  }}	}
t          | j        ||g          \  }}}d |D             }d |	D             }d |
D             }g |||}t          j        ||          } |j        j        j	        | }t          |j
        ||g          \  }}}g ||}t          j                            |          5  t          | j                            g ||          |j        g|R  \  }t          j        ||j
                   d d d            n# 1 swxY w Y    |j        j        j	        | }t          |j
        ||g          \  }}}g |||}t          |||g          \  }}}t          j                            |          5  t          | j                            g ||          |j        g||R  }g |||} | rt          j        |            d d d            n# 1 swxY w Y   t+          |j                  }!|!||z   d          S )N)r  r  r  r  r  c                    g | ]	}|j         
S r>   r   rr  s     r@   r   z(_while_lowering_rule.<locals>.<listcomp>      222af222r?   c                    g | ]	}|j         
S r>   r  rr  s     r@   r   z(_while_lowering_rule.<locals>.<listcomp>  r  r?   c                    g | ]	}|j         
S r>   r  rr  s     r@   r   z(_while_lowering_rule.<locals>.<listcomp>  s    '''A'''r?   r  )r   pattern_match_while_to_fori_loopr  r'   rJ   r   WhileOpbeforeblocksrm  	argumentsr   r  at_block_beginr  rl   rf   r   	conditionafteryield_rr  results_)"rA  r  r  r  r  r  r  errcond_constsr  r  cond_const_block_shapesbody_const_block_shapescarry_block_shapescond_const_typesbody_const_typescarry_types	all_typeswhile_opbefore_blockcond_consts_r   carry_	cond_argscondafter_blockbody_consts_all_argscond_const_argsbody_const_args
carry_argsloop_outall_handlesall_outs"                                     r@   _while_lowering_ruler>    s    !A,
L /*c  	  !!    %/
\<(% %!+{E !L,#?@@ G24F 32k22222k222'''''+B B#3BkB)[D))(.'.	:,&\" ,6 '&v&)	''55 0 0$$H2H5GH 	% 	
 	
 		
 
  FT M$.///0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -%,i8+'1\"( ($,f 5|4l4V4(1;|,2 2./?J 
''44  $$H2H5GH 	% 	
 	
 		
 
 
  H BOAoAAK 	j               "##'	,..	//s&   7AEEEAH$$H(+H(c               l   |^}}t          t          | j                  }t          j        t          j        j        |t          d|j                            j	        }t          j        ||d          }| j                            | j        dd                    }t          j        |j                  5  t%          |          dk    rEt'          | t          j        |t          d|j                            j	        g|R d|dd          i}nt+          ||d         j        g|R  }t          j        |           d d d            n# 1 swxY w Y   t          j        |j                  5  t+          ||d         j        g|R  }t          j        |           d d d            n# 1 swxY w Y   |j        S )Nr   T)hasElser   r  rK  branches)r:  r   rn   r   r  CmpIPredicatener   r   r   r   IfOprl   rf   rJ   r   r  
then_blockrR   _cond_lowering_rulerZ  r  r   r  
else_blockr  )	rA  rA  r  index	out_typesr  if_oprl   r  s	            r@   rF  rF  	  s<   ,%$/3=11)	e[EJ%?%?
 

  (4D
1
1
1%)11#ABB' 2   
)**   8}}q

,uk!UZ88
9
9
@    ABB<	 cc *HQK,=EEEEcK               
)**  
((1+*;
Cd
C
C
CCK               
s%   4BEE
E
*.F$$F(+F(c               h    | j                             | j                  }t          ||j        g|R  S Nr  )rl   rf   rJ   r  r   )rA  r   r  r   rl   s        r@   _pjit_lowering_rulerM  :	  s:    )11s?O1PP	'	<t	<	<	<<r?   
call_jaxprjvp_jaxpr_thunksymbolic_zerosc                   ~|rt           |rt           |j        rt           | j                            | j                  }t          ||j        g|R  S rL  )r   r  rl   rf   rJ   r  r   )rA  rN  rO  r  rP  r  rl   s          r@   _custom_jvp_call_lowering_rulerR  B	  sg     ...***111)11s?O1PP	')9	AD	A	A	AAr?   c                    ~ ~~g S rQ   r>   )rA  r  kwargss      r@   _debug_callback_lowering_rulerU  V	  s    	4	)r?   r  c                   | j         j        t          d| d          t          | j         j                  }d|cxk    r|k     sn t          d| d|           | j         j        |         S )Nzprogram id: z- was passed, but user did not provide a grid.r   %user passed in program id with axis: , but grid only has length: )rl   rI   r   rR   )rA  r  r  s      r@   _program_id_lowering_rulerY  _	  s    +3
JtJJJ   s#566&
t



f




	 	 		 	   
		/	55r?   c               .   t          | j        j                  }d}t          | j        j                  D ]!}|t          ||v          z  }||dz   k    r n-"t          d| dt          | j        j                             t          j	        |          S )Nr   r   rW  rX  )
r(  rl   rH   r  rU   r   r   rR   r)   iteration_bound)rA  r  mapped_axesseen_user_axesrZ   s        r@   _num_programs_lowering_ruler^  m	  s    C(455+.%/00  ac!;.///N!!e " 	: 	: 	:,677	: 	:   
	Q		r?   c               f    | j         \  }t          j        t          |          |||          j        S rQ   )rn   r)   RepeatOpr   r   )rA  r   repeatsr  r  s        r@   _repeat_lowering_rulerb  }	  s,    +8	oh//D'	B	B	IIr?   c               l    | j         \  }t          j        t          |          |||||          j        S )N)r=  stride_dimension)rn   r)   DynamicRotateOpr   r   )rA  r   shiftr  r=  stride_axisr  s          r@   _roll_lowering_rulerh  	  sF     +8		h
"
 
 
 r?   c                    | j         \  }|dgt          |          z  }t          j        |          t          j        |          z
  }t	          j        t          |          ||||          }|j        S )z!Lowers a slice to vector dialect.Nr   )rn   rR   r   r   r   ExtractStridedSliceOpr   r   )rA  r   limit_indicesstart_indicesrI  r  rH  r  s           r@   _slice_lowering_rulerm  	  ss     +8_cC&&&G
(=
!
!BH]$;$;
;%#hM5' " 
r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   r  r   r  s      r@   _xor_lowering_rulero  	  r  r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   ShLIOpr   rA  r   r[   s      r@   _shift_left_lowering_rulers  	  r  r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   ShRSIOpr   rr  s      r@   %_shift_right_arithmetic_lowering_rulerv  	  ;    	1	4s|	4cm	4	4	4$!Q	q!			##r?   c                p    t          ||g| j        | j        R  \  }}t          j        ||          j        S rQ   )rF  rm   rn   r   ShRUIOpr   rr  s      r@   #_shift_right_logical_lowering_rulesrz  	  rw  r?   c                   d}g d}g d}t          j        | |  z             }|dk     }t          j        ||dz
  t          j        |          dz
            }t          j        ||d         |d                   }t	          d|          D ],}t          j        |||         ||                   }|||z  z   }-t          j        t          j        |           d	k    t           j        | z  || z            S )
N	   )	gG,^>gf	>gjfj;goD XkҾg4JY`[,?gA4! Tg8Eqg")`?gܴ?)	gn<_>*g1gav?g@V?g]rng_Ow?gߺ8g0ֿT?g9a?gA@g      @g      @g      @r   r   r  )r   log1pr  sqrtr  absr  )	r   k_degreew_lt_5_constantsw_gt_5_constantsww_lt_5prZ   cs	            r@   _erf_inv_32_helperr  	  s   (  
   yaR!s7&	iC!s!233!	i(+-=a-@AA!H  a	&*1-/?/BCCA	AE	AA	371::$cgk1q5	9	99r?   c                    | j         \  }|j        t          j        k    r  t	          t
          d          | |          S t          rw  )rm   r   r   r   r  r  r   )rA  r   r  s      r@   _erf_inv_lowering_ruler  	  s@    l)6\S[  @9'%@@@aHHH
r?   c               d    ~| j         \  }t          j        t          |          |          j        S rQ   )rn   r)   	BitcastOpr   r   )rA  r   tyr  s       r@   _bitcast_lowering_ruler  	  s+    +8	x00!	4	4	;;r?   c                   | j         \  }| j        \  }|j        j        |j        k    rt	          d          t          j        t          |          |          j        S )Nz!Changing bitwidths not supported.)	rm   rn   r   r  r   r)   r  r   r   )rA  r   r#  in_avalr  s        r@   #_bitcast_convert_type_lowering_ruler  	  sS    +7+8]y111
A
B
BB	x00!	4	4	;;r?   c                   t          | t          j                  rt          j                            d| j         d          }t          j        | j	        t          j                  rK| j        t          j        k    sJ t          | t          j                  }t          j        |          j        S t          j                            | j        t+          | j	        d          |          }t-          j        |g g           j        S t          | t          j                  r4t          | t          j                  }t          j        |          j        S t3          dt5          |            d          )Nrt   ru   r~   Tr   r   r   )r   rw   r   r   r|   r}   ro   r   r   r   r   r   r   r   r   r)   AllocaSemaphoreOpr   r   r   r   r   r   AllocaOpr   r   r   )r   r   memref_typer  s       r@   _alloc_valuer  
  sA   g/00 5|!!"Kt7H"K"K"KLLH
~dj(":;; 	6.":::::#D~7OPPPk";//66""
*
DJ4
@
@
@ # ! !h _Xr2..55$233 5!$^5MNNNK --44<tDzz<<<===r?   c               V   d | j         D             }t          j        |          }d |j        D             }| j                                        5  t          j        |          }d d d            n# 1 swxY w Y   t          j	        |j
                  5  t          t          |          }t          d |D                       }| j                            g | j        |R           } t!          | |g||R  }t          j        |           d d d            n# 1 swxY w Y   |j        S )Nc                ,    g | ]}t          |          S r>   )r   r   s     r@   r   z-_run_scoped_lowering_rule.<locals>.<listcomp>
  s     >>>od##>>>r?   c                    g | ]	}|j         
S r>   r   r  s     r@   r   z-_run_scoped_lowering_rule.<locals>.<listcomp>
  s    +++af+++r?   c              3  \   K   | ]'}t          |t          j                  r|j        nd V  (d S rQ   )r   r   r   r   rr  s     r@   r\   z,_run_scoped_lowering_rule.<locals>.<genexpr>
  sN       , , %/q%2C$D$DN$ , , , , , ,r?   r  )rn   r)   RegionOpr  rl   rd   r  r  r   r  rl  r:  r  r]   rf   rJ   r  r  r  )	rA  r   r  r  regionr  r  rJ   r  s	            r@   _run_scoped_lowering_ruler  
  s   >>>>>(<!!&++el+++(
--// . .&u--E. . . . . . . . . . . . . . .	%%  |X&&D , ,"*, , , , ,L


&
&7s'7,77 '  C U
3V
3d
3
3
3CK               
s%   A11A58A5A8DD Ddevice_id_typetpu_primitives.DeviceIdTypec           	        |t           j        j        u rt          j        |          }| j        j        j        fd}t          | j        t          j
        gt          |          z  t          j
        gdt          |          z            } t          |d          |g|R  S |t           j        j        u r|S t          d|           )Nc                 R    t          d t          |           D                       S )Nc              3  &   K   | ]\  }}||z  V  d S rQ   r>   )rY   r  r   s      r@   r\   zI_device_id_to_logical.<locals>._linearize_mesh_indices.<locals>.<genexpr>0
  s*      >>41aQ>>>>>>r?   )sumr_   )r   r8   s    r@   _linearize_mesh_indicesz6_device_id_to_logical.<locals>._linearize_mesh_indices/
  s*    >>3w#=#=>>>>>>r?   rQ   )rl   rm   rn   rJ   Fr  zUnsupported device id type: )tpu_primitivesDeviceIdTypeMESHr	   r#  rl   rL   r8   rk   rw   r   rR   r  LOGICALr   )rA  	device_idr  r  r  r  r8   s         @r@   _device_id_to_logicalr  (
  s     ~2777&y11J'4AL? ? ? ? ?#--.Z@./s:.	  I F9,uEEE        4<<<K>KKLLLr?   c                   t          j        || j                  \  }}t          j        ||          \  }}t          |||j        |          \  }}t          j        |          j        S rQ   )r	   r!  rm   ri  r   r)   SemaphoreReadOpr   )rA  r   r  sem_avalr   semr'  s          r@   _semaphore_read_lowering_ruler  >
  s_    
 (CLAA+(A*9d;;-#xc8X^X>>&#q		S	!	!	((r?   c                  t          j        || j                  \  }}}}}t          j        ||          \  }}}}	}
t          |||j        |          \  }}|	t          | |	|          }	t          j        |||	|
          j        S )N)r  core_id)	r	   r!  rm   ri  r   r  r)   SemaphoreSignalOpr  )rA  r   r  r  r  r   r  r'  valuer  
core_indexs              r@   _semaphore_signal_lowering_ruler  K
  s     #1)S\JJ(Aq!Q090HTX0Y0Y-#x	:c8X^X>>&#q%c9nEEI			5Iz
 
 
r?   c                   t          j        || j                  \  }}}t          j        ||          \  }}}t          |||j        |          \  }}t          j        ||          j        S rQ   )r	   r!  rm   ri  r   r)   SemaphoreWaitOpr  )rA  r   r  r  r   r  r'  r  s           r@   _semaphore_wait_lowering_ruler  _
  sd    +Is|DD.(Aq"1)TBB#xc8X^X>>&#q		S%	(	(	00r?   c          	     F   t          j        ||          \	  }}}}}}	}
}}t          j        || j                  \	  }}}}}}}}}|j        t          j        k    rt          d          t          j        || j                  }|d         |d         }}t          ||||          \  }}|
t          |
||j	        |          \  }
}t          ||||          \  }}t          |||j	        |	          \  }}|t          | ||          }t          j        ||||
|          j        S )Nz(DMAs with bool dtypes are not supported.r   rK  )source_semaphorer  )r	   r!  rm   r   r   r   r   rJ   ri  r   r  r)   EnqueueDMAOpr  )rA  r&  r  r  src_refsrc_indexersdst_refdst_indexersr  sem_indexerssrc_semsrc_sem_indexersr  src_ref_avalr   dst_ref_avalr  src_sem_avalrJ   src_ref_block_shapedst_ref_block_shapes                        r@   _dma_start_lowering_ruler  f
  sg    tT**
	 tS\22 F<L!Xq,1 39$$
H
I
II)$0@AA,-9!_l1o*|0, *'1 |13CE EJGQ|0, *'1 c8X^\BB&#q%c9nEEI		'7C'$-
/ 
/ 
//67r?   c               L   ~t          j        ||          \  }}}}t          j        || j                  \  }}	}
}	t          j        || j                  }|d         }t	          ||
||          \  }}	t	          |||j        |          \  }}	t          j        ||          j        S r  )	r	   r!  rm   rJ   ri  r   r)   	WaitDMAOpr  )rA  r&  r  r  r  r  r%  r'  r  r   r  rJ   rA  s                r@   _dma_wait_lowering_ruler  
  s    %.%=dD%I%I"#|S(&5dCLII(Ax)$0@AA, O/	8_h &#q c8X^\BB&#q	sC	 	 	((r?   c                2    t          j                    j        S rQ   )r)   
DeviceIdOpr   )rA  s    r@   _device_id_lowering_ruler  
  s    				  r?   	axis_namer   c                  | j         j        }|r(||v r$t          | |                    |                    S t	          j                    j        }| j         j        }|t          d          |j	        }|j
        }|                    |          }t          ||                   }t          t          j        ||dz   d          t          j                            }	t          j        t          j        ||	          |          S )Nr  zMesh context is not set.r   rw  )rl   rG   rY  rH  r)   r  r   rL   r   r6   r7   r   r   prodr+   r   remsidivsi)
rA  r  rG   r  rL   r6   r7   
axis_index	axis_sizeminor_divisors
             r@   _axis_index_ruler  
  s    #.* LI++$Sz/?/?	/J/JKKKKn%)%2,
/
0
00&*&*	****Z011)gja))*"(;;; - 
U[M::I	F	FFr?   c                h    t          | j        d                   }t          j        |          j        S r/  )r   rn   r)   GetBarrierSemaphoreOpr   )rA  r  s     r@   _get_barrier_semaphore_ruler  
  s*    a 011+		";	/	/	66r?   nanosc                4    t          j        |          j        S rQ   )r)   DelayOpr  )rA  r  s     r@   _delay_ruler  
  s    	U			##r?   fmthas_placeholdersc               R   t          j        |g|R   |ryt          d |D                       st          d          d                    d t          t          j                                        |                    D                       }n|}t          j
        |||           dS )Nc              3  t   K   | ]3}t          |j        t          j                  o|j        j        d k    V  4dS )    N)r   r   r   r   r   )rY   r  s     r@   r\   z$_debug_print_rule.<locals>.<genexpr>
  sR         	38R^,,E21E     r?   zAll arguments must be 32-bit integers when using placeholders (`{...}`). If you need to print values of other types, remove placeholders from the format string. c              3  4   K   | ]\  }\  }}}}| d | V  dS )$Nr>   )rY   r  textr   s       r@   r\   z$_debug_print_rule.<locals>.<genexpr>
  sL         C$1a #     r?   )	formattedr>   )r!   check_debug_print_formatr]  	TypeErrorjoinr^   string	Formatterr}   r)   r  )rA  r  r  r  tpu_fmts        r@   _debug_print_ruler  
  s     %c1D1111         9   gg  $-f.>.@.@.F.Fs.K.K$L$L    GG
 G'$#34444	r?   c                L   ~ t          |          dk    r?t          |d         t                    r$t          j        |d         j                  j        S t          d |D                       }|sd |D             }t          d|           t          j        |          j        S )Nr   r   c              3  T   K   | ]#}t          |j        t          j                  V  $d S rQ   )r   r   r   r   rY   seeds     r@   r\   z+_prng_seed_lowering_rule.<locals>.<genexpr>
  s0      MMtZ	2>::MMMMMMr?   c                    g | ]	}|j         
S r>   r  r  s     r@   r   z,_prng_seed_lowering_rule.<locals>.<listcomp>
  s    ...$)...r?   z+All seed data must be scalar integers. Got )	rR   r   r  r)   PRNGSeed32Oprl  r  r]  r   )rA  seedsall_integers
seed_typess       r@   _prng_seed_lowering_ruler  
  s    	 	ZZ1__E!Ho>>_E!H,--55 MMuMMMMM,	 Q.....J
O:OO
P
PP		%	 	 	((r?   c                   t          |          dk    rt          d          | j        d         }t          |          }t	          j        |          j        S )Nr   z*random_bits only supports rank>=2 outputs.r   )rR   r   rn   r   r)   PRNGRandomBitsOpr   )rA  r   r  r  s       r@   _prng_random_bits_lowering_ruler  
  sM    ZZ1__
J
K
KK]1(X&&(		h	'	'	..r?   c               F    t          |j        d          } || |          S rw  )r  r  )rA  r  implseed_lowerings       r@   random_seed_loweringr  
  s/    
i%) ) )-	sE	"	""r?   c                   |dk    s
J d            | j         \  }|j        j        }t          |j        d          } || |||          S )Nr  zOnly 32-bit PRNG supported.Fr  )	bit_widthr   )rm   r   r   r  random_bits)rA  keysr  r   r   r  bits_lowerings          r@   random_bits_loweringr  
  s^    	b7
,%$		$
0 0 0-	sDIU	C	C	CCr?   c                t    | j         \  }}|j        j        }t          |j        d          } || ||          S rw  )rm   r   r   r  fold_in)rA  r  msgs	keys_avalr   r  fold_in_lowerings          r@   random_fold_in_loweringr    sH    ,)Q		$
lU, , ,		#tT	*	**r?   c                j   ~ t          |t                    sJ t          |j                  dk    rt	          d|j                   |j        d         }t          j                            |j        t          t          j        d                              }t          j        ||          j        }|S )Nr  r  r   r+   )r   r  r]   r   r   rl  r   r   r   r   r   r   r   r  r   )rA  keyscalarr  r  s        r@   random_unwrap_loweringr    s    		C	)	))))
3=V##
m    ;q>&]	m&sy'9'9:: ( 	8V,,3#	*r?   c          	        ~ ~t          |j        t          j                  rg }|j        j        }t          |          dk    rt          d          t          |          dk    rt          d|           t          |d                   D ]-}|	                    t          j        |g d|g                     .t          |t          |                    S t          |t                    r|S t          dt          |                     )	NrK  r  r  r  r   r   r  zkey_data wrap )r   r   r   r   r   rR   r   r]   r  rm  r   r0  r  )rA  key_datar  key_data_listkey_data_shaperZ   s         r@   random_wrap_loweringr     s-   	4r}-- @M](N
>a ;<<<^&&	!	! 	!" " " >!$%% C C6+Hb1a&AABBBB~)>)>@ @ @ @/** AO
?tH~~??
@
@@r?   c                   ~~~~~~t          |t          j                  st          d          t	          |j                  dk    rt          d          t          |j                                                  d         \  }
} fd}t	          j                  dk    sJ  t          j	        | t          j
        t          j        j                  gt	                     z  t          j
        t          j        j                  gt	                     z  d t          t	                               D             |
|fft          t          d	          
                    |	 }|dfS )NzMesh must be a PallasMeshr   zMesh must be 1Dr   c                 \    | d t                             }t          j        dg|R   d S r/  )rR   r`   
eval_jaxpr)r  in_refsr  r   s     r@   rl  z'_shard_map_discharge_rule.<locals>.bodyQ  s8    >CMM>"Gr,G,,,,,,r?   r~   c                    i | ]}||S r>   r>   )rY   rZ   s     r@   
<dictcomp>z-_shard_map_discharge_rule.<locals>.<dictcomp>\  s    ???QAq???r?   )r   )r   )mosaic)r  in_specs	out_specsinput_output_aliasesr   compiler_paramsr>   )r   rw   
PallasMeshr   rR   r   rr  itemsrH  r    	BlockSpecr   r   re  r  ru  )r  	out_avalsr   autoin_names	out_namesr   	check_reprewriter  core_axis_name	num_coresrl  r  s   `     `       r@   _shard_map_discharge_ruler+  ?  s    xIw	D',	-	- ;
9
:
::__q
/
0
00"4:#3#3#5#566q9.)- - - - - - 
U]		q	 	 	 	 	
!x/F/JKKKLH"0G0KLLLMH??%H*>*>???Y')-888  	 	 	 
	# 
b.r?   )	shard_map)ro   rp   rq   rr   )F)r   r   r   rN   rq   r   )NNF)ro   rp   r   rN   rQ   )r   r   )NF)rl   r@  rA  rC   r   r   r   r   r   rB  r   r   rO   rN   rq   rC  )rA  rC   r   r   r   r  r#  r5  rT  r   rO   rN   rq   r  )rA  rC   r   r   rT  r   r#  r5  rO   rN   rq   r  )r  r   r  rN   rq   r   )r  r?  r  r?  rq   r  )rA  rB   r   r   r  r  rq   r  r  )r  r7  r5  rN   rq   r8  )r@  r  rA  rB  r5  rN   rq   rC  )
r%  r  r  rO  r@  r  rA  rB  rq   rP  )rA  rk   rq   r  )r  r  rq   r  )r  r  rq   r  )rA  rk   r   r   r;  r  r  r  r  rN   r  r   )
rA  rk   r   r   r;  r   r  r   r  rN   )rA  rk   r   r   r  r  r  r   r  rN   r  r  r  r   r  r   r  rN   )
rA  rk   rN  r   rO  r   r  r   rP  rN   )rA  rk   r  r   )r   r  rq   r  )rA  rk   r  r  )rA  rk   r  r   )rA  rk   r  r   )rA  rk   r  r5  r  rN   (T  rm  
__future__r   collections.abcr   r   rh   re   r<  r  typingr   r   r  r   r	   jax._srcr
   r   r`   r   r   r   r   r  r   mesh_libr   r   r   r   jax._src.interpretersr   r   r  jax._src.laxr  jax._src.lax.control_flowr   jax._src.libr   rg  jax._src.lib.mlirr   jax._src.lib.mlir.dialectsr   r   r   r   r   r   jax._src.pallasr    rw   r!   r"   r  jax._src.pallas.mosaicr   r  jax._src.stater#   r  r$   state_primitivesjax._src.utilr%   r&   r'   r(    jax.experimental.mosaic.dialectsr)   	jax.numpynumpyr   jaxlib.mlir.irr*   r   r  r   rx   rv   rz   r   r   r_  r"  r:  
unsafe_mapr_   r  r  	dataclassr4   rB   rk   r   r   r   r   r  r(  r  r   r   r   r  r  rp  rk  r  r  r  r  r  r  r  r+  get_paddr/  swap_pr3  r6  r?  rN  rg  ri  r  r$  rz  load_pr{  r  r-  r  multiple_of_pr  r  CombiningKindMAXIMUMFr  MAXSIr  MAXUIREDUCE_MAX_KINDSr   r  r+   r  REDUCE_MAX_IDENTITYr  _reduce_max_lowering_rulereduce_max_pMINIMUMFMINSIMINUIREDUCE_MIN_KINDSREDUCE_MIN_IDENTITY_reduce_min_lowering_rulereduce_min_pADDREDUCE_SUM_KINDSREDUCE_SUM_IDENTITYr  _reduce_sum_lowering_rulereduce_sum_pr  reduce_and_pr  reduce_or_pr  broadcast_in_dim_pr  dot_general_pr  r(  convert_element_type_pr-  	reshape_pr1  	squeeze_pr7  concatenate_pr:  iota_pr>  transpose_prF  rL  add_p	add_any_prR  max_prX  min_pr\  sub_pr_  mul_prd  div_prj  rem_prn  abs_prq  neg_pru  ry  sign_pr|  rsqrt_pr  sqrt_pr  exp_pr  pow_pr  integer_pow_pr  exp2_pr  
logistic_pr  sin_pr  tanh_pr  log_pr  log1p_pr  round_peq_pr  lt_ple_pgt_pge_pr  r  r  r  and_pr  or_pr  not_pr  
select_n_pr  r  clamp_pr  for_pr  r  r  scan_pr  r>  while_prF  cond_prM  pjit_prR  custom_jvp_call_prU  debug_callback_prY  program_id_pr^  num_programs_prb  repeat_prh  roll_prm  slice_pro  xor_prs  shift_left_prv  shift_right_arithmetic_prz  shift_right_logical_pr  r  	erf_inv_pr  	bitcast_pr  bitcast_convert_type_pr  r  run_scoped_pr  r  semaphore_read_pr  semaphore_signal_pr  semaphore_wait_pr  dma_start_pr  
dma_wait_pr  device_id_pr  axis_index_pr  get_barrier_semaphore_pr  delay_pr  debug_print_pr  prng_seed_pr  prng_random_bits_pr  random_seed_pr  random_bits_pr  random_fold_in_pr  random_unwrap_pr  random_wrap_pr+  jax.experimentalr,  register_discharge_ruleshard_map_pr>   r?   r@   <module>r     s   B A " " " " " " . . . . . . . .                              



                   % % % % % % ' ' ' ' ' '             & & & & & & % % % % % %             % % % % % %       & & & & & & 4 4 4 4 4 4 , , , , , , . . . . . . 2 2 2 2 2 2             , , , , , , + + + + + + + + + + + + - - - - - - * * * * * * - - - - - - ' ' ' ' ' ' + + + + + + & & & & & & 1 1 1 1 1 1 3 3 3 3 3 3 ? ? ? ? ? ? 7 7 7 7 7 7 # # # # # # 9 9 9 9 9 9 " " " " " " " " " " " " $ $ $ $ $ $             0 0 0 0 0 0       ! ! ! ! ! !    
 	(!N2##28G$$  $

CZCZ BHWxrx''BHX))BHX))BHX))	                         @                B B B B 27    , 7;/4(" (" (" (" ("V% % % %4       , E"""P P P P P P P #"Pd                 "&"G1 G1 G1 G1 G1T= = = =@; ; ; ;|   $    	   P P P P.a a a aH
 
 
& & &*C C C C  *<% &   *0 1 1 1J J J J* +>& '   *1 2 2 29 9 9A A A& & & &0' ' ' 'T       F   d###       $# 77 77 77 77rM M M M@ %8z  !   *+ , , ,I I I I81 1 1 1J J J JZ %?z  !   *+ , , ,    ,Fz' (* * *\ L&&/v+1-3  L%%--xrx))-  10G24 4 #<s   L&&/v+1-3  L%%,,xrx))-  10G24 4 #<s   L&&*v+/-1  L#q  10G24 4 #<s  
+ 
+ 
+ 
+ $=s  
+ 
+ 
+ 
+ #;s "2 "2 "2 "2J *Is% &U U U Up %?s  !M M M673 73 73 73t .Qs) *I I I I !7s} I I I I !7s}     %?s  !: : : :
 1sz = = = = #;s   :, , , , /sy    #) $ $ $$6w  !   '+ , , ,	, 	, 	, 	, /sy    #) $ $ $	, 	, 	, 	, /sy    #) $ $ $, , , , /sy    #) $ $ $, , , , /sy    #) $ $ $	, 	, 	, 	, /sy    #) $ $ $	, 	, 	, 	, /sy    #) $ $ $, , , , /sy I I I I /sy    #) $ $ $  J J J J 1sz         3s{     1sz     /sy " " " " /sy    #) $ $ $   
 %?s  !    1sz    #* % % %
) 
) 
) 
) "9s~     /sy     1sz     /sy         3s{ Q Q Q Q 3s{  HaHaHaHaHaHa  HaHaHaHaHaHa ; ; ; ; -9,-?JJsx ,9,-?JJsx ,9,-?JJsx ,9,-?JJsx ,9,-?JJsx ,9,-?JJsx # # # #
 /sy    #) $ $ $" " " "
 -sx    #( # # #+ + + +, /sy + + + +& "9s~   
K K K K
 3s{    > "4x~ + + + +\   *+ + + +X 1sz    #* % % %   6K0 K0 K0 K0\ 3s{    : 1sz = = = =
 2t{ B B B B" # !3 4    .Ky) *6 6 6 6 +Dz& '        -Hz( )J J J J
 +@~& '    )<~$ %    3s{ # # # #
 /sy    #) $ $ $# # # #
 $=s     #* + + +$ $ $ $
 0Us+ ,   #6 7 7 7$ $ $ $
 -Ps( )   #3 4 4 4: : :4    !7s} < < < <
 ,B~' (< < < < .Qs) *> > > >&   $ +Dz& 'M M M M,) ) ) ) 3P~. /   " $ ~0 11 1 1 1
 3P~. /!7 !7 !7 !7D .F~) *) ) ) ) -D~( )! ! ! !-E~) *G G G G$ $4s  7 7 7 7 :U~5 6$ $ $ $ *5~% &   4 ,=z' () ) ) ) .F~) */ / / / 5T~0 1# # # &:t! "D D D &:t! "+ + + )@t$ %   (>t# $A A A* &:t! "# # #L ' & & & & & > ' '	(= > >    r?   